[SYCL][Docs] Add proposed low-power event extension

This commit adds the sycl_ext_intel_low_power_event extension as
proposed. To support this extension, the barriers in the
sycl_ext_oneapi_enqueue_functions extension are given property arguments
without any current consumers.

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

Author: steffenlarsen

sycl/doc/extensions/experimental/sycl_ext_oneapi_enqueue_functions.asciidoc

@@ -670,9 +670,11 @@ a!
 ----
 namespace sycl::ext::oneapi::experimental {
 
-void barrier(sycl::queue q);
+template <typename Properties = empty_properties_t>
+void barrier(sycl::queue q, Properties properties = {});
 
-void barrier(sycl::handler h);
+template <typename Properties = empty_properties_t>
+void barrier(sycl::handler h, Properties properties = {});
 
 }
 ----
@@ -690,9 +692,13 @@ a!
 ----
 namespace sycl::ext::oneapi::experimental {
 
-void partial_barrier(sycl::queue q, const std::vector<sycl::event>& events);
+template <typename Properties = empty_properties_t>
+void partial_barrier(sycl::queue q, const std::vector<sycl::event>& events,
+                     Properties properties = {});
 
-void partial_barrier(sycl::handler h, const std::vector<sycl::event>& events);
+template <typename Properties = empty_properties_t>
+void partial_barrier(sycl::handler h, const std::vector<sycl::event>& events,
+                     Properties properties = {});
 
 }
 ----

sycl/doc/extensions/proposed/sycl_ext_intel_low_power_event.asciidoc

@@ -0,0 +1,151 @@
+= sycl_ext_intel_low_power_event
+
+:source-highlighter: coderay
+:coderay-linenums-mode: table
+
+// This section needs to be after the document title.
+:doctype: book
+:toc2:
+:toc: left
+:encoding: utf-8
+:lang: en
+:dpcpp: pass:[DPC++]
+:endnote: —{nbsp}end{nbsp}note
+
+// Set the default source code type in this document to C++,
+// for syntax highlighting purposes.  This is needed because
+// docbook uses c++ and html5 uses cpp.
+:language: {basebackend@docbook:c++:cpp}
+
+:common_ref_sem: https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:reference-semantics
+
+== Notice
+
+[%hardbreaks]
+Copyright (C) 2024 Intel Corporation.  All rights reserved.
+
+Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks
+of The Khronos Group Inc.  OpenCL(TM) is a trademark of Apple Inc. used by
+permission by Khronos.
+
+
+== Contact
+
+To report problems with this extension, please open a new issue at:
+
+https://github.com/intel/llvm/issues
+
+
+== Dependencies
+
+This extension is written against the SYCL 2020 revision 9 specification.  All
+references below to the "core SYCL specification" or to section numbers in the
+SYCL specification refer to that revision.
+
+This extension also depends on the following other SYCL extensions:
+
+* link:../experimental/sycl_ext_oneapi_enqueue_functions.asciidoc[
+  sycl_ext_oneapi_enqueue_functions]
+* link:../experimental/sycl_ext_oneapi_properties.asciidoc[
+  sycl_ext_oneapi_enqueue_properties]
+
+
+== Status
+
+This is a proposed extension specification, intended to gather community
+feedback.  Interfaces defined in this specification may not be implemented yet
+or may be in a preliminary state.  The specification itself may also change in
+incompatible ways before it is finalized.  *Shipping software products should
+not rely on APIs defined in this specification.*
+
+
+== Backend support status
+
+This extension is currently implemented in {dpcpp} only for all device targets.
+
+== Overview
+
+On some backends, calling `wait()` on a `event` will synchronize using a
+busy-waiting implementation. Though this comes at a low latency for the
+synchronization of the event, it has the downside of consuming high amounts of
+CPU time for no meaningful work.  This extension introduces a new property for
+`sycl::ext::oneapi::experimental::barrier()` that will produce a "low-power"
+event. These new low-power events will, if possible, yield the thread that the
+`wait()` member function is called on and only wake up occasionally to check if
+the event has finished. This reduces the time the CPU spends checking finish
+condition of the wait, at the cost of latency.
+
+
+== Specification
+
+=== Feature test macro
+
+This extension provides a feature-test macro as described in the core SYCL
+specification.  An implementation supporting this extension must predefine the
+macro `SYCL_EXT_ONEAPI_LOW_POWER_EVENT` to one of the values defined in the table
+below.  Applications can test for the existence of this macro to determine if
+the implementation supports this feature, or applications can test the macro's
+value to determine which of the extension's features the implementation
+supports.
+
+[%header,cols="1,5"]
+|===
+|Value
+|Description
+
+|1
+|The APIs of this experimental extension are not versioned, so the
+ feature-test macro always has this value.
+|===
+
+
+=== Low-power event property
+
+This extension also adds a new property which can be used with the `barrier`
+and `partial_barrier` enqueue free functions from
+link:../experimental/sycl_ext_oneapi_enqueue_functions.asciidoc[sycl_ext_oneapi_enqueue_functions].
+Passing this property to either of these functions will act as a hint to the
+`event` created from the corresponding commands to do low-power synchronization.
+If the backend is able to handle low-power events, calling `event::wait()` or
+`event::wait_and_throw()` will cause the thread to yield and only do occasional
+wake-ups to check the event progress.
+
+```
+namespace sycl::ext::intel::experimental {
+
+struct low_power_event_key {
+  using value_t = 
+      oneapi::experimental::property_value<low_power_event_key>;
+};
+
+inline constexpr low_power_event_key::value_t low_power_event;
+
+} // namespace sycl::ext::intel::experimental
+```
+
+=== New property usage example
+
+As an example of how to use the new `low_power_event` property, see the
+following code:
+
+```
+#include <sycl/sycl.hpp>
+
+namespace oneapiex = sycl::ext::oneapi::experimental;
+namespace intelex = sycl::ext::intel::experimental;
+
+int main() {
+  sycl::queue Q;
+
+  // Submit some work to the queue.
+  oneapiex::submit(Q, [&](sycl::handler &CGH) {...});
+
+  // Submit a barrier with the low-power event property.
+  sycl::event E = oneapiex::submit_with_event(Q, [&](sycl::handler &CGH) {
+    oneapiex::barrier(CGH, oneapiex::properties{intelex::low_power_event});
+  });
+
+  // Waiting for the resulting event will use low-power waiting if possible.
+  E.wait();
+}
+```

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

