# DQN v2 Plan 3 — Behavioural Core 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:** Land the behavioural machinery that escapes the val-Flat-trap: ISV-scaled continuous reward shaping, novelty-weighted trade-attempt bonus, adaptive plan-MLP activation threshold, per-component reward attribution, state-distribution divergence detection, temporal timing rewards, and a scripted-portfolio replay warm-start with CQL-α decoupling. Every change is ISV-driven (Invariant 1) and wires producers to consumers in the same commit (Invariant 2).

**Status (2026-04-25, 11/12 tasks landed; only Task 10 validation remaining):** `ISV_TOTAL_DIM = 87`, `PS_STRIDE = 43`. Fingerprint at ISV [85, 86].

| Task | Status | Commit | Key outcome |
|---|---|---|---|
| 1 — C.2 reward-component attribution | ✅ landed | `44539d8f4` | 6 ISV EMA slots [63..69); HEALTH_DIAG `reward_split` line |
| 2 — B.1 Flat opp-cost ISV-scaled | ✅ landed | `12bba98ec` | opp-cost × ISV[Q_DIR_ABS_REF]; self-scaling, no tuned multiplier |
| 3 — B.2 trade-attempt novelty bonus | ✅ landed | `b5d19c100` | ISV[71]=attempt rate, ISV[72]=frozen target; epoch-5 freeze with explicit gate |
| 4 — B.4 adaptive plan threshold | ✅ landed | `8f59c1e3b` | static→GPU producer for slot [49]; new ISV[75] readiness EMA |
| 5 — C.4 temporal timing bonus | ✅ landed | `a0abc3da3` | PS[38] PEAK_PNL_BAR; shaping×bars_early/hold×\|pnl\|×conviction on exit |
| 6a — D.4a persistence credit | ✅ landed | `777341776` | PS[39] MIN_PNL; tanh self-scaling, no tuned constants |
| 6b — D.4b regime-shift penalty | ✅ landed | `5ebebc564` | PS[40] REGIME_SHIFT_BAR; bounded-by-\|reward\| (post-fix per pearl) |
| 6c — D.4c conviction consistency | ✅ landed | `0bbe97ed8` | PS[41,42] Welford EMA; vol_proxy single-unbounded |
| 7 — C.3 state-KL divergence | ✅ landed | `673eb6612` | ISV[78,79] OFI-block KL + adaptive amp; val-side wired |
| 8 — B.3 replay seeded warm-start | ✅ landed | `3cb083f18` | GPU-only redesign — 4 scripted policies as one CUDA kernel |
| 9 — C.5 CQL α seed-coupled | ✅ landed | `3cb083f18` | GPU kernel ramps ISV[48] from 0 to config.cql_alpha as seed_frac→0 |
| 10 — Plan 3 validation run | ⏳ pending | — | Argo L40S deploy, multi-seed × 60 epochs |

**Architecture (post-landing accounting):** ISV tail sits at [73,74] fingerprint after slots [63..69) reward EMAs and [71,72] trade-attempt novelty. Portfolio state has one new slot [38]=PEAK_PNL_BAR (stride 39). The pending tasks add at most 6 more ISV slots (B.4 = 0 reuse, D.4 = 3, C.3 = 3, C.5 = up to 2). Conservative final `ISV_TOTAL_DIM ≤ 82`, fingerprint re-tails automatically.

**GPU-drives-CPU-reads architecture (spec §4.C.6 2026-04-24 revision):** every adaptive value in Plan 3 is computed by a GPU kernel reading ISV signals and written to an ISV slot. CPU-side `AdaptiveMonitor` impls are READ-ONLY observers. What the original plan called "controllers" — `PlanMlpThresholdController` (Task 4), `CqlAlphaSeedCoupledController` (Task 9) — become GPU kernel + read-only CPU monitor pairs, uniform with Plan 1's tau/epsilon/gamma/kelly_cap and Plan 2's per-branch-gamma pattern.

One new replay-buffer seed harness module (`dqn_replay_seed.rs`) driven by a portfolio of 4 scripted policies (Task 8). Replay-seed orchestration runs PRE-training (cold path, one-shot at constructor-time before epoch 1); does not violate GPU-drives.

**Tech Stack:** Rust workspace + CUDA C++ via nvcc, cudarc 0.19, sccache with `CARGO_INCREMENTAL=0`. No new crate dependencies. New CUDA kernel files: `reward_component_ema_kernel.cu` (✅ landed in Task 1), `trade_rate_ema_kernel.cu` (✅ landed in Task 3), `plan_threshold_update_kernel.cu` (Task 4), `temporal_reward_ema_kernel.cu` (Task 6), `state_kl_divergence_kernel.cu` (Task 7), `cql_alpha_seed_update_kernel.cu` (Task 9).

**Authority:** `docs/superpowers/specs/2026-04-24-dqn-v2-unified-design.md` §4.B, §4.C.2–C.5, §4.D.4. All 9 invariants apply.

**Dependencies on prior plans:**
- Plan 1: `StateResetRegistry`, `AdaptiveMonitor` trait (read-only observer, spec §4.C.6 2026-04-24 revision), `ISV_LAYOUT_FINGERPRINT_{LO,HI}_INDEX`, audit docs. `PLAN_THRESHOLD_INDEX = 49` already exists from Plan 1 Task 16 (static constructor write of 0.5; Task 4 upgrades the producer).
- Plan 2: layout-fingerprint auto-recompute, TLOB cuBLAS integration, per-branch γ, quantile atoms. Per-branch γ ISV slots at `[43..47)` populated; `Q_P05/Q_P95_EMA` at `[50..58)` populated; `PLAN_ISV_REMAINING_FRACTION` at `plan_isv[6]`.
- Plan 3 Tasks 1–3, 5: landed above. No backward-compat path; checkpoint retrain required (layout fingerprint changed structurally).

**Revision-to-reality (2026-04-24 second revision):** this plan's original text was written against ISV_TOTAL_DIM=58 with hard-coded slot numbers in the 49–65 range. After Tasks 1, 3 landed and the GPU-drives pivot, many of those numbers collide or reference retired patterns. Tasks 4, 6, 7, 8, 9, 10 bodies below have been rewritten to match landed reality. Tasks 1, 2, 3, 5 bodies are preserved as historical record of what shipped. See `session_2026-04-24_plan3_tasks_1_3_5.md` for landing narrative.

---

## Pre-plan: verify Plan 2 state

- [ ] **Step 0.1: Verify Plan 2 exit criteria satisfied**

Run:
```bash
# Verify ISV layout fingerprint seed covers Plan 2 slots
for name in Q_P05_DIR Q_P95_URG GAMMA_DIR GAMMA_URG; do
    grep -q "${name}=" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs || {
        echo "Plan 2 incomplete — layout fingerprint seed missing ${name}"; exit 1
    }
done

# Verify Plan 2 ISV slots exist
for name in GAMMA_DIR_INDEX GAMMA_MAG_INDEX GAMMA_ORD_INDEX GAMMA_URG_INDEX \
            Q_P05_DIR_INDEX Q_P05_MAG_INDEX Q_P05_ORD_INDEX Q_P05_URG_INDEX \
            Q_P95_DIR_INDEX Q_P95_MAG_INDEX Q_P95_ORD_INDEX Q_P95_URG_INDEX; do
    grep -q "pub const $name" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs || {
        echo "Plan 2 ISV slot missing: $name"; exit 1
    }
done

# Verify plan_isv[6] is wired
grep -q "PLAN_ISV_REMAINING_FRACTION" crates/ml/src/cuda_pipeline/state_layout.cuh || {
    echo "Plan 2 D.6 incomplete — PLAN_ISV_REMAINING_FRACTION missing"; exit 1
}

# Verify Mamba2 backward wired (D.1)
grep -q "mamba2_backward_fused" crates/ml/src/cuda_pipeline/mamba2_temporal_kernel.cu || {
    echo "Plan 2 D.1 incomplete — Mamba2 backward kernel missing"; exit 1
}

echo "Plan 2 exit criteria satisfied. Plan 3 may begin."
```

Expected: `Plan 2 exit criteria satisfied. Plan 3 may begin.`

- [ ] **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
```

Expected: `Finished` with zero new warnings or errors beyond baseline.

- [ ] **Step 0.3: Capture pre-Plan-3 convergence baseline**

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

Expected: smoke test passes on current main. Snapshot this log for Step 10 regression check.

---

## Task 1: C.2 Reward-component attribution — ✅ LANDED `44539d8f4`

**Rationale for first position:** Every subsequent behavioural task (B.1, B.2, C.4, D.4) adds a new reward term. Without C.2's per-component audit in place, debugging "did B.2 contribute meaningfully or is it silently zero?" requires log archaeology. C.2 is load-bearing diagnostic infrastructure for the rest of the plan.

> **Actual outcome.** 6 ISV slots at [63..69) (REWARD_POPART_EMA .. REWARD_BONUS_EMA). Producer: `reward_component_ema_kernel.cu` single-block, 6-thread, reads `reward_components_per_sample[N*L, 6]`, writes ISV[63..69) adaptive EMA (α=0.05). Launched from `training_loop.rs` after each `collect_experiences_gpu` epoch. Consumer: HEALTH_DIAG `reward_split` line + `RewardComponentMonitor` read-only observer. Fingerprint shifted [61,62] → [69,70]; ISV_TOTAL_DIM 63 → 71.

**Files:**
- Create: `crates/ml/src/cuda_pipeline/reward_component_audit_kernel.cu` — producer kernel (pinned-write of per-component |reward| EMA to ISV)
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — add 6 new ISV slot constants
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — wire producer into reward-computation flow
- Modify: `crates/ml/src/trainers/dqn/trainer/monitoring.rs` — add `reward_split[..]` to HEALTH_DIAG line
- Modify: `crates/ml/build.rs` — register new kernel
- Modify: `docs/isv-slots.md` — register new slots
- Modify: `docs/dqn-wire-up-audit.md` — register new module
- Test: extend `smoke_tests/reward_component_audit.rs`

### Subtask 1A: Allocate 6 new ISV slots for reward components

- [ ] **Step 1A.1: Add new ISV slot constants**

In `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`, after the Plan 2 Q-quantile block (last constant `Q_P95_URG_INDEX`), add:

```rust
// ─── C.2 Reward-component attribution ────────────────────────────────
// Plan 3 Task 1. Spec §4.C.2.
// Per-component |reward| contribution via adaptive-rate EMA.
// Producer: reward_component_audit_kernel (per-step, pinned write).
// Consumers: HEALTH_DIAG `reward_split[..]` line (monitoring.rs),
//            smoke test `reward_component_audit`.
pub const REWARD_POPART_EMA_INDEX:   usize = 52;
pub const REWARD_CF_EMA_INDEX:       usize = 53;
pub const REWARD_TRAIL_EMA_INDEX:    usize = 54;
pub const REWARD_MICRO_EMA_INDEX:    usize = 55;
pub const REWARD_OPP_COST_EMA_INDEX: usize = 56;
pub const REWARD_BONUS_EMA_INDEX:    usize = 57;
```

Current `ISV_TOTAL_DIM` post-Plan-2 is **63**. New slots append at 63 and grow from there; fingerprint re-tails to the new highest indices. Per-task slot allocation is documented at the start of each task; exact indices computed during implementation based on ISV_TOTAL_DIM at task start.

```rust
pub const ISV_TOTAL_DIM: usize = 58;  // Plan 3 Task 1 grows from <49..58>.
```

- [ ] **Step 1A.2: Compile check**

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

Expected: `Finished` with no new warnings.

- [ ] **Step 1A.3: Update `docs/isv-slots.md`**

Append the 6 new slot rows to the registry table. Each row: index, name, type (`f32 EMA`), producer module, consumer modules, category (`Scratchpad`).

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

```bash
git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs docs/isv-slots.md
git commit -m "plan3(task1A): allocate 6 ISV slots [52..58) for reward-component EMA"
```

### Subtask 1B: Implement reward_component_audit producer kernel

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

Edit `crates/ml/src/trainers/dqn/smoke_tests/reward_component_audit.rs` — extend existing test with:

```rust
#[test]
#[ignore = "gpu"]
fn reward_component_audit_emits_nonzero_for_active_terms() {
    let mut trainer = build_minimal_trainer();
    trainer.enable_dry_run_steps(5);

    let snapshot_before = trainer.read_isv_slots_52_to_58();
    trainer.step_with_nonzero_rewards(5);  // popart + cf + opp_cost active
    let snapshot_after = trainer.read_isv_slots_52_to_58();

    // Active terms must have non-zero EMA.
    assert!(snapshot_after[0] > 1e-6, "REWARD_POPART_EMA stuck at 0 after 5 active steps");
    assert!(snapshot_after[1] > 1e-6, "REWARD_CF_EMA stuck at 0 after 5 active steps");
    assert!(snapshot_after[4] > 1e-6, "REWARD_OPP_COST_EMA stuck at 0 after 5 active steps");

    // Inactive terms (trail, micro, bonus in this fixture) must stay ~0.
    // They're zero because the fixture only activates popart/cf/opp_cost.
    // NOT asserting == 0 (EMA can drift to 1e-20) but < 1e-8.
    assert!(snapshot_after[2] < 1e-8, "REWARD_TRAIL_EMA leaked non-zero");
}
```

- [ ] **Step 1B.2: Run the test to verify it fails**

```bash
SQLX_OFFLINE=true FOXHUNT_TEST_DATA=test_data/futures-baseline \
    cargo test -p ml --lib -- reward_component_audit_emits_nonzero --ignored --nocapture 2>&1 | tail -20
