RDNA4 dense WMMA builtins

Wave-Matrix Multiply-Accumulate (WMMA) builtins let you issue hardware dense matrix multiply-accumulate operations directly from HIP device code on RDNA4 GPUs (gfx1200, gfx1201). Each WMMA instruction multiplies a dense \(\pmb{A}\) fragment by a dense \(\pmb{B}\) fragment and accumulates the result into a \(\pmb{C}\) fragment, all within a single 32-wide wavefront.

RDNA4 supports the same FP16, BF16, INT8, and INT4 WMMA shapes as RDNA3 and adds FP8 (E4M3), BF8 (E5M2) input formats and a deeper-K INT4 variant (16x16x32). RDNA4 also provides sparse sparse WMMA variants that halve the \(\pmb{A}\) bandwidth using 2:4 structured sparsity.

CDNA (Instinct) GPUs provide a comparable dense operation through Matrix Fused Multiply-Accumulate (MFMA). The two differ in wavefront size (wave32 for WMMA, wave64 for MFMA) and accumulator storage.

Note

RDNA4 WMMA builtins use a _gfx12 suffix (for example, __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12) to distinguish them from the RDNA3 builtins, which lack this suffix. The two generations have different fragment sizes and accumulator layouts despite sharing the same tile dimensions.

Note

RDNA4 GPUs run all shader programs in wave32 mode by default. The _w32 suffix in each builtin name reflects this: all WMMA builtins on this page require wavefrontsize32.

Architecture availability

The builtins on this page target RDNA4 GPUs. To automatically enable them, pass the Low Level Virtual Machine (LLVM) target architecture flag at compile time:

amdclang++ --offload-arch=gfx1200 ...
amdclang++ --offload-arch=gfx1201 ...

Equivalent dense WMMA builtins for the previous generation are documented on the RDNA3 dense WMMA builtins page.

Naming convention

All dense WMMA builtins on this page follow the pattern:

__builtin_amdgcn_wmma_<out_type>_<M>x<N>x<K>_<in_type>[_<in_type_b>]_w32_gfx12

out_type

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

in_type

Input element type of \(\pmb{A}\) (and \(\pmb{B}\) when both share the same type): f16, bf16, iu8, iu4, fp8, or bf8. The iu prefix means the builtin accepts either signed or unsigned integers, controlled by the a_neg and b_neg parameters.

in_type_b (optional)

Input element type of \(\pmb{B}\) when it differs from \(\pmb{A}\). Used only for mixed FP8 and BF8 variants.

_w32

Wavefront size suffix. All RDNA4 WMMA builtins use wave32.

_gfx12

Architecture suffix distinguishing RDNA4 builtins from the RDNA3 variants.

Fragment layouts

All WMMA builtins on this page use a \(16 \times 16\) output tile computed by one wave32 wavefront. The 32 lanes split into two groups of 16; each group owns a contiguous block of 8 output rows. This is a lane-group-split layout. Unlike RDNA3’s matrix replication, RDNA4 splits the K dimension across lane groups: the two groups cover different K positions in the A and B fragments, interleaved in groups of 4 consecutive K values. The diagram and tables below show the mapping between matrix elements and lane or Vector General-Purpose Register (VGPR) positions for each operand.

Accumulator layout

Each lane holds 8 output elements across VGPRs 0–7. This layout is identical to the SWMMAC accumulator on the same hardware.

Given lane \(L\) and VGPR index \(g\):

\[\begin{split}i &= \lfloor \frac{L}{16} \rfloor \cdot 8 + g \\ j &= L \bmod 16\end{split}\]

Conversely, given output element \((i, j)\):

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

The row-to-lane mapping:

Rows	Lanes	VGPRs
0–7	0–15	0–7
8–15	16–31	0–7

Note

RDNA3 uses an interleaved row assignment (even rows in lanes 0–15, odd rows in lanes 16–31), while RDNA4 uses contiguous 8-row blocks. See Fragment layouts for the RDNA3 layout.

