Skip to content

Commit 52b1d92

Browse files
KrystianChmielewskiCompute-Runtime-Automation
KrystianChmielewski
authored and
Compute-Runtime-Automation
committed
Mock debug program instead of using binaries
Removes usage of precompiled binaries in debug program tests. Related-To: NEO-7383 Signed-off-by: Krystian Chmielewski <[email protected]>
1 parent 69bef97 commit 52b1d92

File tree

8 files changed

+292
-172
lines changed

8 files changed

+292
-172
lines changed

‎opencl/test/unit_test/CMakeLists.txt‎

Lines changed: 0 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -426,13 +426,6 @@ macro(macro_for_each_core_type)
426426
endforeach()
427427
endif()
428428

429-
# Disable debug kernel generation on gen8 - debugger not supported on gen8
430-
if(NOT ("${CORE_TYPE_LOWER}" STREQUAL "gen8"))
431-
foreach(REVISION_ID ${${PLATFORM_TYPE}_${CORE_TYPE}_REVISIONS})
432-
neo_gen_kernel_with_kernel_debug_options(${family_name_with_type} ${PLATFORM_LOWER} ${REVISION_ID} ${family_name_with_type} ${TEST_KERNEL})
433-
endforeach()
434-
endif()
435-
436429
# Gen9lp needs extra -m32 flag
437430
if(("${CORE_TYPE_LOWER}" STREQUAL "gen9") AND ("${PLATFORM_TYPE_LOWER}" STREQUAL "lp"))
438431
foreach(REVISION_ID ${${PLATFORM_TYPE}_${CORE_TYPE}_REVISIONS})

‎opencl/test/unit_test/command_queue/enqueue_debug_kernel_tests.cpp‎

Lines changed: 39 additions & 55 deletions
Original file line numberDiff line numberDiff line change
@@ -13,90 +13,74 @@
1313
#include "shared/test/common/helpers/unit_test_helper.h"
1414
#include "shared/test/common/test_macros/hw_test.h"
1515
#include "shared/test/common/test_macros/mock_method_macros.h"
16+
#include "shared/test/common/test_macros/test.h"
1617

1718
#include "opencl/source/command_queue/command_queue.h"
1819
#include "opencl/source/program/program.h"
1920
#include "opencl/test/unit_test/fixtures/enqueue_handler_fixture.h"
2021
#include "opencl/test/unit_test/mocks/mock_buffer.h"
2122
#include "opencl/test/unit_test/mocks/mock_command_queue.h"
23+
#include "opencl/test/unit_test/mocks/mock_debug_program.h"
2224
#include "opencl/test/unit_test/mocks/mock_kernel.h"
2325
#include "opencl/test/unit_test/program/program_from_binary.h"
2426

2527
using namespace NEO;
26-
using namespace ::testing;
2728

2829
typedef EnqueueHandlerTest EnqueueDebugKernelSimpleTest;
2930

