diff --git a/docs/superpowers/plans/2026-04-29-numerical-stability-sp1-f1-nan-root-cause.md b/docs/superpowers/plans/2026-04-29-numerical-stability-sp1-f1-nan-root-cause.md
new file mode 100644
index 000000000..2aaf737cc
--- /dev/null
+++ b/docs/superpowers/plans/2026-04-29-numerical-stability-sp1-f1-nan-root-cause.md
@@ -0,0 +1,1336 @@
+# Numerical Stability SP1 — F1 NaN Root-Cause Implementation Plan
+
+> **For agentic workers:** REQUIRED SUB-SKILL: Use superpowers:subagent-driven-development (recommended) or superpowers:executing-plans to implement this plan task-by-task. Steps use checkbox (`- [ ]`) syntax for tracking.
+
+**Goal:** Resolve the F1 ep2 NaN explosion in Plan C Phase 2 work — identify the kernel/operation producing the first NaN, fix it (with ISV-driven dynamic bounds), validate via multi-fold L40S smoke that all 3 folds train successfully.
+
+**Architecture:** Two-phase investigation (γ read-only audit + β always-landing instrumentation expanding `nan_flags_buf` 24→48), followed by surgical fix(es) (rich commit per `feedback_no_partial_refactor`), validated by the existing `multi_fold_convergence` L40S smoke harness. All work lands as continuation commits on `plan-c-phase-2-thompson` (latest commit `792812baa`). Diagnostic instrumentation stays as permanent regression sentinel.
+
+**Tech Stack:** Rust 1.85+, CUDA 12.4 (sm_89 L40S), cudarc 0.16, Argo Workflows for L40S smoke runs, MinIO (mc CLI) for log artifact retrieval.
+
+**Spec reference:** `docs/superpowers/specs/2026-04-29-numerical-stability-investigation-design.md`.
+
+---
+
+## Operating Principles (read before starting)
+
+1. **No deferrals.** If investigation reveals additional anomalies (e.g., F2 fails for a different reason than F1), fix them within SP1, not SP2/SP3.
+2. **Combined RELATED fixes — rich commits.** Per `feedback_no_partial_refactor`. Multiple kernel patches addressing one investigation thread land as ONE commit. Unrelated concerns stay separate.
+3. **Mandatory ISV-driven design.** Every dynamic bound, threshold, clamp, or floor MUST derive from an existing or new ISV slot. Numerical-stability ε bounds (e.g., `1e-6` divide-by-zero floor) are the only carve-out (Invariant 1). Hardcoded "tuning" constants for dynamic ranges are explicitly rejected. Existing ISV slots available for reuse: `Q_ABS_REF_INDEX = 16`, `Q_DIR_ABS_REF_INDEX = 21`, `H_S2_RMS_EMA_INDEX`, `Q_DRIFT_RATE_INDEX = 129`, `FOLD_WARMUP_FACTOR_INDEX = 130`.
+4. **F0 paper-review BEFORE smoke.** Audit doc carries `F0 risk: low/medium/high` per proposed fix. High-risk fixes must inspect mathematical effect on F0-typical input ranges before consuming an L40S smoke.
+5. **Permanent diagnostic infrastructure.** The β instrumentation (slots 24-35 + 12 reserved) is NOT a temporary scaffold. It stays in the codebase as production-grade regression sentinel.
+
+---
+
+## File Structure
+
+| Path | Status | Responsibility |
+|------|--------|----------------|
+| `docs/dqn-backward-nan-audit.md` | Create | γ audit findings — per-kernel sections with proposed guards, ISV bound options, F0 risk |
+| `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` | Modify | `nan_flags_buf` 24→48; new `run_nan_checks_post_backward` method; backward-buffer accessors as needed |
+| `crates/ml/src/trainers/dqn/fused_training.rs` | Modify | `read_nan_flags` signature `[i32; 24]→[i32; 48]`; new `run_nan_checks_post_backward` trigger point; possibly new accessors |
+| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | Modify | Name table updated to 48 entries in both consumer sites (`halt_nan` block + `halt_grad_collapse` block at commit `d1808df14`) |
+| `crates/ml/src/cuda_pipeline/batched_backward.rs` | Modify (Task 3 only) | New `pub(crate)` accessors for backward intermediate buffers (`bw_d_h_s2`-related, `cql_grad_scratch` via existing accessor, etc.) |
+| TBD kernel files | Modify (Task 6) | Surgical fix lands wherever γ + β identify — `apply_iqn_trunk_gradient` / `c51_grad_kernel.cu` / `cql_grad_kernel.cu` / `aux_heads_kernel.cu` / `iqn_dual_head_kernel.cu` are candidates |
+| `docs/dqn-wire-up-audit.md` | Modify | Pre-commit Invariant 7 audit doc co-stage entries for every commit touching `crates/ml/src/cuda_pipeline/` |
+
+---
+
+## Phase A — γ Audit (read-only)
+
+### Task 1: Produce backward NaN audit document
+
+**Files:**
+- Create: `docs/dqn-backward-nan-audit.md`
+- Modify: `docs/dqn-wire-up-audit.md` (Invariant 7 co-stage entry)
+
+This task is read-only investigation — no Rust/CUDA changes. The output is a complete audit document that drives all downstream instrumentation and fix decisions. The agent must read each kernel file in full and document findings against the unsafe-pattern checklist.
+
+- [ ] **Step 1: Create audit document with template**
+
+```bash
+cat > docs/dqn-backward-nan-audit.md <<'EOF'
+# DQN Backward-Path NaN Audit — SP1 Phase A (γ)
+
+**Status:** ACTIVE — drives SP1 Phase B instrumentation + Phase C fix decisions.
+**Spec:** `docs/superpowers/specs/2026-04-29-numerical-stability-investigation-design.md`
+**Plan:** `docs/superpowers/plans/2026-04-29-numerical-stability-sp1-f1-nan-root-cause.md`
+**Method:** read-only inventory of every backward-path kernel that writes to `bw_d_h_s2` / `grad_buf` / `save_h_s2` accumulators. Each kernel evaluated against unsafe-pattern checklist. Findings drive Phase B flag-slot allocation and Phase C surgical fix(es).
+
+## Unsafe-Pattern Checklist
+
+For each kernel, scan source for:
+- [ ] `sqrtf(x)` where `x` may be negative (e.g., variance computed via `mean(x²) - mean(x)²` can be slightly negative due to fp roundoff)
+- [ ] `1/x` or `x / y` where divisor may approach 0 (no ε floor present)
+- [ ] `logf(x)` where `x` may be ≤ 0 (no `fmaxf(x, ε)` clamp)
+- [ ] `expf(x)` where `x` may exceed ~88 (single-precision overflow → +inf)
+- [ ] EMA-tracked variance denominators that may approach 0 across folds
+- [ ] `atomicAdd` / `saxpy` accumulators receiving NaN inputs (no `isfinite()` guard at producer)
+- [ ] cuBLAS GEMM with extreme intermediate products (M,N,K combinations producing intermediate overflow)
+- [ ] `fmaxf` / `fminf` on potentially-NaN inputs (NaN propagates per IEEE-754 → masked-zero in downstream)
+
+## Cross-references
+
+- `~/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/session_2026-04-05_nan_investigation.md` — prior NaN root-cause investigation; identifies residual 8% step-2 NaN in `apply_iqn_trunk_gradient`'s cuBLAS backward (NEVER CLOSED — strong candidate).
+- Smoke `smoke-test-m5gxx` (commit `e9096c7be`) confirmed: `save_h_s2` slot 12 + `grad_buf` slot 6 flag at F1 step ~1140; ALL 5 GRN forward sub-stages (slots 16-20) read clean. NaN enters via backward path, not forward.
+
+---
+
+## Per-kernel sections
+
+EOF
+```
+
+Then for each of the kernels listed in Step 2, append a section using this template:
+
+```markdown
+### Kernel: `<KERNEL_NAME>`
+
+**File:** `<full/path/to/file.cu>` (lines `<start>-<end>`)
+**Output buffers:** `<list of buffers this kernel writes to>`
+**Accumulator targets:** `<list of buffers this kernel atomicAdd / saxpy / += into>`
+**Caller:** `<Rust orchestrator method that launches this kernel>`
+
+**Identified unsafe patterns:**
+- (Pattern, file:line, severity: confirmed/suspect/clean, brief description)
+
+**Proposed guard form** (only if any unsafe pattern present):
+- `<guard description, e.g., "wrap sqrtf(x) → sqrtf(fmaxf(x, 0.0f))">` at file:line
+
+**ISV bound option** (only if guard requires a dynamic bound):
+- Option 1: existing slot `<slot name + index>` provides `<what>` — feasible if `<conditions>`
+- Option 2: new ISV slot `<NAME_INDEX = N>` would track `<what>` — required if existing slots don't cover
+- Option 3: numerical-stability ε bound (Invariant 1 carve-out) — applicable only when bound is structural, not dynamic
+
+**F0 risk:** low / medium / high — `<reasoning>`
+
+**Phase B flag-slot allocation:**
+- Slot `<24-35>`: name `<short_name_for_code>` — buffer `<which output buffer to NaN-check>`
+```
+
+- [ ] **Step 2: Audit each kernel in this list**
+
+Read each file end-to-end and fill in the per-kernel section. Use `Read` tool to inspect source; do not modify any code. The list below is the minimum coverage; add additional kernels if they're discovered to write to backward-path buffers.
+
+Kernels to audit (in dependency order — earliest in backward chain first):
+
+1. **`backward_full`** (orchestrator)
+   - File: `crates/ml/src/cuda_pipeline/batched_backward.rs:1665+`
+   - This is the Rust-side orchestrator; document the call sequence (which kernel launches first, which buffers are populated when). Becomes the dependency graph used for triage in error-handling failure mode (3) — multiple kernels flagging.
+
+2. **C51 KL backward**
+   - File: `crates/ml/src/cuda_pipeline/c51_grad_kernel.cu`
+   - Look for: log of softmax, projection of TD targets onto atoms (potential `expf(large)`)
+
+3. **C51 loss kernel**
+   - File: `crates/ml/src/cuda_pipeline/c51_loss_kernel.cu`
+   - Look for: divides by `delta_z` (potential 0), atom interpolation, normalization
+
+4. **`apply_iqn_trunk_gradient`** (HIGH PRIORITY — known residual from `session_2026-04-05_nan_investigation.md`)
+   - File: search via `grep -n "fn apply_iqn_trunk_gradient" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`
+   - Look for: cuBLAS GEMM with extreme dims, residual gradient norm computation, possible accumulator NaN
+
+5. **IQN dual head kernel**
+   - File: `crates/ml/src/cuda_pipeline/iqn_dual_head_kernel.cu`
+   - Look for: `sqrtf` on accumulated sums (per `session_2026-04-05` finding #5: sqrtf-per-block bug), `isfinite` guards in Adam path
+
+6. **CQL gradient**
+   - File: `crates/ml/src/cuda_pipeline/cql_grad_kernel.cu` and `crates/ml/src/cuda_pipeline/cql_logit_grad_kernel.cu` if exists
+   - Look for: log-sum-exp computations (need numerical stability), accumulators
+
+7. **Aux heads backward**
+   - File: `crates/ml/src/cuda_pipeline/aux_heads_kernel.cu`
+   - Look for: regression head MSE backward, classification head softmax-CE backward
+
+8. **Ensemble backward** (currently dormant — instrument anyway)
+   - File: `crates/ml/src/cuda_pipeline/ensemble_kernels.cu`
+   - Look for: ensemble averaging numerics, member loss propagation
+
+9. **MSE / Bellman residual gradient**
+   - File: search via `grep -rn "mse_loss_kernel\|bellman_grad" crates/ml/src/cuda_pipeline/*.cu`
+   - Look for: `sqrtf(loss)` patterns, residual computation
+
+10. **Bottleneck Linear backward** (if separate from `backward_full`)
+    - File: search via `grep -rn "bottleneck.*backward" crates/ml/src/cuda_pipeline/*.cu` `crates/ml/src/cuda_pipeline/*.rs`
+    - Look for: dW / dX GEMM stride mismatches (per `session_2026-04-05` findings #3, #4)
+
+11. **`bw_d_h_s2` accumulator path** (Rust orchestrator, not a kernel)
+    - File: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — search via `grep -n "bw_d_h_s2" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`
+    - Document the saxpy stack: who writes into `bw_d_h_s2`, in what order, with what scaling. The before/after instrumentation (slots 32, 33) needs this dependency order.
+
+For each kernel, the section MUST include:
+- File path + line range
+- Output buffers (what kernel writes)
+- Accumulator targets (what kernel `atomicAdd` / `saxpy` into)
+- Per-pattern severity (confirmed unsafe / suspect / clean)
+- Proposed guard form if any pattern is non-clean
+- ISV bound option (with specific slot name) or numerical-stability ε justification
+- F0 risk assessment
+- Proposed Phase B flag-slot index from the spec table
+
+- [ ] **Step 3: Cross-reference with prior session notes**
+
+Read `~/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/session_2026-04-05_nan_investigation.md` (use `Read` tool with absolute path). For each finding in that note (5 root causes + residual 8% NaN at step 2), check whether the current code still has the original problematic pattern OR whether the documented defense-in-depth guard is in place. Add a "Prior-investigation status" subsection to the audit doc summarizing:
+- Each prior finding: still-present / fixed / superseded
+- Each existing defense-in-depth guard: location + line in current code, still active / removed / weakened
+
+- [ ] **Step 4: Write summary section ranking suspect kernels**
+
+Append to audit doc:
+
+```markdown
+## Summary — kernel suspicion ranking
+
+Ranked by cumulative severity (number of unsafe patterns × severity), with `apply_iqn_trunk_gradient` weighted by prior-investigation note as inherently high-priority regardless of pattern count.
+
+| Rank | Kernel | # confirmed | # suspect | F0 risk of fix | Notes |
+|---|---|---|---|---|---|
+| 1 | `<top suspect>` | N | M | low/med/high | brief |
+| 2 | ... | ... | ... | ... | ... |
+
+## Phase B priority slot order
+
+Based on the suspicion ranking above, instrument these slots first (descending priority):
+1. Slot `<idx>` — `<buffer>` (highest suspicion)
+2. ...
+12. ...
+```
+
+- [ ] **Step 5: Update Invariant 7 audit doc co-stage**
+
+Append to `docs/dqn-wire-up-audit.md`:
+
+```bash
+cat >> docs/dqn-wire-up-audit.md <<'EOF'
+
+SP1 Phase A audit (2026-04-29): produced `docs/dqn-backward-nan-audit.md` — read-only γ inventory of every backward-path kernel writing to `bw_d_h_s2` / `grad_buf` / `save_h_s2` accumulators, cross-referenced against `session_2026-04-05_nan_investigation.md`'s residual 8% step-2 NaN in `apply_iqn_trunk_gradient`. Per-kernel sections include: identified unsafe patterns (sqrtf-neg, 1/0, logf-≤0, expf-large, EMA variance, atomicAdd/saxpy NaN-propagation), proposed guard form, ISV bound option (existing slot leverage or new slot or Invariant 1 ε carve-out), F0 risk assessment (low/medium/high used as paper-review gate before smoke), and Phase B flag-slot allocation. Drives SP1 Phase B instrumentation (12 new slots in `nan_flags_buf` 24→48) and Phase C surgical fix decisions. Becomes durable input artifact for SP2 (framework codification) and SP3 (structural-fix scoping).
+EOF
+```
+
+- [ ] **Step 6: Verify pre-commit hook passes**
+
+```bash
+git add docs/dqn-backward-nan-audit.md docs/dqn-wire-up-audit.md
+SQLX_OFFLINE=true cargo check --workspace --message-format=short 2>&1 | tail -3
+```
+
+Expected output: `Finished` line; no errors. Pre-commit hook will accept the docs-only change.
+
+- [ ] **Step 7: Commit**
+
+```bash
+git commit -m "docs(dqn): SP1 Phase A — backward NaN audit
+
+Read-only γ audit of every backward-path kernel writing to bw_d_h_s2 /
+grad_buf / save_h_s2 accumulators. Per-kernel inventory against
+unsafe-pattern checklist (sqrtf-neg, 1/0, logf-≤0, expf-large, EMA
+variance, atomicAdd/saxpy NaN-propagation). Each section includes:
+identified pattern + line citation, severity, proposed guard form,
+ISV bound option, F0 risk, Phase B flag-slot allocation.
+
+Cross-referenced session_2026-04-05_nan_investigation.md residual 8%
+step-2 NaN in apply_iqn_trunk_gradient — never closed; ranked top
+suspect for SP1.
+
+Output drives SP1 Phase B (12 new flag slots in nan_flags_buf 24→48)
+and Phase C surgical fix(es). Durable artifact for SP2 framework
+codification + SP3 structural-fix scoping."
+```
+
+---
+
+## Phase B — β Instrumentation (always lands)
+
+### Task 2: Expand `nan_flags_buf` from 24 to 48 slots
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (allocation site, `read_nan_flags` signature, internal arrays)
+- Modify: `crates/ml/src/trainers/dqn/fused_training.rs` (`read_nan_flags` signature)
+- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (name tables in two consumer sites)
+- Modify: `docs/dqn-wire-up-audit.md` (Invariant 7 entry)
+
+This is a structural buffer-size change with name-table updates. No new instrumentation logic yet — purely makes room for slots 24-47. Validates compile; smoke not needed (no behavioral change since slots 24-47 stay at zero).
+
+- [ ] **Step 1: Locate allocation site**
+
+```bash
+grep -n "nan_flags_buf = stream.alloc_zeros" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
+```
+
+Expected: one hit at approximately `gpu_dqn_trainer.rs:11178`.
+
+- [ ] **Step 2: Update allocation size**
+
+In `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`, find:
+
+```rust
+let nan_flags_buf = stream.alloc_zeros::<i32>(24)
+    .map_err(|e| MLError::ModelError(format!("nan_flags alloc: {e}")))?;
+```
+
+Change to:
+
+```rust
+// SP1 Phase B: expanded 24 → 48 to make room for backward-path NaN
+// checks (slots 24-35) plus 12 reserved headroom slots (36-47) for
+// SP2 framework checker hooks + SP3 structural observers. Buffer
+// size reviewable by SP2 framework codification.
+let nan_flags_buf = stream.alloc_zeros::<i32>(48)
+    .map_err(|e| MLError::ModelError(format!("nan_flags alloc: {e}")))?;
+```
+
+- [ ] **Step 3: Update `read_nan_flags` signature in `gpu_dqn_trainer.rs`**
+
+Find:
+
+```bash
+grep -n "pub fn read_nan_flags" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
+```
+
+Expected: hit at approximately `gpu_dqn_trainer.rs:14982`. Read the function:
+
+```rust
+pub fn read_nan_flags(&self) -> Result<[i32; 24], MLError> {
+    let mut host = [0i32; 24];
+    self.stream.memcpy_dtoh(&self.nan_flags_buf, &mut host)
+        .map_err(|e| MLError::ModelError(format!("nan_flags read: {e}")))?;
+    Ok(host)
+}
+```
+
+Change to:
+
+```rust
+pub fn read_nan_flags(&self) -> Result<[i32; 48], MLError> {
+    let mut host = [0i32; 48];
+    self.stream.memcpy_dtoh(&self.nan_flags_buf, &mut host)
+        .map_err(|e| MLError::ModelError(format!("nan_flags read: {e}")))?;
+    Ok(host)
+}
+```
+
+- [ ] **Step 4: Update `read_nan_flags` signature in `fused_training.rs`**
+
+```bash
+grep -n "fn read_nan_flags" crates/ml/src/trainers/dqn/fused_training.rs
+```
+
+Expected: pub(crate) accessor that delegates to trainer. Find:
+
+```rust
+pub(crate) fn read_nan_flags(&self) -> Result<[i32; 24], MLError> {
+    self.trainer.read_nan_flags()
+}
+```
+
+Change return type to `Result<[i32; 48], MLError>`. The body delegates unchanged.
+
+- [ ] **Step 5: Locate name table consumer sites in `training_loop.rs`**
+
+```bash
+grep -n 'let names = \[' crates/ml/src/trainers/dqn/trainer/training_loop.rs
+```
+
+Expected: TWO hits — one in the `halt_nan` block (around line ~1985) and one in the `halt_grad_collapse` block (added by commit `d1808df14`, around line ~2017). BOTH must be updated.
+
+- [ ] **Step 6: Update name tables in both consumer sites**
+
+Each `let names = [...]` array needs to grow from 24 entries to 48. The new entries are slots 24-47 per the spec table. Insert AFTER the existing 24 entries:
+
+```rust
+let names = [
+    "states_buf", "on_v_logits", "on_b_logits",
+    "mse_loss_scalar", "params_buf_pre_fwd", "params_ptr_pre_fwd",
+    "grad_buf", "save_current_lp", "save_projected",
+    "moe_gate_softmax", "aux_nb_loss_scalar", "aux_rg_loss_scalar",
+    "save_h_s2",
+    "rsv13", "rsv14", "rsv15",
+    "grn_h_s2_elu_post",      // 16
+    "grn_h_s2_linear_b_out",  // 17
+    "grn_h_s2_glu_sigmoid",   // 18
+    "grn_h_s2_ln_rstd",       // 19
+    "grn_h_s2_ln_normed",     // 20
+    "rsv21", "rsv22", "rsv23",
+    // SP1 Phase B (slots 24-35): backward-path kernel NaN checks
+    "iqn_trunk_grad_out",     // 24
+    "iqn_aux_grad_out",       // 25
+    "c51_d_current_lp",       // 26
+    "c51_d_projected",        // 27
+    "cql_grad",               // 28
+    "ens_grad_out",           // 29
+    "aux_nb_dh_s2",           // 30
+    "aux_rg_dh_s2",           // 31
+    "bw_d_h_s2_pre_saxpy",    // 32
+    "bw_d_h_s2_post_saxpy",   // 33
+    "mse_grad",               // 34
+    "bottleneck_grad",        // 35
+    // SP1 reserved headroom (slots 36-47) for SP2 framework + SP3 observers
+    "rsv36", "rsv37", "rsv38", "rsv39",
+    "rsv40", "rsv41", "rsv42", "rsv43",
+    "rsv44", "rsv45", "rsv46", "rsv47",
+];
+```
+
+Apply to BOTH sites identically (the names must match between the two error-message sites).
+
+- [ ] **Step 7: Verify compile**
+
+```bash
+SQLX_OFFLINE=true cargo check -p ml --lib --message-format=short 2>&1 | tail -10
+```
+
+Expected: clean (no new warnings beyond the pre-existing baseline). Compile-only validates the size + signature changes.
+
+- [ ] **Step 8: Update Invariant 7 audit doc**
+
+```bash
+cat >> docs/dqn-wire-up-audit.md <<'EOF'
+
+SP1 Phase B foundation (2026-04-29): expanded `nan_flags_buf` 24→48 (allocation size in `gpu_dqn_trainer.rs`; `read_nan_flags` signature `[i32; 24]` → `[i32; 48]` in both `gpu_dqn_trainer.rs` and `fused_training.rs`; name tables updated in both `training_loop.rs` consumer sites — `halt_nan` block + `halt_grad_collapse` block from commit `d1808df14`). Slots 24-35 reserved for backward-kernel NaN checks per spec Components C; slots 36-47 reserved as headroom for SP2 framework checker hooks + SP3 structural observers. No behavioral change in this commit (new slots stay at zero until Task 4 wires the check call sites). Buffer size reviewable by SP2 framework codification — if right-size differs (e.g., 36 with no headroom or 64 for more coverage), SP2 may resize.
+EOF
+```
+
+- [ ] **Step 9: Commit**
+
+```bash
+git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \
+        crates/ml/src/trainers/dqn/fused_training.rs \
+        crates/ml/src/trainers/dqn/trainer/training_loop.rs \
+        docs/dqn-wire-up-audit.md
+git commit -m "feat(dqn): SP1 Phase B foundation — nan_flags_buf 24→48
+
+Expands the NaN flag buffer from 24 to 48 slots to make room for
+backward-path NaN checks (slots 24-35 per SP1 spec Components C)
+plus 12 reserved headroom slots (36-47) for SP2 framework + SP3
+observer hooks.
+
+Touches:
+- gpu_dqn_trainer.rs: alloc size 24→48; read_nan_flags signature
+  [i32; 24]→[i32; 48]
+- fused_training.rs: pub(crate) read_nan_flags signature
+- training_loop.rs: BOTH name table sites (halt_nan +
+  halt_grad_collapse) updated to 48 entries with slot 24-35 names
+
+No behavioral change — new slots stay at zero until Task 4 wires
+the kernel-output NaN checks. Buffer size reviewable by SP2."
+```
+
+---
+
+### Task 3: Add backward-buffer accessors
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (new `pub(crate)` accessor methods on relevant struct(s))
+- Modify: `crates/ml/src/cuda_pipeline/batched_backward.rs` (new `pub(crate)` accessors if needed for orchestrator-owned buffers)
+- Modify: `docs/dqn-wire-up-audit.md` (Invariant 7 entry)
+
+The Phase B NaN checks (slots 24-35) need `u64` device pointers to specific output buffers. Some buffers are already accessible (e.g., via `self.bw_d_h_s2.raw_ptr()`); others need new accessors following the pattern from commit `e9096c7be` (`GrnBlock::elu_post_ptr()` etc.).
+
+- [ ] **Step 1: Identify which buffers need new accessors**
+
+For each slot in the table, the buffer pointer source:
+
+| Slot | Buffer | Source location |
+|---|---|---|
+| 24 | `iqn_trunk_grad_out` | Output of `apply_iqn_trunk_gradient` cuBLAS bwd — likely already accessible via `self.cql_grad_scratch` field OR a dedicated buffer; audit doc Task 1 must specify |
+| 25 | `iqn_aux_grad_out` | IQN aux head gradient — likely on `gpu_iqn_head` orchestrator |
+| 26 | `c51_d_current_lp` | C51 KL backward scratch — may be inside `c51_grad_kernel` invocation; could be the same buffer as `save_current_lp` after backward overwrite |
+| 27 | `c51_d_projected` | C51 projection backward — same caveat |
+| 28 | `cql_grad` | Already accessible via `self.cql_grad_scratch.raw_ptr()` |
+| 29 | `ens_grad_out` | Ensemble backward output — `ensemble_kernels.cu` orchestrator |
+| 30 | `aux_nb_dh_s2` | Aux next-bar dh_s2 — `aux_dh_s2_nb_buf` field on trainer |
+| 31 | `aux_rg_dh_s2` | Aux regime dh_s2 — `aux_dh_s2_rg_buf` field on trainer |
+| 32 | `bw_d_h_s2_pre_saxpy` | Snapshot of `bw_d_h_s2` BEFORE saxpy stack runs — needs new scratch buffer + DtoD copy at the right point |
+| 33 | `bw_d_h_s2_post_saxpy` | `bw_d_h_s2` AFTER saxpy — already accessible via `self.bw_d_h_s2.raw_ptr()` |
+| 34 | `mse_grad` | MSE residual gradient — search audit doc for source |
+| 35 | `bottleneck_grad` | Bottleneck Linear bwd output — search audit doc for source |
+
+For each buffer that's already a `pub(crate)` field or has an existing accessor, no new code needed. For buffers that are private to a struct, add a `pub(crate) fn <name>_ptr(&self) -> u64` returning `self.<field>.raw_ptr()`.
+
+- [ ] **Step 2: Add new `bw_d_h_s2_pre_saxpy` scratch buffer**
+
+Slot 32 requires a snapshot of `bw_d_h_s2` BEFORE the saxpy stack runs. This needs a new scratch buffer and a DtoD copy at the right point in the backward orchestration.
+
+In `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`, find the existing `bw_d_h_s2: CudaSlice<f32>` field (around line 3140) and add a sibling field directly after:
+
+```rust
+/// SP1 Phase B (slot 32): snapshot of `bw_d_h_s2` taken BEFORE the
+/// saxpy stack accumulates aux/IQN/etc. dh_s2 contributions. Compared
+/// against `bw_d_h_s2` post-saxpy (slot 33) to identify which saxpy
+/// source first contaminates the trunk gradient accumulator.
+///
+/// Sized identically to `bw_d_h_s2` ([B * SHARED_H2] f32). Populated
+/// via DtoD copy in `launch_cublas_backward_to` at the point right
+/// after the cuBLAS trunk-bwd writes the initial value, before the
+/// saxpy chain runs.
+bw_d_h_s2_pre_saxpy: CudaSlice<f32>,
+```
+
+In the constructor (find via `grep -n "bw_d_h_s2:" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — there are likely two hits, one for the field declaration and one for the construction), add the allocation right after `bw_d_h_s2`:
+
+```rust
+let bw_d_h_s2_pre_saxpy = stream.alloc_zeros::<f32>(b * sh2)
+    .map_err(|e| MLError::ModelError(format!("bw_d_h_s2_pre_saxpy alloc: {e}")))?;
+```
+
+Add to the struct-literal initialization where `bw_d_h_s2` is bound.
+
+Add an accessor:
+
+```rust
+pub(crate) fn bw_d_h_s2_pre_saxpy_ptr(&self) -> u64 {
+    self.bw_d_h_s2_pre_saxpy.raw_ptr()
+}
+```
+
+- [ ] **Step 3: Wire the DtoD snapshot copy at the right point**
+
+In `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`, find where `bw_d_h_s2` is first written by the cuBLAS trunk backward (search `bw_d_h_s2_ptr` for the launch site that writes initially — distinct from the saxpy-accumulation sites). The snapshot must happen AFTER that initial write but BEFORE the first saxpy.
+
+Add a DtoD copy:
+
+```rust
+// SP1 Phase B (slot 32): snapshot bw_d_h_s2 BEFORE saxpy stack.
+// Copy is graph-capture-safe (DtoD via cudaMemcpyAsync); no host sync.
+unsafe {
+    cudarc::driver::result::memcpy_dtod_sync(
+        self.bw_d_h_s2_pre_saxpy_ptr(),
+        self.bw_d_h_s2.raw_ptr(),
+        (b * sh2 * std::mem::size_of::<f32>()) as usize,
+    )
+}.map_err(|e| MLError::ModelError(format!("bw_d_h_s2 snapshot DtoD: {e}")))?;
+```
+
+If the existing codebase uses `stream.memcpy_dtod` or a similar wrapper, prefer that wrapper for consistency with surrounding code (search via `grep -n "memcpy_dtod\|DtoD" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs | head -5` to find existing pattern).
+
+- [ ] **Step 4: Add accessors for any other Phase B target buffers identified by audit**
+
+For each slot 24-35 buffer that isn't already accessible, add `pub(crate) fn <name>_ptr(&self) -> u64 { self.<field>.raw_ptr() }` on the appropriate struct. Likely candidates:
+
+- `iqn_trunk_grad_out_ptr` if not already exposed
+- `ens_grad_out_ptr` for ensemble (if instrumented despite being dormant)
+- `mse_grad_ptr` and `bottleneck_grad_ptr` for those slots
+
+Each accessor follows the GrnBlock pattern from commit `e9096c7be`. The exact set is determined by the audit document — if a buffer is private to a sub-struct (e.g., GpuIqnHead, GpuEnsemble), the accessor may need to land on that sub-struct first, then a delegate accessor on the parent trainer.
+
+- [ ] **Step 5: Verify compile**
+
+```bash
+SQLX_OFFLINE=true cargo check -p ml --lib --message-format=short 2>&1 | tail -10
+```
+
+Expected: clean.
+
+- [ ] **Step 6: Update Invariant 7 audit doc**
+
+```bash
+cat >> docs/dqn-wire-up-audit.md <<'EOF'
+
+SP1 Phase B accessors (2026-04-29): added `pub(crate)` accessor methods on `GpuDqnTrainer` (and sub-structs as needed) exposing backward-path buffer device pointers for the new NaN checks (slots 24-35). New scratch buffer `bw_d_h_s2_pre_saxpy: CudaSlice<f32>` ([B * SHARED_H2] f32) for slot 32 — snapshot of `bw_d_h_s2` taken BEFORE the saxpy stack via graph-capture-safe DtoD copy; compared against post-saxpy state (slot 33) to identify which saxpy source first contaminates the trunk-gradient accumulator. Snapshot copy point: in `launch_cublas_backward_to` after cuBLAS trunk-bwd's initial write, before the first aux/IQN/ensemble saxpy. Pattern follows commit `e9096c7be`'s `GrnBlock::elu_post_ptr()` accessors. Additive — no behavioral change; new accessors consumed by Task 4's NaN check call sites.
+EOF
+```
+
+- [ ] **Step 7: Commit**
+
+```bash
+git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \
+        crates/ml/src/cuda_pipeline/batched_backward.rs \
+        docs/dqn-wire-up-audit.md
+git commit -m "feat(dqn): SP1 Phase B accessors — backward-buffer NaN check pointers
+
+Adds pub(crate) accessor methods on GpuDqnTrainer (and sub-structs as
+needed) exposing backward-path buffer device pointers for the new
+NaN checks (slots 24-35).
+
+New scratch buffer bw_d_h_s2_pre_saxpy [B * SHARED_H2] f32 for
+slot 32 — snapshot of bw_d_h_s2 taken BEFORE the saxpy stack via
+graph-capture-safe DtoD copy. Compared against post-saxpy state
+(slot 33) to identify which saxpy source first contaminates the
+trunk-gradient accumulator.
+
+Pattern follows commit e9096c7be's GrnBlock::elu_post_ptr() accessors.
+Additive — no behavioral change."
+```
+
+---
+
+### Task 4: Wire backward-kernel NaN check call sites
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (new `run_nan_checks_post_backward` method; possible call site insertion)
+- Modify: `crates/ml/src/trainers/dqn/fused_training.rs` (call `run_nan_checks_post_backward` at the right point in the training step)
+- Modify: `docs/dqn-wire-up-audit.md` (Invariant 7 entry)
+
+This is the meat of Phase B — actually firing `check_nan_f32(buf_ptr, len, flag_idx)` for slots 24-35 every training step. Pattern matches the existing `run_nan_checks_post_forward` in `gpu_dqn_trainer.rs:14870`.
+
+- [ ] **Step 1: Add new method `run_nan_checks_post_backward`**
+
+In `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`, immediately after the existing `run_nan_checks_post_forward` method (ends around line 14934), add:
+
+```rust
+/// SP1 Phase B (slots 24-35): NaN checks on backward-path kernel outputs.
+///
+/// Index → buffer mapping (must match training_loop.rs name table sites):
+///   24 iqn_trunk_grad_out      — apply_iqn_trunk_gradient cuBLAS bwd output
+///   25 iqn_aux_grad_out        — IQN aux head gradient
+///   26 c51_d_current_lp        — C51 KL backward (current_lp gradient)
+///   27 c51_d_projected         — C51 projection backward
+///   28 cql_grad                — CQL gradient output (cql_grad_scratch)
+///   29 ens_grad_out            — Ensemble backward output (dormant; instrument anyway)
+///   30 aux_nb_dh_s2            — Aux next-bar backward dh_s2
+///   31 aux_rg_dh_s2            — Aux regime backward dh_s2
+///   32 bw_d_h_s2_pre_saxpy     — bw_d_h_s2 snapshot BEFORE saxpy stack
+///   33 bw_d_h_s2_post_saxpy    — bw_d_h_s2 AFTER saxpy stack (same buffer; before-after)
+///   34 mse_grad                — MSE/Bellman residual gradient
+///   35 bottleneck_grad         — Bottleneck Linear backward output (if exists)
+///
+/// Called from `submit_dqn_step_loop_cublas` after `launch_cublas_backward_to`
+/// completes — the saxpy snapshot (slot 32) was taken inside backward; the
+/// remaining checks read the post-backward state of each kernel's output.
+///
+/// Pattern: identical to run_nan_checks_post_forward (single-block reduce per
+/// buffer, no atomicAdd, no DtoH per step). Fires every training step inside
+/// the captured graph; flags accumulate within a fold and reset on fold
+/// boundary via reset_nan_flags().
+pub fn run_nan_checks_post_backward(&mut self, batch_size: usize) -> Result<(), MLError> {
+    let b = batch_size;
+    let sh2 = self.config.shared_h2;
+    let na = self.config.num_atoms;
+    let b0 = self.config.branch_0_size;
+
+    // Slot 24: IQN trunk gradient cuBLAS bwd output (TBD by audit; placeholder
+    // points at cql_grad_scratch which apply_iqn_trunk_gradient SAXPYs into;
+    // refine to dedicated IQN scratch buffer per audit findings)
+    self.check_nan_f32(<iqn_trunk_grad_out_ptr — see audit Task 1>, <len>, 24)?;
+
+    // Slot 25: IQN aux gradient (TBD by audit)
+    self.check_nan_f32(<iqn_aux_grad_out_ptr>, <len>, 25)?;
+
+    // Slot 26-27: C51 backward outputs (TBD by audit — likely save_current_lp
+    // and save_projected after backward overwrites them in-place)
+    self.check_nan_f32_b(<c51_d_current_lp_ptr>, b * b0 * na, 26)?;
+    self.check_nan_f32(<c51_d_projected_ptr>, b * b0 * na, 27)?;
+
+    // Slot 28: CQL gradient (cql_grad_scratch)
+    self.check_nan_f32(self.cql_grad_scratch_ptr(), self.total_params, 28)?;
+
+    // Slot 29: Ensemble gradient (dormant; instrument anyway)
+    self.check_nan_f32(<ens_grad_out_ptr>, <len>, 29)?;
+
+    // Slot 30-31: Aux head dh_s2 outputs
+    self.check_nan_f32(self.aux_nb_dh_s2_buf_ptr(), b * sh2, 30)?;
+    self.check_nan_f32(self.aux_rg_dh_s2_buf_ptr(), b * sh2, 31)?;
+
+    // Slot 32-33: bw_d_h_s2 before/after saxpy stack
+    self.check_nan_f32(self.bw_d_h_s2_pre_saxpy_ptr(), b * sh2, 32)?;
+    self.check_nan_f32(self.bw_d_h_s2.raw_ptr(), b * sh2, 33)?;
+
+    // Slot 34: MSE/Bellman residual gradient (TBD by audit)
+    self.check_nan_f32(<mse_grad_ptr>, <len>, 34)?;
+
+    // Slot 35: Bottleneck Linear bwd (TBD by audit; may be N/A — if so, leave
+    // as zero and document in audit)
+    self.check_nan_f32(<bottleneck_grad_ptr>, <len>, 35)?;
+
+    Ok(())
+}
+```
+
+**IMPORTANT**: the placeholders `<...>` in the code above must be replaced with the exact buffer pointers and lengths IDENTIFIED IN TASK 1's audit document. The audit doc is the source of truth for which buffer corresponds to each slot. If a slot's buffer is genuinely N/A (e.g., bottleneck Linear backward doesn't exist as a standalone kernel — it's fused into `backward_full`), document that in the audit doc and skip that slot's check (the slot stays at zero — a "not instrumented" indicator, distinct from "no NaN").
+
+- [ ] **Step 2: Wire the call site in `fused_training.rs`**
+
+Find the existing call to `run_nan_checks_post_forward`:
+
+```bash
+grep -n "run_nan_checks_post_forward" crates/ml/src/trainers/dqn/fused_training.rs
+```
+
+Expected: at least one hit at the per-step training loop where the post-forward checks fire.
+
+Immediately AFTER the existing post-forward call AND after the backward orchestration returns (search the same file for `launch_cublas_backward_to` or `backward_full` invocation), add:
+
+```rust
+// SP1 Phase B: NaN checks on backward-path kernel outputs (slots 24-35).
+// Fires every training step; flags accumulate within fold; reset on fold
+// boundary via reset_nan_flags(). Permanent diagnostic infrastructure.
+self.trainer.run_nan_checks_post_backward(batch_size)?;
+```
+
+The exact insertion point depends on where backward kernels actually run — read the surrounding code carefully. The check must happen AFTER all backward kernels have written their outputs but BEFORE the next forward step (where buffers may be reused/overwritten).
+
+- [ ] **Step 3: Verify compile**
+
+```bash
+SQLX_OFFLINE=true cargo check -p ml --lib --message-format=short 2>&1 | tail -10
+```
+
+Expected: clean. Compile validates all the buffer accessors exist and types match.
+
+- [ ] **Step 4: Update Invariant 7 audit doc**
+
+```bash
+cat >> docs/dqn-wire-up-audit.md <<'EOF'
+
+SP1 Phase B instrumentation (2026-04-29): wired backward-path NaN checks for slots 24-35 via new `run_nan_checks_post_backward` method on `GpuDqnTrainer` (mirrors `run_nan_checks_post_forward` pattern from line 14870). Called once per training step from `fused_training.rs` immediately after the backward orchestration completes. Coverage: IQN trunk + aux gradients (24-25), C51 KL/projection backward (26-27), CQL gradient (28), ensemble (29 — dormant but instrumented), aux head dh_s2 (30-31), bw_d_h_s2 before/after saxpy stack (32-33), MSE/Bellman grad (34), bottleneck Linear bwd (35). Each check uses existing `check_nan_f32(buf_ptr, len, flag_idx)` infrastructure (single-block GPU reduce, no atomicAdd, no per-step DtoH). Flags accumulate within fold, reset at fold boundary via existing `reset_nan_flags()`. Readback flow (commit `d1808df14`) consumes them in BOTH `halt_nan` and `halt_grad_collapse` paths. Permanent diagnostic infrastructure — stays as production-grade regression sentinel after surgical fix lands. Slots that map to N/A buffers (e.g., bottleneck if kernel doesn't exist) are documented in the audit doc and skipped (slot stays zero — "not instrumented", distinct from "no NaN").
+EOF
+```
+
+- [ ] **Step 5: Commit**
+
+```bash
+git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \
+        crates/ml/src/trainers/dqn/fused_training.rs \
+        docs/dqn-wire-up-audit.md
+git commit -m "feat(dqn): SP1 Phase B instrumentation — backward NaN checks (slots 24-35)
+
+Wires per-step NaN checks on backward-path kernel outputs. Coverage:
+IQN trunk (24) + aux (25) gradients, C51 KL backward (26-27), CQL
+gradient (28), ensemble (29 — dormant), aux head dh_s2 (30-31),
+bw_d_h_s2 before/after saxpy stack (32-33; before-after delta shows
+which saxpy contaminated), MSE/Bellman grad (34), bottleneck (35).
+
+New method run_nan_checks_post_backward on GpuDqnTrainer; called
+from fused_training.rs after backward orchestration. Mirrors
+run_nan_checks_post_forward pattern.
+
+Permanent diagnostic — stays as regression sentinel after the
+surgical fix lands. Each check uses existing check_nan_f32
+infrastructure (single-block GPU reduce, no atomicAdd, no DtoH per
+step). Flags accumulate within fold; reset at fold boundary.
+Readback (commit d1808df14) consumes flags in BOTH halt_nan and
+halt_grad_collapse paths."
+```
+
+---
+
+### Task 5: Validate β instrumentation captures F1 explosion topology
+
+**Files:** none modified (smoke run + log analysis)
+
+This task validates the Phase B instrumentation actually fires and captures the NaN topology. The smoke is expected to fail (F1 still NaNs — that's the whole point of SP1) but the diagnostic log MUST contain a richer flag list than just `[6, 12]`.
+
+- [ ] **Step 1: Push commits + submit smoke**
+
+```bash
+git log --oneline -5  # Should show Tasks 1, 2, 3, 4 commits + parent
+git push origin plan-c-phase-2-thompson
+SHORT_SHA=$(git rev-parse --short HEAD)
+argo submit -n foxhunt --from=wftmpl/smoke-test \
+  -p commit-ref=$SHORT_SHA \
+  -p test-name=multi_fold_convergence::test_multi_fold_convergence \
+  -p clean-cache=false 2>&1 | head -10
+```
+
+Note the workflow name (e.g., `smoke-test-XXXXX`) for monitoring.
+
+- [ ] **Step 2: Wait for smoke completion**
+
+L40S smoke takes ~16-20 minutes. Monitor via:
+
+```bash
+argo get smoke-test-XXXXX -n foxhunt 2>&1 | head -10
+```
+
+Wait for `Status: Succeeded` (or `Failed` if Argo infra fails — distinct from training failure).
+
+- [ ] **Step 3: Pull smoke log**
+
+```bash
+SMOKE_NAME=smoke-test-XXXXX  # actual workflow name
+mc cp fh/argo-logs/$SMOKE_NAME/$SMOKE_NAME/main.log /tmp/p4t6_revalidate/smoke-${SHORT_SHA}.log
+```
+
+- [ ] **Step 4: Verify Phase B instrumentation fired**
+
+```bash
+grep -E "NaN-CLAMPED-TO-ZERO|NaN SOURCE" /tmp/p4t6_revalidate/smoke-${SHORT_SHA}.log | head -5
+```
+
+Expected output: at least one `NaN-CLAMPED-TO-ZERO` line with a flagged-buffer list including SOME of slots 24-35 (e.g., `flagged=[24=iqn_trunk_grad_out, 32=bw_d_h_s2_pre_saxpy, 33=bw_d_h_s2_post_saxpy, ...]`). The exact buffer set tells us which kernels produced NaN.
+
+If the log shows ONLY the old-coverage flags (e.g., `[6=grad_buf, 12=save_h_s2]`) and NONE of slots 24-35 fired, the instrumentation isn't reaching the explosion path — debug Task 4's call-site insertion point and re-iterate.
+
+- [ ] **Step 5: Capture topology in audit doc**
+
+Append to `docs/dqn-backward-nan-audit.md`:
+
+```bash
+cat >> docs/dqn-backward-nan-audit.md <<EOF
+
+## Phase B smoke result: smoke-${SHORT_SHA}
+
+**Workflow:** ${SMOKE_NAME}
+**Commit:** ${SHORT_SHA}
+**Outcome:** F1 ep2 NaN (expected; SP1 not yet fixed)
+
+### Flagged buffer topology (first NaN-CLAMPED-TO-ZERO line)
+
+(paste the actual flagged=[...] list from the smoke log)
+
+### Per-fold outcome
+
+(paste F0 Best Sharpe, F1 outcome, F2 outcome from smoke log)
+
+### Backward dependency interpretation
+
+Based on the flagged buffer set and the Phase A dependency graph (in
+this audit doc's audit of \`backward_full\`), the FIRST kernel in
+backward order whose output is flagged is the source. Downstream
+flags are propagation. Source kernel: <NAME> at <FILE:LINE>.
+
+This drives Phase C surgical fix in Task 6.
+EOF
+```
+
+- [ ] **Step 6: Commit**
+
+This commit captures the diagnostic finding for posterity (no code change in this task — only audit doc update).
+
+```bash
+git add docs/dqn-backward-nan-audit.md
+git commit -m "docs(dqn): SP1 Phase B smoke result — F1 NaN topology captured
+
+Smoke smoke-${SHORT_SHA} (commit ${SHORT_SHA}) confirms Phase B
+instrumentation fires correctly: backward-path slots 24-35 capture
+the F1 ep2 explosion's NaN topology beyond the prior visible-only
+[6, 12] (grad_buf + save_h_s2).
+
+First flagged buffer in backward dependency order pinpoints the
+source kernel (per audit doc Phase A dependency graph). Downstream
+flags = propagation.
+
+Drives Phase C surgical fix decision."
+```
+
+---
+
+## Phase C — Surgical Fix
+
+### Task 6: Apply surgical fix(es) per audit + Phase B findings
+
+**Files:** TBD — depends on which kernel(s) γ + β identify as source(s).
+
+This task is necessarily content-driven by Tasks 1 + 5 outputs. The plan cannot pre-specify the exact patch. Instead, this task specifies the CONSTRAINTS and FORM of the patch, and provides templates for the most common patterns.
+
+The audit doc + smoke flag topology together identify:
+- Which kernel(s) produce(s) the FIRST NaN (downstream flags = propagation)
+- Which unsafe pattern in that kernel applies (γ catalogue)
+- The proposed guard form + ISV bound option (per audit doc per-kernel section)
+- F0 risk (paper-review gate before consuming an L40S smoke)
+
+- [ ] **Step 1: Confirm fix scope**
+
+Read the latest smoke log + audit doc. Identify:
+
+```
+Source kernel (first in backward order): <NAME at FILE:LINE>
+Unsafe pattern: <sqrtf-neg | 1/0 | logf-≤0 | expf-large | EMA-variance | atomicAdd-NaN | cuBLAS-overflow | other>
+Proposed guard form: <from audit doc>
+ISV bound option chosen: <slot name + index> OR <new slot description> OR <numerical-stability ε with justification>
+F0 risk assessment: <low/medium/high>
+Multi-source: <if smoke flagged additional independent NaN sources, list them — Phase 5 (5) "no escalation" rule applies; fix all in this commit>
+```
+
+Document this analysis in the audit doc as the Phase C section.
+
+- [ ] **Step 2: F0-safety paper review (only if any fix is medium/high risk)**
+
+For each proposed fix marked `F0 risk: medium` or `F0 risk: high`, walk through the math:
+
+- What's the typical input range to this kernel during F0 training? (Check existing HEALTH_DIAG output for F0 from the baseline smoke `smoke-test-m5gxx` log if available; otherwise reason from the kernel's role.)
+- Does the proposed guard fire on F0's typical inputs? If yes, what does F0 training look like with the guard active?
+- Concrete check: pick 3 representative F0 input values (low/typical/high), apply the guard mathematically, verify the output is unchanged (guard is a no-op for F0).
+
+If the paper review reveals the guard would change F0 dynamics, REVISE the guard form before coding. Common revisions: tighten the trigger condition, narrow the clamp range, add an ISV-driven adaptive bound that's wider on F0's data than F1's.
+
+- [ ] **Step 3: Apply the fix (kernel-specific code)**
+
+The exact code depends on the unsafe pattern. Common templates:
+
+**Template (a) — `sqrtf` of potentially-negative value:**
+
+```cuda
+// Before:
+float result = sqrtf(variance);
+
+// After (guard against fp-roundoff producing slightly negative variance):
+float result = sqrtf(fmaxf(variance, 0.0f));
+```
+
+**Template (b) — division by potentially-zero value with ISV-driven bound:**
+
+```cuda
+// Read existing ISV slot for adaptive denominator floor (e.g., Q_DIR_ABS_REF_INDEX = 21):
+const float floor_isv = isv_signals[Q_DIR_ABS_REF_INDEX];
+const float floor_eff = fmaxf(floor_isv, 1e-6f);  // 1e-6 = numerical-stability ε (Invariant 1)
+float result = numerator / fmaxf(fabsf(denominator), floor_eff);
+```
+
+**Template (c) — `logf` of potentially-non-positive value:**
+
+```cuda
+// Before:
+float result = logf(probability);
+
+// After (clamp to numerical-stability ε):
+float result = logf(fmaxf(probability, 1e-6f));  // 1e-6 = numerical-stability ε (Invariant 1)
+```
+
+**Template (d) — `expf` of potentially-large value (overflow):**
+
+```cuda
+// Before:
+float result = expf(value);
+
+// After (subtract max for log-sum-exp stability; bound argument):
+float result = expf(fminf(value, 88.0f));  // 88.0f ≈ log(FLT_MAX); numerical-stability bound
+```
+
+**Template (e) — `isfinite` guard before `atomicAdd` / saxpy:**
+
+```cuda
+// Before:
+atomicAdd(&accumulator[i], contribution);
+
+// After (drop non-finite contributions):
+if (isfinite(contribution)) {
+    atomicAdd(&accumulator[i], contribution);
+}
+```
+
+**Template (f) — wrapper around cuBLAS GEMM with input range guard + output sanitization:**
+
+```rust
+// Before: bare cublas_gemm(...)
+// After: wrapper with pre-call input-range check + post-call NaN clamp on output
+//
+// (Wrapper in Rust; the cuBLAS call itself isn't modified)
+let max_input = max_abs_buffer(&input)?; // existing GPU reduce kernel
+if !max_input.is_finite() || max_input > GEMM_INPUT_MAX_FROM_ISV {
+    // Pre-call: input contains extreme values; sanitize before GEMM
+    launch_clamp_finite(&mut input, GEMM_INPUT_MAX_FROM_ISV)?;
+}
+cublas_gemm(...);
+// Post-call: sanitize output (defense-in-depth — should be unnecessary if pre-call worked)
+launch_clamp_finite(&mut output, GEMM_OUTPUT_MAX_FROM_ISV)?;
+```
+
+The audit doc identifies which template applies. If multiple kernels need fixes (multi-source case), apply each template independently within the SAME commit (per `feedback_no_partial_refactor`).
+
+- [ ] **Step 4: Verify compile**
+
+```bash
+SQLX_OFFLINE=true cargo check -p ml --lib --message-format=short 2>&1 | tail -10
+```
+
+Expected: clean.
+
+- [ ] **Step 5: Update audit doc with fix-applied annotations**
+
+For each kernel patched, update the Phase A audit section of `docs/dqn-backward-nan-audit.md` to add a "Status: FIXED in commit <SHA-TBD>" note. After commit, edit the SHA in via a follow-up amend OR commit the audit doc update separately as a doc-only commit (preferred — keeps the rich-fix commit focused on code).
+
+```bash
+cat >> docs/dqn-backward-nan-audit.md <<'EOF'
+
+## Phase C — Surgical fix(es) applied
+
+### Fix 1: <kernel name>
+
+- **File patched:** <file:line>
+- **Pattern fixed:** <e.g., sqrtf of potentially-negative variance>
+- **Guard form applied:** <e.g., wrapped with fmaxf(x, 0.0f)>
+- **ISV bound source:** <slot name + index, OR Invariant 1 ε with justification>
+- **F0 risk validated:** <paper-review summary if medium/high>
+- **Commit:** <SHA after this commit lands; updated in follow-up>
+
+(Repeat for each kernel patched in the multi-source case.)
+EOF
+```
+
+- [ ] **Step 6: Update Invariant 7 audit doc**
+
+```bash
+cat >> docs/dqn-wire-up-audit.md <<'EOF'
+
+SP1 Phase C surgical fix (2026-04-29): patched <kernel name(s)> at <file:line(s)> to address the NaN source(s) identified by Phase A audit + Phase B smoke topology. Guard form: <e.g., wrapped sqrtf(x) with sqrtf(fmaxf(x, 0))>. ISV bound source: <slot name + index, or Invariant 1 ε justification>. F0 risk: <low/medium/high>; paper-review confirmed guard is no-op on F0-typical input ranges. Multi-source: <none, OR list each kernel with its fix per `feedback_no_partial_refactor` rich-commit principle>. Companion audit doc updates document the fix per kernel for SP2 framework codification.
+EOF
+```
+
+- [ ] **Step 7: Commit (rich commit per `feedback_no_partial_refactor`)**
+
+```bash
+git add <list of kernel files patched> docs/dqn-backward-nan-audit.md docs/dqn-wire-up-audit.md
+git commit -m "fix(dqn): SP1 Phase C — F1 NaN root-cause surgical fix
+
+Patches <KERNEL NAME(S)> to address the NaN source identified by
+Phase A audit + Phase B smoke topology (smoke-test-<SHA>).
+
+Source: <KERNEL at FILE:LINE>
+Pattern: <sqrtf-neg / 1/0 / logf-≤0 / expf-large / EMA-variance / etc.>
+Guard form: <e.g., wrapped sqrtf(x) → sqrtf(fmaxf(x, 0.0f))>
+ISV bound: <slot name + index, or Invariant 1 ε justification>
+F0 risk: <low/medium/high>; paper-review confirmed guard is no-op
+on F0-typical input ranges.
+
+Multi-source: <none, OR list of additional kernels patched in same
+commit per feedback_no_partial_refactor>
+
+Permanent diagnostic instrumentation (slots 24-35 from Phase B)
+stays as regression sentinel — will detect any future regression of
+this NaN class.
+
+SP1 Phase D validation smoke pending."
+```
+
+---
+
+## Phase D — Multi-fold Validation
+
+### Task 7: Multi-fold smoke run + pass-criterion check
+
+**Files:** none modified (smoke run + log analysis only)
+
+- [ ] **Step 1: Submit validation smoke**
+
+```bash
+git log --oneline -3  # Should show Task 6 fix commit on top
+git push origin plan-c-phase-2-thompson
+SHORT_SHA=$(git rev-parse --short HEAD)
+argo submit -n foxhunt --from=wftmpl/smoke-test \
+  -p commit-ref=$SHORT_SHA \
+  -p test-name=multi_fold_convergence::test_multi_fold_convergence \
+  -p clean-cache=false 2>&1 | head -10
+```
+
+(Use `--clean-cache=false` per spec efficiency consideration — same branch as previous smokes; no cross-branch artifacts to fear.)
+
+Note workflow name. Expected duration ~16-20 min.
+
+- [ ] **Step 2: Wait for completion + pull log**
+
+```bash
+argo get smoke-test-XXXXX -n foxhunt 2>&1 | head -10
+# Wait for Succeeded
+mc cp fh/argo-logs/smoke-test-XXXXX/smoke-test-XXXXX/main.log /tmp/p4t6_revalidate/smoke-${SHORT_SHA}.log
+```
+
+- [ ] **Step 3: Verify pass criterion (each item)**
+
+Run each grep separately to verify each criterion:
+
+```bash
+LOG=/tmp/p4t6_revalidate/smoke-${SHORT_SHA}.log
+
+echo "=== Criterion 1: zero NaN-CLAMPED-TO-ZERO (must be empty) ==="
+grep "NaN-CLAMPED-TO-ZERO" $LOG | wc -l
+# Expected: 0
+
+echo "=== Criterion 2: zero NaN/Inf at step (must be empty) ==="
+grep "NaN/Inf at step" $LOG | wc -l
+# Expected: 0
+
+echo "=== Criterion 3: F0 Best Sharpe ≥ 53.08 (95% of 55.87 baseline) ==="
+grep "Best Sharpe" $LOG | head -1
+# Manually check: the first Best Sharpe line is F0 — extract value, must be ≥ 53.08
+
+echo "=== Criterion 4: F1 last-epoch ≥ first-epoch ==="
+grep "Best Sharpe" $LOG | head -2 | tail -1
+# Manually compare: F1 last-epoch Best Sharpe vs first-epoch (need per-epoch logs)
+# If only "Best Sharpe at epoch N" lines visible, find F1's epoch 1 vs epoch 5 values
+
+echo "=== Criterion 5: F2 last-epoch ≥ first-epoch ==="
+grep "Best Sharpe" $LOG | head -3 | tail -1
+# Same comparison for F2
+
+echo "=== Criterion 6: zero flagged slots (all 48 NaN flags remained zero) ==="
+grep -E "flagged=\[" $LOG | wc -l
+# Expected: 0 (no NaN-source readback fired)
+
+echo "=== Criterion 7: all 3 folds completed 5 epochs ==="
+grep -E "Fold [0-2] complete" $LOG | wc -l
+# Expected: 3
+```
+
+Document the results in a per-criterion table:
+
+```
+Criterion 1 (zero NaN-CLAMPED): <PASS/FAIL>
+Criterion 2 (zero NaN/Inf at step): <PASS/FAIL>
+Criterion 3 (F0 ≥ 53.08): <PASS/FAIL — actual value: XX.XX>
+Criterion 4 (F1 monotone): <PASS/FAIL — first XX.XX vs last YY.YY>
+Criterion 5 (F2 monotone): <PASS/FAIL — first XX.XX vs last YY.YY>
+Criterion 6 (zero flagged slots): <PASS/FAIL>
+Criterion 7 (3 folds complete): <PASS/FAIL>
+```
+
+- [ ] **Step 4: Decision based on criterion results**
+
+**If ALL 7 criteria PASS:**
+- SP1 closes successfully. Proceed to Task 8 (closure documentation).
+
+**If ANY criterion FAILS:**
+- Identify which criterion. Common failures and responses:
+  - Criterion 1/2/6 (NaN reappears): the surgical fix didn't address all sources OR introduced a new one. Re-run γ + β analysis on the NEW topology (the smoke log will have `flagged=[...]` showing which kernels NOW produce NaN — different from the previous set, otherwise the fix didn't apply). Return to Task 6 with the new source(s).
+  - Criterion 3 (F0 regression): the fix dropped F0 below 95% baseline. Revise the guard form to be a stricter no-op for F0 inputs. Return to Task 6 step 2 (paper review).
+  - Criterion 4/5 (F1/F2 not monotone improving): the fix prevented NaN but training is degenerate. May need additional Phase A audit (something else is wrong); return to Task 1 with the new failure observation.
+  - Criterion 7 (folds incomplete): some other early-stop fired. Inspect smoke log for the early-stop reason; treat as a new investigation thread under SP1 (no deferrals — no escalation to SP2/SP3).
+- Update audit doc with the failure observation; Return to Task 6.
+- Per `feedback_no_partial_refactor` and the spec's no-deferral rule: continue iterating within SP1 until ALL 7 criteria pass.
+
+- [ ] **Step 5: Update audit doc with smoke result**
+
+```bash
+cat >> docs/dqn-backward-nan-audit.md <<EOF
+
+## Phase D — Multi-fold validation smoke
+
+**Workflow:** smoke-test-XXXXX
+**Commit:** ${SHORT_SHA}
+**Outcome:** <PASS / FAIL with criterion list>
+
+### Per-criterion results
+
+| Criterion | Result | Value |
+|---|---|---|
+| 1. zero NaN-CLAMPED-TO-ZERO | <PASS/FAIL> | <count> |
+| 2. zero NaN/Inf at step | <PASS/FAIL> | <count> |
+| 3. F0 Best Sharpe ≥ 53.08 | <PASS/FAIL> | <actual> |
+| 4. F1 last-epoch ≥ first-epoch | <PASS/FAIL> | first <X>, last <Y> |
+| 5. F2 last-epoch ≥ first-epoch | <PASS/FAIL> | first <X>, last <Y> |
+| 6. zero flagged slots | <PASS/FAIL> | <count> |
+| 7. 3 folds complete 5 epochs | <PASS/FAIL> | <count>/3 |
+
+(If FAIL: include the new flagged buffer topology, drives next iteration of Task 6.)
+EOF
+```
+
+- [ ] **Step 6: Commit (only if criteria all pass — otherwise iterate)**
+
+```bash
+git add docs/dqn-backward-nan-audit.md
+git commit -m "docs(dqn): SP1 Phase D — multi-fold validation passed
+
+Smoke smoke-test-XXXXX (commit ${SHORT_SHA}) satisfies all 7 SP1
+pass criteria:
+  1. zero NaN-CLAMPED-TO-ZERO log lines
+  2. zero NaN/Inf at step errors
+  3. F0 Best Sharpe XX.XX ≥ 53.08 (95% of 55.87 baseline)
+  4. F1 last-epoch <Y> ≥ first-epoch <X>
+  5. F2 last-epoch <Y> ≥ first-epoch <X>
+  6. zero flagged NaN slots throughout run
+  7. all 3 folds complete 5 epochs
+
+SP1 Phase D validation complete. Task 8 closure pending."
+```
+
+(If FAIL: do NOT commit a passing result. Return to Task 6 with the new findings.)
+
+---
+
+### Task 8: SP1 closure — final documentation + handoff prep
+
+**Files:**
+- Modify: `docs/dqn-backward-nan-audit.md` (closure section)
+- Create: `~/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/project_sp1_f1_nan_root_cause_resolved.md` (memory entry — SP2/SP3 reference)
+- Modify: `~/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/MEMORY.md` (index)
+- Modify: `docs/dqn-wire-up-audit.md` (Invariant 7 entry)
+
+This is the SP1 closure commit. It documents the complete investigation arc, captures the kernel(s) fixed and the unsafe pattern(s) addressed, prepares the handoff to SP2 (framework) and SP3 (structural).
+
+- [ ] **Step 1: Write closure section in audit doc**
+
+```bash
+cat >> docs/dqn-backward-nan-audit.md <<'EOF'
+
+## SP1 Closure — Investigation arc summary
+
+**Status:** CLOSED — all 7 pass criteria satisfied at commit <SHA>.
+
+### Root cause(s) identified
+
+(Brief summary: which kernel(s), which unsafe pattern(s), why the fix works.)
+
+### Fix delivered
+
+(Reference to the surgical-fix commit + audit doc Phase C section listing each fix.)
+
+### Permanent infrastructure landed
+
+- `nan_flags_buf` expanded 24 → 48 (Task 2).
+- `run_nan_checks_post_backward` added (Task 4) — covers 12 backward-path kernel outputs (slots 24-35).
+- `bw_d_h_s2_pre_saxpy` snapshot buffer (Task 3) — enables which-saxpy-contaminated diagnosis.
+- 12 reserved headroom slots (36-47) for SP2 framework + SP3 observers.
+
+### Handoff to SP2 (Numerical Stability Framework)
+
+The following patterns identified during SP1's audit + fix should be codified as generic guards in SP2:
+
+(List each unsafe pattern + applicable kernels, e.g.:
+- `sqrtf(potentially-negative)` → SAFE_SQRT macro applied to <list of kernels>
+- `1/(potentially-zero)` → SAFE_DIV with ISV-driven floor applied to <list>
+- `expf(potentially-large)` → SAFE_EXP with arg-clamp applied to <list>
+)
+
+### Handoff to SP3 (Q-learning Structural Stability)
+
+The following structural conditions enabled the F1 NaN that SP1's surgical fix now defends against. SP3 should address them at the source:
+
+(List each, e.g.:
+- C51 atom-position growth at fold boundary creates extreme `delta_z` denominators → SP3 atom-range governance.
+- Target-Q optimism-coupling drives extreme Bellman residuals → SP3 target-Q clipping or pessimistic ensemble.
+)
+EOF
+```
+
+- [ ] **Step 2: Create memory entry**
+
+```bash
+cat > ~/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/project_sp1_f1_nan_root_cause_resolved.md <<'EOF'
+---
+name: project_sp1_f1_nan_root_cause_resolved
+description: SP1 closure — F1 ep2 NaN explosion in Plan C Phase 2 fixed; permanent diagnostic infrastructure landed (nan_flags_buf 24→48); SP2 + SP3 handoff prepared
+type: project
+---
+
+## Status
+
+CLOSED at commit <SHA> on branch `plan-c-phase-2-thompson`.
+
+All 7 SP1 pass criteria satisfied on smoke-test-XXXXX:
+- zero NaN-CLAMPED-TO-ZERO lines, zero NaN/Inf errors
+- F0 Best Sharpe XX.XX ≥ 53.08 (95% of 55.87 baseline)
+- F1 last-epoch ≥ first-epoch (monotone training)
+- F2 last-epoch ≥ first-epoch (monotone training)
+- 3/3 folds complete 5 epochs each
+
+## Root cause
+
+(Kernel + unsafe pattern + fix form)
+
+## Permanent infrastructure
+
+- `nan_flags_buf` 24 → 48 with backward-path coverage at slots 24-35.
+- New `GpuDqnTrainer::run_nan_checks_post_backward` mirrors `run_nan_checks_post_forward` pattern.
+- New `bw_d_h_s2_pre_saxpy` snapshot buffer for which-saxpy-contaminated diagnosis.
+- Slots 36-47 reserved for SP2 framework + SP3 observer hooks.
+
+## Handoff
+
+- **SP2 (numerical stability framework):** codify identified unsafe patterns as kernel-level macros (SAFE_SQRT, SAFE_DIV, SAFE_LOG, SAFE_EXP, ISFINITE_GUARD) + cross-cutting refactor of existing kernels to adopt them + pre-commit invariant checks.
+- **SP3 (Q-learning structural stability):** target-Q clipping, pessimistic ensemble wiring, atom-position growth bounds, conservative target sync at fold boundary.
+
+## References
+
+- Spec: `docs/superpowers/specs/2026-04-29-numerical-stability-investigation-design.md`
+- Plan: `docs/superpowers/plans/2026-04-29-numerical-stability-sp1-f1-nan-root-cause.md`
+- Audit doc: `docs/dqn-backward-nan-audit.md` (durable artifact for SP2/SP3)
+EOF
+```
+
+- [ ] **Step 3: Add to memory index**
+
+Edit `~/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/MEMORY.md` and add to the project pearl section:
+
+```markdown
+- [project_sp1_f1_nan_root_cause_resolved.md](project_sp1_f1_nan_root_cause_resolved.md) — SP1 closure: F1 ep2 NaN in Plan C Phase 2 fixed; permanent diagnostic infrastructure (nan_flags_buf 24→48, slots 24-35 backward coverage); SP2/SP3 handoff prepared
+```
+
+- [ ] **Step 4: Update Invariant 7 audit doc**
+
+```bash
+cat >> docs/dqn-wire-up-audit.md <<'EOF'
+
+SP1 closure (2026-04-29): F1 ep2 NaN root-cause fixed at commit <SHA>. All 7 pass criteria satisfied on smoke-test-XXXXX. Permanent infrastructure: `nan_flags_buf` 24→48 with backward-path coverage at slots 24-35; new `run_nan_checks_post_backward` method; `bw_d_h_s2_pre_saxpy` snapshot buffer; 12 reserved headroom slots for SP2/SP3 hooks. Audit doc `docs/dqn-backward-nan-audit.md` becomes durable artifact for SP2 (framework codification) and SP3 (structural-fix scoping). Memory entry `project_sp1_f1_nan_root_cause_resolved.md` documents handoff. Branch `plan-c-phase-2-thompson` ready for merge OR for SP2's framework work.
+EOF
+```
+
+- [ ] **Step 5: Commit**
+
+```bash
+git add docs/dqn-backward-nan-audit.md docs/dqn-wire-up-audit.md
+git commit -m "docs(dqn): SP1 closure — F1 NaN resolved; SP2/SP3 handoff prepared
+
+SP1 (F1 NaN root-cause) closes successfully at commit <SHA>:
+- All 7 pass criteria satisfied on smoke-test-XXXXX
+- Root cause: <KERNEL + unsafe pattern> patched
+- Permanent diagnostic infrastructure committed
+  - nan_flags_buf 24→48
+  - run_nan_checks_post_backward method
+  - bw_d_h_s2_pre_saxpy snapshot
+  - slots 36-47 reserved for SP2/SP3
+- Audit doc \`docs/dqn-backward-nan-audit.md\` becomes durable
+  artifact for sub-projects 2 and 3
+- Memory entry project_sp1_f1_nan_root_cause_resolved.md documents
+  handoff to SP2 (framework codification) and SP3 (structural Q
+  stability)
+
+Branch plan-c-phase-2-thompson ready for merge OR for SP2's
+framework work."
+```
+
+---
+
+## Self-Review
+
+**Spec coverage check** (each section/requirement → task):
+
+| Spec section | Task |
+|---|---|
+| Architecture: Phase A γ audit | Task 1 |
+| Architecture: Phase B β instrumentation (24→48 + checks) | Tasks 2, 3, 4, 5 |
+| Architecture: Phase C surgical fix | Task 6 |
+| Architecture: Phase D multi-fold validation | Task 7 |
+| Components A: audit doc | Task 1 |
+| Components B: expanded `nan_flags_buf` | Task 2 |
+| Components C: backward NaN check call sites | Task 4 |
+| Components D: surgical fix surface | Task 6 |
+| Components E: validation harness | Task 7 |
+| Operating principles: no deferrals + rich combined commits + ISV-driven | Stated in plan header + reinforced in Task 6 |
+| Success criterion 1 (kernel identified + fixed) | Task 6 |
+| Success criterion 2 (smoke pass criteria) | Task 7 |
+| Success criterion 3 (permanent infrastructure committed) | Tasks 2, 3, 4 |
+| Error handling (1) γ no smoking gun | Task 1 + Task 5 (β fires anyway) |
+| Error handling (2) cuBLAS wrapper case | Task 6 Template (f) |
+| Error handling (3) multi-source | Task 6 Step 1 + 7 commit-message language |
+| Error handling (4) F0 regression | Task 6 Step 2 (paper review) + Task 7 criterion 3 |
+| Error handling (5) F1→F2 different class | Task 7 Step 4 + decision tree |
+| Error handling (6) iteration unbounded | Task 7 Step 4 (return to Task 6) |
+| Closure documentation + memory + handoff | Task 8 |
+
+All spec requirements covered.
+
+**Placeholder scan results:**
+
+The plan contains intentional `<...>` placeholders in two places, both of which are KIND-OF-DATA placeholders (filled by audit/smoke output, not by lazy planning):
+
+- Task 4 step 1: kernel buffer pointers per slot — filled by Task 1 audit output
+- Task 6: kernel files patched, exact guard code — filled by Task 1 audit output + Task 5 smoke topology
+
+These are not "TBD" plan-failures — they are CONTENT-DRIVEN slots where the audit IS the discovery process. The plan provides the templates (Templates a-f in Task 6 Step 3), the constraints (ISV-driven, F0 paper review), and the exact form of the resulting commit. An engineer can execute this plan by:
+1. Running the audit (Task 1) to fill in the per-kernel data
+2. Following the templates in Task 6 once the audit + smoke topology identify the source
+
+This is the correct plan structure for an investigation-driven workflow. Pre-specifying the patch would be lying about what we know.
+
+**Type consistency check:**
+
+- `nan_flags_buf` size: 24 in Task 2 step 2 → 48 in Task 2 step 2 (after change). Used as 48 in all subsequent Task references. Consistent.
+- `read_nan_flags()` return type: `[i32; 24]` → `[i32; 48]` updated in Tasks 2 step 3 + 4. Both files (gpu_dqn_trainer.rs + fused_training.rs) updated. Consistent.
+- Method names: `run_nan_checks_post_backward` (Task 4) — used consistently in Task 4 step 2. Distinct from existing `run_nan_checks_post_forward`.
+- Buffer field names: `bw_d_h_s2_pre_saxpy` (Task 3) → accessor `bw_d_h_s2_pre_saxpy_ptr()` (Task 3 step 2) → used in Task 4 step 1 slot 32. Consistent.
+- Slot indices: 24-35 used consistently across spec table → Task 2 name table → Task 4 method.
+
+No type/name inconsistencies found.
+
+---
+
+## Plan complete and saved to `docs/superpowers/plans/2026-04-29-numerical-stability-sp1-f1-nan-root-cause.md`.
+
+Two execution options:
+
+1. **Subagent-Driven (recommended)** — fresh subagent per task, two-stage review (spec compliance + code quality) between tasks, fast iteration.
+
+2. **Inline Execution** — execute tasks in this session using superpowers:executing-plans, batch execution with checkpoints for review.
+
+Which approach?