GPU kernel curriculum (no shortcuts)

As of 2026-08-18. Deep-study hours, not calendar time. Production bar: Hopper (H100/H200) plus public Blackwell.

LAYER 0 — Computer architecture and C/C++

Deep hours: 220–280. Do not start CUDA until you can explain a cache-miss stall from a perf sample and write lock-free code that is actually correct.

Prerequisites

High-school discrete math, one year of programming, comfort with Linux command line. No GPU required yet.

Resources (10)

1. Bryant & O’Hallaron, Computer Systems: A Programmer’s Perspective, 3rd ed. — https://csapp.cs.cmu.edu/ Data representation, machine code, ILP, memory hierarchy, virtual memory. Do the labs, not just the chapters.
2. Lippman, Lajoie, Moo, C++ Primer, 5th ed. (2013) — https://www.csl.cornell.edu/courses/ece2400/resources.shtml C++11/14 baseline. Then write everything in C++17/20.
3. Scott Meyers, Effective Modern C++ (2014) — same resource list as above. Value categories, move, constexpr, atomics-adjacent lifetime bugs.
4. C++ memory model — https://en.cppreference.com/w/cpp/language/memory_model std::memory_order — https://en.cppreference.com/w/cpp/atomic/memory_order std::atomic — https://en.cppreference.com/w/cpp/atomic/atomic
5. Ulrich Drepper, What Every Programmer Should Know About Memory (2007) — https://www.cs.tufts.edu/cs/40/docs/WhatEveryProgrammerShouldKnowAboutMemory.pdf Still the best cache/TLB/NUMA essay. Pair with CSAPP Ch. 6.
6. Agner Fog, Optimizing software in C++ — https://agner.org/optimize/optimizing_cpp.pdf Plus instruction tables at https://agner.org/optimize/. SIMD, ILP, ABI.
7. Brendan Gregg, Systems Performance, 2nd ed. (2020) — https://www.brendangregg.com/systems-performance-2nd-edition-book.html Methodology (USE), perf, Ftrace, BPF.
8. Linux perf examples — https://www.brendangregg.com/perf.html Official wiki — https://perf.wiki.kernel.org/index.php/Main_Page
9. Compiler Explorer — https://godbolt.org You will live here for both host C++ and later PTX/SASS.
10. Hennessy & Patterson, Computer Architecture: A Quantitative Approach (any recent ed.) — cache-coherence and ILP chapters only. Builder’s view after CSAPP’s programmer view.

Optional but high leverage: Matt Godbolt, “Optimizations in C++ Compilers,” ACM Queue 2019.

Graded exercises (produce artifacts)

1. CSAPP labs repo. Complete Data Lab, Bomb Lab, Attack Lab, Cache Lab, Malloc Lab, Performance Lab. Artifact: public git repo with writeups of every cache-miss and overflow you hit.
2. memcpy / saxpy / reduction from scratch in C++20. Measure with perf stat -e cycles,instructions,cache-misses,L1-dcache-load-misses. Artifact: table of GB/s vs. stride, plus a flame graph.
3. Lock-free SPSC ring buffer using memory_order_acquire/release. Prove it with TSAN and a stress harness that fails under relaxed. Artifact: header-only library + failing/passing tests.
4. SIMD saxpy and GEMM microkernel with AVX2 (and AVX-512 if you have it). Compare against OpenBLAS for 64×64 and 256×256. Artifact: Godbolt links showing FMA throughput vs. dependency-chain latency.
5. Roofline of a CPU GEMM. Compute arithmetic intensity, plot against DRAM and L2 roofs. Artifact: one figure + a 2-page note.
6. Write a tiny allocator (bump + free-list) and measure fragmentation under LLM-shaped allocation traces (many 2–64 KB buffers). Artifact: allocator + benchmark.
7. perf record a real program (e.g. llama.cpp CPU path). Produce annotated source and a flame graph. Artifact: “here is the actual stall.”

Failure modes / interview questions

• Confusing cache coherence with memory ordering.
• Thinking volatile is a mutex.
• Not knowing store-buffer / load-buffer effects; cannot explain why x86 is “almost SC” and ARM is not.
• Cannot compute bytes/FLOP or explain why a reduction is bandwidth-bound.
• Using new/delete in a hot loop; ignoring false sharing.
• Interview: “Draw the cache hierarchy. Where does a 64-byte stride miss? What does memory_order_acquire prevent? Show an ILP-limited loop and how you break the chain.”

Level-up bar

• You can look at perf annotate and say which stall (frontend, backend, bad speculation, retiring).
• Your AVX2 256×256 SGEMM is within 2–3× of OpenBLAS, and you can explain the remaining gap (packing, prefetch, kernel shape).
• You can write a correct acquire/release queue without looking it up.
• You instinctively convert “this kernel is slow” into “bound by X at intensity Y.”

LAYER 1 — GPU architecture

