Neural Accelerators

The M5 generation adds a dedicated matrix unit to every GPU core: the neural accelerators (NAX in MLX's source), reached through Metal 4 tensor ops rather than simdgroup_matrix.

CUDA equivalent: the Tensor Core, finally for real. Earlier pages tell the pre-M5 story on purpose: simdgroup_matrix executes on the regular FP32 pipes, and the reverse-engineered Rigel↗ work confirmed no dedicated matrix datapath as late as M4. On M5 that changed. Apple's own numbers for MLX on M5: ~3.97× time-to-first-token and 1.19-1.27× decode over M4↗, with the prefill gain landing exactly where a matrix unit should land (compute-bound territory; decode stays bandwidth-bound, so it moves with the memory spec instead).

What a CUDA person should map over, and what not to:

• Fragment shape moves from 8×8 to 16×16. MLX's NAX steel is built on a 16×16 fragment held across the 32 threads:

struct BaseNAXFrag {
  STEEL_CONST short kFragRows = 16;
  STEEL_CONST short kFragCols = 16;

  STEEL_CONST short kElemsPerFrag = (kFragRows * kFragCols) / 32;

— MLX steel/gemm/nax.h:27-31↗, which includes <MetalPerformancePrimitives/MetalPerformancePrimitives.h> at the top of the file: the new hardware arrives through a library, not new intrinsics.

• The formats are block-scaled floats. The mxfp4/nvfp4/mxfp8 quantization modes exist because this unit consumes them natively; the fp8-is-emulated caveat is an M4-era fact.
• It is not programmed like a tensor core. No mma.sync descent: you describe a matmul with a descriptor and cooperative tensors, and the compiler and hardware pick the datapath. The Metal 4 tensor-ops page covers the model; the NAX GEMM case study reads real code.

Ecosystem status, honestly: MLX ships a parallel steel fork for it (gemm/nax.h, attn/nax.h, quantized_nax, fp_quantized_nax); llama.cpp has a tensor-API matmul path; and the tuning is young. Recent MLX work is disproportionately NAX-shaped (qmv batch limits raised for M5-class GPUs, NVFP4 QMV optimization, per-expert tile picks in gather_qmm_rhs_nax), which is what a fast-moving fast path looks like. simdgroup_matrix remains the portable primitive and the one most shipped kernels still use; this page is why the glossary keeps teaching both.