[libc-commits] [clang] [libc] [libc][clang] Use __nvvm_nanosleep builtin over inline assembly (PR #198815)

[Alex MacLean via libc-commits]

https://github.com/AlexMaclean created https://github.com/llvm/llvm-project/pull/198815

None

>From fcff1538b383ea70c3edca8b23aad6cf5fe0e714 Mon Sep 17 00:00:00 2001
From: Alex Maclean <amaclean at nvidia.com>
Date: Wed, 20 May 2026 08:45:21 -0700
Subject: [PATCH] [libc][clang] Use __nvvm_nanosleep builtin over inline
 assembly

---
 clang/lib/Headers/nvptxintrin.h    | 2 +-
 libc/shared/rpc_util.h             | 2 +-
 libc/src/__support/threads/sleep.h | 2 +-
 libc/src/time/gpu/nanosleep.cpp    | 2 +-
 4 files changed, 4 insertions(+), 4 deletions(-)

diff --git a/clang/lib/Headers/nvptxintrin.h b/clang/lib/Headers/nvptxintrin.h
index df87cf4eaeaaa..6ccc31ae47d8e 100644
--- a/clang/lib/Headers/nvptxintrin.h
+++ b/clang/lib/Headers/nvptxintrin.h
@@ -199,7 +199,7 @@ _DEFAULT_FN_ATTRS [[noreturn]] static __inline__ void __gpu_exit(void) {
 // Suspend the thread briefly to assist the scheduler during busy loops.
 _DEFAULT_FN_ATTRS static __inline__ void __gpu_thread_suspend(void) {
   if (__nvvm_reflect("__CUDA_ARCH") >= 700)
-    asm("nanosleep.u32 64;" ::: "memory");
+    __nvvm_nanosleep(64);
 }
 
 _Pragma("omp end declare variant");
diff --git a/libc/shared/rpc_util.h b/libc/shared/rpc_util.h
index 09a97c63b562f..da18ad41d6a9d 100644
--- a/libc/shared/rpc_util.h
+++ b/libc/shared/rpc_util.h
@@ -334,7 +334,7 @@ RPC_ATTRS auto apply(F &&f, tuple<Ts...> &t) {
 RPC_ATTRS void sleep_briefly() {
 #if __has_builtin(__nvvm_reflect)
   if (__nvvm_reflect("__CUDA_ARCH") >= 700)
-    asm("nanosleep.u32 64;" ::: "memory");
+    __nvvm_nanosleep(64);
 #elif __has_builtin(__builtin_amdgcn_s_sleep)
   __builtin_amdgcn_s_sleep(2);
 #elif __has_builtin(__builtin_ia32_pause)
diff --git a/libc/src/__support/threads/sleep.h b/libc/src/__support/threads/sleep.h
index 777283356b58e..e05a43170be57 100644
--- a/libc/src/__support/threads/sleep.h
+++ b/libc/src/__support/threads/sleep.h
@@ -18,7 +18,7 @@ namespace LIBC_NAMESPACE_DECL {
 LIBC_INLINE void sleep_briefly() {
 #if defined(LIBC_TARGET_ARCH_IS_NVPTX)
   if (__nvvm_reflect("__CUDA_ARCH") >= 700)
-    LIBC_INLINE_ASM("nanosleep.u32 64;" ::: "memory");
+    __nvvm_nanosleep(64);
 #elif defined(LIBC_TARGET_ARCH_IS_AMDGPU)
   __builtin_amdgcn_s_sleep(2);
 #elif defined(LIBC_TARGET_ARCH_IS_X86)
diff --git a/libc/src/time/gpu/nanosleep.cpp b/libc/src/time/gpu/nanosleep.cpp
index d22d9d6bd8d79..bd479787eea60 100644
--- a/libc/src/time/gpu/nanosleep.cpp
+++ b/libc/src/time/gpu/nanosleep.cpp
@@ -31,7 +31,7 @@ LLVM_LIBC_FUNCTION(int, nanosleep,
   // we will sleep again if we undershot the time.
   while (cur < end) {
     if (__nvvm_reflect("__CUDA_ARCH") >= 700)
-      LIBC_INLINE_ASM("nanosleep.u32 %0;" ::"r"(nsecs));
+      __nvvm_nanosleep(nsecs);
     cur = gpu::fixed_frequency_clock();
     nsecs -= nsecs > cur - start ? cur - start : 0;
   }