From d5c29fb4faecc518282ad35879c7039c9265787a Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sat, 23 May 2026 19:47:00 +0200 Subject: [PATCH] fix(rl): warmup window in plateau-decay LR controller fixes V cold-start MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit `alpha-rl-rzltn` exposed a bug in the plateau-decay design: V head's `best` got bootstrapped to 7.12e-10 (machine epsilon) at step 1 because V regression had no reward signal yet — no trade had closed, the bootstrap V target was 0, so the first V loss was effectively 0. Every subsequent V loss EMA was orders of magnitude higher (4.07 at step 100, 1.15 at step 1000), so the improvement check `loss_ema < best * 0.99` evaluated false FOREVER. The controller then decayed lr_v every 1000 steps purely on the patience clock, not because the model genuinely plateaued. Cross-check across the 50k-step rzltn run: * V best unique values: {0.0, 7.12e-10} — ONLY 2 across 50000 rows * V best max: 7.12e-10 * V best-improvements: 0 (Q: 12, π: 12) * V decays still fired: 7 (one every 1000 steps from step 1001) The plateau-decay mechanics worked correctly — the controller counted to 999 then halved LR exactly as designed. The bug was that "first observation defines best forever" is degenerate for sparse-signal heads whose first loss is a cold-start artifact. ## Fix: LR_WARMUP_STEPS Three new ISV slots (one per head — Q, π, V at 436/437/438) hold a monotonic warmup counter clamped at LR_WARMUP_STEPS = 500. During warmup the controller: * always overwrites `best` with current loss_ema (tracks the EMA as it converges) * holds the plateau counter at 0 (no decay fires during warmup) * increments warmup_counter Once warmup_counter >= LR_WARMUP_STEPS, the controller switches to standard plateau detection — `best` then locks in at the post-warmup loss_ema value (representative of the head's converged loss scale), and patience counting begins. At α=0.05 the EMA half-life is ~14 steps; 500 updates leaves ~35 half-lives, well past convergence. This gives V time to see its first actual losses after trades start closing. ## Slot allocation RL_SLOTS_END: 436 → 439 (adds 3 warmup counter slots). ## Wiring * rl_lr_controller.cu — adds warmup_slot param to plateau_decay_head, kernel takes 12 slot ints (was 9) * isv_slots.rs — 3 new constants, RL_SLOTS_END += 3 * integrated.rs — launch_rl_lr_controller passes 12 slot ints * alpha_rl_train.rs — diag JSONL emits new lr_plateau.{head}.warmup field ## Verified gates (local sm_86) G1 isv_bootstrap ✅ G3 controllers ✅ G4 target_update ✅ G6 r7d_per_wiring ✅ integrated_smoke ✅ Co-Authored-By: Claude Opus 4.7 --- crates/ml-alpha/cuda/rl_lr_controller.cu | 49 ++++++++++++++++------ crates/ml-alpha/examples/alpha_rl_train.rs | 16 ++++--- crates/ml-alpha/src/rl/isv_slots.rs | 22 +++++++++- crates/ml-alpha/src/trainer/integrated.rs | 13 +++++- 4 files changed, 78 insertions(+), 22 deletions(-) diff --git a/crates/ml-alpha/cuda/rl_lr_controller.cu b/crates/ml-alpha/cuda/rl_lr_controller.cu index 164a5a5bd..3897c2614 100644 --- a/crates/ml-alpha/cuda/rl_lr_controller.cu +++ b/crates/ml-alpha/cuda/rl_lr_controller.cu @@ -67,9 +67,21 @@ // 1000 steps ≈ 7 seconds at 145 steps/sec on L40S. // DECAY_FACTOR: shrink LR by this on each plateau fire. // 0.5 = "halve LR" (canonical ReduceLROnPlateau default). +// LR_WARMUP_STEPS: number of EMA updates before `best` is locked in +// and plateau detection turns on. Without this, a head whose +// first non-zero loss observation is a cold-start artifact (e.g. +// V loss ≈ machine-epsilon at step 0 before any reward has +// flowed) pins `best` to an unreachable value, and the controller +// decays on the patience clock alone. At α=0.05 the EMA +// half-life is ≈14 steps; 500 updates leaves ≈35 half-lives, +// which is well past convergence. Canonical incident: +// alpha-rl-rzltn fold0 (commit 13d81dc5e) — V `best` stuck at +// 7.12e-10 across all 50k steps, only 2 unique best values +// across the entire run. #define IMPROVEMENT_THRESHOLD 0.99f #define PLATEAU_PATIENCE 1000.0f #define DECAY_FACTOR 0.5f +#define LR_WARMUP_STEPS 500.0f __device__ __forceinline__ void plateau_decay_head( float* isv, @@ -77,7 +89,8 @@ __device__ __forceinline__ void plateau_decay_head( float observed_loss, int loss_ema_slot, int best_slot, - int counter_slot + int counter_slot, + int warmup_slot ) { const float lr_prev = isv[lr_idx]; @@ -109,16 +122,23 @@ __device__ __forceinline__ void plateau_decay_head( } isv[loss_ema_slot] = loss_ema_new; - // ── 2. Compare to best — improvement or plateau? ──────────────── - const float best_prev = isv[best_slot]; - - if (best_prev == 0.0f) { - // First-observation bootstrap for best (and counter). + // ── 2. Warmup: unconditionally overwrite best while EMA settles. + // + // During the first LR_WARMUP_STEPS observations, `best` tracks + // the current loss_ema. Without this, sparse-signal heads (V + // before any trade closes) lock `best` to a near-zero first + // observation that subsequent EMA values can never beat. + const float warmup_prev = isv[warmup_slot]; + if (warmup_prev < LR_WARMUP_STEPS) { isv[best_slot] = loss_ema_new; isv[counter_slot] = 0.0f; + isv[warmup_slot] = warmup_prev + 1.0f; return; } + // ── 3. Compare to best — improvement or plateau? ──────────────── + const float best_prev = isv[best_slot]; + if (loss_ema_new < best_prev * IMPROVEMENT_THRESHOLD) { // Improvement — update best, reset counter. isv[best_slot] = loss_ema_new; @@ -126,7 +146,7 @@ __device__ __forceinline__ void plateau_decay_head( return; } - // ── 3. Plateau — increment counter; maybe decay LR. ───────────── + // ── 4. Plateau — increment counter; maybe decay LR. ───────────── const float counter_next = isv[counter_slot] + 1.0f; if (counter_next >= PLATEAU_PATIENCE) { // Decay fires. Halve LR (clamped to MIN), reset counter, @@ -150,12 +170,15 @@ extern "C" __global__ void rl_lr_controller( int q_loss_ema_slot, int q_best_slot, int q_counter_slot, + int q_warmup_slot, int pi_loss_ema_slot, int pi_best_slot, int pi_counter_slot, + int pi_warmup_slot, int v_loss_ema_slot, int v_best_slot, - int v_counter_slot + int v_counter_slot, + int v_warmup_slot ) { if (threadIdx.x != 0 || blockIdx.x != 0) return; @@ -163,12 +186,12 @@ extern "C" __global__ void rl_lr_controller( // trainer doesn't compute per-step loss observations for them. // Pass loss_ema_slot=-1 to the plateau function which then // early-returns, holding lr_bce / lr_aux at LR_BOOTSTRAP. - plateau_decay_head(isv, RL_LR_BCE_INDEX, observed_loss_bce, -1, -1, -1); + plateau_decay_head(isv, RL_LR_BCE_INDEX, observed_loss_bce, -1, -1, -1, -1); plateau_decay_head(isv, RL_LR_Q_INDEX, observed_loss_q, - q_loss_ema_slot, q_best_slot, q_counter_slot); + q_loss_ema_slot, q_best_slot, q_counter_slot, q_warmup_slot); plateau_decay_head(isv, RL_LR_PI_INDEX, observed_loss_pi, - pi_loss_ema_slot, pi_best_slot, pi_counter_slot); + pi_loss_ema_slot, pi_best_slot, pi_counter_slot, pi_warmup_slot); plateau_decay_head(isv, RL_LR_V_INDEX, observed_loss_v, - v_loss_ema_slot, v_best_slot, v_counter_slot); - plateau_decay_head(isv, RL_LR_AUX_INDEX, observed_loss_aux, -1, -1, -1); + v_loss_ema_slot, v_best_slot, v_counter_slot, v_warmup_slot); + plateau_decay_head(isv, RL_LR_AUX_INDEX, observed_loss_aux, -1, -1, -1, -1); } diff --git a/crates/ml-alpha/examples/alpha_rl_train.rs b/crates/ml-alpha/examples/alpha_rl_train.rs index 1b601a33b..cf2c0d535 100644 --- a/crates/ml-alpha/examples/alpha_rl_train.rs +++ b/crates/ml-alpha/examples/alpha_rl_train.rs @@ -41,9 +41,10 @@ use ml_alpha::rl::isv_slots::{ RL_ADVANTAGE_VAR_RATIO_EMA_INDEX, RL_ENTROPY_COEF_INDEX, RL_ENTROPY_OBSERVED_EMA_INDEX, RL_GAMMA_INDEX, RL_KL_PI_EMA_INDEX, RL_LR_AUX_INDEX, RL_LR_BCE_INDEX, RL_LR_PI_BEST_LOSS_INDEX, RL_LR_PI_INDEX, RL_LR_PI_LOSS_EMA_INDEX, RL_LR_PI_STEPS_SINCE_BEST_INDEX, - RL_LR_Q_BEST_LOSS_INDEX, RL_LR_Q_INDEX, RL_LR_Q_LOSS_EMA_INDEX, - RL_LR_Q_STEPS_SINCE_BEST_INDEX, RL_LR_V_BEST_LOSS_INDEX, RL_LR_V_INDEX, - RL_LR_V_LOSS_EMA_INDEX, RL_LR_V_STEPS_SINCE_BEST_INDEX, RL_MEAN_ABS_PNL_EMA_INDEX, + RL_LR_PI_WARMUP_COUNTER_INDEX, RL_LR_Q_BEST_LOSS_INDEX, RL_LR_Q_INDEX, + RL_LR_Q_LOSS_EMA_INDEX, RL_LR_Q_STEPS_SINCE_BEST_INDEX, RL_LR_Q_WARMUP_COUNTER_INDEX, + RL_LR_V_BEST_LOSS_INDEX, RL_LR_V_INDEX, RL_LR_V_LOSS_EMA_INDEX, + RL_LR_V_STEPS_SINCE_BEST_INDEX, RL_LR_V_WARMUP_COUNTER_INDEX, RL_MEAN_ABS_PNL_EMA_INDEX, RL_MEAN_TRADE_DURATION_EMA_INDEX, RL_N_ROLLOUT_STEPS_INDEX, RL_PER_ALPHA_INDEX, RL_PI_GRAD_NORM_EMA_INDEX, RL_PPO_CLIP_INDEX, RL_Q_DIVERGENCE_EMA_INDEX, RL_Q_GRAD_NORM_EMA_INDEX, RL_REWARD_SCALE_INDEX, RL_TARGET_TAU_INDEX, RL_TD_KURTOSIS_EMA_INDEX, @@ -528,13 +529,16 @@ fn main() -> Result<()> { "lr_plateau": { "q": { "loss_ema": isv[RL_LR_Q_LOSS_EMA_INDEX], "best": isv[RL_LR_Q_BEST_LOSS_INDEX], - "stale": isv[RL_LR_Q_STEPS_SINCE_BEST_INDEX] }, + "stale": isv[RL_LR_Q_STEPS_SINCE_BEST_INDEX], + "warmup": isv[RL_LR_Q_WARMUP_COUNTER_INDEX] }, "pi": { "loss_ema": isv[RL_LR_PI_LOSS_EMA_INDEX], "best": isv[RL_LR_PI_BEST_LOSS_INDEX], - "stale": isv[RL_LR_PI_STEPS_SINCE_BEST_INDEX] }, + "stale": isv[RL_LR_PI_STEPS_SINCE_BEST_INDEX], + "warmup": isv[RL_LR_PI_WARMUP_COUNTER_INDEX] }, "v": { "loss_ema": isv[RL_LR_V_LOSS_EMA_INDEX], "best": isv[RL_LR_V_BEST_LOSS_INDEX], - "stale": isv[RL_LR_V_STEPS_SINCE_BEST_INDEX] }, + "stale": isv[RL_LR_V_STEPS_SINCE_BEST_INDEX], + "warmup": isv[RL_LR_V_WARMUP_COUNTER_INDEX] }, }, "replay_len": trainer.replay.len(), "rewards": { diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index f31f398d8..f5256f77f 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -175,6 +175,26 @@ pub const RL_LR_V_LOSS_EMA_INDEX: usize = 433; pub const RL_LR_V_BEST_LOSS_INDEX: usize = 434; pub const RL_LR_V_STEPS_SINCE_BEST_INDEX: usize = 435; +// Warmup counters for the plateau-decay controller. Without these, +// a head whose first non-zero loss observation is a cold-start +// artifact (e.g. V loss = 7e-10 at step 0 before any reward has +// flowed) pins `best` to that degenerate value, and the improvement +// check `loss_ema < best * 0.99` evaluates false forever — the +// controller then decays the LR purely on the patience clock, +// not because the model genuinely plateaued. The warmup window +// unconditionally overwrites `best` with the current loss EMA for +// the first `LR_WARMUP_STEPS` observations, locking in a +// representative reference once the EMA has had time to converge +// past cold-start. Canonical incident: alpha-rl-rzltn (commit +// 13d81dc5e) — V head best stuck at 7.12e-10 across all 50k steps, +// only 2 unique best values across the entire run. +// +// Stored as f32 (mantissa precision to 16M — well beyond +// LR_WARMUP_STEPS). +pub const RL_LR_Q_WARMUP_COUNTER_INDEX: usize = 436; +pub const RL_LR_PI_WARMUP_COUNTER_INDEX: usize = 437; +pub const RL_LR_V_WARMUP_COUNTER_INDEX: usize = 438; + /// Last RL-allocated slot index (exclusive). The integrated trainer /// extends `ISV_TOTAL_DIM` to at least this value at trainer init time. -pub const RL_SLOTS_END: usize = 436; +pub const RL_SLOTS_END: usize = 439; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index 63657c709..542d3f1aa 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -2946,7 +2946,7 @@ impl IntegratedTrainer { /// /// Loss observations are passed as scalar floats (host-side /// values from the trainer's last step). Per-head state lives - /// in ISV slots 427..436 (loss_ema, best, counter ×3). + /// in ISV slots 427..439 (loss_ema, best, counter, warmup ×3). fn launch_rl_lr_controller( &self, observed_loss_q: f32, @@ -2966,18 +2966,24 @@ impl IntegratedTrainer { crate::rl::isv_slots::RL_LR_Q_BEST_LOSS_INDEX as i32; let q_counter_slot: i32 = crate::rl::isv_slots::RL_LR_Q_STEPS_SINCE_BEST_INDEX as i32; + let q_warmup_slot: i32 = + crate::rl::isv_slots::RL_LR_Q_WARMUP_COUNTER_INDEX as i32; let pi_loss_ema_slot: i32 = crate::rl::isv_slots::RL_LR_PI_LOSS_EMA_INDEX as i32; let pi_best_slot: i32 = crate::rl::isv_slots::RL_LR_PI_BEST_LOSS_INDEX as i32; let pi_counter_slot: i32 = crate::rl::isv_slots::RL_LR_PI_STEPS_SINCE_BEST_INDEX as i32; + let pi_warmup_slot: i32 = + crate::rl::isv_slots::RL_LR_PI_WARMUP_COUNTER_INDEX as i32; let v_loss_ema_slot: i32 = crate::rl::isv_slots::RL_LR_V_LOSS_EMA_INDEX as i32; let v_best_slot: i32 = crate::rl::isv_slots::RL_LR_V_BEST_LOSS_INDEX as i32; let v_counter_slot: i32 = crate::rl::isv_slots::RL_LR_V_STEPS_SINCE_BEST_INDEX as i32; + let v_warmup_slot: i32 = + crate::rl::isv_slots::RL_LR_V_WARMUP_COUNTER_INDEX as i32; let mut launch = self.stream.launch_builder(&self.rl_lr_controller_fn); launch .arg(&self.isv_d) @@ -2989,12 +2995,15 @@ impl IntegratedTrainer { .arg(&q_loss_ema_slot) .arg(&q_best_slot) .arg(&q_counter_slot) + .arg(&q_warmup_slot) .arg(&pi_loss_ema_slot) .arg(&pi_best_slot) .arg(&pi_counter_slot) + .arg(&pi_warmup_slot) .arg(&v_loss_ema_slot) .arg(&v_best_slot) - .arg(&v_counter_slot); + .arg(&v_counter_slot) + .arg(&v_warmup_slot); unsafe { launch.launch(cfg).context("rl_lr_controller launch")?; }