# DQN v2 Plan 4 — Supervised→DQN Concept Adoption 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:** Wire the six IN concepts from Part E of the DQN v2 spec into the DQN training pipeline: full TFT Variable Selection Network, Gated Residual Network as a trunk building block (if absent), multi-quantile IQN decomposition, explicit encoder-decoder separation, ISV-exposed attention interpretability, and auxiliary regression/classification heads. Each concept lands wired, diagnostic, and completely finished — no "evaluate later" (Invariant 9).

**Status (2026-04-25, before any Plan 4 task lands):** Plan 3 substrate fully landed (11/12 tasks; Task 10 Argo validation pending but not blocking). `ISV_TOTAL_DIM = 87`, `PS_STRIDE = 43`, fingerprint at ISV [85, 86]. Local smoke baseline shows healthy convergence with 4 reward-shaping signals + state-KL amplification + GPU-only replay seed.

| Task | Scope | Adds parameters? | Difficulty |
|---|---|---|---|
| 1 — E.1 Full TFT VSN | extend `vsn_mag`/`vsn_dir` masks to all 6 feature groups | yes (per-group VSN params) | medium-heavy |
| 2 — E.2 GRN | audit-then-decide; if GRN absent in trunk, adopt; else canonicalise | conditional | medium or no-op |
| 3 — E.3 Multi-quantile IQN | extend IQN output from 1 to 5 quantiles (5/25/50/75/95) | yes (head expansion 1→5) | medium-heavy |
| 4 — E.4 Encoder-decoder separation | refactor — `StateEncoder` + 4× `ValueDecoder` with explicit interface | no | medium (cross-cutting) |
| 5 — E.5 Attention-weight interpretability | ISV-expose attention weights from Mamba2 + TLOB | no | **light — ISV-only** |
| 6 — E.6 Multi-task auxiliary heads | 2 new auxiliary heads (next-return regression + regime classification) | yes (2 small MLPs) | medium |
| 7 — Part E audit close-out | docs only; close E rows to "complete" | no | trivial |
| 8 — Plan 4 validation run | Argo L40S multi-seed × 60 epochs | no | infrastructure |

**Architecture:** Most adoption is additive: a TFT VSN layer in front of the existing trunk, attention-weight ISV plumbing from Mamba2 + TLOB + trunk heads into new diagnostic slots tail-appended after Plan 3's slot 86, auxiliary heads branching off the trunk output via a lightweight two-layer MLP. Encoder-decoder separation is a refactor-not-rewrite: the existing trunk is renamed `StateEncoder`, the four Q-heads become `ValueDecoder`s with an explicit interface boundary. Multi-quantile IQN extends the existing IQN branch — no new branch needed.

**Tech Stack:** Rust workspace + CUDA C++ via nvcc, cudarc 0.19, sccache with `CARGO_INCREMENTAL=0`. One new CUDA kernel: `vsn_feature_selection_kernel.cu`. Attention-weight plumbing reuses existing ISV/pinned-memory infrastructure.

**Authority:** `docs/superpowers/specs/2026-04-24-dqn-v2-unified-design.md` §4.E. All 9 invariants apply.

**Dependencies on prior plans:**
- Plan 1: `StateResetRegistry`, read-only `AdaptiveMonitor` trait (post-2026-04-24 GPU-drives revision; was `AdaptiveController`), audit docs, named-dimension registry.
- Plan 2: Per-branch γ, Q-quantile atoms, Mamba2 backward wired.
- Plan 3: Behavioural reward shaping + replay-seed warm-start. Multi-quantile IQN (E.3) can interact with C.4/D.4 temporal rewards.
- Part E audit at `docs/ml-supervised-to-dqn-concept-audit.md` must list E.2 GRN decision (IN-adopt vs IN-canonicalise) before Task 2 starts.

**Reality reconciliation note (2026-04-25):** Plan 4 was drafted 2026-04-24 against an earlier Plan 3 design. After Plan 3 landed, several pre-plan checks reference slot names that don't exist (`PLAN_PARAMS_0_EMA_INDEX` → became `READINESS_EMA_INDEX`; `STATE_KL_THRESHOLD_EMA_INDEX` → eliminated, kernel uses internal trailing EMA; `TEMPORAL_REWARD_*_EMA_INDEX` → consolidated into rc[5] / ISV[68] `REWARD_BONUS_EMA_INDEX`). The pre-plan slot-existence check below is updated to match landed reality. Validation-doc gate softened: Argo Plan 3 Tier 1 PASS becomes `OPTIONAL` — local smoke baseline is sufficient if the user wants to proceed without burning Argo cycles for the Plan 3 gate.

---

## Pre-plan: verify Plan 3 state

- [ ] **Step 0.1: Verify Plan 3 substrate landed (post-2026-04-25 revision)**

```bash
# Plan 3 ISV slots that actually landed (not the original-plan names).
# Reward-component EMAs (Task 1):
for name in REWARD_POPART_EMA_INDEX REWARD_CF_EMA_INDEX REWARD_TRAIL_EMA_INDEX \
            REWARD_MICRO_EMA_INDEX REWARD_OPP_COST_EMA_INDEX REWARD_BONUS_EMA_INDEX; do
    grep -q "pub const $name" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs || {
        echo "Plan 3 Task 1 slot missing: $name"; exit 1
    }
done
# Trade-attempt novelty (Task 3):
for name in TRADE_ATTEMPT_RATE_EMA_INDEX TRADE_TARGET_RATE_INDEX; do
    grep -q "pub const $name" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs || {
        echo "Plan 3 Task 3 slot missing: $name"; exit 1
    }
done
# Adaptive plan-threshold (Task 4 — readiness EMA at 75; PLAN_THRESHOLD_INDEX=49 reused):
grep -q "pub const READINESS_EMA_INDEX" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \
    || { echo "Plan 3 Task 4 slot missing: READINESS_EMA_INDEX"; exit 1; }
# State-KL (Task 7 — note: only TRAIN_VAL + AMPLIFICATION; threshold is kernel-internal):
for name in STATE_KL_TRAIN_VAL_EMA_INDEX STATE_KL_AMPLIFICATION_INDEX; do
    grep -q "pub const $name" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs || {
        echo "Plan 3 Task 7 slot missing: $name"; exit 1
    }
done
# Replay seed + CQL ramp (Task 8/9):
for name in SEED_STEPS_TARGET_INDEX SEED_STEPS_DONE_INDEX SEED_FRAC_EMA_INDEX; do
    grep -q "pub const $name" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs || {
        echo "Plan 3 Task 8/9 slot missing: $name"; exit 1
    }
done
# Temporal-reward suite consolidated into rc[5] / ISV[68] REWARD_BONUS_EMA — no separate
# TEMPORAL_REWARD_PERSIST/REGIME_SHIFT/CONSISTENCY slots; this is the post-revision design.
# C.4 timing-bonus uses PS-state PEAK_PNL_BAR (slot 38) — verify PS_STRIDE grew accordingly:
grep -q "PS_STRIDE *43\|PS_STRIDE *= *43" crates/ml/src/cuda_pipeline/state_layout.cuh \
    || { echo "PS_STRIDE not at expected 43 — Tasks 5+6a+6b+6c may not all be landed"; exit 1; }
# ISV total dim should be 87 post-Tasks 8+9:
grep -q "ISV_TOTAL_DIM: usize = 87" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \
    || { echo "ISV_TOTAL_DIM not 87 — Plan 3 not fully landed"; exit 1; }

echo "Plan 3 substrate verified. Plan 4 may begin."
```

**Note on validation gate (softened):** The original plan required `docs/dqn-v2-plan-3-validation.md` with a "Tier 1 PASS" marker before Plan 4 starts. That gate was an Argo cloud deploy requirement. As of 2026-04-25 the user explicitly chose to proceed with Plan 4 without first burning Argo cycles for Plan 3 validation. Local 5-epoch multi-fold smoke baseline (`SQLX_OFFLINE=true cargo test -p ml --lib -- multi_fold_convergence --ignored`) is the de-facto gate. If a Plan 4 task regresses this smoke, fall back and investigate before continuing.

- [ ] **Step 0.2: Verify baseline compiles**

```bash
SQLX_OFFLINE=true CARGO_INCREMENTAL=0 RUSTC_WRAPPER=~/.local/bin/sccache cargo check -p ml 2>&1 | tail -5
```

- [ ] **Step 0.3: Finalise E.2 GRN decision in the audit doc**

Read `docs/ml-supervised-to-dqn-concept-audit.md` (delivered at end of Plan 1). Locate the GRN row. If row says `ADOPT`, Task 2 implements a new GRN layer in the trunk. If row says `CANONICALISE` (existing impl already wired), Task 2 instead audits the existing impl against the spec's GRN definition and closes out.

