RDNA3 dense WMMA builtins#
Wave-Matrix Multiply-Accumulate (WMMA) builtins let you issue hardware
dense matrix multiply-accumulate operations directly from HIP device code on
RDNA3 and RDNA3.5 GPUs (gfx1100–gfx1103, gfx1150–gfx1153). 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.
WMMA is the dense variant of the RDNA wave-matrix instruction family. RDNA4 extends the family with sparse WMMA variants documented on the RDNA4 sparse WMMA builtins page. 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), accumulator storage (ordinary Vector General-Purpose Registers (VGPRs) for WMMA, dedicated accVGPRs for MFMA on CDNA and CDNA2), and tile shapes.
Note
RDNA3 and RDNA3.5 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 RDNA3 and RDNA3.5 GPUs. To automatically enable them, pass the Low Level Virtual Machine (LLVM) target architecture flag at compile time:
amdclang++ --offload-arch=gfx1100 ...
amdclang++ --offload-arch=gfx1150 ...
Equivalent dense WMMA builtins for RDNA4 are documented on the RDNA4 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
out_typeAccumulator element type (
f32,f16,bf16, ori32).M,N,KTile 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_typeInput element type of \(\pmb{A}\) (and \(\pmb{B}\) when both share the same type):
f16,bf16,iu8, oriu4. Theiuprefix means the builtin accepts either signed or unsigned integers, controlled by thea_negandb_negparameters._w32Wavefront size suffix. All RDNA3 WMMA builtins use wave32.
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. Unlike RDNA4’s contiguous-block layout, RDNA3 interleaves the groups across rows: even-numbered rows are held by lanes 0–15, odd-numbered rows by lanes 16–31. This is the matrix replication pattern: lanes 0–15 and lanes 16–31 carry identical copies of the A and B input fragments, and the hardware fuses their contributions into alternating output rows. The diagrams below show the mapping between matrix elements and lane or VGPR positions for each operand.
Accumulator layout#
Each lane holds 8 output elements across VGPRs 0–7.
Given lane \(L\) and VGPR index \(g\):
Conversely, given output element \((i, j)\):
The row-to-lane mapping:
Rows |
Lanes |
VGPRs |
|---|---|---|
0 |
0–15 |
0 |
1 |
16–31 |
0 |
2 |
0–15 |
1 |
3 |
16–31 |
1 |
4 |
0–15 |
2 |
5 |
16–31 |
2 |
6 |
0–15 |
3 |
7 |
16–31 |
3 |
8 |
0–15 |
4 |
9 |
16–31 |
4 |
10 |
0–15 |
5 |
11 |
16–31 |
5 |
12 |
0–15 |
6 |
13 |
16–31 |
6 |
14 |
0–15 |
7 |
15 |
16–31 |
7 |
srcA and srcB (FP16 and BF16)#
Each lane holds 16 input elements of \(\pmb{A}\) (v16half for FP16
or v16short for BF16; 8 VGPRs × 2 elements), covering one row of the A
fragment; similarly, each lane holds 16 elements of \(\pmb{B}\) covering
one column of the B fragment.
Due to matrix replication, lanes 0–15 and 16–31 must load identical
copies of every row and column. Only lane % 16 determines which row
or column a lane covers – the upper lane-group bit is ignored. This means
your load_a and load_b code must supply the same data to both halves
of the wavefront.
Lane \(L\) covers:
srcA row \(= L \bmod 16\)
srcB column \(= L \bmod 16\)
The row and column-to-lane mapping:
Row (srcA) or Column (srcB) |
Lanes |
VGPRs |
|---|---|---|
0 |
0, 16 |
0–7 |
1 |
1, 17 |
0–7 |
2 |
2, 18 |
0–7 |
3 |
3, 19 |
0–7 |
4 |
4, 20 |
0–7 |
5 |
5, 21 |
0–7 |
6 |
6, 22 |
0–7 |
7 |
7, 23 |
0–7 |
8 |
8, 24 |
0–7 |
9 |
9, 25 |
0–7 |
10 |
10, 26 |
0–7 |
11 |
11, 27 |
0–7 |
12 |
12, 28 |
0–7 |
13 |
13, 29 |
0–7 |
14 |
14, 30 |
0–7 |
15 |
15, 31 |
0–7 |
The 16 FP16 elements are packed contiguously across 8 VGPRs (2 FP16 values per VGPR):
VGPR |
K positions |
|---|---|
0 |
K {0, 1} |
1 |
K {2, 3} |
2 |
K {4, 5} |
3 |
K {6, 7} |
4 |
K {8, 9} |
5 |
K {10, 11} |
6 |
K {12, 13} |
7 |
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. The example
below implements WmmaRdna3F16Policy using
__builtin_amdgcn_wmma_f32_16x16x16_f16_w32.
Each wavefront computes a single \(16 \times 16\) output tile. Due to
matrix replication, lanes 0–15 and 16–31 load identical A and B fragments
(lane_id % 16 selects the row or column), and each lane supplies 16 FP16
values for both operands.
The complete source file is available for download:
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 RDNA3 interleaved accumulator layout.
struct WmmaRdna3F16Policy
{
// -- ComputePolicy constants ----------------------------------------------
// One wavefront (wave32) owns a 16x16 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 effective_lanes = 32;
// v_wmma_f32_16x16x16_f16 consumes 16 K-positions per call (k_step=16).
// Each lane provides 16 FP16 values (v16half) for A and 16 for B.
// Lanes 0--15 and 16--31 mirror each other (matrix replication).
static constexpr int k_step = 16;
// -- Vector types --------------------------------------------------------
using v16half = _Float16 [[clang::ext_vector_type(16)]];
using v8float = float [[clang::ext_vector_type(8)]];
using elem_a = _Float16;
using elem_b = _Float16;
// -- Accumulator ----------------------------------------------------------
// v8float holds 8 FP32 VGPRs per lane (VGPRs 0--7).
struct Accumulator
{
v8float regs;
};
__device__ static void zero(Accumulator& acc)
{
acc.regs = v8float{};
}
// -- thread_tile_offset ---------------------------------------------------
// The entire 16x16 tile belongs to one wavefront (wave32). Multiple
// wavefronts in a block cover different 16x16 sub-tiles.
//
// wavefront id within block: wid = tid / 32
// 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;
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 16 FP16 values from LDS into the fragment
// array. The 16-element _Float16 array has the same layout as v16half
// and is reinterpreted as such in mma().
//
// v_wmma_f32_16x16x16_f16 (wave32) uses matrix replication: lanes 0--15
// and 16--31 index identically via (lane_id mod 16). Each lane loads
// one full row of the 16x16 A-tile.
//
// A-row index : tile_a_row + (lane_id mod 16)
// K-position e: ki + e (e = 0..15)
__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)[16])
{
const int row = tile_a_row + (lane_id % 16);
#pragma unroll
for(int e = 0; e < 16; ++e)
frag[e] = tile_a_ptr[row * k_tile_size + ki + e];
}
// load_b: each lane reads 16 FP16 values from LDS (symmetric to load_a).
//
// tile_b_T is stored transposed: tile_b_T[(tile_b_col + col)][k_off].
// B-col index : tile_b_col + (lane_id mod 16)
// K-position e: ki + e (e = 0..15)
__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)[16])
{
const int col = tile_b_col + (lane_id % 16);
#pragma unroll
for(int e = 0; e < 16; ++e)
frag[e] = tile_b_T_ptr[col * k_tile_size + ki + e];
}
// -- mma -------------------------------------------------------------------
// Issue one v_wmma_f32_16x16x16_f16 instruction.
//
// The _Float16[16] fragment arrays share the same in-register layout as
// v16half; a reinterpret_cast avoids any per-element copy.
__device__ static void mma(Accumulator& acc,
const elem_a (&a_frag)[16],
const elem_b (&b_frag)[16])
{
#if defined(__gfx1100__) || defined(__gfx1101__) || defined(__gfx1102__) || \
defined(__gfx1103__) || defined(__gfx1150__) || defined(__gfx1151__) || \
defined(__gfx1152__) || defined(__gfx1153__)
acc.regs = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32(
reinterpret_cast<const v16half&>(a_frag),
reinterpret_cast<const v16half&>(b_frag),
acc.regs);
#endif
}
// -- store_c ---------------------------------------------------------------
// Scatter the 8 VGPR values to their global-memory positions.
//
// Inverse layout (lane L, VGPR index G in [0,7]):
// row = G * 2 + (L / 16)
// col = 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 < 8; ++G)
{
const int i = G * 2 + (L / 16);
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 WmmaRdna3F16Policy in hand, plug it 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.
// TilePolicy for the FP16 WMMA path (transposed B in LDS).
using WmmaTilePolicy = SingleBufferTilePolicyF16<32, 32, 16>;
// Launch parameters:
//
// RDNA 3/3.5 (gfx11): 2x2 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): 8x8 thread tiles in a 32x32 block → 16 threads per
// block, padded to one full wavefront 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 32x32 output tile.
constexpr int BLOCK_TILE_M = 32;
constexpr int BLOCK_TILE_N = 32;
// Grid: one block per 32x32 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_gfx11)
{
auto launch = [&]()
{
matrix_multiply_generic<WmmaTilePolicy,
WmmaRdna3F16Policy,
_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, WmmaRdna3F16Policy> [gfx11]";
}
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(hipFree(d_A_f32));
HIP_CHECK(hipFree(d_B_f32));
}
Compile and run:
amdclang++ -O3 -std=c++17 --offload-arch=gfx1100 \
matrix_multiply_rdna3_wmma.hip -o mm_rdna3_wmma
./mm_rdna3_wmma
Note
WmmaRdna3F16Policy requires an RDNA3 or RDNA3.5 GPU (gfx1100 or later
gfx11 target). The #if defined(__gfx1100__) || ... 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 v16half = _Float16 [[clang::ext_vector_type(16)]];
using v16short = short [[clang::ext_vector_type(16)]];
using v8float = float [[clang::ext_vector_type(8)]];
using v8half = _Float16 [[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)]];
Common parameters#
The following parameters are shared across multiple WMMA builtins on this page.
Parameter |
Type |
Description |
|---|---|---|
|
|
FP16 and BF16 accumulate variants only. Selects which half of the
16-element accumulator register holds the result. When |
|
|
Integer variants only. When |
|
|
Integer variants only. Same as |
|
|
Integer variants only. When |
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 |
|---|---|---|
|
8192 |
32 |
|
8192 |
32 |
|
8192 |
32 |
|
8192 |
32 |
|
8192 |
32 |
|
8192 |
16 |
Builtin reference#
The following sections list every dense WMMA builtin available on RDNA3 and RDNA3.5, grouped by accumulator type.
FP32-accumulate builtins#
These builtins accumulate into FP32 and accept FP16 or BF16 matrix inputs.
FP16 inputs#
The following builtins use FP16 matrix inputs.
__builtin_amdgcn_wmma_f32_16x16x16_f16_w32#
Signature and parameters for this builtin.
v8float __builtin_amdgcn_wmma_f32_16x16x16_f16_w32(
v16half srcA,
v16half srcB,
v8float srcC);
Computes one step of a dense \(16 \times 16\) FP32 accumulation with FP16
inputs. Each lane supplies 16 FP16 values for both \(\pmb{A}\) and
\(\pmb{B}\). Due to matrix replication, the 16-element fragment is indexed
by lane_id % 16, so lanes 0–15 and 16–31 carry identical copies.
Parameter |
Type |
Description |
|---|---|---|
|
v16half |
Sixteen FP16 elements of \(\pmb{A}\) per lane (one full row of the \(16 \times 16\) A tile, replicated across both lane groups). |
|
v16half |
Sixteen FP16 elements of \(\pmb{B}\) per lane (one full column of the \(16 \times 16\) B tile, replicated across both lane groups). |
|
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#
Signature and parameters for this builtin.
v8float __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32(
v16short srcA,
v16short 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 |
|---|---|---|
|
v16short |
Sixteen BF16 elements of \(\pmb{A}\) per lane (BF16 stored as
|
|
v16short |
Sixteen BF16 elements of \(\pmb{B}\) per lane. |
|
v8float |
Accumulator input: eight FP32 elements per lane. |
Returns v8float – updated accumulator
(\(\text{srcA} \times \text{srcB} + \text{srcC}\)).
FP16-accumulate builtins#
These builtins accumulate into FP16 with FP16 inputs. Both accept an
opsel parameter that selects which half of the 16-element accumulator
register is written (see Common parameters). The
_tied variant constrains srcC and the return value to the same
physical register, preserving the non-selected half in place.
FP16 inputs#
The following builtins use FP16 matrix inputs.
__builtin_amdgcn_wmma_f16_16x16x16_f16_w32#
Signature and parameters for this builtin.
v16half __builtin_amdgcn_wmma_f16_16x16x16_f16_w32(
v16half srcA,
v16half srcB,
v16half srcC,
bool opsel);
Computes one step of a dense \(16 \times 16\) FP16 accumulation. Both
inputs and the accumulator are FP16. The opsel parameter selects the
storage half of the accumulator register (see
Common parameters).
The accumulator is a 16-element vector (v16half, 8 VGPRs × 2 FP16);
the 8 result elements occupy either the low (opsel = false) or high
(opsel = true) half of each VGPR pair.
Parameter |
Type |
Description |
|---|---|---|
|
v16half |
Sixteen FP16 elements of \(\pmb{A}\) per lane. |
|
v16half |
Sixteen FP16 elements of \(\pmb{B}\) per lane. |
|
v16half |
Accumulator input: sixteen FP16 elements per lane (8 VGPRs).
Only the half selected by |
|
|
Selects accumulator half. |
Returns v16half – updated accumulator
(\(\text{srcA} \times \text{srcB} + \text{srcC}\)).
__builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32#
Signature and parameters for this builtin.
v16half __builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32(
v16half srcA,
v16half srcB,
v16half srcC,
bool opsel);
Register-tied variant of
__builtin_amdgcn_wmma_f16_16x16x16_f16_w32. The compiler
constrains srcC and the return value to the same physical register,
guaranteeing that the non-selected half of the accumulator (the half not
written by this opsel setting) is preserved in place.
Use this variant when you chain two WMMA calls with opposite opsel
values into the same v16half accumulator so that both halves remain
live without an extra register copy.
Parameters, types, and return value are identical to the non-tied variant above.
BF16-accumulate builtins#
These builtins accumulate into BF16 with BF16 inputs. Both accept an
opsel parameter and the _tied variant provides register tying, as
described in the FP16 section above.
BF16 inputs#
The following builtins use BF16 matrix inputs.
__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32#
Signature and parameters for this builtin.
v16short __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32(
v16short srcA,
v16short srcB,
v16short srcC,
bool opsel);
Computes one step of a dense \(16 \times 16\) BF16 accumulation. Both
inputs and the accumulator are BF16 (stored as short). The opsel
parameter selects the storage half of the accumulator register (see
Common parameters).
The accumulator is a 16-element vector (v16short, 8 VGPRs × 2 BF16);
the 8 result elements occupy either the low (opsel = false) or high
(opsel = true) half of each VGPR pair.
Parameter |
Type |
Description |
|---|---|---|
|
v16short |
Sixteen BF16 elements of \(\pmb{A}\) per lane. |
|
v16short |
Sixteen BF16 elements of \(\pmb{B}\) per lane. |
|
v16short |
Accumulator input: sixteen BF16 elements per lane (8 VGPRs, stored
as |
|
|
Selects accumulator half. |
Returns v16short – updated accumulator (BF16 stored as short)
(\(\text{srcA} \times \text{srcB} + \text{srcC}\)).
__builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32#
Signature and parameters for this builtin.
v16short __builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32(
v16short srcA,
v16short srcB,
v16short srcC,
bool opsel);
Register-tied variant of
__builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32. The compiler
constrains srcC and the return value to the same physical register,
guaranteeing that the non-selected half of the accumulator (the half not
written by this opsel setting) is preserved in place.
Use this variant when you chain two WMMA calls with opposite opsel
values into the same v16short accumulator so that both halves remain
live without an extra register copy.
Parameters, types, and return value are identical to the non-tied variant above.
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#
The following builtins use INT8 and UINT8 matrix inputs.
__builtin_amdgcn_wmma_i32_16x16x16_iu8_w32#
Signature and parameters for this builtin.
v8int __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(
bool a_neg,
v4int srcA,
bool b_neg,
v4int 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 four int registers per lane; likewise for \(\pmb{B}\).
Parameter |
Type |
Description |
|---|---|---|
|
bool |
|
|
v4int |
Sixteen 8-bit elements of \(\pmb{A}\) per lane, packed into four 32-bit registers (four 8-bit values per register). |
|
bool |
|
|
v4int |
Sixteen 8-bit elements of \(\pmb{B}\) per lane, packed into four 32-bit registers. |
|
v8int |
Accumulator input: eight INT32 elements per lane. |
|
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#
The following builtins use INT4 and UINT4 matrix inputs.
__builtin_amdgcn_wmma_i32_16x16x16_iu4_w32#
Signature and parameters for this builtin.
v8int __builtin_amdgcn_wmma_i32_16x16x16_iu4_w32(
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. Sixteen INT4 elements of \(\pmb{A}\) are packed
into two int registers per lane; likewise for \(\pmb{B}\).
Parameter |
Type |
Description |
|---|---|---|
|
bool |
|
|
v2int |
Sixteen 4-bit elements of \(\pmb{A}\) per lane, packed into two 32-bit registers (eight 4-bit values per register). |
|
bool |
|
|
v2int |
Sixteen 4-bit elements of \(\pmb{B}\) per lane, packed into two 32-bit registers. |
|
v8int |
Accumulator input: eight INT32 elements per lane. |
|
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}\)).