[SYCL] Remove kernel set id to fix kernel lookup

sycl/include/sycl/detail/common.hpp

@@ -369,15 +369,6 @@ size_t getLinearIndex(const T<Dims> &Index, const U<Dims> &Range) {
   return LinearIndex;
 }
 
-// Kernel set ID, used to group kernels (represented by OSModule & kernel name
-// pairs) into disjoint sets based on the kernel distribution among device
-// images.
-using KernelSetId = size_t;
-// Kernel set ID for kernels contained within the SPIR-V file specified via
-// environment.
-constexpr KernelSetId SpvFileKSId = 0;
-constexpr KernelSetId LastKSId = SpvFileKSId;
-
 template <typename T> struct InlineVariableHelper {
   static constexpr T value{};
 };

sycl/source/detail/device_global_map_entry.hpp

@@ -14,6 +14,7 @@
 #include <mutex>
 #include <optional>
 #include <set>
+#include <unordered_set>
 
 #include <detail/pi_utils.hpp>
 #include <sycl/detail/defines_elementary.hpp>
@@ -51,11 +52,11 @@ struct DeviceGlobalMapEntry {
   std::string MUniqueId;
   // Pointer to the device_global on host.
   const void *MDeviceGlobalPtr = nullptr;
+  // Images device_global are used by.
+  std::unordered_set<RTDeviceBinaryImage *> MImages;
   // The image identifiers for the images using the device_global used by in the
   // cache.
   std::set<std::uintptr_t> MImageIdentifiers;
-  // The kernel-set IDs for the images using the device_global.
-  std::set<KernelSetId> MKSIds;
   // Size of the underlying type in the device_global.
   std::uint32_t MDeviceGlobalTSize = 0;
   // True if the device_global has been decorated with device_image_scope.
@@ -68,10 +69,11 @@ struct DeviceGlobalMapEntry {
 
   // Constructor for only initializing ID, type size, and device image scope
   // flag. The pointer to the device global will be initialized later.
-  DeviceGlobalMapEntry(std::string UniqueId, std::uintptr_t ImgId,
-                       KernelSetId KSId, std::uint32_t DeviceGlobalTSize,
+  DeviceGlobalMapEntry(std::string UniqueId, RTDeviceBinaryImage *Img,
+                       std::uint32_t DeviceGlobalTSize,
                        bool IsDeviceImageScopeDecorated)
-      : MUniqueId(UniqueId), MImageIdentifiers{ImgId}, MKSIds{KSId},
+      : MUniqueId(UniqueId), MImages{Img},
+        MImageIdentifiers{reinterpret_cast<uintptr_t>(Img)},
         MDeviceGlobalTSize(DeviceGlobalTSize),
         MIsDeviceImageScopeDecorated(IsDeviceImageScopeDecorated) {}
 
@@ -85,8 +87,7 @@ struct DeviceGlobalMapEntry {
 
   // Initialize the device_global's element type size and the flag signalling
   // if the device_global has the device_image_scope property.
-  void initialize(std::uintptr_t ImgId, KernelSetId KSId,
-                  std::uint32_t DeviceGlobalTSize,
+  void initialize(RTDeviceBinaryImage *Img, std::uint32_t DeviceGlobalTSize,
                   bool IsDeviceImageScopeDecorated) {
     if (MDeviceGlobalTSize != 0) {
       // The device global entry has already been initialized. This can happen
@@ -99,8 +100,8 @@ struct DeviceGlobalMapEntry {
           "Device global intializations disagree on image scope decoration.");
       return;
     }
-    MImageIdentifiers.insert(ImgId);
-    MKSIds.insert(KSId);
+    MImages.insert(Img);
+    MImageIdentifiers.insert(reinterpret_cast<uintptr_t>(Img));
     MDeviceGlobalTSize = DeviceGlobalTSize;
     MIsDeviceImageScopeDecorated = IsDeviceImageScopeDecorated;
   }

sycl/source/detail/memory_manager.cpp

@@ -1133,13 +1133,13 @@ getOrBuildProgramForDeviceGlobal(QueueImplPtr Queue,
   assert(DeviceGlobalEntry->MIsDeviceImageScopeDecorated &&
          "device_global is not device image scope decorated.");
 
-  // If the device global is used in multiple kernel sets we cannot proceed.
-  if (DeviceGlobalEntry->MKSIds.size() > 1)
+  // If the device global is used in multiple device images we cannot proceed.
+  if (DeviceGlobalEntry->MImageIdentifiers.size() > 1)
     throw sycl::exception(make_error_code(errc::invalid),
                           "More than one image exists with the device_global.");
 
   // If there are no kernels using the device_global we cannot proceed.
-  if (DeviceGlobalEntry->MKSIds.size() == 0)
+  if (DeviceGlobalEntry->MImageIdentifiers.size() == 0)
     throw sycl::exception(make_error_code(errc::invalid),
                           "No image exists with the device_global.");
 
@@ -1153,9 +1153,9 @@ getOrBuildProgramForDeviceGlobal(QueueImplPtr Queue,
 
   // If there was no cached program, build one.
   auto Context = createSyclObjFromImpl<context>(ContextImpl);
-  KernelSetId KSId = *DeviceGlobalEntry->MKSIds.begin();
   ProgramManager &PM = ProgramManager::getInstance();
-  RTDeviceBinaryImage &Img = PM.getDeviceImage(KSId, Context, Device);
+  RTDeviceBinaryImage &Img =
+      PM.getDeviceImage(DeviceGlobalEntry->MImages, Context, Device);
   device_image_plain DeviceImage =
       PM.getDeviceImageFromBinaryImage(&Img, Context, Device);
   device_image_plain BuiltImage = PM.build(DeviceImage, {Device}, {});