fix: create temporary main.cl with kernel source when -g debug flag is used

Implements mechanism in ocloc to generate a temporary main.cl file
containing kernel source code when compiling with the debug (-g) flag
and the source file does not exist. This enables proper source code
annotations in generated assembly for SYCL/online compilation workflows.

Related-To: NEO-11900
Signed-off-by: Aleksandra Nizio <aleksandra.nizio@intel.com>

level_zero/core/test/black_box_tests/zello_world_jitc_ocloc.cpp

@@ -41,7 +41,7 @@ void executeKernelAndValidate(ze_context_handle_t &context, ze_device_handle_t &
     memset(dstBuffer, 0, allocSize);
 
     std::string buildLog;
-    auto spirV = LevelZeroBlackBoxTests::compileToSpirV(LevelZeroBlackBoxTests::memcpyBytesTestKernelSrc, "", buildLog);
+    auto spirV = LevelZeroBlackBoxTests::compileToSpirV(LevelZeroBlackBoxTests::memcpyBytesTestKernelSrc, "-g", buildLog);
     LevelZeroBlackBoxTests::printBuildLog(buildLog);
     SUCCESS_OR_TERMINATE((0 == spirV.size()));
 

opencl/test/unit_test/offline_compiler/mock/mock_offline_compiler.cpp

@@ -88,5 +88,8 @@ int MockOfflineCompiler::createDir(const std::string &path) {
         return OfflineCompiler::createDir(path);
     }
 }
-
+void MockOfflineCompiler::createTempSourceFileForDebug() {
+    OfflineCompiler::createTempSourceFileForDebug();
+    createTempSourceFileForDebugCalled++;
+}
 } // namespace NEO

opencl/test/unit_test/offline_compiler/mock/mock_offline_compiler.h

@@ -101,6 +101,7 @@ class MockOfflineCompiler : public OfflineCompiler {
     void clearLog();
 
     int createDir(const std::string &path) override;
+    void createTempSourceFileForDebug() override;
 
     std::map<std::string, std::string> filesMap{};
     int buildToIrBinaryStatus = 0;
@@ -116,6 +117,7 @@ class MockOfflineCompiler : public OfflineCompiler {
     std::optional<int> buildReturnValue{};
     bool interceptCreatedDirs{false};
     std::vector<std::string> createdDirs{};
+    int createTempSourceFileForDebugCalled{0};
 };
 
 } // namespace NEO

opencl/test/unit_test/offline_compiler/offline_compiler_tests.cpp

@@ -20,6 +20,7 @@
 #include "shared/source/device_binary_format/elf/ocl_elf.h"
 #include "shared/source/helpers/api_specific_config.h"
 #include "shared/source/helpers/compiler_product_helper.h"
+#include "shared/source/helpers/hash.h"
 #include "shared/source/helpers/hw_info.h"
 #include "shared/source/helpers/product_config_helper.h"
 #include "shared/test/common/helpers/debug_manager_state_restore.h"
@@ -5833,5 +5834,60 @@ TEST_F(OfflineCompilerTests, givenOneApiPvcSendWarWaEnvSetToFalseWhenInitializin
     std::string internalOptions = mockOfflineCompiler->internalOptions;
     EXPECT_TRUE(hasSubstr(internalOptions, "-ze-opt-disable-sendwarwa"));
 }
+TEST_F(OfflineCompilerTests, GivenDebugFlagWhenBuildingFromSourceThenTemporarySourceFileIsCreated) {
+    MockOfflineCompiler mockOfflineCompiler;
+    mockOfflineCompiler.uniqueHelper->interceptOutput = true;
+    mockOfflineCompiler.uniqueHelper->callBaseFileExists = false;
+    mockOfflineCompiler.uniqueHelper->callBaseLoadDataFromFile = false;
 
+    if (mockOfflineCompiler.cache == nullptr) {
+        mockOfflineCompiler.cache = std::make_unique<CompilerCacheMock>();
+    }
+
+    const char kernelSource[] = R"===(
+    __kernel void simpleKernel() {
+        int gid = get_global_id(0);
+    }
+    )===";
+
+    std::string inputFile = "kernel.cl";
+    mockOfflineCompiler.uniqueHelper->filesMap[inputFile] = kernelSource;
+
+    std::vector<std::string> argv = {
+        "ocloc",
+        "-file",
+        inputFile,
+        "-device",
+        gEnvironment->devicePrefix.c_str(),
+        "-options",
+        "-g"};
+
+    mockOfflineCompiler.buildReturnValue = OCLOC_SUCCESS;
+
+    int retVal = mockOfflineCompiler.parseCommandLine(argv.size(), argv);
+    EXPECT_EQ(OCLOC_SUCCESS, retVal);
+
+    mockOfflineCompiler.inputCodeType = IGC::CodeType::oclC;
+    mockOfflineCompiler.sourceCode = kernelSource;
+    mockOfflineCompiler.inputFile = inputFile;
+
+    uint64_t sourceHash = NEO::Hash::hash(kernelSource, strlen(kernelSource));
+    std::string tempFileName = "main_" + std::to_string(sourceHash) + ".cl";
+    std::filesystem::path expectedTempFilePath = std::filesystem::absolute(tempFileName);
+
+    StreamCapture capture;
+    capture.captureStdout();
+
+    mockOfflineCompiler.createTempSourceFileForDebug();
+
+    std::string capturedOutput = capture.getCapturedStdout();
+    EXPECT_NE(capturedOutput.find("Temporary source file for debug info created: " + expectedTempFilePath.string()), std::string::npos);
+
+    auto it = NEO::virtualFileList.find(expectedTempFilePath.string());
+    ASSERT_NE(NEO::virtualFileList.end(), it);
+    EXPECT_EQ(it->second.str(), kernelSource);
+
+    std::string expectedSOption = "-s \"" + expectedTempFilePath.string() + "\"";
+    EXPECT_TRUE(CompilerOptions::contains(mockOfflineCompiler.options, expectedSOption));
+}
 } // namespace NEO

