[libc-commits] [PATCH] D152294: [libc] Replace use of `asm` in the GPU code with LIBC_INLINE_ASM

[Joseph Huber via Phabricator via libc-commits]

jhuber6 created this revision.
jhuber6 added reviewers: lntue, michaelrj, sivachandra.
Herald added subscribers: libc-commits, mattd, asavonic, tschuett.
Herald added projects: libc-project, All.
jhuber6 requested review of this revision.

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


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D152294

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


Index: libc/src/__support/RPC/rpc_util.h
===================================================================
--- libc/src/__support/RPC/rpc_util.h
+++ libc/src/__support/RPC/rpc_util.h
@@ -23,7 +23,7 @@
 /// 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
Index: libc/src/__support/OSUtil/gpu/quick_exit.cpp
===================================================================
--- libc/src/__support/OSUtil/gpu/quick_exit.cpp
+++ libc/src/__support/OSUtil/gpu/quick_exit.cpp
@@ -24,11 +24,11 @@
   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();
 }
Index: libc/src/__support/GPU/nvptx/utils.h
===================================================================
--- libc/src/__support/GPU/nvptx/utils.h
+++ libc/src/__support/GPU/nvptx/utils.h
@@ -100,7 +100,7 @@
 /// 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;
 }
 


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D152294.528956.patch
Type: text/x-patch
Size: 1662 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/libc-commits/attachments/20230606/3d71ce22/attachment.bin>