[KHR Ext ?] Device side timer (equivalent to OpenGL ARB_shader_clock)

[tdavidcl]

Extension Type

New Feature

Feature Description

Hi,

With this issue, I wanted to suggest the addition of a device-side timer in SYCL. This can be incredibly useful for profiling to diagnose load balancing and starvation issues. Currently, to my knowledge, no such thing is available natively in SYCL.

Ideally I would propose the following:

• sycl::khr::device_globaltimer() -> f64 returns the current timer in nanoseconds (to follow the footsteps of the %%globaltimer special register in PTX), but the choice of the actual unit is open to discussions.
• sycl::khr::device_clock() -> u64 returns the current number of clock cycles (%%clock in PTX)

In term of implementation, I already have a draft supporting both Acpp and intel/llvm on Nvidia & host. For Intel @TApplencourt pointed me to some snippets suggesting that it is feasible quite directly. The only annoying one is AMD GPUs where we must extrapolate the number of seconds from the clock frequency without direct access to a "true" timer.

Here is a header only implementation (Host+CUDA) that I am already using inside Shamrock (although the name aren't aligned with what I'm proposing here)

https://github.com/Shamrock-code/Shamrock/blob/6e8b2e0ac528ae9976ec3903585be26311586676/src/shambackends/include/shambackends/intrisics/get_device_clock.hpp

Related Functionality in C++

high_resolution_clock::now() is the closest, for the actual timestamp counter it is not c++ but the rdtsc instruction

#include <chrono>
#include <iostream>

// GNU C++ x86_64 (thanks wikipedia)
extern inline uint64_t rdtsc() {
    uint64_t a, d;
    __asm__ volatile ("rdtsc" : "=a" (a), "=d" (d));
    return (d<<32) | a;
}

int main(){
    using namespace std::chrono;
    auto device_globaltimer = duration_cast<nanoseconds>(high_resolution_clock::now().time_since_epoch()).count();
    auto device_clock = rdtsc();

    std::cout << device_globaltimer << " " << device_clock << std::endl;
}

Godbolt link

Related Functionality in Other Languages

OpenGL: https://registry.khronos.org/OpenGL/extensions/ARB/ARB_shader_clock.txt
Vulkan: https://docs.vulkan.org/refpages/latest/refpages/source/VK_KHR_shader_clock.html

Related SYCL Extensions

None to my knowledge

[zjin-lcf]

Can your cuda/hip implementions be added to clock.hpp ?

https://github.com/intel/llvm/blob/sycl/sycl/include/sycl/ext/oneapi/experimental/clock.hpp

Reference
intel/llvm#21246

[TApplencourt]

Oh thankz @zjin-lcf , link to the Intel extension: https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_clock.asciidoc

[TApplencourt]

(For OpenCL/Intel clock cycle: https://github.com/intel/intel-graphics-compiler/blob/e72b08379c1402f2d9971053ef6d2cae0e1f7880/IGC/BiFModule/Languages/OpenCL/IBiF_Impl.cl#L391C1-L391C49)

[TApplencourt]

Proposal: Transform https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_clock.asciidoc into a KHR.

• Question: How to implement it in AMD

Are you ok with that @tdavidcl ?

[tdavidcl]

TL;DR: I like the oneapi ext, I would just add the nanotimer notion to it and I have a few nitpicks.

Two part answer:

---- clock ----

I quite like the notion of scope in the oneapi ext. Indeed, especially on CPU the frequency can not be guaranteed to be unique across work items / work groups.

So I'm fine with the following:

namespace sycl::khr {

     enum class aspect : /*unspecified*/ {
       khr_clock_sub_group,
       khr_clock_work_group,
       khr_clock_device
     };

     enum class clock_scope : /* unspecified */ {
         sub_group,
         work_group,
         device
     };

     template <clock_scope scope> uint64_t clock();
}

I would like to understand (and answer a few questions) in practice what would an implementation do if they want to implement the device scope for example on CPU where the frequencies are not matched (not provide khr_clock_device or emulate it ? should we forbid emulation to guarantee that such clock is accurate, as in not skip a cycle ?).

Small nitpick on the oneapi ext:
The clock is defined as an unbounded, unsigned integer counter that monotonically increments over time.
I would add the notion that two subsequent call can not return the same value, and thus enforce that the clock advance, otherwise returning always 0 is valid.

---- nanotimer ----

Here i'd like to precise the point. Initially my use case is that function rather than clock which I included in the issue since it relates to the same concept. The PTX %%globaltimer returns actually a value in nanosecond and thus can not be used to infer the device clock value if we want to guarantee its forward progression.

This function would be similar in definition to clock with the exception that it adds an aspect specifying if the clock is emulated or native and does not guarantee forward progression because of conversion issues. Additionally, we should probably not enforce the accuracy as the frequency is not guaranteed in emulated mode and we have no control over the vendor API.

• emulated = clock() / (average or real frequency) (with a div overhead sure but at least it is available)
• native = directly from vendor API

So I would propose the following addition to the original oneapi ext:

namespace sycl::khr {

     enum class aspect : /*unspecified*/ {
       khr_clock_sub_group,
       khr_clock_work_group,
       khr_clock_device,
       // duplicates of the clock variants
       khr_timer_sub_group,
       khr_timer_work_group,
       khr_timer_device,
       // specify native timer or extrapolated. Should it also be per scope ???
       khr_timer_emulated,
     };

     enum class clock_scope : /* unspecified */ {
         sub_group,
         work_group,
         device
     };

     template <clock_scope scope> uint64_t clock();
     template <clock_scope scope> double nanotimer(); // nanoseconds quite explicitly
}

A precision here is that I would prefer having a raw nanotimer rather than something that feeds into c++ <chrono> API to guarantee no overhead if a native one is available. If people want to convert it they can always feed it to std::chrono. (the PTX %%globaltimer overhead is so low that I can not mesure it when I use it twice in a kernel).
Another advantage is that the API is trivially usable for the user without the need to fiddle with <chrono>.

Lastly I would also add (maybe ?) a default template parameter to simplify usage -> device if available otherwise work-group if avail otherwise sub_group. But I totally get that the WG would not want that to be the case.

Lastly (yeah for the second time :) ), do you think that the notion of timer requires a separate KHR or is it fine to bundle both concepts into a single one. sycl_khr_clocks (to follow the naming of the OpenGL ext) vs sycl_khr_clock & sycl_khr_nanotimer ?

