[libc] Support timing information in libc tests

[jhuber6]

Summary

This patch adds the necessary support to provide timing information in
libc tests. This is useful for determining which tests look what
amount of time. We also can use this as a test basis for providing more
fine-grained timing when implementing things on the GPU.

The main difficulty with this is the fact that the AMDGPU fixed
frequency clock operates at an unknown frequency. We need to read this
on a per-card basis from the driver and then copy it in. NVPTX on the
other hand has a fixed clock at a resolution of 1ns. I have also
increased the resolution of the print-outs as the majority of these are
below a millisecond for me.

[jhuber6]

Format

[jhuber6]

Unify the memcpy into a common function

[JonChesterfield]

Let's take an educated guess that this won't be the one true unique global we end up doing this with and make it a struct up front

libc/startup/gpu/amdgpu/start.cpp

21

This will be in .bss. Is that a good thing? Means we have no option to initialise it before loading the image, i.e. we're stuck with some variant of crossing the bus and waiting to write four bytes

libc/test/UnitTest/LibcTest.cpp

22

Looks sketchy that this is a different type on the various platforms. Safer would be a uint64_t, or at least annotate the type so it isn't a signed integer.

Really not keen on it being a macro pretending to be part of time.h when it isn't in time.h and this is all tied up in testing the implementation of libc.

154

Please parenthesize to indicate what order of operations you want on the multiplication and division here, and if you don't want to give CLOCKS_PER_TYPE an actual type, static cast it to uint64_t at the use site so signed/unsigned integer nonsense can't bite us.

When this is a macro expanding to a non-const global, are we stuck with a separate load for each instance? I _think_ clang will constant fold it anyway because it isn't volatile or atomic but haven't checked.

libc/utils/gpu/loader/amdgpu/Loader.cpp

272

Boilerplate looks ok here, though I'd probably have gone with the synchronous copy instead (no need to mess around with signals then, though I have a vague recollection of there being other hazards with the non-async one)

442

All this stuff? Pretty much why I prefer writing to the global before copying it to the GPU

[jhuber6]

libc/startup/gpu/amdgpu/start.cpp

21

I should probably make this in constant memory.

libc/test/UnitTest/LibcTest.cpp

22

Honestly the whole CLOCKS_PER_SEC thing confuses me in general. POSIX says it's always 1,000,000, other platforms say it's something else. Other's scale the clock so that it matches 1,000,000... We could probably avoid it here if necessary and make it a constant or something.

154

I could probably make it const since it should be set by the loader / whoever's standing up the runtime.

libc/utils/gpu/loader/amdgpu/Loader.cpp

272

Yeah, didn't see any simpler ways to do this in the HSA interface.

442

I didn't want to pass it via the kernel arguments because then it would require changing the NVPTX plugin as well. Plus in the future we'll probably want to abstract this into some general environment constant.

[jhuber6]

Making suggested changes

[JonChesterfield]

libc/test/UnitTest/LibcTest.cpp

159

Overflow hazards abound here. I'm not totally sure what to do about them.

[jhuber6]

libc/test/UnitTest/LibcTest.cpp

159

These are uint64_t, if these overflow I think we'll be fine for just running some unit tests and getting a time that says 0ms instead of a few trillion.

[JonChesterfield]

Seems reasonable. Thanks

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

LIBC_INLINE uint64_t processor_clock() {
   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() {
+LIBC_INLINE uint64_t fixed_frequency_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 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; }
+LIBC_INLINE uint64_t fixed_frequency_clock() { return 0; }
 
 } // namespace gpu
 } // namespace __llvm_libc
 
 #endif

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

 /// 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() {
+LIBC_INLINE uint64_t fixed_frequency_clock() {
   uint64_t nsecs;
   LIBC_INLINE_ASM("mov.u64  %0, %%globaltimer;" : "=l"(nsecs));
   return nsecs;
 }
 
 } // namespace gpu
 } // namespace __llvm_libc
 
 #endif

