[SYCL] Add std::span kernel argument overloads to the other launch functions

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 launch function overloads that take the arguments of a
+ `sycl::kernel` as a `std::span` of `raw_kernel_arg`.
 |===
 
 
@@ -301,19 +310,49 @@ a!
 ----
 namespace sycl::ext::oneapi::experimental {
 
-template <typename Args...>
+template <typename... Args>
 void single_task(sycl::queue q, const sycl::kernel& k, Args&&... args);
 
-template <typename Args...>
+template <typename... Args>
 void single_task(sycl::handler &h, const sycl::kernel& k, Args&&... args);
 
 }
 ----
 !====
+_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 single task. The arguments in `args` are passed to the kernel in the same
 order.
 
+a|
+[frame=all,grid=none]
+!====
+a!
+[source,c++]
+----
+namespace sycl::ext::oneapi::experimental {
+
+void single_task(sycl::queue q, const sycl::kernel& k, // Requires C++20
+                 std::span<const raw_kernel_arg> args);
+
+void single_task(sycl::handler &h, const sycl::kernel& k, // Requires C++20
+                 std::span<const raw_kernel_arg> args);
+
+}
+----
+!====
+_Effects_: Enqueues a kernel object to the `sycl::queue` or `sycl::handler` as
+a single task. 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}_]
+
 |====
 
 
@@ -397,6 +436,9 @@ void parallel_for(sycl::handler &h, sycl::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::range`.
 The arguments in `args` are passed to the kernel in the same order.
@@ -409,24 +451,93 @@ a!
 ----
 namespace sycl::ext::oneapi::experimental {
 
+template <int Dimensions> // Requires C++20
+void parallel_for(sycl::queue q, sycl::range<Dimensions> r,
+                  const sycl::kernel& k,
+                  std::span<const raw_kernel_arg> args);
+
+template <int Dimensions> // Requires C++20
+void parallel_for(sycl::handler &h, sycl::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::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 parallel_for(sycl::queue q,
                   launch_config<sycl::range<Dimensions>, Properties> c,
-                  const sycl::kernel& k, Args&& args...);
+                  const sycl::kernel& k, Args&&... args);
 
 template <int Dimensions, typename Properties, typename... Args>
 void parallel_for(sycl::handler &h,
                   launch_config<sycl::range<Dimensions>, Properties> c,
-                  const sycl::kernel& k, Args&& args...);
+                  const sycl::kernel& k, Args&&... args);
 
 }
 ----
 !====
+_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 launch configuration specified by a
 `launch_config`. The arguments in `args` are passed to the kernel in the same
 order.
 
+a|
+[frame=all,grid=none]
+!====
+a!
+[source,c++]
+----
+namespace sycl::ext::oneapi::experimental {
+
+template <int Dimensions, typename Properties> // Requires C++20
+void parallel_for(sycl::queue q,
+                  launch_config<sycl::range<Dimensions>, Properties> c,
+                  const sycl::kernel& k,
+                  std::span<const raw_kernel_arg> args);
+
+template <int Dimensions, typename Properties> // Requires C++20
+void parallel_for(sycl::handler &h,
+                  launch_config<sycl::range<Dimensions>, Properties> c,
+                  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 launch configuration specified by a
+`launch_config`. 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}_]
+
 |====
 
 
@@ -512,6 +623,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,24 +639,89 @@ 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,
-               const sycl::kernel& k, Args&& args...);
+               const sycl::kernel& k, Args&&... args);
 
 template <int Dimensions, typename Properties, typename... Args>
 void nd_launch(sycl::handler &h,
                launch_config<sycl::nd_range<Dimensions>, Properties> c,
-               const sycl::kernel& k, Args&& args...);
+               const sycl::kernel& k, Args&&... args);
 
 }
 ----
 !====
+_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 launch configuration specified by a
 `launch_config`. The arguments in `args` are passed to the kernel in the same
 order.
 
+a|
+[frame=all,grid=none]
+!====
+a!
+[source,c++]
+----
+namespace sycl::ext::oneapi::experimental {
+
+template <int Dimensions, typename Properties> // Requires C++20
+void nd_launch(sycl::queue q,
+               launch_config<sycl::nd_range<Dimensions>, Properties> c,
+               const sycl::kernel& k, std::span<const raw_kernel_arg> args);
+
+template <int Dimensions, typename Properties> // Requires C++20
+void nd_launch(sycl::handler &h,
+               launch_config<sycl::nd_range<Dimensions>, Properties> c,
+               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 launch configuration specified by a
+`launch_config`. 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}_]
+
 |====
 
 
@@ -771,6 +950,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?