diff --git a/docs/superpowers/plans/2026-04-24-dqn-v2-plan-3-behavioural.md b/docs/superpowers/plans/2026-04-24-dqn-v2-plan-3-behavioural.md
index 2e3401eac..6486b2a43 100644
--- a/docs/superpowers/plans/2026-04-24-dqn-v2-plan-3-behavioural.md
+++ b/docs/superpowers/plans/2026-04-24-dqn-v2-plan-3-behavioural.md
@@ -4,20 +4,37 @@
 
 **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).
 
-**Architecture:** ~15 new ISV slots appended AT TAIL after current `ISV_TOTAL_DIM = 63` (post-Plan-2). Grouped by Part: C.2 = 6 slots (reward-component EMAs), B = 3 slots (Flat-scale, trade-attempt, trade-target), B.4 = 1 slot (plan-threshold-eff), D.4 = 3 slots (temporal reward EMAs), C.3 = 3 slots (state-KL). Layout fingerprint re-tails automatically; final `ISV_TOTAL_DIM ≈ 78`.
+**Status (2026-04-24 after 4 landed tasks):** `ISV_TOTAL_DIM = 75` (up from 63 pre-Plan-3 via Tasks 1+3), `PS_STRIDE = 39` (up from 38 via Task 5). Fingerprint at ISV [73,74].
+
+| 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 | ⏳ pending | — | revise-and-execute; reuse `PLAN_THRESHOLD_INDEX=49` |
+| 5 — C.4 temporal timing bonus | ✅ landed | `a0abc3da3` | PS[38] PEAK_PNL_BAR; shaping×bars_early/hold×\|pnl\|×conviction on exit |
+| 6 — D.4 generalised temporal coupling | ⏳ pending | — | revise to tail-append + adaptive coefficients |
+| 7 — C.3 state-KL divergence | ⏳ pending | — | revise slot allocation (original [58,63,64] collides with Task 1 [63..69)) |
+| 8 — B.3 replay seeded warm-start | ⏳ pending | — | revise with concrete helper prereqs + ISV-driven priority scale |
+| 9 — C.5 CQL α seed-coupled | ⏳ pending | — | revise to GPU-drives-CPU-reads; adaptive EMA α |
+| 10 — Plan 3 validation run | ⏳ pending | — | revise metric bands to actual ISV slot names |
+
+**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_audit_kernel.cu` (host-readable, pinned), `state_kl_divergence_kernel.cu`.
+**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 2: layout-fingerprint auto-recompute, current `ISV_TOTAL_DIM = 63`, TLOB cuBLAS integration, per-branch γ, quantile atoms. New slots in Plan 3 go AT TAIL after index 60 (TLOB_REGIME_FOCUS_EMA); fingerprint (currently 61-62) shifts to new tail.
-- Plan 2: per-branch γ ISV slots at `[37..41)` already populated; `Q_P05/Q_P95_EMA` at `[41..49)` already populated; layout fingerprint updated (seed bytes cover Plan 2 slots); `PLAN_ISV_REMAINING_FRACTION` at `plan_isv[6]`.
+- 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.
 
 ---
 
@@ -77,10 +94,12 @@ Expected: smoke test passes on current main. Snapshot this log for Step 10 regre
 
 ---
 
-## Task 1: C.2 Reward-component attribution
+## 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
@@ -350,7 +369,9 @@ git commit -m "plan3(task1C): HEALTH_DIAG reward_split[..] line wired, audit com
 
 ---
 
-## Task 2: B.1 ISV-driven continuous Flat opp-cost
+## 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.
 
@@ -430,7 +451,9 @@ 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
+## 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
@@ -604,142 +627,201 @@ git commit -m "plan3(task3): B.2 novelty-weighted trade-attempt bonus via ISV sl
 
 ## Task 4: B.4 Adaptive plan-MLP activation threshold
 
+**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` — 1 new ISV slot
-- Modify: `crates/ml/src/cuda_pipeline/backtest_plan_kernel.cu` — read threshold from ISV
-- Modify: `crates/ml/src/cuda_pipeline/experience_kernels.cu` — same read during training
-- Create: `crates/ml/src/trainers/dqn/adaptive_controllers/plan_mlp_threshold.rs` — controller
-- Test: new smoke `plan_activation_scales_with_mlp_output`
+- 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 slot 49**
+### 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
-// ─── B.4 Adaptive plan-MLP activation threshold ──────────────────────
-// Plan 3 Task 4. Spec §4.B.4.
-pub const PLAN_PARAMS_0_EMA_INDEX: usize = 49;
+/// 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;
 ```
 
-- [ ] **Step 4.2: Write failing test**
+Recompute the fingerprint seed and move `ISV_LAYOUT_FINGERPRINT_LO/HI` from `[73,74]` to `[76,77]`. Grow `ISV_TOTAL_DIM` 75 → 78.
 
-Create `crates/ml/src/trainers/dqn/smoke_tests/plan_activation_threshold.rs`:
+### Step 4.2: Implement the producer kernel
 
-```rust
-#[test]
-#[ignore = "gpu"]
-fn plan_activation_scales_with_mlp_output() {
-    let mut trainer = build_minimal_trainer();
+Create `crates/ml/src/cuda_pipeline/plan_threshold_update_kernel.cu`:
 
-    // Scenario: MLP outputs typically near 0.3, not 2.0.
-    // Hardcoded 0.5 threshold would never fire. Adaptive should.
-    trainer.set_isv_slot(PLAN_PARAMS_0_EMA_INDEX, 0.3);
-    let active_count_adaptive = trainer.simulate_plan_activations_over_bars(1000);
+```cuda
+#include <cuda_runtime.h>
 
-    // Reset and run with hardcoded 0.5 (via test harness flag to disable adaptive).
-    trainer.reset_plan_state();
-    trainer.set_adaptive_plan_threshold(false);
-    let active_count_fixed = trainer.simulate_plan_activations_over_bars(1000);
-
-    assert!(active_count_adaptive > 100,
-            "Adaptive threshold should activate ≥10% of plans (got {})",
-            active_count_adaptive);
-    assert!(active_count_fixed < 10,
-            "Fixed 0.5 threshold verification — should almost never activate when MLP ~0.3 (got {})",
-            active_count_fixed);
+/* 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);
 }
 ```
 
-- [ ] **Step 4.3: Run — expect FAIL**
+Register in `crates/ml/build.rs` alongside `trade_rate_ema_kernel.cu`.
 
-Expected: no adaptive path, both counts similar.
+### Step 4.3: Producer — emit `readiness_per_sample`
 
-- [ ] **Step 4.4: Implement `PlanMlpThresholdController`**
+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:
 
-Create `crates/ml/src/trainers/dqn/adaptive_controllers/plan_mlp_threshold.rs`:
+```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
-use super::AdaptiveMonitor;
-
-pub struct PlanMlpThresholdController {
-    alpha: f32,
-    plan_params_0_ema_idx: usize,
+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(())
 }
+```
 
-impl PlanMlpThresholdController {
-    pub fn new(plan_params_0_ema_idx: usize) -> Self {
-        Self { alpha: 0.02, plan_params_0_ema_idx }
+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)
     }
-
-    pub fn update_ema(&mut self, samples: &[f32], isv: &mut [f32]) {
-        let mean = samples.iter().copied().sum::<f32>() / samples.len().max(1) as f32;
-        let prev = isv[self.plan_params_0_ema_idx];
-        isv[self.plan_params_0_ema_idx] = (1.0 - self.alpha) * prev + self.alpha * mean;
-    }
-
-    /// Threshold consumed by both training (experience_kernels.cu) and
-    /// val (backtest_plan_kernel.cu). Read-side is always from ISV slot.
-    pub fn threshold_from_isv(isv: &[f32], plan_params_0_ema_idx: usize) -> f32 {
-        // Floor: if EMA is near zero (cold start), fall back to a hard floor of
-        // 0.1 so we don't activate on arbitrary noise before EMA converges.
-        let ema = isv[plan_params_0_ema_idx];
-        let threshold = 0.5 * ema;
-        threshold.max(0.1)
-    }
-}
-
-impl AdaptiveMonitor for PlanMlpThresholdController {
-    type Signal = PlanParamsBatchSummary;
-    type Control = f32;  // the threshold
-
-    fn read_signals(&self, _isv: &IsvBus) -> Self::Signal {
-        // Filled by the caller with the batch of plan_params[0] values.
-        PlanParamsBatchSummary::empty()
-    }
-
-    fn update(&mut self, signals: Self::Signal) -> Self::Control {
-        // update_ema is the real work; this wrapper is for harness consistency.
-        signals.mean
-    }
-
+    fn observe(&mut self, _value: f32) {}  /* no internal state */
     fn fire_rate(&self) -> FireRateStats { FireRateStats::always() }
-    fn diagnose(&self) -> DiagSnapshot {
-        DiagSnapshot::single("plan_mlp_threshold.ema_alpha", self.alpha as f64)
+    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 StateResetRegistry as `SoftReset(5000)` — Plan 2 Task 4 extended the registry with SoftReset.
+Register in `monitors/mod.rs` and in the `controller_activity` smoke harness.
 
-- [ ] **Step 4.5: Wire producer in training_loop.rs**
+### Step 4.7: Registry entries
 
-After the plan-params MLP call in the training step:
+In `state_reset_registry.rs`, after the existing `isv_plan_threshold` entry:
 
 ```rust
-let plan_params_0_batch = extract_plan_params_column_0(&plan_params_buf, n_samples);
-self.plan_mlp_threshold_controller.update_ema(&plan_params_0_batch, self.isv_host_mut());
+/* 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",
+},
 ```
 
-- [ ] **Step 4.6: Wire consumer in backtest_plan_kernel.cu**
+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`.
 
-Change:
-```cuda
-const float plan_activation_threshold = 0.5f;  // hardcoded
-if (plan_state[0] > plan_activation_threshold) { ... }
+### 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
 ```
 
-to:
+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.
 
-```cuda
-// B.4: adaptive threshold from ISV.
-const float ema = isv[plan_params_0_ema_idx];
-const float plan_activation_threshold = fmaxf(0.1f, 0.5f * ema);
-if (plan_state[0] > plan_activation_threshold) { ... }
-```
+### Step 4.9: Update audits
 
-Thread `plan_params_0_ema_idx` kernel parameter. Apply the same change to `experience_kernels.cu`.
-
-- [ ] **Step 4.7: Run test — expect PASS**
-
-- [ ] **Step 4.8: Update audits and commit**
+- `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} \
@@ -753,7 +835,9 @@ git commit -m "plan3(task4): B.4 adaptive plan-MLP activation threshold via ISV[
 
 ---
 
-## Task 5: C.4 Temporal timing bonus
+## 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
@@ -847,114 +931,109 @@ git commit -m "plan3(task5): C.4 temporal timing bonus on trade exit (peak_bar p
 
 ## 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 (held through drawdown, took profit)
+
 **Files:**
-- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — 3 new ISV slot constants
-- Modify: `crates/ml/src/cuda_pipeline/experience_kernels.cu` — 3 new reward components (persistence, regime-shift, consistency)
-- Create: `crates/ml/src/cuda_pipeline/temporal_reward_ema_kernel.cu` — producer for the 3 ISV slots
-- Test: `temporal_reward_persistence_credit`, `temporal_reward_regime_shift_penalty`, `temporal_reward_conviction_consistency`
+- 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.
 
-- [ ] **Step 6.1: Allocate slots 59, 60, 61**
+**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.
 
-```rust
-// ─── D.4 Generalised temporal reward coupling ────────────────────────
-// Plan 3 Task 6. Spec §4.D.4.
-pub const TEMPORAL_REWARD_PERSIST_EMA_INDEX:      usize = 59;
-pub const TEMPORAL_REWARD_REGIME_SHIFT_EMA_INDEX: usize = 60;
-pub const TEMPORAL_REWARD_CONSISTENCY_EMA_INDEX:  usize = 61;
-```
+### Task 6b: Regime-shift penalty (held past regime flip)
 
-Bump `ISV_TOTAL_DIM` to `62` if it was at `58`.
+**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`.
 
-- [ ] **Step 6.2: Write the 3 failing tests**
+**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).
 
-Create `crates/ml/src/trainers/dqn/smoke_tests/temporal_reward_coupling.rs`:
+### Task 6c: Conviction consistency bonus (stable conviction before entry)
 
-```rust
-#[test]
-#[ignore = "gpu"]
-fn temporal_reward_persistence_credit_rewards_noise_ride_through() {
-    // Scenario: trade entered long at 100, dipped to 99 (noise), recovered to
-    // 102, exited +2. Persistence credit should fire.
-    let bonus = simulate_noisy_profitable_trade();
-    assert!(bonus > 1e-3, "persistence credit missing on noise-ride-through");
-}
+**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`.
 
-#[test]
-#[ignore = "gpu"]
-fn temporal_reward_regime_shift_penalty_punishes_late_exit() {
-    // Scenario: entered long on trend regime, regime flipped at bar 3,
-    // policy held until bar 10, took loss. Penalty proportional to 10-3 = 7.
-    let penalty = simulate_regime_shift_late_exit();
-    assert!(penalty < -1e-3, "regime-shift penalty missing");
-}
+**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.
 
-#[test]
-#[ignore = "gpu"]
-fn temporal_reward_conviction_consistency_rewards_stable_pre_entry() {
-    let bonus_stable   = simulate_stable_conviction_entry();
-    let bonus_volatile = simulate_volatile_conviction_entry();
-    assert!(bonus_stable > 5.0 * bonus_volatile);
-}
-```
+### Step 6.X: EMA producer kernel (shared by 6a, 6b, 6c)
 
-- [ ] **Step 6.3: Run — expect FAIL**
+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.
 
-- [ ] **Step 6.4: Implement the 3 components**
+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].
 
