diff --git a/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu b/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu index dd43cade7..d9bfe7daa 100644 --- a/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu +++ b/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu @@ -224,51 +224,51 @@ extern "C" __global__ void rl_reward_clamp_controller( isv[RL_REWARD_CLAMP_RATIO_INDEX] = ratio_clamped; } - // Step 4 (G.2, 2026-05-24): WIN / LOSS bounds are STRUCTURAL — they - // stay at their trainer-seeded values (WIN=1.0, LOSS=3.0) matching - // the C51 atom span × 3:1 loss-aversion asymmetry. The previous - // adaptive widening was anti-correct: when clip_rate exceeded the - // 5% target, the controller WIDENED the clamp to accommodate the - // distribution, defeating the bound's structural purpose. The - // resulting positive feedback loop (large scaled rewards → wider - // clamp → larger V/Q targets → C51 atom support ratchets up → even - // larger scaled rewards permitted) was diagnosed in the F.5 200-step - // local smoke: WIN drifted 1.0 → 41.3 (41×) in 150 steps, scaled - // rewards of 14.71 flowed through, V regression spiked to l_v=104. - // - // The diagnostic-only state (pos_max_ema, neg_max_ema, RATIO, - // clip_rate_ema, MARGIN) is still maintained above so the diag can - // surface what the system *would* adapt to under an unbounded - // policy — useful for understanding reward-distribution drift. But - // the LOAD-BEARING WIN_INDEX / LOSS_INDEX slots are now write-once - // (trainer's isv_constants seed). Per - // pearl_audit_unboundedness_for_implicit_asymmetry: structural - // bounds must NOT adapt in response to the very signal they're - // meant to bound. - // - // Suppress unused-variable warnings on the diagnostic-only blocks. + // Step 4 (G.2): WIN / LOSS clamp bounds remain STRUCTURAL — the + // G.2 feedback loop required BOTH clamp widening AND span widening. + // With clamp frozen here, only the atom span adapts (Step 5 below). + // Suppress unused warnings on diagnostic-only state. (void) margin; - (void) ema_new; - // ── C51 atom span — STRUCTURAL (G.2, 2026-05-24). ───────────── + // ── Step 5: C51 atom span adaptation from observed reward EMAs. ── // - // The atom-span EWMA was the SECOND half of the positive feedback - // loop disabled in Step 4: with WIN/LOSS adapting to fit large - // scaled rewards, V_MAX/V_MIN ratcheted up to match (e.g., F.5 - // smoke saw V_MAX=2.66 by step 149 from a seed of 1.0). Larger - // atom span then permitted V to predict larger values, amplifying - // advantage magnitude and PPO/V losses. + // G.2 disabled this because atom-span growth + clamp-bound growth + // created a positive feedback loop. With clamp bounds FROZEN (Step + // 4), the loop can't form — the clamp caps the reward magnitude + // regardless of atom span. // - // Fix: the C51 atom span stays at its trainer-seeded value - // (V_MAX=1.0, V_MIN=-1.0). Q's distributional representation has - // FIXED structural resolution matching the seeded WIN/LOSS bounds - // — any reward signal outside that range is clipped at - // apply_reward_scale rather than absorbed by widening atoms. - // Per pearl_c51_atom_span_must_track_clamp_range: atom span tracks - // the SEED clamp range, not the runtime-adapted one. + // The span slowly tracks the observed reward tail EMAs so Q's + // distributional resolution covers the ACTUAL reward range. Without + // this, rewards exceeding the static span project to edge atoms and + // Q loses magnitude discrimination (wr plateaus at ~0.46). // - // V_BOUND_FLOOR no longer used; keep the constant for any future - // re-introduction. - (void) V_BOUND_FLOOR; - (void) V_BOUND_EWMA_ALPHA; + // Floor at V_BOUND_FLOOR (1.0) preserves baseline resolution. + // Ceiling at clamp bounds (WIN=1.0, LOSS=3.0) prevents runaway. + // Slow EWMA (α=0.001, half-life ~700 steps) lets Q's atom mapping + // adapt gradually without whipsawing. + // Atom span anchors on the CLAMP bounds — the post-clamp reward + // range that Q's Bellman targets actually see. Pre-clamp EMAs + // (pos_ema ≈ 50) would blow atoms to [-50, +50] with Δ_z=5, + // destroying resolution. The clamp bounds (WIN, LOSS) are the + // structural ceiling on what rewards enter Q. + const float win_bound = isv[RL_REWARD_CLAMP_WIN_INDEX]; // 1.0 + const float loss_bound = isv[RL_REWARD_CLAMP_LOSS_INDEX]; // 3.0 + { + const float v_max_prev = isv[RL_C51_V_MAX_INDEX]; + const float v_max_target = fmaxf(V_BOUND_FLOOR, win_bound); + const float v_max_new = (v_max_prev == 0.0f) + ? v_max_target + : (1.0f - V_BOUND_EWMA_ALPHA) * v_max_prev + + V_BOUND_EWMA_ALPHA * v_max_target; + isv[RL_C51_V_MAX_INDEX] = fmaxf(V_BOUND_FLOOR, v_max_new); + } + { + const float v_min_prev = isv[RL_C51_V_MIN_INDEX]; + const float v_min_target = fminf(-V_BOUND_FLOOR, -loss_bound); + const float v_min_new = (v_min_prev == 0.0f) + ? v_min_target + : (1.0f - V_BOUND_EWMA_ALPHA) * v_min_prev + + V_BOUND_EWMA_ALPHA * v_min_target; + isv[RL_C51_V_MIN_INDEX] = fminf(-V_BOUND_FLOOR, v_min_new); + } }