No-shortcut MATH + CS foundations for an Anthropic-style kernel / training-systems engineer

Date: 2026-08-18 Goal: be able to sit in a room with people who write FlashAttention-class kernels, argue about Muon vs SOAP vs AdamW, debug a BF16 GEMM that is silently wrong, and read μP / tensor-parallel papers without bluffing. Method: textbooks, problem sets, implementations. Videos are secondary. 3Blue1Brown / popular-science is a weekend supplement, never the path.

Deep hours below mean: read the chapter, do the problems with a pen, implement the non-trivial ones, and be able to re-derive the statement the next week. Watching a lecture at 1.5× does not count.

Honest total for both tracks, after skipping what you already own:

At 20 h/week that is ~4 years. At 30 h/week ~2.6 years. At 40 h/week (full-time study) ~2 years. Anyone quoting “6 months of fundamentals” is selling a shortcut.

How to use this

1. Pick a track at the end. Do not mix tracks week-to-week.
2. A module is not done when you have “read the book.” It is done when you pass the ready-when bar.
3. Write proofs and kernels in a lab notebook / repo. If you cannot reconstruct it from the notebook, you do not know it.
4. After module 5 + CSAPP, start writing tiny CUDA/Triton kernels in parallel with later modules. Theory without a profiler is tourism.

Current (2024–2026) courses kernel people actually mention

These are the applied layer. They assume the foundations below. Do not start here if the math/systems are weak.

Do one of {CS149, 15-418, 6.S894} as a full course, not all three. Then GPU MODE + CS336 A1/A2.

Module 1 — Linear algebra (the real one)

Hard-path pick: Axler, then Trefethen & Bau. Strang is the on-ramp if you cannot yet do Axler exercises. It is not the destination. Axler gives operator structure (what μP, LoRA, and spectral methods actually talk about). Trefethen gives algorithms, conditioning, and the SVD you will implement and debug.

Prerequisites

Single-variable calculus. Comfort with proofs (if not: Hammack Book of Proof, or the first 4 chapters of Abbott first — 40 h, do it).

Primary

1. Sheldon Axler, Linear Algebra Done Right, 4th ed. (2024), Springer, open access. https://linear.axler.net/ New in 4e: expanded SVD, new multilinear / tensor-product chapter. Use that chapter. Do not skip it.
2. Lloyd N. Trefethen & David Bau III, Numerical Linear Algebra, SIAM, 1997 (25th-anniversary reprint 2022). 40 short lectures. Do all of them.

Secondary

• Gilbert Strang, Introduction to Linear Algebra, 6th ed. (2023) or 5th ed. (2016) + MIT 18.06SC — only if Axler is bouncing.
• Gilbert Strang, Linear Algebra and Learning from Data (2019) + 18.065 — after the hard path, for the ML dictionary.
• Horn & Johnson, Matrix Analysis, 2nd ed. (2012) — reference, not a course. Kronecker, norms, variational characterizations.
• Van Loan, “The ubiquitous Kronecker product” (J. Comput. Appl. Math., 2000).

Problem sets (do these, not “some”)

Axler 4e

• Ch 1–3 (vector spaces, linear maps, products): all exercises. This is where most CS people discover they do not know LA.
• Ch 5–8 (eigenvalues, inner products, operators, spectral theorem, SVD): all odd + every exercise that mentions orthonormal, adjoint, or singular values.
• Ch 10 (multilinear / tensors / bilinear forms) in 4e: all of it.

Trefethen & Bau

• Lectures 1–6 (matrix-vector, orthogonality, norms): all exercises.
• Lectures 7–12 (QR, conditioning, stability): all exercises. Lecture 12 (conditioning of least squares / eigenvalues) is non-negotiable.
• Lectures 24–31 (eigenvalue algorithms, SVD): all exercises.
• Lectures 32–40 (iterative methods): selected, but do 32–35 (Arnoldi, GMRES, Lanczos).

Kronecker / layouts (do as a written problem set you invent)

• Prove $(A \otimes B)(C \otimes D) = AC \otimes BD$ and the mixed-product identities.
• Show vec(AXB) = (Bᵀ ⊗ A) vec(X).
• Write the flop count and communication volume of a batched GEMM vs a Kronecker-factored multiply.
• Derive why LoRA’s ΔW = BA is a rank-r factorization, and what the SVD of ΔW tells you about the update.

Estimated hours

You are ready when you can

• State and prove the SVD existence theorem, and compute a thin SVD by hand for a 4×3 matrix.
• Explain, with a condition-number argument, why a least-squares problem with nearly collinear columns is unstable, and what that has to do with a badly scaled attention logit matrix.
• Write vec(AXB) = (Bᵀ ⊗ A) vec(X) on a whiteboard and use it to explain why tensor-parallel all-gather/reduce-scatter layouts look the way they do.
• Given a LoRA pair (B, A), compute the effective rank and the spectral norm of the update, and say whether μP’s “maximal update” is being respected as width grows.
• Implement QR with Householder and compare residual vs a naive Gram-Schmidt in float32 vs float64.

