From 254245f699e52868db23e728c4d2ad71f3ec746b Mon Sep 17 00:00:00 2001
From: "Agarwal, Udit" <udit.agarwal@intel.com>
Date: Thu, 20 Feb 2025 16:52:45 -0800
Subject: [PATCH 01/13] [SYCL] Make SYCL RT compatible with the new offload
 entry type

---
 sycl/source/detail/compiler.hpp               | 41 ++++++++++++++++++-
 sycl/source/detail/device_binary_image.cpp    |  2 +-
 sycl/source/detail/jit_compiler.cpp           | 11 +++--
 sycl/source/detail/jit_device_binaries.cpp    | 21 ++++++----
 sycl/source/detail/jit_device_binaries.hpp    | 34 ++++++++++-----
 .../detail/persistent_device_code_cache.cpp   |  4 +-
 .../program_manager/program_manager.cpp       | 25 +++++------
 sycl/unittests/helpers/MockDeviceImage.hpp    | 10 ++---
 .../PersistentDeviceCodeCache.cpp             | 26 +++++++++---
 9 files changed, 126 insertions(+), 48 deletions(-)

diff --git a/sycl/source/detail/compiler.hpp b/sycl/source/detail/compiler.hpp
index 40bf97299138f..8f16751364287 100644
--- a/sycl/source/detail/compiler.hpp
+++ b/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.
diff --git a/sycl/source/detail/device_binary_image.cpp b/sycl/source/detail/device_binary_image.cpp
index 633a4269e1e78..590086b213e4c 100644
--- a/sycl/source/detail/device_binary_image.cpp
+++ b/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";
diff --git a/sycl/source/detail/jit_compiler.cpp b/sycl/source/detail/jit_compiler.cpp
index 6fc88bb812a20..efcdf76414444 100644
--- a/sycl/source/detail/jit_compiler.cpp
+++ b/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));
     }
 
diff --git a/sycl/source/detail/jit_device_binaries.cpp b/sycl/source/detail/jit_device_binaries.cpp
index f90be2c27ec3a..4be437da4c6ff 100644
--- a/sycl/source/detail/jit_device_binaries.cpp
+++ b/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,
diff --git a/sycl/source/detail/jit_device_binaries.hpp b/sycl/source/detail/jit_device_binaries.hpp
index bca83839f39e8..9f53662991801 100644
--- a/sycl/source/detail/jit_device_binaries.hpp
+++ b/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
diff --git a/sycl/source/detail/persistent_device_code_cache.cpp b/sycl/source/detail/persistent_device_code_cache.cpp
index 213948d526f59..4c7bac4177c95 100644
--- a/sycl/source/detail/persistent_device_code_cache.cpp
+++ b/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;
 }
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index 21ee135074ef0..fbc65bde573de 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/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);
diff --git a/sycl/unittests/helpers/MockDeviceImage.hpp b/sycl/unittests/helpers/MockDeviceImage.hpp
index fea80d6b08c3e..0e986995cf398 100644
--- a/sycl/unittests/helpers/MockDeviceImage.hpp
+++ b/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;
 };
 
diff --git a/sycl/unittests/kernel-and-program/PersistentDeviceCodeCache.cpp b/sycl/unittests/kernel-and-program/PersistentDeviceCodeCache.cpp
index 21fc02c94840a..07612dbbb49c7 100644
--- a/sycl/unittests/kernel-and-program/PersistentDeviceCodeCache.cpp
+++ b/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);

From d7accea5196fd81943e17d3e566be9a5317715ec Mon Sep 17 00:00:00 2001
From: "Agarwal, Udit" <udit.agarwal@intel.com>
Date: Thu, 27 Feb 2025 15:56:52 -0800
Subject: [PATCH 02/13] Fix offload entry iteratoe when using old vs. new
 offload entry format

---
 sycl/source/detail/compiler.hpp               | 12 +++++++----
 .../program_manager/program_manager.cpp       | 20 +++++++++++++++----
 2 files changed, 24 insertions(+), 8 deletions(-)

