[SYCL][2/2] Add nd_launch overloads taking kernel arguments as a span

sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc

@@ -41,6 +41,12 @@ This extension is written against the SYCL 2020 revision 7 specification.  All
 references below to the "core SYCL specification" or to section numbers in the
 SYCL specification refer to that revision.
 
+The `nd_launch` overloads that take the kernel arguments as a `sycl::span`
+depend on
+version 2 of
+link:../experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc[sycl_ext_oneapi_raw_kernel_arg],
+which is what lets an element of that sequence represent a pointer argument.
+
 
 == Status
 
@@ -165,8 +171,13 @@ implementation supports.
 |Description
 
 |1
-|The APIs of this experimental extension are not versioned, so the
- feature-test macro always has this value.
+|Initial version of this extension.
+
+|2
+|Adds the `nd_launch` overloads that take the arguments of a `sycl::kernel` as
+ a `sycl::span` of `raw_kernel_arg`, and specifies that a parameter pack
+ overload handed such a sequence as its only argument passes it on as the
+ argument list.
 |===
 
 
@@ -525,6 +536,51 @@ a!
 ----
 namespace sycl::ext::oneapi::experimental {
 
+template <int Dimensions>
+void nd_launch(sycl::queue q, sycl::nd_range<Dimensions> r,
+               const sycl::kernel& k, sycl::span<const raw_kernel_arg> args);
+
+template <int Dimensions>
+void nd_launch(sycl::handler &h, sycl::nd_range<Dimensions> r,
+               const sycl::kernel& k, sycl::span<const raw_kernel_arg> args);
+
+}
+----
+!====
+_Effects_: Enqueues a kernel object to the `sycl::queue` or `sycl::handler`
+as a basic kernel, using the number of work-items specified by a
+`sycl::nd_range`. The element `args[i]` is passed to the kernel as the argument
+at index `i`, as if it had been passed to `set_arg`. The sequence referenced by
+`args` need only remain valid until the function returns.
+
+If one of the parameter pack overloads above is called with a single argument
+whose type is convertible to `sycl::span<const raw_kernel_arg>`, that argument
+is the kernel argument list and the effects are those of the overload taking a
+`sycl::span`.
+
+[_Note:_ These overloads exist for an application whose kernel argument list is
+only known at run time, which would otherwise have to instantiate a parameter
+pack overload for each number of arguments it may encounter. An application
+holding its arguments in a `std::vector<raw_kernel_arg>`, in a `std::array` of
+them or in a `std::span` of them can therefore pass the container itself, since
+all of those convert to a `sycl::span`. A single `raw_kernel_arg` still selects
+the parameter pack overload, as it is one argument rather than a sequence of
+them. _end note_]
+
+[_Note:_ An element that represents a pointer argument has to be constructed
+with the `pointer_arg` overload of `raw_kernel_arg`, as described in
+link:../experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc[sycl_ext_oneapi_raw_kernel_arg],
+since a `raw_kernel_arg` holding the bytes of a pointer is only guaranteed to
+bind that pointer on the Level Zero backend. _end note_]
+
+a|
+[frame=all,grid=none]
+!====
+a!
+[source,c++]
+----
+namespace sycl::ext::oneapi::experimental {
+
 template <int Dimensions, typename Properties, typename... Args>
 void nd_launch(sycl::queue q,
                launch_config<sycl::nd_range<Dimensions>, Properties> c,

sycl/doc/extensions/experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc

@@ -54,11 +54,17 @@ specification.*
 
 == Backend support status
 
-This extension is currently implemented in {dpcpp} only for GPU devices and
-only when using the Level Zero backend.  Attempting to use this extension in
-kernels that run on other devices or backends may result in undefined
-behavior.  Be aware that the compiler is not able to issue a diagnostic to
-warn you if this happens.
+This extension is currently implemented in {dpcpp} for the Level Zero, OpenCL,
+CUDA and Native CPU backends.  Attempting to use this extension on other
+backends may result in undefined behavior.  Be aware that the compiler is not
+able to issue a diagnostic to warn you if this happens.
+
+A kernel argument that is a pointer must be constructed with the `pointer_arg`
+overload described below.  Passing the byte representation of a pointer to the
+byte overload is only guaranteed to bind that pointer on the Level Zero backend,
+because a backend may take a pointer argument through a different entry point
+than the one that takes plain bytes, and it cannot tell the two apart from the
+bytes alone.
 
 
 == Overview
@@ -85,6 +91,19 @@ h.set_arg(1, sycl::ext::oneapi::experimental::raw_kernel_arg(opaque_type, nbytes
 h.parallel_for(range, kernel);
 ----
 
+An argument that is a pointer says so, since a backend may bind a pointer
+through a different entry point than a sequence of bytes:
+
+[source,c++]
+----
+namespace syclex = sycl::ext::oneapi::experimental;
+
+int* ptr = sycl::malloc_device<int>(n, q);
+...
+h.set_arg(0, syclex::raw_kernel_arg(&ptr, syclex::pointer_arg));
+h.parallel_for(range, kernel);
+----
+
 
 == Specification
 
@@ -104,8 +123,10 @@ implementation supports.
 |Description
 
 |1
-|The APIs of this experimental extension are not versioned, so the
- feature-test macro always has this value.
+|Initial version of this extension.
+
+|2
+|Adds the `raw_kernel_arg` constructor that takes a pointer argument.
 |===
 
 === The `raw_kernel_arg` class
@@ -117,9 +138,15 @@ kernel arguments via a raw byte representation.
 ----
 namespace sycl::ext::oneapi::experimental {
 
+  struct pointer_arg_t {};
+  inline constexpr pointer_arg_t pointer_arg{};
+
   class raw_kernel_arg {
   public:
     raw_kernel_arg(const void* bytes, size_t count);
+
+    template <typename T>
+    raw_kernel_arg(T* const* pointer_location, pointer_arg_t tag);
   };
 
 } // namespace sycl::ext::oneapi::experimental
@@ -131,13 +158,34 @@ raw_kernel_arg(const void* bytes, size_t count);
 ----
 _Preconditions_: `bytes` must point to an array of at least `count` bytes,
 which is the byte representation of a kernel argument that is trivially
-copyable.
+copyable.  If the argument is a pointer, only the Level Zero backend is
+guaranteed to bind it; see _Backend support status_ above.
 
 _Effects_: Constructs a `raw_kernel_arg` representing a view of the `count`
 bytes starting at the address specified by `bytes`. Since the `raw_kernel_arg`
 object is only a view, the caller must ensure that the lifetime of the `bytes`
 memory lasts at least as long as the lifetime of the `raw_kernel_arg` object.
 
+[source,c++]
+----
+template <typename T>
+raw_kernel_arg(T* const* pointer_location, pointer_arg_t tag);
+----
+_Preconditions_: `pointer_location` must point to a pointer that is a valid
+kernel argument, such as a pointer to memory allocated by one of the USM
+allocation functions.
+
+_Effects_: Constructs a `raw_kernel_arg` representing a view of the pointer
+stored at `pointer_location`, to be bound as a pointer argument rather than as
+the `sizeof(T*)` bytes it is made of. Since the `raw_kernel_arg` object is only
+a view, the caller must ensure that the lifetime of the pointer object at
+`pointer_location` lasts at least as long as the lifetime of the
+`raw_kernel_arg` object.
+
+[_Note:_ The constructor takes the address of the pointer, in keeping with the
+byte overload taking the address of the bytes, so that passing the pointer
+itself does not compile. _{endnote}_]
+
 === Using a raw kernel argument
 
 Instances of `raw_kernel_arg` are passed to kernels via the existing `set_arg`
@@ -154,7 +202,8 @@ argument in `args` was passed to `set_arg` ", adding a new overload of
 void set_arg(int argIndex, sycl::ext::oneapi::experimental::raw_kernel_arg&& arg);
 ----
 _Effects_: Sets the kernel argument associated with index `argIndex` using the
-bytes represented by `arg`.
+bytes represented by `arg`, or as a pointer argument if `arg` was constructed
+with the `pointer_arg` overload.
 
 
 == Issues

sycl/include/sycl/detail/kernel_arg_view.hpp

@@ -0,0 +1,35 @@
+//==---- kernel_arg_view.hpp --- SYCL kernel argument as bytes and kind ----==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/detail/kernel_desc.hpp> // for kernel_param_kind_t
+
+#include <stddef.h> // for size_t
+
+namespace sycl {
+inline namespace _V1 {
+namespace detail {
+
+inline namespace kernel_arg_view_v1 {
+
+// A kernel argument reduced to what the runtime needs in order to bind it. Used
+// by the enqueue functions that launch a `sycl::kernel` without building a
+// command group, where the arguments are only known as bytes plus a kind. This
+// is being passed across the ABI boundary, hence the versioned namespace.
+struct KernelArgView {
+  const void *MPtr;
+  size_t MSize;
+  kernel_param_kind_t MKind;
+};
+
+} // namespace kernel_arg_view_v1
+
+} // namespace detail
+} // namespace _V1
+} // namespace sycl

sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp

@@ -15,6 +15,7 @@
 #include <sycl/ext/oneapi/experimental/enqueue_types.hpp>
 #include <sycl/ext/oneapi/experimental/free_function_traits.hpp>
 #include <sycl/ext/oneapi/experimental/graph.hpp>
+#include <sycl/ext/oneapi/experimental/raw_kernel_arg.hpp>
 #include <sycl/ext/oneapi/properties.hpp>
 #include <sycl/handler.hpp>
 #include <sycl/nd_range.hpp>
@@ -98,6 +99,52 @@ template <typename LCRangeT, typename LCPropertiesT> struct LaunchConfigAccess {
   }
 };
 
+// The argument type as the kernel sees it. Deliberately not `std::decay_t`,
+// which turns an array into a pointer: an array has to keep being bound as the
+// bytes it is, which is what `handler::setArgHelper` does with it.
+template <typename T>
+using plain_arg_t = std::remove_cv_t<std::remove_reference_t<T>>;
+
+// An argument that can be bound as plain bytes, i.e. one that carries no
+// requirement for the scheduler to track. Accessors, local accessors, streams
+// and work group memory are deliberately excluded and keep using the command
+// group path; `HasSpecialCaptures` in the runtime draws the same line.
+template <typename T>
+inline constexpr bool is_plain_kernel_arg_v =
+    std::is_arithmetic_v<plain_arg_t<T>> || std::is_enum_v<plain_arg_t<T>> ||
+    std::is_pointer_v<plain_arg_t<T>> ||
+    std::is_same_v<plain_arg_t<T>, raw_kernel_arg>;
+
+// A pointer has to keep its kind. The runtime binds a pointer argument as
+// UR_EXP_KERNEL_ARG_TYPE_POINTER, which the OpenCL adapter passes to
+// clSetKernelArgMemPointerINTEL rather than to clSetKernelArg, so plain bytes
+// are not a substitute. The Native CPU adapter draws the same distinction: it
+// puts a pointer argument straight into the argument slot, whereas a value
+// argument lands there as the address of the adapter's own copy.
+template <typename T>
+sycl::detail::KernelArgView makeKernelArgView(const T &Arg) {
+  using sycl::detail::kernel_param_kind_t;
+  if constexpr (std::is_same_v<plain_arg_t<T>, raw_kernel_arg>)
+    return {RawKernelArgAccess::getData(Arg), RawKernelArgAccess::getSize(Arg),
+            RawKernelArgAccess::isPointer(Arg)
+                ? kernel_param_kind_t::kind_pointer
+                : kernel_param_kind_t::kind_std_layout};
+  else
+    return {&Arg, sizeof(plain_arg_t<T>),
+            std::is_pointer_v<plain_arg_t<T>>
+                ? kernel_param_kind_t::kind_pointer
+                : kernel_param_kind_t::kind_std_layout};
+}
+
+// True when a parameter pack overload was handed the argument list itself, as a
+// container that converts to the span the sibling overload takes. A pack is an
+// exact match and wins overload resolution, so such a call has to be forwarded
+// rather than bound as one argument.
+template <typename... ArgsT>
+inline constexpr bool is_arg_list_container_v =
+    sizeof...(ArgsT) == 1 &&
+    (std::is_convertible_v<ArgsT, sycl::span<const raw_kernel_arg>> && ...);
+
 template <typename CommandGroupFunc, typename PropertiesT>
 void submit_impl(const queue &Q, PropertiesT Props, CommandGroupFunc &&CGF,
                  const sycl::detail::code_location &CodeLoc) {
@@ -398,19 +445,82 @@ void nd_launch(queue Q, launch_config<nd_range<Dimensions>, Properties> Config,
   }
 }
 
+template <int Dimensions>
+void nd_launch(handler &CGH, nd_range<Dimensions> Range,
+               const kernel &KernelObj, sycl::span<const raw_kernel_arg> Args);
+
 template <int Dimensions, typename... ArgsT>
 void nd_launch(handler &CGH, nd_range<Dimensions> Range,
                const kernel &KernelObj, ArgsT &&...Args) {
-  CGH.set_args<ArgsT...>(std::forward<ArgsT>(Args)...);
-  CGH.parallel_for(Range, KernelObj);
+  if constexpr (detail::is_arg_list_container_v<ArgsT...>) {
+    nd_launch(CGH, Range, KernelObj,
+              sycl::span<const raw_kernel_arg>{std::forward<ArgsT>(Args)...});
+  } else {
+    CGH.set_args<ArgsT...>(std::forward<ArgsT>(Args)...);
+    CGH.parallel_for(Range, KernelObj);
+  }
 }
 
+template <int Dimensions>
+void nd_launch(queue Q, nd_range<Dimensions> Range, const kernel &KernelObj,
+               sycl::span<const raw_kernel_arg> Args,
+               const sycl::detail::code_location &CodeLoc =
+                   sycl::detail::code_location::current());
+
 template <int Dimensions, typename... ArgsT>
 void nd_launch(queue Q, nd_range<Dimensions> Range, const kernel &KernelObj,
                ArgsT &&...Args) {
-  submit(std::move(Q), [&](handler &CGH) {
-    nd_launch(CGH, Range, KernelObj, std::forward<ArgsT>(Args)...);
-  });
+  // A container of raw_kernel_arg is the argument list, not one argument, and
+  // the pack is what overload resolution picks for it, so hand it over to the
+  // overload that takes a span.
+  if constexpr (detail::is_arg_list_container_v<ArgsT...>) {
+    nd_launch(std::move(Q), Range, KernelObj,
+              sycl::span<const raw_kernel_arg>{std::forward<ArgsT>(Args)...});
+  } else if constexpr ((detail::is_plain_kernel_arg_v<ArgsT> && ...)) {
+    // Bind the arguments straight from this call, so that neither a handler nor
+    // a command group object has to be created. The array is one element longer
+    // than the pack so that a zero-argument kernel stays well formed.
+    const sycl::detail::KernelArgView ArgViews[sizeof...(ArgsT) + 1] = {
+        detail::makeKernelArgView(Args)...};
+    // An overload that ends in a parameter pack cannot take a trailing
+    // code_location parameter, so the location is the one this header sees,
+    // as it was when this overload went through submit(). Seed the TLS slot
+    // rather than leaving it default-constructed, which the instrumentation
+    // reads as a null file and function name.
+    sycl::detail::tls_code_loc_t TlsCodeLocCapture{
+        sycl::detail::code_location::current()};
+    sycl::submit_kernel_obj_direct_without_event_impl(
+        Q, sycl::detail::nd_range_view(Range), KernelObj,
+        {ArgViews, sizeof...(ArgsT)}, TlsCodeLocCapture.query(),
+        TlsCodeLocCapture.isToplevel());
+  } else {
+    submit(std::move(Q), [&](handler &CGH) {
+      nd_launch(CGH, Range, KernelObj, std::forward<ArgsT>(Args)...);
+    });
+  }
+}
+
+template <int Dimensions>
+void nd_launch(handler &CGH, nd_range<Dimensions> Range,
+               const kernel &KernelObj, sycl::span<const raw_kernel_arg> Args) {
+  // set_arg only takes an rvalue raw_kernel_arg; an lvalue would select the
+  // generic overload and bind the object itself as the argument.
+  for (size_t I = 0; I < Args.size(); ++I)
+    CGH.set_arg(static_cast<int>(I), raw_kernel_arg{Args[I]});
+  CGH.parallel_for(Range, KernelObj);
+}
+
+// Takes the kernel arguments as a contiguous sequence instead of a parameter
+// pack, for a caller that only learns its argument list at run time and would
+// otherwise need one instantiation of the pack overload per argument count.
+template <int Dimensions>
+void nd_launch(queue Q, nd_range<Dimensions> Range, const kernel &KernelObj,
+               sycl::span<const raw_kernel_arg> Args,
+               const sycl::detail::code_location &CodeLoc) {
+  sycl::detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
+  sycl::submit_kernel_obj_direct_without_event_impl(
+      Q, sycl::detail::nd_range_view(Range), KernelObj, Args,
+      TlsCodeLocCapture.query(), TlsCodeLocCapture.isToplevel());
 }
 
 template <int Dimensions, typename Properties, typename... ArgsT>