RPCS3/llvm-mirror
1
0
Fork 0
mirror of https://github.com/RPCS3/llvm-mirror.git synced 2024-10-18 10:32:48 +02:00
Code Issues Projects Releases Wiki Activity

[AMDGPU][MC][DOC] Updated AMD GPU assembler syntax description.

Browse Source 
Summary of changes:
- added description of GFX9 subtargets:
  - gfx900;
  - gfx902;
  - gfx904;
  - gfx906;
  - gfx908;
  - gfx909.
This commit is contained in:
Dmitry Preobrazhensky 2019-12-25 17:51:53 +03:00
parent 603edd8aae
commit d3bcd94780
62 changed files with 1464 additions and 109 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

163
docs/AMDGPU/AMDGPUAsmGFX10.rst
View File

@ -5,20 +5,25 @@
* *
**************************************************
============================
Syntax of GFX10 Instructions
============================
====================================================================================
Syntax of Core GFX10 Instructions
====================================================================================
.. contents::
:local:
Introduction
============
This document describes the syntax of *core* GFX10 instructions.
Notation
========
Notation used in this document is explained :ref:`here<amdgpu_syn_instruction_notation>`.
Introduction
============
Overvew
=======
An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document<amdgpu_syn_instructions>`.
@ -1359,56 +1364,56 @@ VOP2
.. parsed-literal::
**INSTRUCTION** **DST0** **DST1** **SRC0** **SRC1** **SRC2**
\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
v_add_co_ci_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>`
v_add_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_add_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_add_nc_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_and_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_ashrrev_i32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`::ref:`u32<amdgpu_synid10_type_dev>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_cndmask_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>`
v_cvt_pkrtz_f16_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_fmaak_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`imm32<amdgpu_synid10_fimm16>`
v_fmaak_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`imm32<amdgpu_synid10_fimm32>`
v_fmac_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_fmac_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_fmamk_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`imm32<amdgpu_synid10_fimm16>`, :ref:`vsrc2<amdgpu_synid10_vsrc32_0>`
v_fmamk_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`imm32<amdgpu_synid10_fimm32>`, :ref:`vsrc2<amdgpu_synid10_vsrc32_0>`
v_ldexp_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`::ref:`i16<amdgpu_synid10_type_dev>`
v_lshlrev_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`::ref:`u32<amdgpu_synid10_type_dev>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_lshrrev_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`::ref:`u32<amdgpu_synid10_type_dev>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_mac_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_mac_legacy_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_madak_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`imm32<amdgpu_synid10_fimm32>`
v_madmk_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`imm32<amdgpu_synid10_fimm32>`, :ref:`vsrc2<amdgpu_synid10_vsrc32_0>`
v_max_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_max_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_max_i32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_max_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_min_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_min_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_min_i32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_min_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_mul_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_mul_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_mul_hi_i32_i24 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_mul_hi_u32_u24 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_mul_i32_i24 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_mul_legacy_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_mul_u32_u24 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_or_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_pk_fmac_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_sub_co_ci_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>`
v_sub_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_sub_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_sub_nc_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_subrev_co_ci_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>`, :ref:`src0<amdgpu_synid10_src32_2>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>`
v_subrev_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_subrev_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_subrev_nc_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_xnor_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_xor_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
**INSTRUCTION** **DST0** **DST1** **SRC0** **SRC1** **SRC2**
\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
v_add_co_ci_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>`
v_add_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_add_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_add_nc_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_and_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_ashrrev_i32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`::ref:`u32<amdgpu_synid10_type_dev>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_cndmask_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>`
v_cvt_pkrtz_f16_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_fmaak_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`imm32<amdgpu_synid10_fimm16>`
v_fmaak_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`imm32<amdgpu_synid10_fimm32>`
v_fmac_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_fmac_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_fmamk_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`imm32<amdgpu_synid10_fimm16>`, :ref:`vsrc2<amdgpu_synid10_vsrc32_0>`
v_fmamk_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`imm32<amdgpu_synid10_fimm32>`, :ref:`vsrc2<amdgpu_synid10_vsrc32_0>`
v_ldexp_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`::ref:`i16<amdgpu_synid10_type_dev>`
v_lshlrev_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`::ref:`u32<amdgpu_synid10_type_dev>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_lshrrev_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`::ref:`u32<amdgpu_synid10_type_dev>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_mac_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_mac_legacy_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_madak_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`imm32<amdgpu_synid10_fimm32>`
v_madmk_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`imm32<amdgpu_synid10_fimm32>`, :ref:`vsrc2<amdgpu_synid10_vsrc32_0>`
v_max_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_max_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_max_i32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_max_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_min_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_min_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_min_i32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_min_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_mul_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_mul_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_mul_hi_i32_i24 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_mul_hi_u32_u24 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_mul_i32_i24 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_mul_legacy_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_mul_u32_u24 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_or_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_pk_fmac_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`::ref:`f16x2<amdgpu_synid10_type_dev>`, :ref:`src0<amdgpu_synid10_src32_1>`::ref:`f16x2<amdgpu_synid10_type_dev>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`::ref:`f16x2<amdgpu_synid10_type_dev>`
v_sub_co_ci_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>`
v_sub_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_sub_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_sub_nc_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_subrev_co_ci_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>`, :ref:`src0<amdgpu_synid10_src32_2>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`, :ref:`vcc<amdgpu_synid10_vcc_32>`
v_subrev_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_subrev_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_subrev_nc_u32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_xnor_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
v_xor_b32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`vsrc1<amdgpu_synid10_vsrc32_0>`
VOP3
-----------------------
@ -1847,30 +1852,30 @@ VOP3P
.. parsed-literal::
**INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS**
\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
v_fma_mix_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_3>`::ref:`fx<amdgpu_synid10_mad_type_dev>`, :ref:`src1<amdgpu_synid10_src32_0>`::ref:`fx<amdgpu_synid10_mad_type_dev>`, :ref:`src2<amdgpu_synid10_src32_0>`::ref:`fx<amdgpu_synid10_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_fma_mixhi_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_3>`::ref:`fx<amdgpu_synid10_mad_type_dev>`, :ref:`src1<amdgpu_synid10_src32_0>`::ref:`fx<amdgpu_synid10_mad_type_dev>`, :ref:`src2<amdgpu_synid10_src32_0>`::ref:`fx<amdgpu_synid10_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_fma_mixlo_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_3>`::ref:`fx<amdgpu_synid10_mad_type_dev>`, :ref:`src1<amdgpu_synid10_src32_0>`::ref:`fx<amdgpu_synid10_mad_type_dev>`, :ref:`src2<amdgpu_synid10_src32_0>`::ref:`fx<amdgpu_synid10_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_pk_add_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_pk_add_i16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_pk_add_u16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_pk_ashrrev_i16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`::ref:`u16x2<amdgpu_synid10_type_dev>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>`
v_pk_fma_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>`, :ref:`src2<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_pk_lshlrev_b16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`::ref:`u16x2<amdgpu_synid10_type_dev>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>`
v_pk_lshrrev_b16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`::ref:`u16x2<amdgpu_synid10_type_dev>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>`
v_pk_mad_i16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>`, :ref:`src2<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_pk_mad_u16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>`, :ref:`src2<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_pk_max_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_pk_max_i16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>`
v_pk_max_u16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>`
v_pk_min_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_pk_min_i16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>`
v_pk_min_u16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>`
v_pk_mul_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_pk_mul_lo_u16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>`
v_pk_sub_i16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_pk_sub_u16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
**INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS**
\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
v_fma_mix_f32 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_3>`::ref:`m<amdgpu_synid10_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid10_mad_type_dev>`, :ref:`src1<amdgpu_synid10_src32_0>`::ref:`m<amdgpu_synid10_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid10_mad_type_dev>`, :ref:`src2<amdgpu_synid10_src32_0>`::ref:`m<amdgpu_synid10_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid10_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_fma_mixhi_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_3>`::ref:`m<amdgpu_synid10_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid10_mad_type_dev>`, :ref:`src1<amdgpu_synid10_src32_0>`::ref:`m<amdgpu_synid10_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid10_mad_type_dev>`, :ref:`src2<amdgpu_synid10_src32_0>`::ref:`m<amdgpu_synid10_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid10_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_fma_mixlo_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_3>`::ref:`m<amdgpu_synid10_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid10_mad_type_dev>`, :ref:`src1<amdgpu_synid10_src32_0>`::ref:`m<amdgpu_synid10_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid10_mad_type_dev>`, :ref:`src2<amdgpu_synid10_src32_0>`::ref:`m<amdgpu_synid10_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid10_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_pk_add_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_pk_add_i16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_pk_add_u16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_pk_ashrrev_i16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`::ref:`u16x2<amdgpu_synid10_type_dev>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>`
v_pk_fma_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>`, :ref:`src2<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_pk_lshlrev_b16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`::ref:`u16x2<amdgpu_synid10_type_dev>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>`
v_pk_lshrrev_b16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_2>`::ref:`u16x2<amdgpu_synid10_type_dev>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>`
v_pk_mad_i16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>`, :ref:`src2<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_pk_mad_u16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>`, :ref:`src2<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_pk_max_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_pk_max_i16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>`
v_pk_max_u16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>`
v_pk_min_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_pk_min_i16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>`
v_pk_min_u16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>`
v_pk_mul_f16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_pk_mul_lo_u16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>`
v_pk_sub_i16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_pk_sub_u16 :ref:`vdst<amdgpu_synid10_vdst32_0>`, :ref:`src0<amdgpu_synid10_src32_1>`, :ref:`src1<amdgpu_synid10_src32_2>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
VOPC
-----------------------

15
docs/AMDGPU/AMDGPUAsmGFX7.rst
View File

@ -5,20 +5,25 @@
* *
**************************************************
============================
Syntax of GFX7 Instructions
============================
====================================================================================
Syntax of Core GFX7 Instructions
====================================================================================
.. contents::
:local:
Introduction
============
This document describes the syntax of *core* GFX7 instructions.
Notation
========
Notation used in this document is explained :ref:`here<amdgpu_syn_instruction_notation>`.
Introduction
============
Overvew
=======
An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document<amdgpu_syn_instructions>`.

