From 327349c5e21b80e4973baa06a04ad53de42ce62d Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Tue, 7 Apr 2026 23:14:40 +0200 Subject: [PATCH] =?UTF-8?q?docs:=20reward=20v8=20master=20plan=20=E2=80=94?= =?UTF-8?q?=2018-item=20comprehensive=20training=20overhaul?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 5 categories, 18 items: A) v7.1 fixes: aux GEMM, CEA warmup, OFI gate, dead code B) Bootstrap trap: GPU cosine epsilon, n-step→5, dense micro-reward C) Initialization: supervised pre-train, pessimistic Q-init, phase schedule D) Advanced: TD(λ), PopArt normalization, curriculum learning, hindsight relabel E) Dead code removal 4 new CUDA kernels, 3 modified, 12 new config fields, 4 removed. Full GPU, no CPU path, no memory copies. Target: 50%+ win rate. Co-Authored-By: Claude Opus 4.6 (1M context) --- .../plans/2026-04-07-reward-v7-gems-pearls.md | 493 +++++++++++++ .../superpowers/plans/2026-04-07-reward-v7.md | 602 ++++++++++++++++ .../plans/2026-04-07-reward-v71-fixes.md | 647 ++++++++++++++++++ .../specs/2026-04-07-reward-v8-master-plan.md | 526 ++++++++++++++ 4 files changed, 2268 insertions(+) create mode 100644 docs/superpowers/plans/2026-04-07-reward-v7-gems-pearls.md create mode 100644 docs/superpowers/plans/2026-04-07-reward-v7.md create mode 100644 docs/superpowers/plans/2026-04-07-reward-v71-fixes.md create mode 100644 docs/superpowers/specs/2026-04-07-reward-v8-master-plan.md diff --git a/docs/superpowers/plans/2026-04-07-reward-v7-gems-pearls.md b/docs/superpowers/plans/2026-04-07-reward-v7-gems-pearls.md new file mode 100644 index 000000000..6ecaaf825 --- /dev/null +++ b/docs/superpowers/plans/2026-04-07-reward-v7-gems-pearls.md @@ -0,0 +1,493 @@ +# Reward v7 Gems & Pearls — 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:** Enhance reward v7 with 9 additional innovations: 3 confirmed improvements (position entropy, urgency attribution, adaptive CEA), 3 novel gems (OFI-weighted confidence, reward label smoothing, exit timing), and 3 future-research pearls (Kelly sizing signal, PER diversity, multi-resolution reward). + +**Architecture:** All reward enhancements are in the CUDA experience kernel (`experience_kernels.cu`) and Rust config/wiring. No architectural changes — additive modifications to the existing v7 reward pipeline. New kernel args for weights, new portfolio state fields for tracking. + +**Tech Stack:** CUDA (NVRTC), Rust (cudarc), TOML config, `cargo test` + +--- + +### Task 1: Enable Position Entropy Bonus (config only) + +**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` + +The position entropy system is fully implemented (#19) — histogram tracking per episode, entropy computation at episode end, reward bonus. It's just disabled (weight=0.0). + +- [ ] **Step 1: Enable in all TOML configs** + +In each TOML file, find `position_entropy_weight` and change from `0.0` (or `0.01`) to `0.02`: + +`dqn-production.toml`: +```toml +position_entropy_weight = 0.02 +``` + +`dqn-hyperopt.toml` — in `[search_space]`: +```toml +position_entropy_weight = [0.01, 0.1] +``` +And in `[reward]` and `[phase_fast]`: +```toml +position_entropy_weight = 0.02 +``` + +`dqn-smoketest.toml` and `dqn-localdev.toml`: +```toml +position_entropy_weight = 0.02 +``` + +- [ ] **Step 2: Commit** + +```bash +git add config/training/dqn-production.toml config/training/dqn-hyperopt.toml \ + config/training/dqn-smoketest.toml config/training/dqn-localdev.toml +git commit -m "config: enable position entropy bonus (weight=0.02) to combat Flat collapse" +``` + +--- + +### Task 2: Add Urgency Branch Attribution + Exit Timing Quality + Kelly Signal to Config + +**Files:** +- Modify: `crates/ml/src/trainers/dqn/config.rs` +- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` + +Add config fields for the new reward components. We batch all new config fields into one task. + +- [ ] **Step 1: Add new fields to DQNConfig (config.rs)** + +Find the v7 reward fields (near `cea_weight`, `order_credit_weight`, `risk_efficiency_weight`). Add: + +```rust + /// Reward v7 gem: Urgency branch attribution weight. + /// Rewards patient execution that achieves price improvement. + pub urgency_credit_weight: f64, + + /// Reward v7 gem: Exit timing quality weight. + /// Rewards exits where the market reversed direction afterward. + pub exit_timing_weight: f64, + + /// Reward v7 gem: OFI-weighted reward confidence multiplier. + /// Amplifies rewards on trades aligned with informed order flow. + pub ofi_reward_weight: f64, + + /// Reward v7 gem: Kelly-optimal sizing signal weight. + /// Rewards position sizes close to Kelly criterion recommendation. + pub kelly_sizing_weight: f64, + + /// Reward v7 gem: Reward label smoothing noise scale. + /// Adds tiny noise to prevent C51 atom overfitting. + pub reward_noise_scale: f64, +``` + +- [ ] **Step 2: Set defaults in Default impl** + +In the Default impl, find the v7 defaults block and add: + +```rust + // Reward v7 gems & pearls + urgency_credit_weight: 0.1, + exit_timing_weight: 0.05, + ofi_reward_weight: 0.2, + kelly_sizing_weight: 0.1, + reward_noise_scale: 0.05, +``` + +- [ ] **Step 3: Set defaults in conservative() constructor** + +Find the v7 block in `conservative()` and add the same defaults. + +- [ ] **Step 4: Update apply_family_scaling for loss_shaping family** + +Find the loss_shaping block and add scaling for the new params: + +```rust + self.urgency_credit_weight = (self.urgency_credit_weight * li).clamp(0.0, 1.0); + self.exit_timing_weight = (self.exit_timing_weight * li).clamp(0.0, 0.5); + self.ofi_reward_weight = (self.ofi_reward_weight * li).clamp(0.0, 1.0); + self.kelly_sizing_weight = (self.kelly_sizing_weight * li).clamp(0.0, 1.0); + self.reward_noise_scale = (self.reward_noise_scale * li).clamp(0.0, 0.2); +``` + +- [ ] **Step 5: Add fields to ExperienceCollectorConfig (gpu_experience_collector.rs)** + +Find the v7 fields in `ExperienceCollectorConfig` struct and add: + +```rust + pub urgency_credit_weight: f32, + pub exit_timing_weight: f32, + pub ofi_reward_weight: f32, + pub kelly_sizing_weight: f32, + pub reward_noise_scale: f32, +``` + +Add defaults in the `Default` impl: + +```rust + urgency_credit_weight: 0.0, + exit_timing_weight: 0.0, + ofi_reward_weight: 0.0, + kelly_sizing_weight: 0.0, + reward_noise_scale: 0.0, +``` + +- [ ] **Step 6: Wire in training_loop.rs constructor** + +Find where `cea_weight` is wired in the `ExperienceCollectorConfig` constructor in `training_loop.rs` and add: + +```rust + urgency_credit_weight: self.hyperparams.urgency_credit_weight as f32, + exit_timing_weight: self.hyperparams.exit_timing_weight as f32, + ofi_reward_weight: self.hyperparams.ofi_reward_weight as f32, + kelly_sizing_weight: self.hyperparams.kelly_sizing_weight as f32, + reward_noise_scale: self.hyperparams.reward_noise_scale as f32, +``` + +- [ ] **Step 7: Verify build** + +Run: `SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5` + +- [ ] **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 +git commit -m "feat(reward-v7): add urgency, exit timing, OFI, Kelly, noise config fields" +``` + +--- + +### Task 3: Update CUDA Kernel Signature for Gems & Pearls + +**Files:** +- Modify: `crates/ml/src/cuda_pipeline/experience_kernels.cu` +- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` + +Add the 5 new float args to the kernel signature and wire them in Rust. + +- [ ] **Step 1: Add new kernel parameters** + +In `experience_kernels.cu`, find the kernel signature for `experience_env_step`. Add 5 new float params BEFORE the `saboteur_params` (which must remain last): + +After `float position_entropy_weight,` add: + +```cuda + float urgency_credit_weight, /* v7 gem: urgency branch attribution */ + float exit_timing_weight, /* v7 gem: exit timing quality */ + float ofi_reward_weight, /* v7 gem: OFI-weighted reward confidence */ + float kelly_sizing_weight, /* v7 gem: Kelly-optimal sizing signal */ + float reward_noise_scale, /* v7 gem: reward label smoothing noise */ +``` + +- [ ] **Step 2: Wire kernel launch args in gpu_experience_collector.rs** + +Find the `.arg(&config.position_entropy_weight)` line in the env_step launch block. Add AFTER it (before saboteur_params): + +```rust + .arg(&config.urgency_credit_weight) + .arg(&config.exit_timing_weight) + .arg(&config.ofi_reward_weight) + .arg(&config.kelly_sizing_weight) + .arg(&config.reward_noise_scale) +``` + +- [ ] **Step 3: Verify build** + +Run: `SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5` + +- [ ] **Step 4: 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(reward-v7): add 5 gem/pearl kernel args (urgency, timing, OFI, Kelly, noise)" +``` + +--- + +### Task 4: Implement All Gems & Pearls in CUDA Kernel + +**Files:** +- Modify: `crates/ml/src/cuda_pipeline/experience_kernels.cu` + +This is the main implementation. All 5 new reward components are added AFTER the existing v7 reward computation (after risk efficiency, before drawdown penalty). + +- [ ] **Step 1: Add urgency branch attribution (Layer 6)** + +Find the risk efficiency block (Layer 5). Add AFTER it, still inside the `if (segment_complete && segment_hold_time > 0.0f)` block: + +```cuda + /* ── Layer 6: Urgency branch attribution (fill quality) ── */ + if (urgency_credit_weight > 0.0f && fabsf(position) > 0.001f && entry_price > 0.0f) { + /* Price improvement: did we enter at a better price than the bar's midpoint? */ + float mid_price = (raw_close + raw_next) * 0.5f; + float fill_improvement = (position > 0.0f) + ? (mid_price - entry_price) / fmaxf(entry_price, 1.0f) + : (entry_price - mid_price) / fmaxf(entry_price, 1.0f); + /* Only reward improvement, don't penalize — urgency branch learns to be patient */ + float urgency_credit = fmaxf(fill_improvement / vol_proxy, 0.0f); + reward += urgency_credit_weight * fminf(urgency_credit, 2.0f); + } +``` + +- [ ] **Step 2: Add exit timing quality (Layer 7)** + +Add AFTER urgency credit: + +```cuda + /* ── Layer 7: Exit timing quality ── */ + if (exit_timing_weight > 0.0f && (exiting_trade || reversing_trade)) { + /* Good exit = market moved AGAINST our old position after we exited. + * Uses raw_next (1 bar ahead, already available). + * prev_sign holds the pre-trade direction. */ + float post_exit_move = (prev_sign > 0) + ? (raw_next - raw_close) / fmaxf(raw_close, 1.0f) + : (raw_close - raw_next) / fmaxf(raw_close, 1.0f); + /* Negative post_exit_move = market reversed = good exit timing */ + float timing_quality = -post_exit_move / vol_proxy; + reward += exit_timing_weight * asymmetric_soft_clamp(timing_quality * 10.0f) * 0.1f; + } +``` + +- [ ] **Step 3: Add OFI-weighted reward confidence (Layer 8)** + +Add AFTER exit timing, still inside `if (segment_complete)`: + +```cuda + /* ── Layer 8: OFI-weighted reward confidence ── */ + if (ofi_reward_weight > 0.0f && features != NULL && bar_idx < total_bars) { + /* OFI features are embedded in state at known offsets. + * Feature[40] = ADX (trend strength), Feature[41] = CUSUM (regime change). + * For OFI: features are in a SEPARATE buffer if ofi_dim > 0. + * Here we use a simpler proxy: the model's own position alignment + * with the price direction weighted by ATR-normalized move size. */ + float price_move = (raw_next - raw_close) / fmaxf(raw_close, 1.0f); + float move_magnitude = fabsf(price_move) / vol_proxy; + /* Alignment: position direction matches price direction */ + float alignment = 0.0f; + if (position > 0.001f && price_move > 0.0f) alignment = 1.0f; + if (position < -0.001f && price_move < 0.0f) alignment = 1.0f; + /* Only AMPLIFY good trades backed by strong moves, never penalize */ + float confidence_mult = alignment * fminf(move_magnitude, 1.0f); + reward *= (1.0f + ofi_reward_weight * confidence_mult); + } +``` + +- [ ] **Step 4: Add Kelly-optimal sizing signal (Layer 9)** + +Add AFTER OFI confidence, still inside `if (segment_complete)`: + +```cuda + /* ── Layer 9: Kelly-optimal position sizing signal ── */ + if (kelly_sizing_weight > 0.0f && (win_count + loss_count) >= 5.0f) { + /* Compute Kelly fraction from accumulated trade statistics. + * Kelly f* = (W/L * p - (1-p)) / (W/L) + * where p = win_rate, W = avg_win, L = avg_loss */ + float total_trades = win_count + loss_count; + float win_rate = win_count / total_trades; + float avg_win = (win_count > 0.0f) ? (sum_wins / win_count) : 0.001f; + float avg_loss = (loss_count > 0.0f) ? (sum_losses / loss_count) : 0.001f; + float payoff_ratio = avg_win / fmaxf(avg_loss, 0.0001f); + float kelly_f = (payoff_ratio * win_rate - (1.0f - win_rate)) / fmaxf(payoff_ratio, 0.0001f); + kelly_f = fmaxf(fminf(kelly_f, 1.0f), 0.0f); /* clamp [0, 1] */ + + /* How close was our actual position to Kelly-optimal? + * actual_fraction = |position| / max_position + * closeness = 1 - |actual - kelly| (1.0 = perfect, 0.0 = maximally wrong) */ + float actual_fraction = fabsf(position) / fmaxf(max_position, 1.0f); + float kelly_closeness = 1.0f - fabsf(actual_fraction - kelly_f); + kelly_closeness = fmaxf(kelly_closeness, 0.0f); + + /* Small bonus for Kelly-aligned sizing (only for winning trades) */ + if (segment_return > 0.0f) { + reward += kelly_sizing_weight * kelly_closeness * 0.5f; + } + } +``` + +- [ ] **Step 5: Close the segment_complete block and add reward label smoothing** + +The reward label smoothing happens OUTSIDE the `segment_complete` block, applied to ALL rewards (including zero rewards during trade): + +After the `segment_complete` closing brace `}`, and BEFORE the drawdown penalty, add: + +```cuda + /* ── Reward label smoothing (all steps) ── */ + if (reward_noise_scale > 0.0f && reward != 0.0f) { + /* Simple deterministic jitter based on episode + timestep. + * Not truly random but provides sufficient variation to prevent + * C51 atom overfitting. True curand would need per-thread state. */ + unsigned int hash = (unsigned int)(i * 31337 + bar_idx * 7919 + current_t * 1013); + hash ^= (hash >> 16); hash *= 0x45d9f3b; hash ^= (hash >> 16); + float pseudo_noise = ((float)(hash & 0xFFFF) / 65535.0f - 0.5f) * 2.0f; + reward += reward_noise_scale * pseudo_noise; + } +``` + +- [ ] **Step 6: Verify build** + +Run: `SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5` + +- [ ] **Step 7: Commit** + +```bash +git add crates/ml/src/cuda_pipeline/experience_kernels.cu +git commit -m "feat(reward-v7): implement all gems & pearls in CUDA kernel + +Layers 6-9 added to v7 reward: +- Urgency branch attribution (fill price improvement) +- Exit timing quality (market reversal after exit) +- OFI-weighted reward confidence (amplify informed trades) +- Kelly-optimal position sizing signal (reward Kelly-aligned sizing) +- Reward label smoothing (deterministic jitter, anti C51 overfitting)" +``` + +--- + +### Task 5: Implement Adaptive CEA Warmup + +**Files:** +- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs` + +- [ ] **Step 1: Add adaptive CEA warmup logic** + +Find where the `ExperienceCollectorConfig` is constructed (search for `cea_weight: self.hyperparams.cea_weight as f32`). Replace with: + +```rust + cea_weight: { + // Adaptive CEA warmup: full weight (1.0) for first 3 epochs + // to bootstrap exposure branch symmetry breaking, then decay + // to the configured/hyperopt-tuned value. + let base = self.hyperparams.cea_weight as f32; + if self.current_epoch < 3 { 1.0_f32.max(base) } else { base } + }, +``` + +- [ ] **Step 2: Verify build** + +Run: `SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5` + +- [ ] **Step 3: Commit** + +```bash +git add crates/ml/src/trainers/dqn/trainer/training_loop.rs +git commit -m "feat(reward-v7): adaptive CEA warmup — 1.0 for first 3 epochs, then configured value" +``` + +--- + +### Task 6: Update TOML Configs with Gems & Pearls Parameters + +**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: Update all 4 TOMLs** + +Add to the `[reward]` section of each TOML: + +```toml +# Reward v7 gems & pearls +urgency_credit_weight = 0.1 +exit_timing_weight = 0.05 +ofi_reward_weight = 0.2 +kelly_sizing_weight = 0.1 +reward_noise_scale = 0.05 +``` + +For `dqn-hyperopt.toml`, also add to `[search_space]`: + +```toml +# Reward v7 gems & pearls search ranges +urgency_credit_weight = [0.0, 0.5] +exit_timing_weight = [0.0, 0.2] +ofi_reward_weight = [0.0, 0.5] +kelly_sizing_weight = [0.0, 0.5] +reward_noise_scale = [0.01, 0.1] +``` + +And to `[phase_fast]`: + +```toml +urgency_credit_weight = 0.1 +exit_timing_weight = 0.05 +ofi_reward_weight = 0.2 +kelly_sizing_weight = 0.1 +reward_noise_scale = 0.05 +``` + +- [ ] **Step 2: Commit** + +```bash +git add config/training/dqn-production.toml config/training/dqn-hyperopt.toml \ + config/training/dqn-smoketest.toml config/training/dqn-localdev.toml +git commit -m "config(reward-v7): add gems & pearls params to all TOML profiles" +``` + +--- + +### Task 7: Wire Gems in Hyperopt Adapter + +**Files:** +- Modify: `crates/ml/src/hyperopt/adapters/dqn.rs` + +- [ ] **Step 1: Add logging for new params** + +Find where the v7 params are logged (search for `[v7] cea_weight`). Add similar logging for the new gems: + +```rust + info!(" [v7 gems] urgency={:.3} timing={:.3} ofi={:.3} kelly={:.3} noise={:.3}", + hyperparams.urgency_credit_weight, + hyperparams.exit_timing_weight, + hyperparams.ofi_reward_weight, + hyperparams.kelly_sizing_weight, + hyperparams.reward_noise_scale); +``` + +- [ ] **Step 2: Verify build** + +Run: `SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5` + +- [ ] **Step 3: Commit** + +```bash +git add crates/ml/src/hyperopt/adapters/dqn.rs +git commit -m "feat(reward-v7): wire gems & pearls logging in hyperopt adapter" +``` + +--- + +### Task 8: Build Verification, Test Fixes, and Smoke Test + +**Files:** +- Test: existing smoke tests + training_profile tests + +- [ ] **Step 1: Full workspace build** + +Run: `SQLX_OFFLINE=true cargo check --workspace 2>&1 | tail -5` +Expected: `Finished` + +- [ ] **Step 2: Run unit tests** + +Run: `SQLX_OFFLINE=true cargo test -p ml --lib 2>&1 | tail -15` +Fix any test failures from changed config defaults. + +- [ ] **Step 3: Run smoke test if test data available** + +Run: `FOXHUNT_TEST_DATA=test_data/futures-baseline SQLX_OFFLINE=true cargo test -p ml --lib -- smoke_tests --ignored --nocapture 2>&1 | tail -30` + +- [ ] **Step 4: Commit fixes** + +```bash +git add -A && git commit -m "fix(reward-v7): test expectation updates for gems & pearls" +``` diff --git a/docs/superpowers/plans/2026-04-07-reward-v7.md b/docs/superpowers/plans/2026-04-07-reward-v7.md new file mode 100644 index 000000000..cd83772f6 --- /dev/null +++ b/docs/superpowers/plans/2026-04-07-reward-v7.md @@ -0,0 +1,602 @@ +# Reward v7: Counterfactual Branch Attribution — 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:** Replace the v6 reward function (95% breakeven win rate) with v7 featuring asymmetric soft-clamp, counterfactual exposure advantage, order microstructure credit, and intra-trade risk efficiency — dropping breakeven to ~52%. + +**Architecture:** Modify the CUDA experience kernel's reward section, update Rust config/wiring, update all TOML configs. Portfolio state expands from 20→23 fields (add `intra_trade_max_dd`, repurpose 2 reserved slots for new params). Kernel signature changes to replace 4 removed args with 3 new ones. + +**Tech Stack:** CUDA (NVRTC), Rust (cudarc), TOML config, `cargo test` + +--- + +### Task 1: Expand Portfolio State and Add `asymmetric_soft_clamp` Device Function + +**Files:** +- Modify: `crates/ml/src/cuda_pipeline/experience_kernels.cu:23-52` (header + PORTFOLIO_STRIDE) +- Modify: `crates/ml/src/cuda_pipeline/trade_stats_kernel.cu:12` (PORTFOLIO_STRIDE constant) + +- [ ] **Step 1: Update PORTFOLIO_STRIDE from 20 → 23 and document new fields** + +In `experience_kernels.cu`, update the header comment and `#define`: + +```cuda +/* Portfolio state layout used by experience kernels ([N, PORTFOLIO_STRIDE=23]): + * [0] position — current contract position (signed) + * [1] cash — cash balance + * [2] portfolio_value — mark-to-market total value (cash + position * price) + * [3] (reserved) + * [4] (reserved) + * [5] (reserved) + * [6] (reserved) + * [7] peak_equity — high-water mark (init to initial_capital) + * [8] flat_counter — consecutive flat steps + * [9] prev_equity — equity at previous step (init to initial_capital) + * [10] hold_time — consecutive steps with position + * [11] realized_pnl — cumulative realized PnL + * [12] entry_price — trade entry price (raw) + * [13] trade_start_pnl — realized_pnl at trade entry + * [14] win_count — Kelly: number of profitable trade exits + * [15] loss_count — Kelly: number of losing trade exits + * [16] sum_wins — Kelly: cumulative profit from winners + * [17] sum_losses — Kelly: cumulative loss (absolute) from losers + * [18] sum_returns — running sum of segment returns + * [19] sum_sq_returns — running sum of squared segment returns + * [20] intra_trade_max_dd — worst unrealized drawdown during current trade (v7) + * [21] last_trade_t — timestep of last trade exit (for clustering, retained) + * [22] interval_sum — (retained from v6, used internally) + */ +#define PORTFOLIO_STRIDE 23 +``` + +Also update `trade_stats_kernel.cu` line 12: +```cuda +#define PORTFOLIO_STRIDE 23 +``` + +- [ ] **Step 2: Add `asymmetric_soft_clamp` device function** + +Add this device function BEFORE the `experience_env_step` kernel (after the `#include` / device function section, near line 55): + +```cuda +/* ── Reward v7: Asymmetric soft-clamp ────────────────────────────────── + * Gains: linear, capped at +10 (full gradient for positive rewards) + * Losses: exponential compression (smooth natural risk aversion) + * + * f(+5)=+5.0, f(-5)=-3.93, f(-15)=-7.77, f(+15)=+10.0 + * Replaces loss_aversion + tanh from v6. No separate parameter needed. */ +__device__ __forceinline__ float asymmetric_soft_clamp(float x) { + if (x >= 0.0f) return fminf(x, 10.0f); + return -10.0f * (1.0f - expf(x / 10.0f)); +} +``` + +- [ ] **Step 3: Update PORTFOLIO_STRIDE in gpu_experience_collector.rs** + +In `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs`, update the constant at line 54: + +```rust +const PORTFOLIO_STRIDE: usize = 23; +``` + +- [ ] **Step 4: Verify build compiles** + +Run: `SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -3` +Expected: `Finished` (no errors — PORTFOLIO_STRIDE is a constant used in allocation, no signature changes yet) + +- [ ] **Step 5: Commit** + +```bash +git add crates/ml/src/cuda_pipeline/experience_kernels.cu \ + crates/ml/src/cuda_pipeline/trade_stats_kernel.cu \ + crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +git commit -m "feat(reward-v7): expand PORTFOLIO_STRIDE 20→23, add asymmetric_soft_clamp device function" +``` + +--- + +### Task 2: Add v7 Config Fields and Remove v6 Penalty Fields + +**Files:** +- Modify: `crates/ml/src/trainers/dqn/config.rs` + +- [ ] **Step 1: Add new v7 reward fields to DQNConfig** + +Find the existing `loss_aversion` field (line ~774) and the nearby reward fields. Add the 3 new fields in the same region. Search for `pub loss_aversion` and add after the nearby reward fields: + +```rust + /// Reward v7: Counterfactual Exposure Advantage weight. + /// Blends per-branch advantage signal into the scalar reward. + /// 0.0 = disabled, 1.0 = full advantage signal. + pub cea_weight: f64, + + /// Reward v7: Order type microstructure credit weight. + /// Rewards cost-efficient execution (LimitMaker > Market). + pub order_credit_weight: f64, + + /// Reward v7: Intra-trade risk efficiency weight. + /// Rewards clean winning trades with minimal intra-trade drawdown. + pub risk_efficiency_weight: f64, +``` + +- [ ] **Step 2: Set defaults for new fields, zero out removed fields** + +Find the `Default` implementation for DQNConfig (search for the block containing `loss_aversion: 1.5`). Update: + +```rust + // Reward v6 penalties — REMOVED in v7 (zeroed for backward compat) + loss_aversion: 1.0, // 1.0 = no asymmetry (soft-clamp handles it) + regret_blend: 0.0, // replaced by CEA + trade_clustering_penalty: 0.0, // removed + beta_penalty_strength: 0.0, // removed (was already 0) + + // Reward v7 — Counterfactual Branch Attribution + cea_weight: 0.3, + order_credit_weight: 0.1, + risk_efficiency_weight: 0.1, +``` + +- [ ] **Step 3: Update `apply_family_scaling` for loss_shaping family** + +Find the `loss_shaping_intensity` block (line ~1489). Replace: + +```rust + // Loss shaping family — v7: scale CEA + order credit + risk efficiency + let li = self.loss_shaping_intensity; + self.cea_weight = (self.cea_weight * li).clamp(0.0, 2.0); + self.order_credit_weight = (self.order_credit_weight * li).clamp(0.0, 1.0); + self.risk_efficiency_weight = (self.risk_efficiency_weight * li).clamp(0.0, 1.0); + self.position_entropy_weight *= li; + // v6 legacy (zeroed, no-op but kept for TOML backward compat) + self.asymmetric_dd_weight *= li; +``` + +- [ ] **Step 4: Verify build** + +Run: `SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -3` +Expected: `Finished` + +- [ ] **Step 5: Commit** + +```bash +git add crates/ml/src/trainers/dqn/config.rs +git commit -m "feat(reward-v7): add cea_weight, order_credit_weight, risk_efficiency_weight config fields" +``` + +--- + +### Task 3: Update CUDA Kernel Signature and Wire New Args + +**Files:** +- Modify: `crates/ml/src/cuda_pipeline/experience_kernels.cu:812-858` (kernel signature) +- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs:1794-1843` (kernel launch) + +- [ ] **Step 1: Update kernel signature — replace removed args with new ones** + +In `experience_kernels.cu`, find the `experience_env_step` kernel signature (line ~812). Replace the `loss_aversion` parameter and the 4 removed penalty parameters with the 3 new v7 parameters: + +Replace `float loss_aversion,` (line 825) with: +```cuda + float cea_weight, /* v7: counterfactual exposure advantage weight */ +``` + +Replace `float beta_penalty,` (line 845) with: +```cuda + float order_credit_weight, /* v7: order type microstructure credit weight */ +``` + +Replace `float trade_clustering_penalty,` (line 846) with: +```cuda + float risk_efficiency_weight, /* v7: intra-trade risk efficiency weight */ +``` + +Remove `float regret_blend,` (line 853) entirely — this was the last penalty arg before saboteur_params. + +The final kernel args in order (after `position_entropy_weight`): +```cuda + float cea_weight, /* v7: counterfactual exposure advantage */ + float order_credit_weight, /* v7: order microstructure credit */ + float risk_efficiency_weight, /* v7: risk efficiency bonus */ + const float* __restrict__ saboteur_params +``` + +- [ ] **Step 2: Update kernel launch in gpu_experience_collector.rs** + +In `gpu_experience_collector.rs`, find the env_step launch block (line ~1789). Replace the arg wiring: + +Replace `let rw_loss_av = config.loss_aversion;` and its `.arg(&rw_loss_av)` with: +```rust + let cea_w = config.cea_weight as f32; +``` + +Then in the `.arg()` chain, replace `.arg(&rw_loss_av)` (line ~1809) with: +```rust + .arg(&cea_w) +``` + +Replace `.arg(&config.beta_penalty)` (line ~1829) with: +```rust + .arg(&(config.order_credit_weight as f32)) +``` + +Replace `.arg(&config.trade_clustering_penalty)` (line ~1830) with: +```rust + .arg(&(config.risk_efficiency_weight as f32)) +``` + +Remove the `.arg(&config.regret_blend)` line (line ~1834) entirely. + +- [ ] **Step 3: Verify build** + +Run: `SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -3` +Expected: `Finished` + +- [ ] **Step 4: 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(reward-v7): update kernel signature — replace v6 penalty args with v7 CEA/order/risk args" +``` + +--- + +### Task 4: Implement Reward v7 Core Logic in CUDA Kernel + +**Files:** +- Modify: `crates/ml/src/cuda_pipeline/experience_kernels.cu:1198-1410` (reward computation section) + +This is the main implementation task. Replace the v6 reward section with v7 logic. + +- [ ] **Step 1: Add intra-trade max DD tracking (per-bar update)** + +Find the line where `hold_time` is updated (line ~1189: `hold_time = update_hold_time(...)`). Add immediately AFTER it: + +```cuda + /* ── v7: Track intra-trade max drawdown for risk efficiency ── */ + float intra_trade_max_dd = __bfloat162float(ps[20]); + if (entering_trade) { + intra_trade_max_dd = 0.0f; /* reset at trade entry */ + } + if (reversing_trade) { + intra_trade_max_dd = 0.0f; /* reset for new segment */ + } + if (fabsf(position) > 0.001f && entry_price > 0.0f) { + float unrealized = (position > 0.0f) + ? (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); + } +``` + +- [ ] **Step 2: Replace v6 reward section with v7** + +Find the entire reward v6 block starting at line 1198 (`/* ================================================================`) through the end of the trade_clustering_penalty block (line ~1374). Replace the ENTIRE block (from `float reward = 0.0f;` through the end of the clustering penalty) with: + +```cuda + /* ================================================================ + * REWARD v7: Counterfactual Branch Attribution + * + * Fixes: v6 penalty stacking required 95% win rate breakeven. + * v7 drops breakeven to ~52% via: + * 1. Asymmetric soft-clamp (natural risk aversion, no loss_aversion param) + * 2. Counterfactual Exposure Advantage (per-branch gradient signal) + * 3. Order type microstructure credit (execution quality signal) + * 4. Intra-trade risk efficiency (path quality, not just endpoint) + * + * Kept from v6: sparse trade-completion, vol-normalization, drawdown + * penalty, capital floor. + * + * Removed from v6: loss_aversion, regret_blend, hold_scale, + * trade_clustering_penalty, beta_penalty. + * ================================================================ */ + + float reward = 0.0f; + + if (segment_complete && segment_hold_time > 0.0f) { + float segment_pnl = __bfloat162float(ps[11]) + raw_pnl - trade_start_pnl; + float segment_return = segment_pnl / (prev_equity > 1.0f ? prev_equity : 1.0f); + + /* Kelly statistics (unchanged from v6) */ + if (exiting_trade) { + if (segment_return > 0.0f) { + win_count += 1.0f; + sum_wins += segment_return; + } else { + loss_count += 1.0f; + sum_losses += fabsf(segment_return); + } + sum_returns += segment_return; + sum_sq_returns += segment_return * segment_return; + } + if (reversing_trade) { + segment_return = reversal_return; + } + + /* Vol normalization (unchanged from v6) */ + float atr_norm = 0.0f; + if (features != NULL && bar_idx < total_bars && market_dim > 9) { + atr_norm = __bfloat162float(features[(long long)bar_idx * market_dim + 9]); + } + float log_atr = atr_norm * 16.0f - 7.0f; + float atr_pct = expf(log_atr) / fmaxf(raw_close, 1.0f); + float vol_proxy = fmaxf(atr_pct, 0.0001f); + float vol_norm = vol_proxy * sqrtf(fmaxf(segment_hold_time, 1.0f)); + float vol_normalized_return = segment_return / vol_norm; + + /* ── Layer 2: Asymmetric soft-clamp (replaces loss_aversion + tanh) ── */ + float base_reward = 10.0f * vol_normalized_return; + reward = asymmetric_soft_clamp(base_reward); + + /* ── Layer 3: Counterfactual Exposure Advantage (CEA) ── */ + if (cea_weight > 0.0f) { + 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 = reward - mean_alt_reward; + reward += cea_weight * exposure_advantage; + } + + /* ── Layer 4: Order type microstructure credit ── */ + if (order_credit_weight > 0.0f && 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_credit_weight * fminf(order_credit * 100.0f, 2.0f); + } + + /* ── Layer 5: Intra-trade risk efficiency (path quality) ── */ + if (risk_efficiency_weight > 0.0f && segment_return > 0.0f && intra_trade_max_dd < -0.001f) { + float risk_eff = segment_return / fabsf(intra_trade_max_dd); + reward += risk_efficiency_weight * fminf(risk_eff, 5.0f); + } + } + + /* ── Layer 6: Drawdown penalty (dense, every bar — unchanged from v6) ── */ +``` + +Note: Everything AFTER the drawdown penalty (capital floor, position entropy, episode done, portfolio state writeback) stays unchanged from v6. Only the reward computation section (layers 1-5) is replaced. + +- [ ] **Step 3: Save `intra_trade_max_dd` back to portfolio state** + +Find where portfolio state fields are written back (search for `ps_base` or the block that writes `ps[10]`, `ps[11]`, etc.). In the portfolio state writeback section, add: + +```cuda + portfolio_states[ps_base + 20] = __float2bfloat16(intra_trade_max_dd); +``` + +- [ ] **Step 4: Update the kernel docstring** + +Replace the kernel's doc comment (lines ~784-811) header to say v7: + +```cuda +/** + * Portfolio simulation, reward v7 (Counterfactual Branch Attribution), and done detection. +``` + +- [ ] **Step 5: Verify build** + +Run: `SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -3` +Expected: `Finished` + +- [ ] **Step 6: Commit** + +```bash +git add crates/ml/src/cuda_pipeline/experience_kernels.cu +git commit -m "feat(reward-v7): implement CEA, asymmetric soft-clamp, order credit, risk efficiency in CUDA kernel" +``` + +--- + +### Task 5: Update Config TOML Files + +**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: Update dqn-production.toml** + +Find the `[reward]` section and update: + +```toml +[reward] +w_dsr = 1.0 +w_pnl = 0.3 +w_dd = 1.0 +w_idle = 0.0 +dd_threshold = 0.01 +loss_aversion = 1.0 # v7: no-op (soft-clamp handles asymmetry) +time_decay_rate = 0.0005 +q_gap_threshold = 0.1 +# Reward v7 — Counterfactual Branch Attribution +cea_weight = 0.3 +order_credit_weight = 0.1 +risk_efficiency_weight = 0.1 +``` + +- [ ] **Step 2: Update dqn-hyperopt.toml** + +In the `[search_space]` section, add v7 ranges and update removed params: + +```toml +# Reward v7 — Counterfactual Branch Attribution +cea_weight = [0.1, 1.0] +order_credit_weight = [0.0, 0.5] +risk_efficiency_weight = [0.0, 0.5] + +# v6 legacy (fixed at no-op values, not searched) +loss_aversion = [1.0, 1.0] +``` + +In the `[reward]` section: + +```toml +[reward] +w_dsr = 1.0 +w_pnl = 0.3 +w_dd = 1.0 +w_idle = 0.0 +dd_threshold = 0.01 +loss_aversion = 1.0 +time_decay_rate = 0.0005 +q_gap_threshold = 0.1 +cea_weight = 0.3 +order_credit_weight = 0.1 +risk_efficiency_weight = 0.1 +``` + +Also update `[phase_fast]`: + +```toml +loss_aversion = 1.0 +cea_weight = 0.3 +order_credit_weight = 0.1 +risk_efficiency_weight = 0.1 +``` + +- [ ] **Step 3: Update dqn-smoketest.toml** + +Find or add reward section, add v7 params: + +```toml +cea_weight = 0.3 +order_credit_weight = 0.1 +risk_efficiency_weight = 0.1 +loss_aversion = 1.0 +``` + +- [ ] **Step 4: Update dqn-localdev.toml** + +Same changes: + +```toml +cea_weight = 0.3 +order_credit_weight = 0.1 +risk_efficiency_weight = 0.1 +loss_aversion = 1.0 +``` + +- [ ] **Step 5: Commit** + +```bash +git add config/training/dqn-production.toml \ + config/training/dqn-hyperopt.toml \ + config/training/dqn-smoketest.toml \ + config/training/dqn-localdev.toml +git commit -m "config(reward-v7): add cea_weight, order_credit_weight, risk_efficiency_weight to all TOML profiles" +``` + +--- + +### Task 6: Wire Hyperopt Adapter for New Parameters + +**Files:** +- Modify: `crates/ml/src/hyperopt/adapters/dqn.rs` + +- [ ] **Step 1: Find where hyperopt params are read from TOML and applied to config** + +Search for `loss_aversion` in the hyperopt adapter to find where params are extracted. The adapter uses `apply_family_scaling` which we already updated in Task 2. But the hyperopt search space extraction may need updating. + +Search for where `search_space` values are read and applied. The adapter reads ranges from `[search_space]` in the TOML and generates trial params. Find where `loss_aversion` is read as a search param and add the new v7 params in the same pattern: + +```rust + // Reward v7 params + if let Some(v) = trial_params.get("cea_weight") { + config.cea_weight = *v; + } + if let Some(v) = trial_params.get("order_credit_weight") { + config.order_credit_weight = *v; + } + if let Some(v) = trial_params.get("risk_efficiency_weight") { + config.risk_efficiency_weight = *v; + } +``` + +- [ ] **Step 2: Verify build** + +Run: `SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -3` +Expected: `Finished` + +- [ ] **Step 3: Commit** + +```bash +git add crates/ml/src/hyperopt/adapters/dqn.rs +git commit -m "feat(reward-v7): wire cea_weight, order_credit_weight, risk_efficiency_weight in hyperopt adapter" +``` + +--- + +### Task 7: Run Local Smoke Test and Validate + +**Files:** +- Test: existing smoke tests in `crates/ml/src/trainers/dqn/smoke_tests/` + +- [ ] **Step 1: Run cargo check for full workspace** + +Run: `SQLX_OFFLINE=true cargo check --workspace 2>&1 | tail -5` +Expected: `Finished` with no errors + +- [ ] **Step 2: Run ML crate unit tests** + +Run: `SQLX_OFFLINE=true cargo test -p ml --lib 2>&1 | tail -10` +Expected: All tests pass. Any compilation errors from the CUDA kernel changes will surface here. + +- [ ] **Step 3: Run smoke test (if test data available)** + +Run: `FOXHUNT_TEST_DATA=test_data/futures-baseline SQLX_OFFLINE=true cargo test -p ml --lib -- smoke_tests --ignored --nocapture 2>&1 | tail -30` +Expected: Training runs, rewards are non-zero, no panics. + +- [ ] **Step 4: Verify reward v7 diagnostics in smoke test output** + +Check the output for: +- `mean_reward` should NOT be stuck at 0.0 (model is generating trades) +- Action diversity should show >1 exposure level by epoch 2-3 +- No NaN/Inf in rewards + +- [ ] **Step 5: Commit any test fixes if needed** + +```bash +git add -A +git commit -m "fix(reward-v7): address smoke test issues" +``` + +--- + +### Task 8: Deploy and Run Hyperopt + +- [ ] **Step 1: Deploy to Argo** + +```bash +./scripts/argo-compile-deploy.sh --branch main +``` + +- [ ] **Step 2: Launch hyperopt with v7 reward** + +```bash +./scripts/argo-train.sh --model dqn --config dqn-hyperopt +``` + +- [ ] **Step 3: Monitor early metrics** + +Check first epoch of first trial: +- Win rate should be >30% (breaking the random baseline) +- Action diversity should be >1/9 exposure levels +- Q-values should NOT show i%3 pattern +- Rewards should have both positive and negative values + +```bash +kubectl logs -l workflow=train-dqn --tail=50 | grep -E "Epoch.*Sharpe|Action diversity|Per-action Q" +``` diff --git a/docs/superpowers/plans/2026-04-07-reward-v71-fixes.md b/docs/superpowers/plans/2026-04-07-reward-v71-fixes.md new file mode 100644 index 000000000..7ba2f4555 --- /dev/null +++ b/docs/superpowers/plans/2026-04-07-reward-v71-fixes.md @@ -0,0 +1,647 @@ +# Reward v7.1: Six Critical Fixes — 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:** Fix 6 concerns from the v7 hyperopt run: i%3 exposure degeneracy (auxiliary loss), OFI epoch gate, CEA warmup extension, exit timing bug, Kelly prior seeding, and reward noise scaling. Remove all dead v6 code. Zero test regressions (891+ tests must pass). + +**Architecture:** Fixes 2-6 are surgical edits to the experience kernel and training loop. Fix 1 adds a new `exposure_aux_grad_kernel` in `dqn_utility_kernels.cu`, following the IQN/CQL auxiliary gradient pattern. Dead v6 fields (`loss_aversion`, `beta_penalty`, `regret_blend`, `trade_clustering_penalty`) are removed from `ExperienceCollectorConfig` and the kernel launch chain. + +**Tech Stack:** CUDA (nvcc → cubin via build.rs), Rust (cudarc), TOML config + +--- + +### Task 1: Fixes 4, 5, 6 — Surgical CUDA Kernel Edits + +**Files:** +- Modify: `crates/ml/src/cuda_pipeline/experience_kernels.cu` + +These three fixes are localized edits within the reward section of `experience_env_step`. No signature changes, no Rust wiring. + +- [ ] **Step 1: Fix 4 — Add `closing_sign` variable for exit timing** + +Find where `prev_sign` is computed (search for `int prev_sign`). It's computed early in the kernel from `pre_trade_position`. Add `closing_sign` immediately after: + +```cuda + int prev_sign = (pre_trade_position > 0.001f) ? 1 : ((pre_trade_position < -0.001f) ? -1 : 0); + int closing_sign = prev_sign; /* v7.1: capture for exit timing before reversal logic */ +``` + +Then find Layer 7 (exit timing quality, search for `Layer 7: Exit timing`). Replace `prev_sign` with `closing_sign`: + +```cuda + /* ── Layer 7: Exit timing quality ── */ + if (exit_timing_weight > 0.0f && (exiting_trade || reversing_trade)) { + float post_exit_move = (closing_sign > 0) + ? (raw_next - raw_close) / fmaxf(raw_close, 1.0f) + : (raw_close - raw_next) / fmaxf(raw_close, 1.0f); + float timing_quality = -post_exit_move / vol_proxy; + reward += exit_timing_weight * asymmetric_soft_clamp(timing_quality * 10.0f) * 0.1f; + } +``` + +- [ ] **Step 2: Fix 5 — Kelly prior seeding** + +Find Layer 9 (Kelly-optimal, search for `Layer 9: Kelly`). Replace the entire block: + +```cuda + /* ── Layer 9: Kelly-optimal position sizing signal ── */ + if (kelly_sizing_weight > 0.0f) { + /* Bayesian prior (Jeffrey's): 2 pseudo-wins + 2 pseudo-losses. + * Regularizes toward no-bet with no data, converges to empirical + * Kelly as real trades accumulate. Always fires (no threshold). */ + float prior_wins = 2.0f; + float prior_losses = 2.0f; + float prior_sum_wins = 0.01f; + float prior_sum_losses = 0.01f; + + float eff_wins = win_count + prior_wins; + float eff_losses = loss_count + prior_losses; + float eff_total = eff_wins + eff_losses; + + float win_rate_k = eff_wins / eff_total; + float avg_win = (sum_wins + prior_sum_wins) / eff_wins; + float avg_loss = (sum_losses + prior_sum_losses) / eff_losses; + float payoff_ratio = avg_win / fmaxf(avg_loss, 0.0001f); + float kelly_f = (payoff_ratio * win_rate_k - (1.0f - win_rate_k)) + / fmaxf(payoff_ratio, 0.0001f); + kelly_f = fmaxf(fminf(kelly_f, 1.0f), 0.0f); + + float actual_fraction = fabsf(position) / fmaxf(max_position, 1.0f); + float kelly_closeness = 1.0f - fabsf(actual_fraction - kelly_f); + kelly_closeness = fmaxf(kelly_closeness, 0.0f); + + if (segment_return > 0.0f) { + reward += kelly_sizing_weight * kelly_closeness * 0.5f; + } + } +``` + +- [ ] **Step 3: Fix 6 — Reward noise magnitude scaling** + +Find the reward label smoothing block (search for `Reward label smoothing`). Replace: + +```cuda + /* ── Reward label smoothing (magnitude-relative) ── */ + if (reward_noise_scale > 0.0f && reward != 0.0f) { + unsigned int hash = (unsigned int)(i * 31337 + bar_idx * 7919 + current_t * 1013); + hash ^= (hash >> 16); hash *= 0x45d9f3bu; hash ^= (hash >> 16); + float pseudo_noise = ((float)(hash & 0xFFFF) / 65535.0f - 0.5f) * 2.0f; + /* Scale noise to reward_noise_scale% of reward magnitude, floor at 0.01 */ + float noise_magnitude = fmaxf(fabsf(reward) * reward_noise_scale, 0.01f); + reward += noise_magnitude * pseudo_noise; + } +``` + +- [ ] **Step 4: Verify build** + +Run: `SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5` +Expected: `Finished` + +- [ ] **Step 5: Commit** + +```bash +git add crates/ml/src/cuda_pipeline/experience_kernels.cu +git commit -m "fix(reward-v7.1): exit timing closing_sign, Kelly prior seeding, noise magnitude scaling" +``` + +--- + +### Task 2: Fix 1a — Best Exposure Output in Experience Kernel + +**Files:** +- Modify: `crates/ml/src/cuda_pipeline/experience_kernels.cu` (kernel signature + CEA loop) +- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (buffer alloc + arg wiring) + +- [ ] **Step 1: Add `out_best_exposure` parameter to kernel signature** + +Find the kernel signature for `experience_env_step`. Add `int* out_best_exposure` parameter BEFORE `saboteur_params` (saboteur must remain last). Insert after `float reward_noise_scale,`: + +```cuda + int* out_best_exposure, /* v7.1: counterfactual best exposure idx per step (-1 = no trade exit) */ + /* #33 Per-episode saboteur params ... */ + const float* __restrict__ saboteur_params +``` + +- [ ] **Step 2: Track best exposure in the CEA loop and write output** + +Find the CEA loop (search for `Layer 3: Counterfactual Exposure Advantage`). Modify to track `best_k`: + +```cuda + /* ── Layer 3: Counterfactual Exposure Advantage (CEA) ── */ + int best_exposure_k = 0; + float best_exposure_reward = -1e30f; + if (cea_weight > 0.0f) { + 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; + if (alt_reward > best_exposure_reward) { + best_exposure_reward = alt_reward; + best_exposure_k = k; + } + } + float mean_alt_reward = sum_alt_rewards / (float)b0_size; + float exposure_advantage = reward - mean_alt_reward; + reward += cea_weight * exposure_advantage; + } +``` + +Then find where `out_rewards[out_off] = reward;` is written (near end of kernel). Add immediately BEFORE it: + +```cuda + /* v7.1: Write counterfactual best exposure for auxiliary loss */ + if (out_best_exposure != NULL) { + out_best_exposure[out_off] = (segment_complete && segment_hold_time > 0.0f) + ? best_exposure_k : -1; + } +``` + +- [ ] **Step 3: Add buffer and wiring in gpu_experience_collector.rs** + +Add a new field to the collector struct. Find `rewards_out: CudaSlice` (line ~529). Add after `done_out`: + +```rust + best_exposure_out: CudaSlice, // [alloc_episodes * alloc_timesteps] v7.1 aux target +``` + +Allocate in the constructor. Find where `done_out` is allocated (search for `alloc done_out`). Add after: + +```rust + let best_exposure_out = stream + .alloc_zeros::(alloc_episodes * alloc_timesteps) + .map_err(|e| MLError::ModelError(format!("alloc best_exposure_out: {e}")))?; +``` + +Add to the struct literal where `done_out` is assigned: + +```rust + best_exposure_out, +``` + +Wire in the kernel launch. Find the `.arg(&config.reward_noise_scale)` line in the env_step launch. Add AFTER it (before saboteur): + +```rust + .arg(&mut self.best_exposure_out) // v7.1: counterfactual best exposure +``` + +Add a public accessor for the training loop to read the buffer: + +```rust + /// v7.1: Best counterfactual exposure per step (for auxiliary loss). + pub fn best_exposure_buf(&self) -> &CudaSlice { + &self.best_exposure_out + } +``` + +- [ ] **Step 4: Verify build** + +Run: `SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5` +Expected: `Finished` + +- [ ] **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(reward-v7.1): track best_exposure_idx in experience kernel for auxiliary loss" +``` + +--- + +### Task 3: Fix 1b — Exposure Auxiliary Gradient Kernel + +**Files:** +- Modify: `crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu` + +- [ ] **Step 1: Add `exposure_aux_grad_kernel` to the utility kernels file** + +Find the END of `dqn_utility_kernels.cu` (after the last kernel). Append: + +```cuda +/* ══════════════════════════════════════════════════════════════════════ + * EXPOSURE AUXILIARY GRADIENT KERNEL (v7.1) + * + * Breaks the i%3 Q-value degeneracy in branching DQN by giving each + * exposure output a UNIQUE gradient via cross-entropy against the + * counterfactual best exposure. + * + * Follows IQN/CQL pattern: accumulates auxiliary gradients into grad_buf + * between C51 backward and Adam. The C51 gradient only differentiates + * selected vs non-selected (-1/9 for all 8 non-selected). This kernel + * gives each exposure output a target-specific gradient. + * + * Launch: grid=(ceil(batch_size/256)), block=(256). + * ══════════════════════════════════════════════════════════════════════ */ +extern "C" __global__ void exposure_aux_grad_kernel( + const float* __restrict__ adv_logits, /* [batch_size, b0_size * num_atoms] exposure branch logits (f32) */ + const int* __restrict__ targets, /* [batch_size] best exposure idx (-1 = skip) */ + float* __restrict__ grad_buf, /* [total_params] gradient accumulator */ + int adv_offset, /* offset into grad_buf for exposure adv weights */ + int b0_size, /* 9 */ + int num_atoms, /* 51 or 101 */ + float aux_weight, /* scaled weight (decays over epochs) */ + int batch_size +) { + int sample = blockIdx.x * blockDim.x + threadIdx.x; + if (sample >= batch_size) return; + + int target = targets[sample]; + if (target < 0 || target >= b0_size) return; /* skip non-exit steps */ + + const float* logits = adv_logits + (long long)sample * b0_size * num_atoms; + float inv_batch = aux_weight / (float)batch_size; + + for (int z = 0; z < num_atoms; z++) { + /* Stable softmax over b0_size for atom 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); + + /* Cross-entropy gradient: softmax(logit) - one_hot(target) */ + for (int a = 0; a < b0_size; a++) { + float prob = expf(logits[a * num_atoms + z] - max_l) * inv_sum; + float grad = (prob - ((a == target) ? 1.0f : 0.0f)) * inv_batch; + int param_idx = adv_offset + a * num_atoms + z; + atomicAdd(&grad_buf[param_idx], grad); + } + } +} +``` + +- [ ] **Step 2: Verify build** + +Run: `SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5` +Expected: `Finished` (CUDA compiled by build.rs at deploy time) + +- [ ] **Step 3: Commit** + +```bash +git add crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu +git commit -m "feat(reward-v7.1): add exposure_aux_grad_kernel for per-branch gradient signal" +``` + +--- + +### Task 4: Fix 1c — Wire Auxiliary Kernel in Training Pipeline + +**Files:** +- Modify: `crates/ml/src/trainers/dqn/config.rs` +- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` +- Modify: `crates/ml/src/trainers/dqn/fused_training.rs` + +- [ ] **Step 1: Add config fields** + +In `crates/ml/src/trainers/dqn/config.rs`, find the v7 gem fields (near `urgency_credit_weight`). Add: + +```rust + /// v7.1: Exposure auxiliary loss weight (scheduled decay). + pub exposure_aux_weight: f64, + /// v7.1: Epochs over which exposure aux weight decays from base to 10%. + pub exposure_aux_warmup_epochs: usize, +``` + +Add defaults in BOTH `Default` impl and `conservative()`: + +```rust + exposure_aux_weight: 0.5, + exposure_aux_warmup_epochs: 5, +``` + +In `apply_family_scaling`, add in the loss_shaping block: + +```rust + self.exposure_aux_weight = (self.exposure_aux_weight * li).clamp(0.0, 2.0); +``` + +- [ ] **Step 2: Load kernel in GpuDqnTrainer** + +In `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`, find where utility kernels are loaded (search for `load_function("dqn_grad_norm_kernel")`). The utility module is loaded from cubin. Add near the other kernel loads: + +```rust + let exposure_aux_grad = module.load_function("exposure_aux_grad_kernel") + .map_err(|e| MLError::ModelError(format!("exposure_aux_grad_kernel load: {e}")))?; +``` + +Add the function handle to the struct fields (search for `grad_norm_finalize_kernel: CudaFunction`): + +```rust + exposure_aux_grad_kernel: CudaFunction, +``` + +Store it in the constructor's struct literal: + +```rust + exposure_aux_grad_kernel: exposure_aux_grad, +``` + +Add a public method to launch the kernel: + +```rust + /// v7.1: Launch exposure auxiliary gradient kernel. + /// Adds cross-entropy gradient to grad_buf for exposure branch outputs. + pub fn launch_exposure_aux_grad( + &self, + targets: &CudaSlice, + aux_weight: f32, + batch_size: usize, + ) -> Result<(), MLError> { + if aux_weight < 1e-6 { return Ok(()); } + let _evt_guard = EventTrackingGuard::new(self.stream.context()); + + let adv_logits_ptr = self.on_adv_logits_buf.raw_ptr(); // exposure branch logits (f32) + let targets_ptr = targets.raw_ptr(); + let grad_ptr = self.ptrs.grad_buf; + let adv_offset = self.exposure_adv_grad_offset() as i32; + let b0 = self.config.branch_0_size as i32; + let na = self.config.num_atoms as i32; + let bs = batch_size as i32; + + let blocks = ((batch_size + 255) / 256) as u32; + unsafe { + self.stream + .launch_builder(&self.exposure_aux_grad_kernel) + .arg(&adv_logits_ptr) + .arg(&targets_ptr) + .arg(&grad_ptr) + .arg(&adv_offset) + .arg(&b0) + .arg(&na) + .arg(&aux_weight) + .arg(&bs) + .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: {e}")))?; + } + Ok(()) + } + + /// Offset into the flat grad_buf where the exposure advantage weights begin. + fn exposure_adv_grad_offset(&self) -> usize { + // grad_buf layout: [shared_weights..., value_weights..., adv_weights...] + // The advantage section starts after shared + value weights. + // For branching DQN: first b0_size * num_atoms elements of adv are the exposure branch. + self.config.shared_param_count() + self.config.value_param_count() + } +``` + +NOTE: The implementer must verify `on_adv_logits_buf` exists as an f32 buffer (it may be bf16 — check and cast if needed), and `exposure_adv_grad_offset()` computes the correct offset by inspecting the existing `compute_total_params` and weight layout functions. + +- [ ] **Step 3: Wire in fused training** + +In `crates/ml/src/trainers/dqn/fused_training.rs`, find `submit_aux_ops()` where IQN trunk gradient and CQL gradient are applied (search for `apply_iqn_trunk_gradient` or `apply_cql_gradient`). Add the exposure aux gradient call AFTER the CQL block: + +```rust + // v7.1: Exposure auxiliary gradient (breaks i%3 degeneracy) + if self.exposure_aux_weight > 0.0 { + if let Some(ref targets_buf) = self.best_exposure_targets { + self.trainer.launch_exposure_aux_grad( + targets_buf, + self.exposure_aux_weight as f32, + self.batch_size, + ).map_err(|e| anyhow::anyhow!("Exposure aux grad: {e}"))?; + } + } +``` + +Add the `exposure_aux_weight` field and `best_exposure_targets` buffer to `FusedTrainingCtx`: + +```rust + exposure_aux_weight: f64, + best_exposure_targets: Option>, // copied from experience collector per epoch +``` + +Initialize in the constructor from hyperparams: + +```rust + exposure_aux_weight: hyperparams.exposure_aux_weight, + best_exposure_targets: None, +``` + +Add a method to upload targets from the experience collector's buffer: + +```rust + pub fn set_exposure_targets(&mut self, targets: &CudaSlice, batch_size: usize) -> Result<()> { + if self.best_exposure_targets.is_none() { + self.best_exposure_targets = Some( + self.stream.alloc_zeros::(batch_size) + .map_err(|e| anyhow::anyhow!("alloc exposure targets: {e}"))? + ); + } + let dst = self.best_exposure_targets.as_mut().unwrap(); + self.stream.memcpy_dtod(dst, targets) + .map_err(|e| anyhow::anyhow!("copy exposure targets: {e}"))?; + Ok(()) + } +``` + +- [ ] **Step 4: Verify build** + +Run: `SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5` +Expected: `Finished` + +- [ ] **Step 5: Commit** + +```bash +git add crates/ml/src/trainers/dqn/config.rs \ + crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \ + crates/ml/src/trainers/dqn/fused_training.rs +git commit -m "feat(reward-v7.1): wire exposure auxiliary gradient kernel in training pipeline" +``` + +--- + +### Task 5: Fixes 2, 3 — Training Loop Schedule Changes + +**Files:** +- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs` + +- [ ] **Step 1: Fix 3 — Extended CEA warmup (25% linear decay)** + +Find where `cea_weight` is set in the `ExperienceCollectorConfig` constructor (search for `cea_weight:`). Replace: + +```rust + cea_weight: { + // v7.1: Linear warmup over 25% of training — blend from 1.0 to configured value + 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: Fix 2 — OFI epoch gate** + +Find where `ofi_reward_weight` is set (should be near `cea_weight`). Replace: + +```rust + ofi_reward_weight: { + // v7.1: Disable OFI for first 5 epochs — model learns direction first + let base = self.hyperparams.ofi_reward_weight as f32; + if self.current_epoch < 5 { 0.0 } else { base } + }, +``` + +- [ ] **Step 3: Add exposure aux weight decay schedule** + +Find where the `ExperienceCollectorConfig` is constructed. After the config is built, compute the decayed exposure aux weight and pass it to the fused context: + +```rust + // v7.1: Exposure auxiliary loss weight with scheduled decay + let exposure_aux_w = { + 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); + base * (1.0 - 0.9 * progress) // decays from base to 10% of base + }; + if let Some(ref mut fused) = self.fused_ctx { + fused.exposure_aux_weight = exposure_aux_w; + } +``` + +- [ ] **Step 4: Verify build** + +Run: `SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5` +Expected: `Finished` + +- [ ] **Step 5: Commit** + +```bash +git add crates/ml/src/trainers/dqn/trainer/training_loop.rs +git commit -m "fix(reward-v7.1): CEA 25% linear warmup, OFI epoch gate, exposure aux decay schedule" +``` + +--- + +### Task 6: Remove Dead v6 Code + +**Files:** +- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` +- Modify: `crates/ml/src/trainers/dqn/config.rs` +- Modify: `crates/ml/src/training_profile.rs` + +Remove dead v6 config fields that are zeroed and no longer used by the kernel. + +- [ ] **Step 1: Remove dead fields from ExperienceCollectorConfig** + +In `gpu_experience_collector.rs`, find and REMOVE these struct fields: + +```rust + pub loss_aversion: f32, // v6 dead: replaced by asymmetric_soft_clamp + pub beta_penalty: f32, // v6 dead: was disabled (0.0) + pub regret_blend: f32, // v6 dead: replaced by CEA + pub trade_clustering_penalty: f32, // v6 dead: removed +``` + +Remove them from the `Default` impl too: + +```rust + loss_aversion: 1.5, // REMOVE + beta_penalty: 0.0, // REMOVE + regret_blend: 0.0, // REMOVE + trade_clustering_penalty: 0.0, // REMOVE +``` + +- [ ] **Step 2: Remove dead field assignments from training_loop.rs** + +Find where these fields are assigned in the `ExperienceCollectorConfig` constructor: + +```rust + loss_aversion: self.hyperparams.loss_aversion as f32, // REMOVE + beta_penalty: self.hyperparams.beta_penalty_strength as f32, // REMOVE + regret_blend: self.hyperparams.regret_blend as f32, // REMOVE + trade_clustering_penalty: self.hyperparams.trade_clustering_penalty as f32, // REMOVE +``` + +- [ ] **Step 3: Fix any compilation errors** + +After removing fields, search for any remaining references to the removed fields and fix them. Common locations: +- `training_profile.rs`: may assign these fields — change to no-ops or remove +- Other test files or constructors + +Run: `SQLX_OFFLINE=true cargo check -p ml 2>&1 | grep "error"` and fix each one. + +- [ ] **Step 4: Verify build + run tests** + +Run: `SQLX_OFFLINE=true cargo check -p ml 2>&1 | tail -5` +Then: `SQLX_OFFLINE=true cargo test -p ml --lib 2>&1 | grep "^test result:"` +Expected: `ok. 891 passed; 0 failed` + +- [ ] **Step 5: Commit** + +```bash +git add crates/ml/src/cuda_pipeline/gpu_experience_collector.rs \ + crates/ml/src/trainers/dqn/trainer/training_loop.rs \ + crates/ml/src/trainers/dqn/config.rs \ + crates/ml/src/training_profile.rs +git commit -m "cleanup: remove dead v6 reward fields (loss_aversion, beta_penalty, regret_blend, trade_clustering)" +``` + +--- + +### Task 7: Update TOML Configs + Final Test Pass + +**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 v7.1 params to all TOMLs** + +Add to `[reward]` section of each file: + +```toml +exposure_aux_weight = 0.5 +exposure_aux_warmup_epochs = 5 +``` + +For `dqn-hyperopt.toml`, also add to `[search_space]`: + +```toml +exposure_aux_weight = [0.1, 1.0] +``` + +And to `[phase_fast]`: + +```toml +exposure_aux_weight = 0.5 +exposure_aux_warmup_epochs = 5 +``` + +- [ ] **Step 2: Run full test suite** + +Run: `SQLX_OFFLINE=true cargo test -p ml --lib 2>&1 | grep "^test result:"` +Expected: `ok. 891 passed; 0 failed` (or more — zero regressions) + +Fix any test expectation failures from config changes. + +- [ ] **Step 3: Run workspace build** + +Run: `SQLX_OFFLINE=true cargo check --workspace 2>&1 | tail -5` +Expected: `Finished` + +- [ ] **Step 4: Commit** + +```bash +git add config/training/ crates/ml/src/ +git commit -m "config(reward-v7.1): add exposure_aux_weight to all TOML profiles, final test pass" +``` diff --git a/docs/superpowers/specs/2026-04-07-reward-v8-master-plan.md b/docs/superpowers/specs/2026-04-07-reward-v8-master-plan.md new file mode 100644 index 000000000..9b755a629 --- /dev/null +++ b/docs/superpowers/specs/2026-04-07-reward-v8-master-plan.md @@ -0,0 +1,526 @@ +# Reward v8: Comprehensive Training Overhaul — Master Plan + +**Goal:** Break through the 30% win rate ceiling to 50%+ on ES futures by fixing the bootstrapping trap, adding per-branch gradient attribution, and implementing novel training techniques. Full CUDA, no CPU path, no memory copies, zero regressions. + +**Architecture:** 18 items across 5 categories: bug fixes (A), bootstrap trap (B), initialization (C), advanced techniques (D), dead code cleanup (E). All GPU-side. 4 new CUDA kernels, 3 modified kernels, ~12 new config fields, 4 removed dead fields. + +**Tech Stack:** CUDA (nvcc → cubin via build.rs), Rust (cudarc), cuBLAS SGEMM, TOML config + +--- + +## Category A: v7.1 Bug Fixes & Wiring (9 items) + +### A1-A5: Already Committed (Tasks 1-3) + +| Item | Description | Status | +|------|-------------|--------| +| A1 | Exit timing `closing_sign` | DONE (commit 28f8fc74b) | +| A2 | Kelly prior seeding (Bayesian) | DONE (commit 28f8fc74b) | +| A3 | Reward noise magnitude scaling | DONE (commit 28f8fc74b) | +| A4 | Best exposure output buffer (`out_best_exposure`) | DONE (commit 44ca6147a) | +| A5 | Exposure aux grad kernel (cubin) | DONE (commit 3ac6daf3d) | + +### A6: Exposure Auxiliary GEMM Wiring + +Wire the `exposure_aux_grad_kernel` (A5) into the training pipeline with a full backward GEMM for weight gradients. + +**Components:** + +**a) GpuDqnTrainer** (`gpu_dqn_trainer.rs`): +- Load `exposure_aux_grad_kernel` from DQN_UTILITY_CUBIN +- Allocate scratch buffer `exposure_aux_scratch: CudaSlice` [batch × b0 × num_atoms] +- Store `exposure_aux_grad_kernel: CudaFunction` field +- `launch_exposure_aux_grad(&mut self, targets, aux_weight)` method: + 1. Zero scratch + 2. Launch aux kernel → f32 logit gradients in scratch + 3. f32→bf16 cast (reuse `f32_to_bf16_kernel`) + 4. `backward_fc_layer(staging_bf16, save_h_b0, _, grad_w_b0out, grad_b_b0out, 0, b0*na, adv_h, batch)` — cuBLAS SGEMM for weight gradients + bias sum +- `exposure_bias_offset()` helper: `sum(compute_param_sizes()[..11])` (byte offset for b_b0out) + +**b) FusedTrainingCtx** (`fused_training.rs`): +- `exposure_aux_weight: f64` field +- `pending_exposure_targets: Option>` field +- Call `self.trainer.launch_exposure_aux_grad(targets, weight)` in `run_full_step()` between backward and Adam — same injection point as IQN/CQL + +**c) Training loop** (`training_loop.rs`): +- Compute decayed weight per epoch: + ```rust + let warmup = self.hyperparams.exposure_aux_warmup_epochs.max(1) as f64; + let progress = (self.current_epoch as f64 / warmup).min(1.0); + let aux_w = self.hyperparams.exposure_aux_weight * (1.0 - 0.9 * progress); + fused.exposure_aux_weight = aux_w; + ``` + +### A7: CEA 25% Linear Warmup + +In `training_loop.rs`, replace the current CEA warmup (hard cutoff at 3 epochs) with linear decay over 25% of training: + +```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 } +}, +``` + +### A8: OFI Epoch Gate + +Disable OFI reward multiplier for the first 5 epochs. In `training_loop.rs`: + +```rust +ofi_reward_weight: { + let base = self.hyperparams.ofi_reward_weight as f32; + if self.current_epoch < 5 { 0.0 } else { base } +}, +``` + +### A9: Dead v6 Code Removal + +Remove from `ExperienceCollectorConfig` (struct + Default + training_loop constructor): +- `loss_aversion: f32` +- `beta_penalty: f32` +- `regret_blend: f32` +- `trade_clustering_penalty: f32` + +Remove stale references in `training_profile.rs`. Keep the fields in `DQNHyperparameters` for TOML backward compat (zeroed defaults). + +--- + +## Category B: Bootstrap Trap Fixes (3 items) + +### B1: GPU-Side Cosine Epsilon Schedule + +Replace host-provided epsilon with GPU-computed cosine decay. + +**Kernel change** in `experience_action_select` (`experience_kernels.cu`): + +Add 3 new parameters: `float eps_start`, `float eps_end`, `int total_epochs`, `int current_epoch`. + +Inside the kernel, replace `float epsilon` read from param with: + +```cuda +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 * cosine_progress)); +``` + +Remove the old `float epsilon` parameter from `experience_env_step` — the action selection kernel now owns epsilon. + +**Rust wiring**: Pass `eps_start`, `eps_end`, `total_epochs`, `current_epoch` to the action select kernel launch. Update `ExperienceCollectorConfig` to hold these instead of a single `epsilon`. + +**Config**: `epsilon_start` (default 0.3, hyperopt [0.1, 0.5]), `epsilon_end` (default 0.02, hyperopt [0.01, 0.05]). + +### B2: n-step Default Increase + +Config-only change: +- Default `n_steps`: 3 → 5 +- Hyperopt range: [3, 5] → [3, 7] +- `nstep_kernel.cu` already supports arbitrary n. Zero kernel changes. + +### B3: Dense Directional Micro-Reward (Adaptive Vol) + +In `experience_env_step`, OUTSIDE the `segment_complete` block, BEFORE the drawdown penalty: + +```cuda +/* ── Dense micro-reward: directional bootstrap signal ── */ +if (fabsf(position) > 0.001f && micro_reward_scale > 0.0f) { + float price_change = (raw_next - raw_close) / fmaxf(raw_close, 1.0f); + float direction_signal = (position > 0.0f) ? price_change : -price_change; + /* Adaptive: smaller in volatile markets, larger in calm */ + float vol_ratio = vol_proxy / 0.005f; + float adaptive_scale = micro_reward_scale / fmaxf(sqrtf(vol_ratio), 0.5f); + reward += adaptive_scale * direction_signal / vol_proxy; +} +``` + +New kernel param: `float micro_reward_scale` (default 0.001, hyperopt [0.0, 0.005]). + +Note: `vol_proxy` is only computed inside `segment_complete`. For the micro-reward (outside that block), compute a local vol_proxy from ATR: + +```cuda +float micro_vol_proxy = 0.005f; /* default 0.5% */ +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); +} +``` + +--- + +## Category C: Initialization & Pre-Training (3 items) + +### C1: Supervised Pre-Training Kernel (Epoch 0) + +New CUDA kernel `exposure_pretrain_step` in `dqn_utility_kernels.cu`. Functionally identical to `exposure_aux_grad_kernel` but operates on raw bars instead of replay buffer samples: + +```cuda +extern "C" __global__ void exposure_pretrain_step( + const __nv_bfloat16* __restrict__ targets, /* [total_bars, 4] raw prices */ + 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; + + /* Compute target: which exposure maximizes single-bar PnL? */ + /* targets[bar, 2] = raw_close, targets[bar, 3] = raw_next */ + /* (bar index is implicit — caller maps batch samples to bar indices) */ + + /* Same cross-entropy gradient as exposure_aux_grad_kernel */ + /* ... (identical math, different data source) */ +} +``` + +**Training flow** (in `training_loop.rs`, before epoch 0): +1. Sample random batches from the GPU-resident feature/target data +2. cuBLAS forward → exposure branch logits +3. Launch `exposure_pretrain_step` → logit gradients in scratch +4. f32→bf16 cast + `backward_fc_layer` GEMM → weight gradients in grad_buf +5. `replay_adam()` → Adam update +6. Repeat for 100 batches (covers ~800K bars at batch=8192) +7. Normal RL training begins at epoch 1 + +This uses the SAME backward GEMM infrastructure as A6 — identical code path. + +### C2: Pessimistic Q-Value Initialization + +One-time operation in `GpuDqnTrainer::new()`, after weight initialization: + +```rust +// Shift value head output bias (b_v2) to slightly negative +// Makes V(s) ≈ -0.1 for all states → forces exploration +let sizes = compute_param_sizes(&self.config); +let b_v2_offset = sizes[..7].iter().sum::(); // offset to b_v2 +let pessimistic_bias = vec![-0.1_f32; self.config.num_atoms]; +unsafe { + cudarc::driver::sys::cuMemcpyHtoDAsync_v2( + self.grad_buf.raw_ptr() + (b_v2_offset * 4) as u64, // NO — write to params_f32, not grad_buf + pessimistic_bias.as_ptr().cast(), + pessimistic_bias.len() * 4, + self.stream.cu_stream(), + ); +} +// Then f32→bf16 cast to sync params_bf16 +``` + +Actually, write directly to the f32 master params buffer at the b_v2 bias offset, then cast f32→bf16. One HtoD copy (num_atoms × 4 bytes = ~400 bytes) + one cast kernel launch. Happens once at construction. + +### C3: Combined Schedule (Pre-train → Bootstrap → RL) + +All schedules computed GPU-side from `epoch` and `total_epochs`: + +| Phase | Epochs | aux_weight | CEA weight | Epsilon | OFI | Micro-reward | +|-------|--------|-----------|------------|---------|-----|--------------| +| Pre-train | 0 | 1.0 | 0.0 | 1.0 | off | off | +| Bootstrap | 1-25% | 0.8→0.1 | 1.0→base | cosine 0.3→0.05 | off (epoch<5) | on | +| RL-only | 25%+ | 0.05 floor | base | cosine →0.02 | on | on | + +The training loop passes `epoch` and `total_epochs` to both the experience kernel and the action select kernel. The kernels compute their own schedules internally — no CPU scheduling needed. + +--- + +## Category D: Advanced Techniques (4 items) + +### D1: TD(λ) Truncated Returns (Replaces n-step kernel) + +Replace `nstep_kernel.cu` with a TD(λ) kernel that computes geometrically-weighted multi-horizon returns: + +```cuda +extern "C" __global__ void td_lambda_kernel( + const float* __restrict__ rewards, /* [N * L] */ + const float* __restrict__ dones, /* [N * L] */ + const float* __restrict__ q_next, /* [N * L] bootstrapped Q(s') */ + float* __restrict__ out_returns, /* [N * L] overwritten with G^λ */ + float gamma, float lambda, + int max_trace, int L, int N +) { + 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; + + /* Truncated TD(λ): blend n-step returns for n=1..max_trace */ + float G_lambda = 0.0f; + float weight_sum = 0.0f; + float lambda_pow = 1.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++) { + if (dones[base + t + k] > 0.5f && k > 0) { valid = 0; break; } + G_n += gamma_pow * rewards[base + t + k]; + gamma_pow *= gamma; + } + if (!valid || (t + n) >= L) { + /* Terminal or out of bounds — use last valid step */ + G_lambda += lambda_pow * G_n; + weight_sum += lambda_pow; + break; + } + /* Add bootstrap */ + G_n += gamma_pow * q_next[base + t + n - 1]; + + /* Accumulate into λ-weighted sum */ + float w = (n < max_trace) ? (1.0f - lambda) * lambda_pow : lambda_pow; + G_lambda += w * G_n; + weight_sum += w; + lambda_pow *= lambda; + } + + out_returns[idx] = (weight_sum > 0.0f) ? G_lambda / weight_sum : rewards[idx]; +} +``` + +**Config**: `td_lambda` (default 0.9, hyperopt [0.5, 0.99]), `max_trace_length` (default 7, hyperopt [5, 10]). + +**Replaces**: the current `nstep_kernel.cu`. The old kernel is kept but unused (TD(λ) is a strict generalization). The Rust n-step launch code switches to the new kernel. + +### D2: PopArt Reward Normalization + +Adaptive reward normalization with weight rescaling for C51 stability. + +**GPU-resident running statistics** (2 f32 scalars + 1 counter): +```rust +popart_mean: CudaSlice, // [1] running mean of rewards +popart_var: CudaSlice, // [1] running variance +popart_count: CudaSlice, // [1] sample count +``` + +**New kernel** `popart_normalize_kernel`: +```cuda +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 variance */ + float* __restrict__ count_buf, /* [1] sample count */ + int N +) { + /* Thread 0: update running stats with Welford's algorithm */ + /* All threads: normalize rewards in-place */ + float mu = *mean_buf; + float sigma = sqrtf(*var_buf / fmaxf(*count_buf, 1.0f)); + sigma = fmaxf(sigma, 0.01f); + + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i < N) { + rewards[i] = (rewards[i] - mu) / sigma; + } +} +``` + +**Weight rescaling** (once per epoch, after stats update): adjust value head last layer: +``` +W_v2_new = W_v2_old * (sigma_old / sigma_new) +b_v2_new = (b_v2_old * sigma_old + mu_old - mu_new) / sigma_new +``` + +This is a SAXPY-like kernel on the value head weights. Launched once per epoch. Eliminates need for manual v_min/v_max tuning. + +**Config**: `popart_enabled` (default true), `popart_warmup` (default 100 — minimum samples before normalizing). + +### D3: Curriculum Learning (Regime-Ordered Episodes) + +**Pre-computation** (once at init): + +New kernel `compute_difficulty_scores`: +```cuda +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) return; + + /* Difficulty = inverse of directional clarity over 5 bars */ + 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); + /* Low difficulty = strong trend, high difficulty = choppy */ + scores[i] = 1.0f / (ret_5 + 0.001f); +} +``` + +**Episode ordering** (in Rust, once at init): +1. Download difficulty scores from GPU (one-time DtoH) +2. Sort bar indices by difficulty (ascending = easy first) +3. Upload sorted indices as `curriculum_starts: CudaSlice` + +**Per-epoch curriculum** (in training_loop.rs): +```rust +let curriculum_fraction = (self.current_epoch as f32 / self.hyperparams.epochs as f32) + .min(1.0).max(0.0); +// Phase 1 (0-30%): use easiest 30% of bars +// Phase 2 (30-60%): use easiest 60% +// Phase 3 (60%+): use all bars +let usable_bars = ((0.3 + 0.7 * curriculum_fraction) * total_bars as f32) as usize; +``` + +Modify `episode_starts` buffer to only sample from the first `usable_bars` of sorted indices. + +**Config**: `curriculum_enabled` (default true), `curriculum_warmup_fraction` (default 0.6 — fraction of training before full dataset). + +### D4: Hindsight Optimal Exit Relabeling + +New kernel `hindsight_relabel_kernel`: +```cuda +extern "C" __global__ void hindsight_relabel_kernel( + const __nv_bfloat16* __restrict__ targets, /* [total_bars, 4] raw prices */ + const int* __restrict__ actions, /* [batch] taken actions */ + const int* __restrict__ episode_starts, /* [batch] bar index of entry */ + float* __restrict__ rewards, /* [batch] in-place relabel */ + const float* __restrict__ relabel_mask, /* [batch] 1.0 = relabel, 0.0 = keep */ + int b0_size, int b1_size, int b2_size, + float max_position, int lookahead, + int total_bars, float vol_proxy_default, + int batch_size +) { + int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i >= batch_size) return; + if (relabel_mask[i] < 0.5f) return; /* keep original reward */ + + int bar = episode_starts[i]; + int exposure_idx = actions[i] / (b1_size * b2_size); + float pos = compute_target_position(exposure_idx, b0_size, max_position); + if (fabsf(pos) < 0.001f) return; /* flat — nothing to relabel */ + + float entry_price = __bfloat162float(targets[bar * 4 + 2]); + + /* Find optimal exit within lookahead */ + float best_pnl = -1e30f; + 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; + } + + /* Relabel reward with optimal exit PnL (vol-normalized, soft-clamped) */ + float opt_return = best_pnl / fmaxf(entry_price, 1.0f); + float opt_reward = asymmetric_soft_clamp(10.0f * opt_return / vol_proxy_default); + rewards[i] = opt_reward; +} +``` + +**Relabel mask**: Generated per batch using deterministic hash (same as reward noise): +```cuda +float mask = ((hash(i) & 0xFF) < (int)(hindsight_fraction * 255.0f)) ? 1.0f : 0.0f; +``` + +Launched AFTER batch sampling from PER, BEFORE C51 loss computation. Modifies `rewards_buf` in-place for the selected fraction. + +**Config**: `hindsight_fraction` (default 0.1, hyperopt [0.0, 0.3]), `hindsight_lookahead` (default 10, hyperopt [5, 20]). + +--- + +## Category E: Dead Code & Refactoring + +### E1: Remove Dead v6 Fields + +From `ExperienceCollectorConfig`: +- `loss_aversion: f32` → REMOVE (soft-clamp replaces) +- `beta_penalty: f32` → REMOVE (disabled, dead) +- `regret_blend: f32` → REMOVE (CEA replaces) +- `trade_clustering_penalty: f32` → REMOVE (removed in v7) + +From the kernel launch arg chain: remove corresponding `.arg()` lines. + +Keep in `DQNHyperparameters` (for TOML backward compat, zeroed defaults). + +### E2: Stale Code Comments + +Update all references to "reward v6" or "reward v7" to "reward v8" in kernel comments. + +--- + +## New Config Fields Summary + +| Field | Type | Default | Hyperopt Range | Category | +|-------|------|---------|---------------|----------| +| `exposure_aux_weight` | f64 | 0.5 | [0.1, 1.0] | A6 | +| `exposure_aux_warmup_epochs` | usize | 5 | [3, 10] | A6 | +| `epsilon_start` | f64 | 0.3 | [0.1, 0.5] | B1 | +| `epsilon_end` | f64 | 0.02 | [0.01, 0.05] | B1 | +| `micro_reward_scale` | f64 | 0.001 | [0.0, 0.005] | B3 | +| `td_lambda` | f64 | 0.9 | [0.5, 0.99] | D1 | +| `max_trace_length` | usize | 7 | [5, 10] | D1 | +| `popart_enabled` | bool | true | — | D2 | +| `curriculum_enabled` | bool | true | — | D3 | +| `curriculum_warmup_fraction` | f64 | 0.6 | [0.3, 0.8] | D3 | +| `hindsight_fraction` | f64 | 0.1 | [0.0, 0.3] | D4 | +| `hindsight_lookahead` | usize | 10 | [5, 20] | D4 | + +## Removed Fields + +| Field | Reason | +|-------|--------| +| `loss_aversion` | Replaced by asymmetric_soft_clamp | +| `beta_penalty` | Was disabled (0.0) | +| `regret_blend` | Replaced by CEA | +| `trade_clustering_penalty` | Removed in v7 | + +--- + +## Files Modified + +| File | Changes | +|------|---------| +| `experience_kernels.cu` | B1 (epsilon params), B3 (micro-reward), A4 updates | +| `dqn_utility_kernels.cu` | C1 (pretrain kernel), A5 (already done) | +| `nstep_kernel.cu` | D1 (TD(λ) replacement) | +| `gpu_dqn_trainer.rs` | A6 (aux GEMM), C2 (pessimistic init), D2 (PopArt buffers) | +| `gpu_experience_collector.rs` | B1 (epsilon wiring), B3 (micro-reward), A9 (dead fields), D3 (curriculum) | +| `fused_training.rs` | A6 (aux wiring), D4 (hindsight relabel) | +| `training_loop.rs` | A7 (CEA warmup), A8 (OFI gate), C1 (pretrain phase), C3 (schedule), D3 (curriculum) | +| `config.rs` | All new fields, dead field cleanup | +| `training_profile.rs` | Dead field refs cleanup | +| `config/training/*.toml` (×4) | All new params | + +## New CUDA Kernels (4) + +| Kernel | File | Purpose | +|--------|------|---------| +| `exposure_pretrain_step` | dqn_utility_kernels.cu | Supervised direction pre-training | +| `td_lambda_kernel` | nstep_kernel.cu (or new file) | Multi-horizon λ-returns | +| `popart_normalize_kernel` | dqn_utility_kernels.cu | Adaptive reward normalization | +| `compute_difficulty_scores` | experience_kernels.cu | Curriculum learning scores | +| `hindsight_relabel_kernel` | experience_kernels.cu | Optimal exit relabeling | + +## Testing + +- `cargo check --workspace` — zero compilation errors +- `cargo test -p ml --lib` — 891+ pass, 0 fail, 0 regressions +- Explicit v8 unit tests: + - `test_asymmetric_soft_clamp_properties` — f(0)=0, f(5)=5, f(-5)≈-3.93, f(15)=10 + - `test_kelly_prior_neutral` — with no trades, Kelly recommends 0.0 (no bet) + - `test_td_lambda_degenerates_to_nstep` — λ=0 gives 1-step, λ=1 gives MC + - `test_popart_normalization` — mean→0, std→1 after normalization + - `test_cosine_epsilon_schedule` — epoch 0 = eps_start, epoch N = eps_end + - `test_curriculum_ordering` — easy bars (trending) sorted before hard bars (choppy) + - `test_hindsight_relabel_improves_reward` — relabeled reward ≥ original for winners + - `test_micro_reward_adaptive_scale` — smaller in volatile, larger in calm + - `test_exposure_aux_gradient_unique` — each of 9 outputs gets different gradient + +## Performance Budget + +| Item | Per-step cost | Notes | +|------|--------------|-------| +| Aux GEMM (A6) | ~0.1ms | 1 SGEMM + 1 cast + 1 kernel | +| Micro-reward (B3) | ~0.01ms | 1 conditional per thread | +| Cosine epsilon (B1) | ~0ms | 2 float ops per thread | +| TD(λ) (D1) | ~0.5ms | Replaces n-step (~0.1ms), 5× more computation | +| PopArt (D2) | ~0.05ms | Per-batch normalization | +| Hindsight (D4) | ~0.1ms | 10% of batch × 10 bar lookahead | +| Pre-train (C1) | ~2s total | 100 batches × 20ms, epoch 0 only | +| Curriculum (D3) | ~0ms runtime | Sort at init only | +| **Total overhead** | **~0.8ms/step** | From 35s → ~36s per epoch (+3%) |