CDNA3 sparse MFMA builtins

CDNA3 sparse MFMA builtins#

Sparse Matrix Fused Multiply-Accumulate (SMFMAC) builtins let you issue hardware matrix multiply-accumulate operations that exploit 4:2 structured sparsity directly from HIP device code on CDNA3 GPUs (gfx942, MI300 series). Each SMFMAC instruction multiplies a compressed \(\pmb{A}\) fragment by a dense \(\pmb{B}\) fragment and accumulates the result into a \(\pmb{C}\) fragment, all within a single wavefront of 64 lanes. Because the \(\pmb{A}\) operand is stored in compressed form, these builtins halve the storage and memory bandwidth required for \(\pmb{A}\) compared to their dense MFMA counterparts, while the hardware uses a sparsity index to reconstruct the original element positions during the multiply.

CDNA3 adds SMFMAC builtins with FP8 (E4M3) and BF8 (E5M2) inputs in all four A×B type combinations.

Architecture availability#

The builtins on this page target CDNA3 (gfx942, MI300 series). Equivalent builtins for other CDNA generations are documented on their own reference pages:

CDNA and CDNA2 sparse MFMA builtins – CDNA (gfx908, MI100) and CDNA2 (gfx90a, MI200 series)
CDNA4 sparse MFMA builtins – CDNA4 (gfx950, MI350 series)

Note

On CDNA3, all four operands (\(\pmb{A}\), \(\pmb{B}\), \(\pmb{C}\), and \(\pmb{D}\)) can reside in either accumulation VGPRs (accVGPRs) or standard architecture VGPRs (archVGPRs). In HIP device code the compiler selects the appropriate register class automatically.

Naming convention#

All SMFMAC builtins follow the pattern:

__builtin_amdgcn_smfmac_<out_type>_<M>x<N>x<K>_<in_type_a>[_<in_type_b>]

out_type: Accumulator element type (f32 or i32).
M, N, K: Tile dimensions in elements. The instruction computes the contribution of a K-wide compressed panel of \(\pmb{A}\) and a K-wide dense panel of \(\pmb{B}\) to an \(M \times N\) output tile. Each instruction processes one K step; the caller loops over K to accumulate a full matrix product.
in_type_a: Input element type of the \(\pmb{A}\) matrix (fp8 or bf8).
in_type_b: Input element type of the \(\pmb{B}\) matrix. Always present for FP8 and BF8 variants to disambiguate the four possible A×B type combinations (fp8_fp8, fp8_bf8, bf8_fp8, bf8_bf8).

For example:

__builtin_amdgcn_smfmac_f32_16x16x64_fp8_bf8 – a \(16 \times 16\) sparse MMA with \(K=64\) that multiplies FP8 \(\pmb{A}\) by BF8 \(\pmb{B}\) and accumulates into FP32.

Structured sparsity (4:2 pattern)#