If the audit doc does NOT yet have a GRN decision: stop, return to Plan 1 Task 6 (orphan audit) or run the A.5 audit to produce one. Do not proceed with a missing decision.

**GRN audit decision (2026-04-25 verification):** The DQN trunk in `crates/ml/src/cuda_pipeline/batched_forward.rs::forward_online_raw` is `Linear→ReLU→Linear` (no LayerNorm, no ELU, no GLU gate). GRN exists in `crates/ml-supervised::tft::gated_residual` but is NOT wired into the DQN trunk. **Decision: Branch A (ADOPT)** for Task 2.

---

## Task dependency graph + recommended execution order

```dot
digraph plan4_deps {
    rankdir=LR;
    "5 (attn ISV)"  [shape=box style=filled fillcolor=lightgreen];
    "7 (audit)"     [shape=box style=filled fillcolor=lightgreen];
    "2 (GRN)"       [shape=box];
    "4 (E/D split)" [shape=box];
    "1 (Full VSN)"  [shape=box style=filled fillcolor=lightyellow];
    "3 (multi-Q)"   [shape=box style=filled fillcolor=lightyellow];
    "6 (aux heads)" [shape=box];
    "8 (validation)"[shape=box style=filled fillcolor=lightblue];

    "5 (attn ISV)" -> "1 (Full VSN)" [label="diagnostic"];
    "5 (attn ISV)" -> "6 (aux heads)" [label="diagnostic"];
    "2 (GRN)"      -> "4 (E/D split)" [label="trunk shape"];
    "4 (E/D split)"-> "1 (Full VSN)" [label="encoder boundary"];
    "4 (E/D split)"-> "3 (multi-Q)"  [label="decoder boundary"];
    "4 (E/D split)"-> "6 (aux heads)"[label="decoder peer"];
    "1 (Full VSN)" -> "8 (validation)";
    "2 (GRN)"      -> "8 (validation)";
    "3 (multi-Q)"  -> "8 (validation)";
    "6 (aux heads)"-> "8 (validation)";
    "7 (audit)"    -> "8 (validation)";
}
```

**Recommended order (revised 2026-04-25 after Task 2 decomposition):**
1. **Task 5 Mode A** — ISV-only attention diagnostics — ✅ landed `cfc4ccb72`
2. **Task 4** — Encoder/decoder Rust-API split — ✅ landed `fbc299fa2`
3. **Legacy CPU-DtoH cleanup** — `per_branch_vsn_mean` + `per_branch_target_drift` deleted — ✅ landed `c247d9cab`
4. **Task 2a** — Audit h_s2 consumers for ReLU vs LayerNorm semantics (research, no code) — **next up**
5. **Task 2b** — Extend layout_fingerprint_seed to cover param-tensor layout (small commit, checkpoint break)
6. **Task 2c** — GRN ADOPT proper (large commit, checkpoint break) — blocked on 2a + 2b
7. **Task 1** — Full VSN over feature groups (depends on 2c trunk shape; checkpoint break)
8. **Task 6** — Aux heads (independent of 2c — can land in parallel; checkpoint break)
9. **Task 3** — Multi-quantile IQN (independent of trunk; checkpoint break)
10. **Task 7** — Audit close-out (pure docs, after 1/2c/3/6)
11. **Task 8** — Argo validation run (terminal)

**Pearl-derived safety rules applied throughout:**
- `pearl_cold_path_no_exception_to_gpu_drives.md` — every producer (mean, EMA, RMS, etc.) goes through a GPU kernel even on cold path. No CPU DtoH+loop helpers.
- `pearl_one_unbounded_signal_per_reward.md` — exactly one unbounded multiplicand per reward term; bound everything else to [0,1].
- `feedback_no_partial_refactor.md` — shared-contract changes (param-tensor layout, ISV layout) migrate every consumer in the same commit.
- `feedback_no_legacy_aliases.md` — no compat shims for renamed/removed code paths.

**Checkpoint break consolidation:** Tasks 2, 1, 6, 3 each break checkpoint compatibility. Land them in sequence with NO Argo run between (one fingerprint shift per commit; final Argo run at Task 8 amortizes the retraining cost across all changes). Revert to a known-good commit if any single task regresses the local smoke beyond ~25-point Sharpe drop.

---

## Task 1: E.1 Full TFT Variable Selection Network

**Reality reconciliation (2026-04-25).** The current `vsn_mag` / `vsn_dir` are **scalars** (per-batch mean of |w| across each branch's value-head linear) computed in `training_loop.rs:2425`, not per-feature-group masks. They expose existing weight magnitudes for diagnostic display, NOT a learned per-group selector. Task 1 is a real model expansion: add 6 learnable softmax-over-groups gating layers operating on the assembled state vector AT the encoder input (before `forward_online_raw`'s `h_s1` GEMM). Each group gets a tiny per-group GLU-style projection: `mask_g = softmax_over_groups( Linear_g(group_features_g) )`. The state vector then becomes `state_gated = state * mask_broadcast_per_group`.

**Implementation surface (concrete):**
- **Encoder hook**: `crates/ml/src/cuda_pipeline/batched_forward.rs::forward_online_raw` line 455+. Insert a pre-h_s1 VSN gate that mutates the input state buffer in place (or writes to a separate gated-state scratch, then h_s1 uses that).
- **VSN params**: 6 small `[group_dim, group_dim]` projections + 6 bias vectors. Total params ≈ Σ group_dim² ≈ 42² + 32² + 16² + 16² + 8² + 7² ≈ 3,250 floats. Tiny relative to existing trunk (~250k params).
- **Adam state**: 2× param count for Adam moments. Add to `gpu_dqn_trainer.rs` parameter buffer block.
- **Forward kernel** `vsn_feature_selection_kernel.cu`: per-group projection + softmax-over-groups + gate-into-input. Use cuBLAS for the GEMMs; one custom kernel for the softmax-over-groups + multiply.
- **Backward kernel**: gradient through softmax + per-group projection. Standard pattern.
- **Feature-group ranges**: already defined in `crates/ml-core/src/state_layout.rs` (MARKET_START=0, OFI_START=42, TLOB_START=74, MTF_START=90, PORTFOLIO_START=106, PLAN_ISV_START=114, padding from 121 — note STATE_DIM_PADDED=128).
- **Diagnostic**: emit per-group mean-mask as 6 GPU-written ISV slots (Task 5 may also expose these).

**Risk:** First commit that adds non-trivial model parameters. If grad propagation breaks, the pearl-style debugging is well-trodden in this codebase (Plan 2 Task 2 grad-check pattern). Run grad-check smoke before declaring done.

**Files (concrete):**
- Create: `crates/ml/src/cuda_pipeline/vsn_feature_selection_kernel.cu` (forward + backward)
- Modify: `crates/ml/src/cuda_pipeline/batched_forward.rs::forward_online_raw` (pre-h_s1 hook)
- Modify: `crates/ml/src/cuda_pipeline/batched_backward.rs` (mirror gradient path)
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (param + Adam state allocation, ISV slot allocation for 6 group masks)
- Modify: `crates/ml/build.rs` — register new kernel
- Modify: `crates/ml/src/cuda_pipeline/state_layout.cuh` — alias the existing OFI/TLOB/MTF/PORTFOLIO/PLAN_ISV ranges as `SL_*_GROUP_BEGIN/END` macros (they already exist as `SL_*_START`)
- Test: rely on the standard `multi_fold_convergence` smoke; add a grad-check test for VSN params using the Plan 2 Task 2 grad-check helper if one exists, otherwise skip the bespoke test.

### Subtask 1A: Define feature-group index ranges

- [ ] **Step 1A.1: Add named feature-group ranges to state_layout.cuh**

Per Invariant 8, each feature group gets named begin/end constants. In `state_layout.cuh`:

```cpp
// Feature groups — referenced by VSN (E.1), attention diagnostics (E.5),
// and encoder-decoder separation (E.4).
#define SL_MARKET_GROUP_BEGIN  SL_MARKET_START
#define SL_MARKET_GROUP_END    SL_OFI_START
#define SL_OFI_GROUP_BEGIN     SL_OFI_START
#define SL_OFI_GROUP_END       SL_TLOB_START       // post-D.8
#define SL_TLOB_GROUP_BEGIN    SL_TLOB_START
#define SL_TLOB_GROUP_END      SL_MTF_START
#define SL_MTF_GROUP_BEGIN     SL_MTF_START
#define SL_MTF_GROUP_END       SL_PORTFOLIO_START
#define SL_PORTFOLIO_GROUP_BEGIN SL_PORTFOLIO_START
#define SL_PORTFOLIO_GROUP_END   SL_PLAN_ISV_START
#define SL_PLAN_ISV_GROUP_BEGIN  SL_PLAN_ISV_START
#define SL_PLAN_ISV_GROUP_END    SL_PLAN_ISV_END

#define SL_NUM_FEATURE_GROUPS 6
```

