Warp Primitives: Register Exchange Without Shared Memory

Core Concept

__shfl_*_sync, __ballot_sync, and related operations exchange register values or construct masks within a warp without shared memory or a block-wide barrier.

Why It Affects Performance

Warp-level communication can reduce shared traffic, synchronization, and instruction count for reductions, scans, voting, and small rearrangements.

Example or Pseudocode

unsigned mask = __activemask();
for (int offset = 16; offset > 0; offset >>= 1)
  value += __shfl_down_sync(mask, value, offset);

A partial warp must use a mask that represents participating lanes.

Profiling and Performance Judgment

Compare shuffle and shared-memory reductions using instruction count, barrier stalls, shared transactions, and duration.

Five-Minute Exercise

Why can 0xffffffff be unsafe when the final warp has only 13 valid lanes?