[SYCL] Remove fusion scheduler integration (#15185)

Remove the integration of kernel fusion from the SYCL runtime scheduler. This makes the implementation of the scheduler less complex and leaner. As a consequence, the experimental kernel fusion feature is no longer supported, the extension proposal status has been updated. Some stubs of the kernel fusion must remain for now to avoid ABI-break, removing those stubs in the next ABI-breaking window is tracked through #15184. --------- Signed-off-by: Lukas Sommer <[email protected]>
intel · Sep 4, 2024 · 4aee47a · 4aee47a
1 parent 3b6c0d9
commit 4aee47a
Show file tree

Hide file tree

Showing 106 changed files with 133 additions and 6,959 deletions.
diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc
@@ -1921,25 +1921,6 @@ is ignored during queue recording.
 Using this information is something we may look at for future revisions of
 `sycl_ext_oneapi_graph`.
 
-==== sycl_ext_codeplay_kernel_fusion
-
-As the
-link:../experimental/sycl_ext_codeplay_kernel_fusion.asciidoc[sycl_ext_codeplay_kernel_fusion]
-extension also introduces state to a `sycl::queue`, there are restrictions on
-its usage when combined with `sycl_ext_oneapi_graph`. Exceptions with error code
-`invalid` are thrown in the following cases:
-
-* `fusion_wrapper::start_fusion()` is called when its associated queue
-  is in the recording state.
-* `command_graph::begin_recording()` is called passing a queue in fusion mode.
-
-The `sycl::ext::codeplay::experimental::property::queue::enable_fusion` property
-defined by the extension is ignored by queue recording.
-
-To enable kernel fusion in a `command_graph` see the
-link:../proposed/sycl_ext_oneapi_graph_fusion.asciidoc[sycl_ext_oneapi_graph_fusion extension proposal]
-which is layered ontop of `sycl_ext_oneapi_graph`.
-
 ==== sycl_ext_oneapi_kernel_properties
 
 The new handler methods, and queue shortcuts, defined by

diff --git a/.../sycl_ext_codeplay_kernel_fusion.asciidoc → .../sycl_ext_codeplay_kernel_fusion.asciidoc b/.../sycl_ext_codeplay_kernel_fusion.asciidoc → .../sycl_ext_codeplay_kernel_fusion.asciidoc
@@ -42,11 +42,7 @@ SYCL specification refer to that revision.
 
 == Status
 
-This is a proposed extension specification, intended to gather community
-feedback.  Interfaces defined in this specification may not be implemented yet
-or may be in a preliminary state.  The specification itself may also change in
-incompatible ways before it is finalized.  *Shipping software products should
-not rely on APIs defined in this specification.*
+This experimental extension is no longer supported.
 
 [NOTE]
 ====
@@ -713,4 +709,5 @@ with combinations of kernels that should be fused.
 |Rev|Date|Authors|Changes
 |1|2022-10-14|Victor Lomüller, Lukas Sommer and Victor Perez|*Initial draft*
 |2|2022-11-09|Victor Lomüller, Lukas Sommer and Victor Perez|*Separate fusion API into new `fusion_wrapper`*
+|3|2024-08-26|Lukas Sommer|*Mark extension as removed*
 |========================================
diff --git a/sycl/include/sycl/detail/property_helper.hpp b/sycl/include/sycl/detail/property_helper.hpp
@@ -33,6 +33,8 @@ enum DataLessPropKind {
   UseDefaultStream = 8,
   DiscardEvents = 9,
   DeviceReadOnly = 10,
+  // TODO(#15184): Remove the following fusion-related entries in the next
+  // ABI-breaking window.
   FusionPromotePrivate = 11,
   FusionPromoteLocal = 12,
   FusionNoBarrier = 13,

diff --git a/sycl/include/sycl/ext/codeplay/experimental/fusion_properties.hpp b/sycl/include/sycl/ext/codeplay/experimental/fusion_properties.hpp
diff --git a/sycl/include/sycl/ext/codeplay/experimental/fusion_wrapper.hpp b/sycl/include/sycl/ext/codeplay/experimental/fusion_wrapper.hpp
@@ -6,6 +6,8 @@
 //
 //===----------------------------------------------------------------------===//
 
+// TODO(#15184): Delete this file in the next ABI-breaking window.
+
 #pragma once
 
 #include <sycl/detail/export.hpp> // for __SYCL_EXPORT
@@ -18,10 +20,6 @@
 namespace sycl {
 inline namespace _V1 {
 
-namespace detail {
-class fusion_wrapper_impl;
-}
-
 namespace ext::codeplay::experimental {
 
 ///
@@ -92,7 +90,7 @@ class __SYCL_EXPORT fusion_wrapper {
   event complete_fusion(const property_list &propList = {});
 
 private:
-  std::shared_ptr<detail::fusion_wrapper_impl> MImpl;
+  std::shared_ptr<detail::queue_impl> MQueue;
 };
 } // namespace ext::codeplay::experimental
 } // namespace _V1

diff --git a/sycl/include/sycl/ext/oneapi/backend/level_zero.hpp b/sycl/include/sycl/ext/oneapi/backend/level_zero.hpp
@@ -8,29 +8,29 @@
 
 #pragma once
 
-#include <sycl/async_handler.hpp>                    // for async_han...
-#include <sycl/backend.hpp>                          // for backend_i...
-#include <sycl/backend_types.hpp>                    // for backend
-#include <sycl/buffer.hpp>                           // for buffer_al...
-#include <sycl/context.hpp>                          // for context
-#include <sycl/detail/backend_traits.hpp>            // for interop
-#include <sycl/detail/backend_traits_level_zero.hpp> // for ze_comman...
-#include <sycl/detail/defines_elementary.hpp>        // for __SYCL_DE...
-#include <sycl/detail/export.hpp>                    // for __SYCL_EX...
-#include <sycl/detail/impl_utils.hpp>                // for createSyc...
-#include <sycl/detail/ur.hpp>                        // for cast
-#include <sycl/device.hpp>                           // for device
-#include <sycl/event.hpp>                            // for event
-#include <sycl/ext/codeplay/experimental/fusion_properties.hpp> // for buffer
-#include <sycl/ext/oneapi/backend/level_zero_ownership.hpp>     // for ownership
-#include <sycl/image.hpp>                                       // for image
-#include <sycl/kernel.hpp>                                      // for kernel
-#include <sycl/kernel_bundle.hpp>               // for kernel_bu...
-#include <sycl/kernel_bundle_enums.hpp>         // for bundle_state
-#include <sycl/platform.hpp>                    // for platform
-#include <sycl/properties/image_properties.hpp> // for image
-#include <sycl/property_list.hpp>               // for property_...
-#include <sycl/queue.hpp>                       // for queue
+#include <sycl/async_handler.hpp>                           // for async_han...
+#include <sycl/backend.hpp>                                 // for backend_i...
+#include <sycl/backend_types.hpp>                           // for backend
+#include <sycl/buffer.hpp>                                  // for buffer_al...
+#include <sycl/buffer.hpp>                                  // for buffer
+#include <sycl/context.hpp>                                 // for context
+#include <sycl/detail/backend_traits.hpp>                   // for interop
+#include <sycl/detail/backend_traits_level_zero.hpp>        // for ze_comman...
+#include <sycl/detail/defines_elementary.hpp>               // for __SYCL_DE...
+#include <sycl/detail/export.hpp>                           // for __SYCL_EX...
+#include <sycl/detail/impl_utils.hpp>                       // for createSyc...
+#include <sycl/detail/ur.hpp>                               // for cast
+#include <sycl/device.hpp>                                  // for device
+#include <sycl/event.hpp>                                   // for event
+#include <sycl/ext/oneapi/backend/level_zero_ownership.hpp> // for ownership
+#include <sycl/image.hpp>                                   // for image
+#include <sycl/kernel.hpp>                                  // for kernel
+#include <sycl/kernel_bundle.hpp>                           // for kernel_bu...
+#include <sycl/kernel_bundle_enums.hpp>                     // for bundle_state
+#include <sycl/platform.hpp>                                // for platform
+#include <sycl/properties/image_properties.hpp>             // for image
+#include <sycl/property_list.hpp>                           // for property_...
+#include <sycl/queue.hpp>                                   // for queue
 
 #include <memory>      // for shared_ptr
 #include <stdint.h>    // for int32_t

diff --git a/sycl/include/sycl/ext/oneapi/owner_less.hpp b/sycl/include/sycl/ext/oneapi/owner_less.hpp
@@ -8,12 +8,12 @@
 
 #pragma once
 
-#include <sycl/access/access.hpp> // for access_mode
-#include <sycl/accessor.hpp>      // for host_acce...
-#include <sycl/context.hpp>       // for context
-#include <sycl/device.hpp>        // for device
-#include <sycl/event.hpp>         // for event
-#include <sycl/ext/codeplay/experimental/fusion_properties.hpp> // for accessor
+#include <sycl/access/access.hpp>               // for access_mode
+#include <sycl/accessor.hpp>                    // for host_acce...
+#include <sycl/accessor.hpp>                    // for accessor
+#include <sycl/context.hpp>                     // for context
+#include <sycl/device.hpp>                      // for device
+#include <sycl/event.hpp>                       // for event
 #include <sycl/ext/oneapi/weak_object.hpp>      // for weak_object
 #include <sycl/kernel.hpp>                      // for kernel
 #include <sycl/kernel_bundle.hpp>               // for kernel_id

diff --git a/sycl/include/sycl/ext/oneapi/weak_object.hpp b/sycl/include/sycl/ext/oneapi/weak_object.hpp
@@ -8,13 +8,12 @@
 
 #pragma once
 
-#include <sycl/access/access.hpp>     // for target, mode
-#include <sycl/accessor.hpp>          // for accessor
-#include <sycl/buffer.hpp>            // for buffer
-#include <sycl/detail/impl_utils.hpp> // for createSyc...
-#include <sycl/detail/memcpy.hpp>     // for detail
-#include <sycl/exception.hpp>         // for make_erro...
-#include <sycl/ext/codeplay/experimental/fusion_properties.hpp> // for buffer
+#include <sycl/access/access.hpp>               // for target, mode
+#include <sycl/accessor.hpp>                    // for accessor
+#include <sycl/buffer.hpp>                      // for buffer
+#include <sycl/detail/impl_utils.hpp>           // for createSyc...
+#include <sycl/detail/memcpy.hpp>               // for detail
+#include <sycl/exception.hpp>                   // for make_erro...
 #include <sycl/ext/oneapi/weak_object_base.hpp> // for weak_obje...
 #include <sycl/range.hpp>                       // for range
 #include <sycl/stream.hpp>                      // for stream

diff --git a/sycl/include/sycl/info/ext_codeplay_device_traits.def b/sycl/include/sycl/info/ext_codeplay_device_traits.def
@@ -2,6 +2,7 @@
 #define __SYCL_PARAM_TRAITS_TEMPLATE_SPEC_NEEDS_UNDEF
 #define __SYCL_PARAM_TRAITS_TEMPLATE_SPEC __SYCL_PARAM_TRAITS_SPEC
 #endif
+// TODO(#15184): Remove the fusion aspect in the next ABI-breaking window.
 __SYCL_PARAM_TRAITS_SPEC(ext::codeplay::experimental,device, supports_fusion, bool, __SYCL_TRAIT_HANDLED_IN_RT)
 __SYCL_PARAM_TRAITS_SPEC(
     ext::codeplay::experimental, device, max_registers_per_work_group, uint32_t,

diff --git a/sycl/include/sycl/properties/all_properties.hpp b/sycl/include/sycl/properties/all_properties.hpp
@@ -8,7 +8,6 @@
 
 #pragma once
 
-#include <sycl/ext/codeplay/experimental/fusion_properties.hpp>
 #include <sycl/properties/accessor_properties.hpp>
 #include <sycl/properties/buffer_properties.hpp>
 #include <sycl/properties/image_properties.hpp>

diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp
@@ -2576,6 +2576,7 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   /// Equivalent to
   /// `has_property<ext::codeplay::experimental::property::queue::enable_fusion>()`.
   ///
+  // TODO(#15184) Remove this function in the next ABI-breaking window.
   bool ext_codeplay_supports_fusion() const;
 
 // Clean KERNELFUNC macros.

diff --git a/sycl/include/sycl/reduction.hpp b/sycl/include/sycl/reduction.hpp
@@ -8,24 +8,23 @@
 
 #pragma once
 
-#include <sycl/access/access.hpp>              // for address_s...
-#include <sycl/accessor.hpp>                   // for local_acc...
-#include <sycl/aspects.hpp>                    // for aspect
-#include <sycl/atomic.hpp>                     // for IsValidAt...
-#include <sycl/atomic_ref.hpp>                 // for atomic_ref
-#include <sycl/buffer.hpp>                     // for buffer
-#include <sycl/builtins.hpp>                   // for min
-#include <sycl/detail/export.hpp>              // for __SYCL_EX...
-#include <sycl/detail/generic_type_traits.hpp> // for is_sgenfloat
-#include <sycl/detail/impl_utils.hpp>          // for createSyc...
-#include <sycl/detail/item_base.hpp>           // for id
-#include <sycl/detail/reduction_forward.hpp>   // for strategy
-#include <sycl/detail/tuple.hpp>               // for make_tuple
-#include <sycl/device.hpp>                     // for device
-#include <sycl/event.hpp>                      // for event
-#include <sycl/exception.hpp>                  // for make_erro...
-#include <sycl/exception_list.hpp>             // for queue_impl
-#include <sycl/ext/codeplay/experimental/fusion_properties.hpp> // for buffer
+#include <sycl/access/access.hpp>                   // for address_s...
+#include <sycl/accessor.hpp>                        // for local_acc...
+#include <sycl/aspects.hpp>                         // for aspect
+#include <sycl/atomic.hpp>                          // for IsValidAt...
+#include <sycl/atomic_ref.hpp>                      // for atomic_ref
+#include <sycl/buffer.hpp>                          // for buffer
+#include <sycl/builtins.hpp>                        // for min
+#include <sycl/detail/export.hpp>                   // for __SYCL_EX...
+#include <sycl/detail/generic_type_traits.hpp>      // for is_sgenfloat
+#include <sycl/detail/impl_utils.hpp>               // for createSyc...
+#include <sycl/detail/item_base.hpp>                // for id
+#include <sycl/detail/reduction_forward.hpp>        // for strategy
+#include <sycl/detail/tuple.hpp>                    // for make_tuple
+#include <sycl/device.hpp>                          // for device
+#include <sycl/event.hpp>                           // for event
+#include <sycl/exception.hpp>                       // for make_erro...
+#include <sycl/exception_list.hpp>                  // for queue_impl
 #include <sycl/group.hpp>                           // for workGroup...
 #include <sycl/group_algorithm.hpp>                 // for reduce_ov...
 #include <sycl/handler.hpp>                         // for handler

diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt
@@ -153,6 +153,8 @@ function(add_sycl_rt_library LIB_NAME LIB_OBJ_NAME)
     target_include_directories(${LIB_OBJ_NAME} PRIVATE ${SYCL_JIT_INCLUDE_DIRS})
     set_property(GLOBAL APPEND PROPERTY SYCL_TOOLCHAIN_INSTALL_COMPONENTS
        sycl-jit)
+    target_compile_definitions(${LIB_NAME} PRIVATE SYCL_EXT_JIT_ENABLE)
+    target_compile_definitions(${LIB_OBJ_NAME} PRIVATE SYCL_EXT_JIT_ENABLE)
   endif(SYCL_ENABLE_EXTENSION_JIT)
 
   find_package(Threads REQUIRED)
@@ -237,7 +239,6 @@ set(SYCL_COMMON_SOURCES
     "detail/event_impl.cpp"
     "detail/filter_selector_impl.cpp"
     "detail/fusion/fusion_wrapper.cpp"
-    "detail/fusion/fusion_wrapper_impl.cpp"
     "detail/global_handler.cpp"
     "detail/graph_impl.cpp"
     "detail/helpers.cpp"

diff --git a/sycl/source/detail/device_info.hpp b/sycl/source/detail/device_info.hpp
@@ -1173,31 +1173,13 @@ struct get_device_info_impl<bool, info::device::usm_system_allocations> {
 };
 
 // Specialization for kernel fusion support
+// TODO(#15184): Remove this aspect in the next ABI-breaking window.
 template <>
 struct get_device_info_impl<
     bool, ext::codeplay::experimental::info::device::supports_fusion> {
   static bool get(const DeviceImplPtr &Dev) {
-#if SYCL_EXT_CODEPLAY_KERNEL_FUSION
-    // If the JIT library can't be loaded or entry points in the JIT library
-    // can't be resolved, fusion is not available.
-    if (!jit_compiler::get_instance().isAvailable()) {
-      return false;
-    }
-    // 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::ext_oneapi_cuda) ||
-           (Dev->getBackend() == backend::ext_oneapi_hip);
-#else  // SYCL_EXT_CODEPLAY_KERNEL_FUSION
     (void)Dev;
     return false;
-#endif // SYCL_EXT_CODEPLAY_KERNEL_FUSION
   }
 };