Extract common logic from single_task and parallel_for handler-less submit,

add a kernel move unit test for single_task

Author: slawekptak

sycl/include/sycl/khr/free_function_commands.hpp

@@ -319,8 +319,9 @@ void launch_task(handler &h, const KernelType &k) {
   h.single_task(k);
 }
 
-template <typename KernelType>
-void launch_task(const sycl::queue &q, const KernelType &k,
+template <typename KernelType, typename = typename std::enable_if_t<
+                                   enable_kernel_function_overload<KernelType>>>
+void launch_task(const sycl::queue &q, KernelType &&k,
                  const sycl::detail::code_location &codeLoc =
                      sycl::detail::code_location::current()) {
   // TODO The handler-less path does not support kernel function properties
@@ -331,7 +332,8 @@ void launch_task(const sycl::queue &q, const KernelType &k,
                 !(detail::KernelLambdaHasKernelHandlerArgT<KernelType,
                                                            void>::value)) {
     detail::submit_kernel_direct_single_task(
-        q, ext::oneapi::experimental::empty_properties_t{}, k, codeLoc);
+        q, ext::oneapi::experimental::empty_properties_t{},
+        std::forward<KernelType>(k), codeLoc);
   } else {
     submit(q, [&](handler &h) { launch_task<KernelType>(h, k); }, codeLoc);
   }

sycl/include/sycl/queue.hpp

@@ -157,70 +157,12 @@ class __SYCL_EXPORT SubmissionInfo {
 
 } // namespace v1
 
-template <typename KernelName = detail::auto_name, bool EventNeeded = false,
+template <detail::WrapAs WrapAs, typename LambdaArgType,
+          typename KernelName = detail::auto_name, bool EventNeeded = false,
           typename PropertiesT, typename KernelTypeUniversalRef, int Dims>
-auto submit_kernel_direct_parallel_for(
-    const queue &Queue, PropertiesT Props, const nd_range<Dims> &Range,
-    KernelTypeUniversalRef &&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 KernelType =
-      std::remove_const_t<std::remove_reference_t<KernelTypeUniversalRef>>;
-
-  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>;
-
-#ifndef __SYCL_DEVICE_ONLY__
-  detail::checkValueRange<Dims>(Range);
-#endif
-
-  detail::KernelWrapper<detail::WrapAs::parallel_for, NameT, KernelType,
-                        TransformedArgType, PropertiesT>::wrap(KernelFunc);
-
-  HostKernelRef<KernelType, KernelTypeUniversalRef, TransformedArgType, Dims>
-      HostKernel(std::forward<KernelTypeUniversalRef>(KernelFunc));
-
-  // Instantiating the kernel on the host improves debugging.
-  // Passing this pointer to another translation unit prevents optimization.
-#ifndef NDEBUG
-  // TODO: call library to prevent dropping call due to optimization
-  (void)
-      detail::GetInstantiateKernelOnHostPtr<KernelType, LambdaArgType, Dims>();
-#endif
-
-  detail::DeviceKernelInfo *DeviceKernelInfoPtr =
-      &detail::getDeviceKernelInfo<NameT>();
-
-  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());
-  }
-}
-
-template <typename KernelName = detail::auto_name, bool EventNeeded = false,
-          typename PropertiesT, typename KernelTypeUniversalRef>
-auto submit_kernel_direct_single_task(
+auto submit_kernel_direct(
     const queue &Queue, [[maybe_unused]] PropertiesT Props,
-    KernelTypeUniversalRef &&KernelFunc,
+    const nd_range<Dims> &Range, KernelTypeUniversalRef &&KernelFunc,
     const detail::code_location &CodeLoc = detail::code_location::current()) {
   // TODO Properties not supported yet
   static_assert(
@@ -235,17 +177,18 @@ auto submit_kernel_direct_single_task(
   using NameT =
       typename detail::get_kernel_name_t<KernelName, KernelType>::name;
 
-  detail::KernelWrapper<detail::WrapAs::single_task, NameT, KernelType, void,
+  detail::KernelWrapper<WrapAs, NameT, KernelType, LambdaArgType,
                         PropertiesT>::wrap(KernelFunc);
 
-  HostKernelRef<KernelType, KernelTypeUniversalRef, void, 1> HostKernel(
-      std::forward<KernelTypeUniversalRef>(KernelFunc));
+  HostKernelRef<KernelType, KernelTypeUniversalRef, LambdaArgType, Dims>
+      HostKernel(std::forward<KernelTypeUniversalRef>(KernelFunc));
 
   // Instantiating the kernel on the host improves debugging.
   // Passing this pointer to another translation unit prevents optimization.
 #ifndef NDEBUG
   // TODO: call library to prevent dropping call due to optimization.
-  (void)detail::GetInstantiateKernelOnHostPtr<KernelType, void, 1>();
+  (void)
+      detail::GetInstantiateKernelOnHostPtr<KernelType, LambdaArgType, Dims>();
 #endif
 
   detail::DeviceKernelInfo *DeviceKernelInfoPtr =
@@ -269,15 +212,57 @@ auto submit_kernel_direct_single_task(
 
   if constexpr (EventNeeded) {
     return submit_kernel_direct_with_event_impl(
-        Queue, nd_range<1>{1, 1}, HostKernel, DeviceKernelInfoPtr,
+        Queue, Range, HostKernel, DeviceKernelInfoPtr,
         TlsCodeLocCapture.query(), TlsCodeLocCapture.isToplevel());
   } else {
     submit_kernel_direct_without_event_impl(
-        Queue, nd_range<1>{1, 1}, HostKernel, DeviceKernelInfoPtr,
+        Queue, Range, HostKernel, DeviceKernelInfoPtr,
         TlsCodeLocCapture.query(), TlsCodeLocCapture.isToplevel());
   }
 }
 
+template <typename KernelName = detail::auto_name, bool EventNeeded = false,
+          typename PropertiesT, typename KernelTypeUniversalRef, int Dims>
+auto submit_kernel_direct_parallel_for(
+    const queue &Queue, PropertiesT Props, const nd_range<Dims> &Range,
+    KernelTypeUniversalRef &&KernelFunc,
+    const detail::code_location &CodeLoc = detail::code_location::current()) {
+
+  using KernelType =
+      std::remove_const_t<std::remove_reference_t<KernelTypeUniversalRef>>;
+
+  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>;
+
+#ifndef __SYCL_DEVICE_ONLY__
+  detail::checkValueRange<Dims>(Range);
+#endif
+
+  return submit_kernel_direct<detail::WrapAs::parallel_for, TransformedArgType,
+                              KernelName, EventNeeded, PropertiesT,
+                              KernelTypeUniversalRef, Dims>(
+      Queue, Props, Range, std::forward<KernelTypeUniversalRef>(KernelFunc),
+      CodeLoc);
+}
+
+template <typename KernelName = detail::auto_name, bool EventNeeded = false,
+          typename PropertiesT, typename KernelTypeUniversalRef>
+auto submit_kernel_direct_single_task(
+    const queue &Queue, PropertiesT Props, KernelTypeUniversalRef &&KernelFunc,
+    const detail::code_location &CodeLoc = detail::code_location::current()) {
+
+  return submit_kernel_direct<detail::WrapAs::single_task, void, KernelName,
+                              EventNeeded, PropertiesT, KernelTypeUniversalRef,
+                              1>(
+      Queue, Props, nd_range<1>{1, 1},
+      std::forward<KernelTypeUniversalRef>(KernelFunc), CodeLoc);
+}
+
 } // namespace detail
 
 namespace ext ::oneapi ::experimental {

sycl/unittests/Extensions/FreeFunctionCommands/FreeFunctionCommandsEvents.cpp

@@ -78,6 +78,7 @@ class FreeFunctionCommandsEventsTests : public ::testing::Test {
 protected:
   void SetUp() override {
     counter_urEnqueueKernelLaunch = 0;
+    counter_urEnqueueKernelLaunchWithEvent = 0;
     counter_urUSMEnqueueMemcpy = 0;
     counter_urUSMEnqueueFill = 0;
     counter_urUSMEnqueuePrefetch = 0;
@@ -281,6 +282,57 @@ TEST_F(FreeFunctionCommandsEventsTests,
   ASSERT_EQ(counter_urEnqueueKernelLaunchWithEvent, size_t{1});
 }
 
+TEST_F(FreeFunctionCommandsEventsTests, LaunchTaskShortcutMoveKernel) {
+  mock::getCallbacks().set_replace_callback("urEnqueueKernelLaunch",
+                                            &redefined_urEnqueueKernelLaunch);
+
+  TestMoveFunctor::MoveCtorCalls = 0;
+  TestMoveFunctor MoveOnly;
+  std::mutex CvMutex;
+  std::condition_variable Cv;
+  bool ready = false;
+
+  // This kernel submission uses scheduler-bypass path, so the HostKernel
+  // shouldn't be constructed.
+
+  sycl::khr::launch_task(Queue, std::move(MoveOnly));
+
+  ASSERT_EQ(TestMoveFunctor::MoveCtorCalls, 0);
+  ASSERT_EQ(counter_urEnqueueKernelLaunch, size_t{1});
+
+  // Another kernel submission is queued behind a host task,
+  // to force the scheduler-based submission. In this case, the HostKernel
+  // should be constructed.
+
+  // Replace the callback with an event based one, since the scheduler
+  // needs to create an event internally
+  mock::getCallbacks().set_replace_callback(
+      "urEnqueueKernelLaunch", &redefined_urEnqueueKernelLaunchWithEvent);
+
+  Queue.submit([&](sycl::handler &CGH) {
+    CGH.host_task([&] {
+      std::unique_lock<std::mutex> lk(CvMutex);
+      Cv.wait(lk, [&ready] { return ready; });
+    });
+  });
+
+  sycl::khr::launch_task(Queue, std::move(MoveOnly));
+
+  {
+    std::unique_lock<std::mutex> lk(CvMutex);
+    ready = true;
+  }
+  Cv.notify_one();
+
+  Queue.wait();
+
+  // Move ctor for TestMoveFunctor is called during move construction of
+  // HostKernel. Copy ctor is called by InstantiateKernelOnHost, can't delete
+  // it.
+  ASSERT_EQ(TestMoveFunctor::MoveCtorCalls, 1);
+  ASSERT_EQ(counter_urEnqueueKernelLaunchWithEvent, size_t{1});
+}
+
 TEST_F(FreeFunctionCommandsEventsTests, SubmitLaunchGroupedKernelNoEvent) {
   mock::getCallbacks().set_replace_callback("urEnqueueKernelLaunch",
                                             &redefined_urEnqueueKernelLaunch);