Connection to LLM kernels

• Attention: QKᵀ is a GEMM; softmax is a row-wise nonlinear map whose Jacobian is diag(p) − ppᵀ. SVD/eigen of that Jacobian is how you talk about softmax stability. FlashAttention is a blocked, numerically recomputed version of this GEMM+softmax+GEMM.
• LoRA: low-rank factorization. The whole point is the Eckart–Young theorem (best rank-r approx = truncated SVD).
• μP / tensor programs: width is a dimension; features are vectors in that space; the right scaling of init and LR is a statement about operator norms as width → ∞. If you only know “multiply matrices,” μP looks like numerology.
• Tensor layouts / TP: a weight W ∈ R^{d_out × d_in} sharded on the column or row axis is a Kronecker-structured view of a bigger multiply. 2.5D / 3D matmul (Solomonik–Demmel) is the same idea with extra memory to cut communication.
• Muon / SOAP: polar factor UVᵀ of a gradient matrix is the SVD with singular values set to 1. You cannot understand 2025–2026 optimizer conversation without this.

Module 2 — Multivariable calculus + enough real analysis

Not “watch Essence of Calculus.” You need ε-δ, Taylor with remainder, Lipschitz, and derivatives as linear maps.

Prerequisites

Single-variable calculus to the level of being able to prove the mean value theorem (or do Spivak Calculus 4e selected, +80 h, if you cannot).

Primary

1. Stephen Abbott, Understanding Analysis, 2nd ed. (2015), Springer. Whole book except the last chapter if time-crunched.
2. Matrix calculus notes: MIT 18.S096 — https://github.com/mitmath/matrixcalc Treat derivatives as linear maps (Fréchet). This is the notation used in real papers.
3. Hubbard & Hubbard, Vector Calculus, Linear Algebra, and Differential Forms, 5th ed. or Apostol Calculus Vol. 2, 2nd ed. — for multivariable computational fluency if your calc is rusty.

Secondary

• Rudin, Principles of Mathematical Analysis, 3rd ed. (1976) — Ch 1–7 only, after Abbott. Do not start here.
• Nocedal & Wright, Numerical Optimization, 2nd ed. (2006), Ch 2 (fundamentals) — Taylor, Lipschitz, Newton. Bridges to module 4.
• Magnus & Neudecker, Matrix Differential Calculus (3rd ed., 2019) — reference for Kronecker-form Jacobians.

Problem sets

Abbott 2e

• Ch 2–3 (sequences, series, completeness): all odds + every ε-N problem.
• Ch 4–5 (continuity, differentiation): all. Especially uniform continuity and MVT.
• Ch 6 (sequences of functions, uniform convergence): all. This is how you will later understand “interchange gradient and limit / interchange sum and integral.”

Rudin 3e (after Abbott)

• Ch 2 (topology): selected, enough to not fear open/closed/compact.
• Ch 5 (differentiation), Ch 7 (sequences of functions): selected hard problems. Stop. You do not need Baby Rudin cover-to-cover for this job.
• Skip Rudin’s measure chapter; it is the wrong source.

Matrix calculus (write these out in tensor / index / Fréchet notation)

• Jacobian of x ↦ Wx + b, of X ↦ XW, of X ↦ σ(X) elementwise.
• Hessian of a scalar f(W) = ½‖WX − Y‖²_F.
• Chain rule for ℓ(softmax(XW)) in both index notation and vec notation.
• Jacobian of row-wise softmax; show it is symmetric PSD, rows sum to 0, and J 1 = 0.
• Lipschitz constant of softmax, ReLU, GELU, and of x ↦ x/‖x‖.

Estimated hours

You are ready when you can

• Write a first- and second-order Taylor expansion of a scalar loss L(W) with explicit remainder, and say what “locally Lipschitz gradient” means for SGD step-size.
• Derive backprop for a 2-layer MLP in three notations (index, matrix, Fréchet) and get the same answer.
• Prove that a C¹ function on a compact convex set is Lipschitz, and estimate the constant from ‖∇f‖.
• Read an optimization paper that says “assume L-smooth and μ-strongly convex” and translate that into Hessian eigenvalues without looking it up.

Connection to LLM kernels

• Every backward kernel is a Jacobian-vector or vector-Jacobian product. FlashAttention’s backward is a carefully-recomputed VJP of the blocked softmax.
• LoRA / μP step-size arguments are Taylor + Lipschitz stories about how a parameter update moves the function.
• “Convexity” is rare in deep nets, but local smoothness is why Adam’s ε, gradient clipping, and loss spikes exist. You need the language to not treat them as folklore.

Module 3 — Probability + information theory

Start with a hard undergrad, then Cover & Thomas, then just enough measure theory that “almost surely,” “absolutely continuous,” and “Radon–Nikodym” do not make you close the PDF.

