[SYCL][1/2] Submit an already built sycl::kernel without a handler

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 <sycl/detail/kernel_desc.hpp> // for kernel_param_kind_t
+
+#include <stddef.h> // 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

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

@@ -15,6 +15,7 @@
 #include <sycl/ext/oneapi/experimental/enqueue_types.hpp>
 #include <sycl/ext/oneapi/experimental/free_function_traits.hpp>
 #include <sycl/ext/oneapi/experimental/graph.hpp>
+#include <sycl/ext/oneapi/experimental/raw_kernel_arg.hpp>
 #include <sycl/ext/oneapi/properties.hpp>
 #include <sycl/handler.hpp>
 #include <sycl/nd_range.hpp>
@@ -98,6 +99,47 @@ template <typename LCRangeT, typename LCPropertiesT> struct LaunchConfigAccess {
   }
 };
 
+// The argument type as the kernel sees it.
+template <typename T>
+using unqualified_arg_t = std::remove_cv_t<std::remove_reference_t<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 <typename T>
+inline constexpr bool is_scalar_kernel_arg_v =
+    std::is_arithmetic_v<T> || std::is_enum_v<T> || std::is_pointer_v<T>;
+
+// 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 <typename T, typename ArgT = unqualified_arg_t<T>>
+inline constexpr bool is_direct_kernel_arg_v =
+    is_scalar_kernel_arg_v<ArgT> ||
+    (std::is_array_v<ArgT> &&
+     is_scalar_kernel_arg_v<std::remove_all_extents_t<ArgT>>) ||
+    std::is_same_v<ArgT, raw_kernel_arg>;
+
+// 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 <typename T, typename ArgT = unqualified_arg_t<T>>
+inline constexpr sycl::detail::kernel_param_kind_t kernel_arg_kind_v =
+    (std::is_pointer_v<ArgT> && !std::is_same_v<ArgT, sycl::OpenCLMemT>)
+        ? sycl::detail::kernel_param_kind_t::kind_pointer
+        : sycl::detail::kernel_param_kind_t::kind_std_layout;
+
+template <typename T>
+sycl::detail::KernelArgView makeKernelArgView(const T &Arg) {
+  using sycl::detail::kernel_param_kind_t;
+  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};
+  else
+    return {&Arg, sizeof(ArgT), kernel_arg_kind_v<T>};
+}
+
 template <typename CommandGroupFunc, typename PropertiesT>
 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<Dimensions> Range,
 template <int Dimensions, typename... ArgsT>
 void nd_launch(queue Q, nd_range<Dimensions> Range, const kernel &KernelObj,
                ArgsT &&...Args) {
-  submit(std::move(Q), [&](handler &CGH) {
-    nd_launch(CGH, Range, KernelObj, std::forward<ArgsT>(Args)...);
-  });
+  // The handler-less path only takes arguments that can be bound directly,
+  // anything else goes through the handler overload above.
+  if constexpr ((detail::is_direct_kernel_arg_v<ArgsT> && ...)) {
+    // 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<ArgsT>(Args)...);
+    });
+  }
 }
 
 template <int Dimensions, typename Properties, typename... ArgsT>

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

sycl/include/sycl/queue.hpp

@@ -18,6 +18,7 @@
 #include <sycl/detail/common.hpp>             // for code_location
 #include <sycl/detail/defines_elementary.hpp> // for __SYCL2020_DEP...
 #include <sycl/detail/export.hpp>             // for __SYCL_EXPORT
+#include <sycl/detail/kernel_arg_view.hpp>    // for KernelArgView
 #include <sycl/detail/kernel_desc.hpp>        // for KernelInfo
 #include <sycl/detail/nd_range_view.hpp>
 #include <sycl/detail/optional.hpp>
@@ -59,12 +60,21 @@ inline namespace _V1 {
 class context;
 class device;
 class event;
+class kernel;
 class queue;
 
 template <backend BackendName, class SyclObjectT>
 auto get_native(const SyclObjectT &Obj)
     -> backend_return_t<BackendName, SyclObjectT>;
 
+// 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<const detail::KernelArgView> 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,

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<detail::kernel_impl> &KernelImpl,
+    sycl::span<const sycl::detail::KernelArgView> 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<detail::kernel_bundle_impl> 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<EventImplPtr, bool> {
+    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<void *>(Args[I].MPtr);
+      if (!SchedulerBypass) {
+        const char *Bytes = static_cast<const char *>(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<int>(Args[I].MSize),
+                   static_cast<int>(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<detail::CGExecKernel>(
+        KData.getNDRDesc(), /*HostKernel*/ nullptr, KernelImpl,
+        KernelBundleImpl, std::move(CGData), std::move(KData).getArgs(),
+        *KData.getDeviceKernelInfoPtr(),
+        std::vector<std::shared_ptr<detail::stream_impl>>{},
+        std::vector<std::shared_ptr<const void>>{}, 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<ext::oneapi::experimental::detail::exec_graph_impl>
         ExecGraph,

sycl/source/detail/queue_impl.hpp

@@ -378,6 +378,18 @@ class queue_impl : public std::enable_shared_from_this<queue_impl> {
                               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<detail::kernel_impl> &KernelImpl,
+      sycl::span<const sycl::detail::KernelArgView> 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) {

sycl/source/queue.cpp

@@ -13,6 +13,7 @@
 #include <sycl/event.hpp>
 #include <sycl/exception_list.hpp>
 #include <sycl/handler.hpp>
+#include <sycl/kernel.hpp>
 #include <sycl/queue.hpp>
 
 #include <algorithm>
@@ -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<const detail::KernelArgView> 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<

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 <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/experimental/enqueue_functions.hpp>
+#include <sycl/ext/oneapi/free_function_queries.hpp>
+#include <sycl/kernel_bundle.hpp>
+#include <sycl/properties/all_properties.hpp>
+#include <sycl/usm.hpp>
+
+#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 <auto *Func> sycl::kernel getKernel(sycl::queue &Q) {
+  auto Bundle =
+      oneapiext::get_kernel_bundle<Func, sycl::bundle_state::executable>(
+          Q.get_context());
+  return Bundle.template ext_oneapi_get_kernel<Func>();
+}
+
+int main() {
+  sycl::queue Q{sycl::property::queue::in_order{}};
+
+  int *Out = sycl::malloc_shared<int>(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<addFourInts>(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;
+}