Repository navigation
fix(rocm): surface ROCm GPU failures through Event::error instead of NaN or hangs #1804
Description
Part of #1801. Phase 1. Depends on #1802 (closed; landed in PR #1818).
Status 2026-09-29: not started.
Event::erroris still never populated on ROCm: no file undersrc/lib/mlx-cpp/patches-rocm/callsset_error, and the overlay's event, worker, device and eval files have changed only once since they landed with PR #1818 (PR #1824, a compile-time kernel-argument size check indevice.h). The spike failures listed below have since been fixed at their source or no longer reproduce, so this issue no longer has a natural reproducer and its tests must provoke a fault deliberately. See the refresh log.
Context
Upstream ml-explore/mlx#3742 made array::is_available() throw on a failed launch and consume the error, carried through Event::error() (std::atomic<Error*>). The mlxcel bridge relies on this: the drain in src/lib/mlxcel-core/cpp/mlx_cxx_bridge.cpp (around lines 5571-5610) assumes a failed GPU step surfaces as an exception that becomes a Rust error.
Update 2026-09-29 (location moved): the drain is now
drain_pending_verification()atsrc/lib/mlxcel-core/cpp/mlx_cxx_bridge.cpp:5844, and it no longer callsarray::is_available(). It readsstashed_launch_state()(:5812), which reports a launch asFailedonly when the signaled event carries an error pointer. The comment above it (:5795-5809) records that Metal's completion handler and the CPU scheduler store the error before signaling, and that CUDA attaches none. On ROCm a failed launch therefore reads asLandedand its outputs are consumed as if they were written.
The ROCm overlay in mlxcelverse only has storage for the error (added during the retarget in #1801); nothing in the ROCm command encoder, event or completion path ever calls set_error. During the spike, failures looked like this instead:
- mxfp4
quantized_matmulhung until killed; not reproduced since (2026-09-29): PR feat(rocm): vendor the ROCm backend into mlxcelverse, add rocm feature #1818 reports mxfp4quantized_matmulandgather_qmmmatching a dequantized f32 reference and gpt-oss-20b generating, though no record names the change that removed the hang; fix(rocm/quant): mxfp4 qmm hang and GPU quantize launch failure, with an affine 4-bit fallback #1808 re-runs the case. - the unfixed mxfp8 path produced NaN and then
HSA_STATUS_ERROR_MEMORY_FAULTwith a queue hang dump; fixed at the source (2026-09-29): the E8M0 scale-type dispatch fix,patches-rocm/LOCAL_FIXES.mditem 8, landed with PR feat(rocm): vendor the ROCm backend into mlxcelverse, add rocm feature #1818. - a failed
hipLaunchKernel("invalid configuration argument") threw from the launch site, but asynchronous faults did not; fixed at the source (2026-09-29): the instance seen wasarg_reduceexceeding AMD's per-dimension grid limit,LOCAL_FIXES.mditem 11, PR feat(rocm): vendor the ROCm backend into mlxcelverse, add rocm feature #1818. The synchronous-throw versus silent-async-fault distinction still stands.
A user-facing inference server cannot tell these apart from a slow model.
Update 2026-09-29 (other fault classes found since): narrow index dtypes in
Gather/Scattercaused anHSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATIONafter which the process spun instead of failing (#1853, fixed in PR #1856,LOCAL_FIXES.mditem 17). #1876 (open) is a hang with the GPU idle: creating a hipFFT plan blocks forever inside the library once too many plans are alive. That is a host-side block rather than a device fault, so it is out of scope here except that the optional watchdog in step 4 would report it.
Update 2026-09-29 (not a GPU failure class): mlxcel's custom-kernel launchers used to abort the process on ROCm. Since PRs #2018, #2026 and #2029 every launcher resolves its port through
select_kernel_port(src/lib/mlx-cpp/turbo/kernel_port.h), which throws a catchable error naming the entry point and the fallback instead. Those refusals are synchronous and host-side, already reach Rust as errors, and are not part of this issue. What remains here is GPU-side: NaN outputs, hangs, and queue faults.
Scope
Make ROCm launch failures and asynchronous GPU faults reach Event::error so that MLX throws and mlxcel reports an error. Deadlock detection for kernels that spin forever is out of scope beyond a documented watchdog option.
Implementation plan
- Audit the ROCm command encoder (
mlx/backend/rocm/device.*,worker.*,event.hip,eval.cpp) for everyhip*call whosehipError_tis ignored (the build emits-Wunused-valuewarnings for these, e.g.hipGraphDestroyindevice.h). Check them and convert failures into an MLXError. (2026-09-29: still open. Unchecked calls remain, for examplehipGraphDestroyatpatches-rocm/mlx/backend/rocm/device.h:496-503andhipGraphExecDestroy/hipGraphDestroyatdevice.cpp:429-436and:639-644, anddevice.h:504discards the pending error with(void)hipGetLastError(). PR fix(rocm): make scatter apply every update, check kernel arg sizes #1824 added astatic_assertinCommandEncoder::add_kernel_nodethat each kernel argument has its parameter's size, which removes one source of silent wrong-address reads at compile time but does not check any return value.) - On stream completion, query
hipStreamQuery/hipGetLastError(or the completion callback status) and callset_erroron the signaled event, mirroring the Metal completion handler (mlx/backend/metal/device.cpp,event.set_error(error_)) and the scheduler path (mlx/scheduler.cpp). (2026-09-29: still open. The overlay'sEvent::error()accessor is atpatches-rocm/mlx/backend/rocm/event.hip:388and the storage isEventImpl::errorat:319; nothing writes it.) - Treat an HSA memory fault that kills the queue as fatal for the stream: mark pending events with an error so waiters return instead of blocking forever.
- Optional: an
MLXCEL_GPU_WATCHDOG_SECSstyle timeout in the bridge drain that reports a stuck GPU step as an error with the last kernel name. - Keep the change inside
patches-rocm/and record it inpatches-rocm/UPSTREAMas a local fix to upstream later (chore(rocm): mlxcelverse ROCm fork sync script, MLX pin-bump procedure, and upstreaming local fixes #1813). (Corrected 2026-09-29: local changes are recorded inpatches-rocm/LOCAL_FIXES.md, notUPSTREAM, which only pins the fork commit and points atLOCAL_FIXES.md. Item 7 there already names this issue and should be updated when the fix lands.)
Acceptance criteria
- A deliberately invalid launch (for example a kernel launched with an oversized block) returns an error to Rust through the bridge instead of aborting. (
tests/rocm_gpu_faults.rs::oversized_block_launch_fails_the_eval_and_leaves_the_device_usable) - An asynchronous device fault (for example an out-of-bounds read in a test kernel) makes the waiting
evalthrow within a bounded time instead of hanging. (tests/rocm_gpu_faults.rs::out_of_bounds_write_fails_the_eval_within_bounded_time, reported in 195 ms on gfx1151) -
mlxcel-serverreturns an HTTP error for a request whose forward pass fails, and keeps serving later requests if the device is still usable. (The batch scheduler already maps atry_evalerror to a per-request failure; after a queue fault the device is not usable for the rest of the process, see the PR.) - No behavior change on Metal or CUDA. (Overlay-only change plus a ROCm-gated bridge fixture; not run on those hosts.)
- (Added 2026-09-29)
docs/installation.md:480("GPU faults: may show up as NaN output or a hang instead of an error") is updated to the new behavior.
Validation
A small test binary or #[test] behind --features rocm that triggers each failure class, plus a manual run of the pre-fix mxfp4 repro from #1808 showing an error instead of a hang.
Update 2026-09-29: the pre-fix mxfp4 repro no longer hangs per PR #1818, so it cannot serve as the manual check. Use a deliberately faulting test kernel instead. Where the test lands in
mlxcel-core's lib tests, note that until #1806 is fixed the ROCm gate's-p mlxcel-core --librun aborts atffi_tests::compiled_qgelu_mlp_global_scale_native_nvfp4_prefill_matches_reference, and no test sorting after it runs, so run the new test directly. Gate:make verify-rocmon the gfx1151 host.
References
- Upstream contract: Propagate CPU errors to events ml-explore/mlx#3742
- Bridge drain:
src/lib/mlxcel-core/cpp/mlx_cxx_bridge.cpp(around 5571-5610; moved, nowstashed_launch_stateat:5812anddrain_pending_verificationat:5844) - ROCm event storage added in the retarget:
mlx/backend/rocm/event.hip(EventImpl::error; in this reposrc/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/event.hip:319) - Overlay change log:
src/lib/mlx-cpp/patches-rocm/LOCAL_FIXES.md(item 7 tracks this issue)
Refresh log
2026-09-29
- Verified against
origin/mainatd8d34e2b: no ROCm code populatesEvent::error; the only change to the overlay encoder/event files since PR feat(rocm): vendor the ROCm backend into mlxcelverse, add rocm feature #1818 vendored them is PR fix(rocm): make scatter apply every update, check kernel arg sizes #1824's argument-sizestatic_assert. Status stays not started. - Marked the three spike failures as fixed at their source (PR feat(rocm): vendor the ROCm backend into mlxcelverse, add rocm feature #1818,
LOCAL_FIXES.mditems 8 and 11) or, for the mxfp4 hang, no longer reproduced and replaced the mxfp4 manual repro with a deliberate fault, since no natural reproducer remains. - Recorded that custom-kernel launcher refusals now throw a catchable error through
select_kernel_port(PRs fix(rocm): guard the four remaining custom-kernel launchers #2018, refactor(core): choose custom-kernel ports through one helper #2026, refactor: finish the kernel port standardization #2029) and are outside this issue; GPU-side NaN, hangs and queue faults remain in scope. - Corrected the bridge drain location (
mlx_cxx_bridge.cpp:5812,:5844), the overlay record file (LOCAL_FIXES.md, notUPSTREAM), and added fix(rocm): Gather reads narrow index dtypes as int64 and faults the GPU queue #1853 and fix(rocm): hipFFT blocks when too many plans are alive, root cause unknown #1876 as related fault classes.
Activity
fix(rocm): keep CPU-stream fence signals on the scheduler path
fix(rocm): keep recovered allocation failures out of launch checks
fix(rocm): validate MLX_ROCM_GPU_WATCHDOG_SECS and hide fault test su…
fix(rocm): surface GPU failures through Event::error instead of NaN o…
fix(rocm): keep CPU-stream fence signals on the scheduler path
fix(rocm): keep recovered allocation failures out of launch checks
fix(rocm): validate MLX_ROCM_GPU_WATCHDOG_SECS and hide fault test su…
update(ci): make the kernel dtype-key checker's scope fail-closed (#2055
fix(rocm/quant): commit mxfp4 regression tests and log the native rou…
fix(rocm): surface GPU failures through Event::error instead of NaN o…