[SYCL] Implement sycl_ext_oneapi_event_mode extension (#16108)

This commit implements the sycl_ext_oneapi_event_mode extension. Of
particular focus is the low-power event mode.

---------

Signed-off-by: Larsen, Steffen <[email protected]>

Author: steffenlarsen

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

@@ -89,20 +89,27 @@ template <typename LCRangeT, typename LCPropertiesT> struct LaunchConfigAccess {
   }
 };
 
-template <typename CommandGroupFunc>
-void submit_impl(queue &Q, CommandGroupFunc &&CGF,
+template <typename CommandGroupFunc, typename PropertiesT>
+void submit_impl(queue &Q, PropertiesT Props, CommandGroupFunc &&CGF,
                  const sycl::detail::code_location &CodeLoc) {
-  Q.submit_without_event(std::forward<CommandGroupFunc>(CGF), CodeLoc);
+  Q.submit_without_event(Props, std::forward<CommandGroupFunc>(CGF), CodeLoc);
+}
+
+template <typename CommandGroupFunc, typename PropertiesT>
+event submit_with_event_impl(queue &Q, PropertiesT Props,
+                             CommandGroupFunc &&CGF,
+                             const sycl::detail::code_location &CodeLoc) {
+  return Q.submit_with_event(Props, std::forward<CommandGroupFunc>(CGF),
+                             nullptr, CodeLoc);
 }
 } // namespace detail
 
 template <typename CommandGroupFunc, typename PropertiesT>
 void submit(queue Q, PropertiesT Props, CommandGroupFunc &&CGF,
             const sycl::detail::code_location &CodeLoc =
                 sycl::detail::code_location::current()) {
-  std::ignore = Props;
   sycl::ext::oneapi::experimental::detail::submit_impl(
-      Q, std::forward<CommandGroupFunc>(CGF), CodeLoc);
+      Q, Props, std::forward<CommandGroupFunc>(CGF), CodeLoc);
 }
 
 template <typename CommandGroupFunc>
@@ -116,8 +123,8 @@ template <typename CommandGroupFunc, typename PropertiesT>
 event submit_with_event(queue Q, PropertiesT Props, CommandGroupFunc &&CGF,
                         const sycl::detail::code_location &CodeLoc =
                             sycl::detail::code_location::current()) {
-  std::ignore = Props;
-  return Q.submit(std::forward<CommandGroupFunc>(CGF), CodeLoc);
+  return sycl::ext::oneapi::experimental::detail::submit_with_event_impl(
+      Q, Props, std::forward<CommandGroupFunc>(CGF), CodeLoc);
 }
 
 template <typename CommandGroupFunc>

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

@@ -0,0 +1,37 @@
+//==-- cluster_group_prop.hpp --- SYCL extension for event mode property ---==//
+//
+// 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/ext/oneapi/properties/property.hpp>
+
+namespace sycl {
+inline namespace _V1 {
+namespace ext::oneapi::experimental {
+
+enum class event_mode_enum { none, low_power };
+
+struct event_mode
+    : detail::run_time_property_key<event_mode, detail::PropKind::EventMode> {
+  event_mode(event_mode_enum mode) : value(mode) {}
+
+  event_mode_enum value;
+};
+
+using event_mode_key = event_mode;
+
+inline bool operator==(const event_mode &lhs, const event_mode &rhs) {
+  return lhs.value == rhs.value;
+}
+inline bool operator!=(const event_mode &lhs, const event_mode &rhs) {
+  return !(lhs == rhs);
+}
+
+} // namespace ext::oneapi::experimental
+} // namespace _V1
+} // namespace sycl

sycl/include/sycl/ext/oneapi/properties/property.hpp

