[SYCL] Extend the handler-less path to API functions with event dependencies (intel#20461)

The kernel submission APIs with event dependencies are switched to use the handler-less kernel submission path. --------- Co-authored-by: Sergei Vinogradov <sergey.vinogradov@intel.com>

Author: slawekptak

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

@@ -157,7 +157,7 @@ void single_task(queue Q, const KernelType &KernelObj,
  if constexpr (!(detail::KernelLambdaHasKernelHandlerArgT<KernelType,
  void>::value)) {
  detail::submit_kernel_direct_single_task<KernelName>(
- std::move(Q), KernelObj, empty_properties_t{}, CodeLoc);
+ std::move(Q), KernelObj, {}, empty_properties_t{}, CodeLoc);
  } else {
  submit(
  std::move(Q),
@@ -312,7 +312,7 @@ void nd_launch(queue Q, launch_config<nd_range<Dimensions>, Properties> Config,
  LaunchConfigAccess(Config);
 
  detail::submit_kernel_direct_parallel_for<KernelName>(
- std::move(Q), LaunchConfigAccess.getRange(), KernelObj,
+ std::move(Q), LaunchConfigAccess.getRange(), KernelObj, {},
  LaunchConfigAccess.getProperties());
  } else {
  submit(std::move(Q), [&](handler &CGH) {

sycl/include/sycl/khr/free_function_commands.hpp

@@ -317,7 +317,7 @@ void launch_task(const sycl::queue &q, KernelType &&k,
  if constexpr (!(detail::KernelLambdaHasKernelHandlerArgT<KernelType,
  void>::value)) {
  detail::submit_kernel_direct_single_task(
- q, std::forward<KernelType>(k),
+ q, std::forward<KernelType>(k), {},
  ext::oneapi::experimental::empty_properties_t{}, codeLoc);
  } else {
  submit(q, [&](handler &h) { launch_task<KernelType>(h, k); }, codeLoc);

sycl/include/sycl/queue.hpp

@@ -40,6 +40,7 @@
 #include <sycl/nd_range.hpp> // for nd_range
 #include <sycl/property_list.hpp> // for property_list
 #include <sycl/range.hpp> // for range
+#include <sycl/sycl_span.hpp> // for sycl::span
 
 #include <cstddef> // for size_t
 #include <functional> // for function
@@ -68,6 +69,7 @@ event __SYCL_EXPORT submit_kernel_direct_with_event_impl(
  const queue &Queue, const nd_range<Dims> &Range,
  detail::HostKernelRefBase &HostKernel,
  detail::DeviceKernelInfo *DeviceKernelInfo,
+ sycl::span<const event> DepEvents,
  const detail::KernelPropertyHolderStructTy &Props,
  const detail::code_location &CodeLoc, bool IsTopCodeLoc);
 
@@ -76,6 +78,7 @@ void __SYCL_EXPORT submit_kernel_direct_without_event_impl(
  const queue &Queue, const nd_range<Dims> &Range,
  detail::HostKernelRefBase &HostKernel,
  detail::DeviceKernelInfo *DeviceKernelInfo,
+ sycl::span<const event> DepEvents,
  const detail::KernelPropertyHolderStructTy &Props,
  const detail::code_location &CodeLoc, bool IsTopCodeLoc);
 
@@ -165,7 +168,7 @@ template <detail::WrapAs WrapAs, typename LambdaArgType,
  typename KernelTypeUniversalRef, int Dims>
 auto submit_kernel_direct(
  const queue &Queue, const nd_range<Dims> &Range,
- KernelTypeUniversalRef &&KernelFunc,
+ KernelTypeUniversalRef &&KernelFunc, sycl::span<const event> DepEvents,
  const PropertiesT &ExtraProps =
  ext::oneapi::experimental::empty_properties_t{},
  const detail::code_location &CodeLoc = detail::code_location::current()) {
@@ -230,12 +233,14 @@ auto submit_kernel_direct(
 
  if constexpr (EventNeeded) {
  return submit_kernel_direct_with_event_impl(
- Queue, Range, HostKernel, DeviceKernelInfoPtr, ParsedProperties,
- TlsCodeLocCapture.query(), TlsCodeLocCapture.isToplevel());
+ Queue, Range, HostKernel, DeviceKernelInfoPtr, DepEvents,
+ ParsedProperties, TlsCodeLocCapture.query(),
+ TlsCodeLocCapture.isToplevel());
  } else {
  submit_kernel_direct_without_event_impl(
- Queue, Range, HostKernel, DeviceKernelInfoPtr, ParsedProperties,
- TlsCodeLocCapture.query(), TlsCodeLocCapture.isToplevel());
+ Queue, Range, HostKernel, DeviceKernelInfoPtr, DepEvents,
+ ParsedProperties, TlsCodeLocCapture.query(),
+ TlsCodeLocCapture.isToplevel());
  }
 }
 
@@ -244,7 +249,7 @@ template <typename KernelName = detail::auto_name, bool EventNeeded = false,
  typename KernelTypeUniversalRef, int Dims>
 auto submit_kernel_direct_parallel_for(
  const queue &Queue, const nd_range<Dims> &Range,
- KernelTypeUniversalRef &&KernelFunc,
+ KernelTypeUniversalRef &&KernelFunc, sycl::span<const event> DepEvents = {},
  const PropertiesT &Props = ext::oneapi::experimental::empty_properties_t{},
  const detail::code_location &CodeLoc = detail::code_location::current()) {
 
@@ -266,23 +271,25 @@ auto submit_kernel_direct_parallel_for(
  return submit_kernel_direct<detail::WrapAs::parallel_for, TransformedArgType,
  KernelName, EventNeeded, PropertiesT,
  KernelTypeUniversalRef, Dims>(
- Queue, Range, std::forward<KernelTypeUniversalRef>(KernelFunc), Props,
- CodeLoc);
+ Queue, Range, std::forward<KernelTypeUniversalRef>(KernelFunc), DepEvents,
+ Props, CodeLoc);
 }
 
 template <typename KernelName = detail::auto_name, bool EventNeeded = false,
  typename PropertiesT = ext::oneapi::experimental::empty_properties_t,
  typename KernelTypeUniversalRef>
 auto submit_kernel_direct_single_task(
  const queue &Queue, KernelTypeUniversalRef &&KernelFunc,
+ sycl::span<const event> DepEvents = {},
  const PropertiesT &Props = ext::oneapi::experimental::empty_properties_t{},
  const detail::code_location &CodeLoc = detail::code_location::current()) {
 
  return submit_kernel_direct<detail::WrapAs::single_task, void, KernelName,
  EventNeeded, PropertiesT, KernelTypeUniversalRef,
  1>(
  Queue, nd_range<1>{1, 1},
- std::forward<KernelTypeUniversalRef>(KernelFunc), Props, CodeLoc);
+ std::forward<KernelTypeUniversalRef>(KernelFunc), DepEvents, Props,
+ CodeLoc);
 }
 
 } // namespace detail
@@ -2802,7 +2809,7 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
  if constexpr (!(detail::KernelLambdaHasKernelHandlerArgT<KernelType,
  void>::value)) {
  return detail::submit_kernel_direct_single_task<KernelName, true>(
- *this, KernelFunc, Properties, TlsCodeLocCapture.query());
+ *this, KernelFunc, {}, Properties, TlsCodeLocCapture.query());
  } else {
  return submit(
  [&](handler &CGH) {
@@ -2852,13 +2859,23 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
  "Use queue.submit() instead");
 
  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
- return submit(
- [&](handler &CGH) {
- CGH.depends_on(DepEvent);
- CGH.template single_task<KernelName, KernelType, PropertiesT>(
- Properties, KernelFunc);
- },
- TlsCodeLocCapture.query());
+
+ // TODO The handler-less path does not support kernel functions
+ // with the kernel_handler type argument yet.
+ if constexpr (!(detail::KernelLambdaHasKernelHandlerArgT<KernelType,
+ void>::value)) {
+ return detail::submit_kernel_direct_single_task<KernelName, true>(
+ *this, KernelFunc, sycl::span<const event>(&DepEvent, 1), Properties,
+ TlsCodeLocCapture.query());
+ } else {
+ return submit(
+ [&](handler &CGH) {
+ CGH.depends_on(DepEvent);
+ CGH.template single_task<KernelName, KernelType, PropertiesT>(
+ Properties, KernelFunc);
+ },
+ TlsCodeLocCapture.query());
+ }
  }
 
  /// single_task version with a kernel represented as a lambda.
@@ -2903,13 +2920,22 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
  "Use queue.submit() instead");
 
  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
- return submit(
- [&](handler &CGH) {
- CGH.depends_on(DepEvents);
- CGH.template single_task<KernelName, KernelType, PropertiesT>(
- Properties, KernelFunc);
- },
- TlsCodeLocCapture.query());
+
+ // TODO The handler-less path does not support kernel functions
+ // with the kernel_handler type argument yet.
+ if constexpr (!(detail::KernelLambdaHasKernelHandlerArgT<KernelType,
+ void>::value)) {
+ return detail::submit_kernel_direct_single_task<KernelName, true>(
+ *this, KernelFunc, DepEvents, Properties, TlsCodeLocCapture.query());
+ } else {
+ return submit(
+ [&](handler &CGH) {
+ CGH.depends_on(DepEvents);
+ CGH.template single_task<KernelName, KernelType, PropertiesT>(
+ Properties, KernelFunc);
+ },
+ TlsCodeLocCapture.query());
+ }
  }
 
  /// single_task version with a kernel represented as a lambda.
@@ -3348,7 +3374,7 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
  KernelType, sycl::nd_item<Dims>>::value)) {
 
  return detail::submit_kernel_direct_parallel_for<KernelName, true>(
- *this, Range, Rest..., Properties, TlsCodeLocCapture.query());
+ *this, Range, Rest..., {}, Properties, TlsCodeLocCapture.query());
  } else
  return submit(
  [&](handler &CGH) {
@@ -3377,7 +3403,7 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
  !(detail::KernelLambdaHasKernelHandlerArgT<
  KernelType, sycl::nd_item<Dims>>::value)) {
  return detail::submit_kernel_direct_parallel_for<KernelName, true>(
- *this, Range, Rest...,
+ *this, Range, Rest..., {},
  ext::oneapi::experimental::empty_properties_t{},
  TlsCodeLocCapture.query());
  } else {
@@ -3431,12 +3457,25 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
  parallel_for(nd_range<Dims> Range, event DepEvent, RestT &&...Rest) {
  constexpr detail::code_location CodeLoc = getCodeLocation<KernelName>();
  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
- return submit(
- [&](handler &CGH) {
- CGH.depends_on(DepEvent);
- CGH.template parallel_for<KernelName>(Range, Rest...);
- },
- TlsCodeLocCapture.query());
+ using KernelType = std::tuple_element_t<0, std::tuple<RestT...>>;
+
+ // TODO The handler-less path does not support reductions, and
+ // kernel functions with the kernel_handler type argument yet.
+ if constexpr (sizeof...(RestT) == 1 &&
+ !(detail::KernelLambdaHasKernelHandlerArgT<
+ KernelType, sycl::nd_item<Dims>>::value)) {
+ return detail::submit_kernel_direct_parallel_for<KernelName, true>(
+ *this, Range, Rest..., sycl::span<const event>(&DepEvent, 1),
+ ext::oneapi::experimental::empty_properties_t{},
+ TlsCodeLocCapture.query());
+ } else {
+ return submit(
+ [&](handler &CGH) {
+ CGH.depends_on(DepEvent);
+ CGH.template parallel_for<KernelName>(Range, Rest...);
+ },
+ TlsCodeLocCapture.query());
+ }
  }
 
  /// parallel_for version with a kernel represented as a lambda + nd_range that
@@ -3485,12 +3524,25 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
  RestT &&...Rest) {
  constexpr detail::code_location CodeLoc = getCodeLocation<KernelName>();
  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
- return submit(
- [&](handler &CGH) {
- CGH.depends_on(DepEvents);
- CGH.template parallel_for<KernelName>(Range, Rest...);
- },
- TlsCodeLocCapture.query());
+ using KernelType = std::tuple_element_t<0, std::tuple<RestT...>>;
+
+ // TODO The handler-less path does not support reductions, and
+ // kernel functions with the kernel_handler type argument yet.
+ if constexpr (sizeof...(RestT) == 1 &&
+ !(detail::KernelLambdaHasKernelHandlerArgT<
+ KernelType, sycl::nd_item<Dims>>::value)) {
+ return detail::submit_kernel_direct_parallel_for<KernelName, true>(
+ *this, Range, Rest..., DepEvents,
+ ext::oneapi::experimental::empty_properties_t{},
+ TlsCodeLocCapture.query());
+ } else {
+ return submit(
+ [&](handler &CGH) {
+ CGH.depends_on(DepEvents);
+ CGH.template parallel_for<KernelName>(Range, Rest...);
+ },
+ TlsCodeLocCapture.query());
+ }
  }
 
  /// Copies data from a memory region pointed to by a placeholder accessor to

sycl/source/detail/event_deps.hpp

@@ -0,0 +1,125 @@
+//==---------------- event_deps.hpp - SYCL event dependency utils ----------==//
+//
+// 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 <detail/context_impl.hpp>
+#include <detail/device_impl.hpp>
+#include <detail/event_impl.hpp>
+#include <detail/graph/graph_impl.hpp>
+#include <detail/queue_impl.hpp>
+#include <sycl/detail/cg_types.hpp>
+#include <sycl/exception.hpp>
+
+#include <vector>
+
+namespace sycl {
+inline namespace _V1 {
+namespace detail {
+
+/// Adds an event dependency to the list of dependencies, performing
+/// a series of checks.
+///
+/// If the event is associated with a graph, and the queue is not,
+/// the queue will be switched to a recording mode (transitive queue
+/// recording feature).
+///
+/// The LockQueue template argument defines whether the queue lock
+/// should be acquired for the transition to a recording mode. It is
+/// set to false in cases where the event dependencies are set directly
+/// in the command submission flow and the lock is already acquired.
+///
+/// \param EventImpl Event to register as a dependency
+/// \param EventsRegistered A list of already registered events, where
+/// the event will be added.
+/// \param QueueImpl A queue associated with the event dependencies. Can
+/// be nullptr if no associated queue.
+/// \param ContextImpl A context associated with a queue or graph.
+/// \param DeviceImpl A device associated with a queue or graph.
+/// \param GraphImpl A graph associated with a queue or a handler. Can
+/// be nullptr if no associated graph.
+/// \param CommandGroupType Type of command group.
+template <bool LockQueue = true>
+void registerEventDependency(
+ const EventImplPtr &EventImpl, std::vector<EventImplPtr> &EventsRegistered,
+ queue_impl *QueueImpl, const context_impl &ContextImpl,
+ const device_impl &DeviceImpl,
+ const ext::oneapi::experimental::detail::graph_impl *GraphImpl,
+ CGType CommandGroupType) {
+
+ if (!EventImpl)
+ return;
+ if (EventImpl->isDiscarded()) {
+ throw sycl::exception(make_error_code(errc::invalid),
+ "Queue operation cannot depend on discarded event.");
+ }
+
+ // Async alloc calls adapter immediately. Any explicit/implicit dependencies
+ // are handled at that point, including in order queue deps. Further calls to
+ // depends_on after an async alloc are explicitly disallowed.
+ if (CommandGroupType == CGType::AsyncAlloc) {
+ throw sycl::exception(make_error_code(errc::invalid),
+ "Cannot submit a dependency after an asynchronous "
+ "allocation has already been executed!");
+ }
+
+ auto EventGraph = EventImpl->getCommandGraph();
+ if (QueueImpl && EventGraph) {
+ auto QueueGraph = QueueImpl->getCommandGraph();
+
+ if (&EventGraph->getContextImpl() != &ContextImpl) {
+ throw sycl::exception(
+ make_error_code(errc::invalid),
+ "Cannot submit to a queue with a dependency from a graph that is "
+ "associated with a different context.");
+ }
+
+ if (&EventGraph->getDeviceImpl() != &DeviceImpl) {
+ throw sycl::exception(
+ make_error_code(errc::invalid),
+ "Cannot submit to a queue with a dependency from a graph that is "
+ "associated with a different device.");
+ }
+
+ if (QueueGraph && QueueGraph != EventGraph) {
+ throw sycl::exception(sycl::make_error_code(errc::invalid),
+ "Cannot submit to a recording queue with a "
+ "dependency from a different graph.");
+ }
+
+ // If the event dependency has a graph, that means that the queue that
+ // created it was in recording mode. If the current queue is not recording,
+ // we need to set it to recording (implements the transitive queue recording
+ // feature).
+ if (!QueueGraph) {
+ if constexpr (LockQueue) {
+ EventGraph->beginRecording(*QueueImpl);
+ } else {
+ EventGraph->beginRecordingUnlockedQueue(*QueueImpl);
+ }
+ }
+ }
+
+ if (GraphImpl) {
+ if (EventGraph == nullptr) {
+ throw sycl::exception(
+ make_error_code(errc::invalid),
+ "Graph nodes cannot depend on events from outside the graph.");
+ }
+ if (EventGraph.get() != GraphImpl) {
+ throw sycl::exception(
+ make_error_code(errc::invalid),
+ "Graph nodes cannot depend on events from another graph.");
+ }
+ }
+ EventsRegistered.push_back(EventImpl);
+}
+
+} // namespace detail
+} // namespace _V1
+} // namespace sycl