From 2de2275cbdb8d123920f454f34ed4cfb4e1d2dcc Mon Sep 17 00:00:00 2001
From: Dmitry Preobrazhensky <dmitry.preobrazhensky@amd.com>
Date: Fri, 7 Feb 2020 16:20:37 +0300
Subject: [PATCH] [AMDGPU][MC][DOC] Updated AMD GPU assembler syntax
 description.

Summary of changes:
- updated description of gfx906 and gfx908;
- added description of gfx1011 and gfx1012 subtargets.
---
 llvm/docs/AMDGPU/AMDGPUAsmGFX1011.rst         | 92 +++++++++++++++++++++++++++
 llvm/docs/AMDGPU/AMDGPUAsmGFX906.rst          |  8 +--
 llvm/docs/AMDGPU/AMDGPUAsmGFX908.rst          | 70 ++++++++++----------
 llvm/docs/AMDGPU/gfx1011_src32_0.rst          | 17 +++++
 llvm/docs/AMDGPU/gfx1011_src32_1.rst          | 17 +++++
 llvm/docs/AMDGPU/gfx1011_type_dev.rst         | 13 ++++
 llvm/docs/AMDGPU/gfx1011_vdst32_0.rst         | 17 +++++
 llvm/docs/AMDGPU/gfx1011_vsrc32_0.rst         | 17 +++++
 llvm/docs/AMDGPU/gfx908_saddr_flat_global.rst |  2 +-
 llvm/docs/AMDGPUUsage.rst                     | 29 +++++----
 10 files changed, 228 insertions(+), 54 deletions(-)
 create mode 100644 llvm/docs/AMDGPU/AMDGPUAsmGFX1011.rst
 create mode 100644 llvm/docs/AMDGPU/gfx1011_src32_0.rst
 create mode 100644 llvm/docs/AMDGPU/gfx1011_src32_1.rst
 create mode 100644 llvm/docs/AMDGPU/gfx1011_type_dev.rst
 create mode 100644 llvm/docs/AMDGPU/gfx1011_vdst32_0.rst
 create mode 100644 llvm/docs/AMDGPU/gfx1011_vsrc32_0.rst

