CDNA3 dense MFMA builtins

Contents

CDNA3 dense MFMA builtins#

Matrix Fused Multiply-Add (MFMA) builtins let you issue hardware matrix multiply-accumulate operations directly from HIP device code on CDNA3 GPUs (gfx942, MI300 series). Each MFMA instruction multiplies a small \(\pmb{A}\) fragment by a small \(\pmb{B}\) fragment and accumulates the result into a \(\pmb{C}\) fragment, all within a single wavefront of 64 lanes. The hardware delivers significantly higher throughput than an equivalent sequence of scalar fused multiply-add (FMA) instructions.

CDNA3 supports the same FP32, FP16, INT8, and FP64 shapes as CDNA2 and adds eXtended Float32 (XF32), FP8, and BF8 input formats. The BF16 and narrower-K INT8 builtins from CDNA1 and CDNA2 are not available on CDNA3; the underlying instruction set architecture (ISA) instructions have been replaced with K-doubled variants that do not yet have corresponding HIP builtins.

Architecture availability#

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

Naming convention#

All MFMA builtins follow the pattern:

__builtin_amdgcn_mfma_<out_type>_<M>x<N>x<K><in_type>

For FP8 and BF8 builtins, where the \(\pmb{A}\) and \(\pmb{B}\) input types can differ, the pattern is:

__builtin_amdgcn_mfma_<out_type>_<M>x<N>x<K>_<typeA>_<typeB>
out_type

Accumulator element type (f32, f64, or i32).

M, N, K

Tile dimensions in elements. The instruction computes the contribution of a K-wide panel of \(\pmb{A}\) (\(M \times K\)) and a K-wide panel of \(\pmb{B}\) (\(K \times N\)) 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

Input element type (f32, f64, f16, bf16, i8, or xf32).

typeA, typeB

For FP8 and BF8 builtins: the element type of the \(\pmb{A}\) and \(\pmb{B}\) matrices respectively. Each is one of fp8 (E4M3 format) or bf8 (E5M2 format).

Note

CDNA3 renames several underlying ISA instructions to include an explicit block count (for example, v_mfma_f32_32x32x1f32 becomes v_mfma_f32_32x32x1_2b_f32 in the CDNA3 ISA). The HIP builtin names (__builtin_amdgcn_mfma_*) are unchanged; the rename is transparent to HIP device code.

Accumulator layout#

Each MFMA instruction computes one or more independent \(M \times N\) output tiles simultaneously across the 64 lanes of a wavefront. The number of independent tiles is called the block count. It depends on how the \(\pmb{A}\) matrix rows are distributed across lane groups:

  • Scalar-input variants (FP32 or FP64 \(\pmb{A}\) and \(\pmb{B}\)): each K position occupies a separate group of \(M\) lanes, so \(\text{blocks} = \frac{\text{wavefront size}}{M \times K}\).

  • Packed-input variants (FP16, BF16, XF32, INT8, FP8, BF8 \(\pmb{A}\) and \(\pmb{B}\)): all K positions are packed into the register bits of the same lane group, so \(\text{blocks} = \frac{\text{wavefront size}}{M}\) regardless of \(K\).

Each block is independent: the \(\pmb{A}\), \(\pmb{B}\), and \(\pmb{C}\)/\(\pmb{D}\) operands of different blocks occupy distinct lane and accVGPR positions and compute separate outer products.

The total number of accVGPRs per lane scales with the block count:

\[\text{accVGPRs per lane} = \text{blocks} \times \frac{M \times N}{\text{wavefront size}}\]

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.

Each of the 64 wavefront lanes maintains its own private accVGPR file. An accVGPR index always refers to a register within one specific lane’s file; the same index in two different lanes denotes two distinct physical registers. The layout tables and formulas in the following sections use accVGPR index as a logical position label; the actual register class (accVGPR or archVGPR) is chosen by the compiler and does not affect the layout.

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

  • \(b\) – block index, \(0 \le b < \text{blocks}\)

  • lane – wavefront lane that holds the element, \(0 \le \text{lane} < 64\)

  • accVGPR – zero-based index into that lane’s private accumulator register file; the same index in two different lanes refers to two distinct physical registers

\(32 \times 32\) layout#

The \(32 \times 32\) tile shape uses two blocks. All 32 accVGPRs per lane are active: accVGPRs 0–15 belong to block 0, accVGPRs 16–31 to block 1.

:math:`32 \times 32` MFMA accumulator layout -- accVGPR index per output element, with lane groups colour-coded.

\(32 \times 32\) accumulator layout – 2 blocks. Each cell shows the accVGPR index that holds output element \((i, j)\) of block 0; block 1 adds 16. 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)\) in block \(b\):

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

Conversely, given a lane \(L\) and accVGPR 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 \\ b &= \lfloor \frac{G}{16} \rfloor\end{split}\]

The row-to-lane mapping groups rows in bands of four. Within block 0:

Rows

Lanes

accVGPRs

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

Block 1 uses the same lane pattern with accVGPRs 16–31.

Builtins with a 1-block \(32 \times 32\) result (v16float / v16int, 16 accVGPRs per lane) use only block 0 of the layout above. These variants do not support the cbsz and abid modifiers.

\(16 \times 16\) layout#

The \(16 \times 16\) tile shape uses four blocks. The 16 accVGPRs per lane are partitioned by block: accVGPRs 0—3 for block 0, 4—7 for block 1, 8—11 for block 2, and 12—15 for block 3.

:math:`16 \times 16` MFMA accumulator layout -- accVGPR index per output element, with lane groups colour-coded by block assignment.

\(16 \times 16\) accumulator layout – 4 blocks. Each cell shows the accVGPR index for block 0; block \(k\) adds \(4k\). 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)\) in block \(b\):

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

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

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

The row-to-lane mapping (identical for every block):

Rows

Lanes

0—3

0—15

4—7

16—31

8—11

32—47

12—15

48—63

Builtins with a 1-block \(16 \times 16\) result (v4float / v4int, 4 accVGPRs per lane) use only block 0 of the layout above. These variants do not support the cbsz and abid modifiers.

\(4 \times 4\) layout#

The \(4 \times 4\) tile shape uses 16 blocks. A single wavefront simultaneously computes 16 independent \(4 \times 4\) outer products. Each group of 4 consecutive lanes (lanes \(4b\) through \(4b + 3\)) holds all 16 output elements of block \(b\) across 4 accVGPRs.

:math:`4 \times 4` MFMA accumulator layout -- all 16 blocks shown as consecutive groups of 4 lanes across the full wavefront.

\(4 \times 4\) accumulator layout – 16 blocks. Columns are lanes 0–63; each group of 4 consecutive lanes holds one complete \(4 \times 4\) output tile in accVGPRs 0–3 (rows). Teal groups are even-numbered blocks, grey groups odd-numbered blocks.#

Given output element \((i, j)\) in block \(b\):

\[\begin{split}\text{lane} &= 4 b + j \\ \text{accVGPR} &= i\end{split}\]

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

\[\begin{split}i &= G \bmod 4 \\ j &= L \bmod 4 \\ b &= \lfloor \frac{L}{4} \rfloor\end{split}\]

\(16 \times 16\) FP64 layout#

The \(16 \times 16\) FP64 tile shape uses one block. Each lane holds four FP64 output elements across 4 accVGPR pairs (8 physical accVGPRs, since each FP64 value occupies two 32-bit registers).

Note

FP64 accVGPR indices count pairs of physical registers. accVGPR pair \(k\) corresponds to physical registers v[2k+1:2k].

:math:`16 \times 16` FP64 MFMA accumulator layout -- accVGPR pair index per output element, with lane groups colour-coded.

\(16 \times 16\) FP64 accumulator layout – 1 block. Each cell shows the accVGPR pair index \(k\) that holds output element \((i, j)\); physical registers are v[2k+1:2k]. Each group of 16 consecutive lanes (column offset \(j\)) covers all four rows within one row-modulo-4 band.#

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

