[Openmp-commits] [openmp] [OpenMP][NFC] Use `uinc` atomic builtins for this operation (PR #177207)

[Joseph Huber via Openmp-commits]

https://github.com/jhuber6 updated https://github.com/llvm/llvm-project/pull/177207

>From 6f64249ed50c092c83baabcd335fa3b76bd82b83 Mon Sep 17 00:00:00 2001
From: Joseph Huber <huberjn at outlook.com>
Date: Wed, 21 Jan 2026 11:16:52 -0600
Subject: [PATCH] [OpenMP][NFC] Use `uinc` atomic builtins for this operation

Summary:
We support this now, this is 1-to-1 equivalent and simply prevents us
from needing to do it ourselves.
---
 openmp/device/include/Synchronization.h | 11 ++++---
 openmp/device/src/Synchronization.cpp   | 44 -------------------------
 2 files changed, 7 insertions(+), 48 deletions(-)

diff --git a/openmp/device/include/Synchronization.h b/openmp/device/include/Synchronization.h
index 7e7c8eacb9173..f3b47844fa454 100644
--- a/openmp/device/include/Synchronization.h
+++ b/openmp/device/include/Synchronization.h
@@ -34,14 +34,17 @@ enum MemScopeTy {
   single = __MEMORY_SCOPE_SINGLE,
 };
 
-/// Atomically increment \p *Addr and wrap at \p V with \p Ordering semantics.
-uint32_t inc(uint32_t *Addr, uint32_t V, OrderingTy Ordering,
-             MemScopeTy MemScope = MemScopeTy::device);
-
 /// Atomically perform <op> on \p V and \p *Addr with \p Ordering semantics. The
 /// result is stored in \p *Addr;
 /// {
 
+/// Atomically increments with wrapping semantics modulo \p Val.
+template <typename Ty, typename V = utils::remove_addrspace_t<Ty>>
+V inc(Ty *Address, V Val, atomic::OrderingTy Ordering,
+      MemScopeTy MemScope = MemScopeTy::device) {
+  return __scoped_atomic_uinc_wrap(Address, Val, Ordering, MemScope);
+}
+
 template <typename Ty, typename V = utils::remove_addrspace_t<Ty>>
 bool cas(Ty *Address, V ExpectedV, V DesiredV, atomic::OrderingTy OrderingSucc,
          atomic::OrderingTy OrderingFail,
diff --git a/openmp/device/src/Synchronization.cpp b/openmp/device/src/Synchronization.cpp
index 501dc4a291ed1..f574211123aac 100644
--- a/openmp/device/src/Synchronization.cpp
+++ b/openmp/device/src/Synchronization.cpp
@@ -33,40 +33,6 @@ namespace impl {
 ///{
 #ifdef __AMDGPU__
 
-uint32_t atomicInc(uint32_t *A, uint32_t V, atomic::OrderingTy Ordering,
-                   atomic::MemScopeTy MemScope) {
-  // builtin_amdgcn_atomic_inc32 should expand to this switch when
-  // passed a runtime value, but does not do so yet. Workaround here.
-
-#define ScopeSwitch(ORDER)                                                     \
-  switch (MemScope) {                                                          \
-  case atomic::MemScopeTy::system:                                             \
-    return __builtin_amdgcn_atomic_inc32(A, V, ORDER, "");                     \
-  case atomic::MemScopeTy::device:                                             \
-    return __builtin_amdgcn_atomic_inc32(A, V, ORDER, "agent");                \
-  case atomic::MemScopeTy::workgroup:                                          \
-    return __builtin_amdgcn_atomic_inc32(A, V, ORDER, "workgroup");            \
-  case atomic::MemScopeTy::wavefront:                                          \
-    return __builtin_amdgcn_atomic_inc32(A, V, ORDER, "wavefront");            \
-  case atomic::MemScopeTy::single:                                             \
-    return __builtin_amdgcn_atomic_inc32(A, V, ORDER, "singlethread");         \
-  }
-
-#define Case(ORDER)                                                            \
-  case ORDER:                                                                  \
-    ScopeSwitch(ORDER)
-
-  switch (Ordering) {
-    Case(atomic::relaxed);
-    Case(atomic::acquire);
-    Case(atomic::release);
-    Case(atomic::acq_rel);
-    Case(atomic::seq_cst);
-#undef Case
-#undef ScopeSwitch
-  }
-}
-
 [[clang::loader_uninitialized]] Local<uint32_t> namedBarrierTracker;
 
 void namedBarrierInit() {
@@ -184,11 +150,6 @@ void setCriticalLock(omp_lock_t *Lock) {
 ///{
 #ifdef __NVPTX__
 
-uint32_t atomicInc(uint32_t *Address, uint32_t Val, atomic::OrderingTy Ordering,
-                   atomic::MemScopeTy MemScope) {
-  return __nvvm_atom_inc_gen_ui(Address, Val);
-}
-
 void namedBarrierInit() {}
 
 void namedBarrier() {
@@ -281,11 +242,6 @@ void fence::kernel(atomic::OrderingTy Ordering) { impl::fenceKernel(Ordering); }
 
 void fence::system(atomic::OrderingTy Ordering) { impl::fenceSystem(Ordering); }
 
-uint32_t atomic::inc(uint32_t *Addr, uint32_t V, atomic::OrderingTy Ordering,
-                     atomic::MemScopeTy MemScope) {
-  return impl::atomicInc(Addr, V, Ordering, MemScope);
-}
-
 void unsetCriticalLock(omp_lock_t *Lock) { impl::unsetLock(Lock); }
 
 void setCriticalLock(omp_lock_t *Lock) { impl::setLock(Lock); }