The SMFMAC instructions require the \(\pmb{A}\) matrix to obey 4:2 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 \(\pmb{A}\) operand to be stored in a compressed representation that contains 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 idx argument to every SMFMAC builtin. The hardware uses it at execution time to align the compressed \(\pmb{A}\) elements against the correct rows of the \(\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 kernels in this topic adopt a fixed pattern that keeps positions 0 and 2 of every group, producing a compressed buffer of half the original K length.

Accumulator layout#

Every SMFMAC instruction computes a single independent \(M \times N\) output tile (block count = 1). The accumulator (\(\pmb{C}\) / \(\pmb{D}\)) layout across wavefront lanes and VGPRs is identical to the 1-block layout of the corresponding dense MFMA tile shape.

The formulas in the subsections below use the following notation:

\(i\) – zero-based row index within the tile, \(0 \le i < M\)
\(j\) – zero-based column index within the tile, \(0 \le j < N\)
lane – wavefront lane that holds the element, \(0 \le \text{lane} < 64\)
VGPR – zero-based index into that lane’s accumulator register file

\(16 \times 16\) layout#

The \(16 \times 16\) output tile occupies 4 VGPRs per lane (v4float or v4int).

:math:`16 \times 16`, K=64 SMFMAC A (sparse) operand layout. — \(16 \times 16\), **K=64 A (sparse, 2:4) operand layout.**#

:math:`16 \times 16`, K=64 SMFMAC compression index layout. — \(16 \times 16\), **K=64 compression index layout.**#

:math:`16 \times 16`, K=64 SMFMAC B (dense) operand layout. — \(16 \times 16\), **K=64 B (dense) operand layout.**#

:math:`16 \times 16`, K=64 SMFMAC accumulator layout -- VGPR index per output element, with lane groups colour-coded. — \(16 \times 16\), **K=64 accumulator layout.** The output layout is identical to K=32; only the \(\pmb{A}\) and \(\pmb{B}\) input fragment sizes differ.#

Given output element \((i, j)\):

\[\begin{split}\text{lane} &= 16 \lfloor \frac{i}{4} \rfloor + j \\ \text{VGPR} &= i \bmod 4\end{split}\]

Conversely, given lane \(L\) and VGPR index \(G\):

\[\begin{split}i &= 4 \lfloor \frac{L}{16} \rfloor + (G \bmod 4) \\ j &= L \bmod 16\end{split}\]

The row-to-lane mapping:

Rows	Lanes	VGPRs
0—3	0—15	0—3
4—7	16—31	0—3
8—11	32—47	0—3
12—15	48—63	0—3

\(32 \times 32\) layout#

The \(32 \times 32\) output tile occupies 16 VGPRs per lane (v16float or v16int).

:math:`32 \times 32`, K=32 SMFMAC A (sparse) operand layout. — \(32 \times 32\), **K=32 A (sparse, 2:4) operand layout.**#

:math:`32 \times 32`, K=32 SMFMAC compression index layout. — \(32 \times 32\), **K=32 compression index layout.**#

:math:`32 \times 32`, K=32 SMFMAC B (dense) operand layout. — \(32 \times 32\), **K=32 B (dense) operand layout.**#

:math:`32 \times 32`, K=32 SMFMAC accumulator layout -- VGPR index per output element, with lane groups colour-coded. — \(32 \times 32\), **K=32 accumulator layout.** The output layout is identical to K=16; only the \(\pmb{A}\) and \(\pmb{B}\) input fragment sizes differ.#

Given output element \((i, j)\):

\[\begin{split}\text{lane} &= \bigl(32 \cdot \lfloor \frac{i}{4} \rfloor\bigr) \bmod 64 + j \\ \text{VGPR} &= 4 \lfloor \frac{i}{8} \rfloor + (i \bmod 4)\end{split}\]

Conversely, given lane \(L\) and VGPR index \(G\):

\[\begin{split}i &= \bigl(8 \cdot \lfloor \frac{G}{4} \rfloor\bigr) \bmod 32 + 4 \lfloor \frac{L}{32} \rfloor + (G \bmod 4) \\ j &= L \bmod 32\end{split}\]

The row-to-lane mapping:

Rows	Lanes	VGPRs
0—3	0—31	0—3
4—7	32—63	0—3
8—11	0—31	4—7
12—15	32—63	4—7
16—19	0—31	8—11
20—23	32—63	8—11
24—27	0—31	12—15
28—31	32—63	12—15

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:

using v4float  = float    [[clang::ext_vector_type(4)]];
using v16float = float    [[clang::ext_vector_type(16)]];
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 v16int   = int      [[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; the total VGPR count equals the element count multiplied by the element size in 32-bit words.

Common parameters#

The SMFMAC builtins do not use the cbsz, abid, or blgp modifiers from the dense MFMA family. See Common MFMA parameters for a description of those modifiers in the dense context.

Using sparse MFMA builtins as a compute policy#

The following example kernel demonstrates the FP8 SMFMAC builtin in the context of a tiled matrix multiplication. The kernel loads tiles of the compressed \(\pmb{A}\) matrix and the dense \(\pmb{B}\) matrix into LDS, then calls the SMFMAC builtin to replace the inner-product loop of a conventional scalar kernel.

The complete source file is available for download:

matrix_multiply_cdna3_sparse_mfma.hip

FP8 16×16 sparse kernel#

This kernel uses __builtin_amdgcn_smfmac_f32_16x16x64_fp8_fp8 to compute a \(16 \times 16\) sparse matrix multiply-accumulate with \(K=64\) per instruction using FP8 (E4M3) inputs. Each 32-bit register lane packs four FP8 bytes; the kernel constructs the packed v2int and v4int operands from individual bytes loaded from LDS.

struct SmfmacCdna3FP8Policy
{
    static constexpr int thread_tile_m = 16;
    static constexpr int thread_tile_n = 16;

    using v4float = float [[clang::ext_vector_type(4)]];
    using Accumulator = v4float;

    using v2int = int [[clang::ext_vector_type(2)]];
    using v4int = int [[clang::ext_vector_type(4)]];
    using AFrag = v2int;
    using BFrag = v4int;

    __device__ static void zero(Accumulator &d) { d = {0.0f, 0.0f, 0.0f, 0.0f}; }

    __device__ static void load_a(const uint8_t *sA, int warp_m, int lane,
                                  int tk_comp, AFrag &a)
    {
        const int g = lane / 16;
        const int n = lane % 16;
        int a_off = (warp_m * 16 + n) * tk_comp + g * 8;
        for (int j = 0; j < 2; ++j)
        {
            a[j] = static_cast<int>(
                uint32_t(sA[a_off + j*4 + 0])       | (uint32_t(sA[a_off + j*4 + 1]) << 8) |
                (uint32_t(sA[a_off + j*4 + 2]) << 16) | (uint32_t(sA[a_off + j*4 + 3]) << 24));
        }
    }

    __device__ static void load_b(const uint8_t *sB, int warp_n, int lane,
                                  int cta_n, BFrag &b)
    {
        const int g = lane / 16;
        const int n = lane % 16;
        int b_col = warp_n * 16 + n;
        for (int v = 0; v < 4; ++v)
        {
            int kk = g * 16 + v * 4;
            b[v] = static_cast<int>(
                uint32_t(sB[kk * cta_n + b_col])       | (uint32_t(sB[(kk+1) * cta_n + b_col]) << 8) |
                (uint32_t(sB[(kk+2) * cta_n + b_col]) << 16) | (uint32_t(sB[(kk+3) * cta_n + b_col]) << 24));
        }
    }

    __device__ static void mma(Accumulator &d, const AFrag &a, const BFrag &b)
    {
#if defined(__gfx942__) || defined(__gfx950__)
        d = __builtin_amdgcn_smfmac_f32_16x16x64_fp8_fp8(a, b, d, 0x8888, 0, 0);
#endif
    }

    __device__ static void store_c(const Accumulator &d, float *D,
                                   int sub_m, int sub_n, int N, int lane)
    {
        const int g = lane / 16;
        const int n = lane % 16;
        for (int r = 0; r < 4; ++r)
        {
            D[(sub_m + g * 4 + r) * N + sub_n + n] = d[r];
        }
    }
};

The kernel packs four consecutive FP8 bytes into each 32-bit lane using shifts and bitwise OR before passing the vectors to the builtin. The sparsity index is wider than the FP16 variant (0x8888 vs 0x88) because the larger K dimension requires more index bits to cover all element groups.

Note

The FP8 encoding used by gfx942 is FNUZ (Finite, No Unsigned Zero), while later architectures use the standard OCP (Open Compute Project) E4M3 encoding. The example code selects the correct interpretation at runtime based on the device architecture.

Compile and run:

amdclang++ -O3 -std=c++17 --offload-arch=gfx942 \
    matrix_multiply_cdna3_sparse_mfma.hip -o mm_cdna3_sparse_mfma
./mm_cdna3_sparse_mfma

Note

This example requires a CDNA3 GPU (gfx942). Compile with --offload-arch=gfx942 to select the correct architecture.

Instruction throughput#

The cycle count below is the value used to compute theoretical peak throughput: \(\text{peak throughput} = \frac{\text{ops per instruction}}{\text{cycle count}} \times \text{clock frequency}\). All SMFMAC instructions support vector ALU (VALU) co-execution; the VALU co-execution cycle count gives the number of VALU cycles available during the SMFMAC latency window.

Builtin	Ops	Cycle count	VALU co-execution cycles
`__builtin_amdgcn_smfmac_f32_16x16x64_fp8_fp8`	32768	16	8
`__builtin_amdgcn_smfmac_f32_16x16x64_fp8_bf8`	32768	16	8
`__builtin_amdgcn_smfmac_f32_16x16x64_bf8_fp8`	32768	16	8
`__builtin_amdgcn_smfmac_f32_16x16x64_bf8_bf8`	32768	16	8
`__builtin_amdgcn_smfmac_f32_32x32x32_fp8_fp8`	65536	32	24
`__builtin_amdgcn_smfmac_f32_32x32x32_fp8_bf8`	65536	32	24
`__builtin_amdgcn_smfmac_f32_32x32x32_bf8_fp8`	65536	32	24
`__builtin_amdgcn_smfmac_f32_32x32x32_bf8_bf8`	65536	32	24

Builtin reference#

FP32-accumulate builtins#

These builtins accumulate into single-precision (FP32) output fragments.

FP8 and BF8 matrix inputs#

The following builtins accept compressed FP8 (E4M3) or BF8 (E5M2) elements for \(\pmb{A}\) and dense FP8 or BF8 elements for \(\pmb{B}\), accumulating into FP32 output fragments. All four combinations of FP8 and BF8 input types are supported for each tile size.