diff options
62 files changed, 1464 insertions, 109 deletions
diff --git a/llvm/docs/AMDGPU/AMDGPUAsmGFX10.rst b/llvm/docs/AMDGPU/AMDGPUAsmGFX10.rst index 033a1635567..bc2b4f74ecf 100644 --- a/llvm/docs/AMDGPU/AMDGPUAsmGFX10.rst +++ b/llvm/docs/AMDGPU/AMDGPUAsmGFX10.rst @@ -5,20 +5,25 @@ * * ************************************************** -============================ -Syntax of GFX10 Instructions -============================ +==================================================================================== +Syntax of Core GFX10 Instructions +==================================================================================== .. contents:: :local: +Introduction +============ + +This document describes the syntax of *core* GFX10 instructions. + Notation ======== Notation used in this document is explained :ref:`here<amdgpu_syn_instruction_notation>`. -Introduction -============ +Overvew +======= An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document<amdgpu_syn_instructions>`. @@ -1359,56 +1364,56 @@ VOP2 .. parsed-literal:: - **INSTRUCTION** **DST0** **DST1** **SRC0** **SRC1** **SRC2** - \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| - v_add_co_ci_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>` - v_add_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_add_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_add_nc_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_and_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_ashrrev_i32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`::ref:`u32<amdgpu_synid10_type_dev>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_cndmask_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>` - v_cvt_pkrtz_f16_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_fmaak_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`imm32<amdgpu_synid10_fimm16>` - v_fmaak_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`imm32<amdgpu_synid10_fimm32>` - v_fmac_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_fmac_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_fmamk_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`imm32<amdgpu_synid10_fimm16>`, :ref:`vsrc2<amdgpu_synid10_vsrc32_0>` - v_fmamk_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`imm32<amdgpu_synid10_fimm32>`, :ref:`vsrc2<amdgpu_synid10_vsrc32_0>` - v_ldexp_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`::ref:`i16<amdgpu_synid10_type_dev>` - v_lshlrev_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`::ref:`u32<amdgpu_synid10_type_dev>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_lshrrev_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`::ref:`u32<amdgpu_synid10_type_dev>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_mac_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_mac_legacy_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_madak_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`imm32<amdgpu_synid10_fimm32>` - v_madmk_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`imm32<amdgpu_synid10_fimm32>`, :ref:`vsrc2<amdgpu_synid10_vsrc32_0>` - v_max_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_max_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_max_i32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_max_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_min_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_min_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_min_i32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_min_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_mul_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_mul_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_mul_hi_i32_i24 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_mul_hi_u32_u24 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_mul_i32_i24 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_mul_legacy_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_mul_u32_u24 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_or_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_pk_fmac_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_sub_co_ci_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>` - v_sub_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_sub_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_sub_nc_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_subrev_co_ci_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>`, :ref:`src0<amdgpu_synid10_src32_2>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>` - v_subrev_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_subrev_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_subrev_nc_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_xnor_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` - v_xor_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + **INSTRUCTION** **DST0** **DST1** **SRC0** **SRC1** **SRC2** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + v_add_co_ci_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>` + v_add_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_add_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_add_nc_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_and_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_ashrrev_i32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`::ref:`u32<amdgpu_synid10_type_dev>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_cndmask_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>` + v_cvt_pkrtz_f16_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_fmaak_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`imm32<amdgpu_synid10_fimm16>` + v_fmaak_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`imm32<amdgpu_synid10_fimm32>` + v_fmac_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_fmac_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_fmamk_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`imm32<amdgpu_synid10_fimm16>`, :ref:`vsrc2<amdgpu_synid10_vsrc32_0>` + v_fmamk_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`imm32<amdgpu_synid10_fimm32>`, :ref:`vsrc2<amdgpu_synid10_vsrc32_0>` + v_ldexp_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`::ref:`i16<amdgpu_synid10_type_dev>` + v_lshlrev_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`::ref:`u32<amdgpu_synid10_type_dev>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_lshrrev_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`::ref:`u32<amdgpu_synid10_type_dev>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_mac_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_mac_legacy_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_madak_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`imm32<amdgpu_synid10_fimm32>` + v_madmk_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`imm32<amdgpu_synid10_fimm32>`, :ref:`vsrc2<amdgpu_synid10_vsrc32_0>` + v_max_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_max_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_max_i32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_max_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_min_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_min_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_min_i32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_min_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_mul_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_mul_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_mul_hi_i32_i24 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_mul_hi_u32_u24 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_mul_i32_i24 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_mul_legacy_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_mul_u32_u24 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_or_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_pk_fmac_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`::ref:`f16x2<amdgpu_synid10_type_dev>`, :ref:`src0<amdgpu_synid10_src32_1>`::ref:`f16x2<amdgpu_synid10_type_dev>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`::ref:`f16x2<amdgpu_synid10_type_dev>` + v_sub_co_ci_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>` + v_sub_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_sub_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_sub_nc_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_subrev_co_ci_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>`, :ref:`src0<amdgpu_synid10_src32_2>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>` + v_subrev_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_subrev_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_subrev_nc_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_xnor_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` + v_xor_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>` VOP3 ----------------------- @@ -1847,30 +1852,30 @@ VOP3P .. parsed-literal:: - **INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS** - \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| - v_fma_mix_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_3>`::ref:`fx<amdgpu_synid10_mad_type_dev>`, :ref:`src1<amdgpu_synid10_src32_0>`::ref:`fx<amdgpu_synid10_mad_type_dev>`, :ref:`src2<amdgpu_synid10_src32_0>`::ref:`fx<amdgpu_synid10_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` - v_fma_mixhi_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_3>`::ref:`fx<amdgpu_synid10_mad_type_dev>`, :ref:`src1<amdgpu_synid10_src32_0>`::ref:`fx<amdgpu_synid10_mad_type_dev>`, :ref:`src2<amdgpu_synid10_src32_0>`::ref:`fx<amdgpu_synid10_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` - v_fma_mixlo_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_3>`::ref:`fx<amdgpu_synid10_mad_type_dev>`, :ref:`src1<amdgpu_synid10_src32_0>`::ref:`fx<amdgpu_synid10_mad_type_dev>`, :ref:`src2<amdgpu_synid10_src32_0>`::ref:`fx<amdgpu_synid10_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` - v_pk_add_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>` - v_pk_add_i16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` - v_pk_add_u16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` - v_pk_ashrrev_i16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`::ref:`u16x2<amdgpu_synid10_type_dev>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` - v_pk_fma_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>`, :ref:`src2<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>` - v_pk_lshlrev_b16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`::ref:`u16x2<amdgpu_synid10_type_dev>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` - v_pk_lshrrev_b16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`::ref:`u16x2<amdgpu_synid10_type_dev>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` - v_pk_mad_i16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>`, :ref:`src2<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` - v_pk_mad_u16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>`, :ref:`src2<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` - v_pk_max_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>` - v_pk_max_i16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` - v_pk_max_u16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` - v_pk_min_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>` - v_pk_min_i16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` - v_pk_min_u16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` - v_pk_mul_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>` - v_pk_mul_lo_u16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` - v_pk_sub_i16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` - v_pk_sub_u16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` + **INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + v_fma_mix_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_3>`::ref:`m<amdgpu_synid10_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid10_mad_type_dev>`, :ref:`src1<amdgpu_synid10_src32_0>`::ref:`m<amdgpu_synid10_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid10_mad_type_dev>`, :ref:`src2<amdgpu_synid10_src32_0>`::ref:`m<amdgpu_synid10_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid10_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` + v_fma_mixhi_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_3>`::ref:`m<amdgpu_synid10_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid10_mad_type_dev>`, :ref:`src1<amdgpu_synid10_src32_0>`::ref:`m<amdgpu_synid10_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid10_mad_type_dev>`, :ref:`src2<amdgpu_synid10_src32_0>`::ref:`m<amdgpu_synid10_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid10_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` + v_fma_mixlo_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_3>`::ref:`m<amdgpu_synid10_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid10_mad_type_dev>`, :ref:`src1<amdgpu_synid10_src32_0>`::ref:`m<amdgpu_synid10_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid10_mad_type_dev>`, :ref:`src2<amdgpu_synid10_src32_0>`::ref:`m<amdgpu_synid10_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid10_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` + v_pk_add_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>` + v_pk_add_i16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` + v_pk_add_u16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` + v_pk_ashrrev_i16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`::ref:`u16x2<amdgpu_synid10_type_dev>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` + v_pk_fma_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>`, :ref:`src2<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>` + v_pk_lshlrev_b16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`::ref:`u16x2<amdgpu_synid10_type_dev>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` + v_pk_lshrrev_b16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`::ref:`u16x2<amdgpu_synid10_type_dev>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` + v_pk_mad_i16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>`, :ref:`src2<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` + v_pk_mad_u16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>`, :ref:`src2<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` + v_pk_max_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>` + v_pk_max_i16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` + v_pk_max_u16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` + v_pk_min_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>` + v_pk_min_i16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` + v_pk_min_u16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` + v_pk_mul_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>` + v_pk_mul_lo_u16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` + v_pk_sub_i16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` + v_pk_sub_u16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` VOPC ----------------------- diff --git a/llvm/docs/AMDGPU/AMDGPUAsmGFX7.rst b/llvm/docs/AMDGPU/AMDGPUAsmGFX7.rst index 1933aff3bb7..e96862bcf03 100644 --- a/llvm/docs/AMDGPU/AMDGPUAsmGFX7.rst +++ b/llvm/docs/AMDGPU/AMDGPUAsmGFX7.rst @@ -5,20 +5,25 @@ * * ************************************************** -============================ -Syntax of GFX7 Instructions -============================ +==================================================================================== +Syntax of Core GFX7 Instructions +==================================================================================== .. contents:: :local: +Introduction +============ + +This document describes the syntax of *core* GFX7 instructions. + Notation ======== Notation used in this document is explained :ref:`here<amdgpu_syn_instruction_notation>`. -Introduction -============ +Overvew +======= An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document<amdgpu_syn_instructions>`. diff --git a/llvm/docs/AMDGPU/AMDGPUAsmGFX8.rst b/llvm/docs/AMDGPU/AMDGPUAsmGFX8.rst index a0e514e29b2..4dd0439c6fd 100644 --- a/llvm/docs/AMDGPU/AMDGPUAsmGFX8.rst +++ b/llvm/docs/AMDGPU/AMDGPUAsmGFX8.rst @@ -5,20 +5,25 @@ * * ************************************************** -============================ -Syntax of GFX8 Instructions -============================ +==================================================================================== +Syntax of Core GFX8 Instructions +==================================================================================== .. contents:: :local: +Introduction +============ + +This document describes the syntax of *core* GFX8 instructions. + Notation ======== Notation used in this document is explained :ref:`here<amdgpu_syn_instruction_notation>`. -Introduction -============ +Overvew +======= An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document<amdgpu_syn_instructions>`. diff --git a/llvm/docs/AMDGPU/AMDGPUAsmGFX9.rst b/llvm/docs/AMDGPU/AMDGPUAsmGFX9.rst index 8ce056c8caf..7131718df7e 100644 --- a/llvm/docs/AMDGPU/AMDGPUAsmGFX9.rst +++ b/llvm/docs/AMDGPU/AMDGPUAsmGFX9.rst @@ -5,20 +5,25 @@ * * ************************************************** -============================ -Syntax of GFX9 Instructions -============================ +==================================================================================== +Syntax of Core GFX9 Instructions +==================================================================================== .. contents:: :local: +Introduction +============ + +This document describes the syntax of *core* GFX9 instructions. + Notation ======== Notation used in this document is explained :ref:`here<amdgpu_syn_instruction_notation>`. -Introduction -============ +Overvew +======= An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document<amdgpu_syn_instructions>`. @@ -1637,9 +1642,6 @@ VOP3P **INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS** \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| - v_mad_mix_f32 :ref:`vdst<amdgpu_synid9_vdst32_0>`, :ref:`src0<amdgpu_synid9_src32_2>`::ref:`fx<amdgpu_synid9_mad_type_dev>`, :ref:`src1<amdgpu_synid9_src32_3>`::ref:`fx<amdgpu_synid9_mad_type_dev>`, :ref:`src2<amdgpu_synid9_src32_3>`::ref:`fx<amdgpu_synid9_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` - v_mad_mixhi_f16 :ref:`vdst<amdgpu_synid9_vdst32_0>`, :ref:`src0<amdgpu_synid9_src32_2>`::ref:`fx<amdgpu_synid9_mad_type_dev>`, :ref:`src1<amdgpu_synid9_src32_3>`::ref:`fx<amdgpu_synid9_mad_type_dev>`, :ref:`src2<amdgpu_synid9_src32_3>`::ref:`fx<amdgpu_synid9_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` - v_mad_mixlo_f16 :ref:`vdst<amdgpu_synid9_vdst32_0>`, :ref:`src0<amdgpu_synid9_src32_2>`::ref:`fx<amdgpu_synid9_mad_type_dev>`, :ref:`src1<amdgpu_synid9_src32_3>`::ref:`fx<amdgpu_synid9_mad_type_dev>`, :ref:`src2<amdgpu_synid9_src32_3>`::ref:`fx<amdgpu_synid9_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` v_pk_add_f16 :ref:`vdst<amdgpu_synid9_vdst32_0>`, :ref:`src0<amdgpu_synid9_src32_2>`, :ref:`src1<amdgpu_synid9_src32_3>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>` v_pk_add_i16 :ref:`vdst<amdgpu_synid9_vdst32_0>`, :ref:`src0<amdgpu_synid9_src32_2>`, :ref:`src1<amdgpu_synid9_src32_3>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` v_pk_add_u16 :ref:`vdst<amdgpu_synid9_vdst32_0>`, :ref:`src0<amdgpu_synid9_src32_2>`, :ref:`src1<amdgpu_synid9_src32_3>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` @@ -2096,7 +2098,6 @@ VOPC gfx9_vsrc32_0 gfx9_vsrc32_1 gfx9_vsrc64_0 - gfx9_mad_type_dev gfx9_mod_dpp_sdwa_abs_neg gfx9_mod_sdwa_sext gfx9_mod_vop3_abs_neg diff --git a/llvm/docs/AMDGPU/AMDGPUAsmGFX900.rst b/llvm/docs/AMDGPU/AMDGPUAsmGFX900.rst new file mode 100644 index 00000000000..7832e08dbb3 --- /dev/null +++ b/llvm/docs/AMDGPU/AMDGPUAsmGFX900.rst @@ -0,0 +1,58 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +==================================================================================== +Syntax of gfx900, gfx902 and gfx909 Instructions +==================================================================================== + +.. contents:: + :local: + +Introduction +============ + +This document describes the syntax of *instructions specific to gfx900, gfx902 and gfx909*. + +For a description of other gfx900, gfx902 and gfx909 instructions see :doc:`Syntax of Core GFX9 Instructions<AMDGPUAsmGFX9>`. + +Notation +======== + +Notation used in this document is explained :ref:`here<amdgpu_syn_instruction_notation>`. + +Overvew +======= + +An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document<amdgpu_syn_instructions>`. + +Instructions +============ + + +VOP3P +----------------------- + +.. parsed-literal:: + + **INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + v_mad_mix_f32 :ref:`vdst<amdgpu_synid900_vdst32_0>`, :ref:`src0<amdgpu_synid900_src32_0>`::ref:`m<amdgpu_synid900_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid900_mad_type_dev>`, :ref:`src1<amdgpu_synid900_src32_1>`::ref:`m<amdgpu_synid900_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid900_mad_type_dev>`, :ref:`src2<amdgpu_synid900_src32_1>`::ref:`m<amdgpu_synid900_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid900_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` + v_mad_mixhi_f16 :ref:`vdst<amdgpu_synid900_vdst32_0>`, :ref:`src0<amdgpu_synid900_src32_0>`::ref:`m<amdgpu_synid900_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid900_mad_type_dev>`, :ref:`src1<amdgpu_synid900_src32_1>`::ref:`m<amdgpu_synid900_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid900_mad_type_dev>`, :ref:`src2<amdgpu_synid900_src32_1>`::ref:`m<amdgpu_synid900_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid900_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` + v_mad_mixlo_f16 :ref:`vdst<amdgpu_synid900_vdst32_0>`, :ref:`src0<amdgpu_synid900_src32_0>`::ref:`m<amdgpu_synid900_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid900_mad_type_dev>`, :ref:`src1<amdgpu_synid900_src32_1>`::ref:`m<amdgpu_synid900_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid900_mad_type_dev>`, :ref:`src2<amdgpu_synid900_src32_1>`::ref:`m<amdgpu_synid900_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid900_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` + +.. |---| unicode:: U+02014 .. em dash + + +.. toctree:: + :hidden: + + AMDGPUAsmGFX9 + gfx900_src32_0 + gfx900_src32_1 + gfx900_vdst32_0 + gfx900_mad_type_dev + gfx900_mod_vop3_abs_neg diff --git a/llvm/docs/AMDGPU/AMDGPUAsmGFX904.rst b/llvm/docs/AMDGPU/AMDGPUAsmGFX904.rst new file mode 100644 index 00000000000..23f79fc96a4 --- /dev/null +++ b/llvm/docs/AMDGPU/AMDGPUAsmGFX904.rst @@ -0,0 +1,58 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +==================================================================================== +Syntax of gfx904 Instructions +==================================================================================== + +.. contents:: + :local: + +Introduction +============ + +This document describes the syntax of *instructions specific to gfx904*. + +For a description of other gfx904 instructions see :doc:`Syntax of Core GFX9 Instructions<AMDGPUAsmGFX9>`. + +Notation +======== + +Notation used in this document is explained :ref:`here<amdgpu_syn_instruction_notation>`. + +Overvew +======= + +An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document<amdgpu_syn_instructions>`. + +Instructions +============ + + +VOP3P +----------------------- + +.. parsed-literal:: + + **INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + v_fma_mix_f32 :ref:`vdst<amdgpu_synid904_vdst32_0>`, :ref:`src0<amdgpu_synid904_src32_0>`::ref:`m<amdgpu_synid904_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid904_mad_type_dev>`, :ref:`src1<amdgpu_synid904_src32_1>`::ref:`m<amdgpu_synid904_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid904_mad_type_dev>`, :ref:`src2<amdgpu_synid904_src32_1>`::ref:`m<amdgpu_synid904_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid904_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` + v_fma_mixhi_f16 :ref:`vdst<amdgpu_synid904_vdst32_0>`, :ref:`src0<amdgpu_synid904_src32_0>`::ref:`m<amdgpu_synid904_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid904_mad_type_dev>`, :ref:`src1<amdgpu_synid904_src32_1>`::ref:`m<amdgpu_synid904_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid904_mad_type_dev>`, :ref:`src2<amdgpu_synid904_src32_1>`::ref:`m<amdgpu_synid904_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid904_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` + v_fma_mixlo_f16 :ref:`vdst<amdgpu_synid904_vdst32_0>`, :ref:`src0<amdgpu_synid904_src32_0>`::ref:`m<amdgpu_synid904_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid904_mad_type_dev>`, :ref:`src1<amdgpu_synid904_src32_1>`::ref:`m<amdgpu_synid904_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid904_mad_type_dev>`, :ref:`src2<amdgpu_synid904_src32_1>`::ref:`m<amdgpu_synid904_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid904_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` + +.. |---| unicode:: U+02014 .. em dash + + +.. toctree:: + :hidden: + + AMDGPUAsmGFX9 + gfx904_src32_0 + gfx904_src32_1 + gfx904_vdst32_0 + gfx904_mad_type_dev + gfx904_mod_vop3_abs_neg diff --git a/llvm/docs/AMDGPU/AMDGPUAsmGFX906.rst b/llvm/docs/AMDGPU/AMDGPUAsmGFX906.rst new file mode 100644 index 00000000000..8119af4392e --- /dev/null +++ b/llvm/docs/AMDGPU/AMDGPUAsmGFX906.rst @@ -0,0 +1,93 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +==================================================================================== +Syntax of gfx906 Instructions +==================================================================================== + +.. contents:: + :local: + +Introduction +============ + +This document describes the syntax of *instructions specific to gfx906*. + +For a description of other gfx906 instructions see :doc:`Syntax of Core GFX9 Instructions<AMDGPUAsmGFX9>`. + +Notation +======== + +Notation used in this document is explained :ref:`here<amdgpu_syn_instruction_notation>`. + +Overvew +======= + +An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document<amdgpu_syn_instructions>`. + +Instructions +============ + + +VOP2 +----------------------- + +.. parsed-literal:: + + **INSTRUCTION** **DST** **SRC0** **SRC1** **MODIFIERS** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + v_fmac_f32 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_0>`, :ref:`vsrc1<amdgpu_synid906_vsrc32_0>` + v_fmac_f32_dpp :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`vsrc0<amdgpu_synid906_vsrc32_0>`::ref:`m<amdgpu_synid906_mod_dpp_sdwa_abs_neg>`, :ref:`vsrc1<amdgpu_synid906_vsrc32_0>`::ref:`m<amdgpu_synid906_mod_dpp_sdwa_abs_neg>` :ref:`dpp_ctrl<amdgpu_synid_dpp_ctrl>` :ref:`row_mask<amdgpu_synid_row_mask>` :ref:`bank_mask<amdgpu_synid_bank_mask>` :ref:`bound_ctrl<amdgpu_synid_bound_ctrl>` + v_xnor_b32 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_0>`, :ref:`vsrc1<amdgpu_synid906_vsrc32_0>` + v_xnor_b32_dpp :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`vsrc0<amdgpu_synid906_vsrc32_0>`, :ref:`vsrc1<amdgpu_synid906_vsrc32_0>` :ref:`dpp_ctrl<amdgpu_synid_dpp_ctrl>` :ref:`row_mask<amdgpu_synid_row_mask>` :ref:`bank_mask<amdgpu_synid_bank_mask>` :ref:`bound_ctrl<amdgpu_synid_bound_ctrl>` + v_xnor_b32_sdwa :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_0>`::ref:`m<amdgpu_synid906_mod_sdwa_sext>`, :ref:`vsrc1<amdgpu_synid906_vsrc32_0>`::ref:`m<amdgpu_synid906_mod_sdwa_sext>` :ref:`dst_sel<amdgpu_synid_dst_sel>` :ref:`dst_unused<amdgpu_synid_dst_unused>` :ref:`src0_sel<amdgpu_synid_src0_sel>` :ref:`src1_sel<amdgpu_synid_src1_sel>` + +VOP3 +----------------------- + +.. parsed-literal:: + + **INSTRUCTION** **DST** **SRC0** **SRC1** **MODIFIERS** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + v_fmac_f32_e64 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_1>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`, :ref:`src1<amdgpu_synid906_src32_2>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>` :ref:`clamp<amdgpu_synid_clamp>` :ref:`omod<amdgpu_synid_omod>` + v_xnor_b32_e64 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_1>`, :ref:`src1<amdgpu_synid906_src32_2>` + +VOP3P +----------------------- + +.. parsed-literal:: + + **INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + v_dot2_f32_f16 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_1>`::ref:`f16x2<amdgpu_synid906_type_dev>`, :ref:`src1<amdgpu_synid906_src32_2>`::ref:`f16x2<amdgpu_synid906_type_dev>`, :ref:`src2<amdgpu_synid906_src32_2>`::ref:`f32<amdgpu_synid906_type_dev>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>` + v_dot2_i32_i16 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_1>`::ref:`i16x2<amdgpu_synid906_type_dev>`, :ref:`src1<amdgpu_synid906_src32_2>`::ref:`i16x2<amdgpu_synid906_type_dev>`, :ref:`src2<amdgpu_synid906_src32_2>`::ref:`i32<amdgpu_synid906_type_dev>` :ref:`clamp<amdgpu_synid_clamp>` + v_dot2_u32_u16 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_1>`::ref:`u16x2<amdgpu_synid906_type_dev>`, :ref:`src1<amdgpu_synid906_src32_2>`::ref:`u16x2<amdgpu_synid906_type_dev>`, :ref:`src2<amdgpu_synid906_src32_2>`::ref:`u32<amdgpu_synid906_type_dev>` :ref:`clamp<amdgpu_synid_clamp>` + v_dot4_i32_i8 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_1>`::ref:`i8x2<amdgpu_synid906_type_dev>`, :ref:`src1<amdgpu_synid906_src32_2>`::ref:`i8x2<amdgpu_synid906_type_dev>`, :ref:`src2<amdgpu_synid906_src32_2>`::ref:`i32<amdgpu_synid906_type_dev>` :ref:`clamp<amdgpu_synid_clamp>` + v_dot4_u32_u8 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_1>`::ref:`u8x2<amdgpu_synid906_type_dev>`, :ref:`src1<amdgpu_synid906_src32_2>`::ref:`u8x2<amdgpu_synid906_type_dev>`, :ref:`src2<amdgpu_synid906_src32_2>`::ref:`u32<amdgpu_synid906_type_dev>` :ref:`clamp<amdgpu_synid_clamp>` + v_dot8_i32_i4 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_1>`::ref:`i4x2<amdgpu_synid906_type_dev>`, :ref:`src1<amdgpu_synid906_src32_2>`::ref:`i4x2<amdgpu_synid906_type_dev>`, :ref:`src2<amdgpu_synid906_src32_2>`::ref:`i32<amdgpu_synid906_type_dev>` :ref:`clamp<amdgpu_synid_clamp>` + v_dot8_u32_u4 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_1>`::ref:`u4x2<amdgpu_synid906_type_dev>`, :ref:`src1<amdgpu_synid906_src32_2>`::ref:`u4x2<amdgpu_synid906_type_dev>`, :ref:`src2<amdgpu_synid906_src32_2>`::ref:`u32<amdgpu_synid906_type_dev>` :ref:`clamp<amdgpu_synid_clamp>` + v_fma_mix_f32 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_1>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>`, :ref:`src1<amdgpu_synid906_src32_2>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>`, :ref:`src2<amdgpu_synid906_src32_2>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` + v_fma_mixhi_f16 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_1>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>`, :ref:`src1<amdgpu_synid906_src32_2>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>`, :ref:`src2<amdgpu_synid906_src32_2>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` + v_fma_mixlo_f16 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_1>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>`, :ref:`src1<amdgpu_synid906_src32_2>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>`, :ref:`src2<amdgpu_synid906_src32_2>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` + +.. |---| unicode:: U+02014 .. em dash + + +.. toctree:: + :hidden: + + AMDGPUAsmGFX9 + gfx906_src32_0 + gfx906_src32_1 + gfx906_src32_2 + gfx906_vdst32_0 + gfx906_vsrc32_0 + gfx906_mad_type_dev + gfx906_mod_dpp_sdwa_abs_neg + gfx906_mod_sdwa_sext + gfx906_mod_vop3_abs_neg + gfx906_type_dev diff --git a/llvm/docs/AMDGPU/AMDGPUAsmGFX908.rst b/llvm/docs/AMDGPU/AMDGPUAsmGFX908.rst new file mode 100644 index 00000000000..7e90837cdf0 --- /dev/null +++ b/llvm/docs/AMDGPU/AMDGPUAsmGFX908.rst @@ -0,0 +1,165 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +==================================================================================== +Syntax of gfx908 Instructions +==================================================================================== + +.. contents:: + :local: + +Introduction +============ + +This document describes the syntax of *instructions specific to gfx908*. + +For a description of other gfx908 instructions see :doc:`Syntax of Core GFX9 Instructions<AMDGPUAsmGFX9>`. + +Notation +======== + +Notation used in this document is explained :ref:`here<amdgpu_syn_instruction_notation>`. + +Overvew +======= + +An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document<amdgpu_syn_instructions>`. + +Instructions +============ + + +FLAT +----------------------- + +.. parsed-literal:: + + **INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + global_atomic_add_f32 :ref:`vdst<amdgpu_synid908_dst_flat_atomic32>`::ref:`opt<amdgpu_synid908_opt>`, :ref:`vaddr<amdgpu_synid908_vaddr_flat_global>`, :ref:`vdata<amdgpu_synid908_vdata32_0>`, :ref:`saddr<amdgpu_synid908_saddr_flat_global>` :ref:`offset13s<amdgpu_synid_flat_offset13s>` + global_atomic_pk_add_f16 :ref:`vdst<amdgpu_synid908_dst_flat_atomic32>`::ref:`opt<amdgpu_synid908_opt>`::ref:`f16x2<amdgpu_synid908_type_dev>`, :ref:`vaddr<amdgpu_synid908_vaddr_flat_global>`, :ref:`vdata<amdgpu_synid908_vdata32_0>`::ref:`f16x2<amdgpu_synid908_type_dev>`, :ref:`saddr<amdgpu_synid908_saddr_flat_global>` :ref:`offset13s<amdgpu_synid_flat_offset13s>` + +MUBUF +----------------------- + +.. parsed-literal:: + + **INSTRUCTION** **SRC0** **SRC1** **SRC2** **SRC3** **MODIFIERS** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + buffer_atomic_add_f32 :ref:`vdata<amdgpu_synid908_data_buf_atomic32>`::ref:`dst<amdgpu_synid908_ret>`, :ref:`vaddr<amdgpu_synid908_addr_buf>`, :ref:`srsrc<amdgpu_synid908_rsrc_buf>`, :ref:`soffset<amdgpu_synid908_offset_buf>` :ref:`idxen<amdgpu_synid_idxen>` :ref:`offen<amdgpu_synid_offen>` :ref:`offset12<amdgpu_synid_buf_offset12>` :ref:`slc<amdgpu_synid_slc>` + buffer_atomic_pk_add_f16 :ref:`vdata<amdgpu_synid908_data_buf_atomic32>`::ref:`dst<amdgpu_synid908_ret>`::ref:`f16x2<amdgpu_synid908_type_dev>`, :ref:`vaddr<amdgpu_synid908_addr_buf>`, :ref:`srsrc<amdgpu_synid908_rsrc_buf>`, :ref:`soffset<amdgpu_synid908_offset_buf>` :ref:`idxen<amdgpu_synid_idxen>` :ref:`offen<amdgpu_synid_offen>` :ref:`offset12<amdgpu_synid_buf_offset12>` :ref:`slc<amdgpu_synid_slc>` + +VOP2 +----------------------- + +.. parsed-literal:: + + **INSTRUCTION** **DST** **SRC0** **SRC1** **MODIFIERS** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + v_dot2c_f32_f16 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_0>`::ref:`f16x2<amdgpu_synid908_type_dev>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>`::ref:`f16x2<amdgpu_synid908_type_dev>` + v_dot2c_f32_f16_dpp :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`vsrc0<amdgpu_synid908_vsrc32_0>`::ref:`f16x2<amdgpu_synid908_type_dev>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>`::ref:`f16x2<amdgpu_synid908_type_dev>` :ref:`dpp_ctrl<amdgpu_synid_dpp_ctrl>` :ref:`row_mask<amdgpu_synid_row_mask>` :ref:`bank_mask<amdgpu_synid_bank_mask>` :ref:`bound_ctrl<amdgpu_synid_bound_ctrl>` + v_dot2c_i32_i16 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>` + v_dot2c_i32_i16_dpp :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`vsrc0<amdgpu_synid908_vsrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>` :ref:`dpp_ctrl<amdgpu_synid_dpp_ctrl>` :ref:`row_mask<amdgpu_synid_row_mask>` :ref:`bank_mask<amdgpu_synid_bank_mask>` :ref:`bound_ctrl<amdgpu_synid_bound_ctrl>` + v_dot4c_i32_i8 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_0>`::ref:`i8x4<amdgpu_synid908_type_dev>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>`::ref:`i8x4<amdgpu_synid908_type_dev>` + v_dot4c_i32_i8_dpp :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`vsrc0<amdgpu_synid908_vsrc32_0>`::ref:`i8x4<amdgpu_synid908_type_dev>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>`::ref:`i8x4<amdgpu_synid908_type_dev>` :ref:`dpp_ctrl<amdgpu_synid_dpp_ctrl>` :ref:`row_mask<amdgpu_synid_row_mask>` :ref:`bank_mask<amdgpu_synid_bank_mask>` :ref:`bound_ctrl<amdgpu_synid_bound_ctrl>` + v_dot8c_i32_i4 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_0>`::ref:`i4x8<amdgpu_synid908_type_dev>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>`::ref:`i4x8<amdgpu_synid908_type_dev>` + v_dot8c_i32_i4_dpp :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`vsrc0<amdgpu_synid908_vsrc32_0>`::ref:`i4x8<amdgpu_synid908_type_dev>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>`::ref:`i4x8<amdgpu_synid908_type_dev>` :ref:`dpp_ctrl<amdgpu_synid_dpp_ctrl>` :ref:`row_mask<amdgpu_synid_row_mask>` :ref:`bank_mask<amdgpu_synid_bank_mask>` :ref:`bound_ctrl<amdgpu_synid_bound_ctrl>` + v_fmac_f32 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_0>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>` + v_fmac_f32_dpp :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`vsrc0<amdgpu_synid908_vsrc32_0>`::ref:`m<amdgpu_synid908_mod_dpp_sdwa_abs_neg>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>`::ref:`m<amdgpu_synid908_mod_dpp_sdwa_abs_neg>` :ref:`dpp_ctrl<amdgpu_synid_dpp_ctrl>` :ref:`row_mask<amdgpu_synid_row_mask>` :ref:`bank_mask<amdgpu_synid_bank_mask>` :ref:`bound_ctrl<amdgpu_synid_bound_ctrl>` + v_pk_fmac_f16 :ref:`vdst<amdgpu_synid908_vdst32_0>`::ref:`f16x2<amdgpu_synid908_type_dev>`, :ref:`src0<amdgpu_synid908_src32_0>`::ref:`f16x2<amdgpu_synid908_type_dev>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>`::ref:`f16x2<amdgpu_synid908_type_dev>` + v_xnor_b32 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_0>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>` + v_xnor_b32_dpp :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`vsrc0<amdgpu_synid908_vsrc32_0>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>` :ref:`dpp_ctrl<amdgpu_synid_dpp_ctrl>` :ref:`row_mask<amdgpu_synid_row_mask>` :ref:`bank_mask<amdgpu_synid_bank_mask>` :ref:`bound_ctrl<amdgpu_synid_bound_ctrl>` + v_xnor_b32_sdwa :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_0>`::ref:`m<amdgpu_synid908_mod_sdwa_sext>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>`::ref:`m<amdgpu_synid908_mod_sdwa_sext>` :ref:`dst_sel<amdgpu_synid_dst_sel>` :ref:`dst_unused<amdgpu_synid_dst_unused>` :ref:`src0_sel<amdgpu_synid_src0_sel>` :ref:`src1_sel<amdgpu_synid_src1_sel>` + +VOP3 +----------------------- + +.. parsed-literal:: + + **INSTRUCTION** **DST** **SRC0** **SRC1** **MODIFIERS** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + v_fmac_f32_e64 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_1>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`, :ref:`src1<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>` :ref:`clamp<amdgpu_synid_clamp>` :ref:`omod<amdgpu_synid_omod>` + v_xnor_b32_e64 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_1>`, :ref:`src1<amdgpu_synid908_src32_2>` + +VOP3P +----------------------- + +.. parsed-literal:: + + **INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + v_accvgpr_read_b32 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`asrc<amdgpu_synid908_asrc32_0>` + v_accvgpr_write_b32 :ref:`adst<amdgpu_synid908_adst32_0>`, :ref:`src<amdgpu_synid908_src32_3>` + v_dot2_f32_f16 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_1>`::ref:`f16x2<amdgpu_synid908_type_dev>`, :ref:`src1<amdgpu_synid908_src32_2>`::ref:`f16x2<amdgpu_synid908_type_dev>`, :ref:`src2<amdgpu_synid908_src32_2>`::ref:`f32<amdgpu_synid908_type_dev>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>` + v_dot2_i32_i16 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_1>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`src1<amdgpu_synid908_src32_2>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`src2<amdgpu_synid908_src32_2>`::ref:`i32<amdgpu_synid908_type_dev>` :ref:`clamp<amdgpu_synid_clamp>` + v_dot2_u32_u16 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_1>`::ref:`u16x2<amdgpu_synid908_type_dev>`, :ref:`src1<amdgpu_synid908_src32_2>`::ref:`u16x2<amdgpu_synid908_type_dev>`, :ref:`src2<amdgpu_synid908_src32_2>`::ref:`u32<amdgpu_synid908_type_dev>` :ref:`clamp<amdgpu_synid_clamp>` + v_dot4_i32_i8 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_1>`::ref:`i8x2<amdgpu_synid908_type_dev>`, :ref:`src1<amdgpu_synid908_src32_2>`::ref:`i8x2<amdgpu_synid908_type_dev>`, :ref:`src2<amdgpu_synid908_src32_2>`::ref:`i32<amdgpu_synid908_type_dev>` :ref:`clamp<amdgpu_synid_clamp>` + v_dot4_u32_u8 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_1>`::ref:`u8x2<amdgpu_synid908_type_dev>`, :ref:`src1<amdgpu_synid908_src32_2>`::ref:`u8x2<amdgpu_synid908_type_dev>`, :ref:`src2<amdgpu_synid908_src32_2>`::ref:`u32<amdgpu_synid908_type_dev>` :ref:`clamp<amdgpu_synid_clamp>` + v_dot8_i32_i4 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_1>`::ref:`i4x2<amdgpu_synid908_type_dev>`, :ref:`src1<amdgpu_synid908_src32_2>`::ref:`i4x2<amdgpu_synid908_type_dev>`, :ref:`src2<amdgpu_synid908_src32_2>`::ref:`i32<amdgpu_synid908_type_dev>` :ref:`clamp<amdgpu_synid_clamp>` + v_dot8_u32_u4 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_1>`::ref:`u4x2<amdgpu_synid908_type_dev>`, :ref:`src1<amdgpu_synid908_src32_2>`::ref:`u4x2<amdgpu_synid908_type_dev>`, :ref:`src2<amdgpu_synid908_src32_2>`::ref:`u32<amdgpu_synid908_type_dev>` :ref:`clamp<amdgpu_synid_clamp>` + v_fma_mix_f32 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_1>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`, :ref:`src1<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`, :ref:`src2<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` + v_fma_mixhi_f16 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_1>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`, :ref:`src1<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`, :ref:`src2<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` + v_fma_mixlo_f16 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_1>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`, :ref:`src1<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`, :ref:`src2<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>` + v_mfma_f32_16x16x16f16 :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>` + v_mfma_f32_16x16x1f32 :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>` + v_mfma_f32_16x16x2bf16 :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>` + v_mfma_f32_16x16x4f16 :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>` + v_mfma_f32_16x16x4f32 :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>` + v_mfma_f32_16x16x8bf16 :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>` + v_mfma_f32_32x32x1f32 :ref:`adst<amdgpu_synid908_adst1024_0>`::ref:`f32x32<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc1024_0>`::ref:`f32x32<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>` + v_mfma_f32_32x32x2bf16 :ref:`adst<amdgpu_synid908_adst1024_0>`::ref:`f32x32<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc1024_0>`::ref:`f32x32<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>` + v_mfma_f32_32x32x2f32 :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>` + v_mfma_f32_32x32x4bf16 :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>` + v_mfma_f32_32x32x4f16 :ref:`adst<amdgpu_synid908_adst1024_0>`::ref:`f32x32<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc1024_0>`::ref:`f32x32<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>` + v_mfma_f32_32x32x8f16 :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>` + v_mfma_f32_4x4x1f32 :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>` + v_mfma_f32_4x4x2bf16 :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>` + v_mfma_f32_4x4x4f16 :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>` + v_mfma_i32_16x16x16i8 :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`i32x4<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`i32x4<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>` + v_mfma_i32_16x16x4i8 :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`i32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`i32x16<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>` + v_mfma_i32_32x32x4i8 :ref:`adst<amdgpu_synid908_adst1024_0>`::ref:`i32x32<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc1024_0>`::ref:`i32x32<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>` + v_mfma_i32_32x32x8i8 :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`i32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`i32x16<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>` + v_mfma_i32_4x4x4i8 :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`i32x4<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`i32x4<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>` + +.. |---| unicode:: U+02014 .. em dash + + +.. toctree:: + :hidden: + + AMDGPUAsmGFX9 + gfx908_addr_buf + gfx908_adst1024_0 + gfx908_adst128_0 + gfx908_adst32_0 + gfx908_adst512_0 + gfx908_asrc1024_0 + gfx908_asrc128_0 + gfx908_asrc32_0 + gfx908_asrc512_0 + gfx908_data_buf_atomic32 + gfx908_dst_flat_atomic32 + gfx908_offset_buf + gfx908_rsrc_buf + gfx908_saddr_flat_global + gfx908_src32_0 + gfx908_src32_1 + gfx908_src32_2 + gfx908_src32_3 + gfx908_vaddr_flat_global + gfx908_vasrc32_0 + gfx908_vasrc64_0 + gfx908_vdata32_0 + gfx908_vdst32_0 + gfx908_vsrc32_0 + gfx908_mad_type_dev + gfx908_mod_dpp_sdwa_abs_neg + gfx908_mod_sdwa_sext + gfx908_mod_vop3_abs_neg + gfx908_opt + gfx908_ret + gfx908_type_dev diff --git a/llvm/docs/AMDGPU/gfx9_mad_type_dev.rst b/llvm/docs/AMDGPU/gfx900_mad_type_dev.rst index 0602f082176..50c403ef582 100644 --- a/llvm/docs/AMDGPU/gfx9_mad_type_dev.rst +++ b/llvm/docs/AMDGPU/gfx900_mad_type_dev.rst @@ -5,7 +5,7 @@ * * ************************************************** -.. _amdgpu_synid9_mad_type_dev: +.. _amdgpu_synid900_mad_type_dev: fx =========================== diff --git a/llvm/docs/AMDGPU/gfx900_mod_vop3_abs_neg.rst b/llvm/docs/AMDGPU/gfx900_mod_vop3_abs_neg.rst new file mode 100644 index 00000000000..fa0311a9aad --- /dev/null +++ b/llvm/docs/AMDGPU/gfx900_mod_vop3_abs_neg.rst @@ -0,0 +1,14 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid900_mod_vop3_abs_neg: + +m +=========================== + +This operand may be used with floating point operand modifiers :ref:`abs<amdgpu_synid_abs>` and :ref:`neg<amdgpu_synid_neg>`. + diff --git a/llvm/docs/AMDGPU/gfx900_src32_0.rst b/llvm/docs/AMDGPU/gfx900_src32_0.rst new file mode 100644 index 00000000000..9e2e607a056 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx900_src32_0.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid900_src32_0: + +src +=========================== + +Instruction input. + +*Size:* 1 dword. + +*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`s<amdgpu_synid_s>`, :ref:`flat_scratch<amdgpu_synid_flat_scratch>`, :ref:`vcc<amdgpu_synid_vcc>`, :ref:`ttmp<amdgpu_synid_ttmp>`, :ref:`m0<amdgpu_synid_m0>`, :ref:`exec<amdgpu_synid_exec>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`lds_direct<amdgpu_synid_lds_direct>`, :ref:`constant<amdgpu_synid_constant>` diff --git a/llvm/docs/AMDGPU/gfx900_src32_1.rst b/llvm/docs/AMDGPU/gfx900_src32_1.rst new file mode 100644 index 00000000000..82dd50b998d --- /dev/null +++ b/llvm/docs/AMDGPU/gfx900_src32_1.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid900_src32_1: + +src +=========================== + +Instruction input. + +*Size:* 1 dword. + +*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`s<amdgpu_synid_s>`, :ref:`flat_scratch<amdgpu_synid_flat_scratch>`, :ref:`vcc<amdgpu_synid_vcc>`, :ref:`ttmp<amdgpu_synid_ttmp>`, :ref:`m0<amdgpu_synid_m0>`, :ref:`exec<amdgpu_synid_exec>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`constant<amdgpu_synid_constant>` diff --git a/llvm/docs/AMDGPU/gfx900_vdst32_0.rst b/llvm/docs/AMDGPU/gfx900_vdst32_0.rst new file mode 100644 index 00000000000..0ae76da5315 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx900_vdst32_0.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid900_vdst32_0: + +vdst +=========================== + +Instruction output. + +*Size:* 1 dword. + +*Operands:* :ref:`v<amdgpu_synid_v>` diff --git a/llvm/docs/AMDGPU/gfx904_mad_type_dev.rst b/llvm/docs/AMDGPU/gfx904_mad_type_dev.rst new file mode 100644 index 00000000000..6fc3c310880 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx904_mad_type_dev.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid904_mad_type_dev: + +fx +=========================== + +This is an *f32* or *f16* operand depending on instruction modifiers: + +* Operand size is controlled by :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>`. +* Location of 16-bit operand is controlled by :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>`. + diff --git a/llvm/docs/AMDGPU/gfx904_mod_vop3_abs_neg.rst b/llvm/docs/AMDGPU/gfx904_mod_vop3_abs_neg.rst new file mode 100644 index 00000000000..814bccdf3a4 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx904_mod_vop3_abs_neg.rst @@ -0,0 +1,14 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid904_mod_vop3_abs_neg: + +m +=========================== + +This operand may be used with floating point operand modifiers :ref:`abs<amdgpu_synid_abs>` and :ref:`neg<amdgpu_synid_neg>`. + diff --git a/llvm/docs/AMDGPU/gfx904_src32_0.rst b/llvm/docs/AMDGPU/gfx904_src32_0.rst new file mode 100644 index 00000000000..a2930d4c4f9 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx904_src32_0.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid904_src32_0: + +src +=========================== + +Instruction input. + +*Size:* 1 dword. + +*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`s<amdgpu_synid_s>`, :ref:`flat_scratch<amdgpu_synid_flat_scratch>`, :ref:`vcc<amdgpu_synid_vcc>`, :ref:`ttmp<amdgpu_synid_ttmp>`, :ref:`m0<amdgpu_synid_m0>`, :ref:`exec<amdgpu_synid_exec>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`lds_direct<amdgpu_synid_lds_direct>`, :ref:`constant<amdgpu_synid_constant>` diff --git a/llvm/docs/AMDGPU/gfx904_src32_1.rst b/llvm/docs/AMDGPU/gfx904_src32_1.rst new file mode 100644 index 00000000000..23127c6c99b --- /dev/null +++ b/llvm/docs/AMDGPU/gfx904_src32_1.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid904_src32_1: + +src +=========================== + +Instruction input. + +*Size:* 1 dword. + +*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`s<amdgpu_synid_s>`, :ref:`flat_scratch<amdgpu_synid_flat_scratch>`, :ref:`vcc<amdgpu_synid_vcc>`, :ref:`ttmp<amdgpu_synid_ttmp>`, :ref:`m0<amdgpu_synid_m0>`, :ref:`exec<amdgpu_synid_exec>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`constant<amdgpu_synid_constant>` diff --git a/llvm/docs/AMDGPU/gfx904_vdst32_0.rst b/llvm/docs/AMDGPU/gfx904_vdst32_0.rst new file mode 100644 index 00000000000..6eb8bdf5b50 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx904_vdst32_0.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid904_vdst32_0: + +vdst +=========================== + +Instruction output. + +*Size:* 1 dword. + +*Operands:* :ref:`v<amdgpu_synid_v>` diff --git a/llvm/docs/AMDGPU/gfx906_mad_type_dev.rst b/llvm/docs/AMDGPU/gfx906_mad_type_dev.rst new file mode 100644 index 00000000000..eb048b0d2d1 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx906_mad_type_dev.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid906_mad_type_dev: + +fx +=========================== + +This is an *f32* or *f16* operand depending on instruction modifiers: + +* Operand size is controlled by :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>`. +* Location of 16-bit operand is controlled by :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>`. + diff --git a/llvm/docs/AMDGPU/gfx906_mod_dpp_sdwa_abs_neg.rst b/llvm/docs/AMDGPU/gfx906_mod_dpp_sdwa_abs_neg.rst new file mode 100644 index 00000000000..88fdc6f6490 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx906_mod_dpp_sdwa_abs_neg.rst @@ -0,0 +1,14 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid906_mod_dpp_sdwa_abs_neg: + +m +=========================== + +This operand may be used with floating point operand modifiers :ref:`abs<amdgpu_synid_abs>` and :ref:`neg<amdgpu_synid_neg>`. + diff --git a/llvm/docs/AMDGPU/gfx906_mod_sdwa_sext.rst b/llvm/docs/AMDGPU/gfx906_mod_sdwa_sext.rst new file mode 100644 index 00000000000..411251baeb2 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx906_mod_sdwa_sext.rst @@ -0,0 +1,14 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid906_mod_sdwa_sext: + +m +=========================== + +This operand may be used with integer operand modifier :ref:`sext<amdgpu_synid_sext>`. + diff --git a/llvm/docs/AMDGPU/gfx906_mod_vop3_abs_neg.rst b/llvm/docs/AMDGPU/gfx906_mod_vop3_abs_neg.rst new file mode 100644 index 00000000000..ffd6429640d --- /dev/null +++ b/llvm/docs/AMDGPU/gfx906_mod_vop3_abs_neg.rst @@ -0,0 +1,14 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid906_mod_vop3_abs_neg: + +m +=========================== + +This operand may be used with floating point operand modifiers :ref:`abs<amdgpu_synid_abs>` and :ref:`neg<amdgpu_synid_neg>`. + diff --git a/llvm/docs/AMDGPU/gfx906_src32_0.rst b/llvm/docs/AMDGPU/gfx906_src32_0.rst new file mode 100644 index 00000000000..f547beaec23 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx906_src32_0.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid906_src32_0: + +src +=========================== + +Instruction input. + +*Size:* 1 dword. + +*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`s<amdgpu_synid_s>`, :ref:`flat_scratch<amdgpu_synid_flat_scratch>`, :ref:`vcc<amdgpu_synid_vcc>`, :ref:`ttmp<amdgpu_synid_ttmp>`, :ref:`m0<amdgpu_synid_m0>`, :ref:`exec<amdgpu_synid_exec>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`lds_direct<amdgpu_synid_lds_direct>`, :ref:`constant<amdgpu_synid_constant>`, :ref:`literal<amdgpu_synid_literal>` diff --git a/llvm/docs/AMDGPU/gfx906_src32_1.rst b/llvm/docs/AMDGPU/gfx906_src32_1.rst new file mode 100644 index 00000000000..1fc90e163b1 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx906_src32_1.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid906_src32_1: + +src +=========================== + +Instruction input. + +*Size:* 1 dword. + +*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`s<amdgpu_synid_s>`, :ref:`flat_scratch<amdgpu_synid_flat_scratch>`, :ref:`vcc<amdgpu_synid_vcc>`, :ref:`ttmp<amdgpu_synid_ttmp>`, :ref:`m0<amdgpu_synid_m0>`, :ref:`exec<amdgpu_synid_exec>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`lds_direct<amdgpu_synid_lds_direct>`, :ref:`constant<amdgpu_synid_constant>` diff --git a/llvm/docs/AMDGPU/gfx906_src32_2.rst b/llvm/docs/AMDGPU/gfx906_src32_2.rst new file mode 100644 index 00000000000..6c3bb387c4e --- /dev/null +++ b/llvm/docs/AMDGPU/gfx906_src32_2.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid906_src32_2: + +src +=========================== + +Instruction input. + +*Size:* 1 dword. + +*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`s<amdgpu_synid_s>`, :ref:`flat_scratch<amdgpu_synid_flat_scratch>`, :ref:`vcc<amdgpu_synid_vcc>`, :ref:`ttmp<amdgpu_synid_ttmp>`, :ref:`m0<amdgpu_synid_m0>`, :ref:`exec<amdgpu_synid_exec>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`constant<amdgpu_synid_constant>` diff --git a/llvm/docs/AMDGPU/gfx906_type_dev.rst b/llvm/docs/AMDGPU/gfx906_type_dev.rst new file mode 100644 index 00000000000..02f8d98a3dc --- /dev/null +++ b/llvm/docs/AMDGPU/gfx906_type_dev.rst @@ -0,0 +1,14 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid906_type_dev: + +Type deviation +=========================== + +*Type* of this operand differs from *type* :ref:`implied by the opcode<amdgpu_syn_instruction_type>`. This tag specifies actual operand *type*. + diff --git a/llvm/docs/AMDGPU/gfx906_vdst32_0.rst b/llvm/docs/AMDGPU/gfx906_vdst32_0.rst new file mode 100644 index 00000000000..faf0a2d2f9d --- /dev/null +++ b/llvm/docs/AMDGPU/gfx906_vdst32_0.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid906_vdst32_0: + +vdst +=========================== + +Instruction output. + +*Size:* 1 dword. + +*Operands:* :ref:`v<amdgpu_synid_v>` diff --git a/llvm/docs/AMDGPU/gfx906_vsrc32_0.rst b/llvm/docs/AMDGPU/gfx906_vsrc32_0.rst new file mode 100644 index 00000000000..ed350d997eb --- /dev/null +++ b/llvm/docs/AMDGPU/gfx906_vsrc32_0.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid906_vsrc32_0: + +vsrc +=========================== + +Instruction input. + +*Size:* 1 dword. + +*Operands:* :ref:`v<amdgpu_synid_v>` diff --git a/llvm/docs/AMDGPU/gfx908_addr_buf.rst b/llvm/docs/AMDGPU/gfx908_addr_buf.rst new file mode 100644 index 00000000000..9dc593be69e --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_addr_buf.rst @@ -0,0 +1,22 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_addr_buf: + +vaddr +=========================== + +This is an optional operand which may specify offset and/or index. + +*Size:* 0, 1 or 2 dwords. Size is controlled by modifiers :ref:`offen<amdgpu_synid_offen>` and :ref:`idxen<amdgpu_synid_idxen>`: + +* If only :ref:`idxen<amdgpu_synid_idxen>` is specified, this operand supplies an index. Size is 1 dword. +* If only :ref:`offen<amdgpu_synid_offen>` is specified, this operand supplies an offset. Size is 1 dword. +* If both modifiers are specified, index is in the first register and offset is in the second. Size is 2 dwords. +* If none of these modifiers are specified, this operand must be set to :ref:`off<amdgpu_synid_off>`. + +*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`off<amdgpu_synid_off>` diff --git a/llvm/docs/AMDGPU/gfx908_adst1024_0.rst b/llvm/docs/AMDGPU/gfx908_adst1024_0.rst new file mode 100644 index 00000000000..3dfe45ede46 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_adst1024_0.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_adst1024_0: + +adst +=========================== + +Instruction output. + +*Size:* 32 dwords. + +*Operands:* :ref:`a<amdgpu_synid_a>` diff --git a/llvm/docs/AMDGPU/gfx908_adst128_0.rst b/llvm/docs/AMDGPU/gfx908_adst128_0.rst new file mode 100644 index 00000000000..23c7f9db057 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_adst128_0.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_adst128_0: + +adst +=========================== + +Instruction output. + +*Size:* 4 dwords. + +*Operands:* :ref:`a<amdgpu_synid_a>` diff --git a/llvm/docs/AMDGPU/gfx908_adst32_0.rst b/llvm/docs/AMDGPU/gfx908_adst32_0.rst new file mode 100644 index 00000000000..80546fa1c83 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_adst32_0.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_adst32_0: + +adst +=========================== + +Instruction output. + +*Size:* 1 dword. + +*Operands:* :ref:`a<amdgpu_synid_a>` diff --git a/llvm/docs/AMDGPU/gfx908_adst512_0.rst b/llvm/docs/AMDGPU/gfx908_adst512_0.rst new file mode 100644 index 00000000000..446eb02bf03 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_adst512_0.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_adst512_0: + +adst +=========================== + +Instruction output. + +*Size:* 16 dwords. + +*Operands:* :ref:`a<amdgpu_synid_a>` diff --git a/llvm/docs/AMDGPU/gfx908_asrc1024_0.rst b/llvm/docs/AMDGPU/gfx908_asrc1024_0.rst new file mode 100644 index 00000000000..00e1fc7f0f6 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_asrc1024_0.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_asrc1024_0: + +asrc +=========================== + +Instruction input. + +*Size:* 32 dwords. + +*Operands:* :ref:`a<amdgpu_synid_a>` diff --git a/llvm/docs/AMDGPU/gfx908_asrc128_0.rst b/llvm/docs/AMDGPU/gfx908_asrc128_0.rst new file mode 100644 index 00000000000..91af6b7357b --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_asrc128_0.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_asrc128_0: + +asrc +=========================== + +Instruction input. + +*Size:* 4 dwords. + +*Operands:* :ref:`a<amdgpu_synid_a>` diff --git a/llvm/docs/AMDGPU/gfx908_asrc32_0.rst b/llvm/docs/AMDGPU/gfx908_asrc32_0.rst new file mode 100644 index 00000000000..8e46317254d --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_asrc32_0.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_asrc32_0: + +asrc +=========================== + +Instruction input. + +*Size:* 1 dword. + +*Operands:* :ref:`a<amdgpu_synid_a>` diff --git a/llvm/docs/AMDGPU/gfx908_asrc512_0.rst b/llvm/docs/AMDGPU/gfx908_asrc512_0.rst new file mode 100644 index 00000000000..59e22040c22 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_asrc512_0.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_asrc512_0: + +asrc +=========================== + +Instruction input. + +*Size:* 16 dwords. + +*Operands:* :ref:`a<amdgpu_synid_a>` diff --git a/llvm/docs/AMDGPU/gfx908_data_buf_atomic32.rst b/llvm/docs/AMDGPU/gfx908_data_buf_atomic32.rst new file mode 100644 index 00000000000..03163f55fe5 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_data_buf_atomic32.rst @@ -0,0 +1,21 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_data_buf_atomic32: + +vdata +=========================== + +Input data for an atomic instruction. + +Optionally may serve as an output data: + +* If :ref:`glc<amdgpu_synid_glc>` is specified, gets the memory value before the operation. + +*Size:* 1 dword by default. :ref:`tfe<amdgpu_synid_tfe>` adds 1 dword if specified. + +*Operands:* :ref:`v<amdgpu_synid_v>` diff --git a/llvm/docs/AMDGPU/gfx908_dst_flat_atomic32.rst b/llvm/docs/AMDGPU/gfx908_dst_flat_atomic32.rst new file mode 100644 index 00000000000..045058314cb --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_dst_flat_atomic32.rst @@ -0,0 +1,19 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_dst_flat_atomic32: + +vdst +=========================== + +Data returned by a 32-bit atomic flat instruction. + +This is an optional operand. It must be used if and only if :ref:`glc<amdgpu_synid_glc>` is specified. + +*Size:* 1 dword. + +*Operands:* :ref:`v<amdgpu_synid_v>` diff --git a/llvm/docs/AMDGPU/gfx908_mad_type_dev.rst b/llvm/docs/AMDGPU/gfx908_mad_type_dev.rst new file mode 100644 index 00000000000..bbc3ea9dde5 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_mad_type_dev.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_mad_type_dev: + +fx +=========================== + +This is an *f32* or *f16* operand depending on instruction modifiers: + +* Operand size is controlled by :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>`. +* Location of 16-bit operand is controlled by :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>`. + diff --git a/llvm/docs/AMDGPU/gfx908_mod_dpp_sdwa_abs_neg.rst b/llvm/docs/AMDGPU/gfx908_mod_dpp_sdwa_abs_neg.rst new file mode 100644 index 00000000000..028a209ebe5 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_mod_dpp_sdwa_abs_neg.rst @@ -0,0 +1,14 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_mod_dpp_sdwa_abs_neg: + +m +=========================== + +This operand may be used with floating point operand modifiers :ref:`abs<amdgpu_synid_abs>` and :ref:`neg<amdgpu_synid_neg>`. + diff --git a/llvm/docs/AMDGPU/gfx908_mod_sdwa_sext.rst b/llvm/docs/AMDGPU/gfx908_mod_sdwa_sext.rst new file mode 100644 index 00000000000..1afede5e747 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_mod_sdwa_sext.rst @@ -0,0 +1,14 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_mod_sdwa_sext: + +m +=========================== + +This operand may be used with integer operand modifier :ref:`sext<amdgpu_synid_sext>`. + diff --git a/llvm/docs/AMDGPU/gfx908_mod_vop3_abs_neg.rst b/llvm/docs/AMDGPU/gfx908_mod_vop3_abs_neg.rst new file mode 100644 index 00000000000..07bec0c6871 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_mod_vop3_abs_neg.rst @@ -0,0 +1,14 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_mod_vop3_abs_neg: + +m +=========================== + +This operand may be used with floating point operand modifiers :ref:`abs<amdgpu_synid_abs>` and :ref:`neg<amdgpu_synid_neg>`. + diff --git a/llvm/docs/AMDGPU/gfx908_offset_buf.rst b/llvm/docs/AMDGPU/gfx908_offset_buf.rst new file mode 100644 index 00000000000..6f3d6d2b96c --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_offset_buf.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_offset_buf: + +soffset +=========================== + +An unsigned byte offset. + +*Size:* 1 dword. + +*Operands:* :ref:`s<amdgpu_synid_s>`, :ref:`flat_scratch<amdgpu_synid_flat_scratch>`, :ref:`vcc<amdgpu_synid_vcc>`, :ref:`ttmp<amdgpu_synid_ttmp>`, :ref:`m0<amdgpu_synid_m0>`, :ref:`exec<amdgpu_synid_exec>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`constant<amdgpu_synid_constant>` diff --git a/llvm/docs/AMDGPU/gfx908_opt.rst b/llvm/docs/AMDGPU/gfx908_opt.rst new file mode 100644 index 00000000000..efd5a79c8a7 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_opt.rst @@ -0,0 +1,14 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_opt: + +opt +=========================== + +This is an optional operand. It must be used if and only if :ref:`glc<amdgpu_synid_glc>` is specified. + diff --git a/llvm/docs/AMDGPU/gfx908_ret.rst b/llvm/docs/AMDGPU/gfx908_ret.rst new file mode 100644 index 00000000000..dc8a91c5b4c --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_ret.rst @@ -0,0 +1,14 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_ret: + +dst +=========================== + +This is an input operand. It may optionally serve as a destination if :ref:`glc<amdgpu_synid_glc>` is specified. + diff --git a/llvm/docs/AMDGPU/gfx908_rsrc_buf.rst b/llvm/docs/AMDGPU/gfx908_rsrc_buf.rst new file mode 100644 index 00000000000..832adfe0763 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_rsrc_buf.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_rsrc_buf: + +srsrc +=========================== + +Buffer resource constant which defines the address and characteristics of the buffer in memory. + +*Size:* 4 dwords. + +*Operands:* :ref:`s<amdgpu_synid_s>`, :ref:`ttmp<amdgpu_synid_ttmp>` diff --git a/llvm/docs/AMDGPU/gfx908_saddr_flat_global.rst b/llvm/docs/AMDGPU/gfx908_saddr_flat_global.rst new file mode 100644 index 00000000000..a0e450639e0 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_saddr_flat_global.rst @@ -0,0 +1,19 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_saddr_flat_global: + +saddr +=========================== + +An optional 64-bit flat global address. Must be specified as :ref:`off<amdgpu_synid_off>` if not used. + +See :ref:`vaddr<amdgpu_synid908_vaddr_flat_global>` for description of available addressing modes. + +*Size:* 2 dwords. + +*Operands:* :ref:`exec<amdgpu_synid_exec>`, :ref:`off<amdgpu_synid_off>` diff --git a/llvm/docs/AMDGPU/gfx908_src32_0.rst b/llvm/docs/AMDGPU/gfx908_src32_0.rst new file mode 100644 index 00000000000..a6bf415d402 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_src32_0.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_src32_0: + +src +=========================== + +Instruction input. + +*Size:* 1 dword. + +*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`s<amdgpu_synid_s>`, :ref:`flat_scratch<amdgpu_synid_flat_scratch>`, :ref:`vcc<amdgpu_synid_vcc>`, :ref:`ttmp<amdgpu_synid_ttmp>`, :ref:`m0<amdgpu_synid_m0>`, :ref:`exec<amdgpu_synid_exec>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`lds_direct<amdgpu_synid_lds_direct>`, :ref:`constant<amdgpu_synid_constant>`, :ref:`literal<amdgpu_synid_literal>` diff --git a/llvm/docs/AMDGPU/gfx908_src32_1.rst b/llvm/docs/AMDGPU/gfx908_src32_1.rst new file mode 100644 index 00000000000..275e4d796b9 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_src32_1.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_src32_1: + +src +=========================== + +Instruction input. + +*Size:* 1 dword. + +*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`s<amdgpu_synid_s>`, :ref:`flat_scratch<amdgpu_synid_flat_scratch>`, :ref:`vcc<amdgpu_synid_vcc>`, :ref:`ttmp<amdgpu_synid_ttmp>`, :ref:`m0<amdgpu_synid_m0>`, :ref:`exec<amdgpu_synid_exec>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`lds_direct<amdgpu_synid_lds_direct>`, :ref:`constant<amdgpu_synid_constant>` diff --git a/llvm/docs/AMDGPU/gfx908_src32_2.rst b/llvm/docs/AMDGPU/gfx908_src32_2.rst new file mode 100644 index 00000000000..21d0f201f26 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_src32_2.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_src32_2: + +src +=========================== + +Instruction input. + +*Size:* 1 dword. + +*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`s<amdgpu_synid_s>`, :ref:`flat_scratch<amdgpu_synid_flat_scratch>`, :ref:`vcc<amdgpu_synid_vcc>`, :ref:`ttmp<amdgpu_synid_ttmp>`, :ref:`m0<amdgpu_synid_m0>`, :ref:`exec<amdgpu_synid_exec>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`constant<amdgpu_synid_constant>` diff --git a/llvm/docs/AMDGPU/gfx908_src32_3.rst b/llvm/docs/AMDGPU/gfx908_src32_3.rst new file mode 100644 index 00000000000..0e9573d45da --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_src32_3.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_src32_3: + +src +=========================== + +Instruction input. + +*Size:* 1 dword. + +*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`constant<amdgpu_synid_constant>` diff --git a/llvm/docs/AMDGPU/gfx908_type_dev.rst b/llvm/docs/AMDGPU/gfx908_type_dev.rst new file mode 100644 index 00000000000..8d8da92eb1c --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_type_dev.rst @@ -0,0 +1,14 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_type_dev: + +Type deviation +=========================== + +*Type* of this operand differs from *type* :ref:`implied by the opcode<amdgpu_syn_instruction_type>`. This tag specifies actual operand *type*. + diff --git a/llvm/docs/AMDGPU/gfx908_vaddr_flat_global.rst b/llvm/docs/AMDGPU/gfx908_vaddr_flat_global.rst new file mode 100644 index 00000000000..7b53cb3c722 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_vaddr_flat_global.rst @@ -0,0 +1,22 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_vaddr_flat_global: + +vaddr +=========================== + +A 64-bit flat global address or a 32-bit offset depending on addressing mode: + +* Address = :ref:`vaddr<amdgpu_synid908_vaddr_flat_global>` + :ref:`offset13s<amdgpu_synid_flat_offset13s>`. :ref:`vaddr<amdgpu_synid908_vaddr_flat_global>` is a 64-bit address. This mode is indicated by :ref:`saddr<amdgpu_synid908_saddr_flat_global>` set to :ref:`off<amdgpu_synid_off>`. +* Address = :ref:`saddr<amdgpu_synid908_saddr_flat_global>` + :ref:`vaddr<amdgpu_synid908_vaddr_flat_global>` + :ref:`offset13s<amdgpu_synid_flat_offset13s>`. :ref:`vaddr<amdgpu_synid908_vaddr_flat_global>` is a 32-bit offset. This mode is used when :ref:`saddr<amdgpu_synid908_saddr_flat_global>` is not :ref:`off<amdgpu_synid_off>`. + +.. WARNING:: Assembler currently expects a 64-bit *vaddr* regardless of addressing mode. This have to be fixed. + +*Size:* 1 or 2 dwords. + +*Operands:* :ref:`v<amdgpu_synid_v>` diff --git a/llvm/docs/AMDGPU/gfx908_vasrc32_0.rst b/llvm/docs/AMDGPU/gfx908_vasrc32_0.rst new file mode 100644 index 00000000000..52aaaadc580 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_vasrc32_0.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_vasrc32_0: + +vasrc +=========================== + +Instruction input. + +*Size:* 1 dword. + +*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`a<amdgpu_synid_a>` diff --git a/llvm/docs/AMDGPU/gfx908_vasrc64_0.rst b/llvm/docs/AMDGPU/gfx908_vasrc64_0.rst new file mode 100644 index 00000000000..bef80d9637a --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_vasrc64_0.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_vasrc64_0: + +vasrc +=========================== + +Instruction input. + +*Size:* 2 dwords. + +*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`a<amdgpu_synid_a>` diff --git a/llvm/docs/AMDGPU/gfx908_vdata32_0.rst b/llvm/docs/AMDGPU/gfx908_vdata32_0.rst new file mode 100644 index 00000000000..7571818881d --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_vdata32_0.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_vdata32_0: + +vdata +=========================== + +Instruction input. + +*Size:* 1 dword. + +*Operands:* :ref:`v<amdgpu_synid_v>` diff --git a/llvm/docs/AMDGPU/gfx908_vdst32_0.rst b/llvm/docs/AMDGPU/gfx908_vdst32_0.rst new file mode 100644 index 00000000000..ed2214a8fbe --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_vdst32_0.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_vdst32_0: + +vdst +=========================== + +Instruction output. + +*Size:* 1 dword. + +*Operands:* :ref:`v<amdgpu_synid_v>` diff --git a/llvm/docs/AMDGPU/gfx908_vsrc32_0.rst b/llvm/docs/AMDGPU/gfx908_vsrc32_0.rst new file mode 100644 index 00000000000..c990b504515 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_vsrc32_0.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid908_vsrc32_0: + +vsrc +=========================== + +Instruction input. + +*Size:* 1 dword. + +*Operands:* :ref:`v<amdgpu_synid_v>` diff --git a/llvm/docs/AMDGPUModifierSyntax.rst b/llvm/docs/AMDGPUModifierSyntax.rst index 526016d4f4f..5e29004b4d6 100644 --- a/llvm/docs/AMDGPUModifierSyntax.rst +++ b/llvm/docs/AMDGPUModifierSyntax.rst @@ -1337,7 +1337,7 @@ VOP3P Modifiers This section describes modifiers of *regular* VOP3P instructions. -*v_mad_mix_f32*, *v_mad_mixhi_f16* and *v_mad_mixlo_f16* +*v_mad_mix\** and *v_fma_mix\** instructions use these modifiers :ref:`in a special manner<amdgpu_synid_mad_mix>`. GFX9 and GFX10 only. @@ -1494,8 +1494,8 @@ See a description :ref:`here<amdgpu_synid_clamp>`. VOP3P V_MAD_MIX Modifiers ------------------------- -*v_mad_mix_f32*, *v_mad_mixhi_f16* and *v_mad_mixlo_f16* instructions -use *op_sel* and *op_sel_hi* modifiers +*v_mad_mix\** and *v_fma_mix\** +instructions use *op_sel* and *op_sel_hi* modifiers in a manner different from *regular* VOP3P instructions. See a description below. @@ -1581,3 +1581,52 @@ clamp ~~~~~ See a description :ref:`here<amdgpu_synid_clamp>`. + +VOP3P MFMA Modifiers +-------------------- + +.. _amdgpu_synid_cbsz: + +cbsz +~~~~ + + =============================== ================================================================== + Syntax Description + =============================== ================================================================== + cbsz:[{0..7}] TBD + =============================== ================================================================== + +Note: numeric value may be specified as either +an :ref:`integer number<amdgpu_synid_integer_number>` or +an :ref:`absolute expression<amdgpu_synid_absolute_expression>`. + +.. _amdgpu_synid_abid: + +abid +~~~~ + + =============================== ================================================================== + Syntax Description + =============================== ================================================================== + abid:[{0..15}] TBD + =============================== ================================================================== + +Note: numeric value may be specified as either +an :ref:`integer number<amdgpu_synid_integer_number>` or +an :ref:`absolute expression<amdgpu_synid_absolute_expression>`. + +.. _amdgpu_synid_blgp: + +blgp +~~~~ + + =============================== ================================================================== + Syntax Description + =============================== ================================================================== + blgp:[{0..7}] TBD + =============================== ================================================================== + +Note: numeric value may be specified as either +an :ref:`integer number<amdgpu_synid_integer_number>` or +an :ref:`absolute expression<amdgpu_synid_absolute_expression>`. + diff --git a/llvm/docs/AMDGPUOperandSyntax.rst b/llvm/docs/AMDGPUOperandSyntax.rst index c20da004729..842f49fbf73 100644 --- a/llvm/docs/AMDGPUOperandSyntax.rst +++ b/llvm/docs/AMDGPUOperandSyntax.rst @@ -102,6 +102,66 @@ Examples: [v[32],v[1:1],[v2]] [v4,v4,v4,v4] +.. _amdgpu_synid_a: + +a +- + +Accumulator registers. There are 256 32-bit accumulator registers. + +A sequence of *accumulator* registers may be used to operate with more than 32 bits of data. + +Assembler currently supports sequences of 1, 2, 4 and 16 *accumulator* registers. + + =================================================== ========================================================= ==================================================================== + Syntax An Alternative Syntax (SP3) Description + =================================================== ========================================================= ==================================================================== + **a**\<N> **acc**\<N> A single 32-bit *accumulator* register. + + *N* must be a decimal + :ref:`integer number<amdgpu_synid_integer_number>`. + **a[**\ <N>\ **]** **acc[**\ <N>\ **]** A single 32-bit *accumulator* register. + + *N* may be specified as an + :ref:`integer number<amdgpu_synid_integer_number>` + or an :ref:`absolute expression<amdgpu_synid_absolute_expression>`. + **a[**\ <N>:<K>\ **]** **acc[**\ <N>:<K>\ **]** A sequence of (\ *K-N+1*\ ) *accumulator* registers. + + *N* and *K* may be specified as + :ref:`integer numbers<amdgpu_synid_integer_number>` + or :ref:`absolute expressions<amdgpu_synid_absolute_expression>`. + **[a**\ <N>, \ **a**\ <N+1>, ... **a**\ <K>\ **]** **[acc**\ <N>, \ **acc**\ <N+1>, ... **acc**\ <K>\ **]** A sequence of (\ *K-N+1*\ ) *accumulator* registers. + + Register indices must be specified as decimal + :ref:`integer numbers<amdgpu_synid_integer_number>`. + =================================================== ========================================================= ==================================================================== + +Note: *N* and *K* must satisfy the following conditions: + +* *N* <= *K*. +* 0 <= *N* <= 255. +* 0 <= *K* <= 255. +* *K-N+1* must be equal to 1, 2, 4 or 16. + +Examples: + +.. parsed-literal:: + + a255 + a[0] + a[0:1] + a[1:1] + a[0:3] + a[2*2] + a[1-1:2-1] + [a252] + [a252,a253,a254,a255] + + acc0 + acc[1] + [acc250] + [acc2,acc3] + .. _amdgpu_synid_s: s diff --git a/llvm/docs/AMDGPUUsage.rst b/llvm/docs/AMDGPUUsage.rst index ee165f0fc10..d9c64fd83a1 100644 --- a/llvm/docs/AMDGPUUsage.rst +++ b/llvm/docs/AMDGPUUsage.rst @@ -5781,6 +5781,10 @@ Instructions AMDGPU/AMDGPUAsmGFX7 AMDGPU/AMDGPUAsmGFX8 AMDGPU/AMDGPUAsmGFX9 + AMDGPU/AMDGPUAsmGFX900 + AMDGPU/AMDGPUAsmGFX904 + AMDGPU/AMDGPUAsmGFX906 + AMDGPU/AMDGPUAsmGFX908 AMDGPU/AMDGPUAsmGFX10 AMDGPUModifierSyntax AMDGPUOperandSyntax @@ -5792,17 +5796,37 @@ An instruction has the following :doc:`syntax<AMDGPUInstructionSyntax>`: | ``<``\ *opcode*\ ``> <``\ *operand0*\ ``>, <``\ *operand1*\ ``>,... <``\ *modifier0*\ ``> <``\ *modifier1*\ ``>...`` -:doc:`Operands<AMDGPUOperandSyntax>` are normally comma-separated while +:doc:`Operands<AMDGPUOperandSyntax>` are comma-separated while :doc:`modifiers<AMDGPUModifierSyntax>` are space-separated. -The order of *operands* and *modifiers* is fixed. -Most *modifiers* are optional and may be omitted. +The order of operands and modifiers is fixed. +Most modifiers are optional and may be omitted. -See detailed instruction syntax description for -:doc:`GFX7<AMDGPU/AMDGPUAsmGFX7>`, :doc:`GFX8<AMDGPU/AMDGPUAsmGFX8>`, -:doc:`GFX9<AMDGPU/AMDGPUAsmGFX9>`, and :doc:`GFX10<AMDGPU/AMDGPUAsmGFX10>`. +Links to detailed instruction syntax description may be found in the following +table. Note that features under development are not included +in this description. -Note that features under development are not included in this description. + ==================================== ====================================== + Core ISA ISA Extensions + ==================================== ====================================== + :doc:`GFX7<AMDGPU/AMDGPUAsmGFX7>` \- + :doc:`GFX8<AMDGPU/AMDGPUAsmGFX8>` \- + :doc:`GFX9<AMDGPU/AMDGPUAsmGFX9>` :doc:`gfx900<AMDGPU/AMDGPUAsmGFX900>` + + :doc:`gfx902<AMDGPU/AMDGPUAsmGFX900>` + + :doc:`gfx904<AMDGPU/AMDGPUAsmGFX904>` + + :doc:`gfx906<AMDGPU/AMDGPUAsmGFX906>` + + :doc:`gfx908<AMDGPU/AMDGPUAsmGFX908>` + + :doc:`gfx909<AMDGPU/AMDGPUAsmGFX900>` + + :doc:`GFX10<AMDGPU/AMDGPUAsmGFX10>` gfx1011 + + gfx1012 + ==================================== ====================================== For more information about instructions, their semantics and supported combinations of operands, refer to one of instruction set architecture manuals |