@@ -223,8 +223,9 @@ enum PropKind : uint32_t {
   InitializeToIdentity = 78,
   WorkGroupScratchSize = 79,
   Restrict = 80,
+  EventMode = 81,
   // PropKindSize must always be the last value.
-  PropKindSize = 81,
+  PropKindSize = 82,
 };
 
 template <typename PropertyT> struct PropertyToKind {

sycl/include/sycl/queue.hpp

@@ -30,15 +30,16 @@
 #include <sycl/exception_list.hpp>         // for defaultAsyncHa...
 #include <sycl/ext/oneapi/device_global/device_global.hpp> // for device_global
 #include <sycl/ext/oneapi/device_global/properties.hpp> // for device_image_s...
-#include <sycl/ext/oneapi/experimental/graph.hpp>       // for command_graph...
-#include <sycl/ext/oneapi/properties/properties.hpp>    // for empty_properti...
-#include <sycl/handler.hpp>                             // for handler, isDev...
-#include <sycl/id.hpp>                                  // for id
-#include <sycl/kernel.hpp>                              // for auto_name
-#include <sycl/kernel_handler.hpp>                      // for kernel_handler
-#include <sycl/nd_range.hpp>                            // for nd_range
-#include <sycl/property_list.hpp>                       // for property_list
-#include <sycl/range.hpp>                               // for range
+#include <sycl/ext/oneapi/experimental/event_mode_property.hpp>
+#include <sycl/ext/oneapi/experimental/graph.hpp>    // for command_graph...
+#include <sycl/ext/oneapi/properties/properties.hpp> // for empty_properti...
+#include <sycl/handler.hpp>                          // for handler, isDev...
+#include <sycl/id.hpp>                               // for id
+#include <sycl/kernel.hpp>                           // for auto_name
+#include <sycl/kernel_handler.hpp>                   // for kernel_handler
+#include <sycl/nd_range.hpp>                         // for nd_range
+#include <sycl/property_list.hpp>                    // for property_list
+#include <sycl/range.hpp>                            // for range
 
 #include <cstddef>     // for size_t
 #include <functional>  // for function
@@ -99,6 +100,9 @@ class __SYCL_EXPORT SubmissionInfo {
   std::shared_ptr<detail::queue_impl> &SecondaryQueue();
   const std::shared_ptr<detail::queue_impl> &SecondaryQueue() const;
 
+  ext::oneapi::experimental::event_mode_enum &EventMode();
+  const ext::oneapi::experimental::event_mode_enum &EventMode() const;
+
 private:
   std::shared_ptr<SubmissionInfoImpl> impl = nullptr;
 };
@@ -111,9 +115,14 @@ enum class queue_state { executing, recording };
 struct image_descriptor;
 
 namespace detail {
-template <typename CommandGroupFunc>
-void submit_impl(queue &Q, CommandGroupFunc &&CGF,
+template <typename CommandGroupFunc, typename PropertiesT>
+void submit_impl(queue &Q, PropertiesT Props, CommandGroupFunc &&CGF,
                  const sycl::detail::code_location &CodeLoc);
+
+template <typename CommandGroupFunc, typename PropertiesT>
+event submit_with_event_impl(queue &Q, PropertiesT Props,
+                             CommandGroupFunc &&CGF,
+                             const sycl::detail::code_location &CodeLoc);
 } // namespace detail
 } // namespace ext::oneapi::experimental
 
@@ -366,7 +375,9 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   std::enable_if_t<std::is_invocable_r_v<void, T, handler &>, event> submit(
       T CGF,
       const detail::code_location &CodeLoc = detail::code_location::current()) {
-    return submit_with_event(CGF, /*SecondaryQueuePtr=*/nullptr, CodeLoc);
+    return submit_with_event(
+        sycl::ext::oneapi::experimental::empty_properties_t{}, CGF,
+        /*SecondaryQueuePtr=*/nullptr, CodeLoc);
   }
 
   /// Submits a command group function object to the queue, in order to be
@@ -384,7 +395,9 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   std::enable_if_t<std::is_invocable_r_v<void, T, handler &>, event> submit(
       T CGF, queue &SecondaryQueue,
       const detail::code_location &CodeLoc = detail::code_location::current()) {
-    return submit_with_event(CGF, &SecondaryQueue, CodeLoc);
+    return submit_with_event(
+        sycl::ext::oneapi::experimental::empty_properties_t{}, CGF,
+        &SecondaryQueue, CodeLoc);
   }
 
   /// Prevents any commands submitted afterward to this queue from executing
@@ -2747,11 +2760,28 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
                                            const detail::code_location &);
 #endif
 
-  template <typename CommandGroupFunc>
+  template <typename CommandGroupFunc, typename PropertiesT>
   friend void ext::oneapi::experimental::detail::submit_impl(
-      queue &Q, CommandGroupFunc &&CGF,
+      queue &Q, PropertiesT Props, CommandGroupFunc &&CGF,
       const sycl::detail::code_location &CodeLoc);
 
