diff --git a/crates/ml-alpha/cuda/rl_rollout_steps_controller.cu b/crates/ml-alpha/cuda/rl_rollout_steps_controller.cu index 85d2b375e..be623aaa5 100644 --- a/crates/ml-alpha/cuda/rl_rollout_steps_controller.cu +++ b/crates/ml-alpha/cuda/rl_rollout_steps_controller.cu @@ -45,22 +45,29 @@ // updates lag the data so far behind that the importance ratios blow // past the clip band. -#define RL_N_ROLLOUT_STEPS_INDEX 404 -#define ROLLOUT_MIN 256.0f -#define ROLLOUT_MAX 8192.0f -#define ROLLOUT_BOOTSTRAP 2048.0f -#define ADV_VAR_RATIO_TARGET 0.1f -// Below this magnitude the streaming kernel's signal is numerical -// noise / cold-start, not a real divergence-from-target — hold. -#define ADV_VAR_RATIO_NOISE_FLOOR (ADV_VAR_RATIO_TARGET * 0.01f) // = 1e-3 +#define RL_N_ROLLOUT_STEPS_INDEX 404 +#define ROLLOUT_MIN 256.0f +#define ROLLOUT_MAX 8192.0f +#define ROLLOUT_BOOTSTRAP 2048.0f +// ISV slot holding the adaptive regression target. Per +// `feedback_isv_for_adaptive_bounds`: target lives in ISV (seeded at +// trainer init by rl_streaming_clamp_init, visible in diag, +// modifiable at runtime) rather than as a kernel-side `#define`. +// Default 5.0 — matches the b_size=1 streaming regime where var/|mean| +// naturally lives in [1, 10]. Prior `#define = 0.1` was wrong for the +// streaming regime, causing 99.99% WIDEN events in alpha-rl-cvf86. +#define RL_ADV_VAR_RATIO_TARGET_INDEX 449 +// Noise-floor multiplier: input below TARGET × NOISE_FLOOR_FRAC is +// numerical noise — hold. Derived multiplicatively from the +// ISV-resident target so adjusting the target proportionally +// adjusts the floor. +#define ADV_VAR_RATIO_NOISE_FLOOR_FRAC 0.01f // In-band tolerance: input within ±33 % of target → hold. -#define ADV_VAR_RATIO_TOLERANCE 1.5f +#define ADV_VAR_RATIO_TOLERANCE 1.5f // Per-step bounded multiplicative adjustment. 1.5× == 50% growth / -// 33% shrinkage per fire. Prevents one observation from sending -// rollout from BOOTSTRAP (2048) to MAX (8192) in 2 steps as the -// prior 2.0× clamp design did. -#define ADV_VAR_RATIO_ADJUST_RATE 1.5f -#define WIENER_ALPHA_FLOOR 0.4f +// 33% shrinkage per fire. +#define ADV_VAR_RATIO_ADJUST_RATE 1.5f +#define WIENER_ALPHA_FLOOR 0.4f // ───────────────────────────────────────────────────────────────────── // rl_rollout_steps_controller: @@ -104,11 +111,17 @@ extern "C" __global__ void rl_rollout_steps_controller( const float advantage_var_over_abs_mean = isv[input_slot]; if (advantage_var_over_abs_mean == 0.0f) return; - // Noise-floor gate: input below ADV_VAR_RATIO_NOISE_FLOOR is + // ISV-driven regression target — read from + // RL_ADV_VAR_RATIO_TARGET_INDEX (seeded at trainer init). + const float adv_var_target = isv[RL_ADV_VAR_RATIO_TARGET_INDEX]; + + // Noise-floor gate: input below target × NOISE_FLOOR_FRAC is // dominated by numerical noise — hold rollout unchanged. // Mirrors the ppo_clip / target_tau controllers' design after the // alpha-rl-mjzfk multiplicative-blow-up incident. - if (advantage_var_over_abs_mean < ADV_VAR_RATIO_NOISE_FLOOR) return; + const float adv_var_noise_floor = + adv_var_target * ADV_VAR_RATIO_NOISE_FLOOR_FRAC; + if (advantage_var_over_abs_mean < adv_var_noise_floor) return; // Bounded multiplicative adjustment (Schulman-style adaptive). // At most ADV_VAR_RATIO_ADJUST_RATE × shift per step regardless of @@ -117,11 +130,11 @@ extern "C" __global__ void rl_rollout_steps_controller( // MIN/MAX, but a single observation can't slam it there. float scale; if (advantage_var_over_abs_mean - > ADV_VAR_RATIO_TARGET * ADV_VAR_RATIO_TOLERANCE) { + > adv_var_target * ADV_VAR_RATIO_TOLERANCE) { // Noisy → widen. scale = ADV_VAR_RATIO_ADJUST_RATE; } else if (advantage_var_over_abs_mean - < ADV_VAR_RATIO_TARGET / ADV_VAR_RATIO_TOLERANCE) { + < adv_var_target / ADV_VAR_RATIO_TOLERANCE) { // Clean → shrink. scale = 1.0f / ADV_VAR_RATIO_ADJUST_RATE; } else { diff --git a/crates/ml-alpha/cuda/rl_streaming_clamp_init.cu b/crates/ml-alpha/cuda/rl_streaming_clamp_init.cu index 98fcf948f..e7b5801cc 100644 --- a/crates/ml-alpha/cuda/rl_streaming_clamp_init.cu +++ b/crates/ml-alpha/cuda/rl_streaming_clamp_init.cu @@ -1,35 +1,41 @@ // rl_streaming_clamp_init.cu — single-thread device-side seeder for -// the streaming-kernel output clamp ceilings at -// ISV[RL_ADV_VAR_RATIO_CLAMP_INDEX = 447] and -// ISV[RL_TD_KURTOSIS_CLAMP_INDEX = 448]. +// the streaming-kernel output clamp ceilings AND the +// rl_rollout_steps_controller regression target: +// 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 // // Launched once from `with_controllers_bootstrapped` after ISV is -// alloc_zeros'd. Writes design-constant clamp ceilings directly on -// device — no host→device transfer, in line with +// alloc_zeros'd. Writes design-constant values directly on device — +// no host→device transfer, in line with // `feedback_no_htod_htoh_only_mapped_pinned`. // -// The streaming variance and kurtosis kernels then read these slots -// each step and clamp their output before writing to the consumer -// controllers' EMA-input slots. Without the clamp, streaming +// The streaming variance and kurtosis kernels read the clamp slots +// each step and bound their output before writing to the consumer +// controllers' EMA-input slots. Without the clamps streaming // `var/|mean|` reached 3e5 in alpha-rl-gxhr8 fold0 (when streaming -// mean passed through zero — `var / max(|μ|, 1e-6)` blows up), -// pegging the rollout_steps controller at MAX; streaming kurtosis -// reached 50.6, pegging per_α at MAX which over-concentrated PER -// sampling and hurt distributional Q learning. +// mean passed through zero — `var / max(|μ|, 1e-6)` blows up); +// streaming kurtosis reached 50.6. // -// Per `feedback_isv_for_adaptive_bounds`: the ceilings live in ISV -// (visible in diag, modifiable at runtime by re-launching this -// kernel or by a future adaptive controller) rather than as kernel- -// side `#define`s. +// The rollout-controller target slot replaces the prior hardcoded +// `ADV_VAR_RATIO_TARGET = 0.1` `#define` — that value was chosen for +// batched (b_size>1) reductions where var/|mean| naturally sits in +// [0.01, 0.2], but the streaming regime lives in [1, 10] so 0.1 +// triggers WIDEN on 99.99% of steps (alpha-rl-cvf86 fold0). The +// new ISV-driven target defaults to 5.0 (matches streaming median). +// +// Per `feedback_isv_for_adaptive_bounds`: all three values live in +// ISV (visible in diag, modifiable at runtime by re-launching this +// kernel) rather than as kernel-side `#define`s. 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_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 ) { 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_clamp_slot] = adv_var_clamp_val; + isv[td_kurt_clamp_slot] = td_kurt_clamp_val; + isv[adv_var_target_slot] = adv_var_target_val; } diff --git a/crates/ml-alpha/examples/alpha_rl_train.rs b/crates/ml-alpha/examples/alpha_rl_train.rs index a03b9e104..0f189433f 100644 --- a/crates/ml-alpha/examples/alpha_rl_train.rs +++ b/crates/ml-alpha/examples/alpha_rl_train.rs @@ -45,7 +45,8 @@ use ml_alpha::rl::isv_slots::{ 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_ADV_VAR_RATIO_CLAMP_INDEX, RL_ADV_VAR_STREAM_M2_INDEX, RL_ADV_VAR_STREAM_MEAN_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_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,7 +627,9 @@ fn main() -> Result<()> { // controller's behaviour, not adaptive). "controller_branch": { "rollout_steps_input": isv[RL_ADVANTAGE_VAR_RATIO_EMA_INDEX], - "rollout_steps_target": 0.1f32, + // ISV-driven target — seeded by rl_streaming_clamp_init, + // visible here (no longer a hardcoded `#define`). + "rollout_steps_target": isv[RL_ADV_VAR_RATIO_TARGET_INDEX], "ppo_clip_input": isv[RL_KL_PI_EMA_INDEX], "ppo_clip_target": 0.01f32, "target_tau_input": isv[RL_Q_DIVERGENCE_EMA_INDEX], diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index a68967278..04a096ad4 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -311,6 +311,27 @@ pub const RL_ADV_VAR_RATIO_CLAMP_INDEX: usize = 447; /// runaway. pub const RL_TD_KURTOSIS_CLAMP_INDEX: usize = 448; +/// Adaptive target for the advantage-variance ratio that +/// `rl_rollout_steps_controller` regresses toward. Per +/// `feedback_isv_for_adaptive_bounds`: the target lives in ISV +/// (visible in diag, modifiable at runtime by re-seeding) rather than +/// as a kernel-side `#define`. +/// +/// Default 5.0 — matches the b_size=1 streaming-kernel regime where +/// `var/|mean|` naturally lives in [1, 10]. The prior hardcoded +/// `ADV_VAR_RATIO_TARGET = 0.1` was chosen for batched (b_size>1) +/// reductions where var/|mean| sits in [0.01, 0.2]; with the +/// streaming kernel that target is two orders of magnitude too +/// low, so the controller WIDENS 99.99% of steps and pegs +/// n_rollout at MAX (alpha-rl-cvf86 fold0 — commit 708c121f2). +/// +/// Seeded device-side at trainer init by `rl_streaming_clamp_init` +/// (no HtoD per `feedback_no_htod_htoh_only_mapped_pinned`). A +/// future adaptive controller could maintain this slot from the +/// streaming output's own percentile EMA, putting the regression +/// target into a true feedback loop with the signal. +pub const RL_ADV_VAR_RATIO_TARGET_INDEX: usize = 449; + /// 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 = 449; +pub const RL_SLOTS_END: usize = 450; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index c30b948bc..53b758bbc 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -1192,6 +1192,14 @@ impl IntegratedTrainer { let td_kurt_slot: i32 = crate::rl::isv_slots::RL_TD_KURTOSIS_CLAMP_INDEX as i32; let td_kurt_val: f32 = 30.0; + // ISV-driven regression target for rl_rollout_steps_controller. + // 5.0 matches the b_size=1 streaming regime where + // var/|mean| naturally lives in [1, 10]. The prior + // hardcoded #define 0.1 was a b_size>1 design choice and + // caused 99.99% WIDEN events in alpha-rl-cvf86. + 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; let cfg = LaunchConfig { grid_dim: (1, 1, 1), block_dim: (1, 1, 1), @@ -1205,7 +1213,9 @@ impl IntegratedTrainer { .arg(&adv_var_slot) .arg(&adv_var_val) .arg(&td_kurt_slot) - .arg(&td_kurt_val); + .arg(&td_kurt_val) + .arg(&adv_var_target_slot) + .arg(&adv_var_target_val); unsafe { launch .launch(cfg) diff --git a/crates/ml-alpha/tests/isv_bootstrap.rs b/crates/ml-alpha/tests/isv_bootstrap.rs index 0b3c9006f..4d7e54876 100644 --- a/crates/ml-alpha/tests/isv_bootstrap.rs +++ b/crates/ml-alpha/tests/isv_bootstrap.rs @@ -23,10 +23,10 @@ //! `cargo test -p ml-alpha --test isv_bootstrap -- --ignored --nocapture` use ml_alpha::rl::isv_slots::{ - RL_ADV_VAR_RATIO_CLAMP_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_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, }; use ml_alpha::trainer::integrated::{IntegratedTrainer, IntegratedTrainerConfig}; use ml_alpha::trainer::perception::PerceptionTrainerConfig; @@ -144,10 +144,12 @@ fn g1_isv_bootstrap_writes_canonical_values() { // during `with_controllers_bootstrapped`, leaving the slot at // PPO_RATIO_CLAMP_BOOTSTRAP = 10.0 — checked separately below. for slot in RL_MEAN_TRADE_DURATION_EMA_INDEX..RL_SLOTS_END { - // R9 controller-OUTPUT slots that bootstrap to non-zero values. + // R9 controller-OUTPUT / ISV-resident-design-constant slots + // that bootstrap to non-zero values. if slot == RL_PPO_RATIO_CLAMP_MAX_INDEX || slot == RL_ADV_VAR_RATIO_CLAMP_INDEX || slot == RL_TD_KURTOSIS_CLAMP_INDEX + || slot == RL_ADV_VAR_RATIO_TARGET_INDEX { continue; } @@ -175,6 +177,13 @@ fn g1_isv_bootstrap_writes_canonical_values() { "ISV[td_kurtosis_clamp={RL_TD_KURTOSIS_CLAMP_INDEX}] expected 30.0, got {}", isv[RL_TD_KURTOSIS_CLAMP_INDEX] ); + // R9 — ISV-driven advantage-variance-ratio target (replaces the + // prior hardcoded #define = 0.1 in rl_rollout_steps_controller). + assert!( + (isv[RL_ADV_VAR_RATIO_TARGET_INDEX] - 5.0).abs() < EPS, + "ISV[adv_var_ratio_target={RL_ADV_VAR_RATIO_TARGET_INDEX}] expected 5.0, got {}", + isv[RL_ADV_VAR_RATIO_TARGET_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 922591376..eb0762ccc 100644 --- a/crates/ml-alpha/tests/r5_controllers_and_soft_update.rs +++ b/crates/ml-alpha/tests/r5_controllers_and_soft_update.rs @@ -26,12 +26,13 @@ use cudarc::driver::CudaStream; use ml_alpha::rl::isv_slots::{ - RL_ADVANTAGE_VAR_RATIO_EMA_INDEX, RL_ADV_VAR_RATIO_CLAMP_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_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, }; use ml_alpha::trainer::integrated::{IntegratedTrainer, IntegratedTrainerConfig}; use ml_alpha::trainer::perception::PerceptionTrainerConfig; @@ -119,6 +120,7 @@ fn g3_per_step_controllers_move_isv_outputs_when_fed_real_emas() { if slot == RL_PPO_RATIO_CLAMP_MAX_INDEX || slot == RL_ADV_VAR_RATIO_CLAMP_INDEX || slot == RL_TD_KURTOSIS_CLAMP_INDEX + || slot == RL_ADV_VAR_RATIO_TARGET_INDEX { continue; } @@ -127,6 +129,7 @@ fn g3_per_step_controllers_move_isv_outputs_when_fed_real_emas() { assert_eq!(isv_before[RL_PPO_RATIO_CLAMP_MAX_INDEX], 10.0); 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); // Populate each EMA-input slot with a non-zero value via the // R3 ema_update_per_step bootstrap path (sentinel-zero → first @@ -142,7 +145,9 @@ fn g3_per_step_controllers_move_isv_outputs_when_fed_real_emas() { (RL_Q_DIVERGENCE_EMA_INDEX, 0.5), // → rl_target_tau (RL_KL_PI_EMA_INDEX, 0.1), // → rl_ppo_clip (RL_ENTROPY_OBSERVED_EMA_INDEX, 0.5), // → rl_entropy_coef - (RL_ADVANTAGE_VAR_RATIO_EMA_INDEX, 5.0), // → rl_rollout_steps + // Above ADV_VAR_RATIO_TARGET (5.0) × TOLERANCE (1.5) = 7.5 so + // the WIDEN branch fires and rollout_steps moves off bootstrap. + (RL_ADVANTAGE_VAR_RATIO_EMA_INDEX, 20.0), // → rl_rollout_steps (RL_TD_KURTOSIS_EMA_INDEX, 10.0), // → rl_per_alpha (RL_MEAN_ABS_PNL_EMA_INDEX, 50.0), // → rl_reward_scale ];