708c121f20
kc2h9 confirmed: clamping streaming-kernel outputs to [≤100, ≤30]
had ZERO behavioral impact because rl_rollout_steps_controller's
prior design used `scale = clamp(input/target, 0.5, 2.0)` — the
scale saturated to ±2× on the SIGN of (input − target), not the
magnitude. With target=0.1 and typical input=1–10 the controller
slammed to MAX in ≤4 steps regardless of whether input was 4 or
3e5. Bit-identical losses between gxhr8 and kc2h9 confirmed the
saturation.
## Fix 1: rl_rollout_steps_controller — same Schulman pattern as ppo_clip
* input > TARGET × 1.5 → scale = 1.5 (widen)
* input < TARGET / 1.5 → scale = 1/1.5 (shrink)
* in-band → scale = 1.0 (hold)
* input < TARGET × 0.01 → return (noise floor — hold prev)
Per-step adjustment bounded at 1.5×, so rollout_steps drifts
smoothly toward MIN/MAX rather than slamming there. The noise-floor
gate matches the pattern from
`pearl_multiplicative_controllers_need_bounded_step_and_noise_floor`
applied to the ε and τ controllers earlier in R9.
## Fix 2: rl_per_alpha_controller — noise-floor gate (defensive)
per_α uses a LINEAR lift `0.4 + 0.2·(kurt-3)/7` (not multiplicative),
so it doesn't have the saturation bug. But added a noise-floor gate
at KURT_NOISE_FLOOR = 1.0 so a sub-Gaussian kurtosis reading from
the streaming estimator's startup window (when per-step batch-mean
deviations are small before tails develop) doesn't drag α toward
PER_ALPHA_MIN on cold-start.
## Diag bake-in (per user request "bake in diags")
JSONL gains a `controller_branch` block exposing the
multiplicative-controller inputs alongside their design targets:
controller_branch: {
rollout_steps_input: isv[421], rollout_steps_target: 0.1,
ppo_clip_input: isv[419], ppo_clip_target: 0.01,
target_tau_input: isv[418], target_tau_target: 0.01,
per_alpha_input: isv[422], per_alpha_target: 0.6,
}
Post-hoc analysis can compute the branch each step (WIDEN / HOLD /
SHRINK / NOISE) by comparing input/target against the ±33%
tolerance band, revealing whether each controller is being driven
by real signal or sitting in the in-band hold zone. Targets are
reflected from the kernel #defines (synchronised by code review at
the controller-cu file level — there's no ISV slot for these
design constants because they're fundamental to the controller's
behaviour, not adaptive).
## Verified gates (local sm_86)
G1 isv_bootstrap ✅
G3 controllers ✅
G4 target_update ✅
integrated_smoke ✅
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
146 lines
7.5 KiB
Plaintext
146 lines
7.5 KiB
Plaintext
// rl_per_alpha_controller.cu — emits PER (Prioritized Experience Replay)
|
||
// priority exponent to ISV[RL_PER_ALPHA_INDEX=405].
|
||
//
|
||
// Phase E.3b of the integrated RL trainer
|
||
// (docs/superpowers/plans/2026-05-22-integrated-rl-trainer.md).
|
||
//
|
||
// PER's priority exponent α controls how aggressively the replay buffer
|
||
// concentrates sampling on high-TD-error transitions. Per
|
||
// `pearl_controller_anchors_isv_driven` and
|
||
// `feedback_isv_for_adaptive_bounds`, α is NOT a hardcoded constant:
|
||
// it adapts to the tail thickness of the TD-error distribution.
|
||
//
|
||
// Controller logic:
|
||
// * High TD-error kurtosis (heavy tails) → a few transitions
|
||
// dominate the loss → raise α to sharpen the priority distribution
|
||
// and concentrate sampling on the informative tails.
|
||
// * Low kurtosis (light tails, ≈ Gaussian) → TD-errors are uniform →
|
||
// keep α near the standard PER default (0.6) — over-sharpening on a
|
||
// light-tailed distribution wastes samples on noise.
|
||
//
|
||
// Bootstrap discipline (per `pearl_first_observation_bootstrap`): the
|
||
// ISV slot starts at 0.0 sentinel. First emit DERIVES the bootstrap
|
||
// from the current `input_slot` (td_kurtosis EMA) via the same target
|
||
// formula the per-step path uses, instead of writing a hardcoded
|
||
// "canonical PER 0.6 default". The earlier hardcoded-0.6 pattern
|
||
// coincided with `target(kurt=10) = 0.6` — the canonical
|
||
// heavy-tailed-market kurtosis — creating a Wiener-blend fixed point
|
||
// where `prev = target = 0.6` froze the controller in a dead-zone
|
||
// for any input near canonical. The derive-from-input pattern
|
||
// eliminates the coincidence: bootstrap is whatever the target formula
|
||
// emits for the current input (= 0.4 at sentinel-zero input, = 0.6
|
||
// when input has already EMA'd to canonical 10, = 0.886 at kurt=20,
|
||
// etc.) so the controller's per-step Wiener blend always sees a real
|
||
// `prev` vs `target` delta on subsequent inputs.
|
||
//
|
||
// Subsequent emits Wiener-α blend with floor 0.4 (per
|
||
// `pearl_wiener_alpha_floor_for_nonstationary` — the TD-error
|
||
// distribution drifts as Q co-adapts, breaking stationarity).
|
||
//
|
||
// Bounds: α ∈ [0.3, 1.0]. Below 0.3 the priority distribution
|
||
// approaches uniform (PER degenerates into vanilla replay); at 1.0
|
||
// sampling is perfectly proportional to priority.
|
||
|
||
#define RL_PER_ALPHA_INDEX 405
|
||
#define PER_ALPHA_MIN 0.3f
|
||
#define PER_ALPHA_MAX 1.0f
|
||
// 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
|
||
// Width of the kurtosis → α-lift logistic map. At kurt=10 (heavy-tailed
|
||
// market data), the lift fully spans 0.2 of the [PER_ALPHA_MIN,
|
||
// PER_ALPHA_MAX] band. Beyond ~30 the mapping saturates at α_max.
|
||
#define KURT_LIFT_SCALE 7.0f
|
||
// 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
|
||
// after the streaming estimator initialises, a sub-Gaussian kurtosis
|
||
// reading (e.g. 0.5 — near-uniform per-step batch means before tails
|
||
// develop) would yield target = 0.4 = PER_ALPHA_MIN, dragging α toward
|
||
// 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
|
||
#define WIENER_ALPHA_FLOOR 0.4f
|
||
|
||
// ─────────────────────────────────────────────────────────────────────
|
||
// rl_per_alpha_controller:
|
||
// Single-thread controller — writes ONE float to
|
||
// isv[RL_PER_ALPHA_INDEX].
|
||
//
|
||
// Inputs:
|
||
// isv [≥ RL_PER_ALPHA_INDEX+1] — ISV bus
|
||
// alpha_step — Wiener-α from the controller's signal stats;
|
||
// floored at WIENER_ALPHA_FLOOR before the blend.
|
||
// Named `alpha_step` to disambiguate from the
|
||
// output (PER's α).
|
||
// td_kurtosis_ema — caller-computed kurtosis EMA of the per-
|
||
// transition |TD-error| series. Phase F produces
|
||
// this via a 4-moment reduce over the replay
|
||
// priority column.
|
||
//
|
||
// Outputs:
|
||
// isv[RL_PER_ALPHA_INDEX] — PER priority exponent α ∈
|
||
// [PER_ALPHA_MIN, PER_ALPHA_MAX]
|
||
// ─────────────────────────────────────────────────────────────────────
|
||
// Phase R5: scalar input arg replaced with `input_slot` ISV index so the
|
||
// EMA producer (Phase R3 ema_update_on_done targeting
|
||
// ISV[RL_TD_KURTOSIS_EMA_INDEX=422]) feeds this controller without any
|
||
// host roundtrip per `feedback_cpu_is_read_only`.
|
||
extern "C" __global__ void rl_per_alpha_controller(
|
||
float* __restrict__ isv,
|
||
float alpha_step,
|
||
int input_slot
|
||
) {
|
||
if (threadIdx.x != 0 || blockIdx.x != 0) return;
|
||
|
||
const float prev = isv[RL_PER_ALPHA_INDEX];
|
||
|
||
// Noise-floor gate: kurtosis below KURT_NOISE_FLOOR is dominated
|
||
// by streaming-estimator startup noise (per-step batch-mean
|
||
// 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) {
|
||
// 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
|
||
// first-non-zero observation still seeds correctly.)
|
||
if (prev != 0.0f) return;
|
||
}
|
||
|
||
// Compute target from the current input EMA via a piecewise linear
|
||
// lift. Shared between bootstrap and per-step paths so the dead-zone
|
||
// coincidence with a hardcoded bootstrap value cannot recur.
|
||
// kurt ≤ KURT_GAUSSIAN → target = 0.4 (PER_ALPHA_MIN + 0.1)
|
||
// kurt = KURT_GAUSSIAN + KURT_LIFT_SCALE (≈10) → target = 0.6 (canonical PER)
|
||
// kurt → ∞ → target → PER_ALPHA_MAX
|
||
//
|
||
// 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);
|
||
float target = 0.4f + 0.2f * (kurt_excess / KURT_LIFT_SCALE);
|
||
target = fmaxf(PER_ALPHA_MIN, fminf(target, PER_ALPHA_MAX));
|
||
|
||
// Bootstrap on sentinel 0.0 per pearl_first_observation_bootstrap:
|
||
// first emit replaces directly with the computed target. At
|
||
// construction time when input EMA is also sentinel-zero, target =
|
||
// 0.4 (min), which is distinct from EVERY non-sentinel target
|
||
// value the formula can emit (≥ 0.4). The next per-step call with
|
||
// any real input will see prev=0.4 vs target≥0.4 and Wiener-blend
|
||
// away from the sentinel.
|
||
if (prev == 0.0f) {
|
||
isv[RL_PER_ALPHA_INDEX] = target;
|
||
return;
|
||
}
|
||
|
||
// Wiener-α blend with floor per pearl_wiener_alpha_floor_for_nonstationary.
|
||
const float a = fmaxf(alpha_step, WIENER_ALPHA_FLOOR);
|
||
float out = (1.0f - a) * prev + a * target;
|
||
|
||
out = fmaxf(PER_ALPHA_MIN, fminf(out, PER_ALPHA_MAX));
|
||
isv[RL_PER_ALPHA_INDEX] = out;
|
||
}
|