GPU Glossary but for Apple Silicon GPUs

Apple's Metal GPU glossary 'metalworking' targets developers familiar with CUDA but new to Apple Silicon, mapping every concept from GPU cores to MLX's architecture with CUDA equivalents and real kernel code. The hyperlinked guide covers the M-series GPU, Metal API, MLX framework, and optimization techniques, aiming to bridge the gap for performance engineers.

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 .