diff --git a/config/training/dqn-hyperopt.toml b/config/training/dqn-hyperopt.toml
index 43b1aac82..22480fdab 100644
--- a/config/training/dqn-hyperopt.toml
+++ b/config/training/dqn-hyperopt.toml
@@ -5,7 +5,7 @@
 [search_space]
 # Base parameters
 learning_rate = [0.00001, 0.0003]       # log scale in adapter
-batch_size = [64, 512]
+batch_size = [512, 8192]
 gamma = [0.90, 0.99]                    # wider range — v_range computed dynamically from gamma
 buffer_size = [50000, 100000]           # log scale in adapter
 max_leverage = [2.0, 10.0]             # leverage ratio; position computed from capital/price
@@ -95,7 +95,7 @@ reward_noise_scale = [0.01, 0.1]
 exposure_aux_weight = [0.1, 1.0]
 
 # v8 search ranges
-micro_reward_scale = [0.0, 0.005]
+micro_reward_scale = [0.0, 0.05]
 td_lambda = [0.5, 0.99]
 hindsight_fraction = [0.0, 0.3]
 hindsight_lookahead = [5, 20]
@@ -118,7 +118,7 @@ q_clip_max = 200.0
 
 [reward]
 # v8 comprehensive training overhaul
-micro_reward_scale = 0.001
+micro_reward_scale = 0.01
 td_lambda = 0.9
 max_trace_length = 7
 hindsight_fraction = 0.1
@@ -162,7 +162,7 @@ branch_hidden_dim = 64
 dueling_hidden_dim = 128
 v_range = 1.0
 # v8 phase_fast overrides
-micro_reward_scale = 0.001
+micro_reward_scale = 0.01
 td_lambda = 0.9
 hindsight_fraction = 0.1
 hindsight_lookahead = 10
diff --git a/config/training/dqn-localdev.toml b/config/training/dqn-localdev.toml
index 42522c08c..62a98524b 100644
--- a/config/training/dqn-localdev.toml
+++ b/config/training/dqn-localdev.toml
@@ -114,7 +114,7 @@ causal_intensity = 1.0
 
 [reward]
 # v8 comprehensive training overhaul
-micro_reward_scale = 0.001
+micro_reward_scale = 0.01
 td_lambda = 0.9
 max_trace_length = 7
 hindsight_fraction = 0.1
diff --git a/config/training/dqn-production.toml b/config/training/dqn-production.toml
index c4ebe5d2f..b6a4d9b1b 100644
--- a/config/training/dqn-production.toml
+++ b/config/training/dqn-production.toml
@@ -125,7 +125,7 @@ causal_intensity = 1.0
 
 [reward]
 # v8 comprehensive training overhaul
-micro_reward_scale = 0.001
+micro_reward_scale = 0.01
 td_lambda = 0.9
 max_trace_length = 7
 hindsight_fraction = 0.1
diff --git a/config/training/dqn-smoketest.toml b/config/training/dqn-smoketest.toml
index 60c0ca473..68931a4f9 100644
--- a/config/training/dqn-smoketest.toml
+++ b/config/training/dqn-smoketest.toml
@@ -121,7 +121,7 @@ causal_intensity = 1.0
 
 [reward]
 # v8 comprehensive training overhaul
-micro_reward_scale = 0.001
+micro_reward_scale = 0.01
 td_lambda = 0.9
 max_trace_length = 7
 hindsight_fraction = 0.1
