# SP14 Layer C: Separate Aux Trunk Implementation Plan

> **For agentic workers:** REQUIRED SUB-SKILL: Use superpowers:subagent-driven-development to implement this plan task-by-task. Steps use checkbox (`- [ ]`) syntax for tracking.

**Goal:** Give the auxiliary next-bar prediction head its own dedicated trunk (parallel to Q's GRN trunk) so it can shape a representation suited to next-bar prediction, while Q's trunk remains shaped purely by Q-loss. Delete the obsolete EGF α-coupling machinery (Coupling B + scale_wire kernel + Schmitt gate). Keep the forward feature wire (Coupling A: `dir_concat_qaux`) so Q's head still sees aux's directional softmax-diff as a feature.

**Architecture:** Encoder remains shared and Q-shaped. After encoder, the forward path branches: GRN trunk produces `h_s2_q` → Q-heads; new aux trunk (3-layer Linear→ELU MLP) produces `h_s2_aux` → aux head (K=2 softmax). Backward: Q-loss propagates Q-trunk → encoder; aux-loss propagates aux-trunk only (stop-gradient at encoder boundary). ISV is the control plane for aux trunk hyperparameters and adaptive bounds; ISV is NOT used as aux input features.

**Tech Stack:** Rust 1.85, cudarc 0.16+, CUDA 12.4, custom kernels. RTX 3050 Ti (sm_86) for dev tests; L40S (sm_89) for full validation.

**Branch:** Continue on `sp15-phase1-honest-numbers`, fork sub-worktree if scope warrants. Forks from `0ec791734` (4-fix HEAD validated by `train-thtxx`).

---

## Design Decisions Locked

| Question | Decision | Rationale |
|---|---|---|
| Aux trunk topology | 3-layer Linear→ELU→Linear→ELU→Linear (~150K params) | Capacity sufficient for K=2 prediction; GRN gating is Q-specific |
| Aux trunk input | Encoder output (post-attention, pre-GRN) | Reused; no need to re-encode raw features |
| Encoder backward from aux | **Stop-gradient** at encoder boundary | Encoder remains Q-shaped per today's stop-grad philosophy |
| Coupling A (forward feature wire `dir_concat_qaux`) | KEEP unchanged | Aux-as-feature-for-Q is the actual benefit |
| Coupling B (α-gated backward `sp14_scale_wire_col`) | DELETE atomically | Q co-training aux's head defeats "independent predictor" premise |
| α machinery (`alpha_grad_compute_kernel` + variance EMAs + Schmitt gate) | DELETE atomically with Coupling B | No remaining consumer after Coupling B deletion |
| Q-disagreement signal `ISV[383..390)` | KEEP | Useful diagnostic; producer remains active for HEALTH_DIAG |
| ISV role for aux trunk | Control plane only (hyperparams, adaptive bounds, diagnostic outputs) | NOT data plane — per-step broadcast violates per-env granularity |
| Aux loss weight (`aux_w`) | Keep ISV-driven (already exists) | No change |
| New aux trunk hyperparam slots | 6 ISV slots (LR, β1, β2, ε, grad_clip, h_s2_aux_rms_ema) | Mirrors SP5 Pearl 4 Adam pattern + GRN's h_s2_rms_ema |

## File Structure

**Create:**
- `crates/ml/src/cuda_pipeline/aux_trunk_forward_kernel.cu` — 3-layer MLP forward
- `crates/ml/src/cuda_pipeline/aux_trunk_backward_kernel.cu` — 3-layer MLP backward (stop-grad on encoder)
- `crates/ml/src/cuda_pipeline/h_s2_aux_rms_ema_kernel.cu` — adaptive scale producer for aux trunk output magnitude
- `crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs` — Rust wrapper (forward op + backward op + Adam state holders)
- `crates/ml/tests/aux_trunk_oracle_tests.rs` — gradient check + forward output check + stop-grad invariant

**Modify:**
- `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` — add 6 new slots (`AUX_TRUNK_LR_INDEX`, etc.) + mark α slots as deprecated
- `crates/ml/build.rs` — register new cubins
- `crates/ml/src/cuda_pipeline/mod.rs` — module declarations
- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — aux trunk param tensors + Adam state allocation in trainer constructor; delete α launch sites
- `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` — aux trunk forward + backward in collector chain; delete α launch sites; aux head now reads `h_s2_aux` instead of `h_s2`
- `crates/ml/src/trainers/dqn/state_reset_registry.rs` — register new aux trunk slots; remove α slots
- `crates/ml/src/cuda_pipeline/aux_heads_kernel.cu` — aux head input switches from `h_s2` to `h_s2_aux` (parameter rename + dimension verify)
- `docs/dqn-wire-up-audit.md` — close-out entry

**Delete:**
- `crates/ml/src/cuda_pipeline/alpha_grad_compute_kernel.cu`
- `crates/ml/src/cuda_pipeline/sp14_scale_wire_col_kernel.cu`
- `crates/ml/src/cuda_pipeline/q_disagreement_variance_kernel.cu` (if separate file; if inlined into alpha_grad_compute, deletion is in C.7)
- All α-related ISV slot constants in `sp14_isv_slots.rs` (`VAR_AUX_INDEX`, `VAR_Q_INDEX`, `VAR_ALPHA_INDEX`, `ALPHA_RAW_INDEX`, `ALPHA_SMOOTHED_INDEX`, `GATE1_OPEN_STATE_INDEX`, `GATE2_OPEN_STATE_INDEX`)

---

## Phase C.0: Pre-flight Audit

### Task C.0.1: Verify ISV slot vacancy + EGF wire-up snapshot

**Files:**
- Read: `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs`
- Read: `crates/ml/src/cuda_pipeline/sp15_isv_slots.rs`
- Read: `crates/ml/src/cuda_pipeline/alpha_grad_compute_kernel.cu`
- Read: `crates/ml/src/cuda_pipeline/sp14_scale_wire_col_kernel.cu`
- Modify: `docs/dqn-wire-up-audit.md` (audit-only entry, no code change)

- [ ] **Step 1: Confirm ISV vacancy after slot 443 (SP15 took [397..443))**

Run: `grep -nE "INDEX:.*[0-9]+ as usize" crates/ml/src/cuda_pipeline/sp15_isv_slots.rs | tail -5 && grep -nE "ISV_TOTAL_DIM" crates/ml/src/cuda_pipeline/sp15_isv_slots.rs`

Expected: max slot < 443; `ISV_TOTAL_DIM = 443`.

- [ ] **Step 2: Identify exact α-machinery launch sites**

Run: `grep -nE "alpha_grad_compute|scale_wire_col|VARIANCE_REF_AUX|SCHMITT_BAND|gate1_open|gate2_open" crates/ml/src/cuda_pipeline/*.rs crates/ml/src/cuda_pipeline/*.cu crates/ml/src/trainers/dqn/**/*.rs 2>/dev/null`

Capture line numbers. Save list inline in the audit doc.

- [ ] **Step 3: Add audit doc entry**

Append to `docs/dqn-wire-up-audit.md` top:
- Date 2026-05-07
- Title: "SP14 Layer C pre-flight — α machinery deletion targets"
- List of files + line numbers from Step 2
- Note: deletion is atomic in Phase C.7; pre-flight establishes "before" state for review

- [ ] **Step 4: Commit**

```bash
git add docs/dqn-wire-up-audit.md
git commit -m "docs(sp14-c-preflight): catalogue α-machinery launch sites pending atomic deletion"
```

---

## Phase C.1: ATOMIC — Delete α Machinery + Add Aux Trunk ISV Slots

> **CRITICAL: this is ONE atomic commit per `feedback_no_partial_refactor` and `feedback_no_legacy_aliases`** — no DEPRECATED stage. α machinery deletes simultaneously with new slot allocation. Build must be clean before AND after this commit.

### Task C.1.1: Atomic α deletion + 6 new aux trunk ISV slots

**Files** (per C.0 audit `a7a162d1f` — full launch site catalogue):
- Delete: `crates/ml/src/cuda_pipeline/alpha_grad_compute_kernel.cu`
- Delete: `crates/ml/src/cuda_pipeline/sp14_scale_wire_col_kernel.cu`
- Delete: `crates/ml/src/cuda_pipeline/gradient_hack_detect_kernel.cu` (EGF circuit breaker — reads α-machinery products; with EGF gone, breaker has no purpose)
- Modify: `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` (add 6 new + delete 9+ α slots)
- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (9 launch sites)
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (21 reference sites — kernel fields + launcher fns + load sites)
- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (4 sites)
- Modify: `crates/ml/src/trainers/dqn/trainer/mod.rs` (1 site)
- Modify: `crates/ml/src/trainers/dqn/batched_backward.rs` (1 site)
- Modify: `crates/ml/build.rs` (remove cubin entries)
- Modify: `crates/ml/src/cuda_pipeline/mod.rs` (remove module decls)
- Modify: `crates/ml/src/trainers/dqn/state_reset_registry.rs` (remove α slot reset entries; add new aux trunk slot entries)
- Modify: `crates/ml/tests/sp14_oracle_tests.rs` (delete tests for deleted machinery: `alpha_grad_adaptive_beta`, `alpha_grad_schmitt_hysteresis`, `gradient_hack_circuit_breaker_fires`. Preserve `dir_concat_qaux_correct` and `q_disagreement_*`.)
- Modify: `docs/dqn-wire-up-audit.md` (atomic deletion entry)

- [ ] **Step 1: Add 6 new slots after current SP15 max**

C.0 audit (commit `a7a162d1f`) confirmed `ISV_TOTAL_DIM = 444` (SP15 1.3.b-followup occupied 443). New slots start at **444**, not 443.

In `sp14_isv_slots.rs`, append after the SP15 block:

```rust
// SP14 Layer C — separate aux trunk control plane (added 2026-05-08)
pub const AUX_TRUNK_LR_INDEX: usize = 444;            // Adam LR for aux trunk params
pub const AUX_TRUNK_BETA1_INDEX: usize = 445;         // Adam β1
pub const AUX_TRUNK_BETA2_INDEX: usize = 446;         // Adam β2
pub const AUX_TRUNK_EPS_INDEX: usize = 447;           // Adam ε
pub const AUX_TRUNK_GRAD_CLIP_INDEX: usize = 448;     // grad-clip threshold
pub const H_S2_AUX_RMS_EMA_INDEX: usize = 449;        // aux trunk output magnitude EMA (mirrors H_S2_RMS_EMA_INDEX)
```

- [ ] **Step 2: Bump ISV_TOTAL_DIM** to `450` (was 444).

- [ ] **Step 3: DELETE α-related slot constants** (no deprecation step — atomic with kernel deletion in Step 7)

C.0 audit identified **9** α-machinery slots. Real names:
- `K_AUX_ADAPTIVE_INDEX` (385) — written by alpha_grad_compute
- `K_Q_ADAPTIVE_INDEX` (386) — written by alpha_grad_compute
- `BETA_RATE_LIMITER_ADAPTIVE_INDEX` (387) — written by alpha_grad_compute
- `ALPHA_GRAD_RAW_VARIANCE_EMA_INDEX`
- `GATE1_OPEN_STATE_INDEX` (single Schmitt gate; no GATE2)
- `ALPHA_GRAD_RAW_INDEX`
- `ALPHA_SMOOTHED_INDEX`
- Plus any var_aux/var_q-equivalent constants in the file (verify exact names via `grep -nE "VAR_AUX|VAR_Q|AUX_DIR_ACC_POST_OPEN_MIN" crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` — use whatever the file actually has).
- `AUX_DIR_ACC_POST_OPEN_MIN_INDEX` if present (read by gradient_hack_detect_kernel).

Delete the constant definitions outright. Leave the freed indices as RESERVED gap (do NOT compact — preserves checkpoint layout fingerprint compatibility).

- [ ] **Step 3a: Delete α kernel files**

```bash
git rm crates/ml/src/cuda_pipeline/alpha_grad_compute_kernel.cu
git rm crates/ml/src/cuda_pipeline/sp14_scale_wire_col_kernel.cu
git rm crates/ml/src/cuda_pipeline/gradient_hack_detect_kernel.cu
```

- [ ] **Step 3b: Remove from build.rs cubin manifest**

Search build.rs for `alpha_grad_compute`, `sp14_scale_wire_col`, `gradient_hack_detect` — delete those compilation entries.

- [ ] **Step 3c: Remove module decls** in `mod.rs`.

- [ ] **Step 3d: Remove all launch / consumer sites** (use C.0 audit `a7a162d1f` as the work list — 31 reference lines across 7 files):
  - `gpu_experience_collector.rs` (9 sites): remove `launch_alpha_grad_compute(...)`, `launch_sp14_scale_wire_col(...)`, and any `launch_gradient_hack_detect(...)` calls
  - `gpu_dqn_trainer.rs` (21 sites): remove `alpha_grad_compute_kernel: CudaFunction`, `sp14_scale_wire_col_kernel: CudaFunction`, `gradient_hack_detect_kernel: CudaFunction` fields; remove their loads in `compile_training_kernels`; remove the launcher fns
  - `training_loop.rs` (4 sites): remove any `gate1_open` / `α_smoothed` reads or HEALTH_DIAG entries
  - `trainer/mod.rs` (1 site)
  - `batched_backward.rs` (1 site)
  - `sp14_isv_slots.rs` (3 sites — beyond the 9 slot constants, also delete any helper fns / re-exports)

- [ ] **Step 3e: Delete obsolete oracle tests**

In `crates/ml/tests/sp14_oracle_tests.rs`, delete:
- `gpu::alpha_grad_adaptive_beta`
- `gpu::alpha_grad_schmitt_hysteresis`
- `gpu::gradient_hack_circuit_breaker_fires`

Preserve:
- `gpu::dir_concat_qaux_correct` (Coupling A — survives)
- `gpu::q_disagreement_*` (3 tests — diagnostic kernel survives)

- [ ] **Step 3f: Verify q_disagreement signal preserved**

`q_disagreement_short_ema` and `q_disagreement_long_ema` (`ISV[383..390)`) are NOT deleted — they remain useful diagnostics consumed by HEALTH_DIAG. Verify their producer kernel `q_disagreement_update_kernel.cu` is NOT in the deletion list.

Run: `ls crates/ml/src/cuda_pipeline/q_disagreement_update_kernel.cu` — file must still exist.

- [ ] **Step 4: Add reset_registry entries for new slots**

In `state_reset_registry.rs`, follow the existing per-slot pattern (find an existing entry like `H_S2_RMS_EMA_INDEX => StateReset::SentinelF32(0.0)` for reference). New entries:

| Slot | Reset value | Sentinel meaning |
|---|---|---|
| `AUX_TRUNK_LR_INDEX` | `SentinelF32(1e-4)` | Default aux-LR until Pearl-A bootstrap |
| `AUX_TRUNK_BETA1_INDEX` | `SentinelF32(0.9)` | Adam default |
| `AUX_TRUNK_BETA2_INDEX` | `SentinelF32(0.999)` | Adam default |
| `AUX_TRUNK_EPS_INDEX` | `SentinelF32(1e-8)` | Adam default |
| `AUX_TRUNK_GRAD_CLIP_INDEX` | `SentinelF32(1.0)` | Default norm-clip |
| `H_S2_AUX_RMS_EMA_INDEX` | `SentinelF32(0.0)` | Pearl-A first-observation bootstrap (replace, not blend) |

- [ ] **Step 4b: Add dispatch arm for `reset_named_state`** if the registry has a parallel by-name dispatch (mirrors SP5 Layer A bug-fix pattern).

- [ ] **Step 4c: Remove reset_registry entries for the 9 deleted α slots** (atomic with Step 3 deletion).

- [ ] **Step 4d: Update audit doc** — append entry to `docs/dqn-wire-up-audit.md`:
  - Title: "SP14 Layer C — α machinery deleted atomically (Phase C.1)"
  - List of deleted files (3 kernel files) + 9 deleted slot constants
  - Reference to C.0 pre-flight audit + commit SHA `a7a162d1f`
  - Note: `q_disagreement_*` slots preserved as diagnostic-only

- [ ] **Step 5: Verify**

Run: `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --tests --all-targets 2>&1 | tail -15`

Expected: **clean build, all targets**. Compile errors here mean a launch site or import was missed in Step 3d — fix before commit. Build must compile because:
- All α launches deleted simultaneously with the kernels they call
- All α slot reads deleted simultaneously with the slot constants
- Aux head's input still flows from h_s2 (Q's trunk) — that's fine; it just means aux head has slightly different gradient flow now (Q-loss can pull on aux head freely without α scaling). Aux head behavior will be replaced in Phase C.5 anyway.

