FlashAttention 2 vs FlashAttention 3: H100 and H200 Speedups, FP8 Support, and Migration Guide (2026)

Most H100 deployments today run FlashAttention 2. FA3 shipped in 2024 and adds warp-specialized scheduling and FP8 attention support for Hopper, but the actual throughput gain depends heavily on sequence length and whether your stack is current enough to use it. This post covers what changes architecturally, what the throughput numbers look like at 2K-128K sequence lengths on H100 SXM5 and H200 SXM5, and whether migrating from FA2 to FA3 is worth the effort for your workload.

For the next generation beyond FA3, see FlashAttention-4 on Blackwell.

FlashAttention Version Timeline

A quick reference before getting into the architecture details:

FA1 established the core insight: recomputing attention on-chip is faster than reading a materialized attention matrix from HBM. FA2 improved the parallelism strategy so more of the GPU's threads were doing useful work. FA3 added Hopper-specific features that FA2's kernel design couldn't use. FA4 targets Blackwell SM100 and is covered in the FlashAttention-4 guide.

Modern frameworks auto-detect the GPU architecture and select the appropriate version. vLLM v0.6+ and SGLang v0.4+ pick FA3 on H100/H200, FA4 on B200/B300, and FA2 on A100/L40S without any configuration needed. You only need explicit flags when overriding for debugging or benchmarking.

How FlashAttention 2 Works

FA2 made four improvements over the original FlashAttention:

Block-wise softmax tiling. FA2 keeps query, key, and value tiles in SRAM during computation. Rather than loading the full attention matrix from HBM, it computes attention in chunks that fit in the GPU's on-chip memory. This cuts HBM reads by the ratio of sequence length to tile size.

Online softmax. FA2 computes softmax using a running maximum and denominator, which lets it process attention blocks sequentially without materializing the N×N attention matrix. This is what makes tiling possible without numerical instability.

Warp-level parallelism across sequence positions. FA2 assigns different warps to different sequence positions, improving parallelism over FA1's thread-block-level splitting.

GQA and MQA support. FA2 added kernels for grouped-query attention and multi-query attention, used by Llama, Mistral, and most other models released since 2023.

FA2 works on Ampere (A100) and Ada Lovelace (L40S, RTX 4090) through the same kernel path. On Hopper, FA2 runs but does not use any of Hopper's async pipeline hardware.

What FlashAttention 3 Adds for Hopper

FA3 is not a general improvement over FA2. It is a Hopper-specific rewrite that uses three SM90 features FA2 cannot access:

Warp specialization. FA3 splits warps into producers and consumers. Producer warps fetch Q/K/V tiles asynchronously using Hopper's wgmma and cp.async instructions. Consumer warps execute the GEMM and softmax operations concurrently on tiles already in SRAM. In FA2, every warp does both data fetching and computation, which creates a sequential dependency: compute stops and waits for data. FA3 eliminates this stall by overlapping the two.

Ping-pong scheduling. FA3 maintains two SRAM tile buffers. While consumer warps process tiles from buffer A, producer warps load the next tiles into buffer B. When buffer A computation finishes, the roles swap. This keeps the tensor cores fed continuously without pipeline bubbles.

FP8 attention computation. FA3 can execute the Q@K^T and softmax(.)@V matrix multiplications in FP8 instead of BF16. The output is upcast to BF16 before returning. FP8 attention requires Transformer Engine's per-tensor scaling headers and flash-attn 2.7+.

Here is a direct comparison of the two architectures:

H100 SXM5: FA2 vs FA3 Throughput Benchmarks

The table below shows attention kernel throughput in TFLOPS at different sequence lengths on H100 SXM5. Figures are directional, based on the FA3 research paper (Shah et al., 2024) and community benchmarks. Verify against your specific model architecture before making capacity decisions.

The pattern is consistent with FA3's architecture: the gain is modest at short contexts (compute-bound, where better memory pipelining matters less) and grows substantially at long contexts (memory-bandwidth-bound, where eliminating pipeline stalls reduces idle time). The gap widens at longer contexts and shrinks at shorter ones, the same pattern FA4 shows over FA3.