Prerequisites

Module 2 at least through Abbott Ch 2–5. Combinatorics at the “count the ways” level.

Primary

1. Joseph K. Blitzstein & Jessica Hwang, Introduction to Probability, 2nd ed. (2019). Course: Harvard Stat 110 — videos on YouTube, extras at http://stat110.net This is the hard undergrad. Do it. Ross (A First Course in Probability, 10e) is an acceptable substitute if you already own it; Blitzstein has better problems.
2. Thomas M. Cover & Joy A. Thomas, Elements of Information Theory, 2nd ed. (2006), Wiley.

Secondary (measure + concentration)

• Terence Tao, An Introduction to Measure Theory (GSM 126) — Ch 1 and the beginning of Ch 2, or
• Durrett, Probability: Theory and Examples, 5th ed. — Ch 1–2 only.
• Vershynin, High-Dimensional Probability (2018, free PDF) — Ch 2–4 (subgaussian, Hoeffding, Bernstein, Lipschitz concentration).
• Boucheron, Lugosi, Massart, Concentration Inequalities (2013) — reference.

Problem sets

Blitzstein 2e / Stat 110

• Ch 1–5 (counting, conditional, r.v.s, discrete, continuous): all official Stat 110 problem sets. Do not skip story problems.
• Ch 6–9 (moments, joint, transformations, MGFs): all odds + every change-of-variables.
• Ch 10–13 (Markov chains, Markov/Chebyshev/Chernoff, law of large numbers, CLT): all.

Cover & Thomas 2e

• Ch 2 (entropy, KL, mutual information): all exercises. This chapter is the job.
• Ch 3 (AEP / typical sets): all.
• Ch 5 (source coding) selected; Ch 7 (channel capacity) selected — enough to know what a rate is.
• Ch 11 (information theory and statistics) selected — Sanov, hypothesis testing.
• Skip network information theory unless you fall in love.

Concentration (write a 15-page noteset)

• Prove Hoeffding, then Bernstein, then a bounded-difference (McDiarmid) inequality.
• Apply to: (a) average of bounded tokens in a batch; (b) expert-load deviation in a MoE router with noisy gating; (c) why a 1-bit all-reduce of a gradient sign still concentrates.

Measure-theory bridge (do not take a year-long course)

• Define σ-algebra, measure, measurable function, integral, almost-sure vs in-probability.
• State Radon–Nikodym. Recognize that a density is an RN derivative and KL is ∫ p log(p/q).
• That is enough to not be scared. You are not becoming a probabilist.

Estimated hours

You are ready when you can

• Compute H(X), H(X|Y), I(X;Y), D_KL(p‖q) for discrete and simple continuous cases without a formula sheet.
• Prove that KL ≥ 0 from Jensen, and that softmax + cross-entropy is MLE for a categorical.
• Explain typical sets well enough to understand why “most sequences look like the entropy rate” shows up in compression and in why sampling ≠ argmax.
• Write a Hoeffding bound for MoE expert load and say how many tokens you need before a 1% imbalance is surprising.
• Read a paper that writes X_n → X a.s. and not flinch.

Connection to LLM kernels

• Cross-entropy training is KL to the empirical token distribution.
• Softmax numerical issues are probability issues (partition function, log-sum-exp).
• MoE load-balancing losses, auxiliary losses, and expert-choice routing are concentration statements. If you cannot write the bound, you will tune coefficients by folklore.
• Scaling-law “loss vs N, D” fits are statistical estimation; CS336 A3 will punish you if this module is weak.
• μP / feature-learning arguments are high-dimensional probability plus norms.

Module 4 — Optimization

Prerequisites

Modules 1–2. Module 3 at least through expectations and concentration.

Primary

1. Stephen Boyd & Lieven Vandenberghe, Convex Optimization (2004). Free: https://web.stanford.edu/~boyd/cvxbook/ Course: Stanford EE364A — https://ee364a.stanford.edu/ · videos https://see.stanford.edu/Course/EE364A
2. Nocedal & Wright, Numerical Optimization, 2nd ed. (2006) — Ch 2–3, 6–7 (line search, quasi-Newton, CG). The first- and second-order toolkit.

Secondary (deep-learning specific, 2024–2026)

• Bottou, Curtis, Nocedal, “Optimization Methods for Large-Scale Machine Learning” (SIAM Review, 2018).
• Keller Jordan et al., Muon (2024): https://kellerjordan.github.io/posts/muon/
• Jeremy Bernstein, “Deriving Muon” (2025): https://jeremybernste.in/writing/deriving-muon
• Bernstein & Newhouse, “Old Optimizer, New Norm” (2024), arXiv:2409.20325 — SGD / Adam / Shampoo as steepest descent under different norms.
• Gupta et al., Shampoo (2018); Anil et al., scalable Shampoo (2020).
• Vyas et al., SOAP (2024), arXiv:2409.11321.
• Soudry et al., implicit bias of gradient descent (2018) + subsequent implicit-bias literature (Gunasekar, Lyu, Chizat).
• Yang & Hu, μP / Tensor Programs (2021–); read enough to know the optimizer interacts with width.

