learn-metal

Learning Apple Metal from Rust — GPU compute and 3D graphics, from first principles.

Not an engine, not a library. A progression of small, self-contained programs, each isolating one GPU concept, each verifying its own correctness.

cargo run --bin compute_01_saxpy

If that prints ok, you're set up. See docs/00-setup.md if it doesn't.

---

Who this is for

Someone who wrote OpenGL and DirectX 11 back when DrawIndexed() looked like it drew something immediately, and wants to come back to a field where the pipeline is explicit, the GPU is programmable general-purpose hardware, and graphics and machine learning are converging on the same silicon.

The shaders, buffers, and textures you remember are all still here. The submission model is what changed, and it's the first thing these lessons teach.

---

How it works

Each lesson is a folder containing everything it needs — the Rust, the .metal shaders, and a README explaining the concept:

lessons/compute/01_saxpy/
├── README.md      what this teaches, what to notice, what to try
├── main.rs        heavily commented — the comments are the tutorial
└── saxpy.metal    the shader

Every lesson is complete and runs. No exercises, no todo!(), nothing stubbed out. Clone it, run it, read it, change it. The comments in the code and the lesson README are the tutorial.

Where a lesson has a genuine design fork (dispatch style, storage mode, sync primitive), it demonstrates every branch and prints the difference rather than picking one silently — so you see the trade-off happen instead of reading an assertion about it.

Every lesson checks its own answer against a CPU reference. On a GPU, "silently wrong" is the default failure mode — an out-of-bounds read doesn't segfault, it just hands you garbage — so a wrong result must never be able to masquerade as success.

Lessons are also honest about measurement: they warm up before timing, and they report actual GPU execution time (GPUEndTime - GPUStartTime) rather than wall clock, which includes your own CPU work and will happily lie to you about which kernel is faster.

---

Reference docs

Read these when the lessons reference them, or up front if you prefer theory first.

The Submission Model	What command queues actually solve. Why DX11's DrawIndexed lied to you, the three distinct kinds of synchronization, and what modern APIs actually made faster (it isn't shaders). Start here.
TBDR vs. Immediate Mode	Apple GPU architecture vs. NVIDIA/AMD, and why your desktop instincts will mislead you. Hidden surface removal, tile memory, why front-to-back sorting is pointless here, why MSAA is nearly free, and the one call that destroys the whole architecture. The most important doc for an ex-DX11 engineer.
Glossary	Translation tables: OpenGL/DX11 → Metal, CUDA → Metal, and the sync primitives across Metal/Vulkan/D3D12. Plus the objc2-metal idioms that look alarming and aren't.
Setup	Toolchain, and the two genuinely nasty Xcode gotchas — including why xcrun metal insists the compiler is missing when it is right there.
Resources	Where to learn from the internet. Verified July 2026, stale stuff flagged. Also: why the Rust+Metal ground is thin, and what to do about it.

---

The curriculum

Five tracks, 58 lessons. Roughly in order, but the tracks are independent after compute/04 — jump around once you have the submission model.

Legend: ✅ done · ⭐ don't skip this one · ⚠️ hardware caveat

---

Compute — start here (18)

Compute teaches the entire submission model with zero windowing ceremony in the way. No swap chain, no drawable, no window server, no event loop. Just: get work onto the GPU, get the answer back, check it against the CPU. Every concept here reappears in graphics, so you learn it once in the setting with the fewest moving parts.

	Lesson	What it teaches
