test(benchmark): Volta (sm_70) baseline record and build coverage

Part of #1536. Phase 0. The measurement foundation for every other item in the program — each of #1539 through #1545 states its acceptance criteria as a delta against this baseline, so this lands first.

Analogous in role to #624 for the GB10 program (#623).

Context

#1536's diagnosis rests on a one-off audit run on a single V100. Before optimizing anything, that has to become a reproducible record with a committed methodology, otherwise every later "we improved X by N%" is unfalsifiable. There is also no CI or release coverage below sm_80 at all, so a Volta build can break at any time without anyone noticing:

.github/workflows/ci.yml:413,580 pin MLX_CUDA_ARCHITECTURES: "121".
.github/workflows/release.yml:764 ships x86_64 as 80;86;89;90a;100;120.

Volta works today only because build.rs auto-detects the host and falls back to 90a when it cannot — so a container without nvidia-smi silently produces a binary that cannot run on the host it was built for.

Scope

A committed benchmark record plus the minimum build coverage to keep it honest. No kernel changes.

Explicitly in scope: deciding whether sm_70 belongs in CI at all. A no-build-coverage outcome is acceptable if documented, but a Volta build that silently rots is not.

Implementation plan

Reproduce the baseline from a clean build on Tesla V100-PCIE-32GB (sm_70, 900 GB/s HBM2, 14 TFLOPS FP32, 112 TFLOPS FP16 tensor), driver 575.51.03, CUDA 12.9.41, MLX pin 9a79573:
```
MLX_CUDA_ARCHITECTURES=70 make release-cuda
cuobjdump --list-elf target/release/build/mlxcel-core-*/out/build/lib/libmlx.a | grep -oE 'sm_[0-9]+a?' | sort | uniq -c
```
Record the cubin arch histogram as build provenance (the audit saw 96 cubins, all sm_70).
Throughput. Decode rate as the slope of two -n values at a fixed short prompt, so the fixed first-token cost does not contaminate it; prefill as the delta between a ~3-token and a ~600-token prompt at equal -n. Audit numbers to reproduce on qwen3.8-27B-4bit: decode 4.2 tok/s (239 ms/tok), prefill ~7.7 tok/s, TTFT ~13 s.
Roofline attainment, which is the number the program is actually steering: bytes moved per decode step over measured decode time against 900 GB/s (audit: ~70 GB/s, 7.7%), and achieved prefill TFLOPS against 14 TFLOPS FP32 (audit: ~0.41, ~3%).
Kernel profile via nsys profile -t cuda,nvtx --cuda-graph-trace=node. The --cuda-graph-trace=node flag is not optional: without it MLX's graph-captured work is invisible and the report attributes ~100% of GPU time to event_signal_kernel. Record the kernel table and the cuda_api_sum table (the audit's graph-construction costs live there).
Model coverage. At minimum one dense-hybrid 4-bit (qwen3.8-27B-4bit), one MoE 4-bit (gemma-4-26b-a4b-it-4bit, which also exercises the grouped-GEMM path in fix(cuda/moe): grouped GEMM selects cutlass::arch::Sm75 on an sm_70 part #1544), and one 8-bit checkpoint. The 4-bit/8-bit pair is diagnostic, not incidental: qmv selects float accumulators only at bits >= 8 (perf(cuda/quant): use float accumulators in qmv below Ampere at bits < 8 #1539), so an 8-bit checkpoint may decode faster per byte moved than a 4-bit one on Volta. Confirming or refuting that inversion is a deliverable of this issue.
Build coverage. Add sm_70 to a compile-only CI job (no GPU needed — MLX_CUDA_ARCHITECTURES=70 on the existing CUDA build runner catches every arch-conditional compile break, which is the realistic failure mode). Decide and document whether a GPU-backed Volta job is worth it; a compile-only gate plus this document is an acceptable answer.
Doc. docs/benchmark_results/volta-sm70-baseline-2026-08-31.md with the tables, exact commands, host provenance, and the raw nsys reports referenced. Include a post-program comparison table left empty, to be filled as perf(cuda/quant): use float accumulators in qmv below Ampere at bits < 8 #1539-perf(cuda): CUDA graph instantiation and JIT module load dominate Volta TTFT #1545 land.

Acceptance criteria

Baseline doc committed, with every number reproduced from a clean build rather than copied from epic: NVIDIA Volta (sm_70) inference acceleration program #1536.
Methodology stated precisely enough that a second person on a different V100 gets within 10%: slope-based decode rate, prompt-delta prefill, --cuda-graph-trace=node required for nsys.
Roofline attainment reported as a percentage, not just tok/s, for both decode and prefill.
The 4-bit vs 8-bit decode comparison is reported with an explicit verdict on the accumulator-inversion hypothesis.
sm_70 compile coverage in CI, or a documented decision not to, in this issue and in the doc.
make release-cuda on a host without nvidia-smi is covered by the chore(core): expose CUDA compute capability to the mlxcel runtime, build, and diagnostics #1537 mismatch check or, if chore(core): expose CUDA compute capability to the mlxcel runtime, build, and diagnostics #1537 has not landed, called out in the doc as a known trap.
The accumulator deliverable is settled on a dense controlled pair, not on a MoE checkpoint (added by test(benchmark): Volta (sm_70) baseline record and build coverage #1538 (comment) and closed by test(benchmark): Volta (sm_70) baseline record and build coverage #1538 (comment)).
Every nsys absolute time is reconciled against an unprofiled wall-clock measurement, and the check is reported, including where it fails (added by the second comment on this issue).
The document states that make bench-model understates decode on this host, quantifies the error, and reports the slope alongside it (added by the first comment on this issue).

Validation

MLX_CUDA_ARCHITECTURES=70 make release-cuda
for n in 8 40; do ./target/release/mlxcel generate -m ./models/qwen3.8-27B-4bit -p "Hi." -n $n; done
nsys profile -t cuda,nvtx --cuda-graph-trace=node -o volta_base \
  ./target/release/mlxcel generate -m ./models/qwen3.8-27B-4bit -p "Explain what a GPU tensor core does." -n 24
nsys stats --report cuda_gpu_kern_sum --report cuda_api_sum --format csv volta_base.nsys-rep

References

Precedent for a Phase 0 measurement issue: perf(bench): long-prompt prefill benchmark coverage and serving TTFT/decode-rate telemetry #624.
Arch auto-detect and the 90a fallback: src/lib/mlxcel-core/build.rs:358-435.
CI/release arch pins: .github/workflows/ci.yml:413,580; .github/workflows/release.yml:541,764.
Accumulator selection that motivates the 4-bit/8-bit comparison: src/lib/mlx-cpp/patches/mlx/backend/cuda/quantized/qmm/qmv.cu:191.

Measured 2026-08-31: the 4-bit / 8-bit comparison, and a correction to its design

Ran the comparison this issue asks for. The headline prediction holds, but the experimental design in the issue body is wrong for a MoE checkpoint and needs amending.

Result

mlx-community/gemma-4-26b-a4b-it-{4bit,8bit} (Gemma 4 MoE, 26B total / A4B active) on Tesla V100-PCIE-32GB. Identical prompt, --seed 0, decode rate as a slope over -n 40 → -n 120, three repetitions:

Rep	4-bit ms/tok	8-bit ms/tok
1	38.13	31.38
2	37.75	31.50
3	37.00	31.25
mean	37.63 (26.6 tok/s)	31.38 (31.9 tok/s)

8-bit decode is 19.9% faster than 4-bit, reproducibly (under 3% spread within each arm), while carrying 27 GB of weights against 15 GB. The arm moving 1.8x the bytes wins on a 900 GB/s part — independent confirmation that Volta decode is not bandwidth-bound.

Correction: this does not test the `qmv` accumulator

This issue lists the 4/8-bit pair as the way to settle the qmv accumulator inversion (#1539). It does not, on a MoE model. nsys --cuda-graph-trace=node on both arms:

Arm	`qmm_naive` share	`qmv` share
4-bit	44.2 + 16.9 + 12.0 = 73.1%	8.3 + 2.4 + 1.6 + 0.6 = 12.9%
8-bit	49.4 + 26.9 = 76.3%	6.2 + 0.6 = 6.8%

MoE decode goes through GatherQMM, where the expert count makes M * B >= 8, so quantized.cpp:246-252 routes it to qmm_naive, not qmv. qmv is a minor share in both arms and cannot account for a 19.9% gap.

Amendment to this issue's scope: the accumulator test requires a dense checkpoint at matched 4-bit and 8-bit, where decode is genuinely GEMV-shaped. No such pair is available locally (gemma-4-12b-it-8bit is dense but has no 4-bit sibling here). Acquiring one is now a prerequisite for that deliverable.

What was measured instead is a separate finding worth carrying forward: in qmm_naive on Volta, cutlass::integer_subbyte<4> weights lose to unsigned char weights by more than a 1.8x bandwidth advantage. The sub-byte unpack in the inner loop costs more than the traffic it saves — additional evidence for #1543 and a consideration for #1541.

Confound: the 4-bit arm is mixed precision, carrying unsigned char planes too (12.0% of qmm_naive, 2.2% of qmv). The arms differ by less than a full bit-width step, which if anything understates a pure comparison.

Methodology finding for this issue

make bench-model (Makefile:879-896, run_bench at :816) does one run at MAX_TOKENS=100 and reports tokens / wall_time. On GB10 that approximates a decode rate; on Volta it does not, because the fixed first-token cost is ~9-13 s:

Measurement	4-bit	8-bit	Delta
`make bench-model` single-run	3.31 tok/s	3.60 tok/s	+8.8%
Slope (this comment)	26.6 tok/s	31.9 tok/s	+19.9%

On qwen3.8-27B-4bit the harness reports 1.70 tok/s against a true decode rate of 4.2 tok/s — a 2.5x understatement. The harness stays the right artifact producer (it writes the CSV shape the GB10 sweeps in benchmarks/ use), but the Volta baseline must report the slope alongside it, or the recorded numbers will not mean what a reader assumes.

Deliverable closed: the dense controlled pair, and a correction to my MoE reading

Follow-up to #1538 (comment). The dense checkpoint pair that comment named as a prerequisite has been acquired and measured, so the accumulator deliverable of this issue is now satisfied. Two corrections to that earlier comment are also needed.

Deliverable: dense 4-bit vs 8-bit

mlx-community/gemma-4-12B-it-{4bit,8bit} — 33 of 34 config keys identical, sole difference "bits": 4 vs "bits": 8. Full result in #1539 (comment). Headline:

	4-bit	8-bit
Decode (slope, 3 reps)	122.04 ms/tok (8.19 tok/s)	65.46 ms/tok (15.28 tok/s)
`qmv` GPU time (39,151 inst both)	12.84 s	5.99 s
`qmm_naive` GPU time (658 inst both)	18.23 s	19.30 s

The qmv delta (57.1 ms/token) accounts for the wall-clock slope delta (56.6 ms/token) at ~101%, and qmm_naive — which accumulates in float regardless of bit width — moves the other way. The accumulator is isolated.

Correction 1: the MoE pair did partly test it after all

My earlier comment said qmv in the MoE arms "cannot account for a 19.9% gap". That was inferred from percentage shares. In absolute terms the MoE qmv shows the same inversion (28,084 inst both arms: 2.62 s at 4-bit vs 1.49 s at 8-bit, 1.76x). The direction was right in the dense pair and was already present in the MoE data; I under-read it.

Correction 2: do not trust cross-run nsys absolutes on the MoE pair

The MoE profiles do not reconcile with wall clock. Summed GPU time is 20.20 s and 21.91 s for runs that take ~13-14 s unprofiled, and the per-kernel deltas have the wrong sign against the measured slope: qmv favors 8-bit by 9.4 ms/token while qmm_naive favors 4-bit by 16.4 ms/token, netting to 8-bit being slower, which contradicts the measured 19.9% 8-bit win. Graph-node tracing inflates unevenly per kernel type.

The dense pair does reconcile (attribution ~101%), which is why it is the result of record. Methodology note for this issue's baseline doc: --cuda-graph-trace=node absolute times are only trustworthy when they reconcile against an unprofiled wall-clock measurement — check that before drawing kernel-level conclusions, and report the check. The MoE 19.9% wall-clock number stands as a black-box observation; its kernel-level breakdown does not.

Net effect on the guidance this issue produces

Prefer 8-bit over 4-bit on Volta where it fits in 32 GB — confirmed on both a dense (1.86x) and a MoE (1.20x) checkpoint. This is a workaround for #1539, not a substitute; once qmv uses float accumulators below Ampere the 4-bit arm should regain its bandwidth advantage.

Baseline recorded: PR #1556

Doc: docs/benchmark_results/volta-sm70-baseline-2026-08-31.md. Raw rows: benchmarks/cuda_v100_2026-08-31.csv. Build: MLX_CUDA_ARCHITECTURES=70 make release-cuda, 96 cubins all sm_70, libmlx.a 155,355,188 bytes, on c7ef9e4e (after #1551).

Verdict on the accumulator inversion: confirmed

8-bit decodes faster than 4-bit on Volta, 1.886x on the dense controlled pair and 1.13x on the MoE pair. On the dense pair the mechanism is isolated: identical instance counts in both arms, qmv at 12.8366 s (4-bit) against 5.9911 s (8-bit) over 39,151 instances each, a delta of 57.05 ms/token against a measured slope delta of 58.45 ms/token, so qmv accounts for 97.6% of the gap. qmm_naive moves the other way in the same profile (10.0733 s against 11.9215 s over 329 instances each), which is the control: qmv picks its accumulator on bit width at qmv.cu:191 and sm_70 has no bf16 ALU, while qmm_naive accumulates in float regardless. Both dense profiles reconcile against their unprofiled runs at 102.4% and 101.6%.

Baseline numbers

Checkpoint	Decode (slope)	Roofline
`qwen3.8-27B-4bit`	220.33 ms/tok, 4.54 tok/s	65.4 GB/s, 7.27% of 900 GB/s
`gemma-4-12B-it-4bit`	124.41 ms/tok, 8.04 tok/s	49.6 GB/s, 5.51%
`gemma-4-12b-it-8bit`	65.96 ms/tok, 15.16 tok/s	176.6 GB/s, 19.63%
`gemma-4-26b-a4b-it-4bit`	38.61 ms/tok, 25.90 tok/s	MoE, not reported
`gemma-4-26b-a4b-it-8bit`	34.29 ms/tok, 29.16 tok/s	MoE, not reported

Prefill on qwen3.8-27B-4bit, from a five-rung ladder at -n 1: 8.00 tok/s marginal, 0.397 TFLOPS, 2.83% of 14 TFLOPS FP32, with an 18.41 s fixed intercept. TTFT 24.94 s at a 54-token prompt, model loaded, warm cache.

Build coverage decision, per acceptance criterion 5

Added: cuda-sm70-compile in .github/workflows/ci.yml, running cargo check --features cuda --all-targets at MLX_CUDA_ARCHITECTURES: "70" on the existing self-hosted CUDA runner, then asserting with cuobjdump --list-elf that the emitted libmlx.a holds sm_70 and nothing else. Path-filtered on a new cuda_arch filter over src/lib/mlx-cpp/** and the two build scripts, with its own persistent target directory.

Declined, deliberately: a GPU-backed Volta job. The realistic failure mode this epic can introduce is a compile break (a __CUDA_ARCH__ guard not covering cc 7, a CUTLASS type with no pre-Ampere instantiation, a bf16 intrinsic with no sm_70 implementation), all of which the compile gate catches without a Volta card. There is no Volta runner in this repository's pool; no release artifact targets sm_70; #1537 already turns an architecture mismatch into a named startup error instead of an opaque CUDA load failure; and the one Volta machine that exists is this single-GPU development host, where such a job would serialize against the measurements themselves. The cost is explicit and stated in the doc: a change that compiles for sm_70 but computes the wrong answer or regresses throughput is caught by nothing automatic, and what covers it instead is this document plus a re-run of its reproduce commands.

Also declined: adding sm_70 to the release matrix. A cold six-architecture build is already about three hours, and shipping it would be a support commitment for a part these numbers show the current kernels serve badly. Revisit once #1539 to #1545 have moved them.

Two corrections to the measurements in the comments above

MoE decode does not route to qmm_naive. test(benchmark): Volta (sm_70) baseline record and build coverage #1538 (comment) inferred it does from M * B >= 8 reaching GatherQMM. On this build the decode expert path is mlxcel's own fused custom_kernel_moe_gateup_kernel_cu_* and custom_kernel_moe_down_kernel_cu_* at 3,570 instances each (30 layers over 119 steps), while qmm_naive appears with a prefill-sized 326 instances in both arms. The qmv inversion is present in the MoE profile at the same 1.76x that comment found (2.6224 s against 1.4921 s over 28,084 identical instances), but the fused expert kernels move the other way and nearly cancel it, which is why the MoE advantage is 1.13x rather than the dense pair's 1.89x. The qmm_naive half of that comment's "what was measured instead" finding is about prefill, not decode.
TTFT did not reproduce and is recorded as unexplained, not attributed to a cause. Measured 24.94 s against the ~13 s in the issue body. A cold PTX cache in the original is ruled out by direction, and every CUDA API cost in that profile is higher than this record's at the same call counts. Also qmm_naive instance counts came out exactly half the audit's on both models (497/994 and 329/658), which is a counting difference rather than a performance one.

One methodology rule the issue did not anticipate

Both slope runs must reach the token budget. A prompt like "Hi." makes every instruct checkpoint here emit EOS after roughly ten tokens, so -n 40 and -n 120 return the same generation and the slope becomes a difference of two nearly equal numbers over nearly zero. The first attempt at this record did exactly that and produced everything from a division by zero to 193 ms/token on a model whose real figure is 220. The CLI has no --ignore-eos (that flag is on mlxcel-server only), so the fix is a prompt that generates past the budget plus an assertion that generated_tokens == n. Added as a checkbox above and as rule 2 in the doc.

Also confirmed for criterion 6: #1537 has landed, so the nvidia-smi-absent 90a fallback is covered by enforce_cuda_arch_compatibility(), and the doc records it as a trap regardless.

Deferred to GB10

Listed in full in the PR's ## Deferred to GB10 section. In short: the methodology section has not been checked against a GB10 sweep, and cuda-sm70-compile has never executed, since GitHub Actions cannot be run from this host. Nothing here was ticked without being verified on this machine.

Correction: MoE decode does not route to `qmm_naive`

My earlier comment on this issue stated that MoE decode goes through GatherQMM, where the expert count makes M * B >= 8, so quantized.cpp:246-252 routes it to qmm_naive rather than qmv. That mechanism is wrong. #1556's independent profiling caught it, and re-reading my own saved profile confirms it.

The expert path during MoE decode is mlxcel's fused decode-MoE kernel, not qmm_naive. From the audit's own gemma-4-26b-a4b-it-8bit profile, which I had but misread:

Kernel	Instances	GPU time
`custom_kernel_moe_gateup` + `custom_kernel_moe_down`	7,140 (3,570 each)	2.40 s
`qmv`	28,084	1.49 s
`qmm_naive`	652	16.72 s

The instance counts settle it without needing the timings, which is fortunate because the MoE timings are not trustworthy (see below). A 120-token decode over a 30-layer model produces launch counts in the thousands: qmv at 28,084 and the fused MoE kernels at 3,570 each, the latter being exactly 30 layers times 119 steps. qmm_naive at 652 is a single-pass count. It is prefill, not decode.

So the correct statement is: MoE decode runs its experts through the fused custom_kernel_moe_gateup / _down kernels, qmv handles the per-step non-expert projections, and qmm_naive handles prefill.

What this does and does not change. The conclusion I drew from it still holds, for a better reason than the one I gave: the MoE pair was not a clean test of the qmv accumulator, which is why the dense pair in #1539 was needed. It also does not touch the correction I posted on #1543, which correctly described qmm_naive as the prefill kernel.

A caveat on the numbers above. #1556 measured qmm_naive instance counts at exactly half the audit's on both models (326 against 652 here, 329 against 658 on the dense pair), and attributes the factor of two to the audit's report counting a graph node and its kernel separately. Both profiles also fail the wall-clock reconciliation check on the MoE arms (the audit's sums to 20.20 s and 21.91 s for runs taking ~13-14 s unprofiled; #1556's inflates to 144.0% on the 4-bit arm). MoE kernel-level attribution is therefore not reportable from either profile. The routing conclusion above rests on the instance-count shape, which is robust to both problems.

test(benchmark): record the Volta sm_70 baseline and gate sm_70 in CI (…

fix(cuda/moe): tag the pre-Ampere grouped GEMM arm Sm70, not Sm75 (#1562

perf(cuda): explain the Volta fixed first-token cost, and instrument …

Navigation Menu

Repository navigation

test(benchmark): Volta (sm_70) baseline record and build coverage #1538

Context

Scope

Implementation plan

Acceptance criteria

Validation

References

Measured 2026-08-31: the 4-bit / 8-bit comparison, and a correction to its design

Result

Correction: this does not test the `qmv` accumulator

Methodology finding for this issue

Deliverable closed: the dense controlled pair, and a correction to my MoE reading

Deliverable: dense 4-bit vs 8-bit

Correction 1: the MoE pair did partly test it after all

Correction 2: do not trust cross-run nsys absolutes on the MoE pair

Net effect on the guidance this issue produces

Baseline recorded: PR #1556

Verdict on the accumulator inversion: confirmed

Baseline numbers

Build coverage decision, per acceptance criterion 5

Two corrections to the measurements in the comments above

One methodology rule the issue did not anticipate

Deferred to GB10

Correction: MoE decode does not route to `qmm_naive`

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

test(benchmark): Volta (sm_70) baseline record and build coverage #1538

Description

Context

Scope

Implementation plan

Acceptance criteria

Validation

References

Activity

inureyes commented on Aug 31, 2026

Measured 2026-08-31: the 4-bit / 8-bit comparison, and a correction to its design

Result

Correction: this does not test the qmv accumulator

Methodology finding for this issue

inureyes commented on Aug 31, 2026

Deliverable closed: the dense controlled pair, and a correction to my MoE reading

Deliverable: dense 4-bit vs 8-bit

Correction 1: the MoE pair did partly test it after all

Correction 2: do not trust cross-run nsys absolutes on the MoE pair

Net effect on the guidance this issue produces

inureyes commented on Aug 31, 2026

Baseline recorded: PR #1556

Verdict on the accumulator inversion: confirmed

Baseline numbers

Build coverage decision, per acceptance criterion 5

Two corrections to the measurements in the comments above

One methodology rule the issue did not anticipate

Deferred to GB10

inureyes commented on Aug 31, 2026

Correction: MoE decode does not route to qmm_naive

Metadata

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

Issue actions

Correction: this does not test the `qmv` accumulator

Correction: MoE decode does not route to `qmm_naive`