Deep hours: 320–400. This is the longest layer. The output is not “I can write a kernel.” It is “I can predict occupancy, shared-memory bank conflicts, and whether TMA+WGMMA will help before I write the kernel.”

Prerequisites

Layer 0 complete. Comfortable C++17. Access to a CUDA GPU (Ampere+ preferred). Install CUDA Toolkit 12.8+ (13.3 current as of this writing).

Resources (12)

1. CUDA Programming Guide (current, replaces the legacy C++ guide as of CUDA 13.0) — https://docs.nvidia.com/cuda/cuda-programming-guide/index.html SIMT kernels — https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/writing-cuda-kernels.html Async copies / TMA — https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/async-copies.html
2. CUDA C++ Best Practices Guide — https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/
3. CUDA Toolkit portal (13.3) — https://docs.nvidia.com/cuda/ Includes PTX, samples, architecture tuning, CUDA Tile C++.
4. Hwu, Kirk, El Hajj, Programming Massively Parallel Processors, 4th ed. (2022, Ampere) — https://shop.elsevier.com/books/programming-massively-parallel-processors/hwu/978-0-323-91231-0 5th ed. (Feb 2026) adds warp-level programming, CUDA atomics, multi-GPU NCCL/NVSHMEM, advanced GEMM, LLMs — https://shop.elsevier.com/books/programming-massively-parallel-processors/hwu/978-0-443-43900-1 Prefer 5th if you can get it; otherwise 4th + Hopper docs.
5. NVIDIA Hopper Architecture In-Depth — https://developer.nvidia.com/blog/nvidia-hopper-architecture-in-depth/ Hopper Tuning Guide — https://docs.nvidia.com/cuda/hopper-tuning-guide/ Clusters (portable size 8, H100 non-portable 16), TMA 1D–5D, distributed SMEM, 228 KB SMEM/SM with 1 KB reserved.
6. PTX ISA 9.3 — https://docs.nvidia.com/cuda/parallel-thread-execution/index.html PTX interoperability / C++ atomics mapping — https://docs.nvidia.com/cuda/ptx-writers-guide-to-interoperability/index.html
7. Nsight Compute 13.3 — https://docs.nvidia.com/nsight-compute/ Nsight Systems — https://developer.nvidia.com/nsight-systems Occupancy calculator lives inside Nsight Compute.
8. CUDA Samples — https://github.com/NVIDIA/cuda-samples
9. Blackwell, public as of 2025–2026 Official architecture overview — https://resources.nvidia.com/en-us-blackwell-architecture Product page — https://www.nvidia.com/en-us/data-center/technologies/blackwell-architecture/ PTX tcgen05 + TMEM (SM100 only, 256 KB/SM); workstation SM120 does not get TMEM/tcgen05 CTA-pair MMA. Community wiki (useful, not official) — https://0xsero.github.io/blackwell-gpu-wiki/blackwell/tcgen05-and-tmem/ Microbenchmark paper — https://arxiv.org/html/2512.02189v3 Colfax SM100 GEMM/TMEM — https://research.colfax-intl.com/cutlass-tutorial-writing-gemm-kernels-using-tensor-memory-for-nvidia-blackwell-gpus/
10. CUDA Tile C++ (CUDA 13.3) — https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/writing-tile-kernels.html API — https://docs.nvidia.com/cuda/cuda-tile-cpp-api-reference/index.html Blog — https://developer.nvidia.com/blog/develop-high-performance-gpu-kernels-in-cpp-with-nvidia-cuda-tile/
11. GPU MODE lectures (slides + code, 100+ talks) — https://github.com/gpu-mode/lectures Resource stream — https://github.com/gpu-mode/resource-stream Discord — https://discord.gg/gpumode
12. PyTorch TMA unit writeup — https://pytorch.org/blog/hopper-tma-unit/

Graded exercises (artifacts)

1. PMPP / CUDA samples path. Vector add → reduction → scan → histogram → softmax. Every kernel profiled with ncu. Artifact: repo + one Nsight Compute report per kernel.
2. Occupancy vs. performance sweep. Same kernel at 3 register/SMEM points. Show a case where lower occupancy is faster. Artifact: plot + explanation (latency hiding vs. cache).
3. Bank-conflict microbench. Shared-memory 32-way conflict vs. swizzled. Measure with l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld. Artifact: before/after SASS.
4. Warp divergence + predication. Write a kernel that looks branchy in C++ and is predicated in SASS. Artifact: cuobjdump -sass annotated.
5. Roofline with Nsight Compute. Profile a memory-bound and a compute-bound kernel. Use the hierarchical roofline (L1/L2/HBM). Artifact: two roofline charts.
6. PTX-in-CUDA. Write a kernel that uses inline PTX for ldmatrix (Ampere) or a cp.async. Artifact: PTX + matching SASS.
7. Hopper-only (if you have H100): TMA load of a 2D tile into SMEM + mbarrier, no tensor cores yet. Artifact: working cluster-size-1 kernel and an Nsight Systems trace showing async copy overlapping compute.
8. Blackwell reading assignment (even without B200): write a 3-page note contrasting mma.sync / wgmma.async / tcgen05.mma + TMEM. Cite PTX ISA and the Colfax SM100 tutorial. Artifact: the note.