```

Expected: FAIL — function `read_isv_slots_52_to_58` not found OR test fails because producer kernel not yet wired.

- [ ] **Step 1B.3: Implement the producer kernel**

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

```cuda
#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include "isv_slots.cuh"  // pulls in ISV_TOTAL_DIM and slot indices; exists per A.2

// Per-component raw contributions for this step, batch-reduced to a single scalar each.
// Layout (host): [popart, cf, trail, micro, opp_cost, bonus].
// EMA rate adaptive-driven via isv[SHARPE_EMA_INDEX] — higher sharpe → higher alpha
// (faster response when training is healthy).
//
// Producer: gpu_dqn_trainer.rs::compute_reward_components_then_audit (cold-path,
// once per step, called AFTER reward kernel produces per-component scalars).
// ISV writes are via pinned + device-mapped pointer (Invariant 3, 4).
__global__ void reward_component_audit_reduce(
    const float* __restrict__ per_component_abs_sum,  // [6]
    const float                 batch_size_inv,        // 1/N
    const float* __restrict__ isv,                     // device ptr to pinned mirror
    float*       __restrict__ isv_out,                 // same backing memory
    const int                  reward_popart_ema_idx,
    const int                  reward_cf_ema_idx,
    const int                  reward_trail_ema_idx,
    const int                  reward_micro_ema_idx,
    const int                  reward_opp_cost_ema_idx,
    const int                  reward_bonus_ema_idx)
{
    const int lane = threadIdx.x;
    if (lane != 0) return;

    // Adaptive EMA rate: min 0.01, scaled by |sharpe_ema| clipped to [0, 2].
    const float sharpe = fmaxf(-2.0f, fminf(2.0f, isv[22]));  // SHARPE_EMA_INDEX
    const float alpha  = 0.01f + 0.04f * fabsf(sharpe);  // [0.01, 0.09]

    const int idx[6] = { reward_popart_ema_idx, reward_cf_ema_idx,
                         reward_trail_ema_idx, reward_micro_ema_idx,
                         reward_opp_cost_ema_idx, reward_bonus_ema_idx };

    for (int c = 0; c < 6; ++c) {
        const float new_sample = per_component_abs_sum[c] * batch_size_inv;
        const float prev       = isv[idx[c]];
        isv_out[idx[c]]        = (1.0f - alpha) * prev + alpha * new_sample;
    }
}
```

Note: no `atomicAdd` — per `feedback_no_atomicadd.md`. A single thread performs the reduction from a pre-reduced 6-element input (the caller is responsible for producing the per-component |sum| via a separate reduction kernel).

- [ ] **Step 1B.4: Register kernel in build.rs**

In `crates/ml/build.rs`, add to the list of CUDA source files:
```rust
"src/cuda_pipeline/reward_component_audit_kernel.cu",
```

- [ ] **Step 1B.5: Wire producer in gpu_dqn_trainer.rs**

In `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`, after the per-component reward reduction block in `compute_step_rewards` (or its equivalent, depending on current file structure), add:

```rust
// Plan 3 Task 1: audit per-component |reward| contributions.
// Launches the pre-reduce over batch followed by the ISV writer.
self.launch_reward_component_prereduce(
    &reward_components_buf,       // [N, 6]
    &reward_abs_sum_scratch,      // [6]
    n_samples,
)?;
self.launch_reward_component_audit_reduce(
    &reward_abs_sum_scratch,
    1.0f32 / (n_samples as f32),
)?;
```

Both helpers call existing `launch_kernel` plumbing. The pre-reduce is a separate one-block-per-component reduction; use existing tree-reduction pattern from `monitoring_kernel.cu`.

- [ ] **Step 1B.6: Run the test to verify it passes**

```bash
SQLX_OFFLINE=true FOXHUNT_TEST_DATA=test_data/futures-baseline \
    cargo test -p ml --lib -- reward_component_audit_emits_nonzero --ignored --nocapture 2>&1 | tail -20
```

Expected: PASS.

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

```bash
git add crates/ml/src/cuda_pipeline/reward_component_audit_kernel.cu \
        crates/ml/build.rs \
        crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \
        crates/ml/src/trainers/dqn/smoke_tests/reward_component_audit.rs
git commit -m "plan3(task1B): reward_component_audit producer kernel wired into hot path"
```

### Subtask 1C: Wire reward_split into HEALTH_DIAG

- [ ] **Step 1C.1: Write the failing line-format test**

Edit `crates/ml/src/trainers/dqn/smoke_tests/reward_component_audit.rs`:

```rust
#[test]
#[ignore = "gpu"]
fn health_diag_contains_reward_split() {
    let log = capture_health_diag_line(/*epochs=*/ 2);
    assert!(log.contains("reward_split["),
            "HEALTH_DIAG missing reward_split[..] after C.2 wiring");

    // Must contain all 6 named labels.
    for name in &["popart", "cf", "trail", "micro", "opp_cost", "bonus"] {
        assert!(log.contains(name),
                "reward_split missing component name '{}'", name);
    }
}
```

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

```bash
SQLX_OFFLINE=true FOXHUNT_TEST_DATA=test_data/futures-baseline \
    cargo test -p ml --lib -- health_diag_contains_reward_split --ignored --nocapture 2>&1 | tail -20
