Repository navigation
Fix fp quantized matmul corruption when the quantized dim is not a multiple of 32 - #3912
Conversation
|
Coordination note from M5/NAX testing in #3922: This PR remains the authoritative fix for NVFP4's legal 16-wide tails and the required plain-kernel bounded load/store work. I deliberately did not duplicate those changes. The independent M5 matrix clarified the merge interaction:
Independent native M5 tests for #3922 kept affine/MXFP4/MXFP8 tail errors below 4.7e-4; untouched NVFP4 non-64 K remains failing, which confirms the scope split rather than contradicting this PR. Security scope for the related reports: the observed behavior is same-process incorrect tensor values and, in #3856, same-process allocator-pool reuse. We found no cross-process disclosure, arbitrary code execution, sandbox escape, or other trust-boundary crossing; current evidence supports numerical correctness bugs, not cybersecurity vulnerabilities. |
|
Thanks for the heads up, and for running this on M5. That's the piece I couldn't test. Scope split looks right to me. This PR fixes the plain kernels for NVFP4 dims that are We both touch quantized.cpp and test_quantized.py so there's a conflict either way. |
|
Thanks — this confirms both the scope split and the observed M5 NVFP4 path. The cleanest sequence is #3922 first, then rebase #3912 as offered: #3922 supplies the sorted |
| loader_w.next(); | ||
| } | ||
| } | ||
| if (num_k > 0) { |
There was a problem hiding this comment.
You can putting this inside a if constexpr (group_size == 16) to avoid overhead on other quantization modes.
| threadgroup_barrier(mem_flags::mem_threadgroup); | ||
| loader_x.load_safe(short2(BK, num_els)); | ||
| loader_w.load_unsafe(); | ||
| loader_w.load_safe(short2(num_outs, BK)); |
There was a problem hiding this comment.
Similarly here if constexpr (group_size == 16) can be used to choose load_unsafe for faster load.
| // instead of reading past the source (partial tiles happen when the | ||
| // quantized dim is not a multiple of the block, e.g. nvfp4's | ||
| // group_size=16). | ||
| if (bi >= src_tile_dim.y || bj * pack_factor >= src_tile_dim.x) { |
There was a problem hiding this comment.
Should the condition depend on reduction_dim, i.e. bi >= src_tile_dim.y when reduction_dim == 0, and bi >= src_tile_dim.x when reduction_dim == 1?
| const Stream& s, | ||
| const std::string mode) { | ||
| if (metal::is_nax_available() && transpose && | ||
| if (metal::is_nax_available() && transpose && (K % 64 == 0) && |
There was a problem hiding this comment.
I'm not sure on this, should drop it if you can not test it.
9c8d78c to
3cd6f4a
Compare
|
All four addressed, rebased on main.
Same on the qmm_n side, with Whether the mask should depend on The same reasoning is why the NAX dispatch guard. Dropped, as asked. Revalidated on the result: three reproducers clean, 289/289 misalignment sweep, 4515 randomized differential cases across all modes, bit widths and dtypes with no failures, |
…ltiple of 32 nvfp4's group size of 16 is the only mode that lets the quantized dim of a matmul legally be a multiple of 16 but not of 32, and the fp quantized Metal kernels tile that dim by 32 without bounding the tail: - fp_qmm_t_impl looped over K_eff in full-BK steps with unbounded loads, so the final iteration read 16 columns of the next row's weights and scales (past the buffer for the last row) into the accumulator. - fp_qmm_n_impl stored full 32-column tiles for a 16-column tail; the spilled columns land on the next output row, racing the threadgroup that owns it (results differ run to run), and past the end of the output buffer for the last row. - QuantizedBlockLoader::load_safe compared the row index bi against the column bound src_tile_dim.x, so the K-remainder step of fp_gather_qmm_rhs zeroed the threads owning output columns 16..31 of each tile and silently dropped the last 16 K-elements from them. Bound the K loop of fp_qmm_t_impl with a remainder tile, clip the w loads and the stores of fp_qmm_n_impl to the valid columns, and make load_safe mask rows against .y and packed columns against .x, matching the convention the steel BlockLoader already uses. The row half of that mask is not nvfp4 only. With reduction_dim == 1 the old compare was bi >= src_tile_dim.x, that is bi >= BK == 32, while bi ranges over [0, BROWS) == [0, 32), so it never fired and any N that is not a multiple of BN read weight and scale rows past N. The store already clipped those columns so results were correct, but the read was out of bounds, and affine and the mxfp modes reach it too. The affine loader therefore gets the same change. A quantized dim is always a multiple of group_size, so it can only end in a partial tile when group_size does not divide the block. Only nvfp4 does; mxfp4 and mxfp8 use group_size 32 and every affine group size is a multiple of the block. The remainder tile and the column bound are selected with if constexpr so they are discarded at compile time for those modes, which keep load_unsafe on the aligned path.
3cd6f4a to
3a755bf
Compare
|
Hey, please see #4009. Thanks! |
|
Datapoint from M5 Max ( The good news: #4009 closes exactly that gap. With #4009 cherry-picked on top of |
|
This is specifically on metal. For cuda/cpu I don't see it: |
|
Cleaned up comments, removed the partial col filtering and reorganised the tests. The test previously did not cover all cases for all dtypes , and gating transpose logic on the input shape was confusing. |
zcbenz
left a comment
There was a problem hiding this comment.
@RohanGautam Thanks a lot for the review and cleanup!
…ltiple of 32 (ml-explore#3912) Co-authored-by: Rohan Gautam <rohan1gautam@gmail.com>
…tal; SPEC H15-H18; commit the results ADR 0037 records the decision: the MLX pin stays at d9add9d1 with ml-explore/mlx#3912, no revert and no compat patch restoring the old fp_qmm_t loop bound. It rests on correctness, not accuracy — the fixed kernel matches MLX-CUDA's golden delta exactly at 0.0898, and on all 1000 OCRBench v1 items the two kernels are indistinguishable (835 vs 833, exact McNemar p = 0.80). ADR 0036 already accepted the same 4 -> 3 move on 31b's 9px cell as the cost of the faithful computation on the GGUF path; two unrelated correctness fixes on two engines move that cell the same way. SPEC vision-harness-reuse gains four rules this investigation forced, each with an honest conformance row (three are not enforced): H15 one capture of a quality cell is one draw — measure a rate H16 both finetext arms count; the renderer shows only the suite arm (summarize_engine_compare.py falls back to ft_ only when absent) H17 a checkpoint is its MANIFEST digest — not its tag, and not its config digest, which local 31b-nvfp4 and 31b-mlx-bf16 share H18 cross-build accuracy needs a paired test on many items; a tier move needs the unquantized arm on the same binary Promotes the mlx-metal surface: README deploy line and the release matrix's mlx-metal row, generated by release_matrix.py from the build's own preflight run and spliced beside the CUDA row, whose run was never committed — regenerating the whole matrix here would have erased it. The matrix comment records why the two rows come from two hosts and two stamps (0.34.0-maxusai vs 0.34.1-dynres, flagged against ADR 0032). The deploy line also records that the Metal launchd environment does not set OLLAMA_MLX_DRAFT_UNDER_GRAMMAR=0, unlike the CUDA container. Committed results, assembled from the gitignored run captures by script (ADR 0012 rule 8), serialized one record per line: bench-runs/ocrbench-v1-1000-gemma4-31b-nvfp4-0340-vs-0332.json bench-runs/finetext-9px-31b-quant-reps-2026-09-18.json bench-runs/vision-campaign-2026-09-18-mlx8a7ba949.json preflight/runs/preflight-mlx-metal-0340-8a7ba949.json (force-added, as the 0.33.0 fold's closing Metal run was; the invalid --container native run of the same build is deliberately left out) BINARIES.md marks 8a7ba94 promoted and warns that the 0.33.2 rollback target reinstates the defect. The campaign doc records the manifest digest of every checkpoint it measured, per H17. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
…lay (#2298) ## Summary Bumps the pinned upstream `ml-explore/mlx` commit in `src/lib/mlx-cpp/CMakeLists.txt` from `81ba1c6a0e50a9268b931579c2d4f1158b9aab5a` to `59d600b5e64c238427d0f8d897ab7c682ef4d3d2` (14 upstream commits) and establishes a single machine-readable `// Modified from upstream MLX <sha> <upstream/relative/path>` base declaration across all 32 overlay files under `src/lib/mlx-cpp/patches/` and `src/lib/mlx-cpp/patches-cuda/`. ## Changes - **Phase 1 (Base declarations across all 32 overlays):** Determined each overlay's true copy or rebase commit by diffing against upstream `ml-explore/mlx` history (resolving `v0.31.1` to `ce45c525`, fixing `qmm_sm80.cu` where `Patched by mlxcel.` preceded `Modified from upstream MLX`, and recording the copy/rebase SHAs for the remaining files). - **Phase 2 (Rebase conflicting and drifted overlays onto `59d600b5e`):** - 3-way merged the four quantized overlays (`patches/mlx/backend/cuda/device/qmm_naive.cuh`, `patches/mlx/backend/cuda/quantized/qmm/qmv.cu`, `patches/mlx/backend/cuda/quantized/quantized.cpp`, `patches/mlx/backend/metal/quantized.cpp`) and the two `patches-cuda/` overlays (`fast.cpp`, `ops.cpp`) touched between `81ba1c6a` and `59d600b5e`, preserving every documented mlxcel delta. - Confirmed that upstream `dfe17bafb` (ml-explore/mlx#3912, fixing fp-quantized matmul corruption when the quantized dim is not a multiple of 32 in `mlx/backend/metal/kernels/fp_quantized.h` and `mlx/backend/metal/kernels/quantized.h`) is intact: neither header is overlaid in `src/lib/mlx-cpp/patches/mlx/backend/metal/kernels/`, and `patches/mlx/backend/metal/quantized.cpp` is 3-way merged onto `59d600b5e` (incorporating `3eb822cac` and `d9add9d11`) with only the three `MLXCEL_QMV_WIDE` hunks differing from `59d600b5e:mlx/backend/metal/quantized.cpp`. - Refreshed `patches/mlx/backend/cuda/reduce/{all_reduce,col_reduce,row_reduce}.cu` from `ce45c525` (`v0.31.1`) onto `59d600b5e` (folding in upstream `66f58032d` `const_cast` removal and `790be8aa8` 64-bit `all_reduce` planning for large inputs) so the drift check reports 0 upstream commits across all 32 overlays. - **Phase 3 (Pin bump and ROCm overlay retarget):** - Updated `GIT_TAG` in `src/lib/mlx-cpp/CMakeLists.txt` to `59d600b5e64c238427d0f8d897ab7c682ef4d3d2`. - Retargeted `src/lib/mlx-cpp/patches-rocm/` onto `59d600b5e64c238427d0f8d897ab7c682ef4d3d2` via `scripts/mlxcelverse/rocm_overlay.py retarget`, added the `fast::GatedDeltaUpdate` fallback and `NO_GPU_MULTI` stub in `mlx/backend/rocm/primitives.cpp` (documented in `LOCAL_FIXES.md` item 6), and regenerated `CORE_RESIDUAL.diff`. - **Phase 4 (`CONTRIBUTING.md`):** Added the `Reconciling CUDA, Metal, and CPU source overlays` subsection under `Bumping the MLX upstream pin` documenting the base-declaration format, the drift-check loop, and the `git merge-file` 3-way merge procedure. ## ROCm overlay - Fork commit (`UPSTREAM` `fork_commit`): `759159089fc9b17b3e0f1053abebff538f341fa5` (unchanged) - Retargeted MLX pin (`UPSTREAM` `retargeted_to_mlx_pin`): `81ba1c6a0e50a9268b931579c2d4f1053abebff5` -> `59d600b5e64c238427d0f8d897ab7c682ef4d3d2` - Core files merged cleanly by `rocm_overlay.py retarget`: `mlx/fast.cpp`, `mlx/fast.h`, `mlx/fast_primitives.h`, `mlx/ops.cpp` - Backend fix recorded in `LOCAL_FIXES.md` item 6: `fast::GatedDeltaUpdate::use_fallback` and `NO_GPU_MULTI(GatedDeltaUpdate)` in `mlx/backend/rocm/primitives.cpp` ## Verification - `bash scripts/ci/mlx_pinned_commit.sh` prints `59d600b5e64c238427d0f8d897ab7c682ef4d3d2` and `cargo test -p mlxcel-mlx-pin` passes (27/27 tests). - Overlay drift check across all 32 files under `src/lib/mlx-cpp/patches/` and `src/lib/mlx-cpp/patches-cuda/` against `59d600b5e64c238427d0f8d897ab7c682ef4d3d2` produces 0 `NO BASE` lines and 0 `DRIFT` lines. - `make verify-versions verify-kernel-dtype-keys verify-kernel-port-dispatch verify-llama-compat verify-rocm-overlay verify-python-tooling verify-fmt` passes. - Real quantized checkpoint validation (`Qwen3-0.6B-4bit`, greedy `--temp 0`, 32 tokens) on this AMD Ryzen AI MAX+ 395 (`gfx1151`) Linux host: - ROCm GPU (`bash scripts/rocm_gpu_guard.sh -- bash scripts/ci/rocm_smoke.sh`): 32 tokens generated at `251.24 tok/s`, matching the pre-bump greedy output token-for-token (`Okay, the user wants me to write a one-sentence about the ocean. Let me start by brainstorming some key elements of the ocean. The`). - CPU stream (`MLXCEL_DEVICE=cpu`): identical 32-token greedy output. - Cross-host note: the three Apple Silicon Metal turbo launcher tests (`sparse_v_kernel_threshold_zero_matches_graph`, `delegated_fused_kernel_matches_reference_over_200_steps`, `delegated_steel_envelope_matches_cold_only_fused_over_200_steps`) and the GB10 CUDA hardware suite were not run locally because this validation host is a Linux AMD `gfx1151` ROCm machine without Metal or NVIDIA CUDA hardware (CI runs the GB10 CUDA gate). Closes #1932
* Fix Log and Equal is_equivalent ignoring primitive state (ml-explore#4266) Co-authored-by: Cheng <git@zcbenz.com> * Stabilize reduced-precision InstanceNorm (ml-explore#4230) * chore: Normalize negative axes in sort and argsort (ml-explore#4332) * Clean up main thread compile cache before python interpreter shuts down (ml-explore#4373) * chore: Check malformed jaccl hostfile that miss rdma in pairs (ml-explore#4284) Co-authored-by: Cheng <git@zcbenz.com> * Round mxfp8 block scales up to avoid saturation (ml-explore#4353) Co-authored-by: Daniel Hiltgen <daniel.hiltgen@ollama.com> Co-authored-by: Cheng <git@zcbenz.com> * Add support for the __array_namespace_info__ (ml-explore#4334) * Stop a failed CUDA graph commit from poisoning the encoder (ml-explore#4356) Co-authored-by: Cheng <git@zcbenz.com> * Fix quantized kernels in JIT build (ml-explore#4372) Co-authored-by: Cheng <git@zcbenz.com> * Avoid zero work in stride-2 ConvTranspose3d (ml-explore#4343) * [CUDA] Ce fused kernel (ml-explore#3947) * Fix cpu exclusive scan for complex numbers (ml-explore#4272) Co-authored-by: Cheng <git@zcbenz.com> * Support Relocatable CUDA DLLs on Windows (ml-explore#4382) * Use cast_to for fused AsType in compiled Metal kernels (ml-explore#4351) Co-authored-by: katlun-lgtm <katlun@windyviews.com> Co-authored-by: Cheng <zcbenz@gmail.com> * python: Declare DLPackCompatible protocol members as methods (ml-explore#4384) * Fix quantizing sliced arrays (ml-explore#4381) * Fix einsum dropping a trailing empty subscript (ml-explore#4299) Co-authored-by: Cheng <git@zcbenz.com> * Add script to run python tests (ml-explore#4393) * Hold GIL in AttachedData destructor (ml-explore#4391) * Patch bump to 0.32.3 (ml-explore#4406) * Fix scan and sort for a zero-size axis (ml-explore#4340) * python: Add a correction parameter in std and var (ml-explore#4348) Signed-off-by: Pradyot Ranjan <99216956+pradyotRanjan@users.noreply.github.com> Co-authored-by: Pradyot Ranjan <99216956+pradyotRanjan@users.noreply.github.com> Co-authored-by: Cheng <git@zcbenz.com> * Fix tests/run.py stuck in macOS CI (ml-explore#4410) * chore: Update extension example (ml-explore#4412) * Fix sorted gather_qmm NAX row overflow above 32K (ml-explore#3922) Co-authored-by: Cheng <git@zcbenz.com> * Fix deadlock caused by mx.clear_streams() holding GIL (ml-explore#4413) * Fix integer pow zeroing a whole SIMD vector on a negative exponent (ml-explore#4354) * Make concurrency cap on load adaptative so that I/O scales with the machine (ml-explore#4408) * Fix pad with an axes subset and negative axes (ml-explore#4364) * Use sdpa_vector_2pass_1_gqa for GQA size 12 and 16 (ml-explore#4380) * python: Make axis of put_along_axis optional (ml-explore#4360) * python: Make axis default to -1 in take_along_axis (ml-explore#4368) * Propagate NaN in arg reductions (ml-explore#4291) Co-authored-by: Cheng <git@zcbenz.com> * Fix mx.from_fp8 E4M3FN NaN decode with branchless carry (ml-explore#4376) Co-authored-by: Saud <saudkamran@Sauds-MacBook-Pro.local> * Update rules for PR limitation bypass list (ml-explore#4415) * [CUDA] Use native events for GPU fence waits (ml-explore#4401) Co-authored-by: Cheng <git@zcbenz.com> * Validate GGUF tensor dimensions (ml-explore#4378) Co-authored-by: Cheng <git@zcbenz.com> * Fix crash on ellipsis indexing with too many trailing indices (ml-explore#4396) Co-authored-by: Anastasiia Filippova <a_filippova@apple.com> * [CUDA] Cholesky via cuSOLVER (ml-explore#4208) Co-authored-by: Cheng <git@zcbenz.com> * python: Fix __setitem__ with negative index preceded by None (ml-explore#4414) * Update nanobind to 3.0.1 (ml-explore#4417) * chore: Use normalize_axis_index in scan ops (ml-explore#4383) * python: Add explicit index and bytes support (ml-explore#4388) * docs: fix typo recived -> received (ml-explore#4424) Signed-off-by: Vaibhav Srivastava <vaibhavsri1712@gmail.com> * Fix Metal grid sizing for strided scans (ml-explore#4420) Co-authored-by: Cheng <git@zcbenz.com> * Fix SGD and Adafactor weight decay mutating caller arrays (ml-explore#4426) * Widen float16/bfloat16 to float32 in cpu reduce (ml-explore#4387) * Fix sdpa_vector_2pass_1_gqa kernel batch offset for K/V in gqa decode (ml-explore#4431) * Make head-dim-256 prefill with array mask run fused NAX kernel (ml-explore#4416) * python: Make unstack return tuple instead of list (ml-explore#4448) * python: Make iter() throw TypeError for 0-dim array (ml-explore#4425) Co-authored-by: Cheng <git@zcbenz.com> * Fix save and save_safetensors corrupting a lazily loaded source file (ml-explore#4434) * [Metal] Use hypot for complex ops (ml-explore#4345) Co-authored-by: Cheng <git@zcbenz.com> * Add get_array_buffer_size to query buffer size for arrays (ml-explore#4436) * Add M5 ultra tunings for non-quantized matmuls (ml-explore#4447) * [CUDA] Fix completion worker busy loop (ml-explore#4452) * Fix non transposed affine qmm dispatch logic (ml-explore#4392) * Fix sorted gather_qmm on ragged K (ml-explore#4009) Co-authored-by: Cheng <git@zcbenz.com> * Don't restore a thread-affine stream when StreamContext dies on another thread (ml-explore#4462) * Fix Pad vjp with an axes subset and negative axes (ml-explore#4441) Co-authored-by: Cheng <git@zcbenz.com> * Fix integer div-by-zero crash in CPU backend (sort, scan, binary ops) (ml-explore#4442) Co-authored-by: Cheng <git@zcbenz.com> * Use 64-bit file seeks on Windows (ml-explore#4456) Co-authored-by: Cheng <git@zcbenz.com> * Use precise::exp in Sigmoid so compiled and eager sigmoid agree (ml-explore#4461) Co-authored-by: Pierre Lamy <pierrelamy427@gmail.com> * Use NAX attention for long unmasked D72/D80 inputs (ml-explore#4455) * Add D512 support to Metal vector attention (ml-explore#4459) Co-authored-by: Cheng <git@zcbenz.com> * Add a scoped_env utility for tests (ml-explore#4474) * [Metal] global scale for qmm (ml-explore#4458) * Fix re-entering the same stream context manager (ml-explore#4478) * Leak global CommandEncoder on exit (ml-explore#4480) * Update CONTRIBUTING.md (ml-explore#4469) * python: Fix bytes() on non-contiguous arrays (ml-explore#4449) * Break the sibling cycle when an array is released by assignment (ml-explore#4453) Co-authored-by: Cheng <git@zcbenz.com> * win: Use DXGI for accurate WDDM VRAM memory budget (ml-explore#4457) Co-authored-by: Cheng <git@zcbenz.com> * Add type hints and docstrings for ALiBi layer (ml-explore#4472) * Fix fp quantized matmul corruption when the quantized dim is not a multiple of 32 (ml-explore#3912) Co-authored-by: Rohan Gautam <rohan1gautam@gmail.com> * Fix CUDA test synchronization flake (ml-explore#4490) * ci: Skip bypass list updates in forks (ml-explore#4495) * Load global scales in qmm_t kernels (ml-explore#4483) * python: Release the GIL when copying Metal DLPack inputs (ml-explore#4497) Signed-off-by: Ranran Haoran Zhang <ranranhaoranzhang@gmail.com> * Use matrix kernels for global-scale gather_qqmm (ml-explore#4481) * Adding metal kernels for the gated delta nets. (ml-explore#4020) Co-authored-by: Yanzhao Wang <19340816+wyanzhao@users.noreply.github.com> Co-authored-by: Cheng <git@zcbenz.com> * Fix jaccl ring all_gather: direction 1 slice was not mirrored (ml-explore#4443) Co-authored-by: Dave Hind <davehind42@gmail.com> Co-authored-by: Anastasiia Filippova <a_filippova@apple.com> * Skip Metal-only gated delta kernel tests on other backends(fix CI) (ml-explore#4522) Co-authored-by: Anastasiia Filippova <a_filippova@apple.com> * [CUDA] Add global scale support to gather_qmm (ml-explore#4507) Co-authored-by: Anastasiia Filippova <a_filippova@apple.com> * [BUG][Metal] Deadlock in fence (ml-explore#4552) Co-authored-by: Cheng <zcbenz@gmail.com> * Fix ordering of complex64 with NaNs (ml-explore#4519) * Improve metal memory usage for SDPA D256 (ml-explore#4505) Co-authored-by: Cheng <git@zcbenz.com> * Add support for matrix transpose (ml-explore#4402) Co-authored-by: Cheng <git@zcbenz.com> * Improve metal memory usage for SDPA D512 (ml-explore#4487) * Support shapeless compilation of scan operations (ml-explore#4510) * [Metal] Gather mm improvement (ml-explore#4567) * Fix perf regression for upsample `linear` mode align_corners=False (ml-explore#4500) * Add fused Metal kernels for fast.cross_entropy (ml-explore#4520) Co-authored-by: Anastasiia Filippova <a_filippova@apple.com> * python: Add missing defaults to tri, tril, triu, gather_mm signatures (ml-explore#4492) * Add Metal SDPA support for D96/V64 (ml-explore#4499) Co-authored-by: Cheng <git@zcbenz.com> * Raise IndexError for out of bounds axes (ml-explore#4484) * Fix floor_divide for integers (ml-explore#4515) Co-authored-by: Cheng <git@zcbenz.com> * Fix cpu binary ops on large data (ml-explore#4517) Co-authored-by: Cheng <git@zcbenz.com> * Do not throw from CUDA destructors and avoid implicit default streams in eval/compile (ml-explore#4514) Co-authored-by: Cheng <git@zcbenz.com> * [Metal] gather_qmm improvement (ml-explore#4572) * Fix floating-point constant precision in compiled kernels (ml-explore#4511) * chore: Fix assertion in ReduceScatter::eval_cpu (ml-explore#4557) * Leak thread local streams on exit in main thread (ml-explore#4576) * Adding VJP for the scaled dot product attention. (ml-explore#4563) * python: Add __pos__ support to array (ml-explore#4488) Co-authored-by: Cheng <zcbenz@gmail.com> * Add Python 3.15 and 3.15t support (ml-explore#4555) Co-authored-by: Cheng <git@zcbenz.com> * Add fixed size `unique` operation (ml-explore#4501) Co-authored-by: Cheng <zcbenz@gmail.com> * Fix normalizing single-element slice with negative stride (ml-explore#4446) Co-authored-by: Cheng <git@zcbenz.com> * Patch bump to 0.32.4 (ml-explore#4582) * Adding VJP for gated delta net. (ml-explore#4565) * Fix `col_reduce_longcolumn` for negative-stride views (ml-explore#4529) Co-authored-by: Cheng <zcbenz@gmail.com> * Clear CUDA errors once reported and raise eval errors in buffer protocol (ml-explore#4523) Co-authored-by: Cheng <zcbenz@gmail.com> * Improve pre-M5 metal memory usage for SDPA D512 (ml-explore#4518) * Fix split data size span (ml-explore#4540) * Re-introduce Array API Tests suite (ml-explore#4433) Signed-off-by: Pradyot Ranjan <99216956+pradyotRanjan@users.noreply.github.com> Co-authored-by: Pradyot Ranjan <99216956+pradyotRanjan@users.noreply.github.com> Co-authored-by: Cheng <zcbenz@gmail.com> * Fix view inside of a vmap (ml-explore#4586) * Fix cumsum of a non-contiguous bool (ml-explore#4590) * docs: contraction dtype requirements in matmul/einsum (ml-explore#4550) Co-authored-by: Cheng <zcbenz@gmail.com> * Fixing cooperative tensor build issues on Mac OS 27 (ml-explore#4594) * Handle EOF and EINTR in ParallelFileReader::read (ml-explore#4532) Co-authored-by: Cheng <git@zcbenz.com> * Clamp slice bounds when a slice is combined with an array index (ml-explore#4513) Co-authored-by: Cheng <zcbenz@gmail.com> * Fix grouped Conv1d/Conv2d init to use in_channels // groups (ml-explore#4544) Co-authored-by: Cheng <zcbenz@gmail.com> * Fix mx.view misaligned offset (ml-explore#4585) Co-authored-by: Cheng <git@zcbenz.com> * [CUDA] Add fp_gather_qmv to optimize gather_qmm (ml-explore#4508) * Fix eval cleanup deadlock by committing GPU streams before waiting (ml-explore#4556) * Fix shared_buffer_reshape contiguity flags (ml-explore#4589) * Fix view last axis stride (ml-explore#4587) Co-authored-by: Cheng <zcbenz@gmail.com> * Fix empty unsigned sum/prod failing on Metal (ml-explore#4546) Co-authored-by: Cheng <git@zcbenz.com> * [CUDA] Add gather_qmm_rhs_sm80 (ml-explore#4554) Co-authored-by: Cheng <git@zcbenz.com> * python: Fix bare ellipsis __setitem__ by squeezing leading singleton dimensions (ml-explore#4561) * Fix row/col contiguous flags in view (ml-explore#4588) * Do not cache failed distributed backend initialization (ml-explore#4558) * Fix reduction over large arrays (ml-explore#4549) Co-authored-by: Cheng <git@zcbenz.com> * docs: Correct the export callback type in the guide (ml-explore#4578) * Fix MultiOptimizer crash when an optimizer's group is empty (ml-explore#4570) * Fix fft size one (ml-explore#4591) * Sanitize '-' in make_template_hash for metal kernel (ml-explore#4584) Signed-off-by: GokayAI <60583610+gokay-ai@users.noreply.github.com> Co-authored-by: Cheng <git@zcbenz.com> * Fix as_strided contiguity (ml-explore#4575) * Reject non-aligned pointers when converting to mx.array (ml-explore#4592) * Fix conjugation in complex VJPs for inverse trig ops, cos, sqrt, and rsqrt (ml-explore#4606) * Fix complex to bool (ml-explore#4593) * [jaccl] second cable for the mesh (ml-explore#4607) * Fix vmap of scatter placing the batch axis wrongly in the updates (ml-explore#4322) Co-authored-by: Cheng <git@zcbenz.com> * Unroll TK=4 keys per iter in sdpa_vector_2pass_1 (ml-explore#4596) Co-authored-by: Cheng <git@zcbenz.com> * Fix CPU float16 SIMD comparison masks (ml-explore#4602) Co-authored-by: Cheng <git@zcbenz.com> * Add jvp support for logcumsumexp (ml-explore#4611) * [CUDA] Export CommandEncoder and Device (ml-explore#4619) * Use the fast qmv kernel for outputs not divisible by 8 (ml-explore#4516) Co-authored-by: Cheng <git@zcbenz.com> * Fix ResidencySets reading autoreleased error object (ml-explore#4623) * Use the small-channel implicit gemm loader for C_per_group = 8 (ml-explore#4617) * [CUDA] Reduce floor for fp_gather_qmv (ml-explore#4620) * Fix default.metallib lookup under swift test with Swift Build (ml-explore#4527) * Fix metal validation error when passing empty vector (ml-explore#4622) Co-authored-by: Cheng <git@zcbenz.com> * Replace Black formatter with ruff (ml-explore#4600) * Tune M1 quantized matmul dispatch for six-row large shapes (ml-explore#4629) Co-authored-by: deadbydawn <gothravenllm@gmail.com> Co-authored-by: Cheng <git@zcbenz.com> * [metal] Add env var based residency refresh interval (ml-explore#4633) * Promote integer inputs in logcumsumexp (ml-explore#4625) Co-authored-by: Cheng <git@zcbenz.com> * Fix slicing's non contiguous start (ml-explore#4599) * Fix float16 overflow in clip_grad_norm (ml-explore#4626) * Collapse the non sorted axes to pick the contiguous sort kernel (ml-explore#4366) Co-authored-by: Cheng <git@zcbenz.com> * Fix 32 bit dispatch grid overflow in Metal reductions (ml-explore#4437) Co-authored-by: Cheng <git@zcbenz.com> * Fix gpu boolean masking in 0 dim array (ml-explore#4635) Co-authored-by: Cheng <git@zcbenz.com> * Fix vmap of an inverse real FFT to an odd length (ml-explore#4637) * Fix cross-process race in CPU compile cache (ml-explore#4638) Co-authored-by: Cheng <git@zcbenz.com> * Update acknowledgments for past pinv contribution (ml-explore#4645) * Fix foreign buffer ownership (ml-explore#4634) * Add MLX_CONV_WINOGRAD switch to disable winograd for conv2d (ml-explore#4639) * Keep split-K quantized matmul partials in float32 (ml-explore#4641) * Thin N GEMM kernels (ml-explore#4640) Co-authored-by: Cheng <git@zcbenz.com> * Add 1-bit affine quantization support (Metal) (ml-explore#3161) Co-authored-by: Angelos Katharopoulos <a_katharopoulos@apple.com> Co-authored-by: bri-prism <288398250+bri-prism@users.noreply.github.com> * Update Metal Complex division in numpy style to avoid overflow condition (ml-explore#4644) * Fix pickle bfloat16 strides (ml-explore#4649) * Fix metal negative stride index in masked_scatter/scatter/SDPA mask/segmented_gemm (ml-explore#4650) * Fix missing break binary bool (ml-explore#4651) * Fix cpu binary_op_dispatch_dims on large data (ml-explore#4648) * Add FloorDivide primitive for integer floor division (ml-explore#4642) Co-authored-by: Cheng <git@zcbenz.com> * [Metal] Fix SDPA max threadgroup for M1/M2 generation (ml-explore#4643) Co-authored-by: Cheng <git@zcbenz.com> * Fix incorrect output mask computation in `mx.block_masked_mm` (ml-explore#4652) * Add thin M GEMM dispatch (ml-explore#4654) * Fix Python dtpye -> mlx conversion narrowing (ml-explore#4656) * Fix input/output scalar merging in compile (ml-explore#4658) Co-authored-by: Cheng <git@zcbenz.com> * Validate .npy shape in mlx/io/load.cpp (ml-explore#4657) Co-authored-by: Cheng <git@zcbenz.com> * Replace "pull_request_target" event with "pull_request" in update_bypass_list workflow (ml-explore#4667) * [CUDA] Only count consecutive misses in the LRUCache thrashing check (ml-explore#4661) Co-authored-by: Cheng <git@zcbenz.com> * Fix conv input gradient shape with uneven or large padding (ml-explore#4662) * Fix missing stream in layer_norm and sdpa fallbacks (ml-explore#4663) Co-authored-by: Cheng <git@zcbenz.com> * Fix direct CPU conv writing past output (ml-explore#4668) Co-authored-by: Cheng <git@zcbenz.com> * Add stft and istft to mlx.core.fft (ml-explore#3639) * Fix lost rank output in the distributed launcher (ml-explore#4615) Co-authored-by: Cheng <git@zcbenz.com> * Fix split-K test compatibility on M5, M6 (ml-explore#4677) * Fix signedness warning in quantized Metal tail bounds check (ml-explore#4675) * Fix batched row count mismatch in metal qvm_split_k (ml-explore#4673) Co-authored-by: Cheng <git@zcbenz.com> * Fix metal conv bounds checks and integer overflows (ml-explore#4672) Co-authored-by: Cheng <git@zcbenz.com> * python: Make arctan2/radians/degrees support scalar inputs (ml-explore#4679) Co-authored-by: Cheng <git@zcbenz.com> * Fix missing add_temporary in pad_gpu/fill_gpu (ml-explore#4660) Co-authored-by: “Ronald <“Ronaldmannak@me.com”> * Enable non-transposed NAX gather_qmm_rhs (ml-explore#4684) * Add Metal extension build support (ml-explore#4004) * Adding minmax method to MLX to follow Numpy/Jax API (ml-explore#4674) Co-authored-by: Cheng <git@zcbenz.com> * Fix missing stream in gather matmul indices cast (ml-explore#4687) * [BUG] Metal fence deadlock (ml-explore#4689) * Patch bump to 0.32.5 (ml-explore#4707) * docs: Publishing a New Release (ml-explore#4691) * ci: Run update_bypass_list on push instead of pull_request (ml-explore#4671) * [CUDA] Run the vector SDPA kernels at head_dim 256 (ml-explore#4698) Co-authored-by: yanghaojin <haojin.yang@greenbit.ai> * [CUDA] Fix grouped gemm with large data (ml-explore#4696) Co-authored-by: yanghaojin <haojin.yang@greenbit.ai> Co-authored-by: Cheng <git@zcbenz.com> * Fix conv2d OOB reads for partial output-channel tiles (ml-explore#4653) * docs: MetalLibrary in custom extension guide (ml-explore#4701) * [Metal] Fix SDPA f16/f32 d512 max threadgroup for M1/M2 generation (ml-explore#4708) * Add float16 head_dim 256 TK=1 SDPA vector kernel for M1/M2 (ml-explore#4719) Co-authored-by: Cheng <git@zcbenz.com> * Allow configuring the CPU JIT cache directory (ml-explore#4714) Co-authored-by: Cheng <git@zcbenz.com> * Make `broadcast_arrays` return a tuple (ml-explore#4713) Co-authored-by: Cheng <git@zcbenz.com> * Fix unaligned scalar reads and oversized contiguous buffer views (ml-explore#4721) Co-authored-by: Cheng <git@zcbenz.com> * [Metal] Add staged affine_qmm_t_nax kernels with split-k (ml-explore#4682) Signed-off-by: Yanzhao Wang <yzwang1996@gmail.com> Co-authored-by: Cheng <git@zcbenz.com> * [Metal] Add fast route for dequantizing fp4 to half in fp_qmm_t (ml-explore#4709) Co-authored-by: Cheng <git@zcbenz.com> * [Metal] Add qmv_mma/qmv_nax kernels for 6-16 rows with large N/K (ml-explore#4712) Co-authored-by: Cheng <git@zcbenz.com> * docs(forkdiff): move the base to upstream c6abf0d after the merge Set base.hash to the new merge-base. Remove the v0.32.2 sync section and the gather_mm row-tile backport section, because upstream now has those changes. Keep the revert of the ml-explore#4352 simdgroup skip in its own section. Reduce the floor_divide section to the tests the fork keeps. Add the re-deleted bypass-list workflow and scripts to Fork tooling. Keep the revert of the ml-explore#4352 simdgroup skip in fp_gather_qmm_rhs_nax and affine_gather_qmm_rhs_nax. The merge had taken upstream's sg_active guard there. State the new upstream sync policy in FORKDIFF.md: merge upstream when David directs it, and move base.hash at each merge. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * fix: repair the merge in the quantized kernels and the lint The merge put upstream's new "typename U = T" template parameter on the fork's qmm_t_expert_impl instead of qmm_t_impl. Move it back. fp_gather_qmv_fast_tail now calls adjust_matrix_offsets with the global scale, as upstream's fp_gather_qmv_fast does. Remove the extra closing brace that the previous commit put in affine_gather_qmm_rhs_nax. The kernel body is again the fork's body. Import numpy in test_fast_sdpa.py again. Upstream removed the import, and the fork's FP32-partials tests use it. Wrap the long lines in check_forkdiff.py for the ruff E501 rule that upstream turned on. Remove the PR #15 bias-sum sentence from fork.yaml: upstream now promotes the bias-sum operands itself (ml-explore#4516). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * test: follow the upstream allocator and gather_qmm APIs in fork tests Upstream allocator::malloc now returns an owning allocator::Data, and allocator::free is gone. The fork's snapshot and footprint tests need explicit release points, so they use allocator::allocator().malloc and free, as upstream's own allocator tests do. Upstream gather_qmm takes a global_scale argument before sorted_indices. Pass std::nullopt in the GPT-OSS MXFP4 tests. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * fix(cpu): floor_divide of min by -1 must not trap Upstream's FloorDivide divides the minimum signed value by -1 on the CPU. The quotient overflows, and on x86 the division traps (SIGFPE). The fork test "test floor_divide integer edge cases on the CPU" crashed in the Fedora x86_64 job. Divide by 1 there instead, as the PR #32 functors did. The result wraps to min. Describe the guard in the floor_divide section of fork.yaml. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * Compute the JACCL ring wire offsets in 64 bits CodeQL (cpp/integer-multiplication-cast-to-long) flagged the ring's wire offset expressions on PR 38: `lr * n_wires` is an `int` product that is widened to 64 bits only after the multiplication. The same pattern is in three ring_impl.h loops and in the two `size_ * 2 * n_conns_` thresholds in ring.cpp. Widen the first operand to int64_t so every product is computed in 64 bits. The operands are small, so results do not change. fork.yaml describes the new fork change. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * Compute the vector SDPA query layout check in 64 bits CodeQL (cpp/integer-multiplication-cast-to-long, alert 29) fails the PR 38 check on the vector SDPA query layout test: `shape[3] * shape[bidx]` is an `int` product compared with a 64-bit stride. Widen `shape[3]` to int64_t. The shapes are small, so the result does not change. The fork.yaml section for the JACCL ring fix now covers this file too. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * Fix `mx.take`'s negative indices (ml-explore#4729) * Fix normal loc scale shape (ml-explore#4730) * Fix zero-copy numpy arrays dtype mismatch (ml-explore#4731) * ci: skip the 4 GiB reduce test where the runner cannot hold it Upstream's test_init_reduce_large_output (ml-explore#4437) makes a 4 GiB boolean output. The fork CI runs on a GitHub-hosted macOS runner with a 3.5 GiB Metal buffer limit and 7 GB of memory. On the GPU the allocation fails; on the CPU it competes with the other test files for memory. Skip the test when the default device cannot hold the output. Machines with more memory still run it. When a CPU test file fails, print the end of its log at once. The job of the previous run was stopped before the full logs printed, so the cause of the test_eval failure on the CPU is not known. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * Widen affine qmv_wide row tiles to 8 (ml-explore#4728) * Fix/padding jvp vjp (ml-explore#4736) * fix: qqmm refuses a global scale on the CPU; adapt GPU-path tests The CPU QQMatmul kernel has no global scale input and ignores the scales that qqmm passes to it. qqmm with global scales on a CPU stream returned a wrong result with no error: in the fork CPU test pass, test_fp_qmv_fast_unaligned_n_global_scale differed by about 1000. qqmm now raises "Global scale is only supported on the GPU" on a stream that is not a GPU stream, as gather_qmm already does. Upstream tests that check only GPU kernels failed in the CPU pass: test_sdpa_vector_2pass_threadgroup_limit (force_fused=True), test_fp_qmv_fast_unaligned_n_global_scale, test_qqmm_global_scale_matrix_kernels and test_gather_qmm_global_scale. When the default device is the CPU, they now check that the documented error is raised. The GPU pass runs them as before. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * ci: add a temporary diagnostic for the CPU NaN; reword the 4 GiB skip Add one temporary step after the build in metal-build-and-test. It runs the failing case of test_fp_gather_qmm_non_multiple_of_32 (K=512, N=1040, transpose=False, sorted, float32) on the CPU with a fresh allocator, after a warm-up that fills the allocator cache with NaN buffers, and with the buffer cache off, and on the GPU. It prints NaN counts per expert for w, w_hat, y_q and y_hat. The step has continue-on-error, so the full test suite still runs. The next push removes it. State in fork.yaml when test_init_reduce_large_output skips, why it cannot be smaller (ml-explore#4437), and that David ruled on 2026-10-23 to keep the skip. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * ci: move the CPU NaN diagnostic to its own short job Remove the temporary diagnostic step from metal-build-and-test. Its reimplementation found no NaN in any mode. Add a temporary job, diag-cpu-nan, that builds the wheel and runs only the diagnostic, so its log is ready after the build. It runs the real test_fp_gather_qmm_non_multiple_of_32 alone, then after the tests that come before it, and an instrumented copy of the test body (only the failing case, and all cases in the test order, each evaluated as one graph). The findings also go to ::notice annotations. The next push removes the job. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * ci: install numpy for the temporary NaN diagnostic job The tests import numpy through mlx_tests. The diagnostic job installed only the wheel, so every run stopped at the import. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * Fix data race in distributed ring on pending send/recv lists (ml-explore#4737) Co-authored-by: Cheng <git@zcbenz.com> * Use data_size for gather/scatter large-offset checks (ml-explore#4724) Co-authored-by: Cheng <git@zcbenz.com> * fix(cpu): clear C before the BLAS matmul when beta is 0 test_fp_gather_qmm_non_multiple_of_32 (K=512, N=1040, sorted, float32) gave NaN on the CPU. The diagnostic job found the NaN only in the reference y_hat = x @ w_sel, a batched matmul with M = 1. Its inputs had no NaN. The test failed alone too, but only after the earlier subtests of the same test: y_hat is 266240 bytes, the same size as the packed fp4 weights of the previous subtest, so the output reused that buffer from the allocator cache. Accelerate read C even with beta = 0 and computed NaN * 0. Fill C with zeros before cblas_sgemm, cblas_dgemm and cblas_cgemm when beta is 0. MLX's own SIMD gemm already skips C when beta is 0. Remove the temporary diagnostic job. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * fix(cpu): clear C at every BLAS gemm call with beta 0 The previous commit cleared C only in the gemms/cblas.cpp wrappers. Four more call sites call cblas_sgemm directly with beta 0: BlockMaskedMM and GatherMM (the CPU MoE gathered matmul) in masked_mm.cpp, and the two explicit-gemm convolutions in conv.cpp. They can give the same NaN * 0 from a reused output buffer. Add one helper, clear_gemm_output in gemm.h, that fills the exact C region of a call (rows by columns, with its ldc) with zeros. Use it at all five call sites, the wrappers included. No other BLAS gemm or gemv call exists in the CPU backend. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * ci: build once and run the Metal and CPU test passes in parallel Split the fork-tests job into parallel jobs. build-wheel and build-cpp build in parallel. build-cpp also runs the C++ tests on the GPU and on the CPU. python-gpu (3 shards) and python-cpu (4 shards) run the Python tests from the uploaded wheel. metal-build-and-test merges the coverage profiles of all jobs, writes the same coverage report, and fails when another job failed. python_shards.py checks the shard map: each python/tests/test*.py file must be in each pass exactly one time. Split files must partition their test IDs. The last job fails when a test ran in two shards. fork-build-cache restores and saves ccache on pull requests and on main, and keeps a cache of the compiled Metal kernels through an xcrun wrapper keyed on the preprocessed source, the arguments and the toolchain. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * [CUDA] Fix gather_qmm for experts past 2^32 weight elements (ml-explore#4738) Co-authored-by: Cheng <git@zcbenz.com> * ci: split the long CPU test files and skip cache saves with no misses Split test_fast_sdpa and test_quantized into named tests and rest entries in the CPU shards, so that no CPU shard has an entry longer than test_gated_delta_sequential. Put the entries in longest-first order. Save the ccache files and the Metal kernel cache only when the build had misses. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * ci: key the build caches on compiler versions, not the runner image The full --version text of clang and metal names the install folder and the OS, and it changed between three macos-26 runner images, so every job started from an empty cache. Key ccache on the clang major.minor version and fall back to a cache of another clang version; ccache keeps compiler_check=content. Key the Metal kernel cache on the metal compiler version and the kernels hash, and fall back only to a cache of the same metal compiler. Each Metal entry key holds the first line of the metal --version text and the SDK build. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * Add array::is_uniquely_referenced for the fork C API Upstream 5dbe999 (ml-explore#4634) makes is_donatable() also require that MLX owns the buffer, so the allocator never reuses a wrapped external buffer. The fork's mlx-c buffer-info API reported `unique` from is_donatable(), and mlx-swift-lm uses that as a "no other reference" proof for paged KV segments, checkpoint memory accounting and MTP state, including wrapped host pages, which now always read as not unique. Add is_uniquely_referenced(): the same reference-count test without the ownership test. is_donatable() and donation behavior do not change. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * fix: async_eval must not throw a stream task error test_eval.test_async_eval_error_in_synchronize (upstream ml-explore#4338) failed on some runs of the fork's CPU test pass: mx.async_eval raised the LU factorization error that the test expects only from mx.synchronize. Cause: StreamThread::enqueue throws the stream's stored error when it runs on the main thread (ml-explore#3742). async_eval enqueues several tasks. When the stream thread finishes the failing task before async_eval enqueues the next one, that enqueue throws. The order depends on thread timing, so the failure is intermittent. The fork does not change this path. async_eval now holds a scheduler::DeferStreamErrors guard. While it lives, enqueue on that thread does not throw the stored error, and synchronize throws it later. eval does not change. Add a C++ test that stores the error before async_eval, so it fails on every run without the fix. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * Fix GPU pad for outputs larger than 2^31 elements (ml-explore#4744) * ci: format python_shards.py with the pinned ruff Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * ci: run the Python tests in 2 GPU and 3 CPU shards The hosted macOS runners are a shared pool. Pack the GPU pass into 2 shards and the CPU pass into 3 shards, balanced by the durations that run 39902948153 measured. A run now uses 8 macOS jobs, and at most 6 run at the same time (before: 10 jobs, 8 at the same time). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * ci: give the NAX kernels Metal cache hits and lighten CPU shard 0 The preprocessed source of the NAX kernels names the Metal toolchain mount folder, which has a random suffix on each runner. The 10 NAX kernels missed the Metal cache in every warm build (41 of 51 hits). Remove the suffix from the text that the key hashes, and keep the toolchain version. CPU shard 0 was the slowest shard in the cold and warm runs of b72420d (651 s and 470 s of tests). Move test_quantized:rest to CPU shard 1 (221 s). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * docs(forkdiff): move the base to upstream afc03c5 Set base.hash to the new merge-base after the merge of the 9 upstream commits. No fork section changes: the merge has no conflict, and each merged file gets exactly upstream's change. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> --------- Signed-off-by: Pradyot Ranjan <99216956+pradyotRanjan@users.noreply.github.com> Signed-off-by: Vaibhav Srivastava <vaibhavsri1712@gmail.com> Signed-off-by: Ranran Haoran Zhang <ranranhaoranzhang@gmail.com> Signed-off-by: GokayAI <60583610+gokay-ai@users.noreply.github.com> Signed-off-by: Yanzhao Wang <yzwang1996@gmail.com> Co-authored-by: rohith <kapellirohith@gmail.com> Co-authored-by: Cheng <git@zcbenz.com> Co-authored-by: Vladimir Iglovikov <ternaus@users.noreply.github.com> Co-authored-by: Brian C. <94733710+deBrian07@users.noreply.github.com> Co-authored-by: Cheng <zcbenz@gmail.com> Co-authored-by: Erwin Zhang <59893706+erwinzhang7@users.noreply.github.com> Co-authored-by: Daniel Hiltgen <dhiltgen@users.noreply.github.com> Co-authored-by: Daniel Hiltgen <daniel.hiltgen@ollama.com> Co-authored-by: Aaishwarya Mishra <aaishwarymishra@gmail.com> Co-authored-by: YH Yan <strayberry0w0@gmail.com> Co-authored-by: Dwijen Patel <dwijen@gmail.com> Co-authored-by: Anastasiia Filippova <a_filippova@apple.com> Co-authored-by: Ayaan Gazali <ayaangazali.work@gmail.com> Co-authored-by: katlun-lgtm <katlun@gmail.com> Co-authored-by: katlun-lgtm <katlun@windyviews.com> Co-authored-by: Adityaj0 <93090622+Adityaj0@users.noreply.github.com> Co-authored-by: Eyüp Can Akman <eyupcanakman@gmail.com> Co-authored-by: Pradyot Ranjan <99216956+prady0t@users.noreply.github.com> Co-authored-by: Pradyot Ranjan <99216956+pradyotRanjan@users.noreply.github.com> Co-authored-by: Jagrit Digani <digani@apple.com> Co-authored-by: Philip John Basile <PBasile@Basilecom.com> Co-authored-by: Alessio Pollero <alessio.pollero@ventureone.ae> Co-authored-by: Duhyeon, Kim <49020301+dudududukim@users.noreply.github.com> Co-authored-by: Atirna <noctiselara0@gmail.com> Co-authored-by: Saud Kamran <32072820+saud5150@users.noreply.github.com> Co-authored-by: Saud <saudkamran@Sauds-MacBook-Pro.local> Co-authored-by: YH Yan <yh.yukiyan@gmail.com> Co-authored-by: Roshan Sharma <rs127@fordham.edu> Co-authored-by: sashko-zakharchuk <oleksandr.zakharchuk@gmail.com> Co-authored-by: XXXXRT666 <157766680+XXXXRT666@users.noreply.github.com> Co-authored-by: Vaibhav Srivastava <vaibhavsri1712@gmail.com> Co-authored-by: Robin Gieseke <38696388+TheDarkchip@users.noreply.github.com> Co-authored-by: 2sumtech <2sumtech@gmail.com> Co-authored-by: Vedant Singhal <88609765+Ved235@users.noreply.github.com> Co-authored-by: JasonHonKL <148705846+JasonHonKL@users.noreply.github.com> Co-authored-by: Simeet Nayan <78461155+simeetnayan81@users.noreply.github.com> Co-authored-by: Cdimoy <50119145+Cdimoy@users.noreply.github.com> Co-authored-by: Alessio Pollero <alessio.pollero@gmail.com> Co-authored-by: Rohan Gautam <rohan1gautam@gmail.com> Co-authored-by: michalk8 <46717574+michalk8@users.noreply.github.com> Co-authored-by: Nilesh Patil <128893479+nileshpatil6@users.noreply.github.com> Co-authored-by: Marcos Ashton <59975445+MarcosAsh@users.noreply.github.com> Co-authored-by: pierre427 <pierre@userid.org> Co-authored-by: Pierre Lamy <pierrelamy427@gmail.com> Co-authored-by: Yanzhao Wang <19340816+wyanzhao@users.noreply.github.com> Co-authored-by: David Koski <46639364+davidkoski@users.noreply.github.com> Co-authored-by: Isaac Hernández <isadohergar@gmail.com> Co-authored-by: Tudorica Constantin-Alexandru <tudalex@gmail.com> Co-authored-by: Ritaban Mitra <39273308+Ritabanm@users.noreply.github.com> Co-authored-by: Ranran <ranranhaoranzhang@gmail.com> Co-authored-by: Tommaso Pegolotti <31819389+tpegolotti@users.noreply.github.com> Co-authored-by: Drifter4242 <davehind@yahoo.co.uk> Co-authored-by: Dave Hind <davehind42@gmail.com> Co-authored-by: Valentin Roussellet <vroussellet@apple.com> Co-authored-by: Ken Wang <wangxiaotian2007@gmail.com> Co-authored-by: Spass <spass.ni42+github@gmail.com> Co-authored-by: zsun6 <115504839+zsun6@users.noreply.github.com> Co-authored-by: rishabh <rishabhsaiv@gmail.com> Co-authored-by: Yanzhao Wang <yzwang1996@gmail.com> Co-authored-by: devangpratap <115096812+devangpratap@users.noreply.github.com> Co-authored-by: Jasim Kareem <mj9034812@gmail.com> Co-authored-by: Ryan Choi <54298620+Ryan11c@users.noreply.github.com> Co-authored-by: Ronald Mannak <ronaldmannak@me.com> Co-authored-by: Mateusz Rajski <60450261+mateuuszzzzz@users.noreply.github.com> Co-authored-by: Ishaan Potle <ishaanpotle9@gmail.com> Co-authored-by: Moinuddin Shaik <monu.dev@proton.me> Co-authored-by: Moinuddin Shaik <moinuddinmoin1357@gmail.com> Co-authored-by: shaeantoine <42801108+shaeantoine@users.noreply.github.com> Co-authored-by: Nicoló Angileri <nicoloangileri635@gmail.com> Co-authored-by: Sam Small <sam.small@gmail.com> Co-authored-by: Ishti <45158028+ishtihoss@users.noreply.github.com> Co-authored-by: GokayAI <60583610+gokay-ai@users.noreply.github.com> Co-authored-by: Valeriy Sofin <112272443+sofinvalery@users.noreply.github.com> Co-authored-by: Yash Raj Pandey <55940078+devYRPauli@users.noreply.github.com> Co-authored-by: Kaiyu Shi <skyisno.1@gmail.com> Co-authored-by: crabel99 <crabel@mac.com> Co-authored-by: nav <navneet.dagdiya@gmail.com> Co-authored-by: LongYinan <lynweklm@gmail.com> Co-authored-by: Joseph Heck <j_heck@apple.com> Co-authored-by: deadbydawn101 <sanjosepcrepair@gmail.com> Co-authored-by: deadbydawn <gothravenllm@gmail.com> Co-authored-by: Vijay Kapse <vijayskofficial@gmail.com> Co-authored-by: Long Yixing <longyixing331@gmail.com> Co-authored-by: Aditya Dhulipala <adi@dhulipala.xyz> Co-authored-by: Ishaan Mehta <106708534+ishaanzee@users.noreply.github.com> Co-authored-by: Mattt <mattt@me.com> Co-authored-by: Irakli Salia <65120973+Isalia20@users.noreply.github.com> Co-authored-by: Pasha Khosravi <khosravipasha@users.noreply.github.com> Co-authored-by: Angelos Katharopoulos <a_katharopoulos@apple.com> Co-authored-by: bri-prism <288398250+bri-prism@users.noreply.github.com> Co-authored-by: CodeAndCanvas728 <cedar-elk-01@icloud.com> Co-authored-by: SounLabs <159518213+SounLabs@users.noreply.github.com> Co-authored-by: David Dalcu <david.dalcu@gmail.com> Co-authored-by: Akshay Hitesh Shah <109842662+AkshayShah03@users.noreply.github.com> Co-authored-by: Sash <sash@ela.city> Co-authored-by: tc3oliver <57586652+tc3oliver@users.noreply.github.com> Co-authored-by: Magnus Lundstedt <magnus.lundstedt@infidyne.com> Co-authored-by: jrp2014 <jrp2014@users.noreply.github.com> Co-authored-by: Yi <39162855+SleepAviator@users.noreply.github.com> Co-authored-by: “Ronald <“Ronaldmannak@me.com”> Co-authored-by: Ivan Petrukha <93926277+petrukha-ivan@users.noreply.github.com> Co-authored-by: Haojin Yang <yanghaojin@hotmail.com> Co-authored-by: yanghaojin <haojin.yang@greenbit.ai> Co-authored-by: Marcin Markiewicz <75364832+jmmarkiewicz@users.noreply.github.com> Co-authored-by: Corentin Kerisit <corentin.kerisit@gmail.com> Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com> Co-authored-by: Ambisphaeric <270742616+Ambisphaeric@users.noreply.github.com> Co-authored-by: Sankur Kundu <sankur.kundu.tw@gmail.com> Co-authored-by: begcdn <150344566+begcdn@users.noreply.github.com> Co-authored-by: Core <45273146+IamCooller@users.noreply.github.com> Co-authored-by: David Tivris <git@davidtivr.is>
Problem
nvfp4 is the only quantization mode whose group size (16) lets the quantized
dimension of a matmul legally be a multiple of 16 but not of 32.
mx.quantize(w, mode="nvfp4")acceptsK = 1040, and the fp quantized Metalkernels tile that dimension by 32 without bounding the 16-wide tail. Whenever it
happens, every matrix-sized
quantized_matmul/gather_qmmon the GPUsilently returns corrupted results. The CPU backend and the vector (decode)
kernels handle the same shapes correctly, so a model can decode perfectly and
corrupt during prefill, which is about as quiet as corruption gets.
Minimal reproducer (M3 Pro, macOS 25.5, main @ 39d9a8a):
stream=mx.cpu: 1.2e-3.M=1(vector kernel): 1.1e-3. AlignedK=1024: 3e-3.Only the GPU matrix path with
K % 32 == 16is wrong.The three defects
All in
mlx/backend/metal/kernels/fp_quantized.h, introduced with the Metalnvfp4 support in #2946.
fp_qmm_t_implran its K loop pastK_eff(for (int k = 0; k < K_eff; k += BK)with full-BK loads), so the final iteration read 16 columns of thenext output row's packed weights and scales (past the buffer for the last
row) into the accumulator. Metal shader validation on the unpatched kernel:
Invalid device load at offset 2129920 … "nvfp4_qmm_t_float16_t_gs_16_b_4_alN_true_batch_0"(the weight buffer is exactly 2129920 bytes). Also backs the tiled
gather_qmm, so the unsorted MoE path fails identically.fp_qmm_n_implstored full 32-column tiles for a 16-column tail. Thestore clipped rows only (
store_result_safe(y, N, short2(BN, num_els))), sowith
N % 32 == 16the last column tile stored 32 columns into rows only Nlong: the 16 spilled columns land on the next row's first 16 outputs (racing
the neighbouring threadgroup) and past the end of the output buffer for the
last row. Unpatched: 200 identical runs give 61 distinct results,
corruption sits exactly at columns 0 to 15, and shader validation flags
Invalid device store at offset 1064976 … "nvfp4_qmm_n_float16_t_gs_16_b_4_batch_0".QuantizedBlockLoader::load_safemasked the wrong axis (if (reduction_dim == 1 && bi >= src_tile_dim.x)):biis the thread's row,src_tile_dim.xthe column-direction bound. In the K-remainder step of thesorted-MoE kernel
fp_gather_qmm_rhs(tile_w = short2(k_remain=16, tgp_bn)), this zeroed the threads owning output columns 16 to 31 of each tileand silently dropped their last 16 K-elements. Positive ID: on the unpatched
kernel the wrong columns are exactly those ≡ 16..31 (mod 32), and their
values equal
ref - x[:, K-16:] @ w[e][:, K-16:].Tto within 0.05 (fp16rounding): the dot product minus exactly its tail. In the
fp_qmm_tcallsites the same compare is
bi >= 32, dead code, so their intended columnmasking never ran.
Fix
fp_qmm_t_impl: iterateK_eff / BKfull tiles, then one boundedremainder tile, mirroring
fp_qmm_n_impl's existing K handling. The split-Kkernel
fp_qmm_t_splitkpassesK_eff = k_partition_size, which Fix incorrect nvfp4 quantized_matmul through the split-K path #3854 keepsa multiple of 32, so
num_k == 0and split-K is unchanged.fp_qmm_n_impl: takenum_outs = min(BN, N - y_col), clip the w loads toit, and store with
store_result_safe(y, N, short2(num_outs, num_els)).Removes the race and the out-of-bounds store.
QuantizedBlockLoader::load_safe(fp_quantized.hand the affine twin inquantized.h): mask rows against.yand the thread's columns against.x.A thread owns
n_reads * pack_factorcontiguous columns and every legalpartial extent is a whole number of quantization groups, so a thread is never
straddled; the new
static_assert(group_size % (n_reads * pack_factor) == 0)in both loaders pins that invariant.
A quantized dim is always a multiple of
group_size, so it can only end in apartial tile when
group_sizedoes not divide the block. Only nvfp4 does;mxfp4 and mxfp8 have
group_size == 32 == BK == BN, and every affine groupsize is a multiple of the block. The remainder tile and the column bound are
therefore selected with
if constexprongroup_size % BK != 0,group_size % BN != 0andgroup_size % BCOLS != 0, so they are discarded atcompile time for those modes.
This also removes an out-of-bounds read that affine and mxfp hit today
The row half of the mask is not nvfp4-only. In
qmm_t_implthe w loader isinstantiated with
BROWS = BNandreduction_dim = 1, and the call isload_safe(short2(BK, num_outs)). The old compare wasbi >= src_tile_dim.x,that is
bi >= BK == 32, whilebiranges over[0, BROWS) == [0, 32), so itnever fired: whenever
N % BN != 0the last tile read weight and scale rowspast
N. The store already clipped those columns, so results were correct;this is a read past the end of the buffer, not corruption. It reproduces on
affine, which is the most used mode:
K = 32768makes the packed weight buffer an exact multiple of the page size,so the over-read crosses the allocation. Under Metal shader validation on
39d9a8a:
With this change the same run is clean, as are the two nvfp4 shapes below. The
same applies to mxfp4 and mxfp8 with
N % BN != 0.On the loader mask:
biis a row index andbja packed column index, andevery call site passes
src_tile_dimas (valid columns, valid rows), swappingwhich physical dim that is to match the layout.
fp_qmm_t_implpassesshort2(BK, num_outs)withBROWS = BN,fp_qmm_n_implpassesshort2(num_outs, BK)withBROWS = BK, andfp_gather_qmm_rhspicksshort2(k_remain, tgp_bn)orshort2(tgp_bn, k_remain)ontranspose. That isthe convention steel's
BlockLoader::load_safealready uses(
src_tile_dim - short2(bj, bi)), so the row index tests against.yin everycase and the mask does not depend on
reduction_dim.Validation
M3 Pro (applegpu_g15s), macOS 25.5, against main @ 39d9a8a. "ref" is fp64 on
the MLX-dequantized weights (bit-identical CPU vs GPU dequantize, asserted).
fp_qmm_t_splitk), K∈{2048,4096,8192}N % BN != 0The campaign's CPU cross-check uses a looser gate for fp16 and bf16, since the
CPU backend accumulates those differently over a large K; that is not code this
PR touches.
The added tests use M ≥ 33;
get_qmv_batch_limitcaps the vector/matrixthreshold at 32, so the tiled kernels are exercised on every Apple GPU family
(M1 to M4, including Max/Ultra). 17 of the 19 subtests fail on 39d9a8a and all
pass with the fix. The two that pass on both are controls: the vjp with
transpose=Falsehas an aligned quantized dim, and the sorted gather withtranspose=Falseruns the loader withreduction_dim == 0, where the old rowcompare was already correct.
To confirm the
if constexprguards cost the other modes nothing, I compiledfp_quantized.metalagainst the base headers and against this branch andcompared every emitted kernel body after normalizing metadata ids. Of the 306
gs_32kernels, 253 are byte-identical and none gained a singleinstruction; 53 shrank, for a net of −141 instructions, while the +12950
instructions added by the fix land entirely on the nvfp4 kernels. The 53 that
shrank are exactly the
reduction_dim == 1families (qmm_t,qmm_t_splitk,gather_qmm_t,gather_qmm_rhs_nt), where the row mask now short-circuits therows it used to read out of bounds.
Perf (aligned fast path, min-of-medians, ms, before -> after)
Nothing regresses. The mxfp4, mxfp8 and affine rows are the ones that matter
for the
if constexprguards above, and the unaligned gather rows are thekernels where the row mask now short-circuits work the base build was doing on
rows it read out of bounds.
Why fix the kernel rather than reject the dimension
group_size = 16advertises that quantized dimensions which are multiples of 16are supported; silently corrupting a legal, documented input is an API-contract
violation the kernel should honor rather than an input to reject at the op layer.
The inconsistency already shows:
mx.quantize(w, mode="nvfp4", stream=mx.cpu)throws a reshape error when
w.size % 32 != 0while the GPU path accepts thesame tensor and later corrupts. One backend rejects what the other silently
mishandles. Making the kernels correct for every legal group_size=16 dimension
resolves both.
Provenance
Introduced with the Metal nvfp4 kernels in #2946. #3854 narrowed only the
split-K route: for
K ≡ 16 (mod 32)no split_k makesK % (split_k * 32) == 0,so those shapes always take its
split_k = 1fallback straight into plainqmm()->fp_qmm_t_impl, the unbounded loop fixed here. The reproducer above(M=50, K=1040, B=1, transpose=True) is exactly that path, so the shapes #3854
meant to protect were still corrupted.
Related but distinct: #3856 reports affine
gather_qmmcorruption at large unalignedrow counts (
n > 32768,n % 64 != 0) on M5, a row-count (M) bug. This is aquantized-dimension (K/N) bug specific to nvfp4; #3856's own repro does not reproduce
against this branch, and this fix leaves its
K % 64 == 0shapes on the same NAXkernel, so the two do not overlap.
Tests
test_fp_qmm_non_multiple_of_32covers K=1040/528 (transpose=True) and N=1040(transpose=False, run twice and compared bit for bit to catch the store race),
each in fp32, fp16 and bf16, plus mxfp4 and mxfp8 at the same block-unaligned
shapes as a control that the
if constexprguards did not change them.test_fp_gather_qmm_non_multiple_of_32covers sorted and unsortedgather_qmmin the same three dtypes, with the tail on K (transpose=True) and on N
(transpose=False).
test_fp_qmm_non_multiple_of_32_vjpcovers the backward ofboth transposes. Inputs are scaled by
1 / sqrt(K)so the dot products areO(1); at magnitude 32 a single fp16 ulp is 0.03 and would swamp the tolerance.