-In `experience_kernels.cu`:
+### Step 6.Y: Register PS slots
 
-```cuda
-// D.4.a Persistence credit: held through noise then took profit.
-if (is_trade_exit && memo.final_pnl > 0.0f) {
-    const float drawdown_depth = fmaxf(0.0f, memo.entry_price - memo.min_mark_price);
-    const float persist_bonus  = shaping_scale * conviction_core
-                               * drawdown_depth * tanhf(memo.final_pnl / entry_notional);
-    reward += persist_bonus;
-    // accumulate into per_component_abs_sum[PERSIST_IDX] for Task 1 audit.
-}
+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.
 
-// D.4.b Regime-shift penalty: entered on regime X, regime flipped, held too long.
-if (is_trade_exit && memo.regime_at_entry != memo.regime_at_exit) {
-    const int   bars_since_shift = memo.exit_bar - memo.regime_shift_bar;
-    const float late_penalty     = shaping_scale * conviction_core
-                                 * (float)bars_since_shift * 0.5e-4f
-                                 * fabsf(memo.final_pnl);
-    reward -= late_penalty;
-}
+### Step 6.Z: Validate
 
-// D.4.c Conviction consistency: conviction stable N bars before entry.
-if (is_trade_entry && memo.pre_entry_conviction_std < conviction_std_threshold_from_isv) {
-    const float stability = 1.0f - memo.pre_entry_conviction_std / fmaxf(1e-6f, memo.pre_entry_conviction_mean);
-    const float cons_bonus = shaping_scale * fmaxf(0.0f, stability) * conviction_core;
-    reward += cons_bonus;
-}
-```
+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.
 
