chore(mlx): bump pin to e9463bb and rebase CUDA overlays (#625)

[inureyes]

Summary

Bumps the pinned MLX commit from a6ec7123dac814417147e21d4aeed694924ddd4d (2026-06-10) to e9463bbfc1a7cd9e0e6b96aaa3068a316e234a63 (upstream main, 2026-07-01) in both enforcement points, and reconciles all 20 source overlays three-way against the 63 upstream commits between the pins. Part of epic #623; the full GB10 re-baseline sweep and outlier triage verdicts are done separately on top of this (placeholder committed at docs/benchmark_results/mlx-pin-upgrade-2026-07-03.md).

Pin bump

• src/lib/mlx-cpp/CMakeLists.txt GIT_TAG -> e9463bb...
• src/lib/mlxcel-core/build.rs MLX_EXPECTED_COMMIT -> e9463bb... (the mismatch check and .mlx-build-commit purge forced the clean MLX rebuild)
• MLX_CCCL_DIR is now defined on the new mlx_dirs OBJECT target (upstream moved the macro consumer into dirs.cpp to isolate dynamic defines from the compile cache), with a fallback to the mlx target for older trees. Verified in the generated mlx_dirs.dir/flags.make.

Overlay reconciliation

14 of 20 overlay base files are byte-identical between the two pins, so those overlays apply unchanged. The remaining six were reconciled as follows:

Overlay	Decision	Reason
quantized/qmm/qmm.h	dropped	Comment-only delta vs old upstream; pristine e9463bb adds the global_scale param (#3723).
quantized/qmm/qmv.cu	dropped	Comment-only delta (its documented broadcast_w fix had already landed upstream); pristine e9463bb carries qmv global scale (#3723) and the relocated device/cute_dequant.cuh include (#3576).
primitives.cpp	dropped	No-op overlay (annotation comment only).
quantized/quantized.cpp	rebased	Re-applied the ensure_row_contiguous fix for w/scales/biases in QuantizedMatmul::eval_gpu (3D batched MLA weights, upstream-rejected) onto the new base, which now passes std::nullopt global_scale to qmv. GatherQMM path returns to pristine upstream.
jit_module.cpp	rebased	Re-applied /proc/self/exe dir resolution + MLXCEL_CCCL_DIR override + cold-JIT notice onto the new base; adopted upstream's cccl_dir() fallback, extra bundled include path, JIT'd qmm header registration (#3706/#3576), CUTLASS nvrtc flags, and get_jit_module(Device&) signature. MLX_PTX_CACHE_DIR handling is unchanged upstream, so the persistent PTX cache redirect still takes effect (upstream #3587 turned out to be a CI-cache change only).
binary/binary.cuh	rebased (adapted)	Re-applied the mixed-precision bf16/fp32 kernels + trait + dispatch onto the new base. Also drops stale v0.31.1-era reversions the old overlay silently carried (missing __launch_bounds__), and adopts upstream's large-grid index_rest fix (grid.z) and get_launch_args_general in both same-type and mixed general paths.
patches-cuda/ops.cpp	rebased (adapted)	Re-applied bf16_mixed_astype() + its use in add/subtract/multiply/divide/maximum/minimum. Drops stale v0.31.1-era reversions (safe_cast overflow checks, arange zero-step check, use-after-move-prone auto& shape bindings) and picks up upstream's new array-API ops unmodified.

Kept unchanged (base identical between pins): device/binary_ops.cuh, gemms/grouped_gemm.h, gemms/grouped_gemm_unaligned.cu, matmul.cpp (the CUTLASS GatherMM rework intentionally still bypasses upstream's JIT gather_gemm; the re-baseline sweep should confirm this remains the right call on MoE models), reduce/*, metal/compiled.cpp, metal/kernels/utils.h, patches-cuda/{dtype,fast}.cpp.

Quantized kernels moved to runtime JIT (deployment impact)

Upstream #3706/#3576 JIT-compile the CUTLASS qmm/gather_gemm kernels with NVRTC at runtime instead of nvcc at build time. Cold build time dropped from the historical 30+ minutes to ~7.5 minutes on GB10, but the JIT'd kernels include <cute/...>/<cutlass/...>, which a statically linked mlxcel binary could not resolve (first quantized run failed with cannot open source file "cute/numeric/numeric_types.hpp"). Fixed in this PR, mirroring the existing CCCL solution:

• jit_module.cpp overlay resolves CUTLASS headers via MLXCEL_CUTLASS_DIR env override, the bundled <exe>/../include layout, or the compiled-in MLX_CUTLASS_DIR build-tree fallback.
• mlx-cpp/CMakeLists.txt compiles in the _deps/cutlass-src/include fallback for source builds.
• release.yml bundles include/cute/ + include/cutlass/ in both Linux CUDA archive jobs (with completeness checks).
• docs/installation.md documents the new runtime requirement and env var.
• build.rs now watches ../mlx-cpp/patches-cuda with rerun-if-changed (previously unwatched, so edits to the CUDA-only overlays did not retrigger the MLX build).

First quantized-model runs pay a one-time per-kernel-variant NVRTC compile, amortized by the persistent commit-scoped PTX cache.

Test-suite reconciliation

cargo test --release --features cuda on GB10 surfaced two failures, both understood:

• sanitize_tests::load_and_sanitize_weights_dequantizes_nvfp4_gemma4_checkpoint: pre-existing CUDA bug, not a pin regression. The CUDA loader widens F16 safetensors to F32 at load (from_bytes_f16, CUDA branch dating to 2026-03), while the NVFP4 dequant parsed the block scales as raw F16 bytes, reading zeros on CUDA. Fixed by normalizing the scales to F32 via astype before the byte parse (dtype-agnostic on both backends). This also corrects real gemma4-nvfp4 block-scale handling on CUDA, relevant to the epic's gemma-4-31b nvfp4 outlier.
• gemma4_mtp_target::tests::divergent_round_hidden_matches_b1_replay: pin-related property change. Since MLX #3704 (rope without copy) RoPE preserves input layouts instead of canonicalizing them with a copy, so downstream GEMM/SDPA kernels can see different strides for a batched row than for the B=1 replay and pick different reduction orders. Observed max_abs 1.6e-4 on fp32 hidden of magnitude ~8, i.e. reduction-order noise, not a logic bug. The bitwise assertion is now a 2e-3 fp32 tolerance on CUDA (bitwise retained on other backends); the argmax-equality assertion, which is the functional MTP-verify guarantee, is unchanged.

The default parallel cargo test run aborts under the new pin even on warm PTX caches: the JIT'd kernels are compiled and module-loaded on first use, and those first-use paths are not safe against concurrent test threads. Observed as cudaStreamEndCapture ... previous error during capture (module load racing another thread's stream capture) and, with MLX_USE_CUDA_GRAPHS=0, as cuLaunchKernelEx ... invalid argument (kernel-configure race in get_kernel_and_dims). The suite is stable with --test-threads=1; documented in docs/installation.md. Inference binaries are unaffected (parity runs and benches are clean). Worth an upstream issue; also worth watching if mlxcel-server ever evals from multiple threads.

One more pre-existing failure fixed en route: family_order_is_exhaustive (bin target) was failing on main because the recent VLM ports added the GLM VLM, Kimi VLM, and PaddleOCR VLM families without updating FAMILY_ORDER in src/main.rs; the three families are now listed before Other VLM.

Tegra managed-memory finding (#3701)

GB10 reports concurrentManagedAccess=1 (probed with a CUDA attribute query), so upstream's new gate does not trip: managed memory stays enabled on GB10 under both pins, exactly as before. #3701 is a no-op on this platform (it targets Jetson Orin Nano-class boards and WSL, which report 0). The epic hypothesis that the old pin's managed-memory behavior was depressing GB10 numbers is refuted at the gate level.

Test plan

• Cold cargo build --release --features cuda from purged _deps (fetched e9463bb, all overlays applied; verified in the configured tree and mlx_dirs flags)
• cargo test --features cuda locally
• Parity spot-checks (coherent short generations): llama-3.1-8b-4bit, qwen3-8b-4bit, qwen3-30b-a3b-4bit (MoE), qwen2.5-vl-3b-4bit (VLM)
• Quick decode sanity bench on a small subset (no gross regression)
• Full GB10 sweep + long-prompt ladder + outlier verdicts (orchestrator, epic epic: CUDA (MLX backend) prefill and decode acceleration program #623)
• Metal re-bench (follow-up; Metal overlays compile-verified only via CI matrix)

Closes #625