Problem sets

Boyd (selected — do not do the whole book)

• Ch 2–3 (convex sets/functions): EE364A homeworks 1–2, plus Boyd additional exercises on composition rules.
• Ch 4 (convex problems): recognize LP / QP / SDP; you will not solve SDPs for a living, but you must recognize convexity.
• Ch 5 (duality): all of the conceptual exercises; derive duals of SVM and of constrained least squares.
• Ch 9 (unconstrained: gradient, Newton, convergence): do the convergence-rate exercises.
• Skip interior-point implementation unless you want it. Know what a Newton step is.

Derive-and-implement (this is the real work) Write from scratch, in numpy then in a tiny PyTorch, with tests against a known implementation:

1. SGD, SGD+momentum (heavy-ball and Nesterov), write the updates and the effective trust-region / ODE picture.
2. Adam, then AdamW. Derive the bias-correction terms. Show why decoupled weight decay ≠ L2 in Adam.
3. RMSProp, Adafactor (row/column factored second moments).
4. A Newton step on a 2-D Rosenbrock; then a damped Newton; then BFGS on the same.
5. Muon: momentum, then Newton–Schulz orthogonalization of the momentum matrix, then the update. Compare to UVᵀ from an actual SVD. This is the 2025–2026 conversation.
6. SOAP / Shampoo (read + partial implement): Kronecker-factored preconditioner L ⊗ R acting as L^{−1/4} G R^{−1/4}. Implement a toy 2-layer version. You do not need a production Shampoo.

Implicit bias (paper problems)

• Prove, for linear logistic regression on separable data, that GD direction converges to the max-margin separator (Soudry).
• Explain why this is not “the network is convex,” and what it does imply for late-training behavior.

Estimated hours

You are ready when you can

• Derive AdamW from a whiteboard, including why ε is in the denominator and why weight decay is decoupled.
• Explain Muon as “steepest descent under the spectral norm” and as “Shampoo without preconditioner accumulation, polar factor via Newton–Schulz.”
• Look at a SOAP / Shampoo PR and say what is being inverted, how often, at what precision, and why Kronecker structure is there.
• Read an implicit-bias paragraph and not confuse it with explicit regularization.

Connection to LLM kernels

• Optimizer kernels (fused AdamW, Muon Newton–Schulz, SOAP eigenbasis) are production work at Anthropic-scale. Fusing them is a GPU MODE lecture (lec 6).
• Second-order / spectral methods are matrix-valued kernels: you are orthogonalizing or taking matrix powers of gradient tiles. That is numerical LA (module 5) plus this module.
• μP says the learning-rate scaling is an operator-norm statement. Wrong optimizer + wrong μP = you cannot transfer HPs across width.

Module 5 — Numerical linear algebra / floating point

This is more important than people think for kernel work. A kernel that is 5% faster and 0.3 nats worse is a failed kernel. Most of those failures are rounding, accumulation order, and softmax.

Prerequisites

Module 1 (Trefethen done). C or Python enough to poke bits.

Primary

1. Nicholas J. Higham, Accuracy and Stability of Numerical Algorithms, 2nd ed. (2002), SIAM. Not a cover-to-cover. Selected chapters below. This is the grown-up book.
2. Trefethen & Bau (already in module 1) — reread Lectures 12–18 with floating-point eyes.
3. IEEE-754-2019 standard (read the readable parts) + Goldberg, “What Every Computer Scientist Should Know About Floating-Point Arithmetic” (1991) — still mandatory.

Secondary

• Demmel, Applied Numerical Linear Algebra (1997) — especially Ch 1–3.
• Higham’s blog (nhigham.com) — mixed precision, tf32, stochastic rounding.
• NVIDIA docs: PTX FP instructions, mul.rn, Tensor Core accumulation (TF32 / BF16 / FP8 / MXFP8 / NVFP4 as of 2025–26).
• PMPP 4e/5e Appendix on numerical considerations.
• Blanchard, Higham, et al. on mixed-precision and stochastic rounding (2019–2023).

Problem sets / labs (do all of these)

Higham 2e selected

• Ch 1–2 (principles, floating-point arithmetic): all exercises you can do without MATLAB-specific tooling. Re-do them in Python/numpy or C.
• Ch 3 (basics of rounding-error analysis): the model fl(x op y) = (x op y)(1+δ).
• Ch 6–8 (summation, polynomials, nonlinear): compensated summation, pairwise vs sequential sum.
• Ch 10, 13 (LU, perturbation): enough to know backwards stability vs condition.
• Ch 20 (automatic error analysis) optional.

IEEE-754 lab (write C or Rust)

