[libc] Add basic utility support for timing functions on the GPU

[jhuber6]

Summary

This patch adds the utilities for the clocks on the GPU. This is done
prior to exporting it via some other interface and is mainly just done
so they are availible if we wish to do internal testing.

[JonChesterfield]

It's not obvious to me that this is a good thing overall.

First off benchmarking is difficult to do meaningfully and adding easy to access timing functions makes it more likely that they'll be used to make superficial comparisons when changing implementations.

Secondly picking timers that run at different frequencies on different hardware will interact badly with people comparing performance between different machines.

What's the benefit side to this like?

[jhuber6]

In D153388#4436410, @JonChesterfield wrote:

It's not obvious to me that this is a good thing overall.

First off benchmarking is difficult to do meaningfully and adding easy to access timing functions makes it more likely that they'll be used to make superficial comparisons when changing implementations.

Secondly picking timers that run at different frequencies on different hardware will interact badly with people comparing performance between different machines.

What's the benefit side to this like?

This is internal to libc and currently won't be exported as is. We could potentially put some of this into some time.h like functionality but that's tangential. The frequencies are different so we could potentially wrap them into different flags. Another option is to get the loader to look up the frequency and write it to some global so we can measure the nanosecond timer consistently between both architectures. It's useful because we want to be able to measure time elapsed on a much finger-grain than nvprof or rocprof would provide. This makes sure we can embed this analysis into some common tester and check the cycles.

[arsenm]

libc/src/__support/GPU/amdgpu/utils.h

159

Does __has_builtin cover subtarget feature restrictions?

[jhuber6]

Using __has_builtin

[arsenm]

needs some tests targeting the different targets, I don't know if __has_builtin will really work this way

[jhuber6]

In D153388#4436551, @arsenm wrote:

needs some tests targeting the different targets, I don't know if __has_builtin will really work this way

Not sure how to stand up a test just yet. In the future I would like to add a global to normalize the timestamp to a nanosecond resolution. I checked briefly via https://godbolt.org/z/4xjxjTovb and figured it worked.

[sivachandra]

libc/src/__support/GPU/amdgpu/utils.h

155

For consistency, use LIBC_HAS_BUILTIN: https://github.com/llvm/llvm-project/blob/main/libc/src/__support/macros/config.h#L29

[jhuber6]

-> LIBC_HAS_BUILTIN

[arsenm]

libc/src/__support/GPU/amdgpu/utils.h

159

Should check for has s_memtime in case future hardware decides 4th time's the charm on timing instructions

[jhuber6]

Add another check that uses the s_memtime if it's availible. This is because
on RDNA2 we still have s_memtime but also have the read cycle counter. The
read cycle counter implementation is 20-bits so we should favor the 64-bit one
if availible.

[JonChesterfield]

Test could be C to IR I think, check the ASM comes out for ptx and that some known architectures pick the right value for gcn.
I'm not sure we can do much with the generic case.

[JonChesterfield]

@arsenm what do you think of adding a built-in that expands to one of the above? This sort of "try to get the time" isn't libc specific and the switch over different builtins is basically architecture mess that ideally wouldn't get reimplemented in lot of different library code.

[jhuber6]

In D153388#4437661, @JonChesterfield wrote:

Test could be C to IR I think, check the ASM comes out for ptx and that some known architectures pick the right value for gcn.
I'm not sure we can do much with the generic case.

And IR test wouldn't fit in libc, that'd probably be more of a clang thing to check that we specify the built-ins properly.

In D153388#4437695, @JonChesterfield wrote:

@arsenm what do you think of adding a built-in that expands to one of the above? This sort of "try to get the time" isn't libc specific and the switch over different builtins is basically architecture mess that ideally wouldn't get reimplemented in lot of different library code.

It would definitely allow us to remove another cause of divergence for codes that want to be more generic w.r.t. the architecture, but they do slightly different things.

[arsenm]

In D153388#4437695, @JonChesterfield wrote:

@arsenm what do you think of adding a built-in that expands to one of the above? This sort of "try to get the time" isn't libc specific and the switch over different builtins is basically architecture mess that ideally wouldn't get reimplemented in lot of different library code.

That's kind of what I thought __builtin_readcyclecounter was for. I don't really understand the difference between the different time operations (e.g what is the difference between clock and time?)

[jhuber6]

In D153388#4437714, @arsenm wrote:

In D153388#4437695, @JonChesterfield wrote:

@arsenm what do you think of adding a built-in that expands to one of the above? This sort of "try to get the time" isn't libc specific and the switch over different builtins is basically architecture mess that ideally wouldn't get reimplemented in lot of different library code.

That's kind of what I thought __builtin_readcyclecounter was for. I don't really understand the difference between the different time operations (e.g what is the difference between clock and time?)

The boards have a fixed-frequency clock and a cycle counter. The fixed-frequency clock from time is used for measuring absolute time elapsed because it is not matched to any frequency throttling or stalling the processor might do. The cycle counter from clock is the variable frequency clock that corresponds to the actual number of processor cycles that have elapsed since you last called it. In this patch I specifically choose to use memtime over readcyclecounter if available (As it is in RDNA2) because the former is 64-bits while the latter is only 20-bits. Getting the real time from the fixed-frequency clock requires querying the driver unfortunately.

[JonChesterfield]

Let's go very explicit on naming then. fixed_frequency_timer and guess_current_time_in_arbitrary_units or whatever these things actually do.

[jhuber6]

In D153388#4437762, @JonChesterfield wrote:

Let's go very explicit on naming then. fixed_frequency_timer and guess_current_time_in_arbitrary_units or whatever these things actually do.

