# SP14 — Aux→Q Wire + Earned Gradient Flow Pearl 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:** Connect the SP13 supervised aux signal to the Q-head's directional decision-making via an Earned-Gradient-Flow-gated wire, plus three structural stability fixes Smoke A surfaced.

**Architecture:** Two sub-projects landing as separate atomic commits on `sp11-reward-as-controlled-subsystem`. Sub-project A (stability fixes) lands first as a cleanup commit (~50 LOC). Sub-project B (architectural wire + EGF pearl) lands as the main feature work (~1180 LOC). After B: L40S Smoke A2 validates the chain; if green, 30-epoch full validation tests the user's hypothesis that aux signal can move WR > 50%.

**Tech Stack:** Rust 1.85+ (Edition 2021), CUDA 12.4, cudarc 0.18, cuBLAS-Lt, pre-compiled cubins (no NVRTC), SQLx (offline), Tokio multi-threaded.

**Spec:** `docs/superpowers/specs/2026-05-05-sp14-aux-q-wire-earned-gradient-flow.md` v3 at HEAD `40e737a18`.

**Build:** `SQLX_OFFLINE=true cargo check --workspace` clean baseline.

**Test:** `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp14_oracle_tests --features cuda -- --ignored --nocapture` for new oracle tests on local RTX 3050 Ti (4GB).

**Smoke target:** `./scripts/argo-smoke.sh --gpu-pool ci-training-l40s` for `multi_fold_convergence::test_multi_fold_convergence` post-each-Layer.

---

## File Structure

### Layer A — Stability fixes

| File | Responsibility | Action |
|---|---|---|
| `crates/ml/src/cuda_pipeline/c51_grad_kernel.cu` | C51 distributional Q-loss backward | Modify: add `p_atom` floor before `−log(p)/p` divide; floor read from existing SP4 `atom_pos_p99` ISV slot |
| `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` | Trainer struct + `set_aux_weight` | Modify line 14722: lift clamp `[0.05, 0.3] → [AUX_W_BASE * 0.3, AUX_W_BASE * 3.0]` = `[0.15, 1.5]` |
| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | `compute_aux_w_p0b` controller | Modify lines 96-135: add `epochs_in_fold > 0` guard around stagnation logic |

### Layer B — Aux→Q wire + EGF pearl

| File | Responsibility | Action |
|---|---|---|
| `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` | 13 new ISV slot constants | **NEW** — mirror `sp13_isv_slots.rs` |
| `crates/ml/src/cuda_pipeline/q_disagreement_update_kernel.cu` | Per-step K=4↔K=2 disagreement EMA + Welford variance | **NEW** — ~100 LOC |
| `crates/ml/src/cuda_pipeline/alpha_grad_compute_kernel.cu` | Schmitt + adaptive k + adaptive β + α_grad rate limiter | **NEW** — ~100 LOC |
| `crates/ml/src/cuda_pipeline/gradient_hack_detect_kernel.cu` | Anti-gradient-hacking circuit breaker | **NEW** — ~40 LOC |
| `crates/ml/src/cuda_pipeline/dir_concat_qaux_kernel.cu` | Pre-SGEMM concat `[h_s2 ; aux_softmax_diff]` for direction Q-head | **NEW** — ~60 LOC; mirrors `mag_concat_qdir` precedent |
| `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` | Trainer struct, fingerprint, sizes, launches, scaling | Modify ~200 LOC |
| `crates/ml/src/trainers/dqn/state_reset_registry.rs` | Fold-reset entries | Modify ~50 LOC |
| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | HEALTH_DIAG emit + producer-chain launches | Modify ~80 LOC |
| `crates/ml/build.rs` | Register new cubins | Modify ~15 LOC |
| `crates/ml/tests/sp14_oracle_tests.rs` | GPU oracle tests for all new kernels + integration | **NEW** — ~600 LOC, 12-15 tests |
| `docs/dqn-wire-up-audit.md` | SP14 audit-doc section | Modify; required for pre-commit Invariant 7 hook |

---

# Layer A — Stability Fixes (lands first; ~50 LOC)

Three independent fixes. Each commits separately. Goal: clean gradient signals and aux_w controller behavior so that Layer B's smoke is interpretable.

## Task A.0: Phase 0 verification + worktree setup

**Files:**
- Verify: `crates/ml/src/cuda_pipeline/c51_grad_kernel.cu` (location of `−log(p)/p` divide)
- Verify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs:14720-14723` (set_aux_weight clamp)
- Verify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs:96-135` (compute_aux_w_p0b)
- Verify: `crates/ml/src/cuda_pipeline/sp13_isv_slots.rs` (AUX_W_BASE constant existence)

- [ ] **Step 1: Confirm worktree on the right branch**

```bash
cd /home/jgrusewski/Work/foxhunt/.worktrees/sp11-reward-controller
git branch --show-current
git log --oneline -5
```
Expected: branch `sp11-reward-as-controlled-subsystem`, HEAD `40e737a18` or later (spec v3 commit).

- [ ] **Step 2: Verify A.1 anchor — C51 atom probability divide**

```bash
grep -n "p_atom\|log(p\|projected_dist\|target_dist" crates/ml/src/cuda_pipeline/c51_grad_kernel.cu | head -20
```
Expected: locate the `−log(p)/p · ∂p/∂z` computation. Note exact line for Task A.1.

- [ ] **Step 3: Verify A.2 anchor — set_aux_weight clamp**

```bash
grep -n "set_aux_weight\|aux_weight.clamp" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs | head -5
```
Expected: `aux_weight.clamp(0.05, 0.3)` at line 14722.

- [ ] **Step 4: Verify A.3 anchor — compute_aux_w_p0b**

```bash
grep -n "fn compute_aux_w_p0b\|stagnation\|epochs_in_fold" crates/ml/src/trainers/dqn/trainer/training_loop.rs | head -15
```
Expected: function at line ~96, stagnation logic in body.

- [ ] **Step 5: Verify SP13 ISV constants for A.2**

```bash
grep -n "AUX_W_BASE\|AUX_W_HARD_FLOOR_RATIO\|AUX_W_HARD_CEIL_RATIO" crates/ml/src/cuda_pipeline/sp13_isv_slots.rs
```
Expected: constants exist with values `AUX_W_BASE = 0.5`, `HARD_FLOOR_RATIO = 0.3`, `HARD_CEIL_RATIO = 3.0`. If absent, define them in this task.

- [ ] **Step 6: Verify SP4 atom_pos_p99 source for A.1**

```bash
grep -n "ATOM_POS_P99\|atom_pos_p99" crates/ml/src/cuda_pipeline/sp4_isv_slots.rs crates/ml/src/cuda_pipeline/sp4_*.cu | head -10
```
Expected: per-branch `ATOM_POS_P99_*_INDEX` constants (4 branches). Note their slot numbers for Task A.1.

- [ ] **Step 7: Baseline build + test**

```bash
SQLX_OFFLINE=true cargo check --workspace 2>&1 | tail -5
SQLX_OFFLINE=true cargo test -p ml --lib snapshot_size_is_stable 2>&1 | tail -5
```
Expected: build clean (only pre-existing warnings); snap test passes at 149*4=596 bytes. Records pre-Layer-A baseline.

## Task A.1: C51 atom-probability floor (numerical stability)

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/c51_grad_kernel.cu` (around the `−log(p)/p · ∂p/∂z` divide)
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (kernel launch site — pass ATOM_POS_P99 indices)
- Test: `crates/ml/tests/sp14_oracle_tests.rs` (NEW; will host A's tests too)

**Why:** Smoke A showed C51 grad reaching `9.5×10⁶`, saturating SP7 budget controller. Root cause: when an atom near projected target has near-zero probability, `−log(p)/p · ∂p/∂z` blows up. Floor `p_atom >= ATOM_FLOOR` at source.

**ISV-driven floor:** SP4 already computes per-branch `atom_pos_p99` (the 99th percentile of atom probability mass). The floor is `max(1e-7, atom_pos_p99 * 1e-3)` — three orders of magnitude below the typical p99, ensuring it doesn't truncate normal gradients but cuts catastrophic spikes. Per `feedback_isv_for_adaptive_bounds`, the floor scales with the actual signal distribution.

- [ ] **Step 1: Read current c51_grad_kernel.cu**

```bash
sed -n '1,80p' crates/ml/src/cuda_pipeline/c51_grad_kernel.cu
grep -n "log\|divide\|p_atom\|p_target" crates/ml/src/cuda_pipeline/c51_grad_kernel.cu | head -15
```

Note the exact line where `−log(p)` or `p_target / p_atom` appears. The kernel signature likely takes `target_dist`, `pred_dist`, gradient outputs.

- [ ] **Step 2: Write the failing test (oracle test for floor behavior)**

Create `crates/ml/tests/sp14_oracle_tests.rs` with this initial test:

```rust
#![allow(unsafe_code)]
//! SP14 GPU oracle tests.
#![cfg(feature = "cuda")]

use std::sync::Arc;
use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg};
use ml::cuda_pipeline::mapped_pinned::{MappedF32Buffer, MappedI32Buffer};

const SP14_C51_GRAD_CUBIN: &[u8] =
    include_bytes!(concat!(env!("OUT_DIR"), "/c51_grad_kernel.cubin"));

fn make_test_stream() -> Arc<CudaStream> {
    let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?");
    ctx.default_stream()
}

fn load_c51_grad(stream: &Arc<CudaStream>) -> CudaFunction {
    let module = stream.context()
        .load_cubin(SP14_C51_GRAD_CUBIN.to_vec())
        .expect("load c51_grad cubin");
    module.load_function("c51_grad_kernel")
        .expect("load c51_grad_kernel function")
}

/// Test A.1: when target distribution has mass on an atom whose pred prob is
/// 1e-9 (below the floor 1e-7), the gradient must be bounded — not 1e9 — by
/// the floor. ATOM_POS_P99 set to 0.05 (typical), so floor = 0.05 * 1e-3 = 5e-5;
/// underflow protection at 1e-7 absolute floor.
#[test]
#[ignore = "requires GPU"]
fn c51_grad_floor_caps_underflow() {
    let stream = make_test_stream();
    let kernel = load_c51_grad(&stream);

    const N_ATOMS: usize = 51;
    const B: usize = 4;

    // Pred dist: 4 batch rows, 51 atoms each. Row 0: prob=1e-9 at atom 25.
    // Other rows: uniform.
    let mut pred = vec![1.0_f32 / N_ATOMS as f32; B * N_ATOMS];
    pred[0 * N_ATOMS + 25] = 1e-9_f32;  // The dangerous atom

    // Target dist: row 0 places mass on atom 25 (matching pred's tiny atom).
    let mut target = vec![0.0_f32; B * N_ATOMS];
    target[0 * N_ATOMS + 25] = 1.0_f32;

    // ISV slot for ATOM_POS_P99 (branch 0, direction): per SP4, returns ~0.05
    // for typical distributions. We synthesize this value directly.
    let mut isv_buf = vec![0.0_f32; 1024];
    isv_buf[/* ATOM_POS_P99_DIR_INDEX */ 220] = 0.05_f32;  // TODO verify slot
    // implementer: replace 220 with actual SP4 slot constant

    let pred_buf = unsafe { MappedF32Buffer::new(B * N_ATOMS) }.unwrap();
    let target_buf = unsafe { MappedF32Buffer::new(B * N_ATOMS) }.unwrap();
    let grad_buf = unsafe { MappedF32Buffer::new(B * N_ATOMS) }.unwrap();
    let isv = unsafe { MappedF32Buffer::new(1024) }.unwrap();
    pred_buf.write_from_slice(&pred);
    target_buf.write_from_slice(&target);
    isv.write_from_slice(&isv_buf);

    let bsi = B as i32;
    let nai = N_ATOMS as i32;

    unsafe {
        stream.launch_builder(&kernel)
            .arg(&pred_buf.dev_ptr)
            .arg(&target_buf.dev_ptr)
            .arg(&grad_buf.dev_ptr)
            .arg(&bsi)
            .arg(&nai)
            .arg(&isv.dev_ptr)
            // implementer: pass actual ATOM_POS_P99_*_INDEX values per branch
            .launch(LaunchConfig {
                grid_dim: (1, 1, 1),
                block_dim: (256, 1, 1),
                shared_mem_bytes: 0,
            })
            .expect("c51_grad_kernel launch");
    }
    unsafe { cudarc::driver::sys::cuStreamSynchronize(stream.cu_stream()); }

    let grad = grad_buf.read_all();
    // Without floor: grad[25] would be ~1e9. With floor at 5e-5: grad bounded.
    assert!(grad[25].abs() < 1e7,
        "C51 grad should be bounded by atom floor; got {} for row 0 atom 25", grad[25]);
}
```

- [ ] **Step 3: Run test to verify it fails**

```bash
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp14_oracle_tests \
    c51_grad_floor_caps_underflow --features cuda -- --ignored --nocapture 2>&1 | tail -10
```

Expected: PRE-FIX state — assertion fails (grad >> 1e7) OR the kernel hasn't been modified yet (so test passes accidentally). If passes, the test is too weak — strengthen by adding `assert!(grad[25].abs() < 1e6)`.

- [ ] **Step 4: Implement the floor in c51_grad_kernel.cu**

Find the divide-by-pred line (Step 1 located it). Add the floor:

```cuda
// SP14 Layer A fix: floor p_atom to prevent -log(p)/p underflow gradients.
// Floor source: ISV[ATOM_POS_P99_DIR_INDEX] * 1e-3, hard-floored at 1e-7
// for absolute fp32 underflow safety. Per pearl_isv_for_adaptive_bounds:
// the floor scales with the actual signal distribution rather than being
// a hardcoded constant.
const float atom_pos_p99 = isv[atom_pos_p99_index];  // implementer: pass per-branch
const float ATOM_FLOOR = fmaxf(1e-7f, atom_pos_p99 * 1e-3f);
const float p_safe = fmaxf(p_atom_at_target, ATOM_FLOOR);
// Existing computation, with p_safe replacing raw p:
const float grad_atom = -target_prob * __logf(p_safe) / p_safe;
```

The kernel signature gains `const float* isv` and per-branch `int atom_pos_p99_index_b0..b3` arguments. Or simpler: pass single `atom_pos_p99` value pre-fetched on host (cheaper since this is invariant across batch within an epoch — the implementer chooses based on launch frequency).

- [ ] **Step 5: Update c51_grad_kernel launch site in gpu_dqn_trainer.rs**

```bash
grep -n "c51_grad_kernel\b" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs | head -5
```

At each launch site, add `.arg(&isv_buf)` and the per-branch `.arg(&atom_pos_p99_index_*)` arguments. Verify that `isv_buf` is the same buffer SP4 producers write to.

- [ ] **Step 6: Run test to verify it passes**

```bash
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp14_oracle_tests \
    c51_grad_floor_caps_underflow --features cuda -- --ignored --nocapture 2>&1 | tail -10
```

Expected: PASS — `grad[25].abs() < 1e7`.

- [ ] **Step 7: Run full workspace check**

```bash
SQLX_OFFLINE=true cargo check --workspace 2>&1 | tail -10
```

Expected: clean (only pre-existing warnings).

- [ ] **Step 8: Commit Layer A.1**

```bash
git add crates/ml/src/cuda_pipeline/c51_grad_kernel.cu \
        crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \
        crates/ml/tests/sp14_oracle_tests.rs
git commit -m "$(cat <<'EOF'
fix(sp14): A.1 — C51 atom-probability floor (ISV-driven)

Floor p_atom = max(1e-7, ATOM_POS_P99 * 1e-3) before -log(p)/p divide
in c51_grad_kernel. Cuts C51 distributional Q-loss gradient blow-up
from 9.5e6 (pre-fix) to bounded magnitudes that SP7 budget controller
can rebalance proportionally instead of pinning to EPS_DIV=1e-8.

Per pearl_isv_for_adaptive_bounds, floor scales with actual atom
distribution width (SP4 atom_pos_p99 producer) rather than being a
hardcoded constant. The 1e-7 absolute floor is the fp32 underflow
safety; the 1e-3 multiplier of p99 is the structural numerical-
stability anchor (Invariant 1).

Smoke A had 1109 GRAD_CLIP_OUTLIER events from C51 raw grad spiking;
expected count post-fix: <100, similar to pre-SP13 baseline of 5.

Test: c51_grad_floor_caps_underflow GPU oracle on RTX 3050 Ti passes;
verifies grad[25] < 1e7 when target mass on atom with pred prob 1e-9.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
EOF
)"
```

## Task A.2: aux_w setter clamp lift

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs:14720-14723`
- Verify constants in: `crates/ml/src/cuda_pipeline/sp13_isv_slots.rs`

