[OpenMP] Replace __ATOMIC_XYZ with atomic::xyz for style

[jdoerfert]

Summary

Also fixes one ordering argument not used.

[jhuber6]

LG

openmp/libomptarget/DeviceRTL/include/Synchronization.h

 /// instruction (hence all threads in the block are "aligned").
 __attribute__((noinline)) void threadsAligned();
 
 #pragma omp end assumes
 ///}
 
 } // namespace synchronize
 
-namespace fence {
-
-/// Memory fence with \p Ordering semantics for the team.
-void team(int Ordering);
-
-/// Memory fence with \p Ordering semantics for the contention group.
-void kernel(int Ordering);
-
-/// Memory fence with \p Ordering semantics for the system.
-void system(int Ordering);
-
-} // namespace fence
-
 namespace atomic {
 
+enum OrderingTy {
+  relaxed = __ATOMIC_RELAXED,
+  aquire = __ATOMIC_ACQUIRE,
+  release = __ATOMIC_RELEASE,
+  acq_rel = __ATOMIC_ACQ_REL,
+  seq_cst = __ATOMIC_SEQ_CST,
+};
+
 /// Atomically load \p Addr with \p Ordering semantics.
-uint32_t load(uint32_t *Addr, int Ordering);
+uint32_t load(uint32_t *Addr, atomic::OrderingTy Ordering);
 
 /// Atomically store \p V to \p Addr with \p Ordering semantics.
-void store(uint32_t *Addr, uint32_t V, int Ordering);
+void store(uint32_t *Addr, uint32_t V, atomic::OrderingTy Ordering);
 
 /// Atomically increment \p *Addr and wrap at \p V with \p Ordering semantics.
-uint32_t inc(uint32_t *Addr, uint32_t V, int Ordering);
+uint32_t inc(uint32_t *Addr, uint32_t V, atomic::OrderingTy Ordering);
 
 /// Atomically add \p V to \p *Addr with \p Ordering semantics.
-uint32_t add(uint32_t *Addr, uint32_t V, int Ordering);
+uint32_t add(uint32_t *Addr, uint32_t V, atomic::OrderingTy Ordering);
 
 /// Atomically add \p V to \p *Addr with \p Ordering semantics.
-uint64_t add(uint64_t *Addr, uint64_t V, int Ordering);
+uint64_t add(uint64_t *Addr, uint64_t V, atomic::OrderingTy Ordering);
 
 } // namespace atomic
 
+namespace fence {
+
+/// Memory fence with \p Ordering semantics for the team.
+void team(atomic::OrderingTy Ordering);
+
+/// Memory fence with \p Ordering semantics for the contention group.
+void kernel(atomic::OrderingTy Ordering);
+
+/// Memory fence with \p Ordering semantics for the system.
+void system(atomic::OrderingTy Ordering);
+
+} // namespace fence
+
 } // namespace _OMP
 
 #endif

openmp/libomptarget/DeviceRTL/src/Reduction.cpp

