95dcc4e312
cvf86 controller_branch diag (commit 708c121f2) revealed:
rollout_steps: 99.99% WIDEN, 0% HOLD, 0% SHRINK
The bounded-step + noise-floor fix was correctly applied, but the
controller WIDENED 99.99% of steps because the hardcoded
ADV_VAR_RATIO_TARGET = 0.1 (`#define` in the kernel) was a b_size>1
design choice. The streaming-EMA regime at b_size=1 has var/|mean|
naturally living in [1, 10] (median 3.9), so input is ALWAYS >> 0.1
and the controller correctly says "noisy advantages → widen". Result:
n_rollout pegs at MAX=8192 within ~5 steps and stays for 50k steps,
PPO update frequency drops 4×, KL stays in numerical noise (median
1.7e-8), Q can't learn (l_q stuck at ~2.7 vs uniform 3.04).
## Fix: ISV-driven target
Per `feedback_isv_for_adaptive_bounds`: ADV_VAR_RATIO_TARGET now
lives in ISV slot 449 (`RL_ADV_VAR_RATIO_TARGET_INDEX`), seeded at
trainer init to 5.0 (matches streaming-regime median 3.9). The
controller reads `isv[RL_ADV_VAR_RATIO_TARGET_INDEX]` each step
instead of a `#define`.
Expected behavior at TARGET=5.0:
* Median input 3.9 lands in-band [3.33, 7.5] → HOLD
* n_rollout stays near BOOTSTRAP=2048 instead of MAX
* 4× more PPO updates per step → policy actually moves
* KL leaves noise floor → ε controller activates
* Q has gradient signal → can learn
Noise floor is now derived multiplicatively from the ISV target
(`target × ADV_VAR_RATIO_NOISE_FLOOR_FRAC = 0.01`) so adjusting
the target proportionally adjusts the floor — no separate slot
needed.
## Wiring
`rl_streaming_clamp_init.cu` extended to seed all three ISV-resident
design constants (adv_var clamp ceiling, td_kurt clamp ceiling, AND
adv_var regression target). Single kernel call at trainer init —
still no HtoD per `feedback_no_htod_htoh_only_mapped_pinned`.
## Diag bake-in
`controller_branch.rollout_steps_target` now reads from
`isv[RL_ADV_VAR_RATIO_TARGET_INDEX]` instead of the prior hardcoded
`0.1f32` literal. The diag shows the current ISV-resident target
so post-hoc branch analysis uses the actual value the controller
saw, and lets us track whether a future adaptive controller (one
that maintains target from observed-input percentile EMA) is
moving the target correctly.
## Slot allocation
RL_SLOTS_END: 449 → 450 (one new design-constant slot).
## Test updates
G1 (isv_bootstrap) + G3 (r5_controllers) skip slot 449 in the
sentinel-zero loop and assert the seeded value (5.0) separately.
G3's `advantage_var_ratio` input bumped from 5.0 → 20.0 so the
WIDEN branch still fires (input > new target × 1.5 = 7.5) and the
test still validates that the controller moves off bootstrap.
## Verified gates (local sm_86)
G1 isv_bootstrap ✅
G3 controllers ✅ (with updated input)
G4 target_update ✅
integrated_smoke ✅
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
42 lines
2.0 KiB
Plaintext
42 lines
2.0 KiB
Plaintext
// rl_streaming_clamp_init.cu — single-thread device-side seeder for
|
|
// the streaming-kernel output clamp ceilings AND the
|
|
// rl_rollout_steps_controller regression target:
|
|
// ISV[RL_ADV_VAR_RATIO_CLAMP_INDEX = 447] — var/|mean| clamp
|
|
// ISV[RL_TD_KURTOSIS_CLAMP_INDEX = 448] — kurtosis clamp
|
|
// ISV[RL_ADV_VAR_RATIO_TARGET_INDEX = 449] — rollout-controller target
|
|
//
|
|
// Launched once from `with_controllers_bootstrapped` after ISV is
|
|
// alloc_zeros'd. Writes design-constant values directly on device —
|
|
// no host→device transfer, in line with
|
|
// `feedback_no_htod_htoh_only_mapped_pinned`.
|
|
//
|
|
// The streaming variance and kurtosis kernels read the clamp slots
|
|
// each step and bound their output before writing to the consumer
|
|
// controllers' EMA-input slots. Without the clamps streaming
|
|
// `var/|mean|` reached 3e5 in alpha-rl-gxhr8 fold0 (when streaming
|
|
// mean passed through zero — `var / max(|μ|, 1e-6)` blows up);
|
|
// streaming kurtosis reached 50.6.
|
|
//
|
|
// The rollout-controller target slot replaces the prior hardcoded
|
|
// `ADV_VAR_RATIO_TARGET = 0.1` `#define` — that value was chosen for
|
|
// batched (b_size>1) reductions where var/|mean| naturally sits in
|
|
// [0.01, 0.2], but the streaming regime lives in [1, 10] so 0.1
|
|
// triggers WIDEN on 99.99% of steps (alpha-rl-cvf86 fold0). The
|
|
// new ISV-driven target defaults to 5.0 (matches streaming median).
|
|
//
|
|
// Per `feedback_isv_for_adaptive_bounds`: all three values live in
|
|
// ISV (visible in diag, modifiable at runtime by re-launching this
|
|
// kernel) rather than as kernel-side `#define`s.
|
|
|
|
extern "C" __global__ void rl_streaming_clamp_init(
|
|
float* __restrict__ isv,
|
|
int adv_var_clamp_slot, float adv_var_clamp_val,
|
|
int td_kurt_clamp_slot, float td_kurt_clamp_val,
|
|
int adv_var_target_slot, float adv_var_target_val
|
|
) {
|
|
if (threadIdx.x != 0 || blockIdx.x != 0) return;
|
|
isv[adv_var_clamp_slot] = adv_var_clamp_val;
|
|
isv[td_kurt_clamp_slot] = td_kurt_clamp_val;
|
|
isv[adv_var_target_slot] = adv_var_target_val;
|
|
}
|