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..310edd3486fbf 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,12 @@ 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 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 @@ -165,8 +171,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 launch function overloads that take the arguments of a + `sycl::kernel` as a `std::span` of `raw_kernel_arg`. |=== @@ -301,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]. +_{endnote}_] + |==== @@ -397,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. @@ -409,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]. +_{endnote}_] + +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]. +_{endnote}_] + |==== @@ -512,6 +623,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 @@ -525,24 +639,89 @@ a! ---- namespace sycl::ext::oneapi::experimental { +template // Requires C++20 +void nd_launch(sycl::queue q, sycl::nd_range r, + const sycl::kernel& k, std::span args); + +template // Requires C++20 +void nd_launch(sycl::handler &h, sycl::nd_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::nd_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]. +_{endnote}_] + +a| +[frame=all,grid=none] +!==== +a! +[source,c++] +---- +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]. +_{endnote}_] + |==== @@ -771,6 +950,47 @@ _Effects_: Submits an executable command graph to the `sycl::queue` or `sycl::ha |==== +== Example + +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; + +static constexpr size_t N = 1024; + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclex::nd_range_kernel<1>)) +void scale(float *data, int factor) { /* kernel body */ } + +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 . What should `submit_with_event` be called? 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..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 @@ -55,10 +55,10 @@ 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. +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. == Overview @@ -76,12 +76,27 @@ 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); +---- + +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: + +[source,c++] +---- +int* ptr = sycl::malloc_device(n, q); +... +h.set_arg(0, syclex::raw_kernel_arg(&ptr, syclex::pointer_arg)); h.parallel_for(range, kernel); ---- @@ -104,8 +119,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 +134,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 +154,32 @@ 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 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` 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. 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 +196,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 f8d69af1c4df9..66461e316828c 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -12,6 +12,10 @@ #include #include +#if __has_include() +#include +#endif + #include #include #include @@ -219,11 +223,66 @@ 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}; } +#if __cpp_lib_span +inline void setRawKernelArgs(handler &CGH, + std::span Args) { + // 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 + +#if __cpp_lib_span +// A single argument that converts to std::span is the argument list itself, +// which the parameter pack overloads would otherwise take as one argument. +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 +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 + +// 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, + "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, const sycl::detail::code_location &CodeLoc) { @@ -293,19 +352,68 @@ 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); + detail::diagnoseArgListSequence(); + 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()); + +// A sequence comes here rather than to the parameter pack overload below, +// which has no parameter to take the caller's code location. +template >> +void single_task(queue Q, const kernel &KernelObj, ArgListT &&Args, + const sycl::detail::code_location &CodeLoc = + sycl::detail::code_location::current()) { + single_task(std::move(Q), KernelObj, + std::span{std::forward(Args)}, + CodeLoc); +} +#endif + template void single_task(queue Q, const kernel &KernelObj, ArgsT &&...Args) { + detail::diagnoseArgListSequence(); 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 @@ -469,41 +577,161 @@ 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); + detail::diagnoseArgListSequence, ArgsT...>(); + 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()); + +template < + int Dimensions, typename ArgListT, + typename = std::enable_if_t>> +void parallel_for(queue Q, range Range, const kernel &KernelObj, + ArgListT &&Args, + const sycl::detail::code_location &CodeLoc = + sycl::detail::code_location::current()) { + parallel_for(std::move(Q), Range, KernelObj, + std::span{std::forward(Args)}, + CodeLoc); +} +#endif + template void parallel_for(queue Q, range Range, const kernel &KernelObj, ArgsT &&...Args) { + detail::diagnoseArgListSequence, ArgsT...>(); 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) { - ext::oneapi::experimental::detail::LaunchConfigAccess, - Properties> - ConfigAccess(Config); - CGH.set_args(std::forward(Args)...); - sycl::detail::HandlerAccess::parallelForImpl( - CGH, ConfigAccess.getRange(), ConfigAccess.getProperties(), KernelObj); + detail::diagnoseArgListSequence, Properties>, + ArgsT...>(); + 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()); + +template < + int Dimensions, typename Properties, typename ArgListT, + typename = std::enable_if_t>> +void parallel_for(queue Q, launch_config, Properties> Config, + const kernel &KernelObj, ArgListT &&Args, + const sycl::detail::code_location &CodeLoc = + sycl::detail::code_location::current()) { + parallel_for(std::move(Q), Config, KernelObj, + std::span{std::forward(Args)}, + CodeLoc); +} +#endif + template void parallel_for(queue Q, launch_config, Properties> Config, const kernel &KernelObj, ArgsT &&...Args) { + detail::diagnoseArgListSequence, Properties>, + ArgsT...>(); 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); + detail::setRawKernelArgs(CGH, Args); + sycl::detail::HandlerAccess::parallelForImpl( + CGH, ConfigAccess.getRange(), ConfigAccess.getProperties(), KernelObj); +} + +template +void parallel_for(queue Q, launch_config, Properties> Config, + 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 detail::enable_if_not_kernel_function_t @@ -573,16 +801,51 @@ nd_launch(queue Q, launch_config, Properties> Config, } } +#if __cpp_lib_span +template +void nd_launch(handler &CGH, nd_range Range, + const kernel &KernelObj, std::span Args); +#endif + 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); + detail::diagnoseArgListSequence, ArgsT...>(); + if constexpr (detail::is_arg_list_container_v) { +#if __cpp_lib_span + nd_launch(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 nd_launch(queue Q, nd_range Range, const kernel &KernelObj, + std::span Args, + const sycl::detail::code_location &CodeLoc = + sycl::detail::code_location::current()); + +template < + int Dimensions, typename ArgListT, + typename = std::enable_if_t>> +void nd_launch(queue Q, nd_range Range, const kernel &KernelObj, + ArgListT &&Args, + const sycl::detail::code_location &CodeLoc = + sycl::detail::code_location::current()) { + nd_launch(std::move(Q), Range, KernelObj, + std::span{std::forward(Args)}, + CodeLoc); +} +#endif + template void nd_launch(queue Q, nd_range Range, const kernel &KernelObj, ArgsT &&...Args) { + detail::diagnoseArgListSequence, ArgsT...>(); // 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 && ...)) { @@ -593,8 +856,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) { @@ -603,26 +866,106 @@ 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, std::span Args) { + detail::setRawKernelArgs(CGH, Args); + CGH.parallel_for(Range, KernelObj); +} + +template +void nd_launch(queue Q, nd_range Range, const kernel &KernelObj, + 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.data(), + Args.size(), TlsCodeLocCapture.query(), TlsCodeLocCapture.isToplevel()); +} +#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) { - ext::oneapi::experimental::detail::LaunchConfigAccess, - Properties> - ConfigAccess(Config); - CGH.set_args(std::forward(Args)...); - sycl::detail::HandlerAccess::parallelForImpl( - CGH, ConfigAccess.getRange(), ConfigAccess.getProperties(), KernelObj); + detail::diagnoseArgListSequence< + launch_config, Properties>, ArgsT...>(); + 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()); + +template < + int Dimensions, typename Properties, typename ArgListT, + typename = std::enable_if_t>> +void nd_launch(queue Q, launch_config, Properties> Config, + const kernel &KernelObj, ArgListT &&Args, + const sycl::detail::code_location &CodeLoc = + sycl::detail::code_location::current()) { + nd_launch(std::move(Q), Config, KernelObj, + std::span{std::forward(Args)}, + CodeLoc); } +#endif template void nd_launch(queue Q, launch_config, Properties> Config, const kernel &KernelObj, ArgsT &&...Args) { + detail::diagnoseArgListSequence< + launch_config, Properties>, ArgsT...>(); 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); + detail::setRawKernelArgs(CGH, Args); + sycl::detail::HandlerAccess::parallelForImpl( + CGH, ConfigAccess.getRange(), ConfigAccess.getProperties(), KernelObj); +} + +template +void nd_launch(queue Q, launch_config, Properties> Config, + 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 // kernel. Wrapping generated a second, duplicate device kernel and dropped the 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..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,14 +21,22 @@ class dynamic_parameter_impl; struct RawKernelArgAccess; } // namespace detail +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) {} + 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; + bool MIsPointer = false; friend class sycl::handler; // For sycl_ext_oneapi_graph integration @@ -42,6 +50,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 56539906b52d3..1be6bce3d1e4e 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -611,8 +611,9 @@ 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); + 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 3d721f815b6fe..7a88839d575bb 100644 --- a/sycl/include/sycl/queue.hpp +++ b/sycl/include/sycl/queue.hpp @@ -69,14 +69,25 @@ 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. +// 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, sycl::span Args, + const kernel &Kernel, const detail::KernelArgView *Args, size_t NumArgs, const detail::code_location &CodeLoc, bool IsTopCodeLoc); +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); + 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 1983fe81ab310..46350b12413ca 100644 --- a/sycl/source/detail/queue_impl.cpp +++ b/sycl/source/detail/queue_impl.cpp @@ -1013,11 +1013,28 @@ EventImplPtr queue_impl::submit_kernel_direct_impl( /*InsertBarrierForInOrderCommand*/ false); } -void queue_impl::submit_kernel_obj_direct_without_event( +namespace { +inline const 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), + syclex_detail::RawKernelArgAccess::isPointer(Arg) + ? sycl::detail::kernel_param_kind_t::kind_pointer + : 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()); @@ -1053,16 +1070,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)); } @@ -1104,6 +1122,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 ac5e0d28a93bf..f85ff72ae4372 100644 --- a/sycl/source/detail/queue_impl.hpp +++ b/sycl/source/detail/queue_impl.hpp @@ -391,6 +391,12 @@ class queue_impl : public std::enable_shared_from_this { sycl::span Args, const detail::code_location &CodeLoc, bool IsTopCodeLoc); + 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) { @@ -1063,6 +1069,13 @@ class queue_impl : public std::enable_shared_from_this { SubmitCommandFuncType &SubmitCommandFunc, detail::CGType Type, bool InsertBarrierForInOrderCommand); + 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 298f86f1e6196..ba48a8065a916 100644 --- a/sycl/source/feature_test.hpp.in +++ b/sycl/source/feature_test.hpp.in @@ -97,8 +97,8 @@ 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_RAW_KERNEL_ARG 1 +#define SYCL_EXT_ONEAPI_ENQUEUE_FUNCTIONS 2 +#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/source/queue.cpp b/sycl/source/queue.cpp index bdbf0b00b6fe1..af39545a259ba 100644 --- a/sycl/source/queue.cpp +++ b/sycl/source/queue.cpp @@ -321,10 +321,20 @@ 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, 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); } event submit_graph_direct_with_event_impl( 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..ee42e4f52676c --- /dev/null +++ b/sycl/test-e2e/EnqueueFunctions/kernel_obj_span_with_kb.cpp @@ -0,0 +1,162 @@ +// 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. +// 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; + + 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; + + // 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; + 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); + }); + }); + + // A container converts to std::span. + Run(43, "single_task argument list passed as a container", + [&] { oneapiext::single_task(Q, KernelSingleTask, SingleTaskArgs); }); + + // 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)}, + 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 != 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 32ce9adc31218..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,10 +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. 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. 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 3ff144334d385..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 @@ -1,19 +1,9 @@ // REQUIRES: aspect-usm_shared_allocations -// REQUIRES: level_zero // RUN: %{build} -o %t.out // RUN: %{run} %t.out // Tests the nd_launch overload that takes an already built sycl::kernel with -// every argument, the USM pointer included, passed as raw bytes. That is what a -// caller which only knows the signature as sizes has to do. -// -// A `raw_kernel_arg` is bound as UR_EXP_KERNEL_ARG_TYPE_VALUE, which for a -// pointer parameter is only equivalent to UR_EXP_KERNEL_ARG_TYPE_POINTER on -// Level Zero. The OpenCL adapter passes a value argument to clSetKernelArg, -// which rejects a USM pointer with CL_INVALID_MEM_OBJECT, and the Native CPU -// adapter puts the address of its own copy of the bytes into the argument slot -// instead of the pointer itself. Hence the Level Zero requirement above, the -// same restriction the RawKernelArg tests carry. +// every argument passed through raw_kernel_arg, the USM pointer included. #include #include @@ -53,14 +43,16 @@ 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; + 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 new file mode 100644 index 0000000000000..81bb9c27a4751 --- /dev/null +++ b/sycl/test-e2e/EnqueueFunctions/nd_launch_kernel_obj_span.cpp @@ -0,0 +1,165 @@ +// 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 nd_launch overloads that take the arguments of a sycl::kernel as a +// std::span of raw_kernel_arg. + +#include + +#include +#include +#include +#include +#include +#include +#include + +#include "common.hpp" + +#include +#include +#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 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(); + 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; + + std::vector Args; + 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)); + Args.emplace_back(&D, sizeof(D)); + std::span ArgSpan{Args.data(), Args.size()}; + + int Failed = 0; + + // 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) + 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 same arguments through the parameter pack overload. + Q.memset(Memory, 0, N * sizeof(int)); + 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"); + + 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 container converts to std::span. + 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"); + + Q.memset(Memory, 0, N * sizeof(int)); + oneapiext::nd_launch( + Q, Ndr, Kernel, + 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"); + + // 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)); + 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"); + + // An unfinished host task keeps the launch off the scheduler bypass, so the + // arguments have to be copied; they change before the kernel runs. + { + std::atomic Release{false}; + Q.memset(Memory, 0, N * sizeof(int)); + Q.submit([&](sycl::handler &CGH) { + 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 host task"); + } + + sycl::free(Memory, Q); + 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 new file mode 100644 index 0000000000000..42acdd50b2350 --- /dev/null +++ b/sycl/test-e2e/Graph/Update/update_with_raw_kernel_arg_pointer.cpp @@ -0,0 +1,82 @@ +// 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 built as a pointer argument. + +#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{}); + + // Before the update the kernel writes to PtrA only. + 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..487ec5aca0e0d --- /dev/null +++ b/sycl/test-e2e/RawKernelArg/pointer_arg.cpp @@ -0,0 +1,62 @@ +// 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. + +#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; + + 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); + + // 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(); + 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-e2e/Tracing/code_location_kernel_obj_sequence.cpp b/sycl/test-e2e/Tracing/code_location_kernel_obj_sequence.cpp new file mode 100644 index 0000000000000..2e749377e537d --- /dev/null +++ b/sycl/test-e2e/Tracing/code_location_kernel_obj_sequence.cpp @@ -0,0 +1,112 @@ +// REQUIRES: level_zero || opencl +// RUN: %{build} -std=c++20 -o %t.out +// RUN: %{run} sycl-trace --sycl --print-format=verbose %t.out | FileCheck %s + +// Tests that the launch functions report the caller's code location when the +// arguments of a sycl::kernel are passed as a sequence. + +#include + +#include +#include +#include +#include + +#include + +namespace syclex = sycl::ext::oneapi::experimental; + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclex::nd_range_kernel<1>)) +void takeValue(int) {} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclex::single_task_kernel)) +void takeValueOnce(int) {} + +template sycl::kernel getKernel(sycl::queue &Q) { + auto Bundle = syclex::get_kernel_bundle( + Q.get_context()); + return Bundle.template ext_oneapi_get_kernel(); +} + +int main() { + sycl::queue Q; + sycl::kernel Kernel = getKernel(Q); + sycl::kernel SingleTaskKernel = getKernel(Q); + sycl::nd_range<1> Range{1, 1}; + syclex::launch_config RangeConfig{sycl::range<1>{1}}; + syclex::launch_config NdRangeConfig{Range}; + + int Value = 1; + std::vector Args{{&Value, sizeof(Value)}}; + std::span Span{Args}; + + // CHECK: [SYCL] Task begin + // CHECK-DAG: sym_source_file_name : {{.*}}code_location_kernel_obj_sequence.cpp + // CHECK-DAG: sym_line_no : [[# @LINE + 2 ]] + // CHECK: [SYCL] Task end + syclex::nd_launch(Q, Range, Kernel, Args); + Q.wait(); + + // CHECK: [SYCL] Task begin + // CHECK-DAG: sym_source_file_name : {{.*}}code_location_kernel_obj_sequence.cpp + // CHECK-DAG: sym_line_no : [[# @LINE + 2 ]] + // CHECK: [SYCL] Task end + syclex::nd_launch(Q, Range, Kernel, Span); + Q.wait(); + + // CHECK: [SYCL] Task begin + // CHECK-DAG: sym_source_file_name : {{.*}}code_location_kernel_obj_sequence.cpp + // CHECK-DAG: sym_line_no : [[# @LINE + 2 ]] + // CHECK: [SYCL] Task end + syclex::single_task(Q, SingleTaskKernel, Args); + Q.wait(); + + // CHECK: [SYCL] Task begin + // CHECK-DAG: sym_source_file_name : {{.*}}code_location_kernel_obj_sequence.cpp + // CHECK-DAG: sym_line_no : [[# @LINE + 2 ]] + // CHECK: [SYCL] Task end + syclex::single_task(Q, SingleTaskKernel, Span); + Q.wait(); + + // CHECK: [SYCL] Task begin + // CHECK-DAG: sym_source_file_name : {{.*}}code_location_kernel_obj_sequence.cpp + // CHECK-DAG: sym_line_no : [[# @LINE + 2 ]] + // CHECK: [SYCL] Task end + syclex::parallel_for(Q, sycl::range<1>{1}, Kernel, Args); + Q.wait(); + + // CHECK: [SYCL] Task begin + // CHECK-DAG: sym_source_file_name : {{.*}}code_location_kernel_obj_sequence.cpp + // CHECK-DAG: sym_line_no : [[# @LINE + 2 ]] + // CHECK: [SYCL] Task end + syclex::parallel_for(Q, sycl::range<1>{1}, Kernel, Span); + Q.wait(); + + // CHECK: [SYCL] Task begin + // CHECK-DAG: sym_source_file_name : {{.*}}code_location_kernel_obj_sequence.cpp + // CHECK-DAG: sym_line_no : [[# @LINE + 2 ]] + // CHECK: [SYCL] Task end + syclex::parallel_for(Q, RangeConfig, Kernel, Args); + Q.wait(); + + // CHECK: [SYCL] Task begin + // CHECK-DAG: sym_source_file_name : {{.*}}code_location_kernel_obj_sequence.cpp + // CHECK-DAG: sym_line_no : [[# @LINE + 2 ]] + // CHECK: [SYCL] Task end + syclex::parallel_for(Q, RangeConfig, Kernel, Span); + Q.wait(); + + // CHECK: [SYCL] Task begin + // CHECK-DAG: sym_source_file_name : {{.*}}code_location_kernel_obj_sequence.cpp + // CHECK-DAG: sym_line_no : [[# @LINE + 2 ]] + // CHECK: [SYCL] Task end + syclex::nd_launch(Q, NdRangeConfig, Kernel, Args); + Q.wait(); + + // CHECK: [SYCL] Task begin + // CHECK-DAG: sym_source_file_name : {{.*}}code_location_kernel_obj_sequence.cpp + // CHECK-DAG: sym_line_no : [[# @LINE + 2 ]] + // CHECK: [SYCL] Task end + syclex::nd_launch(Q, NdRangeConfig, Kernel, Span); + Q.wait(); +} 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..a51288900fa38 --- /dev/null +++ b/sycl/test/abi/layout_raw_kernel_arg.cpp @@ -0,0 +1,19 @@ +// 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 + +// 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 +// 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/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 90f8c64e451d0..55aa12f1194d9 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3163,7 +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_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 e4674b15c4575..13753c9637dd2 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -4398,7 +4398,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@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..b448eeb594aa3 --- /dev/null +++ b/sycl/test/extensions/enqueue_functions/kernel_arg_sequence_cxx17.cpp @@ -0,0 +1,59 @@ +// 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 + +// 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 + +#include + +#include +#include + +namespace oneapiext = sycl::ext::oneapi::experimental; + +void argument_list_as_a_container(sycl::queue Q, sycl::handler &CGH, + sycl::range<1> Range, + sycl::nd_range<1> NdRange, + 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::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::range<1> Range, sycl::nd_range<1> NdRange, + const sycl::kernel &Kernel) { + // A single raw_kernel_arg and typed arguments are still accepted. + int Value = 1; + int *Pointer = nullptr; + 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, 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 new file mode 100644 index 0000000000000..ce633740a1885 --- /dev/null +++ b/sycl/test/extensions/enqueue_functions/kernel_arg_span.cpp @@ -0,0 +1,148 @@ +// RUN: %clangxx -fsycl -fsyntax-only %s +// RUN: %clangxx -fsycl -fsyntax-only -std=c++20 %s + +// From C++20 the parameter pack overloads forward a sequence of raw_kernel_arg +// to the std::span overloads. + +#include + +#include +#include +#if __cpp_lib_span +#include +#endif + +namespace oneapiext = sycl::ext::oneapi::experimental; +namespace ext_detail = sycl::ext::oneapi::experimental::detail; +using oneapiext::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 + 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 sequences do not form an argument list. +static_assert( + !IsArgList &, std::vector &>); + +void argument_list_spellings(sycl::queue Q, sycl::handler &CGH, + sycl::range<1> Range, sycl::nd_range<1> NdRange, + 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)}}; + std::span Span{Vector.data(), Vector.size()}; + std::span MutableSpan{Vector.data(), + Vector.size()}; + // A sycl::span converts to std::span too. + sycl::span SyclSpan{Vector.data(), + Vector.size()}; + + 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. + 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, NdRange, Kernel, Pointer, Value); + oneapiext::nd_launch(Q, oneapiext::launch_config{NdRange}, Kernel, + oneapiext::raw_kernel_arg{&Value, sizeof(Value)}); + oneapiext::nd_launch(CGH, oneapiext::launch_config{NdRange}, Kernel, Pointer, + Value); +} 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..3c57f02a0a3b3 --- /dev/null +++ b/sycl/test/extensions/raw_kernel_arg/pointer_arg_errors.cpp @@ -0,0 +1,33 @@ +// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s + +// The pointer_arg constructor takes the address of a pointer, not its value. + +#include + +#include + +namespace oneapiext = sycl::ext::oneapi::experimental; + +// 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) { + // 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}; + + // 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'}} + 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; +}