RDNA4 sparse WMMA builtins#
Sparse Wave-Matrix Multiply-Accumulate (SWMMAC) builtins let you issue
hardware sparse matrix multiply-accumulate operations directly from HIP device
code on RDNA4 GPUs (gfx1200, gfx1201). Each SWMMAC instruction
multiplies a compressed sparse \(\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. Because the \(\pmb{A}\)
operand is stored in compressed form, SWMMAC halves the storage and bandwidth
required for \(\pmb{A}\) relative to a dense multiply of the same tile size.
SWMMAC is the sparse variant of the WMMA instruction family, which is used on RDNA (consumer) GPUs. CDNA (Instinct) GPUs provide a comparable operation through Sparse Matrix Fused Multiply-Accumulate (SMFMAC), the sparse variant of Matrix Fused Multiply-Accumulate (MFMA). The two differ in wavefront size (wave32 for SWMMAC, wave64 for SMFMAC) and accumulator storage (ordinary Vector General-Purpose Registers (VGPRs) for SWMMAC, dedicated accVGPRs for SMFMAC). The underlying sparsity model (2:4 structured sparsity on \(\pmb{A}\)) is the same on both architectures.
Note
RDNA4 GPUs run all shader programs in wave32 mode by default. The
_w32 suffix in each builtin name reflects this: all SWMMAC
builtins on this page require wavefrontsize32. wavefrontsize64 is not supported for HIP code.
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=gfx1201 ...
amdclang++ --offload-arch=gfx1200 ...
Naming convention#
All SWMMAC builtins follow the pattern:
__builtin_amdgcn_swmmac_<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. The K dimension refers to the compressed K of the sparse \(\pmb{A}\) operand. The corresponding dense K is twice as large due to 2:4 sparsity.
in_typeInput element type of \(\pmb{A}\) (and \(\pmb{B}\) when both share the same type):
f16,bf16,fp8,bf8,iu8, oriu4. Theiuprefix means the builtin accepts either signed or unsigned integers, controlled by thea_negandb_negparameters.in_type_b(optional)Input element type of \(\pmb{B}\) when it differs from \(\pmb{A}\). Used only for mixed FP8 and BF8 variants.
_w32Wavefront size suffix. All RDNA4 SWMMAC builtins use wave32.
Structured sparsity (2:4 pattern)#
SWMMAC instructions require the \(\pmb{A}\) matrix to obey 2:4 structured sparsity: in every contiguous group of four elements along the K dimension, exactly two are non-zero and the other two are zero. This constraint allows the \(\pmb{A}\) operand to be stored in a compressed representation containing only the non-zero values, reducing the storage to half the original K dimension.
Along with the compressed non-zero values, the hardware requires a sparsity
index that encodes which two of the four positions in each group hold the
non-zeros. This index is passed as the index argument to every SWMMAC
builtin. The hardware uses it at execution time to align the compressed
\(\pmb{A}\) elements against the correct rows of the \(\pmb{B}\)
operand before accumulating the products.
Host-side compression is straightforward: walk the K dimension in groups of four, extract the two non-zero positions, and pack them consecutively into the output buffer. The example kernel in this topic uses a fixed pattern that keeps positions 0 and 2 of every group, producing a compressed buffer of half the original K length.
Fragment layouts#
All SWMMAC 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 half the output rows. The diagrams below show the mapping between matrix elements and lane or VGPR positions for each operand.
The formulas in the subsections below use the following notation:
\(i\) – zero-based row index within the tile, \(0 \le i < 16\)
\(j\) – zero-based column index within the tile, \(0 \le j < 16\)
lane – wavefront lane, \(0 \le \text{lane} < 32\)
VGPR – zero-based index into that lane’s register vector
Accumulator layout#
Each lane holds 8 FP32 output elements across VGPRs 0–7.
Given output element \((i, j)\):
Conversely, given lane \(L\) and VGPR index \(g\):
The row-to-lane mapping:
Rows |
Lanes |
VGPRs |
|---|---|---|
0–7 |
0–15 |
0–7 |
8–15 |
16–31 |
0–7 |
srcA (sparse, FP16 and BF16)#
Each lane holds 8 compressed FP16 values (v8fp16, 4 VGPRs × 2 FP16)
covering one row of the sparse \(\pmb{A}\) matrix. The compressed-K
positions are non-contiguous across the two lane groups.
Lane \(L\) covers matrix row \(L \bmod 16\). The 8 compressed elements are distributed across VGPRs as follows:
Lane group |
VGPR 0 |
VGPR 1 |
VGPR 2 |
VGPR 3 |
|---|---|---|---|---|
0 (lanes 0–15) |
compressed K {0, 1} |
compressed K {2, 3} |
compressed K {8, 9} |
compressed K {10, 11} |
1 (lanes 16–31) |
compressed K {4, 5} |
compressed K {6, 7} |
compressed K {12, 13} |
compressed K {14, 15} |
srcB (dense, FP16 and BF16)#
Each lane holds 16 dense FP16 values (v16fp16, 8 VGPRs × 2 FP16)
covering one column of the dense \(\pmb{B}\) matrix. The K-row positions
are non-contiguous across the two lane groups.
Lane \(L\) covers matrix column \(L \bmod 16\). The 16 dense K rows are distributed across VGPRs as follows:
Lane group |
VGPRs 0–3 |
VGPRs 4–7 |
|---|---|---|
0 (lanes 0–15) |
K rows 0–7 |
K rows 16–23 |
1 (lanes 16–31) |
K rows 8–15 |
K rows 24–31 |
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 v2int = int [[clang::ext_vector_type(2)]];
using v4int = int [[clang::ext_vector_type(4)]];
using v8int = int [[clang::ext_vector_type(8)]];
using v8float = float [[clang::ext_vector_type(8)]];
using v8short = short [[clang::ext_vector_type(8)]]; // BF16 storage
using v16short = short [[clang::ext_vector_type(16)]]; // BF16 storage
using v8fp16 = __fp16 [[clang::ext_vector_type(8)]];
using v16fp16 = __fp16 [[clang::ext_vector_type(16)]];
Each type alias maps one-to-one to the corresponding LLVM vector type used in the builtin definition. The number in the name is the element count per lane.
Note
__fp16 and _Float16 are distinct Clang types. __fp16 is a
storage-only type: arithmetic on __fp16 values promotes to float
before the operation. _Float16 is the C standard FP16 arithmetic type
that supports native half-precision operations without promotion. RDNA4
SWMMAC builtins use __fp16; CDNA MFMA builtins use _Float16.
The two types are not implicitly convertible, so cast explicitly when
sharing FP16 data between code paths that target different architectures.
Note
BF16 matrix inputs use short as the storage type for RDNA4 SWMMAC
(not __bf16). Reinterpret your BF16 data with __builtin_bit_cast
or a union before passing it to the builtin.
Common parameters#
The index parameter is shared by all SWMMAC builtins. The a_neg,
b_neg, and clamp parameters appear only on integer variants.
Parameter |
Type |
Description |
|---|---|---|
|
|
Sparsity index register. Each pair of bits encodes the position (0–3)
of one non-zero element within its block of four consecutive
\(\pmb{A}\) elements along K. The 32-bit |
|
|
Integer variants only. When |
|
|
Integer variants only. Same as |
|
|
Integer variants only. When |
Using sparse 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 SwmmacRdna4F16Policy using
__builtin_amdgcn_swmmac_f32_16x16x32_f16_w32.
Each wavefront computes a single \(16 \times 16\) output tile. The \(\pmb{A}\) operand is pre-sparsified: half the K positions are zero and omitted from storage, so the compressed K dimension is 16 (representing 32 dense K positions).
The complete source file is available for download:
Policy constants
swmmac_f32_16x16x32_f16_w32 consumes 32 dense K positions per call
(compressed to K=16 in \(\pmb{A}\)), so k_step = 32. The
wavefront holds the entire \(16 \times 16\) tile: thread_tile_m =
thread_tile_n = 16 and effective_lanes = 32.
Sparsity index
Each lane’s index register encodes two bits per compressed K position,
identifying which of the four elements in each 2:4 block is non-zero. The
example constructs an even-column sparsity pattern (elements at positions
0 and 2 in each block of four) at the host and passes it to the device.
Accumulator layout
The builtin returns a v8float holding 8 FP32 values per lane. The
store_c() pass maps (lane, VGPR index) back to \((i, j)\)
coordinates using the SWMMAC accumulator layout.
// Vector type aliases required by the intrinsic.
// Use the __attribute__ form; [[clang::ext_vector_type]] is not supported
// for __fp16 in HIP device code.
typedef __fp16 v8fp16 __attribute__((ext_vector_type(8)));
typedef __fp16 v16fp16 __attribute__((ext_vector_type(16)));
typedef float v8float __attribute__((ext_vector_type(8)));
struct SwmmacRdna4F16Policy
{
// -- ComputePolicy constants ----------------------------------------------
// One wave32 wavefront owns a 16×16 output tile.
static constexpr int thread_tile_m = 16;
static constexpr int thread_tile_n = 16;
static constexpr int effective_lanes = 32;
// swmmac_f32_16x16x32_f16_w32 processes 32 dense K-positions per call.
static constexpr int k_step = 32;
using elem_a = __fp16;
using elem_b = __fp16;
// -- Accumulator ----------------------------------------------------------
// v8float holds 8 FP32 output elements per lane.
struct Accumulator
{
v8float regs;
};
__device__ static void zero(Accumulator& acc)
{
acc.regs = v8float{};
}
// -- thread_tile_offset ---------------------------------------------------
// The entire 16×16 tile belongs to one wavefront. Multiple wavefronts in
// a block cover different 16×16 sub-tiles.
//
// wavefront id within block: wid = tid / 32
// waves_n = block_tile_n / thread_tile_n
// wid_x = wid % waves_n
// wid_y = wid / waves_n
__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 8 compressed FP16 values and one index int from
// LDS.
//
// The SWMMAC instruction requires srcA as v8fp16 (8 compressed elements)
// and srcB as v16fp16 (16 dense elements). The index register is a scalar
// int containing 16 two-bit position codes (bits[15:0] used).
//
// Lane L covers A-row = tile_a_row + L % 16. The 8 compressed elements
// are loaded from non-contiguous positions within the row (AMD Matrix
// Instruction Calculator layout for v_swmmac_f32_16x16x32_f16_w32):
// Lanes 0-15: elements[0-3] = compressed k {0-3} (dense k 0- 7)
// elements[4-7] = compressed k {8-11} (dense k 16-23)
// Lanes 16-31: elements[0-3] = compressed k {4-7} (dense k 8-15)
// elements[4-7] = compressed k {12-15} (dense k 24-31)
// Compressed A fragment: 8 FP16 elements per lane.
// Stores both the FP16 data and the matching index int.
struct AFragment
{
v8fp16 data;
int index;
};
__device__ static AFragment load_a_frag(const __fp16* tile_a_ptr,
const int* tile_idx_ptr,
int tile_a_row,
int k_tile_sparse,
int lane_id)
{
AFragment frag;
const int row = tile_a_row + (lane_id % 16);
const int lane_grp = lane_id / 16; // 0 for lanes 0-15, 1 for lanes 16-31
// Lower 4 compressed elements: lanes 0-15 use k {0-3}, lanes 16-31 use k {4-7}.
const int k_lo = lane_grp * 4;
// Upper 4 compressed elements: lanes 0-15 use k {8-11}, lanes 16-31 use k {12-15}.
const int k_hi = 8 + lane_grp * 4;
const __fp16* row_ptr = tile_a_ptr + row * k_tile_sparse;
#pragma unroll
for(int e = 0; e < 4; ++e)
frag.data[e] = row_ptr[k_lo + e];
#pragma unroll
for(int e = 0; e < 4; ++e)
frag.data[4 + e] = row_ptr[k_hi + e];
// One int per row per K-tile (bits[15:0] = 4 × 4-bit groups).
// All lanes in the same output row load the same int; hardware uses
// lane_id / 16 to select the appropriate 8-bit half.
frag.index = tile_idx_ptr[row * (k_tile_sparse / 16)];
return frag;
}
// Dense B fragment: 16 FP16 elements per lane.
// tile_b_T is stored N-major: tile_b_T[n][k], stride = k_tile_dense (=32).
//
// SWMMAC srcB layout (wave32, 16x16x32, AMD Matrix Instruction Calculator):
// Lane L covers one B-column: n = tile_b_col + L % 16.
// The 16 k-positions are non-contiguous within that column:
// Lanes 0-15: elements[0-7] = B[0..7][n] (k 0- 7)
// elements[8-15]= B[16..23][n] (k 16-23)
// Lanes 16-31: elements[0-7] = B[8..15][n] (k 8-15)
// elements[8-15]= B[24..31][n] (k 24-31)
__device__ static v16fp16 load_b_frag(const __fp16* tile_b_T_ptr,
int tile_b_col,
int k_tile_dense,
int lane_id)
{
v16fp16 frag;
const int n = tile_b_col + (lane_id % 16);
const int lane_grp = lane_id / 16; // 0 for lanes 0-15, 1 for lanes 16-31
// Lower 8 k-positions: lanes 0-15 use k {0-7}, lanes 16-31 use k {8-15}.
const int k_lo = lane_grp * 8;
// Upper 8 k-positions: lanes 0-15 use k {16-23}, lanes 16-31 use k {24-31}.
const int k_hi = 16 + lane_grp * 8;
const __fp16* col_ptr = tile_b_T_ptr + n * k_tile_dense;
#pragma unroll
for(int e = 0; e < 8; ++e)
frag[e] = col_ptr[k_lo + e];
#pragma unroll
for(int e = 0; e < 8; ++e)
frag[8 + e] = col_ptr[k_hi + e];
return frag;
}
// -- mma -------------------------------------------------------------------
// Issue one swmmac_f32_16x16x32_f16_w32 instruction.
__device__ static void mma(Accumulator& acc,
const AFragment& a_frag,
const v16fp16& b_frag)
{
acc.regs = __builtin_amdgcn_swmmac_f32_16x16x32_f16_w32(
a_frag.data, b_frag, acc.regs, a_frag.index);
}
// -- store_c ---------------------------------------------------------------
// Scatter the 8 accumulator values to their global-memory positions.
//
// SWMMAC 16x16x32 wave32 output layout (AMD Matrix Instruction Calculator):
// Lane L, VGPR g (0-7):
// row = (L / 16) * 8 + g
// col = 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 col = lane_id % 16;
#pragma unroll
for(int g = 0; g < 8; ++g)
{
const int r = out_row_base + (lane_id / 16) * 8 + g;
const int c = out_col_base + col;
if(r < m && c < n)
C[r * n + c] = acc.regs[g];
}
}
};
Instantiating the kernel
With SwmmacRdna4F16Policy in place, plug it into the generic kernel
alongside a TilePolicy (see Using MFMA builtins as a compute policy) whose
block_tile_m and block_tile_n are multiples of 16 and whose
k_tile_size is a multiple of k_step = 32.
// TilePolicy: 32×32 block tile, 16-element sparse K-strip (= 32 dense).
// SwmmacRdna4F16Policy: one 16×16 tile per wave32 wavefront.
using SwmmacTilePolicy = SingleBufferTilePolicyH<32, 32, 16>;
// Block tile: 32×32 (2×2 wavefronts, each computing a 16×16 SWMMAC tile).
// Threads per block: 2 × 2 wavefronts × 32 lanes = 128.
constexpr int BLOCK_TILE_M = 32;
constexpr int BLOCK_TILE_N = 32;
constexpr int THREADS_PER_BLOCK = 128; // 4 wave32 wavefronts × 32 lanes
// 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);
const dim3 block(THREADS_PER_BLOCK);
auto launch = [&]()
{
matrix_multiply_swmmac<SwmmacTilePolicy>
<<<grid, block>>>(d_A_sparse, reinterpret_cast<const int*>(d_A_index),
d_B, d_C, M, N, K_SPARSE);
HIP_CHECK(hipGetLastError());
};
Compile and run:
amdclang++ -O3 -std=c++17 --offload-arch=gfx1201 \
matrix_multiply_rdna4_swmmac.hip -o mm_rdna4_swmmac
./mm_rdna4_swmmac
amdclang++ -O3 -std=c++17 --offload-arch=gfx1200 \
matrix_multiply_rdna4_swmmac.hip -o mm_rdna4_swmmac
./mm_rdna4_swmmac
Note
SwmmacRdna4F16Policy requires an RDNA4 GPU (gfx1200 or gfx1201).
The #if defined(__gfx1200__) || defined(__gfx1201__)
guard in the example file falls back to ScalarFMASPolicy on other targets,
so the file compiles without modification.
The example uses pre-sparsified input data with a fixed 2:4 pattern. In a production kernel, apply a sparsity pruning pass to the weight matrix offline and store the compressed values and index tensor separately.
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 |
|---|---|---|
|
16384 |
16 |
|
16384 |
16 |
|
16384 |
16 |
|
16384 |
16 |
|
16384 |
8 |
|
16384 |
8 |
|
32768 |
8 |
|
16384 |
8 |
|
16384 |
8 |
|
16384 |
8 |
|
16384 |
8 |
Builtin reference#
The following sections list every SWMMAC 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_swmmac_f32_16x16x32_f16_w32#
Signature and parameters for this builtin.
v8float __builtin_amdgcn_swmmac_f32_16x16x32_f16_w32(
v8fp16 srcA,
v16fp16 srcB,
v8float srcC,
int index);
Computes one step of a sparse \(16 \times 16\) FP32 accumulation with FP16
inputs. \(\pmb{A}\) is a compressed sparse fragment (8 FP16 values per lane
representing 16 K positions after 2:4 expansion). \(\pmb{B}\) is a dense
fragment (16 FP16 values per lane). The index register identifies the
non-zero positions in \(\pmb{A}\).
Parameter |
Type |
Description |
|---|---|---|
|
v8fp16 |
Eight compressed FP16 elements of \(\pmb{A}\) per lane. |
|
v16fp16 |
Sixteen dense FP16 elements of \(\pmb{B}\) per lane. |
|
v8float |
Accumulator input: eight FP32 elements per lane. |
|
int |
Sparsity index, see Common parameters. |
Returns v8float – updated accumulator
(\(\text{expand}(\text{srcA}, \text{index}) \times \text{srcB} + \text{srcC}\)).
BF16 inputs#
The following builtins use BF16 matrix inputs.
__builtin_amdgcn_swmmac_f32_16x16x32_bf16_w32#
Signature and parameters for this builtin.
v8float __builtin_amdgcn_swmmac_f32_16x16x32_bf16_w32(
v8short srcA,
v16short srcB,
v8float srcC,
int index);
Computes one step of a sparse \(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 |
|---|---|---|
|
v8short |
Eight compressed BF16 elements of \(\pmb{A}\) per lane (BF16 stored
as |
|
v16short |
Sixteen dense BF16 elements of \(\pmb{B}\) per lane. |
|
v8float |
Accumulator input: eight FP32 elements per lane. |
|
int |
Sparsity index, see Common parameters. |
Returns v8float – updated accumulator
(\(\text{expand}(\text{srcA}, \text{index}) \times \text{srcB} + \text{srcC}\)).
FP8 and BF8 inputs#
The following builtins use FP8 and BF8 matrix inputs.
__builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w32#
Signature and parameters for this builtin.
v8float __builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w32(
v2int srcA,
v4int srcB,
v8float srcC,
int index);
Computes one step of a sparse \(16 \times 16\) FP32 accumulation with FP8 (E4M3) inputs for both \(\pmb{A}\) and \(\pmb{B}\). Eight compressed FP8 values per lane are packed into two 32-bit registers; sixteen dense FP8 values into four registers.
Parameter |
Type |
Description |
|---|---|---|
|
v2int |
Eight compressed FP8 elements of \(\pmb{A}\) per lane, packed into two 32-bit registers (four 8-bit values per register). |
|
v4int |
Sixteen dense FP8 elements of \(\pmb{B}\) per lane, packed into four 32-bit registers. |
|
v8float |
Accumulator input: eight FP32 elements per lane. |
|
int |
Sparsity index, see Common parameters. |
Returns v8float – updated accumulator
(\(\text{expand}(\text{srcA}, \text{index}) \times \text{srcB} + \text{srcC}\)).
__builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w32#
Signature and parameters for this builtin.
v8float __builtin_amdgcn_swmmac_f32_16x16x32_fp8_bf8_w32(
v2int srcA,
v4int srcB,
v8float srcC,
int index);
Computes one step of a sparse \(16 \times 16\) FP32 accumulation with mixed 8-bit floating-point inputs: FP8 (E4M3) for \(\pmb{A}\) and BF8 (E5M2) for \(\pmb{B}\).
Parameter |
Type |
Description |
|---|---|---|
|
v2int |
Eight compressed FP8 elements of \(\pmb{A}\) per lane, packed into two 32-bit registers (four 8-bit values per register). |
|
v4int |
Sixteen dense BF8 elements of \(\pmb{B}\) per lane, packed into four 32-bit registers. |
|
v8float |
Accumulator input: eight FP32 elements per lane. |
|
int |
Sparsity index, see Common parameters. |
Returns v8float – updated accumulator
(\(\text{expand}(\text{srcA}, \text{index}) \times \text{srcB} + \text{srcC}\)).
__builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w32#
Signature and parameters for this builtin.
v8float __builtin_amdgcn_swmmac_f32_16x16x32_bf8_fp8_w32(
v2int srcA,
v4int srcB,
v8float srcC,
int index);
Computes one step of a sparse \(16 \times 16\) FP32 accumulation with mixed 8-bit floating-point inputs: BF8 (E5M2) for \(\pmb{A}\) and FP8 (E4M3) for \(\pmb{B}\).
Parameter |
Type |
Description |
|---|---|---|
|
v2int |
Eight compressed BF8 elements of \(\pmb{A}\) per lane, packed into two 32-bit registers (four 8-bit values per register). |
|
v4int |
Sixteen dense FP8 elements of \(\pmb{B}\) per lane, packed into four 32-bit registers. |
|
v8float |
Accumulator input: eight FP32 elements per lane. |
|
int |
Sparsity index, see Common parameters. |
Returns v8float – updated accumulator
(\(\text{expand}(\text{srcA}, \text{index}) \times \text{srcB} + \text{srcC}\)).
__builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w32#
Signature and parameters for this builtin.
v8float __builtin_amdgcn_swmmac_f32_16x16x32_bf8_bf8_w32(
v2int srcA,
v4int srcB,
v8float srcC,
int index);
Computes one step of a sparse \(16 \times 16\) FP32 accumulation with BF8 (E5M2) inputs for both \(\pmb{A}\) and \(\pmb{B}\). Eight compressed BF8 values per lane are packed into two 32-bit registers; sixteen dense BF8 values into four registers.
Parameter |
Type |
Description |
|---|---|---|
|
v2int |
Eight compressed BF8 elements of \(\pmb{A}\) per lane, packed into two 32-bit registers (four 8-bit values per register). |
|
v4int |
Sixteen dense BF8 elements of \(\pmb{B}\) per lane, packed into four 32-bit registers. |
|
v8float |
Accumulator input: eight FP32 elements per lane. |
|
int |
Sparsity index, see Common parameters. |
Returns v8float – updated accumulator
(\(\text{expand}(\text{srcA}, \text{index}) \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_swmmac_f16_16x16x32_f16_w32#
Signature and parameters for this builtin.
v8fp16 __builtin_amdgcn_swmmac_f16_16x16x32_f16_w32(
v8fp16 srcA,
v16fp16 srcB,
v8fp16 srcC,
int index);
Computes one step of a sparse \(16 \times 16\) FP16 accumulation. Both inputs and the accumulator are FP16. Operand sizes and the sparsity index are identical to the FP32-accumulate FP16 variant.
Parameter |
Type |
Description |
|---|---|---|
|
v8fp16 |
Eight compressed FP16 elements of \(\pmb{A}\) per lane. |
|
v16fp16 |
Sixteen dense FP16 elements of \(\pmb{B}\) per lane. |
|
v8fp16 |
Accumulator input: eight FP16 elements per lane. |
|
int |
Sparsity index, see Common parameters. |
Returns v8fp16 – updated accumulator
(\(\text{expand}(\text{srcA}, \text{index}) \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_swmmac_bf16_16x16x32_bf16_w32#
Signature and parameters for this builtin.
v8short __builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w32(
v8short srcA,
v16short srcB,
v8short srcC,
int index);
Computes one step of a sparse \(16 \times 16\) BF16 accumulation. Both
inputs and the accumulator are BF16 (stored as short).
Parameter |
Type |
Description |
|---|---|---|
|
v8short |
Eight compressed BF16 elements of \(\pmb{A}\) per lane. |
|
v16short |
Sixteen dense BF16 elements of \(\pmb{B}\) per lane. |
|
v8short |
Accumulator input: eight BF16 elements per lane (stored as |
|
int |
Sparsity index, see Common parameters. |
Returns v8short – updated accumulator (BF16 stored as short)
(\(\text{expand}(\text{srcA}, \text{index}) \times \text{srcB} + \text{srcC}\)).
INT32-accumulate builtins#
Integer SWMMAC 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 (16x16x32)#
The following builtins use INT8 and UINT8 matrix inputs.
__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w32#
Signature and parameters for this builtin.
v8int __builtin_amdgcn_swmmac_i32_16x16x32_iu8_w32(
bool a_neg,
v2int srcA,
bool b_neg,
v4int srcB,
v8int srcC,
int index,
bool clamp);
Computes one step of a sparse \(16 \times 16\) INT32 accumulation with
8-bit integer inputs. \(\pmb{A}\) is packed as two int registers per
lane (8 bytes = 8 INT8 elements after 2:4 expansion to 16 K positions).
\(\pmb{B}\) is packed as four int registers per lane (16 bytes = 16
INT8 elements).
Parameter |
Type |
Description |
|---|---|---|
|
bool |
|
|
v2int |
Eight compressed 8-bit elements of \(\pmb{A}\) per lane, packed into two 32-bit registers. |
|
bool |
|
|
v4int |
Sixteen dense 8-bit elements of \(\pmb{B}\) per lane, packed into four 32-bit registers. |
|
v8int |
Accumulator input: eight INT32 elements per lane. |
|
int |
Sparsity index, see Common parameters. |
|
bool |
Clamp output to input type range on overflow. Compile-time constant, see Common parameters. |
Returns v8int – updated accumulator
(\(\text{expand}(\text{srcA}, \text{index}) \times \text{srcB} + \text{srcC}\)).
INT4 and UINT4 inputs (16x16x32)#
The following builtins use INT4 and UINT4 matrix inputs.
__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w32#
Signature and parameters for this builtin.
v8int __builtin_amdgcn_swmmac_i32_16x16x32_iu4_w32(
bool a_neg,
int srcA,
bool b_neg,
v2int srcB,
v8int srcC,
int index,
bool clamp);
Computes one step of a sparse \(16 \times 16\) INT32 accumulation with
4-bit integer inputs. Eight INT4 elements of \(\pmb{A}\) fit into a single
int per lane; sixteen INT4 elements of \(\pmb{B}\) fit into two
int registers.
Parameter |
Type |
Description |
|---|---|---|
|
bool |
|
|
int |
Eight compressed 4-bit elements of \(\pmb{A}\) per lane, packed into one 32-bit register. |
|
bool |
|
|
v2int |
Sixteen dense 4-bit elements of \(\pmb{B}\) per lane, packed into two 32-bit registers. |
|
v8int |
Accumulator input: eight INT32 elements per lane. |
|
int |
Sparsity index, see Common parameters. |
|
bool |
Clamp output to input type range on overflow. Compile-time constant, see Common parameters. |
Returns v8int – updated accumulator
(\(\text{expand}(\text{srcA}, \text{index}) \times \text{srcB} + \text{srcC}\)).
INT4 and UINT4 inputs (16x16x64)#
The following builtins use INT4 and UINT4 matrix inputs.
__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w32#
Signature and parameters for this builtin.
v8int __builtin_amdgcn_swmmac_i32_16x16x64_iu4_w32(
bool a_neg,
v2int srcA,
bool b_neg,
v4int srcB,
v8int srcC,
int index,
bool clamp);
Computes one step of a sparse \(16 \times 16\) INT32 accumulation with
4-bit integer inputs over a deeper K=64 strip. Sixteen INT4 elements of
\(\pmb{A}\) are packed into two int registers per lane; thirty-two
INT4 elements of \(\pmb{B}\) into four int registers.
Parameter |
Type |
Description |
|---|---|---|
|
bool |
|
|
v2int |
Sixteen compressed 4-bit elements of \(\pmb{A}\) per lane, packed into two 32-bit registers. |
|
bool |
|
|
v4int |
Thirty-two dense 4-bit elements of \(\pmb{B}\) per lane, packed into four 32-bit registers. |
|
v8int |
Accumulator input: eight INT32 elements per lane. |
|
int |
Sparsity index, see Common parameters. |
|
bool |
Clamp output to input type range on overflow. Compile-time constant, see Common parameters. |
Returns v8int – updated accumulator
(\(\text{expand}(\text{srcA}, \text{index}) \times \text{srcB} + \text{srcC}\)).