GB10 really does hit ~1 PFLOP NVFP4 (2:4 sparse) — measured, with an open-source tool to reproduce it

Accelerated Computing DGX Spark / GB10 User Forum DGX Spark / GB10 Projects
, ,
1

GB10 really does hit ~1 PFLOP NVFP4 (2:4 sparse) — measured, with an open-source tool to reproduce it

TL;DR. I built a small CLI, nvfp4bench, to measure peak achievable NVFP4 tensor-core
throughput on DGX Spark (GB10, sm_121a, CUDA 13) and check it against NVIDIA’s headline
figures, which so many times have been called into question.

On my unit the pure tensor-core peak reaches ~511 TFLOPS dense (102% of the 500
spec) and ~1014–1022 TFLOPS for 2:4-sparse NVFP4 (102% of the 1 PFLOP spec). So the 1 PFLOP
number is real silicon — but only via the native packed mxf4nvf4 instruction plus 2:4
sparsity. The byte-padded mxf8f6f4 path that hand-written code often starts from runs at
exactly half rate. Tool + full write-up:

The measurement

The peak is measured with a register-resident microbenchmark: many independent back-to-back
warp MMAs, no global/shared memory traffic, so it isolates the tensor-core issue rate from the
LPDDR5 bandwidth ceiling. Four rungs:

Instruction Format K/instr Measured vs spec
mma.sync … kind::mxf8f6f4 byte-padded FP4, 1 code/byte 32 ~256 TFLOPS 51% of 500
mma.sp … kind::mxf8f6f4 + 2:4 sparse 64 ~511 TFLOPS 51% of 1000
mma.sync … kind::mxf4nvf4 packed FP4, 2 codes/byte 64 ~511 TFLOPS 102% of 500
mma.sp … kind::mxf4nvf4 packed + 2:4 sparse 128 ~1014–1022 TFLOPS 102% of 1000

A clean 256 → 512 → 512 → 1022: packed = 2.00×, sparse = 2.00×, combined = 4.00×. Both
top rungs sit at ~101–102% of spec (running at the observed ~2.6 GHz under load).

The practical catch

The thing worth flagging for anyone writing FP4 kernels by hand: the headline numbers require
the native packed kind::mxf4nvf4 MMA (2 codes/byte, m16n8k64 dense / m16n8k128 sparse).
The mixed kind::mxf8f6f4 path (one FP4 code per byte) is the natural first thing to reach for,
and it tops out at exactly half rate because it moves half the K per instruction. If your
hand-rolled FP4 GEMM is landing around 250 dense / 500 sparse, this is very likely why.

A couple of things that surprised me on GB10

  • GB10 does expose the warp-level sparse block-scaled FP4 MMA (mma.sp::ordered_metadata … block_scale). I had expected it to be a datacenter sm_100/tcgen05 feature, but the raw
    PTX assembles and runs on sm_121a and delivers the full 2× of the dense path.
  • GB10 has no tcgen05/TMEM/2-SM MMA — NVFP4 runs on the GeForce-style warp-level
    mma.sync/mma.sp … block_scale.
  • For kind::mxf4nvf4, ue8m0 scales cap at scale_vec::2X; scale_vec::4X requires
    ue4m3 (the block-16 NVFP4 microscale). ptxas rejects 4X+ue8m0.
  • Toolchain gotcha that cost me a while: you must pass the accelerated arch explicitly,
    -gencode=arch=compute_121a,code=sm_121a. The -arch=sm_121a shorthand (and some CMake
    paths) silently drops the trailing a, producing plain sm_121, where the block-scaled FP4
    MMA either isn’t emitted or aborts at runtime.

Real GEMMs vs the peak (the honest part)

