From 3f953249a9cac4f94177f2417697103e0d7100a1 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 12 Aug 2026 12:49:04 +0000 Subject: [PATCH 01/24] [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 | 46 +++++- 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, 336 insertions(+), 15 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 4652061a55052..93f19737987e6 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -450,11 +450,23 @@ void nd_launch(handler &CGH, nd_range Range, template void nd_launch(queue Q, nd_range Range, const kernel &KernelObj, ArgsT &&...Args) { - // The handler-less path only takes arguments that can be bound directly, - // anything else goes through the handler overload above. - if constexpr ((detail::is_direct_kernel_arg_v && ...)) { - // The array is one element longer than the pack so that a zero-argument - // kernel stays well formed. + // 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_direct_kernel_arg_v && ...)) { + // The handler-less path only takes arguments that can be bound directly, + // anything else goes through the handler overload above. 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)...}; sycl::detail::tls_code_loc_t TlsCodeLocCapture{ @@ -470,6 +482,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 b47c7f3dd0c78..0b2f5f6cf6f6b 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,16 +940,17 @@ 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) { + const sycl::detail::KernelArgView View = makeKernelArgView(Args[I]); // `ArgDesc` holds a `void *` because the kinds that carry an object // rather than bytes hand it out as a mutable pointer. These arguments are // bytes and this path only ever reads them, hence the cast. - void *Value = const_cast(Args[I].MPtr); + 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)); } @@ -974,6 +992,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 a62cb8a99dfc8..b2126a43a06fa 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 7b1344ed5de48..d2c1771696bfa 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 10933bc3c8622..d350b33b5f07e 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -4393,6 +4393,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 293786f2a463d909db21942e5d6eda6e8b34c0f2 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 19 Aug 2026 15:09:56 +0000 Subject: [PATCH 02/24] [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 93f19737987e6..20c06ce9ecf97 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -135,7 +135,9 @@ sycl::detail::KernelArgView makeKernelArgView(const T &Arg) { using ArgT = unqualified_arg_t; if constexpr (std::is_same_v) return {RawKernelArgAccess::getData(Arg), RawKernelArgAccess::getSize(Arg), - kernel_param_kind_t::kind_std_layout}; + RawKernelArgAccess::isPointer(Arg) + ? kernel_param_kind_t::kind_pointer + : kernel_param_kind_t::kind_std_layout}; else return {&Arg, sizeof(ArgT), kernel_arg_kind_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 61fa55d47f1de..004c0cf77f6d2 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 0b2f5f6cf6f6b..fd14b3cb6543a 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 b2126a43a06fa..4a9c995b10005 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 8bcaf7d67390bebba399f831c780897f7e730ea3 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 19 Aug 2026 15:37:56 +0000 Subject: [PATCH 03/24] [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 20c06ce9ecf97..5ba91ed38569b 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -142,6 +142,15 @@ sycl::detail::KernelArgView makeKernelArgView(const T &Arg) { return {&Arg, sizeof(ArgT), kernel_arg_kind_v}; } +// 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) { @@ -442,28 +451,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_direct_kernel_arg_v && ...)) { // The handler-less path only takes arguments that can be bound directly, // anything else goes through the handler overload above. The array is one @@ -486,7 +504,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) @@ -499,9 +517,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)}); -} From 39ac830e07c1a10c65762b7ff32d96bb9d38150e Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 2 Sep 2026 08:33:48 +0000 Subject: [PATCH 04/24] [SYCL] Replace the argument sequence note with an Examples section The note carried two things a specification does not need in that place: why the overloads taking a sequence of arguments exist, and how a container of raw_kernel_arg reaches them. The first is motivation and moves to the Overview, the second is usage and moves to an Examples section at the end of the specification, which is where this extension shows how its APIs are called. --- ...sycl_ext_oneapi_enqueue_functions.asciidoc | 34 ++++++++++++++----- 1 file changed, 25 insertions(+), 9 deletions(-) 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 c03fbc7b814bd..2cafc4739c020 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc @@ -93,6 +93,11 @@ existing queue shortcuts. A future version of this extension may adjust this overload set to include functions compatible with future C++ concepts (e.g, by accepting a scheduler and returning a sender). +The arguments of a `sycl::kernel` can be passed either as a parameter pack or as +a sequence, for an application that only learns its argument list at run time and +would otherwise have to instantiate a parameter pack overload for each number of +arguments it may encounter. + === Usage example @@ -558,15 +563,6 @@ 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. 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 link:../experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc[sycl_ext_oneapi_raw_kernel_arg], @@ -827,6 +823,26 @@ _Effects_: Submits an executable command graph to the `sycl::queue` or `sycl::ha |==== +== Examples + +An argument list that is only known at run time is held in a container and passed +as the sequence of arguments: a `std::vector`, a `std::array` of +them and a `std::span` of them all convert to `sycl::span`. + +[source,c++] +---- +namespace syclex = sycl::ext::oneapi::experimental; + +int count = 1024; +float *data = sycl::malloc_device(count, q); + +std::vector args; +args.emplace_back(&data, syclex::pointer_arg); +args.emplace_back(&count, sizeof(count)); + +syclex::nd_launch(q, sycl::nd_range<1>{count, 256}, k, args); +---- + == Issues . What should `submit_with_event` be called? From 9986af5f047e1436b760be5c9468dd4b2be8ca2d Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 2 Sep 2026 12:48:05 +0000 Subject: [PATCH 05/24] [SYCL] Take the kernel argument sequence as a std::span The overloads that take the arguments of a sycl::kernel as a sequence now take a std::span, which makes them available only with C++20. That is the direction the other extensions take: an implementation cannot turn sycl::span into an alias of std::span without breaking the ABI, so a new API should not add another place that would have to change if SYCL ever switches. The exported entry points take a pointer and a count instead of a view, so the spelling the public overloads use is not part of the ABI either and can be changed again without breaking it. Before C++20 there is no overload taking a sequence, so a container of raw_kernel_arg handed to a parameter pack overload would be bound as a single kernel argument, which compiles for any trivially copyable container and gives wrong results. It is diagnosed instead, with a test that pins both the diagnostic and that the same spellings compile with C++20. --- ...sycl_ext_oneapi_enqueue_functions.asciidoc | 28 ++++--- .../oneapi/experimental/enqueue_functions.hpp | 75 ++++++++++++++++--- sycl/include/sycl/queue.hpp | 8 +- sycl/source/queue.cpp | 10 ++- .../nd_launch_kernel_obj_span.cpp | 20 +++-- sycl/test/abi/sycl_symbols_linux.dump | 4 +- sycl/test/abi/sycl_symbols_windows.dump | 4 +- .../kernel_arg_sequence_cxx17.cpp | 47 ++++++++++++ .../enqueue_functions/kernel_arg_span.cpp | 25 ++++--- 9 files changed, 170 insertions(+), 51 deletions(-) create mode 100644 sycl/test/extensions/enqueue_functions/kernel_arg_sequence_cxx17.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 2cafc4739c020..a4c0d33325fe1 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc @@ -41,11 +41,13 @@ 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 `sycl::span` +The `nd_launch` overloads that take the kernel arguments as a `std::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. +Those overloads use `std::span`, so they are available only when the application +is compiled with C++20 or later. == Status @@ -96,7 +98,8 @@ by accepting a scheduler and returning a sender). The arguments of a `sycl::kernel` can be passed either as a parameter pack or as a sequence, for an application that only learns its argument list at run time and would otherwise have to instantiate a parameter pack overload for each number of -arguments it may encounter. +arguments it may encounter. The sequence is a `std::span`, so those overloads +require C++20. === Usage example @@ -180,9 +183,9 @@ implementation supports. |2 |Adds the `nd_launch` overloads that take the arguments of a `sycl::kernel` as - 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. + a `std::span` of `raw_kernel_arg`, which require C++20, and specifies that a + parameter pack overload handed such a sequence as its only argument passes it + on as the argument list. |=== @@ -541,13 +544,15 @@ a! ---- namespace sycl::ext::oneapi::experimental { +// Requires C++20 template void nd_launch(sycl::queue q, sycl::nd_range r, - const sycl::kernel& k, sycl::span args); + const sycl::kernel& k, std::span args); +// Requires C++20 template void nd_launch(sycl::handler &h, sycl::nd_range r, - const sycl::kernel& k, sycl::span args); + const sycl::kernel& k, std::span args); } ---- @@ -559,9 +564,9 @@ 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. 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`. +whose type is convertible to `std::span`, that argument is +the kernel argument list and the effects are those of the overload taking a +`std::span`. [_Note:_ An element that represents a pointer argument has to be constructed with the `pointer_arg` overload of `raw_kernel_arg`, as described in @@ -827,7 +832,8 @@ _Effects_: Submits an executable command graph to the `sycl::queue` or `sycl::ha An argument list that is only known at run time is held in a container and passed as the sequence of arguments: a `std::vector`, a `std::array` of -them and a `std::span` of them all convert to `sycl::span`. +them and a `std::span` of them all convert to `std::span`. +These overloads require C++20. [source,c++] ---- diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp index 5ba91ed38569b..95b3c3b7e3a23 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -10,6 +10,10 @@ #include +#if __has_include() +#include +#endif + #include #include #include @@ -145,11 +149,39 @@ sycl::detail::KernelArgView makeKernelArgView(const T &Arg) { // 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. +// rather than bound as one argument. Without `std::span` there is no overload +// taking a sequence, so no call can be one. +#if __cpp_lib_span template inline constexpr bool is_arg_list_container_v = sizeof...(ArgsT) == 1 && - (std::is_convertible_v> && ...); + (std::is_convertible_v> && ...); +#else +template +inline constexpr bool is_arg_list_container_v = false; +#endif + +#if !__cpp_lib_span +// A contiguous sequence of `raw_kernel_arg`, recognized without `std::span` so +// that a caller passing its argument list as a container before C++20 is told +// so, rather than having the container bound as a single kernel argument, which +// compiles for any trivially copyable one. +template +inline constexpr bool is_raw_kernel_arg_sequence_v = + std::is_array_v && + std::is_same_v, raw_kernel_arg>; +template +inline constexpr bool is_raw_kernel_arg_sequence_v< + T, std::void_t().size()), + decltype(std::declval().data())>> = + std::is_convertible_v().data()), + const raw_kernel_arg *>; + +template +inline constexpr bool is_arg_list_sequence_v = + sizeof...(ArgsT) == 1 && + (is_raw_kernel_arg_sequence_v> && ...); +#endif template void submit_impl(const queue &Q, PropertiesT Props, CommandGroupFunc &&CGF, @@ -451,27 +483,42 @@ void nd_launch(queue Q, launch_config, Properties> Config, } } +#if __cpp_lib_span template void nd_launch(handler &CGH, nd_range Range, - const kernel &KernelObj, sycl::span Args); + const kernel &KernelObj, std::span Args); +#endif template void nd_launch(handler &CGH, nd_range Range, const kernel &KernelObj, ArgsT &&...Args) { +#if !__cpp_lib_span + // Reached for a container both directly and through the queue overload, which + // falls back to a command group for anything that is not a direct argument. + static_assert(!detail::is_arg_list_sequence_v, + "Passing the arguments of a sycl::kernel as a sequence requires " + "C++20, where the nd_launch overload taking a std::span of " + "raw_kernel_arg is available. Compile with C++20 or pass the " + "arguments as a parameter pack."); +#endif if constexpr (detail::is_arg_list_container_v) { +#if __cpp_lib_span nd_launch(CGH, Range, KernelObj, - sycl::span{std::forward(Args)...}); + std::span{std::forward(Args)...}); +#endif } else { CGH.set_args(std::forward(Args)...); CGH.parallel_for(Range, KernelObj); } } +#if __cpp_lib_span template void nd_launch(queue Q, nd_range Range, const kernel &KernelObj, - sycl::span Args, + std::span Args, const sycl::detail::code_location &CodeLoc = sycl::detail::code_location::current()); +#endif template void nd_launch(queue Q, nd_range Range, const kernel &KernelObj, @@ -480,8 +527,10 @@ void nd_launch(queue Q, nd_range Range, const kernel &KernelObj, // 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) { +#if __cpp_lib_span nd_launch(std::move(Q), Range, KernelObj, - sycl::span{std::forward(Args)...}); + std::span{std::forward(Args)...}); +#endif } else if constexpr ((detail::is_direct_kernel_arg_v && ...)) { // The handler-less path only takes arguments that can be bound directly, // anything else goes through the handler overload above. The array is one @@ -492,8 +541,8 @@ void nd_launch(queue Q, nd_range Range, const kernel &KernelObj, 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(), + Q, sycl::detail::nd_range_view(Range), KernelObj, ArgViews, + sizeof...(ArgsT), TlsCodeLocCapture.query(), TlsCodeLocCapture.isToplevel()); } else { submit(std::move(Q), [&](handler &CGH) { @@ -502,9 +551,10 @@ void nd_launch(queue Q, nd_range Range, const kernel &KernelObj, } } +#if __cpp_lib_span template void nd_launch(handler &CGH, nd_range Range, - const kernel &KernelObj, sycl::span Args) { + const kernel &KernelObj, std::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) @@ -517,13 +567,14 @@ 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, - sycl::span Args, + std::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, - TlsCodeLocCapture.query(), TlsCodeLocCapture.isToplevel()); + Q, sycl::detail::nd_range_view(Range), KernelObj, Args.data(), + Args.size(), TlsCodeLocCapture.query(), TlsCodeLocCapture.isToplevel()); } +#endif // __cpp_lib_span template void nd_launch(handler &CGH, diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp index 5dff032bbdcb8..90c3ef20c1c50 100644 --- a/sycl/include/sycl/queue.hpp +++ b/sycl/include/sycl/queue.hpp @@ -73,10 +73,12 @@ class raw_kernel_arg; // 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. +// submit_kernel_direct_* does for kernel function objects. The arguments are +// passed as a pointer and a count rather than as a view type, so that the +// spelling the public overloads use is not part of the ABI. 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 kernel &Kernel, const detail::KernelArgView *Args, size_t NumArgs, const detail::code_location &CodeLoc, bool IsTopCodeLoc); // As above, for an argument list that is already a contiguous sequence of @@ -85,7 +87,7 @@ void __SYCL_EXPORT submit_kernel_obj_direct_without_event_impl( 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 ext::oneapi::experimental::raw_kernel_arg *Args, size_t NumArgs, const detail::code_location &CodeLoc, bool IsTopCodeLoc); event __SYCL_EXPORT submit_kernel_direct_with_event_impl( diff --git a/sycl/source/queue.cpp b/sycl/source/queue.cpp index 2e1bccd24cb45..74f76d97a460f 100644 --- a/sycl/source/queue.cpp +++ b/sycl/source/queue.cpp @@ -321,19 +321,21 @@ void submit_kernel_direct_without_event_impl( void submit_kernel_obj_direct_without_event_impl( const queue &Queue, const detail::nd_range_view &RangeView, - const kernel &Kernel, sycl::span Args, + const kernel &Kernel, const detail::KernelArgView *Args, size_t NumArgs, const detail::code_location &CodeLoc, bool IsTopCodeLoc) { detail::getSyclObjImpl(Queue)->submit_kernel_obj_direct_without_event( - RangeView, detail::getSyclObjImpl(Kernel), Args, CodeLoc, IsTopCodeLoc); + RangeView, detail::getSyclObjImpl(Kernel), {Args, NumArgs}, 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 ext::oneapi::experimental::raw_kernel_arg *Args, size_t NumArgs, const detail::code_location &CodeLoc, bool IsTopCodeLoc) { detail::getSyclObjImpl(Queue)->submit_kernel_obj_direct_without_event( - RangeView, detail::getSyclObjImpl(Kernel), Args, CodeLoc, IsTopCodeLoc); + RangeView, detail::getSyclObjImpl(Kernel), {Args, NumArgs}, CodeLoc, + IsTopCodeLoc); } event submit_graph_direct_with_event_impl( 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 e69975b2af3b5..2aa044929dc20 100644 --- a/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_span.cpp +++ b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_span.cpp @@ -1,11 +1,17 @@ // REQUIRES: aspect-usm_shared_allocations -// RUN: %{build} -o %t.out + +// DEFINE: %{cpp20} = %if cl_options %{/clang:-std=c++20%} %else %{-std=c++20%} + +// RUN: %{build} %{cpp20} -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. +// std::span of raw_kernel_arg, i.e. an argument list whose length is only known +// at run time. Those overloads take a std::span, hence C++20. 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 @@ -71,7 +77,7 @@ int main() { 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()}; + std::span ArgSpan{Args.data(), Args.size()}; int Failed = 0; @@ -125,7 +131,7 @@ int main() { Q.memset(Memory, 0, N * sizeof(int)); oneapiext::nd_launch( Q, Ndr, Kernel, - sycl::span{Args.data(), Args.size()}); + std::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"); @@ -136,7 +142,7 @@ int main() { {&Memory, oneapiext::pointer_arg}}; Q.memset(Memory, 0, N * sizeof(int)); oneapiext::nd_launch(Q, Ndr, getKernel(Q), - sycl::span{ + std::span{ OneArg.data(), OneArg.size()}); Q.wait(); for (size_t I = 0; I < N; ++I) diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index d2c1771696bfa..b2bce267308e5 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3163,8 +3163,8 @@ _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_V143submit_kernel_obj_direct_without_event_implERKNS0_5queueERKNS0_6detail16nd_range_view_v113nd_range_viewERKNS0_6kernelEPKNS0_3ext6oneapi12experimental14raw_kernel_argEmRKNS4_13code_locationEb +_ZN4sycl3_V143submit_kernel_obj_direct_without_event_implERKNS0_5queueERKNS0_6detail16nd_range_view_v113nd_range_viewERKNS0_6kernelEPKNS4_18kernel_arg_view_v113KernelArgViewEmRKNS4_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 d350b33b5f07e..1f3d4621c0de7 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -4392,8 +4392,8 @@ ?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_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_kernel_obj_direct_without_event_impl@_V1@sycl@@YAXAEBVqueue@12@AEBVnd_range_view@nd_range_view_v1@detail@12@AEBVkernel@12@PEBUKernelArgView@kernel_arg_view_v1@612@_KAEBUcode_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@PEBVraw_kernel_arg@experimental@oneapi@ext@12@_KAEBUcode_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_sequence_cxx17.cpp b/sycl/test/extensions/enqueue_functions/kernel_arg_sequence_cxx17.cpp new file mode 100644 index 0000000000000..8cf32286ac006 --- /dev/null +++ b/sycl/test/extensions/enqueue_functions/kernel_arg_sequence_cxx17.cpp @@ -0,0 +1,47 @@ +// RUN: not %clangxx -fsycl -fsyntax-only -std=c++17 -DSEQUENCE=Vector %s 2>&1 | FileCheck %s +// RUN: not %clangxx -fsycl -fsyntax-only -std=c++17 -DSEQUENCE=Array %s 2>&1 | FileCheck %s +// RUN: not %clangxx -fsycl -fsyntax-only -std=c++17 -DSEQUENCE=Span %s 2>&1 | FileCheck %s +// RUN: %clangxx -fsycl -fsyntax-only -std=c++20 -DSEQUENCE=Vector %s +// RUN: %clangxx -fsycl -fsyntax-only -std=c++20 -DSEQUENCE=Array %s +// RUN: %clangxx -fsycl -fsyntax-only -std=c++20 -DSEQUENCE=Span %s + +// The overloads that take the arguments of a sycl::kernel as a sequence take a +// std::span, so they do not exist before C++20. A caller that passes its +// argument list as a container there would otherwise have the container bound +// as a single kernel argument, which compiles for any trivially copyable one +// and produces wrong results at run time, so it is diagnosed instead. + +// CHECK: Passing the arguments of a sycl::kernel as a sequence requires C++20 + +#include + +#include +#include + +namespace oneapiext = sycl::ext::oneapi::experimental; + +void argument_list_as_a_container(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()}; + + oneapiext::nd_launch(Q, Range, Kernel, SEQUENCE); + oneapiext::nd_launch(CGH, Range, Kernel, SEQUENCE); +} + +void one_argument_at_a_time(sycl::queue Q, sycl::handler &CGH, + sycl::nd_range<1> Range, + const sycl::kernel &Kernel) { + // A single raw_kernel_arg and typed arguments are one argument each, so they + // are unaffected by the standard in use. + int Value = 1; + int *Pointer = nullptr; + oneapiext::nd_launch(Q, Range, Kernel, + oneapiext::raw_kernel_arg{&Value, sizeof(Value)}); + oneapiext::nd_launch(CGH, Range, Kernel, Pointer, Value); +} diff --git a/sycl/test/extensions/enqueue_functions/kernel_arg_span.cpp b/sycl/test/extensions/enqueue_functions/kernel_arg_span.cpp index 208b73f310ead..68af2e1fcfdb5 100644 --- a/sycl/test/extensions/enqueue_functions/kernel_arg_span.cpp +++ b/sycl/test/extensions/enqueue_functions/kernel_arg_span.cpp @@ -1,13 +1,17 @@ // 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 +// An argument list held in a container converts to the std::span that the +// nd_launch sequence 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. +// +// The sequence overloads take a std::span, so they exist only in C++20 and +// later; the first RUN line checks that the parameter pack overloads are +// unaffected without it. #include @@ -23,27 +27,28 @@ void argument_list_spellings(sycl::queue Q, sycl::handler &CGH, sycl::nd_range<1> Range, const sycl::kernel &Kernel) { int Value = 1; +#if __cpp_lib_span std::vector Vector{{&Value, sizeof(Value)}}; std::array Array{ oneapiext::raw_kernel_arg{&Value, sizeof(Value)}}; - sycl::span Span{Vector.data(), + std::span Span{Vector.data(), Vector.size()}; + std::span MutableSpan{Vector.data(), Vector.size()}; - sycl::span MutableSpan{Vector.data(), - Vector.size()}; + // A sycl::span still converts, so a caller holding one does not have to + // change how it stores its arguments. + sycl::span SyclSpan{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(Q, Range, Kernel, SyclSpan); 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); + oneapiext::nd_launch(CGH, Range, Kernel, SyclSpan); #endif // One raw_kernel_arg is one argument, and typed arguments are unaffected. From c23f97539745c13cf36bba4b940ac25862e3a987 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 2 Sep 2026 13:36:57 +0000 Subject: [PATCH 06/24] [SYCL] Apply clang-format to the argument sequence changes Line wrapping only, as git-clang-format produces it for the previous commit. The two test comments also said that a container bound as a single kernel argument produces wrong results at run time. It does not get that far: the launch faults on Level Zero and the OpenCL adapter rejects it, so the comments now say that it only fails once the kernel is launched. --- .../ext/oneapi/experimental/enqueue_functions.hpp | 11 ++++++----- sycl/include/sycl/queue.hpp | 5 ++--- sycl/source/queue.cpp | 5 ++--- .../EnqueueFunctions/nd_launch_kernel_obj_span.cpp | 6 +++--- .../enqueue_functions/kernel_arg_sequence_cxx17.cpp | 2 +- .../extensions/enqueue_functions/kernel_arg_span.cpp | 7 ++++--- 6 files changed, 18 insertions(+), 18 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp index 95b3c3b7e3a23..e56b26b47e1b8 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -495,11 +495,12 @@ void nd_launch(handler &CGH, nd_range Range, #if !__cpp_lib_span // Reached for a container both directly and through the queue overload, which // falls back to a command group for anything that is not a direct argument. - static_assert(!detail::is_arg_list_sequence_v, - "Passing the arguments of a sycl::kernel as a sequence requires " - "C++20, where the nd_launch overload taking a std::span of " - "raw_kernel_arg is available. Compile with C++20 or pass the " - "arguments as a parameter pack."); + static_assert( + !detail::is_arg_list_sequence_v, + "Passing the arguments of a sycl::kernel as a sequence requires " + "C++20, where the nd_launch overload taking a std::span of " + "raw_kernel_arg is available. Compile with C++20 or pass the " + "arguments as a parameter pack."); #endif if constexpr (detail::is_arg_list_container_v) { #if __cpp_lib_span diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp index 90c3ef20c1c50..24e32628a773c 100644 --- a/sycl/include/sycl/queue.hpp +++ b/sycl/include/sycl/queue.hpp @@ -86,9 +86,8 @@ void __SYCL_EXPORT submit_kernel_obj_direct_without_event_impl( // 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, - const ext::oneapi::experimental::raw_kernel_arg *Args, size_t NumArgs, - const detail::code_location &CodeLoc, bool IsTopCodeLoc); + const kernel &Kernel, const ext::oneapi::experimental::raw_kernel_arg *Args, + size_t NumArgs, 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, diff --git a/sycl/source/queue.cpp b/sycl/source/queue.cpp index 74f76d97a460f..af39545a259ba 100644 --- a/sycl/source/queue.cpp +++ b/sycl/source/queue.cpp @@ -330,9 +330,8 @@ void submit_kernel_obj_direct_without_event_impl( void submit_kernel_obj_direct_without_event_impl( const queue &Queue, const detail::nd_range_view &RangeView, - const kernel &Kernel, - const ext::oneapi::experimental::raw_kernel_arg *Args, size_t NumArgs, - const detail::code_location &CodeLoc, bool IsTopCodeLoc) { + const kernel &Kernel, const ext::oneapi::experimental::raw_kernel_arg *Args, + size_t NumArgs, const detail::code_location &CodeLoc, bool IsTopCodeLoc) { detail::getSyclObjImpl(Queue)->submit_kernel_obj_direct_without_event( RangeView, detail::getSyclObjImpl(Kernel), {Args, NumArgs}, CodeLoc, IsTopCodeLoc); 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 2aa044929dc20..3fe3ea195dbc1 100644 --- a/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_span.cpp +++ b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_span.cpp @@ -141,9 +141,9 @@ int main() { std::vector OneArg{ {&Memory, oneapiext::pointer_arg}}; Q.memset(Memory, 0, N * sizeof(int)); - oneapiext::nd_launch(Q, Ndr, getKernel(Q), - std::span{ - OneArg.data(), OneArg.size()}); + oneapiext::nd_launch( + Q, Ndr, getKernel(Q), + std::span{OneArg.data(), OneArg.size()}); Q.wait(); for (size_t I = 0; I < N; ++I) Failed += Check(Memory, 1, I, "one element span"); diff --git a/sycl/test/extensions/enqueue_functions/kernel_arg_sequence_cxx17.cpp b/sycl/test/extensions/enqueue_functions/kernel_arg_sequence_cxx17.cpp index 8cf32286ac006..bbcb109f844dd 100644 --- a/sycl/test/extensions/enqueue_functions/kernel_arg_sequence_cxx17.cpp +++ b/sycl/test/extensions/enqueue_functions/kernel_arg_sequence_cxx17.cpp @@ -9,7 +9,7 @@ // std::span, so they do not exist before C++20. A caller that passes its // argument list as a container there would otherwise have the container bound // as a single kernel argument, which compiles for any trivially copyable one -// and produces wrong results at run time, so it is diagnosed instead. +// and only fails once the kernel is launched, so it is diagnosed instead. // CHECK: Passing the arguments of a sycl::kernel as a sequence requires C++20 diff --git a/sycl/test/extensions/enqueue_functions/kernel_arg_span.cpp b/sycl/test/extensions/enqueue_functions/kernel_arg_span.cpp index 68af2e1fcfdb5..7b4f8ac8670da 100644 --- a/sycl/test/extensions/enqueue_functions/kernel_arg_span.cpp +++ b/sycl/test/extensions/enqueue_functions/kernel_arg_span.cpp @@ -5,9 +5,10 @@ // nd_launch sequence 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. +// compiles for any trivially copyable container and only fails once the kernel +// is launched, with a fault or an adapter error that says nothing about the +// argument list. Check that every spelling of an argument list is accepted, and +// that a single raw_kernel_arg is still one argument. // // The sequence overloads take a std::span, so they exist only in C++20 and // later; the first RUN line checks that the parameter pack overloads are From 82aca3c50abbfdc0ed39731937789dab4dcdd034 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Thu, 3 Sep 2026 12:54:18 +0000 Subject: [PATCH 07/24] [SYCL] Rewrap the argument sequence dependency paragraph The paragraph broke after "depend on" and after "version 2 of", leaving two lines of a few words each. It is filled to the line width now, with the link macro split inside its brackets, the way the other paragraphs in this file that end in a link do it. --- .../sycl_ext_oneapi_enqueue_functions.asciidoc | 12 +++++------- 1 file changed, 5 insertions(+), 7 deletions(-) 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 a4c0d33325fe1..9bcb874afd540 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc @@ -41,13 +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 `std::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. -Those overloads use `std::span`, so they are available only when the application -is compiled with C++20 or later. +The `nd_launch` overloads that take the kernel arguments as a `std::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. Those overloads use `std::span`, so they are +available only when the application is compiled with C++20 or later. == Status From 5502f83c710e005de32f773907134d4b91e96432 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Thu, 3 Sep 2026 12:54:40 +0000 Subject: [PATCH 08/24] [SYCL] Make the C++20 annotations reach the rendered specification A line starting with `//` inside an AsciiDoc table cell is a comment, so the two `// Requires C++20` lines were dropped: neither asciidoctor nor GitHub renders them, while the declarations they annotate are shown. The annotation is a trailing comment now, which is ordinary listing text and is rendered. --- .../experimental/sycl_ext_oneapi_enqueue_functions.asciidoc | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) 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 9bcb874afd540..98259254f3440 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc @@ -542,13 +542,11 @@ a! ---- namespace sycl::ext::oneapi::experimental { -// Requires C++20 -template +template // Requires C++20 void nd_launch(sycl::queue q, sycl::nd_range r, const sycl::kernel& k, std::span args); -// Requires C++20 -template +template // Requires C++20 void nd_launch(sycl::handler &h, sycl::nd_range r, const sycl::kernel& k, std::span args); From 9608d555271bd94d966f83444e62058d3a8a9167 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 23 Sep 2026 11:04:49 +0200 Subject: [PATCH 09/24] Update sycl/doc/extensions/experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc Co-authored-by: Greg Lueck --- .../experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) 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 6826c2c9b719f..4c58d027215cd 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 @@ -176,8 +176,7 @@ 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 +stored at `pointer_location`, to be bound as a pointer argument. 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. From c01fd9a72ce8944f4fdb308b6a489180f246ffe1 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 23 Sep 2026 09:10:21 +0000 Subject: [PATCH 10/24] [SYCL] Rewrap the applied suggestion for the pointer constructor --- .../experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) 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 4c58d027215cd..163870e747797 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 @@ -176,10 +176,10 @@ 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. 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. +stored at `pointer_location`, to be bound as a pointer argument. 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 From 89fb4070b93e46e1489e57f9c1f2c93fadf65170 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 23 Sep 2026 09:10:21 +0000 Subject: [PATCH 11/24] [SYCL] Use the namespace alias in both raw_kernel_arg examples --- .../experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) 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 163870e747797..81128b3e27e86 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 @@ -82,12 +82,14 @@ for which there is no equivalent type defined on the host). [source,c++] ---- +namespace syclex = sycl::ext::oneapi::experimental; + int a; char* opaque_type; int nbytes; ... h.set_arg(0, a); -h.set_arg(1, sycl::ext::oneapi::experimental::raw_kernel_arg(opaque_type, nbytes)); +h.set_arg(1, syclex::raw_kernel_arg(opaque_type, nbytes)); h.parallel_for(range, kernel); ---- @@ -96,8 +98,6 @@ 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)); From 5a2f8b80e9b97bc3e7597f373dbfb066cb7a80a0 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 23 Sep 2026 09:10:22 +0000 Subject: [PATCH 12/24] [SYCL] Say that the byte constructor must not take a pointer argument --- .../experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) 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 81128b3e27e86..69e3824545fe3 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 @@ -158,8 +158,7 @@ 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. If the argument is a pointer, only the Level Zero backend is -guaranteed to bind it; see _Backend support status_ above. +copyable. The `bytes` must not be a pointer argument. _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` From 34bf2f1efe146af8376d365f1f36f173032a31fb Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 23 Sep 2026 09:10:22 +0000 Subject: [PATCH 13/24] [SYCL] Drop the argument sequence paragraph from the Overview --- .../experimental/sycl_ext_oneapi_enqueue_functions.asciidoc | 6 ------ 1 file changed, 6 deletions(-) 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 98259254f3440..7b8f583048675 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc @@ -93,12 +93,6 @@ existing queue shortcuts. A future version of this extension may adjust this overload set to include functions compatible with future C++ concepts (e.g, by accepting a scheduler and returning a sender). -The arguments of a `sycl::kernel` can be passed either as a parameter pack or as -a sequence, for an application that only learns its argument list at run time and -would otherwise have to instantiate a parameter pack overload for each number of -arguments it may encounter. The sequence is a `std::span`, so those overloads -require C++20. - === Usage example From 70e81772c78ac717affb19902145034d5fbc2a53 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 23 Sep 2026 09:10:22 +0000 Subject: [PATCH 14/24] [SYCL] Shorten the version 2 row of the enqueue functions extension --- .../experimental/sycl_ext_oneapi_enqueue_functions.asciidoc | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) 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 7b8f583048675..8547e69195383 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc @@ -175,9 +175,7 @@ implementation supports. |2 |Adds the `nd_launch` overloads that take the arguments of a `sycl::kernel` as - a `std::span` of `raw_kernel_arg`, which require C++20, and specifies that a - parameter pack overload handed such a sequence as its only argument passes it - on as the argument list. + a `std::span` of `raw_kernel_arg`. |=== From b51e275fcce489c505f4b55d3bb54b7e17a6b647 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 23 Sep 2026 09:10:23 +0000 Subject: [PATCH 15/24] [SYCL] Drop the Level Zero justification from the pointer argument note --- .../experimental/sycl_ext_oneapi_enqueue_functions.asciidoc | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) 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 8547e69195383..bff399ce31948 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc @@ -558,9 +558,8 @@ the kernel argument list and the effects are those of the overload taking a [_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_] +link:../experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc[sycl_ext_oneapi_raw_kernel_arg]. +_end note_] a| [frame=all,grid=none] From bd98858a967b31c0d3fca1d716ef1f0b6c4f6017 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 23 Sep 2026 09:19:22 +0000 Subject: [PATCH 16/24] [SYCL] Move the pointer argument requirement to the Overview --- .../sycl_ext_oneapi_raw_kernel_arg.asciidoc | 13 ++++--------- 1 file changed, 4 insertions(+), 9 deletions(-) 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 69e3824545fe3..7253047509bb1 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 @@ -59,13 +59,6 @@ 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 @@ -93,8 +86,10 @@ h.set_arg(1, syclex::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: +A kernel argument that is a pointer must be constructed with the `pointer_arg` +overload described below, 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: [source,c++] ---- From a2576495baa02e03eda8ed531c3aa73809f44935 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 23 Sep 2026 09:19:34 +0000 Subject: [PATCH 17/24] [SYCL] Use the reviewer's wording for the pointer argument requirement --- .../experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) 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 7253047509bb1..7ec02504ae98b 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 @@ -86,8 +86,8 @@ h.set_arg(1, syclex::raw_kernel_arg(opaque_type, nbytes)); h.parallel_for(range, kernel); ---- -A kernel argument that is a pointer must be constructed with the `pointer_arg` -overload described below, because a backend may take a pointer argument through a +When a kernel argument is a pointer, a special overload must be used to indicate +that the argument is a pointer. 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: From 4538a467cc86084f86fde51872682725bb4bc106 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 23 Sep 2026 09:19:48 +0000 Subject: [PATCH 18/24] [SYCL] Limit the backend support status to the tested backends --- .../sycl_ext_oneapi_raw_kernel_arg.asciidoc | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) 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 7ec02504ae98b..5a47d4e8098b4 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,10 +54,10 @@ specification.* == Backend support status -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. +This extension is currently implemented in {dpcpp} for the Level Zero and OpenCL +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. == Overview @@ -88,8 +88,8 @@ h.parallel_for(range, kernel); When a kernel argument is a pointer, a special overload must be used to indicate that the argument is a pointer. 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: +different entry point than the one that takes plain bytes, and it cannot tell +the two apart from the bytes alone: [source,c++] ---- From 5c15db5e71a0ce8fa161e9e1dc471efb039a7ae1 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 23 Sep 2026 12:52:16 +0000 Subject: [PATCH 19/24] [SYCL] Rewrite the argument sequence example as a complete program --- ...sycl_ext_oneapi_enqueue_functions.asciidoc | 42 ++++++++++++++----- 1 file changed, 31 insertions(+), 11 deletions(-) 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 bff399ce31948..edc65b155db3e 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc @@ -815,25 +815,45 @@ _Effects_: Submits an executable command graph to the `sycl::queue` or `sycl::ha |==== -== Examples +== Example -An argument list that is only known at run time is held in a container and passed -as the sequence of arguments: a `std::vector`, a `std::array` of -them and a `std::span` of them all convert to `std::span`. -These overloads require C++20. +The following example launches a kernel whose argument list is assembled at run +time, as a launcher for kernels that are not known at compile time would do, so +the arguments are passed as a sequence rather than as a parameter pack. A +`std::vector`, a `std::array` of them and a `std::span` of them +all convert to `std::span`. [source,c++] ---- +#include +#include namespace syclex = sycl::ext::oneapi::experimental; -int count = 1024; -float *data = sycl::malloc_device(count, q); +static constexpr size_t N = 1024; -std::vector args; -args.emplace_back(&data, syclex::pointer_arg); -args.emplace_back(&count, sizeof(count)); +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclex::nd_range_kernel<1>)) +void scale(float *data, int factor) { /* kernel body */ } -syclex::nd_launch(q, sycl::nd_range<1>{count, 256}, k, args); +int main() { + sycl::queue q; + + auto bundle = + syclex::get_kernel_bundle( + q.get_context()); + sycl::kernel k = bundle.ext_oneapi_get_kernel(); + + float *data = sycl::malloc_device(N, q); + int factor = 2; + + std::vector args; + args.emplace_back(&data, syclex::pointer_arg); + args.emplace_back(&factor, sizeof(factor)); + + syclex::nd_launch(q, sycl::nd_range<1>{{N}, {256}}, k, args); + q.wait(); + + sycl::free(data, q); +} ---- == Issues From 551d09aff574a16bbe112cdc311c850b17485535 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 23 Sep 2026 12:59:43 +0000 Subject: [PATCH 20/24] [SYCL] Constrain the parameter pack overloads instead of describing the forwarding --- .../sycl_ext_oneapi_enqueue_functions.asciidoc | 8 +++----- 1 file changed, 3 insertions(+), 5 deletions(-) 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 edc65b155db3e..e4a8eb9172b17 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc @@ -521,6 +521,9 @@ void nd_launch(sycl::handler &h, sycl::nd_range r, } ---- !==== +_Constraints_: `sizeof...(Args)` is not one, or the single argument in `args` is +not convertible to `std::span`. + _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 arguments in `args` are passed to the kernel in the same @@ -551,11 +554,6 @@ 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. -If one of the parameter pack overloads above is called with a single argument -whose type is convertible to `std::span`, that argument is -the kernel argument list and the effects are those of the overload taking a -`std::span`. - [_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]. From ccb49bb76f0b88c6fb18a81f803d7ef88d7e519b Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Wed, 23 Sep 2026 15:43:25 +0000 Subject: [PATCH 21/24] [SYCL] Keep the parameter pack literal in the rendered constraint MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit AsciiDoc replaces three dots with an ellipsis character, inside a monospace span as well, so the constraint rendered as `sizeof…(Args)`. Pass the text through, as the free function kernels specification already does for `Args...`. --- .../experimental/sycl_ext_oneapi_enqueue_functions.asciidoc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) 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 e4a8eb9172b17..9d5dc7b2b139c 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc @@ -521,8 +521,8 @@ void nd_launch(sycl::handler &h, sycl::nd_range r, } ---- !==== -_Constraints_: `sizeof...(Args)` is not one, or the single argument in `args` is -not convertible to `std::span`. +_Constraints_: `+sizeof...(Args)+` is not one, or the single argument in +`args` is not convertible to `std::span`. _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 From 5317dd6beb4b428c0a08d1818dadab57baafc7a5 Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Fri, 25 Sep 2026 07:19:23 +0000 Subject: [PATCH 22/24] [SYCL] Add std::span overloads to the remaining kernel object launch functions --- ...sycl_ext_oneapi_enqueue_functions.asciidoc | 155 ++++++++- .../oneapi/experimental/enqueue_functions.hpp | 294 ++++++++++++++++-- .../kernel_obj_span_with_kb.cpp | 170 ++++++++++ .../nd_launch_kernel_obj_span.cpp | 6 +- .../kernel_arg_sequence_cxx17.cpp | 32 +- .../enqueue_functions/kernel_arg_span.cpp | 138 ++++++-- 6 files changed, 721 insertions(+), 74 deletions(-) create mode 100644 sycl/test-e2e/EnqueueFunctions/kernel_obj_span_with_kb.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 9d5dc7b2b139c..caf1a110f460a 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc @@ -41,8 +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 `std::span` depend -on version 2 of link:../experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc[ +The overloads that take the kernel arguments as a `std::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. Those overloads use `std::span`, so they are available only when the application is compiled with C++20 or later. @@ -174,7 +174,7 @@ implementation supports. |Initial version of this extension. |2 -|Adds the `nd_launch` overloads that take the arguments of a `sycl::kernel` as +|Adds the overloads that take the arguments of a `sycl::kernel` as a `std::span` of `raw_kernel_arg`. |=== @@ -310,19 +310,49 @@ a! ---- namespace sycl::ext::oneapi::experimental { -template +template void single_task(sycl::queue q, const sycl::kernel& k, Args&&... args); -template +template void single_task(sycl::handler &h, const sycl::kernel& k, Args&&... args); } ---- !==== +_Constraints_: `+sizeof...(Args)+` is not one, or the single argument in +`args` is not convertible to `std::span`. + _Effects_: Enqueues a kernel object to the `sycl::queue` or `sycl::handler` as a single task. The arguments in `args` are passed to the kernel in the same order. +a| +[frame=all,grid=none] +!==== +a! +[source,c++] +---- +namespace sycl::ext::oneapi::experimental { + +void single_task(sycl::queue q, const sycl::kernel& k, // Requires C++20 + std::span args); + +void single_task(sycl::handler &h, const sycl::kernel& k, // Requires C++20 + std::span args); + +} +---- +!==== +_Effects_: Enqueues a kernel object to the `sycl::queue` or `sycl::handler` as +a single task. The element `args[i]` is passed to the kernel as the argument 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. + +[_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]. +_end note_] + |==== @@ -406,6 +436,9 @@ void parallel_for(sycl::handler &h, sycl::range r, } ---- !==== +_Constraints_: `+sizeof...(Args)+` is not one, or the single argument in +`args` is not convertible to `std::span`. + _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::range`. The arguments in `args` are passed to the kernel in the same order. @@ -418,24 +451,93 @@ a! ---- namespace sycl::ext::oneapi::experimental { +template // Requires C++20 +void parallel_for(sycl::queue q, sycl::range r, + const sycl::kernel& k, + std::span args); + +template // Requires C++20 +void parallel_for(sycl::handler &h, sycl::range r, + const sycl::kernel& k, + std::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::range`. +The element `args[i]` is passed to the kernel as the argument 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. + +[_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]. +_end note_] + +a| +[frame=all,grid=none] +!==== +a! +[source,c++] +---- +namespace sycl::ext::oneapi::experimental { + template void parallel_for(sycl::queue q, launch_config, Properties> c, - const sycl::kernel& k, Args&& args...); + const sycl::kernel& k, Args&&... args); template void parallel_for(sycl::handler &h, launch_config, Properties> c, - const sycl::kernel& k, Args&& args...); + const sycl::kernel& k, Args&&... args); } ---- !==== +_Constraints_: `+sizeof...(Args)+` is not one, or the single argument in +`args` is not convertible to `std::span`. + _Effects_: Enqueues a kernel object to the `sycl::queue` or `sycl::handler` as a basic kernel, using the launch configuration specified by a `launch_config`. The arguments in `args` are passed to the kernel in the same order. +a| +[frame=all,grid=none] +!==== +a! +[source,c++] +---- +namespace sycl::ext::oneapi::experimental { + +template // Requires C++20 +void parallel_for(sycl::queue q, + launch_config, Properties> c, + const sycl::kernel& k, + std::span args); + +template // Requires C++20 +void parallel_for(sycl::handler &h, + launch_config, Properties> c, + const sycl::kernel& k, + std::span args); + +} +---- +!==== +_Effects_: Enqueues a kernel object to the `sycl::queue` or `sycl::handler` +as a basic kernel, using the launch configuration specified by a +`launch_config`. The element `args[i]` is passed to the kernel as the argument +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. + +[_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]. +_end note_] + |==== @@ -570,21 +672,56 @@ namespace sycl::ext::oneapi::experimental { template void nd_launch(sycl::queue q, launch_config, Properties> c, - const sycl::kernel& k, Args&& args...); + const sycl::kernel& k, Args&&... args); template void nd_launch(sycl::handler &h, launch_config, Properties> c, - const sycl::kernel& k, Args&& args...); + const sycl::kernel& k, Args&&... args); } ---- !==== +_Constraints_: `+sizeof...(Args)+` is not one, or the single argument in +`args` is not convertible to `std::span`. + _Effects_: Enqueues a kernel object to the `sycl::queue` or `sycl::handler` as a basic kernel, using the launch configuration specified by a `launch_config`. The arguments in `args` are passed to the kernel in the same order. +a| +[frame=all,grid=none] +!==== +a! +[source,c++] +---- +namespace sycl::ext::oneapi::experimental { + +template // Requires C++20 +void nd_launch(sycl::queue q, + launch_config, Properties> c, + const sycl::kernel& k, std::span args); + +template // Requires C++20 +void nd_launch(sycl::handler &h, + launch_config, Properties> c, + const sycl::kernel& k, std::span args); + +} +---- +!==== +_Effects_: Enqueues a kernel object to the `sycl::queue` or `sycl::handler` +as a basic kernel, using the launch configuration specified by a +`launch_config`. The element `args[i]` is passed to the kernel as the argument +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. + +[_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]. +_end note_] + |==== diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp index e56b26b47e1b8..086b94d4b4836 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -146,6 +146,18 @@ sycl::detail::KernelArgView makeKernelArgView(const T &Arg) { return {&Arg, sizeof(ArgT), kernel_arg_kind_v}; } +#if __cpp_lib_span +// Binds an argument list that is held as a sequence, for the overloads that +// enqueue through a handler. +inline void setRawKernelArgs(handler &CGH, + std::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]}); +} +#endif + // 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 @@ -181,6 +193,18 @@ template inline constexpr bool is_arg_list_sequence_v = sizeof...(ArgsT) == 1 && (is_raw_kernel_arg_sequence_v> && ...); + +// Rejects an argument list passed as a container before C++20. The handler +// overloads call it, and the queue overloads reach them for such a call. +// `LaunchT` is the range or launch configuration, `void` for `single_task`, so +// that each launch function called that way is diagnosed, not just the first. +template void diagnoseArgListSequence() { + static_assert(!is_arg_list_sequence_v, + "Passing the arguments of a sycl::kernel as a sequence " + "requires C++20, where the overloads taking a std::span of " + "raw_kernel_arg are available. Compile with C++20 or pass the " + "arguments as a parameter pack."); +} #endif template @@ -252,19 +276,64 @@ void single_task(queue Q, const KernelType &KernelObj, } } +#if __cpp_lib_span +inline void single_task(handler &CGH, const kernel &KernelObj, + std::span Args); +#endif + template void single_task(handler &CGH, const kernel &KernelObj, ArgsT &&...Args) { - CGH.set_args(std::forward(Args)...); - CGH.single_task(KernelObj); +#if !__cpp_lib_span + detail::diagnoseArgListSequence(); +#endif + if constexpr (detail::is_arg_list_container_v) { +#if __cpp_lib_span + single_task(CGH, KernelObj, + std::span{std::forward(Args)...}); +#endif + } else { + CGH.set_args(std::forward(Args)...); + CGH.single_task(KernelObj); + } } +#if __cpp_lib_span +inline void single_task(queue Q, const kernel &KernelObj, + std::span Args, + const sycl::detail::code_location &CodeLoc = + sycl::detail::code_location::current()); +#endif + template void single_task(queue Q, const kernel &KernelObj, ArgsT &&...Args) { - submit(std::move(Q), [&](handler &CGH) { - single_task(CGH, KernelObj, std::forward(Args)...); - }); + if constexpr (detail::is_arg_list_container_v) { +#if __cpp_lib_span + single_task(std::move(Q), KernelObj, + std::span{std::forward(Args)...}); +#endif + } else { + submit(std::move(Q), [&](handler &CGH) { + single_task(CGH, KernelObj, std::forward(Args)...); + }); + } } +#if __cpp_lib_span +inline void single_task(handler &CGH, const kernel &KernelObj, + std::span Args) { + detail::setRawKernelArgs(CGH, Args); + CGH.single_task(KernelObj); +} + +inline void single_task(queue Q, const kernel &KernelObj, + std::span Args, + const sycl::detail::code_location &CodeLoc) { + submit( + std::move(Q), [&](handler &CGH) { single_task(CGH, KernelObj, Args); }, + CodeLoc); +} +#endif // __cpp_lib_span + // Free function kernel single_task enqueue functions. These enqueue the free // function kernel `Func` directly instead of wrapping `Func` in a helper // kernel. Wrapping generated a second, duplicate device kernel and dropped the @@ -382,40 +451,151 @@ void parallel_for(queue Q, launch_config, Properties> Config, } } +#if __cpp_lib_span +template +void parallel_for(handler &CGH, range Range, + const kernel &KernelObj, + std::span Args); +#endif + template void parallel_for(handler &CGH, range Range, const kernel &KernelObj, ArgsT &&...Args) { - CGH.set_args(std::forward(Args)...); - CGH.parallel_for(Range, KernelObj); +#if !__cpp_lib_span + detail::diagnoseArgListSequence, ArgsT...>(); +#endif + if constexpr (detail::is_arg_list_container_v) { +#if __cpp_lib_span + parallel_for(CGH, Range, KernelObj, + std::span{std::forward(Args)...}); +#endif + } else { + CGH.set_args(std::forward(Args)...); + CGH.parallel_for(Range, KernelObj); + } } +#if __cpp_lib_span +template +void parallel_for(queue Q, range Range, const kernel &KernelObj, + std::span Args, + const sycl::detail::code_location &CodeLoc = + sycl::detail::code_location::current()); +#endif + template void parallel_for(queue Q, range Range, const kernel &KernelObj, ArgsT &&...Args) { - submit(std::move(Q), [&](handler &CGH) { - parallel_for(CGH, Range, KernelObj, std::forward(Args)...); - }); + if constexpr (detail::is_arg_list_container_v) { +#if __cpp_lib_span + parallel_for(std::move(Q), Range, KernelObj, + std::span{std::forward(Args)...}); +#endif + } else { + submit(std::move(Q), [&](handler &CGH) { + parallel_for(CGH, Range, KernelObj, std::forward(Args)...); + }); + } +} + +#if __cpp_lib_span +template +void parallel_for(handler &CGH, range Range, + const kernel &KernelObj, + std::span Args) { + detail::setRawKernelArgs(CGH, Args); + CGH.parallel_for(Range, KernelObj); } +template +void parallel_for(queue Q, range Range, const kernel &KernelObj, + std::span Args, + const sycl::detail::code_location &CodeLoc) { + submit( + std::move(Q), + [&](handler &CGH) { parallel_for(CGH, Range, KernelObj, Args); }, + CodeLoc); +} +#endif // __cpp_lib_span + +#if __cpp_lib_span +template +void parallel_for(handler &CGH, + launch_config, Properties> Config, + const kernel &KernelObj, + std::span Args); +#endif + template void parallel_for(handler &CGH, launch_config, Properties> Config, const kernel &KernelObj, ArgsT &&...Args) { +#if !__cpp_lib_span + detail::diagnoseArgListSequence, Properties>, + ArgsT...>(); +#endif + if constexpr (detail::is_arg_list_container_v) { +#if __cpp_lib_span + parallel_for(CGH, Config, KernelObj, + std::span{std::forward(Args)...}); +#endif + } else { + ext::oneapi::experimental::detail::LaunchConfigAccess, + Properties> + ConfigAccess(Config); + CGH.set_args(std::forward(Args)...); + sycl::detail::HandlerAccess::parallelForImpl( + CGH, ConfigAccess.getRange(), ConfigAccess.getProperties(), KernelObj); + } +} + +#if __cpp_lib_span +template +void parallel_for(queue Q, launch_config, Properties> Config, + const kernel &KernelObj, std::span Args, + const sycl::detail::code_location &CodeLoc = + sycl::detail::code_location::current()); +#endif + +template +void parallel_for(queue Q, launch_config, Properties> Config, + const kernel &KernelObj, ArgsT &&...Args) { + if constexpr (detail::is_arg_list_container_v) { +#if __cpp_lib_span + parallel_for(std::move(Q), Config, KernelObj, + std::span{std::forward(Args)...}); +#endif + } else { + submit(std::move(Q), [&](handler &CGH) { + parallel_for(CGH, Config, KernelObj, std::forward(Args)...); + }); + } +} + +#if __cpp_lib_span +template +void parallel_for(handler &CGH, + launch_config, Properties> Config, + const kernel &KernelObj, + std::span Args) { ext::oneapi::experimental::detail::LaunchConfigAccess, Properties> ConfigAccess(Config); - CGH.set_args(std::forward(Args)...); + detail::setRawKernelArgs(CGH, Args); sycl::detail::HandlerAccess::parallelForImpl( CGH, ConfigAccess.getRange(), ConfigAccess.getProperties(), KernelObj); } -template +template void parallel_for(queue Q, launch_config, Properties> Config, - const kernel &KernelObj, ArgsT &&...Args) { - submit(std::move(Q), [&](handler &CGH) { - parallel_for(CGH, Config, KernelObj, std::forward(Args)...); - }); + const kernel &KernelObj, std::span Args, + const sycl::detail::code_location &CodeLoc) { + submit( + std::move(Q), + [&](handler &CGH) { parallel_for(CGH, Config, KernelObj, Args); }, + CodeLoc); } +#endif // __cpp_lib_span template @@ -493,14 +673,7 @@ template void nd_launch(handler &CGH, nd_range Range, const kernel &KernelObj, ArgsT &&...Args) { #if !__cpp_lib_span - // Reached for a container both directly and through the queue overload, which - // falls back to a command group for anything that is not a direct argument. - static_assert( - !detail::is_arg_list_sequence_v, - "Passing the arguments of a sycl::kernel as a sequence requires " - "C++20, where the nd_launch overload taking a std::span of " - "raw_kernel_arg is available. Compile with C++20 or pass the " - "arguments as a parameter pack."); + detail::diagnoseArgListSequence, ArgsT...>(); #endif if constexpr (detail::is_arg_list_container_v) { #if __cpp_lib_span @@ -556,10 +729,7 @@ void nd_launch(queue Q, nd_range Range, const kernel &KernelObj, template void nd_launch(handler &CGH, nd_range Range, const kernel &KernelObj, std::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]}); + detail::setRawKernelArgs(CGH, Args); CGH.parallel_for(Range, KernelObj); } @@ -577,25 +747,81 @@ void nd_launch(queue Q, nd_range Range, const kernel &KernelObj, } #endif // __cpp_lib_span +#if __cpp_lib_span +template +void nd_launch(handler &CGH, + launch_config, Properties> Config, + const kernel &KernelObj, std::span Args); +#endif + template void nd_launch(handler &CGH, launch_config, Properties> Config, const kernel &KernelObj, ArgsT &&...Args) { +#if !__cpp_lib_span + detail::diagnoseArgListSequence< + launch_config, Properties>, ArgsT...>(); +#endif + if constexpr (detail::is_arg_list_container_v) { +#if __cpp_lib_span + nd_launch(CGH, Config, KernelObj, + std::span{std::forward(Args)...}); +#endif + } else { + ext::oneapi::experimental::detail::LaunchConfigAccess, + Properties> + ConfigAccess(Config); + CGH.set_args(std::forward(Args)...); + sycl::detail::HandlerAccess::parallelForImpl( + CGH, ConfigAccess.getRange(), ConfigAccess.getProperties(), KernelObj); + } +} + +#if __cpp_lib_span +template +void nd_launch(queue Q, launch_config, Properties> Config, + const kernel &KernelObj, std::span Args, + const sycl::detail::code_location &CodeLoc = + sycl::detail::code_location::current()); +#endif + +template +void nd_launch(queue Q, launch_config, Properties> Config, + const kernel &KernelObj, ArgsT &&...Args) { + if constexpr (detail::is_arg_list_container_v) { +#if __cpp_lib_span + nd_launch(std::move(Q), Config, KernelObj, + std::span{std::forward(Args)...}); +#endif + } else { + submit(std::move(Q), [&](handler &CGH) { + nd_launch(CGH, Config, KernelObj, std::forward(Args)...); + }); + } +} + +#if __cpp_lib_span +template +void nd_launch(handler &CGH, + launch_config, Properties> Config, + const kernel &KernelObj, std::span Args) { ext::oneapi::experimental::detail::LaunchConfigAccess, Properties> ConfigAccess(Config); - CGH.set_args(std::forward(Args)...); + detail::setRawKernelArgs(CGH, Args); sycl::detail::HandlerAccess::parallelForImpl( CGH, ConfigAccess.getRange(), ConfigAccess.getProperties(), KernelObj); } -template +template void nd_launch(queue Q, launch_config, Properties> Config, - const kernel &KernelObj, ArgsT &&...Args) { - submit(std::move(Q), [&](handler &CGH) { - nd_launch(CGH, Config, KernelObj, std::forward(Args)...); - }); + const kernel &KernelObj, std::span Args, + const sycl::detail::code_location &CodeLoc) { + submit( + std::move(Q), + [&](handler &CGH) { nd_launch(CGH, Config, KernelObj, Args); }, CodeLoc); } +#endif // __cpp_lib_span // Free function kernel nd_launch enqueue functions. These enqueue the free // function kernel `Func` directly instead of wrapping `Func` in a helper diff --git a/sycl/test-e2e/EnqueueFunctions/kernel_obj_span_with_kb.cpp b/sycl/test-e2e/EnqueueFunctions/kernel_obj_span_with_kb.cpp new file mode 100644 index 0000000000000..87a095f212c62 --- /dev/null +++ b/sycl/test-e2e/EnqueueFunctions/kernel_obj_span_with_kb.cpp @@ -0,0 +1,170 @@ +// REQUIRES: aspect-usm_shared_allocations + +// DEFINE: %{cpp20} = %if cl_options %{/clang:-std=c++20%} %else %{-std=c++20%} + +// RUN: %{build} %{cpp20} -o %t.out +// RUN: %{run} %t.out + +// Tests the single_task and parallel_for overloads, and the launch_config forms +// of parallel_for and nd_launch, that take the arguments of a sycl::kernel as a +// std::span of raw_kernel_arg, i.e. an argument list whose length is only known +// at run time. Those overloads take a std::span, hence C++20. They have to bind +// the same arguments in the same order as the parameter pack overloads, on the +// queue and on the handler alike. +// NOTE: This relies on the availability of an OpenCL C compiler. + +#include + +#include +#include +#include +#include +#include + +#include "common.hpp" + +#include + +static_assert(SYCL_EXT_ONEAPI_ENQUEUE_FUNCTIONS >= 2, + "These span overloads require version 2 of the extension"); + +constexpr size_t N = 1024; +constexpr size_t WGSize = 8; + +int main() { + sycl::queue Q{sycl::property::queue::in_order{}}; + + if (!Q.get_device().ext_oneapi_can_build( + oneapiext::source_language::opencl)) { + std::cout + << "Backend does not support OpenCL C source kernel bundle extension: " + << Q.get_backend() << std::endl; + return 0; + } + + auto KB = CreateKB(Q); + sycl::kernel KernelSingleTask = KB.ext_oneapi_get_kernel("KernelSingleTask"); + sycl::kernel Kernel1D = KB.ext_oneapi_get_kernel("Kernel1D"); + sycl::kernel Kernel2D = KB.ext_oneapi_get_kernel("Kernel2D"); + sycl::kernel Kernel3D = KB.ext_oneapi_get_kernel("Kernel3D"); + + int *Memory = sycl::malloc_shared(N, Q); + + int Count = N; + int Value = 0; + + // A pointer argument has to say that it is one, since the byte form of a + // pointer must not be passed to the byte overload of raw_kernel_arg. + std::vector SingleTaskArgs; + SingleTaskArgs.emplace_back(&Count, sizeof(Count)); + SingleTaskArgs.emplace_back(&Value, sizeof(Value)); + SingleTaskArgs.emplace_back(&Memory, oneapiext::pointer_arg); + std::span SingleTaskSpan{ + SingleTaskArgs.data(), SingleTaskArgs.size()}; + + std::vector Args; + Args.emplace_back(&Value, sizeof(Value)); + Args.emplace_back(&Memory, oneapiext::pointer_arg); + std::span ArgSpan{Args.data(), Args.size()}; + + sycl::range<1> R1{N}; + sycl::range<2> R2{8, N / 8}; + sycl::range<3> R3{8, 8, N / 64}; + sycl::nd_range<1> Ndr1{R1, sycl::range<1>{WGSize}}; + sycl::nd_range<2> Ndr2{R2, sycl::range<2>{WGSize, WGSize}}; + + int Failed = 0; + + // Each launch below passes a different value, so a launch that binds a value + // left over from an earlier one, or binds the arguments in the wrong order, + // shows up as a wrong result. + auto Run = [&](int Written, const char *Name, auto &&Launch) { + Q.memset(Memory, 0, N * sizeof(int)); + Value = Written; + Launch(); + Q.wait(); + for (size_t I = 0; I < N; ++I) + Failed += Check(Memory, Written, I, Name); + }; + + Run(41, "single_task span overload", + [&] { oneapiext::single_task(Q, KernelSingleTask, SingleTaskSpan); }); + + Run(42, "single_task span overload on the handler", [&] { + Q.submit([&](sycl::handler &CGH) { + oneapiext::single_task(CGH, KernelSingleTask, SingleTaskSpan); + }); + }); + + // The container holding the arguments converts to that span, so passing it + // has to bind the arguments it holds rather than the container object. + Run(43, "single_task argument list passed as a container", + [&] { oneapiext::single_task(Q, KernelSingleTask, SingleTaskArgs); }); + + // The parameter pack overload has to agree element for element. + Run(44, "single_task parameter pack overload", [&] { + oneapiext::single_task( + Q, KernelSingleTask, oneapiext::raw_kernel_arg{&Count, sizeof(Count)}, + oneapiext::raw_kernel_arg{&Value, sizeof(Value)}, + oneapiext::raw_kernel_arg{&Memory, oneapiext::pointer_arg}); + }); + + Run(45, "1D parallel_for span overload", + [&] { oneapiext::parallel_for(Q, R1, Kernel1D, ArgSpan); }); + + Run(46, "2D parallel_for span overload on the handler", [&] { + Q.submit([&](sycl::handler &CGH) { + oneapiext::parallel_for(CGH, R2, Kernel2D, ArgSpan); + }); + }); + + Run(47, "3D parallel_for argument list passed as a container", + [&] { oneapiext::parallel_for(Q, R3, Kernel3D, Args); }); + + Run(48, "1D parallel_for parameter pack overload", [&] { + oneapiext::parallel_for( + Q, R1, Kernel1D, oneapiext::raw_kernel_arg{&Value, sizeof(Value)}, + oneapiext::raw_kernel_arg{&Memory, oneapiext::pointer_arg}); + }); + + Run(49, "1D parallel_for span overload with launch config", [&] { + oneapiext::parallel_for(Q, oneapiext::launch_config{R1}, Kernel1D, ArgSpan); + }); + + Run(50, "2D parallel_for span overload with launch config on the handler", + [&] { + Q.submit([&](sycl::handler &CGH) { + oneapiext::parallel_for(CGH, oneapiext::launch_config{R2}, Kernel2D, + ArgSpan); + }); + }); + + Run(51, "3D parallel_for container with launch config", [&] { + oneapiext::parallel_for(Q, oneapiext::launch_config{R3}, Kernel3D, Args); + }); + + Run(52, "1D nd_launch span overload with launch config", [&] { + oneapiext::nd_launch(Q, oneapiext::launch_config{Ndr1}, Kernel1D, ArgSpan); + }); + + Run(53, "2D nd_launch span overload with launch config on the handler", [&] { + Q.submit([&](sycl::handler &CGH) { + oneapiext::nd_launch(CGH, oneapiext::launch_config{Ndr2}, Kernel2D, + ArgSpan); + }); + }); + + Run(54, "1D nd_launch container with launch config", [&] { + oneapiext::nd_launch(Q, oneapiext::launch_config{Ndr1}, Kernel1D, Args); + }); + + Run(55, "1D nd_launch parameter pack overload with launch config", [&] { + oneapiext::nd_launch( + Q, oneapiext::launch_config{Ndr1}, Kernel1D, + oneapiext::raw_kernel_arg{&Value, sizeof(Value)}, + oneapiext::raw_kernel_arg{&Memory, oneapiext::pointer_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 3fe3ea195dbc1..1169ce26a8128 100644 --- a/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_span.cpp +++ b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_span.cpp @@ -68,9 +68,9 @@ 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. 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. + // exist for. A pointer argument has to say that it is one, since the byte + // form of a pointer must not be passed to the byte overload of + // raw_kernel_arg. std::vector Args; Args.emplace_back(&Memory, oneapiext::pointer_arg); Args.emplace_back(&A, sizeof(A)); diff --git a/sycl/test/extensions/enqueue_functions/kernel_arg_sequence_cxx17.cpp b/sycl/test/extensions/enqueue_functions/kernel_arg_sequence_cxx17.cpp index bbcb109f844dd..0e9ca85b592d4 100644 --- a/sycl/test/extensions/enqueue_functions/kernel_arg_sequence_cxx17.cpp +++ b/sycl/test/extensions/enqueue_functions/kernel_arg_sequence_cxx17.cpp @@ -9,9 +9,10 @@ // std::span, so they do not exist before C++20. A caller that passes its // argument list as a container there would otherwise have the container bound // as a single kernel argument, which compiles for any trivially copyable one -// and only fails once the kernel is launched, so it is diagnosed instead. +// and only fails once the kernel is launched, so it is diagnosed instead. Each +// of the five launch functions has to say so, hence the count. -// CHECK: Passing the arguments of a sycl::kernel as a sequence requires C++20 +// CHECK-COUNT-5: Passing the arguments of a sycl::kernel as a sequence requires C++20 #include @@ -21,7 +22,8 @@ namespace oneapiext = sycl::ext::oneapi::experimental; void argument_list_as_a_container(sycl::queue Q, sycl::handler &CGH, - sycl::nd_range<1> Range, + sycl::range<1> Range, + sycl::nd_range<1> NdRange, const sycl::kernel &Kernel) { int Value = 1; std::vector Vector{{&Value, sizeof(Value)}}; @@ -30,18 +32,32 @@ void argument_list_as_a_container(sycl::queue Q, sycl::handler &CGH, sycl::span Span{Vector.data(), Vector.size()}; - oneapiext::nd_launch(Q, Range, Kernel, SEQUENCE); - oneapiext::nd_launch(CGH, Range, Kernel, SEQUENCE); + oneapiext::single_task(Q, Kernel, SEQUENCE); + oneapiext::single_task(CGH, Kernel, SEQUENCE); + oneapiext::parallel_for(Q, Range, Kernel, SEQUENCE); + oneapiext::parallel_for(CGH, Range, Kernel, SEQUENCE); + oneapiext::parallel_for(Q, oneapiext::launch_config{Range}, Kernel, SEQUENCE); + oneapiext::parallel_for(CGH, oneapiext::launch_config{Range}, Kernel, + SEQUENCE); + oneapiext::nd_launch(Q, NdRange, Kernel, SEQUENCE); + oneapiext::nd_launch(CGH, NdRange, Kernel, SEQUENCE); + oneapiext::nd_launch(Q, oneapiext::launch_config{NdRange}, Kernel, SEQUENCE); + oneapiext::nd_launch(CGH, oneapiext::launch_config{NdRange}, Kernel, + SEQUENCE); } void one_argument_at_a_time(sycl::queue Q, sycl::handler &CGH, - sycl::nd_range<1> Range, + sycl::range<1> Range, sycl::nd_range<1> NdRange, const sycl::kernel &Kernel) { // A single raw_kernel_arg and typed arguments are one argument each, so they // are unaffected by the standard in use. int Value = 1; int *Pointer = nullptr; - oneapiext::nd_launch(Q, Range, Kernel, + oneapiext::single_task(Q, Kernel, + oneapiext::raw_kernel_arg{&Value, sizeof(Value)}); + oneapiext::parallel_for(CGH, Range, Kernel, Pointer, Value); + oneapiext::nd_launch(Q, NdRange, Kernel, oneapiext::raw_kernel_arg{&Value, sizeof(Value)}); - oneapiext::nd_launch(CGH, Range, Kernel, Pointer, Value); + oneapiext::nd_launch(CGH, oneapiext::launch_config{NdRange}, Kernel, Pointer, + Value); } diff --git a/sycl/test/extensions/enqueue_functions/kernel_arg_span.cpp b/sycl/test/extensions/enqueue_functions/kernel_arg_span.cpp index 7b4f8ac8670da..48f65a4ce8f7a 100644 --- a/sycl/test/extensions/enqueue_functions/kernel_arg_span.cpp +++ b/sycl/test/extensions/enqueue_functions/kernel_arg_span.cpp @@ -2,13 +2,14 @@ // RUN: %clangxx -fsycl -fsyntax-only -std=c++20 %s // An argument list held in a container converts to the std::span that the -// nd_launch sequence overloads take, but a parameter pack is an exact match and -// wins overload resolution. Without the forwarding the pack overloads do, the +// sequence 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 only fails once the kernel // is launched, with a fault or an adapter error that says nothing about the -// argument list. Check that every spelling of an argument list is accepted, and -// that a single raw_kernel_arg is still one argument. +// argument list. Check that every spelling of an argument list is accepted by +// every launch function that takes the arguments of a sycl::kernel, and that a +// single raw_kernel_arg is still one argument. // // The sequence overloads take a std::span, so they exist only in C++20 and // later; the first RUN line checks that the parameter pack overloads are @@ -23,9 +24,45 @@ #endif namespace oneapiext = sycl::ext::oneapi::experimental; +namespace ext_detail = sycl::ext::oneapi::experimental::detail; +using oneapiext::raw_kernel_arg; + +// Pins which single argument a parameter pack overload takes to be the argument +// list: a container of raw_kernel_arg in any value category, which is forwarded +// in C++20 and diagnosed before, but never one raw_kernel_arg. +template +constexpr bool IsArgList = +#if __cpp_lib_span + ext_detail::is_arg_list_container_v; +#else + ext_detail::is_arg_list_sequence_v; +#endif + +static_assert(IsArgList &>); +static_assert(IsArgList &>); +static_assert(IsArgList>); +static_assert(IsArgList &>); +static_assert(IsArgList &>); +static_assert(IsArgList>); +static_assert(IsArgList); +static_assert(IsArgList); +static_assert(IsArgList>); +static_assert(IsArgList>); +#if __cpp_lib_span +static_assert(IsArgList>); +static_assert(IsArgList>); +#endif +static_assert(!IsArgList); +static_assert(!IsArgList); +static_assert(!IsArgList); +static_assert(!IsArgList &>); +static_assert(!IsArgList); +// Two containers are two kernel arguments, not an argument list. +static_assert( + !IsArgList &, std::vector &>); void argument_list_spellings(sycl::queue Q, sycl::handler &CGH, - sycl::nd_range<1> Range, + sycl::range<1> Range, sycl::nd_range<1> NdRange, const sycl::kernel &Kernel) { int Value = 1; #if __cpp_lib_span @@ -40,24 +77,85 @@ void argument_list_spellings(sycl::queue Q, sycl::handler &CGH, sycl::span SyclSpan{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(Q, Range, Kernel, SyclSpan); - 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); - oneapiext::nd_launch(CGH, Range, Kernel, SyclSpan); + oneapiext::single_task(Q, Kernel, Vector); + oneapiext::single_task(Q, Kernel, Array); + oneapiext::single_task(Q, Kernel, Span); + oneapiext::single_task(Q, Kernel, MutableSpan); + oneapiext::single_task(Q, Kernel, SyclSpan); + oneapiext::single_task(CGH, Kernel, Vector); + oneapiext::single_task(CGH, Kernel, Array); + oneapiext::single_task(CGH, Kernel, Span); + oneapiext::single_task(CGH, Kernel, MutableSpan); + oneapiext::single_task(CGH, Kernel, SyclSpan); + + oneapiext::parallel_for(Q, Range, Kernel, Vector); + oneapiext::parallel_for(Q, Range, Kernel, Array); + oneapiext::parallel_for(Q, Range, Kernel, Span); + oneapiext::parallel_for(Q, Range, Kernel, MutableSpan); + oneapiext::parallel_for(Q, Range, Kernel, SyclSpan); + oneapiext::parallel_for(CGH, Range, Kernel, Vector); + oneapiext::parallel_for(CGH, Range, Kernel, Array); + oneapiext::parallel_for(CGH, Range, Kernel, Span); + oneapiext::parallel_for(CGH, Range, Kernel, MutableSpan); + oneapiext::parallel_for(CGH, Range, Kernel, SyclSpan); + + oneapiext::parallel_for(Q, oneapiext::launch_config{Range}, Kernel, Vector); + oneapiext::parallel_for(Q, oneapiext::launch_config{Range}, Kernel, Array); + oneapiext::parallel_for(Q, oneapiext::launch_config{Range}, Kernel, Span); + oneapiext::parallel_for(Q, oneapiext::launch_config{Range}, Kernel, + MutableSpan); + oneapiext::parallel_for(Q, oneapiext::launch_config{Range}, Kernel, SyclSpan); + oneapiext::parallel_for(CGH, oneapiext::launch_config{Range}, Kernel, Vector); + oneapiext::parallel_for(CGH, oneapiext::launch_config{Range}, Kernel, Array); + oneapiext::parallel_for(CGH, oneapiext::launch_config{Range}, Kernel, Span); + oneapiext::parallel_for(CGH, oneapiext::launch_config{Range}, Kernel, + MutableSpan); + oneapiext::parallel_for(CGH, oneapiext::launch_config{Range}, Kernel, + SyclSpan); + + oneapiext::nd_launch(Q, NdRange, Kernel, Vector); + oneapiext::nd_launch(Q, NdRange, Kernel, Array); + oneapiext::nd_launch(Q, NdRange, Kernel, Span); + oneapiext::nd_launch(Q, NdRange, Kernel, MutableSpan); + oneapiext::nd_launch(Q, NdRange, Kernel, SyclSpan); + oneapiext::nd_launch(CGH, NdRange, Kernel, Vector); + oneapiext::nd_launch(CGH, NdRange, Kernel, Array); + oneapiext::nd_launch(CGH, NdRange, Kernel, Span); + oneapiext::nd_launch(CGH, NdRange, Kernel, MutableSpan); + oneapiext::nd_launch(CGH, NdRange, Kernel, SyclSpan); + + oneapiext::nd_launch(Q, oneapiext::launch_config{NdRange}, Kernel, Vector); + oneapiext::nd_launch(Q, oneapiext::launch_config{NdRange}, Kernel, Array); + oneapiext::nd_launch(Q, oneapiext::launch_config{NdRange}, Kernel, Span); + oneapiext::nd_launch(Q, oneapiext::launch_config{NdRange}, Kernel, + MutableSpan); + oneapiext::nd_launch(Q, oneapiext::launch_config{NdRange}, Kernel, SyclSpan); + oneapiext::nd_launch(CGH, oneapiext::launch_config{NdRange}, Kernel, Vector); + oneapiext::nd_launch(CGH, oneapiext::launch_config{NdRange}, Kernel, Array); + oneapiext::nd_launch(CGH, oneapiext::launch_config{NdRange}, Kernel, Span); + oneapiext::nd_launch(CGH, oneapiext::launch_config{NdRange}, Kernel, + MutableSpan); + oneapiext::nd_launch(CGH, oneapiext::launch_config{NdRange}, Kernel, + SyclSpan); #endif // One raw_kernel_arg is one argument, and typed arguments are unaffected. - oneapiext::nd_launch(Q, Range, Kernel, + int *Pointer = nullptr; + oneapiext::single_task(Q, Kernel, + oneapiext::raw_kernel_arg{&Value, sizeof(Value)}); + oneapiext::single_task(CGH, Kernel, Pointer, Value); + oneapiext::parallel_for(Q, Range, Kernel, + oneapiext::raw_kernel_arg{&Value, sizeof(Value)}); + oneapiext::parallel_for(CGH, Range, Kernel, Pointer, Value); + oneapiext::parallel_for(Q, oneapiext::launch_config{Range}, Kernel, + oneapiext::raw_kernel_arg{&Value, sizeof(Value)}); + oneapiext::parallel_for(CGH, oneapiext::launch_config{Range}, Kernel, Pointer, + Value); + oneapiext::nd_launch(Q, NdRange, Kernel, oneapiext::raw_kernel_arg{&Value, sizeof(Value)}); - oneapiext::nd_launch(CGH, Range, Kernel, + oneapiext::nd_launch(CGH, NdRange, Kernel, Pointer, Value); + oneapiext::nd_launch(Q, oneapiext::launch_config{NdRange}, 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); + oneapiext::nd_launch(CGH, oneapiext::launch_config{NdRange}, Kernel, Pointer, + Value); } From cff7c0481d045366eed2e127566ebf48c867cc8a Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Thu, 1 Oct 2026 07:59:34 +0000 Subject: [PATCH 23/24] Review changes --- .../sycl_ext_oneapi_raw_kernel_arg.asciidoc | 7 +++-- .../oneapi/experimental/enqueue_functions.hpp | 28 +++++++++---------- sycl/source/detail/queue_impl.cpp | 4 +-- 3 files changed, 19 insertions(+), 20 deletions(-) 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 5a47d4e8098b4..1f5a1540c5bec 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,8 +54,9 @@ specification.* == Backend support status -This extension is currently implemented in {dpcpp} for the Level Zero and OpenCL -backends. Attempting to use this extension on other backends may result in +This extension is currently implemented in {dpcpp} only for GPU devices and +only when using the Level Zero or OpenCL backends. 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. @@ -153,7 +154,7 @@ 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. The `bytes` must not be a pointer argument. +copyable and is not a pointer. _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` diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp index 086b94d4b4836..d1da282f94cd6 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -193,19 +193,20 @@ template inline constexpr bool is_arg_list_sequence_v = sizeof...(ArgsT) == 1 && (is_raw_kernel_arg_sequence_v> && ...); +#endif -// Rejects an argument list passed as a container before C++20. The handler -// overloads call it, and the queue overloads reach them for such a call. -// `LaunchT` is the range or launch configuration, `void` for `single_task`, so -// that each launch function called that way is diagnosed, not just the first. +// Rejects an argument list passed as a container before C++20. `LaunchT` is the +// range or launch configuration, `void` for `single_task`, so that each launch +// function reports its own error. template void diagnoseArgListSequence() { +#if !__cpp_lib_span static_assert(!is_arg_list_sequence_v, "Passing the arguments of a sycl::kernel as a sequence " "requires C++20, where the overloads taking a std::span of " "raw_kernel_arg are available. Compile with C++20 or pass the " "arguments as a parameter pack."); -} #endif +} template void submit_impl(const queue &Q, PropertiesT Props, CommandGroupFunc &&CGF, @@ -283,9 +284,7 @@ inline void single_task(handler &CGH, const kernel &KernelObj, template void single_task(handler &CGH, const kernel &KernelObj, ArgsT &&...Args) { -#if !__cpp_lib_span detail::diagnoseArgListSequence(); -#endif if constexpr (detail::is_arg_list_container_v) { #if __cpp_lib_span single_task(CGH, KernelObj, @@ -306,6 +305,7 @@ inline void single_task(queue Q, const kernel &KernelObj, template void single_task(queue Q, const kernel &KernelObj, ArgsT &&...Args) { + detail::diagnoseArgListSequence(); if constexpr (detail::is_arg_list_container_v) { #if __cpp_lib_span single_task(std::move(Q), KernelObj, @@ -461,9 +461,7 @@ void parallel_for(handler &CGH, range Range, template void parallel_for(handler &CGH, range Range, const kernel &KernelObj, ArgsT &&...Args) { -#if !__cpp_lib_span detail::diagnoseArgListSequence, ArgsT...>(); -#endif if constexpr (detail::is_arg_list_container_v) { #if __cpp_lib_span parallel_for(CGH, Range, KernelObj, @@ -486,6 +484,7 @@ void parallel_for(queue Q, range Range, const kernel &KernelObj, template void parallel_for(queue Q, range Range, const kernel &KernelObj, ArgsT &&...Args) { + detail::diagnoseArgListSequence, ArgsT...>(); if constexpr (detail::is_arg_list_container_v) { #if __cpp_lib_span parallel_for(std::move(Q), Range, KernelObj, @@ -530,10 +529,8 @@ template void parallel_for(handler &CGH, launch_config, Properties> Config, const kernel &KernelObj, ArgsT &&...Args) { -#if !__cpp_lib_span detail::diagnoseArgListSequence, Properties>, ArgsT...>(); -#endif if constexpr (detail::is_arg_list_container_v) { #if __cpp_lib_span parallel_for(CGH, Config, KernelObj, @@ -560,6 +557,8 @@ void parallel_for(queue Q, launch_config, Properties> Config, template void parallel_for(queue Q, launch_config, Properties> Config, const kernel &KernelObj, ArgsT &&...Args) { + detail::diagnoseArgListSequence, Properties>, + ArgsT...>(); if constexpr (detail::is_arg_list_container_v) { #if __cpp_lib_span parallel_for(std::move(Q), Config, KernelObj, @@ -672,9 +671,7 @@ void nd_launch(handler &CGH, nd_range Range, template void nd_launch(handler &CGH, nd_range Range, const kernel &KernelObj, ArgsT &&...Args) { -#if !__cpp_lib_span detail::diagnoseArgListSequence, ArgsT...>(); -#endif if constexpr (detail::is_arg_list_container_v) { #if __cpp_lib_span nd_launch(CGH, Range, KernelObj, @@ -697,6 +694,7 @@ void nd_launch(queue Q, nd_range Range, const kernel &KernelObj, template void nd_launch(queue Q, nd_range Range, const kernel &KernelObj, ArgsT &&...Args) { + detail::diagnoseArgListSequence, ArgsT...>(); // 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. @@ -758,10 +756,8 @@ template void nd_launch(handler &CGH, launch_config, Properties> Config, const kernel &KernelObj, ArgsT &&...Args) { -#if !__cpp_lib_span detail::diagnoseArgListSequence< launch_config, Properties>, ArgsT...>(); -#endif if constexpr (detail::is_arg_list_container_v) { #if __cpp_lib_span nd_launch(CGH, Config, KernelObj, @@ -788,6 +784,8 @@ void nd_launch(queue Q, launch_config, Properties> Config, template void nd_launch(queue Q, launch_config, Properties> Config, const kernel &KernelObj, ArgsT &&...Args) { + detail::diagnoseArgListSequence< + launch_config, Properties>, ArgsT...>(); if constexpr (detail::is_arg_list_container_v) { #if __cpp_lib_span nd_launch(std::move(Q), Config, KernelObj, diff --git a/sycl/source/detail/queue_impl.cpp b/sycl/source/detail/queue_impl.cpp index fd14b3cb6543a..62b1e25ac1bd6 100644 --- a/sycl/source/detail/queue_impl.cpp +++ b/sycl/source/detail/queue_impl.cpp @@ -887,7 +887,7 @@ namespace { // The two argument forms differ only in how one element yields the bytes to // bind and their kind: a raw_kernel_arg carries plain bytes unless it was built // as a pointer argument. -inline sycl::detail::KernelArgView +inline const sycl::detail::KernelArgView & makeKernelArgView(const sycl::detail::KernelArgView &Arg) { return Arg; } @@ -943,7 +943,7 @@ void queue_impl::submit_kernel_obj_direct_impl( // 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) { - const sycl::detail::KernelArgView View = makeKernelArgView(Args[I]); + const sycl::detail::KernelArgView &View = makeKernelArgView(Args[I]); // `ArgDesc` holds a `void *` because the kinds that carry an object // rather than bytes hand it out as a mutable pointer. These arguments are // bytes and this path only ever reads them, hence the cast. From a47e7535764cb709e8fe7f6855f91fe84b477a0f Mon Sep 17 00:00:00 2001 From: Mieszko Dziadowiec Date: Thu, 1 Oct 2026 10:22:43 +0000 Subject: [PATCH 24/24] comment reduction --- .../oneapi/experimental/enqueue_functions.hpp | 38 +++++--------- .../oneapi/experimental/raw_kernel_arg.hpp | 10 ---- sycl/include/sycl/handler.hpp | 2 - sycl/include/sycl/queue.hpp | 9 ++-- sycl/source/detail/queue_impl.cpp | 3 -- sycl/source/detail/queue_impl.hpp | 6 --- .../kernel_obj_span_with_kb.cpp | 18 ++----- .../nd_launch_kernel_obj_direct.cpp | 5 +- .../nd_launch_kernel_obj_direct_raw_ptr.cpp | 13 +---- .../nd_launch_kernel_obj_span.cpp | 49 +++++++++---------- .../update_with_raw_kernel_arg_pointer.cpp | 5 +- sycl/test-e2e/RawKernelArg/pointer_arg.cpp | 15 ++---- sycl/test/abi/layout_raw_kernel_arg.cpp | 6 +-- .../kernel_arg_sequence_cxx17.cpp | 12 ++--- .../enqueue_functions/kernel_arg_span.cpp | 25 +++------- .../raw_kernel_arg/pointer_arg_errors.cpp | 12 ++--- 16 files changed, 68 insertions(+), 160 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp index d1da282f94cd6..212f7a94eb1d4 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -147,37 +147,28 @@ sycl::detail::KernelArgView makeKernelArgView(const T &Arg) { } #if __cpp_lib_span -// Binds an argument list that is held as a sequence, for the overloads that -// enqueue through a handler. inline void setRawKernelArgs(handler &CGH, std::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. + // An lvalue would pick the generic set_arg and bind the object itself. for (size_t I = 0; I < Args.size(); ++I) CGH.set_arg(static_cast(I), raw_kernel_arg{Args[I]}); } #endif -// 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. Without `std::span` there is no overload -// taking a sequence, so no call can be one. #if __cpp_lib_span +// A single argument that converts to std::span is the argument list itself. The +// pack overload is an exact match for it, so it forwards such a call. template inline constexpr bool is_arg_list_container_v = sizeof...(ArgsT) == 1 && (std::is_convertible_v> && ...); #else +// Without std::span there is no sequence overload to forward to. template inline constexpr bool is_arg_list_container_v = false; #endif #if !__cpp_lib_span -// A contiguous sequence of `raw_kernel_arg`, recognized without `std::span` so -// that a caller passing its argument list as a container before C++20 is told -// so, rather than having the container bound as a single kernel argument, which -// compiles for any trivially copyable one. template inline constexpr bool is_raw_kernel_arg_sequence_v = std::is_array_v && @@ -195,9 +186,9 @@ inline constexpr bool is_arg_list_sequence_v = (is_raw_kernel_arg_sequence_v> && ...); #endif -// Rejects an argument list passed as a container before C++20. `LaunchT` is the -// range or launch configuration, `void` for `single_task`, so that each launch -// function reports its own error. +// Rejects an argument list passed as a sequence before C++20, where a trivially +// copyable one would otherwise bind as a single kernel argument. `LaunchT` +// separates the instantiations, so each launch function reports the error. template void diagnoseArgListSequence() { #if !__cpp_lib_span static_assert(!is_arg_list_sequence_v, @@ -695,19 +686,17 @@ template void nd_launch(queue Q, nd_range Range, const kernel &KernelObj, ArgsT &&...Args) { detail::diagnoseArgListSequence, ArgsT...>(); - // 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. + // A sequence of raw_kernel_arg goes to the std::span overload. + // The handler-less path only takes arguments that can be bound directly, + // anything else goes through the handler overload above. if constexpr (detail::is_arg_list_container_v) { #if __cpp_lib_span nd_launch(std::move(Q), Range, KernelObj, std::span{std::forward(Args)...}); #endif } else if constexpr ((detail::is_direct_kernel_arg_v && ...)) { - // The handler-less path only takes arguments that can be bound directly, - // anything else goes through the handler overload above. The array is one - // element longer than the pack so that a zero-argument kernel stays well - // formed. + // 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)...}; sycl::detail::tls_code_loc_t TlsCodeLocCapture{ @@ -731,9 +720,6 @@ void nd_launch(handler &CGH, nd_range Range, 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, std::span Args, 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 933a3550a1fa3..a29f1622d12de 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/raw_kernel_arg.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/raw_kernel_arg.hpp @@ -21,8 +21,6 @@ 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{}; @@ -31,12 +29,6 @@ class raw_kernel_arg { 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) {} @@ -44,8 +36,6 @@ class raw_kernel_arg { 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; diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 004c0cf77f6d2..8fc619c5979d3 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -603,8 +603,6 @@ class __SYCL_EXPORT handler { void setArgHelper(int ArgIndex, sycl::ext::oneapi::experimental::raw_kernel_arg &&Arg) { auto StoredArg = storeRawArg(Arg); - // 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); diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp index 24e32628a773c..daaed811e3973 100644 --- a/sycl/include/sycl/queue.hpp +++ b/sycl/include/sycl/queue.hpp @@ -73,17 +73,14 @@ class raw_kernel_arg; // 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. The arguments are -// passed as a pointer and a count rather than as a view type, so that the -// spelling the public overloads use is not part of the ABI. +// submit_kernel_direct_* does for kernel function objects. +// A pointer and a count rather than a span, so that the span type the public +// overloads use can change without an ABI break. void __SYCL_EXPORT submit_kernel_obj_direct_without_event_impl( const queue &Queue, const detail::nd_range_view &RangeView, const kernel &Kernel, const detail::KernelArgView *Args, size_t NumArgs, const detail::code_location &CodeLoc, bool IsTopCodeLoc); -// As above, for an argument list that is already a contiguous sequence of -// `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, const ext::oneapi::experimental::raw_kernel_arg *Args, diff --git a/sycl/source/detail/queue_impl.cpp b/sycl/source/detail/queue_impl.cpp index 62b1e25ac1bd6..217d556d0c037 100644 --- a/sycl/source/detail/queue_impl.cpp +++ b/sycl/source/detail/queue_impl.cpp @@ -884,9 +884,6 @@ 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 carries plain bytes unless it was built -// as a pointer argument. inline const sycl::detail::KernelArgView & makeKernelArgView(const sycl::detail::KernelArgView &Arg) { return Arg; diff --git a/sycl/source/detail/queue_impl.hpp b/sycl/source/detail/queue_impl.hpp index e62dc12b9ec76..bcf2bcc1ae35d 100644 --- a/sycl/source/detail/queue_impl.hpp +++ b/sycl/source/detail/queue_impl.hpp @@ -390,8 +390,6 @@ 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 carries its own kind. void submit_kernel_obj_direct_without_event( const detail::nd_range_view &RangeView, const std::shared_ptr &KernelImpl, @@ -1006,10 +1004,6 @@ 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, diff --git a/sycl/test-e2e/EnqueueFunctions/kernel_obj_span_with_kb.cpp b/sycl/test-e2e/EnqueueFunctions/kernel_obj_span_with_kb.cpp index 87a095f212c62..ee42e4f52676c 100644 --- a/sycl/test-e2e/EnqueueFunctions/kernel_obj_span_with_kb.cpp +++ b/sycl/test-e2e/EnqueueFunctions/kernel_obj_span_with_kb.cpp @@ -7,10 +7,7 @@ // Tests the single_task and parallel_for overloads, and the launch_config forms // of parallel_for and nd_launch, that take the arguments of a sycl::kernel as a -// std::span of raw_kernel_arg, i.e. an argument list whose length is only known -// at run time. Those overloads take a std::span, hence C++20. They have to bind -// the same arguments in the same order as the parameter pack overloads, on the -// queue and on the handler alike. +// std::span of raw_kernel_arg. // NOTE: This relies on the availability of an OpenCL C compiler. #include @@ -53,8 +50,6 @@ int main() { int Count = N; int Value = 0; - // A pointer argument has to say that it is one, since the byte form of a - // pointer must not be passed to the byte overload of raw_kernel_arg. std::vector SingleTaskArgs; SingleTaskArgs.emplace_back(&Count, sizeof(Count)); SingleTaskArgs.emplace_back(&Value, sizeof(Value)); @@ -75,9 +70,7 @@ int main() { int Failed = 0; - // Each launch below passes a different value, so a launch that binds a value - // left over from an earlier one, or binds the arguments in the wrong order, - // shows up as a wrong result. + // A different value per launch, so a stale or misordered argument shows up. auto Run = [&](int Written, const char *Name, auto &&Launch) { Q.memset(Memory, 0, N * sizeof(int)); Value = Written; @@ -96,12 +89,11 @@ int main() { }); }); - // The container holding the arguments converts to that span, so passing it - // has to bind the arguments it holds rather than the container object. + // A container converts to std::span. Run(43, "single_task argument list passed as a container", [&] { oneapiext::single_task(Q, KernelSingleTask, SingleTaskArgs); }); - // The parameter pack overload has to agree element for element. + // The same arguments through the parameter pack overload. Run(44, "single_task parameter pack overload", [&] { oneapiext::single_task( Q, KernelSingleTask, oneapiext::raw_kernel_arg{&Count, sizeof(Count)}, @@ -166,5 +158,5 @@ int main() { }); sycl::free(Memory, Q); - return Failed; + return Failed != 0; } 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 299023f94b896..191c2d093808e 100644 --- a/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct.cpp +++ b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_direct.cpp @@ -78,9 +78,8 @@ 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. 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. + // knows the signature as sizes has to pass them. The pointer stays typed; + // nd_launch_kernel_obj_direct_raw_ptr.cpp passes it as a raw_kernel_arg. 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 4d1f4c111161a..cf94a453ce31a 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 @@ -3,16 +3,7 @@ // RUN: %{run} %t.out // Tests the nd_launch overload that takes an already built sycl::kernel with -// 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. -// -// 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. +// every argument passed through raw_kernel_arg, the USM pointer included. #include #include @@ -63,5 +54,5 @@ int main() { Check(Memory, 30, I, "pointer and scalars through raw_kernel_arg"); sycl::free(Memory, Q); - return Failed; + return Failed != 0; } 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 1169ce26a8128..81bb9c27a4751 100644 --- a/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_span.cpp +++ b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_span.cpp @@ -6,10 +6,7 @@ // RUN: %{run} %t.out // Tests the nd_launch overloads that take the arguments of a sycl::kernel as a -// std::span of raw_kernel_arg, i.e. an argument list whose length is only known -// at run time. Those overloads take a std::span, hence C++20. They have to bind -// the same arguments in the same order as the parameter pack overloads, on the -// queue and on the handler alike. +// std::span of raw_kernel_arg. #include @@ -23,6 +20,8 @@ #include "common.hpp" +#include +#include #include namespace syclext = sycl::ext::oneapi; @@ -33,8 +32,7 @@ static_assert(SYCL_EXT_ONEAPI_ENQUEUE_FUNCTIONS >= 2, 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. +// A mixture of argument sizes, so that a wrong size or order shows up. 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(); @@ -67,10 +65,6 @@ int main() { 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. A pointer argument has to say that it is one, since the byte - // form of a pointer must not be passed to the byte overload of - // raw_kernel_arg. std::vector Args; Args.emplace_back(&Memory, oneapiext::pointer_arg); Args.emplace_back(&A, sizeof(A)); @@ -81,8 +75,8 @@ int main() { 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. + // Several launches through one span before a wait, each of which has to bind + // the same arguments. constexpr int Launches = 8; Q.memset(Memory, 0, N * sizeof(int)); for (int I = 0; I < Launches; ++I) @@ -91,7 +85,7 @@ int main() { 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. + // The same arguments through the parameter pack overload. Q.memset(Memory, 0, N * sizeof(int)); oneapiext::nd_launch( Q, Ndr, Kernel, @@ -104,7 +98,6 @@ int main() { 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); @@ -112,8 +105,7 @@ 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. + // A container converts to std::span. Q.memset(Memory, 0, N * sizeof(int)); oneapiext::nd_launch(Q, Ndr, Kernel, Args); Q.wait(); @@ -127,7 +119,6 @@ int main() { 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, @@ -136,8 +127,7 @@ int main() { 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. + // A one element span is still an argument list, not a single argument. std::vector OneArg{ {&Memory, oneapiext::pointer_arg}}; Q.memset(Memory, 0, N * sizeof(int)); @@ -148,21 +138,28 @@ int main() { 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. + // An unfinished host task keeps the launch off the scheduler bypass, so the + // arguments have to be copied; they change before the kernel runs. { - sycl::buffer Buf{sycl::range<1>{N}}; + std::atomic Release{false}; 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); + CGH.host_task([&] { + while (!Release.load()) + std::this_thread::yield(); + }); }); oneapiext::nd_launch(Q, Ndr, Kernel, ArgSpan); + A = 0; + B = 0; + C = 0.0f; + D = 0; + Release = true; Q.wait(); for (size_t I = 0; I < N; ++I) - Failed += Check(Memory, Sum, I, "span form behind a buffer dependency"); + Failed += Check(Memory, Sum, I, "span form behind a host task"); } sycl::free(Memory, Q); - return Failed; + return Failed != 0; } 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 index e899c9d53c0fd..42acdd50b2350 100644 --- 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 @@ -5,8 +5,7 @@ // 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. +// Tests updating a raw_kernel_arg built as a pointer argument. #include "../graph_common.hpp" @@ -52,7 +51,7 @@ int main() { auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{}); - // PtrA is the one the pointer argument was built from. + // Before the update the kernel writes to PtrA only. Queue.ext_oneapi_graph(ExecGraph).wait(); std::vector HostDataA(N); diff --git a/sycl/test-e2e/RawKernelArg/pointer_arg.cpp b/sycl/test-e2e/RawKernelArg/pointer_arg.cpp index e0e894e69ac92..487ec5aca0e0d 100644 --- a/sycl/test-e2e/RawKernelArg/pointer_arg.cpp +++ b/sycl/test-e2e/RawKernelArg/pointer_arg.cpp @@ -4,13 +4,7 @@ // 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. +// Tests the pointer form of raw_kernel_arg. #include #include @@ -40,9 +34,6 @@ int main() { 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)}); @@ -52,8 +43,8 @@ int main() { 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. + // Again through set_args and with a different allocation, so that a stale + // pointer would show up. int *Other = sycl::malloc_shared(N, Q); In = 7; Q.memset(Other, 0xFF, N * sizeof(int)).wait(); diff --git a/sycl/test/abi/layout_raw_kernel_arg.cpp b/sycl/test/abi/layout_raw_kernel_arg.cpp index 0b34c35316a98..a51288900fa38 100644 --- a/sycl/test/abi/layout_raw_kernel_arg.cpp +++ b/sycl/test/abi/layout_raw_kernel_arg.cpp @@ -8,10 +8,8 @@ #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. +// The library reads the members of raw_kernel_arg, so its layout is pinned. +// MIsPointer is last so the other members 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 diff --git a/sycl/test/extensions/enqueue_functions/kernel_arg_sequence_cxx17.cpp b/sycl/test/extensions/enqueue_functions/kernel_arg_sequence_cxx17.cpp index 0e9ca85b592d4..b448eeb594aa3 100644 --- a/sycl/test/extensions/enqueue_functions/kernel_arg_sequence_cxx17.cpp +++ b/sycl/test/extensions/enqueue_functions/kernel_arg_sequence_cxx17.cpp @@ -5,12 +5,9 @@ // RUN: %clangxx -fsycl -fsyntax-only -std=c++20 -DSEQUENCE=Array %s // RUN: %clangxx -fsycl -fsyntax-only -std=c++20 -DSEQUENCE=Span %s -// The overloads that take the arguments of a sycl::kernel as a sequence take a -// std::span, so they do not exist before C++20. A caller that passes its -// argument list as a container there would otherwise have the container bound -// as a single kernel argument, which compiles for any trivially copyable one -// and only fails once the kernel is launched, so it is diagnosed instead. Each -// of the five launch functions has to say so, hence the count. +// Before C++20 a trivially copyable sequence of raw_kernel_arg would bind as +// one kernel argument, so every launch function rejects a sequence. The queue +// and handler forms share an instantiation, hence five errors for ten calls. // CHECK-COUNT-5: Passing the arguments of a sycl::kernel as a sequence requires C++20 @@ -49,8 +46,7 @@ void argument_list_as_a_container(sycl::queue Q, sycl::handler &CGH, void one_argument_at_a_time(sycl::queue Q, sycl::handler &CGH, sycl::range<1> Range, sycl::nd_range<1> NdRange, const sycl::kernel &Kernel) { - // A single raw_kernel_arg and typed arguments are one argument each, so they - // are unaffected by the standard in use. + // A single raw_kernel_arg and typed arguments are still accepted. int Value = 1; int *Pointer = nullptr; oneapiext::single_task(Q, Kernel, diff --git a/sycl/test/extensions/enqueue_functions/kernel_arg_span.cpp b/sycl/test/extensions/enqueue_functions/kernel_arg_span.cpp index 48f65a4ce8f7a..ce633740a1885 100644 --- a/sycl/test/extensions/enqueue_functions/kernel_arg_span.cpp +++ b/sycl/test/extensions/enqueue_functions/kernel_arg_span.cpp @@ -1,19 +1,8 @@ // 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 std::span that the -// sequence 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 only fails once the kernel -// is launched, with a fault or an adapter error that says nothing about the -// argument list. Check that every spelling of an argument list is accepted by -// every launch function that takes the arguments of a sycl::kernel, and that a -// single raw_kernel_arg is still one argument. -// -// The sequence overloads take a std::span, so they exist only in C++20 and -// later; the first RUN line checks that the parameter pack overloads are -// unaffected without it. +// From C++20 the parameter pack overloads forward a sequence of raw_kernel_arg +// to the std::span overloads. #include @@ -27,9 +16,8 @@ namespace oneapiext = sycl::ext::oneapi::experimental; namespace ext_detail = sycl::ext::oneapi::experimental::detail; using oneapiext::raw_kernel_arg; -// Pins which single argument a parameter pack overload takes to be the argument -// list: a container of raw_kernel_arg in any value category, which is forwarded -// in C++20 and diagnosed before, but never one raw_kernel_arg. +// Both standards have to agree on what an argument list is: a sequence of +// raw_kernel_arg in any value category, never a single raw_kernel_arg. template constexpr bool IsArgList = #if __cpp_lib_span @@ -57,7 +45,7 @@ static_assert(!IsArgList); static_assert(!IsArgList); static_assert(!IsArgList &>); static_assert(!IsArgList); -// Two containers are two kernel arguments, not an argument list. +// Two sequences do not form an argument list. static_assert( !IsArgList &, std::vector &>); @@ -72,8 +60,7 @@ void argument_list_spellings(sycl::queue Q, sycl::handler &CGH, std::span Span{Vector.data(), Vector.size()}; std::span MutableSpan{Vector.data(), Vector.size()}; - // A sycl::span still converts, so a caller holding one does not have to - // change how it stores its arguments. + // A sycl::span converts to std::span too. sycl::span SyclSpan{Vector.data(), Vector.size()}; diff --git a/sycl/test/extensions/raw_kernel_arg/pointer_arg_errors.cpp b/sycl/test/extensions/raw_kernel_arg/pointer_arg_errors.cpp index eaad709fe3111..3c57f02a0a3b3 100644 --- a/sycl/test/extensions/raw_kernel_arg/pointer_arg_errors.cpp +++ b/sycl/test/extensions/raw_kernel_arg/pointer_arg_errors.cpp @@ -1,9 +1,6 @@ // 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. +// The pointer_arg constructor takes the address of a pointer, not its value. #include @@ -11,17 +8,16 @@ namespace oneapiext = sycl::ext::oneapi::experimental; -// The graph extension copies a raw_kernel_arg as bytes. +// dynamic_parameter 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. + // Pointers to const and to void are accepted without a cast too. 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. + // An address and a size still select the byte form. oneapiext::raw_kernel_arg Bytes{&Ptr, sizeof(Ptr)}; // expected-error@+1 {{no matching constructor for initialization of 'oneapiext::raw_kernel_arg'}}