FlashAttention-4: Algorithm and Kernel Pipelining for Efficient Attention on Blackwell GPUs

Fields

id

sources/attn-flashattention-4-paper

URL

https://arxiv.org/abs/2501.09294

Publisher

arXiv

Kind

paper

Published

2025-01-17

Accessed

2026-10-04

notes

FA-4 paper. Key claims: up to 1605 TFLOPs/s BF16 on B200 (71% utilization), 1.3x vs cuDNN 9.13, 2.7x vs Triton. Backward bound by shared-memory bandwidth not FLOPs. Uses tcgen05 UMMA with TMEM accumulators, 2-CTA MMA, emulated exponential (Cody-Waite + Horner polynomial on FMA units). Written in CuTeDSL, optimized for Hopper and Blackwell.

Notes

FA-4 paper. Key claims: up to 1605 TFLOPs/s BF16 on B200 (71% utilization), 1.3x vs cuDNN 9.13, 2.7x vs Triton. Backward bound by shared-memory bandwidth not FLOPs. Uses tcgen05 UMMA with TMEM accumulators, 2-CTA MMA, emulated exponential (Cody-Waite + Horner polynomial on FMA units). Written in CuTeDSL, optimized for Hopper and Blackwell.

Cross-references

This record references no other records.