docs: reward v8 master plan — 18-item comprehensive training overhaul
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
493
docs/superpowers/plans/2026-04-07-reward-v7-gems-pearls.md
Normal file
493
docs/superpowers/plans/2026-04-07-reward-v7-gems-pearls.md
Normal file
@@ -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"
|
||||
```
|
||||
602
docs/superpowers/plans/2026-04-07-reward-v7.md
Normal file
602
docs/superpowers/plans/2026-04-07-reward-v7.md
Normal file
@@ -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"
|
||||
```
|
||||
647
docs/superpowers/plans/2026-04-07-reward-v71-fixes.md
Normal file
647
docs/superpowers/plans/2026-04-07-reward-v71-fixes.md
Normal file
@@ -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<f32>` (line ~529). Add after `done_out`:
|
||||
|
||||
```rust
|
||||
best_exposure_out: CudaSlice<i32>, // [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::<i32>(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<i32> {
|
||||
&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<i32>,
|
||||
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<CudaSlice<i32>>, // 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<i32>, batch_size: usize) -> Result<()> {
|
||||
if self.best_exposure_targets.is_none() {
|
||||
self.best_exposure_targets = Some(
|
||||
self.stream.alloc_zeros::<i32>(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"
|
||||
```
|
||||
526
docs/superpowers/specs/2026-04-07-reward-v8-master-plan.md
Normal file
526
docs/superpowers/specs/2026-04-07-reward-v8-master-plan.md
Normal file
@@ -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<f32>` [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<CudaSlice<i32>>` 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::<usize>(); // 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<f32>, // [1] running mean of rewards
|
||||
popart_var: CudaSlice<f32>, // [1] running variance
|
||||
popart_count: CudaSlice<f32>, // [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<i32>`
|
||||
|
||||
**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%) |
|
||||
Reference in New Issue
Block a user