Syntax of gfx908 Instructions¶
Introduction¶
This document describes the syntax of instructions specific to gfx908.
For a description of other gfx908 instructions, see Syntax of Core GFX9 Instructions.
Notation¶
Notation used in this document is explained here.
Overview¶
An overview of generic syntax and other features of AMDGPU instructions may be found in this document.
Instructions¶
FLAT¶
INSTRUCTION SRC0 SRC1 SRC2 MODIFIERS ——————————————————————————————————————————————————————————————————————————————— global_atomic_add_f32 vaddr, vdata, saddr offset13s slc global_atomic_pk_add_f16 vaddr, vdata, saddr offset13s slc
MUBUF¶
INSTRUCTION SRC0 SRC1 SRC2 SRC3 MODIFIERS ———————————————————————————————————————————————————————————————————————————————————————————————————— buffer_atomic_add_f32 vdata, vaddr, srsrc, soffset idxen offen offset12 slc buffer_atomic_pk_add_f16 vdata, vaddr, srsrc, soffset idxen offen offset12 slc
VOP2¶
INSTRUCTION DST SRC0 SRC1 MODIFIERS ———————————————————————————————————————————————————————————————————————————————————————————————————— v_dot2c_f32_f16 vdst, src0:f16x2, vsrc1:f16x2 v_dot2c_f32_f16_dpp vdst, vsrc0:m:f16x2, vsrc1:m:f16x2 dpp_ctrl row_mask bank_mask bound_ctrl v_dot2c_i32_i16 vdst, src0:i16x2, vsrc1:i16x2 v_dot2c_i32_i16_dpp vdst, vsrc0:i16x2, vsrc1:i16x2 dpp_ctrl row_mask bank_mask bound_ctrl v_dot4c_i32_i8 vdst, src0:i8x4, vsrc1:i8x4 v_dot4c_i32_i8_dpp vdst, vsrc0:i8x4, vsrc1:i8x4 dpp_ctrl row_mask bank_mask bound_ctrl v_dot8c_i32_i4 vdst, src0:i4x8, vsrc1:i4x8 v_dot8c_i32_i4_dpp vdst, vsrc0:i4x8, vsrc1:i4x8 dpp_ctrl row_mask bank_mask bound_ctrl v_fmac_f32 vdst, src0, vsrc1 v_fmac_f32_dpp vdst, vsrc0:m, vsrc1:m dpp_ctrl row_mask bank_mask bound_ctrl v_pk_fmac_f16 vdst, src0, vsrc1 v_xnor_b32 vdst, src0, vsrc1 v_xnor_b32_dpp vdst, vsrc0, vsrc1 dpp_ctrl row_mask bank_mask bound_ctrl v_xnor_b32_sdwa vdst, src0:m, src1:m dst_sel dst_unused src0_sel src1_sel
VOP3¶
INSTRUCTION DST SRC0 SRC1 MODIFIERS ——————————————————————————————————————————————————————————————————————————————————— v_dot2c_f32_f16_e64 vdst, src0:m:f16x2, src1:m:f16x2 clamp omod v_dot2c_i32_i16_e64 vdst, src0:i16x2, src1:i16x2 clamp v_dot4c_i32_i8_e64 vdst, src0:i8x4, src1:i8x4 clamp v_dot8c_i32_i4_e64 vdst, src0:i4x8, src1:i4x8 clamp v_fmac_f32_e64 vdst, src0:m, src1:m clamp omod v_xnor_b32_e64 vdst, src0, src1
VOP3P¶
INSTRUCTION DST SRC0 SRC1 SRC2 MODIFIERS ————————————————————————————————————————————————————————————————————————————————————————————————————————— v_accvgpr_read_b32 vdst, vsrc v_accvgpr_write_b32 vdst, src v_dot2_f32_f16 vdst, src0:f16x2, src1:f16x2, src2:f32 neg_lo neg_hi clamp v_dot2_i32_i16 vdst, src0:i16x2, src1:i16x2, src2:i32 clamp v_dot2_u32_u16 vdst, src0:u16x2, src1:u16x2, src2:u32 clamp v_dot4_i32_i8 vdst, src0:i8x4, src1:i8x4, src2:i32 clamp v_dot4_u32_u8 vdst, src0:u8x4, src1:u8x4, src2:u32 clamp v_dot8_i32_i4 vdst, src0:i4x8, src1:i4x8, src2:i32 clamp v_dot8_u32_u4 vdst, src0:u4x8, src1:u4x8, src2:u32 clamp v_fma_mix_f32 vdst, src0:m:fx, src1:m:fx, src2:m:fx m_op_sel m_op_sel_hi clamp v_fma_mixhi_f16 vdst, src0:m:fx, src1:m:fx, src2:m:fx m_op_sel m_op_sel_hi clamp v_fma_mixlo_f16 vdst, src0:m:fx, src1:m:fx, src2:m:fx m_op_sel m_op_sel_hi clamp v_mfma_f32_16x16x16f16 vdst:f32x4, vsrc0:f16x4, vsrc1:f16x4, vsrc2:f32x4 cbsz abid blgp v_mfma_f32_16x16x1f32 vdst:f32x16, vsrc0:f32, vsrc1:f32, vsrc2:f32x16 cbsz abid blgp v_mfma_f32_16x16x2bf16 vdst:f32x16, vsrc0:bf16x2, vsrc1:bf16x2, vsrc2:f32x16 cbsz abid blgp v_mfma_f32_16x16x4f16 vdst:f32x16, vsrc0:f16x4, vsrc1:f16x4, vsrc2:f32x16 cbsz abid blgp v_mfma_f32_16x16x4f32 vdst:f32x4, vsrc0:f32, vsrc1:f32, vsrc2:f32x4 cbsz abid blgp v_mfma_f32_16x16x8bf16 vdst:f32x4, vsrc0:bf16x2, vsrc1:bf16x2, vsrc2:f32x4 cbsz abid blgp v_mfma_f32_32x32x1f32 vdst:f32x32, vsrc0:f32, vsrc1:f32, vsrc2:f32x32 cbsz abid blgp v_mfma_f32_32x32x2bf16 vdst:f32x32, vsrc0:bf16x2, vsrc1:bf16x2, vsrc2:f32x32 cbsz abid blgp v_mfma_f32_32x32x2f32 vdst:f32x16, vsrc0:f32, vsrc1:f32, vsrc2:f32x16 cbsz abid blgp v_mfma_f32_32x32x4bf16 vdst:f32x16, vsrc0:bf16x2, vsrc1:bf16x2, vsrc2:f32x16 cbsz abid blgp v_mfma_f32_32x32x4f16 vdst:f32x32, vsrc0:f16x4, vsrc1:f16x4, vsrc2:f32x32 cbsz abid blgp v_mfma_f32_32x32x8f16 vdst:f32x16, vsrc0:f16x4, vsrc1:f16x4, vsrc2:f32x16 cbsz abid blgp v_mfma_f32_4x4x1f32 vdst:f32x4, vsrc0:f32, vsrc1:f32, vsrc2:f32x4 cbsz abid blgp v_mfma_f32_4x4x2bf16 vdst:f32x4, vsrc0:bf16x2, vsrc1:bf16x2, vsrc2:f32x4 cbsz abid blgp v_mfma_f32_4x4x4f16 vdst:f32x4, vsrc0:f16x4, vsrc1:f16x4, vsrc2:f32x4 cbsz abid blgp v_mfma_i32_16x16x16i8 vdst:i32x4, vsrc0:i8x4, vsrc1:i8x4, vsrc2:i32x4 cbsz abid blgp v_mfma_i32_16x16x4i8 vdst:i32x16, vsrc0:i8x4, vsrc1:i8x4, vsrc2:i32x16 cbsz abid blgp v_mfma_i32_32x32x4i8 vdst:i32x32, vsrc0:i8x4, vsrc1:i8x4, vsrc2:i32x32 cbsz abid blgp v_mfma_i32_32x32x8i8 vdst:i32x16, vsrc0:i8x4, vsrc1:i8x4, vsrc2:i32x16 cbsz abid blgp v_mfma_i32_4x4x4i8 vdst:i32x4, vsrc0:i8x4, vsrc1:i8x4, vsrc2:i32x4 cbsz abid blgp