Media/compute-runtime
1
0
Fork 0
mirror of https://github.com/intel/compute-runtime.git synced 2025-09-15 13:01:45 +08:00
Code Issues Packages Projects Releases Wiki Activity

Mock debug program instead of using binaries

Browse Source 
Removes usage of precompiled binaries in debug program tests.

Related-To: NEO-7383

Signed-off-by: Krystian Chmielewski <krystian.chmielewski@intel.com>
This commit is contained in:
Krystian Chmielewski
2022-11-02 13:36:46 +00:00
committed by Compute-Runtime-Automation
parent 69bef975f0
commit 52b1d92193
8 changed files with 288 additions and 168 deletions
Download Patch File Download Diff File Expand all files Collapse all files

7
opencl/test/unit_test/CMakeLists.txt
View File

@ -426,13 +426,6 @@ macro(macro_for_each_core_type)
endforeach()
endif()
# Disable debug kernel generation on gen8 - debugger not supported on gen8
if(NOT ("${CORE_TYPE_LOWER}" STREQUAL "gen8"))
foreach(REVISION_ID ${${PLATFORM_TYPE}_${CORE_TYPE}_REVISIONS})
neo_gen_kernel_with_kernel_debug_options(${family_name_with_type} ${PLATFORM_LOWER} ${REVISION_ID} ${family_name_with_type} ${TEST_KERNEL})
endforeach()
endif()
# Gen9lp needs extra -m32 flag
if(("${CORE_TYPE_LOWER}" STREQUAL "gen9") AND ("${PLATFORM_TYPE_LOWER}" STREQUAL "lp"))
foreach(REVISION_ID ${${PLATFORM_TYPE}_${CORE_TYPE}_REVISIONS})

86
opencl/test/unit_test/command_queue/enqueue_debug_kernel_tests.cpp
View File