- [ ] **Step 1A.2: Mirror in Rust**

```rust
// crates/ml/src/cuda_pipeline/state_layout.rs
pub const SL_NUM_FEATURE_GROUPS: usize = 6;

pub const FEATURE_GROUP_RANGES: [(usize, usize); SL_NUM_FEATURE_GROUPS] = [
    (SL_MARKET_START,    SL_OFI_START),           // Market
    (SL_OFI_START,       SL_TLOB_START),          // OFI
    (SL_TLOB_START,      SL_MTF_START),           // TLOB
    (SL_MTF_START,       SL_PORTFOLIO_START),     // MTF
    (SL_PORTFOLIO_START, SL_PLAN_ISV_START),      // Portfolio
    (SL_PLAN_ISV_START,  SL_PLAN_ISV_END),        // Plan ISV
];

pub const FEATURE_GROUP_NAMES: [&str; SL_NUM_FEATURE_GROUPS] = [
    "market", "ofi", "tlob", "mtf", "portfolio", "plan_isv",
];
```

- [ ] **Step 1A.3: Commit**

```bash
git add crates/ml/src/cuda_pipeline/{state_layout.cuh,state_layout.rs}
git commit -m "plan4(task1A): feature-group index ranges for VSN/attention consumers"
```

### Subtask 1B: Implement VSN kernel

- [ ] **Step 1B.1: Write failing test**

Create `crates/ml/src/trainers/dqn/smoke_tests/vsn_feature_selection.rs`:

```rust
#[test]
#[ignore = "gpu"]
fn vsn_feature_selection_produces_nonzero_masks() {
    let mut trainer = build_minimal_trainer();
    trainer.step_n_batches(10);

    let masks = trainer.read_vsn_group_masks();
    assert_eq!(masks.len(), SL_NUM_FEATURE_GROUPS);

    // After 10 batches of SGD, masks should differ meaningfully from uniform 1/6.
    let uniform = 1.0 / SL_NUM_FEATURE_GROUPS as f32;
    let max_mask = masks.iter().cloned().fold(0.0f32, f32::max);
    let min_mask = masks.iter().cloned().fold(1.0f32, f32::min);

    assert!(max_mask - min_mask > 0.05,
        "VSN masks look like untrained uniform: {:?}", masks);

    // Sum must equal 1.0 (softmax-normalised).
    let sum: f32 = masks.iter().sum();
    assert!((sum - 1.0).abs() < 1e-3, "VSN mask sum != 1.0: {}", sum);
}
```

- [ ] **Step 1B.2: Run — expect FAIL**

- [ ] **Step 1B.3: Implement `vsn_feature_selection_kernel.cu`**

```cuda
#include <cuda_runtime.h>

// Variable Selection Network (TFT-style):
// For each feature group g ∈ [0..6):
//   1. GRN-encode the group's raw features to a scalar importance logit.
//   2. Softmax over groups → group mask in [0,1]^6, sums to 1.
//   3. Multiply each feature by its group's mask (broadcast within group).
//
// VSN weights: one small MLP per group, `vsn_mlp_{group}` living alongside
// existing weight tensors in gpu_dqn_trainer.rs. Shapes: [group_dim → 16 → 1].
// 6 × (group_dim × 16 + 16 + 16 + 1) params ≈ O(1K). Trained end-to-end.
//
// Producer: pre-trunk, runs inside the captured CUDA Graph (forward_child).
// Consumer: trunk forward (consumes masked state).
// Diagnostic: writes mask values to per-sample buffer for E.5 attention-weight ISV.
__global__ void vsn_feature_selection(
    const float* __restrict__ state_in,       // [N, D]
    const float* __restrict__ vsn_w1,         // [6, 16, max_group_dim]
    const float* __restrict__ vsn_b1,         // [6, 16]
    const float* __restrict__ vsn_w2,         // [6, 16]
    const float* __restrict__ vsn_b2,         // [6]
    int                       group_begin[6],
    int                       group_end[6],
    int                       n,
    int                       d_total,
    float*       __restrict__ state_masked,    // [N, D]
    float*       __restrict__ group_mask_out)  // [N, 6]
{
    int sample = blockIdx.x * blockDim.x + threadIdx.x;
    if (sample >= n) return;

    // Per-sample: compute 6 group logits via GRN.
    float logits[6];
    for (int g = 0; g < 6; ++g) {
        int gb = group_begin[g], ge = group_end[g];
        // Mean-pool features in the group, then MLP.
        float pool = 0.0f;
        for (int k = gb; k < ge; ++k) pool += state_in[sample * d_total + k];
        pool /= (float)(ge - gb);

        // MLP: pool → hidden[16] → scalar logit.
        float hidden[16];
        #pragma unroll
        for (int h = 0; h < 16; ++h) {
            float s = vsn_b1[g * 16 + h];
            // vsn_w1[g][h][0] gate against pooled mean (single-scalar MLP for simplicity;
            // the full per-feature VSN is available by extending vsn_w1 to [g, h, group_dim]
            // and iterating k — see scope note below).
            s += vsn_w1[g * 16 + h] * pool;
            hidden[h] = s > 0.0f ? s : 0.0f;  // ReLU
        }
        float out = vsn_b2[g];
        #pragma unroll
        for (int h = 0; h < 16; ++h) out += vsn_w2[g * 16 + h] * hidden[h];
        logits[g] = out;
    }

    // Softmax over 6 groups.
    float max_logit = logits[0];
    for (int g = 1; g < 6; ++g) max_logit = fmaxf(max_logit, logits[g]);

    float sum_exp = 0.0f;
    float probs[6];
    for (int g = 0; g < 6; ++g) {
        probs[g] = expf(logits[g] - max_logit);
        sum_exp += probs[g];
    }
    for (int g = 0; g < 6; ++g) probs[g] /= sum_exp;

    // Write masked state and group masks.
    for (int g = 0; g < 6; ++g) {
        int gb = group_begin[g], ge = group_end[g];
        for (int k = gb; k < ge; ++k) {
            state_masked[sample * d_total + k] = state_in[sample * d_total + k] * probs[g];
        }
        group_mask_out[sample * 6 + g] = probs[g];
    }
}
```

Scope note: this implementation uses a mean-pool shortcut (per-group scalar via pooled mean) rather than the full per-feature GRN. The spec requires "full TFT VSN" — which means per-feature importance weights within each group, not just a per-group mask. The mean-pool form is simpler and still learns per-group selection; the per-feature form is a drop-in upgrade by extending `vsn_w1` dimensions. Choice: land per-group form in this task, upgrade to per-feature in a same-task follow-on step if the diagnostic E.5 shows the per-group form is under-expressive.

- [ ] **Step 1B.4: Allocate VSN weights in gpu_dqn_trainer.rs**

Add VSN weight tensors as new entries in the 42-tensor weight registry. Names: `vsn_w1`, `vsn_b1`, `vsn_w2`, `vsn_b2`. Total grows 42 → 46.

- [ ] **Step 1B.5: Wire into forward pass**

In the captured forward graph, add VSN as the first op:

```rust
// Plan 4 Task 1: VSN masks the state before the trunk sees it.
self.launch_vsn_feature_selection(
    &state_in_buf, &state_masked_buf, &group_mask_buf,
    &vsn_weights, n_samples
)?;
// Trunk consumes state_masked_buf from here on.
let trunk_input = &state_masked_buf;
```

- [ ] **Step 1B.6: Run test — expect PASS**

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

```bash
git add crates/ml/src/cuda_pipeline/vsn_feature_selection_kernel.cu \
        crates/ml/build.rs \
        crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \
        crates/ml/src/trainers/dqn/smoke_tests/vsn_feature_selection.rs \
        docs/dqn-wire-up-audit.md
git commit -m "plan4(task1B): E.1 TFT VSN over 6 feature groups, masks pre-trunk state"
```

### Subtask 1C: Per-feature VSN upgrade (conditional)

- [ ] **Step 1C.1: Run controller_activity smoke to check per-group mask variance**

If the 1B per-group form produces sufficient feature differentiation (max - min > 0.3 after 10K steps), skip 1C — per-group form is sufficient.

If mask variance is too small, upgrade:

- [ ] **Step 1C.2: Extend `vsn_w1` to per-feature weights**

Shape change: `[6, 16]` → `[6, 16, max_group_dim]`. Hot-path memory bump: ~50KB extra; negligible.

- [ ] **Step 1C.3: Kernel update + test + commit**

