GPU Warps & Warp Shuffle

32 threads, one instruction, zero memory round-trips — press the buttons and watch it happen. (Demos show 8 lanes so everything fits on screen; real warps are 32.)

01 · SIMTAll lanes march to one clock

A GPU runs threads in groups of 32 called a warp. Every thread in the warp executes the same instruction on the same tick, each on its own register. A thread's position in the warp is its lane ID.

One instruction is issued per tick. Watch the clock pulse sweep the warp — every lane updates simultaneously, each with its own data.

v starts as the lane ID, so identical instructions still produce different values per lane — that's the "multiple threads" in SIMT.

02 · DIVERGENCEWhen lanes disagree, everyone pays

If threads in one warp take different branches, the hardware runs both paths one after the other, masking off the lanes that didn't take each path.

if (laneId < 4)  v += 100;   // path A
else            v *= 2;     // path B

PASS 1 · A() · lanes 0–3 PASS 2 · B() · lanes 4–7

Grayed-out lanes are masked: they occupy the hardware but do nothing. One branch, two passes — 2× the time.

03 · SHUFFLELanes read each other's registers

Normally threads trade data through shared memory. Shuffle instructions skip memory entirely: in one instruction, every lane grabs a value straight out of another lane's register. Pick a mode and fire it.

04 · REDUCTIONSum a whole warp in log₂ steps

Chain __shfl_xor_sync with halving offsets and every lane ends up holding the total — no shared memory, no __syncthreads(). Eight lanes need 3 rounds; a full warp of 32 needs just 5.

v += __shfl_xor_sync(0xffffffff, v, 4);
v += __shfl_xor_sync(0xffffffff, v, 2);
v += __shfl_xor_sync(0xffffffff, v, 1);
// every lane now holds the sum of all 8

Each round, partners exchange values and add. Partial sums spread like a butterfly network until the total is everywhere.

05 · WHY IT MATTERSRegisters beat memory

Shuffle is register-to-register. Lower latency than shared memory, zero shared-memory footprint (better occupancy), and no barrier needed — a warp is already in lockstep.

The mask argument (usually 0xffffffff) exists because since Volta, threads in a warp can execute independently — you must declare which lanes participate. And AMD's equivalent group is the wavefront (32 or 64 lanes).