Papers
Topics
Authors
Recent
Search
2000 character limit reached

Hardware-Aware FP4 FlashAttention-4

Published 3 Sep 2026 in cs.LG | (2609.04105v1)

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×\times 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×\times. Matched distributed training retains FP8 probabilities and values; every tested MXFP4 probability/value training trajectory diverges.

Authors (1)

Summary

  • The paper evaluates whether FP4 tensor cores can accelerate attention beyond BF16 through the development of Direct-P, a new method for efficient probability generation Direct-P achieves up to 2.125x speedup in noncausal inference but with a 5% accuracy tradeoff in certain metrics. Noncausal inference with Direct-P is reported to reach 2998 TFLOP/s on a benchmark, though at a slight trade-off of accuracy. Quantized causal backward improves projection-inclusive attention by up to 1.25x and complete single-GPU updates by up to 1.14x, but FP8 P/V is retained for training stability. Direct-Psic XP4 FlashAttention addresses the limitations of Blackwell core tensor memory

Problem formulation and central claims

“Hardware-Aware FP4 FlashAttention-4” (2609.04105) investigates whether Blackwell’s FP4 tensor cores can accelerate attention beyond the gains obtained by simply replacing BF16 matrix products with FP4 operations. The paper’s central argument is that this substitution is insufficient: attention contains two matrix multiplications, QKTQK^\mathsf{T} and PVPV, separated by online softmax, probability quantization, scale publication, synchronization, and tensor-memory ownership. Once the matrix products become faster, these intermediate operations become a critical-path bottleneck.

The work separates two objectives. For noncausal inference, it develops Direct-P, a probability-generation path that maps normalized score fragments directly to MXFP4 E2M1 codes rather than first evaluating a higher-precision exponential and subsequently discarding most of that precision. For causal training, it passes the forward quantization state into backward, allowing the backward kernel to reconstruct probabilities from saved quantized Q/KQ/K payloads, scales, and log-sum-exp statistics.

The principal claims are:

  • Direct-P achieves up to 2.13×2.13\times the BF16 forward throughput on GB200 at favorable D128 attention shapes.
  • Its speed advantage is accompanied by a substantial accuracy trade-off: the mean forward cosine is approximately $0.944$ for the fast policy, compared with approximately $0.990$ for the higher-precision NVFP4-QK/FP8-PV control.
  • Quantized causal backward improves projection-inclusive attention by up to 1.25×1.25\times and a complete single-GPU 8-billion-parameter update by up to 1.14×1.14\times.
  • FP8 P/V is retained for training because every tested MXFP4 P/V trajectory diverges in the reported long-running experiments.
  • The principal hardware limitation is not nominal FP4 tensor throughput but the absence of another legally allocatable score destination in Blackwell’s tensor memory (TMEM).

These results extend the hardware–algorithm co-design perspective established by FlashAttention (Dao et al., 2022), FlashAttention-2 (Dao, 2023), FlashAttention-3 (Shah et al., 2024), and FlashAttention-4 (Zadouri et al., 5 Mar 2026).

Architectural bottleneck: online softmax and TMEM ownership

FlashAttention avoids materializing the quadratic score and probability matrices in HBM by processing key tiles incrementally. For each query row, it maintains a running maximum, denominator, and output accumulator. This eliminates large memory traffic, but introduces a strict dependency chain: a score fragment must be reduced, normalized, converted into the representation consumed by PVPV, and published before the value product can proceed.

The implementation studied in the paper inherits a two-query Blackwell schedule in which one 16-warp CTA processes two query stages. QK and PV are issued through asynchronous tensor-core operations, while specialized warpgroups perform score processing, online correction, and output publication. At D128, each FP32 score or output tile occupies 128 TMEM columns. Two score banks and two persistent output banks consume all 512 available logical TMEM columns.

This layout creates a storage-ownership dependency. A score bank is temporarily reused in the sequence

QK scoreprobability and scale overlayPV consumptionbank release.\text{QK score} \rightarrow \text{probability and scale overlay} \rightarrow \text{PV consumption} \rightarrow \text{bank release}.

The next QK operation cannot overwrite the bank while its probability overlay is still needed by PV. Additional barriers can expose readiness, but they cannot create another legal destination. The paper therefore distinguishes a synchronization problem from an allocation problem: the limiting resource is not merely barrier expressiveness or shared-memory capacity, but the number of usable TMEM score destinations.

The K64 granularity of the scaled-FP4 PV instruction compounds this constraint. Four N32 score quarters are generated, but two adjacent quarters must be complete before the first K64 PV operation can be issued. Consequently, accelerating one probability fragment does not necessarily advance tensor-core work unless its paired fragment is also ready.

Direct-P and the forward FP4 path