Sure, how about processor_clock and fixed_frequency_clock?

[JonChesterfield]

Is that the right basis set? I wonder if we want cycle_count, frequency, width as the controls. It seems like PTX wants to do ns granularity which is useful and GCN needs the loader to write to a global to get the frequency value, and the type wraparound could be anywhere from 20 to 64

[jhuber6]

In D153388#4437780, @JonChesterfield wrote:

Is that the right basis set? I wonder if we want cycle_count, frequency, width as the controls. It seems like PTX wants to do ns granularity which is useful and GCN needs the loader to write to a global to get the frequency value, and the type wraparound could be anywhere from 20 to 64

It describes what the underlying counters are. We could mention that the fixed frequency is 1E9 on PTX and some runtime value on AMDGCN. Relative timings should be fine in either case. The 20 bits is unfortunate but only affects gfx11xx architectures since we still have s_memtime on gfx10xx.

[jhuber6]

Change names.

[arsenm]

libc/src/__support/GPU/amdgpu/utils.h

167

__builtin_amdgcn_s_memfixed_frequrency_clock doesn't exist, and is a typo. If we want to introduce a new readcyclecounter alternative it should be a fully generic intrinsic like readcyclecounter

[jhuber6]

Fix the dangers of using sed to do your edits.

[lntue]

libc/src/__support/GPU/nvptx/utils.h

140

nit: Can you change this to the internal macro LIBC_INLINE_ASM?

[jhuber6]

-> LIBC_INLINE_ASM

[lntue]

LGTM for me, but I'll let others to approve the GPU builtin / inline assembly usage.

libc/src/__support/GPU/amdgpu/utils.h

 //===-------------- AMDGPU implementation of GPU utils ----------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef LLVM_LIBC_SRC_SUPPORT_GPU_AMDGPU_IO_H
 #define LLVM_LIBC_SRC_SUPPORT_GPU_AMDGPU_IO_H
 
 #include "src/__support/common.h"
+#include "src/__support/macros/config.h"
 
 #include <stdint.h>
 
 namespace __llvm_libc {
 namespace gpu {
 
 /// The number of threads that execute in lock-step in a lane.
 constexpr const uint64_t LANE_SIZE = __AMDGCN_WAVEFRONT_SIZE;
[[clang::convergent]] LIBC_INLINE void sync_threads() {
   __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "workgroup");
 }
 
 /// Wait for all threads in the wavefront to converge, this is a noop on AMDGPU.
 [[clang::convergent]] LIBC_INLINE void sync_lane(uint64_t) {
   __builtin_amdgcn_wave_barrier();
 }
 
+/// Returns the current value of the GPU's processor clock.
+/// NOTE: The RDNA3 and RDNA2 architectures use a 20-bit cycle cycle counter.
+LIBC_INLINE uint64_t processor_clock() {
+  if constexpr (LIBC_HAS_BUILTIN(__builtin_amdgcn_s_memtime))
+    return __builtin_amdgcn_s_memtime();
+  else if constexpr (LIBC_HAS_BUILTIN(__builtin_readcyclecounter))
+    return __builtin_readcyclecounter();
+  else
+    return 0;
+}
+
+/// Returns a fixed-frequency timestamp. The actual frequency is dependent on
+/// the card and can only be queried via the driver.
+LIBC_INLINE uint64_t fixed_frequrency_clock() {
+  if constexpr (LIBC_HAS_BUILTIN(__builtin_amdgcn_s_sendmsg_rtnl))
+    return __builtin_amdgcn_s_sendmsg_rtnl(0x83);
+  else if constexpr (LIBC_HAS_BUILTIN(__builtin_amdgcn_s_memrealtime))
+    return __builtin_amdgcn_s_memrealtime();
+  else if constexpr (LIBC_HAS_BUILTIN(__builtin_amdgcn_s_memtime))
+    return __builtin_amdgcn_s_memtime();
+  else
+    return 0;
+}
+
 } // namespace gpu
 } // namespace __llvm_libc
 
 #endif

libc/src/__support/GPU/generic/utils.h

LIBC_INLINE uint64_t ballot(uint64_t lane_mask, bool x) {
   (void)lane_mask;
   return x;
 }
 
 LIBC_INLINE void sync_threads() {}
 
 LIBC_INLINE void sync_lane(uint64_t) {}
 
+LIBC_INLINE uint64_t processor_clock() { return 0; }
+
+LIBC_INLINE uint64_t fixed_frequrency_clock() { return 0; }
+
 } // namespace gpu
 } // namespace __llvm_libc
 
 #endif

libc/src/__support/GPU/nvptx/utils.h

 /// Waits for all the threads in the block to converge and issues a fence.
 [[clang::convergent]] LIBC_INLINE void sync_threads() { __syncthreads(); }
 
 /// Waits for all threads in the warp to reconverge for independent scheduling.
 [[clang::convergent]] LIBC_INLINE void sync_lane(uint64_t mask) {
   __nvvm_bar_warp_sync(mask);
 }
 
+/// Returns the current value of the GPU's processor clock.
+LIBC_INLINE uint64_t processor_clock() {
+  uint64_t timestamp;
+  LIBC_INLINE_ASM("mov.u64  %0, %%clock64;" : "=l"(timestamp));
+  return timestamp;
+}
+
+/// Returns a global fixed-frequency timer at nanosecond frequency.
+LIBC_INLINE uint64_t fixed_frequrency_clock() {
+  uint64_t nsecs;
+  LIBC_INLINE_ASM("mov.u64  %0, %%globaltimer;" : "=l"(nsecs));
+  return nsecs;
+}
+
 } // namespace gpu
 } // namespace __llvm_libc
 
 #endif