-The `conviction_std_threshold_from_isv` is derived at kernel call from `isv[Q_DIR_ABS_REF_INDEX] * 0.1f` — self-scaling, no hardcoded constant.
+### Commit
 
-- [ ] **Step 6.5: Register slots as FoldReset (reset per fold, measured fresh)**
-
-- [ ] **Step 6.6: Wire EMA producer kernel**
-
-Create `temporal_reward_ema_kernel.cu` mirroring `reward_component_audit_kernel.cu` but over the 3 new slots.
-
-- [ ] **Step 6.7: Extend HEALTH_DIAG with temporal_reward_split line**
-
-In `monitoring.rs`:
-
-```rust
-write!(&mut line,
-    " temporal_reward[persist={:.3e} regime_shift={:.3e} consistency={:.3e}]",
-    isv[TEMPORAL_REWARD_PERSIST_EMA_INDEX],
-    isv[TEMPORAL_REWARD_REGIME_SHIFT_EMA_INDEX],
-    isv[TEMPORAL_REWARD_CONSISTENCY_EMA_INDEX])?;
-```
-
-- [ ] **Step 6.8: Run tests, update audits, commit**
+One commit per sub-task. Each commit ships its PS stride bump + all shared-contract consumers in the same commit. Example:
 
 ```bash
-git add [...]
-git commit -m "plan3(task6): D.4 generalised temporal reward coupling via ISV [59..62)"
+git commit -m "feat(dqn-v2): D.4a persistence credit — shaping×conviction×|min_pnl|×tanh(reward/|min_pnl|)"
 ```
 
 ---