Direct-P changes only the interval between a completed score fragment and a legal PV operand. The outer CTA schedule, score-bank reuse, two-query pipeline, and publication protocol remain unchanged. Its design has three components.

First, Direct-P treats probability formation as E2M1 code classification rather than approximate exponential evaluation. Given a score transformed into base-two coordinates, it selects one of the nonnegative E2M1 magnitudes

PVPV0

The implementation uses packed FFMA2 operations followed by Blackwell’s native floating-point conversion instruction. The fitted affine map is selected to place the E2M1 decision boundaries efficiently, not to minimize error against the continuous exponential function. This distinction is important: the value eventually consumed by PV is a four-bit code, so optimizing an intermediate FP32 exponential is computationally wasteful when its precision will be discarded.

Second, Direct-P normalizes using the represented FP4 probabilities actually consumed by PV. For each N32 block, it accumulates the MXFP4 amplitude multiplied by the sum of the emitted E2M1 codes. The numerator and denominator therefore describe the same approximate operator. This avoids a mismatch in which the numerator uses rounded FP4 values while the denominator is computed from an unrelated higher-precision exponential approximation.

Third, the implementation includes a selective guard for extreme model logits. The fast shiftless path is finite on the synthetic benchmark grid but encounters logits above 500 and sometimes 1000 in later Wan layers. Rather than scanning every score row, the guarded path samples fixed key rows to obtain an anchor, applies a compile-time margin and scale floor, and reassociates the denominator computation to avoid subnormal intermediate underflow. This guard is not a general stable-softmax fallback; it is a targeted correction for observed model distributions.

The format choice reflects a latency–range trade-off. Stabilized NVFP4 provides better probability fidelity but requires finer scale handling and additional range correction. MXFP4 offers power-of-two E8M0 scales aligned with the N32 producer granularity, but its E2M1 representation is coarser. In the paper’s Gaussian diagnostic at sequence length 4096, MXFP4 produces a probability relative-PVPV1 error of approximately PVPV2 and a PVPV3 cosine of approximately PVPV4, whereas unscaled NVFP4 exhibits catastrophic underflow. Stabilized NVFP4 achieves a probability relative-PVPV5 error near PVPV6 and PVPV7 cosine near PVPV8.

Forward performance and accuracy

Across the nine-row GB200 D128 benchmark suite, the fast Direct-P policy is reported as PVPV9 faster than the HAO BF16 baseline, reaching a maximum of 2998 TFLOP/s. The accurate policy reaches Q/KQ/K0 and 2416 TFLOP/s. The largest explicitly reported relative speedup is Q/KQ/K1 at S8192/H64, where Direct-P takes 0.758336 ms compared with 1.611488 ms for BF16.

B300 results show 5.6–7.7% lower latency than GB200 on S4096–S8192 rows, with 3116 TFLOP/s at S8192/H64 and 3159 TFLOP/s at the wave-aligned S9472/H64 shape. This improvement is shape-dependent rather than a uniform consequence of the newer GPU. At S32768/H24, B300 reaches 2945 TFLOP/s, below the 2998 TFLOP/s reported on GB200. The paper attributes such variation to launch geometry, persistent-grid width, and the fixed TMEM allocation.

The performance gain is not free. The fast Direct-P policy has mean cosine approximately Q/KQ/K2 and mean relative-Q/KQ/K3 approximately Q/KQ/K4 against BF16 in the principal operator suite. The accurate policy improves these values to approximately Q/KQ/K5 and Q/KQ/K6, respectively, but remains considerably less accurate than the NVFP4-QK/FP8-PV route, whose mean cosine is approximately Q/KQ/K7.

The paper’s accuracy-matched control clarifies the source of this trade-off. An exact local NVFP4-QK/FP8-PV route reaches cosine Q/KQ/K8 but only 1490 TFLOP/s, whereas Direct-P reaches 2945 TFLOP/s at cosine Q/KQ/K9 on the corresponding B300 long-sequence case. Thus the reported speed is not simply a consequence of FP4 QK; it depends materially on reducing the probability-construction path and accepting a coarser P representation.

Model-level fixed-input results are more favorable than the standalone operator metrics. On ViT S4096, fast Direct-P preserves BF16 top-1 accuracy at 88.5%, with 95.5% prediction agreement, while the accurate policy reaches 89.0% and 98.5% agreement. Across 2272 classification examples, 32 fast-policy predictions change, and 31 of those occur in the lowest quartile of BF16 top-two logit margins. This supports a margin-sensitive interpretation of the error, although the evaluation is not sufficient to establish general inference or training safety.