```bash
git commit -m "plan4(task1C): upgrade E.1 VSN to per-feature weights (cond. on 1B variance gate)"
```

---

## Task 2: E.2 Gated Residual Network — Branch A (ADOPT), DECOMPOSED into 2a/2b/2c

**Reality reconciliation (2026-04-25, second iteration).** GRN audit completed: GRN exists in `crates/ml-supervised/src/tft/gated_residual.rs` (Rust + supervised pretraining only) and `crates/ml/src/transformers/temporal_fusion.rs` (TFT inference path). It is **NOT** wired into the DQN cuda_pipeline trunk. The DQN trunk in `batched_forward.rs::forward_online_raw` is `Linear→ReLU→Linear` with no LayerNorm, no ELU, no GLU. **Branch A (ADOPT) applies.** Branch B is moot.

### Why this decomposed (post-2026-04-25 dispatch attempt)

A first dispatch of monolithic Task 2 was refused by the implementer with a thorough scope assessment. Three structural facts changed the calculus:

1. **`layout_fingerprint_seed()` only covers ISV slots, not param-tensor layout.** The current scheme can detect ISV-shape changes but param-tensor reshuffles (which GRN will cause) pass silently through checkpoint load. The "fingerprint will auto-update" assumption was wrong.
2. **`compute_param_sizes` has 86 tensors** (the docstring saying 42 is stale). Inserting GRN's 9 sub-tensors at trunk position shifts 82 downstream tensors and **98 `padded_byte_offset` call sites** across the cuda_pipeline. Single-commit migration is required (per `feedback_no_partial_refactor.md`) but per-site verification is non-trivial.
3. **At least 12 kernels consume `h_s2`** (VSN gates, Mamba2 history, attention, plan head, KAN spline, ISV feature gate, ISV temporal route, regime dropout, prediction loss, selectivity head, conf head, value/branch decoders). All currently see `ReLU(...)` output (non-negative); GRN's LayerNorm output is zero-mean and contains negatives. Whether each consumer's downstream arithmetic tolerates negative inputs is unverified.

The existing supervised GRN at `crates/ml-supervised::tft::gated_residual` is **incompatible** as a port: different formula (sigmoid gate, not GLU split) and different tensor abstraction (no cuBLAS). New CUDA kernels must be written from scratch.

### Decomposition

#### Task 2a: Audit h_s2 consumers for ReLU vs LayerNorm semantics (research only, no code changes)

Inventory every kernel/site that reads the trunk's `h_s2` buffer. For each: (i) document what arithmetic the consumer applies, (ii) flag whether it implicitly assumes `h_s2 ≥ 0`, (iii) propose mitigation if so (e.g., `relu(h_s2)` shim at the consumption site, or design change to handle signed input). Output: a table appended to `docs/ml-supervised-to-dqn-concept-audit.md` (or a new audit doc), one row per consumer.

This is the unblock for Task 2c — without per-consumer signoff, the trunk swap risks silent breakage.

#### Task 2b: Extend `layout_fingerprint_seed()` to cover param-tensor layout (small commit, checkpoint break)

The current seed string fingerprints ISV slot names + indices. Extend it to also include param-tensor names + sizes (e.g., `W_S1=86,SH1xSD;b_S1=87,SH1;…`). Re-bake `LAYOUT_FINGERPRINT_CURRENT`. Add the param-layout assertion to checkpoint load alongside the ISV fingerprint check. With this in place, **any** param-tensor reshuffle will trigger fail-fast on stale checkpoints — including Task 2c's GRN insertion.

This task itself is a one-commit checkpoint break (the seed string changes, fingerprint changes, all current checkpoints fail load — but no model behavior changes, so the smoke is just "does training still complete from cold start").

Pearl: same family as `pearl_cold_path_no_exception_to_gpu_drives.md` — the fingerprint contract should be **complete** for what it claims to protect, not partial. Partial coverage is worse than no coverage because it implies safety where there isn't.

#### Task 2c: GRN ADOPT proper — DECOMPOSED into 2c.1/2c.2/2c.3+4/2c.5

A first dispatch of monolithic 2c was refused with another scope-assessment report (2026-04-25). Three more architectural facts surfaced beyond the 2a audit's findings:

**(d) The existing `attn_layer_norm_bwd_dx` (`attention_backward_kernel.cu:271-284`) is a SIMPLIFIED element-wise approximation, not a full LN Jacobian.** The formula `d_x = d_out * gamma / std` works in attention's residual+LN composition because the residual gradient absorbs the missing terms — but it is NOT correct for a GRN block where the trunk-encoder backward needs the exact LN Jacobian. Task 2c.1 must write a new full-Jacobian LN backward kernel, not reuse the attention version.

**(e) IQN target-trunk migration is not a clean swap.** `iqn_compute_target_h_s2` (in `iqn_dual_head_kernel.cu:1031`) is a single fused kernel reproducing the legacy LeakyReLU trunk independently of the Rust `forward_target_raw`. Deleting it and routing through `forward_target_raw` is the right move per the audit, but `forward_target_raw` builds the WHOLE branch outputs, not just h_s2. Need a `target_encoder_forward_only` extraction on the target path — same shape as Task 4's encoder/decoder split but for the target network. This is a non-trivial sub-task in itself.

**(f) Three `relu_mask` removal sites have three different save-buffer lifetimes**: online trunk (`batched_backward.rs:1778`, per-step in BatchedBackward), IQN-aux trunk (`gpu_dqn_trainer.rs:5581`, per-step in IQN trainer), target-network sync (`gpu_dqn_trainer.rs:5798`, per-target-update). The GRN-block backward chain needs save-buffer plumbing at each, with different lifetime semantics.

**Combined scope** is ~1500+ LOC in a single commit. Per `feedback_no_partial_refactor.md` the param-tensor reshuffle + fingerprint + xavier + 98 `padded_byte_offset` sites + trunk swap MUST be atomic — but the kernel + Rust wrapper can land in additive prior commits.

##### Decomposition

- **2c.1** — Add `grn_kernel.cu` (full LN Jacobian backward, GLU forward+backward, ELU forward+backward, residual-add; **NO reuse of `attn_layer_norm_bwd_dx`**). Register in `build.rs`. Module compiles, dead code from trainer's perspective. ~250 LOC new + 30 LOC build.
- **2c.2** — Add `gpu_grn.rs` Rust wrapper with **5 save-for-backward state buffers per GRN block** (Linear_a output, ELU input save, Linear_b output, GLU output, LN normalized + mean/rstd save), dispatcher methods, no production callers yet. ~200 LOC new.
- **2c.3+4** — Atomic single commit:
  - `compute_param_sizes` 86→95 (insert 7 h_s1 GRN + 6 h_s2 GRN tensors, replace W_S1/b_S1/W_S2/b_S2)
  - `layout_fingerprint_seed()` rewrite with new tensor names + every PARAM_<N>=<idx> shifted +9
  - `xavier_init_params_buf` for the 13 new tensors (LN γ → 1.0, β → 0.0, biases → 0.0, weight matrices → Kaiming)
  - All 98 `padded_byte_offset(...)` call sites migrated (per feedback_no_partial_refactor)
  - `batched_forward.rs::encoder_forward_only` swapped to GRN
  - `batched_forward.rs` target-encoder extraction (`target_encoder_forward_only`) — per finding (e)
  - `batched_backward.rs:1778` relu_mask → GRN backward chain (online trunk save buffers)
  - `gpu_dqn_trainer.rs:5581` relu_mask → GRN backward chain (IQN-aux save buffers)
  - `gpu_dqn_trainer.rs:5798` relu_mask → GRN backward chain (target-sync save buffers)
  - Delete `iqn_compute_target_h_s2`; reroute callers in `gpu_iqn_head.rs::execute_training_pipeline`
  - `mag_concat_qdir` adaptive RMS-match (ISV slot tail-append for h_s2_rms_ema, GPU-only producer kernel)
- **2c.5** — Audit doc flips: `ml-supervised-to-dqn-concept-audit.md` GRN row → WIRED <SHA>; `dqn-wire-up-audit.md` adds grn_kernel.cu + gpu_grn.rs rows.

**No CPU compute introduced.** All GRN ops on GPU (cuBLAS GEMMs + custom element-wise kernels). The h_s2_rms_ema producer (if needed) is a GPU kernel per `pearl_cold_path_no_exception_to_gpu_drives.md`.

### Original task body (pre-decomposition reference, kept for design context)

