[SYCL] Improve is_compatible

sycl/source/kernel_bundle.cpp

@@ -290,12 +290,53 @@ std::vector<kernel_id> get_kernel_ids() {
 }
 
 bool is_compatible(const std::vector<kernel_id> &KernelIDs, const device &Dev) {
-  std::set<detail::RTDeviceBinaryImage *> BinImages =
-      detail::ProgramManager::getInstance().getRawDeviceImages(KernelIDs);
-  return std::all_of(BinImages.begin(), BinImages.end(),
-                     [&Dev](const detail::RTDeviceBinaryImage *Img) {
-                       return doesDevSupportDeviceRequirements(Dev, *Img);
-                     });
+  if (KernelIDs.empty())
+    return false;
+  // TODO: also need to check architectures matching
+  auto doesImageTargetMatchDevice = [](const device &Dev,
+                                       const detail::RTDeviceBinaryImage &Img) {
+    const char *Target = Img.getRawData().DeviceTargetSpec;
+    auto BE = Dev.get_backend();
+    if (strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64) == 0) {
+      return (BE == sycl::backend::opencl ||
+              BE == sycl::backend::ext_oneapi_level_zero ||
+              BE == sycl::backend::ext_intel_esimd_emulator);
+    } else if (strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64_X86_64) ==
+               0) {
+      return Dev.is_cpu();
+    } else if (strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64_GEN) ==
+               0) {
+      return Dev.is_gpu() && (BE == sycl::backend::opencl ||
+                              BE == sycl::backend::ext_oneapi_level_zero);
+    } else if (strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64_FPGA) ==
+               0) {
+      return Dev.is_accelerator();
+    } else if (strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_NVPTX64) == 0) {
+      return BE == sycl::backend::ext_oneapi_cuda;
+    } else if (strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_AMDGCN) == 0) {
+      return BE == sycl::backend::ext_oneapi_hip;
+    }
+
+    return false;
+  };
+
+  // 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<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(Dev, *Img);
+                     }))
+      return false;
+  }
+
+  return true;
 }
 
 } // __SYCL_INLINE_VER_NAMESPACE(_V1)

sycl/test-e2e/OptionalKernelFeatures/is_compatible/Inputs/is_compatible_with_env.cpp

@@ -0,0 +1,14 @@
+#include <sycl/sycl.hpp>
+
+int main() {
+  sycl::device dev;
+  if (sycl::is_compatible<class Kernel>(dev)) {
+    sycl::queue q(dev);
+    q.submit([&](sycl::handler &cgh) {
+       cgh.parallel_for<class Kernel>(sycl::range<1>{1},
+                                      [=](sycl::id<1> Id) { int x = Id[0]; });
+     }).wait_and_throw();
+    return 0;
+  }
+  return 1;
+}

sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_amdgcn.cpp

@@ -0,0 +1,7 @@
+// REQUIRES: hip
+
+// RUN: %clangxx -fsycl -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx906 -fsycl-targets=amdgcn-amd-amdhsa %S/Inputs/is_compatible_with_env.cpp -o %t.out
+
+// RUN: env ONEAPI_DEVICE_SELECTOR=hip:gpu %{run} %t.out
+// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:gpu %{run} not %t.out
+// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run} not %t.out

sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_nvptx64.cpp

@@ -0,0 +1,7 @@
+// REQUIRES: cuda
+
+// RUN: %clangxx -fsycl -fsycl-targets=nvptx64-nvidia-cuda %S/Inputs/is_compatible_with_env.cpp -o %t.out
+
+// RUN: env ONEAPI_DEVICE_SELECTOR=cuda:gpu %{run} %t.out
+// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:gpu %{run} not %t.out
+// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run} not %t.out

sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_several_targets.cpp

@@ -0,0 +1,8 @@
+// REQUIRES: level_zero, accelerator
+
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_fpga,spir64_gen -Xsycl-target-backend "-device *" %S/Inputs/is_compatible_with_env.cpp -o %t.out
+
+// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run} not %t.out
+// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:acc %{run} %t.out
+// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:gpu %{run} %t.out
+// RUN: env ONEAPI_DEVICE_SELECTOR=level_zero:gpu %{run} %t.out

sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_fpga.cpp

@@ -0,0 +1,7 @@
+// REQUIRES: gpu, accelerator
+
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_fpga %S/Inputs/is_compatible_with_env.cpp -o %t.out
+
+// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:fpga %{run} %t.out
+// RUN: env ONEAPI_DEVICE_SELECTOR=*:gpu %{run} not %t.out
+// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run} not %t.out

sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_gen.cpp

@@ -0,0 +1,7 @@
+// REQUIRES: ocloc, gpu, level_zero, opencl
+
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend "-device *" %S/Inputs/is_compatible_with_env.cpp -o %t.out
+
+// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:gpu %{run} %t.out
+// RUN: env ONEAPI_DEVICE_SELECTOR=level_zero:gpu %{run} %t.out
+// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run} not %t.out

sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_x86_64.cpp

@@ -0,0 +1,7 @@
+// REQUIRES: level_zero
+
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64 %S/Inputs/is_compatible_with_env.cpp -o %t.out
+
+// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run} %t.out
+// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:gpu %{run} not %t.out
+// RUN: env ONEAPI_DEVICE_SELECTOR=level_zero:gpu %{run} not %t.out