The 1 PFLOP is a pure tensor-core ceiling. Any real GEMM on GB10 is bandwidth-bound far
below it by the 273 GB/s LPDDR5:

  • A CUTLASS NVFP4 GEMM (CUTLASS ≥ 4.2.1, ArchTag Sm120) peaks around ~375 TFLOPS dense at a favorable shape (e.g. 4096×14336×4096) — that’s its memory ceiling, not the compute ceiling. Large square shapes (≥8192³) also thermally throttle, so the tool reports both burst peak and sustained.
  • My from-scratch packed mxf4nvf4 warp-GEMM validates bit-exact against an FP32 oracle and runs at ~112 TFLOPS at 2048³ (L2-resident) — 2.3× a byte-padded version built from the
    same source. It’s not competitive with CUTLASS (no ldmatrix/TMA-class machinery), but it’s a
    fully validated kernel built directly on the reverse-engineered fragment layout.

What’s in the tool

  • --peak — the register-resident MMA ladder above (this is where the 1 PFLOP shows up).
  • --kernel cutlass — a trusted CUTLASS NVFP4 baseline.
  • --kernel custom — the hand-written, validated packed mxf4nvf4 GEMM.
  • --bandwidth — a STREAM-triad roofline probe (to show where the memory wall is).
  • The reverse-engineering probes that made the custom kernel possible: the (undocumented)
    SM120/sm_121a block-scaled FP4 decode model and fragment layout for both the
    byte-padded and packed formats, extracted by one-hot hardware probing and self-validated.
    The packed model: true E2M1 decode (0, .5, 1, 1.5, 2, 3, 4, 6), ue8m0 scale 2^(E−127)
    (0x7F = unit), two scale bytes per scale_vec::2X each covering one 32-wide K-block.

Reproduce

git clone https://github.com/secYOUre/nvfp4bench
cd nvfp4bench
# headline peak (no CUTLASS needed):
nvcc -gencode=arch=compute_121a,code=sm_121a -O3 -o peak_mma src/peak_mma.cu && ./peak_mma
# or the full tool:
cmake -B build && cmake --build build -j
./build/nvfp4bench --peak

I’d love for others with a DGX Spark to run --peak and post their numbers — curious how much
unit-to-unit and thermal variation there is, and whether anyone sees the ue4m3/scale_vec::4X
path behave differently. Feedback, corrections, and PRs very welcome; some of the layout work is
reverse-engineered and I’d be glad to have it scrutinized.

(Measured on a single GB10 / DGX Spark, CUDA 13.0, driver 580.x. Numbers are from one unit and
will vary with clocks/thermals.)

2

Asus GX10

./build/nvfp4bench --peak --bandwidth  
Device 0: NVIDIA GB10  (sm_121) 
  SMs               : 48 
  Max SM clock      : 3.00 GHz 
  Current SM clock  : 0.21 GHz 
  Total memory      : 121.6 GB 
  L2 cache          : 24.0 MB 
  Mem bandwidth(spec): 273 GB/s 
  NVFP4 peak (spec)  : 500 TFLOPS dense | 1000 TFLOPS 2:4 sparse 
  NVFP4 peak (est)   : 620.4 TFLOPS dense | 1240.8 TFLOPS 2:4 sparse  (clock-scaled) 
 
Memory bandwidth (triad): 205 GB/s  (spec 273 GB/s) 
 
 
Peak NVFP4 tensor-core throughput (register-resident, no memory traffic) 
  device: NVIDIA GB10, 48 SMs, 1536 warps x 16 acc x 700 iters 
  mxf8f6f4 dense     248.7 TFLOPS  ( 49.7% of 500 spec) 
  mxf8f6f4 sparse    496.8 TFLOPS  ( 49.7% of 1000 spec) 
  mxf4nvf4 dense     497.5 TFLOPS  ( 99.5% of 500 spec) 
  mxf4nvf4 sparse    994.8 TFLOPS  ( 99.5% of 1000 spec) 
  ladder: 249 -> packed 497 -> packed+2:4 995 TFLOPS  (packed 2.00x, sparse 2.00x) 
  => GB10 reaches 99% of the 500 TFLOPS dense and 99% of the 1 PFLOP sparse NVFP4 spec.
3

Thanks for reporting. Invoking only with the --peak switch should lead to marginally higher figures. I am anyway happy to see that everything looks solid also on your unit.