15
docs/AMDGPU/AMDGPUAsmGFX8.rst
View File

@ -5,20 +5,25 @@
* *
**************************************************
============================
Syntax of GFX8 Instructions
============================
====================================================================================
Syntax of Core GFX8 Instructions
====================================================================================
.. contents::
:local:
Introduction
============
This document describes the syntax of *core* GFX8 instructions.
Notation
========
Notation used in this document is explained :ref:`here<amdgpu_syn_instruction_notation>`.
Introduction
============
Overvew
=======
An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document<amdgpu_syn_instructions>`.

19
docs/AMDGPU/AMDGPUAsmGFX9.rst
View File

@ -5,20 +5,25 @@
* *
**************************************************
============================
Syntax of GFX9 Instructions
============================
====================================================================================
Syntax of Core GFX9 Instructions
====================================================================================
.. contents::
:local:
Introduction
============
This document describes the syntax of *core* GFX9 instructions.
Notation
========
Notation used in this document is explained :ref:`here<amdgpu_syn_instruction_notation>`.
Introduction
============
Overvew
=======
An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document<amdgpu_syn_instructions>`.
@ -1637,9 +1642,6 @@ VOP3P
**INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS**
\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
v_mad_mix_f32 :ref:`vdst<amdgpu_synid9_vdst32_0>`, :ref:`src0<amdgpu_synid9_src32_2>`::ref:`fx<amdgpu_synid9_mad_type_dev>`, :ref:`src1<amdgpu_synid9_src32_3>`::ref:`fx<amdgpu_synid9_mad_type_dev>`, :ref:`src2<amdgpu_synid9_src32_3>`::ref:`fx<amdgpu_synid9_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_mad_mixhi_f16 :ref:`vdst<amdgpu_synid9_vdst32_0>`, :ref:`src0<amdgpu_synid9_src32_2>`::ref:`fx<amdgpu_synid9_mad_type_dev>`, :ref:`src1<amdgpu_synid9_src32_3>`::ref:`fx<amdgpu_synid9_mad_type_dev>`, :ref:`src2<amdgpu_synid9_src32_3>`::ref:`fx<amdgpu_synid9_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_mad_mixlo_f16 :ref:`vdst<amdgpu_synid9_vdst32_0>`, :ref:`src0<amdgpu_synid9_src32_2>`::ref:`fx<amdgpu_synid9_mad_type_dev>`, :ref:`src1<amdgpu_synid9_src32_3>`::ref:`fx<amdgpu_synid9_mad_type_dev>`, :ref:`src2<amdgpu_synid9_src32_3>`::ref:`fx<amdgpu_synid9_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_pk_add_f16 :ref:`vdst<amdgpu_synid9_vdst32_0>`, :ref:`src0<amdgpu_synid9_src32_2>`, :ref:`src1<amdgpu_synid9_src32_3>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_pk_add_i16 :ref:`vdst<amdgpu_synid9_vdst32_0>`, :ref:`src0<amdgpu_synid9_src32_2>`, :ref:`src1<amdgpu_synid9_src32_3>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_pk_add_u16 :ref:`vdst<amdgpu_synid9_vdst32_0>`, :ref:`src0<amdgpu_synid9_src32_2>`, :ref:`src1<amdgpu_synid9_src32_3>` :ref:`op_sel<amdgpu_synid_op_sel>` :ref:`op_sel_hi<amdgpu_synid_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
@ -2096,7 +2098,6 @@ VOPC
gfx9_vsrc32_0
gfx9_vsrc32_1
gfx9_vsrc64_0
gfx9_mad_type_dev
gfx9_mod_dpp_sdwa_abs_neg
gfx9_mod_sdwa_sext
gfx9_mod_vop3_abs_neg

58
docs/AMDGPU/AMDGPUAsmGFX900.rst Normal file
View File

@ -0,0 +1,58 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
====================================================================================
Syntax of gfx900, gfx902 and gfx909 Instructions
====================================================================================
.. contents::
:local:
Introduction
============
This document describes the syntax of *instructions specific to gfx900, gfx902 and gfx909*.
For a description of other gfx900, gfx902 and gfx909 instructions see :doc:`Syntax of Core GFX9 Instructions<AMDGPUAsmGFX9>`.
Notation
========
Notation used in this document is explained :ref:`here<amdgpu_syn_instruction_notation>`.
Overvew
=======
An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document<amdgpu_syn_instructions>`.
Instructions
============
VOP3P
-----------------------
.. parsed-literal::
**INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS**
\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
v_mad_mix_f32 :ref:`vdst<amdgpu_synid900_vdst32_0>`, :ref:`src0<amdgpu_synid900_src32_0>`::ref:`m<amdgpu_synid900_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid900_mad_type_dev>`, :ref:`src1<amdgpu_synid900_src32_1>`::ref:`m<amdgpu_synid900_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid900_mad_type_dev>`, :ref:`src2<amdgpu_synid900_src32_1>`::ref:`m<amdgpu_synid900_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid900_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_mad_mixhi_f16 :ref:`vdst<amdgpu_synid900_vdst32_0>`, :ref:`src0<amdgpu_synid900_src32_0>`::ref:`m<amdgpu_synid900_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid900_mad_type_dev>`, :ref:`src1<amdgpu_synid900_src32_1>`::ref:`m<amdgpu_synid900_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid900_mad_type_dev>`, :ref:`src2<amdgpu_synid900_src32_1>`::ref:`m<amdgpu_synid900_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid900_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_mad_mixlo_f16 :ref:`vdst<amdgpu_synid900_vdst32_0>`, :ref:`src0<amdgpu_synid900_src32_0>`::ref:`m<amdgpu_synid900_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid900_mad_type_dev>`, :ref:`src1<amdgpu_synid900_src32_1>`::ref:`m<amdgpu_synid900_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid900_mad_type_dev>`, :ref:`src2<amdgpu_synid900_src32_1>`::ref:`m<amdgpu_synid900_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid900_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
.. |---| unicode:: U+02014 .. em dash
.. toctree::
:hidden:
AMDGPUAsmGFX9
gfx900_src32_0
gfx900_src32_1
gfx900_vdst32_0
gfx900_mad_type_dev
gfx900_mod_vop3_abs_neg

58
docs/AMDGPU/AMDGPUAsmGFX904.rst Normal file
View File

