[SYCL] Support sycl::kernel_bundle for multi-device scenario (#15546)

This PR includes:
* Changes in the program manager methods to be able to properly
create/build UR program for multiple devices. So far, we were mostly
using only the first device in the vector to create/build UR program
which made UR program unusable on other devices.

* UR tag update brings the version of urProgramCreateWithBinary which
allows to create UR program from multiple device binaries.

* Our program cache key allowed only a single device. I have changed it
to contain a set of devices. If UR program is created and built for a
set of devices then the same UR program is usable whenver we have any
subset of this set. That's why if we have a program built for a set of
devices then add all subsets to the cache. Before we were adding a
record to the cache for each device from the set which is incorrect. For
example, if someone requests a UR program for {dev2, dev3} from the
cache then it is expected that this UR progam must be usable to submit a
kernel to dev3. But we could get a program for {dev1, dev2} from the
cache which is unusable on dev3.

Author: againull

sycl/cmake/modules/UnifiedRuntimeTag.cmake

@@ -1,7 +1,6 @@
-# commit c742ca49efb12380a35b8b0b467e6577ab8174ce
-# Merge: 3a8bf2c5 504d3b63
-# Author: Kenneth Benzie (Benie) <[email protected]>
-# Date:   Mon Oct 21 11:55:23 2024 +0100
-#     Merge pull request #2131 from Bensuo/ben/command-handle-fix
-#     [EXP][CMDBUF] Make command handle behaviour consistent
-set(UNIFIED_RUNTIME_TAG c742ca49efb12380a35b8b0b467e6577ab8174ce)
+# commit 3db3a5e2d935630f2ffddd93a72ae0aa9af89acb
+# Author: Artur Gainullin <[email protected]>
+# Date:   Tue Oct 22 03:30:08 2024 -0700
+#
+#     Support UR program creation from multiple device binaries (#2147)
+set(UNIFIED_RUNTIME_TAG 3db3a5e2d935630f2ffddd93a72ae0aa9af89acb)

sycl/source/detail/context_impl.cpp

@@ -21,6 +21,7 @@
 #include <sycl/property_list.hpp>
 
 #include <algorithm>
+#include <set>
 
 namespace sycl {
 inline namespace _V1 {
@@ -492,7 +493,7 @@ std::optional<ur_program_handle_t> context_impl::getProgramForDevImgs(
     auto &Cache = LockedCache.get().Cache;
     ur_device_handle_t &DevHandle = getSyclObjImpl(Device)->getHandleRef();
     for (std::uintptr_t ImageIDs : ImgIdentifiers) {
-      auto OuterKey = std::make_pair(ImageIDs, DevHandle);
+      auto OuterKey = std::make_pair(ImageIDs, std::set{DevHandle});
       size_t NProgs = KeyMap.count(OuterKey);
       if (NProgs == 0)
         continue;

sycl/source/detail/helpers.cpp

@@ -65,8 +65,8 @@ retrieveKernelBinary(const QueueImplPtr &Queue, const char *KernelName,
     auto DeviceImpl = Queue->getDeviceImplPtr();
     auto Device = detail::createSyclObjFromImpl<device>(DeviceImpl);
     ur_program_handle_t Program =
-        detail::ProgramManager::getInstance().createURProgram(**DeviceImage,
-                                                              Context, Device);
+        detail::ProgramManager::getInstance().createURProgram(
+            **DeviceImage, Context, {Device});
     return {*DeviceImage, Program};
   }
 
@@ -94,7 +94,7 @@ retrieveKernelBinary(const QueueImplPtr &Queue, const char *KernelName,
     DeviceImage = &detail::ProgramManager::getInstance().getDeviceImage(
         KernelName, Context, Device);
     Program = detail::ProgramManager::getInstance().createURProgram(
-        *DeviceImage, Context, Device);
+        *DeviceImage, Context, {Device});
   }
   return {DeviceImage, Program};
 }

sycl/source/detail/kernel_program_cache.hpp

@@ -20,6 +20,7 @@
 #include <atomic>
 #include <condition_variable>
 #include <mutex>
+#include <set>
 #include <type_traits>
 
 #include <boost/unordered/unordered_flat_map.hpp>
@@ -118,9 +119,10 @@ class KernelProgramCache {
    * when debugging environment variables are set and we can just ignore them
    * since all kernels will have their build options overridden with the same
    * string*/
-  using ProgramCacheKeyT =
-      std::pair<std::pair<SerializedObj, std::uintptr_t>, ur_device_handle_t>;
-  using CommonProgramKeyT = std::pair<std::uintptr_t, ur_device_handle_t>;
+  using ProgramCacheKeyT = std::pair<std::pair<SerializedObj, std::uintptr_t>,
+                                     std::set<ur_device_handle_t>>;
+  using CommonProgramKeyT =
+      std::pair<std::uintptr_t, std::set<ur_device_handle_t>>;
 
   struct ProgramCache {
     ::boost::unordered_map<ProgramCacheKeyT, ProgramBuildResultPtr> Cache;