Run remaining oracle tests to confirm nothing else broke:
`SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp14_oracle_tests --release -- --ignored --nocapture 2>&1 | tail -15`

Expected: 4 tests pass (dir_concat_qaux_correct + 3 q_disagreement_*). Three deleted tests no longer exist.

- [ ] **Step 6: Commit**

```bash
git add -A
git commit -m "$(cat <<'EOF'
refactor(sp14-c): atomic α machinery deletion + aux trunk ISV slot allocation

Phase C.1 of SP14 Layer C separate-aux-trunk refactor. Single atomic
commit per feedback_no_partial_refactor and feedback_no_legacy_aliases —
no DEPRECATED stage.

Deleted (per C.0 audit a7a162d1f):
- alpha_grad_compute_kernel.cu
- sp14_scale_wire_col_kernel.cu
- gradient_hack_detect_kernel.cu (EGF circuit breaker — α-coupled)
- 9 ISV slot constants (K_AUX_ADAPTIVE, K_Q_ADAPTIVE,
  BETA_RATE_LIMITER_ADAPTIVE, ALPHA_GRAD_RAW_VARIANCE_EMA,
  GATE1_OPEN_STATE, ALPHA_GRAD_RAW, ALPHA_SMOOTHED,
  AUX_DIR_ACC_POST_OPEN_MIN, plus var_aux/var_q if present)
- 31 reference sites across 7 files (collector, trainer, training_loop,
  batched_backward, mod.rs, build.rs cubin manifest)
- 3 oracle tests (alpha_grad_adaptive_beta, alpha_grad_schmitt_hysteresis,
  gradient_hack_circuit_breaker_fires)

Added:
- 6 aux trunk control plane ISV slots [444..450):
  AUX_TRUNK_LR (444), AUX_TRUNK_BETA1 (445), AUX_TRUNK_BETA2 (446),
  AUX_TRUNK_EPS (447), AUX_TRUNK_GRAD_CLIP (448), H_S2_AUX_RMS_EMA (449)
- Reset registry entries (Pearl-A first-observation bootstrap)
- ISV_TOTAL_DIM bumped 444 → 450

Preserved:
- dir_concat_qaux_kernel.cu (Coupling A: forward feature wire)
- q_disagreement_update_kernel.cu (diagnostic-only)
- q_disagreement_* ISV slots [383..390) (HEALTH_DIAG consumer)

Build clean; 4 surviving oracle tests pass. ISV layout: α slots left as
RESERVED gap for checkpoint fingerprint compatibility.

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

---

## Phase C.2: Aux Trunk Parameter Tensors + Adam State

### Task C.2.1: Allocate aux trunk params + Adam m/v in trainer constructor

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`
- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs`

**Topology** (sized to ~150K params total):

| Layer | Shape | Init |
|---|---|---|
| `aux_trunk_w1` | `[encoder_out_dim, 256]` | Kaiming-He fan_in=encoder_out_dim |
| `aux_trunk_b1` | `[256]` | zeros |
| `aux_trunk_w2` | `[256, 128]` | Kaiming-He fan_in=256 |
| `aux_trunk_b2` | `[128]` | zeros |
| `aux_trunk_w3` | `[128, AUX_HIDDEN_DIM]` (= existing aux head input dim) | Kaiming-He fan_in=128 |
| `aux_trunk_b3` | `[AUX_HIDDEN_DIM]` | zeros |

`AUX_HIDDEN_DIM` must match the aux head's expected input dimension. Read `aux_heads_kernel.cu` to find the constant (likely `SH2` or a separate aux-specific constant). If aux head currently consumes `SH2`, then `AUX_HIDDEN_DIM = SH2`.

- [ ] **Step 1: Add struct fields to the trainer struct** (`GpuDqnTrainer` in `gpu_dqn_trainer.rs`)

Find an existing param-group struct (e.g., GRN trunk params at the established location — grep for `grn_trunk_w1` or `aux_w1`). Mirror the pattern. Each tensor is a `cudarc::driver::CudaSlice<f32>`. Add 6 param tensors + 12 Adam state tensors (m, v for each).

- [ ] **Step 2: Allocation in `compile_training_kernels` or wherever existing params are allocated**

Look for the existing pattern (e.g., `let aux_w1 = stream.alloc_zeros::<f32>(...)?`). Allocate aux_trunk_w1/b1/w2/b2/w3/b3 + their Adam m/v. Store in trainer struct fields.

- [ ] **Step 3: Kaiming-He init via cuRAND**

Find existing init pattern (e.g., GRN init or VSN init). Kaiming-He fan_in formula: `std = sqrt(2.0 / fan_in)`. Apply to each weight tensor. Biases zero-init (already done by `alloc_zeros`).

- [ ] **Step 4: Mirror in collector** if the collector independently holds copies (per β migration pattern). Check `gpu_experience_collector.rs` for a parallel `aux_trunk_w1` field — if present, the trainer-side allocation must propagate (use `clone_async` or share the underlying `CudaSlice` Arc).

- [ ] **Step 5: Verify**

Run: `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --lib --tests 2>&1 | tail -10`

Expected: clean build. Param tensors allocated but not yet wired to forward/backward.

- [ ] **Step 6: Commit**

```bash
git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
git commit -m "feat(sp14-c): allocate aux trunk params (~150K) + Adam m/v state