@ -0,0 +1,58 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
====================================================================================
Syntax of gfx904 Instructions
====================================================================================
.. contents::
:local:
Introduction
============
This document describes the syntax of *instructions specific to gfx904*.
For a description of other gfx904 instructions see :doc:`Syntax of Core GFX9 Instructions<AMDGPUAsmGFX9>`.
Notation
========
Notation used in this document is explained :ref:`here<amdgpu_syn_instruction_notation>`.
Overvew
=======
An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document<amdgpu_syn_instructions>`.
Instructions
============
VOP3P
-----------------------
.. parsed-literal::
**INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS**
\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
v_fma_mix_f32 :ref:`vdst<amdgpu_synid904_vdst32_0>`, :ref:`src0<amdgpu_synid904_src32_0>`::ref:`m<amdgpu_synid904_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid904_mad_type_dev>`, :ref:`src1<amdgpu_synid904_src32_1>`::ref:`m<amdgpu_synid904_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid904_mad_type_dev>`, :ref:`src2<amdgpu_synid904_src32_1>`::ref:`m<amdgpu_synid904_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid904_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_fma_mixhi_f16 :ref:`vdst<amdgpu_synid904_vdst32_0>`, :ref:`src0<amdgpu_synid904_src32_0>`::ref:`m<amdgpu_synid904_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid904_mad_type_dev>`, :ref:`src1<amdgpu_synid904_src32_1>`::ref:`m<amdgpu_synid904_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid904_mad_type_dev>`, :ref:`src2<amdgpu_synid904_src32_1>`::ref:`m<amdgpu_synid904_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid904_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_fma_mixlo_f16 :ref:`vdst<amdgpu_synid904_vdst32_0>`, :ref:`src0<amdgpu_synid904_src32_0>`::ref:`m<amdgpu_synid904_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid904_mad_type_dev>`, :ref:`src1<amdgpu_synid904_src32_1>`::ref:`m<amdgpu_synid904_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid904_mad_type_dev>`, :ref:`src2<amdgpu_synid904_src32_1>`::ref:`m<amdgpu_synid904_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid904_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
.. |---| unicode:: U+02014 .. em dash
.. toctree::
:hidden:
AMDGPUAsmGFX9
gfx904_src32_0
gfx904_src32_1
gfx904_vdst32_0
gfx904_mad_type_dev
gfx904_mod_vop3_abs_neg

93
docs/AMDGPU/AMDGPUAsmGFX906.rst Normal file
View File

