perf(cuda): extend sdpa_vector decode kernels to head_dim 256/288 via MLX overlay

benchmarks/cuda_gb10_sdpav_675_2026-07-06.csv

@@ -0,0 +1,13 @@
+model,prompt_tokens,generated_tokens,prefill_chunk,prefill_ms,prefill_tok_s,decode_ms,decode_tok_s,mlx_peak_gb,sys_peak_gb,date,hardware,mlx_version,build_type,notes
+gemma-4-12b-it-4bit,512,64,,,,4377.57,14.62,,,2026-07-06,NVIDIA_GB10_122GB,0.3.3,release,sdpa_vector path (#675 overlay); 1-pass kernel; decode_ms derived from tok/s
+gemma-4-12b-it-4bit,512,64,,,,4347.83,14.72,,,2026-07-06,NVIDIA_GB10_122GB,0.3.3,release,materializing fallback (MLXCEL_SDPA_VECTOR_LARGE_D=0); decode_ms derived from tok/s
+gemma-4-12b-it-4bit,2048,64,,,,4466.15,14.33,,,2026-07-06,NVIDIA_GB10_122GB,0.3.3,release,sdpa_vector path (#675 overlay); 2-pass kernel; decode_ms derived from tok/s
+gemma-4-12b-it-4bit,2048,64,,,,4469.27,14.32,,,2026-07-06,NVIDIA_GB10_122GB,0.3.3,release,materializing fallback (MLXCEL_SDPA_VECTOR_LARGE_D=0); decode_ms derived from tok/s
+gemma-4-12b-it-4bit,8192,64,,9768.20,838.64,4481.07,14.28,14.51,,2026-07-06,NVIDIA_GB10_122GB,0.3.3,release,sdpa_vector path (#675 overlay); 2-pass kernel
+gemma-4-12b-it-4bit,8192,64,,10003.80,818.89,4454.60,14.37,14.51,,2026-07-06,NVIDIA_GB10_122GB,0.3.3,release,materializing fallback (MLXCEL_SDPA_VECTOR_LARGE_D=0)
+gemma-4-12b-it-4bit,32768,64,,42398.79,772.85,4643.21,13.78,16.65,,2026-07-07,NVIDIA_GB10_122GB,0.3.3,release,sdpa_vector path (#675 overlay); only 8/48 layers are full-attention
+gemma-4-12b-it-4bit,32768,64,,43317.00,756.47,4678.69,13.68,16.65,,2026-07-07,NVIDIA_GB10_122GB,0.3.3,release,materializing fallback (MLXCEL_SDPA_VECTOR_LARGE_D=0)
+qwen3.5-2b-4bit,8192,64,,2517.63,3253.85,495.89,129.06,5.48,,2026-07-07,NVIDIA_GB10_122GB,0.3.3,release,sdpa_vector path (#675 overlay); full attention every layer; 1.123x vs fallback
+qwen3.5-2b-4bit,8192,64,,2553.53,3208.11,557.06,114.89,5.48,,2026-07-07,NVIDIA_GB10_122GB,0.3.3,release,materializing fallback (MLXCEL_SDPA_VECTOR_LARGE_D=0)
+qwen3.5-9b-4bit,8192,64,,10923.69,749.93,1660.21,38.55,12.30,,2026-07-07,NVIDIA_GB10_122GB,0.3.3,release,sdpa_vector path (#675 overlay); full attention every layer; 1.044x vs fallback
+qwen3.5-9b-4bit,8192,64,,11060.52,740.65,1733.28,36.92,12.30,,2026-07-07,NVIDIA_GB10_122GB,0.3.3,release,materializing fallback (MLXCEL_SDPA_VECTOR_LARGE_D=0)

docs/benchmark_results/mlx-pin-upgrade-2026-07-03.md

@@ -40,6 +40,7 @@ in `src/lib/mlx-cpp/CMakeLists.txt`. Each was reconciled three-way
 | `patches/mlx/backend/cuda/reduce/init_reduce.cu` | kept | Base unchanged. Output buffer typed as T so it matches out.dtype() for bf16. |
 | `patches/mlx/backend/cuda/reduce/reduce_ops.cuh` | kept | Base unchanged. `ReduceResult<Sum/Prod, bf16>` accumulate-in-fp32 specializations. |
 | `patches/mlx/backend/cuda/reduce/row_reduce.cu` | kept | Base unchanged. Same V/U split. |
+| `patches/mlx/backend/cuda/scaled_dot_product_attention.cu` | added (#675) | New overlay, not part of the #625 rebase. Extends the `sdpa_vector` decode kernels and the `supports_sdpa_vector` gate to head_dim 256/288 so head_dim > 128 models (gemma family, qwen3.5/3.6, baichuan-m1, paligemma2) take the fused vector path at decode instead of the unfused materializing SDPA fallback. Gated behind the `MLXCEL_SDPA_VECTOR_LARGE_D` env kill switch (default on); `__launch_bounds__` caps the larger per-thread register footprint of the D=256/288 instantiations. Base is pristine e9463bb, so future pin bumps must rebase this file. |
 | `patches/mlx/backend/metal/compiled.cpp` | kept | Base unchanged. static_cast insertion for mixed-dtype compiled ops. |
 | `patches/mlx/backend/metal/kernels/utils.h` | kept | Base unchanged. `<metal_simdgroup_matrix>` include and `metal::vec` qualification. |
 | `patches-cuda/dtype.cpp` | kept | Base unchanged. bf16+fp32 -> bf16 promotion-table patch (CUDA-only). |

docs/environment-variables.md

@@ -207,6 +207,7 @@ recommended as normal deployment settings.
 | `MLXCEL_DISABLE_SOFTCAP_GQA_DECODE_GROUPED` | `1` disables, `0` enables | unset | Legacy rollback/override for grouped softcap-GQA decode. |
 | `MLXCEL_DISABLE_SINGLE_QUERY_MASKLESS` | truthy disables | maskless path on | Disables the single-query maskless attention path. |
 | `MLXCEL_EXPERIMENTAL_BOOL_CAUSAL_MASK` | truthy enables | off | Enables an experimental boolean causal-mask path. |
+| `MLXCEL_SDPA_VECTOR_LARGE_D` | `0`/`false`/`off`/`no` disable; any other value or unset enables | on | **CUDA only.** Gates whether the CUDA `supports_sdpa_vector` check accepts head_dim 256/288 (gemma family, qwen3.5/3.6, baichuan-m1, paligemma2), routing their decode to the fused `sdpa_vector` kernels instead of the materializing SDPA fallback (issue #675). Disabling restores the prior fallback with no rebuild; used for the A/B in `benchmarks/cuda_gb10_sdpav_675_2026-07-06.csv`. |
 | `MLXCEL_PIPELINE_GRANULARITY` | `off`, `layer`, `block:N` | `off` | Inserts layer-boundary async-eval hints for pipeline experiments. |
 | `MLXCEL_FUSED_MOE` | `0`/`false`/`off`/`no` disable; any other value or unset enables | on | Fused single-token decode-MoE kernel (#268), on by default since #282 (Metal) and #319 (CUDA, via `mx.fast.cuda_kernel`); validated on M1 Ultra, M5, and GB10. Set to `0` to force the proven `gather_qmm`/`SwitchGLU` path. Active for qwen3_moe, qwen3_next, dots.llm1, gemma4, qwen2_moe, mixtral, phimoe, lfm2, qwen3_vl_moe, and olmoe decode. |
 | `MLXCEL_FUSED_MOE_SGY` | `1`-`32` | `8` | Simdgroups (Metal) / warps-per-block (CUDA) per threadgroup for the fused decode-MoE kernel; tune per hardware. |