CDNA4 sparse MFMA builtins

CDNA4 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 CDNA4 GPUs (gfx950, MI350 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.

CDNA4 adds doubled-K SMFMAC variants in all data types (FP16, BF16, INT8, FP8, and BF8), doubling the K depth processed per instruction compared to earlier generations.

Architecture availability#

The builtins on this page target CDNA4 (gfx950 architecture also known as AMD Instinct MI350 series) exclusively. Equivalent builtins for other CDNA generations are documented on their own reference pages:

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 (f16, bf16, i8, fp8, or bf8).

in_type_b (FP8/BF8 variants only)

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). Omitted for FP16, BF16, and INT8 variants where \(\pmb{A}\) and \(\pmb{B}\) always share the same type.

For example:

  • __builtin_amdgcn_smfmac_f32_16x16x128_fp8_fp8 – a \(16 \times 16\) sparse MMA with \(K=128\) that multiplies FP8 \(\pmb{A}\) by FP8 \(\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).

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).

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 v4half   = _Float16 [[clang::ext_vector_type(4)]];
using v8half   = _Float16 [[clang::ext_vector_type(8)]];
using v16half  = _Float16 [[clang::ext_vector_type(16)]];
using v8bf16   = __bf16   [[clang::ext_vector_type(8)]];
using v16bf16  = __bf16   [[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 doubled-K 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:

FP8 16×16 sparse kernel (K=128)#

This kernel uses __builtin_amdgcn_smfmac_f32_16x16x128_fp8_fp8 to compute a \(16 \times 16\) sparse matrix multiply-accumulate with \(K=128\) per instruction, doubling the K depth of the gfx942 variant. The compressed \(\pmb{A}\) operand grows to a v4int vector and the dense \(\pmb{B}\) to a v8int vector.

 1struct SmfmacCdna4FP8Policy
 2{
 3    static constexpr int thread_tile_m = 16;
 4    static constexpr int thread_tile_n = 16;
 5
 6    using v4float = float [[clang::ext_vector_type(4)]];
 7    using Accumulator = v4float;
 8
 9    using v4int = int [[clang::ext_vector_type(4)]];
10    using v8int = int [[clang::ext_vector_type(8)]];
11    using AFrag = v4int;
12    using BFrag = v8int;
13
14    __device__ static void zero(Accumulator &d) { d = {0.0f, 0.0f, 0.0f, 0.0f}; }
15
16    __device__ static void load_a(const uint8_t *sA, int warp_m, int lane,
17                                  int tk_comp, AFrag &a)
18    {
19        const int g = lane / 16;
20        const int n = lane % 16;
21        int a_off = (warp_m * 16 + n) * tk_comp + g * 16;
22        for (int j = 0; j < 4; ++j)
23        {
24            a[j] = static_cast<int>(
25                uint32_t(sA[a_off + j*4 + 0])       | (uint32_t(sA[a_off + j*4 + 1]) << 8) |
26                (uint32_t(sA[a_off + j*4 + 2]) << 16) | (uint32_t(sA[a_off + j*4 + 3]) << 24));
27        }
28    }
29
30    __device__ static void load_b(const uint8_t *sB, int warp_n, int lane,
31                                  int cta_n, BFrag &b)
32    {
33        const int g = lane / 16;
34        const int n = lane % 16;
35        int b_col = warp_n * 16 + n;
36        for (int v = 0; v < 8; ++v)
37        {
38            int kk = (v / 4) * 64 + g * 16 + (v % 4) * 4;
39            b[v] = static_cast<int>(
40                uint32_t(sB[kk * cta_n + b_col])       | (uint32_t(sB[(kk+1) * cta_n + b_col]) << 8) |
41                (uint32_t(sB[(kk+2) * cta_n + b_col]) << 16) | (uint32_t(sB[(kk+3) * cta_n + b_col]) << 24));
42        }
43    }
44
45    __device__ static void mma(Accumulator &d, const AFrag &a, const BFrag &b)
46    {
47#if defined(__gfx950__)
48        d = __builtin_amdgcn_smfmac_f32_16x16x128_fp8_fp8(
49                a, b, d, static_cast<int>(0x88888888u), 0, 0);
50#endif
51    }
52
53    __device__ static void store_c(const Accumulator &d, float *D,
54                                   int sub_m, int sub_n, int N, int lane)
55    {
56        const int g = lane / 16;
57        const int n = lane % 16;
58        for (int r = 0; r < 4; ++r)
59        {
60            D[(sub_m + g * 4 + r) * N + sub_n + n] = d[r];
61        }
62    }
63};

The byte-packing pattern is the same as the K=64 variant, but each lane now loads twice as many bytes. The sparsity index expands to a full 32-bit value (0x88888888) to cover the 128-element K dimension.

Compile and run:

amdclang++ -O3 -std=c++17 --offload-arch=gfx950 \
    matrix_multiply_cdna4_sparse_mfma.hip -o mm_cdna4_sparse_mfma
./mm_cdna4_sparse_mfma

Note

This example requires a CDNA4 GPU (gfx950). Compile with --offload-arch=gfx950 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_f16

16384

16

8

__builtin_amdgcn_smfmac_f32_32x32x32_f16

32768

32

24

__builtin_amdgcn_smfmac_f32_16x16x64_bf16

16384

16

8

__builtin_amdgcn_smfmac_f32_32x32x32_bf16

32768

32

24

__builtin_amdgcn_smfmac_i32_16x16x128_i8

32768

16

8

__builtin_amdgcn_smfmac_i32_32x32x64_i8

65536

32

24

__builtin_amdgcn_smfmac_f32_16x16x128_fp8_fp8

32768

16

8

__builtin_amdgcn_smfmac_f32_16x16x128_fp8_bf8

32768

16

8

__builtin_amdgcn_smfmac_f32_16x16x128_bf8_fp8

32768

16

8

__builtin_amdgcn_smfmac_f32_16x16x128_bf8_bf8

32768

16

8

__builtin_amdgcn_smfmac_f32_32x32x64_fp8_fp8

65536

32

24

__builtin_amdgcn_smfmac_f32_32x32x64_fp8_bf8

65536

32

24

__builtin_amdgcn_smfmac_f32_32x32x64_bf8_fp8

65536

32

24

__builtin_amdgcn_smfmac_f32_32x32x64_bf8_bf8

65536

32

24

Builtin reference#

FP32-accumulate builtins#

These builtins accumulate into single-precision (FP32) output fragments. They differ in the data type of the \(\pmb{A}\) and \(\pmb{B}\) matrix inputs.

FP16 matrix inputs#

The following doubled-K FP16 variants are available on gfx950 only.

__builtin_amdgcn_smfmac_f32_16x16x64_f16#

Signature and parameters for this builtin.

v4float __builtin_amdgcn_smfmac_f32_16x16x64_f16(
    v8half  srcA,
    v16half srcB,
    v4float srcC,
    int     idx,
    int     cbsz,
    int     abid);

Computes one step of a \(16 \times 16\) sparse FP32 outer-product accumulation. Each instruction processes \(K=64\) compressed FP16 elements of \(\pmb{A}\) and \(K=64\) dense FP16 elements of \(\pmb{B}\), accumulating into a \(16 \times 16\) FP32 tile held across 4 accVGPRs per lane. This variant is available on gfx950 only.

Parameter

Type

Description

srcA

v8half

Compressed \(\pmb{A}\) elements per lane (4:2 sparse).

srcB

v16half

Dense \(\pmb{B}\) elements per lane.

srcC

v4float

Accumulator input: 4 FP32 elements per lane.

idx

int

Sparsity index encoding the positions of non-zero elements within each group of four along the K dimension, see Sparse MFMA parameters.

cbsz

int

Control Broadcast Size modifier, see Sparse MFMA parameters.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Sparse MFMA parameters.

Returns v4float – updated accumulator (\(\text{srcA} \times \text{srcB} + \text{srcC}\)).

__builtin_amdgcn_smfmac_f32_32x32x32_f16#

Signature and parameters for this builtin.

v16float __builtin_amdgcn_smfmac_f32_32x32x32_f16(
    v8half   srcA,
    v16half  srcB,
    v16float srcC,
    int      idx,
    int      cbsz,
    int      abid);

Computes one step of a \(32 \times 32\) sparse FP32 outer-product accumulation. Each instruction processes \(K=32\) compressed FP16 elements of \(\pmb{A}\) and \(K=32\) dense FP16 elements of \(\pmb{B}\), accumulating into a \(32 \times 32\) FP32 tile held across 16 accVGPRs per lane. This variant is available on gfx950 only.

Parameter

Type

Description

srcA

v8half

Compressed \(\pmb{A}\) elements per lane (4:2 sparse).

srcB

v16half

Dense \(\pmb{B}\) elements per lane.

srcC

v16float

Accumulator input: 16 FP32 elements per lane.

idx

int

Sparsity index encoding the positions of non-zero elements within each group of four along the K dimension, see Sparse MFMA parameters.

cbsz

int

Control Broadcast Size modifier, see Sparse MFMA parameters.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Sparse MFMA parameters.

Returns v16float – updated accumulator (\(\text{srcA} \times \text{srcB} + \text{srcC}\)).

BF16 matrix inputs#

The following doubled-K BF16 variants are available on gfx950 only. These use native __bf16 register types rather than short storage.

__builtin_amdgcn_smfmac_f32_16x16x64_bf16#

Signature and parameters for this builtin.

v4float __builtin_amdgcn_smfmac_f32_16x16x64_bf16(
    v8bf16  srcA,
    v16bf16 srcB,
    v4float srcC,
    int     idx,
    int     cbsz,
    int     abid);

Computes one step of a \(16 \times 16\) sparse FP32 outer-product accumulation. Each instruction processes \(K=64\) compressed BF16 elements of \(\pmb{A}\) and \(K=64\) dense BF16 elements of \(\pmb{B}\), accumulating into a \(16 \times 16\) FP32 tile held across 4 accVGPRs per lane. This variant is available on gfx950 only.

Parameter

Type

Description

srcA

v8bf16

Compressed \(\pmb{A}\) elements per lane (4:2 sparse, BF16 stored as __bf16).

srcB

v16bf16

Dense \(\pmb{B}\) elements per lane (BF16 stored as __bf16).

srcC

v4float

Accumulator input: 4 FP32 elements per lane.

idx

int

Sparsity index encoding the positions of non-zero elements within each group of four along the K dimension, see Sparse MFMA parameters.

cbsz

int

Control Broadcast Size modifier, see Sparse MFMA parameters.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Sparse MFMA parameters.

Returns v4float – updated accumulator (\(\text{srcA} \times \text{srcB} + \text{srcC}\)).

__builtin_amdgcn_smfmac_f32_32x32x32_bf16#

Signature and parameters for this builtin.

v16float __builtin_amdgcn_smfmac_f32_32x32x32_bf16(
    v8bf16   srcA,
    v16bf16  srcB,
    v16float srcC,
    int      idx,
    int      cbsz,
    int      abid);

Computes one step of a \(32 \times 32\) sparse FP32 outer-product accumulation. Each instruction processes \(K=32\) compressed BF16 elements of \(\pmb{A}\) and \(K=32\) dense BF16 elements of \(\pmb{B}\), accumulating into a \(32 \times 32\) FP32 tile held across 16 accVGPRs per lane. This variant is available on gfx950 only.

Parameter

Type

Description

srcA

v8bf16

Compressed \(\pmb{A}\) elements per lane (4:2 sparse, BF16 stored as __bf16).

srcB

v16bf16

Dense \(\pmb{B}\) elements per lane (BF16 stored as __bf16).

srcC

v16float

Accumulator input: 16 FP32 elements per lane.

idx

int

Sparsity index encoding the positions of non-zero elements within each group of four along the K dimension, see Sparse MFMA parameters.

cbsz

int

Control Broadcast Size modifier, see Sparse MFMA parameters.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Sparse MFMA parameters.

Returns v16float – updated accumulator (\(\text{srcA} \times \text{srcB} + \text{srcC}\)).

FP8 and BF8 matrix inputs#

The following doubled-K FP8 and BF8 variants are available on gfx950 only.

__builtin_amdgcn_smfmac_f32_16x16x128_fp8_fp8#

Signature and parameters for this builtin.

v4float __builtin_amdgcn_smfmac_f32_16x16x128_fp8_fp8(
    v4int     srcA,
    v8int     srcB,
    v4float   srcC,
    int      idx,
    int      cbsz,
    int      abid);

Computes one step of a \(16 \times 16\) sparse FP32 outer-product accumulation. Each instruction processes \(K=128\) compressed FP8 (E4M3) elements of \(\pmb{A}\) and \(K=128\) dense FP8 (E4M3) elements of \(\pmb{B}\), accumulating into a \(16 \times 16\) FP32 tile held across 4 accVGPRs per lane. Each 32-bit register lane packs four 8-bit elements. This variant is available on gfx950 only.

Parameter

Type

Description

srcA

v4int

Compressed \(\pmb{A}\) elements per lane (4:2 sparse, four FP8 (E4M3) values packed per 32-bit lane).

srcB

v8int

Dense \(\pmb{B}\) elements per lane (four FP8 (E4M3) values packed per 32-bit lane).

srcC

v4float

Accumulator input: 4 FP32 elements per lane.

idx

int

Sparsity index encoding the positions of non-zero elements within each group of four along the K dimension, see Sparse MFMA parameters.

cbsz

int

Control Broadcast Size modifier, see Sparse MFMA parameters.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Sparse MFMA parameters.

Returns v4float – updated accumulator (\(\text{srcA} \times \text{srcB} + \text{srcC}\)).

__builtin_amdgcn_smfmac_f32_16x16x128_fp8_bf8#

Signature and parameters for this builtin.

v4float __builtin_amdgcn_smfmac_f32_16x16x128_fp8_bf8(
    v4int     srcA,
    v8int     srcB,
    v4float   srcC,
    int      idx,
    int      cbsz,
    int      abid);

Computes one step of a \(16 \times 16\) sparse FP32 outer-product accumulation. Each instruction processes \(K=128\) compressed FP8 (E4M3) elements of \(\pmb{A}\) and \(K=128\) dense BF8 (E5M2) elements of \(\pmb{B}\), accumulating into a \(16 \times 16\) FP32 tile held across 4 accVGPRs per lane. Each 32-bit register lane packs four 8-bit elements. This variant is available on gfx950 only.

Parameter

Type

Description

srcA

v4int

Compressed \(\pmb{A}\) elements per lane (4:2 sparse, four FP8 (E4M3) values packed per 32-bit lane).

srcB

v8int

Dense \(\pmb{B}\) elements per lane (four BF8 (E5M2) values packed per 32-bit lane).

srcC

v4float

Accumulator input: 4 FP32 elements per lane.

idx

int

Sparsity index encoding the positions of non-zero elements within each group of four along the K dimension, see Sparse MFMA parameters.

cbsz

int

Control Broadcast Size modifier, see Sparse MFMA parameters.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Sparse MFMA parameters.

Returns v4float – updated accumulator (\(\text{srcA} \times \text{srcB} + \text{srcC}\)).

__builtin_amdgcn_smfmac_f32_16x16x128_bf8_fp8#

Signature and parameters for this builtin.

v4float __builtin_amdgcn_smfmac_f32_16x16x128_bf8_fp8(
    v4int     srcA,
    v8int     srcB,
    v4float   srcC,
    int      idx,
    int      cbsz,
    int      abid);

Computes one step of a \(16 \times 16\) sparse FP32 outer-product accumulation. Each instruction processes \(K=128\) compressed BF8 (E5M2) elements of \(\pmb{A}\) and \(K=128\) dense FP8 (E4M3) elements of \(\pmb{B}\), accumulating into a \(16 \times 16\) FP32 tile held across 4 accVGPRs per lane. Each 32-bit register lane packs four 8-bit elements. This variant is available on gfx950 only.

Parameter

Type

Description

srcA

v4int

Compressed \(\pmb{A}\) elements per lane (4:2 sparse, four BF8 (E5M2) values packed per 32-bit lane).

srcB

v8int

Dense \(\pmb{B}\) elements per lane (four FP8 (E4M3) values packed per 32-bit lane).

srcC

v4float

Accumulator input: 4 FP32 elements per lane.

idx

int

Sparsity index encoding the positions of non-zero elements within each group of four along the K dimension, see Sparse MFMA parameters.

cbsz

int

Control Broadcast Size modifier, see Sparse MFMA parameters.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Sparse MFMA parameters.

Returns v4float – updated accumulator (\(\text{srcA} \times \text{srcB} + \text{srcC}\)).

__builtin_amdgcn_smfmac_f32_16x16x128_bf8_bf8#

Signature and parameters for this builtin.

v4float __builtin_amdgcn_smfmac_f32_16x16x128_bf8_bf8(
    v4int     srcA,
    v8int     srcB,
    v4float   srcC,
    int      idx,
    int      cbsz,
    int      abid);

Computes one step of a \(16 \times 16\) sparse FP32 outer-product accumulation. Each instruction processes \(K=128\) compressed BF8 (E5M2) elements of \(\pmb{A}\) and \(K=128\) dense BF8 (E5M2) elements of \(\pmb{B}\), accumulating into a \(16 \times 16\) FP32 tile held across 4 accVGPRs per lane. Each 32-bit register lane packs four 8-bit elements. This variant is available on gfx950 only.

Parameter

Type

Description

srcA

v4int

Compressed \(\pmb{A}\) elements per lane (4:2 sparse, four BF8 (E5M2) values packed per 32-bit lane).

srcB

v8int

Dense \(\pmb{B}\) elements per lane (four BF8 (E5M2) values packed per 32-bit lane).

srcC

v4float

Accumulator input: 4 FP32 elements per lane.

idx

int

Sparsity index encoding the positions of non-zero elements within each group of four along the K dimension, see Sparse MFMA parameters.

cbsz

int

Control Broadcast Size modifier, see Sparse MFMA parameters.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Sparse MFMA parameters.

Returns v4float – updated accumulator (\(\text{srcA} \times \text{srcB} + \text{srcC}\)).

__builtin_amdgcn_smfmac_f32_32x32x64_fp8_fp8#

Signature and parameters for this builtin.

v16float __builtin_amdgcn_smfmac_f32_32x32x64_fp8_fp8(
    v4int     srcA,
    v8int     srcB,
    v16float  srcC,
    int      idx,
    int      cbsz,
    int      abid);

Computes one step of a \(32 \times 32\) sparse FP32 outer-product accumulation. Each instruction processes \(K=64\) compressed FP8 (E4M3) elements of \(\pmb{A}\) and \(K=64\) dense FP8 (E4M3) elements of \(\pmb{B}\), accumulating into a \(32 \times 32\) FP32 tile held across 16 accVGPRs per lane. Each 32-bit register lane packs four 8-bit elements. This variant is available on gfx950 only.

Parameter

Type

Description

srcA

v4int

Compressed \(\pmb{A}\) elements per lane (4:2 sparse, four FP8 (E4M3) values packed per 32-bit lane).

srcB

v8int

Dense \(\pmb{B}\) elements per lane (four FP8 (E4M3) values packed per 32-bit lane).

srcC

v16float

Accumulator input: 16 FP32 elements per lane.

idx

int

Sparsity index encoding the positions of non-zero elements within each group of four along the K dimension, see Sparse MFMA parameters.

cbsz

int

Control Broadcast Size modifier, see Sparse MFMA parameters.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Sparse MFMA parameters.

Returns v16float – updated accumulator (\(\text{srcA} \times \text{srcB} + \text{srcC}\)).

__builtin_amdgcn_smfmac_f32_32x32x64_fp8_bf8#

Signature and parameters for this builtin.

v16float __builtin_amdgcn_smfmac_f32_32x32x64_fp8_bf8(
    v4int     srcA,
    v8int     srcB,
    v16float  srcC,
    int      idx,
    int      cbsz,
    int      abid);

Computes one step of a \(32 \times 32\) sparse FP32 outer-product accumulation. Each instruction processes \(K=64\) compressed FP8 (E4M3) elements of \(\pmb{A}\) and \(K=64\) dense BF8 (E5M2) elements of \(\pmb{B}\), accumulating into a \(32 \times 32\) FP32 tile held across 16 accVGPRs per lane. Each 32-bit register lane packs four 8-bit elements. This variant is available on gfx950 only.

Parameter

Type

Description

srcA

v4int

Compressed \(\pmb{A}\) elements per lane (4:2 sparse, four FP8 (E4M3) values packed per 32-bit lane).

srcB

v8int

Dense \(\pmb{B}\) elements per lane (four BF8 (E5M2) values packed per 32-bit lane).

srcC

v16float

Accumulator input: 16 FP32 elements per lane.

idx

int

Sparsity index encoding the positions of non-zero elements within each group of four along the K dimension, see Sparse MFMA parameters.

cbsz

int

Control Broadcast Size modifier, see Sparse MFMA parameters.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Sparse MFMA parameters.

Returns v16float – updated accumulator (\(\text{srcA} \times \text{srcB} + \text{srcC}\)).

__builtin_amdgcn_smfmac_f32_32x32x64_bf8_fp8#

Signature and parameters for this builtin.

v16float __builtin_amdgcn_smfmac_f32_32x32x64_bf8_fp8(
    v4int     srcA,
    v8int     srcB,
    v16float  srcC,
    int      idx,
    int      cbsz,
    int      abid);

Computes one step of a \(32 \times 32\) sparse FP32 outer-product accumulation. Each instruction processes \(K=64\) compressed BF8 (E5M2) elements of \(\pmb{A}\) and \(K=64\) dense FP8 (E4M3) elements of \(\pmb{B}\), accumulating into a \(32 \times 32\) FP32 tile held across 16 accVGPRs per lane. Each 32-bit register lane packs four 8-bit elements. This variant is available on gfx950 only.

Parameter

Type

Description

srcA

v4int

Compressed \(\pmb{A}\) elements per lane (4:2 sparse, four BF8 (E5M2) values packed per 32-bit lane).

srcB

v8int

Dense \(\pmb{B}\) elements per lane (four FP8 (E4M3) values packed per 32-bit lane).

srcC

v16float

Accumulator input: 16 FP32 elements per lane.

idx

int

Sparsity index encoding the positions of non-zero elements within each group of four along the K dimension, see Sparse MFMA parameters.

cbsz

int

Control Broadcast Size modifier, see Sparse MFMA parameters.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Sparse MFMA parameters.

Returns v16float – updated accumulator (\(\text{srcA} \times \text{srcB} + \text{srcC}\)).

__builtin_amdgcn_smfmac_f32_32x32x64_bf8_bf8#

Signature and parameters for this builtin.

v16float __builtin_amdgcn_smfmac_f32_32x32x64_bf8_bf8(
    v4int     srcA,
    v8int     srcB,
    v16float  srcC,
    int      idx,
    int      cbsz,
    int      abid);

Computes one step of a \(32 \times 32\) sparse FP32 outer-product accumulation. Each instruction processes \(K=64\) compressed BF8 (E5M2) elements of \(\pmb{A}\) and \(K=64\) dense BF8 (E5M2) elements of \(\pmb{B}\), accumulating into a \(32 \times 32\) FP32 tile held across 16 accVGPRs per lane. Each 32-bit register lane packs four 8-bit elements. This variant is available on gfx950 only.

Parameter

Type

Description

srcA

v4int

Compressed \(\pmb{A}\) elements per lane (4:2 sparse, four BF8 (E5M2) values packed per 32-bit lane).

srcB

v8int

Dense \(\pmb{B}\) elements per lane (four BF8 (E5M2) values packed per 32-bit lane).

srcC

v16float

Accumulator input: 16 FP32 elements per lane.

idx

int

Sparsity index encoding the positions of non-zero elements within each group of four along the K dimension, see Sparse MFMA parameters.

cbsz

int

Control Broadcast Size modifier, see Sparse MFMA parameters.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Sparse MFMA parameters.

Returns v16float – updated accumulator (\(\text{srcA} \times \text{srcB} + \text{srcC}\)).

INT32-accumulate builtins#

These builtins accept four signed 8-bit integer elements per 32-bit register lane for both \(\pmb{A}\) and \(\pmb{B}\) inputs, and accumulate into INT32 output fragments.

INT8 matrix inputs#

The following doubled-K INT8 variants are available on gfx950 only.

__builtin_amdgcn_smfmac_i32_16x16x128_i8#

Signature and parameters for this builtin.

v4int __builtin_amdgcn_smfmac_i32_16x16x128_i8(
    v4int  srcA,
    v8int  srcB,
    v4int  srcC,
    int    idx,
    int    cbsz,
    int    abid);

Computes one step of a \(16 \times 16\) sparse INT32 outer-product accumulation. Each instruction processes \(K=128\) compressed INT8 elements of \(\pmb{A}\) and \(K=128\) dense INT8 elements of \(\pmb{B}\), accumulating into a \(16 \times 16\) INT32 tile held across 4 accVGPRs per lane. Each 32-bit register lane packs four INT8 elements. This variant is available on gfx950 only.

Parameter

Type

Description

srcA

v4int

Compressed \(\pmb{A}\) elements per lane (4:2 sparse, four INT8 values packed per 32-bit lane).

srcB

v8int

Dense \(\pmb{B}\) elements per lane (four INT8 values packed per 32-bit lane).

srcC

v4int

Accumulator input: 4 INT32 elements per lane.

idx

int

Sparsity index encoding the positions of non-zero elements within each group of four along the K dimension, see Sparse MFMA parameters.

cbsz

int

Control Broadcast Size modifier, see Sparse MFMA parameters.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Sparse MFMA parameters.

Returns v4int – updated accumulator (\(\text{srcA} \times \text{srcB} + \text{srcC}\)).

__builtin_amdgcn_smfmac_i32_32x32x64_i8#

Signature and parameters for this builtin.

v16int __builtin_amdgcn_smfmac_i32_32x32x64_i8(
    v4int   srcA,
    v8int   srcB,
    v16int  srcC,
    int     idx,
    int     cbsz,
    int     abid);

Computes one step of a \(32 \times 32\) sparse INT32 outer-product accumulation. Each instruction processes \(K=64\) compressed INT8 elements of \(\pmb{A}\) and \(K=64\) dense INT8 elements of \(\pmb{B}\), accumulating into a \(32 \times 32\) INT32 tile held across 16 accVGPRs per lane. Each 32-bit register lane packs four INT8 elements. This variant is available on gfx950 only.

Parameter

Type

Description

srcA

v4int

Compressed \(\pmb{A}\) elements per lane (4:2 sparse, four INT8 values packed per 32-bit lane).

srcB

v8int

Dense \(\pmb{B}\) elements per lane (four INT8 values packed per 32-bit lane).

srcC

v16int

Accumulator input: 16 INT32 elements per lane.

idx

int

Sparsity index encoding the positions of non-zero elements within each group of four along the K dimension, see Sparse MFMA parameters.

cbsz

int

Control Broadcast Size modifier, see Sparse MFMA parameters.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Sparse MFMA parameters.

Returns v16int – updated accumulator (\(\text{srcA} \times \text{srcB} + \text{srcC}\)).

Contents