[SYCL][2/2] Add enqueue function overloads taking kernel arguments as a `std::span` by mieshkiwrk · Pull Request #22971 · intel/llvm

mieshkiwrk · 2026-08-18T18:32:39Z

Follow-up to #22970.

The launch functions taking a const kernel & and args... (single_task, parallel_for and nd_launch, on a queue or a handler) need one instantiation per argument count, so a language runtime that only learns its argument list at run time has to carry a variadic dispatch table with a cap. Add overloads taking the arguments as a std::span<const raw_kernel_arg>: flat build cost, no cap. nd_launch with an nd_range on a queue keeps the same direct submission path as its parameter pack overload; the others bind through the handler, as their parameter pack overloads do. These overloads need C++20; SYCL_EXT_ONEAPI_ENQUEUE_FUNCTIONS becomes 2.

A raw_kernel_arg carried only bytes, so every element of such a sequence bound as a value argument, and a pointer bound that way is not portable: Level Zero happens to accept it, but OpenCL routes value arguments to clSetKernelArg, which rejects a USM pointer. The argument list these overloads exist for is a pointer plus scalars, so the shape they are meant to serve is exactly the one that is not portable. Hence the pointer form of raw_kernel_arg, which takes the address of the pointer the way the byte form takes the address of the bytes; SYCL_EXT_ONEAPI_RAW_KERNEL_ARG becomes 2.

A container of raw_kernel_arg converts to std::span, but a parameter pack is an exact match and wins overload resolution, so the parameter pack overloads now pass such an argument on to the span overload rather than binding the container object as a single argument, which is what the handler overloads used to do silently for a std::array. Before C++20 such a call is rejected with a static_assert.

More context: intel/intel-xpu-backend-for-triton#7737

slawekptak · 2026-08-19T08:13:52Z

@mieshkiwrk thanks for implementing this! Adding @uditagarwal97 as we talked about having handler-less sycl::kernel submission path.

gmlueck · 2026-09-01T15:52:23Z

What do other people on @intel/dpcpp-specification-reviewers think about this? It seems a little weird to exclude certain types like this from the template overload above taking an Args pack. On the other hand, I do see that it makes this API taking span more convenient.

On the other other hand, the span overload seems like a ninja feature, so we could just ask those users to use a cast.

If we did want to add this exclusion, it should be expressed in the spec as a Constraint on the overload taking the Args pack:

Constraints: sizeof(Args...) is not one, or the single Arg is not convertible to std::span<const raw_kernel_arg>.

In case it helps the decision, a few measurements from trying the options locally:

With neither rule, a container is bound as one argument: numArgs = 1, one value argument holding the whole array. It compiles, then faults on Level Zero and errors in the OpenCL adapter at launch.

Constraining the pack overload alone sends a container to the kernel function object overload, not to the span one - it resolves as intended once that overload enforces the constraint it already has, its pack being sycl::reduction objects.

Requiring a cast need not leave that trap open - the container spelling can be diagnosed instead of bound, as it already is in C++17 - but the call site then spells std::span<const raw_kernel_arg>{args}, since std::span{args} deduces a non-const span.

Unaffected in all three: a single raw_kernel_arg, typed arguments, kernel function objects with and without reductions.

slawekptak · 2026-09-02T09:34:36Z

Is the new line needed?

No, it was accidental - rewrapped.

Checking how it renders also showed the // Requires C++20 lines were dropped, since a line starting with // inside an AsciiDoc table cell is a comment - they are trailing comments now, so they show up.

adamfidel

Not commenting on the rest, but the SYCL Graph test LGTM.

The number of arguments is part of the type of a parameter pack, so a caller that only learns its argument list at run time has to instantiate the pack overload once for every count it may encounter. Add queue and handler overloads taking span<const raw_kernel_arg>, where the count is data instead. raw_kernel_arg is already type erased, so the two forms bind the same bytes and measure the same; only the caller's build differs. A container, or a non-const span, converts to span<const raw_kernel_arg>, but a parameter pack is an exact match and wins overload resolution, which would bind the container object itself as a single kernel argument. Diagnose that spelling rather than let it produce wrong results at run time. Bumps SYCL_EXT_ONEAPI_ENQUEUE_FUNCTIONS to 2 and documents the overloads.

