GitHub - secYOUre/nvfp4bench: A command-line utility to measure **peak achievable NVFP4 performance** on the NVIDIA DGX Spark (GB10 Grace Blackwell, `sm_121a`, CUDA 13), and report it against NVIDIA's headline **1 PFLOP NVFP4** (sparse) / **~500 TFLOPS** (dense) figures. · GitHub
/" data-turbo-transient="true" />
Skip to content
Search or jump to...
Search code, repositories, users, issues, pull requests...
-->
Search
Clear
Search syntax tips
Provide feedback
--><br>We read every piece of feedback, and take your input very seriously.
Include my email address so I can be contacted
Cancel
Submit feedback
Saved searches
Use saved searches to filter your results more quickly
-->
Name
Query
To see all available qualifiers, see our documentation.
Cancel
Create saved search
Sign in
/;ref_cta:Sign up;ref_loc:header logged out"}"<br>Sign up
Appearance settings
Resetting focus
You signed in with another tab or window. Reload to refresh your session.<br>You signed out in another tab or window. Reload to refresh your session.<br>You switched accounts on another tab or window. Reload to refresh your session.
Dismiss alert
{{ message }}
secYOUre
nvfp4bench
Public
Notifications<br>You must be signed in to change notification settings
Fork
Star
main
BranchesTags
Go to file
CodeOpen more actions menu
Folders and files<br>NameNameLast commit message<br>Last commit date<br>Latest commit
History<br>7 Commits<br>7 Commits
src
src
CMakeLists.txt
CMakeLists.txt
DESIGN.md
DESIGN.md
LICENSE
LICENSE
README.md
README.md
TODO_ldmatrix.md
TODO_ldmatrix.md
View all files
Repository files navigation
nvfp4bench
A command-line utility to measure peak achievable NVFP4 performance on the<br>NVIDIA DGX Spark (GB10 Grace Blackwell, sm_121a, CUDA 13), and report it against<br>NVIDIA's headline 1 PFLOP NVFP4 (sparse) / ~500 TFLOPS (dense) figures.
See DESIGN.md for the full architecture and rationale.
Layout
CMakeLists.txt<br>src/<br>main.cu CLI + orchestration<br>device_probe.* NVML clock + SM count + theoretical peak (CUDA 13 safe)<br>nvfp4.* E2M1/e4m3 codecs, block-scale packing, dequant<br>reference.* FP32 GEMM oracle for validation<br>bench.* timing, operand setup, CSV/JSON reporting<br>bandwidth.cu STREAM-triad bandwidth roofline probe<br>gemm.h unified backend interface<br>gemm_stub.cu on-the-fly dequant GEMM (correctness/bring-up, slow)<br>gemm_cutlass.cu CUTLASS NVFP4 block-scaled GEMM (trusted baseline, ~375 TFLOPS)<br>gemm_custom.cu from-scratch PACKED mxf4nvf4 warp-MMA GEMM (~112 TFLOPS) -- `custom`<br>gemm_warp.cu standalone byte-padded mxf8f6f4 kernel (~48 TFLOPS, validates + times)<br>gemm_warp_mxf4.cu standalone packed mxf4nvf4 kernel (~112 TFLOPS, validates + times)<br>peak.cu / peak.h pure tensor-core MMA throughput ladder -- the `--peak` mode<br>mma_probe.cu hardware RE: byte-padded mxf8f6f4 MMA semantics + fragment layout<br>mxf4_calib.cu hardware RE: packed mxf4nvf4 decode + ue8m0 scale + nibble model<br>mxf4_layout.cu hardware RE: packed mxf4nvf4 fragment layout (self-validating)<br>peak_mma.cu standalone 4-rung peak ladder (mxf8f6f4/mxf4nvf4, dense+sparse)<br>sparse_probe.cu confirms GB10 exposes the warp-level sparse FP4 MMA<br>mxf4_probe.cu confirms the native packed mxf4nvf4 (k64) MMA<br>mxf4_sparse_probe.cu confirms packed+sparse mxf4nvf4 (k128) -- the 1 PFLOP instruction<br>TODO_ldmatrix.md plan to push the custom kernel further via ldmatrix
Key findings (GB10 / DGX Spark)
The 1 PFLOP is real silicon — and we measured it. The memory-free MMA<br>throughput microbench (--peak) reaches 1022 TFLOPS packed+2:4-sparse NVFP4<br>(102% of the 1 PFLOP spec) and 511 TFLOPS packed dense (102% of the 500 spec).<br>The full ladder is 256 → 512 → 512 → 1022 (packed 2×, sparse 2×, combined 4×).
The catch is the instruction, not the silicon. The headline numbers need the<br>native packed kind::mxf4nvf4 MMA (2 codes/byte: m16n8k64 dense, m16n8k128<br>sparse). The byte-padded kind::mxf8f6f4 path most hand-written code starts from<br>runs at exactly half rate (256/512) — one code per byte, half the K per instr.
GB10 does expose the warp-level sparse block-scaled FP4 MMA (mma.sp … block_scale) — not confined to datacenter sm_100/tcgen05. (ue8m0 scales cap<br>at scale_vec::2X; 4X needs ue4m3.)
Real GEMMs are bandwidth-bound far below this. CUTLASS peaks at ~375 TFLOPS<br>dense (4096×14336×4096) — its memory ceiling, not the compute ceiling. Long<br>square GEMMs (≥8192³) also thermally throttle, so the GEMM path reports both burst<br>peak (min time) and sustained (median).
GB10 has no tcgen05/TMEM/2-SM MMA (those are datacenter sm_100). NVFP4 runs on<br>the GeForce warp-level block-scaled mma.sync...mxf4nvf4/mxf8f6f4.
The from-scratch custom kernel is built on a fully reverse-engineered fragment<br>layout + FP4 decode for both the byte-padded mxf8f6f4 and the native packed<br>mxf4nvf4 formats (undocumented; see DESIGN.md, mma_probe.cu, mxf4_layout.cu).<br>It now uses the packed path and is validated at ~112 TFLOPS (2048³,<br>L2-resident) — 2.3× the byte-padded kernel, with bit-exact correctness vs the FP32<br>oracle...