[tdavidcl]

Can your cuda/hip implementions be added to clock.hpp ?

https://github.com/intel/llvm/blob/sycl/sycl/include/sycl/ext/oneapi/experimental/clock.hpp

Reference intel/llvm#21246

Sure as long as I or Shamrock is credited for the snippet. Also be aware that as mentioned in the comment just above, %%globaltimer returns a value in nanoseconds rather than a number of cycles so the conversion to the global scope of clock() is not guaranteed to have forward progression which would violate the definition change I'm proposing in my comment.

[tdavidcl]

• Question: How to implement it in AMD

For AMD the "official" way of doing it in hipify for example and the few projects I looked at is by emulating the timer and returning clock() / frequency hence why I'm proposing that emulated clock aspect.

In summary for nanotimer:

• Nvidia -> %%globaltimer
• AMD -> (fancy register name) / frequency
• Intel -> already implemented apparently but I don't know how
• CPU -> std::chrono

[tdavidcl]

Yup Nvidia %%clock64 is per SM (https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#time-function).

[TApplencourt]

https://nvidia.github.io/cccl/unstable/libcudacxx/standard_api/time_library.html#implementation-defined-behavior

[gmlueck]

I would like to understand (and answer a few questions) in practice what would an implementation do if they want to implement the device scope for example on CPU where the frequencies are not matched (not provide khr_clock_device or emulate it ? should we forbid emulation to guarantee that such clock is accurate, as in not skip a cycle ?).

The sycl_ext_oneapi_clock extension does not require you to implement the device-scope clock. That's why the extension has the aspects -- it allows an implementation to only provide those clocks that are possible to implement. In your example, where the hardware has no support for a device-scope clock, I'd expect an implementation to return false for aspect::ext_oneapi_clock_device.

Regarding emulation ... the WG has discussed this in the past about other APIs. In general, we don't want to forbid implementations from emulating APIs so long as the emulation satisfies the specification. For one thing, it's hard to precisely define what "emulation" really means. Often hardware features do not exactly match the SYCL API, so you need some software manipulation. Even something as simple as throwing an exception for illegal input is handled in software not hardware. Are these bits of software code "emulation"?

For the same reason, the WG has rejected proposals to add SYCL APIs that query whether an API is "emulated". There's all sorts of grey areas, and it's not clear where native support stops and emulation begins.

I would add the notion that two subsequent call can not return the same value, and thus enforce that the clock advance, otherwise returning always 0 is valid.

I think we should not do this. What if a device has a a real timer-based register (i.e. a register that reads time, not cycles)? Your proposal would make it impossible to implement the extension on that hardware because two reads of the register might return the same value (if the granularity of the device clock is less than the time it takes to make two back-to-back reads of the register).

I understand that an implementation could satisfy the spec by always returning 0, but this would be a stupid implementation. I think it is not the purpose of the spec to prevent implementations that are low-quality. If customers are unhappy with a vendor's implementation, they can complain to the vendor (or use a different vendor's implementation).

I would just add the nanotimer notion to it ...

I'm not sure how many vendors can implement this. The clock on Intel GPUs is not a nanotimer, so I think we would not be able to provide this API. Even Nvidia does not seem to guarantee that %globaltimer will always be available. This is from the PTX specification:

Special registers intended for use by NVIDIA tools. The behavior is target-specific and may change or be removed in future GPUs. When JIT-compiled to other targets, the value of these registers is unspecified.

