From 4ab7af571373c39b54c96f2df3eda25faa52f2cc Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Tue, 7 Apr 2026 22:21:11 +0200 Subject: [PATCH] =?UTF-8?q?docs:=20reward=20v7.1=20spec=20=E2=80=94=206=20?= =?UTF-8?q?critical=20fixes=20for=20exposure=20degeneracy=20and=20training?= =?UTF-8?q?=20quality?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Co-Authored-By: Claude Opus 4.6 (1M context) --- .../2026-04-07-reward-v71-fixes-design.md | 257 ++++++++++++++++++ 1 file changed, 257 insertions(+) create mode 100644 docs/superpowers/specs/2026-04-07-reward-v71-fixes-design.md diff --git a/docs/superpowers/specs/2026-04-07-reward-v71-fixes-design.md b/docs/superpowers/specs/2026-04-07-reward-v71-fixes-design.md new file mode 100644 index 000000000..38b736775 --- /dev/null +++ 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)