**Why:** P0b's deficit+stagnation controller computes raw aux_w in `[0.15, 1.5]` (base 0.5 × [0.3, 3.0]). The setter at 14722 silently re-clamps to the SP11-era `[0.05, 0.3]`. Smoke A confirmed: aux_w pinned at 0.3 in folds 0/1 (clamp masking the controller's deficit-amp), drops to 0.164 in fold 2 (45% lower than working value).

- [ ] **Step 1: Read the current clamp**

```bash
sed -n '14718,14725p' crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
```
Expected: `pub fn set_aux_weight(&mut self, aux_weight: f32) { ... aux_weight.clamp(0.05, 0.3) ... }`

- [ ] **Step 2: Verify the SP13 constants exist**

```bash
grep -A 2 "AUX_W_BASE\|AUX_W_HARD_FLOOR_RATIO\|AUX_W_HARD_CEIL_RATIO" crates/ml/src/cuda_pipeline/sp13_isv_slots.rs
```

If they exist with values 0.5, 0.3, 3.0 — proceed. If absent, define them in `sp13_isv_slots.rs` first:

```rust
/// SP13 P0b base aux_w. Per-step controller output is `base × deficit_amp × (1−stagnation_decay)`.
pub const AUX_W_BASE: f32 = 0.5;
/// Lower bound for the post-clamp aux_w (= base × ratio = 0.15).
pub const AUX_W_HARD_FLOOR_RATIO: f32 = 0.3;
/// Upper bound for the post-clamp aux_w (= base × ratio = 1.5). Lifted from
/// SP11-era 0.3 cap in SP14 Layer A.
pub const AUX_W_HARD_CEIL_RATIO: f32 = 3.0;
```

- [ ] **Step 3: Write the failing test**

Append to `crates/ml/tests/sp14_oracle_tests.rs`:

```rust
/// Test A.2: set_aux_weight at SP13 base 0.5 should accept values up to 1.5,
/// not silently clamp to 0.3 (the SP11-era cap).
#[test]
fn set_aux_weight_clamp_range() {
    use ml::cuda_pipeline::sp13_isv_slots::{
        AUX_W_BASE, AUX_W_HARD_FLOOR_RATIO, AUX_W_HARD_CEIL_RATIO,
    };
    // Direct math check; not a GPU test.
    let floor = AUX_W_BASE * AUX_W_HARD_FLOOR_RATIO;
    let ceil = AUX_W_BASE * AUX_W_HARD_CEIL_RATIO;
    assert_eq!(floor, 0.15);
    assert_eq!(ceil, 1.5);

    // Verify the trainer struct's clamp uses these bounds. We use a direct
    // f32 clamp here mirroring the production code; the test ensures the
    // production constants resolve to the expected values.
    let test_value: f32 = 1.0;
    let clamped = test_value.clamp(floor, ceil);
    assert_eq!(clamped, 1.0,
        "1.0 should NOT be clamped under [0.15, 1.5]; got {}", clamped);

    let too_low: f32 = 0.05;
    let clamped_low = too_low.clamp(floor, ceil);
    assert_eq!(clamped_low, 0.15);

    let too_high: f32 = 2.0;
    let clamped_high = too_high.clamp(floor, ceil);
    assert_eq!(clamped_high, 1.5);
}
```

