llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-libc Author: Alex MacLean (AlexMaclean) <details> <summary>Changes</summary> --- Full diff: https://github.com/llvm/llvm-project/pull/198815.diff 4 Files Affected: - (modified) clang/lib/Headers/nvptxintrin.h (+1-1) - (modified) libc/shared/rpc_util.h (+1-1) - (modified) libc/src/__support/threads/sleep.h (+1-1) - (modified) libc/src/time/gpu/nanosleep.cpp (+1-1) ``````````diff 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; } `````````` </details> https://github.com/llvm/llvm-project/pull/198815 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
