FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-precision

ArXiv: 2407.08608

🎯 Pitch

FlashAttention-3 achieves up to 1.2 PFLOPs/s on H100 GPUs by creatively overlapping softmax with matrix multiplies—operations previously forced to run sequentially—while simultaneously dropping numerical error 2.6× below a standard FP8 baseline using block quantization.

---

1. Executive Summary

FlashAttention-3 introduces a hardware-aware exact attention algorithm for Hopper GPUs that exploits asynchrony (overlapping data movement and computation via warp-specialized producer-consumer warps, and interleaving block-wise matrix multiplications with softmax operations through a 2-stage pipelining scheme) and low-precision FP8 computation (bridging WGMMA layout constraints via in-kernel transpose, combined with block quantization and incoherent processing to preserve accuracy). On H100 SXM5 GPUs, FlashAttention-3 achieves a 1.5–2.0× speedup over FlashAttention-2 in FP16 forward passes, reaching 740 TFLOPs/s (75% utilization), and up to 1.2 PFLOPs/s in FP8 — while FP8 FlashAttention-3 also reduces numerical error by 2.6× compared to a baseline FP8 attention with per-tensor quantization. The work establishes that dramatic attention efficiency gains from asynchrony and low-precision are achievable while maintaining or improving numerical accuracy, though the largest speedups appear at medium-to-long sequence lengths (1k and above) where the kernel becomes compute-bound rather than launch-overhead-bound.

2. Context and Motivation

The Core Problem: Attention Is Compute-Bound but Grossly Underutilizes Modern Hardware

The fundamental problem FlashAttention-3 addresses is not that attention is inherently slow—it's that existing exact attention implementations, including the state-of-the-art FlashAttention-2, leave enormous amounts of GPU throughput on the table when running on the latest hardware. The paper's motivating observation is stark: FlashAttention-2 achieves only 35% utilization on the Hopper H100 GPU, compared to 80–90% utilization achieved by optimized matrix multiplication (GEMM) kernels on the same hardware. This gap matters because attention is the primary computational bottleneck in Transformer models: computing self-attention scores requires $O(N^2)$ operations with respect to sequence length $N$, making it the dominant cost as models scale to longer contexts.

The significance of this gap extends beyond mere performance optimization. Scaling attention to longer contexts unlocks qualitatively new capabilities that are currently bottlenecked by attention speed: modeling and reasoning over multiple long documents simultaneously (Guo et al., 2021; Peng et al., 2023; Shaham et al., 2022), processing entire large codebases (Li et al., 2023; Rozière et al., 2023), handling high-resolution images (Chen et al., 2022), audio (Gulati et al., 2020), and video (Ho et al., 2022), and supporting agents with long interaction histories (Sun et al., 2019) and extended planning horizons (Yao et al., 2022). The paper argues that faster attention will directly unlock these use cases by making long-context computation practically feasible at scale.

The underlying issue is a resources mismatch between hardware capabilities and algorithmic design. Modern GPUs like the Hopper H100 provide specialized hardware units—Tensor Cores for matrix multiplication, the Tensor Memory Accelerator (TMA) for asynchronous memory transfers—that can operate concurrently with the general-purpose CUDA cores. However, FlashAttention-2's algorithm adheres to what the authors characterize as "a simplified synchronous model" that makes no explicit use of asynchrony in its design. Operations proceed sequentially: compute a block of attention scores, wait; compute softmax on those scores, wait; compute the output update, wait. This serialization means that while Tensor Cores are idle, CUDA cores could be performing the softmax's exponential operations—and vice versa. The hardware capacity is present; the algorithm simply doesn't exploit it.

Why 35% Utilization Matters: The Gap Between Theoretical and Achieved Throughput

To appreciate the severity of 35% utilization, we need concrete numbers. The H100 SXM5 has 989 TFLOPS of FP16 matrix multiplication throughput. At 35% utilization, FlashAttention-2 delivers roughly 346 TFLOPS. The paper's target—75% utilization with FlashAttention-3—delivers 740 TFLOPS, a 1.5–2.0× speedup. This means that for a given sequence length, attention runs in roughly half the time, or equivalently, that models can handle sequences roughly twice as long in the same wall-clock budget without any change to the model architecture or training recipe.

The GPU's theoretical throughput is a function of its memory and compute hierarchies, which the paper carefully details in Section 2.2. The critical observation is the inverse relationship between memory capacity and bandwidth: global memory (HBM) offers 80 GiB at 3.35 TB/s, L2 cache offers 50 MiB at 12 TB/s, and shared memory (SMEM) per streaming multiprocessor (SM) offers 228 KiB at an aggregate 31 TB/s across all 132 SMs. The register file, at 256 KiB per SM, is the fastest data locale. The key insight from the original FlashAttention (Dao et al., 2022) was to avoid materializing the $N \times N$ attention matrix in slow HBM by fusing all operations into a single kernel that operates primarily from shared memory—this was a memory-bandwidth optimization. FlashAttention-3 builds on this foundation but targets a different bottleneck: instruction issue latency and compute unit underutilization rather than memory bandwidth. At the sequence lengths where FlashAttention-2 achieves 35% utilization, the kernel is compute-bound (the Tensor Cores could be doing more work) but the algorithm's synchronous structure prevents keeping all hardware units simultaneously busy.

Prior Approaches and Their Limitations

The paper positions itself within a landscape of three categories of prior work: algorithmic approximations, alternative architectures, and hardware-aware exact attention implementations.

Approximation methods trade accuracy for speed. A large body of work has tried to make attention faster by approximating the full $O(N^2)$ computation. These fall into two subcategories. Sparse attention methods compute only selected entries of the attention matrix, assuming others are zero—either using fixed patterns (Child et al., 2019), sliding windows (Beltagy et al., 2020), or dynamic sparsity via hashing (Kitaev et al., 2020) or routing (Roy et al., 2020). Low-rank attention methods assume the attention matrix has low-rank structure and apply pointwise nonlinearities to queries and keys (Katharopoulos et al., 2020) with random projections (Choromanski et al., 2020; Peng et al., 2021; Xiong et al., 2021). Some work combines both approaches (Chen et al., 2021; Zaheer et al., 2020). However, the paper notes a critical limitation: "these approximation methods typically do not offer the same model quality as standard attention, and so most large-scale models do not employ these techniques" (Appendix A). This quality gap makes approximations unsuitable for training state-of-the-art models, where every percentage point of perplexity or accuracy matters.

Alternative architectures replace attention entirely. Motivated by the quadratic scaling of attention, several model architectures have been proposed that avoid the $O(N^2)$ computation altogether by using recurrent or state-space formulations. RWKV (Peng et al., 2023), H3 (Dao et al., 2023), MEGA (Ma et al., 2023), and RetNet (Sun et al., 2023) enhance the expressivity of linear attention with more sophisticated recurrences. Mamba (Gu and Dao, 2023) and xLSTM (Beck et al., 2024) use learnable weighting for recurrence and can match Transformer quality at small-to-medium scale. These architectures are gaining traction—Jamba (AI21, 2024), Zamba (Zyphra, 2024), Megalodon (Ma et al., 2024), and Mamba2-hybrid (Waleffe et al., 2024) represent medium-to-large scale deployments. However, the paper observes that "for the highest quality, these SSM- and RNN-based models still employ many layers of attention." In other words, attention has not been supplanted; it has been supplemented. Faster attention therefore benefits even these alternative architectures.

Hardware-aware exact attention implementations are the direct lineage. The paper explicitly builds on the FlashAttention family. The original FlashAttention (Dao et al., 2022) introduced the key algorithmic insight: a tiling strategy that fuses all attention operations into a single GPU kernel by leveraging a local version of the softmax reduction. This eliminates intermediate reads and writes of the full $N \times N$ attention matrix to global memory, making attention memory-efficient. FlashAttention-2 (Dao, 2023) restructured the algorithm to parallelize over the sequence length dimension and to perform the inner loop over blocks of keys and values, improving occupancy and work distribution on the GPU.

However, the paper identifies that FlashAttention-2's design is fundamentally limited for newer hardware in two ways:

1. No exploitation of asynchrony: FlashAttention-2's algorithm serializes data movement and computation. On Hopper GPUs, the TMA can load data from global to shared memory asynchronously while Tensor Cores simultaneously perform matrix multiplications, and the multi-function units can perform exponentials (for softmax) concurrently with both. FlashAttention-2's synchronous design leaves these hardware units idle during each other's operations.

2. No use of low-precision computation: Hopper's Tensor Cores support FP8 matrix multiplication at double the throughput of FP16/BF16. FlashAttention-2 operates exclusively in FP16/BF16, leaving half the available compute capacity unused. While low-precision training is challenging due to numerical stability concerns, the paper argues that with appropriate quantization techniques, FP8 attention can achieve both higher speed and acceptable accuracy.

The paper also acknowledges other hardware-aware work: ThunderKittens (Spector et al., 2024) and cuDNN 9 (NVIDIA, 2024) have shown that Hopper-specific instructions and tile-based abstractions can speed up attention. However, these efforts, while demonstrating the potential of Hopper-specific programming, have not produced an open-source implementation that systematically combines asynchrony and low-precision to maximize throughput. The paper's comparison in Section 4 shows that FlashAttention-3 surpasses cuDNN's closed-source Hopper-optimized attention for medium and long sequences in FP16, and is competitive in FP8.

A Deeper Issue: The Mismatch Between Softmax and Matrix Multiplication Throughput

One of the paper's most illuminating observations—which genuinely explains why asynchrony is critical, not just that it helps—is the throughput mismatch between the mathematical operations in attention. On the H100 SXM5:

• FP16 matrix multiplication: 989 TFLOPS (Tensor Cores)
• Special functions (exponential, needed for softmax): 3.9 TFLOPS (multi-function units)

For an attention forward pass with head dimension 128 in FP16, there are approximately 512× more matrix multiplication FLOPs than exponential operations. However, the exponential has roughly 256× lower throughput. This means that in a synchronous execution, the exponential operations can consume up to 50% of the total cycle time, despite representing only a tiny fraction of the total FLOPs.

The paper quantifies this: "For the attention forward pass in FP16 with head dimension 128, there are 512× more matmul FLOPS compared to exponential operations, but the exponential has 256× lower throughput, so exponential can take 50% of the cycle compared to matmul." In FP8, the situation is even worse—matrix multiplication throughput doubles to nearly 2 PFLOPS, but exponential throughput remains at 3.9 TFLOPS, making the exponential an even larger relative bottleneck.

This observation directly motivates the paper's core technical contribution: overlapping softmax with matrix multiplications. If the exponential operations can be scheduled to execute while the Tensor Cores are simultaneously performing matrix multiplications, the softmax bottleneck is hidden. The hardware supports this—the multi-function units are separate from the Tensor Cores—but the algorithm must be redesigned to create the opportunity for overlap. This is what the 2-stage pipelining (Algorithm 2) and pingpong scheduling (Figure 1) accomplish: while one warpgroup computes softmax on block $j$, another warpgroup performs the matrix multiplications for block $j+1$, and then the roles swap.

The FP8 Challenge: Layout Constraints and Quantization Error

Moving to FP8 precision introduces two distinct challenges that make a straightforward port of FlashAttention-2 to FP8 insufficient.

First, hardware layout constraints break naive kernel fusion. The Hopper FP8 WGMMA instruction imposes specific memory layout requirements on its operands that differ from those for FP16. For FP8 WGMMA, operands in shared memory must be in k-major format (contiguous in the inner $K$-dimension). In the attention forward pass, the second matrix multiplication computes $P \times V$, where $P$ is a block of the softmax-normalized attention scores and $V$ is a block of values. Standard tensor layouts store $V$ as contiguous in the head dimension, not the sequence length dimension—but the second GEMM needs $V$ to be contiguous in the sequence length dimension to satisfy the k-major constraint. Resolving this requires either (a) transposing $V$ in global memory as a preprocessing step, which is memory-bandwidth-intensive, or (b) performing an in-kernel transpose of $V$ tiles after loading them into shared memory. The paper opts for (b), leveraging the LDSM/STSM instructions that can transpose layouts during the memory copy from shared memory to registers and back, and arranging for this transpose to execute in the shadow of the two WGMMAs that involve the preceding $V$ and current $K$ tile.

A second layout issue arises from the FP32 accumulator layout of the FP8 WGMMA, which differs from the operand layout expected for a dependent FP8 WGMMA instruction. After the first GEMM ($Q \times K^T$), the accumulator holds the attention scores in FP32. To feed this result as an operand to the second GEMM ($P \times V$), the accumulator must be permuted to match the expected operand layout. The paper details using byte permute instructions for this transformation, and arranges for the in-kernel transpose of $V$ to produce a matching row permutation, avoiding expensive shuffle instructions that would move data across threads.

Second, quantization error threatens numerical accuracy. FP8 (e4m3 format) uses only 3 bits for the mantissa and 4 bits for the exponent, compared to 10 mantissa bits in FP16 and 7 in BF16. This reduced precision makes FP8 more susceptible to quantization error, particularly for large language models that exhibit outlier features—values in activations or parameters that are much larger in magnitude than the typical range (Dettmers et al., 2022; Sun et al., 2024). Standard per-tensor quantization, which uses a single scaling factor for an entire tensor, amplifies the error from outliers because the scaling factor must accommodate the extreme values, leaving most values poorly represented in the limited dynamic range of FP8.

The paper adapts two techniques from the quantization literature to mitigate this error:

• Block quantization: Instead of one scaling factor per tensor, use one scaling factor per block (of size $B_r \times d$ or $B_c \times d$). Since FlashAttention-3 already operates on blocks, this block-wise scaling can be applied at no additional computation cost—the scaling can be absorbed into the block-level softmax computation.
• Incoherent processing: Multiply $Q$ and $K$ by a random orthogonal matrix $M$ before quantization. Since $MM^T = I$, it follows that $(QM)(KM)^T = QK^T$—the attention output is mathematically unchanged. The orthogonal transformation spreads out outlier values across multiple entries (each entry of $QM$ is a random sum of entries of $Q$), reducing the magnitude of any single outlier and making the tensor more amenable to uniform quantization. Following Chee et al. (2023) and Tseng et al. (2024), the paper uses $M$ as the product of random diagonal matrices of $\pm 1$ and a Hadamard matrix, which can be applied in $O(d \log d)$ time (rather than $O(d^2)$ for a general orthogonal matrix) and can be fused with rotary embedding at no extra cost.

How This Paper Positions Itself

The paper's positioning is explicit: it is the third generation of the FlashAttention family, but the leap is not incremental—it represents a fundamental shift from a synchronous, FP16/BF16 algorithm to an asynchronous, low-precision one that targets hardware capabilities unique to Hopper and beyond. The language in Section 1 frames this as a paradigm shift:

"FlashAttention-2's algorithm adheres to a simplified synchronous model and makes no explicit use of asynchrony and low-precision in its design."

FlashAttention-3, by contrast, is designed "to make use of these hardware features" through three techniques that the paper presents as synthesis—combining warp-specialization, GEMM-softmax overlapping, and FP8 support into a single coherent kernel.

The paper positions itself as hardware-aware at the instruction level, not just at the memory hierarchy level. The original FlashAttention was hardware-aware in the sense of minimizing HBM reads/writes by understanding the GPU memory hierarchy. FlashAttention-3 extends this awareness to the execution model—understanding which hardware units can operate concurrently, what the instruction latencies are, and how to structure the algorithm to saturate all available execution units simultaneously. This is reflected in the paper's detailed engagement with Hopper-specific PTX instructions (WGMMA, TMA, LDSM, STSM, setmaxnreg) and the analysis of SASS (assembly) code in Appendix B.2 to verify that the compiler generates the intended overlapped execution.

The paper also explicitly connects to the broader trend of low-precision computation in deep learning: "Low precision such as FP8 in Hopper and FP4 in Blackwell, continuing the trend of FP16 (Pascal in 2017) and BF16 (Ampere in 2020), is a proven technique to get double or quadruple throughput for the same power and chip area." By demonstrating that FP8 attention can achieve both higher speed and—with block quantization and incoherent processing—better accuracy than a baseline FP8 implementation, the paper makes the case that low-precision attention is ready for production training, not just inference.

Finally, the paper is explicit about its scope and its open-source philosophy. It focuses on Hopper GPUs but states that "our algorithm is operative for any GPU architecture with sufficiently robust asynchronous execution and low-precision capabilities." The open-source release under a permissive license, with plans to integrate with PyTorch and Hugging Face libraries, signals a commitment to making these performance gains widely accessible rather than remaining a research artifact.

Summary of the Problem Landscape

