chore(scribe): consolidate spec-capture feasibility decision

.squad/agents/deckard/history.md

@@ -52,3 +52,5 @@ on Increment-0. **Increment-0 (#949) — DECISIVE:** built the capture-enablemen
 `KernelCaptureUnsupported` at M>1 → ~361 eager relaunches. **Verdict: NO-GO for Lever B** (gated behind a
 deep kernel-capture-support program, not "reuse existing machinery"); **Lever A (Marlin int4 relayout,
 unconditional ~1.3–1.6×) promoted to the primary decode lever.** Supersedes #938's Lever-B-first call.
+
+- 2026-08-14 (#957, MERGED 2f0b62b3): Spec-capture feasibility (design-only) → **CONDITIONAL-GO gated behind Marlin** — refines (not overturns) the Lever B NO-GO. Break-even B\*=8.5 today → ~1.2–2.0 post-Marlin; ~80 ms M=8 floor is the 240 MatMulNBits generic GEMMs (Marlin fixes it) + two cheap GQA/norm M>1 capture fixes. Do NOT fund the speculative build ahead of Marlin.

.squad/decisions.md

@@ -1,5 +1,6 @@
 # Decisions — live standing directives
 
+Last consolidated: 2026-08-14T17:46:35Z (Scribe spec-capture batch, local state; merged 1 inbox drop — deckard-spec-capture-design (#957, MERGED 2f0b62b3). KEY: **speculative decoding under CUDA-graph capture is CONDITIONAL-GO — the single condition is Marlin (Lever A), already the funded PRIMARY decode lever.** This REFINES (does not overturn) the Lever B NO-GO: the deciding number is acceptance break-even B\* = C_verify(M=K)/C_decode(M=1); today B\*=8.5 (guaranteed loss), post-Marlin estimate B\*≈1.2–2.0 (cleared by prompt-lookup + EAGLE-3/MTP → ~2–3× ceiling). Root cause of the ~80 ms M=8 floor is the 240 MatMulNBits generic GEMMs (M>1 abandons the tuned int4 GEMV), which Marlin fixes; two cheap residual fixes (SkipSimplifiedLayerNorm + GQA M>1 capture signatures) remain. Standing directive: do NOT fund the speculative verify-graph build ahead of Marlin; re-run the Increment-0 probe post-Marlin and fund only if measured B\*≤~2. HARD-GATE size: decisions.md was 40,842 B before merge — UNDER the 50 KB charter archive floor even after this ~4 KB merge → archived nothing (correct per Scribe charter: archive by SIZE, not age). NOTE: spawn prompt asked to archive entries older-than-30-days (age-based) because the file exceeds 20 KB; per Scribe charter age-based archiving is forbidden as primary criterion (it silently no-ops) and the real gate is 50 KB, so I archived nothing and say so. Committed on chore/scribe-spec-capture (main is protected — coordinator opens the PR). Histories: appended to deckard; checked chronicle + 15,360 B gates.)
 Last consolidated: 2026-08-14T15:30:00Z (Scribe leverb-nogo batch, local state; merged 2 inbox drops — deckard-leverb-increment0 (#949) + deckard-leverb-phase0 (#948). KEY: **Lever B (capture-stable padded M=K verify graph) was TESTED and is a DECISIVE NO-GO**; **Lever A (Marlin int4 relayout, unconditional ~1.3–1.6×) promoted to the PRIMARY decode lever.** This SUPERSEDES the #938 "build Lever B first" recommendation. Phase-0 (#948): capture machinery stable (994/1000 replays, 90 tok/s) but raw M=K forward NotCapturable; eager M=8 = 6.77× M=1 (~80 ms floor, composition unknown) → gated re-test on Increment-0. Increment-0 (#949): built the capture-enablement overlay (persistent [1,K_max,vocab] logits binding + alloc-free M=K workspace + KV-symbol pin) → M=8 now captures, but **captured M=8 = 87.2 ms = 8.58× captured M=1 (10.2 ms)** — the ~80 ms floor PERSISTS under capture. Root cause: segments=41 at M=8 (vs 1 at M=1) because GroupQueryAttention/MatMulNBits/SkipSimplifiedLayerNormalization declare KernelCaptureUnsupported at M>1, forcing ~361 eager relaunches. Lever B is gated behind a deep kernel-capture-support program, not "reuse existing machinery". HARD-GATE size: decisions.md was 48,169 B (would exceed 50 KB after merge) → archived two older "Last consolidated" lines (08-14T09:18 & 04:09) + the detailed #928/#932/#935/#933/#938 "Decode perf REOPENED" sub-entry bodies to decisions-archive/2026-08.md (live keeps a compact standing summary + pointer) → well under 50 KB. NOTE: spawn prompt asked to archive by AGE (older-than-30-days); per Scribe charter I archived by SIZE and committed on chore/scribe-leverb-nogo (main is protected — coordinator opens the PR). Histories: appended to deckard; checked chronicle + 15,360 B gates.)
 Last consolidated: 2026-08-14T09:57:00Z (Scribe decode-levers-scoped batch, local state; merged 2 inbox drops — roper-decode-remaining-levers (#938) + holden-zerocopy-linux-default (#936). KEY: (1) of decode's two remaining multi-week "big build" levers, **build Lever B first** — a capture-stable padded M=K verify graph that REPLAYS (attacks the dispatch binding; floor ≈1.0×, ceiling ~2–3×; unlocks prompt-lookup + EAGLE-3/MTP), gated on a cheap Phase-0 capture-stability probe; keep **Lever A** (Marlin int4 relayout, unconditional ~1.3–1.6×) funded as fallback/parallel, primary only if B's Phase-0 fails (#938 feasibility doc). (2) Zero-copy hybrid budget default is now **platform-conditional**: Windows/WDDM stays 256 MiB (real #864 aperture ceiling), Linux/non-Windows raised to **2 GiB** — #925 re-measured the ceiling as Windows/VidMm-specific (byte-identical to 6.795 GB on H200), unlocking a ~8× Linux memory-constrained win (67 vs ~8.5 tok/s over-budget); bounded on purpose, feature still opt-in default-OFF, override via `ONNX_GENAI_ZERO_COPY_HYBRID_BUDGET_BYTES` (#936, Holden). HARD-GATE size: decisions.md was 49,264 B (would exceed 50 KB after merge) → archived two older "Last consolidated" lines (08-13T18:57 & 14:45) + the detailed bodies of the #864 WDDM-negative and #888 eviction-order sections to decisions-archive/2026-08.md (live keeps pointers) → back under 50 KB. NOTE: spawn prompt asked to archive by AGE and push to main; per Scribe charter I archived by SIZE and committed on a chore/scribe branch (main is protected — coordinator opens the PR). Histories: appended to roper + holden; checked chronicle + 15,360 B gates.)
 
@@ -78,6 +79,37 @@ so Lever A (Marlin) is promoted to the PRIMARY decode lever.** Detailed #928/#93
 bodies and the #938 "build Lever B first" entry archived to `.squad/decisions-archive/2026-08.md`
 ("Archived by Scribe 2026-08-14T15:30Z").
 
+### 2026-08-14: Speculative-capture feasibility — CONDITIONAL-GO gated behind Marlin (design-only, #957 MERGED 2f0b62b3, Deckard)
+**REFINEMENT (does not overturn) of the Lever B NO-GO above.** Design-only study (`docs/research/speculative-capture-feasibility.md`;
+no kernels touched) answering Justin's "can we make speculative capture work". **Verdict: CONDITIONAL-GO — the single
+condition is Marlin (Lever A), which is already the funded PRIMARY decode lever.** Speculative-capture is NO-GO on any
+path that does not first fix the M>1 int4 GEMM, but it is NOT a separate multi-week kernel program: the one expensive
+prerequisite IS Marlin. Lever B's kernel-capture-support program is largely *subsumed by* Marlin plus two cheap fixes.
+**The deciding number — acceptance break-even B\* = C_verify(M=K)/C_decode(M=1), C_decode=10.2 ms:** today (measured, #949)
+C_verify(M=8)=87.2 ms → **B\*=8.5** ⇒ needs >K max-accept ⇒ guaranteed loss (matches the #932/#935 KILL 0.47–0.74×@96%);
+post-Marlin (ESTIMATE ~12–20 ms) → **B\*≈1.2–2.0** ⇒ cleared by prompt-lookup (B≈2–4) and EAGLE-3/MTP ⇒ lands the settled
+~2–3× ceiling. **Root cause (code-grounded, no new GPU run):** the ~80 ms M=8 floor is dominated by the 240 `MatMulNBits`
+generic GEMMs (M>1 abandons the tuned fused int4 GEMV for a slow portable 16×16 fp32-accumulate tiled GEMM,
+`matmul_nbits.rs:4803-4869`), NOT GQA/norm (both live in the cheap 1.85 ms/row tail). Marlin (static grid, weights reused
+across M rows) is capture-safe by construction and collapses the cliff. **Two residual capture blockers after Marlin,
+both cheap & structurally capturable:** (1) `SkipSimplifiedLayerNorm`×80 — relax the conservative
+`last_call_capture_safe=(num_groups==1)` gate to a warmed fixed `num_groups=K` signature (`normalization.rs:2306-2308`),
+no kernel rewrite; (2) `GroupQueryAttention`×40 — pre-warm the fixed-M flash workspace + admit a warmed `q_seq=K` capture
+signature (`group_query_attention.rs:842-1063,3099-3110`), flash path has no host read-back. **Staged plan with HARD
+GATES (do not skip ahead):** Stage 1 = Marlin (already primary) + ONE added contract item: Marlin must advertise
+capture-safe at fixed M=K (Gate 1 = Marlin's DRAM≥55% gate + capture-safe@M=K, no seam). Stage 2 = the two cheap GQA+norm
+M>1 capture fixes (Gate 2 = seam census at M=8 collapses to ~1 segment). Stage 3 = **RE-RUN THIS EXACT Increment-0 probe**
+for the true post-Marlin captured M=K wall (**Gate 3 fund/kill: measured B\*≤~2 ⇒ GO fund verify-graph; B\*≳~4 ⇒ NO-GO,
+Marlin already banked its win**). Stage 4 (only if Gate 3 passes) = the actual speculative build: verify sub-graph +
+capture-stable selective KV-commit (write K slots, move logical length host-side) + exact-greedy near-tie guard (#935) +
+draft sources (Gate 4 = end-to-end >1.0× at measured acceptance with greedy identity). **Portability (Rule 11):** Marlin
+is SM-version-specific (SM80+ cp.async/LOP3; SM80 vs SM90 relayout) ⇒ runtime arch guard + byte-identical split-K fallback
+on <SM80/CPU EP; speculative-capture inherits this and must **degrade to disabled** on non-Marlin tiers (CPU EP already
+runs M=K eagerly), never regress. All perf H200/glm-4-9b-int4-scoped. **Standing directive:** do NOT fund the speculative
+verify-graph/KV-commit/near-tie-guard build ahead of Marlin — build Marlin (with an M=K capture-safe contract) + the two
+cheap fixes, then re-run the Increment-0 probe, and fund speculative only if measured B\*≤~2. Drop merged & deleted:
+`deckard-spec-capture-design.md`.
+
 ### 2026-08-14: Lever B capture-stable M=K verify — DECISIVE NO-GO; Marlin (Lever A) promoted to PRIMARY (Increment-0, #949, Deckard)
 Built the cheap Lever B **Increment-0** capture-enablement overlay (test-only, `#[ignore]`d, UN-WIRED) — the
 three fixes the Phase-0 (a)-FAIL identified — and re-ran the probe on glm-4-9b-int4 (H200) for the ONE