foxhunt/crates/ml-alpha/cuda/rl_ppo_clip_controller.cu

// rl_ppo_clip_controller.cu — emits ε to ISV[RL_PPO_CLIP_INDEX=402].
//
// Phase D of the integrated RL trainer
// (docs/superpowers/plans/2026-05-22-integrated-rl-trainer.md).
//
// ε is the PPO clipped-surrogate band half-width — see
// `ppo_clipped_surrogate.cu` for the consumer. Per
// `pearl_controller_anchors_isv_driven` and
// `feedback_isv_for_adaptive_bounds`, ε is NOT a hardcoded constant:
// it adapts to keep KL(π_new ‖ π_old) ≈ 0.01 (the standard PPO
// "trust region" target).
//
// 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
// (the canonical PPO default). Subsequent emits Wiener-α blend with
// floor 0.4 (per `pearl_wiener_alpha_floor_for_nonstationary` — the KL
// target drifts as π_old co-adapts with π_new, breaking stationarity).
//
// Bounds: ε ∈ [0.05, 0.5]. Below 0.05 the policy is effectively frozen;
// above 0.5 the surrogate is so loose that the clip doesn't bite and
// PPO degenerates into vanilla policy gradient (which destabilises
// when ratios run away).

#define RL_PPO_CLIP_INDEX   402
#define EPS_MIN             0.05f
#define EPS_MAX             0.5f
// 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


// ─────────────────────────────────────────────────────────────────────
// rl_ppo_clip_controller:
//   Single-thread controller — writes ONE float to isv[RL_PPO_CLIP_INDEX].
//
// Inputs:
//   isv         [≥ RL_PPO_CLIP_INDEX+1] — ISV bus
//   alpha       — Wiener-α from the controller's signal stats; floored at
//                 WIENER_ALPHA_FLOOR before the blend.
//   kl_ema      — caller-computed EMA of KL(π_new ‖ π_old). Phase E
//                 produces this via a small reduce-norm kernel over the
//                 per-batch log-prob delta from the surrogate fwd kernel.
//
// Outputs:
//   isv[RL_PPO_CLIP_INDEX] — ε ∈ [EPS_MIN, EPS_MAX]
// ─────────────────────────────────────────────────────────────────────
// Phase R5: scalar input arg replaced with `input_slot` ISV index so the
// EMA producer (Phase R3 ema_update_per_step targeting
// ISV[RL_KL_PI_EMA_INDEX=419]) feeds this controller without any host
// roundtrip per `feedback_cpu_is_read_only`.
extern "C" __global__ void rl_ppo_clip_controller(
    float* __restrict__ isv,
    float          alpha,
    int            input_slot
) {
    if (threadIdx.x != 0 || blockIdx.x != 0) return;

    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] = isv[RL_EPS_BOOTSTRAP_INDEX];
        return;
    }

    // Cold-start gate (sentinel-zero EMA — kernel hasn't fired yet).
    const float kl_ema = isv[input_slot];
    if (kl_ema == 0.0f) return;

    // ISV-driven KL target (was hardcoded #define).
    const float kl_target = isv[RL_KL_TARGET_INDEX];
    const float kl_noise_floor = kl_target * KL_NOISE_FLOOR_FRAC;

    // Noise-floor gate: KL below kl_noise_floor is dominated by
    // numerical noise, not real policy divergence. Hold ε unchanged.
    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 * tolerance) {
        ratio = 1.0f / adjust_rate;
    } else if (kl_ema < kl_target / tolerance) {
        ratio = adjust_rate;
    } else {
        ratio = 1.0f;
    }
    float eps_target = eps_prev * ratio;
    eps_target = fmaxf(EPS_MIN, fminf(eps_target, EPS_MAX));

    // First-observation replace-directly per
    // `pearl_first_observation_bootstrap`. See rl_target_tau_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 == isv[RL_EPS_BOOTSTRAP_INDEX]) {
        isv[RL_PPO_CLIP_INDEX] = eps_target;
        return;
    }

    // Wiener-α blend with floor per pearl_wiener_alpha_floor_for_nonstationary.
    const float a       = fmaxf(alpha, WIENER_ALPHA_FLOOR);
    float       eps_new = (1.0f - a) * eps_prev + a * eps_target;

    eps_new = fmaxf(EPS_MIN, fminf(eps_new, EPS_MAX));
    isv[RL_PPO_CLIP_INDEX] = eps_new;
}