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

@@ -259,10 +259,19 @@ 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
+  // TODO The handler-less path does not support reductions yet.
+  if constexpr (sizeof...(ReductionsT) == 0) {
+    detail::submit_kernel_direct<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)...);
+    });
+  }
 }
 
 template <typename KernelName = sycl::detail::auto_name, int Dimensions,
@@ -283,10 +292,23 @@ 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
+  // TODO The handler-less path does not support reductions yet.
+  if constexpr (sizeof...(ReductionsT) == 0) {
+    ext::oneapi::experimental::detail::LaunchConfigAccess<nd_range<Dimensions>,
+                                                          Properties>
+        ConfigAccess(Config);
+    detail::submit_kernel_direct<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)...);
+    });
+  }
 }
 
 template <int Dimensions, typename... ArgsT>

sycl/include/sycl/khr/free_function_commands.hpp

@@ -153,27 +153,45 @@ 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
+  detail::submit_kernel_direct(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
+  detail::submit_kernel_direct(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
+  detail::submit_kernel_direct(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

@@ -62,6 +62,20 @@ template <backend BackendName, class SyclObjectT>
 auto get_native(const SyclObjectT &Obj)
     -> backend_return_t<BackendName, SyclObjectT>;
 
+template <int Dims>
+event __SYCL_EXPORT submit_kernel_direct_with_event_impl(
+    const queue &Queue, const nd_range<Dims> &Range,
+    std::shared_ptr<detail::HostKernelBase> &HostKernel,
+    detail::DeviceKernelInfo *DeviceKernelInfo,
+    const detail::code_location &CodeLoc, bool IsTopCodeLoc);
+
+template <int Dims>
+void __SYCL_EXPORT submit_kernel_direct_without_event_impl(
+    const queue &Queue, const nd_range<Dims> &Range,
+    std::shared_ptr<detail::HostKernelBase> &HostKernel,
+    detail::DeviceKernelInfo *DeviceKernelInfo,
+    const detail::code_location &CodeLoc, bool IsTopCodeLoc);
+
 namespace detail {
 class queue_impl;
 
@@ -141,6 +155,51 @@ class __SYCL_EXPORT SubmissionInfo {
 };
 
 } // namespace v1
+
+template <typename KernelName = detail::auto_name, bool EventNeeded = false,
+          typename PropertiesT, typename KernelType, int Dims>
+auto submit_kernel_direct(
+    const queue &Queue, PropertiesT Props, const nd_range<Dims> &Range,
+    const KernelType &KernelFunc,
+    const detail::code_location &CodeLoc = detail::code_location::current()) {
+  // 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);
+
+  if constexpr (EventNeeded) {
+    return submit_kernel_direct_with_event_impl(
+        Queue, Range, HostKernel, DeviceKernelInfoPtr,
+        TlsCodeLocCapture.query(), TlsCodeLocCapture.isToplevel());
+  } else {
+    submit_kernel_direct_without_event_impl(
+        Queue, Range, HostKernel, DeviceKernelInfoPtr,
+        TlsCodeLocCapture.query(), TlsCodeLocCapture.isToplevel());
+  }
+}
+
 } // namespace detail
 
 namespace ext ::oneapi ::experimental {
@@ -3203,11 +3262,21 @@ 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
+    // TODO The handler-less path does not support reductions yet.
+    if constexpr (sizeof...(RestT) == 1) {
+      return detail::submit_kernel_direct<KernelName, true>(
+          *this, ext::oneapi::experimental::empty_properties_t{}, Range,
+          Rest...);
+    } else
+#endif
+    {
+      return submit(
+          [&](handler &CGH) {
+            CGH.template parallel_for<KernelName>(Range, Rest...);
+          },
+          TlsCodeLocCapture.query());
+    }
   }
 
   /// parallel_for version with a kernel represented as a lambda + nd_range that

sycl/source/detail/queue_impl.cpp

@@ -417,6 +417,106 @@ queue_impl::submit_impl(const detail::type_erased_cgfo_ty &CGF,
   return EventImpl;
 }
 
+detail::EventImplPtr queue_impl::submit_kernel_direct_impl(
+    const NDRDescT &NDRDesc,
+    std::shared_ptr<detail::HostKernelBase> &HostKernel,
+    detail::DeviceKernelInfo *DeviceKernelInfo, bool CallerNeedsEvent,
+    const detail::code_location &CodeLoc, bool IsTopCodeLoc) {
+
+  KernelData KData;
+
+  KData.setDeviceKernelInfoPtr(DeviceKernelInfo);
+  KData.setKernelFunc(HostKernel->getPtr());
+  KData.setNDRDesc(NDRDesc);
+
+  auto SubmitKernelFunc =
+      [&](detail::CG::StorageInitHelper &CGData) -> EventImplPtr {
+    std::unique_ptr<detail::CG> CommandGroup;
+    std::vector<std::shared_ptr<detail::stream_impl>> StreamStorage;
+    std::vector<std::shared_ptr<const void>> AuxiliaryResources;
+
+    KData.extractArgsAndReqsFromLambda();
+
+    CommandGroup.reset(new detail::CGExecKernel(
+        KData.getNDRDesc(), HostKernel,
+        nullptr, // Kernel
+        nullptr, // KernelBundle
+        std::move(CGData), std::move(KData).getArgs(),
+        *KData.getDeviceKernelInfoPtr(), std::move(StreamStorage),
+        std::move(AuxiliaryResources), detail::CGType::Kernel,
+        UR_KERNEL_CACHE_CONFIG_DEFAULT,
+        false, // KernelIsCooperative
+        false, // KernelUsesClusterLaunch
+        0,     // KernelWorkGroupMemorySize
+        CodeLoc));
+    CommandGroup->MIsTopCodeLoc = IsTopCodeLoc;
+
+    return detail::Scheduler::getInstance().addCG(std::move(CommandGroup),
+                                                  *this, true);
+  };
+
+  return submit_direct(CallerNeedsEvent, SubmitKernelFunc);
+}
+
+template <typename SubmitCommandFuncType>
+detail::EventImplPtr
+queue_impl::submit_direct(bool CallerNeedsEvent,
+                          SubmitCommandFuncType &SubmitCommandFunc) {
+  detail::CG::StorageInitHelper CGData;
+  std::unique_lock<std::mutex> Lock(MMutex);
+
+  // Graphs are not supported yet for the no-handler path
+  assert(!hasCommandGraph());
+
+  // Set the No Last Event Mode to false, since the no-handler path
+  // does not support it yet.
+  MNoLastEventMode.store(false, std::memory_order_relaxed);
+
+  // Used by queue_empty() and getLastEvent()
+  MEmpty.store(false, std::memory_order_release);
+
+  // Sync with an external event
+  std::optional<event> ExternalEvent = popExternalEvent();
+  if (ExternalEvent) {
+    CGData.MEvents.push_back(getSyclObjImpl(*ExternalEvent));
+  }
+
+  // Sync with the last event for in order queue
+  EventImplPtr &LastEvent = MDefaultGraphDeps.LastEventPtr;
+  if (isInOrder() && LastEvent) {
+    CGData.MEvents.push_back(LastEvent);
+  }
+
+  // Barrier and un-enqueued commands synchronization for out or order queue
+  if (!isInOrder()) {
+    MMissedCleanupRequests.unset(
+        [&](MissedCleanupRequestsType &MissedCleanupRequests) {
+          for (auto &UpdatedGraph : MissedCleanupRequests)
+            doUnenqueuedCommandCleanup(UpdatedGraph);
+          MissedCleanupRequests.clear();
+        });
+
+    if (MDefaultGraphDeps.LastBarrier &&
+        !MDefaultGraphDeps.LastBarrier->isEnqueued()) {
+      CGData.MEvents.push_back(MDefaultGraphDeps.LastBarrier);
+    }
+  }
+
+  EventImplPtr EventImpl = SubmitCommandFunc(CGData);
+
+  // Sync with the last event for in order queue
+  if (isInOrder() && !EventImpl->isDiscarded()) {
+    LastEvent = EventImpl;
+  }
+
+  // Barrier and un-enqueued commands synchronization for out or order queue
+  if (!isInOrder() && !EventImpl->isEnqueued()) {
+    MDefaultGraphDeps.UnenqueuedCmdEvents.push_back(EventImpl);
+  }
+
+  return CallerNeedsEvent ? EventImpl : nullptr;
+}
+
 template <typename HandlerFuncT>
 event queue_impl::submitWithHandler(const std::vector<event> &DepEvents,
                                     bool CallerNeedsEvent,