Skip to content

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

Merged
uditagarwal97 merged 2 commits into
intel:syclfrom
mieshkiwrk:mdziado/kernel-args-span-nd-launch
Oct 21, 2026
Merged

uditagarwal97 merged 2 commits into
intel:syclfrom
mieshkiwrk:mdziado/kernel-args-span-nd-launch

Conversation

@mieshkiwrk

@mieshkiwrk mieshkiwrk commented Oct 14, 2026 •

Copy link
Copy Markdown
Contributor

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 the pointer_arg form. A sequence passed to the parameter pack overload goes to the span overload, and before C++20 it is rejected with a static_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

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🟡 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::span overloads 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.

Comment thread sycl/include/sycl/handler.hpp
Comment thread sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp Outdated

@adamfidel adamfidel left a comment

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The graph test looks good to me.

@mieshkiwrk

Copy link
Copy Markdown
Contributor Author

@gmlueck friendly reminder bump as there are about 2 weeks till code freeze

@gmlueck

gmlueck commented Oct 19, 2026

Copy link
Copy Markdown
Contributor

@gmlueck friendly reminder bump as there are about 2 weeks till code freeze

I usually review changes to specifications, so I'm probably not the right reviewer for this. However, I have been remis in my review of #23481. I just added my review comments to that.

@iclsrc iclsrc left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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_parameter ignores 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_arg can 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.

Comment thread sycl/source/feature_test.hpp.in
Comment thread sycl/test/abi/sycl_symbols_linux.dump
Comment thread sycl/test/abi/sycl_symbols_windows.dump

@uditagarwal97 uditagarwal97 left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Looks good to me overall, just a few NITs

@mieshkiwrk

Copy link
Copy Markdown
Contributor Author

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_parameter ignores 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_arg can 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.

The spec is in #23479, the graph point is covered in the thread on handler.hpp, and the new member goes after the existing two, so their offsets don't change.

@sys-ce-bb

Copy link
Copy Markdown
Contributor

@intel/llvm-gatekeepers please consider merging

@uditagarwal97
uditagarwal97 merged commit 527edb7 into intel:sycl Oct 21, 2026
30 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

7 participants