Failure modes / interview questions

• Treating a thread block as “a CPU thread pool.”
• Ignoring that warps are the schedulable unit (32 threads).
• Occupancy-maxing as a religion.
• Not knowing SMEM is banked 32 ways; transposing the wrong operand.
• Confusing L1 and SMEM (they carve the same unified cache on many SMs).
• Cannot explain TMA: single thread issues, hardware copies GMEM↔SMEM, multicast inside a cluster, completion via mbarrier.
• Interview: “What is a warp scheduler? How many on H100 SM? Why does register pressure kill occupancy? Show me a coalesced vs. strided load in SASS. What does WGMMA consume from SMEM vs. registers? What is TMEM and which chips have it?”

Level-up bar

• You can look at an ncu report and name the limiter (dram, l2, lsu, tensor, issue, barrier) in one sentence.
• You can compute theoretical occupancy from registers + SMEM + block size without the GUI, then verify.
• You can read PTX of your kernel and point to the instruction you intended (LDG.128, LDSM, WGMMA, TMA).
• On H100: a memory-bound copy kernel within 70%+ of 3.35 TB/s; you know why you are not at 100% (alignment, TLB, L2).

LAYER 2 — GEMM and CUTLASS

Deep hours: 280–360. GEMM is the kernel engineer’s scales-and-arpeggios. Everything else is GEMM plus a fused epilogue or an online reduction.

Prerequisites

Layer 1. You have written a shared-memory tiled SGEMM that you have profiled.

Resources (12)

1. Simon Boehm, “How to Optimize a CUDA Matmul Kernel for cuBLAS-like Performance” — https://siboehm.com/articles/22/CUDA-MMM Code: typically https://github.com/siboehm/SGEMM_CUDA Naive → coalesced → SMEM → 1D/2D tiling → vectorized → autotune → warptile. 93.7% of cuBLAS FP32 on A6000 in the original worklog.
2. Aleksa Gordić, “Inside NVIDIA GPUs: Anatomy of high performance matmul kernels” — https://www.aleksagordic.com/blog/matmul PTX/SASS, warptiling, then Hopper TMA + WGMMA.
3. CUTLASS — https://github.com/NVIDIA/cutlass Docs — https://docs.nvidia.com/cutlass/latest/media/docs/cpp/quickstart.html CuTe layouts — https://github.com/NVIDIA/cutlass/blob/main/media/docs/cpp/cute/00_quickstart.md CuTe GEMM tutorial — https://docs.nvidia.com/cutlass/latest/media/docs/cpp/cute/0x_gemm_tutorial.html CUTLASS 3.x design — https://docs.nvidia.com/cutlass/4.6.2/media/docs/cpp/cutlass_3x_design.html
4. CUTLASS origin blog — https://developer.nvidia.com/blog/cutlass-linear-algebra-cuda/ CUTLASS 3.x composable GEMM — https://developer.nvidia.com/blog/cutlass-3-x-orthogonal-reusable-and-composable-abstractions-for-gemm-kernel-design/ CuTe layout algebra — https://developer.nvidia.com/blog/cutlass-principled-abstractions-for-handling-multidimensional-data-through-tensors-and-spatial-microkernels/
5. Colfax Research (the industrial CUTLASS classroom) Hopper GEMM + TMA/WGMMA — https://research.colfax-intl.com/nvidia-hopper-gemm-cutlass/ WGMMA — https://research.colfax-intl.com/cutlass-tutorial-wgmma-hopper/ Pipelining / warp-specialization — https://research.colfax-intl.com/cutlass-tutorial-design-of-a-gemm-kernel/ Blackwell TMEM GEMM — https://research.colfax-intl.com/cutlass-tutorial-writing-gemm-kernels-using-tensor-memory-for-nvidia-blackwell-gpus/
6. Triton GEMM tutorial — https://triton-lang.org/main/getting-started/tutorials/03-matrix-multiplication.html Grouped GEMM (MoE-shaped) — https://triton-lang.org/main/getting-started/tutorials/08-grouped-gemm.html Persistent + TMA + CLC matmul — https://triton-lang.org/main/getting-started/tutorials/09-persistent-matmul.html
7. cuBLAS on Hopper (how the library thinks: heuristics, fused epilogues, CUTLASS instantiation) — https://developer.nvidia.com/blog/new-cublas-12-0-features-and-matrix-multiplication-performance-on-nvidia-hopper-gpus/
8. DeepGEMM (DeepSeek JIT tensor-core GEMM, grouped MoE layouts) — https://github.com/deepseek-ai/DeepGEMM
9. GPU MODE lecture 15 (CUTLASS) and 36 (CUTLASS + FA3, Jay Shah) — https://github.com/gpu-mode/lectures
10. Lei Mao, “Learn CUTLASS the Hard Way” — https://leimao.github.io/article/Learn-CUTLASS-The-Hard-Way/
11. NVIDIA CUTLASS examples: examples/cute/tutorial/sgemm_1.cu, sgemm_2.cu, Hopper GETT, Blackwell FMHA — under https://github.com/NVIDIA/cutlass/tree/main/examples
12. Anatomy paper (Goto/van de Geijn mental model, still useful) — “Anatomy of high-performance matrix multiplication,” ACM TOMS — https://dl.acm.org/doi/10.1145/1356052.1356053

