shfl GPU Virtual ISA NVIDIA

Shuffle Warp-Level

shfl.mode.b32 d[|p], a, b, c;

Exchange a value directly between lanes of the same warp.

Encoding

PTX is a virtual instruction set. It has no single, stable native binary encoding — the compiler lowers this instruction to different native machine code depending on the selected NVIDIA target architecture (compute capability). This page intentionally shows no bit-diagram; see the target/version requirements below for what governs how this instruction compiles.
PTX ISA Version Introduced PTX ISA 3.0
Minimum Target sm_30
Deprecation PTX ISA 6.0 (legacy non-sync form only)

Syntax Forms

One mnemonic covers many type / state-space / scope / modifier combinations — each row below is an independently valid form.

Syntax Data Types State Space(s) Modifiers Min. Target Description
shfl.mode.b32 d[|p], a, b, c; up, down, bfly, idx sm_30 Legacy (unsynchronized) warp shuffle; deprecated since PTX ISA 6.0 in favor of shfl.sync.
shfl.sync.mode.b32 d[|p], a, b, c, membermask; up, down, bfly, idx sm_30 Warp shuffle that also synchronizes the specified member lanes before exchanging data.

Operands

  • d
    Destination register
  • a
    Value to shuffle
  • b
    Source-lane selector
  • c
    Clamp/width control
  • membermask
    Mask of participating lanes (sync form only)

At a Glance

Data Types —
Modifiers bfly, down, idx, up

Related AMDGPU Concepts

Lane Shuffle ↗
equivalent with restrictions
ds_permute_b32 (AMDGPU)

Reference

NVIDIA PTX ISA

Semantics

d = value of operand a as seen by another lane in the warp, selected by mode/b; optional predicate p reports whether the source lane was valid.

Sources