CDNA dense MFMA builtins

CDNA dense MFMA builtins#

Matrix Fused Multiply-Add (MFMA) builtins let you issue hardware matrix multiply-accumulate operations directly from HIP device code on first-generation CDNA GPUs (gfx908, MI100). 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.

Architecture availability#

The builtins on this page target CDNA (gfx908, MI100) exclusively. Equivalent builtins for later 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>
out_type

Accumulator element type (f32 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, f16, bf16, or i8).

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 \(\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, INT8 \(\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 CDNA GPUs, the \(\pmb{C}\) and \(\pmb{D}\) matrix operands must reside in accumulation VGPRs (accVGPRs). Unlike \(\pmb{A}\) and \(\pmb{B}\), they cannot use standard architecture VGPRs (archVGPRs). In HIP device code the transfer to archVGPRs happens automatically when assigning the builtin return value to a local variable of the appropriate vector type.

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 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 \frac{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. Applies to all \(32 \times 32\) FP32 and INT32 MFMA builtins. See the formulas and row-to-lane table below.#

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. Applies to all \(16 \times 16\) FP32 and INT32 MFMA builtins. 1-block variants use only block 0. See the formulas and row-to-lane table below.#

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. Applies to all \(4 \times 4\) FP32 and INT32 MFMA builtins. See the formulas below.#

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}\]

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

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.

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 MfmaCdnaPolicy using v_mfma_f32_32x32x1f32 – a \(32 \times 32\) FP32 builtin available on all CDNA generations. Each wavefront computes a single \(32 \times 32\) output tile per mma() call; two blocks are active, giving 32 accVGPRs per lane.

The complete source file is available for download:

Policy constants

v_mfma_f32_32x32x1f32 takes one FP32 scalar from A and one from B per call (K=1), so k_step = 1. The entire \(32 \times 32\) tile is owned by a single wavefront, so thread_tile_m = thread_tile_n = 32 and effective_lanes = 64.

Accumulator layout

The builtin returns a v32float holding 32 accVGPR values – 16 for block 0 and 16 for block 1. 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 K=1, each load_a call reads a single scalar: the element at row tile_a_row + (lane mod 32) (since 32 \(\pmb{A}\)-rows map to the lower 32 lanes). In practice the kernel inner loop increments ki in steps of k_step = 1, so each load_a/load_b call loads one element.

// MfmaCdnaPolicy is defined for all compilation passes so that the host and
// device agree on the template instantiation.  The intrinsic body inside mma()
// is guarded by __gfx908__, which is defined only during the device pass.
struct MfmaCdnaPolicy
{
    // -- ComputePolicy constants ----------------------------------------------
    // One wavefront owns a 32×32 output tile.
    // All 64 lanes contribute unique output elements (no lane mirroring).
    static constexpr int thread_tile_m  = 32;
    static constexpr int thread_tile_n  = 32;
    static constexpr int effective_lanes = 64;
    // v_mfma_f32_32x32x1f32 consumes one K-index per call (K=1 per instruction).
    static constexpr int k_step          = 1;
    // MFMA takes one scalar per lane.
    static constexpr int frag_size_m     = 1;
    static constexpr int frag_size_n     = 1;

    using elem_a = float;
    using elem_b = float;

    // -- Accumulator ----------------------------------------------------------
    // v32float holds 32 accVGPRs per lane:
    //   elements [0,15]  → block 0 accVGPRs 0–15
    //   elements [16,31] → block 1 accVGPRs 16–31
    using v32float = float [[clang::ext_vector_type(32)]];

    struct Accumulator
    {
        v32float regs;
    };

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

    // -- thread_tile_offset ---------------------------------------------------
    // The entire 32×32 tile belongs to one wavefront; multiple wavefronts
    // in a block cover different 32×32 sub-tiles.  Each wavefront's top-left
    // corner is determined by its wavefront index within the block.
    //
    //   wavefront id within block: wid = tid / 64
    //   wid_x = wid % (block_tile_n / 32)
    //   wid_y = wid / (block_tile_n / 32)
    //   *thread_row = wid_y * 32
    //   *thread_col = wid_x * 32
    __device__ static void thread_tile_offset(int  tid,
                                               int  /*lane_id*/,
                                               int* thread_row,
                                               int* thread_col)
    {
        // block_tile_n / 32 = number of wavefronts along the N dimension.
        // For a 64×64 block tile this is 64/32 = 2.
        constexpr int waves_n = 2; // block_tile_n / thread_tile_n
        const int wid   = tid / 64;
        const int wid_x = wid % waves_n;
        const int wid_y = wid / waves_n;
        *thread_row = wid_y * 32;
        *thread_col = wid_x * 32;
    }

    // -- Fragment loads --------------------------------------------------------
    // load_a: each lane reads its own A element from LDS.
    //
    // The 32×32 MFMA A-operand distributes 32 rows across the lower 32 lanes:
    //   lane L in [0,31]:  provides A-row (tile_a_row + L)
    //   lane L in [32,63]: mirrors the same 32 rows for the second block
    // Indexing: tile_a[tile_a_row + (lane_id mod 32)][ki].
    __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)[1])
    {
        const int row = tile_a_row + (lane_id % 32);
        frag[0] = tile_a_ptr[row * k_tile_size + ki];
    }

    // load_b: each lane reads its own B element from LDS.
    //
    // The 32×32 MFMA B-operand distributes 32 cols across the lower 32 lanes:
    //   lane L in [0,31]:  provides B-col (tile_b_col + L)
    //   lane L in [32,63]: mirrors the same 32 cols for the second block
    // tile_b_T is column-major: tile_b_T[(tile_b_col + (lane_id mod 32))][ki].
    __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)[1])
    {
        const int col = tile_b_col + (lane_id % 32);
        frag[0] = tile_b_T_ptr[col * k_tile_size + ki];
    }

    // -- mma -------------------------------------------------------------------
    // Issue one v_mfma_f32_32x32x1f32 instruction.
    // cbsz=0: use both blocks (no A-broadcast).
    // abid=0: select block 0 as A source.
    // blgp=0: no B lane-group permutation.
    __device__ static void mma(Accumulator&       acc,
                                const elem_a     (&a_frag)[1],
                                const elem_b     (&b_frag)[1])
    {
#if defined(__gfx908__)
        acc.regs = __builtin_amdgcn_mfma_f32_32x32x1f32(
            a_frag[0], b_frag[0], acc.regs,
            /*cbsz=*/0, /*abid=*/0, /*blgp=*/0);
#endif
    }

    // -- store_c ---------------------------------------------------------------
    // Scatter the 32 accVGPR values to their global-memory positions.
    //
    // Inverse layout (lane L, accVGPR index G in [0,31]):
    //   g_local = G mod 16        -- accVGPR within its block
    //   i = 8*(g_local/4) + 4*(L/32) + (g_local mod 4)
    //   j = L mod 32
    __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 % 32;

        #pragma unroll
        for(int G = 0; G < 32; ++G)
        {
            const int g_local = G % 16;
            const int i = 8 * (g_local / 4)
                        + 4 * (L / 32)
                        + (g_local % 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 MfmaCdnaPolicy in hand, plug it into the generic kernel alongside any TilePolicy whose block_tile_m and block_tile_n are multiples of 32 and whose k_tile_size is a multiple of k_step = 1. The policy aliases and launch configuration from the example file are:

// TilePolicy: 64x64 block tile, 16-element K-strip, single-buffered.
using MfmaTilePolicy = SingleBufferTilePolicy<64, 64, 16>;

// Block tile dimensions shared by both dispatch paths.
constexpr int BLOCK_TILE_M = 64;
constexpr int BLOCK_TILE_N = 64;

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

    float       ms;
    bool        ok;
    const char* policy_label;

    if(is_gfx908)
    {
        // CDNA (MI100): 2×2 wavefronts of 64 lanes = 256 threads per block.
        auto launch = [&]()
        {
            matrix_multiply_generic<MfmaTilePolicy, MfmaCdnaPolicy>
                <<<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, MfmaCdnaPolicy> [gfx908]";
    }
    else
    {
        // Other targets: scalar fallback, 8×8 thread tiles = 64 threads per block.
        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=gfx908 \
    matrix_multiply_cdna_mfma.hip -o mm_cdna_mfma
./mm_cdna_mfma

Note

MfmaCdnaPolicy requires a CDNA GPU (gfx908). Compile with --offload-arch=gfx908 to select the correct architecture. On other targets the #if defined(__gfx908__) 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

56

__builtin_amdgcn_mfma_f32_16x16x1f32

2048

32

24

__builtin_amdgcn_mfma_f32_4x4x1f32

512

8

0

__builtin_amdgcn_mfma_f32_32x32x2f32

4096

64

56

__builtin_amdgcn_mfma_f32_16x16x4f32

2048

32

24

__builtin_amdgcn_mfma_f32_32x32x4f16

16384

64

56

__builtin_amdgcn_mfma_f32_16x16x4f16

8192

32

24

__builtin_amdgcn_mfma_f32_4x4x4f16

2048

8

0

__builtin_amdgcn_mfma_f32_32x32x8f16

16384

64

56

__builtin_amdgcn_mfma_f32_16x16x16f16

8192

32

24

__builtin_amdgcn_mfma_f32_32x32x2bf16

8192

64

56

__builtin_amdgcn_mfma_f32_16x16x2bf16

4096

32

24

__builtin_amdgcn_mfma_f32_4x4x2bf16

1024

8

0

__builtin_amdgcn_mfma_f32_32x32x4bf16

8192

64

56

__builtin_amdgcn_mfma_f32_16x16x8bf16

4096

32

24

__builtin_amdgcn_mfma_i32_32x32x4i8

16384

64

56

__builtin_amdgcn_mfma_i32_16x16x4i8

8192

32

24

__builtin_amdgcn_mfma_i32_4x4x4i8

2048

8

0

__builtin_amdgcn_mfma_i32_32x32x8i8

16384

64

56

__builtin_amdgcn_mfma_i32_16x16x16i8

8192

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.

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#

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

__builtin_amdgcn_mfma_f32_32x32x2bf16#

Signature and parameters for this builtin.

v32float __builtin_amdgcn_mfma_f32_32x32x2bf16(
    v2bfloat srcA,
    v2bfloat 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=2\) BF16 column elements of \(\pmb{A}\) and \(K=2\) 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

v2bfloat

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

srcB

v2bfloat

Two 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_16x16x2bf16#

Signature and parameters for this builtin.

v16float __builtin_amdgcn_mfma_f32_16x16x2bf16(
    v2bfloat srcA,
    v2bfloat 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=2\) BF16 column elements of \(\pmb{A}\) and \(K=2\) 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

v2bfloat

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

srcB

v2bfloat

Two 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_4x4x2bf16#

Signature and parameters for this builtin.

v4float __builtin_amdgcn_mfma_f32_4x4x2bf16(
    v2bfloat srcA,
    v2bfloat 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=2\) BF16 column elements of \(\pmb{A}\) and \(K=2\) 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

v2bfloat

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

srcB

v2bfloat

Two 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_32x32x4bf16#

Signature and parameters for this builtin.

v16float __builtin_amdgcn_mfma_f32_32x32x4bf16(
    v2bfloat srcA,
    v2bfloat 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\) 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 16 accVGPRs per lane. This builtin operates on a single input block.

Parameter

Type

Description

srcA

v2bfloat

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

srcB

v2bfloat

Two 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_16x16x8bf16#

Signature and parameters for this builtin.

v4float __builtin_amdgcn_mfma_f32_16x16x8bf16(
    v2bfloat srcA,
    v2bfloat 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\) BF16 elements of \(\pmb{A}\) and \(K=8\) BF16 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

v2bfloat

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

srcB

v2bfloat

Two 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}\)).

INT32-accumulate builtins#

These builtins accept four signed 8-bit integer elements per lane packed into a single int register 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_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}\)).

__builtin_amdgcn_mfma_i32_32x32x8i8#

Signature and parameters for this builtin.

v16int __builtin_amdgcn_mfma_i32_32x32x8i8(
    int    srcA,
    int    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=8\) INT8 column elements of \(\pmb{A}\) and \(K=8\) 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.

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. 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 v16int – updated accumulator (\(\text{srcA} \times \text{srcB} + \text{srcC}\)).

__builtin_amdgcn_mfma_i32_16x16x16i8#

Signature and parameters for this builtin.

v4int __builtin_amdgcn_mfma_i32_16x16x16i8(
    int   srcA,
    int   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=16\) INT8 column elements of \(\pmb{A}\) and \(K=16\) 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.

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. 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 v4int – updated accumulator (\(\text{srcA} \times \text{srcB} + \text{srcC}\)).

Contents