From cf619c44da3c95ddc7991114f4410fd89f622708 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 12 Aug 2026 12:48:53 +0000 Subject: [PATCH 1/5] [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/5] [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 961867b9e63f2b4b62c05940b936ecfbd8a10504 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 12 Aug 2026 12:49:04 +0000 Subject: [PATCH 3/5] [SYCL] Add nd_launch overloads taking kernel arguments as a span The number of arguments is part of the type of a parameter pack, so a caller that only learns its argument list at run time has to instantiate the pack overload once for every count it may encounter. Add queue and handler overloads taking span, where the count is data instead. raw_kernel_arg is already type erased, so the two forms bind the same bytes and measure the same; only the caller's build differs. A container, or a non-const span, converts to span, but a parameter pack is an exact match and wins overload resolution, which would bind the container object itself as a single kernel argument. Diagnose that spelling rather than let it produce wrong results at run time. Bumps SYCL_EXT_ONEAPI_ENQUEUE_FUNCTIONS to 2 and documents the overloads. --- ...sycl_ext_oneapi_enqueue_functions.asciidoc | 43 +++++- .../oneapi/experimental/enqueue_functions.hpp | 38 ++++- sycl/include/sycl/queue.hpp | 13 ++ sycl/source/detail/queue_impl.cpp | 50 ++++++- sycl/source/detail/queue_impl.hpp | 19 +++ sycl/source/feature_test.hpp.in | 2 +- sycl/source/queue.cpp | 9 ++ .../nd_launch_kernel_obj_span.cpp | 134 ++++++++++++++++++ sycl/test/abi/sycl_symbols_linux.dump | 1 + sycl/test/abi/sycl_symbols_windows.dump | 1 + .../kernel_arg_span_errors.cpp | 33 +++++ 11 files changed, 332 insertions(+), 11 deletions(-) create mode 100644 sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_span.cpp create mode 100644 sycl/test/extensions/enqueue_functions/kernel_arg_span_errors.cpp diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc index a01b9504a5f28..a6eb1ba5ab237 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc @@ -165,8 +165,11 @@ implementation supports. |Description |1 -|The APIs of this experimental extension are not versioned, so the - feature-test macro always has this value. +|Initial version of this extension. + +|2 +|Adds the `nd_launch` overloads that take the arguments of a `sycl::kernel` as + a `span` of `raw_kernel_arg`. |=== @@ -525,6 +528,42 @@ a! ---- namespace sycl::ext::oneapi::experimental { +template +void nd_launch(sycl::queue q, sycl::nd_range r, + const sycl::kernel& k, span args); + +template +void nd_launch(sycl::handler &h, sycl::nd_range r, + const sycl::kernel& k, span args); + +} +---- +!==== +_Effects_: Enqueues a kernel object to the `sycl::queue` or `sycl::handler` +as a basic kernel, using the number of work-items specified by a +`sycl::nd_range`. The element `args[i]` is passed to the kernel as the argument +at index `i`. The sequence referenced by `args` need only remain valid until +the function returns. + +An application that calls one of the parameter pack overloads above with a +single argument whose type is convertible to `span` is +ill-formed. + +[_Note:_ These overloads exist for an application whose kernel argument list is +only known at run time, which would otherwise have to instantiate a parameter +pack overload for each number of arguments it may encounter. Since `span` has +no implicit conversion from a container, an application holding its arguments +in a `std::vector` passes `{args.data(), args.size()}`. Passing +a single `raw_kernel_arg` still selects the parameter pack overload. _end note_] + +a| +[frame=all,grid=none] +!==== +a! +[source,c++] +---- +namespace sycl::ext::oneapi::experimental { + template void nd_launch(sycl::queue q, launch_config, Properties> c, diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp index 1405b41159d0c..653f6da4b2487 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -444,7 +444,19 @@ 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 && ...)) { + // A container of raw_kernel_arg converts to the span the sibling overload + // takes, but a pack is an exact match and wins overload resolution, which + // would bind the container object itself as one argument. Diagnose it here, + // so that no other branch is instantiated for such a call. + constexpr bool ArgListPassedAsContainer = + sizeof...(ArgsT) == 1 && + (std::is_convertible_v> && ...); + if constexpr (ArgListPassedAsContainer) { + static_assert(!ArgListPassedAsContainer, + "The kernel argument list must be passed as " + "sycl::span, e.g. {Args.data(), " + "Args.size()}"); + } else 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. @@ -468,6 +480,30 @@ void nd_launch(queue Q, nd_range Range, const kernel &KernelObj, } } +template +void nd_launch(handler &CGH, nd_range Range, + const kernel &KernelObj, span Args) { + // set_arg only takes an rvalue raw_kernel_arg; an lvalue would select the + // generic overload and bind the object itself as the argument. + for (size_t I = 0; I < Args.size(); ++I) + CGH.set_arg(static_cast(I), raw_kernel_arg{Args[I]}); + CGH.parallel_for(Range, KernelObj); +} + +// Takes the kernel arguments as a contiguous sequence instead of a parameter +// pack, for a caller that only learns its argument list at run time and would +// otherwise need one instantiation of the pack overload per argument count. +template +void nd_launch(queue Q, nd_range Range, const kernel &KernelObj, + span Args, + const sycl::detail::code_location &CodeLoc = + sycl::detail::code_location::current()) { + sycl::detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + sycl::submit_kernel_obj_direct_without_event_impl( + Q, sycl::detail::nd_range_view(Range), KernelObj, Args, + TlsCodeLocCapture.query(), TlsCodeLocCapture.isToplevel()); +} + template void nd_launch(handler &CGH, launch_config, Properties> Config, diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp index ef3a17dbe546b..613307b077972 100644 --- a/sycl/include/sycl/queue.hpp +++ b/sycl/include/sycl/queue.hpp @@ -67,6 +67,10 @@ template auto get_native(const SyclObjectT &Obj) -> backend_return_t; +namespace ext::oneapi::experimental { +class raw_kernel_arg; +} // namespace ext::oneapi::experimental + // 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. @@ -75,6 +79,15 @@ void __SYCL_EXPORT submit_kernel_obj_direct_without_event_impl( const kernel &Kernel, sycl::span Args, const detail::code_location &CodeLoc, bool IsTopCodeLoc); +// As above, for an argument list that is already a contiguous sequence of +// `raw_kernel_arg`. Every element of such a sequence is plain bytes, so no +// per-argument kind has to be carried and the caller needs no conversion step. +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 f34d22baab6eb..e44d1eed9b128 100644 --- a/sycl/source/detail/queue_impl.cpp +++ b/sycl/source/detail/queue_impl.cpp @@ -883,11 +883,28 @@ EventImplPtr queue_impl::submit_kernel_direct_impl( /*InsertBarrierForInOrderCommand*/ false); } -void queue_impl::submit_kernel_obj_direct_without_event( +namespace { +// The two argument forms differ only in how one element yields the bytes to +// bind and their kind: a raw_kernel_arg is always plain bytes. +inline sycl::detail::KernelArgView +makeKernelArgView(const sycl::detail::KernelArgView &Arg) { + return Arg; +} +inline sycl::detail::KernelArgView +makeKernelArgView(const ext::oneapi::experimental::raw_kernel_arg &Arg) { + namespace syclex_detail = ext::oneapi::experimental::detail; + return {syclex_detail::RawKernelArgAccess::getData(Arg), + syclex_detail::RawKernelArgAccess::getSize(Arg), + sycl::detail::kernel_param_kind_t::kind_std_layout}; +} +} // namespace + +template +void queue_impl::submit_kernel_obj_direct_impl( const detail::nd_range_view &RangeView, const std::shared_ptr &KernelImpl, - sycl::span Args, - const detail::code_location &CodeLoc, bool IsTopCodeLoc) { + sycl::span Args, const detail::code_location &CodeLoc, + bool IsTopCodeLoc) { KernelData KData; KData.setDeviceKernelInfoPtr(&KernelImpl->getDeviceKernelInfo()); @@ -923,13 +940,14 @@ 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) { - void *Value = const_cast(Args[I].MPtr); + const sycl::detail::KernelArgView View = makeKernelArgView(Args[I]); + void *Value = const_cast(View.MPtr); if (!SchedulerBypass) { - const char *Bytes = static_cast(Args[I].MPtr); - CGData.MArgsStorage.emplace_back(Bytes, Bytes + Args[I].MSize); + const char *Bytes = static_cast(View.MPtr); + CGData.MArgsStorage.emplace_back(Bytes, Bytes + View.MSize); Value = CGData.MArgsStorage.back().data(); } - KData.addArg(Args[I].MKind, Value, static_cast(Args[I].MSize), + KData.addArg(View.MKind, Value, static_cast(View.MSize), static_cast(I)); } @@ -966,6 +984,24 @@ void queue_impl::submit_kernel_obj_direct_without_event( /*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) { + submit_kernel_obj_direct_impl(RangeView, KernelImpl, Args, CodeLoc, + IsTopCodeLoc); +} + +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) { + submit_kernel_obj_direct_impl(RangeView, KernelImpl, Args, CodeLoc, + IsTopCodeLoc); +} + 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 d1e6db6f9fb80..d3a5253f93a36 100644 --- a/sycl/source/detail/queue_impl.hpp +++ b/sycl/source/detail/queue_impl.hpp @@ -390,6 +390,14 @@ class queue_impl : public std::enable_shared_from_this { sycl::span Args, const detail::code_location &CodeLoc, bool IsTopCodeLoc); + /// As above, for arguments that are already a sequence of `raw_kernel_arg`, + /// each of which is plain bytes. + 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) { @@ -998,6 +1006,17 @@ class queue_impl : public std::enable_shared_from_this { SubmitCommandFuncType &SubmitCommandFunc, detail::CGType Type, bool InsertBarrierForInOrderCommand); + /// Shared implementation of the two submit_kernel_obj_direct_without_event + /// overloads. The element type only decides how one argument yields its bytes + /// and kind, so the submission itself is written once. Instantiated for + /// KernelArgView and for raw_kernel_arg in queue_impl.cpp. + template + void submit_kernel_obj_direct_impl( + const detail::nd_range_view &RangeView, + const std::shared_ptr &KernelImpl, + sycl::span Args, const detail::code_location &CodeLoc, + bool IsTopCodeLoc); + /// Performs barrier submission to the queue. /// /// \param DepEvents is a vector of dependencies of the operation. diff --git a/sycl/source/feature_test.hpp.in b/sycl/source/feature_test.hpp.in index 6aebc773b514c..f1a943f77fc1e 100644 --- a/sycl/source/feature_test.hpp.in +++ b/sycl/source/feature_test.hpp.in @@ -94,7 +94,7 @@ inline namespace _V1 { #define SYCL_EXT_ONEAPI_FORWARD_PROGRESS 1 #define SYCL_EXT_ONEAPI_FREE_FUNCTION_KERNELS 1 #define SYCL_EXT_ONEAPI_PROD 1 -#define SYCL_EXT_ONEAPI_ENQUEUE_FUNCTIONS 1 +#define SYCL_EXT_ONEAPI_ENQUEUE_FUNCTIONS 2 #define SYCL_EXT_ONEAPI_RAW_KERNEL_ARG 1 #define SYCL_EXT_ONEAPI_PROFILING_TAG 1 #define SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND 2 diff --git a/sycl/source/queue.cpp b/sycl/source/queue.cpp index bdbf0b00b6fe1..2e1bccd24cb45 100644 --- a/sycl/source/queue.cpp +++ b/sycl/source/queue.cpp @@ -327,6 +327,15 @@ void submit_kernel_obj_direct_without_event_impl( RangeView, detail::getSyclObjImpl(Kernel), Args, CodeLoc, 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_span.cpp b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_span.cpp new file mode 100644 index 0000000000000..6dbf3b449f325 --- /dev/null +++ b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_span.cpp @@ -0,0 +1,134 @@ +// REQUIRES: aspect-usm_shared_allocations +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +// Tests the nd_launch overloads that take the arguments of a sycl::kernel as a +// span of raw_kernel_arg, i.e. an argument list whose length is only known at +// run time. They have to bind the same arguments in the same order as the +// parameter pack overloads, on the queue and on the handler alike. + +#include +#include +#include +#include +#include +#include +#include + +#include "common.hpp" + +#include + +namespace syclext = sycl::ext::oneapi; + +static_assert(SYCL_EXT_ONEAPI_ENQUEUE_FUNCTIONS >= 2, + "The span overloads require version 2 of the extension"); + +constexpr size_t N = 1024; +constexpr size_t WGSize = 8; + +// A mixture of argument sizes, 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, long B, float C, char D) { + size_t I = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + Ptr[I] += A + static_cast(B) + static_cast(C) + D; +} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((oneapiext::nd_range_kernel<1>)) +void increment(int *Ptr) { + size_t I = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + Ptr[I] += 1; +} + +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{}}; + sycl::kernel Kernel = getKernel(Q); + + int *Memory = sycl::malloc_shared(N, Q); + sycl::nd_range<1> Ndr{sycl::range<1>{N}, sycl::range<1>{WGSize}}; + + int A = 1; + long B = 20; + float C = 300.0f; + char D = 4; + constexpr int Sum = 1 + 20 + 300 + 4; + + // The argument list is built at run time, which is the case these overloads + // exist for. + std::vector Args; + Args.emplace_back(&Memory, sizeof(Memory)); + Args.emplace_back(&A, sizeof(A)); + Args.emplace_back(&B, sizeof(B)); + Args.emplace_back(&C, sizeof(C)); + Args.emplace_back(&D, sizeof(D)); + sycl::span ArgSpan{Args.data(), Args.size()}; + + int Failed = 0; + + // A run of launches through one span, so that an argument bound to storage + // that does not outlive a single call would show up. + constexpr int Launches = 8; + Q.memset(Memory, 0, N * sizeof(int)); + for (int I = 0; I < Launches; ++I) + oneapiext::nd_launch(Q, Ndr, Kernel, ArgSpan); + Q.wait(); + for (size_t I = 0; I < N; ++I) + Failed += Check(Memory, Sum * Launches, I, "span overload"); + + // The parameter pack overload has to agree element for element. + Q.memset(Memory, 0, N * sizeof(int)); + oneapiext::nd_launch(Q, Ndr, Kernel, + oneapiext::raw_kernel_arg{&Memory, sizeof(Memory)}, + oneapiext::raw_kernel_arg{&A, sizeof(A)}, + oneapiext::raw_kernel_arg{&B, sizeof(B)}, + oneapiext::raw_kernel_arg{&C, sizeof(C)}, + oneapiext::raw_kernel_arg{&D, sizeof(D)}); + Q.wait(); + for (size_t I = 0; I < N; ++I) + Failed += Check(Memory, Sum, I, "parameter pack overload"); + + // And so does the handler form of the span overload. + Q.memset(Memory, 0, N * sizeof(int)); + Q.submit([&](sycl::handler &CGH) { + oneapiext::nd_launch(CGH, Ndr, Kernel, ArgSpan); + }).wait(); + for (size_t I = 0; I < N; ++I) + Failed += Check(Memory, Sum, I, "handler form of the span overload"); + + // A one element span is the boundary against the parameter pack overload, + // which a single raw_kernel_arg selects instead. + std::vector OneArg{{&Memory, sizeof(Memory)}}; + Q.memset(Memory, 0, N * sizeof(int)); + oneapiext::nd_launch(Q, Ndr, getKernel(Q), + sycl::span{ + OneArg.data(), OneArg.size()}); + Q.wait(); + for (size_t I = 0; I < N; ++I) + Failed += Check(Memory, 1, I, "one element span"); + + // A dependency the scheduler has to track forces the command group path, + // which the span form has to take as well. + { + sycl::buffer Buf{sycl::range<1>{N}}; + Q.memset(Memory, 0, N * sizeof(int)); + Q.submit([&](sycl::handler &CGH) { + sycl::accessor Acc{Buf, CGH, sycl::write_only, sycl::no_init}; + CGH.fill(Acc, 7); + }); + oneapiext::nd_launch(Q, Ndr, Kernel, ArgSpan); + Q.wait(); + for (size_t I = 0; I < N; ++I) + Failed += Check(Memory, Sum, I, "span form behind a buffer dependency"); + } + + sycl::free(Memory, Q); + return Failed; +} diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index f86720c50889b..a460dec11dea2 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_4spanIKNS0_3ext6oneapi12experimental14raw_kernel_argELm18446744073709551615EEERKNS4_13code_locationEb _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 diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index f7a70af12a53e..13741eb2d21f7 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -4389,6 +4389,7 @@ ?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_kernel_obj_direct_without_event_impl@_V1@sycl@@YAXAEBVqueue@12@AEBVnd_range_view@nd_range_view_v1@detail@12@AEBVkernel@12@V?$span@$$CBVraw_kernel_arg@experimental@oneapi@ext@_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 diff --git a/sycl/test/extensions/enqueue_functions/kernel_arg_span_errors.cpp b/sycl/test/extensions/enqueue_functions/kernel_arg_span_errors.cpp new file mode 100644 index 0000000000000..f5f0dfb8d8305 --- /dev/null +++ b/sycl/test/extensions/enqueue_functions/kernel_arg_span_errors.cpp @@ -0,0 +1,33 @@ +// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s + +// An argument list held in a container converts to the span that the nd_launch +// span overload takes, but a parameter pack is an exact match and wins overload +// resolution. Without a diagnostic the container object itself would be bound +// as a single kernel argument, which compiles for any trivially copyable +// wrapper and produces wrong results at run time. + +#include + +#include + +namespace oneapiext = sycl::ext::oneapi::experimental; + +void argument_list_must_be_a_span(sycl::queue Q, sycl::nd_range<1> Range, + const sycl::kernel &Kernel) { + int Value = 1; + std::vector Args{{&Value, sizeof(Value)}}; + + // expected-error@sycl/ext/oneapi/experimental/enqueue_functions.hpp:* {{The kernel argument list must be passed as sycl::span}} + oneapiext::nd_launch(Q, Range, Kernel, Args); + + sycl::span Mutable{Args.data(), Args.size()}; + // expected-error@sycl/ext/oneapi/experimental/enqueue_functions.hpp:* {{The kernel argument list must be passed as sycl::span}} + oneapiext::nd_launch(Q, Range, Kernel, Mutable); + + // The spelling the diagnostic asks for, and a single argument that happens to + // be a raw_kernel_arg, both have to keep working. + sycl::span AsSpan{Args.data(), Args.size()}; + oneapiext::nd_launch(Q, Range, Kernel, AsSpan); + oneapiext::nd_launch(Q, Range, Kernel, + oneapiext::raw_kernel_arg{&Value, sizeof(Value)}); +} From a861e89eebb744883dfb5b059806c3a3f76f9ff2 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 19 Aug 2026 15:09:56 +0000 Subject: [PATCH 4/5] [SYCL] Let raw_kernel_arg say that an argument is a pointer A raw_kernel_arg carried only bytes, so the span overloads bound every element as a value argument. A pointer bound that way only reaches the kernel 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. The argument list these overloads exist for is a pointer plus scalars, so the shape they are meant to serve was the one that did not work anywhere else. Add a pointer form of raw_kernel_arg. It takes the address of the pointer, the way the byte form takes the address of the bytes, so that passing the pointer itself does not compile. handler::setArgHelper and both makeKernelArgView overloads bind such an argument as a pointer. The graph extension needs no change: a dynamic parameter stores the raw_kernel_arg object itself and an update rewrites the bytes the node holds, which for a pointer argument are the pointer. SYCL_EXT_ONEAPI_RAW_KERNEL_ARG becomes 2. The E2E test of the span overloads no longer needs a Level Zero requirement, and neither does the one that passes every argument through raw_kernel_arg. New tests cover the handler path on OpenCL, a graph dynamic parameter that rebinds a pointer, the layout of raw_kernel_arg, and that passing a pointer by value does not compile. --- ...sycl_ext_oneapi_enqueue_functions.asciidoc | 15 +++- .../sycl_ext_oneapi_raw_kernel_arg.asciidoc | 67 +++++++++++++-- .../oneapi/experimental/enqueue_functions.hpp | 4 +- .../oneapi/experimental/raw_kernel_arg.hpp | 19 +++++ sycl/include/sycl/handler.hpp | 7 +- sycl/include/sycl/queue.hpp | 4 +- sycl/source/detail/queue_impl.cpp | 7 +- sycl/source/detail/queue_impl.hpp | 2 +- sycl/source/feature_test.hpp.in | 2 +- .../nd_launch_kernel_obj_direct.cpp | 7 +- .../nd_launch_kernel_obj_direct_raw_ptr.cpp | 31 +++---- .../nd_launch_kernel_obj_span.cpp | 22 +++-- .../update_with_raw_kernel_arg_pointer.cpp | 83 +++++++++++++++++++ sycl/test-e2e/RawKernelArg/pointer_arg.cpp | 71 ++++++++++++++++ sycl/test/abi/layout_raw_kernel_arg.cpp | 21 +++++ .../raw_kernel_arg/pointer_arg_errors.cpp | 37 +++++++++ 16 files changed, 351 insertions(+), 48 deletions(-) create mode 100644 sycl/test-e2e/Graph/Update/update_with_raw_kernel_arg_pointer.cpp create mode 100644 sycl/test-e2e/RawKernelArg/pointer_arg.cpp create mode 100644 sycl/test/abi/layout_raw_kernel_arg.cpp create mode 100644 sycl/test/extensions/raw_kernel_arg/pointer_arg_errors.cpp diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc index a6eb1ba5ab237..1c16e1f5e0dfd 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc @@ -41,6 +41,11 @@ This extension is written against the SYCL 2020 revision 7 specification. All references below to the "core SYCL specification" or to section numbers in the SYCL specification refer to that revision. +The `nd_launch` overloads that take the kernel arguments as a `span` depend on +version 2 of +link:../experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc[sycl_ext_oneapi_raw_kernel_arg], +which is what lets an element of that sequence represent a pointer argument. + == Status @@ -542,8 +547,8 @@ void nd_launch(sycl::handler &h, sycl::nd_range r, _Effects_: Enqueues a kernel object to the `sycl::queue` or `sycl::handler` as a basic kernel, using the number of work-items specified by a `sycl::nd_range`. The element `args[i]` is passed to the kernel as the argument -at index `i`. The sequence referenced by `args` need only remain valid until -the function returns. +at index `i`, as if it had been passed to `set_arg`. The sequence referenced by +`args` need only remain valid until the function returns. An application that calls one of the parameter pack overloads above with a single argument whose type is convertible to `span` is @@ -556,6 +561,12 @@ no implicit conversion from a container, an application holding its arguments in a `std::vector` passes `{args.data(), args.size()}`. Passing a single `raw_kernel_arg` still selects the parameter pack overload. _end note_] +[_Note:_ An element that represents a pointer argument has to be constructed +with the `pointer_arg` overload of `raw_kernel_arg`, as described in +link:../experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc[sycl_ext_oneapi_raw_kernel_arg], +since a `raw_kernel_arg` holding the bytes of a pointer is only guaranteed to +bind that pointer on the Level Zero backend. _end note_] + a| [frame=all,grid=none] !==== diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc index 6226aa1ec8be9..6826c2c9b719f 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc @@ -54,11 +54,17 @@ specification.* == Backend support status -This extension is currently implemented in {dpcpp} only for GPU devices and -only when using the Level Zero backend. Attempting to use this extension in -kernels that run on other devices or backends may result in undefined -behavior. Be aware that the compiler is not able to issue a diagnostic to -warn you if this happens. +This extension is currently implemented in {dpcpp} for the Level Zero, OpenCL, +CUDA and Native CPU backends. Attempting to use this extension on other +backends may result in undefined behavior. Be aware that the compiler is not +able to issue a diagnostic to warn you if this happens. + +A kernel argument that is a pointer must be constructed with the `pointer_arg` +overload described below. Passing the byte representation of a pointer to the +byte overload is only guaranteed to bind that pointer on the Level Zero backend, +because a backend may take a pointer argument through a different entry point +than the one that takes plain bytes, and it cannot tell the two apart from the +bytes alone. == Overview @@ -85,6 +91,19 @@ h.set_arg(1, sycl::ext::oneapi::experimental::raw_kernel_arg(opaque_type, nbytes h.parallel_for(range, kernel); ---- +An argument that is a pointer says so, since a backend may bind a pointer +through a different entry point than a sequence of bytes: + +[source,c++] +---- +namespace syclex = sycl::ext::oneapi::experimental; + +int* ptr = sycl::malloc_device(n, q); +... +h.set_arg(0, syclex::raw_kernel_arg(&ptr, syclex::pointer_arg)); +h.parallel_for(range, kernel); +---- + == Specification @@ -104,8 +123,10 @@ implementation supports. |Description |1 -|The APIs of this experimental extension are not versioned, so the - feature-test macro always has this value. +|Initial version of this extension. + +|2 +|Adds the `raw_kernel_arg` constructor that takes a pointer argument. |=== === The `raw_kernel_arg` class @@ -117,9 +138,15 @@ kernel arguments via a raw byte representation. ---- namespace sycl::ext::oneapi::experimental { + struct pointer_arg_t {}; + inline constexpr pointer_arg_t pointer_arg{}; + class raw_kernel_arg { public: raw_kernel_arg(const void* bytes, size_t count); + + template + raw_kernel_arg(T* const* pointer_location, pointer_arg_t tag); }; } // namespace sycl::ext::oneapi::experimental @@ -131,13 +158,34 @@ raw_kernel_arg(const void* bytes, size_t count); ---- _Preconditions_: `bytes` must point to an array of at least `count` bytes, which is the byte representation of a kernel argument that is trivially -copyable. +copyable. If the argument is a pointer, only the Level Zero backend is +guaranteed to bind it; see _Backend support status_ above. _Effects_: Constructs a `raw_kernel_arg` representing a view of the `count` bytes starting at the address specified by `bytes`. Since the `raw_kernel_arg` object is only a view, the caller must ensure that the lifetime of the `bytes` memory lasts at least as long as the lifetime of the `raw_kernel_arg` object. +[source,c++] +---- +template +raw_kernel_arg(T* const* pointer_location, pointer_arg_t tag); +---- +_Preconditions_: `pointer_location` must point to a pointer that is a valid +kernel argument, such as a pointer to memory allocated by one of the USM +allocation functions. + +_Effects_: Constructs a `raw_kernel_arg` representing a view of the pointer +stored at `pointer_location`, to be bound as a pointer argument rather than as +the `sizeof(T*)` bytes it is made of. Since the `raw_kernel_arg` object is only +a view, the caller must ensure that the lifetime of the pointer object at +`pointer_location` lasts at least as long as the lifetime of the +`raw_kernel_arg` object. + +[_Note:_ The constructor takes the address of the pointer, in keeping with the +byte overload taking the address of the bytes, so that passing the pointer +itself does not compile. _{endnote}_] + === Using a raw kernel argument Instances of `raw_kernel_arg` are passed to kernels via the existing `set_arg` @@ -154,7 +202,8 @@ argument in `args` was passed to `set_arg` ", adding a new overload of void set_arg(int argIndex, sycl::ext::oneapi::experimental::raw_kernel_arg&& arg); ---- _Effects_: Sets the kernel argument associated with index `argIndex` using the -bytes represented by `arg`. +bytes represented by `arg`, or as a pointer argument if `arg` was constructed +with the `pointer_arg` overload. == Issues diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp index 653f6da4b2487..392924cc242e9 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -126,7 +126,9 @@ 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}; + RawKernelArgAccess::isPointer(Arg) + ? kernel_param_kind_t::kind_pointer + : kernel_param_kind_t::kind_std_layout}; else return {&Arg, sizeof(plain_arg_t), std::is_pointer_v> 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 04ce9fd47c1e8..933a3550a1fa3 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/raw_kernel_arg.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/raw_kernel_arg.hpp @@ -21,14 +21,32 @@ class dynamic_parameter_impl; struct RawKernelArgAccess; } // namespace detail +// Tells the raw_kernel_arg constructor below that the argument it is given is a +// pointer rather than a sequence of bytes to copy. +struct pointer_arg_t {}; +inline constexpr pointer_arg_t pointer_arg{}; + class raw_kernel_arg { public: raw_kernel_arg(const void *bytes, size_t count) : MArgData(bytes), MArgSize(count) {} + // A pointer argument is not interchangeable with the bytes it is made of: a + // backend may have to be told that an argument is a pointer to bind it at + // all, as OpenCL does, where a USM pointer goes to + // clSetKernelArgMemPointerINTEL rather than to clSetKernelArg. Takes the + // address of the pointer, like the byte form takes the address of the bytes, + // so that passing the pointer itself does not compile. + template + raw_kernel_arg(T *const *pointer_location, pointer_arg_t) + : MArgData(pointer_location), MArgSize(sizeof(T *)), MIsPointer(true) {} + private: const void *MArgData; size_t MArgSize; + // Appended last, so that the offsets of the members above, which the library + // reads on the paths that predate this one, stay where they were. + bool MIsPointer = false; friend class sycl::handler; // For sycl_ext_oneapi_graph integration @@ -42,6 +60,7 @@ namespace detail { 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; } + static bool isPointer(const raw_kernel_arg &Arg) { return Arg.MIsPointer; } }; } // namespace detail diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 8354570fd9054..845a3ed0d0948 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -603,8 +603,11 @@ class __SYCL_EXPORT handler { void setArgHelper(int ArgIndex, sycl::ext::oneapi::experimental::raw_kernel_arg &&Arg) { auto StoredArg = storeRawArg(Arg); - addArg(detail::kernel_param_kind_t::kind_std_layout, StoredArg, - Arg.MArgSize, ArgIndex); + // A pointer argument has to be bound as one: a backend may reach for it + // through a different entry point than the one that takes plain bytes. + addArg(Arg.MIsPointer ? detail::kernel_param_kind_t::kind_pointer + : detail::kernel_param_kind_t::kind_std_layout, + StoredArg, Arg.MArgSize, ArgIndex); } /// Registers a dynamic parameter with the handler for later association with diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp index 613307b077972..5dff032bbdcb8 100644 --- a/sycl/include/sycl/queue.hpp +++ b/sycl/include/sycl/queue.hpp @@ -80,8 +80,8 @@ void __SYCL_EXPORT submit_kernel_obj_direct_without_event_impl( const detail::code_location &CodeLoc, bool IsTopCodeLoc); // As above, for an argument list that is already a contiguous sequence of -// `raw_kernel_arg`. Every element of such a sequence is plain bytes, so no -// per-argument kind has to be carried and the caller needs no conversion step. +// `raw_kernel_arg`, which carries its own kind per element, so the caller needs +// no conversion step. void __SYCL_EXPORT submit_kernel_obj_direct_without_event_impl( const queue &Queue, const detail::nd_range_view &RangeView, const kernel &Kernel, diff --git a/sycl/source/detail/queue_impl.cpp b/sycl/source/detail/queue_impl.cpp index e44d1eed9b128..19b76b670f44d 100644 --- a/sycl/source/detail/queue_impl.cpp +++ b/sycl/source/detail/queue_impl.cpp @@ -885,7 +885,8 @@ EventImplPtr queue_impl::submit_kernel_direct_impl( namespace { // The two argument forms differ only in how one element yields the bytes to -// bind and their kind: a raw_kernel_arg is always plain bytes. +// bind and their kind: a raw_kernel_arg carries plain bytes unless it was built +// as a pointer argument. inline sycl::detail::KernelArgView makeKernelArgView(const sycl::detail::KernelArgView &Arg) { return Arg; @@ -895,7 +896,9 @@ makeKernelArgView(const ext::oneapi::experimental::raw_kernel_arg &Arg) { namespace syclex_detail = ext::oneapi::experimental::detail; return {syclex_detail::RawKernelArgAccess::getData(Arg), syclex_detail::RawKernelArgAccess::getSize(Arg), - sycl::detail::kernel_param_kind_t::kind_std_layout}; + syclex_detail::RawKernelArgAccess::isPointer(Arg) + ? sycl::detail::kernel_param_kind_t::kind_pointer + : sycl::detail::kernel_param_kind_t::kind_std_layout}; } } // namespace diff --git a/sycl/source/detail/queue_impl.hpp b/sycl/source/detail/queue_impl.hpp index d3a5253f93a36..e62dc12b9ec76 100644 --- a/sycl/source/detail/queue_impl.hpp +++ b/sycl/source/detail/queue_impl.hpp @@ -391,7 +391,7 @@ class queue_impl : public std::enable_shared_from_this { const detail::code_location &CodeLoc, bool IsTopCodeLoc); /// As above, for arguments that are already a sequence of `raw_kernel_arg`, - /// each of which is plain bytes. + /// each of which carries its own kind. void submit_kernel_obj_direct_without_event( const detail::nd_range_view &RangeView, const std::shared_ptr &KernelImpl, diff --git a/sycl/source/feature_test.hpp.in b/sycl/source/feature_test.hpp.in index f1a943f77fc1e..a8d7396ec122d 100644 --- a/sycl/source/feature_test.hpp.in +++ b/sycl/source/feature_test.hpp.in @@ -95,7 +95,7 @@ inline namespace _V1 { #define SYCL_EXT_ONEAPI_FREE_FUNCTION_KERNELS 1 #define SYCL_EXT_ONEAPI_PROD 1 #define SYCL_EXT_ONEAPI_ENQUEUE_FUNCTIONS 2 -#define SYCL_EXT_ONEAPI_RAW_KERNEL_ARG 1 +#define SYCL_EXT_ONEAPI_RAW_KERNEL_ARG 2 #define SYCL_EXT_ONEAPI_PROFILING_TAG 1 #define SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND 2 #define SYCL_EXT_ONEAPI_GET_KERNEL_INFO 1 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 32ce9adc31218..299023f94b896 100644 --- a/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct.cpp +++ b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct.cpp @@ -78,10 +78,9 @@ int main() { Failed += Check(Memory, 7, I, "typed arguments"); // 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. + // knows the signature as sizes has to pass them. Mixing them with a typed + // pointer has to work; nd_launch_kernel_obj_direct_raw_ptr.cpp covers an + // argument list that is raw throughout. int A = 10, B = 20; Q.memset(Memory, 0, N * sizeof(int)); oneapiext::nd_launch(Q, Ndr, ScalarsKernel, Memory, 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 index 3ff144334d385..4d1f4c111161a 100644 --- 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 @@ -1,19 +1,18 @@ // 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. +// every argument passed through raw_kernel_arg, the USM pointer included. 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. +// The pointer is passed through the pointer form of raw_kernel_arg. Passing the +// bytes of a pointer instead would bind it as a value argument, which only +// reaches the kernel 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. The RawKernelArg tests cover +// the byte form of a pointer, on Level Zero only. #include #include @@ -53,13 +52,15 @@ int main() { 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)}); + oneapiext::nd_launch( + Q, Ndr, ScalarsKernel, + oneapiext::raw_kernel_arg{&Memory, oneapiext::pointer_arg}, + 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"); + Failed += + Check(Memory, 30, I, "pointer and scalars through raw_kernel_arg"); sycl::free(Memory, Q); return Failed; diff --git a/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_span.cpp b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_span.cpp index 6dbf3b449f325..844fc3ba54e00 100644 --- a/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_span.cpp +++ b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_span.cpp @@ -62,9 +62,11 @@ int main() { constexpr int Sum = 1 + 20 + 300 + 4; // The argument list is built at run time, which is the case these overloads - // exist for. + // exist for. The pointer says that it is one, so that it is bound as a + // pointer rather than as the bytes it is made of, which only the Level Zero + // backend binds as a pointer. std::vector Args; - Args.emplace_back(&Memory, sizeof(Memory)); + Args.emplace_back(&Memory, oneapiext::pointer_arg); Args.emplace_back(&A, sizeof(A)); Args.emplace_back(&B, sizeof(B)); Args.emplace_back(&C, sizeof(C)); @@ -85,12 +87,13 @@ int main() { // The parameter pack overload has to agree element for element. Q.memset(Memory, 0, N * sizeof(int)); - oneapiext::nd_launch(Q, Ndr, Kernel, - oneapiext::raw_kernel_arg{&Memory, sizeof(Memory)}, - oneapiext::raw_kernel_arg{&A, sizeof(A)}, - oneapiext::raw_kernel_arg{&B, sizeof(B)}, - oneapiext::raw_kernel_arg{&C, sizeof(C)}, - oneapiext::raw_kernel_arg{&D, sizeof(D)}); + oneapiext::nd_launch( + Q, Ndr, Kernel, + oneapiext::raw_kernel_arg{&Memory, oneapiext::pointer_arg}, + oneapiext::raw_kernel_arg{&A, sizeof(A)}, + oneapiext::raw_kernel_arg{&B, sizeof(B)}, + oneapiext::raw_kernel_arg{&C, sizeof(C)}, + oneapiext::raw_kernel_arg{&D, sizeof(D)}); Q.wait(); for (size_t I = 0; I < N; ++I) Failed += Check(Memory, Sum, I, "parameter pack overload"); @@ -105,7 +108,8 @@ int main() { // A one element span is the boundary against the parameter pack overload, // which a single raw_kernel_arg selects instead. - std::vector OneArg{{&Memory, sizeof(Memory)}}; + std::vector OneArg{ + {&Memory, oneapiext::pointer_arg}}; Q.memset(Memory, 0, N * sizeof(int)); oneapiext::nd_launch(Q, Ndr, getKernel(Q), sycl::span{ diff --git a/sycl/test-e2e/Graph/Update/update_with_raw_kernel_arg_pointer.cpp b/sycl/test-e2e/Graph/Update/update_with_raw_kernel_arg_pointer.cpp new file mode 100644 index 0000000000000..e899c9d53c0fd --- /dev/null +++ b/sycl/test-e2e/Graph/Update/update_with_raw_kernel_arg_pointer.cpp @@ -0,0 +1,83 @@ +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out +// Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG +// RUN: %if level_zero %{%{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %} + +// REQUIRES: ocloc && level_zero + +// Tests updating a raw_kernel_arg that was built as a pointer argument, which +// has to rebind the pointer rather than the bytes one is made of. + +#include "../graph_common.hpp" + +auto constexpr CLSource = R"===( +__kernel void RawArgKernel(int scalar, __global int *out) { + size_t id = get_global_id(0); + out[id] = id + scalar; +} +)==="; + +int main() { + queue Queue{}; + + auto SourceKB = + sycl::ext::oneapi::experimental::create_kernel_bundle_from_source( + Queue.get_context(), + sycl::ext::oneapi::experimental::source_language::opencl, CLSource); + auto ExecKB = sycl::ext::oneapi::experimental::build(SourceKB); + + exp_ext::command_graph Graph{Queue}; + + const size_t N = 1024; + int32_t *PtrA = malloc_device(N, Queue); + int32_t *PtrB = malloc_device(N, Queue); + Queue.memset(PtrA, 0, N * sizeof(int32_t)); + Queue.memset(PtrB, 0, N * sizeof(int32_t)); + Queue.wait(); + + int32_t Scalar = 42; + exp_ext::raw_kernel_arg RawScalar(&Scalar, sizeof(int32_t)); + exp_ext::raw_kernel_arg RawPtrA(&PtrA, exp_ext::pointer_arg); + exp_ext::raw_kernel_arg RawPtrB(&PtrB, exp_ext::pointer_arg); + + exp_ext::dynamic_parameter ScalarParam(RawScalar); + exp_ext::dynamic_parameter PtrParam(RawPtrA); + + auto KernelNode = Graph.add([&](handler &cgh) { + cgh.set_arg(0, ScalarParam); + cgh.set_arg(1, PtrParam); + cgh.parallel_for(sycl::range<1>{N}, + ExecKB.ext_oneapi_get_kernel("RawArgKernel")); + }); + + auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{}); + + // PtrA is the one the pointer argument was built from. + Queue.ext_oneapi_graph(ExecGraph).wait(); + + std::vector HostDataA(N); + std::vector HostDataB(N); + Queue.copy(PtrA, HostDataA.data(), N); + Queue.copy(PtrB, HostDataB.data(), N); + Queue.wait(); + for (size_t i = 0; i < N; i++) { + assert(HostDataA[i] == static_cast(i + Scalar)); + assert(HostDataB[i] == 0); + } + + // Rebind the pointer argument to the other allocation. + PtrParam.update(RawPtrB); + ExecGraph.update(KernelNode); + Queue.ext_oneapi_graph(ExecGraph).wait(); + + Queue.copy(PtrA, HostDataA.data(), N).wait(); + Queue.copy(PtrB, HostDataB.data(), N).wait(); + for (size_t i = 0; i < N; i++) { + assert(HostDataA[i] == static_cast(i + Scalar)); + assert(HostDataB[i] == static_cast(i + Scalar)); + } + + sycl::free(PtrA, Queue); + sycl::free(PtrB, Queue); + return 0; +} diff --git a/sycl/test-e2e/RawKernelArg/pointer_arg.cpp b/sycl/test-e2e/RawKernelArg/pointer_arg.cpp new file mode 100644 index 0000000000000..e0e894e69ac92 --- /dev/null +++ b/sycl/test-e2e/RawKernelArg/pointer_arg.cpp @@ -0,0 +1,71 @@ +// REQUIRES: aspect-usm_shared_allocations +// REQUIRES: ocloc && (opencl || level_zero) + +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +// Tests the pointer form of raw_kernel_arg, which says that an argument is a +// pointer instead of leaving the runtime to bind the bytes it is made of. That +// is what makes a pointer argument reach the kernel on a backend which takes a +// pointer through a different entry point than a value: OpenCL passes a value +// argument to clSetKernelArg, which rejects a USM pointer, and Native CPU puts +// the address of its own copy of the bytes into the argument slot. Hence no +// Level Zero requirement here, unlike the tests of the byte form. + +#include +#include +#include +#include + +#include + +namespace oneapiext = sycl::ext::oneapi::experimental; + +auto constexpr CLSource = R"===( +__kernel void WriteScalar(int in, __global int *out) { + out[get_global_id(0)] = in; +} +)==="; + +constexpr size_t N = 8; + +int main() { + sycl::queue Q; + + auto SourceKB = oneapiext::create_kernel_bundle_from_source( + Q.get_context(), oneapiext::source_language::opencl, CLSource); + auto ExecKB = oneapiext::build(SourceKB); + sycl::kernel Kernel = ExecKB.ext_oneapi_get_kernel("WriteScalar"); + + int *Out = sycl::malloc_shared(N, Q); + int In = 42; + + // Both arguments are raw, the pointer as a pointer and the scalar as bytes. + // The queue is out of order, so the sentinel has to be waited for rather than + // left to race with the kernel that overwrites it. + Q.memset(Out, 0xFF, N * sizeof(int)).wait(); + Q.submit([&](sycl::handler &CGH) { + CGH.set_arg(0, oneapiext::raw_kernel_arg{&In, sizeof(In)}); + CGH.set_arg(1, oneapiext::raw_kernel_arg{&Out, oneapiext::pointer_arg}); + CGH.parallel_for(sycl::range<1>{N}, Kernel); + }).wait(); + for (size_t I = 0; I < N; ++I) + assert(Out[I] == In); + + // The same through set_args, and with the pointer bound to a different + // allocation, so that a stale argument would show up. + int *Other = sycl::malloc_shared(N, Q); + In = 7; + Q.memset(Other, 0xFF, N * sizeof(int)).wait(); + Q.submit([&](sycl::handler &CGH) { + CGH.set_args(oneapiext::raw_kernel_arg{&In, sizeof(In)}, + oneapiext::raw_kernel_arg{&Other, oneapiext::pointer_arg}); + CGH.parallel_for(sycl::range<1>{N}, Kernel); + }).wait(); + for (size_t I = 0; I < N; ++I) + assert(Other[I] == In); + + sycl::free(Other, Q); + sycl::free(Out, Q); + return 0; +} diff --git a/sycl/test/abi/layout_raw_kernel_arg.cpp b/sycl/test/abi/layout_raw_kernel_arg.cpp new file mode 100644 index 0000000000000..0b34c35316a98 --- /dev/null +++ b/sycl/test/abi/layout_raw_kernel_arg.cpp @@ -0,0 +1,21 @@ +// 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 // for SYCL_EXTERNAL +#include + +// A raw_kernel_arg crosses the ABI boundary as an element of the span that the +// nd_launch overloads take, and the graph extension copies one as bytes, so its +// layout is fixed here. MIsPointer comes last, so that the two members the +// library read before it existed keep their offsets. +SYCL_EXTERNAL void takeRawKernelArg(sycl::ext::oneapi::experimental::raw_kernel_arg) {} +// CHECK: 0 | class sycl::ext::oneapi::experimental::raw_kernel_arg +// CHECK-NEXT: 0 | const void * MArgData +// CHECK-NEXT: 8 | size_t MArgSize +// CHECK-NEXT: 16 | _Bool MIsPointer +// CHECK-NEXT: | [sizeof=24, dsize=17, align=8, +// CHECK-NEXT: | nvsize=17, nvalign=8] diff --git a/sycl/test/extensions/raw_kernel_arg/pointer_arg_errors.cpp b/sycl/test/extensions/raw_kernel_arg/pointer_arg_errors.cpp new file mode 100644 index 0000000000000..eaad709fe3111 --- /dev/null +++ b/sycl/test/extensions/raw_kernel_arg/pointer_arg_errors.cpp @@ -0,0 +1,37 @@ +// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s + +// The pointer form of raw_kernel_arg takes the address of the pointer, in +// keeping with the byte form taking the address of the bytes. Passing the +// pointer itself would name an address the runtime would then read a pointer +// from, which is why it does not compile. + +#include + +#include + +namespace oneapiext = sycl::ext::oneapi::experimental; + +// The graph extension copies a raw_kernel_arg as bytes. +static_assert(std::is_trivially_copyable_v); + +void pointer_form(int *Ptr, const float *ConstPtr, void *VoidPtr) { + // Every pointer type binds through the address form without a cast. + oneapiext::raw_kernel_arg Typed{&Ptr, oneapiext::pointer_arg}; + oneapiext::raw_kernel_arg Const{&ConstPtr, oneapiext::pointer_arg}; + oneapiext::raw_kernel_arg Void{&VoidPtr, oneapiext::pointer_arg}; + + // The byte form is unchanged, including for the bytes of a pointer, which + // only bind as a pointer on the Level Zero backend. + oneapiext::raw_kernel_arg Bytes{&Ptr, sizeof(Ptr)}; + + // expected-error@+1 {{no matching constructor for initialization of 'oneapiext::raw_kernel_arg'}} + oneapiext::raw_kernel_arg PointerItself{Ptr, oneapiext::pointer_arg}; + + // expected-error@+1 {{no matching constructor for initialization of 'oneapiext::raw_kernel_arg'}} + oneapiext::raw_kernel_arg NotAPointer{42, oneapiext::pointer_arg}; + + (void)Typed; + (void)Const; + (void)Void; + (void)Bytes; +} From 62fa8b3199c124cddecba8df95b06494473f8a12 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 19 Aug 2026 15:37:56 +0000 Subject: [PATCH 5/5] [SYCL] Pass an argument list given as a container to the span overload The overloads taking the kernel arguments as a sequence take a sycl::span, which a std::vector, a std::array or a std::span of raw_kernel_arg all convert to, but a parameter pack is an exact match and wins overload resolution, so a container reached the parameter pack overload instead. On the queue that was diagnosed. On the handler it was not: a container that is trivially copyable, std::array among them, compiled and bound the container object as a single kernel argument, and one that is not produced an error from inside handler.hpp. Have both parameter pack overloads pass such an argument on to the overload taking a span, so that passing the container itself binds the arguments it holds. The specification loses the clause that made such a call ill-formed, along with a note claiming that sycl::span has no implicit conversion from a container, which it does have. Spell out sycl::span rather than span in the specification, since std::span is a different type, and replace the test of the diagnostic with one that checks that every spelling of an argument list compiles. --- ...sycl_ext_oneapi_enqueue_functions.asciidoc | 28 +++++---- .../oneapi/experimental/enqueue_functions.hpp | 53 +++++++++++------ .../nd_launch_kernel_obj_span.cpp | 24 ++++++++ .../enqueue_functions/kernel_arg_span.cpp | 57 +++++++++++++++++++ .../kernel_arg_span_errors.cpp | 33 ----------- 5 files changed, 133 insertions(+), 62 deletions(-) create mode 100644 sycl/test/extensions/enqueue_functions/kernel_arg_span.cpp delete mode 100644 sycl/test/extensions/enqueue_functions/kernel_arg_span_errors.cpp diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc index 1c16e1f5e0dfd..c03fbc7b814bd 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc @@ -41,7 +41,8 @@ This extension is written against the SYCL 2020 revision 7 specification. All references below to the "core SYCL specification" or to section numbers in the SYCL specification refer to that revision. -The `nd_launch` overloads that take the kernel arguments as a `span` depend on +The `nd_launch` overloads that take the kernel arguments as a `sycl::span` +depend on version 2 of link:../experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc[sycl_ext_oneapi_raw_kernel_arg], which is what lets an element of that sequence represent a pointer argument. @@ -174,7 +175,9 @@ implementation supports. |2 |Adds the `nd_launch` overloads that take the arguments of a `sycl::kernel` as - a `span` of `raw_kernel_arg`. + a `sycl::span` of `raw_kernel_arg`, and specifies that a parameter pack + overload handed such a sequence as its only argument passes it on as the + argument list. |=== @@ -535,11 +538,11 @@ namespace sycl::ext::oneapi::experimental { template void nd_launch(sycl::queue q, sycl::nd_range r, - const sycl::kernel& k, span args); + const sycl::kernel& k, sycl::span args); template void nd_launch(sycl::handler &h, sycl::nd_range r, - const sycl::kernel& k, span args); + const sycl::kernel& k, sycl::span args); } ---- @@ -550,16 +553,19 @@ as a basic kernel, using the number of work-items specified by a at index `i`, as if it had been passed to `set_arg`. The sequence referenced by `args` need only remain valid until the function returns. -An application that calls one of the parameter pack overloads above with a -single argument whose type is convertible to `span` is -ill-formed. +If one of the parameter pack overloads above is called with a single argument +whose type is convertible to `sycl::span`, that argument +is the kernel argument list and the effects are those of the overload taking a +`sycl::span`. [_Note:_ These overloads exist for an application whose kernel argument list is only known at run time, which would otherwise have to instantiate a parameter -pack overload for each number of arguments it may encounter. Since `span` has -no implicit conversion from a container, an application holding its arguments -in a `std::vector` passes `{args.data(), args.size()}`. Passing -a single `raw_kernel_arg` still selects the parameter pack overload. _end note_] +pack overload for each number of arguments it may encounter. An application +holding its arguments in a `std::vector`, in a `std::array` of +them or in a `std::span` of them can therefore pass the container itself, since +all of those convert to a `sycl::span`. A single `raw_kernel_arg` still selects +the parameter pack overload, as it is one argument rather than a sequence of +them. _end note_] [_Note:_ An element that represents a pointer argument has to be constructed with the `pointer_arg` overload of `raw_kernel_arg`, as described in diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp index 392924cc242e9..c3520ef60795b 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -136,6 +136,15 @@ sycl::detail::KernelArgView makeKernelArgView(const T &Arg) { : kernel_param_kind_t::kind_std_layout}; } +// True when a parameter pack overload was handed the argument list itself, as a +// container that converts to the span the sibling overload takes. A pack is an +// exact match and wins overload resolution, so such a call has to be forwarded +// rather than bound as one argument. +template +inline constexpr bool is_arg_list_container_v = + sizeof...(ArgsT) == 1 && + (std::is_convertible_v> && ...); + template void submit_impl(const queue &Q, PropertiesT Props, CommandGroupFunc &&CGF, const sycl::detail::code_location &CodeLoc) { @@ -436,28 +445,37 @@ void nd_launch(queue Q, launch_config, Properties> Config, } } +template +void nd_launch(handler &CGH, nd_range Range, + const kernel &KernelObj, sycl::span Args); + template void nd_launch(handler &CGH, nd_range Range, const kernel &KernelObj, ArgsT &&...Args) { - CGH.set_args(std::forward(Args)...); - CGH.parallel_for(Range, KernelObj); + if constexpr (detail::is_arg_list_container_v) { + nd_launch(CGH, Range, KernelObj, + sycl::span{std::forward(Args)...}); + } else { + CGH.set_args(std::forward(Args)...); + CGH.parallel_for(Range, KernelObj); + } } +template +void nd_launch(queue Q, nd_range Range, const kernel &KernelObj, + sycl::span Args, + const sycl::detail::code_location &CodeLoc = + sycl::detail::code_location::current()); + template void nd_launch(queue Q, nd_range Range, const kernel &KernelObj, ArgsT &&...Args) { - // A container of raw_kernel_arg converts to the span the sibling overload - // takes, but a pack is an exact match and wins overload resolution, which - // would bind the container object itself as one argument. Diagnose it here, - // so that no other branch is instantiated for such a call. - constexpr bool ArgListPassedAsContainer = - sizeof...(ArgsT) == 1 && - (std::is_convertible_v> && ...); - if constexpr (ArgListPassedAsContainer) { - static_assert(!ArgListPassedAsContainer, - "The kernel argument list must be passed as " - "sycl::span, e.g. {Args.data(), " - "Args.size()}"); + // A container of raw_kernel_arg is the argument list, not one argument, and + // the pack is what overload resolution picks for it, so hand it over to the + // overload that takes a span. + if constexpr (detail::is_arg_list_container_v) { + nd_launch(std::move(Q), Range, KernelObj, + sycl::span{std::forward(Args)...}); } else 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 @@ -484,7 +502,7 @@ void nd_launch(queue Q, nd_range Range, const kernel &KernelObj, template void nd_launch(handler &CGH, nd_range Range, - const kernel &KernelObj, span Args) { + const kernel &KernelObj, sycl::span Args) { // set_arg only takes an rvalue raw_kernel_arg; an lvalue would select the // generic overload and bind the object itself as the argument. for (size_t I = 0; I < Args.size(); ++I) @@ -497,9 +515,8 @@ void nd_launch(handler &CGH, nd_range Range, // otherwise need one instantiation of the pack overload per argument count. template void nd_launch(queue Q, nd_range Range, const kernel &KernelObj, - span Args, - const sycl::detail::code_location &CodeLoc = - sycl::detail::code_location::current()) { + sycl::span Args, + const sycl::detail::code_location &CodeLoc) { sycl::detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); sycl::submit_kernel_obj_direct_without_event_impl( Q, sycl::detail::nd_range_view(Range), KernelObj, Args, diff --git a/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_span.cpp b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_span.cpp index 844fc3ba54e00..e69975b2af3b5 100644 --- a/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_span.cpp +++ b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_span.cpp @@ -106,6 +106,30 @@ int main() { for (size_t I = 0; I < N; ++I) Failed += Check(Memory, Sum, I, "handler form of the span overload"); + // The container holding the arguments converts to that span, so passing it + // has to bind the arguments it holds rather than the container object. + Q.memset(Memory, 0, N * sizeof(int)); + oneapiext::nd_launch(Q, Ndr, Kernel, Args); + Q.wait(); + for (size_t I = 0; I < N; ++I) + Failed += Check(Memory, Sum, I, "argument list passed as a container"); + + Q.memset(Memory, 0, N * sizeof(int)); + Q.submit([&](sycl::handler &CGH) { + oneapiext::nd_launch(CGH, Ndr, Kernel, Args); + }).wait(); + for (size_t I = 0; I < N; ++I) + Failed += Check(Memory, Sum, I, "container through the handler form"); + + // A mutable span of the same sequence names the same argument list. + Q.memset(Memory, 0, N * sizeof(int)); + oneapiext::nd_launch( + Q, Ndr, Kernel, + sycl::span{Args.data(), Args.size()}); + Q.wait(); + for (size_t I = 0; I < N; ++I) + Failed += Check(Memory, Sum, I, "argument list as a mutable span"); + // A one element span is the boundary against the parameter pack overload, // which a single raw_kernel_arg selects instead. std::vector OneArg{ diff --git a/sycl/test/extensions/enqueue_functions/kernel_arg_span.cpp b/sycl/test/extensions/enqueue_functions/kernel_arg_span.cpp new file mode 100644 index 0000000000000..208b73f310ead --- /dev/null +++ b/sycl/test/extensions/enqueue_functions/kernel_arg_span.cpp @@ -0,0 +1,57 @@ +// RUN: %clangxx -fsycl -fsyntax-only %s +// RUN: %clangxx -fsycl -fsyntax-only -std=c++20 %s + +// An argument list held in a container converts to the sycl::span that the +// nd_launch span overloads take, but a parameter pack is an exact match and +// wins overload resolution. Without the forwarding the pack overloads do, the +// container object itself would be bound as a single kernel argument, which +// compiles for any trivially copyable container and produces wrong results at +// run time. Check that every spelling of an argument list is accepted, and that +// a single raw_kernel_arg is still one argument. + +#include + +#include +#include +#if __cpp_lib_span +#include +#endif + +namespace oneapiext = sycl::ext::oneapi::experimental; + +void argument_list_spellings(sycl::queue Q, sycl::handler &CGH, + sycl::nd_range<1> Range, + const sycl::kernel &Kernel) { + int Value = 1; + std::vector Vector{{&Value, sizeof(Value)}}; + std::array Array{ + oneapiext::raw_kernel_arg{&Value, sizeof(Value)}}; + sycl::span Span{Vector.data(), + Vector.size()}; + sycl::span MutableSpan{Vector.data(), + Vector.size()}; + + oneapiext::nd_launch(Q, Range, Kernel, Vector); + oneapiext::nd_launch(Q, Range, Kernel, Array); + oneapiext::nd_launch(Q, Range, Kernel, Span); + oneapiext::nd_launch(Q, Range, Kernel, MutableSpan); + oneapiext::nd_launch(CGH, Range, Kernel, Vector); + oneapiext::nd_launch(CGH, Range, Kernel, Array); + oneapiext::nd_launch(CGH, Range, Kernel, Span); + oneapiext::nd_launch(CGH, Range, Kernel, MutableSpan); +#if __cpp_lib_span + std::span StdSpan{Vector.data(), + Vector.size()}; + oneapiext::nd_launch(Q, Range, Kernel, StdSpan); + oneapiext::nd_launch(CGH, Range, Kernel, StdSpan); +#endif + + // One raw_kernel_arg is one argument, and typed arguments are unaffected. + oneapiext::nd_launch(Q, Range, Kernel, + oneapiext::raw_kernel_arg{&Value, sizeof(Value)}); + oneapiext::nd_launch(CGH, Range, Kernel, + oneapiext::raw_kernel_arg{&Value, sizeof(Value)}); + int *Pointer = nullptr; + oneapiext::nd_launch(Q, Range, Kernel, Pointer, Value); + oneapiext::nd_launch(CGH, Range, Kernel, Pointer, Value); +} diff --git a/sycl/test/extensions/enqueue_functions/kernel_arg_span_errors.cpp b/sycl/test/extensions/enqueue_functions/kernel_arg_span_errors.cpp deleted file mode 100644 index f5f0dfb8d8305..0000000000000 --- a/sycl/test/extensions/enqueue_functions/kernel_arg_span_errors.cpp +++ /dev/null @@ -1,33 +0,0 @@ -// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s - -// An argument list held in a container converts to the span that the nd_launch -// span overload takes, but a parameter pack is an exact match and wins overload -// resolution. Without a diagnostic the container object itself would be bound -// as a single kernel argument, which compiles for any trivially copyable -// wrapper and produces wrong results at run time. - -#include - -#include - -namespace oneapiext = sycl::ext::oneapi::experimental; - -void argument_list_must_be_a_span(sycl::queue Q, sycl::nd_range<1> Range, - const sycl::kernel &Kernel) { - int Value = 1; - std::vector Args{{&Value, sizeof(Value)}}; - - // expected-error@sycl/ext/oneapi/experimental/enqueue_functions.hpp:* {{The kernel argument list must be passed as sycl::span}} - oneapiext::nd_launch(Q, Range, Kernel, Args); - - sycl::span Mutable{Args.data(), Args.size()}; - // expected-error@sycl/ext/oneapi/experimental/enqueue_functions.hpp:* {{The kernel argument list must be passed as sycl::span}} - oneapiext::nd_launch(Q, Range, Kernel, Mutable); - - // The spelling the diagnostic asks for, and a single argument that happens to - // be a raw_kernel_arg, both have to keep working. - sycl::span AsSpan{Args.data(), Args.size()}; - oneapiext::nd_launch(Q, Range, Kernel, AsSpan); - oneapiext::nd_launch(Q, Range, Kernel, - oneapiext::raw_kernel_arg{&Value, sizeof(Value)}); -}