opencl/test/unit_test/offline_compiler/offline_compiler_tests.h

@@ -18,7 +18,6 @@
 
 #include "environment.h"
 #include "gtest/gtest.h"
-#include "mock/mock_argument_helper.h"
 
 #include <cstdint>
 #include <memory>

shared/offline_compiler/source/offline_compiler.cpp

@@ -15,6 +15,7 @@
 #include "shared/offline_compiler/source/ocloc_supported_devices_helper.h"
 #include "shared/offline_compiler/source/queries.h"
 #include "shared/offline_compiler/source/utilities/get_git_version_info.h"
+#include "shared/source/compiler_interface/compiler_cache.h"
 #include "shared/source/compiler_interface/compiler_options.h"
 #include "shared/source/compiler_interface/compiler_options_extra.h"
 #include "shared/source/compiler_interface/default_cache_config.h"
@@ -28,6 +29,8 @@
 #include "shared/source/helpers/compiler_options_parser.h"
 #include "shared/source/helpers/compiler_product_helper.h"
 #include "shared/source/helpers/debug_helpers.h"
+#include "shared/source/helpers/file_io.h"
+#include "shared/source/helpers/hash.h"
 #include "shared/source/helpers/string.h"
 #include "shared/source/helpers/validators.h"
 #include "shared/source/os_interface/debug_env_reader.h"
@@ -37,6 +40,7 @@
 #include "neo_aot_platforms.h"
 #include "offline_compiler_ext.h"
 
+#include <filesystem>
 #include <iomanip>
 #include <iterator>
 #include <list>
@@ -729,6 +733,10 @@ int OfflineCompiler::build() {
         return igcInitializationResult;
     }
 
+    if (!inputFileSpirV()) {
+        createTempSourceFileForDebug();
+    }
+
     int retVal = OCLOC_SUCCESS;
     if (isOnlySpirV()) {
         retVal = buildToIrBinary();
@@ -1546,6 +1554,31 @@ bool OfflineCompiler::generateElfBinary() {
     return true;
 }
 
+void OfflineCompiler::createTempSourceFileForDebug() {
+    if (!CompilerOptions::contains(options, CompilerOptions::generateDebugInfo) || CompilerOptions::contains(options, CompilerOptions::useCMCompiler)) {
+        return;
+    }
+
+    std::string sourcePathOption = CompilerOptions::generateSourcePath.str();
+    size_t sourcePosInOpts = options.find(sourcePathOption);
+
+    if (sourcePosInOpts != std::string::npos) {
+        return;
+    }
+
+    uint64_t sourceHash = NEO::Hash::hash(sourceCode.c_str(), sourceCode.size());
+    std::string tempFilePath = "main_" + std::to_string(sourceHash) + ".cl";
+    std::filesystem::path absTempFilePath = std::filesystem::absolute(tempFilePath);
+
+    writeDataToFile(absTempFilePath.string().c_str(), std::string_view(sourceCode.c_str(), sourceCode.size()));
+
+    if (argHelper && !isQuiet()) {
+        argHelper->printf("Temporary source file for debug info created: %s\n", absTempFilePath.string().c_str());
+    }
+    std::string sourcePathWithFile = sourcePathOption + " " + CompilerOptions::wrapInQuotes(absTempFilePath.string());
+    options = CompilerOptions::concatenate(options, sourcePathWithFile);
+}
+
 void OfflineCompiler::writeOutAllFiles() {
     std::string fileBase;
     std::string fileTrunk = getFileNameTrunk(inputFile);

shared/offline_compiler/source/offline_compiler.h

@@ -269,6 +269,7 @@ All supported acronyms: %s.
     IGC::CodeType::CodeType_t intermediateRepresentation = IGC::CodeType::undefined;
 
     OclocArgHelper *argHelper = nullptr;
+    MOCKABLE_VIRTUAL void createTempSourceFileForDebug();
 };
 
 static_assert(NEO::NonCopyableAndNonMovable<OfflineCompiler>);