H200 SXM5: FA2 vs FA3 Throughput Benchmarks

H200 SXM5 has the same compute capacity as H100 SXM5 in FP16 TFLOPS, but 4.8 TB/s HBM3e bandwidth versus H100's 3.35 TB/s. That 43% bandwidth increase matters most at long context lengths where attention is memory-bound.

Figures are directional based on published benchmarks. Verify against your workload before using these numbers in architecture decisions.

Two things stand out from this table. H200's absolute throughput at 8K+ context is substantially higher than H100's, reflecting the bandwidth advantage. FA3's relative gain over FA2 on H200 is slightly lower than on H100 at 128K (2.35x vs 2.44x). This is expected: H200's higher baseline bandwidth already compensates for some of the memory latency that FA3 hides through async pipelining. FA3 is still the right choice on H200, but H200 leaves less performance on the table with FA2 than H100 does.

FP8 Attention on FA3

FA3 exposes a second compute path where the Q@K^T and softmax(.)@V matrix multiplications run in FP8 instead of BF16. This is different from what people usually mean by "FP8 inference," which refers to weight quantization at the model level. FP8 attention is about the precision of the attention operation itself.

The precision split in FA3 FP8 attention:

• Q, K, V projections: BF16 (set at the model/vLLM level, not by FA3)
• Q@K^T matmul: FP8
• Softmax: FP8 numerics with scaling
• Softmax(.)@V matmul: FP8
• Output: upcast to BF16 before returning to the model

FA3 FP8 attention requires NVIDIA Transformer Engine's per-tensor scaling infrastructure to handle the dynamic range calibration that keeps FP8 numerics stable.

Throughput gain. FA3 FP8 attention adds roughly 15-25% throughput on top of FA3 BF16 at 8K+ context. The gain is additive with FP8 weight quantization, since they affect different parts of the computation.

Accuracy. At standard context lengths (2K-32K), the quality gap between FA3 FP8 attention and FA3 BF16 is within measurement noise on common benchmarks. At very long context (128K+), numerical precision matters more and FP8 attention can degrade output quality. Always validate on your target task before enabling FP8 attention in production.

For a full breakdown of FP8 formats (E4M3 vs E5M2) and how they interact with weight quantization, see FP8 quantization explained.

vLLM and SGLang: End-to-End Throughput Impact

Attention kernel throughput does not translate 1:1 to end-to-end serving throughput, since prefill and decode both involve non-attention operations like linear layers, normalization, and sampling. But FA3 does produce measurable improvements in serving metrics, especially at larger batch sizes and longer input lengths.

Directional figures for Llama 3.3 70B, BF16, H100 SXM5, 100 concurrent requests:

Time to first token (TTFT):

• At 2K input: FA3 and FA2 are roughly equivalent (less than 5% difference). Prefill at short context is compute-bound by the linear layers.
• At 8K input: FA3 reduces TTFT by 15-20%. The prefill phase becomes attention-heavy at longer inputs.

Inter-token latency (ITL):

• At batch size 8: The difference is small (less than 5%). Memory bandwidth at small batch sizes is not attention-dominated.
• At batch size 32+: FA3 reduces ITL by 10-15%. Larger batch sizes increase the attention memory footprint, where FA3's async pipelining helps.

SGLang shows a similar pattern. SGLang's RadixAttention cache layer works with both FA2 and FA3 without modification.

These are directional figures based on published community benchmarks. Your numbers will differ based on model, hardware SKU, and serving configuration. For a full framework comparison across vLLM, TensorRT-LLM, and SGLang, see vLLM vs TensorRT-LLM vs SGLang benchmarks.

When FA2 Is Still the Right Call

FA3 is better than FA2 on H100 and H200. That said, there are cases where staying on FA2 is the practical choice:

Older CUDA or PyTorch stacks. FA3 requires CUDA 12.3+ and PyTorch 2.2+. Running PyTorch 2.0 or CUDA 12.2 means enabling FA3 requires a full stack upgrade. The upgrade is worth doing for a new deployment, but not necessarily mid-flight for a stable production service.

