Consolidate the event-based and event-less functions

Author: slawekptak

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

@@ -262,8 +262,8 @@ void nd_launch(queue Q, nd_range<Dimensions> Range, const KernelType &KernelObj,
                ReductionsT &&...Reductions) {
 #ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
   if constexpr (sizeof...(ReductionsT) == 0) {
-    detail::submit_kernel_direct_without_event<KernelName>(
-        std::move(Q), empty_properties_t{}, Range, KernelObj);
+    detail::submit_kernel_direct<KernelName>(std::move(Q), empty_properties_t{},
+                                             Range, KernelObj);
   } else {
 #endif
     submit(std::move(Q), [&](handler &CGH) {
@@ -298,7 +298,7 @@ void nd_launch(queue Q, launch_config<nd_range<Dimensions>, Properties> Config,
     ext::oneapi::experimental::detail::LaunchConfigAccess<nd_range<Dimensions>,
                                                           Properties>
         ConfigAccess(Config);
-    detail::submit_kernel_direct_without_event<KernelName>(
+    detail::submit_kernel_direct<KernelName>(
         std::move(Q), ConfigAccess.getProperties(), ConfigAccess.getRange(),
         KernelObj);
   } else {

sycl/include/sycl/khr/free_function_commands.hpp

@@ -154,9 +154,9 @@ void launch_grouped(const queue &q, range<1> r, range<1> size,
                     const sycl::detail::code_location &codeLoc =
                         sycl::detail::code_location::current()) {
 #ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
-  detail::submit_kernel_direct_without_event(
-      q, ext::oneapi::experimental::empty_properties_t{}, nd_range<1>(r, size),
-      k);
+  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); },
@@ -169,9 +169,9 @@ void launch_grouped(const queue &q, range<2> r, range<2> size,
                     const sycl::detail::code_location &codeLoc =
                         sycl::detail::code_location::current()) {
 #ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
-  detail::submit_kernel_direct_without_event(
-      q, ext::oneapi::experimental::empty_properties_t{}, nd_range<2>(r, size),
-      k);
+  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); },
@@ -184,9 +184,9 @@ void launch_grouped(const queue &q, range<3> r, range<3> size,
                     const sycl::detail::code_location &codeLoc =
                         sycl::detail::code_location::current()) {
 #ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
-  detail::submit_kernel_direct_without_event(
-      q, ext::oneapi::experimental::empty_properties_t{}, nd_range<3>(r, size),
-      k);
+  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); },

sycl/include/sycl/queue.hpp

@@ -142,16 +142,9 @@ class __SYCL_EXPORT SubmissionInfo {
 
 } // namespace v1
 
-template <typename KernelName = detail::auto_name, typename PropertiesT,
-          typename KernelType, int Dims>
-event submit_kernel_direct_with_event(
-    const queue &Queue, PropertiesT Props, nd_range<Dims> Range,
-    const KernelType &KernelFunc,
-    const detail::code_location &CodeLoc = detail::code_location::current());
-
-template <typename KernelName = detail::auto_name, typename PropertiesT,
-          typename KernelType, int Dims>
-void submit_kernel_direct_without_event(
+template <typename KernelName = detail::auto_name, bool EventNeeded = false,
+          typename PropertiesT, typename KernelType, int Dims>
+auto submit_kernel_direct(
     const queue &Queue, PropertiesT Props, nd_range<Dims> Range,
     const KernelType &KernelFunc,
     const detail::code_location &CodeLoc = detail::code_location::current());
@@ -3228,7 +3221,7 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
     detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
 #ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
     if constexpr (sizeof...(RestT) == 1) {
-      return detail::submit_kernel_direct_with_event<KernelName>(
+      return detail::submit_kernel_direct<KernelName, true>(
           *this, ext::oneapi::experimental::empty_properties_t{}, Range,
           Rest...);
     } else {
@@ -3619,15 +3612,9 @@ 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 sycl::detail::submit_kernel_direct_without_event(
-      const queue &Queue, PropertiesT Props, nd_range<Dims> Range,
-      const KernelType &KernelFunc, const detail::code_location &CodeLoc);
-
-  template <typename KernelName, typename PropertiesT, typename KernelType,
-            int Dims>
-  friend event sycl::detail::submit_kernel_direct_with_event(
+  template <typename KernelName, bool EventNeeded, typename PropertiesT,
+            typename KernelType, int Dims>
+  friend auto sycl::detail::submit_kernel_direct(
       const queue &Queue, PropertiesT Props, nd_range<Dims> Range,
       const KernelType &KernelFunc, const detail::code_location &CodeLoc);
 
@@ -3913,12 +3900,11 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
 
 namespace detail {
 
-template <typename KernelName, typename PropertiesT, typename KernelType,
-          int Dims>
-event submit_kernel_direct_with_event(const queue &Queue, PropertiesT Props,
-                                      nd_range<Dims> Range,
-                                      const KernelType &KernelFunc,
-                                      const detail::code_location &CodeLoc) {
+template <typename KernelName, bool EventNeeded, typename PropertiesT,
+          typename KernelType, int Dims>
+auto submit_kernel_direct(const queue &Queue, PropertiesT Props,
+                          nd_range<Dims> Range, const KernelType &KernelFunc,
+                          const detail::code_location &CodeLoc) {
   // TODO Properties not supported yet
   (void)Props;
   static_assert(
@@ -3946,48 +3932,17 @@ event submit_kernel_direct_with_event(const queue &Queue, PropertiesT Props,
   detail::KernelWrapper<detail::WrapAs::parallel_for, NameT, KernelType,
                         TransformedArgType, PropertiesT>::wrap(KernelFunc);
 
-  return Queue.submit_kernel_direct_with_event_impl(
-      Range, HostKernel, DeviceKernelInfoPtr, TlsCodeLocCapture.query(),
-      TlsCodeLocCapture.isToplevel());
+  if constexpr (EventNeeded) {
+    return Queue.submit_kernel_direct_with_event_impl(
+        Range, HostKernel, DeviceKernelInfoPtr, TlsCodeLocCapture.query(),
+        TlsCodeLocCapture.isToplevel());
+  } else {
+    Queue.submit_kernel_direct_without_event_impl(
+        Range, HostKernel, DeviceKernelInfoPtr, TlsCodeLocCapture.query(),
+        TlsCodeLocCapture.isToplevel());
+  }
 }
 
-template <typename KernelName, typename PropertiesT, typename KernelType,
-          int Dims>
-void submit_kernel_direct_without_event(const queue &Queue, PropertiesT Props,
-                                        nd_range<Dims> Range,
-                                        const KernelType &KernelFunc,
-                                        const detail::code_location &CodeLoc) {
-  // 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);
-
-  Queue.submit_kernel_direct_without_event_impl(
-      Range, HostKernel, DeviceKernelInfoPtr, TlsCodeLocCapture.query(),
-      TlsCodeLocCapture.isToplevel());
-}
 } // namespace detail
 
 } // namespace _V1