[libc-commits] [libc] 27a80fc - [libc] Replace use of `asm` in the GPU code with LIBC_INLINE_ASM

[Joseph Huber via libc-commits]

Author: Joseph Huber
Date: 2023-06-06T14:24:52-05:00
New Revision: 27a80fc946186257899bce5256951b4e35b7e8c3

URL: https://github.com/llvm/llvm-project/commit/27a80fc946186257899bce5256951b4e35b7e8c3
DIFF: https://github.com/llvm/llvm-project/commit/27a80fc946186257899bce5256951b4e35b7e8c3.diff

LOG: [libc] Replace use of `asm` in the GPU code with LIBC_INLINE_ASM

We should more consistently use inline assembly using the LIBC wrappers.
It's much safer to mark all of these volatile as well.

Reviewed By: lntue

Differential Revision: https://reviews.llvm.org/D152294

Added: 
    

Modified: 
    libc/src/__support/GPU/nvptx/utils.h
    libc/src/__support/OSUtil/gpu/quick_exit.cpp
    libc/src/__support/RPC/rpc_util.h

Removed: 
    


################################################################################
diff  --git a/libc/src/__support/GPU/nvptx/utils.h b/libc/src/__support/GPU/nvptx/utils.h
index 5460c3da2885e..560f2c0733a04 100644
--- a/libc/src/__support/GPU/nvptx/utils.h
+++ b/libc/src/__support/GPU/nvptx/utils.h
@@ -100,7 +100,7 @@ LIBC_INLINE uint32_t get_lane_size() { return LANE_SIZE; }
 /// Returns the bit-mask of active threads in the current warp.
 [[clang::convergent]] LIBC_INLINE uint64_t get_lane_mask() {
   uint32_t mask;
-  asm volatile("activemask.b32 %0;" : "=r"(mask));
+  LIBC_INLINE_ASM("activemask.b32 %0;" : "=r"(mask));
   return mask;
 }
 

diff  --git a/libc/src/__support/OSUtil/gpu/quick_exit.cpp b/libc/src/__support/OSUtil/gpu/quick_exit.cpp
index 5cee9010bbf6c..6432965a741f0 100644
--- a/libc/src/__support/OSUtil/gpu/quick_exit.cpp
+++ b/libc/src/__support/OSUtil/gpu/quick_exit.cpp
@@ -24,11 +24,11 @@ void quick_exit(int status) {
   port.close();
 
 #if defined(LIBC_TARGET_ARCH_IS_NVPTX)
-  asm("exit;" ::: "memory");
+  LIBC_INLINE_ASM("exit;" ::: "memory");
 #elif defined(LIBC_TARGET_ARCH_IS_AMDGPU)
   // This will terminate the entire wavefront, may not be valid with divergent
   // work items.
-  asm("s_endpgm" ::: "memory");
+  __builtin_amdgcn_endpgm();
 #endif
   __builtin_unreachable();
 }

diff  --git a/libc/src/__support/RPC/rpc_util.h b/libc/src/__support/RPC/rpc_util.h
index 67a509c6499b7..d8cb88d2a23fb 100644
--- a/libc/src/__support/RPC/rpc_util.h
+++ b/libc/src/__support/RPC/rpc_util.h
@@ -23,7 +23,7 @@ constexpr uint64_t MAX_LANE_SIZE = 64;
 /// Suspend the thread briefly to assist the thread scheduler during busy loops.
 LIBC_INLINE void sleep_briefly() {
 #if defined(LIBC_TARGET_ARCH_IS_NVPTX) && __CUDA_ARCH__ >= 700
-  asm("nanosleep.u32 64;" ::: "memory");
+  LIBC_INLINE_ASM("nanosleep.u32 64;" ::: "memory");
 #elif defined(LIBC_TARGET_ARCH_IS_AMDGPU)
   __builtin_amdgcn_s_sleep(2);
 #else