epic: NVIDIA Volta (sm_70) inference acceleration program

## Summary

Program of work to make the CUDA backend usable on NVIDIA **Volta (sm_70)**. The headline finding from a first-principles audit plus an nsys/measured baseline on a Tesla V100-PCIE-32GB: **mlxcel never issues a single tensor-core instruction on Volta.** Every quantized GEMM falls into a scalar-FMA CuTe path, and the activation dtype (bf16) is one that Volta has no hardware for at all — no ALU, no tensor core, no cuBLAS GEMM.

This is the pre-Ampere counterpart to #623 (which targeted GB10/sm_121 and treated sm_80 as the floor). #623 explicitly scoped Blackwell; nothing in that program looked below Ampere, and the arch matrix shipped in the release workflow (`80;86;89;90a;100;120` for x86_64) starts at sm_80. Volta is reachable today only because `build.rs` auto-detects the host compute capability.

## Baseline (measured)

Host: 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, mlxcel 0.6.0, MLX pin `9a79573`. Build: `make release-cuda` with `MLX_CUDA_ARCHITECTURES` auto-detected to `70` (verified: `cuobjdump --list-elf libmlx.a` reports 96 cubins, all `sm_70`).

Model: `qwen3.8-27B-4bit` (15 GB, affine 4-bit / group 64; all non-quantized tensors BF16; hybrid GDN + full attention every 4 layers, 64 layers, hidden 5120, head_dim 256).

| Phase | Measured | Roofline | Attained |
|---|---|---|---|
| Decode | **4.2 tok/s** (239 ms/tok) | ~60 tok/s (900 GB/s / 15 GB) | **7%** |
| Prefill | **~7.7 tok/s** (~600-token prompt in 78.2 s) | ~0.41 TFLOPS achieved vs 14 TFLOPS FP32 peak | **~3%** |
| TTFT (3-token prompt) | **~13 s** | — | — |

Decode rate is the slope of two runs (`-n 8` → 14.89 s, `-n 40`/31 generated → 20.39 s); the ~13 s intercept is the fixed first-token cost. Prefill is the delta between a 3-token and a ~600-token prompt at equal `-n`.

nsys kernel profile (`-t cuda,nvtx --cuda-graph-trace=node`, 24-token run):

| Kernel | GPU time | Instances | Avg |
|---|---|---|---|
| `qmm_naive_kernel<..., cutlass::bfloat16_t, integer_subbyte<4>, ..., tile 64x64x64>` | **82.0%** | 994 | 29.3 ms |
| `qmv_kernel<..., cutlass::bfloat16_t, integer_subbyte<4>, ...>` | **14.7%** | 11431 | 457 us |
| `event_signal_kernel` | 0.4% | 1847 | 74 us |
| `volta_sgemm_64x64_nn` (cuBLAS **FP32**, not tensor core) | 0.3% | 6192 | 19 us |
| `naive_grouped_unfold_transpose_nd<__nv_bfloat16>` (GDN causal conv im2col) | 0.3% | 1200 | 97 us |

**The decisive datum: `qmv` moves ~15 GB of weights per decode step in 218 ms = ~70 GB/s, which is 7.7% of the V100's 900 GB/s.** Decode on Volta is not bandwidth-bound. It is ALU-bound, which for a 4-bit GEMV is only explicable by the arithmetic itself being wrong for the hardware.

## Diagnosis

**1. No tensor cores, anywhere.** `mlx/backend/cuda/device/gemm_sm70.cuh:42-63`:

```cpp
template
inline constexpr auto make_tiled_mma(CtaTiler cta_tiler) {
using Atom = cuda::std::conditional_t>; // <- every cc < 8 lands here
```

Despite the file name, the non-SM80 branch is `UniversalFMA` — scalar FP32 FMA on CUDA cores. Both `qmm_naive` (`device/qmm_naive.cuh:63,233`) and `gather_gemm` (`gemms/gather_gemm.cu:54`) select it via `compute_capability_major() >= 8`. That is an 8x theoretical gap (14 vs 112 TFLOPS) before accounting for the measured 3%-of-FP32-peak.

**2. bf16 is the worst possible dtype on Volta.** sm_70 has no bf16 ALU, no bf16 tensor-core variant (`cute/arch/mma_sm70.hpp` provides only `SM70_8x8x4_F16F16F16F16_*` and `SM70_8x8x4_F32F16F16F32_*`), and no cuBLAS bf16 GEMM. `src/lib/mlx-cpp/patches-cuda/dtype.cpp` *deliberately preserves* bf16 through type promotion — correct and load-bearing on GB10 ("bf16 is the native compute type for LLM inference"), exactly inverted on Volta. Every `__nv_bfloat16` elementwise, norm, and binary kernel in the profile pays emulation.

