fix(rl): bound multiplicative controllers + add KL noise-floor gate

mjzfk diag (commit 53aeef099) showed PPO clip ε pegged at MAX=0.5
for 100% of the 50k-step run, despite kl_pi_ema median = 5.3e-9 and
max = 3.4e-4 (well below KL_TARGET=0.01). The widened clip band is
why ratio_clamp settled at (1+0.5)*10=15 instead of 12, and why π
took no meaningful updates — KL was essentially zero meaning the
policy wasn't moving.

## Root cause

The controller's adaptation was `ratio = KL_TARGET / max(kl_ema, 1e-6)`
— a multiplicative formula with no per-step bound. With kl_ema=1e-11
the kernel sees:

  ratio = 0.01 / max(1e-11, 1e-6) = 0.01 / 1e-6 = 10000
  target = eps_prev * 10000 = clamped to EPS_MAX

First-observation replace-directly then locks ε at MAX immediately,
and the Wiener α-floor=0.4 blend keeps it there forever.

The cold-start `if (kl_ema == 0.0f) return;` gate only caught EXACT
zero — first observable but tiny KL (typical: cold-start LR not yet
producing measurable policy drift) blows past the gate and saturates
the multiplier.

Same dangerous pattern existed in `rl_target_tau_controller` (uses
`q_div / DIV_TARGET` ratio with no per-step bound) — hadn't bitten
because q_divergence_norm naturally lives in the [0.01, 0.1] range,
but a quiet initialisation could hit it the same way.

## Fix: Schulman-style bounded adaptive KL

Both controllers now use a discrete-step adjustment:

  * input > target × TOLERANCE (1.5)   → ratio = ADJUST_RATE (1.5)
  * input < target / TOLERANCE          → ratio = 1/ADJUST_RATE
  * in-band                              → ratio = 1.0 (hold)

Per-step adjustment is bounded at 1.5× (50% expansion / 33%
shrinkage), so no single observation can swing the output across the
[MIN, MAX] range regardless of how outlier-tiny or outlier-huge it
is. After several consecutive out-of-band observations the output
drifts smoothly toward MIN/MAX, but the response is dampened.

## Noise-floor gate

In addition to the bounded step, both controllers now hold their
output when the input EMA is below a noise floor:

  KL_NOISE_FLOOR  = KL_TARGET  × 0.01 = 1e-4  (ppo_clip)
  DIV_NOISE_FLOOR = DIV_TARGET × 0.01 = 1e-4  (target_tau)

Two orders of magnitude below the design target = "policy isn't
actually updating" / "Q hasn't started learning" / numerical noise.
Reacting to this signal can only mis-tune the controller — we'd
rather hold a sane default than chase noise.

The TARGET-derived floors (rather than absolute constants) mean
adjusting KL_TARGET / DIV_TARGET shifts the floors proportionally —
consistent with the existing pattern.

## ISV discipline

Per `feedback_isv_for_adaptive_bounds`: KL_TARGET and DIV_TARGET
themselves are PPO/DQN design constants (like REWARD_CLAMP_WIN =
1.0 in apply_reward_scale, or LR_BOOTSTRAP = 1e-3 in
rl_lr_controller). The NOISE_FLOOR derives from them, and the
TOLERANCE / ADJUST_RATE constants are structural Schulman-recipe
parameters that don't adapt at runtime. Existing diag exposes both
controllers' outputs (isv_out.ppo_clip_eps, isv_out.target_tau) and
inputs (isv_ema_in.kl_pi, isv_ema_in.q_divergence) so the
controller behaviour is fully observable from the JSONL — no new
slots needed.

## Verified gates (local sm_86)

  G1 isv_bootstrap   ✅
  G3 controllers     ✅ (controllers still move outputs when fed real EMAs)
  G4 target_update   ✅
  integrated_smoke   ✅

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>

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,19 +95,34 @@ 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);
+
+    // 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));
 

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,22 +76,34 @@ 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);
+
+    // 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));