From 0c8cb6ad5b9e07ddf8c073a68009f6daa15dbc48 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sun, 31 May 2026 01:52:29 +0200 Subject: [PATCH] fix(rl): v9 defensive eval-boundary calibration MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Implements docs/superpowers/specs/2026-05-31-v9-defensive-eval-boundary-calibration.md. Closes the [[pearl_adaptive_carryover_discipline]] gap: every adaptive EMA must reset OR re-bootstrap at regime boundaries, never selectively preserve train-acquired statistics. Layer 1 — full EMA reset in `reset_session_state`: - win_rate, avg_win/loss EMAs → neutral sentinels - cumulative_dones → 0 (re-enter Kelly bootstrap) - inventory_beta + inventory_variance EMAs → 0 - reward_clamp pos/neg max EMAs + clip_rate EMA → 0 - Kelly fraction → 1.0 (bootstrap) Lifts Fix D's deliberate carryover (commit 7064c9269), which was diagnosed in v8 fold-1 eval as the primary -$507k driver: agent entered eval with train wr=0.34 EMA and took aggressively-sized losing trades while EMA decayed to true eval wr=0.25. Layer 2 — defensive warmup window: - New `rl_eval_warmup_decay` CUDA kernel: single-thread single-block, no atomics, mapped-pinned only. Runs after fused controllers each step. Overrides 4 risk-sizing floors (Kelly safety_frac, IQN τ_min, entropy_coef_min, PPO clip ε_min) with conservative defensive values for the first 500 steps post-boundary, with linear decay over the final 200 steps back to normal targets. - 11 new ISV slots (685-695): remaining counter, configured durations, defensive overrides (0.25/0.30/0.05/0.10), normal targets (0.50/0.10/0.01/0.05). - Sentinel counter = -1 at boot → kernel is no-op until reset arms it. All boundary semantics live in ISV; no hardcoded constants in the kernel. Build verified, all 13 risk-stack invariants + 5 controller adaptive-floor tests + integrated-trainer smoke pass on RTX 3050. --- crates/ml-alpha/build.rs | 1 + crates/ml-alpha/cuda/rl_eval_warmup_decay.cu | 105 +++++++++++++++++++ crates/ml-alpha/src/rl/isv_slots.rs | 48 ++++++++- crates/ml-alpha/src/trainer/integrated.rs | 94 ++++++++++++++++- 4 files changed, 244 insertions(+), 4 deletions(-) create mode 100644 crates/ml-alpha/cuda/rl_eval_warmup_decay.cu diff --git a/crates/ml-alpha/build.rs b/crates/ml-alpha/build.rs index a377e00fe..763fdb057 100644 --- a/crates/ml-alpha/build.rs +++ b/crates/ml-alpha/build.rs @@ -111,6 +111,7 @@ const KERNELS: &[&str] = &[ "rl_iqn_action_tau_controller", // Adaptive risk management (2026-05-30): Layer 2 — adapts IQN action-selection quantile τ from session drawdown signal "rl_inventory_beta_controller", // Adaptive risk management (2026-05-30): Layer 3 — adapts inventory penalty β from observed inventory variance vs reward magnitude "rl_kelly_fraction_controller", // Adaptive risk management (2026-05-30): Layer 4 — half-Kelly position-size multiplier from observed win_rate × R-multiple; warmup-gated + "rl_eval_warmup_decay", // v9 (2026-05-31): defensive eval-boundary calibration — overrides Kelly/τ_min/entropy_min/ε_min during warmup, linearly decays back to normal; runs every step "rl_ensemble_action_value", // C51+IQN ensemble: E_ensemble = α×E_C51 + (1-α)×E_IQN; α from ISV[544] "rl_noisy_linear_forward", // NoisyNet: factored noisy linear forward — y = (mu_w + sigma_w ⊙ eps_w) × x + (mu_b + sigma_b ⊙ eps_b); state-dependent exploration for C51/IQN final projection "rl_noisy_linear_backward", // NoisyNet: factored noisy linear backward — grad_mu_w/sigma_w/mu_b/sigma_b per-batch scratch for reduce_axis0 diff --git a/crates/ml-alpha/cuda/rl_eval_warmup_decay.cu b/crates/ml-alpha/cuda/rl_eval_warmup_decay.cu new file mode 100644 index 000000000..e21f3ba1c --- /dev/null +++ b/crates/ml-alpha/cuda/rl_eval_warmup_decay.cu @@ -0,0 +1,105 @@ +// rl_eval_warmup_decay.cu — defensive eval-boundary calibration (v9). +// +// Spec: docs/superpowers/specs/2026-05-31-v9-defensive-eval-boundary-calibration.md +// Pearl: pearl_adaptive_carryover_discipline +// +// At every regime boundary (train→eval, fold transition), the trainer +// calls `reset_session_state` which sets `RL_EVAL_WARMUP_REMAINING_INDEX` +// to the configured warmup duration. This kernel runs once per step +// thereafter; while the counter is positive, it overrides four +// risk-sizing controller floors with conservative defensive values. +// +// Schedule (let R = warmup_remaining, W = warmup_steps, D = decay_steps): +// +// R > D full defensive overrides written +// 0 < R ≤ D linear interpolation: defensive (1) → normal (0) +// R = 0 normal values written once, counter goes to -1 +// R < 0 no-op (warmup completed, controllers run unimpeded) +// +// The override slots affected are read-only by their consuming controllers +// (Kelly safety_frac, IQN τ_min, entropy_coef_min, PPO clip ε_min), so +// last-writer-wins applies cleanly when this kernel runs AFTER the +// controller batch. +// +// Per `feedback_no_atomicadd`: single-thread single-block, no atomics. +// Per `feedback_cpu_is_read_only`: pure device kernel. +// Per `feedback_isv_for_adaptive_bounds`: defensive override values, normal +// target values, warmup/decay durations are all ISV-driven. + +#include + +#define RL_EVAL_WARMUP_REMAINING_INDEX 685 +#define RL_EVAL_WARMUP_STEPS_CONFIG_INDEX 686 +#define RL_EVAL_WARMUP_DECAY_STEPS_CONFIG_INDEX 687 +#define RL_EVAL_KELLY_SAFETY_DEFENSIVE_INDEX 688 +#define RL_EVAL_IQN_TAU_MIN_DEFENSIVE_INDEX 689 +#define RL_EVAL_ENTROPY_COEF_MIN_DEFENSIVE_INDEX 690 +#define RL_EVAL_PPO_CLIP_EPS_MIN_DEFENSIVE_INDEX 691 +#define RL_EVAL_KELLY_SAFETY_NORMAL_INDEX 692 +#define RL_EVAL_IQN_TAU_MIN_NORMAL_INDEX 693 +#define RL_EVAL_ENTROPY_COEF_MIN_NORMAL_INDEX 694 +#define RL_EVAL_PPO_CLIP_EPS_MIN_NORMAL_INDEX 695 + +// Consumer slots (controller floors that this kernel overrides during warmup). +#define RL_KELLY_SAFETY_FRAC_INDEX 680 +#define RL_IQN_ACTION_TAU_MIN_INDEX 672 +#define RL_ENTROPY_COEF_MIN_INDEX 645 +#define RL_PPO_CLIP_EPS_MIN_INDEX 640 + +extern "C" __global__ void rl_eval_warmup_decay( + float* __restrict__ isv +) { + if (threadIdx.x != 0 || threadIdx.y != 0 || threadIdx.z != 0) return; + if (blockIdx.x != 0 || blockIdx.y != 0 || blockIdx.z != 0) return; + + const float remaining = isv[RL_EVAL_WARMUP_REMAINING_INDEX]; + + // Post-warmup steady state: do nothing. The normal floors stay at + // their bootstrap values (or whatever the last-warmup-step wrote). + if (remaining < 0.0f) { + return; + } + + // Read configured durations + override / normal values. + const float warmup_steps = isv[RL_EVAL_WARMUP_STEPS_CONFIG_INDEX]; + const float decay_steps = isv[RL_EVAL_WARMUP_DECAY_STEPS_CONFIG_INDEX]; + + const float kelly_def = isv[RL_EVAL_KELLY_SAFETY_DEFENSIVE_INDEX]; + const float tau_def = isv[RL_EVAL_IQN_TAU_MIN_DEFENSIVE_INDEX]; + const float entropy_def = isv[RL_EVAL_ENTROPY_COEF_MIN_DEFENSIVE_INDEX]; + const float ppo_def = isv[RL_EVAL_PPO_CLIP_EPS_MIN_DEFENSIVE_INDEX]; + + const float kelly_norm = isv[RL_EVAL_KELLY_SAFETY_NORMAL_INDEX]; + const float tau_norm = isv[RL_EVAL_IQN_TAU_MIN_NORMAL_INDEX]; + const float entropy_norm = isv[RL_EVAL_ENTROPY_COEF_MIN_NORMAL_INDEX]; + const float ppo_norm = isv[RL_EVAL_PPO_CLIP_EPS_MIN_NORMAL_INDEX]; + + // Phase selection: + // pure warmup remaining > decay_steps → blend = 1.0 (full defensive) + // decay phase 0 < remaining ≤ decay_steps → blend = remaining/decay_steps + // final boundary remaining == 0 → blend = 0.0 (full normal) + float blend; + if (remaining > decay_steps) { + blend = 1.0f; + } else if (decay_steps > 0.0f) { + blend = remaining / decay_steps; // 1.0 → 0.0 as remaining → 0 + } else { + // Defensive: decay_steps configured to 0 means no decay phase. + blend = (remaining > 0.0f) ? 1.0f : 0.0f; + } + + // Linear interpolation: blend × defensive + (1−blend) × normal. + isv[RL_KELLY_SAFETY_FRAC_INDEX] = blend * kelly_def + (1.0f - blend) * kelly_norm; + isv[RL_IQN_ACTION_TAU_MIN_INDEX] = blend * tau_def + (1.0f - blend) * tau_norm; + isv[RL_ENTROPY_COEF_MIN_INDEX] = blend * entropy_def + (1.0f - blend) * entropy_norm; + isv[RL_PPO_CLIP_EPS_MIN_INDEX] = blend * ppo_def + (1.0f - blend) * ppo_norm; + + // Decrement counter for next step. When it crosses 0, the next step + // sees remaining = -1 and returns early (controllers run unimpeded). + isv[RL_EVAL_WARMUP_REMAINING_INDEX] = remaining - 1.0f; + // Spec note: warmup_steps is used only by the trainer to set the + // initial counter value; this kernel reads warmup_steps only to make + // the `remaining > decay_steps` comparison meaningful for non-default + // configurations (e.g., shorter warmup with same decay). + (void) warmup_steps; +} diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index b0c0c950a..536eef567 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -1460,6 +1460,48 @@ pub const RL_TRAIL_LOOSEN_FACTOR_INDEX: usize = 683; /// the [[feedback_no_partial_refactor]] per-batch fix. pub const RL_SESSION_PNL_WORST_INDEX: usize = 684; -/// Last RL-allocated slot index (exclusive). Pre-risk-stack: 662. -/// Post-risk-stack (Fix B IQN-τ split): 685. -pub const RL_SLOTS_END: usize = 685; +// ============================================================ +// v9 — Defensive eval-boundary calibration (spec 2026-05-31-v9) +// +// Implements [[pearl_adaptive_carryover_discipline]]: at every regime +// boundary (fold/eval transition), defensive overrides force conservative +// behavior for the first N steps while EMAs re-bootstrap from new-regime +// data. After N steps, override values linearly decay back to the normal +// (bootstrap-configured) controller floors over `decay_steps` steps. +// ============================================================ + +/// Counter — set to WARMUP_STEPS at every regime boundary, decremented each +/// step. >0 = warmup/decay active; ≤0 = inactive (normal controller floors). +pub const RL_EVAL_WARMUP_REMAINING_INDEX: usize = 685; +/// Configured warmup duration (steps with full defensive overrides). +/// Bootstrap 500 (~10% of a 5000-step eval phase). +pub const RL_EVAL_WARMUP_STEPS_CONFIG_INDEX: usize = 686; +/// Configured decay duration (steps over which override → normal interpolation). +/// Bootstrap 200. Total transition: warmup_steps + decay_steps = 700 steps. +pub const RL_EVAL_WARMUP_DECAY_STEPS_CONFIG_INDEX: usize = 687; + +/// Defensive override: Kelly safety_frac. Bootstrap 0.25 (quarter-Kelly). +/// Active during warmup; decays to NORMAL (slot 692) over decay phase. +pub const RL_EVAL_KELLY_SAFETY_DEFENSIVE_INDEX: usize = 688; +/// Defensive override: IQN τ_min. Bootstrap 0.30 (pessimistic action floor). +/// Active during warmup; decays to NORMAL (slot 693) over decay phase. +pub const RL_EVAL_IQN_TAU_MIN_DEFENSIVE_INDEX: usize = 689; +/// Defensive override: entropy_coef_min. Bootstrap 0.05 (more exploration). +/// Active during warmup; decays to NORMAL (slot 694) over decay phase. +pub const RL_EVAL_ENTROPY_COEF_MIN_DEFENSIVE_INDEX: usize = 690; +/// Defensive override: PPO clip ε_min. Bootstrap 0.10 (larger policy steps). +/// Active during warmup; decays to NORMAL (slot 695) over decay phase. +pub const RL_EVAL_PPO_CLIP_EPS_MIN_DEFENSIVE_INDEX: usize = 691; + +/// Normal-value saved slots — read by warmup_decay kernel as the linear- +/// interpolation target during the decay phase, and as the final restored +/// value after warmup completes. Bootstrap to match the existing controller +/// floor constants. These are NOT modified after bootstrap (read-only refs). +pub const RL_EVAL_KELLY_SAFETY_NORMAL_INDEX: usize = 692; // 0.50 +pub const RL_EVAL_IQN_TAU_MIN_NORMAL_INDEX: usize = 693; // 0.10 +pub const RL_EVAL_ENTROPY_COEF_MIN_NORMAL_INDEX: usize = 694; // 0.01 +pub const RL_EVAL_PPO_CLIP_EPS_MIN_NORMAL_INDEX: usize = 695; // 0.05 + +/// Last RL-allocated slot index (exclusive). +/// Pre-risk-stack: 662. Post-Fix-B: 685. Post-v9 (warmup boundary): 696. +pub const RL_SLOTS_END: usize = 696; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index 38a079d06..fc8130c37 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -214,6 +214,8 @@ const RL_INVENTORY_BETA_CONTROLLER_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_inventory_beta_controller.cubin")); const RL_KELLY_FRACTION_CONTROLLER_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_kelly_fraction_controller.cubin")); +const RL_EVAL_WARMUP_DECAY_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_eval_warmup_decay.cubin")); const RL_WIN_RATE_EMA_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_win_rate_ema_update.cubin")); const RL_AVG_WIN_LOSS_EMA_UPDATE_CUBIN: &[u8] = @@ -678,6 +680,9 @@ pub struct IntegratedTrainer { rl_inventory_beta_controller_fn: CudaFunction, _rl_kelly_fraction_controller_module: Arc, rl_kelly_fraction_controller_fn: CudaFunction, + // v9 (2026-05-31): defensive eval-boundary calibration kernel. + _rl_eval_warmup_decay_module: Arc, + rl_eval_warmup_decay_fn: CudaFunction, // EMA producers feeding the new controllers. _rl_win_rate_ema_update_module: Arc, rl_win_rate_ema_update_fn: CudaFunction, @@ -1595,6 +1600,13 @@ impl IntegratedTrainer { let rl_kelly_fraction_controller_fn = rl_kelly_fraction_controller_module .load_function("rl_kelly_fraction_controller") .context("load rl_kelly_fraction_controller")?; + // v9 — defensive eval-boundary calibration kernel. + let rl_eval_warmup_decay_module = ctx + .load_cubin(RL_EVAL_WARMUP_DECAY_CUBIN.to_vec()) + .context("load rl_eval_warmup_decay cubin")?; + let rl_eval_warmup_decay_fn = rl_eval_warmup_decay_module + .load_function("rl_eval_warmup_decay") + .context("load rl_eval_warmup_decay")?; let rl_win_rate_ema_update_module = ctx .load_cubin(RL_WIN_RATE_EMA_UPDATE_CUBIN.to_vec()) .context("load rl_win_rate_ema_update cubin")?; @@ -2671,6 +2683,8 @@ impl IntegratedTrainer { rl_inventory_beta_controller_fn, _rl_kelly_fraction_controller_module: rl_kelly_fraction_controller_module, rl_kelly_fraction_controller_fn, + _rl_eval_warmup_decay_module: rl_eval_warmup_decay_module, + rl_eval_warmup_decay_fn, _rl_win_rate_ema_update_module: rl_win_rate_ema_update_module, rl_win_rate_ema_update_fn, _rl_avg_win_loss_ema_update_module: rl_avg_win_loss_ema_update_module, @@ -2970,12 +2984,43 @@ impl IntegratedTrainer { /// inheriting all the dead/cooldown accounts from training. pub fn reset_session_state(&self) -> Result<()> { // Summary slots — direct ISV writes. + // + // v9 (2026-05-31, [[pearl_adaptive_carryover_discipline]]): reset + // EVERY adaptive EMA driving downstream behavior, including the + // ones Fix D previously preserved. The rationale for the prior + // selective preservation ("training-acquired statistics should + // carry into eval") proved wrong at the v8 fold-1 boundary: + // Kelly entered eval with train wr=0.34 vs eval wr=0.25, took + // overconfident positions, and lost $500k before EMAs could + // re-converge. Pure adaptive principle now: every signal that + // drives behavior must be re-bootstrapped from new-regime data. + // Defensive warmup duration. Matches RL_EVAL_WARMUP_STEPS_CONFIG_INDEX + // bootstrap (see ISV bootstrap block in `with_controllers_bootstrapped`). + // Hardcoded here because reset_session_state runs before any ISV read + // would observe a host-visible value; the kernel reads its own config. + let warmup_steps = 500.0_f32; for (slot, value) in [ + // Layer 1 CMDP summary slots (kept from Fix D). (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), + // Layer 4 (Kelly) EMAs — NEW in v9, was preserved by Fix D. + (crate::rl::isv_slots::RL_WIN_RATE_EMA_INDEX, 0.5_f32), // neutral sentinel + (crate::rl::isv_slots::RL_AVG_WIN_USD_EMA_INDEX, 0.0_f32), // sentinel + (crate::rl::isv_slots::RL_AVG_LOSS_USD_EMA_INDEX, 0.0_f32), // sentinel + (crate::rl::isv_slots::RL_CUMULATIVE_DONES_INDEX, 0.0_f32), // re-enter warmup + (crate::rl::isv_slots::RL_KELLY_FRACTION_INDEX, 1.0_f32), // bootstrap + // Layer 3 (Inventory) EMAs — NEW in v9. + (crate::rl::isv_slots::RL_INVENTORY_PENALTY_BETA_INDEX, 0.0_f32), // sentinel + (crate::rl::isv_slots::RL_INVENTORY_VARIANCE_EMA_INDEX, 0.0_f32), // sentinel + // Reward-clamp EMAs — NEW in v9. + (crate::rl::isv_slots::RL_POS_SCALED_REWARD_MAX_EMA_INDEX, 0.0_f32), + (crate::rl::isv_slots::RL_NEG_SCALED_REWARD_MAX_EMA_INDEX, 0.0_f32), + (crate::rl::isv_slots::RL_REWARD_CLAMP_CLIP_RATE_EMA_INDEX, 0.0_f32), + // v9 — arm the defensive warmup window. + (crate::rl::isv_slots::RL_EVAL_WARMUP_REMAINING_INDEX, warmup_steps), ] { let mut args = RawArgs::new(); args.push_ptr(self.isv_dev_ptr); @@ -3091,7 +3136,7 @@ impl IntegratedTrainer { // (slot, value) pair — pure device write, no HtoD per // `feedback_no_htod_htoh_only_mapped_pinned`. { - let isv_constants: [(usize, f32); 168] = [ + let isv_constants: [(usize, f32); 179] = [ // Static seeds for the adaptive reward-clamp controller — // these are the initial values that // `rl_reward_clamp_controller` will replace once it @@ -3369,6 +3414,23 @@ impl IntegratedTrainer { // Layer D (TrailTighten/Loosen factors). (crate::rl::isv_slots::RL_TRAIL_TIGHTEN_FACTOR_INDEX, 0.9), (crate::rl::isv_slots::RL_TRAIL_LOOSEN_FACTOR_INDEX, 1.1), + // v9 (defensive eval-boundary calibration, [[pearl_adaptive_carryover_discipline]]): + // 500-step warmup with 200-step linear decay phase. During warmup, + // risk-sizing floors are temporarily overridden with defensive + // values; reset_session_state arms the counter at each boundary. + (crate::rl::isv_slots::RL_EVAL_WARMUP_REMAINING_INDEX, -1.0), // sentinel: no active warmup at boot + (crate::rl::isv_slots::RL_EVAL_WARMUP_STEPS_CONFIG_INDEX, 500.0), // total warmup duration + (crate::rl::isv_slots::RL_EVAL_WARMUP_DECAY_STEPS_CONFIG_INDEX, 200.0), // linear decay window + // Defensive overrides (active for ~first 300 steps post-boundary). + (crate::rl::isv_slots::RL_EVAL_KELLY_SAFETY_DEFENSIVE_INDEX, 0.25), // quarter-Kelly + (crate::rl::isv_slots::RL_EVAL_IQN_TAU_MIN_DEFENSIVE_INDEX, 0.30), // raise pessimism floor + (crate::rl::isv_slots::RL_EVAL_ENTROPY_COEF_MIN_DEFENSIVE_INDEX, 0.05), // more exploration + (crate::rl::isv_slots::RL_EVAL_PPO_CLIP_EPS_MIN_DEFENSIVE_INDEX, 0.10), // wider trust region + // Normal target values (restored post-decay; mirror Phase-3/4 bootstrap floors). + (crate::rl::isv_slots::RL_EVAL_KELLY_SAFETY_NORMAL_INDEX, 0.50), // half-Kelly (Thorp) + (crate::rl::isv_slots::RL_EVAL_IQN_TAU_MIN_NORMAL_INDEX, 0.10), // existing IQN floor + (crate::rl::isv_slots::RL_EVAL_ENTROPY_COEF_MIN_NORMAL_INDEX, 0.01), // existing entropy floor + (crate::rl::isv_slots::RL_EVAL_PPO_CLIP_EPS_MIN_NORMAL_INDEX, 0.05), // existing PPO ε floor ]; for (slot, value) in isv_constants.iter() { let slot_i32 = *slot as i32; @@ -3710,6 +3772,27 @@ impl IntegratedTrainer { Ok(()) } + /// v9 — defensive eval-boundary calibration. Launch once per step + /// AFTER all risk-stack controllers run (so this kernel's overrides + /// of Kelly_safety_frac / IQN τ_min / entropy_coef_min / PPO ε_min + /// take effect for the next step's consumers). Returns early if + /// `RL_EVAL_WARMUP_REMAINING_INDEX` ≤ 0, so the runtime cost when + /// warmup is inactive is just a single mapped-pinned ISV read. + pub fn launch_rl_eval_warmup_decay(&self) -> Result<()> { + let mut args = RawArgs::new(); + args.push_ptr(self.isv_dev_ptr); + let mut ptrs = args.build_arg_ptrs(); + unsafe { + raw_launch( + self.rl_eval_warmup_decay_fn.cu_function(), + (1, 1, 1), (1, 1, 1), 0, + self.raw_stream, + &mut ptrs[..args.len()], + ).map_err(|e| anyhow::anyhow!("rl_eval_warmup_decay: {:?}", e))?; + } + Ok(()) + } + /// Adaptive risk management — Layer 4 input EMA producer (win_rate). /// Reads `rewards_d`, `dones_d` [b_size] f32; outcome derived inline. /// Writes ISV[RL_WIN_RATE_EMA_INDEX = 677]. @@ -6994,6 +7077,15 @@ impl IntegratedTrainer { self.launch_rl_fused_controllers(b_size) .context("launch_rl_fused_controllers")?; + // v9 defensive eval-boundary calibration ([[pearl_adaptive_carryover_discipline]]). + // MUST run AFTER fused controllers: last-writer-wins on the four risk-sizing + // floors (Kelly safety_frac, IQN τ_min, entropy_coef_min, PPO clip ε_min). + // Steady-state (counter < 0): kernel returns early — zero behavioral effect. + // After `reset_session_state` arms the counter, kernel overrides floors with + // defensive defaults for warmup_steps, with a linear decay over decay_steps. + self.launch_rl_eval_warmup_decay() + .context("launch_rl_eval_warmup_decay")?; + // PopArt: normalize rewards in-place (replaces apply_reward_scale). // Welford-EMA on batch mean+variance → whitened rewards → ISV // stats for V-correct. Grid=(1), Block=(min(B, 256)). Shared mem