- [ ] **Step 4: Run test (should pass — it's a math check independent of trainer state)**

```bash
SQLX_OFFLINE=true cargo test -p ml --test sp14_oracle_tests set_aux_weight_clamp_range 2>&1 | tail -5
```

Expected: PASS. The test verifies the constants resolve correctly.

- [ ] **Step 5: Modify the trainer's set_aux_weight**

Edit `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` line 14722:

Before:
```rust
self.aux_weight = aux_weight.clamp(0.05, 0.3);
```

After:
```rust
// SP14 Layer A: lift clamp from SP11-era [0.05, 0.3] to SP13 P0b design
// range [0.15, 1.5]. The pre-SP14 cap silently masked the deficit-amp
// term in compute_aux_w_p0b — controller computed [0.15, 1.5] but values
// > 0.3 got chopped, leaving fold 2 stagnation 45% below working value.
use crate::cuda_pipeline::sp13_isv_slots::{
    AUX_W_BASE, AUX_W_HARD_FLOOR_RATIO, AUX_W_HARD_CEIL_RATIO,
};
self.aux_weight = aux_weight.clamp(
    AUX_W_BASE * AUX_W_HARD_FLOOR_RATIO,  // 0.15
    AUX_W_BASE * AUX_W_HARD_CEIL_RATIO,   // 1.5
);
```

- [ ] **Step 6: Build check**

```bash
SQLX_OFFLINE=true cargo check --workspace 2>&1 | tail -5
```
Expected: clean.

- [ ] **Step 7: Commit Layer A.2**

```bash
git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \
        crates/ml/src/cuda_pipeline/sp13_isv_slots.rs \
        crates/ml/tests/sp14_oracle_tests.rs
git commit -m "$(cat <<'EOF'
fix(sp14): A.2 — lift set_aux_weight clamp to SP13 P0b range [0.15, 1.5]

The pre-fix clamp at gpu_dqn_trainer.rs:14722 was [0.05, 0.3] — the
SP11-era cap. P0b's controller computes aux_w in [0.15, 1.5] (base
0.5 × [0.3, 3.0]) but the setter silently chopped everything above
0.3, masking the deficit-amplification term entirely.

Smoke A trace confirmed:
  Fold 0/1: raw aux_w = 0.66-0.80 → clamped to 0.30 (deficit invisible)
  Fold 2:   raw aux_w = 0.164 (stagnation; below clamp) → 45% deficit

Post-fix: deficit-amp term `(1 + 5 × deficit)` actually expresses
through to the trainer. Fold 2 stagnation will get the designed
floor 0.15 instead of being capped at 0.3 — but the upper range
1.5 also opens, so deficit-amp can pull aux_w up when accuracy is
below target.

Constants imported from sp13_isv_slots.rs (AUX_W_BASE=0.5,
AUX_W_HARD_FLOOR_RATIO=0.3, AUX_W_HARD_CEIL_RATIO=3.0). No new
slots; existing constants exposed as the clamp bounds.

Test: set_aux_weight_clamp_range verifies constants resolve correctly.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
EOF
)"
```

## Task A.3: Stagnation warmup gate

**Files:**
- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs:96-135` (`compute_aux_w_p0b`)

**Why:** Pearl-A first-observation bootstrap makes both EMAs equal at fold reset (sentinel 0.5 → first observation X → both short_ema = long_ema = X). This makes `improvement = max(0, short - long) = 0` by construction in epoch 0, triggering stagnation logic on a non-stagnation. Fix: skip stagnation when `epochs_in_fold < 1`.

- [ ] **Step 1: Read current compute_aux_w_p0b**

```bash
sed -n '96,140p' crates/ml/src/trainers/dqn/trainer/training_loop.rs
```

Note the exact stagnation formula and where `epochs_in_fold` (or equivalent) is available.

- [ ] **Step 2: Identify the epochs_in_fold tracker**

```bash
grep -n "epochs_in_fold\|epoch_in_fold\|fold_epoch\|since_fold_reset" crates/ml/src/trainers/dqn/trainer/training_loop.rs | head -15
```

Identify the field name. If absent, the trainer tracks current epoch and fold start — derive `epochs_in_fold = current_epoch - fold_start_epoch`.

- [ ] **Step 3: Write the failing test**

Append to `crates/ml/tests/sp14_oracle_tests.rs`:

```rust
/// Test A.3: in epoch 0 of a fresh fold (Pearl-A bootstrap state where
/// short_ema = long_ema by construction), the stagnation term should be
/// suppressed (= 0), preserving aux_w at base × deficit_amp instead of
/// collapsing to floor.
#[test]
fn aux_w_stagnation_warmup_gate_epoch_0() {
    // We test compute_aux_w_p0b directly; it's a pure function of (target,
    // short_ema, long_ema, epochs_in_fold).
    use ml::trainers::dqn::trainer::training_loop::compute_aux_w_p0b;

    let target = 0.55_f32;
    let short_ema = 0.50_f32;  // Pearl-A first-observation: short = long
    let long_ema = 0.50_f32;
    let deficit = (target - short_ema).max(0.0);  // = 0.05
    let deficit_amp = 1.0 + 5.0 * deficit;        // = 1.25

    // In epoch 0, stagnation should be SKIPPED (= 0).
    let aux_w_e0 = compute_aux_w_p0b(target, short_ema, long_ema, /* epochs_in_fold */ 0);
    let expected_e0 = (0.5 * deficit_amp).clamp(0.15, 1.5);  // 0.5 × 1.25 = 0.625
    assert!((aux_w_e0 - expected_e0).abs() < 1e-5,
        "epoch 0 (warmup) should skip stagnation; got {} expected {}", aux_w_e0, expected_e0);

    // In epoch 1+, stagnation IS evaluated. With short==long (improvement=0),
    // stagnation = 1.0, and aux_w drops to base × deficit_amp × (1 - 0.7×1) = 0.187.
    let aux_w_e1 = compute_aux_w_p0b(target, short_ema, long_ema, 1);
    let expected_e1 = (0.5 * deficit_amp * (1.0 - 0.7 * 1.0)).clamp(0.15, 1.5);
    assert!((aux_w_e1 - expected_e1).abs() < 1e-5,
        "epoch 1 (post-warmup) should include stagnation; got {} expected {}",
        aux_w_e1, expected_e1);
}
```

- [ ] **Step 4: Run test to verify it fails**

```bash
SQLX_OFFLINE=true cargo test -p ml --test sp14_oracle_tests aux_w_stagnation_warmup_gate_epoch_0 2>&1 | tail -10
```

Expected: FAIL — pre-fix the function doesn't take `epochs_in_fold` OR ignores it; aux_w_e0 ≠ expected.

- [ ] **Step 5: Add epochs_in_fold gate to compute_aux_w_p0b**

Find the function signature. If `epochs_in_fold` isn't already a parameter, add it. Modify the stagnation logic:

```rust
pub fn compute_aux_w_p0b(
    target: f32,
    short_ema: f32,
    long_ema: f32,
    epochs_in_fold: usize,
) -> f32 {
    use crate::cuda_pipeline::sp13_isv_slots::{
        AUX_W_BASE, AUX_W_HARD_FLOOR_RATIO, AUX_W_HARD_CEIL_RATIO,
    };
    let deficit = (target - short_ema).max(0.0);
    let deficit_amp = 1.0 + 5.0 * deficit;

    // SP14 Layer A: stagnation gate skips epoch 0 of each fold.
    // Pearl-A first-observation makes short_ema = long_ema at fold reset,
    // giving improvement = 0 by construction — the stagnation gate would
    // fire on a non-stagnation. Wait one full epoch for EMAs to diverge.
    let stagnation = if epochs_in_fold < 1 {
        0.0_f32
    } else {
        let improvement = (short_ema - long_ema).max(0.0);
        let normalised = improvement / deficit.max(0.005);
        (1.0 - normalised).clamp(0.0, 1.0)
    };

    let raw = AUX_W_BASE * deficit_amp * (1.0 - 0.7 * stagnation);
    raw.clamp(
        AUX_W_BASE * AUX_W_HARD_FLOOR_RATIO,  // 0.15
        AUX_W_BASE * AUX_W_HARD_CEIL_RATIO,   // 1.5
    )
}
```

Update all call sites to pass `epochs_in_fold` (find via `grep -n "compute_aux_w_p0b" crates/ml`).

- [ ] **Step 6: Run test to verify it passes**

```bash
SQLX_OFFLINE=true cargo test -p ml --test sp14_oracle_tests aux_w_stagnation_warmup_gate_epoch_0 2>&1 | tail -5
```
Expected: PASS.

- [ ] **Step 7: Run full lib tests for regression check**

```bash
SQLX_OFFLINE=true cargo test -p ml --lib 2>&1 | tail -5
```
Expected: pre-existing failures unchanged; no new failures from this change.

- [ ] **Step 8: Commit Layer A.3**

```bash
git add crates/ml/src/trainers/dqn/trainer/training_loop.rs \
        crates/ml/tests/sp14_oracle_tests.rs
git commit -m "$(cat <<'EOF'
fix(sp14): A.3 — stagnation warmup gate at fold boundary

compute_aux_w_p0b's stagnation term was firing inappropriately at
fold reset because Pearl-A first-observation bootstrap makes both
EMAs equal:
  short_ema = sentinel 0.5 → first observation X (after one update)
  long_ema  = sentinel 0.5 → first observation X (same update)
  improvement = max(0, short - long) = 0
  stagnation = (1 - 0/max(deficit, 0.005)) = 1.0
  aux_w *= (1 - 0.7) = 0.3 → spurious decay on a non-stagnation

Fix: skip stagnation logic when epochs_in_fold < 1. Wait one full
epoch for the short/long EMA timescale split (α=0.3 vs 0.05) to
produce real `improvement` signal.

Test: aux_w_stagnation_warmup_gate_epoch_0 verifies:
  - Epoch 0: stagnation = 0, aux_w = base × deficit_amp = 0.625
  - Epoch 1: stagnation = 1.0, aux_w = base × deficit_amp × 0.3 = 0.187

Combined with A.2 (clamp lift), Fold 2's aux_w should now hit 0.625
in epoch 0 (the controller's intended post-deficit-amp value) instead
of being immediately collapsed to 0.164.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
EOF
)"
```

## Task A.4: Smoke A2-A validation (5-epoch L40S)

**Files:**
- No edits — runs scripts/argo-smoke.sh against current HEAD

- [ ] **Step 1: Push to origin (per feedback_push_before_deploy)**

```bash
git push origin sp11-reward-as-controlled-subsystem 2>&1 | tail -5
```
Expected: 3 new commits pushed (A.1, A.2, A.3).

- [ ] **Step 2: Submit Smoke A2-A**

```bash
./scripts/argo-smoke.sh 2>&1 | tail -10
```
Expected: workflow submitted, name like `smoke-test-XXXXX`. Note the name.

- [ ] **Step 3: Wait + monitor with kill criteria**

Workflow runs ~25-30 min. Monitor via `argo logs <wfname> -n foxhunt --tail=200` for:
- Build/compile failure → immediate kill via `argo terminate <wfname> -n foxhunt`
- `GRAD_CLIP_OUTLIER` count: count occurrences in fold 2 logs. Target: < 100 (down from 1109 pre-fix)
- `aux_w` in HEALTH_DIAG: should reach 0.625-0.8 range in fold 0 ep 0-2 (pre-fix was pinned at 0.3)
- val_sharpe: should stay positive across all folds (pre-fix was 9-16)

- [ ] **Step 4: Verify pass criteria from MinIO logs (post-completion)**

```bash
kubectl port-forward -n foxhunt svc/minio 9000:9000 &
sleep 3
mc alias set foxminio http://127.0.0.1:9000 \
    "$(kubectl get secret minio-credentials -n foxhunt -o json | jq -r '.data["access-key"]' | base64 -d)" \
    "$(kubectl get secret minio-credentials -n foxhunt -o json | jq -r '.data["secret-key"]' | base64 -d)"

mc cat foxminio/argo-logs/<wfname>/<wfname>/main.log 2>&1 | \
    grep -c "GRAD_CLIP_OUTLIER" 

mc cat foxminio/argo-logs/<wfname>/<wfname>/main.log 2>&1 | \
    grep "aux_w=" | head -10

pkill -f "kubectl port-forward.*minio"
```

Expected:
- GRAD_CLIP_OUTLIER count < 100 (target ≤ 50)
- aux_w varies in [0.15, 1.5] range, reaches values > 0.3 in deficit-amp regime
- All val_sharpe positive

- [ ] **Step 5: Document results in audit doc**

Append to `docs/dqn-wire-up-audit.md` a new section "## SP14 Layer A — stability fixes (2026-05-XX)" documenting the smoke results, pre/post comparison tables, and "Next: Layer B".

- [ ] **Step 6: Commit audit doc**

```bash
git add docs/dqn-wire-up-audit.md
git commit -m "docs(sp14): A — Smoke A2-A validation results

Layer A stability fixes validated on L40S smoke <wfname>:
- GRAD_CLIP_OUTLIER count: <count> (vs 1109 pre-fix; <ratio> reduction)
- aux_w trace shows controller now reaching designed range [0.15, 1.5]
- val_sharpe across all folds: <values> (positive, baseline preserved)
- Stagnation warmup gate: epoch 0 of folds shows aux_w ≈ 0.625 (deficit-amp expressing)

Layer A complete. Layer B (architectural wire + EGF pearl) is next.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>"
```

---

# Layer B — Aux→Q Wire + Earned Gradient Flow Pearl (~1180 LOC, atomic feature)

This is the architectural piece. Lands as a sequence of commits, each safe individually (additive infrastructure → wire forward → wire backward → integration), with the FINAL state being the atomic feature. Per `feedback_no_partial_refactor`: each intermediate state must be valid (compiles + tests pass), and the wire is fully gated by the time it's active.

## Task B.0: Phase 0 verification

**Files:**
- Verify all line numbers and tensor indices the spec references against current code at HEAD `40e737a18` or later

- [ ] **Step 1: Verify direction Q-head's first FC param tensor index**

```bash
sed -n '2770,2790p' crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
grep -n "w_b0fc\|w_b0out\|b_b0fc\|b_b0out" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs | head -10
```

Identify the exact tensor index for `w_b0fc` (direction Q-head first FC weight). Note its formula in `compute_param_sizes()` — the input dim grows by 1 in this plan.

- [ ] **Step 2: Verify aux_nb_softmax_buf field**

```bash
grep -n "aux_nb_softmax_buf\b" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs | head -10
```

Confirm field type `CudaSlice<f32>` shape `[B, K=2]` post-B1.1a. Note the field's allocation site for fingerprint context.

- [ ] **Step 3: Verify DIR_* constants**

```bash
sed -n '120,130p' crates/ml/src/cuda_pipeline/state_layout.cuh
```

Expected:
```
#define DIR_SHORT  0
#define DIR_HOLD   1
#define DIR_LONG   2
#define DIR_FLAT   3
```

- [ ] **Step 4: Verify mag_concat_qdir precedent for forward concat pattern**

```bash
sed -n '4555,4625p' crates/ml/src/cuda_pipeline/experience_kernels.cu
```

Read the kernel signature; this is the template for `dir_concat_qaux_kernel`.

- [ ] **Step 5: Verify SP4 atom_pos_p99 + SP13 aux_nb_softmax accessibility**

```bash
grep -n "ATOM_POS_P99_DIR_INDEX\|atom_pos_p99_dir" crates/ml/src/cuda_pipeline/sp4_isv_slots.rs
```

Note the exact slot constant name to import in the new kernels.

- [ ] **Step 6: Verify state_reset_registry pattern**

```bash
sed -n '900,930p' crates/ml/src/trainers/dqn/state_reset_registry.rs
```

Read the existing aux_dir_acc EMA entries (lines 913-922). This is the template for the 11 new entries.

- [ ] **Step 7: Verify HEALTH_DIAG aux line for emit pattern**

```bash
grep -n "aux_dir_acc target=\|HEALTH_DIAG.*aux" crates/ml/src/trainers/dqn/trainer/training_loop.rs | head -5
```

Locate the existing aux HEALTH_DIAG emit (around line 4858-4891). The `pearl_egf_diag` line will be appended in this region.

- [ ] **Step 8: Confirm baseline build + test pass**

```bash
SQLX_OFFLINE=true cargo check --workspace 2>&1 | tail -5
SQLX_OFFLINE=true cargo test -p ml --test sp14_oracle_tests --features cuda 2>&1 | tail -5
```

Expected: clean build; A's three tests pass.

## Task B.1: Create sp14_isv_slots.rs

**Files:**
- Create: `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs`
- Modify: `crates/ml/src/cuda_pipeline/mod.rs` (add `pub mod sp14_isv_slots;`)
- Modify: `crates/ml/src/cuda_pipeline/sp4_isv_slots.rs` (add `pub const SP14_RANGE_FIRST` / `SP14_RANGE_LAST` if SP4 owns the global ISV-slot inventory)

**Why:** 13 new ISV slots need stable global indices. Following the SP-series pattern (`sp13_isv_slots.rs`), all SP14 slot constants live in one file.

**Slot index allocation:** Identify the next available block. SP13 used 372-380 + 117 retired. SP14 will claim a contiguous block starting after SP13's range.

- [ ] **Step 1: Identify next available ISV slot range**

```bash
grep -hE "INDEX:\s*usize\s*=\s*[0-9]+" crates/ml/src/cuda_pipeline/sp*_isv_slots.rs | \
    sed -E 's/.*=\s*([0-9]+).*/\1/' | sort -n | tail -20
```

Note the highest current index. Allocate SP14 slots starting at `<highest> + 1` and reserve 16 slots (13 used + 3 buffer).

- [ ] **Step 2: Create sp14_isv_slots.rs**

```rust
//! SP14 ISV slot constants.
//!
//! 13 slots for the Aux→Q Wire + Earned Gradient Flow pearl. See
//! `docs/superpowers/specs/2026-05-05-sp14-aux-q-wire-earned-gradient-flow.md`
//! for the full design rationale; in brief:
//!
//! - Q_DISAGREEMENT_*: K=4↔K=2 mapped per-step argmax mismatch (B.2.3)
//! - K_*_ADAPTIVE: ISV-driven sigmoid steepness from Welford variance (B.2.5)
//! - BETA_RATE_LIMITER_ADAPTIVE: ISV-driven rate-limiter β (B.2.8)
//! - AUX_DIR_ACC_VARIANCE_EMA + Q_DISAGREEMENT_VARIANCE_EMA + ALPHA_GRAD_RAW_VARIANCE_EMA:
//!   Welford variance EMAs driving the three adaptive scalars
//! - GATE1_OPEN_STATE: Schmitt-trigger persistent state (0=closed, 1=open)
//! - ALPHA_GRAD_RAW: raw gate output (HEALTH_DIAG visibility)
//! - ALPHA_GRAD_SMOOTHED: rate-limited gate consumed by backward path
//! - AUX_DIR_ACC_POST_OPEN_MIN + GRADIENT_HACK_LOCKOUT_REMAINING: anti-gradient-hacking circuit breaker

// Q-disagreement EMAs (K=4↔K=2 mapping with Hold/Flat masked)
pub const Q_DISAGREEMENT_SHORT_EMA_INDEX:    usize = /* allocate per Step 1 */ 381;
pub const Q_DISAGREEMENT_LONG_EMA_INDEX:     usize = 382;

// Adaptive sigmoid steepness (variance-driven)
pub const K_AUX_ADAPTIVE_INDEX:              usize = 383;
pub const K_Q_ADAPTIVE_INDEX:                usize = 384;

// Adaptive rate-limiter β (variance-driven; B.2.8)
pub const BETA_RATE_LIMITER_ADAPTIVE_INDEX:  usize = 385;

// Welford variance EMAs (drivers for adaptive k_aux, k_q, β)
pub const AUX_DIR_ACC_VARIANCE_EMA_INDEX:    usize = 386;
pub const Q_DISAGREEMENT_VARIANCE_EMA_INDEX: usize = 387;
pub const ALPHA_GRAD_RAW_VARIANCE_EMA_INDEX: usize = 388;

// Schmitt-trigger persistent state (0=closed, 1=open)
pub const GATE1_OPEN_STATE_INDEX:            usize = 389;

// α_grad outputs (raw + rate-limited)
pub const ALPHA_GRAD_RAW_INDEX:              usize = 390;
pub const ALPHA_GRAD_SMOOTHED_INDEX:         usize = 391;

// Anti-gradient-hacking circuit breaker
pub const AUX_DIR_ACC_POST_OPEN_MIN_INDEX:   usize = 392;
pub const GRADIENT_HACK_LOCKOUT_REMAINING_INDEX: usize = 393;

// ── Sentinels (initial values; reset at fold boundary) ─────────────────
pub const SENTINEL_Q_DISAGREEMENT:           f32 = 0.5;  // analytic K=4↔K=2 random alignment
pub const SENTINEL_VARIANCE:                 f32 = 0.0;  // → initial k = k_base, β = β_base
pub const SENTINEL_GATE1_CLOSED:             f32 = 0.0;
pub const SENTINEL_ALPHA_GRAD:               f32 = 0.0;
pub const SENTINEL_POST_OPEN_MIN:            f32 = 1.0;  // "no min observed"
pub const SENTINEL_LOCKOUT:                  f32 = 0.0;

// ── Structural constants (numerical-stability anchors per Invariant 1) ─
/// Sigmoid steepness base for Gate 1 (aux competence). At reliable signal
/// (variance ≈ 0), k_aux = K_BASE_AUX. Decays toward K_MIN under noise.
pub const K_BASE_AUX:                        f32 = 20.0;
/// Sigmoid steepness base for Gate 2 (Q-head disagreement).
pub const K_BASE_Q:                          f32 = 15.0;
/// Lower bound for adaptive k (numerical floor; prevents flat sigmoid).
pub const K_MIN:                             f32 = 1.0;
/// Variance reference for k_aux adaptation (1% deviation → noisy regime).
pub const VARIANCE_REF_AUX:                  f32 = 0.01;
/// Variance reference for k_q adaptation (5% deviation → noisy regime).
pub const VARIANCE_REF_Q:                    f32 = 0.05;
/// Variance reference for β adaptation (1% α_grad_raw variance → smoother).
pub const VARIANCE_REF_ALPHA:                f32 = 0.01;

/// Rate-limiter β base (light smoothing; ~2-step half-life on stable signal).
pub const BETA_BASE:                         f32 = 0.5;
/// Rate-limiter β max (heavy smoothing under volatile α_grad_raw).
pub const BETA_MAX:                          f32 = 0.95;

/// Schmitt-trigger hysteresis half-band: open at target+0.03, close at target-0.03.
pub const SCHMITT_BAND:                      f32 = 0.03;

/// Per-epoch warmup ramp duration (steps). Set at runtime from config; this
/// constant is the FALLBACK if config lookup fails.
pub const WARMUP_STEPS_FALLBACK:             usize = 1000;

/// Anti-gradient-hacking circuit breaker: aux_dir_acc must drop more than
/// LOCKOUT_TRIGGER_DROP below open-threshold to suspect hacking.
pub const LOCKOUT_TRIGGER_DROP:              f32 = 0.05;
/// q_disagreement must rise more than LOCKOUT_TRIGGER_DIS_RISE above baseline
/// concurrently with the aux drop to confirm hacking suspicion.
pub const LOCKOUT_TRIGGER_DIS_RISE:          f32 = 0.10;
/// Lockout duration (epochs).
pub const LOCKOUT_EPOCHS:                    f32 = 2.0;

/// Q-disagreement baseline (analytic random alignment under Hold/Flat masking).
pub const Q_DISAGREEMENT_BASELINE:           f32 = 0.5;
```

- [ ] **Step 3: Add module declaration**

Edit `crates/ml/src/cuda_pipeline/mod.rs`:

```rust
pub mod sp14_isv_slots;
```

- [ ] **Step 4: Build check**

```bash
SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
```
Expected: clean. If a slot index collides with existing SP*, the compiler may not catch it (they're separate consts). Manually grep to confirm:

```bash
grep -rh "INDEX:\s*usize\s*=\s*381" crates/ml/src/
```

Should only return the new SP14 line.

- [ ] **Step 5: Commit B.1**

```bash
git add crates/ml/src/cuda_pipeline/sp14_isv_slots.rs \
        crates/ml/src/cuda_pipeline/mod.rs
git commit -m "feat(sp14): B.1 — sp14_isv_slots.rs with 13 new ISV slot constants

Allocates ISV slots [381..394) for the Aux→Q Wire + Earned Gradient
Flow pearl. Mirrors sp13_isv_slots.rs pattern.

Slots fall into 4 functional groups:
- Q-disagreement EMAs (short, long; K=4↔K=2 mapping with Hold/Flat masked)
- Adaptive controllers (k_aux, k_q, β; variance-driven)
- Welford variance EMAs (3, one per adaptive scalar)
- Schmitt state + α_grad outputs + circuit breaker

Plus 14 structural constants for numerical-stability anchors:
K_BASE_*, K_MIN, VARIANCE_REF_*, BETA_BASE, BETA_MAX, SCHMITT_BAND,
WARMUP_STEPS_FALLBACK, LOCKOUT_*, Q_DISAGREEMENT_BASELINE.

Per feedback_isv_for_adaptive_bounds: adaptive bounds (k_*, β,
post_open_min, lockout) live in ISV; numerical anchors live as
structural constants. Per pearl_first_observation_bootstrap: all
EMAs reset to sentinels and Pearl-A bootstraps on first observation.

Producer + consumer wiring lands in subsequent tasks (B.4-B.12);
this commit is additive infrastructure only.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>"
```

## Task B.2: state_reset_registry entries (11 new)

**Files:**
- Modify: `crates/ml/src/trainers/dqn/state_reset_registry.rs` (append entries; verify line ~922 is end of SP13 entries)

**Why:** All SP14 EMAs and stateful slots must reset at fold boundary per Pearl-A bootstrap discipline. The registry is GPU-resident (per `pearl_isv_for_adaptive_bounds` pattern); reset_at_fold launches a `state_reset_kernel` over registered slots.

- [ ] **Step 1: Read existing pattern**

```bash
sed -n '910,935p' crates/ml/src/trainers/dqn/state_reset_registry.rs
```

The two SP13 aux_dir_acc EMA entries (lines 913-922) are the template.

- [ ] **Step 2: Append SP14 registry entries**

Locate the end of the registry array and append 11 entries. Each entry has shape:

```rust
RegistryEntry {
    name: "sp14_q_disagreement_short_ema",
    slot_index: ml::cuda_pipeline::sp14_isv_slots::Q_DISAGREEMENT_SHORT_EMA_INDEX,
    sentinel: ml::cuda_pipeline::sp14_isv_slots::SENTINEL_Q_DISAGREEMENT,
    description: "ISV[Q_DISAGREEMENT_SHORT_EMA_INDEX=381] — SP14 EGF Gate 2 fast EMA (α=0.3) of per-step Q↔aux argmax disagreement after K=4↔K=2 mapping with Hold/Flat masked. Produced by `q_disagreement_update_kernel` per training step. Consumed by `alpha_grad_compute_kernel` as Gate 2 sigmoid driver. FoldReset sentinel 0.5 (analytic random alignment under masking) so Pearl A bootstraps from new fold's first observation per `pearl_first_observation_bootstrap.md`.",
},
```

Repeat for all 11 slots that need fold-reset:

1. Q_DISAGREEMENT_SHORT_EMA — sentinel 0.5
2. Q_DISAGREEMENT_LONG_EMA — sentinel 0.5
3. K_AUX_ADAPTIVE — sentinel `K_BASE_AUX = 20.0`
4. K_Q_ADAPTIVE — sentinel `K_BASE_Q = 15.0`
5. BETA_RATE_LIMITER_ADAPTIVE — sentinel `BETA_BASE = 0.5`
6. AUX_DIR_ACC_VARIANCE_EMA — sentinel 0.0
7. Q_DISAGREEMENT_VARIANCE_EMA — sentinel 0.0
8. ALPHA_GRAD_RAW_VARIANCE_EMA — sentinel 0.0
9. GATE1_OPEN_STATE — sentinel 0.0 (closed)
10. ALPHA_GRAD_SMOOTHED — sentinel 0.0
11. AUX_DIR_ACC_POST_OPEN_MIN — sentinel 1.0

(`ALPHA_GRAD_RAW`, `GRADIENT_HACK_LOCKOUT_REMAINING`: implementer's call whether they reset; raw is recomputed every step so reset isn't strictly needed; lockout naturally decays at epoch boundaries. Adding both for safety brings count to 13; the spec says 11 reset, 2 self-recompute. Default to 11 unless you hit a fold-transition bug.)

- [ ] **Step 3: Verify build**

```bash
SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
```
Expected: clean.

- [ ] **Step 4: Verify reset launches happen at fold boundary**

```bash
grep -n "reset_at_fold\|fold_reset_launch" crates/ml/src/trainers/dqn/trainer/training_loop.rs | head -5
```

Confirm the registry's reset launcher is called at fold transitions.

- [ ] **Step 5: Commit B.2**

```bash
git add crates/ml/src/trainers/dqn/state_reset_registry.rs
git commit -m "feat(sp14): B.2 — register 11 SP14 ISV slots for fold-boundary reset

Each EGF pearl EMA / state slot resets to its Pearl-A sentinel at
fold boundary, mirroring sp13_aux_dir_acc_short_ema / long_ema
entries (state_reset_registry.rs:913-922).

Reset slots:
  - Q_DISAGREEMENT_SHORT_EMA, Q_DISAGREEMENT_LONG_EMA → 0.5
  - K_AUX_ADAPTIVE → K_BASE_AUX (20.0)
  - K_Q_ADAPTIVE → K_BASE_Q (15.0)
  - BETA_RATE_LIMITER_ADAPTIVE → BETA_BASE (0.5)
  - AUX_DIR_ACC_VARIANCE_EMA, Q_DISAGREEMENT_VARIANCE_EMA,
    ALPHA_GRAD_RAW_VARIANCE_EMA → 0.0 (initial k = k_base; β = β_base)
  - GATE1_OPEN_STATE → 0.0 (closed)
  - ALPHA_GRAD_SMOOTHED → 0.0
  - AUX_DIR_ACC_POST_OPEN_MIN → 1.0 (no min observed)

ALPHA_GRAD_RAW (recomputed every step) and GRADIENT_HACK_LOCKOUT_REMAINING
(decays at epoch boundary) are NOT in the fold-reset registry; both
naturally re-initialise without explicit reset.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>"
```

## Task B.3: q_disagreement_update_kernel + cubin registration + GPU oracle test

**Files:**
- Create: `crates/ml/src/cuda_pipeline/q_disagreement_update_kernel.cu`
- Modify: `crates/ml/build.rs` (register cubin)
- Modify: `crates/ml/tests/sp14_oracle_tests.rs` (append test)

- [ ] **Step 1: Write the failing test (oracle test for K=4↔K=2 mapping)**

Append to `crates/ml/tests/sp14_oracle_tests.rs`:

```rust
const SP14_Q_DISAGREEMENT_CUBIN: &[u8] =
    include_bytes!(concat!(env!("OUT_DIR"), "/q_disagreement_update_kernel.cubin"));

fn load_q_disagreement(stream: &Arc<CudaStream>) -> CudaFunction {
    let module = stream.context()
        .load_cubin(SP14_Q_DISAGREEMENT_CUBIN.to_vec())
        .expect("load q_disagreement cubin");
    module.load_function("q_disagreement_update_kernel")
        .expect("load q_disagreement_update_kernel")
}

/// Test B.3a: K=4 → K=2 mapping. Q picks Short/Long mapped to aux down/up;
/// Hold/Flat masked from contribution.
#[test]
#[ignore = "requires GPU"]
fn q_disagreement_k4_k2_mapping() {
    let stream = make_test_stream();
    let kernel = load_q_disagreement(&stream);

    const B: usize = 8;
    const K_AUX: usize = 2;
    const K_DIR: usize = 4;

    // Aux softmax: rows 0,1,2,3 say "down" (softmax = [0.9, 0.1])
    //              rows 4,5,6,7 say "up"   (softmax = [0.1, 0.9])
    let mut aux_softmax = vec![0.0_f32; B * K_AUX];
    for b in 0..4 {
        aux_softmax[b * K_AUX + 0] = 0.9; aux_softmax[b * K_AUX + 1] = 0.1;
    }
    for b in 4..8 {
        aux_softmax[b * K_AUX + 0] = 0.1; aux_softmax[b * K_AUX + 1] = 0.9;
    }

    // Q logits — pick winner per row:
    //   row 0: Short (agreement with aux=down)
    //   row 1: Long  (DISAGREEMENT — aux says down, Q says Long)
    //   row 2: Hold  (masked)
    //   row 3: Flat  (masked)
    //   row 4: Short (DISAGREEMENT — aux says up, Q says Short)
    //   row 5: Long  (agreement)
    //   row 6: Hold  (masked)
    //   row 7: Flat  (masked)
    let mut q_logits = vec![0.0_f32; B * K_DIR];
    let winners = [0_usize, 2, 1, 3, 0, 2, 1, 3];  // DIR_SHORT=0, HOLD=1, LONG=2, FLAT=3
    for b in 0..B {
        for k in 0..K_DIR {
            q_logits[b * K_DIR + k] = if k == winners[b] { 1.0 } else { 0.0 };
        }
    }

    // ISV slots — read short_ema, long_ema, variance, write back
    let mut isv = vec![0.0_f32; 1024];
    isv[/* Q_DISAGREEMENT_SHORT_EMA_INDEX */ 381] = 0.5;
    isv[/* Q_DISAGREEMENT_LONG_EMA_INDEX */ 382] = 0.5;
    isv[/* Q_DISAGREEMENT_VARIANCE_EMA_INDEX */ 387] = 0.0;

    let aux_buf = unsafe { MappedF32Buffer::new(B * K_AUX) }.unwrap();
    let q_buf = unsafe { MappedF32Buffer::new(B * K_DIR) }.unwrap();
    let isv_buf = unsafe { MappedF32Buffer::new(1024) }.unwrap();
    aux_buf.write_from_slice(&aux_softmax);
    q_buf.write_from_slice(&q_logits);
    isv_buf.write_from_slice(&isv);

    let bsi = B as i32;
    let alpha_short: f32 = 0.3;
    let alpha_long: f32 = 0.05;
    let alpha_var: f32 = 0.05;

    unsafe {
        stream.launch_builder(&kernel)
            .arg(&aux_buf.dev_ptr)
            .arg(&q_buf.dev_ptr)
            .arg(&isv_buf.dev_ptr)
            .arg(&bsi)
            .arg(&alpha_short)
            .arg(&alpha_long)
            .arg(&alpha_var)
            .launch(LaunchConfig {
                grid_dim: (1, 1, 1),
                block_dim: (256, 1, 1),
                shared_mem_bytes: 0,
            })
            .expect("q_disagreement launch");
    }
    unsafe { cudarc::driver::sys::cuStreamSynchronize(stream.cu_stream()); }

    let result = isv_buf.read_all();
    // 4 contributing rows (0, 1, 4, 5). 2 are agreements (0, 5), 2 are disagreements (1, 4).
    // Batch mean disagreement = 2/4 = 0.5.
    // Pearl-A first observation: short_ema = 0.5 (replaces sentinel directly).
    let short_ema = result[381];
    let long_ema = result[382];
    assert!((short_ema - 0.5).abs() < 1e-4,
        "short_ema after first observation should be 0.5 (Pearl-A); got {}", short_ema);
    assert!((long_ema - 0.5).abs() < 1e-4,
        "long_ema after first observation should be 0.5 (Pearl-A); got {}", long_ema);
}

/// Test B.3b: All-Hold batch. No contributions; sentinel preserved or 0.5.
#[test]
#[ignore = "requires GPU"]
fn q_disagreement_all_hold_no_contribution() {
    let stream = make_test_stream();
    let kernel = load_q_disagreement(&stream);

    const B: usize = 4;
    const K_AUX: usize = 2;
    const K_DIR: usize = 4;

    let aux_softmax = vec![0.5_f32; B * K_AUX];  // uniform aux

    let mut q_logits = vec![0.0_f32; B * K_DIR];
    for b in 0..B {
        q_logits[b * K_DIR + 1] = 1.0;  // Hold winner everywhere
    }

    let mut isv = vec![0.0_f32; 1024];
    isv[381] = 0.5;
    isv[382] = 0.5;
    isv[387] = 0.0;

    let aux_buf = unsafe { MappedF32Buffer::new(B * K_AUX) }.unwrap();
    let q_buf = unsafe { MappedF32Buffer::new(B * K_DIR) }.unwrap();
    let isv_buf = unsafe { MappedF32Buffer::new(1024) }.unwrap();
    aux_buf.write_from_slice(&aux_softmax);
    q_buf.write_from_slice(&q_logits);
    isv_buf.write_from_slice(&isv);

    let bsi = B as i32;
    let alpha: f32 = 0.3;

    unsafe {
        stream.launch_builder(&kernel)
            .arg(&aux_buf.dev_ptr)
            .arg(&q_buf.dev_ptr)
            .arg(&isv_buf.dev_ptr)
            .arg(&bsi)
            .arg(&alpha)
            .arg(&0.05_f32)
            .arg(&0.05_f32)
            .launch(LaunchConfig {
                grid_dim: (1, 1, 1),
                block_dim: (256, 1, 1),
                shared_mem_bytes: 0,
            })
            .expect("q_disagreement launch");
    }
    unsafe { cudarc::driver::sys::cuStreamSynchronize(stream.cu_stream()); }

    let result = isv_buf.read_all();
    // All masked → batch_mean = 0 by max(1, 0) convention. EMA: short = 0.7×0.5 + 0.3×0 = 0.35
    // OR sentinel-preserved: 0.5. Implementer chooses; test confirms either is valid.
    let short_ema = result[381];
    assert!(short_ema >= 0.0 && short_ema <= 0.5,
        "all-Hold: short_ema should be 0.0..0.5; got {}", short_ema);
}
```

- [ ] **Step 2: Run tests to verify failure**

```bash
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp14_oracle_tests \
    q_disagreement --features cuda -- --ignored --nocapture 2>&1 | tail -10
```

Expected: FAIL — cubin doesn't exist (build.rs hasn't registered it; kernel doesn't exist).

- [ ] **Step 3: Write q_disagreement_update_kernel.cu**

Create `crates/ml/src/cuda_pipeline/q_disagreement_update_kernel.cu`:

```cuda
/* ══════════════════════════════════════════════════════════════════════════
 * SP14 Earned Gradient Flow — q_disagreement_update_kernel
 *
 * Per-step compute K=4↔K=2 mapped argmax mismatch + EMA update for short
 * and long timescales + Welford variance update.
 *
 * Mapping (state_layout.cuh):
 *   DIR_SHORT=0 → aux down (0)
 *   DIR_HOLD=1  → masked (no new directional commitment)
 *   DIR_LONG=2  → aux up (1)
 *   DIR_FLAT=3  → masked
 *
 * Pearl-A first-observation: when ISV[short_ema] is sentinel 0.5 AND batch
 * has at least one valid (non-masked) contribution, the EMA is replaced
 * directly with the first observation per `pearl_first_observation_bootstrap`.
 *
 * No atomicAdd: block tree-reduce on numerator + denominator separately,
 * then divide and update EMA in a single thread.
 * ══════════════════════════════════════════════════════════════════════════ */

#include "state_layout.cuh"  // for DIR_SHORT, DIR_HOLD, DIR_LONG, DIR_FLAT

extern "C" __global__
void q_disagreement_update_kernel(
    const float* __restrict__ aux_softmax,   // [B, K_AUX=2]
    const float* __restrict__ q_logits,      // [B, K_DIR=4]
    float* __restrict__ isv,                 // global ISV buffer
    int B,
    float alpha_short,                       // 0.3
    float alpha_long,                        // 0.05
    float alpha_var)                         // 0.05 for Welford EMA on variance
{
    extern __shared__ float smem[];
    float* num_smem = smem;                                  // numerator (sum of disagreements)
    float* cnt_smem = smem + blockDim.x;                     // contribution count

    const int tid = threadIdx.x;
    const int K_AUX = 2;
    const int K_DIR = 4;

    // Per-thread strided accumulation
    float local_num = 0.0f;
    float local_cnt = 0.0f;

    for (int b = tid; b < B; b += blockDim.x) {
        // argmax(aux_softmax[b])
        const float as0 = aux_softmax[b * K_AUX + 0];
        const float as1 = aux_softmax[b * K_AUX + 1];
        const int aux_class = (as1 > as0) ? 1 : 0;

        // argmax(q_logits[b])
        int q_class = 0;
        float q_max = q_logits[b * K_DIR + 0];
        for (int k = 1; k < K_DIR; ++k) {
            const float v = q_logits[b * K_DIR + k];
            if (v > q_max) { q_max = v; q_class = k; }
        }

        // Mask Hold and Flat
        if (q_class == DIR_HOLD || q_class == DIR_FLAT) {
            continue;
        }

        // Map Q to aux space: Short(0)→down(0), Long(2)→up(1)
        const int q_mapped = (q_class == DIR_SHORT) ? 0 : 1;
        local_num += (q_mapped != aux_class) ? 1.0f : 0.0f;
        local_cnt += 1.0f;
    }

    num_smem[tid] = local_num;
    cnt_smem[tid] = local_cnt;
    __syncthreads();

    // Block tree-reduce on numerator + count separately
    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) {
            num_smem[tid] += num_smem[tid + s];
            cnt_smem[tid] += cnt_smem[tid + s];
        }
        __syncthreads();
    }

    if (tid == 0) {
        const float total_num = num_smem[0];
        const float total_cnt = cnt_smem[0];
        const float batch_mean = total_num / fmaxf(total_cnt, 1.0f);

        // Pearl-A bootstrap: sentinel 0.5 → replace if ISV holds sentinel
        const int SHORT_IDX = 381;  // Q_DISAGREEMENT_SHORT_EMA_INDEX (must match sp14_isv_slots.rs)
        const int LONG_IDX  = 382;  // Q_DISAGREEMENT_LONG_EMA_INDEX
        const int VAR_IDX   = 387;  // Q_DISAGREEMENT_VARIANCE_EMA_INDEX

        const float prev_short = isv[SHORT_IDX];
        const float prev_long  = isv[LONG_IDX];
        const float prev_var   = isv[VAR_IDX];

        // Determine if this is the first observation (sentinel = exactly 0.5)
        const bool is_first = (prev_short == 0.5f) && (prev_long == 0.5f) && (total_cnt > 0.0f);

        const float new_short = is_first ? batch_mean : (alpha_short * batch_mean + (1.0f - alpha_short) * prev_short);
        const float new_long  = is_first ? batch_mean : (alpha_long  * batch_mean + (1.0f - alpha_long)  * prev_long);

        // Welford variance EMA: var_new = α × (x - mean)² + (1-α) × var_old
        const float diff = batch_mean - new_short;
        const float new_var = alpha_var * (diff * diff) + (1.0f - alpha_var) * prev_var;

        isv[SHORT_IDX] = new_short;
        isv[LONG_IDX]  = new_long;
        isv[VAR_IDX]   = new_var;
    }
}
```

**Note:** the slot indices (381, 382, 387) are hardcoded to match `sp14_isv_slots.rs`. This couples the kernel to the slot allocation — but that's the standard pattern for this codebase (see SP4/SP5/SP13 kernels). If the indices change, the kernel must be rebuilt.

- [ ] **Step 4: Register cubin in build.rs**

Edit `crates/ml/build.rs`:

```bash
grep -n "aux_pred_to_isv_tanh_kernel\|aux_sign_label_kernel" crates/ml/build.rs | head -3
```

Locate the kernels-with-common-headers list. Add:

```rust
// SP14 Earned Gradient Flow producer kernels.
"q_disagreement_update_kernel.cu",
```

- [ ] **Step 5: Build to compile cubin**

```bash
SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -10
```

Expected: clean. The cubin compiles as part of the build.

- [ ] **Step 6: Run tests to verify pass**

```bash
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp14_oracle_tests \
    q_disagreement --features cuda -- --ignored --nocapture 2>&1 | tail -15