@@ -968,144 +1047,114 @@ git commit -m "plan3(task6): D.4 generalised temporal reward coupling via ISV [5
 - Modify: `experience_kernels.cu` — consume STATE_KL_AMPLIFICATION for reward scaling
 - Test: `state_kl_fires_on_distribution_shift`
 
-- [ ] **Step 7.1: Allocate slots 58, 63, 64**
+**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. Spec §4.C.3.
-pub const STATE_KL_TRAIN_VAL_EMA_INDEX: usize = 58;
-pub const STATE_KL_THRESHOLD_EMA_INDEX: usize = 63;
-pub const STATE_KL_AMPLIFICATION_INDEX: usize = 64;
+// 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 */;
 ```
 
-Note: slot `[58]` was already "reserved for C.3" per the spec §6 allocation table. `[63]` and `[64]` were notionally "reserved for Part E" but C.3's description requires them — allocating here. The Part E reservation in spec §6 is advisory; actual allocations are driven by the per-slot producer requirements.
+Grow `ISV_TOTAL_DIM` by +2. Fingerprint re-tails.
 
-Bump `ISV_TOTAL_DIM` to `65`.
+### Step 7.2: Kernel — adaptive coefficients, no tuned multipliers
 
-- [ ] **Step 7.2: Write failing test**
-
-Create `crates/ml/src/trainers/dqn/smoke_tests/state_kl_divergence.rs`:
-
-```rust
-#[test]
-#[ignore = "gpu"]
-fn state_kl_fires_on_distribution_shift() {
-    let mut trainer = build_minimal_trainer();
-    let train_states = build_normal_distribution_states(n = 1000, mean = 0.0, std = 1.0);
-    let val_states   = build_normal_distribution_states(n = 1000, mean = 3.0, std = 1.0);
-
-    trainer.upload_train_states_for_kl(&train_states);
-    trainer.upload_val_states_for_kl(&val_states);
-    trainer.compute_state_kl();
-
-    let kl = trainer.read_isv_slot(STATE_KL_TRAIN_VAL_EMA_INDEX);
-    assert!(kl > 0.1, "KL should fire on mean-shifted distributions, got {}", kl);
-
-    // After ~10 more calls with same shift, AMPLIFICATION should move above 1.0.
-    for _ in 0..10 {
-        trainer.compute_state_kl();
-    }
-    let amp = trainer.read_isv_slot(STATE_KL_AMPLIFICATION_INDEX);
-    assert!(amp > 1.05, "AMPLIFICATION should increase on sustained divergence, got {}", amp);
-}
-
-#[test]
-#[ignore = "gpu"]
-fn state_kl_stays_calm_on_matched_distributions() {
-    let mut trainer = build_minimal_trainer();
-    let train_states = build_normal_distribution_states(n = 1000, mean = 0.0, std = 1.0);
-    let val_states   = build_normal_distribution_states(n = 1000, mean = 0.0, std = 1.0);
-
-    trainer.upload_train_states_for_kl(&train_states);
-    trainer.upload_val_states_for_kl(&val_states);
-    for _ in 0..20 { trainer.compute_state_kl(); }
-
-    let amp = trainer.read_isv_slot(STATE_KL_AMPLIFICATION_INDEX);
-    assert!(amp < 1.05, "AMPLIFICATION should stay near 1.0 on matched dists, got {}", amp);
-}
-```
-
-- [ ] **Step 7.3: Run — expect FAIL**
-
-- [ ] **Step 7.4: Implement state_kl_divergence_kernel.cu**
+Create `crates/ml/src/cuda_pipeline/state_kl_divergence_kernel.cu`:
 
 ```cuda
 #include <cuda_runtime.h>
 
