[llvm-branch-commits] [llvm] [libsycl] add single_task (PR #188797)

[Kseniya Tikhomirova via llvm-branch-commits]

================
@@ -138,12 +161,87 @@ class _LIBSYCL_EXPORT queue {
   template <typename Param>
   typename Param::return_type get_backend_info() const;
 
+  /// Defines and invokes a SYCL kernel function as a lambda expression or a
+  /// named function object type.
+  ///
+  /// \param kernelFunc is the kernel functor or lambda.
+  /// \return an event that represents the status of the submitted kernel.
+  template <typename KernelName, typename KernelType>
+  event single_task(const KernelType &kernelFunc) {
+    return single_task<KernelName, KernelType>({}, kernelFunc);
+  }
+
+  /// Defines and invokes a SYCL kernel function as a lambda expression or a
+  /// named function object type.
+  ///
+  /// \param depEvent is an event that specifies the kernel dependency.
+  /// \param kernelFunc is the kernel functor or lambda.
+  /// \return an event that represents the status of the submitted kernel.
+  template <typename KernelName, typename KernelType>
+  event single_task(event depEvent, const KernelType &kernelFunc) {
+    return single_task<KernelName, KernelType>({depEvent}, kernelFunc);
+  }
+
+  /// Defines and invokes a SYCL kernel function as a lambda expression or a
+  /// named function object type.
+  ///
+  /// \param depEvents is a collection of events that specify the kernel
+  /// dependencies.
+  /// \param kernelFunc is the kernel functor or lambda.
+  /// \return an event that represents the status of the submitted kernel.
+  template <typename KernelName, typename KernelType>
+  event single_task(const std::vector<event> &depEvents,
+                    const KernelType &kernelFunc) {
+    static_assert(
+        (detail::CheckFunctionSignature<std::remove_reference_t<KernelType>,
+                                        void()>::value),
+        "sycl::queue::single_task() requires a kernel instead of a command "
+        "group");
+
+    setKernelParameters(depEvents);
+    submitSingleTask<KernelName, KernelType>(kernelFunc);
+    return getLastEvent();
+  }
+
   /// Blocks the calling thread until all commands previously submitted to this
   /// queue have completed. Synchronous errors are reported through SYCL
   /// exceptions.
   void wait();
 
 private:
+  // Name of this function is defined by compiler. It generates call to this
+  // function in the host implementation of KernelFunc in submitSingleTask.
+  template <typename, typename... Args>
+  void sycl_kernel_launch(const char *KernelName, Args &&...args) {
+    static_assert(
+        (sizeof...(args) == 1) &&
+        "sycl_kernel_launch expects only 2 arguments now: name of kernel and "
+        "callable object passed to kernel invocation by the user.");
+    detail::ArgCollection TypelessArgs;
+    (TypelessArgs.addArg(args), ...);
+
+    submitKernelImpl(KernelName, TypelessArgs);
+  }
----------------
KseniyaTikhomirova wrote:

@tahonermann 
I marked that as open to discuss at our last meeting but there was no time to do that. I think it worths discussion all together.

in the current state if we need to pass >1 argument - they expect us to do a copy of arguments into a contiguous buffer and pass sizes of each argument separately. This part is not implemented since I expect liboffload team to reiterate discussion to introduce API to avoid this copy, Piotr said that they have tried to bring more efficient API to liboffload but their proposal was rejected. Although I don't think this copy can ever be acceptable from runtime side.

as an example what is expected, let's assume we have 3 arguments - lambda capture (Obj1), usm pointer for buffer (ptr2), and usm buffer size (size3). We must pass it to liboffload as the following:
ArgumentsData as a pointer to
struct Data {
  Obj1,
  ptr2,
  size3
};
ArgumentsSize as sizeof (Data) (sizeof(obj1) + sizeof(ptr2) + sizeof(size3)).

Plus (this PR is not merged) we need to pass 
 ol_kernel_launch_prop_t Props[2];
 Props[0].type = OL_KERNEL_LAUNCH_PROP_TYPE_SIZE;
 Props[0].data = &ArgumentsSize;
 Props[1] = OL_KERNEL_LAUNCH_PROP_END;

where ArgumentsSize is size_t[3] with {sizeof(obj1), sizeof(ptr2), sizeof(size3)}.

@pbalcer please correct my understanding if there is any issues.

https://github.com/llvm/llvm-project/pull/188797