[libomptarget][nvptx] Replace cuda atomic primitives with clang intrinsics

[JonChesterfield]

Summary

[libomptarget][nvptx] Replace cuda atomic primitives with clang intrinsics

Tested by diff of IR generated for target_impl.cu before and after. NFC. Part
of removing deviceRTL build time dependency on cuda SDK.

[tianshilei1992]

openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.cu

142

what about using NVVM atomic intrinsics directly? We don't need the memory order then.

[tianshilei1992]

openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.cu

30–31

__syncwarp is left. It can be also simply replaced by __nvvm_bar_warp_sync(mask).

[JonChesterfield]

openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.cu

30–31

Syncwarp is one of the cuda_version intrinsics, will handle it with the rest of them.

142

Exposing memory order is a feature. Makes it clear we're using the slow one, gives the hook to change that if we wish.

Also gives an option to use the same clang intrinsics on amdgpu and nvptx if we wish.

[tianshilei1992]

openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.cu

142

That sounds appealing. Maybe we could use this patch to move atomic operations back to common part, and create another patch to rewrite other CUDA intrinsics related functions.

[JonChesterfield]

openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.cu

142

I'd rather land this separately to starting to merge parts of the two target_impl files.

I'm not sure whether seq_cst is the right order for nvptx. Amdgcn should probably be device scope + acquire/release, though it's presently seq_cst as that's what hip did by default. So we may want slightly different Intrinsics again later.

[tianshilei1992]

openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.cu

142

JFYI, I'm working on cleaning up code to remove target_impl.h now.

[JonChesterfield]

openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.cu

142

This is presently orthogonal to D95298 and removes about half of the remaining uses of the cuda sdk. I'd prefer land it independently of hashing out consensus on what the library structure should change to.

[tianshilei1992]

LG

openmp/libomptarget/deviceRTLs/nvptx/src/target_impl.cu

 unsigned int __activemask();
 unsigned int __ballot(unsigned);
 // The default argument here is based on NVIDIA's website
 // https://developer.nvidia.com/blog/using-cuda-warp-level-primitives/
 int __shfl_sync(unsigned mask, int val, int src_line, int width = WARPSIZE);
 int __shfl(int val, int src_line, int width = WARPSIZE);
 int __shfl_down(int var, unsigned detla, int width);
 int __shfl_down_sync(unsigned mask, int var, unsigned detla, int width);
 void __syncwarp(int mask);
-void __threadfence();
-void __threadfence_block();
-void __threadfence_system();
 }
 
 DEVICE void __kmpc_impl_unpack(uint64_t val, uint32_t &lo, uint32_t &hi) {
   asm volatile("mov.b64 {%0,%1}, %2;" : "=r"(lo), "=r"(hi) : "l"(val));
 }
 
 DEVICE uint64_t __kmpc_impl_pack(uint32_t lo, uint32_t hi) {
   uint64_t val;
   asm volatile("mov.b64 %0, {%1,%2};" : "=l"(val) : "r"(lo), "r"(hi));
DEVICE void __kmpc_impl_named_sync(uint32_t num_threads) {
   // region to synchronize with each other.
   int barrier = 1;
   asm volatile("bar.sync %0, %1;"
                :
                : "r"(barrier), "r"(num_threads)
                : "memory");
 }
 
-DEVICE void __kmpc_impl_threadfence() { __threadfence(); }
-DEVICE void __kmpc_impl_threadfence_block() { __threadfence_block(); }
-DEVICE void __kmpc_impl_threadfence_system() { __threadfence_system(); }
+DEVICE void __kmpc_impl_threadfence() { __nvvm_membar_gl(); }
+DEVICE void __kmpc_impl_threadfence_block() { __nvvm_membar_cta(); }
+DEVICE void __kmpc_impl_threadfence_system() { __nvvm_membar_sys(); }
 
 // Calls to the NVPTX layer (assuming 1D layout)
 DEVICE int GetThreadIdInBlock() { return __nvvm_read_ptx_sreg_tid_x(); }
 DEVICE int GetBlockIdInKernel() { return __nvvm_read_ptx_sreg_ctaid_x(); }
 DEVICE int GetNumberOfBlocksInKernel() {
   return __nvvm_read_ptx_sreg_nctaid_x();
 }
 DEVICE int GetNumberOfThreadsInBlock() { return __nvvm_read_ptx_sreg_ntid_x(); }
 DEVICE unsigned GetWarpId() { return GetThreadIdInBlock() / WARPSIZE; }
 DEVICE unsigned GetLaneId() { return GetThreadIdInBlock() & (WARPSIZE - 1); }
 
-// Forward declaration of atomics. Although they're template functions, we
-// already have definitions for different types in CUDA internal headers with
-// the right mangled names.
-template <typename T> DEVICE T atomicAdd(T *address, T val);
-template <typename T> DEVICE T atomicInc(T *address, T val);
-template <typename T> DEVICE T atomicMax(T *address, T val);
-template <typename T> DEVICE T atomicExch(T *address, T val);
-template <typename T> DEVICE T atomicCAS(T *address, T compare, T val);
-
+// Atomics
 DEVICE uint32_t __kmpc_atomic_add(uint32_t *Address, uint32_t Val) {
-  return atomicAdd(Address, Val);
+  return __atomic_fetch_add(Address, Val, __ATOMIC_SEQ_CST);
 }
 DEVICE uint32_t __kmpc_atomic_inc(uint32_t *Address, uint32_t Val) {
-  return atomicInc(Address, Val);
+  return __nvvm_atom_inc_gen_ui(Address, Val);
 }
+
 DEVICE uint32_t __kmpc_atomic_max(uint32_t *Address, uint32_t Val) {
-  return atomicMax(Address, Val);
+  return __atomic_fetch_max(Address, Val, __ATOMIC_SEQ_CST);
 }
+
 DEVICE uint32_t __kmpc_atomic_exchange(uint32_t *Address, uint32_t Val) {
-  return atomicExch(Address, Val);
+  uint32_t R;
+  __atomic_exchange(Address, &Val, &R, __ATOMIC_SEQ_CST);
+  return R;
 }
+
 DEVICE uint32_t __kmpc_atomic_cas(uint32_t *Address, uint32_t Compare,
                                   uint32_t Val) {
-  return atomicCAS(Address, Compare, Val);
+  (void)__atomic_compare_exchange(Address, &Compare, &Val, false,
+                                  __ATOMIC_SEQ_CST, __ATOMIC_SEQ_CST);
+  return Compare;
 }
 
 DEVICE unsigned long long __kmpc_atomic_exchange(unsigned long long *Address,
                                                  unsigned long long Val) {
-  return atomicExch(Address, Val);
+  unsigned long long R;
+  __atomic_exchange(Address, &Val, &R, __ATOMIC_SEQ_CST);
+  return R;
 }
+
 DEVICE unsigned long long __kmpc_atomic_add(unsigned long long *Address,
                                             unsigned long long Val) {
-  return atomicAdd(Address, Val);
+  return __atomic_fetch_add(Address, Val, __ATOMIC_SEQ_CST);
 }
 
 #define __OMP_SPIN 1000
 #define UNSET 0u
 #define SET 1u
 
 DEVICE void __kmpc_impl_init_lock(omp_lock_t *lock) {
   __kmpc_impl_unset_lock(lock);