30-
class EnqueueDebugKernelTest : public ProgramSimpleFixture,
31-
public ::testing::Test {
31+
class EnqueueDebugKernelFixture {
3232
public:
33-
void SetUp() override {
34-
ProgramSimpleFixture::setUp();
35-
device = pClDevice;
36-
pDevice->executionEnvironment->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->debugger.reset(new SourceLevelDebugger(nullptr));
37-
38-
auto sipType = SipKernel::getSipKernelType(*pDevice);
39-
SipKernel::initSipKernel(sipType, *pDevice);
40-
41-
if (pDevice->getHardwareInfo().platform.eRenderCoreFamily >= IGFX_GEN9_CORE) {
42-
pDevice->deviceInfo.debuggerActive = true;
43-
std::string filename;
44-
std::string kernelOption(CompilerOptions::debugKernelEnable);
45-
KernelFilenameHelper::getKernelFilenameFromInternalOption(kernelOption, filename);
46-
47-
kbHelper = new KernelBinaryHelper(filename, false);
48-
createProgramWithSource(
49-
pContext,
50-
"copybuffer.cl");
51-
pProgram->enableKernelDebug();
52-
53-
cl_int retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
33+
void setUp() {
34+
clDevice = context.getDevice(0);
35+
device = &clDevice->getDevice();
36+
37+
device->getExecutionEnvironment()->rootDeviceEnvironments[device->getRootDeviceIndex()]->debugger.reset(new SourceLevelDebugger(nullptr));
38+
39+
auto sipType = SipKernel::getSipKernelType(*device);
40+
SipKernel::initSipKernel(sipType, *device);
41+
42+
if (device->getHardwareInfo().platform.eRenderCoreFamily >= IGFX_GEN9_CORE) {
43+
const_cast<DeviceInfo &>(device->getDeviceInfo()).debuggerActive = true;
44+
45+
program = std::make_unique<MockDebugProgram>(context.getDevices());
46+
cl_int retVal = program->build(program->getDevices(), nullptr, false);
5447
ASSERT_EQ(CL_SUCCESS, retVal);
5548

56-
// create a kernel
57-
pMultiDeviceKernel = MultiDeviceKernel::create(
58-
pProgram,
59-
pProgram->getKernelInfosForKernel("CopyBuffer"),
49+
multiDeviceKernel = MultiDeviceKernel::create(
50+
static_cast<NEO::Program *>(program.get()),
51+
MockKernel::toKernelInfoContainer(*program->getKernelInfo("kernel", 0), device->getRootDeviceIndex()),
6052
&retVal);
61-
debugKernel = pMultiDeviceKernel->getKernel(rootDeviceIndex);
53+
debugKernel = multiDeviceKernel->getKernel(device->getRootDeviceIndex());
6254

6355
ASSERT_EQ(CL_SUCCESS, retVal);
6456
ASSERT_NE(nullptr, debugKernel);
65-
66-
cl_mem src = &bufferSrc;
67-
cl_mem dst = &bufferDst;
68-
retVal = debugKernel->setArg(
69-
0,
70-
sizeof(cl_mem),
71-
&src);
72-
retVal = debugKernel->setArg(
73-
1,
74-
sizeof(cl_mem),
75-
&dst);
7657
}
7758
}
7859

79-
void TearDown() override {
80-
if (pDevice->getHardwareInfo().platform.eRenderCoreFamily >= IGFX_GEN9_CORE) {
81-
delete kbHelper;
82-
pMultiDeviceKernel->release();
60+
void tearDown() {
61+
if (multiDeviceKernel != nullptr) {
62+
multiDeviceKernel->release();
8363
}
84-
ProgramSimpleFixture::tearDown();
8564
}
86-
cl_device_id device;
65+
66+
std::unique_ptr<char[]> ssh = nullptr;
67+
std::unique_ptr<MockDebugProgram> program = nullptr;
68+
NEO::ClDevice *clDevice = nullptr;
69+
NEO::Device *device = nullptr;
8770
Kernel *debugKernel = nullptr;
88-
MultiDeviceKernel *pMultiDeviceKernel = nullptr;
89-
KernelBinaryHelper *kbHelper = nullptr;
71+
MultiDeviceKernel *multiDeviceKernel = nullptr;
9072
MockContext context;
9173
MockBuffer bufferSrc;
9274
MockBuffer bufferDst;
9375
};
9476

77+
using EnqueueDebugKernelTest = Test<EnqueueDebugKernelFixture>;
78+
9579
HWTEST_F(EnqueueDebugKernelTest, givenDebugKernelWhenEnqueuedThenSSHAndBtiAreCorrectlySet) {
96-
if (pDevice->isDebuggerActive()) {
80+
if (device->isDebuggerActive()) {
9781
using BINDING_TABLE_STATE = typename FamilyType::BINDING_TABLE_STATE;
9882
using RENDER_SURFACE_STATE = typename FamilyType::RENDER_SURFACE_STATE;
99-
std::unique_ptr<MockCommandQueueHw<FamilyType>> mockCmdQ(new MockCommandQueueHw<FamilyType>(&context, pClDevice, 0));
83+
std::unique_ptr<MockCommandQueueHw<FamilyType>> mockCmdQ(new MockCommandQueueHw<FamilyType>(&context, clDevice, 0));
10084

10185
size_t gws[] = {1, 1, 1};
10286
auto &ssh = mockCmdQ->getIndirectHeap(IndirectHeap::Type::SURFACE_STATE, 4096u);
@@ -118,10 +102,10 @@ HWTEST_F(EnqueueDebugKernelTest, givenDebugKernelWhenEnqueuedThenSSHAndBtiAreCor
118102
}
119103

120104
HWTEST_F(EnqueueDebugKernelTest, givenDebugKernelWhenEnqueuedThenSurfaceStateForDebugSurfaceIsSetAtBindlessOffsetZero) {
121-
if (pDevice->isDebuggerActive()) {
105+
if (device->isDebuggerActive()) {
122106
using BINDING_TABLE_STATE = typename FamilyType::BINDING_TABLE_STATE;
123107
using RENDER_SURFACE_STATE = typename FamilyType::RENDER_SURFACE_STATE;
124-
std::unique_ptr<MockCommandQueueHw<FamilyType>> mockCmdQ(new MockCommandQueueHw<FamilyType>(&context, pClDevice, 0));
108+
std::unique_ptr<MockCommandQueueHw<FamilyType>> mockCmdQ(new MockCommandQueueHw<FamilyType>(&context, clDevice, 0));
125109

126110
size_t gws[] = {1, 1, 1};
127111
auto &ssh = mockCmdQ->getIndirectHeap(IndirectHeap::Type::SURFACE_STATE, 4096u);
@@ -219,9 +203,9 @@ HWTEST_F(EnqueueDebugKernelSimpleTest, givenKernelFromProgramWithoutDebugEnabled
219203
using ActiveDebuggerTest = EnqueueDebugKernelTest;
220204

221205
HWTEST_F(ActiveDebuggerTest, givenKernelFromProgramWithoutDebugEnabledAndActiveDebuggerWhenEnqueuedThenDebugSurfaceIsSetup) {
222-
MockProgram program(&context, false, toClDeviceVector(*pClDevice));
223-
std::unique_ptr<MockDebugKernel> kernel(MockKernel::create<MockDebugKernel>(*pDevice, &program));
224-
std::unique_ptr<CommandQueueHw<FamilyType>> cmdQ(new CommandQueueHw<FamilyType>(&context, pClDevice, nullptr, false));
206+
MockProgram program(&context, false, toClDeviceVector(*clDevice));
207+
std::unique_ptr<MockDebugKernel> kernel(MockKernel::create<MockDebugKernel>(*device, &program));
208+
std::unique_ptr<CommandQueueHw<FamilyType>> cmdQ(new CommandQueueHw<FamilyType>(&context, clDevice, nullptr, false));
225209

226210
size_t gws[] = {1, 1, 1};
227211
cmdQ->enqueueKernel(kernel.get(), 1, nullptr, gws, nullptr, 0, nullptr, nullptr);

‎opencl/test/unit_test/mocks/CMakeLists.txt‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -15,6 +15,8 @@ set(IGDRCL_SRCS_tests_mocks
1515
${CMAKE_CURRENT_SOURCE_DIR}/mock_command_queue.h
1616
${CMAKE_CURRENT_SOURCE_DIR}/mock_context.cpp
1717
${CMAKE_CURRENT_SOURCE_DIR}/mock_context.h
18+
${CMAKE_CURRENT_SOURCE_DIR}/mock_debug_program.cpp
19+
${CMAKE_CURRENT_SOURCE_DIR}/mock_debug_program.h
1820
${CMAKE_CURRENT_SOURCE_DIR}/mock_event.h
1921
${CMAKE_CURRENT_SOURCE_DIR}/mock_gmm_resource_info_ocl.cpp
2022
${CMAKE_CURRENT_SOURCE_DIR}/mock_image.h

‎opencl/test/unit_test/mocks/mock_debug_program.cpp‎

Lines changed: 107 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,107 @@
1+
/*
2+
* Copyright (C) 2022 Intel Corporation
3+
*
4+
* SPDX-License-Identifier: MIT
5+
*
6+
*/
7+
8+
#include "opencl/test/unit_test/mocks/mock_debug_program.h"
9+
10+
#include "shared/source/device/device.h"
11+
#include "shared/source/helpers/ptr_math.h"
12+
#include "shared/source/program/kernel_info.h"
13+
#include "shared/test/common/mocks/mock_compiler_interface.h"
14+
15+
#include "opencl/source/cl_device/cl_device.h"
16+
17+
#include "program_debug_data.h"
18+
19+
MockDebugProgram::MockDebugProgram(const NEO::ClDeviceVector &deviceVector) : NEO::Program(nullptr, false, deviceVector) {
20+
createdFrom = CreatedFrom::SOURCE;
21+
sourceCode = "__kernel void kernel(){}";
22+
kernelDebugEnabled = true;
23+
prepareMockCompilerInterface(deviceVector[0]->getDevice());
24+
}
25+
26+
void MockDebugProgram::debugNotify(const NEO::ClDeviceVector &deviceVector, std::unordered_map<uint32_t, BuildPhase> &phasesReached) {
27+
Program::debugNotify(deviceVector, phasesReached);
28+
wasDebuggerNotified = true;
29+
}
30+
31+
void MockDebugProgram::createDebugZebin(uint32_t rootDeviceIndex) {
32+
Program::createDebugZebin(rootDeviceIndex);
33+
wasCreateDebugZebinCalled = true;
34+
}
35+
36+
void MockDebugProgram::addKernelInfo(NEO::KernelInfo *inInfo, uint32_t rootDeviceIndex) {
37+
buildInfos[rootDeviceIndex].kernelInfoArray.push_back(inInfo);
38+
}
39+
40+
void MockDebugProgram::processDebugData(uint32_t rootDeviceIndex) {
41+
Program::processDebugData(rootDeviceIndex);
42+
wasProcessDebugDataCalled = true;
43+
}
44+
45+
cl_int MockDebugProgram::processGenBinary(const NEO::ClDevice &clDevice) {
46+
auto &kernelInfoArray = buildInfos[0].kernelInfoArray;
47+
kernelInfoArray.resize(1);
48+
if (kernelInfo == nullptr) {
49+
prepareKernelInfo();
50+
}
51+
kernelInfoArray[0] = kernelInfo;
52+
return CL_SUCCESS;
53+
}
54+
55+
void MockDebugProgram::prepareKernelInfo() {
56+
kernelInfo = new NEO::KernelInfo;
57+
kernelInfo->kernelDescriptor.kernelMetadata.kernelName = "kernel";
58+
kernelInfo->kernelDescriptor.kernelAttributes.simdSize = 32U;
59+
prepareSSHForDebugSurface();
60+
}
61+
62+
void MockDebugProgram::prepareSSHForDebugSurface() {
63+
kernelInfo->heapInfo.SurfaceStateHeapSize = static_cast<uint32_t>(alignUp(64U + sizeof(int), 64U));
64+
kernelSsh = std::make_unique<char[]>(kernelInfo->heapInfo.SurfaceStateHeapSize);
65+
memset(kernelSsh.get(), 0U, kernelInfo->heapInfo.SurfaceStateHeapSize);
66+
kernelInfo->heapInfo.pSsh = kernelSsh.get();
67+
68+
kernelInfo->kernelDescriptor.payloadMappings.implicitArgs.systemThreadSurfaceAddress.bindful = 0U;
69+
kernelInfo->kernelDescriptor.payloadMappings.bindingTable.numEntries = 1U;
70+
kernelInfo->kernelDescriptor.payloadMappings.bindingTable.tableOffset = 64U;
71+
}
72+
73+
void MockDebugProgram::prepareMockCompilerInterface(NEO::Device &device) {
74+
auto mockCompilerInterface = std::make_unique<NEO::MockCompilerInterfaceCaptureBuildOptions>();
75+
this->compilerInterface = mockCompilerInterface.get();
76+
device.getRootDevice()->getExecutionEnvironment()->rootDeviceEnvironments[0]->compilerInterface = std::move(mockCompilerInterface);
77+
78+
compilerInterface->output.intermediateRepresentation.size = 32;
79+
compilerInterface->output.intermediateRepresentation.mem = std::make_unique<char[]>(32);
80+
81+
compilerInterface->output.deviceBinary.size = 32;
82+
compilerInterface->output.deviceBinary.mem = std::make_unique<char[]>(32);
83+
84+
constexpr char kernelName[] = "kernel";
85+
constexpr size_t isaSize = 8;
86+
constexpr size_t visaSize = 8;
87+
auto &debugData = compilerInterface->output.debugData;
88+
debugData.size = sizeof(iOpenCL::SProgramDebugDataHeaderIGC) + sizeof(iOpenCL::SKernelDebugDataHeaderIGC) + sizeof(kernelName) + isaSize + visaSize;
89+
debugData.mem = std::make_unique<char[]>(debugData.size);
90+
91+
auto programDebugHeader = reinterpret_cast<iOpenCL::SProgramDebugDataHeaderIGC *>(debugData.mem.get());
92+
programDebugHeader->NumberOfKernels = 1;
93+
94+
auto kernelDebugHeader = reinterpret_cast<iOpenCL::SKernelDebugDataHeaderIGC *>(ptrOffset(programDebugHeader, sizeof(iOpenCL::SProgramDebugDataHeaderIGC)));
95+
kernelDebugHeader->KernelNameSize = sizeof(kernelName);
96+
kernelDebugHeader->SizeGenIsaDbgInBytes = isaSize;
97+
kernelDebugHeader->SizeVisaDbgInBytes = visaSize;
98+
99+
auto kernelNameDst = reinterpret_cast<char *>(ptrOffset(kernelDebugHeader, sizeof(iOpenCL::SKernelDebugDataHeader)));
100+
std::memcpy(kernelNameDst, kernelName, sizeof(kernelName));
101+
102+
auto visa = ptrOffset(kernelNameDst, sizeof(kernelName));
103+
std::memset(visa, 0x10, visaSize);
104+
105+
auto isa = ptrOffset(visa, visaSize);
106+
std::memset(isa, 0x20, isaSize);
107+
}

‎opencl/test/unit_test/mocks/mock_debug_program.h‎

Lines changed: 49 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,49 @@
1+
/*
2+
* Copyright (C) 2022 Intel Corporation
3+
*
4+
* SPDX-License-Identifier: MIT
5+
*
6+
*/
7+
8+
#include "opencl/source/program/program.h"
9+
10+
#include <memory>
11+
12+
namespace NEO {
13+
class ClDevice;
14+
class ClDeviceVector;
15+
class Device;
16+
struct KernelInfo;
17+
struct MockCompilerInterfaceCaptureBuildOptions;
18+
} // namespace NEO
19+
20+
class MockDebugProgram : public NEO::Program {
21+
public:
22+
using Base = NEO::Program;
23+
using Base::Base;
24+
using Base::buildInfos;
25+
using Base::irBinary;
26+
using Base::irBinarySize;
27+
using Base::kernelDebugEnabled;
28+
29+
MockDebugProgram(const NEO::ClDeviceVector &deviceVector);
30+
31+
void debugNotify(const NEO::ClDeviceVector &deviceVector, std::unordered_map<uint32_t, BuildPhase> &phasesReached) override;
32+
void createDebugZebin(uint32_t rootDeviceIndex) override;
33+
void processDebugData(uint32_t rootDeviceIndex) override;
34+
cl_int processGenBinary(const NEO::ClDevice &clDevice) override;
35+
36+
void addKernelInfo(NEO::KernelInfo *inInfo, uint32_t rootDeviceIndex);
37+
38+
NEO::KernelInfo *kernelInfo = nullptr;
39+
std::unique_ptr<char[]> kernelSsh;
40+
NEO::MockCompilerInterfaceCaptureBuildOptions *compilerInterface;
41+
bool wasDebuggerNotified = false;
42+
bool wasCreateDebugZebinCalled = false;
43+
bool wasProcessDebugDataCalled = false;
44+
45+
protected:
46+
void prepareKernelInfo();
47+
void prepareSSHForDebugSurface();
48+
void prepareMockCompilerInterface(NEO::Device &device);
49+
};

‎opencl/test/unit_test/mocks/mock_program.h‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -210,6 +210,7 @@ class MockProgramAppendKernelDebugOptions : public Program {
210210
public:
211211
using Program::Program;
212212
ADDMETHOD_NOBASE(appendKernelDebugOptions, bool, true, (ClDevice & clDevice, std::string &internalOptions));
213+
ADDMETHOD_NOBASE(processGenBinary, cl_int, CL_SUCCESS, (const ClDevice &clDevice));
213214
};
214215

215216
} // namespace NEO

0 commit comments

Comments
 (0)