How cuBLAS actually works (high level, not a reverse-engineer)

cuBLAS is a heuristic dispatcher over many compiled kernels, not one GEMM. Boehm traced square SGEMM on Ampere and saw different kernels at 128 / 256 (split-K + reduce) / 512 / 1024+. Many modern kernels are CUTLASS instantiations with a large binary and a shape/dtype/layout lookup. cuBLASLt adds fused epilogues (bias, ReLU/GELU, dBias). You will not beat it on vanilla large GEMM; you beat it on odd shapes, fusions, and persistent/grouped cases.

Graded exercises (artifacts)

1. Rebuild Boehm’s ladder on your GPU (A100 or H100). Same kernel numbers. Artifact: table of TFLOPS vs. cuBLAS for M=N=K ∈ {1024, 2048, 4096, 8192}.
2. Tensor-core GEMM (Ampere mma.sync or Hopper wgmma). FP16 in, FP32 accum. Artifact: kernel + ncu showing tensor-pipe utilization.
3. CuTe sgemm_1 then sgemm_2 from scratch (type them, do not copy-paste blindly). Artifact: your own layouts for NT and TN.
4. CUTLASS 3.x Hopper GEMM using Collective mainloop + TMA + WGMMA. Beat your hand kernel. Artifact: instantiated kernel + occupancy/pipeline notes.
5. Triton tutorial GEMM + autotune. Compare to cuBLAS and your CUDA. Artifact: Triton config space and a plot.
6. Grouped GEMM (Triton tutorial 08 or CUTLASS grouped). MoE-shaped: many small M, fixed N/K. Artifact: tokens-per-expert sweep.
7. Persistent GEMM (Triton 09 or CUTLASS persistent/Stream-K). Show why small-grid GEMMs under-occupy H100. Artifact: Stream-K vs. data-parallel tile schedule.
8. Epilogue fusion: GEMM + bias + GELU in one kernel (CUTLASS epilogue or hand). Measure vs. two kernels. Artifact: bandwidth saved.
9. Read DeepGEMM m_grouped_fp8_gemm_*_contiguous and write a 2-page “how the M-grouped layout works.” Artifact: the note + a toy clone of the layout.

Failure modes / interview questions

• Stopping at SMEM tiling and calling it “optimized.”
• Wrong swizzle → silent bank conflicts (FA3 authors have called this out in other kernels).
• Ignoring K-padding / alignment (ldmatrix 8×8 / TMA 16-byte).
• Confusing CuTe logical layout with physical SMEM swizzle.
• Interview: “Walk naive → tiled → tensor-core GEMM. Why square output tiles? What is a warpgroup? Draw a 2-stage TMA/WGMMA pipeline. How does Stream-K fix wave quantization? When is split-K a win?”

Level-up bar (H100 SXM, large square, FP16 tensor)

• Hand CUDA/CUTLASS GEMM ≥ 80% of cuBLAS at 4096³+ (roughly ≥ 790 TFLOPS if cuBLAS is near the 989 roof; measure your cuBLAS, do not assume).
• Triton GEMM ≥ 70% of cuBLAS on the same shape after autotune.
• You can instantiate a CUTLASS 3 Hopper GEMM without a tutorial open.
• You can explain, from an ncu report, whether you are TMA-bound, WGMMA-bound, or epilogue-bound.

LAYER 3 — LLM kernels

Deep hours: 360–450. This is the job. GEMM is necessary; attention/norm/fused-epilogue is what labs hire for.

Prerequisites

Layer 2. You have a tensor-core GEMM and understand online softmax on paper before you write FlashAttention.

Resources (12)