**3. `qmv` accumulates in bf16 at 4 bits.** `patches/mlx/backend/cuda/quantized/qmm/qmv.cu:191`:

```cpp
cuda::std::conditional_t<(bits >= 8), float, T> sums[elems_per_thread] = {};
```

The float-accumulation specialization already exists (`qmv.cu:77`, `:145`) and is simply not selected for `bits < 8`. On sm_70 the `T` (bf16) branch is emulated per element. This is the 7.7%-of-bandwidth result.

**4. The prefill tile is halved on Volta for no reason.** `quantized/qmm/qmm_naive.cu:18`: `enough_smem = sm80 && itemsize <= 2 && group_size <= 64` forces `tile_n = 64` instead of 128. The 128-wide tile needs ~18-36 KB of shared memory; V100 offers 96 KB per block via `cudaFuncAttributeMaxDynamicSharedMemorySize`.

**5. MoE grouped GEMM claims Turing on a Volta part.** `patches/mlx/backend/cuda/gemms/grouped_gemm_unaligned.cu:340` maps `compute_capability_major() < 8` to `cutlass::arch::Sm75`, whose `m16n8k8` MMA does not exist on sm_70. A MoE checkpoint does generate output today, so this is a verification item rather than a confirmed defect.

**6. mlxcel itself is arch-blind.** `grep -rn compute_capability src/ --include=*.rs` returns nothing. Every architecture decision is delegated to MLX's C++ gates; mlxcel has no way to reason about, log, or override them. This is the structural gap underneath items 1-5.

**7. ~~Host-side graph construction dominates TTFT.~~ CORRECTED by #1545: the fixed first-token cost is lazy weight materialization, and it is not Volta-specific.** This diagnosis was wrong, and wrong in an instructive way. It read `cudaGraphInstantiate` at 2.14 s / 208 calls and `cudaGraphAddKernelNode` at 1.95 s / 72,539 calls off `cuda_api_sum` and concluded graph construction dominated. #1545 measured the phases directly and found weight materialization is 12.08 s of a 15.54 s fixed cost (77.8%): host staging of 15.13 GB at 7.40 s plus a pageable H2D copy at 4.19 s. Graph construction is 6.2% and kernel execution 4.8%. **The dominant phase is invisible to `cuda_api_sum` entirely**, because host-side file reads and buffer staging issue no CUDA calls, which is exactly why an API-time profile could not close the gap. There is also no graph thrashing: 196 distinct graphs, saturating well under the 2,000-entry cache, with the large add-node count being a per-decode-token cost rather than a startup one. The verdict is general rather than Volta-specific, settled locally by comparing loaders on the same hardware: the Gemma 4 loader calls `eval_all` and reaches a first token in 2.4 s, the Qwen 3.5 loader defers and takes 15.4 s, so which phase is charged is a loader property. Tracked outside this epic as #1564.

