diff --git a/level_zero/core/source/module/module_imp.cpp b/level_zero/core/source/module/module_imp.cpp index 3d93f700e0..07311564b4 100644 --- a/level_zero/core/source/module/module_imp.cpp +++ b/level_zero/core/source/module/module_imp.cpp @@ -475,6 +475,7 @@ bool ModuleImp::linkBinary() { isaSegmentsForPatching[segmentId].segmentSize); } } + DBG_LOG(PrintRelocations, NEO::constructRelocationsDebugMessage(this->symbols)); return true; } diff --git a/opencl/source/program/process_device_binary.cpp b/opencl/source/program/process_device_binary.cpp index cf138d4c28..fec9ec30a4 100644 --- a/opencl/source/program/process_device_binary.cpp +++ b/opencl/source/program/process_device_binary.cpp @@ -117,6 +117,7 @@ cl_int Program::linkBinary() { kernHeapInfo.pKernelHeader->KernelHeapSize); } } + DBG_LOG(PrintRelocations, NEO::constructRelocationsDebugMessage(this->symbols)); return CL_SUCCESS; } diff --git a/opencl/test/unit_test/test_files/igdrcl.config b/opencl/test/unit_test/test_files/igdrcl.config index 8e44b684e2..8a20d8ce01 100644 --- a/opencl/test/unit_test/test_files/igdrcl.config +++ b/opencl/test/unit_test/test_files/igdrcl.config @@ -67,6 +67,7 @@ PrintEMDebugInformation = 0 PrintLWSSizes = 0 PrintDispatchParameters = 0 PrintProgramBinaryProcessingTime = 0 +PrintRelocations = 0 WddmResidencyLogger = 0 PrintDriverDiagnostics = -1 EnableDirectSubmission = -1 diff --git a/shared/source/compiler_interface/linker.cpp b/shared/source/compiler_interface/linker.cpp index 1e9e298002..9520045e81 100644 --- a/shared/source/compiler_interface/linker.cpp +++ b/shared/source/compiler_interface/linker.cpp @@ -290,4 +290,19 @@ std::string constructLinkerErrorMessage(const Linker::UnresolvedExternals &unres return errorStream.str(); } +std::string constructRelocationsDebugMessage(const Linker::RelocatedSymbolsMap &relocatedSymbols) { + if (relocatedSymbols.empty()) { + return ""; + } + std::stringstream stream; + stream << "Relocations debug informations :\n"; + for (const auto &symbol : relocatedSymbols) { + stream << " * \"" << symbol.first << "\" [" << symbol.second.symbol.size << " bytes]"; + stream << " " << asString(symbol.second.symbol.segment) << "_SEGMENT@" << symbol.second.symbol.offset; + stream << " -> " << std::hex << std::showbase << symbol.second.gpuAddress << " GPUVA" << std::dec; + stream << "\n"; + } + return stream.str(); +} + } // namespace NEO diff --git a/shared/source/compiler_interface/linker.h b/shared/source/compiler_interface/linker.h index 131beec7ba..8059ed63a5 100644 --- a/shared/source/compiler_interface/linker.h +++ b/shared/source/compiler_interface/linker.h @@ -188,5 +188,6 @@ struct Linker { }; std::string constructLinkerErrorMessage(const Linker::UnresolvedExternals &unresolvedExternals, const std::vector &instructionsSegmentsNames); +std::string constructRelocationsDebugMessage(const Linker::RelocatedSymbolsMap &relocatedSymbols); } // namespace NEO diff --git a/shared/source/debug_settings/debug_variables_base.inl b/shared/source/debug_settings/debug_variables_base.inl index 670589b227..23c1eb8940 100644 --- a/shared/source/debug_settings/debug_variables_base.inl +++ b/shared/source/debug_settings/debug_variables_base.inl @@ -79,6 +79,7 @@ DECLARE_DEBUG_VARIABLE(bool, PrintEMDebugInformation, false, "prints execution m DECLARE_DEBUG_VARIABLE(bool, PrintLWSSizes, false, "prints driver choosen local workgroup sizes") DECLARE_DEBUG_VARIABLE(bool, PrintDispatchParameters, false, "prints dispatch paramters of kernels passed to clEnqueueNDRangeKernel") DECLARE_DEBUG_VARIABLE(bool, PrintProgramBinaryProcessingTime, false, "prints execution time of Program::processGenBinary() method during program building") +DECLARE_DEBUG_VARIABLE(bool, PrintRelocations, false, "prints relocations debug information") DECLARE_DEBUG_VARIABLE(bool, WddmResidencyLogger, false, "gather Wddm residency statistics to file") DECLARE_DEBUG_VARIABLE(int32_t, PrintDriverDiagnostics, -1, "prints driver diagnostics messages to standard output, value corresponds to hint level") diff --git a/shared/test/unit_test/compiler_interface/linker_tests.cpp b/shared/test/unit_test/compiler_interface/linker_tests.cpp index 890e40c01e..29dbbb788a 100644 --- a/shared/test/unit_test/compiler_interface/linker_tests.cpp +++ b/shared/test/unit_test/compiler_interface/linker_tests.cpp @@ -986,3 +986,45 @@ TEST(LinkerErrorMessageTests, givenListOfUnresolvedExternalsThenSymbolNameOrSymb EXPECT_THAT(err.c_str(), ::testing::HasSubstr(NEO::asString(NEO::SegmentType::Unknown))); EXPECT_THAT(err.c_str(), ::testing::HasSubstr(std::to_string(unresolvedExternal.unresolvedRelocation.offset).c_str())); } + +TEST(RelocationsDebugMessageTests, givenEmptyListOfRelocatedSymbolsTheReturnsEmptyString) { + auto message = NEO::constructRelocationsDebugMessage({}); + EXPECT_EQ(0U, message.size()) << message; +} + +TEST(RelocationsDebugMessageTests, givenListOfRelocatedSymbolsTheReturnsProperDebugMessage) { + NEO::Linker::RelocatedSymbolsMap symbols; + + auto &funcSymbol = symbols["foo"]; + auto &constDataSymbol = symbols["constInt"]; + auto &globalVarSymbol = symbols["intX"]; + funcSymbol.symbol.segment = NEO::SegmentType::Instructions; + funcSymbol.symbol.offset = 64U; + funcSymbol.symbol.size = 1024U; + funcSymbol.gpuAddress = 4096U; + + constDataSymbol.symbol.segment = NEO::SegmentType::GlobalConstants; + constDataSymbol.symbol.offset = 32U; + constDataSymbol.symbol.size = 16U; + constDataSymbol.gpuAddress = 8U; + + globalVarSymbol.symbol.segment = NEO::SegmentType::GlobalVariables; + globalVarSymbol.symbol.offset = 72U; + globalVarSymbol.symbol.size = 8U; + globalVarSymbol.gpuAddress = 256U; + + auto message = NEO::constructRelocationsDebugMessage(symbols); + + std::stringstream expected; + expected << "Relocations debug informations :\n"; + for (const auto &symbol : symbols) { + if (symbol.first == "foo") { + expected << " * \"foo\" [1024 bytes] INSTRUCTIONS_SEGMENT@64 -> 0x1000 GPUVA\n"; + } else if (symbol.first == "constInt") { + expected << " * \"constInt\" [16 bytes] GLOBAL_CONSTANTS_SEGMENT@32 -> 0x8 GPUVA\n"; + } else { + expected << " * \"intX\" [8 bytes] GLOBAL_VARIABLES_SEGMENT@72 -> 0x100 GPUVA\n"; + } + } + EXPECT_STREQ(expected.str().c_str(), message.c_str()); +}