@ -0,0 +1,93 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
====================================================================================
Syntax of gfx906 Instructions
====================================================================================
.. contents::
:local:
Introduction
============
This document describes the syntax of *instructions specific to gfx906*.
For a description of other gfx906 instructions see :doc:`Syntax of Core GFX9 Instructions<AMDGPUAsmGFX9>`.
Notation
========
Notation used in this document is explained :ref:`here<amdgpu_syn_instruction_notation>`.
Overvew
=======
An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document<amdgpu_syn_instructions>`.
Instructions
============
VOP2
-----------------------
.. parsed-literal::
**INSTRUCTION** **DST** **SRC0** **SRC1** **MODIFIERS**
\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
v_fmac_f32 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_0>`, :ref:`vsrc1<amdgpu_synid906_vsrc32_0>`
v_fmac_f32_dpp :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`vsrc0<amdgpu_synid906_vsrc32_0>`::ref:`m<amdgpu_synid906_mod_dpp_sdwa_abs_neg>`, :ref:`vsrc1<amdgpu_synid906_vsrc32_0>`::ref:`m<amdgpu_synid906_mod_dpp_sdwa_abs_neg>` :ref:`dpp_ctrl<amdgpu_synid_dpp_ctrl>` :ref:`row_mask<amdgpu_synid_row_mask>` :ref:`bank_mask<amdgpu_synid_bank_mask>` :ref:`bound_ctrl<amdgpu_synid_bound_ctrl>`
v_xnor_b32 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_0>`, :ref:`vsrc1<amdgpu_synid906_vsrc32_0>`
v_xnor_b32_dpp :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`vsrc0<amdgpu_synid906_vsrc32_0>`, :ref:`vsrc1<amdgpu_synid906_vsrc32_0>` :ref:`dpp_ctrl<amdgpu_synid_dpp_ctrl>` :ref:`row_mask<amdgpu_synid_row_mask>` :ref:`bank_mask<amdgpu_synid_bank_mask>` :ref:`bound_ctrl<amdgpu_synid_bound_ctrl>`
v_xnor_b32_sdwa :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_0>`::ref:`m<amdgpu_synid906_mod_sdwa_sext>`, :ref:`vsrc1<amdgpu_synid906_vsrc32_0>`::ref:`m<amdgpu_synid906_mod_sdwa_sext>` :ref:`dst_sel<amdgpu_synid_dst_sel>` :ref:`dst_unused<amdgpu_synid_dst_unused>` :ref:`src0_sel<amdgpu_synid_src0_sel>` :ref:`src1_sel<amdgpu_synid_src1_sel>`
VOP3
-----------------------
.. parsed-literal::
**INSTRUCTION** **DST** **SRC0** **SRC1** **MODIFIERS**
\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
v_fmac_f32_e64 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_1>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`, :ref:`src1<amdgpu_synid906_src32_2>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>` :ref:`clamp<amdgpu_synid_clamp>` :ref:`omod<amdgpu_synid_omod>`
v_xnor_b32_e64 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_1>`, :ref:`src1<amdgpu_synid906_src32_2>`
VOP3P
-----------------------
.. parsed-literal::
**INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS**
\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
v_dot2_f32_f16 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_1>`::ref:`f16x2<amdgpu_synid906_type_dev>`, :ref:`src1<amdgpu_synid906_src32_2>`::ref:`f16x2<amdgpu_synid906_type_dev>`, :ref:`src2<amdgpu_synid906_src32_2>`::ref:`f32<amdgpu_synid906_type_dev>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_dot2_i32_i16 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_1>`::ref:`i16x2<amdgpu_synid906_type_dev>`, :ref:`src1<amdgpu_synid906_src32_2>`::ref:`i16x2<amdgpu_synid906_type_dev>`, :ref:`src2<amdgpu_synid906_src32_2>`::ref:`i32<amdgpu_synid906_type_dev>` :ref:`clamp<amdgpu_synid_clamp>`
v_dot2_u32_u16 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_1>`::ref:`u16x2<amdgpu_synid906_type_dev>`, :ref:`src1<amdgpu_synid906_src32_2>`::ref:`u16x2<amdgpu_synid906_type_dev>`, :ref:`src2<amdgpu_synid906_src32_2>`::ref:`u32<amdgpu_synid906_type_dev>` :ref:`clamp<amdgpu_synid_clamp>`
v_dot4_i32_i8 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_1>`::ref:`i8x2<amdgpu_synid906_type_dev>`, :ref:`src1<amdgpu_synid906_src32_2>`::ref:`i8x2<amdgpu_synid906_type_dev>`, :ref:`src2<amdgpu_synid906_src32_2>`::ref:`i32<amdgpu_synid906_type_dev>` :ref:`clamp<amdgpu_synid_clamp>`
v_dot4_u32_u8 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_1>`::ref:`u8x2<amdgpu_synid906_type_dev>`, :ref:`src1<amdgpu_synid906_src32_2>`::ref:`u8x2<amdgpu_synid906_type_dev>`, :ref:`src2<amdgpu_synid906_src32_2>`::ref:`u32<amdgpu_synid906_type_dev>` :ref:`clamp<amdgpu_synid_clamp>`
v_dot8_i32_i4 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_1>`::ref:`i4x2<amdgpu_synid906_type_dev>`, :ref:`src1<amdgpu_synid906_src32_2>`::ref:`i4x2<amdgpu_synid906_type_dev>`, :ref:`src2<amdgpu_synid906_src32_2>`::ref:`i32<amdgpu_synid906_type_dev>` :ref:`clamp<amdgpu_synid_clamp>`
v_dot8_u32_u4 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_1>`::ref:`u4x2<amdgpu_synid906_type_dev>`, :ref:`src1<amdgpu_synid906_src32_2>`::ref:`u4x2<amdgpu_synid906_type_dev>`, :ref:`src2<amdgpu_synid906_src32_2>`::ref:`u32<amdgpu_synid906_type_dev>` :ref:`clamp<amdgpu_synid_clamp>`
v_fma_mix_f32 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_1>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>`, :ref:`src1<amdgpu_synid906_src32_2>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>`, :ref:`src2<amdgpu_synid906_src32_2>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_fma_mixhi_f16 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_1>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>`, :ref:`src1<amdgpu_synid906_src32_2>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>`, :ref:`src2<amdgpu_synid906_src32_2>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_fma_mixlo_f16 :ref:`vdst<amdgpu_synid906_vdst32_0>`, :ref:`src0<amdgpu_synid906_src32_1>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>`, :ref:`src1<amdgpu_synid906_src32_2>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>`, :ref:`src2<amdgpu_synid906_src32_2>`::ref:`m<amdgpu_synid906_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid906_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
.. |---| unicode:: U+02014 .. em dash
.. toctree::
:hidden:
AMDGPUAsmGFX9
gfx906_src32_0
gfx906_src32_1
gfx906_src32_2
gfx906_vdst32_0
gfx906_vsrc32_0
gfx906_mad_type_dev
gfx906_mod_dpp_sdwa_abs_neg
gfx906_mod_sdwa_sext
gfx906_mod_vop3_abs_neg
gfx906_type_dev

165
docs/AMDGPU/AMDGPUAsmGFX908.rst Normal file
View File

@ -0,0 +1,165 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
====================================================================================
Syntax of gfx908 Instructions
====================================================================================
.. contents::
:local:
Introduction
============
This document describes the syntax of *instructions specific to gfx908*.
For a description of other gfx908 instructions see :doc:`Syntax of Core GFX9 Instructions<AMDGPUAsmGFX9>`.
Notation
========
Notation used in this document is explained :ref:`here<amdgpu_syn_instruction_notation>`.
Overvew
=======
An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document<amdgpu_syn_instructions>`.
Instructions
============
FLAT
-----------------------
.. parsed-literal::
**INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS**
\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
global_atomic_add_f32 :ref:`vdst<amdgpu_synid908_dst_flat_atomic32>`::ref:`opt<amdgpu_synid908_opt>`, :ref:`vaddr<amdgpu_synid908_vaddr_flat_global>`, :ref:`vdata<amdgpu_synid908_vdata32_0>`, :ref:`saddr<amdgpu_synid908_saddr_flat_global>` :ref:`offset13s<amdgpu_synid_flat_offset13s>`
global_atomic_pk_add_f16 :ref:`vdst<amdgpu_synid908_dst_flat_atomic32>`::ref:`opt<amdgpu_synid908_opt>`::ref:`f16x2<amdgpu_synid908_type_dev>`, :ref:`vaddr<amdgpu_synid908_vaddr_flat_global>`, :ref:`vdata<amdgpu_synid908_vdata32_0>`::ref:`f16x2<amdgpu_synid908_type_dev>`, :ref:`saddr<amdgpu_synid908_saddr_flat_global>` :ref:`offset13s<amdgpu_synid_flat_offset13s>`
MUBUF
-----------------------
.. parsed-literal::
**INSTRUCTION** **SRC0** **SRC1** **SRC2** **SRC3** **MODIFIERS**
\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
buffer_atomic_add_f32 :ref:`vdata<amdgpu_synid908_data_buf_atomic32>`::ref:`dst<amdgpu_synid908_ret>`, :ref:`vaddr<amdgpu_synid908_addr_buf>`, :ref:`srsrc<amdgpu_synid908_rsrc_buf>`, :ref:`soffset<amdgpu_synid908_offset_buf>` :ref:`idxen<amdgpu_synid_idxen>` :ref:`offen<amdgpu_synid_offen>` :ref:`offset12<amdgpu_synid_buf_offset12>` :ref:`slc<amdgpu_synid_slc>`
buffer_atomic_pk_add_f16 :ref:`vdata<amdgpu_synid908_data_buf_atomic32>`::ref:`dst<amdgpu_synid908_ret>`::ref:`f16x2<amdgpu_synid908_type_dev>`, :ref:`vaddr<amdgpu_synid908_addr_buf>`, :ref:`srsrc<amdgpu_synid908_rsrc_buf>`, :ref:`soffset<amdgpu_synid908_offset_buf>` :ref:`idxen<amdgpu_synid_idxen>` :ref:`offen<amdgpu_synid_offen>` :ref:`offset12<amdgpu_synid_buf_offset12>` :ref:`slc<amdgpu_synid_slc>`
VOP2
-----------------------
.. parsed-literal::
**INSTRUCTION** **DST** **SRC0** **SRC1** **MODIFIERS**
\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
v_dot2c_f32_f16 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_0>`::ref:`f16x2<amdgpu_synid908_type_dev>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>`::ref:`f16x2<amdgpu_synid908_type_dev>`
v_dot2c_f32_f16_dpp :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`vsrc0<amdgpu_synid908_vsrc32_0>`::ref:`f16x2<amdgpu_synid908_type_dev>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>`::ref:`f16x2<amdgpu_synid908_type_dev>` :ref:`dpp_ctrl<amdgpu_synid_dpp_ctrl>` :ref:`row_mask<amdgpu_synid_row_mask>` :ref:`bank_mask<amdgpu_synid_bank_mask>` :ref:`bound_ctrl<amdgpu_synid_bound_ctrl>`
v_dot2c_i32_i16 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`
v_dot2c_i32_i16_dpp :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`vsrc0<amdgpu_synid908_vsrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>` :ref:`dpp_ctrl<amdgpu_synid_dpp_ctrl>` :ref:`row_mask<amdgpu_synid_row_mask>` :ref:`bank_mask<amdgpu_synid_bank_mask>` :ref:`bound_ctrl<amdgpu_synid_bound_ctrl>`
v_dot4c_i32_i8 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_0>`::ref:`i8x4<amdgpu_synid908_type_dev>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>`::ref:`i8x4<amdgpu_synid908_type_dev>`
v_dot4c_i32_i8_dpp :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`vsrc0<amdgpu_synid908_vsrc32_0>`::ref:`i8x4<amdgpu_synid908_type_dev>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>`::ref:`i8x4<amdgpu_synid908_type_dev>` :ref:`dpp_ctrl<amdgpu_synid_dpp_ctrl>` :ref:`row_mask<amdgpu_synid_row_mask>` :ref:`bank_mask<amdgpu_synid_bank_mask>` :ref:`bound_ctrl<amdgpu_synid_bound_ctrl>`
v_dot8c_i32_i4 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_0>`::ref:`i4x8<amdgpu_synid908_type_dev>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>`::ref:`i4x8<amdgpu_synid908_type_dev>`
v_dot8c_i32_i4_dpp :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`vsrc0<amdgpu_synid908_vsrc32_0>`::ref:`i4x8<amdgpu_synid908_type_dev>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>`::ref:`i4x8<amdgpu_synid908_type_dev>` :ref:`dpp_ctrl<amdgpu_synid_dpp_ctrl>` :ref:`row_mask<amdgpu_synid_row_mask>` :ref:`bank_mask<amdgpu_synid_bank_mask>` :ref:`bound_ctrl<amdgpu_synid_bound_ctrl>`
v_fmac_f32 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_0>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>`
v_fmac_f32_dpp :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`vsrc0<amdgpu_synid908_vsrc32_0>`::ref:`m<amdgpu_synid908_mod_dpp_sdwa_abs_neg>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>`::ref:`m<amdgpu_synid908_mod_dpp_sdwa_abs_neg>` :ref:`dpp_ctrl<amdgpu_synid_dpp_ctrl>` :ref:`row_mask<amdgpu_synid_row_mask>` :ref:`bank_mask<amdgpu_synid_bank_mask>` :ref:`bound_ctrl<amdgpu_synid_bound_ctrl>`
v_pk_fmac_f16 :ref:`vdst<amdgpu_synid908_vdst32_0>`::ref:`f16x2<amdgpu_synid908_type_dev>`, :ref:`src0<amdgpu_synid908_src32_0>`::ref:`f16x2<amdgpu_synid908_type_dev>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>`::ref:`f16x2<amdgpu_synid908_type_dev>`
v_xnor_b32 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_0>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>`
v_xnor_b32_dpp :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`vsrc0<amdgpu_synid908_vsrc32_0>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>` :ref:`dpp_ctrl<amdgpu_synid_dpp_ctrl>` :ref:`row_mask<amdgpu_synid_row_mask>` :ref:`bank_mask<amdgpu_synid_bank_mask>` :ref:`bound_ctrl<amdgpu_synid_bound_ctrl>`
v_xnor_b32_sdwa :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_0>`::ref:`m<amdgpu_synid908_mod_sdwa_sext>`, :ref:`vsrc1<amdgpu_synid908_vsrc32_0>`::ref:`m<amdgpu_synid908_mod_sdwa_sext>` :ref:`dst_sel<amdgpu_synid_dst_sel>` :ref:`dst_unused<amdgpu_synid_dst_unused>` :ref:`src0_sel<amdgpu_synid_src0_sel>` :ref:`src1_sel<amdgpu_synid_src1_sel>`
VOP3
-----------------------
.. parsed-literal::
**INSTRUCTION** **DST** **SRC0** **SRC1** **MODIFIERS**
\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
v_fmac_f32_e64 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_1>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`, :ref:`src1<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>` :ref:`clamp<amdgpu_synid_clamp>` :ref:`omod<amdgpu_synid_omod>`
v_xnor_b32_e64 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_1>`, :ref:`src1<amdgpu_synid908_src32_2>`
VOP3P
-----------------------
.. parsed-literal::
**INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS**
\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
v_accvgpr_read_b32 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`asrc<amdgpu_synid908_asrc32_0>`
v_accvgpr_write_b32 :ref:`adst<amdgpu_synid908_adst32_0>`, :ref:`src<amdgpu_synid908_src32_3>`
v_dot2_f32_f16 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_1>`::ref:`f16x2<amdgpu_synid908_type_dev>`, :ref:`src1<amdgpu_synid908_src32_2>`::ref:`f16x2<amdgpu_synid908_type_dev>`, :ref:`src2<amdgpu_synid908_src32_2>`::ref:`f32<amdgpu_synid908_type_dev>` :ref:`neg_lo<amdgpu_synid_neg_lo>` :ref:`neg_hi<amdgpu_synid_neg_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_dot2_i32_i16 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_1>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`src1<amdgpu_synid908_src32_2>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`src2<amdgpu_synid908_src32_2>`::ref:`i32<amdgpu_synid908_type_dev>` :ref:`clamp<amdgpu_synid_clamp>`
v_dot2_u32_u16 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_1>`::ref:`u16x2<amdgpu_synid908_type_dev>`, :ref:`src1<amdgpu_synid908_src32_2>`::ref:`u16x2<amdgpu_synid908_type_dev>`, :ref:`src2<amdgpu_synid908_src32_2>`::ref:`u32<amdgpu_synid908_type_dev>` :ref:`clamp<amdgpu_synid_clamp>`
v_dot4_i32_i8 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_1>`::ref:`i8x2<amdgpu_synid908_type_dev>`, :ref:`src1<amdgpu_synid908_src32_2>`::ref:`i8x2<amdgpu_synid908_type_dev>`, :ref:`src2<amdgpu_synid908_src32_2>`::ref:`i32<amdgpu_synid908_type_dev>` :ref:`clamp<amdgpu_synid_clamp>`
v_dot4_u32_u8 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_1>`::ref:`u8x2<amdgpu_synid908_type_dev>`, :ref:`src1<amdgpu_synid908_src32_2>`::ref:`u8x2<amdgpu_synid908_type_dev>`, :ref:`src2<amdgpu_synid908_src32_2>`::ref:`u32<amdgpu_synid908_type_dev>` :ref:`clamp<amdgpu_synid_clamp>`
v_dot8_i32_i4 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_1>`::ref:`i4x2<amdgpu_synid908_type_dev>`, :ref:`src1<amdgpu_synid908_src32_2>`::ref:`i4x2<amdgpu_synid908_type_dev>`, :ref:`src2<amdgpu_synid908_src32_2>`::ref:`i32<amdgpu_synid908_type_dev>` :ref:`clamp<amdgpu_synid_clamp>`
v_dot8_u32_u4 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_1>`::ref:`u4x2<amdgpu_synid908_type_dev>`, :ref:`src1<amdgpu_synid908_src32_2>`::ref:`u4x2<amdgpu_synid908_type_dev>`, :ref:`src2<amdgpu_synid908_src32_2>`::ref:`u32<amdgpu_synid908_type_dev>` :ref:`clamp<amdgpu_synid_clamp>`
v_fma_mix_f32 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_1>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`, :ref:`src1<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`, :ref:`src2<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_fma_mixhi_f16 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_1>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`, :ref:`src1<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`, :ref:`src2<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_fma_mixlo_f16 :ref:`vdst<amdgpu_synid908_vdst32_0>`, :ref:`src0<amdgpu_synid908_src32_1>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`, :ref:`src1<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>`, :ref:`src2<amdgpu_synid908_src32_2>`::ref:`m<amdgpu_synid908_mod_vop3_abs_neg>`::ref:`fx<amdgpu_synid908_mad_type_dev>` :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>` :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>` :ref:`clamp<amdgpu_synid_clamp>`
v_mfma_f32_16x16x16f16 :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
v_mfma_f32_16x16x1f32 :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
v_mfma_f32_16x16x2bf16 :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
v_mfma_f32_16x16x4f16 :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
v_mfma_f32_16x16x4f32 :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
v_mfma_f32_16x16x8bf16 :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
v_mfma_f32_32x32x1f32 :ref:`adst<amdgpu_synid908_adst1024_0>`::ref:`f32x32<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc1024_0>`::ref:`f32x32<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
v_mfma_f32_32x32x2bf16 :ref:`adst<amdgpu_synid908_adst1024_0>`::ref:`f32x32<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc1024_0>`::ref:`f32x32<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
v_mfma_f32_32x32x2f32 :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
v_mfma_f32_32x32x4bf16 :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
v_mfma_f32_32x32x4f16 :ref:`adst<amdgpu_synid908_adst1024_0>`::ref:`f32x32<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc1024_0>`::ref:`f32x32<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
v_mfma_f32_32x32x8f16 :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`f32x16<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
v_mfma_f32_4x4x1f32 :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`f32<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
v_mfma_f32_4x4x2bf16 :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i16x2<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
v_mfma_f32_4x4x4f16 :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc64_0>`::ref:`f16x4<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`f32x4<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
v_mfma_i32_16x16x16i8 :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`i32x4<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`i32x4<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
v_mfma_i32_16x16x4i8 :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`i32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`i32x16<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
v_mfma_i32_32x32x4i8 :ref:`adst<amdgpu_synid908_adst1024_0>`::ref:`i32x32<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc1024_0>`::ref:`i32x32<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
v_mfma_i32_32x32x8i8 :ref:`adst<amdgpu_synid908_adst512_0>`::ref:`i32x16<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc512_0>`::ref:`i32x16<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
v_mfma_i32_4x4x4i8 :ref:`adst<amdgpu_synid908_adst128_0>`::ref:`i32x4<amdgpu_synid908_type_dev>`, :ref:`vasrc0<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`, :ref:`vasrc1<amdgpu_synid908_vasrc32_0>`::ref:`i32<amdgpu_synid908_type_dev>`, :ref:`asrc2<amdgpu_synid908_asrc128_0>`::ref:`i32x4<amdgpu_synid908_type_dev>` :ref:`cbsz<amdgpu_synid_cbsz>` :ref:`abid<amdgpu_synid_abid>` :ref:`blgp<amdgpu_synid_blgp>`
.. |---| unicode:: U+02014 .. em dash
.. toctree::
:hidden:
AMDGPUAsmGFX9
gfx908_addr_buf
gfx908_adst1024_0
gfx908_adst128_0
gfx908_adst32_0
gfx908_adst512_0
gfx908_asrc1024_0
gfx908_asrc128_0
gfx908_asrc32_0
gfx908_asrc512_0
gfx908_data_buf_atomic32
gfx908_dst_flat_atomic32
gfx908_offset_buf
gfx908_rsrc_buf
gfx908_saddr_flat_global
gfx908_src32_0
gfx908_src32_1
gfx908_src32_2
gfx908_src32_3
gfx908_vaddr_flat_global
gfx908_vasrc32_0
gfx908_vasrc64_0
gfx908_vdata32_0
gfx908_vdst32_0
gfx908_vsrc32_0
gfx908_mad_type_dev
gfx908_mod_dpp_sdwa_abs_neg
gfx908_mod_sdwa_sext
gfx908_mod_vop3_abs_neg
gfx908_opt
gfx908_ret
gfx908_type_dev

2
docs/AMDGPU/gfx9_mad_type_dev.rst → docs/AMDGPU/gfx900_mad_type_dev.rst
View File

@ -5,7 +5,7 @@
* *
**************************************************
.. _amdgpu_synid9_mad_type_dev:
.. _amdgpu_synid900_mad_type_dev:
fx
===========================

14
docs/AMDGPU/gfx900_mod_vop3_abs_neg.rst Normal file
View File

@ -0,0 +1,14 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid900_mod_vop3_abs_neg:
m
===========================
This operand may be used with floating point operand modifiers :ref:`abs<amdgpu_synid_abs>` and :ref:`neg<amdgpu_synid_neg>`.

17
docs/AMDGPU/gfx900_src32_0.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid900_src32_0:
src
===========================
Instruction input.
*Size:* 1 dword.
*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`s<amdgpu_synid_s>`, :ref:`flat_scratch<amdgpu_synid_flat_scratch>`, :ref:`vcc<amdgpu_synid_vcc>`, :ref:`ttmp<amdgpu_synid_ttmp>`, :ref:`m0<amdgpu_synid_m0>`, :ref:`exec<amdgpu_synid_exec>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`lds_direct<amdgpu_synid_lds_direct>`, :ref:`constant<amdgpu_synid_constant>`