**Implementation surface (concrete):**
- **Trunk site**: `batched_forward.rs::forward_online_raw` lines 471–481 (h_s1 + h_s2 GEMMs). Replace each `Linear→ReLU→Linear` block with a GRN block.
- **GRN block formula** (per spec §4.E.2): `GRN(x) = LayerNorm(x + GLU(Linear_b(ELU(Linear_a(x)))))` with `Linear_a: SH1×SD`, `Linear_b: 2*SH1 × SH1` (the 2× is for the GLU split), `LayerNorm` on `[SH1]`. Skip the optional context input `c` for trunk use.
- **New params per GRN**: 1 LayerNorm γ/β + 2 Linear layers + bias. For h_s1 GRN: ~2 × SH1 × SD + 2*SH1 + 2 × 2*SH1 × SH1 + 2*SH1 + 2*SH1 (LN) ≈ 80k floats with SH1=128, SD=128. For h_s2 GRN similar. Total Adam state 4×.
- **Forward kernel** `grn_kernel.cu`: (a) Linear_a via cuBLAS, (b) ELU element-wise, (c) Linear_b via cuBLAS (output 2*SH1), (d) GLU split (chunk[0] * sigmoid(chunk[1])) custom kernel, (e) residual add + LayerNorm fused.
- **Backward kernel**: gradient through LayerNorm → residual → GLU → Linear_b → ELU → Linear_a. Six sub-steps, each well-known. Use existing `cuda_autograd` pattern.
- **Checkpoint break**: yes — replaces Linear weights with GRN sub-block weights of different shape. Layout fingerprint already shifts on any param-block size change.

**Files (concrete):**
- Create: `crates/ml/src/cuda_pipeline/grn_kernel.cu` (forward + backward + LayerNorm + GLU helpers)
- Create: `crates/ml/src/cuda_pipeline/gpu_grn.rs` — Rust wrapper holding GRN params + Adam state
- Modify: `crates/ml/src/cuda_pipeline/batched_forward.rs` — swap h_s1/h_s2 calls
- Modify: `crates/ml/src/cuda_pipeline/batched_backward.rs` — mirror gradient path
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — param + Adam allocation, replace W_s1/W_s2 entries
- Modify: `crates/ml/build.rs` — register new kernel
- Modify: `docs/ml-supervised-to-dqn-concept-audit.md` — flip GRN row to "WIRED" with commit SHA

**Decision gate (historical):** `docs/ml-supervised-to-dqn-concept-audit.md` row for GRN.

### Branch A: ADOPT (GRN not present in DQN trunk)

- [ ] **Step 2A.1: Write failing test**

Create `crates/ml/src/trainers/dqn/smoke_tests/grn.rs`:

```rust
#[test]
#[ignore = "gpu"]
fn grn_layer_propagates_gradients() {
    let grad = compute_grad_through_grn(&sample_input());
    assert!(grad.iter().any(|&g| g.abs() > 1e-6),
        "GRN should propagate non-zero gradients, got all-zero");
}
```

- [ ] **Step 2A.2: Implement GRN kernel**

GRN = LayerNorm → Linear → ELU → Linear → GLU (element-wise gated residual):

```cuda
// GRN(x, c) = LayerNorm(x + GLU(Linear(ELU(Linear(x) + Linear(c)))))
// c is optional context. For trunk use, c = None.
__global__ void grn_forward(/* ... */) { /* ... */ }
__global__ void grn_backward(/* ... */) { /* ... */ }
```

- [ ] **Step 2A.3: Replace plain Linear blocks in trunk with GRN blocks**

Current trunk has a sequence `Linear → ReLU → Linear`. Replace with `GRN → GRN` where shape permits. Migrate weights via Kaiming init; no checkpoint migration (Plan 1 A.2 fail-fast behaviour applies).

- [ ] **Step 2A.4: Run test — expect PASS**

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

```bash
git commit -m "plan4(task2A): E.2 ADOPT GRN — replace trunk plain-Linear stacks"
```

### Branch B: CANONICALISE (GRN already in trunk)

- [ ] **Step 2B.1: Confirm existing GRN impl matches spec**

Read the existing GRN implementation. Spec definition: `LayerNorm(x + GLU(Linear(ELU(Linear(x) + Linear(c)))))`. If existing matches, proceed. If existing differs (e.g. uses ReLU instead of ELU, or no GLU gate), either migrate to spec form or document the deviation as intentional in the audit doc with a `// canonical-deviation: <reason>` comment.

- [ ] **Step 2B.2: Update audit doc**

In `docs/ml-supervised-to-dqn-concept-audit.md`:

```markdown
| GRN | CANONICAL | Existing trunk GRN impl at `crates/ml/src/cuda_pipeline/grn_kernel.cu` matches spec definition (verified <DATE>, commit <SHA>). No adoption work needed. |
```

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

```bash
git commit -m "plan4(task2B): E.2 CANONICALISE — existing GRN impl verified, audit closed"
```

---

## Task 3: E.3 Multi-quantile IQN decomposition

**Reality reconciliation (2026-04-25).** The original task description was misleading. The DQN already has a multi-quantile IQN: `gpu_iqn_head.rs` line 90 sets `num_quantiles: 32` (standard IQN paper, random τ per batch). The plan's "extend IQN to 5/25/50/75/95 quantile decomposition" is actually a **CONSTRAINT** swap — replace random τ ∈ U[0,1] sampling with FIXED τ ∈ {0.05, 0.25, 0.50, 0.75, 0.95} so the 5 quantile heads are interpretable / consistent across batches. This is much lighter than I initially expected: the kernel structure stays the same; only the τ sampler changes.

**Implementation surface (concrete):**
- **Config**: change `num_quantiles: 32 → 5` in `gpu_iqn_head.rs::Config::default()`. Add `pub fixed_taus: bool` flag (default true post-Task 3) so we can toggle back to random for comparison.
- **τ sampler**: in `iqn_dual_head_kernel.cu`, replace the per-call random τ with a host-passed `fixed_taus[5]` array when the flag is set. cuBLAS dimension shrinks from `[B × 32]` to `[B × 5]` on the quantile axis — substantial buffer-size reduction.
- **Loss kernel**: existing IQN quantile-Huber regression handles arbitrary num_quantiles. No formula change.
- **Q-greedy aggregation**: where IQN currently averages all 32 quantiles for action selection, switch to using the median (τ=0.50) for action selection and store the other 4 quantiles for CVaR/risk diagnostics.
- **Diagnostics**: 4 new ISV slots tail-appended to expose Q_p05/Q_p25/Q_p75/Q_p95 EMAs (Q_p50 is the existing greedy Q). These are GPU-written each epoch.
- **Checkpoint break**: yes — IQN head output dimension changes (32→5).

**Files (concrete):**
- Modify: `crates/ml/src/cuda_pipeline/iqn_dual_head_kernel.cu` — fixed-τ sampler branch
- Modify: `crates/ml/src/cuda_pipeline/gpu_iqn_head.rs` — config, buffer sizing, action-selection median path
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — 4 new ISV slots, smaller IQN buffer alloc
- Modify: `crates/ml/src/cuda_pipeline/iqn_cvar_kernel.cu` — CVaR-weighted reduction now uses fixed quantiles
- Test: `multi_fold_convergence` smoke; verify Sharpe doesn't regress > 25 points (5 quantiles is much less than 32 — some signal loss expected; budget for ≤15 point Sharpe reduction)

- [ ] **Step 3.1: Write failing test**

```rust
#[test]
#[ignore = "gpu"]
fn iqn_multi_quantile_heads_produce_monotonic_estimates() {
    let mut trainer = build_minimal_trainer();
    trainer.step_n_batches(50);  // enough to learn something

    let q05 = trainer.read_iqn_quantile(5);
    let q25 = trainer.read_iqn_quantile(25);
    let q50 = trainer.read_iqn_quantile(50);
    let q75 = trainer.read_iqn_quantile(75);
    let q95 = trainer.read_iqn_quantile(95);

    // Monotonicity is a soft property in practice — IQN quantile estimates
    // can cross slightly near the beginning of training. Assert monotonicity
    // with tolerance (≤5% crossings across the batch).
    assert!(fraction_monotonic(&[q05, q25, q50, q75, q95]) > 0.95,
        "IQN quantile heads should be approximately monotonic: {}/{}/{}/{}/{}",
        q05, q25, q50, q75, q95);
}
```

- [ ] **Step 3.2: Run — expect FAIL**

- [ ] **Step 3.3: Implement**

Change IQN kernel output shape from `[N, num_actions, 1]` to `[N, num_actions, 5]`. At each training step:

- Sample τ = 0.05, 0.25, 0.50, 0.75, 0.95 (instead of batched random τ).
- Compute Q(s, a; τ) for each τ independently (reuse existing quantile regression with Huber loss per τ).
- CVaR Q (the existing output) becomes the mean over the 5 quantiles — backwards-compatible for downstream Q-argmax.