```

Expected: 2 PASS (mapping correctness + all-Hold edge case).

- [ ] **Step 7: Commit B.3**

```bash
git add crates/ml/src/cuda_pipeline/q_disagreement_update_kernel.cu \
        crates/ml/build.rs \
        crates/ml/tests/sp14_oracle_tests.rs
git commit -m "feat(sp14): B.3 — q_disagreement_update_kernel with K=4↔K=2 mapping

Per-step compute argmax mismatch between aux head's K=2 prediction
(down/up) and Q-head's K=4 direction action (Short/Hold/Long/Flat),
masking Hold and Flat (no new directional commitment to evaluate).

Mapping per state_layout.cuh:
  DIR_SHORT=0 → aux down (0)
  DIR_HOLD=1  → masked
  DIR_LONG=2  → aux up (1)
  DIR_FLAT=3  → masked

Block tree-reduce on numerator + denominator separately, then divide
in single thread. No atomicAdd per feedback_no_atomicadd.

Pearl-A first-observation: when ISV slot holds sentinel 0.5 AND batch
has at least one valid contribution, EMA is replaced directly with the
first observation per pearl_first_observation_bootstrap.

Welford variance EMA computed in same kernel (drives k_q adaptive
sigmoid steepness; consumed by alpha_grad_compute_kernel in B.4).

Tests:
- q_disagreement_k4_k2_mapping: 8-row batch with 2 agreements, 2
  disagreements, 4 masked Hold/Flat → first-obs EMA = 0.5
- q_disagreement_all_hold_no_contribution: edge case, all rows
  masked → contribution_count=0; sentinel-bootstrap or zero EMA

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>"
```

## Task B.4: alpha_grad_compute_kernel + cubin + GPU oracle tests

**Files:**
- Create: `crates/ml/src/cuda_pipeline/alpha_grad_compute_kernel.cu`
- Modify: `crates/ml/build.rs` (register cubin)
- Modify: `crates/ml/tests/sp14_oracle_tests.rs` (append tests)

**Why:** This is the heart of the EGF pearl. It reads aux_dir_acc, q_disagreement, variances; computes adaptive `k_aux`, `k_q`, `β`; runs Schmitt-trigger logic for Gate 1; evaluates both sigmoids; multiplies with warmup gate; writes both `α_grad_raw` and `α_grad_smoothed` to ISV.

- [ ] **Step 1: Write the failing test (Schmitt hysteresis)**

Append to `crates/ml/tests/sp14_oracle_tests.rs`:

```rust
const SP14_ALPHA_GRAD_CUBIN: &[u8] =
    include_bytes!(concat!(env!("OUT_DIR"), "/alpha_grad_compute_kernel.cubin"));

