-
Notifications
You must be signed in to change notification settings - Fork 791
[SYCL][DOC] Add new extension sycl_ext_oneapi_clock #19842
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
Changes from 5 commits
a5f0e9d
571c693
eeb908b
0843199
42ab657
4fd1334
a01aeb3
d1f70e7
File filter
Filter by extension
Conversations
Jump to
Diff view
Diff view
There are no files selected for viewing
Original file line number | Diff line number | Diff line change |
---|---|---|
@@ -0,0 +1,161 @@ | ||
= sycl_ext_oneapi_clock | ||
|
||
: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} | ||
|
||
|
||
== Notice | ||
|
||
[%hardbreaks] | ||
Copyright (C) 2025 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 10 specification. All | ||
references below to the "core SYCL specification" or to section numbers in the | ||
SYCL specification refer to that revision. | ||
|
||
== 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 | ||
|
||
The APIs in this extension may be used only on a device that has | ||
`aspect::ext_oneapi_clock`. The application must check that the device has | ||
this aspect before submitting a kernel using any of the APIs in this | ||
extension. If the application fails to do this, the implementation throws | ||
a synchronous exception with the `errc::kernel_not_supported` error code | ||
when the kernel is submitted to the queue. | ||
|
||
== Overview | ||
|
||
This extension introduces a new free function `clock(scope)`. This function | ||
allows the user to sample the value from one of three clocks provided by the | ||
compute units, depending on the value of the scope argument. `scope` is an | ||
enumeration constant of the new `clock_scope` enum. It should be passed to the | ||
function to define the clock source; e.g., `clock(clock_scope::sub_group)` | ||
samples the value from a clock shared by all work-items executing in the same | ||
sub-group. | ||
|
||
This extension also adds a new aspect, `ext_oneapi_clock`, indicating whether | ||
the feature is supported by the device. | ||
|
||
== 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_CLOCK` 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. | ||
|=== | ||
|
||
=== New device aspect | ||
|
||
This extension adds new device aspect: | ||
|
||
```c++ | ||
namespace sycl { | ||
|
||
enum class aspect : /*unspecified*/ { | ||
ext_oneapi_clock | ||
}; | ||
|
||
} // namespace sycl | ||
``` | ||
|
||
[width="100%",%header,cols="50%,50%"] | ||
|=== | ||
|Aspect | ||
|Description | ||
|
||
|`ext_oneapi_clock` | ||
|Indicates that the device supports the `sycl::ext::oneapi::experimental::clock()` call. | ||
|=== | ||
|
||
=== New free function | ||
|
||
```c++ | ||
|
||
namespace sycl::ext::oneap::experimental { | ||
|
||
enum class clock_scope : /* unspecified */ { | ||
sub_group, | ||
work_group, | ||
device | ||
}; | ||
steffenlarsen marked this conversation as resolved.
Show resolved
Hide resolved
|
||
``` | ||
|
||
KornevNikita marked this conversation as resolved.
Show resolved
Hide resolved
|
||
An enumerator from `clock_scope` should be passed to the `clock()` function to | ||
define the clock source. There are 3 clock sources: `sub_group`, `work_group` | ||
and `device`. Values from the first one (`sub_group`) are shared by all | ||
work-items executing in the same sub-group. Values from the second one | ||
(`work_group`) are are shared by all work-items executing in the same work-group | ||
etc. | ||
steffenlarsen marked this conversation as resolved.
Outdated
Show resolved
Hide resolved
|
||
|
||
```c++ | ||
uint64_t clock(clock_scope scope = clock_scope::sub_group); | ||
|
||
|
||
} // namespace sycl::ext::oneapi::experimental | ||
``` | ||
|
||
This function may only be called from within a SYCL kernel function. | ||
|
||
All work-items within the `scope` read from the same source clock. There is no | ||
guarantee that two work-items get the same value. | ||
|
||
steffenlarsen marked this conversation as resolved.
Show resolved
Hide resolved
|
||
_Returns:_ The sample value of a clock as seen by the work-item. | ||
The clock is defined as an unbounded, unsigned integer counter that | ||
monotonically increments over time. The rate at which the clock advances is not | ||
guaranteed to be constant: it may vary over the lifetime of a work-item, differ | ||
between separate executions of the program, and be affected by conditions | ||
outside the control of the programmer. The value returned by this instruction | ||
corresponds to the least significant bits of the clock counter at the time of | ||
execution. Consequently, the sampled value may wrap around zero. | ||
There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. Can we provide a guarantee that the clock counter starts at zero when the kernel starts executing? If so, that might help with the wrap-around problem. There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. My 2 cents:
(the above is not based on any hardware details; just trying to make some estimates) There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. The spec does not provide enough guarantees to do what you want:
To be clear, I don't think SYCL can provide any of these guarantees. The SYCL spec layers on top of SPV_KHR_shader_clock, which provides very few guarantees. @bashbaug do you have any insight here? The SPIR-V spec (and the OpenCL C cl_khr_kernel_clock extension) don't seem to provide enough guarantees to do any useful timing operations. Do vendors provide some additional guarantees in some separate specification? Do applications just make assumptions about the clock that are not guaranteed by the spec, and the code just happens to work? There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more.
I made an assumption of this kind when estimating the upper limit on the wrap-around time. I don't see how this assumption is required for computing time difference (provided wrap-around is handled, but, unlike clock reset, it can be handled). For
Yes, I pointed this out earlier in the discussion (and the fact that the units are not even guaranteed to be uniform). For the wrap-around estimation, 1 tick = 1 nanosecond was an assumption indeed, but based on the conversion constant used in some PTI-GPU samples. For the practical use, there is an issue recorded in the spec below, so hopefully it will be addressed. There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more.
Yes, good point. So, is the only remaining open the one about the units of the counter? There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more.
I think it can be split into two parts, where the second one can be done without the first one:
EDIT: From the application perspective, having the capability in the API is nice even if some hardware don't support it. At the very least, CPUs can support steady counters with known conversion factors. There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. From my perspective, this comment is resolved. I've been convinced that the API is useful even though the values it returns cannot be converted to time units. I have not clicked the "Resolve conversation" button only because @al42and has made some comments here too. Addressing your comments ... I think we cannot provide the behavior you request on GPU because the hardware just doesn't work that way. We can count cycles (not time), and the frequency may change dynamically as the kernel executes. I agree that we could do better on CPU, but I think the motivation for adding this extension is to support GPU, and I do not think we want to invest the effort right now to implement the feature just for CPU. There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. You're right about the current Intel GPU hardware limitations. My concern, however, is designing a future-proof API rather than one constrained by the initial target's lowest common denominator. I believe the API, even an experimental one, should reflect the capabilities of the entire DPC++ ecosystem. We shouldn't permanently limit the API due to the constraints of the first backend to support it, when other backends have broader capabilities. A more flexible API doesn't require implementing it for every backend immediately. It simply prevents breaking changes later when (if) we do add support for more capable hardware. Many SYCL extensions are rolled out incrementally across backends. To be clear, this is about the API design, not a demand for an immediate implementation on all platforms :) There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. It's much easier to loosen constraints in the future than it is to add new constraints. If we have new GPU devices in the future that can provide more guarantees about the clock, then we can update this extension to a new revision and add a new aspect that provides those additional guarantees. Therefore, I don't think the extension's current wording will tie our hands in the future. There was a problem hiding this comment. Choose a reason for hiding this commentThe reason will be displayed to describe this comment to others. Learn more. I agree that it's easier to loosen constraints than to add them. However, I see this as less about adding constraints and more about designing an API that accurately communicates diverse hardware capabilities from the start. Deferring this decision creates unnecessary friction for developers:
A more flexible API from the beginning would avoid these issues and allow developers to fully use all the hardware supported by DPC++. And this extension already seems to accommodate functionality beyond what is supported by off-the-shelf Intel GPUs (#20131). |
||
|
||
== Issues | ||
|
||
. How to convert the result of the function to seconds? | ||
gmlueck marked this conversation as resolved.
Show resolved
Hide resolved
|
Uh oh!
There was an error while loading. Please reload this page.