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

[jhuber6]

Summary

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

libc/src/__support/GPU/nvptx/utils.h

 /// Returns the id of the thread inside of a CUDA warp executing together.
 [[clang::convergent]] LIBC_INLINE uint32_t get_lane_id() {
   return get_thread_id() & (get_lane_size() - 1);
 }
 
 /// 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;
 }
 
 /// Copies the value from the first active thread in the warp to the rest.
 [[clang::convergent]] LIBC_INLINE uint32_t broadcast_value(uint32_t x) {
   // NOTE: This is not sufficient in all cases on Volta hardware or later. The
   // lane mask returned here is not always the true lane mask used by the
   // intrinsics in cases of incedental or enforced divergence by the user.

libc/src/__support/OSUtil/gpu/quick_exit.cpp

 void quick_exit(int status) {
   rpc::Client::Port port = rpc::client.open<rpc::EXIT>();
   port.send([&](rpc::Buffer *buffer) {
     reinterpret_cast<uint32_t *>(buffer->data)[0] = 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();
 }
 
 } // namespace __llvm_libc
 
 #endif // LLVM_LIBC_SRC_SUPPORT_OSUTIL_GPU_QUICK_EXIT_H

libc/src/__support/RPC/rpc_util.h

 namespace rpc {
 
 /// Maximum amount of data a single lane can use.
 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
   // Simply do nothing if sleeping isn't supported on this platform.
 #endif
 }
 
 /// Get the first active thread inside the lane.