FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-precision

On Hopper, FlashAttention-2's synchronous schedule wastes the tensor cores: the softmax and the matmuls cannot overlap, so the FP8 tensor cores sit idle waiting on non-matmul work.

Hopper-specific async rewrite. (1) Warp-specialisation: partition warps into producer and consumer groups so Tensor Memory Accelerator (TMA) data movement overlaps matmul and softmax. (2) Interleave block-wise matmul with softmax so one runs while the other waits. (3) Block quantisation and incoherent processing to exploit hardware FP8 support. Requires Hopper's wgmma/TMA instructions; the async pipeline has no meaning on Ampere.

Paper's own claim (abstract): 1.5-2.0x speedup on H100 GPUs, FP16 reaching up to 740 TFLOPs/s (75% utilisation) and FP8 reaching close to 1.2 PFLOPs/s; FP8 FlashAttention-3 achieves 2.6x LOWER numerical error than a baseline FP8 attention. Baselines are FlashAttention-2 (FP16/FP8) and cuDNN.

arXiv ID 2407.08608 verified via abs-page meta tags on 2026-10-03; title matches. CORRECTION TO THE SOURCE RECORD: pap-attn-flashattention-3 says 'no venue claim in the arXiv comment' and the dead agent was presumably about to leave venue as arxiv-preprint. OpenReview NeurIPS.cc/2024/Conference records venue='NeurIPS 2024 spotlight' for the exact title. venue=neurips, year=2024. adoption=in-upstream-engine: the flash_attn Python package (Dao-AILab/flash-attention, actively pushed 2026-10-03) ships the FA3 kernels and vLLM's flash_attn backend dispatches to them; vLLM v1 backend list also carries fa_utils.py. The 740 TFLOPs/s figure is Hopper-only, which is why hardware_relevance is H100/H200/H800 and nothing else. VENUE_TRACK BACKFILL 2026-10-04 (paper-fields pass, agent w3p): venue_track 'main'. This is STRUCTURAL rather than researched: neurips has no Findings volume and no workshop track in which a paper could appear, so 'main' is the only possible honest value and no per-paper evidence was needed. The field is annotated here so that 'findings' and 'workshop' remain MEANINGFUL - a filter for findings=true must return only the papers that actually are findings, which requires every other paper to be explicitly NOT findings.