• Print bit patterns of +0, -0, denorm min, min-normal, 1, inf, qNaN, sNaN in f32 and f16.
• Demonstrate flush-to-zero vs IEEE denorms. Time a denorm-heavy loop. This still bites GPU kernels.
• Rounding modes: RN, RZ, RU, RD. Show a GEMM-like accumulation that changes with mode.
• FMA vs a*b + c : construct a case where they differ, and a case where FMA saves you.

Kernel-relevant numerical labs

1. Softmax conditioning. Implement naive softmax, max-subtracted softmax, and online/streaming softmax (Milakov / FlashAttention-style). Compare in f32, fp16, bf16 against an f64 reference on rows with range 0, 20, 80, 200. Plot max error vs row max.
2. GEMM accumulation. Naive ijk matmul in f32 vs blocked, vs Kahan, vs pairwise. Then the same in mixed precision (fp16 multiply, f32 acc). Measure residual ‖AB − C_ref‖. Change reduction tree order. This is why Tensor Core “big enough accumulator” matters.
3. Log-sum-exp and cross-entropy fusion. Show the fused kernel is not just faster — it is more accurate (no materializing the softmax vector).
4. Newton–Schulz in bf16 vs f32. Muon’s orthogonalization is a polynomial iteration. Reproduce the “must run in f32” claim from the Muon writeup.

Estimated hours

You are ready when you can

• Look at an Nsight Compute report showing imma / Tensor Core pipes and say whether the accumulator width is the reason the loss does not match the f32 reference.
• Explain, with numbers, why bf16 GEMM + f32 accumulate is the training default and when fp8/MXFP8 breaks it.
• Write a numerically stable online softmax and prove it is equivalent to the naive one in exact arithmetic.
• Diagnose a “training loss is 0.01 higher after I fused the kernel” bug as rounding, not math.

Connection to LLM kernels

This is the job. FlashAttention exists because the naive attention materializes an S matrix that (a) does not fit and (b) is a numerical disaster if done in fp16 without the right rescale. Every fused CE, RMSNorm, and AdamW kernel is a numerical-analysis problem wearing a CUDA hat. Denorms, rounding mode, and reduction order are why two “correct” kernels disagree in the 4th decimal of the loss.

Module 6 — Algorithms + data structures (interview + real)

Prerequisites

Discrete math (proofs, induction, recurrences). If missing: Lehman/Leighton/Meyer Mathematics for Computer Science (MIT 6.042, free) — 80 h.

Primary

Cormen, Leiserson, Rivest, Stein, Introduction to Algorithms, 4th ed. (2022), MIT Press. Do not read all 1,300 pages. Selected chapters below.

Secondary

• Erickson, Algorithms (free) — cleaner for graphs.
• CLRS 4e new chapters on online algorithms and ML: skim, do not live there.
• Frigo, Leiserson, Prokop, Ramachandran, “Cache-Oblivious Algorithms” (FOCS 1999 / ACM Trans. 2012).
• Aggarwal & Vitter, I/O model (1988) — external memory.
• Aho / Ullman dragon-book graph chapters, or Cooper & Torczon Ch 9 (data-flow as graphs).
• For interviews: not a book. A disciplined problem log (see below).

Problem sets

CLRS 4e (selected)

• Ch 2–4 (analyzing algorithms, divide-and-conquer, recurrences): all problems, not just exercises. Master Master theorem and recursion trees and Akra–Bazzi enough to not fear it.
• Ch 6–8 (heaps, qsort, linear-time sort): implement + prove.
• Ch 11–13 (hashing, BSTs, RB-trees): implement a hash table and a treap or RB-tree. Know amortized vs worst-case.
• Ch 14 (augmenting DS) selected.
• Ch 15–16 (DP, greedy): all problems in 15; selected 16.
• Ch 20–23 (B-trees, graphs BFS/DFS, MST, shortest paths): all. Implement Dijkstra with a heap and with a radix/bucket heap; know when each wins.
• Ch 24 (max flow) selected — enough for “min-cut shows up in partitioning.”
• Ch 25 (matching, new in 4e) optional.
• Parallel / external: read the cache-oblivious paper and work the matmul / FFT recurrences.

Cache-oblivious / external memory (written + coded)

• Derive the I/O complexity of tiled vs recursive matmul. Implement both, sweep cache size (use perf or a cache simulator).
• This is the CPU version of “why we tile GEMMs and FlashAttention.”

Graphs for compiler / IR

• Implement CFG construction from a tiny 3-address IR.
• Dominators (Denis / Lengauer–Tarjan at least at the level of understanding).
• Liveness via data-flow equations on the graph. This is CS6120 work; do it here so compilers is not magic.

Interview-hard (real, timed)

• 80 problems, not 800. Distribution: 25 graphs, 20 DP, 15 data structures, 10 strings / hashing, 10 bits+math. Half should be “Hard.”
• After each, write a 10-line “what was the invariant.”
• This is for the interview loop and because kernel work is full of “I need a lock-free queue / a radix sort of indices / a union-find for clustering tokens.”