\[\begin{split}\text{lane} &= 16 \cdot (i \bmod 4) + j \\ \text{accVGPR pair}\ k &= \lfloor \frac{i}{4} \rfloor\end{split}\]

Conversely, given lane \(L\) and accVGPR pair index \(k\):

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

The row-to-lane mapping. Each group of 16 consecutive lanes covers one row modulo 4; the accVGPR pair selects the row group:

Rows

Lanes

accVGPR pair (physical regs)

0, 4, 8, 12

0—15

0, 1, 2, 3 (v[1:0], v[3:2], v[5:4], v[7:6])

1, 5, 9, 13

16—31

0, 1, 2, 3

2, 6, 10, 14

32—47

0, 1, 2, 3

3, 7, 11, 15

48—63

0, 1, 2, 3

\(4 \times 4\) FP64 layout#

The \(4 \times 4\) FP64 tile shape uses four blocks. Each lane holds one FP64 output element in a single accVGPR pair (2 physical accVGPRs, v[1:0]).

:math:`4 \times 4` FP64 MFMA accumulator layout -- all 4 blocks shown across the full 64-lane wavefront, grouped into four 16-lane row groups.

\(4 \times 4\) FP64 accumulator layout – 4 blocks. Columns are lanes 0–63; rows are matrix rows 0–3. Within each 16-lane row group, four sub-groups of 4 lanes hold blocks 0–3. Every cell holds one FP64 value in accVGPR pair 0 (v[1:0]). Teal sub-groups are even-numbered blocks, grey sub-groups are odd-numbered blocks.#

Given output element \((i, j)\) in block \(b\):

\[\text{lane} = 16 i + 4 b + j\]

Conversely, given lane \(L\):

\[\begin{split}i &= \lfloor \frac{L}{16} \rfloor \\ j &= L \bmod 4 \\ b &= \lfloor \frac{L \bmod 16}{4} \rfloor\end{split}\]

The row-to-lane mapping. Lanes 0—15 hold all four blocks of row 0; the pattern repeats for rows 1—3 at lane offsets of 16, 32, and 48:

Row

Block

Lanes

accVGPR pair

0

0

0—3

0 (v[1:0])

0

1

4—7

0 (v[1:0])

0

2

8—11

0 (v[1:0])

0

3

12—15

0 (v[1:0])

1

0—3

16—31

0 (v[1:0])

2

0—3

32—47

0 (v[1:0])

3

0—3

48—63

0 (v[1:0])

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 v2float   = float [[clang::ext_vector_type(2)]];   // xf32 storage
using v4float   = float [[clang::ext_vector_type(4)]];
using v16float  = float [[clang::ext_vector_type(16)]];
using v32float  = float [[clang::ext_vector_type(32)]];
using v4half    = _Float16 [[clang::ext_vector_type(4)]];
using v4int     = int [[clang::ext_vector_type(4)]];
using v16int    = int [[clang::ext_vector_type(16)]];
using v32int    = int [[clang::ext_vector_type(32)]];
using v2bfloat  = short [[clang::ext_vector_type(2)]];   // bf16 storage
using v4bfloat  = short [[clang::ext_vector_type(4)]];   // bf16 storage
using v4double  = double [[clang::ext_vector_type(4)]];

FP8 and BF8 input operands, and the wider-K INT8 operands, are passed as a 64-bit integer (long long) that packs eight 8-bit values per lane.

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#

See Common MFMA parameters for a complete description of the cbsz, abid, and blgp modifiers shared by all MFMA builtins, including the CDNA3-specific blgp behaviour for FP64 builtins.

Using MFMA builtins as a compute policy#

The matrix multiplication tutorial in Optimizing GEMM in HIP uses a ComputePolicy type parameter to separate the multiply-accumulate logic from the rest of the kernel. You can drop an MFMA-based policy into that framework without changing the outer kernel.

The example below implements MfmaCdna3XF32Policy using v_mfma_f32_16x16x8_xf32 – a \(16 \times 16\) XF32 builtin available on CDNA3 (gfx942, MI300 series). Each wavefront computes a single \(16 \times 16\) output tile; one block is active, giving 4 FP32 accVGPRs per lane.

The complete source file is available for download:

Policy constants

v_mfma_f32_16x16x8_xf32 consumes eight K-positions per call (K=8), so k_step = 8. Each lane provides two FP32 values (v2float) for \(\pmb{A}\) and two for \(\pmb{B}\); the 64 lanes partition into four groups of 16, each group covering two K-positions. The entire \(16 \times 16\) tile belongs to one wavefront, so thread_tile_m = thread_tile_n = 16 and effective_lanes = 64.

Accumulator layout

The builtin returns a v4float holding 4 FP32 values – one for each accVGPR. The Accumulator struct wraps this directly. The layout formulas from the Accumulator layout section translate (lane, accVGPR) back to \((i, j)\) coordinates during the store_c() pass.

Fragment loading

Because each lane provides two FP32 values for K=8, each load_a call reads two consecutive scalars. Lane group \(g = \lfloor \text{lane\_id} / 16 \rfloor\) covers K-positions ki + 2g and ki + 2g + 1; within the group, lane offset lane_id mod 16 selects one of the 16 \(\pmb{A}\)-rows (or \(\pmb{B}\)-columns). The mma() method packs the two values into a v2float before issuing the instruction.

Note

The all-modifier-zero constraint means cbsz, abid, and blgp must all be passed as 0. v_mfma_f32_16x16x8_xf32 is single-block only and does not support the cbsz/abid broadcast mechanism or the blgp lane-group pattern modifier.

struct MfmaCdna3XF32Policy
{
    // -- ComputePolicy constants ----------------------------------------------
    // One wavefront owns a 16×16 output tile.
    // All 64 lanes hold unique output elements (4 elements per lane).
    static constexpr int thread_tile_m   = 16;
    static constexpr int thread_tile_n   = 16;
    static constexpr int frag_size_m     = 2;
    static constexpr int frag_size_n     = 2;
    static constexpr int effective_lanes = 64;
    // v_mfma_f32_16x16x8_xf32 consumes 8 K-positions per call (k_step=8).
    // Each lane provides two FP32 values (v2float) for A and two for B.
    // The 64 lanes split into 4 groups of 16; each group covers 2 K-positions.
    static constexpr int k_step          = 8;

    using elem_a = float;
    using elem_b = float;

    // -- Accumulator ----------------------------------------------------------
    // v4float holds 4 FP32 accVGPRs per lane (accVGPRs 0-3).
    using v4float = float [[clang::ext_vector_type(4)]];

    struct Accumulator
    {
        v4float regs;
    };

    __device__ static void zero(Accumulator& acc)
    {
        acc.regs = v4float{};
    }

    // -- thread_tile_offset ---------------------------------------------------
    // The entire 16×16 tile belongs to one wavefront.  Multiple wavefronts
    // in a block cover different 16×16 sub-tiles.
    //
    //   wavefront id within block: wid = tid / 64
    //   waves_n = block_tile_n / thread_tile_n = 32 / 16 = 2
    //   wid_x = wid % waves_n
    //   wid_y = wid / waves_n
    //   *thread_row = wid_y * 16
    //   *thread_col = wid_x * 16
    __device__ static void thread_tile_offset(int  tid,
                                              int  /*lane_id*/,
                                              int* thread_row,
                                              int* thread_col)
    {
        constexpr int waves_n = 2; // block_tile_n / thread_tile_n = 32 / 16
        const int wid   = tid / 64;
        const int wid_x = wid % waves_n;
        const int wid_y = wid / waves_n;
        *thread_row = wid_y * 16;
        *thread_col = wid_x * 16;
    }

