Distinction between compilation ID and prefix is no longer necessary.

Signed-off-by: Julian Oppermann <[email protected]>

Author: jopperm

sycl/source/detail/jit_compiler.cpp

@@ -1142,7 +1142,7 @@ sycl_device_binaries jit_compiler::createPIDeviceBinary(
 
 sycl_device_binaries jit_compiler::createDeviceBinaryImage(
     const ::jit_compiler::RTCBundleInfo &BundleInfo,
-    const std::string &OffloadEntryPrefix) {
+    const std::string &Prefix) {
   DeviceBinariesCollection Collection;
 
   for (const auto &DevImgInfo : BundleInfo) {
@@ -1153,7 +1153,7 @@ sycl_device_binaries jit_compiler::createDeviceBinaryImage(
       // entrypoints remain unchanged.
       // It seems to be OK to set zero for most of the information here, at
       // least that is the case for compiled SPIR-V binaries.
-      std::string PrefixedName = OffloadEntryPrefix + Symbol.c_str();
+      std::string PrefixedName = Prefix + Symbol.c_str();
       OffloadEntryContainer Entry{PrefixedName, /*Addr=*/nullptr, /*Size=*/0,
                                   /*Flags=*/0, /*Reserved=*/0};
       Binary.addOffloadEntry(std::move(Entry));
@@ -1236,10 +1236,11 @@ std::vector<uint8_t> jit_compiler::encodeReqdWorkGroupSize(
 }
 
 sycl_device_binaries jit_compiler::compileSYCL(
-    const std::string &CompilationID, const std::string &SYCLSource,
+    const std::string &SYCLSource,
     const std::vector<std::pair<std::string, std::string>> &IncludePairs,
     const std::vector<std::string> &UserArgs, std::string *LogPtr,
-    const std::vector<std::string> &RegisteredKernelNames) {
+    const std::vector<std::string> &RegisteredKernelNames,
+    const std::string &Prefix) {
   auto appendToLog = [LogPtr](const char *Msg) {
     if (LogPtr) {
       LogPtr->append(Msg);
@@ -1309,8 +1310,7 @@ sycl_device_binaries jit_compiler::compileSYCL(
     PersistentDeviceCodeCache::putDeviceCodeIRToDisc(CacheKey, SavedIR);
   }
 
-  return createDeviceBinaryImage(Result.getBundleInfo(),
-                                 /*OffloadEntryPrefix=*/CompilationID + '$');
+  return createDeviceBinaryImage(Result.getBundleInfo(), Prefix);
 }
 
 } // namespace detail

sycl/source/detail/jit_compiler.hpp

@@ -48,10 +48,11 @@ class jit_compiler {
                            const std::vector<unsigned char> &SpecConstBlob);
 
   sycl_device_binaries compileSYCL(
-      const std::string &CompilationID, const std::string &SYCLSource,
+      const std::string &SYCLSource,
       const std::vector<std::pair<std::string, std::string>> &IncludePairs,
       const std::vector<std::string> &UserArgs, std::string *LogPtr,
-      const std::vector<std::string> &RegisteredKernelNames);
+      const std::vector<std::string> &RegisteredKernelNames,
+      const std::string &Prefix);
 
   bool isAvailable() { return Available; }
 
@@ -74,7 +75,7 @@ class jit_compiler {
 
   sycl_device_binaries
   createDeviceBinaryImage(const ::jit_compiler::RTCBundleInfo &BundleInfo,
-                          const std::string &OffloadEntryPrefix);
+                          const std::string &Prefix);
 
   std::vector<uint8_t>
   encodeArgUsageMask(const ::jit_compiler::ArgUsageMask &Mask) const;

sycl/source/detail/kernel_bundle_impl.hpp

@@ -497,7 +497,7 @@ class kernel_bundle_impl {
     if (MLanguage == syclex::source_language::sycl_jit) {
       // Build device images via the program manager.
       const std::string &SourceStr = std::get<std::string>(MSource);
-      auto [Binaries, CompilationID] = syclex::detail::SYCL_JIT_to_SPIRV(
+      auto [Binaries, Prefix] = syclex::detail::SYCL_JIT_to_SPIRV(
           SourceStr, MIncludePairs, BuildOptions, LogPtr,
           RegisteredKernelNames);
 
@@ -506,9 +506,6 @@ class kernel_bundle_impl {
 
       std::vector<kernel_id> KernelIDs;
       std::vector<std::string> KernelNames;
-      // `jit_compiler::compileSYCL(..)` uses `CompilationID + '$'` as prefix
-      // for offload entry names.
-      std::string Prefix = CompilationID + '$';
       for (const auto &KernelID : PM.getAllSYCLKernelIDs()) {
         std::string_view KernelName{KernelID.get_name()};
         if (KernelName.find(Prefix) == 0) {

sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp

@@ -305,18 +305,18 @@ bool SYCL_JIT_Compilation_Available() {
 
 std::pair<sycl_device_binaries, std::string> SYCL_JIT_to_SPIRV(
     [[maybe_unused]] const std::string &SYCLSource,
-    [[maybe_unused]] include_pairs_t IncludePairs,
+    [[maybe_unused]] const include_pairs_t& IncludePairs,
     [[maybe_unused]] const std::vector<std::string> &UserArgs,
     [[maybe_unused]] std::string *LogPtr,
     [[maybe_unused]] const std::vector<std::string> &RegisteredKernelNames) {
 #if SYCL_EXT_JIT_ENABLE
-  static std::atomic_uintptr_t CompilationCounter;
-  std::string CompilationID = "rtc_" + std::to_string(CompilationCounter++);
+  static std::atomic_uintptr_t CompilationID;
+  std::string Prefix = "rtc_" + std::to_string(CompilationID++) + "$";
   sycl_device_binaries Binaries =
       sycl::detail::jit_compiler::get_instance().compileSYCL(
-          CompilationID, SYCLSource, IncludePairs, UserArgs, LogPtr,
-          RegisteredKernelNames);
-  return std::make_pair(Binaries, std::move(CompilationID));
+          SYCLSource, IncludePairs, UserArgs, LogPtr,
+          RegisteredKernelNames, Prefix);
+  return std::make_pair(Binaries, std::move(Prefix));
 #else
   throw sycl::exception(sycl::errc::build,
                         "kernel_compiler via sycl-jit is not available");

sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp

@@ -41,7 +41,8 @@ std::string userArgsAsString(const std::vector<std::string> &UserArguments);
 // Returns a pointer to the image (owned by the `jit_compiler` class), and the
 // bundle-specific prefix used for loading the kernels.
 std::pair<sycl_device_binaries, std::string>
-SYCL_JIT_to_SPIRV(const std::string &Source, include_pairs_t IncludePairs,
+SYCL_JIT_to_SPIRV(const std::string &Source,
+                  const include_pairs_t &IncludePairs,
                   const std::vector<std::string> &UserArgs, std::string *LogPtr,
                   const std::vector<std::string> &RegisteredKernelNames);