Hardware-Aware FP4 FlashAttention-4
cs.LG
Quick verdict
Reports 2.13x BF16 FP4 attention speed on GB200 and 1.14x 8B training; well-bounded results, but notable accuracy cost and single-run training evidence.
Summary
This Graphcore technical report targets the gap between Blackwell's fast FP4 matrix cores and the serial softmax work inside tiled attention. For noncausal inference it proposes Direct-P, which maps log-scores directly to MXFP4 (E2M1) probability codes through a fitted affine map and normalizes the output using the same rounded probabilities the value product consumes, avoiding a separate exponential-denominator pass. On a GB200, across nine D128 shapes it achieves a 2.023x geometric speedup over BF16 FlashAttention-4 (peak 2998 TFLOP/s; 2.13x at S8192/H64), at output cosine ~0.944 versus 0.990 for an FP8-P/V control. Fixed-input ViT/BERT/SST-2, ViT-MAE and Wan video evaluations remain finite with modest task-score changes. For causal training, the forward pass ships quantized Q/K, scales and LSE to backward, which reconstructs probabilities and uses FP8 gradient operands; the complete 8B update speeds up up to 1.14x at batch 4, but the retained route keeps FP8 P/V because every tested MXFP4-P/V distributed trajectory diverges. A matched 64-GPU 100B-token run shows 1.112x median throughput with a +0.090 validation-loss gap. Main limitations: sizeable operator error (relative-L2 ~0.33), one trajectory per arm with no variance estimate, shape-dependent speedups, and B300 records not from a same-binary hardware A/B.
Review complete
Assessment at a glance
4
risk flags identified
● high 0
● medium 3
● low 1
5 strengths · 5 weaknesses · 5 author questions
AI text signal = Binoculars zero-shot detector score over 20 sampled passages (Fast-DetectGPT co-reported) — absence does not prove human authorship, and a flag does not imply misconduct.
Levels reflect verification complexity — not a quality score. Evolve never rates your paper.
Key claims under review
- Direct-P maps log-scores directly to MXFP4 probability codes and reaches up to 2.13x BF16 FlashAttention-4 forward throughput on GB200 (2.023x geometric mean across nine D128 shapes, peak 2998 TFLOP/s).
- Passing the forward quantization (NVFP4 Q/K payload, scales, LSE) into causal backward accelerates projection-inclusive attention by 1.245x and a complete single-GPU 8B update by up to 1.14x at batch 4.
- Every tested MXFP4-P/V training trajectory diverges while FP8-P/V arms remain stable, so the retained training route keeps FP8 P/V with a +0.090 same-update validation gap over a completed 100B-token matched run at 1.112x median throughput.
- At D128, tensor-memory ownership (two 128-column score banks plus two 128-column output banks consume all 512 TMEM columns) and operand readiness - not arithmetic throughput - limit attention overlap on Blackwell.
Strengths
Measurement boundaries are pre-declared and never violated
Table 6 separates kernel / projection-inclusive / single-GPU 8B update / distributed boundaries; Section 7 reports isolated backward (1.405x core), projection-inclusive (1.245x), and full-update (1.14x) results strictly within their own boundary and explicitly refuses to transfer timings between recipe generations (Section 5.3).
Negative results are reported with unusual candor
The isolated-backward saving is consumed by the E5M2 dO publisher (Table 12: 1.405x core falls to 0.986x); every MXFP4-P/V training trajectory diverges (Section 7.6, Appendix G.8); the fixed-P speed-of-light ceiling shows only 5.23% headroom remains (Table 17); Appendix B catalogs ~15 rejected kernel directions with failure mechanisms.
Factorial diagnostic isolates the divergence cause
Reproducibility infrastructure exceeds field norms
Accuracy cost is quantified against a stronger-accuracy control, not hidden
Weaknesses
The strongest baseline comparison is cross-run, and the local FP8 control is weak
Operator error is large and downstream coverage thin
Single trajectory per arm in the 100B-token matched run
B300 rows come from a tuning campaign, not a controlled A/B
MXFP4-P/V divergence is observed but not mechanistically explained
Full Review Report
$1.99 · one-timeUnlock the complete preliminary analysis Evolve produced for this paper.
Report preview
Scope & Claims
This single-author Graphcore technical report (arXiv 2609.04105, Sep 2026) attacks a real bottleneck: on Blackwell, FP4 matrix products (QK, PV) accelerate but the serial softmax between them does not. It contributes (1) **Direct-P**, a noncausal-inference path that maps log-scores directly to MXFP4 (E2M1) probability codes via a fitted affine map and normalizes using the same rounded probabilities the PV product consumes; (2) a **causal-training path** that ships the forward's quantized Q/K payload, block/global scales and LSE into backward; and (3) a hardware analysis attributing residual latency to tensor-memory (TMEM) ownership. Headline claims: up to 2.13x BF16 forward throughput on GB200; 1.245x projection-inclusive attention and up to 1.14x complete 8B update; all tested MXFP4-P/V training trajectories diverge, so the retained route uses FP8 P/V. Claim discipline is unusually good - the abstract hedges with 'up to' and 'favorable shapes' consistently with Section 6.
The remaining sections cover Methodology Assessment, Statistical Evidence, Results & Reproducibility, Limitations & Risks, and Verdict & Recommendations.
- All 5 strengths & 5 weaknesses, analyzed in detail
- 4 risk flags with severity, section & reviewer questions
- Methods & statistical evidence assessment
- 5 questions a peer reviewer would ask
- Full six-dimension evaluation report & recommendations
One-time purchase · this paper only
One-time purchase, this paper only. AI-generated preliminary analysis — not peer review.
Frequently asked questions
How much faster is FP4 FlashAttention than BF16?
On an NVIDIA GB200, the paper's Direct-P full-FP4 forward attention reaches a 2.023x geometric-mean speedup over BF16 FlashAttention-4 across nine D128 shapes, peaking at 2998 TFLOP/s and up to 2.13x at sequence 8192 with 64 heads; B300 records reach 3116-3159 TFLOP/s on large-sequence shapes. End-to-end, a complete 8B-model training step speeds up up to 1.14x on one GPU.
Can attention run in FP4 during training, or only inference?
Only inference runs fully FP4. For training, every tested MXFP4 probability/value (P/V) trajectory diverged in distributed experiments, so the retained training route keeps FP8 P/V with FP4 (NVFP4) queries/keys and quantized-backward reconstruction; it stayed stable and descending over a completed 100-billion-token matched run, with a +0.090 validation-loss gap versus BF16 and 1.112x median throughput on 64 GPUs.
How accurate is the FP4 attention output?
On the paper's Gaussian operator tests the fast FP4 route reaches cosine ~0.944 and relative-L2 ~0.336 versus BF16 - noticeably less accurate than an FP8 P/V control (cosine 0.990). Downstream, ViT/BERT/SST-2 task scores changed little (e.g., ViT S4096 top-1 matched BF16 at 88.5%), and Wan video-diffusion latents stayed finite but drifted over 20 steps (cosine 0.85-0.91).
Is there code to reproduce this paper?
Yes. The implementation and reproduction materials are at github.com/MrHuff/fp4-fa4, including ThunderKittens forward/backward kernels, experiment configs, SHA256-hashed evidence receipts, and a planner that regenerates measurement commands. Actual re-measurement requires Blackwell-class GPUs (GB200/B300).
About this review
This is an AI-generated preliminary review by Evolve.Pub, produced from the openly available text of the paper. It is not peer review and does not represent the authors' views. The canonical record of the paper is at the publisher link above.
- The provided text is a sampled excerpt with [...]sampling gap[...] markers: parts of Section 4.2, the end of Section 6.4, portions of Sections 8.3-8.5, and parts of Appendices C, D and G were skipped by sampling; judgments here rely only on visible content and do not assert absence of anything behind the gaps.
- Figures are represented only by extracted captions and axis-label text; visual plot content could not be inspected.
Go deeper on this paper
Wrote a paper? Stress-test it before reviewers do.
Upload your manuscript for the same structured scrutiny — reasoning, statistics, evidence, figures, and reproducibility.