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 module

int8_matmul

Apple M5 int8 W8A8 matmul: out = dequant(A_int8 @ B_int8^T).

Apple silicon GPU (Metal 4, compute_capability == 5). The int32-accumulate sibling of the W4A16 kernel (matmul2d_fp4.mojo / fp4_matmul.mojo): where the FP4 path decodes a packed 4-bit weight to bf16 and feeds the float simdgroup MMA, this path feeds int8 A and int8 B straight to the M5 integer widening simdgroup MMA (_mma_apple with int8 inputs -> int32 accumulator, emitting air.simdgroup_matrix_16x16x16_widening_multiply_accumulate). That widening op is a genuinely faster M5 datapath than the float MMA (~1.7-1.9x at the FLUX.2 projection/MLP shapes; the dequant epilogue nets ~1.5-1.7x). The dense AppleM5MatMul cannot be reused: it hardcodes a float accumulator (MmaOpApple[float32, in_type]) and asserts a float c_type.

Quantization (W8A8): both operands are symmetric-absmax int8 (scale = absmax / 127), A per-row (per-token), B per-column (per-output-channel). The kernel accumulates the raw int8xint8 products in int32, then the epilogue DEQUANTIZES each output element by a_scale[row] * b_scale[col] and casts to c_type (bf16), with an optional per-column bias added after dequant. A is quantized online by enqueue_apple_int8_quantize_activation (a separate one-pass kernel, reusing the _quantize_a_block-style absmax); B is pre-quantized at weight load.

Structure mirrors AppleM5MatMul._run_gemm_body: 64x64 threadgroup block, 4 simdgroups (128 threads), each a 32x32 (2x2 MmaOpApple) subtile; rectangular-Morton tile schedule (wide-N shapes are Morton-critical -- a flat grid measured 50-70 Tops/s vs 96-99 with Morton); DRAM->register loads (no SMEM staging -- the M5 rule). The K-loop feeds the widening MMA per BK=64 strip: the K % 16 == 0 interior takes a width-16 int8 K-repartition; the K-tail and K % 16 != 0 strips take an AGX3 edge-masked width-4 path; genuinely ragged M/N tiles take the bounded path. See KB kernels/apple-m5-int8-matmul.

Structs

AppleInt8ActQuant: Per-row symmetric-absmax int8 quantization of the activation.
AppleM5Int8MatMul: W8A8 GEMM: int8 A x int8 B^T -> int32 accum -> per-row/col dequant.
Int8DequantWriter: Owns the C-output write of the W8A8 epilogue (the write-side counterpart to the FP4 Fp4WeightLoader).

Functions

enqueue_apple_int8_matmul: Enqueue the W8A8 matmul: out = dequant(a @ b^T).
enqueue_apple_int8_quantize_activation: Quantize a (M, K) to int8 q + per-row fp32 a_scale (M,).

Structs​

Functions​

Structs

Functions