CDNA and CDNA2 sparse MFMA builtins

CDNA and CDNA2 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 CDNA GPUs (gfx908, MI100) and CDNA2 GPUs (gfx90a, MI200 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.

Architecture availability#

The builtins on this page target CDNA (gfx908, MI100) and CDNA2 (gfx90a, MI200 series). Equivalent builtins for later 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, or i8).

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_16x16x32_f16 – a \(16 \times 16\) sparse MMA with \(K=32\) that multiplies FP16 \(\pmb{A}\) by FP16 \(\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=32 SMFMAC A (sparse) operand layout.

\(16 \times 16\), K=32 A (sparse, 2:4) operand layout.#

:math:`16 \times 16`, K=32 SMFMAC compression index layout.

\(16 \times 16\), K=32 compression index layout.#

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

\(16 \times 16\), K=32 B (dense) operand layout.#

:math:`16 \times 16`, K=32 SMFMAC accumulator layout -- VGPR index per output element, with lane groups colour-coded.

\(16 \times 16\), K=32 accumulator layout. Each cell shows the VGPR index that holds output element \((i, j)\). Rows 0—3 (teal, lanes 0—15), rows 4—7 (grey, lanes 16—31), rows 8—11 (teal, lanes 32—47), rows 12—15 (grey, lanes 48—63). Column \(j\) gives the lane offset within the group.#

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=16 SMFMAC A (sparse) operand layout.

\(32 \times 32\), K=16 A (sparse, 2:4) operand layout.#

:math:`32 \times 32`, K=16 SMFMAC compression index layout.

\(32 \times 32\), K=16 compression index layout.#

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

\(32 \times 32\), K=16 B (dense) operand layout.#

:math:`32 \times 32`, K=16 SMFMAC accumulator layout -- VGPR index per output element, with lane groups colour-coded.

\(32 \times 32\), K=16 accumulator layout. Each cell shows the VGPR index that holds output element \((i, j)\). Teal cells (rows where \(\lfloor i/4 \rfloor\) is even) belong to lanes 0—31; grey cells to lanes 32—63. Column \(j\) gives the lane offset within the group.#

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 v4short  = short    [[clang::ext_vector_type(4)]];   // BF16 storage
using v8short  = short    [[clang::ext_vector_type(8)]];   // BF16 storage
using v2int    = int      [[clang::ext_vector_type(2)]];
using v4int    = int      [[clang::ext_vector_type(4)]];
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 kernels demonstrate SMFMAC builtins in the context of a tiled matrix multiplication. Each 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:

FP16 16×16 sparse kernel#

This kernel uses __builtin_amdgcn_smfmac_f32_16x16x32_f16 to compute a \(16 \times 16\) sparse matrix multiply-accumulate with \(K=32\) per instruction. The compressed \(\pmb{A}\) operand is loaded as a v4half vector, the dense \(\pmb{B}\) as a v8half vector, and the accumulator is a v4float vector.

 1struct SmfmacCdna16x16F16Policy
 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 v4half = _Float16 [[clang::ext_vector_type(4)]];
10    using v8half = _Float16 [[clang::ext_vector_type(8)]];
11    using AFrag = v4half;
12    using BFrag = v8half;
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 _Float16 *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 * 4;
22        for (int i = 0; i < 4; ++i)
23        {
24            a[i] = sA[a_off + i];
25        }
26    }
27
28    __device__ static void load_b(const _Float16 *sB, int warp_n, int lane,
29                                  int cta_n, BFrag &b)
30    {
31        const int g = lane / 16;
32        const int n = lane % 16;
33        int b_col = warp_n * 16 + n;
34        for (int v = 0; v < 4; ++v)
35        {
36            int kk = g * 8 + v * 2;
37            b[v * 2 + 0] = sB[kk * cta_n + b_col];
38            b[v * 2 + 1] = sB[(kk + 1) * cta_n + b_col];
39        }
40    }
41
42    __device__ static void mma(Accumulator &d, const AFrag &a, const BFrag &b)
43    {
44#if defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx940__) || \
45    defined(__gfx942__) || defined(__gfx950__)
46        d = __builtin_amdgcn_smfmac_f32_16x16x32_f16(a, b, d, 0x88, 0, 0);
47#endif
48    }
49
50    __device__ static void store_c(const Accumulator &d, float *D,
51                                   int sub_m, int sub_n, int N, int lane)
52    {
53        const int g = lane / 16;
54        const int n = lane % 16;
55        for (int r = 0; r < 4; ++r)
56        {
57            D[(sub_m + g * 4 + r) * N + sub_n + n] = d[r];
58        }
59    }
60};

