Repository navigation
test(benchmark): Volta (sm_70) baseline record and build coverage #1538
Description
Activity
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
qmvaccumulatorThis issue lists the 4/8-bit pair as the way to settle the
qmvaccumulator inversion (#1539). It does not, on a MoE model. nsys--cuda-graph-trace=nodeon both arms:Arm qmm_naiveshareqmvshare4-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 makesM * B >= 8, soquantized.cpp:246-252routes it toqmm_naive, notqmv.qmvis 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-8bitis 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_naiveon Volta,cutlass::integer_subbyte<4>weights lose tounsigned charweights 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 charplanes too (12.0% ofqmm_naive, 2.2% ofqmv). 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_benchat:816) does one run atMAX_TOKENS=100and reportstokens / 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-modelsingle-run3.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-4bitthe 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 inbenchmarks/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": 4vs"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) qmvGPU time (39,151 inst both)12.84 s 5.99 s qmm_naiveGPU time (658 inst both)18.23 s 19.30 s The
qmvdelta (57.1 ms/token) accounts for the wall-clock slope delta (56.6 ms/token) at ~101%, andqmm_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
qmvin the MoE arms "cannot account for a 19.9% gap". That was inferred from percentage shares. In absolute terms the MoEqmvshows 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:
qmvfavors 8-bit by 9.4 ms/token whileqmm_naivefavors 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=nodeabsolute 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
qmvuses 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 allsm_70,libmlx.a155,355,188 bytes, onc7ef9e4e(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,
qmvat 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, soqmvaccounts for 97.6% of the gap.qmm_naivemoves the other way in the same profile (10.0733 s against 11.9215 s over 329 instances each), which is the control:qmvpicks its accumulator on bit width atqmv.cu:191and sm_70 has no bf16 ALU, whileqmm_naiveaccumulates 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-4bit220.33 ms/tok, 4.54 tok/s 65.4 GB/s, 7.27% of 900 GB/s gemma-4-12B-it-4bit124.41 ms/tok, 8.04 tok/s 49.6 GB/s, 5.51% gemma-4-12b-it-8bit65.96 ms/tok, 15.16 tok/s 176.6 GB/s, 19.63% gemma-4-26b-a4b-it-4bit38.61 ms/tok, 25.90 tok/s MoE, not reported gemma-4-26b-a4b-it-8bit34.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-compilein.github/workflows/ci.yml, runningcargo check --features cuda --all-targetsatMLX_CUDA_ARCHITECTURES: "70"on the existing self-hosted CUDA runner, then asserting withcuobjdump --list-elfthat the emittedlibmlx.aholdssm_70and nothing else. Path-filtered on a newcuda_archfilter oversrc/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 targetssm_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 forsm_70but 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_70to 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 fromM * B >= 8reachingGatherQMM. On this build the decode expert path is mlxcel's own fusedcustom_kernel_moe_gateup_kernel_cu_*andcustom_kernel_moe_down_kernel_cu_*at 3,570 instances each (30 layers over 119 steps), whileqmm_naiveappears with a prefill-sized 326 instances in both arms. Theqmvinversion 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. Theqmm_naivehalf 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_naiveinstance 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 40and-n 120return 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 onmlxcel-serveronly), so the fix is a prompt that generates past the budget plus an assertion thatgenerated_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-absent90afallback is covered byenforce_cuda_arch_compatibility(), and the doc records it as a trap regardless.Deferred to GB10
Listed in full in the PR's
## Deferred to GB10section. In short: the methodology section has not been checked against a GB10 sweep, andcuda-sm70-compilehas never executed, since GitHub Actions cannot be run from this host. Nothing here was ticked without being verified on this machine.- MoE decode does not route to
Correction: MoE decode does not route to
qmm_naiveMy earlier comment on this issue stated that MoE decode goes through
GatherQMM, where the expert count makesM * B >= 8, soquantized.cpp:246-252routes it toqmm_naiverather thanqmv. 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 owngemma-4-26b-a4b-it-8bitprofile, which I had but misread:Kernel Instances GPU time custom_kernel_moe_gateup+custom_kernel_moe_down7,140 (3,570 each) 2.40 s qmv28,084 1.49 s qmm_naive652 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:
qmvat 28,084 and the fused MoE kernels at 3,570 each, the latter being exactly 30 layers times 119 steps.qmm_naiveat 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/_downkernels,qmvhandles the per-step non-expert projections, andqmm_naivehandles 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
qmvaccumulator, which is why the dense pair in #1539 was needed. It also does not touch the correction I posted on #1543, which correctly describedqmm_naiveas the prefill kernel.A caveat on the numbers above. #1556 measured
qmm_naiveinstance 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.- added 3 commits that reference this issue
on Aug 31, 2026
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,580pinMLX_CUDA_ARCHITECTURES: "121"..github/workflows/release.yml:764ships x86_64 as80;86;89;90a;100;120.Volta works today only because
build.rsauto-detects the host and falls back to90awhen it cannot — so a container withoutnvidia-smisilently 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
9a79573:sm_70).-nvalues 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 onqwen3.8-27B-4bit: decode 4.2 tok/s (239 ms/tok), prefill ~7.7 tok/s, TTFT ~13 s.nsys profile -t cuda,nvtx --cuda-graph-trace=node. The--cuda-graph-trace=nodeflag is not optional: without it MLX's graph-captured work is invisible and the report attributes ~100% of GPU time toevent_signal_kernel. Record the kernel table and thecuda_api_sumtable (the audit's graph-construction costs live there).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:qmvselects float accumulators only atbits >= 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.MLX_CUDA_ARCHITECTURES=70on 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.docs/benchmark_results/volta-sm70-baseline-2026-08-31.mdwith 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
--cuda-graph-trace=noderequired for nsys.make release-cudaon a host withoutnvidia-smiis 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.make bench-modelunderstates decode on this host, quantifies the error, and reports the slope alongside it (added by the first comment on this issue).Validation
References
90afallback:src/lib/mlxcel-core/build.rs:358-435..github/workflows/ci.yml:413,580;.github/workflows/release.yml:541,764.src/lib/mlx-cpp/patches/mlx/backend/cuda/quantized/qmm/qmv.cu:191.