[SYCL] Make SYCL RT compatible with the new offload entry type

sycl/source/detail/compiler.hpp

@@ -76,14 +76,53 @@
 
 #define __SYCL_PROGRAM_METADATA_TAG_NEED_FINALIZATION "Requires finalization"
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGE
 // Entry type, matches OpenMP for compatibility
-struct _sycl_offload_entry_struct {
+struct _sycl_offload_entry_struct_legacy {
   void *addr;
   char *name;
   size_t size;
   int32_t flags;
   int32_t reserved;
 };
+using sycl_offload_entry_legacy = _sycl_offload_entry_struct_legacy *;
+#endif
+
+// New entry type after
+// https://github.com/llvm/llvm-project/pull/124018
+struct _sycl_offload_entry_struct {
+  /// Reserved bytes used to detect an older version of the struct, always zero.
+  uint64_t Reserved;
+  /// The current version of the struct for runtime forward compatibility.
+  uint16_t Version;
+  /// The expected consumer of this entry, e.g. CUDA or OpenMP.
+  uint16_t Kind;
+  /// Flags associated with the global.
+  uint32_t Flags;
+  /// The address of the global to be registered by the runtime.
+  void *Address;
+  /// The name of the symbol in the device image.
+  char *SymbolName;
+  /// The number of bytes the symbol takes.
+  uint64_t Size;
+  /// Extra generic data used to register this entry.
+  uint64_t Data;
+  /// An extra pointer, usually null.
+  void *AuxAddr;
+
+  // Name is the only field that's used in SYCL.
+  inline char *GetName() {
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGE
+    // Check if the first 64 bits of this struct are not zero, if so, this is an
+    // older version of the struct.
+    if (*(uint64_t *)(this)) {
+      // This is an older version of the struct, use the old name field.
+      return reinterpret_cast<_sycl_offload_entry_struct_legacy *>(this)->name;
+    }
+#endif
+    return SymbolName;
+  }
+};
 using sycl_offload_entry = _sycl_offload_entry_struct *;
 
 // A type of a binary image property.

sycl/source/detail/device_binary_image.cpp

@@ -122,7 +122,7 @@ void RTDeviceBinaryImage::print() const {
   std::cerr << "    Entries  : ";
   for (sycl_offload_entry EntriesIt = Bin->EntriesBegin;
        EntriesIt != Bin->EntriesEnd; ++EntriesIt)
-    std::cerr << EntriesIt->name << " ";
+    std::cerr << EntriesIt->GetName() << " ";
   std::cerr << "\n";
   std::cerr << "    Properties [" << Bin->PropertySetsBegin << "-"
             << Bin->PropertySetsEnd << "]:\n";

sycl/source/detail/jit_compiler.cpp

@@ -1081,7 +1081,10 @@ sycl_device_binaries jit_compiler::createPIDeviceBinary(
   // Create an offload entry for the fused kernel.
   // 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.
-  OffloadEntryContainer Entry{FusedKernelName, nullptr, 0, 0, 0};
+  OffloadEntryContainer Entry{
+      0 /*Reserved*/, 0 /*Version*/,       0 /*Kind*/,
+      0 /*Flags*/,    nullptr /*Address*/, FusedKernelName /*Name*/,
+      0 /*Size*/,     0 /*Data*/,          nullptr /*AuxAddr*/};
   Binary.addOffloadEntry(std::move(Entry));
 
   // Create a property entry for the argument usage mask for the fused kernel.
@@ -1154,8 +1157,10 @@ sycl_device_binaries jit_compiler::createDeviceBinaries(
       // 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();
-      OffloadEntryContainer Entry{PrefixedName, /*Addr=*/nullptr, /*Size=*/0,
-                                  /*Flags=*/0, /*Reserved=*/0};
+      OffloadEntryContainer Entry{
+          0 /*Reserved*/, 0 /*Version*/,       0 /*Kind*/,
+          0 /*Flags*/,    nullptr /*Address*/, PrefixedName /*Name*/,
+          0 /*Size*/,     0 /*Data*/,          nullptr /*AuxAddr*/};
       Binary.addOffloadEntry(std::move(Entry));
     }
 

sycl/source/detail/jit_device_binaries.cpp

@@ -14,17 +14,22 @@ namespace sycl {
 inline namespace _V1 {
 namespace detail {
 
-OffloadEntryContainer::OffloadEntryContainer(const std::string &Name,
-                                             void *Addr, size_t Size,
-                                             int32_t Flags, int32_t Reserved)
-    : KernelName{new char[Name.length() + 1]}, Address{Addr}, EntrySize{Size},
-      EntryFlags{Flags}, EntryReserved{Reserved} {
-  std::memcpy(KernelName.get(), Name.c_str(), Name.length() + 1);
+OffloadEntryContainer::OffloadEntryContainer(uint64_t _Reserved,
+                                             uint16_t _Version, uint16_t _Kind,
+                                             uint32_t _Flags, void *_Addr,
+                                             const std::string &_Name,
+                                             uint64_t _Size, uint64_t _Data,
+                                             void *_AuxAddr)
+    : Reserved{_Reserved}, Version{_Version}, Kind{_Kind}, EntryFlags{_Flags},
+      Address{_Addr}, KernelName{new char[_Name.length() + 1]},
+      EntrySize{_Size}, Data{_Data}, AuxAddr{_AuxAddr} {
+  std::memcpy(KernelName.get(), _Name.c_str(), _Name.length() + 1);
 }
 
 _sycl_offload_entry_struct OffloadEntryContainer::getPIOffloadEntry() {
-  return _sycl_offload_entry_struct{Address, KernelName.get(), EntrySize,
-                                    EntryFlags, EntryReserved};
+  return _sycl_offload_entry_struct{Reserved,   Version, Kind,
+                                    EntryFlags, Address, KernelName.get(),
+                                    EntrySize,  Data,    AuxAddr};
 }
 
 PropertyContainer::PropertyContainer(const std::string &Name, const void *Data,

sycl/source/detail/jit_device_binaries.hpp

@@ -18,13 +18,14 @@ namespace sycl {
 inline namespace _V1 {
 namespace detail {
 
-/// Representation of _sycl_offload_entry_struct for creation of JIT device
-/// binaries at runtime. Owns the necessary data and provides raw pointers for
-/// the UR struct.
+/// Representation of _sycl_offload_entry_struct for creation of JIT
+/// device binaries at runtime. Owns the necessary data and provides raw
+/// pointers for the UR struct.
 class OffloadEntryContainer {
 public:
-  OffloadEntryContainer(const std::string &Name, void *Addr, size_t Size,
-                        int32_t Flags, int32_t Reserved);
+  OffloadEntryContainer(uint64_t _Reserved, uint16_t _Version, uint16_t _Kind,
+                        uint32_t _Flags, void *_Addr, const std::string &_Name,
+                        uint64_t _Size, uint64_t _Data, void *_AuxAddr);
 
   OffloadEntryContainer(OffloadEntryContainer &&) = default;
   OffloadEntryContainer &operator=(OffloadEntryContainer &&) = default;
@@ -36,12 +37,25 @@ class OffloadEntryContainer {
   _sycl_offload_entry_struct getPIOffloadEntry();
 
 private:
-  std::unique_ptr<char[]> KernelName;
-
+  // Reserved bytes used to detect an older version of the struct, always zero.
+  uint64_t Reserved;
+  // The current version of the struct for runtime forward compatibility.
+  uint16_t Version;
+  // The expected consumer of this entry, e.g. CUDA or OpenMP.
+  uint16_t Kind;
+  // Flags associated with the global.
+  uint32_t EntryFlags;
+  // The address of the global to be registered by the runtime.
   void *Address;
-  size_t EntrySize;
-  int32_t EntryFlags;
-  int32_t EntryReserved;
+  // The name of the symbol in the device image. This is the only field that's
+  // used in SYCL.
+  std::unique_ptr<char[]> KernelName;
+  // The number of bytes the symbol takes.
+  uint64_t EntrySize;
+  // Extra generic data used to register this entry.
+  uint64_t Data;
+  // An extra pointer, usually null.
+  void *AuxAddr;
 };
 
 /// Representation of _sycl_device_binary_property_struct for creation of JIT

sycl/source/detail/persistent_device_code_cache.cpp

@@ -103,8 +103,8 @@ getSortedImages(const std::vector<const RTDeviceBinaryImage *> &Imgs) {
             [](const RTDeviceBinaryImage *A, const RTDeviceBinaryImage *B) {
               // All entry names are unique among these images, so comparing the
               // first ones is enough.
-              return std::strcmp(A->getRawData().EntriesBegin->name,
-                                 B->getRawData().EntriesBegin->name) < 0;
+              return std::strcmp(A->getRawData().EntriesBegin->GetName(),
+                                 B->getRawData().EntriesBegin->GetName()) < 0;
             });
   return SortedImgs;
 }

sycl/source/detail/program_manager/program_manager.cpp

@@ -1869,31 +1869,32 @@ void ProgramManager::addImages(sycl_device_binaries DeviceBinary) {
     for (sycl_offload_entry EntriesIt = EntriesB; EntriesIt != EntriesE;
          ++EntriesIt) {
 
+      auto name = EntriesIt->GetName();
+
       // Skip creating unique kernel ID if it is a service kernel.
       // SYCL service kernels are identified by having
       // __sycl_service_kernel__ in the mangled name, primarily as part of
       // the namespace of the name type.
-      if (std::strstr(EntriesIt->name, "__sycl_service_kernel__")) {
-        m_ServiceKernels.insert(std::make_pair(EntriesIt->name, Img.get()));
+      if (std::strstr(name, "__sycl_service_kernel__")) {
+        m_ServiceKernels.insert(std::make_pair(name, Img.get()));
         continue;
       }
 
       // Skip creating unique kernel ID if it is an exported device
       // function. Exported device functions appear in the offload entries
       // among kernels, but are identifiable by being listed in properties.
-      if (m_ExportedSymbolImages.find(EntriesIt->name) !=
-          m_ExportedSymbolImages.end())
+      if (m_ExportedSymbolImages.find(name) != m_ExportedSymbolImages.end())
         continue;
 
       // ... and create a unique kernel ID for the entry
-      auto It = m_KernelName2KernelIDs.find(EntriesIt->name);
+      auto It = m_KernelName2KernelIDs.find(name);
       if (It == m_KernelName2KernelIDs.end()) {
         std::shared_ptr<detail::kernel_id_impl> KernelIDImpl =
-            std::make_shared<detail::kernel_id_impl>(EntriesIt->name);
+            std::make_shared<detail::kernel_id_impl>(name);
         sycl::kernel_id KernelID =
             detail::createSyclObjFromImpl<sycl::kernel_id>(KernelIDImpl);
 
-        It = m_KernelName2KernelIDs.emplace_hint(It, EntriesIt->name, KernelID);
+        It = m_KernelName2KernelIDs.emplace_hint(It, name, KernelID);
       }
       m_KernelIDs2BinImage.insert(std::make_pair(It->second, Img.get()));
       m_BinImg2KernelIDs[Img.get()]->push_back(It->second);
@@ -2008,7 +2009,7 @@ void ProgramManager::removeImages(sycl_device_binaries DeviceBinary) {
       continue;
 
     // Retrieve RTDeviceBinaryImage by looking up the first offload entry
-    kernel_id FirstKernelID = getSYCLKernelID(RawImg->EntriesBegin->name);
+    kernel_id FirstKernelID = getSYCLKernelID(RawImg->EntriesBegin->GetName());
     auto RTDBImages = getRawDeviceImages({FirstKernelID});
     assert(RTDBImages.size() == 1);
 
@@ -2025,18 +2026,18 @@ void ProgramManager::removeImages(sycl_device_binaries DeviceBinary) {
          ++EntriesIt) {
 
       // Drop entry for service kernel
-      if (std::strstr(EntriesIt->name, "__sycl_service_kernel__")) {
-        m_ServiceKernels.erase(EntriesIt->name);
+      if (std::strstr(EntriesIt->GetName(), "__sycl_service_kernel__")) {
+        m_ServiceKernels.erase(EntriesIt->GetName());
         continue;
       }
 
       // Exported device functions won't have a kernel ID
-      if (m_ExportedSymbolImages.find(EntriesIt->name) !=
+      if (m_ExportedSymbolImages.find(EntriesIt->GetName()) !=
           m_ExportedSymbolImages.end()) {
         continue;
       }
 
-      auto It = m_KernelName2KernelIDs.find(EntriesIt->name);
+      auto It = m_KernelName2KernelIDs.find(EntriesIt->GetName());
       assert(It != m_KernelName2KernelIDs.end());
       m_KernelName2KernelIDs.erase(It);
       m_KernelIDs2BinImage.erase(It->second);

sycl/unittests/helpers/MockDeviceImage.hpp

@@ -80,7 +80,7 @@ class MockOffloadEntry {
   using NativeType = _sycl_offload_entry_struct;
 
   MockOffloadEntry(const std::string &Name, std::vector<char> Data,
-                   int32_t Flags)
+                   uint32_t Flags)
       : MName(Name), MData(std::move(Data)), MFlags(Flags) {
     updateNativeType();
   }
@@ -104,13 +104,13 @@ class MockOffloadEntry {
 private:
   void updateNativeType() {
     MNative = NativeType{
-        const_cast<char *>(MData.data()), MName.data(), MData.size(), MFlags,
-        0 // Reserved
-    };
+        0 /*Reserved*/, 0 /*Version*/,       0 /*Kind*/,
+        MFlags,         nullptr /*Address*/, const_cast<char *>(MName.data()),
+        MData.size(),   0 /*Data*/,          nullptr /*AuxAddr*/};
   }
   std::string MName;
   std::vector<char> MData;
-  int32_t MFlags;
+  uint32_t MFlags;
   NativeType MNative;
 };
 

sycl/unittests/kernel-and-program/PersistentDeviceCodeCache.cpp

@@ -261,8 +261,15 @@ class PersistentDeviceCodeCache
   device Dev;
   const char *EntryName = "Entry";
   _sycl_offload_entry_struct EntryStruct = {
-      /*addr*/ nullptr, const_cast<char *>(EntryName), strlen(EntryName),
-      /*flags*/ 0, /*reserved*/ 0};
+      /*Reserved*/ 0,
+      /*Version*/ 0,
+      /*Kind*/ 0,
+      /*Flags*/ 0,
+      /*Address*/ nullptr,
+      /*Name*/ const_cast<char *>(EntryName),
+      /*Size*/ strlen(EntryName),
+      /*Data*/ 0,
+      /*AuxAddr*/ nullptr};
   sycl_device_binary_struct BinStruct{/*Version*/ 1,
                                       /*Kind*/ 4,
                                       /*Format*/ GetParam(),
@@ -311,8 +318,15 @@ TEST_P(PersistentDeviceCodeCache, KeysWithNullTermSymbol) {
 TEST_P(PersistentDeviceCodeCache, MultipleImages) {
   const char *ExtraEntryName = "ExtraEntry";
   _sycl_offload_entry_struct ExtraEntryStruct = {
-      /*addr*/ nullptr, const_cast<char *>(ExtraEntryName),
-      strlen(ExtraEntryName), /*flags*/ 0, /*reserved*/ 0};
+      /*Reserved*/ 0,
+      /*Version*/ 0,
+      /*Kind*/ 0,
+      /*Flags*/ 0,
+      /*Address*/ nullptr,
+      /*Name*/ const_cast<char *>(ExtraEntryName),
+      /*Size*/ strlen(ExtraEntryName),
+      /*Data*/ 0,
+      /*AuxAddr*/ nullptr};
   sycl_device_binary_struct ExtraBinStruct{/*Version*/ 1,
                                            /*Kind*/ 4,
                                            /*Format*/ GetParam(),
@@ -336,8 +350,8 @@ TEST_P(PersistentDeviceCodeCache, MultipleImages) {
   std::sort(Imgs.begin(), Imgs.end(),
             [](const detail::RTDeviceBinaryImage *A,
                const detail::RTDeviceBinaryImage *B) {
-              return std::strcmp(A->getRawData().EntriesBegin->name,
-                                 B->getRawData().EntriesBegin->name) < 0;
+              return std::strcmp(A->getRawData().EntriesBegin->GetName(),
+                                 B->getRawData().EntriesBegin->GetName()) < 0;
             });
   std::string ItemDir = detail::PersistentDeviceCodeCache::getCacheItemPath(
       Dev, Imgs, {}, BuildOptions);