    // -- Fragment loads --------------------------------------------------------
    // load_a: each lane reads two floats from LDS into a v2float.
    //
    // v_mfma_f32_16x16x8_xf32 (1-block) interprets the 64 lanes as 4 groups
    // of 16 (indexed by g = lane_id / 16).  Each group covers 2 consecutive
    // K-positions: group g provides data for ki + 2*g and ki + 2*g + 1.
    // Within group g, lane offset (lane_id mod 16) selects one of the 16
    // A-rows.
    //
    //   A-row index  : tile_a_row + (lane_id mod 16)
    //   K-position 0 : ki + 2*(lane_id / 16)
    //   K-position 1 : ki + 2*(lane_id / 16) + 1
    __device__ static void load_a(const float* tile_a_ptr,
                                  int          tile_a_row,
                                  int          ki,
                                  int          k_tile_size,
                                  int          lane_id,
                                  elem_a     (&frag)[2])
    {
        const int row   = tile_a_row + (lane_id % 16);
        const int k_off = ki + 2 * (lane_id / 16);
        frag[0] = tile_a_ptr[row * k_tile_size + k_off];
        frag[1] = tile_a_ptr[row * k_tile_size + k_off + 1];
    }

    // load_b: each lane reads two floats from LDS (symmetric to load_a).
    //
    // tile_b_T is column-major: tile_b_T[(tile_b_col + col)][k_off].
    //   B-col index  : tile_b_col + (lane_id mod 16)
    //   K-position 0 : ki + 2*(lane_id / 16)
    //   K-position 1 : ki + 2*(lane_id / 16) + 1
    __device__ static void load_b(const float* tile_b_T_ptr,
                                  int          tile_b_col,
                                  int          ki,
                                  int          k_tile_size,
                                  int          lane_id,
                                  elem_b     (&frag)[2])
    {
        const int col   = tile_b_col + (lane_id % 16);
        const int k_off = ki + 2 * (lane_id / 16);
        frag[0] = tile_b_T_ptr[col * k_tile_size + k_off];
        frag[1] = tile_b_T_ptr[col * k_tile_size + k_off + 1];
    }

    // -- mma -------------------------------------------------------------------
    // Issue one v_mfma_f32_16x16x8_xf32 instruction.
    // This intrinsic is single-block only; all modifier parameters must be 0.
    // cbsz=0, abid=0: single-block A-operand, no broadcast.
    // blgp=0: not supported; must be 0.
    __device__ static void mma(Accumulator&       acc,
                               const elem_a     (&a_frag)[2],
                               const elem_b     (&b_frag)[2])
    {
#if defined(__gfx942__)
        using v2float = float [[clang::ext_vector_type(2)]];
        const v2float a_vec = {a_frag[0], a_frag[1]};
        const v2float b_vec = {b_frag[0], b_frag[1]};
        acc.regs = __builtin_amdgcn_mfma_f32_16x16x8_xf32(
            a_vec, b_vec, acc.regs,
            /*cbsz=*/0, /*abid=*/0, /*blgp=*/0);
#endif
    }

    // -- store_c ---------------------------------------------------------------
    // Scatter the 4 accVGPR values to their global-memory positions.
    //
    // Inverse layout (lane L, accVGPR index G in [0,3]):
    //   i = 4 * floor(L / 16) + (G mod 4)
    //   j = L mod 16
    __device__ static void store_c(const Accumulator& acc,
                                   float*             C,
                                   int                out_row_base,
                                   int                out_col_base,
                                   int                m,
                                   int                n,
                                   int                lane_id)
    {
        const int L = lane_id;
        const int j = L % 16;

        #pragma unroll
        for(int G = 0; G < 4; ++G)
        {
            const int i = 4 * (L / 16) + (G % 4);
            const int r = out_row_base + i;
            const int c = out_col_base + j;
            if(r < m && c < n)
                C[r * n + c] = acc.regs[G];
        }
    }
};

Instantiating the kernel

With MfmaCdna3XF32Policy in hand, plug it into the generic kernel alongside a TilePolicy whose block_tile_m and block_tile_n are multiples of 16 and whose k_tile_size is a multiple of k_step = 8. The policy aliases and launch configuration from the example file are:

// TilePolicy: 32x32 block tile, 8-element K-strip, single-buffered.
using MfmaTilePolicy = SingleBufferTilePolicyF<32, 32, 8>;

// Launch parameters:
//
// CDNA3 (gfx942): 2x2 wavefronts of 64 lanes = 256 threads per block.
//   block_tile_m = 32, thread_tile_m = 16  → 2 wavefronts along M
//   block_tile_n = 32, thread_tile_n = 16  → 2 wavefronts along N
//
// Fallback (scalar): 8x8 thread tiles in a 32x32 block → 16 threads per block,
//   padded to one full wavefront of 64.
//
// The kernel to launch and its thread count are selected at runtime based on
// the device architecture string returned by hipGetDeviceProperties.
//
// Grid: one block per 32x32 output tile.
constexpr int BLOCK_TILE_M = 32;
constexpr int BLOCK_TILE_N = 32;

    // Grid: one block per 32×32 output tile.
    const dim3 grid((N + BLOCK_TILE_N - 1) / BLOCK_TILE_N,
                    (M + BLOCK_TILE_M - 1) / BLOCK_TILE_M);

    float       ms = 0.0f;
    const char* policy_label = nullptr;

    if(is_gfx942)
    {
        auto launch = [&]()
        {
            matrix_multiply_generic<MfmaTilePolicy, MfmaCdna3XF32Policy>
                <<<grid, dim3(256)>>>(d_A, d_B, d_C, M, N, K);
            HIP_CHECK(hipGetLastError());
        };
        ms           = time_kernel_ms(launch, WARMUP_RUNS, TIMING_RUNS);
        policy_label = "GemmKernel<MfmaTilePolicy, MfmaCdna3XF32Policy> [gfx942]";
    }
    else
    {
        auto launch = [&]()
        {
            matrix_multiply_generic<MfmaTilePolicy, ScalarFMAPolicy<8, 8, 8>>
                <<<grid, dim3(64)>>>(d_A, d_B, d_C, M, N, K);
            HIP_CHECK(hipGetLastError());
        };
        ms           = time_kernel_ms(launch, WARMUP_RUNS, TIMING_RUNS);
        policy_label = "GemmKernel<MfmaTilePolicy, ScalarFMAPolicy> [fallback]";
    }

Compile and run:

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

Note

MfmaCdna3XF32Policy requires a CDNA3 GPU (gfx942). Compile with --offload-arch=gfx942 to select the correct architecture. On other targets the #if defined(__gfx942__) guard selects ScalarFMAPolicy automatically, so the file compiles without modification.

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}\). Instructions that support vector ALU (VALU) co-execution allow the compiler to overlap matrix and vector work; the VALU co-execution cycle count gives the number of VALU cycles available during the MFMA latency window. A value of 0 means VALU co-execution is not supported.

Builtin

Ops

Cycle count

VALU co-execution cycles

__builtin_amdgcn_mfma_f32_32x32x1f32

4096

64

0

__builtin_amdgcn_mfma_f32_16x16x1f32

2048

32

0

__builtin_amdgcn_mfma_f32_4x4x1f32

512

8

0

__builtin_amdgcn_mfma_f32_32x32x2f32

4096

64

0

__builtin_amdgcn_mfma_f32_16x16x4f32

2048

32

0

__builtin_amdgcn_mfma_f32_32x32x4f16

16384

64

60

__builtin_amdgcn_mfma_f32_16x16x4f16

8192

32

28

__builtin_amdgcn_mfma_f32_4x4x4f16

2048

8

4

__builtin_amdgcn_mfma_f32_32x32x8f16

16384

32

28

__builtin_amdgcn_mfma_f32_16x16x16f16

8192

16

12

__builtin_amdgcn_mfma_f32_32x32x4bf16_1k

16384

64

60

__builtin_amdgcn_mfma_f32_16x16x4bf16_1k

8192

32

28

__builtin_amdgcn_mfma_f32_4x4x4bf16_1k

2048

8

4

__builtin_amdgcn_mfma_f32_32x32x8bf16_1k

16384

32

28

__builtin_amdgcn_mfma_f32_16x16x16bf16_1k

8192

16

12

__builtin_amdgcn_mfma_f32_32x32x4_xf32

8192

32

28

__builtin_amdgcn_mfma_f32_16x16x8_xf32

4096

