From 7064c9269e59de13d7a6da4d9d91d3d35890163e Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sat, 30 May 2026 23:19:20 +0200 Subject: [PATCH] fix(rl): un-freeze adaptive WIN/LOSS clamp + reset per-batch state on fold MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Two compounding bugs found via cluster diag at alpha-rl-gwmkf step 4083: Fix C — reward-clamp writeback un-freeze: rewards.clip_rate_ema = 0.9999 — essentially every win clamped to +1. With R-multiple 43 trade outcomes ($7k avg win / $163 avg loss), V regression saw identical targets for +$500 wins and +$42k wins; PPO advantage lost magnitude signal; policy stagnated at qpa = -0.95. Root cause: a 2026-05-24 "G.2" patch in rl_reward_clamp_controller.cu added `(void) margin;` and froze WIN=1.0, LOSS=3.0 over concern that simultaneously adapting clamp bounds + atom span could create a positive feedback loop. Empirically refuted — MARGIN is bounded [1.0, 5.0] and pos_max_ema is bounded by reward_scale × raw_PnL, so WIN has natural structural ceiling. The freeze was over-conservative. Restore the writeback: WIN = max(MIN_WIN, MARGIN × pos_max_ema), LOSS = max(MIN_WIN, WIN × RATIO). Atom span Step 5 then follows via the slow EWMA already in place. Local b=128 1k smoke confirms: clip_rate_ema: 0.9999 → 0.094 (target 5%, near hit) WIN clamp: 1.0 → 245.6 V_max atom: 1.0 → 1929 qpa: -0.95 → +0.71 (Q and π aligned!) mean_active: +$47k stays positive Fix D — reset_session_state extends to per-batch buffers: The 39efacf77 per-batch CMDP refactor added 4 device buffers but reset_session_state() was only zeroing the 4 ISV summary slots. On train→eval fold transition, eval inherited all the accumulated dead/cooldown accounts from training — eval started with a poisoned fleet. reset_session_state now also memsets session_pnl_per_batch_d, session_dd_triggered_per_batch_d, consec_loss_per_batch_d, and cooldown_remaining_per_batch_d to zero. Plus also resets the new RL_SESSION_PNL_WORST_INDEX slot. Validation: 13/13 risk_stack_invariants pass, 20/20 trade_management pass, integrated_trainer_smoke end-to-end passes. --- .../cuda/rl_reward_clamp_controller.cu | 34 ++++++++++++++++--- crates/ml-alpha/src/trainer/integrated.rs | 29 ++++++++++++++++ 2 files changed, 58 insertions(+), 5 deletions(-) diff --git a/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu b/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu index 133d33dac..97fec4492 100644 --- a/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu +++ b/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu @@ -235,11 +235,35 @@ extern "C" __global__ void rl_reward_clamp_controller( isv[RL_REWARD_CLAMP_RATIO_INDEX] = ratio_clamped; } - // Step 4 (G.2): WIN / LOSS clamp bounds remain STRUCTURAL — the - // G.2 feedback loop required BOTH clamp widening AND span widening. - // With clamp frozen here, only the atom span adapts (Step 5 below). - // Suppress unused warnings on diagnostic-only state. - (void) margin; + // ── Step 4: Adaptive WIN / LOSS clamp writeback. ────────────── + // + // History: a 2026-05-24 G.2 patch froze the clamp bounds at the + // trainer-seeded WIN=1.0, LOSS=3.0 over concern that simultaneously + // adapting clamps + atom span could form a positive feedback loop. + // + // Empirical refutation (alpha-rl-gwmkf step 4083, 2026-05-30): + // with R-multiple 43 trade distributions (avg_win $7062, avg_loss + // $163) the frozen WIN=1.0 caused clip_rate_ema = 0.9999 — + // essentially every win gets clamped, V regression sees identical + // targets for +$500 wins and +$42k wins, and PPO advantage loses + // magnitude signal. Diagnosis: starved policy gradient. + // + // Safety argument: the MARGIN controller above already saturates + // at MAX_MARGIN_ISV (bounded ≤ 5.0). pos_max_ema is bounded by + // `reward_scale × raw_PnL`, both finite. So WIN ≤ MAX_MARGIN × + // MAX_REWARD_SCALE × MAX_RAW_PNL is structurally bounded. + // + // Atom span (Step 5) follows the new WIN via slow EWMA (α=0.001, + // half-life ~700 steps), so resolution adjusts gradually as the + // reward distribution evolves rather than whipsawing. + if (ema_new > 0.0f) { + const float min_win_isv = isv[RL_REWARD_CLAMP_MIN_WIN_INDEX]; + const float win_new = fmaxf(min_win_isv, margin * ema_new); + const float ratio = isv[RL_REWARD_CLAMP_RATIO_INDEX]; + const float loss_new = fmaxf(min_win_isv, win_new * ratio); + isv[RL_REWARD_CLAMP_WIN_INDEX] = win_new; + isv[RL_REWARD_CLAMP_LOSS_INDEX] = loss_new; + } // ── Step 5: C51 atom span adaptation from observed reward EMAs. ── // diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index 12b8a5869..38a079d06 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -2961,9 +2961,18 @@ impl IntegratedTrainer { /// are intentionally NOT reset: training-acquired statistics should /// carry into eval (matches "policy learned during train, applied /// in eval" intent). + /// + /// Fix D (2026-05-30): also zero the per-batch state buffers added + /// by the CMDP per-batch refactor (39efacf77). Without this, the + /// summary slots are reset to 0 but per-batch arrays carry over, + /// so on the next launch the CMDP kernel re-overwrites the summary + /// slots with the stale aggregates. Net effect was eval phases + /// inheriting all the dead/cooldown accounts from training. pub fn reset_session_state(&self) -> Result<()> { + // Summary slots — direct ISV writes. for (slot, value) in [ (crate::rl::isv_slots::RL_SESSION_PNL_USD_INDEX, 0.0_f32), + (crate::rl::isv_slots::RL_SESSION_PNL_WORST_INDEX, 0.0_f32), (crate::rl::isv_slots::RL_SESSION_DD_TRIGGERED_INDEX, 0.0_f32), (crate::rl::isv_slots::RL_CONSEC_LOSS_COUNT_INDEX, 0.0_f32), (crate::rl::isv_slots::RL_COOLDOWN_REMAINING_STEPS_INDEX, 0.0_f32), @@ -2982,6 +2991,26 @@ impl IntegratedTrainer { ).map_err(|e| anyhow::anyhow!("rl_isv_write (reset_session_state slot {slot}): {:?}", e))?; } } + // Per-batch buffers — bulk memset to zero. Each is [b_size] f32. + for (buf_ptr, buf_bytes, name) in [ + (self.session_pnl_per_batch_d.raw_ptr(), + self.session_pnl_per_batch_d.num_bytes(), + "session_pnl_per_batch_d"), + (self.session_dd_triggered_per_batch_d.raw_ptr(), + self.session_dd_triggered_per_batch_d.num_bytes(), + "session_dd_triggered_per_batch_d"), + (self.consec_loss_per_batch_d.raw_ptr(), + self.consec_loss_per_batch_d.num_bytes(), + "consec_loss_per_batch_d"), + (self.cooldown_remaining_per_batch_d.raw_ptr(), + self.cooldown_remaining_per_batch_d.num_bytes(), + "cooldown_remaining_per_batch_d"), + ] { + unsafe { + raw_memset_d8_zero(buf_ptr, buf_bytes, self.raw_stream) + .map_err(|e| anyhow::anyhow!("zero {name}: {:?}", e))?; + } + } Ok(()) }