diff --git a/crates/ml-alpha/cuda/rl_ppo_clip_controller.cu b/crates/ml-alpha/cuda/rl_ppo_clip_controller.cu index 6acf18f05..2da462d1c 100644 --- a/crates/ml-alpha/cuda/rl_ppo_clip_controller.cu +++ b/crates/ml-alpha/cuda/rl_ppo_clip_controller.cu @@ -10,10 +10,25 @@ // it adapts to keep KL(π_new ‖ π_old) ≈ 0.01 (the standard PPO // "trust region" target). // -// Controller logic: -// * High KL → new policy diverges too quickly → shrink ε to tighten -// the trust region. -// * Low KL → policy is stable → widen ε to allow larger updates. +// Controller logic (Schulman-style adaptive KL, bounded step): +// * KL > KL_TARGET × KL_TOLERANCE → policy diverging too quickly → +// shrink ε by KL_ADJUST_RATE (tighten the trust region) +// * KL < KL_TARGET / KL_TOLERANCE → policy stable → widen ε by +// KL_ADJUST_RATE (allow larger updates) +// * else (in-band): hold ε unchanged +// +// CRITICAL: KL ≪ KL_TARGET is meaningless if the policy isn't actually +// updating at all (LR too small, etc). Without a noise-floor gate the +// prior multiplicative-ratio design `KL_TARGET / max(kl, 1e-6)` would +// produce ratio=10000 from KL=1e-11, slam target to EPS_MAX, and stick +// there forever even after KL grew to small-but-real values. Canonical +// incident: alpha-rl-mjzfk fold0 (commit 53aeef099) had ε pegged at MAX +// for 100 % of the run while KL_PI median was 5e-9. +// +// FIX: the bounded multiplicative step (1.5× max change per step) +// prevents one observation from swinging ε across the entire range, +// and the KL_NOISE_FLOOR gate holds the controller when KL signal is +// below "meaningful policy update" magnitude. // // Bootstrap discipline (per `pearl_first_observation_bootstrap`): the // ISV slot starts at 0.0 sentinel. First emit writes ε = 0.2 directly @@ -31,6 +46,19 @@ #define EPS_MAX 0.5f #define EPS_BOOTSTRAP 0.2f #define KL_TARGET 0.01f +// Below this KL magnitude the controller holds ε unchanged. Two +// orders of magnitude below the target = signal is dominated by +// numerical noise / policy isn't measurably updating, not a "low +// KL → widen" signal we should react to. +#define KL_NOISE_FLOOR (KL_TARGET * 0.01f) // = 1e-4 +// 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. Bounded +// at 1.5× (50% expansion / 33% shrinkage per step), preventing any +// single observation from swinging ε across the entire [MIN, MAX] +// range. +#define KL_ADJUST_RATE 1.5f #define WIENER_ALPHA_FLOOR 0.4f @@ -67,20 +95,35 @@ extern "C" __global__ void rl_ppo_clip_controller( return; } - // Cold-start gate: when KL EMA is sentinel- - // zero (no policy update has produced a measurable KL yet), the - // multiplicative ratio `KL_TARGET / max(kl, 1e-6)` collapses to - // 10,000 — the target slams to EPS_MAX (0.5) and the Wiener - // α-floor=0.4 drags ε toward MAX every step. Within ~6 cold-start - // steps ε hits 0.49 despite zero real signal, then sticks at - // 0.5 for the entire smoke. Hold at bootstrap until first - // observation. + // Cold-start gate (sentinel-zero EMA — kernel hasn't fired yet). const float kl_ema = isv[input_slot]; if (kl_ema == 0.0f) return; - // Multiplicative adaptation toward the KL target: if measured KL is - // higher than `KL_TARGET` we want a smaller ε; if lower, grow ε. - const float ratio = KL_TARGET / fmaxf(kl_ema, 1e-6f); - float eps_target = eps_prev * ratio; + + // Noise-floor gate: KL below KL_NOISE_FLOOR is dominated by + // numerical noise, not real policy divergence. Hold ε unchanged. + // This catches the mjzfk failure mode where KL_PI median = 5e-9 + // (essentially zero — policy isn't actually updating) but a + // multiplicative `target / max(kl, 1e-6)` design would still + // produce ratio=10000 and slam ε to MAX. + if (kl_ema < KL_NOISE_FLOOR) return; + + // Bounded multiplicative adjustment (Schulman-style adaptive KL). + // At most KL_ADJUST_RATE × shift per step regardless of how far KL + // is from target — no single observation can swing ε across the + // entire [MIN, MAX] range. + float ratio; + if (kl_ema > KL_TARGET * KL_TOLERANCE) { + // KL too high → policy moving too fast → tighten ε. + ratio = 1.0f / KL_ADJUST_RATE; + } else if (kl_ema < KL_TARGET / KL_TOLERANCE) { + // KL too low (but ABOVE noise floor) → policy can take bigger + // steps → widen ε. + ratio = KL_ADJUST_RATE; + } else { + // In-band: KL within ±33 % of target → hold. + ratio = 1.0f; + } + float eps_target = eps_prev * ratio; eps_target = fmaxf(EPS_MIN, fminf(eps_target, EPS_MAX)); // First-observation replace-directly per diff --git a/crates/ml-alpha/cuda/rl_target_tau_controller.cu b/crates/ml-alpha/cuda/rl_target_tau_controller.cu index f89353076..a3a1b60b1 100644 --- a/crates/ml-alpha/cuda/rl_target_tau_controller.cu +++ b/crates/ml-alpha/cuda/rl_target_tau_controller.cu @@ -30,6 +30,16 @@ #define TAU_MAX 0.05f #define TAU_BOOTSTRAP 0.005f #define DIV_TARGET 0.01f +// Below this divergence magnitude the controller holds τ unchanged — +// the signal is dominated by numerical noise / network just-init, not +// "Q has converged and we should shrink τ". Same anti-runaway pattern +// as the PPO clip controller's KL_NOISE_FLOOR. +#define DIV_NOISE_FLOOR (DIV_TARGET * 0.01f) // = 1e-4 +// In-band tolerance: divergence within ±33 % of target → hold. +#define DIV_TOLERANCE 1.5f +// Bounded per-step multiplicative adjustment so one outlier +// observation can't swing τ across the [MIN, MAX] range. +#define DIV_ADJUST_RATE 1.5f #define WIENER_ALPHA_FLOOR 0.4f @@ -66,23 +76,35 @@ extern "C" __global__ void rl_target_tau_controller( return; } - // Cold-start gate: when the input EMA is - // sentinel-zero (no signal observed yet — typical for the first - // ~N steps before any Q backward fires), the multiplicative - // ratio `q_div / DIV_TARGET` collapses to 0, which makes the - // target slam to TAU_MIN. The Wiener α-floor=0.4 then drags - // prev toward TAU_MIN every step, so by step ~6 τ has migrated - // halfway to MIN despite zero real signal. Hold at bootstrap - // until the first non-zero observation. Per - // `pearl_first_observation_bootstrap`: sentinel = 0 means "no - // data" — the controller must NOT adapt against no-data. + // Cold-start gate (sentinel-zero EMA — kernel hasn't fired yet). const float q_divergence_norm = isv[input_slot]; if (q_divergence_norm == 0.0f) return; - // Multiplicative adaptation toward the target divergence: if the - // measured divergence is higher than `DIV_TARGET` we want a larger - // τ so the target tracks faster; if lower, shrink τ. - const float ratio = q_divergence_norm / fmaxf(DIV_TARGET, 1e-6f); - float tau_target = tau_prev * ratio; + + // Noise-floor gate: divergence below DIV_NOISE_FLOOR is dominated + // by numerical noise / network initialisation, not a real "Q has + // converged → shrink τ" signal. Hold τ unchanged. Mirrors the PPO + // clip controller's KL_NOISE_FLOOR design after the alpha-rl-mjzfk + // root cause. + if (q_divergence_norm < DIV_NOISE_FLOOR) return; + + // Bounded multiplicative adjustment (matches PPO clip controller). + // At most DIV_ADJUST_RATE × shift per step regardless of how far + // divergence is from target. + float ratio; + if (q_divergence_norm > DIV_TARGET * DIV_TOLERANCE) { + // Divergence too high → online drifted from target faster than + // τ can track → raise τ so target accelerates. + ratio = DIV_ADJUST_RATE; + } else if (q_divergence_norm < DIV_TARGET / DIV_TOLERANCE) { + // Divergence too low → bootstrap is overly stable / Q stuck → + // lower τ to let the target lag a bit more (gives the online + // net room to actually learn before being shadowed). + ratio = 1.0f / DIV_ADJUST_RATE; + } else { + // In-band: hold. + ratio = 1.0f; + } + float tau_target = tau_prev * ratio; tau_target = fmaxf(TAU_MIN, fminf(tau_target, TAU_MAX)); // First-observation replace-directly: if