Flash Attention v2 made attention computationally tractable for long sequences by keeping the working set in SRAM rather than writing intermediate attention matrices back to HBM. That eliminated the O(n^2) memory access pattern that made naive attention prohibitively slow for sequences beyond a few thousand tokens. Flash Attention v3 takes the memory efficiency further by overlapping the SRAM-to-HBM movement with the computation itself, using warp specialization and asynchronous DMA to keep the compute units busy while KV cache data is being loaded from HBM.
The practical question for inference operators: what does FA3 actually change about your bandwidth utilization and when does it matter? The answer depends more on your hardware generation and sequence length distribution than most documentation makes clear.
What FA3's warp specialization changes
FA2 uses a standard blocked attention pattern: load a tile of queries, compute attention scores against all key tiles, reduce into the output. During the key tile loading phase, compute units wait for HBM reads to complete. On H100 with HBM3 (3.35 TB/s bandwidth), this waiting is less severe than on A100 with HBM2e (2.0 TB/s). But for long sequences where many key tiles must be loaded, the memory latency accumulates.
FA3 restructures the warp assignments to overlap these loads with useful computation. Producer warps handle the asynchronous memory loads from HBM into SRAM. Consumer warps handle the matrix multiplications using data already in SRAM. When the pipeline is balanced, the compute units are executing GEMM operations continuously while the producer warps fetch the next tile. Stall cycles decrease because the hardware's prefetcher is pulling ahead of the consumer warps' demand.
This only works when the pipeline can stay ahead: if your KV cache tiles are fragmented in VRAM (as happens with paged attention when pages for a single sequence are non-contiguous in memory), the async prefetch cannot predict the next address to load. The memory access pattern becomes irregular, and FA3's overlap advantage partially collapses back toward FA2 behavior.
H100 benefits more than A100
FA3 was designed around H100's warp-specialized tensor core architecture. H100's Hopper architecture has dedicated warp groups for tensor core operations and separate warp groups for memory movement, managed through TMA (Tensor Memory Accelerator) instructions. FA3 maps cleanly onto this by assigning producer warps to TMA loads and consumer warps to tensor cores.
On A100 (Ampere architecture), FA3 still provides improvements over FA2, but the improvement is smaller because A100's memory movement and compute pipelines share resources differently. A100 does not have TMA instructions. The async copy operations available on A100 (cp.async) provide some overlap capability but are less capable than TMA.
In our testing, moving from FA2 to FA3 on H100 with 4K-token sequences produced roughly 18-24% improvement in attention kernel throughput. On A100 with the same sequence length, the improvement was 8-11%. On H100 with 32K-token sequences, the improvement was larger: 28-35%, because longer sequences have more KV tiles to load and the overlap benefit compounds. On A100 at 32K tokens, the improvement was 14-18%.
If your current fleet is A100-only, FA3 is worth enabling but the gains are less dramatic. Plan for larger improvements when you transition to H100 nodes.
Verifying FA3 is actually active in your setup
This is where the documentation gaps are most painful. Installing the flash-attn package from PyPI and calling flash_attn_func does not guarantee FA3 is being used. The dispatch depends on your CUDA version, your GPU architecture, your PyTorch version, and whether the FA3-specific CUDA extensions compiled correctly during package installation.
The reliable verification method: profile a long-context attention computation with Nsight Systems and look for flash_fwd_block_sm90 in the kernel names (for H100, SM 9.0) versus flash_fwd_block_sm80 (for A100, SM 8.0). The FA3 kernel names for H100 will include references to TMA and async copy pipeline stages. If you are seeing the SM80 kernels on an H100 node, FA3 is not being used despite being installed.
A simpler but less definitive check: measure HBM bandwidth utilization during a long-sequence attention computation. With FA2 on H100, attention at 8K sequence length typically runs at 60-70% of theoretical HBM bandwidth. With FA3, the async prefetch keeps the bandwidth utilization higher and more consistent, typically 75-85% for the same sequence length. Bandwidth utilization that looks FA2-like on H100 hardware suggests FA3 is not active.
KV cache layout and FA3 compatibility
Paged attention manages KV cache in fixed-size blocks that are allocated non-contiguously. A sequence of 8192 tokens with a block size of 16 tokens requires 512 blocks, and those blocks may be scattered across GPU memory due to allocations and frees from other sequences. FA3's async prefetch works best with contiguous memory access patterns.
The interaction between paged attention and FA3 is manageable but requires attention to block size selection. Larger block sizes reduce fragmentation (fewer, larger contiguous blocks per sequence) at the cost of VRAM granularity (a partly-used block wastes its empty pages). For FA3 on H100, we find that block sizes of 64-128 tokens produce better attention kernel performance than the 16-token default, with a modest increase in average VRAM waste per active sequence.
The KV cache access pattern during decode (adding one token per step to each sequence's cached context) is somewhat irregular by nature. FA3's prefetch cannot predict which token position's KV entries will be needed next at the cache access level, because this depends on the attention mask which is computed during the forward pass. What FA3 can prefetch efficiently is the full KV cache for a given sequence in order, which is what the attention computation requires.
Bandwidth numbers to expect
Calibrated expectations help you determine whether your FA3 deployment is working correctly. These figures are from our internal profiling on 8xH100 80GB SXM5 with NVLink, Llama-3 70B BF16, paged attention with 64-token block size.
At 4K sequence length, decode phase: HBM bandwidth 2.8-3.1 TB/s (out of 3.35 TB/s theoretical), attention kernel utilization 82-88%.
At 16K sequence length, decode phase: HBM bandwidth 2.9-3.2 TB/s, attention kernel utilization 86-91%. Longer sequences allow the prefetch pipeline to run more steadily.
At 32K sequence length, decode phase: HBM bandwidth 3.0-3.25 TB/s, attention kernel utilization 88-93%. The asymptote is from non-attention operations (FFN layers) that have different memory access patterns.
If you are seeing sustained HBM bandwidth below 2.5 TB/s on H100 during long-sequence decode, either FA3 is not active or there is significant VRAM access fragmentation. Both are worth diagnosing before adding more hardware.
What FA3 does not change
FA3 improves the attention kernel. It does not reduce the total amount of KV cache data that must be stored and accessed. If you are running at the limit of your KV cache capacity and seeing frequent evictions, FA3 will not help: the data is being evicted not because the attention kernel is slow, but because there are more active sequences than available KV cache pages. That is a memory capacity problem, not a bandwidth problem.
FA3 also does not improve prefill performance to the same degree as decode. Prefill attention is already compute-bound on large batch sizes (many input tokens, large matrix multiplications), and the memory movement is a smaller fraction of total time. The FA3 overlap benefit is most pronounced in the bandwidth-bound regime that characterizes decode.
Running long-context workloads on H100?
Inferact integrates FA3 with paged attention and provides per-operation bandwidth metrics so you can verify the attention kernels are working at expected efficiency. We work directly with early access partners on hardware-specific tuning.
Request Early Access