CUDA Benchmarking - imp / sm_120 / RTX 5090

Pair with sm120-cuda-expert (levers) and docs/internals/BENCHMARKING.md (measurement contract).

STOP: what is a real signal on this box

#	Fact	Consequence
1	pp512 spread across process starts is a property of the MODEL: 0.6-1.2% on Qwen3-8B Q8_0, 37.6% on a resident NVFP4 MoE; cuBLAS algo re-selection alone 3.50% over nine starts	Decode is the A/B signal. Prefill-kernel deltas <=5% resolve ONLY in nsys per-kernel sums (tools/analysis/prefill_kernel_ab.sh), never in e2e pp. Gate key is tg128; ad-hoc runs print tg256
2	The GPU is water-cooled, idles ~30 C, never throttles	No temperature cooldowns. The 15 s cooldown in gen_perf_baseline.sh resets cuBLAS algo state
3	Idle downclock: clocks ramp ~1 s; between bench shapes they fall to 210 MHz	Discard >1 s warmup before timing; warm per shape in isolated benches. A cold shot once read -42% that re-measured +20%
4	Decode reads 8-15% low for a whole day (host/driver state, #526)	Sample clocks DURING the bench: healthy 2850 MHz SM / 13801 MHz mem / ~500 W. Never refresh a baseline or trust a cross-day delta without them
5	Back-to-back sweeps read 6-10% low	One model per process
6	--bench prompt is tokens[i] = i % vocab (512 distinct ids, tools/imp-cli/main.cpp; no 6-gram repeats at speculative.min_match=6); n-gram drafts come from the GENERATION looping on it, quantisation-dependent: Qwen3-14B NVFP4 accepts 504/504, Q6_K 6/96 and bistable	Decode-kernel A/Bs on dense models: --set speculative.ngram=false in BOTH arms, or a real +2-3% GEMV win is invisible
7	imp-cli --bench pins speculative.moe=false, speculative.suffix=false, speculative.hybrid=false, speculative.mtp_k=0, server.recurrent_snapshot_mb=0 (tools/imp-cli/args.cpp, apply_config_pins); an explicit --set wins	Every new auto-default needs a bench pin or gen_perf_baseline.sh bakes speculation into the gate (mtp auto engaged the head before its pin, 71.7 vs 86.8 tok/s)
8	A number is comparable only to a run with the same flags	--bench disables prefix caching since #1061 (--set server.prefix_cache=true restores the old band); tg128_at_ctx_2048 needs --bench-pp 2048; pp16/pp512 read ~+13%/+10% on the same build; verify-fast runs fewer trials than gen-perf-baseline. Log the imp.conf loaded from, dtype=, attn_decode= lines per arm (a stray repo-root imp.conf made two profile arms identical before #1784)
9	No-graphs profiles overstate tiny-kernel classes ~1.8x at batch=1 on MoE; on Qwen3.8 M=1 the decode graph is strictly serial (union == sum, #1797)	A batch=1 lever must hold bytes or critical-path math. At 32 streams the class pays (RMSNorm +6.8% #1769, act-quantize +4.6% #1771, producer fusion +2.6% #1773) but residual-accumulate lost -0.9% (#1793). Measure the regime you claim
10	Isolated kernel benches measure L2: every single Qwen3.8 weight fits the 96 MB L2 (>1792 GB/s = self-disqualified; the GQA-tile decode kernel built on that misread measured -9% e2e, #1785); balanced per-expert inputs flip the sign vs real routing (grouped GEMM mt32 -8.7% isolated, +4% in situ)	Rotate >=4-8 x 100 MB slabs; sub-wave shapes are bistable 1.6x between runs; verdicts from in-process nsys kernel sums; when the in-situ sweep ranks like the isolated one, record that too (#1768)
11	Capture costs under nsys are CUPTI-inflated (27.8 ms gaps; >1 ms gaps at the wave ramp = graph captures)	Never price graph capture from a profile
12	"Neutral" on a gate-based feature is usually a dead path	Prove activity by kernel launch counts (stream-K workspace refused every launch: 1960 -> 1400 CUTLASS launches, 22k -> 4.2k tok/s; sparse needs the sparse decode attention ACTIVE line in one arm only). A control MODEL can be dead too: a 32-stream A/B of an F16 paged-attention change on Nemotron-3-Nano ran FP8 KV (attention_paged_fp8*), a different kernel (#2074); read KV cache dtype: first
13	Batch=1 roofline: 1628 GB/s resident (8 GiB sweep, #1797), spilled ~237 GB/s; Qwen3.8 spec-off ceiling ~112 tok/s, measured 87.4	Re-measure the sweep before quoting bandwidth; the old 1530 pin is stale

Methodology (every A/B)

• CUBLAS_WORKSPACE_CONFIG=:4096:8, 10 reps, 3+ trials, one model per process, warm >1 s.
• GPU free first: make check-gpu (scripts/require_free_gpu.sh) or ~/.claude/skills/gpu-stats/gpu-busy-check.sh (utilisation + memory; a Windows tenant is invisible to docker ps and --query-compute-apps). gpu-busy-check && (nohup ...) ; echo launched prints "launched" on BUSY too; test the exit code.
• Server A/Bs: nothing besides the harness's own server; host 1-min load <= nproc/4 (bench_competitive.sh guard, IMP_BENCH_MAX_LOAD). An idle container is not load.
• Everything runs in Docker (imp:test, models from $HOME/models). imp-cli has no --ctx; in --bench mode the pool follows --bench-pp + --max-tokens + max(256, 12.5%) and --max-seq-len is overridden (an explicit --set kv_cache.max_blocks=N still wins); outside bench mode the ceiling is --max-seq-len. 32k prompts via --prompt-file (argv caps at ~128 KiB).
• Bench from binary COPIES (-v dir:/bin_arm): make dev rebuilds build-dev/ under a running chain.
• dtype/KV A/Bs: force emitted tokens equal (max_tokens below need); library-reserve state equal in both arms (a --rm container plans with the 3900 MiB constant, a mounted /home/imp/.cache/imp with its measurement: 716 blocks apart on Qwen3.8).
• PPL verdicts: --set runtime.deterministic=true both arms (0.35% run-to-run otherwise), tools/analysis/ppl_corpus_45k.txt (untracked, rebuild with tools/analysis/make_ppl_corpus.sh), --set speculative.mtp_k=0 (auto loads the head, +0.79 GiB). Qwen3.6-35B PPL moves +-0.2..0.5% between fp32-equivalent kernels (routing flips): numerics judge is Qwen3.8-27B-NVFP4-vllm (fused GDN 4.6283) plus the unit-test state diff and tools/analysis/layer_ab_diff.py added divergence.
• Every verdict ships its harness in tools/analysis/ or its md5 in the PROV block (scripts/bench_competitive.sh re-execs from a mktemp copy and prints harness: md5=).
• Check ad-hoc grep/awk against a known case first (grep -c 'hero$' also matched nonhero).

Pick the right tool

Goal	Tool
End-to-end engine perf	make bench (imp-cli --bench --bench-pp 512 --bench-reps 5)
Single model quick check	make test-perf (Qwen3-8B Q8_0)
Refresh perf baseline	make gen-perf-baseline [MODEL=/models/...]: 5 trials x 5 reps, 15 s cooldown, median; writes tests/perf_baseline.json incl. own_peak_mb
Guarded re-pin (refuses top-range / volatile days)	scripts/repin_baselines_if_median.sh
Regression gate	make verify-fast (scripts/verify.sh): 8% decode / 8% prefill / 10% peak VRAM
Paired regression gate, CODE change vs origin/main	make verify-ab (scripts/verify_ab.sh): imp:ab-<sha> vs imp:test, alternating pairs, mean paired decode delta vs thresholds.paired_decode_regression_pct; the pre-push hook runs it after verify-fast
VRAM attribution	the VRAM audit table every imp-cli run prints at init and shutdown (own_peak= is what the gate parses)
North-star gate	make verify-north-star (tests/perf_baseline_north_star.json)
Long-context decode A/B	scripts/bench_longctx_ab.sh; capacity vs decode tools/analysis/ctx_capacity_decode_sweep.sh
Prefill-kernel A/B for one config key (nsys kernel sum per arm)	tools/analysis/prefill_kernel_ab.sh; FA2 hd=256 variant fa2_hd256_bkv_ab.sh
32-stream aggregate A/B, config key	tools/analysis/smallm_v2_conc_ab.sh + tools/analysis/conc_client.py
32-stream aggregate A/B, CODE change	tools/analysis/two_image_conc_ab.sh (two prebuilt images, alternating)
Long-context serving A/B (tg8/tg520 differential)	tools/analysis/serving_sparse_ab.sh + longctx_conc_client.py
Serving idle attribution	tools/analysis/serving_idle_profile.sh + nsys_gap_attribution.py
MTP / speculation	tools/analysis/mtp_adaptive_ab.sh, mtp_k_sweep.sh, token_recycling_ab.sh, scripts/mtp_accuracy_bench.sh
Decoder ITL under concurrent ingest	scripts/bench_prefill_latency.py
Serving KPIs per concurrency, per-token ITL p50/p90/p95/p99/max	tools/analysis/serving_kpi.py --levels 1,8,32 [--process-workers]
Cold start: container start, load start, ready, first token (ms, N repeats)	scripts/bench_cold_start.sh
Per-config sweep MBU/MFU/TTFT/TBT	tools/bench/bench.py
Single kernel wall-clock	cudaEvent in the launcher (below)
Per-kernel metrics, stalls	ncu (below); wrapper tools/bench/profile.sh
Timeline, launch gaps, graphs	nsys (below)
Full roofline sweep	make roofline-measure (tools/roofline/, README there); pin make roofline-pin, roofline-regress; history tools/roofline/history/
imp vs llama.cpp	make bench-competitive (scripts/bench_competitive.sh, writes /tmp/bench_competitive.tsv)
Hero scoreboard	bash scripts/scoreboard.sh

Phantom: gemma-3-12b --bench prints bogus tok/s (#514); trust its PPL only.

Aggregate throughput (the serving regime)

• N-stream burst vs a fresh imp-server per arm; aggregate = sum completion_tokens / wall per wave; median over waves, 3+ alternating trials, all pairs same sign.
• conc_client.py PORT CONC WAVES [TAG]: unique prompts per stream, 300-token greedy. Harness baseline ~1714-1780 tok/s on 40-token prompts vs ~1030-1050 published on 130-token prompts; never mix.
• Pin runtime.max_batch_size, runtime.max_seq_len, kv_cache.max_blocks in BOTH arms (free-VRAM swing ~1.6 GB re-resolves auto batch: 5 -> 28 once moved 224.68 -> 630.19 tok/s with no code change).
• MEDIAN 0.0 = broken client (exceptions become (0, elapsed), model name hardcoded), not a regression. Read stderr and per-wave token counts.
• Discard wave 1 (graph captures + ramp: 629 vs 954-991 steady).
• Default-ON knobs that define "same flags": runtime.prefill_batch, gdn.state_bf16, gemm.nvfp4_smallm (+_impl=2, _pair), speculative.mtp_adaptive_k, runtime.prefill_chunk_decode_cap (1024; 4096 = +~10% burst lever), runtime.graph_prewarm.
• Profile the right binary: make build tags only imp:test; imp:builder can be days old.
• Method behind #1750-#1793: two-image A/B, all pairs same sign. Prefill-bound bursts cannot see KV capacity (32x8k with 60-token completions: identical walls while the pool grew 3.2x); capacity levers need max_tokens >= 512.
• Per-kernel time inside replayed decode graphs: nsys --cuda-graph-trace=node (#856). On a server: no --delay/--duration (flush hangs on cudaProfilerStop, SIGKILL loses all); capture fully, docker stop gracefully, filter the window by timestamp.

Step 1: cudaEvent (quick A/B)

N_ITER >= 100 after >1 s busy warmup (STOP #3); sample clocks during the run.

Step 2: ncu

ncu is not in the runtime image; mount the host install (/opt/nvidia/nsight-compute/2026.2.1, also 2026.2.0, 2025.4.0) or build tools/Dockerfile.ncu:

docker run --rm --gpus all -v $HOME/models:/models \
  -v /opt/nvidia/nsight-compute/2026.2.1:/ncu -v /tmp/out:/out --user root \
  imp:test /ncu/ncu --kernel-name "regex:my_kernel.*" --launch-skip 3 --launch-count 10 \
  -o /out/profile imp-bench ...      # chmod 777 /tmp/out first

Canonical metric set: .claude/skills/benchmark-cuda/ncu-basic.sh "<kernel-regex>" <binary> [args].

ncu traps: --launch-skip above the launch count waits forever; ncu + CUDA graphs hangs in the async graph loop (always --no-cuda-graphs, as the roofline harness does); the "Available Kernels" list shows base names, template regexes do not match; --clock-control base does NOT stop idle downclock between replays (a launch at 0.31 GHz read 99.7% of roofline; tools/roofline/config.json ncu.clock_floor_ghz=1.2 now drops such launches as n_launches_dropped_clock); multi-pass replay dies on TMA kernels on this WSL2 driver (the roofline runs single-pass metric groups). A low memory-stall ratio next to low bandwidth means "not bandwidth-bound", not "too few warps".

Step 3: nsys

nsys profile --sample=none --cpuctxsw=none --backtrace=none -t cuda,nvtx \
    --stats=true --cuda-memory-usage=true -o timeline --force-overwrite=true imp-bench ...
nsys stats timeline.nsys-rep

• WSL2 needs sampling off or nsys hangs.
• nsys in imp:test (no toolchain image needed): -v /opt/nvidia/nsight-systems/2026.1.3:/nsys --entrypoint /nsys/target-linux-x64/nsys; export --type sqlite the same way; read the sqlite in python:3.13-slim (imp:test has no python/sqlite3). Scratchpad files mounted into the container need chmod 755 dir / 644 file (--prompt-file: cannot open).
• Phase split: kernels with start in [end - tg_ms - pp_ms, end - tg_ms) = prefill; end - tg_ms onward = decode (pp/tg walls from the same run's log).
• e2e delta larger than kernel-sum delta = host-bound phase: sum CUPTI_ACTIVITY_KIND_RUNTIME per API name and CUPTI_ACTIVITY_KIND_MEMCPY by bytes/srcKind (1 = pageable) in that window. Qwen3.8-Flash-Next pp13863: kernels 8.9 s of a 63.8 s wall, 837k pageable cudaMemcpyAsync = 28 s of API time (legacy host-expert prefill, moe_prefill=legacy_fallback in Resolved dispatch).
• Graphs hide captured kernels: --no-cuda-graphs for the mix, --cuda-graph-trace=node when graphs must stay on.
• In imp:toolchain the qdstrm -> nsys-rep conversion fails silently (libcap.so.2, libdw.so.1 missing): apt-get install libcap2 libdw1t64, then /opt/nvidia/nsight-compute/*/host/*/QdstrmImporter -i x.qdstrm -o x.nsys-rep.
• nsys prints template args as <(int)64, (int)256, ...>; grep accordingly.
• Class maps come from the nsys steady-state sum (n >= 100 launches, decode vs prefill split by per-launch duration from the sqlite), never from the roofline's 120-launch ncu window: that window showed no GDN scan class while the scan was 42% of the hybrid pp512 wall.
• On Qwen3.6-35B total_kernel_ms scatters 1203-1255 (MoE routing); judge by class sums plus pp.
• .nsys-rep/.sqlite and diagnostics.dump_hidden_dir dumps (3.6 GB per arm at 300 tokens) fill the 40 GB tmpfs; delete after reading.
• compute-sanitizer (make sanitize) does not work on WSL2 (WDDM).

Step 4: roofline

Peaks: 1792 GB/s DRAM, FP16 838 TFLOPS, FP8 1677, FP4 3354 (datasheet), L2 96 MB. Measured FP4 mma.sync = 2019 TOPS (~1/2 datasheet), f32-accumulate 1/4 rate; GeForce TF32 = 1/2 the FP16-fp32acc rate. Use measured peaks or every FP4 kernel reads falsely bad. Roofline cells (tools/roofline/config.json): q8-dense Qwen3-8B-Q8_0, nvfp4-dense Qwen3-14B-NVFP4, nvfp4-moe Qwen3-30B-A3B-NVFP4, q4k-moe, q4k-dense-hd256 gemma-3-12b, nvfp4-hybrid = Qwen3.6-35B (NOT Qwen3.8-27B). Pinned baseline tools/roofline/history/BASELINE (run 1d5b9230 pinned #1835).

Report template

Kernel: <name>, config: <block=X, grid=Y, smem=Z>
  Wall:        <us> us (N=<iters>, warmup >1s)
  DRAM:        <pct>% of 1792 GB/s
  SM:          <pct>% of peak
  Occup:       <pct>%
  TC util:     <pct>%
  Clocks live: <MHz SM>/<MHz mem>/<W>   (healthy: 2850/13801/~500)
  Bound by:    <memory|compute|latency|stalls>  reason: <top stall>
  vs baseline: <+-X%> on tg (decode), pairs <n/n same sign>

Publishing numbers

• tests/perf_baseline.json is the gate: read values THERE, never from a skill. Two gates: throughput (8%/8% since #1400) and peak VRAM (metrics.memory_mb.own_peak_mb, 10%). Refresh only for an intentional perf or VRAM move: make gen-perf-baseline on a healthy day (STOP #4), second cold-median run agreeing, and say so in the PR. The gate measures spec-OFF decode.
• Before bisecting a red gate: (0) VRAM-resident? At ~0 MiB free WDDM spills into host memory and every cudaMalloc still succeeds (28 GiB succeeds at 22.6 GiB reported free); bandwidth is the discriminator (#1103: 55 vs 391 tok/s). (1) Anything else on the GPU? Read nvidia-smi --query-gpu=memory.used,utilization.gpu, not the process list (16.4 GiB held with a blank --query-compute-apps read -5.5%; a forgotten server ~-12%; a killed docker run keeps its commitment for tens of minutes). (2) Can the diff reach the measured code (git diff --stat main -- src/ include/ tools/imp-cli/)? (3) Does a cold-median run reproduce it?
• docs/BENCHMARKS.md rows are SHA-anchored; update them and the README block in the same commit; scripts/check-release.sh gates release PRs. docs/PERF.md owns every number (docs-sync).

A published verdict expires when its path is fixed

Verdict as written	Re-measured	Between
"MTP loses: 84.7-85.8 vs ~88 tok/s"	+21.3% at k=1 (#1481)	ea547a53, 3 weeks
"token_recycling net-negative, -7%"	-0.27%, neutral (#1483)	same commit
"+21.3% at k=1" (LIMITATIONS) and "+15% k=2" (/health)	measured without template/think; on think traffic MTP was dead (#1796)	spec verdicts need a think arm
"GDN scan: must be FP32"	layout bug; BF16 +12.5% (#1776)	-
"QSA: attention is 0.3 % of the decode step, structural" (#2072)	computed, not measured: dense decode attention 81.6 us/layer, the select kernel was 128.4 us on one SM; fixed, QSA tg +11 % at 13.9k (#2074)	same day; a breakdown without a profile is a hypothesis

1. git log --oneline <PROV commit>..HEAD -- <files of the measured path> (all provenance blocks have perf-path commits behind them; only THIS path matters). Renamed TUs: engine_scheduler.cpp split into engine_prefill.cpp, engine_prefill_ragged.cpp, engine_decode_pipeline.cpp (#1782).
2. Re-run the harness instead of reasoning about the delta.
3. Check the level, not only the delta (token_recycling re-read at 156 tok/s where the original said 99.37; #1102 sat between). A verdict can also be right for the wrong mechanism: the MR row cost was register pressure, both obvious fixes refuted (#1482).

Red flags

• pp512 delta without a decode delta (MoE spread ~38% across starts).
• nvidia-smi process list as "GPU free" (Windows-side load invisible; 12.9 GiB at 96% util with no container, 2026-08-14).
• Cold single shot; cross-day delta without live clocks; baseline refresh on a depressed day.
• Measuring after a non-default CMake build (verify-fast does not rebuild; an IMP_ALLOC_INTERPOSE=ON image reproduced -3% four times).
• Wrong peak dtype; A/B with graphs ON only; multi-model back-to-back sweeps.
• Dense decode-kernel A/B without speculative.ngram=false (STOP #6).
• "perf-neutral" without a SASS diff (cuobjdump -sass; byte-identical SASS is proof, a bench is not).
• Gate figures typed into a PR body from memory: capture the output first, paste second.