cuda: step 8 - CUDA Graph capture+replay for routed-MoE decode

Capture the mmvq routed-MoE decode kernel sequence into a
cudaGraphExec_t on first encounter with a given (layer-shape,
buffer-pointer) tuple; replay via cudaGraphLaunch on subsequent calls
with the same key.  Replay eliminates the ~8 per-kernel CPU<->driver
round trips per MoE layer, the dominant overhead at decode where
individual kernels are small.

Graph cache: 256-entry direct-mapped hash table keyed by all
(gate_offset, up_offset, down_offset, n_tokens, q4k_path, shape ints,
buffer pointers).  Hash collisions tear down the prior exec.  For V4
Flash decode (43 layers, all same shape-class) we keep 43 entries
warm after the first forward pass.

Stream plumbing: ds4 historically used stream=0 (legacy default null
stream) for kernel launches and cudaStreamPerThread for the
ds4_naive_pool's cudaMallocAsync.  cudaStream capture requires (a) a
non-default stream and (b) all stream-ordered ops on that stream.
This commit:
  - Creates g_moe_stream (cudaStreamNonBlocking) lazily on first use
  - Adds a thread_local current-pool-stream in ds4_naive_pool plus
    ds4_pool_set_stream() that ds4_mmq_*_moe_vec wrappers call on
    entry to redirect cudaMallocAsync onto the caller's stream
  - Plumbs moe_stream through all kernel launches in the mmvq decode
    branch (4 ds4_mmq_*_moe_vec calls, swiglu kernel, moe_sum kernel)
  - Uses cudaStreamCaptureModeThreadLocal - strict mode that
    correctly fences out the legacy default stream

Failure handling: a mmvq_decode_bail label in the captured region
calls cudaStreamEndCapture and discards the partial graph on any
launch failure, ensuring the stream isn't left in a broken capture
state.  Caller falls through to the un-captured mmq path.

Opt-in: DS4_CUDA_MOE_GRAPHS=1 (default off until validated
end-to-end).  AGENT.md updated with the new env-var doc.

Bench on PRO 6000 Blackwell (sm_120) at ctx=2048, gen-tokens=128,
N=5 runs per config:
  - graphs OFF (post-Step-6 baseline): gen 43.20 tok/s (IQR 0.02)
  - graphs ON                        : gen 45.12 tok/s (IQR 0.20)
  - delta:                             gen +4.44% (Mann-Whitney U
                                       p=0.0079, B wins)
  - prefill:                           1009.75 -> 1009.31, p=0.69
                                       (inconclusive, no real change)

The Step 8 win is smaller than the plan's projected +15-30% range
because this commit captures ONLY the routed-MoE block.  Attention
projections (also significant launch overhead at decode) are still
on the un-captured legacy path.  Future work: extend capture to the
dense Q8_0 vec path and to the layer-norm + RoPE preamble.

Cumulative trajectory on PRO 6000:
  - Pre-lift baseline:        gen 38.0 tok/s
  - Session 1 (mmq default):  gen 38.0 (within noise)
  - Session 2 + Step 6:       gen 43.2 (+13.7% vs baseline)
  - Session 2 + Step 8:       gen 45.1 (+18.7% vs baseline)

AGENT.md

@@ -113,6 +113,21 @@ to opt in to the newer paths.
   range 0-8.  0 disables (same as `DS4_CUDA_NO_MMVQ_DECODE=1` for the
   MoE path).  Values 2-8 extend mmvq coverage to short-prefill batches
   but require n_tokens * n_expert_used ≤ 8 for the down matmul.
+- `DS4_CUDA_MOE_GRAPHS=1` (default off): opt-in CUDA Graph capture for
+  the mmvq routed-MoE decode block.  Step 8 of the optimization plan.
+  Each per-layer kernel sequence (gate + up + swiglu + down + sum =
+  ~8 launches) is captured into a `cudaGraphExec_t` on first execution
+  and replayed via `cudaGraphLaunch` on subsequent calls with the same
+  (gate_offset, up_offset, down_offset, n_tokens, q4k_path, buffer
+  pointers) tuple.  The cache holds up to 256 entries (one per layer
+  shape-class).  Replay eliminates ~5-15µs of CPU↔driver
+  round-trip per kernel launch, the dominant overhead at decode where
+  individual kernels are small.  Requires an explicit non-default
+  stream (`g_moe_stream`) for capture; routes the ds4_mmq pool's
+  `cudaMallocAsync` through the same stream via the thread-local
+  `ds4_pool_set_stream()` so allocations don't invalidate capture.
+  Falls through to the un-captured mmvq decode path on any capture
+  error; opt-in until validated end-to-end.
 
 ## Testing
 

