[libc] Implement the 'nanosleep' function on the GPU

[jhuber6]

Summary

The GPU has the ability to sleep for very short periods of time. We can
map this to the existing nanosleep utility. This patch maps the
nanosleep utility to the existing hardware instructions as best as
possible.

Depends on D159118

[jhuber6]

Add docs

[sivachandra]

OK for structuring changes.

[JonChesterfield]

We should probably create llvm intrinsics for the ad hoc ptx assembly remaining. I don't think there's any inline asm for intrinsics in amdgpu, though there may still be some as control flow hacks

[jhuber6]

In D159225#4629508, @JonChesterfield wrote:

We should probably create llvm intrinsics for the ad hoc ptx assembly remaining. I don't think there's any inline asm for intrinsics in amdgpu, though there may still be some as control flow hacks

Definitely, probably a good job to round them all up and implement them since they're pretty simple. Which actually reminds me that I forgot to check for sm_70 in this so thanks for that.

[jhuber6]

Add default and change the unit test. The test requires C++20 but doesn't seem to actually use it so I removed it such that we can run it here.

libc/config/gpu/api.td

 include "config/public_api.td"
 
 include "spec/stdc.td"
+include "spec/posix.td"
 include "spec/gpu_ext.td"
 
 def StringAPI : PublicAPI<"string.h"> {
   let Types = ["size_t"];
 }
 
 def StdlibAPI : PublicAPI<"stdlib.h"> {
   let Types = [

 
 def IntTypesAPI : PublicAPI<"inttypes.h"> {
   let Types = ["imaxdiv_t"];
 }
 
 def TimeAPI : PublicAPI<"time.h"> {
   let Types = [
     "clock_t",
+    "time_t",
+    "struct timespec",
   ];
 }

libc/config/gpu/entrypoints.txt

set(TARGET_LIBC_ENTRYPOINTS
     # inttypes.h entrypoints
     libc.src.inttypes.imaxabs
     libc.src.inttypes.imaxdiv
     libc.src.inttypes.strtoimax
     libc.src.inttypes.strtoumax
 
     # time.h entrypoints
     libc.src.time.clock
+    libc.src.time.nanosleep
 
     # gpu/rpc.h entrypoints
     libc.src.gpu.rpc_reset
     libc.src.gpu.rpc_host_call
 )
 
 set(TARGET_LIBM_ENTRYPOINTS
     # math.h entrypoints

libc/src/time/CMakeLists.txt

HDRS
     mktime.h
   DEPENDS
     .time_utils
     libc.include.time
     libc.src.errno.errno
 )
 
 add_entrypoint_object(
-  nanosleep
-  SRCS
-    nanosleep.cpp
-  HDRS
-    nanosleep.h
-  DEPENDS
-    libc.include.time
-    libc.include.sys_syscall
-    libc.src.__support.OSUtil.osutil
-    libc.src.errno.errno
-)
-
-add_entrypoint_object(
   time
   ALIAS
   DEPENDS
     .${LIBC_TARGET_OS}.time
 )
 
 add_entrypoint_object(
   clock
   ALIAS
   DEPENDS
     .${LIBC_TARGET_OS}.clock
 )
+
+add_entrypoint_object(
+  nanosleep
+  ALIAS
+  DEPENDS
+    .${LIBC_TARGET_OS}.nanosleep
+)

libc/src/time/gpu/CMakeLists.txt

SRCS
     clock.cpp
   HDRS
     ../clock.h
   DEPENDS
     libc.include.time
     libc.src.__support.GPU.utils
     .time_utils
 )
+
+add_entrypoint_object(
+  nanosleep
+  SRCS
+    nanosleep.cpp
+  HDRS
+    ../nanosleep.h
+  DEPENDS
+    libc.include.time
+    libc.src.__support.GPU.utils
+    .time_utils
+)

libc/src/time/gpu/nanosleep.cpp

+//===-- GPU implementation of the nanosleep function ----------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#include "src/time/nanosleep.h"
+
+#include "time_utils.h"
+
+namespace __llvm_libc {
+
+constexpr uint64_t TICKS_PER_NS = 1000000000UL;
+
+LLVM_LIBC_FUNCTION(int, nanosleep,
+                   (const struct timespec *req, struct timespec *rem)) {
+  if (!GPU_CLOCKS_PER_SEC || !req)
+    return -1;
+
+  uint64_t nsecs = req->tv_nsec + req->tv_sec * TICKS_PER_NS;
+
+  uint64_t start = gpu::fixed_frequency_clock();
+#if defined(LIBC_TARGET_ARCH_IS_NVPTX)
+  LIBC_INLINE_ASM("nanosleep.u32 %0;" ::"r"(nsecs));
+#elif defined(LIBC_TARGET_ARCH_IS_AMDGPU)
+  uint64_t end = start + nsecs / (TICKS_PER_NS / GPU_CLOCKS_PER_SEC);
+  uint64_t cur = gpu::fixed_frequency_clock();
+  // The AMDGPU architecture does not provide a sleep implementation with a
+  // known delay so we simply repeatedly sleep with a large value of ~960 clock
+  // cycles and check until we've passed the time using the known frequency.
+  __builtin_amdgcn_s_sleep(2);
+  while (cur < end) {
+    __builtin_amdgcn_s_sleep(15);
+    cur = gpu::fixed_frequency_clock();
+  }
+#endif
+  uint64_t stop = gpu::fixed_frequency_clock();
+
+  // Check to make sure we slept for at least the desired duration and set the
+  // remaining time if not.
+  uint64_t elapsed = (stop - start) * (TICKS_PER_NS / GPU_CLOCKS_PER_SEC);
+  if (elapsed < nsecs) {
+    if (rem) {
+      rem->tv_sec = (nsecs - elapsed) / TICKS_PER_NS;
+      rem->tv_nsec = (nsecs - elapsed) % TICKS_PER_NS;
+    }
+    return -1;
+  }
+
+  return 0;
+}
+
+} // namespace __llvm_libc

libc/src/time/gpu/time_utils.h

 //===-- Generic utilities for GPU timing ----------------------------------===//
 //
 // 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_TIME_GPU_TIME_UTILS_H
 #define LLVM_LIBC_SRC_TIME_GPU_TIME_UTILS_H
 
 #include "src/__support/GPU/utils.h"
 
 namespace __llvm_libc {
 
-#if defined(LIBC_TARGET_ARCH_IS_GPU)
+#if defined(LIBC_TARGET_ARCH_IS_AMDGPU)
 // AMDGPU does not have a single set frequency. Different architectures and
 // cards can have vary values. Here we default to a few known values, but for
 // complete support the frequency needs to be read from the kernel driver.
 #if defined(__gfx1010__) || defined(__gfx1011__) || defined(__gfx1012__) ||    \
     defined(__gfx1013__) || defined(__gfx1030__) || defined(__gfx1031__) ||    \
     defined(__gfx1032__) || defined(__gfx1033__) || defined(__gfx1034__) ||    \
     defined(__gfx1035__) || defined(__gfx1036__) || defined(__gfx1100__) ||    \
     defined(__gfx1101__) || defined(__gfx1102__) || defined(__gfx1103__) ||    \

libc/src/time/gpu/time_utils.cpp

 //===-- Generic utilities for GPU timing ----------------------------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #include "time_utils.h"
 
 namespace __llvm_libc {
 
-#if defined(LIBC_TARGET_ARCH_IS_GPU)
+#if defined(LIBC_TARGET_ARCH_IS_AMDGPU)
 // This is expected to be initialized by the runtime if the default value is
 // insufficient.
 // TODO: Once we have another use-case for this we should put it in a common
 // device environment struct.
 extern "C" [[gnu::visibility("protected")]] uint64_t
     [[clang::address_space(4)]] __llvm_libc_clock_freq = clock_freq;
 #endif
 
 } // namespace __llvm_libc

libc/src/time/linux/CMakeLists.txt

HDRS
     ../clock.h
   DEPENDS
     libc.include.time
     libc.include.sys_syscall
     libc.src.__support.CPP.limits
     libc.src.__support.OSUtil.osutil
     libc.src.errno.errno
 )
+
+add_entrypoint_object(
+  nanosleep
+  SRCS
+    nanosleep.cpp
+  HDRS
+    ../nanosleep.h
+  DEPENDS
+    libc.include.time
+    libc.include.sys_syscall
+    libc.src.__support.CPP.limits
+    libc.src.__support.OSUtil.osutil
+    libc.src.errno.errno
+)

libc/src/time/linux/nanosleep.cpp

-//===-- Implementation of nanosleep function ------------------------------===//
+//===-- Linux implementation of nanosleep function ------------------------===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 
 #include "src/time/nanosleep.h"
 
 #include "src/__support/OSUtil/syscall.h" // For syscall functions.
 #include "src/__support/common.h"
 #include "src/errno/libc_errno.h"
 
 #include <sys/syscall.h> // For syscall numbers.
 
 namespace __llvm_libc {
 
-// TODO(michaelrj): Move this into time/linux with the other syscalls.
 LLVM_LIBC_FUNCTION(int, nanosleep,
                    (const struct timespec *req, struct timespec *rem)) {
 #if SYS_nanosleep
   int ret = __llvm_libc::syscall_impl<int>(SYS_nanosleep, req, rem);
 #elif defined(SYS_clock_nanosleep_time64)
   int ret =
       __llvm_libc::syscall_impl<int>(SYS_clock_nanosleep_time64, req, rem);
 #else