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