libc/startup/gpu/amdgpu/start.cpp

 #include "src/__support/RPC/rpc_client.h"
 #include "src/stdlib/atexit.h"
 #include "src/stdlib/exit.h"
 
 extern "C" int main(int argc, char **argv, char **envp);
 
 namespace __llvm_libc {
 
+// The AMDGPU architecture provides a fixed frequency clock used for obtaining
+// real time. However, the frequency of this clock varies between cards and can
+// only be obtained via the driver. The loader will set this so we can use it.
+extern "C" [[gnu::visibility("protected")]] uint64_t
+    [[clang::address_space(4)]] __llvm_libc_clock_freq = 0;
+
 extern "C" uintptr_t __init_array_start[];
 extern "C" uintptr_t __init_array_end[];
 extern "C" uintptr_t __fini_array_start[];
 extern "C" uintptr_t __fini_array_end[];
 
 using InitCallback = void(int, char **, char **);
 using FiniCallback = void(void);
 

libc/test/UnitTest/LibcTest.cpp

 
 #include "src/__support/CPP/string.h"
 #include "src/__support/CPP/string_view.h"
 #include "src/__support/UInt128.h"
 #include "test/UnitTest/TestLogger.h"
 
 #if __STDC_HOSTED__
 #include <time.h>
+#elif defined(LIBC_TARGET_ARCH_IS_GPU)
+#include "src/__support/GPU/utils.h"
+static long clock() { return __llvm_libc::gpu::fixed_frequency_clock(); }
+#if LIBC_TARGET_ARCH_IS_NVPTX
+uint64_t CLOCKS_PER_SEC = 1000000000UL;
+#else
+// The AMDGPU loader needs to initialize this at runtime by querying the driver.
+extern "C" [[gnu::visibility("protected")]] uint64_t __llvm_libc_clock_freq;
+uint64_t CLOCKS_PER_SEC = __llvm_libc_clock_freq;
+#endif
 #else
 static long clock() { return 0; }
 #define CLOCKS_PER_SEC 1
 #endif
 
 namespace __llvm_libc {
 namespace testing {
 
for (Test *T = Start; T != nullptr; T = T->Next) {
     [[maybe_unused]] const auto end_time = clock();
     switch (Ctx.status()) {
     case RunContext::RunResult::Fail:
       tlog << RED << "[  FAILED  ] " << RESET << TestName << '\n';
       ++FailCount;
       break;
     case RunContext::RunResult::Pass:
       tlog << GREEN << "[       OK ] " << RESET << TestName;
-#if __STDC_HOSTED__
+#if __STDC_HOSTED__ || defined(LIBC_TARGET_ARCH_IS_GPU)
       tlog << " (took ";
       if (start_time > end_time) {
         tlog << "unknown - try rerunning)\n";
       } else {
         const auto duration = end_time - start_time;
-        const uint64_t duration_ms = duration * 1000 / CLOCKS_PER_SEC;
+        const uint64_t duration_ms = (duration * 1000) / CLOCKS_PER_SEC;
+        const uint64_t duration_us = (duration * 1000 * 1000) / CLOCKS_PER_SEC;
+        const uint64_t duration_ns =
+            (duration * 1000 * 1000 * 1000) / CLOCKS_PER_SEC;
+        if (duration_ms != 0)
           tlog << duration_ms << " ms)\n";
+        else if (duration_us != 0)
+          tlog << duration_us << " us)\n";
+        else
+          tlog << duration_ns << " ns)\n";
       }
 #else
       tlog << '\n';
 #endif
       break;
     }
     ++TestCount;
   }

libc/utils/gpu/loader/amdgpu/Loader.cpp

hsa_status_t launch_kernel(hsa_agent_t dev_agent, hsa_executable_t executable,
   if (hsa_status_t err = hsa_amd_memory_pool_free(args))
     handle_error(err);
   if (hsa_status_t err = hsa_signal_destroy(packet->completion_signal))
     handle_error(err);
 
   return HSA_STATUS_SUCCESS;
 }
 
+/// Copies data from the source agent to the destination agent. The source
+/// memory must first be pinned explicitly or allocated via HSA.
+static hsa_status_t hsa_memcpy(void *dst, hsa_agent_t dst_agent,
+                               const void *src, hsa_agent_t src_agent,
+                               uint64_t size) {
+  // Create a memory signal to copy information between the host and device.
+  hsa_signal_t memory_signal;
+  if (hsa_status_t err = hsa_signal_create(1, 0, nullptr, &memory_signal))
+    return err;
+
+  if (hsa_status_t err = hsa_amd_memory_async_copy(
+          dst, dst_agent, src, src_agent, size, 0, nullptr, memory_signal))
+    return err;
+
+  while (hsa_signal_wait_scacquire(memory_signal, HSA_SIGNAL_CONDITION_EQ, 0,
+                                   UINT64_MAX, HSA_WAIT_STATE_ACTIVE) != 0)
+    ;
+
+  if (hsa_status_t err = hsa_signal_destroy(memory_signal))
+    return err;
+
+  return HSA_STATUS_SUCCESS;
+}
+
 int load(int argc, char **argv, char **envp, void *image, size_t size,
          const LaunchParameters &params) {
   // Initialize the HSA runtime used to communicate with the device.
   if (hsa_status_t err = hsa_init())
     handle_error(err);
 
   // Register a callback when the device encounters a memory fault.
   if (hsa_status_t err = hsa_amd_register_system_event_handler(
if (hsa_status_t err = hsa_amd_memory_pool_allocate(finegrained_pool, size,
       handle_error(err);
     hsa_amd_agents_allow_access(1, &dev_agent, nullptr, dev_ptr);
     return dev_ptr;
   };
   if (rpc_status_t err = rpc_server_init(device_id, RPC_MAXIMUM_PORT_COUNT,
                                          wavefront_size, rpc_alloc, &tuple))
     handle_error(err);
 
+  // Obtain the GPU's fixed-frequency clock rate and copy it to the GPU.
+  void *host_clock_freq;
+  if (hsa_status_t err =
+          hsa_amd_memory_pool_allocate(finegrained_pool, sizeof(uint64_t),
+                                       /*flags=*/0, &host_clock_freq))
+    handle_error(err);
+  hsa_amd_agents_allow_access(1, &dev_agent, nullptr, host_clock_freq);
+
+  if (hsa_status_t err = hsa_agent_get_info(
+          dev_agent,
+          static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_TIMESTAMP_FREQUENCY),
+          host_clock_freq))
+    handle_error(err);
+
+  hsa_executable_symbol_t freq_sym;
+  if (hsa_status_t err = hsa_executable_get_symbol_by_name(
+          executable, "__llvm_libc_clock_freq", &dev_agent, &freq_sym))
+    handle_error(err);
+
+  void *freq_addr;
+  if (hsa_status_t err = hsa_executable_symbol_get_info(
+          freq_sym, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS, &freq_addr))
+    handle_error(err);
+
+  if (hsa_status_t err = hsa_memcpy(freq_addr, dev_agent, host_clock_freq,
+                                    host_agent, sizeof(uint64_t)))
+    handle_error(err);
+
   // Obtain a queue with the minimum (power of two) size, used to send commands
   // to the HSA runtime and launch execution on the device.
   uint64_t queue_size;
   if (hsa_status_t err = hsa_agent_get_info(
           dev_agent, HSA_AGENT_INFO_QUEUE_MIN_SIZE, &queue_size))
     handle_error(err);
   hsa_queue_t *queue = nullptr;
   if (hsa_status_t err =
if (hsa_status_t err = launch_kernel(
     handle_error(err);
 
   start_args_t args = {argc, dev_argv, dev_envp, dev_ret};
   if (hsa_status_t err =
           launch_kernel(dev_agent, executable, kernargs_pool,
                         coarsegrained_pool, queue, params, "_start.kd", args))
     handle_error(err);
 
-  // Create a memory signal and copy the return value back from the device into
-  // a new buffer.
-  hsa_signal_t memory_signal;
-  if (hsa_status_t err = hsa_signal_create(1, 0, nullptr, &memory_signal))
-    handle_error(err);
-
   void *host_ret;
   if (hsa_status_t err =
           hsa_amd_memory_pool_allocate(finegrained_pool, sizeof(int),
                                        /*flags=*/0, &host_ret))
     handle_error(err);
   hsa_amd_agents_allow_access(1, &dev_agent, nullptr, host_ret);
 
   if (hsa_status_t err =
-          hsa_amd_memory_async_copy(host_ret, host_agent, dev_ret, dev_agent,
-                                    sizeof(int), 0, nullptr, memory_signal))
+          hsa_memcpy(host_ret, host_agent, dev_ret, dev_agent, sizeof(int)))
     handle_error(err);
 
-  while (hsa_signal_wait_scacquire(memory_signal, HSA_SIGNAL_CONDITION_EQ, 0,
-                                   UINT64_MAX, HSA_WAIT_STATE_ACTIVE) != 0)
-    ;
-
   // Save the return value and perform basic clean-up.
   int ret = *static_cast<int *>(host_ret);
 
   end_args_t fini_args = {ret};
   if (hsa_status_t err = launch_kernel(
           dev_agent, executable, kernargs_pool, coarsegrained_pool, queue,
           single_threaded_params, "_end.kd", fini_args))
     handle_error(err);
 
   if (rpc_status_t err = rpc_server_shutdown(
           device_id, [](void *ptr, void *) { hsa_amd_memory_pool_free(ptr); },
           nullptr))
     handle_error(err);
 
   // Free the memory allocated for the device.
   if (hsa_status_t err = hsa_amd_memory_pool_free(dev_argv))
     handle_error(err);
   if (hsa_status_t err = hsa_amd_memory_pool_free(dev_ret))
     handle_error(err);
   if (hsa_status_t err = hsa_amd_memory_pool_free(host_ret))
     handle_error(err);
 
-  if (hsa_status_t err = hsa_signal_destroy(memory_signal))
-    handle_error(err);
   if (hsa_status_t err = hsa_queue_destroy(queue))
     handle_error(err);
 
   if (hsa_status_t err = hsa_executable_destroy(executable))
     handle_error(err);
 
   if (hsa_status_t err = hsa_code_object_destroy(object))
     handle_error(err);
 
   if (rpc_status_t err = rpc_shutdown())
     handle_error(err);
   if (hsa_status_t err = hsa_shut_down())
     handle_error(err);
 
   return ret;
 }