02 — Kernels Papers¶

Phase alignment: Months 3–9 (Phase 2/3). Read these while writing the SGEMM ladder and your Triton FA2.

These papers explain why kernels are the way they are, and how the current bag of tricks was assembled. Do not read them before you’ve written a naive attention kernel and felt how slow it is — the papers only make emotional sense after that pain.

1. Roofline: An Insightful Visual Performance Model (Williams, Waterman, Patterson)¶

No arxiv — CACM 2009 · [EASY] · Prereqs: none

URL: https://dl.acm.org/doi/10.1145/1498765.1498785 (paywalled) · free author copy at https://people.eecs.berkeley.edu/~kubitron/cs252/handouts/papers/RooflineVyNoYellow.pdf

Key trick: Any kernel’s max achievable performance is min(peak_flops, arithmetic_intensity × peak_bandwidth) — a two-line graph that tells you whether to optimize compute or memory.

What to extract:

Arithmetic intensity = FLOPs / bytes moved. Compute it for GEMM (M·N·K FLOPs vs M·K+K·N+M·N bytes), for elementwise ops (1 FLOP / ~4 bytes), and for attention.
The ridge point of the H100 SXM: 989 TFLOPs BF16 ÷ 3.35 TB/s ≈ 295 FLOPs/byte. Anything under 295 is memory-bound.
The 3090’s ridge (142 TF FP16 ÷ 936 GB/s) ≈ 152. The 4090’s ridge (330 TF FP16 ÷ 1008 GB/s) ≈ 327.
This paper is the whole roadmap in 12 pages. Reread it after each phase.

2. Online normalizer calculation for softmax (Milakov & Gimelshein)¶

arxiv:1805.02867 · NVIDIA, 2018 · [EASY] · Prereqs: high-school algebra

Key trick: Compute softmax in one streaming pass over the input by maintaining a running max m and running sum d, rescaling d by exp(m_old − m_new) whenever the max updates.

What to extract:

Derive the rescaling identity on paper. This is the mathematical seed of FlashAttention.
Understand what “streaming” buys: you no longer need N elements in memory at once, only 2 scalars.
Numerical stability: without the max subtraction, exp(500) overflows fp16 (max ~65504) and bf16 (max ~3.4e38). Fp8-e4m3 max is 448. You will burn yourself on this later.
Implement it in numpy in 15 minutes, then in Triton in an hour.

3. FlashAttention (v1): Fast and Memory-Efficient Exact Attention with IO-Awareness¶

arxiv:<phone_number_or_numberic_id_or_random_id_21> · Dao et al., 2022 · [HARD] · Prereqs: papers #1, #2, and a working naive attention kernel

Key trick: Tile Q/K/V into SRAM-sized blocks and fuse QKᵀ → softmax → PV in one kernel using online-softmax rescaling, cutting HBM traffic from O(N²) to O(N²·d/M) where M is SRAM size.

What to extract:

The IO complexity analysis. Section 3.2. Derive it yourself: naive attention reads Q,K,V,S once each (4·N·d + N² bytes); FA reads Q once, K and V O(N·d/M) times = O(N²·d²/M) total bytes when d < M. This is where the speedup comes from — not FLOPs, IO.
The tiling scheme: block sizes Bc (columns of K) and Br (rows of Q), constraints Bc·d ≤ M/4.
Recomputation in backward. You don’t care about backward for inference, but understand why storing only the O(N) softmax stats (m, ℓ) is enough.
The algorithm (Algorithm 1) fits on a page. Write it out from memory before continuing.

Implement it: you will do this in Triton in Phase 3. Do not read FA2 until you’ve tried FA1.

4. FlashAttention-2: Faster Attention with Better Parallelism and Work Partitioning¶

arxiv:<phone_number_or_numberic_id_or_random_id_22> · Tri Dao, 2023 · [HARD] · Prereqs: #3

Key trick: Swap the outer loop (Q outside, K/V inside), parallelize across the sequence dimension of Q (not just batch·heads), and reduce non-matmul FLOPs by delaying the rescaling by 1/ℓ to the epilogue.

What to extract:

