metalworkingGitHub ↗
/machine/amx

AMX

AMX is Apple's undocumented CPU matrix coprocessor. Not part of the GPU at all, but the other half of the Apple matmul story and the answer when small-matrix CPU throughput seems impossible.

CUDA equivalent: none. The closest analogy is Intel's AMX or ARM's SME (into which Apple's AMX has been evolving), but the important point for a GPU person is negative: when you benchmark Apple's Accelerate/BNNS/MPS small-matrix performance and the numbers look too good for a CPU, you're not looking at the CPU's NEON units. You're looking at a dedicated 32×32 multiply-accumulate grid hanging off the CPU cluster, fed by its own X/Y register pools, executing an outer product per instruction.

Why it appears in a GPU glossary:

  • Dispatch decisions. For small matrices, the fastest processor on the chip is often AMX, not the GPU: kernel launch overhead plus bandwidth round trips swamp the GPU's advantage below a size threshold. Apple's own libraries route accordingly. If your profiling shows MPS beating your GPU kernel at 256×256, this is frequently why.
  • It shares the memory pool. Unified memory means CPU-side AMX work and GPU work compete for the same bandwidth; a background BNNS workload dents your GPU kernel's roofline.

ONE DIE, THREE MATRIX ENGINES CPU cluster AMX 32×32 MAC grid outer product / instr wins small matrices: no launch overhead, no round trips GPU · N cores simdgroup_matrix this glossary's territory Neural Engine (ANE) ~6.6 TFLOPS/watt (vs ~1 for the GPU) quantized-conv-shaped inference at low power Apple's frameworks may route here without telling you unified memory fabric · one pool, one bandwidth budget: the three engines compete for it the API that reaches each: Accelerate / BNNS (sanctioned) · corsix/amx (docs) Metal · MSL · MLX Core ML only · private APIs otherwise

Where AMX sits. Three matrix engines share one bandwidth budget; Metal talks only to the middle one, and the API rows at the bottom are the whole accessibility story.

Apple has never documented the instruction set. The definitive reference is corsix/amx↗ (MIT): per-instruction docs with an executable software emulator as the specification, plus per-generation deltas (M2 adds bf16; M3 adds new load/int modes; M4-era hardware moves toward standard SME). Using AMX directly means issuing undocumented instructions, and Accelerate/BNNS is the sanctioned interface. But reading corsix's docs is the only way to understand the performance you observe through those libraries.

The ANE war story is the same genre one chip region over: the Neural Engine, also undocumented, also reachable only through private interfaces, at ~6.6 TFLOPS/watt. The GPU is one of three matrix engines on this die, and Metal only talks to one of them.

Edit this page on GitHub ↗