[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/docs/gpu/support.rst

 
 stdio.h
 --------
 
 =============  =========  ============
 Function Name  Available  RPC Required
 =============  =========  ============
 clock          |check|
+nanosleep      |check|
 =============  =========  ============

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) && __CUDA_ARCH__ >= 700
+  uint64_t end = start + nsecs / (TICKS_PER_NS / GPU_CLOCKS_PER_SEC);
+  uint64_t cur = gpu::fixed_frequency_clock();
+  // The NVPTX architecture supports sleeping and guaruntees the actual time
+  // slept will be somewhere between zero and twice the requested amount. Here
+  // we will sleep again if we undershot the time.
+  while (cur < end) {
+    LIBC_INLINE_ASM("nanosleep.u32 %0;" ::"r"(nsecs));
+    cur = gpu::fixed_frequency_clock();
+    nsecs -= nsecs > cur - start ? cur - start : 0;
+  }
+#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();
+  }
+#else
+  // Sleeping is not supported.
+  if (rem) {
+    rem->tv_sec = req->tv_sec;
+    rem->tv_nsec = req->tv_nsec;
+  }
+  return -1;
+#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/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

libc/test/src/time/CMakeLists.txt

HDRS
     TmHelper.h
     TmMatcher.h
   CXX_STANDARD
     20
   DEPENDS
     libc.src.time.mktime
 )
 
-add_libc_unittest(
+add_libc_test(
   nanosleep_test
   SUITE
     libc_time_unittests
   SRCS
     nanosleep_test.cpp
-  HDRS
-    TmHelper.h
-    TmMatcher.h
-  CXX_STANDARD
-    20
   DEPENDS
     libc.include.time
     libc.src.time.nanosleep
     libc.src.errno.errno
 )
 
 add_libc_unittest(
   time_test