Skip to content

fix(rocm): validate MLX_ROCM_FFT_CACHE_SIZE (0 is UB, junk throws) - #2073

Merged
inureyes merged 3 commits into
mainfrom
fix/issue-2051-rocm-fft-cache-env
Sep 30, 2026
Merged

inureyes merged 3 commits into
mainfrom
fix/issue-2051-rocm-fft-cache-env

Conversation

@inureyes

Copy link
Copy Markdown
Member

Closes #2051

MLX_ROCM_FFT_CACHE_SIZE was parsed with an unchecked std::stoul in the overlay's lru_cache.h. Now capacity_from_env() applies the strtol rules of gpu_watchdog_seconds(): 1 to INT_MAX is used, unset gives the default silently, and anything else prints one line, [ROCm] ignoring invalid MLX_ROCM_FFT_CACHE_SIZE="<value>" (expected a positive integer); using the default 128, then uses the default. LRUBytesKeyCache and LRUCache(size_t) throw LRUCache requires capacity > 0. on 0, as upstream CUDA's LRUCache does.

  • fft.hip: the fft_plan_cache() comment states the rule. Editing it also forces a warm build to recompile fft.hip: the overlay's CMakeLists.txt makes each HIP object depend only on its own .hip file, so a header-only change leaves hip_objs/fft.o stale (seen here: the first run with the fix still failed until fft.hip changed).
  • tests/rocm_fft_cache_env.rs: one child process per value (0, -1, abc, empty, 8abc, 0x10, 99999999999999999999999, 2147483648, 16, 1, 2147483647, unset), run one at a time. Each child runs rfft on the GPU twice over two lengths and compares it to the CPU stream; the parent checks exit status and that stderr has exactly one warning line for invalid values and none for valid ones.
  • LOCAL_FIXES.md item 25 (24 is reserved for fix(rocm): SliceUpdate and DynamicSliceUpdate donate a source array that is still referenced #2052), marked as an upstreaming candidate for chore(rocm): mlxcelverse ROCm fork sync script, MLX pin-bump procedure, and upstreaming local fixes #1813. docs/environment-variables.md and docs/installation.md give the valid range.

Verification

On gfx1151 (ROCm 7.15):

  • cargo test --features rocm --test rocm_fft_cache_env: passes, all 12 cases ok.
  • The same test with lru_cache.h reverted to main fails 8 of 8 invalid cases: 0 and 0x10 fail the first GPU rfft with std::bad_alloc, abc, empty and the overflow with stoul, and -1, 8abc and 2147483648 print no warning.
  • cargo test --features rocm --test rocm_fft_plan_cache and --test rocm_gpu_faults: pass.
  • cargo clippy --features rocm --test rocm_fft_cache_env -- -D warnings, cargo fmt --check, make verify-versions verify-kernel-dtype-keys verify-kernel-port-dispatch verify-llama-compat verify-fmt, cargo test --test dead_doc_pointers: pass.
  • A host g++ 14 probe of the header with -D_GLIBCXX_ASSERTIONS covered the same values plus 16 and +16 (accepted as 16).

Not verified: Metal and CUDA (not available on this host). The change is limited to patches-rocm/, which only ROCm builds copy, plus docs and a #![cfg(feature = "rocm")] test.

The ROCm overlay parsed MLX_ROCM_FFT_CACHE_SIZE with an unchecked std::stoul: 0 (and 0x10, read as 0) made the first plan insert pop an empty list, abc, empty and overflowing values threw a bare "stoul" out of every FFT, and -1 and 8abc were accepted silently as SIZE_MAX and 8.

capacity_from_env() in lru_cache.h now parses it with the strtol rules of MLX_ROCM_GPU_WATCHDOG_SECS, accepts 1 to INT_MAX, and otherwise prints one stderr line naming the variable and the default (128) and uses the default. Both LRU constructors throw CUDA's "LRUCache requires capacity > 0." for 0. The fft.hip comment states the rule; touching it also makes warm builds recompile fft.hip, since HIP objects do not track header dependencies.

tests/rocm_fft_cache_env.rs runs one child per value and checks the warning count and a GPU rfft against the CPU stream; with the fix reverted 8 of 8 invalid cases fail on gfx1151. LOCAL_FIXES item 25 records it as an upstreaming candidate (#1813), and the valid range is documented.

Closes #2051
@inureyes inureyes added type:bug Bug fixes, error corrections, or issue resolutions priority:low Low priority area:core mlxcel-core: MLX FFI, primitives, KV cache, layers platform:linux Linux (CUDA / packaging) specific status:review Under review labels Sep 30, 2026
@inureyes inureyes added status:done Completed and removed status:review Under review labels Sep 30, 2026
@inureyes
inureyes merged commit 0d17303 into main Sep 30, 2026
22 of 25 checks passed
@inureyes
inureyes deleted the fix/issue-2051-rocm-fft-cache-env branch September 30, 2026 05:04
inureyes added a commit that referenced this pull request Oct 7, 2026
#2193)

## Summary

The ROCm overlay's integer GEMM knobs used `strtol` plus a bare `static_cast<int>` (or `atoi`), so out-of-range values wrapped silently: `MLX_ROCM_WMMA_QMM_MAX_M=4294967297` became a ceiling of 1 and moved every bf16 GEMM of two or more rows to dequantize + hipBLASLt. They now go through one helper with the #2073 rule.

- New overlay header `mlx/backend/rocm/env_int.h` with `env_int_or_default(name, default, min, max, expected[, default_description])`: whole-string base-10, `ERANGE`, junk and out-of-range values print `[ROCm] ignoring invalid NAME="value" (expected ...); using ...` once and take the default; unset or empty is silent. The optional last argument names sentinel defaults ("the per-device default") instead of printing `-1`.
- `parse_positive_int_env` and the three `parse_non_negative_int_env` copies are deleted. Routed through the helper from statics: `MLX_ROCM_QMM_DEQUANT_M_THRESHOLD` and `MLX_ROCM_WMMA_QMM_MAX_M` (1 to `INT_MAX`; the threshold is now read once instead of per call), the four `MLX_ROCM_GEMM_{F32,BF16}{,_BATCHED}_SOLUTION_INDEX` (0 to `INT_MAX`; the three files' copies of their readers are now one inline pair in `gemms/rocblas_gemm.h`, so a bad index warns once per process, not once per file), `MLX_ROCM_QMM_DEQUANT_CACHE_SIZE` (0 to `INT_MAX`, 0 still off), and from the #2180 filing `MLX_ROCM_QMV_TILE_N` (1 to `TILE_N_MAX` = 32, the tiled kernels' launch bound; now read once), `MLX_ROCM_GROUPED_PREFILL_MIN_B`, `MLX_ROCM_MOE_SEG_MIN`, `MLX_GRIDX_MULT` (1 to `INT_MAX`). `MLX_GRIDX_MULT` used to clamp `0`/negatives to 1; it now warns and uses 4, and the grid size it scales is computed in 64 bits.
- `use_hip_graphs()` comment no longer names a prefill GEMM.
- LOCAL_FIXES item 34 (not an upstreaming candidate), README backend count 109, `docs/environment-variables.md` states ranges and the warning.

## Verification

On gfx1151 (Radeon 8060S), rebased on bbd0509:

- `cargo test --release --features rocm --test rocm_qmm_env -- --test-threads=1`: pass, 15 child cases (ceiling unset, `128`, empty, `32`, `4294967297`, `2147483648`, `-1`, `0`, `12abc`; threshold `0`; cache `4294967304`, `0`; tile width `16`, `abc`, `33`), checking warning lines, the route `quantized_matmul_matches_dense_gemm` reports for bf16 `[64, 4096]` x 4-bit `[4096, 4096]` g64, and GPU GEMM and GEMV against an f32 CPU reference.
- Fails with the fix reverted: I restored `qmm.hip`, `matmul.cpp`, `gemms/rocblas_gemm.{cpp,h}` and `device.cpp` from main, rebuilt, and 9 of 15 cases failed (`4294967297` reported the dense route, and none of the nine invalid values warned). With only the threshold read made per call again, the threshold case fails with one warning per GEMM.
- `cargo clippy --release --features rocm --test rocm_qmm_env -- -D warnings`, `make verify-rocm-overlay verify-versions verify-kernel-dtype-keys verify-kernel-port-dispatch verify-llama-compat verify-fmt`, `cargo test --test dead_doc_pointers`: pass.
- Full `make verify-rocm` with `MLXCEL_ROCM_SMOKE_MODEL=models/mlx/Qwen3-0.6B-4bit` at `01a28a10` on `bbd05099` (the later commit only edits the technical report): `[verify-rocm] OK`, 153 cargo test suites, 11997 passed, 0 failed, 384 ignored; the ROCm smoke generated 32 tokens on the GPU.
- Review: pr-reviewer found no CRITICAL or HIGH issues; its two MEDIUM findings (solution indexes warning once per file, no test for the threshold caching) are fixed in `e54ffc1e`. Security and performance check: values are parsed with `strtol` into `long` and range-checked before narrowing, the grid size multiplication is done in 64 bits, and every read is a function-local static, so there is no per-call cost.

Not verified: Metal and CUDA (not available on this host). The change touches only `patches-rocm/`, which Metal and CUDA builds never copy, plus a `#![cfg(feature = "rocm")]` test.

## Left for #2180

Behavior unchanged for boolean and presence parsers: `MLX_ROCM_USE_ASYNC_POOL`, `MLX_ROCM_FORCE_ASYNC_POOL`, `MLX_ROCM_NO_ASYNC_POOL`, `MLX_ROCM_ALLOW_MANAGED_FALLBACK`, `MLX_ROCM_NO_MANAGED_FALLBACK`, `MLX_ROCM_QMV_NO_TILED`, `MLX_ROCM_QMV_6BIT_SLOW`, `MLX_ROCM_NO_HIPBLASLT_EPILOGUE`, `MLX_ROCM_HIPBLASLT_TUNE`, `MLX_ROCM_GEMM_DEBUG`, `MLX_ROCM_MOE_PIN_STATS`, `MLX_ROCM_GATHER_QMV_USE_TILED` (presence), `MLX_ROCM_GEMM_BF16` and the `MLX_ROCM_MOE_*` first-character flags, and `MLX_ROCM_FINEGRAINED` and `MLX_ROCM_GROUPED_PREFILL` (`atoi` nonzero). #2180's table should describe the new ranges for `MLX_ROCM_QMV_TILE_N` and `MLX_ROCM_MOE_SEG_MIN`. Also untouched: `MLX_GRAPH_REPLAY_SLOTS` in `device.cpp`, whose `std::max<size_t>(2, atoi(e))` turns `-1` into `SIZE_MAX` (out of scope per #2180; needs its own issue).

Closes #2152
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:low Low 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): validate MLX_ROCM_FFT_CACHE_SIZE (0 is UB, junk throws)

1 participant