For the complete documentation index, see llms.txt. Markdown versions of all pages are available by appending .md to any URL (e.g. /max/get-started.md).

Mojo function

matmul2d_mma_regc_bt_native

def matmul2d_mma_regc_bt_native(a_frag: SIMD[DType.bfloat16, SIMDLength(8)], b_frag: SIMD[DType.bfloat16, SIMDLength(16)], mut c_acc: SIMD[DType.float32, SIMDLength(16)])

Pure-Mojo transpose_right=1 16x32x16 MMA (native simdgroup_matrix).

Fragment ABI: this lane's 8-element A, 16-element transpose_right B, and 16-element f32 C accumulator; (N, K) right-operand layout; fp32 reduction. Calls the native _mma_apple_transposable intrinsic (llvm.air.simdgroup_matrix_16x16x16_multiply_accumulate) directly -- no external library or codegen dependency.

The 16x32 tile decomposes into two side-by-side 16x16x16 MMAs along N:

half 0 (N 0..15): A @ B[0:8] -> C[0:8]
half 1 (N 16..31): A @ B[8:16] -> C[8:16] Each half's fragments are in the native _apple_frag_layout order: the per-lane element map (bt_frag_coord / bc_frag_coord) restricted to one N-half is bit-for-bit the native 16x16 fragment layout (frag_row_base / frag_col_base == _apple_frag_layout for all 32 lanes; the two C halves stitch back into exactly bc_frag_coord, so the C-store map is UNCHANGED). transpose_a=False, transpose_b=True reads B as (N, K), matching the coalesced bt_frag_coord feed.

Args:

a_frag (SIMD[DType.bfloat16, SIMDLength(8)]): This lane's 8-element A fragment (16x16 tile).
b_frag (SIMD[DType.bfloat16, SIMDLength(16)]): This lane's 16-element B fragment (op transpose_right layout; elements 0-7 = N 0..15, 8-15 = N 16..31).
c_acc (SIMD[DType.float32, SIMDLength(16)]): This lane's 16-element C accumulator, updated in place.