Repository navigation
[SYCL][1/2] Submit an already built sycl::kernel without a handler - #22970
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.
`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.
There was a problem hiding this comment.
Pull request overview
Adds direct submission for prebuilt sycl::kernel objects while preserving scheduler fallback for special arguments and dependencies.
Changes:
- Introduces ABI-stable kernel argument views and a direct runtime submission path.
- Classifies and binds plain arguments without constructing a handler.
- Adds ABI and end-to-end coverage for typed, raw, accessor, scratch-memory, and array arguments.
Reviewed changes
Copilot reviewed 13 out of 13 changed files in this pull request and generated 1 comment.
Show a summary per file
| File | Description |
|---|---|
sycl/test/abi/sycl_symbols_windows.dump |
Records the new Windows ABI symbol. |
sycl/test/abi/sycl_symbols_linux.dump |
Records the new Linux ABI symbol. |
sycl/test/abi/layout_kernel_arg_view.cpp |
Verifies argument-view ABI layout. |
sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct.cpp |
Tests direct and fallback launches. |
sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct_raw_ptr.cpp |
Tests raw pointer arguments on Level Zero. |
sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_array_arg.cpp |
Tests array argument handling. |
sycl/source/queue.cpp |
Bridges the exported API to queue_impl. |
sycl/source/detail/queue_impl.hpp |
Declares kernel-object direct submission. |
sycl/source/detail/queue_impl.cpp |
Implements bypass and scheduler fallback. |
sycl/include/sycl/queue.hpp |
Exposes the internal submission entry point. |
sycl/include/sycl/ext/oneapi/experimental/raw_kernel_arg.hpp |
Adds raw-argument access helpers. |
sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp |
Routes eligible launches directly. |
sycl/include/sycl/detail/kernel_arg_view.hpp |
Defines the ABI argument representation. |
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
iclsrc
left a comment
iclsrc
left a comment
There was a problem hiding this comment.
Clean implementation of the handler-less nd_launch fast path. The two-commit structure is clear: commit 1 wires up the infrastructure (submit_kernel_obj_direct_without_event + the span-based arg view), commit 2 fixes the array-as-pointer misclassification by switching from std::decay_t to plain_arg_t. ABI versioning via kernel_arg_view_v1 inline namespace is the right approach and the layout test pins the struct correctly.
One pattern inconsistency flagged below; everything else looks solid.
`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`. `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.
I know this is also a problem in the existing code, but why aren't these functions in the detail namespace? In general, we shouldn't define identifiers directly in the sycl namespace unless they are part of the public API.
This might be a question that is more for @intel/llvm-reviewers-runtime .
There was a problem hiding this comment.
I agree, we should move these functions to the detail namespace. I've opened a JIRA, CMPLRLLVM-77982, to review and fix all such cases in the next ABI breaking window.
|
Friendly reminder bump |
uditagarwal97
left a comment
uditagarwal97
left a comment
There was a problem hiding this comment.
SYCL changes LGTM overall, just a few NITs
There was a problem hiding this comment.
Removed, along with the two lines left over after it - the trait now says in three lines what it accepts.
There was a problem hiding this comment.
I think the name "plain" kernel arguments is ambiguous - can we have a more descriptive name instead? Something that suggests that this trait filters out arguments of type Accessors, Streams, etc.
There was a problem hiding this comment.
Renamed to is_direct_kernel_arg_v, with what it accepts spelled out above it. Not named after special captures, since the trait is also false for a plain struct that has none.
Also renamed plain_arg_kind_v to kernel_arg_kind_v, so no name is left calling an argument plain - say if you would rather keep it.
Co-authored-by: Udit Kumar Agarwal <udit.agarwal@intel.com>
Applying the suggestion replaced only its last line, so the comment kept a sentence that ended mid-way and repeated the first one. It is a single line now, as the suggestion intended, while still saying that an array is not decayed, which is the reason this alias exists.
Rename `plain_arg_t` to `unqualified_arg_t` and `is_plain_kernel_arg_v` to `is_direct_kernel_arg_v`, which says what the trait selects for, and rename `plain_arg_kind_v` to `kernel_arg_kind_v` so that no name is left describing an argument as plain. Shorten the comments on the traits and on the pointer kind, say above the `if constexpr` what each of its branches is for, and drop the notes on binding the arguments in place and on seeding the code location, both of which the code says on its own.
There was a problem hiding this comment.
NIT: Can we simplify this type trait by calling unqualified_arg_t<T> once and then reusing the unqualified type?
Similar simplifications can be done in the following traits as well.
There was a problem hiding this comment.
Done - a defaulted template parameter in the traits, a local alias in makeKernelArgView.
The traits spelled `unqualified_arg_t<T>` again for every use of the unqualified type. They take it as a defaulted template parameter now, and makeKernelArgView keeps it in a local alias, so each of them names it once.
|
The three red CI checks are known Arc/Windows flakes, two of them green on the previous commit, which differs only by comments and an alias refactor. |
Yes, CI failures are unrelated. Merging the PR. |
|
Thanks for the review! |
… `std::span` (#23480) 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
nd_launch(queue, nd_range, const kernel &, args...)expands tosubmit()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_bypasswith a realkernel_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 lineHasSpecialCapturesdraws in the runtime, and a dependency the scheduler has to track still falls back to a command group.KernelArgViewis passed across the ABI boundary, so it gets its own header in an inline versioned namespace with a layout test, followingnd_range_view.Arrays were classified through
std::decay_t, so an array argument was bound as a pointer and the runtime read its first bytes as an address. Classification no longer decays, andmakeKernelArgViewbinds an array as the bytes it is, which is whathandler::setArgHelperdoes with it.A USM pointer passed as
raw_kernel_argbinds as a value argument, which only works on Level Zero - pre-existing, it reproduces on the released handler path. The test now passes the pointer typed, with the all-raw case gated on Level Zero. The bundle travels to the direct submission as it does through the handler.Review fixes.
cl_memis a pointer typedef, so it was bound as a pointer argument and the handle reachedclSetKernelArgMemPointerINTELas if it were an address, which faults on the OpenCL backend; it now binds as the bytes of the handle, the exceptionhandler::setArgHelpermakes forOpenCLMemT.An array of scalars now reaches the fast path, which no trait let it do, so the binding above is the one it gets. 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
kernel_arg_classification.cpp.More context: intel/intel-xpu-backend-for-triton#7737