diff --git a/docs/superpowers/plans/2026-04-07-reward-v8-master.md b/docs/superpowers/plans/2026-04-07-reward-v8-master.md
new file mode 100644
index 000000000..a187989d7
--- /dev/null
+++ b/docs/superpowers/plans/2026-04-07-reward-v8-master.md
@@ -0,0 +1,1359 @@
+# Reward v8: Comprehensive Training Overhaul — 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:** Break through the 30% win rate ceiling to 50%+ on ES futures via 18 improvements across bug fixes, bootstrap trap fixes, initialization, and advanced training techniques. Full CUDA, no CPU path, zero regressions.
+
+**Architecture:** Sequential implementation in 13 tasks. Each task produces a compiling codebase. CUDA kernels compiled via build.rs (nvcc → cubin). Category A (bug fixes) first to establish foundation, then B (bootstrap trap), C (initialization), D (advanced), E (cleanup). All config fields added in Task 1 to avoid repeated struct modification.
+
+**Tech Stack:** CUDA (nvcc → cubin), Rust (cudarc), cuBLAS SGEMM, TOML config, `cargo test`
+
+---
+
+### Task 1: Add ALL New Config Fields + Remove Dead v6 Fields (A6, A9, B1-B3, C1-C2, D1-D4, E1)
+
+**Files:**
+- Modify: `crates/ml/src/trainers/dqn/config.rs`
+- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs`
+- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs`
+- Modify: `crates/ml/src/training_profile.rs`
+
+This task does ALL config changes at once — adding 12 new fields and removing 4 dead fields — so we never revisit these structs.
+
+- [ ] **Step 1: Add 12 new fields to DQNHyperparameters (config.rs)**
+
+Find the v7 gem fields (near `reward_noise_scale`). Add AFTER them:
+
+```rust
+    /// v8: Exposure auxiliary loss base weight (scheduled decay).
+    pub exposure_aux_weight: f64,
+    /// v8: Epochs over which exposure aux weight decays to 10%.
+    pub exposure_aux_warmup_epochs: usize,
+    /// v8: Dense directional micro-reward scale (adaptive vol).
+    pub micro_reward_scale: f64,
+    /// v8: TD(λ) trace parameter. 0=1-step, 1=Monte Carlo.
+    pub td_lambda: f64,
+    /// v8: Maximum trace length for truncated TD(λ).
+    pub max_trace_length: usize,
+    /// v8: PopArt reward normalization enabled.
+    pub popart_enabled: bool,
+    /// v8: Curriculum learning enabled.
+    pub curriculum_enabled: bool,
+    /// v8: Fraction of training before full dataset (curriculum).
+    pub curriculum_warmup_fraction: f64,
+    /// v8: Fraction of replay batch relabeled with hindsight optimal exit.
+    pub hindsight_fraction: f64,
+    /// v8: Lookahead bars for hindsight optimal exit.
+    pub hindsight_lookahead: usize,
+```
+
+Note: `epsilon_start` and `epsilon_end` already exist in DQNHyperparameters. Change their defaults:
+- `epsilon_start`: 1.0 → 0.3
+- `epsilon_end`: 0.01 → 0.02
+
+- [ ] **Step 2: Set defaults in BOTH Default impl and conservative()**
+
+```rust
+            // v8 — comprehensive training overhaul
+            exposure_aux_weight: 0.5,
+            exposure_aux_warmup_epochs: 5,
+            micro_reward_scale: 0.001,
+            td_lambda: 0.9,
+            max_trace_length: 7,
+            popart_enabled: true,
+            curriculum_enabled: true,
+            curriculum_warmup_fraction: 0.6,
+            hindsight_fraction: 0.1,
+            hindsight_lookahead: 10,
+```
+
+Set `n_steps` default from 3 → 5 in BOTH constructors.
+
+- [ ] **Step 3: Update apply_family_scaling**
+
+In the loss_shaping block, add:
+
+```rust
+        self.exposure_aux_weight = (self.exposure_aux_weight * li).clamp(0.0, 2.0);
+        self.micro_reward_scale = (self.micro_reward_scale * li).clamp(0.0, 0.01);
+```
+
+- [ ] **Step 4: Remove dead v6 fields from ExperienceCollectorConfig**
+
+In `gpu_experience_collector.rs`, remove these struct fields:
+- `pub loss_aversion: f32`
+- `pub beta_penalty: f32`
+- `pub regret_blend: f32`
+- `pub trade_clustering_penalty: f32`
+
+Remove from Default impl. Add new v8 fields:
+
+```rust
+    pub micro_reward_scale: f32,
+    pub current_epoch: i32,
+    pub total_epochs: i32,
+    pub eps_start: f32,
+    pub eps_end: f32,
+```
+
+With defaults:
+```rust
+            micro_reward_scale: 0.0,
+            current_epoch: 0,
+            total_epochs: 20,
+            eps_start: 0.3,
+            eps_end: 0.02,
+```
+
+- [ ] **Step 5: Fix training_loop.rs — remove dead refs, add new fields**
+
+Remove references to removed fields:
+```rust
+            // DELETE THESE LINES:
+            // loss_aversion: self.hyperparams.loss_aversion as f32,
+            // beta_penalty: self.hyperparams.beta_penalty_strength as f32,
+            // regret_blend: self.hyperparams.regret_blend as f32,
+            // trade_clustering_penalty: self.hyperparams.trade_clustering_penalty as f32,
+```
+
+Add new field assignments:
+```rust
+            micro_reward_scale: self.hyperparams.micro_reward_scale as f32,
+            current_epoch: self.current_epoch as i32,
+            total_epochs: self.hyperparams.epochs as i32,
+            eps_start: self.hyperparams.epsilon_start as f32,
+            eps_end: self.hyperparams.epsilon_end as f32,
+```
+
+- [ ] **Step 6: Fix training_profile.rs references**
+
+Search for `loss_aversion`, `regret_blend`, `trade_clustering_penalty`, `beta_penalty` in training_profile.rs. For any assignment to `ExperienceCollectorConfig` fields that were removed, delete the line. Keep assignments to `DQNHyperparameters` fields (those are retained for TOML compat).
+
+- [ ] **Step 7: Verify build + run tests**
+
+```bash
+SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
+SQLX_OFFLINE=true cargo test -p ml --lib 2>&1 | grep "^test result:"
+```
+
+Expected: `Finished`, `891 passed; 0 failed` (fix any test expectation failures from default changes).
+
+- [ ] **Step 8: Commit**
+
+```bash
+git add crates/ml/src/trainers/dqn/config.rs \
+       crates/ml/src/cuda_pipeline/gpu_experience_collector.rs \
+       crates/ml/src/trainers/dqn/trainer/training_loop.rs \
+       crates/ml/src/training_profile.rs
+git commit -m "feat(v8): add 12 config fields, remove 4 dead v6 fields, update defaults
+
+New: exposure_aux, micro_reward, td_lambda, popart, curriculum, hindsight
+Removed: loss_aversion, beta_penalty, regret_blend, trade_clustering_penalty
+Default changes: epsilon_start 1.0→0.3, n_steps 3→5"
+```
+
+---
+
+### Task 2: CUDA Kernel Changes — Epsilon Schedule + Micro-Reward + Kernel Sig (B1, B3)
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/experience_kernels.cu`
+- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs`
+
+- [ ] **Step 1: Update action_select kernel for GPU-side cosine epsilon (B1)**
+
+Find `experience_action_select` kernel signature (line ~681). Replace `float epsilon,` with:
+
+```cuda
+    float eps_start,
+    float eps_end,
+    int current_epoch,
+    int total_epochs,
+```
+
+Inside the kernel, BEFORE the first use of `epsilon`, add:
+
+```cuda
+    /* GPU-side cosine epsilon schedule — no CPU scheduling needed */
+    float cosine_progress = (float)current_epoch / fmaxf((float)total_epochs, 1.0f);
+    float epsilon = eps_end + 0.5f * (eps_start - eps_end)
+                  * (1.0f + cosf(3.14159265f * fminf(cosine_progress, 1.0f)));
+```
+
+Also update the PORTFOLIO_STRIDE reference in this kernel from `20` to `23` (line ~715):
+```cuda
+        int ps_base = i * 23;  /* PORTFOLIO_STRIDE = 23 */
+```
+
+- [ ] **Step 2: Add micro-reward + micro_reward_scale to env_step kernel (B3)**
+
+Find the `experience_env_step` kernel signature. Add `float micro_reward_scale,` BEFORE `out_best_exposure`:
+
+```cuda
+    float micro_reward_scale,       /* v8: dense directional micro-reward scale */
+    int*  out_best_exposure,        /* v7.1: counterfactual best exposure idx */
+```
+
+In the kernel body, find the reward label smoothing block. Add the micro-reward BEFORE it (OUTSIDE `segment_complete`, BEFORE drawdown penalty):
+
+```cuda
+    /* ── Dense micro-reward: directional bootstrap signal (v8) ── */
+    if (fabsf(position) > 0.001f && micro_reward_scale > 0.0f) {
+        float micro_vol_proxy = 0.005f;
+        if (features != NULL && bar_idx < total_bars && market_dim > 9) {
+            float atr_n = __bfloat162float(features[(long long)bar_idx * market_dim + 9]);
+            float log_a = atr_n * 16.0f - 7.0f;
+            micro_vol_proxy = fmaxf(expf(log_a) / fmaxf(raw_close, 1.0f), 0.0001f);
+        }
+        float price_change = (raw_next - raw_close) / fmaxf(raw_close, 1.0f);
+        float direction_signal = (position > 0.0f) ? price_change : -price_change;
+        float vol_ratio = micro_vol_proxy / 0.005f;
+        float adaptive_scale = micro_reward_scale / fmaxf(sqrtf(vol_ratio), 0.5f);
+        reward += adaptive_scale * direction_signal / micro_vol_proxy;
+    }
+```
+
+- [ ] **Step 3: Wire new kernel args in gpu_experience_collector.rs**
+
+Find the action_select kernel launch. Replace `.arg(&config.epsilon)` with:
+
+```rust
+                    .arg(&config.eps_start)
+                    .arg(&config.eps_end)
+                    .arg(&config.current_epoch)
+                    .arg(&config.total_epochs)
+```
+
+Find the env_step kernel launch. Add `.arg(&config.micro_reward_scale)` BEFORE `.arg(&mut self.best_exposure_out)`.
+
+- [ ] **Step 4: Verify build**
+
+```bash
+SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
+```
+
+- [ ] **Step 5: Commit**
+
+```bash
+git add crates/ml/src/cuda_pipeline/experience_kernels.cu \
+       crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
+git commit -m "feat(v8): GPU cosine epsilon schedule + dense directional micro-reward"
+```
+
+---
+
+### Task 3: Exposure Aux GEMM Wiring (A6)
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`
+- Modify: `crates/ml/src/trainers/dqn/fused_training.rs`
+
+- [ ] **Step 1: Load kernel + allocate scratch in GpuDqnTrainer**
+
+Add struct fields (near other kernel fields):
+
+```rust
+    exposure_aux_grad_kernel: CudaFunction,
+    exposure_aux_scratch: CudaSlice<f32>,
+```
+
+In the constructor, find where utility kernels are loaded from DQN_UTILITY_CUBIN. Add:
+
+```rust
+    let exposure_aux_grad = module.load_function("exposure_aux_grad_kernel")
+        .map_err(|e| MLError::ModelError(format!("exposure_aux_grad_kernel load: {e}")))?;
+```
+
+Allocate scratch:
+
+```rust
+    let exposure_aux_scratch = stream.alloc_zeros::<f32>(b * config.branch_0_size * config.num_atoms)
+        .map_err(|e| MLError::ModelError(format!("alloc exposure_aux_scratch: {e}")))?;
+```
+
+Store in struct literal:
+
+```rust
+            exposure_aux_grad_kernel: exposure_aux_grad,
+            exposure_aux_scratch,
+```
+
+The function is returned from the utility loader function `load_utility_kernels()` — add it to the return tuple and destructure.
+
+- [ ] **Step 2: Add launch method with backward GEMM**
+
+```rust
+    /// v8: Exposure auxiliary gradient with backward GEMM.
+    /// Phase 1: aux kernel → f32 logit gradients in scratch
+    /// Phase 2: f32→bf16 cast
+    /// Phase 3: backward_fc_layer GEMM → weight + bias gradients in grad_buf
+    pub fn launch_exposure_aux_grad(
+        &mut self,
+        targets: &CudaSlice<i32>,
+        aux_weight: f32,
+    ) -> Result<(), MLError> {
+        if aux_weight < 1e-6 { return Ok(()); }
+        let _evt_guard = EventTrackingGuard::new(self.stream.context());
+
+        let b = self.config.batch_size;
+        let b0 = self.config.branch_0_size;
+        let na = self.config.num_atoms;
+        let ah = self.config.adv_h;
+
+        // Phase 1: Aux kernel → f32 logit gradients in scratch
+        self.stream.memset_zeros(&mut self.exposure_aux_scratch)
+            .map_err(|e| MLError::ModelError(format!("zero exposure_aux_scratch: {e}")))?;
+
+        let adv_logits_ptr = self.on_b_logits_buf.raw_ptr();
+        let targets_ptr = targets.raw_ptr();
+        let scratch_ptr = self.exposure_aux_scratch.raw_ptr();
+
+        let blocks = ((b + 255) / 256) as u32;
+        unsafe {
+            self.stream
+                .launch_builder(&self.exposure_aux_grad_kernel)
+                .arg(&adv_logits_ptr)
+                .arg(&targets_ptr)
+                .arg(&scratch_ptr)
+                .arg(&0i32)
+                .arg(&(b0 as i32))
+                .arg(&(na as i32))
+                .arg(&aux_weight)
+                .arg(&(b as i32))
+                .launch(LaunchConfig {
+                    grid_dim: (blocks, 1, 1),
+                    block_dim: (256, 1, 1),
+                    shared_mem_bytes: 0,
+                })
+                .map_err(|e| MLError::ModelError(format!("exposure_aux_grad kernel: {e}")))?;
+        }
+
+        // Phase 2: f32→bf16 cast for backward GEMM
+        let n_logits = (b * b0 * na) as i32;
+        let cast_blocks = ((n_logits as usize + 255) / 256) as u32;
+        let staging_ptr = self.bw_dy_bf16_staging.raw_ptr();
+        unsafe {
+            self.stream
+                .launch_builder(&self.f32_to_bf16_kernel)
+                .arg(&scratch_ptr)
+                .arg(&staging_ptr)
+                .arg(&n_logits)
+                .launch(LaunchConfig {
+                    grid_dim: (cast_blocks, 1, 1),
+                    block_dim: (256, 1, 1),
+                    shared_mem_bytes: 0,
+                })
+                .map_err(|e| MLError::ModelError(format!("f32_to_bf16 exposure_aux: {e}")))?;
+        }
+
+        // Phase 3: backward_fc_layer GEMM
+        let sizes = compute_param_sizes(&self.config);
+        let f32_bytes = std::mem::size_of::<f32>() as u64;
+        let goff_w_b0out: u64 = sizes[..10].iter().sum::<usize>() as u64 * f32_bytes;
+        let goff_b_b0out: u64 = sizes[..11].iter().sum::<usize>() as u64 * f32_bytes;
+        let grad_base = self.grad_buf.raw_ptr();
+        let h_b0_ptr = self.save_h_b0.raw_ptr();
+
+        self.cublas_backward.backward_fc_layer(
+            &self.stream,
+            staging_ptr,
+            h_b0_ptr,
+            0u64,  // W — not needed (dx=0)
+            grad_base + goff_w_b0out,
+            grad_base + goff_b_b0out,
+            0u64,  // dx = 0 (no upstream propagation)
+            b0 * na,
+            ah,
+            b,
+        )?;
+
+        Ok(())
+    }
+```
+
+- [ ] **Step 3: Wire in FusedTrainingCtx**
+
+Add fields:
+
+```rust
+    pub(crate) exposure_aux_weight: f64,
+    pub(crate) exposure_targets: CudaSlice<i32>,
+```
+
+Initialize in constructor:
+
+```rust
+            exposure_aux_weight: hyperparams.exposure_aux_weight,
+            exposure_targets: stream.alloc_zeros::<i32>(batch_size)
+                .map_err(|e| anyhow::anyhow!("alloc exposure_targets: {e}"))?,
+```
+
+In `run_full_step()`, find the section between auxiliary ops and pruning/Adam (after CQL gradient, before `apply_pruning_mask`). Add:
+
+```rust
+        // v8: Exposure auxiliary gradient (breaks i%3 degeneracy)
+        if self.exposure_aux_weight > 1e-6 {
+            if let Err(e) = self.trainer.launch_exposure_aux_grad(
+                &self.exposure_targets,
+                self.exposure_aux_weight as f32,
+            ) {
+                tracing::warn!("Exposure aux grad failed (non-fatal): {e}");
+            }
+        }
+```
+
+- [ ] **Step 4: Verify build**
+
+```bash
+SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
+```
+
+- [ ] **Step 5: Commit**
+
+```bash
+git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \
+       crates/ml/src/trainers/dqn/fused_training.rs
+git commit -m "feat(v8): wire exposure aux GEMM — kernel + f32→bf16 cast + backward_fc_layer"
+```
+
+---
+
+### Task 4: Training Loop Schedules (A7, A8, C3)
+
+**Files:**
+- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs`
+
+- [ ] **Step 1: CEA 25% linear warmup (A7)**
+
+Find `cea_weight:` in the ExperienceCollectorConfig constructor. Replace with:
+
+```rust
+            cea_weight: {
+                let base = self.hyperparams.cea_weight as f32;
+                let warmup_end = (self.hyperparams.epochs as f32 * 0.25).max(3.0);
+                if (self.current_epoch as f32) < warmup_end {
+                    let progress = self.current_epoch as f32 / warmup_end;
+                    1.0_f32 * (1.0 - progress) + base * progress
+                } else { base }
+            },
+```
+
+- [ ] **Step 2: OFI epoch gate (A8)**
+
+Find `ofi_reward_weight:`. Replace with:
+
+```rust
+            ofi_reward_weight: {
+                let base = self.hyperparams.ofi_reward_weight as f32;
+                if self.current_epoch < 5 { 0.0 } else { base }
+            },
+```
+
+- [ ] **Step 3: Exposure aux weight decay (C3)**
+
+After the ExperienceCollectorConfig is built, add:
+
+```rust
+            // v8: Exposure aux weight decay schedule
+            if let Some(ref mut fused) = self.fused_ctx {
+                let base = self.hyperparams.exposure_aux_weight;
+                let warmup = self.hyperparams.exposure_aux_warmup_epochs.max(1) as f64;
+                let progress = (self.current_epoch as f64 / warmup).min(1.0);
+                fused.exposure_aux_weight = base * (1.0 - 0.9 * progress);
+            }
+```
+
+- [ ] **Step 4: Verify build**
+
+```bash
+SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
+```
+
+- [ ] **Step 5: Commit**
+
+```bash
+git add crates/ml/src/trainers/dqn/trainer/training_loop.rs
+git commit -m "feat(v8): CEA 25% warmup, OFI epoch gate, exposure aux decay schedule"
+```
+
+---
+
+### Task 5: TD(λ) Kernel (D1)
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/nstep_kernel.cu`
+- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs`
+
+- [ ] **Step 1: Add td_lambda_kernel to nstep_kernel.cu**
+
+Append after the existing `nstep_accumulate_kernel`:
+
+```cuda
+/**
+ * TD(λ) truncated return kernel — strict generalization of n-step.
+ *
+ * Computes G^λ = (1-λ) Σ_{n=1}^{N-1} λ^{n-1} G^(n) + λ^{N-1} G^(N)
+ * where G^(n) = r_t + γr_{t+1} + ... + γ^{n-1}r_{t+n-1} + γ^n Q(s_{t+n})
+ *
+ * When λ=0: reduces to 1-step TD.
+ * When λ=1: reduces to Monte Carlo (N-step return, no intermediate bootstrap).
+ *
+ * Requires bootstrapped Q(s') values in q_next buffer.
+ *
+ * Launch: grid=(ceil(N*L/256)), block=(256).
+ */
+extern "C" __global__ void td_lambda_kernel(
+    const float* __restrict__ raw_rewards,   /* [N * L] original 1-step rewards */
+    const float* __restrict__ raw_dones,     /* [N * L] original 1-step dones */
+    const float* __restrict__ q_next,        /* [N * L] bootstrapped Q(s_{t+1}) */
+    float*       __restrict__ out_rewards,   /* [N * L] overwritten with G^λ */
+    float*       __restrict__ out_dones,     /* [N * L] overwritten with done_n */
+    float gamma,
+    float lambda,
+    int max_trace,   /* N in the formula, typically 5-10 */
+    int L,           /* timesteps per episode */
+    int N            /* number of episodes */
+) {
+    int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    if (idx >= N * L) return;
+
+    int ep = idx / L;
+    int t  = idx % L;
+    int base = ep * L;
+
+    float G_lambda = 0.0f;
+    float weight_sum = 0.0f;
+    float lambda_pow = 1.0f;
+    float any_done = 0.0f;
+
+    for (int n = 1; n <= max_trace; n++) {
+        /* Compute n-step return G^(n) */
+        float G_n = 0.0f;
+        float gamma_pow = 1.0f;
+        int valid = 1;
+
+        for (int k = 0; k < n && (t + k) < L; k++) {
+            int off = base + t + k;
+            float d_k = raw_dones[off];
+            if (d_k > 0.5f && k > 0) { valid = 0; any_done = 1.0f; break; }
+            G_n += gamma_pow * raw_rewards[off];
+            gamma_pow *= gamma;
+            if (d_k > 0.5f) { any_done = 1.0f; break; }
+        }
+
+        if (!valid || (t + n) >= L) {
+            /* Terminal or out of bounds — use partial return */
+            G_lambda += lambda_pow * G_n;
+            weight_sum += lambda_pow;
+            break;
+        }
+
+        /* Add bootstrap: G^(n) += γ^n * Q(s_{t+n}) */
+        G_n += gamma_pow * q_next[base + t + n - 1];
+
+        /* Accumulate: w_n = (1-λ)λ^{n-1} for n < N, λ^{N-1} for n = N */
+        float w = (n < max_trace) ? (1.0f - lambda) * lambda_pow : lambda_pow;
+        G_lambda += w * G_n;
+        weight_sum += w;
+        lambda_pow *= lambda;
+    }
+
+    int out_idx = base + t;
+    out_rewards[out_idx] = (weight_sum > 1e-8f) ? G_lambda / weight_sum : raw_rewards[out_idx];
+    out_dones[out_idx] = any_done;
+}
+```
+
+- [ ] **Step 2: Wire TD(λ) kernel launch in gpu_experience_collector.rs**
+
+Find where `nstep_accumulate_kernel` is loaded and launched. The kernel is loaded from a cubin. Add the new kernel load:
+
+```rust
+    let td_lambda_fn = nstep_module.load_function("td_lambda_kernel")
+        .map_err(|e| MLError::ModelError(format!("td_lambda_kernel load: {e}")))?;
+```
+
+Store as a field. Find the nstep launch site (search for `nstep_accumulate_kernel` or `nstep_fn`). The launch needs the additional `q_next` buffer. If the collector doesn't have Q(s') bootstrapped values readily available, use the existing n-step kernel as fallback and gate TD(λ) behind a config flag.
+
+For now, launch TD(λ) by passing `rewards_out` as `q_next` (self-bootstrap approximation that improves as Q-values improve):
+
+```rust
+    // TD(λ) — uses rewards as initial Q bootstrap approximation
+    // In practice, Q(s') comes from the target network forward pass
+    let lambda = self.config.td_lambda;
+    let max_trace = self.config.max_trace_length;
+    if lambda > 0.01 && max_trace > 1 {
+        // Launch td_lambda_kernel instead of nstep_accumulate_kernel
+        // ... (same launch pattern, additional args: q_next, lambda, max_trace)
+    }
+```
+
+- [ ] **Step 3: Verify build**
+
+```bash
+SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
+```
+
+- [ ] **Step 4: Commit**
+
+```bash
+git add crates/ml/src/cuda_pipeline/nstep_kernel.cu \
+       crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
+git commit -m "feat(v8): TD(λ) truncated return kernel — replaces fixed n-step"
+```
+
+---
+
+### Task 6: Pessimistic Q-Value Initialization (C2)
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`
+
+- [ ] **Step 1: Add pessimistic bias initialization after weight setup**
+
+Find `params_initialized` or the weight initialization section in `GpuDqnTrainer::new()`. After weights are initialized (Xavier/He), add:
+
+```rust
+        // v8: Pessimistic Q-value initialization — shift value head bias to -0.1
+        // Forces exploration: model starts believing all states are slightly negative,
+        // must DISCOVER positive-value states through experience.
+        {
+            let sizes = compute_param_sizes(&config);
+            let b_v2_offset: usize = sizes[..7].iter().sum(); // offset to b_v2 in flat params
+            let pessimistic_val = -0.1_f32;
+            let bias_data = vec![pessimistic_val; config.num_atoms];
+            let bias_bf16: Vec<half::bf16> = bias_data.iter().map(|&v| half::bf16::from_f32(v)).collect();
+
+            // Write to params_f32 master buffer
+            let f32_offset = b_v2_offset * std::mem::size_of::<f32>();
+            unsafe {
+                cudarc::driver::sys::cuMemcpyHtoDAsync_v2(
+                    params_f32.raw_ptr() + f32_offset as u64,
+                    bias_data.as_ptr().cast(),
+                    bias_data.len() * std::mem::size_of::<f32>(),
+                    stream.cu_stream(),
+                );
+            }
+
+            // Write to params_bf16 shadow
+            let bf16_offset = b_v2_offset * std::mem::size_of::<half::bf16>();
+            unsafe {
+                cudarc::driver::sys::cuMemcpyHtoDAsync_v2(
+                    params_bf16.raw_ptr() + bf16_offset as u64,
+                    bias_bf16.as_ptr().cast(),
+                    bias_bf16.len() * std::mem::size_of::<half::bf16>(),
+                    stream.cu_stream(),
+                );
+            }
+            tracing::info!("v8: Pessimistic Q-init: b_v2 set to {pessimistic_val} ({} atoms)", config.num_atoms);
+        }
+```
+
+- [ ] **Step 2: Verify build**
+
+```bash
+SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
+```
+
+- [ ] **Step 3: Commit**
+
+```bash
+git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
+git commit -m "feat(v8): pessimistic Q-value initialization — b_v2 bias set to -0.1"
+```
+
+---
+
+### Task 7: Supervised Pre-Training Kernel (C1)
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu`
+- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`
+- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs`
+
+- [ ] **Step 1: Add exposure_pretrain_step kernel**
+
+Append to `dqn_utility_kernels.cu`:
+
+```cuda
+/* ══════════════════════════════════════════════════════════════════════
+ * SUPERVISED PRE-TRAINING KERNEL (v8 C1)
+ *
+ * Computes cross-entropy logit gradients for the exposure branch
+ * against the optimal single-bar direction. Runs at epoch 0 to
+ * initialize exposure weights with directional knowledge before RL.
+ *
+ * Target: argmax_k(position_k * (raw_next - raw_close))
+ * Same math as exposure_aux_grad_kernel but target computed from
+ * raw prices instead of from the CEA loop.
+ *
+ * Launch: grid=(ceil(batch_size/256)), block=(256).
+ * ══════════════════════════════════════════════════════════════════════ */
+extern "C" __global__ void exposure_pretrain_step(
+    const __nv_bfloat16* __restrict__ targets,  /* [total_bars, 4] raw prices */
+    const int*   __restrict__ bar_indices,       /* [batch] which bar each sample maps to */
+    const float* __restrict__ adv_logits,        /* [batch, b0_size * num_atoms] from forward */
+    float*       __restrict__ grad_scratch,      /* [batch, b0_size * num_atoms] output */
+    int b0_size, int num_atoms, int batch_size,
+    float max_position
+) {
+    int sample = blockIdx.x * blockDim.x + threadIdx.x;
+    if (sample >= batch_size) return;
+
+    int bar = bar_indices[sample];
+
+    /* Compute optimal direction from raw prices */
+    float raw_close = __bfloat162float(targets[bar * 4 + 2]);
+    float raw_next  = __bfloat162float(targets[bar * 4 + 3]);
+    float price_delta = raw_next - raw_close;
+
+    /* Find best exposure: which position maximizes single-bar PnL? */
+    int best_k = b0_size / 2;  /* default: Flat */
+    float best_pnl = 0.0f;     /* Flat PnL = 0 */
+    for (int k = 0; k < b0_size; k++) {
+        float step_size = (b0_size > 1) ? 2.0f / (float)(b0_size - 1) : 0.0f;
+        float fraction = -1.0f + (float)k * step_size;
+        float pos = fraction * max_position;
+        float pnl = pos * price_delta;
+        if (pnl > best_pnl) { best_pnl = pnl; best_k = k; }
+    }
+
+    /* Cross-entropy gradient (same math as exposure_aux_grad_kernel) */
+    const float* logits = adv_logits + (long long)sample * b0_size * num_atoms;
+    float inv_batch = 1.0f / (float)batch_size;
+
+    for (int z = 0; z < num_atoms; z++) {
+        float max_l = -1e30f;
+        for (int a = 0; a < b0_size; a++) {
+            float l = logits[a * num_atoms + z];
+            if (l > max_l) max_l = l;
+        }
+        float sum_exp = 0.0f;
+        for (int a = 0; a < b0_size; a++) {
+            sum_exp += expf(logits[a * num_atoms + z] - max_l);
+        }
+        float inv_sum = 1.0f / fmaxf(sum_exp, 1e-10f);
+
+        for (int a = 0; a < b0_size; a++) {
+            float prob = expf(logits[a * num_atoms + z] - max_l) * inv_sum;
+            float grad = (prob - ((a == best_k) ? 1.0f : 0.0f)) * inv_batch;
+            grad_scratch[(long long)sample * b0_size * num_atoms + a * num_atoms + z] = grad;
+        }
+    }
+}
+```
+
+- [ ] **Step 2: Load kernel + add pre-training method to GpuDqnTrainer**
+
+Load from DQN_UTILITY_CUBIN alongside `exposure_aux_grad_kernel`. Add a `run_pretrain_step()` method that:
+1. Launches `exposure_pretrain_step` with random bar indices
+2. Does f32→bf16 cast on the scratch output
+3. Calls `backward_fc_layer` for weight + bias gradients
+4. Returns Ok(())
+
+The method follows the same pattern as `launch_exposure_aux_grad` but with different inputs.
+
+- [ ] **Step 3: Wire pre-training phase in training_loop.rs**
+
+Find the epoch loop start. Add BEFORE epoch 0:
+
+```rust
+        // v8: Supervised pre-training phase (epoch 0 only)
+        if self.hyperparams.exposure_aux_weight > 0.0 {
+            info!("v8: Running supervised pre-training (100 batches)...");
+            if let Some(ref mut fused) = self.fused_ctx {
+                for batch_idx in 0..100 {
+                    // Sample random bar indices, run forward, launch pretrain kernel + backward GEMM + Adam
+                    // ... (uses fused.trainer.run_pretrain_step() + replay_adam())
+                }
+                info!("v8: Pre-training complete — exposure branch initialized with directional knowledge");
+            }
+        }
+```
+
+This is a substantial integration that requires careful wiring. The implementer should study how `train_step_gpu()` works and replicate a simplified version for pre-training.
+
+- [ ] **Step 4: Verify build + commit**
+
+```bash
+SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
+git add crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu \
+       crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \
+       crates/ml/src/trainers/dqn/trainer/training_loop.rs
+git commit -m "feat(v8): supervised pre-training kernel + epoch 0 direction initialization"
+```
+
+---
+
+### Task 8: PopArt Reward Normalization (D2)
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu`
+- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`
+- Modify: `crates/ml/src/trainers/dqn/fused_training.rs`
+
+- [ ] **Step 1: Add PopArt normalize kernel**
+
+Append to `dqn_utility_kernels.cu`:
+
+```cuda
+/* ══════════════════════════════════════════════════════════════════════
+ * POPART REWARD NORMALIZATION KERNEL (v8 D2)
+ *
+ * Normalizes rewards in-place using running mean/variance (Welford).
+ * Also updates running statistics from the current batch.
+ * Eliminates manual v_min/v_max tuning for C51.
+ *
+ * Launch: grid=(ceil(N/256)), block=(256). Thread 0 updates stats.
+ * ══════════════════════════════════════════════════════════════════════ */
+extern "C" __global__ void popart_normalize_kernel(
+    float* __restrict__ rewards,       /* [N] in-place normalize */
+    float* __restrict__ mean_buf,      /* [1] running mean */
+    float* __restrict__ var_buf,       /* [1] running M2 (Welford) */
+    float* __restrict__ count_buf,     /* [1] sample count */
+    int N, int warmup_count
+) {
+    /* Thread 0: update running stats with batch mean */
+    __shared__ float batch_sum;
+    __shared__ float batch_sq_sum;
+
+    int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+    /* Compute batch statistics via parallel reduction */
+    float local_sum = (i < N) ? rewards[i] : 0.0f;
+    float local_sq  = (i < N) ? rewards[i] * rewards[i] : 0.0f;
+
+    /* Warp reduce */
+    for (int offset = 16; offset > 0; offset >>= 1) {
+        local_sum += __shfl_xor_sync(0xFFFFFFFF, local_sum, offset);
+        local_sq  += __shfl_xor_sync(0xFFFFFFFF, local_sq, offset);
+    }
+
+    if (threadIdx.x == 0) {
+        batch_sum = local_sum;
+        batch_sq_sum = local_sq;
+    }
+    __syncthreads();
+
+    /* Thread 0: update running stats */
+    if (threadIdx.x == 0 && blockIdx.x == 0) {
+        float count = *count_buf + (float)N;
+        float new_mean = (*mean_buf * *count_buf + batch_sum) / fmaxf(count, 1.0f);
+        float new_var = (*var_buf * *count_buf + batch_sq_sum) / fmaxf(count, 1.0f)
+                      - new_mean * new_mean;
+        new_var = fmaxf(new_var, 0.0001f);
+        *mean_buf = new_mean;
+        *var_buf = new_var;
+        *count_buf = count;
+    }
+    __syncthreads();
+
+    /* All threads: normalize rewards in-place */
+    float mu = *mean_buf;
+    float sigma = sqrtf(fmaxf(*var_buf, 0.0001f));
+    float total_count = *count_buf;
+
+    if (i < N && total_count >= (float)warmup_count) {
+        rewards[i] = (rewards[i] - mu) / sigma;
+    }
+}
+```
+
+- [ ] **Step 2: Add PopArt buffers and launch in GpuDqnTrainer**
+
+Add fields:
+```rust
+    popart_mean: CudaSlice<f32>,
+    popart_var: CudaSlice<f32>,
+    popart_count: CudaSlice<f32>,
+    popart_normalize_kernel: CudaFunction,
+```
+
+Allocate + load. Add `normalize_rewards_popart(&mut self, rewards, n)` method.
+
+- [ ] **Step 3: Wire in fused training — normalize rewards before C51 loss**
+
+In `run_full_step()`, after batch upload and before C51 forward, call:
+```rust
+        if self.popart_enabled {
+            self.trainer.normalize_rewards_popart(&mut batch.rewards, batch_size)?;
+        }
+```
+
+- [ ] **Step 4: Verify build + commit**
+
+```bash
+SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
+git add crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu \
+       crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \
+       crates/ml/src/trainers/dqn/fused_training.rs
+git commit -m "feat(v8): PopArt reward normalization — adaptive running mean/var + in-place normalize"
+```
+
+---
+
+### Task 9: Curriculum Learning (D3)
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/experience_kernels.cu`
+- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs`
+- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs`
+
+- [ ] **Step 1: Add difficulty scoring kernel**
+
+Append to `experience_kernels.cu`:
+
+```cuda
+/* ══════════════════════════════════════════════════════════════════════
+ * CURRICULUM DIFFICULTY SCORING KERNEL (v8 D3)
+ *
+ * Computes per-bar difficulty: inverse of 5-bar directional clarity.
+ * Low score = strong trend (easy). High score = choppy (hard).
+ * Used to order episode starts from easy → hard across training.
+ *
+ * Launch: grid=(ceil(total_bars/256)), block=(256).
+ * ══════════════════════════════════════════════════════════════════════ */
+extern "C" __global__ void compute_difficulty_scores(
+    const __nv_bfloat16* __restrict__ targets,  /* [total_bars, 4] */
+    float* __restrict__ scores,                  /* [total_bars] output */
+    int total_bars
+) {
+    int i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i >= total_bars - 5) {
+        if (i < total_bars) scores[i] = 1e6f;  /* sentinel for last 5 bars */
+        return;
+    }
+
+    float close_0 = __bfloat162float(targets[i * 4 + 2]);
+    float close_5 = __bfloat162float(targets[(i + 5) * 4 + 2]);
+    float ret_5 = fabsf(close_5 - close_0) / fmaxf(close_0, 1.0f);
+    scores[i] = 1.0f / (ret_5 + 0.001f);
+}
+```
+
+- [ ] **Step 2: Wire scoring + sorted indices in collector**
+
+At init, launch `compute_difficulty_scores`. Download scores (one-time DtoH), sort indices, upload sorted starts. Add `curriculum_starts: CudaSlice<i32>` field.
+
+- [ ] **Step 3: Per-epoch curriculum in training_loop**
+
+```rust
+            if self.hyperparams.curriculum_enabled {
+                let frac = (self.current_epoch as f32 / self.hyperparams.epochs as f32).min(1.0);
+                let warmup = self.hyperparams.curriculum_warmup_fraction as f32;
+                let usable = ((warmup + (1.0 - warmup) * (frac / warmup).min(1.0)) * total_bars as f32) as usize;
+                // Restrict episode_starts to first `usable` of sorted indices
+            }
+```
+
+- [ ] **Step 4: Verify build + commit**
+
+```bash
+SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
+git add crates/ml/src/cuda_pipeline/experience_kernels.cu \
+       crates/ml/src/cuda_pipeline/gpu_experience_collector.rs \
+       crates/ml/src/trainers/dqn/trainer/training_loop.rs
+git commit -m "feat(v8): curriculum learning — difficulty-ordered episode starts"
+```
+
+---
+
+### Task 10: Hindsight Optimal Exit Relabeling (D4)
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/experience_kernels.cu`
+- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`
+- Modify: `crates/ml/src/trainers/dqn/fused_training.rs`
+
+- [ ] **Step 1: Add hindsight_relabel_kernel**
+
+Append to `experience_kernels.cu` (after the difficulty kernel):
+
+```cuda
+/* ══════════════════════════════════════════════════════════════════════
+ * HINDSIGHT OPTIMAL EXIT RELABELING KERNEL (v8 D4)
+ *
+ * For a fraction of replay samples, replaces the reward with the
+ * optimal exit PnL within a lookahead window. Teaches the model
+ * "what was the best possible outcome for this entry?"
+ *
+ * Launch: grid=(ceil(batch_size/256)), block=(256).
+ * ══════════════════════════════════════════════════════════════════════ */
+extern "C" __global__ void hindsight_relabel_kernel(
+    const __nv_bfloat16* __restrict__ targets,
+    const int*   __restrict__ actions,
+    const int*   __restrict__ bar_indices,
+    float*       __restrict__ rewards,
+    int b0_size, int b1_size, int b2_size,
+    float max_position, int lookahead,
+    int total_bars, float hindsight_fraction,
+    int batch_size
+) {
+    int i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i >= batch_size) return;
+
+    /* Deterministic mask: relabel only hindsight_fraction of samples */
+    unsigned int hash = (unsigned int)(i * 48271 + bar_indices[i] * 7919);
+    hash ^= (hash >> 16); hash *= 0x45d9f3bu; hash ^= (hash >> 16);
+    if ((float)(hash & 0xFFFF) / 65535.0f >= hindsight_fraction) return;
+
+    int bar = bar_indices[i];
+    if (bar < 0 || bar + lookahead >= total_bars) return;
+
+    int exposure_idx = actions[i] / (b1_size * b2_size);
+    float step_size = (b0_size > 1) ? 2.0f / (float)(b0_size - 1) : 0.0f;
+    float fraction = -1.0f + (float)exposure_idx * step_size;
+    float pos = fraction * max_position;
+    if (fabsf(pos) < 0.001f) return;
+
+    float entry_price = __bfloat162float(targets[bar * 4 + 2]);
+    if (entry_price < 1.0f) return;
+
+    float best_pnl = 0.0f;
+    for (int k = 1; k <= lookahead && (bar + k) < total_bars; k++) {
+        float future_price = __bfloat162float(targets[(bar + k) * 4 + 2]);
+        float pnl = pos * (future_price - entry_price);
+        if (pnl > best_pnl) best_pnl = pnl;
+    }
+
+    if (best_pnl > 0.0f) {
+        float opt_return = best_pnl / entry_price;
+        float opt_reward = asymmetric_soft_clamp(10.0f * opt_return / 0.005f);
+        rewards[i] = opt_reward;
+    }
+}
+```
+
+- [ ] **Step 2: Wire kernel launch in training pipeline**
+
+Load kernel from experience cubin. In `run_full_step()`, launch AFTER batch upload, BEFORE C51 loss:
+
+```rust
+        if self.hindsight_fraction > 0.0 {
+            self.trainer.launch_hindsight_relabel(
+                &batch.actions, &batch.bar_indices, &mut batch.rewards,
+                self.hindsight_fraction, self.hindsight_lookahead,
+            )?;
+        }
+```
+
+- [ ] **Step 3: Verify build + commit**
+
+```bash
+SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5
+git add crates/ml/src/cuda_pipeline/experience_kernels.cu \
+       crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \
+       crates/ml/src/trainers/dqn/fused_training.rs
+git commit -m "feat(v8): hindsight optimal exit relabeling — 10% of samples get best-case reward"
+```
+
+---
+
+### Task 11: TOML Configs + Hyperopt Ranges
+
+**Files:**
+- Modify: `config/training/dqn-production.toml`
+- Modify: `config/training/dqn-hyperopt.toml`
+- Modify: `config/training/dqn-smoketest.toml`
+- Modify: `config/training/dqn-localdev.toml`
+
+- [ ] **Step 1: Add v8 params to all 4 TOMLs**
+
+Add to `[reward]` section of each:
+
+```toml
+# v8 comprehensive training overhaul
+exposure_aux_weight = 0.5
+exposure_aux_warmup_epochs = 5
+micro_reward_scale = 0.001
+td_lambda = 0.9
+max_trace_length = 7
+popart_enabled = true
+curriculum_enabled = true
+curriculum_warmup_fraction = 0.6
+hindsight_fraction = 0.1
+hindsight_lookahead = 10
+```
+
+Update `epsilon_start = 0.3` and `n_steps = 5` in all profiles.
+
+For `dqn-hyperopt.toml`, add to `[search_space]`:
+
+```toml
+# v8 search ranges
+exposure_aux_weight = [0.1, 1.0]
+micro_reward_scale = [0.0, 0.005]
+td_lambda = [0.5, 0.99]
+max_trace_length = [5, 10]
+curriculum_warmup_fraction = [0.3, 0.8]
+hindsight_fraction = [0.0, 0.3]
+hindsight_lookahead = [5, 20]
+epsilon_start = [0.1, 0.5]
+epsilon_end = [0.01, 0.05]
+n_steps = [3, 7]
+```
+
+- [ ] **Step 2: Commit**
+
+```bash
+git add config/training/
+git commit -m "config(v8): add all v8 params to TOML profiles + hyperopt ranges"
+```
+
+---
+
+### Task 12: Explicit v8 Unit Tests
+
+**Files:**
+- Create: `crates/ml/src/trainers/dqn/smoke_tests/reward_v8.rs`
+- Modify: `crates/ml/src/trainers/dqn/smoke_tests/mod.rs`
+
+- [ ] **Step 1: Create reward_v8.rs test module**
+
+```rust
+//! Reward v8 unit tests — verify mathematical properties of each component.
+//! These are CPU-side tests that validate the CUDA kernel logic would produce
+//! correct results. The actual CUDA kernels are tested via smoke tests on GPU.
+
+#[cfg(test)]
+mod tests {
+    use std::f64::consts::PI;
+
+    /// asymmetric_soft_clamp: f(x>=0) = min(x, 10), f(x<0) = -10*(1-exp(x/10))
+    fn asymmetric_soft_clamp(x: f32) -> f32 {
+        if x >= 0.0 { x.min(10.0) } else { -10.0 * (1.0 - (x / 10.0).exp()) }
+    }
+
+    #[test]
+    fn test_asymmetric_soft_clamp_properties() {
+        assert!((asymmetric_soft_clamp(0.0)).abs() < 1e-6, "f(0) = 0");
+        assert!((asymmetric_soft_clamp(5.0) - 5.0).abs() < 1e-6, "f(5) = 5 (linear)");
+        assert!((asymmetric_soft_clamp(15.0) - 10.0).abs() < 1e-6, "f(15) = 10 (capped)");
+        assert!((asymmetric_soft_clamp(-5.0) - (-3.935)).abs() < 0.01, "f(-5) ≈ -3.93");
+        assert!((asymmetric_soft_clamp(-15.0) - (-7.769)).abs() < 0.01, "f(-15) ≈ -7.77");
+        // Monotonicity
+        assert!(asymmetric_soft_clamp(3.0) > asymmetric_soft_clamp(2.0));
+        assert!(asymmetric_soft_clamp(-2.0) > asymmetric_soft_clamp(-3.0));
+        // Asymmetry: |f(-5)| < |f(5)|
+        assert!(asymmetric_soft_clamp(-5.0).abs() < asymmetric_soft_clamp(5.0).abs());
+    }
+
+    #[test]
+    fn test_kelly_prior_neutral() {
+        // With Jeffrey's prior only (no real trades): kelly_f = 0.0
+        let prior_wins = 2.0_f64;
+        let prior_losses = 2.0;
+        let prior_sum_wins = 0.01;
+        let prior_sum_losses = 0.01;
+
+        let eff_total = prior_wins + prior_losses;
+        let win_rate = prior_wins / eff_total;
+        let avg_win = prior_sum_wins / prior_wins;
+        let avg_loss = prior_sum_losses / prior_losses;
+        let payoff = avg_win / avg_loss;
+        let kelly_f = (payoff * win_rate - (1.0 - win_rate)) / payoff;
+
+        assert!(kelly_f.abs() < 0.01, "Kelly with prior only should be ~0 (no bet), got {kelly_f}");
+    }
+
+    #[test]
+    fn test_cosine_epsilon_schedule() {
+        let eps_start = 0.3_f64;
+        let eps_end = 0.02;
+        let total_epochs = 20;
+
+        // Epoch 0: should be near eps_start
+        let e0 = eps_end + 0.5 * (eps_start - eps_end) * (1.0 + (PI * 0.0 / total_epochs as f64).cos());
+        assert!((e0 - eps_start).abs() < 0.001, "Epoch 0 epsilon should be eps_start, got {e0}");
+
+        // Final epoch: should be near eps_end
+        let ef = eps_end + 0.5 * (eps_start - eps_end) * (1.0 + (PI * 1.0).cos());
+        assert!((ef - eps_end).abs() < 0.001, "Final epoch epsilon should be eps_end, got {ef}");
+
+        // Mid-epoch: should be between start and end
+        let em = eps_end + 0.5 * (eps_start - eps_end) * (1.0 + (PI * 0.5).cos());
+        assert!(em > eps_end && em < eps_start, "Mid epsilon {em} should be between {eps_end} and {eps_start}");
+    }
+
+    #[test]
+    fn test_td_lambda_degenerates() {
+        // λ=0 should give 1-step return (just r_0 + γ*Q(s_1))
+        let rewards = [1.0_f64, 2.0, 3.0];
+        let q_next = [10.0, 20.0, 30.0];
+        let gamma = 0.95;
+        let lambda = 0.0;
+
+        // G^λ with λ=0: only n=1 term survives
+        // G^(1) = r_0 + γ * Q(s_1) = 1.0 + 0.95 * 10.0 = 10.5
+        let g1 = rewards[0] + gamma * q_next[0];
+        assert!((g1 - 10.5).abs() < 0.01, "1-step return should be 10.5, got {g1}");
+
+        // λ=1 should give full Monte Carlo (N-step, no intermediate bootstrap)
+        // With max_trace=3: G^(3) = r_0 + γ*r_1 + γ^2*r_2 + γ^3*Q(s_3)
+        let g_mc = rewards[0] + gamma * rewards[1] + gamma * gamma * rewards[2]
+                 + gamma * gamma * gamma * q_next[2];
+        let expected = 1.0 + 0.95 * 2.0 + 0.9025 * 3.0 + 0.857375 * 30.0;
+        assert!((g_mc - expected).abs() < 0.01, "MC return should be {expected}, got {g_mc}");
+    }
+
+    #[test]
+    fn test_micro_reward_adaptive_scale() {
+        let base_scale = 0.001_f64;
+
+        // Calm market (vol = 0.25% ATR, half of baseline 0.5%)
+        let vol_calm = 0.0025;
+        let vol_ratio_calm = vol_calm / 0.005;
+        let adaptive_calm = base_scale / vol_ratio_calm.sqrt().max(0.5);
+        assert!(adaptive_calm > base_scale, "Calm market scale ({adaptive_calm}) should be > base ({base_scale})");
+
+        // Volatile market (vol = 2% ATR, 4x baseline)
+        let vol_volatile = 0.02;
+        let vol_ratio_volatile = vol_volatile / 0.005;
+        let adaptive_volatile = base_scale / vol_ratio_volatile.sqrt().max(0.5);
+        assert!(adaptive_volatile < base_scale, "Volatile market scale ({adaptive_volatile}) should be < base ({base_scale})");
+    }
+
+    #[test]
+    fn test_exposure_aux_gradient_unique() {
+        // With 9 exposures and a target of k=3, each exposure should get a different gradient
+        let b0_size = 9;
+        let target = 3;
+
+        // Simulate softmax over uniform logits
+        let logits = vec![0.0_f64; b0_size];
+        let prob = 1.0 / b0_size as f64;
+
+        let grads: Vec<f64> = (0..b0_size)
+            .map(|a| prob - if a == target { 1.0 } else { 0.0 })
+            .collect();
+
+        // Target exposure gets negative gradient (pull toward it)
+        assert!(grads[target] < 0.0, "Target exposure should get negative gradient");
+
+        // All non-targets get positive gradient (push away)
+        for a in 0..b0_size {
+            if a != target {
+                assert!(grads[a] > 0.0, "Non-target exposure {a} should get positive gradient");
+            }
+        }
+
+        // All non-target gradients are equal (uniform logits)
+        for a in 1..b0_size {
+            if a != target {
+                assert!((grads[a] - grads[0]).abs() < 1e-10,
+                    "Non-target gradients should be equal with uniform logits");
+            }
+        }
+
+        // Target gradient is different from non-target
+        assert!((grads[target] - grads[0]).abs() > 0.1,
+            "Target gradient should differ from non-target");
+    }
+
+    #[test]
+    fn test_hindsight_relabel_improves_reward() {
+        // If we have a winning trade with entry=100, and future prices [101, 103, 102],
+        // the hindsight optimal is exit at 103 (bar+2)
+        let entry = 100.0_f64;
+        let futures = [101.0, 103.0, 102.0];
+        let position = 1.0; // long
+
+        let actual_pnl = position * (futures[0] - entry); // exited at 101 → +1
+        let best_pnl = futures.iter()
+            .map(|&p| position * (p - entry))
+            .fold(0.0_f64, f64::max);
+
+        assert!((best_pnl - 3.0).abs() < 0.01, "Best PnL should be 3.0 (at 103)");
+        assert!(best_pnl >= actual_pnl, "Hindsight reward >= actual reward");
+    }
+
+    #[test]
+    fn test_curriculum_ordering() {
+        // Trending bars (large |return|) should have LOW difficulty score
+        // Choppy bars (small |return|) should have HIGH difficulty score
+        let trending_return = 0.02_f64; // 2% move in 5 bars
+        let choppy_return = 0.001;       // 0.1% move in 5 bars
+
+        let score_trending = 1.0 / (trending_return + 0.001);
+        let score_choppy = 1.0 / (choppy_return + 0.001);
+
+        assert!(score_trending < score_choppy,
+            "Trending (easy) should have lower score than choppy (hard)");
+    }
+
+    #[test]
+    fn test_popart_normalization() {
+        // After normalization, rewards should have mean≈0, std≈1
+        let rewards = vec![1.0_f64, 2.0, 3.0, 4.0, 5.0];
+        let mean = rewards.iter().sum::<f64>() / rewards.len() as f64;
+        let var = rewards.iter().map(|r| (r - mean).powi(2)).sum::<f64>() / rewards.len() as f64;
+        let std = var.sqrt();
+
+        let normalized: Vec<f64> = rewards.iter().map(|r| (r - mean) / std).collect();
+        let norm_mean = normalized.iter().sum::<f64>() / normalized.len() as f64;
+        let norm_var = normalized.iter().map(|r| (r - norm_mean).powi(2)).sum::<f64>() / normalized.len() as f64;
+
+        assert!(norm_mean.abs() < 1e-10, "Normalized mean should be 0, got {norm_mean}");
+        assert!((norm_var - 1.0).abs() < 0.01, "Normalized variance should be 1, got {norm_var}");
+    }
+}
+```
+
+- [ ] **Step 2: Register module in mod.rs**
+
+Find `crates/ml/src/trainers/dqn/smoke_tests/mod.rs`. Add:
+
+```rust
+pub mod reward_v8;
+```
+
+- [ ] **Step 3: Run tests**
+
+```bash
+SQLX_OFFLINE=true cargo test -p ml --lib -- reward_v8 --nocapture 2>&1 | tail -15
+```
+
+Expected: 9 passed, 0 failed.
+
+- [ ] **Step 4: Commit**
+
+```bash
+git add crates/ml/src/trainers/dqn/smoke_tests/reward_v8.rs \
+       crates/ml/src/trainers/dqn/smoke_tests/mod.rs
+git commit -m "test(v8): 9 explicit unit tests for all v8 reward components"
+```
+
+---
+
+### Task 13: Final Build Verification + Full Test Pass
+
+- [ ] **Step 1: Full workspace build**
+
+```bash
+SQLX_OFFLINE=true cargo check --workspace 2>&1 | tail -5
+```
+Expected: `Finished`
+
+- [ ] **Step 2: Full ML test suite**
+
+```bash
+SQLX_OFFLINE=true cargo test -p ml --lib 2>&1 | grep "^test result:"
+```
+Expected: `900+ passed; 0 failed` (891 existing + 9 new v8 tests)
+
+- [ ] **Step 3: Fix any regressions**
+
+If any tests fail, fix them. Common issues:
+- Config default changes affecting test expectations
+- Missing field initializations in test constructors
+
+- [ ] **Step 4: Update stale comments to v8**
+
+Search for "reward v6" and "reward v7" in kernel comments. Update to "reward v8".
+
+- [ ] **Step 5: Final commit**
+
+```bash
+git add -A
+git commit -m "fix(v8): final test pass + comment updates — 900+ tests pass, 0 regressions"
+```
diff --git a/docs/superpowers/specs/2026-04-07-reward-v7-design.md b/docs/superpowers/specs/2026-04-07-reward-v7-design.md
new file mode 100644
index 000000000..9f03ede76
--- /dev/null
+++ b/docs/superpowers/specs/2026-04-07-reward-v7-design.md
@@ -0,0 +1,297 @@
+# Reward v7: Counterfactual Branch Attribution Design
+
+**Goal:** Fix the 95% breakeven win rate caused by v6 penalty stacking and introduce novel per-branch credit assignment for the branching DQN's factored 81-action space.
+
+**Architecture:** Sparse trade-completion reward with asymmetric soft-clamp, counterfactual exposure advantage (CEA), order-type microstructure credit, and intra-trade risk efficiency. Removes 5 penalty layers, adds 3 novel reward components.
+
+**Tech Stack:** CUDA kernel (`experience_kernels.cu`), f32 arithmetic, branching DQN (9 exposure × 3 order × 3 urgency = 81 actions), C51 distributional RL.
+
+---
+
+## 1. Problem Statement
+
+Reward v6 applies 6+ penalty layers that compound to require a **94.9% win rate** for positive expected reward — mathematically impossible on ES futures. The model rationally converges to a "don't trade" (Flat) policy across ALL hyperopt trials. The exposure branch shows degenerate `i%3` Q-value patterns, confirming it receives no meaningful gradient signal.
+
+### Root Causes
+
+| Issue | Impact |
+|-------|--------|
+| Loss aversion (1.5×) applied BEFORE tanh | 2× asymmetry: wins +2.62, losses -3.80 |
+| Counterfactual regret always ≤ 0 | Systematically drags expected reward negative |
+| Hold scale redundant with min_hold_bars | Double-penalizes short trades |
+| Trade clustering penalty | Adds noise without meaningful signal |
+| Beta penalty | Already disabled (0.0), dead code |
+| Single scalar reward for 81 factored actions | Exposure branch can't learn action semantics |
+
+### Evidence
+
+- 50-trial hyperopt: ALL trials collapse to 100% Flat within 2-3 epochs
+- Win rate stuck at exactly ~30% (random epsilon baseline) across every configuration
+- Q-values: S100≈S25≈L50, S75≈Flat≈L75, S50≈L25≈L100 (i%3 degeneracy)
+- Research: stacking >3 reward penalties causes signal collapse (per RL literature review)
+
+## 2. Design
+
+### 2.1 Reward Pipeline Overview
+
+```
+AT TRADE EXIT:
+  1. segment_return = segment_pnl / prev_equity
+  2. vol_normalize(segment_return, atr, hold_time)
+  3. scale: base_reward = 10.0 * vol_normalized_return
+  4. squash: asymmetric_soft_clamp(base_reward)    ← replaces loss_aversion + tanh
+  5. + cea_weight * exposure_advantage              ← THE GEM (novel)
+  6. + order_weight * order_credit                  ← microstructure (novel)
+  7. + risk_eff_weight * risk_efficiency             ← path quality (novel)
+  8. + drawdown_penalty (dense, every bar)           ← kept from v6
+  9. capital_floor → reward = -10                    ← kept from v6
+
+DURING TRADE / WHEN FLAT:
+  reward = 0.0 + drawdown_penalty (if DD > threshold)
+```
+
+### 2.2 Layer 1: Sparse Trade-Completion Base (kept from v6)
+
+No changes to the core sparse reward. Validated by ETDQN (Takara et al., 2023) — outperforms dense rewards by 1.46-7.13× on intraday trading.
+
+```cuda
+float segment_return = segment_pnl / (prev_equity > 1.0f ? prev_equity : 1.0f);
+float vol_normalized_return = segment_return / vol_norm;
+float base_reward = 10.0f * vol_normalized_return;
+```
+
+Vol normalization uses ATR(14) proxy: `vol_norm = atr_pct * sqrt(hold_time)`. Makes rewards comparable across trending and ranging regimes.
+
+### 2.3 Layer 2: Asymmetric Soft-Clamp (replaces loss_aversion + tanh)
+
+**Removes:** `loss_aversion` parameter, separate `tanh` squash.
+
+The key insight: loss aversion applied BEFORE a nonlinear squash creates compound asymmetry that makes expected reward negative even for a fair game. Instead, use a single function that IS the risk aversion:
+
+```cuda
+__device__ float asymmetric_soft_clamp(float x) {
+    // Gains: linear with hard cap at +10 (full gradient preserved)
+    // Losses: exponential compression (smooth, natural risk aversion)
+    if (x >= 0.0f) return fminf(x, 10.0f);
+    return -10.0f * (1.0f - expf(x / 10.0f));
+}
+```
+
+**Properties:**
+- `f(+5.0) = +5.0` — full value, no compression
+- `f(-5.0) = -3.93` — compressed 21%
+- `f(-15.0) = -7.77` — compressed 48%
+- `f(+15.0) = +10.0` — capped (C51 bound)
+- Continuous, differentiable everywhere
+- No separate parameter needed — the function shape IS the risk preference
+- Gradient for gains: 1.0 (full). Gradient for losses at x=-5: `exp(-0.5) = 0.61`
+
+**Breakeven impact:** Expected reward for 50/50 win/loss at ±5 magnitude:
+- v6: `0.5 × 4.62 + 0.5 × (-9.05) = -2.22` (need 66% wins)
+- v7: `0.5 × 5.0 + 0.5 × (-3.93) = +0.54` (breakeven at ~44%)
+
+### 2.4 Layer 3: Counterfactual Exposure Advantage — CEA (novel)
+
+**Replaces:** `regret_blend` (always ≤ 0, pure punishment).
+
+**Core insight:** The branching DQN decomposes `Q(s,a) = V(s) + A_exposure(i) + A_order(j) + A_urgency(k)`. With one scalar reward, the network must implicitly discover which branch caused the outcome. CEA provides **explicit advantage signal** for the exposure branch.
+
+**Method:** At trade exit, compute the squashed reward for all 9 exposure levels using the same bar's price movement:
+
+```cuda
+float sum_alt_rewards = 0.0f;
+float price_delta = raw_next - raw_close;
+for (int k = 0; k < b0_size; k++) {
+    float alt_pos = compute_target_position(k, b0_size, max_position);
+    float alt_pnl = alt_pos * price_delta;
+    float alt_return = alt_pnl / (prev_equity > 1.0f ? prev_equity : 1.0f);
+    float alt_vol_norm = alt_return / vol_norm;
+    float alt_reward = asymmetric_soft_clamp(10.0f * alt_vol_norm);
+    sum_alt_rewards += alt_reward;
+}
+float mean_alt_reward = sum_alt_rewards / (float)b0_size;
+float exposure_advantage = squashed_reward - mean_alt_reward;
+
+reward += cea_weight * exposure_advantage;
+```
+
+**Properties:**
+- Can be **positive** (model outperformed average exposure) or **negative** (underperformed)
+- **Zero-mean by construction** — no systematic drag on expected reward
+- Zero additional compute cost — we already have the counterfactual loop from v6 regret
+- Directly breaks i%3 Q-value degeneracy: S100 has different advantage than L50 because the position sizes differ
+- For Flat action: exposure_advantage is always ≤ 0 when any directional exposure would have profited, naturally encouraging the model to explore directional actions
+
+**Why this is novel:** Counterfactual advantage decomposition has not been applied to branching DQN in the RL trading literature. It gives the exposure branch an **exact, per-step attribution signal** rather than forcing implicit credit assignment through shared scalar rewards.
+
+**Hyperopt parameter:** `cea_weight` in [0.1, 1.0], default 0.3.
+
+### 2.5 Layer 4: Order Type Microstructure Credit (novel)
+
+**New component:** Rewards cost-efficient execution by comparing the taken order type's transaction cost against the worst-case (MarketOrder).
+
+```cuda
+if (segment_complete && fabsf(position) > 0.001f) {
+    float market_cost = compute_tx_cost(position, raw_close, tx_cost_multiplier,
+                                        spread_cost, max_position, 0/*Market*/, -1.0f);
+    float taken_cost  = compute_tx_cost(position, raw_close, tx_cost_multiplier,
+                                        spread_cost, max_position, order_type_idx, spread_scale);
+    float order_credit = (market_cost - taken_cost) / (prev_equity > 1.0f ? prev_equity : 1.0f);
+    reward += order_weight * fminf(order_credit * 100.0f, 2.0f);  // cap at +2.0
+}
+```
+
+**Properties:**
+- Always ≥ 0 (LimitMaker saves spread, Market is baseline)
+- On ES: LimitMaker saves ~$3.12/trade → `order_credit ≈ 0.009%` → scaled to ~0.09
+- Gives the ORDER branch its own learning signal
+- Capped at +2.0 to prevent dominating the sparse exit reward
+
+**Hyperopt parameter:** `order_credit_weight` in [0.0, 0.5], default 0.1.
+
+### 2.6 Layer 5: Intra-Trade Risk Efficiency (novel)
+
+**New component:** Rewards clean winning trades (minimal intra-trade drawdown) over volatile ones.
+
+**Requires:** New portfolio state field `ps[22] = intra_trade_max_dd` (bf16), tracking the worst unrealized drawdown during the current trade. Reset to 0 at trade entry.
+
+**Per-bar update (during trade):**
+```cuda
+if (fabsf(position) > 0.001f) {
+    float unrealized = (position > 0) ? (raw_close - entry_price) / entry_price
+                                      : (entry_price - raw_close) / entry_price;
+    float current_dd = fminf(unrealized, 0.0f);  // negative when underwater
+    intra_trade_max_dd = fminf(intra_trade_max_dd, current_dd);
+}
+```
+
+**At trade exit (winners only):**
+```cuda
+if (segment_return > 0.0f && intra_trade_max_dd < -0.001f) {
+    float risk_eff = segment_return / fabsf(intra_trade_max_dd);
+    reward += risk_eff_weight * fminf(risk_eff, 5.0f);
+}
+```
+
+**Properties:**
+- Only fires for winning trades with intra-trade drawdown > 0.1%
+- A +2% trade that never went negative: no bonus (perfect entry doesn't need it)
+- A +2% trade that dipped -1% first: `risk_eff = 2.0` → bonus = +0.2 (at default weight)
+- A +2% trade that dipped -5% first: `risk_eff = 0.4` → bonus = +0.04 (less reward for volatility)
+- Captures **path quality** — novel in RL trading
+- Capped at 5.0 to prevent outlier domination
+
+**Hyperopt parameter:** `risk_efficiency_weight` in [0.0, 0.5], default 0.1.
+
+### 2.7 Layer 6: Drawdown Penalty (kept from v6)
+
+No changes. The only dense reward component. Fires every bar when drawdown exceeds `dd_threshold`.
+
+```cuda
+reward += compute_drawdown_penalty(equity, peak_equity, dd_threshold, w_dd);
+```
+
+Essential for risk management. Validated by research as the only dense component worth keeping.
+
+### 2.8 Layer 7: Capital Floor (kept from v6)
+
+No changes. Regulatory requirement (PDT $25K rule). Terminal penalty.
+
+```cuda
+if (check_capital_floor(portfolio_value, peak_equity)) {
+    reward = -10.0f;
+}
+```
+
+## 3. Removed Components
+
+| Component | Reason for Removal |
+|-----------|-------------------|
+| `loss_aversion` multiplier | Replaced by asymmetric_soft_clamp function shape |
+| `regret_blend` | Replaced by CEA (symmetric, can be positive) |
+| `hold_scale` / hold-time reward scaling | Redundant with min_hold_bars enforcement in trade physics |
+| `trade_clustering_penalty` | Weak impact (0.05 × CV), adds noise, redundant with HER temporal diversity |
+| `beta_penalty` | Already disabled (0.0), dead code removed |
+
+## 4. New Portfolio State Field
+
+Add `ps[22] = intra_trade_max_dd` (bf16):
+- Reset to 0.0 at trade entry (`entering_trade` or `reversing_trade` new segment)
+- Updated every bar during active trade: `min(current_unrealized_dd, intra_trade_max_dd)`
+- Read at trade exit for risk efficiency computation
+
+This increases portfolio state from 22 to 23 bf16 fields. No layout changes elsewhere — the portfolio state buffer is dynamically sized from `PORTFOLIO_STATE_DIM`.
+
+## 5. Configuration Changes
+
+### New Parameters (experience kernel args)
+
+| Parameter | Type | Default | Hyperopt Range | Purpose |
+|-----------|------|---------|---------------|---------|
+| `cea_weight` | f32 | 0.3 | [0.1, 1.0] | Counterfactual exposure advantage blend |
+| `order_credit_weight` | f32 | 0.1 | [0.0, 0.5] | Order type microstructure bonus |
+| `risk_efficiency_weight` | f32 | 0.1 | [0.0, 0.5] | Intra-trade path quality bonus |
+
+### Removed Parameters
+
+| Parameter | Previously | Status |
+|-----------|-----------|--------|
+| `loss_aversion` | 1.5 | REMOVED — function shape provides risk aversion |
+| `regret_blend` | 0.3 | REMOVED — replaced by CEA |
+| `trade_clustering_penalty` | 0.05 | REMOVED |
+| `beta_penalty` | 0.0 | REMOVED (was disabled) |
+| `position_entropy_weight` | 0.0 | KEPT (already disabled, may enable later) |
+
+### TOML Config Updates
+
+All 4 config files need updates:
+- `config/training/dqn-production.toml`
+- `config/training/dqn-hyperopt.toml`
+- `config/training/dqn-smoketest.toml`
+- `config/training/dqn-localdev.toml`
+
+## 6. Expected Breakeven Analysis
+
+| Configuration | Win Rate for E[R] = 0 |
+|---------------|----------------------|
+| v6 (current, all penalties) | 94.9% |
+| v7 base (soft-clamp only) | ~44% |
+| v7 + CEA (advantage boosts good picks) | ~42% |
+| v7 + CEA + order credit (saves ~$3/trade) | ~40% |
+| Realistic ES target with tx costs | ~52-55% |
+
+The ~12% reduction in breakeven win rate (from impossible 95% to achievable 52-55%) is the critical fix.
+
+## 7. Files Modified
+
+| File | Changes |
+|------|---------|
+| `crates/ml/src/cuda_pipeline/experience_kernels.cu` | Reward v7 logic, new kernel args, portfolio state field |
+| `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` | Wire new kernel args, update PORTFOLIO_STATE_DIM |
+| `config/training/dqn-production.toml` | New params, remove old ones |
+| `config/training/dqn-hyperopt.toml` | New hyperopt ranges, remove old params |
+| `config/training/dqn-smoketest.toml` | New params |
+| `config/training/dqn-localdev.toml` | New params |
+| `crates/ml/src/trainers/dqn/config.rs` | Add new config fields, remove old ones |
+| `crates/ml/src/hyperopt/adapters/dqn.rs` | Wire new hyperopt params |
+
+## 8. Testing
+
+### Smoke Test Validation
+- Run `cargo test -p ml --lib -- smoke_tests --ignored` with v7 reward
+- Verify: reward distribution is NOT all-zero (model trades)
+- Verify: exposure_advantage has both positive and negative values
+- Verify: win rate > 30% by epoch 3 (exceeds random baseline)
+- Verify: Q-values do NOT show i%3 degeneracy pattern
+
+### Numerical Sanity Checks
+- Breakeven trade (0 PnL): reward ≈ 0 + small CEA adjustment
+- +1% trade, 5-bar hold, 0.5% ATR: reward ≈ +4.5 (before CEA)
+- -1% trade, 5-bar hold, 0.5% ATR: reward ≈ -3.5 (after soft-clamp compression)
+- Capital floor breach: reward = -10.0 exactly
+
+### Regression Tests
+- Kelly statistics accumulation unchanged (wins/losses/returns tracking)
+- Drawdown penalty unchanged
+- Capital floor behavior unchanged
+- Episode termination unchanged