1//===-- GPU implementation of the nanosleep function ----------------------===//
2//
3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4// See https://llvm.org/LICENSE.txt for license information.
5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6//
7//===----------------------------------------------------------------------===//
8
9#include "src/time/nanosleep.h"
10
11#include "time_utils.h"
12
13namespace LIBC_NAMESPACE {
14
15constexpr uint64_t TICKS_PER_SEC = 1000000000UL;
16
17LLVM_LIBC_FUNCTION(int, nanosleep,
18 (const struct timespec *req, struct timespec *rem)) {
19 if (!GPU_CLOCKS_PER_SEC || !req)
20 return -1;
21
22 uint64_t nsecs = req->tv_nsec + req->tv_sec * TICKS_PER_SEC;
23 uint64_t tick_rate = TICKS_PER_SEC / GPU_CLOCKS_PER_SEC;
24
25 uint64_t start = gpu::fixed_frequency_clock();
26#if defined(LIBC_TARGET_ARCH_IS_NVPTX)
27 uint64_t end = start + (nsecs + tick_rate - 1) / tick_rate;
28 uint64_t cur = gpu::fixed_frequency_clock();
29 // The NVPTX architecture supports sleeping and guaruntees the actual time
30 // slept will be somewhere between zero and twice the requested amount. Here
31 // we will sleep again if we undershot the time.
32 while (cur < end) {
33 if (__nvvm_reflect("__CUDA_ARCH") >= 700)
34 LIBC_INLINE_ASM("nanosleep.u32 %0;" ::"r"(nsecs));
35 cur = gpu::fixed_frequency_clock();
36 nsecs -= nsecs > cur - start ? cur - start : 0;
37 }
38#elif defined(LIBC_TARGET_ARCH_IS_AMDGPU)
39 uint64_t end = start + (nsecs + tick_rate - 1) / tick_rate;
40 uint64_t cur = gpu::fixed_frequency_clock();
41 // The AMDGPU architecture does not provide a sleep implementation with a
42 // known delay so we simply repeatedly sleep with a large value of ~960 clock
43 // cycles and check until we've passed the time using the known frequency.
44 __builtin_amdgcn_s_sleep(2);
45 while (cur < end) {
46 __builtin_amdgcn_s_sleep(15);
47 cur = gpu::fixed_frequency_clock();
48 }
49#else
50 // Sleeping is not supported.
51 if (rem) {
52 rem->tv_sec = req->tv_sec;
53 rem->tv_nsec = req->tv_nsec;
54 }
55 return -1;
56#endif
57 uint64_t stop = gpu::fixed_frequency_clock();
58
59 // Check to make sure we slept for at least the desired duration and set the
60 // remaining time if not.
61 uint64_t elapsed = (stop - start) * tick_rate;
62 if (elapsed < nsecs) {
63 if (rem) {
64 rem->tv_sec = (nsecs - elapsed) / TICKS_PER_SEC;
65 rem->tv_nsec = (nsecs - elapsed) % TICKS_PER_SEC;
66 }
67 return -1;
68 }
69
70 return 0;
71}
72
73} // namespace LIBC_NAMESPACE
74

source code of libc/src/time/gpu/nanosleep.cpp