Handler-less kernel submit path (parallel_for with nd_range)

sycl/cmake/modules/AddSYCLUnitTest.cmake

@@ -1,6 +1,6 @@
 # Internal function to create SYCL unit tests with code reuse
-# add_sycl_unittest_internal(test_dirname SHARED|OBJECT is_preview file1.cpp, file2.cpp ...)
-function(add_sycl_unittest_internal test_dirname link_variant is_preview)
+# add_sycl_unittest_internal(test_dirname SHARED|OBJECT is_preview is_no_cgh file1.cpp, file2.cpp ...)
+function(add_sycl_unittest_internal test_dirname link_variant is_preview is_no_cgh)
   # Enable exception handling for these unit tests
   set(LLVM_REQUIRES_EH ON)
   set(LLVM_REQUIRES_RTTI ON)
@@ -34,7 +34,11 @@ function(add_sycl_unittest_internal test_dirname link_variant is_preview)
   # Chaning CMAKE_CURRENT_BINARY_DIR should not affect this variable in its
   # parent scope.
   if (${is_preview})
-    set(CMAKE_CURRENT_BINARY_DIR "${CMAKE_CURRENT_BINARY_DIR}/Preview")
+      set(CMAKE_CURRENT_BINARY_DIR "${CMAKE_CURRENT_BINARY_DIR}/Preview")
+  endif()
+
+  if (${is_no_cgh})
+    set(CMAKE_CURRENT_BINARY_DIR "${CMAKE_CURRENT_BINARY_DIR}/NoCGH")
   endif()
 
   if ("${link_variant}" MATCHES "SHARED")
@@ -65,6 +69,18 @@ function(add_sycl_unittest_internal test_dirname link_variant is_preview)
     set(sycl_cache_suffix "_preview")
   endif()
 
+  if (${is_no_cgh})
+    set(sycl_cache_suffix "_no_cgh")
+  endif()
+
+  if (${is_no_cgh})
+    target_compile_definitions(
+      ${test_dirname}
+      PRIVATE
+        __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
+    )
+  endif()
+
   if (SYCL_ENABLE_XPTI_TRACING)
     target_compile_definitions(${test_dirname}
       PRIVATE XPTI_ENABLE_INSTRUMENTATION XPTI_STATIC_LIBRARY)
@@ -150,7 +166,6 @@ function(add_sycl_unittest_internal test_dirname link_variant is_preview)
         -Wno-inconsistent-missing-override
     )
   endif()
-
   target_compile_definitions(${test_dirname} PRIVATE SYCL_DISABLE_FSYCL_SYCLHPP_WARNING)
 endfunction()
 
@@ -160,6 +175,7 @@ endfunction()
 # the SYCL preview features enabled.
 # Produces two binaries, named `basename(test_name_prefix_non_preview)` and `basename(test_name_prefix_preview)`
 macro(add_sycl_unittest test_name_prefix link_variant)
-  add_sycl_unittest_internal(${test_name_prefix}_non_preview ${link_variant} FALSE ${ARGN})
-  add_sycl_unittest_internal(${test_name_prefix}_preview ${link_variant} TRUE ${ARGN})
+  add_sycl_unittest_internal(${test_name_prefix}_non_preview ${link_variant} FALSE FALSE ${ARGN})
+  add_sycl_unittest_internal(${test_name_prefix}_no_cgh ${link_variant} FALSE TRUE ${ARGN})
+  add_sycl_unittest_internal(${test_name_prefix}_preview ${link_variant} TRUE FALSE ${ARGN})
 endmacro()

sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp

@@ -109,6 +109,17 @@ event submit_with_event_impl(const queue &Q, PropertiesT Props,
                              const sycl::detail::code_location &CodeLoc) {
   return Q.submit_with_event(Props, detail::type_erased_cgfo_ty{CGF}, CodeLoc);
 }
+
+template <typename KernelName, typename PropertiesT, typename KernelType,
+          int Dims>
+void submit_kernel_direct_impl(const queue &Q, PropertiesT Props,
+                               nd_range<Dims> Range,
+                               const KernelType &KernelFunc,
+                               const sycl::detail::code_location &CodeLoc) {
+  Q.submit_kernel_direct_without_event<KernelName, PropertiesT, KernelType,
+                                       Dims>(Props, Range, KernelFunc, CodeLoc);
+}
+
 } // namespace detail
 
 template <typename CommandGroupFunc, typename PropertiesT>