3-layer MLP: encoder_out_dim → 256 → 128 → AUX_HIDDEN_DIM. Kaiming-He
weights, zero biases, separate Adam m/v buffers. Allocated in trainer
constructor + collector mirror per β-migration pattern. Not yet wired
to forward/backward — pure allocation."
```

---

## Phase C.3: Aux Trunk Forward Kernel + Oracle Test

### Task C.3.1: Forward kernel `aux_trunk_forward_kernel.cu`

**Files:**
- Create: `crates/ml/src/cuda_pipeline/aux_trunk_forward_kernel.cu`
- Modify: `crates/ml/build.rs` (register cubin)
- Modify: `crates/ml/src/cuda_pipeline/mod.rs` (module decl)
- Create: `crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs` (Rust wrapper, forward only — backward in C.4)

**Kernel signature** (mirrors `aux_next_bar_forward` in `aux_heads_kernel.cu`):

```c
extern "C" __global__ void aux_trunk_forward(
    const float* __restrict__ x_in,        // [B, encoder_out_dim] post-encoder
    const float* __restrict__ w1,          // [encoder_out_dim, 256]
    const float* __restrict__ b1,          // [256]
    const float* __restrict__ w2,          // [256, 128]
    const float* __restrict__ b2,          // [128]
    const float* __restrict__ w3,          // [128, AUX_HIDDEN_DIM]
    const float* __restrict__ b3,          // [AUX_HIDDEN_DIM]
    float* __restrict__ h_aux1_out,        // [B, 256] saved for backward
    float* __restrict__ h_aux2_out,        // [B, 128] saved for backward
    float* __restrict__ h_s2_aux_out,      // [B, AUX_HIDDEN_DIM] final aux trunk output
    int B,
    int ENCODER_OUT_DIM,
    int AUX_HIDDEN_DIM
);
```

Body: standard 3-layer MLP. Linear → ELU → Linear → ELU → Linear (no activation on final output — let aux head's own activation handle it). Use shared memory for intermediate hidden vectors per existing aux_heads kernel pattern.

**Reference:** `aux_heads_kernel.cu:117-189` for the 2-layer pattern; extend to 3 layers.

- [ ] **Step 1: Write the kernel file**

Follow `aux_heads_kernel.cu`'s style: thread-block per batch row, threadIdx.x sweeps the output dimension. ELU device function (already exists at `aux_heads_kernel.cu:38-44` — use the same).

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

Find the existing pattern (grep `aux_next_bar` in `build.rs`). Add `aux_trunk_forward` to the cubin compilation list.

- [ ] **Step 3: Add module decl in `mod.rs`** (if not auto-discovered).

- [ ] **Step 4: Write Rust wrapper `gpu_aux_trunk.rs`**

Pattern: `pub struct AuxTrunkForwardOps { kernel: CudaFunction, ... }`. `pub fn launch(stream, x_in, w1, b1, w2, b2, w3, b3, scratch_h_aux1, scratch_h_aux2, h_s2_aux_out, B, encoder_out_dim, aux_hidden_dim) -> Result<(), MLError>`. Mirror `gpu_aux_heads.rs` if exists, else `gpu_grn.rs`.

**Critical:** pre-load `CudaFunction` once at construction; do NOT `load_cubin` in the launch hot path (per `pearl_no_host_branches_in_captured_graph`).

- [ ] **Step 5: Write oracle test (forward output check)**

Create `crates/ml/tests/aux_trunk_oracle_tests.rs`:

```rust
//! Oracle tests for aux trunk forward+backward.

use cudarc::driver::*;
use ml::cuda_pipeline::gpu_aux_trunk::AuxTrunkForwardOps;

#[test]
#[ignore]  // GPU test
fn aux_trunk_forward_matches_numpy_reference() {
    // Setup: deterministic seed, B=4, enc=64, h1=256, h2=128, aux=128
    // 1. Allocate input + weights with fixed values
    // 2. Launch aux_trunk_forward
    // 3. Compute reference on host: h_aux1 = elu(x @ w1 + b1); h_aux2 = elu(h_aux1 @ w2 + b2); out = h_aux2 @ w3 + b3
    // 4. Compare element-wise; tol 1e-4 (f32 SGEMM precision)
}
```

Full body using `cudarc::driver::CudaContext` plus `ndarray` for the host reference.

- [ ] **Step 6: Run test, verify pass**

Run: `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test aux_trunk_oracle_tests --release -- --ignored aux_trunk_forward_matches_numpy_reference --nocapture 2>&1 | tail -10`

Expected: PASS.

- [ ] **Step 7: Commit**

```bash
git add crates/ml/src/cuda_pipeline/aux_trunk_forward_kernel.cu crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs crates/ml/src/cuda_pipeline/mod.rs crates/ml/build.rs crates/ml/tests/aux_trunk_oracle_tests.rs
git commit -m "feat(sp14-c): aux trunk forward kernel + Rust wrapper + oracle test

3-layer MLP forward (Linear→ELU→Linear→ELU→Linear). Pre-loaded
CudaFunction for graph-capture safety. Oracle test verifies bit-for-bit
match against numpy reference within 1e-4 tol."
```

---

## Phase C.4: Aux Trunk Backward Kernel + Oracle Test

### Task C.4.1: Backward kernel + stop-grad on encoder boundary

**Files:**
- Create: `crates/ml/src/cuda_pipeline/aux_trunk_backward_kernel.cu`
- Modify: `crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs` (add backward op)
- Modify: `crates/ml/src/cuda_pipeline/mod.rs`
- Modify: `crates/ml/build.rs`
- Modify: `crates/ml/tests/aux_trunk_oracle_tests.rs` (add gradient-check test)

**Kernel signature:**

```c
extern "C" __global__ void aux_trunk_backward(
    const float* __restrict__ dh_s2_aux_in,   // [B, AUX_HIDDEN_DIM] gradient from aux head
    const float* __restrict__ x_in,           // [B, ENCODER_OUT_DIM] saved fwd input
    const float* __restrict__ h_aux1,         // [B, 256] saved fwd
    const float* __restrict__ h_aux2,         // [B, 128] saved fwd
    const float* __restrict__ w1,             // [ENCODER_OUT_DIM, 256]
    const float* __restrict__ w2,             // [256, 128]
    const float* __restrict__ w3,             // [128, AUX_HIDDEN_DIM]
    float* __restrict__ dW1,                  // [ENCODER_OUT_DIM, 256] grad accum
    float* __restrict__ db1,                  // [256]
    float* __restrict__ dW2,                  // [256, 128]
    float* __restrict__ db2,                  // [128]
    float* __restrict__ dW3,                  // [128, AUX_HIDDEN_DIM]
    float* __restrict__ db3,                  // [AUX_HIDDEN_DIM]
    int B,
    int ENCODER_OUT_DIM,
    int AUX_HIDDEN_DIM
);
```

**Critical: NO `dx_in_out` parameter.** The kernel does NOT compute or write a gradient back to `x_in` (the encoder output). This is the stop-gradient at the encoder boundary — aux trunk's gradient terminates here.

**ELU derivative reminder:** `d_elu(x) = 1 if x > 0 else exp(x)` (continuous version: `elu(x) + 1 if x ≤ 0 else 1`).

**Reference:** `aux_heads_kernel.cu:481-614` for the 2-layer backward pattern; extend to 3 layers, omit dx propagation.

- [ ] **Step 1: Write backward kernel**

Three steps in reverse order:
1. d_logits = dh_s2_aux_in directly (output is linear, no activation)
2. dh_aux2 = (d_logits @ w3.T) ⊙ d_elu(h_aux2 pre-activation); accumulate dW3, db3
3. dh_aux1 = (dh_aux2 @ w2.T) ⊙ d_elu(h_aux1 pre-activation); accumulate dW2, db2
4. (NO Step 4 dx propagation — stop at w1 grad accumulation): accumulate dW1, db1

`atomicAdd` on dW gradients is forbidden per `feedback_no_atomicadd`. Use block-tree-reduce per `aux_heads_kernel.cu` precedent (lines ~520-570 in `aux_next_bar_backward`).

- [ ] **Step 2: Add backward op to Rust wrapper**

In `gpu_aux_trunk.rs`, add `pub struct AuxTrunkBackwardOps { kernel: CudaFunction, ... }` with `launch(...)`. Pre-load CudaFunction.

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

- [ ] **Step 4: Write gradient-check oracle test**

```rust
#[test]
#[ignore]
fn aux_trunk_backward_gradient_check() {
    // Setup deterministic input + weights
    // 1. Forward to get h_s2_aux
    // 2. Compute analytic gradient via aux_trunk_backward (using identity dh_s2_aux for clarity)
    // 3. Compute numerical gradient via finite differences:
    //    grad[i,j] ≈ (loss(W + ε·e_ij) - loss(W - ε·e_ij)) / (2ε)
    //    where loss = sum(h_s2_aux²)/2 (so dloss/dh_s2_aux = h_s2_aux directly)
    // 4. Compare analytic vs numerical for each of dW1/dW2/dW3, db1/db2/db3
    // 5. Tolerance: 1e-3 (numerical gradient noisy in f32)
}
```

- [ ] **Step 5: Write stop-grad invariant test**

```rust
#[test]
#[ignore]
fn aux_trunk_backward_does_not_write_dx() {
    // Allocate dx_in buffer with known sentinel pattern (e.g., 0xDEADBEEF / -1.0)
    // Call aux_trunk_backward
    // Verify dx_in unchanged — kernel must NOT touch encoder gradient buffer
}
```

This is critical: encoder gradient must remain Q-shaped only.

- [ ] **Step 6: Run tests**

Run: `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test aux_trunk_oracle_tests --release -- --ignored --nocapture 2>&1 | tail -15`

Expected: 3 tests pass (forward + gradient check + stop-grad invariant).

- [ ] **Step 7: Commit**

```bash
git add crates/ml/src/cuda_pipeline/aux_trunk_backward_kernel.cu crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs crates/ml/build.rs crates/ml/tests/aux_trunk_oracle_tests.rs
git commit -m "feat(sp14-c): aux trunk backward kernel + gradient check + stop-grad test

