[SYCL][Fusion] Error if fusion unsupported on device (#12171)

The kernel fusion extension document says that an error will be thrown
if a `queue` is constructed with
`sycl::ext::codeplay::experimental::property::queue::enable_fusion`, but
the underlying device does not support fusion (i.e., returns false for
`get_info<sycl::ext::codeplay::experimental::info::device::supports_fusion>()`).
This change implements this behavior.

Also, 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. Therefore, mark all end-to-end tests related to
kernel fusion as unsupported for devices of type accelerator.

In addition, use the newly created file to make the fact that kernel
fusion is currently unsupported on Windows more explicit.

---------

Signed-off-by: Lukas Sommer <[email protected]>

Author: sommerlukas

sycl/source/detail/device_info.hpp

@@ -1009,8 +1009,13 @@ 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->is_cpu() || Dev->is_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

sycl/source/detail/queue_impl.hpp

@@ -25,6 +25,7 @@
 #include <sycl/event.hpp>
 #include <sycl/exception.hpp>
 #include <sycl/exception_list.hpp>
+#include <sycl/ext/codeplay/experimental/fusion_properties.hpp>
 #include <sycl/handler.hpp>
 #include <sycl/properties/context_properties.hpp>
 #include <sycl/properties/queue_properties.hpp>
@@ -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(

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

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<experimental::graph_state::modifiable> Graph{
       Q.get_context(), Q.get_device()};