..
**************************************************
* *
* Automatically generated file, do not edit! *
* *
**************************************************
====================================================================================
Syntax of gfx908 Instructions
====================================================================================
.. contents::
:local:
Introduction
============
This document describes the syntax of *instructions specific to gfx908*.
For a description of other gfx908 instructions see :doc:`Syntax of Core GFX9 Instructions`.
Notation
========
Notation used in this document is explained :ref:`here`.
Overview
========
An overview of generic syntax and other features of AMDGPU instructions may be found :ref:`in this document`.
Instructions
============
FLAT
-----------------------
.. parsed-literal::
**INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS**
\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
global_atomic_add_f32 :ref:`vdst`::ref:`opt`, :ref:`vaddr`, :ref:`vdata`, :ref:`saddr` :ref:`offset13s` :ref:`slc`
global_atomic_pk_add_f16 :ref:`vdst`::ref:`opt`::ref:`f16x2`, :ref:`vaddr`, :ref:`vdata`::ref:`f16x2`, :ref:`saddr` :ref:`offset13s` :ref:`slc`
MUBUF
-----------------------
.. parsed-literal::
**INSTRUCTION** **SRC0** **SRC1** **SRC2** **SRC3** **MODIFIERS**
\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
buffer_atomic_add_f32 :ref:`vdata`::ref:`dst`, :ref:`vaddr`, :ref:`srsrc`, :ref:`soffset` :ref:`idxen` :ref:`offen` :ref:`offset12` :ref:`slc`
buffer_atomic_pk_add_f16 :ref:`vdata`::ref:`dst`::ref:`f16x2`, :ref:`vaddr`, :ref:`srsrc`, :ref:`soffset` :ref:`idxen` :ref:`offen` :ref:`offset12` :ref:`slc`
VOP2
-----------------------
.. parsed-literal::
**INSTRUCTION** **DST** **SRC0** **SRC1** **MODIFIERS**
\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
v_dot2c_f32_f16 :ref:`vdst`, :ref:`src0`::ref:`f16x2`, :ref:`vsrc1`::ref:`f16x2`
v_dot2c_f32_f16_dpp :ref:`vdst`, :ref:`vsrc0`::ref:`f16x2`, :ref:`vsrc1`::ref:`f16x2` :ref:`dpp_ctrl` :ref:`row_mask` :ref:`bank_mask` :ref:`bound_ctrl`
v_dot2c_i32_i16 :ref:`vdst`, :ref:`src0`::ref:`i16x2`, :ref:`vsrc1`::ref:`i16x2`
v_dot2c_i32_i16_dpp :ref:`vdst`, :ref:`vsrc0`::ref:`i16x2`, :ref:`vsrc1`::ref:`i16x2` :ref:`dpp_ctrl` :ref:`row_mask` :ref:`bank_mask` :ref:`bound_ctrl`
v_dot4c_i32_i8 :ref:`vdst`, :ref:`src0`::ref:`i8x4`, :ref:`vsrc1`::ref:`i8x4`
v_dot4c_i32_i8_dpp :ref:`vdst`, :ref:`vsrc0`::ref:`i8x4`, :ref:`vsrc1`::ref:`i8x4` :ref:`dpp_ctrl` :ref:`row_mask` :ref:`bank_mask` :ref:`bound_ctrl`
v_dot8c_i32_i4 :ref:`vdst`, :ref:`src0`::ref:`i4x8`, :ref:`vsrc1`::ref:`i4x8`
v_dot8c_i32_i4_dpp :ref:`vdst`, :ref:`vsrc0`::ref:`i4x8`, :ref:`vsrc1`::ref:`i4x8` :ref:`dpp_ctrl` :ref:`row_mask` :ref:`bank_mask` :ref:`bound_ctrl`
v_fmac_f32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1`
v_fmac_f32_dpp :ref:`vdst`, :ref:`vsrc0`::ref:`m`, :ref:`vsrc1`::ref:`m` :ref:`dpp_ctrl` :ref:`row_mask` :ref:`bank_mask` :ref:`bound_ctrl`
v_pk_fmac_f16 :ref:`vdst`::ref:`f16x2`, :ref:`src0`::ref:`f16x2`, :ref:`vsrc1`::ref:`f16x2`
v_xnor_b32 :ref:`vdst`, :ref:`src0`, :ref:`vsrc1`
v_xnor_b32_dpp :ref:`vdst`, :ref:`vsrc0`, :ref:`vsrc1` :ref:`dpp_ctrl` :ref:`row_mask` :ref:`bank_mask` :ref:`bound_ctrl`
v_xnor_b32_sdwa :ref:`vdst`, :ref:`src0`::ref:`m`, :ref:`vsrc1`::ref:`m` :ref:`dst_sel` :ref:`dst_unused` :ref:`src0_sel` :ref:`src1_sel`
VOP3
-----------------------
.. parsed-literal::
**INSTRUCTION** **DST** **SRC0** **SRC1** **MODIFIERS**
\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
v_fmac_f32_e64 :ref:`vdst`, :ref:`src0`::ref:`m`, :ref:`src1`::ref:`m` :ref:`clamp` :ref:`omod`
v_xnor_b32_e64 :ref:`vdst`, :ref:`src0`, :ref:`src1`
VOP3P
-----------------------
.. parsed-literal::
**INSTRUCTION** **DST** **SRC0** **SRC1** **SRC2** **MODIFIERS**
\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|\ |---|
v_accvgpr_read_b32 :ref:`vdst`, :ref:`asrc`
v_accvgpr_write_b32 :ref:`adst`, :ref:`src`
v_dot2_f32_f16 :ref:`vdst`, :ref:`src0`::ref:`f16x2`, :ref:`src1`::ref:`f16x2`, :ref:`src2`::ref:`f32` :ref:`neg_lo` :ref:`neg_hi` :ref:`clamp`
v_dot2_i32_i16 :ref:`vdst`, :ref:`src0`::ref:`i16x2`, :ref:`src1`::ref:`i16x2`, :ref:`src2`::ref:`i32` :ref:`clamp`
v_dot2_u32_u16 :ref:`vdst`, :ref:`src0`::ref:`u16x2`, :ref:`src1`::ref:`u16x2`, :ref:`src2`::ref:`u32` :ref:`clamp`
v_dot4_i32_i8 :ref:`vdst`, :ref:`src0`::ref:`i8x4`, :ref:`src1`::ref:`i8x4`, :ref:`src2`::ref:`i32` :ref:`clamp`
v_dot4_u32_u8 :ref:`vdst`, :ref:`src0`::ref:`u8x4`, :ref:`src1`::ref:`u8x4`, :ref:`src2`::ref:`u32` :ref:`clamp`
v_dot8_i32_i4 :ref:`vdst`, :ref:`src0`::ref:`i4x8`, :ref:`src1`::ref:`i4x8`, :ref:`src2`::ref:`i32` :ref:`clamp`
v_dot8_u32_u4 :ref:`vdst`, :ref:`src0`::ref:`u4x8`, :ref:`src1`::ref:`u4x8`, :ref:`src2`::ref:`u32` :ref:`clamp`
v_fma_mix_f32 :ref:`vdst`, :ref:`src0`::ref:`m`::ref:`fx`, :ref:`src1`::ref:`m`::ref:`fx`, :ref:`src2`::ref:`m`::ref:`fx` :ref:`m_op_sel` :ref:`m_op_sel_hi` :ref:`clamp`
v_fma_mixhi_f16 :ref:`vdst`, :ref:`src0`::ref:`m`::ref:`fx`, :ref:`src1`::ref:`m`::ref:`fx`, :ref:`src2`::ref:`m`::ref:`fx` :ref:`m_op_sel` :ref:`m_op_sel_hi` :ref:`clamp`
v_fma_mixlo_f16 :ref:`vdst`, :ref:`src0`::ref:`m`::ref:`fx`, :ref:`src1`::ref:`m`::ref:`fx`, :ref:`src2`::ref:`m`::ref:`fx` :ref:`m_op_sel` :ref:`m_op_sel_hi` :ref:`clamp`
v_mfma_f32_16x16x16f16 :ref:`adst`::ref:`f32x4`, :ref:`vasrc0`::ref:`f16x4`, :ref:`vasrc1`::ref:`f16x4`, :ref:`asrc2`::ref:`f32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp`
v_mfma_f32_16x16x1f32 :ref:`adst`::ref:`f32x16`, :ref:`vasrc0`::ref:`f32`, :ref:`vasrc1`::ref:`f32`, :ref:`asrc2`::ref:`f32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp`
v_mfma_f32_16x16x2bf16 :ref:`adst`::ref:`f32x16`, :ref:`vasrc0`::ref:`bf16x2`, :ref:`vasrc1`::ref:`bf16x2`, :ref:`asrc2`::ref:`f32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp`
v_mfma_f32_16x16x4f16 :ref:`adst`::ref:`f32x16`, :ref:`vasrc0`::ref:`f16x4`, :ref:`vasrc1`::ref:`f16x4`, :ref:`asrc2`::ref:`f32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp`
v_mfma_f32_16x16x4f32 :ref:`adst`::ref:`f32x4`, :ref:`vasrc0`::ref:`f32`, :ref:`vasrc1`::ref:`f32`, :ref:`asrc2`::ref:`f32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp`
v_mfma_f32_16x16x8bf16 :ref:`adst`::ref:`f32x4`, :ref:`vasrc0`::ref:`bf16x2`, :ref:`vasrc1`::ref:`bf16x2`, :ref:`asrc2`::ref:`f32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp`
v_mfma_f32_32x32x1f32 :ref:`adst`::ref:`f32x32`, :ref:`vasrc0`::ref:`f32`, :ref:`vasrc1`::ref:`f32`, :ref:`asrc2`::ref:`f32x32` :ref:`cbsz` :ref:`abid` :ref:`blgp`
v_mfma_f32_32x32x2bf16 :ref:`adst`::ref:`f32x32`, :ref:`vasrc0`::ref:`bf16x2`, :ref:`vasrc1`::ref:`bf16x2`, :ref:`asrc2`::ref:`f32x32` :ref:`cbsz` :ref:`abid` :ref:`blgp`
v_mfma_f32_32x32x2f32 :ref:`adst`::ref:`f32x16`, :ref:`vasrc0`::ref:`f32`, :ref:`vasrc1`::ref:`f32`, :ref:`asrc2`::ref:`f32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp`
v_mfma_f32_32x32x4bf16 :ref:`adst`::ref:`f32x16`, :ref:`vasrc0`::ref:`bf16x2`, :ref:`vasrc1`::ref:`bf16x2`, :ref:`asrc2`::ref:`f32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp`
v_mfma_f32_32x32x4f16 :ref:`adst`::ref:`f32x32`, :ref:`vasrc0`::ref:`f16x4`, :ref:`vasrc1`::ref:`f16x4`, :ref:`asrc2`::ref:`f32x32` :ref:`cbsz` :ref:`abid` :ref:`blgp`
v_mfma_f32_32x32x8f16 :ref:`adst`::ref:`f32x16`, :ref:`vasrc0`::ref:`f16x4`, :ref:`vasrc1`::ref:`f16x4`, :ref:`asrc2`::ref:`f32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp`
v_mfma_f32_4x4x1f32 :ref:`adst`::ref:`f32x4`, :ref:`vasrc0`::ref:`f32`, :ref:`vasrc1`::ref:`f32`, :ref:`asrc2`::ref:`f32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp`
v_mfma_f32_4x4x2bf16 :ref:`adst`::ref:`f32x4`, :ref:`vasrc0`::ref:`bf16x2`, :ref:`vasrc1`::ref:`bf16x2`, :ref:`asrc2`::ref:`f32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp`
v_mfma_f32_4x4x4f16 :ref:`adst`::ref:`f32x4`, :ref:`vasrc0`::ref:`f16x4`, :ref:`vasrc1`::ref:`f16x4`, :ref:`asrc2`::ref:`f32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp`
v_mfma_i32_16x16x16i8 :ref:`adst`::ref:`i32x4`, :ref:`vasrc0`::ref:`i8x4`, :ref:`vasrc1`::ref:`i8x4`, :ref:`asrc2`::ref:`i32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp`
v_mfma_i32_16x16x4i8 :ref:`adst`::ref:`i32x16`, :ref:`vasrc0`::ref:`i8x4`, :ref:`vasrc1`::ref:`i8x4`, :ref:`asrc2`::ref:`i32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp`
v_mfma_i32_32x32x4i8 :ref:`adst`::ref:`i32x32`, :ref:`vasrc0`::ref:`i8x4`, :ref:`vasrc1`::ref:`i8x4`, :ref:`asrc2`::ref:`i32x32` :ref:`cbsz` :ref:`abid` :ref:`blgp`
v_mfma_i32_32x32x8i8 :ref:`adst`::ref:`i32x16`, :ref:`vasrc0`::ref:`i8x4`, :ref:`vasrc1`::ref:`i8x4`, :ref:`asrc2`::ref:`i32x16` :ref:`cbsz` :ref:`abid` :ref:`blgp`
v_mfma_i32_4x4x4i8 :ref:`adst`::ref:`i32x4`, :ref:`vasrc0`::ref:`i8x4`, :ref:`vasrc1`::ref:`i8x4`, :ref:`asrc2`::ref:`i32x4` :ref:`cbsz` :ref:`abid` :ref:`blgp`
.. |---| unicode:: U+02014 .. em dash
.. toctree::
:hidden:
AMDGPUAsmGFX9
gfx908_addr_buf
gfx908_adst1024_0
gfx908_adst128_0
gfx908_adst32_0
gfx908_adst512_0
gfx908_asrc1024_0
gfx908_asrc128_0
gfx908_asrc32_0
gfx908_asrc512_0
gfx908_data_buf_atomic32
gfx908_dst_flat_atomic32
gfx908_offset_buf
gfx908_rsrc_buf
gfx908_saddr_flat_global
gfx908_src32_0
gfx908_src32_1
gfx908_src32_2
gfx908_src32_3
gfx908_src32_4
gfx908_src32_5
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