You signed in with another tab or window. Reload to refresh your session.You signed out in another tab or window. Reload to refresh your session.You switched accounts on another tab or window. Reload to refresh your session.Dismiss alert
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.
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_matmul hung until killed;
the unfixed mxfp8 path produced NaN and then HSA_STATUS_ERROR_MEMORY_FAULT with a queue hang dump;
a failed hipLaunchKernel ("invalid configuration argument") threw from the launch site, but asynchronous faults did not.
A user-facing inference server cannot tell these apart from a slow model.
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 every hip* call whose hipError_t is ignored (the build emits -Wunused-value warnings for these, e.g. hipGraphDestroy in device.h). Check them and convert failures into an MLX Error.
On stream completion, query hipStreamQuery/hipGetLastError (or the completion callback status) and call set_error on the signaled event, mirroring the Metal completion handler (mlx/backend/metal/device.cpp, event.set_error(error_)) and the scheduler path (mlx/scheduler.cpp).
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_SECS style timeout in the bridge drain that reports a stuck GPU step as an error with the last kernel name.
A deliberately invalid launch (for example a kernel launched with an oversized block) returns an error to Rust through the bridge instead of aborting.
An asynchronous device fault (for example an out-of-bounds read in a test kernel) makes the waiting eval throw within a bounded time instead of hanging.
mlxcel-server returns an HTTP error for a request whose forward pass fails, and keeps serving later requests if the device is still usable.
No behavior change on Metal or CUDA.
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.
Part of #1801. Phase 1. Depends on #1802.
Context
Upstream ml-explore/mlx#3742 made
array::is_available()throw on a failed launch and consume the error, carried throughEvent::error()(std::atomic<Error*>). The mlxcel bridge relies on this: the drain insrc/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.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:quantized_matmulhung until killed;HSA_STATUS_ERROR_MEMORY_FAULTwith a queue hang dump;hipLaunchKernel("invalid configuration argument") threw from the launch site, but asynchronous faults did not.A user-facing inference server cannot tell these apart from a slow model.
Scope
Make ROCm launch failures and asynchronous GPU faults reach
Event::errorso 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
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.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).MLXCEL_GPU_WATCHDOG_SECSstyle timeout in the bridge drain that reports a stuck GPU step as an error with the last kernel name.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).Acceptance criteria
evalthrow within a bounded time instead of hanging.mlxcel-serverreturns an HTTP error for a request whose forward pass fails, and keeps serving later requests if the device is still usable.Validation
A small test binary or
#[test]behind--features rocmthat triggers each failure class, plus a manual run of the pre-fix mxfp4 repro from #1808 showing an error instead of a hang.References
src/lib/mlxcel-core/cpp/mlx_cxx_bridge.cpp(around 5571-5610)mlx/backend/rocm/event.hip(EventImpl::error)