@@ -126,6 +137,17 @@ void submit(const queue &Q, CommandGroupFunc &&CGF,
   submit(Q, empty_properties_t{}, std::forward<CommandGroupFunc>(CGF), CodeLoc);
 }
 
+template <typename KernelName = sycl::detail::auto_name, typename PropertiesT,
+          typename KernelType, int Dims>
+void submit(const queue &Q, PropertiesT Props, nd_range<Dims> Range,
+            const KernelType &KernelFunc,
+            const sycl::detail::code_location &CodeLoc =
+                sycl::detail::code_location::current()) {
+  sycl::ext::oneapi::experimental::detail::submit_kernel_direct_impl<
+      KernelName, PropertiesT, KernelType, Dims>(Q, Props, Range, KernelFunc,
+                                                 CodeLoc);
+}
+
 template <typename CommandGroupFunc, typename PropertiesT>
 event submit_with_event(const queue &Q, PropertiesT Props,
                         CommandGroupFunc &&CGF,
@@ -259,10 +281,18 @@ template <typename KernelName = sycl::detail::auto_name, int Dimensions,
           typename KernelType, typename... ReductionsT>
 void nd_launch(queue Q, nd_range<Dimensions> Range, const KernelType &KernelObj,
                ReductionsT &&...Reductions) {
-  submit(std::move(Q), [&](handler &CGH) {
-    nd_launch<KernelName>(CGH, Range, KernelObj,
-                          std::forward<ReductionsT>(Reductions)...);
-  });
+#ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
+  if constexpr (sizeof...(ReductionsT) == 0) {
+    submit<KernelName>(std::move(Q), empty_properties_t{}, Range, KernelObj);
+  } else {
+#endif
+    submit(std::move(Q), [&](handler &CGH) {
+      nd_launch<KernelName>(CGH, Range, KernelObj,
+                            std::forward<ReductionsT>(Reductions)...);
+    });
+#ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
+  }
+#endif
 }
 
 template <typename KernelName = sycl::detail::auto_name, int Dimensions,
@@ -283,10 +313,22 @@ template <typename KernelName = sycl::detail::auto_name, int Dimensions,
           typename Properties, typename KernelType, typename... ReductionsT>
 void nd_launch(queue Q, launch_config<nd_range<Dimensions>, Properties> Config,
                const KernelType &KernelObj, ReductionsT &&...Reductions) {
-  submit(std::move(Q), [&](handler &CGH) {
-    nd_launch<KernelName>(CGH, Config, KernelObj,
-                          std::forward<ReductionsT>(Reductions)...);
-  });
+#ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
+  if constexpr (sizeof...(ReductionsT) == 0) {
+    ext::oneapi::experimental::detail::LaunchConfigAccess<nd_range<Dimensions>,
+                                                          Properties>
+        ConfigAccess(Config);
+    submit<KernelName>(std::move(Q), ConfigAccess.getProperties(),
+                       ConfigAccess.getRange(), KernelObj);
+  } else {
+#endif
+    submit(std::move(Q), [&](handler &CGH) {
+      nd_launch<KernelName>(CGH, Config, KernelObj,
+                            std::forward<ReductionsT>(Reductions)...);
+    });
+#ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
+  }
+#endif
 }
 
 template <int Dimensions, typename... ArgsT>

sycl/include/sycl/khr/free_function_commands.hpp

@@ -153,27 +153,42 @@ void launch_grouped(const queue &q, range<1> r, range<1> size,
                     const KernelType &k,
                     const sycl::detail::code_location &codeLoc =
                         sycl::detail::code_location::current()) {
+#ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
+  submit(q, ext::oneapi::experimental::empty_properties_t{},
+         nd_range<1>(r, size), k);
+#else
   submit(
       q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
       codeLoc);
+#endif
 }
 template <typename KernelType>
 void launch_grouped(const queue &q, range<2> r, range<2> size,
                     const KernelType &k,
                     const sycl::detail::code_location &codeLoc =
                         sycl::detail::code_location::current()) {
+#ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
+  submit(q, ext::oneapi::experimental::empty_properties_t{},
+         nd_range<2>(r, size), k);
+#else
   submit(
       q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
       codeLoc);
+#endif
 }
 template <typename KernelType>
 void launch_grouped(const queue &q, range<3> r, range<3> size,
                     const KernelType &k,
                     const sycl::detail::code_location &codeLoc =
                         sycl::detail::code_location::current()) {
+#ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
+  submit(q, ext::oneapi::experimental::empty_properties_t{},
+         nd_range<3>(r, size), k);
+#else
   submit(
       q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
       codeLoc);
+#endif
 }
 
 template <typename... Args>

sycl/include/sycl/queue.hpp

@@ -158,6 +158,14 @@ template <typename CommandGroupFunc, typename PropertiesT>
 event submit_with_event_impl(const queue &Q, PropertiesT Props,
                              CommandGroupFunc &&CGF,
                              const sycl::detail::code_location &CodeLoc);
+
+template <typename KernelName, typename PropertiesT, typename KernelType,
+          int Dims>
+void submit_kernel_direct_impl(const queue &Q, PropertiesT Props,
+                               nd_range<Dims> Range,
+                               const KernelType &KernelFunc,
+                               const sycl::detail::code_location &CodeLoc);
+
 } // namespace detail
 } // namespace ext::oneapi::experimental
 