17
docs/AMDGPU/gfx900_src32_1.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid900_src32_1:
src
===========================
Instruction input.
*Size:* 1 dword.
*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`s<amdgpu_synid_s>`, :ref:`flat_scratch<amdgpu_synid_flat_scratch>`, :ref:`vcc<amdgpu_synid_vcc>`, :ref:`ttmp<amdgpu_synid_ttmp>`, :ref:`m0<amdgpu_synid_m0>`, :ref:`exec<amdgpu_synid_exec>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`constant<amdgpu_synid_constant>`

17
docs/AMDGPU/gfx900_vdst32_0.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid900_vdst32_0:
vdst
===========================
Instruction output.
*Size:* 1 dword.
*Operands:* :ref:`v<amdgpu_synid_v>`

17
docs/AMDGPU/gfx904_mad_type_dev.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid904_mad_type_dev:
fx
===========================
This is an *f32* or *f16* operand depending on instruction modifiers:
* Operand size is controlled by :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>`.
* Location of 16-bit operand is controlled by :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>`.

14
docs/AMDGPU/gfx904_mod_vop3_abs_neg.rst Normal file
View File

@ -0,0 +1,14 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid904_mod_vop3_abs_neg:
m
===========================
This operand may be used with floating point operand modifiers :ref:`abs<amdgpu_synid_abs>` and :ref:`neg<amdgpu_synid_neg>`.

17
docs/AMDGPU/gfx904_src32_0.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid904_src32_0:
src
===========================
Instruction input.
*Size:* 1 dword.
*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`s<amdgpu_synid_s>`, :ref:`flat_scratch<amdgpu_synid_flat_scratch>`, :ref:`vcc<amdgpu_synid_vcc>`, :ref:`ttmp<amdgpu_synid_ttmp>`, :ref:`m0<amdgpu_synid_m0>`, :ref:`exec<amdgpu_synid_exec>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`lds_direct<amdgpu_synid_lds_direct>`, :ref:`constant<amdgpu_synid_constant>`

