FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-precision
Fields
sources/kern-flashattention-3-blogNotes
Author blog for arXiv 2407.08608. Mechanism claims sourced here: wgmma matters because without it mma.sync reaches only about 2/3 of Hopper tensor-core peak; TMA is a hardware unit that does index calculation and out-of-bounds predication for global->shared copies, freeing registers so tiles can be larger; softmax cannot overlap matmul automatically because H100 SXM5 does 989 TFLOPS FP16 matmul but only 3.9 TFLOPS of special functions (16 ops/SM/clock * 132 SM * 1.83 GHz), and at head_dim 128 there are 512x more matmul than exponential flops; the Hadamard transform with random signs is O(d log d) and can be fused into the RoPE pass for free. Numbers (FA-2 at 35% of peak, FA3 1.5-2.0x, 740 TFLOPS = 75%, ~1.2 PFLOPS FP8, 2.6x lower quantization error) are author-reported, not third-party measured. | AUDIT 2026-10-03 RE-VERIFIED: live. Title 'FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-precision', authors Jay Shah (Colfax), Ganesh Bikshandi (Colfax), Ying Zhang (Meta), Vijay Thakkar (NVIDIA/Georgia Tech), Pradeep Ramani (NVIDIA), Tri Dao (Princeton/Together AI); page states Published July 11, 2024 and links arXiv 2407.08608. Published field 2024-07-11 confirmed against the page.
Cross-references
This record references no other records.