[SYCL][2/2] Add nd_launch overloads taking kernel arguments as a span - #22971
Draft
mieshkiwrk wants to merge 5 commits into
Draft
[SYCL][2/2] Add nd_launch overloads taking kernel arguments as a span#22971mieshkiwrk wants to merge 5 commits into
nd_launch overloads taking kernel arguments as a span#22971mieshkiwrk wants to merge 5 commits into
Conversation
nd_launch(queue, nd_range, const kernel &, args...) expands to submit() plus a handler, so a kernel object never reaches the direct submission path that a kernel function object already takes. For a language runtime that loads kernels from a binary, the handler-less enqueue functions therefore save nothing. Bind the arguments straight from the call and reuse queue_impl::submit_kernel_scheduler_bypass with a real kernel_impl *, which that function already accepts, whenever every argument can be bound as plain bytes. Accessors, local accessors, streams and work group memory keep the command group path, the same line HasSpecialCaptures draws in the runtime, and a dependency the scheduler has to track still falls back to a command group. KernelArgView is passed across the ABI boundary, so it gets its own header in an inline versioned namespace with a layout test, following nd_range_view.
Contributor
|
@mieshkiwrk thanks for implementing this! Adding @uditagarwal97 as we talked about having handler-less sycl::kernel submission path. |
slawekptak
reviewed
Aug 19, 2026
Contributor
There was a problem hiding this comment.
I think it would be good to specify if this is sycl::span or std::span. The same applies to some other places.
Contributor
Author
There was a problem hiding this comment.
It is sycl::span, qualified now here and in the other places.
sycl::span is what SYCL uses as its own vocabulary type today - queue.hpp takes sycl::span<const event> all over, and the other extensions spell it that way too - so these overloads follow suit.
If SYCL ever moves to std::span, these overloads get aligned along with the rest of the API.
`is_plain_kernel_arg_v` classified through `std::decay_t`, which turns an array into a pointer, so `nd_launch(queue, range, kernel, array, ...)` bound the array as UR_EXP_KERNEL_ARG_TYPE_POINTER and the runtime read its first bytes as an address. That binds neither the bytes nor the array: wrong results on Level Zero and an abort inside the OpenCL driver, where the handler path binds the array as plain bytes. Classify without decaying, so an array keeps using the command group path, and cover both paths with a test. The E2E test also passed a USM pointer through `raw_kernel_arg`, which binds as a value argument and therefore only reaches the kernel on Level Zero. The pointer is now passed typed, and the all-raw case moved to a Level Zero gated test, the same restriction the RawKernelArg tests carry. Pass the kernel bundle to the direct submission the way the handler path passes it, so the device globals a bundle keeps to itself can be initialized.
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.
mieshkiwrk
force-pushed
the
kernel-fastpath-span
branch
from
August 19, 2026 15:39
6423856 to
62fa8b3
Compare
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
More context: intel/intel-xpu-backend-for-triton#7737