.. meta::
   :description: Reference for RDNA4 sparse wave-matrix multiply-accumulate
      intrinsics, covering all supported __builtin_amdgcn_swmmac variants, parameters, and layouts.
   :keywords: RDNA4, gfx1200, gfx1201, SWMMAC, sparse matrix, wave-matrix, HIP intrinsics, FP32, FP16, BF16, FP8, BF8, INT8, INT4, __builtin_amdgcn_swmmac

.. _rdna4-sparse-wmma-builtins:

********************************************************************************
RDNA4 sparse WMMA builtins
********************************************************************************

Sparse Wave-Matrix Multiply-Accumulate (SWMMAC) builtins let you issue
hardware sparse matrix multiply-accumulate operations directly from HIP device
code on RDNA4 GPUs (``gfx1200``, ``gfx1201``). Each SWMMAC instruction
multiplies a compressed sparse :math:`\pmb{A}` fragment by a dense
:math:`\pmb{B}` fragment and accumulates the result into a :math:`\pmb{C}`
fragment, all within a single 32-wide wavefront. Because the :math:`\pmb{A}`
operand is stored in compressed form, SWMMAC halves the storage and bandwidth
required for :math:`\pmb{A}` relative to a dense multiply of the same tile size.

SWMMAC is the sparse variant of the WMMA instruction family, which is used on
RDNA (consumer) GPUs. CDNA (Instinct) GPUs provide a comparable operation
through :doc:`Sparse Matrix Fused Multiply-Accumulate (SMFMAC) <../cdna/sparse-mfma-builtins>`, the sparse variant of
Matrix Fused Multiply-Accumulate (MFMA). The two differ in wavefront size (wave32 for SWMMAC, wave64 for SMFMAC)
and accumulator storage (ordinary Vector General-Purpose Registers (VGPRs) for SWMMAC, dedicated accVGPRs for
SMFMAC). The underlying sparsity model (2:4 structured sparsity on
:math:`\pmb{A}`) is the same on both architectures.

.. note::

   RDNA4 GPUs run all shader programs in ``wave32`` mode by default. The
   ``_w32`` suffix in each builtin name reflects this: all SWMMAC
   builtins on this page require ``wavefrontsize32``. ``wavefrontsize64`` is not supported for HIP code.

Architecture availability
=========================

The builtins on this page target RDNA4 GPUs. To automatically enable them, pass the Low Level Virtual Machine (LLVM) target architecture flag at compile time:

.. code-block:: bash

   amdclang++ --offload-arch=gfx1201 ...
   amdclang++ --offload-arch=gfx1200 ...

Naming convention
=================

All SWMMAC builtins follow the pattern:

.. code-block:: text

   __builtin_amdgcn_swmmac_<out_type>_<M>x<N>x<K>_<in_type>[_<in_type_b>]_w32

``out_type``
    Accumulator element type (``f32``, ``f16``, ``bf16``, or ``i32``).

``M``, ``N``, ``K``
    Tile dimensions. The K dimension refers to the *compressed* K of the sparse
    :math:`\pmb{A}` operand. The corresponding dense K is twice as large due to
    2:4 sparsity.

``in_type``
    Input element type of :math:`\pmb{A}` (and :math:`\pmb{B}` when both share
    the same type): ``f16``, ``bf16``, ``fp8``, ``bf8``, ``iu8``, or ``iu4``.
    The ``iu`` prefix means the builtin accepts either signed or unsigned
    integers, controlled by the ``a_neg`` and ``b_neg`` parameters.

``in_type_b`` (optional)
    Input element type of :math:`\pmb{B}` when it differs from :math:`\pmb{A}`.
    Used only for mixed FP8 and BF8 variants.

``_w32``
    Wavefront size suffix. All RDNA4 SWMMAC builtins use wave32.

Structured sparsity (2:4 pattern)
=================================

