Skip to content
Merged
Show file tree
Hide file tree
Changes from 5 commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Original file line number Diff line number Diff line change
Expand Up @@ -233,6 +233,9 @@ namespace sycl::ext::oneapi::experimental {
template <typename CommandGroupFunc>
void submit(sycl::queue q, CommandGroupFunc&& cgf);

template <typename CommandGroupFunc, typename Properties>
void submit(sycl::queue q, Properties properties, CommandGroupFunc&& cgf);

}
----
!====
Expand All @@ -250,6 +253,10 @@ namespace sycl::ext::oneapi::experimental {
template <typename CommandGroupFunc>
sycl::event submit_with_event(sycl::queue q, CommandGroupFunc&& cgf);

template <typename CommandGroupFunc, typename Properties>
sycl::event submit_with_event(sycl::queue q, Properties properties,
CommandGroupFunc&& cgf);

}
----
!====
Expand Down
153 changes: 153 additions & 0 deletions sycl/doc/extensions/proposed/sycl_ext_intel_low_power_event.asciidoc
Original file line number Diff line number Diff line change
@@ -0,0 +1,153 @@
= 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: &#8212;{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_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.*


== Overview

On some backends, calling `wait()` on an `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 commands 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_INTEL_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 adds a new property which can be used with the
`submit_with_event` free function from
link:../experimental/sycl_ext_oneapi_enqueue_functions.asciidoc[sycl_ext_oneapi_enqueue_functions].
Passing this property to this function 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.

[_Note:_ The property currently only has an effect on `barrier` and
`partial_barrier` commands enqueued on queues that return
`backend::ext_oneapi_level_zero` from `queue::get_backend()`.
_{endnote}_]

```
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;
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think this should be a runtime property that takes a bool parameter. This will allow usage like:

bool low_power = /* check some option */;
oneapiex::properties props{intelex::low_power_event{low_power}});

Code like this is difficult with the API as you have it now because the existence of the low_power_event property affects the type of the props variable. As a result, you can not do this:

bool low_power = /* check some option */;
oneapiex::properties props;
if (low_power)
  /* no way to add 'low_power_event' property now */

If you make it a runtime property, the property's default constructor can set the property's value to true.

Alternatively, the property could take an enum instead of a bool. I think someone suggested in one of our meetings that there could be other types of low-power events in the future? I can't remember whether we decided that was realistic, though.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

That's a good point. Would it make sense to change the property to an event_mode property and have the value as:

enum event_mode_enum {
  default_mode,
  low_power_mode
};

If we think future modes might overlap, we can also do it as a bitmap.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Am I right in thinking that we can define it as an enum now, and convert it to a bitmask later just by defining the relevant operators?

Either way, I like this direction. I'd be tempted to change the names, though, to avoid some redundancy:

enum class event_mode {
  none, // we can't use"default" because it's a keyword
  low_power,
};

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Am I right in thinking that we can define it as an enum now, and convert it to a bitmask later just by defining the relevant operators?

Absolutely. As long as it doesn't cross the library boundary, we can change it later if we want.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think the enum direction makes sense, but I'm not sure what names you are proposing. It seems like the name event_mode has been proposed as both the name of the property and the name of the enumeration. Can you write out an example usage to illustrate the naming that you propose?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I'm definitely not married to the name of the enum, but I just took inspiration from the naming of the compile-time property enums.

Adjusting the example in the text, the property would be used as

oneapiex::properties Props{intelex::event_mode{intelex::event_mode_enum::low_power}};

Note, I agree with @Pennycook in that it should be an enum class, mainly so it doesn't bleed its contents.


} // 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 command with the low-power event property.
oneapiex::properties Props{intelex::low_power_event};
sycl::event E = oneapiex::submit_with_event(Q, Props, [&](sycl::handler &CGH) {
...
});

// Waiting for the resulting event will use low-power waiting if possible.
E.wait();
}
```
23 changes: 20 additions & 3 deletions sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -95,21 +95,38 @@ void submit_impl(queue &Q, CommandGroupFunc &&CGF,
}
} // namespace detail

template <typename CommandGroupFunc>
void submit(queue Q, CommandGroupFunc &&CGF,
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);
}

template <typename CommandGroupFunc>
event submit_with_event(queue Q, CommandGroupFunc &&CGF,
void submit(queue Q, CommandGroupFunc &&CGF,
const sycl::detail::code_location &CodeLoc =
sycl::detail::code_location::current()) {
submit(Q, empty_properties_t{}, std::forward<CommandGroupFunc>(CGF), CodeLoc);
}

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);
}

template <typename CommandGroupFunc>
event submit_with_event(queue Q, CommandGroupFunc &&CGF,
const sycl::detail::code_location &CodeLoc =
sycl::detail::code_location::current()) {
return submit_with_event(Q, empty_properties_t{},
std::forward<CommandGroupFunc>(CGF), CodeLoc);
}

template <typename KernelName = sycl::detail::auto_name, typename KernelType>
void single_task(handler &CGH, const KernelType &KernelObj) {
CGH.single_task<KernelName>(KernelObj);
Expand Down
Loading