Why the outer-loop-Q ordering: one Q tile is loaded once and streams all K/V. The output tile is finalized in-place. Contrast to FA1 which had K/V outside and needed to update all of O.
Sequence-dimension parallelism unlocks utilization on long-context single-batch decoding. This is why FA2 lifted the ceiling for the local-LLM crowd.
Warp partitioning: each warp handles a slice of the K/V dimension (“split-K”), not the sequence dimension — avoids expensive inter-warp reduction.
Non-matmul FLOP reduction: keep ℓ, m; only divide once at the end. Small in count, big in perf because non-matmul FLOPs on Ampere/Hopper use non-tensor-core paths at ~16x lower peak.
The A100 numbers: FA2 reaches 50–73% of theoretical max FLOPs. That’s the ceiling to compare your Triton impl to.

5. FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-Precision¶

arxiv:<phone_number_or_numberic_id_or_random_id_23> · Shah, Bikshandi, Zhang, Thakkar, Ramani, Ré, Dao, 2024 · [HARD] · Prereqs: #4, some Hopper knowledge

Key trick: On H100, use TMA async copies + wgmma async matmuls + warp specialization (producer/consumer) with a 2-stage pingpong scheduler to overlap softmax with the next matmul, and add FP8 support with incoherent processing for outlier control.

What to extract:

Warp specialization: consumer warps compute matmuls; producer warps issue TMA loads. They synchronize via bar.sync and mbarrier. This is the Hopper programming model.
Pingpong scheduling: two warpgroups take turns doing softmax vs. GEMM — while warpgroup A does softmax on tile i, warpgroup B does GEMM on tile i+1. Overlaps the non-matmul stall.
FP8 support with incoherent processing: apply a random Hadamard rotation before quantization to spread outliers — this is a numerics trick borrowed from QuIP# and applied to attention.
The H100 achieved TFLOPs: 740 fp16, 1.2 PFLOPs fp8. Approaches 75% of peak.
If you have no H100, read this paper anyway. It teaches you what Hopper actually does.