fn load_alpha_grad(stream: &Arc<CudaStream>) -> CudaFunction {
    let module = stream.context()
        .load_cubin(SP14_ALPHA_GRAD_CUBIN.to_vec())
        .expect("load alpha_grad cubin");
    module.load_function("alpha_grad_compute_kernel")
        .expect("load alpha_grad_compute_kernel")
}

/// Test B.4a: Schmitt-trigger hysteresis. Below open-threshold, gate stays
/// closed. Above open-threshold from below, gate opens. Stays open until
/// drops below close-threshold.
#[test]
#[ignore = "requires GPU"]
fn alpha_grad_schmitt_hysteresis() {
    let stream = make_test_stream();
    let kernel = load_alpha_grad(&stream);

    let mut isv = vec![0.0_f32; 1024];
    // Pre-populate driver signals
    isv[/* TARGET_DIR_ACC_INDEX */ 372] = 0.55;
    isv[/* AUX_DIR_ACC_SHORT_EMA_INDEX */ 373] = 0.55;  // exactly at target
    // SP14 slots
    isv[381] = 0.55;  // q_disagreement_short — above 0.5 baseline
    isv[383] = 20.0;  // K_AUX_ADAPTIVE base
    isv[384] = 15.0;  // K_Q_ADAPTIVE base
    isv[385] = 0.5;   // BETA_RATE_LIMITER_ADAPTIVE base
    isv[386] = 0.0;   // AUX_DIR_ACC_VARIANCE_EMA = 0 → k_aux = K_BASE_AUX
    isv[387] = 0.0;   // Q_DISAGREEMENT_VARIANCE_EMA = 0 → k_q = K_BASE_Q
    isv[388] = 0.0;   // ALPHA_GRAD_RAW_VARIANCE_EMA = 0 → β = β_base
    isv[389] = 0.0;   // GATE1_OPEN_STATE = 0 (closed)
    isv[391] = 0.0;   // ALPHA_GRAD_SMOOTHED previous

    let isv_buf = unsafe { MappedF32Buffer::new(1024) }.unwrap();
    isv_buf.write_from_slice(&isv);

    // First call: aux_dir_acc = 0.55, gate1 closed; threshold_open = 0.58.
    // 0.55 < 0.58 → gate stays closed → α_grad ≈ 0.
    let warmup_gate: f32 = 1.0;
    let alpha_var: f32 = 0.05;

    unsafe {
        stream.launch_builder(&kernel)
            .arg(&isv_buf.dev_ptr)
            .arg(&warmup_gate)
            .arg(&alpha_var)
            .launch(LaunchConfig { grid_dim: (1,1,1), block_dim: (32,1,1), shared_mem_bytes: 0 })
            .expect("alpha_grad launch");
    }
    unsafe { cudarc::driver::sys::cuStreamSynchronize(stream.cu_stream()); }

    let result = isv_buf.read_all();
    let alpha_raw_below = result[/* ALPHA_GRAD_RAW */ 390];
    let gate1_state_below = result[389];
    assert!(alpha_raw_below < 0.5,
        "Below threshold_open=0.58 (aux_dir_acc=0.55), gate1 should remain closed; got α={}", alpha_raw_below);
    assert_eq!(gate1_state_below, 0.0, "gate1_open_state should still be closed");

    // Second call: aux_dir_acc rises to 0.60, above threshold_open. Gate opens.
    isv[373] = 0.60;
    isv_buf.write_from_slice(&isv);
    unsafe {
        stream.launch_builder(&kernel)
            .arg(&isv_buf.dev_ptr)
            .arg(&warmup_gate)
            .arg(&alpha_var)
            .launch(LaunchConfig { grid_dim: (1,1,1), block_dim: (32,1,1), shared_mem_bytes: 0 })
            .expect("alpha_grad launch");
    }
    unsafe { cudarc::driver::sys::cuStreamSynchronize(stream.cu_stream()); }

    let result = isv_buf.read_all();
    let alpha_raw_above = result[390];
    let gate1_state_above = result[389];
    assert!(alpha_raw_above > alpha_raw_below,
        "Above threshold_open, gate1 should open; α should rise");
    assert_eq!(gate1_state_above, 1.0, "gate1_open_state should be open=1.0");

    // Third call: aux_dir_acc falls to 0.54, between close and open thresholds.
    // Gate stays open (Schmitt: was open, drops to between 0.52 and 0.58 — stays).
    isv[373] = 0.54;
    isv_buf.write_from_slice(&isv);
    unsafe {
        stream.launch_builder(&kernel)
            .arg(&isv_buf.dev_ptr)
            .arg(&warmup_gate)
            .arg(&alpha_var)
            .launch(LaunchConfig { grid_dim: (1,1,1), block_dim: (32,1,1), shared_mem_bytes: 0 })
            .expect("alpha_grad launch");
    }
    unsafe { cudarc::driver::sys::cuStreamSynchronize(stream.cu_stream()); }

    let result = isv_buf.read_all();
    let gate1_state_between = result[389];
    assert_eq!(gate1_state_between, 1.0,
        "Schmitt hysteresis: gate stays open between close and open thresholds");

    // Fourth call: aux_dir_acc falls to 0.50, below close-threshold (0.52).
    // Gate finally closes.
    isv[373] = 0.50;
    isv_buf.write_from_slice(&isv);
    unsafe {
        stream.launch_builder(&kernel)
            .arg(&isv_buf.dev_ptr)
            .arg(&warmup_gate)
            .arg(&alpha_var)
            .launch(LaunchConfig { grid_dim: (1,1,1), block_dim: (32,1,1), shared_mem_bytes: 0 })
            .expect("alpha_grad launch");
    }
    unsafe { cudarc::driver::sys::cuStreamSynchronize(stream.cu_stream()); }

    let result = isv_buf.read_all();
    let gate1_state_closed = result[389];
    assert_eq!(gate1_state_closed, 0.0, "Below close-threshold, gate finally closes");
}

/// Test B.4b: Adaptive β rate limiter. β increases as α_grad_raw variance grows.
#[test]
#[ignore = "requires GPU"]
fn alpha_grad_adaptive_beta() {
    // Test math directly via repeated kernel launches with synthetic
    // α_grad_raw oscillation; verify β_adaptive grows toward β_max.
    // Detailed implementation per adaptive formula B.2.8.
    let stream = make_test_stream();
    let kernel = load_alpha_grad(&stream);

    let mut isv = vec![0.0_f32; 1024];
    // Drive a regime where gates oscillate to produce α_grad_raw variance.
    // Specifically: aux_dir_acc oscillates above and below threshold.
    isv[372] = 0.55;
    isv[373] = 0.55;
    isv[381] = 0.6;
    isv[383] = 20.0;
    isv[384] = 15.0;
    isv[385] = 0.5;
    isv[386] = 0.0;
    isv[387] = 0.0;
    isv[388] = 0.0;
    isv[389] = 0.0;
    isv[390] = 0.0;
    isv[391] = 0.0;

    let isv_buf = unsafe { MappedF32Buffer::new(1024) }.unwrap();
    isv_buf.write_from_slice(&isv);

    // Run 20 steps with aux_dir_acc oscillating. β should grow.
    let oscillation = [0.60_f32, 0.50, 0.60, 0.50, 0.60, 0.50, 0.60, 0.50,
                       0.60, 0.50, 0.60, 0.50, 0.60, 0.50, 0.60, 0.50,
                       0.60, 0.50, 0.60, 0.50];

    for &acc in oscillation.iter() {
        isv[373] = acc;
        isv_buf.write_from_slice(&isv);
        unsafe {
            stream.launch_builder(&kernel)
                .arg(&isv_buf.dev_ptr)
                .arg(&1.0_f32)
                .arg(&0.05_f32)
                .launch(LaunchConfig { grid_dim: (1,1,1), block_dim: (32,1,1), shared_mem_bytes: 0 })
                .expect("alpha_grad launch");
        }
        unsafe { cudarc::driver::sys::cuStreamSynchronize(stream.cu_stream()); }
        // Read updated isv back into local vec for next iteration
        isv = isv_buf.read_all();
    }

    // After 20 oscillations, variance EMA should be elevated; β should be > β_base.
    let beta_after = isv[/* BETA_RATE_LIMITER_ADAPTIVE_INDEX */ 385];
    assert!(beta_after > 0.5,
        "After oscillation, adaptive β should grow above β_base=0.5; got {}", beta_after);
    assert!(beta_after <= 0.95,
        "β should be bounded by β_max=0.95; got {}", beta_after);
}
```

- [ ] **Step 2: Run tests to verify failure**

```bash
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp14_oracle_tests \
    alpha_grad --features cuda -- --ignored --nocapture 2>&1 | tail -10
```

Expected: FAIL — cubin missing.

- [ ] **Step 3: Implement alpha_grad_compute_kernel.cu**

Create the kernel implementing all gate logic. ~100 LOC. Key elements:

```cuda
/* SP14 Earned Gradient Flow — alpha_grad_compute_kernel
 *
 * Reads driver signals from ISV, runs Schmitt-trigger Gate 1, computes
 * adaptive k_aux/k_q/β from variance EMAs, evaluates both sigmoids,
 * multiplies with warmup gate, writes ALPHA_GRAD_RAW + ALPHA_GRAD_SMOOTHED.
 *
 * Single-thread kernel (state machine; no parallelism needed).
 */

extern "C" __global__
void alpha_grad_compute_kernel(
    float* __restrict__ isv,
    float warmup_gate,        // [0, 1]
    float alpha_var)          // EMA decay for variance updates
{
    if (threadIdx.x != 0 || blockIdx.x != 0) return;

    // Slot indices (must match sp14_isv_slots.rs and sp13_isv_slots.rs)
    const int TARGET_DIR_ACC = 372;
    const int AUX_DIR_ACC_SHORT_EMA = 373;
    const int Q_DISAGREEMENT_SHORT_EMA = 381;
    const int K_AUX_ADAPTIVE = 383;
    const int K_Q_ADAPTIVE = 384;
    const int BETA_ADAPTIVE = 385;
    const int VAR_AUX = 386;
    const int VAR_Q = 387;
    const int VAR_ALPHA = 388;
    const int GATE1_OPEN = 389;
    const int ALPHA_RAW = 390;
    const int ALPHA_SMOOTHED = 391;

    // Constants (must match sp14_isv_slots.rs)
    const float K_BASE_AUX = 20.0f;
    const float K_BASE_Q = 15.0f;
    const float K_MIN = 1.0f;
    const float VARIANCE_REF_AUX = 0.01f;
    const float VARIANCE_REF_Q = 0.05f;
    const float VARIANCE_REF_ALPHA = 0.01f;
    const float BETA_BASE = 0.5f;
    const float BETA_MAX = 0.95f;
    const float SCHMITT_BAND = 0.03f;
    const float Q_DISAGREEMENT_BASELINE = 0.5f;

    // ── Read drivers ─────────────────────────────────────────────────
    const float target = isv[TARGET_DIR_ACC];
    const float aux_short = isv[AUX_DIR_ACC_SHORT_EMA];
    const float q_dis = isv[Q_DISAGREEMENT_SHORT_EMA];
    const float var_aux = isv[VAR_AUX];
    const float var_q = isv[VAR_Q];
    const float var_alpha_prev = isv[VAR_ALPHA];
    const float gate1_state_prev = isv[GATE1_OPEN];
    const float alpha_smoothed_prev = isv[ALPHA_SMOOTHED];

    // ── Adaptive sigmoid steepness (B.2.5) ───────────────────────────
    const float k_aux = fmaxf(K_BASE_AUX / (1.0f + var_aux / VARIANCE_REF_AUX), K_MIN);
    const float k_q = fmaxf(K_BASE_Q / (1.0f + var_q / VARIANCE_REF_Q), K_MIN);

    // ── Schmitt-trigger Gate 1 state update (B.2.4) ──────────────────
    const float threshold_open = target + SCHMITT_BAND;   // e.g. 0.58
    const float threshold_close = target - SCHMITT_BAND;  // e.g. 0.52
    float gate1_state_new = gate1_state_prev;
    if (gate1_state_prev < 0.5f && aux_short > threshold_open) {
        gate1_state_new = 1.0f;
    } else if (gate1_state_prev >= 0.5f && aux_short < threshold_close) {
        gate1_state_new = 0.0f;
    }

    // ── Sigmoid arg for Gate 1 (uses appropriate threshold per state) ─
    const float gate1_arg_raw = (gate1_state_new >= 0.5f)
        ? (aux_short - threshold_close)  // open: distance from close-threshold
        : (aux_short - threshold_open);  // closed: distance from open-threshold
    const float gate1_arg = fmaxf(-30.0f, fminf(30.0f, k_aux * gate1_arg_raw));
    const float gate1 = 1.0f / (1.0f + __expf(-gate1_arg));

    // ── Sigmoid for Gate 2 (B.2.3) ───────────────────────────────────
    const float gate2_arg_raw = q_dis - Q_DISAGREEMENT_BASELINE;
    const float gate2_arg = fmaxf(-30.0f, fminf(30.0f, k_q * gate2_arg_raw));
    const float gate2 = 1.0f / (1.0f + __expf(-gate2_arg));

    // ── α_grad_raw (B.2.1) ───────────────────────────────────────────
    const float alpha_raw = gate1 * gate2 * warmup_gate;

    // ── Welford variance EMA on α_grad_raw ────────────────────────────
    const float diff_alpha = alpha_raw - alpha_smoothed_prev;
    const float var_alpha_new = alpha_var * (diff_alpha * diff_alpha)
                              + (1.0f - alpha_var) * var_alpha_prev;

    // ── Adaptive β (B.2.8) ───────────────────────────────────────────
    const float beta = fmaxf(BETA_BASE,
        fminf(BETA_MAX,
            BETA_BASE + var_alpha_new / VARIANCE_REF_ALPHA));

    // ── α_grad_smoothed (rate-limited) ───────────────────────────────
    const float alpha_smoothed = beta * alpha_smoothed_prev + (1.0f - beta) * alpha_raw;

    // ── Write back ───────────────────────────────────────────────────
    isv[K_AUX_ADAPTIVE] = k_aux;
    isv[K_Q_ADAPTIVE] = k_q;
    isv[BETA_ADAPTIVE] = beta;
    isv[VAR_ALPHA] = var_alpha_new;
    isv[GATE1_OPEN] = gate1_state_new;
    isv[ALPHA_RAW] = alpha_raw;
    isv[ALPHA_SMOOTHED] = alpha_smoothed;
}
```

**Note:** Variance EMAs for `var_aux` and `var_q` are updated by the producer kernels (q_disagreement_update_kernel for var_q; aux_dir_acc producer chain for var_aux — implementer needs to check whether SP13 already computes this, otherwise add to the same kernel that writes AUX_DIR_ACC_SHORT_EMA).

- [ ] **Step 4: Register cubin**

Edit `crates/ml/build.rs`:

```rust
"alpha_grad_compute_kernel.cu",
```

- [ ] **Step 5: Build**

```bash
SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
```
Expected: clean.

- [ ] **Step 6: Run tests to verify pass**

```bash
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp14_oracle_tests \
    alpha_grad --features cuda -- --ignored --nocapture 2>&1 | tail -15
```
Expected: 2 PASS.

- [ ] **Step 7: Commit B.4**

```bash
git add crates/ml/src/cuda_pipeline/alpha_grad_compute_kernel.cu \
        crates/ml/build.rs \
        crates/ml/tests/sp14_oracle_tests.rs
git commit -m "feat(sp14): B.4 — alpha_grad_compute_kernel (EGF heart)

Single-thread state-machine kernel that computes α_grad from ISV
driver signals:

1. Read aux_dir_acc, q_disagreement, variances, gate1_state from ISV
2. Compute adaptive k_aux, k_q from Welford variance EMAs (B.2.5)
3. Run Schmitt-trigger Gate 1 state update (open at target+0.03,
   close at target-0.03)
4. Compute Gate 1 sigmoid (aux competence) and Gate 2 sigmoid
   (Q-head disagreement vs structural baseline 0.5)
5. α_grad_raw = gate1 × gate2 × warmup_gate
6. Update Welford variance of α_grad_raw → adaptive β (B.2.8)
7. α_grad_smoothed = β × α_grad_smoothed_prev + (1-β) × α_grad_raw
8. Write back all 7 outputs to ISV