✅	01_saxpy	Device, queue, command buffer, encoder, dispatch, pipeline state. Unified memory. Threadgroups vs. non-uniform threadgroups. Honest GPU timing.
	02_buffers_and_storage	StorageModeShared / Private / Memoryless, newBufferWithBytesNoCopy (wrap a Vec with zero copy), alignment, didModifyRange. What unified memory does and does not give you for free.
	03_image_kernels	2-D grids. MTLTexture in compute, read()/write(), sampler access. Box blur, Sobel edge detect. Why a 16×16 threadgroup beats 256×1 for image work (cache locality).
	04_threadgroup_memory	⭐ On-chip scratchpad — CUDA's __shared__. threadgroup_barrier. Separable blur: load a tile once, read it 9× from on-chip memory instead of 9× from DRAM. The single most important GPU optimization there is.
	05_frames_in_flight	⭐ The CPU/GPU timeline. Triple buffering, dispatch_semaphore, addCompletedHandler. Why waitUntilCompleted in a loop halves your throughput. The lesson that answers "what did command queues actually solve?"
	06_metal4_model	⭐ The modern API. MTL4CommandQueue, MTL4CommandAllocator, MTL4ArgumentTable, MTL4Compiler, residency sets. Lesson 01 rebuilt, side by side. Sits here on purpose: Metal 4 makes you own the memory commands are recorded into, and you must not reset it while the GPU is still executing from it — which is exactly the hazard you just met in lesson 05, applied to command memory. See why we don't start here.
	07_simd_groups	The warp, up close. simd_sum, simd_shuffle, simd_ballot, simd_prefix_exclusive_sum. Lockstep execution, and what divergence actually costs.
	08_reduction	Sum a million floats. Naive → threadgroup tree → SIMD intrinsics → two-pass. Each version measured. Where "just parallelize it" stops being obvious.
	09_scan	Prefix sum. The algorithm that looks inherently sequential and isn't. Hillis-Steele vs. Blelloch, work-efficiency vs. depth.
	10_atomics_histogram	atomic_uint, atomic_fetch_add_explicit, memory orderings. Contention: watch a histogram collapse when 10,000 threads hammer one bin, then fix it with per-threadgroup privatization.
	11_sorting	Bitonic sort, then a radix sort built on your own scan from lesson 09. Multi-pass compute, ping-pong buffers. And the crossover point where the CPU still wins.
	12_matmul_naive	The baseline. Correct, simple, and leaving ~95% of the hardware on the table. Roofline analysis: is this compute-bound or memory-bound? (It's memory-bound, and the number proves it.)
	13_matmul_tiled	⭐ Threadgroup-memory tiling, register blocking, occupancy vs. register pressure. The classic GPU optimization story, start to finish, each step measured against the last.
	14_simdgroup_matrix	simdgroup_float8x8 — Apple's hardware matrix units. The step from "good tiled matmul" to "within reach of MPS". Where the real FLOPs live.
	15_function_constants	[[function_constant(n)]] — specialize one shader source into many pipeline variants without #ifdef hell or runtime branches. How real engines manage shader permutations.
	16_indirect_dispatch	dispatchThreadgroupsWithIndirectBuffer. The GPU decides its own workload — no CPU readback, no stall. Foundation for GPU-driven rendering.
	17_heaps_and_aliasing	MTLHeap, sub-allocation, memory aliasing for transient resources. Manual memory management, and an honest answer on when it's worth it.
	18_aot_metallib	Ahead-of-time shader compilation: xcrun metal → .air → .metallib, via build.rs. Binary archives and pipeline caching. Kill the first-frame hitch.

---

Graphics (21)

	Lesson	What it teaches
	01_offscreen_triangle	Render pipeline state, MSL vertex + fragment shaders, MTLRenderPassDescriptor. Renders to a texture, writes a PNG. Still no window — the render pipeline without the windowing distraction.
	02_window	winit + CAMetalLayer. The drawable/present cycle, presentDrawable, resize handling. A live triangle at last.
	03_vertex_buffers	MTLVertexDescriptor, interleaved vs. planar layouts, [[stage_in]]. Sharing struct definitions between Rust and MSL without them silently drifting apart.
	04_uniforms_and_camera	MVP matrices, glam, the uniform buffer. Row-major vs. column-major, and the depth-range difference (Metal is [0,1], OpenGL was [-1,1]) that will bite you exactly once.
	05_indexed_cube	Index buffers, depth testing, MTLDepthStencilState, winding order, back-face culling. A cube that actually looks like a cube.
	06_textures_samplers	MTLTexture, MTLSamplerState, filtering, wrap modes, mipmaps, sRGB (and why getting gamma wrong makes everything look muddy).
	07_tbdr_and_tile_memory	⭐⭐ The most important graphics lesson here. Apple GPUs are tile-based deferred renderers. Load/store actions, memoryless attachments, hidden surface removal. Your DX11 instincts are actively wrong here — this is where you unlearn them, with numbers.
	08_lighting	Blinn-Phong. Normals, the normal matrix, and why non-uniform scale breaks your lighting.
	09_normal_mapping	Tangent space, TBN matrix, and how to make a flat wall look like brick.
	10_pbr	Cook-Torrance. Metallic/roughness workflow, the NDF/geometry/Fresnel terms, energy conservation.
	11_ibl	Image-based lighting. Cubemaps, irradiance convolution, prefiltered specular, the BRDF LUT — all precomputed with compute kernels from the compute track.
	12_shadow_mapping	Two render passes, one command buffer. Depth bias, peter-panning, PCF, cascades. The classic "why is my shadow acne everywhere" tour.
	13_instancing	[[instance_id]], per-instance buffers. 100,000 cubes in one draw call.
	14_msaa	Multisampling, and the TBDR twist: on tile-based hardware you resolve inside tile memory and never write the multisampled buffer to DRAM at all. Nearly free antialiasing.
	15_deferred_shading	G-buffer, light accumulation. Then the Apple version: imageblocks and tile shading, keeping the entire G-buffer in tile memory and never touching DRAM. This is the payoff for lesson 07.
	16_post_processing	Bloom, tone mapping, HDR. Compute shaders operating on render targets — the two tracks meeting.
	17_compute_skinning	Skeletal animation. Compute skins the mesh, the vertex shader renders it.
	18_tessellation	Tessellation shaders, displacement mapping, LOD.
	19_mesh_shaders	Object + mesh shaders. The modern replacement for the vertex/geometry pipeline.
	20_argument_buffers	Bindless. Thousands of textures addressable from a shader without rebinding. Residency sets, useResource.
	21_metalfx	Temporal upscaling. Render at 60%, present at 100%.

---

Hybrid — compute and graphics in one frame (7)

Where it gets interesting, and where Metal's automatic hazard tracking earns its keep.

	Lesson	What it teaches
	01_gpu_particles	⭐ Compute writes a position buffer, the vertex shader reads it — in the same command buffer. Cross-encoder dependencies, and how Metal inserts the barrier that Vulkan would make you write by hand.
	02_explicit_barriers	Now turn the automatic tracking off. MTLFence, memoryBarrier, untracked resources, hazardTrackingMode. What Metal was doing for you, and what it costs.
	03_gpu_frustum_culling	Compute culls, then builds the draw list. The CPU never learns how many objects survived.
	04_indirect_command_buffers	GPU-driven rendering. The GPU encodes its own draw calls. MTLIndirectCommandBuffer, executeCommandsInBuffer.
	05_multiple_queues	Two MTLCommandQueues, MTLEvent between them. Async compute — and an honest measurement of whether it actually helped (often: no).
	06_ray_tracing	MTLAccelerationStructure, intersection functions, ray-traced shadows and reflections.
	07_path_tracer	A real one. Progressive accumulation, importance sampling, denoising. The capstone.

---

ML — where this is all going (7)

	Lesson	What it teaches
	01_mlx_custom_kernel	mx.fast.metal_kernel() — write MSL, call it from Python, zero host boilerplate. The fastest possible loop for learning MSL compute semantics.
	02_rmsnorm_softmax	Implement them in MLX primitives, then as one fused Metal kernel. Benchmark both. Watch memory traffic collapse.
	03_quantized_gemv	4-bit weights, the kernel that actually makes local LLM inference fast. Unpacking, dequant in registers, bandwidth as the whole ballgame.
	04_flash_attention	⭐ Tiled attention, online softmax. The algorithm that made long contexts affordable — built on the tiling and reduction skills from the compute track.
	05_kv_cache	KV-cache update kernels, paged attention. Where inference actually spends its time.
	06_metal4_tensors	MTLTensor, MTL4MachineLearningCommandEncoder — ML inference scheduled on the GPU timeline, interleaved with rendering. ⚠️ In-shader Shader ML (TensorOps) needs an M5; it will not run on your M4.
	07_coreml_and_ane	The Neural Engine. Why you can't write ANE kernels, how Core ML decides CPU vs. GPU vs. ANE, and how to measure where your model actually ran.

---

Tooling — a track, not an afterthought (5)

Most people never learn the GPU debugger properly and pay for it forever.

	Lesson	What it teaches
	01_frame_capture	⭐ Xcode GPU frame capture on a Rust process (it works — you capture the process, not the language). Inspect every resource, every draw, every pipeline. This is RenderDoc/PIX, and it's better.
	02_shader_profiler	Per-line shader cost. Find the one pow() eating your frame.
	03_gpu_counters	Occupancy, bandwidth, ALU utilization, stall reasons. Reading a roofline and knowing whether to optimize memory or math.
	04_validation	Metal API validation, shader validation, MTL_DEBUG_LAYER. Catching the silent out-of-bounds write from compute/01 before it corrupts your data.
	05_perf_hud	Metal Performance HUD, metalperftrace, Instruments. Profiling without a debugger attached.

---

Suggested path

You don't need to go in order after the first five.

1. compute/01 → compute/05. Non-negotiable. This is the submission model, and everything else assumes it.
2. graphics/07 (TBDR) — read it early even if you skip ahead to get there. It reframes every graphics lesson that follows, and it's where ex-desktop-GL engineers leave the most performance on the floor.
3. tooling/01 (frame capture) — as soon as you have something on screen. Debugging blind is a choice.
4. Then follow whatever you're actually curious about. compute/12 (tiled matmul) and hybrid/01 (GPU particles) are the two most satisfying single lessons in the list.

---

Why this stack

	Choice	Why
CPU-side API	objc2-metal	Direct, complete, current. Not metal-rs — deprecated by its own maintainers, who now point here.
Shaders	Metal Shading Language	The real thing. C++-based; immediately familiar after GLSL/HLSL.
Windowing	winit + CAMetalLayer	Introduced at graphics/02, not before.

Deliberately not wgpu. It's a fine library and it compiles down to Metal — but the entire point here is to see the machinery, not an abstraction over it. If you don't know what's under the abstraction, you'll spend your life wondering what it's hiding.

The cost of objc2-metal is verbosity and some unsafe, because the API is Objective-C underneath. That's the price of admission, and it's cheaper than it looks — see the glossary.

---

Hardware assumption

Apple Silicon (developed on an M4 Max, 40 GPU cores, unified memory, Metal 4). No discrete-GPU code paths, no staging buffers, no explicit host↔device copies — because on this hardware there is nothing to copy.