Skip to content

chore(deps): bump the MLX pin to 2c46b953 and re-derive the three conflicting source overlays #1042

Closed

Description

@inureyes

Problem / Background

Bump mlx from b7c3dd6d27f45b5365b08a840310187dc503f1db (2026-07-17) to 2c46b953db88965c4270cc7306eda6887a3247f2 (2026-08-05, upstream main), and re-derive the source overlays that upstream touched in that window. The range covers 58 commits and 156 changed files.

Why now

Fixes in range that land on code paths this repo actually ships:

  • 6c0ea7fb0 Fix incorrect nvfp4 quantized_matmul through the split-K path (#3854). NVFP4 direct transcode is mlxcel's default path on all builds, so this is a correctness fix on a shipped code path.
  • 973e27f82 [CUDA] Fix grid overflow in gemm conv unfold kernels for >= 65,536 output positions (#3893). Same class as the grid-dimension crash mlxcel already carries a chunking overlay for.

Performance and coverage:

  • b400c6ced metal: reduce NVFP4 scales per 16-lane group (#3934)
  • 0ebcee8db metal: gemv_wide for fp16/bf16 few-row matmuls (#3888)
  • ab946f90e head dimension 96 in Metal full attention (#3943)
  • 8584610b3 pad 2D conv input channels to reach the specialized Metal kernel (#3904)

Assorted correctness:

  • 9b40c9d1b prod dtype promotion on a size-1 axis (#3898)
  • 6dfd81406 shapeless matmul with dynamic batch dims (#3813)
  • ce3073389 captured random state in compile (#3828)
  • 2c46b953d mx.remainder floored-mod for float16/bfloat16 on CPU (#3976)

The risk this issue exists to manage

mlxcel applies whole-file source overlays via configure_file(... COPYONLY) in src/lib/mlx-cpp/CMakeLists.txt. An overlay whose upstream counterpart has changed will silently revert that upstream change. This is not hypothetical: it is exactly how #824 / #829 shipped a broken RMSNorm kernel by overlaying a pre-fix version of rms_norm.cu, which took until #830 / #831 to detect and undo. Every pin bump must therefore audit all 21 overlay files against the upstream diff, not just move the tag.

Audit result: 3 of 21 overlays conflict

Overlay Upstream change in range Resolution
src/lib/mlx-cpp/patches/mlx/backend/cuda/quantized/quantized.cpp qmm_naive gained a const std::optional<array>& global_scale parameter at position 5, from a79332de6 Add gather_qqmm (#3757) mlxcel's overlay wraps both qmm_naive call sites in run_row_chunked / run_batch_chunked, so git merge-file cannot reconcile it. Resolved by hand, inserting std::nullopt after biases at both sites and verifying both calls are 13 arguments in declaration order.
src/lib/mlx-cpp/patches/mlx/backend/metal/kernels/utils.h c4111f8c3 Template Metal complex scalar lanes (#3970) and ef5fc0fab explicit thread address space qualifier for metal 4.1 (#3963) Clean 3-way merge, no overlap with mlxcel's delta (#include <metal_simdgroup_matrix> plus two metal::vec qualifications).
src/lib/mlx-cpp/patches-cuda/ops.cpp 189 changed lines Clean 3-way merge, mlxcel's 57-line delta preserved.

mlx/backend/cuda/quantized/qmm/qmm.h also changed upstream but is not overlaid, so mlxcel picks up the new declaration. The full diff of that header in range is the single global_scale line, so the overlay's other calls (qmm_sm80, qmv, gather_qmv, fp_qmv, qmm_sm90) are unaffected. The qmm_sm80.cuh race-fix overlay (from #910) is untouched upstream and is still required.

Second trap: the pin has two sources of truth

  • src/lib/mlx-cpp/CMakeLists.txt:95 (GIT_TAG), which controls what gets fetched.
  • src/lib/mlxcel-core/build.rs:217 (MLX_EXPECTED_COMMIT), which drives purge_stale_mlx_cache and is exported as MLXCEL_MLX_COMMIT.

Updating only the CMakeLists leaves the cache marker matching MLX_EXPECTED_COMMIT, so _deps/ is judged valid and never purged, and the build silently compiles the old MLX while appearing to have bumped. Both must move together.

Suggested follow-up (out of scope here): collapse these into a single source of truth, or add a build-time consistency assert that fails when the two disagree.

Acceptance Criteria

  • Both pin locations updated to 2c46b953db88965c4270cc7306eda6887a3247f2 and verified identical.
  • The three conflicting overlays re-derived, with each one's upstream delta and mlxcel delta both demonstrably present.
  • A release CUDA build completes, and the fetched _deps/mlx-src is confirmed at the new commit (guard against the silent stale-cache mode above).
  • Test suites for the touched surfaces pass, plus a real-model smoke on GB10.
  • Any behavior change from the upstream nvfp4 split-K fix (#3854) is checked rather than assumed, since NVFP4 is a default path here.

Technical Considerations

Work is in progress on branch chore/mlx-pin-2c46b953.

Refs: follow-up to the overlay-audit lesson from #830 / #831. Related: #821 (the MLX graph-cache abort tracked separately), which is still unfixed at this new pin, so apply_cuda_graph_cache_default's 2000 default and its tracking comment stay.

Activity

  1. added a commit that references this issue on Aug 6, 2026
    94a8608
  2. self-assigned this
    on Aug 31, 2026
  3. added a commit that references this issue on Oct 27, 2026
    b0b8960
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

Labels

Type

No type

Projects

No projects

    Milestone

    No milestone

    Relationships

    None yet

    Development

    No branches or pull requests

    Issue actions