Repository navigation
test(rocm): add the gfx1151 correctness matrix and its gate - #1826
Merged
Merged
Conversation
Metal reference traces for the cross-backend correctness matrix in #1809, measured at bec6474 so the ROCm traces can come from the same commit. Four affine 4-bit checkpoints (qwen3-0.6b, llama-3.1-8b-instruct, qwen3-30b-a3b, mixtral-8x7b-instruct) are traced over tests/fixtures/wikitext2_excerpt.txt at width 1, width 8 behind 512 tokens of context, and width 256. The two MoE checkpoints are traced with fused MoE on and off. The directory README gives the arguments and the comparison command, METADATA.txt records the host, OS, MLX pin and the metallib, binary and corpus hashes, and SHA256SUMS covers every trace. All 18 runs exited 0 with the expected row counts and no NaN. On this host fused MoE and gather_qmm agree at every position except one undecided position for qwen3-30b-a3b at width 1. Absolute local paths in the trace headers were rewritten to repository-relative ones, which the comparer does not read. Refs #1809
The ROCm backend had no evidence that it decodes the same text as Metal for the same checkpoint and input. This adds that evidence and the gate that produced the defects it did not find. Twelve model x width pairs (qwen3-0.6b, llama-3.1-8b-instruct, qwen3-30b-a3b, mixtral-8x7b-instruct at widths 1, 8 behind 512 tokens of context, and 256) traced on a Radeon 8060S at bec6474 against the Metal traces from the same commit. Every pair with decided positions has zero decided-position disagreements. Of the 114 top-1 disagreements across all pairs, 100 pick the reference's rank-2 token and every one falls where the reference's top-2 gap is under 2.0. The two MoE checkpoints ran with MLXCEL_FUSED_MOE=0 on both sides because the fused kernel has no ROCm port (#1803). docs/benchmark_results/rocm-correctness-gfx1151-2026-09-12.md records the method, why byte-identity is not the criterion, why no noise floor was measured, and the known gaps. make verify-test-rocm and make test-fast-rocm mirror the existing CUDA targets so the workspace suite can be run against a ROCm build. That gate, not the traces, found the two defects fixed in #1823 and #1824: both sit off the decode path the traces exercise. compare_logit_traces.py reported the undecided-only verdict when the reference had no decided position at all, which reads as a pass on no evidence. It now says inconclusive for that case, which is what the two 128-position width-1 runs hit. Refs #1809
Under `--features rocm` the test binaries of mlxcel-surgery and mlxcel-xla built and then failed to start with "libamdhip64.so.7: cannot open shared object file", so six targets in `make verify-test-rocm` never ran a test. mlxcel-core links the ROCm shared libraries and emits the rpath for them, but a dependency's `cargo:rustc-link-arg` does not reach the crate that links it, and ROCm installs do not always register their library directory with the dynamic loader. The root `mlxcel` crate worked only because #1802 gave it its own copy of the rpath block. The emission moves to `src/lib/mlxcel-core/build_support/rocm_rpath.rs`, shared by `#[path]` the way `mlx_pin.rs` already is, and the root build script, a new mlxcel-surgery build script and the mlxcel-xla build script all call it. Both crates gain a `rocm` feature forwarding to `mlxcel-core/rocm` so `--workspace --features rocm` reaches their build scripts; neither compiles any ROCm code of its own. With the fix mlxcel-surgery runs 143 tests and mlxcel-xla 280, all passing, and the ROCm gate goes from 12 failing targets to 5. The remaining 5 are the missing primitives in #1825 and the CUDA-only kernel routing in #1803. Refs #1802, #1809
inureyes
added
type:test
The 1.125 logit figure the results doc leans on had no reproduction path: `compare_logit_traces.py` printed only bucketed rates, so the number came from a one-off script. It now prints the largest reference top-two gap at which the two arms still disagree, which is the threshold-free form of the decided-position metric and states directly that any `--decided` above that value is all-zero. Pairs that agree everywhere say so instead of printing a number. Review notes from the Metal side, all in text: the Metal host's OS is macOS 27.0 (26A428), not just "macOS"; the checkpoint directory name is host-local and differs between the two trace sets, so the reproduction snippet says which name belongs to which host and that the Hugging Face revision is the thing that has to match; the Metal trace README records why the six `default` MoE traces are unused by this matrix and what they are kept for. The MoE routing explanation gains its same-backend control: swapping the fused MoE kernel for `gather_qmm` within Metal moves only `qwen3-30b-a3b` `w1` (1 of 128 top-1, 0 of 74 decided) and leaves the other five pairs identical, so the model that reacts to a kernel change is the one that reacts to a backend change. Refs #1809
The same-backend control paragraph said which model moves when the MoE kernel is swapped within Metal, but not by how much, so a reader had nothing to size the 1.125 cross-backend figure against. The one disagreement in that control sits at a reference gap of 0.062, measured with the new largest-gap line in compare_logit_traces.py on the Metal default and fused0 traces. Refs #1809
This was referenced Sep 12, 2026
inureyes
added a commit
that referenced
this pull request
Sep 29, 2026
) Removes the ROCm overlay's `get_launch_args` (both overloads in `mlx/backend/rocm/kernel_utils.hpp`), which capped the grid at 65535 blocks of 256 threads without saying the kernel had to be grid-stride. It had no caller, but upstream CUDA's helper of the same name does not cap x and is called throughout `mlx/backend/cuda`, so a kernel ported from there would have compiled here and silently left everything past 65535 * 256 elements unwritten. In its place the header declares `template <typename... Args> void get_launch_args(Args&&...) = delete;` with a comment giving the reason, so any call, in either upstream's or the fork's signature, fails to compile at the explanation rather than as an undeclared name that invites restoring the old helper. No kernel's launch geometry changes. The decision (option 1 of the three in the issue) and its reason are recorded as `LOCAL_FIXES.md` item 19: renaming keeps a helper that only looks like upstream's, and spilling the overflow into y would not make a one-index-per-thread kernel correct unless it also reads `blockIdx.y`. Item 19 is an upstreaming candidate for #1813. ## Found during review, not fixed here `SliceUpdate::eval_gpu`'s reduce-op path in `indexing.hip` clamps its grid at 65535 blocks and launches `slice_update_op_kernel`, which handles at most `NWORK` elements per thread and never reads `gridDim`. An update larger than 65535 x 256 x `NWORK` elements leaves its tail unapplied. This is the defect #1874 describes at a live call site; item 19 records it, and it needs its own issue since this PR changes no kernel geometry. ## Verification On gfx1151 (Radeon 8060S, ROCm): - `cargo build --release --features rocm`: passes, cold and again incrementally after the second commit (the build tree's copy of the header matches the overlay). - Compile-fail probe with `hipcc -fsyntax-only --offload-arch=gfx1151` on a TU calling `get_launch_args(arr, false)` and `get_launch_args(n, shape, strides, false, 4)`: both rejected as "call to deleted function" with the patched header; the same TU compiles with the header from `origin/main`. - `grep -rn get_launch_args src/lib/mlx-cpp/patches-rocm/`: only the deleted declaration and its comment remain. - `make verify-rocm-smoke` with Qwen3-0.6B-4bit: OK. - `make verify-versions verify-kernel-dtype-keys verify-kernel-port-dispatch verify-llama-compat verify-fmt`: pass. Not run here: the full `make verify-rocm` gate (left to the central run) and the #1826 correctness matrix. The removed function had no caller, so no kernel code changes. Not verified: Metal and CUDA (not available on this host). Neither build copies `patches-rocm/`, so neither path is touched. Closes #1874
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.
The experimental ROCm backend had no evidence that it decodes the same text as Metal. This adds that evidence, the gate that produced it, and a build fix the gate exposed.
Twelve model x width pairs (qwen3-0.6b, llama-3.1-8b-instruct, qwen3-30b-a3b, mixtral-8x7b-instruct at widths 1, 8 behind 512 tokens of context, and 256) traced on a Radeon 8060S (
gfx1151) atbec64748against Metal traces from the same commit, same corpus, same arguments and same checkpoint revisions. The two backends disagree on the top-1 token at 114 of 5120 positions, and the largest reference gap at which they ever disagree is 1.125 logits: wherever the reference preferred its own choice by more than that, ROCm made the same choice. 100 of the 114 disagreements take the reference's rank-2 token.docs/benchmark_results/rocm-correctness-gfx1151-2026-09-12.mdhas the full matrix, a threshold sweep at 0.5 / 1.0 / 2.0 so the--decidedline is inspectable rather than asserted, the rank distribution, why byte-identity is the right criterion where it holds but not across backends, and why no noise floor was measured.One caveat worth reading before the table: the Mixtral ROCm traces come from
bec64748plus only the f16gather_qmmdispatch change merged as #1822, which lives entirely insrc/lib/mlx-cpp/patches-rocm/and touches no Rust or bridge code. Without it that f16 checkpoint spends about 205 ms per MoEgather_qmmcall and the traces take hours. Every other ROCm trace and all of the Metal side are plainbec64748.make verify-test-rocmandmake test-fast-rocmmirror the CUDA targets. That gate, not the traces, found the defects fixed in #1823 and #1824, plus the one fixed here: mlxcel-surgery and mlxcel-xla test binaries never received the ROCm rpath, because a dependency'scargo:rustc-link-argdoes not reach the crate that links it, so they built and then failed to start with "libamdhip64.so.7: cannot open shared object file". The emission now lives inbuild_support/rocm_rpath.rsand all three crates call it. The gate goes from 12 failing targets to 5, and the remaining 5 are all #1825 (FFT, Hadamard, SearchSorted stubs) and #1803 (CUDA-only kernel routing).compare_logit_traces.pyreported the undecided-only verdict when the reference had no decided position at all, which reads as a pass on no evidence; it now saysinconclusive, which is what the two 128-position width-1 runs hit.Validated on
gfx1151:cargo fmt --all -- --check,cargo clippy --workspace --all-targets --features rocm -- -D warnings,scripts/ci/check_cross_repo_refs.py,scripts/insert_apache_header.py --check, andmake verify-test-rocmwith the results recorded in the doc. The Metal and CUDA paths compile no file this PR touches except the Makefile and the comparison script.Refs #1809, #1802