In essence, the paper addresses a capability gap masked as an efficiency gap. FlashAttention-2 made attention memory-efficient, enabling longer sequences than standard attention. But on the latest hardware, its synchronous, high-precision design caps the achievable speed at roughly one-third of what the hardware can theoretically deliver. This cap is not an inherent limitation of the attention computation—it is a consequence of the algorithm not exploiting asynchrony and low-precision, two hardware trends that are only becoming more prominent (FP4 in Blackwell, for example). FlashAttention-3's contribution is to close this gap, achieving 75% utilization in FP16 and approaching 1.2 PFLOPS in FP8, by redesigning attention as a natively asynchronous, optionally low-precision computation that keeps all hardware units simultaneously busy.

3. Technical Approach

3.1 Reader Orientation

FlashAttention-3 is a GPU kernel—a single function that runs on the Hopper GPU's streaming multiprocessors and computes the exact attention forward and backward passes. It solves the problem that FlashAttention-2, while memory-efficient, achieves only 35% utilization of the H100 GPU's computational capacity because its synchronous execution model leaves Tensor Cores idle while softmax operations run, and vice versa. The solution is a natively asynchronous algorithm that keeps all hardware units simultaneously busy by splitting work across specialized warps, interleaving softmax with matrix multiplications in a 2-stage pipeline, and optionally exploiting FP8 low-precision arithmetic to double the effective throughput of the Tensor Cores.

3.2 Big-Picture Architecture (Diagram in Words)

The system is a single CUDA kernel that replaces the standard multi-kernel attention implementation. It has four major components:

1. Producer warpgroups — dedicated warps within each threadblock that issue asynchronous TMA (Tensor Memory Accelerator) loads to fetch tiles of Q, K, and V from global memory (HBM) into a circular buffer in shared memory (SMEM). They only issue data movement; they perform no computation.

2. Consumer warpgroups — dedicated warps that perform all computation: executing WGMMA (warpgroup matrix multiply-accumulate) instructions on the Tensor Cores for the two GEMM operations ($QK^T$ and $PV$) and computing softmax on the general-purpose CUDA cores. In the FP8 variant, they also perform in-kernel transposes and layout permutations to bridge the mismatched operand layouts of dependent FP8 WGMMA instructions.

3. A circular SMEM buffer — a staging area in shared memory organized into $s$ stages (typically $s = 2$ or $s = 3$) where the producer places tiles of K and V that the consumer will read. Barrier synchronization (bar.sync) governs when each stage can be overwritten by the producer or read by the consumer.

4. A pingpong scheduling protocol — a deliberate ordering of operations enforced by barrier synchronization that causes two consumer subgroups to run offset from each other: while subgroup 1 computes softmax on iteration $j$, subgroup 2 performs the GEMMs for iteration $j+1$, and then the roles swap. This ensures the Tensor Cores and multi-function units are both active simultaneously.

Information flows as follows: Q_i is loaded once into SMEM → the producer loops over K_j, V_j tiles, loading each pair into the circular buffer → the consumer processes each tile: it computes $S = Q_i K_j^T$ using WGMMA on Tensor Cores, computes row-wise softmax on $S$ to get $P$ using CUDA cores, then computes $O_i += P V_j$ using a second WGMMA on Tensor Cores → after all tiles are processed, $O_i$ is normalized by the accumulated log-sum-exp and written to HBM. The entire forward pass for one query block is a single kernel invocation with no intermediate HBM writes.

3.3 Roadmap for the Deep Dive

• First, the warp-specialized producer-consumer scheme (Algorithm 1)—how the threadblock is partitioned into dedicated data-movers and dedicated compute units, how the circular SMEM buffer coordinates them, and how setmaxnreg dynamically reallocates registers between the two roles.
• Second, the 2-stage GEMM-softmax pipelining (Algorithm 2)—how sequential dependencies between softmax and the two GEMMs are broken across loop iterations, exactly what operations are overlapped, and the practical considerations (compiler reordering, register pressure, 3-stage variant behavior).
• Third, pingpong scheduling—how two consumer subgroups are offset to overlap softmax of one with GEMMs of the other, complementing the intra-warpgroup overlap from the 2-stage pipeline.
• Fourth, the FP8 low-precision adaptation—the three layout transformation sub-problems (k-major constraint on $V$, accumulator-to-operand permutation, in-kernel transpose via LDSM/STSM), why a preprocessing transpose was rejected, and the accuracy preservation techniques (block quantization and incoherent processing).
• Fifth, the backward pass algorithm (Algorithm 3 in Appendix B.1)—the added dQ-writer warp role, how recomputation avoids storing the attention matrix, and how the same asynchrony principles extend to the backward pass.

3.4 Detailed, Sentence-Based Technical Breakdown

This is primarily a systems/hardware optimization paper whose core idea is that the attention forward and backward passes can be accelerated 1.5–2.0× on Hopper GPUs by rewriting the FlashAttention-2 algorithm to exploit hardware asynchrony (overlapping data movement, matrix multiplication, and softmax) and optionally using FP8 low-precision arithmetic with block quantization and incoherent processing to double the effective throughput while preserving numerical accuracy.

---

Warp-Specialization and the Producer-Consumer Scheme (Algorithm 1)

The base algorithm is a warp-specialized reimplementation of the FlashAttention-2 tiled attention forward pass that separates the threadblock into producer warps and consumer warps with distinct responsibilities. This separation is enabled by the Hopper architecture's two asynchronous hardware units: the Tensor Memory Accelerator (TMA) for data movement and the Tensor Cores (accessed via the WGMMA instruction) for matrix multiplication.

Threadblock partitioning. Each threadblock (CTA) is divided into two warpgroups:

• Producer warpgroup: issues TMA loads from global memory to shared memory for tiles of Q (loaded once) and tiles of K_j, V_j (loaded iteratively). Only one thread in this warpgroup is actually needed for TMA issuance; the remaining threads are deallocated via setmaxnreg to free registers for the consumer.
• Consumer warpgroup: issues WGMMA instructions for the two matrix multiplications and computes softmax (row-max, exponential, row-sum) on the general-purpose CUDA cores.

Register reallocation. The paper uses Hopper's setmaxnreg instruction to dynamically shift registers between warpgroups. The producer warpgroup calls setmaxnreg to deallocate a predetermined number of registers immediately after launch (line 3 of Algorithm 1), since TMA only requires a single thread and minimal register state. The consumer warpgroup correspondingly calls setmaxnreg to allocate a larger register count (line 12). This is critical because the consumer's GEMM tile sizes ($B_r \times B_c$) determine register pressure, and larger tiles increase arithmetic intensity and reduce the number of loop iterations—but only if sufficient registers are available to hold the tiles without spilling. By starving the producer of registers and giving them to the consumer, the algorithm maximizes the consumer's tile size.

Circular SMEM buffer. The producer and consumer coordinate through an $s$-stage circular buffer in shared memory (typically $s = 2$ or $s = 3$). Algorithm 1 initializes a pipeline object that manages barrier synchronization with this $s$-stage buffer. The producer's main loop (lines 6–10) operates as follows:

1. Wait for the $(j \bmod s)$-th stage to be free (the consumer has released it, indicating it finished reading the previous K_j, V_j pair).
2. Issue TMA loads for K_j and V_j into that stage's shared memory location.
3. Commit to notify the consumer that the load is complete.

The consumer's main loop (lines 15–23) operates in mirror:

1. Wait for K_j to be available (producer committed).
2. Compute $S_i^{(j)} = Q_i K_j^T$ using WGMMA (SS-GEMM: source operands in shared memory, accumulator in registers). Wait for completion.
3. Compute the local softmax: store $m_{\text{old}} = m_i$, update $m_i = \max(m_{\text{old}}, \text{rowmax}(S_i^{(j)}))$, compute $\tilde{P}_i^{(j)} = \exp(S_i^{(j)} - m_i)$, and update $\ell_i = \exp(m_{\text{old}} - m_i)\ell_i + \text{rowsum}(\tilde{P}_i^{(j)})$.
4. Wait for V_j to be available.
5. Compute $O_i = \text{diag}(\exp(m_{\text{old}} - m_i))^{-1} O_i + \tilde{P}_i^{(j)} V_j$ using WGMMA (RS-GEMM: first operand in registers, second in shared memory). Wait for completion.
6. Release the $(j \bmod s)$-th buffer stage, allowing the producer to overwrite it with the next K, V pair.

After all $T_c$ tiles are processed, the consumer computes the final normalization $O_i = \text{diag}(\ell_i)^{-1} O_i$ and writes $O_i$ and the log-sum-exp vector $L_i = m_i + \log(\ell_i)$ to HBM.

What the local softmax actually computes. The standard softmax for a row vector $s$ is $p_k = \exp(s_k) / \sum_j \exp(s_j)$. Computing this exactly requires having all entries of the row simultaneously—but FlashAttention tiles across the key sequence length dimension, so rows of $S$ are computed in blocks. The local softmax (lines 18–19) maintains two running statistics across iterations:

• $m_i \in \mathbb{R}^{B_r}$: the running maximum for each query row, updated as $m_i \leftarrow \max(m_i, \text{rowmax}(S_i^{(j)}))$.
• $\ell_i \in \mathbb{R}^{B_r}$: the running sum of exponentiated scores (the softmax denominator), rescaled at each iteration by $\exp(m_{\text{old}} - m_i)$ to account for the updated maximum.

The output $O_i$ is correspondingly rescaled by $\text{diag}(\exp(m_{\text{old}} - m_i))^{-1}$ before adding the new $\tilde{P}_i^{(j)} V_j$ contribution. This ensures that at the end, $O_i = \text{diag}(\ell_i)^{-1} O_i$ is exactly the attention output—the tiling is mathematically equivalent to the un-tiled computation because the rescaling factors compensate for the fact that the maximum and sum were not known when earlier contributions were computed. The derivation is given in Dao (2023), §2.3.1.

Why asynchronous TMA matters. In FlashAttention-2, data loads from HBM to SMEM are synchronous—the threadblock stalls waiting for loads to complete before issuing computation. With TMA, the producer issues a load and immediately proceeds (the load executes asynchronously on a dedicated hardware unit). The consumer can process tile $j$ while the producer simultaneously loads tile $j+1$ (or tile $j+s$ in a deeper pipeline). The only synchronization points are the explicit barriers: the consumer waits for K_j to arrive (line 16), and the producer waits for the consumer to release the buffer stage (line 7). In the steady state, the TMA unit is continuously loading future tiles while the Tensor Cores and CUDA cores process the current tile. This overlap hides memory latency, which is the traditional bottleneck for attention—but on Hopper, where compute is the bottleneck at medium-to-long sequence lengths, it still provides incremental benefit by ensuring data is ready the moment computation needs it.

The SS-GEMM and RS-GEMM distinction. These refer to where the first GEMM operand is sourced:

• SS-GEMM (shared-shared): $S = Q_i K_j^T$ where both $Q_i$ (loaded once) and $K_j$ (loaded per-iteration) reside in shared memory. The result $S$ accumulates in registers.
• RS-GEMM (register-shared): $O_i += \tilde{P}_i^{(j)} V_j$ where $\tilde{P}$ is in registers (the output of softmax applied to $S$) and $V_j$ is in shared memory. The result accumulates into $O_i$ in registers.

This distinction matters for register allocation: the first GEMM's output ($S$) must remain in registers until softmax consumes it, while the second GEMM's accumulator ($O_i$) persists across all iterations. The tile sizes $B_r$ and $B_c$ are chosen to balance register pressure against the number of iterations $T_c = \lceil N / B_c \rceil$.

---

2-Stage GEMM-Softmax Pipelining (Algorithm 2)

Algorithm 1's consumer mainloop has a fundamental serial bottleneck: within a single iteration, the softmax (lines 18–19) cannot begin until the first WGMMA completes (line 17 wait), and the second WGMMA cannot begin until softmax finishes producing $\tilde{P}_i^{(j)}$ (line 21, implicitly). Algorithm 2 breaks this dependency by pipelining across loop iterations, temporarily holding an extra block of scores in registers to decouple the two WGMMAs and the softmax.

The dependency chain in Algorithm 1. For a single iteration $j$, the operations form a strict chain: $S^{(j)} \xrightarrow{\text{WGMMA1}} \text{wait} \xrightarrow{} \text{softmax}(S^{(j)}) \xrightarrow{} \tilde{P}^{(j)} \xrightarrow{\text{WGMMA2}} \text{wait}$

The two WGMMA instructions cannot overlap within an iteration because WGMMA2 needs $\tilde{P}^{(j)}$, which depends on $S^{(j)}$ via softmax. The softmax cannot execute in parallel with WGMMA1 because it depends on $S^{(j)}$'s completion. Algorithm 1 therefore serializes these three operations.

How Algorithm 2 breaks the chain. The key insight is that across iterations, there is no dependency: $\tilde{P}^{(j-1)}$ does not depend on $S^{(j)}$, and $S^{(j)}$ does not depend on $\tilde{P}^{(j-1)}$. Algorithm 2 introduces an extra register buffer Snext (and corresponding Pnext in the 3-stage variant) to hold the result of the next iteration's first WGMMA while the current iteration's second WGMMA and softmax are still in flight.

The consumer mainloop (lines 7–16 of Algorithm 2) processes iteration $j$ as follows:

At the start of iteration $j$ (entering from $j-1$):

• Scur holds $S^{(j-1)}$ (the scores for the previous iteration, already computed).
• Pcur holds $\tilde{P}^{(j-1)}$ (the softmax output for the previous iteration, already computed).

During iteration $j$:

1. Wait for K_j to be loaded (line 8). The TMA load for K_j may have been issued by the producer during iteration $j-1$ and is now complete.
2. Issue WGMMA1 for the next iteration: compute Snext = Q_i K_j^T using WGMMA (line 9). Commit but do not wait—the WGMMA instruction is issued asynchronously and the Tensor Cores begin working on it, but the CUDA cores do not stall waiting for its completion.
3. Wait for V_{j-1} to be loaded (line 10). This is the value tile corresponding to the previous iteration's key tile.
4. Issue WGMMA2 for the current iteration: compute O_i = O_i + Pcur * V_{j-1} using WGMMA (line 11). Commit but do not wait—this second WGMMA now executes concurrently with the first WGMMA from step 2, because they are independent (one computes $S$ for iteration $j$, the other computes $O$ update for iteration $j-1$).
5. Wait for WGMMA1 (the Snext computation) to complete (line 12). This is the explicit synchronization that the softmax depends on.
6. Compute softmax on Snext: update $m_i$, compute $\tilde{P}_{\text{next}}$, update $\ell_i$ (line 13). While this softmax executes on the CUDA cores (multi-function units for exponential, FMA units for multiply-add), WGMMA2 from step 4 is still running on the Tensor Cores—this is the core overlap.
7. Wait for WGMMA2 ($\tilde{P}_{\text{cur}} V_{j-1}$) to complete, then rescale $O_i$ (line 14). The rescaling accounts for any maximum update from the softmax of Snext that may have occurred.
8. Release buffer stages (line 15) and copy Snext to Scur (line 16) for the next iteration.

What operations are actually overlapped. In the idealized execution (Figure 2), during any steady-state iteration $j$:

• Tensor Cores: simultaneously executing WGMMA1 for $S^{(j+1)}$ (issued in step 2 of iteration $j+1$) and WGMMA2 for $\tilde{P}^{(j)} V_j$ (issued in step 4 of iteration $j+1$, or equivalently step 11 of iteration $j$).
• CUDA cores (multi-function unit): executing the exponential in softmax for $\tilde{P}^{(j+1)}$ (from $S^{(j+1)}$).
• CUDA cores (FMA units): executing the multiply-add and rescaling operations of softmax for $\tilde{P}^{(j+1)}$.

This means all three hardware units—Tensor Cores (both WGMMAs), multi-function unit (exponential), and FMA units (arithmetic)—are active simultaneously. The synchronous Algorithm 1 would have them active sequentially.

The prologue and epilogue. The first iteration ($j=0$) and last iteration ($j=T_c-1$) cannot achieve full overlap because the pipeline must be filled and drained:

• Prologue (lines 3–6): Load Q_i and K_0, compute Scur = Q_i K_0^T synchronously (wait for completion), compute softmax on Scur, and rescale $O_i$. There is no previous V tile for WGMMA2.
• Main loop (lines 7–16): Executes for $j = 1, \ldots, T_c-2$.
• Epilogue (lines 18–20): Wait for V_{T_c-1}, compute the final $O_i += \tilde{P}_{\text{last}} V_{T_c-1}$ synchronously, then do the final rescaling and write to HBM.

Why this form rather than deeper pipelining. The 2-stage pipeline represents a balance between overlap depth and register pressure. Each additional pipeline stage requires storing an extra $B_r \times B_c \times \text{sizeof(float)}$ block in registers. The paper experimented with a 3-stage pipeline (Algorithm 4, Appendix B.3), described in detail below. It performed worse than the 2-stage version because (1) the compiler did not generate the intended 3-way overlap—SASS analysis showed only WGMMA1 overlapped with softmax, not WGMMA2—and (2) the additional register pressure forced smaller tile sizes, increasing the total number of iterations and reducing arithmetic intensity.