+  template <typename CommandGroupFunc, typename PropertiesT>
+  friend event ext::oneapi::experimental::detail::submit_with_event_impl(
+      queue &Q, PropertiesT Props, CommandGroupFunc &&CGF,
+      const sycl::detail::code_location &CodeLoc);
+
+  template <typename PropertiesT>
+  void ProcessSubmitProperties(PropertiesT Props, detail::SubmissionInfo &SI) {
+    if constexpr (Props.template has_property<
+                      ext::oneapi::experimental::event_mode_key>()) {
+      ext::oneapi::experimental::event_mode EventModeProp =
+          Props.template get_property<ext::oneapi::experimental::event_mode>();
+      if (EventModeProp.value !=
+          ext::oneapi::experimental::event_mode_enum::none)
+        SI.EventMode() = EventModeProp.value;
+    }
+  }
+
 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   /// TODO: Unused. Remove these when ABI-break window is open.
   event submit_impl(std::function<void(handler &)> CGH,
@@ -2800,16 +2830,18 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   /// Submits a command group function object to the queue, in order to be
   /// scheduled for execution on the device.
   ///
+  /// \param Props is a property list with submission properties.
   /// \param CGF is a function object containing command group.
   /// \param CodeLoc is the code location of the submit call (default argument)
   /// \return a SYCL event object for the submitted command group.
-  template <typename T>
+  template <typename T, typename PropertiesT>
   std::enable_if_t<std::is_invocable_r_v<void, T, handler &>, event>
   submit_with_event(
-      T CGF, queue *SecondaryQueuePtr,
+      PropertiesT Props, T CGF, queue *SecondaryQueuePtr,
       const detail::code_location &CodeLoc = detail::code_location::current()) {
     detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
     detail::SubmissionInfo SI{};
+    ProcessSubmitProperties(Props, SI);
     if (SecondaryQueuePtr)
       SI.SecondaryQueue() = detail::getSyclObjImpl(*SecondaryQueuePtr);
 #if __SYCL_USE_FALLBACK_ASSERT
@@ -2834,18 +2866,21 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
   /// Submits a command group function object to the queue, in order to be
   /// scheduled for execution on the device.
   ///
+  /// \param Props is a property list with submission properties.
   /// \param CGF is a function object containing command group.
   /// \param CodeLoc is the code location of the submit call (default argument)
-  template <typename T>
+  template <typename T, typename PropertiesT>
   std::enable_if_t<std::is_invocable_r_v<void, T, handler &>, void>
-  submit_without_event(T CGF, const detail::code_location &CodeLoc) {
+  submit_without_event(PropertiesT Props, T CGF,
+                       const detail::code_location &CodeLoc) {
 #if __SYCL_USE_FALLBACK_ASSERT
     // If post-processing is needed, fall back to the regular submit.
     // TODO: Revisit whether we can avoid this.
-    submit_with_event(CGF, nullptr, CodeLoc);
+    submit_with_event(Props, CGF, nullptr, CodeLoc);
 #else
     detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
     detail::SubmissionInfo SI{};
+    ProcessSubmitProperties(Props, SI);
     submit_without_event_impl(CGF, SI, TlsCodeLocCapture.query(),
                               TlsCodeLocCapture.isToplevel());
 #endif // __SYCL_USE_FALLBACK_ASSERT
@@ -3072,8 +3107,12 @@ event submitAssertCapture(queue &Self, event &Event, queue *SecondaryQueue,
     });
   };
 
-  CopierEv = Self.submit_with_event(CopierCGF, SecondaryQueue, CodeLoc);
-  CheckerEv = Self.submit_with_event(CheckerCGF, SecondaryQueue, CodeLoc);
+  CopierEv = Self.submit_with_event(
+      sycl::ext::oneapi::experimental::empty_properties_t{}, CopierCGF,
+      SecondaryQueue, CodeLoc);
+  CheckerEv = Self.submit_with_event(
+      sycl::ext::oneapi::experimental::empty_properties_t{}, CheckerCGF,
+      SecondaryQueue, CodeLoc);
 
   return CheckerEv;
 }

sycl/source/detail/cg.hpp

@@ -15,8 +15,9 @@
 #include <sycl/detail/ur.hpp>       // for ur_rect_region_t, ur_rect_offset_t
 #include <sycl/event.hpp>           // for event_impl
 #include <sycl/exception_list.hpp>  // for queue_impl
-#include <sycl/kernel.hpp>          // for kernel_impl
-#include <sycl/kernel_bundle.hpp>   // for kernel_bundle_impl
+#include <sycl/ext/oneapi/experimental/event_mode_property.hpp>
+#include <sycl/kernel.hpp>        // for kernel_impl
+#include <sycl/kernel_bundle.hpp> // for kernel_bundle_impl
 
 #include <assert.h> // for assert
 #include <memory>   // for shared_ptr, unique_ptr
@@ -425,12 +426,16 @@ class CGAdviseUSM : public CG {
 class CGBarrier : public CG {
 public:
   std::vector<detail::EventImplPtr> MEventsWaitWithBarrier;
+  ext::oneapi::experimental::event_mode_enum MEventMode =
+      ext::oneapi::experimental::event_mode_enum::none;
 
   CGBarrier(std::vector<detail::EventImplPtr> EventsWaitWithBarrier,
+            ext::oneapi::experimental::event_mode_enum EventMode,
             CG::StorageInitHelper CGData, CGType Type,
             detail::code_location loc = {})
       : CG(Type, std::move(CGData), std::move(loc)),
-        MEventsWaitWithBarrier(std::move(EventsWaitWithBarrier)) {}
+        MEventsWaitWithBarrier(std::move(EventsWaitWithBarrier)),
+        MEventMode(EventMode) {}
 };
 
 class CGProfilingTag : public CG {

sycl/source/detail/handler_impl.hpp

@@ -196,7 +196,12 @@ class handler_impl {
   bool MIsTopCodeLoc = true;
 
   /// List of work group memory objects associated with this handler
-  std::vector<std::shared_ptr<detail::work_group_memory_impl>> MWorkGroupMemoryObjects;
+  std::vector<std::shared_ptr<detail::work_group_memory_impl>>
+      MWorkGroupMemoryObjects;
+
+  /// Potential event mode for the result event of the command.
+  ext::oneapi::experimental::event_mode_enum MEventMode =
+      ext::oneapi::experimental::event_mode_enum::none;
 };
 
 } // namespace detail

sycl/source/detail/queue_impl.cpp

@@ -358,6 +358,7 @@ event queue_impl::submit_impl(const std::function<void(handler &)> &CGF,
                               bool IsTopCodeLoc,
                               const SubmissionInfo &SubmitInfo) {
   handler Handler(Self, PrimaryQueue, SecondaryQueue, CallerNeedsEvent);
+  auto HandlerImpl = detail::getSyclObjImpl(Handler);
   Handler.saveCodeLoc(Loc, IsTopCodeLoc);
 
   {
@@ -368,13 +369,15 @@ event queue_impl::submit_impl(const std::function<void(handler &)> &CGF,
   // Scheduler will later omit events, that are not required to execute tasks.
   // Host and interop tasks, however, are not submitted to low-level runtimes
   // and require separate dependency management.
-  const CGType Type = detail::getSyclObjImpl(Handler)->MCGType;
+  const CGType Type = HandlerImpl->MCGType;
   event Event = detail::createSyclObjFromImpl<event>(
       std::make_shared<detail::event_impl>());
   std::vector<StreamImplPtr> Streams;
   if (Type == CGType::Kernel)
     Streams = std::move(Handler.MStreamStorage);
 
+  HandlerImpl->MEventMode = SubmitInfo.EventMode();
+
   if (SubmitInfo.PostProcessorFunc()) {
     auto &PostProcess = *SubmitInfo.PostProcessorFunc();
 

sycl/source/detail/queue_impl.hpp

@@ -71,6 +71,8 @@ enum QueueOrder { Ordered, OOO };
 struct SubmissionInfoImpl {
   optional<detail::SubmitPostProcessF> MPostProcessorFunc = std::nullopt;
   std::shared_ptr<detail::queue_impl> MSecondaryQueue = nullptr;
+  ext::oneapi::experimental::event_mode_enum MEventMode =
+      ext::oneapi::experimental::event_mode_enum::none;
 };
 
 class queue_impl {