1. Milakov & Gimelshein, Online normalizer calculation for softmax (2018) — https://arxiv.org/abs/1805.02867 Reference impl mentioned in the paper: https://github.com/NVIDIA/online-softmax
2. Zihao Ye, “From Online Softmax to FlashAttention” — https://courses.cs.washington.edu/courses/cse599m/23sp/notes/flashattn.pdf
3. FlashAttention-1 — https://arxiv.org/abs/2205.14135 FlashAttention-2 — https://tridao.me/publications/flash2/flash2.pdf (also https://arxiv.org/abs/2307.08691) FlashAttention-3 — https://tridao.me/publications/flash3/flash3.pdf (https://arxiv.org/abs/2407.08608) Blog — https://dao-lab.ai/blog/2024/flash3/ and https://pytorch.org/blog/flashattention-3/ Code — https://github.com/Dao-AILab/flash-attention
4. PagedAttention / vLLM — paper https://arxiv.org/pdf/2309.06180 (https://arxiv.org/abs/2309.06180) Repo — https://github.com/vllm-project/vllm Kernel design — https://docs.vllm.ai/en/latest/design/paged_attention/ (also https://docs.vllm.ai/en/v0.9.2/design/kernel/paged_attention.html)
5. FlashInfer (production inference attention/GEMM/MoE generator) — https://github.com/flashinfer-ai/flashinfer Paper — https://arxiv.org/abs/2501.01005
6. RoPE: Su et al., RoFormer — https://arxiv.org/abs/2104.09864 LayerNorm: Ba et al. (2016) — https://arxiv.org/abs/1607.06450 RMSNorm: Zhang & Sennrich (2019) — https://arxiv.org/abs/1910.07467 SwiGLU: Shazeer, GLU Variants Improve Transformer — https://arxiv.org/abs/2002.05202
7. Liger-Kernel (Triton RMSNorm, RoPE, SwiGLU, fused linear-CE) — https://github.com/linkedin/Liger-Kernel Paper — https://arxiv.org/abs/2410.10989
8. Transformer Engine (FP8/MXFP8/NVFP4 production path) — https://docs.nvidia.com/deeplearning/transformer-engine/user-guide/index.html FP8/FP4 primer — https://docs.nvidia.com/deeplearning/transformer-engine/user-guide/examples/fp8_primer.html NVFP4 — https://docs.nvidia.com/deeplearning/transformer-engine/user-guide/features/low_precision_training/nvfp4/nvfp4.html MXFP8 — https://docs.nvidia.com/deeplearning/transformer-engine/user-guide/features/low_precision_training/mxfp8/mxfp8.html Code — https://github.com/NVIDIA/TransformerEngine
9. DeepGEMM grouped + Mega-MoE — https://github.com/deepseek-ai/DeepGEMM
10. Speculative decoding (kernel-relevant: tree attention, verify GEMM) Medusa — https://github.com/FasterDecoding/Medusa (paper https://arxiv.org/abs/2401.10774) EAGLE — https://github.com/SafeAILab/EAGLE (EAGLE-1 https://arxiv.org/abs/2401.15077, EAGLE-3 NeurIPS 2025)
11. Educational FA implementations: https://github.com/gitctrlx/flash-attention-tutorial and https://github.com/shreyansh26/FlashAttention-PyTorch
12. GPU MODE lectures on FA, FA3, quantization, speculative decoding — https://github.com/gpu-mode/lectures (FA through FA4 appears in later lecture numbers; lecture 36 is Colfax FA3)

Kernel map you must be able to reimplement

Graded exercises (artifacts)

1. Softmax ladder. Naive, safe, online, then a CUDA warp-level online softmax. Match PyTorch to 1e-5 rel. Artifact: kernel + bandwidth table.
2. RMSNorm + fused residual+RMSNorm in CUDA or Triton. Match TE or Liger within 5% time at hidden=4096 and 8192. Artifact: repo.
3. RoPE kernel (Triton or CUDA). Fuse into Q/K. Artifact: vs. HuggingFace vector RoPE.
4. PyTorch tiled FlashAttention (forward only, then backward). No CUDA. Must match torch.nn.functional.scaled_dot_product_attention numerically. Artifact: notebook that never allocates N×N.
5. Triton FlashAttention-2-style (forward + causal + GQA). Benchmark vs. flash_attn on your GPU. Artifact: kernel + speed/% of FA2.
6. CUDA FA2-style on Ampere (mma.sync) or Hopper CUTLASS FA3-style if you have H100. Start with head_dim=64 or 128, seq=1024–4096, causal. Artifact: this is the portfolio piece.
7. Fused linear + SwiGLU and fused linear + cross-entropy (chunked). Compare to Liger. Artifact: memory peak plot (the point is not materializing [B,S,V]).
8. Paged decode attention. Implement block-table gather + online softmax for decode (seqlen_q=1, paged KV). Compare to vLLM’s documented layout (k_cache [num_blocks, num_kv_heads, head_size/x, block_size, x]). Artifact: kernel + a tiny block allocator.
9. Grouped GEMM for MoE. Variable tokens/expert, fixed N/K. Compare contiguous vs. padded. Artifact: tokens-per-expert histogram vs. TFLOPS.
10. FP8 attention or GEMM using TE recipes and a from-scratch scale/swizzle path if you have Hopper+. Artifact: numerical-error table vs. BF16 (FA3 claims 2.6× lower error than naive FP8 attn).

Speculative decoding: implement a tree-attention mask verify kernel (not a full EAGLE stack). The kernel skill is “one target forward over a tree of draft tokens,” which is a block-sparse / packed attention variant.

