[SYCL] Add nd_launch overloads taking sycl::kernel arguments as a std::span

sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp

@@ -12,6 +12,10 @@
 #include <type_traits>
 #include <utility>
 
+#if __has_include(<span>)
+#include <span>
+#endif
+
 #include <sycl/detail/common.hpp>
 #include <sycl/event.hpp>
 #include <sycl/ext/oneapi/experimental/enqueue_types.hpp>
@@ -219,11 +223,65 @@ sycl::detail::KernelArgView makeKernelArgView(const T &Arg) {
   using ArgT = unqualified_arg_t<T>;
   if constexpr (std::is_same_v<ArgT, raw_kernel_arg>)
     return {RawKernelArgAccess::getData(Arg), RawKernelArgAccess::getSize(Arg),
-            kernel_param_kind_t::kind_std_layout};
+            RawKernelArgAccess::isPointer(Arg)
+                ? kernel_param_kind_t::kind_pointer
+                : kernel_param_kind_t::kind_std_layout};
   else
     return {&Arg, sizeof(ArgT), kernel_arg_kind_v<T>};
 }
 
+#if __cpp_lib_span
+inline void setRawKernelArgs(handler &CGH,
+                             std::span<const raw_kernel_arg> 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<int>(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 <typename... ArgsT>
+inline constexpr bool is_arg_list_container_v =
+    sizeof...(ArgsT) == 1 &&
+    (std::is_convertible_v<ArgsT, std::span<const raw_kernel_arg>> && ...);
+#else
+// Without std::span there is no sequence overload to forward to.
+template <typename... ArgsT>
+inline constexpr bool is_arg_list_container_v = false;
+#endif
+
+#if !__cpp_lib_span
+template <typename T, typename = void>
+inline constexpr bool is_raw_kernel_arg_sequence_v =
+    std::is_array_v<T> &&
+    std::is_same_v<std::remove_all_extents_t<T>, raw_kernel_arg>;
+template <typename T>
+inline constexpr bool is_raw_kernel_arg_sequence_v<
+    T, std::void_t<decltype(std::declval<const T &>().size()),
+                   decltype(std::declval<const T &>().data())>> =
+    std::is_convertible_v<decltype(std::declval<const T &>().data()),
+                          const raw_kernel_arg *>;
+
+template <typename... ArgsT>
+inline constexpr bool is_arg_list_sequence_v =
+    sizeof...(ArgsT) == 1 &&
+    (is_raw_kernel_arg_sequence_v<unqualified_arg_t<ArgsT>> && ...);
+#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.
+template <typename... ArgsT> void diagnoseArgListSequence() {
+#if !__cpp_lib_span
+  static_assert(!is_arg_list_sequence_v<ArgsT...>,
+                "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 <typename CommandGroupFunc, typename PropertiesT>
 void submit_impl(const queue &Q, PropertiesT Props, CommandGroupFunc &&CGF,
                  const sycl::detail::code_location &CodeLoc) {
@@ -573,16 +631,53 @@ nd_launch(queue Q, launch_config<nd_range<Dimensions>, Properties> Config,
   }
 }
 
+#if __cpp_lib_span
+template <int Dimensions>
+void nd_launch(handler &CGH, nd_range<Dimensions> Range,
+               const kernel &KernelObj, std::span<const raw_kernel_arg> Args);
+#endif
+
 template <int Dimensions, typename... ArgsT>
 void nd_launch(handler &CGH, nd_range<Dimensions> Range,
                const kernel &KernelObj, ArgsT &&...Args) {
-  CGH.set_args<ArgsT...>(std::forward<ArgsT>(Args)...);
-  CGH.parallel_for(Range, KernelObj);
+  detail::diagnoseArgListSequence<ArgsT...>();
+  if constexpr (detail::is_arg_list_container_v<ArgsT...>) {
+#if __cpp_lib_span
+    nd_launch(CGH, Range, KernelObj,
+              std::span<const raw_kernel_arg>{std::forward<ArgsT>(Args)...});
+#endif
+  } else {
+    CGH.set_args<ArgsT...>(std::forward<ArgsT>(Args)...);
+    CGH.parallel_for(Range, KernelObj);
+  }
 }
 
+#if __cpp_lib_span
+template <int Dimensions>
+void nd_launch(queue Q, nd_range<Dimensions> Range, const kernel &KernelObj,
+               std::span<const raw_kernel_arg> 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 <
+    int Dimensions, typename ArgListT,
+    typename = std::enable_if_t<detail::is_arg_list_container_v<ArgListT>>>
+void nd_launch(queue Q, nd_range<Dimensions> 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<const raw_kernel_arg>{std::forward<ArgListT>(Args)},
+            CodeLoc);
+}
+#endif
+
 template <int Dimensions, typename... ArgsT>
 void nd_launch(queue Q, nd_range<Dimensions> 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<ArgsT> && ...)) {
@@ -593,8 +688,8 @@ void nd_launch(queue Q, nd_range<Dimensions> 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,6 +698,25 @@ void nd_launch(queue Q, nd_range<Dimensions> Range, const kernel &KernelObj,
   }
 }
 
+#if __cpp_lib_span
+template <int Dimensions>
+void nd_launch(handler &CGH, nd_range<Dimensions> Range,
+               const kernel &KernelObj, std::span<const raw_kernel_arg> Args) {
+  detail::setRawKernelArgs(CGH, Args);
+  CGH.parallel_for(Range, KernelObj);
+}
+
+template <int Dimensions>
+void nd_launch(queue Q, nd_range<Dimensions> Range, const kernel &KernelObj,
+               std::span<const raw_kernel_arg> 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
+
 template <int Dimensions, typename Properties, typename... ArgsT>
 void nd_launch(handler &CGH,
                launch_config<nd_range<Dimensions>, Properties> Config,

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 <typename T>
+  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
 

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

sycl/include/sycl/queue.hpp

@@ -69,14 +69,25 @@ template <backend BackendName, class SyclObjectT>
 auto get_native(const SyclObjectT &Obj)
     -> backend_return_t<BackendName, SyclObjectT>;
 
+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<const detail::KernelArgView> 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,

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 <typename ArgT>
+void queue_impl::submit_kernel_obj_direct_impl(
     const detail::nd_range_view &RangeView,
     const std::shared_ptr<detail::kernel_impl> &KernelImpl,
-    sycl::span<const sycl::detail::KernelArgView> Args,
-    const detail::code_location &CodeLoc, bool IsTopCodeLoc) {
+    sycl::span<const ArgT> 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<void *>(Args[I].MPtr);
+      void *Value = const_cast<void *>(View.MPtr);
       if (!SchedulerBypass) {
-        const char *Bytes = static_cast<const char *>(Args[I].MPtr);
-        CGData.MArgsStorage.emplace_back(Bytes, Bytes + Args[I].MSize);
+        const char *Bytes = static_cast<const char *>(View.MPtr);
+        CGData.MArgsStorage.emplace_back(Bytes, Bytes + View.MSize);
         Value = CGData.MArgsStorage.back().data();
       }
-      KData.addArg(Args[I].MKind, Value, static_cast<int>(Args[I].MSize),
+      KData.addArg(View.MKind, Value, static_cast<int>(View.MSize),
                    static_cast<int>(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<detail::kernel_impl> &KernelImpl,
+    sycl::span<const sycl::detail::KernelArgView> 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<detail::kernel_impl> &KernelImpl,
+    sycl::span<const ext::oneapi::experimental::raw_kernel_arg> 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<ext::oneapi::experimental::detail::exec_graph_impl>
         ExecGraph,

sycl/source/detail/queue_impl.hpp

@@ -391,6 +391,12 @@ class queue_impl : public std::enable_shared_from_this<queue_impl> {
       sycl::span<const sycl::detail::KernelArgView> 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<detail::kernel_impl> &KernelImpl,
+      sycl::span<const ext::oneapi::experimental::raw_kernel_arg> Args,
+      const detail::code_location &CodeLoc, bool IsTopCodeLoc);
+
   event submit_barrier_direct_with_event(sycl::span<const event> DepEvents,
                                          detail::CGType BarrierType,
                                          const detail::code_location &CodeLoc) {
@@ -1063,6 +1069,13 @@ class queue_impl : public std::enable_shared_from_this<queue_impl> {
                 SubmitCommandFuncType &SubmitCommandFunc, detail::CGType Type,
                 bool InsertBarrierForInOrderCommand);
 
+  template <typename ArgT>
+  void submit_kernel_obj_direct_impl(
+      const detail::nd_range_view &RangeView,
+      const std::shared_ptr<detail::kernel_impl> &KernelImpl,
+      sycl::span<const ArgT> Args, const detail::code_location &CodeLoc,
+      bool IsTopCodeLoc);
+
   /// Performs barrier submission to the queue.
   ///
   /// \param DepEvents is a vector of dependencies of the operation.

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

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<const detail::KernelArgView> 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(

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,