From 3d35bf4291bb7c22152bd39ec2f92c254f1b5dcc Mon Sep 17 00:00:00 2001 From: "Hoppe, Mateusz" Date: Wed, 5 Dec 2018 10:58:08 +0100 Subject: [PATCH] Fix Source Level Debugger scenario - when Program is compiled and linked, kernel debug options must be added when linking - when linking by CompilerInterface, store debugData in Program Change-Id: Ie93a8fa7586681b94307a30c109c103f78ec861a --- .../compiler_interface/compiler_interface.cpp | 4 + runtime/program/build.cpp | 36 +++++--- runtime/program/compile.cpp | 16 +--- runtime/program/link.cpp | 7 ++ runtime/program/program.h | 3 + .../compiler_interface_tests.cpp | 92 +++++++++++-------- unit_tests/mocks/mock_program.h | 7 ++ .../program/process_debug_data_tests.cpp | 16 ++++ .../program_with_kernel_debug_tests.cpp | 22 ++++- 9 files changed, 142 insertions(+), 61 deletions(-) diff --git a/runtime/compiler_interface/compiler_interface.cpp b/runtime/compiler_interface/compiler_interface.cpp index 6c58fa2c8a..d87fc07837 100644 --- a/runtime/compiler_interface/compiler_interface.cpp +++ b/runtime/compiler_interface/compiler_interface.cpp @@ -260,6 +260,10 @@ cl_int CompilerInterface::link( program.storeGenBinary(currOut->GetOutput()->GetMemory(), currOut->GetOutput()->GetSizeRaw()); program.updateBuildLog(&device, currOut->GetBuildLog()->GetMemory(), currOut->GetBuildLog()->GetSizeRaw()); + + if (currOut->GetDebugData()->GetSizeRaw() != 0) { + program.storeDebugData(currOut->GetDebugData()->GetMemory(), currOut->GetDebugData()->GetSizeRaw()); + } } return CL_SUCCESS; diff --git a/runtime/program/build.cpp b/runtime/program/build.cpp index b790b6f714..b930f89997 100644 --- a/runtime/program/build.cpp +++ b/runtime/program/build.cpp @@ -65,24 +65,19 @@ cl_int Program::build( } TranslationArgs inputArgs = {}; + if (strcmp(sourceCode.c_str(), "") == 0) { retVal = CL_INVALID_PROGRAM; break; } if (isKernelDebugEnabled()) { - internalOptions.append(CompilerOptions::debugKernelEnable); - options.append(" -g "); - if (pDevice->getSourceLevelDebugger()) { - if (pDevice->getSourceLevelDebugger()->isOptimizationDisabled()) { - options.append("-cl-opt-disable "); - } - std::string filename; - pDevice->getSourceLevelDebugger()->notifySourceCode(sourceCode.c_str(), sourceCode.size(), filename); - if (!filename.empty()) { - // Add "-s" flag first so it will be ignored by clang in case the options already have this flag set. - options = std::string("-s ") + filename + " " + options; - } + std::string filename; + appendKernelDebugOptions(); + notifyDebuggerWithSourceCode(filename); + if (!filename.empty()) { + // Add "-s" flag first so it will be ignored by clang in case the options already have this flag set. + options = std::string("-s ") + filename + " " + options; } } @@ -143,6 +138,23 @@ cl_int Program::build( return retVal; } +bool Program::appendKernelDebugOptions() { + internalOptions.append(CompilerOptions::debugKernelEnable); + options.append(" -g "); + if (pDevice->getSourceLevelDebugger()) { + if (pDevice->getSourceLevelDebugger()->isOptimizationDisabled()) { + options.append("-cl-opt-disable "); + } + } + return true; +} + +void Program::notifyDebuggerWithSourceCode(std::string &filename) { + if (pDevice->getSourceLevelDebugger()) { + pDevice->getSourceLevelDebugger()->notifySourceCode(sourceCode.c_str(), sourceCode.size(), filename); + } +} + cl_int Program::build(const cl_device_id device, const char *buildOptions, bool enableCaching, std::unordered_map &builtinsMap) { auto ret = this->build(1, &device, buildOptions, nullptr, nullptr, enableCaching); diff --git a/runtime/program/compile.cpp b/runtime/program/compile.cpp index 4930f5336e..ee1b2d3ad2 100644 --- a/runtime/program/compile.cpp +++ b/runtime/program/compile.cpp @@ -131,17 +131,11 @@ cl_int Program::compile( internalOptions.append(platform()->peekCompilerExtensions()); if (isKernelDebugEnabled()) { - internalOptions.append(CompilerOptions::debugKernelEnable); - options.append(" -g "); - if (pDevice->getSourceLevelDebugger()) { - if (pDevice->getSourceLevelDebugger()->isOptimizationDisabled()) { - options.append("-cl-opt-disable "); - } - std::string filename; - pDevice->getSourceLevelDebugger()->notifySourceCode(sourceCode.c_str(), sourceCode.size(), filename); - if (!filename.empty()) { - options = std::string("-s ") + filename + " " + options; - } + std::string filename; + appendKernelDebugOptions(); + notifyDebuggerWithSourceCode(filename); + if (!filename.empty()) { + options = std::string("-s ") + filename + " " + options; } } diff --git a/runtime/program/link.cpp b/runtime/program/link.cpp index 05630447d1..ba565cc5b0 100644 --- a/runtime/program/link.cpp +++ b/runtime/program/link.cpp @@ -11,6 +11,9 @@ #include "runtime/source_level_debugger/source_level_debugger.h" #include "program.h" #include "elf/writer.h" + +#include "runtime/compiler_interface/compiler_options.h" + #include namespace OCLRT { @@ -59,6 +62,10 @@ cl_int Program::link( options = (buildOptions != nullptr) ? buildOptions : ""; + if (isKernelDebugEnabled()) { + appendKernelDebugOptions(); + } + isCreateLibrary = (strstr(options.c_str(), "-create-library") != nullptr); buildStatus = CL_BUILD_IN_PROGRESS; diff --git a/runtime/program/program.h b/runtime/program/program.h index 0aec88a753..2a5c4b31d6 100644 --- a/runtime/program/program.h +++ b/runtime/program/program.h @@ -267,6 +267,9 @@ class Program : public BaseObject<_cl_program> { void extractInternalOptions(std::string &options); + MOCKABLE_VIRTUAL bool appendKernelDebugOptions(); + void notifyDebuggerWithSourceCode(std::string &filename); + static const std::string clOptNameClVer; static const std::string clOptNameUniformWgs; // clang-format off diff --git a/unit_tests/compiler_interface/compiler_interface_tests.cpp b/unit_tests/compiler_interface/compiler_interface_tests.cpp index 9d74fe0669..413fe82780 100644 --- a/unit_tests/compiler_interface/compiler_interface_tests.cpp +++ b/unit_tests/compiler_interface/compiler_interface_tests.cpp @@ -24,6 +24,7 @@ #include "unit_tests/mocks/mock_program.h" #include "gmock/gmock.h" +#include using namespace OCLRT; @@ -93,43 +94,43 @@ class CompilerInterfaceTest : public DeviceFixture, cl_int retVal = CL_SUCCESS; }; +class MyCompilerInterface : public CompilerInterface { + public: + static MyCompilerInterface *allocate() { + + auto compilerInterface = new MyCompilerInterface(); + if (!compilerInterface->initializePub()) { + delete compilerInterface; + compilerInterface = nullptr; + } + + for (size_t n = 0; n < sizeof(compilerInterface->mockDebugData); n++) { + compilerInterface->mockDebugData[n] = (char)n; + } + + auto vars = OCLRT::getIgcDebugVars(); + vars.debugDataToReturn = compilerInterface->mockDebugData; + vars.debugDataToReturnSize = sizeof(compilerInterface->mockDebugData); + OCLRT::setIgcDebugVars(vars); + + return compilerInterface; + } + + ~MyCompilerInterface() override { + auto vars = OCLRT::getIgcDebugVars(); + vars.debugDataToReturn = nullptr; + vars.debugDataToReturnSize = 0; + OCLRT::setIgcDebugVars(vars); + } + + bool initializePub() { + return initialize(); + } + + char mockDebugData[32]; +}; + TEST_F(CompilerInterfaceTest, BuildWithDebugData) { - class MyCompilerInterface : public CompilerInterface { - public: - static MyCompilerInterface *allocate() { - - auto compilerInterface = new MyCompilerInterface(); - if (!compilerInterface->initializePub()) { - delete compilerInterface; - compilerInterface = nullptr; - } - - for (size_t n = 0; n < sizeof(compilerInterface->mockDebugData); n++) { - compilerInterface->mockDebugData[n] = (char)n; - } - - auto vars = OCLRT::getIgcDebugVars(); - vars.debugDataToReturn = compilerInterface->mockDebugData; - vars.debugDataToReturnSize = sizeof(compilerInterface->mockDebugData); - OCLRT::setIgcDebugVars(vars); - - return compilerInterface; - } - - ~MyCompilerInterface() override { - auto vars = OCLRT::getIgcDebugVars(); - vars.debugDataToReturn = nullptr; - vars.debugDataToReturnSize = 0; - OCLRT::setIgcDebugVars(vars); - } - - bool initializePub() { - return initialize(); - } - - char mockDebugData[32]; - }; - // Build a regular program cl_device_id device = pDevice; char *kernel = (char *)"__kernel void\nCB(\n__global unsigned int* src, __global unsigned int* dst)\n{\nint id = (int)get_global_id(0);\ndst[id] = src[id];\n}\n"; @@ -138,7 +139,7 @@ TEST_F(CompilerInterfaceTest, BuildWithDebugData) { EXPECT_EQ(CL_SUCCESS, retVal); // Inject DebugData during this build - class MyCompilerInterface *cip = MyCompilerInterface::allocate(); + MyCompilerInterface *cip = MyCompilerInterface::allocate(); EXPECT_NE(nullptr, cip); retVal = cip->build(*pProgram, inputArgs, false); EXPECT_EQ(CL_SUCCESS, retVal); @@ -189,6 +190,23 @@ TEST_F(CompilerInterfaceTest, BuildWithDebugData) { delete cip; } +TEST_F(CompilerInterfaceTest, GivenDebugDataAvailableWhenLinkingProgramThenDebugDataIsStoredInProgram) { + cl_device_id device = pDevice; + char *kernel = (char *)"__kernel void\nCB(\n__global unsigned int* src, __global unsigned int* dst)\n{\nint id = (int)get_global_id(0);\ndst[id] = src[id];\n}\n"; + pProgram->setSource(kernel); + retVal = pProgram->compile(1, &device, nullptr, 0, nullptr, nullptr, nullptr, nullptr); + EXPECT_EQ(CL_SUCCESS, retVal); + + // Inject DebugData during this link + auto cip = std::unique_ptr(MyCompilerInterface::allocate()); + EXPECT_NE(nullptr, cip); + retVal = cip->link(*pProgram, inputArgs); + EXPECT_EQ(CL_SUCCESS, retVal); + + EXPECT_EQ(sizeof(cip->mockDebugData), pProgram->getDebugDataSize()); + EXPECT_NE(nullptr, pProgram->getDebugData()); +} + TEST_F(CompilerInterfaceTest, CompileClToIsa) { // build from .cl to gen ISA retVal = pCompilerInterface->build(*pProgram, inputArgs, false); diff --git a/unit_tests/mocks/mock_program.h b/unit_tests/mocks/mock_program.h index 39c41b74ee..1515bc6518 100644 --- a/unit_tests/mocks/mock_program.h +++ b/unit_tests/mocks/mock_program.h @@ -10,6 +10,7 @@ #include "runtime/helpers/options.h" #include "runtime/helpers/string.h" #include "runtime/program/program.h" +#include "gmock/gmock.h" #include @@ -153,4 +154,10 @@ class GlobalMockSipProgram : public Program { static ExecutionEnvironment executionEnvironment; }; +class GMockProgram : public Program { + public: + using Program::Program; + MOCK_METHOD0(appendKernelDebugOptions, bool(void)); +}; + } // namespace OCLRT diff --git a/unit_tests/program/process_debug_data_tests.cpp b/unit_tests/program/process_debug_data_tests.cpp index 833c4a2bae..a30cfc0929 100644 --- a/unit_tests/program/process_debug_data_tests.cpp +++ b/unit_tests/program/process_debug_data_tests.cpp @@ -78,3 +78,19 @@ TEST_F(ProgramTests, GivenProgramWithDebugDataForTwoKernelsWhenPorcessedThenDebu EXPECT_EQ(ptrDiff(vIsa2, debugData.get()), ptrDiff(kernelInfo2->debugData.vIsa, program->getDebugDataBinary(programDebugDataSize))); EXPECT_EQ(ptrDiff(genIsa2, debugData.get()), ptrDiff(kernelInfo2->debugData.genIsa, program->getDebugDataBinary(programDebugDataSize))); } + +TEST_F(ProgramTests, GivenProgramWithoutDebugDataWhenPorcessedThenDebugDataIsNotSetInKernelInfo) { + const char kernelName1[] = "kernel1"; + + auto kernelInfo1 = new KernelInfo(); + kernelInfo1->name = kernelName1; + auto program = std::make_unique(*pDevice->getExecutionEnvironment()); + + program->addKernelInfo(kernelInfo1); + program->processDebugData(); + + size_t programDebugDataSize = 0; + EXPECT_EQ(0u, kernelInfo1->debugData.genIsaSize); + EXPECT_EQ(0u, kernelInfo1->debugData.vIsaSize); + EXPECT_EQ(nullptr, program->getDebugDataBinary(programDebugDataSize)); +} diff --git a/unit_tests/program/program_with_kernel_debug_tests.cpp b/unit_tests/program/program_with_kernel_debug_tests.cpp index 77e680bcfd..03aa1cdcc0 100644 --- a/unit_tests/program/program_with_kernel_debug_tests.cpp +++ b/unit_tests/program/program_with_kernel_debug_tests.cpp @@ -174,6 +174,26 @@ TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsBuilt } } +TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsLinkedThenKernelDebugOptionsAreAppended) { + if (pDevice->getHardwareInfo().pPlatform->eRenderCoreFamily >= IGFX_GEN9_CORE) { + + MockActiveSourceLevelDebugger *sourceLevelDebugger = new MockActiveSourceLevelDebugger; + pDevice->executionEnvironment->sourceLevelDebugger.reset(sourceLevelDebugger); + + cl_int retVal = pProgram->compile(1, &device, nullptr, 0, nullptr, nullptr, nullptr, nullptr); + EXPECT_EQ(CL_SUCCESS, retVal); + + auto program = std::unique_ptr(new GMockProgram(*pContext->getDevice(0)->getExecutionEnvironment(), pContext, false)); + program->enableKernelDebug(); + + EXPECT_CALL(*program, appendKernelDebugOptions()).Times(1); + + cl_program clProgramToLink = pProgram; + retVal = program->link(1, &device, nullptr, 1, &clProgramToLink, nullptr, nullptr); + EXPECT_EQ(CL_SUCCESS, retVal); + } +} + TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsBuiltThenDebuggerIsNotifiedWithKernelDebugData) { if (pDevice->getHardwareInfo().pPlatform->eRenderCoreFamily >= IGFX_GEN9_CORE) { @@ -200,7 +220,7 @@ TEST_F(ProgramWithKernelDebuggingTest, givenEnabledKernelDebugWhenProgramIsLinke ON_CALL(*sourceLevelDebugger, notifySourceCode(::testing::_, ::testing::_, ::testing::_)).WillByDefault(::testing::Return(false)); ON_CALL(*sourceLevelDebugger, isOptimizationDisabled()).WillByDefault(::testing::Return(false)); - EXPECT_CALL(*sourceLevelDebugger, isOptimizationDisabled()).Times(1); + EXPECT_CALL(*sourceLevelDebugger, isOptimizationDisabled()).Times(2); EXPECT_CALL(*sourceLevelDebugger, notifySourceCode(::testing::_, ::testing::_, ::testing::_)).Times(1); EXPECT_CALL(*sourceLevelDebugger, notifyKernelDebugData(::testing::_)).Times(1);