srcA and srcB (FP16 and BF16)

Each lane holds 8 input elements of \(\pmb{A}\) (v8half, 4 VGPRs × 2 FP16), covering one row of the \(16 \times 16\) A fragment; similarly, each lane holds 8 elements of \(\pmb{B}\) covering one column of the \(16 \times 16\) B fragment. For srcA, lane % 16 gives the matrix row; for srcB, lane % 16 gives the matrix column. The VGPR-to-K mapping is the same for both operands (transposed orientation).

Unlike RDNA3, where both lane groups carry identical copies of A and B (matrix replication), RDNA4 splits the K dimension across lane groups. The two groups cover different K positions in an interleaved pattern: groups of 4 consecutive K values alternate between lane groups. Lane group 0 (lanes 0–15) covers K {0–3, 8–11}; lane group 1 (lanes 16–31) covers K {4–7, 12–15}. Both groups cover all 16 rows (srcA) or columns (srcB).

The row and column-to-lane mapping:

Row (srcA) or Column (srcB)	Lanes	VGPRs
0	0, 16	0–3
1	1, 17	0–3
2	2, 18	0–3
3	3, 19	0–3
4	4, 20	0–3
5	5, 21	0–3
6	6, 22	0–3
7	7, 23	0–3
8	8, 24	0–3
9	9, 25	0–3
10	10, 26	0–3
11	11, 27	0–3
12	12, 28	0–3
13	13, 29	0–3
14	14, 30	0–3
15	15, 31	0–3

The 8 FP16 elements per lane are distributed across VGPRs:

Lane group	VGPR 0	VGPR 1	VGPR 2	VGPR 3
0 (lanes 0–15)	K {0, 1}	K {2, 3}	K {8, 9}	K {10, 11}
1 (lanes 16–31)	K {4, 5}	K {6, 7}	K {12, 13}	K {14, 15}

Using WMMA builtins as a compute policy

Using MFMA builtins as a compute policy explains the ComputePolicy pattern used to separate the multiply-accumulate logic from the rest of a kernel. A baseline policy is provided for RDNA4; a second, hardware-transpose policy that loads B via a transpose-load builtin instead of a scalar loop is also defined in the same example file and covered on that page.

The complete source file is available for download:

• matrix_multiply_rdna4_wmma.hip

Baseline policy

WmmaRdna4F16Policy uses scalar Local Data Share (LDS) loads for both A and B fragments. B is cooperatively transposed during the tile load (in the TilePolicy).

Each wavefront computes a single \(16 \times 16\) output tile. The 32 lanes form 2 groups of 16 (laneGroup = lane_id / 16); each group covers 8 consecutive K-positions, so each lane supplies 8 FP16 values for both A and B.

Policy constants

v_wmma_f32_16x16x16_f16 consumes 16 FP16 K positions per call, so k_step = 16. The wavefront holds the entire \(16 \times 16\) tile: thread_tile_m = thread_tile_n = 16 and effective_lanes = 32.

Accumulator layout

The builtin returns a v8float holding 8 VGPR values per lane. The store_c() pass maps (lane, VGPR index) back to \((i, j)\) coordinates using the RDNA4 lane-group-split layout.

struct WmmaRdna4F16Policy
{
    // -- ComputePolicy constants ----------------------------------------------
    // One wave32 wavefront owns a 16×16 output tile.
    // All 32 lanes hold unique output elements (8 elements per lane).
    static constexpr int thread_tile_m   = 16;
    static constexpr int thread_tile_n   = 16;
    static constexpr int frag_size_m     = 8;
    static constexpr int frag_size_n     = 8;
    // RDNA4 uses wave32: 32 lanes per wavefront (vs 64 on CDNA4).
    static constexpr int effective_lanes = 32;
    // wmma_f32_16x16x16_f16_w32 consumes 16 K-positions per call (k_step=16).
    // Each lane provides eight FP16 values (half8) for A and eight for B.
    // The 32 lanes split into 2 groups of 16; each group covers 8 K-positions.
    static constexpr int k_step          = 16;

