[SYCL] Check if kernel_bundle contains compatible kernels (#7691)

According to the SYCL2020 get_kernel_bundle() and
has_kernel_bundle() should check that kernels are compatible
with device.

sycl/source/detail/global_handler.cpp

@@ -148,7 +148,7 @@ ThreadPool &GlobalHandler::getHostTaskThreadPool() {
   return TP;
 }
 
-void releaseDefaultContexts() {
+void GlobalHandler::releaseDefaultContexts() {
   // Release shared-pointers to SYCL objects.
 #ifndef _WIN32
   GlobalHandler::instance().MPlatformToDefaultContextCache.Inst.reset(nullptr);
@@ -163,7 +163,9 @@ void releaseDefaultContexts() {
 }
 
 struct DefaultContextReleaseHandler {
-  ~DefaultContextReleaseHandler() { releaseDefaultContexts(); }
+  ~DefaultContextReleaseHandler() {
+    GlobalHandler::instance().releaseDefaultContexts();
+  }
 };
 
 void GlobalHandler::registerDefaultContextReleaseHandler() {
@@ -210,7 +212,7 @@ void shutdown() {
   // prior to closing the plugins.
   // Note: Releasing a default context here may cause failures in plugins with
   // global state as the global state may have been released.
-  releaseDefaultContexts();
+  GlobalHandler::instance().releaseDefaultContexts();
 
   // First, release resources, that may access plugins.
   GlobalHandler::instance().MPlatformCache.Inst.reset(nullptr);

sycl/source/detail/global_handler.hpp

@@ -75,13 +75,13 @@ class GlobalHandler {
   static void registerDefaultContextReleaseHandler();
 
   void unloadPlugins();
+  void releaseDefaultContexts();
   void drainThreadPool();
 
   // For testing purposes only
   void attachScheduler(Scheduler *Scheduler);
 
 private:
-  friend void releaseDefaultContexts();
   friend void shutdown();
 
   // Constructor and destructor are declared out-of-line to allow incomplete

sycl/source/detail/kernel_bundle_impl.hpp

@@ -236,8 +236,6 @@ class kernel_bundle_impl {
                      bundle_state State)
       : MContext(std::move(Ctx)), MDevices(std::move(Devs)), MState(State) {
 
-    // TODO: Add a check that all kernel ids are compatible with at least one
-    // device in Devs
     common_ctor_checks(State);
 
     MDeviceImages = detail::ProgramManager::getInstance().getSYCLDeviceImages(

sycl/source/detail/program_manager/program_manager.cpp

@@ -1605,15 +1605,16 @@ void ProgramManager::addOrInitDeviceGlobalEntry(const void *DeviceGlobalPtr,
   m_Ptr2DeviceGlobal.insert({DeviceGlobalPtr, NewEntry.first->second.get()});
 }
 
-void ProgramManager::getRawDeviceImages(
-    const std::vector<kernel_id> &KernelIDs,
-    std::set<RTDeviceBinaryImage *> &BinImages) {
+std::set<RTDeviceBinaryImage *>
+ProgramManager::getRawDeviceImages(const std::vector<kernel_id> &KernelIDs) {
+  std::set<RTDeviceBinaryImage *> BinImages;
   std::lock_guard<std::mutex> KernelIDsGuard(m_KernelIDsMutex);
   for (const kernel_id &KID : KernelIDs) {
     auto Range = m_KernelIDs2BinImage.equal_range(KID);
     for (auto It = Range.first, End = Range.second; It != End; ++It)
       BinImages.insert(It->second);
   }
+  return BinImages;
 }
 
 std::vector<device_image_plain>
@@ -1625,7 +1626,17 @@ ProgramManager::getSYCLDeviceImagesWithCompatibleState(
   // TODO: Can we avoid repacking?
   std::set<RTDeviceBinaryImage *> BinImages;
   if (!KernelIDs.empty()) {
-    getRawDeviceImages(KernelIDs, BinImages);
+    for (const auto &KID : KernelIDs) {
+      bool isCompatibleWithAtLeastOneDev =
+          std::any_of(Devs.begin(), Devs.end(), [&KID](const auto &Dev) {
+            return sycl::is_compatible({KID}, Dev);
+          });
+      if (!isCompatibleWithAtLeastOneDev)
+        throw sycl::exception(
+            make_error_code(errc::invalid),
+            "Kernel is incompatible with all devices in devs");
+    }
+    BinImages = getRawDeviceImages(KernelIDs);
   } else {
     std::lock_guard<std::mutex> Guard(Sync::getGlobalLock());
     for (auto &ImagesSets : m_DeviceImages) {
@@ -1653,7 +1664,8 @@ ProgramManager::getSYCLDeviceImagesWithCompatibleState(
       continue;
 
     for (const sycl::device &Dev : Devs) {
-      if (!compatibleWithDevice(BinImage, Dev))
+      if (!compatibleWithDevice(BinImage, Dev) ||
+          !doesDevSupportImgAspects(Dev, *BinImage))
         continue;
 
       std::shared_ptr<std::vector<sycl::kernel_id>> KernelIDs;
@@ -1732,7 +1744,7 @@ ProgramManager::getSYCLDeviceImages(const context &Ctx,
   // Collect device images with compatible state
   std::vector<device_image_plain> DeviceImages =
       getSYCLDeviceImagesWithCompatibleState(Ctx, Devs, TargetState);
-  // Brind device images with compatible state to desired state
+  // Bring device images with compatible state to desired state.
   bringSYCLDeviceImagesToState(DeviceImages, TargetState);
   return DeviceImages;
 }
@@ -1779,7 +1791,7 @@ std::vector<device_image_plain> ProgramManager::getSYCLDeviceImages(
   std::vector<device_image_plain> DeviceImages =
       getSYCLDeviceImagesWithCompatibleState(Ctx, Devs, TargetState, KernelIDs);
 
-  // Brind device images with compatible state to desired state
+  // Bring device images with compatible state to desired state.
   bringSYCLDeviceImagesToState(DeviceImages, TargetState);
   return DeviceImages;
 }
@@ -2118,6 +2130,29 @@ std::pair<RT::PiKernel, std::mutex *> ProgramManager::getOrCreateKernel(
                         &(BuildResult->MBuildResultMutex));
 }
 
+bool doesDevSupportImgAspects(const device &Dev,
+                              const RTDeviceBinaryImage &Img) {
+  const RTDeviceBinaryImage::PropertyRange &PropRange =
+      Img.getDeviceRequirements();
+  RTDeviceBinaryImage::PropertyRange::ConstIterator PropIt = std::find_if(
+      PropRange.begin(), PropRange.end(),
+      [](RTDeviceBinaryImage::PropertyRange::ConstIterator &&Prop) {
+        using namespace std::literals;
+        return (*Prop)->Name == "aspects"sv;
+      });
+  if (PropIt == PropRange.end())
+    return true;
+  ByteArray Aspects = DeviceBinaryProperty(*PropIt).asByteArray();
+  // Drop 8 bytes describing the size of the byte array.
+  Aspects.dropBytes(8);
+  while (!Aspects.empty()) {
+    aspect Aspect = Aspects.consume<aspect>();
+    if (!Dev.has(Aspect))
+      return false;
+  }
+  return true;
+}
+
 } // namespace detail
 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
 } // namespace sycl

sycl/source/detail/program_manager/program_manager.hpp

@@ -45,6 +45,9 @@ __SYCL_INLINE_VER_NAMESPACE(_V1) {
 class context;
 namespace detail {
 
+bool doesDevSupportImgAspects(const device &Dev,
+                              const RTDeviceBinaryImage &BinImages);
+
 // This value must be the same as in libdevice/device_itt.h.
 // See sycl/doc/design/ITTAnnotations.md for more info.
 static constexpr uint32_t inline ITTSpecConstId = 0xFF747469;
@@ -255,8 +258,8 @@ class ProgramManager {
 
   bool kernelUsesAssert(OSModuleHandle M, const std::string &KernelName) const;
 
-  void getRawDeviceImages(const std::vector<kernel_id> &KernelIDs,
-                          std::set<RTDeviceBinaryImage *> &BinImages);
+  std::set<RTDeviceBinaryImage *>
+  getRawDeviceImages(const std::vector<kernel_id> &KernelIDs);
 
 private:
   ProgramManager(ProgramManager const &) = delete;

sycl/source/kernel_bundle.cpp

@@ -184,9 +184,6 @@ bool has_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
       detail::ProgramManager::getInstance()
           .getSYCLDeviceImagesWithCompatibleState(Ctx, Devs, State);
 
-  // TODO: Add a check that all kernel ids are compatible with at least one
-  // device in Devs
-
   return (bool)DeviceImages.size();
 }
 
@@ -236,9 +233,6 @@ bool has_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
                     return CombinedKernelIDs.count(KernelID);
                   });
 
-  // TODO: Add a check that all kernel ids are compatible with at least one
-  // device in Devs
-
   return AllKernelIDsRepresented;
 }
 
@@ -296,28 +290,12 @@ std::vector<kernel_id> get_kernel_ids() {
 }
 
 bool is_compatible(const std::vector<kernel_id> &KernelIDs, const device &Dev) {
-  using namespace detail;
-  std::set<RTDeviceBinaryImage *> BinImages;
-  ProgramManager::getInstance().getRawDeviceImages(KernelIDs, BinImages);
-  for (RTDeviceBinaryImage *Img : BinImages) {
-    const RTDeviceBinaryImage::PropertyRange &PropRange =
-        Img->getDeviceRequirements();
-    for (RTDeviceBinaryImage::PropertyRange::ConstIterator It : PropRange) {
-      using namespace std::literals;
-      if ((*It)->Name != "aspects"sv)
-        continue;
-      ByteArray Aspects = DeviceBinaryProperty(*It).asByteArray();
-      // Drop 8 bytes describing the size of the byte array
-      Aspects.dropBytes(8);
-      while (!Aspects.empty()) {
-        aspect Aspect = Aspects.consume<aspect>();
-        if (!Dev.has(Aspect))
-          return false;
-      }
-    }
-  }
-
-  return true;
+  std::set<detail::RTDeviceBinaryImage *> BinImages =
+      detail::ProgramManager::getInstance().getRawDeviceImages(KernelIDs);
+  return std::all_of(BinImages.begin(), BinImages.end(),
+                     [&Dev](const detail::RTDeviceBinaryImage *Img) {
+                       return doesDevSupportImgAspects(Dev, *Img);
+                     });
 }
 
 } // __SYCL_INLINE_VER_NAMESPACE(_V1)

sycl/unittests/SYCL2020/KernelBundle.cpp

@@ -17,6 +17,7 @@
 
 class TestKernel;
 class TestKernelExeOnly;
+class TestKernelWithAspects;
 
 namespace sycl {
 __SYCL_INLINE_VER_NAMESPACE(_V1) {
@@ -47,17 +48,33 @@ template <> struct KernelInfo<TestKernelExeOnly> {
   static constexpr int64_t getKernelSize() { return 1; }
 };
 
+template <> struct KernelInfo<TestKernelWithAspects> {
+  static constexpr unsigned getNumParams() { return 0; }
+  static const kernel_param_desc_t &getParamDesc(int) {
+    static kernel_param_desc_t Dummy;
+    return Dummy;
+  }
+  static constexpr const char *getName() { return "TestKernelWithAspects"; }
+  static constexpr bool isESIMD() { return false; }
+  static constexpr bool callsThisItem() { return false; }
+  static constexpr bool callsAnyThisFreeFunction() { return false; }
+  static constexpr int64_t getKernelSize() { return 1; }
+};
+
 } // namespace detail
 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
 } // namespace sycl
 
 static sycl::unittest::PiImage
 generateDefaultImage(std::initializer_list<std::string> KernelNames,
                      pi_device_binary_type BinaryType,
-                     const char *DeviceTargetSpec) {
+                     const char *DeviceTargetSpec,
+                     const std::vector<sycl::aspect> &Aspects = {}) {
   using namespace sycl::unittest;
 
   PiPropertySet PropSet;
+  if (!Aspects.empty())
+    addAspects(PropSet, Aspects);
 
   std::vector<unsigned char> Bin{0, 1, 2, 3, 4, 5}; // Random data
 
@@ -74,16 +91,30 @@ generateDefaultImage(std::initializer_list<std::string> KernelNames,
   return Img;
 }
 
-static sycl::unittest::PiImage Imgs[3] = {
+static sycl::unittest::PiImage Imgs[] = {
     generateDefaultImage({"TestKernel"}, PI_DEVICE_BINARY_TYPE_SPIRV,
                          __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64),
     generateDefaultImage({"TestKernelExeOnly"}, PI_DEVICE_BINARY_TYPE_NATIVE,
                          __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64_X86_64),
     // A device image without entires
-    generateDefaultImage({},
-                         PI_DEVICE_BINARY_TYPE_NATIVE,
-                         __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64_X86_64)};
-static sycl::unittest::PiImageArray<3> ImgArray{Imgs};
+    generateDefaultImage({}, PI_DEVICE_BINARY_TYPE_NATIVE,
+                         __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64_X86_64),
+    generateDefaultImage(
+        {"TestKernelWithAspects"}, PI_DEVICE_BINARY_TYPE_NATIVE,
+        __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64, {sycl::aspect::gpu})};
+static sycl::unittest::PiImageArray<std::size(Imgs)> ImgArray{Imgs};
+
+static pi_result redefinedDeviceGetInfoCPU(pi_device device,
+                                           pi_device_info param_name,
+                                           size_t param_value_size,
+                                           void *param_value,
+                                           size_t *param_value_size_ret) {
+  if (param_name == PI_DEVICE_INFO_TYPE) {
+    auto *Result = reinterpret_cast<_pi_device_type *>(param_value);
+    *Result = PI_DEVICE_TYPE_CPU;
+  }
+  return PI_SUCCESS;
+}
 
 TEST(KernelBundle, GetKernelBundleFromKernel) {
   sycl::unittest::PiMock Mock;
@@ -537,3 +568,94 @@ TEST(KernelBundle, DescendentDevice) {
 
   EXPECT_EQ(KernelBundle, RetKernelBundle);
 }
+
+TEST(KernelBundle, CheckIfBundleHasIncompatibleKernel) {
+  sycl::unittest::PiMock Mock;
+  // TestKernelWithAspects has GPU aspect, so it shouldn't be compatible with
+  // the CPU device and hence shouldn't be in the kernel bundle.
+  Mock.redefineAfter<sycl::detail::PiApiKind::piDeviceGetInfo>(
+      redefinedDeviceGetInfoCPU);
+  sycl::platform Plt = Mock.getPlatform();
+  const sycl::device Dev = Plt.get_devices()[0];
+  EXPECT_TRUE(Dev.is_cpu());
+
+  auto Bundle = sycl::get_kernel_bundle<sycl::bundle_state::executable>(
+      sycl::context(Dev), {Dev});
+  auto KernelId1 = sycl::get_kernel_id<TestKernelWithAspects>();
+  auto KernelId2 = sycl::get_kernel_id<TestKernel>();
+
+  EXPECT_FALSE(Bundle.has_kernel(KernelId1));
+  EXPECT_TRUE(Bundle.has_kernel(KernelId2));
+}
+
+TEST(KernelBundle, CheckIfBundleHasCompatibleKernel) {
+  sycl::unittest::PiMock Mock;
+  sycl::platform Plt = Mock.getPlatform();
+  // GPU by default.
+  const sycl::device Dev = Plt.get_devices()[0];
+  EXPECT_TRUE(Dev.is_gpu());
+
+  auto Bundle = sycl::get_kernel_bundle<sycl::bundle_state::executable>(
+      sycl::context(Dev), {Dev});
+  auto KernelId1 = sycl::get_kernel_id<TestKernelWithAspects>();
+  auto KernelId2 = sycl::get_kernel_id<TestKernel>();
+
+  EXPECT_TRUE(Bundle.has_kernel(KernelId1));
+  EXPECT_TRUE(Bundle.has_kernel(KernelId2));
+}
+
+TEST(KernelBundle, CheckIfIncompatibleBundleExists) {
+  sycl::unittest::PiMock Mock;
+  // TestKernelWithAspects has GPU aspect, so it shouldn't be compatible with
+  // the CPU device and hence shouldn't be in the kernel bundle.
+  Mock.redefineAfter<sycl::detail::PiApiKind::piDeviceGetInfo>(
+      redefinedDeviceGetInfoCPU);
+  sycl::platform Plt = Mock.getPlatform();
+  const sycl::device Dev = Plt.get_devices()[0];
+  EXPECT_TRUE(Dev.is_cpu());
+
+  auto KernelId1 = sycl::get_kernel_id<TestKernelWithAspects>();
+  auto KernelId2 = sycl::get_kernel_id<TestKernel>();
+
+  EXPECT_FALSE(sycl::has_kernel_bundle<sycl::bundle_state::executable>(
+      sycl::context(Dev), {KernelId1, KernelId2}));
+  EXPECT_FALSE(sycl::has_kernel_bundle<sycl::bundle_state::executable>(
+      sycl::context(Dev), {KernelId1}));
+  EXPECT_TRUE(sycl::has_kernel_bundle<sycl::bundle_state::executable>(
+      sycl::context(Dev), {KernelId2}));
+}
+
+TEST(KernelBundle, CheckIfCompatibleBundleExists2) {
+  sycl::unittest::PiMock Mock;
+  sycl::platform Plt = Mock.getPlatform();
+  // GPU by default.
+  const sycl::device Dev = Plt.get_devices()[0];
+  EXPECT_TRUE(Dev.is_gpu());
+
+  auto KernelId1 = sycl::get_kernel_id<TestKernelWithAspects>();
+  auto KernelId2 = sycl::get_kernel_id<TestKernel>();
+
+  EXPECT_TRUE(sycl::has_kernel_bundle<sycl::bundle_state::executable>(
+      sycl::context(Dev), {KernelId1, KernelId2}));
+}
+
+TEST(KernelBundle, CheckExceptionIfKernelIncompatible) {
+  sycl::unittest::PiMock Mock;
+  // TestKernelWithAspects has GPU aspect, so it shouldn't be compatible with
+  // the CPU device and hence shouldn't be in the kernel bundle.
+  Mock.redefineAfter<sycl::detail::PiApiKind::piDeviceGetInfo>(
+      redefinedDeviceGetInfoCPU);
+  sycl::platform Plt = Mock.getPlatform();
+  const sycl::device Dev = Plt.get_devices()[0];
+  EXPECT_TRUE(Dev.is_cpu());
+
+  auto KernelId = sycl::get_kernel_id<TestKernelWithAspects>();
+  std::string msg = "";
+  try {
+    auto Bundle = sycl::get_kernel_bundle<sycl::bundle_state::executable>(
+        sycl::context(Dev), {Dev}, {KernelId});
+  } catch (sycl::exception &e) {
+    msg = e.what();
+  }
+  EXPECT_EQ(msg, "Kernel is incompatible with all devices in devs");
+}