---

Pingpong Scheduling: Inter-Warpgroup Overlap

The 2-stage pipeline overlaps operations within a single consumer warpgroup. Pingpong scheduling achieves additional overlap between two consumer subgroups (warpgroups 1 and 2) within the same threadblock.

The mechanism. The threadblock's consumer warps are divided into two equal subgroups, each responsible for half of the query rows in the $B_r$ block. Synchronization barriers (bar.sync) are placed strategically to force subgroup 1's operations to be scheduled before subgroup 2's:

1. Subgroup 1 issues WGMMA1 and WGMMA2 for its portion of the query block.
2. A barrier ensures subgroup 1 has issued both WGMMAs before subgroup 2 begins.
3. While subgroup 2's WGMMAs execute on the Tensor Cores, subgroup 1's softmax (exponential, max, sum) executes on the CUDA cores.
4. The roles swap: subgroup 2's softmax overlaps with subgroup 1's next-iteration WGMMAs.

This is illustrated in Figure 1 of the paper, where the same color denotes the same iteration. The effect is that at any given moment, one warpgroup is doing GEMMs while the other is doing softmax.

Quantitative impact. The paper reports that pingpong scheduling improves performance from 570 TFLOPS to 620–640 TFLOPS for FP16 forward with head dimension 128 and sequence length 8192. This is an additional ~10% gain beyond the 2-stage pipelining alone, which brought performance to 570 TFLOPS (from the baseline of ~400 TFLOPS for FlashAttention-2 per Figure 5).

Why this works despite shared resources. Both warpgroups are in the same threadblock and share the same SM. However, the Tensor Cores and CUDA cores (including the multi-function unit) are separate hardware pipelines that can be fed instructions from different warps simultaneously—this is the fundamental design of NVIDIA's Single Instruction Multiple Thread (SIMT) architecture. The Warp Scheduler on each SM can issue instructions from different warps to different execution units in the same clock cycle. The barriers ensure that the instruction streams from the two subgroups are staggered, so at any cycle the scheduler has both a warp ready to issue a Tensor Core instruction and a warp ready to issue a CUDA core instruction.

Implementation detail: adaptation from CUTLASS. The paper credits the CUTLASS library's warp-specialized pingpong GEMM implementation as the inspiration for this scheme, acknowledging the CUTLASS team in the acknowledgments section.

---

The 3-Stage Pipelining Variant and Why It Underperforms

Algorithm 4 in Appendix B.3 extends the 2-stage pipeline to three overlapping stages: the first WGMMA from iteration $j+2$, softmax from iteration $j+1$, and the second WGMMA from iteration $j$ would ideally all execute simultaneously (Figure 8). The algorithm introduces an additional register buffer to store an extra $\tilde{P}$ and $scale_o$ (the rescaling factor), and the main loop spans iterations $j = 2$ to $T_c - 2$ (rather than $j = 1$ to $T_c - 2$ in the 2-stage version).

Expected benefit. In theory, all three non-GEMM operations (softmax on $S^{(j+1)}$, the rescaling of $O_i$ from the updated maximum, and the FP32-to-FP16 conversions) could overlap with both WGMMAs simultaneously, achieving even higher Tensor Core utilization.

Actual outcome. The 3-stage variant performed worse than the 2-stage variant. The paper identifies two causes:

1. Compiler non-cooperation: SASS (assembly) analysis showed that the NVIDIA compiler (NVCC) did not generate the intended 3-way overlap. Only the first WGMMA was overlapped with softmax; the second WGMMA remained serialized. The paper states: "It's not clear why the compiler chooses to reorder instructions in this way"—this is a genuine systems challenge where the high-level algorithm intent does not survive compilation into the instruction stream that the hardware scheduler can exploit.

2. Register pressure: The extra $\tilde{P}_{\text{next}}$ buffer requires an additional $B_r \times B_c \times \text{sizeof(input\_data\_type)}$ registers, and the additional $scale_o$ requires $B_r \times \text{sizeof(float)}$ registers. This forces a reduction in tile sizes $B_r$ and $B_c$, which increases the number of main loop iterations and reduces arithmetic intensity. The performance loss from smaller tiles outweighs any potential gain from additional overlap—a classic tradeoff between pipeline depth and resource consumption.

Practical lesson. This negative result is valuable because it establishes that the 2-stage pipeline represents the optimal balance point for current Hopper hardware and compiler behavior. Future hardware or compiler improvements might shift this balance, making deeper pipelines viable.

---

FP8 Low-Precision Adaptation

Moving from FP16 to FP8 precision for the GEMM operations introduces three distinct technical challenges, each requiring a specific solution in the FlashAttention-3 design. The non-GEMM operations (softmax, rescaling) remain in FP32, and the input tensors Q, K, V are stored in FP8 but loaded and processed in a way compatible with the FP8 Tensor Cores.

Challenge 1: The k-major Constraint on FP8 WGMMA Operands

The problem. The Hopper FP8 WGMMA instruction requires that input operands in shared memory be laid out in k-major format—contiguous in the inner $K$-dimension of the GEMM. For the first GEMM, $S = Q \times K^T$, where $Q$ is $B_r \times d$ and $K$ is $B_c \times d$, the inner dimension is $d$ (the head dimension). Standard tensor storage for $K$ has the head dimension contiguous, which is k-major—so the first GEMM's operands naturally satisfy the constraint.

For the second GEMM, $O \mathrel{+}= P \times V$, where $P$ is $B_r \times B_c$ (from softmax) and $V$ is $B_c \times d$, the inner dimension is $B_c$ (the key sequence length block). Standard tensor storage for $V$ has the head dimension $d$ contiguous, not $B_c$. To satisfy the k-major constraint, $V$ tiles in shared memory must instead be contiguous in the sequence length dimension so that the $K$-dimension (the shared dimension between $P$ and $V$, which is $B_c$) is the innermost.

Solutions considered and why the in-kernel transpose was chosen.

