diff --git a/crates/ml-alpha/cuda/rl_kurtosis_streaming.cu b/crates/ml-alpha/cuda/rl_kurtosis_streaming.cu index 26ebbf7d1..f30fdb65e 100644 --- a/crates/ml-alpha/cuda/rl_kurtosis_streaming.cu +++ b/crates/ml-alpha/cuda/rl_kurtosis_streaming.cu @@ -43,8 +43,9 @@ // // Per `pearl_no_atomicadd`: single-thread kernel. -#define STREAM_ALPHA 0.05f -#define M2_SQ_FLOOR 1e-12f +// ISV-driven streaming-EMA α (was 0.05f #define). +#define RL_STREAM_ALPHA_INDEX 470 +#define M2_SQ_FLOOR 1e-12f // clamp_slot: ISV slot holding the output ceiling // (RL_TD_KURTOSIS_CLAMP_INDEX = 448). Seeded once at @@ -71,14 +72,17 @@ extern "C" __global__ void rl_kurtosis_streaming( for (int b = 0; b < b_size; ++b) batch_mean += x[b]; batch_mean /= (float)b_size; + // ISV-driven streaming-EMA α (shared with var/|mean| kernel). + const float stream_alpha = isv[RL_STREAM_ALPHA_INDEX]; + // Streaming mean with first-observation bootstrap. const float mean_prev = isv[mean_slot]; float mean_new; if (mean_prev == 0.0f) { mean_new = batch_mean; } else { - mean_new = (1.0f - STREAM_ALPHA) * mean_prev - + STREAM_ALPHA * batch_mean; + mean_new = (1.0f - stream_alpha) * mean_prev + + stream_alpha * batch_mean; } isv[mean_slot] = mean_new; @@ -90,20 +94,20 @@ extern "C" __global__ void rl_kurtosis_streaming( const float m2_prev = isv[m2_slot]; float m2_new; if (m2_prev == 0.0f && mean_prev == 0.0f) { - m2_new = dev2; // first observation + m2_new = dev2; } else { - m2_new = (1.0f - STREAM_ALPHA) * m2_prev - + STREAM_ALPHA * dev2; + m2_new = (1.0f - stream_alpha) * m2_prev + + stream_alpha * dev2; } isv[m2_slot] = m2_new; const float m4_prev = isv[m4_slot]; float m4_new; if (m4_prev == 0.0f && mean_prev == 0.0f) { - m4_new = dev4; // first observation + m4_new = dev4; } else { - m4_new = (1.0f - STREAM_ALPHA) * m4_prev - + STREAM_ALPHA * dev4; + m4_new = (1.0f - stream_alpha) * m4_prev + + stream_alpha * dev4; } isv[m4_slot] = m4_new; diff --git a/crates/ml-alpha/cuda/rl_per_alpha_controller.cu b/crates/ml-alpha/cuda/rl_per_alpha_controller.cu index a0279b654..51e0e9d33 100644 --- a/crates/ml-alpha/cuda/rl_per_alpha_controller.cu +++ b/crates/ml-alpha/cuda/rl_per_alpha_controller.cu @@ -47,11 +47,11 @@ // Kurtosis of a standard normal = 3.0 ("excess kurtosis 0" with the // alternative convention). Used as the breakpoint above which we start // raising α. -#define KURT_GAUSSIAN 3.0f -// ISV-driven kurtosis-to-α lift scale per `feedback_isv_for_adaptive_bounds`. -// Default 7.0 — at kurt=10, lift = 0.2 of [PER_ALPHA_MIN, MAX]. Seeded -// by rl_isv_write at init. -#define RL_KURT_LIFT_SCALE_INDEX 459 +// ISV-driven kurtosis interpretation constants per +// `feedback_isv_for_adaptive_bounds`. +#define RL_KURT_GAUSSIAN_INDEX 471 +#define RL_KURT_NOISE_FLOOR_INDEX 472 +#define RL_KURT_LIFT_SCALE_INDEX 459 // Noise-floor gate: if the streaming kurtosis estimator emits a value // below this magnitude, treat it as "no signal" (sentinel-zero proxy) // and hold α at the prior value. Without this, on the first few steps @@ -61,7 +61,7 @@ // MIN despite zero real signal. Matches the defensive noise-floor // pattern on the other controllers (per // pearl_multiplicative_controllers_need_bounded_step_and_noise_floor). -#define KURT_NOISE_FLOOR 1.0f +// (KURT_NOISE_FLOOR — was 1.0f #define — now isv[RL_KURT_NOISE_FLOOR_INDEX]) #define WIENER_ALPHA_FLOOR 0.4f // ───────────────────────────────────────────────────────────────────── @@ -102,7 +102,7 @@ extern "C" __global__ void rl_per_alpha_controller( // differences before tails accumulate). Hold α at prev to avoid // dragging toward PER_ALPHA_MIN on cold-start. const float td_kurtosis_ema = isv[input_slot]; - if (td_kurtosis_ema > 0.0f && td_kurtosis_ema < KURT_NOISE_FLOOR) { + if (td_kurtosis_ema > 0.0f && td_kurtosis_ema < isv[RL_KURT_NOISE_FLOOR_INDEX]) { // Real signal present but below the noise floor — hold prev. // (Strict sentinel zero handled by the prev==0 bootstrap path // below, which derives target from the current EMA so the @@ -120,7 +120,7 @@ extern "C" __global__ void rl_per_alpha_controller( // The 0.4-0.6 baseline keeps the steady-state output near PER's // canonical 0.6 (when input EMA stabilises at kurt=10) while // leaving headroom to lift toward 1.0 under heavy tails. - const float kurt_excess = fmaxf(0.0f, td_kurtosis_ema - KURT_GAUSSIAN); + const float kurt_excess = fmaxf(0.0f, td_kurtosis_ema - isv[RL_KURT_GAUSSIAN_INDEX]); const float kurt_lift_scale = isv[RL_KURT_LIFT_SCALE_INDEX]; float target = 0.4f + 0.2f * (kurt_excess / kurt_lift_scale); target = fmaxf(PER_ALPHA_MIN, fminf(target, PER_ALPHA_MAX)); diff --git a/crates/ml-alpha/cuda/rl_ppo_clip_controller.cu b/crates/ml-alpha/cuda/rl_ppo_clip_controller.cu index 350b33d29..457a1f5ab 100644 --- a/crates/ml-alpha/cuda/rl_ppo_clip_controller.cu +++ b/crates/ml-alpha/cuda/rl_ppo_clip_controller.cu @@ -44,22 +44,14 @@ #define RL_PPO_CLIP_INDEX 402 #define EPS_MIN 0.05f #define EPS_MAX 0.5f -#define EPS_BOOTSTRAP 0.2f -// ISV-driven KL target per `feedback_isv_for_adaptive_bounds`. Default -// 0.01 (canonical PPO Schulman 2017). Tunable at runtime by re-seeding -// via `rl_isv_write`. The NOISE_FLOOR is derived multiplicatively from -// this (target × 0.01) so adjusting the target proportionally adjusts -// the floor — no separate slot needed. -#define RL_KL_TARGET_INDEX 454 -// Noise-floor fraction: KL below `target × this` is dominated by -// numerical noise — hold ε unchanged. -#define KL_NOISE_FLOOR_FRAC 0.01f -// In-band tolerance: KL within ±33% of target is "good enough" and -// the controller holds. Outside this band it adjusts. -#define KL_TOLERANCE 1.5f -// Per-step multiplicative adjustment when KL is out-of-band. -#define KL_ADJUST_RATE 1.5f -#define WIENER_ALPHA_FLOOR 0.4f +// ISV-driven bootstrap + KL target per `feedback_isv_for_adaptive_bounds`. +#define RL_EPS_BOOTSTRAP_INDEX 474 +#define RL_KL_TARGET_INDEX 454 +// Schulman pattern parameters — global slots shared by 4 controllers. +#define RL_SCHULMAN_TOLERANCE_INDEX 468 +#define RL_SCHULMAN_ADJUST_RATE_INDEX 469 +#define KL_NOISE_FLOOR_FRAC 0.01f +#define WIENER_ALPHA_FLOOR 0.4f // ───────────────────────────────────────────────────────────────────── @@ -91,7 +83,7 @@ extern "C" __global__ void rl_ppo_clip_controller( const float eps_prev = isv[RL_PPO_CLIP_INDEX]; // Bootstrap on sentinel 0.0 per pearl_first_observation_bootstrap. if (eps_prev == 0.0f) { - isv[RL_PPO_CLIP_INDEX] = EPS_BOOTSTRAP; + isv[RL_PPO_CLIP_INDEX] = isv[RL_EPS_BOOTSTRAP_INDEX]; return; } @@ -108,15 +100,15 @@ extern "C" __global__ void rl_ppo_clip_controller( if (kl_ema < kl_noise_floor) return; // Bounded multiplicative adjustment (Schulman-style adaptive KL). + // Schulman params from ISV — shared with target_tau, rollout_steps. + const float tolerance = isv[RL_SCHULMAN_TOLERANCE_INDEX]; + const float adjust_rate = isv[RL_SCHULMAN_ADJUST_RATE_INDEX]; float ratio; - if (kl_ema > kl_target * KL_TOLERANCE) { - // KL too high → tighten ε. - ratio = 1.0f / KL_ADJUST_RATE; - } else if (kl_ema < kl_target / KL_TOLERANCE) { - // KL too low (but above noise floor) → widen ε. - ratio = KL_ADJUST_RATE; + if (kl_ema > kl_target * tolerance) { + ratio = 1.0f / adjust_rate; + } else if (kl_ema < kl_target / tolerance) { + ratio = adjust_rate; } else { - // In-band: hold. ratio = 1.0f; } float eps_target = eps_prev * ratio; @@ -127,7 +119,7 @@ extern "C" __global__ void rl_ppo_clip_controller( // for the rationale — the Wiener blend's bootstrap contamination // would otherwise leave 60% of the canonical default in the // output on the first real KL observation. - if (eps_prev == EPS_BOOTSTRAP) { + if (eps_prev == isv[RL_EPS_BOOTSTRAP_INDEX]) { isv[RL_PPO_CLIP_INDEX] = eps_target; return; } diff --git a/crates/ml-alpha/cuda/rl_ppo_ratio_clamp_controller.cu b/crates/ml-alpha/cuda/rl_ppo_ratio_clamp_controller.cu index 091bc0517..210bd7b7c 100644 --- a/crates/ml-alpha/cuda/rl_ppo_ratio_clamp_controller.cu +++ b/crates/ml-alpha/cuda/rl_ppo_ratio_clamp_controller.cu @@ -47,7 +47,8 @@ #define RL_PPO_CLIP_INDEX 402 #define RL_PPO_RATIO_CLAMP_MAX_INDEX 440 -#define PPO_RATIO_CLAMP_BOOTSTRAP 10.0f +// ISV-driven bootstrap (was 10.0f #define). +#define RL_PPO_RATIO_CLAMP_BOOTSTRAP_INDEX 477 // ISV-driven margin multiplier per `feedback_isv_for_adaptive_bounds`. // Default 10.0 — clamp_max = (1+ε) × this. Seeded by rl_isv_write // at init. Tuning lower tightens the importance-ratio bound. @@ -87,7 +88,7 @@ extern "C" __global__ void rl_ppo_ratio_clamp_controller( // Bootstrap on sentinel 0.0 per pearl_first_observation_bootstrap. if (prev == 0.0f) { - isv[RL_PPO_RATIO_CLAMP_MAX_INDEX] = PPO_RATIO_CLAMP_BOOTSTRAP; + isv[RL_PPO_RATIO_CLAMP_MAX_INDEX] = isv[RL_PPO_RATIO_CLAMP_BOOTSTRAP_INDEX]; return; } @@ -106,7 +107,7 @@ extern "C" __global__ void rl_ppo_ratio_clamp_controller( // step and this is the first real ε observation"; blending with the // bootstrap would leave 60% of the 10.0 constant in the output even // when ε says target should be much smaller (or larger). - if (prev == PPO_RATIO_CLAMP_BOOTSTRAP) { + if (prev == isv[RL_PPO_RATIO_CLAMP_BOOTSTRAP_INDEX]) { isv[RL_PPO_RATIO_CLAMP_MAX_INDEX] = target; return; } diff --git a/crates/ml-alpha/cuda/rl_reward_scale_controller.cu b/crates/ml-alpha/cuda/rl_reward_scale_controller.cu index 2ed960365..fe07aff8d 100644 --- a/crates/ml-alpha/cuda/rl_reward_scale_controller.cu +++ b/crates/ml-alpha/cuda/rl_reward_scale_controller.cu @@ -38,7 +38,8 @@ #define RL_REWARD_SCALE_INDEX 406 #define REWARD_SCALE_MIN 1e-3f #define REWARD_SCALE_MAX 1e3f -#define REWARD_SCALE_BOOTSTRAP 1.0f +// ISV-driven bootstrap (was 1.0f #define). +#define RL_REWARD_SCALE_BOOTSTRAP_INDEX 476 // Floor for the mean_abs_pnl_ema denominator to guard against div-by-zero // when the EMA hasn't accumulated any closed trades yet. #define EPS_PNL 1e-3f @@ -77,7 +78,7 @@ extern "C" __global__ void rl_reward_scale_controller( const float prev = isv[RL_REWARD_SCALE_INDEX]; // Bootstrap on sentinel 0.0 per pearl_first_observation_bootstrap. if (prev == 0.0f) { - isv[RL_REWARD_SCALE_INDEX] = REWARD_SCALE_BOOTSTRAP; + isv[RL_REWARD_SCALE_INDEX] = isv[RL_REWARD_SCALE_BOOTSTRAP_INDEX]; return; } @@ -114,7 +115,7 @@ extern "C" __global__ void rl_reward_scale_controller( // feeds V regression at magnitude 500× the atom support's // ±1 expectation. Replacing directly with target eliminates // this cold-start contamination. - if (prev == REWARD_SCALE_BOOTSTRAP) { + if (prev == isv[RL_REWARD_SCALE_BOOTSTRAP_INDEX]) { isv[RL_REWARD_SCALE_INDEX] = target; return; } diff --git a/crates/ml-alpha/cuda/rl_rollout_steps_controller.cu b/crates/ml-alpha/cuda/rl_rollout_steps_controller.cu index be623aaa5..ad546a823 100644 --- a/crates/ml-alpha/cuda/rl_rollout_steps_controller.cu +++ b/crates/ml-alpha/cuda/rl_rollout_steps_controller.cu @@ -45,29 +45,17 @@ // 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 -// 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 -// Per-step bounded multiplicative adjustment. 1.5× == 50% growth / -// 33% shrinkage per fire. -#define ADV_VAR_RATIO_ADJUST_RATE 1.5f -#define WIENER_ALPHA_FLOOR 0.4f +#define RL_N_ROLLOUT_STEPS_INDEX 404 +#define ROLLOUT_MIN 256.0f +#define ROLLOUT_MAX 8192.0f +// ISV-driven bootstrap + target. +#define RL_ROLLOUT_BOOTSTRAP_INDEX 475 +#define RL_ADV_VAR_RATIO_TARGET_INDEX 449 +// Schulman params from shared global slots (same as ppo_clip, target_tau). +#define RL_SCHULMAN_TOLERANCE_INDEX 468 +#define RL_SCHULMAN_ADJUST_RATE_INDEX 469 +#define ADV_VAR_RATIO_NOISE_FLOOR_FRAC 0.01f +#define WIENER_ALPHA_FLOOR 0.4f // ───────────────────────────────────────────────────────────────────── // rl_rollout_steps_controller: @@ -102,7 +90,7 @@ extern "C" __global__ void rl_rollout_steps_controller( const float prev = isv[RL_N_ROLLOUT_STEPS_INDEX]; // Bootstrap on sentinel 0.0 per pearl_first_observation_bootstrap. if (prev == 0.0f) { - isv[RL_N_ROLLOUT_STEPS_INDEX] = ROLLOUT_BOOTSTRAP; + isv[RL_N_ROLLOUT_STEPS_INDEX] = isv[RL_ROLLOUT_BOOTSTRAP_INDEX]; return; } @@ -128,17 +116,14 @@ extern "C" __global__ void rl_rollout_steps_controller( // how far input is from target. After several consecutive // out-of-band observations the output drifts smoothly toward // MIN/MAX, but a single observation can't slam it there. + const float tolerance = isv[RL_SCHULMAN_TOLERANCE_INDEX]; + const float adjust_rate = isv[RL_SCHULMAN_ADJUST_RATE_INDEX]; float scale; - if (advantage_var_over_abs_mean - > adv_var_target * ADV_VAR_RATIO_TOLERANCE) { - // Noisy → widen. - scale = ADV_VAR_RATIO_ADJUST_RATE; - } else if (advantage_var_over_abs_mean - < adv_var_target / ADV_VAR_RATIO_TOLERANCE) { - // Clean → shrink. - scale = 1.0f / ADV_VAR_RATIO_ADJUST_RATE; + if (advantage_var_over_abs_mean > adv_var_target * tolerance) { + scale = adjust_rate; + } else if (advantage_var_over_abs_mean < adv_var_target / tolerance) { + scale = 1.0f / adjust_rate; } else { - // In-band: hold. scale = 1.0f; } float target = prev * scale; @@ -146,7 +131,7 @@ extern "C" __global__ void rl_rollout_steps_controller( // First-observation replace-directly per // `pearl_first_observation_bootstrap`. - if (prev == ROLLOUT_BOOTSTRAP) { + if (prev == isv[RL_ROLLOUT_BOOTSTRAP_INDEX]) { isv[RL_N_ROLLOUT_STEPS_INDEX] = target; return; } diff --git a/crates/ml-alpha/cuda/rl_target_tau_controller.cu b/crates/ml-alpha/cuda/rl_target_tau_controller.cu index 7290a9dc8..f1a9f9b7e 100644 --- a/crates/ml-alpha/cuda/rl_target_tau_controller.cu +++ b/crates/ml-alpha/cuda/rl_target_tau_controller.cu @@ -28,17 +28,14 @@ #define RL_TARGET_TAU_INDEX 401 #define TAU_MIN 0.001f #define TAU_MAX 0.05f -#define TAU_BOOTSTRAP 0.005f -// ISV-driven Q-divergence target per `feedback_isv_for_adaptive_bounds`. -// Default 0.01 (seeded by rl_isv_write at init). Noise floor derives -// multiplicatively from this value. -#define RL_DIV_TARGET_INDEX 457 -#define DIV_NOISE_FLOOR_FRAC 0.01f -// In-band tolerance: divergence within ±33 % of target → hold. -#define DIV_TOLERANCE 1.5f -// Bounded per-step multiplicative adjustment. -#define DIV_ADJUST_RATE 1.5f -#define WIENER_ALPHA_FLOOR 0.4f +// ISV-driven bootstrap + target. +#define RL_TAU_BOOTSTRAP_INDEX 473 +#define RL_DIV_TARGET_INDEX 457 +// Schulman params from shared global slots (same as ppo_clip). +#define RL_SCHULMAN_TOLERANCE_INDEX 468 +#define RL_SCHULMAN_ADJUST_RATE_INDEX 469 +#define DIV_NOISE_FLOOR_FRAC 0.01f +#define WIENER_ALPHA_FLOOR 0.4f // ───────────────────────────────────────────────────────────────────── @@ -70,7 +67,7 @@ extern "C" __global__ void rl_target_tau_controller( const float tau_prev = isv[RL_TARGET_TAU_INDEX]; // Bootstrap on sentinel 0.0 per pearl_first_observation_bootstrap. if (tau_prev == 0.0f) { - isv[RL_TARGET_TAU_INDEX] = TAU_BOOTSTRAP; + isv[RL_TARGET_TAU_INDEX] = isv[RL_TAU_BOOTSTRAP_INDEX]; return; } @@ -87,10 +84,12 @@ extern "C" __global__ void rl_target_tau_controller( // Bounded multiplicative adjustment. float ratio; - if (q_divergence_norm > div_target * DIV_TOLERANCE) { - ratio = DIV_ADJUST_RATE; - } else if (q_divergence_norm < div_target / DIV_TOLERANCE) { - ratio = 1.0f / DIV_ADJUST_RATE; + const float tolerance = isv[RL_SCHULMAN_TOLERANCE_INDEX]; + const float adjust_rate = isv[RL_SCHULMAN_ADJUST_RATE_INDEX]; + if (q_divergence_norm > div_target * tolerance) { + ratio = adjust_rate; + } else if (q_divergence_norm < div_target / tolerance) { + ratio = 1.0f / adjust_rate; } else { ratio = 1.0f; } @@ -106,7 +105,7 @@ extern "C" __global__ void rl_target_tau_controller( // the controller's first emit. Write target directly instead. // Subsequent steps (where prev has drifted off bootstrap via // earlier blends) take the Wiener path below. - if (tau_prev == TAU_BOOTSTRAP) { + if (tau_prev == isv[RL_TAU_BOOTSTRAP_INDEX]) { isv[RL_TARGET_TAU_INDEX] = tau_target; return; } diff --git a/crates/ml-alpha/cuda/rl_var_over_abs_mean_streaming.cu b/crates/ml-alpha/cuda/rl_var_over_abs_mean_streaming.cu index 2e1430a20..903d2aa18 100644 --- a/crates/ml-alpha/cuda/rl_var_over_abs_mean_streaming.cu +++ b/crates/ml-alpha/cuda/rl_var_over_abs_mean_streaming.cu @@ -38,7 +38,8 @@ // Per `pearl_no_atomicadd`: single-thread kernel (all state lives in // ISV slots updated by thread 0, output also written by thread 0). -#define STREAM_ALPHA 0.05f +// ISV-driven streaming-EMA α (was 0.05f #define). +#define RL_STREAM_ALPHA_INDEX 470 #define ABS_MEAN_FLOOR 1e-6f // out_slot: ISV slot to write var/|mean| to (e.g. @@ -68,14 +69,17 @@ extern "C" __global__ void rl_var_over_abs_mean_streaming( for (int b = 0; b < b_size; ++b) batch_mean += x[b]; batch_mean /= (float)b_size; + // ISV-driven streaming-EMA α (shared with kurtosis kernel). + const float stream_alpha = isv[RL_STREAM_ALPHA_INDEX]; + // Streaming mean update with first-observation bootstrap. const float mean_prev = isv[mean_slot]; float mean_new; if (mean_prev == 0.0f) { mean_new = batch_mean; } else { - mean_new = (1.0f - STREAM_ALPHA) * mean_prev - + STREAM_ALPHA * batch_mean; + mean_new = (1.0f - stream_alpha) * mean_prev + + stream_alpha * batch_mean; } isv[mean_slot] = mean_new; @@ -87,12 +91,10 @@ extern "C" __global__ void rl_var_over_abs_mean_streaming( const float m2_prev = isv[m2_slot]; float m2_new; if (m2_prev == 0.0f && mean_prev == 0.0f) { - // First observation: M2 starts at the current squared dev - // (which is 0 since batch_mean = mean_new on the first step). m2_new = dev2; } else { - m2_new = (1.0f - STREAM_ALPHA) * m2_prev - + STREAM_ALPHA * dev2; + m2_new = (1.0f - stream_alpha) * m2_prev + + stream_alpha * dev2; } isv[m2_slot] = m2_new; diff --git a/crates/ml-alpha/examples/alpha_rl_train.rs b/crates/ml-alpha/examples/alpha_rl_train.rs index 394ffb138..cc65f19b7 100644 --- a/crates/ml-alpha/examples/alpha_rl_train.rs +++ b/crates/ml-alpha/examples/alpha_rl_train.rs @@ -47,12 +47,16 @@ use ml_alpha::rl::isv_slots::{ 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_DIV_TARGET_INDEX, RL_ENTROPY_TARGET_FRAC_INDEX, - RL_IMPROVEMENT_THRESHOLD_INDEX, RL_KL_TARGET_INDEX, RL_KURT_LIFT_SCALE_INDEX, + RL_EPS_BOOTSTRAP_INDEX, RL_IMPROVEMENT_THRESHOLD_INDEX, RL_KL_TARGET_INDEX, + RL_KURT_GAUSSIAN_INDEX, RL_KURT_LIFT_SCALE_INDEX, RL_KURT_NOISE_FLOOR_INDEX, RL_K_LOOP_DIVISOR_INDEX, RL_K_LOOP_MAX_INDEX, RL_LOSS_LAMBDA_AUX_INDEX, RL_LR_BOOTSTRAP_INDEX, RL_LR_DECAY_FACTOR_INDEX, RL_LR_LOSS_EMA_ALPHA_INDEX, RL_LR_MAX_INDEX, RL_LR_MIN_INDEX, RL_LR_WARMUP_STEPS_INDEX, - RL_PLATEAU_PATIENCE_INDEX, RL_PPO_CLAMP_MARGIN_INDEX, RL_Q_ARG_VS_PI_AGREE_INDEX, - RL_REWARD_CLAMP_LOSS_INDEX, RL_REWARD_CLAMP_WIN_INDEX, + RL_PLATEAU_PATIENCE_INDEX, RL_PPO_CLAMP_MARGIN_INDEX, + RL_PPO_RATIO_CLAMP_BOOTSTRAP_INDEX, RL_Q_ARG_VS_PI_AGREE_INDEX, + RL_REWARD_CLAMP_LOSS_INDEX, RL_REWARD_CLAMP_WIN_INDEX, RL_REWARD_SCALE_BOOTSTRAP_INDEX, + RL_ROLLOUT_BOOTSTRAP_INDEX, RL_SCHULMAN_ADJUST_RATE_INDEX, RL_SCHULMAN_TOLERANCE_INDEX, + RL_STREAM_ALPHA_INDEX, RL_TAU_BOOTSTRAP_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, @@ -404,6 +408,15 @@ fn main() -> Result<()> { // Per-step scratch (replaced each step via mapped-pinned helpers). // ── Training loop. ─────────────────────────────────────────────── + // Windowed action distribution — EMA-smoothed across recent steps + // so action_entropy is meaningful at b_size=1 (per-step act_hist + // is one-hot, H=0; windowed dist captures actual policy variety). + // α = 1/1000 → half-life ≈ 690 steps. Tracks the recent ~1k steps' + // action distribution; entropy of normalised windowed_act_hist is + // a real exploration signal. + let mut windowed_act_hist: [f32; 9] = [0.0; 9]; + const WINDOWED_ACT_ALPHA: f32 = 1.0 / 1000.0; + let t_start = std::time::Instant::now(); for step in 0..cli.n_steps { let pair = loader @@ -478,18 +491,29 @@ fn main() -> Result<()> { act_hist[a as usize] += 1; } } - // audit — action entropy H(action_dist) per step. - // Computed from act_hist; high entropy = exploring, low entropy = - // collapsed to a few actions. Already computable post-hoc from - // action_hist but emitting saves analysis cost + makes - // exploration-collapse visible at a glance. + // audit — action entropy H(action_dist) windowed. + // + // Per-batch act_hist is one-hot at b_size=1 so per-step entropy + // is structurally 0. Maintain an EMA-smoothed action histogram + // across recent steps (α=1/1000, half-life ≈ 690 steps) and + // compute entropy on the normalised window. This captures the + // actual policy exploration variety over the recent run. let total_actions: u32 = act_hist.iter().sum(); - let action_entropy: f32 = if total_actions > 0 { - act_hist + if total_actions > 0 { + for i in 0..9 { + let p_step = (act_hist[i] as f32) / (total_actions as f32); + windowed_act_hist[i] = (1.0 - WINDOWED_ACT_ALPHA) + * windowed_act_hist[i] + + WINDOWED_ACT_ALPHA * p_step; + } + } + let win_sum: f32 = windowed_act_hist.iter().sum(); + let action_entropy: f32 = if win_sum > 1e-9 { + windowed_act_hist .iter() - .filter(|&&c| c > 0) - .map(|&c| { - let p = (c as f32) / (total_actions as f32); + .filter(|&&p| p > 1e-9) + .map(|&p_raw| { + let p = p_raw / win_sum; -p * p.ln() }) .sum() @@ -693,6 +717,16 @@ fn main() -> Result<()> { "lr_loss_ema_alpha": isv[RL_LR_LOSS_EMA_ALPHA_INDEX], "lr_decay_factor": isv[RL_LR_DECAY_FACTOR_INDEX], "loss_lambda_aux": isv[RL_LOSS_LAMBDA_AUX_INDEX], + "schulman_tolerance": isv[RL_SCHULMAN_TOLERANCE_INDEX], + "schulman_adjust_rate": isv[RL_SCHULMAN_ADJUST_RATE_INDEX], + "stream_alpha": isv[RL_STREAM_ALPHA_INDEX], + "kurt_gaussian": isv[RL_KURT_GAUSSIAN_INDEX], + "kurt_noise_floor": isv[RL_KURT_NOISE_FLOOR_INDEX], + "tau_bootstrap": isv[RL_TAU_BOOTSTRAP_INDEX], + "eps_bootstrap": isv[RL_EPS_BOOTSTRAP_INDEX], + "rollout_bootstrap": isv[RL_ROLLOUT_BOOTSTRAP_INDEX], + "reward_scale_bootstrap":isv[RL_REWARD_SCALE_BOOTSTRAP_INDEX], + "ppo_ratio_clamp_bootstrap": isv[RL_PPO_RATIO_CLAMP_BOOTSTRAP_INDEX], }, // audit — Q vs π action agreement EMA at slot 407 // (previously dead). 1.0 = perfect ranking consistency, diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index 15d594d0a..c3aa77618 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -35,9 +35,11 @@ //! | 447-448 | Streaming output clamp ceilings | rl_streaming_clamp_init (seed) | //! | 449 | adv_var_ratio target (controller) | rl_streaming_clamp_init (seed) | //! | 450-451 | K-loop divisor + max | rl_streaming_clamp_init (seed) | -//! | 452-461 | 10 ISV-driven design constants | rl_isv_write (seed loop) | +//! | 452-461 | 10 ISV-driven design constants (first batch) | rl_isv_write (seed loop) | +//! | 462-467 | 6 LR-controller numerics + aux λ | rl_isv_write (seed loop) | +//! | 468-477 | 10 design constants (Schulman + bootstraps + streaming α + kurt refs) | rl_isv_write | //! -//! Total: 62 slots; `RL_SLOTS_END = 462`. +//! Total: 78 slots; `RL_SLOTS_END = 478`. /// Discount factor γ. Controller input: mean trade duration / anchor. /// Bootstrap 0.99. @@ -503,6 +505,49 @@ pub const RL_LR_DECAY_FACTOR_INDEX: usize = 466; // separate path. Adding aux to ISV for full coverage: pub const RL_LOSS_LAMBDA_AUX_INDEX: usize = 467; +// ───────────────────────────────────────────────────────────────────── +// Schulman bounded-step pattern parameters (shared by 4 multiplicative +// controllers — ppo_clip, target_tau, rollout_steps, per_α). Made +// global because all 4 use the same algorithm with identical defaults. + +/// In-band tolerance for Schulman controllers: input within ±33 % of +/// target → hold. Default 1.5. +pub const RL_SCHULMAN_TOLERANCE_INDEX: usize = 468; + +/// Per-step bounded multiplicative adjustment for Schulman controllers. +/// Default 1.5 (50 % growth / 33 % shrinkage per fire). +pub const RL_SCHULMAN_ADJUST_RATE_INDEX: usize = 469; + +/// Streaming-EMA α shared by var/|mean| and kurtosis streaming kernels. +/// Default 0.05 (half-life ≈ 14 steps). +pub const RL_STREAM_ALPHA_INDEX: usize = 470; + +/// Kurtosis of standard normal (Pearson). Default 3.0. Tuning permits +/// domain-specific overrides. +pub const RL_KURT_GAUSSIAN_INDEX: usize = 471; + +/// Kurtosis noise-floor for per_α controller. Default 1.0. +pub const RL_KURT_NOISE_FLOOR_INDEX: usize = 472; + +// Controller bootstrap values — first-emit defaults previously held +// in kernel-side `#define`s. + +/// τ bootstrap (target-net soft update Polyak step). Default 0.005. +pub const RL_TAU_BOOTSTRAP_INDEX: usize = 473; + +/// ε bootstrap (PPO clip half-width canonical default). Default 0.2. +pub const RL_EPS_BOOTSTRAP_INDEX: usize = 474; + +/// Rollout-steps bootstrap. Default 2048 (matches K-loop divisor so +/// K=1 at controller bootstrap). +pub const RL_ROLLOUT_BOOTSTRAP_INDEX: usize = 475; + +/// Reward-scale bootstrap (raw passthrough). Default 1.0. +pub const RL_REWARD_SCALE_BOOTSTRAP_INDEX: usize = 476; + +/// PPO ratio-clamp bootstrap. Default 10.0. +pub const RL_PPO_RATIO_CLAMP_BOOTSTRAP_INDEX: usize = 477; + /// 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 = 468; +pub const RL_SLOTS_END: usize = 478; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index da1792b0b..2fd6c6825 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -1170,7 +1170,7 @@ impl IntegratedTrainer { // (slot, value) pair — pure device write, no HtoD per // `feedback_no_htod_htoh_only_mapped_pinned`. { - let isv_constants: [(usize, f32); 16] = [ + let isv_constants: [(usize, f32); 26] = [ (crate::rl::isv_slots::RL_REWARD_CLAMP_WIN_INDEX, 1.0), (crate::rl::isv_slots::RL_REWARD_CLAMP_LOSS_INDEX, 3.0), (crate::rl::isv_slots::RL_KL_TARGET_INDEX, 0.01), @@ -1181,13 +1181,23 @@ impl IntegratedTrainer { (crate::rl::isv_slots::RL_KURT_LIFT_SCALE_INDEX, 7.0), (crate::rl::isv_slots::RL_PPO_CLAMP_MARGIN_INDEX, 10.0), (crate::rl::isv_slots::RL_LR_WARMUP_STEPS_INDEX, 2000.0), - // audit additions: (crate::rl::isv_slots::RL_LR_BOOTSTRAP_INDEX, 1e-3), (crate::rl::isv_slots::RL_LR_MIN_INDEX, 1e-4), (crate::rl::isv_slots::RL_LR_MAX_INDEX, 1e-2), (crate::rl::isv_slots::RL_LR_LOSS_EMA_ALPHA_INDEX, 0.05), (crate::rl::isv_slots::RL_LR_DECAY_FACTOR_INDEX, 0.5), (crate::rl::isv_slots::RL_LOSS_LAMBDA_AUX_INDEX, 1.0), + // Batch 3 — Schulman pattern + streaming α + kurt refs + bootstraps. + (crate::rl::isv_slots::RL_SCHULMAN_TOLERANCE_INDEX, 1.5), + (crate::rl::isv_slots::RL_SCHULMAN_ADJUST_RATE_INDEX, 1.5), + (crate::rl::isv_slots::RL_STREAM_ALPHA_INDEX, 0.05), + (crate::rl::isv_slots::RL_KURT_GAUSSIAN_INDEX, 3.0), + (crate::rl::isv_slots::RL_KURT_NOISE_FLOOR_INDEX, 1.0), + (crate::rl::isv_slots::RL_TAU_BOOTSTRAP_INDEX, 0.005), + (crate::rl::isv_slots::RL_EPS_BOOTSTRAP_INDEX, 0.2), + (crate::rl::isv_slots::RL_ROLLOUT_BOOTSTRAP_INDEX, 2048.0), + (crate::rl::isv_slots::RL_REWARD_SCALE_BOOTSTRAP_INDEX, 1.0), + (crate::rl::isv_slots::RL_PPO_RATIO_CLAMP_BOOTSTRAP_INDEX, 10.0), ]; let cfg_isv = LaunchConfig { grid_dim: (1, 1, 1), diff --git a/crates/ml-alpha/tests/isv_bootstrap.rs b/crates/ml-alpha/tests/isv_bootstrap.rs index 8e138c46a..a12299ca5 100644 --- a/crates/ml-alpha/tests/isv_bootstrap.rs +++ b/crates/ml-alpha/tests/isv_bootstrap.rs @@ -24,15 +24,19 @@ use ml_alpha::rl::isv_slots::{ RL_ADV_VAR_RATIO_CLAMP_INDEX, RL_ADV_VAR_RATIO_TARGET_INDEX, RL_DIV_TARGET_INDEX, - RL_ENTROPY_COEF_INDEX, RL_ENTROPY_TARGET_FRAC_INDEX, RL_GAMMA_INDEX, - RL_IMPROVEMENT_THRESHOLD_INDEX, RL_KL_TARGET_INDEX, RL_KURT_LIFT_SCALE_INDEX, + RL_ENTROPY_COEF_INDEX, RL_ENTROPY_TARGET_FRAC_INDEX, RL_EPS_BOOTSTRAP_INDEX, + RL_GAMMA_INDEX, RL_IMPROVEMENT_THRESHOLD_INDEX, RL_KL_TARGET_INDEX, + RL_KURT_GAUSSIAN_INDEX, RL_KURT_LIFT_SCALE_INDEX, RL_KURT_NOISE_FLOOR_INDEX, RL_K_LOOP_DIVISOR_INDEX, RL_K_LOOP_MAX_INDEX, RL_LOSS_LAMBDA_AUX_INDEX, RL_LR_BOOTSTRAP_INDEX, RL_LR_DECAY_FACTOR_INDEX, RL_LR_LOSS_EMA_ALPHA_INDEX, RL_LR_MAX_INDEX, RL_LR_MIN_INDEX, RL_LR_WARMUP_STEPS_INDEX, RL_MEAN_TRADE_DURATION_EMA_INDEX, RL_N_ROLLOUT_STEPS_INDEX, RL_PER_ALPHA_INDEX, RL_PLATEAU_PATIENCE_INDEX, RL_PPO_CLAMP_MARGIN_INDEX, RL_PPO_CLIP_INDEX, - RL_PPO_RATIO_CLAMP_MAX_INDEX, RL_REWARD_CLAMP_LOSS_INDEX, RL_REWARD_CLAMP_WIN_INDEX, - RL_REWARD_SCALE_INDEX, RL_SLOTS_END, RL_TARGET_TAU_INDEX, RL_TD_KURTOSIS_CLAMP_INDEX, + RL_PPO_RATIO_CLAMP_BOOTSTRAP_INDEX, RL_PPO_RATIO_CLAMP_MAX_INDEX, + RL_REWARD_CLAMP_LOSS_INDEX, RL_REWARD_CLAMP_WIN_INDEX, RL_REWARD_SCALE_BOOTSTRAP_INDEX, + RL_REWARD_SCALE_INDEX, RL_ROLLOUT_BOOTSTRAP_INDEX, RL_SCHULMAN_ADJUST_RATE_INDEX, + RL_SCHULMAN_TOLERANCE_INDEX, RL_SLOTS_END, RL_STREAM_ALPHA_INDEX, RL_TARGET_TAU_INDEX, + RL_TAU_BOOTSTRAP_INDEX, RL_TD_KURTOSIS_CLAMP_INDEX, }; use ml_alpha::trainer::integrated::{IntegratedTrainer, IntegratedTrainerConfig}; use ml_alpha::trainer::perception::PerceptionTrainerConfig; @@ -174,6 +178,16 @@ fn g1_isv_bootstrap_writes_canonical_values() { || slot == RL_LR_LOSS_EMA_ALPHA_INDEX || slot == RL_LR_DECAY_FACTOR_INDEX || slot == RL_LOSS_LAMBDA_AUX_INDEX + || slot == RL_SCHULMAN_TOLERANCE_INDEX + || slot == RL_SCHULMAN_ADJUST_RATE_INDEX + || slot == RL_STREAM_ALPHA_INDEX + || slot == RL_KURT_GAUSSIAN_INDEX + || slot == RL_KURT_NOISE_FLOOR_INDEX + || slot == RL_TAU_BOOTSTRAP_INDEX + || slot == RL_EPS_BOOTSTRAP_INDEX + || slot == RL_ROLLOUT_BOOTSTRAP_INDEX + || slot == RL_REWARD_SCALE_BOOTSTRAP_INDEX + || slot == RL_PPO_RATIO_CLAMP_BOOTSTRAP_INDEX { continue; } @@ -236,6 +250,16 @@ fn g1_isv_bootstrap_writes_canonical_values() { assert!((isv[RL_LR_LOSS_EMA_ALPHA_INDEX] - 0.05).abs() < EPS); assert!((isv[RL_LR_DECAY_FACTOR_INDEX] - 0.5).abs() < EPS); assert!((isv[RL_LOSS_LAMBDA_AUX_INDEX] - 1.0).abs() < EPS); + assert!((isv[RL_SCHULMAN_TOLERANCE_INDEX] - 1.5).abs() < EPS); + assert!((isv[RL_SCHULMAN_ADJUST_RATE_INDEX] - 1.5).abs() < EPS); + assert!((isv[RL_STREAM_ALPHA_INDEX] - 0.05).abs() < EPS); + assert!((isv[RL_KURT_GAUSSIAN_INDEX] - 3.0).abs() < EPS); + assert!((isv[RL_KURT_NOISE_FLOOR_INDEX] - 1.0).abs() < EPS); + assert!((isv[RL_TAU_BOOTSTRAP_INDEX] - 0.005).abs() < EPS); + assert!((isv[RL_EPS_BOOTSTRAP_INDEX] - 0.2).abs() < EPS); + assert!((isv[RL_ROLLOUT_BOOTSTRAP_INDEX] - 2048.0).abs() < EPS); + assert!((isv[RL_REWARD_SCALE_BOOTSTRAP_INDEX] - 1.0).abs() < EPS); + assert!((isv[RL_PPO_RATIO_CLAMP_BOOTSTRAP_INDEX] - 10.0).abs() < EPS); 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 d86367c08..6ebec92e0 100644 --- a/crates/ml-alpha/tests/r5_controllers_and_soft_update.rs +++ b/crates/ml-alpha/tests/r5_controllers_and_soft_update.rs @@ -28,16 +28,20 @@ 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_DIV_TARGET_INDEX, RL_ENTROPY_COEF_INDEX, - RL_ENTROPY_OBSERVED_EMA_INDEX, RL_ENTROPY_TARGET_FRAC_INDEX, RL_GAMMA_INDEX, - RL_IMPROVEMENT_THRESHOLD_INDEX, RL_KL_PI_EMA_INDEX, RL_KL_TARGET_INDEX, - RL_KURT_LIFT_SCALE_INDEX, RL_K_LOOP_DIVISOR_INDEX, RL_K_LOOP_MAX_INDEX, + RL_ENTROPY_OBSERVED_EMA_INDEX, RL_ENTROPY_TARGET_FRAC_INDEX, RL_EPS_BOOTSTRAP_INDEX, + RL_GAMMA_INDEX, RL_IMPROVEMENT_THRESHOLD_INDEX, RL_KL_PI_EMA_INDEX, + RL_KL_TARGET_INDEX, RL_KURT_GAUSSIAN_INDEX, RL_KURT_LIFT_SCALE_INDEX, + RL_KURT_NOISE_FLOOR_INDEX, RL_K_LOOP_DIVISOR_INDEX, RL_K_LOOP_MAX_INDEX, RL_LOSS_LAMBDA_AUX_INDEX, RL_LR_BOOTSTRAP_INDEX, RL_LR_DECAY_FACTOR_INDEX, RL_LR_LOSS_EMA_ALPHA_INDEX, RL_LR_MAX_INDEX, RL_LR_MIN_INDEX, RL_LR_WARMUP_STEPS_INDEX, RL_MEAN_ABS_PNL_EMA_INDEX, RL_MEAN_TRADE_DURATION_EMA_INDEX, RL_N_ROLLOUT_STEPS_INDEX, RL_PER_ALPHA_INDEX, RL_PLATEAU_PATIENCE_INDEX, RL_PPO_CLAMP_MARGIN_INDEX, - RL_PPO_CLIP_INDEX, RL_PPO_RATIO_CLAMP_MAX_INDEX, RL_Q_DIVERGENCE_EMA_INDEX, - RL_REWARD_CLAMP_LOSS_INDEX, RL_REWARD_CLAMP_WIN_INDEX, RL_REWARD_SCALE_INDEX, - RL_SLOTS_END, RL_TARGET_TAU_INDEX, RL_TD_KURTOSIS_CLAMP_INDEX, RL_TD_KURTOSIS_EMA_INDEX, + RL_PPO_CLIP_INDEX, RL_PPO_RATIO_CLAMP_BOOTSTRAP_INDEX, RL_PPO_RATIO_CLAMP_MAX_INDEX, + RL_Q_DIVERGENCE_EMA_INDEX, RL_REWARD_CLAMP_LOSS_INDEX, RL_REWARD_CLAMP_WIN_INDEX, + RL_REWARD_SCALE_BOOTSTRAP_INDEX, RL_REWARD_SCALE_INDEX, RL_ROLLOUT_BOOTSTRAP_INDEX, + RL_SCHULMAN_ADJUST_RATE_INDEX, RL_SCHULMAN_TOLERANCE_INDEX, RL_SLOTS_END, + RL_STREAM_ALPHA_INDEX, RL_TARGET_TAU_INDEX, RL_TAU_BOOTSTRAP_INDEX, + RL_TD_KURTOSIS_CLAMP_INDEX, RL_TD_KURTOSIS_EMA_INDEX, }; use ml_alpha::trainer::integrated::{IntegratedTrainer, IntegratedTrainerConfig}; use ml_alpha::trainer::perception::PerceptionTrainerConfig; @@ -144,6 +148,16 @@ fn g3_per_step_controllers_move_isv_outputs_when_fed_real_emas() { || slot == RL_LR_LOSS_EMA_ALPHA_INDEX || slot == RL_LR_DECAY_FACTOR_INDEX || slot == RL_LOSS_LAMBDA_AUX_INDEX + || slot == RL_SCHULMAN_TOLERANCE_INDEX + || slot == RL_SCHULMAN_ADJUST_RATE_INDEX + || slot == RL_STREAM_ALPHA_INDEX + || slot == RL_KURT_GAUSSIAN_INDEX + || slot == RL_KURT_NOISE_FLOOR_INDEX + || slot == RL_TAU_BOOTSTRAP_INDEX + || slot == RL_EPS_BOOTSTRAP_INDEX + || slot == RL_ROLLOUT_BOOTSTRAP_INDEX + || slot == RL_REWARD_SCALE_BOOTSTRAP_INDEX + || slot == RL_PPO_RATIO_CLAMP_BOOTSTRAP_INDEX { continue; } @@ -171,6 +185,16 @@ fn g3_per_step_controllers_move_isv_outputs_when_fed_real_emas() { assert!((isv_before[RL_LR_LOSS_EMA_ALPHA_INDEX] - 0.05).abs() < 1e-7); assert_eq!(isv_before[RL_LR_DECAY_FACTOR_INDEX], 0.5); assert_eq!(isv_before[RL_LOSS_LAMBDA_AUX_INDEX], 1.0); + assert_eq!(isv_before[RL_SCHULMAN_TOLERANCE_INDEX], 1.5); + assert_eq!(isv_before[RL_SCHULMAN_ADJUST_RATE_INDEX], 1.5); + assert!((isv_before[RL_STREAM_ALPHA_INDEX] - 0.05).abs() < 1e-7); + assert_eq!(isv_before[RL_KURT_GAUSSIAN_INDEX], 3.0); + assert_eq!(isv_before[RL_KURT_NOISE_FLOOR_INDEX], 1.0); + assert!((isv_before[RL_TAU_BOOTSTRAP_INDEX] - 0.005).abs() < 1e-7); + assert!((isv_before[RL_EPS_BOOTSTRAP_INDEX] - 0.2).abs() < 1e-7); + assert_eq!(isv_before[RL_ROLLOUT_BOOTSTRAP_INDEX], 2048.0); + assert_eq!(isv_before[RL_REWARD_SCALE_BOOTSTRAP_INDEX], 1.0); + assert_eq!(isv_before[RL_PPO_RATIO_CLAMP_BOOTSTRAP_INDEX], 10.0); // Populate each EMA-input slot with a non-zero value via the // R3 ema_update_per_step bootstrap path (sentinel-zero → first