[SYCL][1/2] Submit an already built `sycl::kernel` without a handler by mieshkiwrk · Pull Request #22970 · intel/llvm

mieshkiwrk · 2026-08-18T18:31:41Z

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.

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, and makeKernelArgView binds an array as the bytes it is, which is what handler::setArgHelper does with it.

A USM pointer passed as raw_kernel_arg binds 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_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; it now binds as the bytes of the handle, the exception handler::setArgHelper makes for OpenCLMemT.

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

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.

Copilot

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.

mieshkiwrk · 2026-09-01T08:54:41Z

+            std::is_pointer_v<plain_arg_t<T>>
+                ? kernel_param_kind_t::kind_pointer
+                : kernel_param_kind_t::kind_std_layout};


Good catch, and it is worse than a behavior change: fixed in 73dbc05. handler::ShouldEnableSetArg whitelists OpenCLMemT, so nd_launch(queue, range, kernel, cl_mem) is a supported call today, and cl_mem being _cl_mem * did reach this path. The handle then went to clSetKernelArgMemPointerINTEL as if it were an address, which faults the GPU context on the OpenCL backend rather than misbehaving quietly.

The kind selection moved into plain_arg_kind_v, which makes the same exception handler::setArgHelper makes for OpenCLMemT, so a cl_mem binds as the bytes of the handle via clSetKernelArg. Verified on a Data Center GPU Max 1100 against an OpenCL C source kernel taking a cl_mem: correct on both the handler-less and the command group path with the fix, GPU page fault without it. #22971 carries the same change.

iclsrc

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.

iclsrc · 2026-08-31T15:19:52Z

+    auto CommandGroup = std::make_unique<detail::CGExecKernel>(
+        KData.getNDRDesc(), /*HostKernel*/ nullptr, KernelImpl,
+        KernelBundleImpl, std::move(CGData), std::move(KData).getArgs(),
+        *KData.getDeviceKernelInfoPtr(),
+        std::vector<std::shared_ptr<detail::stream_impl>>{},
+        std::vector<std::shared_ptr<const void>>{}, detail::CGType::Kernel,
+        KData.getKernelCacheConfig(), KData.isCooperative(),
+        KData.usesClusterLaunch(), KData.getKernelWorkGroupMemorySize(),
+        CodeLoc);


🟡 Important: missing "Extract data to move KData" pre-extraction step

The sibling function submit_kernel_direct_impl (lines 853-857) has an explicit comment and extracts KernelCacheConfig, IsCooperative, UsesClusterLaunch, and KernelWorkGroupMemorySize into local variables before the CGExecKernel constructor call that also contains std::move(KData).getArgs(). The reason is that C++ function-argument evaluation order is unspecified, so KData.getKernelCacheConfig() etc. and std::move(KData).getArgs() could be evaluated in any order. While the four accessors and getArgs() && touch different KData fields (so the code is safe today), it diverges from the defensive pattern established in the codebase and will silently break if a future refactor moves additional fields out of KData in getArgs(). Please apply the same extraction pattern here.

Suggested change

auto CommandGroup = std::make_unique<detail::CGExecKernel>(

KData.getNDRDesc(), /*HostKernel*/ nullptr, KernelImpl,

KernelBundleImpl, std::move(CGData), std::move(KData).getArgs(),

*KData.getDeviceKernelInfoPtr(),

std::vector<std::shared_ptr<detail::stream_impl>>{},

std::vector<std::shared_ptr<const void>>{}, detail::CGType::Kernel,

KData.getKernelCacheConfig(), KData.isCooperative(),

KData.usesClusterLaunch(), KData.getKernelWorkGroupMemorySize(),

CodeLoc);

// Extract data to move KData.

auto KernelCacheConfig = KData.getKernelCacheConfig();

bool IsCooperative = KData.isCooperative();

bool UsesClusterLaunch = KData.usesClusterLaunch();

auto KernelWorkGroupMemorySize = KData.getKernelWorkGroupMemorySize();

auto CommandGroup = std::make_unique<detail::CGExecKernel>(

KData.getNDRDesc(), /*HostKernel*/ nullptr, KernelImpl,

KernelBundleImpl, std::move(CGData), std::move(KData).getArgs(),

*KData.getDeviceKernelInfoPtr(),

std::vector<std::shared_ptr<detail::stream_impl>>{},

std::vector<std::shared_ptr<const void>>{}, detail::CGType::Kernel,

KernelCacheConfig, IsCooperative, UsesClusterLaunch,

KernelWorkGroupMemorySize, CodeLoc);

Applied in 73dbc05, with the sibling's explicit types rather than auto, so the two functions stay textually comparable. Agreed it was safe as written - getArgs() && only moves MArgs and the four accessors read disjoint members - but keeping the pattern identical to submit_kernel_direct_impl is worth more than the saved lines.

iclsrc · 2026-08-31T15:19:53Z

