A hyperlinked glossary of Apple Metal GPU performance: the machine, the Metal stack, MLX's architecture, the techniques, and real production kernels read line by line.
Who it's for: you know GPU fundamentals, say at the level of Modal's GPU Glossary↗, but you've never worked with Metal or MLX. This is the companion volume for the other hardware: every page opens with the CUDA equivalent where one exists, then spends its words on what's different here. Concepts Modal already covers well are linked, not re-taught.
How to read it: like a glossary or like a book. Every page stands alone and
links to what it assumes, and every page ends with a Next
link, so front to back also works. Front to back is a deliberate arc: what is this machine → how do you talk to it → what does the framework do on your behalf → why are fast kernels shaped this way → read five real ones → learn from the community's scars.
All code excerpts are quoted verbatim from pinned commits and permalink back to
their exact lines, so you can always zoom out from a fragment to the full source.
No setup is needed to read; ./fetch.sh
(optional, ~240 MB) clones every referenced repo at its pinned commit for browsing beyond the excerpts.
The Machine #
The M-series GPU, as a diff against the GPU you already know.
GPU Core- the SM analogue, and the numbers that differSimdgroup- the warp, renamedRegisters- ~208 KB per core, the real budget, and the 10× spill cliffThreadgroup Memory- 32 KB; shared memory demoted to staging bufferUnified Memory- no transfers, modest bandwidth, and the platform's one lawOccupancy- saturation at ~24 simdgroups, and occupancy as a currencyF16- faster for stall and register reasons, not throughput reasonsSpecial Paths- the fast exp2; the emulated float atomicsAMX- the other matrix engine on the dieNeural Accelerators- the M5's real tensor unit, and what it changes
Metal, the Stack #
What replaces the CUDA driver, runtime, nvcc, and PTX.
Metal, the API- orientation and the full mapping tableMSL- Metal Shading Language, as a CUDA C++ accentDispatch Geometry- grids and threadgroups, plus two API trapsCompilation Pipeline- MSL → AIR → metallib → pipeline stateFunction Constants- specialization without a kernel explosionCommand Buffers- the batching model, and why it's free performanceSynchronization- barriers cost ~2 cycles; design accordinglysimdgroup_matrix- the tensor-core analogue that isn't onesimdgroup_async_copy- the deadcp.async
, and why kernels look the way they doMTLTensor and MPP- Metal 4's tensor model: descriptors, cooperative tensorsMPS- the cuBLAS analogue, and why it's beatableProfiling- the honest page: there is no Nsight (yet)Disassembly- reading what the compiler actually emitted
MLX #
The framework layer, where CUDA-land has no single equivalent.
MLX, an Overview- arrays with no device, and the layer mapLazy Evaluation- the graph,mx.eval
, and the one-eval-per-step ruleHow an Op Becomes a Kernel- the readable dispatcher, tile tables and allSteel- the CUTLASS of Apple Siliconmx.fast- fused ops, the SDPA dispatch gate, and the custom-kernel escape hatchQuantization- group-wise affine, QMV/QMM, and bandwidth arithmeticmx.compile- elementwise fusion, deliberately scopedDistributed- Mac clusters over Thunderbolt 5 RDMA
Techniques #
Why fast kernels are shaped the way they are.
Arithmetic Intensity- the number that decides everything hereRoofline- build your own; no profiler will do it for youTiling- the reuse pyramid, Apple proportionsCooperative Load- how tiles move, now that DMA is goneRegister Blocking- the accumulator grid and its cliff edgesDouble Buffering- overlap via ILP, and a trade that flips signFusion and Epilogues- the platform's most profitable techniqueOnline Softmax- the five-line algorithm behind flash attentionFlash Attention- never materialize the score matrixThe KV Cache- the structure decode streams, and its K/V asymmetryDecode vs Prefill- one API, two workloads, separate kernels
Kernels: the Case Studies #
Real production code, quoted and read.
The Tiled GEMM- 140 lines that beat MPS (m5-gemm)The Double-Buffered GEMM- and the benchmark table with three winnersThe Async-Copy Ghost- the DMA original, and its death on Metal 4Steel's Block- the cooperative load as a templateSteel's BlockMMA- register blocking as a componentThe Fused GEMM Kernel- when the kernel is just wiringSteel Attention- the whole flash-attention algorithm, readablemetal-flash-attention- the codegen schoolllama.cpp Attention- the enumeration school, quantized to the boneThe NAX GEMM- the fourth school: describe the matmul, let the hardware place it
War Stories #
What happened when real people fought the defaults.
The Three Questions- ask before writing any kernelSparse-V- the canonical win: +22.8% from three linesThe Failures- published negative results, and the checklist they teachCheap Tricks- sysctls, spec sheets, and fans
Hands-on #
The glossary is reading; if you want your hands on a keyboard: Metal-Puzzles↗ (14 progressive MSL kernels, checked from Python) is the on-ramp, and m5-gemm↗ runs the GEMM case studies' benchmarks on your own machine with nothing but Python installed.
Layout #
glossary/ the six sections; start anywhere, or at the top
code/ fetched third-party repos (created by ./fetch.sh, gitignored, optional)
SOURCES.md the full annotated source list: repos, blogs, papers, people to follow
fetch.sh reproduces code/ at the exact pinned commits
Licensing #
The glossary text and fetch.sh
are MIT (LICENSE↗). Code excerpts are quoted from their original repositories under their own licenses (MIT and BSD-3-Clause throughout), verbatim, with attribution and a permalink at every quotation site. Full attribution in SOURCES.md↗.