Enable move semantic for kernels.

Author: Alexandr-Konovalov

sycl/include/sycl/detail/cg_types.hpp

@@ -239,10 +239,36 @@ class HostKernel : public HostKernelBase {
 class HostKernelRefBase : public HostKernelBase {
 public:
   virtual std::shared_ptr<HostKernelBase> takeOrCopyOwnership() const = 0;
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
+  // This function can't be called from old user code, because there is no
+  // HostKernelRef in old user code. So, make it empty.
+  void InstantiateKernelOnHost() override {}
+#endif
 };
 
-template <class KernelType, class KernelArgType, int Dims>
+// primary template for movable lambda objects
+template <class KernelType, class KernelTypeUniversalRef, class KernelArgType, int Dims>
 class HostKernelRef : public HostKernelRefBase {
+  KernelType &&MKernel;
+
+public:
+  HostKernelRef(KernelType &&Kernel) : MKernel(std::move(Kernel)) {}
+
+  virtual char *getPtr() override {
+    return const_cast<char *>(reinterpret_cast<const char *>(&MKernel));
+  }
+  virtual std::shared_ptr<HostKernelBase> takeOrCopyOwnership() const override {
+    std::shared_ptr<HostKernelBase> Kernel;
+    Kernel.reset(new HostKernel<KernelType, KernelArgType, Dims>(std::move(MKernel)));
+    return Kernel;
+  }
+
+  ~HostKernelRef() noexcept override = default;
+};
+
+// specialization for copyable lambda objects
+template <class KernelType, class KernelTypeUniversalRef, class KernelArgType, int Dims>
+class HostKernelRef<KernelType, KernelTypeUniversalRef&, KernelArgType, Dims> : public HostKernelRefBase {
   const KernelType &MKernel;
 
 public:

sycl/include/sycl/khr/free_function_commands.hpp

@@ -148,45 +148,56 @@ void launch_grouped(handler &h, range<3> r, range<3> size,
   h.parallel_for(nd_range<3>(r, size), k);
 }
 
-template <typename KernelType>
+template <typename KernelType,
+          // overload of launch_grouped for sycl::kernel must be preferred
+          typename = typename std::enable_if<!std::is_same<typename std::decay<KernelType>::type,
+                                                           sycl::kernel>::value
+                                             >::type>
 void launch_grouped(const queue &q, range<1> r, range<1> size,
-                    const KernelType &k,
+                    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);
+                               nd_range<1>(r, size), std::forward<KernelType>(k));
 #else
   submit(
       q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
       codeLoc);
 #endif
 }
-template <typename KernelType>
+template <typename KernelType,
+          typename = typename std::enable_if<!std::is_same<typename std::decay<KernelType>::type,
+                                                             sycl::kernel>::value
+                                             >::type>
 void launch_grouped(const queue &q, range<2> r, range<2> size,
-                    const KernelType &k,
+                    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);
+                               nd_range<2>(r, size),
+                               std::forward<KernelType>(k));
 #else
   submit(
       q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
       codeLoc);
 #endif
 }
-template <typename KernelType>
+template <typename KernelType,
+          typename = typename std::enable_if<!std::is_same<typename std::decay<KernelType>::type,
+                                                             sycl::kernel>::value
+                                             >::type>
 void launch_grouped(const queue &q, range<3> r, range<3> size,
-                    const KernelType &k,
+                    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);
+                               nd_range<3>(r, size), std::forward<KernelType>(k));
 #else
   submit(
       q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },

sycl/include/sycl/queue.hpp

@@ -157,10 +157,10 @@ class __SYCL_EXPORT SubmissionInfo {
 } // namespace v1
 
 template <typename KernelName = detail::auto_name, bool EventNeeded = false,
-          typename PropertiesT, typename KernelType, int Dims>
+          typename PropertiesT, typename KernelTypeUniversalRef, int Dims>
 auto submit_kernel_direct(
     const queue &Queue, PropertiesT Props, const nd_range<Dims> &Range,
-    const KernelType &KernelFunc,
+    KernelTypeUniversalRef &&KernelFunc,
     const detail::code_location &CodeLoc = detail::code_location::current()) {
   // TODO Properties not supported yet
   (void)Props;
@@ -170,6 +170,9 @@ auto submit_kernel_direct(
       "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 =
@@ -180,7 +183,11 @@ auto submit_kernel_direct(
       "must be either sycl::nd_item or be convertible from sycl::nd_item");
   using TransformedArgType = sycl::nd_item<Dims>;
 
-  HostKernelRef<KernelType, TransformedArgType, Dims> HostKernel(KernelFunc);
+  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.
@@ -193,9 +200,6 @@ auto submit_kernel_direct(
   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,