@@ -3203,11 +3211,20 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   parallel_for(nd_range<Dims> Range, RestT &&...Rest) {
     constexpr detail::code_location CodeLoc = getCodeLocation<KernelName>();
     detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
-    return submit(
-        [&](handler &CGH) {
-          CGH.template parallel_for<KernelName>(Range, Rest...);
-        },
-        TlsCodeLocCapture.query());
+#ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
+    if constexpr (sizeof...(RestT) == 1) {
+      return submit_kernel_direct_with_event<KernelName>(
+          ext::oneapi::experimental::empty_properties_t{}, Range, Rest...);
+    } else {
+#endif
+      return submit(
+          [&](handler &CGH) {
+            CGH.template parallel_for<KernelName>(Range, Rest...);
+          },
+          TlsCodeLocCapture.query());
+#ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
+    }
+#endif
   }
 
   /// parallel_for version with a kernel represented as a lambda + nd_range that
@@ -3586,6 +3603,12 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
       const queue &Q, PropertiesT Props, CommandGroupFunc &&CGF,
       const sycl::detail::code_location &CodeLoc);
 
+  template <typename KernelName, typename PropertiesT, typename KernelType,
+            int Dims>
+  friend void ext::oneapi::experimental::detail::submit_kernel_direct_impl(
+      const queue &Q, PropertiesT Props, nd_range<Dims> Range,
+      const KernelType &KernelFunc, const sycl::detail::code_location &CodeLoc);
+
   template <typename PropertiesT>
   void ProcessSubmitProperties(PropertiesT Props,
                                detail::v1::SubmissionInfo &SI) const {
@@ -3670,6 +3693,36 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
                                const detail::code_location &CodeLoc,
                                bool IsTopCodeLoc) const;
 
+  event submit_kernel_direct_with_event_impl(
+      nd_range<1> Range, std::shared_ptr<detail::HostKernelBase> &HostKernel,
+      detail::DeviceKernelInfo *DeviceKernelInfo,
+      const detail::code_location &CodeLoc, bool IsTopCodeLoc) const;
+
+  event submit_kernel_direct_with_event_impl(
+      nd_range<2> Range, std::shared_ptr<detail::HostKernelBase> &HostKernel,
+      detail::DeviceKernelInfo *DeviceKernelInfo,
+      const detail::code_location &CodeLoc, bool IsTopCodeLoc) const;
+
+  event submit_kernel_direct_with_event_impl(
+      nd_range<3> Range, std::shared_ptr<detail::HostKernelBase> &HostKernel,
+      detail::DeviceKernelInfo *DeviceKernelInfo,
+      const detail::code_location &CodeLoc, bool IsTopCodeLoc) const;
+
+  void submit_kernel_direct_without_event_impl(
+      nd_range<1> Range, std::shared_ptr<detail::HostKernelBase> &HostKernel,
+      detail::DeviceKernelInfo *DeviceKernelInfo,
+      const detail::code_location &CodeLoc, bool IsTopCodeLoc) const;
+
+  void submit_kernel_direct_without_event_impl(
+      nd_range<2> Range, std::shared_ptr<detail::HostKernelBase> &HostKernel,
+      detail::DeviceKernelInfo *DeviceKernelInfo,
+      const detail::code_location &CodeLoc, bool IsTopCodeLoc) const;
+
+  void submit_kernel_direct_without_event_impl(
+      nd_range<3> Range, std::shared_ptr<detail::HostKernelBase> &HostKernel,
+      detail::DeviceKernelInfo *DeviceKernelInfo,
+      const detail::code_location &CodeLoc, bool IsTopCodeLoc) const;
+
   /// A template-free version of submit_without_event as const member function.
   void submit_without_event_impl(const detail::type_erased_cgfo_ty &CGH,
                                  const detail::v1::SubmissionInfo &SubmitInfo,
@@ -3699,6 +3752,83 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
                                   TlsCodeLocCapture.isToplevel());
   }
 
