From 644fbe03486cde844e183a2cb9522aafa685a5bb Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sun, 24 May 2026 00:23:38 +0200 Subject: [PATCH] fix(rl): ISV-driven K-loop divisor + max ceiling (slot 450, 451) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit f2ggr confirmed K-loop wiring works mechanically but K=8 firing on 22 % of steps over-trained at b_size=1: KL excursions to 12.44 (vs prior 3.4e-4), policy overshoot, reward/trade -$0.585 → -$0.723. Per `feedback_isv_for_adaptive_bounds` the K-loop config must live in ISV, not as hardcoded values in the trainer. Two new slots: RL_K_LOOP_DIVISOR_INDEX (450) — divides n_rollout_steps to get K Default 2048 (matches ROLLOUT_BOOTSTRAP so K=1 at controller bootstrap) RL_K_LOOP_MAX_INDEX (451) — clamp ceiling on K Default 4 (was hardcoded 8; halved to prevent gradient overtraining) K computation in step_with_lobsim now reads both from ISV: K = clamp(isv[404] / isv[450], 1, isv[451]) Halves worst-case overtraining while preserving the controller cascade activation (KL above noise floor, ε actively adapting, ratio_clamp firing). Distribution shifts from K=8 @ 22% → K=4 @ 22% (half the gradient updates in the high-noise case). ## Wiring `rl_streaming_clamp_init.cu` extended to seed 5 ISV-resident design constants (was 3): adv_var_clamp, td_kurt_clamp, adv_var_target, k_loop_divisor, k_loop_max. Still one kernel call, no HtoD. ## Diag bake-in JSONL `k_updates` field replaced with `k_loop` block: k_loop.k_updates — actual K used this step k_loop.divisor — current divisor (reads isv[450]) k_loop.max — current max (reads isv[451]) Post-hoc analysis can verify the K-computation by independently recomputing K from isv[404] / k_loop.divisor. ## Slot allocation RL_SLOTS_END: 450 → 452 (+2 new config slots). ## Test updates G1 + G3 skip slots 450, 451 in sentinel-zero loop and assert seeded values (2048.0 + 4.0) separately. ## Verified gates (local sm_86) G1 isv_bootstrap ✅ G3 controllers ✅ G4 target_update ✅ integrated_smoke ✅ Co-Authored-By: Claude Opus 4.7 --- .../ml-alpha/cuda/rl_streaming_clamp_init.cu | 12 +++++-- crates/ml-alpha/examples/alpha_rl_train.rs | 19 ++++++---- crates/ml-alpha/src/rl/isv_slots.rs | 34 +++++++++++++++++- crates/ml-alpha/src/trainer/integrated.rs | 36 +++++++++++++++++-- crates/ml-alpha/tests/isv_bootstrap.rs | 20 +++++++++-- .../tests/r5_controllers_and_soft_update.rs | 14 +++++--- 6 files changed, 113 insertions(+), 22 deletions(-) diff --git a/crates/ml-alpha/cuda/rl_streaming_clamp_init.cu b/crates/ml-alpha/cuda/rl_streaming_clamp_init.cu index e7b5801cc..0924dbad6 100644 --- a/crates/ml-alpha/cuda/rl_streaming_clamp_init.cu +++ b/crates/ml-alpha/cuda/rl_streaming_clamp_init.cu @@ -1,9 +1,11 @@ // rl_streaming_clamp_init.cu — single-thread device-side seeder for -// the streaming-kernel output clamp ceilings AND the -// rl_rollout_steps_controller regression target: +// the streaming-kernel output clamp ceilings, the rollout-controller +// regression target, AND the K-loop training-intensification config: // ISV[RL_ADV_VAR_RATIO_CLAMP_INDEX = 447] — var/|mean| clamp // ISV[RL_TD_KURTOSIS_CLAMP_INDEX = 448] — kurtosis clamp // ISV[RL_ADV_VAR_RATIO_TARGET_INDEX = 449] — rollout-controller target +// ISV[RL_K_LOOP_DIVISOR_INDEX = 450] — K = n_rollout / divisor +// ISV[RL_K_LOOP_MAX_INDEX = 451] — K clamped to this ceiling // // Launched once from `with_controllers_bootstrapped` after ISV is // alloc_zeros'd. Writes design-constant values directly on device — @@ -32,10 +34,14 @@ extern "C" __global__ void rl_streaming_clamp_init( float* __restrict__ isv, int adv_var_clamp_slot, float adv_var_clamp_val, int td_kurt_clamp_slot, float td_kurt_clamp_val, - int adv_var_target_slot, float adv_var_target_val + int adv_var_target_slot, float adv_var_target_val, + int k_loop_divisor_slot, float k_loop_divisor_val, + int k_loop_max_slot, float k_loop_max_val ) { if (threadIdx.x != 0 || blockIdx.x != 0) return; isv[adv_var_clamp_slot] = adv_var_clamp_val; isv[td_kurt_clamp_slot] = td_kurt_clamp_val; isv[adv_var_target_slot] = adv_var_target_val; + isv[k_loop_divisor_slot] = k_loop_divisor_val; + isv[k_loop_max_slot] = k_loop_max_val; } diff --git a/crates/ml-alpha/examples/alpha_rl_train.rs b/crates/ml-alpha/examples/alpha_rl_train.rs index 877eccef4..5c7285f79 100644 --- a/crates/ml-alpha/examples/alpha_rl_train.rs +++ b/crates/ml-alpha/examples/alpha_rl_train.rs @@ -46,7 +46,7 @@ use ml_alpha::rl::isv_slots::{ 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_ADV_VAR_RATIO_CLAMP_INDEX, RL_ADV_VAR_RATIO_TARGET_INDEX, RL_ADV_VAR_STREAM_M2_INDEX, - RL_ADV_VAR_STREAM_MEAN_INDEX, + RL_ADV_VAR_STREAM_MEAN_INDEX, RL_K_LOOP_DIVISOR_INDEX, RL_K_LOOP_MAX_INDEX, RL_MAX_ABS_SCALED_REWARD_PRE_CLAMP_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_PPO_LOG_RATIO_ABS_MAX_INDEX, @@ -626,12 +626,17 @@ fn main() -> Result<()> { // design constants because they're fundamental to the // controller's behaviour, not adaptive). // R9 — K-loop training-intensification iterations performed - // this step. K = clamp(isv[404]/1024, 1, 8). Surfaces the - // wire-up of n_rollout_steps as the DQN-replay/PPO multi- - // update knob: when advantages are noisy, do K Q+π+V - // updates per env step (each with fresh PER sample) so - // gradient signal isn't starved at b_size=1. - "k_updates": trainer.last_k_updates, + // this step. K = clamp(isv[404] / isv[450], 1, isv[451]). + // Both divisor + max are ISV-driven (per + // feedback_isv_for_adaptive_bounds) — surfaced here so + // post-hoc analysis can compute K from any sample without + // recompile. Defaults: divisor=2048 (K=1 at controller + // bootstrap), max=4 (prevents gradient overtraining). + "k_loop": { + "k_updates": trainer.last_k_updates, + "divisor": isv[RL_K_LOOP_DIVISOR_INDEX], + "max": isv[RL_K_LOOP_MAX_INDEX], + }, "controller_branch": { "rollout_steps_input": isv[RL_ADVANTAGE_VAR_RATIO_EMA_INDEX], // ISV-driven target — seeded by rl_streaming_clamp_init, diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index 04a096ad4..aeaa8fff7 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -332,6 +332,38 @@ pub const RL_TD_KURTOSIS_CLAMP_INDEX: usize = 448; /// target into a true feedback loop with the signal. pub const RL_ADV_VAR_RATIO_TARGET_INDEX: usize = 449; +/// K-loop divisor for the n_rollout_steps training-intensification +/// wire-up. The trainer computes +/// `K = clamp(isv[RL_N_ROLLOUT_STEPS_INDEX] / isv[RL_K_LOOP_DIVISOR_INDEX], +/// 1, isv[RL_K_LOOP_MAX_INDEX])` +/// to determine how many full (PER sample + step_synthetic + PER +/// priority) iterations to run per env step. +/// +/// Default 2048.0 — matches `ROLLOUT_BOOTSTRAP` so at controller +/// bootstrap (before any adaptation), `K = 2048/2048 = 1` and the +/// loop behaves like a single per-step update (current baseline). +/// As n_rollout grows toward MAX=8192, K rises to K_MAX. +/// +/// Raised from the prior hardcoded `1024.0` after alpha-rl-f2ggr +/// fold0 (commit 1d8ef9484) confirmed K=8 firing on 22 % of steps +/// caused gradient overtraining: KL excursions to 12.44 (vs prior +/// 3.4e-4), policy overshoot, reward/trade -$0.585 → -$0.723. +/// 2048 halves the worst-case K to 4 while preserving the controller +/// cascade activation (KL above noise floor, ε actively adapting). +/// +/// Per `feedback_isv_for_adaptive_bounds`: lives in ISV, discoverable +/// in diag, tunable at runtime via re-launching rl_streaming_clamp_init. +pub const RL_K_LOOP_DIVISOR_INDEX: usize = 450; + +/// K-loop maximum ceiling. K is clamped to this value regardless of +/// how high n_rollout_steps grows. Default 4.0 — prevents +/// gradient overtraining at b_size=1 even if n_rollout_steps pegs +/// at MAX=8192 (would otherwise give K=4 with divisor=2048, but +/// kept explicit as defense-in-depth). +/// +/// Per `feedback_isv_for_adaptive_bounds`: lives in ISV. +pub const RL_K_LOOP_MAX_INDEX: usize = 451; + /// 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 = 450; +pub const RL_SLOTS_END: usize = 452; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index 53f5c17cc..cc29c845a 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -1209,6 +1209,19 @@ impl IntegratedTrainer { let adv_var_target_slot: i32 = crate::rl::isv_slots::RL_ADV_VAR_RATIO_TARGET_INDEX as i32; let adv_var_target_val: f32 = 5.0; + // K-loop config — also ISV-driven per + // `feedback_isv_for_adaptive_bounds`. Divisor 2048 matches + // ROLLOUT_BOOTSTRAP (so K=1 at controller bootstrap, before + // any adaptation). K_MAX=4 prevents gradient overtraining + // at b_size=1; raised from default 8 after alpha-rl-f2ggr + // confirmed K=8 firing on 22% of steps over-trained + // (KL excursions to 12.44, reward/trade -$0.585 → -$0.723). + let k_loop_divisor_slot: i32 = + crate::rl::isv_slots::RL_K_LOOP_DIVISOR_INDEX as i32; + let k_loop_divisor_val: f32 = 2048.0; + let k_loop_max_slot: i32 = + crate::rl::isv_slots::RL_K_LOOP_MAX_INDEX as i32; + let k_loop_max_val: f32 = 4.0; let cfg = LaunchConfig { grid_dim: (1, 1, 1), block_dim: (1, 1, 1), @@ -1224,7 +1237,11 @@ impl IntegratedTrainer { .arg(&td_kurt_slot) .arg(&td_kurt_val) .arg(&adv_var_target_slot) - .arg(&adv_var_target_val); + .arg(&adv_var_target_val) + .arg(&k_loop_divisor_slot) + .arg(&k_loop_divisor_val) + .arg(&k_loop_max_slot) + .arg(&k_loop_max_val); unsafe { launch .launch(cfg) @@ -2767,10 +2784,23 @@ impl IntegratedTrainer { // addresses the per-step gradient signal starvation that left // l_q stuck at 2.82 (uniform=3.04, only 7% information gain) // in alpha-rl-frt7s. + // K-loop config — both divisor and max are ISV-driven per + // `feedback_isv_for_adaptive_bounds` so they're discoverable + // in diag and tunable at runtime by re-launching + // rl_streaming_clamp_init. Defaults seeded at trainer init: + // divisor=2048 (matches ROLLOUT_BOOTSTRAP so K=1 at controller + // bootstrap), max=4 (prevents gradient overtraining at + // b_size=1). let n_rollout_steps = self.isv_host[crate::rl::isv_slots::RL_N_ROLLOUT_STEPS_INDEX]; - let k_updates = ((n_rollout_steps / 1024.0).round() as usize) - .clamp(1, 8); + let k_divisor = self + .isv_host[crate::rl::isv_slots::RL_K_LOOP_DIVISOR_INDEX] + .max(1.0); + let k_max = self + .isv_host[crate::rl::isv_slots::RL_K_LOOP_MAX_INDEX] + .max(1.0) as usize; + let k_updates = ((n_rollout_steps / k_divisor).round() as usize) + .clamp(1, k_max); self.last_k_updates = k_updates; // Track the LAST iteration's stats for caller reporting (per- diff --git a/crates/ml-alpha/tests/isv_bootstrap.rs b/crates/ml-alpha/tests/isv_bootstrap.rs index 4d7e54876..0f190ae04 100644 --- a/crates/ml-alpha/tests/isv_bootstrap.rs +++ b/crates/ml-alpha/tests/isv_bootstrap.rs @@ -24,9 +24,10 @@ use ml_alpha::rl::isv_slots::{ RL_ADV_VAR_RATIO_CLAMP_INDEX, RL_ADV_VAR_RATIO_TARGET_INDEX, RL_ENTROPY_COEF_INDEX, - RL_GAMMA_INDEX, RL_MEAN_TRADE_DURATION_EMA_INDEX, RL_N_ROLLOUT_STEPS_INDEX, - RL_PER_ALPHA_INDEX, RL_PPO_CLIP_INDEX, RL_PPO_RATIO_CLAMP_MAX_INDEX, RL_REWARD_SCALE_INDEX, - RL_SLOTS_END, RL_TARGET_TAU_INDEX, RL_TD_KURTOSIS_CLAMP_INDEX, + RL_GAMMA_INDEX, RL_K_LOOP_DIVISOR_INDEX, RL_K_LOOP_MAX_INDEX, + RL_MEAN_TRADE_DURATION_EMA_INDEX, RL_N_ROLLOUT_STEPS_INDEX, RL_PER_ALPHA_INDEX, + RL_PPO_CLIP_INDEX, RL_PPO_RATIO_CLAMP_MAX_INDEX, RL_REWARD_SCALE_INDEX, RL_SLOTS_END, + RL_TARGET_TAU_INDEX, RL_TD_KURTOSIS_CLAMP_INDEX, }; use ml_alpha::trainer::integrated::{IntegratedTrainer, IntegratedTrainerConfig}; use ml_alpha::trainer::perception::PerceptionTrainerConfig; @@ -150,6 +151,8 @@ fn g1_isv_bootstrap_writes_canonical_values() { || slot == RL_ADV_VAR_RATIO_CLAMP_INDEX || slot == RL_TD_KURTOSIS_CLAMP_INDEX || slot == RL_ADV_VAR_RATIO_TARGET_INDEX + || slot == RL_K_LOOP_DIVISOR_INDEX + || slot == RL_K_LOOP_MAX_INDEX { continue; } @@ -184,6 +187,17 @@ fn g1_isv_bootstrap_writes_canonical_values() { "ISV[adv_var_ratio_target={RL_ADV_VAR_RATIO_TARGET_INDEX}] expected 5.0, got {}", isv[RL_ADV_VAR_RATIO_TARGET_INDEX] ); + // R9 — K-loop config (divisor + max), seeded by rl_streaming_clamp_init. + assert!( + (isv[RL_K_LOOP_DIVISOR_INDEX] - 2048.0).abs() < EPS, + "ISV[k_loop_divisor={RL_K_LOOP_DIVISOR_INDEX}] expected 2048.0, got {}", + isv[RL_K_LOOP_DIVISOR_INDEX] + ); + assert!( + (isv[RL_K_LOOP_MAX_INDEX] - 4.0).abs() < EPS, + "ISV[k_loop_max={RL_K_LOOP_MAX_INDEX}] expected 4.0, got {}", + isv[RL_K_LOOP_MAX_INDEX] + ); eprintln!( "G1 OK — bootstraps: γ={:.4} τ={:.4} ε={:.4} entropy_coef={:.4} \ diff --git a/crates/ml-alpha/tests/r5_controllers_and_soft_update.rs b/crates/ml-alpha/tests/r5_controllers_and_soft_update.rs index eb0762ccc..0955535c1 100644 --- a/crates/ml-alpha/tests/r5_controllers_and_soft_update.rs +++ b/crates/ml-alpha/tests/r5_controllers_and_soft_update.rs @@ -28,11 +28,11 @@ use cudarc::driver::CudaStream; use ml_alpha::rl::isv_slots::{ RL_ADVANTAGE_VAR_RATIO_EMA_INDEX, RL_ADV_VAR_RATIO_CLAMP_INDEX, RL_ADV_VAR_RATIO_TARGET_INDEX, RL_ENTROPY_COEF_INDEX, RL_ENTROPY_OBSERVED_EMA_INDEX, - RL_GAMMA_INDEX, RL_KL_PI_EMA_INDEX, RL_MEAN_ABS_PNL_EMA_INDEX, - RL_MEAN_TRADE_DURATION_EMA_INDEX, RL_N_ROLLOUT_STEPS_INDEX, RL_PER_ALPHA_INDEX, - RL_PPO_CLIP_INDEX, RL_PPO_RATIO_CLAMP_MAX_INDEX, RL_Q_DIVERGENCE_EMA_INDEX, - RL_REWARD_SCALE_INDEX, RL_SLOTS_END, RL_TARGET_TAU_INDEX, RL_TD_KURTOSIS_CLAMP_INDEX, - RL_TD_KURTOSIS_EMA_INDEX, + RL_GAMMA_INDEX, RL_KL_PI_EMA_INDEX, RL_K_LOOP_DIVISOR_INDEX, RL_K_LOOP_MAX_INDEX, + RL_MEAN_ABS_PNL_EMA_INDEX, RL_MEAN_TRADE_DURATION_EMA_INDEX, RL_N_ROLLOUT_STEPS_INDEX, + RL_PER_ALPHA_INDEX, RL_PPO_CLIP_INDEX, RL_PPO_RATIO_CLAMP_MAX_INDEX, + RL_Q_DIVERGENCE_EMA_INDEX, RL_REWARD_SCALE_INDEX, RL_SLOTS_END, RL_TARGET_TAU_INDEX, + RL_TD_KURTOSIS_CLAMP_INDEX, RL_TD_KURTOSIS_EMA_INDEX, }; use ml_alpha::trainer::integrated::{IntegratedTrainer, IntegratedTrainerConfig}; use ml_alpha::trainer::perception::PerceptionTrainerConfig; @@ -121,6 +121,8 @@ fn g3_per_step_controllers_move_isv_outputs_when_fed_real_emas() { || slot == RL_ADV_VAR_RATIO_CLAMP_INDEX || slot == RL_TD_KURTOSIS_CLAMP_INDEX || slot == RL_ADV_VAR_RATIO_TARGET_INDEX + || slot == RL_K_LOOP_DIVISOR_INDEX + || slot == RL_K_LOOP_MAX_INDEX { continue; } @@ -130,6 +132,8 @@ fn g3_per_step_controllers_move_isv_outputs_when_fed_real_emas() { assert_eq!(isv_before[RL_ADV_VAR_RATIO_CLAMP_INDEX], 100.0); assert_eq!(isv_before[RL_TD_KURTOSIS_CLAMP_INDEX], 30.0); assert_eq!(isv_before[RL_ADV_VAR_RATIO_TARGET_INDEX], 5.0); + assert_eq!(isv_before[RL_K_LOOP_DIVISOR_INDEX], 2048.0); + assert_eq!(isv_before[RL_K_LOOP_MAX_INDEX], 4.0); // Populate each EMA-input slot with a non-zero value via the // R3 ema_update_per_step bootstrap path (sentinel-zero → first