Repository navigation
chore(deps): bump the MLX pin to 2c46b953 and re-derive three overlays - #1046
Merged
Merged
Conversation
Move the MLX pin from b7c3dd6d (2026-07-17) to 2c46b953 (2026-08-05 upstream main), 58 commits. The range carries fixes that land on paths this repo ships: incorrect nvfp4 quantized_matmul through the split-K path (#3854), a CUDA gemm conv unfold grid overflow past 65,535 output positions (#3893), Metal NVFP4 scale reduction per 16-lane group (#3934), gemv_wide for few-row fp16/bf16 matmuls (#3888), head dimension 96 in Metal full attention (#3943), plus prod dtype promotion on a size-1 axis (#3898) and shapeless matmul with dynamic batch dims (#3813). mlxcel applies whole-file source overlays, so an overlay whose upstream counterpart moved silently reverts that upstream change. That is exactly how a broken RMSNorm kernel shipped in #824/#829 and took until #830/#831 to undo, so all 21 overlays were audited against the upstream diff. Three conflicted: - quantized.cpp: #3757 (gather_qqmm) inserted a global_scale parameter into qmm_naive at position 5. mlxcel wraps both call sites in run_row_chunked and run_batch_chunked, so git merge-file cannot reconcile the two; resolved by hand and verified as 13 arguments in declaration order at both sites. The full diff of quantized/qmm/qmm.h across the range is that one line, so the other dispatch entry points consumed here are unaffected. - metal/kernels/utils.h: clean three-way merge against #3970 and #3963; no overlap with mlxcel's include and metal::vec qualifications. - patches-cuda/ops.cpp: clean three-way merge, mlxcel's 57-line delta intact. The pin has two sources of truth: GIT_TAG in mlx-cpp/CMakeLists.txt and MLX_EXPECTED_COMMIT in mlxcel-core/build.rs, the latter driving purge_stale_mlx_cache. Updating only the first leaves the cache marker matching and the build silently compiles the old MLX. Both moved together, and the result was checked three ways: the binary's embedded commit, the cache marker, and the fetched _deps/mlx-src HEAD, with no trace of the old SHA. Validation on GB10: release CUDA build clean, fused_norm_parity 8/8, turbo_tests 97/97, cargo fmt clean, and greedy decode on llama-3.2-1b-4bit, qwen3-4b-4bit and deepseek-v2-lite-4bit coherent and byte-identical run to run, which also confirms the qmm_sm80 race-fix overlay still holds on the new base. Two pre-existing failures were separated from this change by a control build at the old pin. fused_rope_parity fails 2 of 9 at offsets 131069 and 131071, with normalized rms 2.497e-3 and 2.684e-3 against a 2.0e-3 tolerance; the old pin reproduces those digits exactly, and upstream touched no rope file in the range. The full lib suite aborts under CUDA graph capture with a nondeterministic site and message (cuGraphAddKernelNode invalid argument, then cudaStreamEndCapture previous error during capture); the old pin aborts the same way. Neither is caused by this bump and both are left for separate issues. Refs #1042
inureyes
added
type:dependency
This was referenced Aug 5, 2026
8 tasks done
inureyes
added a commit
that referenced
this pull request
Aug 6, 2026
PR #1051 (issue #1047) made the GIT_TAG line in src/lib/mlx-cpp/CMakeLists.txt the single source of truth for the pinned MLX commit, but review left several loose ends that this closes. The comment above GIT_TAG in src/lib/mlx-cpp/CMakeLists.txt told a future bumper to keep the literal shape but not why; it now states the awk parser's actual constraints: GIT_TAG must be the first token on its line, that line must come after the GIT_REPOSITORY line naming ml-explore/mlx, and no closing paren may appear between them. src/lib/mlxcel-core/build.rs and src/lib/mlx-cpp/turbo/CMakeLists.txt pointed readers at a CLAUDE.md that does not exist in this repository. Both now point at CONTRIBUTING.md's "Bumping the MLX upstream pin" section, which already covers the same checklist. docs/upstream/mlx-cuda-graph-cache-lifetime-miss-abort.md and docs/upstream/mlx-cuda-rmsnorm-small-axis-regression.md claimed a commit that went stale at #1046 was still mlxcel's current pin, including a line reference into CMakeLists.txt that had also drifted. Both now describe that commit as the pin at the time each report was written and point at the GIT_TAG argument for the pin in effect today, without restating a value that can go stale again. Added src/lib/mlxcel-mlx-pin/tests/cross_parser.rs, which shells out to mlx_pinned_commit.sh and asserts that whenever it accepts an input, the Rust parser resolves to the identical commit for that same input. This surfaced a real, previously untested asymmetry: the Rust parser accepts a GIT_TAG line placed before GIT_REPOSITORY in the declaration block, while the shell parser rejects it. A fixture for that case is kept in the new test so the asymmetry stays a documented, deliberate property (the shell parser is the stricter of the two by design) rather than an untested assumption. Wired cargo test -p mlxcel-mlx-pin into the mlx-pin job in .github/workflows/ci.yml, and added the crate's own sources to the job's path filter, so both parsers are actually exercised on every PR that touches them instead of only the shell-side negative-path script. Validation: - cargo test -p mlxcel-mlx-pin: 27 tests pass, including the new cross-parser test - cargo clippy -p mlxcel-mlx-pin --all-targets -- -D warnings: clean - cargo fmt --all -- --check: clean - scripts/ci/mlx_pinned_commit.sh and scripts/ci/mlx_pinned_commit_test.sh: both pass against the edited CMakeLists.txt - .github/workflows/ci.yml validated with python3's yaml.safe_load Refs #1047
inureyes
added a commit
that referenced
this pull request
Aug 6, 2026
## Summary The pinned MLX commit was written out in three places with nothing checking that they agreed, so bumping one produced a build that looked bumped and compiled the old MLX. This makes `GIT_TAG` in `src/lib/mlx-cpp/CMakeLists.txt` the only place the value lives (the issue's option 1), adds a post-build check that the checkout which actually landed on disk is that commit (option 3), and removes the third location, which had already gone stale. The pin value is unchanged at `2c46b953db88965c4270cc7306eda6887a3247f2`. ## The third location, and why it matters The issue names two locations. There were three, and `.github/workflows/release.yml` was already a generation behind: it carried `b7c3dd6d27f45b5365b08a840310187dc503f1db` while the tree had moved to `2c46b953` in #1046, and the comment above it called this "the three-place pin" and pointed at a documentation file that does not exist in this repository (the checklist is in `CONTRIBUTING.md`). The effect there was fail-safe rather than silent, since the mismatched literal made every marker comparison fail, but the practical result was that every release build purged and rebuilt MLX from scratch while the workflow reported a commit the tree no longer used. So the silent-divergence failure the issue predicts had already happened once, in the fail-safe direction. ## What changed - `src/lib/mlxcel-core/build.rs`: `MLX_EXPECTED_COMMIT` is deleted. `main()` resolves the pin once from `../mlx-cpp/CMakeLists.txt` before anything compares against it, and threads that one value through `build_mlx`, `purge_stale_mlx_cache`, the new `verify_fetched_mlx_head`, `mark_mlx_cache_valid` and the `cargo:rustc-env=MLXCEL_MLX_COMMIT` export. A parse failure panics with a message naming the file and what it expected. No duplicate `rerun-if-changed` was added; the existing one at what is now line 204 is unchanged. - `src/lib/mlxcel-core/build_support/mlx_pin.rs` (new): the pure logic. Parsing strips CMake comments first, then scopes itself to the `FetchContent_Declare` block whose `GIT_REPOSITORY` names `ml-explore/mlx`, and treats zero blocks, two blocks, zero `GIT_TAG`s, two `GIT_TAG`s, and any value that is not a 40-character lowercase hex SHA as hard failures with distinct actionable messages. Uppercase is rejected too: `git rev-parse` prints lowercase and the marker is compared byte for byte, so an uppercase pin would make every cache look stale forever. - Post-build HEAD verification: after CMake returns and **before** the marker is written, `_deps/mlx-src`'s HEAD is read and a disagreement fails the build, telling the reader to delete `_deps/` to force a refetch. Verifying before marking is the point: a tree that failed verification is never blessed as valid. A missing `.git` or a host without `git` emits a `cargo:warning` and skips, so vendored and offline trees still build. The `.git` existence probe is load-bearing rather than a fast path, because `git -C` walks up to the nearest enclosing repository and `_deps/` sits inside the mlxcel checkout, so without it a metadata-free tree would answer with mlxcel's own HEAD and be reported as a pin mismatch. There is a unit test for exactly that. - `.github/workflows/release.yml`: the workflow-scope `env: MLX_EXPECTED_COMMIT` literal is gone. Both "Validate MLX build cache" steps derive the pin from the CMake file through the new `scripts/ci/mlx_pinned_commit.sh`, and refuse to run the comparison against an empty value, since an empty pin would mark every cache stale and purge all of them rather than validate any. Purge behaviour is otherwise byte-identical. The stale documentation pointer and the stale SBOM comment at the bottom of the file are fixed. - `scripts/ci/mlx_pinned_commit.sh` (new): the shell half of the parser. Same scoping (comments stripped, matched on the MLX repository), exits non-zero unless exactly one 40-character lowercase hex SHA resolves. Deliberately free of regex interval syntax, which is not portable across the awk implementations on the GitHub-hosted Linux and self-hosted macOS runners; the SHA shape is validated in shell instead. `MLX_PIN_CMAKE_FILE` overrides the input file for testing. - `src/lib/mlx-cpp/CMakeLists.txt`: a comment above `GIT_TAG` naming both parsers and stating that the literal `GIT_TAG <40-char lowercase hex sha>` shape must stay parseable. Verified that CMake accepts a comment inside the command invocation and drops it from `ARGN`. - Docs: `CONTRIBUTING.md`'s "Bumping the MLX upstream pin" section is rewritten for the one-file reality (the Metal kernel re-validation checklist after it is kept verbatim), and `docs/benchmark_results/model_tests.md` now points at the CMake `GIT_TAG` instead of the deleted constant. ## Test mechanism, and why it is not a build-dependency The logic lives in `mlxcel-core/build_support/mlx_pin.rs`, which `build.rs` includes with `#[path]` and a new leaf crate `mlxcel-mlx-pin` includes as a library. That crate has no production role and nothing depends on it; it exists purely so `cargo test -p mlxcel-mlx-pin` can cover the logic in about a second without compiling `mlxcel-core` and therefore without building MLX C++. The more conventional shape, a `[build-dependencies]` entry, was rejected on purpose: a new build-dependency feeds the build-script unit's metadata hash, that hash names `mlxcel-core`'s `OUT_DIR` (`target/<profile>/build/mlxcel-core-<hash>/out`), and moving it strands every warm `_deps/` MLX checkout under the old hash. Including the source by path leaves the hash untouched. This is reasoning from Cargo's metadata model rather than an experiment; testing it would have required the rebuild it is meant to avoid. Adding a fifth workspace member made four "four members" statements stale, in `Cargo.toml`, `Makefile`, `CONTRIBUTING.md` and `docs/installation.md`. Those are corrected in the same commit rather than left to rot, which is the same defect class as the issue itself. ## Test plan 25 unit tests, all passing, all runnable without an MLX build: - [x] `cargo test -p mlxcel-mlx-pin` (25 passed): happy path, quoted value, comment inside the block, commented-out declaration, missing `GIT_TAG`, branch name, short SHA, uppercase SHA, two MLX declarations, two `GIT_TAG`s in one declaration, an unrelated dependency carrying its own `GIT_TAG` (the wrong-block case), no MLX declaration, unreadable file, error-message content, marker-matches, marker-matches-with-trailing-whitespace, marker-differs, marker-absent, HEAD-matches, HEAD-differs, no-git-metadata, missing directory, no-walk-up-into-enclosing-repository, and the recovery text of the mismatch message. The HEAD tests run against real `git init` repositories, not a mocked `git`. - [x] One test parses the actual in-tree `CMakeLists.txt` and asserts the result is a well-formed SHA. It deliberately does **not** assert the literal, because doing so would put a second copy of the pin in the tree and recreate this issue inside the test suite. - [x] `cargo clippy -p mlxcel-mlx-pin --all-targets -- -D warnings`: clean. - [x] `cargo fmt --all -- --check`: clean. - [x] `build.rs` type-checks and passes `clippy -D warnings` with and without the `cuda` feature. Checked by compiling it verbatim in an isolated scratch crate with `cmake` and `cxx-build` as ordinary dependencies, since checking it in place re-enters CMake. This also confirms the build.rs view is free of `dead_code`, which the `#![allow(dead_code)]` in the shared module exists to guarantee for helpers only the tests call. - [x] `scripts/ci/mlx_pinned_commit.sh` exercised against the real file plus eight synthetic fixtures via `MLX_PIN_CMAKE_FILE`: correct SHA on the real file and on the other-dependency and commented-out cases, non-zero with a specific message on missing tag, branch name, short SHA, two MLX declarations, no MLX declaration and missing file. - [x] Both parsers cross-checked: the Rust parser, the awk parser and the value committed on `main` all resolve to `2c46b953db88965c4270cc7306eda6887a3247f2`, so this change is value-preserving. - [x] `release.yml` parses as YAML and both "Validate MLX build cache" steps still sit immediately after `Checkout code`, which is what makes the repo available to the extraction script. ## What is not verified here - The full `cargo build --release --features cuda` is run separately, outside this PR's local verification, because every profile/feature pair that is not already warm rebuilds MLX from source. - The issue's second acceptance criterion, bumping to a different commit and confirming the fetched `_deps/mlx-src` HEAD, the `.mlx-build-commit` marker and the embedded `MLXCEL_MLX_COMMIT` all move together, was not executed as a deliberate two-rebuild experiment. What supports it: the five caches present on the development host before this change spanned both pin generations (two at `2c46b953`, three at `b7c3dd6d`) with the marker and the `mlx-src` HEAD agreeing in every one, which is the invariant the new HEAD check enforces rather than assumes. - A note on how the cold build became available: while validating this change I ran the `Validate MLX build cache` step body extracted from `release.yml` against a scratch `CARGO_TARGET_DIR`, from the repository root. The first loop honoured the scratch directory; the second loop is written against the literal relative path `target` and consults no marker, so it deleted all five `_deps/` trees under the real `target/` (about 3.8 GB) with `rm -rf`. They are unrecoverable. The upside is incidental, not planned: with `_deps/` absent, the verification build is a genuine cold fetch, so the fetched HEAD, the marker and the embedded `MLXCEL_MLX_COMMIT` are all observed being produced from the single source in one run. That is what turned the second criterion from prohibitively expensive into a by-product, at a cost nobody chose to pay. - The second commit documents that unconditional loop in place. It is deliberate for CI, where the persistent self-hosted `$CARGO_TARGET_DIR` is the only authoritative cache and anything a cache restore leaves in the ephemeral workspace `target/` is discarded outright, so the behaviour is left alone and only the intent and the local-execution hazard are now written down. Closes #1047
7 tasks done
inureyes
added a commit
that referenced
this pull request
Aug 7, 2026
## Summary `loading::vlm::gemma_unified::tests::unified_sanitize_quantized_split_dequant_equivalence` was the only failing test on `main`. The root cause is a real defect, but not in the gemma4_unified sanitize split: `mlx::core::dequantize` silently returns wrong values whenever its `biases` input is not row-contiguous. The test expectation (exact equality, `assert_eq!(max_abs_diff(...), 0.0)`) was sound and is unchanged. ## Root cause The sanitize split is correct. Slicing the fused quantized `gate_up_proj` on the output axis selects exactly the right `weight` / `scales` / `biases` bytes: dequantizing a row-contiguous copy of the sanitize output reproduces `split(dequantize(W))` bit for bit on both legs. Isolating the three inputs one at a time shows the result is wrong if and only if `biases` is strided, regardless of the layout of `weight` and `scales`. `quantize_impl` on the MLX Metal backend (<https://github.com/ml-explore/mlx/blob/main/mlx/backend/metal/quantized.cpp>) binds `w` (buffer 0) and `scales` (buffer 1) on the shared compute encoder and only then calls `ensure_row_contiguous` on `biases`. For a strided `biases`, that call enqueues a copy kernel on the same encoder, and the copy rebinds buffers 0 and 1 to its own operands. The dequantize dispatch that follows reads `w` and `scales` from the copy's operands instead of its own, so the affine kernel computes `biases * (q + 1)` with `q` decoded from the raw bias bytes. That model matches the observed numbers exactly: the gate leg's minimum was `-3.344`, which is `min(biases) * 16`, and its maximum was `-0.1793`, which is `max(biases) * 1`. This is an upstream ordering defect, not a contract on the caller. `QuantizedMatmul::eval_gpu` in the same file hoists every `ensure_row_contiguous*` above its first binding, and so did `fast::Quantize::eval_gpu` until ml-explore/mlx#3757 folded it into the shared `quantize_impl`. ## Regression window The test has been red since `94a8608a` (#1046, 2026-08-06), which bumped the MLX pin from `b7c3dd6d` (2026-07-17) to `2c46b953` (2026-08-05). The upstream commit that reordered the bindings, `a79332de6` (ml-explore/mlx#3757), is an ancestor of `2c46b953` and is not an ancestor of `b7c3dd6d`; the pre-bump tree still has all three contiguity copies hoisted above the first `set_input_array`. The test passed from the day it was added in #156 until that pin bump, so it was red for two days, not silently red for months. The window was established by reading `fast::Quantize::eval_gpu` in both pinned MLX trees (both are present in the local FetchContent clone) rather than by rebuilding mlxcel at the old pin, which would have cost a cold MLX C++ build. ## What changed - `src/lib/mlxcel-core/cpp/mlx_cxx_bridge.cpp`: the `dequantize` shim forces `biases` row-contiguous before calling `mlx::core::dequantize`. `biases` is the only input whose contiguity copy is enqueued after a binding, so it is the only one guarded, and `mlx::core::contiguous` shares the buffer instead of copying when the input is already row-contiguous, so the common path is unchanged. The comment records the mechanism and the upstream reference for the next MLX bump. - `src/lib/mlxcel-core/src/ffi_tests.rs`: adds `dequantize_commutes_with_output_axis_slice_on_strided_inputs`, which pins the guard at the FFI boundary where the defect actually lives, covering both halves of the split. - `src/loading/vlm_gemma_unified.rs`: doc-comment only. `split_fused_gate_up_component` keeps emitting strided views rather than materializing copies, because `gather_qmm` and `quantized_matmul` receive per-input batch strides explicitly and consume the views zero-copy; materializing would add a second copy of every fused expert tensor at load time for no correctness gain. The comment now says consumers must not assume row-contiguity there. ## Notes - Both legs confirmed. The `up_proj` assertion that the earlier failure masked also passes; before the fix it was wrong by 4.0477 against an expected 0. - `mistral4` and `qwen3_5_moe` split fused MoE experts the same way and therefore also hand strided quantized triplets downstream. They are not changed here: the guard is at the `dequantize` boundary, so it covers every caller. - The other in-tree `mlx::core::dequantize` call sites (`mlx_cxx_nemotron.cpp`, the fused-MoE parity reference in `mlx_cxx_kernels.cpp`, and `quantized_embedding`) feed `biases` from `take` or `astype`, both of which produce row-contiguous output, so none of them were exposed. ## Test plan - [x] `cargo test --profile test-fast --features metal,accelerate -p mlxcel --lib unified_sanitize` (6 passed, 0 failed; was 5 passed, 1 failed) - [x] `cargo test --profile test-fast --features metal,accelerate -p mlxcel-core --lib dequantize_commutes_with_output_axis_slice_on_strided_inputs` (passed) - [x] Negative control: the same test fails with a 3.71 max absolute difference when the guard is reverted, so it is not vacuous - [x] `cargo test --profile test-fast --features metal,accelerate -p mlxcel --test granite4_vision_parity` (3 passed, real 4-bit checkpoint) - [x] `cargo fmt --check` - [x] `cargo clippy --profile test-fast --features metal,accelerate -p mlxcel-core --lib --tests -- -D warnings` - [x] `make verify-test` Closes #1074 ## Review follow-ups A second commit adds comments only, no behavior change: the three call sites that reach `mlx::core::dequantize` directly (`quantized_embedding` in the bridge, the nemotron embedding path, and the fused-MoE parity reference) now record inline that they are safe only because `biases` comes from `take` / `astype`; the shim's cost claim is corrected (`mlx::core::contiguous` always adds a `Contiguous` node, and the elision happens at eval and only for a standalone row-contiguous input); the shim records why the guard is unconditional rather than gated on `flags().row_contiguous` (flags are not meaningful before eval, so the obvious short-circuit would skip the copy exactly when it is needed); and the Rust declaration states the normalization where callers look. ## Not covered - **`gather_qmm` on strided expert weights is not pinned by a test.** This PR deliberately keeps `split_fused_gate_up_component` emitting strided views, and that decision rests on `QuantizedMatmul::eval_gpu` and `GatherQMM::eval_gpu` hoisting every `ensure_row_contiguous_matrix` above their first binding and passing per-input batch strides explicitly. That was verified by reading the pinned MLX source, not by a test. A future MLX bump that regresses it the way ml-explore/mlx#3757 regressed `dequantize` would corrupt quantized MoE inference for gemma4_unified, mistral4, and qwen3_5_moe silently. A companion FFI test comparing `gather_qmm` on sliced expert weights against `gather_qmm` on materialized copies would close that gap and is worth a follow-up issue. - **The MLX bump checklist was not updated.** `CLAUDE.md` carries the post-bump re-validation list and would be the right home for "re-check the `dequantize` biases guard, regression test `dequantize_commutes_with_output_axis_slice_on_strided_inputs`". It is gitignored in this repository, so the entry cannot land through this PR and is left to the maintainer. ## Observed while verifying `make verify-test` was run twice on this branch. Both runs were green except for one real-model VLM parity test asserting `text-only logits must be finite`, and it was a **different** test each time (`granite4_vision_parity` in the first run, `hunyuan_vl_parity` in the second). Each passes in isolation, including its full test binary. That looks like load-dependent flakiness in the real-checkpoint parity tests when the whole workspace runs concurrently, not a regression from this change; worth its own issue.
be-free-snowboy
pushed a commit
to freestylejs/snowbrain-mlxcel
that referenced
this pull request
Oct 27, 2026
lablup#1046) ## Summary Bumps the MLX pin from `b7c3dd6d` (2026-07-17) to `2c46b953` (2026-08-05 upstream main), 58 commits, and re-derives the three source overlays that upstream touched in that window. Closes lablup#1042 ## Why Fixes in range that land on paths this repo ships: | Upstream | What | |---|---| | #3854 | incorrect nvfp4 `quantized_matmul` through the split-K path (NVFP4 direct transcode is mlxcel's default on all builds) | | #3893 | CUDA gemm conv unfold grid overflow past 65,535 output positions | | #3934 | Metal NVFP4 scale reduction per 16-lane group | | #3888 | Metal `gemv_wide` for few-row fp16/bf16 matmuls | | #3943 | head dimension 96 in Metal full attention | | #3898, #3813 | `prod` dtype promotion on a size-1 axis, shapeless matmul with dynamic batch dims | ## Overlay audit mlxcel applies whole-file overlays via `configure_file(... COPYONLY)`, so an overlay whose upstream counterpart moved silently reverts that upstream change. That is how a broken RMSNorm kernel shipped in lablup#824 / lablup#829 and took until lablup#830 / lablup#831 to undo. All 21 overlays were diffed against the upstream range; three conflicted. | Overlay | Upstream change | Resolution | |---|---|---| | `patches/mlx/backend/cuda/quantized/quantized.cpp` | #3757 inserted `const std::optional<array>& global_scale` into `qmm_naive` at position 5 | resolved by hand: mlxcel wraps both call sites in `run_row_chunked` / `run_batch_chunked`, which `git merge-file` cannot reconcile with an inline insertion. Both sites verified as 13 arguments in declaration order | | `patches/mlx/backend/metal/kernels/utils.h` | #3970 templated complex scalar lanes, #3963 explicit `thread` address space for metal 4.1 | clean three-way merge, no overlap with mlxcel's delta | | `patches-cuda/ops.cpp` | 189 lines | clean three-way merge, mlxcel's 57-line delta intact | `quantized/qmm/qmm.h` also changed but is not overlaid, so mlxcel takes the new declaration. Its full diff across the range is the single `global_scale` line, so `qmm_sm90`, `qmm_sm80`, `qmv`, `gather_qmv` and `fp_qmv` are unaffected. The `qmm_sm80.cuh` race-fix overlay from lablup#910 is untouched upstream and still required. The stale `Synced to upstream e9463bb` header comment is updated to the real base, with an explicit instruction to re-verify these call sites on the next bump. ## The pin has two sources of truth `GIT_TAG` in `src/lib/mlx-cpp/CMakeLists.txt:95` and `MLX_EXPECTED_COMMIT` in `src/lib/mlxcel-core/build.rs:217`. The latter drives `purge_stale_mlx_cache`. Updating only the CMakeLists leaves the cache marker matching `MLX_EXPECTED_COMMIT`, so `_deps/` is judged valid, never purged, and the build silently compiles the OLD MLX while looking bumped. Both moved together, verified three ways with no trace of the old SHA: - binary embedded commit: `2c46b953…` - `_deps/.mlx-build-commit` marker: `2c46b953…` - fetched `_deps/mlx-src` HEAD: `2c46b953…` Worth a follow-up to collapse this to one source of truth or add a build-time assert. ## Validation (GB10, sm_121) - release CUDA build clean - `cargo fmt --all -- --check` clean - `fused_norm_parity` 8/8, `turbo_tests` 97/97 - greedy smoke on `llama-3.2-1b-4bit`, `qwen3-4b-4bit`, `deepseek-v2-lite-4bit`: coherent, and byte-identical run to run with a warm kernel cache. That determinism also confirms the `qmm_sm80` race-fix overlay still holds on the new base. ## Two pre-existing failures, separated by a control build A control build at the old pin (`_deps` re-fetched to `b7c3dd6d`, verified) reproduces both, so neither is caused by this bump. | Failure | New pin | Old pin (control) | |---|---|---| | `fused_rope_parity` 2 of 9 | offset 131069 rms 2.497e-3, offset 131071 rms 2.684e-3, tol 2.0e-3 | same digits, same 7 pass / 2 fail | | full lib suite abort under graph capture | `cuGraphAddKernelNode` invalid argument, then `cudaStreamEndCapture` previous error during capture (site and message vary per run) | aborts the same way | For the rope case the numbers match bit for bit, which also shows that path carries no exposure to this change. Upstream touched no rope file in the range. Both are left for separate issues rather than folded in here.
be-free-snowboy
pushed a commit
to freestylejs/snowbrain-mlxcel
that referenced
this pull request
Oct 27, 2026
## Summary The pinned MLX commit was written out in three places with nothing checking that they agreed, so bumping one produced a build that looked bumped and compiled the old MLX. This makes `GIT_TAG` in `src/lib/mlx-cpp/CMakeLists.txt` the only place the value lives (the issue's option 1), adds a post-build check that the checkout which actually landed on disk is that commit (option 3), and removes the third location, which had already gone stale. The pin value is unchanged at `2c46b953db88965c4270cc7306eda6887a3247f2`. ## The third location, and why it matters The issue names two locations. There were three, and `.github/workflows/release.yml` was already a generation behind: it carried `b7c3dd6d27f45b5365b08a840310187dc503f1db` while the tree had moved to `2c46b953` in lablup#1046, and the comment above it called this "the three-place pin" and pointed at a documentation file that does not exist in this repository (the checklist is in `CONTRIBUTING.md`). The effect there was fail-safe rather than silent, since the mismatched literal made every marker comparison fail, but the practical result was that every release build purged and rebuilt MLX from scratch while the workflow reported a commit the tree no longer used. So the silent-divergence failure the issue predicts had already happened once, in the fail-safe direction. ## What changed - `src/lib/mlxcel-core/build.rs`: `MLX_EXPECTED_COMMIT` is deleted. `main()` resolves the pin once from `../mlx-cpp/CMakeLists.txt` before anything compares against it, and threads that one value through `build_mlx`, `purge_stale_mlx_cache`, the new `verify_fetched_mlx_head`, `mark_mlx_cache_valid` and the `cargo:rustc-env=MLXCEL_MLX_COMMIT` export. A parse failure panics with a message naming the file and what it expected. No duplicate `rerun-if-changed` was added; the existing one at what is now line 204 is unchanged. - `src/lib/mlxcel-core/build_support/mlx_pin.rs` (new): the pure logic. Parsing strips CMake comments first, then scopes itself to the `FetchContent_Declare` block whose `GIT_REPOSITORY` names `ml-explore/mlx`, and treats zero blocks, two blocks, zero `GIT_TAG`s, two `GIT_TAG`s, and any value that is not a 40-character lowercase hex SHA as hard failures with distinct actionable messages. Uppercase is rejected too: `git rev-parse` prints lowercase and the marker is compared byte for byte, so an uppercase pin would make every cache look stale forever. - Post-build HEAD verification: after CMake returns and **before** the marker is written, `_deps/mlx-src`'s HEAD is read and a disagreement fails the build, telling the reader to delete `_deps/` to force a refetch. Verifying before marking is the point: a tree that failed verification is never blessed as valid. A missing `.git` or a host without `git` emits a `cargo:warning` and skips, so vendored and offline trees still build. The `.git` existence probe is load-bearing rather than a fast path, because `git -C` walks up to the nearest enclosing repository and `_deps/` sits inside the mlxcel checkout, so without it a metadata-free tree would answer with mlxcel's own HEAD and be reported as a pin mismatch. There is a unit test for exactly that. - `.github/workflows/release.yml`: the workflow-scope `env: MLX_EXPECTED_COMMIT` literal is gone. Both "Validate MLX build cache" steps derive the pin from the CMake file through the new `scripts/ci/mlx_pinned_commit.sh`, and refuse to run the comparison against an empty value, since an empty pin would mark every cache stale and purge all of them rather than validate any. Purge behaviour is otherwise byte-identical. The stale documentation pointer and the stale SBOM comment at the bottom of the file are fixed. - `scripts/ci/mlx_pinned_commit.sh` (new): the shell half of the parser. Same scoping (comments stripped, matched on the MLX repository), exits non-zero unless exactly one 40-character lowercase hex SHA resolves. Deliberately free of regex interval syntax, which is not portable across the awk implementations on the GitHub-hosted Linux and self-hosted macOS runners; the SHA shape is validated in shell instead. `MLX_PIN_CMAKE_FILE` overrides the input file for testing. - `src/lib/mlx-cpp/CMakeLists.txt`: a comment above `GIT_TAG` naming both parsers and stating that the literal `GIT_TAG <40-char lowercase hex sha>` shape must stay parseable. Verified that CMake accepts a comment inside the command invocation and drops it from `ARGN`. - Docs: `CONTRIBUTING.md`'s "Bumping the MLX upstream pin" section is rewritten for the one-file reality (the Metal kernel re-validation checklist after it is kept verbatim), and `docs/benchmark_results/model_tests.md` now points at the CMake `GIT_TAG` instead of the deleted constant. ## Test mechanism, and why it is not a build-dependency The logic lives in `mlxcel-core/build_support/mlx_pin.rs`, which `build.rs` includes with `#[path]` and a new leaf crate `mlxcel-mlx-pin` includes as a library. That crate has no production role and nothing depends on it; it exists purely so `cargo test -p mlxcel-mlx-pin` can cover the logic in about a second without compiling `mlxcel-core` and therefore without building MLX C++. The more conventional shape, a `[build-dependencies]` entry, was rejected on purpose: a new build-dependency feeds the build-script unit's metadata hash, that hash names `mlxcel-core`'s `OUT_DIR` (`target/<profile>/build/mlxcel-core-<hash>/out`), and moving it strands every warm `_deps/` MLX checkout under the old hash. Including the source by path leaves the hash untouched. This is reasoning from Cargo's metadata model rather than an experiment; testing it would have required the rebuild it is meant to avoid. Adding a fifth workspace member made four "four members" statements stale, in `Cargo.toml`, `Makefile`, `CONTRIBUTING.md` and `docs/installation.md`. Those are corrected in the same commit rather than left to rot, which is the same defect class as the issue itself. ## Test plan 25 unit tests, all passing, all runnable without an MLX build: - [x] `cargo test -p mlxcel-mlx-pin` (25 passed): happy path, quoted value, comment inside the block, commented-out declaration, missing `GIT_TAG`, branch name, short SHA, uppercase SHA, two MLX declarations, two `GIT_TAG`s in one declaration, an unrelated dependency carrying its own `GIT_TAG` (the wrong-block case), no MLX declaration, unreadable file, error-message content, marker-matches, marker-matches-with-trailing-whitespace, marker-differs, marker-absent, HEAD-matches, HEAD-differs, no-git-metadata, missing directory, no-walk-up-into-enclosing-repository, and the recovery text of the mismatch message. The HEAD tests run against real `git init` repositories, not a mocked `git`. - [x] One test parses the actual in-tree `CMakeLists.txt` and asserts the result is a well-formed SHA. It deliberately does **not** assert the literal, because doing so would put a second copy of the pin in the tree and recreate this issue inside the test suite. - [x] `cargo clippy -p mlxcel-mlx-pin --all-targets -- -D warnings`: clean. - [x] `cargo fmt --all -- --check`: clean. - [x] `build.rs` type-checks and passes `clippy -D warnings` with and without the `cuda` feature. Checked by compiling it verbatim in an isolated scratch crate with `cmake` and `cxx-build` as ordinary dependencies, since checking it in place re-enters CMake. This also confirms the build.rs view is free of `dead_code`, which the `#![allow(dead_code)]` in the shared module exists to guarantee for helpers only the tests call. - [x] `scripts/ci/mlx_pinned_commit.sh` exercised against the real file plus eight synthetic fixtures via `MLX_PIN_CMAKE_FILE`: correct SHA on the real file and on the other-dependency and commented-out cases, non-zero with a specific message on missing tag, branch name, short SHA, two MLX declarations, no MLX declaration and missing file. - [x] Both parsers cross-checked: the Rust parser, the awk parser and the value committed on `main` all resolve to `2c46b953db88965c4270cc7306eda6887a3247f2`, so this change is value-preserving. - [x] `release.yml` parses as YAML and both "Validate MLX build cache" steps still sit immediately after `Checkout code`, which is what makes the repo available to the extraction script. ## What is not verified here - The full `cargo build --release --features cuda` is run separately, outside this PR's local verification, because every profile/feature pair that is not already warm rebuilds MLX from source. - The issue's second acceptance criterion, bumping to a different commit and confirming the fetched `_deps/mlx-src` HEAD, the `.mlx-build-commit` marker and the embedded `MLXCEL_MLX_COMMIT` all move together, was not executed as a deliberate two-rebuild experiment. What supports it: the five caches present on the development host before this change spanned both pin generations (two at `2c46b953`, three at `b7c3dd6d`) with the marker and the `mlx-src` HEAD agreeing in every one, which is the invariant the new HEAD check enforces rather than assumes. - A note on how the cold build became available: while validating this change I ran the `Validate MLX build cache` step body extracted from `release.yml` against a scratch `CARGO_TARGET_DIR`, from the repository root. The first loop honoured the scratch directory; the second loop is written against the literal relative path `target` and consults no marker, so it deleted all five `_deps/` trees under the real `target/` (about 3.8 GB) with `rm -rf`. They are unrecoverable. The upside is incidental, not planned: with `_deps/` absent, the verification build is a genuine cold fetch, so the fetched HEAD, the marker and the embedded `MLXCEL_MLX_COMMIT` are all observed being produced from the single source in one run. That is what turned the second criterion from prohibitively expensive into a by-product, at a cost nobody chose to pay. - The second commit documents that unconditional loop in place. It is deliberate for CI, where the persistent self-hosted `$CARGO_TARGET_DIR` is the only authoritative cache and anything a cache restore leaves in the ephemeral workspace `target/` is discarded outright, so the behaviour is left alone and only the intent and the local-execution hazard are now written down. Closes lablup#1047
be-free-snowboy
pushed a commit
to freestylejs/snowbrain-mlxcel
that referenced
this pull request
Oct 27, 2026
## Summary `loading::vlm::gemma_unified::tests::unified_sanitize_quantized_split_dequant_equivalence` was the only failing test on `main`. The root cause is a real defect, but not in the gemma4_unified sanitize split: `mlx::core::dequantize` silently returns wrong values whenever its `biases` input is not row-contiguous. The test expectation (exact equality, `assert_eq!(max_abs_diff(...), 0.0)`) was sound and is unchanged. ## Root cause The sanitize split is correct. Slicing the fused quantized `gate_up_proj` on the output axis selects exactly the right `weight` / `scales` / `biases` bytes: dequantizing a row-contiguous copy of the sanitize output reproduces `split(dequantize(W))` bit for bit on both legs. Isolating the three inputs one at a time shows the result is wrong if and only if `biases` is strided, regardless of the layout of `weight` and `scales`. `quantize_impl` on the MLX Metal backend (<https://github.com/ml-explore/mlx/blob/main/mlx/backend/metal/quantized.cpp>) binds `w` (buffer 0) and `scales` (buffer 1) on the shared compute encoder and only then calls `ensure_row_contiguous` on `biases`. For a strided `biases`, that call enqueues a copy kernel on the same encoder, and the copy rebinds buffers 0 and 1 to its own operands. The dequantize dispatch that follows reads `w` and `scales` from the copy's operands instead of its own, so the affine kernel computes `biases * (q + 1)` with `q` decoded from the raw bias bytes. That model matches the observed numbers exactly: the gate leg's minimum was `-3.344`, which is `min(biases) * 16`, and its maximum was `-0.1793`, which is `max(biases) * 1`. This is an upstream ordering defect, not a contract on the caller. `QuantizedMatmul::eval_gpu` in the same file hoists every `ensure_row_contiguous*` above its first binding, and so did `fast::Quantize::eval_gpu` until ml-explore/mlx#3757 folded it into the shared `quantize_impl`. ## Regression window The test has been red since `94a8608a` (lablup#1046, 2026-08-06), which bumped the MLX pin from `b7c3dd6d` (2026-07-17) to `2c46b953` (2026-08-05). The upstream commit that reordered the bindings, `a79332de6` (ml-explore/mlx#3757), is an ancestor of `2c46b953` and is not an ancestor of `b7c3dd6d`; the pre-bump tree still has all three contiguity copies hoisted above the first `set_input_array`. The test passed from the day it was added in lablup#156 until that pin bump, so it was red for two days, not silently red for months. The window was established by reading `fast::Quantize::eval_gpu` in both pinned MLX trees (both are present in the local FetchContent clone) rather than by rebuilding mlxcel at the old pin, which would have cost a cold MLX C++ build. ## What changed - `src/lib/mlxcel-core/cpp/mlx_cxx_bridge.cpp`: the `dequantize` shim forces `biases` row-contiguous before calling `mlx::core::dequantize`. `biases` is the only input whose contiguity copy is enqueued after a binding, so it is the only one guarded, and `mlx::core::contiguous` shares the buffer instead of copying when the input is already row-contiguous, so the common path is unchanged. The comment records the mechanism and the upstream reference for the next MLX bump. - `src/lib/mlxcel-core/src/ffi_tests.rs`: adds `dequantize_commutes_with_output_axis_slice_on_strided_inputs`, which pins the guard at the FFI boundary where the defect actually lives, covering both halves of the split. - `src/loading/vlm_gemma_unified.rs`: doc-comment only. `split_fused_gate_up_component` keeps emitting strided views rather than materializing copies, because `gather_qmm` and `quantized_matmul` receive per-input batch strides explicitly and consume the views zero-copy; materializing would add a second copy of every fused expert tensor at load time for no correctness gain. The comment now says consumers must not assume row-contiguity there. ## Notes - Both legs confirmed. The `up_proj` assertion that the earlier failure masked also passes; before the fix it was wrong by 4.0477 against an expected 0. - `mistral4` and `qwen3_5_moe` split fused MoE experts the same way and therefore also hand strided quantized triplets downstream. They are not changed here: the guard is at the `dequantize` boundary, so it covers every caller. - The other in-tree `mlx::core::dequantize` call sites (`mlx_cxx_nemotron.cpp`, the fused-MoE parity reference in `mlx_cxx_kernels.cpp`, and `quantized_embedding`) feed `biases` from `take` or `astype`, both of which produce row-contiguous output, so none of them were exposed. ## Test plan - [x] `cargo test --profile test-fast --features metal,accelerate -p mlxcel --lib unified_sanitize` (6 passed, 0 failed; was 5 passed, 1 failed) - [x] `cargo test --profile test-fast --features metal,accelerate -p mlxcel-core --lib dequantize_commutes_with_output_axis_slice_on_strided_inputs` (passed) - [x] Negative control: the same test fails with a 3.71 max absolute difference when the guard is reverted, so it is not vacuous - [x] `cargo test --profile test-fast --features metal,accelerate -p mlxcel --test granite4_vision_parity` (3 passed, real 4-bit checkpoint) - [x] `cargo fmt --check` - [x] `cargo clippy --profile test-fast --features metal,accelerate -p mlxcel-core --lib --tests -- -D warnings` - [x] `make verify-test` Closes lablup#1074 ## Review follow-ups A second commit adds comments only, no behavior change: the three call sites that reach `mlx::core::dequantize` directly (`quantized_embedding` in the bridge, the nemotron embedding path, and the fused-MoE parity reference) now record inline that they are safe only because `biases` comes from `take` / `astype`; the shim's cost claim is corrected (`mlx::core::contiguous` always adds a `Contiguous` node, and the elision happens at eval and only for a standalone row-contiguous input); the shim records why the guard is unconditional rather than gated on `flags().row_contiguous` (flags are not meaningful before eval, so the obvious short-circuit would skip the copy exactly when it is needed); and the Rust declaration states the normalization where callers look. ## Not covered - **`gather_qmm` on strided expert weights is not pinned by a test.** This PR deliberately keeps `split_fused_gate_up_component` emitting strided views, and that decision rests on `QuantizedMatmul::eval_gpu` and `GatherQMM::eval_gpu` hoisting every `ensure_row_contiguous_matrix` above their first binding and passing per-input batch strides explicitly. That was verified by reading the pinned MLX source, not by a test. A future MLX bump that regresses it the way ml-explore/mlx#3757 regressed `dequantize` would corrupt quantized MoE inference for gemma4_unified, mistral4, and qwen3_5_moe silently. A companion FFI test comparing `gather_qmm` on sliced expert weights against `gather_qmm` on materialized copies would close that gap and is worth a follow-up issue. - **The MLX bump checklist was not updated.** `CLAUDE.md` carries the post-bump re-validation list and would be the right home for "re-check the `dequantize` biases guard, regression test `dequantize_commutes_with_output_axis_slice_on_strided_inputs`". It is gitignored in this repository, so the entry cannot land through this PR and is left to the maintainer. ## Observed while verifying `make verify-test` was run twice on this branch. Both runs were green except for one real-model VLM parity test asserting `text-only logits must be finite`, and it was a **different** test each time (`granite4_vision_parity` in the first run, `hunyuan_vl_parity` in the second). Each passes in isolation, including its full test binary. That looks like load-dependent flakiness in the real-checkpoint parity tests when the whole workspace runs concurrently, not a regression from this change; worth its own issue.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Bumps the MLX pin from
b7c3dd6d(2026-07-17) to2c46b953(2026-08-05 upstream main), 58 commits, and re-derives the three source overlays that upstream touched in that window.Closes #1042
Why
Fixes in range that land on paths this repo ships:
quantized_matmulthrough the split-K path (NVFP4 direct transcode is mlxcel's default on all builds)gemv_widefor few-row fp16/bf16 matmulsproddtype promotion on a size-1 axis, shapeless matmul with dynamic batch dimsOverlay audit
mlxcel applies whole-file overlays via
configure_file(... COPYONLY), so an overlay whose upstream counterpart moved silently reverts that upstream change. That is how a broken RMSNorm kernel shipped in #824 / #829 and took until #830 / #831 to undo. All 21 overlays were diffed against the upstream range; three conflicted.patches/mlx/backend/cuda/quantized/quantized.cppconst std::optional<array>& global_scaleintoqmm_naiveat position 5run_row_chunked/run_batch_chunked, whichgit merge-filecannot reconcile with an inline insertion. Both sites verified as 13 arguments in declaration orderpatches/mlx/backend/metal/kernels/utils.hthreadaddress space for metal 4.1patches-cuda/ops.cppquantized/qmm/qmm.halso changed but is not overlaid, so mlxcel takes the new declaration. Its full diff across the range is the singleglobal_scaleline, soqmm_sm90,qmm_sm80,qmv,gather_qmvandfp_qmvare unaffected. Theqmm_sm80.cuhrace-fix overlay from #910 is untouched upstream and still required.The stale
Synced to upstream e9463bbheader comment is updated to the real base, with an explicit instruction to re-verify these call sites on the next bump.The pin has two sources of truth
GIT_TAGinsrc/lib/mlx-cpp/CMakeLists.txt:95andMLX_EXPECTED_COMMITinsrc/lib/mlxcel-core/build.rs:217. The latter drivespurge_stale_mlx_cache. Updating only the CMakeLists leaves the cache marker matchingMLX_EXPECTED_COMMIT, so_deps/is judged valid, never purged, and the build silently compiles the OLD MLX while looking bumped. Both moved together, verified three ways with no trace of the old SHA:2c46b953…_deps/.mlx-build-commitmarker:2c46b953…_deps/mlx-srcHEAD:2c46b953…Worth a follow-up to collapse this to one source of truth or add a build-time assert.
Validation (GB10, sm_121)
cargo fmt --all -- --checkcleanfused_norm_parity8/8,turbo_tests97/97llama-3.2-1b-4bit,qwen3-4b-4bit,deepseek-v2-lite-4bit: coherent, and byte-identical run to run with a warm kernel cache. That determinism also confirms theqmm_sm80race-fix overlay still holds on the new base.Two pre-existing failures, separated by a control build
A control build at the old pin (
_depsre-fetched tob7c3dd6d, verified) reproduces both, so neither is caused by this bump.fused_rope_parity2 of 9cuGraphAddKernelNodeinvalid argument, thencudaStreamEndCaptureprevious error during capture (site and message vary per run)For the rope case the numbers match bit for bit, which also shows that path carries no exposure to this change. Upstream touched no rope file in the range. Both are left for separate issues rather than folded in here.