Flash Attention
Online softmax, IO-awareness, and the kernel that defined the era.
The UW note beats the FlashAttention paper as the entry point because the whole kernel falls out of one algebraic trick, the online-softmax rescaling, and the note derives it in a few pages where the paper assumes it. Do the algebra by hand; the Triton kernel is that algebra transcribed, plus tiling. The CS149 CPU rung in between is the best pedagogical bridge anyone has built. Watching the N×N intermediate shrink from megabytes to kilobytes turns “IO-awareness” from a slogan into something you saw happen, and that is the paper's actual claim: fewer HBM reads, not fewer FLOPs.
In the lineage, read FA2 the most carefully (the work-partitioning fixes are where the practical speedup lives) and FA3 for what Hopper-era hardware demands: warp specialization and TMA. The back half of the quest is built to leave artifacts. Your kernel inside your engine with the probe still green, a Boehm-style worklog, a leaderboard submission — these are what kernel-team interviews actually ask to see.
- PyTorch Developer PodcastEdward Z. Yangpodcast↗
The cleanest derivation of the online-softmax trick. Do the algebra yourself.
Stanford CS149's ladder: naive → blocked matmul → fused rows (watch the N×N intermediate shrink from MBs to KBs) — the teaching moment between the derivation and the Triton kernel.
FA2 most carefully (work partitioning); FA3 for Hopper warp-specialization. It's about HBM reads, not FLOPs.
Causal, tiled, online-softmax — never materializing the score matrix. Graded on correctness (incl. a seq-8192 case that OOMs naive approaches) and ≥1.5× speedup over materialized attention.
kernel+300 XPauto-verifiedFlashInfer (MLSys '25 best paper — the attention library inside vLLM/SGLang), Colfax's CUTLASS tutorials, and the “GPUs Go Brrr” ThunderKittens post.
TMA + warp specialization (Colfax), CuTe DSL (the GPU MODE leaderboard's NVIDIA competitions are CuTeDSL-shaped now), Marlin/Machete-style W4A16 GEMM, and a taste of ROCm/HIP — three leaderboard comps run on MI300X.
Swap your Triton kernel into the Phase 2 engine's attention path. It isn't done until the OpenAI-conformance probe still passes with YOUR kernel serving the tokens.
build+150 XPBoehm-style: the kernel, the profiler evidence, each optimization step with measured numbers. This genre of post is a known door-opener.
write+200 XPauto-verifiedLeaderboard placements are a real 2026 hiring signal — NVIDIA teams compete on it themselves.