# GPU Glossary but for Apple Silicon GPUs > Source: > Published: 2026-08-25 14:33:25+00:00 # 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↗](https://modal.com/gpu-glossary/readme), 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](/machine/gpu-core/)- the SM analogue, and the numbers that differ[Simdgroup](/machine/simdgroup/)- the warp, renamed[Registers](/machine/registers/)- ~208 KB per core, the real budget, and the 10× spill cliff[Threadgroup Memory](/machine/threadgroup-memory/)- 32 KB; shared memory demoted to staging buffer[Unified Memory](/machine/unified-memory/)- no transfers, modest bandwidth, and the platform's one law[Occupancy](/machine/occupancy/)- saturation at ~24 simdgroups, and occupancy as a currency[F16](/machine/f16/)- faster for stall and register reasons, not throughput reasons[Special Paths](/machine/special-paths/)- the fast exp2; the emulated float atomics[AMX](/machine/amx/)- the other matrix engine on the die[Neural Accelerators](/machine/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](/metal/metal-the-api/)- orientation and the full mapping table[MSL](/metal/msl/)- Metal Shading Language, as a CUDA C++ accent[Dispatch Geometry](/metal/dispatch-geometry/)- grids and threadgroups, plus two API traps[Compilation Pipeline](/metal/compilation-pipeline/)- MSL → AIR → metallib → pipeline state[Function Constants](/metal/function-constants/)- specialization without a kernel explosion[Command Buffers](/metal/command-buffers/)- the batching model, and why it's free performance[Synchronization](/metal/synchronization/)- barriers cost ~2 cycles; design accordingly[simdgroup_matrix](/metal/simdgroup-matrix/)- the tensor-core analogue that isn't one[simdgroup_async_copy](/metal/simdgroup-async-copy/)- the dead`cp.async` , and why kernels look the way they do[MTLTensor and MPP](/metal/mtltensor-and-mpp/)- Metal 4's tensor model: descriptors, cooperative tensors[MPS](/metal/mps/)- the cuBLAS analogue, and why it's beatable[Profiling](/metal/profiling/)- the honest page: there is no Nsight (yet)[Disassembly](/metal/disassembly/)- reading what the compiler actually emitted ## MLX *The framework layer, where CUDA-land has no single equivalent.* [MLX, an Overview](/mlx/mlx-overview/)- arrays with no device, and the layer map[Lazy Evaluation](/mlx/lazy-evaluation/)- the graph,`mx.eval` , and the one-eval-per-step rule[How an Op Becomes a Kernel](/mlx/how-an-op-becomes-a-kernel/)- the readable dispatcher, tile tables and all[Steel](/mlx/steel/)- the CUTLASS of Apple Silicon[mx.fast](/mlx/mx-fast/)- fused ops, the SDPA dispatch gate, and the custom-kernel escape hatch[Quantization](/mlx/quantization/)- group-wise affine, QMV/QMM, and bandwidth arithmetic[mx.compile](/mlx/mx-compile/)- elementwise fusion, deliberately scoped[Distributed](/mlx/distributed/)- Mac clusters over Thunderbolt 5 RDMA ## Techniques *Why fast kernels are shaped the way they are.* [Arithmetic Intensity](/techniques/arithmetic-intensity/)- the number that decides everything here[Roofline](/techniques/roofline/)- build your own; no profiler will do it for you[Tiling](/techniques/tiling/)- the reuse pyramid, Apple proportions[Cooperative Load](/techniques/cooperative-load/)- how tiles move, now that DMA is gone[Register Blocking](/techniques/register-blocking/)- the accumulator grid and its cliff edges[Double Buffering](/techniques/double-buffering/)- overlap via ILP, and a trade that flips sign[Fusion and Epilogues](/techniques/fusion-and-epilogues/)- the platform's most profitable technique[Online Softmax](/techniques/online-softmax/)- the five-line algorithm behind flash attention[Flash Attention](/techniques/flash-attention/)- never materialize the score matrix[The KV Cache](/techniques/kv-cache/)- the structure decode streams, and its K/V asymmetry[Decode vs Prefill](/techniques/decode-vs-prefill/)- one API, two workloads, separate kernels ## Kernels: the Case Studies *Real production code, quoted and read.* [The Tiled GEMM](/kernels/gemm-tiled/)- 140 lines that beat MPS (m5-gemm)[The Double-Buffered GEMM](/kernels/gemm-double-buffered/)- and the benchmark table with three winners[The Async-Copy Ghost](/kernels/gemm-async-ghost/)- the DMA original, and its death on Metal 4[Steel's BlockLoader](/kernels/steel-blockloader/)- the cooperative load as a template[Steel's BlockMMA](/kernels/steel-blockmma/)- register blocking as a component[The Fused GEMM Kernel](/kernels/steel-gemm-fused/)- when the kernel is just wiring[Steel Attention](/kernels/steel-attention/)- the whole flash-attention algorithm, readable[metal-flash-attention](/kernels/mfa-codegen/)- the codegen school[llama.cpp Attention](/kernels/llamacpp-attention/)- the enumeration school, quantized to the bone[The NAX GEMM](/kernels/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](/war-stories/three-questions/)- ask before writing any kernel[Sparse-V](/war-stories/sparse-v/)- the canonical win: +22.8% from three lines[The Failures](/war-stories/the-failures/)- published negative results, and the checklist they teach[Cheap Tricks](/war-stories/cheap-tricks/)- sysctls, spec sheets, and fans ## Hands-on The glossary is reading; if you want your hands on a keyboard: [Metal-Puzzles↗](https://github.com/abeleinin/Metal-Puzzles) (14 progressive MSL kernels, checked from Python) is the on-ramp, and [m5-gemm↗](https://github.com/yaroslavvb/m5-gemm) runs the [GEMM case studies'](/kernels/gemm-tiled/) 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↗](https://github.com/Exorust/metalworking/blob/main/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↗](https://github.com/Exorust/metalworking/blob/main/SOURCES.md).