    using elem_a = _Float16;
    using elem_b = _Float16;

    // -- Accumulator ----------------------------------------------------------
    // wmma_f8 holds 8 FP32 values per lane (8 output elements per lane in
    // the 16×16 tile, vs 4 per lane in CDNA4's MFMA with wave64).
    using wmma_f8 = float [[clang::ext_vector_type(8)]];

    struct Accumulator
    {
        wmma_f8 regs;
    };

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

    // -- thread_tile_offset ---------------------------------------------------
    // The entire 16×16 tile belongs to one wave32 wavefront.  Multiple
    // wavefronts in a block cover different 16×16 sub-tiles.
    //
    //   wavefront id within block: wid = tid / 32  (wave32, not 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 / 32; // wave32: divide by 32, not 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 eight FP16 values from LDS into a half8.
    //
    // wmma_f32_16x16x16_f16_w32 interprets the 32 lanes as 2 groups of 16
    // (indexed by laneGroup = lane_id / 16).  Each group covers 8 consecutive
    // K-positions: group g provides data for ki + 8*g through ki + 8*g + 7.
    // Within a group, lane offset (lane_id % 16) selects one of the 16
    // A-rows.
    //
    //   A-row index : tile_a_row + (lane_id mod 16)
    //   K-position e: ki + 8*(lane_id / 16) + e   (e = 0..7)
    //
    // Note: the LDS tile is indexed as tile_a[k_local][m_local] (K rows,
    // BlockTileM cols), so tile_a_ptr[k * k_tile_size + row] becomes
    // tile_a_ptr[row * k_tile_size + k] when stored as tile_a[M][K].
    // We use the M×K (row-major A) layout: tile_a[m_local][k_local].
    __device__ static void load_a(const _Float16* tile_a_ptr,
                                  int             tile_a_row,
                                  int             ki,
                                  int             k_tile_size,
                                  int             lane_id,
                                  elem_a        (&frag)[8])
    {
        const int row   = tile_a_row + (lane_id % 16);
        const int k_off = ki + 8 * (lane_id / 16);
        #pragma unroll
        for(int e = 0; e < 8; ++e)
            frag[e] = tile_a_ptr[row * k_tile_size + k_off + e];
    }

    // load_b: each lane reads eight FP16 values from transposed LDS tile.
    //
    // tile_b_T is stored transposed: tile_b_T[n_local][k_local].
    //   B-col index : tile_b_col + (lane_id mod 16)
    //   K-position e: ki + 8*(lane_id / 16) + e   (e = 0..7)
    __device__ static void load_b(const _Float16* tile_b_T_ptr,
                                  int             tile_b_col,
                                  int             ki,
                                  int             k_tile_size,
                                  int             lane_id,
                                  elem_b        (&frag)[8])
    {
        const int col   = tile_b_col + (lane_id % 16);
        const int k_off = ki + 8 * (lane_id / 16);
        #pragma unroll
        for(int e = 0; e < 8; ++e)
            frag[e] = tile_b_T_ptr[col * k_tile_size + k_off + e];
    }

    // -- mma -------------------------------------------------------------------
    // Issue one wmma_f32_16x16x16_f16_w32_gfx12 instruction.
    __device__ static void mma(Accumulator&       acc,
                               const elem_a     (&a_frag)[8],
                               const elem_b     (&b_frag)[8])
    {
#if defined(__gfx1200__) || defined(__gfx1201__)
        using half8 = __attribute__((ext_vector_type(8))) __fp16;
        const half8 a_vec = {a_frag[0], a_frag[1], a_frag[2], a_frag[3],
                             a_frag[4], a_frag[5], a_frag[6], a_frag[7]};
        const half8 b_vec = {b_frag[0], b_frag[1], b_frag[2], b_frag[3],
                             b_frag[4], b_frag[5], b_frag[6], b_frag[7]};
        acc.regs = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12(
            a_vec, b_vec, acc.regs);
#endif
    }

