The FlashAttention lineage — FA1 / FA2 / FA3 / FA4

FlashAttention isn’t a single algorithm; it’s a lineage of increasingly hardware-specific refinements. Each generation targets a specific NVIDIA architecture and squeezes ~1.5–2× out of the previous. Understanding the lineage tells you what “attention performance” actually means on each hardware.

The one-glance table

Version

arXiv

Author(s)

Year

Target HW

Peak achieved (BF16)

Key innovation

FA1

<phone_number_or_numberic_id_or_random_id_58>

Dao, Fu, Ermon, Rudra, Ré

2022

Ampere (A100)

~ 3× stdatt

Tiling + online softmax + recompute-in-backward

FA2

<phone_number_or_numberic_id_or_random_id_59>

Dao

2023

Ampere/Hopper

~ 230 TFLOPs (A100)

Parallelize seq dim + fewer non-matmul FLOPs + better warp partitioning

FA3

<phone_number_or_numberic_id_or_random_id_60>

Shah, Bikshandi, Dao et al.

2024

Hopper (H100)

840 TFLOPs (85% util)

Warp-specialization + async TMA/wgmma + pingpong overlap + FP8 (1.3 PFLOPs)

FA4

<phone_number_or_numberic_id_or_random_id_61>

(Colfax + NVIDIA + Meta)

2026

Blackwell (B200)

1613 TFLOPs (71% util)

Fully async MMA + software exp + 2-CTA MMA + TMEM; 2.7× Triton, 1.3× cuDNN; written in Python via CuTe-DSL

FA1 (2022) — the founding paper

arXiv: <phone_number_or_numberic_id_or_random_id_53> — FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness. Authors: Tri Dao, Daniel Y. Fu, Stefano Ermon, Atri Rudra, Christopher Ré. Venue: NeurIPS 2022.

The three ideas:

  1. Tiling. Compute attention block-by-block so intermediate S = QK^T never lives in HBM — only the current tile lives in SMEM.

  2. Online softmax. Enables (1) by maintaining (m, d, o) state per Q-row across tiles. See 01_online_softmax.md.

  3. Recomputation in backward. Instead of saving the N×N attention matrix for the backward pass (which would defeat the point), recompute it on the fly. Trades FLOPs for HBM.

Result: ~ 3× faster attention on BERT, 15% faster GPT-2 training end-to-end. Fits 2–4× longer sequences. IO-complexity O(N²d²/M) beats standard O(N²d + Nd²) when d < M.

Prerequisites for reading: you must have derived online softmax on paper first (01_online_softmax.md). Otherwise the recurrence in §3 will feel like magic.

When to read: week 1 of Phase 3.

FA2 (2023) — the algorithm you will actually implement

arXiv: <phone_number_or_numberic_id_or_random_id_54> — FlashAttention-2: Faster Attention with Better Parallelism and Work Partitioning. Author: Tri Dao. PDF: https://tridao.me/publications/flash2/flash2.pdf

Three refinements over FA1:

  1. Parallelize across the sequence dimension. FA1 parallelized over batch × heads only. For long context with small batch, that starved the SMs. FA2 also parallelizes across Q-blocks, so an SM can be working on a different chunk of the same sequence. Huge for long-context.

  2. Reduce non-matmul FLOPs. In FA1, softmax renormalization ran per-tile with divides. FA2 defers the final divide by d to the end of the loop, saving flops proportional to number of tiles. On tensor-core hardware, non-matmul FLOPs are the bottleneck once matmul is fast.

  3. Better warp partitioning. FA1 used a “split-K” partitioning inside each Q-block that required inter-warp SMEM communication. FA2 uses “split-Q” partitioning — each warp owns Q-rows independently, no cross-warp softmax stitching. Cuts SMEM traffic significantly.

Result: ~2× FA1, 230 TFLOPs on A100 (72% of matmul peak). This is the algorithm you will re-implement in Triton in 04_writing_fa_in_triton.md.

When to read: the moment your FA1-shaped Triton kernel works. Understand the three refinements and integrate them into your kernel. Then your kernel is legitimately “FA2 in Triton.”

FA3 (2024) — the Hopper masterclass

arXiv: <phone_number_or_numberic_id_or_random_id_55> — FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-precision. Authors: Jay Shah, Ganesh Bikshandi, Ying Zhang, Vijay Thakkar, Pradeep Ramani, Tri Dao. Venue: NeurIPS 2024. Blog: https://tridao.me/blog/2024/flash3

