InferQuest
← Quest map
Phase 4 · Kernelcraft

Flash Attention

Online softmax, IO-awareness, and the kernel that defined the era.

0/9 · 1230 XP
You’re previewing the curriculum. Sign in (top right) to track progress, take knowledge checks, and unlock the verifiers.
Tasks
  • The cleanest derivation of the online-softmax trick. Do the algebra yourself.

    read+60 XPresource ↗
  • 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.

    build+100 XPresource ↗
  • FA2 most carefully (work partitioning); FA3 for Hopper warp-specialization. It's about HBM reads, not FLOPs.

    paper+100 XPresource ↗
  • 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-verified
  • FlashInfer (MLSys '25 best paper — the attention library inside vLLM/SGLang), Colfax's CUTLASS tutorials, and the “GPUs Go Brrr” ThunderKittens post.

    read+80 XPresource ↗
  • 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.

    read+90 XPresource ↗
  • 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 XP
  • Boehm-style: the kernel, the profiler evidence, each optimization step with measured numbers. This genre of post is a known door-opener.

    write+200 XPauto-verified
  • Leaderboard placements are a real 2026 hiring signal — NVIDIA teams compete on it themselves.

    kernel+150 XPresource ↗