Backward propagates dh_s2_aux through w3/w2/w1 with block-tree-reduce
(no atomicAdd per feedback_no_atomicadd). Critical: kernel does NOT
write dx_in — encoder gradient remains Q-shaped only. Stop-grad
invariant verified by sentinel-pattern test."
```

---

## Phase C.4b: ISV-Driven Aux Prediction Horizon (multi-bar pivot)

> **Added 2026-05-08 per user design correction.** Aux's original label `(p_{t+1} > p_t)` from `aux_sign_label_kernel.cu` predicts next-bar at HFT scale — pure microstructure noise. User trades at HFT-MFT scale (multi-bar holds), so aux must predict over a horizon matching actual hold times. Without this, aux flatlines at ln(2) regardless of trunk separation: the *target itself* is unlearnable.

> **Scope:** label production + 1 new ISV slot + adaptive producer. Trunk math (C.2/C.3/C.4) is label-agnostic and unchanged. Cold-start sentinel H=60 (≈1hr @ 1min); first winning-trade observation triggers Pearl-A bootstrap (replace, not blend); thereafter Wiener-α EMA tracks avg winning hold time.

### Task C.4b.1: Migrate aux label kernels to ISV-driven H-bar horizon

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` (add 1 new slot, bump ISV_TOTAL_DIM 450→451; possibly add 1-2 hold-time tracker slots if not already present)
- Modify: `crates/ml/src/trainers/dqn/state_reset_registry.rs` (Pearl-A bootstrap entry)
- Modify: `crates/ml/src/cuda_pipeline/aux_sign_label_kernel.cu` (use H from ISV, mask `t+H >= total_bars`)
- Modify: `crates/ml/src/cuda_pipeline/aux_sign_label_per_step_kernel.cu` (same)
- Create: `crates/ml/src/cuda_pipeline/aux_horizon_update_kernel.cu` (adaptive producer reading avg winning hold time, writing to AUX_PRED_HORIZON_BARS_INDEX)
- Modify: `crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs` (add launcher op for producer)
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (kernel field + load + launch site)
- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (per-epoch boundary launch)
- Modify: `crates/ml/build.rs` (cubin entry)
- Modify: `crates/ml/src/cuda_pipeline/mod.rs` (module decl)
- Modify: `crates/ml/tests/aux_trunk_oracle_tests.rs` (4 tests: H-bar correctness, lookahead mask, Pearl-A bootstrap, Wiener-α EMA convergence)
- Modify: `docs/dqn-wire-up-audit.md` (Phase C.4b entry)

- [ ] **Step 1: Add ISV slot**

In `sp14_isv_slots.rs`:

```rust
// SP14 Layer C — aux prediction horizon (added 2026-05-08, multi-bar pivot)
pub const AUX_PRED_HORIZON_BARS_INDEX: usize = 450;
```

Bump `ISV_TOTAL_DIM` from 450 to **451**.

- [ ] **Step 2: Reset registry entry**

In `state_reset_registry.rs`:

```rust
AUX_PRED_HORIZON_BARS_INDEX => StateReset::SentinelF32(60.0),
```

Sentinel 60.0 = 60 bars ≈ 1 hour @ 1-min bars. This is the cold-start default. Pearl-A bootstrap will replace on first observation when adaptive producer lands (deferred follow-up).

Add the dispatch arm to `reset_named_state` if applicable.

- [ ] **Step 3: Modify `aux_sign_label_kernel.cu`**

Find the current label computation (likely a single line: `out_labels[i] = (p_fut > p_now) ? 1 : 0`). Replace with:

```c
extern "C" __global__ void aux_sign_label(
    const float* __restrict__ prices,    // [B*T, 1] price column from fxcache
    const float* __restrict__ isv,       // [ISV_TOTAL_DIM]
    int* __restrict__ out_labels,        // [B*T] -1 sentinel for masked
    int total_bars,
    int B,
    int T
) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx >= B * T) return;
    
    int b = idx / T;
    int t = idx % T;
    
    int H = (int)__float2int_rn(isv[AUX_PRED_HORIZON_BARS_INDEX]);
    if (H < 1) H = 1;        // floor: can't predict past
    if (H > 240) H = 240;    // ceiling: rollout buffer guard
    
    int t_future = t + H;
    if (t_future >= total_bars) {
        // Lookahead mask: no label available for this position
        out_labels[idx] = -1;  // sentinel for "ignore in loss"
        return;
    }
    
    float p_now = prices[b * total_bars + t];
    float p_fut = prices[b * total_bars + t_future];
    out_labels[idx] = (p_fut > p_now) ? 1 : 0;
}
```

Adjust signature to match the EXISTING kernel's parameter list. The key changes:
- Read `H` from ISV (read once per thread, fine — broadcast value)
- Index `t_future = t + H` instead of `t + 1`
- Mask `out_labels[idx] = -1` when `t_future >= total_bars`

**Critical:** the loss reduction kernel (`aux_next_bar_loss_reduce` in `aux_heads_kernel.cu`) must be checked — if it reads label as `int` and casts to softmax index, a `-1` sentinel will go out of bounds. Verify the consumer handles `-1` (skip the sample) or change the masking convention. The existing `out_labels` is `i32` per SP13 B1.1a (K=2 softmax CE).

If the loss-reduce kernel does NOT support a -1 mask, add masking handling:
- Easier alternative: instead of -1 mask, FORWARD-FILL with a deterministic but safe label (e.g., 0). The loss term will be incorrect but bounded; lookahead loss is ~6% of training data so signal is preserved.
- OR: add a `valid_mask` parameter that the loss-reduce kernel multiplies by.

