diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index e124f0316932a..44b172c60d12c 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -6,6 +6,7 @@ // //===----------------------------------------------------------------------===// +#include #include #include #include @@ -290,12 +291,59 @@ std::vector get_kernel_ids() { } bool is_compatible(const std::vector &KernelIDs, const device &Dev) { - std::set 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 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) diff --git a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/Inputs/is_compatible_with_env.cpp b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/Inputs/is_compatible_with_env.cpp new file mode 100644 index 0000000000000..e919a6a3bf001 --- /dev/null +++ b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/Inputs/is_compatible_with_env.cpp @@ -0,0 +1,14 @@ +#include + +int main() { + sycl::device dev; + if (sycl::is_compatible(dev)) { + sycl::queue q(dev); + q.submit([&](sycl::handler &cgh) { + cgh.parallel_for(sycl::range<1>{1}, + [=](sycl::id<1> Id) { int x = Id[0]; }); + }).wait_and_throw(); + return 0; + } + return 1; +} diff --git a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_amdgcn.cpp b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_amdgcn.cpp new file mode 100644 index 0000000000000..5e240044483b7 --- /dev/null +++ b/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 diff --git a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_esimd_emulator.cpp b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_esimd_emulator.cpp new file mode 100644 index 0000000000000..1042b983167cc --- /dev/null +++ b/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 +#include + +int main() { + sycl::device dev; + if (sycl::is_compatible(dev)) { + float *A = malloc_shared(Size, q); + float *B = malloc_shared(Size, q); + float *C = malloc_shared(Size, q); + + for (unsigned i = 0; i != Size; i++) { + A[i] = B[i] = i; + } + + q.submit([&](handler &cgh) { + cgh.parallel_for(Size / VL, + [=](id<1> i) [[intel::sycl_explicit_simd]] { + auto offset = i * VL; + // pointer arithmetic, so offset is in + // elements: + simd va(A + offset); + simd vb(B + offset); + simd vc = va + vb; + vc.copy_to(C + offset); + }); + }).wait_and_throw(); + return 0; + } + return 1; +} diff --git a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_nvptx64.cpp b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_nvptx64.cpp new file mode 100644 index 0000000000000..ccfa829293c3f --- /dev/null +++ b/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 diff --git a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_several_targets.cpp b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_several_targets.cpp new file mode 100644 index 0000000000000..6dcc4690880d6 --- /dev/null +++ b/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 diff --git a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64.cpp b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64.cpp new file mode 100644 index 0000000000000..465a79056906a --- /dev/null +++ b/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 diff --git a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_fpga.cpp b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_fpga.cpp new file mode 100644 index 0000000000000..57366482e7082 --- /dev/null +++ b/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 diff --git a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_gen.cpp b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_gen.cpp new file mode 100644 index 0000000000000..5adb27e0ae697 --- /dev/null +++ b/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 diff --git a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_x86_64.cpp b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_x86_64.cpp new file mode 100644 index 0000000000000..0a6f2c39df8af --- /dev/null +++ b/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 diff --git a/sycl/test-e2e/OptionalKernelFeatures/is_compatible.cpp b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_with_aspects.cpp similarity index 100% rename from sycl/test-e2e/OptionalKernelFeatures/is_compatible.cpp rename to sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_with_aspects.cpp