Estimated hours

You are ready when you can

• Derive the work/span of a parallel reduce and a tiled matmul, then implement both.
• Walk a CFG, compute dominators and liveness, and say which loads a compiler may sink.
• Solve a random LeetCode Hard graph/DP in 45 minutes and explain the I/O complexity of your solution.
• Look at a fused kernel’s index arithmetic and recognize a prefix-sum / histogram / radix-sort underneath.

Connection to LLM kernels

• Reductions, scans, histograms, segmented softmax — these are the parallel primitives in PMPP and GPU MODE lectures 9 and 24.
• Compiler IRs (Triton AST → TTGIR → LLVM → PTX/SASS; CuTe layouts) are graphs plus data-flow.
• Cache-oblivious thinking is how you invent a FlashAttention-style algorithm before you know the cache size, then tune tiles to the real SRAM.

Module 7 — Computer systems (non-negotiable)

If you skip this, you will never be a kernel engineer. You will be a person who pastes Triton.

Prerequisites

C, including pointers, malloc, and reading man 2. K&R 2e if needed (+40 h). Linux userland comfort.

7a. CSAPP — cover to cover

Primary

Randal E. Bryant & David R. O’Hallaron, Computer Systems: A Programmer’s Perspective, 3rd ed. (2016). ISBN 0-13-409266-X (hardcover). Avoid the international paperback — different, often-wrong exercises. Site: https://csapp.cs.cmu.edu/ Course: CMU 15-213 https://www.cs.cmu.edu/~213/

Problem sets / labs

Do every chapter’s practice + homework problems you can get, and all of these labs (student-site / course offerings; do not skip):

• Data lab (bit puzzles)
• Bomb lab (x86-64 reading)
• Attack lab (stack)
• Cache lab
• Shell lab
• Malloc lab
• Proxy lab

Read every chapter, including linking, exceptional control flow, virtual memory, I/O, network, concurrency. “Cover to cover” means cover to cover.

Hours: 350

Ready when: you can read a page of compiler-generated x86-64, predict cache misses for a blocked matmul, write a malloc that is not a toy, and explain a page fault from a perf trace.

7b. Operating systems

Primary

Remzi & Andrea Arpaci-Dusseau, Operating Systems: Three Easy Pieces (OSTEP), free. https://ostep.org/

Secondary

• xv6 (MIT 6.1810 / 6.S081) — https://pdos.csail.mit.edu/6.828/ · read the book, do a few labs (syscalls, traps, lazy alloc, or networking).
• Linux Kernel Development (Love, 3e) or LWN articles — only after OSTEP.

Problem sets

OSTEP chapters on processes, scheduling, VM, page tables, TLBs, locking, concurrency. Do the homework questions in the book. One xv6 lab that touches page tables.

Hours: 160

Ready when: you can draw what happens on a CUDA kernel launch from userspace through the driver, and you know which of that is analogy vs literal (GPU page tables, UVM, cudaMalloc vs mmap).

LLM connection: multi-process training, CUDA IPC, NCCL, pinned memory, GPUDirect, and “why is my DataLoader stalled” are OS problems. Virtual memory intuition is how you understand Unified Memory and why it is usually the wrong default.

7c. Compilers — Engineering a Compiler, plus LLVM mental model

Pick: Cooper & Torczon, Engineering a Compiler, 3rd ed. (2022) as primary. Dragon Book (Aho, Lam, Sethi, Ullman, Compilers, 2nd ed., 2006) as a reference for parsing and classical data-flow. Do not make Dragon Book the course — the backend chapters are thinner and older.

Course

Cornell CS6120 self-guided — https://www.cs.cornell.edu/courses/cs6120/2025fa/self-guided/ Do the lessons through LLVM, and write an LLVM pass.

Problem sets

• EaC 3e: Ch 1–4 (overview, scanners, parsers, IR) — selected; you need IR more than you need LR(1) by hand.
• Ch 8–10 (optimization, data-flow, scalar opt): problems + implement liveness and a simple DCE on a toy IR.
• Ch 11–13 (instruction selection, scheduling, register allocation) — read carefully; do selected.
• CS6120: Bril tasks + one LLVM pass that changes a program (even a simple one).

Hours: 200

Ready when: you can look at torch.compile / Triton / XLA / NVFuser output and say: this is an IR, these are the loops, this is the allocation, this is why the scheduler fused these two elementwise ops and refused the third.

LLM connection: kernel compilers are compilers. CuTe layout algebra, Triton’s blocked IR, ThunderKittens tiles, Helion/Gluon, torch.compile FX → Inductor — if you only know CUDA C++, you are programming in 2015.

7d. Computer architecture

Primary

John L. Hennessy & David A. Patterson, Computer Architecture: A Quantitative Approach, 6th ed. (2017). (A 7e is signaled as available in 2026 catalogs; if you buy now, 6e is still the one people have problem sets for. Prefer 7e if you can confirm the chapter map.)

