Hardware-Aware FP4 FlashAttention-4
Abstract: Blackwell's 4-bit floating-point (FP4) tensor cores do not automatically make attention faster because softmax conversion and on-chip dependencies dominate once its matrix products shrink. We address this with \emph{Direct-P} for noncausal inference and a causal path that passes the forward quantization directly into backward. Direct-P maps scores directly to FP4 probabilities and reaches up to 2.13 the bfloat16 (BF16) forward throughput on an NVIDIA GB200. The causal path reconstructs probabilities from saved quantized queries and keys and uses 8-bit floating-point (FP8) gradient operands, accelerating a complete single-GPU 8-billion-parameter update by up to 1.14. Matched distributed training retains FP8 probabilities and values; every tested MXFP4 probability/value training trajectory diverges.
Paper Prompts
Sign up for free to create and run prompts on this paper.
Top Community Prompts
Explain it Like I'm 14
1. What is this paper about?
This paper studies how to make transformer attention faster on newer NVIDIA Blackwell GPUs.
Attention is the part of a transformer that helps the model decide which words, image pieces, or video pieces are important to each other. It uses several large calculations. New Blackwell GPUs can do calculations using very small numbers called FP4, which use only 4 bits. Smaller numbers can make calculations much faster and use less memory.
However, simply changing everything to FP4 does not automatically make attention faster. Some other steps—especially softmax, which turns scores into probabilities—can become the slowest part.
The paper introduces a method called Direct-P. It changes how attention probabilities are created so that they can be sent directly to the GPU’s fast FP4 hardware.
The paper also studies whether low-precision numbers can be reused during the backward pass used for training.
2. What questions does the research ask?
The researchers focus on several main questions:
- Can all important attention inputs use FP4 during inference? In other words, can the queries, keys, probabilities, and values all use very small 4-bit numbers while still producing useful answers?
- Can FP4 make attention faster in practice? The researchers want to know whether the whole attention process becomes faster—not just one matrix multiplication.
- How much accuracy is lost when probabilities use FP4? FP4 has very few possible values, so it cannot represent numbers as precisely as larger formats such as BF16 or FP8.
- Can the low-precision information from the forward pass help training? During training, the model must perform a backward pass to calculate how its weights should change. The paper asks whether the backward pass can reuse the quantized information from the forward pass.
- Will models still train successfully with these low-precision calculations? A method can be fast but useless if the model’s training becomes unstable or produces bad results.
3. How did the researchers do the study?
Attention in simple terms
Attention can be thought of as a system for deciding how much each piece of information should “listen” to every other piece.
It uses three main inputs:
- Queries (
Q): what each item is looking for. - Keys (
K): labels describing what each item contains. - Values (
V): the information that will actually be passed along.
The calculation first compares queries and keys:
1 |
scores = Q × K |
A process called softmax changes these scores into probabilities. These probabilities say how much attention each item should give to other items:
1 |
probabilities = softmax(scores) |
Finally, the probabilities are used to combine the values:
1 |
output = probabilities × V |
The problem is that the GPU may perform the two matrix multiplications very quickly, but softmax and the movement of data between steps can slow everything down.
Tiling and FlashAttention
The researchers use a technique related to FlashAttention. Instead of handling the entire attention table at once, FlashAttention divides it into small pieces called tiles.
This is similar to solving a giant puzzle one small section at a time. The GPU keeps only the needed pieces nearby, rather than storing the entire puzzle in slower memory.
This saves memory and can improve speed. However, each tile still has to be processed in the correct order:
- Calculate a score tile.
- Find its largest score.
- Convert scores into probabilities.
- Quantize and pack the probabilities.
- Use them in the second matrix multiplication.
The researchers studied how to shorten this chain.
FP4, FP8, BF16, and quantization
The paper compares several number formats:
- BF16: a relatively precise 16-bit format.
- FP8: an 8-bit format with less precision.
- FP4: a very small 4-bit format.
Using fewer bits is like writing numbers with fewer digits. It saves space and can be faster, but it also creates more rounding errors.
The process of changing a precise number into a smaller format is called quantization. It is similar to rounding prices to the nearest dollar instead of keeping every cent.
The paper mainly uses:
- NVFP4 for queries and keys.
- MXFP4 for probabilities and values.
Because FP4 has so few possible values, the numbers are stored together with scales. A scale tells the hardware how large the stored FP4 values should be interpreted as being.
The Direct-P method
Normally, the GPU might calculate a fairly accurate exponential value during softmax and then round it to FP4. Direct-P skips much of that unnecessary precision.
Instead, it asks:
Which FP4 code should this score become?
This is like deciding directly whether a number should be rounded to 1, 2, 3, or 4, rather than first calculating many decimal places that will immediately be thrown away.
Direct-P also calculates the normalization value—the number used to make probabilities add up correctly—from the same rounded FP4 values that the GPU actually uses. This keeps the calculation internally consistent.
The method uses:
- A fast approximate rule for converting scores into FP4 codes.
- Occasional hardware exponential calculations when they are useful.
- Special safeguards for unusually large model scores.
- Careful reuse of on-chip memory so that scores and probabilities do not overwrite each other too early.
Experiments
The researchers tested the method on NVIDIA Blackwell GPUs, especially the GB200 and B300.
They measured:
- Attention kernel speed.
- Speed compared with BF16.
- Similarity to higher-precision results.
- Performance on vision and LLMs.
- Behavior on video-generation models.
- Backward-pass speed.
- Complete model-update speed.
- Training stability in distributed experiments.
They used measures such as:
- Cosine similarity: whether two outputs point in nearly the same direction.
- Relative error: how far the approximate answer is from the reference answer.
- Throughput: how much computation the GPU completes per second.
4. What were the main findings?
Direct-P made forward attention much faster
On the tested GB200 shapes, the fast Direct-P method was about 2.02 times faster on average than the BF16 comparison in the main D128 test set.
Its best reported result was about 2.13 times faster than BF16 for one of the tested shapes.
On the B300, the method reached about 3.1 PFLOP/s on some long-sequence examples. A PFLOP is an extremely large number of calculations per second.
These results show that FP4 can provide a real speed benefit when the entire probability path is designed for FP4—not merely when the matrix multiplication is changed to FP4.
The method trades some accuracy for speed
The fastest FP4 method was less accurate than the FP8 probability path.
The paper reports approximately:
- Fast Direct-P: average cosine similarity around 0.944 in a difficult standalone test.
- A more accurate version: around 0.952.
- The FP8 probability path: around 0.990.
This means Direct-P’s outputs can differ noticeably from higher-precision outputs, especially in artificial or demanding tests.
However, errors were often smaller in complete model tests. For example, later model layers, normalization, and residual connections sometimes reduced the effect of attention errors.
Many complete model tasks still worked well
In fixed-input tests involving ViT and BERT:
- Classification accuracy was often close to BF16.
- Some tests showed the same task accuracy as BF16.
- The fast method changed only a small number of predictions.
- Changed predictions were usually examples where the model was already uncertain.
This suggests that a noticeable numerical difference inside attention does not always lead to a different final answer.
However, the researchers are careful not to claim that this proves the method is safe for every model or every training task.
The backward pass could reuse forward quantization
For causal training, the paper passes quantized information from the forward pass directly into the backward pass.
This avoids rebuilding a separate higher-precision probability path. The backward calculation uses FP8 gradient operands in important places.
According to the abstract, this approach accelerated a complete single-GPU update of an 8-billion-parameter model by up to 1.14 times.
This is a smaller improvement than the forward-inference speedup, because full training includes many other operations besides attention.
FP8 was more stable than MXFP4 for distributed training
The researchers tested different formats during training.
They found that using FP8 for probabilities and values was more reliable. In contrast, every tested training run using MXFP4 probabilities and values diverged. Divergence means that training became unstable and the model stopped learning properly or produced unusable values.
This is an important result: the fastest format for inference is not necessarily suitable for training.
Hardware memory organization was a major limit
The paper found that the GPU’s on-chip tensor memory, called TMEM, was a major bottleneck.
The GPU needs separate areas to hold:
- Current score tiles.
- Other score tiles being processed.
- Output accumulators.
- Temporary probability data.
At the tested tile size, all available storage was already being used. This made it difficult to keep more work ready in advance.
The researchers conclude that future hardware could improve attention performance by providing another usable temporary storage area with suitable access rules.
5. Why are these results important?
The main lesson is that low-precision hardware alone is not enough.
A GPU may have extremely fast FP4 matrix units, but the whole attention algorithm can still be slow if:
- Softmax takes too long.
- Probabilities must be converted through unnecessary intermediate formats.
- Data has to wait for a memory location to become available.
- Synchronization prevents different parts of the GPU from working at the same time.
Direct-P tries to redesign the surrounding algorithm so that the fast FP4 hardware is actually used effectively.
The results suggest a practical division of labor:
- FP4 may be very useful for fast inference when a small loss of numerical precision is acceptable.
- FP8 may be a better choice for training, especially for probabilities and values.
- Larger formats may still be needed for sensitive parts of the model.
Simple conclusion
This paper presents a way to make transformer attention faster on new NVIDIA GPUs by sending probability data directly into FP4 hardware instead of carefully calculating values that will immediately be rounded away.
The method can produce roughly two times the forward-attention speed in favorable cases and usually keeps model-level results reasonably close to higher-precision versions. It also offers a smaller speedup for complete training updates.
However, the method is not universally better. Its fastest version introduces more numerical error, and MXFP4 training was unstable in the distributed experiments. Therefore, the most useful future systems may use a mixed-precision design: FP4 for speed-critical inference calculations, FP8 for more sensitive training calculations, and higher precision where accuracy matters most.
Knowledge Gaps
Knowledge gaps, limitations, and open questions
- Generality beyond NVIDIA Blackwell: The method is evaluated primarily on GB200 and B300/SM103 systems, so its performance, numerical behavior, and scheduling assumptions on other GPU generations, vendors, or future Blackwell variants remain unknown.
- Dependence on a specific hardware schedule: Direct-P inherits HAO’s two-query CTA layout, TMEM allocation, barrier protocol, and K64 PV issue pattern. It is unclear whether the method remains advantageous for alternative CTA organizations, head dimensions, tile sizes, or hardware with different TMEM capacities and issue semantics.
- Limited head-dimension coverage: The main results focus on D128, while D64 is treated separately and other dimensions are not evaluated. Performance and accuracy for D32, D equal to 96, D256, grouped-query attention, and architectures with heterogeneous head dimensions remain unresolved.
- Incomplete shape coverage: The benchmark grid does not systematically vary batch size, query length, key length, number of query heads, number of key/value heads, or causal versus noncausal masks. The method’s behavior for decode-time attention, variable-length batches, ragged sequences, and long-context settings beyond the tested cases is unknown.
- Limited comparison against equivalent FP8 baselines: The strongest FP8 comparison relies partly on published HAO results from different hardware and harnesses, while the local FP8 control is weaker. A fully matched, same-binary, same-hardware comparison across all formats is needed to establish the true speed–accuracy advantage of Direct-P.
- End-to-end latency is not established for inference: Timed kernels exclude dynamic Q/K/V quantization, scale generation, layout conversion, K/V permutation, and potentially required upstream fusion. The net latency and memory-bandwidth impact when these costs are included in a production transformer remain uncertain.
- Offline scale preparation may limit deployment flexibility: Direct-P assumes prequantized operands and offline-folded adjacent NVFP4 Q/K scales. The cost and feasibility of producing these layouts dynamically for changing inputs, autoregressive decoding, retrieval-augmented inputs, or continuously changing model weights are not quantified.
- Reliability of the sampled logit guard is unproven: The guard estimates a row maximum from 128 sampled or permuted key rows rather than performing an exact scan. Its failure probability across distributions, models, sequence lengths, adversarial inputs, and attention patterns has not been measured.
- No formal error bound for sampled anchoring: The paper does not provide a bound relating the sampled anchor to the true row maximum or quantify the resulting output error when the sampled maximum is substantially below the actual maximum.
- Global affine quantizer parameters may not generalize: The retained map is selected from Wan-related evaluations, but its robustness across architectures, layers, modalities, training checkpoints, logit distributions, and attention temperatures is not established.
- Calibration methodology is underdeveloped: The paper rejects layer-wise affine calibration based on isolated substitutions failing to predict composed trajectories, but it does not determine whether joint calibration, distribution-aware calibration, per-head calibration, or learned quantizer parameters could improve the speed–accuracy frontier.
- MXFP4 probability error remains substantial: The fast path reports approximately $0.944$ cosine similarity and relative- error around $0.34$ in the synthetic operator tests. The causes of this error, its dependence on attention entropy and sequence length, and whether alternative block sizes or stochastic/biased rounding can reduce it without sacrificing throughput remain open.
- Synthetic Gaussian diagnostics may not represent real attention distributions: The probability-range study uses exact Gaussian softmax probabilities. The extent to which its zero-payload, lost-mass, and output-error statistics predict real transformer attention distributions is not established.
- No systematic analysis by attention entropy or sparsity: The method’s quantization error likely depends strongly on peaked versus diffuse attention. Results are not stratified by entropy, effective support size, temperature, or number of dominant keys, leaving the favorable and unfavorable regimes unclear.
- Masking and causal boundaries are insufficiently characterized: The method is mainly evaluated for noncausal forward attention, while causal training uses a separate path. Behavior near causal diagonals, padding boundaries, block-sparse masks, sliding-window masks, prefix-LM masks, and arbitrary additive masks is not fully evaluated.
- Normalization changes the operator being computed: Direct-P normalizes the rounded FP4 probabilities rather than approximating the exact softmax denominator. The consequences for calibration, attention entropy, gradient magnitude, and model behavior across depth are not theoretically characterized.
- Backward-gradient fidelity is not comprehensively reported: The paper states that causal backward uses reconstructed probabilities and FP8 gradient operands, but broad gradient-error statistics, per-parameter error distributions, gradient covariance, and effects on optimization are not provided in the supplied text.
- The interaction between forward quantization and backward reconstruction remains unclear: It is unknown whether reconstructing probabilities from saved quantized Q/K and normalizers introduces biased or inconsistent gradients relative to the actual forward operator, especially when scale guards, blockwise normalization, or underflow correction are activated.
- Training claims are based on narrow model and duration coverage: The reported single-GPU 8-billion-parameter update and distributed trajectories do not establish long-horizon convergence, final task quality, robustness across optimizers, or behavior across substantially different architectures and datasets.
- Distributed training evidence is insufficiently broad: The paper reports that tested MXFP4 P/V trajectories diverge and retains FP8 P/V, but the number of nodes, communication topologies, global batch sizes, seeds, training durations, and divergence criteria are not sufficiently characterized to determine whether MXFP4 can be stabilized.
- Potential remedies for MXFP4 training are unexplored: Quantization-aware training, loss scaling, gradient clipping, delayed quantization, stochastic rounding, per-layer formats, selective FP8 fallback, and mixed-precision schedules are not systematically investigated.
- Projection precision remains a major confounder: Learned projections are kept at a separate precision boundary, and the reported throughput arm uses NVFP4 projections. Consequently, the contribution of attention quantization cannot be cleanly separated from projection quantization error and projection-kernel speedups.
- Full-transformer FP4 coverage is not demonstrated: The claim of “full FP4” applies only to Q, K, P, and V inside attention. Layer normalization, projections, residual paths, activations, embeddings, logits, and optimizer states remain at other precisions, so end-to-end FP4 feasibility remains unresolved.
- Model-quality evaluation is limited: ViT, BERT, and Wan fixed-input evaluations provide useful evidence but do not cover language-model perplexity, generative quality, instruction following, translation, summarization, speech, multimodal tasks, or long-context reasoning.
- Fixed-input evaluations do not establish robustness to distribution shift: The tests use predetermined examples, prompts, seeds, and masks. Sensitivity to unseen data, rare-token distributions, adversarial inputs, domain shifts, and changing generation temperatures is unknown.
- Diffusion evaluation is narrow: The Wan results use a small number of prompts, seeds, and diffusion step counts. Effects on perceptual quality, temporal consistency, human preference, failure rates, and other diffusion architectures remain unexplored.
- Accuracy metrics are not aligned consistently with downstream objectives: Operator cosine and relative- error can obscure changes in attention distributions, logits, calibration, or generated outputs. More task-specific metrics and statistically powered quality comparisons are needed.
- Statistical uncertainty is not fully reported: Many results appear to rely on individual deterministic records or short timing windows. Confidence intervals, run-to-run variance, thermal effects, clock variability, and statistical tests for both performance and accuracy are not systematically provided.
- Cross-hardware comparisons are not causal: B300 and GB200 results use different systems and, in some cases, different binaries or harnesses. The contributions of hardware, compiler version, tuning, clock settings, and implementation changes cannot be cleanly separated.
- Compiler and software-version sensitivity is unknown: The paper does not establish whether the reported instruction schedules and speedups persist across CUDA, PTX, compiler, driver, ThunderKittens, and FlashAttention revisions.
- Resource and occupancy trade-offs are incompletely quantified: TMEM ownership is identified as a bottleneck, but the effects of register pressure, shared-memory usage, occupancy, warp specialization, CTA residency, and contention with other kernels are not fully modeled.
- The proposed extra score destination is not experimentally validated: The paper motivates another allocatable score destination with compatible issue semantics, but does not demonstrate a hardware or simulated implementation showing the achievable gain.
- Energy efficiency and total cost are not measured: Throughput improvements are reported, but power consumption, energy per token, thermal behavior, and total deployment cost are not evaluated.
- Memory savings are not quantified end to end: FP4 operands may reduce storage and bandwidth, but the additional scale pages, packed layouts, temporary buffers, guards, and projection metadata could offset these benefits. A complete memory-footprint analysis is missing.
- Numerical corner cases remain incompletely tested: Subnormals, zero scales, extreme positive and negative logits, negative or additive masks, NaNs/Infs, very long rows, and pathological scale distributions require systematic stress testing beyond the reported layer-specific guard.
- The relationship between speed and accuracy operating points is not optimized globally: The fast and accurate policies use fixed choices for native
EX2, anchors, and denominator correction. A dynamic controller that selects policies by layer, shape, entropy, or runtime pressure could offer a better Pareto frontier but is not investigated. - Reproducibility is vulnerable to the incomplete scope of the supplied manuscript: The text references appendices, figures, generated tables, and a public implementation, but the provided content ends before the full results and evidence boundaries are presented. Independent verification of the complete claims therefore requires access to all omitted artifacts and appendices.
Practical Applications
Immediate Applications
- Blackwell GPU inference acceleration for long-context transformers (software, AI infrastructure; deployable now)
- long-context language-model encoding;
- bidirectional encoder inference;
- vision transformers with large patch sequences;
- video and multimodal transformer encoders;
- diffusion-model self-attention.
Dependencies: NVIDIA Blackwell hardware, compatible CUDA/PTX and tensor-memory instructions, D128-oriented kernels, NVFP4 Q/K and MXFP4 P/V operand layouts, and upstream quantization support.
- Production video-diffusion inference optimization (generative AI, media, cloud inference; deployable now with validation) Apply the guarded Direct-P policy to video diffusion models such as Wan-like architectures. The sampled-logit guard is specifically intended for layers with extreme logits and avoids a full second score scan. This could reduce latency and GPU cost for text-to-video generation, particularly at long spatial-temporal sequence lengths.
Dependencies: model-specific validation across prompts, seeds, guidance scales, and diffusion steps; correct routing of guarded layers; preservation of the required scale and key/value permutations.
- Quantized ViT and BERT inference services (computer vision, NLP, enterprise AI; deployable now) Replace eligible attention calls in classification, masked-language-modeling, and vision workloads with the fast or accurate NV/MX policy. The paper reports near-baseline task scores in fixed-input ViT and BERT evaluations, while attention speedups range from roughly 1.17× to 1.78×, depending on sequence length.
Potential products/workflows: low-latency document classification, semantic search encoders, image classification APIs, high-throughput batch embedding services, and edge or cloud inference endpoints.
Dependencies: fixed-input results do not establish universal safety; production deployment should include calibration, task-level regression tests, prediction-agreement monitoring, and fallback to FP8/BF16 for sensitive layers.
- Accuracy–latency selectable attention kernels (ML compiler/runtime engineering; deployable now)
fast: maximum throughput and lower numerical fidelity;accurate: additional exponential evaluation and guard logic for improved fidelity.
Frameworks such as PyTorch, TensorRT, JAX, and custom inference runtimes could select the policy based on sequence length, model layer, hardware, and application tolerance.
Dependencies: kernel autotuning must account for head count, sequence length, D64 versus D128 tile regimes, and whether attention is causal or noncausal.
- Quantization-aware model deployment pipelines (MLOps, compiler tooling, model optimization; deployable now)
- quantize Q/K to NVFP4;
- produce MXFP4-compatible P/V scales;
- fold adjacent Q/K scales for K64 access;
- emit the layouts expected by the fused attention kernel;
- route extreme-logit layers through the sampled guard.
This could become a model-conversion pass or compiler intermediate representation for Blackwell deployment.
Dependencies: the paper emphasizes that ordinary block-16 scales are an invalid operand contract for the reported binary; quantization and layout generation must therefore be co-designed with the kernel.
- Hardware-aware GPU kernel design methodology (academia and industrial GPU performance engineering; deployable now)
- treat on-chip tensor-memory allocation as a resource-ownership problem;
- model producer–consumer publication events explicitly;
- reuse retired score storage only after the consumer has completed;
- optimize work before the first legal consumer instruction;
- balance special-function-unit and arithmetic-pipeline use.
These methods are applicable beyond attention to fused normalization, mixture-of-experts routing, recurrent kernels, and other asynchronous tensor-core workloads.
- Benchmarking and regression-testing tools for low-precision attention (academia, hardware vendors, software engineering; deployable now)
- kernel latency and throughput;
- cosine similarity;
- relative- error;
- RMSE;
- final hidden-state and logit error;
- task accuracy and prediction agreement;
- end-to-end model step time.
This avoids treating a kernel-level speedup as evidence of end-to-end model quality or training stability.
- Selective precision policies for safety-critical inference (healthcare, finance, public-sector AI; deployable now with conservative controls)
- low-margin classification layers;
- final logits;
- retrieval or ranking stages where small score changes affect ordering;
- medically or financially consequential predictions.
A practical workflow would monitor logit margins and automatically fall back to higher precision when the margin is below a calibrated threshold.
Dependencies: application-specific risk assessment, formal validation, reproducibility controls, and regulatory requirements. The paper does not establish safety for clinical or financial decisions.
Long-Term Applications
- Full low-precision transformer training with quantized causal backward (AI training infrastructure; requires further research and scaling) Extend the causal path so that forward-produced quantized Q/K state, softmax normalizers, and FP8 gradient views are reused during backward. The reported complete single-GPU 8-billion-parameter update speedup of up to 1.14× suggests a path toward reducing training memory traffic and attention cost.
Dependencies: long-horizon convergence studies, larger models, diverse datasets, optimizer compatibility, loss-scaling strategies, and validation beyond short fixed-token updates.
- Distributed training with mixed-precision attention (large-scale AI systems; long-term)
- FP4 Q/K for matrix-product efficiency;
- FP8 P/V for stability;
- BF16 or FP8 learned projections;
- higher precision for loss and optimizer states.
Dependencies: stable distributed trajectories, communication precision, checkpoint compatibility, optimizer behavior, and model-family-specific calibration. Full MXFP4 P/V training should not be deployed based on the current evidence.
- End-to-end FP4 or near-FP4 transformer training (AI research and accelerator design; long-term)
- adaptive block sizes;
- learned probability codebooks;
- stochastic or error-feedback quantization;
- hybrid FP4/FP8 probability representations;
- higher-precision accumulation or correction paths;
- layer- and head-specific precision selection.
Dependencies: the present E2M1/MXFP4 probability path exhibits materially higher operator error than FP8 and failed distributed training tests, so new numerical mechanisms are required.
- Hardware support for a third score or probability destination (GPU architecture; long-term) Modify tensor-memory allocation or issue semantics to provide another 128-column destination for D128 attention. This could allow deeper look-ahead, reduce score/PV storage conflicts, and improve overlap among QK, softmax, and PV.
Potential hardware innovations: allocatable TMEM banks, more flexible score-overlay semantics, additional asynchronous operand queues, or instructions that transfer quantized probabilities without consuming the score bank.
Dependencies: silicon area, power, compiler exposure, bank conflicts, and evidence that additional storage improves end-to-end workloads rather than only isolated kernels.
- Compiler-generated hardware-specialized attention (ML compilers and domain-specific languages; long-term)
- GPU generations;
- head dimensions;
- causal versus noncausal masks;
- quantization formats;
- sequence-length regimes;
- TMEM capacities.
Dependencies: accurate cost models, access to low-level hardware semantics, reliable register/TMEM allocation, and automated numerical validation.
- Dynamic per-layer or per-token precision routing (adaptive inference and training; long-term)
- observed logit range;
- sampled row maxima;
- quantization error;
- output-margin estimates;
- layer sensitivity;
- sequence length and hardware occupancy.
This could provide most of the FP4 speedup while avoiding high-error or unstable cases.
Dependencies: routing overhead must remain lower than the saved computation; sampled guards must be shown reliable across substantially broader model and input distributions.
- Low-power and high-throughput transformer deployment (robotics, autonomous systems, edge data centers, energy efficiency; long-term)
- onboard robotic perception;
- autonomous-vehicle scene understanding;
- real-time multimodal assistants;
- industrial inspection;
- wearable or mobile vision-language systems.
Dependencies: current results target data-center GPUs; power, memory capacity, thermal constraints, and latency behavior on embedded hardware remain untested.
- Attention-specific numerical verification and certification tools (policy, regulated AI, academia; long-term)
- operator;
- layer;
- model output;
- task metric;
- training trajectory.
Such tools could support deployment documentation and risk controls for healthcare, finance, education, and public-sector systems.
Dependencies: agreed tolerances, representative test sets, robustness testing under distribution shift, and standards for documenting quantization-induced behavior.
- New FP4 probability formats and hardware instructions (semiconductor research; long-term) The results motivate formats designed specifically for softmax probabilities rather than general floating-point values. Potential innovations include asymmetric probability codebooks, denormal-preserving scales, fused normalized-PV instructions, or instructions that directly consume log-score fragments and produce scaled probability operands.
Dependencies: compatibility with online softmax, stable normalization, efficient scale publication, and demonstrated training as well as inference benefits.
Glossary
- Attention: A mechanism that computes weighted combinations of value vectors using query–key similarity scores. “Attention combines a query matrix , key matrix , and value matrix ”
- Affine classifier: A classifier that maps an input using a linear transformation plus an offset. “We therefore fit an affine classifier in value space”
- Asynchronous matrix operation: A matrix computation issued without synchronously blocking other GPU work. “Blackwell's asynchronous matrix hardware”
- Bfloat16 (BF16): A 16-bit floating-point format with a wide exponent range and reduced precision. “bfloat16 (BF16) forward throughput”
- Block scale: A shared scaling factor used to represent a group of quantized values. “each N32 probability block uses one MXFP4 E8M0 scale”
- Cooperative thread array (CTA): A CUDA thread block whose threads cooperate on a GPU task. “A cooperative thread array (CTA) is a CUDA thread block scheduled on an SM.”
- Cosine similarity: A measure of angular agreement between two vectors. “Cosine can hide magnitude error”
- Critical path: The sequence of dependent operations that determines minimum execution latency. “P lies on the critical path”
- Cross-attention: Attention in which queries and keys or values originate from different sequences or representations. “cross-attention and the rest of the model remain BF16”
- CUDA: NVIDIA’s parallel-computing programming platform and execution model. “A cooperative thread array (CTA) is a CUDA thread block scheduled on an SM.”
- CuTe: A CUDA library and domain-specific abstraction for composing tensor layouts and GPU operations. “HAO's generated CuTe domain-specific-language (DSL) BF16 kernel”
- Denominator: The normalization term in a softmax-weighted sum. “Direct-P instead builds the denominator from the exact codes and block scales consumed by PV.”
- Domain-specific language (DSL): A programming language or abstraction specialized for a particular computational domain. “CuTe domain-specific-language (DSL) BF16 kernel”
- E2M1: A 4-bit floating-point encoding with two exponent-related bits and one explicit fraction bit. “E2M1 payloads”
- E4M3: An 8-bit floating-point format with four exponent bits and three explicit fraction bits. “NVFP4 uses fine-grained data-dependent scales”
- E8M0: An exponent-only 8-bit floating-point scale format with no explicit fraction bits. “MXFP4 shares one power-of-two scale across each 32-value block.”
- Epilogue: The final computation or transformation performed after a main GPU matrix operation. “The scaled matrix instruction uses ; its output epilogue applies the fixed $1/36$ correction”
- Exponentiation: The computation of a number raised to a power, such as evaluating in softmax. “FP4 accelerates QK and PV, but not score reduction, exponentiation”
- FP4: A 4-bit floating-point numerical representation used for compact storage and high-throughput computation. “Blackwell's 4-bit floating-point (FP4) tensor cores”
- FP8: An 8-bit floating-point representation used for lower-precision computation and storage. “The causal path reconstructs probabilities from saved quantized queries and keys and uses 8-bit floating-point (FP8) gradient operands”
- FlashAttention: A memory-efficient tiled attention algorithm that avoids materializing full score and probability matrices. “FlashAttention evaluates these equations in tiles”
- Fused multiply–add (FMA): An operation that computes multiplication and addition in a single instruction. “A packed two-lane FMA instruction”
- Gradient epilogue: The final stage of a gradient computation that formats or stores gradient outputs. “The projection and gradient epilogues also publish the row- or column-oriented FP8 views”
- High-bandwidth memory (HBM): GPU memory designed to provide very high data-transfer bandwidth. “the quadratic score and probability matrices need not be stored in high-bandwidth memory (HBM)”
- Logit: An unnormalized scalar score that is commonly converted into a probability distribution. “a few late Wan layers produce BF16 logits above 500”
- Lookup table (LUT): A table used to replace repeated computation with indexed retrieval. “Integer threshold trees, lookup tables (LUTs), and custom nibble packing”
- Matrix multiply–accumulate (MMA): An operation that multiplies matrices and accumulates the result into an accumulator. “a matrix multiply--accumulate (MMA) instruction performs the tensor-core product”
- MXFP4: A block-scaled FP4 format that uses a shared power-of-two scale for each block of values. “MXFP4 shares one power-of-two scale across each 32-value block.”
- NVFP4: An FP4 format using fine-grained, data-dependent scaling. “NVFP4 uses fine-grained data-dependent scales”
- Online softmax: A streaming softmax algorithm that updates normalization statistics as tiles arrive. “Why FlashAttention uses an online softmax”
- Operand: A value or tensor supplied to an arithmetic or matrix operation. “all four operands inside attention---, , , and ---can use FP4”
- Quantization: The process of mapping higher-precision numerical values to a lower-precision representation. “online P quantization and scale movement as central costs”
- Quantization-aware training (QAT): Training that simulates or incorporates quantization effects during optimization. “an attention quantization-aware-training study”
- Quantizer: A rule or algorithm that maps continuous or high-precision values to discrete representable levels. “They are quantizer-calibration parameters, not changes to exact softmax.”
- Relative- error: The Euclidean error between an approximation and reference, normalized by the reference norm. “relative- is reported alongside both”
- Root-mean-square error (RMSE): The square root of the mean squared difference between predicted and reference values. “root-mean-square error (RMSE) depends on output scale”
- SASS: NVIDIA’s low-level assembly language for GPU machine instructions. “generated more NVIDIA machine code (SASS)”
- Softmax: A function that converts a vector of scores into normalized exponential probabilities. “Softmax must still reduce each score tile, evaluate exponentials”
- Special-function unit (SFU): A GPU execution unit optimized for functions such as exponentials and logarithms. “FA4 routes some exponential work from special-function units (SFUs)”
- Streaming multiprocessor (SM): A major parallel execution unit within an NVIDIA GPU. “A graphics processing unit (GPU) contains streaming multiprocessors (SMs).”
- Tensor core: Specialized GPU hardware for high-throughput matrix operations. “Blackwell's 4-bit floating-point (FP4) tensor cores”
- Tensor Memory Accelerator (TMA): Hardware that asynchronously transfers tensor tiles between memory levels. “The Tensor Memory Accelerator (TMA) moves tiles”
- Tensor memory (TMEM): On-chip GPU storage used for accumulators in asynchronous tensor operations. “Tensor memory: Blackwell's on-chip accumulator scratchpad for asynchronous matrix operations.”
- Tile: A submatrix processed as a unit to reduce memory usage and expose parallelism. “FlashAttention evaluates these equations in tiles”
- Underflow: A numerical condition in which a value becomes too small to be represented and is rounded to zero. “unscaled NVFP4 exposes underflow”
- Warp: A group of GPU threads that execute instructions together. “A load warp keeps K and V ahead in a TMA-fed shared-memory ring.”
- Warpgroup: A collection of GPU warps cooperating on a larger operation. “Two softmax warpgroups (WGs), each made of four warps”
- Zero-point-free power-of-two scaling: Scaling based on powers of two without an additive zero-point, enabling efficient exponent-based representation. “MXFP4 instead stores a power-of-two block amplitude near ”