17
docs/AMDGPU/gfx904_src32_1.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid904_src32_1:
src
===========================
Instruction input.
*Size:* 1 dword.
*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`s<amdgpu_synid_s>`, :ref:`flat_scratch<amdgpu_synid_flat_scratch>`, :ref:`vcc<amdgpu_synid_vcc>`, :ref:`ttmp<amdgpu_synid_ttmp>`, :ref:`m0<amdgpu_synid_m0>`, :ref:`exec<amdgpu_synid_exec>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`constant<amdgpu_synid_constant>`

17
docs/AMDGPU/gfx904_vdst32_0.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid904_vdst32_0:
vdst
===========================
Instruction output.
*Size:* 1 dword.
*Operands:* :ref:`v<amdgpu_synid_v>`

17
docs/AMDGPU/gfx906_mad_type_dev.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid906_mad_type_dev:
fx
===========================
This is an *f32* or *f16* operand depending on instruction modifiers:
* Operand size is controlled by :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>`.
* Location of 16-bit operand is controlled by :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>`.

14
docs/AMDGPU/gfx906_mod_dpp_sdwa_abs_neg.rst Normal file
View File

@ -0,0 +1,14 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid906_mod_dpp_sdwa_abs_neg:
m
===========================
This operand may be used with floating point operand modifiers :ref:`abs<amdgpu_synid_abs>` and :ref:`neg<amdgpu_synid_neg>`.

14
docs/AMDGPU/gfx906_mod_sdwa_sext.rst Normal file
View File

@ -0,0 +1,14 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid906_mod_sdwa_sext:
m
===========================
This operand may be used with integer operand modifier :ref:`sext<amdgpu_synid_sext>`.

14
docs/AMDGPU/gfx906_mod_vop3_abs_neg.rst Normal file
View File

@ -0,0 +1,14 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid906_mod_vop3_abs_neg:
m
===========================
This operand may be used with floating point operand modifiers :ref:`abs<amdgpu_synid_abs>` and :ref:`neg<amdgpu_synid_neg>`.

17
docs/AMDGPU/gfx906_src32_0.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid906_src32_0:
src
===========================
Instruction input.
*Size:* 1 dword.
*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`s<amdgpu_synid_s>`, :ref:`flat_scratch<amdgpu_synid_flat_scratch>`, :ref:`vcc<amdgpu_synid_vcc>`, :ref:`ttmp<amdgpu_synid_ttmp>`, :ref:`m0<amdgpu_synid_m0>`, :ref:`exec<amdgpu_synid_exec>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`lds_direct<amdgpu_synid_lds_direct>`, :ref:`constant<amdgpu_synid_constant>`, :ref:`literal<amdgpu_synid_literal>`

17
docs/AMDGPU/gfx906_src32_1.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid906_src32_1:
src
===========================
Instruction input.
*Size:* 1 dword.
*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`s<amdgpu_synid_s>`, :ref:`flat_scratch<amdgpu_synid_flat_scratch>`, :ref:`vcc<amdgpu_synid_vcc>`, :ref:`ttmp<amdgpu_synid_ttmp>`, :ref:`m0<amdgpu_synid_m0>`, :ref:`exec<amdgpu_synid_exec>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`lds_direct<amdgpu_synid_lds_direct>`, :ref:`constant<amdgpu_synid_constant>`

17
docs/AMDGPU/gfx906_src32_2.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid906_src32_2:
src
===========================
Instruction input.
*Size:* 1 dword.
*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`s<amdgpu_synid_s>`, :ref:`flat_scratch<amdgpu_synid_flat_scratch>`, :ref:`vcc<amdgpu_synid_vcc>`, :ref:`ttmp<amdgpu_synid_ttmp>`, :ref:`m0<amdgpu_synid_m0>`, :ref:`exec<amdgpu_synid_exec>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`constant<amdgpu_synid_constant>`

14
docs/AMDGPU/gfx906_type_dev.rst Normal file
View File

@ -0,0 +1,14 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid906_type_dev:
Type deviation
===========================
*Type* of this operand differs from *type* :ref:`implied by the opcode<amdgpu_syn_instruction_type>`. This tag specifies actual operand *type*.

17
docs/AMDGPU/gfx906_vdst32_0.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid906_vdst32_0:
vdst
===========================
Instruction output.
*Size:* 1 dword.
*Operands:* :ref:`v<amdgpu_synid_v>`

17
docs/AMDGPU/gfx906_vsrc32_0.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid906_vsrc32_0:
vsrc
===========================
Instruction input.
*Size:* 1 dword.
*Operands:* :ref:`v<amdgpu_synid_v>`

22
docs/AMDGPU/gfx908_addr_buf.rst Normal file
View File

@ -0,0 +1,22 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_addr_buf:
vaddr
===========================
This is an optional operand which may specify offset and/or index.
*Size:* 0, 1 or 2 dwords. Size is controlled by modifiers :ref:`offen<amdgpu_synid_offen>` and :ref:`idxen<amdgpu_synid_idxen>`:
* If only :ref:`idxen<amdgpu_synid_idxen>` is specified, this operand supplies an index. Size is 1 dword.
* If only :ref:`offen<amdgpu_synid_offen>` is specified, this operand supplies an offset. Size is 1 dword.
* If both modifiers are specified, index is in the first register and offset is in the second. Size is 2 dwords.
* If none of these modifiers are specified, this operand must be set to :ref:`off<amdgpu_synid_off>`.
*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`off<amdgpu_synid_off>`

17
docs/AMDGPU/gfx908_adst1024_0.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_adst1024_0:
adst
===========================
Instruction output.
*Size:* 32 dwords.
*Operands:* :ref:`a<amdgpu_synid_a>`

17
docs/AMDGPU/gfx908_adst128_0.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_adst128_0:
adst
===========================
Instruction output.
*Size:* 4 dwords.
*Operands:* :ref:`a<amdgpu_synid_a>`

17
docs/AMDGPU/gfx908_adst32_0.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_adst32_0:
adst
===========================
Instruction output.
*Size:* 1 dword.
*Operands:* :ref:`a<amdgpu_synid_a>`

17
docs/AMDGPU/gfx908_adst512_0.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_adst512_0:
adst
===========================
Instruction output.
*Size:* 16 dwords.
*Operands:* :ref:`a<amdgpu_synid_a>`

17
docs/AMDGPU/gfx908_asrc1024_0.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_asrc1024_0:
asrc
===========================
Instruction input.
*Size:* 32 dwords.
*Operands:* :ref:`a<amdgpu_synid_a>`

17
docs/AMDGPU/gfx908_asrc128_0.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_asrc128_0:
asrc
===========================
Instruction input.
*Size:* 4 dwords.
*Operands:* :ref:`a<amdgpu_synid_a>`

