[SYCL][Doc] Add std::span kernel arguments to nd_launch

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 overloads that take the kernel arguments as a `std::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.  Those overloads use `std::span`, so they are
+available only when the application is compiled with C++20 or later.
+
 
 == Status
 
@@ -165,8 +171,11 @@ 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 a `sycl::nd_range` and the
+ arguments of a `sycl::kernel` as a `std::span` of `raw_kernel_arg`.
 |===
 
 
@@ -512,6 +521,9 @@ void nd_launch(sycl::handler &h, sycl::nd_range<Dimensions> r,
 }
 ----
 !====
+_Constraints_: `+sizeof...(Args)+` is not one, or the single argument in
+`args` is not convertible to `std::span<const raw_kernel_arg>`.
+
 _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 arguments in `args` are passed to the kernel in the same
@@ -525,6 +537,36 @@ a!
 ----
 namespace sycl::ext::oneapi::experimental {
 
+template <int Dimensions> // Requires C++20
+void nd_launch(sycl::queue q, sycl::nd_range<Dimensions> r,
+               const sycl::kernel& k, std::span<const raw_kernel_arg> args);
+
+template <int Dimensions> // Requires C++20
+void nd_launch(sycl::handler &h, sycl::nd_range<Dimensions> r,
+               const sycl::kernel& k, std::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.
+
+[_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].
+_{endnote}_]
+
+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,
@@ -771,6 +813,47 @@ _Effects_: Submits an executable command graph to the `sycl::queue` or `sycl::ha
 
 |====
 
+== Example
+
+The following example launches a kernel whose argument list is assembled at run
+time, as a launcher for kernels that are not known at compile time would do, so
+the arguments are passed as a sequence rather than as a parameter pack.  A
+`std::vector<raw_kernel_arg>`, a `std::array` of them and a `std::span` of them
+all convert to `std::span<const raw_kernel_arg>`.
+
+[source,c++]
+----
+#include <sycl/sycl.hpp>
+#include <vector>
+namespace syclex = sycl::ext::oneapi::experimental;
+
+static constexpr size_t N = 1024;
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclex::nd_range_kernel<1>))
+void scale(float *data, int factor) { /* kernel body */ }
+
+int main() {
+  sycl::queue q;
+
+  auto bundle =
+      syclex::get_kernel_bundle<scale, sycl::bundle_state::executable>(
+          q.get_context());
+  sycl::kernel k = bundle.ext_oneapi_get_kernel<scale>();
+
+  float *data = sycl::malloc_device<float>(N, q);
+  int factor = 2;
+
+  std::vector<syclex::raw_kernel_arg> args;
+  args.emplace_back(&data, syclex::pointer_arg);
+  args.emplace_back(&factor, sizeof(factor));
+
+  syclex::nd_launch(q, sycl::nd_range<1>{{N}, {256}}, k, args);
+  q.wait();
+
+  sycl::free(data, q);
+}
+----
+
 == Issues
 
 . What should `submit_with_event` be called?

sycl/doc/extensions/experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc

@@ -55,10 +55,10 @@ 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.
+only when using the Level Zero or OpenCL backends.  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.
 
 
 == Overview
@@ -76,12 +76,27 @@ for which there is no equivalent type defined on the host).
 
 [source,c++]
 ----
+namespace syclex = sycl::ext::oneapi::experimental;
+
 int a;
 char* opaque_type;
 int nbytes;
 ...
 h.set_arg(0, a);
-h.set_arg(1, sycl::ext::oneapi::experimental::raw_kernel_arg(opaque_type, nbytes));
+h.set_arg(1, syclex::raw_kernel_arg(opaque_type, nbytes));
+h.parallel_for(range, kernel);
+----
+
+When a kernel argument is a pointer, a special overload must be used to indicate
+that the argument is a pointer.  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:
+
+[source,c++]
+----
+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);
 ----
 
@@ -104,8 +119,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 +134,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 +154,32 @@ 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 and is not a pointer.
 
 _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. 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 +196,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