A raw_kernel_arg carried only bytes, so the span overloads bound every element as a value argument. A pointer bound that way only reaches the kernel on Level Zero: the OpenCL adapter passes a value argument to clSetKernelArg, which rejects a USM pointer with CL_INVALID_MEM_OBJECT, and the Native CPU adapter puts the address of its own copy of the bytes into the argument slot. The argument list these overloads exist for is a pointer plus scalars, so the shape they are meant to serve was the one that did not work anywhere else. Add a pointer form of raw_kernel_arg. It takes the address of the pointer, the way the byte form takes the address of the bytes, so that passing the pointer itself does not compile. handler::setArgHelper and both makeKernelArgView overloads bind such an argument as a pointer. The graph extension needs no change: a dynamic parameter stores the raw_kernel_arg object itself and an update rewrites the bytes the node holds, which for a pointer argument are the pointer. SYCL_EXT_ONEAPI_RAW_KERNEL_ARG becomes 2. The E2E test of the span overloads no longer needs a Level Zero requirement, and neither does the one that passes every argument through raw_kernel_arg. New tests cover the handler path on OpenCL, a graph dynamic parameter that rebinds a pointer, the layout of raw_kernel_arg, and that passing a pointer by value does not compile.

The overloads taking the kernel arguments as a sequence take a sycl::span, which a std::vector, a std::array or a std::span of raw_kernel_arg all convert to, but a parameter pack is an exact match and wins overload resolution, so a container reached the parameter pack overload instead. On the queue that was diagnosed. On the handler it was not: a container that is trivially copyable, std::array among them, compiled and bound the container object as a single kernel argument, and one that is not produced an error from inside handler.hpp. Have both parameter pack overloads pass such an argument on to the overload taking a span, so that passing the container itself binds the arguments it holds. The specification loses the clause that made such a call ill-formed, along with a note claiming that sycl::span has no implicit conversion from a container, which it does have. Spell out sycl::span rather than span in the specification, since std::span is a different type, and replace the test of the diagnostic with one that checks that every spelling of an argument list compiles.

The note carried two things a specification does not need in that place: why the overloads taking a sequence of arguments exist, and how a container of raw_kernel_arg reaches them. The first is motivation and moves to the Overview, the second is usage and moves to an Examples section at the end of the specification, which is where this extension shows how its APIs are called.

The overloads that take the arguments of a sycl::kernel as a sequence now take a std::span, which makes them available only with C++20. That is the direction the other extensions take: an implementation cannot turn sycl::span into an alias of std::span without breaking the ABI, so a new API should not add another place that would have to change if SYCL ever switches. The exported entry points take a pointer and a count instead of a view, so the spelling the public overloads use is not part of the ABI either and can be changed again without breaking it. Before C++20 there is no overload taking a sequence, so a container of raw_kernel_arg handed to a parameter pack overload would be bound as a single kernel argument, which compiles for any trivially copyable container and gives wrong results. It is diagnosed instead, with a test that pins both the diagnostic and that the same spellings compile with C++20.

Line wrapping only, as git-clang-format produces it for the previous commit. The two test comments also said that a container bound as a single kernel argument produces wrong results at run time. It does not get that far: the launch faults on Level Zero and the OpenCL adapter rejects it, so the comments now say that it only fails once the kernel is launched.

The paragraph broke after "depend on" and after "version 2 of", leaving two lines of a few words each. It is filled to the line width now, with the link macro split inside its brackets, the way the other paragraphs in this file that end in a link do it.

A line starting with `//` inside an AsciiDoc table cell is a comment, so the two `// Requires C++20` lines were dropped: neither asciidoctor nor GitHub renders them, while the declarations they annotate are shown. The annotation is a trailing comment now, which is ordinary listing text and is rendered.

mieshkiwrk · 2026-09-21T07:13:22Z

Friendly reminder bump

gmlueck

I have some specific feedback below, but I also wanted to direct your attention to KhronosGroup/SYCL-Docs#922. That PR will add a more official "KHR" extension with much of the same functionality we have in "sycl_ext_oneapi_enqueue_functions". Eventually, we will redirect users to that KHR instead of this "oneapi" one.

In order to provide this new functionality in that KHR, we will probably want to add raw_kernel_arg also as a KHR extension and then augment KhronosGroup/SYCL-Docs#922 with new overloads like you are doing here.

There is no harm is adding the PR you have here, and this is the fastest way to make these new APIs available. However, I wanted to let you know of the long term direction.

gmlueck · 2026-09-22T20:11:25Z

+args.emplace_back(&data, syclex::pointer_arg);
+args.emplace_back(&count, sizeof(count));
+
+syclex::nd_launch(q, sycl::nd_range<1>{count, 256}, k, args);