int32_t __kmpc_nvptx_teams_reduce_nowait_v2(
   static unsigned SHARED(Bound);
   static unsigned SHARED(ChunkTeamCount);
 
   // Block progress for teams greater than the current upper
   // limit. We always only allow a number of teams less or equal
   // to the number of slots in the buffer.
   bool IsMaster = (ThreadId == 0);
   while (IsMaster) {
-    Bound = atomic::load(&IterCnt, __ATOMIC_SEQ_CST);
+    Bound = atomic::load(&IterCnt, atomic::seq_cst);
     if (TeamId < Bound + num_of_records)
       break;
   }
 
   if (IsMaster) {
     int ModBockId = TeamId % num_of_records;
     if (TeamId < num_of_records) {
       lgcpyFct(GlobalBuffer, ModBockId, reduce_data);
     } else
       lgredFct(GlobalBuffer, ModBockId, reduce_data);
 
-    fence::system(__ATOMIC_SEQ_CST);
+    fence::system(atomic::seq_cst);
 
     // Increment team counter.
     // This counter is incremented by all teams in the current
     // BUFFER_SIZE chunk.
-    ChunkTeamCount = atomic::inc(&Cnt, num_of_records - 1u, __ATOMIC_SEQ_CST);
+    ChunkTeamCount = atomic::inc(&Cnt, num_of_records - 1u, atomic::seq_cst);
   }
   // Synchronize
   if (mapping::isSPMDMode())
     __kmpc_barrier(Loc, TId);
 
   // reduce_data is global or shared so before being reduced within the
   // warp we need to bring it in local memory:
   // local_reduce_data = reduce_data[i]
if (IsMaster) {
       IterCnt = 0;
       return 1;
     }
     return 0;
   }
   if (IsMaster && ChunkTeamCount == num_of_records - 1) {
     // Allow SIZE number of teams to proceed writing their
     // intermediate results to the global buffer.
-    atomic::add(&IterCnt, uint32_t(num_of_records), __ATOMIC_SEQ_CST);
+    atomic::add(&IterCnt, uint32_t(num_of_records), atomic::seq_cst);
   }
 
   return 0;
 }
 
 void __kmpc_nvptx_end_reduce(int32_t TId) { FunctionTracingRAII(); }
 
 void __kmpc_nvptx_end_reduce_nowait(int32_t TId) { FunctionTracingRAII(); }
 }
 
 #pragma omp end declare target