diff --git a/llvm/docs/AMDGPU/AMDGPUAsmGFX1011.rst b/llvm/docs/AMDGPU/AMDGPUAsmGFX1011.rst
new file mode 100644
index 0000000..5e825c6
--- /dev/null
+++ b/llvm/docs/AMDGPU/AMDGPUAsmGFX1011.rst
@@ -0,0 +1,92 @@
+..
+    **************************************************
+    *                                                *
+    *   Automatically generated file, do not edit!   *
+    *                                                *
+    **************************************************
+
+====================================================================================
+Syntax of gfx1011 and gfx1012 Instructions
+====================================================================================
+
+.. contents::
+  :local:
+
+Introduction
+============
+
+This document describes the syntax of *instructions specific to gfx1011 and gfx1012*.
+
+For a description of other gfx1011 and gfx1012 instructions see :doc:`Syntax of Core GFX10 Instructions<AMDGPUAsmGFX10>`.
+
+Notation
+========
+
+Notation used in this document is explained :ref:`here<amdgpu_syn_instruction_notation>`.
+
+Overview
+========
+
+An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document<amdgpu_syn_instructions>`.
+
+Instructions
+============
+
+
+DPP16
+-----------------------
+
+.. parsed-literal::
+
+    **INSTRUCTION**           **DST**      **SRC0**         **SRC1**          **MODIFIERS**
+    \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
+    v_dot2c_f32_f16_dpp   :ref:`vdst<amdgpu_synid1011_vdst32_0>`,    :ref:`vsrc0<amdgpu_synid1011_vsrc32_0>`::ref:`f16x2<amdgpu_synid1011_type_dev>`, :ref:`vsrc1<amdgpu_synid1011_vsrc32_0>`::ref:`f16x2<amdgpu_synid1011_type_dev>`   :ref:`dpp16_ctrl<amdgpu_synid_dpp16_ctrl>` :ref:`row_mask<amdgpu_synid_row_mask>` :ref:`bank_mask<amdgpu_synid_bank_mask>` :ref:`bound_ctrl<amdgpu_synid_bound_ctrl>` :ref:`fi<amdgpu_synid_fi16>`
+    v_dot4c_i32_i8_dpp    :ref:`vdst<amdgpu_synid1011_vdst32_0>`,    :ref:`vsrc0<amdgpu_synid1011_vsrc32_0>`::ref:`i8x4<amdgpu_synid1011_type_dev>`,  :ref:`vsrc1<amdgpu_synid1011_vsrc32_0>`::ref:`i8x4<amdgpu_synid1011_type_dev>`    :ref:`dpp16_ctrl<amdgpu_synid_dpp16_ctrl>` :ref:`row_mask<amdgpu_synid_row_mask>` :ref:`bank_mask<amdgpu_synid_bank_mask>` :ref:`bound_ctrl<amdgpu_synid_bound_ctrl>` :ref:`fi<amdgpu_synid_fi16>`
+
+DPP8
+-----------------------
+
+.. parsed-literal::
+
+    **INSTRUCTION**                    **DST**       **SRC0**         **SRC1**             **MODIFIERS**
+    \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
+    v_dot2c_f32_f16_dpp            :ref:`vdst<amdgpu_synid1011_vdst32_0>`,     :ref:`vsrc0<amdgpu_synid1011_vsrc32_0>`::ref:`f16x2<amdgpu_synid1011_type_dev>`, :ref:`vsrc1<amdgpu_synid1011_vsrc32_0>`::ref:`f16x2<amdgpu_synid1011_type_dev>`      :ref:`dpp8_sel<amdgpu_synid_dpp8_sel>` :ref:`fi<amdgpu_synid_fi8>`
+    v_dot4c_i32_i8_dpp             :ref:`vdst<amdgpu_synid1011_vdst32_0>`,     :ref:`vsrc0<amdgpu_synid1011_vsrc32_0>`::ref:`i8x4<amdgpu_synid1011_type_dev>`,  :ref:`vsrc1<amdgpu_synid1011_vsrc32_0>`::ref:`i8x4<amdgpu_synid1011_type_dev>`       :ref:`dpp8_sel<amdgpu_synid_dpp8_sel>` :ref:`fi<amdgpu_synid_fi8>`
+
+VOP2
+-----------------------
+
+.. parsed-literal::
+
+    **INSTRUCTION**                    **DST**       **SRC0**        **SRC1**
+    \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
+    v_dot2c_f32_f16                :ref:`vdst<amdgpu_synid1011_vdst32_0>`,     :ref:`src0<amdgpu_synid1011_src32_0>`::ref:`f16x2<amdgpu_synid1011_type_dev>`, :ref:`vsrc1<amdgpu_synid1011_vsrc32_0>`::ref:`f16x2<amdgpu_synid1011_type_dev>`
+    v_dot4c_i32_i8                 :ref:`vdst<amdgpu_synid1011_vdst32_0>`,     :ref:`src0<amdgpu_synid1011_src32_0>`::ref:`i8x4<amdgpu_synid1011_type_dev>`,  :ref:`vsrc1<amdgpu_synid1011_vsrc32_0>`::ref:`i8x4<amdgpu_synid1011_type_dev>`
+
+VOP3P
+-----------------------
+
+.. parsed-literal::
+
+    **INSTRUCTION**                    **DST**       **SRC0**        **SRC1**        **SRC2**           **MODIFIERS**
+    \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
+    v_dot2_f32_f16                 :ref:`vdst<amdgpu_synid1011_vdst32_0>`,     :ref:`src0<amdgpu_synid1011_src32_0>`::ref:`f16x2<amdgpu_synid1011_type_dev>`, :ref:`src1<amdgpu_synid1011_src32_1>`::ref:`f16x2<amdgpu_synid1011_type_dev>`, :ref:`src2<amdgpu_synid1011_src32_1>`::ref:`f32<amdgpu_synid1011_type_dev>`       :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>`
+    v_dot2_i32_i16                 :ref:`vdst<amdgpu_synid1011_vdst32_0>`,     :ref:`src0<amdgpu_synid1011_src32_0>`::ref:`i16x2<amdgpu_synid1011_type_dev>`, :ref:`src1<amdgpu_synid1011_src32_1>`::ref:`i16x2<amdgpu_synid1011_type_dev>`, :ref:`src2<amdgpu_synid1011_src32_1>`::ref:`i32<amdgpu_synid1011_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
+    v_dot2_u32_u16                 :ref:`vdst<amdgpu_synid1011_vdst32_0>`,     :ref:`src0<amdgpu_synid1011_src32_0>`::ref:`u16x2<amdgpu_synid1011_type_dev>`, :ref:`src1<amdgpu_synid1011_src32_1>`::ref:`u16x2<amdgpu_synid1011_type_dev>`, :ref:`src2<amdgpu_synid1011_src32_1>`::ref:`u32<amdgpu_synid1011_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
+    v_dot4_i32_i8                  :ref:`vdst<amdgpu_synid1011_vdst32_0>`,     :ref:`src0<amdgpu_synid1011_src32_0>`::ref:`i8x4<amdgpu_synid1011_type_dev>`,  :ref:`src1<amdgpu_synid1011_src32_1>`::ref:`i8x4<amdgpu_synid1011_type_dev>`,  :ref:`src2<amdgpu_synid1011_src32_1>`::ref:`i32<amdgpu_synid1011_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
+    v_dot4_u32_u8                  :ref:`vdst<amdgpu_synid1011_vdst32_0>`,     :ref:`src0<amdgpu_synid1011_src32_0>`::ref:`u8x4<amdgpu_synid1011_type_dev>`,  :ref:`src1<amdgpu_synid1011_src32_1>`::ref:`u8x4<amdgpu_synid1011_type_dev>`,  :ref:`src2<amdgpu_synid1011_src32_1>`::ref:`u32<amdgpu_synid1011_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
+    v_dot8_i32_i4                  :ref:`vdst<amdgpu_synid1011_vdst32_0>`,     :ref:`src0<amdgpu_synid1011_src32_0>`::ref:`i4x8<amdgpu_synid1011_type_dev>`,  :ref:`src1<amdgpu_synid1011_src32_1>`::ref:`i4x8<amdgpu_synid1011_type_dev>`,  :ref:`src2<amdgpu_synid1011_src32_1>`::ref:`i32<amdgpu_synid1011_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
+    v_dot8_u32_u4                  :ref:`vdst<amdgpu_synid1011_vdst32_0>`,     :ref:`src0<amdgpu_synid1011_src32_0>`::ref:`u4x8<amdgpu_synid1011_type_dev>`,  :ref:`src1<amdgpu_synid1011_src32_1>`::ref:`u4x8<amdgpu_synid1011_type_dev>`,  :ref:`src2<amdgpu_synid1011_src32_1>`::ref:`u32<amdgpu_synid1011_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
+
+.. |---| unicode:: U+02014 .. em dash
+
+
+.. toctree::
+    :hidden:
+
+    AMDGPUAsmGFX10
+    gfx1011_src32_0
+    gfx1011_src32_1
+    gfx1011_vdst32_0
+    gfx1011_vsrc32_0
+    gfx1011_type_dev
diff --git a/llvm/docs/AMDGPU/AMDGPUAsmGFX906.rst b/llvm/docs/AMDGPU/AMDGPUAsmGFX906.rst
index ee9c45d..385917d 100644
--- a/llvm/docs/AMDGPU/AMDGPUAsmGFX906.rst
+++ b/llvm/docs/AMDGPU/AMDGPUAsmGFX906.rst
@@ -66,10 +66,10 @@ VOP3P
     v_dot2_f32_f16                 :ref:`vdst<amdgpu_synid906_vdst32_0>`,     :ref:`src0<amdgpu_synid906_src32_1>`::ref:`f16x2<amdgpu_synid906_type_dev>`, :ref:`src1<amdgpu_synid906_src32_2>`::ref:`f16x2<amdgpu_synid906_type_dev>`, :ref:`src2<amdgpu_synid906_src32_2>`::ref:`f32<amdgpu_synid906_type_dev>`       :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>`
     v_dot2_i32_i16                 :ref:`vdst<amdgpu_synid906_vdst32_0>`,     :ref:`src0<amdgpu_synid906_src32_1>`::ref:`i16x2<amdgpu_synid906_type_dev>`, :ref:`src1<amdgpu_synid906_src32_2>`::ref:`i16x2<amdgpu_synid906_type_dev>`, :ref:`src2<amdgpu_synid906_src32_2>`::ref:`i32<amdgpu_synid906_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
     v_dot2_u32_u16                 :ref:`vdst<amdgpu_synid906_vdst32_0>`,     :ref:`src0<amdgpu_synid906_src32_1>`::ref:`u16x2<amdgpu_synid906_type_dev>`, :ref:`src1<amdgpu_synid906_src32_2>`::ref:`u16x2<amdgpu_synid906_type_dev>`, :ref:`src2<amdgpu_synid906_src32_2>`::ref:`u32<amdgpu_synid906_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
-    v_dot4_i32_i8                  :ref:`vdst<amdgpu_synid906_vdst32_0>`,     :ref:`src0<amdgpu_synid906_src32_1>`::ref:`i8x2<amdgpu_synid906_type_dev>`,  :ref:`src1<amdgpu_synid906_src32_2>`::ref:`i8x2<amdgpu_synid906_type_dev>`,  :ref:`src2<amdgpu_synid906_src32_2>`::ref:`i32<amdgpu_synid906_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
-    v_dot4_u32_u8                  :ref:`vdst<amdgpu_synid906_vdst32_0>`,     :ref:`src0<amdgpu_synid906_src32_1>`::ref:`u8x2<amdgpu_synid906_type_dev>`,  :ref:`src1<amdgpu_synid906_src32_2>`::ref:`u8x2<amdgpu_synid906_type_dev>`,  :ref:`src2<amdgpu_synid906_src32_2>`::ref:`u32<amdgpu_synid906_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
-    v_dot8_i32_i4                  :ref:`vdst<amdgpu_synid906_vdst32_0>`,     :ref:`src0<amdgpu_synid906_src32_1>`::ref:`i4x2<amdgpu_synid906_type_dev>`,  :ref:`src1<amdgpu_synid906_src32_2>`::ref:`i4x2<amdgpu_synid906_type_dev>`,  :ref:`src2<amdgpu_synid906_src32_2>`::ref:`i32<amdgpu_synid906_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
-    v_dot8_u32_u4                  :ref:`vdst<amdgpu_synid906_vdst32_0>`,     :ref:`src0<amdgpu_synid906_src32_1>`::ref:`u4x2<amdgpu_synid906_type_dev>`,  :ref:`src1<amdgpu_synid906_src32_2>`::ref:`u4x2<amdgpu_synid906_type_dev>`,  :ref:`src2<amdgpu_synid906_src32_2>`::ref:`u32<amdgpu_synid906_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
+    v_dot4_i32_i8                  :ref:`vdst<amdgpu_synid906_vdst32_0>`,     :ref:`src0<amdgpu_synid906_src32_1>`::ref:`i8x4<amdgpu_synid906_type_dev>`,  :ref:`src1<amdgpu_synid906_src32_2>`::ref:`i8x4<amdgpu_synid906_type_dev>`,  :ref:`src2<amdgpu_synid906_src32_2>`::ref:`i32<amdgpu_synid906_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
+    v_dot4_u32_u8                  :ref:`vdst<amdgpu_synid906_vdst32_0>`,     :ref:`src0<amdgpu_synid906_src32_1>`::ref:`u8x4<amdgpu_synid906_type_dev>`,  :ref:`src1<amdgpu_synid906_src32_2>`::ref:`u8x4<amdgpu_synid906_type_dev>`,  :ref:`src2<amdgpu_synid906_src32_2>`::ref:`u32<amdgpu_synid906_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
+    v_dot8_i32_i4                  :ref:`vdst<amdgpu_synid906_vdst32_0>`,     :ref:`src0<amdgpu_synid906_src32_1>`::ref:`i4x8<amdgpu_synid906_type_dev>`,  :ref:`src1<amdgpu_synid906_src32_2>`::ref:`i4x8<amdgpu_synid906_type_dev>`,  :ref:`src2<amdgpu_synid906_src32_2>`::ref:`i32<amdgpu_synid906_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
+    v_dot8_u32_u4                  :ref:`vdst<amdgpu_synid906_vdst32_0>`,     :ref:`src0<amdgpu_synid906_src32_1>`::ref:`u4x8<amdgpu_synid906_type_dev>`,  :ref:`src1<amdgpu_synid906_src32_2>`::ref:`u4x8<amdgpu_synid906_type_dev>`,  :ref:`src2<amdgpu_synid906_src32_2>`::ref:`u32<amdgpu_synid906_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
     v_fma_mix_f32                  :ref:`vdst<amdgpu_synid906_vdst32_0>`,     :ref:`src0<amdgpu_synid906_src32_1>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>`,  :ref:`src1<amdgpu_synid906_src32_2>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>`,  :ref:`src2<amdgpu_synid906_src32_2>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>`      :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
     v_fma_mixhi_f16                :ref:`vdst<amdgpu_synid906_vdst32_0>`,     :ref:`src0<amdgpu_synid906_src32_1>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>`,  :ref:`src1<amdgpu_synid906_src32_2>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>`,  :ref:`src2<amdgpu_synid906_src32_2>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>`      :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
     v_fma_mixlo_f16                :ref:`vdst<amdgpu_synid906_vdst32_0>`,     :ref:`src0<amdgpu_synid906_src32_1>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>`,  :ref:`src1<amdgpu_synid906_src32_2>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>`,  :ref:`src2<amdgpu_synid906_src32_2>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>`      :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
diff --git a/llvm/docs/AMDGPU/AMDGPUAsmGFX908.rst b/llvm/docs/AMDGPU/AMDGPUAsmGFX908.rst
index bae8f77..27da017 100644
--- a/llvm/docs/AMDGPU/AMDGPUAsmGFX908.rst
+++ b/llvm/docs/AMDGPU/AMDGPUAsmGFX908.rst
@@ -39,9 +39,9 @@ FLAT
 .. parsed-literal::
 
     **INSTRUCTION**                    **DST**             **SRC0**      **SRC1**         **SRC2**           **MODIFIERS**
-    \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
-    global_atomic_add_f32          :ref:`vdst<amdgpu_synid908_dst_flat_atomic32>`::ref:`opt<amdgpu_synid908_opt>`,       :ref:`vaddr<amdgpu_synid908_vaddr_flat_global>`,    :ref:`vdata<amdgpu_synid908_vdata32_0>`,       :ref:`saddr<amdgpu_synid908_saddr_flat_global>`          :ref:`offset13s<amdgpu_synid_flat_offset13s>`
-    global_atomic_pk_add_f16       :ref:`vdst<amdgpu_synid908_dst_flat_atomic32>`::ref:`opt<amdgpu_synid908_opt>`::ref:`f16x2<amdgpu_synid908_type_dev>`, :ref:`vaddr<amdgpu_synid908_vaddr_flat_global>`,    :ref:`vdata<amdgpu_synid908_vdata32_0>`::ref:`f16x2<amdgpu_synid908_type_dev>`, :ref:`saddr<amdgpu_synid908_saddr_flat_global>`          :ref:`offset13s<amdgpu_synid_flat_offset13s>`
+    \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
+    global_atomic_add_f32          :ref:`vdst<amdgpu_synid908_dst_flat_atomic32>`::ref:`opt<amdgpu_synid908_opt>`,       :ref:`vaddr<amdgpu_synid908_vaddr_flat_global>`,    :ref:`vdata<amdgpu_synid908_vdata32_0>`,       :ref:`saddr<amdgpu_synid908_saddr_flat_global>`          :ref:`offset13s<amdgpu_synid_flat_offset13s>` :ref:`slc<amdgpu_synid_slc>`
+    global_atomic_pk_add_f16       :ref:`vdst<amdgpu_synid908_dst_flat_atomic32>`::ref:`opt<amdgpu_synid908_opt>`::ref:`f16x2<amdgpu_synid908_type_dev>`, :ref:`vaddr<amdgpu_synid908_vaddr_flat_global>`,    :ref:`vdata<amdgpu_synid908_vdata32_0>`::ref:`f16x2<amdgpu_synid908_type_dev>`, :ref:`saddr<amdgpu_synid908_saddr_flat_global>`          :ref:`offset13s<amdgpu_synid_flat_offset13s>` :ref:`slc<amdgpu_synid_slc>`
 
 MUBUF
 -----------------------
@@ -90,40 +90,40 @@ VOP3P
 
 .. parsed-literal::
 
-    **INSTRUCTION**            **DST**          **SRC0**          **SRC1**          **SRC2**           **MODIFIERS**
-    \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
+    **INSTRUCTION**            **DST**          **SRC0**           **SRC1**           **SRC2**           **MODIFIERS**
+    \ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
     v_accvgpr_read_b32     :ref:`vdst<amdgpu_synid908_vdst32_0>`,        :ref:`asrc<amdgpu_synid908_asrc32_0>`
     v_accvgpr_write_b32    :ref:`adst<amdgpu_synid908_adst32_0>`,        :ref:`src<amdgpu_synid908_src32_3>`
-    v_dot2_f32_f16         :ref:`vdst<amdgpu_synid908_vdst32_0>`,        :ref:`src0<amdgpu_synid908_src32_1>`::ref:`f16x2<amdgpu_synid908_type_dev>`,   :ref:`src1<amdgpu_synid908_src32_2>`::ref:`f16x2<amdgpu_synid908_type_dev>`,   :ref:`src2<amdgpu_synid908_src32_2>`::ref:`f32<amdgpu_synid908_type_dev>`       :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>`
-    v_dot2_i32_i16         :ref:`vdst<amdgpu_synid908_vdst32_0>`,        :ref:`src0<amdgpu_synid908_src32_1>`::ref:`i16x2<amdgpu_synid908_type_dev>`,   :ref:`src1<amdgpu_synid908_src32_2>`::ref:`i16x2<amdgpu_synid908_type_dev>`,   :ref:`src2<amdgpu_synid908_src32_2>`::ref:`i32<amdgpu_synid908_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
-    v_dot2_u32_u16         :ref:`vdst<amdgpu_synid908_vdst32_0>`,        :ref:`src0<amdgpu_synid908_src32_1>`::ref:`u16x2<amdgpu_synid908_type_dev>`,   :ref:`src1<amdgpu_synid908_src32_2>`::ref:`u16x2<amdgpu_synid908_type_dev>`,   :ref:`src2<amdgpu_synid908_src32_2>`::ref:`u32<amdgpu_synid908_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
-    v_dot4_i32_i8          :ref:`vdst<amdgpu_synid908_vdst32_0>`,        :ref:`src0<amdgpu_synid908_src32_1>`::ref:`i8x2<amdgpu_synid908_type_dev>`,    :ref:`src1<amdgpu_synid908_src32_2>`::ref:`i8x2<amdgpu_synid908_type_dev>`,    :ref:`src2<amdgpu_synid908_src32_2>`::ref:`i32<amdgpu_synid908_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
-    v_dot4_u32_u8          :ref:`vdst<amdgpu_synid908_vdst32_0>`,        :ref:`src0<amdgpu_synid908_src32_1>`::ref:`u8x2<amdgpu_synid908_type_dev>`,    :ref:`src1<amdgpu_synid908_src32_2>`::ref:`u8x2<amdgpu_synid908_type_dev>`,    :ref:`src2<amdgpu_synid908_src32_2>`::ref:`u32<amdgpu_synid908_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
-    v_dot8_i32_i4          :ref:`vdst<amdgpu_synid908_vdst32_0>`,        :ref:`src0<amdgpu_synid908_src32_1>`::ref:`i4x2<amdgpu_synid908_type_dev>`,    :ref:`src1<amdgpu_synid908_src32_2>`::ref:`i4x2<amdgpu_synid908_type_dev>`,    :ref:`src2<amdgpu_synid908_src32_2>`::ref:`i32<amdgpu_synid908_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
-    v_dot8_u32_u4          :ref:`vdst<amdgpu_synid908_vdst32_0>`,        :ref:`src0<amdgpu_synid908_src32_1>`::ref:`u4x2<amdgpu_synid908_type_dev>`,    :ref:`src1<amdgpu_synid908_src32_2>`::ref:`u4x2<amdgpu_synid908_type_dev>`,    :ref:`src2<amdgpu_synid908_src32_2>`::ref:`u32<amdgpu_synid908_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
-    v_fma_mix_f32          :ref:`vdst<amdgpu_synid908_vdst32_0>`,        :ref:`src0<amdgpu_synid908_src32_1>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`,    :ref:`src1<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`,    :ref:`src2<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`      :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
-    v_fma_mixhi_f16        :ref:`vdst<amdgpu_synid908_vdst32_0>`,        :ref:`src0<amdgpu_synid908_src32_1>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`,    :ref:`src1<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`,    :ref:`src2<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`      :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
-    v_fma_mixlo_f16        :ref:`vdst<amdgpu_synid908_vdst32_0>`,        :ref:`src0<amdgpu_synid908_src32_1>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`,    :ref:`src1<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`,    :ref:`src2<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`      :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
-    v_mfma_f32_16x16x16f16 :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`,  :ref:`vasrc0<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`    :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
-    v_mfma_f32_16x16x1f32  :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`,   :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`,   :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`   :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
-    v_mfma_f32_16x16x2bf16 :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`   :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
-    v_mfma_f32_16x16x4f16  :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`   :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
-    v_mfma_f32_16x16x4f32  :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`,  :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`,   :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`,   :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`    :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
-    v_mfma_f32_16x16x8bf16 :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`,  :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`    :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
-    v_mfma_f32_32x32x1f32  :ref:`adst<amdgpu_synid908_adst1024_0>`::ref:`f32x32<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`,   :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`,   :ref:`asrc2<amdgpu_synid908_asrc1024_0>`::ref:`f32x32<amdgpu_synid908_type_dev>`   :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
-    v_mfma_f32_32x32x2bf16 :ref:`adst<amdgpu_synid908_adst1024_0>`::ref:`f32x32<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc1024_0>`::ref:`f32x32<amdgpu_synid908_type_dev>`   :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
-    v_mfma_f32_32x32x2f32  :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`,   :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`,   :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`   :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
-    v_mfma_f32_32x32x4bf16 :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`   :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
-    v_mfma_f32_32x32x4f16  :ref:`adst<amdgpu_synid908_adst1024_0>`::ref:`f32x32<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc1024_0>`::ref:`f32x32<amdgpu_synid908_type_dev>`   :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
-    v_mfma_f32_32x32x8f16  :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`   :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
-    v_mfma_f32_4x4x1f32    :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`,  :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`,   :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`,   :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`    :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
-    v_mfma_f32_4x4x2bf16   :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`,  :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`    :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
-    v_mfma_f32_4x4x4f16    :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`,  :ref:`vasrc0<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`    :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
-    v_mfma_i32_16x16x16i8  :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`i32x4<amdgpu_synid908_type_dev>`,  :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`,   :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`,   :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`i32x4<amdgpu_synid908_type_dev>`    :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
-    v_mfma_i32_16x16x4i8   :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`i32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`,   :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`,   :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`i32x16<amdgpu_synid908_type_dev>`   :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
-    v_mfma_i32_32x32x4i8   :ref:`adst<amdgpu_synid908_adst1024_0>`::ref:`i32x32<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`,   :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`,   :ref:`asrc2<amdgpu_synid908_asrc1024_0>`::ref:`i32x32<amdgpu_synid908_type_dev>`   :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
-    v_mfma_i32_32x32x8i8   :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`i32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`,   :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`,   :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`i32x16<amdgpu_synid908_type_dev>`   :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
-    v_mfma_i32_4x4x4i8     :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`i32x4<amdgpu_synid908_type_dev>`,  :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`,   :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`,   :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`i32x4<amdgpu_synid908_type_dev>`    :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
+    v_dot2_f32_f16         :ref:`vdst<amdgpu_synid908_vdst32_0>`,        :ref:`src0<amdgpu_synid908_src32_1>`::ref:`f16x2<amdgpu_synid908_type_dev>`,    :ref:`src1<amdgpu_synid908_src32_2>`::ref:`f16x2<amdgpu_synid908_type_dev>`,    :ref:`src2<amdgpu_synid908_src32_2>`::ref:`f32<amdgpu_synid908_type_dev>`       :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>`
+    v_dot2_i32_i16         :ref:`vdst<amdgpu_synid908_vdst32_0>`,        :ref:`src0<amdgpu_synid908_src32_1>`::ref:`i16x2<amdgpu_synid908_type_dev>`,    :ref:`src1<amdgpu_synid908_src32_2>`::ref:`i16x2<amdgpu_synid908_type_dev>`,    :ref:`src2<amdgpu_synid908_src32_2>`::ref:`i32<amdgpu_synid908_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
+    v_dot2_u32_u16         :ref:`vdst<amdgpu_synid908_vdst32_0>`,        :ref:`src0<amdgpu_synid908_src32_1>`::ref:`u16x2<amdgpu_synid908_type_dev>`,    :ref:`src1<amdgpu_synid908_src32_2>`::ref:`u16x2<amdgpu_synid908_type_dev>`,    :ref:`src2<amdgpu_synid908_src32_2>`::ref:`u32<amdgpu_synid908_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
+    v_dot4_i32_i8          :ref:`vdst<amdgpu_synid908_vdst32_0>`,        :ref:`src0<amdgpu_synid908_src32_1>`::ref:`i8x4<amdgpu_synid908_type_dev>`,     :ref:`src1<amdgpu_synid908_src32_2>`::ref:`i8x4<amdgpu_synid908_type_dev>`,     :ref:`src2<amdgpu_synid908_src32_2>`::ref:`i32<amdgpu_synid908_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
+    v_dot4_u32_u8          :ref:`vdst<amdgpu_synid908_vdst32_0>`,        :ref:`src0<amdgpu_synid908_src32_1>`::ref:`u8x4<amdgpu_synid908_type_dev>`,     :ref:`src1<amdgpu_synid908_src32_2>`::ref:`u8x4<amdgpu_synid908_type_dev>`,     :ref:`src2<amdgpu_synid908_src32_2>`::ref:`u32<amdgpu_synid908_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
+    v_dot8_i32_i4          :ref:`vdst<amdgpu_synid908_vdst32_0>`,        :ref:`src0<amdgpu_synid908_src32_1>`::ref:`i4x8<amdgpu_synid908_type_dev>`,     :ref:`src1<amdgpu_synid908_src32_2>`::ref:`i4x8<amdgpu_synid908_type_dev>`,     :ref:`src2<amdgpu_synid908_src32_2>`::ref:`i32<amdgpu_synid908_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
+    v_dot8_u32_u4          :ref:`vdst<amdgpu_synid908_vdst32_0>`,        :ref:`src0<amdgpu_synid908_src32_1>`::ref:`u4x8<amdgpu_synid908_type_dev>`,     :ref:`src1<amdgpu_synid908_src32_2>`::ref:`u4x8<amdgpu_synid908_type_dev>`,     :ref:`src2<amdgpu_synid908_src32_2>`::ref:`u32<amdgpu_synid908_type_dev>`       :ref:`clamp<amdgpu_synid_clamp>`
+    v_fma_mix_f32          :ref:`vdst<amdgpu_synid908_vdst32_0>`,        :ref:`src0<amdgpu_synid908_src32_1>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`,     :ref:`src1<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`,     :ref:`src2<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`      :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
+    v_fma_mixhi_f16        :ref:`vdst<amdgpu_synid908_vdst32_0>`,        :ref:`src0<amdgpu_synid908_src32_1>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`,     :ref:`src1<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`,     :ref:`src2<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`      :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
+    v_fma_mixlo_f16        :ref:`vdst<amdgpu_synid908_vdst32_0>`,        :ref:`src0<amdgpu_synid908_src32_1>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`,     :ref:`src1<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`,     :ref:`src2<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`      :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
+    v_mfma_f32_16x16x16f16 :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`,  :ref:`vasrc0<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`,  :ref:`vasrc1<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`,  :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`    :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
+    v_mfma_f32_16x16x1f32  :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`,    :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`,    :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`   :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
+    v_mfma_f32_16x16x2bf16 :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`bf16x2<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`bf16x2<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`   :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
+    v_mfma_f32_16x16x4f16  :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`,  :ref:`vasrc1<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`,  :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`   :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
+    v_mfma_f32_16x16x4f32  :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`,  :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`,    :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`,    :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`    :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
+    v_mfma_f32_16x16x8bf16 :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`,  :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`bf16x2<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`bf16x2<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`    :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
+    v_mfma_f32_32x32x1f32  :ref:`adst<amdgpu_synid908_adst1024_0>`::ref:`f32x32<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`,    :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`,    :ref:`asrc2<amdgpu_synid908_asrc1024_0>`::ref:`f32x32<amdgpu_synid908_type_dev>`   :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
+    v_mfma_f32_32x32x2bf16 :ref:`adst<amdgpu_synid908_adst1024_0>`::ref:`f32x32<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`bf16x2<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`bf16x2<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc1024_0>`::ref:`f32x32<amdgpu_synid908_type_dev>`   :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
+    v_mfma_f32_32x32x2f32  :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`,    :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`,    :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`   :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
+    v_mfma_f32_32x32x4bf16 :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`bf16x2<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`bf16x2<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`   :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
+    v_mfma_f32_32x32x4f16  :ref:`adst<amdgpu_synid908_adst1024_0>`::ref:`f32x32<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`,  :ref:`vasrc1<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`,  :ref:`asrc2<amdgpu_synid908_asrc1024_0>`::ref:`f32x32<amdgpu_synid908_type_dev>`   :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
+    v_mfma_f32_32x32x8f16  :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`,  :ref:`vasrc1<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`,  :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`   :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
+    v_mfma_f32_4x4x1f32    :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`,  :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`,    :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`,    :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`    :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
+    v_mfma_f32_4x4x2bf16   :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`,  :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`bf16x2<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`bf16x2<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`    :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
+    v_mfma_f32_4x4x4f16    :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`,  :ref:`vasrc0<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`,  :ref:`vasrc1<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`,  :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`    :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
+    v_mfma_i32_16x16x16i8  :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`i32x4<amdgpu_synid908_type_dev>`,  :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i8x4<amdgpu_synid908_type_dev>`,   :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i8x4<amdgpu_synid908_type_dev>`,   :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`i32x4<amdgpu_synid908_type_dev>`    :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
+    v_mfma_i32_16x16x4i8   :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`i32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i8x4<amdgpu_synid908_type_dev>`,   :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i8x4<amdgpu_synid908_type_dev>`,   :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`i32x16<amdgpu_synid908_type_dev>`   :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
+    v_mfma_i32_32x32x4i8   :ref:`adst<amdgpu_synid908_adst1024_0>`::ref:`i32x32<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i8x4<amdgpu_synid908_type_dev>`,   :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i8x4<amdgpu_synid908_type_dev>`,   :ref:`asrc2<amdgpu_synid908_asrc1024_0>`::ref:`i32x32<amdgpu_synid908_type_dev>`   :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
+    v_mfma_i32_32x32x8i8   :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`i32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i8x4<amdgpu_synid908_type_dev>`,   :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i8x4<amdgpu_synid908_type_dev>`,   :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`i32x16<amdgpu_synid908_type_dev>`   :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
+    v_mfma_i32_4x4x4i8     :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`i32x4<amdgpu_synid908_type_dev>`,  :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i8x4<amdgpu_synid908_type_dev>`,   :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i8x4<amdgpu_synid908_type_dev>`,   :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`i32x4<amdgpu_synid908_type_dev>`    :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
 
 .. |---| unicode:: U+02014 .. em dash
 
