Skip to content

perf(cuda/quant): accumulate qmv in float below Ampere for bf16 - #1557

Merged
inureyes merged 2 commits into
mainfrom
perf/issue-1539-qmv-float-accum-pre-ampere
Aug 31, 2026
Merged

inureyes merged 2 commits into
mainfrom
perf/issue-1539-qmv-float-accum-pre-ampere

Conversation

@inureyes

Copy link
Copy Markdown
Member

Summary

qmv, the CUDA quantized GEMV behind every single-stream decode step, chose its accumulator from the weight bit width alone: float at bits >= 8, the element type below that. For a bf16 checkpoint at 4 bits that is a bf16 accumulator, and no pre-Ampere part has a bf16 ALU to run it on, so every FMA in the k-loop ran as convert-to-float, fma, convert-back. Below Ampere the kernel now accumulates bf16 in float at every bit width. On a Tesla V100 that makes qmv 2.01x faster and qwen3.8-27B-4bit decode 1.87x faster end to end.

Route taken, and why

The __CUDA_ARCH__ guard the issue preferred, not a host-side dispatch. qmv is AOT-compiled: the host takes &qmv_kernel<...> as a plain function pointer (qmv.cu:617, :662) and launches it through add_kernel_node_raw (:790, :845), and the accumulator appears in neither the kernel's template signature nor its mangled name. nvcc emits one device pass per entry in MLX_CUDA_ARCHITECTURES, so a guard inside device code gives exactly per-architecture behavior inside a fat binary with nothing for the single host pass to disagree with. The fallback bool FloatAccum template parameter from #1537's capability probe was not needed.

One deliberate narrowing against the issue text, which says "regardless of bit width". The promotion is conditioned on the element type being bfloat16_t, not applied to every type below Ampere. sm_70 and sm_75 do have a native fp16 FMA at twice the fp32 rate, so promoting f16 accumulators there would spend throughput to buy precision, the opposite of the trade this makes for bf16, and f16 policy below Ampere is #1542's subject. Every checkpoint measured here is bf16, so this costs nothing measured and avoids pre-empting #1542. The narrowing is on element type; the "regardless of bit width" part is implemented exactly as written.

What changed

  • src/lib/mlx-cpp/patches/mlx/backend/cuda/quantized/qmm/qmv.cu: new qmv_accumulator<bits, T> trait selecting float below sm_80 for bfloat16_t and keeping upstream's bits >= 8 rule everywhere else, used by both qmv_kernel_impl and qmv_multirow_kernel_impl. gather_qmv_kernel inherits it through qmv_kernel_impl. The multirow kernel had to move with the single-row one: perf(cuda/quant): amortizing small-M quantized GEMM for batched decode on Blackwell #725 guarantees a multirow launch is bit-identical to the per-row launches it replaces.
  • src/lib/mlxcel-core/src/ffi_tests.rs: qmv_matches_qmm_across_bits_and_group_sizes, comparing qmv at M = 1 against qmm at M = 8 on identical operands across bits 4 and 8 and group sizes 32, 64 and 128.
  • docs/benchmark_results/qmv-float-accum-v100-2026-08-31.md: the full before and after record.
  • docs/benchmark_results/volta-sm70-baseline-2026-08-31.md: the #1539 row of the post-program comparison table is filled in, and the "prefer 8-bit on Volta" guidance is re-measured, which that section had asked for.
  • CHANGELOG.md: entry recording the pre-Ampere output change explicitly.

Measured

Every number follows the six methodology rules of the baseline record: decode rate is a slope over -n 40 to -n 120 at a fixed prompt with generated_tokens == n asserted, warm PTX cache, --cuda-graph-trace=node, profiles reconciled against unprofiled wall clock, and absolute kernel times compared only at matching instance counts. The before arm reproduces the committed baseline to within 1.2% on every checkpoint and to four significant figures on qmv kernel time.

