Lane Shuffle Warp/Wavefront-Level

Exchange a value between lanes of the same warp/wavefront without going through memory.

Vendor-Neutral Definition

d[lane] = src[shuffle_source_lane(lane)], evaluated across the active lanes of one warp (NVIDIA) or wavefront (AMD); out-of-range or inactive source lanes are handled per the specific instruction's rules.
⚠️ Cross-vendor mappings describe semantic relationships and are not guaranteed one-to-one compiler translations.
equivalent with restrictions Both exchange a value between lanes of the same warp/wavefront without going through addressable memory.
Property NVIDIA PTX AMDGPU
Instruction shfl ds_permute_b32, v_readfirstlane_b32
ISA Layer Virtual Native
Data Types — b32
Version / Target Introduced PTX ISA 3.0 gfx942/gfx1100, gfx942/gfx1100

Important Differences

  • NVIDIA warps are fixed at 32 lanes; AMD wavefronts are 32 or 64 lanes depending on GFX target and compile mode (wave32 vs wave64), which changes the valid shuffle-index range.
  • PTX shfl.sync takes an explicit member mask argument; AMDGPU's permute instructions operate relative to the EXEC mask rather than a per-call mask argument.

Source Evidence

Verification method: documentation