Repository navigation
Conversation
inureyes
added a commit
that referenced
this pull request
Oct 6, 2026
This was referenced Oct 6, 2026
`GatherQMM::eval_gpu` passes `x.size() / K` to `gather_qmm_rhs` as its row count. That is the row count of x before `gather_qmm_rhs` broadcasts it against the indices, so when one activation row is gathered against B sorted slots (M == 1, B >= 16, B / E >= 4) the grid and the kernel row bound cover one row and rows 1..B-1 of the freshly allocated output are never written. That is the nightly failure on the M1 Ultra runner since 2026-09-30: the `SortedSharedActivation` case of `mxfp4/mxfp8_gather_qmm_matches_host_reference` reads uninitialized memory (relative L2 7.4e32, non-finite values), while the same case passes on MLX's CPU backend in the same run. The overlay now passes `B * M`, which equals the old value whenever x is already expanded (every `gather_sort` caller, so production MoE paths are unchanged) and covers the NAX variant too. Upstream ml-explore/mlx main carries the same line. Not verified on Metal: this host is CUDA, where the overlay is not compiled. Refs #1599
inureyes
force-pushed
the
fix/issue-1599-metal-gather-qmm-rhs-rows
branch
from
October 7, 2026 14:34
838936a to
2520219
Compare
Member
Author
|
Metal validation passed on Apple M1 Ultra (128 GB), macOS 27.0 (26A428), full Xcode 27 RC, after rebasing this PR onto main
This satisfies #2158 item 4's local gate. Keep #1599 open until a post-merge Metal nightly |
inureyes
marked this pull request as ready for review
October 7, 2026 14:39
Keep the broadcast-row Metal gather fix and make split-MTP staging tests follow canonical macOS paths and the host filesystem's invalid-byte filename support. Filesystems that reject the name must now return the same I/O error without leaving staging residue. Update the bilingual report with the exact historical M1 evidence and current Linux build, focused-test, contract, and independent-review results. A fresh Metal nightly remains required before #1599 closes. Refs #1599
Include the merged Qwen3 Next prefill contract update before validating the Metal correction against the current default branch. The tested Metal kernel and portability test sources remain unchanged.
…gather-qmm-rhs-rows # Conflicts: # src/lib/mlx-cpp/patches/mlx/backend/metal/quantized.cpp
Normalize local cooperative-tensor operand types before passing them to the MPP destination factory. Register the pinned-source overlay so the new MLX gated-delta kernel compiles under Xcode 27.
Carry the upstream Apple copyright in the whole-file overlay so the repository license contract recognizes its source provenance.
This branch has not been deployed
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.
Fix Metal
gather_qmm_rhsto use the broadcast output row count, so shared activations gathered across sorted slots write every output row. Make two split-MTP tests respect canonical macOS paths and filesystems that reject invalid-byte names; production path and symlink safeguards are unchanged.The branch includes main
4b0c4519and MLX pin59d600b5e. The #1599quantized.cppoverlay preserves the pin's global-scale changes and theB * Mrow-count fix.Candidate Metal nightly 38430123889 failed while compiling the pinned MLX sources at source head
4323eb39, before the #1599 numerical tests ran. Xcode 27 rejected thethreadaddress-space qualifiers carried bydecltype(ct_a)anddecltype(ct_b)into three gated-delta MPP destination cooperative-tensor factories. The follow-up overlay strips those qualifiers withmetal::remove_addrspace_t, matching the pinned tree's working steel NAX helpers and the compatibility issue addressed by upstream MLX #4594. It does not change the cooperative-tensor dtype, descriptor, mode, or algorithm. Independent source review of the functional fix at1b0ec371found no correctness, security, or performance issue; head7d545981adds the preserved upstream Apple copyright required by the repository license contract. A new candidate Metal nightly is required; no current-head Metal pass is claimed yet.The earlier candidate Metal nightly 38418557399 passed at source head
2dbffe6d: formatting, Clippy, 12,283 tests (390 ignored), and 51 video tests. That evidence predates the material MLX pin update from81ba1c6ato59d600b5e.Linux focused portability tests pass 2/2, with 123 Python contract tests and 49 WebUI fixtures. These runtime results also predate the material MLX pin update. Independent source review of the rebased row-count overlay at
4323eb39and the gated-delta compatibility overlay at1b0ec371found no source or security issue. The full Python tooling contract passes 123/123 at7d545981, including the source-license check.Historical checkpoint evidence remains separate: at exact head
25202198on M1 Ultra, all five focused Metal mxfp tests passed and Qwen3-30B-A3B-4bit greedy 64-token output matched base and repeated base byte for byte (SHA-256d2f40b2d5831c3d7790d467fb51123f0d1f615fb7300aa2f92d86a9d8831a0de). The checkpoint smoke was not rerun at the current head.Keep #1599 open through merge. Close it only after the final integrated main Metal nightly passes.
Refs #1599