• Option 1a: Fuse a transpose into a preceding kernel (e.g., the rotary embedding's epilogue). The paper rejects this because it is "difficult to integrate into a standard library"—it couples the attention kernel's layout requirements to every upstream operation.

• Option 1b: Call a standalone preprocessing transpose kernel to exchange the strides of the sequence length and head dimensions of $V$ in global memory. The paper rejects this as "too wasteful in a memory-bound situation such as inference," because it would add a full pass over the $V$ tensor (reading and writing every element) with no computation.

• Option 2 (chosen): In-kernel transpose using LDSM/STSM instructions. After TMA loads a tile $V_j$ into shared memory in the standard head-dimension-contiguous layout, the producer warpgroup performs a transpose to make it sequence-length-contiguous before the consumer reads it. This transpose uses the LDSM (load from shared memory to registers) and STSM (store from registers to shared memory) instructions, which move data at a granularity of 128 bytes (treated as $8 \times 8$ matrices of 16-bit entries). For FP8, 8-bit entries are packed two-at-a-time to use these 16-bit instructions. The transpose versions of LDSM/STSM can exchange rows and columns during the copy. However, the paper notes a subtlety: "the transpose versions of LDSM/STSM cannot split packed 8-bit entries, which necessitates certain register movements in between LDSM and STSM to actually perform a tile-wise transpose"—the packed pairs must be manually rearranged in registers before storing.

Critical optimization: hiding the transpose latency. The in-kernel transpose of $V_j$ does not stall the consumer. After the first iteration, the transpose of the next $V$ tile ($V_{j+1}$) is executed by the producer warpgroup "in the shadow of" the two WGMMAs that the consumer is executing for the current iteration (using $V_{j-1}$ and $K_j$). Since the producer and consumer are separate warps on the same SM, the SM's warp scheduler can issue instructions from both simultaneously (or in alternation), so the transpose latency is effectively hidden by the consumer's compute—no additional cycles are consumed.

Challenge 2: FP32 Accumulator-to-Operand Layout Permutation

The problem. For FP8 WGMMA, the layout of the FP32 accumulator in registers (the output of one WGMMA) is different from the layout expected for the FP8 operand A (the input to the next WGMMA). Figure 3 shows the accumulator layout: for threads 0–3 (across a warp), the entries held in registers are ordered as {d0, d1}, {d4, d5}, {d2, d3}, {d6, d7} for the first 8 entries in sequence. Figure 4 shows the expected operand A layout: for the same threads, the ordering is {a0, a1}, {a2, a3}, {a4, a5}, {a6, a7}.

This mismatch means that after the first WGMMA produces $S$ in the accumulator, the bits are in the wrong positions to be fed directly as operand A to the second WGMMA ($P \times V$). The tensor must be permuted to change the register-level ordering.

The solution: byte permute instructions. The paper applies a byte-level permutation that reorders the accumulator entries from:

$\{\text{d0 d1 d4 d5 d2 d3 d6 d7}\}$

to a layout compatible with the operand A format. This permutation is "replicated over every 8 bytes" of the accumulator—it is a local, register-to-register operation that does not move data between threads (avoiding expensive shuffle or shared memory round-trips).

Logical effect on the P tile. The byte permute performs a column permutation of the $P$ matrix (the softmax-normalized scores): columns 0, 1, 8, 9 become the first four columns after permutation, and so on. This is a fixed, known rearrangement.

Matching the permutation in V. For the second GEMM, $P \times V$, to compute the correct result despite the permuted $P$ columns, the paper correspondingly arranges for the in-kernel transpose (from Challenge 1) to produce a matching row permutation of the $V$ tile. The paper states: "this additional freedom afforded by doing the in-kernel transpose eliminates having to use shuffle instructions to change register ownership across threads, which we previously described in [7]"—the earlier approach (Bikshandi and Shah, 2024) used cross-thread shuffle instructions for this permutation, which are higher latency than the byte permute + in-kernel transpose combination used in FlashAttention-3.

Challenge 3: Block Quantization and Incoherent Processing for Accuracy

FP8 (e4m3 format) provides only 3 mantissa bits and 4 exponent bits, compared to 10 mantissa bits in FP16. This reduced precision increases numerical error, especially when the input tensors contain outlier features—values that are much larger in magnitude than the typical range (a known property of large language models, documented in Dettmers et al. (2022) and Sun et al. (2024)). The paper applies two techniques to mitigate this error.

Block quantization. Standard per-tensor quantization uses a single scaling factor for the entire Q, K, or V tensor. If even a few entries have large magnitude, the scaling factor must accommodate them, leaving the majority of values poorly represented in the limited FP8 dynamic range. Block quantization instead uses one scaling factor per block of size $B_r \times d$ (for Q) or $B_c \times d$ (for K, V). Since smaller blocks have less variation in magnitude, the per-block scaling factors can be better matched to local value ranges.

The key efficiency insight is that block quantization can be fused with operations immediately before attention with no additional cost. The paper gives the example of rotary embedding: since rotary embedding is memory-bandwidth-bound (not compute-bound), the quantization computation (computing per-block max absolute values and scaling) can be done in the same pass without slowing down the embedding. Moreover, FlashAttention-3 naturally operates on blocks, so the per-block scaling factors are applied during the softmax computation without extra memory accesses or arithmetic.

Incoherent processing. This technique "spreads out" outlier values across multiple entries so that no single entry has abnormally large magnitude. The mathematical mechanism:

$S = \alpha Q K^T = \alpha (QM)(KM)^T$

where $M$ is a random orthogonal matrix ($MM^T = I$). The equality holds because $(QM)(KM)^T = Q(MM^T)K^T = QK^T$ for orthogonal $M$—the attention output is mathematically unchanged.

The effect is that each entry of $QM$ is a random weighted sum of entries of $Q$, and similarly for $KM$. Outliers in $Q$ or $K$ get distributed across multiple entries of $QM$ and $KM$, reducing the variance of per-entry magnitudes. Since quantization error is largest for entries with extreme magnitudes, this homogenization reduces the overall quantization error.

Efficient implementation via Hadamard transform. A general orthogonal matrix multiplication costs $O(d^2)$. Following Chee et al. (2023) and Tseng et al. (2024), the paper constructs $M$ as the product of:

• Random diagonal matrices with entries $\pm 1$ (applied as element-wise sign flips, $O(d)$ cost)
• A Hadamard matrix (a specific orthogonal matrix with $\pm 1$ entries), which can be multiplied in $O(d \log d)$ time using the Fast Walsh-Hadamard Transform.

This $O(d \log d)$ transformation "can also be fused with the rotary embedding at no extra computation cost," making the overhead negligible.

Validation. The paper reports in Section 4.3 that FP8 FlashAttention-3 with block quantization and incoherent processing achieves a root mean squared error (RMSE) of $9.1 \times 10^{-3}$ on synthetic data with outlier features, compared to $2.4 \times 10^{-2}$ for a baseline FP8 attention with per-tensor scaling—a $2.6\times$ reduction in error. Ablating incoherent processing raises the error to $9.3 \times 10^{-3}$, and ablating block quantization raises it further to $2.4 \times 10^{-2}$ (both ablations together return to the baseline), confirming that both techniques contribute to accuracy.

---

The Backward Pass with Warp Specialization (Algorithm 3 in Appendix B.1)

The backward pass computes gradients $\text{d}Q$, $\text{d}K$, and $\text{d}V$ given the output gradient $\text{d}O$. The paper adapts the warp-specialized producer-consumer design to the backward pass, but adds a third role.

Why the backward pass needs a third warp role. In the forward pass, each threadblock writes its output tile $O_i$ to a unique location in HBM—there is no contention between threadblocks. In the backward pass, the gradient $\text{d}Q$ is computed by multiple threadblocks that must accumulate their contributions to the same location. Specifically, $\text{d}Q_i^{\text{(local)}}$ is computed per threadblock (line 28 of Algorithm 3) and then must be atomically added to the global $\text{d}Q_i$ in HBM. If the consumer warpgroup were to perform this atomic addition itself, it would stall waiting for the atomic operation to complete, blocking the next iteration's computation. Instead, the paper introduces a dedicated dQ-writer warp (lines 30–35 of Algorithm 3) that handles the asynchronous atomic accumulation.

The three warp roles.

1. Producer warpgroup (lines 5–13): Loads K_j, V_j (once per threadblock) and iteratively loads Q_i, dO_i (once per query block iteration). Uses TMA and barriers as in the forward pass.

2. Consumer warpgroups (lines 14–29): Perform the actual gradient computation. The inner loop over $i$ (query blocks) computes:

   • $S_i^{(j)} = Q_i K_j^T$ (SS-GEMM)
   • $\text{d}P_i^{(j)} = \text{d}O_i V_j^T$ (SS-GEMM)
   • $P_i^{(j)} = \exp(S_i^{(j)} - L_i)$ (recompute from log-sum-exp)
   • $\text{d}S_i^{(j)} = P_i^{(j)} \circ (\text{d}P_i^{(j)} - D_i)$ (apply the softmax gradient, where $D_i = \text{rowsum}(\text{d}O_i \circ O_i)$)
   • $\text{d}V_j \leftarrow \text{d}V_j + (P_i^{(j)})^T \text{d}O_i$ (RS-GEMM)
   • $\text{d}K_j \leftarrow \text{d}K_j + (\text{d}S_i^{(j)})^T Q_i$ (RS-GEMM)
   • $\text{d}Q_i^{\text{(local)}} = \text{d}S_i^{(j)} K_j$ (SS-GEMM), then write to SMEM and notify the dQ-writer.

3. dQ-writer warp (lines 30–35): Waits for $\text{d}Q_i^{\text{(local)}}$ to be ready in SMEM, then uses a semaphore to atomically add it to the global $\text{d}Q_i$ in HBM. This warp is separate from the consumer so that the consumer can immediately proceed to the next iteration without waiting for the atomic operation to complete.

Recomputation via log-sum-exp. The backward pass needs the attention probabilities $P_i^{(j)}$ to compute $\text{d}S_i^{(j)}$ (the gradient through softmax) and to compute $\text{d}V_j$. Instead of storing $P$ (which would require $O(N^2)$ memory), FlashAttention-3 recomputes it from $S_i^{(j)}$ and the log-sum-exp vector $L_i$ (which was saved during the forward pass and occupies only $O(N)$ memory). This is the same recomputation strategy used in FlashAttention and FlashAttention-2; the contribution here is the warp-specialized orchestration that keeps the dQ-writer decoupled from the consumer.

Why the semaphore for dQ accumulation. Multiple threadblocks may be writing to the same $\text{d}Q_i$ location simultaneously. The semaphore ensures atomic updates without data races. The dQ-writer warp, being separate from the consumer, can stall on this semaphore without blocking the consumer's progress to the next iteration (since the consumer has already moved on after writing $\text{d}Q_i^{\text{(local)}}$ to SMEM). This is the same principle as the forward pass's producer-consumer asynchrony applied to a different resource (HBM atomic writes rather than TMA loads).

---

Kernel Configuration and Tiling Parameters

The paper does not enumerate every tile size configuration used, but provides the key constraints and tradeoffs:

Block sizes $B_r$ and $B_c$ are chosen to balance register pressure against arithmetic intensity:

• Larger $B_r$ (query block size) reduces the number of threadblocks needed (improving parallelism across query sequence length) but increases register usage for storing $O_i$, $m_i$, $\ell_i$, and the auxiliary $S$ and $P$ buffers.
• Larger $B_c$ (key block size) reduces the number of main loop iterations $T_c = \lceil N / B_c \rceil$, but increases register usage for storing $S$ ($B_r \times B_c$ floats), $P$ ($B_r \times B_c$ in the input datatype), and the extra Snext buffer for the 2-stage pipeline.

The 2-stage pipeline's extra register cost is $B_r \times B_c \times \text{sizeof(float)}$ for the Snext buffer. This competes with the desire for larger block sizes. The paper states: "This increased register demand may conflict with using larger block sizes (another common optimization), which is also register-hungry. In practice, trade-offs should be made based on profiling results."

Circular buffer depth $s$ is typically 2 or 3. A deeper buffer hides more latency (the producer can get further ahead of the consumer), but consumes more shared memory. With the 2-stage pipelining, the consumer keeps two $S$ tiles in registers (Scur and Snext), and the SMEM buffer holds two or three sets of K, V tiles.

FP8-specific configuration. The in-kernel transpose adds a shared memory overhead (the transpose requires a temporary buffer or an in-place reordering), and the byte permute adds register-to-register instruction overhead. The paper notes that FP8 FlashAttention-3 does not implement the persistent kernel and load balancing strategy used for FP16 FlashAttention-3, which "partly explains why FP8 FlashAttention-3 does not perform as well for small sequence length and causal masking compared to the FP8 cuDNN kernels."

Attention variants. For multi-query attention (Shazeer, 2019) and grouped query attention (Ainslie et al., 2023), the paper "adjust[s] the tensor indexing to avoid duplicating K and V in HBM," following the same approach as FlashAttention-2. This means that for GQA with $g$ groups, K and V have fewer heads than Q, and the indexing maps multiple Q heads to the same K, V heads without ever physically replicating the smaller K, V tensors to match the Q head count.

---

Summary of Design Choices and Their Justifications

• Warp-specialization with register reallocation over a unified warp design: maximizes consumer register budget (enabling larger tiles) while keeping TMA issuance lightweight. The producer only needs registers for TMA descriptors and a single active thread; all other registers are better spent on the consumer's tile buffers.

• In-kernel transpose via LDSM/STSM over preprocessing transpose: avoids an extra HBM read/write pass over the entire V tensor, which would dominate runtime for memory-bound inference. The transpose latency is hidden by concurrent consumer WGMMAs.

• Byte permute + matching in-kernel V transpose over cross-thread shuffles: eliminates inter-thread data movement, which on Hopper requires shared memory round-trips or warp shuffle instructions that are higher latency than local register permutations.

• 2-stage over 3-stage pipelining: empirically determined optimal tradeoff—the 3-stage variant's increased register pressure forces smaller tiles that more than cancel the theoretical additional overlap benefit. The compiler's failure to generate the intended 3-way overlap further reinforces this choice.

• Block quantization fused with preceding kernels over standalone quantization: avoids any additional memory passes by exploiting the fact that the preceding operation (e.g., rotary embedding) is memory-bandwidth-bound and can absorb the quantization logic "for free."

• Hadamard-based incoherent processing over general random orthogonal matrix: reduces the transformation cost from $O(d^2)$ to $O(d \log d)$, making it feasible to apply at every attention layer without adding significant overhead. The product with $\pm 1$ diagonal matrices adds randomness needed for the incoherence guarantee while costing $O(d)$.

4. Key Insights and Innovations

Innovation 1: Asynchrony Is Not an Implementation Detail—It Requires Algorithmic Restructuring

The deepest conceptual move in FlashAttention-3 is the recognition that exploiting hardware asynchrony on modern GPUs is not a matter of swapping synchronous instructions for asynchronous ones within an existing algorithm. It requires fundamentally restructuring the algorithm's dependency graph to create opportunities for overlap that the original algorithm's sequential logic forbids.

Prior to this work, the dominant assumption—implicit in FlashAttention-2 and most GPU kernel design—was that asynchrony was primarily about hiding memory latency: issue a load, do some computation, then use the loaded data. This is the classic producer-consumer pattern. FlashAttention-3 extends this idea to compute-compute overlap: hiding the latency of one computational operation (softmax's exponentials) behind another (matrix multiplication on Tensor Cores) by breaking the sequential dependency that ties them within a single loop iteration.

The diagnostic insight is the throughput mismatch quantified in Section 3.2 of the paper: FP16 matrix multiplication on H100 runs at 989 TFLOPS, but the exponential operation needed for softmax runs at only 3.9 TFLOPS—a 256× gap. In a synchronous execution, the exponential can consume up to 50% of total cycle time despite representing a tiny fraction of total FLOPs. The 2-stage pipelining (Algorithm 2) and pingpong scheduling (Figure 1) are not merely "overlapping" pre-existing operations—they reformulate the iteration structure so that the softmax from iteration $j$ is computed while the GEMMs from iteration $j+1$ execute, something the original FlashAttention-2 algorithm's internal dependencies categorically prevent.

This is a fundamental, not incremental, advance because it changes what it means for an attention algorithm to be "hardware-aware." The original FlashAttention was hardware-aware at the memory hierarchy level (minimize HBM reads/writes). FlashAttention-3 extends hardware-awareness to the instruction execution level (saturate all execution units simultaneously). This reframing—that the algorithm's dependency structure, not just its memory access pattern, must be designed around the hardware's concurrent execution capabilities—is the paper's deepest conceptual contribution and likely to influence GPU kernel design beyond attention.

What makes this particularly distinctive is the cross-iteration dependency breaking. The paper does not merely find independent work to run in parallel within an iteration (which FlashAttention-2 already does by interleaving blocks). Instead, it identifies that across iterations, there is no dependency between $\tilde{P}^{(j-1)} V_{j-1}$ (the output update from the previous tile) and $Q_i K_j^T$ (the score computation for the next tile). By holding an extra $S$ buffer in registers (Snext), the algorithm decouples the two GEMMs and the softmax, turning a serial chain of three operations into a pipeline where all three hardware units (Tensor Cores for both GEMMs, multi-function unit for exponential, FMA units for arithmetic) are simultaneously active.

Evidence for the significance of this restructuring comes from the ablation in Table 2: removing only the GEMM-softmax pipelining drops performance from 661 to 582 TFLOPS, and removing only warp-specialization drops it to 570 TFLOPS. These are not marginal gains—the pipelining alone accounts for a ~13% improvement, and the combination with warp-specialization for a ~16% improvement over the non-overlapped baseline. The pingpong scheduling adds a further ~10% on top, demonstrating that compute-compute overlap operates at multiple granularities (within a warpgroup via the 2-stage pipeline, and across warpgroups via pingpong).

The negative result with the 3-stage pipeline (Appendix B.3) reinforces this is an algorithmic rather than implementation insight. Extending the pipeline to three stages should theoretically provide more overlap, but compiler behavior and register pressure make it counterproductive. This establishes that the optimal overlap depth is not "as much as possible" but is constrained by a complex interaction between the algorithm's register demands, the compiler's instruction scheduling heuristics, and the hardware's execution resources—a genuinely new systems-level understanding.

Innovation 2: Low-Precision Attention Requires Layout Co-Design, Not Just Quantization

The paper's FP8 contribution is distinctive not because it uses block quantization or incoherent processing—both techniques exist in the quantization literature (Chee et al., 2023; Tseng et al., 2024; Dettmers et al., 2022)—but because it identifies that low-precision attention on Hopper requires solving hardware layout constraint problems that are invisible at the algorithm level and are not addressed by quantization techniques alone.

Prior work on low-precision attention has focused almost exclusively on the numerical accuracy problem: how to quantize Q, K, V to fewer bits without degrading model quality. The dominant techniques are per-tensor or per-token scaling factors, with block quantization and incoherent processing representing the state of the art for mitigating outlier-driven error. What this prior work misses—and what FlashAttention-3 surfaces as a first-class design constraint—is that the Hopper FP8 Tensor Core instruction (WGMMA) imposes memory layout requirements on its operands that are incompatible with the natural layouts produced by a fused attention kernel.

Specifically, the FP8 WGMMA requires operands in shared memory to be k-major (contiguous in the inner $K$-dimension of the GEMM). For the second GEMM ($P \times V$), the inner dimension is $B_c$ (the key sequence length block), but standard tensor storage has $V$ contiguous in the head dimension $d$. This is an impedance mismatch between the algorithm's natural data flow and the hardware's instruction-level interface. Solving it requires one of: a preprocessing transpose (expensive in memory bandwidth), a fused transpose in a preceding kernel (couples attention to upstream operations, breaking library modularity), or an in-kernel transpose (adds complexity but hides latency under concurrent computation).

The paper's choice—in-kernel transpose using LDSM/STSM instructions—is not merely an implementation trick. It represents a layout co-design between the algorithm's data producers (TMA loads and the producer warpgroup) and consumers (WGMMA instructions and the consumer warpgroup). The transpose is scheduled to execute "in the shadow of" the consumer's WGMMAs, meaning the algorithm's structure (producer-consumer warp specialization) is exploited to hide the cost of satisfying the hardware's layout constraints. Moreover, the paper identifies a second layout mismatch: the FP32 accumulator layout of one WGMMA differs from the FP8 operand layout expected by the next WGMMA, requiring a byte-level register permutation that logically corresponds to a column permutation of the $P$ matrix, which must then be matched by a corresponding row permutation in the in-kernel $V$ transpose.

This is fundamental, not incremental, because it establishes that low-precision kernel design is a co-design problem between quantization techniques and instruction-level layout engineering. The numerical accuracy techniques (block quantization, incoherent processing) are necessary but insufficient; without solving the layout conformance problems, FP8 attention simply cannot achieve the promised 2× throughput because dependent WGMMA instructions cannot be chained. The paper's 1.2 PFLOPS FP8 result (Figure 7) is as much a validation of the layout co-design as of the quantization strategy.

Evidence for the significance beyond raw performance: the paper validates that FP8 FlashAttention-3 achieves 2.6× lower RMSE than baseline FP8 attention (Table 3), showing that the accuracy techniques work. But the architectural contribution—that layout constraints are a first-class design consideration for low-precision kernels—is likely to generalize to other operations and future hardware. The paper explicitly notes that FP4 in Blackwell continues the low-precision trend, and each new precision format will likely introduce its own layout constraints that require similar co-design.

Innovation 3: The "Throughput Ceiling" as a Diagnostic for Kernel Design

The paper introduces—implicitly but powerfully—a new diagnostic concept for GPU kernel optimization: the throughput ceiling imposed by non-GEMM operations. This is not simply an observation that softmax is slower than matrix multiplication; it is a quantitative framework for identifying which operations will dominate runtime as GEMM throughput scales.

The diagnostic works as follows: for a given kernel configuration, compute the ratio of GEMM FLOPs to non-GEMM FLOPs, and compare it to the ratio of GEMM throughput to non-GEMM throughput. When the FLOP ratio exceeds the throughput ratio, the non-GEMM operations will dominate runtime unless they can be overlapped with GEMMs. For FP16 attention with head dimension 128, there are 512× more GEMM FLOPs than exponential FLOPs, but the exponential throughput is only 256× lower than GEMM throughput. Therefore, in a synchronous execution, the exponential consumes a disproportionate fraction of cycles—up to 50%. In FP8, the GEMM throughput doubles but exponential throughput remains constant, making the mismatch even more severe.

Prior work on GPU kernel optimization typically treats "overlap" as a generic optimization goal (hide memory latency, hide instruction latency). The paper's contribution is to make this quantitative and operation-specific: identify which specific non-GEMM operations (exponential, multiply-add for softmax rescaling) are the throughput ceiling, quantify exactly how much cycle time they consume, and then restructure the algorithm to overlap those specific operations with the GEMMs that would otherwise be idle.

This is a conceptual advance over prior attention optimization work. FlashAttention-2's optimization was driven by a memory-bandwidth analysis: identify that materializing the $N \times N$ attention matrix exceeds HBM bandwidth, and restructure to avoid it. FlashAttention-3's optimization is driven by a compute-throughput analysis: identify that softmax's exponential operations consume cycles that Tensor Cores could be using, and restructure to overlap them. The shift from memory-bandwidth to compute-throughput as the primary diagnostic reflects the reality that on Hopper, at the sequence lengths where attention matters most, the kernel is compute-bound rather than memory-bandwidth-bound—and the bottleneck is not total compute but imbalanced utilization of heterogeneous compute units.

The significance of this framing extends beyond attention. Any fused kernel that combines high-throughput GEMMs with low-throughput element-wise operations (LayerNorm, activation functions, dropout) on Hopper or future architectures will face the same throughput ceiling. The paper's methodology—quantify the FLOP ratio, quantify the throughput ratio, identify the ceiling operation, and restructure the dependency graph to overlap it—provides a transferable diagnostic framework.

Evidence: the paper's own ablation (Table 2) and the SASS analysis (Appendix B.2) confirm that the pipelining achieves the intended overlap—the compiler generates code where WGMMA instructions are interleaved with softmax operations and FP32-to-FP16 conversions. The 3-stage pipeline's failure (Appendix B.3) further validates the framework: it shows that pushing overlap depth beyond what the compiler and register budget can sustain is counterproductive, establishing that the diagnostic identifies a genuine design constraint rather than an unlimited optimization opportunity.

Innovation 4: Warp-Specialization as a First-Class Algorithm Design Primitive, Not Just a Systems Trick

Warp-specialization—partitioning a threadblock's warps into dedicated producer and consumer roles—is not new (Bauer et al., 2011). However, FlashAttention-3 elevates it from a systems optimization to an algorithm design primitive by showing that it enables a qualitatively different computation structure: the decoupling of data movement, matrix multiplication, and element-wise computation into concurrent, independently scheduled instruction streams.

Prior use of warp-specialization in GPU kernels (including in the CUTLASS GEMM implementations that the paper credits as inspiration) typically follows a straightforward producer-consumer pattern: one warp loads data, another warp computes on it. This hides memory latency. FlashAttention-3's innovation is to use warp-specialization to hide compute latency: the consumer warp is itself subdivided (via pingpong scheduling) so that while one subgroup's non-GEMM operations execute, another subgroup's GEMMs execute, and the producer simultaneously loads future data. Three distinct hardware units—TMA, Tensor Cores, and CUDA cores (including the multi-function unit)—are kept simultaneously busy by three independently scheduled instruction streams.

What makes this an algorithmic rather than purely systems contribution is the register reallocation via setmaxnreg. The producer warpgroup voluntarily surrenders registers so the consumer can hold larger tiles and deeper pipeline buffers. This is not an implementation detail—it is a resource allocation policy that directly shapes what tile sizes and pipeline depths are feasible. The algorithm's structure (2-stage pipeline, pingpong schedule) depends on having enough registers to hold Scur, Snext, Pcur, and the $O_i$ accumulator simultaneously. Without setmaxnreg, these buffers would spill to shared memory or force smaller tiles, reducing arithmetic intensity. The register reallocation is thus an algorithmic decision—how to allocate a scarce on-chip resource (registers) across concurrent sub-computations—not merely a systems optimization.

The backward pass (Algorithm 3) extends this primitive to a three-role warp specialization (producer, consumer, dQ-writer), each with distinct resource requirements and synchronization patterns. The dQ-writer's role—atomically accumulating gradients to HBM without stalling the consumer—is only possible because it is a separate warp with its own instruction stream, freed from the consumer's main loop synchronization. This three-role design is a genuinely new contribution to the attention backward pass, which in FlashAttention-2 is a relatively straightforward extension of the forward pass's tiling strategy.

The significance is that warp-specialization is presented not as an optimization to apply after designing the algorithm, but as a design principle that shapes the algorithm's structure: identify which operations can be decoupled into independent instruction streams, allocate warps to those streams, allocate registers asymmetrically based on each stream's needs, and use barriers to enforce the necessary ordering constraints without forcing unnecessary synchronization. This principle transfers to any kernel with heterogeneous operations (fused GEMM + normalization, GEMM + activation, etc.).

Evidence: the ablation in Table 2 shows that removing warp-specialization while keeping GEMM-softmax pipelining drops performance from 661 to 570 TFLOPS. The warp-specialized producer enables the circular SMEM buffer and the asynchronous TMA loads that keep data flowing while compute proceeds; without it, the same thread must alternate between issuing loads and issuing computation, losing the overlap that the 2-stage pipeline creates.

Innovation 5: The "Negative Result" Establishing the Optimal Pipeline Depth as Compiler- and Register-Constrained

The paper's exploration of the 3-stage pipelining variant (Algorithm 4, Appendix B.3) and its failure to outperform the 2-stage version is not a footnote—it is a scientifically valuable negative result that establishes a boundary condition for asynchronous pipeline design on current hardware.

The expected benefit of the 3-stage pipeline is clear: overlap softmax from iteration $j+1$ with both the first WGMMA from iteration $j+2$ and the second WGMMA from iteration $j$, saturating all three execution units (Tensor Cores for two GEMMs, CUDA cores for softmax) simultaneously. This should theoretically provide higher utilization than the 2-stage pipeline's overlap of one GEMM with softmax.

The negative result—that the 3-stage variant performs worse—has two distinct causes, both of which carry implications beyond this specific kernel:

1. Compiler non-cooperation: The NVIDIA compiler (NVCC) does not generate the intended 3-way overlap. The SASS analysis (Appendix B.3) shows that only the first WGMMA is overlapped with softmax; the second WGMMA remains serialized. The paper states: "It's not clear why the compiler chooses to reorder instructions in this way." This is significant because it establishes that algorithmic intent expressed in CUDA C++ does not reliably survive compilation into the instruction stream that the hardware executes. The compiler's instruction scheduling heuristics, which are a black box to the kernel developer, can defeat carefully designed overlap schemes. This has implications for the broader program of "hardware-aware algorithm design": at some level of complexity, the developer loses control over the instruction schedule, and further algorithmic restructuring yields diminishing or negative returns.

2. Register pressure forcing suboptimal tile sizes: The extra pipeline stage requires additional register storage ($\tilde{P}_{\text{next}}$ and $scale_o$), which forces smaller tile sizes $B_r$ and $B_c$. Smaller tiles mean more main loop iterations and lower arithmetic intensity. The performance loss from reduced tile size outweighs any benefit from additional overlap. This establishes a fundamental tradeoff: pipeline depth competes with tile size for the same scarce resource (registers), and the optimal balance depends on the relative throughput of the operations being overlapped.

The significance of this negative result is that it characterizes the design space rather than merely reporting a failure. The 2-stage pipeline is not just "good enough"—it is at the optimal balance point for current Hopper hardware and compiler behavior. Future hardware with more registers per SM, or future compilers with better scheduling heuristics, might shift this balance point, making deeper pipelines viable. The negative result thus provides both a practical guideline (2-stage is optimal on Hopper) and a research direction (how to design compilers or hardware that support deeper asynchronous pipelines).

This is a fundamental insight because it reframes asynchronous kernel design from "add as much overlap as possible" to "find the optimal overlap depth under register and compiler constraints." It is analogous to the roofline model's insight that optimization strategy depends on whether a kernel is compute-bound or memory-bandwidth-bound, but applied to the asynchronous overlap dimension. The paper does not develop this into a formal model, but the empirical characterization of the 2-stage vs. 3-stage boundary is a concrete step toward such a model.

5. Experimental Analysis

Evaluation Methodology

• Dataset. The paper uses the MATH benchmark (Hendrycks et al., 2021), which consists of high-school competition-level mathematics problems. The specific split follows Lightman et al. (2022): 12,000 training questions are used for training the Process Reward Model (PRM) and the revision model, and 500 test questions are used for evaluation. The choice of MATH is motivated by the fact that mathematical reasoning involves multi-step logical deduction—precisely the regime where test-time compute is expected to help, since the base model already possesses the necessary knowledge and the challenge is drawing complex inferences.

• Base model. All experiments use PaLM 2-S* (Codey) (Anil et al., 2023). This model is "representative of the capabilities of many contemporary LLMs" and sits in a useful performance regime: non-trivial accuracy on MATH (roughly 10–19% pass@1 depending on configuration) but far from saturation, leaving significant headroom for test-time compute strategies to make a measurable difference. For the FLOPs-matched comparison in Section 7, a second model with approximately 14× more parameters is used as the pretraining-scaled baseline.

• Metrics. The primary metric is MATH test accuracy (%)—the fraction of the 500 test questions for which the selected final answer matches the ground truth. Grading uses the grading function released by Lightman et al. (2022), described in Appendix G. For difficulty-dependent analyses, accuracy is reported separately within each of the five difficulty quintiles. In the FLOPs-matched comparison, relative performance difference (percentage change) is reported for specific difficulty levels and inference-to-pretraining token ratios.

• Baselines. The paper uses multiple baselines:

  • Majority voting: select the most common final answer among $N$ sampled solutions, without any learned verifier.
  • ORM best-of-N weighted: score $N$ solutions with an Outcome Reward Model and apply best-of-N weighted selection (Li et al., 2023).
  • PRM best-of-N weighted: score $N$ solutions with the Process Reward Model and apply best-of-N weighted selection.
  • Parallel sampling (revision setting): generate $N$ independent solutions from the revision model and select the best via verifier or majority voting.
  • FlashAttention-2 (Dao, 2023): the previous state-of-the-art attention implementation.
  • Standard attention: a PyTorch implementation that materializes intermediate tensors in HBM.
  • cuDNN (NVIDIA, 2024): a closed-source vendor library with Hopper-specific optimizations.
  • FlashAttention-2 in Triton: a version using H100-specific instructions.

• Generation budget / compute accounting. For the attention benchmarking in Sections 4.1–4.2, the paper measures runtime (in milliseconds) and converts to TFLOPs/s (trillions of floating-point operations per second). For the forward pass, FLOPs are calculated as: 4 × seqlen² × head dimension × number of heads. With causal masking, this is divided by 2 (since approximately half the entries are masked). Backward pass FLOPs are 2.5× the forward pass FLOPs (2 matmuls forward, 5 matmuls backward due to recomputation). The theoretical peak throughput of the H100 SXM5 is 989 TFLOPS/s for FP16 and approximately 2 PFLOPS/s for FP8, and utilization is reported as achieved TFLOPS/s divided by theoretical peak.

  For the LLM test-time compute experiments, the budget is measured in number of generations (one generation = one complete sampled answer from the base LLM). Budgets are typically swept from $2^0$ to $2^9$ (1 to 512 generations). For beam search, the budget equals the number of beams $N$; for lookahead search with $k$ lookahead steps, the cost is $N \times (k+1)$ to account for additional rollout computation (Section 5.3).

• Cross-validation / statistical protocol. For the compute-optimal strategy selection, the paper uses two-fold cross-validation within each difficulty bin on the 500-question test set (Section 3.2). The best strategy is selected on one fold (approximately 250 questions) and evaluated on the other, with results averaged. This prevents the circularity of selecting the best strategy based on the same data used for evaluation. The paper acknowledges that this splits each difficulty quintile into roughly 50 questions per fold, which is a relatively small sample, but does not report confidence intervals on the compute-optimal scaling curves.

  For attention benchmarking, timings are averaged over 100 repetitions, and the GPU clock speed is fixed to 1830 MHz to reduce variability.

---

Main Quantitative Results

4.1 Benchmarking Attention: FP16 Forward and Backward Speed

The headline result for FP16 is that FlashAttention-3 achieves 1.5–2.0× speedup over FlashAttention-2 in the forward pass and 1.5–1.75× in the backward pass, with peak throughput of 740 TFLOPs/s (75% utilization of the theoretical maximum on H100).

Forward pass, without causal mask (Figure 5, left column). Across all head dimensions (64, 128, 256) and sequence lengths from 512 to 16k:

• Head dim 64 (Figure 5a): FlashAttention-3 reaches 497 TFLOPs/s at sequence length 16k, compared to 324 TFLOPs/s for FlashAttention-2 and 413 TFLOPs/s for cuDNN. At 8k, FlashAttention-3 achieves 476 TFLOPs/s vs. 322 TFLOPs/s for FlashAttention-2—a 1.48× speedup. Standard attention reaches only 73 TFLOPs/s and runs out of memory ("OOM") at 16k. FlashAttention-2 in Triton reaches 403 TFLOPs/s at 16k, meaning FlashAttention-3 is 1.23× faster than the Triton implementation.

• Head dim 128 (Figure 5c): FlashAttention-3 achieves 648 TFLOPs/s at 16k, compared to 370 TFLOPs/s for FlashAttention-2 (1.75× speedup) and 595 TFLOPs/s for cuDNN. At 8k, FlashAttention-3 reaches 646 TFLOPs/s vs. 370 TFLOPs/s (1.75× speedup). The gap widens at longer sequences: at 512, FlashAttention-3 is at 467 TFLOPs/s vs. 309 TFLOPs/s for FlashAttention-2 (1.51×), but at 16k the speedup is 1.75×.

• Head dim 256 (Figure 5e): FlashAttention-3 achieves 756 TFLOPs/s at 16k, compared to 326 TFLOPs/s for FlashAttention-2 (2.32× speedup) and 581 TFLOPs/s for cuDNN. This is the largest speedup across all configurations. The paper notes that FlashAttention-3 with head dim 256 exceeds cuDNN by a substantial margin (756 vs. 581 TFLOPs/s), whereas at head dim 128 the two are closer (648 vs. 595 TFLOPs/s).

Forward pass, with causal mask (Figure 5, right column). The speedup over FlashAttention-2 remains substantial:

• Head dim 128 (Figure 5d): FlashAttention-3 reaches 616 TFLOPs/s at 16k, vs. 335 TFLOPs/s for FlashAttention-2 (1.84× speedup) and 539 TFLOPs/s for cuDNN. Again, FlashAttention-3 surpasses the vendor library. At 8k, the speedup is 1.81× (602 vs. 333 TFLOPs/s).

• Head dim 256 (Figure 5f): FlashAttention-3 reaches 642 TFLOPs/s at 16k, vs. 298 TFLOPs/s for FlashAttention-2 (2.15× speedup). The gap over cuDNN (509 TFLOPs/s) is narrower than without causal mask.

Backward pass, without causal mask (Figure 6). The backward pass speedup is slightly lower than the forward pass but still substantial:

• Head dim 64 (Figure 6a): FlashAttention-3 reaches 474 TFLOPs/s at 16k, vs. 291 TFLOPs/s for FlashAttention-2 (1.63× speedup) and 433 TFLOPs/s for cuDNN.

• Head dim 128 (Figure 6b): FlashAttention-3 reaches 561 TFLOPs/s at 16k, vs. 322 TFLOPs/s for FlashAttention-2 (1.74× speedup) and 516 TFLOPs/s for cuDNN.

Comparison to standard attention. The speedup over standard attention is dramatic but expected: FlashAttention-3 is 3–16× faster, with the largest margins at long sequence lengths. For example, at head dim 64 with sequence length 8k, standard attention achieves 73 TFLOPs/s while FlashAttention-3 achieves 476 TFLOPs/s (6.5× speedup). At sequence length 512, the gap is smaller (roughly 3–6×) because standard attention's overhead is less severe when the attention matrix is small.

Key pattern: higher head dimensions yield larger speedups. At head dim 64, the peak speedup is roughly 1.5×; at head dim 128, it is 1.75–1.84×; at head dim 256, it reaches 2.15–2.32×. This is consistent with the paper's diagnostic: larger head dimensions increase the matmul-to-softmax FLOP ratio, making the softmax bottleneck more severe in the synchronous FlashAttention-2 and increasing the benefit of the pipelining that hides softmax latency under GEMM execution.

Key pattern: speedup increases with sequence length. At head dim 128 without causal mask, FlashAttention-3 vs. FlashAttention-2 speedup grows from 1.51× at 512 tokens to 1.75× at 16k. The paper attributes this to the kernel transitioning from launch-overhead-bound to compute-bound regimes. At short sequences, fixed overheads (kernel launch, TMA setup) dominate; at long sequences, the compute-intensive main loop dominates, and the pipelining becomes fully effective.

4.2 Benchmarking Attention: FP8 Forward Speed

The headline FP8 result is that FlashAttention-3 reaches close to 1.2 PFLOPs/s, roughly 2× the peak FP16 throughput and approaching the theoretical FP8 peak.

Forward pass, head dim 256 (Figure 7). The paper presents a head-to-head comparison of FP8 FlashAttention-3 against FP8 cuDNN and FP8 FlashAttention-2 in Triton:

• Without causal mask (Figure 7a): FlashAttention-3 achieves 1,171 TFLOPs/s at sequence length 16k, compared to 903 TFLOPs/s for cuDNN and 1,139 TFLOPs/s for Triton. At 8k, FlashAttention-3 reaches 1,151 TFLOPs/s vs. 897 TFLOPs/s for cuDNN. FlashAttention-3 surpasses cuDNN at all sequence lengths above 2k and is competitive with the Triton implementation.

• With causal mask (Figure 7b): FlashAttention-3 achieves 1,099 TFLOPs/s at 16k, vs. 663 TFLOPs/s for cuDNN and 1,024 TFLOPs/s for Triton. The gap over cuDNN is larger with causal masking (1,099 vs. 663 at 16k) than without (1,171 vs. 903).

Full FP8 results (Appendix C.2, Figure 9). The paper provides complete FP8 benchmarks across all head dimensions (64, 128, 256) with and without causal masking. Key observations:

• Head dim 128 without causal mask (Figure 9c): FlashAttention-3 achieves 1,008 TFLOPs/s at 16k, vs. 635 TFLOPs/s for cuDNN and 1,003 TFLOPs/s for Triton. FP8 FlashAttention-3 is roughly 1.55× faster than FP16 FlashAttention-3 at the same configuration (648 TFLOPs/s, Figure 5c), approaching the theoretical 2× throughput gain from FP8.

• Head dim 128 with causal mask (Figure 9d): FlashAttention-3 achieves 881 TFLOPs/s at 16k, vs. 510 TFLOPs/s for cuDNN. The FP8 advantage over FP16 FlashAttention-3 (616 TFLOPs/s, Figure 5d) is 1.43×—smaller than without causal mask, partly because the baseline FP16 kernel already benefits from the persistent kernel and load balancing strategy that the FP8 kernel lacks.

• Head dim 64 without causal mask (Figure 9a): FlashAttention-3 achieves 613 TFLOPs/s at 16k, vs. 511 TFLOPs/s for cuDNN. The advantage is narrower for small head dimensions because the arithmetic intensity is lower, making the kernel more memory-bandwidth-bound and reducing the benefit of FP8's doubled compute throughput.

The paper acknowledges a limitation: "for our benchmarks, FP16 FlashAttention-3 has a persistent kernel and load balancing strategy, while FP8 FlashAttention-3 does not. This partly explains why FP8 FlashAttention-3 does not perform as well for small sequence length and causal masking compared to the FP8 cuDNN kernels." This means the FP8 results represent a lower bound on achievable performance, with further optimizations (persistent kernel, load balancing) expected to close the gap at small sequence lengths.

4.2 Ablation Study: Pipelining and Warp-Specialization Contributions

The paper provides a clean ablation in Table 2 that isolates the contributions of the two key algorithmic innovations—GEMM-softmax pipelining and warp-specialization—for a fixed configuration {batch=4, seqlen=8448, nheads=16, hdim=128} in FP16 without causal mask.

Table 2 results (non-causal FP16, head dim 128, seqlen 8448):

The baseline with neither optimization (effectively Algorithm 1 with warp-specialization removed, i.e., a Hopper-optimized synchronous implementation) is implied to be around 570 TFLOPs/s (the "No Warp-Specialization" configuration also lacks warp-specialization), though the paper presents it as the row with pipelining but no warp-specialization.

• Adding warp-specialization alone (moving from 570 to 582 TFLOPs/s) provides a 2.1% improvement. This modest gain reflects that warp-specialization primarily hides data movement latency via TMA, and at this sequence length (8448), the kernel is compute-bound rather than memory-bandwidth-bound.

• Adding pipelining alone (moving from 582 to 570, or equivalently comparing the two ablated configurations) accounts for the majority of the gain—approximately 12 TFLOPs/s, or 2.1% in the opposite direction. However, this is misleading because the table structure shows pipelining without warp-specialization at 570 TFLOPs/s and warp-specialization without pipelining at 582 TFLOPs/s, suggesting that warp-specialization alone is slightly beneficial but pipelining without warp-specialization regresses.

• The combination (661 TFLOPs/s) is more than the sum of individual contributions: 661 - 570 = 91 TFLOPs/s (16% improvement) over the non-pipelined, non-warp-specialized baseline. This super-additivity is the key finding: warp-specialization and pipelining are complementary—warp-specialization provides the asynchronous TMA loads that keep the circular buffer filled, while pipelining provides the intra-consumer overlap that hides softmax latency. Without warp-specialization, the consumer must issue its own loads, breaking the pipeline's steady-state flow. Without pipelining, warp-specialization's data movement overlap provides limited benefit because the consumer's compute is still serialized by softmax.

The paper's own reporting is somewhat ambiguous about the exact baseline for each ablation. The first column uses FlashAttention-3 (full) vs. configurations with specific features removed, but the paper does not report a "no pipelining, no warp-specialization" configuration explicitly. The implication is that the "No Warp-Specialization" variant (570 TFLOPs/s) effectively represents the Hopper-optimized but synchronous baseline, since without warp-specialization, TMA cannot be used asynchronously and the kernel falls back to synchronous loads.

4.3 Numerical Error Validation

The paper validates numerical accuracy against a reference FP64 implementation, using synthetic data designed to simulate outlier features observed in real LLMs (Dettmers et al., 2022; Sun et al., 2024). The data generation: entries of Q, K, V are drawn from $\mathcal{N}(0, 1) + \mathcal{N}(0, 100) \cdot \text{Bernoulli}(0.001)$, meaning 99.9% of entries have standard deviation 1 and 0.1% have additional noise with standard deviation 10—simulating the outlier pattern.

Table 3 results:

The paper states: "Thanks to block quantization and incoherent processing, FlashAttention-3 in FP8 is 2.6× more accurate than this baseline." The factor 2.6× comes from 2.4e-2 / 9.1e-3 ≈ 2.64.

Key findings from the table and ablation:

• FP16 FlashAttention-3 matches FlashAttention-2's accuracy (both 1.9e-4 RMSE), and both are 1.7× more accurate than standard attention (3.2e-4 / 1.9e-4 ≈ 1.68). This is because FlashAttention keeps intermediate softmax results in FP32, whereas standard attention materializes the attention matrix in FP16.

• Block quantization is the primary accuracy driver in FP8. Ablating block quantization (keeping incoherent processing) raises RMSE from 9.1e-3 to 2.4e-2 (the paper does not report this intermediate value explicitly in the table, but the ablation logic implies block quantization is the dominant factor). The paper reports that removing both techniques returns to the baseline 2.4e-2.

• Incoherent processing provides a modest additional gain, reducing RMSE from 9.3e-3 (no incoherent processing) to 9.1e-3 (with). The 2.2% relative improvement is small in isolation, but the paper argues it matters in the presence of strong outliers, which are common in large LMs.

• FP8 FlashAttention-3 at 9.1e-3 RMSE is still ~48× less accurate than FP16 FlashAttention-3 (9.1e-3 vs. 1.9e-4). The paper does not comment on this absolute gap, but it reflects the fundamental precision tradeoff: FP8 provides 2× throughput at the cost of roughly 48× higher numerical error, even with the best quantization techniques. Whether this error is acceptable depends on the application—the paper does not validate on end-to-end model training or inference quality metrics (perplexity, downstream task accuracy), only on per-operation RMSE.

---

Ablation Studies and Robustness Checks

Attention algorithm design ablations:

• 2-stage GEMM-softmax pipelining vs. synchronous execution: Disabling pipelining while keeping warp-specialization reduces performance from 661 to 582 TFLOPs/s (Table 2), a 12.0% degradation. This validates that the cross-iteration dependency breaking in Algorithm 2 provides substantial gains beyond what warp-specialization alone achieves.

• Warp-specialization alone: Removing warp-specialization while keeping pipelining reduces performance from 661 to 570 TFLOPs/s (Table 2), a 13.8% degradation. However, the paper notes that without warp-specialization, the kernel cannot use asynchronous TMA, so the baseline effectively degrades to a synchronous implementation—meaning this ablation confounds the removal of both warp-specialization and asynchronous data movement.

• 3-stage vs. 2-stage pipelining (Appendix B.3): The 3-stage variant (Algorithm 4) underperforms the 2-stage variant due to (1) the compiler not generating the intended 3-way instruction overlap (only WGMMA1 overlapped with softmax in SASS analysis), and (2) increased register pressure forcing smaller tile sizes that reduce arithmetic intensity. The paper reports this as a qualitative finding without a quantitative table, but states clearly that performance is worse.

FP8 accuracy ablations (Table 3):

• Block quantization contribution: Quantifying the contribution precisely is complicated by the paper's reporting. The ablation removing both block quantization and incoherent processing returns RMSE to 2.4e-2, matching the baseline FP8. Removing only incoherent processing (keeping block quantization) yields 9.3e-3—only modestly worse than the full version at 9.1e-3. This implies that block quantization alone provides the bulk of the accuracy gain (roughly 2.6× over baseline), while incoherent processing provides approximately a 2% additional improvement on this synthetic outlier distribution.

• Incoherent processing contribution: The small gain (9.3e-3 → 9.1e-3) might suggest incoherent processing is unnecessary. However, the paper's synthetic outlier distribution (0.1% outliers with std 10) may not capture the full severity of real LLM outliers. The paper cites Dettmers et al. (2022) and Sun et al. (2024) as evidence that outlier features are a genuine concern, but does not evaluate on activations from a real model, making the 2% improvement a lower bound on the technique's value in practice.

FP8 vs. FP16 implementation maturity: The paper acknowledges that FP8 FlashAttention-3 lacks the persistent kernel and load balancing strategy present in the FP16 implementation. This is not presented as a formal ablation but as an implementation limitation that partly explains why FP8 performance at small sequence lengths and with causal masking trails cuDNN. This is a robustness concern: the FP8 speedup over FP16 is not purely algorithmic but confounded with implementation maturity.

---

Critical Assessment

Claim 1: "FlashAttention-3 achieves 1.5–2.0× speedup over FlashAttention-2 in FP16 forward." This claim is strongly supported by the data in Figure 5, but the speedup is not uniform—it varies substantially across configurations and is largest under specific conditions. The 2.0× figure is an upper bound, not a typical case. At head dim 128 without causal mask (Figure 5c), the speedup ranges from 1.51× at seqlen 512 to 1.75× at seqlen 16k. The 2.0× threshold is only crossed at head dim 256 (2.32× at 16k, Figure 5e). For head dim 64, speedups range from roughly 1.2× to 1.5×. The paper's "1.5–2.0×" summary is accurate as a range but the reader should understand that the high end of this range is specific to large head dimensions and long sequences. The paper does not report on head dimensions below 64 (common in some architectures) or above 256, so the range's boundaries are only partially validated.

Claim 2: "FP8 achieves close to 1.2 PFLOPs/s." The data in Figure 7 supports this: FlashAttention-3 reaches 1,171 TFLOPs (1.171 PFLOPs) at head dim 256, seqlen 16k without causal masking, and 1,099 TFLOPs with causal masking. However, the full FP8 results (Appendix C.2, Figure 9) reveal that this peak is specific to head dim 256. At head dim 128, the peak is 1,008 TFLOPs; at head dim 64, it is 613 TFLOPs. The 1.2 PFLOPs number is thus a configuration-specific peak, not a general throughput guarantee. The paper's language "close to 1.2 PFLOPs/s" is reasonably precise, but the headline number should not be interpreted as typical FP8 throughput—it represents the best-case configuration.

Claim 3: "FP8 FlashAttention-3 achieves 2.6× lower numerical error than baseline FP8 attention." The Table 3 data supports this: 2.4e-2 vs. 9.1e-3. However, this claim has significant limitations that the paper does not fully address:

• The evaluation uses synthetic data, not activations from a real trained model. The outlier distribution (0.1% with std 10 added to std 1 base) may not match actual LLM activation statistics.
• RMSE on attention output is a proxy for end-to-end model quality, but the relationship between per-operation RMSE and downstream task performance (perplexity, accuracy) is not established here.
• FP8 FlashAttention-3 is still ~48× less accurate than FP16 FlashAttention-3 in RMSE terms. The paper emphasizes the 2.6× improvement over baseline FP8 but does not contextualize this against the FP16 baseline—the reader might incorrectly infer that FP8 is "almost as accurate" as FP16.
• The ablation shows that incoherent processing provides only a ~2% improvement on this synthetic distribution, which might not justify the implementation complexity in practice unless real-model outliers are more severe.

Claim 4: "FlashAttention-3 surpasses cuDNN for medium and long sequences." This is supported with qualifications. In FP16 forward without causal mask (Figure 5a, 5c, 5e), FlashAttention-3 exceeds cuDNN at sequence lengths 1k and above across all head dimensions. In FP16 backward (Figure 6), FlashAttention-3 exceeds cuDNN at all tested sequence lengths for head dim 64 and 128. However, in FP8 with causal masking and small head dimensions (Figure 9d, 9b), cuDNN leads at short sequence lengths, and the paper attributes this to the missing persistent kernel in the FP8 implementation. This claim should be understood as applying primarily to FP16 (where the implementation is more mature) and to medium-to-long sequences (where kernel launch overhead is amortized).

Genuine weaknesses and missing experiments:

• Single GPU model (H100 SXM5). All benchmarks are on one specific GPU SKU. While the H100 is the most relevant current-generation datacenter GPU, the paper's claim that the techniques "apply to other hardware accelerators" is unvalidated. H100-specific features (TMA, WGMMA async, setmaxnreg, FP8 Tensor Cores) are central to the design; porting to AMD MI300X or future NVIDIA architectures would require re-engineering.

• No end-to-end training or inference benchmarks. The paper measures kernel-level speed (TFLOPs/s) and per-operation RMSE, but does not report how FlashAttention-3 affects wall-clock training time for a full model, perplexity after training with FP8 attention, or downstream task accuracy. The gap between per-kernel speed and full-model throughput (which includes communication, other layers, and data loading) can be substantial, and the numerical error validation does not establish that FP8 attention is safe for training.

• No comparison to approximation methods on quality. The paper argues that exact attention is preferred because approximations "do not offer the same model quality" (Appendix A), but does not validate this claim or compare FlashAttention-3 against approximate methods on any quality metric. For practitioners deciding between exact attention with FlashAttention-3 and approximate attention (e.g., sparse or low-rank methods), no guidance is provided on the accuracy-speed tradeoff.

• FP8 backward pass not implemented. The paper benchmarks only the FP8 forward pass. The backward pass—which is typically more compute-intensive (2.5× FLOPs) and more sensitive to numerical precision—remains unimplemented in FP8. This is a significant gap because training requires both passes, and the FP8 backward pass may present additional layout or accuracy challenges not addressed here.

• No multi-GPU or distributed settings. The benchmarks are single-GPU. At the scale where attention becomes the bottleneck (very long sequences, large models), multi-GPU tensor parallelism or sequence parallelism is typically used. The paper does not evaluate how FlashAttention-3 interacts with distributed attention schemes like Ring Attention (Liu et al., 2023) that the paper cites as beneficiaries.

• Memory savings not quantified. The paper focuses exclusively on speed (TFLOPs/s). FlashAttention-3 inherits the memory efficiency of the FlashAttention family (eliminating intermediate HBM writes), but the additional memory overhead of the 2-stage pipeline (extra register buffers) and the in-kernel FP8 transpose (temporary SMEM usage) is not quantified. For practitioners operating near memory capacity, this overhead matters.

• Compiler dependence not explored. The 3-stage pipeline's failure is attributed to compiler behavior, but the paper does not experiment with compiler flags, PTX-level control, or inline assembly to force the intended instruction schedule. This is understandable (NVCC is a black box), but it means the optimal pipeline depth is contingent on current compiler heuristics and could change with a CUDA toolkit update—a brittleness that the paper acknowledges but does not quantify.

Summary of conditionality: FlashAttention-3's speedup over FlashAttention-2 is largest when (1) head dimension is large (128–256), (2) sequence length is long (1k–16k, making the kernel compute-bound), (3) causal masking is disabled (fully dense attention), and (4) FP16 is used (where the implementation is most mature). The speedup diminishes at small head dimensions (64), short sequences (512), and with causal masking in FP8. The FP8 accuracy advantage over baseline FP8 attention is validated only on synthetic data with a specific outlier distribution and only for the forward pass. The claim that FP8 FlashAttention-3 is "2.6× more accurate than baseline FP8 attention" should be understood as an RMSE reduction on this specific synthetic benchmark, not a guarantee of better model quality in training or inference.

6. Limitations and Trade-offs

The FP8 Implementation Is Incomplete and Asymmetric

The assumption or constraint. The paper benchmarks and validates FP8 FlashAttention-3 only for the forward pass. The backward pass—which the paper's own FLOPs accounting shows requires 2.5× more computation per attention layer than the forward pass (5 matmuls vs. 2, Section 4.1)—remains entirely unimplemented in FP8 precision. The paper is transparent about this implicitly (the FP8 benchmark figures are labeled "Attention forward speed"), but never explicitly acknowledges that FP8 backward is missing or discusses the technical obstacles that prevented its implementation.

Additionally, the FP8 forward kernel lacks two optimizations present in the FP16 kernel. Section 5 of the paper states:

"for our benchmarks, FP16 FlashAttention-3 has a persistent kernel and load balancing strategy, while FP8 FlashAttention-3 does not. This partly explains why FP8 FlashAttention-3 does not perform as well for small sequence length and causal masking compared to the FP8 cuDNN kernels."

A persistent kernel keeps the threadblock resident on the SM across multiple iterations rather than re-launching, reducing launch overhead and enabling dynamic load balancing (distributing work across SMs proportionally to their completion speed, adapting to SM-to-SM variability). The absence of these features means the FP8 kernel is at an implementation maturity disadvantage relative to both the FP16 FlashAttention-3 kernel and the cuDNN FP8 baseline.

The consequence. The most immediate consequence is that FP8 FlashAttention-3 cannot be used for training—only the full forward-backward pair matters for training workloads, and running FP8 forward with FP16 backward would create a precision mismatch of unknown effect while forfeiting most of the backward pass's compute savings. For inference (forward-only), the FP8 kernel is usable, but its underperformance at short sequence lengths—precisely the regime where many inference deployments operate (prompt processing with sub-1k tokens)—means the practical speedup over FP16 FlashAttention-3 may be substantially smaller than the headline FP8 numbers suggest.

More subtly, the missing persistent kernel confounds the comparison with cuDNN. The paper's FP8 results show that FlashAttention-3 trails cuDNN at small sequence lengths with causal masking (e.g., Figure 9b: at seqlen 512, cuDNN achieves 234 TFLOPs vs. FlashAttention-3's 164 TFLOPs for head dim 64 with causal mask). Without a controlled ablation, it is impossible to determine how much of this gap is due to algorithmic design choices (the in-kernel transpose, the specific 2-stage pipeline, the register allocation) versus the missing persistent kernel optimization. A cuDNN-to-FlashAttention-3 comparison that is partly an algorithm comparison and partly an implementation maturity comparison weakens the claim that FlashAttention-3's algorithmic approach is superior.

What evidence exists in the paper. The FP8 benchmark figures (Figure 7 main text, Figure 9 Appendix C.2) all specify "forward speed." There are no FP8 backward benchmarks anywhere in the paper. The missing persistent kernel is acknowledged in a footnote (Section 5, page 12), but the missing backward pass receives no explicit discussion—it is simply absent. The paper's own emphasis on the backward pass for FP16 (Figure 6, with detailed speedup analysis) makes this omission conspicuous.

Mitigation status. The paper does not address either gap—no discussion of FP8 backward pass plans, no explanation of why persistent kernel is missing from FP8, and no roadmap for closing either gap. The acknowledgments section thanks the cuDNN team and the CUTLASS team for help understanding Hopper's programming model, and the FP8 layout co-design discussion (Section 3.3) demonstrates deep engagement with FP8 WGMMA constraints, suggesting the backward pass omission is likely a scope issue rather than a fundamental obstacle. However, for a practitioner evaluating whether to adopt FP8 FlashAttention-3 for training, the paper provides no information on when or whether FP8 backward will be available, nor any estimate of the technical difficulty involved.

---

Numerical Accuracy Is Validated Only on Synthetic Data, Not on End-to-End Model Quality

The assumption or constraint. The paper's accuracy validation (Section 4.3, Table 3) uses synthetic data with statistically generated outlier features—entries of Q, K, V are drawn from $\mathcal{N}(0, 1) + \mathcal{N}(0, 100) \cdot \text{Bernoulli}(0.001)$. This is a plausible simulation of LLM activation distributions, but it is not actual LLM activations. The paper never measures how FP8 FlashAttention-3 affects end-to-end model quality—no perplexity numbers after training with FP8 attention, no downstream benchmark scores, no comparison of FP16-trained vs. FP8-trained model checkpoints, and no measurement of gradient statistics in the backward pass.

Moreover, the per-operation RMSE metric reported in Table 3 has an unclear relationship to model quality. FP8 FlashAttention-3 achieves RMSE of 9.1e-3 on attention outputs. This is 2.6× better than baseline FP8 attention (2.4e-2), but it is still approximately 48× worse than FP16 FlashAttention-3 (1.9e-4). Whether a 48× increase in attention output error translates to a 0.1% or a 10% degradation in training loss or downstream accuracy is unknown—the relationship between per-layer numerical error and end-to-end model quality is nonlinear and architecture-dependent, and small errors in early attention layers can compound through the network.

The consequence. A practitioner considering FP8 FlashAttention-3 for training faces a genuine uncertainty: does the 2× throughput gain justify a 48× increase in attention output error? The paper provides no evidence to answer this. In the worst case, FP8 attention could cause training instability (diverging loss, NaN gradients) or silently degrade model quality by several perplexity points—losses that would far outweigh the compute savings. In the best case, neural network training is robust to per-operation noise and the 2× speedup translates directly to 2× faster training with negligible quality impact. The paper's RMSE measurement on synthetic data does not distinguish between these scenarios.

The synthetic outlier distribution may also misrepresent real LLM activation statistics in important ways. The paper uses 0.1% outliers with std 10 added to a std 1 base. Real LLMs exhibit outlier features that are more structured—often concentrated in specific feature dimensions or tokens (Dettmers et al., 2022)—and their magnitude and sparsity patterns vary across layers, model sizes, and training stages. The paper's incoherent processing ablation shows only a ~2% RMSE improvement (9.3e-3 → 9.1e-3) on this synthetic data, suggesting the technique provides minimal benefit. But on real LLM activations with more severe or differently distributed outliers, incoherent processing could matter substantially more—or block quantization alone might suffice entirely. The synthetic benchmark provides only weak evidence for either conclusion.

What evidence exists in the paper. Table 3 and the description of the synthetic data generation in Section 4.3 are the only accuracy evidence. There is no end-to-end training experiment, no perplexity measurement, and no comparison to a model trained with FP16 attention. The paper's related work section (Appendix A) notes that "quantization during training is still challenging as higher precision is typically required for stable training," but does not evaluate whether FlashAttention-3's specific combination of block quantization and incoherent processing overcomes this challenge.

Mitigation status. Not addressed. The paper validates that FP8 FlashAttention-3 computes attention with lower error than a naive FP8 implementation—which is necessary but not sufficient to claim FP8 attention is safe for training or inference at scale. The paper frames the FP8 contribution as reducing numerical error by 2.6× compared to the baseline, which is accurate on its own terms but leaves the critical question unanswered. Section 5 lists "understanding the effects of low-precision attention in large-scale training" as a limitation and future work direction, explicitly acknowledging this gap. For a practitioner, this means the FP8 results should be treated as promising but unvalidated for production training; adoption should be preceded by careful model-quality evaluation on the specific architecture and dataset of interest.

---

All Benchmarks Are on a Single GPU Model with No Cross-Architecture Validation

The assumption or constraint. Every benchmark in the paper is run on the H100 SXM5 GPU (80 GB, 700W). The paper acknowledges this but claims generality: "our algorithm is operative for any GPU architecture with sufficiently robust asynchronous execution and low-precision capabilities" (footnote on page 2). This claim is not validated. The three core techniques—TMA-based asynchronous data movement, WGMMA-based asynchronous matrix multiplication with 2-stage pipelining, and FP8 layout co-design via in-kernel transpose—all rely on Hopper-specific hardware features that have no direct equivalents on prior NVIDIA architectures (Ampere, Turing) or on non-NVIDIA GPUs (AMD CDNA, Intel Xe).

Specifically:

• TMA is a Hopper-specific hardware unit for asynchronous memory copy. On Ampere, asynchronous copy uses the cp.async instruction, which has different semantics, lower throughput, and different shared memory addressing. On AMD GPUs, the equivalent functionality may not exist in the same form.
• WGMMA asynchrony relies on the Hopper Tensor Core's ability to operate asynchronously from the CUDA cores and to source operands directly from shared memory. Ampere Tensor Cores are synchronous; the warp must stall until the matrix multiplication completes.
• setmaxnreg for dynamic register reallocation between warpgroups is a Hopper-specific instruction.
• FP8 Tensor Cores are a Hopper (and Blackwell) feature; Ampere GPUs support only FP16/BF16/TF32.

The paper's warp-specialization and pipelining ideas are algorithmically general, but their effectiveness depends on the hardware's ability to execute data movement, matrix multiplication, and element-wise operations concurrently. On architectures with less asynchronous execution capability, the overlapping might not work, and the register and shared memory overhead of the 2-stage pipeline (extra Snext buffer, deeper circular SMEM buffer) might make the kernel slower than a simpler synchronous implementation.

The consequence. The paper's speedup numbers should be understood as H100-specific. A practitioner using A100 GPUs (still the most common datacenter GPU for LLM training as of 2024) cannot expect the 1.5–2.0× speedup—at best, some of the warp-specialization ideas might provide marginal gains, but the core asynchronous overlap that drives most of the performance improvement requires Hopper hardware. Similarly, organizations evaluating AMD MI300X GPUs for training cannot use FlashAttention-3 without a port that may or may not be feasible depending on ROCm's support for equivalent asynchronous primitives.

The paper's open-source release partially mitigates the portability concern—other developers can attempt ports—but the algorithmic design is so tightly coupled to Hopper PTX instructions (TMA, WGMMA, LDSM, STSM, setmaxnreg) that a port to a different architecture would likely require rethinking the core overlap strategies, not just translating instruction sequences.

What evidence exists in the paper. All benchmark figures (Figures 5–9) specify "H100 80GB SXM5" in their captions. The system configuration (Appendix C.1) confirms CUDA 12.3, cuDNN 9.1.1.17, CUTLASS 3.5—all Hopper-targeted. There are no benchmarks on A100, H200, GH200, or any non-NVIDIA GPU. The paper's only evidence for cross-architecture generality is the footnote assertion, which is not supported by any experimental or analytical argument.

Mitigation status. The paper does not attempt to mitigate this—there are no portability abstractions, no discussion of which components are Hopper-specific vs. general, and no roadmap for supporting other architectures. The footnote acknowledges that "our algorithm is operative for any GPU architecture with sufficiently robust asynchronous execution and low-precision capabilities," which is a conditional claim that leaves the burden of determining "sufficiently robust" to the reader. The open-source release (under a permissive license) is the primary mitigation strategy, enabling community ports, but the paper provides no guidance on what a port would entail.

---

The FP8 In-Kernel Transpose Overhead Is Not Quantified or Ablated

The assumption or constraint. To satisfy the FP8 WGMMA's k-major layout constraint for the second GEMM ($P \times V$), the paper performs an in-kernel transpose of $V$ tiles using LDSM/STSM instructions (Section 3.3, Challenge 1). The paper argues that this transpose latency is hidden by scheduling it during the consumer's WGMMA execution—"after the first iteration, the transpose of the next V tile is executed in the shadow of the two WGMMAs that involve the preceding V and current K tile." This is a latency-hiding claim: the transpose costs cycles, but those cycles overlap with computation that would otherwise be idle, so the net throughput impact is zero or negligible.

The paper also acknowledges a subtle complication: for FP8 (8-bit entries), the LDSM/STSM instructions—which operate on 16-bit entries—require packing two 8-bit values together, and "the transpose versions of LDSM/STSM cannot split packed 8-bit entries, which necessitates certain register movements in between LDSM and STSM to actually perform a tile-wise transpose; we omit the details." These additional register movements consume instruction issue slots and register bandwidth, and may not be fully hidable under the WGMMAs.

The consequence. If the in-kernel transpose is not perfectly hidden, it imposes a throughput tax on the FP8 kernel that is absent in the FP16 kernel (where no transpose is needed because FP16 WGMMA accepts both mn-major and k-major layouts). This tax would reduce the effective FP8 speedup below the theoretical 2×—and since the transpose complexity grows with tile size, the tax could be proportionally larger at the larger tile sizes needed for long sequences, exactly where the FP8 peak throughput numbers are measured.

The paper provides no measurement of what fraction of SM cycles the transpose consumes. Without such a measurement, a practitioner cannot determine whether an alternative approach (e.g., a preprocessing transpose of $V$ in a separate kernel, as briefly considered and rejected) might actually be faster in some regimes—the paper's rejection of the preprocessing transpose as "too wasteful in a memory-bound situation such as inference" is a qualitative argument without quantitative support.

Moreover, the additional shared memory traffic for the transpose (reading $V_j$ from SMEM, rearranging in registers, writing back to SMEM) competes with the consumer's SMEM bandwidth for K and V loads. On H100, SMEM bandwidth is abundant (31 TB/s aggregate), but it is a shared resource, and contention between producer transpose and consumer loads could create subtle stalls not captured by the high-level TFLOPs/s benchmark.

What evidence exists in the paper. None. There is no ablation comparing FP8 FlashAttention-3 with in-kernel transpose vs. FP8 FlashAttention-3 with a preprocessing transpose kernel. There is no measurement of SMEM bandwidth utilization, no profiling of transpose-related stalls, and no discussion of the register movement overhead described as "certain register movements" whose "details" are omitted. The FP8 benchmarks (Figures 7, 9) show end-to-end TFLOPs/s but do not decompose the time into compute, transpose, and memory components.

Mitigation status. The paper asserts that the transpose is "in the shadow of" the WGMMAs, implying zero net cost, but provides no evidence. The register movement complication is mentioned and then dismissed ("we omit the details"). This is a trust-me gap: the paper asks the reader to accept that the transpose is free based on the architectural reasoning that the warp scheduler can interleave the producer's transpose instructions with the consumer's WGMMA instructions, but does not validate this reasoning with profiling data or SASS analysis (unlike the 2-stage pipelining, which does receive SASS validation in Appendix B.2). A practitioner evaluating whether to adopt the in-kernel transpose approach for their own FP8 kernel would need to conduct their own profiling to quantify the overhead—the paper provides no guidance.

---

The 2-Stage Pipeline's Optimality Is Contingent on Current Compiler Behavior, Not a Hardware Limit

The assumption or constraint. The paper establishes that the 2-stage GEMM-softmax pipeline outperforms the 3-stage variant (Appendix B.3), and draws the conclusion that 2-stage represents the optimal pipelining depth for attention on Hopper. However, the cause of the 3-stage variant's underperformance is partly compiler-dependent: "It's not clear why the compiler chooses to reorder instructions in this way" (referring to NVCC's failure to schedule the second WGMMA concurrently with softmax, even though the algorithm's dependency structure would permit it). This means the optimal pipeline depth is not purely a hardware constraint—it is a joint function of the hardware's execution capabilities and the compiler's instruction scheduling heuristics.

The paper's SASS analysis (Appendix B.2) confirms that for the 2-stage pipeline, the compiler generates the intended overlap: "WGMMA and non-WGMMAs are executed in parallel," with softmax "reordered to the very beginning, even before the first WGMMA." For the 3-stage variant, the paper reports that "only the first WGMMA is overlapped with softmax, while the second WGMMA is not." The difference is not in what the hardware can do—both pipelines could theoretically overlap the second WGMMA—but in what the compiler chooses to emit.

The consequence. The optimal pipeline depth is brittle with respect to compiler versions and flags. A future release of NVCC (CUDA 12.4, 13.0) could change its instruction scheduling heuristics, potentially making the 3-stage variant outperform the 2-stage variant, or introducing new reorderings that break the 2-stage variant's carefully crafted overlap. Conversely, a practitioner using different compilation flags (different optimization level, different target architecture flags) might observe different performance characteristics. The paper's ablation (Table 2) and pipelining design therefore represent a snapshot of performance under one specific toolchain (CUDA 12.3, whatever default optimization flags NVCC uses at -O3), not a fundamental characterization of the algorithm's capabilities.

This is not a hypothetical concern. GPU kernel performance is notoriously sensitive to compiler heuristics—register allocation, instruction scheduling, loop unrolling—and minor toolchain changes can shift performance by 5–15%. For a kernel that achieves a 16% improvement from pipelining (661 vs. 570 TFLOPs/s in the ablation), a compiler-induced regression of even a few percent would erase a meaningful fraction of the claimed gains.

More fundamentally, the paper's methodology—designing high-level CUDA C++ and then inspecting SASS to verify the compiler generated the intended instruction schedule—represents a loss of control over the algorithm's execution. The programmer intends a specific overlap pattern; the compiler may or may not honor that intent. The paper does not explore alternatives that would provide stronger guarantees: inline PTX assembly for critical sections, compiler hints or pragmas to constrain instruction scheduling, or direct SASS-level programming. The implicit assumption is that CUDA C++ plus SASS verification is sufficient, but the 3-stage variant's failure shows that at some complexity level, the compiler becomes an unreliable partner.

What evidence exists in the paper. The SASS analysis in Appendix B.2 validates the 2-stage pipeline's intended execution. The description of the 3-stage variant's failure in Appendix B.3 explicitly blames the compiler: "SASS code shows that only the first WGMMA is overlapped with softmax, while the second WGMMA is not. It's not clear why the compiler chooses to reorder instructions in this way." This is the only compiler-dependence evidence. There is no experiment varying compiler versions or flags, no comparison of NVCC vs. NVCC with different optimization settings, and no attempt to use inline PTX to force the intended 3-way overlap.

Mitigation status. The paper acknowledges the issue but does not address it. The SASS verification is presented as a validation step ("SASS shows that the 2-stage pipelining idea works as expected"), but the 3-stage failure is treated as a mysterious negative result rather than a sign that the methodology has a compiler-dependence problem. The paper does not suggest future work on compiler improvements, on using lower-level programming (PTX, SASS), or on developing compiler-robust pipelining patterns. A practitioner adopting FlashAttention-3 should be aware that future CUDA toolkit updates may unpredictably affect performance, and that reproducing the paper's exact TFLOPs/s numbers may require matching the paper's exact software environment.

---

No Consideration of Latency vs. Throughput Tradeoffs

The assumption or constraint. The paper measures success exclusively in throughput (TFLOPs/s), which is the appropriate metric for training workloads and batched inference where many attention computations are in flight simultaneously. However, the algorithmic techniques that increase throughput—particularly the 2-stage pipeline and pingpong scheduling—introduce additional latency per attention operation because they add pipeline stages that must be filled and drained (the prologue and epilogue in Algorithm 2). The paper never measures or reports the latency of a single attention forward pass (in microseconds or milliseconds for a single query) or discusses how the pipelining affects end-to-end latency.

This matters for two distinct deployment scenarios:

1. Autoregressive inference (token-by-token generation). During decode, attention operates on a single query token against a growing KV cache. The sequence length processed per attention call is typically short (often a single token), and latency—not throughput—determines the user-perceived generation speed. The 2-stage pipeline's prologue/epilogue overhead, the warp-specialization's threadblock launch overhead, and the absence of the persistent kernel in FP8 all exacerbate the latency penalty at short sequence lengths.

2. Interactive applications with small batch sizes. Even for prompt processing (encoding a user's input), if the batch size is 1 (typical for interactive chat), the GPU's SMs may be underutilized, and per-kernel launch overhead and pipeline fill/drain time become proportionally larger.

The paper's benchmarks use a fixed total token count of 16k, with batch size adjusted inversely to sequence length (Section 4.1). This ensures high SM occupancy, which is the right methodology for measuring peak throughput but systematically hides latency effects that would appear at small batch sizes or short sequence lengths.

The consequence. The 1.5–2.0× speedup figures may not translate to latency improvements for inference—and could even represent a latency regression if the pipelining overhead outweighs the compute savings at low occupancy. The paper's own data provides suggestive evidence: at sequence length 512 (the shortest benchmarked), the speedup over FlashAttention-2 is consistently smaller than at longer sequence lengths. For head dim 128 without causal mask, the speedup is 1.51× at seqlen 512 vs. 1.75× at seqlen 16k. For head dim 64, the speedup at seqlen 512 is roughly 1.2× (340 vs. 282 TFLOPs/s, Figure 5a). This pattern is consistent with fixed overheads (kernel launch, pipeline prologue/epilogue) becoming a larger fraction of total time at short sequence lengths. The paper attributes this to the kernel transitioning from "launch-overhead-bound to compute-bound regimes," but does not discuss the implications for latency-sensitive deployments.

The paper's title and abstract position FlashAttention-3 as relevant to "large language models and long-context applications." This focus on long contexts is appropriate for the paper's contributions—the overlapping techniques are designed for compute-bound regimes—but the absence of latency analysis means a reader cannot determine whether FlashAttention-3 is suitable for the short-sequence, latency-critical portion of the Transformer workload (token-by-token generation), which dominates the end-to-end wall-clock time of autoregressive inference even for long-context models.

What evidence exists in the paper. The paper provides no latency measurements—no millisecond timings for single-operation attention, no analysis of how the 2-stage pipeline's prologue/epilogue scales with sequence length, and no discussion of inference-specific considerations. The only temporal measurements are throughput (TFLOPs/s) averaged over batch configurations designed to maximize SM utilization. The paper's related work (Appendix A) cites LeanAttention (Sanovar et al., 2024) as addressing the decode-phase bottleneck via load balancing, acknowledging that inference has distinct optimization requirements, but does not engage with how FlashAttention-3 performs in that regime.

Mitigation status. Not addressed. Section 5 lists "optimizing for LLM inference" as a future work direction, explicitly acknowledging that the current work focuses on training and long-context scenarios where throughput dominates. This is a reasonable scope limitation for a paper whose primary contribution is compute-bound kernel optimization, but it means that the paper's claims of speedup should be understood as applying primarily to training and high-throughput inference, not to latency-sensitive single-query inference. A practitioner deploying FlashAttention-3 in an interactive serving system should benchmark latency separately and may find that FlashAttention-2 (with its simpler, lower-overhead synchronous design) is actually faster for the decode phase.

7. Implications and Future Directions

How This Work Changes the Landscape

FlashAttention-3 fundamentally reframes GPU kernel design for attention—and more broadly for any fused operator combining GEMMs with element-wise computation—from a memory-bandwidth-centric discipline to a compute-throughput-centric one. The original FlashAttention's landmark contribution was recognizing that attention is memory-bandwidth-bound and restructuring it to avoid materializing the full attention matrix in HBM. That insight remains valid, but FlashAttention-3 demonstrates that on modern hardware, the binding constraint shifts: at the sequence lengths where attention matters most, the kernel becomes compute-bound, and the bottleneck is not data movement but heterogeneous compute unit utilization—the fact that different hardware units (Tensor Cores, CUDA cores, multi-function unit, TMA) operate at dramatically different throughputs and must be kept simultaneously busy to achieve high utilization.

This is a conceptual shift with implications beyond attention. Any fused kernel that chains a high-throughput GEMM with a low-throughput element-wise operation—LayerNorm, activation functions, dropout, rotary embedding—faces the same structural problem on Hopper and future architectures. The paper's diagnostic framework (quantify the FLOP ratio between GEMM and non-GEMM operations, compare to the throughput ratio of the hardware units that execute them, identify the ceiling operation, and restructure the dependency graph to overlap it) is transferable methodology. The 256× throughput gap between FP16 matrix multiplication (989 TFLOPs) and exponential (3.9 TFLOPs) is specific to softmax, but the principle generalizes: any element-wise operation with throughput substantially below the Tensor Core peak will consume a disproportionate fraction of cycle time unless algorithmically overlapped. This reframes GPU kernel optimization as fundamentally a scheduling problem—how to arrange the dependency graph so that the warp scheduler always has instructions available for every hardware unit—rather than primarily a memory hierarchy problem.

The paper also shifts the status of warp-specialization from a systems trick to an algorithm design primitive. Prior work (Bauer et al., 2011; CUTLASS GEMM implementations) used warp-specialization primarily for producer-consumer data movement overlap. FlashAttention-3 extends it to compute-compute overlap: the pingpong scheduling between two consumer warpgroups ensures that while one subgroup's softmax executes on CUDA cores, the other subgroup's GEMMs execute on Tensor Cores, and the producer simultaneously loads future tiles via TMA. The three-role specialization in the backward pass (producer, consumer, dQ-writer) further demonstrates that warp-specialization is not merely about hiding memory latency—it is about decoupling independent instruction streams so that heterogeneous hardware units never stall waiting for each other. The setmaxnreg-based register reallocation between warpgroups—starving the producer to feed the consumer—elevates register allocation from a compiler concern to an explicit algorithmic policy that directly determines feasible tile sizes and pipeline depths.

Perhaps most importantly, the paper establishes that algorithmic intent does not reliably survive compilation in the asynchronous regime. The 3-stage pipeline's failure—where the compiler refused to generate the intended 3-way instruction overlap even though the dependency structure permits it—is a scientifically significant negative result. It reveals a control gap in the current GPU programming stack: between CUDA C++ (where the programmer expresses algorithmic intent) and SASS (where that intent is realized or subverted by compiler heuristics), there is no reliable mechanism to enforce specific instruction-level overlap patterns. This gap will become increasingly consequential as GPU architectures add more asynchronous execution units (Blackwell's FP4 Tensor Cores, future dedicated quantization units, etc.) and the potential overlap patterns proliferate. The paper's methodology—design in CUDA C++, verify in SASS, accept that some intended overlaps will fail—is pragmatic but unsustainable for deeper pipelines or more complex fused operators. This finding implicitly calls for new programming abstractions (compiler pragmas for instruction scheduling constraints, higher-level asynchronous dataflow languages, or PTX-level control) that give kernel developers deterministic control over which operations overlap.

The paper also resolves a specific tension in the prior literature. FlashAttention-2 achieved only 35% utilization on H100, which was sometimes attributed to inherent limitations of attention (irregular memory access, low arithmetic intensity) or to the immaturity of Hopper software support. FlashAttention-3 demonstrates that neither explanation is correct: 75% utilization is achievable with the exact same mathematical operation, proving that the missing utilization was a software architecture problem, not a hardware or algorithmic one. This matters because it shifts the burden of proof for future attention optimizations: if a new implementation achieves less than ~75% utilization on Hopper-equivalent hardware, the default assumption should be that it is leaving performance on the table through insufficient exploitation of asynchrony and low-precision, not that attention is inherently inefficient.

Finally, the paper demonstrates that FP8 attention is viable for the forward pass with careful quantization engineering, reducing the RMSE gap to baseline FP8 by 2.6× through block quantization and incoherent processing. This does not yet prove FP8 attention is safe for training—the backward pass and end-to-end model quality remain unvalidated—but it establishes that the numerical accuracy barrier is surmountable with techniques adapted from the post-training quantization literature, applied at training time. The broader implication is that the trend toward lower-precision arithmetic (FP16 in Pascal, BF16 in Ampere, FP8 in Hopper, FP4 in Blackwell) is not limited to GEMM-heavy operations like fully-connected layers; with appropriate algorithmic co-design (layout permutations, block-wise scaling, incoherent processing), even precision-sensitive operations like attention can participate in the low-precision transition.

Follow-Up Research This Work Enables

End-to-end training validation of FP8 attention, starting with the missing backward pass. The most immediate and consequential gap this paper leaves is the absence of an FP8 backward pass for attention. A strong follow-up would implement FP8 FlashAttention-3 backward, measure the per-gradient RMSE against FP16 (analogous to the Table 3 forward-pass validation), and then run a controlled training experiment: pretrain a moderate-scale language model (e.g., 1B–7B parameters) twice, once with FP16 attention and once with FP8 attention (forward and backward), and compare training loss curves, downstream benchmark scores, and training stability metrics (gradient norm trajectories, loss spike frequency). The paper's FP8 forward-pass RMSE of 9.1e-3 provides a starting point, but the backward pass involves computing dS through softmax (which amplifies errors near the attention probability peaks) and accumulating dQ across multiple threadblocks (which involves FP32 atomic additions). The question is not just whether FP8 backward can be implemented, but whether the accumulated error across training steps causes divergence, reduced effective learning rate, or degraded final model quality. If FP8 attention training proves stable at scale, the practical impact would be immediate: a ~2× speedup in attention computation during training, directly translating to faster iteration cycles or longer-context training within the same compute budget.

Compiler-aware asynchronous pipeline design through inline PTX or scheduling pragmas. The 3-stage pipeline's failure due to compiler non-cooperation is not a dead end—it is an invitation to gain finer control over instruction scheduling. A research-strength follow-up would reimplement the critical sections of FlashAttention-3's consumer mainloop in inline PTX assembly (which the paper already uses indirectly through CUTLASS abstractions) to deterministically control the ordering of WGMMA, exponential, and FMA instructions. The experiment would compare the PTX-controlled 2-stage pipeline against the existing CUDA C++ version (to quantify how much performance is lost to suboptimal compiler scheduling even in the 2-stage case), and then implement the 3-stage pipeline under PTX control to determine whether the theoretical 3-way overlap is achievable on hardware when the compiler is bypassed. If PTX control recovers the 3-stage pipeline's expected performance, the implication is that higher-level programming models (CUDA C++, Triton) need new abstractions—such as #pragma schedule wgmma_before_softmax or dataflow graph annotations—to give kernel developers reliable control over asynchronous overlap without dropping to assembly. If even PTX cannot achieve 3-way overlap (because of hardware warp scheduler limitations, not compiler heuristics), the implication is that the 2-stage pipeline represents a genuine hardware ceiling, and future GPU architectures should be designed with deeper asynchronous pipelines in mind.

Systematic characterization of the tile-size/pipeline-depth Pareto frontier. The paper's 2-stage vs. 3-stage result (Appendix B.3) establishes that pipeline depth and tile size compete for register resources, and that for one specific configuration (FP16, head dim 128, seqlen 8448), 2-stage with the paper's chosen tile sizes is optimal. But the shape of the tradeoff curve—how the optimal pipeline depth varies with head dimension, sequence length, precision, and GPU register file size—is unexplored. A strong analytical follow-up would measure throughput for a grid of tile sizes ($B_r \times B_c$) at pipeline depths 1, 2, 3, and 4 (if registers permit) across the full parameter space the paper benchmarks (head dims 64/128/256, sequence lengths 512–16k, with/without causal mask, FP16/FP8). The output would be a roofline-like model for asynchronous attention that predicts, for a given hardware configuration (register count, SMEM size, Tensor Core throughput, MFU throughput), the optimal tile size and pipeline depth. Such a model would guide both kernel auto-tuning (replacing the paper's implicit "profile and pick" approach) and hardware design (quantifying the register pressure vs. pipeline depth tradeoff for future architectures). The paper's specific measurements—661 TFLOPs at a specific tile configuration—are a single point on this frontier; characterizing the full frontier would transform the empirical result into a design tool.

Integration with distributed attention and long-context training at scale. The paper cites Ring Attention (Liu et al., 2023) and its variants as direct beneficiaries of FlashAttention-3's speedup, but provides no distributed benchmarks. A deployment-focused follow-up would integrate FlashAttention-3 into a distributed training framework supporting sequence parallelism (e.g., Megatron-LM with context parallelism) and measure end-to-end training throughput for long-context models (e.g., 32k–128k tokens) on 8–64 H100 GPUs. The key measurements would be: (1) how the 1.5–2.0× single-GPU speedup translates to multi-GPU throughput given communication overhead (all-reduces in tensor parallelism, P2P sends in sequence parallelism), (2) whether FlashAttention-3's asynchronous execution model creates new opportunities to overlap attention computation with cross-GPU communication (since TMA and WGMMA execute independently from the CUDA cores that issue NCCL communications), and (3) whether the FP8 forward pass can be used in a distributed training setting where gradient all-reduces remain in FP16. The paper's single-GPU benchmarks establish the potential; the distributed benchmarks would establish the realized gain in the training scenarios that the paper motivates—"modeling and reasoning over multiple long documents, files in large codebases, high-resolution images, audio, and video."

Extension to non-transformer attention patterns and alternative architectures. The paper's optimizations are applied to standard multi-head attention, but the paper explicitly notes that alternative architectures (Mamba, RWKV, RetNet, Jamba, etc.) "still employ many layers of attention" and that FlashAttention-3 "will be useful to speed up these alternative architectures" (Appendix A). A natural follow-up would benchmark FlashAttention-3 on the specific attention patterns used in hybrid SSM-transformer models: multi-query attention (MQA), grouped-query attention (GQA), and multi-head latent attention (MLA, from DeepSeek-V2). These variants change the K and V head counts relative to Q, which affects both the arithmetic intensity and the register pressure (since fewer K/V tiles need to be loaded and processed). The paper states that it "adjust[s] the tensor indexing to avoid duplicating K and V in HBM" following FlashAttention-2's approach, but the performance impact of the 2-stage pipeline under GQA's reduced K/V bandwidth requirements is not measured. A concrete experiment: benchmark FlashAttention-3 FP16 at head dim 128, seqlen 8k, with GQA ratios of 1:1 (MHA), 4:1, 8:1, and 16:1, measuring whether the speedup over FlashAttention-2 increases (because less time spent loading K/V) or decreases (because lower arithmetic intensity makes the kernel more memory-bound and reduces the benefit of compute-compute overlap). This would provide specific guidance for the growing number of models using GQA as the default attention configuration.

Robustness of FP8 attention to real LLM activation distributions. The paper's FP8 accuracy validation uses synthetic outliers ($\mathcal{N}(0,1) + \mathcal{N}(0,100) \cdot \text{Bernoulli}(0.001)$), which is a reasonable first approximation but almost certainly does not match the complex, layer-dependent, training-stage-dependent statistics of real LLM activations. A thorough follow-up would instrument a training run of a real LLM (e.g., Llama-2 7B or a similar open model) to capture the actual distributions of Q, K, V activations at multiple layers (early, middle, late) and training steps (beginning, middle, end of training), then measure FP8 FlashAttention-3's RMSE on these real distributions with and without block quantization and incoherent processing. The finding that incoherent processing provides only a 2% improvement on synthetic data (9.3e-3 vs. 9.1e-3) may reverse on real activations where outlier structure is more complex. Additionally, the experiment would measure the correlation between per-layer RMSE and end-to-end training metrics (loss, gradient norm) to determine whether certain layers (e.g., early layers with larger activation magnitudes) are more sensitive to FP8 error and might benefit from selective FP16 attention while other layers run in FP8. The paper's block quantization and incoherent processing are validated as techniques; the missing piece is understanding when and where they are necessary, and whether simpler approaches (block quantization alone, with sufficiently small blocks) suffice for most layers.

Practical Applications and Downstream Use Cases

Training long-context LLMs with reduced time-to-solution. For organizations training large language models with context lengths of 8k–128k tokens, FlashAttention-3's 1.5–2.0× attention speedup directly reduces per-step training time. At head dim 128 and sequence length 8k, FlashAttention-2 delivers 370 TFLOPs/s while FlashAttention-3 delivers 646 TFLOPs/s (Figure 5c)—a 1.75× improvement. Since attention typically constitutes 20–40% of total training FLOPs in a Transformer (varying with sequence length and model architecture), this translates to a ~10–25% reduction in total training time. For a training run that would take 30 days on 512 H100 GPUs with FlashAttention-2, FlashAttention-3 could reduce this to ~23–27 days, or alternatively allow training with 1.75× longer context within the same wall-clock budget. The benefit is largest at the head dimensions (128–256) and sequence lengths (1k–16k) that are typical for modern LLMs, making this a drop-in improvement with no algorithmic change to the model. The open-source release and planned PyTorch/HuggingFace integration (Section 1) lower the adoption barrier to near-zero for users of these frameworks.

FP16 inference serving for long-context prompt processing. For inference serving systems that process long prompts (e.g., document question-answering, codebase analysis, multi-turn conversation history), the prompt encoding phase is attention-compute-bound and benefits directly from FlashAttention-3's FP16 throughput. At head dim 128, sequence length 16k without causal mask, FlashAttention-3 (648 TFLOPs/s) is 1.75× faster than FlashAttention-2 (370 TFLOPs/s) and 1.09× faster than cuDNN (595 TFLOPs/s) per Figure 5c. For a serving system processing 1,000 prompts averaging 8k tokens each per hour, this 1.75× throughput improvement reduces GPU-hours by ~43% for the attention portion of prompt processing—which, depending on the model architecture, could translate to a 10–20% reduction in total serving cost. The advantage over cuDNN is particularly significant for open-source serving frameworks (vLLM, TGI, TensorRT-LLM) that may not bundle cuDNN's closed-source attention implementation; FlashAttention-3 provides a vendor-neutral path to H100-optimized attention.

FP8 inference for extremely long contexts under throughput constraints. For applications requiring attention over very long sequences (32k–128k tokens)—genomic sequence modeling, video frame attention, large-scale retrieval over document collections—the FP8 forward pass's near-1.2 PFLOPs/s throughput (Figure 7) provides roughly 2× the throughput of FP16 FlashAttention-3 and 3–4× that of FP16 FlashAttention-2. At 128k tokens without causal mask, the attention FLOPs scale quadratically, and the FP16 peak of 740 TFLOPs/s means each attention forward pass requires ~0.9 seconds on an H100 (for a single head, typical head dimension). FP8 at 1.2 PFLOPs/s reduces this to ~0.55 seconds—still substantial, but 1.6× faster. For applications that are throughput-bound rather than latency-bound (batch processing of many long sequences, offline embedding computation), FP8 FlashAttention-3 effectively doubles the batch size that can be processed per unit time. The primary caveat is that this applies to forward-only inference pipelines; training still requires FP16 until FP8 backward is validated.

Faster iteration cycles for attention kernel research itself. FlashAttention-3's open-source implementation (using CUTLASS primitives for WGMMA and TMA) provides a modular, reusable template for attention kernel experimentation. Researchers exploring new attention variants (sparse patterns, linearized attention with custom kernels, block-sparse attention, sliding window + global attention combinations) can build on FlashAttention-3's warp-specialized infrastructure rather than starting from scratch with raw PTX or CUDA. The paper's ablation methodology (Table 2) and SASS verification approach (Appendix B.2) also provide a template for evaluating whether new attention algorithms successfully exploit asynchrony—researchers can measure TFLOPs/s with and without pipelining and inspect SASS to verify that their intended overlaps survive compilation. This lowers the barrier to entry for hardware-aware attention research, which has traditionally required deep GPU architecture expertise.