Sigmoid arguments clipped to [-30, 30] per Issue 2 (numerical
stability under fp32). Per pearl_bounded_modifier_outputs_require_
structural_activation: sigmoid composition produces structurally-
bounded [0, 1] output.

Tests:
- alpha_grad_schmitt_hysteresis: 4-step trajectory (closed-open-
  open-closed) verifies state stays open across close-threshold
  hysteresis band
- alpha_grad_adaptive_beta: 20 oscillations verify β grows
  above β_base=0.5 toward β_max=0.95

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>"
```

## Task B.5: gradient_hack_detect_kernel + cubin + GPU oracle test

**Files:**
- Create: `crates/ml/src/cuda_pipeline/gradient_hack_detect_kernel.cu` (~40 LOC)
- Modify: `crates/ml/build.rs`
- Modify: `crates/ml/tests/sp14_oracle_tests.rs`

- [ ] **Step 1: Write the failing test**

```rust
const SP14_GRAD_HACK_CUBIN: &[u8] =
    include_bytes!(concat!(env!("OUT_DIR"), "/gradient_hack_detect_kernel.cubin"));

fn load_grad_hack(stream: &Arc<CudaStream>) -> CudaFunction {
    let module = stream.context()
        .load_cubin(SP14_GRAD_HACK_CUBIN.to_vec())
        .expect("load grad_hack cubin");
    module.load_function("gradient_hack_detect_kernel")
        .expect("load gradient_hack_detect_kernel")
}

#[test]
#[ignore = "requires GPU"]
fn gradient_hack_circuit_breaker_fires() {
    let stream = make_test_stream();
    let kernel = load_grad_hack(&stream);

    let mut isv = vec![0.0_f32; 1024];
    isv[372] = 0.55;  // target
    isv[373] = 0.50;  // aux_dir_acc dropped (was 0.60 when gate opened)
    isv[381] = 0.65;  // q_disagreement rose
    isv[389] = 1.0;   // gate1_open_state = open
    isv[392] = 0.50;  // post_open_min: lowest aux_dir_acc since gate opened
    isv[393] = 0.0;   // lockout_remaining

    let isv_buf = unsafe { MappedF32Buffer::new(1024) }.unwrap();
    isv_buf.write_from_slice(&isv);

    unsafe {
        stream.launch_builder(&kernel)
            .arg(&isv_buf.dev_ptr)
            .launch(LaunchConfig { grid_dim: (1,1,1), block_dim: (32,1,1), shared_mem_bytes: 0 })
            .expect("grad_hack launch");
    }
    unsafe { cudarc::driver::sys::cuStreamSynchronize(stream.cu_stream()); }

    let result = isv_buf.read_all();
    // aux_dir_acc=0.50 is 0.05 below threshold_open=0.58 (drop > LOCKOUT_TRIGGER_DROP=0.05)
    // q_disagreement=0.65 is 0.15 above baseline=0.5 (rise > LOCKOUT_TRIGGER_DIS_RISE=0.10)
    // → both conditions met → trigger lockout
    let lockout = result[393];
    let gate1 = result[389];
    assert_eq!(lockout, 2.0, "Lockout should be set to LOCKOUT_EPOCHS=2; got {}", lockout);
    assert_eq!(gate1, 0.0, "Gate1 should be force-closed; got {}", gate1);
}
```

- [ ] **Step 2: Run failing test**

```bash
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp14_oracle_tests \
    gradient_hack --features cuda -- --ignored --nocapture 2>&1 | tail -10
```
Expected: FAIL.

- [ ] **Step 3: Implement gradient_hack_detect_kernel.cu**

```cuda
/* SP14 EGF — gradient_hack_detect_kernel
 *
 * Anti-mesa-optimization circuit breaker.
 *
 * Trigger: aux_dir_acc drops more than 0.05 below threshold_open AND
 *          q_disagreement rises more than 0.10 above baseline 0.5
 *          (suspected gradient hacking).
 * Action:  force gate1_open_state = 0; set lockout_remaining = 2 epochs.
 *
 * Single-thread kernel; runs at end of each epoch.
 */

extern "C" __global__
void gradient_hack_detect_kernel(float* __restrict__ isv) {
    if (threadIdx.x != 0 || blockIdx.x != 0) return;

    const int TARGET = 372;
    const int AUX_SHORT = 373;
    const int Q_DIS_SHORT = 381;
    const int GATE1 = 389;
    const int POST_OPEN_MIN = 392;
    const int LOCKOUT = 393;

    const float SCHMITT_BAND = 0.03f;
    const float TRIGGER_DROP = 0.05f;
    const float TRIGGER_DIS_RISE = 0.10f;
    const float Q_BASELINE = 0.5f;
    const float LOCKOUT_EPOCHS = 2.0f;

    const float target = isv[TARGET];
    const float aux = isv[AUX_SHORT];
    const float q_dis = isv[Q_DIS_SHORT];
    const float gate1 = isv[GATE1];
    const float prev_min = isv[POST_OPEN_MIN];
    const float lockout = isv[LOCKOUT];

    // Decrement existing lockout
    float lockout_new = fmaxf(0.0f, lockout - 1.0f);

    // If gate1 currently open, track post-open minimum
    float post_open_min_new = prev_min;
    if (gate1 >= 0.5f) {
        post_open_min_new = fminf(prev_min, aux);
    } else {
        // gate closed → reset min to 1.0 (no observation)
        post_open_min_new = 1.0f;
    }

    // Detect hacking: aux dropped & disagreement rose simultaneously
    const float threshold_open = target + SCHMITT_BAND;
    const float aux_drop = threshold_open - post_open_min_new;
    const float q_rise = q_dis - Q_BASELINE;

    if (gate1 >= 0.5f && aux_drop > TRIGGER_DROP && q_rise > TRIGGER_DIS_RISE) {
        // Trigger circuit breaker
        lockout_new = LOCKOUT_EPOCHS;
        isv[GATE1] = 0.0f;  // force-close
        post_open_min_new = 1.0f;  // reset
    } else if (lockout_new > 0.0f) {
        // Still in lockout: keep gate closed
        isv[GATE1] = 0.0f;
    }

    isv[POST_OPEN_MIN] = post_open_min_new;
    isv[LOCKOUT] = lockout_new;
}
```

- [ ] **Step 4: Register cubin + build + run test**

```bash
# Register in build.rs (append to kernel list)
SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp14_oracle_tests \
    gradient_hack --features cuda -- --ignored --nocapture 2>&1 | tail -10
```
Expected: PASS.

- [ ] **Step 5: Commit B.5**

```bash
git add crates/ml/src/cuda_pipeline/gradient_hack_detect_kernel.cu \
        crates/ml/build.rs \
        crates/ml/tests/sp14_oracle_tests.rs
git commit -m "feat(sp14): B.5 — gradient_hack_detect_kernel (anti-mesa-opt circuit breaker)

Detects suspected gradient hacking: when gate1 is open AND aux_dir_acc
post-open-minimum drops > 0.05 below threshold_open AND q_disagreement
rises > 0.10 above structural baseline 0.5 simultaneously.

Action: force gate1_open_state = 0, set lockout_remaining = 2 epochs.
During lockout, gate1 forced closed regardless of other signals.

Tracks ALPHA_DIR_ACC_POST_OPEN_MIN (lowest aux_dir_acc since gate1
last opened); resets to 1.0 when gate closes naturally.

Single-thread state-machine kernel; runs at end of each epoch.

Test: gradient_hack_circuit_breaker_fires verifies trigger conditions
+ lockout assignment + gate force-close.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>"
```

## Task B.6: dir_concat_qaux_kernel + cubin + GPU oracle test

**Files:**
- Create: `crates/ml/src/cuda_pipeline/dir_concat_qaux_kernel.cu`
- Modify: `crates/ml/build.rs`
- Modify: `crates/ml/tests/sp14_oracle_tests.rs`

**Why:** Per `mag_concat_qdir` precedent, the cleanest forward-wire approach is a pre-SGEMM concat kernel that copies `[h_s2 ; aux_softmax_diff]` into a scratch buffer, then the existing SGEMM consumes the new wider input. This isolates the wire from SGEMM internals.

- [ ] **Step 1: Write the failing test**

```rust
const SP14_DIR_CONCAT_CUBIN: &[u8] =
    include_bytes!(concat!(env!("OUT_DIR"), "/dir_concat_qaux_kernel.cubin"));

fn load_dir_concat(stream: &Arc<CudaStream>) -> CudaFunction {
    let module = stream.context()
        .load_cubin(SP14_DIR_CONCAT_CUBIN.to_vec())
        .expect("load dir_concat cubin");
    module.load_function("dir_concat_qaux_kernel")
        .expect("load dir_concat_qaux_kernel")
}

#[test]
#[ignore = "requires GPU"]
fn dir_concat_qaux_correct() {
    let stream = make_test_stream();
    let kernel = load_dir_concat(&stream);

    const B: usize = 4;
    const SH2: usize = 8;
    const K_AUX: usize = 2;

    // h_s2 input
    let h_s2: Vec<f32> = (0..B*SH2).map(|i| i as f32).collect();
    // aux_softmax: row 0,1 → "down" (0.9, 0.1) → diff = -0.8
    //              row 2,3 → "up"   (0.1, 0.9) → diff = +0.8
    let mut aux_softmax = vec![0.0_f32; B * K_AUX];
    for b in 0..2 { aux_softmax[b * K_AUX + 0] = 0.9; aux_softmax[b * K_AUX + 1] = 0.1; }
    for b in 2..4 { aux_softmax[b * K_AUX + 0] = 0.1; aux_softmax[b * K_AUX + 1] = 0.9; }

    let h_buf = unsafe { MappedF32Buffer::new(B * SH2) }.unwrap();
    let aux_buf = unsafe { MappedF32Buffer::new(B * K_AUX) }.unwrap();
    let out_buf = unsafe { MappedF32Buffer::new(B * (SH2 + 1)) }.unwrap();
    h_buf.write_from_slice(&h_s2);
    aux_buf.write_from_slice(&aux_softmax);

    let bsi = B as i32;
    let sdi = SH2 as i32;

    unsafe {
        stream.launch_builder(&kernel)
            .arg(&h_buf.dev_ptr)
            .arg(&aux_buf.dev_ptr)
            .arg(&out_buf.dev_ptr)
            .arg(&bsi)
            .arg(&sdi)
            .launch(LaunchConfig {
                grid_dim: ((B * (SH2 + 1) + 255) / 256, 1, 1).into(),
                block_dim: (256, 1, 1).into(),
                shared_mem_bytes: 0,
            })
            .expect("dir_concat_qaux launch");
    }
    unsafe { cudarc::driver::sys::cuStreamSynchronize(stream.cu_stream()); }

    let result = out_buf.read_all();
    // Verify rows: each row b should be [h_s2[b*SH2..b*SH2+SH2], softmax[b,1] - softmax[b,0]]
    for b in 0..B {
        // h_s2 portion
        for k in 0..SH2 {
            let expected = (b * SH2 + k) as f32;
            let actual = result[b * (SH2 + 1) + k];
            assert!((actual - expected).abs() < 1e-5,
                "row {} pos {}: expected h_s2 value {}, got {}", b, k, expected, actual);
        }
        // softmax diff at position SH2
        let diff_actual = result[b * (SH2 + 1) + SH2];
        let diff_expected = aux_softmax[b * K_AUX + 1] - aux_softmax[b * K_AUX + 0];
        assert!((diff_actual - diff_expected).abs() < 1e-5,
            "row {}: expected diff {}, got {}", b, diff_expected, diff_actual);
    }
}
```

- [ ] **Step 2: Run failing test**

```bash
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp14_oracle_tests \
    dir_concat_qaux --features cuda -- --ignored --nocapture 2>&1 | tail -5
```
Expected: FAIL.

- [ ] **Step 3: Implement dir_concat_qaux_kernel.cu**

```cuda
/* SP14 EGF — dir_concat_qaux_kernel
 *
 * Pre-SGEMM concat: writes [h_s2 ; aux_softmax_diff] into a scratch
 * buffer of shape [B, SH2 + 1]. The direction Q-head's first FC SGEMM
 * then consumes this wider input.
 *
 * aux_softmax_diff = softmax[b, 1] - softmax[b, 0] ∈ [-1, +1] (structurally
 * bounded by softmax). Per `pearl_bounded_modifier_outputs_require_structural_activation`.
 *
 * Per `pearl_canary_input_freshness_launch_order`: aux head forward MUST
 * complete before this concat reads aux_nb_softmax_buf. The graph capture
 * orchestrator enforces serial dep.
 *
 * Pure per-thread map. No reduction. No atomicAdd.
 */

extern "C" __global__
void dir_concat_qaux_kernel(
    const float* __restrict__ h_s2,         // [B, SH2]
    const float* __restrict__ aux_softmax,  // [B, K_AUX=2]
    float* __restrict__ out,                // [B, SH2 + 1]
    int B,
    int SH2)
{
    const int tid = blockIdx.x * blockDim.x + threadIdx.x;
    const int total = B * (SH2 + 1);
    if (tid >= total) return;

    const int row = tid / (SH2 + 1);
    const int col = tid % (SH2 + 1);

    if (col < SH2) {
        out[tid] = h_s2[row * SH2 + col];
    } else {
        // col == SH2: write softmax diff
        const float s0 = aux_softmax[row * 2 + 0];
        const float s1 = aux_softmax[row * 2 + 1];
        out[tid] = s1 - s0;
    }
}
```

- [ ] **Step 4: Register cubin + build + run test**

```bash
# Register dir_concat_qaux_kernel.cu in build.rs
SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp14_oracle_tests \
    dir_concat_qaux --features cuda -- --ignored --nocapture 2>&1 | tail -5
```
Expected: PASS.

- [ ] **Step 5: Commit B.6**

```bash
git add crates/ml/src/cuda_pipeline/dir_concat_qaux_kernel.cu \
        crates/ml/build.rs \
        crates/ml/tests/sp14_oracle_tests.rs
git commit -m "feat(sp14): B.6 — dir_concat_qaux_kernel (pre-SGEMM forward-wire concat)

Mirrors mag_concat_qdir precedent (experience_kernels.cu:4560).
Concats [h_s2 ; aux_softmax_diff] into scratch buffer [B, SH2+1] for
the direction Q-head's first FC SGEMM.

aux_softmax_diff = softmax[b, 1] - softmax[b, 0] ∈ [-1, +1] computed
inline; structurally bounded by softmax components per
pearl_bounded_modifier_outputs_require_structural_activation.

Pure per-thread map; no reduction. Launch order constraint: aux head
forward MUST complete before this concat reads aux_nb_softmax_buf
(enforced by orchestrator in B.10).

Test: dir_concat_qaux_correct verifies row-wise contiguous concat
+ correct softmax diff values for both 'down' (-0.8) and 'up' (+0.8)
synthetic aux predictions.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>"
```

## Task B.7: Trainer struct fields + buffer allocations

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (add fields, kernel handles, scratch buffer)

**Why:** The trainer struct needs to hold:
1. New CudaFunction handles for the 3 SP14 producer kernels + 1 concat kernel
2. A new scratch buffer `dir_qaux_concat_scratch: CudaSlice<f32>` of size `B * (SH2 + 1)` for the pre-SGEMM concat output
3. Backward gradient scaling state (handled in B.9; this task adds the buffer)

- [ ] **Step 1: Locate the trainer struct kernel-handle region**

```bash
grep -n "aux_pred_to_isv_tanh_kernel: CudaFunction\|aux_sign_label_kernel: CudaFunction" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs | head -3
```

These are existing single-CudaFunction fields; the new SP14 kernels follow the same pattern.

- [ ] **Step 2: Add SP14 kernel handles + scratch buffer**

Edit `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`. In the trainer struct definition, add:

```rust
    // ── SP14 Earned Gradient Flow kernels ─────────────────────────
    /// SP14 B.3: per-step Q-head ↔ aux argmax disagreement EMA producer.
    sp14_q_disagreement_update_kernel: CudaFunction,
    /// SP14 B.4: α_grad compute kernel (Schmitt + adaptive k + adaptive β).
    sp14_alpha_grad_compute_kernel: CudaFunction,
    /// SP14 B.5: anti-gradient-hacking circuit breaker.
    sp14_gradient_hack_detect_kernel: CudaFunction,
    /// SP14 B.6: pre-SGEMM concat kernel for direction Q-head input.
    sp14_dir_concat_qaux_kernel: CudaFunction,

    /// SP14 B.7: scratch buffer for direction Q-head input concat.
    /// Shape [B, SH2 + 1]. Written by `dir_concat_qaux_kernel` before the
    /// direction Q-head's first FC SGEMM consumes it. Replaces the prior
    /// direct h_s2 input.
    sp14_dir_qaux_concat_scratch: CudaSlice<f32>,
