From 84e8ad1910d417479aa4c045ab90f17b296539fc Mon Sep 17 00:00:00 2001
From: Jeongkyu Shin <jshin@lablup.com>
Date: Wed, 30 Sep 2026 00:40:07 +0900
Subject: [PATCH 1/4] fix(rocm): set blocking-sync before first HIP queue, FFT
 cache to 128

The hipFFT plan cache was capped at 8 because the first plan eviction hung in hipfftDestroy's hipFree. The cause was not rocFFT but ordering: the allocator's first host read called hipStreamQuery(nullptr), creating the null stream's queue with non-interrupt completion signals, and only then did rocm::device() call hipSetDeviceFlags(hipDeviceScheduleBlockingSync). After the flip CLR no longer swaps those signals for interrupt ones, ROCr rejects the async handler hipFree's device sync attaches, and the wait never returns. Any hipFree or device-wide sync with pending null-stream work could hit it.

rocm::ensure_device_flags() sets the flag once per device index (thread-safe, only that device, caller's current device restored) and now runs before every path that can create a queue ahead of rocm::device(): unified_malloc, both raw_ptr() branches, the async-pool free streams, gpu::init() and staged_write's sync. The plan cache goes back to the CUDA value of 128, and the execute_fft comment no longer claims the captured plan outlives the call.

tests/rocm_fft_plan_cache.rs runs 48 round trips and GPU-vs-CPU checks in a child with a 16-entry cache and a 120 s budget; it passes, and times out with the ordering reverted. LOCAL_FIXES entry 18 updated, entry 20 added, docs updated.

Closes #1876
---
 docs/environment-variables.md                 |   9 +-
 docs/installation.md                          |   2 +-
 src/lib/mlx-cpp/patches-rocm/LOCAL_FIXES.md   |   3 +-
 .../mlx/backend/rocm/allocator.cpp            |  12 +
 .../patches-rocm/mlx/backend/rocm/device.cpp  |  72 ++++-
 .../patches-rocm/mlx/backend/rocm/device.h    |  14 +
 .../patches-rocm/mlx/backend/rocm/eval.cpp    |   2 +
 .../patches-rocm/mlx/backend/rocm/fft.hip     |  30 +-
 .../mlx/backend/rocm/host_stage.cpp           |   4 +
 tests/rocm_fft_plan_cache.rs                  | 279 ++++++++++++++++++
 10 files changed, 400 insertions(+), 27 deletions(-)
 create mode 100644 tests/rocm_fft_plan_cache.rs

diff --git a/docs/environment-variables.md b/docs/environment-variables.md
index 70cc07cc6..7a73b791c 100644
--- a/docs/environment-variables.md
+++ b/docs/environment-variables.md
@@ -189,10 +189,11 @@ CUDA builds also use non-`MLXCEL_*` variables such as `CUDA_HOME` and
 ROCm builds (`--features rocm`) read `ROCM_PATH` (default `/opt/rocm`) and
 `MLX_ROCM_ARCHITECTURES` (the `gfx` targets, default detected with `rocminfo`); see
 [Installation](installation.md#linux-with-amd-rocm-experimental). At runtime they
-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. `MLX_ROCM_GPU_WATCHDOG_SECS` (default `0`,
+also read `MLX_ROCM_FFT_CACHE_SIZE` (default 128, as on CUDA), the number of
+hipFFT plans the plan cache keeps; a smaller value only trades memory for plan
+rebuilds. It used to default to 8 to dodge a hang at the first plan eviction,
+which turned out to be a device-flag ordering bug in the backend, now fixed
+(lablup/mlxcel#1876). `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
diff --git a/docs/installation.md b/docs/installation.md
index edef3a3d6..b55eeb82c 100644
--- a/docs/installation.md
+++ b/docs/installation.md
@@ -481,7 +481,7 @@ thinking model that spends its whole budget inside the thinking block reports
 | 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`. |
-| Audio (speech to text, text to speech) | Works. The FFT primitive runs on hipFFT; `MLX_ROCM_FFT_CACHE_SIZE` bounds how many transform plans stay alive (lablup/mlxcel#1825). |
+| Audio (speech to text, text to speech) | Works. The FFT primitive runs on hipFFT; plans are cached up to `MLX_ROCM_FFT_CACHE_SIZE` (default 128, as on CUDA; lablup/mlxcel#1825, #1876). |
 | Windows, multiple GPUs, distributed inference | Not supported. |
 
 Decode throughput measured on the tested configuration, for orientation only
diff --git a/src/lib/mlx-cpp/patches-rocm/LOCAL_FIXES.md b/src/lib/mlx-cpp/patches-rocm/LOCAL_FIXES.md
index cb24cc56a..844001c20 100644
--- a/src/lib/mlx-cpp/patches-rocm/LOCAL_FIXES.md
+++ b/src/lib/mlx-cpp/patches-rocm/LOCAL_FIXES.md
@@ -26,7 +26,8 @@ The fork branched from upstream on 2026-06-26 (`39886de4`). Moving its ROCm code
 15. **`SearchSorted` implemented.** The fork leaves `SearchSorted` a `NO_GPU` stub, and the exception crosses the cxx bridge into a `noexcept` extern, so it ends in `std::terminate`. mlxcel has no call site of its own; it is reached from MLX internals, because `residual_resample` in speculative decoding calls `fused_sample`, which falls back to `random::categorical` where no Gumbel-max kernel exists, and `categorical_impl` takes its inverse-CDF branch for 1-D logits, which is mlxcel's single-row sampling case. `mlx/backend/rocm/sort.hip` now carries the kernel, ported from `mlx/backend/cuda/sort.cu` and reusing this backend's `LessThan` so NaN ordering matches the CPU. Unlike the CUDA original it is grid-stride: the ROCm launch geometry caps the grid, so one index per thread would silently leave the tail of a large output unwritten. Verified against the CPU stream on both sides for duplicate runs, strictly increasing sequences, out-of-range queries, NaN, a single element, and 17M queries (past the 65535 x 256 grid extent). Applies to the fork; to be proposed there.
 16. **`Hadamard` implemented.** The fork leaves `Hadamard` a `NO_GPU` stub, which aborted every TurboQuant KV cache mode and `tests/wht_op.rs`. `mlx/backend/rocm/hadamard.hip` adds a single-stage radix-2 butterfly in shared memory for `m == 1` and a power-of-two last axis up to 8192, which is the whole of what mlxcel reaches: `wht()` asserts a power-of-two last axis and the KV cache advisor demotes Turbo modes for the models that would not have one, so the sizes seen here are head dims, chiefly 64, 128 and 256. Larger `n` and the `m` in {12, 20, 28} families throw a message naming the limit rather than returning a wrong answer; the CUDA path covers them by generating a radix-m codelet through NVRTC. The scale is the primitive's `scale_`, which MLX defaults to `1/sqrt(n)`, so the transform stays the involution mlxcel relies on; accumulation is in float regardless of element type. Like the CUDA original the grid is clamped at 65535 blocks and the kernel is grid-stride. Measured with `examples/wht_numeric_probe` against the M1 Ultra Metal reference: norm ratio 1.000000 to within 1e-5 at head_dim 64, 128 and 256, fp16 relative RMS about 2.0e-4 against Metal's 3.8e-4 to 5.0e-4, f32 relative RMS about 1e-7 matching Metal. The fp16 figure is lower than Metal's because this kernel keeps the butterfly in float in shared memory and rounds once on the way out, where Metal rounds between stages; more accurate than the reference is still inside the band, but it is worth knowing before someone reads the gap as a measurement error. Applies to the fork; to be proposed there.
 17. **Narrow index dtypes in `Gather` and `Scatter`.** Five dispatch sites in `mlx/backend/rocm/indexing.hip` picked the kernel's index type with `int32_t` for int32 and uint32 and `int64_t` for everything else, so int8, uint8, int16, uint16 and bool index arrays reached a kernel that read eight bytes per element from a one or two byte buffer. On ROCm that is an `HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION`: the GPU queue aborts and the process spins instead of failing. mlxcel reaches it with uint8, because the TurboQuant 4-bit codebook lookup is a `take` with uint8 indices, which is why every TurboQuant KV cache test faulted once the `Hadamard` stub was gone. Narrow index arrays are now widened to int32 once before the dispatch. The alternative, giving each of the eight integer dtypes its own instantiation, was implemented and measured first: it is also correct but takes this translation unit from 26 s to 348 s to compile, because the index type multiplies an existing product of twelve value dtypes and five to six index-count branches, so widening was kept at the cost of one copy pass on the paths that are broken today. Upstream does instantiate per index dtype, in both Metal (`backend/metal/indexing.cpp` puts the index dtype in the kernel name and builds the source from a template) and CUDA (`backend/cuda/indexing.cpp` passes it as an NVRTC template argument), but both compile the kernel at runtime and cache it, so the combinations cost nothing at build time. This backend compiles `indexing.hip` ahead of time, so the same strategy is paid by every build. The asymmetry is the compile model, not the strategy, and it is the argument for proposing widening rather than instantiation to the fork. Upstream already uses the same move for the offset type, choosing int32 or int64 by size rather than instantiating every combination. Verified against the CPU stream for int8, uint8, int16, uint16, int32, uint32, int64 and uint64 at four index shapes. Applies to the fork; to be proposed there.
-18. **`FFT` implemented through hipFFT.** The fork leaves `FFT` a `NO_GPU` stub, which aborted Kokoro text to speech and the Phi-4-multimodal audio front end, the only two mlxcel paths that reach a device FFT (the other audio front ends use the host-side transform in `src/audio/fft.rs`). `mlx/backend/rocm/fft.hip` ports `mlx/backend/cuda/fft.cu`, with two differences. This hipFFT has no `hipfftXtMakePlanMany` or `hipfftXtExec`, so the typed API is used instead: one `hipfftType` per plan and one exec entry point per transform, which is enough because MLX only asks for float32 and complex64. And hipFFT work cannot be a HIP graph kernel node, so execution goes through `CommandEncoder::launch_kernel`, the same path the rocBLAS and hipBLASLt GEMMs take. The plan cache defaults to 8 entries rather than the CUDA file's 128: with more hipFFT plans alive at once, creating the next one blocks forever inside the library with the GPU idle, reproduced deterministically at capacities 16, 24, 32 and 128 and never at 8, by a probe cycling 26 transform shapes. A plan outlives its cache entry while a pending execution still holds it, so the cap has to leave headroom. The root cause inside rocFFT is not identified; `MLX_ROCM_FFT_CACHE_SIZE` raises the cap for anyone measuring it. Verified against the CPU stream for rfft, irfft, round trips and complex fft at lengths 2, 8, 400, 512, 1024, 1200 and 2048 and batches up to 512, all within 1e-5 relative. Applies to the fork; to be proposed there.
+18. **`FFT` implemented through hipFFT.** The fork leaves `FFT` a `NO_GPU` stub, which aborted Kokoro text to speech and the Phi-4-multimodal audio front end, the only two mlxcel paths that reach a device FFT (the other audio front ends use the host-side transform in `src/audio/fft.rs`). `mlx/backend/rocm/fft.hip` ports `mlx/backend/cuda/fft.cu`, with two differences. This hipFFT has no `hipfftXtMakePlanMany` or `hipfftXtExec`, so the typed API is used instead: one `hipfftType` per plan and one exec entry point per transform, which is enough because MLX only asks for float32 and complex64. And hipFFT work cannot be a HIP graph kernel node, so execution goes through `CommandEncoder::launch_kernel`, the same path the rocBLAS and hipBLASLt GEMMs take. The plan cache holds 128 entries, as in the CUDA file (`MLX_ROCM_FFT_CACHE_SIZE` overrides it). It was 8 until lablup/mlxcel#1876: at 16 and above the first plan eviction hung inside `hipfftDestroy`, whose `hipFree` never returned. The cause was not rocFFT or the number of live plans but the order of `hipSetDeviceFlags` against the first HIP queue, fixed in entry 20. Verified against the CPU stream for rfft, irfft, round trips and complex fft at lengths 2, 8, 400, 512, 1024, 1200 and 2048 and batches up to 512, all within 1e-5 relative. Applies to the fork; to be proposed there.
+20. **Blocking-sync flag set before the first HIP queue.** The fork set `hipSetDeviceFlags(hipDeviceScheduleBlockingSync)` in `rocm::device()`, but the allocator runs first: the first unified allocation's host read calls `hipStreamQuery(nullptr)` in `allocator::Buffer::raw_ptr()`, which creates the null stream's queue, and CLR gives that queue non-interrupt completion signals because active wait is still on. After the flag flips, CLR (`HwQueueTracker::ActiveSignal` in `rocvirtual.cpp`) no longer swaps such a signal for an interrupt one when a command needs a completion handler, and ROCr rejects async handlers on non-interrupt signals (`hsa_amd_signal_async_handler() failed to set the handler!` under `AMD_LOG_LEVEL=4`). The marker a device-wide sync enqueues on the null stream then never signals its waiter: `hipFree` hung in `std::condition_variable::wait` below `Device::SyncAllStreams`. rocFFT's plan destroy was the first thing to hit it (entry 18), but any `hipFree` or device-wide sync with pending null-stream work could. `rocm::ensure_device_flags()` in `device.cpp` now sets the flag once per device index, thread-safe and touching only that device, and runs before every path that can create a queue ahead of `rocm::device()`: `unified_malloc`, both branches of `raw_ptr()`, the async-pool free streams in the `RocmAllocator` constructor, `gpu::init()` and `staged_write`'s device sync; `rocm::device()` uses it too. With it, `hipSetDeviceFlags` precedes the first `Created SWq` line in the `AMD_LOG_LEVEL=3` log, the 40-shape plan probe completes at plan-cache capacities 16, 32 and 128, and `tests/rocm_fft_plan_cache.rs` guards it (it times out with the ordering reverted). A standalone HIP reproducer without MLX or hipFFT (null-stream touch, then the flag, then `hipMalloc`, sync `hipMemcpy`, a null-stream kernel and `hipFree` in a loop) hangs at its first `hipFree` and completes without the early touch. Measured on gfx1151 with HIP 7.15.26333 and rocFFT 1.0.39 only. Applies to the fork; to be proposed there, and the reproducer is a candidate for a CLR report.
 
 ## Dropped from the vendored set
 
diff --git a/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/allocator.cpp b/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/allocator.cpp
index 4489295f1..51928b7fd 100644
--- a/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/allocator.cpp
+++ b/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/allocator.cpp
@@ -130,6 +130,9 @@ inline void* unified_malloc(size_t size, bool& is_managed) {
   void* data = nullptr;
   hipError_t err;
   ensure_mlx_device_current();
+  // Often the first HIP work in the process, ahead of rocm::device(): put the
+  // device in blocking-sync mode before any queue exists on it (#1876).
+  ensure_current_device_flags();
 
   if (use_finegrained()) {
     err = hipExtMallocWithFlags(&data, size, hipDeviceMallocFinegrained);
@@ -326,6 +329,9 @@ RocmAllocator::RocmAllocator()
         }
       }
       (void)hipSetDevice(i);
+      // The free stream below is a queue on device i; flag the device first
+      // (#1876). This loop already switches to every pooled device.
+      ensure_device_flags(i);
       hipMemPool_t pool = nullptr;
       if (hipDeviceGetDefaultMemPool(&pool, i) == hipSuccess) {
         mem_pools_[i] = pool;
@@ -925,11 +931,17 @@ void* Buffer::raw_ptr() {
   auto& cbuf = *static_cast<rocm::RocmBuffer*>(ptr_);
 
   if (cbuf.device == -1) {
+    // The null-stream query creates the null stream's queue on first use, and
+    // a host read can come before rocm::device() ever runs. Flag the device
+    // the query is about to touch first, or that queue keeps completion
+    // signals a later hipFree waits on forever (#1876).
+    rocm::ensure_current_device_flags();
     if (hipStreamQuery(nullptr) != hipSuccess) {
       (void)hipStreamSynchronize(nullptr);
     }
   } else {
     // Discrete: host shadow (CUDA: move_to_unified_memory).
+    rocm::ensure_current_device_flags();
     (void)hipDeviceSynchronize();
     rocm::allocator().ensure_host_shadow(cbuf);
     return cbuf.host_shadow;
diff --git a/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/device.cpp b/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/device.cpp
index ea18dbe27..f563c336c 100644
--- a/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/device.cpp
+++ b/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/device.cpp
@@ -1345,19 +1345,73 @@ std::unordered_map<int, Device>& get_devices() {
   return devices;
 }
 
+void ensure_device_flags(int device_index) {
+  if (device_index < 0) {
+    return;
+  }
+  // Lock-free fast path for the common indices: this runs on every unified
+  // allocation and every host read of a unified buffer.
+  static std::atomic<uint64_t> flagged_low{0};
+  const bool low = device_index < 64;
+  const uint64_t bit = low ? (uint64_t{1} << device_index) : 0;
+  if (low && (flagged_low.load(std::memory_order_acquire) & bit)) {
+    return;
+  }
+  static std::mutex mu;
+  static std::vector<int> flagged_high;
+  std::lock_guard<std::mutex> lock(mu);
+  if (low ? (flagged_low.load(std::memory_order_relaxed) & bit) != 0
+          : std::find(flagged_high.begin(), flagged_high.end(), device_index) !=
+              flagged_high.end()) {
+    return;
+  }
+  // Per index, not one process-wide bool: if device 0 were flagged first a
+  // global gate would leave device 1 unflagged. hipSetDeviceFlags applies to
+  // the current device, so switch only when the caller is on another one and
+  // switch back after. Never iterate every device: creating a context or queue
+  // on the other GPU of a multi-GPU host is what wedges the discrete GPU's
+  // queue over a TB5 link, so touch only this device.
+  int prev = -1;
+  const bool have_prev = hipGetDevice(&prev) == hipSuccess;
+  const bool switched = !have_prev || prev != device_index;
+  if (switched) {
+    (void)hipSetDevice(device_index);
+  }
+  if (hipSetDeviceFlags(hipDeviceScheduleBlockingSync) != hipSuccess) {
+    // Not fatal (the device then keeps its default wait mode, as it would
+    // have before this helper existed); clear the error so it does not
+    // surface from an unrelated later call. Recorded as done either way so a
+    // failing device is not retried on every allocation.
+    (void)hipGetLastError();
+  }
+  if (switched && have_prev) {
+    (void)hipSetDevice(prev);
+  }
+  if (low) {
+    flagged_low.fetch_or(bit, std::memory_order_release);
+  } else {
+    flagged_high.push_back(device_index);
+  }
+}
+
+void ensure_current_device_flags() {
+  int current = -1;
+  if (hipGetDevice(&current) == hipSuccess) {
+    ensure_device_flags(current);
+  }
+}
+
 Device& device(mlx::core::Device device) {
   auto& devices = get_devices();
   auto it = devices.find(device.index);
   if (it == devices.end()) {
-    // Set blocking sync flags on THIS device (per index, not a single global
-    // bool: if device 0 were touched first the global gate would leave device 1
-    // unflagged). Must happen while this device is current and before its
-    // context is created — i.e. before the Device is constructed. Iterating
-    // every device would create a context/queue on the other GPU too; on a
-    // multi-GPU host that cross-device coexistence is what wedges the discrete
-    // GPU's queue over a TB5 link, so touch only this device.
-    hipSetDevice(device.index);
-    hipSetDeviceFlags(hipDeviceScheduleBlockingSync);
+    // Bind this device (callers rely on it being current afterwards) and make
+    // sure it is in blocking-sync mode before the Device and its streams are
+    // constructed. The first unified allocation usually got there first, since
+    // the allocator runs before the first stream exists; see
+    // ensure_device_flags() for why the order matters.
+    (void)hipSetDevice(device.index);
+    ensure_device_flags(device.index);
     it = devices.try_emplace(device.index, device.index).first;
   }
   return it->second;
diff --git a/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/device.h b/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/device.h
index 6d0b70dda..e65403fb5 100644
--- a/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/device.h
+++ b/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/device.h
@@ -482,6 +482,20 @@ class Device {
   std::mutex encoders_mtx_;
 };
 
+// Put `device_index` in blocking-sync wait mode (hipDeviceScheduleBlockingSync)
+// once per process, before MLX creates the first HIP queue on it. Thread-safe;
+// touches only that device, and leaves the calling thread's current device as
+// it found it. Call it before any HIP call that can create a queue on a device
+// (a null-stream query, sync, copy or launch, or a stream creation) on a path
+// that can run before rocm::device(): the order is what matters. CLR types a
+// queue's completion signals by the wait mode in force when the queue is
+// created and never re-types them, so a queue created before the flag keeps
+// signals that cannot carry the completion handlers a later hipFree or device
+// sync attaches, and that wait never returns (lablup/mlxcel#1876).
+void ensure_device_flags(int device_index);
+// ensure_device_flags() for the calling thread's current HIP device.
+void ensure_current_device_flags();
+
 Device& device(mlx::core::Device device);
 CommandEncoder& get_command_encoder(Stream s);
 void clear_all_encoders();
diff --git a/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/eval.cpp b/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/eval.cpp
index b47844847..45c2bd3a0 100644
--- a/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/eval.cpp
+++ b/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/eval.cpp
@@ -24,6 +24,8 @@ void init() {
   auto d = mlx::core::default_device();
   if (d.type == mlx::core::Device::gpu) {
     (void)hipSetDevice(d.index);
+    // Before the hipFree below or anything else can create a queue (#1876).
+    rocm::ensure_device_flags(d.index);
   }
   hipFree(nullptr);
 }
diff --git a/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/fft.hip b/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/fft.hip
index 8ed3eee85..6e4b3308d 100644
--- a/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/fft.hip
+++ b/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/fft.hip
@@ -125,19 +125,19 @@ struct OrderedArray {
   std::vector<int> order;
 };
 
-// Deliberately much smaller than the CUDA file's 128. Measured on gfx1151 with
-// ROCm 10.0.0: once enough hipFFT plans are alive at the same time, creating the
-// next one blocks forever inside the library, with the GPU idle and the calling
-// thread asleep. A probe cycling 26 distinct transform shapes reproduces it at
-// capacity 16, 24, 32 and 128, and passes every time at 8. A plan also outlives
-// its cache entry while an execution that captured it is still pending, so the
-// live count can exceed the capacity and the cap has to leave headroom. The
-// root cause inside rocFFT is not identified; this bounds it. Raise it with
-// MLX_ROCM_FFT_CACHE_SIZE when measuring.
+// Same capacity as the CUDA file. This was once 8: past that, the first plan
+// eviction hung inside hipfftDestroy (rocFFT's hipFree never returned). The
+// cause was not rocFFT or the number of live plans but a late
+// hipSetDeviceFlags(hipDeviceScheduleBlockingSync): the allocator had already
+// created the null stream's queue with completion signals that cannot carry
+// the handler hipFree's device-wide sync attaches, so that wait never woke.
+// rocm::ensure_device_flags() now runs before the first queue exists
+// (lablup/mlxcel#1876, LOCAL_FIXES.md entry 18 and the device-flags entry).
+// MLX_ROCM_FFT_CACHE_SIZE overrides the capacity.
 auto& fft_plan_cache() {
   static rocm::LRUBytesKeyCache<std::string, std::shared_ptr<HipFFTPlan>> cache(
       "MLX_ROCM_FFT_CACHE_SIZE",
-      /* default_capacity */ 8);
+      /* default_capacity */ 128);
   return cache;
 }
 
@@ -306,8 +306,14 @@ void execute_fft(
   hipfftHandle handle = plan->handle;
 
   // hipFFT is not a graph kernel node; launch_kernel splits the graph and runs
-  // this on the encoder stream, ordered after what came before. The captured
-  // plan keeps the cache entry alive for the duration of the call.
+  // this on the encoder stream, ordered after what came before. launch_kernel
+  // calls the lambda synchronously and drops it on return, so the captured
+  // plan lives only as long as the enqueue, not the asynchronous execution.
+  // An eviction later is still safe for that execution: hipfftDestroy frees
+  // the plan's device memory through hipFree, which synchronizes the device
+  // first. That argument covers the eager launch; the opt-in
+  // MLX_GRAPH_PREFILL_REPLAY=1 path, which records the call into a graph that
+  // launches later, was not examined.
   encoder.launch_kernel([handle,
                          plan,
                          in_ptr,
diff --git a/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/host_stage.cpp b/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/host_stage.cpp
index f0b555aa9..0606d81cf 100644
--- a/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/host_stage.cpp
+++ b/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/host_stage.cpp
@@ -8,6 +8,7 @@
 #include <stdexcept>
 
 #include "mlx/array.h"
+#include "mlx/backend/rocm/device.h"
 #include "mlx/io/load.h"
 
 namespace mlx::core::rocm {
@@ -80,6 +81,9 @@ bool staged_write(io::Writer& out, const array& a) {
   // stream is drained. Sync the device so the DMA cannot observe a half-written
   // buffer: MLX ROCm streams are hipStreamNonBlocking, so a plain hipMemcpy on
   // the NULL stream does NOT order against them.
+  // A checkpoint save can be the first sync on this thread's device; flag it
+  // before the sync touches the null stream (#1876).
+  rocm::ensure_current_device_flags();
   if (hipDeviceSynchronize() != hipSuccess) {
     (void)hipGetLastError();
     return false;
diff --git a/tests/rocm_fft_plan_cache.rs b/tests/rocm_fft_plan_cache.rs
new file mode 100644
index 000000000..65a11e728
--- /dev/null
+++ b/tests/rocm_fft_plan_cache.rs
@@ -0,0 +1,279 @@
+// Copyright 2025-2026 Lablup Inc.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+//! The ROCm FFT survives hipFFT plan evictions and matches the CPU stream
+//! (issue #1876).
+//!
+//! The hipFFT plan cache used to be capped at 8 because the first eviction
+//! hung: `hipfftDestroy` frees the plan's device memory with `hipFree`, and
+//! that `hipFree` waited forever. The cause was the backend setting
+//! `hipDeviceScheduleBlockingSync` after the allocator had already created the
+//! null stream's HIP queue, which left that queue with completion signals the
+//! runtime cannot attach a handler to (`rocm::ensure_device_flags` in the
+//! overlay's `device.cpp` fixes the order).
+//!
+//! The work runs in a child process (this binary, re-invoked on an ignored
+//! test) with `MLX_ROCM_FFT_CACHE_SIZE=16`, so evictions start after eight
+//! round trips and repeat dozens of times, all through the production
+//! `mlxcel_core` FFT ops. The variable has to be in the environment before the
+//! first FFT, which is why it is set on a fresh process rather than here. The
+//! parent kills a child that does not finish in time, which is what the hang
+//! looks like: with the ordering fix reverted the child never gets past the
+//! first eviction (checked by reverting it).
+//!
+//! Skips on any other backend. Run on a ROCm host with:
+//!
+//! ```sh
+//! cargo test --features rocm --test rocm_fft_plan_cache
+//! ```
+
+#![cfg(feature = "rocm")]
+
+use std::io::Read;
+use std::process::{Command, Stdio};
+use std::thread;
+use std::time::{Duration, Instant};
+
+use mlxcel_core::hardware::{GpuBackendKind, gpu_backend_kind};
+use mlxcel_core::streams::DefaultDeviceGuard;
+use mlxcel_core::{MlxArray, UniquePtr, dtype};
+
+/// Set by the parent on the child it spawns; the child body runs only when it
+/// is present, so `--include-ignored` sweeps do not run it in a shared process.
+const CHILD_ENV: &str = "MLXCEL_ROCM_FFT_CHILD";
+const CHILD_TEST: &str = "child_fft_plan_cache";
+/// Small enough that the first eviction comes early (16 hung at the first
+/// eviction before the fix), and every later transform evicts again.
+const CHILD_CACHE_SIZE: &str = "16";
+/// Measured at about 10 s for the whole child on gfx1151 including rocFFT's
+/// kernel compilation on a cold cache; the pre-fix behaviour is an unbounded
+/// wait, so the margin only has to absorb a slow host.
+const CHILD_BUDGET: Duration = Duration::from_secs(120);
+/// Distinct round-trip lengths: two plans (R2C and C2R) each, so 96 plans
+/// through a 16-entry cache.
+const ROUND_TRIP_SHAPES: i32 = 48;
+const TOLERANCE: f32 = 1e-5;
+/// The lengths mlxcel reaches: 2 and 8 are degenerate, 400 is Whisper's frame
+/// (not a power of two), 512 the Phi-4-multimodal front end, 1200 Kokoro's
+/// STFT, and 1024/2048 bracket them.
+const LENGTHS: &[i32] = &[2, 8, 400, 512, 1024, 1200, 2048];
+const DONE_MARKER: &str = "FFT_PLAN_CACHE done";
+
+fn on_rocm() -> bool {
+    gpu_backend_kind() == GpuBackendKind::Rocm
+}
+
+#[test]
+fn fft_survives_plan_evictions_and_matches_cpu() {
+    if !on_rocm() {
+        eprintln!("skipping: not running on a ROCm device");
+        return;
+    }
+    let exe = std::env::current_exe().expect("path of this test binary");
+    let mut child = Command::new(exe)
+        .args([
+            CHILD_TEST,
+            "--exact",
+            "--ignored",
+            "--nocapture",
+            "--test-threads=1",
+        ])
+        .env(CHILD_ENV, "1")
+        .env("MLX_ROCM_FFT_CACHE_SIZE", CHILD_CACHE_SIZE)
+        .stdout(Stdio::piped())
+        .stderr(Stdio::piped())
+        .spawn()
+        .expect("spawn the child test process");
+
+    // Drain both pipes on their own threads so the child can never block on a
+    // full pipe while the parent waits for it.
+    let mut stdout_pipe = child.stdout.take().expect("child stdout");
+    let mut stderr_pipe = child.stderr.take().expect("child stderr");
+    let stdout_reader = thread::spawn(move || {
+        let mut text = String::new();
+        let _ = stdout_pipe.read_to_string(&mut text);
+        text
+    });
+    let stderr_reader = thread::spawn(move || {
+        let mut text = String::new();
+        let _ = stderr_pipe.read_to_string(&mut text);
+        text
+    });
+
+    let deadline = Instant::now() + CHILD_BUDGET;
+    let status = loop {
+        if let Some(status) = child.try_wait().expect("poll the child") {
+            break Some(status);
+        }
+        if Instant::now() > deadline {
+            let _ = child.kill();
+            let _ = child.wait();
+            break None;
+        }
+        thread::sleep(Duration::from_millis(50));
+    };
+    let stdout = stdout_reader.join().expect("stdout reader thread");
+    let stderr = stderr_reader.join().expect("stderr reader thread");
+
+    let status = status.unwrap_or_else(|| {
+        panic!(
+            "the child did not finish within {CHILD_BUDGET:?}: an FFT plan eviction hung (issue #1876)\n--- child stdout ---\n{stdout}\n--- child stderr ---\n{stderr}"
+        )
+    });
+    assert!(
+        status.success(),
+        "child exited with {status}\n--- child stdout ---\n{stdout}\n--- child stderr ---\n{stderr}"
+    );
+    assert!(
+        stdout.contains(DONE_MARKER),
+        "the child exited without finishing its checks\n--- child stdout ---\n{stdout}\n--- child stderr ---\n{stderr}"
+    );
+}
+
+fn pseudo_random(count: usize, seed: u32) -> Vec<f32> {
+    let mut state = seed;
+    (0..count)
+        .map(|_| {
+            state = state.wrapping_mul(1_664_525).wrapping_add(1_013_904_223);
+            ((state >> 8) as f32 / (1u32 << 24) as f32) * 2.0 - 1.0
+        })
+        .collect()
+}
+
+fn input(batch: i32, n: i32, seed: u32) -> UniquePtr<MlxArray> {
+    let values = pseudo_random((batch * n) as usize, seed);
+    mlxcel_core::from_slice_f32(&values, &[batch, n])
+}
+
+fn max_abs(a: &MlxArray) -> f32 {
+    let f = mlxcel_core::astype(a, dtype::FLOAT32);
+    let m = mlxcel_core::max_all(&mlxcel_core::abs(&f));
+    mlxcel_core::eval(&m);
+    mlxcel_core::item_f32(&m)
+}
+
+fn max_abs_diff(a: &MlxArray, b: &MlxArray) -> f32 {
+    max_abs(&mlxcel_core::subtract(
+        &mlxcel_core::astype(a, dtype::FLOAT32),
+        &mlxcel_core::astype(b, dtype::FLOAT32),
+    ))
+}
+
+fn real_rel(got: &MlxArray, reference: &MlxArray) -> f32 {
+    let scale = max_abs(reference);
+    max_abs_diff(got, reference) / if scale > 0.0 { scale } else { 1.0 }
+}
+
+/// Real and imaginary differences summed, scaled by the reference magnitude.
+fn complex_rel(got: &MlxArray, reference: &MlxArray) -> f32 {
+    let re = |a: &MlxArray| mlxcel_core::real_part(a);
+    let im = |a: &MlxArray| mlxcel_core::imag_part(a);
+    let scale = max_abs(&re(reference)) + max_abs(&im(reference));
+    let diff = max_abs_diff(&re(got), &re(reference)) + max_abs_diff(&im(got), &im(reference));
+    diff / if scale > 0.0 { scale } else { 1.0 }
+}
+
+/// Evaluate on the GPU with the default device pinned, before the guard drops.
+fn on_gpu<F: FnOnce() -> UniquePtr<MlxArray>>(f: F) -> UniquePtr<MlxArray> {
+    let _guard = DefaultDeviceGuard::gpu();
+    let out = f();
+    mlxcel_core::eval(&out);
+    out
+}
+
+fn on_cpu<F: FnOnce() -> UniquePtr<MlxArray>>(f: F) -> UniquePtr<MlxArray> {
+    let _guard = DefaultDeviceGuard::cpu();
+    let out = f();
+    mlxcel_core::eval(&out);
+    out
+}
+
+/// The body of `fft_survives_plan_evictions_and_matches_cpu`, run in the child
+/// process it spawns with a 16-entry plan cache.
+#[test]
+#[ignore = "spawned by fft_survives_plan_evictions_and_matches_cpu with a small plan cache"]
+fn child_fft_plan_cache() {
+    if std::env::var_os(CHILD_ENV).is_none() {
+        eprintln!("skipping: run through fft_survives_plan_evictions_and_matches_cpu");
+        return;
+    }
+    if !on_rocm() {
+        eprintln!("skipping: not running on a ROCm device");
+        return;
+    }
+
+    let mut failures = Vec::new();
+    let mut check = |case: &str, batch: i32, n: i32, rel: f32| {
+        println!("{case:<10} batch={batch:<4} n={n:<5} rel={rel:.3e}");
+        // Written so a NaN fails too.
+        let within = rel < TOLERANCE;
+        if !within {
+            failures.push(format!("{case} batch={batch} n={n}: rel {rel:.3e}"));
+        }
+    };
+
+    // Eviction stress first: every length is a new pair of plans, so from the
+    // ninth length on, each transform evicts and destroys an older plan.
+    for i in 0..ROUND_TRIP_SHAPES {
+        let n = 64 + i * 8;
+        let x = input(4, n, 0x5EED ^ n as u32);
+        let back = on_gpu(|| {
+            let spec = mlxcel_core::rfft(&x, n, 1);
+            mlxcel_core::irfft(&spec, n, 1)
+        });
+        check("stress", 4, n, real_rel(&back, &x));
+    }
+
+    // GPU against the CPU stream, still through the small cache.
+    for &n in LENGTHS {
+        let x = input(4, n, 0xC0FFEE ^ n as u32);
+        let gpu = on_gpu(|| mlxcel_core::rfft(&x, n, 1));
+        let cpu = on_cpu(|| mlxcel_core::rfft(&x, n, 1));
+        check("rfft", 4, n, complex_rel(&gpu, &cpu));
+    }
+    for &n in LENGTHS {
+        let x = input(4, n, 0xBEEF00 ^ n as u32);
+        // The spectrum comes from the CPU so both inverses see the same input.
+        let spec = on_cpu(|| mlxcel_core::rfft(&x, n, 1));
+        let gpu = on_gpu(|| mlxcel_core::irfft(&spec, n, 1));
+        let cpu = on_cpu(|| mlxcel_core::irfft(&spec, n, 1));
+        check("irfft", 4, n, real_rel(&gpu, &cpu));
+    }
+    for &n in LENGTHS {
+        let x = input(3, n, 0x1234 ^ n as u32);
+        let back = on_gpu(|| {
+            let spec = mlxcel_core::rfft(&x, n, 1);
+            mlxcel_core::irfft(&spec, n, 1)
+        });
+        check("roundtrip", 3, n, real_rel(&back, &x));
+    }
+    for &n in LENGTHS {
+        let x = mlxcel_core::astype(&input(4, n, 0x77 ^ n as u32), dtype::COMPLEX64);
+        let gpu = on_gpu(|| mlxcel_core::fft(&x, n, 1));
+        let cpu = on_cpu(|| mlxcel_core::fft(&x, n, 1));
+        check("fft", 4, n, complex_rel(&gpu, &cpu));
+    }
+    // A batch large enough that the launch geometry is not the trivial one.
+    let x = input(512, 512, 0xABCD);
+    let gpu = on_gpu(|| mlxcel_core::rfft(&x, 512, 1));
+    let cpu = on_cpu(|| mlxcel_core::rfft(&x, 512, 1));
+    check("rfft", 512, 512, complex_rel(&gpu, &cpu));
+
+    assert!(
+        failures.is_empty(),
+        "FFT results outside {TOLERANCE:e} relative of the reference:\n{}",
+        failures.join("\n")
+    );
+    println!("{DONE_MARKER}");
+}

From 6e4ac7d62259d37daf7c4d594ea096e348f13041 Mon Sep 17 00:00:00 2001
From: Jeongkyu Shin <jshin@lablup.com>
Date: Wed, 30 Sep 2026 01:01:48 +0900
Subject: [PATCH 2/4] docs(rocm): name the LOCAL_FIXES entries in the FFT cache
 comment

Entry 20 now has its number after merging main, which added entry 19 in PR #2046.
---
 src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/fft.hip | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/fft.hip b/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/fft.hip
index 6e4b3308d..c73c5814f 100644
--- a/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/fft.hip
+++ b/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/fft.hip
@@ -132,7 +132,7 @@ struct OrderedArray {
 // created the null stream's queue with completion signals that cannot carry
 // the handler hipFree's device-wide sync attaches, so that wait never woke.
 // rocm::ensure_device_flags() now runs before the first queue exists
-// (lablup/mlxcel#1876, LOCAL_FIXES.md entry 18 and the device-flags entry).
+// (lablup/mlxcel#1876, LOCAL_FIXES.md entries 18 and 20).
 // MLX_ROCM_FFT_CACHE_SIZE overrides the capacity.
 auto& fft_plan_cache() {
   static rocm::LRUBytesKeyCache<std::string, std::shared_ptr<HipFFTPlan>> cache(

From 9a9c1b0616ffeb5a25cd278852703edde928aae5 Mon Sep 17 00:00:00 2001
From: Jeongkyu Shin <jshin@lablup.com>
Date: Wed, 30 Sep 2026 01:16:03 +0900
Subject: [PATCH 3/4] fix(rocm): report device flag failures and clean up child
 on poll error

ensure_device_flags now returns without recording the device as flagged when hipSetDevice fails (the error is cleared so a later call retries), and prints a one-time stderr line naming the device index and HIP error string when hipSetDeviceFlags fails, so a regression of #1876 is visible. It still does not retry the flags call.

tests/rocm_fft_plan_cache.rs now kills and waits the child before panicking when try_wait returns an error. LOCAL_FIXES.md entry 20 describes the new failure handling.

Validation:
- cargo test --features rocm --test rocm_fft_plan_cache: passed
- cargo clippy --features rocm --test rocm_fft_plan_cache -- -D warnings: clean
- cargo fmt --all -- --check and make verify-fmt: clean

Refs #1876
---
 src/lib/mlx-cpp/patches-rocm/LOCAL_FIXES.md   |  2 +-
 .../patches-rocm/mlx/backend/rocm/device.cpp  | 21 +++++++++++++++----
 tests/rocm_fft_plan_cache.rs                  | 10 +++++++--
 3 files changed, 26 insertions(+), 7 deletions(-)

diff --git a/src/lib/mlx-cpp/patches-rocm/LOCAL_FIXES.md b/src/lib/mlx-cpp/patches-rocm/LOCAL_FIXES.md
index 5c13a149b..7e8b58f2f 100644
--- a/src/lib/mlx-cpp/patches-rocm/LOCAL_FIXES.md
+++ b/src/lib/mlx-cpp/patches-rocm/LOCAL_FIXES.md
@@ -28,7 +28,7 @@ The fork branched from upstream on 2026-06-26 (`39886de4`). Moving its ROCm code
 17. **Narrow index dtypes in `Gather` and `Scatter`.** Five dispatch sites in `mlx/backend/rocm/indexing.hip` picked the kernel's index type with `int32_t` for int32 and uint32 and `int64_t` for everything else, so int8, uint8, int16, uint16 and bool index arrays reached a kernel that read eight bytes per element from a one or two byte buffer. On ROCm that is an `HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION`: the GPU queue aborts and the process spins instead of failing. mlxcel reaches it with uint8, because the TurboQuant 4-bit codebook lookup is a `take` with uint8 indices, which is why every TurboQuant KV cache test faulted once the `Hadamard` stub was gone. Narrow index arrays are now widened to int32 once before the dispatch. The alternative, giving each of the eight integer dtypes its own instantiation, was implemented and measured first: it is also correct but takes this translation unit from 26 s to 348 s to compile, because the index type multiplies an existing product of twelve value dtypes and five to six index-count branches, so widening was kept at the cost of one copy pass on the paths that are broken today. Upstream does instantiate per index dtype, in both Metal (`backend/metal/indexing.cpp` puts the index dtype in the kernel name and builds the source from a template) and CUDA (`backend/cuda/indexing.cpp` passes it as an NVRTC template argument), but both compile the kernel at runtime and cache it, so the combinations cost nothing at build time. This backend compiles `indexing.hip` ahead of time, so the same strategy is paid by every build. The asymmetry is the compile model, not the strategy, and it is the argument for proposing widening rather than instantiation to the fork. Upstream already uses the same move for the offset type, choosing int32 or int64 by size rather than instantiating every combination. Verified against the CPU stream for int8, uint8, int16, uint16, int32, uint32, int64 and uint64 at four index shapes. Applies to the fork; to be proposed there.
 18. **`FFT` implemented through hipFFT.** The fork leaves `FFT` a `NO_GPU` stub, which aborted Kokoro text to speech and the Phi-4-multimodal audio front end, the only two mlxcel paths that reach a device FFT (the other audio front ends use the host-side transform in `src/audio/fft.rs`). `mlx/backend/rocm/fft.hip` ports `mlx/backend/cuda/fft.cu`, with two differences. This hipFFT has no `hipfftXtMakePlanMany` or `hipfftXtExec`, so the typed API is used instead: one `hipfftType` per plan and one exec entry point per transform, which is enough because MLX only asks for float32 and complex64. And hipFFT work cannot be a HIP graph kernel node, so execution goes through `CommandEncoder::launch_kernel`, the same path the rocBLAS and hipBLASLt GEMMs take. The plan cache holds 128 entries, as in the CUDA file (`MLX_ROCM_FFT_CACHE_SIZE` overrides it). It was 8 until lablup/mlxcel#1876: at 16 and above the first plan eviction hung inside `hipfftDestroy`, whose `hipFree` never returned. The cause was not rocFFT or the number of live plans but the order of `hipSetDeviceFlags` against the first HIP queue, fixed in entry 20. Verified against the CPU stream for rfft, irfft, round trips and complex fft at lengths 2, 8, 400, 512, 1024, 1200 and 2048 and batches up to 512, all within 1e-5 relative. Applies to the fork; to be proposed there.
 19. **`get_launch_args` removed.** The fork's `mlx/backend/rocm/kernel_utils.hpp` defined `get_launch_args` (a sized overload and an `array` overload) that fixed the block at 256 threads, ignored `shape`, `strides` and `large`, and capped the grid at 65535 blocks, returning that grid with nothing in its name or comment saying the kernel had to be grid-stride. A kernel written one index per thread with an early return past the end would leave every element past 65535 x 256 unwritten, with the output keeping whatever the allocation held (the failure shape of item 13). Upstream CUDA's helper of the same name (`mlx/backend/cuda/kernel_utils.cu` at the pin) does not cap x, takes a `max_block_dim`, and honours `large` through `get_2d_grid_dims`, so a kernel ported from the CUDA backend with its call intact would compile here and inherit the truncation. Nothing called it, in this overlay or in the fork at `75915908`; launch sites compute their own geometry. Most sites that clamp at 65535 blocks launch a grid-stride kernel (`binary.hip`, `unary.hip`, `hadamard.hip`, `sort.hip` among them), but not all: `SliceUpdate::eval_gpu`'s reduce-op path in `indexing.hip` clamps and launches `slice_update_op_kernel`, which does at most `NWORK` elements per thread and never reads `gridDim`, so an update larger than 65535 x 256 x `NWORK` elements leaves its tail unapplied. That is the same defect at a real call site, left for a separate fix because this item changes no kernel's geometry. Of the three options weighed in lablup/mlxcel#1874 (delete it, rename it to state the grid-stride requirement, or spill the overflow into another grid dimension the way upstream's `get_launch_args_general` spills y into z), deletion was taken: it cannot regress, and the other two keep a helper that only looks like upstream's. Spilling into y would not even make a one-index-per-thread kernel correct unless that kernel also reads `blockIdx.y`, and AMD caps each grid dimension at 2^32 - 1 threads (item 11), so a faithful port of upstream's geometry is not a drop-in either. In place of the definitions the header now declares `template <typename... Args> void get_launch_args(Args&&...) = delete;` with a comment giving this reason, so a ported call fails to compile at the explanation rather than as an undeclared name someone might "fix" by restoring the old helper. A fork sync that brings the old definitions back must drop them again. Applies to the fork; to be proposed there.
-20. **Blocking-sync flag set before the first HIP queue.** The fork set `hipSetDeviceFlags(hipDeviceScheduleBlockingSync)` in `rocm::device()`, but the allocator runs first: the first unified allocation's host read calls `hipStreamQuery(nullptr)` in `allocator::Buffer::raw_ptr()`, which creates the null stream's queue, and CLR gives that queue non-interrupt completion signals because active wait is still on. After the flag flips, CLR (`HwQueueTracker::ActiveSignal` in `rocvirtual.cpp`) no longer swaps such a signal for an interrupt one when a command needs a completion handler, and ROCr rejects async handlers on non-interrupt signals (`hsa_amd_signal_async_handler() failed to set the handler!` under `AMD_LOG_LEVEL=4`). The marker a device-wide sync enqueues on the null stream then never signals its waiter: `hipFree` hung in `std::condition_variable::wait` below `Device::SyncAllStreams`. rocFFT's plan destroy was the first thing to hit it (entry 18), but any `hipFree` or device-wide sync with pending null-stream work could. `rocm::ensure_device_flags()` in `device.cpp` now sets the flag once per device index, thread-safe and touching only that device, and runs before every path that can create a queue ahead of `rocm::device()`: `unified_malloc`, both branches of `raw_ptr()`, the async-pool free streams in the `RocmAllocator` constructor, `gpu::init()` and `staged_write`'s device sync; `rocm::device()` uses it too. With it, `hipSetDeviceFlags` precedes the first `Created SWq` line in the `AMD_LOG_LEVEL=3` log, the 40-shape plan probe completes at plan-cache capacities 16, 32 and 128, and `tests/rocm_fft_plan_cache.rs` guards it (it times out with the ordering reverted). A standalone HIP reproducer without MLX or hipFFT (null-stream touch, then the flag, then `hipMalloc`, sync `hipMemcpy`, a null-stream kernel and `hipFree` in a loop) hangs at its first `hipFree` and completes without the early touch. Measured on gfx1151 with HIP 7.15.26333 and rocFFT 1.0.39 only. Applies to the fork; to be proposed there, and the reproducer is a candidate for a CLR report.
+20. **Blocking-sync flag set before the first HIP queue.** The fork set `hipSetDeviceFlags(hipDeviceScheduleBlockingSync)` in `rocm::device()`, but the allocator runs first: the first unified allocation's host read calls `hipStreamQuery(nullptr)` in `allocator::Buffer::raw_ptr()`, which creates the null stream's queue, and CLR gives that queue non-interrupt completion signals because active wait is still on. After the flag flips, CLR (`HwQueueTracker::ActiveSignal` in `rocvirtual.cpp`) no longer swaps such a signal for an interrupt one when a command needs a completion handler, and ROCr rejects async handlers on non-interrupt signals (`hsa_amd_signal_async_handler() failed to set the handler!` under `AMD_LOG_LEVEL=4`). The marker a device-wide sync enqueues on the null stream then never signals its waiter: `hipFree` hung in `std::condition_variable::wait` below `Device::SyncAllStreams`. rocFFT's plan destroy was the first thing to hit it (entry 18), but any `hipFree` or device-wide sync with pending null-stream work could. `rocm::ensure_device_flags()` in `device.cpp` now sets the flag once per device index, thread-safe and touching only that device, and runs before every path that can create a queue ahead of `rocm::device()`: `unified_malloc`, both branches of `raw_ptr()`, the async-pool free streams in the `RocmAllocator` constructor, `gpu::init()` and `staged_write`'s device sync; `rocm::device()` uses it too. If `hipSetDevice` fails it returns without recording the device, so a later call retries; if `hipSetDeviceFlags` fails it prints one line to stderr naming the device index and the HIP error and does not retry, so a regression of #1876 is visible. With it, `hipSetDeviceFlags` precedes the first `Created SWq` line in the `AMD_LOG_LEVEL=3` log, the 40-shape plan probe completes at plan-cache capacities 16, 32 and 128, and `tests/rocm_fft_plan_cache.rs` guards it (it times out with the ordering reverted). A standalone HIP reproducer without MLX or hipFFT (null-stream touch, then the flag, then `hipMalloc`, sync `hipMemcpy`, a null-stream kernel and `hipFree` in a loop) hangs at its first `hipFree` and completes without the early touch. Measured on gfx1151 with HIP 7.15.26333 and rocFFT 1.0.39 only. Applies to the fork; to be proposed there, and the reproducer is a candidate for a CLR report.
 
 ## Dropped from the vendored set
 
diff --git a/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/device.cpp b/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/device.cpp
index f563c336c..2aae2c09d 100644
--- a/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/device.cpp
+++ b/src/lib/mlx-cpp/patches-rocm/mlx/backend/rocm/device.cpp
@@ -1374,15 +1374,28 @@ void ensure_device_flags(int device_index) {
   int prev = -1;
   const bool have_prev = hipGetDevice(&prev) == hipSuccess;
   const bool switched = !have_prev || prev != device_index;
-  if (switched) {
-    (void)hipSetDevice(device_index);
+  if (switched && hipSetDevice(device_index) != hipSuccess) {
+    // Could not bind the device, so the flags cannot be applied to it. Clear
+    // the error and do not record the device as flagged, so a later call can
+    // try again.
+    (void)hipGetLastError();
+    return;
   }
-  if (hipSetDeviceFlags(hipDeviceScheduleBlockingSync) != hipSuccess) {
+  const hipError_t flags_err = hipSetDeviceFlags(hipDeviceScheduleBlockingSync);
+  if (flags_err != hipSuccess) {
     // Not fatal (the device then keeps its default wait mode, as it would
     // have before this helper existed); clear the error so it does not
     // surface from an unrelated later call. Recorded as done either way so a
-    // failing device is not retried on every allocation.
+    // failing device is not retried on every allocation. Say so once on
+    // stderr: without blocking-sync a regression of #1876 would otherwise be
+    // silent.
     (void)hipGetLastError();
+    std::fprintf(
+        stderr,
+        "[mlxcel rocm] hipSetDeviceFlags(hipDeviceScheduleBlockingSync) failed "
+        "for device %d: %s\n",
+        device_index,
+        hipGetErrorString(flags_err));
   }
   if (switched && have_prev) {
     (void)hipSetDevice(prev);
diff --git a/tests/rocm_fft_plan_cache.rs b/tests/rocm_fft_plan_cache.rs
index 65a11e728..fc232f765 100644
--- a/tests/rocm_fft_plan_cache.rs
+++ b/tests/rocm_fft_plan_cache.rs
@@ -113,8 +113,14 @@ fn fft_survives_plan_evictions_and_matches_cpu() {
 
     let deadline = Instant::now() + CHILD_BUDGET;
     let status = loop {
-        if let Some(status) = child.try_wait().expect("poll the child") {
-            break Some(status);
+        match child.try_wait() {
+            Ok(Some(status)) => break Some(status),
+            Ok(None) => {}
+            Err(err) => {
+                let _ = child.kill();
+                let _ = child.wait();
+                panic!("poll the child: {err}");
+            }
         }
         if Instant::now() > deadline {
             let _ = child.kill();

From 1e37dabe4d15cc5824cd95e03f624510d686d191 Mon Sep 17 00:00:00 2001
From: Jeongkyu Shin <jshin@lablup.com>
Date: Wed, 30 Sep 2026 01:57:26 +0900
Subject: [PATCH 4/4] docs: add technical report for PR #2049

---
 ...m-blocking-sync-first-queue-20260930.en.md | 142 ++++++++++++++++++
 ...m-blocking-sync-first-queue-20260930.ko.md | 142 ++++++++++++++++++
 2 files changed, 284 insertions(+)
 create mode 100644 TECHNICAL_REPORTS/2049-rocm-blocking-sync-first-queue-20260930.en.md
 create mode 100644 TECHNICAL_REPORTS/2049-rocm-blocking-sync-first-queue-20260930.ko.md

diff --git a/TECHNICAL_REPORTS/2049-rocm-blocking-sync-first-queue-20260930.en.md b/TECHNICAL_REPORTS/2049-rocm-blocking-sync-first-queue-20260930.en.md
new file mode 100644
index 000000000..f69946821
--- /dev/null
+++ b/TECHNICAL_REPORTS/2049-rocm-blocking-sync-first-queue-20260930.en.md
@@ -0,0 +1,142 @@
+# Technical Report: PR #2049 - Set ROCm Blocking-Sync Before the First HIP Queue, Restore the FFT Plan Cache to 128
+
+**Date**: 2026-09-30
+
+**Status**: Implemented and validated on the gfx1151 host; pending merge.
+
+**Languages**: C++ (HIP runtime calls in the ROCm overlay), Rust (integration test), Markdown
+
+**Risk Level**: Medium (changes when every ROCm process sets its device wait mode, on the allocator's hot path; ROCm-only overlay plus a `#![cfg(feature = "rocm")]` test, Metal and CUDA never copy `patches-rocm/`)
+
+## Executive Summary
+
+Issue #1876 (part of epic #1801) asked why the ROCm hipFFT plan cache had to be capped at 8 entries when the CUDA file uses 128. Above 8, an FFT-heavy workload hung forever with the GPU idle, and the cause had been written down as "somewhere inside rocFFT". The investigation in this PR shows the hang was not in rocFFT and did not depend on the number of live plans. It was an ordering bug in the ROCm backend: `hipSetDeviceFlags(hipDeviceScheduleBlockingSync)` ran in `rocm::device()`, after the allocator had already created the null stream's HIP queue through `hipStreamQuery(nullptr)`. CLR had given that queue non-interrupt completion signals, which cannot carry the async handler that a later `hipFree` relies on, so the first `hipFree` that had to synchronize the null stream never returned. rocFFT's `hipfftDestroy` at the first plan eviction was simply the first code in mlxcel that did null-stream work followed by `hipFree`.
+
+A 30-line standalone HIP program with no MLX and no hipFFT reproduces the hang: touch the null stream, set the flag, then loop over `hipMalloc`, `hipMemcpy`, a null-stream kernel and `hipFree`. It hangs at the first `hipFree`, and completes when the early touch is removed.
+
+The fix adds `rocm::ensure_device_flags(int)` and `ensure_current_device_flags()`, which set the flag once per device before any queue exists on it, and calls them from every path that can reach HIP before `rocm::device()`. The plan cache goes back to 128, and a new integration test, `tests/rocm_fft_plan_cache.rs`, drives 48 evicting round trips at cache size 16 in a child process and fails if the hang returns. The defect was never specific to FFT: any device-wide sync with pending null-stream work was exposed, so the fix is wider than the issue that found it.
+
+## 1. Problem Statement
+
+### What the issue recorded
+
+The FFT port from #1825 (PRs #1856 and #1861) set `default_capacity` to 8 on the hipFFT `LRUBytesKeyCache` in `fft.hip`, with a comment saying the root cause inside rocFFT was unknown. The measurements were consistent and misleading: a probe cycling 26 transform shapes completed at capacity 8 and hung at 16, 24, 32 and 128, always at the same position in the sequence. Reordering the cases moved the hang to whatever ran at that position, and 60 repeats of a single shape never hung. That pattern pointed at "too many live plans", and the issue's proposed hypotheses followed from it: rocFFT runtime compilation, the manual work area set by `hipfftSetAutoAllocation(handle, 0)`, and plans outliving their cache entries through `shared_ptr` captures. A standalone program had also created 64 hipFFT plans without trouble, so the limit was known to depend on how MLX used the library, but not how.
+
+The cost of the cap was real. Kokoro text to speech and the Phi-4-multimodal audio front end, the two mlxcel paths that reach a device FFT, rebuilt plans far more often than on CUDA.
+
+### What the investigation found
+
+This is the core of the PR, and the reasoning is worth keeping because the symptom pointed away from the cause at every step.
+
+**Step 1: locate the hang in time.** The probe was rebuilt as `fft_numeric_probe --plans N` (originally from PR #1878). At cap 16 it hung after `plan 7`; at cap 32, after `plan 15`. Each probe iteration creates two plans (R2C and C2R), so the iteration after plan 7 creates the 17th plan and the iteration after plan 15 creates the 33rd: in both cases the first plan that forces an eviction. The hang was not at a live-plan threshold but at the first `hipfftDestroy`. Because the LRU cache evicts while inserting the new entry, the stall looked like "creating the next plan blocks", which is how the issue described it.
+
+**Step 2: locate the hang in the stack.** rocgdb on the blocked thread gave `hipfftDestroy -> rocfft_plan_destroy -> ~ExecPlan -> ~TreeNode -> gpubuf_t::free -> hipFree -> ... -> std::condition_variable::wait`. The GPU was idle and no rocFFT RTC helper process was running, which ruled out the runtime-compilation hypothesis. rocFFT was only freeing a device buffer; the wait was inside HIP's `hipFree`, which synchronizes all streams on the device (`Device::SyncAllStreams`) before freeing.
+
+**Step 3: find why the wait never wakes.** With `AMD_LOG_LEVEL=4`, the last line before the hang was `rocvirtual.cpp:844 hsa_amd_signal_async_handler() failed to set the handler!`. ROCr refuses async handlers on signals that have no event mailbox, that is, non-interrupt signals. `SyncAllStreams` enqueues a marker on the null stream and waits for the marker's completion handler to fire. The handler was never installed, so the condition variable was never notified.
+
+**Step 4: find where the non-interrupt signal came from.** CLR types a queue's ring of completion signals according to the device wait mode in force when the queue is created. In the default active-wait mode, the host spins on signals, so CLR allocates cheap `HSA_AMD_SIGNAL_AMD_GPU_ONLY` (non-interrupt) signals. While active wait is on, `HwQueueTracker::ActiveSignal` swaps in an interrupt signal whenever a command needs a handler. Once `hipDeviceScheduleBlockingSync` turns active wait off, that swap stops, on the assumption that the queue was created with interrupt signals in the first place. The HIP log showed which queue violated that assumption: the null stream's queue was created by `hipStreamQuery(nullptr)` in `allocator::Buffer::raw_ptr()`, on the first host read of a unified buffer, and only afterwards did `rocm::device()` call `hipSetDeviceFlags`. The null stream's queue therefore kept non-interrupt signals for the life of the process, and the first marker that landed on one of them and needed a handler hung its waiter.
+
+**Step 5: prove it outside MLX.** The standalone reproducer in the PR body strips everything else away (built with `hipcc --offload-arch=gfx1151 -O2`):
+
+```cpp
+if (early_touch) (void)hipStreamQuery(nullptr);        // creates the null-stream queue now
+(void)hipSetDeviceFlags(hipDeviceScheduleBlockingSync); // flips active wait off
+// loop: hipMalloc, hipMemcpy H2D, kernel on the null stream,
+//       keep 8 buffers live, hipFree the oldest
+```
+
+`timeout 30 ./repro 1 200` prints `iter 7` and exits 124, hanging at its first `hipFree`. `./repro 0 200`, which only drops the early touch, prints `done`. No hipFFT, no MLX, no plan cache: the whole defect is the order of two HIP calls.
+
+### Why the issue's hypotheses were wrong, and what that rules out
+
+- **Live plan count.** The hang depends on the first eviction, not on how many plans are alive. The standalone reproducer has no plans at all.
+- **rocFFT runtime compilation.** No RTC helper was running, and the stack was in `hipFree`, not in a compile or cache lock.
+- **Plans outliving their cache entries.** The `execute_fft` comment claimed the captured `shared_ptr` kept the plan alive for the asynchronous execution. It does not: `CommandEncoder::launch_kernel` calls the lambda synchronously and drops it on return. That premise was the reason the cap "had to leave headroom", and it was false.
+- **rocFFT itself.** rocFFT's only role was to be the first caller of `hipFree` after null-stream work. Nothing needs to be reported to rocFFT. The reproducer is a candidate for a CLR report instead (not filed yet).
+
+One part is not explained: why capacities of 12 or below completed while 16 and above hung at the first eviction, given that evictions also happen at capacity 8 with 26 shapes. The PR's working guess is the position in the null stream's signal ring at the time of the first `hipFree` (the reproducer also first hangs after eight kernels), but this was not characterized.
+
+## 2. Change Summary
+
+- **`device.cpp` / `device.h`: `ensure_device_flags(int)` and `ensure_current_device_flags()`.** Sets `hipDeviceScheduleBlockingSync` once per device index. A lock-free atomic bitmask covers indices below 64, because this runs on every unified allocation and every host read of a unified buffer; a mutex and a vector cover anything above. The helper switches to the target device only when the caller is on another one and restores the caller's device afterwards. If `hipSetDevice` fails, the device is not recorded, so a later call retries. If `hipSetDeviceFlags` fails, the error is cleared, one stderr line names the device and the HIP error, and the device is recorded so the failure is not retried on every allocation. `rocm::device()` now calls the helper instead of setting the flag inline. The flag value is unchanged.
+- **Call sites that can create a queue before `rocm::device()`.** `unified_malloc` in `allocator.cpp` (often the first HIP work in the process); both branches of `Buffer::raw_ptr()` (the null-stream query on unified memory and the `hipDeviceSynchronize` on discrete memory); the opt-in async-pool free streams in the `RocmAllocator` constructor, per pooled device; `gpu::init()` in `eval.cpp` before its `hipFree(nullptr)`; and `staged_write` in `host_stage.cpp` before its device sync, because a checkpoint save can be the first sync on a thread's device.
+- **`fft.hip`.** `default_capacity` returns to 128, matching `mlx/backend/cuda/fft.cu`. The cache comment now gives the real cause and points to `LOCAL_FIXES.md` entries 18 and 20. The `execute_fft` comment is corrected: the captured plan lives only as long as the enqueue, eviction is still safe for the eager launch because `hipFree` synchronizes the device first, and the opt-in `MLX_GRAPH_PREFILL_REPLAY=1` graph path was not examined.
+- **`tests/rocm_fft_plan_cache.rs` (new, 285 lines).** The parent test re-invokes this binary on an ignored child test with `MLX_ROCM_FFT_CACHE_SIZE=16` and a 120 s budget, and kills the child if it does not finish. The child runs 48 rfft/irfft round trips at distinct lengths (96 plans through a 16-entry cache, so evictions start after eight round trips and repeat), then compares GPU against the CPU stream within 1e-5 for rfft, irfft, round trip and complex fft at lengths 2, 8, 400, 512, 1024, 1200 and 2048, plus a 512-batch rfft. The child body runs only when `MLXCEL_ROCM_FFT_CHILD` is set, so `--include-ignored` sweeps do not run it in a shared process.
+- **`LOCAL_FIXES.md`.** Entry 18 (the FFT port) now describes the 128 capacity and the history. New entry 20 records the ordering defect, the CLR mechanism, every call site, the failure handling, the evidence and the reproducer, as a fork-side upstreaming candidate.
+- **`docs/environment-variables.md`, `docs/installation.md`.** `MLX_ROCM_FFT_CACHE_SIZE` is documented with default 128 and the hang is described as fixed (#1876).
+
+Commit history: `84e8ad19` is the fix; `6e4ac7d6` names the `LOCAL_FIXES.md` entries in the `fft.hip` comment; `9a9c1b06` makes the helper's failure paths visible and makes the test kill its child if polling it fails; `78a53b15` and `3ccab4e5` merge origin/main (the first to pick up `LOCAL_FIXES.md` entry 19 from PR #2046).
+
+## 3. Technical Decisions
+
+### Fix the order, not the cap
+
+Keeping the cap at 8 would have left the defect in place for every other `hipFree` or device-wide sync that follows null-stream work. FFT was only the first caller to trip it; a future op that frees device memory at runtime could have hung the same way with no FFT involved. Fixing the order removes the class of hangs, and the cache capacity goes back to a memory-versus-rebuild tradeoff rather than a safety limit.
+
+### Flag early at every entry point instead of once at process start
+
+HIP has no hook that runs before the first HIP call made by any code, and the device index is not known until MLX picks a device. The flag has to be set on the right device before the first queue on it, so the helper runs at each place that can create a queue ahead of `rocm::device()`. The per-call cost after the first time is one atomic load and a bit test, which is why the fast path is lock-free: it sits on the allocation path.
+
+### Keep blocking-sync rather than dropping it
+
+Removing `hipSetDeviceFlags` altogether would also make the mode consistent (active wait everywhere), but it would change host CPU usage during every GPU wait, which is the reason the fork sets blocking-sync in the first place. The PR keeps the flag value and changes only when it is applied.
+
+### Per-device, current-device-only
+
+The helper tracks each index separately, so flagging device 0 first cannot leave device 1 unflagged, and it never iterates every device. The existing `rocm::device()` comment records why: on a multi-GPU host, creating a context or queue on the other GPU is what wedged the discrete GPU's queue over a TB5 link. The helper preserves that constraint and restores the caller's current device so it can be called from the allocator without side effects.
+
+### Make failure visible, retry only what can succeed
+
+A `hipSetDevice` failure means the flag was never attempted, so the device stays unrecorded and the next call retries. A `hipSetDeviceFlags` failure is recorded and reported once on stderr. Without that line, a device left in active-wait mode would reproduce #1876 with no trace, which is the silent hard hang this PR exists to remove.
+
+### Test in a child process with a timeout
+
+The hang has no timeout of its own and the plan cache capacity is read from the environment on the first FFT, so the test cannot run in-process. The child-process pattern gives both: a fresh environment and an external kill. The test was confirmed to fail (`the child did not finish within 120s`, after 8 of 48 round trips) with the early calls commented out, so it guards the ordering and not just FFT correctness.
+
+## 4. Validation
+
+Author's runs (gfx1151, Radeon 8060S, HIP 7.15.26333, ROCm 10.0 core, rocFFT 1.0.39):
+
+- Probe at cap 128 (`--plans 40`), three consecutive runs: all exit 0, 40/40 round trips ok. This run evicts nothing, so the eviction runs below are the meaningful ones.
+- Probe at caps 16 and 32, which hung before the fix after `plan 7` and `plan 15`: both exit 0, 40/40 ok. Cap 16 with `--plans 100` also completes.
+- Probe correctness mode: all 25 comparisons within 1e-5 of the CPU stream; worst 4.2e-7 (round trip at n=2048 and n=1200).
+- Ordering proof under `AMD_LOG_LEVEL=3`: `hipSetDeviceFlags ( 4 )` on line 26, before `hipExtMallocWithFlags` (line 32), `hipStreamQuery ( <null> )` (line 50) and the first `Created SWq` (line 52). Before the fix the flag came after the null-stream `Created SWq`.
+- `cargo test --features rocm --test rocm_fft_plan_cache`: passes (9.97 s and 6.91 s); fails with the early calls commented out.
+- Standalone reproducer: exit 124 at `iter 7` with the early touch, `done` without it.
+- Clippy on the new test with `-D warnings`, `dead_doc_pointers`, and the fast `make verify-*` gates: pass.
+
+Orchestrator verification (gfx1151, branch head `3ccab4e5`, which merges origin/main `ac02b2dc`):
+
+- The regression test and `dead_doc_pointers` pass (child finishes in about 7 s); the regression test fails with the early calls commented out; the probe at cap 128 three times, at caps 16 and 32, and in correctness mode all complete; `scripts/ci/rocm_smoke.sh` with Qwen3-0.6B-4bit generates 32 tokens.
+- `make verify-rocm` ran every step. Versions, kernel dtype keys, kernel port dispatch, llama-compat, fmt, workspace clippy with `--features rocm`, and the ROCm smoke (32 tokens) passed.
+- `verify-test-rocm` failed in three targets with exactly the 37 known baseline failures and nothing else:
+  - `-p mlxcel-core --lib`: 35 failures. 34 are fused paged-attention tests without a ROCm port (#1814); the other is the bf16 `prefill_dense_gemm_matches_qmm_bytes_where_eligible`.
+  - `gemma3_backbone_tests::gelu_approx_matches_mlx_nn_bit_for_bit` and `tests::family_order_is_exhaustive`, both from #2037.
+- The new `tests/rocm_fft_plan_cache.rs` passed inside the full suite.
+
+## 5. Learning Points
+
+- **A threshold that moves with the configuration is often an event, not a limit.** "Hangs above N live plans" became "hangs at the first eviction" as soon as the probe printed plan indices at two different caps. Measuring the position of the hang at more than one setting was what broke the live-count theory.
+- **Get the stack before forming hypotheses.** The issue's four hypotheses were all reasonable and all wrong. One rocgdb backtrace moved the problem from rocFFT to `hipFree`, and one `AMD_LOG_LEVEL=4` line moved it from `hipFree` to signal typing.
+- **Some HIP device flags must be set before the first queue, not just before the first stream.** CLR types a queue's completion signals by the wait mode at creation and does not re-type them. Any HIP call that touches the null stream (a query, a sync, a copy, a launch) can create a queue. Code that sets device flags "when the device is first used" has to define first use as the first HIP call of any kind, including those made by the allocator.
+- **Reduce to the runtime before blaming a library.** The standalone reproducer shows the defect with two HIP calls in the wrong order. Without it, the natural next step would have been a rocFFT bug report about something rocFFT does not do.
+- **Re-check comments that justify a limit.** The "captured plan outlives the call" comment was the stated reason for cache headroom. It described behavior `launch_kernel` does not have.
+
+## 6. What Is Not Verified
+
+- **Other GPUs and ROCm versions.** Only gfx1151 with HIP 7.15.26333 and rocFFT 1.0.39 was measured. The CLR behavior is generic, but a different HIP could type signals differently.
+- **The threshold at small capacities.** Why capacities of 12 or below completed at all is not characterized; the signal-ring position is a guess.
+- **The `MLX_GRAPH_PREFILL_REPLAY=1` path for FFT.** It records the hipFFT call into a graph that launches later, and the eviction-safety argument in `execute_fft` covers only the eager launch.
+- **The `hipSetDeviceFlags` failure path.** The stderr message is reached only if the call fails, and no test forces that.
+- **Metal and CUDA.** Not available on this host. The change touches only `patches-rocm/` and a `#![cfg(feature = "rocm")]` test, so neither build is affected by construction.
+- **Multi-GPU hosts.** The helper touches only the requested device by design, but no multi-GPU ROCm host was tested.
+
+## 7. Remaining Work
+
+- File the standalone reproducer as a CLR report. None has been filed.
+- Propose `LOCAL_FIXES.md` entry 20 (and the updated entry 18) to the ROCm fork as part of #1813.
+- Pre-existing issues outside this diff: `lru_cache.h` evicts from an empty list when `MLX_ROCM_FFT_CACHE_SIZE=0`, and `std::stoul` throws on a non-numeric value.
+- Examine FFT under `MLX_GRAPH_PREFILL_REPLAY=1`.
+- Baseline test failures: #1814 (fused paged-attention ports, 34 failures), the bf16 `prefill_dense_gemm_matches_qmm_bytes_where_eligible` failure, and the two #2037 failures.
+
+Refs: #1876 (closed by this PR), #1801, #1825, #1878, #2033, #2046.
diff --git a/TECHNICAL_REPORTS/2049-rocm-blocking-sync-first-queue-20260930.ko.md b/TECHNICAL_REPORTS/2049-rocm-blocking-sync-first-queue-20260930.ko.md
new file mode 100644
index 000000000..ede286ed7
--- /dev/null
+++ b/TECHNICAL_REPORTS/2049-rocm-blocking-sync-first-queue-20260930.ko.md
@@ -0,0 +1,142 @@
+# 기술 보고서: PR #2049 - 첫 HIP 큐 이전에 ROCm blocking-sync 설정, FFT plan 캐시를 128로 복원
+
+**날짜**: 2026-09-30
+
+**상태**: gfx1151 호스트에서 구현 및 검증 완료, 머지 대기 중.
+
+**언어**: C++ (ROCm overlay의 HIP 런타임 호출), Rust (통합 테스트), Markdown
+
+**위험도**: 중간 (모든 ROCm 프로세스가 장치 대기 모드를 설정하는 시점을 allocator의 hot path에서 바꿉니다. ROCm 전용 overlay와 `#![cfg(feature = "rocm")]` 테스트만 바뀌며, Metal과 CUDA 빌드는 `patches-rocm/`을 복사하지 않습니다)
+
+## 요약
+
+이슈 #1876 (에픽 #1801의 일부)은 CUDA 파일이 128을 쓰는데 왜 ROCm hipFFT plan 캐시는 8로 제한해야 하는지 물었습니다. 8을 넘기면 FFT를 많이 쓰는 작업이 GPU가 쉬는 상태로 영원히 멈췄고, 원인은 "rocFFT 내부 어딘가"로만 기록되어 있었습니다. 이 PR의 조사 결과, 멈춤은 rocFFT 안에 있지 않았고 살아 있는 plan 수와도 무관했습니다. ROCm 백엔드의 순서 버그였습니다. `hipSetDeviceFlags(hipDeviceScheduleBlockingSync)`는 `rocm::device()`에서 실행되었는데, 그 전에 allocator가 `hipStreamQuery(nullptr)`로 null stream의 HIP 큐를 이미 만들어 두었습니다. CLR은 그 큐에 non-interrupt 완료 signal을 주었고, 이 signal에는 이후 `hipFree`가 의존하는 async handler를 붙일 수 없습니다. 그래서 null stream을 동기화해야 하는 첫 `hipFree`가 돌아오지 않았습니다. 첫 plan eviction 때의 rocFFT `hipfftDestroy`는 mlxcel에서 null stream 작업 뒤에 `hipFree`를 부른 첫 코드였을 뿐입니다.
+
+MLX도 hipFFT도 없는 30줄짜리 독립 HIP 프로그램이 이 멈춤을 재현합니다. null stream을 건드리고, 플래그를 설정한 다음, `hipMalloc`, `hipMemcpy`, null stream 커널, `hipFree`를 반복합니다. 첫 `hipFree`에서 멈추고, 앞의 null stream 접근만 빼면 끝까지 실행됩니다.
+
+수정은 `rocm::ensure_device_flags(int)`와 `ensure_current_device_flags()`를 추가해, 장치에 큐가 하나도 없을 때 장치마다 한 번 플래그를 설정하고, `rocm::device()`보다 먼저 HIP에 닿을 수 있는 모든 경로에서 이를 호출합니다. plan 캐시는 128로 돌아가고, 새 통합 테스트 `tests/rocm_fft_plan_cache.rs`가 자식 프로세스에서 캐시 크기 16으로 eviction이 일어나는 round trip 48회를 돌려 멈춤이 돌아오면 실패합니다. 이 결함은 처음부터 FFT에 한정된 것이 아니었습니다. null stream에 대기 중인 작업이 있는 상태의 모든 장치 전체 동기화가 노출되어 있었으므로, 수정 범위는 이를 발견한 이슈보다 넓습니다.
+
+## 1. 문제 정의
+
+### 이슈에 기록된 내용
+
+#1825의 FFT 이식(PR #1856, #1861)은 `fft.hip`의 hipFFT `LRUBytesKeyCache`에 `default_capacity` 8을 주고, rocFFT 내부의 근본 원인은 모른다는 주석을 달았습니다. 측정은 일관되었지만 오해를 부르는 형태였습니다. 26개 변환 형태를 돌리는 probe는 용량 8에서 끝까지 실행되고 16, 24, 32, 128에서는 항상 순서상 같은 위치에서 멈췄습니다. 케이스 순서를 바꾸면 그 위치에서 실행되는 것이 멈췄고, 한 형태를 60번 반복하면 멈추지 않았습니다. 이 패턴은 "살아 있는 plan이 너무 많다"를 가리켰고, 이슈가 제시한 가설들도 여기서 나왔습니다. rocFFT 런타임 컴파일, `hipfftSetAutoAllocation(handle, 0)`로 설정한 수동 work area, 그리고 `shared_ptr` 캡처 때문에 plan이 캐시 항목보다 오래 사는 문제입니다. 독립 프로그램은 hipFFT plan 64개를 문제없이 만들었으므로, 한계가 MLX의 사용 방식에 달려 있다는 것까지는 알았지만 어떻게인지는 몰랐습니다.
+
+제한의 비용은 실제로 있었습니다. 장치 FFT를 쓰는 mlxcel 경로 두 개인 Kokoro 음성 합성과 Phi-4-multimodal 오디오 front end는 CUDA보다 훨씬 자주 plan을 다시 만들었습니다.
+
+### 조사에서 찾은 것
+
+이 부분이 PR의 핵심이며, 증상이 매 단계 원인과 다른 쪽을 가리켰기 때문에 추론 과정을 남길 가치가 있습니다.
+
+**1단계: 멈춤의 시점 찾기.** probe를 `fft_numeric_probe --plans N`(원래 PR #1878)으로 다시 만들었습니다. 캐시 16에서는 `plan 7` 다음에, 32에서는 `plan 15` 다음에 멈췄습니다. probe 한 반복은 plan 두 개(R2C, C2R)를 만들므로, plan 7 다음 반복은 17번째 plan을, plan 15 다음 반복은 33번째 plan을 만듭니다. 두 경우 모두 eviction을 처음 일으키는 plan입니다. 멈춤은 살아 있는 plan 수의 문턱이 아니라 첫 `hipfftDestroy`에서 일어났습니다. LRU 캐시는 새 항목을 넣는 도중에 eviction을 하므로, 멈춤이 "다음 plan 생성이 막힌다"처럼 보였고 이슈도 그렇게 기술했습니다.
+
+**2단계: 멈춤의 스택 찾기.** 멈춘 스레드에 rocgdb를 붙이면 `hipfftDestroy -> rocfft_plan_destroy -> ~ExecPlan -> ~TreeNode -> gpubuf_t::free -> hipFree -> ... -> std::condition_variable::wait`가 나왔습니다. GPU는 쉬고 있었고 rocFFT RTC 헬퍼 프로세스도 없었으므로 런타임 컴파일 가설은 배제되었습니다. rocFFT는 장치 버퍼를 해제하고 있었을 뿐이고, 대기는 HIP의 `hipFree` 안에 있었습니다. `hipFree`는 해제 전에 장치의 모든 stream을 동기화합니다(`Device::SyncAllStreams`).
+
+**3단계: 대기가 깨어나지 않는 이유 찾기.** `AMD_LOG_LEVEL=4`에서 멈추기 직전 마지막 줄은 `rocvirtual.cpp:844 hsa_amd_signal_async_handler() failed to set the handler!`였습니다. ROCr는 event mailbox가 없는 signal, 즉 non-interrupt signal에는 async handler를 거부합니다. `SyncAllStreams`는 null stream에 marker를 넣고 그 marker의 완료 handler가 호출되기를 기다립니다. handler가 설치되지 않았으니 condition variable은 끝내 notify되지 않았습니다.
+
+**4단계: non-interrupt signal의 출처 찾기.** CLR은 큐를 만들 때 적용 중인 장치 대기 모드에 따라 큐의 완료 signal ring 종류를 정합니다. 기본값인 active wait 모드에서는 호스트가 signal을 spin으로 기다리므로 CLR은 가벼운 `HSA_AMD_SIGNAL_AMD_GPU_ONLY`(non-interrupt) signal을 할당합니다. active wait가 켜져 있는 동안에는 `HwQueueTracker::ActiveSignal`이 handler가 필요한 명령마다 interrupt signal로 바꿔 줍니다. `hipDeviceScheduleBlockingSync`가 active wait를 끄면, 큐가 애초에 interrupt signal로 만들어졌다고 가정하고 이 교체를 멈춥니다. HIP 로그가 이 가정을 어긴 큐를 보여 주었습니다. null stream의 큐는 unified 버퍼를 처음 호스트에서 읽을 때 `allocator::Buffer::raw_ptr()`의 `hipStreamQuery(nullptr)`가 만들었고, 그 뒤에야 `rocm::device()`가 `hipSetDeviceFlags`를 호출했습니다. 그래서 null stream 큐는 프로세스가 끝날 때까지 non-interrupt signal을 유지했고, 그중 하나에 떨어지면서 handler가 필요한 첫 marker가 대기자를 멈추게 했습니다.
+
+**5단계: MLX 밖에서 증명하기.** PR 본문의 독립 재현 프로그램은 나머지를 모두 걷어 냅니다(`hipcc --offload-arch=gfx1151 -O2`로 빌드).
+
+```cpp
+if (early_touch) (void)hipStreamQuery(nullptr);        // creates the null-stream queue now
+(void)hipSetDeviceFlags(hipDeviceScheduleBlockingSync); // flips active wait off
+// loop: hipMalloc, hipMemcpy H2D, kernel on the null stream,
+//       keep 8 buffers live, hipFree the oldest
+```
+
+`timeout 30 ./repro 1 200`은 `iter 7`을 출력하고 exit 124로 끝나며, 첫 `hipFree`에서 멈춥니다. 앞의 null stream 접근만 뺀 `./repro 0 200`은 `done`을 출력합니다. hipFFT도, MLX도, plan 캐시도 없습니다. 결함 전체가 HIP 호출 두 개의 순서입니다.
+
+### 이슈의 가설이 틀린 이유와 이로써 배제되는 것
+
+- **살아 있는 plan 수.** 멈춤은 살아 있는 plan 수가 아니라 첫 eviction에 달려 있습니다. 독립 재현 프로그램에는 plan이 아예 없습니다.
+- **rocFFT 런타임 컴파일.** RTC 헬퍼는 실행 중이 아니었고, 스택은 컴파일이나 캐시 lock이 아니라 `hipFree`에 있었습니다.
+- **plan이 캐시 항목보다 오래 사는 문제.** `execute_fft` 주석은 캡처된 `shared_ptr`가 비동기 실행 동안 plan을 살려 둔다고 했습니다. 그렇지 않습니다. `CommandEncoder::launch_kernel`은 lambda를 동기적으로 호출하고 반환할 때 버립니다. 캐시가 "여유를 남겨야 한다"는 근거가 이 전제였고, 전제는 틀렸습니다.
+- **rocFFT 자체.** rocFFT의 역할은 null stream 작업 뒤에 `hipFree`를 부른 첫 호출자였다는 것뿐입니다. rocFFT에 보고할 것은 없습니다. 대신 재현 프로그램은 CLR 보고 후보입니다(아직 제출하지 않았습니다).
+
+설명되지 않은 부분이 하나 있습니다. 26개 형태에서는 용량 8에서도 eviction이 일어나는데, 왜 용량 12 이하는 끝까지 실행되고 16 이상은 첫 eviction에서 멈췄는지입니다. PR의 작업 가설은 첫 `hipFree` 시점의 null stream signal ring 위치입니다(재현 프로그램도 커널 8개 뒤에 처음 멈춥니다). 하지만 이 부분은 규명하지 않았습니다.
+
+## 2. 변경 요약
+
+- **`device.cpp` / `device.h`: `ensure_device_flags(int)`와 `ensure_current_device_flags()`.** 장치 인덱스마다 한 번 `hipDeviceScheduleBlockingSync`를 설정합니다. 모든 unified 할당과 unified 버퍼의 모든 호스트 읽기에서 실행되므로, 64 미만 인덱스는 lock 없는 atomic 비트마스크로 처리하고, 그 이상은 mutex와 vector로 처리합니다. 호출자가 다른 장치에 있을 때만 대상 장치로 전환하고 끝나면 원래 장치로 되돌립니다. `hipSetDevice`가 실패하면 장치를 기록하지 않아 다음 호출이 다시 시도합니다. `hipSetDeviceFlags`가 실패하면 오류를 지우고, 장치와 HIP 오류를 적은 한 줄을 stderr에 출력하며, 할당마다 재시도하지 않도록 장치를 기록합니다. `rocm::device()`는 이제 플래그를 직접 설정하지 않고 이 헬퍼를 호출합니다. 플래그 값 자체는 바뀌지 않았습니다.
+- **`rocm::device()`보다 먼저 큐를 만들 수 있는 호출 지점.** `allocator.cpp`의 `unified_malloc`(프로세스의 첫 HIP 작업인 경우가 많습니다), `Buffer::raw_ptr()`의 두 분기(unified 메모리의 null stream query와 discrete 메모리의 `hipDeviceSynchronize`), `RocmAllocator` 생성자에서 pool 장치마다 만드는 opt-in async pool 해제 stream, `eval.cpp`의 `gpu::init()`에서 `hipFree(nullptr)` 전, 그리고 `host_stage.cpp`의 `staged_write`에서 장치 동기화 전입니다. 체크포인트 저장이 스레드 장치의 첫 동기화일 수 있기 때문입니다.
+- **`fft.hip`.** `default_capacity`는 `mlx/backend/cuda/fft.cu`와 같은 128로 돌아갑니다. 캐시 주석은 실제 원인을 적고 `LOCAL_FIXES.md` 항목 18과 20을 가리킵니다. `execute_fft` 주석도 바로잡았습니다. 캡처된 plan은 enqueue 동안만 살고, `hipFree`가 먼저 장치를 동기화하므로 eager launch에서는 eviction이 여전히 안전하며, opt-in `MLX_GRAPH_PREFILL_REPLAY=1` 그래프 경로는 검토하지 않았다고 적습니다.
+- **`tests/rocm_fft_plan_cache.rs` (신규, 285줄).** 부모 테스트는 `MLX_ROCM_FFT_CACHE_SIZE=16`과 120초 제한으로 같은 바이너리의 ignored 자식 테스트를 다시 실행하고, 시간 안에 끝나지 않으면 자식을 종료합니다. 자식은 서로 다른 길이로 rfft/irfft round trip 48회를 실행합니다(16칸 캐시에 plan 96개이므로 round trip 8회 뒤부터 eviction이 계속 일어납니다). 그다음 길이 2, 8, 400, 512, 1024, 1200, 2048에서 rfft, irfft, round trip, complex fft와 512 batch rfft를 CPU stream과 1e-5 이내로 비교합니다. 자식 본문은 `MLXCEL_ROCM_FFT_CHILD`가 설정되었을 때만 실행되므로 `--include-ignored` 실행이 공유 프로세스에서 이를 돌리지 않습니다.
+- **`LOCAL_FIXES.md`.** 항목 18(FFT 이식)은 용량 128과 그 경위를 기술합니다. 새 항목 20은 순서 결함, CLR 메커니즘, 모든 호출 지점, 실패 처리, 근거와 재현 프로그램을 fork 쪽 upstream 반영 후보로 기록합니다.
+- **`docs/environment-variables.md`, `docs/installation.md`.** `MLX_ROCM_FFT_CACHE_SIZE`의 기본값을 128로 적고, 멈춤이 해결되었다고 기술합니다(#1876).
+
+커밋 이력: `84e8ad19`가 수정 본체입니다. `6e4ac7d6`은 `fft.hip` 주석에 `LOCAL_FIXES.md` 항목 번호를 적고, `9a9c1b06`은 헬퍼의 실패 경로를 드러내고 테스트가 자식 polling 실패 시 자식을 종료하도록 합니다. `78a53b15`와 `3ccab4e5`는 origin/main을 머지합니다(첫 번째는 PR #2046의 `LOCAL_FIXES.md` 항목 19를 가져오기 위한 것입니다).
+
+## 3. 기술적 선택과 그 이유
+
+### 캐시 제한이 아니라 순서를 고칩니다
+
+캐시를 8로 두면 null stream 작업 뒤에 오는 다른 모든 `hipFree`와 장치 전체 동기화에 결함이 그대로 남습니다. FFT는 처음 걸린 호출자였을 뿐이고, 런타임에 장치 메모리를 해제하는 이후의 연산이 FFT 없이도 같은 식으로 멈출 수 있었습니다. 순서를 고치면 이런 멈춤 전체가 사라지고, 캐시 용량은 안전 한계가 아니라 메모리와 plan 재생성 사이의 절충으로 돌아갑니다.
+
+### 프로세스 시작 시 한 번이 아니라 모든 진입점에서 일찍 설정합니다
+
+HIP에는 어떤 코드든 첫 HIP 호출 이전에 실행되는 hook이 없고, MLX가 장치를 고르기 전에는 장치 인덱스를 알 수 없습니다. 플래그는 해당 장치의 첫 큐 이전에 그 장치에 설정되어야 하므로, `rocm::device()`보다 먼저 큐를 만들 수 있는 각 지점에서 헬퍼를 실행합니다. 첫 호출 이후의 비용은 atomic load 한 번과 비트 검사이며, 할당 경로에 있기 때문에 fast path를 lock 없이 만들었습니다.
+
+### blocking-sync를 없애지 않고 유지합니다
+
+`hipSetDeviceFlags`를 아예 빼도 모드는 일관됩니다(전부 active wait). 하지만 모든 GPU 대기 동안의 호스트 CPU 사용량이 바뀌고, 이는 fork가 처음에 blocking-sync를 설정한 이유이기도 합니다. PR은 플래그 값을 유지하고 적용 시점만 바꿉니다.
+
+### 장치별로, 현재 장치만
+
+헬퍼는 인덱스별로 따로 기록하므로 장치 0을 먼저 설정해도 장치 1이 누락되지 않고, 모든 장치를 순회하지도 않습니다. 기존 `rocm::device()` 주석이 그 이유를 기록하고 있습니다. 다중 GPU 호스트에서 다른 GPU에 context나 큐를 만드는 것이 TB5 링크 너머 discrete GPU의 큐를 멈추게 한 원인이었습니다. 헬퍼는 이 제약을 지키고 호출자의 현재 장치를 복원하므로 allocator에서 부작용 없이 부를 수 있습니다.
+
+### 실패를 드러내고, 성공할 수 있는 것만 재시도합니다
+
+`hipSetDevice` 실패는 플래그 설정을 시도조차 못 했다는 뜻이므로 장치를 기록하지 않고 다음 호출이 다시 시도합니다. `hipSetDeviceFlags` 실패는 기록하고 stderr에 한 번 알립니다. 이 한 줄이 없으면 active wait 모드로 남은 장치가 아무 흔적 없이 #1876을 재현하게 되고, 이는 이 PR이 없애려는 조용한 완전 멈춤 그대로입니다.
+
+### 제한 시간이 있는 자식 프로세스에서 테스트합니다
+
+멈춤 자체에는 제한 시간이 없고, plan 캐시 용량은 첫 FFT 때 환경 변수에서 읽으므로 같은 프로세스 안에서는 테스트할 수 없습니다. 자식 프로세스 방식은 새 환경과 외부 종료를 모두 제공합니다. 앞쪽 호출들을 주석 처리하면 테스트가 실패함(`the child did not finish within 120s`, round trip 48회 중 8회 뒤)을 확인했으므로, 이 테스트는 FFT 정확도뿐 아니라 순서 자체를 지킵니다.
+
+## 4. 검증
+
+작성자 실행 (gfx1151, Radeon 8060S, HIP 7.15.26333, ROCm 10.0 core, rocFFT 1.0.39):
+
+- 캐시 128에서 probe(`--plans 40`) 연속 3회: 모두 exit 0, round trip 40/40 정상. 이 실행은 eviction이 없으므로 의미 있는 것은 아래 eviction 실행입니다.
+- 수정 전 `plan 7`, `plan 15` 뒤에 멈췄던 캐시 16과 32에서 probe: 둘 다 exit 0, 40/40 정상. 캐시 16에 `--plans 100`도 끝까지 실행됩니다.
+- probe 정확도 모드: 25개 비교 모두 CPU stream과 1e-5 이내, 최악은 4.2e-7(n=2048과 n=1200 round trip).
+- `AMD_LOG_LEVEL=3`에서 순서 증명: `hipSetDeviceFlags ( 4 )`가 26번째 줄로, `hipExtMallocWithFlags`(32번째 줄), `hipStreamQuery ( <null> )`(50번째 줄), 첫 `Created SWq`(52번째 줄)보다 앞섭니다. 수정 전에는 플래그가 null stream의 `Created SWq` 뒤에 나왔습니다.
+- `cargo test --features rocm --test rocm_fft_plan_cache`: 통과(9.97초, 6.91초). 앞쪽 호출을 주석 처리하면 실패합니다.
+- 독립 재현 프로그램: 앞의 null stream 접근이 있으면 `iter 7`에서 exit 124, 없으면 `done`.
+- 새 테스트에 대한 `-D warnings` clippy, `dead_doc_pointers`, 빠른 `make verify-*` 게이트: 통과.
+
+오케스트레이터 검증 (gfx1151, origin/main `ac02b2dc`를 머지한 브랜치 head `3ccab4e5`):
+
+- 회귀 테스트와 `dead_doc_pointers` 통과(자식은 약 7초에 끝남). 앞쪽 호출을 주석 처리하면 회귀 테스트가 실패합니다. 캐시 128에서 probe 3회, 캐시 16과 32, 정확도 모드 모두 끝까지 실행됩니다. Qwen3-0.6B-4bit으로 `scripts/ci/rocm_smoke.sh`가 32 토큰을 생성합니다.
+- `make verify-rocm`의 모든 단계를 실행했습니다. versions, kernel dtype keys, kernel port dispatch, llama-compat, fmt, `--features rocm` workspace clippy, ROCm smoke(32 토큰)가 통과했습니다.
+- `verify-test-rocm`은 세 타깃에서 실패했고, 알려진 기준 실패 37개와 정확히 같았으며 그 외는 없었습니다.
+  - `-p mlxcel-core --lib`: 35개 실패. 34개는 ROCm 이식이 없는 fused paged-attention 테스트(#1814)이고, 나머지 하나는 bf16 `prefill_dense_gemm_matches_qmm_bytes_where_eligible`입니다.
+  - `gemma3_backbone_tests::gelu_approx_matches_mlx_nn_bit_for_bit`와 `tests::family_order_is_exhaustive`, 둘 다 #2037에서 온 것입니다.
+- 새 `tests/rocm_fft_plan_cache.rs`는 전체 스위트 안에서 통과했습니다.
+
+## 5. 학습 포인트
+
+- **설정에 따라 움직이는 문턱은 한계가 아니라 사건인 경우가 많습니다.** 두 캐시 크기에서 probe가 plan 번호를 출력하자 "살아 있는 plan이 N개를 넘으면 멈춘다"가 "첫 eviction에서 멈춘다"로 바뀌었습니다. 멈춤 위치를 여러 설정에서 잰 것이 plan 수 이론을 무너뜨렸습니다.
+- **가설을 세우기 전에 스택부터 확보합니다.** 이슈의 가설 네 개는 모두 그럴듯했고 모두 틀렸습니다. rocgdb backtrace 하나가 문제를 rocFFT에서 `hipFree`로 옮겼고, `AMD_LOG_LEVEL=4` 한 줄이 `hipFree`에서 signal 종류로 옮겼습니다.
+- **일부 HIP 장치 플래그는 첫 stream이 아니라 첫 큐 이전에 설정해야 합니다.** CLR은 큐를 만들 때의 대기 모드로 완료 signal 종류를 정하고 다시 바꾸지 않습니다. null stream을 건드리는 모든 HIP 호출(query, sync, copy, launch)이 큐를 만들 수 있습니다. "장치를 처음 쓸 때" 플래그를 설정하는 코드는 allocator가 하는 호출까지 포함해 어떤 종류든 첫 HIP 호출을 첫 사용으로 정의해야 합니다.
+- **라이브러리를 탓하기 전에 런타임 수준까지 줄입니다.** 독립 재현 프로그램은 잘못된 순서의 HIP 호출 두 개로 결함을 보여 줍니다. 이것이 없었다면 자연스러운 다음 단계는 rocFFT가 하지도 않는 일에 대한 rocFFT 버그 보고였을 것입니다.
+- **한계를 정당화하는 주석은 다시 확인합니다.** "캡처된 plan이 호출보다 오래 산다"는 주석이 캐시 여유의 근거였는데, `launch_kernel`에는 그런 동작이 없었습니다.
+
+## 6. 검증되지 않은 부분
+
+- **다른 GPU와 ROCm 버전.** gfx1151, HIP 7.15.26333, rocFFT 1.0.39만 측정했습니다. CLR 동작은 일반적이지만 다른 HIP 버전은 signal 종류를 다르게 정할 수 있습니다.
+- **작은 용량에서의 문턱.** 용량 12 이하가 왜 끝까지 실행되었는지는 규명하지 않았고, signal ring 위치는 추정입니다.
+- **FFT의 `MLX_GRAPH_PREFILL_REPLAY=1` 경로.** 이 경로는 hipFFT 호출을 나중에 실행되는 그래프에 기록하며, `execute_fft`의 eviction 안전성 논증은 eager launch만 다룹니다.
+- **`hipSetDeviceFlags` 실패 경로.** stderr 메시지는 호출이 실패할 때만 나오며, 이를 강제하는 테스트는 없습니다.
+- **Metal과 CUDA.** 이 호스트에서는 사용할 수 없습니다. 변경은 `patches-rocm/`과 `#![cfg(feature = "rocm")]` 테스트에만 닿으므로 두 빌드는 구조적으로 영향을 받지 않습니다.
+- **다중 GPU 호스트.** 헬퍼는 설계상 요청된 장치만 건드리지만, 다중 GPU ROCm 호스트에서는 테스트하지 않았습니다.
+
+## 7. 남은 작업
+
+- 독립 재현 프로그램으로 CLR 보고를 제출합니다. 아직 제출하지 않았습니다.
+- `LOCAL_FIXES.md` 항목 20(과 갱신된 항목 18)을 #1813의 일부로 ROCm fork에 제안합니다.
+- 이 diff 밖의 기존 문제: `MLX_ROCM_FFT_CACHE_SIZE=0`일 때 `lru_cache.h`가 빈 리스트에서 eviction을 하고, 숫자가 아닌 값에서는 `std::stoul`이 예외를 던집니다.
+- `MLX_GRAPH_PREFILL_REPLAY=1`에서의 FFT를 검토합니다.
+- 기준 테스트 실패: #1814(fused paged-attention 이식, 실패 34개), bf16 `prefill_dense_gemm_matches_qmm_bytes_where_eligible` 실패, #2037의 실패 두 개.
+
+참고: #1876 (이 PR로 닫힘), #1801, #1825, #1878, #2033, #2046.