Checkpoint Before After Speedup
qwen3.8-27B-4bit 219.96 ms/tok (4.55 tok/s) 117.83 ms/tok (8.49 tok/s) 1.87x
gemma-4-12B-it-4bit 124.26 ms/tok (8.05 tok/s) 71.66 ms/tok (13.95 tok/s) 1.73x
gemma-4-12b-it-8bit (control) 66.75 ms/tok 66.73 ms/tok 1.00x
gemma-4-26b-a4b-it-4bit (MoE) 35.45 ms/tok 29.56 ms/tok 1.20x, ordinal only

qmv in nsys at identical launch counts: 5.2165 to 2.5984 s over 11,431 instances on the qwen run (2.01x) and 12.8369 to 6.5654 s over 39,151 on the dense run (1.96x). Reconciliation 101.3% to 101.9% and 100.7% to 103.3%. Attribution closes: the qmv delta is 52.26 ms/token against a measured slope delta of 52.60 on the dense pair, that is 99.4%. qmm_naive, which already accumulated in float and cannot be reached by the guard, moves by -0.12% and +0.55%, and the 8-bit checkpoint is unmoved inside its own repeat spread. Those two controls are what confine the result to the accumulator.

The issue's 25%-of-roofline acceptance criterion is not met, and that is reported rather than worked around. qmv on qwen3.8-27B-4bit goes from 7.37% to 14.80% of the 900 GB/s roofline. Reaching 25% would need another 1.69x on top. The criterion was written before the controlled measurement existed; that measurement put the ceiling at 2.14x and said in terms it was an upper bound, and 2.01x is 94% of it. The two figures are not consistent with each other: the 8-bit arm itself only reaches 26.7% of roofline while reading 1.9x the bytes, so no accumulator change alone could put a 4-bit arm at the same fraction. The remaining headroom belongs to the later items in #1536.

The 4-bit against 8-bit inversion the baseline record found is now nearly closed: the 8-bit arm's advantage falls from 1.886x to 1.074x while still reading 1.9x the bytes.

Occupancy

cuobjdump -res-usage on the sm_70 object, before and after, for the <8, 16, 64> bf16 / uint4b instantiation the profile runs:

Kernel REG Blocks/SM Warps/SM
qmv_kernel, has_residue_k = false 55 -> 61 4 -> 4 32 -> 32
qmv_kernel, has_residue_k = true 59 -> 66 4 -> 3 32 -> 24
qmv_multirow_kernel, max_x_rows = 8 110 -> 168 2 -> 1 16 -> 8