We generally prefer an example to be a whole program that the user can compile and run. It's common to make the kernel body entirely empty for examples that illustrate host APIs (like this). See, for example:

https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_profiling_tag.asciidoc#example

Right, rewritten as a complete program.

It uses a free function kernel so the sycl::kernel comes from a bundle, and it compiles and runs as it stands. I also dropped the "requires C++20", since the dependency is already stated above the API.

gmlueck · 2026-09-22T20:15:28Z

+If one of the parameter pack overloads above is called with a single argument
+whose type is convertible to `std::span<const raw_kernel_arg>`, that argument is
+the kernel argument list and the effects are those of the overload taking a
+`std::span`.


Rather than adding this paragraph, we should add a constraint to the previous overloads like:

Constraints: sizeof(Args...) is not one, or the single Arg is not convertible to std::span<const raw_kernel_arg>.

Is this sufficient to replace the intent of the paragraph?

Yes, replaced - a single raw_kernel_arg is not convertible to std::span<const raw_kernel_arg>, so only a container is excluded.

In the headers the pack overload simply passes such a call on to the span one, so the behaviour is the same, and in C++17 it still gives a clear error message.

gmlueck · 2026-09-22T20:21:46Z

+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_]
+


Just adding this comment here as a convenient place ... There are several launch functions that take a pack of Args, but you only added a raw_kernel_arg overload for one set. The others are:

single_task taking const sycl::kernel& k, Args&&... args.

parallel_for taking sycl::range<Dimensions> r, const sycl::kernel& k, Args&&... args

parallel_for taking launch_config<sycl::range<Dimensions>, Properties> c, const sycl::kernel& k, Args&& args...

nd_launch taking launch_config<sycl::nd_range<Dimensions>, Properties> c, const sycl::kernel& k, Args&& args...

Can you add raw_kernel_arg overloads for those too?

Would you prefer them in this PR or a separate one?
It is eight more overloads plus spec and tests, and they would go through the handler - the fast path from #22970 only covers nd_launch with an nd_range.

I'm OK either way. Maybe members of @intel/llvm-reviewers-runtime have a preference?

If you do make a separate PR, I assume you will prioritize it, so it comes soon? I'd like the APIs taking std::span to be consistent before we make our next release.

I'll start looking into exact needed changes and let you know when it's ready about it's size.
For me it would be easier to keep it within this PR but it would get kinda big for review.

Let me update about the changes size and then we'll see how to handle that

Added them here in 5317dd6: single_task, parallel_for with a range, and both launch_config forms, on the queue and on the handler, with spec entries and tests. They bind through the handler, since the handler-less path only covers nd_launch with an nd_range, and add no exported symbols; the extension stays at version 2 as it is not released yet.

It is one big commit, let me know if you prefer as a separate PR.

Friendly reminder bump

this PR is huge already, I would suggest creating a follow-up PR for extra overloads

I've seen you did review already, not sure to split it now

Let me know if still needed then I'll do that

…g.asciidoc Co-authored-by: Greg Lueck <gregory.m.lueck@intel.com>

…he forwarding

AsciiDoc replaces three dots with an ellipsis character, inside a monospace span as well, so the constraint rendered as `sizeof…(Args)`. Pass the text through, as the free function kernels specification already does for `Args...`.

…functions

KseniyaTikhomirova

haven't reviewed tests yet

KseniyaTikhomirova · 2026-09-30T15:14:46Z

-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 and OpenCL


where this change came from? UR/driver feature is implemented? Does OpenCL here covers both GPU and CPU?

It came with the pointer_arg form, OpenCL already supports USM pointer arguments, so nothing new was needed in UR or the driver.

Reverted to mention only for GPU devices (FYI - tests passes on OpenCL CPU too)

KseniyaTikhomirova · 2026-09-30T15:18:56Z

 _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.  The `bytes` must not be a pointer argument.


Not sure if I understand what this comment is intended to say. the first sentence relates to the first public contructor of raw_kernel_arg, what is the second one stands for?

It meant to say it must not be a pointer as these go through pointer_arg

Agree it was confusing changed:
copyable. The bytes must not be a pointer argument.
to
copyable and is not a pointer.

