AMDGPU Opcode Map
Reverse lookup instructions by their hexadecimal encoding.
| Opcode | Mnemonic | Syntax | Format | Name |
|---|---|---|---|---|
| v_max_u16 | v_max_u16 | VOP2 | V MAX U16 | |
| 0x28000000 | v_max_u32 | v_max_u32 | VOP2 | V MAX U32 |
| 0x00000000D7190000 | v_max_u64 | v_max_u64 | VOP3 | V MAX U64 |
| 0x00000000D6300000 | v_maximum3_f16 | v_maximum3_f16 | VOP3 | V MAXIMUM3 F16 |
| 0x00000000D62E0000 | v_maximum3_f32 | v_maximum3_f32 | VOP3 | V MAXIMUM3 F32 |
| 0x00000000D7680000 | v_maximum_f16 | v_maximum_f16 | VOP3 | V MAXIMUM F16 |
| 0x00000000D7660000 | v_maximum_f32 | v_maximum_f32 | VOP3 | V MAXIMUM F32 |
| 0x00000000D7420000 | v_maximum_f64 | v_maximum_f64 | VOP3 | V MAXIMUM F64 |
| 0x00000000D66F0000 | v_maximumminimum_f16 | v_maximumminimum_f16 | VOP3 | V MAXIMUMMINIMUM F16 |
| 0x00000000D66D0000 | v_maximumminimum_f32 | v_maximumminimum_f32 | VOP3 | V MAXIMUMMINIMUM F32 |
| 0x00000000D6600000 | v_maxmin_f16 | v_maxmin_f16 | VOP3 | V MAXMIN F16 |
| 0x00000000D65E0000 | v_maxmin_f32 | v_maxmin_f32 | VOP3 | V MAXMIN F32 |
| 0x00000000D6640000 | v_maxmin_i32 | v_maxmin_i32 | VOP3 | V MAXMIN I32 |
| 0x00000000D66B0000 | v_maxmin_num_f16 | v_maxmin_num_f16 | VOP3 | V MAXMIN NUM F16 |
| 0x00000000D6690000 | v_maxmin_num_f32 | v_maxmin_num_f32 | VOP3 | V MAXMIN NUM F32 |
| 0x00000000D6620000 | v_maxmin_u32 | v_maxmin_u32 | VOP3 | V MAXMIN U32 |
| v_mbcnt_hi_u32_b32 | v_mbcnt_hi_u32_b32 | VOP2 | V MBCNT HI U32 B32 | |
| v_mbcnt_lo_u32_b32 | v_mbcnt_lo_u32_b32 | VOP2 | V MBCNT LO U32 B32 | |
| 0x00000000D1FA0000 | v_med3_f16 | v_med3_f16 | VOP3 | V MED3 F16 |
| 0x00000000D1D60000 | v_med3_f32 | v_med3_f32 | VOP3 | V MED3 F32 |
| 0x00000000D6500000 | v_med3_i16 | v_med3_i16 | VOP3 | V MED3 I16 |
| 0x00000000D6200000 | v_med3_i32 | v_med3_i32 | VOP3 | V MED3 I32 |
| 0x00000000D6320000 | v_med3_num_f16 | v_med3_num_f16 | VOP3 | V MED3 NUM F16 |
| 0x00000000D6310000 | v_med3_num_f32 | v_med3_num_f32 | VOP3 | V MED3 NUM F32 |
| 0x00000000D6510000 | v_med3_u16 | v_med3_u16 | VOP3 | V MED3 U16 |
| 0x00000000D6210000 | v_med3_u32 | v_med3_u32 | VOP3 | V MED3 U32 |
| 0x00000000D3AD0000 | v_mfma_f32_16x16x128_f8f6f4 | v_mfma_f32_16x16x128_f8f6f4 | VOP3P | V MFMA F32 16X16X128 F8F6F4 |
| 0x00000000D3E10000 | v_mfma_f32_16x16x16_bf16 | v_mfma_f32_16x16x16_bf16 | VOP3P | V MFMA F32 16X16X16 BF16 |
| 0x00000000D3CD0000 | v_mfma_f32_16x16x16_f16 | v_mfma_f32_16x16x16_f16 | VOP3P | V MFMA F32 16X16X16 F16 |
| 0x00000000D3E70000 | v_mfma_f32_16x16x16bf16_1k | v_mfma_f32_16x16x16bf16_1k | VOP3P | V MFMA F32 16X16X16BF16 1K |
| 0x00000000D3CD0000 | v_mfma_f32_16x16x16f16 | v_mfma_f32_16x16x16f16 D, A, B, C | VOP3P-MAI | V MFMA F32 16X16X16F16 |
| 0x00000000D3C10000 | v_mfma_f32_16x16x1_4b_f32 | v_mfma_f32_16x16x1_4b_f32 | VOP3P | V MFMA F32 16X16X1 4B F32 |
| 0x00000000D3C10000 | v_mfma_f32_16x16x1f32 | v_mfma_f32_16x16x1f32 | VOP3P | V MFMA F32 16X16X1F32 |
| 0x00000000D3E90000 | v_mfma_f32_16x16x2bf16 | v_mfma_f32_16x16x2bf16 | VOP3P | V MFMA F32 16X16X2BF16 |
| 0x00000000D3B50000 | v_mfma_f32_16x16x32_bf16 | v_mfma_f32_16x16x32_bf16 | VOP3P | V MFMA F32 16X16X32 BF16 |
| 0x00000000D3F00000 | v_mfma_f32_16x16x32_bf8_bf8 | v_mfma_f32_16x16x32_bf8_bf8 | VOP3P | V MFMA F32 16X16X32 BF8 BF8 |
| 0x00000000D3F10000 | v_mfma_f32_16x16x32_bf8_fp8 | v_mfma_f32_16x16x32_bf8_fp8 | VOP3P | V MFMA F32 16X16X32 BF8 FP8 |
| 0x00000000D3D40000 | v_mfma_f32_16x16x32_f16 | v_mfma_f32_16x16x32_f16 | VOP3P | V MFMA F32 16X16X32 F16 |
| 0x00000000D3F20000 | v_mfma_f32_16x16x32_fp8_bf8 | v_mfma_f32_16x16x32_fp8_bf8 | VOP3P | V MFMA F32 16X16X32 FP8 BF8 |
| 0x00000000D3F30000 | v_mfma_f32_16x16x32_fp8_fp8 | v_mfma_f32_16x16x32_fp8_fp8 | VOP3P | V MFMA F32 16X16X32 FP8 FP8 |
| 0x00000000D3DE0000 | v_mfma_f32_16x16x4_4b_bf16 | v_mfma_f32_16x16x4_4b_bf16 | VOP3P | V MFMA F32 16X16X4 4B BF16 |
| 0x00000000D3C90000 | v_mfma_f32_16x16x4_4b_f16 | v_mfma_f32_16x16x4_4b_f16 | VOP3P | V MFMA F32 16X16X4 4B F16 |
| 0x00000000D3C50000 | v_mfma_f32_16x16x4_f32 | v_mfma_f32_16x16x4_f32 | VOP3P | V MFMA F32 16X16X4 F32 |
| 0x00000000D3E40000 | v_mfma_f32_16x16x4bf16_1k | v_mfma_f32_16x16x4bf16_1k | VOP3P | V MFMA F32 16X16X4BF16 1K |
| 0x00000000D3C90000 | v_mfma_f32_16x16x4f16 | v_mfma_f32_16x16x4f16 | VOP3P | V MFMA F32 16X16X4F16 |
| 0x00000000D3C50000 | v_mfma_f32_16x16x4f32 | v_mfma_f32_16x16x4f32 | VOP3P | V MFMA F32 16X16X4F32 |
| 0x00000000D3BE0000 | v_mfma_f32_16x16x8_xf32 | v_mfma_f32_16x16x8_xf32 | VOP3P | V MFMA F32 16X16X8 XF32 |
| 0x00000000D3ED0000 | v_mfma_f32_16x16x8bf16 | v_mfma_f32_16x16x8bf16 | VOP3P | V MFMA F32 16X16X8BF16 |
| 0x00000000D3B70000 | v_mfma_f32_32x32x16_bf16 | v_mfma_f32_32x32x16_bf16 | VOP3P | V MFMA F32 32X32X16 BF16 |
| 0x00000000D3F40000 | v_mfma_f32_32x32x16_bf8_bf8 | v_mfma_f32_32x32x16_bf8_bf8 | VOP3P | V MFMA F32 32X32X16 BF8 BF8 |
| 0x00000000D3F50000 | v_mfma_f32_32x32x16_bf8_fp8 | v_mfma_f32_32x32x16_bf8_fp8 | VOP3P | V MFMA F32 32X32X16 BF8 FP8 |
| 0x00000000D3D50000 | v_mfma_f32_32x32x16_f16 | v_mfma_f32_32x32x16_f16 | VOP3P | V MFMA F32 32X32X16 F16 |
| 0x00000000D3F60000 | v_mfma_f32_32x32x16_fp8_bf8 | v_mfma_f32_32x32x16_fp8_bf8 | VOP3P | V MFMA F32 32X32X16 FP8 BF8 |
| 0x00000000D3F70000 | v_mfma_f32_32x32x16_fp8_fp8 | v_mfma_f32_32x32x16_fp8_fp8 | VOP3P | V MFMA F32 32X32X16 FP8 FP8 |
| 0x00000000D3C00000 | v_mfma_f32_32x32x1_2b_f32 | v_mfma_f32_32x32x1_2b_f32 | VOP3P | V MFMA F32 32X32X1 2B F32 |
| 0x00000000D3C00000 | v_mfma_f32_32x32x1f32 | v_mfma_f32_32x32x1f32 | VOP3P | V MFMA F32 32X32X1F32 |
| 0x00000000D3C40000 | v_mfma_f32_32x32x2_f32 | v_mfma_f32_32x32x2_f32 | VOP3P | V MFMA F32 32X32X2 F32 |
| 0x00000000D3E80000 | v_mfma_f32_32x32x2bf16 | v_mfma_f32_32x32x2bf16 | VOP3P | V MFMA F32 32X32X2BF16 |
| 0x00000000D3C40000 | v_mfma_f32_32x32x2f32 | v_mfma_f32_32x32x2f32 | VOP3P | V MFMA F32 32X32X2F32 |
| 0x00000000D3DD0000 | v_mfma_f32_32x32x4_2b_bf16 | v_mfma_f32_32x32x4_2b_bf16 | VOP3P | V MFMA F32 32X32X4 2B BF16 |
| 0x00000000D3C80000 | v_mfma_f32_32x32x4_2b_f16 | v_mfma_f32_32x32x4_2b_f16 | VOP3P | V MFMA F32 32X32X4 2B F16 |
| 0x00000000D3BF0000 | v_mfma_f32_32x32x4_xf32 | v_mfma_f32_32x32x4_xf32 | VOP3P | V MFMA F32 32X32X4 XF32 |
| 0x00000000D3EC0000 | v_mfma_f32_32x32x4bf16 | v_mfma_f32_32x32x4bf16 | VOP3P | V MFMA F32 32X32X4BF16 |
| 0x00000000D3E30000 | v_mfma_f32_32x32x4bf16_1k | v_mfma_f32_32x32x4bf16_1k | VOP3P | V MFMA F32 32X32X4BF16 1K |
| 0x00000000D3C80000 | v_mfma_f32_32x32x4f16 | v_mfma_f32_32x32x4f16 | VOP3P | V MFMA F32 32X32X4F16 |
| 0x00000000D3AE0000 | v_mfma_f32_32x32x64_f8f6f4 | v_mfma_f32_32x32x64_f8f6f4 | VOP3P | V MFMA F32 32X32X64 F8F6F4 |
| 0x00000000D3E00000 | v_mfma_f32_32x32x8_bf16 | v_mfma_f32_32x32x8_bf16 | VOP3P | V MFMA F32 32X32X8 BF16 |
| 0x00000000D3CC0000 | v_mfma_f32_32x32x8_f16 | v_mfma_f32_32x32x8_f16 | VOP3P | V MFMA F32 32X32X8 F16 |
| 0x00000000D3E60000 | v_mfma_f32_32x32x8bf16_1k | v_mfma_f32_32x32x8bf16_1k | VOP3P | V MFMA F32 32X32X8BF16 1K |
| 0x00000000D3CC0000 | v_mfma_f32_32x32x8f16 | v_mfma_f32_32x32x8f16 | VOP3P | V MFMA F32 32X32X8F16 |
| 0x00000000D3C20000 | v_mfma_f32_4x4x1_16b_f32 | v_mfma_f32_4x4x1_16b_f32 | VOP3P | V MFMA F32 4X4X1 16B F32 |
| 0x00000000D3C20000 | v_mfma_f32_4x4x1f32 | v_mfma_f32_4x4x1f32 | VOP3P | V MFMA F32 4X4X1F32 |
| 0x00000000D3EB0000 | v_mfma_f32_4x4x2bf16 | v_mfma_f32_4x4x2bf16 | VOP3P | V MFMA F32 4X4X2BF16 |
| 0x00000000D3DF0000 | v_mfma_f32_4x4x4_16b_bf16 | v_mfma_f32_4x4x4_16b_bf16 | VOP3P | V MFMA F32 4X4X4 16B BF16 |
| 0x00000000D3CA0000 | v_mfma_f32_4x4x4_16b_f16 | v_mfma_f32_4x4x4_16b_f16 | VOP3P | V MFMA F32 4X4X4 16B F16 |
| 0x00000000D3E50000 | v_mfma_f32_4x4x4bf16_1k | v_mfma_f32_4x4x4bf16_1k | VOP3P | V MFMA F32 4X4X4BF16 1K |
| 0x00000000D3CA0000 | v_mfma_f32_4x4x4f16 | v_mfma_f32_4x4x4f16 | VOP3P | V MFMA F32 4X4X4F16 |
| 0x00000000D3EE0000 | v_mfma_f64_16x16x4_f64 | v_mfma_f64_16x16x4_f64 | VOP3P | V MFMA F64 16X16X4 F64 |
| 0x00000000D3EE0000 | v_mfma_f64_16x16x4f64 | v_mfma_f64_16x16x4f64 | VOP3P | V MFMA F64 16X16X4F64 |
| 0x00000000D3EF0000 | v_mfma_f64_4x4x4_4b_f64 | v_mfma_f64_4x4x4_4b_f64 | VOP3P | V MFMA F64 4X4X4 4B F64 |
| 0x00000000D3EF0000 | v_mfma_f64_4x4x4f64 | v_mfma_f64_4x4x4f64 | VOP3P | V MFMA F64 4X4X4F64 |
| 0x00000000D3D50000 | v_mfma_i32_16x16x16i8 | v_mfma_i32_16x16x16i8 | VOP3P | V MFMA I32 16X16X16I8 |
| 0x00000000D3D70000 | v_mfma_i32_16x16x32_i8 | v_mfma_i32_16x16x32_i8 | VOP3P | V MFMA I32 16X16X32 I8 |
| 0x00000000D3D10000 | v_mfma_i32_16x16x4_4b_i8 | v_mfma_i32_16x16x4_4b_i8 | VOP3P | V MFMA I32 16X16X4 4B I8 |
| 0x00000000D3D10000 | v_mfma_i32_16x16x4i8 | v_mfma_i32_16x16x4i8 | VOP3P | V MFMA I32 16X16X4I8 |
| 0x00000000D3B60000 | v_mfma_i32_16x16x64_i8 | v_mfma_i32_16x16x64_i8 | VOP3P | V MFMA I32 16X16X64 I8 |
| 0x00000000D3D60000 | v_mfma_i32_32x32x16_i8 | v_mfma_i32_32x32x16_i8 | VOP3P | V MFMA I32 32X32X16 I8 |
| 0x00000000D3B80000 | v_mfma_i32_32x32x32_i8 | v_mfma_i32_32x32x32_i8 | VOP3P | V MFMA I32 32X32X32 I8 |
| 0x00000000D3D00000 | v_mfma_i32_32x32x4_2b_i8 | v_mfma_i32_32x32x4_2b_i8 | VOP3P | V MFMA I32 32X32X4 2B I8 |
| 0x00000000D3D00000 | v_mfma_i32_32x32x4i8 | v_mfma_i32_32x32x4i8 | VOP3P | V MFMA I32 32X32X4I8 |
| 0x00000000D3D40000 | v_mfma_i32_32x32x8i8 | v_mfma_i32_32x32x8i8 | VOP3P | V MFMA I32 32X32X8I8 |
| 0x00000000D3D20000 | v_mfma_i32_4x4x4_16b_i8 | v_mfma_i32_4x4x4_16b_i8 | VOP3P | V MFMA I32 4X4X4 16B I8 |
| 0x00000000D3D20000 | v_mfma_i32_4x4x4i8 | v_mfma_i32_4x4x4i8 | VOP3P | V MFMA I32 4X4X4I8 |
| v_mfma_ld_scale_b32 | v_mfma_ld_scale_b32 | VOP3P | V MFMA LD SCALE B32 | |
| v_mfma_scale_f32_16x16x128_f8f6f4 | v_mfma_scale_f32_16x16x128_f8f6f4 | VOP3P | V MFMA SCALE F32 16X16X128 F8F6F4 | |
| v_mfma_scale_f32_32x32x64_f8f6f4 | v_mfma_scale_f32_32x32x64_f8f6f4 | VOP3P | V MFMA SCALE F32 32X32X64 F8F6F4 | |
| 0x00000000D1F40000 | v_min3_f16 | v_min3_f16 | VOP3 | V MIN3 F16 |
| 0x00000000D1D00000 | v_min3_f32 | v_min3_f32 | VOP3 | V MIN3 F32 |
| 0x00000000D64A0000 | v_min3_i16 | v_min3_i16 | VOP3 | V MIN3 I16 |
| 0x00000000D61A0000 | v_min3_i32 | v_min3_i32 | VOP3 | V MIN3 I32 |