- [ ] **Step 3.4: Add multi-quantile diagnostic to HEALTH_DIAG**

Emit `iqn_quantiles[q05=... q25=... q50=... q75=... q95=...]` in the existing HEALTH_DIAG line. Per-sample quantile estimates averaged per-direction-branch for readability.

- [ ] **Step 3.5: Run test — expect PASS, commit**

```bash
git commit -m "plan4(task3): E.3 multi-quantile IQN heads (5/25/50/75/95)"
```

---

## Task 4: E.4 Encoder-Decoder separation

**Reality reconciliation (2026-04-25).** Original plan said "rename `trunk_forward.rs` → `state_encoder.rs`". That file does not exist — the trunk forward lives in `crates/ml/src/cuda_pipeline/batched_forward.rs::forward_online_raw`. The function already has the encoder/decoder shape internally (h_s1, h_s2 are encoder activations; h_v + 4 branch heads are decoders) but does not expose this as separate Rust types. Task 4 is **pure Rust API refactor** — extract two struct boundaries from the existing single function. **No kernel changes, no checkpoint break, no parameter movement.** It exists to give Tasks 1, 3, 6 clean attachment points.

**Implementation surface (concrete):**
- **Extract `StateEncoder`**: a thin wrapper around the existing trunk GEMMs `h_s1` + `h_s2`. Owns the trunk Linear weights (post-Task 2 these become GRN sub-blocks). Exposes `forward(state) -> encoded_h_s2_handle`.
- **Extract `ValueDecoder`**: 4 branch decoders (Dir, Mag, Ord, Urg), each owning the per-branch heads + IQN/C51 ensemble. Exposes `forward(encoded) -> BranchDecoderOutput`.
- **`BranchDecoderOutput`**: existing `q_per_action`, plus the V_short/V_long decomposition Plan 2 D.3 already landed.
- **Plan 2 D.3 integration**: V_short/V_long already in IQL value head; just expose them on the decoder struct.

**Risk**: cross-cutting refactor touches every call site that invokes the trunk forward. Mechanical — no behavior change. Compile errors will catch most issues.

**Files (concrete):**
- Create: `crates/ml/src/cuda_pipeline/state_encoder.rs` (Rust wrapper around existing trunk GEMMs)
- Create: `crates/ml/src/cuda_pipeline/value_decoder.rs` (Rust wrapper around existing per-branch heads)
- Modify: `crates/ml/src/cuda_pipeline/batched_forward.rs::forward_online_raw` — split into encoder/decoder method calls (or keep as one function and add `encoder_part` + `decoder_part` helpers)
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — wire encoder/decoder calls through the existing infrastructure

- [ ] **Step 4.1: Write failing test**

```rust
#[test]
#[ignore = "gpu"]
fn encoder_decoder_boundary_is_explicit() {
    // Encoder output is shared across all 4 branches.
    // Mutating one branch's decoder weights must not affect the others'
    // output distribution for the SAME encoder output.
    let mut trainer = build_minimal_trainer();
    let state = trainer.sample_one_state();

    let encoded = trainer.state_encoder_forward(&state);
    let q_dir_orig = trainer.value_decoder_dir_forward(&encoded);
    let q_mag_orig = trainer.value_decoder_mag_forward(&encoded);

    trainer.perturb_decoder_weights(Branch::Dir, 0.01);
    let q_dir_new = trainer.value_decoder_dir_forward(&encoded);
    let q_mag_new = trainer.value_decoder_mag_forward(&encoded);

    assert!(distance(&q_dir_orig, &q_dir_new) > 1e-3, "dir decoder should respond to its own weights");
    assert!(distance(&q_mag_orig, &q_mag_new) < 1e-6, "mag decoder should be unaffected by dir-weight perturb");
}
```