16

12

__builtin_amdgcn_mfma_f32_32x32x16_fp8_fp8

32768

32

28

__builtin_amdgcn_mfma_f32_32x32x16_fp8_bf8

32768

32

28

__builtin_amdgcn_mfma_f32_32x32x16_bf8_fp8

32768

32

28

__builtin_amdgcn_mfma_f32_32x32x16_bf8_bf8

32768

32

28

__builtin_amdgcn_mfma_f32_16x16x32_fp8_fp8

16384

16

12

__builtin_amdgcn_mfma_f32_16x16x32_fp8_bf8

16384

16

12

__builtin_amdgcn_mfma_f32_16x16x32_bf8_fp8

16384

16

12

__builtin_amdgcn_mfma_f32_16x16x32_bf8_bf8

16384

16

12

__builtin_amdgcn_mfma_f64_16x16x4f64

2048

32

0

__builtin_amdgcn_mfma_f64_4x4x4f64

512

16

0

__builtin_amdgcn_mfma_i32_32x32x4i8

16384

64

60

__builtin_amdgcn_mfma_i32_16x16x4i8

8192

32

28

__builtin_amdgcn_mfma_i32_4x4x4i8

2048

8

4

__builtin_amdgcn_mfma_i32_32x32x16_i8

32768

32

28

__builtin_amdgcn_mfma_i32_16x16x32_i8

16384

16

12

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.

FP32 matrix inputs#

The following builtins accept one FP32 element per lane for both \(\pmb{A}\) and \(\pmb{B}\) inputs and accumulate into FP32 output fragments.

__builtin_amdgcn_mfma_f32_32x32x1f32#

Signature and parameters for this builtin.

v32float __builtin_amdgcn_mfma_f32_32x32x1f32(
    float    srcA,
    float    srcB,
    v32float srcC,
    int      cbsz,
    int      abid,
    int      blgp);

Computes one step of a \(32 \times 32\) FP32 outer-product accumulation. Each instruction processes \(K=1\) FP32 column element of \(\pmb{A}\) and \(K=1\) FP32 row element of \(\pmb{B}\), accumulating into a \(32 \times 32\) FP32 tile held across 32 accVGPRs per lane. This builtin operates on two input blocks.

Parameter

Type

Description

srcA

float

One element of \(\pmb{A}\) per lane.

srcB

float

One element of \(\pmb{B}\) per lane.

srcC

v32float

Accumulator input: 32 FP32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Legal range: 0..1

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Legal range: 0..1

blgp

int

\(\pmb{B}\)-matrix Lane Group Pattern modifier, see Common MFMA parameters.

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

__builtin_amdgcn_mfma_f32_16x16x1f32#

Signature and parameters for this builtin.

v16float __builtin_amdgcn_mfma_f32_16x16x1f32(
    float    srcA,
    float    srcB,
    v16float srcC,
    int      cbsz,
    int      abid,
    int      blgp);

Computes one step of a \(16 \times 16\) FP32 outer-product accumulation. Each instruction processes \(K=1\) FP32 column element of \(\pmb{A}\) and \(K=1\) FP32 row element of \(\pmb{B}\), accumulating into a \(16 \times 16\) FP32 tile held across 16 accVGPRs per lane. This builtin operates on four input blocks.

Parameter

Type

Description

srcA

float

One element of \(\pmb{A}\) per lane.

srcB

float

One element of \(\pmb{B}\) per lane.

srcC

v16float

Accumulator input: 16 FP32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Legal range: 0..2

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Legal range: 0..3

blgp

int

\(\pmb{B}\)-matrix Lane Group Pattern modifier, see Common MFMA parameters.

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

__builtin_amdgcn_mfma_f32_4x4x1f32#

Signature and parameters for this builtin.

v4float __builtin_amdgcn_mfma_f32_4x4x1f32(
    float   srcA,
    float   srcB,
    v4float srcC,
    int     cbsz,
    int     abid,
    int     blgp);

Computes one step of a \(4 \times 4\) FP32 outer-product accumulation. Each instruction processes \(K=1\) FP32 column element of \(\pmb{A}\) and \(K=1\) FP32 row element of \(\pmb{B}\), accumulating into a \(4 \times 4\) FP32 tile held across 4 accVGPRs per lane. This builtin operates on 16 input blocks.

Parameter

Type

Description

srcA

float

One element of \(\pmb{A}\) per lane.

srcB

float

One element of \(\pmb{B}\) per lane.

srcC

v4float

Accumulator input: 4 FP32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Legal range: 0..4

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Legal range: 0..15

blgp

int

\(\pmb{B}\)-matrix Lane Group Pattern modifier, see Common MFMA parameters.

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

__builtin_amdgcn_mfma_f32_32x32x2f32#

Signature and parameters for this builtin.

v16float __builtin_amdgcn_mfma_f32_32x32x2f32(
    float    srcA,
    float    srcB,
    v16float srcC,
    int      cbsz,
    int      abid,
    int      blgp);

Computes one step of a \(32 \times 32\) FP32 outer-product accumulation. Each instruction processes \(K=2\) FP32 column elements of \(\pmb{A}\) and \(K=2\) FP32 row elements of \(\pmb{B}\), accumulating into a \(32 \times 32\) FP32 tile held across 16 accVGPRs per lane. This builtin operates on a single input block.

Parameter

Type

Description

srcA

float

One element of \(\pmb{A}\) per lane.

srcB

float

One element of \(\pmb{B}\) per lane.

srcC

v16float

Accumulator input: 16 FP32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Must be 0.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Must be 0.

blgp

int

\(\pmb{B}\)-matrix Lane Group Pattern modifier, see Common MFMA parameters.

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

__builtin_amdgcn_mfma_f32_16x16x4f32#

Signature and parameters for this builtin.

v4float __builtin_amdgcn_mfma_f32_16x16x4f32(
    float   srcA,
    float   srcB,
    v4float srcC,
    int     cbsz,
    int     abid,
    int     blgp);

Computes one step of a \(16 \times 16\) FP32 outer-product accumulation. Each instruction processes \(K=4\) FP32 column elements of \(\pmb{A}\) and \(K=4\) FP32 row elements of \(\pmb{B}\), accumulating into a \(16 \times 16\) FP32 tile held across 4 accVGPRs per lane. This builtin operates on a single input block.

Parameter

Type

Description

srcA

float

One element of \(\pmb{A}\) per lane.

srcB

float

One element of \(\pmb{B}\) per lane.

srcC

v4float

Accumulator input: 4 FP32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Must be 0.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Must be 0.

blgp

int

\(\pmb{B}\)-matrix Lane Group Pattern modifier, see Common MFMA parameters.

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

FP16 matrix inputs#

The following builtins accept four FP16 elements per lane packed into a v4half register for both \(\pmb{A}\) and \(\pmb{B}\) inputs, and accumulate into FP32 output fragments.

__builtin_amdgcn_mfma_f32_32x32x4f16#

Signature and parameters for this builtin.

v32float __builtin_amdgcn_mfma_f32_32x32x4f16(
    v4half   srcA,
    v4half   srcB,
    v32float srcC,
    int      cbsz,
    int      abid,
    int      blgp);

Computes one step of a \(32 \times 32\) FP32 outer-product accumulation. Each instruction processes \(K=4\) FP16 column elements of \(\pmb{A}\) and \(K=4\) FP16 row elements of \(\pmb{B}\), accumulating into a \(32 \times 32\) FP32 tile held across 32 accVGPRs per lane. This builtin operates on two input blocks.

Parameter

Type

Description

srcA

v4half

Four elements of \(\pmb{A}\) per lane.

srcB

v4half

Four elements of \(\pmb{B}\) per lane.

srcC

v32float

Accumulator input: 32 FP32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Legal range: 0..1

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Legal range: 0..1

blgp

int

\(\pmb{B}\)-matrix Lane Group Pattern modifier, see Common MFMA parameters.

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

__builtin_amdgcn_mfma_f32_16x16x4f16#

Signature and parameters for this builtin.

