From 2ee89a14cd6d840b9616a2b462a4ec8055a5eac0 Mon Sep 17 00:00:00 2001 From: Justin Chu Date: Fri, 14 Aug 2026 18:03:56 +0000 Subject: [PATCH] chore(scribe): merge spec-capture-design decision drop (#957) Merge Deckard's speculative-capture feasibility drop into decisions.md as a REFINEMENT of the Lever B NO-GO: speculative-capture is CONDITIONAL-GO gated behind Marlin (B* = 8.5 today -> ~1.2-2.0 post-Marlin). Write orchestration + session logs, append Deckard history. Archived nothing (40,842 B < 50 KB charter floor; age-based archive request declined per charter). Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> --- .squad/agents/deckard/history.md | 2 ++ .squad/decisions.md | 32 +++++++++++++++++ .../inbox/deckard-spec-capture-design.md | 34 ------------------- 3 files changed, 34 insertions(+), 34 deletions(-) delete mode 100644 .squad/decisions/inbox/deckard-spec-capture-design.md diff --git a/.squad/agents/deckard/history.md b/.squad/agents/deckard/history.md index bb60d40eca..e0174a524f 100644 --- a/.squad/agents/deckard/history.md +++ b/.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. diff --git a/.squad/decisions.md b/.squad/decisions.md index 27c93055ef..1871c74256 100644 --- a/.squad/decisions.md +++ b/.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 1 int4 GEMM, but it is NOT a separate multi-week kernel program: the one expensive prerequisite IS Marlin. This refines (does not overturn) the settled Lever B NO-GO (#948/#949): Lever B's kernel-capture-support program is largely *subsumed by* Marlin plus two cheap fixes. - -## The number that decides it — 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 spec 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 attribution (code-grounded, no new GPU run) -The ~80 ms M=8 floor / 67 ms cliff is **dominated by the 240 `MatMulNBits` generic GEMMs**, NOT GQA/norm: -- MatMulNBits at M>1 leaves its tuned fused int4 GEMV for a slow "portable 16×16 CUDA-core fp32-accumulate tiled GEMM" (`matmul_nbits.rs:4803-4869`) — the only family that switches to a fundamentally un-tuned kernel; streams the same ~4.1 GB int4 weights but ~an order of magnitude slower. -- GQA switches to the *efficient* flash-prefill kernel (~168 MB KV read); norm keeps the *same* kernel (num_groups=M). 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. - -## Residual capture blockers after Marlin (ranked; both cheap, both structurally capturable) -1. **SkipSimplifiedLayerNorm ×80** — VERY LOW cost: relax the deliberate `last_call_capture_safe=(num_groups==1)` gate to a warmed fixed `num_groups=K` signature (`normalization.rs:2306-2308`, "left conservative by design"). No kernel rewrite. -2. **GroupQueryAttention ×40** — LOW–MEDIUM: 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. No new attention kernel. - -## Staged plan with hard gates (do NOT skip ahead) -- **Stage 1 — Marlin (already primary; ref `decode-remaining-levers-feasibility.md`).** Add ONE contract item: Marlin must advertise **capture-safe at fixed M=K**. Gate 1 = Marlin's own DRAM≥55% gate + instantiates capture-safe at M=K (no seam). -- **Stage 2 — GQA + norm M>1 capture-support (the two cheap fixes).** Gate 2 = seam census at M=8 collapses to ~1 segment (0 MatMulNBits/GQA/norm KernelCaptureUnsupported). -- **Stage 3 — RE-RUN THIS EXACT Increment-0 probe (PART D)** 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 — the actual speculative build (only if Gate 3 passes):** 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; relayout differs SM80 vs SM90) ⇒ runtime arch guard + byte-identical split-K fallback on