From 08ff2dc0231e4d2a79b6c103e184b5acd755819d Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Tue, 6 Jun 2023 10:05:59 -0700 Subject: [PATCH 1/9] [WIP][SYCL] Improve is_compatible Modify is_compatible to check if specific target 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. https://github.com/KhronosGroup/SYCL-Docs/pull/381 --- sycl/source/kernel_bundle.cpp | 31 +++++++++++++++++-- .../Inputs/is_compatible_with_env.cpp | 15 +++++++++ .../is_compatible_spir64_fpga.cpp | 6 ++++ .../is_compatible_spir64_gen.cpp | 6 ++++ .../is_compatible_spir64_x86_64.cpp | 6 ++++ .../is_compatible_with_aspects.cpp} | 0 6 files changed, 62 insertions(+), 2 deletions(-) create mode 100644 sycl/test-e2e/OptionalKernelFeatures/is_compatible/Inputs/is_compatible_with_env.cpp create mode 100644 sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_fpga.cpp create mode 100644 sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_gen.cpp create mode 100644 sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_x86_64.cpp rename sycl/test-e2e/OptionalKernelFeatures/{is_compatible.cpp => is_compatible/is_compatible_with_aspects.cpp} (100%) diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index e124f0316932a..79d3ae1cdc31e 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -292,9 +292,36 @@ std::vector get_kernel_ids() { bool is_compatible(const std::vector &KernelIDs, const device &Dev) { std::set BinImages = detail::ProgramManager::getInstance().getRawDeviceImages(KernelIDs); + + auto imageTargetMatchDevice = [](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); + } 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; + }; return std::all_of(BinImages.begin(), BinImages.end(), - [&Dev](const detail::RTDeviceBinaryImage *Img) { - return doesDevSupportDeviceRequirements(Dev, *Img); + [&](const detail::RTDeviceBinaryImage *Img) { + return doesDevSupportDeviceRequirements(Dev, *Img) && + imageTargetMatchDevice(Dev, *Img); }); } 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..8e721e660e6b0 --- /dev/null +++ b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/Inputs/is_compatible_with_env.cpp @@ -0,0 +1,15 @@ +#include + +int main() { + sycl::device dev; + // Should not throw any exception as it should only run on the specific + // target device, defined during compilation. + 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; +} 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..4ec1bc73e63df --- /dev/null +++ b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_fpga.cpp @@ -0,0 +1,6 @@ +// REQUIRES: cpu, 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=opencl:gpu %{run} %t.out +// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run} %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..6ca3223296c83 --- /dev/null +++ b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_gen.cpp @@ -0,0 +1,6 @@ +// 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_zerp:gpu %{run} %t.out +// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run} %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..8f9681c88f051 --- /dev/null +++ b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_x86_64.cpp @@ -0,0 +1,6 @@ +// REQUIRES: gpu, level_zero, opencl + +// 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=level_zero:gpu %{run} %t.out +// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:gpu %{run} %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 From b0c54efca941c6840b7463d74e5eae1e858a04ee Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Wed, 7 Jun 2023 09:21:25 -0700 Subject: [PATCH 2/9] Modify logic --- sycl/source/kernel_bundle.cpp | 29 +++++++++++++++++++++-------- 1 file changed, 21 insertions(+), 8 deletions(-) diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index 79d3ae1cdc31e..4401604057979 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -290,11 +290,11 @@ std::vector get_kernel_ids() { } bool is_compatible(const std::vector &KernelIDs, const device &Dev) { - std::set BinImages = - detail::ProgramManager::getInstance().getRawDeviceImages(KernelIDs); - - auto imageTargetMatchDevice = [](const device &Dev, - const detail::RTDeviceBinaryImage &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) { @@ -318,11 +318,24 @@ bool is_compatible(const std::vector &KernelIDs, const device &Dev) { return false; }; - return std::all_of(BinImages.begin(), BinImages.end(), + + // 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) && - imageTargetMatchDevice(Dev, *Img); - }); + doesImageTargetMatchDevice(Dev, *Img); + })) + return false; + } + + return true; } } // __SYCL_INLINE_VER_NAMESPACE(_V1) From 44d5c89980e3879df1784c952e554b61ff8f09a1 Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Mon, 12 Jun 2023 07:11:24 -0700 Subject: [PATCH 3/9] Add tests for CUDA & HIP --- .../is_compatible/is_compatible_amdgcn.cpp | 7 +++++++ .../is_compatible/is_compatible_nvptx64.cpp | 7 +++++++ 2 files changed, 14 insertions(+) create mode 100644 sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_amdgcn.cpp create mode 100644 sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_nvptx64.cpp 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..63f672364237d --- /dev/null +++ b/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 %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_nvptx64.cpp b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_nvptx64.cpp new file mode 100644 index 0000000000000..1f2669883f2a3 --- /dev/null +++ b/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 From 1bafaa7139c797f9dd1cb0c42bad2630964be6f8 Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Mon, 12 Jun 2023 07:12:14 -0700 Subject: [PATCH 4/9] Add test with multiple bin images --- .../is_compatible/is_compatible_several_targets.cpp | 8 ++++++++ 1 file changed, 8 insertions(+) create mode 100644 sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_several_targets.cpp 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..c8627e9cc2906 --- /dev/null +++ b/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 From 7cce7060fc40cb10a9cf3921b09727a33923da6c Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Mon, 12 Jun 2023 07:12:48 -0700 Subject: [PATCH 5/9] Improve base test & add esimd backend --- sycl/source/kernel_bundle.cpp | 3 ++- .../is_compatible/Inputs/is_compatible_with_env.cpp | 5 ++--- .../is_compatible/is_compatible_amdgcn.cpp | 2 +- .../is_compatible/is_compatible_spir64_fpga.cpp | 7 ++++--- .../is_compatible/is_compatible_spir64_gen.cpp | 5 +++-- .../is_compatible/is_compatible_spir64_x86_64.cpp | 7 ++++--- 6 files changed, 16 insertions(+), 13 deletions(-) diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index 4401604057979..bedd42e707a31 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -299,7 +299,8 @@ bool is_compatible(const std::vector &KernelIDs, const device &Dev) { 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_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(); 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 index 8e721e660e6b0..e919a6a3bf001 100644 --- 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 @@ -2,14 +2,13 @@ int main() { sycl::device dev; - // Should not throw any exception as it should only run on the specific - // target device, defined during compilation. 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 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 index 63f672364237d..ddea75cb7e8cb 100644 --- a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_amdgcn.cpp +++ b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_amdgcn.cpp @@ -1,6 +1,6 @@ // REQUIRES: hip -// RUN: %clangxx -fsycl -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx906 %S/Inputs/is_compatible_with_env.cpp -o %t.out +// 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 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 index 4ec1bc73e63df..c8ee54bcfebde 100644 --- a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_fpga.cpp +++ b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_fpga.cpp @@ -1,6 +1,7 @@ -// REQUIRES: cpu, gpu, accelerator +// 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=opencl:gpu %{run} %t.out -// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:cpu %{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 index 6ca3223296c83..d524197b1b230 100644 --- a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_gen.cpp +++ b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_gen.cpp @@ -1,6 +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_zerp:gpu %{run} %t.out -// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:cpu %{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 index 8f9681c88f051..c0c4b4185931f 100644 --- 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 @@ -1,6 +1,7 @@ -// REQUIRES: gpu, level_zero, opencl +// 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=level_zero:gpu %{run} %t.out -// RUN: env ONEAPI_DEVICE_SELECTOR=opencl:gpu %{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 From b4cb2dd49fd673d7082a7e6edcfa5f4f18ba64c3 Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Mon, 12 Jun 2023 10:29:03 -0700 Subject: [PATCH 6/9] Update tests requirements & add spir64 test --- .../is_compatible/is_compatible_amdgcn.cpp | 2 +- .../is_compatible/is_compatible_nvptx64.cpp | 2 +- .../is_compatible/is_compatible_several_targets.cpp | 2 +- .../is_compatible/is_compatible_spir64.cpp | 7 +++++++ .../is_compatible/is_compatible_spir64_fpga.cpp | 2 +- .../is_compatible/is_compatible_spir64_gen.cpp | 2 +- .../is_compatible/is_compatible_spir64_x86_64.cpp | 2 +- 7 files changed, 13 insertions(+), 6 deletions(-) create mode 100644 sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64.cpp diff --git a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_amdgcn.cpp b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_amdgcn.cpp index ddea75cb7e8cb..5e240044483b7 100644 --- a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_amdgcn.cpp +++ b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_amdgcn.cpp @@ -1,4 +1,4 @@ -// REQUIRES: hip +// 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 diff --git a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_nvptx64.cpp b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_nvptx64.cpp index 1f2669883f2a3..ccfa829293c3f 100644 --- a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_nvptx64.cpp +++ b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_nvptx64.cpp @@ -1,4 +1,4 @@ -// REQUIRES: cuda +// REQUIRES: cuda, opencl, gpu, cpu // RUN: %clangxx -fsycl -fsycl-targets=nvptx64-nvidia-cuda %S/Inputs/is_compatible_with_env.cpp -o %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 index c8627e9cc2906..6dcc4690880d6 100644 --- a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_several_targets.cpp +++ b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_several_targets.cpp @@ -1,4 +1,4 @@ -// REQUIRES: level_zero, accelerator +// 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 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 index c8ee54bcfebde..57366482e7082 100644 --- a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_fpga.cpp +++ b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_fpga.cpp @@ -1,4 +1,4 @@ -// REQUIRES: gpu, accelerator +// REQUIRES: opencl-aot, accelerator, gpu, cpu // RUN: %clangxx -fsycl -fsycl-targets=spir64_fpga %S/Inputs/is_compatible_with_env.cpp -o %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 index d524197b1b230..5adb27e0ae697 100644 --- a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_gen.cpp +++ b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_spir64_gen.cpp @@ -1,4 +1,4 @@ -// REQUIRES: ocloc, gpu, level_zero, opencl +// 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 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 index c0c4b4185931f..0a6f2c39df8af 100644 --- 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 @@ -1,4 +1,4 @@ -// REQUIRES: level_zero +// 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 From 187bf48c79615cd1d6ffa53e382021cc32a391b1 Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Tue, 13 Jun 2023 05:58:37 -0700 Subject: [PATCH 7/9] Improve esimd BE handling & add esimd test --- sycl/source/kernel_bundle.cpp | 13 +++++-- .../is_compatible_esimd_emulator.cpp | 39 +++++++++++++++++++ 2 files changed, 49 insertions(+), 3 deletions(-) create mode 100644 sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_esimd_emulator.cpp diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index bedd42e707a31..01dedc0026469 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -292,15 +292,22 @@ std::vector get_kernel_ids() { bool is_compatible(const std::vector &KernelIDs, const device &Dev) { if (KernelIDs.empty()) return false; - // TODO: also need to check architectures matching + // 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 is only compatible with esimd kernels. + if (BE == sycl::backend::ext_intel_esimd_emulator) { + pi_device_binary_property Prop = Img.getProperty("isEsimdImage"); + return(Prop && (DeviceBinaryProperty(Prop).asUint32() != 0)); + } if (strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64) == 0) { + if (sycl::detail::KernelInfo) return (BE == sycl::backend::opencl || - BE == sycl::backend::ext_oneapi_level_zero || - BE == sycl::backend::ext_intel_esimd_emulator); + BE == sycl::backend::ext_oneapi_level_zero); } else if (strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64_X86_64) == 0) { return Dev.is_cpu(); 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..e4eeaac5199e8 --- /dev/null +++ b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_esimd_emulator.cpp @@ -0,0 +1,39 @@ +// 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; +} \ No newline at end of file From d3ea554af670710a22acbc15d551dada79589c49 Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Tue, 13 Jun 2023 07:06:23 -0700 Subject: [PATCH 8/9] Fix build --- sycl/source/kernel_bundle.cpp | 4 ++-- .../is_compatible_esimd_emulator.cpp | 24 ++++++++++--------- 2 files changed, 15 insertions(+), 13 deletions(-) diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index 01dedc0026469..978b1a36e8a6e 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -6,6 +6,7 @@ // //===----------------------------------------------------------------------===// +#include #include #include #include @@ -302,10 +303,9 @@ bool is_compatible(const std::vector &KernelIDs, const device &Dev) { // ESIMD emulator is only compatible with esimd kernels. if (BE == sycl::backend::ext_intel_esimd_emulator) { pi_device_binary_property Prop = Img.getProperty("isEsimdImage"); - return(Prop && (DeviceBinaryProperty(Prop).asUint32() != 0)); + return (Prop && (detail::DeviceBinaryProperty(Prop).asUint32() != 0)); } if (strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64) == 0) { - if (sycl::detail::KernelInfo) return (BE == sycl::backend::opencl || BE == sycl::backend::ext_oneapi_level_zero); } else if (strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64_X86_64) == 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 index e4eeaac5199e8..5e356fa60fad3 100644 --- a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_esimd_emulator.cpp +++ b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_esimd_emulator.cpp @@ -6,7 +6,8 @@ // 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 +// Just an example from +// https://github.com/intel/llvm/tree/sycl/sycl/doc/extensions/experimental/sycl_ext_intel_esimd #include #include @@ -23,16 +24,17 @@ int main() { } 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(); + 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; From 712f2efb0c582e34b25c5ebfd7291a633a695e4a Mon Sep 17 00:00:00 2001 From: KornevNikita Date: Wed, 14 Jun 2023 03:12:51 -0700 Subject: [PATCH 9/9] Apply suggestions --- sycl/source/kernel_bundle.cpp | 2 +- .../is_compatible/is_compatible_esimd_emulator.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index 978b1a36e8a6e..44b172c60d12c 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -300,7 +300,7 @@ bool is_compatible(const std::vector &KernelIDs, const device &Dev) { const detail::RTDeviceBinaryImage &Img) { const char *Target = Img.getRawData().DeviceTargetSpec; auto BE = Dev.get_backend(); - // ESIMD emulator is only compatible with esimd kernels. + // 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)); 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 index 5e356fa60fad3..1042b983167cc 100644 --- a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_esimd_emulator.cpp +++ b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_esimd_emulator.cpp @@ -38,4 +38,4 @@ int main() { return 0; } return 1; -} \ No newline at end of file +}