Each lane loads its \(\pmb{A}\) fragment from the compressed LDS tile, where the K dimension is halved. The \(\pmb{B}\) fragment is loaded from the full dense \(\pmb{B}\) tile in LDS. The builtin call replaces the entire inner-product loop of a baseline kernel with a single instruction. The 4-element FP32 result vector maps to four rows of the \(16 \times 16\) output sub-tile, with the lane index selecting the column.

FP16 32×32 sparse kernel#

This kernel uses __builtin_amdgcn_smfmac_f32_32x32x16_f16 to compute a \(32 \times 32\) sparse matrix multiply-accumulate with \(K=16\) per instruction. The Cooperative Thread Array (CTA) tile grows to 64×64 to accommodate the larger 32×32 wavefront tiles. The accumulator is a 16-element FP32 vector.

 1struct SmfmacCdna32x32F16Policy
 2{
 3    static constexpr int thread_tile_m = 32;
 4    static constexpr int thread_tile_n = 32;
 5
 6    using v16float = float [[clang::ext_vector_type(16)]];
 7    using Accumulator = v16float;
 8
 9    using v4half = _Float16 [[clang::ext_vector_type(4)]];
10    using v8half = _Float16 [[clang::ext_vector_type(8)]];
11    using AFrag = v4half;
12    using BFrag = v8half;
13
14    __device__ static void zero(Accumulator &d) { d = {}; }
15
16    __device__ static void load_a(const _Float16 *sA, int warp_m, int lane,
17                                  int tk_comp, AFrag &a)
18    {
19        const int g = lane / 16;
20        int a_row = warp_m * 32 + (g % 2) * 16 + (lane % 16);
21        int a_col = (g / 2) * 4;
22        for (int i = 0; i < 4; ++i)
23        {
24            a[i] = sA[a_row * tk_comp + a_col + i];
25        }
26    }
27
28    __device__ static void load_b(const _Float16 *sB, int warp_n, int lane,
29                                  int cta_n, BFrag &b)
30    {
31        const int g = lane / 16;
32        int b_col = warp_n * 32 + (g % 2) * 16 + (lane % 16);
33        int k_off = (g / 2) * 8;
34        for (int v = 0; v < 4; ++v)
35        {
36            int kk = k_off + v * 2;
37            b[v * 2 + 0] = sB[kk * cta_n + b_col];
38            b[v * 2 + 1] = sB[(kk + 1) * cta_n + b_col];
39        }
40    }
41
42    __device__ static void mma(Accumulator &d, const AFrag &a, const BFrag &b)
43    {
44#if defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx940__) || \
45    defined(__gfx942__) || defined(__gfx950__)
46        d = __builtin_amdgcn_smfmac_f32_32x32x16_f16(a, b, d, 0x88, 0, 0);
47#endif
48    }
49
50    __device__ static void store_c(const Accumulator &d, float *D,
51                                   int sub_m, int sub_n, int N, int lane)
52    {
53        int j_base = lane % 32;
54        int grp    = lane / 32;
55        for (int r = 0; r < 16; ++r)
56        {
57            int i_out = (r / 4) * 8 + grp * 4 + (r % 4);
58            D[(sub_m + i_out) * N + sub_n + j_base] = d[r];
59        }
60    }
61};

The lane-to-element mapping for the 16-element accumulator differs from the \(16 \times 16\) variant: each lane’s 16 result elements are distributed across four groups of four consecutive rows, with the group stride determined by the lane’s position within the wavefront.

Compile and run:

amdclang++ -O3 -std=c++17 --offload-arch=gfx908 \
    matrix_multiply_cdna_sparse_mfma.hip -o mm_cdna_sparse_mfma
./mm_cdna_sparse_mfma

Note

This example requires a CDNA or CDNA2 GPU (gfx908 or gfx90a). Compile with --offload-arch=gfx908 or --offload-arch=gfx90a 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_16x16x32_f16

16384

16

8

__builtin_amdgcn_smfmac_f32_32x32x16_f16

32768

32

24

__builtin_amdgcn_smfmac_f32_16x16x32_bf16

16384

16

8

__builtin_amdgcn_smfmac_f32_32x32x16_bf16

32768

32

24

__builtin_amdgcn_smfmac_i32_16x16x64_i8

32768

16

8

__builtin_amdgcn_smfmac_i32_32x32x32_i8

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 builtins accept compressed FP16 elements for \(\pmb{A}\) and dense FP16 elements for \(\pmb{B}\), accumulating into FP32 output fragments.

__builtin_amdgcn_smfmac_f32_16x16x32_f16#

Signature and parameters for this builtin.

v4float __builtin_amdgcn_smfmac_f32_16x16x32_f16(
    v4half  srcA,
    v8half  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=32\) compressed FP16 elements of \(\pmb{A}\) and \(K=32\) dense FP16 elements of \(\pmb{B}\), accumulating into a \(16 \times 16\) FP32 tile held across 4 accVGPRs per lane.

Parameter

Type

Description

srcA

v4half

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

srcB

v8half

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_32x32x16_f16#

Signature and parameters for this builtin.

v16float __builtin_amdgcn_smfmac_f32_32x32x16_f16(
    v4half   srcA,
    v8half   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=16\) compressed FP16 elements of \(\pmb{A}\) and \(K=16\) dense FP16 elements of \(\pmb{B}\), accumulating into a \(32 \times 32\) FP32 tile held across 16 accVGPRs per lane.

Parameter

Type

Description

srcA

v4half

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

srcB

v8half

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 builtins accept compressed BF16 elements for \(\pmb{A}\) and dense BF16 elements for \(\pmb{B}\), accumulating into FP32 output fragments. BF16 values are stored as short in register.

__builtin_amdgcn_smfmac_f32_16x16x32_bf16#

Signature and parameters for this builtin.

v4float __builtin_amdgcn_smfmac_f32_16x16x32_bf16(
    v4short  srcA,
    v8short  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=32\) compressed BF16 elements of \(\pmb{A}\) and \(K=32\) dense BF16 elements of \(\pmb{B}\), accumulating into a \(16 \times 16\) FP32 tile held across 4 accVGPRs per lane.

Parameter

Type

Description

srcA

v4short

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

srcB

v8short

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

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_32x32x16_bf16#

Signature and parameters for this builtin.

v16float __builtin_amdgcn_smfmac_f32_32x32x16_bf16(
    v4short   srcA,
    v8short   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=16\) compressed BF16 elements of \(\pmb{A}\) and \(K=16\) dense BF16 elements of \(\pmb{B}\), accumulating into a \(32 \times 32\) FP32 tile held across 16 accVGPRs per lane.

Parameter

Type

Description

srcA

v4short

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

srcB

v8short

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

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 builtins use INT8 matrix inputs.

__builtin_amdgcn_smfmac_i32_16x16x64_i8#

Signature and parameters for this builtin.

v4int __builtin_amdgcn_smfmac_i32_16x16x64_i8(
    v2int  srcA,
    v4int  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=64\) compressed INT8 elements of \(\pmb{A}\) and \(K=64\) 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.

Parameter

Type

Description

srcA

v2int

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

srcB

v4int

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_32x32x32_i8#

Signature and parameters for this builtin.

v16int __builtin_amdgcn_smfmac_i32_32x32x32_i8(
    v2int   srcA,
    v4int   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=32\) compressed INT8 elements of \(\pmb{A}\) and \(K=32\) 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.

Parameter

Type

Description

srcA

v2int

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

srcB

v4int

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