@ -13,90 +13,74 @@
#include "shared/test/common/helpers/unit_test_helper.h"
#include "shared/test/common/test_macros/hw_test.h"
#include "shared/test/common/test_macros/mock_method_macros.h"
#include "shared/test/common/test_macros/test.h"
#include "opencl/source/command_queue/command_queue.h"
#include "opencl/source/program/program.h"
#include "opencl/test/unit_test/fixtures/enqueue_handler_fixture.h"
#include "opencl/test/unit_test/mocks/mock_buffer.h"
#include "opencl/test/unit_test/mocks/mock_command_queue.h"
#include "opencl/test/unit_test/mocks/mock_debug_program.h"
#include "opencl/test/unit_test/mocks/mock_kernel.h"
#include "opencl/test/unit_test/program/program_from_binary.h"
using namespace NEO;
using namespace ::testing;
typedef EnqueueHandlerTest EnqueueDebugKernelSimpleTest;
class EnqueueDebugKernelTest : public ProgramSimpleFixture,
public ::testing::Test {
class EnqueueDebugKernelFixture {
public:
void SetUp() override {
ProgramSimpleFixture::setUp();
device = pClDevice;
pDevice->executionEnvironment->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->debugger.reset(new SourceLevelDebugger(nullptr));
void setUp() {
clDevice = context.getDevice(0);
device = &clDevice->getDevice();
auto sipType = SipKernel::getSipKernelType(*pDevice);
SipKernel::initSipKernel(sipType, *pDevice);
device->getExecutionEnvironment()->rootDeviceEnvironments[device->getRootDeviceIndex()]->debugger.reset(new SourceLevelDebugger(nullptr));
if (pDevice->getHardwareInfo().platform.eRenderCoreFamily >= IGFX_GEN9_CORE) {
pDevice->deviceInfo.debuggerActive = true;
std::string filename;
std::string kernelOption(CompilerOptions::debugKernelEnable);
KernelFilenameHelper::getKernelFilenameFromInternalOption(kernelOption, filename);
auto sipType = SipKernel::getSipKernelType(*device);
SipKernel::initSipKernel(sipType, *device);
kbHelper = new KernelBinaryHelper(filename, false);
createProgramWithSource(
pContext,
"copybuffer.cl");
pProgram->enableKernelDebug();
if (device->getHardwareInfo().platform.eRenderCoreFamily >= IGFX_GEN9_CORE) {
const_cast<DeviceInfo &>(device->getDeviceInfo()).debuggerActive = true;
cl_int retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
program = std::make_unique<MockDebugProgram>(context.getDevices());
cl_int retVal = program->build(program->getDevices(), nullptr, false);
ASSERT_EQ(CL_SUCCESS, retVal);
// create a kernel
pMultiDeviceKernel = MultiDeviceKernel::create(
pProgram,
pProgram->getKernelInfosForKernel("CopyBuffer"),
multiDeviceKernel = MultiDeviceKernel::create(
static_cast<NEO::Program *>(program.get()),
MockKernel::toKernelInfoContainer(*program->getKernelInfo("kernel", 0), device->getRootDeviceIndex()),
&retVal);
debugKernel = pMultiDeviceKernel->getKernel(rootDeviceIndex);
debugKernel = multiDeviceKernel->getKernel(device->getRootDeviceIndex());
ASSERT_EQ(CL_SUCCESS, retVal);
ASSERT_NE(nullptr, debugKernel);
cl_mem src = &bufferSrc;
cl_mem dst = &bufferDst;
retVal = debugKernel->setArg(
0,
sizeof(cl_mem),
&src);
retVal = debugKernel->setArg(
1,
sizeof(cl_mem),
&dst);
}
}
void TearDown() override {
if (pDevice->getHardwareInfo().platform.eRenderCoreFamily >= IGFX_GEN9_CORE) {
delete kbHelper;
pMultiDeviceKernel->release();
void tearDown() {
if (multiDeviceKernel != nullptr) {
multiDeviceKernel->release();
}
ProgramSimpleFixture::tearDown();
}
cl_device_id device;
std::unique_ptr<char[]> ssh = nullptr;
std::unique_ptr<MockDebugProgram> program = nullptr;
NEO::ClDevice *clDevice = nullptr;
NEO::Device *device = nullptr;
Kernel *debugKernel = nullptr;
MultiDeviceKernel *pMultiDeviceKernel = nullptr;
KernelBinaryHelper *kbHelper = nullptr;
MultiDeviceKernel *multiDeviceKernel = nullptr;
MockContext context;
MockBuffer bufferSrc;
MockBuffer bufferDst;
};
using EnqueueDebugKernelTest = Test<EnqueueDebugKernelFixture>;
HWTEST_F(EnqueueDebugKernelTest, givenDebugKernelWhenEnqueuedThenSSHAndBtiAreCorrectlySet) {
if (pDevice->isDebuggerActive()) {
if (device->isDebuggerActive()) {
using BINDING_TABLE_STATE = typename FamilyType::BINDING_TABLE_STATE;
using RENDER_SURFACE_STATE = typename FamilyType::RENDER_SURFACE_STATE;
std::unique_ptr<MockCommandQueueHw<FamilyType>> mockCmdQ(new MockCommandQueueHw<FamilyType>(&context, pClDevice, 0));
std::unique_ptr<MockCommandQueueHw<FamilyType>> mockCmdQ(new MockCommandQueueHw<FamilyType>(&context, clDevice, 0));
size_t gws[] = {1, 1, 1};
auto &ssh = mockCmdQ->getIndirectHeap(IndirectHeap::Type::SURFACE_STATE, 4096u);
@ -118,10 +102,10 @@ HWTEST_F(EnqueueDebugKernelTest, givenDebugKernelWhenEnqueuedThenSSHAndBtiAreCor
}
HWTEST_F(EnqueueDebugKernelTest, givenDebugKernelWhenEnqueuedThenSurfaceStateForDebugSurfaceIsSetAtBindlessOffsetZero) {
if (pDevice->isDebuggerActive()) {
if (device->isDebuggerActive()) {
using BINDING_TABLE_STATE = typename FamilyType::BINDING_TABLE_STATE;
using RENDER_SURFACE_STATE = typename FamilyType::RENDER_SURFACE_STATE;
std::unique_ptr<MockCommandQueueHw<FamilyType>> mockCmdQ(new MockCommandQueueHw<FamilyType>(&context, pClDevice, 0));
std::unique_ptr<MockCommandQueueHw<FamilyType>> mockCmdQ(new MockCommandQueueHw<FamilyType>(&context, clDevice, 0));
size_t gws[] = {1, 1, 1};
auto &ssh = mockCmdQ->getIndirectHeap(IndirectHeap::Type::SURFACE_STATE, 4096u);
@ -219,9 +203,9 @@ HWTEST_F(EnqueueDebugKernelSimpleTest, givenKernelFromProgramWithoutDebugEnabled
using ActiveDebuggerTest = EnqueueDebugKernelTest;
HWTEST_F(ActiveDebuggerTest, givenKernelFromProgramWithoutDebugEnabledAndActiveDebuggerWhenEnqueuedThenDebugSurfaceIsSetup) {
MockProgram program(&context, false, toClDeviceVector(*pClDevice));
std::unique_ptr<MockDebugKernel> kernel(MockKernel::create<MockDebugKernel>(*pDevice, &program));
std::unique_ptr<CommandQueueHw<FamilyType>> cmdQ(new CommandQueueHw<FamilyType>(&context, pClDevice, nullptr, false));
MockProgram program(&context, false, toClDeviceVector(*clDevice));
std::unique_ptr<MockDebugKernel> kernel(MockKernel::create<MockDebugKernel>(*device, &program));
std::unique_ptr<CommandQueueHw<FamilyType>> cmdQ(new CommandQueueHw<FamilyType>(&context, clDevice, nullptr, false));
size_t gws[] = {1, 1, 1};
cmdQ->enqueueKernel(kernel.get(), 1, nullptr, gws, nullptr, 0, nullptr, nullptr);

2
opencl/test/unit_test/mocks/CMakeLists.txt
View File

@ -15,6 +15,8 @@ set(IGDRCL_SRCS_tests_mocks
${CMAKE_CURRENT_SOURCE_DIR}/mock_command_queue.h
${CMAKE_CURRENT_SOURCE_DIR}/mock_context.cpp
${CMAKE_CURRENT_SOURCE_DIR}/mock_context.h
${CMAKE_CURRENT_SOURCE_DIR}/mock_debug_program.cpp
${CMAKE_CURRENT_SOURCE_DIR}/mock_debug_program.h
${CMAKE_CURRENT_SOURCE_DIR}/mock_event.h
${CMAKE_CURRENT_SOURCE_DIR}/mock_gmm_resource_info_ocl.cpp
${CMAKE_CURRENT_SOURCE_DIR}/mock_image.h

107
opencl/test/unit_test/mocks/mock_debug_program.cpp Normal file
View File

@ -0,0 +1,107 @@
/*
* Copyright (C) 2022 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
*/
#include "opencl/test/unit_test/mocks/mock_debug_program.h"
#include "shared/source/device/device.h"
#include "shared/source/helpers/ptr_math.h"
#include "shared/source/program/kernel_info.h"
#include "shared/test/common/mocks/mock_compiler_interface.h"
#include "opencl/source/cl_device/cl_device.h"
#include "program_debug_data.h"
MockDebugProgram::MockDebugProgram(const NEO::ClDeviceVector &deviceVector) : NEO::Program(nullptr, false, deviceVector) {
createdFrom = CreatedFrom::SOURCE;
sourceCode = "__kernel void kernel(){}";
kernelDebugEnabled = true;
prepareMockCompilerInterface(deviceVector[0]->getDevice());
}
void MockDebugProgram::debugNotify(const NEO::ClDeviceVector &deviceVector, std::unordered_map<uint32_t, BuildPhase> &phasesReached) {
Program::debugNotify(deviceVector, phasesReached);
wasDebuggerNotified = true;
}
void MockDebugProgram::createDebugZebin(uint32_t rootDeviceIndex) {
Program::createDebugZebin(rootDeviceIndex);
wasCreateDebugZebinCalled = true;
}
void MockDebugProgram::addKernelInfo(NEO::KernelInfo *inInfo, uint32_t rootDeviceIndex) {
buildInfos[rootDeviceIndex].kernelInfoArray.push_back(inInfo);
}
void MockDebugProgram::processDebugData(uint32_t rootDeviceIndex) {
Program::processDebugData(rootDeviceIndex);
wasProcessDebugDataCalled = true;
}
cl_int MockDebugProgram::processGenBinary(const NEO::ClDevice &clDevice) {
auto &kernelInfoArray = buildInfos[0].kernelInfoArray;
kernelInfoArray.resize(1);
if (kernelInfo == nullptr) {
prepareKernelInfo();
}
kernelInfoArray[0] = kernelInfo;
return CL_SUCCESS;
}
void MockDebugProgram::prepareKernelInfo() {
kernelInfo = new NEO::KernelInfo;
kernelInfo->kernelDescriptor.kernelMetadata.kernelName = "kernel";
kernelInfo->kernelDescriptor.kernelAttributes.simdSize = 32U;
prepareSSHForDebugSurface();
}
void MockDebugProgram::prepareSSHForDebugSurface() {
kernelInfo->heapInfo.SurfaceStateHeapSize = static_cast<uint32_t>(alignUp(64U + sizeof(int), 64U));
kernelSsh = std::make_unique<char[]>(kernelInfo->heapInfo.SurfaceStateHeapSize);
memset(kernelSsh.get(), 0U, kernelInfo->heapInfo.SurfaceStateHeapSize);
kernelInfo->heapInfo.pSsh = kernelSsh.get();
kernelInfo->kernelDescriptor.payloadMappings.implicitArgs.systemThreadSurfaceAddress.bindful = 0U;
kernelInfo->kernelDescriptor.payloadMappings.bindingTable.numEntries = 1U;
kernelInfo->kernelDescriptor.payloadMappings.bindingTable.tableOffset = 64U;
}
void MockDebugProgram::prepareMockCompilerInterface(NEO::Device &device) {
auto mockCompilerInterface = std::make_unique<NEO::MockCompilerInterfaceCaptureBuildOptions>();
this->compilerInterface = mockCompilerInterface.get();
device.getRootDevice()->getExecutionEnvironment()->rootDeviceEnvironments[0]->compilerInterface = std::move(mockCompilerInterface);
compilerInterface->output.intermediateRepresentation.size = 32;
compilerInterface->output.intermediateRepresentation.mem = std::make_unique<char[]>(32);
compilerInterface->output.deviceBinary.size = 32;
compilerInterface->output.deviceBinary.mem = std::make_unique<char[]>(32);
constexpr char kernelName[] = "kernel";
constexpr size_t isaSize = 8;
constexpr size_t visaSize = 8;
auto &debugData = compilerInterface->output.debugData;
debugData.size = sizeof(iOpenCL::SProgramDebugDataHeaderIGC) + sizeof(iOpenCL::SKernelDebugDataHeaderIGC) + sizeof(kernelName) + isaSize + visaSize;
debugData.mem = std::make_unique<char[]>(debugData.size);
auto programDebugHeader = reinterpret_cast<iOpenCL::SProgramDebugDataHeaderIGC *>(debugData.mem.get());
programDebugHeader->NumberOfKernels = 1;
auto kernelDebugHeader = reinterpret_cast<iOpenCL::SKernelDebugDataHeaderIGC *>(ptrOffset(programDebugHeader, sizeof(iOpenCL::SProgramDebugDataHeaderIGC)));
kernelDebugHeader->KernelNameSize = sizeof(kernelName);
kernelDebugHeader->SizeGenIsaDbgInBytes = isaSize;
kernelDebugHeader->SizeVisaDbgInBytes = visaSize;
auto kernelNameDst = reinterpret_cast<char *>(ptrOffset(kernelDebugHeader, sizeof(iOpenCL::SKernelDebugDataHeader)));
std::memcpy(kernelNameDst, kernelName, sizeof(kernelName));
auto visa = ptrOffset(kernelNameDst, sizeof(kernelName));
std::memset(visa, 0x10, visaSize);
auto isa = ptrOffset(visa, visaSize);
std::memset(isa, 0x20, isaSize);
}

49
opencl/test/unit_test/mocks/mock_debug_program.h Normal file
View File

@ -0,0 +1,49 @@
/*
* Copyright (C) 2022 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
*/
#include "opencl/source/program/program.h"
#include <memory>
namespace NEO {
class ClDevice;
class ClDeviceVector;
class Device;
struct KernelInfo;
struct MockCompilerInterfaceCaptureBuildOptions;
} // namespace NEO
class MockDebugProgram : public NEO::Program {
public:
using Base = NEO::Program;
using Base::Base;
using Base::buildInfos;
using Base::irBinary;
using Base::irBinarySize;
using Base::kernelDebugEnabled;
MockDebugProgram(const NEO::ClDeviceVector &deviceVector);
void debugNotify(const NEO::ClDeviceVector &deviceVector, std::unordered_map<uint32_t, BuildPhase> &phasesReached) override;
void createDebugZebin(uint32_t rootDeviceIndex) override;
void processDebugData(uint32_t rootDeviceIndex) override;
cl_int processGenBinary(const NEO::ClDevice &clDevice) override;
void addKernelInfo(NEO::KernelInfo *inInfo, uint32_t rootDeviceIndex);
NEO::KernelInfo *kernelInfo = nullptr;
std::unique_ptr<char[]> kernelSsh;
NEO::MockCompilerInterfaceCaptureBuildOptions *compilerInterface;
bool wasDebuggerNotified = false;
bool wasCreateDebugZebinCalled = false;
bool wasProcessDebugDataCalled = false;
protected:
void prepareKernelInfo();
void prepareSSHForDebugSurface();
void prepareMockCompilerInterface(NEO::Device &device);
};

1
opencl/test/unit_test/mocks/mock_program.h
View File

@ -210,6 +210,7 @@ class MockProgramAppendKernelDebugOptions : public Program {
public:
using Program::Program;
ADDMETHOD_NOBASE(appendKernelDebugOptions, bool, true, (ClDevice & clDevice, std::string &internalOptions));
ADDMETHOD_NOBASE(processGenBinary, cl_int, CL_SUCCESS, (const ClDevice &clDevice));
};
} // namespace NEO

170
opencl/test/unit_test/program/program_with_kernel_debug_tests.cpp
View File

@ -14,18 +14,19 @@
#include "shared/test/common/helpers/kernel_binary_helper.h"
#include "shared/test/common/helpers/kernel_filename_helper.h"
#include "shared/test/common/libult/global_environment.h"
#include "shared/test/common/mocks/mock_compiler_interface.h"
#include "shared/test/common/mocks/mock_modules_zebin.h"
#include "shared/test/common/mocks/mock_source_level_debugger.h"
#include "shared/test/common/test_macros/hw_test.h"
#include "shared/test/common/test_macros/test.h"
#include "opencl/test/unit_test/fixtures/program_fixture.h"
#include "opencl/test/unit_test/mocks/mock_debug_program.h"
#include "opencl/test/unit_test/mocks/mock_program.h"
#include "opencl/test/unit_test/program/program_from_binary.h"
#include "opencl/test/unit_test/program/program_tests.h"
#include "gtest/gtest.h"
#include "program_debug_data.h"
#include <algorithm>
#include <memory>
@ -95,10 +96,9 @@ HWTEST_F(ZebinFallbackToPatchtokensLegacyDebugger, WhenCreatingProgramFromNonBui
EXPECT_EQ(CL_INVALID_BINARY, retVal);
}
class ProgramWithKernelDebuggingTest : public ProgramFixture,
public ::testing::Test {
class ProgramWithKernelDebuggingFixture {
public:
void SetUp() override {
void setUp() {
pDevice = static_cast<MockDevice *>(&mockContext.getDevice(0)->getDevice());
if (!pDevice->getHardwareInfo().capabilityTable.debuggerSupported) {
@ -109,42 +109,31 @@ class ProgramWithKernelDebuggingTest : public ProgramFixture,
std::string kernelOption(CompilerOptions::debugKernelEnable);
KernelFilenameHelper::getKernelFilenameFromInternalOption(kernelOption, filename);
kbHelper = std::make_unique<KernelBinaryHelper>(filename, false);
createProgramWithSource(
&mockContext,
"copybuffer.cl");
pProgram->enableKernelDebug();
program = std::make_unique<MockDebugProgram>(mockContext.getDevices());
}
void TearDown() override {
ProgramFixture::tearDown();
}
std::unique_ptr<KernelBinaryHelper> kbHelper;
void tearDown() {}
std::unique_ptr<MockDebugProgram> program = nullptr;
MockUnrestrictiveContext mockContext;
MockDevice *pDevice = nullptr;
};
using ProgramWithKernelDebuggingTest = Test<ProgramWithKernelDebuggingFixture>;
TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsCompiledThenInternalOptionsIncludeDebugFlag) {
std::string receivedInternalOptions;
auto debugVars = NEO::getFclDebugVars();
debugVars.receivedInternalOptionsOutput = &receivedInternalOptions;
gEnvironment->fclPushDebugVars(debugVars);
cl_int retVal = pProgram->compile(pProgram->getDevices(), nullptr,
cl_int retVal = program->compile(program->getDevices(), nullptr,
0, nullptr, nullptr);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_TRUE(CompilerOptions::contains(receivedInternalOptions, CompilerOptions::debugKernelEnable)) << receivedInternalOptions;
gEnvironment->fclPopDebugVars();
EXPECT_TRUE(CompilerOptions::contains(program->compilerInterface->buildInternalOptions, CompilerOptions::debugKernelEnable));
}
TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsCompiledThenInternalOptionsIncludeDashGFlag) {
cl_int retVal = pProgram->compile(pProgram->getDevices(), nullptr,
cl_int retVal = program->compile(program->getDevices(), nullptr,
0, nullptr, nullptr);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_TRUE(hasSubstr(pProgram->getOptions(), "-g"));
EXPECT_TRUE(hasSubstr(program->getOptions(), "-g"));
}
TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugAndOptDisabledWhenProgramIsCompiledThenOptionsIncludeClOptDisableFlag) {
@ -152,10 +141,10 @@ TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugAndOptDisabledWhen
sourceLevelDebugger->isOptDisabled = true;
pDevice->executionEnvironment->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->debugger.reset(sourceLevelDebugger);
cl_int retVal = pProgram->compile(pProgram->getDevices(), nullptr,
cl_int retVal = program->compile(program->getDevices(), nullptr,
0, nullptr, nullptr);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_TRUE(hasSubstr(pProgram->getOptions(), CompilerOptions::optDisable.data()));
EXPECT_TRUE(hasSubstr(program->getOptions(), CompilerOptions::optDisable.data()));
}
TEST_F(ProgramWithKernelDebuggingTest, GivenDebugVarDebuggerOptDisableZeroWhenOptDisableIsTrueFromDebuggerThenOptDisableIsNotAdded) {
@ -166,10 +155,10 @@ TEST_F(ProgramWithKernelDebuggingTest, GivenDebugVarDebuggerOptDisableZeroWhenOp
sourceLevelDebugger->isOptDisabled = true;
pDevice->executionEnvironment->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->debugger.reset(sourceLevelDebugger);
cl_int retVal = pProgram->compile(pProgram->getDevices(), nullptr,
cl_int retVal = program->compile(program->getDevices(), nullptr,
0, nullptr, nullptr);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_FALSE(hasSubstr(pProgram->getOptions(), CompilerOptions::optDisable.data()));
EXPECT_FALSE(hasSubstr(program->getOptions(), CompilerOptions::optDisable.data()));
}
TEST_F(ProgramWithKernelDebuggingTest, GivenDebugVarDebuggerOptDisableOneWhenOptDisableIsFalseFromDebuggerThenOptDisableIsAdded) {
@ -180,10 +169,10 @@ TEST_F(ProgramWithKernelDebuggingTest, GivenDebugVarDebuggerOptDisableOneWhenOpt
sourceLevelDebugger->isOptDisabled = false;
pDevice->executionEnvironment->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->debugger.reset(sourceLevelDebugger);
cl_int retVal = pProgram->compile(pProgram->getDevices(), nullptr,
cl_int retVal = program->compile(program->getDevices(), nullptr,
0, nullptr, nullptr);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_TRUE(hasSubstr(pProgram->getOptions(), CompilerOptions::optDisable.data()));
EXPECT_TRUE(hasSubstr(program->getOptions(), CompilerOptions::optDisable.data()));
}
TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsCompiledThenOptionsStartsWithDashSFilename) {
@ -191,10 +180,10 @@ TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsCompi
sourceLevelDebugger->sourceCodeFilename = "debugFileName";
pDevice->executionEnvironment->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->debugger.reset(sourceLevelDebugger);
cl_int retVal = pProgram->compile(pProgram->getDevices(), nullptr,
cl_int retVal = program->compile(program->getDevices(), nullptr,
0, nullptr, nullptr);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_TRUE(startsWith(pProgram->getOptions(), "-s \"debugFileName\""));
EXPECT_TRUE(startsWith(program->getOptions(), "-s \"debugFileName\""));
}
TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsCompiledWithCmCOptionThenDashSFilenameIsNotPrepended) {
@ -203,31 +192,23 @@ TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsCompi
pDevice->executionEnvironment->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->debugger.reset(sourceLevelDebugger);
char options[] = "-cmc -cl-opt-disable";
cl_int retVal = pProgram->compile(pProgram->getDevices(), options,
cl_int retVal = program->compile(program->getDevices(), options,
0, nullptr, nullptr);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_FALSE(startsWith(pProgram->getOptions(), "-s debugFileName"));
EXPECT_TRUE(hasSubstr(pProgram->getOptions(), CompilerOptions::optDisable.data()));
EXPECT_FALSE(startsWith(program->getOptions(), "-s debugFileName"));
EXPECT_TRUE(hasSubstr(program->getOptions(), CompilerOptions::optDisable.data()));
}
TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsBuiltThenInternalOptionsIncludeDebugFlag) {
std::string receivedInternalOptions;
auto debugVars = NEO::getFclDebugVars();
debugVars.receivedInternalOptionsOutput = &receivedInternalOptions;
gEnvironment->fclPushDebugVars(debugVars);
cl_int retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
cl_int retVal = program->build(program->getDevices(), nullptr, false);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_TRUE(CompilerOptions::contains(receivedInternalOptions, CompilerOptions::debugKernelEnable)) << receivedInternalOptions;
gEnvironment->fclPopDebugVars();
EXPECT_TRUE(CompilerOptions::contains(program->compilerInterface->buildInternalOptions, CompilerOptions::debugKernelEnable));
}
TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsBuiltThenOptionsIncludeDashGFlag) {
cl_int retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
cl_int retVal = program->build(program->getDevices(), nullptr, false);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_TRUE(hasSubstr(pProgram->getOptions(), "-g"));
EXPECT_TRUE(hasSubstr(program->getOptions(), "-g"));
}
TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugAndOptDisabledWhenProgramIsBuiltThenOptionsIncludeClOptDisableFlag) {
@ -235,9 +216,9 @@ TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugAndOptDisabledWhen
sourceLevelDebugger->isOptDisabled = true;
pDevice->executionEnvironment->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->debugger.reset(sourceLevelDebugger);
cl_int retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
cl_int retVal = program->build(program->getDevices(), nullptr, false);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_TRUE(hasSubstr(pProgram->getOptions(), CompilerOptions::optDisable.data()));
EXPECT_TRUE(hasSubstr(program->getOptions(), CompilerOptions::optDisable.data()));
}
TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsBuiltThenOptionsStartsWithDashSFilename) {
@ -245,9 +226,9 @@ TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsBuilt
sourceLevelDebugger->sourceCodeFilename = "debugFileName";
pDevice->executionEnvironment->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->debugger.reset(sourceLevelDebugger);
cl_int retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
cl_int retVal = program->build(program->getDevices(), nullptr, false);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_TRUE(startsWith(pProgram->getOptions(), "-s \"debugFileName\""));
EXPECT_TRUE(startsWith(program->getOptions(), "-s \"debugFileName\""));
}
TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsBuiltWithCmCOptionThenDashSFilenameIsNotPrepended) {
@ -256,25 +237,28 @@ TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsBuilt
pDevice->executionEnvironment->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->debugger.reset(sourceLevelDebugger);
char options[] = "-cmc -cl-opt-disable";
cl_int retVal = pProgram->build(pProgram->getDevices(), options, false);
cl_int retVal = program->build(program->getDevices(), options, false);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_FALSE(startsWith(pProgram->getOptions(), "-s debugFileName"));
EXPECT_FALSE(startsWith(program->getOptions(), "-s debugFileName"));
}
TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsLinkedThenKernelDebugOptionsAreAppended) {
MockActiveSourceLevelDebugger *sourceLevelDebugger = new MockActiveSourceLevelDebugger;
pDevice->executionEnvironment->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->debugger.reset(sourceLevelDebugger);
cl_int retVal = pProgram->compile(pProgram->getDevices(), nullptr, 0, nullptr, nullptr);
program->compilerInterface->output.debugData.size = 0;
program->compilerInterface->output.debugData.mem = nullptr;
cl_int retVal = program->compile(program->getDevices(), nullptr, 0, nullptr, nullptr);
EXPECT_EQ(CL_SUCCESS, retVal);
auto program = std::unique_ptr<MockProgramAppendKernelDebugOptions>(new MockProgramAppendKernelDebugOptions(&mockContext, false, mockContext.getDevices()));
program->enableKernelDebug();
cl_program clProgramToLink = program.get();
auto &devices = program->getDevices();
cl_program clProgramToLink = pProgram;
retVal = program->link(pProgram->getDevices(), nullptr, 1, &clProgramToLink);
auto newProgram = std::unique_ptr<MockProgramAppendKernelDebugOptions>(new MockProgramAppendKernelDebugOptions(&mockContext, false, mockContext.getDevices()));
newProgram->enableKernelDebug();
retVal = newProgram->link(devices, nullptr, 1, &clProgramToLink);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_EQ(static_cast<unsigned int>(mockContext.getRootDeviceIndices().size()), program->appendKernelDebugOptionsCalled);
EXPECT_EQ(static_cast<unsigned int>(mockContext.getRootDeviceIndices().size()), newProgram->appendKernelDebugOptionsCalled);
}
TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsBuiltThenDebuggerIsNotifiedWithKernelDebugData) {
@ -289,7 +273,7 @@ TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsBuilt
i++;
}
cl_int retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
cl_int retVal = program->build(program->getDevices(), nullptr, false);
EXPECT_EQ(CL_SUCCESS, retVal);
for (auto &el : sourceLevelDebugger) {
@ -315,13 +299,13 @@ TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsLinke
i++;
}
cl_int retVal = pProgram->compile(pProgram->getDevices(), nullptr,
cl_int retVal = program->compile(program->getDevices(), nullptr,
0, nullptr, nullptr);
EXPECT_EQ(CL_SUCCESS, retVal);
cl_program program = pProgram;
retVal = pProgram->link(pProgram->getDevices(), nullptr,
1, &program);
cl_program clprogram = program.get();
retVal = program->link(program->getDevices(), nullptr,
1, &clprogram);
EXPECT_EQ(CL_SUCCESS, retVal);
for (auto &el : sourceLevelDebugger) {
@ -335,27 +319,13 @@ TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsLinke
}
}
TEST_F(ProgramWithKernelDebuggingTest, givenProgramWithKernelDebugEnabledWhenBuiltThenPatchTokenAllocateSipSurfaceHasSizeGreaterThanZero) {
auto &devBinary = pProgram->buildInfos[pDevice->getRootDeviceIndex()].packedDeviceBinary;
auto devBinarySize = pProgram->buildInfos[pDevice->getRootDeviceIndex()].packedDeviceBinarySize;
if (NEO::isDeviceBinaryFormat<NEO::DeviceBinaryFormat::Zebin>(ArrayRef<const uint8_t>::fromAny(devBinary.get(), devBinarySize))) {
GTEST_SKIP();
}
auto retVal = pProgram->build(pProgram->getDevices(), CompilerOptions::debugKernelEnable.data(), false);
EXPECT_EQ(CL_SUCCESS, retVal);
auto kernelInfo = pProgram->getKernelInfo("CopyBuffer", pDevice->getRootDeviceIndex());
EXPECT_NE(0u, kernelInfo->kernelDescriptor.kernelAttributes.perThreadSystemThreadSurfaceSize);
}
TEST_F(ProgramWithKernelDebuggingTest, givenGtpinInitializedWhenCreatingProgramFromBinaryThenDebugDataIsAvailable) {
bool gtpinInitializedBackup = NEO::isGTPinInitialized;
NEO::isGTPinInitialized = true;
auto retVal = pProgram->build(pProgram->getDevices(), CompilerOptions::debugKernelEnable.data(), false);
auto retVal = program->build(program->getDevices(), CompilerOptions::debugKernelEnable.data(), false);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_TRUE(pProgram->wasDebuggerNotified);
EXPECT_TRUE(program->wasDebuggerNotified);
NEO::isGTPinInitialized = gtpinInitializedBackup;
}
@ -363,22 +333,22 @@ TEST_F(ProgramWithKernelDebuggingTest, givenGtpinInitializedWhenCreatingProgramF
TEST_F(ProgramWithKernelDebuggingTest, givenGtpinNotInitializedWhenCreatingProgramFromBinaryThenDebugDataINullptr) {
bool gtpinInitializedBackup = NEO::isGTPinInitialized;
NEO::isGTPinInitialized = false;
pProgram->kernelDebugEnabled = false;
auto retVal = pProgram->build(pProgram->getDevices(), CompilerOptions::debugKernelEnable.data(), false);
program->kernelDebugEnabled = false;
auto retVal = program->build(program->getDevices(), CompilerOptions::debugKernelEnable.data(), false);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_FALSE(pProgram->wasDebuggerNotified);
EXPECT_FALSE(program->wasDebuggerNotified);
NEO::isGTPinInitialized = gtpinInitializedBackup;
}
TEST_F(ProgramWithKernelDebuggingTest, givenKernelDebugEnabledWhenProgramIsBuiltThenDebugDataIsStored) {
auto retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
auto retVal = program->build(program->getDevices(), nullptr, false);
EXPECT_EQ(CL_SUCCESS, retVal);
auto debugData = pProgram->getDebugData(pDevice->getRootDeviceIndex());
auto debugData = program->getDebugData(pDevice->getRootDeviceIndex());
EXPECT_NE(nullptr, debugData);
EXPECT_NE(0u, pProgram->getDebugDataSize(pDevice->getRootDeviceIndex()));
EXPECT_NE(0u, program->getDebugDataSize(pDevice->getRootDeviceIndex()));
}
TEST_F(ProgramWithKernelDebuggingTest, givenProgramWithKernelDebugEnabledWhenProcessDebugDataIsCalledThenKernelInfosAreFilledWithDebugData) {
@ -396,7 +366,7 @@ TEST_F(ProgramWithKernelDebuggingTest, givenProgramWithKernelDebugEnabledWhenPro
KernelInfo *mockKernelInfo = new KernelInfo{};
mockKernelInfo->kernelDescriptor.kernelMetadata.kernelName = "CopyBuffer";
pProgram->addKernelInfo(mockKernelInfo, pDevice->getRootDeviceIndex());
program->addKernelInfo(mockKernelInfo, pDevice->getRootDeviceIndex());
constexpr size_t mockDebugDataSize = sizeof(iOpenCL::SProgramDebugDataHeaderIGC) + sizeof(PatchTokenBinary::KernelFromPatchtokens) + sizeof(mockKernelName) + mockKernelDebugDataSize;
@ -410,10 +380,10 @@ TEST_F(ProgramWithKernelDebuggingTest, givenProgramWithKernelDebugEnabledWhenPro
memcpy_s(dataPtr, mockDebugDataSize, &mockKernelName, sizeof(mockKernelName));
dataPtr = ptrOffset(dataPtr, sizeof(mockKernelName));
memcpy_s(dataPtr, mockDebugDataSize, mockKerneDebugData, mockKernelDebugDataSize);
pProgram->buildInfos[pDevice->getRootDeviceIndex()].debugData.reset(mockDebugData);
program->buildInfos[pDevice->getRootDeviceIndex()].debugData.reset(mockDebugData);
pProgram->processDebugData(pDevice->getRootDeviceIndex());
auto receivedKernelInfo = pProgram->getKernelInfo("CopyBuffer", pDevice->getRootDeviceIndex());
program->processDebugData(pDevice->getRootDeviceIndex());
auto receivedKernelInfo = program->getKernelInfo("CopyBuffer", pDevice->getRootDeviceIndex());
EXPECT_NE(0u, receivedKernelInfo->debugData.vIsaSize);
EXPECT_NE(nullptr, receivedKernelInfo->debugData.vIsa);
@ -422,27 +392,27 @@ TEST_F(ProgramWithKernelDebuggingTest, givenProgramWithKernelDebugEnabledWhenPro
TEST_F(ProgramWithKernelDebuggingTest, givenProgramWithNonZebinaryFormatAndKernelDebugEnabledWhenProgramIsBuiltThenProcessDebugDataIsCalledAndDebuggerNotified) {
MockSourceLevelDebugger *sourceLevelDebugger = new MockSourceLevelDebugger;
pDevice->executionEnvironment->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->debugger.reset(sourceLevelDebugger);
pProgram->enableKernelDebug();
program->enableKernelDebug();
auto mockElf = std::make_unique<MockElfBinaryPatchtokens<>>(pDevice->getHardwareInfo());
auto mockElfSize = mockElf->storage.size();
auto mockElfData = mockElf->storage.data();
pProgram->buildInfos[pDevice->getRootDeviceIndex()].unpackedDeviceBinarySize = mockElfSize;
pProgram->buildInfos[pDevice->getRootDeviceIndex()].unpackedDeviceBinary.reset(new char[mockElfSize]);
memcpy_s(pProgram->buildInfos[pDevice->getRootDeviceIndex()].unpackedDeviceBinary.get(), pProgram->buildInfos[pDevice->getRootDeviceIndex()].unpackedDeviceBinarySize,
program->buildInfos[pDevice->getRootDeviceIndex()].unpackedDeviceBinarySize = mockElfSize;
program->buildInfos[pDevice->getRootDeviceIndex()].unpackedDeviceBinary.reset(new char[mockElfSize]);
memcpy_s(program->buildInfos[pDevice->getRootDeviceIndex()].unpackedDeviceBinary.get(), program->buildInfos[pDevice->getRootDeviceIndex()].unpackedDeviceBinarySize,
mockElfData, mockElfSize);
KernelInfo *mockKernelInfo = new KernelInfo{};
mockKernelInfo->kernelDescriptor.kernelMetadata.kernelName = "CopyBuffer";
pProgram->addKernelInfo(mockKernelInfo, pDevice->getRootDeviceIndex());
program->addKernelInfo(mockKernelInfo, pDevice->getRootDeviceIndex());
auto counter = 0u;
for (const auto &device : pProgram->getDevices()) {
pProgram->notifyDebuggerWithDebugData(device);
for (const auto &device : program->getDevices()) {
program->notifyDebuggerWithDebugData(device);
EXPECT_FALSE(pProgram->wasCreateDebugZebinCalled);
EXPECT_TRUE(pProgram->wasProcessDebugDataCalled);
EXPECT_FALSE(program->wasCreateDebugZebinCalled);
EXPECT_TRUE(program->wasProcessDebugDataCalled);
EXPECT_EQ(++counter, sourceLevelDebugger->notifyKernelDebugDataCalled);
}
}

18
shared/test/common/mocks/mock_compiler_interface.h
View File

@ -1,5 +1,5 @@
/*
* Copyright (C) 2020-2021 Intel Corporation
* Copyright (C) 2020-2022 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
@ -173,7 +173,7 @@ inline std::map<const Device *, MockCompilerInterface::fclDevCtxUptr> &MockCompi
}
struct MockCompilerInterfaceCaptureBuildOptions : CompilerInterface {
TranslationOutput::ErrorCode compile(const NEO::Device &device, const TranslationInput &input, TranslationOutput &) override {
TranslationOutput::ErrorCode compile(const NEO::Device &device, const TranslationInput &input, TranslationOutput &out) override {
buildOptions.clear();
if ((input.apiOptions.size() > 0) && (input.apiOptions.begin() != nullptr)) {
buildOptions.assign(input.apiOptions.begin(), input.apiOptions.end());
@ -182,6 +182,19 @@ struct MockCompilerInterfaceCaptureBuildOptions : CompilerInterface {
if ((input.internalOptions.size() > 0) && (input.internalOptions.begin() != nullptr)) {
buildInternalOptions.assign(input.internalOptions.begin(), input.internalOptions.end());
}
auto copy = [](TranslationOutput::MemAndSize &dst, TranslationOutput::MemAndSize &src) {
if (src.size > 0) {
dst.size = src.size;
dst.mem.reset(new char[src.size]);
std::memcpy(dst.mem.get(), src.mem.get(), src.size);
}
};
copy(out.debugData, output.debugData);
copy(out.intermediateRepresentation, output.intermediateRepresentation);
copy(out.deviceBinary, output.intermediateRepresentation);
out.intermediateCodeType = output.intermediateCodeType;
return TranslationOutput::ErrorCode::Success;
}
@ -195,6 +208,7 @@ struct MockCompilerInterfaceCaptureBuildOptions : CompilerInterface {
return this->MockCompilerInterfaceCaptureBuildOptions::compile(device, input, output);
}
TranslationOutput output;
std::string buildOptions;
std::string buildInternalOptions;
};