Skip to content

fix(rocm): refuse wave32-only kernel ports on wave64 devices - #2179

Merged
inureyes merged 5 commits into
mainfrom
fix/issue-2147-wave-size-port-check
Oct 6, 2026
Merged

inureyes merged 5 commits into
mainfrom
fix/issue-2147-wave-size-port-check

Conversation

@inureyes

@inureyes inureyes commented Oct 6, 2026 •

Copy link
Copy Markdown
Member

Summary

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 macro spelling (rechecked here with hipcc -E -dM for 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 when rocm_port_allowed(ports.rocm_any_wave_size, rocm_port_warp_size()), so has_kernel_port, every *_available() predicate and select_kernel_port agree. The refusal message names the width; stderr notes a non-32 wavefront once per process.
  • The width is the hardware value from a new overlay call, mlx::core::rocm::device_warp_size() (hipDeviceAttributeWarpSize, ignores MLX_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 in no_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_available now 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 .rocm getter 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 seam set_rocm_port_warp_size_for_tests out 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_scan is now -> Result<()>, since select_kernel_port can refuse there; Mamba, Jamba and its parity test expect the gated launch. BitLinear's refusal names the BitNet load-time refusal instead of a graph fallback it lacks.
  • Parity tests that assert a ROCm port exists (fused MoE, relu2, SSM, fused norm, Mamba1, add3) skip on a wave64 device through test_support::kernel_ports::skip_for_wave32_only_rocm_port instead of failing.
  • build.rs now tracks kernel_port.* and gpu_backend.*; edits to them did not rebuild the bridge before (found when a negative check passed on a stale library).
  • Guard comments now name the host-side check as the protection; docs/installation.md gains 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_allowed truth table, runs on every backend).
  • tests/rocm_wave_size.rs: device_warp_size() == 32 on gfx1151; with MLX_ROCM_FORCE_WARP_SIZE=64 in 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_norm and fused_moe_expert_kernel (moe_gateup) refuse naming the width; layers::fused_add_rms_norm (fusion on) and residual_add3_layer_norm return the graph pair's bytes. With the hold disabled in port_for, this test fails listing all 11 tables.
  • Existing port tests (debug profile): fused norm, RoPE, MoE, relu2, SSM, Mamba1, Gumbel, rejection, fixed-key, paged v2, MLA, residual add3: 357 passed; 2 cache::paged_detach tests fail on a debug_assert in the dev profile only (the gate's test-fast profile 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-rocm on 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.
  • Meta-Llama-3.1-8B-Instruct-4bit, release binary under scripts/rocm_gpu_guard.sh (clean attempt): greedy 64 tokens with MLXCEL_FUSED_ADD_RMSNORM=0 and =1 give identical text (33.88 and 33.90 tok/s, single run, not a performance claim); rocprofv3 --kernel-trace of a sampled run shows custom_kernel_mlxcel_fused_add_rms_norm (wave32-only, 544 launches) and custom_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 KernelPorts struct and port_for, whose Metal and CUDA cases are unchanged, and bitlinear_kernel_available, which answers the same on Metal and CUDA (both tables have ports).

Closes #2147

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
@inureyes inureyes added status:review Under review type:bug Bug fixes, error corrections, or issue resolutions priority:medium Medium priority area:core mlxcel-core: MLX FFI, primitives, KV cache, layers platform:linux Linux (CUDA / packaging) specific labels Oct 6, 2026
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 inureyes added status:done Completed and removed status:review Under review labels Oct 6, 2026
@inureyes
inureyes merged commit b7116d1 into main Oct 6, 2026
26 checks passed
@inureyes
inureyes deleted the fix/issue-2147-wave-size-port-check branch October 6, 2026 18:08
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
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

area:core mlxcel-core: MLX FFI, primitives, KV cache, layers platform:linux Linux (CUDA / packaging) specific priority:medium Medium priority status:done Completed type:bug Bug fixes, error corrections, or issue resolutions

Projects

None yet

Development

Successfully merging this pull request may close these issues.

fix(rocm): refuse wave32-only kernel ports on wave64 devices at runtime

1 participant