Repository navigation
[SYCL] Add nd_launch overloads taking sycl::kernel arguments as a std::span - #23480
Conversation
There was a problem hiding this comment.
🟡 Changes recommended
Dynamic graph updates can retain the wrong argument kind, and container forwarding records an incorrect source location.
2 open findings
What changed in this PR
Adds runtime-sized kernel argument lists for nd_launch while retaining direct submission.
Changes:
- Adds pointer-aware
raw_kernel_arg. - Adds C++20
std::spanoverloads and C++17 diagnostics. - Updates direct-submission ABI and coverage.
| File | Description |
|---|---|
sycl/test/extensions/raw_kernel_arg/pointer_arg_errors.cpp |
Tests pointer-form constraints. |
sycl/test/extensions/enqueue_functions/kernel_arg_span.cpp |
Tests span-compatible argument lists. |
sycl/test/extensions/enqueue_functions/kernel_arg_sequence_cxx17.cpp |
Tests pre-C++20 diagnostics. |
sycl/test/abi/sycl_symbols_windows.dump |
Updates Windows ABI symbols. |
sycl/test/abi/sycl_symbols_linux.dump |
Updates Linux ABI symbols. |
sycl/test/abi/layout_raw_kernel_arg.cpp |
Pins the expanded layout. |
sycl/test-e2e/RawKernelArg/pointer_arg.cpp |
Tests pointer binding end-to-end. |
sycl/test-e2e/Graph/Update/update_with_raw_kernel_arg_pointer.cpp |
Tests graph pointer updates. |
sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_span.cpp |
Tests span launches. |
sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct.cpp |
Updates test documentation. |
sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct_raw_ptr.cpp |
Makes raw pointers portable. |
sycl/source/queue.cpp |
Exports pointer/count entry points. |
sycl/source/feature_test.hpp.in |
Advances extension versions. |
sycl/source/detail/queue_impl.hpp |
Declares raw-argument submission. |
sycl/source/detail/queue_impl.cpp |
Converts raw arguments for submission. |
sycl/include/sycl/queue.hpp |
Revises direct-submission declarations. |
sycl/include/sycl/handler.hpp |
Binds pointer-tagged raw arguments. |
sycl/include/sycl/ext/oneapi/experimental/raw_kernel_arg.hpp |
Adds the pointer argument form. |
sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp |
Adds span overloads and forwarding. |
🧠 Review effort: Balanced
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
adamfidel
left a comment
There was a problem hiding this comment.
The graph test looks good to me.
|
@gmlueck friendly reminder bump as there are about 2 weeks till code freeze |
iclsrc
left a comment
There was a problem hiding this comment.
This PR adds std::span<const raw_kernel_arg> overloads of nd_launch for the queue and handler forms. It adds a pointer form of raw_kernel_arg (pointer_arg), which binds as kind_pointer. It switches the exported submit_kernel_obj_direct_without_event_impl to take a pointer and a count, and bumps the SYCL_EXT_ONEAPI_ENQUEUE_FUNCTIONS and SYCL_EXT_ONEAPI_RAW_KERNEL_ARG feature macros to 2.
I found no correctness problems in the implementation. The overload resolution works out on paper: the span overload is more specialized than the container forwarder, which is more specialized than the pack overload. Argument lifetime is handled on both the scheduler-bypass and copy paths, and the ABI dump and layout test are updated to match. I did not build or run anything, and I did not read the extension spec documents.
The findings are:
- The extension spec documents are not updated, although the feature macros are bumped to 2.
- Graph
dynamic_parameterignores the new pointer flag. sizeof(raw_kernel_arg)grows from 16 to 24 bytes.
The PR touches no SYCL compiler pass code, so the sycl-passes checklist did not apply.
Review comments (lines not in diff — posted as body):
sycl/source/feature_test.hpp.in:100: 🟡 Important: Both feature-test macros are bumped to 2, but this PR does not update sycl_ext_oneapi_enqueue_functions.asciidoc or sycl_ext_oneapi_raw_kernel_arg.asciidoc. Neither document describes the new raw_kernel_arg(T *const *, pointer_arg_t) constructor, pointer_arg, or the std::span<const raw_kernel_arg> nd_launch overloads. Nothing in the diff documents what version 2 means. Please add the spec changes and a revision-history entry in this PR, or in a linked one.
sycl/include/sycl/ext/oneapi/experimental/raw_kernel_arg.hpp:199: 🟡 Important: raw_kernel_arg now carries MIsPointer, but the graph path in dynamic_impl.hpp and dynamic_impl.cpp never reads it. Both the dynamic_parameter constructor and updateValue(const raw_kernel_arg *, size_t) copy only MArgData and MArgSize. The new graph test only passes because the node's original kind is kept and the bytes are memcpy'd over it.
This has two consequences:
- A pointer-form
raw_kernel_argcan be used to update a byte-form parameter, or the reverse, with no diagnostic. The two forms reach the backend as different argument types, so the update is silently a different kind. sizeof(raw_kernel_arg)grows from 16 to 24 bytes, so any code that stores or copies the type by value now sees a different layout. The layout test pins this.
Consider checking that updateValue receives the same kind as the node's existing argument, or document that the kind is fixed at node creation.
sycl/include/sycl/ext/oneapi/experimental/raw_kernel_arg.hpp:192: 🔵 Suggestion: The new constructor stores the address of the user's pointer variable (pointer_location), not the pointer value. The handler path copies the bytes immediately, but nd_launch(queue, ...) and the span overloads read *pointer_location later. That read is synchronous, but only because the arguments are copied before return. Please state in the constructor's documentation, and in the spec, that the pointer variable only has to stay alive until the launching call returns. As it stands, the lifetime rule is easy to get wrong for the span case, where the raw_kernel_arg objects are often stored in a container.
uditagarwal97
left a comment
There was a problem hiding this comment.
Looks good to me overall, just a few NITs
The spec is in #23479, the graph point is covered in the thread on |
|
@intel/llvm-gatekeepers please consider merging |


Implements #23479. Follow-up to #22970, earlier review in #22971.
The parameter pack overload needs one instantiation per argument count, so a runtime that builds the argument list at run time, like Triton, needs a dispatch table with a cap. The new overload takes a
std::span<const raw_kernel_arg>and keeps the direct submission path from #22970; pointer arguments use thepointer_argform. A sequence passed to the parameter pack overload goes to the span overload, and before C++20 it is rejected with astatic_assert. The entry point exported by #22970 now takes a pointer and a count, so no span type is part of the ABI.More context: intel/intel-xpu-backend-for-triton#7737