Non-Hopper hardware. FA3 is SM90-only. A100, L40S, and RTX 4090 all run FA2 regardless of what flash-attn version you have installed. If you're deploying across a mixed GPU fleet, FA3 activates only on the Hopper nodes; Ampere/Ada nodes run FA2 automatically. You don't need separate container images.

Non-power-of-2 head dimensions. FA3 as of flash-attn 2.7 has coverage gaps for non-standard head dimensions. If your model uses head dimensions that aren't powers of 2 (such as 96 or 192), test FA3 compatibility before deploying.

Custom masking patterns. FA3 did not ship all of FA2's custom mask options at launch. If you rely on non-causal masking modes or sliding window attention with unusual patterns, test the specific mask configuration before switching.

For most Hopper deployments running standard Llama, Mistral, or Qwen model families with 128-dim heads, FA3 works without issues and the upgrade path is straightforward.

Migrating from FA2 to FA3

Install

# FA3 support requires flash-attn >= 2.7.0 and CUDA 12.3+
pip install "flash-attn>=2.7.0" --no-build-isolation

# Verify FA3 is available
python -c "import flash_attn
print(flash_attn.__version__)
from flash_attn.flash_attn_interface import flash_attn_func_v3
print('FA3 available')
"

vLLM

# Default: auto-selects FA3 on Hopper (vLLM 0.6+)
docker run --gpus all vllm/vllm-openai:latest --model <model> --dtype bfloat16

# Explicit FA3
vllm serve <model> --attention-backend flash_attn_v3

# Force FA2 for debugging or comparison
vllm serve <model> --attention-backend flash_attn

SGLang

# SGLang v0.4+ auto-selects FA3 on H100/H200
python -m sglang.launch_server --model-path <model> --tp 1 --dtype bfloat16

# Explicit FA3
python -m sglang.launch_server --model-path <model> --attention-backend flashattn3

Common pitfalls

• flash_attn version below 2.7 silently falls back to FA2 on Hopper without error or warning.
• FA3 requires --dtype bfloat16 or --dtype float16. Running --dtype float32 disables FA3.
• If you see CUDA error: invalid device function, your CUDA version is below 12.3.
• FA3 is the default in vLLM 0.6+ on Hopper. If your logs don't show the FA3 backend, your flash-attn version is likely the issue.

For a full production deployment guide covering tensor parallelism, memory optimization, and serving configuration alongside FA3, see the vLLM production deployment guide.

Cost Impact: Tokens/sec/$ on Spheron H100 and H200

The attention kernel choice affects throughput directly, which changes cost per token. Using live pricing from Spheron as of 25 May 2026 (H100 SXM5: $3.84/hr, H200 SXM5: $4.56/hr on-demand):

Formula: Cost per 1M tokens = ($/hr) / (tokens/sec × 3,600) × 1,000,000

Tokens/sec figures are directional estimates for Llama 3.3 70B at 8K context, batch 32. Actual throughput depends on model architecture, batch size, and input length distribution.

Switching from FA2 to FA3 BF16 on H100 SXM5 drops cost from $0.089 to $0.059 per million tokens with the same hardware at the same hourly rate. FA3 + FP8 attention cuts it further to $0.051. On H200 SXM5, FA3 BF16 at $0.055/1M tokens is competitive with H100 FA3 + FP8 attention ($0.051), with H200's 141 GB VRAM headroom as the additional advantage for larger models or longer KV caches.

You can benchmark these configurations on H100 GPU rental or try H200 SXM5 for long-context inference on Spheron, both available on-demand with no commitment. Running both back-to-back on identical hardware takes under an hour and gives you real numbers for your specific workload.

Pricing fluctuates based on GPU availability. The prices above are based on 25 May 2026 and may have changed. Check current GPU pricing → for live rates.

FA3 is active by default on H100 and H200 in vLLM 0.6+ and SGLang 0.4+. If you want to benchmark FA2 vs FA3 side-by-side on identical hardware, H100 SXM5 and H200 SXM5 are both available on Spheron at listed per-hour rates with no commitment.

Rent H100 → | Rent H200 → | View all GPU pricing →