4

Output from my HP ZGX Nano G1N (the HP variant of the Spark):

zgx-spark:~/tests/nvfp4bench$ ./build/nvfp4bench --peak --bandwidth
Device 0: NVIDIA GB10  (sm_121)
  SMs               : 48
  Max SM clock      : 3.00 GHz
  Current SM clock  : 0.78 GHz
  Total memory      : 121.6 GB
  L2 cache          : 24.0 MB
  Mem bandwidth(spec): 273 GB/s
  NVFP4 peak (spec)  : 500 TFLOPS dense | 1000 TFLOPS 2:4 sparse
  NVFP4 peak (est)   : 620.4 TFLOPS dense | 1240.8 TFLOPS 2:4 sparse  (clock-scaled)

Memory bandwidth (triad): 207 GB/s  (spec 273 GB/s)


Peak NVFP4 tensor-core throughput (register-resident, no memory traffic)
  device: NVIDIA GB10, 48 SMs, 1536 warps x 16 acc x 700 iters
  mxf8f6f4 dense     251.9 TFLOPS  ( 50.4% of 500 spec)
  mxf8f6f4 sparse    503.8 TFLOPS  ( 50.4% of 1000 spec)
  mxf4nvf4 dense     503.9 TFLOPS  (100.8% of 500 spec)
  mxf4nvf4 sparse   1007.8 TFLOPS  (100.8% of 1000 spec)
  ladder: 252 -> packed 504 -> packed+2:4 1008 TFLOPS  (packed 2.00x, sparse 2.00x)
  => GB10 reaches 101% of the 500 TFLOPS dense and 101% of the 1 PFLOP sparse NVFP4 spec.

Output from my FE Spark unit:

dgx-spark:~/tests/nvfp4bench$ ./build/nvfp4bench --peak --bandwidth
Device 0: NVIDIA GB10  (sm_121)
  SMs               : 48
  Max SM clock      : 3.00 GHz
  Current SM clock  : 0.70 GHz
  Total memory      : 121.7 GB
  L2 cache          : 24.0 MB
  Mem bandwidth(spec): 273 GB/s
  NVFP4 peak (spec)  : 500 TFLOPS dense | 1000 TFLOPS 2:4 sparse
  NVFP4 peak (est)   : 620.4 TFLOPS dense | 1240.8 TFLOPS 2:4 sparse  (clock-scaled)

Memory bandwidth (triad): 207 GB/s  (spec 273 GB/s)


Peak NVFP4 tensor-core throughput (register-resident, no memory traffic)
  device: NVIDIA GB10, 48 SMs, 1536 warps x 16 acc x 700 iters
  mxf8f6f4 dense     243.5 TFLOPS  ( 48.7% of 500 spec)
  mxf8f6f4 sparse    486.8 TFLOPS  ( 48.7% of 1000 spec)
  mxf4nvf4 dense     486.3 TFLOPS  ( 97.3% of 500 spec)
  mxf4nvf4 sparse    973.2 TFLOPS  ( 97.3% of 1000 spec)
  ladder: 243 -> packed 486 -> packed+2:4 973 TFLOPS  (packed 2.00x, sparse 2.00x)
  => GB10 reaches 97% of the 500 TFLOPS dense and 97% of the 1 PFLOP sparse NVFP4 spec.
5

Thanks for your reports. Everything looks generally in very good order. Minor variations in the figures are related to the exact clock and thermals. Overall the marketing claim holds up.

6

Jenson be like: pfeww, I totally winged it but those guys really found a way how to get there!

7

You made me laugh – thank you! – and hope at NVIDIA our colleagues feel relieved too. No class action! 😂

As the old saying goes: Trust, but verify! In fact vendor performance claims are checkable – This is the takeaway I keep coming back to.

8

Do we know approximately how much are we currently squeezing when using NVFP4 models?

If so, how much more performance in practical tok/s should we be able to get when (if) we get the full FP4 speed?

