cuda: step 8.2 - graphs for dense Q8_0 vec + flip default to ON

Two related changes that land together:

1. cudaGraph capture+replay for the n_tok=1 dense Q8_0 mmvq path
   (cuda_matmul_q8_0_tensor_labeled).  Each per-projection
   (weight_offset, x_ptr, out_ptr) tuple gets its own cached
   cudaGraphExec_t.  For V4 Flash decode there are ~5 such
   projections per layer (q/k/v + output + hc-expand) * 43 layers =
   ~215 cache keys per gen.  Cache sized at 1024 entries to fit
   comfortably and tolerate routine MTP/decode-2 add-ons.  Same
   g_moe_stream + ds4_pool_set_stream infrastructure as the
   routed-MoE graph branch from Step 8.

2. DS4_CUDA_MOE_GRAPHS default flipped from off to ON after
   determinism validation:
   - Round 1 (Step 8 v1, MoE-only): 10/10 greedy decode runs (128
     tokens, temp=0) bit-identical vs graphs-OFF baseline.
   - Round 2 (with Step 8.2 dense capture added): 10/10 runs
     bit-identical again.
   - Opt-out preserved: DS4_CUDA_MOE_GRAPHS=0 (or off/no/false)
     restores the un-captured paths.

Combined effect on PRO 6000 Blackwell (sm_120), ctx=2048,
gen-tokens=128, N=5 runs per config:

  - graphs OFF (DS4_CUDA_MOE_GRAPHS=0): gen 43.15 tok/s (IQR 0.03)
  - graphs ON  (default after this commit): gen 47.01 tok/s (IQR 0.05)
  - delta: gen +8.95% (Mann-Whitney U p=0.0079, B wins)
  - prefill: 1009.58 -> 1009.43 (-0.01%, p=1.00, inconclusive)

Step 8.2 adds ~+4.4% gen on top of Step 8 v1 (which was MoE-only at
~+4.4% on top of Step 6 mmvq).  Combined Step 6 + 8 + 8.2 lands gen
at 47.01 tok/s vs the pre-lift baseline of 38.0 = +23.7% cumulative.

AGENT.md updated:
  - DS4_CUDA_MOE_GRAPHS doc reflects default-ON status and the
    Step 8.2 dense Q8_0 cache.

Notes:
  - Step 8.2 reuses g_moe_stream across MoE and dense captures.  A
    single capture stream simplifies synchronisation with the rest
    of the forward pass (which still uses stream=0; legacy default
    stream serialises with all named streams so no explicit
    cudaStreamWaitEvent is needed).
  - The dense graph cache is structurally identical to the MoE one
    (FNV-1a hash, direct-mapped, collision-overwrites prior exec),
    just with a smaller key type and a larger array.
  - Both caches share the DS4_CUDA_MOE_GRAPHS opt-out switch -
    consumers don't need a separate knob for the dense path.

AGENT.md