-// Per-dim moment-match KL approximation between train-state and val-state.
-// Computes summed KL over a subset of state dimensions (portfolio + plan + ISV
-// feature block — ranges exposed via `monitored_dims_begin/end` kernel params).
-//
-// Pinned device-mapped writes only. No copy.
-__global__ void state_kl_moment_match(
-    const float* __restrict__ train_state,   // [N_train, D]
-    const float* __restrict__ val_state,     // [N_val, D]
-    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_threshold_ema_idx,
-    int         kl_amplification_idx,
-    float       alpha)
+/* 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)
 {
-    // Compute per-dim mean/var for train and val, then
-    // KL_gauss(train ‖ val) = 0.5 × (log σ_val/σ_train + (σ_train²+μdiff²)/σ_val² − 1)
-    // per-dim, summed over monitored range.
-    //
-    // Single-block kernel (D_monitored ≤ 32). Block reduces via shared mem.
-    // One thread writes EMA updates at end.
-    extern __shared__ float smem[];
-    // Implementation detail: standard block-reduction per dim, thread 0 writes.
-    // ...
+    /* 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   = /* computed KL sum */;
-        const float prev_kl  = isv_out[kl_ema_idx];
-        const float new_kl_ema = (1.0f - alpha) * prev_kl + alpha * new_kl;
-        isv_out[kl_ema_idx] = new_kl_ema;
+        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;
 
