Making FlashAttention-4 faster for inference
Jun 11, 2026

Summary
This post details contributions to the FlashAttention-4 kernel to improve LLM inference performance, specifically for decode-heavy workloads. Key changes include adding support for FP8 inputs to reduce memory and arithmetic bandwidth demands, and enabling arbitrary KV page sizes by implementing a path using `cp.async` instructions to handle irregular global memory accesses, which is more efficient for smaller page sizes than the original TMA-based approach. The post also discusses the benefits of adjusting parallelism strategies and the potential for tile-based programming models.