Nothing spills: STACK and LOCAL are zero across all 378 qmv-family instantiations in both arms. The residue-k variant is the one this workload runs (gemma-4-12B-it-4bit has k = 3840 and 3840 % 512 = 256), so the measured 1.96x already includes its lost block per SM. elems_per_thread was not re-tuned, because the issue conditions that on occupancy dropping enough to eat the win and 24 warps per SM did not come close. The multirow kernel halves its occupancy and is the one unquantified risk here: it serves batched decode with 2 to 8 rows, which no measurement in this PR exercises, it is the most weight-load-amortized member of the family, and its row window is already autotuner-tunable (#906).

Numerics

Float accumulation is strictly more accurate than the emulated bf16 accumulation it replaces, so quality improves on pre-Ampere parts. It also changes them: greedy decode on a Volta host is no longer token-identical to a build from before this change. No published Volta baseline exists for it to regress against, and the changelog entry states it explicitly.

Determinism run to run is unaffected: three repetitions of a greedy 32-token generation (--temp 0 --seed 0) give one distinct output on the before build and one on the after build, and the two builds produce the same tokens as each other on that prompt.

tests/cuda_qmm_determinism.rs does not pass on this host, and it did not pass before this change either. Pointed at gemma-4-12B-it-4bit (its default models/llama-3.2-1b-4bit is not present here, or in CI, and the test skips rather than fails when the directory is missing, so every CI run of it so far has been a skip) it diverges at step 0, the prefill. That was verified against an origin/main build with qmv.cu reverted: same test, same checkpoint, same divergence at the same step. Prefill is M = 64, so it dispatches to qmm_naive and never reaches qmv. The divergence is between iterations inside one process rather than between the two builds, and iteration 0's prefill logits hash to the same value (5078372879123904733) on both builds, so prefill output is bit-identical across this change. Quantized prefill on sm_70 not being bitwise reproducible is the same class of defect #910 fixed for qmm_sm80 on sm_121 and deserves its own issue; it is out of scope here and is written up in the benchmark record so it is not lost.

Test plan

  • Repeated greedy decode byte-identical run to run on the V100, three repetitions on the before build and three on the after build
  • MLXCEL_TEST_QMM_MODEL_DIR=models/mlx-community/gemma-4-12B-it-4bit cargo test --release --features cuda --test cuda_qmm_determinism on the V100: fails, and fails identically on origin/main with qmv.cu reverted, at the prefill step this change cannot reach. See the numerics section
  • cargo test --release --features cuda -p mlxcel-core --lib qmv_ on the V100, 3 passed, including qmv_multirow_matches_per_row_qmv_bitwise (perf(cuda/quant): amortizing small-M quantized GEMM for batched decode on Blackwell #725 bit-identity) and the new qmv_matches_qmm_across_bits_and_group_sizes
  • MLX_CUDA_ARCHITECTURES=70 make release-cuda, and cuobjdump --list-elf reports 96 cubins, every one sm_70
  • Decode slopes on four checkpoints, three repetitions each (five for MoE), before and after
  • nsys kernel attribution at matching instance counts, both runs reconciled against unprofiled wall clock
  • cuobjdump -res-usage occupancy comparison at sm_70
  • cuobjdump --dump-sass diff at sm_80 and sm_121, byte-identical before and after
  • cargo fmt -p mlxcel-core -- --check

Deferred to GB10

No sm_80-or-later device exists on this host, so three of the issue's acceptance criteria cannot be closed here. They belong to the ## GB10 (sm_121) continuation section of epic #1536.

  • Byte-identical greedy output on GB10 for a fixed prompt and seed. Not run: there is no sm_80+ part here.
  • GB10 throughput unmoved. Not run, same reason.
  • cargo test --features cuda green on an sm_121 host. Only the sm_70 half was run here.

The third criterion, that the emitted SASS for the sm_121 pass is unchanged, was verified and is not deferred. nvcc does not need the target part to emit code for it, so qmv.cu was compiled from both the unpatched and the patched source at one architecture with the production flag set and the disassembly diffed: cuobjdump --dump-sass is byte-identical at sm_80 (389,998 lines) and at sm_121 (437,519 lines), and cuobjdump -res-usage is identical at both. Identical SASS is strong evidence for the two runtime criteria above as well, since a device executing the same instructions on the same inputs produces the same outputs, but it is evidence and not the measurement, so they stay open.

Closes #1539

qmv is the CUDA quantized GEMV behind every single-stream decode step, and it chose its accumulator from the weight bit width alone: float at bits >= 8, the element type below that. For a bf16 checkpoint at 4 bits that is a bf16 accumulator, and no pre-Ampere part has a bf16 ALU to run it on, so every FMA in the k-loop ran as convert-to-float, fma, convert-back. Below Ampere the kernel now accumulates bf16 in float at every bit width. The float specializations it selects already existed and were exercised by the bits >= 8 path, so nothing new is instantiated.

The selection is a __CUDA_ARCH__ < 800 guard in device code rather than a host-side dispatch. qmv is AOT-compiled: the host takes &qmv_kernel<...> as a plain function pointer and launches it through add_kernel_node_raw, and the accumulator appears in neither the template signature nor the mangled name, so nvcc's per-architecture device pass gives exactly per-arch behavior inside a fat binary with nothing for the single host pass to disagree with. The multirow kernel moves with the single-row one because #725 guarantees a multirow launch is bit-identical to the per-row launches it replaces.

The promotion is conditioned on the element type being bfloat16_t rather than applied to every type below Ampere. sm_70 and sm_75 do have a native fp16 FMA at twice the fp32 rate, so promoting f16 there would spend throughput to buy precision, and f16 policy below Ampere belongs to #1542.

Measured on a Tesla V100-PCIE-32GB, following the six methodology rules of the committed sm_70 baseline: qwen3.8-27B-4bit decodes at 117.83 ms/token against 219.96 (4.55 to 8.49 tok/s, 1.87x) and gemma-4-12B-it-4bit at 71.66 against 124.26 (8.05 to 13.95 tok/s, 1.73x). In nsys at identical launch counts qmv itself runs 2.01x faster and accounts for 99.4% of the end-to-end gain, while qmm_naive, which already accumulated in float, moves by less than 0.6% and the 8-bit sibling checkpoint is unmoved inside its own repeat spread. qmv roofline attainment rises from 7.37% to 14.80% of 900 GB/s, short of the 25% the issue asked for and at 94% of the 2.14x ceiling the controlled 4-bit against 8-bit measurement had put on it.

Nothing spills: STACK and LOCAL are zero across all 378 qmv-family instantiations at sm_70 in both arms. The residue-k single-row variant loses one block per SM (32 to 24 warps) and is the variant this workload runs, so the measured 1.96x already includes that cost; elems_per_thread was not re-tuned.

Pre-Ampere output changes. Float accumulation is strictly more accurate than the emulated bf16 accumulation it replaces, so quality improves, but greedy decode on a Volta host is no longer token-identical to a build from before this change. No published Volta baseline exists for it to regress against. Repeated greedy decode stays byte-identical run to run. sm_80 and later are untouched: cuobjdump --dump-sass is byte-identical before and after for the sm_80 and sm_121 device passes.

Separately, tests/cuda_qmm_determinism.rs fails on this host both before and after, at the prefill step, which is M = 64 and therefore dispatches to qmm_naive rather than qmv. That was confirmed against an origin/main build with qmv.cu reverted, and iteration 0's prefill logits hash identically on both builds. Quantized prefill on sm_70 not being bitwise reproducible is a pre-existing defect outside this change's reach and is written up in the benchmark record.

Adds qmv_matches_qmm_across_bits_and_group_sizes, comparing qmv at M = 1 against qmm at M = 8 on identical operands across bits 4 and 8 and group sizes 32, 64 and 128. Full measurement record in docs/benchmark_results/qmv-float-accum-v100-2026-08-31.md, and the post-program comparison row of the sm_70 baseline is filled in.

Refs #1536

Closes #1539
@inureyes inureyes added status:review Under review type:performance Performance improvements priority:high High priority area:core mlxcel-core: MLX FFI, primitives, KV cache, layers platform:linux Linux (CUDA / packaging) specific labels Aug 31, 2026
@inureyes inureyes added status:done Completed and removed status:review Under review labels Aug 31, 2026
@inureyes
inureyes merged commit 8b4a25c into main Aug 31, 2026
13 checks passed
inureyes added a commit that referenced this pull request Sep 1, 2026
)

## Summary

`dispatch_cutlass_arch` in `grouped_gemm_unaligned.cu` mapped every device below compute capability 8.0 to `cutlass::arch::Sm75`, which names the `m16n8k8` MMA Turing introduced and an sm_70 part does not have. The pre-Ampere arm now selects `cutlass::arch::Sm70`, and the `Sm75` placeholder that initialised `fun` in `get_grouped_mm_funcion` is gone.

#1544 was written verification-first and said a negative result would be a complete answer. The answer is not negative, and not the positive one either.

## The branch is live, on the checkpoint this epic benchmarks

The issue expected `grouped_gemm_v2` to be unreachable because mlxcel's fused MoE path bypasses it. That is right about decode and wrong about prefill. #629's sorted-MoE prefill fast path (`quantized.cpp:359`) routes a quantized `GatherQMM` into `cutlass_grouped_gemm_unaligned` once `B >= min_rows * num_experts`; for `gemma-4-26b-a4b-it-4bit` that is `E = 128` and `top_k = 8`, so the gate is a 128-token prompt. Three nsys runs, same binary, same checkpoint:

| Run | Prompt | `B` | `cutlass::Kernel<GemmGrouped>` | `prepare_grouped_mm_data` |
|---|---|---|---|---|
| default | 573 tokens | 4584 | **180 inst, 853.8 ms, 3.8% of GPU time** | 180 inst |
| `MLXCEL_GATHER_QMM_GROUPED=0` | 573 tokens | 4584 | absent | absent |
| default | 46 tokens | 368 | absent | absent |

The 46-token run reproduces the #1538 baseline's MoE profile exactly at 326 `qmm_naive` instances, which is the cross-check that it is the same measurement the baseline made. That baseline profiled MoE below the gate, which is why no CUTLASS kernel appears in it; its MoE section is annotated here so the omission does not keep reading as absence.

Also reachable, without any gate, by every non-quantized MoE checkpoint through `SwitchLinear::Regular::forward` (`gpt_oss`, `deepseek`, `qwen3_next`, `jamba`, `nemotron_h`, `exaone_moe` and the rest).

## Nothing was computing wrong, and that was measured

The pre-Ampere arm resolves to `GemmConfiguration`'s primary template, which is `OpClassSimt` with `InstructionShape<1, 1, 1>`. Both tensor-core specializations are constrained on `Arch::kMinComputeCapability >= 80`, so no pre-Ampere tag can select an MMA atom of any shape, with or without `MLX_ENABLE_TF32`. CUTLASS therefore erases the tag, and the kernel that actually ran says so in its own name: `MmaPipelined` / `MmaSimt` / `OpMultiplyAdd`, with no `Sm70`, `Sm75` or `Sm80` token anywhere in it.

Correctness was checked rather than inferred, because a mismatched CUTLASS path can return plausible numbers that are wrong and greedy decoding will not reveal it. Two independent checks: the grouped path and the legacy `qmm_naive` reference path produce a **byte-identical 1,793-character greedy continuation** on a 285-token prompt, and this holds for all four combinations of {before, after} x {grouped, legacy}; and new unit tests compare `gather_mm` against an `f64` dense per-expert reference on the model's real expert dims. The pre-existing `test_gather_mm` asserted the output shape and never looked at a value, which is how this went unchecked.

## The retag moves no device code, on any architecture

Compiling the translation unit before and after with the production flags, and comparing `cuobjdump --dump-sass` per symbol:

| Target | Symbols before | Symbols after | Only in one | Bodies differing | SASS compared |
|---|---|---|---|---|---|
| `compute_70` | 51 | 51 | 0 | **0** | 58,211,476 bytes |
| `compute_80` | 51 | 51 | 0 | **0** | 35,077,812 bytes |
| `compute_121` | 51 | 51 | 0 | **0** | 50,952,400 bytes |

This is #1539's technique, and unlike #1541's case it transfers: this file compiles its CUTLASS kernels ahead of time, so its object holds real device code and an identical dump is a result rather than a tautology. Whole-file dumps differ only in cubin emission order, so the comparison is per symbol.

A control settled the design against giving Turing its own arm: adding an `Sm70` arm *alongside* `Sm75` emits the identical 51 device symbols and a byte-identical 493,538-line dump, for 26 extra host-side instantiations and 194,704 more bytes of object. The pre-Ampere arm therefore stays one arm, and its tag names the floor of the range it covers.

## What changed

- `patches/mlx/backend/cuda/gemms/grouped_gemm_arch.h` (new): the architecture decision as a pure `constexpr` function of the compute capability major version, with the evidence for why one arm covers Volta and Turing and what has to change before that stops holding.
- `patches/mlx/backend/cuda/gemms/grouped_gemm_unaligned.cu`: `dispatch_cutlass_arch` switches on that function; `get_grouped_mm_funcion` returns a named `GroupedGemmFn` starting from `nullptr` with a loud guard instead of an `Sm75` placeholder instantiation; two `static_assert`s pin the pre-Ampere configuration as SIMT (including under `kEnableTF32`) and its stage count at 2.
- `src/lib/mlxcel-core/cpp/grouped_gemm_arch_probe.cpp` (new) and `build.rs`: C shim over the shipped decision, compiled unconditionally rather than behind `cuda`, following the #1541 precedent.
- `src/lib/mlxcel-core/src/grouped_gemm_arch_tests.rs` (new): enumerates the tag mapping over every compute capability, on the shipped function, with no GPU.
- `src/lib/mlxcel-core/src/grouped_gemm_numeric_tests.rs` (new): `gather_mm` against an `f64` dense per-expert reference at `k = 2816, n = 704` and its transpose, across both entry points, both `kAlignmentC` arms, both operand layouts and f32/bf16/f16, plus a constant-per-expert case that catches a mis-gathered index.
- `docs/benchmark_results/grouped-gemm-arch-v100-2026-08-31.md` (new), the #1538 baseline's reserved post-program row and its MoE section, and a CHANGELOG entry.
- `TECHNICAL_REPORTS/1562-grouped-gemm-arch-tag-sm70-20260901.{en,ko}.md` (new), the bilingual report.

## `kStages` and `cp.async`

The issue asked whether `kStages = 3; // use SM80_CP_ASYNC` can reach a pre-Ampere arch tag. It cannot, structurally: that member belongs to `GemmConfiguration<float, cutlass::arch::Sm80, kAlignmentC, true>`, an explicit full specialization on `Sm80` that no other tag can name. The pre-Ampere arm gets `kStages = 2` and runs `MmaPipelined`, which the profiled kernel name confirms. Now asserted at compile time rather than left as reasoning.

## The `cuFuncSetAttribute` gap handed over from #1541: recorded, not fixed

#1541 handed this issue the missing dynamic shared-memory opt-in in `gather_gemm.cu`. Two findings, and the conclusion is to record.

The gap is in the shared encoder rather than one launch site: `CommandEncoder::add_kernel_node_raw` sets `sharedMemBytes` and calls `cudaGraphAddKernelNode` without ever calling `cudaFuncSetAttribute`, so every launch site must opt in itself. But the reachable Volta configuration is nowhere near the ceiling, measured from the profile rather than computed: the grouped GEMM asks for **10,320 bytes** of dynamic shared memory against the 49,152-byte non-opt-in limit, 21% of it. `gather_gemm.cu` also remains unreachable, which reproduces here (`nm -C` finds zero undefined references to `mlx::core::gather_mm(bool, bool, ...)` in the built `libmlx.a`, while all three `grouped_gemm_unaligned.cu` entry points carry one from `matmul.cpp.o`).

What is genuinely open and new: the sm_80 tensor-core configurations in this same file use much larger tiles and, in the tf32 arm, three stages. That is the plausible place for this path to trip the ceiling, and it needs an Ampere-or-later part to answer. Left as a follow-up rather than fixed blind.

## MoE decode and prefill on Volta, before and after

Five repetitions each, warm PTX cache, `nvidia-smi --query-compute-apps` asserted empty before every run, decode as a slope over `-n 40` to `-n 120` per the baseline's rule 1, prefill at `-n 1` on the 285-token prompt that crosses the grouped-GEMM gate.

| Checkpoint | Metric | Before | After | Delta | Before-arm spread |
|---|---|---|---|---|---|
| `gemma-4-26b-a4b-it-4bit` | decode | 29.66 ms/tok (33.72 tok/s) | 29.73 ms/tok (33.63 tok/s) | +0.26% | 10.42% |
| `gemma-4-26b-a4b-it-8bit` | decode | 31.68 ms/tok (31.56 tok/s) | 32.71 ms/tok (30.57 tok/s) | +3.24% | 11.11% |
| `gemma-4-26b-a4b-it-4bit` | prefill | 12,720 ms (22.41 tok/s) | 12,644 ms (22.54 tok/s) | -0.59% | 0.75% |
| `gemma-4-26b-a4b-it-8bit` | prefill | 15,212 ms (18.73 tok/s) | 15,606 ms (18.26 tok/s) | +2.59% | 14.15% |

Every delta is smaller than the before arm's own repetition spread, which is the only outcome byte-identical device code permits. The 8-bit prefill before cell's 14.15% spread is one cold first repetition that is left in rather than dropped; comparing repetitions 2 to 5 on both sides gives 15,633 ms against 15,606 ms, a -0.17% delta.

The before column is re-measured on this host from this worktree rather than quoted from #1538, whose 38.61 and 34.29 ms/token predate #1539. It also inverts that record's 4-bit against 8-bit MoE finding, in the same direction #1539's own record reports for the dense pair.

## Deferred to GB10

No sm_80-or-later device exists on this host, so the following are left unticked on #1544 rather than claimed. See epic #1536's `## GB10 (sm_121) continuation`.

- **GB10 MoE output byte-identical**: not run. What is recovered locally is the mechanism rather than the measurement: the tag mapping is provably unchanged for every compute capability at or above 8, and the device code that mapping selects is byte-identical before and after at `compute_80` and `compute_121`, per symbol. That leaves no route by which GB10 output could move, but identity is a measurement and this is an argument.
- **GB10 MoE throughput unmoved**: same, and for the same reason.
- **`cargo test --features cuda` green on sm_121**: needs a GB10 host. On sm_70 the `mlxcel-core` library suite is 1672 passed against 1 pre-existing failure that this change is shown not to cause; see the test plan.
- **The sm_80 tensor-core grouped-GEMM configurations against the 48 KB dynamic shared-memory ceiling**: newly raised here, genuinely open, and worth its own issue.

Note that the `CUDA sm_70 compile` CI check is not evidence for any of this. CUDA 13 removed Volta support and cannot compile `compute_70`, so that job passes in about 11 seconds by skipping; the local build is the only real coverage.

## Test plan

- [x] `MLX_CUDA_ARCHITECTURES=70 cargo build --release --features cuda`, both arms; `cuobjdump --list-elf libmlx.a` reports 96 cubins, every one sm_70.
- [x] `cargo test -p mlxcel-core --release --features cuda --lib -- --test-threads=1` on the V100: **1672 passed, 1 failed, 1 ignored** in 496 s. The one failure is `sampling::tests::temperature_one_support_unchanged`, and it is pre-existing on `c2e54939` rather than caused by this change. That was isolated rather than assumed: reverting `grouped_gemm_unaligned.cu` to its base version, rebuilding, and re-running the test reproduces a byte-identical failure. It asserts bit-exactness of `fused_sample_probs` at `T = 1.0` and fails by 1 ULP on 5 of 64 entries, the same class of sm_70 float-reduction non-determinism #1557 recorded for `tests/cuda_qmm_determinism.rs`. This branch touches no sampling code. Worth its own issue; flagged here rather than folded into this one. (`--test-threads=1` is what #1048 requires of the CUDA suite.)
- [x] `cargo clippy -p mlxcel-core --release --features cuda --lib --tests -- -D warnings` clean.
- [x] `cargo fmt -p mlxcel-core -- --check` clean.
- [x] Reachability: three nsys profiles with `--cuda-graph-trace=node`, table above.
- [x] Correctness: byte-identical greedy continuation, grouped against legacy, before against after, all four combinations.
- [x] Device-code delta: per-symbol SASS comparison at `compute_70`, `compute_80` and `compute_121`.
- [x] Throughput: five repetitions per cell, both MoE arms, decode slope and prefill, table above.

Closes #1544
@inureyes
inureyes deleted the perf/issue-1539-qmv-float-accum-pre-ampere branch October 6, 2026 12:12
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

area:core mlxcel-core: MLX FFI, primitives, KV cache, layers platform:linux Linux (CUDA / packaging) specific priority:high High priority status:done Completed type:performance Performance improvements

Projects

None yet

Development

Successfully merging this pull request may close these issues.

perf(cuda/quant): use float accumulators in qmv below Ampere at bits < 8

1 participant