SWMMAC instructions require the :math:`\pmb{A}` matrix to obey 2:4 structured
sparsity: in every contiguous group of four elements along the K dimension,
exactly two are non-zero and the other two are zero. This constraint allows the
:math:`\pmb{A}` operand to be stored in a compressed representation containing
only the non-zero values, reducing the storage to half the original K
dimension.

Along with the compressed non-zero values, the hardware requires a sparsity
index that encodes which two of the four positions in each group hold the
non-zeros. This index is passed as the ``index`` argument to every SWMMAC
builtin. The hardware uses it at execution time to align the compressed
:math:`\pmb{A}` elements against the correct rows of the :math:`\pmb{B}`
operand before accumulating the products.

Host-side compression is straightforward: walk the K dimension in groups of
four, extract the two non-zero positions, and pack them consecutively into the
output buffer. The example kernel in this topic uses a fixed pattern that keeps
positions 0 and 2 of every group, producing a compressed buffer of half the
original K length.

.. _rdna4-swmmac-accumulator-layout:

Fragment layouts
================

All SWMMAC builtins on this page use a :math:`16 \times 16` output tile
computed by one wave32 wavefront. The 32 lanes split into two groups of 16;
each group owns half the output rows. The diagrams below show the mapping
between matrix elements and lane or VGPR positions for each operand.

The formulas in the subsections below use the following notation:

* :math:`i` -- zero-based row index within the tile, :math:`0 \le i < 16`
* :math:`j` -- zero-based column index within the tile, :math:`0 \le j < 16`
* **lane** -- wavefront lane, :math:`0 \le \text{lane} < 32`
* **VGPR** -- zero-based index into that lane's register vector

Accumulator layout
------------------

Each lane holds 8 FP32 output elements across VGPRs 0--7.

.. figure:: ../../data/compiler-builtins/rdna/swmmac-builtins/swmmac-layout-d-16x16.svg
   :alt: 16×16 SWMMAC accumulator layout. Rows 0–7 (teal) are held by
         lanes 0–15; rows 8–15 (grey) by lanes 16–31. Each cell shows
         the VGPR index g (0–7) that holds element (i, j). Column j
         equals lane % 16.
   :align: center
   :width: 80%

Given output element :math:`(i, j)`:

.. math::

   \text{lane} &= \lfloor \frac{i}{8} \rfloor \cdot 16 + j \\
   \text{VGPR} &= i \bmod 8

Conversely, given lane :math:`L` and VGPR index :math:`g`:

.. math::

   i &= \lfloor \frac{L}{16} \rfloor \cdot 8 + g \\
   j &= L \bmod 16

The row-to-lane mapping:

.. list-table::
   :header-rows: 1
   :widths: auto

   * - Rows
     - Lanes
     - VGPRs
   * - 0--7
     - 0--15
     - 0--7
   * - 8--15
     - 16--31
     - 0--7

srcA (sparse, FP16 and BF16)
----------------------------

Each lane holds 8 compressed FP16 values (``v8fp16``, 4 VGPRs × 2 FP16)
covering one row of the sparse :math:`\pmb{A}` matrix. The compressed-K
positions are non-contiguous across the two lane groups.

.. figure:: ../../data/compiler-builtins/rdna/swmmac-builtins/swmmac-layout-a-16x16x32.svg
   :alt: 16×16×32 SWMMAC srcA sparse fragment layout. Rose columns
         (compressed K 0–3 and 8–11) are held by lane group 0 (lanes
         0–15); grey columns (compressed K 4–7 and 12–15) by lane group
         1 (lanes 16–31). Each cell shows the VGPR index (0–3).
   :align: center
   :width: 80%

Lane :math:`L` covers matrix row :math:`L \bmod 16`. The 8 compressed
elements are distributed across VGPRs as follows:

.. list-table::
   :header-rows: 1
   :widths: auto

   * - Lane group
     - VGPR 0
     - VGPR 1
     - VGPR 2
     - VGPR 3
   * - 0 (lanes 0--15)
     - compressed K {0, 1}
     - compressed K {2, 3}
     - compressed K {8, 9}
     - compressed K {10, 11}
   * - 1 (lanes 16--31)
     - compressed K {4, 5}
     - compressed K {6, 7}
     - compressed K {12, 13}
     - compressed K {14, 15}

