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

[Piotr Balcer 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);
+  }
----------------
pbalcer wrote:

> @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?

That depends. There's already a host task API in liboffload: https://github.com/llvm/llvm-project/blob/c1fd7d7a1f573477d86a716bacda01ed1d8c7a43/offload/liboffload/API/Queue.td#L113
We have a L0 plugin implementation here: https://github.com/llvm/llvm-project/pull/189443

However, [SYCL host tasks](https://github.khronos.org/SYCL_Reference/iface/host-task-basic.html) are very permissive, allowing:
> it does not have the same restrictions as a SYCL kernel function, and can therefore contain any arbitrary C++ code.

Also, SYCL host tasks are explicitly called out as a mechanism for native interop:

> A host task can optionally be used to interoperate with the native backend objects associated with the [sycl::queue](https://github.khronos.org/SYCL_Reference/iface/queue.html#queue) executing the host task, the context that the [sycl::queue](https://github.khronos.org/SYCL_Reference/iface/queue.html#queue) is associated with, the [sycl::device](https://github.khronos.org/SYCL_Reference/iface/device.html#device) that the [sycl::queue](https://github.khronos.org/SYCL_Reference/iface/queue.html#queue) is associated with and the accessors that have been captured in the callable, via an optional sycl::interop_handle parameter.

Liboffload host tasks have additional implementation-defined restrictions. Most importantly, both CUDA and L0 (and I imagine HSA as well) do not allow calling their APIs from within a host function.

https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__EXECUTION.html
> The host function must not make any CUDA API calls.


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