warp shuffle
**GPU Warp Shuffle Operations** are the **hardware-supported intrinsic instructions that allow threads within a warp (group of 32 threads executing in lockstep) to directly exchange register values without using shared memory** — enabling ultra-fast intra-warp communication with single-cycle latency and zero memory bandwidth consumption, making shuffle operations the fastest primitive for warp-level reductions, prefix scans, and data rearrangement in CUDA and GPU compute kernels.
**Why Shuffle Exists**
- Without shuffle: Thread 0 needs value from Thread 5 → write to shared memory → __syncthreads() → read.
- Cost: 2 memory transactions + synchronization barrier → ~20-30 cycles.
- With shuffle: __shfl_sync(mask, val, srcLane) → direct register-to-register transfer.
- Cost: ~1 cycle, no memory, no barrier.
**Shuffle Variants**
| Instruction | What It Does | Use Case |
|------------|-------------|----------|
| __shfl_sync(mask, val, srcLane) | Read val from specific lane | Broadcast, gather |
| __shfl_up_sync(mask, val, delta) | Read from lane (myLane - delta) | Prefix scan (inclusive) |
| __shfl_down_sync(mask, val, delta) | Read from lane (myLane + delta) | Reduction |
| __shfl_xor_sync(mask, val, laneMask) | Read from lane (myLane ^ mask) | Butterfly reduction |
**Warp Reduction (Sum)**
```cuda
__device__ float warpReduceSum(float val) {
// Butterfly reduction using XOR shuffle
val += __shfl_xor_sync(0xFFFFFFFF, val, 16); // Lanes 0-15 ↔ 16-31
val += __shfl_xor_sync(0xFFFFFFFF, val, 8); // Lanes 0-7 ↔ 8-15, etc.
val += __shfl_xor_sync(0xFFFFFFFF, val, 4);
val += __shfl_xor_sync(0xFFFFFFFF, val, 2);
val += __shfl_xor_sync(0xFFFFFFFF, val, 1);
return val; // All lanes have the sum
}
```
- 5 shuffle instructions → complete 32-element reduction.
- Shared memory reduction: ~10 transactions + barriers → 4-6× slower.
**Warp Prefix Scan**
```cuda
__device__ float warpPrefixSum(float val) {
float n;
n = __shfl_up_sync(0xFFFFFFFF, val, 1); if (threadIdx.x >= 1) val += n;
n = __shfl_up_sync(0xFFFFFFFF, val, 2); if (threadIdx.x >= 2) val += n;
n = __shfl_up_sync(0xFFFFFFFF, val, 4); if (threadIdx.x >= 4) val += n;
n = __shfl_up_sync(0xFFFFFFFF, val, 8); if (threadIdx.x >= 8) val += n;
n = __shfl_up_sync(0xFFFFFFFF, val, 16); if (threadIdx.x >= 16) val += n;
return val; // Inclusive prefix sum
}
```
**Broadcast**
```cuda
// All 32 lanes receive the value from lane 0
float shared_val = __shfl_sync(0xFFFFFFFF, my_val, 0);
```
**Performance Comparison**
| Operation | Shared Memory | Shuffle | Speedup |
|-----------|-------------|---------|--------|
| 32-element reduction | ~40 cycles | ~10 cycles | 4× |
| 32-element prefix scan | ~60 cycles | ~15 cycles | 4× |
| Broadcast from lane 0 | ~25 cycles | ~1 cycle | 25× |
**Practical Applications**
- **Softmax kernel**: Warp-level max reduction + sum reduction → fast attention.
- **LayerNorm**: Mean and variance computed via warp shuffle → fused kernel.
- **Histogram**: Warp-level partial histograms → reduce across warps.
- **Matrix transpose**: Shuffle for register-level data rearrangement.
- **FlashAttention**: Uses shuffle for warp-level coordination in tiled attention.
Warp shuffle operations are **the lowest-latency communication primitive available on GPUs** — by enabling direct register-to-register data exchange within a warp at single-cycle cost, shuffle instructions are the building block of every high-performance reduction, scan, and broadcast operation in modern GPU kernels, making them essential knowledge for anyone writing custom CUDA kernels for ML or scientific computing.