intel/compute-runtime
1
0
Fork 0
mirror of https://github.com/intel/compute-runtime.git synced 2025-12-23 11:03:02 +08:00
Code Issues Packages Projects Releases Wiki Activity

Fix Source Level Debugger scenario

Browse Source 
- 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
This commit is contained in:
Hoppe, Mateusz
2018-12-05 10:58:08 +01:00
committed by sys_ocldev
parent a39660de92
commit 3d35bf4291
9 changed files with 142 additions and 61 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
runtime/compiler_interface/compiler_interface.cpp
View File

@@ -260,6 +260,10 @@ cl_int CompilerInterface::link(
program.storeGenBinary(currOut->GetOutput()->GetMemory<char>(), currOut->GetOutput()->GetSizeRaw());
program.updateBuildLog(&device, currOut->GetBuildLog()->GetMemory<char>(), currOut->GetBuildLog()->GetSizeRaw());
if (currOut->GetDebugData()->GetSizeRaw() != 0) {
program.storeDebugData(currOut->GetDebugData()->GetMemory<char>(), currOut->GetDebugData()->GetSizeRaw());
}
}
return CL_SUCCESS;

36
runtime/program/build.cpp
View File

@@ -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<std::string, BuiltinDispatchInfoBuilder *> &builtinsMap) {
auto ret = this->build(1, &device, buildOptions, nullptr, nullptr, enableCaching);

16
runtime/program/compile.cpp
View File

@@ -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;
}
}

7
runtime/program/link.cpp
View File

@@ -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 <cstring>
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;

3
runtime/program/program.h
View File

@@ -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

92
unit_tests/compiler_interface/compiler_interface_tests.cpp
View File

@@ -24,6 +24,7 @@
#include "unit_tests/mocks/mock_program.h"
#include "gmock/gmock.h"
#include <memory>
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>(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);

7
unit_tests/mocks/mock_program.h
View File

@@ -10,6 +10,7 @@
#include "runtime/helpers/options.h"
#include "runtime/helpers/string.h"
#include "runtime/program/program.h"
#include "gmock/gmock.h"
#include <string>
@@ -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

16
unit_tests/program/process_debug_data_tests.cpp
View File

@@ -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<MockProgram>(*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));
}

22
unit_tests/program/program_with_kernel_debug_tests.cpp
View File

@@ -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<GMockProgram>(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);