Read alongside: ThunderKittens “GPUs Go Brrr” blog (https://hazyresearch.stanford.edu/blog/2024-05-12-tk) — same crowd, same lessons, more accessible than the paper.

6. Flash-Decoding for long-context inference (Tri Dao blog + FlashAttention repo)¶

No paper. Blog: https://pytorch.org/blog/flash-decoding/ (Oct 2023) · [MEDIUM] · Prereqs: #4

Key trick: At decode time Q is a single token, so parallelizing over the Q sequence (FA2’s trick) does nothing. Instead split the KV dimension across many CUDA blocks, each producing a partial softmax, then reduce.

What to extract:

Why the FA2 kernel is bad at batch-1 decode with long context: only 1 Q row × few heads gives too few blocks to fill the SMs.
The split-KV algorithm: chunk K/V into S splits, run FA-like fused attention per split, keep (o_i, m_i, ℓ_i) per split, do one reduction kernel.
The choice of S: memory-bandwidth-utilization sweet spot. FlashInfer generalizes this.

Implement it (later): after your Triton FA2, add a split-KV path. This is the actual kernel your mini-engine will use for decode.

7. FlashInfer: Efficient and Customizable Attention Engine for LLM Inference Serving¶

arxiv:<phone_number_or_numberic_id_or_random_id_24> · Ye et al., 2025 · MLSys 2025 Best Paper · [HARD] · Prereqs: #6, PagedAttention (Phase 4)

Key trick: A single kernel library serving prefill, decode, append, paged, and ragged attention across sm70–sm90+, driven by a compile-time policy language that picks tile shapes and split strategies per shape at plan time.

What to extract:

The “plan-and-run” API: plan(kv_indices, kv_indptr, ...) builds a schedule (which block gets which KV chunk) once; run(q) executes it every step. Amortize the planning cost across decode steps.
Sparse and ragged KV: not every request has the same length — the paged storage means each request’s KV is a scatter of blocks. FlashInfer handles this natively.
Composable formats: paged, ragged, sliding-window, tree (for speculative decoding verification), custom masks.
Why vLLM and SGLang both adopted it as their attention backend — having one team maintain the attention kernel frees the engine teams.

Read the code: https://github.com/flashinfer-ai/flashinfer. Study include/flashinfer/attention/scheduler.cuh.

8. Triton: An Intermediate Language and Compiler for Tiled Neural Network Computations (Tillet et al.)¶

MAPL 2019 · https://dl.acm.org/doi/10.1145/3315508.3329973 · [MEDIUM] · Prereqs: some CUDA

Key trick: A tile-level programming model — you write block-level programs (single-instruction on tiles) and the compiler handles the warp / thread / shared-memory / coalescing details.

What to extract:

The mental model: pid = program_id(0), offs = pid*BLOCK + arange(BLOCK), mask = offs < N, tl.load(ptr+offs, mask=mask). You write index arithmetic; the compiler emits SASS.
What Triton cannot do (as of 2026): fine-grained thread-level control (e.g., non-standard warp-specialization patterns — though Triton is closing this gap fast with Gluon).
Autotuning: @triton.autotune searches over BLOCK_M, BLOCK_N, BLOCK_K, num_warps, num_stages for you. Learn to use it after you can reason about the choices.
Read the Triton FA2 tutorial (github.com/triton-lang/triton/blob/main/python/tutorials/06-fused-attention.py) with this paper open.

Companion (much better than the original paper for pedagogy): https://triton-lang.org/main/getting-started/tutorials/index.html — the official tutorial series is the actual curriculum.

9. Volkov — Better Performance at Lower Occupancy (GTC 2010)¶

No arxiv — GTC talk · [EASY] · Prereqs: some CUDA

URL: https://www.nvidia.com/content/GTC-2010/pdfs/2238_GTC2010.pdf

Key trick: More registers per thread + fewer threads per SM (lower occupancy) can increase throughput because instruction-level parallelism hides latency better than thread-level parallelism.

Why it’s in the canon: Every CUDA tutorial tells you to maximize occupancy. Every actual fast kernel author knows that’s often wrong. Volkov’s slides are the classical debunking. FlashAttention uses lots of registers per warp; so does Marlin. That’s not accident.

What to extract:

The register pressure vs. occupancy tradeoff.
Why hiding memory latency needs either many warps or many independent instructions per warp — not necessarily many warps.
Compile with --ptxas-options=-v to see register usage; watch what it costs you in occupancy.

10. CUTLASS 3.x / CuTe layout algebra (documentation, not paper)¶

No paper. Docs: https://github.com/NVIDIA/cutlass/blob/main/media/docs/cpp/cute/00_quickstart.md · [HARD] · Prereqs: papers #1–8

Not a paper, but you must read the CuTe layout algebra docs. Every serious Hopper/Blackwell kernel is written in CuTe. The tensor/layout algebra is the language.

What to extract:

Layout = (Shape, Stride). Tensors have layouts, not just shapes.
Composition, complement, coalescing operations on layouts.
How MMA atoms describe tensor-core operations.
Read the examples/cute/tutorial/ files in the CUTLASS repo.

Budget a week of evenings. This is the hardest reading in the kernel section, and it pays only in Phase 5+.

Optional / awareness-level¶

StreamK (arxiv:<phone_number_or_numberic_id_or_random_id_25>) — load-balanced GEMM decomposition across SMs. Reference for the Marlin/Machete lineage.
ThunderKittens (arxiv:<phone_number_or_numberic_id_or_random_id_26>) — 100-line FA2 in a tile DSL. Read the blog first, paper only if curious.
Mirage / Mega Kernels (arxiv:<phone_number_or_numberic_id_or_random_id_27>) — fusing a whole transformer layer into one kernel via superoptimization. 2026 frontier.

The whiteboard test for this set¶

After finishing this section you should be able to, in 10 minutes on a whiteboard:

Draw the roofline of an H100 with the ridge point marked, and place naive attention, FA2 prefill, FA2 decode, and matmul on it.
Derive the online softmax rescaling identity from scratch.
Draw the FA2 tiling scheme with correct outer/inner loop and identify where the O tile is finalized.
Explain why decode needs split-KV (flash-decoding) and prefill does not.
Write a Triton pseudo-kernel for a fused RMSNorm+residual and identify its arithmetic intensity.

If any of those five fails, you’re not done. Reread.