barrier synchronization

**Barrier Synchronization** is the **fundamental parallel synchronization primitive where all participating threads or processes must arrive at a designated program point before any are allowed to proceed past it — ensuring that all work from the previous phase is complete before the next phase begins, which is essential for phased parallel algorithms but creates a performance bottleneck proportional to the slowest thread's arrival time**. **Why Barriers Are Necessary** In phased parallel computations (iterative solvers, stencil codes, BSP algorithms), each phase depends on results from the previous phase. Without a barrier between phases, fast threads in phase K+1 would read stale data from slow threads still in phase K, producing incorrect results. The barrier guarantees consistency at the cost of forcing all threads to wait for the slowest. **Implementation Approaches** - **Centralized Counter Barrier**: An atomic counter initialized to N (thread count). Each arriving thread decrements it. The last thread (counter → 0) signals all others to proceed. Simple but creates contention on the counter — O(N) serialized atomic operations on the same cache line. - **Tree Barrier**: Threads are organized in a binary tree. Each pair synchronizes locally (leaf level), then representatives synchronize at the next level, up to the root. The root signals completion back down the tree. Total steps: O(log N). Reduces contention by distributing synchronization across the tree. - **Butterfly Barrier**: In round k, each thread i synchronizes with thread i XOR 2^k. After log2(N) rounds, all threads have transitively synchronized. O(log N) steps with good locality properties for hardware with neighbor communication. - **Sense-Reversing Barrier**: Uses a shared "sense" flag that alternates between true and false at each barrier. Threads spin on their local sense copy, which is updated when the barrier completes. Avoids the "early arrival" race where a thread from barrier K+1 arrives before barrier K has fully released. **GPU Barriers** - **Block Barrier (`__syncthreads()`)**: Synchronizes all threads within a thread block. Implemented in hardware — ~20 cycles. Required after shared memory writes that other threads will read. - **Grid Barrier (Cooperative Groups)**: Synchronizes all thread blocks in a grid. Requires cooperative launch and is limited to grids that fit simultaneously on the GPU (one block per SM maximum). Used for persistent kernels. - **No Inter-Block Sync**: CUDA deliberately provides no inter-block barrier in the normal programming model because blocks may not execute concurrently. Algorithms requiring global sync must use kernel boundaries. **Performance Impact** The cost of a barrier has two components: the synchronization mechanism overhead (~100 ns for a good tree barrier on multi-core CPU) and the load imbalance cost (time the fastest thread waits for the slowest). The imbalance cost often dominates by 10-100x — making load balancing far more important than barrier algorithm optimization. Barrier Synchronization is **the metronome of phased parallel computing** — enforcing lockstep progress that guarantees correctness but imposes a speed limit equal to the slowest participant in each phase.