+  template <typename KernelName = detail::auto_name, typename PropertiesT,
+            typename KernelType, int Dims>
+  event
+  submit_kernel_direct_with_event(PropertiesT Props, nd_range<Dims> Range,
+                                  const KernelType &KernelFunc,
+                                  const detail::code_location &CodeLoc =
+                                      detail::code_location::current()) const {
+    // TODO Properties not supported yet
+    (void)Props;
+    static_assert(
+        std::is_same_v<PropertiesT,
+                       ext::oneapi::experimental::empty_properties_t>,
+        "Setting properties not supported yet for no-CGH kernel submit.");
+    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
+
+    using NameT =
+        typename detail::get_kernel_name_t<KernelName, KernelType>::name;
+    using LambdaArgType =
+        sycl::detail::lambda_arg_type<KernelType, nd_item<Dims>>;
+    static_assert(
+        std::is_convertible_v<sycl::nd_item<Dims>, LambdaArgType>,
+        "Kernel argument of a sycl::parallel_for with sycl::nd_range "
+        "must be either sycl::nd_item or be convertible from sycl::nd_item");
+    using TransformedArgType = sycl::nd_item<Dims>;
+
+    std::shared_ptr<detail::HostKernelBase> HostKernel = std::make_shared<
+        detail::HostKernel<KernelType, TransformedArgType, Dims>>(KernelFunc);
+
+    detail::DeviceKernelInfo *DeviceKernelInfoPtr =
+        &detail::getDeviceKernelInfo<NameT>();
+
+    detail::KernelWrapper<detail::WrapAs::parallel_for, NameT, KernelType,
+                          TransformedArgType, PropertiesT>::wrap(KernelFunc);
+
+    return submit_kernel_direct_with_event_impl(
+        Range, HostKernel, DeviceKernelInfoPtr, TlsCodeLocCapture.query(),
+        TlsCodeLocCapture.isToplevel());
+  }
+
+  template <typename KernelName = detail::auto_name, typename PropertiesT,
+            typename KernelType, int Dims>
+  void submit_kernel_direct_without_event(
+      PropertiesT Props, nd_range<Dims> Range, const KernelType &KernelFunc,
+      const detail::code_location &CodeLoc =
+          detail::code_location::current()) const {
+    // TODO Properties not supported yet
+    (void)Props;
+    static_assert(
+        std::is_same_v<PropertiesT,
+                       ext::oneapi::experimental::empty_properties_t>,
+        "Setting properties not supported yet for no-CGH kernel submit.");
+    detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
+
+    using NameT =
+        typename detail::get_kernel_name_t<KernelName, KernelType>::name;
+    using LambdaArgType =
+        sycl::detail::lambda_arg_type<KernelType, nd_item<Dims>>;
+    static_assert(
+        std::is_convertible_v<sycl::nd_item<Dims>, LambdaArgType>,
+        "Kernel argument of a sycl::parallel_for with sycl::nd_range "
+        "must be either sycl::nd_item or be convertible from sycl::nd_item");
+    using TransformedArgType = sycl::nd_item<Dims>;
+
+    std::shared_ptr<detail::HostKernelBase> HostKernel = std::make_shared<
+        detail::HostKernel<KernelType, TransformedArgType, Dims>>(KernelFunc);
+
+    detail::DeviceKernelInfo *DeviceKernelInfoPtr =
+        &detail::getDeviceKernelInfo<NameT>();
+
+    detail::KernelWrapper<detail::WrapAs::parallel_for, NameT, KernelType,
+                          TransformedArgType, PropertiesT>::wrap(KernelFunc);
+
+    submit_kernel_direct_without_event_impl(
+        Range, HostKernel, DeviceKernelInfoPtr, TlsCodeLocCapture.query(),
+        TlsCodeLocCapture.isToplevel());
+  }
+
   /// Submits a command group function object to the queue, in order to be
   /// scheduled for execution on the device.
   ///