Failure modes / interview questions

• Reimplementing FA2 but writing S back to HBM “just to debug.”
• Wrong online-softmax rescale of O (the exp(m_old - m_new) factor).
• Causal mask load-imbalance; forgetting zigzag when you later do ring.
• Ignoring GQA (repeat KV heads) and wasting SMEM.
• Paged attention: treating the block table as optional; uncoalesced gather.
• FP8: using tensor-wise scale when the hardware wants block scale; forgetting columnwise/rowwise dual quant for training.
• Interview: “Derive online softmax. Write FA2 forward in 20 lines of numpy. Why is FA2 only 35% on H100? What three things does FA3 add? How does paged KV change the inner loop? Why is decode memory-bound and prefill compute-bound? How do you fuse residual+RMSNorm without an extra HBM write? What is NVFP4’s scale hierarchy (E2M1 values, 16-element blocks, FP8 E4M3 scales, FP32 tensor scale)?”

Level-up bar

• Triton FA-style within 20% of FA2 on A100/H100 for seq=2048–4096, head=128, causal, FP16.
• CUDA/CUTLASS attention within 30% of FA2 on Ampere, or within 25% of FA3 on Hopper for a restricted shape (fixed head, power-of-two seq). Full FA3 parity is a team artifact; interviews look for the algorithm + a serious kernel, not 740 TFLOPS on day one.
• RMSNorm/RoPE/SwiGLU ≥ 80% of Liger or TE at hidden≥4096.
• Fused CE: peak memory ≪ B·S·V; you can state the chunk size math.
• You can read csrc/ in FlashAttention and csrc/attention/ in vLLM and narrate the pipeline.

LAYER 4 — Compilers and DSLs

Deep hours: 160–220. The frontier skill is choosing the layer, not worshipping one.

Prerequisites

You have written both a CUDA GEMM and a Triton GEMM. You have profiled both.

Resources (10)

1. OpenAI Triton announcement (compiler pipeline: Python AST → Triton-IR → LLVM → PTX) — https://openai.com/index/triton/ Language/repo — https://github.com/triton-lang/triton Docs/tutorials — https://triton-lang.org/main/getting-started/tutorials/03-matrix-multiplication.html
2. torch.compile / Inductor — https://docs.pytorch.org/docs/stable/user_guide/torch_compiler/torch.compiler.md Getting started (Triton dump via TORCH_COMPILE_DEBUG=1) — https://docs.pytorch.org/docs/stable/user_guide/torch_compiler/torch.compiler_get_started.md CUDA Graph trees — https://pytorch.org/docs/stable/torch.compiler_cudagraph_trees.html
3. CUDA Graphs (runtime) — https://pytorch.org/docs/stable/notes/cuda.html#cuda-graphs
4. ThunderKittens 2.0 (Jan 2026; Blackwell, MXFP8, NVFP4) — https://github.com/HazyResearch/thunderkittens Paper — https://arxiv.org/abs/2410.20399 Blog index on the repo README (FA3-matching GEMM/attn claims).
5. Helion (PyTorch DSL → autotuned Triton; TileIR backend for Blackwell) — https://github.com/pytorch/helion Blog — https://pytorch.org/blog/helion/ Docs mentioned as helionlang in the 2025-10-22 beta post.
6. TileLang (TVM-based; CuTe DSL backend as of 2025-12) — https://github.com/tile-ai/tilelang
7. CUDA Tile C++ / cuTile — https://docs.nvidia.com/cuda/cuda-tile-cpp-api-reference/index.html
8. GPU MODE Triton internals talks (Kapil Sharma) — listed in https://github.com/gpu-mode/resource-stream
9. Colfax + CUTLASS Python / CuTe DSL quickstart — https://docs.nvidia.com/cutlass/latest/media/docs/pythonDSL/quick_start.html
10. When-to-use decision (from production practice, not a single paper): see vLLM 2025 CMU talk notes on “manual CUDA vs torch.compile” — https://llmsystem.github.io/llmsystem2025spring/assets/files/llmsys-22-vLLM_woosuk_kwon-1f34697dbb1a1fb5b798daf6eff14b67.pdf

When to write CUDA vs Triton vs CUTLASS vs a DSL

Graded exercises (artifacts)

1. Implement softmax + RMSNorm + a tiny GEMM in Triton. Dump TTIR/TTGIR/LLVM/PTX (TRITON_KERNEL_DUMP / compile-debug). Artifact: annotated IR dump.
2. Autotune study. Same Triton GEMM, 20 configs. Show the compiler is not magic. Artifact: heatmap of BLOCK_M/N/K × num_warps × num_stages.
3. torch.compile a 2-layer MLP + RMSNorm. TORCH_COMPILE_DEBUG=1, identify fused Triton kernels, force a graph break, fix it. Artifact: before/after kernel count.
4. CUDA Graph a decode-step (fixed batch, padded seq). Measure launch overhead vs. eager. Artifact: Nsight Systems with/without graph.
5. Port your RMSNorm to Helion and TileLang. Same unit tests. Artifact: three implementations, one bench table.
6. ThunderKittens GEMM or attn tutorial kernel. Read the 16×16 tile + warp-group template. Artifact: modified kernel (e.g. fused bias).
7. Write a one-page RFC: “We should implement X in Y, not Z,” with ncu evidence.