Secondary

• Patterson & Hennessy Computer Organization and Design (RISC-V ed.) if 6e CA:AQA is too abrupt.
• NVIDIA whitepapers: V100, A100, H100, B200 — treat as chapter 7 of this module.
• Onur Mutlu’s lecture notes (ETH / CMU) for memory systems.

Problem sets

• Ch 1 (fundamentals, Amdahl, energy): all exercises.
• Ch 2 (memory hierarchy): all. This is caches, then DRAM, then the reason HBM exists.
• Ch 3–4 (ILP, data-level parallelism): selected; SIMD/vector is the CPU cousin of warps.
• Ch 5 (multiprocessors / coherence): enough to understand why GPUs are not coherent the way CPUs are.
• Ch 7 (DSA) in 6e: read. GPUs and TPUs are the case study.
• Then: draw an H100/B200: HBM → L2 → SM shared mem / TMEM → tensor core. Label bandwidths. This is a required artifact.

Hours: 150

Ready when: you can look at a roofline, name the ridge point, and predict whether a kernel is HBM-bound or tensor-core-bound before you write it. You can explain why FlashAttention is an arithmetic-intensity play, and why Blackwell TMEM changes the FlashAttention-4 / TK 2.0 design.

7e. Networks enough for RDMA / InfiniBand intuition

You do not need to become a network engineer. You do need to know why all-reduce has a ring and why IB verbs look like post_send.

Primary

• Kurose & Ross, Computer Networking, 8th ed. — Ch 1–3, 6 (enough).
• RDMA/IB: NVIDIA/Mellanox “RDMA Aware Programming” user manual; a short IB verbs tutorial (e.g. RDMAmojo, or the NCCL / MSCCL paper related work).
• NCCL paper + GPU MODE lectures 17 and 67 (NCCL, NVSHMEM).
• ThunderKittens PGL / ParallelKittens posts (2025): https://hazyresearch.stanford.edu/blog/2025-09-22-pgl

Problem sets

• Compute time for a 1.2 GB all-reduce on 8 GPUs over 400 Gb/s IB vs NVLink, with a ring vs a tree, including a 2 μs latency term. Then compare to a GEMM of the same bytes.
• Read an NCCL debug log and identify the algorithm (ring / tree / NVLS).
• Write a 2-process CPU RDMA ping-pong or, if you have no IB, a correct-but-simulated version over sockets that uses the same verb structure (reg mr, post recv, post send, poll cq).

Hours: 60

Ready when: you can explain GPUDirect RDMA, why NCCL uses topology-aware rings, and why a “fused all-reduce + GEMM” (async TP) is a networks problem as much as a kernel problem.

Module 7 total: 920 hours

Module 8 — Parallel computing theory

Prerequisites

CSAPP done (at least through memory hierarchy + concurrency). Module 1. CLRS recurrences.

Primary

1. Cilk / work-span theory:

• Blumofe & Leiserson, “Scheduling Multithreaded Computations by Work Stealing” (1999).
• MIT 6.172 / 6.1060 lectures on work and span.
• Brent’s theorem: T_P ≤ T_1 / P + T_∞. Live this formula.

1. Williams, Waterman, Patterson, “Roofline” (CACM 2009) + the GPU roofline papers (Yang, et al.; NVIDIA Nsight Compute roofline).
2. Communication-avoiding algorithms:

• Demmel SC12 tutorial: https://people.eecs.berkeley.edu/~demmel/SC12_tutorial/Demmel_SC12_Tutorial_final.pdf
• Solomonik & Demmel, “Communication-optimal parallel 2.5D matrix multiplication and LU” (2011).
• Ballard, Demmel, Holtz, Schwartz, communication lower bounds (2010–).
• CS267 lecture notes: https://sites.google.com/lbl.gov/cs267-spr2025

1. A full parallel course, one of:

• Stanford CS149 Fall 2025: https://gfxcourses.stanford.edu/cs149/fall25/
• CMU 15-418 Spring 2026: https://www.cs.cmu.edu/~418/
• MIT 6.S894 Fall 2025: https://accelerated-computing.academy/fall25/

Secondary

• PMPP 4e or 5e (2026) — this is the CUDA patterns book; put the deep reading in the applied layer, but start Ch 1–6 here.
• Culler / Singh / Gupta Parallel Computer Architecture — historical, useful for PRAM-vs-reality.
• Blelloch, “Prefix Sums and Their Applications.”
• McColl / Tiskin, PRAM and BSP.

Problem sets

