[SYCL][2/2] Add nd_launch overloads taking kernel arguments as a span - #22971
[SYCL][2/2] Add nd_launch overloads taking kernel arguments as a span#22971mieshkiwrk wants to merge 9 commits into
nd_launch overloads taking kernel arguments as a span#22971Conversation
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.
|
@mieshkiwrk thanks for implementing this! Adding @uditagarwal97 as we talked about having handler-less sycl::kernel submission path. |
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.
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.
There was a problem hiding this comment.
I'd like to use std::span here instead of sycl::span, and require applications to compile in C++20 mode in order to use these APIs. This is the direction we've been going for other extensions such as:
- sycl_ext_oneapi_inter_process_communication
- sycl_ext_oneapi_device_image_backend_content
- sycl_ext_oneapi_syclbin
We have a similar direction also in the KHR extensions. For example, we rely on C++23 std::dextents here:
https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:khr-group-interface
I believe it was a mistake for SYCL to "pre adopt" sycl::span because implementations will never be able to change this to be an alias to std::span without making an ABI break. TBH, I'm not sure how we are going to resolve this, but I'd like to stop relying on it in order to minimize the number of places that will be impacted if we do switch.
There was a problem hiding this comment.
Changed, both overloads take std::span<const raw_kernel_arg> behind __cpp_lib_span now, the spec marks them as C++20, and the exported entry points take a pointer and a count so sycl::span is out of the ABI as well. In C++17 a container would otherwise bind as a single kernel argument silently, so that is diagnosed.
One question: the feature-test macro still reports 2 in C++17, where version 2 of the extension does not exist - should it report 1 without __cpp_lib_span?
`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.
6423856 to
62fa8b3
Compare
`cl_mem` is a pointer typedef, so it was bound as a pointer argument and the handle reached clSetKernelArgMemPointerINTEL as if it were an address, which faults on the OpenCL backend. Bind it as the bytes of the handle, the exception `handler::setArgHelper` makes for `OpenCLMemT`. A `raw_kernel_arg` keeps saying at run time which of the two kinds it carries. `makeKernelArgView` already bound an array as the bytes it is, but no trait let one reach the fast path, so an array argument still went through the handler. Class types, and arrays of them, keep the command group path: those may be structs with special types inside. Neither which arguments take the fast path nor the kind they are bound with is observable end to end, hence the new test. Extract the `KData` fields before the `CGExecKernel` construction that moves it, the way `submit_kernel_direct_impl` does, and say why the argument pointer can be cast.
There was a problem hiding this comment.
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>.
There was a problem hiding this comment.
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
spanone - it resolves as intended once that overload enforces the constraint it already has, its pack beingsycl::reductionobjects. - 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}, sincestd::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.
There was a problem hiding this comment.
Personally, I think this note is not needed in the specification. If you want to show the usage (like the automatic conversion from containers like std::array), add an Examples section to the bottom of the spec and add an example.
There was a problem hiding this comment.
Removed the whole note. The usage it described moved to an Examples section at the end of the spec,
with the motivation kept as one line in the Overview.
There was a problem hiding this comment.
This seems like as good a place as any to add this comment ... Do we really need to call a special API on OpenCL to pass a kernel argument that is a pointer? What happens if you call the API that just sets raw bytes?
It seems weird that it wouldn't work because you can wrap the pointer in a struct, and then pass the struct by value. In that case, would wouldn't call the special "pointer" API to set the kernel argument. Why would this case be different from the case when the pointer is not wrapped in a struct?
CC: @bashbaug
There was a problem hiding this comment.
We ran into this passing a USM pointer through raw_kernel_arg, so I tested the spellings side by side with released headers, i.e. before this PR, on a Data Center GPU Max 1100:
backend typed pointer pointer as its bytes pointer in a struct, by value
Level Zero correct correct correct
OpenCL correct CL_INVALID_MEM_OBJECT correct
So the struct case does work, as you say, and on Level Zero even the byte form does. What fails is the plain pointer argument on OpenCL - a struct by value is a value argument, so nothing is bound as a pointer there, whereas setting a pointer parameter goes through a different entry point than setting bytes. Native CPU differs again, and there the byte form would not even report an error.
The fix in this PR is to let an element say it is a pointer, since a raw_kernel_arg otherwise carries only an address and a length, which a pointer and a same-sized scalar share. The spelling is open - arg_kind on the byte constructor would work just as well - if you would prefer it that way.
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.
3 participants
*Depends on #22970
nd_launch(queue|handler, nd_range, const kernel &, args...)needs 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 an overload taking the arguments as asycl::span<const raw_kernel_arg>: flat build cost, no cap, same direct submission path as the parameter pack overload.A
raw_kernel_argcarried only bytes, so every element of such a sequence bound as a value argument, and a pointer bound that way only reaches the kernel on Level Zero: OpenCL routes value arguments toclSetKernelArg, which rejects a USM pointer, and Native CPU stores the address of its own copy of the bytes. 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. Hence the pointer form ofraw_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_ARGbecomes 2.A container of
raw_kernel_argconverts tosycl::span, but a parameter pack is an exact match and wins overload resolution, so both 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 overload used to do silently for astd::array.The review fixes from #22970 are carried here too, so the two stay aligned:
cl_memkeeps the kind the handler gives it, and an array of scalars is bound as bytes on the fast path. Araw_kernel_argstays outside that classification, since it says at run time which of the two kinds it carries.More context: intel/intel-xpu-backend-for-triton#7737