Warp-Specialized FP8 FlashAttention-3 Kernels Achieve 1.32-PetaFLOPs Sustained Attention Throughput on 256K-Token Contexts

CUDA kernel architects and deep-learning systems researchers across Santa Clara, San Jose, and Stanford University released production open-source FP8 FlashAttention-3 kernel updates on Sunday morning through the Dao-AILab FlashAttention Repository, demonstrating 1.32 PetaFLOPs per GPU of sustained forward-pass attention throughput (74 percent of theoretical FP8 Tensor Core peak) and a 2.4x speedup over BF16 FlashAttention-2 on 256-thousand-token sequences. In synchronous attention kernels, streaming multiprocessors (SMs) stall for up to 40 percent of execution cycles waiting on shared-memory barrier synchronizations and Special Function Unit (SFU) exponential softmax evaluations between consecutive QK^T and PV matrix multiplications.

The new warp-specialized Ping-Pong Producer-Consumer schedule partitions each threadblock into dedicated Tensor Memory Accelerator (TMA) producer warps and asynchronous Warpgroup Matrix-Multiply-Accumulate (wgmma.mma_async) consumer warps coordinated through circular hardware transaction barriers in shared memory. While consumer warpgroup A executes the register-resident softmax rescaling and exponential polynomial approximation on tile i, consumer warpgroup B simultaneously issues wgmma FP8 E4M3 tensor-core instructions for tile i+1, keeping the matrix pipelines saturated above 88 percent utilization.

To eliminate quantization accuracy loss when casting attention probabilities and Value (V) matrices to 8-bit E4M3, the kernel applies per-head block-wise incoherent Hadamard rotation inside the register file. An in-shared-memory hardware byte-permute (prmt.b32) transpose then prepares column-major V tiles for FP8 wgmma consumption at zero global-memory bandwidth overhead.

Benchmarked across a 64-GPU H200 SXM cluster at 2800 Scott Boulevard in Santa Clara funded by a $390,000 open-kernel research program, 256K-token context prefill latency dropped from 3.8 seconds to 1.54 seconds while matching BF16 perplexity within 0.03 percent. CUTLASS 3.6 CuTe templates and PyTorch bindings are available in the Dao-AILab FlashAttention GitHub Repository.