-        // Threshold EMA tracks KL EMA itself.
-        const float prev_thresh = isv_out[kl_threshold_ema_idx];
-        isv_out[kl_threshold_ema_idx] = (1.0f - alpha * 0.5f) * prev_thresh + (alpha * 0.5f) * new_kl_ema;
-
-        // Amplification: bump toward 2.0 when KL > 2× threshold; decay toward 1.0 otherwise.
-        const float prev_amp = isv_out[kl_amplification_idx];
-        const float fire     = (new_kl_ema > 2.0f * isv_out[kl_threshold_ema_idx]) ? 1.0f : 0.0f;
-        const float target   = fire * 2.0f + (1.0f - fire) * 1.0f;
-        isv_out[kl_amplification_idx] = prev_amp + 0.02f * (target - prev_amp);
+        /* 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.5: Wire into val backtest (cold-path)**
+### Step 7.3: Wire into val backtest
 
-In `gpu_backtest_evaluator.rs`, per-epoch boundary after val-window pass:
+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.
 
-```rust
-self.launch_state_kl_divergence(
-    &train_state_sample_buf,  // size N_train
-    &val_state_buf,           // N_val
-    self.n_train_kl_samples,
-    self.n_val,
-    STATE_D_TOTAL,
-    SL_PORTFOLIO_START,       // monitored_dims_begin
-    SL_ISV_END,               // monitored_dims_end
-    alpha = 0.05,
-)?;
+### 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;
 ```
 
-- [ ] **Step 7.6: Consume AMPLIFICATION in reward shaping**
+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).
 
-In `experience_kernels.cu`, multiply the Flat-trap escape reward components (B.1 opp-cost, B.2 bonus) by `isv[STATE_KL_AMPLIFICATION_INDEX]` — when KL fires, the model gets more incentive to escape.
+### Step 7.5: Registry + monitor
 
-- [ ] **Step 7.7: Register slots as FoldReset**
+`FoldReset` entries for both slots. Read-only `StateKLMonitor` follows the `PlanThresholdMonitor` template from Task 4.
 
-- [ ] **Step 7.8: Run, update audits, commit**
+### 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 "plan3(task7): C.3 state-distribution KL divergence signal via ISV [58, 63, 64]"
+git commit -m "feat(dqn-v2): C.3 state-KL divergence + adaptive amp ISV slots"
 ```
 
 ---
@@ -1121,239 +1170,233 @@ git commit -m "plan3(task7): C.3 state-distribution KL divergence signal via ISV
 - Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — seed step before epoch 1
 - Test: `replay_seed_fills_buffer_with_diverse_experiences`
 
-- [ ] **Step 8.1: Define scripted-policy trait**
+**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.
 
-Create `crates/ml/src/trainers/dqn/scripted_policies/mod.rs`:
+**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
-pub trait ScriptedPolicy {
+// crates/ml/src/trainers/dqn/scripted_policies/mod.rs
+pub trait ScriptedPolicy: Send + Sync {
     fn name(&self) -> &'static str;
-    /// Returns (factored_action_idx, conviction_in_[0,1]).
+    /// Returns (factored_action_idx ∈ [0, 108), conviction ∈ [0, 1]).
     fn choose_action(&self, state: &StateVector, features: &FeatureWindow) -> (u32, f32);
 }
 ```
 
-- [ ] **Step 8.2: Write failing test**
+Implement `UniformRandomPolicy`, `MomentumPolicy`, `MeanReversionPolicy`, `VwapDeviationPolicy` as separate files. Each is pure-CPU, no mutable state, no allocation on hot path.
 
-Create `crates/ml/src/trainers/dqn/smoke_tests/replay_seed.rs`:
+### Step 8.2: Seed harness
 
 ```rust
-#[test]
-#[ignore = "gpu"]
-fn replay_seed_fills_buffer_with_diverse_experiences() {
-    let mut trainer = build_minimal_trainer();
-    let stats = trainer.seed_replay_buffer(
-        target_experiences = 10_000,
-        policy_portfolio = default_portfolio(),
-    );
-
-    assert_eq!(stats.total_experiences, 10_000);
-
-    // Per-policy counts match 40/20/20/20.
-    assert!((stats.per_policy["uniform"]        - 4000).abs() < 200);
-    assert!((stats.per_policy["momentum"]       - 2000).abs() < 200);
-    assert!((stats.per_policy["mean_reversion"] - 2000).abs() < 200);
-    assert!((stats.per_policy["vwap_deviation"] - 2000).abs() < 200);
-
-    // Action diversity: at least 60 of 108 factored actions observed.
-    assert!(stats.distinct_actions_observed >= 60,
-        "seed portfolio should cover >=60/108 actions, got {}",
-        stats.distinct_actions_observed);
+pub struct ReplaySeedHarness {
+    portfolio: Vec<(Box<dyn ScriptedPolicy>, f32)>,  // (policy, fraction)
 }
-```
 
-- [ ] **Step 8.3: Run — expect FAIL**
-
-- [ ] **Step 8.4: Implement each scripted policy**
-
-Lightweight, pure-CPU. Each policy takes the current state + feature window and returns `(action_idx, conviction)`. Use OFI features and bars already in the state.
-
-Example — `momentum.rs`:
-
-```rust
-pub struct MomentumPolicy { pub lookback_bars: u32 }
-
-impl ScriptedPolicy for MomentumPolicy {
-    fn name(&self) -> &'static str { "momentum" }
-
-    fn choose_action(&self, state: &StateVector, features: &FeatureWindow) -> (u32, f32) {
-        let ret = features.return_over(self.lookback_bars as usize);
-        let conviction = ret.abs().min(1.0) as f32;
-
-        let dir = if ret > 0.001 { DIR_LONG }
-                  else if ret < -0.001 { DIR_SHORT }
-                  else { DIR_HOLD };
-
-        // Fixed magnitude (Half); order market; urgency normal.
-        let action = (dir * 27) + (MAG_HALF * 9) + (ORD_MARKET * 3) + URG_NORMAL;
-        (action as u32, conviction)
-    }
+pub struct SeedStats {
+    pub total_experiences: usize,
+    pub per_policy: std::collections::HashMap<&'static str, usize>,
+    pub distinct_actions: std::collections::BTreeSet<u32>,
 }
-```
-
-Same pattern for mean_reversion (contrarian on short-horizon z-score), vwap_deviation (crossover of close vs. VWAP).
-
-- [ ] **Step 8.5: Implement seed harness**
-
-```rust
-pub struct ReplaySeedHarness { portfolio: Vec<(Box<dyn ScriptedPolicy>, f32)> }
 
 impl ReplaySeedHarness {
-    pub fn run(&self, trainer: &mut DqnTrainer, target: usize) -> SeedStats {
-        let mut stats = SeedStats::default();
-        let mut remaining = target;
+    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 (state, features) = trainer.sample_state_feature_pair();
+                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.step_scripted(&state, action, conviction);
+                let (next_state, reward, done) = trainer.simulate_scripted_transition(bar, action, conviction);
+                trainer.replay_insert(state, action, reward, next_state, done, priority_scale);
 
-                // Insert into PER with priority = 0.5 × baseline (Task 9 seed-decay).
-                trainer.replay_insert_with_priority_scale(
-                    state, action, reward, next_state, done, /*priority_scale=*/ 0.5);
-
-                stats.per_policy.entry(policy.name().to_string()).or_insert(0) += 1;
-                stats.distinct_actions_observed_mut().insert(action);
+                *stats.per_policy.entry(policy.name()).or_insert(0) += 1;
+                stats.distinct_actions.insert(action);
+                stats.total_experiences += 1;
             }
