[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,61 @@ template <typename LCRangeT, typename LCPropertiesT> struct LaunchConfigAccess {
   }
 };
 
+// The argument type as the kernel sees it. Deliberately not `std::decay_t`,
+// which turns an array into a pointer: an array has to keep being bound as the
+// bytes it is, which is what `handler::setArgHelper` does with it.
+template <typename T>
+using plain_arg_t = std::remove_cv_t<std::remove_reference_t<T>>;
+
+// An argument that is bound as its own bytes with no further interpretation.
+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 as plain bytes, i.e. one that carries no
+// requirement for the scheduler to track. Accessors, local accessors, streams
+// and work group memory are deliberately excluded and keep using the command
+// group path; `HasSpecialCaptures` in the runtime draws the same line. So is
+// every other class type, which may be a struct with special types inside and
+// then needs `kind_struct_with_special_type` instead.
+//
+// An array of scalars is bound as the bytes it is, which is what
+// `handler::setArgHelper` does with it, so it belongs on this path.
+template <typename T>
+inline constexpr bool is_plain_kernel_arg_v =
+    is_scalar_kernel_arg_v<plain_arg_t<T>> ||
+    (std::is_array_v<plain_arg_t<T>> &&
+     is_scalar_kernel_arg_v<std::remove_all_extents_t<plain_arg_t<T>>>) ||
+    std::is_same_v<plain_arg_t<T>, raw_kernel_arg>;
+
+// The kind a plain argument has to carry. A pointer has to keep its kind: the
+// runtime binds a pointer argument as UR_EXP_KERNEL_ARG_TYPE_POINTER, which the
+// OpenCL adapter passes to clSetKernelArgMemPointerINTEL rather than to
+// clSetKernelArg, so plain bytes are not a substitute. The Native CPU adapter
+// draws the same distinction: it puts a pointer argument straight into the
+// argument slot, whereas a value argument lands there as the address of the
+// adapter's own copy.
+//
+// `cl_mem` is the one pointer that is not an address: it names a memory object
+// and has to be bound as the bytes of the handle, which is the exception
+// `handler::setArgHelper` makes for `OpenCLMemT`.
+template <typename T>
+inline constexpr sycl::detail::kernel_param_kind_t plain_arg_kind_v =
+    (std::is_pointer_v<plain_arg_t<T>> &&
+     !std::is_same_v<plain_arg_t<T>, 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;
+  if constexpr (std::is_same_v<plain_arg_t<T>, raw_kernel_arg>)
+    return {RawKernelArgAccess::getData(Arg), RawKernelArgAccess::getSize(Arg),
+            kernel_param_kind_t::kind_std_layout};
+  else
+    return {&Arg, sizeof(plain_arg_t<T>), plain_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 +464,28 @@ 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)...);
-  });
+  if constexpr ((detail::is_plain_kernel_arg_v<ArgsT> && ...)) {
+    // Bind the arguments straight from this call, so that neither a handler nor
+    // a command group object has to be created. 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)...};
+    // An overload that ends in a parameter pack cannot take a trailing
+    // code_location parameter, so the location is the one this header sees,
+    // as it was when this overload went through submit(). Seed the TLS slot
+    // rather than leaving it default-constructed, which the instrumentation
+    // reads as a null file and function name.
+    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<