diff --git a/llvm/docs/AMDGPU/gfx1011_src32_0.rst b/llvm/docs/AMDGPU/gfx1011_src32_0.rst
new file mode 100644
index 0000000..8d82f41
--- /dev/null
+++ b/llvm/docs/AMDGPU/gfx1011_src32_0.rst
@@ -0,0 +1,17 @@
+..
+    **************************************************
+    *                                                *
+    *   Automatically generated file, do not edit!   *
+    *                                                *
+    **************************************************
+
+.. _amdgpu_synid1011_src32_0:
+
+src
+===========================
+
+Instruction input.
+
+*Size:* 1 dword.
+
+*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`s<amdgpu_synid_s>`, :ref:`vcc<amdgpu_synid_vcc>`, :ref:`ttmp<amdgpu_synid_ttmp>`, :ref:`m0<amdgpu_synid_m0>`, :ref:`exec<amdgpu_synid_exec>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`lds_direct<amdgpu_synid_lds_direct>`, :ref:`constant<amdgpu_synid_constant>`, :ref:`literal<amdgpu_synid_literal>`
diff --git a/llvm/docs/AMDGPU/gfx1011_src32_1.rst b/llvm/docs/AMDGPU/gfx1011_src32_1.rst
new file mode 100644
index 0000000..b923912
--- /dev/null
+++ b/llvm/docs/AMDGPU/gfx1011_src32_1.rst
@@ -0,0 +1,17 @@
+..
+    **************************************************
+    *                                                *
+    *   Automatically generated file, do not edit!   *
+    *                                                *
+    **************************************************
+
+.. _amdgpu_synid1011_src32_1:
+
+src
+===========================
+
+Instruction input.
+
+*Size:* 1 dword.
+
+*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`s<amdgpu_synid_s>`, :ref:`vcc<amdgpu_synid_vcc>`, :ref:`ttmp<amdgpu_synid_ttmp>`, :ref:`m0<amdgpu_synid_m0>`, :ref:`exec<amdgpu_synid_exec>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`constant<amdgpu_synid_constant>`, :ref:`literal<amdgpu_synid_literal>`
diff --git a/llvm/docs/AMDGPU/gfx1011_type_dev.rst b/llvm/docs/AMDGPU/gfx1011_type_dev.rst
new file mode 100644
index 0000000..c7ddb82
--- /dev/null
+++ b/llvm/docs/AMDGPU/gfx1011_type_dev.rst
@@ -0,0 +1,13 @@
+..
+    **************************************************
+    *                                                *
+    *   Automatically generated file, do not edit!   *
+    *                                                *
+    **************************************************
+
+.. _amdgpu_synid1011_type_dev:
+
+Type deviation
+===========================
+
+*Type* of this operand differs from *type* :ref:`implied by the opcode<amdgpu_syn_instruction_type>`. This tag specifies actual operand *type*.
diff --git a/llvm/docs/AMDGPU/gfx1011_vdst32_0.rst b/llvm/docs/AMDGPU/gfx1011_vdst32_0.rst
new file mode 100644
index 0000000..115b852
--- /dev/null
+++ b/llvm/docs/AMDGPU/gfx1011_vdst32_0.rst
@@ -0,0 +1,17 @@
+..
+    **************************************************
+    *                                                *
+    *   Automatically generated file, do not edit!   *
+    *                                                *
+    **************************************************
+
+.. _amdgpu_synid1011_vdst32_0:
+
+vdst
+===========================
+
+Instruction output.
+
+*Size:* 1 dword.
+
+*Operands:* :ref:`v<amdgpu_synid_v>`
diff --git a/llvm/docs/AMDGPU/gfx1011_vsrc32_0.rst b/llvm/docs/AMDGPU/gfx1011_vsrc32_0.rst
new file mode 100644
index 0000000..be63c65
--- /dev/null
+++ b/llvm/docs/AMDGPU/gfx1011_vsrc32_0.rst
@@ -0,0 +1,17 @@
+..
+    **************************************************
+    *                                                *
+    *   Automatically generated file, do not edit!   *
+    *                                                *
+    **************************************************
+
+.. _amdgpu_synid1011_vsrc32_0:
+
+vsrc
+===========================
+
+Instruction input.
+
+*Size:* 1 dword.
+
+*Operands:* :ref:`v<amdgpu_synid_v>`
diff --git a/llvm/docs/AMDGPU/gfx908_saddr_flat_global.rst b/llvm/docs/AMDGPU/gfx908_saddr_flat_global.rst
index a0e4506..fe9864c 100644
--- a/llvm/docs/AMDGPU/gfx908_saddr_flat_global.rst
+++ b/llvm/docs/AMDGPU/gfx908_saddr_flat_global.rst
@@ -16,4 +16,4 @@ See :ref:`vaddr<amdgpu_synid908_vaddr_flat_global>` for description of available
 
 *Size:* 2 dwords.
 
