From 67f6bba2537dbacf91f023caf830d9dbb7536f33 Mon Sep 17 00:00:00 2001
From: Nikita <nikita.kornev@intel.com>
Date: Mon, 28 Nov 2022 18:46:39 +0100
Subject: [PATCH] [SYCL] Implement sycl::is_compatible() function (#7510)

---
 sycl/include/sycl/kernel_bundle.hpp     |  3 ++-
 sycl/source/kernel_bundle.cpp           | 29 +++++++++++++++++++++++++
 sycl/test/abi/sycl_symbols_linux.dump   |  1 +
 sycl/test/abi/sycl_symbols_windows.dump |  1 +
 4 files changed, 33 insertions(+), 1 deletion(-)

diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp
index 23b540cc810b7..791eee545052c 100644
--- a/sycl/include/sycl/kernel_bundle.hpp
+++ b/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);
diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp
index db94b79561a7e..826170316f124 100644
--- a/sycl/source/kernel_bundle.cpp
+++ b/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
diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump
index 1a3f21f63425d..8a452fefbe141 100644
--- a/sycl/test/abi/sycl_symbols_linux.dump
+++ b/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
diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump
index b1b892c6e24ac..36657b27b85ce 100644
--- a/sycl/test/abi/sycl_symbols_windows.dump
+++ b/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