+    // so they can be bound where the caller keeps them. Otherwise the command
+    // group outlives the call and they have to be copied into its storage.
+    for (size_t I = 0; I < Args.size(); ++I) {
+      void *Value = const_cast<void *>(Args[I].MPtr);


🔵 Suggestion: document why const_cast is safe here

KernelArgView::MPtr is const void* (pointing to caller-owned data), but ArgDesc::MPtr is void*, making the cast necessary. The block-level comment above already explains the bypass invariant ("argument values are read before this call returns"), but a brief inline note on the cast itself would make the const-correctness reasoning immediately obvious to a reader who lands here without reading the preceding comment.

Longer-term, making ArgDesc::MPtr a const void* would eliminate the cast entirely.

Note added in 73dbc05. On making ArgDesc::MPtr a const void *: it is deliberately mutable for the kinds that carry an object rather than bytes - the scheduler casts it to a non-const Requirement * and to sampler *, and the stream path writes through it - so the cast would move to those call sites rather than disappear. Worth doing separately, if at all.

iclsrc · 2026-08-31T15:19:53Z

+// requirement for the scheduler to track. Accessors, local accessors, streams
+// and work group memory are deliberately excluded and keep using the command
+// group path; `HasSpecialCaptures` in the runtime draws the same line.
+template <typename T>
+inline constexpr bool is_plain_kernel_arg_v =
+    std::is_arithmetic_v<plain_arg_t<T>> || std::is_enum_v<plain_arg_t<T>> ||
+    std::is_pointer_v<plain_arg_t<T>> ||
+    std::is_same_v<plain_arg_t<T>, raw_kernel_arg>;


🔵 Suggestion: note that arrays are intentionally excluded from the fast path in [1/2]

The plain_arg_t comment correctly explains why std::decay_t is avoided (decay would turn int[4] into int* and classify it as kind_pointer). However, with the current is_plain_kernel_arg_v, an int[4] argument satisfies none of the four traits — is_arithmetic, is_enum, is_pointer, is_same<raw_kernel_arg> — so arrays still fall through to the command-group path. A brief note (e.g. "Arrays fall through to the command-group path for now; std::is_array_v will be added in the follow-up.") would prevent a reader from concluding that arrays silently misfire on the fast path.

The observation is right, the conclusion is not, so this became a fix rather than a note (73dbc05).

#22971 does not add std::is_array_v either, and the exclusion was not intentional: commit 2 is titled "Bind array arguments as bytes on the kernel object fast path", and nd_launch_kernel_obj_array_arg.cpp checks its first launch under "array without a handler" - which was in fact taking the handler, so the test covered the command group path twice.

makeKernelArgView already bound an array the way the handler does, {&Arg, sizeof(int[4]), kind_std_layout}; only the trait kept it away. An array of scalars is now is_plain_kernel_arg_v, while class types, and arrays of them, keep the command group path since those may be structs with special types inside. For the record the sizeof(T) -> sizeof(plain_arg_t<T>) part of commit 2 was a no-op: T is deduced from const T &, so it was already cv/ref-free. #22971 carries the same change.

`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.

gmlueck · 2026-09-01T14:22:30Z

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 .

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.

mieshkiwrk requested a review from a team as a code owner August 18, 2026 18:31

mieshkiwrk requested a review from cperkinsintel August 18, 2026 18:31

mieshkiwrk marked this pull request as draft August 18, 2026 18:31

mieshkiwrk mentioned this pull request Aug 18, 2026

[SYCL] Enable handler-less launcher intel/intel-xpu-backend-for-triton#7737

Open

mieshkiwrk mentioned this pull request Aug 20, 2026

[SYCL][2/2] Add nd_launch overloads taking kernel arguments as a span #22971

Open

mieshkiwrk marked this pull request as ready for review August 31, 2026 14:55

uditagarwal97 requested a balanced review from Copilot August 31, 2026 14:55

uditagarwal97 added the claude-review label Aug 31, 2026

Copilot started reviewing on behalf of uditagarwal97 August 31, 2026 14:56 View session

Copilot AI reviewed Aug 31, 2026

View reviewed changes

iclsrc reviewed Aug 31, 2026

View reviewed changes

iclsrc removed the claude-review label Aug 31, 2026

gmlueck reviewed Sep 1, 2026

View reviewed changes

Navigation Menu

Uh oh!

[SYCL][1/2] Submit an already built `sycl::kernel` without a handler - #22970

[SYCL][1/2] Submit an already built `sycl::kernel` without a handler#22970
mieshkiwrk wants to merge 3 commits into
intel:syclfrom
mieshkiwrk:kernel-fastpath-variadic

mieshkiwrk commented Aug 18, 2026 •

edited

Loading

Uh oh!

Copilot AI left a comment

Uh oh!

mieshkiwrk Sep 1, 2026

Uh oh!

iclsrc left a comment

Uh oh!

iclsrc Aug 31, 2026

Uh oh!

mieshkiwrk Sep 1, 2026

Uh oh!

iclsrc Aug 31, 2026

Uh oh!

mieshkiwrk Sep 1, 2026

Uh oh!

iclsrc Aug 31, 2026

Uh oh!

mieshkiwrk Sep 1, 2026

Uh oh!

gmlueck Sep 1, 2026

Uh oh!

uditagarwal97 Sep 1, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

Uh oh!

Conversation

mieshkiwrk commented Aug 18, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Reviewed changes

Uh oh!

Choose a reason for hiding this comment

Uh oh!

iclsrc left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

mieshkiwrk commented Aug 18, 2026 •

edited

Loading