(i.e.: Qwen3.6-35b-A3B-NVFP4 with MTP is currently producing 100-120 tok/s. How would this translate in “flops” and how short are we from the full power so we can extrapolate this to end-user performance?)

Thanks! I’m just curious here :)

9

Great question — and the answer is a bit counterintuitive: for single-stream token generation (decode), the NVFP4 compute peak is almost irrelevant. Decode is memory-bandwidth bound, not compute bound, so getting “full FP4 speed” on the tensor cores would buy you ~0 extra tok/s. Here’s the back-of-the-envelope, which you can redo with your exact numbers.

Why decode is memory-bound

Generating one token requires reading (essentially) all the active weights once. For an A3B MoE that’s ~3B active params/token. NVFP4 is ~0.5 byte/param (4-bit + block scale), so:

bytes/token ≈ 3e9 params × 0.5 B ≈ 1.5 GB   (+ KV-cache reads, context-dependent)
ceiling     ≈ 273 GB/s ÷ 1.5 GB ≈ ~180 tok/s   (single stream, no speculation)

You’re seeing 100–120 tok/s with MTP, i.e. roughly 65–80% of the bandwidth ceiling — that’s actually a healthy number. (MTP/speculative decoding is exactly the right lever here, because it amortizes one weight-read across several candidate tokens.)

How much compute is that, in FLOPS?

Decode does ~2 FLOP per active param per token:

2 × 3e9 × 120 tok/s ≈ 0.7 TFLOPS

…against a measured NVFP4 peak of ~500 TFLOPS dense / ~1000 TFLOPS 2:4-sparse. So during decode the tensor cores run at well under 0.1% of peak — they’re sitting idle waiting on LPDDR5. That’s the real meaning of “how much are we squeezing”: of memory bandwidth, a lot (~70%); of compute, almost nothing — and that’s expected, not a bug.

The roofline crossover (the actually-useful number)

Single-stream decode has an arithmetic intensity of ~2 FLOP / 0.5 B = 4 FLOP/byte. To become compute-bound on GB10 you’d need intensity above ~500e12 / 273e9 ≈ 1,830 FLOP/byte (dense) or~3,700 (sparse). You get there by reusing each weight across many tokens, i.e. batching:

need batch ≈ 1,830 / 4 ≈ ~450 concurrent tokens to saturate the FP4 tensor cores

So the FP4 compute peak pays off in exactly two regimes, neither of which is single-user decode:

  1. Prefill / long prompts — many tokens processed in parallel, genuinely compute-bound.
    Bigger prompts = more of that 500/1000 TFLOPS actually used.
  2. Concurrency / batched serving — aggregate tok/s across many simultaneous requests scales
    toward the compute peak (this is why throughput benchmarks at high concurrency look so
    different from single-stream reviews).

Bottom line

  • Single-stream decode tok/s is set by bandwidth ÷ bytes-per-token, not FLOPS. You’re already near that ceiling (~180 tok/s for this model); “full FP4 speed” won’t move it.
  • NVFP4 does already help your decode — vs FP8 it halves the bytes/token, so ~2× the tok/s.
    That benefit is from the smaller weights, not the faster math.
  • To go faster on decode: more aggressive speculation (you’re doing MTP — tune acceptance),
    smaller active footprint, or KV-cache quantization for long contexts.
  • To see the petaflop: prefill-heavy and/or batched/concurrent workloads.

(All approximate — swap in your real active-param count, NVFP4 byte overhead, and context length. If it helps, the FP4 peak + a STREAM bandwidth roofline are both measurable directly with the little tool I posted; the --bandwidth probe shows where the 273 GB/s wall actually sits on your unit.)

I hope this helps!

10

Same ASUS GX10 with
NVIDIA-SMI 610.43.02 CUDA UMD Version: 13.3

