Optimizing FlashAttention-4 for Efficient Inference: Enhancements and Performance Metrics

The release of the FlashAttention-4 kernel last year sparked significant interest, particularly regarding its utility for large language model (LLM) inference tasks. Recent developments have focused on enhancing this kernel's performance, specifically targeting scenarios characterized by memory bandwidth limitations and variable workload factors typical in decoding phases.

Key Enhancements for Inference Optimization

Our contributions can be broadly categorized into two strategies: optimizing parallel execution strategies and improving handling of global memory accesses. Both strategies are critical for increasing the throughput of inference tasks.

• Adjusting Parallelism Strategy: This entails modifying the number of query tiles that can be processed concurrently and transitioning from query parallelism to key/value (KV) parallelism. The impact of this change is substantial, especially in environments that benefit from modern massively parallel architectures.
• Supporting Irregular Global Memory Access: The approach includes employing cp.async loads instead of cp.async.bulk loads to make full use of the Tensor Memory Accelerator (TMA) capabilities present in the architecture.

The following figures illustrate the improvements made through these optimizations:

Optimizing parallelism strategies often results in the most significant performance gains. A flexible approach to parallelization allows for adjustments that can shift workload distribution within the algorithm’s components, which is generally more effective than merely speeding up a fixed parallel component.

Evaluation of Kernel Contributions

We organized our enhancements by specific pull requests (PRs), each associated with a "Figure of Merit"—a specific measurement indicating performance improvements. Below is a summary of key contributions:

PR 2109: Support FP8 Inputs (Merged April 17, 2026)

Figure of Merit: Up to 1.16x throughput relative to bf16 baseline

| Batch Size / Seq Len | BF 16 TFLOP/s | FP8 TFLOP/s | Speedup |
| -------------------- | ------------- | ----------- | ------- |
| 1 / 16384            | 1569          | 1818        | 1.13x   |
| 32 / 512             | 962           | 1090        | 1.16x   |

Typically, training requires higher precision for accumulating small gradient changes. However, during inference, lower precision can yield significant efficiency gains. For instance, utilizing FP8 instead of bf16 results in reduced memory and arithmetic demands, directly resulting in improved throughput metrics. This change is beneficial specifically for certain layers like MLPs or MoEs, which may still need higher precision for gradients but can afford lower precision during inference.

PR 1999 and PR 2104: Support Arbitrary KV Page Sizes (Merged November 13, 2025) and Performance Optimization (Merged January 15, 2026)

Figure of Merit: Up to 2.40x throughput for small page sizes

| Page Size | Added in PR 1999? | TFLOP/s, PR 1999 | TFLOP/s, PR 2104 | Speedup |
| --------- | ----------------- | ---------------- | ---------------- | ------- |
| 1         | y                 | 18.56            | 44.57            | 2.40x   |
| 8         | y                 | 31.21            | 42.58            | 1.37x   |
| 32        | y                 | 34.98            | 42.47            | 1.21x   |
| 128       | n                 | 42.11            | 41.96            | -       |

In FlashAttention-4, the original implementation required that KV cache pages were the same size as the processing tiles, which limited flexibility and efficiency. By updating this requirement, we enabled a broader range of page sizes, thereby reducing waste in memory bandwidth during inference. Arbitrary page sizes also allowed for better memory cache usage, reducing internal fragmentation and improving performance overall.

PR 1940: Add Parallelism Across the KV Dimension (Merged November 4, 2025)

Figure of Merit: Up to 4.37x greater throughput for small query lengths

| Number of KV splits | Memory throughput (TB/s) |
| ------------------- | ------------------------ |
| 1 (baseline)        | 0.83                     |
| 2                   | 2.65                     |
| 4                   | 4.30                     |
| 8                   | 4.27                     |
| 16                  | 4.22                     |
| 32                  | 4.37                     |
| 64                  | 4.17                     |
| 128                 | 3.83                     |

Normal inference requests typically process one or a few tokens at a time, often using cached KV values. The original kernel's design focused primarily on query parallelism, resulting in underutilization of GPU resources. By incorporating a mechanism to parallelize across the KV dimension, we increased output generation capabilities significantly, ensuring enhanced throughput during inferencing, especially for low-latency applications.

PR 1993: Reduce Wasted Work for Small Query Sizes (Merged January 8, 2026)

Figure of Merit: Up to 3.06x throughput for single-token decode

| Number of Splits | TFLOP/s | Speedup vs baseline |
| ---------------- | ------- | ------------------- |
| 1                | 1.79    | 1.00x               |
| 2                | 3.39    | 1.89x               |
| 4                | 5.47    | 3.06x               |
| 8                | 5.23    | 2.92x               |
| 16               | 5.08    | 2.84x               |
| 32               | 5.12    | 2.86x               |

Recognizing that many attention operations yield far fewer than the maximum number of expected queries, this enhancement optimizes tile-level processing. The kernel can now efficiently handle scenarios with limited queries, eliminating redundant computations and boosting throughput significantly, especially for single-token decoding tasks.

PR 2186: Speed Up Irregular Q::KV Head Ratios by Extending GQA Packing (Merged March 20, 2026)

Figure of Merit: 2.92x throughput increase for single-token decode

| Pack GQA | TFLOP/s |
| -------- | ------- |
| OFF      | 7.1     |
| ON       | 20.7    |

Our adaptation of Grouped-Query Attention (GQA) allows multiple queries to be processed against each KV entry. However, the original kernel's structure posed challenges in efficiently loading KV values, as it required redundant loads. By implementing GQA packing, we could consolidate query processing under a single tile, significantly enhancing arithmetic intensity while maintaining efficient memory loads.

Looking Ahead: The Future of FlashAttention Enhancements

These refinements to FlashAttention-4 illustrate a concerted effort to optimize the kernel specifically for inference, with performance gains that translate to faster, more efficient processing in real-world applications. With continued exploration and implementation of novel programming paradigms and architectural adjustments, we anticipate further improvements in LLM processing capabilities.

At Modal, our commitment to advancing open-source software extends beyond just the FlashAttention-4 kernel. We are actively collaborating on various projects aimed at reshaping the landscape of inference and training frameworks. As we integrate these software developments into our infrastructure for optimized deployment, we invite others in the field to contribute and explore opportunities with us.