Failure modes / interview questions

• “Triton is as fast as CUTLASS” as a slogan (true for some ops, false for FA3-class Hopper attn).
• Autotune on the wrong keys (forgot M,N,K in triton.autotune).
• CUDA graphs + dynamic shapes = silent recapture explosion.
• Interview: “Trace Triton from @jit to SASS. What does num_stages mean? When do you drop to CuTe? Why did vLLM delete many custom kernels in favor of compile + a few remaining CUDA kernels?”

Level-up bar

• You can read generated Triton PTX and say whether TMA or ld.global was used.
• You can take an Inductor-generated kernel and beat it by ≥15% with a 50-line Triton change or explain why you cannot.
• You have a written decision tree (your RFC) you would actually use at work.

LAYER 5 — Distributed / multi-GPU kernels

Deep hours: 200–280. Single-GPU peak is table stakes. Training/inference at frontier scale is overlap.

Prerequisites

Layer 3. You understand FA online softmax (you will merge LSE across ranks). Access to ≥2 GPUs; 8×H100 + NVLink is the real classroom.

Resources (12)

1. NCCL overview — https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/overview.html Developer page — https://developer.nvidia.com/nccl Source — https://github.com/NVIDIA/nccl “Demystifying NCCL” — https://arxiv.org/html/2507.04786v1
2. Megatron-LM / Megatron Core — https://github.com/NVIDIA/Megatron-LM Original paper — https://arxiv.org/abs/1909.08053 Sequence-parallel / less rematerialization — https://arxiv.org/abs/2205.05198 Parallelism guide — https://docs.nvidia.com/megatron-core/developer-guide/latest/user-guide/parallelism-guide.html
3. Ring Attention — https://arxiv.org/abs/2310.01889 Implementation — https://github.com/zhuzilin/ring-flash-attention
4. DeepSpeed-Ulysses — https://arxiv.org/abs/2309.14509 Blog/code — https://github.com/deepspeedai/DeepSpeed/blob/master/blogs/deepspeed-ulysses/README.md
5. USP (Ulysses × Ring) — https://arxiv.org/abs/2405.07719 Code — https://github.com/feifeibear/long-context-attention
6. DeepEP (MoE dispatch/combine; V2 = NCCL Gin, less NVSHMEM) — https://github.com/deepseek-ai/DeepEP
7. Transformer Engine TP/SP/CP + quantized all-gather — TE docs above; NVFP4 columnwise AG notes.
8. FlashInfer parallel attention (Ulysses + Ring wrappers) — https://github.com/flashinfer-ai/flashinfer
9. NVLink / NVSwitch product context — Hopper in-depth blog (above); Blackwell NVLink 5 on https://www.nvidia.com/en-us/data-center/technologies/blackwell-architecture/
10. GPU MODE NCCL lectures — https://github.com/gpu-mode/lectures
11. Llama-3-style CP (in ring-flash-attention as llama3_flash_attn_varlen_func) — Llama 3 report https://arxiv.org/abs/2407.21783
12. PMPP 5th ed. multi-GPU / NCCL / NVSHMEM chapters (if you have it).

Concepts you must own

• Collectives: all-reduce = reduce-scatter + all-gather. TP matmul uses that identity to shard activations (Megatron sequence-parallel).
• TP attention/MLP: column-parallel QKV / row-parallel out; 2 all-reduces per layer fwd (or RS+AG with SP).
• CP / sequence-parallel attention:
• Ulysses: all-to-all heads ↔ sequence; degree ≤ #heads.
• Ring: P2P KV + online-softmax merge; unlimited degree, causal imbalance.
• Zigzag: rank r owns chunk r and 2·cp−1−r.
• USP: 2D mesh, Ulysses rows × ring columns.
• Overlap: CUDA_DEVICE_MAX_CONNECTIONS=1 for Megatron SP; NCCL on a side stream; DeepEP SM-limited comm (V2: 4–6 SMs vs 24).
• MoE: dispatch/combine all-to-all; grouped GEMM; Mega-MoE fuses dispatch+lin1+SwiGLU+lin2+combine.

Graded exercises (artifacts)