@@ -113,21 +113,25 @@ 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
+- `DS4_CUDA_MOE_GRAPHS` (default ON after 2026-05-16 determinism
+  validation - 10/10 greedy decode runs bit-identical to graphs-off
+  baseline): CUDA Graph capture+replay for the mmvq routed-MoE decode
+  block and (Step 8.2) for the n_tok=1 dense Q8_0 vec path used by
+  attention projections.  Each kernel sequence is captured into a
+  `cudaGraphExec_t` on first execution with a given
+  (layer-shape, buffer-pointer) tuple and replayed via
+  `cudaGraphLaunch` on subsequent calls.  The MoE cache holds 256
+  entries (one per layer shape-class); the dense Q8_0 cache holds
+  1024 entries (one per attention projection per layer).  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`,
+  shared across both caches) 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.
+  Each cache slot falls through cleanly to the un-captured path on
+  any capture error.  Opt-out: `DS4_CUDA_MOE_GRAPHS=0` (or `off`,
+  `no`, `false`).
 
 ## Testing
 

ds4_cuda.cu

@@ -6716,15 +6716,21 @@ static cudaStream_t ds4_cuda_moe_stream(void) {
 }
 
 static int ds4_cuda_moe_graphs_enabled(void) {
+    /* Default ON after determinism validation (10/10 greedy decode runs
+     * bit-identical vs graphs OFF, 2026-05-16).  Opt-out:
+     *   DS4_CUDA_MOE_GRAPHS=0  (or off / no / false) */
     static int init = 0;
-    static int enabled = 0;
+    static int enabled = 1;
     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");
+        if (s && *s &&
+            (strcmp(s, "0") == 0 ||
+             strcmp(s, "off") == 0 || strcmp(s, "OFF") == 0 ||
+             strcmp(s, "no") == 0 || strcmp(s, "NO") == 0 ||
+             strcmp(s, "false") == 0 || strcmp(s, "FALSE") == 0)) {
+            enabled = 0;
+            fprintf(stderr, "ds4: DS4_CUDA_MOE_GRAPHS=%s - graph capture disabled\n", s);
         }
     }
     return enabled;
@@ -6746,6 +6752,45 @@ static struct moe_graph_entry *moe_graph_slot(const struct moe_graph_key *key) {
     return &g_moe_graphs[h % DS4_MOE_GRAPH_CACHE_SIZE];
 }
 
+/* Step 8.2: dense Q8_0 vec graph cache.  Each n_tok=1 attention-side
+ * matmul (q/k/v/output projections and HC-expand variants) becomes its
+ * own cached cudaGraphExec_t.  V4 Flash has ~5 such projections per
+ * layer * 43 layers = ~215 keys per gen.  Sized at 1024 to fit
+ * comfortably and tolerate routine MTP/decode-2 add-ons. */
+struct dense_graph_key {
+    uint64_t weight_offset;
+    uint32_t in_dim;
+    uint32_t out_dim;
+    uint32_t n_tok;
+    void *x_ptr;
+    void *out_ptr;
+};
+
+struct dense_graph_entry {
+    struct dense_graph_key key;
+    cudaGraphExec_t exec;
+    int valid;
+    uint64_t hits;
+};
+
+#define DS4_DENSE_GRAPH_CACHE_SIZE 1024
+
+static struct dense_graph_entry g_dense_graphs[DS4_DENSE_GRAPH_CACHE_SIZE];
+
+static uint64_t dense_graph_hash(const struct dense_graph_key *k) {
+    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 dense_graph_entry *dense_graph_slot(const struct dense_graph_key *key) {
+    return &g_dense_graphs[dense_graph_hash(key) % DS4_DENSE_GRAPH_CACHE_SIZE];
+}
+
 static int ds4_cuda_use_mmq() {
     if (!g_ds4_use_mmq_init) {
         g_ds4_use_mmq_init = 1;
@@ -6788,12 +6833,95 @@ static int cuda_matmul_q8_0_tensor_labeled(ds4_gpu_tensor *out, const void *mode
      * first; on failure falls through to mmq dense, which itself falls
      * through to the legacy native kernels at the bottom of this function.
      *
-     * Opt-out: DS4_CUDA_NO_MMVQ_DECODE=1 (same flag as routed_moe_launch). */
+     * Opt-out: DS4_CUDA_NO_MMVQ_DECODE=1 (same flag as routed_moe_launch).
+     *
+     * Step 8.2: when DS4_CUDA_MOE_GRAPHS=1 also enables graph capture
+     * around this branch.  Each (weight_offset, x_ptr, out_ptr) tuple
+     * gets its own cached cudaGraphExec_t.  On cache hit, a single
+     * cudaGraphLaunch replaces the alloc + quantize + mmvq + free
+     * launches.  Same g_moe_stream + ds4_pool_set_stream plumbing as
+     * the routed-MoE graph branch. */
     if (ds4_cuda_use_mmq() && (in_dim % 256u == 0) && n_tok == 1u &&
         getenv("DS4_CUDA_NO_MMVQ_DECODE") == NULL) {
+
+        struct dense_graph_entry *dslot = NULL;
+        int dcapturing = 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 dense_graph_key dkey;
+            memset(&dkey, 0, sizeof(dkey));
+            dkey.weight_offset = weight_offset;
+            dkey.in_dim        = (uint32_t)in_dim;
+            dkey.out_dim       = (uint32_t)out_dim;
+            dkey.n_tok         = (uint32_t)n_tok;
+            dkey.x_ptr         = x->ptr;
+            dkey.out_ptr       = out->ptr;
+            dslot = dense_graph_slot(&dkey);
+            if (dslot->valid && memcmp(&dslot->key, &dkey, sizeof(dkey)) == 0) {
+                cudaError_t ge = cudaGraphLaunch(dslot->exec, moe_stream);
+                if (ge == cudaSuccess) {
+                    dslot->hits++;
+                    return 1;
+                }
+                fprintf(stderr, "ds4: cudaGraphLaunch (dense) failed: %s; recapturing\n",
+                        cudaGetErrorString(ge));
+                cudaGraphExecDestroy(dslot->exec);
+                dslot->valid = 0;
+            }
+            memcpy(&dslot->key, &dkey, sizeof(dkey));
+            if (dslot->valid) {
+                cudaGraphExecDestroy(dslot->exec);
+                dslot->valid = 0;
+                dslot->hits = 0;
+            }
+            cudaError_t ge = cudaStreamBeginCapture(moe_stream, cudaStreamCaptureModeThreadLocal);
+            if (ge == cudaSuccess) {
+                dcapturing = 1;
+            } else {
+                fprintf(stderr, "ds4: cudaStreamBeginCapture (dense) failed: %s\n",
+                        cudaGetErrorString(ge));
+                dslot = NULL;
+            }
+        }
+
         int rc = ds4_mmq_q8_0_dense_vec(wptr, (const float *)x->ptr, (float *)out->ptr,
                                         (int)out_dim, (int)n_tok, (int)in_dim,
-                                        /*stream=*/0);
+                                        moe_stream);
+
+        if (dcapturing && dslot) {
+            cudaGraph_t graph;
+            cudaError_t ge = cudaStreamEndCapture(moe_stream, &graph);
+            dcapturing = 0;
+            if (ge == cudaSuccess) {
+                if (rc == 0) {
+                    cudaGraphExec_t exec;
+                    ge = cudaGraphInstantiate(&exec, graph, NULL, NULL, 0);
+                    if (ge == cudaSuccess) {
+                        dslot->exec = exec;
+                        dslot->valid = 1;
+                        dslot->hits = 0;
+                        ge = cudaGraphLaunch(exec, moe_stream);
+                        if (ge != cudaSuccess) {
+                            fprintf(stderr, "ds4: cudaGraphLaunch (dense first) failed: %s\n",
+                                    cudaGetErrorString(ge));
+                        }
+                    } else {
+                        fprintf(stderr, "ds4: cudaGraphInstantiate (dense) failed: %s\n",
+                                cudaGetErrorString(ge));
+                    }
+                }
+                cudaGraphDestroy(graph);
+            } else {
+                fprintf(stderr, "ds4: cudaStreamEndCapture (dense) failed: %s\n",
+                        cudaGetErrorString(ge));
+            }
+        } else if (dcapturing) {
+            /* Slot was nulled out for some reason; bail capture cleanly. */
+            cudaGraph_t partial;
+            cudaError_t ge = cudaStreamEndCapture(moe_stream, &partial);
+            if (ge == cudaSuccess) cudaGraphDestroy(partial);
+        }
+
         if (rc == 0) return 1;
         fprintf(stderr, "ds4: ds4_mmq_q8_0_dense_vec returned %d (label='%s' in=%llu out=%llu); falling back to mmq\n",
                 rc, label ? label : "", (unsigned long long)in_dim, (unsigned long long)out_dim);