    // -- store_c ---------------------------------------------------------------
    // Scatter the 8 accumulator values to their global-memory positions.
    //
    // Inverse layout (lane L, accumulator index G in [0,7]):
    //   i = laneGroup * 8 + G    (laneGroup = L / 16)
    //   j = laneWrapped          (laneWrapped = L % 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 laneWrapped = lane_id % 16;
        const int laneGroup   = lane_id / 16;

        #pragma unroll
        for(int G = 0; G < 8; ++G)
        {
            const int i = laneGroup * 8 + G;
            const int j = laneWrapped;
            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

Plug this policy into the generic kernel alongside any TilePolicy whose block_tile_m and block_tile_n are multiples of 16 and whose k_tile_size is a multiple of k_step = 16. The launch code below also instantiates WmmaTilePolicyTrB, the hardware-transpose tile policy explained on RDNA4 WMMA transpose load builtins, and runs both back to back for comparison.

// TilePolicy for the baseline FP16 path (row-major B, scalar-transposed into LDS).
using WmmaTilePolicy       = SingleBufferTilePolicyF16<32, 32, 16>;
// TilePolicy for the global_load_tr path (row-major B, hardware-transposed into LDS).
using WmmaTilePolicyTrB    = SingleBufferTilePolicyF16TrB<32, 32, 16>;

// Launch parameters:
//
// RDNA4 (gfx12xx): 2×2 wave32 wavefronts of 32 lanes = 128 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): 8×8 thread tiles in a 32×32 block → 16 threads per
//   block, padded to one full wave32 of 32.
//
// 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 32×32 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_gfx12)
    {
        // Baseline path: row-major B, scalar transpose into LDS.
        auto launch = [&]()
        {
            matrix_multiply_generic<WmmaTilePolicy,
                                    WmmaRdna4F16Policy,
                                    _Float16>
                <<<grid, dim3(128)>>>(d_A, d_B, d_C, M, N, K);
            HIP_CHECK(hipGetLastError());
        };
        ms           = time_kernel_ms(launch, WARMUP_RUNS, TIMING_RUNS);
        policy_label = "GemmKernel<WmmaTilePolicy, WmmaRdna4F16Policy> [gfx12]";

        HIP_CHECK(hipMemcpy(h_C.data(), d_C, sizeof(float) * M * N,
                            hipMemcpyDeviceToHost));
        const bool ok_baseline = verify_result(h_C, h_A, M, N, K);
        print_metrics(policy_label, ok_baseline, ms, sizeof(_Float16), M, N, K);

        // global_load_tr path: row-major B, hardware transpose into LDS.
        HIP_CHECK(hipMemset(d_C, 0, sizeof(float) * M * N));
        auto launch_tr = [&]()
        {
            matrix_multiply_generic<WmmaTilePolicyTrB,
                                    WmmaRdna4F16TrPolicy,
                                    _Float16>
                <<<grid, dim3(128)>>>(d_A, d_B, d_C, M, N, K);
            HIP_CHECK(hipGetLastError());
        };
        const float ms_tr = time_kernel_ms(launch_tr, WARMUP_RUNS, TIMING_RUNS);
        HIP_CHECK(hipMemcpy(h_C.data(), d_C, sizeof(float) * M * N,
                            hipMemcpyDeviceToHost));
        const bool ok_tr = verify_result(h_C, h_A, M, N, K);
        print_metrics(
            "GemmKernel<WmmaTilePolicyTrB, WmmaRdna4F16TrPolicy> [gfx12]",
            ok_tr, ms_tr, sizeof(_Float16), M, N, K);
    }
    else
    {
        // Scalar FP32 fallback: upload A and B as FP32.
        std::vector<float> h_A_f32(M * K);
        for(int i = 0; i < M * K; ++i)
            h_A_f32[i] = static_cast<float>(h_A[i]);

        std::vector<float> h_B_f32(K * N, 0.0f);
        for(int i = 0; i < std::min(K, N); ++i)
            h_B_f32[i * N + i] = 1.0f;

        float* d_A_f32;
        float* d_B_f32;
        HIP_CHECK(hipMalloc(&d_A_f32, sizeof(float) * M * K));
        HIP_CHECK(hipMalloc(&d_B_f32, sizeof(float) * K * N));
        HIP_CHECK(hipMemcpy(d_A_f32, h_A_f32.data(), sizeof(float) * M * K,
                            hipMemcpyHostToDevice));
        HIP_CHECK(hipMemcpy(d_B_f32, h_B_f32.data(), sizeof(float) * K * N,
                            hipMemcpyHostToDevice));

        auto launch = [&]()
        {
            matrix_multiply_generic<SingleBufferTilePolicyF<32, 32, 8>,
                                    ScalarFMAPolicy<8, 8, 4>,
                                    float>
                <<<grid, dim3(32)>>>(d_A_f32, d_B_f32, d_C, M, N, K);
            HIP_CHECK(hipGetLastError());
        };
        ms           = time_kernel_ms(launch, WARMUP_RUNS, TIMING_RUNS);
        policy_label = "GemmKernel<SingleBufferTilePolicyF, ScalarFMAPolicy> [fallback]";

        HIP_CHECK(hipMemcpy(h_C.data(), d_C, sizeof(float) * M * N,
                            hipMemcpyDeviceToHost));
        const bool ok = verify_result(h_C, h_A, M, N, K);
        print_metrics(policy_label, ok, ms, sizeof(_Float16), M, N, K);

        HIP_CHECK(hipFree(d_A_f32));
        HIP_CHECK(hipFree(d_B_f32));
    }

Compile and run:

amdclang++ -O3 -std=c++17 --offload-arch=gfx1200 \
    matrix_multiply_rdna4_wmma.hip -o mm_rdna4_wmma
./mm_rdna4_wmma

amdclang++ -O3 -std=c++17 --offload-arch=gfx1201 \
    matrix_multiply_rdna4_wmma.hip -o mm_rdna4_wmma
./mm_rdna4_wmma

Note

WmmaRdna4F16Policy and WmmaRdna4F16TrPolicy require an RDNA4 GPU (gfx1200 or gfx1201). The #if defined(__gfx1200__) || defined(__gfx1201__) guard in the example file falls back to ScalarFMAPolicy on other targets, so the file compiles without modification.

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 v8half   = _Float16 [[clang::ext_vector_type(8)]];
using v8float  = float [[clang::ext_vector_type(8)]];
using v8short  = short [[clang::ext_vector_type(8)]];
using v8int    = int [[clang::ext_vector_type(8)]];
using v4int    = int [[clang::ext_vector_type(4)]];
using v2int    = int [[clang::ext_vector_type(2)]];

Common parameters

The a_neg, b_neg, and clamp parameters appear only on integer WMMA variants.

Parameter	Type	Description
a_neg	bool (compile-time constant)	Integer variants only. When true, the \(\pmb{A}\) elements are treated as signed integers; when false, as unsigned.
b_neg	bool (compile-time constant)	Integer variants only. Same as a_neg but for \(\pmb{B}\).
clamp	bool (compile-time constant)	Integer variants only. When true, the INT32 accumulator output is clamped to the representable range of the input type on overflow.

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

Builtin	Ops	Cycle count
__builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12	8192	16
__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12	8192	16
__builtin_amdgcn_wmma_f16_16x16x16_f16_w32_gfx12	8192	16
__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32_gfx12	8192	16
__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12	8192	8
__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32_gfx12	8192	8
__builtin_amdgcn_wmma_i32_16x16x32_iu4_w32_gfx12	16384	8
__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w32_gfx12	8192	8
__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w32_gfx12	8192	8
__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w32_gfx12	8192	8
__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w32_gfx12	8192	8

Builtin reference

The following sections list every dense WMMA builtin available on RDNA4, grouped by accumulator type.

FP32-accumulate builtins

These builtins accumulate into FP32 and accept FP16, BF16, FP8, or BF8 matrix inputs.

FP16 inputs

The following builtins use FP16 matrix inputs.

__builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12

Signature and parameters for this builtin.

v8float __builtin_amdgcn_wmma_f32_16x16x16_f16_w32_gfx12(
    v8half  srcA,
    v8half  srcB,
    v8float srcC);

Computes one step of a dense \(16 \times 16\) FP32 accumulation with FP16 inputs. Each lane supplies 8 FP16 values for both \(\pmb{A}\) and \(\pmb{B}\). The 32 lanes form two groups of 16; each group covers 8 of the 16 K-positions.

Parameter	Type	Description
srcA	v8half	Eight FP16 elements of \(\pmb{A}\) per lane.
srcB	v8half	Eight FP16 elements of \(\pmb{B}\) per lane.
srcC	v8float	Accumulator input: eight FP32 elements per lane.

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

BF16 inputs

The following builtins use BF16 matrix inputs.

__builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12

Signature and parameters for this builtin.

v8float __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32_gfx12(
    v8short srcA,
    v8short srcB,
    v8float srcC);

Computes one step of a dense \(16 \times 16\) FP32 accumulation with BF16 inputs. BF16 values are passed as short (16-bit storage). Otherwise identical in structure to the FP16 variant.

Parameter	Type	Description
srcA	v8short	Eight BF16 elements of \(\pmb{A}\) per lane (BF16 stored as short).
srcB	v8short	Eight BF16 elements of \(\pmb{B}\) per lane.
srcC	v8float	Accumulator input: eight FP32 elements per lane.

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

FP8 and BF8 inputs

The following builtins use FP8 and BF8 matrix inputs.

__builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w32_gfx12

Signature and parameters for this builtin.

v8float __builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w32_gfx12(
    v2int   srcA,
    v2int   srcB,
    v8float srcC);

Computes one step of a dense \(16 \times 16\) FP32 accumulation with FP8 (E4M3) inputs for both \(\pmb{A}\) and \(\pmb{B}\). Eight FP8 values are packed into two 32-bit registers per lane (four 8-bit values per register).

Parameter	Type	Description
srcA	v2int	Eight FP8 elements of \(\pmb{A}\) per lane, packed into two 32-bit registers (four 8-bit values per register).
srcB	v2int	Eight FP8 elements of \(\pmb{B}\) per lane, packed into two 32-bit registers.
srcC	v8float	Accumulator input: eight FP32 elements per lane.

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

__builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w32_gfx12

Signature and parameters for this builtin.

v8float __builtin_amdgcn_wmma_f32_16x16x16_fp8_bf8_w32_gfx12(
    v2int   srcA,
    v2int   srcB,
    v8float srcC);

Computes one step of a dense \(16 \times 16\) FP32 accumulation with mixed 8-bit floating-point inputs: FP8 (E4M3) for \(\pmb{A}\) and BF8 (E5M2) for \(\pmb{B}\). Eight elements per lane are packed into two 32-bit registers.

Parameter	Type	Description
srcA	v2int	Eight FP8 elements of \(\pmb{A}\) per lane, packed into two 32-bit registers (four 8-bit values per register).
srcB	v2int	Eight BF8 elements of \(\pmb{B}\) per lane, packed into two 32-bit registers.
srcC	v8float	Accumulator input: eight FP32 elements per lane.

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

__builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w32_gfx12

Signature and parameters for this builtin.

v8float __builtin_amdgcn_wmma_f32_16x16x16_bf8_fp8_w32_gfx12(
    v2int   srcA,
    v2int   srcB,
    v8float srcC);

Computes one step of a dense \(16 \times 16\) FP32 accumulation with mixed 8-bit floating-point inputs: BF8 (E5M2) for \(\pmb{A}\) and FP8 (E4M3) for \(\pmb{B}\). Eight elements per lane are packed into two 32-bit registers.

Parameter	Type	Description
srcA	v2int	Eight BF8 elements of \(\pmb{A}\) per lane, packed into two 32-bit registers (four 8-bit values per register).
srcB	v2int	Eight FP8 elements of \(\pmb{B}\) per lane, packed into two 32-bit registers.
srcC	v8float	Accumulator input: eight FP32 elements per lane.

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

__builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w32_gfx12

Signature and parameters for this builtin.

v8float __builtin_amdgcn_wmma_f32_16x16x16_bf8_bf8_w32_gfx12(
    v2int   srcA,
    v2int   srcB,
    v8float srcC);

Computes one step of a dense \(16 \times 16\) FP32 accumulation with BF8 (E5M2) inputs for both \(\pmb{A}\) and \(\pmb{B}\). Eight BF8 values are packed into two 32-bit registers per lane (four 8-bit values per register).

Parameter	Type	Description
srcA	v2int	Eight BF8 elements of \(\pmb{A}\) per lane, packed into two 32-bit registers (four 8-bit values per register).
srcB	v2int	Eight BF8 elements of \(\pmb{B}\) per lane, packed into two 32-bit registers.
srcC	v8float	Accumulator input: eight FP32 elements per lane.

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

FP16-accumulate builtins

This builtin accumulates into FP16 with FP16 inputs.

FP16 inputs

The following builtin uses FP16 matrix inputs.

__builtin_amdgcn_wmma_f16_16x16x16_f16_w32_gfx12

Signature and parameters for this builtin.

v8half __builtin_amdgcn_wmma_f16_16x16x16_f16_w32_gfx12(
    v8half srcA,
    v8half srcB,
    v8half srcC);

Computes one step of a dense \(16 \times 16\) FP16 accumulation. Both inputs and the accumulator are FP16.

Parameter	Type	Description
srcA	v8half	Eight FP16 elements of \(\pmb{A}\) per lane.
srcB	v8half	Eight FP16 elements of \(\pmb{B}\) per lane.
srcC	v8half	Accumulator input: eight FP16 elements per lane.

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

BF16-accumulate builtins

This builtin accumulates into BF16 with BF16 inputs.

BF16 inputs

The following builtin uses BF16 matrix inputs.

__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32_gfx12

Signature and parameters for this builtin.

v8short __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32_gfx12(
    v8short srcA,
    v8short srcB,
    v8short srcC);

Computes one step of a dense \(16 \times 16\) BF16 accumulation. Both inputs and the accumulator are BF16 (stored as short).

Parameter	Type	Description
srcA	v8short	Eight BF16 elements of \(\pmb{A}\) per lane.
srcB	v8short	Eight BF16 elements of \(\pmb{B}\) per lane.
srcC	v8short	Accumulator input: eight BF16 elements per lane (stored as short).

Returns v8short – updated accumulator (BF16 stored as short) (\(\text{srcA} \times \text{srcB} + \text{srcC}\)).

INT32-accumulate builtins

Integer WMMA builtins accept either signed or unsigned 8-bit or 4-bit integer inputs, controlled by the a_neg and b_neg compile-time constants.

INT8 and UINT8 inputs (16x16x16)

The following builtins use INT8 and UINT8 matrix inputs.

__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12

Signature and parameters for this builtin.

v8int __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32_gfx12(
    bool  a_neg,
    v2int srcA,
    bool  b_neg,
    v2int srcB,
    v8int srcC,
    bool  clamp);

Computes one step of a dense \(16 \times 16\) INT32 accumulation with 8-bit integer inputs. Sixteen INT8 elements of \(\pmb{A}\) are packed into two int registers per lane (four 8-bit values per register); likewise for \(\pmb{B}\).

Parameter	Type	Description
a_neg	bool	true for signed INT8, false for unsigned UINT8. Compile-time constant, see Common parameters.
srcA	v2int	Sixteen 8-bit elements of \(\pmb{A}\) per lane, packed into two 32-bit registers.
b_neg	bool	true for signed INT8, false for unsigned UINT8 in \(\pmb{B}\). Compile-time constant, see Common parameters.
srcB	v2int	Sixteen 8-bit elements of \(\pmb{B}\) per lane, packed into two 32-bit registers.
srcC	v8int	Accumulator input: eight INT32 elements per lane.
clamp	bool	Clamp output to input type range on overflow. Compile-time constant, see Common parameters.

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

INT4 and UINT4 inputs (16x16x16)

The following builtins use INT4 and UINT4 matrix inputs.

__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32_gfx12

Signature and parameters for this builtin.

v8int __builtin_amdgcn_wmma_i32_16x16x16_iu4_w32_gfx12(
    bool  a_neg,
    int   srcA,
    bool  b_neg,
    int   srcB,
    v8int srcC,
    bool  clamp);

Computes one step of a dense \(16 \times 16\) INT32 accumulation with 4-bit integer inputs. Sixteen INT4 elements of \(\pmb{A}\) fit into a single int per lane; likewise for \(\pmb{B}\).

Parameter	Type	Description
a_neg	bool	true for signed INT4, false for unsigned UINT4 in \(\pmb{A}\). Compile-time constant, see Common parameters.
srcA	int	Sixteen 4-bit elements of \(\pmb{A}\) per lane, packed into one 32-bit register.
b_neg	bool	true for signed INT4, false for unsigned UINT4 in \(\pmb{B}\). Compile-time constant, see Common parameters.
srcB	int	Sixteen 4-bit elements of \(\pmb{B}\) per lane, packed into one 32-bit register.
srcC	v8int	Accumulator input: eight INT32 elements per lane.
clamp	bool	Clamp output to input type range on overflow. Compile-time constant, see Common parameters.

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

INT4 and UINT4 inputs (16x16x32)

The following builtins use INT4 and UINT4 matrix inputs.

__builtin_amdgcn_wmma_i32_16x16x32_iu4_w32_gfx12

Signature and parameters for this builtin.

v8int __builtin_amdgcn_wmma_i32_16x16x32_iu4_w32_gfx12(
    bool  a_neg,
    v2int srcA,
    bool  b_neg,
    v2int srcB,
    v8int srcC,
    bool  clamp);

Computes one step of a dense \(16 \times 16\) INT32 accumulation with 4-bit integer inputs over a deeper K=32 strip. Thirty-two INT4 elements of \(\pmb{A}\) are packed into two int registers per lane; likewise for \(\pmb{B}\).

Parameter	Type	Description
a_neg	bool	true for signed INT4, false for unsigned UINT4 in \(\pmb{A}\). Compile-time constant, see Common parameters.
srcA	v2int	Thirty-two 4-bit elements of \(\pmb{A}\) per lane, packed into two 32-bit registers (sixteen 4-bit values per register).
b_neg	bool	true for signed INT4, false for unsigned UINT4 in \(\pmb{B}\). Compile-time constant, see Common parameters.
srcB	v2int	Thirty-two 4-bit elements of \(\pmb{B}\) per lane, packed into two 32-bit registers.
srcC	v8int	Accumulator input: eight INT32 elements per lane.
clamp	bool	Clamp output to input type range on overflow. Compile-time constant, see Common parameters.

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