GB10 Dissection

A reproducible empirical map of NVIDIA GB10 and SM121.

This repository tests the parts of GB10 that spec sheets do not answer cleanly: sustained bandwidth, cache geometry, coherent-memory behavior, CUDA-core and tensor-core throughput, FP4/FP8 instruction support, launch overhead, concurrency, occupancy, and real batch-1 LLM decode.

The final clean run completed Phases 0 through 4 with zero verifier errors.

The five results worth knowing

1. “1 PFLOP” is not one dense, dtype-neutral number

Path	Verified median
FP16 cuBLAS	99.14 TFLOP/s
BF16 cuBLAS	97.69 TFLOP/s
FP8 E4M3 to BF16	175.81 TFLOP/s
FP8 E5M2 to BF16	176.69 TFLOP/s
NVFP4 to BF16	314.12 TFLOP/s
NVFP4 to NVFP4	380.25 TFLOP/s

Every CUTLASS cell passed reference verification. The generic 1 PFLOP comparison lacks the dtype, sparsity, scaling, accumulator, output, and shape conventions needed for an apples-to-apples denominator.

2. Qwen2.5-7B decode follows measured bandwidth

Metric	Result
Large-read plateau	236.62 GB/s
Qwen2.5-7B Q4_K_M runtime bytes	4.67712 GB
Weight-stream roofline	50.592 tok/s
Actual llama.cpp decode	44.146 tok/s
Actual / predicted	87.26%
Effective weight bandwidth	206.48 GB/s

The model revision, GGUF shard hashes, llama.cpp commit, CUDA backend, 99-layer offload, and ten retained samples are pinned in the result files.

3. One high-capacity vector path collapses, chunking restores it

Treatment	Useful bandwidth
Single float4, 31.5 GiB/array	218.05 GB/s
Single float4, 31.75 GiB/array	15.87 GB/s
Single float4, exact 32 GiB/array	15.84 GB/s
Scalar, exact 32 GiB/array	213.66 GB/s
Four chunks, exact 32 GiB/array	217.83 GB/s
Eight chunks	219.22 GB/s
Sixteen chunks	220.24 GB/s

Exact 32 GiB and minus 16 bytes were both slow, so this is not an exact 2^31 boundary. The onset moved between independent processes. The behavior is confirmed; the compiler, mapping, and hardware explanations remain open.

4. Low-precision throughput depends heavily on K

E4M3 treatment	Median
4096 square	175.95 TFLOP/s
5120 square	77.90 TFLOP/s
6144 square	54.81 TFLOP/s
8192 square	53.28 TFLOP/s
Double M only	170.52 TFLOP/s
Double N only	170.33 TFLOP/s
Double K only	78.39 TFLOP/s

Nsight Systems reported the same kernel specialization, grid, block, registers, shared memory, and zero local memory for matched K=4096 and K=8192 controls. The internal cause remains open.

5. SM121 supports key Blackwell paths, not every tested path

Feature	Result
FP16, BF16, TF32, INT8 MMA	Correct runtime plus matching SASS
E2M1, E4M3, E5M2 MMA	Correct runtime plus matching QMMA
Base cp.async.bulk	Works
Bulk L2 prefetch	Works
TMA	Works, 30 UTMA lines in verified CUTLASS SASS
Cluster launch and distributed shared memory	Work
Cluster tile::gather4	Rejected by CUDA 13 ptxas for sm_121a
tile::scatter4	Rejected by CUDA 13 ptxas for sm_121a
setmaxnreg.dec	Accepted, reported ignored, absent from SASS

Additional measured results

Platform

• NVIDIA GB10, compute capability 12.1
• 48 SMs
• Native sm_121 kernel correct in 10/10 runs
• CUDA 13.0, driver 580.159.03
• 128,452,841,472 CUDA-visible bytes, exactly equal to Linux MemTotal
• 24 MiB L2
• 101,376-byte launchable dynamic shared-memory ceiling
• 1,536 threads/SM and 24 blocks/SM confirmed empirically

Memory