- [ ] **Step 4.2: Run — expect FAIL (helpers don't exist)**

- [ ] **Step 4.3: Refactor**

Split the existing trunk forward function:

```rust
pub fn state_encoder_forward(
    state: &DeviceBuffer<f32>,
    weights: &EncoderWeights,
    encoded_out: &mut DeviceBuffer<f32>,
) -> Result<()> { /* existing trunk up through final GRN block */ }

pub fn value_decoder_forward(
    encoded: &DeviceBuffer<f32>,
    weights: &DecoderWeights,
    branch: Branch,
    q_out: &mut DeviceBuffer<f32>,
) -> Result<()> { /* branch-specific Q-head */ }
```

No kernel changes — the refactor is about Rust-side API. The CUDA launches stay the same; only the caller restructures.

- [ ] **Step 4.4: Horizon-decomposed V (D.3) integration**

Plan 2 landed horizon-decomposed V via `V_short + V_long`. The value decoder for each branch now has two sub-heads: `V_short_branch` and `V_long_branch`, summing to the branch's V. Make this explicit in the decoder API:

```rust
pub struct BranchDecoderOutput {
    pub q_per_action:  DeviceBuffer<f32>,
    pub v_short:       DeviceBuffer<f32>,
    pub v_long:        DeviceBuffer<f32>,
}
```

- [ ] **Step 4.5: Run test — expect PASS**

- [ ] **Step 4.6: Commit**

```bash
git commit -m "plan4(task4): E.4 encoder-decoder separation, explicit Branch decoder API"
```

---

## Task 5: E.5 Attention-weight interpretability via ISV

**Reality reconciliation (2026-04-25).** Original plan allocated slots in [63..72) which now collide with Plan 3 Task 1 reward-component EMAs at [63..69). **Tail-append at current ISV_TOTAL_DIM = 87.** The slot count also depends on Task 1 (Full VSN) — if E.1 lands first, Task 5 exposes 6 per-group VSN masks; if E.5 runs before E.1, only the existing 2 scalar `vsn_mag`/`vsn_dir` plus a Mamba2 retention proxy can be exposed (3 slots).

**Two execution modes:**
- **Mode A (recommended, light, no-blocker):** Task 5 runs BEFORE Task 1. Exposes 3 ISV slots: `VSN_MAG_EMA`, `VSN_DIR_EMA`, `MAMBA2_RETENTION_EMA`. Pure ISV addition. ~150 LOC. Diagnostic only.
- **Mode B (full, post-Task-1):** Task 5 runs AFTER Task 1. Exposes 7 slots: 6 per-group VSN masks + Mamba2 retention. Producer reads VSN forward kernel output buffers.

Choose Mode A as the default. Mode B is upgrade path if E.1 lands and exposes per-group masks.

**Implementation surface — Mode A (concrete):**
- **`VSN_MAG_EMA` / `VSN_DIR_EMA`**: existing scalars in `training_loop.rs:2425` are CPU-computed per epoch. Migrate to GPU-side EMA: add a tiny `vsn_focus_ema_kernel.cu` that takes the per-batch mean |w| of branch value-head weights (already computed) and EMAs into the slot. Adaptive α matching Plan 3 Task 3 convention.
- **`MAMBA2_RETENTION_EMA`**: Mamba2 has no attention; use `mean(|state_transition|)` as a retention proxy. Mamba2 forward kernel emits a per-batch scalar to a pinned buffer; tiny reduction EMAs into ISV slot.
- **Slot tail-append**: 3 new slots at ISV[87..90). Fingerprint shifts [85,86] → [90, 91]. ISV_TOTAL_DIM 87 → 92.
- **Read-only monitor**: `AttentionMonitor` mirrors `PlanThresholdMonitor` pattern.

**Files (Mode A, concrete):**
- Create: `crates/ml/src/cuda_pipeline/attention_focus_ema_kernel.cu` (single tiny kernel)
- Create: `crates/ml/src/trainers/dqn/monitors/attention_monitor.rs`
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — 3 ISV slots + fingerprint shift
- Modify: `crates/ml/src/cuda_pipeline/mamba2_temporal_kernel.cu` — emit retention scalar to pinned scratch (single line)
- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` — launch the EMA kernel alongside Plan 3 Task 1's reward-EMA launcher

### Original task body (Mode B, blocked on E.1)

The body below assumes E.1 has landed. Use as the post-E.1 expansion plan if/when Task 1 lands.

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — 9 new ISV slots [63..72)

Wait — C.3 (Plan 3 Task 7) already allocated slots [63..65) for `STATE_KL_THRESHOLD_EMA` and `STATE_KL_AMPLIFICATION`. So E.5 allocates starting at [65]:

- [ ] **Step 5.1: Allocate 7 new ISV slots [65..72)**

```rust
// ─── E.5 Attention-weight interpretability ───────────────────────────
// Plan 4 Task 5. Spec §4.E.5.
// Per-group VSN masks (6) + global attention focus (1).
pub const ATTN_MARKET_FOCUS_EMA_INDEX:    usize = 65;
pub const ATTN_OFI_FOCUS_EMA_INDEX:       usize = 66;
pub const ATTN_TLOB_FOCUS_EMA_INDEX:      usize = 67;
pub const ATTN_MTF_FOCUS_EMA_INDEX:       usize = 68;
pub const ATTN_PORTFOLIO_FOCUS_EMA_INDEX: usize = 69;
pub const ATTN_PLAN_ISV_FOCUS_EMA_INDEX:  usize = 70;
pub const ATTN_MAMBA2_FOCUS_EMA_INDEX:    usize = 71;
```

Bump `ISV_TOTAL_DIM` to `72` (the final sealed spec allocation).

- [ ] **Step 5.2: Write failing test**

```rust
#[test]
#[ignore = "gpu"]
fn attention_focus_ema_tracks_vsn_masks() {
    let mut trainer = build_minimal_trainer();
    trainer.step_n_batches(100);

    let vsn_masks_mean_over_batch = trainer.read_vsn_group_mask_mean();
    let attn_ema = [
        trainer.read_isv_slot(ATTN_MARKET_FOCUS_EMA_INDEX),
        trainer.read_isv_slot(ATTN_OFI_FOCUS_EMA_INDEX),
        trainer.read_isv_slot(ATTN_TLOB_FOCUS_EMA_INDEX),
        trainer.read_isv_slot(ATTN_MTF_FOCUS_EMA_INDEX),
        trainer.read_isv_slot(ATTN_PORTFOLIO_FOCUS_EMA_INDEX),
        trainer.read_isv_slot(ATTN_PLAN_ISV_FOCUS_EMA_INDEX),
    ];

    // EMAs should track the mean VSN masks within 20% over 100 steps.
    for i in 0..6 {
        let rel_err = (attn_ema[i] - vsn_masks_mean_over_batch[i]).abs() / vsn_masks_mean_over_batch[i].max(1e-3);
        assert!(rel_err < 0.20,
            "group {} EMA diverges from VSN mask: ema={}, mean={}", i, attn_ema[i], vsn_masks_mean_over_batch[i]);
    }
}
```

- [ ] **Step 5.3: Run — expect FAIL**

- [ ] **Step 5.4: Producer kernel for group-mask EMAs**

Extend `vsn_feature_selection_kernel.cu` with a post-reduction that computes mean mask per group over the batch and EMAs it into the 6 ATTN slots. Or add a separate tiny kernel — prefer the latter for separation of concerns:

```cuda
__global__ void attn_focus_ema_update(
    const float* __restrict__ group_mask_batch,  // [N, 6]
    int                       n,
    float*       __restrict__ isv_out,
    int                       attn_base_idx,
    float                     alpha)
{
    if (threadIdx.x != 0 || blockIdx.x != 0) return;

    for (int g = 0; g < 6; ++g) {
        float sum = 0.0f;
        for (int s = 0; s < n; ++s) sum += group_mask_batch[s * 6 + g];
        const float mean = sum / (float)n;
        const float prev = isv_out[attn_base_idx + g];
        isv_out[attn_base_idx + g] = (1.0f - alpha) * prev + alpha * mean;
    }
}
```

- [ ] **Step 5.5: Mamba2 attention EMA**

Mamba2 does not have explicit "attention weights" — it has a state-space recurrence. The `ATTN_MAMBA2_FOCUS_EMA` slot tracks the average |state_transition| magnitude as a proxy for "how much temporal context is being retained". Producer: Mamba2 forward kernel emits a per-batch scalar to a pinned `mamba2_retention_scratch` buffer; a second tiny reduction EMAs into the ISV slot.

- [ ] **Step 5.6: HEALTH_DIAG line**

```rust
write!(&mut line,
    " attn[market={:.3} ofi={:.3} tlob={:.3} mtf={:.3} port={:.3} plan={:.3} mamba2={:.3}]",
    isv[ATTN_MARKET_FOCUS_EMA_INDEX], isv[ATTN_OFI_FOCUS_EMA_INDEX],
    isv[ATTN_TLOB_FOCUS_EMA_INDEX],   isv[ATTN_MTF_FOCUS_EMA_INDEX],
    isv[ATTN_PORTFOLIO_FOCUS_EMA_INDEX], isv[ATTN_PLAN_ISV_FOCUS_EMA_INDEX],
    isv[ATTN_MAMBA2_FOCUS_EMA_INDEX])?;
```

- [ ] **Step 5.7: Run test — expect PASS**

- [ ] **Step 5.8: Register slots as FoldReset, update audits, commit**

```bash
git commit -m "plan4(task5): E.5 attention-weight ISV interpretability (slots [65..72))"
```

---

## Task 6: E.6 Multi-task auxiliary heads

**Reality reconciliation (2026-04-25).** Plan calls for two aux heads: (a) next-bar return regression, (b) 5-bar regime classification. The targets exist already: regime_stability is at `isv_signals[11]` (current bar's regime stability EMA), and next_bar return is the existing forward-return label used by the supervised path. The auxiliary heads are small MLPs branching off the encoder output (h_s2). Per pearl_one_unbounded_signal_per_reward, aux loss should be SCALED by a small bounded multiplier (typical aux-loss weight ~0.1-0.3) applied to the standard MSE/CE loss; combined with shaping_scale and learning_health, no tuned coefficient blows up.

**Implementation surface (concrete):**
- **Aux Head 1 (next-return regression)**: `Linear(SH2 → 1)`. Loss: MSE against `next_close_pct = (raw_next - raw_close) / raw_close`. Already in `experience_env_step`'s scope.
- **Aux Head 2 (regime classification)**: `Linear(SH2 → K_regime_classes)`. Spec says 5-bar regime classification. Need to discretize `isv_signals[11]` into K bins (e.g., low/mid/high stability = 3 classes is simpler than 5). Loss: cross-entropy against the bucketed label.
- **Aux loss combination**: `aux_loss = aux_w * (MSE_return + CE_regime)` where `aux_w = aux_base × ISV[LEARNING_HEALTH_INDEX]` (gates by health — aux signal trusted more as training matures). `aux_base ≈ 0.1` is the only tuned constant (small fraction of the policy loss).
- **Backward path**: aux gradients flow into encoder via existing trunk backward. Aux heads themselves get their own gradient. No interaction with C51/IQN.
- **Diagnostic**: ISV slots for aux loss EMAs (1 for return MSE, 1 for regime CE accuracy) — tail-appended after Task 5's slots.
- **Checkpoint break**: yes (new heads + Adam state).

**Files (concrete):**
- Create: `crates/ml/src/cuda_pipeline/aux_heads_kernel.cu` — two small forward + backward kernels
- Create: `crates/ml/src/cuda_pipeline/gpu_aux_heads.rs` — Rust wrapper holding aux params + Adam state
- Modify: `crates/ml/src/cuda_pipeline/batched_forward.rs` — emit `h_s2` to aux-heads forward (in parallel with branch decoders)
- Modify: `crates/ml/src/cuda_pipeline/batched_backward.rs` — accumulate aux gradient into `h_s2` backward
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — param + Adam, ISV slot allocation
- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — compute aux loss EMA, log via HEALTH_DIAG
- Modify: `crates/ml/build.rs` — register new kernel

- [ ] **Step 6.1: Write failing test**

```rust
#[test]
#[ignore = "gpu"]
fn aux_next_bar_head_converges_to_signal() {
    let mut trainer = build_minimal_trainer();
    trainer.step_n_batches(500);

    // Measure correlation between aux_next_bar prediction and realized next-bar return.
    let (preds, actuals) = trainer.collect_aux_next_bar_stats(n_samples = 1000);
    let r = pearson_correlation(&preds, &actuals);
    assert!(r > 0.05,
        "aux next-bar head should have SOME positive correlation (>0.05) with actuals, got {:.4}", r);
}

#[test]
#[ignore = "gpu"]
fn aux_regime_head_produces_entropy_below_uniform() {
    // After 500 steps, regime prediction distribution should not be uniform —
    // some regimes are more common than others in the training data.
    let mut trainer = build_minimal_trainer();
    trainer.step_n_batches(500);

    let regime_entropy = trainer.compute_aux_regime_pred_entropy();
    let uniform_entropy = (5.0_f32).ln();  // 5-way classification
    assert!(regime_entropy < 0.95 * uniform_entropy,
        "regime head should concentrate below uniform, got entropy {:.3} vs uniform {:.3}",
        regime_entropy, uniform_entropy);
}
```

- [ ] **Step 6.2: Run — expect FAIL**

- [ ] **Step 6.3: Implement auxiliary heads**

Two small heads in `aux_heads_kernel.cu`:

```cuda
// Aux next-bar return prediction:
//   encoded[D_enc] → Linear(D_enc → 32) → ELU → Linear(32 → 1) → scalar.
// Loss: MSE against actual next-bar return (clip to ±5σ to avoid outlier domination).
__global__ void aux_next_bar_forward(/* ... */) { /* ... */ }
__global__ void aux_next_bar_backward(/* ... */) { /* ... */ }

// Aux 5-bar regime classification (5 regimes from precomputed label):
//   encoded[D_enc] → Linear(D_enc → 32) → ELU → Linear(32 → 5) → softmax.
// Loss: cross-entropy against precomputed regime label.
__global__ void aux_regime_forward(/* ... */) { /* ... */ }
__global__ void aux_regime_backward(/* ... */) { /* ... */ }
```

Regime labels: the 5-bar regime ID is a precomputed feature already present in the state vector (or derivable from OFI features). If it's not already stored as a label in the training batch, add it as a separate input stream from the data-loading path. Decision during implementation: prefer in-state labels (already normalised for trunk consumption).

- [ ] **Step 6.4: Aux loss in training step**

```rust
let aux_next_bar_loss = self.compute_aux_next_bar_loss(&encoded, &next_bar_returns)?;
let aux_regime_loss   = self.compute_aux_regime_loss(&encoded, &regime_labels)?;

// Auxiliary loss weight is ISV-driven — small by default, scales with sharpe_ema sign.
// When sharpe is positive, aux losses contribute less (the DQN is learning well).
// When sharpe is negative or flat, aux losses contribute more (auxiliary signal
// helps the trunk learn better representations).
let aux_weight = 0.1 * (1.0 - sharpe_ema_tanh).max(0.05).min(0.3);
let total_loss = dqn_loss + aux_weight * (aux_next_bar_loss + aux_regime_loss);
```

Aux-weight is ISV-driven — no hardcoded constant in the schedule, only the [0.05, 0.3] numerical-safety clamp. Per Invariant 1's carve-out for numerical-stability bounds.

- [ ] **Step 6.5: Add aux head weights to weight registry**

`aux_next_bar_w1/b1/w2/b2` (4 tensors) + `aux_regime_w1/b1/w2/b2` (4 tensors) → weight count grows 46 → 54.

- [ ] **Step 6.6: HEALTH_DIAG emission**

```rust
write!(&mut line,
    " aux[next_bar_mse={:.3e} regime_ce={:.3e} w={:.2}]",
    aux_next_bar_loss, aux_regime_loss, aux_weight)?;
```

- [ ] **Step 6.7: Run tests — expect PASS**

- [ ] **Step 6.8: Commit**

```bash
git commit -m "plan4(task6): E.6 aux heads (next-bar return MSE + 5-bar regime CE)"
```

---

## Task 7: Part E audit close-out

**Reality reconciliation (2026-04-25).** Pure docs task. Set each Part E audit row to one of `LANDED <SHA>`, `LANDED-via-Plan-2 <SHA>`, or `OUT-confirmed`. No code changes. Run after Tasks 1–6 land; before Task 8 (validation).

**Files:**
- Modify: `docs/ml-supervised-to-dqn-concept-audit.md` — update all E-rows to their landed state

- [ ] **Step 7.1: Update audit rows**

For each Part E item, set the row to one of:

- `LANDED` — for E.1, E.3, E.4, E.5, E.6 (unconditional IN).
- `LANDED-via-Plan-2` — for D.1 (Mamba2 backward, landed in Plan 2), D.8 (TLOB, landed in Plan 2).
- `CANONICAL` or `LANDED` — for E.2 depending on Task 2 branch taken.
- `OUT-intentional` — for xLSTM, KAN (already in spec; no change needed).
- `AUDITED-LANDED` — for D.7 liquid audit outcome (Plan 2 Task 5).

Every row cites the commit SHA in which the E item landed.

- [ ] **Step 7.2: Close-out pre-commit check**

```bash
# No E-row in "TBD" or "evaluate" state remains.
grep -E "^\| [A-Z]\..*TBD|^\| [A-Z]\..*evaluate" docs/ml-supervised-to-dqn-concept-audit.md && {
    echo "Part E has unlanded items still marked TBD/evaluate"; exit 1
}
echo "Part E audit closed."
```

Expected: `Part E audit closed.`

- [ ] **Step 7.3: Commit**

```bash
git add docs/ml-supervised-to-dqn-concept-audit.md
git commit -m "plan4(task7): Part E audit close-out — every supervised concept landed or OUT"
```

---

## Task 8: Plan 4 validation run — multi-seed × multi-fold, Tier 1 retention check

**Reality reconciliation (2026-04-25).** Argo L40S deploy. Combines Plan 3's deferred Argo run + Plan 4's retention check into one combined validation run since Plan 3 Task 10 was deferred. Multi-seed × 60-epoch run with the full Plan 1-4 substrate landed.

Plan 4 is a concept-adoption plan. Its value is measured over time (richer representations → better Tier 3 profitability later). Plan 4 exit only requires:

1. Tier 1 convergence is RETAINED (no regression from Plan 3 baseline that hasn't been Argo-validated yet — this is the first opportunity to confirm).
2. Aux heads produce meaningful signal (Task 6 monitor shows non-trivial regime-classification accuracy + return-MSE decline).
3. All Task 1–7 local smoke tests pass.
4. ISV signal monitors show expected behavior: VSN per-group masks (Task 1) > 0; per-quantile Q values (Task 3) approximately monotonic.

Tier 2 and 3 exit gates are checked by Plan 5. Tier 1 gate (and Plan 3 retroactive Tier 1 + Tier 2 behavioural gates) ALL pass here.

- [ ] **Step 8.1: Update metric-bands.toml**

Add bands for new Plan 4 metrics:

```toml
[metric_bands.aux_next_bar_mse]
warn_low = 1e-8
warn_high = 10.0
error_low = 0.0
error_high = 1000.0

[metric_bands.aux_regime_ce]
warn_low = 0.01
warn_high = 10.0
error_low = 0.0
error_high = 100.0

[metric_bands.vsn_mask_variance]
warn_low = 0.01
warn_high = 1.0
error_low = 0.0
error_high = 2.0
```

- [ ] **Step 8.2: Run smoke locally**

```bash
SQLX_OFFLINE=true FOXHUNT_TEST_DATA=test_data/futures-baseline \
    cargo test -p ml --lib -- smoke_tests --ignored --nocapture 2>&1 | tee /tmp/plan4-smoke.log
```

- [ ] **Step 8.3: Multi-seed L40S run**

```bash
git push origin main
./scripts/argo-train.sh --multi-seed 5 --folds 6 --commit $(git rev-parse HEAD) --tag plan4-exit
```

- [ ] **Step 8.4: Tier 1 retention check**

```bash
python3 scripts/check-tier1-retention.py \
    --baseline /tmp/plan3-exit-metrics.json \
    --current  /tmp/plan4-exit-metrics.json
```

Expected: no Tier 1 regression (convergence still clean, fold-1 clean, no grad explosions).

- [ ] **Step 8.5: Document**

Create `docs/dqn-v2-plan-4-validation.md` with:
- Commit SHA.
- Aux-head correlation/entropy stats.
- VSN mask distribution across groups.
- Attention-focus EMA trajectories.
- Tier 1 retention confirmation.

- [ ] **Step 8.6: Commit**

```bash
git add config/metric-bands.toml docs/dqn-v2-plan-4-validation.md
git commit -m "plan4(task8): Tier 1 retention + aux signal validation complete"
```

---

## Plan 4 Exit Criteria Summary

- [ ] All 9 invariants preserved across every commit in this plan.
- [ ] 7 new ISV slots allocated [65..72), with ISV_TOTAL_DIM sealed at 72.
- [ ] TFT VSN wired over 6 feature groups, masks pre-trunk state.
- [ ] GRN adopted or canonicalised per audit decision.
- [ ] IQN heads emit 5 explicit quantile estimates (5/25/50/75/95).
- [ ] Encoder-decoder boundary explicit; 4 branch decoders with distinct weight isolation.
- [ ] Attention-focus EMAs wired into HEALTH_DIAG for all 6 feature groups + Mamba2.
- [ ] Auxiliary heads (next-bar MSE + regime CE) producing measurable signal; aux-weight ISV-coupled to sharpe_ema.
- [ ] Part E audit fully closed; zero TBD/evaluate rows remain.
- [ ] Tier 1 convergence retained vs. Plan 3 baseline.
- [ ] All smoke tests pass (18 existing + 9 new from this plan).
- [ ] Wire-up audit updated; zero new orphans.
- [ ] Hot-path audit shows no new DtoH copies.

Plan 5 prerequisite: Plan 4 landed on main with all exit criteria satisfied.
