Making FlashAttention-4 faster for inference (modal.com)

Recent enhancements to the FlashAttention-4 kernel have significantly optimized its performance for large language model (LLM) inference tasks, particularly in the decode-heavy phase where memory bandwidth limitations can slow down processing. The refinements focus on adjusting parallelism strategies and accommodating irregular global memory accesses, which are crucial as inference workloads often involve variable batch sizes and sequence lengths. By shifting from query parallelism to key/value parallelism and introducing cp.async loads, the developers have improved throughput notably, with some configurations achieving up to 2.40x speedup for small page sizes. The introduction of support for 8-bit floating-point inputs reduces memory and arithmetic demands during inference, allowing models to maintain quality while enhancing speed. The changes mean more efficient usage of the KV cache, particularly important for speculative decoding where short sequences are common. Additionally, the optimizations streamline the memory access patterns by decoupling address generation from their utilization, resulting in more efficient computations and enabling simpler integration of new algorithms into the inference process. Overall, these developments position FlashAttention-4 as a more powerful and flexible tool for AI applications, facilitating faster and more efficient processing of LLM queries.