srcB (dense, FP16 and BF16)
---------------------------

Each lane holds 16 dense FP16 values (``v16fp16``, 8 VGPRs × 2 FP16)
covering one column of the dense :math:`\pmb{B}` matrix. The K-row positions
are non-contiguous across the two lane groups.

.. figure:: ../../data/compiler-builtins/rdna/swmmac-builtins/swmmac-layout-b-16x16x32.svg
   :alt: 16×16×32 SWMMAC srcB dense fragment layout. Rose rows (K 0–7
         and K 16–23) are held by lane group 0 (lanes 0–15); grey rows
         (K 8–15 and K 24–31) by lane group 1 (lanes 16–31). Each cell
         shows the VGPR index (0–7).
   :align: center
   :width: 80%

Lane :math:`L` covers matrix column :math:`L \bmod 16`. The 16 dense
K rows are distributed across VGPRs as follows:

.. list-table::
   :header-rows: 1
   :widths: auto

   * - Lane group
     - VGPRs 0--3
     - VGPRs 4--7
   * - 0 (lanes 0--15)
     - K rows 0--7
     - K rows 16--23
   * - 1 (lanes 16--31)
     - K rows 8--15
     - K rows 24--31

Register types used in this reference
=====================================

The signatures below use the following type aliases, which you can declare with
C++ attributes in any HIP translation unit:

.. code-block:: cpp

   using v2int    = int    [[clang::ext_vector_type(2)]];
   using v4int    = int    [[clang::ext_vector_type(4)]];
   using v8int    = int    [[clang::ext_vector_type(8)]];
   using v8float  = float  [[clang::ext_vector_type(8)]];
   using v8short  = short  [[clang::ext_vector_type(8)]];   // BF16 storage
   using v16short = short  [[clang::ext_vector_type(16)]];  // BF16 storage
   using v8fp16   = __fp16 [[clang::ext_vector_type(8)]];
   using v16fp16  = __fp16 [[clang::ext_vector_type(16)]];

Each type alias maps one-to-one to the corresponding LLVM vector type used in
the builtin definition. The number in the name is the element count per lane.

.. note::

   ``__fp16`` and ``_Float16`` are distinct Clang types. ``__fp16`` is a
   storage-only type: arithmetic on ``__fp16`` values promotes to ``float``
   before the operation. ``_Float16`` is the C standard FP16 arithmetic type
   that supports native half-precision operations without promotion. RDNA4
   SWMMAC builtins use ``__fp16``; CDNA MFMA builtins use ``_Float16``.
   The two types are not implicitly convertible, so cast explicitly when
   sharing FP16 data between code paths that target different architectures.

.. note::

   BF16 matrix inputs use ``short`` as the storage type for RDNA4 SWMMAC
   (not ``__bf16``). Reinterpret your BF16 data with ``__builtin_bit_cast``
   or a union before passing it to the builtin.

.. _rdna4-swmmac-common-parameters:

Common parameters
=================

The ``index`` parameter is shared by all SWMMAC builtins. The ``a_neg``,
``b_neg``, and ``clamp`` parameters appear only on integer variants.

.. list-table::
   :header-rows: 1
   :widths: auto

   * - Parameter
     - Type
     - Description
   * - ``index``
     - ``int``
     - Sparsity index register. Each pair of bits encodes the position (0--3)
       of one non-zero element within its block of four consecutive
       :math:`\pmb{A}` elements along K. The 32-bit ``int`` holds indices for
       all elements owned by one lane. Must satisfy the 2:4 constraint: exactly
       two of the four elements in each block must be selected.
   * - ``a_neg``
     - ``bool`` (compile-time constant)
     - Integer variants only. When ``true``, the :math:`\pmb{A}` elements are
       treated as signed integers; when ``false``, as unsigned.
   * - ``b_neg``
     - ``bool`` (compile-time constant)
     - Integer variants only. Same as ``a_neg`` but for :math:`\pmb{B}`.
   * - ``clamp``
     - ``bool`` (compile-time constant)
     - Integer variants only. When ``true``, the INT32 accumulator output is
       clamped to the representable range of the input type on overflow.