diff --git a/sycl/source/detail/compiler.hpp b/sycl/source/detail/compiler.hpp
index 8f16751364287..9cd795f332bf8 100644
--- a/sycl/source/detail/compiler.hpp
+++ b/sycl/source/detail/compiler.hpp
@@ -110,16 +110,20 @@ struct _sycl_offload_entry_struct {
   /// An extra pointer, usually null.
   void *AuxAddr;
 
+  inline bool IsLegacy() {
+    // Check if first 128 bits of this struct are zero, if so, this is a newer
+    // version of the struct.
+    return *(uint64_t *)(this) || *(((uint64_t *)this) + 1);
+  }
+
   // 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)) {
+    if (IsLegacy()) {
       // This is an older version of the struct, use the old name field.
-      return reinterpret_cast<_sycl_offload_entry_struct_legacy *>(this)->name;
+      return reinterpret_cast<sycl_offload_entry_legacy>(this)->name;
     }
-#endif
     return SymbolName;
   }
 };
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index d16a6a188459b..1ae81b4567576 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -1866,8 +1866,7 @@ void ProgramManager::addImages(sycl_device_binaries DeviceBinary) {
 
     m_BinImg2KernelIDs[Img.get()].reset(new std::vector<kernel_id>);
 
-    for (sycl_offload_entry EntriesIt = EntriesB; EntriesIt != EntriesE;
-         ++EntriesIt) {
+    for (sycl_offload_entry EntriesIt = EntriesB; EntriesIt != EntriesE;) {
 
       auto name = EntriesIt->GetName();
 
@@ -1898,6 +1897,13 @@ void ProgramManager::addImages(sycl_device_binaries DeviceBinary) {
       }
       m_KernelIDs2BinImage.insert(std::make_pair(It->second, Img.get()));
       m_BinImg2KernelIDs[Img.get()]->push_back(It->second);
+
+      // Increment iterator.
+      if (EntriesIt->IsLegacy())
+        EntriesIt = reinterpret_cast<sycl_offload_entry>(
+            reinterpret_cast<sycl_offload_entry_legacy>(EntriesIt) + 1);
+      else
+        EntriesIt++;
     }
 
     cacheKernelUsesAssertInfo(*Img);
@@ -2020,8 +2026,7 @@ void ProgramManager::removeImages(sycl_device_binaries DeviceBinary) {
     std::lock_guard<std::mutex> KernelIDsGuard(m_KernelIDsMutex);
 
     // Unmap the unique kernel IDs for the offload entries
-    for (sycl_offload_entry EntriesIt = EntriesB; EntriesIt != EntriesE;
-         ++EntriesIt) {
+    for (sycl_offload_entry EntriesIt = EntriesB; EntriesIt != EntriesE;) {
 
       // Drop entry for service kernel
       if (std::strstr(EntriesIt->GetName(), "__sycl_service_kernel__")) {
@@ -2044,6 +2049,13 @@ void ProgramManager::removeImages(sycl_device_binaries DeviceBinary) {
         m_KernelName2KernelIDs.erase(It);
         m_KernelIDs2BinImage.erase(It->second);
       }
+
+      // Increment iterator.
+      if (EntriesIt->IsLegacy())
+        EntriesIt = reinterpret_cast<sycl_offload_entry>(
+            reinterpret_cast<sycl_offload_entry_legacy>(EntriesIt) + 1);
+      else
+        EntriesIt++;
     }
 
     // Drop reverse mapping

From 23bbab409f992329c0ebd94b4810f2cb63b524da Mon Sep 17 00:00:00 2001
From: "Agarwal, Udit" <udit.agarwal@intel.com>
Date: Fri, 28 Feb 2025 12:17:47 -0800
Subject: [PATCH 03/13] Fix iterator increment

---
 .../program_manager/program_manager.cpp       | 36 ++++++++++---------
 1 file changed, 20 insertions(+), 16 deletions(-)

diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index 1ae81b4567576..64341049bf3fd 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -1866,7 +1866,16 @@ void ProgramManager::addImages(sycl_device_binaries DeviceBinary) {
 
     m_BinImg2KernelIDs[Img.get()].reset(new std::vector<kernel_id>);
 
-    for (sycl_offload_entry EntriesIt = EntriesB; EntriesIt != EntriesE;) {
+    sycl_offload_entry EntriesIt;
+    auto IncrementEntriesIt = [&]() {
+      if (EntriesIt->IsLegacy())
+        EntriesIt = reinterpret_cast<sycl_offload_entry>(
+            reinterpret_cast<sycl_offload_entry_legacy>(EntriesIt) + 1);
+      else
+        EntriesIt++;
+    };
+
+    for (EntriesIt = EntriesB; EntriesIt != EntriesE; IncrementEntriesIt()) {
 
       auto name = EntriesIt->GetName();
 
@@ -1897,13 +1906,6 @@ void ProgramManager::addImages(sycl_device_binaries DeviceBinary) {
       }
       m_KernelIDs2BinImage.insert(std::make_pair(It->second, Img.get()));
       m_BinImg2KernelIDs[Img.get()]->push_back(It->second);
-
-      // Increment iterator.
-      if (EntriesIt->IsLegacy())
-        EntriesIt = reinterpret_cast<sycl_offload_entry>(
-            reinterpret_cast<sycl_offload_entry_legacy>(EntriesIt) + 1);
-      else
-        EntriesIt++;
     }
 
     cacheKernelUsesAssertInfo(*Img);
@@ -2025,8 +2027,17 @@ void ProgramManager::removeImages(sycl_device_binaries DeviceBinary) {
     // Acquire lock to modify maps for kernel bundles
     std::lock_guard<std::mutex> KernelIDsGuard(m_KernelIDsMutex);
 
+    sycl_offload_entry EntriesIt;
+    auto IncrementEntriesIt = [&]() {
+      if (EntriesIt->IsLegacy())
+        EntriesIt = reinterpret_cast<sycl_offload_entry>(
+            reinterpret_cast<sycl_offload_entry_legacy>(EntriesIt) + 1);
+      else
+        EntriesIt++;
+    };
+
     // Unmap the unique kernel IDs for the offload entries
-    for (sycl_offload_entry EntriesIt = EntriesB; EntriesIt != EntriesE;) {
+    for (EntriesIt = EntriesB; EntriesIt != EntriesE; IncrementEntriesIt()) {
 
       // Drop entry for service kernel
       if (std::strstr(EntriesIt->GetName(), "__sycl_service_kernel__")) {
@@ -2049,13 +2060,6 @@ void ProgramManager::removeImages(sycl_device_binaries DeviceBinary) {
         m_KernelName2KernelIDs.erase(It);
         m_KernelIDs2BinImage.erase(It->second);
       }
-
-      // Increment iterator.
-      if (EntriesIt->IsLegacy())
-        EntriesIt = reinterpret_cast<sycl_offload_entry>(
-            reinterpret_cast<sycl_offload_entry_legacy>(EntriesIt) + 1);
-      else
-        EntriesIt++;
     }
 
     // Drop reverse mapping

From ecf90ce56e1bbd8562f53223c697343e916233d0 Mon Sep 17 00:00:00 2001
From: "Agarwal, Udit" <udit.agarwal@intel.com>
Date: Sun, 2 Mar 2025 14:26:47 -0800
Subject: [PATCH 04/13] Fix warning

---
 sycl/source/detail/program_manager/program_manager.cpp | 6 +++---
 1 file changed, 3 insertions(+), 3 deletions(-)

diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index 64341049bf3fd..3347c940d9e0d 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -1869,13 +1869,13 @@ void ProgramManager::addImages(sycl_device_binaries DeviceBinary) {
     sycl_offload_entry EntriesIt;
     auto IncrementEntriesIt = [&]() {
       if (EntriesIt->IsLegacy())
-        EntriesIt = reinterpret_cast<sycl_offload_entry>(
+        return reinterpret_cast<sycl_offload_entry>(
             reinterpret_cast<sycl_offload_entry_legacy>(EntriesIt) + 1);
       else
-        EntriesIt++;
+        return EntriesIt + 1;
     };
 
-    for (EntriesIt = EntriesB; EntriesIt != EntriesE; IncrementEntriesIt()) {
+    for (EntriesIt = EntriesB; EntriesIt != EntriesE; EntriesIt = IncrementEntriesIt()) {
 
       auto name = EntriesIt->GetName();
 

From 8907ff306cb89543412bb8040761cb2237ddbc24 Mon Sep 17 00:00:00 2001
From: "Agarwal, Udit" <udit.agarwal@intel.com>
Date: Sun, 2 Mar 2025 14:39:28 -0800
Subject: [PATCH 05/13] Fix more warnings

---
 sycl/source/detail/program_manager/program_manager.cpp | 6 +++---
 1 file changed, 3 insertions(+), 3 deletions(-)

diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index 3347c940d9e0d..ac6d3f9c32133 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -2030,14 +2030,14 @@ void ProgramManager::removeImages(sycl_device_binaries DeviceBinary) {
     sycl_offload_entry EntriesIt;
     auto IncrementEntriesIt = [&]() {
       if (EntriesIt->IsLegacy())
-        EntriesIt = reinterpret_cast<sycl_offload_entry>(
+        return reinterpret_cast<sycl_offload_entry>(
             reinterpret_cast<sycl_offload_entry_legacy>(EntriesIt) + 1);
       else
-        EntriesIt++;
+        return EntriesIt + 1;
     };
 
     // Unmap the unique kernel IDs for the offload entries
-    for (EntriesIt = EntriesB; EntriesIt != EntriesE; IncrementEntriesIt()) {
+    for (EntriesIt = EntriesB; EntriesIt != EntriesE; EntriesIt = IncrementEntriesIt()) {
 
       // Drop entry for service kernel
       if (std::strstr(EntriesIt->GetName(), "__sycl_service_kernel__")) {

From 3f05144d7f237cb09ac44f49fa95edaf2920d51f Mon Sep 17 00:00:00 2001
From: "Agarwal, Udit" <udit.agarwal@intel.com>
Date: Sun, 2 Mar 2025 14:41:05 -0800
Subject: [PATCH 06/13] clang format

---
 sycl/source/detail/program_manager/program_manager.cpp | 6 ++++--
 1 file changed, 4 insertions(+), 2 deletions(-)

diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index ac6d3f9c32133..8f8d5cdbf6445 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -1875,7 +1875,8 @@ void ProgramManager::addImages(sycl_device_binaries DeviceBinary) {
         return EntriesIt + 1;
     };
 
-    for (EntriesIt = EntriesB; EntriesIt != EntriesE; EntriesIt = IncrementEntriesIt()) {
+    for (EntriesIt = EntriesB; EntriesIt != EntriesE;
+         EntriesIt = IncrementEntriesIt()) {
 
       auto name = EntriesIt->GetName();
 
@@ -2037,7 +2038,8 @@ void ProgramManager::removeImages(sycl_device_binaries DeviceBinary) {
     };
 
     // Unmap the unique kernel IDs for the offload entries
-    for (EntriesIt = EntriesB; EntriesIt != EntriesE; EntriesIt = IncrementEntriesIt()) {
+    for (EntriesIt = EntriesB; EntriesIt != EntriesE;
+         EntriesIt = IncrementEntriesIt()) {
 
       // Drop entry for service kernel
       if (std::strstr(EntriesIt->GetName(), "__sycl_service_kernel__")) {

From 7d5c1509f524e3bc72034dfa0f2fbe7ac5e64d55 Mon Sep 17 00:00:00 2001
From: "Agarwal, Udit" <udit.agarwal@intel.com>
Date: Mon, 3 Mar 2025 09:13:58 -0800
Subject: [PATCH 07/13] Revert "[SYCL] Use legacy offload entry type for now"

This reverts commit 0f19fbd63746ad4019febfc9726bbde3dc83025f.
---
 .../lib/Frontend/Offloading/SYCLOffloadWrapper.cpp | 14 +-------------
 1 file changed, 1 insertion(+), 13 deletions(-)

diff --git a/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp b/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp
index 7ad1c24f27b30..5cd54df70e3ed 100644
--- a/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp
+++ b/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp
@@ -62,18 +62,6 @@ int8_t binaryImageFormatToInt8(SYCLBinaryImageFormat Format) {
   }
 }
 
-StructType* getLegacyOffloadEntryTy(Module &M) {
-  LLVMContext &C = M.getContext();
-  StructType *EntryTy =
-      StructType::getTypeByName(C, "struct.__tgt_offload_entry");
-  if (!EntryTy)
-    EntryTy = StructType::create(
-        "struct.__tgt_offload_entry", PointerType::getUnqual(C),
-        PointerType::getUnqual(C), M.getDataLayout().getIntPtrType(C),
-        Type::getInt32Ty(C), Type::getInt32Ty(C));
-  return EntryTy;
-}
-
 /// Wrapper helper class that creates all LLVM IRs wrapping given images.
 /// Note: All created structures, "_pi_device_*", "__sycl_*" and "__tgt*" names
 /// in this implementation are aligned with "sycl/include/sycl/detail/pi.h".
@@ -95,7 +83,7 @@ struct Wrapper {
 
     SyclPropTy = getSyclPropTy();
     SyclPropSetTy = getSyclPropSetTy();
-    EntryTy = getLegacyOffloadEntryTy(M);
+    EntryTy = offloading::getEntryTy(M);
     SyclDeviceImageTy = getSyclDeviceImageTy();
     SyclBinDescTy = getSyclBinDescTy();
   }

From bc61a95a31698802a9c86be15bdc710728fecef6 Mon Sep 17 00:00:00 2001
From: "Agarwal, Udit" <udit.agarwal@intel.com>
Date: Mon, 3 Mar 2025 16:09:41 -0800
Subject: [PATCH 08/13] Fix how we detect newer/older version of the struct

---
 .../Offloading/SYCLOffloadWrapper.cpp         | 20 ++++++++++++++-----
 sycl/source/detail/compiler.hpp               | 12 +++++++----
 2 files changed, 23 insertions(+), 9 deletions(-)

diff --git a/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp b/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp
index 5cd54df70e3ed..3d227d0c2e050 100644
--- a/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp
+++ b/llvm/lib/Frontend/Offloading/SYCLOffloadWrapper.cpp
@@ -387,16 +387,26 @@ struct Wrapper {
       return std::pair<Constant *, Constant *>(NullPtr, NullPtr);
     }
 
-    auto *Zero = ConstantInt::get(getSizeTTy(), 0);
+    auto *I64Zero = ConstantInt::get(Type::getInt64Ty(C), 0);
     auto *I32Zero = ConstantInt::get(Type::getInt32Ty(C), 0);
     auto *NullPtr = Constant::getNullValue(PointerType::getUnqual(C));
 
     SmallVector<Constant *> EntriesInits;
     std::unique_ptr<MemoryBuffer> MB = MemoryBuffer::getMemBuffer(Entries);
-    for (line_iterator LI(*MB); !LI.is_at_eof(); ++LI)
-      EntriesInits.push_back(ConstantStruct::get(
-          EntryTy, NullPtr, addStringToModule(*LI, "__sycl_offload_entry_name"),
-          Zero, I32Zero, I32Zero));
+    for (line_iterator LI(*MB); !LI.is_at_eof(); ++LI) {
+      Constant *EntryData[] = {
+          ConstantExpr::getNullValue(Type::getInt64Ty(C)),
+          ConstantInt::get(Type::getInt16Ty(C), 1),
+          ConstantInt::get(Type::getInt16Ty(C), object::OffloadKind::OFK_SYCL),
+          I32Zero,
+          NullPtr,
+          addStringToModule(*LI, "__sycl_offload_entry_name"),
+          I64Zero,
+          I64Zero,
+          NullPtr};
+
+      EntriesInits.push_back(ConstantStruct::get(EntryTy, EntryData));
+    }
 
     auto *Arr = ConstantArray::get(ArrayType::get(EntryTy, EntriesInits.size()),
                                    EntriesInits);
diff --git a/sycl/source/detail/compiler.hpp b/sycl/source/detail/compiler.hpp
index a636c17b3d815..15668cf469a2b 100644
--- a/sycl/source/detail/compiler.hpp
+++ b/sycl/source/detail/compiler.hpp
@@ -112,10 +112,14 @@ struct _sycl_offload_entry_struct {
   /// An extra pointer, usually null.
   void *AuxAddr;
 
-  inline bool IsLegacy() {
-    // Check if first 128 bits of this struct are zero, if so, this is a newer
-    // version of the struct.
-    return *(uint64_t *)(this) || *(((uint64_t *)this) + 1);
+  bool IsLegacy() {
+    // Chek if first 64 bits is 0, next 16 bits is equal to 1, next 16 bits
+    // is equal to 4 (OK_SYCL), and Flags should be zero. If all these
+    // conditions are met, then this is a newer version of the struct.
+    // We can not just rely on checking the first 64 bits, because even for the
+    // older version of the struct, the first 64 bits (void* addr) are zero.
+    return !(this->Reserved == 0 && this->Version == 1 && this->Kind == 4 &&
+             this->Flags == 0);
   }
 
   // Name is the only field that's used in SYCL.

From 374a4bcc002728b1c9a91e837b12cbca8420f34b Mon Sep 17 00:00:00 2001
From: "Agarwal, Udit" <udit.agarwal@intel.com>
Date: Mon, 3 Mar 2025 17:39:35 -0800
Subject: [PATCH 09/13] Fix test failures

---
 sycl/source/detail/compiler.hpp               | 44 +++++++++----------
 sycl/source/detail/device_binary_image.cpp    | 13 +++++-
 sycl/source/detail/jit_compiler.cpp           | 11 ++---
 sycl/source/detail/jit_device_binaries.cpp    | 21 ++++-----
 sycl/source/detail/jit_device_binaries.hpp    | 34 +++++---------
 .../program_manager/program_manager.cpp       |  8 ++--
 sycl/unittests/helpers/MockDeviceImage.hpp    | 10 ++---
 .../PersistentDeviceCodeCache.cpp             | 22 ++--------
 8 files changed, 66 insertions(+), 97 deletions(-)

diff --git a/sycl/source/detail/compiler.hpp b/sycl/source/detail/compiler.hpp
index 15668cf469a2b..102c56df7b9b8 100644
--- a/sycl/source/detail/compiler.hpp
+++ b/sycl/source/detail/compiler.hpp
@@ -78,21 +78,9 @@
 
 #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_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 {
+struct _sycl_offload_entry_struct_new {
   /// 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.
@@ -111,26 +99,36 @@ struct _sycl_offload_entry_struct {
   uint64_t Data;
   /// An extra pointer, usually null.
   void *AuxAddr;
+};
+using sycl_offload_entry_new = _sycl_offload_entry_struct_new *;
+
+// Entry type, matches OpenMP for compatibility
+struct _sycl_offload_entry_struct {
+  void *addr;
+  char *name;
+  size_t size;
+  int32_t flags;
+  int32_t reserved;
+
+  bool IsNewOffloadEntryType() {
+    // Assume this is the new version of the struct.
+    auto newStruct = reinterpret_cast<sycl_offload_entry_new>(this);
 
-  bool IsLegacy() {
     // Chek if first 64 bits is 0, next 16 bits is equal to 1, next 16 bits
     // is equal to 4 (OK_SYCL), and Flags should be zero. If all these
     // conditions are met, then this is a newer version of the struct.
     // We can not just rely on checking the first 64 bits, because even for the
     // older version of the struct, the first 64 bits (void* addr) are zero.
-    return !(this->Reserved == 0 && this->Version == 1 && this->Kind == 4 &&
-             this->Flags == 0);
+    return newStruct->Reserved == 0 && newStruct->Version == 1 &&
+           newStruct->Kind == 4 && newStruct->Flags == 0;
   }
 
   // Name is the only field that's used in SYCL.
   inline char *GetName() {
-    // Check if the first 64 bits of this struct are not zero, if so, this is an
-    // older version of the struct.
-    if (IsLegacy()) {
-      // This is an older version of the struct, use the old name field.
-      return reinterpret_cast<sycl_offload_entry_legacy>(this)->name;
-    }
-    return SymbolName;
+    if (IsNewOffloadEntryType())
+      return reinterpret_cast<sycl_offload_entry_new>(this)->SymbolName;
+
+    return name;
   }
 };
 using sycl_offload_entry = _sycl_offload_entry_struct *;
diff --git a/sycl/source/detail/device_binary_image.cpp b/sycl/source/detail/device_binary_image.cpp
index c5db61ff81a64..57792d3cee74e 100644
--- a/sycl/source/detail/device_binary_image.cpp
+++ b/sycl/source/detail/device_binary_image.cpp
@@ -120,8 +120,17 @@ void RTDeviceBinaryImage::print() const {
   std::cerr << "    Link options    : "
             << (Bin->LinkOptions ? Bin->LinkOptions : "NULL") << "\n";
   std::cerr << "    Entries  : ";
-  for (sycl_offload_entry EntriesIt = Bin->EntriesBegin;
-       EntriesIt != Bin->EntriesEnd; ++EntriesIt)
+
+  sycl_offload_entry EntriesIt;
+  auto IncrementEntriesIt = [&]() {
+    if (EntriesIt->IsNewOffloadEntryType())
+      return reinterpret_cast<sycl_offload_entry>(
+          reinterpret_cast<sycl_offload_entry_new>(EntriesIt) + 1);
+    else
+      return EntriesIt + 1;
+  };
+  for (EntriesIt = Bin->EntriesBegin; EntriesIt != Bin->EntriesEnd;
+       EntriesIt = IncrementEntriesIt())
     std::cerr << EntriesIt->GetName() << " ";
   std::cerr << "\n";
   std::cerr << "    Properties [" << Bin->PropertySetsBegin << "-"
diff --git a/sycl/source/detail/jit_compiler.cpp b/sycl/source/detail/jit_compiler.cpp
index bf143c2fd1407..c155022e40896 100644
--- a/sycl/source/detail/jit_compiler.cpp
+++ b/sycl/source/detail/jit_compiler.cpp
@@ -1091,10 +1091,7 @@ 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{
-      0 /*Reserved*/, 0 /*Version*/,       0 /*Kind*/,
-      0 /*Flags*/,    nullptr /*Address*/, FusedKernelName /*Name*/,
-      0 /*Size*/,     0 /*Data*/,          nullptr /*AuxAddr*/};
+  OffloadEntryContainer Entry{FusedKernelName, nullptr, 0, 0, 0};
   Binary.addOffloadEntry(std::move(Entry));
 
   // Create a property entry for the argument usage mask for the fused kernel.
@@ -1167,10 +1164,8 @@ 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 = Prefix + Symbol.c_str();
-      OffloadEntryContainer Entry{
-          0 /*Reserved*/, 0 /*Version*/,       0 /*Kind*/,
-          0 /*Flags*/,    nullptr /*Address*/, PrefixedName /*Name*/,
-          0 /*Size*/,     0 /*Data*/,          nullptr /*AuxAddr*/};
+      OffloadEntryContainer Entry{PrefixedName, /*Addr=*/nullptr, /*Size=*/0,
+                                  /*Flags=*/0, /*Reserved=*/0};
       Binary.addOffloadEntry(std::move(Entry));
     }
 
diff --git a/sycl/source/detail/jit_device_binaries.cpp b/sycl/source/detail/jit_device_binaries.cpp
index 4be437da4c6ff..f90be2c27ec3a 100644
--- a/sycl/source/detail/jit_device_binaries.cpp
+++ b/sycl/source/detail/jit_device_binaries.cpp
@@ -14,22 +14,17 @@ namespace sycl {
 inline namespace _V1 {
 namespace detail {
 
-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);
+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);
 }
 
 _sycl_offload_entry_struct OffloadEntryContainer::getPIOffloadEntry() {
-  return _sycl_offload_entry_struct{Reserved,   Version, Kind,
-                                    EntryFlags, Address, KernelName.get(),
-                                    EntrySize,  Data,    AuxAddr};
+  return _sycl_offload_entry_struct{Address, KernelName.get(), EntrySize,
+                                    EntryFlags, EntryReserved};
 }
 
 PropertyContainer::PropertyContainer(const std::string &Name, const void *Data,
diff --git a/sycl/source/detail/jit_device_binaries.hpp b/sycl/source/detail/jit_device_binaries.hpp
index 9f53662991801..bca83839f39e8 100644
--- a/sycl/source/detail/jit_device_binaries.hpp
+++ b/sycl/source/detail/jit_device_binaries.hpp
@@ -18,14 +18,13 @@ 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(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(const std::string &Name, void *Addr, size_t Size,
+                        int32_t Flags, int32_t Reserved);
 
   OffloadEntryContainer(OffloadEntryContainer &&) = default;
   OffloadEntryContainer &operator=(OffloadEntryContainer &&) = default;
@@ -37,25 +36,12 @@ class OffloadEntryContainer {
   _sycl_offload_entry_struct getPIOffloadEntry();
 
 private:
-  // 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;
-  // 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;
+
+  void *Address;
+  size_t EntrySize;
+  int32_t EntryFlags;
+  int32_t EntryReserved;
 };
 
 /// Representation of _sycl_device_binary_property_struct for creation of JIT
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index 8f8d5cdbf6445..d2006e4ddd527 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -1868,9 +1868,9 @@ void ProgramManager::addImages(sycl_device_binaries DeviceBinary) {
 
     sycl_offload_entry EntriesIt;
     auto IncrementEntriesIt = [&]() {
-      if (EntriesIt->IsLegacy())
+      if (EntriesIt->IsNewOffloadEntryType())
         return reinterpret_cast<sycl_offload_entry>(
-            reinterpret_cast<sycl_offload_entry_legacy>(EntriesIt) + 1);
+            reinterpret_cast<sycl_offload_entry_new>(EntriesIt) + 1);
       else
         return EntriesIt + 1;
     };
@@ -2030,9 +2030,9 @@ void ProgramManager::removeImages(sycl_device_binaries DeviceBinary) {
 
     sycl_offload_entry EntriesIt;
     auto IncrementEntriesIt = [&]() {
-      if (EntriesIt->IsLegacy())
+      if (EntriesIt->IsNewOffloadEntryType())
         return reinterpret_cast<sycl_offload_entry>(
-            reinterpret_cast<sycl_offload_entry_legacy>(EntriesIt) + 1);
+            reinterpret_cast<sycl_offload_entry_new>(EntriesIt) + 1);
       else
         return EntriesIt + 1;
     };
diff --git a/sycl/unittests/helpers/MockDeviceImage.hpp b/sycl/unittests/helpers/MockDeviceImage.hpp
index 332541859df93..9c8fdef642893 100644
--- a/sycl/unittests/helpers/MockDeviceImage.hpp
+++ b/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,
-                   uint32_t Flags)
+                   int32_t Flags)
       : MName(Name), MData(std::move(Data)), MFlags(Flags) {
     updateNativeType();
   }
@@ -104,13 +104,13 @@ class MockOffloadEntry {
 private:
   void updateNativeType() {
     MNative = NativeType{
-        0 /*Reserved*/, 0 /*Version*/,       0 /*Kind*/,
-        MFlags,         nullptr /*Address*/, const_cast<char *>(MName.data()),
-        MData.size(),   0 /*Data*/,          nullptr /*AuxAddr*/};
+        const_cast<char *>(MData.data()), MName.data(), MData.size(), MFlags,
+        0 // Reserved
+    };
   }
   std::string MName;
   std::vector<char> MData;
-  uint32_t MFlags;
+  int32_t MFlags;
   NativeType MNative;
 };
 
diff --git a/sycl/unittests/kernel-and-program/PersistentDeviceCodeCache.cpp b/sycl/unittests/kernel-and-program/PersistentDeviceCodeCache.cpp
index 07612dbbb49c7..5d75b63875fda 100644
--- a/sycl/unittests/kernel-and-program/PersistentDeviceCodeCache.cpp
+++ b/sycl/unittests/kernel-and-program/PersistentDeviceCodeCache.cpp
@@ -261,15 +261,8 @@ class PersistentDeviceCodeCache
   device Dev;
   const char *EntryName = "Entry";
   _sycl_offload_entry_struct EntryStruct = {
-      /*Reserved*/ 0,
-      /*Version*/ 0,
-      /*Kind*/ 0,
-      /*Flags*/ 0,
-      /*Address*/ nullptr,
-      /*Name*/ const_cast<char *>(EntryName),
-      /*Size*/ strlen(EntryName),
-      /*Data*/ 0,
-      /*AuxAddr*/ nullptr};
+      /*addr*/ nullptr, const_cast<char *>(EntryName), strlen(EntryName),
+      /*flags*/ 0, /*reserved*/ 0};
   sycl_device_binary_struct BinStruct{/*Version*/ 1,
                                       /*Kind*/ 4,
                                       /*Format*/ GetParam(),
@@ -318,15 +311,8 @@ TEST_P(PersistentDeviceCodeCache, KeysWithNullTermSymbol) {
 TEST_P(PersistentDeviceCodeCache, MultipleImages) {
   const char *ExtraEntryName = "ExtraEntry";
   _sycl_offload_entry_struct ExtraEntryStruct = {
-      /*Reserved*/ 0,
-      /*Version*/ 0,
-      /*Kind*/ 0,
-      /*Flags*/ 0,
-      /*Address*/ nullptr,
-      /*Name*/ const_cast<char *>(ExtraEntryName),
-      /*Size*/ strlen(ExtraEntryName),
-      /*Data*/ 0,
-      /*AuxAddr*/ nullptr};
+      /*addr*/ nullptr, const_cast<char *>(ExtraEntryName),
+      strlen(ExtraEntryName), /*flags*/ 0, /*reserved*/ 0};
   sycl_device_binary_struct ExtraBinStruct{/*Version*/ 1,
                                            /*Kind*/ 4,
                                            /*Format*/ GetParam(),

From e75c615b1abe6eefa04b9a75bdd53a4425b19765 Mon Sep 17 00:00:00 2001
From: "Agarwal, Udit" <udit.agarwal@intel.com>
Date: Mon, 3 Mar 2025 19:51:53 -0800
Subject: [PATCH 10/13] Fix driver LIT failure

---
 clang/test/Driver/sycl-linker-wrapper-image.cpp | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/clang/test/Driver/sycl-linker-wrapper-image.cpp b/clang/test/Driver/sycl-linker-wrapper-image.cpp
index 8c88047e9e649..57bab9cb7800e 100644
--- a/clang/test/Driver/sycl-linker-wrapper-image.cpp
+++ b/clang/test/Driver/sycl-linker-wrapper-image.cpp
@@ -37,7 +37,7 @@ int main() {
 
 // CHECK-DAG: %_pi_device_binary_property_struct = type { ptr, ptr, i32, i64 }
 // CHECK-DAG: %_pi_device_binary_property_set_struct = type { ptr, ptr, ptr }
-// CHECK-DAG: %struct.__tgt_offload_entry = type { ptr, ptr, i64, i32, i32 }
+// CHECK-DAG: %struct.__tgt_offload_entry = type { i64, i16, i16, i32, ptr, ptr, i64, i64, ptr }
 // CHECK-DAG: %__sycl.tgt_device_image = type { i16, i8, i8, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr, ptr }
 // CHECK-DAG: %__sycl.tgt_bin_desc = type { i16, i16, ptr, ptr, ptr }
 
@@ -55,7 +55,7 @@ int main() {
 // CHECK-DAG: @__sycl_offload_prop_sets_arr.5 = internal constant [3 x %_pi_device_binary_property_set_struct] [%_pi_device_binary_property_set_struct { ptr @SYCL_PropSetName, ptr @__sycl_offload_prop_sets_arr, ptr getelementptr ([1 x %_pi_device_binary_property_struct], ptr @__sycl_offload_prop_sets_arr, i64 0, i64 1) }, %_pi_device_binary_property_set_struct { ptr @SYCL_PropSetName.3, ptr @__sycl_offload_prop_sets_arr.2, ptr getelementptr ([1 x %_pi_device_binary_property_struct], ptr @__sycl_offload_prop_sets_arr.2, i64 0, i64 1) }, %_pi_device_binary_property_set_struct { ptr @SYCL_PropSetName.4, ptr null, ptr null }]
 // CHECK-DAG: @.sycl_offloading.0.data = internal unnamed_addr constant [912 x i8]
 // CHECK-DAG: @__sycl_offload_entry_name = internal unnamed_addr constant [25 x i8] c"_ZTSZ4mainE11fake_kernel\00"
-// CHECK-DAG: @__sycl_offload_entries_arr = internal constant [1 x %struct.__tgt_offload_entry] [%struct.__tgt_offload_entry { ptr null, ptr @__sycl_offload_entry_name, i64 0, i32 0, i32 0 }]
+// CHECK-DAG: @__sycl_offload_entries_arr = internal constant [1 x %struct.__tgt_offload_entry] [%struct.__tgt_offload_entry { i64 0, i16 1, i16 4, i32 0, ptr null, ptr @__sycl_offload_entry_name, i64 0, i64 0, ptr null }]
 // CHECK-DAG: @.sycl_offloading.0.info = internal local_unnamed_addr constant [2 x i64] [i64 ptrtoint (ptr @.sycl_offloading.0.data to i64), i64 912], section ".tgtimg", align 16
 // CHECK-DAG: @llvm.used = appending global [1 x ptr] [ptr @.sycl_offloading.0.info], section "llvm.metadata"
 // CHECK-DAG: @.sycl_offloading.device_images = internal unnamed_addr constant [1 x %__sycl.tgt_device_image] [%__sycl.tgt_device_image { i16 2, i8 4, i8 0, ptr @.sycl_offloading.target.0, ptr @.sycl_offloading.opts.compile.0, ptr @.sycl_offloading.opts.link.0, ptr null, ptr null, ptr @.sycl_offloading.0.data, ptr getelementptr ([912 x i8], ptr @.sycl_offloading.0.data, i64 0, i64 912), ptr @__sycl_offload_entries_arr, ptr getelementptr ([1 x %struct.__tgt_offload_entry], ptr @__sycl_offload_entries_arr, i64 0, i64 1), ptr @__sycl_offload_prop_sets_arr.5, ptr getelementptr ([3 x %_pi_device_binary_property_set_struct], ptr @__sycl_offload_prop_sets_arr.5, i64 0, i64 3) }]

From 8d706016b4e55d792453dff255b23ee05c883348 Mon Sep 17 00:00:00 2001
From: "Agarwal, Udit" <udit.agarwal@intel.com>
Date: Wed, 5 Mar 2025 20:54:44 -0800
Subject: [PATCH 11/13] Address reviews

---
 .../program_manager/program_manager.cpp       | 20 +++++++++++--------
 1 file changed, 12 insertions(+), 8 deletions(-)

diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index d2006e4ddd527..8fb8508cd0c27 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -1866,17 +1866,19 @@ void ProgramManager::addImages(sycl_device_binaries DeviceBinary) {
 
     m_BinImg2KernelIDs[Img.get()].reset(new std::vector<kernel_id>);
 
-    sycl_offload_entry EntriesIt;
+    sycl_offload_entry EntriesIt = EntriesB;
+    // Assuming that there isn't a mix of new and old offload entry types in the
+    // image.
+    bool isNewOffloadEntryType = EntriesIt->IsNewOffloadEntryType();
     auto IncrementEntriesIt = [&]() {
-      if (EntriesIt->IsNewOffloadEntryType())
+      if (isNewOffloadEntryType)
         return reinterpret_cast<sycl_offload_entry>(
             reinterpret_cast<sycl_offload_entry_new>(EntriesIt) + 1);
       else
         return EntriesIt + 1;
     };
 
-    for (EntriesIt = EntriesB; EntriesIt != EntriesE;
-         EntriesIt = IncrementEntriesIt()) {
+    for (; EntriesIt != EntriesE; EntriesIt = IncrementEntriesIt()) {
 
       auto name = EntriesIt->GetName();
 
@@ -2028,9 +2030,12 @@ void ProgramManager::removeImages(sycl_device_binaries DeviceBinary) {
     // Acquire lock to modify maps for kernel bundles
     std::lock_guard<std::mutex> KernelIDsGuard(m_KernelIDsMutex);
 
-    sycl_offload_entry EntriesIt;
+    sycl_offload_entry EntriesIt = EntriesB;
+    // Assuming that there isn't a mix of new and old offload entry types in the
+    // image.
+    bool isNewOffloadEntryType = EntriesIt->IsNewOffloadEntryType();
     auto IncrementEntriesIt = [&]() {
-      if (EntriesIt->IsNewOffloadEntryType())
+      if (isNewOffloadEntryType)
         return reinterpret_cast<sycl_offload_entry>(
             reinterpret_cast<sycl_offload_entry_new>(EntriesIt) + 1);
       else
@@ -2038,8 +2043,7 @@ void ProgramManager::removeImages(sycl_device_binaries DeviceBinary) {
     };
 
     // Unmap the unique kernel IDs for the offload entries
-    for (EntriesIt = EntriesB; EntriesIt != EntriesE;
-         EntriesIt = IncrementEntriesIt()) {
+    for (; EntriesIt != EntriesE; EntriesIt = IncrementEntriesIt()) {
 
       // Drop entry for service kernel
       if (std::strstr(EntriesIt->GetName(), "__sycl_service_kernel__")) {

From 4365f4a875fb4a5c6f05378c04ef984e5e3a54ed Mon Sep 17 00:00:00 2001
From: Udit Kumar Agarwal <udit.agarwal@intel.com>
Date: Thu, 6 Mar 2025 07:52:55 -0800
Subject: [PATCH 12/13] Apply suggestions from code review

Co-authored-by: Sergey Semenov <sergey.semenov@intel.com>
---
 sycl/source/detail/program_manager/program_manager.cpp | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index 8fb8508cd0c27..e00900efbe1cb 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -1869,7 +1869,7 @@ void ProgramManager::addImages(sycl_device_binaries DeviceBinary) {
     sycl_offload_entry EntriesIt = EntriesB;
     // Assuming that there isn't a mix of new and old offload entry types in the
     // image.
-    bool isNewOffloadEntryType = EntriesIt->IsNewOffloadEntryType();
+    const bool isNewOffloadEntryType = EntriesIt->IsNewOffloadEntryType();
     auto IncrementEntriesIt = [&]() {
       if (isNewOffloadEntryType)
         return reinterpret_cast<sycl_offload_entry>(
@@ -2033,7 +2033,7 @@ void ProgramManager::removeImages(sycl_device_binaries DeviceBinary) {
     sycl_offload_entry EntriesIt = EntriesB;
     // Assuming that there isn't a mix of new and old offload entry types in the
     // image.
-    bool isNewOffloadEntryType = EntriesIt->IsNewOffloadEntryType();
+    const bool isNewOffloadEntryType = EntriesIt->IsNewOffloadEntryType();
     auto IncrementEntriesIt = [&]() {
       if (isNewOffloadEntryType)
         return reinterpret_cast<sycl_offload_entry>(

From 9ec31163d97262a168cb522d1acb72a33b85c8ff Mon Sep 17 00:00:00 2001
From: "Agarwal, Udit" <udit.agarwal@intel.com>
Date: Fri, 7 Mar 2025 18:10:46 -0800
Subject: [PATCH 13/13] Adress reviews

---
 sycl/source/detail/compiler.hpp               | 18 +++++++++--
 sycl/source/detail/device_binary_image.cpp    | 12 ++------
 .../program_manager/program_manager.cpp       | 30 +++----------------
 3 files changed, 21 insertions(+), 39 deletions(-)

diff --git a/sycl/source/detail/compiler.hpp b/sycl/source/detail/compiler.hpp
index 102c56df7b9b8..50d67800e446d 100644
--- a/sycl/source/detail/compiler.hpp
+++ b/sycl/source/detail/compiler.hpp
@@ -80,6 +80,8 @@
 
 // New entry type after
 // https://github.com/llvm/llvm-project/pull/124018
+// This is a replica of the EntryTy data structure in
+// llvm/include/llvm/Frontend/Offloading/Utility.h.
 struct _sycl_offload_entry_struct_new {
   /// Reserved bytes used to detect an older version of the struct, always zero.
   uint64_t Reserved;
@@ -110,12 +112,12 @@ struct _sycl_offload_entry_struct {
   int32_t flags;
   int32_t reserved;
 
-  bool IsNewOffloadEntryType() {
+  inline bool IsNewOffloadEntryType() {
     // Assume this is the new version of the struct.
     auto newStruct = reinterpret_cast<sycl_offload_entry_new>(this);
 
-    // Chek if first 64 bits is 0, next 16 bits is equal to 1, next 16 bits
-    // is equal to 4 (OK_SYCL), and Flags should be zero. If all these
+    // Check if first 64 bits is equal to 0, next 16 bits is equal to 1, next 16
+    // bits is equal to 4 (OK_SYCL), and check if Flags are zero. If all these
     // conditions are met, then this is a newer version of the struct.
     // We can not just rely on checking the first 64 bits, because even for the
     // older version of the struct, the first 64 bits (void* addr) are zero.
@@ -130,6 +132,16 @@ struct _sycl_offload_entry_struct {
 
     return name;
   }
+
+  // Increment the pointer to the next entry. A mix of old and new offload entry
+  // types is not supported.
+  inline _sycl_offload_entry_struct *Increment() {
+    if (IsNewOffloadEntryType())
+      return reinterpret_cast<_sycl_offload_entry_struct *>(
+          reinterpret_cast<sycl_offload_entry_new>(this) + 1);
+
+    return this + 1;
+  }
 };
 using sycl_offload_entry = _sycl_offload_entry_struct *;
 
diff --git a/sycl/source/detail/device_binary_image.cpp b/sycl/source/detail/device_binary_image.cpp
index 57792d3cee74e..89042b10cb48c 100644
--- a/sycl/source/detail/device_binary_image.cpp
+++ b/sycl/source/detail/device_binary_image.cpp
@@ -121,16 +121,8 @@ void RTDeviceBinaryImage::print() const {
             << (Bin->LinkOptions ? Bin->LinkOptions : "NULL") << "\n";
   std::cerr << "    Entries  : ";
 
-  sycl_offload_entry EntriesIt;
-  auto IncrementEntriesIt = [&]() {
-    if (EntriesIt->IsNewOffloadEntryType())
-      return reinterpret_cast<sycl_offload_entry>(
-          reinterpret_cast<sycl_offload_entry_new>(EntriesIt) + 1);
-    else
-      return EntriesIt + 1;
-  };
-  for (EntriesIt = Bin->EntriesBegin; EntriesIt != Bin->EntriesEnd;
-       EntriesIt = IncrementEntriesIt())
+  for (sycl_offload_entry EntriesIt = Bin->EntriesBegin;
+       EntriesIt != Bin->EntriesEnd; EntriesIt = EntriesIt->Increment())
     std::cerr << EntriesIt->GetName() << " ";
   std::cerr << "\n";
   std::cerr << "    Properties [" << Bin->PropertySetsBegin << "-"
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index e00900efbe1cb..8148b413f0ecc 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -1866,19 +1866,8 @@ void ProgramManager::addImages(sycl_device_binaries DeviceBinary) {
 
     m_BinImg2KernelIDs[Img.get()].reset(new std::vector<kernel_id>);
 
-    sycl_offload_entry EntriesIt = EntriesB;
-    // Assuming that there isn't a mix of new and old offload entry types in the
-    // image.
-    const bool isNewOffloadEntryType = EntriesIt->IsNewOffloadEntryType();
-    auto IncrementEntriesIt = [&]() {
-      if (isNewOffloadEntryType)
-        return reinterpret_cast<sycl_offload_entry>(
-            reinterpret_cast<sycl_offload_entry_new>(EntriesIt) + 1);
-      else
-        return EntriesIt + 1;
-    };
-
-    for (; EntriesIt != EntriesE; EntriesIt = IncrementEntriesIt()) {
+    for (sycl_offload_entry EntriesIt = EntriesB; EntriesIt != EntriesE;
+         EntriesIt = EntriesIt->Increment()) {
 
       auto name = EntriesIt->GetName();
 
@@ -2030,20 +2019,9 @@ void ProgramManager::removeImages(sycl_device_binaries DeviceBinary) {
     // Acquire lock to modify maps for kernel bundles
     std::lock_guard<std::mutex> KernelIDsGuard(m_KernelIDsMutex);
 
-    sycl_offload_entry EntriesIt = EntriesB;
-    // Assuming that there isn't a mix of new and old offload entry types in the
-    // image.
-    const bool isNewOffloadEntryType = EntriesIt->IsNewOffloadEntryType();
-    auto IncrementEntriesIt = [&]() {
-      if (isNewOffloadEntryType)
-        return reinterpret_cast<sycl_offload_entry>(
-            reinterpret_cast<sycl_offload_entry_new>(EntriesIt) + 1);
-      else
-        return EntriesIt + 1;
-    };
-
     // Unmap the unique kernel IDs for the offload entries
-    for (; EntriesIt != EntriesE; EntriesIt = IncrementEntriesIt()) {
+    for (sycl_offload_entry EntriesIt = EntriesB; EntriesIt != EntriesE;
+         EntriesIt = EntriesIt->Increment()) {
 
       // Drop entry for service kernel
       if (std::strstr(EntriesIt->GetName(), "__sycl_service_kernel__")) {