Three ideas, all Hopper-specific:

  1. Warp-specialization + async TMA/wgmma. Instead of all warps doing all work, dedicate a producer warp group to issuing TMA loads and a consumer warp group to running wgmma matmuls. Producer/consumer pipeline overlapped through Hopper’s async barrier machinery.

  2. Pingpong GEMM/softmax overlap. Two Q-block computations run “in phase” on the same SM — while one is doing the matmul (tensor cores busy), the other is doing softmax (CUDA cores busy). Hides softmax latency behind matmul.

  3. FP8 with block quantization + incoherent processing. Attention with FP8 (E4M3) inputs, block-quantized. Uses “incoherent processing” (random rotation via Hadamard) to spread outliers, reducing quant error 2.6× vs baseline FP8.

Result:

  • BF16: 1.5–2× FA2 → 840 TFLOPs/s on H100 (85% utilization).

  • FP8: 1.3 PFLOPs/s at 2.6× lower error than naive FP8.

When to read: after your FA2-in-Triton is done and you’ve done a Hopper GEMM read (Colfax notes + a CUTLASS SM90 example). The paper is a case study in Hopper primitives; you’ll get almost nothing from it without that hardware context first.

Not to implement. You will read + understand FA3, not re-derive it. Reference implementation lives at https://github.com/Dao-AILab/flash-attention (hopper/ subdirectory).

FA4 (2026) — Blackwell, written in Python

arXiv: <phone_number_or_numberic_id_or_random_id_56> — FlashAttention-4 (Colfax + NVIDIA + Meta authorship). Preview at Hot Chips 2025; formal paper March 2026. Reverse-engineering breakdown: Modal, https://modal.com/blog/reverse-engineer-flash-attention-4. Lambda writeup: https://lambda.ai/blog/flashattention-4-gives-the-nvidia-blackwell-platform-its-most-optimized-attention-kernel-yet. GPU MODE lecture: How FlashAttention 4 Works, Charles Frye, 2 Oct 2025, https://www.youtube.com/watch?v=ZIEq-WTquy4 — mandatory.

The five big changes (each Blackwell-specific):

  1. Fully async MMA pipeline. Every matrix multiply is async (tcgen05.mma on Blackwell); loads and stores overlap fully with compute. Larger tile sizes than FA3.

  2. Software-emulated exp + conditional softmax rescale. Non-matmul FLOP reduction pushed further — approximates exp with a polynomial in software when accuracy allows, and only rescales the online-softmax state when the running max actually changes.

  3. TMEM (Tensor Memory) exploitation. Blackwell’s dedicated 256 KB/SM tensor memory reduces SMEM pressure and lets tile sizes grow.

  4. 2-CTA MMA mode in backward. Blackwell’s ability to have two cooperating CTAs share a single MMA operation is used in FA4’s backward pass to double effective tile size.

  5. Written in Python via CuTe-DSL. Not CUDA C++. Frontier attention kernel authored in Python. Signals the future of kernel authoring.

Result:

  • 1613 TFLOPs/s on B200 BF16 (71% utilization).

  • 2.7× faster than a Triton FA implementation on the same hardware.

  • 1.3× faster than cuDNN 9.13’s attention kernel.

The bigger “signal”:

  • CuTe-DSL is the new frontier tool. Not Triton, not CUTLASS C++. If Blackwell/next-gen kernels are your ceiling ambition, learn CuTe-DSL.

  • Non-matmul reduction still pays. Even at 5th-gen tensor cores, the softmax rescale can dominate; software exp + conditional rescale still matters.

  • Blackwell primitives (TMEM, 2-CTA, tcgen05) are the new load-bearing set — wgmma/TMA are now “previous generation.”

When to read: end of Phase 3. Read Modal’s blog + watch Charles Frye’s lecture. Read the arxiv paper only if you have Blackwell HW access; otherwise it’s aspirational.

The mental model — what each generation actually optimizes

Generation

The one thing it exploits

FA1

HBM traffic (via tiling + online softmax + recompute)

FA2

SM utilization (via seq-dim parallelism + non-matmul FLOP reduction)

FA3

Async pipelines (via warp specialization + pingpong + FP8)

FA4

Blackwell-native primitives (TMEM + 2-CTA + fully-async MMA + polynomial exp)

Each step is: “exploit a specific new hardware primitive.” Reading them in order teaches you the modern GPU story.

Reference implementations to read (in order)

  1. Triton fused-attention tutorial — https://triton-lang.org/main/getting-started/tutorials/06-fused-attention.html — shortest FA2-style kernel you’ll find. Read this.

  2. Dao-AILab flash-attention repo — https://github.com/Dao-AILab/flash-attention — canonical CUDA/CUTLASS implementations. Look at the csrc/flash_attn/kernels/ tree for FA2 and hopper/ for FA3.

  3. CUTLASS Hopper attention examples — in examples/ of the CUTLASS repo.

  4. Modal FA4 reverse-engineering — https://modal.com/blog/reverse-engineer-flash-attention-4 — walks through the CuTe-DSL source line by line.

References