warp level primitives cuda

**Warp-Level Primitives** are **the specialized CUDA intrinsics that enable efficient communication and synchronization among the 32 threads within a warp — leveraging the SIMT execution model where warp threads execute in lockstep to perform shuffle operations, collective votes, and reductions without shared memory or atomics, achieving single-cycle data exchange and enabling high-performance algorithms like warp-level reductions and parallel scans**. **Warp Shuffle Operations:** - **__shfl_sync(mask, var, srcLane)**: thread receives the value of var from thread srcLane within the warp; mask specifies which threads participate (0xffffffff for all 32); single instruction, zero latency data exchange — no shared memory required; enables efficient broadcast, rotation, and butterfly exchange patterns - **__shfl_up_sync(mask, var, delta)**: thread i receives var from thread i-delta; threads 0 to delta-1 receive their own value; used for prefix sum (scan) operations; delta=1,2,4,8,16 sequence implements log₂(32) parallel scan across the warp - **__shfl_down_sync(mask, var, delta)**: thread i receives var from thread i+delta; threads 32-delta to 31 receive their own value; used for suffix operations and reverse scans; complementary to shfl_up - **__shfl_xor_sync(mask, var, laneMask)**: thread i receives var from thread i^laneMask; implements butterfly exchange patterns; laneMask=1,2,4,8,16 sequence performs parallel reduction or broadcast in log₂(32) steps; critical for FFT and bitonic sort algorithms **Warp Vote Functions:** - **__all_sync(mask, predicate)**: returns true if predicate is true for all threads in mask; single instruction evaluates collective condition; used for early exit (if all threads finished, exit loop) and validation (assert all threads agree) - **__any_sync(mask, predicate)**: returns true if predicate is true for any thread in mask; detects if any thread needs special handling; enables divergence detection and conditional execution optimization - **__ballot_sync(mask, predicate)**: returns 32-bit integer where bit i is set if thread i's predicate is true; provides complete information about which threads satisfy condition; enables compact encoding of thread states and efficient work distribution - **__activemask()**: returns mask of currently active threads in the warp; threads that have exited or are in divergent branches are inactive; critical for correct synchronization in divergent code paths **Warp-Level Reductions:** - **Sum Reduction**: sum = val; for (int offset = 16; offset > 0; offset /= 2) sum += __shfl_down_sync(0xffffffff, sum, offset); — reduces 32 values in 5 shuffle operations (log₂(32)); 10-20× faster than shared memory reduction for small reductions - **Max/Min Reduction**: identical pattern using max/min instead of addition; single warp reduces 32 elements to maximum in 5 instructions; critical for finding global extrema in parallel algorithms - **Warp Aggregated Atomics**: reduce 32 values within warp using shuffle, then single thread performs atomic to global memory; reduces atomic contention by 32× compared to per-thread atomics; essential for high-performance histograms and scatter operations - **Segmented Reduction**: use __ballot_sync to identify segment boundaries; perform reduction within each segment using masked shuffle operations; enables variable-length reductions within a single warp **Cooperative Groups Warp Interface:** - **thread_block_tile<32> warp = tiled_partition<32>(this_thread_block())**: creates explicit warp object; provides .shfl(), .any(), .all() member functions with cleaner syntax than intrinsics - **Subwarp Tiles**: tiled_partition<16> or tiled_partition<8> creates sub-warp groups; enables fine-grained parallelism for small problems; each tile operates independently with its own shuffle and vote operations - **warp.sync()**: explicit warp synchronization; required on Volta+ where independent thread scheduling allows warp threads to diverge; replaces implicit warp-synchronous assumptions from pre-Volta architectures - **warp.match_any(value)**: returns mask of threads with the same value; enables efficient grouping and work distribution based on data values; used in hash table lookups and dynamic parallelism **Performance Characteristics:** - **Latency**: shuffle operations complete in 1-2 cycles; vote operations complete in 1 cycle; 10-100× faster than shared memory access (20-30 cycles) for small data exchanges - **Bandwidth**: warp shuffles provide 32 × 4 bytes × GPU_clock bandwidth per SM; at 1.4 GHz, this is ~180 GB/s per SM — comparable to shared memory bandwidth but without occupying shared memory capacity - **Occupancy Independence**: shuffle operations don't consume shared memory; enables high occupancy even with complex per-thread state; critical for latency-hiding in memory-bound kernels - **Register Pressure**: shuffle operates on registers; excessive register usage limits occupancy; balance between using shuffles (register-to-register) vs shared memory (register-to-memory-to-register) based on register availability **Common Patterns:** - **Warp-Level Matrix Multiply**: each warp computes a small tile (32×32 or 16×16) using shuffle to broadcast matrix elements; eliminates shared memory for small GEMM operations; used in Tensor Core warp-level matrix fragments - **Parallel Scan (Prefix Sum)**: Hillis-Steele scan using shuffle_up in log₂(32) iterations; each iteration doubles the stride; produces inclusive scan of 32 elements in 5 steps; building block for larger scans - **Compact/Stream Compaction**: use __ballot_sync to identify valid elements; __popc (population count) on ballot result gives compaction offset; shuffle valid elements to compact positions; single-warp compaction without shared memory - **Warp-Aggregated Loads**: threads cooperatively load data using shuffle to distribute addresses; reduces load instructions and improves cache utilization; particularly effective for irregular access patterns Warp-level primitives are **the low-level building blocks that enable the highest-performance GPU algorithms — by exploiting the SIMT execution model to perform single-cycle data exchange and collective operations, expert CUDA programmers achieve 2-10× speedups over shared memory implementations for fine-grained parallel patterns, making warp primitives essential for extracting maximum performance from modern GPUs**.