Home / AMDGPU / Table

AMDGPU Opcode Map

Reverse lookup instructions by their hexadecimal encoding.

← Back to List
Opcode Mnemonic Syntax Format Name
buffer_load_dword buffer_load_dword VDST, VADDR, SRSRC, offset MUBUF BUFFER LOAD DWORD
ds_add_u32 ds_add_u32 ADDR, DATA, offset DS DS ADD U32
ds_permute_b32 ds_permute_b32 VDST, ADDR, DATA DS DS PERMUTE B32
ds_read_b32 ds_read_b32 VDST, ADDR, offset DS DS READ B32
ds_write_b32 ds_write_b32 ADDR, DATA, offset DS DS WRITE B32
flat_load_dword flat_load_dword VDST, VADDR FLAT FLAT LOAD DWORD
global_atomic_add global_atomic_add VDST, VADDR, VDATA, SADDR GLOBAL GLOBAL ATOMIC ADD
global_load_dword global_load_dword VDST, VADDR, SADDR, offset GLOBAL GLOBAL LOAD DWORD
global_store_dword global_store_dword VADDR, VDATA, SADDR, offset GLOBAL GLOBAL STORE DWORD
s_add_u32 s_add_u32 SDST, S0, S1 SOP2 S ADD U32
s_barrier s_barrier SOPP S BARRIER
s_branch s_branch offset SOPP S BRANCH
s_cbranch_scc1 s_cbranch_scc1 offset SOPP S CBRANCH SCC1
s_cmp_eq_i32 s_cmp_eq_i32 S0, S1 SOPC S CMP EQ I32
s_load_dword s_load_dword SDST, SBASE, offset SMEM S LOAD DWORD
s_mul_i32 s_mul_i32 SDST, S0, S1 SOP2 S MUL I32
s_sub_u32 s_sub_u32 SDST, S0, S1 SOP2 S SUB U32
v_add_f32 v_add_f32 VDST, S0, S1 VOP2 V ADD F32
v_add_u32 v_add_u32 VDST, S0, S1 VOP2 V ADD U32
v_bcnt_u32_b32 v_bcnt_u32_b32 VDST, S0, S1 VOP3 V BCNT U32 B32
v_cmp_eq_f32 v_cmp_eq_f32 VCC, S0, S1 VOPC V CMP EQ F32
v_cvt_f32_i32 v_cvt_f32_i32 VDST, S0 VOP1 V CVT F32 I32
v_fma_f32 v_fma_f32 VDST, S0, S1, S2 VOP3 V FMA F32
v_mfma_f32_16x16x16f16 v_mfma_f32_16x16x16f16 D, A, B, C VOP3P-MAI V MFMA F32 16X16X16F16
v_mul_f32 v_mul_f32 VDST, S0, S1 VOP2 V MUL F32
v_mul_lo_u32 v_mul_lo_u32 VDST, S0, S1 VOP3 V MUL LO U32
v_readfirstlane_b32 v_readfirstlane_b32 SDST, S0 VOP1 V READFIRSTLANE B32
v_rsq_f32 v_rsq_f32 VDST, S0 VOP1 V RSQ F32
v_sin_f32 v_sin_f32 VDST, S0 VOP1 V SIN F32
v_sub_f32 v_sub_f32 VDST, S0, S1 VOP2 V SUB F32