Device 0: NVIDIA GB10  (sm_121)
  SMs               : 48
  Max SM clock      : 3.00 GHz
  Current SM clock  : 0.21 GHz
  Total memory      : 121.6 GB
  L2 cache          : 24.0 MB
  Mem bandwidth(spec): 273 GB/s
  NVFP4 peak (spec)  : 500 TFLOPS dense | 1000 TFLOPS 2:4 sparse
  NVFP4 peak (est)   : 620.4 TFLOPS dense | 1240.8 TFLOPS 2:4 sparse  (clock-scaled)

Memory bandwidth (triad): 200 GB/s  (spec 273 GB/s)


Peak NVFP4 tensor-core throughput (register-resident, no memory traffic)
  device: NVIDIA GB10, 48 SMs, 1536 warps x 16 acc x 700 iters
  mxf8f6f4 dense     247.5 TFLOPS  ( 49.5% of 500 spec)
  mxf8f6f4 sparse    494.8 TFLOPS  ( 49.5% of 1000 spec)
  mxf4nvf4 dense     494.9 TFLOPS  ( 99.0% of 500 spec)
  mxf4nvf4 sparse    989.8 TFLOPS  ( 99.0% of 1000 spec)
  ladder: 247 -> packed 495 -> packed+2:4 990 TFLOPS  (packed 2.00x, sparse 2.00x)
  => GB10 reaches 99% of the 500 TFLOPS dense and 99% of the 1 PFLOP sparse NVFP4 spec.
11

I’ve posted on LinkedIn a more high-level introduction amenable to be shared, alongside a histogram of the observed ladder – and, of course, you are more than welcome to connect there:

https://www.linkedin.com/feed/update/urn:li:activity:7473012333426937856/

nvfp4bench_ladder
nvfp4bench_ladder1312×743 65.6 KB

12

Everyone, I am moving my announcement to GB10 Projects for the sake of classification/ease of find.

13

Extra compute would transfer to faster PP speed. This is why the DGX spark is significantly faster in PP speed than a strix halo or a mac studio. Actually being able to use the 1 PF during training would speed up training as well but the issue is how the hell do you use those cores lol.

14

Lenovo PGX

NVIDIA-SMI 580.159.03 Driver Version: 580.159.03 CUDA Version: 13.0

mike@pgx:~/nvfp4bench$ ./build/nvfp4bench --peak --bandwidth
Device 0: NVIDIA GB10 (sm_121)
SMs : 48
Max SM clock : 3.00 GHz
Current SM clock : 0.41 GHz
Total memory : 121.6 GB
L2 cache : 24.0 MB
Mem bandwidth(spec): 273 GB/s
NVFP4 peak (spec) : 500 TFLOPS dense | 1000 TFLOPS 2:4 sparse
NVFP4 peak (est) : 620.4 TFLOPS dense | 1240.8 TFLOPS 2:4 sparse (clock-scaled)

Memory bandwidth (triad): 213 GB/s (spec 273 GB/s)

Peak NVFP4 tensor-core throughput (register-resident, no memory traffic)
device: NVIDIA GB10, 48 SMs, 1536 warps x 16 acc x 700 iters
mxf8f6f4 dense 249.2 TFLOPS ( 49.8% of 500 spec)
mxf8f6f4 sparse 498.4 TFLOPS ( 49.8% of 1000 spec)
mxf4nvf4 dense 498.4 TFLOPS ( 99.7% of 500 spec)
mxf4nvf4 sparse 996.7 TFLOPS ( 99.7% of 1000 spec)
ladder: 249 → packed 498 → packed+2:4 997 TFLOPS (packed 2.00x, sparse 2.00x)
=> GB10 reaches 100% of the 500 TFLOPS dense and 100% of the 1 PFLOP sparse NVFP4 spec.

15

Thank you for your report; actually I was wondering about the Lenovo PGX, because I am considering buying one. :-)

16

Am I off kilter in wanting everything that support MTP to use the compute engine more and offset some of deficiency of the memory bandwidth? This is way outside my knowledge, but are there model design patterns that reduce memory pressure?

17

At my current underclock on Asus GX10 (1.99 GHz / 3.00 GHz max):

