AMDGPU / GFX Instructions GPU Native ISA AMD
1960 instructions - showing 100 per page, page 17 of 20 - 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.
| Mnemonic | Syntax | Format | GFX Targets | Unit | Summary |
|---|---|---|---|---|---|
| v_mad_u64_u32 | v_mad_u64_u32 | VOP3 | Multiply two unsigned integer inputs, add a third unsigned integer input, store the result into a 64-bit vector register and store the… | ||
| v_madak_f16 | v_madak_f16 | VOP2 | Multiply two floating point inputs and add a literal constant, and store the result into a vector register. Implements IEEE rules. | ||
| v_madak_f32 | v_madak_f32 | VOP2 | Multiply two floating point inputs and add a literal constant, and store the result into a vector register. | ||
| v_madmk_f16 | v_madmk_f16 | VOP2 | Multiply a floating point input with a literal constant and add a second floating point input, and store the result into a vector register. | ||
| v_madmk_f32 | v_madmk_f32 | VOP2 | Multiply a floating point input with a literal constant and add a second floating point input, and store the result into a vector register. | ||
| v_max3_f16 | v_max3_f16 | VOP3 | gfx1100 | Select the maximum of three half-precision float inputs and store the selected value into a vector register. | |
| v_max3_f32 | v_max3_f32 | VOP3 | gfx1100 | Select the maximum of three single-precision float inputs and store the selected value into a vector register. | |
| v_max3_i16 | v_max3_i16 | VOP3 | gfx1100 | Select the maximum of three signed 16-bit integer inputs and store the selected value into a vector register. | |
| v_max3_i32 | v_max3_i32 | VOP3 | gfx1100 | Select the maximum of three signed 32-bit integer inputs and store the selected value into a vector register. | |
| v_max3_num_f16 | v_max3_num_f16 | VOP3 | Select the IEEE maximumNumber() of three half-precision float inputs and store the selected value into a vector register. | ||
| v_max3_num_f32 | v_max3_num_f32 | VOP3 | Select the IEEE maximumNumber() of three single-precision float inputs and store the selected value into a vector register. | ||
| v_max3_u16 | v_max3_u16 | VOP3 | gfx1100 | Select the maximum of three unsigned 16-bit integer inputs and store the selected value into a vector register. | |
| v_max3_u32 | v_max3_u32 | VOP3 | gfx1100 | Select the maximum of three unsigned 32-bit integer inputs and store the selected value into a vector register. | |
| v_max_bf16 | v_max_bf16 | VOP3 | AMDGPU VOP3 vector instruction. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) | ||
| v_max_f16 | v_max_f16 | VOP2 | gfx1100 | Select the maximum of two half-precision float inputs and store the result into a vector register. | |
| v_max_f32 | v_max_f32 | VOP2 | gfx1100 | Select the maximum of two single-precision float inputs and store the result into a vector register. | |
| v_max_f64 | v_max_f64 | VOP2 | Select the maximum of two double-precision float inputs and store the selected value into a vector register. | ||
| v_max_i16 | v_max_i16 | VOP2 | gfx1100 | Select the maximum of two signed 16-bit integer inputs and store the selected value into a vector register. | |
| v_max_i32 | v_max_i32 | VOP2 | gfx1100 | Select the maximum of two signed 32-bit integer inputs and store the selected value into a vector register. | |
| v_max_i64 | v_max_i64 | VOP3 | AMDGPU VOP3 vector instruction operating on i64 data. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) | ||
| v_max_legacy_f32 | v_max_legacy_f32 | VOP2 | AMDGPU VOP2 vector instruction operating on f32 data. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) | ||
| v_max_num_f16 | v_max_num_f16 | VOP2 | Select the IEEE maximumNumber() of two half-precision float inputs and store the selected value into a vector register. | ||
| v_max_num_f32 | v_max_num_f32 | VOP2 | Select the IEEE maximumNumber() of two single-precision float inputs and store the selected value into a vector register. | ||
| v_max_num_f64 | v_max_num_f64 | VOP2 | Select the IEEE maximumNumber() of two double-precision float inputs and store the selected value into a vector register. | ||
| v_max_u16 | v_max_u16 | VOP2 | gfx1100 | Select the maximum of two unsigned 16-bit integer inputs and store the selected value into a vector register. | |
| v_max_u32 | v_max_u32 | VOP2 | gfx1100 | Select the maximum of two unsigned 32-bit integer inputs and store the selected value into a vector register. | |
| v_max_u64 | v_max_u64 | VOP3 | AMDGPU VOP3 vector instruction operating on u64 data. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) | ||
| v_maximum3_f16 | v_maximum3_f16 | VOP3 | gfx1100 | Select the IEEE maximum() of three half-precision float inputs and store the selected value into a vector register. | |
| v_maximum3_f32 | v_maximum3_f32 | VOP3 | gfx1100 | Select the IEEE maximum() of three single-precision float inputs and store the result into a vector register. | |
| v_maximum_f16 | v_maximum_f16 | VOP3 | gfx1100 | Select the IEEE maximum() of two half-precision float inputs and store the selected value into a vector register. | |
| v_maximum_f32 | v_maximum_f32 | VOP3 | gfx1100 | Select the IEEE maximum() of two single-precision float inputs and store the selected value into a vector register. | |
| v_maximum_f64 | v_maximum_f64 | VOP3 | gfx1100 | Select the IEEE maximum() of two double-precision float inputs and store the selected value into a vector register. | |
| v_maximumminimum_f16 | v_maximumminimum_f16 | VOP3 | gfx1100 | Select the IEEE maximum() of the first two half-precision float inputs and then select the IEEE minimum() of that result and third half-precision… | |
| v_maximumminimum_f32 | v_maximumminimum_f32 | VOP3 | gfx1100 | Select the IEEE maximum() of the first two single-precision float inputs and then select the IEEE minimum() of that result and third single-precision… | |
| v_maxmin_f16 | v_maxmin_f16 | VOP3 | gfx1100 | Select the maximum of the first two half-precision float inputs and then select the minimum of that result and third half-precision float input. | |
| v_maxmin_f32 | v_maxmin_f32 | VOP3 | gfx1100 | Select the maximum of the first two single-precision float inputs and then select the minimum of that result and third single-precision float input. | |
| v_maxmin_i32 | v_maxmin_i32 | VOP3 | gfx1100 | Select the maximum of the first two signed 32-bit integer inputs and then select the minimum of that result and third signed 32-bit integer input. | |
| v_maxmin_num_f16 | v_maxmin_num_f16 | VOP3 | Select the IEEE maximumNumber() of the first two half-precision float inputs and then select the IEEE minimumNumber() of that result and third… | ||
| v_maxmin_num_f32 | v_maxmin_num_f32 | VOP3 | Select the IEEE maximumNumber() of the first two single-precision float inputs and then select the IEEE minimumNumber() of that result and third… | ||
| v_maxmin_u32 | v_maxmin_u32 | VOP3 | gfx1100 | Select the maximum of the first two unsigned 32-bit integer inputs and then select the minimum of that result and third unsigned 32-bit integer input. | |
| v_mbcnt_hi_u32_b32 | v_mbcnt_hi_u32_b32 | VOP2 | gfx1100 | For each lane 32 <= N < 64, examine the N least significant bits of the first input and count how many of those bits are "1". | |
| v_mbcnt_lo_u32_b32 | v_mbcnt_lo_u32_b32 | VOP2 | gfx1100 | For each lane 0 <= N < 32, examine the N least significant bits of the first input and count how many of those bits are "1". | |
| v_med3_f16 | v_med3_f16 | VOP3 | gfx1100 | Select the median of three half-precision float values and store the selected value into a vector register. | |
| v_med3_f32 | v_med3_f32 | VOP3 | gfx1100 | Select the median of three single-precision float values and store the selected value into a vector register. | |
| v_med3_i16 | v_med3_i16 | VOP3 | gfx1100 | Select the median of three signed 16-bit integer values and store the selected value into a vector register. | |
| v_med3_i32 | v_med3_i32 | VOP3 | gfx1100 | Select the median of three signed 32-bit integer values and store the selected value into a vector register. | |
| v_med3_num_f16 | v_med3_num_f16 | VOP3 | Select the median of three half-precision float inputs and store the selected value into a vector register. | ||
| v_med3_num_f32 | v_med3_num_f32 | VOP3 | Select the median of three single-precision float inputs and store the selected value into a vector register. | ||
| v_med3_u16 | v_med3_u16 | VOP3 | gfx1100 | Select the median of three unsigned 16-bit integer values and store the selected value into a vector register. | |
| v_med3_u32 | v_med3_u32 | VOP3 | gfx1100 | Select the median of three unsigned 32-bit integer values and store the selected value into a vector register. | |
| v_mfma_f32_16x16x128_f8f6f4 | v_mfma_f32_16x16x128_f8f6f4 | VOP3P | Multiply the 16x128 matrix in the first input by the 128x16 matrix in the second input and add the 16x16 matrix in the third input using fused… | ||
| v_mfma_f32_16x16x16_bf16 | v_mfma_f32_16x16x16_bf16 | VOP3P | gfx942 | Multiply the 16x16 matrix in the first input by the 16x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… | |
| v_mfma_f32_16x16x16_f16 | v_mfma_f32_16x16x16_f16 | VOP3P | Multiply the 16x16 matrix in the first input by the 16x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… | ||
| 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_mfma_f32_16x16x1_4b_f32 | v_mfma_f32_16x16x1_4b_f32 | VOP3P | Multiply the 16x1 matrix in the first input by the 1x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… | ||
| v_mfma_f32_16x16x32_bf16 | v_mfma_f32_16x16x32_bf16 | VOP3P | Multiply the 16x32 matrix in the first input by the 32x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… | ||
| v_mfma_f32_16x16x32_f16 | v_mfma_f32_16x16x32_f16 | VOP3P | Multiply the 16x32 matrix in the first input by the 32x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… | ||
| v_mfma_f32_16x16x4_4b_bf16 | v_mfma_f32_16x16x4_4b_bf16 | VOP3P | gfx942 | Multiply the 16x4 matrix in the first input by the 4x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… | |
| v_mfma_f32_16x16x4_4b_f16 | v_mfma_f32_16x16x4_4b_f16 | VOP3P | Multiply the 16x4 matrix in the first input by the 4x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… | ||
| v_mfma_f32_16x16x4_f32 | v_mfma_f32_16x16x4_f32 | VOP3P | Multiply the 16x4 matrix in the first input by the 4x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… | ||
| v_mfma_f32_16x16x8_xf32 | v_mfma_f32_16x16x8_xf32 | VOP3P | gfx942 | Multiply the 16x8 matrix in the first input by the 8x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… | |
| v_mfma_f32_32x32x16_bf16 | v_mfma_f32_32x32x16_bf16 | VOP3P | Multiply the 32x16 matrix in the first input by the 16x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… | ||
| v_mfma_f32_32x32x16_f16 | v_mfma_f32_32x32x16_f16 | VOP3P | Multiply the 32x16 matrix in the first input by the 16x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… | ||
| v_mfma_f32_32x32x1_2b_f32 | v_mfma_f32_32x32x1_2b_f32 | VOP3P | Multiply the 32x1 matrix in the first input by the 1x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… | ||
| v_mfma_f32_32x32x2_f32 | v_mfma_f32_32x32x2_f32 | VOP3P | Multiply the 32x2 matrix in the first input by the 2x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… | ||
| v_mfma_f32_32x32x4_2b_bf16 | v_mfma_f32_32x32x4_2b_bf16 | VOP3P | gfx942 | Multiply the 32x4 matrix in the first input by the 4x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… | |
| v_mfma_f32_32x32x4_2b_f16 | v_mfma_f32_32x32x4_2b_f16 | VOP3P | Multiply the 32x4 matrix in the first input by the 4x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… | ||
| v_mfma_f32_32x32x4_xf32 | v_mfma_f32_32x32x4_xf32 | VOP3P | gfx942 | Multiply the 32x4 matrix in the first input by the 4x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… | |
| v_mfma_f32_32x32x64_f8f6f4 | v_mfma_f32_32x32x64_f8f6f4 | VOP3P | Multiply the 32x64 matrix in the first input by the 64x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… | ||
| v_mfma_f32_32x32x8_bf16 | v_mfma_f32_32x32x8_bf16 | VOP3P | gfx942 | Multiply the 32x8 matrix in the first input by the 8x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… | |
| v_mfma_f32_32x32x8_f16 | v_mfma_f32_32x32x8_f16 | VOP3P | Multiply the 32x8 matrix in the first input by the 8x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… | ||
| v_mfma_f32_4x4x1_16b_f32 | v_mfma_f32_4x4x1_16b_f32 | VOP3P | Multiply the 4x1 matrix in the first input by the 1x4 matrix in the second input and add the 4x4 matrix in the third input using fused multiply add. | ||
| v_mfma_f32_4x4x4_16b_bf16 | v_mfma_f32_4x4x4_16b_bf16 | VOP3P | gfx942 | Multiply the 4x4 matrix in the first input by the 4x4 matrix in the second input and add the 4x4 matrix in the third input using fused multiply add. | |
| v_mfma_f32_4x4x4_16b_f16 | v_mfma_f32_4x4x4_16b_f16 | VOP3P | Multiply the 4x4 matrix in the first input by the 4x4 matrix in the second input and add the 4x4 matrix in the third input using fused multiply add. | ||
| v_mfma_f64_16x16x4_f64 | v_mfma_f64_16x16x4_f64 | VOP3P | gfx942 | Multiply the 16x4 matrix in the first input by the 4x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… | |
| v_mfma_f64_4x4x4_4b_f64 | v_mfma_f64_4x4x4_4b_f64 | VOP3P | gfx942 | Multiply the 4x4 matrix in the first input by the 4x4 matrix in the second input and add the 4x4 matrix in the third input using fused multiply add. | |
| v_mfma_i32_16x16x32_i8 | v_mfma_i32_16x16x32_i8 | VOP3P | gfx942 | Multiply the 16x32 matrix in the first input by the 32x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… | |
| v_mfma_i32_16x16x4_4b_i8 | v_mfma_i32_16x16x4_4b_i8 | VOP3P | Multiply the 16x4 matrix in the first input by the 4x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… | ||
| v_mfma_i32_16x16x64_i8 | v_mfma_i32_16x16x64_i8 | VOP3P | Multiply the 16x64 matrix in the first input by the 64x16 matrix in the second input and add the 16x16 matrix in the third input using fused multiply… | ||
| v_mfma_i32_32x32x16_i8 | v_mfma_i32_32x32x16_i8 | VOP3P | gfx942 | Multiply the 32x16 matrix in the first input by the 16x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… | |
| v_mfma_i32_32x32x32_i8 | v_mfma_i32_32x32x32_i8 | VOP3P | Multiply the 32x32 matrix in the first input by the 32x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… | ||
| v_mfma_i32_32x32x4_2b_i8 | v_mfma_i32_32x32x4_2b_i8 | VOP3P | Multiply the 32x4 matrix in the first input by the 4x32 matrix in the second input and add the 32x32 matrix in the third input using fused multiply… | ||
| v_mfma_i32_4x4x4_16b_i8 | v_mfma_i32_4x4x4_16b_i8 | VOP3P | Multiply the 4x4 matrix in the first input by the 4x4 matrix in the second input and add the 4x4 matrix in the third input using fused multiply add. | ||
| v_mfma_ld_scale_b32 | v_mfma_ld_scale_b32 | VOP3P | AMDGPU VOP3P matrix instruction operating on b32 data. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) | ||
| v_min3_f16 | v_min3_f16 | VOP3 | gfx1100 | Select the minimum of three half-precision float inputs and store the selected value into a vector register. | |
| v_min3_f32 | v_min3_f32 | VOP3 | gfx1100 | Select the minimum of three single-precision float inputs and store the selected value into a vector register. | |
| v_min3_i16 | v_min3_i16 | VOP3 | gfx1100 | Select the minimum of three signed 16-bit integer inputs and store the selected value into a vector register. | |
| v_min3_i32 | v_min3_i32 | VOP3 | gfx1100 | Select the minimum of three signed 32-bit integer inputs and store the selected value into a vector register. | |
| v_min3_num_f16 | v_min3_num_f16 | VOP3 | Select the IEEE minimumNumber() of three half-precision float inputs and store the selected value into a vector register. | ||
| v_min3_num_f32 | v_min3_num_f32 | VOP3 | Select the IEEE minimumNumber() of three single-precision float inputs and store the selected value into a vector register. | ||
| v_min3_u16 | v_min3_u16 | VOP3 | gfx1100 | Select the minimum of three unsigned 16-bit integer inputs and store the selected value into a vector register. | |
| v_min3_u32 | v_min3_u32 | VOP3 | gfx1100 | Select the minimum of three unsigned 32-bit integer inputs and store the selected value into a vector register. | |
| v_min_f16 | v_min_f16 | VOP2 | gfx1100 | Select the minimum of two half-precision float inputs and store the result into a vector register. | |
| v_min_f32 | v_min_f32 | VOP2 | gfx1100 | Select the minimum of two single-precision float inputs and store the result into a vector register. | |
| v_min_f64 | v_min_f64 | VOP2 | Select the minimum of two double-precision float inputs and store the selected value into a vector register. | ||
| v_min_i16 | v_min_i16 | VOP2 | gfx1100 | Select the minimum of two signed 16-bit integer inputs and store the selected value into a vector register. | |
| v_min_i32 | v_min_i32 | VOP2 | gfx1100 | Select the minimum of two signed 32-bit integer inputs and store the selected value into a vector register. | |
| v_min_i64 | v_min_i64 | VOP3 | AMDGPU VOP3 vector instruction operating on i64 data. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) | ||
| v_min_legacy_f32 | v_min_legacy_f32 | VOP2 | AMDGPU VOP2 vector instruction operating on f32 data. (Format and name extracted from LLVM's AMDGPU backend source - semantics not yet curated.) | ||
| v_min_num_f16 | v_min_num_f16 | VOP2 | Select the IEEE minimumNumber() of two half-precision float inputs and store the selected value into a vector register. |
Source
Normalized from AMD's official ROCm documentation, with the LLVM AMDGPU backend documentation as supplementary compiler-target information. ROCm documentation ↗