KseniyaTikhomirova · 2026-09-30T15:41:22Z

                  const kernel &KernelObj, ArgsT &&...Args) {
-  CGH.set_args<ArgsT...>(std::forward<ArgsT>(Args)...);
-  CGH.parallel_for(Range, KernelObj);
+#if !__cpp_lib_span


to avoid addition of #if/#endif every time - can this condition be a part of diagnoseArgListSequence impl?

Yes thanks, moved

KseniyaTikhomirova · 2026-09-30T15:42:18Z

-  submit(std::move(Q), [&](handler &CGH) {
-    parallel_for(CGH, Range, KernelObj, std::forward<ArgsT>(Args)...);
-  });
+  if constexpr (detail::is_arg_list_container_v<ArgsT...>) {


diagnoseArgListSequence?

It was already caught through handler overload but adding here make it clearer.

Added

KseniyaTikhomirova · 2026-09-30T15:43:50Z

-  if constexpr ((detail::is_direct_kernel_arg_v<ArgsT> && ...)) {
-    // The array is one element longer than the pack so that a zero-argument
-    // kernel stays well formed.
+  // A container of raw_kernel_arg is the argument list, not one argument, and


diagnoseArgListSequence?

KseniyaTikhomirova · 2026-09-30T15:45:03Z

+template <int Dimensions, typename Properties, typename... ArgsT>
+void nd_launch(queue Q, launch_config<nd_range<Dimensions>, Properties> Config,
+               const kernel &KernelObj, ArgsT &&...Args) {
+  if constexpr (detail::is_arg_list_container_v<ArgsT...>) {


diagnoseArgListSequence?

KseniyaTikhomirova · 2026-09-30T15:47:54Z

 private:
  const void *MArgData;
  size_t MArgSize;
+  // Appended last, so that the offsets of the members above, which the library


assume it is AI generated and redundant,
please read through comment in this PR and remove redundant ones or anything that doesn't add value or extra clarification to the code.

Went through and minimized/removed them (+68/-160)

KseniyaTikhomirova · 2026-09-30T15:51:42Z

    // so they can be bound where the caller keeps them. Otherwise the command
    // group outlives the call and they have to be copied into its storage.
    for (size_t I = 0; I < Args.size(); ++I) {
+      const sycl::detail::KernelArgView View = makeKernelArgView(Args[I]);


is there any better approaches instead of doing 2 extra copies per iteration? 1 in makeView, another 1 is here

Thanks good catch, changed to return reference

KseniyaTikhomirova · 2026-09-30T16:10:05Z

Hi @gmlueck and @mieshkiwrk, I am curious why are we adding new overloads with std::span (unsupported with C++17) instead of using sycl::span that is supported by intel/llvm implementation and that is:
"Derived from libcxx span. Original _LIBCPP macros replaced with _SYCL_SPAN to avoid collisions."

Is there anything missing in sycl::span impl? I understand that usage of STL classes is more common but:

usually nobody creates std::span to pass to some API explicitly - just passes container to some function and then type deduction does its job
sycl::span doesn't require C++20.

gmlueck · 2026-09-30T19:28:15Z

Hi @gmlueck and @mieshkiwrk, I am curious why are we adding new overloads with std::span (unsupported with C++17) instead of using sycl::span that is supported by intel/llvm implementation and that is: "Derived from libcxx span. Original _LIBCPP macros replaced with _SYCL_SPAN to avoid collisions."

See #22971 (comment)

mieshkiwrk · 2026-10-05T07:49:10Z

Friendly reminder bump

mieshkiwrk · 2026-10-12T07:13:11Z

Friendly reminder bump

mieshkiwrk · 2026-10-14T09:08:18Z

Split into three PRs, as suggested and marking this PR as draft:

[SYCL][Doc] Add std::span kernel arguments to nd_launch #23479: the spec, nd_launch with a std::span and the pointer_arg form of raw_kernel_arg (version 2 of both extensions).
[SYCL] Add nd_launch overloads taking sycl::kernel arguments as a std::span #23480: the runtime part Triton needs, implementing [SYCL][Doc] Add std::span kernel arguments to nd_launch #23479.
[SYCL] Add std::span kernel argument overloads to the other launch functions #23481: the same overloads for the other launch functions, with their spec as version 3, so a release with only [SYCL][Doc] Add std::span kernel arguments to nd_launch #23479 and [SYCL] Add nd_launch overloads taking sycl::kernel arguments as a std::span #23480 doesn't describe overloads it doesn't have.

Together they are this PR rebased onto current sycl, apart from version 3 in #23481, so the review changes from here are all in. Keeping this one as a draft for the review history and will close it once they land.

@gmlueck you wanted the span APIs consistent before the next release - with this split the other overloads may come after it (depends on review), is that OK?

gmlueck · 2026-10-19T23:00:27Z

@gmlueck you wanted the span APIs consistent before the next release - with this split the other overloads may come after it (depends on review), is that OK?

I'd prefer to get them all into the same release if possible.

Spec part of #22971, split out so it can be reviewed on its own: - `sycl_ext_oneapi_enqueue_functions` version 2: `nd_launch` overloads taking a `sycl::nd_range` and the arguments of a `sycl::kernel` as a `std::span<const raw_kernel_arg>` (C++20). - `sycl_ext_oneapi_raw_kernel_arg` version 2: a `pointer_arg` constructor, since the byte form cannot pass a pointer portably. The other launch functions get the same overloads in #23481.

mieshkiwrk requested review from a team as code owners August 18, 2026 18:32

mieshkiwrk requested a review from KseniyaTikhomirova August 18, 2026 18:32

mieshkiwrk marked this pull request as draft August 18, 2026 18:32

mieshkiwrk mentioned this pull request Aug 18, 2026

[SYCL] Enable handler-less launcher intel/intel-xpu-backend-for-triton#7737

Open

slawekptak reviewed Aug 19, 2026

View reviewed changes

Comment thread sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc Outdated

mieshkiwrk force-pushed the kernel-fastpath-span branch from 6423856 to 62fa8b3 Compare August 19, 2026 15:39

mieshkiwrk marked this pull request as ready for review August 31, 2026 14:55

mieshkiwrk requested a review from a team as a code owner August 31, 2026 14:55

mieshkiwrk requested a review from adamfidel August 31, 2026 14:55

mieshkiwrk mentioned this pull request Sep 1, 2026

[SYCL][1/2] Submit an already built sycl::kernel without a handler #22970

Merged

gmlueck reviewed Sep 1, 2026

View reviewed changes

slawekptak reviewed Sep 3, 2026

View reviewed changes

adamfidel approved these changes Sep 8, 2026

View reviewed changes

mieshkiwrk added 8 commits September 14, 2026 16:44

mieshkiwrk force-pushed the kernel-fastpath-span branch from 6511561 to 5502f83 Compare September 14, 2026 16:59

gmlueck reviewed Sep 22, 2026

View reviewed changes

mieshkiwrk and others added 4 commits September 23, 2026 11:04

Update sycl/doc/extensions/experimental/sycl_ext_oneapi_raw_kernel_ar…

9608d55

…g.asciidoc Co-authored-by: Greg Lueck <gregory.m.lueck@intel.com>

[SYCL] Rewrap the applied suggestion for the pointer constructor

c01fd9a

[SYCL] Use the namespace alias in both raw_kernel_arg examples

89fb407

[SYCL] Say that the byte constructor must not take a pointer argument

5a2f8b8

mieshkiwrk added 10 commits September 23, 2026 09:10

[SYCL] Drop the argument sequence paragraph from the Overview

34bf2f1

[SYCL] Shorten the version 2 row of the enqueue functions extension

70e8177

[SYCL] Drop the Level Zero justification from the pointer argument note

b51e275

[SYCL] Move the pointer argument requirement to the Overview

bd98858

[SYCL] Use the reviewer's wording for the pointer argument requirement

a257649

[SYCL] Limit the backend support status to the tested backends

4538a46

[SYCL] Rewrite the argument sequence example as a complete program

5c15db5

[SYCL] Constrain the parameter pack overloads instead of describing t…

551d09a

…he forwarding

[SYCL] Add std::span overloads to the remaining kernel object launch …

5317dd6

…functions

mieshkiwrk changed the title ~~[SYCL][2/2] Add nd_launch overloads taking kernel arguments as a span~~ [SYCL][2/2] Add enqueue function overloads taking kernel arguments as a std::span Sep 25, 2026

KseniyaTikhomirova reviewed Sep 30, 2026

View reviewed changes

mieshkiwrk added 2 commits October 1, 2026 07:59

Review changes

cff7c04

comment reduction

a47e753

This was referenced Oct 14, 2026

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

Merged

[SYCL] Add nd_launch overloads taking sycl::kernel arguments as a std::span #23480

Open

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

Open

mieshkiwrk marked this pull request as draft October 14, 2026 09:08

Conversation

mieshkiwrk commented Aug 18, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

slawekptak commented Aug 19, 2026

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

adamfidel left a comment

Choose a reason for hiding this comment

Uh oh!

mieshkiwrk commented Sep 21, 2026

Uh oh!

gmlueck left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

mieshkiwrk Sep 25, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

KseniyaTikhomirova left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

mieshkiwrk commented Aug 18, 2026 •

edited

Loading

mieshkiwrk Sep 25, 2026 •

edited

Loading