-            remaining -= count;
         }
-        stats.total_experiences = target - remaining;
-        stats
+        Ok(stats)
     }
 }
 ```
 
-- [ ] **Step 8.6: Invoke in training_loop.rs**
+### 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
-// Plan 3 Task 8: B.3 Replay warm-start.
-// Runs unconditionally on fresh training (Invariant: no feature flags).
-// Checkpoint restart implicitly skips because replay buffer is restored
-// from checkpoint already.
 if !self.resumed_from_checkpoint {
-    let seed_target = self.config.replay_seed_experiences;  // 100_000 default
-    let stats = ReplaySeedHarness::new(default_portfolio())
-        .run(self, seed_target);
-    tracing::info!(target: "dqn::B3", "Replay seed complete: {:?}", stats);
+    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());
 }
 ```
 
-The config field `replay_seed_experiences: u32` lands in the same commit. Default 100_000. Cold-path only — this runs once before epoch 1.
+Add `replay_seed_experiences: u32` to `DqnTrainerConfig` with default `100_000`.
 
-- [ ] **Step 8.7: Run test — expect PASS**
+### Step 8.5: Seed-origin vs. dqn-origin tagging (for Task 9)
 
-- [ ] **Step 8.8: Commit**
+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 add crates/ml/src/trainers/dqn/dqn_replay_seed.rs \
-        crates/ml/src/trainers/dqn/scripted_policies/ \
-        crates/ml/src/trainers/dqn/trainer/training_loop.rs \
-        crates/ml/src/trainers/dqn/smoke_tests/replay_seed.rs \
-        docs/dqn-wire-up-audit.md
-git commit -m "plan3(task8): B.3 replay-buffer seeded warm-start (40/20/20/20 portfolio)"
+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
 
+**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/trainers/dqn/adaptive_controllers/cql_alpha_seed_coupled.rs`
-- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs` — use controller
-- Modify: existing `cql_alpha` regime-gate logic to route through the controller
-- Test: `cql_alpha_ramps_as_seed_displaces`
+- 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: Write failing test**
+### Step 9.1: Kernel — GPU producer
 
-```rust
-#[test]
-#[ignore = "gpu"]
-fn cql_alpha_ramps_as_seed_displaces() {
-    let mut trainer = build_minimal_trainer_with_replay_seed(target = 1000);
+```cuda
+#include <cuda_runtime.h>
 
