fix(rl): warmup window in plateau-decay LR controller fixes V cold-start

`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 <noreply@anthropic.com>

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? ────────────────
+    // ── 2. Warmup: unconditionally overwrite best while EMA settles.
-    const float best_prev = isv[best_slot];
+    //
-
+    // During the first LR_WARMUP_STEPS observations, `best` tracks
-    if (best_prev == 0.0f) {
+    // the current loss_ema. Without this, sparse-signal heads (V
-        // First-observation bootstrap for best (and counter).
+    // 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);
+                       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);
+    plateau_decay_head(isv, RL_LR_AUX_INDEX, observed_loss_aux, -1, -1, -1, -1);
 }

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_PI_WARMUP_COUNTER_INDEX, RL_LR_Q_BEST_LOSS_INDEX, RL_LR_Q_INDEX,
-    RL_LR_Q_STEPS_SINCE_BEST_INDEX, RL_LR_V_BEST_LOSS_INDEX, RL_LR_V_INDEX,
+    RL_LR_Q_LOSS_EMA_INDEX, RL_LR_Q_STEPS_SINCE_BEST_INDEX, RL_LR_Q_WARMUP_COUNTER_INDEX,
-    RL_LR_V_LOSS_EMA_INDEX, RL_LR_V_STEPS_SINCE_BEST_INDEX, RL_MEAN_ABS_PNL_EMA_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": {

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;

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")?;
         }