**8. cuDNN SDPA is off below Ampere** (`scaled_dot_product_attention.cpp:327-330`). Immaterial for this model (head_dim 256 exceeds cuDNN's 128 limit regardless) but relevant to long-context serving.

## Non-goals

- Adding sm_70 to the release arch matrix. Volta stays a source-build target unless the throughput work lands and justifies the binary size.
- Changing any behavior on sm_80 and later. Every item here is gated on `cc < 8` and must be a no-op above it, verified by an unchanged GB10 baseline.
- bf16 numerical parity between Volta and Ampere+. Once activations are f16 on Volta the arithmetic differs by construction; the contract is task quality, not bit-identity across architectures.

## Sub-issues

Phase groups below map to execution waves. Explicit `depends on` annotations carry the real producer/consumer edges and override the phase default, so items with no ordering constraint stay parallel.

### Phase 0

Foundation. Gates everything else: five of the six later items need a runtime or build-time answer to "what compute capability is this?", and every item states its acceptance criteria as a delta against the baseline.

- [x] #1537 Expose CUDA compute capability to the mlxcel runtime, build, and diagnostics
- [x] #1538 Volta (sm_70) baseline record and build coverage

### Phase 1

Low-risk kernel fixes plus the two independent tracks. All four are mutually independent.

- [x] #1539 `qmv`: float accumulators below Ampere at `bits < 8` (depends on #1537, #1538)
- [x] #1541 `qmm_naive`: size the tile from the device shared-memory budget, not an sm80 flag (depends on #1537, #1538)
- [ ] #1542 f16 activation policy below Ampere (depends on #1537, #1538)
- [x] #1544 MoE grouped GEMM selects `cutlass::arch::Sm75` on an sm_70 part (depends on #1537, #1538)
- [x] #1545 CUDA graph instantiation and JIT module load dominate Volta TTFT (depends on #1538)

### Phase 2

Tensor cores. #1542 gates this on a hardware fact, not a preference: Volta MMA has no bf16 variant, so no tensor-core path exists until activations are f16. #1541 is also an edge because both overlay `gemm_sm70.cuh` and touch the same launch path.

- [ ] #1543 `qmm_sm70`: Volta MMA path for quantized GEMM and `gather_gemm` (depends on #1541, #1542)

## GB10 (sm_121) continuation

**Read this section first if you are picking this epic up on a GB10 host.**

Implementation and merge for this epic happen on a Tesla V100-PCIE-32GB (sm_70), which is the only device on the development machine. Every sub-issue's sm_70-side criteria are verified there before merge. The **sm_80-and-later non-regression criteria cannot be verified on that machine and are deliberately left open** rather than ticked unverified. This section is the handoff.

Each merged sub-issue PR states, in its own body, which of its criteria were verified on sm_70 and which were deferred here. Nothing below has been checked on Ampere-or-later silicon.

### What remains to verify on GB10

| Issue | Deferred criterion |
|---|---|
| #1537 | No behavior change on any existing platform with `MLXCEL_TRACE_ARCH` unset; capability probe returns `Some((12, 1))`; the arch-mismatch error does not fire on a matching build |
| #1538 | Nothing device-specific; confirm the baseline doc's methodology section still describes the GB10 sweep correctly |
| #1539 | Greedy output byte-identical on GB10 for a fixed prompt and seed; emitted SASS for the sm_121 pass unchanged (`cuobjdump --dump-sass` diff), or a documented explanation if the `__CUDA_ARCH__` guard perturbs it; GB10 baseline throughput unmoved |
| #1541 | GB10 throughput unmoved; the tiler's sm_80+ selection provably identical for every `(itemsize, group_size, m)` combination |
| #1542 | GB10 greedy output byte-identical; GB10 baseline throughput unmoved; `patches-cuda/dtype.cpp` behavior on sm_80+ bit-for-bit unchanged |
| #1543 | GB10 greedy output byte-identical; GB10 baseline throughput unmoved; the new sm_70 branch provably unreachable above cc 7 |
| #1544 | GB10 MoE output byte-identical; GB10 MoE throughput unmoved; the `== 8` and `> 8` dispatch branches untouched |
| #1545 | GB10 throughput and TTFT unmoved by any change made; plus the open question the issue raises — whether the graph-construction cost is Volta-specific or general, which needs a GB10 comparison to answer |

### How to close it out

1. Build for the release arch matrix and confirm it still compiles: `MLX_CUDA_ARCHITECTURES="90a;100;121" make release-cuda`, and separately the x86_64 matrix `80;86;89;90a;100;120`.
2. Run the GB10 baseline sweep and compare against the pre-epic numbers in `benchmarks/`.
3. Run `cargo test --features cuda` on the GB10 host — the sm_121 half of the program-level criterion below.
4. For the byte-identical checks, generate with a fixed prompt and seed on the merged code and on the pre-epic commit, and diff the token streams.
5. Tick the matching program-level criteria below, and comment the results on each sub-issue.

If a regression appears, the offending change is by construction reachable above cc 7 and its `cc < 8` gate is wrong. File a fix issue referencing this epic and the sub-issue that introduced it, rather than reopening the merged sub-issue.

## Acceptance criteria (program level)

- [x] Volta baseline doc in `docs/benchmark_results/` with the measurements above reproduced from a clean build, plus a post-program comparison table.
- [ ] Decode on `qwen3.8-27B-4bit` reaches >= 25% of the bandwidth roofline (from the measured 7%).
- [ ] Prefill reaches >= 20% of FP32 peak, or >= 10% of tensor-core peak if #1543 lands (from the measured ~3%).
- [ ] No regression on GB10 (sm_121) or the x86_64 release arch matrix: every change gated on `cc < 8`, verified by re-running the GB10 baseline.
- [ ] `cargo test --features cuda` green on both an sm_70 and an sm_121 host.

## References

- Predecessor program, Ampere-and-later: #623.
- Arch selection and auto-detect: `src/lib/mlxcel-core/build.rs:358-435`; `docs/installation.md` (CUDA architecture selection); release matrix `.github/workflows/release.yml:541,764`.
- Overlay inventory that this program will extend: `src/lib/mlx-cpp/patches/mlx/backend/cuda/`.
- Raw profiles from this audit are attached to #1538.