[NFCI][SYCL] Change device_image_impl::MDevices to store raw device_impl * (#19459)

#18251 extended `device_impl`s'
lifetimes until shutdown and #18270
started to pass devices as raw pointers in some of the APIs.

This PR builds on top of that and extends usage of raw
pointers/references/`device_range` as the devices are known to be alive
and extra `std::shared_ptr`'s atomic increments aren't necessary and
could be avoided.

Since we change the type of `device_image_impl::MDevices`, other APIs in
that class and in `program_manager` don't need to operate in terms of
`sycl::device` or `std::shared_ptr<device_impl>` and we can switch them
to use `devices_range` instead. A small number of other modifications
are caused by these APIs' changes and are necessary to keep the code
buildable.

One extra change is the addition of a minor
`devices_range::to<std::vector<ur_device_handle_t>>()` helper that we
can use now that most of the arguments are `device_range`. Technically,
could go in another PR but then we'd just be modifying the exact same
lines two times, so I decided to fuse it here.

Author: aelovikov-intel

sycl/source/detail/context_impl.cpp

@@ -326,13 +326,13 @@ void context_impl::removeAssociatedDeviceGlobal(const void *DeviceGlobalPtr) {
 }
 
 void context_impl::addDeviceGlobalInitializer(
-    ur_program_handle_t Program, const std::vector<device> &Devs,
+    ur_program_handle_t Program, devices_range Devs,
     const RTDeviceBinaryImage *BinImage) {
   if (BinImage->getDeviceGlobals().empty())
     return;
   std::lock_guard<std::mutex> Lock(MDeviceGlobalInitializersMutex);
-  for (const device &Dev : Devs) {
-    auto Key = std::make_pair(Program, getSyclObjImpl(Dev)->getHandleRef());
+  for (device_impl &Dev : Devs) {
+    auto Key = std::make_pair(Program, Dev.getHandleRef());
     auto [Iter, Inserted] = MDeviceGlobalInitializers.emplace(Key, BinImage);
     if (Inserted && !Iter->second.MDeviceGlobalsFullyInitialized)
       ++MDeviceGlobalNotInitializedCnt;

sycl/source/detail/context_impl.hpp

@@ -216,7 +216,7 @@ class context_impl : public std::enable_shared_from_this<context_impl> {
 
   /// Adds a device global initializer.
   void addDeviceGlobalInitializer(ur_program_handle_t Program,
-                                  const std::vector<device> &Devs,
+                                  devices_range Devs,
                                   const RTDeviceBinaryImage *BinImage);
 
   /// Initializes device globals for a program on the associated queue.

sycl/source/detail/device_image_impl.hpp

@@ -2297,16 +2297,31 @@ struct devices_deref_impl {
 using devices_iterator =
     variadic_iterator<devices_deref_impl, device,
                       std::vector<std::shared_ptr<device_impl>>::const_iterator,
-                      std::vector<device>::const_iterator, device_impl *>;
+                      std::vector<device>::const_iterator,
+                      std::vector<device_impl *>::const_iterator,
+                      device_impl *>;
 
 class devices_range : public iterator_range<devices_iterator> {
 private:
   using Base = iterator_range<devices_iterator>;
 
 public:
   using Base::Base;
-  devices_range(const device &Dev)
-      : devices_range(&*getSyclObjImpl(Dev), (&*getSyclObjImpl(Dev) + 1), 1) {}
+  template <typename Container>
+  decltype(std::declval<Base>().to<Container>()) to() const {
+    return this->Base::to<Container>();
+  }
+
+  template <typename Container>
+  std::enable_if_t<std::is_same_v<Container, std::vector<ur_device_handle_t>>,
+                   Container>
+  to() const {
+    std::vector<ur_device_handle_t> DeviceHandles;
+    DeviceHandles.reserve(size());
+    std::transform(begin(), end(), std::back_inserter(DeviceHandles),
+                   [](device_impl &Dev) { return Dev.getHandleRef(); });
+    return DeviceHandles;
+  }
 };
 
 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES

sycl/source/detail/device_impl.hpp

@@ -110,6 +110,11 @@ template <typename iterator> class iterator_range {
   iterator_range(const ContainerTy &Container)
       : iterator_range(Container.begin(), Container.end(), Container.size()) {}
 
+  iterator_range(value_type &Obj) : iterator_range(&Obj, &Obj + 1, 1) {}
+
+  iterator_range(const sycl_type &Obj)
+      : iterator_range(&*getSyclObjImpl(Obj), (&*getSyclObjImpl(Obj) + 1), 1) {}
+
   iterator begin() const { return Begin; }
   iterator end() const { return End; }
   size_t size() const { return Size; }

sycl/source/detail/helpers.hpp

@@ -1138,6 +1138,29 @@ class kernel_bundle_impl
   DeviceGlobalMap MDeviceGlobals{/*OwnerControlledCleanup=*/false};
 };
 
+inline bool is_compatible(const std::vector<kernel_id> &KernelIDs,
+                          device_impl &Dev) {
+  if (KernelIDs.empty())
+    return true;
+  // One kernel may be contained in several binary images depending on the
+  // number of targets. This kernel is compatible with the device if there is
+  // at least one image (containing this kernel) whose aspects are supported by
+  // the device and whose target matches the device.
+  for (const auto &KernelID : KernelIDs) {
+    std::set<const detail::RTDeviceBinaryImage *> BinImages =
+        detail::ProgramManager::getInstance().getRawDeviceImages({KernelID});
+
+    if (std::none_of(BinImages.begin(), BinImages.end(),
+                     [&](const detail::RTDeviceBinaryImage *Img) {
+                       return doesDevSupportDeviceRequirements(Dev, *Img) &&
+                              doesImageTargetMatchDevice(*Img, Dev);
+                     }))
+      return false;
+  }
+
+  return true;
+}
+
 } // namespace detail
 } // namespace _V1
 } // namespace sycl