diff --git a/llvm/docs/AMDGPU/AMDGPUAsmGFX10.rst b/llvm/docs/AMDGPU/AMDGPUAsmGFX10.rst index 033a16355672b..bc2b4f74ecfef 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`. -Introduction -============ +Overvew +======= An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document`. @@ -1359,56 +1364,56 @@ VOP2 .. parsed-literal:: - **INSTRUCTION** **DST0** **DST1** **SRC0** **SRC1** **SRC2** - \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| - v_add_co_ci_u32 :ref:`vdst`, :ref:`vcc`, :ref:`src0`, :ref:`vsrc1`, :ref:`vcc` - v_add_f16 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_add_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_add_nc_u32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_and_b32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_ashrrev_i32 :ref:`vdst`, :ref:`src0`::ref:`u32`, :ref:`vsrc1` - v_cndmask_b32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1`, :ref:`vcc` - v_cvt_pkrtz_f16_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_fmaak_f16 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1`, :ref:`imm32` - v_fmaak_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1`, :ref:`imm32` - v_fmac_f16 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_fmac_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_fmamk_f16 :ref:`vdst`, :ref:`src0`, :ref:`imm32`, :ref:`vsrc2` - v_fmamk_f32 :ref:`vdst`, :ref:`src0`, :ref:`imm32`, :ref:`vsrc2` - v_ldexp_f16 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1`::ref:`i16` - v_lshlrev_b32 :ref:`vdst`, :ref:`src0`::ref:`u32`, :ref:`vsrc1` - v_lshrrev_b32 :ref:`vdst`, :ref:`src0`::ref:`u32`, :ref:`vsrc1` - v_mac_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_mac_legacy_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_madak_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1`, :ref:`imm32` - v_madmk_f32 :ref:`vdst`, :ref:`src0`, :ref:`imm32`, :ref:`vsrc2` - v_max_f16 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_max_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_max_i32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_max_u32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_min_f16 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_min_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_min_i32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_min_u32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_mul_f16 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_mul_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_mul_hi_i32_i24 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_mul_hi_u32_u24 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_mul_i32_i24 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_mul_legacy_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_mul_u32_u24 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_or_b32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_pk_fmac_f16 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_sub_co_ci_u32 :ref:`vdst`, :ref:`vcc`, :ref:`src0`, :ref:`vsrc1`, :ref:`vcc` - v_sub_f16 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_sub_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_sub_nc_u32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_subrev_co_ci_u32 :ref:`vdst`, :ref:`vcc`, :ref:`src0`, :ref:`vsrc1`, :ref:`vcc` - v_subrev_f16 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_subrev_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_subrev_nc_u32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_xnor_b32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` - v_xor_b32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + **INSTRUCTION** **DST0** **DST1** **SRC0** **SRC1** **SRC2** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + v_add_co_ci_u32 :ref:`vdst`, :ref:`vcc`, :ref:`src0`, :ref:`vsrc1`, :ref:`vcc` + v_add_f16 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_add_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_add_nc_u32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_and_b32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_ashrrev_i32 :ref:`vdst`, :ref:`src0`::ref:`u32`, :ref:`vsrc1` + v_cndmask_b32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1`, :ref:`vcc` + v_cvt_pkrtz_f16_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_fmaak_f16 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1`, :ref:`imm32` + v_fmaak_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1`, :ref:`imm32` + v_fmac_f16 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_fmac_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_fmamk_f16 :ref:`vdst`, :ref:`src0`, :ref:`imm32`, :ref:`vsrc2` + v_fmamk_f32 :ref:`vdst`, :ref:`src0`, :ref:`imm32`, :ref:`vsrc2` + v_ldexp_f16 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1`::ref:`i16` + v_lshlrev_b32 :ref:`vdst`, :ref:`src0`::ref:`u32`, :ref:`vsrc1` + v_lshrrev_b32 :ref:`vdst`, :ref:`src0`::ref:`u32`, :ref:`vsrc1` + v_mac_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_mac_legacy_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_madak_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1`, :ref:`imm32` + v_madmk_f32 :ref:`vdst`, :ref:`src0`, :ref:`imm32`, :ref:`vsrc2` + v_max_f16 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_max_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_max_i32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_max_u32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_min_f16 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_min_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_min_i32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_min_u32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_mul_f16 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_mul_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_mul_hi_i32_i24 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_mul_hi_u32_u24 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_mul_i32_i24 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_mul_legacy_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_mul_u32_u24 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_or_b32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_pk_fmac_f16 :ref:`vdst`::ref:`f16x2`, :ref:`src0`::ref:`f16x2`, :ref:`vsrc1`::ref:`f16x2` + v_sub_co_ci_u32 :ref:`vdst`, :ref:`vcc`, :ref:`src0`, :ref:`vsrc1`, :ref:`vcc` + v_sub_f16 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_sub_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_sub_nc_u32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_subrev_co_ci_u32 :ref:`vdst`, :ref:`vcc`, :ref:`src0`, :ref:`vsrc1`, :ref:`vcc` + v_subrev_f16 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_subrev_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_subrev_nc_u32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_xnor_b32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_xor_b32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` VOP3 ----------------------- @@ -1847,30 +1852,30 @@ VOP3P .. parsed-literal:: - **INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS** - \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| - v_fma_mix_f32 :ref:`vdst`, :ref:`src0`::ref:`fx`, :ref:`src1`::ref:`fx`, :ref:`src2`::ref:`fx` :ref:`m_op_sel` :ref:`m_op_sel_hi` :ref:`clamp` - v_fma_mixhi_f16 :ref:`vdst`, :ref:`src0`::ref:`fx`, :ref:`src1`::ref:`fx`, :ref:`src2`::ref:`fx` :ref:`m_op_sel` :ref:`m_op_sel_hi` :ref:`clamp` - v_fma_mixlo_f16 :ref:`vdst`, :ref:`src0`::ref:`fx`, :ref:`src1`::ref:`fx`, :ref:`src2`::ref:`fx` :ref:`m_op_sel` :ref:`m_op_sel_hi` :ref:`clamp` - v_pk_add_f16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` :ref:`neg_lo` :ref:`neg_hi` :ref:`clamp` - v_pk_add_i16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` :ref:`clamp` - v_pk_add_u16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` :ref:`clamp` - v_pk_ashrrev_i16 :ref:`vdst`, :ref:`src0`::ref:`u16x2`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` - v_pk_fma_f16 :ref:`vdst`, :ref:`src0`, :ref:`src1`, :ref:`src2` :ref:`op_sel` :ref:`op_sel_hi` :ref:`neg_lo` :ref:`neg_hi` :ref:`clamp` - v_pk_lshlrev_b16 :ref:`vdst`, :ref:`src0`::ref:`u16x2`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` - v_pk_lshrrev_b16 :ref:`vdst`, :ref:`src0`::ref:`u16x2`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` - v_pk_mad_i16 :ref:`vdst`, :ref:`src0`, :ref:`src1`, :ref:`src2` :ref:`op_sel` :ref:`op_sel_hi` :ref:`clamp` - v_pk_mad_u16 :ref:`vdst`, :ref:`src0`, :ref:`src1`, :ref:`src2` :ref:`op_sel` :ref:`op_sel_hi` :ref:`clamp` - v_pk_max_f16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` :ref:`neg_lo` :ref:`neg_hi` :ref:`clamp` - v_pk_max_i16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` - v_pk_max_u16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` - v_pk_min_f16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` :ref:`neg_lo` :ref:`neg_hi` :ref:`clamp` - v_pk_min_i16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` - v_pk_min_u16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` - v_pk_mul_f16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` :ref:`neg_lo` :ref:`neg_hi` :ref:`clamp` - v_pk_mul_lo_u16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` - v_pk_sub_i16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` :ref:`clamp` - v_pk_sub_u16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` :ref:`clamp` + **INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + v_fma_mix_f32 :ref:`vdst`, :ref:`src0`::ref:`m`::ref:`fx`, :ref:`src1`::ref:`m`::ref:`fx`, :ref:`src2`::ref:`m`::ref:`fx` :ref:`m_op_sel` :ref:`m_op_sel_hi` :ref:`clamp` + v_fma_mixhi_f16 :ref:`vdst`, :ref:`src0`::ref:`m`::ref:`fx`, :ref:`src1`::ref:`m`::ref:`fx`, :ref:`src2`::ref:`m`::ref:`fx` :ref:`m_op_sel` :ref:`m_op_sel_hi` :ref:`clamp` + v_fma_mixlo_f16 :ref:`vdst`, :ref:`src0`::ref:`m`::ref:`fx`, :ref:`src1`::ref:`m`::ref:`fx`, :ref:`src2`::ref:`m`::ref:`fx` :ref:`m_op_sel` :ref:`m_op_sel_hi` :ref:`clamp` + v_pk_add_f16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` :ref:`neg_lo` :ref:`neg_hi` :ref:`clamp` + v_pk_add_i16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` :ref:`clamp` + v_pk_add_u16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` :ref:`clamp` + v_pk_ashrrev_i16 :ref:`vdst`, :ref:`src0`::ref:`u16x2`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` + v_pk_fma_f16 :ref:`vdst`, :ref:`src0`, :ref:`src1`, :ref:`src2` :ref:`op_sel` :ref:`op_sel_hi` :ref:`neg_lo` :ref:`neg_hi` :ref:`clamp` + v_pk_lshlrev_b16 :ref:`vdst`, :ref:`src0`::ref:`u16x2`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` + v_pk_lshrrev_b16 :ref:`vdst`, :ref:`src0`::ref:`u16x2`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` + v_pk_mad_i16 :ref:`vdst`, :ref:`src0`, :ref:`src1`, :ref:`src2` :ref:`op_sel` :ref:`op_sel_hi` :ref:`clamp` + v_pk_mad_u16 :ref:`vdst`, :ref:`src0`, :ref:`src1`, :ref:`src2` :ref:`op_sel` :ref:`op_sel_hi` :ref:`clamp` + v_pk_max_f16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` :ref:`neg_lo` :ref:`neg_hi` :ref:`clamp` + v_pk_max_i16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` + v_pk_max_u16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` + v_pk_min_f16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` :ref:`neg_lo` :ref:`neg_hi` :ref:`clamp` + v_pk_min_i16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` + v_pk_min_u16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` + v_pk_mul_f16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` :ref:`neg_lo` :ref:`neg_hi` :ref:`clamp` + v_pk_mul_lo_u16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` + v_pk_sub_i16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` :ref:`clamp` + v_pk_sub_u16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` :ref:`clamp` VOPC ----------------------- diff --git a/llvm/docs/AMDGPU/AMDGPUAsmGFX7.rst b/llvm/docs/AMDGPU/AMDGPUAsmGFX7.rst index 1933aff3bb7c4..e96862bcf0323 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`. -Introduction -============ +Overvew +======= An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document`. diff --git a/llvm/docs/AMDGPU/AMDGPUAsmGFX8.rst b/llvm/docs/AMDGPU/AMDGPUAsmGFX8.rst index a0e514e29b2c1..4dd0439c6fdb8 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`. -Introduction -============ +Overvew +======= An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document`. diff --git a/llvm/docs/AMDGPU/AMDGPUAsmGFX9.rst b/llvm/docs/AMDGPU/AMDGPUAsmGFX9.rst index 8ce056c8caf8e..7131718df7e7d 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`. -Introduction -============ +Overvew +======= An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document`. @@ -1637,9 +1642,6 @@ VOP3P **INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS** \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| - v_mad_mix_f32 :ref:`vdst`, :ref:`src0`::ref:`fx`, :ref:`src1`::ref:`fx`, :ref:`src2`::ref:`fx` :ref:`m_op_sel` :ref:`m_op_sel_hi` :ref:`clamp` - v_mad_mixhi_f16 :ref:`vdst`, :ref:`src0`::ref:`fx`, :ref:`src1`::ref:`fx`, :ref:`src2`::ref:`fx` :ref:`m_op_sel` :ref:`m_op_sel_hi` :ref:`clamp` - v_mad_mixlo_f16 :ref:`vdst`, :ref:`src0`::ref:`fx`, :ref:`src1`::ref:`fx`, :ref:`src2`::ref:`fx` :ref:`m_op_sel` :ref:`m_op_sel_hi` :ref:`clamp` v_pk_add_f16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` :ref:`neg_lo` :ref:`neg_hi` :ref:`clamp` v_pk_add_i16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` :ref:`clamp` v_pk_add_u16 :ref:`vdst`, :ref:`src0`, :ref:`src1` :ref:`op_sel` :ref:`op_sel_hi` :ref:`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 0000000000000..7832e08dbb374 --- /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`. + +Notation +======== + +Notation used in this document is explained :ref:`here`. + +Overvew +======= + +An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document`. + +Instructions +============ + + +VOP3P +----------------------- + +.. parsed-literal:: + + **INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + v_mad_mix_f32 :ref:`vdst`, :ref:`src0`::ref:`m`::ref:`fx`, :ref:`src1`::ref:`m`::ref:`fx`, :ref:`src2`::ref:`m`::ref:`fx` :ref:`m_op_sel` :ref:`m_op_sel_hi` :ref:`clamp` + v_mad_mixhi_f16 :ref:`vdst`, :ref:`src0`::ref:`m`::ref:`fx`, :ref:`src1`::ref:`m`::ref:`fx`, :ref:`src2`::ref:`m`::ref:`fx` :ref:`m_op_sel` :ref:`m_op_sel_hi` :ref:`clamp` + v_mad_mixlo_f16 :ref:`vdst`, :ref:`src0`::ref:`m`::ref:`fx`, :ref:`src1`::ref:`m`::ref:`fx`, :ref:`src2`::ref:`m`::ref:`fx` :ref:`m_op_sel` :ref:`m_op_sel_hi` :ref:`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 0000000000000..23f79fc96a43d --- /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`. + +Notation +======== + +Notation used in this document is explained :ref:`here`. + +Overvew +======= + +An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document`. + +Instructions +============ + + +VOP3P +----------------------- + +.. parsed-literal:: + + **INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + v_fma_mix_f32 :ref:`vdst`, :ref:`src0`::ref:`m`::ref:`fx`, :ref:`src1`::ref:`m`::ref:`fx`, :ref:`src2`::ref:`m`::ref:`fx` :ref:`m_op_sel` :ref:`m_op_sel_hi` :ref:`clamp` + v_fma_mixhi_f16 :ref:`vdst`, :ref:`src0`::ref:`m`::ref:`fx`, :ref:`src1`::ref:`m`::ref:`fx`, :ref:`src2`::ref:`m`::ref:`fx` :ref:`m_op_sel` :ref:`m_op_sel_hi` :ref:`clamp` + v_fma_mixlo_f16 :ref:`vdst`, :ref:`src0`::ref:`m`::ref:`fx`, :ref:`src1`::ref:`m`::ref:`fx`, :ref:`src2`::ref:`m`::ref:`fx` :ref:`m_op_sel` :ref:`m_op_sel_hi` :ref:`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 0000000000000..8119af4392eb3 --- /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`. + +Notation +======== + +Notation used in this document is explained :ref:`here`. + +Overvew +======= + +An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document`. + +Instructions +============ + + +VOP2 +----------------------- + +.. parsed-literal:: + + **INSTRUCTION** **DST** **SRC0** **SRC1** **MODIFIERS** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + v_fmac_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_fmac_f32_dpp :ref:`vdst`, :ref:`vsrc0`::ref:`m`, :ref:`vsrc1`::ref:`m` :ref:`dpp_ctrl` :ref:`row_mask` :ref:`bank_mask` :ref:`bound_ctrl` + v_xnor_b32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_xnor_b32_dpp :ref:`vdst`, :ref:`vsrc0`, :ref:`vsrc1` :ref:`dpp_ctrl` :ref:`row_mask` :ref:`bank_mask` :ref:`bound_ctrl` + v_xnor_b32_sdwa :ref:`vdst`, :ref:`src0`::ref:`m`, :ref:`vsrc1`::ref:`m` :ref:`dst_sel` :ref:`dst_unused` :ref:`src0_sel` :ref:`src1_sel` + +VOP3 +----------------------- + +.. parsed-literal:: + + **INSTRUCTION** **DST** **SRC0** **SRC1** **MODIFIERS** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + v_fmac_f32_e64 :ref:`vdst`, :ref:`src0`::ref:`m`, :ref:`src1`::ref:`m` :ref:`clamp` :ref:`omod` + v_xnor_b32_e64 :ref:`vdst`, :ref:`src0`, :ref:`src1` + +VOP3P +----------------------- + +.. parsed-literal:: + + **INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + v_dot2_f32_f16 :ref:`vdst`, :ref:`src0`::ref:`f16x2`, :ref:`src1`::ref:`f16x2`, :ref:`src2`::ref:`f32` :ref:`neg_lo` :ref:`neg_hi` :ref:`clamp` + v_dot2_i32_i16 :ref:`vdst`, :ref:`src0`::ref:`i16x2`, :ref:`src1`::ref:`i16x2`, :ref:`src2`::ref:`i32` :ref:`clamp` + v_dot2_u32_u16 :ref:`vdst`, :ref:`src0`::ref:`u16x2`, :ref:`src1`::ref:`u16x2`, :ref:`src2`::ref:`u32` :ref:`clamp` + v_dot4_i32_i8 :ref:`vdst`, :ref:`src0`::ref:`i8x2`, :ref:`src1`::ref:`i8x2`, :ref:`src2`::ref:`i32` :ref:`clamp` + v_dot4_u32_u8 :ref:`vdst`, :ref:`src0`::ref:`u8x2`, :ref:`src1`::ref:`u8x2`, :ref:`src2`::ref:`u32` :ref:`clamp` + v_dot8_i32_i4 :ref:`vdst`, :ref:`src0`::ref:`i4x2`, :ref:`src1`::ref:`i4x2`, :ref:`src2`::ref:`i32` :ref:`clamp` + v_dot8_u32_u4 :ref:`vdst`, :ref:`src0`::ref:`u4x2`, :ref:`src1`::ref:`u4x2`, :ref:`src2`::ref:`u32` :ref:`clamp` + v_fma_mix_f32 :ref:`vdst`, :ref:`src0`::ref:`m`::ref:`fx`, :ref:`src1`::ref:`m`::ref:`fx`, :ref:`src2`::ref:`m`::ref:`fx` :ref:`m_op_sel` :ref:`m_op_sel_hi` :ref:`clamp` + v_fma_mixhi_f16 :ref:`vdst`, :ref:`src0`::ref:`m`::ref:`fx`, :ref:`src1`::ref:`m`::ref:`fx`, :ref:`src2`::ref:`m`::ref:`fx` :ref:`m_op_sel` :ref:`m_op_sel_hi` :ref:`clamp` + v_fma_mixlo_f16 :ref:`vdst`, :ref:`src0`::ref:`m`::ref:`fx`, :ref:`src1`::ref:`m`::ref:`fx`, :ref:`src2`::ref:`m`::ref:`fx` :ref:`m_op_sel` :ref:`m_op_sel_hi` :ref:`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 0000000000000..7e90837cdf0b1 --- /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`. + +Notation +======== + +Notation used in this document is explained :ref:`here`. + +Overvew +======= + +An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document`. + +Instructions +============ + + +FLAT +----------------------- + +.. parsed-literal:: + + **INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + global_atomic_add_f32 :ref:`vdst`::ref:`opt`, :ref:`vaddr`, :ref:`vdata`, :ref:`saddr` :ref:`offset13s` + global_atomic_pk_add_f16 :ref:`vdst`::ref:`opt`::ref:`f16x2`, :ref:`vaddr`, :ref:`vdata`::ref:`f16x2`, :ref:`saddr` :ref:`offset13s` + +MUBUF +----------------------- + +.. parsed-literal:: + + **INSTRUCTION** **SRC0** **SRC1** **SRC2** **SRC3** **MODIFIERS** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + buffer_atomic_add_f32 :ref:`vdata`::ref:`dst`, :ref:`vaddr`, :ref:`srsrc`, :ref:`soffset` :ref:`idxen` :ref:`offen` :ref:`offset12` :ref:`slc` + buffer_atomic_pk_add_f16 :ref:`vdata`::ref:`dst`::ref:`f16x2`, :ref:`vaddr`, :ref:`srsrc`, :ref:`soffset` :ref:`idxen` :ref:`offen` :ref:`offset12` :ref:`slc` + +VOP2 +----------------------- + +.. parsed-literal:: + + **INSTRUCTION** **DST** **SRC0** **SRC1** **MODIFIERS** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + v_dot2c_f32_f16 :ref:`vdst`, :ref:`src0`::ref:`f16x2`, :ref:`vsrc1`::ref:`f16x2` + v_dot2c_f32_f16_dpp :ref:`vdst`, :ref:`vsrc0`::ref:`f16x2`, :ref:`vsrc1`::ref:`f16x2` :ref:`dpp_ctrl` :ref:`row_mask` :ref:`bank_mask` :ref:`bound_ctrl` + v_dot2c_i32_i16 :ref:`vdst`, :ref:`src0`::ref:`i16x2`, :ref:`vsrc1`::ref:`i16x2` + v_dot2c_i32_i16_dpp :ref:`vdst`, :ref:`vsrc0`::ref:`i16x2`, :ref:`vsrc1`::ref:`i16x2` :ref:`dpp_ctrl` :ref:`row_mask` :ref:`bank_mask` :ref:`bound_ctrl` + v_dot4c_i32_i8 :ref:`vdst`, :ref:`src0`::ref:`i8x4`, :ref:`vsrc1`::ref:`i8x4` + v_dot4c_i32_i8_dpp :ref:`vdst`, :ref:`vsrc0`::ref:`i8x4`, :ref:`vsrc1`::ref:`i8x4` :ref:`dpp_ctrl` :ref:`row_mask` :ref:`bank_mask` :ref:`bound_ctrl` + v_dot8c_i32_i4 :ref:`vdst`, :ref:`src0`::ref:`i4x8`, :ref:`vsrc1`::ref:`i4x8` + v_dot8c_i32_i4_dpp :ref:`vdst`, :ref:`vsrc0`::ref:`i4x8`, :ref:`vsrc1`::ref:`i4x8` :ref:`dpp_ctrl` :ref:`row_mask` :ref:`bank_mask` :ref:`bound_ctrl` + v_fmac_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_fmac_f32_dpp :ref:`vdst`, :ref:`vsrc0`::ref:`m`, :ref:`vsrc1`::ref:`m` :ref:`dpp_ctrl` :ref:`row_mask` :ref:`bank_mask` :ref:`bound_ctrl` + v_pk_fmac_f16 :ref:`vdst`::ref:`f16x2`, :ref:`src0`::ref:`f16x2`, :ref:`vsrc1`::ref:`f16x2` + v_xnor_b32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1` + v_xnor_b32_dpp :ref:`vdst`, :ref:`vsrc0`, :ref:`vsrc1` :ref:`dpp_ctrl` :ref:`row_mask` :ref:`bank_mask` :ref:`bound_ctrl` + v_xnor_b32_sdwa :ref:`vdst`, :ref:`src0`::ref:`m`, :ref:`vsrc1`::ref:`m` :ref:`dst_sel` :ref:`dst_unused` :ref:`src0_sel` :ref:`src1_sel` + +VOP3 +----------------------- + +.. parsed-literal:: + + **INSTRUCTION** **DST** **SRC0** **SRC1** **MODIFIERS** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + v_fmac_f32_e64 :ref:`vdst`, :ref:`src0`::ref:`m`, :ref:`src1`::ref:`m` :ref:`clamp` :ref:`omod` + v_xnor_b32_e64 :ref:`vdst`, :ref:`src0`, :ref:`src1` + +VOP3P +----------------------- + +.. parsed-literal:: + + **INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + v_accvgpr_read_b32 :ref:`vdst`, :ref:`asrc` + v_accvgpr_write_b32 :ref:`adst`, :ref:`src` + v_dot2_f32_f16 :ref:`vdst`, :ref:`src0`::ref:`f16x2`, :ref:`src1`::ref:`f16x2`, :ref:`src2`::ref:`f32` :ref:`neg_lo` :ref:`neg_hi` :ref:`clamp` + v_dot2_i32_i16 :ref:`vdst`, :ref:`src0`::ref:`i16x2`, :ref:`src1`::ref:`i16x2`, :ref:`src2`::ref:`i32` :ref:`clamp` + v_dot2_u32_u16 :ref:`vdst`, :ref:`src0`::ref:`u16x2`, :ref:`src1`::ref:`u16x2`, :ref:`src2`::ref:`u32` :ref:`clamp` + v_dot4_i32_i8 :ref:`vdst`, :ref:`src0`::ref:`i8x2`, :ref:`src1`::ref:`i8x2`, :ref:`src2`::ref:`i32` :ref:`clamp` + v_dot4_u32_u8 :ref:`vdst`, :ref:`src0`::ref:`u8x2`, :ref:`src1`::ref:`u8x2`, :ref:`src2`::ref:`u32` :ref:`clamp` + v_dot8_i32_i4 :ref:`vdst`, :ref:`src0`::ref:`i4x2`, :ref:`src1`::ref:`i4x2`, :ref:`src2`::ref:`i32` :ref:`clamp` + v_dot8_u32_u4 :ref:`vdst`, :ref:`src0`::ref:`u4x2`, :ref:`src1`::ref:`u4x2`, :ref:`src2`::ref:`u32` :ref:`clamp` + v_fma_mix_f32 :ref:`vdst`, :ref:`src0`::ref:`m`::ref:`fx`, :ref:`src1`::ref:`m`::ref:`fx`, :ref:`src2`::ref:`m`::ref:`fx` :ref:`m_op_sel` :ref:`m_op_sel_hi` :ref:`clamp` + v_fma_mixhi_f16 :ref:`vdst`, :ref:`src0`::ref:`m`::ref:`fx`, :ref:`src1`::ref:`m`::ref:`fx`, :ref:`src2`::ref:`m`::ref:`fx` :ref:`m_op_sel` :ref:`m_op_sel_hi` :ref:`clamp` + v_fma_mixlo_f16 :ref:`vdst`, :ref:`src0`::ref:`m`::ref:`fx`, :ref:`src1`::ref:`m`::ref:`fx`, :ref:`src2`::ref:`m`::ref:`fx` :ref:`m_op_sel` :ref:`m_op_sel_hi` :ref:`clamp` + v_mfma_f32_16x16x16f16 :ref:`adst`::ref:`f32x4`, :ref:`vasrc0`::ref:`f16x4`, :ref:`vasrc1`::ref:`f16x4`, :ref:`asrc2`::ref:`f32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_16x16x1f32 :ref:`adst`::ref:`f32x16`, :ref:`vasrc0`::ref:`f32`, :ref:`vasrc1`::ref:`f32`, :ref:`asrc2`::ref:`f32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_16x16x2bf16 :ref:`adst`::ref:`f32x16`, :ref:`vasrc0`::ref:`i16x2`, :ref:`vasrc1`::ref:`i16x2`, :ref:`asrc2`::ref:`f32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_16x16x4f16 :ref:`adst`::ref:`f32x16`, :ref:`vasrc0`::ref:`f16x4`, :ref:`vasrc1`::ref:`f16x4`, :ref:`asrc2`::ref:`f32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_16x16x4f32 :ref:`adst`::ref:`f32x4`, :ref:`vasrc0`::ref:`f32`, :ref:`vasrc1`::ref:`f32`, :ref:`asrc2`::ref:`f32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_16x16x8bf16 :ref:`adst`::ref:`f32x4`, :ref:`vasrc0`::ref:`i16x2`, :ref:`vasrc1`::ref:`i16x2`, :ref:`asrc2`::ref:`f32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_32x32x1f32 :ref:`adst`::ref:`f32x32`, :ref:`vasrc0`::ref:`f32`, :ref:`vasrc1`::ref:`f32`, :ref:`asrc2`::ref:`f32x32` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_32x32x2bf16 :ref:`adst`::ref:`f32x32`, :ref:`vasrc0`::ref:`i16x2`, :ref:`vasrc1`::ref:`i16x2`, :ref:`asrc2`::ref:`f32x32` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_32x32x2f32 :ref:`adst`::ref:`f32x16`, :ref:`vasrc0`::ref:`f32`, :ref:`vasrc1`::ref:`f32`, :ref:`asrc2`::ref:`f32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_32x32x4bf16 :ref:`adst`::ref:`f32x16`, :ref:`vasrc0`::ref:`i16x2`, :ref:`vasrc1`::ref:`i16x2`, :ref:`asrc2`::ref:`f32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_32x32x4f16 :ref:`adst`::ref:`f32x32`, :ref:`vasrc0`::ref:`f16x4`, :ref:`vasrc1`::ref:`f16x4`, :ref:`asrc2`::ref:`f32x32` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_32x32x8f16 :ref:`adst`::ref:`f32x16`, :ref:`vasrc0`::ref:`f16x4`, :ref:`vasrc1`::ref:`f16x4`, :ref:`asrc2`::ref:`f32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_4x4x1f32 :ref:`adst`::ref:`f32x4`, :ref:`vasrc0`::ref:`f32`, :ref:`vasrc1`::ref:`f32`, :ref:`asrc2`::ref:`f32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_4x4x2bf16 :ref:`adst`::ref:`f32x4`, :ref:`vasrc0`::ref:`i16x2`, :ref:`vasrc1`::ref:`i16x2`, :ref:`asrc2`::ref:`f32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_4x4x4f16 :ref:`adst`::ref:`f32x4`, :ref:`vasrc0`::ref:`f16x4`, :ref:`vasrc1`::ref:`f16x4`, :ref:`asrc2`::ref:`f32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_i32_16x16x16i8 :ref:`adst`::ref:`i32x4`, :ref:`vasrc0`::ref:`i32`, :ref:`vasrc1`::ref:`i32`, :ref:`asrc2`::ref:`i32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_i32_16x16x4i8 :ref:`adst`::ref:`i32x16`, :ref:`vasrc0`::ref:`i32`, :ref:`vasrc1`::ref:`i32`, :ref:`asrc2`::ref:`i32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_i32_32x32x4i8 :ref:`adst`::ref:`i32x32`, :ref:`vasrc0`::ref:`i32`, :ref:`vasrc1`::ref:`i32`, :ref:`asrc2`::ref:`i32x32` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_i32_32x32x8i8 :ref:`adst`::ref:`i32x16`, :ref:`vasrc0`::ref:`i32`, :ref:`vasrc1`::ref:`i32`, :ref:`asrc2`::ref:`i32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_i32_4x4x4i8 :ref:`adst`::ref:`i32x4`, :ref:`vasrc0`::ref:`i32`, :ref:`vasrc1`::ref:`i32`, :ref:`asrc2`::ref:`i32x4` :ref:`cbsz` :ref:`abid` :ref:`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 similarity index 94% rename from llvm/docs/AMDGPU/gfx9_mad_type_dev.rst rename to llvm/docs/AMDGPU/gfx900_mad_type_dev.rst index 0602f08217627..50c403ef582d7 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 0000000000000..fa0311a9aadfc --- /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` and :ref:`neg`. + diff --git a/llvm/docs/AMDGPU/gfx900_src32_0.rst b/llvm/docs/AMDGPU/gfx900_src32_0.rst new file mode 100644 index 0000000000000..9e2e607a0565e --- /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`, :ref:`s`, :ref:`flat_scratch`, :ref:`vcc`, :ref:`ttmp`, :ref:`m0`, :ref:`exec`, :ref:`vccz`, :ref:`execz`, :ref:`scc`, :ref:`lds_direct`, :ref:`constant` diff --git a/llvm/docs/AMDGPU/gfx900_src32_1.rst b/llvm/docs/AMDGPU/gfx900_src32_1.rst new file mode 100644 index 0000000000000..82dd50b998d67 --- /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`, :ref:`s`, :ref:`flat_scratch`, :ref:`vcc`, :ref:`ttmp`, :ref:`m0`, :ref:`exec`, :ref:`vccz`, :ref:`execz`, :ref:`scc`, :ref:`constant` diff --git a/llvm/docs/AMDGPU/gfx900_vdst32_0.rst b/llvm/docs/AMDGPU/gfx900_vdst32_0.rst new file mode 100644 index 0000000000000..0ae76da5315fe --- /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` 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 0000000000000..6fc3c310880d3 --- /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`. +* Location of 16-bit operand is controlled by :ref:`m_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 0000000000000..814bccdf3a471 --- /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` and :ref:`neg`. + diff --git a/llvm/docs/AMDGPU/gfx904_src32_0.rst b/llvm/docs/AMDGPU/gfx904_src32_0.rst new file mode 100644 index 0000000000000..a2930d4c4f900 --- /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`, :ref:`s`, :ref:`flat_scratch`, :ref:`vcc`, :ref:`ttmp`, :ref:`m0`, :ref:`exec`, :ref:`vccz`, :ref:`execz`, :ref:`scc`, :ref:`lds_direct`, :ref:`constant` diff --git a/llvm/docs/AMDGPU/gfx904_src32_1.rst b/llvm/docs/AMDGPU/gfx904_src32_1.rst new file mode 100644 index 0000000000000..23127c6c99b85 --- /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`, :ref:`s`, :ref:`flat_scratch`, :ref:`vcc`, :ref:`ttmp`, :ref:`m0`, :ref:`exec`, :ref:`vccz`, :ref:`execz`, :ref:`scc`, :ref:`constant` diff --git a/llvm/docs/AMDGPU/gfx904_vdst32_0.rst b/llvm/docs/AMDGPU/gfx904_vdst32_0.rst new file mode 100644 index 0000000000000..6eb8bdf5b506d --- /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` 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 0000000000000..eb048b0d2d1f9 --- /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`. +* Location of 16-bit operand is controlled by :ref:`m_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 0000000000000..88fdc6f6490eb --- /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` and :ref:`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 0000000000000..411251baeb275 --- /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`. + 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 0000000000000..ffd6429640d1a --- /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` and :ref:`neg`. + diff --git a/llvm/docs/AMDGPU/gfx906_src32_0.rst b/llvm/docs/AMDGPU/gfx906_src32_0.rst new file mode 100644 index 0000000000000..f547beaec23b7 --- /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`, :ref:`s`, :ref:`flat_scratch`, :ref:`vcc`, :ref:`ttmp`, :ref:`m0`, :ref:`exec`, :ref:`vccz`, :ref:`execz`, :ref:`scc`, :ref:`lds_direct`, :ref:`constant`, :ref:`literal` diff --git a/llvm/docs/AMDGPU/gfx906_src32_1.rst b/llvm/docs/AMDGPU/gfx906_src32_1.rst new file mode 100644 index 0000000000000..1fc90e163b1d8 --- /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`, :ref:`s`, :ref:`flat_scratch`, :ref:`vcc`, :ref:`ttmp`, :ref:`m0`, :ref:`exec`, :ref:`vccz`, :ref:`execz`, :ref:`scc`, :ref:`lds_direct`, :ref:`constant` diff --git a/llvm/docs/AMDGPU/gfx906_src32_2.rst b/llvm/docs/AMDGPU/gfx906_src32_2.rst new file mode 100644 index 0000000000000..6c3bb387c4ed0 --- /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`, :ref:`s`, :ref:`flat_scratch`, :ref:`vcc`, :ref:`ttmp`, :ref:`m0`, :ref:`exec`, :ref:`vccz`, :ref:`execz`, :ref:`scc`, :ref:`constant` diff --git a/llvm/docs/AMDGPU/gfx906_type_dev.rst b/llvm/docs/AMDGPU/gfx906_type_dev.rst new file mode 100644 index 0000000000000..02f8d98a3dc4f --- /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`. 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 0000000000000..faf0a2d2f9dd1 --- /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` diff --git a/llvm/docs/AMDGPU/gfx906_vsrc32_0.rst b/llvm/docs/AMDGPU/gfx906_vsrc32_0.rst new file mode 100644 index 0000000000000..ed350d997eb70 --- /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` diff --git a/llvm/docs/AMDGPU/gfx908_addr_buf.rst b/llvm/docs/AMDGPU/gfx908_addr_buf.rst new file mode 100644 index 0000000000000..9dc593be69eb8 --- /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` and :ref:`idxen`: + +* If only :ref:`idxen` is specified, this operand supplies an index. Size is 1 dword. +* If only :ref:`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`. + +*Operands:* :ref:`v`, :ref:`off` diff --git a/llvm/docs/AMDGPU/gfx908_adst1024_0.rst b/llvm/docs/AMDGPU/gfx908_adst1024_0.rst new file mode 100644 index 0000000000000..3dfe45ede469e --- /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` diff --git a/llvm/docs/AMDGPU/gfx908_adst128_0.rst b/llvm/docs/AMDGPU/gfx908_adst128_0.rst new file mode 100644 index 0000000000000..23c7f9db057b6 --- /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` diff --git a/llvm/docs/AMDGPU/gfx908_adst32_0.rst b/llvm/docs/AMDGPU/gfx908_adst32_0.rst new file mode 100644 index 0000000000000..80546fa1c8326 --- /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` diff --git a/llvm/docs/AMDGPU/gfx908_adst512_0.rst b/llvm/docs/AMDGPU/gfx908_adst512_0.rst new file mode 100644 index 0000000000000..446eb02bf0355 --- /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` diff --git a/llvm/docs/AMDGPU/gfx908_asrc1024_0.rst b/llvm/docs/AMDGPU/gfx908_asrc1024_0.rst new file mode 100644 index 0000000000000..00e1fc7f0f6e0 --- /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` diff --git a/llvm/docs/AMDGPU/gfx908_asrc128_0.rst b/llvm/docs/AMDGPU/gfx908_asrc128_0.rst new file mode 100644 index 0000000000000..91af6b7357bb7 --- /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` diff --git a/llvm/docs/AMDGPU/gfx908_asrc32_0.rst b/llvm/docs/AMDGPU/gfx908_asrc32_0.rst new file mode 100644 index 0000000000000..8e46317254df7 --- /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` diff --git a/llvm/docs/AMDGPU/gfx908_asrc512_0.rst b/llvm/docs/AMDGPU/gfx908_asrc512_0.rst new file mode 100644 index 0000000000000..59e22040c2226 --- /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` 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 0000000000000..03163f55fe541 --- /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` is specified, gets the memory value before the operation. + +*Size:* 1 dword by default. :ref:`tfe` adds 1 dword if specified. + +*Operands:* :ref:`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 0000000000000..045058314cbb6 --- /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` is specified. + +*Size:* 1 dword. + +*Operands:* :ref:`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 0000000000000..bbc3ea9dde555 --- /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`. +* Location of 16-bit operand is controlled by :ref:`m_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 0000000000000..028a209ebe5fc --- /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` and :ref:`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 0000000000000..1afede5e74742 --- /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`. + 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 0000000000000..07bec0c687110 --- /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` and :ref:`neg`. + diff --git a/llvm/docs/AMDGPU/gfx908_offset_buf.rst b/llvm/docs/AMDGPU/gfx908_offset_buf.rst new file mode 100644 index 0000000000000..6f3d6d2b96c17 --- /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`, :ref:`flat_scratch`, :ref:`vcc`, :ref:`ttmp`, :ref:`m0`, :ref:`exec`, :ref:`vccz`, :ref:`execz`, :ref:`scc`, :ref:`constant` diff --git a/llvm/docs/AMDGPU/gfx908_opt.rst b/llvm/docs/AMDGPU/gfx908_opt.rst new file mode 100644 index 0000000000000..efd5a79c8a7d7 --- /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` is specified. + diff --git a/llvm/docs/AMDGPU/gfx908_ret.rst b/llvm/docs/AMDGPU/gfx908_ret.rst new file mode 100644 index 0000000000000..dc8a91c5b4cac --- /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` 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 0000000000000..832adfe0763ea --- /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`, :ref:`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 0000000000000..a0e450639e0ea --- /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` if not used. + +See :ref:`vaddr` for description of available addressing modes. + +*Size:* 2 dwords. + +*Operands:* :ref:`exec`, :ref:`off` diff --git a/llvm/docs/AMDGPU/gfx908_src32_0.rst b/llvm/docs/AMDGPU/gfx908_src32_0.rst new file mode 100644 index 0000000000000..a6bf415d402e4 --- /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`, :ref:`s`, :ref:`flat_scratch`, :ref:`vcc`, :ref:`ttmp`, :ref:`m0`, :ref:`exec`, :ref:`vccz`, :ref:`execz`, :ref:`scc`, :ref:`lds_direct`, :ref:`constant`, :ref:`literal` diff --git a/llvm/docs/AMDGPU/gfx908_src32_1.rst b/llvm/docs/AMDGPU/gfx908_src32_1.rst new file mode 100644 index 0000000000000..275e4d796b962 --- /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`, :ref:`s`, :ref:`flat_scratch`, :ref:`vcc`, :ref:`ttmp`, :ref:`m0`, :ref:`exec`, :ref:`vccz`, :ref:`execz`, :ref:`scc`, :ref:`lds_direct`, :ref:`constant` diff --git a/llvm/docs/AMDGPU/gfx908_src32_2.rst b/llvm/docs/AMDGPU/gfx908_src32_2.rst new file mode 100644 index 0000000000000..21d0f201f2629 --- /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`, :ref:`s`, :ref:`flat_scratch`, :ref:`vcc`, :ref:`ttmp`, :ref:`m0`, :ref:`exec`, :ref:`vccz`, :ref:`execz`, :ref:`scc`, :ref:`constant` diff --git a/llvm/docs/AMDGPU/gfx908_src32_3.rst b/llvm/docs/AMDGPU/gfx908_src32_3.rst new file mode 100644 index 0000000000000..0e9573d45da46 --- /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`, :ref:`vccz`, :ref:`execz`, :ref:`scc`, :ref:`constant` diff --git a/llvm/docs/AMDGPU/gfx908_type_dev.rst b/llvm/docs/AMDGPU/gfx908_type_dev.rst new file mode 100644 index 0000000000000..8d8da92eb1cf9 --- /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`. 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 0000000000000..7b53cb3c72267 --- /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` + :ref:`offset13s`. :ref:`vaddr` is a 64-bit address. This mode is indicated by :ref:`saddr` set to :ref:`off`. +* Address = :ref:`saddr` + :ref:`vaddr` + :ref:`offset13s`. :ref:`vaddr` is a 32-bit offset. This mode is used when :ref:`saddr` is not :ref:`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` diff --git a/llvm/docs/AMDGPU/gfx908_vasrc32_0.rst b/llvm/docs/AMDGPU/gfx908_vasrc32_0.rst new file mode 100644 index 0000000000000..52aaaadc5808f --- /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`, :ref:`a` diff --git a/llvm/docs/AMDGPU/gfx908_vasrc64_0.rst b/llvm/docs/AMDGPU/gfx908_vasrc64_0.rst new file mode 100644 index 0000000000000..bef80d9637a9d --- /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`, :ref:`a` diff --git a/llvm/docs/AMDGPU/gfx908_vdata32_0.rst b/llvm/docs/AMDGPU/gfx908_vdata32_0.rst new file mode 100644 index 0000000000000..7571818881d85 --- /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` diff --git a/llvm/docs/AMDGPU/gfx908_vdst32_0.rst b/llvm/docs/AMDGPU/gfx908_vdst32_0.rst new file mode 100644 index 0000000000000..ed2214a8fbed2 --- /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` diff --git a/llvm/docs/AMDGPU/gfx908_vsrc32_0.rst b/llvm/docs/AMDGPU/gfx908_vsrc32_0.rst new file mode 100644 index 0000000000000..c990b504515b4 --- /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` diff --git a/llvm/docs/AMDGPUModifierSyntax.rst b/llvm/docs/AMDGPUModifierSyntax.rst index 526016d4f4f81..5e29004b4d6ed 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`. GFX9 and GFX10 only. @@ -1494,8 +1494,8 @@ See a description :ref:`here`. 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`. + +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` or +an :ref:`absolute expression`. + +.. _amdgpu_synid_abid: + +abid +~~~~ + + =============================== ================================================================== + Syntax Description + =============================== ================================================================== + abid:[{0..15}] TBD + =============================== ================================================================== + +Note: numeric value may be specified as either +an :ref:`integer number` or +an :ref:`absolute expression`. + +.. _amdgpu_synid_blgp: + +blgp +~~~~ + + =============================== ================================================================== + Syntax Description + =============================== ================================================================== + blgp:[{0..7}] TBD + =============================== ================================================================== + +Note: numeric value may be specified as either +an :ref:`integer number` or +an :ref:`absolute expression`. + diff --git a/llvm/docs/AMDGPUOperandSyntax.rst b/llvm/docs/AMDGPUOperandSyntax.rst index c20da0047296b..842f49fbf735a 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**\ **acc**\ A single 32-bit *accumulator* register. + + *N* must be a decimal + :ref:`integer number`. + **a[**\ \ **]** **acc[**\ \ **]** A single 32-bit *accumulator* register. + + *N* may be specified as an + :ref:`integer number` + or an :ref:`absolute expression`. + **a[**\ :\ **]** **acc[**\ :\ **]** A sequence of (\ *K-N+1*\ ) *accumulator* registers. + + *N* and *K* may be specified as + :ref:`integer numbers` + or :ref:`absolute expressions`. + **[a**\ , \ **a**\ , ... **a**\ \ **]** **[acc**\ , \ **acc**\ , ... **acc**\ \ **]** A sequence of (\ *K-N+1*\ ) *accumulator* registers. + + Register indices must be specified as decimal + :ref:`integer numbers`. + =================================================== ========================================================= ==================================================================== + +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 ee165f0fc1057..d9c64fd83a199 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`: | ``<``\ *opcode*\ ``> <``\ *operand0*\ ``>, <``\ *operand1*\ ``>,... <``\ *modifier0*\ ``> <``\ *modifier1*\ ``>...`` -:doc:`Operands` are normally comma-separated while +:doc:`Operands` are comma-separated while :doc:`modifiers` 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`, :doc:`GFX8`, -:doc:`GFX9`, and :doc:`GFX10`. +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` \- + :doc:`GFX8` \- + :doc:`GFX9` :doc:`gfx900` + + :doc:`gfx902` + + :doc:`gfx904` + + :doc:`gfx906` + + :doc:`gfx908` + + :doc:`gfx909` + + :doc:`GFX10` gfx1011 + + gfx1012 + ==================================== ====================================== For more information about instructions, their semantics and supported combinations of operands, refer to one of instruction set architecture manuals