Read Lane to Scalar Register Warp/Wavefront-Level

Move a value out of one lane of a vector register into a scalar register shared by the whole wavefront.

Vendor-Neutral Definition

s = v[lane], where s is uniform across the wavefront. The point is to convert a per-lane value the compiler can prove is uniform into one the scalar unit can use.
⚠️ Cross-vendor mappings describe semantic relationships and are not guaranteed one-to-one compiler translations.
equivalent with restrictions AMDGPU has a real scalar unit and dedicated instructions to feed it. PTX has no scalar register file at all, so the nearest equivalent broadcasts within the warp instead.
Property NVIDIA PTX AMDGPU
Instruction shfl v_readfirstlane_b32, v_readlane_b32
ISA Layer Virtual Native
Data Types - b32
Version / Target Introduced PTX ISA 6.0 gfx942/gfx1100, gfx1100

Important Differences

  • This is an architectural asymmetry rather than a naming difference. AMDGPU wavefronts have SGPRs executed by a separate scalar ALU, and v_readfirstlane_b32 is how a value crosses from vector to scalar. NVIDIA has no such register file exposed in PTX.
  • The closest PTX construction is shfl.sync.idx to broadcast one lane's value to all lanes, which leaves the result in a per-lane register rather than a uniform one, so it saves no register pressure and no issue slots.
  • v_readfirstlane_b32 reads the lowest active lane; v_readlane_b32 reads a lane the instruction names.
  • The inverse direction is equally asymmetric: AMDGPU's v_writelane_b32 writes a scalar into one named lane, and PTX has nothing comparable. A thread can only write its own registers, so the effect has to go through shared memory or a shuffle every lane joins.

Source Evidence

Verification method: documentation