From f4ddb3aa03f6def2833ffafa14e0e884d55a64a2 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 14 Oct 2026 07:53:17 +0000 Subject: [PATCH 1/4] [SYCL][Doc] Specify std::span kernel argument overloads of nd_launch and the pointer form of raw_kernel_arg --- ...sycl_ext_oneapi_enqueue_functions.asciidoc | 87 ++++++++++++++++++- .../sycl_ext_oneapi_raw_kernel_arg.asciidoc | 61 +++++++++++-- 2 files changed, 137 insertions(+), 11 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc index a01b9504a5f28..8a846c7e32a1c 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc +++ b/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 r, } ---- !==== +_Constraints_: `+sizeof...(Args)+` is not one, or the single argument in +`args` is not convertible to `std::span`. + _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 // Requires C++20 +void nd_launch(sycl::queue q, sycl::nd_range r, + const sycl::kernel& k, std::span args); + +template // Requires C++20 +void nd_launch(sycl::handler &h, sycl::nd_range r, + const sycl::kernel& k, std::span 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]. +_end note_] + +a| +[frame=all,grid=none] +!==== +a! +[source,c++] +---- +namespace sycl::ext::oneapi::experimental { + template void nd_launch(sycl::queue q, launch_config, 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`, a `std::array` of them and a `std::span` of them +all convert to `std::span`. + +[source,c++] +---- +#include +#include +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( + q.get_context()); + sycl::kernel k = bundle.ext_oneapi_get_kernel(); + + float *data = sycl::malloc_device(N, q); + int factor = 2; + + std::vector 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? diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc index 6226aa1ec8be9..1f5a1540c5bec 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc +++ b/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(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 + 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 +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 From c6b09f2acf0900e9bb369ea1d46033a7ec5c1b8c Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Thu, 15 Oct 2026 10:26:10 +0000 Subject: [PATCH 2/4] Use the {endnote} attribute in the pointer_arg note --- .../experimental/sycl_ext_oneapi_enqueue_functions.asciidoc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc index 8a846c7e32a1c..749630d19e927 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc @@ -557,7 +557,7 @@ at index `i`, as if it had been passed to `set_arg`. The sequence referenced by [_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]. -_end note_] +_{endnote}_] a| [frame=all,grid=none] From f9f737488bf722ed08ce1804fdd8841c7a9de4c1 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Tue, 20 Oct 2026 08:40:31 +0000 Subject: [PATCH 3/4] Define the example kernel as a named function object --- ...sycl_ext_oneapi_enqueue_functions.asciidoc | 29 ++++++++++++------- 1 file changed, 18 insertions(+), 11 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc index 749630d19e927..9cc789ce48954 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc @@ -815,9 +815,10 @@ _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 +The following example launches `scale` the usual way, which is also what makes +it a kernel that `get_kernel_id` can find. It then launches it through its +`sycl::kernel` object with the argument list assembled at run time, as a +launcher for kernels that are not known at compile time would do. A `std::vector`, a `std::array` of them and a `std::span` of them all convert to `std::span`. @@ -829,25 +830,31 @@ 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 */ } +struct scale { + float *data; + int factor; + void operator()(sycl::nd_item<1>) const { /* kernel body */ } +}; int main() { sycl::queue q; - - auto bundle = - syclex::get_kernel_bundle( - q.get_context()); - sycl::kernel k = bundle.ext_oneapi_get_kernel(); + sycl::nd_range<1> range{{N}, {256}}; float *data = sycl::malloc_device(N, q); int factor = 2; + syclex::nd_launch(q, range, scale{data, factor}); + q.wait(); + + auto bundle = sycl::get_kernel_bundle( + q.get_context(), {sycl::get_kernel_id()}); + sycl::kernel k = bundle.get_kernel(); + std::vector 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); + syclex::nd_launch(q, range, k, args); q.wait(); sycl::free(data, q); From 9107be2aab0c3fcb76a4aae8a725969b0e8b78b7 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Tue, 20 Oct 2026 12:12:57 +0000 Subject: [PATCH 4/4] Revert "Define the example kernel as a named function object" This reverts commit f9f737488bf7. A kernel defined as a named function object cannot be launched through a kernel object, so the example needs the free function kernel. --- ...sycl_ext_oneapi_enqueue_functions.asciidoc | 29 +++++++------------ 1 file changed, 11 insertions(+), 18 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc index 9cc789ce48954..749630d19e927 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc @@ -815,10 +815,9 @@ _Effects_: Submits an executable command graph to the `sycl::queue` or `sycl::ha == Example -The following example launches `scale` the usual way, which is also what makes -it a kernel that `get_kernel_id` can find. It then launches it through its -`sycl::kernel` object with the argument list assembled at run time, as a -launcher for kernels that are not known at compile time would do. A +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`, a `std::array` of them and a `std::span` of them all convert to `std::span`. @@ -830,31 +829,25 @@ namespace syclex = sycl::ext::oneapi::experimental; static constexpr size_t N = 1024; -struct scale { - float *data; - int factor; - void operator()(sycl::nd_item<1>) const { /* kernel body */ } -}; +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclex::nd_range_kernel<1>)) +void scale(float *data, int factor) { /* kernel body */ } int main() { sycl::queue q; - sycl::nd_range<1> range{{N}, {256}}; + + auto bundle = + syclex::get_kernel_bundle( + q.get_context()); + sycl::kernel k = bundle.ext_oneapi_get_kernel(); float *data = sycl::malloc_device(N, q); int factor = 2; - syclex::nd_launch(q, range, scale{data, factor}); - q.wait(); - - auto bundle = sycl::get_kernel_bundle( - q.get_context(), {sycl::get_kernel_id()}); - sycl::kernel k = bundle.get_kernel(); - std::vector args; args.emplace_back(&data, syclex::pointer_arg); args.emplace_back(&factor, sizeof(factor)); - syclex::nd_launch(q, range, k, args); + syclex::nd_launch(q, sycl::nd_range<1>{{N}, {256}}, k, args); q.wait(); sycl::free(data, q);