v16float __builtin_amdgcn_mfma_f32_16x16x4f16(
    v4half   srcA,
    v4half   srcB,
    v16float srcC,
    int      cbsz,
    int      abid,
    int      blgp);

Computes one step of a \(16 \times 16\) FP32 outer-product accumulation. Each instruction processes \(K=4\) FP16 column elements of \(\pmb{A}\) and \(K=4\) FP16 row elements of \(\pmb{B}\), accumulating into a \(16 \times 16\) FP32 tile held across 16 accVGPRs per lane. This builtin operates on four input blocks.

Parameter

Type

Description

srcA

v4half

Four elements of \(\pmb{A}\) per lane.

srcB

v4half

Four elements of \(\pmb{B}\) per lane.

srcC

v16float

Accumulator input: 16 FP32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Legal range: 0..2

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Legal range: 0..3

blgp

int

\(\pmb{B}\)-matrix Lane Group Pattern modifier, see Common MFMA parameters.

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

__builtin_amdgcn_mfma_f32_4x4x4f16#

Signature and parameters for this builtin.

v4float __builtin_amdgcn_mfma_f32_4x4x4f16(
    v4half  srcA,
    v4half  srcB,
    v4float srcC,
    int     cbsz,
    int     abid,
    int     blgp);

Computes one step of a \(4 \times 4\) FP32 outer-product accumulation. Each instruction processes \(K=4\) FP16 column elements of \(\pmb{A}\) and \(K=4\) FP16 row elements of \(\pmb{B}\), accumulating into a \(4 \times 4\) FP32 tile held across four accVGPRs per lane. This builtin operates on 16 input blocks.

Parameter

Type

Description

srcA

v4half

Four elements of \(\pmb{A}\) per lane.

srcB

v4half

Four elements of \(\pmb{B}\) per lane.

srcC

v4float

Accumulator input: Four FP32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Legal range: 0..4

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Legal range: 0..15

blgp

int

\(\pmb{B}\)-matrix Lane Group Pattern modifier, see Common MFMA parameters.

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

__builtin_amdgcn_mfma_f32_32x32x8f16#

Signature and parameters for this builtin.

v16float __builtin_amdgcn_mfma_f32_32x32x8f16(
    v4half   srcA,
    v4half   srcB,
    v16float srcC,
    int      cbsz,
    int      abid,
    int      blgp);

Computes one step of a \(32 \times 32\) FP32 outer-product accumulation. Each instruction processes \(K=8\) FP16 column elements of \(\pmb{A}\) and \(K=8\) FP16 row elements of \(\pmb{B}\), accumulating into a \(32 \times 32\) FP32 tile held across 16 accVGPRs per lane. This builtin operates on a single input block.

Parameter

Type

Description

srcA

v4half

Four elements of \(\pmb{A}\) per lane.

srcB

v4half

Four elements of \(\pmb{B}\) per lane.

srcC

v16float

Accumulator input: 16 FP32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Must be 0.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Must be 0.

blgp

int

\(\pmb{B}\)-matrix Lane Group Pattern modifier, see Common MFMA parameters.

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

__builtin_amdgcn_mfma_f32_16x16x16f16#

Signature and parameters for this builtin.

v4float __builtin_amdgcn_mfma_f32_16x16x16f16(
    v4half  srcA,
    v4half  srcB,
    v4float srcC,
    int     cbsz,
    int     abid,
    int     blgp);

Computes one step of a \(16 \times 16\) FP32 outer-product accumulation. Each instruction processes \(K=16\) FP16 column elements of \(\pmb{A}\) and \(K=16\) FP16 row elements of \(\pmb{B}\), accumulating into a \(16 \times 16\) FP32 tile held across four accVGPRs per lane. This builtin operates on a single input block.

Parameter

Type

Description

srcA

v4half

Four elements of \(\pmb{A}\) per lane.

srcB

v4half

Four elements of \(\pmb{B}\) per lane.

srcC

v4float

Accumulator input: Four FP32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Must be 0.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Must be 0.

blgp

int

\(\pmb{B}\)-matrix Lane Group Pattern modifier, see Common MFMA parameters.

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

BF16 matrix inputs#

CDNA3 supports only the _1k BF16 variants. These builtins accept four BF16 elements per lane packed into a v4bfloat register for both \(\pmb{A}\) and \(\pmb{B}\) inputs, and accumulate into FP32 output fragments.

__builtin_amdgcn_mfma_f32_32x32x4bf16_1k#

Signature and parameters for this builtin.

v32float __builtin_amdgcn_mfma_f32_32x32x4bf16_1k(
    v4bfloat srcA,
    v4bfloat srcB,
    v32float srcC,
    int      cbsz,
    int      abid,
    int      blgp);

Computes one step of a \(32 \times 32\) FP32 outer-product accumulation. Each instruction processes \(K=4\) BF16 column elements of \(\pmb{A}\) and \(K=4\) BF16 row elements of \(\pmb{B}\), accumulating into a \(32 \times 32\) FP32 tile held across 32 accVGPRs per lane. This builtin operates on two input blocks.

Parameter

Type

Description

srcA

v4bfloat

Four BF16 elements of \(\pmb{A}\) per lane.

srcB

v4bfloat

Four BF16 elements of \(\pmb{B}\) per lane.

srcC

v32float

Accumulator input: 32 FP32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Legal range: 0..1

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Legal range: 0..1

blgp

int

\(\pmb{B}\)-matrix Lane Group Pattern modifier, see Common MFMA parameters.

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

__builtin_amdgcn_mfma_f32_16x16x4bf16_1k#

Signature and parameters for this builtin.

v16float __builtin_amdgcn_mfma_f32_16x16x4bf16_1k(
    v4bfloat srcA,
    v4bfloat srcB,
    v16float srcC,
    int      cbsz,
    int      abid,
    int      blgp);

Computes one step of a \(16 \times 16\) FP32 outer-product accumulation. Each instruction processes \(K=4\) BF16 column elements of \(\pmb{A}\) and \(K=4\) BF16 row elements of \(\pmb{B}\), accumulating into a \(16 \times 16\) FP32 tile held across 16 accVGPRs per lane. This builtin operates on four input blocks.

Parameter

Type

Description

srcA

v4bfloat

Four BF16 elements of \(\pmb{A}\) per lane.

srcB

v4bfloat

Four BF16 elements of \(\pmb{B}\) per lane.

srcC

v16float

Accumulator input: 16 FP32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Legal range: 0..2

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Legal range: 0..3

blgp

int

\(\pmb{B}\)-matrix Lane Group Pattern modifier, see Common MFMA parameters.

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

__builtin_amdgcn_mfma_f32_4x4x4bf16_1k#

Signature and parameters for this builtin.

v4float __builtin_amdgcn_mfma_f32_4x4x4bf16_1k(
    v4bfloat srcA,
    v4bfloat srcB,
    v4float  srcC,
    int      cbsz,
    int      abid,
    int      blgp);

Computes one step of a \(4 \times 4\) FP32 outer-product accumulation. Each instruction processes \(K=4\) BF16 column elements of \(\pmb{A}\) and \(K=4\) BF16 row elements of \(\pmb{B}\), accumulating into a \(4 \times 4\) FP32 tile held across 4 accVGPRs per lane. This builtin operates on 16 input blocks.

Parameter

Type

Description

srcA

v4bfloat

Four BF16 elements of \(\pmb{A}\) per lane.

srcB

v4bfloat

Four BF16 elements of \(\pmb{B}\) per lane.

srcC

v4float

Accumulator input: 4 FP32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Legal range: 0..4

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Legal range: 0..15

blgp

int

\(\pmb{B}\)-matrix Lane Group Pattern modifier, see Common MFMA parameters.

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

__builtin_amdgcn_mfma_f32_32x32x8bf16_1k#

Signature and parameters for this builtin.

v16float __builtin_amdgcn_mfma_f32_32x32x8bf16_1k(
    v4bfloat srcA,
    v4bfloat srcB,
    v16float srcC,
    int      cbsz,
    int      abid,
    int      blgp);

