AMDGPU / GFX Instructions GPU Native ISA AMD

30 instructions — AMDGPU/GFX is AMD's native, low-level GPU instruction set family, assembled directly with no virtual intermediate layer. Supported instructions and their exact encodings vary by GFX compatibility target (e.g. gfx942, gfx1100). Scalar instructions execute once per wavefront on the Scalar ALU; vector instructions execute per-lane on the Vector ALU, gated by the EXEC mask.

30 instructions
Mnemonic ⇅ Syntax Format ⇅ GFX Targets Unit ⇅ Summary
buffer_load_dword buffer_load_dword VDST, VADDR, SRSRC, offset MUBUF gfx942, gfx1100 Vector Memory Unit Load one 32-bit dword per lane through a buffer (raw/structured) resource descriptor.
ds_add_u32 ds_add_u32 ADDR, DATA, offset DS gfx942, gfx1100 LDS Unit Atomically add a per-lane value to an LDS location.
ds_permute_b32 ds_permute_b32 VDST, ADDR, DATA DS gfx942, gfx1100 LDS Unit Forward-permute: each lane sends its value to a lane index computed by another lane, via the LDS crossbar (no LDS storage consumed).
ds_read_b32 ds_read_b32 VDST, ADDR, offset DS gfx942, gfx1100 LDS Unit Read one 32-bit value per lane from the Local Data Share (LDS).
ds_write_b32 ds_write_b32 ADDR, DATA, offset DS gfx942, gfx1100 LDS Unit Write one 32-bit value per lane to the Local Data Share (LDS).
flat_load_dword flat_load_dword VDST, VADDR FLAT gfx942, gfx1100 Vector Memory Unit Load one 32-bit dword per lane through the flat (generic) address space, resolved to global/scratch/LDS at runtime.
global_atomic_add global_atomic_add VDST, VADDR, VDATA, SADDR GLOBAL gfx942, gfx1100 Vector Memory Unit Atomically add a per-lane value to a global-memory location.
global_load_dword global_load_dword VDST, VADDR, SADDR, offset GLOBAL gfx942, gfx1100 Vector Memory Unit Load one 32-bit dword per lane from the global address space using a 64-bit per-lane address.
global_store_dword global_store_dword VADDR, VDATA, SADDR, offset GLOBAL gfx942, gfx1100 Vector Memory Unit Store one 32-bit dword per lane to the global address space using a 64-bit per-lane address.
s_add_u32 s_add_u32 SDST, S0, S1 SOP2 gfx942, gfx1100 Scalar ALU Add two 32-bit unsigned scalar operands, wavefront-uniform.
s_barrier s_barrier SOPP gfx942, gfx1100 Scalar ALU Synchronize all waves of the executing workgroup at this point.
s_branch s_branch offset SOPP gfx942, gfx1100 Scalar ALU Unconditional relative branch.
s_cbranch_scc1 s_cbranch_scc1 offset SOPP gfx942, gfx1100 Scalar ALU Conditional relative branch, taken when the SCC (scalar condition code) flag is set.
s_cmp_eq_i32 s_cmp_eq_i32 S0, S1 SOPC gfx942, gfx1100 Scalar ALU Scalar signed-32-bit equality compare, result written to SCC.
s_load_dword s_load_dword SDST, SBASE, offset SMEM gfx942, gfx1100 Scalar Memory Unit Load one 32-bit dword from memory into a scalar register, wavefront-uniform.
s_mul_i32 s_mul_i32 SDST, S0, S1 SOP2 gfx942, gfx1100 Scalar ALU Multiply two 32-bit signed scalar operands, wavefront-uniform, low 32 bits of the product.
s_sub_u32 s_sub_u32 SDST, S0, S1 SOP2 gfx942, gfx1100 Scalar ALU Subtract two 32-bit unsigned scalar operands, wavefront-uniform.
v_add_f32 v_add_f32 VDST, S0, S1 VOP2 gfx942, gfx1100 Vector ALU Per-lane single-precision floating-point add.
v_add_u32 v_add_u32 VDST, S0, S1 VOP2 gfx942, gfx1100 Vector ALU Per-lane add of two 32-bit unsigned vector operands.
v_bcnt_u32_b32 v_bcnt_u32_b32 VDST, S0, S1 VOP3 gfx942, gfx1100 Vector ALU Per-lane accumulating population count.
v_cmp_eq_f32 v_cmp_eq_f32 VCC, S0, S1 VOPC gfx942, gfx1100 Vector ALU Per-lane single-precision equality compare, result written as an execution-mask-width bitmask.
v_cvt_f32_i32 v_cvt_f32_i32 VDST, S0 VOP1 gfx942, gfx1100 Vector ALU Per-lane conversion from signed 32-bit integer to single-precision float.
v_fma_f32 v_fma_f32 VDST, S0, S1, S2 VOP3 gfx942, gfx1100 Vector ALU Per-lane single-precision fused multiply-add.
v_mfma_f32_16x16x16f16 v_mfma_f32_16x16x16f16 D, A, B, C VOP3P-MAI gfx942 Matrix Core Unit Matrix-fused-multiply-add: cooperative 16x16x16 matrix-multiply-accumulate on matrix-core hardware, fp16 inputs, fp32 accumulate.
v_mul_f32 v_mul_f32 VDST, S0, S1 VOP2 gfx942, gfx1100 Vector ALU Per-lane single-precision floating-point multiply.
v_mul_lo_u32 v_mul_lo_u32 VDST, S0, S1 VOP3 gfx942, gfx1100 Vector ALU Per-lane 32-bit unsigned multiply, low half of the product.
v_readfirstlane_b32 v_readfirstlane_b32 SDST, S0 VOP1 gfx942, gfx1100 Vector ALU Read the value of a VGPR from the first active lane into a scalar register.
v_rsq_f32 v_rsq_f32 VDST, S0 VOP1 gfx942, gfx1100 Vector ALU Per-lane fast approximate reciprocal square root.
v_sin_f32 v_sin_f32 VDST, S0 VOP1 gfx942, gfx1100 Vector ALU Per-lane fast approximate sine.
v_sub_f32 v_sub_f32 VDST, S0, S1 VOP2 gfx942, gfx1100 Vector ALU Per-lane single-precision floating-point subtract.

Source

Normalized from AMD's official ROCm documentation, with the LLVM AMDGPU backend documentation as supplementary compiler-target information. ROCm documentation ↗