[SYCL][DOC] Add new extension sycl_ext_oneapi_clock #19842
Conversation
|
@intel/llvm-gatekeepers please consider merging |
| 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.
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.
My 2 cents:
- Hardware counter wrapping is a common theme on CPUs.
- The prospective usecase I have in mind requires comparing counters between two kernels; resetting the device-scope counter for each kernel would break that.
- 64-bit nanosecond counter would need a few centuries to overflow, so wrap-around is more of a theoretical possibility. But the spec seem to assume that the counters can have less bits. A 32-bit nanosecond counter will overflow after ~4 seconds, so an overflow can occur within a kernel's lifetime.
(the above is not based on any hardware details; just trying to make some estimates)
There was a problem hiding this comment.
The spec does not provide enough guarantees to do what you want:
-
You seem to be assuming that the clocks start at zero when the process first starts (or maybe when the process submits its first kernel). You are also assuming that the clocks are not reset between kernel submissions. However, none of this is guaranteed by the spec.
-
You are asserting that the clocks count in units of nanoseconds, but this is also not guaranteed by the spec.
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.
You seem to be assuming that the clocks start at zero when the process first starts (or maybe when the process submits its first kernel).
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 clock_scope::device, the spec say: "clock() gets values shared by all work-items executing on the device.". To me, it reads like the clocks are not reset for the lifetime of the device object and also that two kernels executing simultaneously use the same clock.
You are asserting that the clocks count in units of nanoseconds, but this is also not guaranteed by the spec.
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.
For
clock_scope::device, the spec say: "clock()gets values shared by all work-items executing on the device.". To me, it reads like the clocks are not reset for the lifetime of thedeviceobject and also that two kernels executing simultaneously use the same clock.
Yes, good point.
So, is the only remaining open the one about the units of the counter?
There was a problem hiding this comment.
So, is the only remaining open the one about the units of the counter?
I think it can be split into two parts, where the second one can be done without the first one:
-
Defining the counter units: whether the ticks can be converted into seconds (and how).
-
Guaranteeing a "steady" clock: a weaker guarantee that the clock is at least steady (each tick represents a consistent, though unknown, duration). That would still allow, e.g., benchmarking and load balancing. See, e.g.,
GL_EXT_shader_realtime_clock.
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.
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.
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.
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.
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:
-
Preventing developers from using steady clocks, even when the underlying hardware supports them, will slow down the extension adoption.
-
Experimental extensions are unversioned, so adding a new symbol (e.g.,
ext_oneapi_clock_device_steadyaspect) forces developers to write complex code with build-system checks and#ifdefsto manage API differences.
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).
| ```c++ | ||
| namespace sycl::ext::oneap::experimental { | ||
|
|
||
| uint64_t clock(clock_scope scope = clock_scope::sub_group); |
There was a problem hiding this comment.
is it possible to use the clock() to implement the usleep(usecond) function in kernel? or just provide a device function usleep()? Thanks.
There was a problem hiding this comment.
There is no way to convert these clocks to seconds, so unfortunately the answer is no.
updated the overview added new aspects fix some typos
|
@intel/llvm-gatekeepers please consider merging |
|
@intel/llvm-gatekeepers yep, I believe this one is ready. Please merge |
|
@aelovikov-intel updated |
7 participants
No description provided.