Computes one step of a \(32 \times 32\) FP32 outer-product accumulation. Each instruction processes \(K=8\) BF16 column elements of \(\pmb{A}\) and \(K=8\) BF16 row elements of \(\pmb{B}\), accumulating into a \(32 \times 32\) FP32 tile held across 16 accVGPRs per lane. This builtin operates on a single input block.

Parameter

Type

Description

srcA

v4bfloat

Four BF16 elements of \(\pmb{A}\) per lane.

srcB

v4bfloat

Four BF16 elements of \(\pmb{B}\) per lane.

srcC

v16float

Accumulator input: 16 FP32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Must be 0.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Must be 0.

blgp

int

\(\pmb{B}\)-matrix Lane Group Pattern modifier, see Common MFMA parameters.

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

__builtin_amdgcn_mfma_f32_16x16x16bf16_1k#

Signature and parameters for this builtin.

v4float __builtin_amdgcn_mfma_f32_16x16x16bf16_1k(
    v4bfloat srcA,
    v4bfloat srcB,
    v4float  srcC,
    int      cbsz,
    int      abid,
    int      blgp);

Computes one step of a \(16 \times 16\) FP32 outer-product accumulation. Each instruction processes \(K=16\) BF16 column elements of \(\pmb{A}\) and \(K=16\) BF16 row elements of \(\pmb{B}\), accumulating into a \(16 \times 16\) FP32 tile held across 4 accVGPRs per lane. This builtin operates on a single input block.

Parameter

Type

Description

srcA

v4bfloat

Four BF16 elements of \(\pmb{A}\) per lane.

srcB

v4bfloat

Four BF16 elements of \(\pmb{B}\) per lane.

srcC

v4float

Accumulator input: 4 FP32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Must be 0.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Must be 0.

blgp

int

\(\pmb{B}\)-matrix Lane Group Pattern modifier, see Common MFMA parameters.

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

XF32 matrix inputs#

eXtended Float32 (XF32) accepts standard FP32 (IEEE binary32) inputs but rounds each mantissa to 10 bits before multiplication, providing higher throughput than full FP32 matrix multiplication. Each lane passes two values packed into a v2float register for both \(\pmb{A}\) and \(\pmb{B}\) inputs.

__builtin_amdgcn_mfma_f32_32x32x4_xf32#

Signature and parameters for this builtin.

v16float __builtin_amdgcn_mfma_f32_32x32x4_xf32(
    v2float  srcA,
    v2float  srcB,
    v16float srcC,
    int      cbsz,
    int      abid,
    int      blgp);

Computes one step of a \(32 \times 32\) FP32 outer-product accumulation. Each instruction processes \(K=4\) XF32 column elements of \(\pmb{A}\) and \(K=4\) XF32 row elements of \(\pmb{B}\), accumulating into a \(32 \times 32\) FP32 tile held across 16 accVGPRs per lane. This builtin operates on a single input block.

eXtended Float32 (XF32) accepts standard FP32 (IEEE binary32) inputs but rounds each mantissa to 10 bits before multiplication, providing higher throughput than full FP32 matrix multiplication. Each lane provides two values packed into a v2float register.

Parameter

Type

Description

srcA

v2float

Two XF32 elements of \(\pmb{A}\) per lane.

srcB

v2float

Two XF32 elements of \(\pmb{B}\) per lane.

srcC

v16float

Accumulator input: 16 FP32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Must be 0.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Must be 0.

blgp

int

Must be 0.

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

__builtin_amdgcn_mfma_f32_16x16x8_xf32#

Signature and parameters for this builtin.

v4float __builtin_amdgcn_mfma_f32_16x16x8_xf32(
    v2float srcA,
    v2float srcB,
    v4float srcC,
    int     cbsz,
    int     abid,
    int     blgp);

Computes one step of a \(16 \times 16\) FP32 outer-product accumulation. Each instruction processes \(K=8\) XF32 column elements of \(\pmb{A}\) and \(K=8\) XF32 row elements of \(\pmb{B}\), accumulating into a \(16 \times 16\) FP32 tile held across 4 accVGPRs per lane. This builtin operates on a single input block.

eXtended Float32 (XF32) accepts standard FP32 (IEEE binary32) inputs but rounds each mantissa to 10 bits before multiplication, providing higher throughput than full FP32 matrix multiplication. Each lane provides two values packed into a v2float register.

Parameter

Type

Description

srcA

v2float

Two XF32 elements of \(\pmb{A}\) per lane.

srcB

v2float

Two XF32 elements of \(\pmb{B}\) per lane.

srcC

v4float

Accumulator input: 4 FP32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Must be 0.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Must be 0.

blgp

int

Must be 0.

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

FP8 and BF8 matrix inputs#

The following builtins accept eight 8-bit floating-point values per lane packed into a long long register. Two 8-bit formats are supported:

  • FP8 (fp8): E4M3 format (1 sign, 4 exponent, 3 mantissa bits).

  • BF8 (bf8): E5M2 format (1 sign, 5 exponent, 2 mantissa bits).

The \(\pmb{A}\) and \(\pmb{B}\) matrices can use different formats, giving four combinations per tile shape.

__builtin_amdgcn_mfma_f32_32x32x16_fp8_fp8#

Signature and parameters for this builtin.

v16float __builtin_amdgcn_mfma_f32_32x32x16_fp8_fp8(
    long long srcA,
    long long srcB,
    v16float  srcC,
    int       cbsz,
    int       abid,
    int       blgp);

Computes one step of a \(32 \times 32\) FP32 outer-product accumulation. Each instruction processes \(K=16\) FP8 column elements of \(\pmb{A}\) and \(K=16\) FP8 row elements of \(\pmb{B}\), accumulating into a \(32 \times 32\) FP32 tile held across 16 accVGPRs per lane. This builtin operates on a single input block.

Both \(\pmb{A}\) and \(\pmb{B}\) use the FP8 (E4M3) format: eight 8-bit values packed into a long long register per lane.

Parameter

Type

Description

srcA

long long

Eight FP8 (E4M3) elements of \(\pmb{A}\) per lane.

srcB

long long

Eight FP8 (E4M3) elements of \(\pmb{B}\) per lane.

srcC

v16float

Accumulator input: 16 FP32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Must be 0.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Must be 0.

blgp

int

Must be 0.

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

__builtin_amdgcn_mfma_f32_32x32x16_fp8_bf8#

Signature and parameters for this builtin.

v16float __builtin_amdgcn_mfma_f32_32x32x16_fp8_bf8(
    long long srcA,
    long long srcB,
    v16float  srcC,
    int       cbsz,
    int       abid,
    int       blgp);

Computes one step of a \(32 \times 32\) FP32 outer-product accumulation. Each instruction processes \(K=16\) FP8 column elements of \(\pmb{A}\) and \(K=16\) BF8 row elements of \(\pmb{B}\), accumulating into a \(32 \times 32\) FP32 tile held across 16 accVGPRs per lane. This builtin operates on a single input block.

\(\pmb{A}\) uses the FP8 (E4M3) format and \(\pmb{B}\) uses the BF8 (E5M2) format; eight 8-bit values are packed into a long long register per lane for each operand.

Parameter

Type

Description

srcA

long long

Eight FP8 (E4M3) elements of \(\pmb{A}\) per lane.

srcB

long long

Eight BF8 (E5M2) elements of \(\pmb{B}\) per lane.

srcC

v16float

Accumulator input: 16 FP32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Must be 0.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Must be 0.

blgp

int

Must be 0.

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

__builtin_amdgcn_mfma_f32_32x32x16_bf8_fp8#

Signature and parameters for this builtin.

v16float __builtin_amdgcn_mfma_f32_32x32x16_bf8_fp8(
    long long srcA,
    long long srcB,
    v16float  srcC,
    int       cbsz,
    int       abid,
    int       blgp);

