[SYCL] Improve is_compatible (#9769)

Modify `is_compatible` to check if specific target is defined with
`-fsycl-targets` and change the result. Previously there was a situation
when kernel is compatible with the device by aspects, but actually it
fails to run on this device as it was compiled for another target
device.

Related spec change: KhronosGroup/SYCL-Docs#381

Resolves #7561

Author: KornevNikita

sycl/source/kernel_bundle.cpp

@@ -6,6 +6,7 @@
 //
 //===----------------------------------------------------------------------===//
 
+#include <detail/device_binary_image.hpp>
 #include <detail/kernel_bundle_impl.hpp>
 #include <detail/kernel_id_impl.hpp>
 #include <detail/program_manager/program_manager.hpp>
@@ -290,12 +291,59 @@ 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 that the architecture specified by the
+  // "-fsycl-targets" flag matches the device when we are able to get the
+  // device's arch.
+  auto doesImageTargetMatchDevice = [](const device &Dev,
+                                       const detail::RTDeviceBinaryImage &Img) {
+    const char *Target = Img.getRawData().DeviceTargetSpec;
+    auto BE = Dev.get_backend();
+    // ESIMD emulator backend is only compatible with esimd kernels.
+    if (BE == sycl::backend::ext_intel_esimd_emulator) {
+      pi_device_binary_property Prop = Img.getProperty("isEsimdImage");
+      return (Prop && (detail::DeviceBinaryProperty(Prop).asUint32() != 0));
+    }
+    if (strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64) == 0) {
+      return (BE == sycl::backend::opencl ||
+              BE == sycl::backend::ext_oneapi_level_zero);
+    } 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_amd, opencl, gpu, cpu
+
+// 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_esimd_emulator.cpp

@@ -0,0 +1,41 @@
+// REQUIRES: esimd_emulator
+
+// RUN: %clangxx -fsycl %S/Inputs/is_compatible_with_env.cpp %t_negative_case.out
+// RUN: env ONEAPI_DEVICE_SELECTOR=ext_intel_esimd_emulator:gpu %{run} not %t_negative_case.out
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// Just an example from
+// https://github.com/intel/llvm/tree/sycl/sycl/doc/extensions/experimental/sycl_ext_intel_esimd
+
+#include <sycl/ext/intel/esimd.hpp>
+#include <sycl/sycl.hpp>
+
+int main() {
+  sycl::device dev;
+  if (sycl::is_compatible<class Test>(dev)) {
+    float *A = malloc_shared<float>(Size, q);
+    float *B = malloc_shared<float>(Size, q);
+    float *C = malloc_shared<float>(Size, q);
+
+    for (unsigned i = 0; i != Size; i++) {
+      A[i] = B[i] = i;
+    }
+
+    q.submit([&](handler &cgh) {
+       cgh.parallel_for<class Test>(Size / VL,
+                                    [=](id<1> i) [[intel::sycl_explicit_simd]] {
+                                      auto offset = i * VL;
+                                      // pointer arithmetic, so offset is in
+                                      // elements:
+                                      simd<float, VL> va(A + offset);
+                                      simd<float, VL> vb(B + offset);
+                                      simd<float, VL> vc = va + vb;
+                                      vc.copy_to(C + offset);
+                                    });
+     }).wait_and_throw();
+    return 0;
+  }
+  return 1;
+}

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

@@ -0,0 +1,7 @@
+// REQUIRES: cuda, opencl, gpu, cpu
+
+// 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: ocloc, level_zero, gpu, cpu
+
+// 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.cpp

@@ -0,0 +1,7 @@
+// REQUIRES: cuda, opencl, gpu, cpu
+
+// RUN: %clangxx -fsycl -fsycl-targets=spir64 %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} %t.out
+// RUN: env ONEAPI_DEVICE_SELECTOR=cuda:gpu %{run} not %t.out

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

@@ -0,0 +1,7 @@
+// REQUIRES: opencl-aot, accelerator, gpu, cpu
+
+// 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, cpu
+
+// 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: opencl-aot, cpu, gpu, 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