Repository navigation
fix(rocm): refuse wave32-only kernel ports on wave64 devices - #2179
Merged
Merged
Conversation
The shuffle-based HIP ports have only run on 32-lane wavefronts, and their `#error` guards on `__AMDGCN_WAVEFRONT_SIZE` are inert because AMD clang 23 defines neither spelling of the macro (rechecked with `hipcc -E -dM` for gfx1151, gfx942 and gfx90a). On a CDNA device every one of them would run unvalidated. `port_for` now holds a ROCm port to devices whose hardware wavefront is 32 lanes unless its table sets `KernelPorts::rocm_any_wave_size`, so `has_kernel_port`, every `*_available()` predicate and `select_kernel_port` answer "no port" there and callers take their graph fallbacks. The width comes from a new overlay call, `mlx::core::rocm::device_warp_size()` (LOCAL_FIXES item 31), which ignores `MLX_ROCM_FORCE_WARP_SIZE`. The refusal message names the width, and stderr notes it once per process. Five tables have no lane-level operation and are marked any-wave: RoPE append, paged merge, Gumbel, rejection and xIELU (elementwise; its inert guard is dropped like the samplers'). `bitlinear_kernel_available` now reads its table, so BitNet is refused at load on wave64. `verify-kernel-port-dispatch` pins the marked tables, checks their HIP sources for lane intrinsics, and keeps the test seam `set_rocm_port_warp_size_for_tests` out of production code; a new script covers those rules. build.rs now tracks `kernel_port.*` and `gpu_backend.*`. Tests: the pure rule in mlxcel-core's default set, `tests/rocm_wave_size.rs` (hardware width, override ignored in a child, every port still selected on gfx1151), and `tests/rocm_wave64_port_refusal.rs` (simulated 64-lane report: wave32-only predicates false, launchers refuse, callers match the graph bytes, any-wave tables still selected). Closes #2147
5 tasks done
Review of the wavefront hold found gaps around it rather than in it. - `verify-kernel-port-dispatch` rule 5 now follows each marked table's `.rocm` getter to its holder and requires it to compile exactly the pinned source with no header argument, scans for lane builtins (`__builtin_amdgcn_ballot_w32`, `permlane`, `mbcnt`, DPP), wavefront macros and 32-lane index arithmetic, and rejects a source split across literals. Rule 6 allows an exact count of seam uses in its definition files instead of exempting them. New negative cases cover each. - `mamba1_selective_scan` is declared `-> Result<()>`, since `select_kernel_port` can now refuse there; Mamba, Jamba and the parity test `expect` it, as the SSM callers do. - Parity tests that assert a ROCm port exists (fused MoE, relu2, SSM, fused norm, Mamba1, add3) skip on a wave64 device instead of failing. - BitLinear's refusal names the BitNet load-time refusal rather than a graph fallback it does not have; the stderr note says so too. - `device_warp_size()` clears only an error it raised itself. LOCAL_FIXES item 31 moves to an "Additions to the fork's API" section. Refs #2147
inureyes
added a commit
that referenced
this pull request
Oct 6, 2026
…2181) ## Summary The ROCm overlay cached compiled `fast::hip_kernel` modules under device index plus kernel name, and the name covers only `template_args`. `build_kernel` also writes each input's and output's dtype and each input's 0-d-ness into the source, so a later call that differed only there reused the first module and returned wrong values (or, for a 0-d input after a 1-d one, launched with an argument list the module does not declare). `CustomKernel::eval_gpu` now passes `name_ + "_" + hex(std::hash<std::string>{}(source_))` as the module name, copied from upstream CUDA at the pin (`mlx/backend/cuda/custom_kernel.cpp` around line 313, upstream ml-explore/mlx#4273). The kernel symbol and `get_jit_module`'s device-index prefix are unchanged, so `rocprofv3` traces keep their `custom_kernel_*` names. ## What changed - Overlay fix in `patches-rocm/mlx/backend/rocm/custom_kernel.cpp`, recorded as LOCAL_FIXES item 33 (#2179 took 32), kept in mlxcelverse under the fork policy. - Test-only bridge fixture `rocm_jit_key_probe(input, f16_output)` (wrapper `mlxcel_core::rocm_faults::jit_key_probe_array`): one kernel name, no template args, an f32 or f16 output of `inp_shape[0]` (or a -1 sentinel for a 0-d input) followed by the input's values. - `tests/rocm_custom_kernel_jit_key.rs`: f32 then f16 then bf16 inputs, f32 then f16 outputs at one input dtype, and 1-d then 0-d input in a child process as `tests/rocm_gpu_faults.rs` does. - `scripts/ci/check_kernel_dtype_keys.py` and `docs/code-guidelines.md`: the CUDA paragraph now says CUDA keys by source hash since upstream 9f5f7931 (present at pin 81ba1c6a), ROCm does too after this change, and the check stays as defense in depth against a re-vendored fork, pin bump or new backend bringing a name-only key back. No launch drops its explicit dtype keys: the keys cost nothing, several are used as type aliases in kernel bodies, and the rest are the insurance the check exists for. The checker's rule, scope and pinned count (9) are unchanged. The same stale CUDA description in the keying comments of `turbo/paged_attention.cpp` and `turbo/paged_attention_v2.cpp` is corrected (comment-only). ## Verification On gfx1151 (ROCm, this host): - `cargo test --profile test-fast --features rocm --test rocm_custom_kernel_jit_key -- --test-threads=1`: 3 passed. With the name-only key restored in the overlay, all 3 fail: the f16-input call returns `[8, -0.0313, 131336, 3.5e13, 8, 0, ...]` for input `[0.5, -1.25, 3, 8, -0.75, ...]`, the f16-output call returns `[0, 2.5, 0, 1.75, ...]`, and the 0-d child dies with SIGSEGV after its 1-d call. - `cargo clippy --features rocm --test rocm_custom_kernel_jit_key -- -D warnings` and `cargo clippy -p mlxcel-core --features rocm --lib -- -D warnings`: clean. - `make verify-versions verify-kernel-dtype-keys verify-kernel-port-dispatch verify-llama-compat verify-fmt verify-rocm-overlay` after rebasing onto b7116d1: pass (checker negative-case suite included). - `make verify-rocm` on the final head b11b4cc: OK, 152 test suites, 11972 passed, 0 failed, 382 ignored (the new file: 3 passed, 1 ignored child); clippy, smoke, overlay and script gates pass. Not verified: Metal and CUDA are not available on this host. The change touches only the ROCm overlay, a bridge function compiled to a throwing stub off ROCm, a ROCm-only test file (`#![cfg(feature = "rocm")]`), and comments and docs; no Metal or CUDA code path changes. Closes #2149
This was referenced Oct 7, 2026
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
The shuffle-based HIP ports have only run on 32-lane wavefronts, and their
#errorguards on__AMDGCN_WAVEFRONT_SIZEare inert because AMD clang 23 defines neither macro spelling (rechecked here withhipcc -E -dMfor gfx1151, gfx942 and gfx90a). This PR holds those ports to 32-lane devices at port selection, so a wave64 (CDNA) device takes the graph fallbacks instead of running unvalidated kernels. Nothing changes on gfx1151.What changed
port_for(kernel_port.cpp) returns a ROCm port only whenrocm_port_allowed(ports.rocm_any_wave_size, rocm_port_warp_size()), sohas_kernel_port, every*_available()predicate andselect_kernel_portagree. The refusal message names the width; stderr notes a non-32 wavefront once per process.mlx::core::rocm::device_warp_size()(hipDeviceAttributeWarpSize, ignoresMLX_ROCM_FORCE_WARP_SIZE, 0 on error), recorded as LOCAL_FIXES item 32 (new "Additions to the fork's API" section; perf(rocm): batch sorted mxfp4 gather_qmm prefill by expert #2164 took 31) and stubbed inno_rocm.cpp. It clears only an HIP error it raised itself.KernelPorts::rocm_any_wave_size(default false) is set on the five tables whose HIP body has no lane-level operation: RoPE append, paged merge, Gumbel, rejection, and xIELU. xIELU is not in the issue's list, but checking the current code shows its body is purely elementwise; its inert guard is dropped, as the sampler and RoPE bodies already have none. The other ten ROCm tables (bitlinear, ssm, mamba1_scan, moe_gateup, moe_down, moe_fc1_relu2, add3_layer_norm, fused_norm, paged_attention, paged_v2_partial) stay wave32-only. update(rocm): port fused_add_rms_norm and fused_rope_qk_append to HIP #2107's fused norm is wave32-only (shuffles) and its RoPE port is any-wave.bitlinear_kernel_availablenow reads its port table (it returned "any GPU backend" before), so BitNet is refused at load on wave64 rather than failing mid-request.verify-kernel-port-dispatch: rule 5 pins the any-wave tables with their HIP source, follows each table's.rocmgetter to its holder and requires exactly that source (no header argument, one raw literal), and fails on lane intrinsics, AMDGCN lane builtins, wavefront macros or 32-lane index arithmetic in it; rule 6 keeps the test seamset_rocm_port_warp_size_for_testsout of non-test code, with an exact use count in its definition files.scripts/ci/check_kernel_port_dispatch_test.sh(run by the make target and CI job) proves each rule fails on a mutated copy (13 cases).mamba1_selective_scanis now-> Result<()>, sinceselect_kernel_portcan refuse there; Mamba, Jamba and its parity testexpectthe gated launch. BitLinear's refusal names the BitNet load-time refusal instead of a graph fallback it lacks.test_support::kernel_ports::skip_for_wave32_only_rocm_portinstead of failing.build.rsnow trackskernel_port.*andgpu_backend.*; edits to them did not rebuild the bridge before (found when a negative check passed on a stale library).docs/installation.mdgains a "Wave64 devices" row; CONTRIBUTING and the paged-test skip message mention the hold.Verification (gfx1151)
cargo test --features rocm -p mlxcel-core --lib kernel_port: 2 passed (rocm_port_allowedtruth table, runs on every backend).tests/rocm_wave_size.rs:device_warp_size() == 32on gfx1151; withMLX_ROCM_FORCE_WARP_SIZE=64in a child (the child's device log confirms the override applied) it still reads 32; all 16 port predicates are true.tests/rocm_wave64_port_refusal.rs, with the reported width forced to 64 through the test seam: all 11 wave32-only predicates false, the five any-wave ones true;bitlinear_matmul,fused_add_rms_normandfused_moe_expert_kernel(moe_gateup) refuse naming the width;layers::fused_add_rms_norm(fusion on) andresidual_add3_layer_normreturn the graph pair's bytes. With the hold disabled inport_for, this test fails listing all 11 tables.cache::paged_detachtests fail on adebug_assertin the dev profile only (the gate'stest-fastprofile has no debug assertions; unrelated to ports). Apertus, BitNet and Cohere2 tests in the mlxcel lib: 26 passed.cpu_device_custom_kernel_gates: passed.make verify-versions verify-kernel-dtype-keys verify-kernel-port-dispatch verify-llama-compat verify-fmt verify-rocm-overlay: pass.make verify-rocmon the final head (merged with origin/main 08e2269): OK, 151 test suites, 11969 passed, 0 failed, 381 ignored; clippy, smoke, overlay and script gates pass. The targeted tests above were re-run after the review fixes (29 mlxcel-core, 30 Mamba/Jamba, the three integration files): all pass.scripts/rocm_gpu_guard.sh(clean attempt): greedy 64 tokens withMLXCEL_FUSED_ADD_RMSNORM=0and=1give identical text (33.88 and 33.90 tok/s, single run, not a performance claim);rocprofv3 --kernel-traceof a sampled run showscustom_kernel_mlxcel_fused_add_rms_norm(wave32-only, 544 launches) andcustom_kernel_mlxcel_gumbel_max_sample(any-wave) still selected, and no wavefront notice on stderr.Not verified: no wave64 device is available, so the refusal is proven only through the test seam. Metal and CUDA are not available on this host; the change touches the shared
KernelPortsstruct andport_for, whose Metal and CUDA cases are unchanged, andbitlinear_kernel_available, which answers the same on Metal and CUDA (both tables have ports).Closes #2147