FlashAttention-4's Direct-P Finally Unlocks 2.13x FP4 Speedup on NVIDIA GB200

Blackwell's shiny new 4-bit floating point tensor cores were supposed to make attention screamingly fast. In practice, they didn't. A new paper, Hardware-Aware FP4 FlashAttention-4, explains why the naive port stalls and introduces two kernels that finally cash in the FP4 speedup for both inference and training.

The headline result: a technique called Direct-P delivers up to 2.13x the BF16 forward throughput on an NVIDIA GB200 for noncausal attention. For causal training, a matched forward/backward path speeds up a complete single-GPU 8-billion-parameter update by up to 1.14x using FP8 gradient operands.

Why FP4 alone leaves throughput on the table

To understand the problem, you need to know how attention runs on Blackwell. FlashAttention keeps the score and probability matrices out of HBM by walking through keys in tiles, running a small softmax on each tile, and correcting the running output as it goes. On a B200, asynchronous tcgen05 operations accumulate 32-bit score and output tiles in tensor memory (TMEM), the Tensor Memory Accelerator supplies shared-memory operands, and separate warps issue matrix operations, compute softmax, and correct the online output.

Here is the catch. Blackwell's tensor cores got much faster, but the surrounding units did not scale at the same rate. Analysis of the base FlashAttention-4 kernel found that although Blackwell B200 doubles tensor core throughput compared to Hopper H100 (2.25 vs 1 PFLOPS for FP16/BF16), other functional units scale more slowly or remain unchanged, so non-MMA resources emerge as bottlenecks, with shared memory traffic and exponential operations exceeding MMA compute by 25 to 60%. Switching QK from BF16 to FP4 shrinks the matmul portion further, so softmax conversion and on-chip dependencies dominate the remaining runtime. You get the peak FLOPs, but nothing to feed them.