[SYCL][1/2] Submit an already built sycl::kernel without a handler - #22970
Draft
mieshkiwrk wants to merge 2 commits into
Draft
[SYCL][1/2] Submit an already built sycl::kernel without a handler#22970mieshkiwrk wants to merge 2 commits into
sycl::kernel without a handler#22970mieshkiwrk wants to merge 2 commits into
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.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
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 and an array falls back to the command group path.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.More context: intel/intel-xpu-backend-for-triton#7737