From 288e39c27c61e0fe0886b51ff3893b1f73f1e092 Mon Sep 17 00:00:00 2001 From: Lukas Sommer Date: Thu, 14 Dec 2023 10:18:59 +0000 Subject: [PATCH 1/5] [SYCL][Fusion] Mark unsupported on accelerator Kernel fusion is currently not supported for SYCL devices of device type accelerator, as devices such as FPGAs do not support JIT compilation at runtime. Mark all end-to-end tests related to kernel fusion as unsupported for devices of type accelerator. Also use the newly created file to make the fact that kernel fusion is currently unsupported on Windows more explicit. Signed-off-by: Lukas Sommer --- sycl/test-e2e/KernelFusion/lit.local.cfg | 7 +++++++ 1 file changed, 7 insertions(+) create mode 100644 sycl/test-e2e/KernelFusion/lit.local.cfg diff --git a/sycl/test-e2e/KernelFusion/lit.local.cfg b/sycl/test-e2e/KernelFusion/lit.local.cfg new file mode 100644 index 0000000000000..218cd14243702 --- /dev/null +++ b/sycl/test-e2e/KernelFusion/lit.local.cfg @@ -0,0 +1,7 @@ +import platform + +config.unsupported_features += ['accelerator'] + +# TODO: enable on Windows once kernel fusion is supported on Windows. +if platform.system() != "Linux": + config.unsupported = True From fdec77883e8011c846a153b086b17c9d988559a7 Mon Sep 17 00:00:00 2001 From: Lukas Sommer Date: Thu, 14 Dec 2023 17:06:00 +0000 Subject: [PATCH 2/5] [SYCL][Fusion] Throw error if fusion unsupported Signed-off-by: Lukas Sommer --- sycl/source/detail/queue_impl.hpp | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/sycl/source/detail/queue_impl.hpp b/sycl/source/detail/queue_impl.hpp index ea5542d23decf..a5d556a7f60a5 100644 --- a/sycl/source/detail/queue_impl.hpp +++ b/sycl/source/detail/queue_impl.hpp @@ -25,6 +25,7 @@ #include #include #include +#include #include #include #include @@ -146,6 +147,14 @@ class queue_impl { "Queue compute index must be a non-negative number less than " "device's number of available compute queue indices."); } + if (has_property< + ext::codeplay::experimental::property::queue::enable_fusion>() && + !MDevice->get_info< + ext::codeplay::experimental::info::device::supports_fusion>()) { + throw sycl::exception( + make_error_code(errc::invalid), + "Cannot enable fusion if device does not support fusion"); + } if (!Context->isDeviceValid(Device)) { if (!Context->is_host() && Context->getBackend() == backend::opencl) throw sycl::invalid_object_error( From f2d2a6ab844ae9996a339d8e68f079a6be156598 Mon Sep 17 00:00:00 2001 From: Lukas Sommer Date: Fri, 15 Dec 2023 07:15:50 +0000 Subject: [PATCH 3/5] [SYCL][Fusion] Refine fusion device info Refine fusion device info to not report support for accelerator devices on the OpenCL backend. Signed-off-by: Lukas Sommer --- sycl/source/detail/device_info.hpp | 8 +++++++- 1 file changed, 7 insertions(+), 1 deletion(-) diff --git a/sycl/source/detail/device_info.hpp b/sycl/source/detail/device_info.hpp index c2932e0ebe9d9..dbeb26f8fc416 100644 --- a/sycl/source/detail/device_info.hpp +++ b/sycl/source/detail/device_info.hpp @@ -1009,8 +1009,14 @@ struct get_device_info_impl< #if SYCL_EXT_CODEPLAY_KERNEL_FUSION // Currently fusion is only supported for SPIR-V based backends, // CUDA and HIP. + if (Dev->getBackend() == backend::opencl) { + // Exclude all non-CPU or non-GPU devices on OpenCL, in particular + // accelerators. + return Dev->get_device_type() == pi::PiDeviceType::PI_DEVICE_TYPE_CPU || + Dev->get_device_type() == pi::PiDeviceType::PI_DEVICE_TYPE_GPU; + } + return (Dev->getBackend() == backend::ext_oneapi_level_zero) || - (Dev->getBackend() == backend::opencl) || (Dev->getBackend() == backend::ext_oneapi_cuda) || (Dev->getBackend() == backend::ext_oneapi_hip); #else // SYCL_EXT_CODEPLAY_KERNEL_FUSION From 06f5ee12f2646c42fae18a273856e12faa2a6b90 Mon Sep 17 00:00:00 2001 From: Lukas Sommer Date: Fri, 15 Dec 2023 09:26:52 +0000 Subject: [PATCH 4/5] Skip tests if fusion is not supported by device Signed-off-by: Lukas Sommer --- sycl/unittests/Extensions/CommandGraph.cpp | 10 +++++++++- 1 file changed, 9 insertions(+), 1 deletion(-) diff --git a/sycl/unittests/Extensions/CommandGraph.cpp b/sycl/unittests/Extensions/CommandGraph.cpp index 24a55475121e6..02d2e1e3c922e 100644 --- a/sycl/unittests/Extensions/CommandGraph.cpp +++ b/sycl/unittests/Extensions/CommandGraph.cpp @@ -1525,7 +1525,15 @@ TEST_F(CommandGraphTest, DependencyLeavesKeyword4) { } TEST_F(CommandGraphTest, FusionExtensionExceptionCheck) { - queue Q{ext::codeplay::experimental::property::queue::enable_fusion{}}; + device D; + if (!D.get_info< + ext::codeplay::experimental::info::device::supports_fusion>()) { + // Skip this test if the device does not support fusion. Otherwise, the + // queue construction in the next step would fail. + GTEST_SKIP(); + } + + queue Q{D, ext::codeplay::experimental::property::queue::enable_fusion{}}; experimental::command_graph Graph{ Q.get_context(), Q.get_device()}; From 364d790cf0eb8122ade27c2da5db61e3281e93bd Mon Sep 17 00:00:00 2001 From: Lukas Sommer Date: Fri, 15 Dec 2023 11:20:26 +0000 Subject: [PATCH 5/5] Address PR feedback Signed-off-by: Lukas Sommer --- sycl/source/detail/device_info.hpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/sycl/source/detail/device_info.hpp b/sycl/source/detail/device_info.hpp index dbeb26f8fc416..0d06befc71198 100644 --- a/sycl/source/detail/device_info.hpp +++ b/sycl/source/detail/device_info.hpp @@ -1012,8 +1012,7 @@ struct get_device_info_impl< if (Dev->getBackend() == backend::opencl) { // Exclude all non-CPU or non-GPU devices on OpenCL, in particular // accelerators. - return Dev->get_device_type() == pi::PiDeviceType::PI_DEVICE_TYPE_CPU || - Dev->get_device_type() == pi::PiDeviceType::PI_DEVICE_TYPE_GPU; + return Dev->is_cpu() || Dev->is_gpu(); } return (Dev->getBackend() == backend::ext_oneapi_level_zero) ||