Using sparse WMMA builtins as a compute policy
================================================

:ref:`mfma-compute-policy` explains the ``ComputePolicy`` pattern used to
separate the multiply-accumulate logic from the rest of a kernel. The example
below implements ``SwmmacRdna4F16Policy`` using
``__builtin_amdgcn_swmmac_f32_16x16x32_f16_w32``.

Each wavefront computes a single :math:`16 \times 16` output tile. The
:math:`\pmb{A}` operand is pre-sparsified: half the K positions are zero and
omitted from storage, so the compressed K dimension is 16 (representing 32
dense K positions).

The complete source file is available for download:

* :download:`matrix_multiply_rdna4_swmmac.hip <../../tools/example_codes/matrix_multiply_rdna4_swmmac.hip>`

.. rubric:: Policy constants

``swmmac_f32_16x16x32_f16_w32`` consumes 32 dense K positions per call
(compressed to K=16 in :math:`\pmb{A}`), so ``k_step = 32``. The
wavefront holds the entire :math:`16 \times 16` tile: ``thread_tile_m =
thread_tile_n = 16`` and ``effective_lanes = 32``.

.. rubric:: Sparsity index

Each lane's ``index`` register encodes two bits per compressed K position,
identifying which of the four elements in each 2:4 block is non-zero. The
example constructs an even-column sparsity pattern (elements at positions
0 and 2 in each block of four) at the host and passes it to the device.

.. rubric:: Accumulator layout

The builtin returns a ``v8float`` holding 8 FP32 values per lane. The
``store_c()`` pass maps ``(lane, VGPR index)`` back to :math:`(i, j)`
coordinates using the SWMMAC accumulator layout.

.. literalinclude:: ../../tools/example_codes/matrix_multiply_rdna4_swmmac.hip
   :language: cpp
   :start-after: [Sphinx swmmac rdna4 policy start]
   :end-before: [Sphinx swmmac rdna4 policy end]

.. rubric:: Instantiating the kernel

With ``SwmmacRdna4F16Policy`` in place, plug it into the generic kernel
alongside a ``TilePolicy`` (see :ref:`mfma-compute-policy`) whose
``block_tile_m`` and ``block_tile_n`` are multiples of 16 and whose
``k_tile_size`` is a multiple of ``k_step = 32``.

.. literalinclude:: ../../tools/example_codes/matrix_multiply_rdna4_swmmac.hip
   :language: cpp
   :start-after: [Sphinx swmmac policy aliases start]
   :end-before: [Sphinx swmmac policy aliases end]

.. literalinclude:: ../../tools/example_codes/matrix_multiply_rdna4_swmmac.hip
   :language: cpp
   :start-after: [Sphinx swmmac launch config start]
   :end-before: [Sphinx swmmac launch config end]

.. literalinclude:: ../../tools/example_codes/matrix_multiply_rdna4_swmmac.hip
   :language: cpp
   :start-after: [Sphinx swmmac kernel launch start]
   :end-before: [Sphinx swmmac kernel launch end]

**Compile and run:**

.. code-block:: bash

   amdclang++ -O3 -std=c++17 --offload-arch=gfx1201 \
       matrix_multiply_rdna4_swmmac.hip -o mm_rdna4_swmmac
   ./mm_rdna4_swmmac

   amdclang++ -O3 -std=c++17 --offload-arch=gfx1200 \
       matrix_multiply_rdna4_swmmac.hip -o mm_rdna4_swmmac
   ./mm_rdna4_swmmac

