CDNA4 MFMA transpose load builtins#
The CDNA4 (gfx950, MI350 series) architecture introduces a set of LDS
(Local Data Share) load builtins that perform a hardware-assisted transpose
as part of the load. These builtins let you store matrix \(\pmb{A}\) tiles
in LDS in column-major order or matrix \(\pmb{B}\) tiles in row-major order,
then load them directly into the per-lane fragment layout expected by the
CDNA4 dense MFMA builtins, without a software shuffle step.
ds_read_tr loads B from LDS in row-major order and distributes each
column to the corresponding lane’s VGPRs. The simplified example uses
\(K = 4\) and \(N = 8\); colors identify each \(N\)-column.#
Architecture availability#
The builtins on this page target CDNA4 (gfx950, MI350 series)
exclusively. For the MFMA compute builtins they are designed to feed, see
CDNA4 dense MFMA builtins.
Naming convention#
All transpose load builtins follow the pattern:
__builtin_amdgcn_ds_read_tr<N>_b<M>_v<K><type>
NTranspose group width: the number of elements exchanged between adjacent lanes during the transpose operation.
MBits loaded per lane from LDS.
KNumber of elements in the return vector.
typeElement type of the return vector:
i32,i16,f16, orbf16.
Register types used in this reference#
The signatures below use the following type aliases, which you can declare with Clang vector attributes in any HIP translation unit:
using v2int = int [[clang::ext_vector_type(2)]]; // FP4 / FP8 / BF8
using v3int = int [[clang::ext_vector_type(3)]]; // FP6
using v4short = short [[clang::ext_vector_type(4)]]; // INT8 (2 × INT8 per short)
using v4half = _Float16 [[clang::ext_vector_type(4)]]; // FP16
using v4bfloat = short [[clang::ext_vector_type(4)]]; // BF16 (1 × BF16 per short)
Builtin reference#
The following builtins are available on CDNA4. Each takes a single
pointer argument in LDS address space (__shared__) and returns a vector
holding the lane’s share of the transposed tile. All builtins are
const and have no side effects.
Hardware constraints#
The following constraints apply to every builtin in this family:
Pair of calls required. One call loads only half of the matrix tile’s K positions. A second call with a different
ptrvalue and a different destination register pair completes the tile. The K-position split varies by element size and is summarized in the table below.LDS address alignment. The
ptraddress must be aligned to the load data size: 8 bytes forb64loads and 12 bytes forb96loads.VGPR alignment. All
b64loads write to an even-aligned VGPR pair. Theb96load (ds_read_tr6_b96_v3i32) does not require even-VGPR alignment.
Builtin |
Call 1 K positions |
Call 2 K positions |
|---|---|---|
|
0—15, 32—47 |
16—31, 48—63 |
|
0—15, 32—47 |
16—31, 48—63 |
|
0—7, 16—23, 32—39, 48—55 |
8—15, 24—31, 40—47, 56—63 |
|
0—3, 8—11 |
4—7, 12—15 |
|
0—3, 8—11 |
4—7, 12—15 |
|
0—3, 8—11 |
4—7, 12—15 |
K-position split for each ds_read_tr builtin. Teal cells are
covered by the first call; grey cells by the second. White cells indicate
K positions beyond the builtin’s total K depth.#
Sub-byte operand loads#
These builtins load and transpose operands for the scaled sub-byte MFMA
family (__builtin_amdgcn_mfma_scale_f32_*).
__builtin_amdgcn_ds_read_tr4_b64_v2i32#
v2int __builtin_amdgcn_ds_read_tr4_b64_v2i32(__shared__ v2int* ptr);
Loads 64 bits per lane from LDS with a 4-bit element transpose and returns
the result as a v2int (two 32-bit words).
Parameter |
Type |
Description |
|---|---|---|
|
|
Pointer to the lane’s portion of the FP4 tile in LDS. |
Returns v2int – two 32-bit words holding the lane’s FP4 fragment
after transposition.
__builtin_amdgcn_ds_read_tr6_b96_v3i32#
v3int __builtin_amdgcn_ds_read_tr6_b96_v3i32(__shared__ v3int* ptr);
Loads 96 bits per lane from LDS with a 6-bit element transpose and returns
the result as a v3int (three 32-bit words).
Parameter |
Type |
Description |
|---|---|---|
|
|
Pointer to the lane’s portion of the FP6 tile in LDS. |
Returns v3int – three 32-bit words holding the lane’s FP6 fragment
after transposition.
8-bit operand loads#
These builtins load and transpose operands for the FP8, BF8, and INT8 MFMA families.
__builtin_amdgcn_ds_read_tr8_b64_v2i32#
v2int __builtin_amdgcn_ds_read_tr8_b64_v2i32(__shared__ v2int* ptr);
Loads 64 bits per lane from LDS with an 8-bit element transpose and returns
the result as a v2int (two 32-bit words).
Parameter |
Type |
Description |
|---|---|---|
|
|
Pointer to the lane’s portion of the FP8 or BF8 tile in LDS. |
Returns v2int – two 32-bit words holding the lane’s FP8 or BF8
fragment after transposition.
__builtin_amdgcn_ds_read_tr16_b64_v4i16#
v4short __builtin_amdgcn_ds_read_tr16_b64_v4i16(__shared__ v4short* ptr);
Loads 64 bits per lane from LDS with a 16-bit element transpose and returns
the result as a v4short (four 16-bit words, each holding two packed INT8
values). Each lane receives four consecutive values along the M or N
dimension of the matrix.
Parameter |
Type |
Description |
|---|---|---|
|
|
Pointer to the lane’s portion of the INT8 tile in LDS. |
Returns v4short – four 16-bit words (eight packed INT8 values)
holding the lane’s INT8 fragment after transposition.
16-bit operand loads#
These builtins load and transpose operands for the FP16 and BF16 MFMA families.
__builtin_amdgcn_ds_read_tr16_b64_v4f16#
v4half __builtin_amdgcn_ds_read_tr16_b64_v4f16(__shared__ v4half* ptr);
Loads 64 bits per lane from LDS with a 16-bit element transpose and returns
the result as a v4half (four FP16 values). Each lane receives four
consecutive values along the M or N dimension of the matrix.
Parameter |
Type |
Description |
|---|---|---|
|
|
Pointer to the lane’s portion of the FP16 tile in LDS. |
Returns v4half – four FP16 values holding the lane’s FP16 fragment
after transposition.
Note
This builtin requires the __fp16 type (GCC-style half-precision),
not _Float16 (ISO C half-precision). When calling the builtin,
cast pointers and vector types to __fp16-based variants. The standard
MFMA compute builtins accept _Float16, so a cast is needed only at
the ds_read_tr call site.
__builtin_amdgcn_ds_read_tr16_b64_v4bf16#
v4bfloat __builtin_amdgcn_ds_read_tr16_b64_v4bf16(__shared__ v4bfloat* ptr);
Loads 64 bits per lane from LDS with a 16-bit element transpose and returns
the result as a v4bfloat (four BF16 values). Each lane receives four
consecutive values along the M or N dimension of the matrix.
Parameter |
Type |
Description |
|---|---|---|
|
|
Pointer to the lane’s portion of the BF16 tile in LDS. |
Returns v4bfloat – four BF16 values holding the lane’s BF16
fragment after transposition.
FP16 GEMM example#
The following excerpts from
docs/tools/example_codes/matrix_multiply_cdna4_mfma.hip show how to use
__builtin_amdgcn_ds_read_tr16_b64_v4f16 in a complete FP16 GEMM kernel.
Two cooperating structures are required: a tile policy that stores the B tile
in LDS in row-major order, and a compute policy that issues the transpose load
calls in place of the scalar load loop.
Step 1: store B in row-major order in LDS#
The standard FP16 tile policy stores B transposed (tile_b_T[N][K]) so
that scalar lane indexing can stride along the K dimension. The
ds_read_tr path instead requires B in row-major order (tile_b[K][N])
so the hardware can scatter each N-column to its lane.
template<int BlockTileM_, int BlockTileN_, int KTileSize_>
struct SingleBufferTilePolicyF16RowB
{
static constexpr int num_buffers = 1;
static constexpr int block_tile_m = BlockTileM_;
static constexpr int block_tile_n = BlockTileN_;
static constexpr int k_tile_size = KTileSize_;
struct SharedStorage
{
_Float16 tile_a[BlockTileM_][KTileSize_];
// B stored row-major: tile_b[k][n], so that ds_read_tr can transpose
// K-rows into per-lane N-column fragments in hardware.
_Float16 tile_b[KTileSize_][BlockTileN_];
};
__device__ static void prologue(SharedStorage& /*smem*/,
const _Float16* /*A*/, const _Float16* /*B*/,
int /*block_row_start*/, int /*block_col_start*/,
int /*m*/, int /*n*/, int /*k*/,
int /*tid*/, int /*num_threads*/) {}
__device__ static void prefetch(SharedStorage& smem,
const _Float16* A,
const _Float16* B,
int t_next,
int block_row_start, int block_col_start,
int m, int n, int k,
int tid, int num_threads)
{
const int t = t_next;
// Load A tile: row-major, M rows × K cols.
const int tile_a_elems = block_tile_m * k_tile_size;
for(int idx = tid; idx < tile_a_elems; idx += num_threads)
{
const int tile_row = idx / k_tile_size;
const int tile_col = idx % k_tile_size;
const int global_row = block_row_start + tile_row;
const int global_col = t * k_tile_size + tile_col;
smem.tile_a[tile_row][tile_col] =
(global_row < m && global_col < k)
? A[global_row * k + global_col]
: static_cast<_Float16>(0.0f);
}
// Load B tile: row-major K × N layout in LDS.
const int tile_b_elems = k_tile_size * block_tile_n;
for(int idx = tid; idx < tile_b_elems; idx += num_threads)
{
const int tile_k_idx = idx / block_tile_n;
const int tile_col_idx = idx % block_tile_n;
const int global_row = t * k_tile_size + tile_k_idx;
const int global_col = block_col_start + tile_col_idx;
smem.tile_b[tile_k_idx][tile_col_idx] =
(global_row < k && global_col < n)
? B[global_row * n + global_col]
: static_cast<_Float16>(0.0f);
}
}
__device__ static void acquire(int /*t*/) { __syncthreads(); }
__device__ static void release(int /*t*/) { __syncthreads(); }
__device__ static int buf_idx(int /*t*/) { return 0; }
__device__ static _Float16* get_tile_a(SharedStorage& smem, int /*buf*/)
{
return &smem.tile_a[0][0];
}
// Returns the row-major B tile base pointer; the ds_read_tr policy will
// compute per-lane pointers into this region.
__device__ static _Float16* get_tile_b_T(SharedStorage& smem, int /*buf*/)
{
return &smem.tile_b[0][0];
}
};
Step 2: load B with hardware-transposed calls#
MfmaCdna4F16TrPolicy::load_b replaces the scalar load loop with two
ds_read_tr16_b64_v4f16 calls. Each call loads 64 bits (four FP16 values)
per lane and performs the 16-element transpose in hardware.
v_mfma_f32_16x16x32_f16 consumes 32 K-positions per call
(k_step = 32). The 64 lanes form four groups of 16
(g = lane_id / 16), each group covering 8 consecutive K-positions
starting at ki + 8 * g. Each ds_read_tr16_b64_v4f16 call fills half
of a group’s 8 K-positions (see the K-position split table above), so two
calls per load_b invocation produce the full 8-element B fragment:
Call 1 (
r0): each group loads the first 4 of its 8 K-positions. Wavefront-wide this covers K positionski + 0toki + 3,ki + 8toki + 11,ki + 16toki + 19, andki + 24toki + 27.Call 2 (
r1): each group loads the remaining 4 K-positions. Wavefront-wide:ki + 4toki + 7,ki + 12toki + 15,ki + 20toki + 23, andki + 24toki + 31.
Each lane addresses its own N-column:
column = tile_b_col + (lane_id mod 16).
The pointer stride between the two calls is
4 * BlockTileN * sizeof(__fp16) bytes (four K-rows of the B tile in LDS).
struct MfmaCdna4F16TrPolicy
{
// -- ComputePolicy constants ----------------------------------------------
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;
static constexpr int effective_lanes = 64;
// k_step=32: one v_mfma_f32_16x16x32_f16 call consumes 32 K-positions,
// filled by two ds_read_tr16_b64_v4f16 calls.
static constexpr int k_step = 32;
using elem_a = _Float16;
using elem_b = _Float16;
// Accumulator and zero() are identical to MfmaCdna4F16Policy.
using v4float = float [[clang::ext_vector_type(4)]];
struct Accumulator
{
v4float regs;
};
__device__ static void zero(Accumulator& acc)
{
acc.regs = v4float{};
}
// thread_tile_offset is identical to MfmaCdna4F16Policy.
__device__ static void thread_tile_offset(int tid,
int /*lane_id*/,
int* thread_row,
int* thread_col)
{
constexpr int waves_n = 2;
const int wid = tid / 64;
const int wid_x = wid % waves_n;
const int wid_y = wid / waves_n;
*thread_row = wid_y * 16;
*thread_col = wid_x * 16;
}
// -- load_a ---------------------------------------------------------------
// Scalar load, identical to MfmaCdna4F16Policy::load_a.
__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 ---------------------------------------------------------------
// Hardware-transposed load using ds_read_tr16_b64_v4f16.
//
// tile_b_ptr points to the base of the row-major B tile in LDS:
// tile_b[k][n], dimensions KTileSize × BlockTileN.
//
// Each lane belongs to one of four 16-lane groups (g = lane_id / 16),
// handling K-positions ki + 8*g through ki + 8*g + 7. Two ds_read_tr
// calls fill the eight-element B fragment per lane:
//
// local_tx = lane_id % 16
// k_base = ki + 8 * (lane_id / 16)
//
// Call 1: ptr = &tile_b[k_base + local_tx/4][tile_b_col + (local_tx%4)*4]
// -> frag[0..3]
// Call 2: ptr = &tile_b[k_base + 4 + local_tx/4][tile_b_col + (local_tx%4)*4]
// -> frag[4..7]
//
// The stride between the two calls is 4 rows of the B tile (4 * BlockTileN
// elements). The hardware transpose gathers strided K elements across
// adjacent lanes so that each lane receives its own column's data.
__device__ static void load_b(const _Float16* tile_b_ptr,
int tile_b_col,
int ki,
int block_tile_n,
int lane_id,
elem_b (&frag)[8])
{
#if defined(__gfx950__)
using v4half = __fp16 [[clang::ext_vector_type(4)]];
const int local_tx = lane_id % 16;
const int k_base = ki + 8 * (lane_id / 16);
__fp16* base0 =
const_cast<__fp16*>(reinterpret_cast<const __fp16*>(tile_b_ptr))
+ (k_base + local_tx / 4) * block_tile_n
+ tile_b_col + (local_tx % 4) * 4;
__fp16* base1 =
const_cast<__fp16*>(reinterpret_cast<const __fp16*>(tile_b_ptr))
+ (k_base + 4 + local_tx / 4) * block_tile_n
+ tile_b_col + (local_tx % 4) * 4;
const v4half r0 = __builtin_amdgcn_ds_read_tr16_b64_v4f16(
(__attribute__((address_space(3))) v4half*)(base0));
const v4half r1 = __builtin_amdgcn_ds_read_tr16_b64_v4f16(
(__attribute__((address_space(3))) v4half*)(base1));
frag[0] = r0[0]; frag[1] = r0[1]; frag[2] = r0[2]; frag[3] = r0[3];
frag[4] = r1[0]; frag[5] = r1[1]; frag[6] = r1[2]; frag[7] = r1[3];
#endif
}
// mma and store_c are identical to MfmaCdna4F16Policy.
__device__ static void mma(Accumulator& acc,
const elem_a (&a_frag)[8],
const elem_b (&b_frag)[8])
{
#if defined(__gfx950__)
using v8half = _Float16 [[clang::ext_vector_type(8)]];
const v8half 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 v8half 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_mfma_f32_16x16x32_f16(
a_vec, b_vec, acc.regs,
/*cbsz=*/0, /*abid=*/0, /*blgp=*/0);
#endif
}
__device__ static void store_c(const Accumulator& acc,
float* C,
int out_row_base,
int out_col_base,
int m,
int n,
int lane_id)
{
const int L = lane_id;
const int j = L % 16;
#pragma unroll
for(int G = 0; G < 4; ++G)
{
const int i = 4 * (L / 16) + (G % 4);
const int r = out_row_base + i;
const int c = out_col_base + j;
if(r < m && c < n)
C[r * n + c] = acc.regs[G];
}
}
};