378 lines · cpp
1//===- Synchronization.cpp - OpenMP Device synchronization API ---- c++ -*-===//2//3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.4// See https://llvm.org/LICENSE.txt for license information.5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception6//7//===----------------------------------------------------------------------===//8//9// Include all synchronization.10//11//===----------------------------------------------------------------------===//12 13#include "Synchronization.h"14 15#include "Debug.h"16#include "DeviceTypes.h"17#include "DeviceUtils.h"18#include "Interface.h"19#include "Mapping.h"20#include "State.h"21 22using namespace ompx;23 24namespace impl {25 26/// Atomics27///28///{29///}30 31/// AMDGCN Implementation32///33///{34#ifdef __AMDGPU__35 36uint32_t atomicInc(uint32_t *A, uint32_t V, atomic::OrderingTy Ordering,37 atomic::MemScopeTy MemScope) {38 // builtin_amdgcn_atomic_inc32 should expand to this switch when39 // passed a runtime value, but does not do so yet. Workaround here.40 41#define ScopeSwitch(ORDER) \42 switch (MemScope) { \43 case atomic::MemScopeTy::system: \44 return __builtin_amdgcn_atomic_inc32(A, V, ORDER, ""); \45 case atomic::MemScopeTy::device: \46 return __builtin_amdgcn_atomic_inc32(A, V, ORDER, "agent"); \47 case atomic::MemScopeTy::workgroup: \48 return __builtin_amdgcn_atomic_inc32(A, V, ORDER, "workgroup"); \49 case atomic::MemScopeTy::wavefront: \50 return __builtin_amdgcn_atomic_inc32(A, V, ORDER, "wavefront"); \51 case atomic::MemScopeTy::single: \52 return __builtin_amdgcn_atomic_inc32(A, V, ORDER, "singlethread"); \53 }54 55#define Case(ORDER) \56 case ORDER: \57 ScopeSwitch(ORDER)58 59 switch (Ordering) {60 Case(atomic::relaxed);61 Case(atomic::acquire);62 Case(atomic::release);63 Case(atomic::acq_rel);64 Case(atomic::seq_cst);65#undef Case66#undef ScopeSwitch67 }68}69 70[[clang::loader_uninitialized]] Local<uint32_t> namedBarrierTracker;71 72void namedBarrierInit() {73 // Don't have global ctors, and shared memory is not zero init74 atomic::store(&namedBarrierTracker, 0u, atomic::release);75}76 77void namedBarrier() {78 uint32_t NumThreads = omp_get_num_threads();79 // assert(NumThreads % 32 == 0);80 81 uint32_t WarpSize = mapping::getWarpSize();82 uint32_t NumWaves = NumThreads / WarpSize;83 84 fence::team(atomic::acquire);85 86 // named barrier implementation for amdgcn.87 // Uses two 16 bit unsigned counters. One for the number of waves to have88 // reached the barrier, and one to count how many times the barrier has been89 // passed. These are packed in a single atomically accessed 32 bit integer.90 // Low bits for the number of waves, assumed zero before this call.91 // High bits to count the number of times the barrier has been passed.92 93 // precondition: NumWaves != 0;94 // invariant: NumWaves * WarpSize == NumThreads;95 // precondition: NumWaves < 0xffffu;96 97 // Increment the low 16 bits once, using the lowest active thread.98 if (mapping::isLeaderInWarp()) {99 uint32_t load = atomic::add(&namedBarrierTracker, 1,100 atomic::relaxed); // commutative101 102 // Record the number of times the barrier has been passed103 uint32_t generation = load & 0xffff0000u;104 105 if ((load & 0x0000ffffu) == (NumWaves - 1)) {106 // Reached NumWaves in low bits so this is the last wave.107 // Set low bits to zero and increment high bits108 load += 0x00010000u; // wrap is safe109 load &= 0xffff0000u; // because bits zeroed second110 111 // Reset the wave counter and release the waiting waves112 atomic::store(&namedBarrierTracker, load, atomic::relaxed);113 } else {114 // more waves still to go, spin until generation counter changes115 do {116 __builtin_amdgcn_s_sleep(0);117 load = atomic::load(&namedBarrierTracker, atomic::relaxed);118 } while ((load & 0xffff0000u) == generation);119 }120 }121 fence::team(atomic::release);122}123 124void fenceTeam(atomic::OrderingTy Ordering) {125 return __scoped_atomic_thread_fence(Ordering, atomic::workgroup);126}127 128void fenceKernel(atomic::OrderingTy Ordering) {129 return __scoped_atomic_thread_fence(Ordering, atomic::device);130}131 132void fenceSystem(atomic::OrderingTy Ordering) {133 return __scoped_atomic_thread_fence(Ordering, atomic::system);134}135 136void syncWarp(__kmpc_impl_lanemask_t) {137 // This is a no-op on current AMDGPU hardware but it is used by the optimizer138 // to enforce convergent behaviour between control flow graphs.139 __builtin_amdgcn_wave_barrier();140}141 142void syncThreads(atomic::OrderingTy Ordering) {143 if (Ordering != atomic::relaxed)144 fenceTeam(Ordering == atomic::acq_rel ? atomic::release : atomic::seq_cst);145 146 __builtin_amdgcn_s_barrier();147 148 if (Ordering != atomic::relaxed)149 fenceTeam(Ordering == atomic::acq_rel ? atomic::acquire : atomic::seq_cst);150}151void syncThreadsAligned(atomic::OrderingTy Ordering) { syncThreads(Ordering); }152 153// TODO: Don't have wavefront lane locks. Possibly can't have them.154void unsetLock(omp_lock_t *) { __builtin_trap(); }155int testLock(omp_lock_t *) { __builtin_trap(); }156void initLock(omp_lock_t *) { __builtin_trap(); }157void destroyLock(omp_lock_t *) { __builtin_trap(); }158void setLock(omp_lock_t *) { __builtin_trap(); }159 160constexpr uint32_t UNSET = 0;161constexpr uint32_t SET = 1;162 163void unsetCriticalLock(omp_lock_t *Lock) {164 (void)atomicExchange((uint32_t *)Lock, UNSET, atomic::acq_rel);165}166 167void setCriticalLock(omp_lock_t *Lock) {168 uint64_t LowestActiveThread = utils::ffs(mapping::activemask()) - 1;169 if (mapping::getThreadIdInWarp() == LowestActiveThread) {170 fenceKernel(atomic::release);171 while (172 !cas((uint32_t *)Lock, UNSET, SET, atomic::relaxed, atomic::relaxed)) {173 __builtin_amdgcn_s_sleep(32);174 }175 fenceKernel(atomic::acquire);176 }177}178 179#endif180///}181 182/// NVPTX Implementation183///184///{185#ifdef __NVPTX__186 187uint32_t atomicInc(uint32_t *Address, uint32_t Val, atomic::OrderingTy Ordering,188 atomic::MemScopeTy MemScope) {189 return __nvvm_atom_inc_gen_ui(Address, Val);190}191 192void namedBarrierInit() {}193 194void namedBarrier() {195 uint32_t NumThreads = omp_get_num_threads();196 ASSERT(NumThreads % 32 == 0, nullptr);197 198 // The named barrier for active parallel threads of a team in an L1 parallel199 // region to synchronize with each other.200 constexpr int BarrierNo = 7;201 __nvvm_barrier_sync_cnt(BarrierNo, NumThreads);202}203 204void fenceTeam(atomic::OrderingTy) { __nvvm_membar_cta(); }205 206void fenceKernel(atomic::OrderingTy) { __nvvm_membar_gl(); }207 208void fenceSystem(atomic::OrderingTy) { __nvvm_membar_sys(); }209 210void syncWarp(__kmpc_impl_lanemask_t Mask) { __nvvm_bar_warp_sync(Mask); }211 212void syncThreads(atomic::OrderingTy Ordering) {213 constexpr int BarrierNo = 8;214 __nvvm_barrier_sync(BarrierNo);215}216 217void syncThreadsAligned(atomic::OrderingTy Ordering) { __syncthreads(); }218 219constexpr uint32_t OMP_SPIN = 1000;220constexpr uint32_t UNSET = 0;221constexpr uint32_t SET = 1;222 223// TODO: This seems to hide a bug in the declare variant handling. If it is224// called before it is defined225// here the overload won't happen. Investigate lalter!226void unsetLock(omp_lock_t *Lock) {227 (void)atomicExchange((uint32_t *)Lock, UNSET, atomic::seq_cst);228}229 230int testLock(omp_lock_t *Lock) {231 return atomic::add((uint32_t *)Lock, 0u, atomic::seq_cst);232}233 234void initLock(omp_lock_t *Lock) { unsetLock(Lock); }235 236void destroyLock(omp_lock_t *Lock) { unsetLock(Lock); }237 238void setLock(omp_lock_t *Lock) {239 // TODO: not sure spinning is a good idea here..240 while (atomic::cas((uint32_t *)Lock, UNSET, SET, atomic::seq_cst,241 atomic::seq_cst) != UNSET) {242 int32_t start = __nvvm_read_ptx_sreg_clock();243 int32_t now;244 for (;;) {245 now = __nvvm_read_ptx_sreg_clock();246 int32_t cycles = now > start ? now - start : now + (0xffffffff - start);247 if (cycles >= OMP_SPIN * mapping::getBlockIdInKernel()) {248 break;249 }250 }251 } // wait for 0 to be the read value252}253 254void unsetCriticalLock(omp_lock_t *Lock) { unsetLock(Lock); }255 256void setCriticalLock(omp_lock_t *Lock) { setLock(Lock); }257 258#endif259///}260 261} // namespace impl262 263void synchronize::init(bool IsSPMD) {264 if (!IsSPMD)265 impl::namedBarrierInit();266}267 268void synchronize::warp(LaneMaskTy Mask) { impl::syncWarp(Mask); }269 270void synchronize::threads(atomic::OrderingTy Ordering) {271 impl::syncThreads(Ordering);272}273 274void synchronize::threadsAligned(atomic::OrderingTy Ordering) {275 impl::syncThreadsAligned(Ordering);276}277 278void fence::team(atomic::OrderingTy Ordering) { impl::fenceTeam(Ordering); }279 280void fence::kernel(atomic::OrderingTy Ordering) { impl::fenceKernel(Ordering); }281 282void fence::system(atomic::OrderingTy Ordering) { impl::fenceSystem(Ordering); }283 284uint32_t atomic::inc(uint32_t *Addr, uint32_t V, atomic::OrderingTy Ordering,285 atomic::MemScopeTy MemScope) {286 return impl::atomicInc(Addr, V, Ordering, MemScope);287}288 289void unsetCriticalLock(omp_lock_t *Lock) { impl::unsetLock(Lock); }290 291void setCriticalLock(omp_lock_t *Lock) { impl::setLock(Lock); }292 293extern "C" {294void __kmpc_ordered(IdentTy *Loc, int32_t TId) {}295 296void __kmpc_end_ordered(IdentTy *Loc, int32_t TId) {}297 298int32_t __kmpc_cancel_barrier(IdentTy *Loc, int32_t TId) {299 __kmpc_barrier(Loc, TId);300 return 0;301}302 303void __kmpc_barrier(IdentTy *Loc, int32_t TId) {304 if (mapping::isSPMDMode())305 return __kmpc_barrier_simple_spmd(Loc, TId);306 307 // Generic parallel regions are run with multiple of the warp size or single308 // threaded, in the latter case we need to stop here.309 if (omp_get_num_threads() == 1)310 return __kmpc_flush(Loc);311 312 impl::namedBarrier();313}314 315[[clang::noinline]] void __kmpc_barrier_simple_spmd(IdentTy *Loc, int32_t TId) {316 synchronize::threadsAligned(atomic::OrderingTy::seq_cst);317}318 319[[clang::noinline]] void __kmpc_barrier_simple_generic(IdentTy *Loc,320 int32_t TId) {321 synchronize::threads(atomic::OrderingTy::seq_cst);322}323 324int32_t __kmpc_master(IdentTy *Loc, int32_t TId) {325 return omp_get_thread_num() == 0;326}327 328void __kmpc_end_master(IdentTy *Loc, int32_t TId) {}329 330int32_t __kmpc_masked(IdentTy *Loc, int32_t TId, int32_t Filter) {331 return omp_get_thread_num() == Filter;332}333 334void __kmpc_end_masked(IdentTy *Loc, int32_t TId) {}335 336int32_t __kmpc_single(IdentTy *Loc, int32_t TId) {337 return __kmpc_master(Loc, TId);338}339 340void __kmpc_end_single(IdentTy *Loc, int32_t TId) {341 // The barrier is explicitly called.342}343 344void __kmpc_flush(IdentTy *Loc) { fence::kernel(atomic::seq_cst); }345 346uint64_t __kmpc_warp_active_thread_mask(void) { return mapping::activemask(); }347 348void __kmpc_syncwarp(uint64_t Mask) { synchronize::warp(Mask); }349 350void __kmpc_critical(IdentTy *Loc, int32_t TId, CriticalNameTy *Name) {351 impl::setCriticalLock(reinterpret_cast<omp_lock_t *>(Name));352}353 354void __kmpc_end_critical(IdentTy *Loc, int32_t TId, CriticalNameTy *Name) {355 impl::unsetCriticalLock(reinterpret_cast<omp_lock_t *>(Name));356}357 358void omp_init_lock(omp_lock_t *Lock) { impl::initLock(Lock); }359 360void omp_destroy_lock(omp_lock_t *Lock) { impl::destroyLock(Lock); }361 362void omp_set_lock(omp_lock_t *Lock) { impl::setLock(Lock); }363 364void omp_unset_lock(omp_lock_t *Lock) { impl::unsetLock(Lock); }365 366int omp_test_lock(omp_lock_t *Lock) { return impl::testLock(Lock); }367 368void ompx_sync_block(int Ordering) {369 impl::syncThreadsAligned(atomic::OrderingTy(Ordering));370}371void ompx_sync_block_acq_rel() {372 impl::syncThreadsAligned(atomic::OrderingTy::acq_rel);373}374void ompx_sync_block_divergent(int Ordering) {375 impl::syncThreads(atomic::OrderingTy(Ordering));376}377} // extern "C"378