metalworking

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 differ
• Simdgroup - the warp, renamed
• Registers - ~208 KB per core, the real budget, and the 10× spill cliff
• Threadgroup Memory - 32 KB; shared memory demoted to staging buffer
• Unified Memory - no transfers, modest bandwidth, and the platform's one law
• Occupancy - saturation at ~24 simdgroups, and occupancy as a currency
• F16 - faster for stall and register reasons, not throughput reasons
• Special Paths - the fast exp2; the emulated float atomics
• AMX - the other matrix engine on the die
• Neural 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 table
• MSL - Metal Shading Language, as a CUDA C++ accent
• Dispatch Geometry - grids and threadgroups, plus two API traps
• Compilation Pipeline - MSL → AIR → metallib → pipeline state
• Function Constants - specialization without a kernel explosion
• Command Buffers - the batching model, and why it's free performance
• Synchronization - barriers cost ~2 cycles; design accordingly
• simdgroup_matrix - the tensor-core analogue that isn't one
• simdgroup_async_copy - the dead cp.async, and why kernels look the way they do
• MTLTensor and MPP - Metal 4's tensor model: descriptors, cooperative tensors
• MPS - the cuBLAS analogue, and why it's beatable
• Profiling - 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 map
• Lazy Evaluation - the graph, mx.eval, and the one-eval-per-step rule
• How an Op Becomes a Kernel - the readable dispatcher, tile tables and all
• Steel - the CUTLASS of Apple Silicon
• mx.fast - fused ops, the SDPA dispatch gate, and the custom-kernel escape hatch
• Quantization - group-wise affine, QMV/QMM, and bandwidth arithmetic
• mx.compile - elementwise fusion, deliberately scoped
• Distributed - Mac clusters over Thunderbolt 5 RDMA

Techniques

Why fast kernels are shaped the way they are.

• Arithmetic Intensity - the number that decides everything here
• Roofline - build your own; no profiler will do it for you
• Tiling - the reuse pyramid, Apple proportions
• Cooperative Load - how tiles move, now that DMA is gone
• Register Blocking - the accumulator grid and its cliff edges
• Double Buffering - overlap via ILP, and a trade that flips sign
• Fusion and Epilogues - the platform's most profitable technique
• Online Softmax - the five-line algorithm behind flash attention
• Flash Attention - never materialize the score matrix
• The KV Cache - the structure decode streams, and its K/V asymmetry
• Decode 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 winners
• The Async-Copy Ghost - the DMA original, and its death on Metal 4
• Steel's BlockLoader - the cooperative load as a template
• Steel's BlockMMA - register blocking as a component
• The Fused GEMM Kernel - when the kernel is just wiring
• Steel Attention - the whole flash-attention algorithm, readable
• metal-flash-attention - the codegen school
• llama.cpp Attention - the enumeration school, quantized to the bone
• The 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 kernel
• Sparse-V - the canonical win: +22.8% from three lines
• The Failures - published negative results, and the checklist they teach
• Cheap 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↗.