-    // At t=0, seed = 100% of buffer.
-    let alpha_0 = trainer.current_cql_alpha();
-    assert!(alpha_0 < 0.001, "cql_alpha should be ~0 when buffer is 100% seed");
-
-    // Simulate 15 epochs of DQN trajectories displacing seed (500 → 0 seed-origin).
-    trainer.simulate_seed_displacement(final_fraction = 0.0, epochs = 15);
-
-    let alpha_15 = trainer.current_cql_alpha();
-    let final_alpha = trainer.config.cql_alpha_final;  // 0.014
-    assert!((alpha_15 - final_alpha).abs() < 0.001,
-        "cql_alpha should ramp to final {:.4} at t=15, got {:.4}",
-        final_alpha, alpha_15);
+/* 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;
 }
 ```
 
-- [ ] **Step 9.2: Run — expect FAIL**
+Register in `build.rs`.
 
-- [ ] **Step 9.3: Implement `CqlAlphaSeedCoupledController`**
+### Step 9.2: Launcher
+
+In `gpu_experience_collector.rs`:
 
 ```rust
-pub struct CqlAlphaSeedCoupledController {
-    final_alpha:        f32,
-    warmup_time_scale:  f32,  // in step-units
-    step_count:         u64,
-    seed_fraction_ema:  f32,
-}
-
-impl AdaptiveMonitor for CqlAlphaSeedCoupledController {
-    type Signal = f32;  // current seed fraction in replay buffer
-    type Control = f32; // cql_alpha
-
-    fn read_signals(&self, isv: &IsvBus) -> Self::Signal {
-        // Producer: replay buffer tracks seed-origin vs. dqn-origin count
-        // and writes the current seed fraction into ISV slot
-        // SEED_FRACTION_EMA_INDEX (a Plan 1 reserved controller-signal slot).
-        // Here we just read it back.
-        isv.read(SEED_FRACTION_EMA_INDEX)
-    }
-
-    fn update(&mut self, seed_fraction: f32) -> f32 {
-        self.seed_fraction_ema = 0.9 * self.seed_fraction_ema + 0.1 * seed_fraction;
-
-        // Ramp: cql_alpha = final × (1 − seed_fraction_ema) × saturating(t / warmup_time_scale)
-        let saturation = (self.step_count as f32 / self.warmup_time_scale).min(1.0);
-        let displace   = (1.0 - self.seed_fraction_ema).max(0.0);
-        self.step_count += 1;
-
-        self.final_alpha * displace * saturation
-    }
-
-    fn fire_rate(&self) -> FireRateStats { FireRateStats::always() }
-    fn diagnose(&self) -> DiagSnapshot {
-        DiagSnapshot::pairs(&[
-            ("cql_alpha_seed_coupled.seed_frac_ema", self.seed_fraction_ema as f64),
-            ("cql_alpha_seed_coupled.warmup_saturation",
-                (self.step_count as f32 / self.warmup_time_scale).min(1.0) as f64),
-        ])
+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.4: Wire controller**
+### Step 9.3: Training-loop wiring
 
-In `training_loop.rs`, each training step:
+In `training_loop.rs`, in the per-epoch Plan 3 kernel block (same location as `launch_trade_attempt_rate_ema_inplace`):
 
 ```rust
-let seed_frac = self.replay_buffer.seed_fraction();
-let alpha = self.cql_alpha_controller.update(seed_frac);
-self.cuda_config.cql_alpha = alpha;  // consumed by next CQL kernel launch
+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 */
+    )?;
+}
 ```
 
-The existing `cql_alpha regime-gate` from task #7 (Bug #7 fix in `90e1e3dbb`) preserves its null-ISV guard — that guard wraps this controller's output.
+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.5: Run test — expect PASS**
+### Step 9.4: Read-only monitor + registry
 
-- [ ] **Step 9.6: Update audits**
+- `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.7: Commit**
+### Step 9.5: Validate + commit
 
 ```bash
-git commit -m "plan3(task9): C.5 cql_alpha schedule coupled to replay seed-fraction decay"
+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"
 ```
 
 ---
@@ -1368,21 +1411,51 @@ Plan 3 exit requires passing **both** Tier 1 (convergence) and a subset of Tier
 
 - [ ] **Step 10.1: Extend metric-bands.toml**
 
-Append bands for new metrics:
+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_plan_params_0_ema]
-warn_low = 0.05
-warn_high = 10.0
-error_low = 0.01
-error_high = 50.0
-
 [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
@@ -1403,7 +1476,7 @@ 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, including the 9 new tests from Tasks 1–9.
+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**
 
@@ -1462,15 +1535,18 @@ git commit -m "plan3(task10): Tier 1 + Tier 2 (behavioural) validation passes 
 
 ## Plan 3 Exit Criteria Summary
 
-- [ ] All 9 invariants preserved across every commit in this plan.
-- [ ] 15 new ISV slots allocated [49..65), all with named constants, all in `docs/isv-slots.md`.
-- [ ] 4 scripted policies (uniform, momentum, mean_reversion, vwap_deviation) wired into replay seed.
-- [ ] 2 new AdaptiveMonitor impls (`PlanMlpThresholdController`, `CqlAlphaSeedCoupledController`) registered with the controller harness.
-- [ ] HEALTH_DIAG emits `reward_split[..]` + `temporal_reward[..]` lines.
-- [ ] Multi-seed × multi-fold Tier 1 + Tier 2 (behavioural subset) exit pass.
-- [ ] No stubs, no TODO/FIXME, no deferred work (Invariant 9).
-- [ ] Wire-up audit has zero orphans added by this plan.
-- [ ] Hot-path audit shows no new DtoH copies (Invariant 3).
-- [ ] All worktrees cleaned up if any were used during implementation.
+Progress as of 2026-04-24 post-revision:
 
-Plan 4 prerequisite: Plan 3 landed on main with all exit criteria satisfied.
+- [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).