Implementer choice: pick the simpler one given the existing kernel structure. Prefer the valid_mask approach (correctness > simplicity here since we're already touching label production).

- [ ] **Step 4: Modify `aux_sign_label_per_step_kernel.cu`**

Same migration. The per-step kernel runs at every rollout step (not at end-of-rollout); H is read from ISV at that step. Per-step kernel has access to forward bars in the rollout buffer; mask if `t + H` exceeds rollout length.

- [ ] **Step 5: Update launch sites**

Find launch sites in `gpu_dqn_trainer.rs` and `gpu_experience_collector.rs`. If they don't already pass the ISV pointer to the label kernel, add it. ISV ptr is already passed to dozens of other kernels — find an example pattern (`launch_aux_horizon_compute` or similar from neighboring launches) and mirror.

- [ ] **Step 5b: Identify avg winning hold time source in existing infrastructure**

Run:
```bash
grep -nE "winning_hold|win_hold|hold_time|HOLD_TIME|trade_close|trade_dur|HOLDING|holding_bars" \
    crates/ml/src/cuda_pipeline/*.rs crates/ml/src/cuda_pipeline/*.cu \
    crates/ml/src/cuda_pipeline/sp14_isv_slots.rs \
    crates/ml/src/cuda_pipeline/sp15_isv_slots.rs 2>/dev/null
```

SP12/SP15 introduced trade-close events (per-trade reward emission). Expected to find: per-env "ticks since position opened" tracker + on-close emission of duration to a hold-time EMA.

Three possible findings:

**Case A: Existing slot tracks avg winning hold time directly** (e.g., `AVG_WIN_HOLD_TIME_EMA_INDEX`).
→ Producer reads from that slot. Done.

**Case B: Existing infrastructure tracks per-env open-tick + emits on close, but no aggregate slot.**
→ Add a small aggregate kernel: averaged-on-close winning hold time → new ISV slot `AVG_WIN_HOLD_TIME_BARS_INDEX = 451`. Producer reads from that.
→ Bumps ISV_TOTAL_DIM 451→452. Document.

**Case C: No existing hold-time infrastructure.**
→ Out of scope for this phase; report BLOCKED. (We don't want to introduce trade-close tracking infrastructure here — that's SP15 territory.)

Document which case you found in the audit doc.

- [ ] **Step 5c: Write `aux_horizon_update_kernel.cu` (adaptive producer)**

```c
extern "C" __global__ void aux_horizon_update(
    float* __restrict__ isv,            // [ISV_TOTAL_DIM]
    int isv_h_idx,                      // AUX_PRED_HORIZON_BARS_INDEX = 450
    int isv_h_target_idx,               // index of avg winning hold time slot (Case A or B above)
    int isv_h_target_var_idx,           // optional: variance EMA of target for Wiener-α (-1 if not available)
    float h_min,                        // 1.0
    float h_max,                        // 240.0
    float sentinel_h                    // 60.0 (cold-start default)
) {
    if (blockIdx.x != 0 || threadIdx.x != 0) return;  // single-thread compute
    
    float h_target = isv[isv_h_target_idx];
    float h_current = isv[isv_h_idx];
    
    // No winning trade observed yet — keep sentinel
    if (h_target <= 0.0f) {
        isv[isv_h_idx] = sentinel_h;  // keep sentinel
        return;
    }
    
    // Pearl-A first-observation bootstrap: replace sentinel directly
    if (fabsf(h_current - sentinel_h) < 1e-6f) {
        // First valid observation: replace sentinel directly
        float h_clamped = fmaxf(h_min, fminf(h_max, h_target));
        isv[isv_h_idx] = h_clamped;
        return;
    }
    
    // Steady-state Wiener-α EMA blend
    float alpha;
    if (isv_h_target_var_idx >= 0) {
        // Wiener-optimal: α = sample_var / (sample_var + diff_var + ε)
        // For horizon (slow signal): use signal-derived α with floor for stability
        float sample_var = isv[isv_h_target_var_idx];
        float diff = h_target - h_current;
        float diff_var = diff * diff;  // single-sample approximation
        alpha = sample_var / (sample_var + diff_var + 1e-6f);
        alpha = fmaxf(0.001f, fminf(0.1f, alpha));  // bounds: slow horizon
    } else {
        alpha = 0.01f;  // fixed slow blend for stable horizon
    }
    
    float h_blended = (1.0f - alpha) * h_current + alpha * h_target;
    isv[isv_h_idx] = fmaxf(h_min, fminf(h_max, h_blended));
}
```

Single-thread kernel (one block, one thread) — compute is trivial; just need to write a single ISV slot. No atomicity concerns.

Per `pearl_first_observation_bootstrap`: sentinel = 60.0; first observation replaces directly (not blends).

Per `pearl_wiener_optimal_adaptive_alpha`: if target variance is available, use Wiener-α; else fixed 0.01 (slow blend).

Per `feedback_isv_for_adaptive_bounds`: bounds h_min=1, h_max=240 are constants (acceptable — fundamental floors/ceilings, not tuning parameters).

- [ ] **Step 5d: Wire producer**

In `gpu_aux_trunk.rs`, add `AuxHorizonUpdateOps` struct mirroring forward/backward op pattern. Pre-load CudaFunction. Launch sig matches kernel.

In `gpu_experience_collector.rs`, add launch at **per-epoch boundary** (horizon is slow-moving — epoch-level cadence is appropriate, not per-step). Find the existing per-epoch ISV update block (mirrors SP4 producer launches at end of epoch).

In `gpu_dqn_trainer.rs`, add `aux_horizon_update_ops: AuxHorizonUpdateOps` field; load in trainer constructor.

- [ ] **Step 6: Add label correctness oracle test**

In `aux_trunk_oracle_tests.rs`:

```rust
#[test]
#[ignore]
fn aux_sign_label_h_bar_horizon() {
    // 1. Synthetic price series: linearly increasing for B=2, T=200
    // 2. Set ISV[AUX_PRED_HORIZON_BARS_INDEX] = 30.0
    // 3. Launch aux_sign_label
    // 4. Verify out_labels[0..170] == 1 (price 30 bars ahead is higher)
    // 5. Verify out_labels[170..200] == -1 (lookahead masked)
    
    // Then with descending series: out_labels[0..170] == 0
    
    // Then with sine wave: verify labels match analytic (p(t+30) > p(t))
}
```

- [ ] **Step 7: Add lookahead mask invariant test**

```rust
#[test]
#[ignore]
fn aux_sign_label_lookahead_mask() {
    // Set H = 60. Set total_bars = 100.
    // Launch.
    // Verify out_labels[40..100] == -1 (cannot look 60 bars past total_bars - 60 = 40)
    // Verify out_labels[0..40] are valid (∈ {0, 1})
}
```

- [ ] **Step 7b: Pearl-A bootstrap test for horizon producer**

```rust
#[test]
#[ignore]
fn aux_horizon_pearl_a_bootstrap() {
    // 1. ISV[AUX_PRED_HORIZON_BARS_INDEX] = 60.0 (sentinel)
    // 2. ISV[h_target_idx] = 90.0 (first observation)
    // 3. Launch aux_horizon_update
    // 4. Assert ISV[AUX_PRED_HORIZON_BARS_INDEX] == 90.0 (REPLACE, not blend)
    
    // 5. ISV[h_target_idx] = 100.0 (subsequent observation)
    // 6. Launch
    // 7. Assert ISV[AUX_PRED_HORIZON_BARS_INDEX] in (90, 100) — blended, not replaced
}
```

- [ ] **Step 7c: Wiener-α convergence test**

```rust
#[test]
#[ignore]
fn aux_horizon_converges_to_steady_target() {
    // 1. Bootstrap ISV[AUX_PRED_HORIZON_BARS_INDEX] = 50.0 (post-bootstrap state)
    // 2. ISV[h_target_idx] = 100.0 (steady target)
    // 3. Launch aux_horizon_update 100 times
    // 4. Assert ISV[AUX_PRED_HORIZON_BARS_INDEX] converges to within 5% of 100.0
    //    (with α=0.01 fixed blend, after 100 iters: (1-0.01)^100 ≈ 0.366
    //     so distance halves every ~70 iters; 100 iters → ~63% closure → ~78.5)
    // 5. Adjust assertion to (1 - (1-α)^100) * 100 + (1-α)^100 * 50 ≈ 81.7
}
```

Test the actual convergence math, not a hand-wavy "approaches 100".

- [ ] **Step 7d: Cold-start "no winning trades" guard test**

```rust
#[test]
#[ignore]
fn aux_horizon_holds_sentinel_with_no_winning_trades() {
    // 1. ISV[AUX_PRED_HORIZON_BARS_INDEX] = 60.0 (sentinel)
    // 2. ISV[h_target_idx] = 0.0 (no winning trades observed yet)
    // 3. Launch
    // 4. Assert ISV[AUX_PRED_HORIZON_BARS_INDEX] == 60.0 (unchanged — kernel guard)
}
```

- [ ] **Step 8: Verify**

```bash
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --tests --all-targets 2>&1 | tail -10
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test aux_trunk_oracle_tests --release -- --ignored --nocapture 2>&1 | tail -20
```

Expected: clean build; all aux_trunk tests pass + 2 new label tests pass.

- [ ] **Step 9: Audit doc**

Append to `docs/dqn-wire-up-audit.md`:
- Title: "SP14 Layer C Phase C.4b — aux label horizon ISV-driven"
- Date 2026-05-08
- Rationale: HFT-scale next-bar prediction is unlearnable noise; user trades multi-bar
- Scope: 1 new ISV slot (450), 2 label kernels migrated, lookahead masking
- Cold-start: H=60 bars sentinel, adaptive producer deferred
- Open follow-up: ISV producer for `AUX_PRED_HORIZON_BARS_INDEX` from observed winning hold time (likely sourced from existing trade-close event tracking infrastructure introduced in SP15)

- [ ] **Step 10: Commit**

```bash
git add -A
git commit -m "$(cat <<'EOF'
feat(sp14-c): aux prediction horizon ISV-driven (multi-bar pivot)

Aux's original label was (p_{t+1} > p_t) — pure HFT-scale microstructure
noise that's unlearnable at our HFT-MFT trading frequency. Migrated to
(p_{t+H} > p_t) where H is read from ISV[AUX_PRED_HORIZON_BARS_INDEX].

Adaptive producer drives H from observed avg winning hold time:
- Pearl-A first-observation bootstrap: replace sentinel H=60 directly
  on first valid observation
- Steady-state Wiener-α EMA blend, slow (α≈0.01) for stable horizon
- "No winning trades yet" guard keeps sentinel until first valid observation

Lookahead truncation: labels at t where t+H >= total_bars are masked
(sentinel -1, loss-reduce skips). ~H/T fraction of training data lost
per fold; for H=60, T=1000 → ~6% acceptable.

Trunk math (C.2/C.3/C.4) unchanged — separate aux trunk is label-
agnostic. Validation in C.10 will use H=60 cold-start; if aux learns,
adaptive H lands as follow-up.

Phase C.4b of SP14 Layer C separate-aux-trunk refactor.

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

DO NOT push.

## Hard rules

1. **`feedback_no_partial_refactor`**: both label kernels (regular + per_step) migrate together. If only one is migrated, validation will be inconsistent (training-time labels vs eval-time labels diverge).
2. **Lookahead masking is correctness-critical**: a label at `t + H >= total_bars` reads garbage memory. Mask MUST be applied or the kernel produces undefined labels.
3. **Loss-reduce consumer**: verify the existing aux loss-reduce kernel handles the chosen mask convention. If `-1` causes OOB, use `valid_mask` parameter approach instead.
4. **No host branches in graph capture**: all logic in the label kernel runs on-device.

## Falsification gate (post-validation)

If C.10 L40S validation shows aux loss STILL flatlines at ln(2) with H=60:
- Check label distribution (rolling p(up) — if extremely imbalanced, that's the bug)
- Check encoder output normalization (regime-spike instability)
- Escalate to "aux pulls from PRE-attention encoder layer" or "aux trunk needs LayerNorm" — out of scope for this plan

---

## Phase C.5: Wire-Up — Split into C.5a (additive infra) + C.5b (atomic contract migration)

> Split authorized 2026-05-08 after C.5 implementer reported ~600-800 LOC scope across 4 files with subtle gradient-correctness windows. Split is architecturally legitimate (mirrors SP14 Layer A→B pattern):
> - **C.5a** is *additive* — allocates buffers + adds Adam launcher as dead code. No contract change; nothing in production calls the new code yet. Cannot break the build's behavior.
> - **C.5b** is *atomic* — kernel rename + revert zero-fills + wire-up redirects in ONE commit. ~300 LOC. This is the genuine contract migration.

## Phase C.5a: Additive Infrastructure — Buffers + Adam Launcher (DEAD CODE)

> No contract change. Build clean before, build clean after. Not yet wired.

### Task C.5a.1: Allocate fwd/bwd buffers + Adam launcher

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (allocate `h_s2_aux`, `h_aux1`, `h_aux2` saved-fwd; allocate aux trunk grad buffers if separate from params; add `launch_aux_trunk_adam_update` fn — dead, no caller yet)
- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (allocate rollout-sized `h_s2_aux`, `h_aux1`, `h_aux2`)
- Modify: `docs/dqn-wire-up-audit.md` (audit entry — additive infra, no behavior change)

- [ ] **Step 1: Allocate trainer-side saved-fwd buffers**

In `gpu_dqn_trainer.rs`, alongside existing aux buffers (find `aux_next_bar_logits` or similar), allocate:
- `h_s2_aux: CudaSlice<f32>` — `[B_max, AUX_HIDDEN_DIM]`
- `h_aux1: CudaSlice<f32>` — `[B_max, H1=256]` saved-for-backward intermediate
- `h_aux2: CudaSlice<f32>` — `[B_max, H2=128]` saved-for-backward intermediate
- `dh_s2_aux_accum: CudaSlice<f32>` — `[B_max, AUX_HIDDEN_DIM]` accumulator for upstream gradient (where aux_next_bar_backward + aux_regime_backward will write SAXPY-style after C.5b's revert)

These are dead until C.5b wires them up.

- [ ] **Step 2: Allocate collector-side rollout-sized buffers**

In `gpu_experience_collector.rs`, mirror trainer buffers but rollout-sized (`B * T_rollout` * dim).

- [ ] **Step 3: Allocate aux trunk gradient buffers**

C.2 allocated 6 separate weight/bias param tensors + 12 separate Adam m/v tensors. C.5a needs corresponding gradient buffers:
- `aux_trunk_w1_grad: CudaSlice<f32>` matching `aux_trunk_w1` shape
- ... same for b1, w2, b2, w3, b3 (6 grad tensors total)

OR, if you prefer the flat-buffer Adam pattern that OFI embed uses: introduce flat `aux_trunk_params_flat`, `aux_trunk_grad_flat`, `aux_trunk_m_flat`, `aux_trunk_v_flat` (each ~132K f32) with byte-offset views into w1/b1/w2/b2/w3/b3 sub-regions. This requires re-declaring C.2's separate buffers as views into the flat buffer.

**Implementer choice:** whichever fits cleanly into existing patterns. If 6-buffer approach is straightforward, prefer it (smaller refactor of C.2's work). If existing Adam launcher pattern strictly requires flat buffer, restructure to flat. Document the choice in audit entry.

- [ ] **Step 4: Write `launch_aux_trunk_adam_update`** (dead, no caller yet)

Mirror existing Adam launcher pattern. Search for `OFI` embed Adam path or similar single-param-group launcher (e.g., `launch_ofi_embed_adam`). The aux trunk lives outside the SP4-canonical 4-way Adam taxonomy (Q/Target/CQL/Aux per-branch), so a dedicated launcher is appropriate.

Reads from ISV:
- `lr` from `ISV[AUX_TRUNK_LR_INDEX]` (444)
- `β1` from `ISV[AUX_TRUNK_BETA1_INDEX]` (445)
- `β2` from `ISV[AUX_TRUNK_BETA2_INDEX]` (446)
- `ε` from `ISV[AUX_TRUNK_EPS_INDEX]` (447)
- `grad_clip_norm` from `ISV[AUX_TRUNK_GRAD_CLIP_INDEX]` (448)

Pre-step: global L2 norm of all 6 grad tensors → if exceeds clip threshold, scale all by `clip_norm / norm`. Then Adam step.

- [ ] **Step 5: Verify**

```bash
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --tests --all-targets 2>&1 | tail -10
```

Expected: clean build. Dead code may warn about unused fields/functions; that's OK because C.5b consumes them. Add `#[allow(dead_code)]` only if clippy promotes the warning to an error.

Run all oracle tests to confirm nothing regressed:
```bash
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test aux_trunk_oracle_tests --test sp14_oracle_tests --release -- --ignored --nocapture 2>&1 | tail -25
```

Expected: 12 tests pass (no regression — the new buffers + launcher are inert).

- [ ] **Step 6: Audit doc + commit**

```bash
git add -A
git commit -m "$(cat <<'EOF'
feat(sp14-c.5a): allocate aux trunk fwd/bwd buffers + Adam launcher (dead code)

Phase C.5a — additive infrastructure for the aux trunk wire-up.
Allocates saved-fwd buffers (h_s2_aux, h_aux1, h_aux2),
gradient buffers (6× aux_trunk_*_grad), accumulator
(dh_s2_aux_accum), and dedicated Adam launcher
(launch_aux_trunk_adam_update) reading β1/β2/ε/LR/grad-clip from
ISV[444..449).

No contract change. No call sites for the new launcher yet.
C.5b atomically wires these in.

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

DO NOT push.

## Phase C.5a-fixup: Add Missing Scaffolding (DEAD CODE)

> Authorized 2026-05-08 after C.5b implementer flagged 5 specific gaps (most importantly a graph-capture hazard) plus 2 missing buffer allocations from C.5a. C.5a-fixup is **additive-only** — same discipline as C.5a. Goal: by end of fixup, C.5b is a genuine ~300 LOC atomic flip with zero scaffolding work mixed in.

### Task C.5a-fixup.1: Allocate missing scratch buffers + collector ptr scaffolding + graph-capture audit

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (allocate `dh_aux1_pre_scratch`, `dh_aux2_pre_scratch`; add `aux_trunk_param_ptrs()` accessor)
- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (6× `u64` aux_trunk param ptr fields + `set_trainer_aux_trunk_param_ptrs` setter; new `exp_aux_trunk_forward_ops: AuxTrunkForwardOps` field; constructor populates it; **leave wire sites still calling `save_h_s2` — no semantic change**)
- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (call new `set_trainer_aux_trunk_param_ptrs` setter at the existing setup point near `set_trainer_params_ptr` — ~lines 1703 + 2317 per implementer report)
- Modify: `docs/dqn-wire-up-audit.md` (audit entry + graph-capture verification result)

- [ ] **Step 1: Allocate missing scratch buffers in trainer**

`aux_trunk_backward.launch` requires `dh_aux1_pre_scratch_ptr [B_max, AUX_TRUNK_H1=256]` and `dh_aux2_pre_scratch_ptr [B_max, AUX_TRUNK_H2=128]` per `gpu_aux_trunk.rs:266-267`. C.5a missed these. Add as `CudaSlice<f32>` zero-init alongside existing C.5a buffers.

- [ ] **Step 2: Collector aux trunk ptr scaffolding** (mirrors existing `set_trainer_params_ptr` pattern — find via grep)

```rust
// In GpuExperienceCollector struct:
aux_trunk_w1_ptr: u64,
aux_trunk_b1_ptr: u64,
aux_trunk_w2_ptr: u64,
aux_trunk_b2_ptr: u64,
aux_trunk_w3_ptr: u64,
aux_trunk_b3_ptr: u64,
exp_aux_trunk_forward_ops: AuxTrunkForwardOps,

// Setter:
pub fn set_trainer_aux_trunk_param_ptrs(&mut self, w1: u64, b1: u64, w2: u64, b2: u64, w3: u64, b3: u64) {
    self.aux_trunk_w1_ptr = w1;
    // ... etc
}
```

- [ ] **Step 3: Trainer-side accessor**

```rust
// In GpuDqnTrainer:
pub fn aux_trunk_param_ptrs(&self) -> (u64, u64, u64, u64, u64, u64) {
    use cudarc::driver::DevicePtr;
    (
        *self.aux_trunk_w1.device_ptr(),
        *self.aux_trunk_b1.device_ptr(),
        *self.aux_trunk_w2.device_ptr(),
        *self.aux_trunk_b2.device_ptr(),
        *self.aux_trunk_w3.device_ptr(),
        *self.aux_trunk_b3.device_ptr(),
    )
}
```

(Adjust to match the actual cudarc `device_ptr()` API — find an existing accessor like `embed_w_ptr()` and mirror.)

- [ ] **Step 4: Wire setter call in training_loop**

Find the existing call to `set_trainer_params_ptr` (~lines 1703 + 2317 per implementer report). Insert call to `set_trainer_aux_trunk_param_ptrs` at the same site, passing the trainer's accessor result.

- [ ] **Step 5: Graph-capture audit** (the showstopper concern)

Read `gpu_dqn_trainer.rs` aux_heads_backward orchestrator (around line 29865 per implementer report). Trace: does the call sit INSIDE a `cuStreamBeginCapture`/`end_capture` region?

If INSIDE: the existing aux backward CANNOT have host-side writes (step counter, pinned buffer). Verify the existing code by inspecting it — the existing aux_heads_backward should have a graph-capture-compatible structure (probably a separate function called pre-capture for any host writes). Document the existing pattern.

If OUTSIDE: the new aux_trunk Adam launch can sit alongside existing aux_heads_backward; host writes are fine.

Document the verification result in audit doc — this determines C.5b's wiring strategy.

- [ ] **Step 6: Verify**

```bash
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --tests --all-targets 2>&1 | tail -10
SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test aux_trunk_oracle_tests --test sp14_oracle_tests --release -- --ignored --nocapture 2>&1 | tail -15
```

Expected: clean build; 12 tests still pass (no regression — pure scaffolding).

- [ ] **Step 7: Commit**

```bash
git add -A
git commit -m "$(cat <<'EOF'
feat(sp14-c.5a-fixup): missing scratch buffers + collector ptr scaffolding (dead code)

Phase C.5a-fixup — completes C.5a's additive infrastructure:
- Allocate dh_aux1_pre_scratch + dh_aux2_pre_scratch (missed in C.5a)
- Collector struct gains 6× aux_trunk param ptr fields + setter
- Trainer gains aux_trunk_param_ptrs() accessor
- training_loop wires the setter at the existing set_trainer_params_ptr site
- Collector constructor populates AuxTrunkForwardOps

NO contract change. Wire sites still call save_h_s2. The setter and ptrs
are populated but unused by aux head consumers (still wired through Q's
trunk).

Graph-capture audit: <result documented in audit doc>

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

DO NOT push.

---

## Phase C.5b: Atomic Contract Migration — Aux Head Input Switches h_s2 → h_s2_aux

> ONE atomic commit per `feedback_no_partial_refactor`. With C.5a + C.5a-fixup providing all scaffolding (buffers, ptrs, accessor, setter, AuxTrunkForwardOps in collector), C.5b is now a genuine ~300 LOC atomic flip: 4 kernel param renames + 2 zero-fill reverts + 4 wire sites + Adam launch + ISV pinned-buffer population.

### Task C.5b.1: Atomic kernel rename + revert zero-fills + wire-up

> **CRITICAL: this is ONE atomic commit per `feedback_no_partial_refactor`.** All consumers of aux head's input migrate from `h_s2` to `h_s2_aux` simultaneously with the new kernel launches in collector. C.5a already allocated the buffers + Adam launcher (dead). C.5b makes them live atomically.

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (insert aux_trunk_forward post-graph; redirect aux_heads_fwd input ptr to h_s2_aux)
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (insert aux_trunk_forward in trainer fwd; insert aux_trunk_backward + Adam launch in bwd; redirect aux_heads_fwd/bwd input/grad ptrs)
- Modify: `crates/ml/src/cuda_pipeline/aux_heads_kernel.cu` (param rename `h_s2` → `h_s2_aux` in 4 kernels; **REVERT** zero-fills in `aux_next_bar_backward:599-638` and `aux_regime_backward:757-790` from commits `872bd7392` + `411a30473` — restore proper SAXPY-back-to-input-grad)
- Modify: `crates/ml/src/cuda_pipeline/gpu_aux_heads.rs` (Rust wrapper input/grad plumbing)
- Modify: `docs/dqn-wire-up-audit.md` (atomic migration entry; note C.0 stop-grad audit comments are now stale and removed)

- [ ] **Step 1: Locate aux head's current forward call site in collector**

Run: `grep -nE "aux_next_bar_forward|aux_regime_forward|launch_aux_forward" crates/ml/src/cuda_pipeline/gpu_experience_collector.rs`

Save line numbers.

- [ ] **Step 2: Locate encoder output buffer**

The encoder output is the input to GRN trunk. Find it via: `grep -nE "encoder_out|h_s1|grn_input|grn_trunk_forward" crates/ml/src/cuda_pipeline/gpu_experience_collector.rs | head -20`

Identify the buffer that GRN reads as input. That's also aux trunk's input.

- [ ] **Step 3: Allocate `h_s2_aux` buffer** (rollout-sized, mirrors aux's existing forward buffers in collector)

Find where `aux_next_bar_logits` or similar is allocated; allocate `h_s2_aux` same way + scratch buffers for `h_aux1`, `h_aux2` saved-for-backward.

- [ ] **Step 4: Insert aux trunk forward launch** AFTER encoder forward, BEFORE aux head forward.

Pseudocode order in the collector forward chain:
```
1. encoder_forward → encoder_out
2. grn_trunk_forward(encoder_out) → h_s2          // Q's path, unchanged
3. aux_trunk_forward(encoder_out) → h_s2_aux      // NEW
4. dir_concat_qaux(h_s2, aux_softmax_diff) → x_concat  // Coupling A unchanged, but reads aux's softmax which now comes from h_s2_aux
5. q_head_forward(x_concat) → q_logits
6. aux_head_forward(h_s2_aux) → aux_logits        // INPUT CHANGED from h_s2 to h_s2_aux
```

- [ ] **Step 5: Update aux head input plumbing**

In `aux_heads_kernel.cu`, the `aux_next_bar_forward` and `aux_regime_forward` kernels currently take `const float* h_s2` as input. Rename parameter to `h_s2_aux` (semantic rename — interface contract unchanged at call site since both buffers have same shape `[B, AUX_HIDDEN_DIM]`).

In `gpu_aux_heads.rs` Rust wrapper, update the launch call to pass `h_s2_aux` buffer ptr instead of `h_s2`.

- [ ] **Step 6: Wire backward chain**

After aux loss reduce produces `dh_s2_aux` (currently it produces `dh_s2`):
1. **OLD path (delete):** `dh_s2_aux` was zero-filled by today's stop-grad; no further propagation.
2. **NEW path:** `dh_s2_aux` → `aux_trunk_backward` → `dW1/db1/dW2/db2/dW3/db3` for aux trunk Adam update. NO propagation back to encoder (stop-grad).

Find aux backward call site in trainer (likely `gpu_dqn_trainer.rs`, search for `aux_next_bar_backward`). Insert `aux_trunk_backward` launch immediately after.

Also: the aux head's backward kernels (`aux_next_bar_backward`, `aux_regime_backward`) currently zero-fill `dh_s2_out`. With this commit, they should write the actual `dh_s2_aux` gradient (no zero-fill) — since the consumer is now aux_trunk_backward, not Q's trunk. **Revert the zero-fills landed today** (commits 872bd7392 and 411a30473) — they're no longer needed because aux's gradient now flows into aux_trunk, not h_s2.

- [ ] **Step 7: Add Adam update launch for aux trunk params**

Find existing Adam launch pattern (e.g., `apply_adam_update` for GRN params). Add aux trunk param-group with ISV-driven β1/β2/ε/LR (slots from C.1).

- [ ] **Step 8: Verify**

Run: `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --tests --all-targets 2>&1 | tail -15`

Expected: clean build. Build failure here means the contract migration is non-atomic — fix before proceeding.

- [ ] **Step 9: Run all oracle tests**

Run: `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test aux_trunk_oracle_tests --test sp14_oracle_tests --release -- --ignored --nocapture 2>&1 | tail -20`

Expected: all pre-existing tests still pass + 3 new aux_trunk tests pass.

- [ ] **Step 10: Commit**

```bash
git add crates/ml/src/cuda_pipeline/gpu_experience_collector.rs crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs crates/ml/src/cuda_pipeline/aux_heads_kernel.cu crates/ml/src/cuda_pipeline/gpu_aux_heads.rs
git commit -m "feat(sp14-c): atomic aux trunk wire-up — fwd+bwd+aux head input switch

Aux trunk inserted between encoder and aux head. Aux head input migrates
h_s2 → h_s2_aux atomically. Aux backward chain: aux CE → dh_s2_aux →
aux_trunk_backward → param Adam updates. Encoder gradient untouched
(stop-grad at trunk boundary). Today's stop-grad zero-fills in
aux_next_bar_backward + aux_regime_backward reverted — gradient now has
a valid consumer (aux_trunk_backward) instead of being structurally
discarded. Coupling A (dir_concat_qaux) unchanged — still feeds aux's
softmax-diff to Q's x_concat as forward feature."
```

---

## Phase C.6: h_s2_aux_rms_ema Producer Kernel

### Task C.6.1: Adaptive scale producer for aux trunk output magnitude

**Files:**
- Create: `crates/ml/src/cuda_pipeline/h_s2_aux_rms_ema_kernel.cu`
- Modify: `crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs` (add producer launch op)
- Modify: `crates/ml/src/cuda_pipeline/mod.rs`
- Modify: `crates/ml/build.rs`
- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (launch site)

**Goal:** mirror `h_s2_rms_ema` (existing slot 96 ISV producer for GRN trunk's output magnitude) but for aux trunk's output `h_s2_aux`. Used downstream for adaptive bound computation.

**Reference:** find existing `h_s2_rms_ema` kernel via `grep -rln "H_S2_RMS_EMA_INDEX\|h_s2_rms_ema" crates/ml/src/cuda_pipeline/`. Copy structure verbatim, retarget input buffer.

- [ ] **Step 1: Write kernel** (mirrors h_s2_rms_ema):
  - Input: `h_s2_aux [B, AUX_HIDDEN_DIM]`
  - Output: `ISV[H_S2_AUX_RMS_EMA_INDEX]` (single float)
  - Math: per-batch RMS → reduce-mean across batch → blend with EMA via Pearl-A bootstrap (sentinel-0 → batch_mean replace; subsequent: Welford EMA blend)

- [ ] **Step 2: Wrapper, cubin registration, mod decl** — standard.

- [ ] **Step 3: Wire into collector** AFTER `aux_trunk_forward`, BEFORE consumers (Adam launches that read this slot for adaptive bounds).

- [ ] **Step 4: Unit test**

```rust
#[test]
#[ignore]
fn h_s2_aux_rms_ema_pearl_a_bootstrap() {
    // 1. Set ISV[H_S2_AUX_RMS_EMA_INDEX] = 0.0 (sentinel)
    // 2. Launch with known h_s2_aux (controlled RMS)
    // 3. Assert ISV slot == known RMS (replace, not blend)
    // 4. Launch again with different RMS
    // 5. Assert ISV slot is Welford-EMA blend of the two
}
```

- [ ] **Step 5: Verify + commit**

```bash
git add crates/ml/src/cuda_pipeline/h_s2_aux_rms_ema_kernel.cu crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs crates/ml/src/cuda_pipeline/gpu_experience_collector.rs crates/ml/build.rs crates/ml/src/cuda_pipeline/mod.rs crates/ml/tests/aux_trunk_oracle_tests.rs
git commit -m "feat(sp14-c): h_s2_aux_rms_ema producer — adaptive scale for aux trunk

Mirrors GRN trunk's h_s2_rms_ema (slot 96) for aux trunk output.
Pearl-A first-observation bootstrap, Welford EMA blend thereafter.
Wired into collector chain after aux_trunk_forward."
```

---

## Phase C.7: ~~Atomic Deletion of α Machinery~~ — MERGED INTO C.1

α machinery deletion folded into Phase C.1 (no DEPRECATED stage; atomic with new ISV slot allocation per `feedback_no_legacy_aliases`). Phase C.7 task #380 should be marked completed once C.1 lands.

---

## Phase C.8: ISV-Driven Aux Trunk Hyperparameters

### Task C.8.1: Wire Adam β1/β2/ε/LR/grad-clip from ISV slots

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (Adam launch site for aux trunk)
- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (if Adam launches live there)

- [ ] **Step 1: Find existing per-group Adam pattern**

SP5 Pearl 4 wired per-group Adam β1/β2/ε from ISV. Reference: search for `adaptive_adam_beta1` or `ISV_BETA1_BASE` in production launches. Find the existing pattern.

- [ ] **Step 2: Wire aux trunk param-group**

Aux trunk's Adam launch reads:
- `β1` from `ISV[AUX_TRUNK_BETA1_INDEX]`
- `β2` from `ISV[AUX_TRUNK_BETA2_INDEX]`
- `ε` from `ISV[AUX_TRUNK_EPS_INDEX]`
- `lr` from `ISV[AUX_TRUNK_LR_INDEX]`
- `grad_clip_norm` from `ISV[AUX_TRUNK_GRAD_CLIP_INDEX]` applied before Adam step

Each is read from device-mapped ISV (no DtoH copy).

- [ ] **Step 3: Cold-start defaults**

Sentinel values from C.1 reset registry serve as bootstrapping defaults. Verify Pearl-A first-observation behavior: the producer that adapts these slots (if any — could be a controller TBD in a future phase) replaces sentinel on first valid observation.

**Note:** this plan does NOT include adaptive controllers driving these slots — only the static sentinel→Adam path. Adaptive driving is a follow-up (mirrors how SP5 Pearl 4 was extended over multiple sub-tasks).

- [ ] **Step 4: Verify**

Run: `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --lib 2>&1 | tail -10`

- [ ] **Step 5: Commit**

```bash
git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
git commit -m "feat(sp14-c): aux trunk Adam reads β1/β2/ε/LR/grad-clip from ISV

Per-group Adam pattern (mirrors SP5 Pearl 4). Sentinel defaults from
C.1 reset registry serve as bootstrap. Adaptive controllers driving
these slots deferred to follow-up phase."
```

---

## Phase C.9: Local Validation + Audit Doc + Memory Pearl

### Task C.9.1: Synthetic-data smoke test (RTX 3050 Ti)

**Files:**
- Modify: `crates/ml/tests/aux_trunk_oracle_tests.rs` (add smoke test)

- [ ] **Step 1: Synthetic up-trend smoke**

```rust
#[test]
#[ignore]
fn aux_trunk_learns_synthetic_uptrend() {
    // 1. Generate synthetic batch: prices monotone increasing
    // 2. Labels: all UP (label=1, K=2)
    // 3. Run 100 forward+backward+Adam steps with deterministic seed
    // 4. Assert: aux_loss < 0.1 by step 100 (ln(2)≈0.693 = random; <0.1 = clear signal of learning)
    // 5. Assert: aux_dir_acc > 0.95 on the same synthetic data
}
```

- [ ] **Step 2: Run + verify**

Run: `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test aux_trunk_oracle_tests --release -- --ignored aux_trunk_learns_synthetic_uptrend --nocapture 2>&1 | tail -20`

Expected: aux_loss < 0.1 by step 100, dir_acc > 0.95.

If the smoke test FAILS (aux can't learn synthetic up-trend), the architectural premise is wrong. Stop and investigate before L40S dispatch.

### Task C.9.2: Audit doc close-out + memory pearl

- [ ] **Step 1: Audit doc final entry**

Append to `docs/dqn-wire-up-audit.md`:
- Title: "SP14 Layer C — separate aux trunk landed"
- Phase summary (C.1 through C.8 commit SHAs)
- Decision rationale: stop-grad at encoder boundary, ISV as control-plane-only, α machinery deleted
- Open follow-ups: ISV-driven adaptive controllers for aux LR/β/grad-clip (deferred to next phase)

- [ ] **Step 2: Memory pearl candidate**

Create `/home/jgrusewski/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/pearl_separate_aux_trunk_when_shared_starves.md`:

```
---
name: separate aux trunk when shared trunk starves
description: When auxiliary head can't learn from a shared trunk despite gradient stop-grad, the cause is shared-trunk task incompatibility (Q-loss-shaped representation lacks aux-relevant features). Solution is separate aux trunk with stop-grad at encoder boundary, NOT more stop-grad / regularization tweaks
type: pearl
---

When an auxiliary head's loss flatlines at random baseline (ln(2) for K=2 CE)
even after stop-gradient prevents the aux loss from corrupting the shared trunk:
the bug is structural — Q-loss alone shapes the trunk to support value
estimation, which doesn't surface aux's task-relevant features. The aux head
reads the trunk forward but cannot reshape it backward (stop-grad). Result:
aux is starving for input.

Fix is **separate aux trunk** parallel to Q's trunk, both reading shared
encoder output. Aux trunk receives gradient only from aux loss; encoder
remains Q-shaped via additional stop-grad at encoder boundary.

Identification signal: aux_loss converges to and STAYS at the random-prediction
baseline (ln(2) ≈ 0.693 for binary CE; ln(K) generally) for several epochs
despite hyperparameter tuning. Pure stop-grad, regularization changes, or
loss weight adjustments will not fix this — they're treating the symptom of a
structural decoupling.

Anti-pattern: invent ever-cleverer α coupling formulas (Schmitt gates,
variance-driven scaling) to bridge the gap. The gap is structural; bridging
it via gradient-flow tricks fails because the representation gap can't be
crossed by gradient scaling.

Canonical: SP14 Layer C (2026-05-07), separate aux trunk for next-bar K=2
prediction in DQN. Replaced EGF α machinery (Coupling B + Schmitt gate +
variance EMAs) with simple "aux gets its own representation pipeline."
Net delete ~300 LOC, net add ~750 LOC for trunk + tests.
```

Add index entry to `MEMORY.md` under "Architectural constraints":

```markdown
- [pearl_separate_aux_trunk_when_shared_starves.md](pearl_separate_aux_trunk_when_shared_starves.md) — aux head flatlining at random baseline despite stop-grad means structural representation mismatch; fix is separate trunk, not gradient-flow cleverness; canonical SP14 Layer C (2026-05-07)
```

- [ ] **Step 3: Commit**

```bash
git add docs/dqn-wire-up-audit.md
git commit -m "docs(sp14-c): close-out audit + memory pearl candidate

Audit doc entry summarizing C.1-C.8 phases, decision rationale, and
deferred follow-ups (adaptive controllers for aux trunk hyperparams).
Memory pearl candidate: 'separate aux trunk when shared trunk starves'
— diagnostic for aux-loss-flatlines-at-ln(2) class of bugs."
```

(Memory file lives outside repo; commit only the audit doc.)

---

## Phase C.10: L40S 30-Epoch Validation

### Task C.10.1: Push + dispatch + monitor

**Acceptance criteria:**
- aux `next_bar_mse` (CE) drops below **0.65** by epoch 5 (vs. ln(2)=0.693 random baseline)
- aux `next_bar_mse` continues dropping (not flat) through epoch 10
- val_Sharpe trajectory does NOT collapse to <30 by epoch 15 (current bqtnd/thtxx baseline shows Sharpe→40 by ep10 with Hold-collapse)
- val_dir_dist `hold` stays below 50% through epoch 15
- No NaN/Inf in any metric

- [ ] **Step 1: Push**

```bash
git push origin sp15-phase1-honest-numbers
```

- [ ] **Step 2: Verify HEAD on remote**

```bash
git ls-remote origin sp15-phase1-honest-numbers
```

Expected: SHA matches local HEAD.

- [ ] **Step 3: Dispatch**

```bash
./scripts/argo-train.sh dqn --branch sp15-phase1-honest-numbers --gpu-pool ci-training-l40s --epochs 30 --baseline --tag sp14-c-aux-trunk
```

- [ ] **Step 4: Monitor with tightened filter**

Arm Monitor on log stream filtering for: aux_loss trajectory, val_Sharpe, val_dir_dist hold%, train_active_frac, NaN/Inf, FATAL/panic, workflow Succeeded/Failed.

Kill-fast triggers (per `feedback_kill_runs_on_anomaly_quickly`):
- aux next_bar_mse > 0.69 at epoch 5 (no learning — separate trunk hypothesis falsified, escalate)
- val_Sharpe collapse to <20 by epoch 10 (model worse than baseline)
- NaN/Inf anywhere
- CUDA_ERROR_ILLEGAL_ADDRESS or panic

- [ ] **Step 5: On success — close out**

If aux_loss < 0.65 by ep5 and trajectory is healthy through ep30:
- Mark memory pearl as canonical (move from "candidate" status)
- Promote SP14 Layer C task in MEMORY.md project state
- Begin scoping ISV-driven adaptive controllers for aux trunk hyperparams (deferred phase)

- [ ] **Step 6: On failure — escalate**

If aux_loss STILL flatlines at ln(2) with separate trunk:
- The deeper hypothesis (next-bar prediction is intrinsically too hard at HFT scale) is confirmed
- Options: (a) drop aux entirely, replace EGF intent with q_disagreement-only signal; (b) keep aux as feature-only via Coupling A but accept it's untrainable; (c) change aux's task (next-bar → multi-bar return, regime → trend phase)
- Brainstorm separately; not in this plan

---

## Self-Review Checklist (run after writing complete)

- [ ] **Spec coverage:** every locked decision has a phase that implements it
- [ ] **No placeholders:** zero `TODO`, `tbd`, `// fill in later`, `(implementation per existing pattern)` without a specific file:line reference
- [ ] **Type consistency:** `AUX_HIDDEN_DIM` referenced uniformly; `h_s2_aux` buffer name consistent across phases; `CudaFunction` pre-load discipline consistent
- [ ] **Atomic refactor rule:** Phase C.5 wire-up and Phase C.7 deletion are each ONE commit (not split)
- [ ] **Hard rules met:** no atomicAdd in any new kernel; ISV slot consumers never read pre-bootstrap sentinels without Pearl-A first-observation handling; no host branches in graph-captured paths
- [ ] **Test coverage:** forward oracle, backward gradient check, stop-grad invariant, Pearl-A bootstrap test, synthetic learning smoke
- [ ] **Validation gate:** acceptance criteria in C.10 are measurable and falsifiable

---

## Execution Handoff

**Recommended execution path:** subagent-driven-development.

Reasoning: 10 phases, each phase is one commit (mostly mechanical implementation given the locked design), independent enough that reviewer subagents can verify per-phase. Subagent-driven beats inline because:
- Each phase requires reading several files to find existing patterns — subagents can do that exploration without polluting controller context
- Spec-compliance review catches scope creep (e.g., adding adaptive controllers in C.8 when they're explicitly deferred)
- Code-quality review catches missing CudaFunction pre-loads, atomicAdd violations, missing reset_registry entries

**REQUIRED SUB-SKILL:** Use superpowers:subagent-driven-development to execute.
