From 3a4d9b6a6856c5f8cb006f61e93bebb01f956d62 Mon Sep 17 00:00:00 2001 From: Dmitry Preobrazhensky Date: Fri, 1 Jul 2022 12:34:59 +0300 Subject: [PATCH] [AMDGPU][GFX908][DOC][NFC] Update assembler syntax description Summary of changes: - Remove dst for global_atomic_add_f32, global_atomic_pk_add_f16. - Make vdata input-only for buffer_atomic_add_f32, buffer_atomic_pk_add_f16. - Other minor improvements. --- llvm/docs/AMDGPU/AMDGPUAsmGFX908.rst | 173 ++++++++++----------- llvm/docs/AMDGPU/gfx908_dst.rst | 13 -- .../AMDGPU/{gfx908_m_1.rst => gfx908_m_254bcb.rst} | 2 +- .../AMDGPU/{gfx908_m.rst => gfx908_m_f5d306.rst} | 2 +- llvm/docs/AMDGPU/gfx908_opt.rst | 13 -- llvm/docs/AMDGPU/gfx908_saddr.rst | 2 +- .../{gfx908_src_2.rst => gfx908_src_4e78e6.rst} | 2 +- .../{gfx908_src_3.rst => gfx908_src_58d119.rst} | 2 +- .../{gfx908_src.rst => gfx908_src_73ab34.rst} | 2 +- .../{gfx908_src_4.rst => gfx908_src_955b45.rst} | 2 +- .../{gfx908_src_1.rst => gfx908_src_d578c4.rst} | 2 +- .../{gfx908_src_5.rst => gfx908_src_d95796.rst} | 2 +- llvm/docs/AMDGPU/gfx908_vaddr.rst | 20 --- llvm/docs/AMDGPU/gfx908_vaddr_0212e3.rst | 20 +++ ...{gfx908_vaddr_1.rst => gfx908_vaddr_b73dc0.rst} | 2 +- llvm/docs/AMDGPU/gfx908_vdata_1.rst | 21 --- .../{gfx908_vdata.rst => gfx908_vdata_6802ce.rst} | 2 +- llvm/docs/AMDGPU/gfx908_vdata_fe1edf.rst | 17 ++ llvm/docs/AMDGPU/gfx908_vdst.rst | 19 --- .../{gfx908_vdst_4.rst => gfx908_vdst_0c4ef8.rst} | 2 +- .../{gfx908_vdst_5.rst => gfx908_vdst_2c8d1e.rst} | 2 +- .../{gfx908_vdst_2.rst => gfx908_vdst_78dd0a.rst} | 2 +- .../{gfx908_vdst_1.rst => gfx908_vdst_89680f.rst} | 2 +- .../{gfx908_vdst_3.rst => gfx908_vdst_bcee7a.rst} | 2 +- .../{gfx908_vsrc_6.rst => gfx908_vsrc_036abe.rst} | 2 +- .../{gfx908_vsrc_1.rst => gfx908_vsrc_1027ca.rst} | 2 +- .../{gfx908_vsrc_5.rst => gfx908_vsrc_2d4632.rst} | 2 +- .../{gfx908_vsrc.rst => gfx908_vsrc_6802ce.rst} | 2 +- .../{gfx908_vsrc_2.rst => gfx908_vsrc_9ad749.rst} | 2 +- .../{gfx908_vsrc_4.rst => gfx908_vsrc_be4895.rst} | 2 +- .../{gfx908_vsrc_3.rst => gfx908_vsrc_f3d248.rst} | 2 +- 31 files changed, 145 insertions(+), 197 deletions(-) delete mode 100644 llvm/docs/AMDGPU/gfx908_dst.rst rename llvm/docs/AMDGPU/{gfx908_m_1.rst => gfx908_m_254bcb.rst} (91%) rename llvm/docs/AMDGPU/{gfx908_m.rst => gfx908_m_f5d306.rst} (92%) delete mode 100644 llvm/docs/AMDGPU/gfx908_opt.rst rename llvm/docs/AMDGPU/{gfx908_src_2.rst => gfx908_src_4e78e6.rst} (95%) rename llvm/docs/AMDGPU/{gfx908_src_3.rst => gfx908_src_58d119.rst} (92%) rename llvm/docs/AMDGPU/{gfx908_src.rst => gfx908_src_73ab34.rst} (95%) rename llvm/docs/AMDGPU/{gfx908_src_4.rst => gfx908_src_955b45.rst} (95%) rename llvm/docs/AMDGPU/{gfx908_src_1.rst => gfx908_src_d578c4.rst} (95%) rename llvm/docs/AMDGPU/{gfx908_src_5.rst => gfx908_src_d95796.rst} (95%) delete mode 100644 llvm/docs/AMDGPU/gfx908_vaddr.rst create mode 100644 llvm/docs/AMDGPU/gfx908_vaddr_0212e3.rst rename llvm/docs/AMDGPU/{gfx908_vaddr_1.rst => gfx908_vaddr_b73dc0.rst} (96%) delete mode 100644 llvm/docs/AMDGPU/gfx908_vdata_1.rst rename llvm/docs/AMDGPU/{gfx908_vdata.rst => gfx908_vdata_6802ce.rst} (90%) create mode 100644 llvm/docs/AMDGPU/gfx908_vdata_fe1edf.rst delete mode 100644 llvm/docs/AMDGPU/gfx908_vdst.rst rename llvm/docs/AMDGPU/{gfx908_vdst_4.rst => gfx908_vdst_0c4ef8.rst} (90%) rename llvm/docs/AMDGPU/{gfx908_vdst_5.rst => gfx908_vdst_2c8d1e.rst} (90%) rename llvm/docs/AMDGPU/{gfx908_vdst_2.rst => gfx908_vdst_78dd0a.rst} (90%) rename llvm/docs/AMDGPU/{gfx908_vdst_1.rst => gfx908_vdst_89680f.rst} (90%) rename llvm/docs/AMDGPU/{gfx908_vdst_3.rst => gfx908_vdst_bcee7a.rst} (90%) rename llvm/docs/AMDGPU/{gfx908_vsrc_6.rst => gfx908_vsrc_036abe.rst} (90%) rename llvm/docs/AMDGPU/{gfx908_vsrc_1.rst => gfx908_vsrc_1027ca.rst} (90%) rename llvm/docs/AMDGPU/{gfx908_vsrc_5.rst => gfx908_vsrc_2d4632.rst} (90%) rename llvm/docs/AMDGPU/{gfx908_vsrc.rst => gfx908_vsrc_6802ce.rst} (90%) rename llvm/docs/AMDGPU/{gfx908_vsrc_2.rst => gfx908_vsrc_9ad749.rst} (91%) rename llvm/docs/AMDGPU/{gfx908_vsrc_4.rst => gfx908_vsrc_be4895.rst} (91%) rename llvm/docs/AMDGPU/{gfx908_vsrc_3.rst => gfx908_vsrc_f3d248.rst} (90%) diff --git a/llvm/docs/AMDGPU/AMDGPUAsmGFX908.rst b/llvm/docs/AMDGPU/AMDGPUAsmGFX908.rst index 7185a2d..be0ae03 100644 --- a/llvm/docs/AMDGPU/AMDGPUAsmGFX908.rst +++ b/llvm/docs/AMDGPU/AMDGPUAsmGFX908.rst @@ -34,131 +34,128 @@ 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` :ref:`slc` - global_atomic_pk_add_f16 :ref:`vdst`::ref:`opt`, :ref:`vaddr`, :ref:`vdata`, :ref:`saddr` :ref:`offset13s` :ref:`slc` + **INSTRUCTION** **SRC0** **SRC1** **SRC2** **MODIFIERS** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + global_atomic_add_f32 :ref:`vaddr`, :ref:`vdata`, :ref:`saddr` :ref:`offset13s` :ref:`slc` + global_atomic_pk_add_f16 :ref:`vaddr`, :ref:`vdata`, :ref:`saddr` :ref:`offset13s` :ref:`slc` 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:`vaddr`, :ref:`srsrc`, :ref:`soffset` :ref:`idxen` :ref:`offen` :ref:`offset12` :ref:`slc` + **INSTRUCTION** **SRC0** **SRC1** **SRC2** **SRC3** **MODIFIERS** + \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---| + buffer_atomic_add_f32 :ref:`vdata`, :ref:`vaddr`, :ref:`srsrc`, :ref:`soffset` :ref:`idxen` :ref:`offen` :ref:`offset12` :ref:`slc` + buffer_atomic_pk_add_f16 :ref:`vdata`, :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:`src0`, :ref:`vsrc1` - 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:`src1`::ref:`m` :ref:`dst_sel` :ref:`dst_unused` :ref:`src0_sel` :ref:`src1_sel` + 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:`src0`, :ref:`vsrc1` + 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:`src1`::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` + 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:`vsrc` - v_accvgpr_write_b32 :ref:`vdst`, :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:`i8x4`, :ref:`src1`::ref:`i8x4`, :ref:`src2`::ref:`i32` :ref:`clamp` - v_dot4_u32_u8 :ref:`vdst`, :ref:`src0`::ref:`u8x4`, :ref:`src1`::ref:`u8x4`, :ref:`src2`::ref:`u32` :ref:`clamp` - v_dot8_i32_i4 :ref:`vdst`, :ref:`src0`::ref:`i4x8`, :ref:`src1`::ref:`i4x8`, :ref:`src2`::ref:`i32` :ref:`clamp` - v_dot8_u32_u4 :ref:`vdst`, :ref:`src0`::ref:`u4x8`, :ref:`src1`::ref:`u4x8`, :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:`vdst`::ref:`f32x4`, :ref:`vsrc0`::ref:`f16x4`, :ref:`vsrc1`::ref:`f16x4`, :ref:`vsrc2`::ref:`f32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp` - v_mfma_f32_16x16x1f32 :ref:`vdst`::ref:`f32x16`, :ref:`vsrc0`::ref:`f32`, :ref:`vsrc1`::ref:`f32`, :ref:`vsrc2`::ref:`f32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp` - v_mfma_f32_16x16x2bf16 :ref:`vdst`::ref:`f32x16`, :ref:`vsrc0`::ref:`bf16x2`, :ref:`vsrc1`::ref:`bf16x2`, :ref:`vsrc2`::ref:`f32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp` - v_mfma_f32_16x16x4f16 :ref:`vdst`::ref:`f32x16`, :ref:`vsrc0`::ref:`f16x4`, :ref:`vsrc1`::ref:`f16x4`, :ref:`vsrc2`::ref:`f32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp` - v_mfma_f32_16x16x4f32 :ref:`vdst`::ref:`f32x4`, :ref:`vsrc0`::ref:`f32`, :ref:`vsrc1`::ref:`f32`, :ref:`vsrc2`::ref:`f32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp` - v_mfma_f32_16x16x8bf16 :ref:`vdst`::ref:`f32x4`, :ref:`vsrc0`::ref:`bf16x2`, :ref:`vsrc1`::ref:`bf16x2`, :ref:`vsrc2`::ref:`f32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp` - v_mfma_f32_32x32x1f32 :ref:`vdst`::ref:`f32x32`, :ref:`vsrc0`::ref:`f32`, :ref:`vsrc1`::ref:`f32`, :ref:`vsrc2`::ref:`f32x32` :ref:`cbsz` :ref:`abid` :ref:`blgp` - v_mfma_f32_32x32x2bf16 :ref:`vdst`::ref:`f32x32`, :ref:`vsrc0`::ref:`bf16x2`, :ref:`vsrc1`::ref:`bf16x2`, :ref:`vsrc2`::ref:`f32x32` :ref:`cbsz` :ref:`abid` :ref:`blgp` - v_mfma_f32_32x32x2f32 :ref:`vdst`::ref:`f32x16`, :ref:`vsrc0`::ref:`f32`, :ref:`vsrc1`::ref:`f32`, :ref:`vsrc2`::ref:`f32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp` - v_mfma_f32_32x32x4bf16 :ref:`vdst`::ref:`f32x16`, :ref:`vsrc0`::ref:`bf16x2`, :ref:`vsrc1`::ref:`bf16x2`, :ref:`vsrc2`::ref:`f32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp` - v_mfma_f32_32x32x4f16 :ref:`vdst`::ref:`f32x32`, :ref:`vsrc0`::ref:`f16x4`, :ref:`vsrc1`::ref:`f16x4`, :ref:`vsrc2`::ref:`f32x32` :ref:`cbsz` :ref:`abid` :ref:`blgp` - v_mfma_f32_32x32x8f16 :ref:`vdst`::ref:`f32x16`, :ref:`vsrc0`::ref:`f16x4`, :ref:`vsrc1`::ref:`f16x4`, :ref:`vsrc2`::ref:`f32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp` - v_mfma_f32_4x4x1f32 :ref:`vdst`::ref:`f32x4`, :ref:`vsrc0`::ref:`f32`, :ref:`vsrc1`::ref:`f32`, :ref:`vsrc2`::ref:`f32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp` - v_mfma_f32_4x4x2bf16 :ref:`vdst`::ref:`f32x4`, :ref:`vsrc0`::ref:`bf16x2`, :ref:`vsrc1`::ref:`bf16x2`, :ref:`vsrc2`::ref:`f32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp` - v_mfma_f32_4x4x4f16 :ref:`vdst`::ref:`f32x4`, :ref:`vsrc0`::ref:`f16x4`, :ref:`vsrc1`::ref:`f16x4`, :ref:`vsrc2`::ref:`f32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp` - v_mfma_i32_16x16x16i8 :ref:`vdst`::ref:`i32x4`, :ref:`vsrc0`::ref:`i8x4`, :ref:`vsrc1`::ref:`i8x4`, :ref:`vsrc2`::ref:`i32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp` - v_mfma_i32_16x16x4i8 :ref:`vdst`::ref:`i32x16`, :ref:`vsrc0`::ref:`i8x4`, :ref:`vsrc1`::ref:`i8x4`, :ref:`vsrc2`::ref:`i32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp` - v_mfma_i32_32x32x4i8 :ref:`vdst`::ref:`i32x32`, :ref:`vsrc0`::ref:`i8x4`, :ref:`vsrc1`::ref:`i8x4`, :ref:`vsrc2`::ref:`i32x32` :ref:`cbsz` :ref:`abid` :ref:`blgp` - v_mfma_i32_32x32x8i8 :ref:`vdst`::ref:`i32x16`, :ref:`vsrc0`::ref:`i8x4`, :ref:`vsrc1`::ref:`i8x4`, :ref:`vsrc2`::ref:`i32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp` - v_mfma_i32_4x4x4i8 :ref:`vdst`::ref:`i32x4`, :ref:`vsrc0`::ref:`i8x4`, :ref:`vsrc1`::ref:`i8x4`, :ref:`vsrc2`::ref:`i32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_accvgpr_read_b32 :ref:`vdst`, :ref:`vsrc` + v_accvgpr_write_b32 :ref:`vdst`, :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:`i8x4`, :ref:`src1`::ref:`i8x4`, :ref:`src2`::ref:`i32` :ref:`clamp` + v_dot4_u32_u8 :ref:`vdst`, :ref:`src0`::ref:`u8x4`, :ref:`src1`::ref:`u8x4`, :ref:`src2`::ref:`u32` :ref:`clamp` + v_dot8_i32_i4 :ref:`vdst`, :ref:`src0`::ref:`i4x8`, :ref:`src1`::ref:`i4x8`, :ref:`src2`::ref:`i32` :ref:`clamp` + v_dot8_u32_u4 :ref:`vdst`, :ref:`src0`::ref:`u4x8`, :ref:`src1`::ref:`u4x8`, :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:`vdst`::ref:`f32x4`, :ref:`vsrc0`::ref:`f16x4`, :ref:`vsrc1`::ref:`f16x4`, :ref:`vsrc2`::ref:`f32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_16x16x1f32 :ref:`vdst`::ref:`f32x16`, :ref:`vsrc0`::ref:`f32`, :ref:`vsrc1`::ref:`f32`, :ref:`vsrc2`::ref:`f32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_16x16x2bf16 :ref:`vdst`::ref:`f32x16`, :ref:`vsrc0`::ref:`bf16x2`, :ref:`vsrc1`::ref:`bf16x2`, :ref:`vsrc2`::ref:`f32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_16x16x4f16 :ref:`vdst`::ref:`f32x16`, :ref:`vsrc0`::ref:`f16x4`, :ref:`vsrc1`::ref:`f16x4`, :ref:`vsrc2`::ref:`f32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_16x16x4f32 :ref:`vdst`::ref:`f32x4`, :ref:`vsrc0`::ref:`f32`, :ref:`vsrc1`::ref:`f32`, :ref:`vsrc2`::ref:`f32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_16x16x8bf16 :ref:`vdst`::ref:`f32x4`, :ref:`vsrc0`::ref:`bf16x2`, :ref:`vsrc1`::ref:`bf16x2`, :ref:`vsrc2`::ref:`f32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_32x32x1f32 :ref:`vdst`::ref:`f32x32`, :ref:`vsrc0`::ref:`f32`, :ref:`vsrc1`::ref:`f32`, :ref:`vsrc2`::ref:`f32x32` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_32x32x2bf16 :ref:`vdst`::ref:`f32x32`, :ref:`vsrc0`::ref:`bf16x2`, :ref:`vsrc1`::ref:`bf16x2`, :ref:`vsrc2`::ref:`f32x32` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_32x32x2f32 :ref:`vdst`::ref:`f32x16`, :ref:`vsrc0`::ref:`f32`, :ref:`vsrc1`::ref:`f32`, :ref:`vsrc2`::ref:`f32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_32x32x4bf16 :ref:`vdst`::ref:`f32x16`, :ref:`vsrc0`::ref:`bf16x2`, :ref:`vsrc1`::ref:`bf16x2`, :ref:`vsrc2`::ref:`f32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_32x32x4f16 :ref:`vdst`::ref:`f32x32`, :ref:`vsrc0`::ref:`f16x4`, :ref:`vsrc1`::ref:`f16x4`, :ref:`vsrc2`::ref:`f32x32` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_32x32x8f16 :ref:`vdst`::ref:`f32x16`, :ref:`vsrc0`::ref:`f16x4`, :ref:`vsrc1`::ref:`f16x4`, :ref:`vsrc2`::ref:`f32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_4x4x1f32 :ref:`vdst`::ref:`f32x4`, :ref:`vsrc0`::ref:`f32`, :ref:`vsrc1`::ref:`f32`, :ref:`vsrc2`::ref:`f32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_4x4x2bf16 :ref:`vdst`::ref:`f32x4`, :ref:`vsrc0`::ref:`bf16x2`, :ref:`vsrc1`::ref:`bf16x2`, :ref:`vsrc2`::ref:`f32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_f32_4x4x4f16 :ref:`vdst`::ref:`f32x4`, :ref:`vsrc0`::ref:`f16x4`, :ref:`vsrc1`::ref:`f16x4`, :ref:`vsrc2`::ref:`f32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_i32_16x16x16i8 :ref:`vdst`::ref:`i32x4`, :ref:`vsrc0`::ref:`i8x4`, :ref:`vsrc1`::ref:`i8x4`, :ref:`vsrc2`::ref:`i32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_i32_16x16x4i8 :ref:`vdst`::ref:`i32x16`, :ref:`vsrc0`::ref:`i8x4`, :ref:`vsrc1`::ref:`i8x4`, :ref:`vsrc2`::ref:`i32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_i32_32x32x4i8 :ref:`vdst`::ref:`i32x32`, :ref:`vsrc0`::ref:`i8x4`, :ref:`vsrc1`::ref:`i8x4`, :ref:`vsrc2`::ref:`i32x32` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_i32_32x32x8i8 :ref:`vdst`::ref:`i32x16`, :ref:`vsrc0`::ref:`i8x4`, :ref:`vsrc1`::ref:`i8x4`, :ref:`vsrc2`::ref:`i32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp` + v_mfma_i32_4x4x4i8 :ref:`vdst`::ref:`i32x4`, :ref:`vsrc0`::ref:`i8x4`, :ref:`vsrc1`::ref:`i8x4`, :ref:`vsrc2`::ref:`i32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp` .. |---| unicode:: U+02014 .. em dash .. toctree:: :hidden: - gfx908_dst gfx908_fx_operand - gfx908_m - gfx908_m_1 - gfx908_opt + gfx908_m_254bcb + gfx908_m_f5d306 gfx908_saddr gfx908_soffset - gfx908_src - gfx908_src_1 - gfx908_src_2 - gfx908_src_3 - gfx908_src_4 - gfx908_src_5 + gfx908_src_4e78e6 + gfx908_src_58d119 + gfx908_src_73ab34 + gfx908_src_955b45 + gfx908_src_d578c4 + gfx908_src_d95796 gfx908_srsrc gfx908_type_deviation - gfx908_vaddr - gfx908_vaddr_1 - gfx908_vdata - gfx908_vdata_1 - gfx908_vdst - gfx908_vdst_1 - gfx908_vdst_2 - gfx908_vdst_3 - gfx908_vdst_4 - gfx908_vdst_5 - gfx908_vsrc - gfx908_vsrc_1 - gfx908_vsrc_2 - gfx908_vsrc_3 - gfx908_vsrc_4 - gfx908_vsrc_5 - gfx908_vsrc_6 + gfx908_vaddr_0212e3 + gfx908_vaddr_b73dc0 + gfx908_vdata_6802ce + gfx908_vdata_fe1edf + gfx908_vdst_0c4ef8 + gfx908_vdst_2c8d1e + gfx908_vdst_78dd0a + gfx908_vdst_89680f + gfx908_vdst_bcee7a + gfx908_vsrc_036abe + gfx908_vsrc_1027ca + gfx908_vsrc_2d4632 + gfx908_vsrc_6802ce + gfx908_vsrc_9ad749 + gfx908_vsrc_be4895 + gfx908_vsrc_f3d248 diff --git a/llvm/docs/AMDGPU/gfx908_dst.rst b/llvm/docs/AMDGPU/gfx908_dst.rst deleted file mode 100644 index 80d8a40..0000000 --- a/llvm/docs/AMDGPU/gfx908_dst.rst +++ /dev/null @@ -1,13 +0,0 @@ -.. - ************************************************** - * * - * Automatically generated file, do not edit! * - * * - ************************************************** - -.. _amdgpu_synid_gfx908_dst: - -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_m_1.rst b/llvm/docs/AMDGPU/gfx908_m_254bcb.rst similarity index 91% rename from llvm/docs/AMDGPU/gfx908_m_1.rst rename to llvm/docs/AMDGPU/gfx908_m_254bcb.rst index 0e697e5..cb0ae72 100644 --- a/llvm/docs/AMDGPU/gfx908_m_1.rst +++ b/llvm/docs/AMDGPU/gfx908_m_254bcb.rst @@ -5,7 +5,7 @@ * * ************************************************** -.. _amdgpu_synid_gfx908_m_1: +.. _amdgpu_synid_gfx908_m_254bcb: m = diff --git a/llvm/docs/AMDGPU/gfx908_m.rst b/llvm/docs/AMDGPU/gfx908_m_f5d306.rst similarity index 92% rename from llvm/docs/AMDGPU/gfx908_m.rst rename to llvm/docs/AMDGPU/gfx908_m_f5d306.rst index 3f6c0d3..77c1481 100644 --- a/llvm/docs/AMDGPU/gfx908_m.rst +++ b/llvm/docs/AMDGPU/gfx908_m_f5d306.rst @@ -5,7 +5,7 @@ * * ************************************************** -.. _amdgpu_synid_gfx908_m: +.. _amdgpu_synid_gfx908_m_f5d306: m = diff --git a/llvm/docs/AMDGPU/gfx908_opt.rst b/llvm/docs/AMDGPU/gfx908_opt.rst deleted file mode 100644 index 343f0c1..0000000 --- a/llvm/docs/AMDGPU/gfx908_opt.rst +++ /dev/null @@ -1,13 +0,0 @@ -.. - ************************************************** - * * - * Automatically generated file, do not edit! * - * * - ************************************************** - -.. _amdgpu_synid_gfx908_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_saddr.rst b/llvm/docs/AMDGPU/gfx908_saddr.rst index 950d709..35a1de8 100644 --- a/llvm/docs/AMDGPU/gfx908_saddr.rst +++ b/llvm/docs/AMDGPU/gfx908_saddr.rst @@ -12,7 +12,7 @@ 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. +See :ref:`vaddr` for description of available addressing modes. *Size:* 2 dwords. diff --git a/llvm/docs/AMDGPU/gfx908_src_2.rst b/llvm/docs/AMDGPU/gfx908_src_4e78e6.rst similarity index 95% rename from llvm/docs/AMDGPU/gfx908_src_2.rst rename to llvm/docs/AMDGPU/gfx908_src_4e78e6.rst index 6de8806..f08acf5 100644 --- a/llvm/docs/AMDGPU/gfx908_src_2.rst +++ b/llvm/docs/AMDGPU/gfx908_src_4e78e6.rst @@ -5,7 +5,7 @@ * * ************************************************** -.. _amdgpu_synid_gfx908_src_2: +.. _amdgpu_synid_gfx908_src_4e78e6: src === diff --git a/llvm/docs/AMDGPU/gfx908_src_3.rst b/llvm/docs/AMDGPU/gfx908_src_58d119.rst similarity index 92% rename from llvm/docs/AMDGPU/gfx908_src_3.rst rename to llvm/docs/AMDGPU/gfx908_src_58d119.rst index 588d535..6a966e1 100644 --- a/llvm/docs/AMDGPU/gfx908_src_3.rst +++ b/llvm/docs/AMDGPU/gfx908_src_58d119.rst @@ -5,7 +5,7 @@ * * ************************************************** -.. _amdgpu_synid_gfx908_src_3: +.. _amdgpu_synid_gfx908_src_58d119: src === diff --git a/llvm/docs/AMDGPU/gfx908_src.rst b/llvm/docs/AMDGPU/gfx908_src_73ab34.rst similarity index 95% rename from llvm/docs/AMDGPU/gfx908_src.rst rename to llvm/docs/AMDGPU/gfx908_src_73ab34.rst index c9463fd..024448a 100644 --- a/llvm/docs/AMDGPU/gfx908_src.rst +++ b/llvm/docs/AMDGPU/gfx908_src_73ab34.rst @@ -5,7 +5,7 @@ * * ************************************************** -.. _amdgpu_synid_gfx908_src: +.. _amdgpu_synid_gfx908_src_73ab34: src === diff --git a/llvm/docs/AMDGPU/gfx908_src_4.rst b/llvm/docs/AMDGPU/gfx908_src_955b45.rst similarity index 95% rename from llvm/docs/AMDGPU/gfx908_src_4.rst rename to llvm/docs/AMDGPU/gfx908_src_955b45.rst index c098de8..bf3b01d 100644 --- a/llvm/docs/AMDGPU/gfx908_src_4.rst +++ b/llvm/docs/AMDGPU/gfx908_src_955b45.rst @@ -5,7 +5,7 @@ * * ************************************************** -.. _amdgpu_synid_gfx908_src_4: +.. _amdgpu_synid_gfx908_src_955b45: src === diff --git a/llvm/docs/AMDGPU/gfx908_src_1.rst b/llvm/docs/AMDGPU/gfx908_src_d578c4.rst similarity index 95% rename from llvm/docs/AMDGPU/gfx908_src_1.rst rename to llvm/docs/AMDGPU/gfx908_src_d578c4.rst index 42d9451..194e776 100644 --- a/llvm/docs/AMDGPU/gfx908_src_1.rst +++ b/llvm/docs/AMDGPU/gfx908_src_d578c4.rst @@ -5,7 +5,7 @@ * * ************************************************** -.. _amdgpu_synid_gfx908_src_1: +.. _amdgpu_synid_gfx908_src_d578c4: src === diff --git a/llvm/docs/AMDGPU/gfx908_src_5.rst b/llvm/docs/AMDGPU/gfx908_src_d95796.rst similarity index 95% rename from llvm/docs/AMDGPU/gfx908_src_5.rst rename to llvm/docs/AMDGPU/gfx908_src_d95796.rst index 0b95e87..7364956 100644 --- a/llvm/docs/AMDGPU/gfx908_src_5.rst +++ b/llvm/docs/AMDGPU/gfx908_src_d95796.rst @@ -5,7 +5,7 @@ * * ************************************************** -.. _amdgpu_synid_gfx908_src_5: +.. _amdgpu_synid_gfx908_src_d95796: src === diff --git a/llvm/docs/AMDGPU/gfx908_vaddr.rst b/llvm/docs/AMDGPU/gfx908_vaddr.rst deleted file mode 100644 index 3b785ff..0000000 --- a/llvm/docs/AMDGPU/gfx908_vaddr.rst +++ /dev/null @@ -1,20 +0,0 @@ -.. - ************************************************** - * * - * Automatically generated file, do not edit! * - * * - ************************************************** - -.. _amdgpu_synid_gfx908_vaddr: - -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`. - -*Size:* 1 or 2 dwords. - -*Operands:* :ref:`v` diff --git a/llvm/docs/AMDGPU/gfx908_vaddr_0212e3.rst b/llvm/docs/AMDGPU/gfx908_vaddr_0212e3.rst new file mode 100644 index 0000000..5ecbb08 --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_vaddr_0212e3.rst @@ -0,0 +1,20 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid_gfx908_vaddr_0212e3: + +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`. + +*Size:* 1 or 2 dwords. + +*Operands:* :ref:`v` diff --git a/llvm/docs/AMDGPU/gfx908_vaddr_1.rst b/llvm/docs/AMDGPU/gfx908_vaddr_b73dc0.rst similarity index 96% rename from llvm/docs/AMDGPU/gfx908_vaddr_1.rst rename to llvm/docs/AMDGPU/gfx908_vaddr_b73dc0.rst index c749bd4..867ac36 100644 --- a/llvm/docs/AMDGPU/gfx908_vaddr_1.rst +++ b/llvm/docs/AMDGPU/gfx908_vaddr_b73dc0.rst @@ -5,7 +5,7 @@ * * ************************************************** -.. _amdgpu_synid_gfx908_vaddr_1: +.. _amdgpu_synid_gfx908_vaddr_b73dc0: vaddr ===== diff --git a/llvm/docs/AMDGPU/gfx908_vdata_1.rst b/llvm/docs/AMDGPU/gfx908_vdata_1.rst deleted file mode 100644 index 7d57811..0000000 --- a/llvm/docs/AMDGPU/gfx908_vdata_1.rst +++ /dev/null @@ -1,21 +0,0 @@ -.. - ************************************************** - * * - * Automatically generated file, do not edit! * - * * - ************************************************** - -.. _amdgpu_synid_gfx908_vdata_1: - -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_vdata.rst b/llvm/docs/AMDGPU/gfx908_vdata_6802ce.rst similarity index 90% rename from llvm/docs/AMDGPU/gfx908_vdata.rst rename to llvm/docs/AMDGPU/gfx908_vdata_6802ce.rst index 6db6e57..8c42320 100644 --- a/llvm/docs/AMDGPU/gfx908_vdata.rst +++ b/llvm/docs/AMDGPU/gfx908_vdata_6802ce.rst @@ -5,7 +5,7 @@ * * ************************************************** -.. _amdgpu_synid_gfx908_vdata: +.. _amdgpu_synid_gfx908_vdata_6802ce: vdata ===== diff --git a/llvm/docs/AMDGPU/gfx908_vdata_fe1edf.rst b/llvm/docs/AMDGPU/gfx908_vdata_fe1edf.rst new file mode 100644 index 0000000..dc5d48a --- /dev/null +++ b/llvm/docs/AMDGPU/gfx908_vdata_fe1edf.rst @@ -0,0 +1,17 @@ +.. + ************************************************** + * * + * Automatically generated file, do not edit! * + * * + ************************************************** + +.. _amdgpu_synid_gfx908_vdata_fe1edf: + +vdata +===== + +Input data for an atomic instruction. + +*Size:* 1 dword. + +*Operands:* :ref:`v` diff --git a/llvm/docs/AMDGPU/gfx908_vdst.rst b/llvm/docs/AMDGPU/gfx908_vdst.rst deleted file mode 100644 index f3b952b..0000000 --- a/llvm/docs/AMDGPU/gfx908_vdst.rst +++ /dev/null @@ -1,19 +0,0 @@ -.. - ************************************************** - * * - * Automatically generated file, do not edit! * - * * - ************************************************** - -.. _amdgpu_synid_gfx908_vdst: - -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_vdst_4.rst b/llvm/docs/AMDGPU/gfx908_vdst_0c4ef8.rst similarity index 90% rename from llvm/docs/AMDGPU/gfx908_vdst_4.rst rename to llvm/docs/AMDGPU/gfx908_vdst_0c4ef8.rst index dab8c5b..c18c764 100644 --- a/llvm/docs/AMDGPU/gfx908_vdst_4.rst +++ b/llvm/docs/AMDGPU/gfx908_vdst_0c4ef8.rst @@ -5,7 +5,7 @@ * * ************************************************** -.. _amdgpu_synid_gfx908_vdst_4: +.. _amdgpu_synid_gfx908_vdst_0c4ef8: vdst ==== diff --git a/llvm/docs/AMDGPU/gfx908_vdst_5.rst b/llvm/docs/AMDGPU/gfx908_vdst_2c8d1e.rst similarity index 90% rename from llvm/docs/AMDGPU/gfx908_vdst_5.rst rename to llvm/docs/AMDGPU/gfx908_vdst_2c8d1e.rst index 1112a44..6e4fa0f 100644 --- a/llvm/docs/AMDGPU/gfx908_vdst_5.rst +++ b/llvm/docs/AMDGPU/gfx908_vdst_2c8d1e.rst @@ -5,7 +5,7 @@ * * ************************************************** -.. _amdgpu_synid_gfx908_vdst_5: +.. _amdgpu_synid_gfx908_vdst_2c8d1e: vdst ==== diff --git a/llvm/docs/AMDGPU/gfx908_vdst_2.rst b/llvm/docs/AMDGPU/gfx908_vdst_78dd0a.rst similarity index 90% rename from llvm/docs/AMDGPU/gfx908_vdst_2.rst rename to llvm/docs/AMDGPU/gfx908_vdst_78dd0a.rst index 4eb4231..1616b03 100644 --- a/llvm/docs/AMDGPU/gfx908_vdst_2.rst +++ b/llvm/docs/AMDGPU/gfx908_vdst_78dd0a.rst @@ -5,7 +5,7 @@ * * ************************************************** -.. _amdgpu_synid_gfx908_vdst_2: +.. _amdgpu_synid_gfx908_vdst_78dd0a: vdst ==== diff --git a/llvm/docs/AMDGPU/gfx908_vdst_1.rst b/llvm/docs/AMDGPU/gfx908_vdst_89680f.rst similarity index 90% rename from llvm/docs/AMDGPU/gfx908_vdst_1.rst rename to llvm/docs/AMDGPU/gfx908_vdst_89680f.rst index 985f754..f4ba151 100644 --- a/llvm/docs/AMDGPU/gfx908_vdst_1.rst +++ b/llvm/docs/AMDGPU/gfx908_vdst_89680f.rst @@ -5,7 +5,7 @@ * * ************************************************** -.. _amdgpu_synid_gfx908_vdst_1: +.. _amdgpu_synid_gfx908_vdst_89680f: vdst ==== diff --git a/llvm/docs/AMDGPU/gfx908_vdst_3.rst b/llvm/docs/AMDGPU/gfx908_vdst_bcee7a.rst similarity index 90% rename from llvm/docs/AMDGPU/gfx908_vdst_3.rst rename to llvm/docs/AMDGPU/gfx908_vdst_bcee7a.rst index 9661df5..e3d35ab 100644 --- a/llvm/docs/AMDGPU/gfx908_vdst_3.rst +++ b/llvm/docs/AMDGPU/gfx908_vdst_bcee7a.rst @@ -5,7 +5,7 @@ * * ************************************************** -.. _amdgpu_synid_gfx908_vdst_3: +.. _amdgpu_synid_gfx908_vdst_bcee7a: vdst ==== diff --git a/llvm/docs/AMDGPU/gfx908_vsrc_6.rst b/llvm/docs/AMDGPU/gfx908_vsrc_036abe.rst similarity index 90% rename from llvm/docs/AMDGPU/gfx908_vsrc_6.rst rename to llvm/docs/AMDGPU/gfx908_vsrc_036abe.rst index f11d027..2ca4f19 100644 --- a/llvm/docs/AMDGPU/gfx908_vsrc_6.rst +++ b/llvm/docs/AMDGPU/gfx908_vsrc_036abe.rst @@ -5,7 +5,7 @@ * * ************************************************** -.. _amdgpu_synid_gfx908_vsrc_6: +.. _amdgpu_synid_gfx908_vsrc_036abe: vsrc ==== diff --git a/llvm/docs/AMDGPU/gfx908_vsrc_1.rst b/llvm/docs/AMDGPU/gfx908_vsrc_1027ca.rst similarity index 90% rename from llvm/docs/AMDGPU/gfx908_vsrc_1.rst rename to llvm/docs/AMDGPU/gfx908_vsrc_1027ca.rst index 373cb26..a15fb4e 100644 --- a/llvm/docs/AMDGPU/gfx908_vsrc_1.rst +++ b/llvm/docs/AMDGPU/gfx908_vsrc_1027ca.rst @@ -5,7 +5,7 @@ * * ************************************************** -.. _amdgpu_synid_gfx908_vsrc_1: +.. _amdgpu_synid_gfx908_vsrc_1027ca: vsrc ==== diff --git a/llvm/docs/AMDGPU/gfx908_vsrc_5.rst b/llvm/docs/AMDGPU/gfx908_vsrc_2d4632.rst similarity index 90% rename from llvm/docs/AMDGPU/gfx908_vsrc_5.rst rename to llvm/docs/AMDGPU/gfx908_vsrc_2d4632.rst index b51ca84..8e8fabc 100644 --- a/llvm/docs/AMDGPU/gfx908_vsrc_5.rst +++ b/llvm/docs/AMDGPU/gfx908_vsrc_2d4632.rst @@ -5,7 +5,7 @@ * * ************************************************** -.. _amdgpu_synid_gfx908_vsrc_5: +.. _amdgpu_synid_gfx908_vsrc_2d4632: vsrc ==== diff --git a/llvm/docs/AMDGPU/gfx908_vsrc.rst b/llvm/docs/AMDGPU/gfx908_vsrc_6802ce.rst similarity index 90% rename from llvm/docs/AMDGPU/gfx908_vsrc.rst rename to llvm/docs/AMDGPU/gfx908_vsrc_6802ce.rst index 0c3ad8a..edfdb85 100644 --- a/llvm/docs/AMDGPU/gfx908_vsrc.rst +++ b/llvm/docs/AMDGPU/gfx908_vsrc_6802ce.rst @@ -5,7 +5,7 @@ * * ************************************************** -.. _amdgpu_synid_gfx908_vsrc: +.. _amdgpu_synid_gfx908_vsrc_6802ce: vsrc ==== diff --git a/llvm/docs/AMDGPU/gfx908_vsrc_2.rst b/llvm/docs/AMDGPU/gfx908_vsrc_9ad749.rst similarity index 91% rename from llvm/docs/AMDGPU/gfx908_vsrc_2.rst rename to llvm/docs/AMDGPU/gfx908_vsrc_9ad749.rst index d70ed03..7535c35 100644 --- a/llvm/docs/AMDGPU/gfx908_vsrc_2.rst +++ b/llvm/docs/AMDGPU/gfx908_vsrc_9ad749.rst @@ -5,7 +5,7 @@ * * ************************************************** -.. _amdgpu_synid_gfx908_vsrc_2: +.. _amdgpu_synid_gfx908_vsrc_9ad749: vsrc ==== diff --git a/llvm/docs/AMDGPU/gfx908_vsrc_4.rst b/llvm/docs/AMDGPU/gfx908_vsrc_be4895.rst similarity index 91% rename from llvm/docs/AMDGPU/gfx908_vsrc_4.rst rename to llvm/docs/AMDGPU/gfx908_vsrc_be4895.rst index e3c8c3f..2b9a75f 100644 --- a/llvm/docs/AMDGPU/gfx908_vsrc_4.rst +++ b/llvm/docs/AMDGPU/gfx908_vsrc_be4895.rst @@ -5,7 +5,7 @@ * * ************************************************** -.. _amdgpu_synid_gfx908_vsrc_4: +.. _amdgpu_synid_gfx908_vsrc_be4895: vsrc ==== diff --git a/llvm/docs/AMDGPU/gfx908_vsrc_3.rst b/llvm/docs/AMDGPU/gfx908_vsrc_f3d248.rst similarity index 90% rename from llvm/docs/AMDGPU/gfx908_vsrc_3.rst rename to llvm/docs/AMDGPU/gfx908_vsrc_f3d248.rst index 2456afb..98f94d4 100644 --- a/llvm/docs/AMDGPU/gfx908_vsrc_3.rst +++ b/llvm/docs/AMDGPU/gfx908_vsrc_f3d248.rst @@ -5,7 +5,7 @@ * * ************************************************** -.. _amdgpu_synid_gfx908_vsrc_3: +.. _amdgpu_synid_gfx908_vsrc_f3d248: vsrc ==== -- 2.7.4