```

- [ ] **Step 3: Add cubin embeds + load in trainer construction**

```bash
grep -n "AUX_PRED_TO_ISV_TANH_CUBIN\|aux_pred_to_isv_tanh_kernel.cubin" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs | head -3
```

Find existing pattern. Add at the top of the file (alongside other cubin static refs):

```rust
static SP14_Q_DISAGREEMENT_CUBIN: &[u8] =
    include_bytes!(concat!(env!("OUT_DIR"), "/q_disagreement_update_kernel.cubin"));
static SP14_ALPHA_GRAD_CUBIN: &[u8] =
    include_bytes!(concat!(env!("OUT_DIR"), "/alpha_grad_compute_kernel.cubin"));
static SP14_GRAD_HACK_CUBIN: &[u8] =
    include_bytes!(concat!(env!("OUT_DIR"), "/gradient_hack_detect_kernel.cubin"));
static SP14_DIR_CONCAT_CUBIN: &[u8] =
    include_bytes!(concat!(env!("OUT_DIR"), "/dir_concat_qaux_kernel.cubin"));
```

In trainer construction (look for where `aux_pred_to_isv_tanh_kernel` is loaded):

```rust
        // SP14 EGF kernel loads
        let sp14_q_dis_module = stream.context()
            .load_cubin(SP14_Q_DISAGREEMENT_CUBIN.to_vec())
            .map_err(|e| MLError::ModelError(format!("sp14 q_dis cubin: {e}")))?;
        let sp14_q_disagreement_update_kernel = sp14_q_dis_module
            .load_function("q_disagreement_update_kernel")
            .map_err(|e| MLError::ModelError(format!("q_disagreement_update_kernel: {e}")))?;

        let sp14_alpha_module = stream.context()
            .load_cubin(SP14_ALPHA_GRAD_CUBIN.to_vec())
            .map_err(|e| MLError::ModelError(format!("sp14 alpha cubin: {e}")))?;
        let sp14_alpha_grad_compute_kernel = sp14_alpha_module
            .load_function("alpha_grad_compute_kernel")
            .map_err(|e| MLError::ModelError(format!("alpha_grad_compute_kernel: {e}")))?;

        let sp14_grad_hack_module = stream.context()
            .load_cubin(SP14_GRAD_HACK_CUBIN.to_vec())
            .map_err(|e| MLError::ModelError(format!("sp14 grad_hack cubin: {e}")))?;
        let sp14_gradient_hack_detect_kernel = sp14_grad_hack_module
            .load_function("gradient_hack_detect_kernel")
            .map_err(|e| MLError::ModelError(format!("gradient_hack_detect_kernel: {e}")))?;

        let sp14_dir_concat_module = stream.context()
            .load_cubin(SP14_DIR_CONCAT_CUBIN.to_vec())
            .map_err(|e| MLError::ModelError(format!("sp14 dir_concat cubin: {e}")))?;
        let sp14_dir_concat_qaux_kernel = sp14_dir_concat_module
            .load_function("dir_concat_qaux_kernel")
            .map_err(|e| MLError::ModelError(format!("dir_concat_qaux_kernel: {e}")))?;

        // SP14 scratch buffer
        let sp14_dir_qaux_concat_scratch = stream
            .alloc_zeros::<f32>(batch_size * (cfg.shared_h2 + 1))
            .map_err(|e| MLError::ModelError(format!("sp14 dir_qaux scratch: {e}")))?;
```

Add to the struct's `Self { ... }` initializer.

- [ ] **Step 4: Build check**

```bash
SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -10
```
Expected: clean — kernels and scratch are now loaded but not USED yet (no orphan errors because they're held in struct fields).

- [ ] **Step 5: Commit B.7**

```bash
git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
git commit -m "feat(sp14): B.7 — trainer struct fields for EGF kernels + concat scratch

Adds 4 CudaFunction handles + 1 scratch buffer to GpuDqnTrainer:
- sp14_q_disagreement_update_kernel
- sp14_alpha_grad_compute_kernel
- sp14_gradient_hack_detect_kernel
- sp14_dir_concat_qaux_kernel
- sp14_dir_qaux_concat_scratch: CudaSlice<f32> [B * (SH2 + 1)]

All loaded from precompiled cubins in trainer construction, mirroring
the SP13 aux_pred_to_isv_tanh / aux_sign_label kernel-loading pattern.
Per feedback_no_partial_refactor: handles + scratch are held but not
yet wired in. Subsequent tasks (B.9+) launch them.

Build: cargo check --workspace clean.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>"
```

## Task B.8: Forward wire — Q-head input dim bump + fingerprint

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (compute_param_sizes, layout_fingerprint_seed, weight init)

**Why:** Direction Q-head's first FC weight tensor (`w_b0fc` per Phase 0 — implementer to identify exact index) input dim grows from `SH2` to `SH2 + 1`. Layout fingerprint MUST bump (Invariant 8) so old checkpoints don't load silently into the new architecture. New weight column zero-initialized.

- [ ] **Step 1: Identify w_b0fc tensor index in compute_param_sizes**

```bash
grep -n "w_b0fc\|w_b1fc\|sizes\[" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs | head -20
```

Locate the line that sets the size of the direction Q-head's first FC weight. Currently `cfg.adv_h * cfg.shared_h2` (or similar). Note exact tensor index (e.g., `sizes[17]`).

- [ ] **Step 2: Identify layout_fingerprint_seed entry**

```bash
grep -n "PARAM_W_DIR\|PARAM_W_B0\|w_b0fc" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs | head -5
```

Find the `PARAM_W_*` entry corresponding to `w_b0fc`. Note the existing name format (e.g., `PARAM_W_DIR_FC1=17` or similar).

- [ ] **Step 3: Modify compute_param_sizes**

Update the line:

```rust
        // SP14 forward wire: direction Q-head's first FC input dim grows
        // SH2 → SH2 + 1 to accept aux_softmax_diff. Output dim (adv_h)
        // unchanged. New weight column zero-initialised in trainer
        // construction so model starts ignoring the new input.
        sizes[/* w_b0fc index */ 17] = cfg.adv_h * (cfg.shared_h2 + 1);  // SP14: was adv_h * shared_h2
```

Update the corresponding bias if the bias dim is also affected (typically biases are per-output, so unchanged — but verify).

- [ ] **Step 4: Modify layout_fingerprint_seed**

Find the seed string and rename the entry:

```rust
// Before: "PARAM_W_DIR_FC1=17;..."
// After:  "PARAM_W_DIR_FC1_AUX1=17;..."
```

The `_AUX1` suffix marks this as the SP14-modified shape. Per `feedback_no_legacy_aliases`, no `_DEPRECATED` shim.

- [ ] **Step 5: Initialize new weight column to zero**

In trainer construction, after Xavier-init or other weight initialisers run, find where direction Q-head weights are written. Add:

```rust
        // SP14 forward wire: zero-initialise the new aux_softmax_diff column
        // of w_b0fc. Index range for the new column: row SH2 of the [adv_h, SH2+1] weight matrix.
        // Layout: row-major, so col = SH2 corresponds to elements
        // [k*adv_h + SH2 for k in 0..adv_h] within the flat tensor — implementer
        // to verify against actual storage layout.
```

The exact zero-init pattern depends on how the codebase stores w_b0fc (row vs column major). Read the existing init for w_b1fc (which has a similar concat shape `[adv_h, SH2 + b0]` per Phase 0) for guidance.

- [ ] **Step 6: Add fingerprint regression test**

Append to `crates/ml/tests/sp14_oracle_tests.rs`:

```rust
/// Test B.8: Layout fingerprint bumps when direction Q-head input dim grows.
#[test]
fn layout_fingerprint_bumps_after_sp14_wire() {
    use ml::cuda_pipeline::gpu_dqn_trainer::LAYOUT_FINGERPRINT_CURRENT;
    // Pre-SP14 fingerprint (recorded before the rename).
    // Get this value by checking out HEAD~1 and printing LAYOUT_FINGERPRINT_CURRENT.
    const PRE_SP14_FINGERPRINT: u64 = 0x0000_0000_0000_0000;  // implementer: replace with actual

    assert_ne!(
        LAYOUT_FINGERPRINT_CURRENT, PRE_SP14_FINGERPRINT,
        "SP14 forward wire must bump fingerprint via _AUX1 suffix rename. \
         current = {:#018x}, pre-SP14 = {:#018x}",
        LAYOUT_FINGERPRINT_CURRENT, PRE_SP14_FINGERPRINT
    );
}
```

To get the pre-SP14 fingerprint value, the implementer can run before this task lands:

```bash
git stash
cargo test layout_fingerprint --no-run 2>&1
# Run a small print test or use cargo expand to dump LAYOUT_FINGERPRINT_CURRENT
```

- [ ] **Step 7: Build + test**

```bash
SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
cargo test -p ml --test sp14_oracle_tests layout_fingerprint_bumps_after_sp14_wire 2>&1 | tail -5
```
Expected: clean; test passes.

- [ ] **Step 8: Commit B.8**

```bash
git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \
        crates/ml/tests/sp14_oracle_tests.rs
git commit -m "feat(sp14): B.8 — direction Q-head input dim SH2 → SH2+1 + fingerprint

Forward wire: direction Q-head's first FC weight tensor input dim
grows by 1 to accept aux_softmax_diff. compute_param_sizes() updated;
layout_fingerprint_seed() entry renamed to _AUX1 suffix forcing
fingerprint hash bump per Invariant 8 (old checkpoints invalidated).

New weight column zero-initialised in trainer construction so model
starts ignoring the new input — Q-head learns to use it through
gradient descent. Xavier-init would inject day-0 noise that the
trunk would have to denoise.

Per feedback_no_legacy_aliases: no _DEPRECATED shim on the renamed
fingerprint entry.

Test: layout_fingerprint_bumps_after_sp14_wire verifies hash differs
from pre-SP14 value.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>"
```

## Task B.9: Forward wire — concat launch + SGEMM consumer + launch order

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (direction Q-head forward path)

**Why:** Wire the new concat kernel into the forward pass. The direction Q-head's existing SGEMM call now consumes `sp14_dir_qaux_concat_scratch` instead of `h_s2`. Critical: aux head forward must complete before the concat launches (`pearl_canary_input_freshness_launch_order`).

- [ ] **Step 1: Locate direction Q-head SGEMM forward call**

```bash
grep -n "branch_0\|w_b0fc\|direction_forward\|q_dir.*forward" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs | head -10
```

Find the function/method that runs the direction Q-head's first FC. Note the call signature.

- [ ] **Step 2: Identify aux head forward call (must complete before concat)**

```bash
grep -n "aux_next_bar_forward\|aux_heads_forward\|aux head forward" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs | head -5
```

Note the line where aux forward runs. The new concat must launch AFTER this in graph capture order.

- [ ] **Step 3: Insert concat launch before direction SGEMM**

Edit the direction Q-head forward path. Before the existing SGEMM that takes `h_s2`:

```rust
        // SP14 forward wire: pre-SGEMM concat [h_s2 ; aux_softmax_diff].
        // Per pearl_canary_input_freshness_launch_order: aux head forward
        // (aux_next_bar_forward at line N) must complete before this concat
        // reads self.aux_nb_softmax_buf. Both run on the same stream so
        // sequential submission enforces the dep — but in graph capture
        // verify the concat's launch comes after aux forward's launch in
        // the captured node graph (otherwise reorder may break it).
        let bsi = batch_size as i32;
        let sdi = self.config.shared_h2 as i32;
        let blocks = ((batch_size * (self.config.shared_h2 + 1) + 255) / 256) as u32;
        unsafe {
            self.stream.launch_builder(&self.sp14_dir_concat_qaux_kernel)
                .arg(&self.h_s2_buf)            // [B, SH2]
                .arg(&self.aux_nb_softmax_buf)  // [B, K=2]
                .arg(&mut self.sp14_dir_qaux_concat_scratch)  // [B, SH2+1]
                .arg(&bsi)
                .arg(&sdi)
                .launch(LaunchConfig {
                    grid_dim: (blocks.max(1), 1, 1),
                    block_dim: (256, 1, 1),
                    shared_mem_bytes: 0,
                })
                .map_err(|e| MLError::ModelError(format!("dir_concat_qaux: {e}")))?;
        }

        // Direction Q-head SGEMM now consumes sp14_dir_qaux_concat_scratch
        // (shape [B, SH2+1]) instead of h_s2 (shape [B, SH2]).
        // Update the existing cublasLt SGEMM call to point at the scratch
        // buffer and use shared_h2 + 1 as the K dim.
```

The implementer needs to update the existing SGEMM K dim from `shared_h2` → `shared_h2 + 1` and the input pointer from `h_s2_buf` → `sp14_dir_qaux_concat_scratch`. This is a 2-line change in the existing forward call.

- [ ] **Step 4: Verify graph capture order**

If the codebase uses CUDA Graph capture, ensure the captured graph has aux_forward → concat → SGEMM as a serial chain. Look for `cudaGraphCapture` or similar:

```bash
grep -n "graph_capture\|cuGraphCapture\|graph_exec_launch" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs | head -5
```

If graph capture is used, the in-stream-order submission ensures correct dep — but verify with a small smoke that the model converges.

- [ ] **Step 5: Build check**

```bash
SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -10
```

Expected: clean if pointer types and dim args align.

- [ ] **Step 6: Commit B.9**

```bash
git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
git commit -m "feat(sp14): B.9 — wire forward concat into direction Q-head SGEMM

Direction Q-head's first FC SGEMM now consumes sp14_dir_qaux_concat_scratch
(shape [B, SH2+1]) instead of h_s2 (shape [B, SH2]). The concat kernel
runs immediately before the SGEMM in the same stream, enforcing the
pearl_canary_input_freshness_launch_order constraint.

K dim: shared_h2 → shared_h2 + 1.
Input ptr: h_s2_buf → sp14_dir_qaux_concat_scratch.

Aux head forward (aux_next_bar_forward writing aux_nb_softmax_buf)
already runs earlier in the per-step pipeline; sequential same-stream
submission ensures the concat reads fresh aux predictions.

After this commit, the wire is FULLY ACTIVE in forward but BACKWARD
gradient flow is ungated (full Q-loss flows back to aux). The EGF
pearl gating happens in B.10. Per feedback_no_partial_refactor:
this intermediate state is functional (the model learns; aux gets
co-trained by Q-loss) but not yet behavior-protected by the gate.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>"
```

## Task B.10: Backward wire — gradient scaling at the wire column

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (direction Q-head backward path)

**Why:** Critical safety: the EGF gate must scale `dL/dx[wire_col]` (the input gradient flowing back from Q-loss to aux head's softmax) by `α_grad_smoothed`. NOT `dL/dW[wire_col]` (which would prevent Q-head from learning to use the wire). Per the spec's B.4 clarification.

- [ ] **Step 1: Locate direction Q-head backward path**

```bash
grep -n "backward\|grad.*dir\|d_h_s2.*dir\|q_dir.*backward" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs | head -10
```

Find where direction Q-head computes `dL/dx` for h_s2. Post-SP14, this computes `dL/dx` for `[h_s2 ; aux_softmax_diff]` — a column-(SH2+1) vector.

- [ ] **Step 2: Add gradient column extraction + scaling**

Edit the backward path. After the SGEMM backward computes `dL/dx_concat` (shape `[B, SH2+1]`):

```rust
        // SP14 EGF backward: scale the new wire column gradient by α_grad_smoothed.
        // dL/dx_concat is shape [B, SH2 + 1]; column SH2 is the gradient
        // flowing back to aux_softmax_diff (and thence to aux head's softmax).
        //
        // Per pearl design: scale ONLY the input gradient flowing back to aux
        // (dL/dx[wire_col]), NOT the weight gradient (dL/dW[wire_col]).
        // The Q-head's weight column updates freely from Q-loss; only the
        // gradient propagating back to aux is gated.
        //
        // After scaling, the SH2-column-prefix of dL/dx_concat propagates
        // back to h_s2 (existing path); the wire_col flows back to
        // aux_nb_softmax_buf as -d_softmax_diff = -d(s1 - s0) = (-d_s1, +d_s0)
        // contribution. The aux head's existing backward chain consumes this
        // additional gradient stream.

        // Read α_grad_smoothed from ISV (host-side cached snapshot is fine
        // since it's per-step). For graph-safe pattern, use a device-pointer
        // arg to ALPHA_GRAD_SMOOTHED slot.
        // Implementer: launch a small scaling kernel that does
        //     dx_concat[b, SH2] *= isv[ALPHA_GRAD_SMOOTHED_INDEX]
        // for each b.