-*Operands:* :ref:`exec<amdgpu_synid_exec>`, :ref:`off<amdgpu_synid_off>`
+*Operands:* :ref:`s<amdgpu_synid_s>`, :ref:`flat_scratch<amdgpu_synid_flat_scratch>`, :ref:`vcc<amdgpu_synid_vcc>`, :ref:`ttmp<amdgpu_synid_ttmp>`, :ref:`exec<amdgpu_synid_exec>`, :ref:`off<amdgpu_synid_off>`
diff --git a/llvm/docs/AMDGPUUsage.rst b/llvm/docs/AMDGPUUsage.rst
index c153363..84761dc 100644
--- a/llvm/docs/AMDGPUUsage.rst
+++ b/llvm/docs/AMDGPUUsage.rst
@@ -5786,6 +5786,7 @@ Instructions
    AMDGPU/AMDGPUAsmGFX906
    AMDGPU/AMDGPUAsmGFX908
    AMDGPU/AMDGPUAsmGFX10
+   AMDGPU/AMDGPUAsmGFX1011
    AMDGPUModifierSyntax
    AMDGPUOperandSyntax
    AMDGPUInstructionSyntax
@@ -5806,27 +5807,27 @@ Links to detailed instruction syntax description may be found in the following
 table. Note that features under development are not included
 in this description.
 