In the Wan2.1 diffusion evaluation, all self-attention layers use Direct-P while the remainder of the model remains BF16. Fast Direct-P is 1.75× faster for the 1.3-billion-parameter model and 2.09× faster for the 14-billion-parameter model. However, error accumulates over diffusion steps: for Wan2.1-14B, the 20-step output has cosine 0.8496 and relative-2.13×2.13\times0 0.5337 against BF16. The selective guard enables all 1600 attention calls to complete, but guarded layers are 21–23% slower individually; because only a minority of layers are guarded, the aggregate penalty is approximately 2.2–2.3%.

The ViT-MAE reconstruction experiment shows a smaller end-task displacement. After replacing all twelve encoder attention layers, fast Direct-P obtains reconstruction cosine 0.99973 and relative-2.13×2.13\times1 0.0203, with a PSNR decrease of 2.13×2.13\times2 dB. The result indicates that residual attention error can be attenuated by downstream network structure in some vision workloads, but this does not contradict the larger accumulated drift observed in long diffusion trajectories.

Quantized causal backward and training

The causal training method saves the forward NVFP4 Q/K payloads, block and global scales, and per-row LSE values. Backward reconstructs the represented scores and probabilities rather than generating a separate BF16 score path. Projection epilogues publish row- and column-oriented FP8 views needed by different gradient products, avoiding standalone transpose and quantization kernels.

The backward computation retains the standard dependency structure:

2.13×2.13\times3

followed by softmax centering and the 2.13×2.13\times4 and 2.13×2.13\times5 products. The implementation reuses the represented P state, publishes both physical 2.13×2.13\times6 layouts, and uses E5M2 rather than E4M3 for 2.13×2.13\times7. This choice is motivated by range rather than precision: E4M3 rounded approximately 97% of observed 2.13×2.13\times8 values to zero in a failing diagnostic, while E5M2 reduced the zero fraction to approximately 14% with less than 1% publisher overhead.

The isolated reconstruction core reduces D128 causal-backward latency from 0.501 ms for BF16 to 0.356 ms, a 2.13×2.13\times9 speedup. Once the E5M2 publisher and row-statistics publisher are included, latency increases to 0.508 ms, slightly slower than BF16. This result is important because it demonstrates that a faster inner backward kernel does not imply a faster training path; producer and publication overhead can erase the kernel-level gain.

The gain reappears when the method is integrated with projections and output gradients. At B1/S4096/D128, the projection-inclusive attention sublayer reaches:

Boundary BF16 Quantized route Speedup
Backward only 1.572 ms 1.397 ms $0.944$0
Forward plus backward 2.656 ms 2.133 ms $0.944$1

For a complete 8.03-billion-parameter update at S4096, speedup increases with local batch size: $0.944$2 at batch 1, $0.944$3 at batch 2, and $0.944$4 for FP8 P/V at batch 4. The MXFP4 P/V arm reaches $0.944$5 at batch 4, but FP8 and MXFP4 complete updates differ by at most 0.31%, so the isolated MXFP4 forward advantage does not produce a material end-to-end improvement.

The paper explicitly separates these timing results from training quality. Initial-logit cosine is only 0.416–0.426 for FP8 P/V and 0.373–0.374 for MXFP4 P/V in the short update experiment, and fixed-token timing does not establish convergence.

Why FP8 P/V is retained for training

The strongest negative result concerns MXFP4 P/V in longer training. In the reported four-arm diagnostic, both FP8-P/V trajectories remain non-divergent, while both MXFP4-P/V trajectories separate from their FP8 controls near 0.1 billion tokens and later exhibit rising loss and very large pre-clipping gradient norms. A matched B4 experiment with NVFP4 projections and MXFP4 P/V shows loss rising from 7.01 at update 325 to 16.25 at update 350.

The factorial structure strengthens the interpretation. The experiments vary learned projection precision while holding the attention backward path fixed. Both projection formats fail with MXFP4 P/V, whereas both remain stable with FP8 P/V. The evidence identifies P/V representation, or state changes induced by it, as the common separator. It does not establish whether the cause is forward probability quantization, the saved V payload, backward use of the representation, or an interaction among them.

The retained matched distributed study uses 64 GPUs, local batch four, four accumulation steps, and effective global batch 1024. The BF16 and NVFP4-projection/FP8-PV routes share model, optimizer, tokenizer, sample order, and token schedule over approximately 100 billion tokens. At the final scheduled training report, BF16 loss is 2.3095 and FP8-route loss is 2.3613. At the final same-update held-out validation report, losses are 2.3048 and 2.3948, respectively, yielding a gap of 0.0900.

The FP8 route is stable and descending but not numerically equivalent to BF16. Its median throughput is 24,303 tokens/s/GPU versus 21,853 for BF16, corresponding to a $0.944$6 median speedup. The paired throughput ratio has a 10th–90th percentile range of 1.080–1.114×.