17
docs/AMDGPU/gfx908_asrc32_0.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_asrc32_0:
asrc
===========================
Instruction input.
*Size:* 1 dword.
*Operands:* :ref:`a<amdgpu_synid_a>`

17
docs/AMDGPU/gfx908_asrc512_0.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_asrc512_0:
asrc
===========================
Instruction input.
*Size:* 16 dwords.
*Operands:* :ref:`a<amdgpu_synid_a>`

21
docs/AMDGPU/gfx908_data_buf_atomic32.rst Normal file
View File

@ -0,0 +1,21 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_data_buf_atomic32:
vdata
===========================
Input data for an atomic instruction.
Optionally may serve as an output data:
* If :ref:`glc<amdgpu_synid_glc>` is specified, gets the memory value before the operation.
*Size:* 1 dword by default. :ref:`tfe<amdgpu_synid_tfe>` adds 1 dword if specified.
*Operands:* :ref:`v<amdgpu_synid_v>`

19
docs/AMDGPU/gfx908_dst_flat_atomic32.rst Normal file
View File

@ -0,0 +1,19 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_dst_flat_atomic32:
vdst
===========================
Data returned by a 32-bit atomic flat instruction.
This is an optional operand. It must be used if and only if :ref:`glc<amdgpu_synid_glc>` is specified.
*Size:* 1 dword.
*Operands:* :ref:`v<amdgpu_synid_v>`

17
docs/AMDGPU/gfx908_mad_type_dev.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_mad_type_dev:
fx
===========================
This is an *f32* or *f16* operand depending on instruction modifiers:
* Operand size is controlled by :ref:`m_op_sel_hi<amdgpu_synid_mad_mix_op_sel_hi>`.
* Location of 16-bit operand is controlled by :ref:`m_op_sel<amdgpu_synid_mad_mix_op_sel>`.

14
docs/AMDGPU/gfx908_mod_dpp_sdwa_abs_neg.rst Normal file
View File

@ -0,0 +1,14 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_mod_dpp_sdwa_abs_neg:
m
===========================
This operand may be used with floating point operand modifiers :ref:`abs<amdgpu_synid_abs>` and :ref:`neg<amdgpu_synid_neg>`.

14
docs/AMDGPU/gfx908_mod_sdwa_sext.rst Normal file
View File

@ -0,0 +1,14 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_mod_sdwa_sext:
m
===========================
This operand may be used with integer operand modifier :ref:`sext<amdgpu_synid_sext>`.

14
docs/AMDGPU/gfx908_mod_vop3_abs_neg.rst Normal file
View File

@ -0,0 +1,14 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_mod_vop3_abs_neg:
m
===========================
This operand may be used with floating point operand modifiers :ref:`abs<amdgpu_synid_abs>` and :ref:`neg<amdgpu_synid_neg>`.

17
docs/AMDGPU/gfx908_offset_buf.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_offset_buf:
soffset
===========================
An unsigned byte offset.
*Size:* 1 dword.
*Operands:* :ref:`s<amdgpu_synid_s>`, :ref:`flat_scratch<amdgpu_synid_flat_scratch>`, :ref:`vcc<amdgpu_synid_vcc>`, :ref:`ttmp<amdgpu_synid_ttmp>`, :ref:`m0<amdgpu_synid_m0>`, :ref:`exec<amdgpu_synid_exec>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`constant<amdgpu_synid_constant>`

14
docs/AMDGPU/gfx908_opt.rst Normal file
View File

@ -0,0 +1,14 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_opt:
opt
===========================
This is an optional operand. It must be used if and only if :ref:`glc<amdgpu_synid_glc>` is specified.

14
docs/AMDGPU/gfx908_ret.rst Normal file
View File

@ -0,0 +1,14 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_ret:
dst
===========================
This is an input operand. It may optionally serve as a destination if :ref:`glc<amdgpu_synid_glc>` is specified.

17
docs/AMDGPU/gfx908_rsrc_buf.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_rsrc_buf:
srsrc
===========================
Buffer resource constant which defines the address and characteristics of the buffer in memory.
*Size:* 4 dwords.
*Operands:* :ref:`s<amdgpu_synid_s>`, :ref:`ttmp<amdgpu_synid_ttmp>`

19
docs/AMDGPU/gfx908_saddr_flat_global.rst Normal file
View File

@ -0,0 +1,19 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_saddr_flat_global:
saddr
===========================
An optional 64-bit flat global address. Must be specified as :ref:`off<amdgpu_synid_off>` if not used.
See :ref:`vaddr<amdgpu_synid908_vaddr_flat_global>` for description of available addressing modes.
*Size:* 2 dwords.
*Operands:* :ref:`exec<amdgpu_synid_exec>`, :ref:`off<amdgpu_synid_off>`

17
docs/AMDGPU/gfx908_src32_0.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_src32_0:
src
===========================
Instruction input.
*Size:* 1 dword.
*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`s<amdgpu_synid_s>`, :ref:`flat_scratch<amdgpu_synid_flat_scratch>`, :ref:`vcc<amdgpu_synid_vcc>`, :ref:`ttmp<amdgpu_synid_ttmp>`, :ref:`m0<amdgpu_synid_m0>`, :ref:`exec<amdgpu_synid_exec>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`lds_direct<amdgpu_synid_lds_direct>`, :ref:`constant<amdgpu_synid_constant>`, :ref:`literal<amdgpu_synid_literal>`

17
docs/AMDGPU/gfx908_src32_1.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_src32_1:
src
===========================
Instruction input.
*Size:* 1 dword.
*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`s<amdgpu_synid_s>`, :ref:`flat_scratch<amdgpu_synid_flat_scratch>`, :ref:`vcc<amdgpu_synid_vcc>`, :ref:`ttmp<amdgpu_synid_ttmp>`, :ref:`m0<amdgpu_synid_m0>`, :ref:`exec<amdgpu_synid_exec>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`lds_direct<amdgpu_synid_lds_direct>`, :ref:`constant<amdgpu_synid_constant>`

17
docs/AMDGPU/gfx908_src32_2.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_src32_2:
src
===========================
Instruction input.
*Size:* 1 dword.
*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`s<amdgpu_synid_s>`, :ref:`flat_scratch<amdgpu_synid_flat_scratch>`, :ref:`vcc<amdgpu_synid_vcc>`, :ref:`ttmp<amdgpu_synid_ttmp>`, :ref:`m0<amdgpu_synid_m0>`, :ref:`exec<amdgpu_synid_exec>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`constant<amdgpu_synid_constant>`

17
docs/AMDGPU/gfx908_src32_3.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_src32_3:
src
===========================
Instruction input.
*Size:* 1 dword.
*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`vccz<amdgpu_synid_vccz>`, :ref:`execz<amdgpu_synid_execz>`, :ref:`scc<amdgpu_synid_scc>`, :ref:`constant<amdgpu_synid_constant>`

14
docs/AMDGPU/gfx908_type_dev.rst Normal file
View File

@ -0,0 +1,14 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_type_dev:
Type deviation
===========================
*Type* of this operand differs from *type* :ref:`implied by the opcode<amdgpu_syn_instruction_type>`. This tag specifies actual operand *type*.

22
docs/AMDGPU/gfx908_vaddr_flat_global.rst Normal file
View File

@ -0,0 +1,22 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_vaddr_flat_global:
vaddr
===========================
A 64-bit flat global address or a 32-bit offset depending on addressing mode:
* Address = :ref:`vaddr<amdgpu_synid908_vaddr_flat_global>` + :ref:`offset13s<amdgpu_synid_flat_offset13s>`. :ref:`vaddr<amdgpu_synid908_vaddr_flat_global>` is a 64-bit address. This mode is indicated by :ref:`saddr<amdgpu_synid908_saddr_flat_global>` set to :ref:`off<amdgpu_synid_off>`.
* Address = :ref:`saddr<amdgpu_synid908_saddr_flat_global>` + :ref:`vaddr<amdgpu_synid908_vaddr_flat_global>` + :ref:`offset13s<amdgpu_synid_flat_offset13s>`. :ref:`vaddr<amdgpu_synid908_vaddr_flat_global>` is a 32-bit offset. This mode is used when :ref:`saddr<amdgpu_synid908_saddr_flat_global>` is not :ref:`off<amdgpu_synid_off>`.
.. WARNING:: Assembler currently expects a 64-bit *vaddr* regardless of addressing mode. This have to be fixed.
*Size:* 1 or 2 dwords.
*Operands:* :ref:`v<amdgpu_synid_v>`

17
docs/AMDGPU/gfx908_vasrc32_0.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_vasrc32_0:
vasrc
===========================
Instruction input.
*Size:* 1 dword.
*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`a<amdgpu_synid_a>`