1. NCCL playground. all-reduce, all-gather, reduce-scatter, send/recv on 2–8 GPUs. Sweep bytes. Artifact: busbw vs. algobw table; identify ring vs. tree.
2. Megatron-style TP Linear. Implement ColumnParallel + RowParallel in PyTorch + NCCL. Match a single-GPU Linear. Artifact: repo.
3. Sequence-parallel RMSNorm. Shard sequence, no extra comm vs. a naive replicated norm. Artifact: memory table.
4. Ring attention (forward) on 2–4 GPUs using FA as the block kernel + LSE merge. Causal + zigzag. Artifact: equal work per rank (nsys).
5. Ulysses attention (all-to-all on QKV, local FA, all-to-all back). Artifact: compare to ring at the same world size.
6. USP 2×2 (2 Ulysses × 2 ring). Artifact: one script + a comm-volume calculation that matches nsys.
7. Overlap study. GEMM on compute stream, all-reduce on comm stream, vs. serial. Artifact: timeline showing hidden comm or an honest failure (too-small tensor).
8. MoE dispatch toy. Token permutation + grouped GEMM + unpermute; then read DeepEP’s dispatch/combine and write a 2-page map of the API. Artifact: note.

Failure modes / interview questions

• Using all-reduce where reduce-scatter suffices (2× traffic).
• Ring causal without zigzag → last rank does ~2× work.
• Ulysses degree > #heads (or > #KV heads for GQA).
• NCCL + compute on the same stream; “overlap” that is fictional.
• Forgetting that CP attention must merge global LSE, not concatenate outputs.
• Interview: “Derive that all-reduce = RS + AG. Draw TP for a Llama MLP. Why CUDA_DEVICE_MAX_CONNECTIONS=1? Compare Ulysses vs ring vs USP communication volume. How does zigzag rebalance causal FA? What does DeepEP overlap with?”

Level-up bar

• TP=2 Linear matches single-GPU numerically (1e-4) and is not 2× slower than the shard math predicts (you hid the collective).
• Ring/zigzag FA on 4 GPUs, seq=8k+, ≥ 70% of single-GPU FA throughput × 4 after accounting for causal density (or you can explain the residual with nsys: wait on P2P, not kernel).
• You can read Megatron tensor_parallel/ and zhuzilin zigzag varlen and narrate the LSE merge.

Sequence (do not rearrange)

L0  C++/arch/perf          220–280h
L1  CUDA/PTX/Hopper/ncu    320–400h
L2  GEMM/CUTLASS/Triton    280–360h
L3  FA/norm/RoPE/paged/MoE 360–450h
L4  compilers/DSLs         160–220h   (start Triton in L2; L4 is internals + choice)
L5  NCCL/TP/CP/EP          200–280h   (start NCCL playground during L3)
────────────────────────────────────
Total                      ~1540–1990 deep hours

Calendar: ~12–18 months at 25–30 deep h/week with GPU access; faster if you already have L0. Triton GEMM (L2) and NCCL hello-world (L5) overlap; do not skip L1 SASS for L4 DSLs.

Weekly rhythm that works: 60% writing kernels, 25% reading (papers + other people’s SASS), 15% profiling notes. If you are only reading, you are not becoming a kernel engineer.

The 20 repos to clone and study (not black-box)

Study means: build, run a unit test, cuobjdump one kernel, write a note on the pipeline.

Honorable (21–26, study if you have time): https://github.com/NVIDIA/TensorRT-LLM , https://github.com/Dao-AILab/quack , https://github.com/FasterDecoding/Medusa , https://github.com/SafeAILab/EAGLE , https://github.com/siboehm/SGEMM_CUDA , https://github.com/NVIDIA/cuda-python .

Interview loop (what kernel teams actually ask)

Expect a whiteboard algorithm (online softmax / FA2 tiling / TP split) plus a take-home or live kernel (RMSNorm, GEMM tile, or decode attention). They will ask you to read someone else’s SASS or ncu and talk. Portfolio that gets a callback:

1. Boehm-style GEMM worklog on H100 with TFLOPS table.
2. A Triton or CUDA attention that you can defend against FA2.
3. One fused memory-bound kernel (RMSNorm or CE) with a roofline.
4. A short note on TMA/WGMMA or tcgen05 showing you read the PTX.

If you cannot get an H100, do Ampere thoroughly and write the Hopper note from docs + Colfax + FA3 paper. Say so. Fake H100 numbers are worse than no numbers.

What “ready for Anthropic / OpenAI / DeepMind kernel team” means

Not “I completed a MOOC.” All of the following:

• You have written (not wrapped) a tensor-core GEMM within ~20% of cuBLAS on a datacenter GPU.
• You can re-derive FA2 and explain FA3’s three ideas (warp-specialized TMA overlap, GEMM/softmax interleave, FP8 block quant) and have a non-toy implementation of at least FA2-style forward.
• You can profile with Nsight Compute/Systems and change the kernel in the direction the counters indicate.
• You can read CUTLASS CuTe and FlashAttention C++ without panic.
• You know when not to write CUDA.
• You can reason about TP/CP/EP traffic and overlap.

That is the no-shortcut path. Start Layer 0 this week with CSAPP Chapter 6 + a perf-instrumented saxpy. Every later layer is the same loop: correctness harness → roofline hypothesis → kernel → ncu → next bottleneck.