• Compute work, span, and parallelism of: reduce, scan, filter, matmul, FlashAttention-style blocked attention, a data-parallel Adam step.
• Prove Brent’s theorem and use it to decide when adding SMs stops helping.
• Draw a roofline for an H100 (HBM ~3 TB/s, dense BF16 tensor ~990 TFLOP/s — look up the current datasheet and use real numbers). Place: naive attention, FlashAttention, RMSNorm, AdamW, a 4096³ GEMM.
• Derive the 2D SUMMA communication volume. Then the 2.5D improvement with c extra copies. This is tensor-parallelism’s grandfather.
• PRAM: write a O(log n) span reduce. Then write the paragraph on why PRAM is a lie (bandwidth, latency, synchronization, control divergence, memory coalescing).
• Do the programming assignments of whichever course you picked. CS149/15-418 without the labs is a reading group.

Estimated hours

You are ready when you can

• Given a kernel, estimate arithmetic intensity, place it on a roofline, and say whether tiling, fusion, or a better algorithm is the move.
• Explain why FlashAttention is a communication-avoiding algorithm in the Demmel sense (recompute to avoid HBM traffic).
• Explain 2.5D matmul and then point at ZeRO / TP / SP / context-parallel ring attention as the same family.
• Say, precisely, what PRAM gets wrong about a warp.

Connection to LLM kernels

This module is the theory of the job. Tensor parallel, sequence parallel, expert parallel, ring attention, and Megatron-style overlapping are communication-avoiding (or communication-hiding) algorithms on a particular topology (NVLink / NVSwitch / IB). Work/span tells you why a tiny sequential piece (one expert’s token sort, a lock, a CPU-side GIL) wrecks a 256-GPU run.

Applied GPU / LM-systems layer (after foundations)

Not a “foundations” module, but you asked for the 2024–2026 courses. Do this after modules 1, 5, 7a, and 8-theory. Overlap with 7c–7e and 4 is fine.

You are ready for an Anthropic-style kernel interview when you can, in one week, take a slow PyTorch op, write a Triton or CUDA kernel, beat the baseline on a real GPU, explain the roofline, the numerical error vs f32, and the communication if it is multi-GPU — and you can talk about Muon / FA / μP without reading from a card.

Hour roll-up (foundations)

Interview-loop extra (company-specific + more LeetCode) is not counted; budget another 80–120 h when you are actually interviewing.

Two minimum-viable orders

Same destination. Different first year. Do not “do a bit of everything.”

Track A — Strong undergrad CS, weak math

You can write C++, you have seen an OS course or CSAPP chapters, you have done some LeetCode. Matrices are “things numpy multiplies.” Proofs make you itch.

Principle: fix analysis and linear algebra before you touch CUDA in anger. Otherwise every kernel paper is a fog of symbols. But use CSAPP immediately so you do not lose your systems identity.

Track A first 400 hours if you must cut: A0 + Axler Ch 1–8 + CSAPP through cache lab + IEEE-754/softmax lab. That is the minimum that makes GPU MODE lec 1–5 not theater.

Do not start CS336 A2 until A2, A5, and A6-theory are done. You will otherwise copy Triton without knowing why the tl.dot is in that precision.

Track B — Strong math, weak systems

You have a math or physics degree. Axler is easy, you have seen some analysis and probability. You have never written a malloc, you treat the GPU as a cloud API, and volatile is a personality trait.

Principle: CSAPP is week 1, not “later.” Your comparative advantage is already modules 1–4. Spend it on becoming dangerous with machines. Do not take another analysis course to procrastinate gdb.

Track B first 400 hours if you must cut: B0 + CSAPP through cache + malloc labs + module 5 softmax/GEMM labs + roofline notes. Then write a CUDA GEMM. Your math will carry the rest of LA/opt; the failure mode is remaining unable to read ncu.

What “no shortcuts” means in practice

• If a chapter has exercises, the exercises are the chapter.
• If a course has labs, the labs are the course. CS149 without CUDA assignments is a podcast.
• Implement: Householder QR, online softmax, AdamW, Muon Newton–Schulz, a tiled GEMM, a Triton FA2, a malloc, an LLVM pass, a reduce-scan. If that list is not in your git log, you are not done.
• Re-derive, on paper, once a quarter: SVD, softmax Jacobian, Hoeffding, Brent, AdamW, IEEE-754 encoding of 0.1f.
• Popular visualizations (3B1B, etc.) are allowed after you have done the proof, as a memory aid.

What this curriculum is not

It is not “how to get hired at Anthropic in 2026.” That also needs a public artifact trail (kernels, blog posts, GPU MODE contest, CS336-quality repo), taste, and the ability to work with other people on a huge stack. It is the foundations those people are assuming when they say “we need someone who can actually do the math and the machine.”

It is also not a substitute for reading the live stack: Megatron / TorchTitan, NCCL, CUTLASS, ThunderKittens, torch.compile, and whatever Anthropic’s internal kernel lib is. Foundations make those readable. They do not replace them.

Bottom line: ~4,000 deep hours, two tracks, Axler+Trefethen (not Strang-only), CSAPP cover-to-cover, Higham-level floating point, Boyd-selected + derive-Muon, Blitzstein-then-Cover, one of {CS149, 15-418, 6.S894}, then GPU MODE + CS336 A1/A2. That is the hard way.