17
docs/AMDGPU/gfx908_vasrc64_0.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_vasrc64_0:
vasrc
===========================
Instruction input.
*Size:* 2 dwords.
*Operands:* :ref:`v<amdgpu_synid_v>`, :ref:`a<amdgpu_synid_a>`

17
docs/AMDGPU/gfx908_vdata32_0.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_vdata32_0:
vdata
===========================
Instruction input.
*Size:* 1 dword.
*Operands:* :ref:`v<amdgpu_synid_v>`

17
docs/AMDGPU/gfx908_vdst32_0.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_vdst32_0:
vdst
===========================
Instruction output.
*Size:* 1 dword.
*Operands:* :ref:`v<amdgpu_synid_v>`

17
docs/AMDGPU/gfx908_vsrc32_0.rst Normal file
View File

@ -0,0 +1,17 @@
..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
.. _amdgpu_synid908_vsrc32_0:
vsrc
===========================
Instruction input.
*Size:* 1 dword.
*Operands:* :ref:`v<amdgpu_synid_v>`

55
docs/AMDGPUModifierSyntax.rst
View File

@ -1337,7 +1337,7 @@ VOP3P Modifiers
This section describes modifiers of *regular* VOP3P instructions.
*v_mad_mix_f32*, *v_mad_mixhi_f16* and *v_mad_mixlo_f16*
*v_mad_mix\** and *v_fma_mix\**
instructions use these modifiers :ref:`in a special manner<amdgpu_synid_mad_mix>`.
GFX9 and GFX10 only.
@ -1494,8 +1494,8 @@ See a description :ref:`here<amdgpu_synid_clamp>`.
VOP3P V_MAD_MIX Modifiers
-------------------------
*v_mad_mix_f32*, *v_mad_mixhi_f16* and *v_mad_mixlo_f16* instructions
use *op_sel* and *op_sel_hi* modifiers
*v_mad_mix\** and *v_fma_mix\**
instructions use *op_sel* and *op_sel_hi* modifiers
in a manner different from *regular* VOP3P instructions.
See a description below.
@ -1581,3 +1581,52 @@ clamp
~~~~~
See a description :ref:`here<amdgpu_synid_clamp>`.
VOP3P MFMA Modifiers
--------------------
.. _amdgpu_synid_cbsz:
cbsz
~~~~
=============================== ==================================================================
Syntax Description
=============================== ==================================================================
cbsz:[{0..7}] TBD
=============================== ==================================================================
Note: numeric value may be specified as either
an :ref:`integer number<amdgpu_synid_integer_number>` or
an :ref:`absolute expression<amdgpu_synid_absolute_expression>`.
.. _amdgpu_synid_abid:
abid
~~~~
=============================== ==================================================================
Syntax Description
=============================== ==================================================================
abid:[{0..15}] TBD
=============================== ==================================================================
Note: numeric value may be specified as either
an :ref:`integer number<amdgpu_synid_integer_number>` or
an :ref:`absolute expression<amdgpu_synid_absolute_expression>`.
.. _amdgpu_synid_blgp:
blgp
~~~~
=============================== ==================================================================
Syntax Description
=============================== ==================================================================
blgp:[{0..7}] TBD
=============================== ==================================================================
Note: numeric value may be specified as either
an :ref:`integer number<amdgpu_synid_integer_number>` or
an :ref:`absolute expression<amdgpu_synid_absolute_expression>`.

60
docs/AMDGPUOperandSyntax.rst
View File

@ -102,6 +102,66 @@ Examples:
[v[32],v[1:1],[v2]]
[v4,v4,v4,v4]
.. _amdgpu_synid_a:
a
-
Accumulator registers. There are 256 32-bit accumulator registers.
A sequence of *accumulator* registers may be used to operate with more than 32 bits of data.
Assembler currently supports sequences of 1, 2, 4 and 16 *accumulator* registers.
=================================================== ========================================================= ====================================================================
Syntax An Alternative Syntax (SP3) Description
=================================================== ========================================================= ====================================================================
**a**\<N> **acc**\<N> A single 32-bit *accumulator* register.
*N* must be a decimal
:ref:`integer number<amdgpu_synid_integer_number>`.
**a[**\ <N>\ **]** **acc[**\ <N>\ **]** A single 32-bit *accumulator* register.
*N* may be specified as an
:ref:`integer number<amdgpu_synid_integer_number>`
or an :ref:`absolute expression<amdgpu_synid_absolute_expression>`.
**a[**\ <N>:<K>\ **]** **acc[**\ <N>:<K>\ **]** A sequence of (\ *K-N+1*\ ) *accumulator* registers.
*N* and *K* may be specified as
:ref:`integer numbers<amdgpu_synid_integer_number>`
or :ref:`absolute expressions<amdgpu_synid_absolute_expression>`.
**[a**\ <N>, \ **a**\ <N+1>, ... **a**\ <K>\ **]** **[acc**\ <N>, \ **acc**\ <N+1>, ... **acc**\ <K>\ **]** A sequence of (\ *K-N+1*\ ) *accumulator* registers.
Register indices must be specified as decimal
:ref:`integer numbers<amdgpu_synid_integer_number>`.
=================================================== ========================================================= ====================================================================
Note: *N* and *K* must satisfy the following conditions:
* *N* <= *K*.
* 0 <= *N* <= 255.
* 0 <= *K* <= 255.
* *K-N+1* must be equal to 1, 2, 4 or 16.
Examples:
.. parsed-literal::
a255
a[0]
a[0:1]
a[1:1]
a[0:3]
a[2*2]
a[1-1:2-1]
[a252]
[a252,a253,a254,a255]
acc0
acc[1]
[acc250]
[acc2,acc3]
.. _amdgpu_synid_s:
s

38
docs/AMDGPUUsage.rst
View File

@ -5781,6 +5781,10 @@ Instructions
AMDGPU/AMDGPUAsmGFX7
AMDGPU/AMDGPUAsmGFX8
AMDGPU/AMDGPUAsmGFX9
AMDGPU/AMDGPUAsmGFX900
AMDGPU/AMDGPUAsmGFX904
AMDGPU/AMDGPUAsmGFX906
AMDGPU/AMDGPUAsmGFX908
AMDGPU/AMDGPUAsmGFX10
AMDGPUModifierSyntax
AMDGPUOperandSyntax
@ -5792,17 +5796,37 @@ An instruction has the following :doc:`syntax<AMDGPUInstructionSyntax>`:
| ``<``\ *opcode*\ ``> <``\ *operand0*\ ``>, <``\ *operand1*\ ``>,...
<``\ *modifier0*\ ``> <``\ *modifier1*\ ``>...``
:doc:`Operands<AMDGPUOperandSyntax>` are normally comma-separated while
:doc:`Operands<AMDGPUOperandSyntax>` are comma-separated while
:doc:`modifiers<AMDGPUModifierSyntax>` are space-separated.
The order of *operands* and *modifiers* is fixed.
Most *modifiers* are optional and may be omitted.
The order of operands and modifiers is fixed.
Most modifiers are optional and may be omitted.
See detailed instruction syntax description for
:doc:`GFX7<AMDGPU/AMDGPUAsmGFX7>`, :doc:`GFX8<AMDGPU/AMDGPUAsmGFX8>`,
:doc:`GFX9<AMDGPU/AMDGPUAsmGFX9>`, and :doc:`GFX10<AMDGPU/AMDGPUAsmGFX10>`.
Links to detailed instruction syntax description may be found in the following
table. Note that features under development are not included
in this description.
Note that features under development are not included in this description.
==================================== ======================================
Core ISA ISA Extensions
==================================== ======================================
:doc:`GFX7<AMDGPU/AMDGPUAsmGFX7>` \-
:doc:`GFX8<AMDGPU/AMDGPUAsmGFX8>` \-
:doc:`GFX9<AMDGPU/AMDGPUAsmGFX9>` :doc:`gfx900<AMDGPU/AMDGPUAsmGFX900>`
:doc:`gfx902<AMDGPU/AMDGPUAsmGFX900>`
:doc:`gfx904<AMDGPU/AMDGPUAsmGFX904>`
:doc:`gfx906<AMDGPU/AMDGPUAsmGFX906>`
:doc:`gfx908<AMDGPU/AMDGPUAsmGFX908>`
:doc:`gfx909<AMDGPU/AMDGPUAsmGFX900>`
:doc:`GFX10<AMDGPU/AMDGPUAsmGFX10>` gfx1011
gfx1012
==================================== ======================================
For more information about instructions, their semantics and supported
combinations of operands, refer to one of instruction set architecture manuals