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.