(https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#special-registers-globaltimer)

[gmlueck]

I think no one linked the SPIR-V extension yet:

https://github.khronos.org/SPIRV-Registry/extensions/KHR/SPV_KHR_shader_clock.html

The DPC++ sycl_ext_oneapi_clock extension is quite similar to this SPIR-V extension. In particular, the SPIR-V extension is also a monotonic counter, not a nanosecond timer.

[tdavidcl]

I asked around to people involved in Vulkan and OpenGL they rightfully pointed me to
https://github.com/KhronosGroup/GLSL/blob/main/extensions/ext/GL_EXT_shader_realtime_clock.txt

The clockRealtimeEXT() function returns a 64-bit value representing a
    real-time clock that is globally coherent by all invocations on the GPU.
    clockRealtime2x32EXT() returns the same value encoded as a two-component
    vector of 32-bit unsigned integers with the first component containing
    the 32 least significant bits and the second component containing the 
    32 most significant bits.

    The units of time are not defined for either of these operations and will
    wrap after exceeding the maximum value representable in 64 bits. These
    functions serve as code motion barriers

We could provide a device property to query the ratio if known, and as such call the function clock_realtime() for consistency

In my head that would yield

namespace sycl::info::device {
    // return the ratio relative to the second if known
    struct khr_clock_realtime_ratio {
        using return_type = uint64_t;  
    };
}

namespace sycl::khr {

     enum class aspect : /*unspecified*/ {
       khr_clock_sub_group,
       khr_clock_work_group,
       khr_clock_device,

       khr_clock_realtime,
       // specify if provided directly with the backend or by the SYCL implementation
       khr_vendor_provided_clock_realtime,
     };

     enum class clock_scope : /* unspecified */ {
         sub_group,
         work_group,
         device
     };

     template <clock_scope scope> uint64_t clock();
     
     // the scope define the common reference time point now
     template <clock_scope scope> uint64_t clock_realtime(); 
     template <clock_scope scope> sycl::vec<uint32_t,2> clock_realtime_2x32(); 
}

and say in the spec

The units of time can be provided via the khr_clock_realtime_ratio device property 
by the implementation if known. Either operations and will  wrap after exceeding 
the maximum value representable in 64 bits.

Also in OpenGL they are two distinct ext whereas there are combined in VK and Spir from what i understood. So there is probably a question of wether we want them combined or separated in SYCL (i would be leaning toward combined since that is what is done in Spir-V Source).

This extension adds new capabilities to support the GL_ARB_shader_clock
and GL_EXT_shader_realtime_clock GLSL extensions in SPIR-V.

To answer your comments :

I think we should not do this. What if a device has a a real timer-based register (i.e. a register that reads time, not cycles)? Your proposal would make it impossible to implement the extension on that hardware because two reads of the register might return the same value

Fair point i agree

For the same reason, the WG has rejected proposals to add SYCL APIs that query whether an API is "emulated". There's all sorts of grey areas, and it's not clear where native support stops and emulation begins.

Sure that could add confusion, what about vendor provided vs implementation provided ? While for other features emulation might be unsuitable, when it come to profiling, it is the kind of thing where one would hope for a fallback mechanism if possible. I'm especially thinking of computer graphics people wanting to mesure time of flight for the kernels work groups.

I'm not sure how many vendors can implement this. The clock on Intel GPUs is not a nanotimer, so I think we would not be able to provide this API. Even Nvidia does not seem to guarantee that %globaltimer will always be available.

Nvidia (source) yes and the range is from ptx3.0 so all version supported by SYCL compilers.
I do not have enough knowledge to speak about the intel stack.
For AMD while having a quick look in rccl they provide the following:
https://github.com/ROCm/rccl/blob/57e58688f44c77076ad536ef1f6b68741fc6e694/src/device/common.h#L411-415
(so realtime is aliased to clock64 which is constant 1GHz from what i understood when looking around), in that case we can provide clock_realtime but not define the ratio

[gmlueck]

I looked at the two OpenGL extensions you referenced:

• https://registry.khronos.org/OpenGL/extensions/ARB/ARB_shader_clock.txt
• https://github.com/KhronosGroup/GLSL/blob/main/extensions/ext/GL_EXT_shader_realtime_clock.txt

It looks to me like these are both monotonic counters, not nanosecond timers. Therefore, they are consistent with the SPIR-V extension (SPV_KHR_shader_clock) and the DPC++ extension (sycl_ext_oneapi_clock). None of these provide a way to measure time. Instead, they provide a way to count some abstract units, which I presume would be clock cycles on most implementations.

In particular, see this statement in ARB_shader_clock:

The units of time are not defined and need not be constant.

I think this is the fundamental issue with many devices. The clock frequency is not fixed, so it can vary even when a kernel is executing. As a result, there's no fixed ratio between clock cycles and time.