@@ -349,21 +349,34 @@ __SYCL_EXPORT void mem_advise(queue Q, void *Ptr, size_t NumBytes, int Advice,
                               const sycl::detail::code_location &CodeLoc =
                                   sycl::detail::code_location::current());
 
-inline void barrier(handler &CGH) { CGH.ext_oneapi_barrier(); }
+template <typename PropertiesT = empty_properties_t>
+inline void barrier(handler &CGH, PropertiesT Properties = {}) {
+  std::ignore = Properties;
+  CGH.ext_oneapi_barrier();
+}
 
-inline void barrier(queue Q, const sycl::detail::code_location &CodeLoc =
-                                 sycl::detail::code_location::current()) {
-  submit(Q, [&](handler &CGH) { barrier(CGH); }, CodeLoc);
+template <typename PropertiesT = empty_properties_t>
+inline void barrier(queue Q, PropertiesT Properties = {},
+                    const sycl::detail::code_location &CodeLoc =
+                        sycl::detail::code_location::current()) {
+  submit(Q, [&](handler &CGH) { barrier(CGH, Properties); }, CodeLoc);
 }
 
-inline void partial_barrier(handler &CGH, const std::vector<event> &Events) {
+template <typename PropertiesT = empty_properties_t>
+inline void partial_barrier(handler &CGH, const std::vector<event> &Events,
+                            PropertiesT Properties = {}) {
+  std::ignore = Properties;
   CGH.ext_oneapi_barrier(Events);
 }
 
+template <typename PropertiesT = empty_properties_t>
 inline void partial_barrier(queue Q, const std::vector<event> &Events,
+                            PropertiesT Properties = {},
                             const sycl::detail::code_location &CodeLoc =
                                 sycl::detail::code_location::current()) {
-  submit(Q, [&](handler &CGH) { partial_barrier(CGH, Events); }, CodeLoc);
+  submit(
+      Q, [&](handler &CGH) { partial_barrier(CGH, Events, Properties); },
+      CodeLoc);
 }
 
 } // namespace ext::oneapi::experimental