mirror of
https://github.com/llvm/llvm-project.git
synced 2025-04-26 08:36:07 +00:00
[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.
This commit is contained in:
Dmitry Preobrazhensky
parent
5a21d2bbb9
commit
2de2275cbd
92
llvm/docs/AMDGPU/AMDGPUAsmGFX1011.rst
Normal file
92
llvm/docs/AMDGPU/AMDGPUAsmGFX1011.rst
Normal file
@ -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
|
||||||
@ -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_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_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_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_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:`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_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:`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_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:`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_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_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_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>`
|
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>`
|
||||||
|
|||||||
@ -39,9 +39,9 @@ FLAT
|
|||||||
.. parsed-literal::
|
.. parsed-literal::
|
||||||
|
|
||||||
**INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS**
|
**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_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>`
|
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
|
MUBUF
|
||||||
-----------------------
|
-----------------------
|
||||||
@ -90,40 +90,40 @@ VOP3P
|
|||||||
|
|
||||||
.. parsed-literal::
|
.. 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_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_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_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_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_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_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:`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_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:`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_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:`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_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_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_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_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_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_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_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_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_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_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_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_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_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_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_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_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_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_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_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_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:`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_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:`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_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:`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_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:`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_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
|
.. |---| unicode:: U+02014 .. em dash
|
||||||
|
|
||||||
|
|||||||
17
llvm/docs/AMDGPU/gfx1011_src32_0.rst
Normal file
17
llvm/docs/AMDGPU/gfx1011_src32_0.rst
Normal file
@ -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>`
|
||||||
17
llvm/docs/AMDGPU/gfx1011_src32_1.rst
Normal file
17
llvm/docs/AMDGPU/gfx1011_src32_1.rst
Normal file
@ -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>`
|
||||||
13
llvm/docs/AMDGPU/gfx1011_type_dev.rst
Normal file
13
llvm/docs/AMDGPU/gfx1011_type_dev.rst
Normal file
@ -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*.
|
||||||
17
llvm/docs/AMDGPU/gfx1011_vdst32_0.rst
Normal file
17
llvm/docs/AMDGPU/gfx1011_vdst32_0.rst
Normal file
@ -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>`
|
||||||
17
llvm/docs/AMDGPU/gfx1011_vsrc32_0.rst
Normal file
17
llvm/docs/AMDGPU/gfx1011_vsrc32_0.rst
Normal file
@ -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>`
|
||||||
@ -16,4 +16,4 @@ See :ref:`vaddr<amdgpu_synid908_vaddr_flat_global>` for description of available
|
|||||||
|
|
||||||
*Size:* 2 dwords.
|
*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>`
|
||||||
|
|||||||
@ -5786,6 +5786,7 @@ Instructions
|
|||||||
AMDGPU/AMDGPUAsmGFX906
|
AMDGPU/AMDGPUAsmGFX906
|
||||||
AMDGPU/AMDGPUAsmGFX908
|
AMDGPU/AMDGPUAsmGFX908
|
||||||
AMDGPU/AMDGPUAsmGFX10
|
AMDGPU/AMDGPUAsmGFX10
|
||||||
|
AMDGPU/AMDGPUAsmGFX1011
|
||||||
AMDGPUModifierSyntax
|
AMDGPUModifierSyntax
|
||||||
AMDGPUOperandSyntax
|
AMDGPUOperandSyntax
|
||||||
AMDGPUInstructionSyntax
|
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
|
table. Note that features under development are not included
|
||||||
in this description.
|
in this description.
|
||||||
|
|
||||||
==================================== ======================================
|
=================================== =======================================
|
||||||
Core ISA ISA Extensions
|
Core ISA ISA Extensions
|
||||||
==================================== ======================================
|
=================================== =======================================
|
||||||
:doc:`GFX7<AMDGPU/AMDGPUAsmGFX7>` \-
|
:doc:`GFX7<AMDGPU/AMDGPUAsmGFX7>` \-
|
||||||
:doc:`GFX8<AMDGPU/AMDGPUAsmGFX8>` \-
|
:doc:`GFX8<AMDGPU/AMDGPUAsmGFX8>` \-
|
||||||
:doc:`GFX9<AMDGPU/AMDGPUAsmGFX9>` :doc:`gfx900<AMDGPU/AMDGPUAsmGFX900>`
|
: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
|
For more information about instructions, their semantics and supported
|
||||||
combinations of operands, refer to one of instruction set architecture manuals
|
combinations of operands, refer to one of instruction set architecture manuals
|
||||||
|
|||||||
Loading…
x
Reference in New Issue
Block a user