.. note::

   ``SwmmacRdna4F16Policy`` requires an RDNA4 GPU (``gfx1200`` or ``gfx1201``).
   The ``#if defined(__gfx1200__) || defined(__gfx1201__)``
   guard in the example file falls back to ``ScalarFMASPolicy`` on other targets,
   so the file compiles without modification.

   The example uses pre-sparsified input data with a fixed 2:4 pattern. In a
   production kernel, apply a sparsity pruning pass to the weight matrix offline
   and store the compressed values and index tensor separately.

.. _rdna4-swmmac-instruction-throughput:

Instruction throughput
======================

The cycle count below is the value used to compute theoretical peak
throughput: :math:`\text{peak throughput} =
\frac{\text{ops per instruction}}{\text{cycle count}} \times
\text{clock frequency}`.

.. list-table::
   :header-rows: 1
   :widths: auto

   * - Builtin
     - Ops
     - Cycle count
   * - ``__builtin_amdgcn_swmmac_f32_16x16x32_f16_w32``
     - 16384
     - 16
   * - ``__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w32``
     - 16384
     - 16
   * - ``__builtin_amdgcn_swmmac_f16_16x16x32_f16_w32``
     - 16384
     - 16
   * - ``__builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w32``
     - 16384
     - 16
   * - ``__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w32``
     - 16384
     - 8
   * - ``__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w32``
     - 16384
     - 8
   * - ``__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w32``
     - 32768
     - 8
   * - ``__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w32``
     - 16384
     - 8
   * - ``__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w32``
     - 16384
     - 8
   * - ``__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w32``
     - 16384
     - 8
   * - ``__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w32``
     - 16384
     - 8

.. _rdna4-swmmac-builtin-reference:

Builtin reference
===================

The following sections list every SWMMAC builtin available on RDNA4,
grouped by accumulator type.

FP32-accumulate builtins
--------------------------

These builtins accumulate into FP32 and accept FP16, BF16, FP8, or BF8
matrix inputs.

FP16 inputs
^^^^^^^^^^^

The following builtins use FP16 matrix inputs.

.. include:: swmmac-ref/f32-16x16x32f16.rst

BF16 inputs
^^^^^^^^^^^

The following builtins use BF16 matrix inputs.

.. include:: swmmac-ref/f32-16x16x32bf16.rst

FP8 and BF8 inputs
^^^^^^^^^^^^^^^^^^

The following builtins use FP8 and BF8 matrix inputs.

.. include:: swmmac-ref/f32-16x16x32fp8-fp8.rst

.. include:: swmmac-ref/f32-16x16x32fp8-bf8.rst

.. include:: swmmac-ref/f32-16x16x32bf8-fp8.rst

.. include:: swmmac-ref/f32-16x16x32bf8-bf8.rst

FP16-accumulate builtins
--------------------------

This builtin accumulates into FP16 with FP16 inputs.

FP16 inputs
^^^^^^^^^^^

The following builtin uses FP16 matrix inputs.

.. include:: swmmac-ref/f16-16x16x32f16.rst

BF16-accumulate builtins
--------------------------

This builtin accumulates into BF16 with BF16 inputs.

BF16 inputs
^^^^^^^^^^^

The following builtin uses BF16 matrix inputs.

.. include:: swmmac-ref/bf16-16x16x32bf16.rst

INT32-accumulate builtins
---------------------------

Integer SWMMAC builtins accept either signed or unsigned 8-bit or 4-bit
integer inputs, controlled by the ``a_neg`` and ``b_neg`` compile-time
constants.

INT8 and UINT8 inputs (16x16x32)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^

The following builtins use INT8 and UINT8 matrix inputs.

.. include:: swmmac-ref/i32-16x16x32iu8.rst

INT4 and UINT4 inputs (16x16x32)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^

The following builtins use INT4 and UINT4 matrix inputs.

.. include:: swmmac-ref/i32-16x16x32iu4.rst

INT4 and UINT4 inputs (16x16x64)
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^

The following builtins use INT4 and UINT4 matrix inputs.

.. include:: swmmac-ref/i32-16x16x64iu4.rst