cuda/mmq/ds4_ggml_stubs.cu

@@ -100,6 +100,22 @@ int64_t ggml_time_us() {
 
 namespace {
 
+/* Thread-local stream that ds4_naive_pool uses for cudaMallocAsync /
+ * cudaFreeAsync.  Defaults to cudaStreamPerThread (preserves prior
+ * behaviour).  Step 8 / CUDA Graphs sets this to the capture stream
+ * before allocating, so the alloc node lives on the captured stream
+ * and capture is not invalidated.  Set via ds4_pool_set_stream() from
+ * ds4_mmq.cu wrappers. */
+static thread_local cudaStream_t t_ds4_pool_stream = cudaStreamPerThread;
+
+extern "C" void ds4_pool_set_stream(cudaStream_t stream) {
+    t_ds4_pool_stream = stream ? stream : cudaStreamPerThread;
+}
+
+extern "C" cudaStream_t ds4_pool_get_stream(void) {
+    return t_ds4_pool_stream;
+}
+
 struct ds4_naive_pool : public ggml_cuda_pool {
     int device;
 
@@ -108,7 +124,7 @@ struct ds4_naive_pool : public ggml_cuda_pool {
     void * alloc(size_t size, size_t * actual_size) override {
         ggml_cuda_set_device(device);
         void * ptr = nullptr;
-        CUDA_CHECK(cudaMallocAsync(&ptr, size, cudaStreamPerThread));
+        CUDA_CHECK(cudaMallocAsync(&ptr, size, t_ds4_pool_stream));
         if (actual_size) *actual_size = size;
         return ptr;
     }
@@ -116,7 +132,7 @@ struct ds4_naive_pool : public ggml_cuda_pool {
     void free(void * ptr, size_t /*size*/) override {
         if (!ptr) return;
         ggml_cuda_set_device(device);
-        CUDA_CHECK(cudaFreeAsync(ptr, cudaStreamPerThread));
+        CUDA_CHECK(cudaFreeAsync(ptr, t_ds4_pool_stream));
     }
 };
 

cuda/mmq/ds4_mmq.cu

@@ -738,6 +738,12 @@ int ds4_mmq_moe_vec_impl(
         return -1;
     }
 
+    // Route the pool's cudaMallocAsync / cudaFreeAsync through the same
+    // stream the caller uses for kernel launches.  Required for Step 8
+    // (CUDA Graph capture): pool allocations on a different stream than
+    // the capture stream would invalidate the capture.
+    ds4_pool_set_stream(stream);
+
     // 1. Quantize X into CANONICAL Q8_1 (NOT the MMQ-interleaved variant).
     //    Layout: [ne13=1, ne12=n_tokens, ne11=1, ne10_padded blocks].
     const int64_t ne10_padded = GGML_PAD((int64_t)K, MATRIX_ROW_PADDING);
@@ -856,6 +862,10 @@ int ds4_mmq_moe_pair_vec_impl(
         return -1;
     }
 
+    // Route the pool's cudaMallocAsync through the caller-supplied stream
+    // for Step 8 / CUDA Graph compatibility.  See ds4_mmq_moe_vec_impl.
+    ds4_pool_set_stream(stream);
+
     const int n_tokens = 1;  // fusion only supported at ncols_dst=1.
 
     // Quantize X (single token) into canonical Q8_1.
@@ -958,6 +968,10 @@ int ds4_mmq_dense_vec_impl(
         return -1;
     }
 
+    // Route the pool's cudaMallocAsync through the caller-supplied stream
+    // for Step 8 / CUDA Graph compatibility.  See ds4_mmq_moe_vec_impl.
+    ds4_pool_set_stream(stream);
+
     // Dense: no MoE, ids=null. Layout [K, N, 1, 1] for src1.
     const int64_t ne10_padded = GGML_PAD((int64_t)K, MATRIX_ROW_PADDING);
     const size_t  nbytes_q8_1 = (size_t)N * ne10_padded *

cuda/mmq/ds4_mmq.h

@@ -338,6 +338,13 @@ int ds4_mmq_q8_0_dense_vec(
     int           K,
     cudaStream_t  stream);
 
+// Set the thread-local stream that the internal cuda pool uses for
+// cudaMallocAsync / cudaFreeAsync.  Defaults to cudaStreamPerThread.
+// Step 8 (CUDA Graphs) calls this with the capture stream so pool
+// allocations land on the captured stream and don't invalidate capture.
+// Pass NULL to reset to cudaStreamPerThread.
+void ds4_pool_set_stream(cudaStream_t stream);
+
 #ifdef __cplusplus
 } // extern "C"
 #endif

ds4_cuda.cu

@@ -6639,6 +6639,113 @@ extern "C" int ds4_gpu_dsv4_topk_mask_tensor(
  * mmq on.  Cached on first use. */
 static int g_ds4_use_mmq_init = 0;
 static int g_ds4_use_mmq = 0;
+/* -----------------------------------------------------------------------
+ * Step 8: CUDA Graph cache for the routed-MoE decode block.
+ *
+ * The mmvq decode branch (Step 6) launches ~8 kernels per MoE layer.  At
+ * decode (n_tokens=1) each layer's kernel sequence is repeated identically
+ * across the model's 43 layers, and the BUFFER POINTERS are stable across
+ * forward passes - the model is mmapped and ds4 reuses the same scratch
+ * tensors every call.  This makes the kernel sequence a perfect candidate
+ * for cudaGraph capture+replay: capture once per (layer-shape, weight-offset)
+ * tuple, then on every subsequent call replay the captured graph in a
+ * single cudaGraphLaunch.  Each replay eliminates ~8 CPU<->driver round
+ * trips, worth ~5-15us per launch.  At 350+ launches per token and a 23ms
+ * total token time today, eliminating launch overhead is the largest
+ * remaining gen-tok/s lever.
+ *
+ * Cache key: hash of all input pointers, weight offsets, and shape ints
+ * that affect the captured graph's bound parameters.  If any pointer or
+ * shape differs, we capture a new graph (cheap one-time cost).
+ *
+ * Capture mode: cudaStreamCaptureModeRelaxed on stream=0.  Relaxed is
+ * the permissive mode - tolerates cross-stream ops that strict Global
+ * would reject.  Some pool internals (cudaMallocAsync on the per-thread
+ * default stream) live outside the captured stream; in Relaxed mode they
+ * still complete normally and the cached pool returns the same memory
+ * on subsequent calls.  The captured graph thus contains the kernel
+ * launches and the alloc/free nodes for the malloc-async calls in scope.
+ *
+ * Opt-in: DS4_CUDA_MOE_GRAPHS=1 (default off until proven).
+ * --------------------------------------------------------------------- */
+
+struct moe_graph_key {
+    uint64_t gate_offset;
+    uint64_t up_offset;
+    uint64_t down_offset;
+    uint32_t n_tokens;
+    uint32_t q4k_path;
+    uint32_t expert_in_dim;
+    uint32_t expert_mid_dim;
+    uint32_t out_dim;
+    void *gate_ptr;
+    void *up_ptr;
+    void *mid_ptr;
+    void *down_ptr;
+    void *out_ptr;
+    void *x_ptr;
+    void *sel_ptr;
+    void *w_ptr;
+};
+
+struct moe_graph_entry {
+    struct moe_graph_key key;
+    cudaGraphExec_t exec;
+    int valid;
+    uint64_t hits;
+};
+
+#define DS4_MOE_GRAPH_CACHE_SIZE 256
+
+static struct moe_graph_entry g_moe_graphs[DS4_MOE_GRAPH_CACHE_SIZE];
+
+/* Explicit stream used for the captured routed_moe_launch sequence.
+ * Lazily created at first use.  Must be non-default so capture is legal. */
+static cudaStream_t g_moe_stream = NULL;
+
+static cudaStream_t ds4_cuda_moe_stream(void) {
+    if (!g_moe_stream) {
+        cudaError_t ge = cudaStreamCreateWithFlags(&g_moe_stream, cudaStreamNonBlocking);
+        if (ge != cudaSuccess) {
+            fprintf(stderr, "ds4: cudaStreamCreate (moe) failed: %s\n",
+                    cudaGetErrorString(ge));
+            g_moe_stream = NULL;
+        }
+    }
+    return g_moe_stream;
+}
+
+static int ds4_cuda_moe_graphs_enabled(void) {
+    static int init = 0;
+    static int enabled = 0;
+    if (!init) {
+        init = 1;
+        const char *s = getenv("DS4_CUDA_MOE_GRAPHS");
+        if (s && *s && strcmp(s, "0") != 0 &&
+            strcmp(s, "off") != 0 && strcmp(s, "no") != 0 && strcmp(s, "false") != 0) {
+            enabled = 1;
+            fprintf(stderr, "ds4: DS4_CUDA_MOE_GRAPHS enabled - capturing routed-MoE decode\n");
+        }
+    }
+    return enabled;
+}
+
+static uint64_t moe_graph_hash(const struct moe_graph_key *k) {
+    /* FNV-1a over the key bytes. */
+    uint64_t h = 0xcbf29ce484222325ULL;
+    const uint8_t *p = (const uint8_t *)k;
+    for (size_t i = 0; i < sizeof(*k); i++) {
+        h ^= p[i];
+        h *= 0x100000001b3ULL;
+    }
+    return h;
+}
+
+static struct moe_graph_entry *moe_graph_slot(const struct moe_graph_key *key) {
+    uint64_t h = moe_graph_hash(key);
+    return &g_moe_graphs[h % DS4_MOE_GRAPH_CACHE_SIZE];
+}
+
 static int ds4_cuda_use_mmq() {
     if (!g_ds4_use_mmq_init) {
         g_ds4_use_mmq_init = 1;
@@ -11319,6 +11426,68 @@ static int routed_moe_launch(
         if (n_assignments > 8u) {
             /* Outside the mmvq vec path's batch envelope; fall through. */
         } else {
+            /* Step 8: graph cache fast path.  If enabled, check the cache
+             * for a captured graph matching the current shape + pointers.
+             * On hit: replay (cudaGraphLaunch).  On miss: capture the
+             * kernel sequence below and store the exec for next time. */
+            struct moe_graph_entry *graph_slot = NULL;
+            int graph_capturing = 0;
+            cudaStream_t moe_stream = ds4_cuda_moe_graphs_enabled() ? ds4_cuda_moe_stream() : (cudaStream_t)0;
+            if (ds4_cuda_moe_graphs_enabled() && moe_stream) {
+                struct moe_graph_key key;
+                memset(&key, 0, sizeof(key));
+                key.gate_offset     = gate_offset;
+                key.up_offset       = up_offset;
+                key.down_offset     = down_offset;
+                key.n_tokens        = n_tokens;
+                key.q4k_path        = (uint32_t)q4k_path;
+                key.expert_in_dim   = expert_in_dim;
+                key.expert_mid_dim  = expert_mid_dim;
+                key.out_dim         = out_dim;
+                key.gate_ptr        = gate->ptr;
+                key.up_ptr          = up->ptr;
+                key.mid_ptr         = mid->ptr;
+                key.down_ptr        = down->ptr;
+                key.out_ptr         = out->ptr;
+                key.x_ptr           = x->ptr;
+                key.sel_ptr         = selected->ptr;
+                key.w_ptr           = weights->ptr;
+                graph_slot = moe_graph_slot(&key);
+                if (graph_slot->valid &&
+                    memcmp(&graph_slot->key, &key, sizeof(key)) == 0) {
+                    /* HIT: replay the cached graph and return. */
+                    cudaError_t ge = cudaGraphLaunch(graph_slot->exec, moe_stream);
+                    if (ge != cudaSuccess) {
+                        fprintf(stderr, "ds4: cudaGraphLaunch failed: %s; recapturing\n",
+                                cudaGetErrorString(ge));
+                        cudaGraphExecDestroy(graph_slot->exec);
+                        graph_slot->valid = 0;
+                        /* fall through to capture path below */
+                    } else {
+                        graph_slot->hits++;
+                        return 1;
+                    }
+                }
+                /* MISS: begin capture, run the sequence, then end capture
+                 * and instantiate.  Save the exec at graph_slot. */
+                memcpy(&graph_slot->key, &key, sizeof(key));
+                if (graph_slot->valid) {
+                    /* Hash collision with a different key - tear down the
+                     * old exec before recapturing for the new one. */
+                    cudaGraphExecDestroy(graph_slot->exec);
+                    graph_slot->valid = 0;
+                    graph_slot->hits = 0;
+                }
+                cudaError_t ge = cudaStreamBeginCapture(moe_stream, cudaStreamCaptureModeThreadLocal);
+                if (ge == cudaSuccess) {
+                    graph_capturing = 1;
+                } else {
+                    fprintf(stderr, "ds4: cudaStreamBeginCapture failed: %s; graphs disabled this call\n",
+                            cudaGetErrorString(ge));
+                    graph_slot = NULL;
+                }
+            }
+
             int rc = -1;
             /* 1. Two separate gate/up matmuls through mmvq.  Each call
              *    re-quantizes the activation - acceptable for n_tokens=1
@@ -11330,50 +11499,50 @@ static int routed_moe_launch(
                                           (float *)gate->ptr,
                                           (int)expert_mid_dim, (int)expert_in_dim,
                                           (int)n_tokens, (int)n_experts_total,
-                                          (int)n_expert_used, /*stream=*/0);
+                                          (int)n_expert_used, moe_stream);
                 if (rc != 0) {
                     fprintf(stderr, "ds4: ds4_mmq_q4_K_moe_vec (gate) returned %d; falling back\n", rc);
-                    goto mmq_moe_fallback;
+                    goto mmvq_decode_bail;
                 }
                 rc = ds4_mmq_q4_K_moe_vec(up_w, (const float *)x->ptr,
                                           (const int32_t *)selected->ptr,
                                           (float *)up->ptr,
                                           (int)expert_mid_dim, (int)expert_in_dim,
                                           (int)n_tokens, (int)n_experts_total,
-                                          (int)n_expert_used, /*stream=*/0);
+                                          (int)n_expert_used, moe_stream);
             } else {
                 rc = ds4_mmq_iq2_xxs_moe_vec(gate_w, (const float *)x->ptr,
                                              (const int32_t *)selected->ptr,
                                              (float *)gate->ptr,
                                              (int)expert_mid_dim, (int)expert_in_dim,
                                              (int)n_tokens, (int)n_experts_total,
-                                             (int)n_expert_used, /*stream=*/0);
+                                             (int)n_expert_used, moe_stream);
                 if (rc != 0) {
                     fprintf(stderr, "ds4: ds4_mmq_iq2_xxs_moe_vec (gate) returned %d; falling back\n", rc);
-                    goto mmq_moe_fallback;
+                    goto mmvq_decode_bail;
                 }
                 rc = ds4_mmq_iq2_xxs_moe_vec(up_w, (const float *)x->ptr,
                                              (const int32_t *)selected->ptr,
                                              (float *)up->ptr,
                                              (int)expert_mid_dim, (int)expert_in_dim,
                                              (int)n_tokens, (int)n_experts_total,
-                                             (int)n_expert_used, /*stream=*/0);
+                                             (int)n_expert_used, moe_stream);
             }
             if (rc != 0) {
                 fprintf(stderr, "ds4: ds4_mmq_<>_moe_vec (up) returned %d; falling back\n", rc);
-                goto mmq_moe_fallback;
+                goto mmvq_decode_bail;
             }
 
             /* 2. SwiGLU + clamp + router_weight.  Same kernel as the mmq
              *    path uses - applies the V4 clamp to gate/up BEFORE silu. */
             {
                 const uint64_t mid_floats = n_assignments * expert_mid_dim;
-                moe_mmq_swiglu_weighted_clamp_kernel<<<(uint32_t)((mid_floats + 255) / 256), 256>>>(
+                moe_mmq_swiglu_weighted_clamp_kernel<<<(uint32_t)((mid_floats + 255) / 256), 256, 0, moe_stream>>>(
                     (float *)mid->ptr, /*gate_out_dbg=*/nullptr, /*up_out_dbg=*/nullptr,
                     (const float *)gate->ptr, (const float *)up->ptr,
                     (const float *)weights->ptr,
                     expert_mid_dim, n_tokens, n_expert_used, clamp);
-                if (!cuda_ok(cudaGetLastError(), "mmvq routed_moe swiglu launch")) goto mmq_moe_fallback;
+                if (!cuda_ok(cudaGetLastError(), "mmvq routed_moe swiglu launch")) goto mmvq_decode_bail;
             }
 
             /* 3. Down matmul: same reinterpretation trick the mmq path uses -
@@ -11386,29 +11555,86 @@ static int routed_moe_launch(
                                           (float *)down->ptr,
                                           (int)out_dim, (int)expert_mid_dim,
                                           (int)n_assignments, (int)n_experts_total,
-                                          /*n_expert_used=*/1, /*stream=*/0);
+                                          /*n_expert_used=*/1, moe_stream);
             } else {
                 rc = ds4_mmq_q2_K_moe_vec(down_w, (const float *)mid->ptr,
                                           (const int32_t *)selected->ptr,
                                           (float *)down->ptr,
                                           (int)out_dim, (int)expert_mid_dim,
                                           (int)n_assignments, (int)n_experts_total,
-                                          /*n_expert_used=*/1, /*stream=*/0);
+                                          /*n_expert_used=*/1, moe_stream);
             }
             if (rc != 0) {
                 fprintf(stderr, "ds4: ds4_mmq_<>_moe_vec (down) returned %d; falling back\n", rc);
-                goto mmq_moe_fallback;
+                goto mmvq_decode_bail;
             }
 
