Cold-Air-Networks/mirror-llvm-project
2
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions 7

[OpenMP][NFC] Use uinc atomic builtins for this operation (#177207)

Browse source 
Summary:
We support this now, this is 1-to-1 equivalent and simply prevents us
from needing to do it ourselves.
This commit is contained in:
Joseph Huber 2026-01-21 11:31:23 -06:00 committed by GitHub
parent ac5e1b7a96
commit f534711729
No known key found for this signature in database
GPG key ID: B5690EEEBB952194
2 changed files with 7 additions and 48 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

11
openmp/device/include/Synchronization.h
View file

@ -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,

44
openmp/device/src/Synchronization.cpp
View file

@ -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); }