Computes one step of a \(32 \times 32\) FP32 outer-product accumulation. Each instruction processes \(K=16\) BF8 column elements of \(\pmb{A}\) and \(K=16\) FP8 row elements of \(\pmb{B}\), accumulating into a \(32 \times 32\) FP32 tile held across 16 accVGPRs per lane. This builtin operates on a single input block.

\(\pmb{A}\) uses the BF8 (E5M2) format and \(\pmb{B}\) uses the FP8 (E4M3) format; eight 8-bit values are packed into a long long register per lane for each operand.

Parameter

Type

Description

srcA

long long

Eight BF8 (E5M2) elements of \(\pmb{A}\) per lane.

srcB

long long

Eight FP8 (E4M3) elements of \(\pmb{B}\) per lane.

srcC

v16float

Accumulator input: 16 FP32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Must be 0.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Must be 0.

blgp

int

Must be 0.

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

__builtin_amdgcn_mfma_f32_32x32x16_bf8_bf8#

Signature and parameters for this builtin.

v16float __builtin_amdgcn_mfma_f32_32x32x16_bf8_bf8(
    long long srcA,
    long long srcB,
    v16float  srcC,
    int       cbsz,
    int       abid,
    int       blgp);

Computes one step of a \(32 \times 32\) FP32 outer-product accumulation. Each instruction processes \(K=16\) BF8 column elements of \(\pmb{A}\) and \(K=16\) BF8 row elements of \(\pmb{B}\), accumulating into a \(32 \times 32\) FP32 tile held across 16 accVGPRs per lane. This builtin operates on a single input block.

Both \(\pmb{A}\) and \(\pmb{B}\) use the BF8 (E5M2) format: eight 8-bit values packed into a long long register per lane.

Parameter

Type

Description

srcA

long long

Eight BF8 (E5M2) elements of \(\pmb{A}\) per lane.

srcB

long long

Eight BF8 (E5M2) elements of \(\pmb{B}\) per lane.

srcC

v16float

Accumulator input: 16 FP32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Must be 0.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Must be 0.

blgp

int

Must be 0.

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

__builtin_amdgcn_mfma_f32_16x16x32_fp8_fp8#

Signature and parameters for this builtin.

v4float __builtin_amdgcn_mfma_f32_16x16x32_fp8_fp8(
    long long srcA,
    long long srcB,
    v4float   srcC,
    int       cbsz,
    int       abid,
    int       blgp);

Computes one step of a \(16 \times 16\) FP32 outer-product accumulation. Each instruction processes \(K=32\) FP8 column elements of \(\pmb{A}\) and \(K=32\) FP8 row elements of \(\pmb{B}\), accumulating into a \(16 \times 16\) FP32 tile held across 4 accVGPRs per lane. This builtin operates on a single input block.

Both \(\pmb{A}\) and \(\pmb{B}\) use the FP8 (E4M3) format: eight 8-bit values packed into a long long register per lane.

Parameter

Type

Description

srcA

long long

Eight FP8 (E4M3) elements of \(\pmb{A}\) per lane.

srcB

long long

Eight FP8 (E4M3) elements of \(\pmb{B}\) per lane.

srcC

v4float

Accumulator input: 4 FP32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Must be 0.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Must be 0.

blgp

int

Must be 0.

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

__builtin_amdgcn_mfma_f32_16x16x32_fp8_bf8#

Signature and parameters for this builtin.

v4float __builtin_amdgcn_mfma_f32_16x16x32_fp8_bf8(
    long long srcA,
    long long srcB,
    v4float   srcC,
    int       cbsz,
    int       abid,
    int       blgp);

Computes one step of a \(16 \times 16\) FP32 outer-product accumulation. Each instruction processes \(K=32\) FP8 column elements of \(\pmb{A}\) and \(K=32\) BF8 row elements of \(\pmb{B}\), accumulating into a \(16 \times 16\) FP32 tile held across 4 accVGPRs per lane. This builtin operates on a single input block.

\(\pmb{A}\) uses the FP8 (E4M3) format and \(\pmb{B}\) uses the BF8 (E5M2) format; eight 8-bit values are packed into a long long register per lane for each operand.

Parameter

Type

Description

srcA

long long

Eight FP8 (E4M3) elements of \(\pmb{A}\) per lane.

srcB

long long

Eight BF8 (E5M2) elements of \(\pmb{B}\) per lane.

srcC

v4float

Accumulator input: 4 FP32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Must be 0.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Must be 0.

blgp

int

Must be 0.

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

__builtin_amdgcn_mfma_f32_16x16x32_bf8_fp8#

Signature and parameters for this builtin.

v4float __builtin_amdgcn_mfma_f32_16x16x32_bf8_fp8(
    long long srcA,
    long long srcB,
    v4float   srcC,
    int       cbsz,
    int       abid,
    int       blgp);

Computes one step of a \(16 \times 16\) FP32 outer-product accumulation. Each instruction processes \(K=32\) BF8 column elements of \(\pmb{A}\) and \(K=32\) FP8 row elements of \(\pmb{B}\), accumulating into a \(16 \times 16\) FP32 tile held across 4 accVGPRs per lane. This builtin operates on a single input block.

\(\pmb{A}\) uses the BF8 (E5M2) format and \(\pmb{B}\) uses the FP8 (E4M3) format; eight 8-bit values are packed into a long long register per lane for each operand.

Parameter

Type

Description

srcA

long long

Eight BF8 (E5M2) elements of \(\pmb{A}\) per lane.

srcB

long long

Eight FP8 (E4M3) elements of \(\pmb{B}\) per lane.

srcC

v4float

Accumulator input: 4 FP32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Must be 0.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Must be 0.

blgp

int

Must be 0.

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

__builtin_amdgcn_mfma_f32_16x16x32_bf8_bf8#

Signature and parameters for this builtin.

v4float __builtin_amdgcn_mfma_f32_16x16x32_bf8_bf8(
    long long srcA,
    long long srcB,
    v4float   srcC,
    int       cbsz,
    int       abid,
    int       blgp);

Computes one step of a \(16 \times 16\) FP32 outer-product accumulation. Each instruction processes \(K=32\) BF8 column elements of \(\pmb{A}\) and \(K=32\) BF8 row elements of \(\pmb{B}\), accumulating into a \(16 \times 16\) FP32 tile held across 4 accVGPRs per lane. This builtin operates on a single input block.

Both \(\pmb{A}\) and \(\pmb{B}\) use the BF8 (E5M2) format: eight 8-bit values packed into a long long register per lane.

Parameter

Type

Description

srcA

long long

Eight BF8 (E5M2) elements of \(\pmb{A}\) per lane.

srcB

long long

Eight BF8 (E5M2) elements of \(\pmb{B}\) per lane.

srcC

v4float

Accumulator input: 4 FP32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Must be 0.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Must be 0.

blgp

int

Must be 0.

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

FP64-accumulate builtins#

CDNA3 retains native double-precision (FP64) matrix accumulation from CDNA2. These builtins accept one FP64 element per lane for both \(\pmb{A}\) and \(\pmb{B}\) inputs and accumulate into FP64 output fragments held in v4double registers.

Note

On CDNA3, the blgp parameter of FP64 builtins is repurposed as a 3-bit source negation modifier rather than a lane-group pattern selector. See the Common parameters section for details.

FP64 matrix inputs#

The following builtins accept one FP64 element per lane for both \(\pmb{A}\) and \(\pmb{B}\) inputs and accumulate into FP64 output fragments.

__builtin_amdgcn_mfma_f64_16x16x4f64#

Signature and parameters for this builtin.

v4double __builtin_amdgcn_mfma_f64_16x16x4f64(
    double   srcA,
    double   srcB,
    v4double srcC,
    int      cbsz,
    int      abid,
    int      blgp);

Computes one step of a \(16 \times 16\) FP64 outer-product accumulation. Each instruction processes \(K=4\) FP64 column elements of \(\pmb{A}\) and \(K=4\) FP64 row elements of \(\pmb{B}\), accumulating into a \(16 \times 16\) FP64 tile held across 4 accVGPRs per lane (8 physical accVGPRs, since each FP64 value occupies two 32-bit registers). This builtin operates on a single input block.

