fix(rocm): surface GPU failures through Event::error instead of NaN or hangs

docs/environment-variables.md

@@ -192,7 +192,17 @@ ROCm builds (`--features rocm`) read `ROCM_PATH` (default `/opt/rocm`) and
 also read `MLX_ROCM_FFT_CACHE_SIZE` (default 8), the number of hipFFT plans kept
 alive. The default is low on purpose: past roughly a dozen live plans, creating
 the next one blocks inside hipFFT with the GPU idle (lablup/mlxcel#1825). Raise
-it only while measuring that limit.
+it only while measuring that limit. `MLX_ROCM_GPU_WATCHDOG_SECS` (default `0`,
+off) is the longest, in whole seconds, a single host wait for GPU work may
+block before it fails with an error (lablup/mlxcel#1804). A value that is not a
+non-negative integer (for example `10s`) is ignored with a stderr warning and
+leaves the watchdog off. A GPU fault is reported without it, in
+about a second; the watchdog is for a kernel that never finishes, which cannot
+be told from a slow one, so it is opt-in and must exceed the longest prefill
+the server accepts. The stream stays wedged behind the kernel either way, so a
+watchdog failure still means restarting the process. It does not apply under
+`MLX_EVENT_BLOCKING`, whose blocking waits return promptly on a fault but have
+no poll loop to time.
 
 ## OpenXLA / StableHLO backend variables
 

docs/installation.md

@@ -477,7 +477,7 @@ thinking model that spends its whole budget inside the thinking block reports
 | NVFP4 checkpoints | No native kernel. ModelOpt NVFP4 checkpoints (for example the Gemma 4 NVFP4 exports) are converted to affine 4-bit at load with no environment variables, and the load log names the route and the reason. A layer that cannot be converted fails the load with its name and the reason. MLX-native NVFP4 exports (`"mode": "nvfp4"` in `config.json`, such as `mlx-community/*-nvfp4`) have no load-time conversion and are refused at load; use an affine export instead (lablup/mlxcel#1806). |
 | Affine MoE models (for example Qwen3-30B-A3B) | Run natively. The fused MoE path used to abort on ROCm; lablup/mlxcel#1803 routes it to the MLX graph fallback, so `MLXCEL_FUSED_MOE=0` is no longer needed. |
 | mlxcel's fused kernels (sampling, fused norm, RoPE + KV append, paged attention) | Run as MLX graph fallbacks on the paths that have one (lablup/mlxcel#1803); ROCm ports are lablup/mlxcel#1814. |
-| GPU faults | May show up as NaN output or a hang instead of an error (lablup/mlxcel#1804). |
+| GPU faults | Reported as errors (lablup/mlxcel#1804). A launch HIP rejects (an oversized block, no code object for the device) fails the evaluation that issued it and the device stays usable. An asynchronous fault (an out-of-bounds access) fails the evaluation waiting on it within about a second instead of hanging or returning NaN; the server fails that request and answers later ones with the same error, because after a queue fault the HIP runtime rejects every call for the rest of the process, so the process has to restart, and its shutdown may need a SIGKILL (HIP's teardown waits on callbacks the faulted queue never runs). A kernel that never finishes is not detected; `MLX_ROCM_GPU_WATCHDOG_SECS` (default off) fails any single host wait that outlives it. |
 | Memory estimation on UMA hosts | Correct. Measured on the tested configuration: the ROCm allocator reports a nonzero cap (76.80 GiB of the 96 GiB carve-out), which the estimator reads before it would ever reach host RAM, so nothing that fits the carve-out is refused for that reason (lablup/mlxcel#1805). |
 | Diagnostics | Report the AMD vendor, device name, `gfx` target and device memory, and no longer print a CUDA compute capability for it (lablup/mlxcel#1805). A binary whose compiled `gfx` list does not cover the device refuses to start rather than failing at the first kernel launch. |
 | `mlxcel-server` chat completions | Work for dense and affine MoE checkpoints, streaming and non-streaming; verified with `scripts/server_chat_smoke.sh`. |

src/lib/mlx-cpp/patches-rocm/LOCAL_FIXES.md

@@ -12,7 +12,7 @@ The fork branched from upstream on 2026-06-26 (`39886de4`). Moving its ROCm code
 4. **Custom kernels.** `mlx/backend/rocm/custom_kernel.cpp` passes an empty `CompileOptions::Data` before the aliases argument. `output_input_aliases_` stays out of `CustomKernel::state()` because `mlx/export.cpp` serializes that tuple and has no encoding for it (the fork's `state()` did not include it either).
 5. **SDPA.** `ScaledDotProductAttention::use_fallback` takes upstream's new `force_fused` argument, with the CUDA semantics: throw if fused is forced and no fused kernel applies.
 6. **New upstream primitives.** `GatherQQMM`, `SearchSorted` and `fast::CrossEntropy` (with its VJP, using the graph fallback) get `NO_GPU` stubs in `mlx/backend/rocm/primitives.cpp`.
-7. **Event errors.** `Event::error()` (ml-explore/mlx#3742) has storage in the ROCm `EventImpl`. ROCm GPU failures do not populate it yet; tracked in lablup/mlxcel#1804.
+7. **Event errors.** `Event::error()` (ml-explore/mlx#3742) is populated on ROCm since lablup/mlxcel#1804, on the Metal model: each `CommandEncoder` owns an `Error`, and events signaled on a stream point at that stream's `Error` once the stream fails. Measured on gfx1151 with ROCm 7.15 (a probe kernel writing 2^40 bytes past its buffer): the HIP runtime does not abort the process on a queue fault, with or without `HIP_SKIP_ABORT_ON_GPU_ERROR` (the variable made no observable difference, so mlxcel does not set it); it prints the fault and the kernel name to stderr, then every HIP call on every thread returns `hipErrorIllegalAddress` for the rest of the process, `hipDeviceReset` included, and no queued `hipLaunchHostFunc` callback ever runs. That defeated all three of the backend's waits: `HipEvent::wait()` spun while `hipEventQuery` was not `hipSuccess`, `AtomicEvent::wait()` polled a counter that only a host callback lands, and `CommandEncoder::synchronize()` waited on a promise a host callback resolves. `HipEvent::wait()` now ends on any status other than `hipErrorNotReady`; `AtomicEvent` keeps the stream it was signaled from and polls `hipStreamQuery` on it about once a millisecond while waiting; `Worker::commit()` returns the `hipLaunchHostFunc` status and the encoder records a failure; `synchronize()` checks every status and polls the stream while it waits for the handlers. `Event::wait()` calls `check_error()` before and after the wait as Metal and CUDA do, `Event::is_signaled()` reports a failed event as signaled with the error attached (so `array::is_available()` detaches it and throws instead of pending forever), and `Event::wait(Stream)`/`Event::signal(Stream)` go through `scheduler::wait_event`/`signal_event` on CPU streams so the CPU-stream error path is upstream's. Synchronous launch failures: `launch_module_kernel` checks `hipModuleLaunchKernel`, and `gpu::eval` reads the thread's pending HIP error after each `eval_gpu` (kernels launched with `hipLaunchKernelGGL` or `<<<>>>` report a rejected launch no other way) and throws it as that primitive's error; a pending `hipErrorOutOfMemory` is skipped there, because it is what a recovered allocation failure leaves behind (the allocator's cache-release retry and managed fallback, hipBLASLt's workspace fallback) and an unrecovered one has already thrown. Under `MLX_GRAPH_PREFILL_REPLAY=1`, a launch that throws inside `launch_kernel`'s capture ends the capture and falls back to the eager launch, so a now-checked launch failure cannot leave the stream capturing. Destructor-path `hipGraphDestroy`/`hipGraphExecDestroy` calls discard their status explicitly, and `HipHandle::reset()` reports rather than throws, since on a faulted device every destroy fails and a throwing destructor would terminate the process at exit. `MLX_ROCM_GPU_WATCHDOG_SECS` (default off) bounds a single host wait for a kernel that never finishes; it is parsed with `strtol` and must be a whole non-negative integer, otherwise it warns once on stderr and stays off. Not covered: after a queue fault the device is dead for the process and its shutdown may hang, because HIP's teardown waits on the callbacks that never run (measured: the probe process hung at exit whenever a host callback was pending); the process has to be killed and restarted. Applies to the fork; to be proposed there.
 
 ## Fixes to the fork's kernels
 

src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/device.cpp

@@ -8,11 +8,13 @@
 #include "mlx/utils.h"
 
 #include <chrono>
+#include <cerrno>
 #include <cstdio>
 #include <cstdlib>
 #include <cstring>
 #include <future>
 #include <iostream>
+#include <limits>
 #include <map>
 #include <mutex>
 #include <sstream>
@@ -411,6 +413,12 @@ void Device::clear_encoders() {
   encoders_.clear();
 }
 
+CommandEncoder* Device::find_encoder(Stream s) {
+  std::lock_guard<std::mutex> lk(encoders_mtx_);
+  auto it = encoders_.find(s.index);
+  return it == encoders_.end() ? nullptr : it->second.get();
+}
+
 CommandEncoder::CommandEncoder(Device& d)
     : device_(d),
       stream_(d),
@@ -424,20 +432,137 @@ CommandEncoder::CommandEncoder(Device& d)
 }
 
 CommandEncoder::~CommandEncoder() {
+  // Destructor path: a failed destroy has nowhere to go, and on a device that
+  // has faulted every one of these returns the fault.
   for (auto& [key, pool] : exec_pool_) {
     for (auto& slot : pool) {
-      hipGraphExecDestroy(slot.exec);
+      (void)hipGraphExecDestroy(slot.exec);
       if (slot.source_graph) {
-        hipGraphDestroy(slot.source_graph);
+        (void)hipGraphDestroy(slot.source_graph);
       }
     }
   }
   if (build_graph_) {
-    hipGraphDestroy(build_graph_);
+    (void)hipGraphDestroy(build_graph_);
     build_graph_ = nullptr;
   }
 }
 
+// --- GPU failure reporting (lablup/mlxcel#1804) -----------------------------
+
+int gpu_watchdog_seconds() {
+  static const int secs = [] {
+    const char* e = std::getenv("MLX_ROCM_GPU_WATCHDOG_SECS");
+    if (!e || !*e) {
+      return 0;
+    }
+    // The whole string must be a non-negative decimal integer that fits an
+    // int. Anything else (trailing junk such as "10s", overflow, a negative
+    // value) leaves the watchdog off rather than guessing what was meant.
+    char* end = nullptr;
+    errno = 0;
+    long v = std::strtol(e, &end, 10);
+    if (end == e || *end != '\0' || errno == ERANGE || v < 0 ||
+        v > std::numeric_limits<int>::max()) {
+      std::fprintf(
+          stderr,
+          "[ROCm] ignoring invalid MLX_ROCM_GPU_WATCHDOG_SECS=\"%s\" "
+          "(expected a non-negative integer); the GPU watchdog stays off\n",
+          e);
+      return 0;
+    }
+    return static_cast<int>(v);
+  }();
+  return secs;
+}
+
+std::string describe_device_error(hipError_t status, const char* where) {
+  std::ostringstream oss;
+  if (status == kWatchdogExpired) {
+    oss << "[ROCm] GPU watchdog: a host wait (" << where << ") exceeded "
+        << "MLX_ROCM_GPU_WATCHDOG_SECS=" << gpu_watchdog_seconds()
+        << " seconds. The GPU stream is still busy or stuck; work queued "
+        << "behind it will not complete, so restart the process to use the "
+        << "GPU again.";
+    return oss.str();
+  }
+  oss << "[ROCm] GPU stream failed while " << where << ": "
+      << hipGetErrorString(status) << " (" << hipGetErrorName(status) << ", "
+      << static_cast<int>(status) << ").";
+  // After a queue fault the runtime fails every HIP call on every thread,
+  // hipGetDevice included; say so rather than let the caller retry.
+  int dev = 0;
+  bool context_gone = hipGetDevice(&dev) != hipSuccess;
+  (void)hipGetLastError();
+  if (context_gone) {
+    oss << " The HIP runtime reported the faulting kernel on stderr. The "
+        << "device context is unusable for the rest of the process (every "
+        << "HIP call now returns this error), so restart the process to use "
+        << "the GPU again.";
+  }
+  return oss.str();
+}
+
+void CommandEncoder::set_device_error(hipError_t status, const char* where) {
+  // Keep the earliest error, as the Metal completion handler does.
+  if (error_.valid()) {
+    return;
+  }
+  error_.set_message(
+      std::make_shared<std::string>(describe_device_error(status, where)));
+}
+
+void CommandEncoder::check_launch(const char* primitive) {
+  hipError_t status = hipGetLastError();
+  if (status == hipSuccess) {
+    return;
+  }
+  // A failed allocation that was recovered from leaves hipErrorOutOfMemory
+  // pending too: the allocator retries after releasing its cache or falls
+  // back to managed memory, and hipBLASLt runs without a workspace when its
+  // hipMalloc fails. An allocation that was not recovered from has already
+  // thrown, and a kernel launch does not report this status, so it is not
+  // this primitive's launch failure.
+  if (status == hipErrorOutOfMemory) {
+    return;
+  }
+  int dev = 0;
+  bool context_gone = hipGetDevice(&dev) != hipSuccess;
+  (void)hipGetLastError();
+  if (context_gone) {
+    std::string where = std::string("launching ") + primitive;
+    set_device_error(status, where.c_str());
+    error_.check();
+  }
+  std::ostringstream oss;
+  oss << "[ROCm] launching " << primitive
+      << " failed: " << hipGetErrorString(status) << " ("
+      << hipGetErrorName(status) << ", " << static_cast<int>(status) << ").";
+  throw std::runtime_error(oss.str());
+}
+
+std::unordered_map<int, Device>& get_devices();
+
+Error& record_stream_error(Stream s, hipError_t status, const char* where) {
+  if (s.device.type == mlx::core::Device::gpu) {
+    auto& devices = get_devices();
+    if (auto it = devices.find(s.device.index); it != devices.end()) {
+      if (auto* encoder = it->second.find_encoder(s)) {
+        encoder->set_device_error(status, where);
+        return encoder->error();
+      }
+    }
+  }
+  // Leaked on purpose: Event::set_error stores a raw pointer, so the object
+  // must outlive every event that could point at it.
+  static Error* fallback = new Error();
+  if (!fallback->valid()) {
+    fallback->set_message(
+        std::make_shared<std::string>(describe_device_error(status, where)));
+  }
+  return *fallback;
+}
+
 void CommandEncoder::add_temporary(const array& arr) {
   auto data = arr.data_shared_ptr();
   const array::Data* ptr = data.get();
@@ -904,7 +1029,9 @@ bool CommandEncoder::decode_capture_end_record(int slot) {
   // Stream capture records WITHOUT executing — run the exec once to actually
   // compute the record token's logits/state.
   CHECK_HIP_ERROR(hipGraphLaunch(exec, stream_));
-  worker_->commit(stream_);
+  if (hipError_t st = worker_->commit(stream_); st != hipSuccess) {
+    set_device_error(st, "queuing the completion callback of a graph launch");
+  }
   return true;
 }
 
@@ -914,7 +1041,9 @@ bool CommandEncoder::decode_capture_replay(int slot) {
     return false;
   device_.make_current();
   CHECK_HIP_ERROR(hipGraphLaunch(decode_cap_exec_[slot], stream_));
-  worker_->commit(stream_);
+  if (hipError_t st = worker_->commit(stream_); st != hipSuccess) {
+    set_device_error(st, "queuing the completion callback of a graph replay");
+  }
   return true;
 }
 
@@ -1121,18 +1250,52 @@ void CommandEncoder::commit() {
 
   node_count_ = 0;
 
-  // Put completion handlers in a batch.
-  worker_->commit(stream_);
+  // Put completion handlers in a batch. On a stream that has failed the
+  // callback cannot be queued and the handlers never run; record that so
+  // synchronize() and the event waiters fail instead of blocking on them.
+  if (hipError_t st = worker_->commit(stream_); st != hipSuccess) {
+    set_device_error(st, "queuing the completion callback of a commit");
+  }
 }
 
 void CommandEncoder::synchronize() {
-  (void)hipStreamSynchronize(stream_);
+  using namespace std::chrono_literals;
+  // Mirrors the Metal encoder: wait, then throw the stream's error. Every
+  // HIP status is checked because a failed stream returns its fault from
+  // each of these calls and never runs the completion handlers, so the
+  // promise below would otherwise be waited on forever (lablup/mlxcel#1804).
+  auto fail = [this](hipError_t st, const char* where) {
+    set_device_error(st, where);
+    error_.check();
+    throw std::runtime_error(describe_device_error(st, where));
+  };
+  if (hipError_t st = hipStreamSynchronize(stream_); st != hipSuccess) {
+    fail(st, "synchronizing the stream");
+  }
   auto p = std::make_shared<std::promise<void>>();
   std::future<void> f = p->get_future();
   add_completed_handler([p = std::move(p)]() { p->set_value(); });
   commit();
-  f.wait();
-  (void)hipStreamSynchronize(stream_);
+  if (error_.valid()) {
+    error_.check(); // the callback could not be queued
+  }
+  // The handler fires right after the stream drains, so poll the future at a
+  // fine grain and the stream (in case it faults meanwhile) about once a
+  // millisecond.
+  for (int polls = 1; f.wait_for(100us) != std::future_status::ready;
+       polls++) {
+    if (polls % 10 != 0) {
+      continue;
+    }
+    hipError_t st = hipStreamQuery(stream_);
+    if (st != hipSuccess && st != hipErrorNotReady) {
+      fail(st, "waiting for the stream's completion handlers");
+    }
+  }
+  if (hipError_t st = hipStreamSynchronize(stream_); st != hipSuccess) {
+    fail(st, "synchronizing the stream");
+  }
+  error_.check();
   // Stream is fully drained. Non-cached (no-reuse) execs reference these Packs
   // until now; cached-exec Packs live in their ExecSlot (clr#138) and are NOT
   // in these vectors, so clearing here is safe.