openmp/libomptarget/DeviceRTL/src/Synchronization.cpp

 using namespace _OMP;
 
 namespace impl {
 
 /// Atomics
 ///
 ///{
 /// NOTE: This function needs to be implemented by every target.
-uint32_t atomicInc(uint32_t *Address, uint32_t Val, int Ordering);
+uint32_t atomicInc(uint32_t *Address, uint32_t Val,
+                   atomic::OrderingTy Ordering);
 
-uint32_t atomicLoad(uint32_t *Address, int Ordering) {
-  return __atomic_fetch_add(Address, 0U, __ATOMIC_SEQ_CST);
+uint32_t atomicLoad(uint32_t *Address, atomic::OrderingTy Ordering) {
+  return __atomic_fetch_add(Address, 0U, Ordering);
 }
 
-void atomicStore(uint32_t *Address, uint32_t Val, int Ordering) {
+void atomicStore(uint32_t *Address, uint32_t Val, atomic::OrderingTy Ordering) {
   __atomic_store_n(Address, Val, Ordering);
 }
 
-uint32_t atomicAdd(uint32_t *Address, uint32_t Val, int Ordering) {
+uint32_t atomicAdd(uint32_t *Address, uint32_t Val,
+                   atomic::OrderingTy Ordering) {
   return __atomic_fetch_add(Address, Val, Ordering);
 }
-uint32_t atomicMax(uint32_t *Address, uint32_t Val, int Ordering) {
+uint32_t atomicMax(uint32_t *Address, uint32_t Val,
+                   atomic::OrderingTy Ordering) {
   return __atomic_fetch_max(Address, Val, Ordering);
 }
 
-uint32_t atomicExchange(uint32_t *Address, uint32_t Val, int Ordering) {
+uint32_t atomicExchange(uint32_t *Address, uint32_t Val,
+                        atomic::OrderingTy Ordering) {
   uint32_t R;
   __atomic_exchange(Address, &Val, &R, Ordering);
   return R;
 }
 uint32_t atomicCAS(uint32_t *Address, uint32_t Compare, uint32_t Val,
-                   int Ordering) {
+                   atomic::OrderingTy Ordering) {
   (void)__atomic_compare_exchange(Address, &Compare, &Val, false, Ordering,
                                   Ordering);
   return Compare;
 }
 
-uint64_t atomicAdd(uint64_t *Address, uint64_t Val, int Ordering) {
+uint64_t atomicAdd(uint64_t *Address, uint64_t Val,
+                   atomic::OrderingTy Ordering) {
   return __atomic_fetch_add(Address, Val, Ordering);
 }
 ///}
 
 // Forward declarations defined to be defined for AMDGCN and NVPTX.
-uint32_t atomicInc(uint32_t *A, uint32_t V, int Ordering);
+uint32_t atomicInc(uint32_t *A, uint32_t V, atomic::OrderingTy Ordering);
 void namedBarrierInit();
 void namedBarrier();
-void fenceTeam(int Ordering);
-void fenceKernel(int Ordering);
-void fenceSystem(int Ordering);
+void fenceTeam(atomic::OrderingTy Ordering);
+void fenceKernel(atomic::OrderingTy Ordering);
+void fenceSystem(atomic::OrderingTy Ordering);
 void syncWarp(__kmpc_impl_lanemask_t);
 void syncThreads();
 void syncThreadsAligned() { syncThreads(); }
 void unsetLock(omp_lock_t *);
 int testLock(omp_lock_t *);
 void initLock(omp_lock_t *);
 void destroyLock(omp_lock_t *);
 void setLock(omp_lock_t *);
 
 /// AMDGCN Implementation
 ///
 ///{
 #pragma omp begin declare variant match(device = {arch(amdgcn)})
 
-uint32_t atomicInc(uint32_t *A, uint32_t V, int Ordering) {
+uint32_t atomicInc(uint32_t *A, uint32_t V, atomic::OrderingTy Ordering) {
   // builtin_amdgcn_atomic_inc32 should expand to this switch when
   // passed a runtime value, but does not do so yet. Workaround here.
   switch (Ordering) {
   default:
     __builtin_unreachable();
-  case __ATOMIC_RELAXED:
-    return __builtin_amdgcn_atomic_inc32(A, V, __ATOMIC_RELAXED, "");
-  case __ATOMIC_ACQUIRE:
-    return __builtin_amdgcn_atomic_inc32(A, V, __ATOMIC_ACQUIRE, "");
-  case __ATOMIC_RELEASE:
-    return __builtin_amdgcn_atomic_inc32(A, V, __ATOMIC_RELEASE, "");
-  case __ATOMIC_ACQ_REL:
-    return __builtin_amdgcn_atomic_inc32(A, V, __ATOMIC_ACQ_REL, "");
-  case __ATOMIC_SEQ_CST:
-    return __builtin_amdgcn_atomic_inc32(A, V, __ATOMIC_SEQ_CST, "");
+  case atomic::relaxed:
+    return __builtin_amdgcn_atomic_inc32(A, V, atomic::relaxed, "");
+  case atomic::aquire:
+    return __builtin_amdgcn_atomic_inc32(A, V, atomic::aquire, "");
+  case atomic::release:
+    return __builtin_amdgcn_atomic_inc32(A, V, atomic::release, "");
+  case atomic::acq_rel:
+    return __builtin_amdgcn_atomic_inc32(A, V, atomic::acq_rel, "");
+  case atomic::seq_cst:
+    return __builtin_amdgcn_atomic_inc32(A, V, atomic::seq_cst, "");
   }
 }
 
 uint32_t SHARED(namedBarrierTracker);
 
 void namedBarrierInit() {
   // Don't have global ctors, and shared memory is not zero init
-  atomic::store(&namedBarrierTracker, 0u, __ATOMIC_RELEASE);
+  atomic::store(&namedBarrierTracker, 0u, atomic::release);
 }
 
 void namedBarrier() {
   uint32_t NumThreads = omp_get_num_threads();
   // assert(NumThreads % 32 == 0);
 
   uint32_t WarpSize = mapping::getWarpSize();
   uint32_t NumWaves = NumThreads / WarpSize;
 
-  fence::team(__ATOMIC_ACQUIRE);
+  fence::team(atomic::aquire);
 
   // named barrier implementation for amdgcn.
   // Uses two 16 bit unsigned counters. One for the number of waves to have
   // reached the barrier, and one to count how many times the barrier has been
   // passed. These are packed in a single atomically accessed 32 bit integer.
   // Low bits for the number of waves, assumed zero before this call.
   // High bits to count the number of times the barrier has been passed.
 
   // precondition: NumWaves != 0;
   // invariant: NumWaves * WarpSize == NumThreads;
   // precondition: NumWaves < 0xffffu;
 
   // Increment the low 16 bits once, using the lowest active thread.
   if (mapping::isLeaderInWarp()) {
     uint32_t load = atomic::add(&namedBarrierTracker, 1,
-                                __ATOMIC_RELAXED); // commutative
+                                atomic::relaxed); // commutative
 
     // Record the number of times the barrier has been passed
     uint32_t generation = load & 0xffff0000u;
 
     if ((load & 0x0000ffffu) == (NumWaves - 1)) {
       // Reached NumWaves in low bits so this is the last wave.
       // Set low bits to zero and increment high bits
       load += 0x00010000u; // wrap is safe
       load &= 0xffff0000u; // because bits zeroed second
 
       // Reset the wave counter and release the waiting waves
-      atomic::store(&namedBarrierTracker, load, __ATOMIC_RELAXED);
+      atomic::store(&namedBarrierTracker, load, atomic::relaxed);
     } else {
       // more waves still to go, spin until generation counter changes
       do {
         __builtin_amdgcn_s_sleep(0);
-        load = atomic::load(&namedBarrierTracker, __ATOMIC_RELAXED);
+        load = atomic::load(&namedBarrierTracker, atomic::relaxed);
       } while ((load & 0xffff0000u) == generation);
     }
   }
-  fence::team(__ATOMIC_RELEASE);
+  fence::team(atomic::release);
 }
 
 // sema checking of amdgcn_fence is aggressive. Intention is to patch clang
 // so that it is usable within a template environment and so that a runtime
 // value of the memory order is expanded to this switch within clang/llvm.
-void fenceTeam(int Ordering) {
+void fenceTeam(atomic::OrderingTy Ordering) {
   switch (Ordering) {
   default:
     __builtin_unreachable();
-  case __ATOMIC_ACQUIRE:
-    return __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "workgroup");
-  case __ATOMIC_RELEASE:
-    return __builtin_amdgcn_fence(__ATOMIC_RELEASE, "workgroup");
-  case __ATOMIC_ACQ_REL:
-    return __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, "workgroup");
-  case __ATOMIC_SEQ_CST:
-    return __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "workgroup");
+  case atomic::aquire:
+    return __builtin_amdgcn_fence(atomic::aquire, "workgroup");
+  case atomic::release:
+    return __builtin_amdgcn_fence(atomic::release, "workgroup");
+  case atomic::acq_rel:
+    return __builtin_amdgcn_fence(atomic::acq_rel, "workgroup");
+  case atomic::seq_cst:
+    return __builtin_amdgcn_fence(atomic::seq_cst, "workgroup");
   }
 }
-void fenceKernel(int Ordering) {
+void fenceKernel(atomic::OrderingTy Ordering) {
   switch (Ordering) {
   default:
     __builtin_unreachable();
-  case __ATOMIC_ACQUIRE:
-    return __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "agent");
-  case __ATOMIC_RELEASE:
-    return __builtin_amdgcn_fence(__ATOMIC_RELEASE, "agent");
-  case __ATOMIC_ACQ_REL:
-    return __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, "agent");
-  case __ATOMIC_SEQ_CST:
-    return __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "agent");
+  case atomic::aquire:
+    return __builtin_amdgcn_fence(atomic::aquire, "agent");
+  case atomic::release:
+    return __builtin_amdgcn_fence(atomic::release, "agent");
+  case atomic::acq_rel:
+    return __builtin_amdgcn_fence(atomic::acq_rel, "agent");
+  case atomic::seq_cst:
+    return __builtin_amdgcn_fence(atomic::seq_cst, "agent");
   }
 }
-void fenceSystem(int Ordering) {
+void fenceSystem(atomic::OrderingTy Ordering) {
   switch (Ordering) {
   default:
     __builtin_unreachable();
-  case __ATOMIC_ACQUIRE:
-    return __builtin_amdgcn_fence(__ATOMIC_ACQUIRE, "");
-  case __ATOMIC_RELEASE:
-    return __builtin_amdgcn_fence(__ATOMIC_RELEASE, "");
-  case __ATOMIC_ACQ_REL:
-    return __builtin_amdgcn_fence(__ATOMIC_ACQ_REL, "");
-  case __ATOMIC_SEQ_CST:
-    return __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, "");
+  case atomic::aquire:
+    return __builtin_amdgcn_fence(atomic::aquire, "");
+  case atomic::release:
+    return __builtin_amdgcn_fence(atomic::release, "");
+  case atomic::acq_rel:
+    return __builtin_amdgcn_fence(atomic::acq_rel, "");
+  case atomic::seq_cst:
+    return __builtin_amdgcn_fence(atomic::seq_cst, "");
   }
 }
 
 void syncWarp(__kmpc_impl_lanemask_t) {
   // AMDGCN doesn't need to sync threads in a warp
 }
 
 void syncThreads() { __builtin_amdgcn_s_barrier(); }

 ///}
 
 /// NVPTX Implementation
 ///
 ///{
 #pragma omp begin declare variant match(                                       \
     device = {arch(nvptx, nvptx64)}, implementation = {extension(match_any)})
 
-uint32_t atomicInc(uint32_t *Address, uint32_t Val, int Ordering) {
+uint32_t atomicInc(uint32_t *Address, uint32_t Val,
+                   atomic::OrderingTy Ordering) {
   return __nvvm_atom_inc_gen_ui(Address, Val);
 }
 
 void namedBarrierInit() {}
 
 void namedBarrier() {
   uint32_t NumThreads = omp_get_num_threads();
   ASSERT(NumThreads % 32 == 0);

 constexpr uint32_t OMP_SPIN = 1000;
 constexpr uint32_t UNSET = 0;
 constexpr uint32_t SET = 1;
 
 // TODO: This seems to hide a bug in the declare variant handling. If it is
 // called before it is defined
 //       here the overload won't happen. Investigate lalter!
 void unsetLock(omp_lock_t *Lock) {
-  (void)atomicExchange((uint32_t *)Lock, UNSET, __ATOMIC_SEQ_CST);
+  (void)atomicExchange((uint32_t *)Lock, UNSET, atomic::seq_cst);
 }
 
 int testLock(omp_lock_t *Lock) {
-  return atomicAdd((uint32_t *)Lock, 0u, __ATOMIC_SEQ_CST);
+  return atomicAdd((uint32_t *)Lock, 0u, atomic::seq_cst);
 }
 
 void initLock(omp_lock_t *Lock) { unsetLock(Lock); }
 
 void destroyLock(omp_lock_t *Lock) { unsetLock(Lock); }
 
 void setLock(omp_lock_t *Lock) {
   // TODO: not sure spinning is a good idea here..
-  while (atomicCAS((uint32_t *)Lock, UNSET, SET, __ATOMIC_SEQ_CST) != UNSET) {
+  while (atomicCAS((uint32_t *)Lock, UNSET, SET, atomic::seq_cst) != UNSET) {
     int32_t start = __nvvm_read_ptx_sreg_clock();
     int32_t now;
     for (;;) {
       now = __nvvm_read_ptx_sreg_clock();
       int32_t cycles = now > start ? now - start : now + (0xffffffff - start);
       if (cycles >= OMP_SPIN * mapping::getBlockId()) {
         break;
       }

 }
 
 void synchronize::warp(LaneMaskTy Mask) { impl::syncWarp(Mask); }
 
 void synchronize::threads() { impl::syncThreads(); }
 
 void synchronize::threadsAligned() { impl::syncThreadsAligned(); }
 
-void fence::team(int Ordering) { impl::fenceTeam(Ordering); }
+void fence::team(atomic::OrderingTy Ordering) { impl::fenceTeam(Ordering); }
 
-void fence::kernel(int Ordering) { impl::fenceKernel(Ordering); }
+void fence::kernel(atomic::OrderingTy Ordering) { impl::fenceKernel(Ordering); }
 
-void fence::system(int Ordering) { impl::fenceSystem(Ordering); }
+void fence::system(atomic::OrderingTy Ordering) { impl::fenceSystem(Ordering); }
 
-uint32_t atomic::load(uint32_t *Addr, int Ordering) {
+uint32_t atomic::load(uint32_t *Addr, atomic::OrderingTy Ordering) {
   return impl::atomicLoad(Addr, Ordering);
 }
 
-void atomic::store(uint32_t *Addr, uint32_t V, int Ordering) {
+void atomic::store(uint32_t *Addr, uint32_t V, atomic::OrderingTy Ordering) {
   impl::atomicStore(Addr, V, Ordering);
 }
 
-uint32_t atomic::inc(uint32_t *Addr, uint32_t V, int Ordering) {
+uint32_t atomic::inc(uint32_t *Addr, uint32_t V, atomic::OrderingTy Ordering) {
   return impl::atomicInc(Addr, V, Ordering);
 }
 
-uint32_t atomic::add(uint32_t *Addr, uint32_t V, int Ordering) {
+uint32_t atomic::add(uint32_t *Addr, uint32_t V, atomic::OrderingTy Ordering) {
   return impl::atomicAdd(Addr, V, Ordering);
 }
 
-uint64_t atomic::add(uint64_t *Addr, uint64_t V, int Ordering) {
+uint64_t atomic::add(uint64_t *Addr, uint64_t V, atomic::OrderingTy Ordering) {
   return impl::atomicAdd(Addr, V, Ordering);
 }
 
 extern "C" {
 void __kmpc_ordered(IdentTy *Loc, int32_t TId) { FunctionTracingRAII(); }
 
 void __kmpc_end_ordered(IdentTy *Loc, int32_t TId) { FunctionTracingRAII(); }
 

 
 void __kmpc_end_single(IdentTy *Loc, int32_t TId) {
   FunctionTracingRAII();
   // The barrier is explicitly called.
 }
 
 void __kmpc_flush(IdentTy *Loc) {
   FunctionTracingRAII();
-  fence::kernel(__ATOMIC_SEQ_CST);
+  fence::kernel(atomic::seq_cst);
 }
 
 uint64_t __kmpc_warp_active_thread_mask(void) {
   FunctionTracingRAII();
   return mapping::activemask();
 }
 
 void __kmpc_syncwarp(uint64_t Mask) {

openmp/libomptarget/DeviceRTL/src/Workshare.cpp

if (schedule == kmp_sched_static_chunk) {
       if (chunk < 1)
         chunk = 1;
       DST->Chunk = chunk;
       DST->LoopUpperBound = ub;
       DST->NextLowerBound = lb;
       __kmpc_barrier(loc, threadId);
       if (tid == 0) {
         Cnt = 0;
-        fence::team(__ATOMIC_SEQ_CST);
+        fence::team(atomic::seq_cst);
       }
       __kmpc_barrier(loc, threadId);
     }
   }
 
   ////////////////////////////////////////////////////////////////////////////////
   // Support for dispatch next
 
   static uint64_t NextIter() {
     __kmpc_impl_lanemask_t active = mapping::activemask();
     uint32_t leader = utils::ffs(active) - 1;
     uint32_t change = utils::popc(active);
     __kmpc_impl_lanemask_t lane_mask_lt = mapping::lanemaskLT();
     unsigned int rank = utils::popc(active & lane_mask_lt);
     uint64_t warp_res = 0;
     if (rank == 0) {
-      warp_res = atomic::add(&Cnt, change, __ATOMIC_SEQ_CST);
+      warp_res = atomic::add(&Cnt, change, atomic::seq_cst);
     }
     warp_res = utils::shuffle(active, warp_res, leader);
     return warp_res + rank;
   }
 
   static int DynamicNextChunk(T &lb, T &ub, T chunkSize, T loopLowerBound,
                               T loopUpperBound) {
     T N = NextIter();