epic: NVIDIA Volta (sm_70) inference acceleration program #1536

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).

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):

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:

template <bool SM80 = false, typename Element = float, typename CtaTiler>
inline constexpr auto make_tiled_mma(CtaTiler cta_tiler) {
  using Atom = cuda::std::conditional_t<SM80,
      /* SM80_16x8x16_F32F16F16F32_TN / SM80_16x8x16_F32BF16BF16F32_TN */,
      UniversalFMA<float, Element, Element>>;   // <- 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:

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.

• chore(core): expose CUDA compute capability to the mlxcel runtime, build, and diagnostics #1537 Expose CUDA compute capability to the mlxcel runtime, build, and diagnostics
• test(benchmark): Volta (sm_70) baseline record and build coverage #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.

• perf(cuda/quant): use float accumulators in qmv below Ampere at bits < 8 #1539 qmv: float accumulators below Ampere at bits < 8 (depends on chore(core): expose CUDA compute capability to the mlxcel runtime, build, and diagnostics #1537, test(benchmark): Volta (sm_70) baseline record and build coverage #1538)
• perf(cuda/quant): size the qmm_naive tile from the device shared-memory budget, not an sm80 flag #1541 qmm_naive: size the tile from the device shared-memory budget, not an sm80 flag (depends on chore(core): expose CUDA compute capability to the mlxcel runtime, build, and diagnostics #1537, test(benchmark): Volta (sm_70) baseline record and build coverage #1538)
• perf(cuda): f16 activation policy below Ampere (bf16 has no Volta hardware path) #1542 f16 activation policy below Ampere (depends on chore(core): expose CUDA compute capability to the mlxcel runtime, build, and diagnostics #1537, test(benchmark): Volta (sm_70) baseline record and build coverage #1538)
• fix(cuda/moe): grouped GEMM selects cutlass::arch::Sm75 on an sm_70 part #1544 MoE grouped GEMM selects cutlass::arch::Sm75 on an sm_70 part (depends on chore(core): expose CUDA compute capability to the mlxcel runtime, build, and diagnostics #1537, test(benchmark): Volta (sm_70) baseline record and build coverage #1538)
• perf(cuda): CUDA graph instantiation and JIT module load dominate Volta TTFT #1545 CUDA graph instantiation and JIT module load dominate Volta TTFT (depends on test(benchmark): Volta (sm_70) baseline record and build coverage #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.

• perf(cuda/quant): qmm_sm70 — Volta tensor-core MMA path for quantized GEMM #1543 qmm_sm70: Volta MMA path for quantized GEMM and gather_gemm (depends on perf(cuda/quant): size the qmm_naive tile from the device shared-memory budget, not an sm80 flag #1541, perf(cuda): f16 activation policy below Ampere (bf16 has no Volta hardware path) #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

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)

• 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 perf(cuda/quant): qmm_sm70 — Volta tensor-core MMA path for quantized GEMM #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: epic: CUDA (MLX backend) prefill and decode acceleration program #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 test(benchmark): Volta (sm_70) baseline record and build coverage #1538.