| Mode                | TFLOPS | % of spec |
|---------------------|--------|-----------|
| mxf8f6f4 dense      | 194.9  | 39%       |
| mxf8f6f4 sparse     | 389.8  | 39%       |
| mxf4nvf4 dense      | 389.8  | 78%       |
| mxf4nvf4 sparse     | 779.6  | 78%       |
| Custom GEMM (2048³) | 74.6   | —         |
18

You’re not off-kilter at all — that’s exactly the right instinct, and it has a name: you’re trying to raise arithmetic intensity (FLOPs per byte). Decode sits way down on the memory-bound side of the roofline, and there are two ways to climb toward the compute headroom:

  • raise the numerator → do more useful compute per weight-read (your former question), or
  • shrink the denominator → read fewer bytes per token (your latter question).

Both move you off the 273 GB/s wall toward the tensor cores that are otherwise ~99.9% idle during decode. They’re the same lever from two ends.

A. The numerator: “Use the compute engine to offset bandwidth”:

MTP / speculative decoding is precisely a spend-FLOPs-to-save-bytes trick. You read the weights once and use that single read to evaluate several candidate future tokens; every accepted draft is a token you emitted without paying for another full weight pass. That’s how you can beat the naive bandwidth ÷ bytes-per-token ceiling — you raised intensity by doing more compute per byte.

Same family, same principle:

  • Batching / concurrency — the cleanest version: each weight read serves B sequences, so intensity scales ~B×. This is what eventually saturates the FP4 tensor cores (you need a batch
    of a few hundred to get there on GB10).
  • Draft-model / self-speculation (Medusa, EAGLE, n-gram/prompt-lookup, and MTP heads).

The honest caveat: speculation helps most for low-batch, latency-sensitive decode — exactly where you are. Its gain is bounded by draft acceptance rate (rejected guesses are wasted compute), and once you’re already compute-bound (high batch) it stops helping and can even hurt. So “more MTP everywhere” is right for single-stream/interactive, less so for a saturated server.

B. The denominator: Patterns that reduce memory pressure

During decode you read two things each step: weights and the KV cache. Different tricks attack different terms (your list covers most of them):

Cut weight bytes/token

  • MoE / sparse activation — the big one, and why your A3B model already decodes fast for its
    size: total params huge, but only the active subset is read per token.
  • Quantization (NVFP4 etc.) — fewer bytes/param; 4-bit ≈ 2× the tok/s of 8-bit, purely from
    smaller weights (not faster math).
  • Smaller / distilled models (SLMs), weight tying.

Cut KV-cache bytes/token (dominates at long context)

  • GQA / MQA — fewer KV heads → smaller cache → less read per step (now basically standard).
  • MLA (Multi-head Latent Attention) — compresses KV into a low-rank latent; large KV savings.
  • Sliding-window / local attention — bounded KV reads regardless of context length (often
    mixed with a few global layers).
  • KV-cache quantization (int8/fp8/4-bit KV), cross-layer KV sharing (YOCO-style).

Change the scaling regime entirely

  • State-space models / Mamba / linear attention — replace the growing KV cache with a
    fixed-size recurrent state, so per-token memory is ~O(1) in context length instead of O(L). Hybrids (Jamba, Zamba, Mamba-Transformer) keep some attention layers for quality. This is the
    most fundamental attack on the long-context memory wall.

TL;DR

Your instinct is correct: you can’t add bandwidth, but you can need less of it per emitted token – either by doing more compute per weight-read (MTP, batching) or by designing the model to touch fewer bytes (MoE, GQA/MLA, sliding-window, SSMs, quantization). Both raise FLOPs/byte and walk decode up the roofline toward the compute that’s just sitting there. The exception is prefill and high-concurrency serving, which are already intensity-rich and where that FP4 peak shows up directly.

That being said, more than inference itself the DGX Spark is best suited to do fine-tuning and data science work…

I hope this helps!

19

Ah, interesting data point and tradeoff: 2/3 clock buys the 78% of peak NVFP4 performance.

20

Great news, and thank you for the investigations! Glad to see this is real.