e87e0b0774
Two architectural fixes from rljzl in-flight analysis (ultrathink
deep dive on actions a7/a8 + per-action calibration):
(1) λ_distill: static → controller-driven via Schulman bounded step
wwcsz showed Q→π KL EMA dropped 2.10 → 0.30 with λ=0.01, then
rljzl bumped to 0.05. Static λ is design intuition; KL is the
natural feedback signal:
if KL > target × 1.5 → λ *= 1.2 (Q not landing, pull harder)
if KL < target / 1.5 → λ /= 1.2 (Q absorbed, relax)
Bounds [MIN=0.001, MAX=1.0]. Target KL seeded 0.1 (slot 491).
New kernel `rl_q_distill_lambda_controller.cu`. Runs after the
distill kernel writes KL_EMA each step.
(2) REWARD_SCALE_MIN: hardcoded 1e-3 → ISV-driven 1e-4
wwcsz audit (mean_abs_pnl_ema mean=920, max=49437, p99=high):
the controller wanted scale ≈ 3.5e-4 when EMA spiked to 2871
but pegged at 1e-3, letting scaled rewards exceed unit support
and wasting C51 atom resolution on outliers. ISV slot 492
permits runtime re-tuning; default 1e-4 admits one more order
of magnitude before pegging. Per user-stated "floors and clamp
bounds" exemption — ISV-resident for tunability, not because
required.
Diag exposes q_distill_kl_target + reward_scale_min so the new
adaptation chains are observable.
Investigation (ultrathink): actions 7/8 (TrailTighten/TrailLoosen)
have ZERO consumers across the codebase. Spec'd as "ISV mutation"
in actions_to_market_targets.cu header but no slot, no mutation
kernel, no stop-check kernel. ~10% of wwcsz policy mass goes to
dead no-ops. Documented in
`pearl_dead_trail_stop_actions_a7_a8.md` — implementation
deferred to its own SP (per-batch trail_distance buffer +
mutation kernel + stop-check integration with LobSim apply_fill_to_pos
per `pearl-stop-checks-run-at-deadline-cadence`). N_ACTIONS=9
preserved; alternative refactor to 7 actions captured as
"Path B" in the pearl.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
68 lines
2.8 KiB
Plaintext
68 lines
2.8 KiB
Plaintext
// rl_q_distill_lambda_controller.cu — adaptive λ_distill controller.
|
||
//
|
||
// Audit 2026-05-24 (rljzl followup design): λ_distill (slot 486) was
|
||
// previously seeded statically (0.01 in wwcsz, 0.05 in rljzl). Static
|
||
// λ leaves Q's pull strength at design intuition — too low and Q's
|
||
// preferences don't reach π; too high and PPO surrogate gets drowned.
|
||
// Make λ adaptive driven by the natural feedback signal:
|
||
// KL(π_target || π_new) EMA at slot 488.
|
||
//
|
||
// Controller logic (Schulman bounded step):
|
||
//
|
||
// target_kl = isv[491] // RL_Q_DISTILL_KL_TARGET (seed 0.1)
|
||
// observed = isv[488] // RL_Q_DISTILL_KL_EMA
|
||
// upper = target × TOLERANCE // 0.15 at default
|
||
// lower = target / TOLERANCE // 0.067 at default
|
||
//
|
||
// if observed > upper → λ *= ADJUST_RATE (Q signal not landing,
|
||
// increase pull)
|
||
// if observed < lower → λ /= ADJUST_RATE (Q absorbed, can relax)
|
||
// else → λ unchanged (dead-zone)
|
||
//
|
||
// Bounds: λ ∈ [MIN_LAMBDA=0.001, MAX_LAMBDA=1.0]. MIN preserves a
|
||
// minimal Q-pull always-on; MAX caps the surrogate dominance (any
|
||
// higher and distill swamps PPO).
|
||
//
|
||
// Bootstrap: on sentinel ema=0 or λ=0, defer adaptation (let static
|
||
// seed take effect). Once both have non-zero values the controller
|
||
// starts adjusting.
|
||
//
|
||
// Per `feedback_no_atomicadd`: single-thread kernel, no atomics.
|
||
// Per `feedback_cpu_is_read_only`: all state in ISV.
|
||
|
||
#define RL_Q_DISTILL_LAMBDA_INDEX 486
|
||
#define RL_Q_DISTILL_KL_EMA_INDEX 488
|
||
#define RL_Q_DISTILL_KL_TARGET_INDEX 491
|
||
|
||
#define MIN_LAMBDA 0.001f
|
||
#define MAX_LAMBDA 1.0f
|
||
#define KL_TOLERANCE 1.5f
|
||
#define LAMBDA_ADJUST_RATE 1.2f
|
||
|
||
extern "C" __global__ void rl_q_distill_lambda_controller(
|
||
float* __restrict__ isv
|
||
) {
|
||
if (threadIdx.x != 0 || blockIdx.x != 0) return;
|
||
|
||
const float kl_observed = isv[RL_Q_DISTILL_KL_EMA_INDEX];
|
||
const float kl_target = isv[RL_Q_DISTILL_KL_TARGET_INDEX];
|
||
float lambda = isv[RL_Q_DISTILL_LAMBDA_INDEX];
|
||
|
||
// Defer if either input slot is sentinel-zero.
|
||
if (kl_observed <= 0.0f || kl_target <= 0.0f || lambda <= 0.0f) return;
|
||
|
||
const float upper = kl_target * KL_TOLERANCE;
|
||
const float lower = kl_target / KL_TOLERANCE;
|
||
|
||
if (kl_observed > upper) {
|
||
// Q signal not landing strongly — raise λ to pull π harder.
|
||
lambda = fminf(MAX_LAMBDA, lambda * LAMBDA_ADJUST_RATE);
|
||
} else if (kl_observed < lower) {
|
||
// π already aligned with Q — can ease off (preserve PPO autonomy).
|
||
lambda = fmaxf(MIN_LAMBDA, lambda / LAMBDA_ADJUST_RATE);
|
||
}
|
||
// Dead-zone: kl_observed in [lower, upper] → no change.
|
||
|
||
isv[RL_Q_DISTILL_LAMBDA_INDEX] = lambda;
|
||
}
|