Figure 1

Figure 1: Token-aligned training and held-out validation for the matched B4 experiment over the completed 100-billion-token schedule; the plot shows stable but non-identical FP8-P/V and BF16 trajectories without uncertainty estimates.

The training result therefore supports a narrower claim than “FP4 training.” The retained route uses NVFP4 in learned projections and Q/K attention operands, but FP8 P/V and FP8 gradient operands remain necessary under the reported stability evidence.

Hardware implications

The paper’s hardware analysis argues that Direct-P has already removed most of the exposed probability arithmetic. A historical matched diagnostic executes the same 98,304 tensor instructions with either real or fixed probability construction. Tensor-pipe activity rises from 18.8% to 26.6% when probability work is removed, but the tensor instruction count is unchanged. Source sampling attributes most not-issued samples to long-scoreboard dependencies involving final statistics, score readiness, and output publication.

A final speed-of-light diagnostic reinforces this conclusion. Relative to a valid 0.092448-ms kernel, retaining only raw score packing saves 1.38%, retaining row maxima saves 2.53%, and using a fixed probability tile saves 5.23%. Thus even eliminating nearly all probability construction cannot approach the nominal fourfold FP4 matrix-throughput ratio.

The paper proposes several hardware and instruction-set directions, but does not measure them as improvements:

  • an additional allocatable score bank with compatible issue semantics;
  • K32 scaled-FP4 PV instructions so each N32 probability fragment can be consumed immediately;
  • scale storage outside TMEM;
  • wider tiles with a compatible score, probability, and output lifecycle.

The first proposal is the most consequential. The paper’s analysis indicates that a larger usable overlap window would matter more than another polynomial approximation once Direct-P has shortened probability generation.

Limitations and open questions

The noncausal evidence consists primarily of fixed-input inference evaluations. It does not establish finetuning behavior, pretraining stability, or broad task robustness. The Wan results show substantial multi-step drift, especially for the 14-billion-parameter model, despite successful finite execution.

The distributed training result is based on one trajectory per route. It supports stability and throughput over the measured 100-billion-token schedule but provides no run-to-run variance estimate or statistical-equivalence claim. The FP8-versus-MXFP4 comparison also does not isolate the precise source of MXFP4 divergence. Determining whether the failure arises from forward P quantization, V quantization, backward reconstruction, or their interaction remains open.

The comparison boundaries also matter. Learned projection precision is separate from attention precision, and the reported end-to-end training route is not pure FP4. D64 and other head dimensions use different tile and TMEM ownership regimes, so the D128 conclusions cannot be transferred directly. Several B300 comparisons use different binaries or cross-run published baselines, which limits causal interpretation of hardware-generation differences. Finally, the paper’s hardware proposals are hypotheses: no additional score bank, K32 PV instruction, or external scale path is evaluated.

Conclusion

The paper establishes that FP4 attention performance on Blackwell is governed by the full score-to-PV pipeline rather than by tensor-core matrix throughput alone. Direct-P obtains up to $0.944$7 BF16 forward speed by mapping scores directly to MXFP4 probability codes and normalizing the represented probabilities consumed by PV. The resulting accuracy is materially below FP8-PV controls, although several fixed-input model evaluations show limited task-level impact.

For causal training, reusing the forward quantized state enables $0.944$8 projection-inclusive attention speedup and up to $0.944$9 complete-update speedup. However, the long-running experiments reject MXFP4 P/V for the reported training recipe and retain FP8 P/V. The dominant unresolved systems issue is TMEM ownership: at D128, the existing score and output allocation leaves insufficient space for a larger overlap window.

Paper to Video (Beta)

No one has generated a video about this paper yet.

Whiteboard

No one has generated a whiteboard explanation for this paper yet.

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:

  1. 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?
  2. Can FP4 make attention faster in practice? The researchers want to know whether the whole attention process becomes faster—not just one matrix multiplication.
  3. 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.
  4. 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.
  5. 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:

  1. Calculate a score tile.
  2. Find its largest score.
  3. Convert scores into probabilities.
  4. Quantize and pack the probabilities.
  5. 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 L2L_2 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 A=1.60,B=0.95A=1.60,B=0.95 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-L2L_2 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-L2L_2 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-L2L_2 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 QQ, key matrix KK, and value matrix VV
  • 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 αBqij\alpha_Bq_{ij}; its output epilogue applies the fixed $1/36$ correction”
  • Exponentiation: The computation of a number raised to a power, such as evaluating exe^x 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---QQ, KK, PP, and VV---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-L2L_2 error: The Euclidean error between an approximation and reference, normalized by the reference norm. “relative-L2L_2 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 aBa_B

Tweets

Sign up for free to view the 9 tweets with 215 likes about this paper.