docs: reward v7.1 spec — 6 critical fixes for exposure degeneracy and training quality

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
2026-04-07 22:21:11 +02:00
parent b7c79f8a5d
commit 4ab7af5713
1 changed files with 257 additions and 0 deletions
--- a/docs/superpowers/specs/2026-04-07-reward-v71-fixes-design.md
+++ b/docs/superpowers/specs/2026-04-07-reward-v71-fixes-design.md
@@ -0,0 +1,257 @@
+# Reward v7.1: Six Critical Fixes Design
+
+**Goal:** Fix 6 concerns identified in the initial v7 hyperopt run: i%3 Q-value degeneracy (exposure auxiliary loss), win rate below random (OFI epoch gate), CEA warmup too short, exit timing prev_sign bug, Kelly threshold with prior seeding, and reward noise magnitude scaling.
+
+**Architecture:** Exposure auxiliary loss follows the existing IQN/CQL auxiliary gradient pattern — new lightweight CUDA kernel launched between C51 backward and Adam. Other 5 fixes are surgical changes to the experience kernel and Rust training loop config gating.
+
+**Tech Stack:** CUDA (nvcc → cubin), Rust (cudarc), TOML config
+
+---
+
+## Fix 1: Exposure Auxiliary Loss with Scheduled Decay
+
+### Problem
+C51 backward gives identical `-1/9 × dL` gradient to all 8 non-selected exposure outputs. This causes them to converge into degenerate i%3 clusters. The pattern recurs even after CEA warmup breaks it temporarily.
+
+### Root Cause
+The dueling advantage mean subtraction `A'[i] = A[i] - mean(A)` creates gradient `dL/dA[selected] = 8/9 × dL` and `dL/dA[other] = -1/9 × dL`. With b0_size=9, the non-selected gradient is only 12.5% of the selected — too weak to break weight clusters from random initialization.
+
+### Solution: Three Components
+
+**1a. Experience kernel output** — In the CEA loop (already computing rewards for all 9 exposures), track `best_exposure_idx` and write to a new output buffer `out_best_exposure[N*L]`:
+
+```cuda
+int best_k = 0;
+float best_alt = -1e30f;
+for (int k = 0; k < b0_size; k++) {
+    // ...existing CEA loop body...
+    if (alt_reward > best_alt) { best_alt = alt_reward; best_k = k; }
+}
+if (segment_complete && segment_hold_time > 0.0f) {
+    out_best_exposure[out_off] = best_k;
+} else {
+    out_best_exposure[out_off] = -1;  // no target (not a trade exit)
+}
+```
+
+New kernel parameter: `int* out_best_exposure` added to `experience_env_step` signature.
+
+**1b. New CUDA kernel** — `exposure_aux_grad_kernel` in `dqn_utility_kernels.cu`:
+
+```cuda
+extern "C" __global__ void exposure_aux_grad_kernel(
+    const float* __restrict__ adv_logits,  // [batch, b0_size * num_atoms] exposure advantage logits
+    const int*   __restrict__ targets,     // [batch] best exposure idx (-1 = skip)
+    float*       __restrict__ grad_buf,    // [total_params] accumulate gradient
+    int          adv_offset,               // offset into grad_buf for exposure adv weights
+    int          b0_size,
+    int          num_atoms,
+    float        aux_weight,
+    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
+
+    // Cross-entropy gradient for exposure branch:
+    // For each atom z: grad[a][z] = softmax(logits[a][z]) - 1{a == target}
+    const float* logits = adv_logits + (long long)sample * b0_size * num_atoms;
+
+    for (int z = 0; z < num_atoms; z++) {
+        // Compute softmax over b0_size for this atom
+        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);
+        }
+
+        // Write gradients
+        for (int a = 0; a < b0_size; a++) {
+            float prob = expf(logits[a * num_atoms + z] - max_l) / sum_exp;
+            float grad = prob - ((a == target) ? 1.0f : 0.0f);
+            // Accumulate into grad_buf at the exposure advantage weight offset
+            int param_idx = adv_offset + a * num_atoms + z;
+            atomicAdd(&grad_buf[param_idx], aux_weight * grad / (float)batch_size);
+        }
+    }
+}
+```
+
+**1c. Fused training pipeline** — Launch between C51 backward and Adam:
+- Store `best_exposure` buffer in GPU replay buffer alongside actions/rewards
+- In `submit_aux_ops()` or as a standalone post-backward step, launch `exposure_aux_grad_kernel`
+- Weight schedule in training_loop.rs:
+  ```rust
+  let warmup = self.hyperparams.exposure_aux_warmup_epochs.max(1) as f32;
+  let progress = (self.current_epoch as f32 / warmup).min(1.0);
+  let aux_w = self.hyperparams.exposure_aux_weight as f32 * (1.0 - 0.9 * progress);
+  // Decays from base (0.5) to 10% of base (0.05) over warmup_epochs
+  ```
+
+### Config
+- `exposure_aux_weight: f64` — default 0.5, hyperopt [0.1, 1.0]
+- `exposure_aux_warmup_epochs: usize` — default 5, hyperopt [3, 10]
+
+---
+
+## Fix 2: OFI Confidence Epoch Gate
+
+### Problem
+OFI multiplier amplifies wrong-direction trades in early training, reducing win rate below random (24% vs 30%).
+
+### Solution
+Gate OFI weight by epoch 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 }
+},
+```
+
+No kernel changes. Model learns direction first (epochs 0-4), then OFI amplifies informed trades (epoch 5+).
+
+---
+
+## Fix 3: Extended CEA Warmup
+
+### Problem
+3-epoch hard cutoff warmup breaks symmetry but model reverts when CEA drops from 1.0 to 0.3 at epoch 4.
+
+### Solution
+Linear blend 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 * (1.0 - progress) + base * progress
+    } else {
+        base
+    }
+},
+```
+
+For 20 epochs: warmup spans epochs 0-4, linearly blending 1.0→base. Smooth transition, no cliff.
+
+---
+
+## Fix 4: Exit Timing `prev_sign` Bug
+
+### Problem
+After a reversal (S100→L50), `prev_sign` holds the pre-trade direction. Exit timing quality uses `prev_sign` to determine which direction was closed, but for reversals, it references the OLD position, not the closing segment.
+
+### Solution
+`prev_sign` is computed at the top of the kernel from the pre-trade position — it IS the closing segment's direction. The variable name is confusing but the VALUE is correct for exits. For reversals, `prev_sign` is the direction being closed (the old position), which IS what we want.
+
+However, the current code runs the timing block AFTER the reversal block has already updated `position` and `entry_price`. The `prev_sign` variable itself is still correct (computed from `pre_trade_position` earlier). But for clarity, save it explicitly:
+
+```cuda
+int closing_sign = prev_sign;  // capture before any trade logic
+
+// ... later, in Layer 7 ...
+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;
+}
+```
+
+This makes the intent explicit and guards against future refactors that might move `prev_sign` computation.
+
+---
+
+## Fix 5: Kelly Threshold with Bayesian Prior
+
+### Problem
+Kelly signal requires 5+ completed trades before computing. Early in training, episodes rarely accumulate enough trades for this threshold.
+
+### Solution
+Seed with a Bayesian prior (Jeffrey's prior: 2 pseudo-wins + 2 pseudo-losses with 1% average return each). This makes Kelly available from the first real trade while regularizing toward "no bet" with limited data:
+
+```cuda
+if (kelly_sizing_weight > 0.0f) {
+    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;
+
+    // Always fires (eff_total >= 4.0 from priors)
+    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;
+    }
+}
+```
+
+With prior only: Kelly `f* = (1.0 × 0.5 - 0.5) / 1.0 = 0.0` (recommends no bet). As real trades accumulate, converges to empirical Kelly.
+
+---
+
+## Fix 6: Reward Noise Magnitude Scaling
+
+### Problem
+Fixed noise scale (±0.05) is 10-50% of small reward signals, drowning meaningful differences.
+
+### Solution
+Scale noise to a fixed PERCENTAGE of reward magnitude with a floor:
+
+```cuda
+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;
+    float noise_magnitude = fmaxf(fabsf(reward) * reward_noise_scale, 0.01f);
+    reward += noise_magnitude * pseudo_noise;
+}
+```
+
+`reward_noise_scale` becomes a percentage (0.05 = 5% of reward magnitude). A +5.0 reward gets ±0.25 noise. A +0.1 reward gets ±0.01 noise (floor). Constant SNR across magnitudes.
+
+---
+
+## Files Modified
+
+| File | Changes |
+|------|---------|
+| `crates/ml/src/cuda_pipeline/experience_kernels.cu` | Fixes 1a, 4, 5, 6 + new `out_best_exposure` arg |
+| `crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu` | Fix 1b: new `exposure_aux_grad_kernel` |
+| `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` | Fix 1c: load + launch aux kernel |
+| `crates/ml/src/trainers/dqn/fused_training.rs` | Fix 1c: wire aux kernel in training step |
+| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | Fixes 2, 3: OFI gate + CEA warmup schedule |
+| `crates/ml/src/trainers/dqn/config.rs` | New config fields |
+| `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` | New `out_best_exposure` buffer + arg wiring |
+| `config/training/dqn-*.toml` (all 4) | New params |
+
+## Testing
+
+- `cargo check --workspace` — build verification
+- `cargo test -p ml --lib` — all 891+ tests pass
+- Smoke test: verify exposure Q-values do NOT show i%3 pattern by epoch 3
+- Smoke test: verify win rate ≥ 30% by epoch 5 (above random baseline)
+- Smoke test: verify Kelly signal fires from epoch 1 (prior seeding)