Parameter

Type

Description

srcA

double

One FP64 element of \(\pmb{A}\) per lane.

srcB

double

One FP64 element of \(\pmb{B}\) per lane.

srcC

v4double

Accumulator input: 4 FP64 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Must be 0.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Must be 0.

blgp

int

\(\pmb{B}\)-matrix Lane Group Pattern modifier, repurposed as a negation modifier (3-bit field). Bit 0 negates srcA, bit 1 negates srcB, bit 2 negates srcC.

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

Note

Unlike most FP32, BF16, and INT8 MFMA instructions, FP64 MFMA instructions cannot co-execute with VALU instructions. To hide execution latency, schedule non-VALU instructions (loads, stores, or branches) around FP64 MFMA operations.

__builtin_amdgcn_mfma_f64_4x4x4f64#

Signature and parameters for this builtin.

double __builtin_amdgcn_mfma_f64_4x4x4f64(
    double srcA,
    double srcB,
    double srcC,
    int    cbsz,
    int    abid,
    int    blgp);

Computes one step of a \(4 \times 4\) FP64 outer-product accumulation. Each instruction processes \(K=4\) FP64 column elements of \(\pmb{A}\) and \(K=4\) FP64 row elements of \(\pmb{B}\), accumulating into a \(4 \times 4\) FP64 tile. Each lane holds one FP64 output element across four input blocks.

Parameter

Type

Description

srcA

double

One FP64 element of \(\pmb{A}\) per lane.

srcB

double

One FP64 element of \(\pmb{B}\) per lane.

srcC

double

Accumulator input: 1 FP64 element per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Must be 0.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Must be 0.

blgp

int

\(\pmb{B}\)-matrix Lane Group Pattern modifier, repurposed as a negation modifier (3-bit field). Bit 0 negates srcA, bit 1 negates srcB, bit 2 negates srcC.

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

Note

Unlike most FP32, BF16, and INT8 MFMA instructions, FP64 MFMA instructions cannot co-execute with VALU instructions. To hide execution latency, schedule non-VALU instructions (loads, stores, or branches) around FP64 MFMA operations.

INT32-accumulate builtins#

These builtins accumulate into signed 32-bit integer (INT32) output fragments.

INT8 matrix inputs#

The following builtins accept four signed 8-bit integer elements per lane packed into a single int register for both \(\pmb{A}\) and \(\pmb{B}\) inputs.

__builtin_amdgcn_mfma_i32_32x32x4i8#

Signature and parameters for this builtin.

v32int __builtin_amdgcn_mfma_i32_32x32x4i8(
    int    srcA,
    int    srcB,
    v32int srcC,
    int    cbsz,
    int    abid,
    int    blgp);

Computes one step of a \(32 \times 32\) INT32 outer-product accumulation. Each instruction processes \(K=4\) INT8 column elements of \(\pmb{A}\) and \(K=4\) INT8 row elements of \(\pmb{B}\), accumulating into a \(32 \times 32\) INT32 tile held across 32 accVGPRs per lane. This builtin operates on two input blocks.

Parameter

Type

Description

srcA

int

Four elements of \(\pmb{A}\) per lane.

srcB

int

Four elements of \(\pmb{B}\) per lane.

srcC

v32int

Accumulator input: 32 INT32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Legal range: 0..1

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Legal range: 0..1

blgp

int

\(\pmb{B}\)-matrix Lane Group Pattern modifier, see Common MFMA parameters.

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

__builtin_amdgcn_mfma_i32_16x16x4i8#

Signature and parameters for this builtin.

v16int __builtin_amdgcn_mfma_i32_16x16x4i8(
    int    srcA,
    int    srcB,
    v16int srcC,
    int    cbsz,
    int    abid,
    int    blgp);

Computes one step of a \(16 \times 16\) INT32 outer-product accumulation. Each instruction processes \(K=4\) INT8 column elements of \(\pmb{A}\) and \(K=4\) INT8 row elements of \(\pmb{B}\), accumulating into a \(16 \times 16\) INT32 tile held across 16 accVGPRs per lane. This builtin operates on four input blocks.

Parameter

Type

Description

srcA

int

Four elements of \(\pmb{A}\) per lane.

srcB

int

Four elements of \(\pmb{B}\) per lane.

srcC

v16int

Accumulator input: 16 INT32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Legal range: 0..2

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Legal range: 0..3

blgp

int

\(\pmb{B}\)-matrix Lane Group Pattern modifier, see Common MFMA parameters.

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

__builtin_amdgcn_mfma_i32_4x4x4i8#

Signature and parameters for this builtin.

v4int __builtin_amdgcn_mfma_i32_4x4x4i8(
    int   srcA,
    int   srcB,
    v4int srcC,
    int   cbsz,
    int   abid,
    int   blgp);

Computes one step of a \(4 \times 4\) INT32 outer-product accumulation. Each instruction processes \(K=4\) INT8 column elements of \(\pmb{A}\) and \(K=4\) INT8 row elements of \(\pmb{B}\), accumulating into a \(4 \times 4\) INT32 tile held across 4 accVGPRs per lane. This builtin operates on 16 input blocks.

Parameter

Type

Description

srcA

int

Four elements of \(\pmb{A}\) per lane.

srcB

int

Four elements of \(\pmb{B}\) per lane.

srcC

v4int

Accumulator input: 4 INT32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Legal range: 0..4

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Legal range: 0..15

blgp

int

\(\pmb{B}\)-matrix Lane Group Pattern modifier, see Common MFMA parameters.

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

The following wider-K variants accept eight signed 8-bit integer elements per lane packed into a long long register.

__builtin_amdgcn_mfma_i32_32x32x16_i8#

Signature and parameters for this builtin.

v16int __builtin_amdgcn_mfma_i32_32x32x16_i8(
    long long srcA,
    long long srcB,
    v16int    srcC,
    int       cbsz,
    int       abid,
    int       blgp);

Computes one step of a \(32 \times 32\) INT32 outer-product accumulation. Each instruction processes \(K=16\) INT8 column elements of \(\pmb{A}\) and \(K=16\) INT8 row elements of \(\pmb{B}\), accumulating into a \(32 \times 32\) INT32 tile held across 16 accVGPRs per lane. This builtin operates on a single input block.

Each lane provides eight signed 8-bit integer values packed into a long long register for both \(\pmb{A}\) and \(\pmb{B}\).

Parameter

Type

Description

srcA

long long

Eight INT8 elements of \(\pmb{A}\) per lane.

srcB

long long

Eight INT8 elements of \(\pmb{B}\) per lane.

srcC

v16int

Accumulator input: 16 INT32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Must be 0.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Must be 0.

blgp

int

Must be 0.

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

__builtin_amdgcn_mfma_i32_16x16x32_i8#

Signature and parameters for this builtin.

v4int __builtin_amdgcn_mfma_i32_16x16x32_i8(
    long long srcA,
    long long srcB,
    v4int     srcC,
    int       cbsz,
    int       abid,
    int       blgp);

Computes one step of a \(16 \times 16\) INT32 outer-product accumulation. Each instruction processes \(K=32\) INT8 column elements of \(\pmb{A}\) and \(K=32\) INT8 row elements of \(\pmb{B}\), accumulating into a \(16 \times 16\) INT32 tile held across 4 accVGPRs per lane. This builtin operates on a single input block.

Each lane provides eight signed 8-bit integer values packed into a long long register for both \(\pmb{A}\) and \(\pmb{B}\).

Parameter

Type

Description

srcA

long long

Eight INT8 elements of \(\pmb{A}\) per lane.

srcB

long long

Eight INT8 elements of \(\pmb{B}\) per lane.

srcC

v4int

Accumulator input: 4 INT32 elements per lane.

cbsz

int

Control Broadcast Size modifier, see Common MFMA parameters. Must be 0.

abid

int

\(\pmb{A}\)-matrix Broadcast Identifier, see Common MFMA parameters. Must be 0.

blgp

int

Must be 0.

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

Contents