-    ==================================== ======================================
-    Core ISA                             ISA Extensions
-    ==================================== ======================================
-    :doc:`GFX7<AMDGPU/AMDGPUAsmGFX7>`    \-
-    :doc:`GFX8<AMDGPU/AMDGPUAsmGFX8>`    \-
-    :doc:`GFX9<AMDGPU/AMDGPUAsmGFX9>`    :doc:`gfx900<AMDGPU/AMDGPUAsmGFX900>`
+    =================================== =======================================
+    Core ISA                            ISA Extensions
+    =================================== =======================================
+    :doc:`GFX7<AMDGPU/AMDGPUAsmGFX7>`   \-
+    :doc:`GFX8<AMDGPU/AMDGPUAsmGFX8>`   \-
+    :doc:`GFX9<AMDGPU/AMDGPUAsmGFX9>`   :doc:`gfx900<AMDGPU/AMDGPUAsmGFX900>`
 
-                                         :doc:`gfx902<AMDGPU/AMDGPUAsmGFX900>`
+                                        :doc:`gfx902<AMDGPU/AMDGPUAsmGFX900>`
 
-                                         :doc:`gfx904<AMDGPU/AMDGPUAsmGFX904>`
+                                        :doc:`gfx904<AMDGPU/AMDGPUAsmGFX904>`
 
-                                         :doc:`gfx906<AMDGPU/AMDGPUAsmGFX906>`
+                                        :doc:`gfx906<AMDGPU/AMDGPUAsmGFX906>`
 
-                                         :doc:`gfx908<AMDGPU/AMDGPUAsmGFX908>`
+                                        :doc:`gfx908<AMDGPU/AMDGPUAsmGFX908>`
 
-                                         :doc:`gfx909<AMDGPU/AMDGPUAsmGFX900>`
+                                        :doc:`gfx909<AMDGPU/AMDGPUAsmGFX900>`
 
-    :doc:`GFX10<AMDGPU/AMDGPUAsmGFX10>`  gfx1011
+    :doc:`GFX10<AMDGPU/AMDGPUAsmGFX10>` :doc:`gfx1011<AMDGPU/AMDGPUAsmGFX1011>`
 
-                                         gfx1012
-    ==================================== ======================================
+                                        :doc:`gfx1012<AMDGPU/AMDGPUAsmGFX1011>`
+    =================================== =======================================
 
 For more information about instructions, their semantics and supported
 combinations of operands, refer to one of instruction set architecture manuals
-- 
2.7.4