From cf619c44da3c95ddc7991114f4410fd89f622708 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 12 Aug 2026 12:48:53 +0000 Subject: [PATCH 1/7] [SYCL] Submit an already built sycl::kernel without a handler 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. --- sycl/include/sycl/detail/kernel_arg_view.hpp | 35 +++++ .../oneapi/experimental/enqueue_functions.hpp | 53 ++++++- .../oneapi/experimental/raw_kernel_arg.hpp | 11 ++ sycl/include/sycl/queue.hpp | 10 ++ sycl/source/detail/queue_impl.cpp | 77 +++++++++ sycl/source/detail/queue_impl.hpp | 12 ++ sycl/source/queue.cpp | 9 ++ .../nd_launch_kernel_obj_direct.cpp | 147 ++++++++++++++++++ sycl/test/abi/layout_kernel_arg_view.cpp | 17 ++ sycl/test/abi/sycl_symbols_linux.dump | 1 + sycl/test/abi/sycl_symbols_windows.dump | 1 + 11 files changed, 370 insertions(+), 3 deletions(-) create mode 100644 sycl/include/sycl/detail/kernel_arg_view.hpp create mode 100644 sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct.cpp create mode 100644 sycl/test/abi/layout_kernel_arg_view.cpp diff --git a/sycl/include/sycl/detail/kernel_arg_view.hpp b/sycl/include/sycl/detail/kernel_arg_view.hpp new file mode 100644 index 0000000000000..6062e6ed2ec1d --- /dev/null +++ b/sycl/include/sycl/detail/kernel_arg_view.hpp @@ -0,0 +1,35 @@ +//==---- kernel_arg_view.hpp --- SYCL kernel argument as bytes and kind ----==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include // for kernel_param_kind_t + +#include // for size_t + +namespace sycl { +inline namespace _V1 { +namespace detail { + +inline namespace kernel_arg_view_v1 { + +// A kernel argument reduced to what the runtime needs in order to bind it. Used +// by the enqueue functions that launch a `sycl::kernel` without building a +// command group, where the arguments are only known as bytes plus a kind. This +// is being passed across the ABI boundary, hence the versioned namespace. +struct KernelArgView { + const void *MPtr; + size_t MSize; + kernel_param_kind_t MKind; +}; + +} // namespace kernel_arg_view_v1 + +} // namespace detail +} // namespace _V1 +} // namespace sycl diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp index efc355875b6a5..f4012ee928e9e 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -15,6 +15,7 @@ #include #include #include +#include #include #include #include @@ -98,6 +99,33 @@ template struct LaunchConfigAccess { } }; +// An argument that can be bound as plain bytes, i.e. one that carries no +// 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 +inline constexpr bool is_plain_kernel_arg_v = + std::is_arithmetic_v> || std::is_enum_v> || + std::is_pointer_v> || + std::is_same_v, raw_kernel_arg>; + +// A pointer has to keep its kind. The runtime binds a pointer argument as +// UR_EXP_KERNEL_ARG_TYPE_POINTER, which the OpenCL adapter passes to +// clSetKernelArgMemPointerINTEL rather than to clSetKernelArg, so plain bytes +// are not a substitute. +template +sycl::detail::KernelArgView makeKernelArgView(const T &Arg) { + using sycl::detail::kernel_param_kind_t; + if constexpr (std::is_same_v, raw_kernel_arg>) + return {RawKernelArgAccess::getData(Arg), RawKernelArgAccess::getSize(Arg), + kernel_param_kind_t::kind_std_layout}; + else + return {&Arg, sizeof(T), + std::is_pointer_v> + ? kernel_param_kind_t::kind_pointer + : kernel_param_kind_t::kind_std_layout}; +} + template void submit_impl(const queue &Q, PropertiesT Props, CommandGroupFunc &&CGF, const sycl::detail::code_location &CodeLoc) { @@ -408,9 +436,28 @@ void nd_launch(handler &CGH, nd_range Range, template void nd_launch(queue Q, nd_range Range, const kernel &KernelObj, ArgsT &&...Args) { - submit(std::move(Q), [&](handler &CGH) { - nd_launch(CGH, Range, KernelObj, std::forward(Args)...); - }); + if constexpr ((detail::is_plain_kernel_arg_v && ...)) { + // Bind the arguments straight from this call, so that neither a handler nor + // a command group object has to be created. The array is one element longer + // than the pack so that a zero-argument kernel stays well formed. + const sycl::detail::KernelArgView ArgViews[sizeof...(ArgsT) + 1] = { + detail::makeKernelArgView(Args)...}; + // An overload that ends in a parameter pack cannot take a trailing + // code_location parameter, so the location is the one this header sees, + // as it was when this overload went through submit(). Seed the TLS slot + // rather than leaving it default-constructed, which the instrumentation + // reads as a null file and function name. + sycl::detail::tls_code_loc_t TlsCodeLocCapture{ + sycl::detail::code_location::current()}; + sycl::submit_kernel_obj_direct_without_event_impl( + Q, sycl::detail::nd_range_view(Range), KernelObj, + {ArgViews, sizeof...(ArgsT)}, TlsCodeLocCapture.query(), + TlsCodeLocCapture.isToplevel()); + } else { + submit(std::move(Q), [&](handler &CGH) { + nd_launch(CGH, Range, KernelObj, std::forward(Args)...); + }); + } } template diff --git a/sycl/include/sycl/ext/oneapi/experimental/raw_kernel_arg.hpp b/sycl/include/sycl/ext/oneapi/experimental/raw_kernel_arg.hpp index 4ca41f315a43f..04ce9fd47c1e8 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/raw_kernel_arg.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/raw_kernel_arg.hpp @@ -18,6 +18,7 @@ namespace ext::oneapi::experimental { namespace detail { class dynamic_parameter_impl; +struct RawKernelArgAccess; } // namespace detail class raw_kernel_arg { @@ -32,8 +33,18 @@ class raw_kernel_arg { friend class sycl::handler; // For sycl_ext_oneapi_graph integration friend class detail::dynamic_parameter_impl; + // For the enqueue paths that bind arguments without a handler + friend struct detail::RawKernelArgAccess; }; +namespace detail { +// Helper for accessing the members of raw_kernel_arg. +struct RawKernelArgAccess { + static const void *getData(const raw_kernel_arg &Arg) { return Arg.MArgData; } + static size_t getSize(const raw_kernel_arg &Arg) { return Arg.MArgSize; } +}; +} // namespace detail + } // namespace ext::oneapi::experimental } // namespace _V1 } // namespace sycl diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp index 629b73e8258f3..ef3a17dbe546b 100644 --- a/sycl/include/sycl/queue.hpp +++ b/sycl/include/sycl/queue.hpp @@ -18,6 +18,7 @@ #include // for code_location #include // for __SYCL2020_DEP... #include // for __SYCL_EXPORT +#include // for KernelArgView #include // for KernelInfo #include #include @@ -59,12 +60,21 @@ inline namespace _V1 { class context; class device; class event; +class kernel; class queue; template auto get_native(const SyclObjectT &Obj) -> backend_return_t; +// Launches an already built `sycl::kernel` with an explicit argument list, +// bypassing the handler and the command group object the same way +// submit_kernel_direct_* does for kernel function objects. +void __SYCL_EXPORT submit_kernel_obj_direct_without_event_impl( + const queue &Queue, const detail::nd_range_view &RangeView, + const kernel &Kernel, sycl::span Args, + const detail::code_location &CodeLoc, bool IsTopCodeLoc); + event __SYCL_EXPORT submit_kernel_direct_with_event_impl( const queue &Queue, const detail::nd_range_view &RangeView, detail::HostKernelRefBase &HostKernel, diff --git a/sycl/source/detail/queue_impl.cpp b/sycl/source/detail/queue_impl.cpp index a8c51a03b710e..f7f443f16d98e 100644 --- a/sycl/source/detail/queue_impl.cpp +++ b/sycl/source/detail/queue_impl.cpp @@ -883,6 +883,83 @@ EventImplPtr queue_impl::submit_kernel_direct_impl( /*InsertBarrierForInOrderCommand*/ false); } +void queue_impl::submit_kernel_obj_direct_without_event( + const detail::nd_range_view &RangeView, + const std::shared_ptr &KernelImpl, + sycl::span Args, + const detail::code_location &CodeLoc, bool IsTopCodeLoc) { + + KernelData KData; + KData.setDeviceKernelInfoPtr(&KernelImpl->getDeviceKernelInfo()); + KData.setNDRDesc(NDRDescT(RangeView)); + KData.getArgs().reserve(Args.size()); + + // This overload carries no properties, so a kernel that needs work group + // scratch memory can never have been given a size. The handler path reports + // that in handler.cpp, so report it here too rather than launching a kernel + // whose scratch allocation is missing. + if (KData.getDeviceKernelInfoPtr()->getWorkGroupDynamicLocalMem()) + throw sycl::exception( + sycl::make_error_code(sycl::errc::memory_allocation), + "Kernel allocates work group scratch memory but an allocation size " + "has not been specified through the work_group_scratch_size property!"); + + auto SubmitKernelFunc = [&](detail::CG::StorageInitHelper &&CGData) + -> std::pair { + bool SchedulerBypass = + (CGData.MEvents.size() > 0 + ? detail::Scheduler::areEventsSafeForSchedulerBypass( + CGData.MEvents, getContextImpl()) + : true) && + !hasCommandGraph(); + + // On the bypass path the argument values are read before this call returns, + // 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(Args[I].MPtr); + if (!SchedulerBypass) { + const char *Bytes = static_cast(Args[I].MPtr); + CGData.MArgsStorage.emplace_back(Bytes, Bytes + Args[I].MSize); + Value = CGData.MArgsStorage.back().data(); + } + KData.addArg(Args[I].MKind, Value, static_cast(Args[I].MSize), + static_cast(I)); + } + + if (SchedulerBypass) + return {submit_kernel_scheduler_bypass( + KData, CGData.MEvents, /*EventNeeded*/ false, + KernelImpl.get(), /*KernelBundleImpPtr*/ nullptr, CodeLoc, + IsTopCodeLoc), + /*SchedulerBypass*/ true}; + + auto CommandGroup = std::make_unique( + KData.getNDRDesc(), /*HostKernel*/ nullptr, KernelImpl, + /*KernelBundle*/ nullptr, std::move(CGData), std::move(KData).getArgs(), + *KData.getDeviceKernelInfoPtr(), + std::vector>{}, + std::vector>{}, detail::CGType::Kernel, + KData.getKernelCacheConfig(), KData.isCooperative(), + KData.usesClusterLaunch(), KData.getKernelWorkGroupMemorySize(), + CodeLoc); + CommandGroup->MIsTopCodeLoc = IsTopCodeLoc; + + if (auto GraphImpl = getCommandGraph(); GraphImpl) + return {submit_command_to_graph(*GraphImpl, std::move(CommandGroup), + detail::CGType::Kernel), + /*SchedulerBypass*/ false}; + + return {detail::Scheduler::getInstance().addCG(std::move(CommandGroup), + *this, true), + /*SchedulerBypass*/ false}; + }; + + submit_direct(/*CallerNeedsEvent*/ false, /*DepEvents*/ {}, SubmitKernelFunc, + detail::CGType::Kernel, + /*InsertBarrierForInOrderCommand*/ false); +} + EventImplPtr queue_impl::submit_graph_direct_impl( std::shared_ptr ExecGraph, diff --git a/sycl/source/detail/queue_impl.hpp b/sycl/source/detail/queue_impl.hpp index 2128ff4be7b0d..d1e6db6f9fb80 100644 --- a/sycl/source/detail/queue_impl.hpp +++ b/sycl/source/detail/queue_impl.hpp @@ -378,6 +378,18 @@ class queue_impl : public std::enable_shared_from_this { CodeLoc, IsTopCodeLoc); } + /// Submits an already built kernel with an explicit argument list, without + /// creating a handler or a command group object. + /// + /// \param RangeView is the execution range. + /// \param KernelImpl is the kernel to launch. + /// \param Args are the kernel arguments, as bytes plus their kind. + void submit_kernel_obj_direct_without_event( + const detail::nd_range_view &RangeView, + const std::shared_ptr &KernelImpl, + sycl::span Args, + const detail::code_location &CodeLoc, bool IsTopCodeLoc); + event submit_barrier_direct_with_event(sycl::span DepEvents, detail::CGType BarrierType, const detail::code_location &CodeLoc) { diff --git a/sycl/source/queue.cpp b/sycl/source/queue.cpp index 91e5aa5f4f82a..bdbf0b00b6fe1 100644 --- a/sycl/source/queue.cpp +++ b/sycl/source/queue.cpp @@ -13,6 +13,7 @@ #include #include #include +#include #include #include @@ -318,6 +319,14 @@ void submit_kernel_direct_without_event_impl( IsTopCodeLoc); } +void submit_kernel_obj_direct_without_event_impl( + const queue &Queue, const detail::nd_range_view &RangeView, + const kernel &Kernel, sycl::span Args, + const detail::code_location &CodeLoc, bool IsTopCodeLoc) { + detail::getSyclObjImpl(Queue)->submit_kernel_obj_direct_without_event( + RangeView, detail::getSyclObjImpl(Kernel), Args, CodeLoc, IsTopCodeLoc); +} + event submit_graph_direct_with_event_impl( const queue &Queue, ext::oneapi::experimental::command_graph< diff --git a/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct.cpp b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct.cpp new file mode 100644 index 0000000000000..50205cbd5ff1d --- /dev/null +++ b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct.cpp @@ -0,0 +1,147 @@ +// REQUIRES: aspect-usm_shared_allocations +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +// Tests the nd_launch overloads that take an already built sycl::kernel and its +// arguments as a parameter pack, which bypass the handler when every argument +// can be bound as plain bytes. Arguments that the scheduler has to track, an +// accessor here, must still reach the kernel through the command group path. + +#include +#include +#include +#include +#include +#include +#include +#include + +#include "common.hpp" + +namespace syclext = sycl::ext::oneapi; + +constexpr size_t N = 1024; +constexpr size_t WGSize = 8; + +enum class Sign : int { Plus = 1 }; + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((oneapiext::nd_range_kernel<1>)) +void addScalars(int *Ptr, int A, int B) { + size_t I = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + Ptr[I] += A + B; +} + +// A mixture of argument kinds, so that a wrong size or a wrong order shows up +// as a wrong result rather than as a silent pass. +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((oneapiext::nd_range_kernel<1>)) +void addMixed(int *Ptr, int A, unsigned long B, float C, Sign S) { + size_t I = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + Ptr[I] += + static_cast(S) * (A + static_cast(B) + static_cast(C)); +} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((oneapiext::nd_range_kernel<1>)) +void addViaAccessor(sycl::accessor Acc, int A) { + size_t I = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + Acc[I] += A; +} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((oneapiext::nd_range_kernel<1>)) +void usesScratch(int *Ptr) { + size_t I = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + int *Scratch = + reinterpret_cast(oneapiext::get_work_group_scratch_memory()); + Scratch[I % WGSize] = static_cast(I); + Ptr[I] = Scratch[I % WGSize]; +} + +template sycl::kernel getKernel(sycl::queue &Q) { + auto Bundle = + oneapiext::get_kernel_bundle( + Q.get_context()); + return Bundle.template ext_oneapi_get_kernel(); +} + +int main() { + sycl::queue Q{sycl::property::queue::in_order{}}; + int *Memory = sycl::malloc_shared(N, Q); + sycl::nd_range<1> Ndr{sycl::range<1>{N}, sycl::range<1>{WGSize}}; + + int Failed = 0; + + // Typed arguments: no handler is created for these. + sycl::kernel ScalarsKernel = getKernel(Q); + Q.memset(Memory, 0, N * sizeof(int)); + oneapiext::nd_launch(Q, Ndr, ScalarsKernel, Memory, 3, 4); + Q.wait(); + for (size_t I = 0; I < N; ++I) + Failed += Check(Memory, 7, I, "typed arguments"); + + // The same arguments as raw bytes, which is how a caller that only knows the + // signature as sizes has to pass them. + int A = 10, B = 20; + Q.memset(Memory, 0, N * sizeof(int)); + oneapiext::nd_launch(Q, Ndr, ScalarsKernel, + oneapiext::raw_kernel_arg{&Memory, sizeof(Memory)}, + oneapiext::raw_kernel_arg{&A, sizeof(A)}, + oneapiext::raw_kernel_arg{&B, sizeof(B)}); + Q.wait(); + for (size_t I = 0; I < N; ++I) + Failed += Check(Memory, 30, I, "raw_kernel_arg arguments"); + + // Mixed argument kinds, including an enum and a float. + Q.memset(Memory, 0, N * sizeof(int)); + oneapiext::nd_launch(Q, Ndr, getKernel(Q), Memory, 1, 2ul, 3.0f, + Sign::Plus); + Q.wait(); + for (size_t I = 0; I < N; ++I) + Failed += Check(Memory, 6, I, "mixed argument kinds"); + + // Ordering against the preceding commands on an in-order queue has to hold + // for a launch that bypasses the scheduler. + Q.memset(Memory, 0, N * sizeof(int)); + for (int I = 0; I < 8; ++I) + oneapiext::nd_launch(Q, Ndr, ScalarsKernel, Memory, 1, 0); + Q.wait(); + for (size_t I = 0; I < N; ++I) + Failed += Check(Memory, 8, I, "in-order accumulation"); + + // An accessor cannot be bound as plain bytes, so this has to fall back to the + // command group path and still produce the right result. + { + std::vector Data(N, 5); + sycl::buffer Buf{Data.data(), sycl::range<1>{N}}; + sycl::kernel AccessorKernel = getKernel(Q); + Q.submit([&](sycl::handler &CGH) { + sycl::accessor Acc{Buf, CGH, sycl::read_write}; + oneapiext::nd_launch(CGH, Ndr, AccessorKernel, Acc, 4); + }); + Q.wait(); + sycl::host_accessor Host{Buf}; + for (size_t I = 0; I < N; ++I) + Failed += + Check(&Host[0], 9, I, "accessor through the command group path"); + } + + // These overloads cannot carry a work_group_scratch_size property, so a + // kernel that allocates work group scratch memory has to be reported rather + // than launched, exactly as the command group path reports it. + { + bool Reported = false; + try { + oneapiext::nd_launch(Q, Ndr, getKernel(Q), Memory); + Q.wait_and_throw(); + } catch (const sycl::exception &E) { + Reported = E.code() == sycl::errc::memory_allocation; + } + if (!Reported) { + std::cout << "Failed: work group scratch memory without a size property " + "was not reported" + << std::endl; + ++Failed; + } + } + + sycl::free(Memory, Q); + return Failed; +} diff --git a/sycl/test/abi/layout_kernel_arg_view.cpp b/sycl/test/abi/layout_kernel_arg_view.cpp new file mode 100644 index 0000000000000..5ebd25d87eaea --- /dev/null +++ b/sycl/test/abi/layout_kernel_arg_view.cpp @@ -0,0 +1,17 @@ +// RUN: %clangxx -fsycl -c -fno-color-diagnostics -Xclang -fdump-record-layouts %s -o %t.out | FileCheck %s +// RUN: %clangxx -fsycl -fsycl-device-only -c -fno-color-diagnostics -Xclang -fdump-record-layouts %s -o %t.out | FileCheck %s +// REQUIRES: linux +// UNSUPPORTED: libcxx + +// clang-format off + +#include + + +SYCL_EXTERNAL void kernel_arg_view(sycl::detail::kernel_arg_view_v1::KernelArgView) {} +// CHECK: 0 | struct sycl::detail::KernelArgView +// CHECK-NEXT: 0 | const void * MPtr +// CHECK-NEXT: 8 | size_t MSize +// CHECK-NEXT: 16 | kernel_param_kind_t MKind +// CHECK-NEXT: | [sizeof=24, dsize=24, align=8, +// CHECK-NEXT: | nvsize=24, nvalign=8] diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 01137ab1ab81c..f86720c50889b 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3163,6 +3163,7 @@ _ZN4sycl3_V13ext6oneapi15filter_selectorC1ENS0_6detail11string_viewE _ZN4sycl3_V13ext6oneapi15filter_selectorC2ENS0_6detail11string_viewE _ZN4sycl3_V13ext6oneapilsERSoRKNS2_8bfloat16E _ZN4sycl3_V13ext6oneapirsERSiRNS2_8bfloat16E +_ZN4sycl3_V143submit_kernel_obj_direct_without_event_implERKNS0_5queueERKNS0_6detail16nd_range_view_v113nd_range_viewERKNS0_6kernelENS0_4spanIKNS4_18kernel_arg_view_v113KernelArgViewELm18446744073709551615EEERKNS4_13code_locationEb _ZN4sycl3_V14freeEPvRKNS0_5queueERKNS0_6detail13code_locationE _ZN4sycl3_V14freeEPvRKNS0_7contextERKNS0_6detail13code_locationE _ZN4sycl3_V15event13get_wait_listEv diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 70a39a2743b10..f7a70af12a53e 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -4388,6 +4388,7 @@ ?submit_graph_direct_without_event_impl@_V1@sycl@@YAXAEBVqueue@12@AEAV?$command_graph@$00@experimental@oneapi@ext@12@V?$span@$$CBVevent@_V1@sycl@@$0?0@12@AEBUcode_location@detail@12@@Z ?submit_kernel_direct_with_event_impl@_V1@sycl@@YA?AVevent@12@AEBVqueue@12@AEBVnd_range_view@nd_range_view_v1@detail@12@AEAVHostKernelRefBase@712@PEAVDeviceKernelInfo@712@V?$span@$$CBVevent@_V1@sycl@@$0?0@12@AEBU?$PropsHolder@Uwork_group_scratch_size@experimental@oneapi@ext@_V1@sycl@@Ucache_config@2intel@456@Uuse_root_sync_key@23456@Uwork_group_progress_key@23456@Usub_group_progress_key@23456@Uwork_item_progress_key@23456@U?$cluster_size@$00@cuda@23456@U?$cluster_size@$01@cuda@23456@U?$cluster_size@$02@cuda@23456@@kernel_launch_properties_v1@712@AEBUcode_location@712@_N@Z ?submit_kernel_direct_without_event_impl@_V1@sycl@@YAXAEBVqueue@12@AEBVnd_range_view@nd_range_view_v1@detail@12@AEAVHostKernelRefBase@612@PEAVDeviceKernelInfo@612@V?$span@$$CBVevent@_V1@sycl@@$0?0@12@AEBU?$PropsHolder@Uwork_group_scratch_size@experimental@oneapi@ext@_V1@sycl@@Ucache_config@2intel@456@Uuse_root_sync_key@23456@Uwork_group_progress_key@23456@Usub_group_progress_key@23456@Uwork_item_progress_key@23456@U?$cluster_size@$00@cuda@23456@U?$cluster_size@$01@cuda@23456@U?$cluster_size@$02@cuda@23456@@kernel_launch_properties_v1@612@AEBUcode_location@612@_N@Z +?submit_kernel_obj_direct_without_event_impl@_V1@sycl@@YAXAEBVqueue@12@AEBVnd_range_view@nd_range_view_v1@detail@12@AEBVkernel@12@V?$span@$$CBUKernelArgView@kernel_arg_view_v1@detail@_V1@sycl@@$0?0@12@AEBUcode_location@612@_N@Z ?submit_with_event_impl@queue@_V1@sycl@@AEBA?AVevent@23@AEBVtype_erased_cgfo_ty@detail@23@AEBVSubmissionInfo@2623@AEBUcode_location@623@_N@Z ?submit_without_event_impl@queue@_V1@sycl@@AEBAXAEBVtype_erased_cgfo_ty@detail@23@AEBVSubmissionInfo@2523@AEBUcode_location@523@_N@Z ?supportsUSMFill2D@handler@_V1@sycl@@AEAA_NXZ From ea6c051c3982d89f3ad819a8c2cb71a65d1efc3c Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 19 Aug 2026 11:09:03 +0000 Subject: [PATCH 2/7] [SYCL] Bind array arguments as bytes on the kernel object fast path `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. --- .../oneapi/experimental/enqueue_functions.hpp | 22 ++++-- sycl/source/detail/queue_impl.cpp | 10 ++- .../nd_launch_kernel_obj_array_arg.cpp | 78 +++++++++++++++++++ .../nd_launch_kernel_obj_direct.cpp | 10 ++- .../nd_launch_kernel_obj_direct_raw_ptr.cpp | 66 ++++++++++++++++ 5 files changed, 173 insertions(+), 13 deletions(-) create mode 100644 sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_array_arg.cpp create mode 100644 sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct_raw_ptr.cpp diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp index f4012ee928e9e..1405b41159d0c 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -99,29 +99,37 @@ template struct LaunchConfigAccess { } }; +// The argument type as the kernel sees it. Deliberately not `std::decay_t`, +// which turns an array into a pointer: an array has to keep being bound as the +// bytes it is, which is what `handler::setArgHelper` does with it. +template +using plain_arg_t = std::remove_cv_t>; + // An argument that can be bound as plain bytes, i.e. one that carries no // 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 inline constexpr bool is_plain_kernel_arg_v = - std::is_arithmetic_v> || std::is_enum_v> || - std::is_pointer_v> || - std::is_same_v, raw_kernel_arg>; + std::is_arithmetic_v> || std::is_enum_v> || + std::is_pointer_v> || + std::is_same_v, raw_kernel_arg>; // A pointer has to keep its kind. The runtime binds a pointer argument as // UR_EXP_KERNEL_ARG_TYPE_POINTER, which the OpenCL adapter passes to // clSetKernelArgMemPointerINTEL rather than to clSetKernelArg, so plain bytes -// are not a substitute. +// are not a substitute. The Native CPU adapter draws the same distinction: it +// puts a pointer argument straight into the argument slot, whereas a value +// argument lands there as the address of the adapter's own copy. template sycl::detail::KernelArgView makeKernelArgView(const T &Arg) { using sycl::detail::kernel_param_kind_t; - if constexpr (std::is_same_v, raw_kernel_arg>) + if constexpr (std::is_same_v, raw_kernel_arg>) return {RawKernelArgAccess::getData(Arg), RawKernelArgAccess::getSize(Arg), kernel_param_kind_t::kind_std_layout}; else - return {&Arg, sizeof(T), - std::is_pointer_v> + return {&Arg, sizeof(plain_arg_t), + std::is_pointer_v> ? kernel_param_kind_t::kind_pointer : kernel_param_kind_t::kind_std_layout}; } diff --git a/sycl/source/detail/queue_impl.cpp b/sycl/source/detail/queue_impl.cpp index f7f443f16d98e..f34d22baab6eb 100644 --- a/sycl/source/detail/queue_impl.cpp +++ b/sycl/source/detail/queue_impl.cpp @@ -894,6 +894,12 @@ void queue_impl::submit_kernel_obj_direct_without_event( KData.setNDRDesc(NDRDescT(RangeView)); KData.getArgs().reserve(Args.size()); + // The kernel may have come from a bundle, and the bundle has to travel with + // it the way the handler path lets it: it is what lets enqueueImpKernel + // initialize the device globals a bundle keeps to itself. + std::shared_ptr KernelBundleImpl = + KernelImpl->get_kernel_bundle(); + // This overload carries no properties, so a kernel that needs work group // scratch memory can never have been given a size. The handler path reports // that in handler.cpp, so report it here too rather than launching a kernel @@ -930,13 +936,13 @@ void queue_impl::submit_kernel_obj_direct_without_event( if (SchedulerBypass) return {submit_kernel_scheduler_bypass( KData, CGData.MEvents, /*EventNeeded*/ false, - KernelImpl.get(), /*KernelBundleImpPtr*/ nullptr, CodeLoc, + KernelImpl.get(), KernelBundleImpl.get(), CodeLoc, IsTopCodeLoc), /*SchedulerBypass*/ true}; auto CommandGroup = std::make_unique( KData.getNDRDesc(), /*HostKernel*/ nullptr, KernelImpl, - /*KernelBundle*/ nullptr, std::move(CGData), std::move(KData).getArgs(), + KernelBundleImpl, std::move(CGData), std::move(KData).getArgs(), *KData.getDeviceKernelInfoPtr(), std::vector>{}, std::vector>{}, detail::CGType::Kernel, diff --git a/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_array_arg.cpp b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_array_arg.cpp new file mode 100644 index 0000000000000..44caf806aab04 --- /dev/null +++ b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_array_arg.cpp @@ -0,0 +1,78 @@ +// REQUIRES: aspect-usm_shared_allocations +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +// Tests that an array argument reaches a `sycl::kernel` the same way through +// the handler-less nd_launch overload as it does through the handler. +// +// An array is bound as the bytes it is, which is what `handler::set_args` does +// with it, so the kernel below reads those 16 bytes as its parameter. The +// handler-less overload must not classify the array as a pointer argument: the +// runtime would then read the first bytes of the array as an address, which +// binds neither the bytes nor the array itself. +// +// Kernel arguments are sticky on the backend side and a bundle hands out the +// same underlying kernel for the same kernel id, so a correct launch would mask +// a wrong one. The handler-less path therefore goes first. + +#include +#include +#include +#include +#include +#include + +#include "common.hpp" + +namespace syclext = sycl::ext::oneapi; + +constexpr size_t N = 8; + +struct FourInts { + int A, B, C, D; +}; + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((oneapiext::nd_range_kernel<1>)) +void addFourInts(FourInts Values, int *Out) { + size_t I = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + Out[I] = Values.A + Values.B + Values.C + Values.D; +} + +template sycl::kernel getKernel(sycl::queue &Q) { + auto Bundle = + oneapiext::get_kernel_bundle( + Q.get_context()); + return Bundle.template ext_oneapi_get_kernel(); +} + +int main() { + sycl::queue Q{sycl::property::queue::in_order{}}; + + int *Out = sycl::malloc_shared(N, Q); + sycl::nd_range<1> Ndr{sycl::range<1>{N}, sycl::range<1>{N}}; + int Values[4] = {1, 2, 3, 4}; + constexpr int Expected = 1 + 2 + 3 + 4; + + int Failed = 0; + + sycl::kernel Kernel = getKernel(Q); + + // The handler-less path, which must bind the array as plain bytes. + Q.memset(Out, 0, N * sizeof(int)); + oneapiext::nd_launch(Q, Ndr, Kernel, Values, Out); + Q.wait(); + for (size_t I = 0; I < N; ++I) + Failed += Check(Out, Expected, I, "array without a handler"); + + // The command group path has to agree with it. + Q.memset(Out, 0, N * sizeof(int)); + Q.submit([&](sycl::handler &CGH) { + oneapiext::nd_launch(CGH, Ndr, Kernel, Values, Out); + }); + Q.wait(); + for (size_t I = 0; I < N; ++I) + Failed += Check(Out, Expected, I, "array through the command group path"); + + sycl::free(Out, Q); + return Failed; +} diff --git a/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct.cpp b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct.cpp index 50205cbd5ff1d..32ce9adc31218 100644 --- a/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct.cpp +++ b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct.cpp @@ -77,12 +77,14 @@ int main() { for (size_t I = 0; I < N; ++I) Failed += Check(Memory, 7, I, "typed arguments"); - // The same arguments as raw bytes, which is how a caller that only knows the - // signature as sizes has to pass them. + // The same scalar arguments as raw bytes, which is how a caller that only + // knows the signature as sizes has to pass them. The pointer stays typed: + // `raw_kernel_arg` always binds as plain bytes, and a USM pointer bound that + // way only reaches the kernel on Level Zero, hence the separate + // nd_launch_kernel_obj_direct_raw_ptr.cpp for that case. int A = 10, B = 20; Q.memset(Memory, 0, N * sizeof(int)); - oneapiext::nd_launch(Q, Ndr, ScalarsKernel, - oneapiext::raw_kernel_arg{&Memory, sizeof(Memory)}, + oneapiext::nd_launch(Q, Ndr, ScalarsKernel, Memory, oneapiext::raw_kernel_arg{&A, sizeof(A)}, oneapiext::raw_kernel_arg{&B, sizeof(B)}); Q.wait(); diff --git a/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct_raw_ptr.cpp b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct_raw_ptr.cpp new file mode 100644 index 0000000000000..3ff144334d385 --- /dev/null +++ b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct_raw_ptr.cpp @@ -0,0 +1,66 @@ +// REQUIRES: aspect-usm_shared_allocations +// REQUIRES: level_zero +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +// Tests the nd_launch overload that takes an already built sycl::kernel with +// every argument, the USM pointer included, passed as raw bytes. That is what a +// caller which only knows the signature as sizes has to do. +// +// A `raw_kernel_arg` is bound as UR_EXP_KERNEL_ARG_TYPE_VALUE, which for a +// pointer parameter is only equivalent to UR_EXP_KERNEL_ARG_TYPE_POINTER on +// Level Zero. The OpenCL adapter passes a value argument to clSetKernelArg, +// which rejects a USM pointer with CL_INVALID_MEM_OBJECT, and the Native CPU +// adapter puts the address of its own copy of the bytes into the argument slot +// instead of the pointer itself. Hence the Level Zero requirement above, the +// same restriction the RawKernelArg tests carry. + +#include +#include +#include +#include +#include +#include +#include + +#include "common.hpp" + +namespace syclext = sycl::ext::oneapi; + +constexpr size_t N = 1024; +constexpr size_t WGSize = 8; + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((oneapiext::nd_range_kernel<1>)) +void addScalars(int *Ptr, int A, int B) { + size_t I = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + Ptr[I] += A + B; +} + +template sycl::kernel getKernel(sycl::queue &Q) { + auto Bundle = + oneapiext::get_kernel_bundle( + Q.get_context()); + return Bundle.template ext_oneapi_get_kernel(); +} + +int main() { + sycl::queue Q{sycl::property::queue::in_order{}}; + int *Memory = sycl::malloc_shared(N, Q); + sycl::nd_range<1> Ndr{sycl::range<1>{N}, sycl::range<1>{WGSize}}; + + int Failed = 0; + + int A = 10, B = 20; + sycl::kernel ScalarsKernel = getKernel(Q); + Q.memset(Memory, 0, N * sizeof(int)); + oneapiext::nd_launch(Q, Ndr, ScalarsKernel, + oneapiext::raw_kernel_arg{&Memory, sizeof(Memory)}, + oneapiext::raw_kernel_arg{&A, sizeof(A)}, + oneapiext::raw_kernel_arg{&B, sizeof(B)}); + Q.wait(); + for (size_t I = 0; I < N; ++I) + Failed += Check(Memory, 30, I, "pointer and scalars as raw bytes"); + + sycl::free(Memory, Q); + return Failed; +} From 73dbc05fd15ddcdddb4fdc8c0bd04873920100c6 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Tue, 1 Sep 2026 08:44:11 +0000 Subject: [PATCH 3/7] [SYCL] Keep the handler's argument kinds on the kernel object fast path `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. --- .../oneapi/experimental/enqueue_functions.hpp | 46 ++++++++++----- sycl/source/detail/queue_impl.cpp | 14 ++++- .../kernel_arg_classification.cpp | 56 +++++++++++++++++++ 3 files changed, 100 insertions(+), 16 deletions(-) create mode 100644 sycl/test/extensions/enqueue_functions/kernel_arg_classification.cpp diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp index 1405b41159d0c..bf78cfe30b8b3 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -105,22 +105,45 @@ template struct LaunchConfigAccess { template using plain_arg_t = std::remove_cv_t>; +// An argument that is bound as its own bytes with no further interpretation. +template +inline constexpr bool is_scalar_kernel_arg_v = + std::is_arithmetic_v || std::is_enum_v || std::is_pointer_v; + // An argument that can be bound as plain bytes, i.e. one that carries no // 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. +// group path; `HasSpecialCaptures` in the runtime draws the same line. So is +// every other class type, which may be a struct with special types inside and +// then needs `kind_struct_with_special_type` instead. +// +// An array of scalars is bound as the bytes it is, which is what +// `handler::setArgHelper` does with it, so it belongs on this path. template inline constexpr bool is_plain_kernel_arg_v = - std::is_arithmetic_v> || std::is_enum_v> || - std::is_pointer_v> || + is_scalar_kernel_arg_v> || + (std::is_array_v> && + is_scalar_kernel_arg_v>>) || std::is_same_v, raw_kernel_arg>; -// A pointer has to keep its kind. The runtime binds a pointer argument as -// UR_EXP_KERNEL_ARG_TYPE_POINTER, which the OpenCL adapter passes to -// clSetKernelArgMemPointerINTEL rather than to clSetKernelArg, so plain bytes -// are not a substitute. The Native CPU adapter draws the same distinction: it -// puts a pointer argument straight into the argument slot, whereas a value -// argument lands there as the address of the adapter's own copy. +// The kind a plain argument has to carry. A pointer has to keep its kind: the +// runtime binds a pointer argument as UR_EXP_KERNEL_ARG_TYPE_POINTER, which the +// OpenCL adapter passes to clSetKernelArgMemPointerINTEL rather than to +// clSetKernelArg, so plain bytes are not a substitute. The Native CPU adapter +// draws the same distinction: it puts a pointer argument straight into the +// argument slot, whereas a value argument lands there as the address of the +// adapter's own copy. +// +// `cl_mem` is the one pointer that is not an address: it names a memory object +// and has to be bound as the bytes of the handle, which is the exception +// `handler::setArgHelper` makes for `OpenCLMemT`. +template +inline constexpr sycl::detail::kernel_param_kind_t plain_arg_kind_v = + (std::is_pointer_v> && + !std::is_same_v, sycl::OpenCLMemT>) + ? sycl::detail::kernel_param_kind_t::kind_pointer + : sycl::detail::kernel_param_kind_t::kind_std_layout; + template sycl::detail::KernelArgView makeKernelArgView(const T &Arg) { using sycl::detail::kernel_param_kind_t; @@ -128,10 +151,7 @@ sycl::detail::KernelArgView makeKernelArgView(const T &Arg) { return {RawKernelArgAccess::getData(Arg), RawKernelArgAccess::getSize(Arg), kernel_param_kind_t::kind_std_layout}; else - return {&Arg, sizeof(plain_arg_t), - std::is_pointer_v> - ? kernel_param_kind_t::kind_pointer - : kernel_param_kind_t::kind_std_layout}; + return {&Arg, sizeof(plain_arg_t), plain_arg_kind_v}; } template diff --git a/sycl/source/detail/queue_impl.cpp b/sycl/source/detail/queue_impl.cpp index f34d22baab6eb..b47c7f3dd0c78 100644 --- a/sycl/source/detail/queue_impl.cpp +++ b/sycl/source/detail/queue_impl.cpp @@ -923,6 +923,9 @@ void queue_impl::submit_kernel_obj_direct_without_event( // 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) { + // `ArgDesc` holds a `void *` because the kinds that carry an object + // rather than bytes hand it out as a mutable pointer. These arguments are + // bytes and this path only ever reads them, hence the cast. void *Value = const_cast(Args[I].MPtr); if (!SchedulerBypass) { const char *Bytes = static_cast(Args[I].MPtr); @@ -940,15 +943,20 @@ void queue_impl::submit_kernel_obj_direct_without_event( IsTopCodeLoc), /*SchedulerBypass*/ true}; + // Extract data to move KData + ur_kernel_cache_config_t KernelCacheConfig = KData.getKernelCacheConfig(); + bool IsCooperative = KData.isCooperative(); + bool UsesClusterLaunch = KData.usesClusterLaunch(); + size_t KernelWorkGroupMemorySize = KData.getKernelWorkGroupMemorySize(); + auto CommandGroup = std::make_unique( KData.getNDRDesc(), /*HostKernel*/ nullptr, KernelImpl, KernelBundleImpl, std::move(CGData), std::move(KData).getArgs(), *KData.getDeviceKernelInfoPtr(), std::vector>{}, std::vector>{}, detail::CGType::Kernel, - KData.getKernelCacheConfig(), KData.isCooperative(), - KData.usesClusterLaunch(), KData.getKernelWorkGroupMemorySize(), - CodeLoc); + KernelCacheConfig, IsCooperative, UsesClusterLaunch, + KernelWorkGroupMemorySize, CodeLoc); CommandGroup->MIsTopCodeLoc = IsTopCodeLoc; if (auto GraphImpl = getCommandGraph(); GraphImpl) diff --git a/sycl/test/extensions/enqueue_functions/kernel_arg_classification.cpp b/sycl/test/extensions/enqueue_functions/kernel_arg_classification.cpp new file mode 100644 index 0000000000000..6276ecbb7c14e --- /dev/null +++ b/sycl/test/extensions/enqueue_functions/kernel_arg_classification.cpp @@ -0,0 +1,56 @@ +// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s +// expected-no-diagnostics + +// Pins which arguments the handler-less `nd_launch` overloads bind themselves +// and with which kind, since neither is observable from an end to end test. +// Every case here has to agree with `handler::setArgHelper`. + +#include + +namespace oneapiext = sycl::ext::oneapi::experimental; +namespace ext_detail = sycl::ext::oneapi::experimental::detail; + +using sycl::detail::kernel_param_kind_t; + +enum class Sign : int { Plus = 1 }; +struct Trivial { + int A, B; +}; + +// Bound without a handler: no requirement for the scheduler to track. +static_assert(ext_detail::is_plain_kernel_arg_v); +static_assert(ext_detail::is_plain_kernel_arg_v); +static_assert(ext_detail::is_plain_kernel_arg_v); +static_assert(ext_detail::is_plain_kernel_arg_v); +static_assert(ext_detail::is_plain_kernel_arg_v); +static_assert(ext_detail::is_plain_kernel_arg_v); +static_assert(ext_detail::is_plain_kernel_arg_v); +// An array of scalars is bound as the bytes it is, so it belongs on this path. +static_assert(ext_detail::is_plain_kernel_arg_v); +static_assert(ext_detail::is_plain_kernel_arg_v); + +// Kept on the command group path. +static_assert(!ext_detail::is_plain_kernel_arg_v>); +static_assert(!ext_detail::is_plain_kernel_arg_v>); +// A class type may be a struct with special types inside, and an array of them +// would hide that just as well, so both keep using the handler. +static_assert(!ext_detail::is_plain_kernel_arg_v); +static_assert(!ext_detail::is_plain_kernel_arg_v); + +// A pointer is an address, and has to be bound as one. +static_assert(ext_detail::plain_arg_kind_v == + kernel_param_kind_t::kind_pointer); +static_assert(ext_detail::plain_arg_kind_v == + kernel_param_kind_t::kind_pointer); +// `cl_mem` is a pointer typedef but names a memory object, so it is bound as +// the bytes of the handle, the way the handler binds it. +static_assert(ext_detail::plain_arg_kind_v == + kernel_param_kind_t::kind_std_layout); +static_assert(ext_detail::plain_arg_kind_v == + kernel_param_kind_t::kind_std_layout); +static_assert(ext_detail::plain_arg_kind_v == + kernel_param_kind_t::kind_std_layout); +static_assert(ext_detail::plain_arg_kind_v == + kernel_param_kind_t::kind_std_layout); + +int main() { return 0; } From 38de0ed1c4e7a9657a38144d59cd3144be5f99e6 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 9 Sep 2026 09:56:23 +0200 Subject: [PATCH 4/7] Apply suggestion from @uditagarwal97 Co-authored-by: Udit Kumar Agarwal --- sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp index bf78cfe30b8b3..8bedc17af07e7 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -101,7 +101,7 @@ template struct LaunchConfigAccess { // The argument type as the kernel sees it. Deliberately not `std::decay_t`, // which turns an array into a pointer: an array has to keep being bound as the -// bytes it is, which is what `handler::setArgHelper` does with it. +// The argument type as the kernel sees it. template using plain_arg_t = std::remove_cv_t>; From 088a5d2d65b0b729d133bc3d2b2107fa4a4e70ca Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 9 Sep 2026 08:00:58 +0000 Subject: [PATCH 5/7] [SYCL] Fix up the applied suggestion for the plain_arg_t comment 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. --- .../sycl/ext/oneapi/experimental/enqueue_functions.hpp | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp index 8bedc17af07e7..7cd75dff95358 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -99,9 +99,7 @@ template struct LaunchConfigAccess { } }; -// The argument type as the kernel sees it. Deliberately not `std::decay_t`, -// which turns an array into a pointer: an array has to keep being bound as the -// The argument type as the kernel sees it. +// The argument type as the kernel sees it, with an array kept as an array. template using plain_arg_t = std::remove_cv_t>; From bb242a912ef7b129eca73058476226a59469834f Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 9 Sep 2026 08:50:29 +0000 Subject: [PATCH 6/7] [SYCL] Address review comments on the kernel argument classification 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. --- .../oneapi/experimental/enqueue_functions.hpp | 70 +++++++------------ .../kernel_arg_classification.cpp | 39 ++++++----- 2 files changed, 47 insertions(+), 62 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp index 7cd75dff95358..93aebaab7f8ea 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -99,57 +99,45 @@ template struct LaunchConfigAccess { } }; -// The argument type as the kernel sees it, with an array kept as an array. +// The argument type as the kernel sees it. template -using plain_arg_t = std::remove_cv_t>; +using unqualified_arg_t = std::remove_cv_t>; -// An argument that is bound as its own bytes with no further interpretation. +// An argument that is bound as its own bytes with no further interpretation or +// special captures. Excludes arguments of types Accessors, Stream, etc. template inline constexpr bool is_scalar_kernel_arg_v = std::is_arithmetic_v || std::is_enum_v || std::is_pointer_v; -// An argument that can be bound as plain bytes, i.e. one that carries no -// 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. So is -// every other class type, which may be a struct with special types inside and -// then needs `kind_struct_with_special_type` instead. -// -// An array of scalars is bound as the bytes it is, which is what -// `handler::setArgHelper` does with it, so it belongs on this path. +// An argument that can be bound without a handler: a scalar, an array of +// scalars, or a `raw_kernel_arg`. Anything else keeps using the command group +// path. template -inline constexpr bool is_plain_kernel_arg_v = - is_scalar_kernel_arg_v> || - (std::is_array_v> && - is_scalar_kernel_arg_v>>) || - std::is_same_v, raw_kernel_arg>; - -// The kind a plain argument has to carry. A pointer has to keep its kind: the -// runtime binds a pointer argument as UR_EXP_KERNEL_ARG_TYPE_POINTER, which the -// OpenCL adapter passes to clSetKernelArgMemPointerINTEL rather than to -// clSetKernelArg, so plain bytes are not a substitute. The Native CPU adapter -// draws the same distinction: it puts a pointer argument straight into the -// argument slot, whereas a value argument lands there as the address of the -// adapter's own copy. -// -// `cl_mem` is the one pointer that is not an address: it names a memory object -// and has to be bound as the bytes of the handle, which is the exception -// `handler::setArgHelper` makes for `OpenCLMemT`. +inline constexpr bool is_direct_kernel_arg_v = + is_scalar_kernel_arg_v> || + (std::is_array_v> && + is_scalar_kernel_arg_v>>) || + std::is_same_v, raw_kernel_arg>; + +// The kind such an argument is bound with. A pointer keeps its kind, since a +// backend may bind a pointer through a different entry point than bytes. +// `cl_mem` is the exception: it names a memory object, so it is bound as the +// bytes of the handle, as `handler::setArgHelper` does. template -inline constexpr sycl::detail::kernel_param_kind_t plain_arg_kind_v = - (std::is_pointer_v> && - !std::is_same_v, sycl::OpenCLMemT>) +inline constexpr sycl::detail::kernel_param_kind_t kernel_arg_kind_v = + (std::is_pointer_v> && + !std::is_same_v, sycl::OpenCLMemT>) ? sycl::detail::kernel_param_kind_t::kind_pointer : sycl::detail::kernel_param_kind_t::kind_std_layout; template sycl::detail::KernelArgView makeKernelArgView(const T &Arg) { using sycl::detail::kernel_param_kind_t; - if constexpr (std::is_same_v, raw_kernel_arg>) + if constexpr (std::is_same_v, raw_kernel_arg>) return {RawKernelArgAccess::getData(Arg), RawKernelArgAccess::getSize(Arg), kernel_param_kind_t::kind_std_layout}; else - return {&Arg, sizeof(plain_arg_t), plain_arg_kind_v}; + return {&Arg, sizeof(unqualified_arg_t), kernel_arg_kind_v}; } template @@ -462,17 +450,13 @@ void nd_launch(handler &CGH, nd_range Range, template void nd_launch(queue Q, nd_range Range, const kernel &KernelObj, ArgsT &&...Args) { - if constexpr ((detail::is_plain_kernel_arg_v && ...)) { - // Bind the arguments straight from this call, so that neither a handler nor - // a command group object has to be created. The array is one element longer - // than the pack so that a zero-argument kernel stays well formed. + // The handler-less path only takes arguments that can be bound directly, + // anything else goes through the handler overload above. + if constexpr ((detail::is_direct_kernel_arg_v && ...)) { + // The array is one element longer than the pack so that a zero-argument + // kernel stays well formed. const sycl::detail::KernelArgView ArgViews[sizeof...(ArgsT) + 1] = { detail::makeKernelArgView(Args)...}; - // An overload that ends in a parameter pack cannot take a trailing - // code_location parameter, so the location is the one this header sees, - // as it was when this overload went through submit(). Seed the TLS slot - // rather than leaving it default-constructed, which the instrumentation - // reads as a null file and function name. sycl::detail::tls_code_loc_t TlsCodeLocCapture{ sycl::detail::code_location::current()}; sycl::submit_kernel_obj_direct_without_event_impl( diff --git a/sycl/test/extensions/enqueue_functions/kernel_arg_classification.cpp b/sycl/test/extensions/enqueue_functions/kernel_arg_classification.cpp index 6276ecbb7c14e..956d3ed6f6817 100644 --- a/sycl/test/extensions/enqueue_functions/kernel_arg_classification.cpp +++ b/sycl/test/extensions/enqueue_functions/kernel_arg_classification.cpp @@ -18,39 +18,40 @@ struct Trivial { }; // Bound without a handler: no requirement for the scheduler to track. -static_assert(ext_detail::is_plain_kernel_arg_v); -static_assert(ext_detail::is_plain_kernel_arg_v); -static_assert(ext_detail::is_plain_kernel_arg_v); -static_assert(ext_detail::is_plain_kernel_arg_v); -static_assert(ext_detail::is_plain_kernel_arg_v); -static_assert(ext_detail::is_plain_kernel_arg_v); -static_assert(ext_detail::is_plain_kernel_arg_v); +static_assert(ext_detail::is_direct_kernel_arg_v); +static_assert(ext_detail::is_direct_kernel_arg_v); +static_assert(ext_detail::is_direct_kernel_arg_v); +static_assert(ext_detail::is_direct_kernel_arg_v); +static_assert(ext_detail::is_direct_kernel_arg_v); +static_assert(ext_detail::is_direct_kernel_arg_v); +static_assert(ext_detail::is_direct_kernel_arg_v); // An array of scalars is bound as the bytes it is, so it belongs on this path. -static_assert(ext_detail::is_plain_kernel_arg_v); -static_assert(ext_detail::is_plain_kernel_arg_v); +static_assert(ext_detail::is_direct_kernel_arg_v); +static_assert(ext_detail::is_direct_kernel_arg_v); // Kept on the command group path. -static_assert(!ext_detail::is_plain_kernel_arg_v>); -static_assert(!ext_detail::is_plain_kernel_arg_v>); +static_assert(!ext_detail::is_direct_kernel_arg_v>); +static_assert( + !ext_detail::is_direct_kernel_arg_v>); // A class type may be a struct with special types inside, and an array of them // would hide that just as well, so both keep using the handler. -static_assert(!ext_detail::is_plain_kernel_arg_v); -static_assert(!ext_detail::is_plain_kernel_arg_v); +static_assert(!ext_detail::is_direct_kernel_arg_v); +static_assert(!ext_detail::is_direct_kernel_arg_v); // A pointer is an address, and has to be bound as one. -static_assert(ext_detail::plain_arg_kind_v == +static_assert(ext_detail::kernel_arg_kind_v == kernel_param_kind_t::kind_pointer); -static_assert(ext_detail::plain_arg_kind_v == +static_assert(ext_detail::kernel_arg_kind_v == kernel_param_kind_t::kind_pointer); // `cl_mem` is a pointer typedef but names a memory object, so it is bound as // the bytes of the handle, the way the handler binds it. -static_assert(ext_detail::plain_arg_kind_v == +static_assert(ext_detail::kernel_arg_kind_v == kernel_param_kind_t::kind_std_layout); -static_assert(ext_detail::plain_arg_kind_v == +static_assert(ext_detail::kernel_arg_kind_v == kernel_param_kind_t::kind_std_layout); -static_assert(ext_detail::plain_arg_kind_v == +static_assert(ext_detail::kernel_arg_kind_v == kernel_param_kind_t::kind_std_layout); -static_assert(ext_detail::plain_arg_kind_v == +static_assert(ext_detail::kernel_arg_kind_v == kernel_param_kind_t::kind_std_layout); int main() { return 0; } From dd09f40a5c6d89ad40c381234f8d285bfd2ef782 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Thu, 10 Sep 2026 08:32:22 +0000 Subject: [PATCH 7/7] [SYCL] Name the unqualified argument type once per trait The traits spelled `unqualified_arg_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. --- .../oneapi/experimental/enqueue_functions.hpp | 20 +++++++++---------- 1 file changed, 10 insertions(+), 10 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp index 93aebaab7f8ea..4652061a55052 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -112,32 +112,32 @@ inline constexpr bool is_scalar_kernel_arg_v = // An argument that can be bound without a handler: a scalar, an array of // scalars, or a `raw_kernel_arg`. Anything else keeps using the command group // path. -template +template > inline constexpr bool is_direct_kernel_arg_v = - is_scalar_kernel_arg_v> || - (std::is_array_v> && - is_scalar_kernel_arg_v>>) || - std::is_same_v, raw_kernel_arg>; + is_scalar_kernel_arg_v || + (std::is_array_v && + is_scalar_kernel_arg_v>) || + std::is_same_v; // The kind such an argument is bound with. A pointer keeps its kind, since a // backend may bind a pointer through a different entry point than bytes. // `cl_mem` is the exception: it names a memory object, so it is bound as the // bytes of the handle, as `handler::setArgHelper` does. -template +template > inline constexpr sycl::detail::kernel_param_kind_t kernel_arg_kind_v = - (std::is_pointer_v> && - !std::is_same_v, sycl::OpenCLMemT>) + (std::is_pointer_v && !std::is_same_v) ? sycl::detail::kernel_param_kind_t::kind_pointer : sycl::detail::kernel_param_kind_t::kind_std_layout; template sycl::detail::KernelArgView makeKernelArgView(const T &Arg) { using sycl::detail::kernel_param_kind_t; - if constexpr (std::is_same_v, raw_kernel_arg>) + using ArgT = unqualified_arg_t; + if constexpr (std::is_same_v) return {RawKernelArgAccess::getData(Arg), RawKernelArgAccess::getSize(Arg), kernel_param_kind_t::kind_std_layout}; else - return {&Arg, sizeof(unqualified_arg_t), kernel_arg_kind_v}; + return {&Arg, sizeof(ArgT), kernel_arg_kind_v}; } template