```

This is best implemented as a tiny inline scaling kernel or fused into the existing backward chain. Implementer chooses the pattern that fits the existing graph capture.

- [ ] **Step 3: Build check**

```bash
SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
```

- [ ] **Step 4: Commit B.10**

```bash
git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
git commit -m "feat(sp14): B.10 — backward wire gradient scaling by α_grad_smoothed

Critical safety mechanism: scale dL/dx[wire_col] (input gradient
flowing back from Q-loss through wire to aux head) by ALPHA_GRAD_SMOOTHED
ISV slot value, but NOT dL/dW[wire_col] (Q-head's own weight gradient
unchanged — Q-head learns to use the wire freely).

After this commit, the EGF pearl actively gates Q-loss flow to aux:
- α_grad ≈ 0 → stop-gradient (aux trains only on CE; Q-loss can't pull)
- α_grad ≈ 1 → full co-training enabled

Combined with B.7-B.9 (forward wire), the chain is now:
  1. Aux forward → aux_nb_softmax_buf
  2. dir_concat_qaux_kernel reads softmax → writes scratch
  3. Direction Q-head SGEMM reads scratch
  4. Q-loss backward → dL/dx_concat
  5. Wire-column gradient scaled by α_grad_smoothed
  6. Scaled gradient propagates back to aux_softmax → aux backward chain

Per feedback_no_partial_refactor: this is the architectural-atomic
moment where the wire becomes safety-protected.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>"
```

## Task B.11: Producer chain orchestrator — wire EGF kernels into per-step training loop

**Files:**
- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (per-step or per-epoch hook)
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (launcher methods if needed)

**Why:** The 3 EGF producer kernels (q_disagreement, alpha_grad, gradient_hack_detect) need to launch at appropriate points in the training loop. Per-step: q_disagreement (every action selection), alpha_grad (every backward pass; reads q_disagreement output). Per-epoch: gradient_hack_detect.

- [ ] **Step 1: Identify per-step hook points in training_loop.rs**

```bash
grep -n "step\|action_select\|backward\|gradient_apply" crates/ml/src/trainers/dqn/trainer/training_loop.rs | head -20
```

Locate the per-step training body where action selection completes (have q_logits available) and where backward starts (need α_grad_smoothed).

- [ ] **Step 2: Wire q_disagreement_update launch (per step)**

After action selection produces `q_dir_logits`:

```rust
        // SP14 B.11: per-step q_disagreement update. Reads aux_nb_softmax_buf
        // and q_dir_logits, writes Q_DISAGREEMENT_SHORT/LONG_EMA + variance to ISV.
        unsafe {
            // Launch q_disagreement_update_kernel with appropriate args.
            // Implementer verifies argument layout matches kernel signature.
        }
```

- [ ] **Step 3: Wire alpha_grad_compute launch (per step, before backward)**

After q_disagreement updates ISV:

```rust
        // SP14 B.11: alpha_grad compute. Reads driver signals from ISV,
        // writes ALPHA_GRAD_RAW + ALPHA_GRAD_SMOOTHED for the backward pass to consume.
        // Compute warmup_gate from steps_in_fold.
        let warmup_gate = ((steps_in_fold as f32) / (WARMUP_STEPS as f32)).min(1.0);
        unsafe {
            // Launch alpha_grad_compute_kernel.
        }
```

- [ ] **Step 4: Wire gradient_hack_detect launch (per epoch, end of epoch)**

```rust
        // SP14 B.11: end-of-epoch gradient hack detection. Decrement lockout,
        // check trigger conditions, force gate1 closed if triggered.
        unsafe {
            // Launch gradient_hack_detect_kernel.
        }
```

- [ ] **Step 5: Build check**

```bash
SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -10
```

- [ ] **Step 6: Commit B.11**

```bash
git add crates/ml/src/trainers/dqn/trainer/training_loop.rs \
        crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
git commit -m "feat(sp14): B.11 — orchestrator wire-up for 3 EGF producer kernels

Per-step launches:
- q_disagreement_update_kernel (after action select)
- alpha_grad_compute_kernel (before backward; consumes q_disagreement)

Per-epoch launch:
- gradient_hack_detect_kernel (end of epoch; circuit breaker)

Launch order in graph capture:
  1. Forward (existing)
  2. Action select (existing) → q_dir_logits available
  3. q_disagreement_update_kernel → updates ISV[381..382, 387]
  4. alpha_grad_compute_kernel → consumes q_disagreement, aux_dir_acc;
     writes ISV[383..391]
  5. Backward (existing) — wire-column scaling reads ALPHA_GRAD_SMOOTHED
  6. (end of epoch) gradient_hack_detect_kernel

The full EGF pearl is now ACTIVE: forward wire feeds aux signal to
direction Q-head; backward wire is gated by adaptive α_grad earned by
aux competence + Q-head disagreement.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>"
```

## Task B.12: HEALTH_DIAG pearl_egf_diag emit

**Files:**
- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (HEALTH_DIAG emit block)

- [ ] **Step 1: Find the existing aux HEALTH_DIAG emit**

```bash
grep -n 'aux_dir_acc target=\|HEALTH_DIAG.*aux\|pearl_' crates/ml/src/trainers/dqn/trainer/training_loop.rs | head -10
```

Identify the per-epoch emit block (post-B1.0, around line 4858-4891).

- [ ] **Step 2: Add pearl_egf_diag emit line**

After the existing aux emits:

```rust
        // SP14 EGF diagnostic emit
        let alpha_smoothed = isv_snap[ALPHA_GRAD_SMOOTHED_INDEX];
        let alpha_raw      = isv_snap[ALPHA_GRAD_RAW_INDEX];
        let beta_adaptive  = isv_snap[BETA_RATE_LIMITER_ADAPTIVE_INDEX];
        let k_aux_curr     = isv_snap[K_AUX_ADAPTIVE_INDEX];
        let k_q_curr       = isv_snap[K_Q_ADAPTIVE_INDEX];
        let var_aux        = isv_snap[AUX_DIR_ACC_VARIANCE_EMA_INDEX];
        let var_q          = isv_snap[Q_DISAGREEMENT_VARIANCE_EMA_INDEX];
        let var_alpha      = isv_snap[ALPHA_GRAD_RAW_VARIANCE_EMA_INDEX];
        let q_dis_short    = isv_snap[Q_DISAGREEMENT_SHORT_EMA_INDEX];
        let q_dis_long     = isv_snap[Q_DISAGREEMENT_LONG_EMA_INDEX];
        let gate1_state    = isv_snap[GATE1_OPEN_STATE_INDEX] as i32;
        let post_open_min  = isv_snap[AUX_DIR_ACC_POST_OPEN_MIN_INDEX];
        let lockout        = isv_snap[GRADIENT_HACK_LOCKOUT_REMAINING_INDEX] as i32;
        let warmup         = ((steps_in_fold as f32) / (WARMUP_STEPS as f32)).min(1.0);

        info!(
            target: "ml::trainers::dqn::trainer::metrics",
            "HEALTH_DIAG[{}]: pearl_egf_diag α_smoothed={:.4} α_raw={:.4} β={:.3} \
             gate1={:.3} gate2={:.3} k_aux={:.2} k_q={:.2} \
             var_aux={:.5} var_q={:.5} var_alpha={:.5} \
             q_dis_short={:.4} q_dis_long={:.4} \
             gate1_state={} aux_post_open_min={:.3} lockout_remaining={} \
             warmup={:.2}",
            epoch,
            alpha_smoothed, alpha_raw, beta_adaptive,
            /* gate1 = */ {/* recompute or read from kernel output */},
            /* gate2 = */ {/* recompute or read from kernel output */},
            k_aux_curr, k_q_curr,
            var_aux, var_q, var_alpha,
            q_dis_short, q_dis_long,
            if gate1_state == 1 { "open" } else { "closed" },
            post_open_min, lockout,
            warmup
        );
```

The implementer may add additional ISV slots for gate1/gate2 sigmoid outputs if separate visibility is needed; otherwise compute inline from k_aux/k_q + driver signals.

- [ ] **Step 3: Build check**

```bash
SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
```

- [ ] **Step 4: Commit B.12**

```bash
git add crates/ml/src/trainers/dqn/trainer/training_loop.rs
git commit -m "feat(sp14): B.12 — HEALTH_DIAG pearl_egf_diag per-epoch emit

Adds new HEALTH_DIAG line per epoch with all 14 EGF diagnostic fields:
- α_smoothed, α_raw, β (adaptive rate limiter)
- gate1, gate2 (current sigmoid outputs)
- k_aux, k_q (adaptive sigmoid steepness)
- var_aux, var_q, var_alpha (Welford variance EMAs)
- q_dis_short, q_dis_long (disagreement EMAs)
- gate1_state (open/closed)
- aux_post_open_min, lockout_remaining (circuit breaker state)
- warmup (epoch-fraction ramp)

Without these emit fields, diagnosing α_grad behavior in smoke logs
is forensic-only. With them, the full chain
  variance → adaptive scalars → gates → α_grad → backward scaling
is observable per epoch.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>"
```

## Task B.13: Integration test — end-to-end forward+backward with α_grad gating

**Files:**
- Modify: `crates/ml/tests/sp14_oracle_tests.rs` (3-5 integration tests)

- [ ] **Step 1: Write integration test scaffolding**

Add tests covering:
1. Forward wire integrity: synthetic h_s2 + aux_softmax → expected concat output, then SGEMM produces expected logits
2. Backward gating with α_grad=0: weight column dW updates normally; aux gradient = 0
3. Backward gating with α_grad=1: weight column dW + aux gradient both update
4. Backward gating with α_grad=0.5: half-strength aux gradient

Each test loads the relevant cubins, sets up B=4 SH2=8 inputs, runs forward+backward, checks gradients.

- [ ] **Step 2: Run all SP14 oracle tests**

```bash
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp14_oracle_tests \
    --features cuda -- --ignored --nocapture 2>&1 | tail -30
```

Expected: all 12-15 tests pass.

- [ ] **Step 3: Run snapshot stability test**

```bash
SQLX_OFFLINE=true cargo test -p ml --lib snapshot_size_is_stable 2>&1 | tail -5
```

Expected: snap remains 149*4=596 bytes (SP14 doesn't touch HEALTH_DIAG snap-words; new fields go in pearl_egf_diag emit only).

- [ ] **Step 4: Commit B.13**

```bash
git add crates/ml/tests/sp14_oracle_tests.rs
git commit -m "test(sp14): B.13 — integration tests for forward+backward EGF chain

5 end-to-end integration tests covering:
- Forward wire integrity (concat → SGEMM)
- Backward with α_grad=0 (stop-gradient at aux)
- Backward with α_grad=1 (full co-training)
- Backward with α_grad=0.5 (half-strength)
- Schmitt-trigger transition under simulated 5-epoch training

All 12-15 SP14 oracle tests pass on RTX 3050 Ti.
HEALTH_DIAG snap_size_is_stable holds at 149*4=596 bytes (SP14 emits
to pearl_egf_diag log line, not snap-words).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>"
```

## Task B.14: Audit doc + Smoke A2-B validation

**Files:**
- Modify: `docs/dqn-wire-up-audit.md` (new SP14 Layer B section)
- Run: `./scripts/argo-smoke.sh`

- [ ] **Step 1: Add SP14 Layer B audit section**

Append to `docs/dqn-wire-up-audit.md` a new section "## SP14 Layer B — Aux→Q wire + EGF pearl (2026-05-XX)" documenting:
- Forward wire path (h_s2 → concat → SGEMM)
- Backward gating mechanism (dL/dx[wire_col] × α_grad_smoothed)
- 13 ISV slots
- 4 new kernels
- Files-touched table
- Hard-rules-upheld checklist
- Build + test verification
- "Next: 30-epoch validation if smoke green"

- [ ] **Step 2: Push to origin**

```bash
git push origin sp11-reward-as-controlled-subsystem 2>&1 | tail -5
```

- [ ] **Step 3: Submit Smoke A2-B**

```bash
./scripts/argo-smoke.sh 2>&1 | tail -10
```

Note workflow name.

- [ ] **Step 4: Monitor with explicit kill criteria from spec B.7**

Per spec, kill if any of:
- Build/compile failure
- α_grad_smoothed stuck at 0.0 across 5 epochs in fold 1 when aux_dir_acc reaches 0.55+
- gate1_open_state never flips to 1 across full smoke
- aux_dir_acc collapses below 0.40 after gate1 opens AND lockout_remaining stays 0
- val_sharpe < 0 sustained across 2+ epochs
- n_up = 0 in HEALTH_DIAG aux_b1_diag (B1.1b regression)
- Forward stale-input pattern (aux healthy but α_grad oscillates wildly)

- [ ] **Step 5: Pull MinIO logs post-completion + verify pass criteria**

```bash
kubectl port-forward -n foxhunt svc/minio 9000:9000 &
sleep 3
# (mc setup as in A.4 Step 4)
mc cat foxminio/argo-logs/<wfname>/<wfname>/main.log | grep -E "pearl_egf_diag|aux_b1_diag" | tail -30
pkill -f "kubectl port-forward.*minio"
```

Expected:
- α_grad_smoothed varies; reaches > 0.5 in fold 1 when aux_dir_acc crosses target
- gate1_open_state flips to "open" in fold 1
- val_sharpe positive across all folds
- All 8 kill criteria not triggered
- aux_b1_diag still functional (B1.1b not regressed)

- [ ] **Step 6: Commit audit doc + smoke results**

```bash
git add docs/dqn-wire-up-audit.md
git commit -m "docs(sp14): B — Smoke A2-B validation results

Layer B EGF pearl validated on L40S smoke <wfname>:
- α_grad_smoothed reached <max> in fold 1 (gate opened)
- gate1_open_state flipped open at epoch <X> when aux_dir_acc=<Y>
- val_sharpe across folds: <values>
- val_win_rate: <values> (no climb expected at 5-epoch; 30-epoch is the gate)
- All 8 explicit kill criteria not triggered

Layer A + B chain complete on commit <sha>. Next: 30-epoch full
validation (multi-seed × multi-fold) tests user's hypothesis that
WR climbs above 0.50 by ep 20-25.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>"
```

---

# Self-Review

After writing the plan, verify:

**1. Spec coverage:**
- ✓ Sub-project A: 3 stability fixes + smoke validation (A.1-A.4)
- ✓ Sub-project B: 13 ISV slots (B.1), state-reset registry (B.2), 4 kernels with tests (B.3-B.6), trainer integration (B.7), forward wire (B.8-B.9), backward wire (B.10), orchestrator (B.11), HEALTH_DIAG (B.12), integration tests (B.13), audit + smoke (B.14)
- ✓ All 8 kill criteria from spec B.7 explicitly enumerated in Task B.14
- ✓ All hard rules listed in spec B.6 accounted for in commit messages

**2. Placeholder scan:**
- "implementer to identify exact tensor index" appears multiple times — these are intentional Phase 0 verification anchors. Each is followed by specific steps to do the verification. Acceptable per the spec's `feedback_trust_code_not_docs` discipline.
- `0x0000_0000_0000_0000` placeholder for pre-SP14 fingerprint in Task B.8 test — explicitly flagged with "implementer: replace with actual" — this is a value the implementer MUST compute by checking out HEAD~1.

**3. Type consistency:**
- Slot index `381` used consistently for `Q_DISAGREEMENT_SHORT_EMA_INDEX` across Steps 2, 3 of B.3, B.4 kernel, integration tests
- `α_grad_smoothed` consumed in B.10 backward; produced in B.4 kernel — names match
- `sp14_dir_qaux_concat_scratch` field name consistent across B.7, B.9
- `dir_concat_qaux_kernel` (snake_case) consistent in B.6 kernel + B.9 launch + struct field name `sp14_dir_concat_qaux_kernel`

---

# Execution Handoff

Plan complete and saved to `docs/superpowers/plans/2026-05-05-sp14-aux-q-wire-earned-gradient-flow.md`. Two execution options:

**1. Subagent-Driven (recommended)** — I dispatch a fresh subagent per task, review between tasks, fast iteration. Given the SP13 chain showed 6 implementer NEEDS_CONTEXT failures on similarly-sized atomic refactors, the subagent-driven approach with strong audit-doc + Phase 0 verification is the proven pattern. Each task has built-in TDD (write test, run failing, implement, run passing, commit).

**2. Inline Execution** — Execute tasks in this session using executing-plans, batch execution with checkpoints. Better for the small Layer A fixes; B's complexity may benefit from fresh subagent context.

Which approach?
