[SYCL] Implement sycl::is_compatible() function (#7510)

Author: KornevNikita

sycl/include/sycl/kernel_bundle.hpp

@@ -564,7 +564,8 @@ bool has_kernel_bundle(const context &Ctx, const std::vector<device> &Devs) {
 
 /// \returns true if all of the kernels identified by KernelIDs are compatible
 /// with the device Dev.
-bool is_compatible(const std::vector<kernel_id> &KernelIDs, const device &Dev);
+__SYCL_EXPORT bool is_compatible(const std::vector<kernel_id> &KernelIDs,
+                                 const device &Dev);
 
 template <typename KernelName> bool is_compatible(const device &Dev) {
   return is_compatible({get_kernel_id<KernelName>()}, Dev);

sycl/source/kernel_bundle.cpp

@@ -8,6 +8,7 @@
 
 #include <detail/kernel_bundle_impl.hpp>
 #include <detail/kernel_id_impl.hpp>
+#include <detail/program_manager/program_manager.hpp>
 
 #include <set>
 
@@ -294,5 +295,33 @@ std::vector<kernel_id> get_kernel_ids() {
   return detail::ProgramManager::getInstance().getAllSYCLKernelIDs();
 }
 
+bool is_compatible(const std::vector<kernel_id> &KernelIDs, const device &Dev) {
+  for (const auto &KernelId : KernelIDs) {
+    const detail::RTDeviceBinaryImage &Img =
+        detail::ProgramManager::getInstance().getDeviceImage(
+            detail::OSUtil::ExeModuleHandle, KernelId.get_name(), context(Dev),
+            Dev);
+    const detail::RTDeviceBinaryImage::PropertyRange &ARange =
+        Img.getDeviceRequirements();
+    for (detail::RTDeviceBinaryImage::PropertyRange::ConstIterator It :
+         ARange) {
+      using namespace std::literals;
+      if ((*It)->Name != "aspects"sv)
+        continue;
+      detail::ByteArray Aspects =
+          detail::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;
+}
+
 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
 } // namespace sycl

sycl/test/abi/sycl_symbols_linux.dump

@@ -3619,6 +3619,7 @@ _ZN4sycl3_V113aligned_allocEmmRKNS0_5queueENS0_3usm5allocERKNS0_13property_listE
 _ZN4sycl3_V113aligned_allocEmmRKNS0_5queueENS0_3usm5allocERKNS0_6detail13code_locationE
 _ZN4sycl3_V113aligned_allocEmmRKNS0_6deviceERKNS0_7contextENS0_3usm5allocERKNS0_13property_listERKNS0_6detail13code_locationE
 _ZN4sycl3_V113aligned_allocEmmRKNS0_6deviceERKNS0_7contextENS0_3usm5allocERKNS0_6detail13code_locationE
+_ZN4sycl3_V113is_compatibleERKSt6vectorINS0_9kernel_idESaIS2_EERKNS0_6deviceE
 _ZN4sycl3_V113malloc_deviceEmRKNS0_5queueERKNS0_13property_listERKNS0_6detail13code_locationE
 _ZN4sycl3_V113malloc_deviceEmRKNS0_5queueERKNS0_6detail13code_locationE
 _ZN4sycl3_V113malloc_deviceEmRKNS0_6deviceERKNS0_7contextERKNS0_13property_listERKNS0_6detail13code_locationE

sycl/test/abi/sycl_symbols_windows.dump

@@ -948,6 +948,7 @@
 ?isValidModeForSourceAccessor@handler@_V1@sycl@@CA_NW4mode@access@23@@Z
 ?isValidTargetForExplicitOp@handler@_V1@sycl@@CA_NW4target@access@23@@Z
 ?is_accelerator@device@_V1@sycl@@QEBA_NXZ
+?is_compatible@_V1@sycl@@YA_NAEBV?$vector@Vkernel_id@_V1@sycl@@V?$allocator@Vkernel_id@_V1@sycl@@@std@@@std@@AEBVdevice@12@@Z
 ?is_cpu@device@_V1@sycl@@QEBA_NXZ
 ?is_gpu@device@_V1@sycl@@QEBA_NXZ
 ?is_host@context@_V1@sycl@@QEBA_NXZ