Operation	Peak	1-32 GiB median
Read	239.62 GB/s	236.62 GB/s
Write	196.67 GB/s	195.38 GB/s
Copy	218.15 GB/s	215.99 GB/s
Triad	221.26 GB/s	219.05 GB/s
D2D cudaMemcpyAsync	228.05 GB/s	223.63 GB/s

Cache controls found:

• L1 transition between 64 and 128 KiB
• L2 transition between 16 and 32 MiB, consistent with the 24 MiB runtime value
• 32-byte sector transition
• 128-byte line transition
• 7.362 TB/s useful shared load+store bandwidth
• stride-32 shared access at 95 cycles versus 33 cycles for stride 1

At 8 GiB, explicit pinned VMM host-NUMA placement was only 0.226% below explicit device placement. Allocation class mattered more: ordinary malloc and managed memory were roughly 29% to 32% slower than cudaMalloc for the matched 1 GiB GPU-read treatment.

CUDA cores

Operation	Median
FP32 FMA	24.931 TFLOP/s
FP64 FMA	0.38937 TFLOP/s
INT32 IMAD, two-op convention	14.843 TOP/s

Measured FP64:FP32 ratio: 1:64.03.

Launches, occupancy, and overlap

• Ungraphed empty sequence: 2.019 us/kernel on the GPU timeline
• One 1,024-node graph: 0.434 us/node
• Device-sequence reduction: 78.50%, or 4.65x
• Two partial grids: 99.98% overlap
• Two saturated grids: 10.27% slower than their isolated serial sum
• Pinned H2D copy plus compute: 94.41% overlap
• Register use reduced 256-thread residency from 6 blocks/SM at 38 registers/thread to 2 blocks/SM at 103 registers/thread

Evidence contract

A positive instruction claim requires:

1. Successful compilation
2. Successful launch and synchronization
3. Correct output
4. Matching native SASS

Timed cells retain warmups, ten recorded iterations, raw values, medians, minima, maxima, correctness, and telemetry. FMA counts as two operations. GEMM counts as 2*M*N*K. Bandwidth is useful application-byte bandwidth, not a claim about directly counted physical LPDDR traffic.

See REPORT.md for the complete methodology, caveats, results, and implications.

Repository map

src/        CUDA benchmarks and instruction probes
scripts/    benchmark runners and phase verifiers
results/    final machine-readable JSON and manifests
evidence/   selected SASS, compiler diagnostics, telemetry, and raw controls
REPORT.md   complete standalone technical summary

The repository intentionally excludes:

• Qwen model shards
• compiled binaries
• build directories
• repetitive raw logs
• large Nsight databases
• machine credentials and connection configuration

Reproduce

Requirements:

• NVIDIA GB10 / SM121
• CUDA 13 with sm_121 and sm_121a support
• Python 3
• CMake and Git for CUTLASS
• Nsight Systems for the optional matched timeline controls
• CUDA-enabled llama.cpp for Phase 4

git clone https://github.com/sxuff/gb10-dissection.git
cd gb10-dissection
make clean
make -j4
./run_all.sh

Phase-specific runners:

./run_phase1.sh
./run_phase2.sh
./run_phase3.sh
./run_phase4.sh

CUTLASS is pinned to:

dcf215af68a2d08d305076c152a06f201728cd53

Qwen2.5-7B-Instruct-GGUF is pinned to:

bb5d59e06d9551d752d08b292a50eb208b07ab1f

llama.cpp is pinned to:

62061f91088281e65071cc38c5f69ee95c39f14e

Set CUTLASS_ROOT, LLAMA_ROOT, or LLAMA_BENCH if those dependencies are not in their default locations. Model files are fetched by scripts/fetch_qwen25_7b_model.sh and verified against fixed hashes.

Scope

These results characterize one GB10 system under the preserved driver, toolkit, source, library, shape, and runtime configuration. Library-path throughput is not a universal silicon ceiling. Compiler failures apply to the exact CUDA 13.0.88 probes, not every future toolkit.

License

Original benchmark code and documentation are MIT licensed. Vendored NVIDIA CUDA Samples and adapted CUTLASS examples retain their original BSD-3-Clause notices.