-            /* 4. Sum across n_expert_used dim - same kernel as the mmq path. */
+            /* 4. Sum across n_expert_used dim - same kernel as the mmq path.
+             *    Launched on moe_stream so Step 8 capture is consistent. */
             {
                 uint64_t n = (uint64_t)n_tokens * out_dim;
-                moe_sum_kernel<<<(uint32_t)((n + 255) / 256), 256>>>(
+                moe_sum_kernel<<<(uint32_t)((n + 255) / 256), 256, 0, moe_stream>>>(
                     (float *)out->ptr, (const float *)down->ptr,
                     out_dim, n_expert_used, n_tokens);
-                if (!cuda_ok(cudaGetLastError(), "mmvq routed_moe sum launch")) goto mmq_moe_fallback;
+                if (!cuda_ok(cudaGetLastError(), "mmvq routed_moe sum launch")) goto mmvq_decode_bail;
+            }
+            /* Step 8: if we were capturing, finalize the graph and replay
+             * it.  cudaStreamEndCapture builds the graph object; the
+             * captured kernels have NOT executed yet (stream capture
+             * records but doesn't run).  cudaGraphInstantiate builds the
+             * exec; cudaGraphLaunch fires it.  Subsequent calls with the
+             * same key replay via cudaGraphLaunch at the top of this
+             * branch, avoiding the ~8 per-kernel CPU<->driver round
+             * trips. */
+            if (graph_capturing && graph_slot) {
+                cudaGraph_t graph;
+                cudaError_t ge = cudaStreamEndCapture(moe_stream, &graph);
+                graph_capturing = 0;
+                if (ge == cudaSuccess) {
+                    cudaGraphExec_t exec;
+                    ge = cudaGraphInstantiate(&exec, graph, NULL, NULL, 0);
+                    if (ge == cudaSuccess) {
+                        graph_slot->exec = exec;
+                        graph_slot->valid = 1;
+                        graph_slot->hits = 0;
+                        /* Execute the just-captured graph for THIS call.
+                         * Capture itself does not run the kernels; we must
+                         * launch the exec to actually do the work. */
+                        ge = cudaGraphLaunch(exec, moe_stream);
+                        if (ge != cudaSuccess) {
+                            fprintf(stderr, "ds4: cudaGraphLaunch (first) failed: %s\n",
+                                    cudaGetErrorString(ge));
+                        }
+                    } else {
+                        fprintf(stderr, "ds4: cudaGraphInstantiate failed: %s; graphs disabled this layer\n",
+                                cudaGetErrorString(ge));
+                    }
+                    cudaGraphDestroy(graph);
+                } else {
+                    fprintf(stderr, "ds4: cudaStreamEndCapture failed: %s\n",
+                            cudaGetErrorString(ge));
+                }
             }
             return 1;
+        mmvq_decode_bail:
+            /* Failure inside the mmvq decode branch: end any in-flight
+             * capture cleanly (discard the partial graph) before falling
+             * through to the mmq path. */
+            if (graph_capturing) {
+                cudaGraph_t partial;
+                cudaError_t ge = cudaStreamEndCapture(0, &partial);
+                if (ge == cudaSuccess) {
+                    cudaGraphDestroy(partial);
+                } else {
+                    /* Stream may be left in a bad capture state; log so we
+                     * notice but proceed - the mmq fallback will reset. */
+                    fprintf(stderr, "ds4: cudaStreamEndCapture (bail) failed: %s\n",
+                            cudaGetErrorString(ge));
+                }
+                graph_capturing = 0;
+                if (graph_slot) graph_slot->valid = 0;
+            }
+            goto mmq_moe_fallback;
         }
     }