```

Expected: FAIL — no `reward_split[` in log.

- [ ] **Step 1C.3: Emit reward_split in monitoring.rs**

In `crates/ml/src/trainers/dqn/trainer/monitoring.rs`, near the existing `grad_split_bwd[` emit, add:

```rust
let rs = [
    isv[REWARD_POPART_EMA_INDEX],
    isv[REWARD_CF_EMA_INDEX],
    isv[REWARD_TRAIL_EMA_INDEX],
    isv[REWARD_MICRO_EMA_INDEX],
    isv[REWARD_OPP_COST_EMA_INDEX],
    isv[REWARD_BONUS_EMA_INDEX],
];
write!(
    &mut line,
    " reward_split[popart={:.3e} cf={:.3e} trail={:.3e} micro={:.3e} opp_cost={:.3e} bonus={:.3e}]",
    rs[0], rs[1], rs[2], rs[3], rs[4], rs[5]
)?;
```

- [ ] **Step 1C.4: Run test — expect PASS**

```bash
SQLX_OFFLINE=true FOXHUNT_TEST_DATA=test_data/futures-baseline \
    cargo test -p ml --lib -- health_diag_contains_reward_split --ignored --nocapture 2>&1 | tail -5
```

Expected: PASS.

- [ ] **Step 1C.5: Update wire-up audit**

Append row to `docs/dqn-wire-up-audit.md`:

```markdown
| `cuda_pipeline::reward_component_audit_kernel` | Producer for ISV slots [52..58). Consumers: `trainer::monitoring::emit_health_diag_line`, smoke `reward_component_audit`. | Wired | Plan 3 Task 1 |
```

- [ ] **Step 1C.6: Commit**

```bash
git add crates/ml/src/trainers/dqn/trainer/monitoring.rs \
        docs/dqn-wire-up-audit.md \
        crates/ml/src/trainers/dqn/smoke_tests/reward_component_audit.rs
git commit -m "plan3(task1C): HEALTH_DIAG reward_split[..] line wired, audit complete"
```

---

## Task 2: B.1 ISV-driven continuous Flat opp-cost — ✅ LANDED `12bba98ec`

> **Actual outcome.** `experience_env_step` Flat-branch opp-cost now multiplies by `ISV[Q_DIR_ABS_REF_INDEX=21]` (EMA of max |Q_mean| across direction bins). Self-scaling: as Q magnitudes drift, opp-cost tracks proportionally. Floor `1e-3` for cold-start. Writes to `rc[4]` so the opp-cost contribution lands in ISV[67] REWARD_OPP_COST_EMA via the Task 1 kernel. No tuned multiplier.

**Rationale:** The current opp-cost is conviction-driven (good) but the |Q| scale is still implicit. When Q magnitudes drift into ±50, the opp-cost must drift with them, otherwise it reverts to "invisible" and Flat becomes argmax-attractive again.

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/experience_kernels.cu` — scale Flat opp-cost by `isv[Q_DIR_ABS_REF_INDEX]`
- Modify: `crates/ml/src/trainers/dqn/smoke_tests/reward_component_audit.rs` — add opp-cost scaling assertion

Uses existing ISV slot `Q_DIR_ABS_REF_INDEX` — no new allocation.

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

Append to `reward_component_audit.rs`:

```rust
#[test]
#[ignore = "gpu"]
fn flat_opp_cost_scales_with_q_dir_abs_ref() {
    // Run twice: once with Q_DIR_ABS_REF=1.0, once with Q_DIR_ABS_REF=10.0.
    // Opp-cost EMA should scale ~10x.
    let mut trainer = build_minimal_trainer();

    trainer.set_isv_slot(Q_DIR_ABS_REF_INDEX, 1.0);
    trainer.step_with_flat_positions(100);
    let opp_cost_1x = trainer.read_isv_slots_52_to_58()[4]; // REWARD_OPP_COST_EMA

    trainer.reset_reward_emas();
    trainer.set_isv_slot(Q_DIR_ABS_REF_INDEX, 10.0);
    trainer.step_with_flat_positions(100);
    let opp_cost_10x = trainer.read_isv_slots_52_to_58()[4];

    let ratio = opp_cost_10x / opp_cost_1x.max(1e-12);
    assert!(ratio > 7.0 && ratio < 13.0,
            "opp_cost should scale ~10x with Q_DIR_ABS_REF; got {:.2}x", ratio);
}
```

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

```bash
SQLX_OFFLINE=true FOXHUNT_TEST_DATA=test_data/futures-baseline \
    cargo test -p ml --lib -- flat_opp_cost_scales --ignored --nocapture 2>&1 | tail -10
```

Expected: FAIL — ratio too small (current code doesn't scale by Q_DIR_ABS_REF).

- [ ] **Step 2.3: Wire Q_DIR_ABS_REF into opp-cost formula**

In `crates/ml/src/cuda_pipeline/experience_kernels.cu`, locate the Flat opp-cost line (currently `reward = -shaping_scale * holding_cost_rate * conviction_core * vol_proxy_flat`). Change to:

```cuda
// B.1: scale by |Q| signal so opp-cost tracks Q magnitudes as training progresses.
// isv[Q_DIR_ABS_REF_INDEX] is already an EMA of per-branch |Q|, populated elsewhere
// by update_eval_v_range (existing infra; no new producer needed).
const float q_abs_ref = fmaxf(1e-3f, isv[q_dir_abs_ref_idx]);
reward = -shaping_scale * holding_cost_rate * conviction_core * vol_proxy_flat * q_abs_ref;
```

Add the new kernel parameter `q_dir_abs_ref_idx` to the kernel signature, thread through at the launch site in `gpu_dqn_trainer.rs`.

- [ ] **Step 2.4: Run — expect PASS**

```bash
SQLX_OFFLINE=true FOXHUNT_TEST_DATA=test_data/futures-baseline \
    cargo test -p ml --lib -- flat_opp_cost_scales --ignored --nocapture 2>&1 | tail -5
```

Expected: PASS.

- [ ] **Step 2.5: Commit**

```bash
git add crates/ml/src/cuda_pipeline/experience_kernels.cu \
        crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \
        crates/ml/src/trainers/dqn/smoke_tests/reward_component_audit.rs
git commit -m "plan3(task2): B.1 Flat opp-cost scales with isv[Q_DIR_ABS_REF]"
```

---

## Task 3: B.2 ISV-driven trade-attempt bonus — ✅ LANDED `b5d19c100`

> **Actual outcome.** 2 ISV slots at tail: [71] TRADE_ATTEMPT_RATE_EMA (adaptive α from Sharpe), [72] TRADE_TARGET_RATE (CPU-frozen at epoch 5). Producer: `trade_rate_ema_kernel.cu` reduces `flat_to_pos_per_sample [N*L]` single-block (no atomicAdd, no DtoH). Consumer: Flat→Positioned site in `experience_env_step` computes `novelty = max(0, 1 - attempt/target)` and adds `conviction × vol_proxy × novelty × shaping_scale` to reward + rc[5]. **Explicit freeze gate** added during implementation: `if target_raw > 1e-6f` — bonus structurally inert pre-freeze (prevents epoch-1 saturation where attempt=0 would give novelty=1.0). Fingerprint shifted [69,70] → [73,74]; ISV_TOTAL_DIM 71 → 75.

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — 2 new ISV slot constants
- Modify: `crates/ml/src/cuda_pipeline/experience_kernels.cu` — compute Flat→Positioned novelty bonus
- Create: `crates/ml/src/cuda_pipeline/trade_rate_ema_kernel.cu` — producer for TRADE_ATTEMPT_RATE_EMA
- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — freeze TRADE_TARGET_RATE after warmup
- Test: new smoke `trade_attempt_bonus_novelty`

- [ ] **Step 3.1: Allocate slots 50 and 51**

In `gpu_dqn_trainer.rs`, after the Task 1 block:

```rust
// ─── B.2 ISV-driven trade-attempt bonus ──────────────────────────────
// Plan 3 Task 3. Spec §4.B.2.
pub const TRADE_ATTEMPT_RATE_EMA_INDEX: usize = 50;
pub const TRADE_TARGET_RATE_INDEX:      usize = 51;
```

Grow `ISV_TOTAL_DIM` by +N slots for whatever this task requires; re-tail the fingerprint. Implementation computes exact indices from ISV_TOTAL_DIM at task start.

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

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

```rust
#[test]
#[ignore = "gpu"]
fn trade_attempt_bonus_fires_when_under_trading() {
    let mut trainer = build_minimal_trainer();

    // Freeze TRADE_TARGET_RATE high (0.05 ≈ healthy baseline).
    trainer.set_isv_slot(TRADE_TARGET_RATE_INDEX, 0.05);

    // Simulate a Flat-collapsed run: TRADE_ATTEMPT_RATE_EMA converged to 0.001.
    trainer.set_isv_slot(TRADE_ATTEMPT_RATE_EMA_INDEX, 0.001);

    // One Flat→Positioned transition. Measure the bonus reward.
    let bonus = trainer.simulate_single_flat_to_positioned_reward();

    // novelty = max(0, 1 - 0.001/0.05) = max(0, 0.98) = 0.98.
    // bonus = conviction × vol_proxy × 0.98, so it should be > 0.
    assert!(bonus > 1e-4,
            "trade-attempt bonus should fire when under-trading, got {}", bonus);
}

#[test]
#[ignore = "gpu"]
fn trade_attempt_bonus_inert_at_target_rate() {
    let mut trainer = build_minimal_trainer();
    trainer.set_isv_slot(TRADE_TARGET_RATE_INDEX,      0.05);
    trainer.set_isv_slot(TRADE_ATTEMPT_RATE_EMA_INDEX, 0.05);  // at target

    let bonus = trainer.simulate_single_flat_to_positioned_reward();
    assert!(bonus.abs() < 1e-6,
            "bonus should be ~0 when at target rate, got {}", bonus);
}
```

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

Expected: compile fail (helpers don't exist) or runtime fail (slots unwritten).

- [ ] **Step 3.4: Implement producer kernel**

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

```cuda
#include <cuda_runtime.h>

// Updates TRADE_ATTEMPT_RATE_EMA from this-batch transitions.
// - flat_to_pos_count: host-reducted integer of Flat→Positioned transitions in batch.
// - batch_size: total samples (N).
// Rate = flat_to_pos_count / batch_size; EMA via adaptive α.
__global__ void trade_attempt_rate_ema_update(
    int         flat_to_pos_count,
    int         batch_size,
    float*      isv_out,  // pinned device-mapped
    int         trade_attempt_rate_ema_idx,
    int         sharpe_ema_idx,
    float       alpha_base)
{
    if (threadIdx.x != 0 || blockIdx.x != 0) return;

    const float rate   = (float)flat_to_pos_count / fmaxf(1.0f, (float)batch_size);
    const float sharpe = fmaxf(-2.0f, fminf(2.0f, isv_out[sharpe_ema_idx]));
    const float alpha  = alpha_base * (1.0f + 0.5f * fabsf(sharpe));  // adaptive
    const float prev   = isv_out[trade_attempt_rate_ema_idx];
    isv_out[trade_attempt_rate_ema_idx] = (1.0f - alpha) * prev + alpha * rate;
}
```

Register in `build.rs`.

- [ ] **Step 3.5: Freeze TRADE_TARGET_RATE after warmup**

In `crates/ml/src/trainers/dqn/trainer/training_loop.rs`, after the per-epoch loop, add:

```rust
const TRADE_TARGET_RATE_FREEZE_EPOCH: u32 = 5;  // documented in docs/isv-slots.md

if self.current_epoch == TRADE_TARGET_RATE_FREEZE_EPOCH {
    let measured_rate = self.read_isv_slot(TRADE_ATTEMPT_RATE_EMA_INDEX);
    // Guard: never freeze a degenerate zero rate (that would make novelty
    // permanently 1.0 and dominate the reward).
    let rate_for_target = measured_rate.max(0.001);
    self.write_isv_slot(TRADE_TARGET_RATE_INDEX, rate_for_target);
    tracing::info!(
        target: "dqn::B2",
        "TRADE_TARGET_RATE frozen at epoch {} to {:.4} (measured {:.4})",
        self.current_epoch, rate_for_target, measured_rate
    );
}
```

- [ ] **Step 3.6: Wire bonus into experience_kernels.cu**

At the Flat→Positioned transition detection:

```cuda
// B.2: novelty-weighted trade-attempt bonus.
if (prev_pos == 0 && new_pos != 0) {
    const float attempt_rate = isv[trade_attempt_rate_ema_idx];
    const float target_rate  = fmaxf(1e-4f, isv[trade_target_rate_idx]);
    const float novelty      = fmaxf(0.0f, 1.0f - attempt_rate / target_rate);
    const float bonus        = conviction_core * vol_proxy * novelty;
    reward += shaping_scale * bonus;
}
```

Thread new kernel params `trade_attempt_rate_ema_idx`, `trade_target_rate_idx`.

- [ ] **Step 3.7: Register in StateResetRegistry**

`TRADE_TARGET_RATE` is `FoldReset` (re-measures per fold). `TRADE_ATTEMPT_RATE_EMA` is also `FoldReset` (per-fold attempt rate).

In `state_reset_registry.rs`:

```rust
registry.register(IsvSlot::TradeAttemptRateEma, Category::FoldReset);
registry.register(IsvSlot::TradeTargetRate,    Category::FoldReset);
```

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

```bash
SQLX_OFFLINE=true FOXHUNT_TEST_DATA=test_data/futures-baseline \
    cargo test -p ml --lib -- trade_attempt_bonus --ignored --nocapture 2>&1 | tail -10
```

Expected: both tests PASS.

- [ ] **Step 3.9: Update audits**

- `docs/isv-slots.md` — add 2 rows.
- `docs/dqn-wire-up-audit.md` — add `trade_rate_ema_kernel` row.

- [ ] **Step 3.10: Commit**

```bash
git add crates/ml/src/cuda_pipeline/{gpu_dqn_trainer.rs,experience_kernels.cu,trade_rate_ema_kernel.cu} \
        crates/ml/build.rs \
        crates/ml/src/trainers/dqn/trainer/training_loop.rs \
        crates/ml/src/trainers/dqn/state_reset_registry.rs \
        crates/ml/src/trainers/dqn/smoke_tests/trade_attempt_bonus.rs \
        docs/isv-slots.md docs/dqn-wire-up-audit.md
git commit -m "plan3(task3): B.2 novelty-weighted trade-attempt bonus via ISV slots [50,51]"
```

---

## Task 4: B.4 Adaptive plan-MLP activation threshold — ✅ LANDED `8f59c1e3b`

> **Actual outcome.** New ISV slot `READINESS_EMA_INDEX = 75` (batch-mean readiness EMA, GPU-written). Existing `PLAN_THRESHOLD_INDEX = 49` producer upgraded from static constructor write to GPU kernel `plan_threshold_update_kernel.cu` (single-block, derives `max(0.1, 0.5 × ema)` from slot 75 into slot 49). Adaptive α = α_base × (1 + 0.5 × |sharpe|) matches Task 3's convention. New `readiness_per_sample [N*L]` buffer written by `experience_env_step`. Constructor cold-start writes preserved (0.5 → slot 49, 1.0 → slot 75) so consumers have sensible threshold before first kernel fire. Registry: PLAN_THRESHOLD flipped SchemaContract → FoldReset; READINESS_EMA registered FoldReset. Read-only `PlanThresholdMonitor` added (HEALTH_DIAG hookup deferred — matches existing 7 monitors' state). Fingerprint shifted [73,74] → [76,77]; `ISV_TOTAL_DIM` 75 → 78. Smoke fold-2 best Sharpe 93.72; HEALTH_DIAG `reward_split bonus=17.11`.

**Reality reconciliation (2026-04-24 second revision).** The original task text had two incorrect premises, both fixed here:

1. **"Allocate slot 49" was a collision.** Slot 49 already exists as `PLAN_THRESHOLD_INDEX` (Plan 1 Task 16 constructor-writes `0.5f` there, consumed by 4 sites in `experience_kernels.cu` + 1 site in `backtest_plan_kernel.cu` via `ISV_PLAN_THRESHOLD_IDX` in `state_layout.cuh`). Task 4 does NOT allocate a new slot — it upgrades the producer of slot 49 from static constructor write to GPU-driven EMA. Consumers unchanged.
2. **"EMA of plan_params[0]" was wrong-input.** Consumer sites compare the threshold against `readiness` and `exploration_scale` (both in `[0,1]`), NOT against `plan_params[0]` (which is `TARGET_BARS`, a raw bar count, unbounded). The adaptive signal should be readiness — EMA of the per-batch mean readiness, threshold = `0.5 × readiness_mean`. That way the plan activates when readiness is in the upper half of its own distribution, tracking the policy's actual confidence budget rather than a hardcoded 0.5.

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — add `READINESS_EMA_INDEX` constant (new slot at tail) + producer wiring for the existing `PLAN_THRESHOLD_INDEX = 49`; update constructor to stop writing the static 0.5 (now GPU-driven).
- Modify: `crates/ml/src/trainers/dqn/state_reset_registry.rs` — flip `PLAN_THRESHOLD_INDEX=49` category from `SchemaContract` (current, because it was static) to `FoldReset`; register new `READINESS_EMA` slot.
- Create: `crates/ml/src/cuda_pipeline/plan_threshold_update_kernel.cu` — producer.
- Create: `crates/ml/src/trainers/dqn/monitors/plan_threshold_monitor.rs` — read-only monitor implementing `AdaptiveMonitor`.
- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` — launch new kernel alongside reward/trade-rate EMA kernels in the per-epoch block.
- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — call `launch_plan_threshold_update_inplace` after experience collection (CPU never reads the value — consumer is the next epoch's kernel run).
- Test: rely on `multi_fold_convergence` smoke; no new unit test harness helpers to invent.

### Step 4.1: Allocate `READINESS_EMA_INDEX` at tail

In `gpu_dqn_trainer.rs`, after the Plan 3 Task 3 block, add the new slot and shift the fingerprint one slot further:

```rust
/// ISV slot [75] — Plan 3 Task 4 B.4: per-epoch mean readiness EMA
/// (α = α_base × (1 + 0.5 × |clamp(sharpe,−2,2)|)). Adaptive threshold
/// derived from this EMA is written into `PLAN_THRESHOLD_INDEX=49`
/// by the same kernel (no intermediate CPU step). Not a new threshold —
/// upgrades the producer of the existing threshold slot from static
/// constructor write to GPU-driven.
pub const READINESS_EMA_INDEX: usize = 75;
```

Recompute the fingerprint seed and move `ISV_LAYOUT_FINGERPRINT_LO/HI` from `[73,74]` to `[76,77]`. Grow `ISV_TOTAL_DIM` 75 → 78.

### Step 4.2: Implement the producer kernel

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

```cuda
#include <cuda_runtime.h>

/* plan_threshold_update — GPU-driven readiness EMA + derived threshold.
 *
 * Plan 3 Task 4 B.4. Reads `readiness_per_sample [n_samples]` (values in
 * [0,1] emitted by the readiness head during experience collection),
 * computes the batch mean on-GPU, EMAs it into ISV[READINESS_EMA_INDEX],
 * then derives `threshold = fmaxf(0.1f, 0.5f * ema)` and writes it into
 * ISV[PLAN_THRESHOLD_INDEX=49]. Consumers (experience_env_step,
 * backtest_plan_kernel) continue reading ISV[49] unchanged.
 *
 * Single-block, single-thread. No atomicAdd, no DtoH. */
extern "C" __global__ void plan_threshold_update(
    const float* __restrict__ readiness_per_sample, /* [n_samples] ∈ [0,1] */
    int          n_samples,
    float*       __restrict__ isv_out,
    int          readiness_ema_idx,         /* = READINESS_EMA_INDEX */
    int          plan_threshold_idx,        /* = PLAN_THRESHOLD_INDEX = 49 */
    int          sharpe_ema_idx,            /* = SHARPE_EMA_INDEX = 22 */
    float        alpha_base
) {
    if (threadIdx.x != 0 || blockIdx.x != 0) return;
    float sum = 0.0f;
    for (int idx = 0; idx < n_samples; ++idx) sum += readiness_per_sample[idx];
    const float mean   = sum / fmaxf(1.0f, (float)n_samples);
    const float sharpe = fmaxf(-2.0f, fminf(2.0f, isv_out[sharpe_ema_idx]));
    const float alpha  = alpha_base * (1.0f + 0.5f * fabsf(sharpe));
    const float prev   = isv_out[readiness_ema_idx];
    const float ema    = (1.0f - alpha) * prev + alpha * mean;
    isv_out[readiness_ema_idx] = ema;
    /* Derived threshold — floor at 0.1 for cold-start robustness. */
    isv_out[plan_threshold_idx] = fmaxf(0.1f, 0.5f * ema);
}
```

Register in `crates/ml/build.rs` alongside `trade_rate_ema_kernel.cu`.

### Step 4.3: Producer — emit `readiness_per_sample`

In `experience_kernels.cu::experience_env_step`, the kernel already reads `readiness_ptr[0]` (grep `readiness_ptr`). What's missing is a per-sample output buffer analogous to `flat_to_pos_per_sample`. Add:

```cuda
/* B.4 Plan 3 Task 4: readiness value this (i,t) slot saw. Written
 * unconditionally every reached slot (race-free per (i,t)). Consumer:
 * plan_threshold_update GPU kernel. */
float* __restrict__ readiness_per_sample,   /* [N*L] */
```

Writer block (early in the kernel, before any early-return):

```cuda
readiness_per_sample[out_off] = readiness;
```

Default in the pre-early-return block: `readiness_per_sample[out_off] = 0.0f;` (same pattern B.2 used).

Allocate the buffer in `gpu_experience_collector.rs` the same way `flat_to_pos_per_sample` is allocated; thread through the `experience_env_step` launch args.

### Step 4.4: Wire the launcher

In `gpu_experience_collector.rs`, alongside `launch_trade_attempt_rate_ema_inplace`, add:

```rust
pub fn launch_plan_threshold_update_inplace(&mut self, n_samples: usize, alpha_base: f32)
    -> Result<(), crate::MLError>
{
    use crate::cuda_pipeline::gpu_dqn_trainer::{
        READINESS_EMA_INDEX, PLAN_THRESHOLD_INDEX, SHARPE_EMA_INDEX,
    };
    if n_samples == 0 || self.isv_signals_dev_ptr == 0 { return Ok(()); }
    let n_i32 = n_samples as i32;
    let readiness_slot = READINESS_EMA_INDEX as i32;
    let thr_slot       = PLAN_THRESHOLD_INDEX as i32;
    let sharpe_slot    = SHARPE_EMA_INDEX as i32;
    let buf_ptr        = self.readiness_per_sample.raw_ptr();
    let isv_ptr        = self.isv_signals_dev_ptr;
    unsafe {
        self.stream.launch_builder(&self.plan_threshold_update_kernel)
            .arg(&buf_ptr).arg(&n_i32).arg(&isv_ptr)
            .arg(&readiness_slot).arg(&thr_slot).arg(&sharpe_slot)
            .arg(&alpha_base)
            .launch(LaunchConfig { grid_dim: (1,1,1), block_dim: (1,1,1), shared_mem_bytes: 0 })
            .map_err(|e| crate::MLError::ModelError(format!("plan_threshold_update: {e}")))?;
    }
    self.stream.memset_zeros(&mut self.readiness_per_sample)
        .map_err(|e| crate::MLError::ModelError(format!("memset readiness_per_sample: {e}")))?;
    Ok(())
}
```

In `training_loop.rs`, call it next to the existing Plan 3 kernel launches (after `launch_trade_attempt_rate_ema_inplace`).

### Step 4.5: Constructor — stop writing the static 0.5

In `gpu_dqn_trainer.rs`, find the constructor line `*sig_ptr.add(PLAN_THRESHOLD_INDEX) = 0.5_f32;` (Plan 1 Task 16). Replace with `*sig_ptr.add(PLAN_THRESHOLD_INDEX) = 0.5_f32;  // cold-start value; overwritten by plan_threshold_update on first epoch`. Leave the constructor-write in place so the cold start is sane, but add the comment documenting the upgrade.

Also update the ISV layout comment (~line 202) to reflect the new producer.

### Step 4.6: Read-only monitor

Create `crates/ml/src/trainers/dqn/monitors/plan_threshold_monitor.rs`:

```rust
use super::adaptive_monitor::{AdaptiveMonitor, DiagSnapshot, FireRateStats, IsvBus};

pub struct PlanThresholdMonitor;

impl AdaptiveMonitor for PlanThresholdMonitor {
    type Signal = f32;
    fn name(&self) -> &'static str { "plan_threshold" }
    fn read(&self, isv: &IsvBus) -> f32 {
        isv.get(crate::cuda_pipeline::gpu_dqn_trainer::PLAN_THRESHOLD_INDEX)
    }
    fn observe(&mut self, _value: f32) {}  /* no internal state */
    fn fire_rate(&self) -> FireRateStats { FireRateStats::always() }
    fn diagnose(&self, isv: &IsvBus) -> DiagSnapshot {
        DiagSnapshot::pairs(&[
            ("plan_threshold.eff",
                isv.get(crate::cuda_pipeline::gpu_dqn_trainer::PLAN_THRESHOLD_INDEX) as f64),
            ("plan_threshold.readiness_ema",
                isv.get(crate::cuda_pipeline::gpu_dqn_trainer::READINESS_EMA_INDEX) as f64),
        ])
    }
}
```

Register in `monitors/mod.rs` and in the `controller_activity` smoke harness.

### Step 4.7: Registry entries

In `state_reset_registry.rs`, after the existing `isv_plan_threshold` entry:

```rust
/* B.4 Plan 3 Task 4: upgrade PLAN_THRESHOLD_INDEX producer from static
 * constructor write to GPU-driven; add READINESS_EMA_INDEX for the
 * derived source. Both reset per fold so each fold measures its own
 * readiness distribution. */
RegistryEntry {
    name: "isv_readiness_ema",
    category: ResetCategory::FoldReset,
    description: "ISV[75] — batch-mean readiness EMA; GPU plan_threshold_update fills; drives PLAN_THRESHOLD_INDEX derivation",
},
```

Update the existing `isv_plan_threshold` row's description and category to reflect the new GPU producer (change from SchemaContract to FoldReset if it was SchemaContract). Add a matching fold-reset dispatch arm in `training_loop.rs`.

### Step 4.8: Verify + commit

```bash
SQLX_OFFLINE=true cargo check --workspace 2>&1 | tail -5                          # expect 11 warnings
SQLX_OFFLINE=true FOXHUNT_TEST_DATA=test_data/futures-baseline \
    cargo test -p ml --lib -- multi_fold_convergence --ignored --nocapture 2>&1 | tail -40   # expect pass, no kernel errors
```

Expected behavioral signal in HEALTH_DIAG: a new line or an extension of an existing line showing `plan_threshold_eff=X.XXX readiness_ema=Y.YYY`. Value of `plan_threshold_eff` should be roughly `0.5 × readiness_ema` with 0.1 floor.

### Step 4.9: Update audits

- `docs/isv-slots.md`: add row for `READINESS_EMA_INDEX=75`; update row for `PLAN_THRESHOLD_INDEX=49` to show new producer (was: constructor; now: `plan_threshold_update` kernel).
- `docs/dqn-wire-up-audit.md`: add row for `plan_threshold_update_kernel.cu`; bump Wired count.

```bash
git add crates/ml/src/cuda_pipeline/{gpu_dqn_trainer.rs,backtest_plan_kernel.cu,experience_kernels.cu} \
        crates/ml/src/trainers/dqn/adaptive_controllers/plan_mlp_threshold.rs \
        crates/ml/src/trainers/dqn/adaptive_controllers/mod.rs \
        crates/ml/src/trainers/dqn/trainer/training_loop.rs \
        crates/ml/src/trainers/dqn/smoke_tests/plan_activation_threshold.rs \
        docs/isv-slots.md docs/dqn-wire-up-audit.md
git commit -m "plan3(task4): B.4 adaptive plan-MLP activation threshold via ISV[49]"
```

---

## Task 5: C.4 Temporal timing bonus — ✅ LANDED `a0abc3da3`

> **Actual outcome.** New portfolio-state slot `PS_PEAK_PNL_BAR = 38` (hold_time at which MAX_PNL was last updated); `PS_STRIDE` grew 38 → 39. Shared-contract migration completed in same commit: `trade_stats_kernel.cu::PORTFOLIO_STRIDE`, `gpu_experience_collector.rs` allocator, `ml-core::state_layout.rs` mirror, and `launch_kelly_cap_update` hardcoded `ps_stride=38` all bumped to 39. Consumer at `segment_complete` block: `timing_bonus = shaping_scale × (bars_early / hold_time) × |reward| × conviction_core`, accumulated into `reward` and `rc[5]` (uses `+=` because Task 3 B.2 also touches rc[5] on a different (i,t)). No new ISV slot. Smoke Sharpe 84.44.

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/experience_kernels.cu` — on-exit temporal_timing_bonus
- Modify: per-sample trade-memo buffer to track `peak_unrealized_pnl_bar_offset`
- Test: `temporal_timing_bonus_rewards_early_peak`

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

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

```rust
#[test]
#[ignore = "gpu"]
fn temporal_timing_bonus_rewards_early_peak() {
    let mut trainer = build_minimal_trainer();

    // Entry at bar 0, peak unrealized P&L at bar 2, exit at bar 10.
    // hold_time = 10; bars_early = 10 - 2 = 8.
    // bonus = shaping_scale × (8/10) × |final_pnl| × conviction.
    let bonus_early = trainer.simulate_trade_exit_bonus(
        /*entry_bar=*/ 0,
        /*peak_bar=*/  2,
        /*exit_bar=*/  10,
        /*final_pnl=*/ 100.0,
        /*conviction=*/ 0.5,
    );

    // Same trade but peak at bar 9 (late).
    let bonus_late = trainer.simulate_trade_exit_bonus(0, 9, 10, 100.0, 0.5);

    assert!(bonus_early > 5.0 * bonus_late,
            "early peak should produce much larger bonus: early={}, late={}",
            bonus_early, bonus_late);
}
```

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

- [ ] **Step 5.3: Extend per-sample trade memo**

In the trade-memo struct (likely `crates/ml/src/cuda_pipeline/trade_memo.cuh` or defined inline in experience_kernels.cu), add:

```cuda
struct TradeMemo {
    // ... existing fields ...
    int    peak_unrealized_pnl_bar_offset;  // 0-based bars since entry
    float  peak_unrealized_pnl;
};
```

Per-bar updater:
```cuda
if (unrealized_pnl > memo.peak_unrealized_pnl) {
    memo.peak_unrealized_pnl            = unrealized_pnl;
    memo.peak_unrealized_pnl_bar_offset = bar_offset_since_entry;
}
```

- [ ] **Step 5.4: Compute bonus on exit**

```cuda
// C.4 Temporal timing bonus (on trade exit only).
if (is_trade_exit) {
    const int   hold_time_i = memo.exit_bar - memo.entry_bar;
    const float hold_time   = fmaxf(1.0f, (float)hold_time_i);
    const int   bars_early_i = hold_time_i - memo.peak_unrealized_pnl_bar_offset;
    const float bars_early  = fmaxf(0.0f, (float)bars_early_i);

    const float timing_bonus = shaping_scale
                             * (bars_early / hold_time)
                             * fabsf(memo.final_pnl)
                             * conviction_core;
    reward += timing_bonus;
}
```

- [ ] **Step 5.5: Feed into reward_component_audit as a new component**

Task 1 allocated 6 components (popart, cf, trail, micro, opp_cost, bonus). Use `bonus` to carry the timing bonus. The `bonus` slot is named generically on purpose so later components can share it with the existing B.2 trade-attempt bonus — both are conceptually "exploration credit" adjustments. HEALTH_DIAG comment in monitoring.rs should indicate this (add a one-line comment).

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

```bash
git add crates/ml/src/cuda_pipeline/experience_kernels.cu \
        crates/ml/src/trainers/dqn/smoke_tests/temporal_timing.rs \
        crates/ml/src/trainers/dqn/trainer/monitoring.rs
git commit -m "plan3(task5): C.4 temporal timing bonus on trade exit (peak_bar proximity)"
```

---

## Task 6: D.4 Generalised temporal reward coupling

**Reality reconciliation (2026-04-24 second revision).** Original slot numbers 59/60/61 are now occupied by Plan 2 Q-quantile slots (`[50..58)`) and Plan 3 Task 1's reward-component EMAs (`[63..69)`). Original reward formulae referenced PS/memo fields that don't exist (`min_mark_price`, `regime_at_exit`, `regime_shift_bar`, `pre_entry_conviction_std/mean`, `entry_notional`), would require a 6-field portfolio-state expansion, and included a tuned `0.5e-4f` penalty coefficient that violates `feedback_adaptive_not_tuned.md`. This revision splits Task 6 into three sub-tasks (6a, 6b, 6c) that each land independently with minimal PS growth and ISV-driven coefficients, all at tail after Task 4's slots.

Post-revision slot layout (assumes Task 4 lands first, so it starts at [76..79)):

| Slot | Name | Producer | Consumer |
|---|---|---|---|
| 76 | TEMPORAL_PERSIST_EMA_INDEX (6a) | `temporal_reward_ema_kernel` | HEALTH_DIAG `temporal_reward` |
| 77 | TEMPORAL_REGIME_SHIFT_EMA_INDEX (6b) | same kernel | same |
| 78 | TEMPORAL_CONSISTENCY_EMA_INDEX (6c) | same kernel | same |

Fingerprint shifts to `[79,80]`; `ISV_TOTAL_DIM` grows from 78 (post-Task-4) to 81.

> **IMPORTANT — task ordering**. Task 6 must land AFTER Task 4. If Task 4 is skipped, Task 6's tail-start index drops to 75. Recompute slot constants from the current `ISV_TOTAL_DIM` at implementation time — never hard-code from this plan.

### Task 6a: Persistence credit — ✅ LANDED `777341776`

> **Actual outcome.** `PS_INTRA_TRADE_MIN_PNL = 39` added; `PS_STRIDE` grew 39 → 40 with 6 hardcoded-stride sites + 7 comment mentions migrated in lockstep. MIN_PNL tracked per bar alongside MAX_PNL; reset at 5 MAX_PNL reset sites. Consumer fires on `segment_complete && reward > 0 && drawdown_depth > 1e-6f`, adds `shaping × conviction × |min_pnl| × tanh(reward/|min_pnl|)` to `reward` and `rc[5]`. Smoke fold-2 best Sharpe jumped 84.44 → 100.10; HEALTH_DIAG `reward_split bonus=17.21` confirms firing.

#### Original design (for reference)

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/state_layout.cuh` — add 1 PS slot `PS_INTRA_TRADE_MIN_PNL`; grow `PS_STRIDE` 39 → 40.
- Modify: `crates/ml/src/cuda_pipeline/trade_stats_kernel.cu` — bump hardcoded `PORTFOLIO_STRIDE` 39 → 40.
- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` — bump allocator `PORTFOLIO_STRIDE` 39 → 40.
- Modify: `crates/ml-core/src/state_layout.rs` — mirror the PS_STRIDE change.
- Modify: `crates/ml/src/cuda_pipeline/experience_kernels.cu` — track MIN_PNL analogous to MAX_PNL; compute persistence credit on exit.

**Design.**
- Producer: every bar in an active trade, update `ps[PS_INTRA_TRADE_MIN_PNL] = fminf(ps[PS_INTRA_TRADE_MIN_PNL], pnl_pct)`. Reset to 0 at every site that resets `PS_INTRA_TRADE_MAX_PNL` (grep `PS_INTRA_TRADE_MAX_PNL] = 0.0f` — Task 5 already established the reset-site list).
- Consumer (at `segment_complete && reward > 0.0f`):
  ```cuda
  const float drawdown_depth = fmaxf(0.0f, -ps[PS_INTRA_TRADE_MIN_PNL]);  // |min PnL|
  const float persist_bonus  = shaping_scale * conviction_core
                             * drawdown_depth * tanhf(reward / fmaxf(1e-4f, drawdown_depth));
  reward               += persist_bonus;
  reward_components_per_sample[out_off * 6 + 5] += persist_bonus;  // rc[5] bonus, += for co-existence with B.2/C.4
  ```
- Self-scaling rationale: `tanh(reward / drawdown_depth)` saturates at ±1 when reward ≫ drawdown (full credit) and approaches 0 when reward ≈ drawdown (no surprise). No tuned coefficients.

### Task 6b: Regime-shift penalty — ✅ LANDED `5ebebc564`

> **Actual outcome.** `PS_REGIME_SHIFT_BAR = 40` added; `PS_STRIDE` grew 40 → 41 via 6 stride sites. Detector fires first bar where `|regime_now - PS_PLAN_ENTRY_REGIME|` exceeds `clamp(0.25 × |clamp(sharpe, −2, 2)|, 0.05, 0.5)`. Consumer on exit: `penalty = shaping × conviction × bars_late_frac × |reward|` (bounded by |reward|). Smoke fold-2 best Sharpe: 117.92. **Iteration note:** first formula multiplied by `ISV[Q_DIR_ABS_REF] × |reward|` simultaneously — absolute Q-magnitude (5–50) amplified |reward| by 20×, flipping trade outcomes and destabilising training. Fixed by dropping `q_scale` multiplier. Pearl: pick exactly ONE unbounded scaling signal per reward term; bound everything else.

### Task 6b: Regime-shift penalty (held past regime flip)

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/state_layout.cuh` — add 1 PS slot `PS_REGIME_SHIFT_BAR` (hold_time at which regime_stability crossed the entry value); grow `PS_STRIDE` 40 → 41.
- Shared-contract migration: bump `PORTFOLIO_STRIDE` in `trade_stats_kernel.cu`, `gpu_experience_collector.rs`, `ml-core::state_layout.rs`, `launch_kelly_cap_update` hardcoded copy. Same pattern as Task 5.
- Modify: `crates/ml/src/cuda_pipeline/experience_kernels.cu`.

**Design.**
- `PS_PLAN_ENTRY_REGIME` already captures regime at entry (Task 3 Plan B.2 landing). Per bar during an active trade: if `|current_regime - ps[PS_PLAN_ENTRY_REGIME]| > regime_shift_eps` AND `ps[PS_REGIME_SHIFT_BAR] == 0.0f`, set `ps[PS_REGIME_SHIFT_BAR] = ps[PS_HOLD_TIME]`. `regime_shift_eps` is derived adaptively: `0.25 × |ISV[SHARPE_EMA_INDEX]|` (clamped to `[0.05, 0.5]`) — the threshold tightens when the model is confidently trading.
- Consumer on exit:
  ```cuda
  if (ps[PS_REGIME_SHIFT_BAR] > 0.0f) {
      const float bars_late = fmaxf(0.0f, (float)saved_hold_time - ps[PS_REGIME_SHIFT_BAR]);
      /* Scale coefficient by ISV[Q_DIR_ABS_REF_INDEX] — the same self-scaling the
       * B.1 opp-cost uses. No tuned 0.5e-4f. */
      const float q_scale   = isv_signals_ptr[Q_DIR_ABS_REF_INDEX];
      const float penalty   = shaping_scale * conviction_core
                            * (bars_late / fmaxf(1.0f, (float)saved_hold_time))
                            * q_scale * fabsf(reward);
      reward                -= penalty;
      reward_components_per_sample[out_off * 6 + 5] -= penalty;  /* rc[5], same as persist */
  }
  ps[PS_REGIME_SHIFT_BAR] = 0.0f;  /* reset for next trade */
  ```
- Reset `PS_REGIME_SHIFT_BAR = 0` at every trade-entry + fold-boundary reset (same list as MIN_PNL).

### Task 6c: Conviction consistency bonus — ✅ LANDED `0bbe97ed8`

> **Actual outcome.** PS slots 41/42 (`PS_PRE_ENTRY_CONVICTION_EMA`, `PS_PRE_ENTRY_CONVICTION_VAR_EMA`) added; `PS_STRIDE` grew 41 → 43 via the same 6 hardcoded-stride migration pattern. Producer updates Welford-EMA of `conviction_core` each Flat bar (α=0.05). Consumer at Flat→Positioned entry fires iff `stddev/mean < 0.2` (the single permitted tuned constant), adds `shaping × vol_proxy × stability × conviction_core` to reward + rc[5] — max ≈ 0.01, mirroring B.2 novelty-bonus shape. Single unbounded multiplicand is `vol_proxy` per the one-unbounded-signal pearl. Smoke fold-2 best Sharpe 92.40 at epoch 3; HEALTH_DIAG `reward_split bonus=17.95`. 25-pt swing below T6b's 117 is at the edge of noise — flag for L40S validation watch.

### Task 6c: Conviction consistency bonus (stable conviction before entry)

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/state_layout.cuh` — add 2 PS slots: `PS_PRE_ENTRY_CONVICTION_EMA`, `PS_PRE_ENTRY_CONVICTION_VAR_EMA` (Welford variance via EMA). Grow `PS_STRIDE` 41 → 43.
- Same shared-contract migration as 6a/6b.
- Modify: `crates/ml/src/cuda_pipeline/experience_kernels.cu`.

**Design.**
- Every Flat bar: update running mean-var of `conviction_core` via EMA (α = 0.05). Store mean in slot A, running var in slot B.
- On Flat→Positioned entry: if `sqrt(var) < 0.2 × mean` (stable conviction), apply bonus:
  ```cuda
  const float cons_stddev = sqrtf(fmaxf(0.0f, ps[PS_PRE_ENTRY_CONVICTION_VAR_EMA]));
  const float cons_mean   = fmaxf(1e-6f, ps[PS_PRE_ENTRY_CONVICTION_EMA]);
  if (cons_stddev < 0.2f * cons_mean) {   /* stable */
      const float stability = 1.0f - cons_stddev / cons_mean;  /* ∈ [0,1] */
      const float bonus     = shaping_scale * fmaxf(0.0f, stability) * conviction_core;
      reward                += bonus;
      reward_components_per_sample[out_off * 6 + 5] += bonus;
  }
  ```
- Reset both EMA slots to 0 at trade entry (conviction state is per-Flat-epoch only). The `0.2f` stability threshold is the only tuned constant in Task 6c — if the reviewer prefers ISV-driven, derive `stability_thresh = 0.2 × clamp(ISV[Q_DIR_ABS_REF_INDEX]/5.0, 0.5, 2.0)`. Defer that hardening to a follow-up if it complicates the first cut.

### Step 6.X: EMA producer kernel (shared by 6a, 6b, 6c)

Create `crates/ml/src/cuda_pipeline/temporal_reward_ema_kernel.cu` mirroring `reward_component_ema_kernel.cu` but aggregating 3 per-sample accumulators (one per sub-task). Land in the first of 6a/6b/6c to actually implement; subsequent sub-tasks just add their accumulator. Alternatively, reuse `rc[5]` alone and drop this kernel — Task 5 already feeds `rc[5]` through `reward_component_ema_kernel` into ISV[68] REWARD_BONUS_EMA, so persist/regime-shift/consistency contributions all land there automatically. **Prefer the rc[5]-only path** unless HEALTH_DIAG needs per-sub-task breakdown (simpler, fewer slots, matches Task 5's precedent). If per-sub-task breakdown is required later, introduce it as a Plan 4 follow-up.

With rc[5]-only, slots 76/77/78 aren't needed. Task 6 keeps `ISV_TOTAL_DIM` unchanged from post-Task-4 = 78. Fingerprint stays at [76,77].

### Step 6.Y: Register PS slots

In `state_reset_registry.rs`, add one `FoldReset` entry per new PS slot (names: `ps_intra_trade_min_pnl`, `ps_regime_shift_bar`, `ps_pre_entry_conviction_ema`, `ps_pre_entry_conviction_var_ema`). Follow the existing pattern.

### Step 6.Z: Validate

Run the standard multi-fold smoke after each sub-task (6a, 6b, 6c) — each should pass with roughly baseline Sharpe. Because all three touch `rc[5]` with `+=`, watch for sign cancellation in HEALTH_DIAG `reward_split bonus=X` — if bonus EMA is near 0 but individual sub-tasks report activity, that's probably persist + regime-shift cancelling. Add a `persist_fired` and `regime_penalty_fired` per-sample counter (one-bit) if the cancellation becomes hard to audit.

### Commit

One commit per sub-task. Each commit ships its PS stride bump + all shared-contract consumers in the same commit. Example:

```bash
git commit -m "feat(dqn-v2): D.4a persistence credit — shaping×conviction×|min_pnl|×tanh(reward/|min_pnl|)"
```

---

## Task 7: C.3 State-distribution divergence signal

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — 3 new ISV slot constants
- Create: `crates/ml/src/cuda_pipeline/state_kl_divergence_kernel.cu` — producer
- Modify: val-side backtest to call the producer on the val-state batch
- Modify: `experience_kernels.cu` — consume STATE_KL_AMPLIFICATION for reward scaling
- Test: `state_kl_fires_on_distribution_shift`

### ✅ LANDED `673eb6612`

> **Actual outcome.** ISV slots [78] STATE_KL_TRAIN_VAL_EMA + [79] STATE_KL_AMPLIFICATION; fingerprint shifted [76,77] → [80,81]; ISV_TOTAL_DIM 78 → 82. New kernel `state_kl_divergence_kernel.cu` does single-block per-dim Gaussian moment-match KL on the OFI block (32 dims), kernel-internal trailing-EMA-of-self for the amp ratio (no extra threshold slot). Val-side launch via `GpuBacktestEvaluator::val_state_sample()` (returns `chunked_states_buf` device pointer + sample count). Consumer wires kl_amp ∈ [1,2] into B.1 opp_cost and B.2 bonus via `fmaxf(1.0, …)` guard against cold-start 0. Smoke 3-fold convergence: best at ep5/ep2/ep3 — *diverse* epoch positions, addressing the earlier "always ep1" concern. Per-fold best Sharpe 105.13 / 83.28 / 103.49.

**Reality reconciliation.** Original slot numbers 58/63/64 collide with Plan 2 Q-quantile [50..58) and Plan 3 Task 1 reward-component EMAs [63..69). Original fire-trigger used a tuned `2.0× threshold` multiplier and amplification-decay constants `0.02f`, `2.0f`, `1.0f` that violate `feedback_adaptive_not_tuned.md`. This revision tail-appends after Task 6 and makes the amplification formula ISV-derived.

Post-revision slot layout (assumes Tasks 4+6a+6b+6c land first — actual ISV_TOTAL_DIM ≥ 78 at Task 7 start):

| Slot | Name | Semantics |
|---|---|---|
| N | STATE_KL_TRAIN_VAL_EMA_INDEX | per-epoch mean moment-match KL(train ‖ val) |
| N+1 | STATE_KL_AMPLIFICATION_INDEX | bonus multiplier ∈ [1, 2], grows as KL stays elevated |

Threshold EMA is inlined into the kernel (no separate slot) — compared against trailing-EMA-of-self which is a stateful kernel computation, not an ISV slot. Saves 1 slot vs. the original design. Recompute `N` from current `ISV_TOTAL_DIM` at implementation time.

### Step 7.1: Allocate slots at tail

```rust
// ─── C.3 State-distribution divergence signal ────────────────────────
// Plan 3 Task 7. Tail-appended — indices recomputed from ISV_TOTAL_DIM at task start.
pub const STATE_KL_TRAIN_VAL_EMA_INDEX: usize = /* N */;
pub const STATE_KL_AMPLIFICATION_INDEX: usize = /* N+1 */;
```

Grow `ISV_TOTAL_DIM` by +2. Fingerprint re-tails.

### Step 7.2: Kernel — adaptive coefficients, no tuned multipliers

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

```cuda
#include <cuda_runtime.h>

/* State-distribution moment-match KL: KL_gauss(train ‖ val) summed over
 * monitored dims. Cold-path — runs once per val-window boundary. Single
 * block (D_monitored ≤ 32 in practice); block reduces per-dim means/vars
 * via shared mem.
 *
 * Amplification formula is ISV-driven:
 *   ratio = new_kl_ema / max(prev_threshold_ema, 1e-6f)
 *   target = 1.0f + fminf(1.0f, fmaxf(0.0f, ratio - 1.0f))   // ∈ [1,2]
 *   amp_alpha = alpha_base * (1.0f + 0.5f * |clamp(sharpe, -2, 2)|)
 *   new_amp = (1 - amp_alpha) * prev_amp + amp_alpha * target
 *
 * No tuned 2.0× trigger, no tuned 0.02 decay — everything follows alpha
 * and the ratio-based target. When KL matches baseline (ratio ≤ 1) target
 * is 1.0 (no amplification); when KL is double baseline, target is 2.0. */
extern "C" __global__ void state_kl_moment_match(
    const float* __restrict__ train_state,
    const float* __restrict__ val_state,
    int          n_train,
    int          n_val,
    int          d_total,
    int          monitored_dims_begin,
    int          monitored_dims_end,
    float*       isv_out,
    int          kl_ema_idx,
    int          kl_amplification_idx,
    int          sharpe_ema_idx,
    float        alpha_base)
{
    /* Standard per-dim reduction on smem; final KL is summed across monitored
     * dims. Implementation detail: one warp per dim, tree-reduce, thread 0
     * aggregates; then thread 0 writes the EMA updates. Full implementation
     * mirrors `trade_rate_ema_kernel.cu` structure. */
    ...
    if (threadIdx.x == 0 && blockIdx.x == 0) {
        const float new_kl      = /* summed KL */;
        const float prev_kl_ema = isv_out[kl_ema_idx];
        const float sharpe      = fmaxf(-2.0f, fminf(2.0f, isv_out[sharpe_ema_idx]));
        const float kl_alpha    = alpha_base * (1.0f + 0.5f * fabsf(sharpe));
        const float new_kl_ema  = (1.0f - kl_alpha) * prev_kl_ema + kl_alpha * new_kl;
        isv_out[kl_ema_idx]     = new_kl_ema;

        /* Implicit threshold = the slower-EMA of KL itself (half alpha). */
        const float threshold   = (prev_kl_ema > 0.0f) ? prev_kl_ema : new_kl_ema;
        const float ratio       = new_kl_ema / fmaxf(1e-6f, threshold);
        const float target      = 1.0f + fminf(1.0f, fmaxf(0.0f, ratio - 1.0f));
        const float prev_amp    = isv_out[kl_amplification_idx];
        const float amp_alpha   = alpha_base * (1.0f + 0.5f * fabsf(sharpe));
        isv_out[kl_amplification_idx] = (1.0f - amp_alpha) * prev_amp + amp_alpha * target;
    }
}
```

### Step 7.3: Wire into val backtest

In `gpu_backtest_evaluator.rs` at the per-epoch val-window boundary, launch the kernel with pointers to a sampled train-state batch and the val-state batch. Sample N_train ≤ 4096 train states per fold (not per step — cold path). Monitored dims: the ISV-feature block and the portfolio block — skip OFI since it has different normalization.

### Step 7.4: Consumer — multiply B.1/B.2 rewards by amplification

In `experience_kernels.cu`, at the B.1 opp-cost site AND the B.2 novelty-bonus site:

```cuda
const float amp = (isv_signals_ptr != NULL)
    ? fmaxf(1.0f, isv_signals_ptr[STATE_KL_AMPLIFICATION_INDEX])
    : 1.0f;
// B.1: opp_cost *= amp; B.2: bonus_scaled *= amp;
```

When KL fires (amp ≈ 2), flat-trap escape pressure doubles. When distributions match (amp = 1), no-op. Gated `fmaxf(1.0f, ...)` guards against pre-populated slot (amp=0 → silent drop in reward).

### Step 7.5: Registry + monitor

`FoldReset` entries for both slots. Read-only `StateKLMonitor` follows the `PlanThresholdMonitor` template from Task 4.

### Step 7.6: Verify + commit

Validate via `multi_fold_convergence` smoke; no bespoke test helpers. Expected: no regression, `STATE_KL_AMPLIFICATION` slot reads ~1.0 in smoke (train and val distributions are adjacent folds of the same instrument; KL is small). On an intentionally-shifted val (later Argo validation run), amp should climb toward 1.5–2.0.

```bash
git commit -m "feat(dqn-v2): C.3 state-KL divergence + adaptive amp ISV slots"
```

---

## Task 8: B.3 Replay Buffer Seeded Warm-Start — ✅ LANDED `3cb083f18`

> **Actual outcome (full GPU-only redesign per user direction "fully gpu driven no cpu involvement").** The original CPU-trait `ScriptedPolicy` design + `simulate_scripted_transition` CPU physics mirror were **rejected and replaced** with a single CUDA kernel. ISV slots [82] SEED_STEPS_TARGET, [83] SEED_STEPS_DONE, [84] SEED_FRAC_EMA. Fingerprint shifted [80,81] → [85,86]; ISV_TOTAL_DIM 82 → 87. New kernels: `scripted_policy_kernel.cu` (4 policies in one kernel — 40% uniform LCG / 20% momentum / 20% mean-reversion / 20% vwap-deviation, deterministic per-sample policy assignment via `i % 5`), `seed_step_counter_update_kernel.cu` (cold-path counter + EMA). Action source switched at launch boundary: CPU reads ISV[83] vs ISV[82] once per epoch (cold-path orchestration; action computation itself is 100% GPU) and dispatches scripted_policy_select OR experience_action_select. No CPU physics mirror — `experience_env_step` runs unchanged. `replay_seed_steps` config field added to `DQNHyperparameters` + TOML override in `dqn-smoketest.toml` (1000 steps for smoke visibility vs. 100k production). Read-only `SeedMonitor` registered. Smoke fold-2 best Sharpe 92.49 at epoch 1.

## Task 8 (original CPU-based design — superseded)

**Files:**
- Create: `crates/ml/src/trainers/dqn/dqn_replay_seed.rs` — seed harness
- Create: `crates/ml/src/trainers/dqn/scripted_policies/momentum.rs` — 20% momentum
- Create: `crates/ml/src/trainers/dqn/scripted_policies/mean_reversion.rs` — 20% contrarian
- Create: `crates/ml/src/trainers/dqn/scripted_policies/vwap_deviation.rs` — 20% VWAP
- Create: `crates/ml/src/trainers/dqn/scripted_policies/uniform_random.rs` — 40% random
- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — seed step before epoch 1
- Test: `replay_seed_fills_buffer_with_diverse_experiences`

**Reality reconciliation.** Original text assumed trainer helpers (`sample_state_feature_pair`, `step_scripted`, `replay_insert_with_priority_scale`) that don't exist, and a tuned `priority_scale=0.5` that violates `feedback_adaptive_not_tuned.md`. This revision enumerates the prerequisite trainer helpers that must land in the same commit, replaces `priority_scale` with an ISV-driven value, and clarifies the cold-path-only lifecycle.

**Prerequisite helpers (must land in Task 8's commit — none exist today):**
1. `DqnTrainer::sample_feature_window(bar_idx: usize) -> FeatureWindow` — returns a window of the preprocessed feature matrix at an arbitrary bar. The trainer already has this matrix in the fxcache-loaded feature buffer; expose a read-only view with a defined lookback (e.g. 64 bars).
2. `DqnTrainer::simulate_scripted_transition(&mut self, bar_idx: usize, action: u32, conviction: f32) -> (StateVector, f32, bool)` — CPU-side single-step environment simulation reusing the same physics (trade execution, Kelly cap, slippage, fee) as the GPU `experience_env_step` kernel. Returns `(next_state, reward, done)`. Must produce identical reward-formula output as the GPU kernel would for matching inputs (invariant: scripted and live experiences flow through the same reward code-path).
3. `DqnTrainer::replay_insert(&mut self, state, action, reward, next_state, done, priority_weight: f32)` — wraps the PER insert with a per-insertion priority weight in `[0,1]` (existing PER insert already supports weighted priority).

**Scripted-policy portfolio** (40% uniform, 20% each of momentum/mean-reversion/vwap-deviation) is the same as original.

### Step 8.1: Scripted-policy trait

```rust
// crates/ml/src/trainers/dqn/scripted_policies/mod.rs
pub trait ScriptedPolicy: Send + Sync {
    fn name(&self) -> &'static str;
    /// Returns (factored_action_idx ∈ [0, 108), conviction ∈ [0, 1]).
    fn choose_action(&self, state: &StateVector, features: &FeatureWindow) -> (u32, f32);
}
```

Implement `UniformRandomPolicy`, `MomentumPolicy`, `MeanReversionPolicy`, `VwapDeviationPolicy` as separate files. Each is pure-CPU, no mutable state, no allocation on hot path.

### Step 8.2: Seed harness

```rust
pub struct ReplaySeedHarness {
    portfolio: Vec<(Box<dyn ScriptedPolicy>, f32)>,  // (policy, fraction)
}

pub struct SeedStats {
    pub total_experiences: usize,
    pub per_policy: std::collections::HashMap<&'static str, usize>,
    pub distinct_actions: std::collections::BTreeSet<u32>,
}

impl ReplaySeedHarness {
    pub fn run(&self, trainer: &mut DqnTrainer, target: usize) -> anyhow::Result<SeedStats> {
        let mut stats = SeedStats { total_experiences: 0, per_policy: Default::default(), distinct_actions: Default::default() };
        /* Read seed priority scale from ISV — allocated at tail as SEED_PRIORITY_SCALE_INDEX.
         * Constructor-writes initial value (cold start); Task 9 will migrate to GPU kernel
         * update based on seed-displacement signal. */
        let priority_scale = trainer.read_isv_signal_at(SEED_PRIORITY_SCALE_INDEX);
        let priority_scale = priority_scale.clamp(0.1, 1.0);  /* safe range */

        for (policy, frac) in &self.portfolio {
            let count = (target as f32 * frac) as usize;
            for _ in 0..count {
                let bar = trainer.sample_seedable_bar()?;  /* uniform over train-fold bars */
                let features = trainer.sample_feature_window(bar);
                let state    = trainer.snapshot_state_at(bar);  /* zero-positioned entry state */
                let (action, conviction) = policy.choose_action(&state, &features);
                let (next_state, reward, done) = trainer.simulate_scripted_transition(bar, action, conviction);
                trainer.replay_insert(state, action, reward, next_state, done, priority_scale);

                *stats.per_policy.entry(policy.name()).or_insert(0) += 1;
                stats.distinct_actions.insert(action);
                stats.total_experiences += 1;
            }
        }
        Ok(stats)
    }
}
```

### Step 8.3: Allocate `SEED_PRIORITY_SCALE_INDEX`

Tail-append after Task 7's state-KL slots. Constructor writes `0.5` (same starting value as original plan, but now it's an ISV slot so Task 9 can make it adaptive). Registry entry: `FoldReset`.

### Step 8.4: Invoke in training_loop

```rust
if !self.resumed_from_checkpoint {
    let target = self.config.replay_seed_experiences as usize;  /* 100_000 default */
    let stats  = ReplaySeedHarness::new(default_portfolio()).run(self, target)?;
    tracing::info!(target: "dqn::B3",
        "Replay seed: total={} per_policy={:?} distinct_actions={}/108",
        stats.total_experiences, stats.per_policy, stats.distinct_actions.len());
}
```

Add `replay_seed_experiences: u32` to `DqnTrainerConfig` with default `100_000`.

### Step 8.5: Seed-origin vs. dqn-origin tagging (for Task 9)

The PER entry needs a single-bit tag distinguishing seed-origin from DQN-origin experiences. Task 9's CQL-α ramp reads an aggregated seed-fraction from this tag. Add a `seed_origin: bool` field to the replay entry; update `replay_insert` signatures; emit the GPU-side aggregation in a new `seed_fraction_kernel.cu` (Task 9 builds on this foundation).

### Step 8.6: Validate + commit

Without a dedicated test harness (the trainer helpers need real GPU state), validate via:
```bash
SQLX_OFFLINE=true FOXHUNT_TEST_DATA=test_data/futures-baseline \
    cargo test -p ml --lib -- multi_fold_convergence --ignored --nocapture 2>&1 | grep -E "Replay seed:|distinct_actions"
```
Expected log line showing `distinct_actions ≥ 60/108` and `per_policy` ratios near 40/20/20/20. Smoke Sharpe should land in the same range; seeded warm-start typically converges faster in the first 5 epochs (+1–3 Sharpe vs. cold-start baseline at fold-2 ep-1).

```bash
git commit -m "feat(dqn-v2): B.3 scripted-policy replay warm-start + ISV-driven priority scale"
```

---

## Task 9: C.5 CQL alpha schedule coupled to seed-phase decay — ✅ LANDED `3cb083f18`

> **Actual outcome (landed in same commit as Task 8).** New kernel `cql_alpha_seed_update_kernel.cu`: `target = config.cql_alpha × max(0, 1 - SEED_FRAC_EMA)`; adaptive α matching Task 3 convention. Updates ISV[CQL_ALPHA_INDEX=48] each epoch. CQL loss kernel consumer was already wired to ISV[48] via pinned device-mapped read (Plan 1 Task 12 / Bug #7 fix `90e1e3dbb`) — Step 7 was a verified no-op. CQL_ALPHA registry entry flipped SchemaContract → FoldReset; cold-start preserved. Read-only `CqlAlphaMonitor` registered. During seed phase (frac≈1) cql_alpha decays toward 0 (no pessimism on exploration data); as frac→0 it ramps to `config.cql_alpha`. Verified in smoke via HEALTH_DIAG `cql_alpha_eff` showing the ramping behavior consistent with seed-phase progress.

## Task 9 (original CPU-controller design — superseded)

**Reality reconciliation.** Original design used the pre-pivot `AdaptiveMonitor` trait with CPU-side `update()` returning a control signal — doesn't match the GPU-drives-CPU-reads pivot. Also used tuned EMA rates (`0.9`, `0.1`) that violate `feedback_adaptive_not_tuned.md`. This revision reshapes Task 9 into the canonical GPU-kernel-producer + read-only-monitor pattern, with an adaptive EMA rate coupled to Sharpe.

**Dependency:** Task 8 must land first — specifically Step 8.5's `seed_origin: bool` tagging and `seed_fraction_kernel.cu`. Task 9 consumes the `SEED_FRACTION_EMA_INDEX` ISV slot that Task 8 populates.

**Existing `PLAN_1_T12 CQL_ALPHA_INDEX = 48` (static constructor write of `config.cql_alpha`).** Task 9 upgrades this from static to GPU-driven, analogous to Task 4's upgrade of PLAN_THRESHOLD_INDEX. No new cql_alpha slot.

**Files:**
- Create: `crates/ml/src/cuda_pipeline/cql_alpha_seed_update_kernel.cu` — GPU-drives producer reading `SEED_FRACTION_EMA_INDEX` and writing `CQL_ALPHA_INDEX=48`.
- Create: `crates/ml/src/trainers/dqn/monitors/cql_alpha_monitor.rs` — read-only, follows `PlanThresholdMonitor` template.
- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` — launcher method.
- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — call launcher in the per-epoch Plan 3 kernel block; remove any static `cql_alpha` push to CUDA config (it now reads from ISV).
- Modify: `state_reset_registry.rs` — flip CQL_ALPHA category from `SchemaContract` to `FoldReset`.

### Step 9.1: Kernel — GPU producer

```cuda
#include <cuda_runtime.h>

/* cql_alpha_seed_update — couple CQL α to replay seed displacement.
 *
 * Ramp formula:
 *   seed_disp  = max(0, 1 - ISV[SEED_FRACTION_EMA_INDEX])  /* DQN-origin fraction */
 *   target     = config_final_alpha * seed_disp * saturation
 *     where saturation = fminf(1, step / warmup_time_scale)
 *   alpha_rate = alpha_base * (1 + 0.5 * |clamp(sharpe, -2, 2)|)   /* adaptive EMA */
 *   new_alpha  = (1 - alpha_rate) * prev_alpha + alpha_rate * target
 *
 * Writes into ISV[CQL_ALPHA_INDEX=48]. The existing Bug #7 regime-gate
 * (`cql_alpha_regime_gate` in the CQL loss kernel) continues to wrap
 * this ISV read — no semantic change for null-ISV paths.
 *
 * No tuned 0.9/0.1 EMA rate — alpha_rate tracks training health via Sharpe. */
extern "C" __global__ void cql_alpha_seed_update(
    float*       __restrict__ isv_out,
    int          seed_fraction_idx,     /* = SEED_FRACTION_EMA_INDEX */
    int          cql_alpha_idx,         /* = CQL_ALPHA_INDEX = 48 */
    int          sharpe_ema_idx,        /* = SHARPE_EMA_INDEX = 22 */
    int          step_count,
    int          warmup_time_scale,
    float        final_alpha,
    float        alpha_base
) {
    if (threadIdx.x != 0 || blockIdx.x != 0) return;
    const float seed_frac = isv_out[seed_fraction_idx];
    const float seed_disp = fmaxf(0.0f, 1.0f - seed_frac);
    const float sat       = fminf(1.0f, (float)step_count / fmaxf(1.0f, (float)warmup_time_scale));
    const float target    = final_alpha * seed_disp * sat;
    const float sharpe    = fmaxf(-2.0f, fminf(2.0f, isv_out[sharpe_ema_idx]));
    const float alpha     = alpha_base * (1.0f + 0.5f * fabsf(sharpe));
    const float prev      = isv_out[cql_alpha_idx];
    isv_out[cql_alpha_idx] = (1.0f - alpha) * prev + alpha * target;
}
```

Register in `build.rs`.

### Step 9.2: Launcher

In `gpu_experience_collector.rs`:

```rust
pub fn launch_cql_alpha_seed_update_inplace(
    &mut self,
    step_count: u64,
    warmup_time_scale: u64,
    final_alpha: f32,
    alpha_base: f32,
) -> Result<(), crate::MLError> {
    use crate::cuda_pipeline::gpu_dqn_trainer::{
        SEED_FRACTION_EMA_INDEX, CQL_ALPHA_INDEX, SHARPE_EMA_INDEX,
    };
    if self.isv_signals_dev_ptr == 0 { return Ok(()); }
    let step    = step_count.min(i32::MAX as u64) as i32;
    let warmup  = warmup_time_scale.min(i32::MAX as u64) as i32;
    let seed_s  = SEED_FRACTION_EMA_INDEX as i32;
    let cql_s   = CQL_ALPHA_INDEX as i32;
    let sharpe_s= SHARPE_EMA_INDEX as i32;
    let isv_ptr = self.isv_signals_dev_ptr;
    unsafe {
        self.stream.launch_builder(&self.cql_alpha_seed_update_kernel)
            .arg(&isv_ptr).arg(&seed_s).arg(&cql_s).arg(&sharpe_s)
            .arg(&step).arg(&warmup).arg(&final_alpha).arg(&alpha_base)
            .launch(LaunchConfig { grid_dim: (1,1,1), block_dim: (1,1,1), shared_mem_bytes: 0 })
            .map_err(|e| crate::MLError::ModelError(format!("cql_alpha_seed_update: {e}")))?;
    }
    Ok(())
}
```

### Step 9.3: Training-loop wiring

In `training_loop.rs`, in the per-epoch Plan 3 kernel block (same location as `launch_trade_attempt_rate_ema_inplace`):

```rust
if let Some(ref mut c) = self.gpu_experience_collector {
    c.launch_cql_alpha_seed_update_inplace(
        self.total_steps_completed,
        self.config.cql_alpha_warmup_steps as u64,  /* new config field */
        self.config.cql_alpha_final,                 /* existing field, e.g. 0.014 */
        0.05,                                         /* α_base */
    )?;
}
```

Update the CQL loss kernel call site (`cql_loss_kernel.cu` or wherever `cql_alpha` is consumed) to read `ISV[CQL_ALPHA_INDEX]` rather than `config.cql_alpha`. The existing null-ISV guard ("Bug #7" in commit `90e1e3dbb`) already handles the fallback path correctly — no changes there.

### Step 9.4: Read-only monitor + registry

- `CqlAlphaMonitor` reads ISV[48] and surfaces it in HEALTH_DIAG next to the existing `cql_alpha=X.XXXX` line (rename current readout to `cql_alpha_eff` to match the `tau_eff`, `epsilon_eff` convention).
- Registry entry for CQL_ALPHA_INDEX shifts category `SchemaContract → FoldReset` (per-fold recomputes from step 0 ramp).
- Add fold-reset dispatch arm in `training_loop.rs`.

### Step 9.5: Validate + commit

```bash
SQLX_OFFLINE=true cargo check --workspace 2>&1 | tail -5
SQLX_OFFLINE=true FOXHUNT_TEST_DATA=test_data/futures-baseline \
    cargo test -p ml --lib -- multi_fold_convergence --ignored --nocapture 2>&1 | grep -E "cql_alpha|HEALTH_DIAG"
```

Expected HEALTH_DIAG: `cql_alpha_eff` starts near 0 at epoch 1 (seed dominates buffer), ramps toward `config.cql_alpha_final` as epochs progress. On the 5-epoch local smoke, ramp may not fully complete — that's fine; the Argo validation run at Task 10 is where the full ramp gets verified.

```bash
git commit -m "feat(dqn-v2): C.5 CQL α GPU-driven + seed-displacement coupled"
```

---

## Task 10: Plan 3 validation run — multi-seed × multi-fold, Tier 1 + Tier 2 exit

**Files:**
- Create (or modify): `config/metric-bands.toml` — expanded with Plan 3's new metrics
- Modify: `docs/dqn-v2-plan-3-validation.md` — log final run results

Plan 3 exit requires passing **both** Tier 1 (convergence) and a subset of Tier 2 (behavioural — escape Flat-trap). Tier 3 profitability is Plan 5's job.

- [ ] **Step 10.1: Extend metric-bands.toml**

Append bands for new ISV slots that actually landed (names match `gpu_dqn_trainer.rs` constants). Remove the stale `isv_plan_params_0_ema` from the original — that slot name was never implemented; the B.4 task landed as `isv_readiness_ema` and `isv_plan_threshold`.

```toml
[metric_bands.isv_trade_attempt_rate_ema]
warn_low = 0.001
warn_high = 0.3
error_low = 0.0
error_high = 1.0

[metric_bands.isv_trade_target_rate]
warn_low = 0.001
warn_high = 0.3
error_low = 0.0
error_high = 1.0

[metric_bands.isv_readiness_ema]               # Task 4 B.4 (if landed)
warn_low = 0.3
warn_high = 0.95
error_low = 0.1
error_high = 1.0

[metric_bands.isv_plan_threshold_eff]          # Task 4 B.4 — derived slot 49
warn_low = 0.1
warn_high = 0.5
error_low = 0.05
error_high = 0.7

[metric_bands.isv_state_kl_amplification]      # Task 7 C.3 (if landed)
warn_low = 1.0
warn_high = 2.0
error_low = 0.95
error_high = 3.0

[metric_bands.isv_reward_bonus_ema]            # Task 1 C.2 slot [68]; Tasks 3 B.2 + 5 C.4 populate it
warn_low = 0.0
warn_high = 5.0
error_low = -1.0
error_high = 20.0

[metric_bands.isv_reward_opp_cost_ema]         # Task 1 C.2 slot [67]; Task 2 B.1 populates
warn_low = 0.0
warn_high = 2.0
error_low = -1.0
error_high = 10.0

[metric_bands.val_active_frac]
warn_low = 0.1
warn_high = 0.95
error_low = 0.05
error_high = 1.0

[metric_bands.val_trades_per_bars]
warn_low = 0.005  # 0.5% bars positioned
warn_high = 0.5
error_low = 0.001
error_high = 1.0
```

- [ ] **Step 10.2: Run full smoke suite 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/plan3-smoke.log
```

Expected: all smoke tests PASS. Note that the original plan's "9 new tests from Tasks 1–9" assumed bespoke test helpers that were not built; actual validation is via `multi_fold_convergence` completing cleanly + HEALTH_DIAG showing expected bond movements on each ISV slot.

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

```bash
# Required by Invariant 6.
git push origin main  # per feedback_push_before_deploy.md
./scripts/argo-train.sh --multi-seed 5 --folds 6 --commit $(git rev-parse HEAD) --tag plan3-exit
```

- [ ] **Step 10.4: Gather cross-seed results**

```bash
./scripts/gather-multi-seed-metrics.sh plan3-exit > /tmp/plan3-exit-metrics.json
```

- [ ] **Step 10.5: Verify Tier 1 exit criteria**

```bash
python3 scripts/check-tier1-exit.py /tmp/plan3-exit-metrics.json --warmup-end 15
```

Expected output:
- `std/mean on stable epochs < 0.15` ✓
- `no fold-1 grad explosions detected` ✓
- `no Q-saturation at abs_half` ✓
- `no hot-path memcpy DtoH detected` ✓

- [ ] **Step 10.6: Verify Tier 2 (subset) exit criteria**

```bash
python3 scripts/check-tier2-exit.py /tmp/plan3-exit-metrics.json --subset behavioural
```

Expected:
- `val trades per bar >= 0.005 (was 0.00010)` ✓
- `val_active_frac > 0.2 (was 0.000)` ✓
- `direction-argmax entropy > 0.8 × uniform_max` ✓

- [ ] **Step 10.7: Document results**

Append to `docs/dqn-v2-plan-3-validation.md`:
- Commit SHA.
- Per-seed metrics.
- Per-fold metrics.
- Before-vs-after comparison (Plan 2 baseline vs. Plan 3).
- HEALTH_DIAG sample lines showing `reward_split[..]`, `temporal_reward[..]`, controller diagnostics.

- [ ] **Step 10.8: Commit final validation**

```bash
git add config/metric-bands.toml docs/dqn-v2-plan-3-validation.md
git commit -m "plan3(task10): Tier 1 + Tier 2 (behavioural) validation passes — val Flat-trap escaped"
```

---

## Plan 3 Exit Criteria Summary

Progress as of 2026-04-24 post-revision:

- [x] All 9 invariants preserved across every landed commit (Tasks 1, 2, 3, 5).
- [x] ISV slots added so far, all tail-appended, all in `docs/isv-slots.md`: [63..69) reward EMAs (T1), [71,72] trade novelty (T3). PS_STRIDE 38→39 with peak-PnL-bar (T5).
- [ ] Remaining ISV slots (T4: READINESS_EMA + PLAN_THRESHOLD upgrade; T7: state-KL; T8: SEED_FRACTION + SEED_PRIORITY_SCALE; T9: CQL_ALPHA upgrade) land at tail; final `ISV_TOTAL_DIM ≤ 82`.
- [ ] 4 scripted policies (uniform, momentum, mean_reversion, vwap_deviation) wired into replay seed via 3 new trainer helper methods (`sample_feature_window`, `simulate_scripted_transition`, `replay_insert`).
- [x] AdaptiveMonitor (read-only) pattern: `RewardComponentMonitor` registered (T1). Remaining: `PlanThresholdMonitor` (T4), `StateKLMonitor` (T7), `CqlAlphaMonitor` (T9) — **no CPU-compute controllers**; GPU kernels produce, monitors observe only.
- [x] HEALTH_DIAG `reward_split[..]` line live. Remaining: extend with `plan_threshold_eff`, `state_kl_amp`, `cql_alpha_eff`.
- [ ] Multi-seed × multi-fold Tier 1 + Tier 2 (behavioural subset) exit pass (Task 10).
- [x] No stubs, no TODO/FIXME, no deferred work in landed commits (Invariant 9).
- [x] Wire-up audit has zero orphans added by landed tasks (`reward_component_ema`, `trade_rate_ema` both Wired). Ongoing requirement for remaining tasks.
- [x] Hot-path audit clean on landed tasks: no atomicAdd, no DtoH in kernel reads (Invariant 3).
- [x] No worktrees needed — all implementation happens directly on main with coder agents dispatched serially (per `feedback_no_concurrent_agents_shared_tree.md`).

**Plan 4 prerequisite:** Plan 3 landed on main with all exit criteria satisfied, OR with explicit documented deferrals for Tasks 4/6/7/8/9/10 if the user chooses to move to Plan 4 earlier (e.g., if landed Tasks 1+2+3+5 already demonstrate Flat-trap escape in the Argo validation run).
