FP8 GEMM with FP32-accumulation accuracy at long K without a throughput penalty
Difficulty 34: getting promoted-FP32 accuracy at fast-accumulate speed requires a new way to hide promotion cost against a 13-bit hardware accumulator, harder than plain parity (24) and close to the parked accumulator-width bound problem (35.5). Application value 47: weight-gradient GEMMs with K equal to the token count decide FP8 training stability. Leverage 26: informs FP8 recipes and kernels.
The problem
Can an H100 FP8 matrix multiply add very long sums as accurately as 32-bit, at full fast-mode speed?
- 1Long sums pile up rounding errors
- 2H100 FP8 adds with 13-bit precision
- 3Moving sums to FP32 costs time
- 4Fast mode skips it, loses accuracy
- 5Want both: accurate and fast
What is asked. Each entry of a matrix product is a long sum: . In neural-network training, can be tens of thousands, for example when the gradient for a weight matrix is summed over every token in a batch. On an H100, FP8 matrix multiplies are fast, but the hardware adds the products with limited precision. The problem asks for an FP8 matrix multiply, for up to 32768, whose error is at most twice that of an ideal version that adds in full 32-bit precision. It must also run within 5% of the speed of cuBLAS's mode, which skips the usual accuracy fix.
Path to understanding
foundations → problemThe research entries this problem rests on, in the order to read them. Each links to its explainer.
- Start with
- Step 2
- Step 3
- Then
- FP8 GEMM with FP32-accumulation accuracy at long K without a throughput penaltyThe problem
State of the art
Khattak and Mikaitis (arXiv:2512.07004) model the H100 FP8 tensor-core path (wgmma) as adding groups of 32 products with only 13 fractional bits, truncating aligned significands, whereas B200 keeps 25 bits; a PyTorch engineering post reached a similar conclusion from measurements. The DeepSeek-V3 report works around this by promoting partial sums to FP32 on CUDA cores at fixed intervals, and cuBLAS does the same by default, periodically promoting unless fast accumulation is switched on. Promotion costs issue slots and registers, and its interval trades accuracy against throughput. No primary source publishes the error-versus-throughput frontier over K, or a kernel that gets promoted-level accuracy at fast-accumulate speed.
Realistic applications
Long-K reductions appear in the backward weight-gradient GEMM (K equal to tokens) and in large-hidden-size layers; accumulation error there decides whether FP8 training is stable.
What counts as solved
A public kernel and a published error and throughput table over K at a pinned commit, reproduced by an independent party.
How it gets verified
GPU benchmarkA kernel run and timed on stated hardware against a baseline.
Hardware: one H100 SXM 80 GB. Dtypes: E4M3 A and B with block scaling, BF16 output. Shapes: M=N in {1024, 4096}, K in {1024, 4096, 8192, 16384, 32768}. Inputs: standard normal and a heavy-tailed mixture (1% of entries scaled by 100) with fixed seeds. Reference: FP64 GEMM of the dequantised inputs. Baselines: cuBLAS FP8 GEMM (CUDA 12.9 or later) with fast accumulation on, for throughput (default-accumulation cuBLAS also reported, for error and speed); an emulated exact-FP32-accumulation FP8 GEMM for the error yardstick. Metric: max and RMS relative error, and median TFLOP/s over at least 100 iterations. Pass: error at most 2x the emulated-FP32 yardstick and throughput at least 95% of fast-accumulate cuBLAS for every listed shape. Harness: public script at a pinned commit.
Claims under verification
+ add a claimNo claimed results on the Verification board for this problem yet.
Where the work is
Mostly kernel engineering to hide promotion cost, plus error analysis against a now-documented accumulator model and measurement of error versus K.
Bottlenecks
ranked by votes- #1EngineeringHiding the cost of promoting partial sums
Moving partial sums from the tensor cores to FP32 registers on CUDA cores consumes issue slots and registers; cuBLAS fast-accumulate mode is faster precisely because it skips this.
Community grades
Sign in to add your own grade.