[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:

"My expectation is that the raw argument buffer only needs to persist until olLaunchKernel() is called; " 
"The ideal ABI surface for (most of) kernel submission is the liboffload API."

The problems/features are:
- with host task as dependency olLaunchKernel call can be done outside of sycl_kernel_launch call. To support this we need to copy arguments that we should pass to backend to be able to use them later when all dependencies are ready. backends doesn't have native support of host tasks so dependency handling between kernels and host tasks is pretty complex and is done on SYCL rt side. At least it is how we handle them in intel/llvm. @pbalcer. @lplewa is there any chance we can have support of host tasks in liboffload itself so that SYCL RT can work with them as with other commands?
- once buffer & accessors are supported we will have to do more preparation work in RT. This is where our "Scheduler" become needed. It helps to build graph of memory usage and build proper dependencies for kernels using that data. I doubt this amount of work should be done purely in headers. So we need to have calls to SYCL RT for this. That will extend RT ABI with a lot of implementation details methods that will make RT ABI fragile. the only way to keep clear ABI with the current FE approach (at least, the only way that I see) is to collect all required data about args from FE and pass it to RT without type specifics. Then RT can decide if we need to add extra dependencies to kernel (buffer usage + extra stuff) and if we can submit kernel directly to liboffload (no blocking dependencies). If we can submit - we use wrapped arg data and don't do any copy of data itself. If we can't - we do a complete copy of arguments to use them later.



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