diff --git a/crates/ml-alpha/build.rs b/crates/ml-alpha/build.rs index 5be249f05..9e911c1c9 100644 --- a/crates/ml-alpha/build.rs +++ b/crates/ml-alpha/build.rs @@ -103,6 +103,7 @@ const KERNELS: &[&str] = &[ "rl_v_blend", // Phase 4.4 (2026-05-30): elementwise V_used = α V_scalar + (1−α) V_dq, α from ISV "rl_v_blend_alpha_controller", // Phase 4.4: ISV-adaptive Schulman-bounded controller on α from observed |V_dq − V_scalar| / |V_scalar| tracking ratio "rl_advantage_normalize", // Phase 4.5 (2026-05-30): per-batch advantage normalization (A − mean)/std with ε² variance floor — standard PPO practice, self-adaptive (no tuned params) + "rl_signal_variance_update", // Adaptive controller floors (2026-05-30): Welford online variance for controller input EMAs; replaces hardcoded NOISE_FLOOR_FRAC constants in 8 controllers — sample_var = M² / (count-1); see spec 2026-05-30-adaptive-controller-floor-design "rl_ensemble_action_value", // C51+IQN ensemble: E_ensemble = α×E_C51 + (1-α)×E_IQN; α from ISV[544] "rl_noisy_linear_forward", // NoisyNet: factored noisy linear forward — y = (mu_w + sigma_w ⊙ eps_w) × x + (mu_b + sigma_b ⊙ eps_b); state-dependent exploration for C51/IQN final projection "rl_noisy_linear_backward", // NoisyNet: factored noisy linear backward — grad_mu_w/sigma_w/mu_b/sigma_b per-batch scratch for reduce_axis0 diff --git a/crates/ml-alpha/cuda/rl_advantage_normalize.cu b/crates/ml-alpha/cuda/rl_advantage_normalize.cu index 6f497b304..b8c057711 100644 --- a/crates/ml-alpha/cuda/rl_advantage_normalize.cu +++ b/crates/ml-alpha/cuda/rl_advantage_normalize.cu @@ -30,8 +30,17 @@ #define BLOCK_X 1024 #define ADVANTAGE_VAR_FLOOR 1e-6f +// Adaptive controller floors (spec 2026-05-30-adaptive-controller-floor-design): +// emit raw per-batch advantage variance to ISV[612] BEFORE the in-place +// normalize. `rl_rollout_steps_controller` consumes this as its driving +// signal — post-normalization variance is definitionally ~ε² floor under +// Phase 4.5 and useless as a controller input. Pre-norm variance is the +// real signal of "how much advantage spread the policy is producing". +#define RL_ADV_VAR_PRE_NORM_INDEX 612 + extern "C" __global__ void rl_advantage_normalize( float* __restrict__ advantage, // [B] in-place + float* __restrict__ isv, // ISV bus — only ADV_VAR_PRE_NORM written int B ) { const int tid = threadIdx.x; @@ -72,6 +81,10 @@ extern "C" __global__ void rl_advantage_normalize( __shared__ float s_inv_std; if (tid == 0) { const float var = s_var[0] / (float)B; + // Emit raw pre-normalization variance for rl_rollout_steps_controller. + // Done BEFORE we add the ε² floor — the controller wants the true + // signal, not the numerical-stability-adjusted divisor. + isv[RL_ADV_VAR_PRE_NORM_INDEX] = var; s_inv_std = rsqrtf(var + ADVANTAGE_VAR_FLOOR); // 1/sqrt(var + ε²) } __syncthreads(); diff --git a/crates/ml-alpha/cuda/rl_entropy_coef_controller.cu b/crates/ml-alpha/cuda/rl_entropy_coef_controller.cu index 62fe52073..15205f407 100644 --- a/crates/ml-alpha/cuda/rl_entropy_coef_controller.cu +++ b/crates/ml-alpha/cuda/rl_entropy_coef_controller.cu @@ -46,13 +46,35 @@ #define RL_ENTROPY_COEF_INDEX 403 #define N_ACTIONS 11 -#define COEF_MIN 0.0f -#define COEF_MAX 0.5f +// COEF MIN/MAX clamp bounds are now ISV-driven per the 2026-05-30 +// clamp-bound extension. Runtime-tunable + visible in diag. +#define RL_ENTROPY_COEF_MIN_INDEX 645 +#define RL_ENTROPY_COEF_MAX_INDEX 646 // ISV-driven entropy-target fraction per `feedback_isv_for_adaptive_bounds`. // Default 0.7 (70% of ln(N_ACTIONS) = "explore but not too randomly"). // Seeded by rl_isv_write at trainer init. #define RL_ENTROPY_TARGET_FRAC_INDEX 458 -#define WIENER_ALPHA_FLOOR 0.4f +// Wiener-α floor — shared across 9 controllers (slot 659). +#define RL_WIENER_ALPHA_FLOOR_INDEX 659 + +// Adaptive controller floors (spec 2026-05-30-adaptive-controller-floor-design): +// noise floor derives from observed entropy_observed_ema variance via +// Welford triples. Replaces the hardcoded `0.5f` emergency-bypass fraction +// with `max(h_target × 0.5, h_target − 2σ)` — under Phase 4.5 the entropy +// distribution narrows around the SAC target and the hardcoded 0.5×h_target +// emergency gate never fires, leaving the Wiener blend to drag coef to MIN +// (alpha-rl-... fold 0 confirmed entropy_coef stuck at 0.01). +// Asymmetric Schulman semantics: coef RAISES on a single below-target +// observation (entropy collapse = safety signal, act fast); coef DROPS +// only after N consecutive above-target observations (healthy entropy = +// drift coef down patiently). +#define RL_ENTROPY_OBS_VAR_COUNT_INDEX 600 +#define RL_ENTROPY_OBS_VAR_M2_INDEX 602 +#define RL_ENTROPY_OBS_BELOW_COUNT_INDEX 603 +#define RL_SCHULMAN_TOLERANCE_INDEX 468 +#define NOISE_FLOOR_TARGET_FRAC 0.5f // floor ≥ 50% of target +#define NOISE_FLOOR_STD_MULTIPLIER 2.0f // floor ≥ 2σ of observed signal +#define WIDEN_PATIENCE_CONSECUTIVE 3.0f // descent requires N above-band steps // ───────────────────────────────────────────────────────────────────── @@ -101,10 +123,12 @@ extern "C" __global__ void rl_entropy_coef_controller( // At h_obs = 0 (total collapse): coef = COEF_MAX (emergency) // At h_obs > h_target: coef = COEF_FLOOR (no penalty for exploring) const float COEF_FLOOR = 0.01f; + const float coef_min = isv[RL_ENTROPY_COEF_MIN_INDEX]; + const float coef_max = isv[RL_ENTROPY_COEF_MAX_INDEX]; const float deficit_frac = fmaxf(0.0f, fminf((h_target - entropy_observed_ema) / fmaxf(h_target, 1e-6f), 1.0f)); - float coef_target = COEF_FLOOR + (COEF_MAX - COEF_FLOOR) * deficit_frac; - coef_target = fmaxf(COEF_MIN, fminf(coef_target, COEF_MAX)); + float coef_target = COEF_FLOOR + (coef_max - COEF_FLOOR) * deficit_frac; + coef_target = fmaxf(coef_min, fminf(coef_target, coef_max)); // Bootstrap on sentinel 0.0 per pearl_first_observation_bootstrap: // first emit replaces directly with the computed target. At cold @@ -117,17 +141,56 @@ extern "C" __global__ void rl_entropy_coef_controller( return; } - // Emergency: bypass Wiener blend when entropy is below 50% of target. - // The slow blend can't react fast enough to prevent Hold collapse. - if (entropy_observed_ema < h_target * 0.5f && entropy_observed_ema > 0.0f) { + // Adaptive noise floor — signal-driven per + // `pearl_zscore_normalization_for_magnitude_asymmetric_signals` and + // `feedback_adaptive_not_tuned`. Replaces hardcoded `h_target × 0.5` + // emergency-bypass threshold with adaptive `max(h_target × 0.5, + // h_target − 2σ_observed)` so the emergency path still fires when + // entropy drops more than 2σ below its observed mean even if the + // distribution is so narrow that 50% × h_target is unreachable. + // Welford sample variance = M² / (count − 1) when count > 1. + const float h_count = isv[RL_ENTROPY_OBS_VAR_COUNT_INDEX]; + const float h_var = (h_count > 1.0f) + ? isv[RL_ENTROPY_OBS_VAR_M2_INDEX] / (h_count - 1.0f) + : 0.0f; + const float h_std = sqrtf(h_var); + const float emergency_floor = fmaxf(h_target * NOISE_FLOOR_TARGET_FRAC, + h_target - h_std * NOISE_FLOOR_STD_MULTIPLIER); + + // Asymmetric Schulman (spec 2026-05-30): RAISE coef on a single + // below-emergency-floor observation — entropy collapse is a safety + // signal we act on fast. DROP coef only after N consecutive + // above-band observations so a single noisy spike of healthy + // entropy can't drag coef toward MIN. Below-counter is on the + // "healthy entropy" direction here (entropy_observed > h_target × + // tolerance), distinct from the canonical Schulman semantics where + // below-counter means "input below band." + if (entropy_observed_ema > 0.0f && entropy_observed_ema < emergency_floor) { isv[RL_ENTROPY_COEF_INDEX] = coef_target; + isv[RL_ENTROPY_OBS_BELOW_COUNT_INDEX] = 0.0f; return; } - // Wiener-α blend with floor per pearl_wiener_alpha_floor_for_nonstationary. - const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR); - float coef_new = (1.0f - a) * coef_prev + a * coef_target; + const float tolerance = isv[RL_SCHULMAN_TOLERANCE_INDEX]; + const float wiener_a = fmaxf(alpha, isv[RL_WIENER_ALPHA_FLOOR_INDEX]); + float coef_new = (1.0f - wiener_a) * coef_prev + wiener_a * coef_target; + coef_new = fmaxf(coef_min, fminf(coef_new, coef_max)); + + if (coef_new < coef_prev && entropy_observed_ema > h_target * tolerance) { + // Healthy-entropy direction — require N consecutive observations + // before letting the Wiener blend drag coef downward. Single noisy + // high-entropy spike must not erode the entropy bonus. + const float new_count = isv[RL_ENTROPY_OBS_BELOW_COUNT_INDEX] + 1.0f; + isv[RL_ENTROPY_OBS_BELOW_COUNT_INDEX] = new_count; + if (new_count < WIDEN_PATIENCE_CONSECUTIVE) { + // Hold coef at previous value until patience accumulates. + isv[RL_ENTROPY_COEF_INDEX] = coef_prev; + return; + } + } else { + // Either entropy unhealthy (coef rising) or in-band — reset patience. + isv[RL_ENTROPY_OBS_BELOW_COUNT_INDEX] = 0.0f; + } - coef_new = fmaxf(COEF_MIN, fminf(coef_new, COEF_MAX)); isv[RL_ENTROPY_COEF_INDEX] = coef_new; } diff --git a/crates/ml-alpha/cuda/rl_fused_controllers.cu b/crates/ml-alpha/cuda/rl_fused_controllers.cu index ee05695e1..d8b81f816 100644 --- a/crates/ml-alpha/cuda/rl_fused_controllers.cu +++ b/crates/ml-alpha/cuda/rl_fused_controllers.cu @@ -2,15 +2,20 @@ // // Eliminates 9 kernel-launch overheads (~40-80μs/step) by running all // controllers sequentially in one thread. Each controller's logic is -// copied verbatim from its standalone .cu file. The standalone files -// are retained for documentation and individual testing. +// MIRRORED from its standalone .cu file (post-2026-05-30 adaptive +// controller-floor refactor): adaptive noise floor derived from Welford +// variance, asymmetric Schulman patience counters, ISV-driven clamp +// bounds + Wiener-α floor. Per `feedback_no_partial_refactor` and the +// "fused kernel must be semantically equivalent to running all individual +// controllers in sequence" contract, every branch below matches its +// standalone counterpart's ISV reads, math, and counter semantics. // // Controllers fused (in execution order): // 1. rl_gamma_controller ISV[400] ← ISV[input_slots[0]] // 2. rl_target_tau_controller ISV[401] ← ISV[input_slots[1]] // 3. rl_ppo_clip_controller ISV[402] ← ISV[input_slots[2]] // 4. rl_entropy_coef_controller ISV[403] ← ISV[input_slots[3]] -// 5. rl_rollout_steps_controller ISV[404] ← ISV[input_slots[4]] +// 5. rl_rollout_steps_controller ISV[404] ← ISV[input_slots[4]] (caller passes ADV_VAR_PRE_NORM=612) // 6. rl_per_alpha_controller ISV[405] ← ISV[input_slots[5]] // 7. rl_reward_scale_controller ISV[406] ← ISV[input_slots[6]] // 8. rl_ppo_ratio_clamp_controller ISV[440] ← ISV[402] (reads ε just written) @@ -24,65 +29,104 @@ // ═══════════════════════════════════════════════════════════════════════ // ISV slot indices — mirrored from individual controller headers. +// All MIN/MAX/threshold/rate constants are ISV-resident per the 2026-05-30 +// adaptive controller-floor design. The few remaining `#define`s below are +// architectural exemptions (physics/math constants, structural safeguards). // ═══════════════════════════════════════════════════════════════════════ -// --- 1. Gamma controller --- +// --- Output slots (controller emits) --- #define RL_GAMMA_INDEX 400 -#define GAMMA_MIN 0.995f -#define GAMMA_MAX 0.999f +#define RL_TARGET_TAU_INDEX 401 +#define RL_PPO_CLIP_INDEX 402 +#define RL_ENTROPY_COEF_INDEX 403 +#define RL_N_ROLLOUT_STEPS_INDEX 404 +#define RL_PER_ALPHA_INDEX 405 +#define RL_REWARD_SCALE_INDEX 406 +#define RL_PPO_RATIO_CLAMP_MAX_INDEX 440 +#define RL_Q_DISTILL_LAMBDA_INDEX 486 + +// --- 1. Gamma controller (Special G) --- +// γ MIN now adaptive via RL_GAMMA_MIN_ADAPTIVE_INDEX (slot 613) derived +// from the Welford mean of trade duration. γ MAX is ISV-resident. +#define RL_GAMMA_MAX_INDEX 649 +#define RL_GAMMA_MIN_ADAPTIVE_INDEX 613 +#define RL_TRADE_DUR_VAR_COUNT_INDEX 608 +#define RL_TRADE_DUR_VAR_MEAN_INDEX 609 +#define RL_MEAN_TRADE_DURATION_EMA_INDEX 417 +#define GAMMA_MIN_ABSOLUTE 0.9f // architectural: don't degenerate to bandit +#define HORIZON_MULTIPLIER 2.0f // architectural: effective horizon = 2× transaction horizon +#define WELFORD_WARMUP_OBS 100.0f // architectural: confidence threshold, not magnitude calibration +#define HORIZON_FLOOR 10.0f // architectural: `1/d` numerical-stability floor // --- 2. Target tau controller --- -#define RL_TARGET_TAU_INDEX 401 -#define TAU_MIN 0.001f +#define RL_TARGET_TAU_MIN_INDEX 642 #define RL_TARGET_TAU_MAX_INDEX 573 #define RL_TAU_BOOTSTRAP_INDEX 473 #define RL_DIV_TARGET_INDEX 457 -#define DIV_NOISE_FLOOR_FRAC 0.01f +#define RL_Q_DIV_VAR_COUNT_INDEX 592 +#define RL_Q_DIV_VAR_M2_INDEX 594 +#define RL_Q_DIV_BELOW_COUNT_INDEX 595 -// --- 3. PPO clip controller --- -#define RL_PPO_CLIP_INDEX 402 -#define EPS_MIN 0.05f -#define EPS_MAX 0.5f +// --- 3. PPO clip controller (canonical pattern) --- +#define RL_PPO_CLIP_EPS_MIN_INDEX 640 +#define RL_PPO_CLIP_EPS_MAX_INDEX 641 #define RL_EPS_BOOTSTRAP_INDEX 474 #define RL_KL_TARGET_INDEX 454 -#define KL_NOISE_FLOOR_FRAC 0.01f +#define RL_KL_PI_VAR_COUNT_INDEX 588 +#define RL_KL_PI_VAR_M2_INDEX 590 +#define RL_KL_PI_BELOW_COUNT_INDEX 591 // --- 4. Entropy coef controller --- -#define RL_ENTROPY_COEF_INDEX 403 -#define N_ACTIONS 11 -#define COEF_MIN 0.0f -#define COEF_MAX 0.5f +#define N_ACTIONS 11 // architectural: matches Q head / action space dim +#define RL_ENTROPY_COEF_MIN_INDEX 645 +#define RL_ENTROPY_COEF_MAX_INDEX 646 #define RL_ENTROPY_TARGET_FRAC_INDEX 458 +#define RL_ENTROPY_OBS_VAR_COUNT_INDEX 600 +#define RL_ENTROPY_OBS_VAR_M2_INDEX 602 +#define RL_ENTROPY_OBS_BELOW_COUNT_INDEX 603 // --- 5. Rollout steps controller --- -#define RL_N_ROLLOUT_STEPS_INDEX 404 -#define ROLLOUT_MIN 256.0f -#define ROLLOUT_MAX 8192.0f +// Input slot is now RL_ADV_VAR_PRE_NORM_INDEX=612 (caller-supplied via +// input_slots[4]) since post-Phase-4.5 advantage variance is definitionally +// near-zero and useless. See `rl_rollout_steps_controller.cu` comments. +#define RL_ROLLOUT_MIN_INDEX 643 +#define RL_ROLLOUT_MAX_INDEX 644 #define RL_ROLLOUT_BOOTSTRAP_INDEX 475 #define RL_ADV_VAR_RATIO_TARGET_INDEX 449 -#define ADV_VAR_RATIO_NOISE_FLOOR_FRAC 0.01f +#define RL_ADV_VAR_VAR_COUNT_INDEX 596 +#define RL_ADV_VAR_VAR_M2_INDEX 598 +#define RL_ADV_VAR_BELOW_COUNT_INDEX 599 // --- 6. PER alpha controller --- -#define RL_PER_ALPHA_INDEX 405 -#define PER_ALPHA_MIN 0.3f -#define PER_ALPHA_MAX 1.0f +#define RL_PER_ALPHA_MIN_INDEX 647 +#define RL_PER_ALPHA_MAX_INDEX 648 #define RL_KURT_GAUSSIAN_INDEX 471 #define RL_KURT_NOISE_FLOOR_INDEX 472 #define RL_KURT_LIFT_SCALE_INDEX 459 +#define RL_TD_KURT_VAR_COUNT_INDEX 604 +#define RL_TD_KURT_VAR_M2_INDEX 606 +#define RL_TD_KURT_BELOW_COUNT_INDEX 607 -// --- 7. Reward scale controller --- -#define RL_REWARD_SCALE_INDEX 406 +// --- 7. Reward scale controller (Special R) --- #define RL_REWARD_SCALE_MIN_INDEX 492 -#define REWARD_SCALE_MAX 1e3f #define RL_REWARD_SCALE_BOOTSTRAP_INDEX 476 -#define EPS_PNL 1e-3f +#define RL_REWARD_MAGNITUDE_EMA_INDEX 614 +#define EPS_PNL 1e-3f // architectural: div-by-zero guard +#define REWARD_SCALE_MAX 1e3f // architectural: 5-order-of-magnitude runaway guard +#define BOOTSTRAP_FRACTION_FLOOR 0.1f // architectural: early-training spiral guard (spec §Special R) +#define MIN_TRADES_FOR_RELEASE 100.0f // architectural: confidence threshold for releasing bootstrap floor +#define ASYM_DECREASE_RATE 1.05f // architectural: asymmetric per-step decrease cap // --- 8. PPO ratio clamp controller --- -#define RL_PPO_RATIO_CLAMP_MAX_INDEX 440 -#define RL_PPO_RATIO_CLAMP_BOOTSTRAP_INDEX 477 -#define RL_PPO_CLAMP_MARGIN_INDEX 460 -#define PPO_RATIO_CLAMP_MIN_OUT 2.0f -#define PPO_RATIO_CLAMP_MAX_OUT 1000.0f +// MIN_OUT now adaptive from observed ratio variance; MAX_OUT is ISV-resident. +#define RL_PPO_RATIO_CLAMP_MAX_ADAPTIVE_INDEX 629 +#define RL_PPO_RATIO_CLAMP_BOOTSTRAP_INDEX 477 +#define RL_PPO_CLAMP_MARGIN_INDEX 460 +#define RL_PPO_RATIO_VAR_COUNT_INDEX 630 +#define RL_PPO_RATIO_VAR_M2_INDEX 632 +#define PPO_RATIO_MIN_ABSOLUTE 2.0f // architectural exemption: vanilla-PG safeguard +#define PPO_RATIO_MIN_STD_SCALE 3.0f // architectural: 3σ healthy-update coverage +#define PPO_RATIO_MAX_OVER_MIN_FACTOR 5.0f // architectural: ensure max ≥ 5× min // --- 9. Gate threshold controller --- #define RL_CONF_GATE_THRESHOLD_INDEX 512 @@ -96,23 +140,32 @@ #define RL_GATE_FRD_MIN_INDEX 529 #define RL_GATE_FRD_MAX_INDEX 530 #define RL_GATE_ADJUST_RATE_INDEX 531 +#define RL_GATE_EMA_ALPHA_INDEX 638 -// --- 10. Q distill lambda controller --- -#define RL_Q_DISTILL_LAMBDA_INDEX 486 +// --- 10. Q distill lambda controller (Special Q) --- #define RL_Q_DISTILL_KL_EMA_INDEX 488 #define RL_Q_DISTILL_KL_TARGET_INDEX 491 -#define MIN_LAMBDA 0.05f -#define MAX_LAMBDA 1.0f -#define KL_TOLERANCE 3.0f -#define LAMBDA_RAMP_RATE 1.2f -#define LAMBDA_DECAY_RATE 0.998f +#define RL_Q_DISTILL_KL_VAR_COUNT_INDEX 618 +#define RL_Q_DISTILL_KL_VAR_M2_INDEX 620 +#define RL_Q_DISTILL_LAMBDA_MIN_ADAPTIVE_INDEX 639 +#define RL_Q_DISTILL_LAMBDA_MAX_INDEX 650 +#define RL_Q_DISTILL_KL_TOLERANCE_INDEX 651 +#define RL_Q_DISTILL_LAMBDA_RAMP_RATE_INDEX 652 +#define RL_Q_DISTILL_LAMBDA_DECAY_RATE_INDEX 653 +#define ADAPTIVE_MIN_ABSOLUTE 0.001f // architectural: q_distill_lambda safety floor +#define ADAPTIVE_MIN_STD_SCALE 0.05f // architectural: signal-noise proportional scale // --- Shared Schulman params --- #define RL_SCHULMAN_TOLERANCE_INDEX 468 #define RL_SCHULMAN_ADJUST_RATE_INDEX 469 -// --- Shared Wiener --- -#define WIENER_ALPHA_FLOOR 0.4f +// --- Shared adaptive-floor constants (spec 2026-05-30) --- +#define NOISE_FLOOR_TARGET_FRAC 0.5f // floor ≥ 50% of target +#define NOISE_FLOOR_STD_MULTIPLIER 2.0f // floor ≥ 2σ of observed signal +#define WIDEN_PATIENCE_CONSECUTIVE 3.0f // widen/descend requires N below-band steps + +// --- Shared Wiener-α floor --- +#define RL_WIENER_ALPHA_FLOOR_INDEX 659 // --- Step counter (device-resident) --- #define RL_STEP_COUNTER_ISV_INDEX 548 @@ -126,7 +179,7 @@ // [1] = RL_Q_DIVERGENCE_EMA_INDEX (418) — target_tau // [2] = RL_KL_PI_EMA_INDEX (419) — ppo_clip // [3] = RL_ENTROPY_OBSERVED_EMA_INDEX (420) — entropy_coef -// [4] = RL_ADVANTAGE_VAR_RATIO_EMA_INDEX (421) — rollout_steps +// [4] = RL_ADV_VAR_PRE_NORM_INDEX (612) — rollout_steps (post-Phase-4.5) // [5] = RL_TD_KURTOSIS_EMA_INDEX (422) — per_alpha // [6] = RL_MEAN_ABS_PNL_EMA_INDEX (423) — reward_scale // ═══════════════════════════════════════════════════════════════════════ @@ -140,29 +193,44 @@ extern "C" __global__ void rl_fused_controllers( if (threadIdx.x != 0 || blockIdx.x != 0) return; const int current_step = (int)isv[RL_STEP_COUNTER_ISV_INDEX]; + const float wiener_floor = isv[RL_WIENER_ALPHA_FLOOR_INDEX]; // ═══════════════════════════════════════════════════════════════════ - // 1. Gamma controller → ISV[400] + // 1. Gamma controller → ISV[400] (Special G) // ═══════════════════════════════════════════════════════════════════ { const float gamma_prev = isv[RL_GAMMA_INDEX]; + + // Adaptive GAMMA_MIN: Welford mean of trade duration after warmup, + // EMA before warmup. See rl_gamma_controller.cu Special case G. + const float d_welford_count = isv[RL_TRADE_DUR_VAR_COUNT_INDEX]; + const float d_welford_mean = isv[RL_TRADE_DUR_VAR_MEAN_INDEX]; + const float d_smoothed = (d_welford_count >= WELFORD_WARMUP_OBS) + ? d_welford_mean + : isv[RL_MEAN_TRADE_DURATION_EMA_INDEX]; + const float adaptive_min = fmaxf(GAMMA_MIN_ABSOLUTE, + 1.0f - 1.0f / fmaxf(d_smoothed * HORIZON_MULTIPLIER, + HORIZON_FLOOR)); + isv[RL_GAMMA_MIN_ADAPTIVE_INDEX] = adaptive_min; + const float mean_trade_duration_events = isv[input_slots[0]]; const float d = fmaxf(mean_trade_duration_events, 1.0f); + const float gamma_max = isv[RL_GAMMA_MAX_INDEX]; float gamma_target = powf(0.5f, 1.0f / d); - gamma_target = fmaxf(GAMMA_MIN, fminf(gamma_target, GAMMA_MAX)); + gamma_target = fmaxf(adaptive_min, fminf(gamma_target, gamma_max)); if (gamma_prev == 0.0f) { isv[RL_GAMMA_INDEX] = gamma_target; } else { - const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR); + const float a = fmaxf(alpha, wiener_floor); float gamma_new = (1.0f - a) * gamma_prev + a * gamma_target; - gamma_new = fmaxf(GAMMA_MIN, fminf(gamma_new, GAMMA_MAX)); + gamma_new = fmaxf(adaptive_min, fminf(gamma_new, gamma_max)); isv[RL_GAMMA_INDEX] = gamma_new; } } // ═══════════════════════════════════════════════════════════════════ - // 2. Target tau controller → ISV[401] + // 2. Target tau controller → ISV[401] (canonical pattern) // ═══════════════════════════════════════════════════════════════════ { const float tau_prev = isv[RL_TARGET_TAU_INDEX]; @@ -172,28 +240,44 @@ extern "C" __global__ void rl_fused_controllers( const float q_divergence_norm = isv[input_slots[1]]; if (q_divergence_norm != 0.0f) { const float div_target = isv[RL_DIV_TARGET_INDEX]; - const float div_noise_floor = div_target * DIV_NOISE_FLOOR_FRAC; + + // Adaptive noise floor — signal-driven via Welford variance. + const float div_count = isv[RL_Q_DIV_VAR_COUNT_INDEX]; + const float div_var = (div_count > 1.0f) + ? isv[RL_Q_DIV_VAR_M2_INDEX] / (div_count - 1.0f) + : 0.0f; + const float div_std = sqrtf(div_var); + const float div_noise_floor = fmaxf(div_target * NOISE_FLOOR_TARGET_FRAC, + div_std * NOISE_FLOOR_STD_MULTIPLIER); + if (q_divergence_norm >= div_noise_floor) { - float ratio; + // Asymmetric Schulman: raise τ on single observation, + // lower τ requires N consecutive below-band observations. const float tolerance = isv[RL_SCHULMAN_TOLERANCE_INDEX]; const float adjust_rate = isv[RL_SCHULMAN_ADJUST_RATE_INDEX]; + float ratio; if (q_divergence_norm > div_target * tolerance) { ratio = adjust_rate; + isv[RL_Q_DIV_BELOW_COUNT_INDEX] = 0.0f; } else if (q_divergence_norm < div_target / tolerance) { - ratio = 1.0f / adjust_rate; + const float new_count = isv[RL_Q_DIV_BELOW_COUNT_INDEX] + 1.0f; + isv[RL_Q_DIV_BELOW_COUNT_INDEX] = new_count; + ratio = (new_count >= WIDEN_PATIENCE_CONSECUTIVE) ? (1.0f / adjust_rate) : 1.0f; } else { + isv[RL_Q_DIV_BELOW_COUNT_INDEX] = 0.0f; ratio = 1.0f; } float tau_target = tau_prev * ratio; - const float tau_max_isv = isv[RL_TARGET_TAU_MAX_INDEX]; - tau_target = fmaxf(TAU_MIN, fminf(tau_target, tau_max_isv)); + const float tau_min = isv[RL_TARGET_TAU_MIN_INDEX]; + const float tau_max = isv[RL_TARGET_TAU_MAX_INDEX]; + tau_target = fmaxf(tau_min, fminf(tau_target, tau_max)); if (tau_prev == isv[RL_TAU_BOOTSTRAP_INDEX]) { isv[RL_TARGET_TAU_INDEX] = tau_target; } else { - const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR); + const float a = fmaxf(alpha, wiener_floor); float tau_new = (1.0f - a) * tau_prev + a * tau_target; - tau_new = fmaxf(TAU_MIN, fminf(tau_new, tau_max_isv)); + tau_new = fmaxf(tau_min, fminf(tau_new, tau_max)); isv[RL_TARGET_TAU_INDEX] = tau_new; } } @@ -202,7 +286,7 @@ extern "C" __global__ void rl_fused_controllers( } // ═══════════════════════════════════════════════════════════════════ - // 3. PPO clip controller → ISV[402] + // 3. PPO clip controller → ISV[402] (canonical pattern — reference) // ═══════════════════════════════════════════════════════════════════ { const float eps_prev = isv[RL_PPO_CLIP_INDEX]; @@ -212,27 +296,44 @@ extern "C" __global__ void rl_fused_controllers( const float kl_ema = isv[input_slots[2]]; if (kl_ema != 0.0f) { const float kl_target = isv[RL_KL_TARGET_INDEX]; - const float kl_noise_floor = kl_target * KL_NOISE_FLOOR_FRAC; + + // Adaptive noise floor — signal-driven via Welford variance. + const float kl_count = isv[RL_KL_PI_VAR_COUNT_INDEX]; + const float kl_var = (kl_count > 1.0f) + ? isv[RL_KL_PI_VAR_M2_INDEX] / (kl_count - 1.0f) + : 0.0f; + const float kl_std = sqrtf(kl_var); + const float kl_noise_floor = fmaxf(kl_target * NOISE_FLOOR_TARGET_FRAC, + kl_std * NOISE_FLOOR_STD_MULTIPLIER); + if (kl_ema >= kl_noise_floor) { + // Asymmetric Schulman: tighten on single observation, + // widen requires N consecutive below-band observations. 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; + isv[RL_KL_PI_BELOW_COUNT_INDEX] = 0.0f; } else if (kl_ema < kl_target / tolerance) { - ratio = adjust_rate; + const float new_count = isv[RL_KL_PI_BELOW_COUNT_INDEX] + 1.0f; + isv[RL_KL_PI_BELOW_COUNT_INDEX] = new_count; + ratio = (new_count >= WIDEN_PATIENCE_CONSECUTIVE) ? adjust_rate : 1.0f; } else { + isv[RL_KL_PI_BELOW_COUNT_INDEX] = 0.0f; ratio = 1.0f; } + const float eps_min = isv[RL_PPO_CLIP_EPS_MIN_INDEX]; + const float eps_max = isv[RL_PPO_CLIP_EPS_MAX_INDEX]; float eps_target = eps_prev * ratio; - eps_target = fmaxf(EPS_MIN, fminf(eps_target, EPS_MAX)); + eps_target = fmaxf(eps_min, fminf(eps_target, eps_max)); if (eps_prev == isv[RL_EPS_BOOTSTRAP_INDEX]) { isv[RL_PPO_CLIP_INDEX] = eps_target; } else { - const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR); + const float a = fmaxf(alpha, wiener_floor); float eps_new = (1.0f - a) * eps_prev + a * eps_target; - eps_new = fmaxf(EPS_MIN, fminf(eps_new, EPS_MAX)); + eps_new = fmaxf(eps_min, fminf(eps_new, eps_max)); isv[RL_PPO_CLIP_INDEX] = eps_new; } } @@ -242,6 +343,8 @@ extern "C" __global__ void rl_fused_controllers( // ═══════════════════════════════════════════════════════════════════ // 4. Entropy coef controller → ISV[403] + // Asymmetric: raise coef on single below-floor observation, + // drop coef requires N consecutive above-band observations. // ═══════════════════════════════════════════════════════════════════ { const float coef_prev = isv[RL_ENTROPY_COEF_INDEX]; @@ -249,26 +352,57 @@ extern "C" __global__ void rl_fused_controllers( const float h_max = logf((float)N_ACTIONS); const float target_frac = isv[RL_ENTROPY_TARGET_FRAC_INDEX]; const float h_target = target_frac * h_max; - const float COEF_FLOOR = 0.01f; + const float COEF_FLOOR = 0.01f; // architectural: maintenance pressure baseline + const float coef_min = isv[RL_ENTROPY_COEF_MIN_INDEX]; + const float coef_max = isv[RL_ENTROPY_COEF_MAX_INDEX]; const float deficit_frac = fmaxf(0.0f, fminf((h_target - entropy_observed_ema) / fmaxf(h_target, 1e-6f), 1.0f)); - float coef_target = COEF_FLOOR + (COEF_MAX - COEF_FLOOR) * deficit_frac; - coef_target = fmaxf(COEF_MIN, fminf(coef_target, COEF_MAX)); + float coef_target = COEF_FLOOR + (coef_max - COEF_FLOOR) * deficit_frac; + coef_target = fmaxf(coef_min, fminf(coef_target, coef_max)); if (coef_prev == 0.0f) { isv[RL_ENTROPY_COEF_INDEX] = coef_target; - } else if (entropy_observed_ema < h_target * 0.5f && entropy_observed_ema > 0.0f) { - isv[RL_ENTROPY_COEF_INDEX] = coef_target; } else { - const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR); - float coef_new = (1.0f - a) * coef_prev + a * coef_target; - coef_new = fmaxf(COEF_MIN, fminf(coef_new, COEF_MAX)); - isv[RL_ENTROPY_COEF_INDEX] = coef_new; + // Adaptive emergency floor (signal-driven via Welford variance). + const float h_count = isv[RL_ENTROPY_OBS_VAR_COUNT_INDEX]; + const float h_var = (h_count > 1.0f) + ? isv[RL_ENTROPY_OBS_VAR_M2_INDEX] / (h_count - 1.0f) + : 0.0f; + const float h_std = sqrtf(h_var); + const float emergency_floor = fmaxf(h_target * NOISE_FLOOR_TARGET_FRAC, + h_target - h_std * NOISE_FLOOR_STD_MULTIPLIER); + + if (entropy_observed_ema > 0.0f && entropy_observed_ema < emergency_floor) { + // Single-observation emergency raise; reset patience. + isv[RL_ENTROPY_COEF_INDEX] = coef_target; + isv[RL_ENTROPY_OBS_BELOW_COUNT_INDEX] = 0.0f; + } else { + const float tolerance = isv[RL_SCHULMAN_TOLERANCE_INDEX]; + const float a = fmaxf(alpha, wiener_floor); + float coef_new = (1.0f - a) * coef_prev + a * coef_target; + coef_new = fmaxf(coef_min, fminf(coef_new, coef_max)); + + if (coef_new < coef_prev && entropy_observed_ema > h_target * tolerance) { + // Healthy-entropy descent direction — require N consecutive observations + // before letting Wiener drag coef downward. + const float new_count = isv[RL_ENTROPY_OBS_BELOW_COUNT_INDEX] + 1.0f; + isv[RL_ENTROPY_OBS_BELOW_COUNT_INDEX] = new_count; + if (new_count < WIDEN_PATIENCE_CONSECUTIVE) { + isv[RL_ENTROPY_COEF_INDEX] = coef_prev; + } else { + isv[RL_ENTROPY_COEF_INDEX] = coef_new; + } + } else { + isv[RL_ENTROPY_OBS_BELOW_COUNT_INDEX] = 0.0f; + isv[RL_ENTROPY_COEF_INDEX] = coef_new; + } + } } } // ═══════════════════════════════════════════════════════════════════ // 5. Rollout steps controller → ISV[404] + // Input now is RL_ADV_VAR_PRE_NORM_INDEX=612 (caller via input_slots[4]). // ═══════════════════════════════════════════════════════════════════ { const float prev = isv[RL_N_ROLLOUT_STEPS_INDEX]; @@ -278,28 +412,44 @@ extern "C" __global__ void rl_fused_controllers( const float advantage_var_over_abs_mean = isv[input_slots[4]]; if (advantage_var_over_abs_mean != 0.0f) { const float adv_var_target = isv[RL_ADV_VAR_RATIO_TARGET_INDEX]; - const float adv_var_noise_floor = - adv_var_target * ADV_VAR_RATIO_NOISE_FLOOR_FRAC; + + // Adaptive noise floor — signal-driven via Welford variance. + const float adv_count = isv[RL_ADV_VAR_VAR_COUNT_INDEX]; + const float adv_var = (adv_count > 1.0f) + ? isv[RL_ADV_VAR_VAR_M2_INDEX] / (adv_count - 1.0f) + : 0.0f; + const float adv_std = sqrtf(adv_var); + const float adv_var_noise_floor = fmaxf(adv_var_target * NOISE_FLOOR_TARGET_FRAC, + adv_std * NOISE_FLOOR_STD_MULTIPLIER); + if (advantage_var_over_abs_mean >= adv_var_noise_floor) { + // Asymmetric Schulman: widen rollout on single above-band observation, + // shrink requires N consecutive below-band observations. const float tolerance = isv[RL_SCHULMAN_TOLERANCE_INDEX]; const float adjust_rate = isv[RL_SCHULMAN_ADJUST_RATE_INDEX]; float scale; if (advantage_var_over_abs_mean > adv_var_target * tolerance) { scale = adjust_rate; + isv[RL_ADV_VAR_BELOW_COUNT_INDEX] = 0.0f; } else if (advantage_var_over_abs_mean < adv_var_target / tolerance) { - scale = 1.0f / adjust_rate; + const float new_count = isv[RL_ADV_VAR_BELOW_COUNT_INDEX] + 1.0f; + isv[RL_ADV_VAR_BELOW_COUNT_INDEX] = new_count; + scale = (new_count >= WIDEN_PATIENCE_CONSECUTIVE) ? (1.0f / adjust_rate) : 1.0f; } else { + isv[RL_ADV_VAR_BELOW_COUNT_INDEX] = 0.0f; scale = 1.0f; } + const float rollout_min = isv[RL_ROLLOUT_MIN_INDEX]; + const float rollout_max = isv[RL_ROLLOUT_MAX_INDEX]; float target = prev * scale; - target = fmaxf(ROLLOUT_MIN, fminf(target, ROLLOUT_MAX)); + target = fmaxf(rollout_min, fminf(target, rollout_max)); if (prev == isv[RL_ROLLOUT_BOOTSTRAP_INDEX]) { isv[RL_N_ROLLOUT_STEPS_INDEX] = target; } else { - const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR); + const float a = fmaxf(alpha, wiener_floor); float out = (1.0f - a) * prev + a * target; - out = fmaxf(ROLLOUT_MIN, fminf(out, ROLLOUT_MAX)); + out = fmaxf(rollout_min, fminf(out, rollout_max)); isv[RL_N_ROLLOUT_STEPS_INDEX] = out; } } @@ -309,35 +459,70 @@ extern "C" __global__ void rl_fused_controllers( // ═══════════════════════════════════════════════════════════════════ // 6. PER alpha controller → ISV[405] + // Asymmetric Schulman patience on descent direction (light tails). // ═══════════════════════════════════════════════════════════════════ { const float prev = isv[RL_PER_ALPHA_INDEX]; const float td_kurtosis_ema = isv[input_slots[5]]; - if (td_kurtosis_ema > 0.0f && td_kurtosis_ema < isv[RL_KURT_NOISE_FLOOR_INDEX]) { + // Adaptive noise floor: combines absolute ISV floor with Welford-derived + // adaptive component (std-scaled). + const float kurt_count = isv[RL_TD_KURT_VAR_COUNT_INDEX]; + const float kurt_var = (kurt_count > 1.0f) + ? isv[RL_TD_KURT_VAR_M2_INDEX] / (kurt_count - 1.0f) + : 0.0f; + const float kurt_std = sqrtf(kurt_var); + const float adaptive_noise_floor = fmaxf(isv[RL_KURT_NOISE_FLOOR_INDEX], + kurt_std * NOISE_FLOOR_STD_MULTIPLIER); + + // Hold-at-prev when signal present but below adaptive noise floor + // (matches standalone semantics: cold-start prev==0 takes the + // bootstrap path below, otherwise return). + if (td_kurtosis_ema > 0.0f && td_kurtosis_ema < adaptive_noise_floor) { if (prev != 0.0f) goto per_alpha_done; } { + const float per_alpha_min = isv[RL_PER_ALPHA_MIN_INDEX]; + const float per_alpha_max = isv[RL_PER_ALPHA_MAX_INDEX]; const float kurt_excess = fmaxf(0.0f, td_kurtosis_ema - isv[RL_KURT_GAUSSIAN_INDEX]); const float kurt_lift_scale = isv[RL_KURT_LIFT_SCALE_INDEX]; float target = 0.4f + 0.2f * (kurt_excess / kurt_lift_scale); - target = fmaxf(PER_ALPHA_MIN, fminf(target, PER_ALPHA_MAX)); + target = fmaxf(per_alpha_min, fminf(target, per_alpha_max)); if (prev == 0.0f) { isv[RL_PER_ALPHA_INDEX] = target; } else { - const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR); + const float a = fmaxf(alpha, wiener_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; + out = fmaxf(per_alpha_min, fminf(out, per_alpha_max)); + + // Asymmetric Schulman patience on the DESCENT direction: + // α rises on a single above-band observation, falls only + // after N consecutive below-band observations. + const float kurt_gaussian = isv[RL_KURT_GAUSSIAN_INDEX]; + const float tolerance = isv[RL_SCHULMAN_TOLERANCE_INDEX]; + if (out < prev && td_kurtosis_ema < kurt_gaussian / tolerance) { + const float new_count = isv[RL_TD_KURT_BELOW_COUNT_INDEX] + 1.0f; + isv[RL_TD_KURT_BELOW_COUNT_INDEX] = new_count; + if (new_count < WIDEN_PATIENCE_CONSECUTIVE) { + isv[RL_PER_ALPHA_INDEX] = prev; + } else { + isv[RL_PER_ALPHA_INDEX] = out; + } + } else { + isv[RL_TD_KURT_BELOW_COUNT_INDEX] = 0.0f; + isv[RL_PER_ALPHA_INDEX] = out; + } } } per_alpha_done:; } // ═══════════════════════════════════════════════════════════════════ - // 7. Reward scale controller → ISV[406] + // 7. Reward scale controller → ISV[406] (Special R) + // Asymmetric per-step decrease cap (5%) + bootstrap-fraction floor + // (10% of bootstrap) until N=100 closed trades. // ═══════════════════════════════════════════════════════════════════ { const float prev = isv[RL_REWARD_SCALE_INDEX]; @@ -346,23 +531,50 @@ extern "C" __global__ void rl_fused_controllers( } else { const float mean_abs_pnl_ema = isv[input_slots[6]]; if (mean_abs_pnl_ema != 0.0f) { + // Emit per-batch pnl magnitude EMA to the public slot for + // downstream Welford-variance kernel consumption. + isv[RL_REWARD_MAGNITUDE_EMA_INDEX] = mean_abs_pnl_ema; + const float denom = fmaxf(mean_abs_pnl_ema, EPS_PNL); float target = 1.0f / denom; const float scale_min = isv[RL_REWARD_SCALE_MIN_INDEX]; target = fmaxf(scale_min, fminf(target, REWARD_SCALE_MAX)); - if (prev == isv[RL_REWARD_SCALE_BOOTSTRAP_INDEX]) { - isv[RL_REWARD_SCALE_INDEX] = target; + // Bootstrap-fraction floor: until N=100 trades have closed, + // scale cannot drop below 10% of bootstrap. The trade count + // comes from the Welford trade-duration counter (incremented + // once per closed trade). + const float trade_count = isv[RL_TRADE_DUR_VAR_COUNT_INDEX]; + const float boot = isv[RL_REWARD_SCALE_BOOTSTRAP_INDEX]; + const float boot_floor = (trade_count < MIN_TRADES_FOR_RELEASE) + ? boot * BOOTSTRAP_FRACTION_FLOOR + : scale_min; + + if (prev == boot) { + // First-observation replace-directly with asymmetric DECREASE + // rate-limit applied (per spec §Special R bullet (a)). + float first = target; + if (first < prev) { + first = fmaxf(first, prev / ASYM_DECREASE_RATE); + } + first = fmaxf(boot_floor, fminf(first, REWARD_SCALE_MAX)); + isv[RL_REWARD_SCALE_INDEX] = first; } else { - const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR); + const float a = fmaxf(alpha, wiener_floor); float out = (1.0f - a) * prev + a * target; - // Per-step ±2% movement clamp. + // Per-step ±2% symmetric movement clamp. const float max_move = prev * 1.02f; const float min_move = prev * 0.98f; out = fmaxf(min_move, fminf(out, max_move)); - out = fmaxf(scale_min, fminf(out, REWARD_SCALE_MAX)); + // Asymmetric DECREASE rate limit (belt-and-braces — documents + // the controller invariant even if 2% symmetric clamp relaxes). + if (out < prev) { + out = fmaxf(out, prev / ASYM_DECREASE_RATE); + } + + out = fmaxf(boot_floor, fminf(out, REWARD_SCALE_MAX)); isv[RL_REWARD_SCALE_INDEX] = out; } } @@ -372,6 +584,7 @@ extern "C" __global__ void rl_fused_controllers( // ═══════════════════════════════════════════════════════════════════ // 8. PPO ratio clamp controller → ISV[440] // Reads ε from ISV[402] which was just written above. + // MIN_OUT adaptive from observed ratio variance; MAX_OUT ISV-driven. // ═══════════════════════════════════════════════════════════════════ { const float prev = isv[RL_PPO_RATIO_CLAMP_MAX_INDEX]; @@ -381,16 +594,26 @@ extern "C" __global__ void rl_fused_controllers( const float eps = isv[RL_PPO_CLIP_INDEX]; const float clamp_margin = isv[RL_PPO_CLAMP_MARGIN_INDEX]; float target = (1.0f + eps) * clamp_margin; - target = fmaxf(PPO_RATIO_CLAMP_MIN_OUT, - fminf(target, PPO_RATIO_CLAMP_MAX_OUT)); + + // Adaptive MIN/MAX from observed log-ratio variance. + const float ratio_count = isv[RL_PPO_RATIO_VAR_COUNT_INDEX]; + const float ratio_var = (ratio_count > 1.0f) + ? isv[RL_PPO_RATIO_VAR_M2_INDEX] / (ratio_count - 1.0f) + : 0.0f; + const float ratio_std = sqrtf(ratio_var); + const float adaptive_min = fmaxf(PPO_RATIO_MIN_ABSOLUTE, + 1.0f + ratio_std * PPO_RATIO_MIN_STD_SCALE); + const float adaptive_max = fmaxf(adaptive_min * PPO_RATIO_MAX_OVER_MIN_FACTOR, + isv[RL_PPO_RATIO_CLAMP_MAX_ADAPTIVE_INDEX]); + + target = fmaxf(adaptive_min, fminf(target, adaptive_max)); if (prev == isv[RL_PPO_RATIO_CLAMP_BOOTSTRAP_INDEX]) { isv[RL_PPO_RATIO_CLAMP_MAX_INDEX] = target; } else { - const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR); + const float a = fmaxf(alpha, wiener_floor); float out = (1.0f - a) * prev + a * target; - out = fmaxf(PPO_RATIO_CLAMP_MIN_OUT, - fminf(out, PPO_RATIO_CLAMP_MAX_OUT)); + out = fmaxf(adaptive_min, fminf(out, adaptive_max)); isv[RL_PPO_RATIO_CLAMP_MAX_INDEX] = out; } } @@ -399,7 +622,7 @@ extern "C" __global__ void rl_fused_controllers( // ═══════════════════════════════════════════════════════════════════ // 9. Gate threshold controller → ISV[512,516,517] // Scans dones[B] to compute done_rate EMA → ISV[525]. - // Adjusts conf + FRD thresholds via Schulman bounded step. + // EMA-α is now ISV-resident (slot 638) per spec §Group B Task 13. // ═══════════════════════════════════════════════════════════════════ { const int warmup = (int)isv[RL_GATE_WARMUP_STEPS_INDEX]; @@ -409,7 +632,7 @@ extern "C" __global__ void rl_fused_controllers( if (dones[b] > 0.5f) done_count += 1.0f; } const float done_rate = done_count / (float)b_size; - const float gate_alpha = 0.01f; + const float gate_alpha = isv[RL_GATE_EMA_ALPHA_INDEX]; const float prev_ema = isv[RL_GATE_DONES_EMA_INDEX]; const float dones_ema = (prev_ema == 0.0f && done_rate > 0.0f) ? done_rate @@ -445,8 +668,9 @@ extern "C" __global__ void rl_fused_controllers( } // ═══════════════════════════════════════════════════════════════════ - // 10. Q→π distill lambda controller → ISV[486] - // Asymmetric bounded step on KL_EMA vs target. + // 10. Q→π distill lambda controller → ISV[486] (Special Q) + // Adaptive MIN from Welford-var of q_distill_kl_ema; MAX + rates + // ISV-driven (slots 650/651/652/653). // ═══════════════════════════════════════════════════════════════════ { const float kl_observed = isv[RL_Q_DISTILL_KL_EMA_INDEX]; @@ -454,13 +678,28 @@ extern "C" __global__ void rl_fused_controllers( float lambda = isv[RL_Q_DISTILL_LAMBDA_INDEX]; if (kl_observed > 0.0f && kl_target > 0.0f && lambda > 0.0f) { - const float upper = kl_target * KL_TOLERANCE; - const float lower = kl_target / KL_TOLERANCE; + // Adaptive MIN: max(0.001, sqrt(var) × 0.05). + const float kl_var_count = isv[RL_Q_DISTILL_KL_VAR_COUNT_INDEX]; + const float kl_var = (kl_var_count > 1.0f) + ? isv[RL_Q_DISTILL_KL_VAR_M2_INDEX] / (kl_var_count - 1.0f) + : 0.0f; + const float kl_std = sqrtf(kl_var); + const float adaptive_min = fmaxf(ADAPTIVE_MIN_ABSOLUTE, + kl_std * ADAPTIVE_MIN_STD_SCALE); + isv[RL_Q_DISTILL_LAMBDA_MIN_ADAPTIVE_INDEX] = adaptive_min; + + const float max_lambda = isv[RL_Q_DISTILL_LAMBDA_MAX_INDEX]; + const float kl_tolerance = isv[RL_Q_DISTILL_KL_TOLERANCE_INDEX]; + const float lambda_ramp_rate = isv[RL_Q_DISTILL_LAMBDA_RAMP_RATE_INDEX]; + const float lambda_decay_rate = isv[RL_Q_DISTILL_LAMBDA_DECAY_RATE_INDEX]; + + const float upper = kl_target * kl_tolerance; + const float lower = kl_target / kl_tolerance; if (kl_observed > upper) { - lambda = fminf(MAX_LAMBDA, lambda * LAMBDA_RAMP_RATE); + lambda = fminf(max_lambda, lambda * lambda_ramp_rate); } else if (kl_observed < lower) { - lambda = fmaxf(MIN_LAMBDA, lambda * LAMBDA_DECAY_RATE); + lambda = fmaxf(adaptive_min, lambda * lambda_decay_rate); } isv[RL_Q_DISTILL_LAMBDA_INDEX] = lambda; diff --git a/crates/ml-alpha/cuda/rl_gamma_controller.cu b/crates/ml-alpha/cuda/rl_gamma_controller.cu index 2b14d9b7f..e12cb3465 100644 --- a/crates/ml-alpha/cuda/rl_gamma_controller.cu +++ b/crates/ml-alpha/cuda/rl_gamma_controller.cu @@ -45,9 +45,35 @@ // the policy myopic). #define RL_GAMMA_INDEX 400 -#define GAMMA_MIN 0.995f -#define GAMMA_MAX 0.999f -#define WIENER_ALPHA_FLOOR 0.4f +// γ MAX clamp ceiling is now ISV-driven per the 2026-05-30 clamp-bound +// extension. γ MIN already adaptive via RL_GAMMA_MIN_ADAPTIVE_INDEX (613). +#define RL_GAMMA_MAX_INDEX 649 +// Wiener-α floor — shared across 9 controllers (slot 659). +#define RL_WIENER_ALPHA_FLOOR_INDEX 659 + +// Adaptive controller floors (spec 2026-05-30 Special case G): +// hardcoded `GAMMA_MIN = 0.995f` replaced with an adaptive bound derived +// from the Welford mean of trade duration. With short-horizon trade +// regimes (d_ema = 15.6 in fold 0) the static 0.995 floor pinned gamma +// for the entire run; adaptive bound `1 - 1/(d_smoothed × 2)` lets the +// controller drop gamma low enough for the current regime while keeping +// a hard `GAMMA_MIN_ABSOLUTE = 0.9` to prevent bandit-mode degenerate. +// +// Smoothed-duration choice: Welford mean (slot 609) is used after a +// 100-observation warmup so the gamma↔trade_duration feedback loop is +// damped by the natural N-smoothing of running mean. Before warmup, the +// EMA at slot 417 is used directly (less stable but available +// immediately at cold-start). +#define RL_GAMMA_MIN_ADAPTIVE_INDEX 613 +#define RL_TRADE_DUR_VAR_COUNT_INDEX 608 +#define RL_TRADE_DUR_VAR_MEAN_INDEX 609 +#define RL_MEAN_TRADE_DURATION_EMA_INDEX 417 +#define GAMMA_MIN_ABSOLUTE 0.9f +#define HORIZON_MULTIPLIER 2.0f +#define WELFORD_WARMUP_OBS 100.0f +// Floor for `d × HORIZON_MULTIPLIER` to prevent the `1/d` term from +// blowing up when trade duration is very small (cold start, no closes). +#define HORIZON_FLOOR 10.0f // ───────────────────────────────────────────────────────────────────── @@ -84,6 +110,27 @@ extern "C" __global__ void rl_gamma_controller( const float gamma_prev = isv[RL_GAMMA_INDEX]; + // Adaptive GAMMA_MIN (spec 2026-05-30 Special case G). + // Replaces hardcoded `GAMMA_MIN = 0.995f`. With d_ema = 15.6 (fold 0) + // the static 0.995 floor pinned γ for the entire run — the target + // formula `0.5^(1/d)` produces γ_target = 0.936 at d = 15.6, which + // clamps to 0.995 = floor and never moves. Adaptive bound + // `1 - 1/(d × 2)` gives γ_min ≈ 0.968 at d = 15.6 — leaving room + // for the controller to track the actual trade horizon. + // + // After 100 Welford observations the running mean is used (more + // stable than EMA against per-batch outliers); before that the EMA + // at slot 417 is used directly so cold-start has a usable signal. + const float d_welford_count = isv[RL_TRADE_DUR_VAR_COUNT_INDEX]; + const float d_welford_mean = isv[RL_TRADE_DUR_VAR_MEAN_INDEX]; + const float d_smoothed = (d_welford_count >= WELFORD_WARMUP_OBS) + ? d_welford_mean + : isv[RL_MEAN_TRADE_DURATION_EMA_INDEX]; + const float adaptive_min = fmaxf(GAMMA_MIN_ABSOLUTE, + 1.0f - 1.0f / fmaxf(d_smoothed * HORIZON_MULTIPLIER, + HORIZON_FLOOR)); + isv[RL_GAMMA_MIN_ADAPTIVE_INDEX] = adaptive_min; + // Compute target from the current input EMA. Shared between // bootstrap and per-step paths so the dead-zone coincidence with // a hardcoded bootstrap cannot recur. @@ -91,26 +138,25 @@ extern "C" __global__ void rl_gamma_controller( // single-event trade doesn't push γ to 0.5. const float mean_trade_duration_events = isv[input_slot]; const float d = fmaxf(mean_trade_duration_events, 1.0f); + const float gamma_max = isv[RL_GAMMA_MAX_INDEX]; float gamma_target = powf(0.5f, 1.0f / d); - gamma_target = fmaxf(GAMMA_MIN, fminf(gamma_target, GAMMA_MAX)); + gamma_target = fmaxf(adaptive_min, fminf(gamma_target, gamma_max)); // Bootstrap on sentinel 0.0 per pearl_first_observation_bootstrap: // first emit replaces directly with the computed target. At cold // start (input EMA also sentinel-zero), clamped d=1, target=0.5, - // clamped to GAMMA_MIN = 0.90 (the floor). Any non-sentinel input - // produces target ≥ 0.90 → ≤ 0.999, distinct from the floor so - // the per-step Wiener blend on subsequent calls always sees a - // prev vs target delta. + // clamped to adaptive_min (≥ GAMMA_MIN_ABSOLUTE = 0.90). Any + // non-sentinel input produces target ≥ adaptive_min → ≤ 0.999. if (gamma_prev == 0.0f) { isv[RL_GAMMA_INDEX] = gamma_target; return; } // Wiener-α blend with floor per pearl_wiener_alpha_floor_for_nonstationary. - const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR); + const float a = fmaxf(alpha, isv[RL_WIENER_ALPHA_FLOOR_INDEX]); float gamma_new = (1.0f - a) * gamma_prev + a * gamma_target; // Clamp into the bounded range; runaway protection. - gamma_new = fmaxf(GAMMA_MIN, fminf(gamma_new, GAMMA_MAX)); + gamma_new = fmaxf(adaptive_min, fminf(gamma_new, gamma_max)); isv[RL_GAMMA_INDEX] = gamma_new; } diff --git a/crates/ml-alpha/cuda/rl_gate_threshold_controller.cu b/crates/ml-alpha/cuda/rl_gate_threshold_controller.cu index 76b4c5a36..b91bf2e55 100644 --- a/crates/ml-alpha/cuda/rl_gate_threshold_controller.cu +++ b/crates/ml-alpha/cuda/rl_gate_threshold_controller.cu @@ -23,6 +23,12 @@ #define RL_GATE_FRD_MAX_INDEX 530 #define RL_GATE_ADJUST_RATE_INDEX 531 #define RL_STEP_COUNTER_ISV_INDEX 548 +// Adaptive controller floors (spec 2026-05-30-adaptive-controller-floor-design): +// EMA smoothing α was hardcoded 0.01f; ISV slot 638 makes it tunable +// at runtime so the dones-rate EMA half-life can adapt to data volume +// (e.g. small-batch smoke runs may want a faster α to surface signal +// before warmup completes). +#define RL_GATE_EMA_ALPHA_INDEX 638 extern "C" __global__ void rl_gate_threshold_controller( float* __restrict__ isv, @@ -41,7 +47,7 @@ extern "C" __global__ void rl_gate_threshold_controller( if (dones[b] > 0.5f) done_count += 1.0f; } const float done_rate = done_count / (float)b_size; - const float alpha = 0.01f; + const float alpha = isv[RL_GATE_EMA_ALPHA_INDEX]; const float prev_ema = isv[RL_GATE_DONES_EMA_INDEX]; const float dones_ema = (prev_ema == 0.0f && done_rate > 0.0f) ? done_rate diff --git a/crates/ml-alpha/cuda/rl_per_alpha_controller.cu b/crates/ml-alpha/cuda/rl_per_alpha_controller.cu index 51e0e9d33..5a45c2836 100644 --- a/crates/ml-alpha/cuda/rl_per_alpha_controller.cu +++ b/crates/ml-alpha/cuda/rl_per_alpha_controller.cu @@ -42,8 +42,10 @@ // sampling is perfectly proportional to priority. #define RL_PER_ALPHA_INDEX 405 -#define PER_ALPHA_MIN 0.3f -#define PER_ALPHA_MAX 1.0f +// PER α MIN/MAX clamp bounds are now ISV-driven per the 2026-05-30 +// clamp-bound extension. Runtime-tunable + visible in diag. +#define RL_PER_ALPHA_MIN_INDEX 647 +#define RL_PER_ALPHA_MAX_INDEX 648 // Kurtosis of a standard normal = 3.0 ("excess kurtosis 0" with the // alternative convention). Used as the breakpoint above which we start // raising α. @@ -52,6 +54,8 @@ #define RL_KURT_GAUSSIAN_INDEX 471 #define RL_KURT_NOISE_FLOOR_INDEX 472 #define RL_KURT_LIFT_SCALE_INDEX 459 +// Schulman tolerance from the shared global slot. +#define RL_SCHULMAN_TOLERANCE_INDEX 468 // 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 @@ -62,7 +66,23 @@ // pattern on the other controllers (per // pearl_multiplicative_controllers_need_bounded_step_and_noise_floor). // (KURT_NOISE_FLOOR — was 1.0f #define — now isv[RL_KURT_NOISE_FLOOR_INDEX]) -#define WIENER_ALPHA_FLOOR 0.4f +// Wiener-α floor — shared across 9 controllers (slot 659). +#define RL_WIENER_ALPHA_FLOOR_INDEX 659 + +// Adaptive controller floors (spec 2026-05-30-adaptive-controller-floor-design): +// noise floor augmented with Welford-derived adaptive component (replaces +// reliance on the ISV-driven RL_KURT_NOISE_FLOOR_INDEX absolute floor alone). +// Asymmetric Schulman semantics: α RISES on a single above-band kurtosis +// observation (heavy tails = safety signal, sharpen PER fast); α FALLS only +// after N consecutive below-band observations (light tails = patient drift). +// Replaces the regime where the controller stuck at MIN 0.4 because every +// step's kurtosis read landed just above the absolute floor but well below +// the band — the Wiener blend dragged α down even on single observations. +#define RL_TD_KURT_VAR_COUNT_INDEX 604 +#define RL_TD_KURT_VAR_M2_INDEX 606 +#define RL_TD_KURT_BELOW_COUNT_INDEX 607 +#define NOISE_FLOOR_STD_MULTIPLIER 2.0f // floor ≥ 2σ of observed kurtosis +#define WIDEN_PATIENCE_CONSECUTIVE 3.0f // α descent requires N below-band steps // ───────────────────────────────────────────────────────────────────── // rl_per_alpha_controller: @@ -97,12 +117,23 @@ extern "C" __global__ void rl_per_alpha_controller( 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 + // Noise-floor gate: kurtosis below the adaptive 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. + // + // Adaptive component (spec 2026-05-30): floor scales with observed + // kurtosis std. Replaces reliance on the absolute ISV floor alone. + // Welford sample variance = M² / (count − 1) when count > 1. const float td_kurtosis_ema = isv[input_slot]; - if (td_kurtosis_ema > 0.0f && td_kurtosis_ema < isv[RL_KURT_NOISE_FLOOR_INDEX]) { + const float kurt_count = isv[RL_TD_KURT_VAR_COUNT_INDEX]; + const float kurt_var = (kurt_count > 1.0f) + ? isv[RL_TD_KURT_VAR_M2_INDEX] / (kurt_count - 1.0f) + : 0.0f; + const float kurt_std = sqrtf(kurt_var); + const float adaptive_noise_floor = fmaxf(isv[RL_KURT_NOISE_FLOOR_INDEX], + kurt_std * NOISE_FLOOR_STD_MULTIPLIER); + if (td_kurtosis_ema > 0.0f && td_kurtosis_ema < adaptive_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 @@ -120,10 +151,12 @@ extern "C" __global__ void rl_per_alpha_controller( // 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 per_alpha_min = isv[RL_PER_ALPHA_MIN_INDEX]; + const float per_alpha_max = isv[RL_PER_ALPHA_MAX_INDEX]; const float kurt_excess = fmaxf(0.0f, td_kurtosis_ema - isv[RL_KURT_GAUSSIAN_INDEX]); const float kurt_lift_scale = isv[RL_KURT_LIFT_SCALE_INDEX]; float target = 0.4f + 0.2f * (kurt_excess / kurt_lift_scale); - target = fmaxf(PER_ALPHA_MIN, fminf(target, PER_ALPHA_MAX)); + 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 @@ -138,9 +171,29 @@ extern "C" __global__ void rl_per_alpha_controller( } // Wiener-α blend with floor per pearl_wiener_alpha_floor_for_nonstationary. - const float a = fmaxf(alpha_step, WIENER_ALPHA_FLOOR); + const float a = fmaxf(alpha_step, isv[RL_WIENER_ALPHA_FLOOR_INDEX]); float out = (1.0f - a) * prev + a * target; - out = fmaxf(PER_ALPHA_MIN, fminf(out, PER_ALPHA_MAX)); + out = fmaxf(per_alpha_min, fminf(out, per_alpha_max)); + + // Asymmetric Schulman patience on the DESCENT direction (spec 2026-05-30): + // PER α rises on a single above-band kurtosis observation (heavy tails + // = act fast to concentrate sampling on informative tails). PER α + // falls only after N consecutive below-band observations — a single + // light-tailed step must not drag α toward MIN. + const float kurt_gaussian = isv[RL_KURT_GAUSSIAN_INDEX]; + const float tolerance = isv[RL_SCHULMAN_TOLERANCE_INDEX]; + if (out < prev && td_kurtosis_ema < kurt_gaussian / tolerance) { + const float new_count = isv[RL_TD_KURT_BELOW_COUNT_INDEX] + 1.0f; + isv[RL_TD_KURT_BELOW_COUNT_INDEX] = new_count; + if (new_count < WIDEN_PATIENCE_CONSECUTIVE) { + // Hold at prev until patience accumulates. + isv[RL_PER_ALPHA_INDEX] = prev; + return; + } + } else { + isv[RL_TD_KURT_BELOW_COUNT_INDEX] = 0.0f; + } + isv[RL_PER_ALPHA_INDEX] = out; } diff --git a/crates/ml-alpha/cuda/rl_ppo_clip_controller.cu b/crates/ml-alpha/cuda/rl_ppo_clip_controller.cu index 457a1f5ab..2f75d3afd 100644 --- a/crates/ml-alpha/cuda/rl_ppo_clip_controller.cu +++ b/crates/ml-alpha/cuda/rl_ppo_clip_controller.cu @@ -42,16 +42,36 @@ // when ratios run away). #define RL_PPO_CLIP_INDEX 402 -#define EPS_MIN 0.05f -#define EPS_MAX 0.5f +// ε MIN/MAX are now ISV-driven per the 2026-05-30 clamp-bound extension — +// the prior hardcoded `[0.05, 0.5]` calibration was a fixed-regime choice +// that contributed to fold 0/1 saturation when Phase 4.5 narrowed the KL +// signal distribution. ISV-residence makes the bound runtime-tunable +// (visible in diag JSONL, re-seedable without recompile). +#define RL_PPO_CLIP_EPS_MIN_INDEX 640 +#define RL_PPO_CLIP_EPS_MAX_INDEX 641 // 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 +// Wiener-α floor — shared across 9 controllers per +// `pearl_wiener_alpha_floor_for_nonstationary`. ISV-resident at slot 659 +// so all consumers stay aligned without per-file `#define` drift. +#define RL_WIENER_ALPHA_FLOOR_INDEX 659 + +// Adaptive controller floors (spec 2026-05-30-adaptive-controller-floor-design): +// noise floor derives from observed kl_pi_ema variance via Welford triples, +// asymmetric Schulman widening requires N consecutive below-band observations. +// Replaces hardcoded KL_NOISE_FLOOR_FRAC=0.01f which was the root cause of +// the fold 0/1 saturation (ε locked at MAX 0.50 within 50 steps because the +// 1%-of-target floor was 100× too low for the post-Phase-4.5 KL regime). +#define RL_KL_PI_VAR_COUNT_INDEX 588 +#define RL_KL_PI_VAR_M2_INDEX 590 +#define RL_KL_PI_BELOW_COUNT_INDEX 591 +#define NOISE_FLOOR_TARGET_FRAC 0.5f // floor ≥ 50% of target +#define NOISE_FLOOR_STD_MULTIPLIER 2.0f // floor ≥ 2σ of observed signal +#define WIDEN_PATIENCE_CONSECUTIVE 3.0f // widen requires N below-band steps // ───────────────────────────────────────────────────────────────────── @@ -93,26 +113,49 @@ extern "C" __global__ void rl_ppo_clip_controller( // 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. + // Adaptive noise floor — signal-driven per + // `pearl_zscore_normalization_for_magnitude_asymmetric_signals` and + // `feedback_adaptive_not_tuned`. Floor must scale with observed signal + // magnitude AND have an absolute minimum tied to the target so a + // momentarily-noisy signal can't trigger spurious adjustments. + // Welford sample variance = M² / (count − 1) when count > 1. + const float kl_count = isv[RL_KL_PI_VAR_COUNT_INDEX]; + const float kl_var = (kl_count > 1.0f) + ? isv[RL_KL_PI_VAR_M2_INDEX] / (kl_count - 1.0f) + : 0.0f; + const float kl_std = sqrtf(kl_var); + const float kl_noise_floor = fmaxf(kl_target * NOISE_FLOOR_TARGET_FRAC, + kl_std * NOISE_FLOOR_STD_MULTIPLIER); + + // Noise-floor gate: KL below the adaptive floor is dominated by + // numerical noise OR signal stays naturally below target under the + // current architecture (e.g., Phase 4.5 reduces KL by ~50×). Hold ε. if (kl_ema < kl_noise_floor) return; - // Bounded multiplicative adjustment (Schulman-style adaptive KL). - // Schulman params from ISV — shared with target_tau, rollout_steps. + // Asymmetric Schulman (spec 2026-05-30 Section "Approach C, folded in"): + // tightening fires on a single above-band observation — real policy + // divergence is a safety signal we act on fast. Widening requires + // N consecutive below-band observations so a single noisy step can't + // drive ε past bootstrap. Below-band counter is per-controller in ISV. 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; + ratio = 1.0f / adjust_rate; // tighten — single observation + isv[RL_KL_PI_BELOW_COUNT_INDEX] = 0.0f; // reset patience } else if (kl_ema < kl_target / tolerance) { - ratio = adjust_rate; + const float new_count = isv[RL_KL_PI_BELOW_COUNT_INDEX] + 1.0f; + isv[RL_KL_PI_BELOW_COUNT_INDEX] = new_count; + ratio = (new_count >= WIDEN_PATIENCE_CONSECUTIVE) ? adjust_rate : 1.0f; } else { + isv[RL_KL_PI_BELOW_COUNT_INDEX] = 0.0f; // in-band → reset ratio = 1.0f; } + const float eps_min = isv[RL_PPO_CLIP_EPS_MIN_INDEX]; + const float eps_max = isv[RL_PPO_CLIP_EPS_MAX_INDEX]; float eps_target = eps_prev * ratio; - eps_target = fmaxf(EPS_MIN, fminf(eps_target, EPS_MAX)); + eps_target = fmaxf(eps_min, fminf(eps_target, eps_max)); // First-observation replace-directly per // `pearl_first_observation_bootstrap`. See rl_target_tau_controller @@ -125,9 +168,9 @@ extern "C" __global__ void rl_ppo_clip_controller( } // Wiener-α blend with floor per pearl_wiener_alpha_floor_for_nonstationary. - const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR); + const float a = fmaxf(alpha, isv[RL_WIENER_ALPHA_FLOOR_INDEX]); float eps_new = (1.0f - a) * eps_prev + a * eps_target; - eps_new = fmaxf(EPS_MIN, fminf(eps_new, EPS_MAX)); + eps_new = fmaxf(eps_min, fminf(eps_new, eps_max)); isv[RL_PPO_CLIP_INDEX] = eps_new; } diff --git a/crates/ml-alpha/cuda/rl_ppo_ratio_clamp_controller.cu b/crates/ml-alpha/cuda/rl_ppo_ratio_clamp_controller.cu index 210bd7b7c..222d644b6 100644 --- a/crates/ml-alpha/cuda/rl_ppo_ratio_clamp_controller.cu +++ b/crates/ml-alpha/cuda/rl_ppo_ratio_clamp_controller.cu @@ -53,9 +53,21 @@ // Default 10.0 — clamp_max = (1+ε) × this. Seeded by rl_isv_write // at init. Tuning lower tightens the importance-ratio bound. #define RL_PPO_CLAMP_MARGIN_INDEX 460 -#define PPO_RATIO_CLAMP_MIN_OUT 2.0f -#define PPO_RATIO_CLAMP_MAX_OUT 1000.0f -#define WIENER_ALPHA_FLOOR 0.4f + +// Adaptive controller floors (spec 2026-05-30-adaptive-controller-floor-design): +// the MIN_OUT floor derives from observed log-ratio variance (Welford +// triple at slots 630/631/632) rather than the static `2.0f` baseline. +// The MAX ceiling is ISV-resident at slot 629 so it can be tuned at +// runtime. The hardcoded `2.0f` survives as an architectural minimum +// per the spec's "physics-constant" exemption ("don't degenerate to +// vanilla policy gradient when σ_ratio is tiny"). +#define RL_PPO_RATIO_CLAMP_MAX_ADAPTIVE_INDEX 629 +#define RL_PPO_RATIO_VAR_COUNT_INDEX 630 +#define RL_PPO_RATIO_VAR_M2_INDEX 632 +#define RL_WIENER_ALPHA_FLOOR_INDEX 659 +#define PPO_RATIO_MIN_ABSOLUTE 2.0f // architectural exemption: vanilla-PG safeguard +#define PPO_RATIO_MIN_STD_SCALE 3.0f // adaptive min = max(2.0, 1 + 3σ) +#define PPO_RATIO_MAX_OVER_MIN_FACTOR 5.0f // adaptive max ≥ 5× adaptive min // ───────────────────────────────────────────────────────────────────── // rl_ppo_ratio_clamp_controller: @@ -98,9 +110,25 @@ extern "C" __global__ void rl_ppo_ratio_clamp_controller( const float clamp_margin = isv[RL_PPO_CLAMP_MARGIN_INDEX]; float target = (1.0f + eps) * clamp_margin; + // ── Adaptive MIN/MAX from observed log-ratio variance ─────────── + // Welford sample variance = M² / (count − 1) when count > 1. + // Adaptive min: max(2.0, 1 + 3σ_ratio). The 2.0 floor is the + // architectural-minimum exemption ("don't degenerate to vanilla + // policy gradient") — kept as a #define per the spec. The 3σ scale + // gives the clamp window roughly the same coverage of healthy + // policy updates as the prior static 2.0 baseline did pre-Phase 4.5. + const float ratio_count = isv[RL_PPO_RATIO_VAR_COUNT_INDEX]; + const float ratio_var = (ratio_count > 1.0f) + ? isv[RL_PPO_RATIO_VAR_M2_INDEX] / (ratio_count - 1.0f) + : 0.0f; + const float ratio_std = sqrtf(ratio_var); + const float adaptive_min = fmaxf(PPO_RATIO_MIN_ABSOLUTE, + 1.0f + ratio_std * PPO_RATIO_MIN_STD_SCALE); + const float adaptive_max = fmaxf(adaptive_min * PPO_RATIO_MAX_OVER_MIN_FACTOR, + isv[RL_PPO_RATIO_CLAMP_MAX_ADAPTIVE_INDEX]); + // Permanent floor / ceiling per pearl_blend_formulas_must_have_permanent_floor. - target = fmaxf(PPO_RATIO_CLAMP_MIN_OUT, - fminf(target, PPO_RATIO_CLAMP_MAX_OUT)); + target = fmaxf(adaptive_min, fminf(target, adaptive_max)); // First-observation replace-directly per pearl_first_observation_bootstrap. // prev == PPO_RATIO_CLAMP_BOOTSTRAP means "we just bootstrapped last @@ -113,10 +141,9 @@ extern "C" __global__ void rl_ppo_ratio_clamp_controller( } // Wiener-α blend with floor. - const float a = fmaxf(alpha_step, WIENER_ALPHA_FLOOR); + const float a = fmaxf(alpha_step, isv[RL_WIENER_ALPHA_FLOOR_INDEX]); float out = (1.0f - a) * prev + a * target; - out = fmaxf(PPO_RATIO_CLAMP_MIN_OUT, - fminf(out, PPO_RATIO_CLAMP_MAX_OUT)); + out = fmaxf(adaptive_min, fminf(out, adaptive_max)); isv[RL_PPO_RATIO_CLAMP_MAX_INDEX] = out; } diff --git a/crates/ml-alpha/cuda/rl_q_distill_lambda_controller.cu b/crates/ml-alpha/cuda/rl_q_distill_lambda_controller.cu index 86ec6ee83..aaf483358 100644 --- a/crates/ml-alpha/cuda/rl_q_distill_lambda_controller.cu +++ b/crates/ml-alpha/cuda/rl_q_distill_lambda_controller.cu @@ -23,11 +23,27 @@ #define RL_Q_DISTILL_KL_EMA_INDEX 488 #define RL_Q_DISTILL_KL_TARGET_INDEX 491 -#define MIN_LAMBDA 0.05f -#define MAX_LAMBDA 1.0f -#define KL_TOLERANCE 3.0f -#define LAMBDA_RAMP_RATE 1.2f -#define LAMBDA_DECAY_RATE 0.998f +// Adaptive controller floors (spec 2026-05-30 Special case Q): +// hardcoded `MIN_LAMBDA = 0.05f` floor replaced with adaptive bound +// derived from Welford variance on q_distill_kl_ema. Phase 4.5 reduces +// KL signal magnitudes which this controller consumes; the "below target" +// path fires continuously, decaying λ to MIN. Adaptive bound +// `max(0.001, sqrt(var) × 0.05)` lets λ decay to a level proportional to +// the actual signal noise rather than the hardcoded 0.05 floor. +#define RL_Q_DISTILL_KL_VAR_COUNT_INDEX 618 +#define RL_Q_DISTILL_KL_VAR_M2_INDEX 620 +#define RL_Q_DISTILL_LAMBDA_MIN_ADAPTIVE_INDEX 639 + +// MAX_LAMBDA, KL_TOLERANCE, LAMBDA_RAMP_RATE, LAMBDA_DECAY_RATE are now +// ISV-driven per the 2026-05-30 clamp-bound extension. ADAPTIVE_MIN_* +// remain hardcoded — they're new constants from the noise-floor design +// itself per spec exemption ("not clamp bounds, new pattern constants"). +#define RL_Q_DISTILL_LAMBDA_MAX_INDEX 650 +#define RL_Q_DISTILL_KL_TOLERANCE_INDEX 651 +#define RL_Q_DISTILL_LAMBDA_RAMP_RATE_INDEX 652 +#define RL_Q_DISTILL_LAMBDA_DECAY_RATE_INDEX 653 +#define ADAPTIVE_MIN_ABSOLUTE 0.001f +#define ADAPTIVE_MIN_STD_SCALE 0.05f extern "C" __global__ void rl_q_distill_lambda_controller( float* __restrict__ isv @@ -40,13 +56,32 @@ extern "C" __global__ void rl_q_distill_lambda_controller( 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; + // Adaptive MIN bound (spec 2026-05-30 Special case Q): replaces + // hardcoded `MIN_LAMBDA = 0.05f`. Welford sample variance = + // M² / (count − 1) when count > 1. Adaptive min scales with signal + // std so λ can decay proportional to actual KL noise rather than + // pegging at an arbitrary 0.05 floor. + const float kl_var_count = isv[RL_Q_DISTILL_KL_VAR_COUNT_INDEX]; + const float kl_var = (kl_var_count > 1.0f) + ? isv[RL_Q_DISTILL_KL_VAR_M2_INDEX] / (kl_var_count - 1.0f) + : 0.0f; + const float kl_std = sqrtf(kl_var); + const float adaptive_min = fmaxf(ADAPTIVE_MIN_ABSOLUTE, + kl_std * ADAPTIVE_MIN_STD_SCALE); + isv[RL_Q_DISTILL_LAMBDA_MIN_ADAPTIVE_INDEX] = adaptive_min; + + const float max_lambda = isv[RL_Q_DISTILL_LAMBDA_MAX_INDEX]; + const float kl_tolerance = isv[RL_Q_DISTILL_KL_TOLERANCE_INDEX]; + const float lambda_ramp_rate = isv[RL_Q_DISTILL_LAMBDA_RAMP_RATE_INDEX]; + const float lambda_decay_rate = isv[RL_Q_DISTILL_LAMBDA_DECAY_RATE_INDEX]; + + const float upper = kl_target * kl_tolerance; + const float lower = kl_target / kl_tolerance; if (kl_observed > upper) { - lambda = fminf(MAX_LAMBDA, lambda * LAMBDA_RAMP_RATE); + lambda = fminf(max_lambda, lambda * lambda_ramp_rate); } else if (kl_observed < lower) { - lambda = fmaxf(MIN_LAMBDA, lambda * LAMBDA_DECAY_RATE); + lambda = fmaxf(adaptive_min, lambda * lambda_decay_rate); } isv[RL_Q_DISTILL_LAMBDA_INDEX] = lambda; diff --git a/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu b/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu index d9bfe7daa..133d33dac 100644 --- a/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu +++ b/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu @@ -88,25 +88,21 @@ #define RL_NEG_SCALED_REWARD_MAX_INDEX 489 #define RL_NEG_SCALED_REWARD_MAX_EMA_INDEX 490 -#define MIN_WIN 1.0f -#define MIN_RATIO 1.0f // no inverted asymmetry (loss never < win) -#define MAX_RATIO 3.0f // no worse than original loss-aversion -#define V_BOUND_FLOOR 1.0f // |V_MIN| and V_MAX both ≥ 1.0 -#define V_BOUND_EWMA_ALPHA 0.001f // slow — half-life ~700 steps -#define WIENER_ALPHA_FLOOR 0.4f - -// MARGIN Schulman bounded-step bounds + adjust rate per -// `pearl_multiplicative_controllers_need_bounded_step_and_noise_floor`. -// TOLERANCE=1.5 → dead-zone is [target/1.5, target×1.5] = [0.033, 0.075]; -// ADJUST_RATE=1.2 → MARGIN moves by ≤20% per controller call (one per -// trainer step). Bounds [MIN_MARGIN=1.0, MAX_MARGIN=5.0] — MIN preserves -// the initial passthrough behaviour; MAX × (mean pos_max_ema≈3-5) → -// WIN_eff up to ~15-25 which is well within MAX_WIN=50 ceiling. -#define MIN_MARGIN 1.0f -#define MAX_MARGIN 5.0f -#define MARGIN_TOLERANCE 1.5f -#define MARGIN_ADJUST_RATE 1.2f -#define CLIP_RATE_EMA_ALPHA 0.05f +// Adaptive controller floors (spec 2026-05-30-adaptive-controller-floor-design): +// every prior `#define` clamp-bound constant becomes ISV-driven so it can +// be tuned at runtime via re-seed instead of recompile. Defaults preserve +// the historic behaviour exactly. +#define RL_REWARD_CLAMP_V_BOUND_FLOOR_INDEX 633 +#define RL_REWARD_CLAMP_V_BOUND_EWMA_ALPHA_INDEX 634 +#define RL_REWARD_CLAMP_MIN_WIN_INDEX 635 +#define RL_REWARD_CLAMP_MIN_RATIO_INDEX 636 +#define RL_REWARD_CLAMP_MAX_RATIO_INDEX 637 +#define RL_REWARD_CLAMP_MIN_MARGIN_INDEX 654 +#define RL_REWARD_CLAMP_MAX_MARGIN_INDEX 655 +#define RL_REWARD_CLAMP_MARGIN_TOLERANCE_INDEX 656 +#define RL_REWARD_CLAMP_MARGIN_ADJUST_RATE_INDEX 657 +#define RL_REWARD_CLAMP_CLIP_RATE_EMA_ALPHA_INDEX 658 +#define RL_WIENER_ALPHA_FLOOR_INDEX 659 extern "C" __global__ void rl_reward_clamp_controller( float* __restrict__ isv, @@ -114,6 +110,21 @@ extern "C" __global__ void rl_reward_clamp_controller( ) { if (threadIdx.x != 0 || blockIdx.x != 0) return; + // ── ISV-driven design constants (spec 2026-05-30) ── + // Loaded once at the top of the kernel so they're consistent across + // all sub-controllers (EMA → MARGIN → RATIO → C51 span) within this + // step. Default values preserve historic behaviour exactly. + const float v_bound_floor = isv[RL_REWARD_CLAMP_V_BOUND_FLOOR_INDEX]; + const float v_bound_ewma_alpha = isv[RL_REWARD_CLAMP_V_BOUND_EWMA_ALPHA_INDEX]; + const float min_ratio_isv = isv[RL_REWARD_CLAMP_MIN_RATIO_INDEX]; + const float max_ratio_isv = isv[RL_REWARD_CLAMP_MAX_RATIO_INDEX]; + const float min_margin_isv = isv[RL_REWARD_CLAMP_MIN_MARGIN_INDEX]; + const float max_margin_isv = isv[RL_REWARD_CLAMP_MAX_MARGIN_INDEX]; + const float margin_tolerance = isv[RL_REWARD_CLAMP_MARGIN_TOLERANCE_INDEX]; + const float margin_adjust_rate = isv[RL_REWARD_CLAMP_MARGIN_ADJUST_RATE_INDEX]; + const float clip_rate_ema_alpha = isv[RL_REWARD_CLAMP_CLIP_RATE_EMA_ALPHA_INDEX]; + const float wiener_floor = isv[RL_WIENER_ALPHA_FLOOR_INDEX]; + const float pos_max = isv[RL_POS_SCALED_REWARD_MAX_INDEX]; const float ema_prev = isv[RL_POS_SCALED_REWARD_MAX_EMA_INDEX]; @@ -136,7 +147,7 @@ extern "C" __global__ void rl_reward_clamp_controller( if (ema_prev == 0.0f) { ema_new = pos_max; // first-observation bootstrap } else { - const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR); + const float a = fmaxf(alpha, wiener_floor); ema_new = (1.0f - a) * ema_prev + a * pos_max; } isv[RL_POS_SCALED_REWARD_MAX_EMA_INDEX] = ema_new; @@ -171,8 +182,8 @@ extern "C" __global__ void rl_reward_clamp_controller( if (clip_rate_prev == 0.0f) { clip_rate_new = clip_indicator; // bootstrap } else { - clip_rate_new = (1.0f - CLIP_RATE_EMA_ALPHA) * clip_rate_prev - + CLIP_RATE_EMA_ALPHA * clip_indicator; + clip_rate_new = (1.0f - clip_rate_ema_alpha) * clip_rate_prev + + clip_rate_ema_alpha * clip_indicator; } isv[RL_REWARD_CLAMP_CLIP_RATE_EMA_INDEX] = clip_rate_new; } @@ -184,12 +195,12 @@ extern "C" __global__ void rl_reward_clamp_controller( float margin = isv[RL_REWARD_CLAMP_MARGIN_INDEX]; if (pos_max > 0.0f) { const float clip_target = isv[RL_REWARD_CLAMP_CLIP_RATE_TARGET_INDEX]; - const float upper = clip_target * MARGIN_TOLERANCE; - const float lower = clip_target / MARGIN_TOLERANCE; + const float upper = clip_target * margin_tolerance; + const float lower = clip_target / margin_tolerance; if (clip_rate_new > upper) { - margin = fminf(MAX_MARGIN, margin * MARGIN_ADJUST_RATE); + margin = fminf(max_margin_isv, margin * margin_adjust_rate); } else if (clip_rate_new < lower) { - margin = fmaxf(MIN_MARGIN, margin / MARGIN_ADJUST_RATE); + margin = fmaxf(min_margin_isv, margin / margin_adjust_rate); } isv[RL_REWARD_CLAMP_MARGIN_INDEX] = margin; } @@ -210,7 +221,7 @@ extern "C" __global__ void rl_reward_clamp_controller( if (neg_prev == 0.0f) { neg_new = neg_max; } else { - const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR); + const float a = fmaxf(alpha, wiener_floor); neg_new = (1.0f - a) * neg_prev + a * neg_max; } isv[RL_NEG_SCALED_REWARD_MAX_EMA_INDEX] = neg_new; @@ -219,8 +230,8 @@ extern "C" __global__ void rl_reward_clamp_controller( // statically-seeded RATIO=3.0 until both EMAs are non-zero. if (ema_new > 0.0f && neg_new > 0.0f) { const float ratio_target = neg_new / ema_new; - const float ratio_clamped = fmaxf(MIN_RATIO, - fminf(MAX_RATIO, ratio_target)); + const float ratio_clamped = fmaxf(min_ratio_isv, + fminf(max_ratio_isv, ratio_target)); isv[RL_REWARD_CLAMP_RATIO_INDEX] = ratio_clamped; } @@ -255,20 +266,20 @@ extern "C" __global__ void rl_reward_clamp_controller( 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_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); + : (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_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); + : (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); } } diff --git a/crates/ml-alpha/cuda/rl_reward_scale_controller.cu b/crates/ml-alpha/cuda/rl_reward_scale_controller.cu index 58e586e56..f60d95a49 100644 --- a/crates/ml-alpha/cuda/rl_reward_scale_controller.cu +++ b/crates/ml-alpha/cuda/rl_reward_scale_controller.cu @@ -48,7 +48,30 @@ // Floor for the mean_abs_pnl_ema denominator to guard against div-by-zero // when the EMA hasn't accumulated any closed trades yet. #define EPS_PNL 1e-3f -#define WIENER_ALPHA_FLOOR 0.4f +// Wiener-α floor — shared across 9 controllers (slot 659). +#define RL_WIENER_ALPHA_FLOOR_INDEX 659 + +// Adaptive controller floors (spec 2026-05-30 Special case R): +// asymmetric rate limit on DECREASE (new ≥ prev / 1.05) + bootstrap-fraction +// floor until N=100 trades have closed. Together these prevent the +// catastrophic 1.0 → 0.004 crash in 100 steps observed in fold 0 (where a +// single large pnl signal — pre-trade — would crater the scale before any +// closed-trade signal could anchor it). +// +// Trade-count signal: RL_TRADE_DUR_VAR_COUNT_INDEX (608) is the Welford +// count over per-batch trade durations — incremented once per closed trade. +// Reusing it as the trade-count signal avoids a parallel counter. +// +// Emit: RL_REWARD_MAGNITUDE_EMA_INDEX (614) mirrors the input EMA so the +// downstream Welford kernel can track variance for noise-floor derivation +// elsewhere in the controller cascade. Single-source-of-truth: this +// controller's input IS the pnl magnitude EMA; slot 614 is the public +// emit-slot consumed by the Welford-variance kernel. +#define RL_TRADE_DUR_VAR_COUNT_INDEX 608 +#define RL_REWARD_MAGNITUDE_EMA_INDEX 614 +#define BOOTSTRAP_FRACTION_FLOOR 0.1f +#define MIN_TRADES_FOR_RELEASE 100.0f +#define ASYM_DECREASE_RATE 1.05f // ───────────────────────────────────────────────────────────────────── // rl_reward_scale_controller: @@ -105,6 +128,15 @@ extern "C" __global__ void rl_reward_scale_controller( // noise. const float mean_abs_pnl_ema = isv[input_slot]; if (mean_abs_pnl_ema == 0.0f) return; + + // Emit per-batch pnl magnitude EMA to the public slot so the Welford + // variance kernel and any other downstream consumer can track it + // without re-reading the input slot through indirect plumbing. + // Single-source-of-truth: this controller's input IS the pnl magnitude + // EMA — slot 614 mirrors it as the public emit (spec 2026-05-30 + // Special case R bullet (c)). + isv[RL_REWARD_MAGNITUDE_EMA_INDEX] = mean_abs_pnl_ema; + // target_scale = 1.0 / max(mean_abs_pnl_ema, EPS_PNL). // Larger typical PnL → smaller scale (reward is divided down toward ±1). // Smaller typical PnL → larger scale (reward is amplified toward ±1). @@ -113,6 +145,20 @@ extern "C" __global__ void rl_reward_scale_controller( const float scale_min = isv[RL_REWARD_SCALE_MIN_INDEX]; target = fmaxf(scale_min, fminf(target, REWARD_SCALE_MAX)); + // Bootstrap-fraction floor (spec 2026-05-30 Special case R bullet (b)): + // until N=MIN_TRADES_FOR_RELEASE trades have closed, the scale cannot + // drop below 10% of bootstrap. The trade count comes from the Welford + // trade-duration counter (incremented once per closed trade). This + // protects against the early-training spiral where a single large + // mean_abs_pnl pre-trade signal cratered the scale before any closed- + // trade ground truth could anchor it (fold 0 crash 1.0 → 0.004 in 100 + // steps). + const float trade_count = isv[RL_TRADE_DUR_VAR_COUNT_INDEX]; + const float boot = isv[RL_REWARD_SCALE_BOOTSTRAP_INDEX]; + const float boot_floor = (trade_count < MIN_TRADES_FOR_RELEASE) + ? boot * BOOTSTRAP_FRACTION_FLOOR + : scale_min; + // First-observation replace-directly per // `pearl_first_observation_bootstrap`. With prev=1.0 (bootstrap) // and a target far from 1.0 on the first real PnL signal, the @@ -121,13 +167,25 @@ extern "C" __global__ void rl_reward_scale_controller( // feeds V regression at magnitude 500× the atom support's // ±1 expectation. Replacing directly with target eliminates // this cold-start contamination. - if (prev == isv[RL_REWARD_SCALE_BOOTSTRAP_INDEX]) { - isv[RL_REWARD_SCALE_INDEX] = target; + // + // Asymmetric DECREASE rate limit (spec 2026-05-30 bullet (a)): even + // on the first-observation path, cap per-step decrease at 5% + // (new ≥ prev / 1.05) so a tiny initial target can't crater scale + // before downstream signal accumulates. The existing 2% per-step + // symmetric clamp below remains for subsequent steps (stricter than + // 5%, so it stays load-bearing on the slow path). + if (prev == boot) { + float first = target; + if (first < prev) { + first = fmaxf(first, prev / ASYM_DECREASE_RATE); + } + first = fmaxf(boot_floor, fminf(first, REWARD_SCALE_MAX)); + isv[RL_REWARD_SCALE_INDEX] = first; return; } // Wiener-α blend with floor per pearl_wiener_alpha_floor_for_nonstationary. - const float a = fmaxf(alpha_step, WIENER_ALPHA_FLOOR); + const float a = fmaxf(alpha_step, isv[RL_WIENER_ALPHA_FLOOR_INDEX]); float out = (1.0f - a) * prev + a * target; // Per-step movement clamp: scale moves at most ±2% from previous. @@ -139,6 +197,16 @@ extern "C" __global__ void rl_reward_scale_controller( const float min_move = prev * 0.98f; out = fmaxf(min_move, fminf(out, max_move)); - out = fmaxf(scale_min, fminf(out, REWARD_SCALE_MAX)); + // Asymmetric DECREASE rate limit (spec 2026-05-30 bullet (a)): + // belt-and-braces — the symmetric 2% clamp above is stricter than + // the 5% asymmetric cap on the slow path, but the explicit + // asymmetric clamp here documents the controller invariant + // (decreases cannot exceed 5% per step) and survives any future + // relaxation of the 2% symmetric clamp. + if (out < prev) { + out = fmaxf(out, prev / ASYM_DECREASE_RATE); + } + + out = fmaxf(boot_floor, fminf(out, REWARD_SCALE_MAX)); isv[RL_REWARD_SCALE_INDEX] = out; } diff --git a/crates/ml-alpha/cuda/rl_rollout_steps_controller.cu b/crates/ml-alpha/cuda/rl_rollout_steps_controller.cu index ad546a823..634384f85 100644 --- a/crates/ml-alpha/cuda/rl_rollout_steps_controller.cu +++ b/crates/ml-alpha/cuda/rl_rollout_steps_controller.cu @@ -46,16 +46,33 @@ // past the clip band. #define RL_N_ROLLOUT_STEPS_INDEX 404 -#define ROLLOUT_MIN 256.0f -#define ROLLOUT_MAX 8192.0f +// ROLLOUT MIN/MAX clamp bounds are now ISV-driven per the 2026-05-30 +// clamp-bound extension. Runtime-tunable + visible in diag. +#define RL_ROLLOUT_MIN_INDEX 643 +#define RL_ROLLOUT_MAX_INDEX 644 // ISV-driven bootstrap + target. #define RL_ROLLOUT_BOOTSTRAP_INDEX 475 #define RL_ADV_VAR_RATIO_TARGET_INDEX 449 // Schulman params from shared global slots (same as ppo_clip, target_tau). #define RL_SCHULMAN_TOLERANCE_INDEX 468 #define RL_SCHULMAN_ADJUST_RATE_INDEX 469 -#define ADV_VAR_RATIO_NOISE_FLOOR_FRAC 0.01f -#define WIENER_ALPHA_FLOOR 0.4f +// Wiener-α floor — shared across 9 controllers (slot 659). +#define RL_WIENER_ALPHA_FLOOR_INDEX 659 + +// Adaptive controller floors (spec 2026-05-30-adaptive-controller-floor-design): +// noise floor derives from observed advantage-variance signal via Welford +// triples, asymmetric Schulman widening requires N consecutive below-band +// observations. Replaces hardcoded ADV_VAR_RATIO_NOISE_FLOOR_FRAC=0.01f +// (calibrated against the pre-Phase-4.5 advantage regime; under Phase 4.5 +// post-norm advantage variance is definitionally near-zero, so the +// controller is also rewired to consume RL_ADV_VAR_PRE_NORM_INDEX=612 — +// the pre-normalization variance emitted by rl_advantage_normalize.cu). +#define RL_ADV_VAR_VAR_COUNT_INDEX 596 +#define RL_ADV_VAR_VAR_M2_INDEX 598 +#define RL_ADV_VAR_BELOW_COUNT_INDEX 599 +#define NOISE_FLOOR_TARGET_FRAC 0.5f // floor ≥ 50% of target +#define NOISE_FLOOR_STD_MULTIPLIER 2.0f // floor ≥ 2σ of observed signal +#define WIDEN_PATIENCE_CONSECUTIVE 3.0f // widen requires N below-band steps // ───────────────────────────────────────────────────────────────────── // rl_rollout_steps_controller: @@ -103,31 +120,43 @@ extern "C" __global__ void rl_rollout_steps_controller( // RL_ADV_VAR_RATIO_TARGET_INDEX (seeded at trainer init). const float adv_var_target = isv[RL_ADV_VAR_RATIO_TARGET_INDEX]; - // Noise-floor gate: input below target × NOISE_FLOOR_FRAC is - // dominated by numerical noise — hold rollout unchanged. - // Mirrors the ppo_clip / target_tau controllers' design after the - // alpha-rl-mjzfk multiplicative-blow-up incident. - const float adv_var_noise_floor = - adv_var_target * ADV_VAR_RATIO_NOISE_FLOOR_FRAC; + // Adaptive noise floor — signal-driven per + // `pearl_zscore_normalization_for_magnitude_asymmetric_signals` and + // `feedback_adaptive_not_tuned`. Welford sample variance = + // M² / (count − 1) when count > 1. + const float adv_count = isv[RL_ADV_VAR_VAR_COUNT_INDEX]; + const float adv_var = (adv_count > 1.0f) + ? isv[RL_ADV_VAR_VAR_M2_INDEX] / (adv_count - 1.0f) + : 0.0f; + const float adv_std = sqrtf(adv_var); + const float adv_var_noise_floor = fmaxf(adv_var_target * NOISE_FLOOR_TARGET_FRAC, + adv_std * NOISE_FLOOR_STD_MULTIPLIER); if (advantage_var_over_abs_mean < adv_var_noise_floor) return; - // Bounded multiplicative adjustment (Schulman-style adaptive). - // At most ADV_VAR_RATIO_ADJUST_RATE × shift per step regardless of - // how far input is from target. After several consecutive - // out-of-band observations the output drifts smoothly toward - // MIN/MAX, but a single observation can't slam it there. + // Asymmetric Schulman (spec 2026-05-30 Section "Approach C, folded in"): + // widening rollout fires on a single above-band observation — noisy + // advantages are a safety signal we act on fast (collect more samples + // before updating). Shrinking rollout requires N consecutive below-band + // observations so a single quiet step can't push the buffer too small. + // Below-band counter is per-controller in ISV. const float tolerance = isv[RL_SCHULMAN_TOLERANCE_INDEX]; const float adjust_rate = isv[RL_SCHULMAN_ADJUST_RATE_INDEX]; float scale; if (advantage_var_over_abs_mean > adv_var_target * tolerance) { - scale = adjust_rate; + scale = adjust_rate; // widen — single observation + isv[RL_ADV_VAR_BELOW_COUNT_INDEX] = 0.0f; // reset patience } else if (advantage_var_over_abs_mean < adv_var_target / tolerance) { - scale = 1.0f / adjust_rate; + const float new_count = isv[RL_ADV_VAR_BELOW_COUNT_INDEX] + 1.0f; + isv[RL_ADV_VAR_BELOW_COUNT_INDEX] = new_count; + scale = (new_count >= WIDEN_PATIENCE_CONSECUTIVE) ? (1.0f / adjust_rate) : 1.0f; } else { + isv[RL_ADV_VAR_BELOW_COUNT_INDEX] = 0.0f; // in-band → reset scale = 1.0f; } + const float rollout_min = isv[RL_ROLLOUT_MIN_INDEX]; + const float rollout_max = isv[RL_ROLLOUT_MAX_INDEX]; float target = prev * scale; - target = fmaxf(ROLLOUT_MIN, fminf(target, ROLLOUT_MAX)); + target = fmaxf(rollout_min, fminf(target, rollout_max)); // First-observation replace-directly per // `pearl_first_observation_bootstrap`. @@ -137,9 +166,9 @@ extern "C" __global__ void rl_rollout_steps_controller( } // Wiener-α blend with floor per pearl_wiener_alpha_floor_for_nonstationary. - const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR); + const float a = fmaxf(alpha, isv[RL_WIENER_ALPHA_FLOOR_INDEX]); float out = (1.0f - a) * prev + a * target; - out = fmaxf(ROLLOUT_MIN, fminf(out, ROLLOUT_MAX)); + out = fmaxf(rollout_min, fminf(out, rollout_max)); isv[RL_N_ROLLOUT_STEPS_INDEX] = out; } diff --git a/crates/ml-alpha/cuda/rl_signal_variance_update.cu b/crates/ml-alpha/cuda/rl_signal_variance_update.cu new file mode 100644 index 000000000..578c2e8c4 --- /dev/null +++ b/crates/ml-alpha/cuda/rl_signal_variance_update.cu @@ -0,0 +1,73 @@ +// rl_signal_variance_update.cu — Welford online variance for controller inputs (2026-05-30). +// +// Per `feedback_adaptive_not_tuned`, `feedback_isv_for_adaptive_bounds`, +// `pearl_controller_anchors_isv_driven`: every adaptive controller's noise +// floor / target / threshold must derive from observed signal statistics, +// not from hardcoded constants. This kernel maintains a running mean and +// M² (sum of squared deviations from mean) for each controller-input EMA, +// using Welford's online algorithm for numerical stability with float32. +// +// Welford recurrence (numerically stable, single-pass): +// n_new = n_prev + 1 +// delta = x - mean_prev +// mean = mean_prev + delta / n_new +// delta2 = x - mean // recomputed against NEW mean +// m2 = m2_prev + delta * delta2 +// +// Sample variance = m2 / (n - 1) when n > 1; undefined when n ≤ 1. +// +// Per `pearl_first_observation_bootstrap`: skip the update when input == +// sentinel 0.0. The first non-zero observation initializes mean directly +// (delta against mean_prev=0 produces mean = x for n=1; delta against the +// new mean is 0; m2 stays 0; subsequent observations build real variance). +// +// Per `feedback_cpu_is_read_only`: pure device kernel. ISV is the contract +// surface; no host computation, no host control. Caller threads ISV bus +// pointer + four slot indices (input + count/mean/M² triple). Caller is +// responsible for launching once per controller input per step (usually +// via fused-kernel orchestration). +// +// Per `feedback_no_atomicadd`, `feedback_nvidia_grade_perf_for_kernels`: +// single-thread block. Each Welford triple is owned by exactly one input +// slot, accessed sequentially across steps — no concurrent writers, no +// reduction, no atomic. +// +// Launch config: +// grid = (1, 1, 1) +// block = (1, 1, 1) +// smem = 0 +// stream = main RL stream (sequenced after EMA producers and before +// controllers consume the variance) + +extern "C" __global__ void rl_signal_variance_update( + float* __restrict__ isv, + int input_slot, + int count_slot, + int mean_slot, + int m2_slot +) { + // Single-thread guard. Block shape is (1,1,1) by contract but defend + // against a mis-configured launcher. + if (threadIdx.x != 0 || threadIdx.y != 0 || threadIdx.z != 0) return; + if (blockIdx.x != 0 || blockIdx.y != 0 || blockIdx.z != 0) return; + + // Sentinel-zero skip per pearl_first_observation_bootstrap. + // Producers (EMA kernels) hold their slot at 0.0 until the first + // real observation. Welford should NOT count those zero reads as + // genuine observations of "the signal is zero". + const float x = isv[input_slot]; + if (x == 0.0f) return; + + // Welford online update. + const float n_prev = isv[count_slot]; + const float n_new = n_prev + 1.0f; + const float mean_prev = isv[mean_slot]; + const float delta = x - mean_prev; + const float mean_new = mean_prev + delta / n_new; + const float delta2 = x - mean_new; // against new mean — load-bearing for stability + const float m2_new = isv[m2_slot] + delta * delta2; + + isv[count_slot] = n_new; + isv[mean_slot] = mean_new; + isv[m2_slot] = m2_new; +} diff --git a/crates/ml-alpha/cuda/rl_target_tau_controller.cu b/crates/ml-alpha/cuda/rl_target_tau_controller.cu index 2885cca8a..a01bfc9b2 100644 --- a/crates/ml-alpha/cuda/rl_target_tau_controller.cu +++ b/crates/ml-alpha/cuda/rl_target_tau_controller.cu @@ -26,16 +26,34 @@ // the target-net design exists to provide. #define RL_TARGET_TAU_INDEX 401 -#define TAU_MIN 0.001f -#define RL_TARGET_TAU_MAX_INDEX 573 +// τ MIN clamp bound is now ISV-driven per the 2026-05-30 clamp-bound +// extension. ISV-residence makes the floor runtime-tunable and aligns +// with the broader principle that every controller bound is observed- +// signal driven rather than calibrated against a prior regime. +#define RL_TARGET_TAU_MIN_INDEX 642 +#define RL_TARGET_TAU_MAX_INDEX 573 // ISV-driven bootstrap + target. #define RL_TAU_BOOTSTRAP_INDEX 473 #define RL_DIV_TARGET_INDEX 457 // Schulman params from shared global slots (same as ppo_clip). #define RL_SCHULMAN_TOLERANCE_INDEX 468 #define RL_SCHULMAN_ADJUST_RATE_INDEX 469 -#define DIV_NOISE_FLOOR_FRAC 0.01f -#define WIENER_ALPHA_FLOOR 0.4f +// Wiener-α floor — shared across 9 controllers (slot 659). +#define RL_WIENER_ALPHA_FLOOR_INDEX 659 + +// Adaptive controller floors (spec 2026-05-30-adaptive-controller-floor-design): +// noise floor derives from observed q_divergence_ema variance via Welford +// triples, asymmetric Schulman widening requires N consecutive below-band +// observations. Replaces hardcoded DIV_NOISE_FLOOR_FRAC=0.01f which was +// calibrated against the pre-Phase-4.5 divergence regime — under Phase 4.5 +// the 1%-of-target floor is ~100× too low and the controller never moved +// from bootstrap (alpha-rl-... fold 0 confirmed τ stuck at 0.005). +#define RL_Q_DIV_VAR_COUNT_INDEX 592 +#define RL_Q_DIV_VAR_M2_INDEX 594 +#define RL_Q_DIV_BELOW_COUNT_INDEX 595 +#define NOISE_FLOOR_TARGET_FRAC 0.5f // floor ≥ 50% of target +#define NOISE_FLOOR_STD_MULTIPLIER 2.0f // floor ≥ 2σ of observed signal +#define WIDEN_PATIENCE_CONSECUTIVE 3.0f // widen requires N below-band steps // ───────────────────────────────────────────────────────────────────── @@ -77,25 +95,50 @@ extern "C" __global__ void rl_target_tau_controller( // ISV-driven divergence target (was hardcoded #define). const float div_target = isv[RL_DIV_TARGET_INDEX]; - const float div_noise_floor = div_target * DIV_NOISE_FLOOR_FRAC; - // Noise-floor gate. + // Adaptive noise floor — signal-driven per + // `pearl_zscore_normalization_for_magnitude_asymmetric_signals` and + // `feedback_adaptive_not_tuned`. Floor must scale with observed signal + // magnitude AND have an absolute minimum tied to the target so a + // momentarily-noisy signal can't trigger spurious adjustments. + // Welford sample variance = M² / (count − 1) when count > 1. + const float div_count = isv[RL_Q_DIV_VAR_COUNT_INDEX]; + const float div_var = (div_count > 1.0f) + ? isv[RL_Q_DIV_VAR_M2_INDEX] / (div_count - 1.0f) + : 0.0f; + const float div_std = sqrtf(div_var); + const float div_noise_floor = fmaxf(div_target * NOISE_FLOOR_TARGET_FRAC, + div_std * NOISE_FLOOR_STD_MULTIPLIER); + + // Noise-floor gate: divergence below the adaptive floor is dominated by + // numerical noise OR signal stays naturally below target under the + // current architecture. Hold τ. if (q_divergence_norm < div_noise_floor) return; - // Bounded multiplicative adjustment. - float ratio; + // Asymmetric Schulman (spec 2026-05-30 Section "Approach C, folded in"): + // raising τ fires on a single above-band observation — real Q + // divergence is a safety signal we act on fast (target net falling + // behind = raise the bootstrap rate). Lowering τ requires N consecutive + // below-band observations so a single noisy step can't push τ toward + // the floor. Below-band counter is per-controller in ISV. const float tolerance = isv[RL_SCHULMAN_TOLERANCE_INDEX]; const float adjust_rate = isv[RL_SCHULMAN_ADJUST_RATE_INDEX]; + float ratio; if (q_divergence_norm > div_target * tolerance) { - ratio = adjust_rate; + ratio = adjust_rate; // raise τ — single observation + isv[RL_Q_DIV_BELOW_COUNT_INDEX] = 0.0f; // reset patience } else if (q_divergence_norm < div_target / tolerance) { - ratio = 1.0f / adjust_rate; + const float new_count = isv[RL_Q_DIV_BELOW_COUNT_INDEX] + 1.0f; + isv[RL_Q_DIV_BELOW_COUNT_INDEX] = new_count; + ratio = (new_count >= WIDEN_PATIENCE_CONSECUTIVE) ? (1.0f / adjust_rate) : 1.0f; } else { + isv[RL_Q_DIV_BELOW_COUNT_INDEX] = 0.0f; // in-band → reset ratio = 1.0f; } float tau_target = tau_prev * ratio; + const float tau_min = isv[RL_TARGET_TAU_MIN_INDEX]; const float tau_max = isv[RL_TARGET_TAU_MAX_INDEX]; - tau_target = fmaxf(TAU_MIN, fminf(tau_target, tau_max)); + tau_target = fmaxf(tau_min, fminf(tau_target, tau_max)); // First-observation replace-directly: if // prev is still exactly the hardcoded bootstrap value, this is @@ -112,9 +155,9 @@ extern "C" __global__ void rl_target_tau_controller( } // Wiener-α blend with floor per pearl_wiener_alpha_floor_for_nonstationary. - const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR); + const float a = fmaxf(alpha, isv[RL_WIENER_ALPHA_FLOOR_INDEX]); float tau_new = (1.0f - a) * tau_prev + a * tau_target; - tau_new = fmaxf(TAU_MIN, fminf(tau_new, tau_max)); + tau_new = fmaxf(tau_min, fminf(tau_new, tau_max)); isv[RL_TARGET_TAU_INDEX] = tau_new; } diff --git a/crates/ml-alpha/cuda/rl_v_blend_alpha_controller.cu b/crates/ml-alpha/cuda/rl_v_blend_alpha_controller.cu index a6e6e3d91..34e323161 100644 --- a/crates/ml-alpha/cuda/rl_v_blend_alpha_controller.cu +++ b/crates/ml-alpha/cuda/rl_v_blend_alpha_controller.cu @@ -20,22 +20,38 @@ // does, controller re-bootstraps harmlessly. // // Per pearl_blend_formulas_must_have_permanent_floor: dead-signal -// guard — if EMA(|V_scalar|) < FLOOR, hold α (no V signal to -// calibrate against; the trainer hasn't seen meaningful rewards yet). +// guard — if EMA(|V_scalar|) < adaptive floor, hold α (no V signal +// to calibrate against; the trainer hasn't seen meaningful rewards +// yet). // // Per feedback_cpu_is_read_only: pure device kernel; reads V_scalar, // V_dq, ISV; emits α + EMAs to ISV. No host control. // +// Adaptive controller floors (spec 2026-05-30-adaptive-controller-floor-design): +// the 5 design constants below become ISV-driven so they can be tuned +// at runtime without recompile. The dead-signal floor additionally +// adapts to observed |V_scalar| variance via Welford slots 626/627/628 +// (Welford updates wired by the trainer pass alongside the other +// controller Welford triples). +// // Block layout: grid=(1, 1, 1), block=(BLOCK_X=1024, 1, 1). Single // block does parallel reduction over batch up to B=1024. Thread 0 // performs the controller update. #define BLOCK_X 1024 -#define EMA_ALPHA 0.01f -#define TARGET_TRACK_RATIO 0.10f -#define SCHULMAN_STEP 0.01f -#define DEAD_SIGNAL_FLOOR 1e-4f -#define BOOTSTRAP_ALPHA 1.0f +// ISV-driven slots — replace prior hardcoded design constants per the +// 2026-05-30 adaptive-controller-floor spec. +#define RL_V_BLEND_DEAD_SIGNAL_FLOOR_INDEX 621 +#define RL_V_BLEND_TARGET_TRACK_RATIO_INDEX 622 +#define RL_V_BLEND_SCHULMAN_STEP_INDEX 623 +#define RL_V_BLEND_EMA_ALPHA_INDEX 624 +#define RL_V_BLEND_BOOTSTRAP_ALPHA_INDEX 625 +#define RL_V_BLEND_SCALAR_VAR_COUNT_INDEX 626 +#define RL_V_BLEND_SCALAR_VAR_M2_INDEX 628 +// Adaptive-dead-floor scaling for the Welford std term — keeps the +// floor at least one σ_observed × ADAPTIVE_DEAD_FLOOR_STD_SCALE above +// numerical noise even when the ISV-resident absolute floor is tighter. +#define ADAPTIVE_DEAD_FLOOR_STD_SCALE 0.1f extern "C" __global__ void rl_v_blend_alpha_controller( const float* __restrict__ v_scalar, // [B] @@ -79,10 +95,33 @@ extern "C" __global__ void rl_v_blend_alpha_controller( const float track_mean = s_t[0] * inv_B; const float mag_mean = s_m[0] * inv_B; + // ── ISV-driven design constants (spec 2026-05-30) ── + const float ema_alpha = isv[RL_V_BLEND_EMA_ALPHA_INDEX]; + const float target_track_ratio = isv[RL_V_BLEND_TARGET_TRACK_RATIO_INDEX]; + const float schulman_step = isv[RL_V_BLEND_SCHULMAN_STEP_INDEX]; + const float bootstrap_alpha = isv[RL_V_BLEND_BOOTSTRAP_ALPHA_INDEX]; + + // ── Adaptive dead-signal floor ───────────────────────────────── + // Combines the ISV-resident absolute floor with an observed + // |V_scalar| variance term (Welford sample variance = + // M² / (count − 1) when count > 1). Cold-start safe: count=0 + // → var=0 → adaptive = isv[621] absolute default. Once Welford + // accumulates observations, the floor scales with σ_observed so + // the dead-signal gate adapts to whatever magnitude regime + // |V_scalar| settles into post-Phase-4.5. + const float vmag_count = isv[RL_V_BLEND_SCALAR_VAR_COUNT_INDEX]; + const float vmag_var = (vmag_count > 1.0f) + ? isv[RL_V_BLEND_SCALAR_VAR_M2_INDEX] / (vmag_count - 1.0f) + : 0.0f; + const float vmag_std = sqrtf(vmag_var); + const float dead_floor_abs = isv[RL_V_BLEND_DEAD_SIGNAL_FLOOR_INDEX]; + const float adaptive_dead_floor = fmaxf(dead_floor_abs, + vmag_std * ADAPTIVE_DEAD_FLOOR_STD_SCALE); + // ── Bootstrap α on sentinel ── float alpha = isv[alpha_slot]; if (alpha == 0.0f) { - alpha = BOOTSTRAP_ALPHA; + alpha = bootstrap_alpha; } // ── First-observation bootstrap on EMAs ── @@ -90,13 +129,13 @@ extern "C" __global__ void rl_v_blend_alpha_controller( float prev_mag_ema = isv[v_scalar_mag_ema_slot]; const float track_ema = (prev_track_ema == 0.0f) ? track_mean - : (1.0f - EMA_ALPHA) * prev_track_ema + EMA_ALPHA * track_mean; + : (1.0f - ema_alpha) * prev_track_ema + ema_alpha * track_mean; const float mag_ema = (prev_mag_ema == 0.0f) ? mag_mean - : (1.0f - EMA_ALPHA) * prev_mag_ema + EMA_ALPHA * mag_mean; + : (1.0f - ema_alpha) * prev_mag_ema + ema_alpha * mag_mean; // ── Dead-signal guard: V_scalar magnitude too small to calibrate against ── - if (mag_ema < DEAD_SIGNAL_FLOOR) { + if (mag_ema < adaptive_dead_floor) { isv[trackerr_ema_slot] = track_ema; isv[v_scalar_mag_ema_slot] = mag_ema; isv[alpha_slot] = alpha; // hold (write bootstrap if needed) @@ -105,12 +144,12 @@ extern "C" __global__ void rl_v_blend_alpha_controller( // ── Track ratio + Schulman-bounded step on α ── const float track_ratio = track_ema / mag_ema; - if (track_ratio > 1.5f * TARGET_TRACK_RATIO) { + if (track_ratio > 1.5f * target_track_ratio) { // V_dq diverged from V_scalar → raise α toward V_scalar - alpha = fminf(alpha + SCHULMAN_STEP, 1.0f); - } else if (track_ratio < TARGET_TRACK_RATIO / 1.5f) { + alpha = fminf(alpha + schulman_step, 1.0f); + } else if (track_ratio < target_track_ratio / 1.5f) { // V_dq tracks V_scalar well → lower α toward V_dq - alpha = fmaxf(alpha - SCHULMAN_STEP, 0.0f); + alpha = fmaxf(alpha - schulman_step, 0.0f); } // else: hold α (within band) diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index 2efb7450b..456b5c4b6 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -1131,5 +1131,209 @@ pub const RL_V_TRACK_ERR_EMA_INDEX: usize = 586; /// If `EMA < 1e-4` the controller holds α (no V signal to calibrate against). pub const RL_V_SCALAR_MAG_EMA_INDEX: usize = 587; -/// Last RL-allocated slot index (exclusive). -pub const RL_SLOTS_END: usize = 588; +// ============================================================ +// Adaptive controller floors — spec 2026-05-30-adaptive-controller-floor-design +// +// Replaces hardcoded thresholds across 12 RL controllers with observed- +// signal-driven ISV-slot bounds. Per `feedback_adaptive_not_tuned` and +// `pearl_controller_anchors_isv_driven`: every threshold derives from +// observed signal statistics (Welford online variance), not from a +// constant calibrated against a prior signal regime. +// +// Group A (8 controllers that saturated under Phase 4.5): +// - Welford variance triple (count, mean, M²) per controller's input EMA +// - Below-band counter for asymmetric Schulman (5 of 8 use it) +// +// Group B (4 not-yet-saturated controllers with hardcoded thresholds): +// - Direct ISV-slot replacements for the hardcoded values +// +// All slots are written by `rl_signal_variance_update` kernel or by +// controllers themselves; consumed by individual controller kernels and +// by the fused rl_fused_controllers kernel. +// ============================================================ + +// ─── Group A — Welford variance triples + asymmetric-Schulman counters ─── + +/// PPO clip controller — Welford count for kl_pi_ema (`RL_KL_PI_EMA_INDEX`). +pub const RL_KL_PI_VAR_COUNT_INDEX: usize = 588; +/// PPO clip controller — Welford running mean for kl_pi_ema. +pub const RL_KL_PI_VAR_MEAN_INDEX: usize = 589; +/// PPO clip controller — Welford M² (sum of squared deviations from mean). +/// Sample variance = M² / (count - 1). +pub const RL_KL_PI_VAR_M2_INDEX: usize = 590; +/// PPO clip controller — N-consecutive-below-band counter for asymmetric Schulman. +/// Tightening fires on single observation; widening requires 3 consecutive. +pub const RL_KL_PI_BELOW_COUNT_INDEX: usize = 591; + +/// target_tau controller — Welford count for q_divergence_ema. +pub const RL_Q_DIV_VAR_COUNT_INDEX: usize = 592; +pub const RL_Q_DIV_VAR_MEAN_INDEX: usize = 593; +pub const RL_Q_DIV_VAR_M2_INDEX: usize = 594; +pub const RL_Q_DIV_BELOW_COUNT_INDEX: usize = 595; + +/// rollout_steps controller — Welford for pre-normalization advantage variance +/// (`RL_ADV_VAR_PRE_NORM_INDEX`). Post-normalization variance is definitionally +/// near zero under Phase 4.5 and useless as a controller signal. +pub const RL_ADV_VAR_VAR_COUNT_INDEX: usize = 596; +pub const RL_ADV_VAR_VAR_MEAN_INDEX: usize = 597; +pub const RL_ADV_VAR_VAR_M2_INDEX: usize = 598; +pub const RL_ADV_VAR_BELOW_COUNT_INDEX: usize = 599; + +/// entropy_coef controller — Welford for entropy_observed_ema. +pub const RL_ENTROPY_OBS_VAR_COUNT_INDEX: usize = 600; +pub const RL_ENTROPY_OBS_VAR_MEAN_INDEX: usize = 601; +pub const RL_ENTROPY_OBS_VAR_M2_INDEX: usize = 602; +pub const RL_ENTROPY_OBS_BELOW_COUNT_INDEX: usize = 603; + +/// per_α controller — Welford for td_kurtosis_ema. +pub const RL_TD_KURT_VAR_COUNT_INDEX: usize = 604; +pub const RL_TD_KURT_VAR_MEAN_INDEX: usize = 605; +pub const RL_TD_KURT_VAR_M2_INDEX: usize = 606; +pub const RL_TD_KURT_BELOW_COUNT_INDEX: usize = 607; + +/// gamma controller — Welford for mean_trade_duration_ema. Mean is used directly +/// by `rl_gamma_controller` to compute `RL_GAMMA_MIN_ADAPTIVE_INDEX` after a +/// 100-observation warmup (spec Q6 resolution: Welford mean's natural N-smoothing +/// breaks the gamma↔trade_duration feedback loop without explicit step gating). +pub const RL_TRADE_DUR_VAR_COUNT_INDEX: usize = 608; +pub const RL_TRADE_DUR_VAR_MEAN_INDEX: usize = 609; +pub const RL_TRADE_DUR_VAR_M2_INDEX: usize = 610; + +/// q_distill_lambda controller — placeholder for asymmetric counter (currently +/// unused since q_distill uses ramp/decay rather than Schulman; reserved for future). +pub const RL_RESERVED_611_INDEX: usize = 611; + +// ─── New input-signal slots (consumed by controllers + Welford kernel) ─── + +/// Emitted by `rl_advantage_normalize` BEFORE in-place normalize. Carries the raw +/// variance of advantages across the batch — used by `rl_rollout_steps_controller` +/// as the input signal that Phase 4.5 no longer provides via post-norm advantages. +pub const RL_ADV_VAR_PRE_NORM_INDEX: usize = 612; + +/// Adaptive lower bound for gamma, derived from `RL_TRADE_DUR_VAR_MEAN_INDEX` +/// (Welford mean of trade duration). See `rl_gamma_controller.cu` Special case G. +/// Replaces hardcoded `GAMMA_MIN = 0.995f` with signal-driven bound: +/// `max(0.9, 1 - 1/(d_mean × 2))`. +pub const RL_GAMMA_MIN_ADAPTIVE_INDEX: usize = 613; + +/// Emitted by `rl_reward_scale_controller` — observed per-batch pnl magnitude +/// EMA. Input to the Welford kernel; the Welford variance becomes the noise-floor +/// reference for reward_scale's asymmetric rate limit. +pub const RL_REWARD_MAGNITUDE_EMA_INDEX: usize = 614; + +/// reward_scale controller — Welford triple for pnl magnitude. +pub const RL_REWARD_MAG_VAR_COUNT_INDEX: usize = 615; +pub const RL_REWARD_MAG_VAR_MEAN_INDEX: usize = 616; +pub const RL_REWARD_MAG_VAR_M2_INDEX: usize = 617; + +/// q_distill_lambda controller — Welford triple for q_distill_kl_ema. +/// Adaptive `MIN_LAMBDA` derives from `sqrt(var) × 0.05`. +pub const RL_Q_DISTILL_KL_VAR_COUNT_INDEX: usize = 618; +pub const RL_Q_DISTILL_KL_VAR_MEAN_INDEX: usize = 619; +pub const RL_Q_DISTILL_KL_VAR_M2_INDEX: usize = 620; + +// ─── Group B — Adaptive bound slots replacing Phase 4.4/4.5-era hardcoded constants ─── + +/// V-blend alpha controller (Phase 4.4) — adaptive dead-signal floor (replaces `1e-4f`). +pub const RL_V_BLEND_DEAD_SIGNAL_FLOOR_INDEX: usize = 621; +/// V-blend alpha controller — target track ratio (replaces `0.10f`). +pub const RL_V_BLEND_TARGET_TRACK_RATIO_INDEX: usize = 622; +/// V-blend alpha controller — Schulman step magnitude (replaces `0.01f`). +pub const RL_V_BLEND_SCHULMAN_STEP_INDEX: usize = 623; +/// V-blend alpha controller — EMA smoothing alpha (replaces `0.01f`). +pub const RL_V_BLEND_EMA_ALPHA_INDEX: usize = 624; +/// V-blend alpha controller — bootstrap alpha on sentinel input (replaces `1.0f`). +pub const RL_V_BLEND_BOOTSTRAP_ALPHA_INDEX: usize = 625; +/// V-blend alpha controller — Welford count for V_scalar magnitude. +pub const RL_V_BLEND_SCALAR_VAR_COUNT_INDEX: usize = 626; +pub const RL_V_BLEND_SCALAR_VAR_MEAN_INDEX: usize = 627; +pub const RL_V_BLEND_SCALAR_VAR_M2_INDEX: usize = 628; + +/// PPO ratio clamp controller — adaptive max output (replaces `1000.0f`). +/// Min derives directly from observed ratio variance: `max(1.5, 1 + 3·σ_ratio)`. +pub const RL_PPO_RATIO_CLAMP_MAX_ADAPTIVE_INDEX: usize = 629; +/// PPO ratio clamp controller — Welford triple for observed log-ratio magnitudes. +pub const RL_PPO_RATIO_VAR_COUNT_INDEX: usize = 630; +pub const RL_PPO_RATIO_VAR_MEAN_INDEX: usize = 631; +pub const RL_PPO_RATIO_VAR_M2_INDEX: usize = 632; + +/// Reward clamp controller — V-bound floor (replaces hardcoded `V_BOUND_FLOOR = 1.0f`). +pub const RL_REWARD_CLAMP_V_BOUND_FLOOR_INDEX: usize = 633; +/// Reward clamp controller — V-bound EWMA alpha (replaces `0.001f`). +pub const RL_REWARD_CLAMP_V_BOUND_EWMA_ALPHA_INDEX: usize = 634; +/// Reward clamp controller — minimum win clip magnitude (replaces `MIN_WIN = 1.0f`). +pub const RL_REWARD_CLAMP_MIN_WIN_INDEX: usize = 635; +/// Reward clamp controller — minimum win/loss ratio (replaces `MIN_RATIO = 1.0f`). +pub const RL_REWARD_CLAMP_MIN_RATIO_INDEX: usize = 636; +/// Reward clamp controller — maximum win/loss ratio (replaces `MAX_RATIO = 3.0f`). +pub const RL_REWARD_CLAMP_MAX_RATIO_INDEX: usize = 637; + +/// Gate threshold controller — EMA smoothing alpha (replaces hardcoded `alpha = 0.01f`). +pub const RL_GATE_EMA_ALPHA_INDEX: usize = 638; + +/// q_distill_lambda controller — adaptive MIN bound (replaces hardcoded `MIN_LAMBDA = 0.05f`). +/// Derives from `sqrt(q_distill_kl_var) × 0.05` with absolute floor 0.001. +pub const RL_Q_DISTILL_LAMBDA_MIN_ADAPTIVE_INDEX: usize = 639; + +// ============================================================ +// Extension (2026-05-30, same spec): all hardcoded MIN/MAX clamp bounds + +// ramp/decay rates across the 12 controllers become ISV-driven. Same +// `feedback_adaptive_not_tuned` principle applied consistently — clamp +// bounds calibrated against a prior signal regime are the next layer of +// the same architectural bug that saturated the 8 noise-floor controllers +// in fold 0/1. Bootstrap defaults preserve current behavior. +// ============================================================ + +// ─── PPO clip clamp bounds (rl_ppo_clip_controller) ─── +pub const RL_PPO_CLIP_EPS_MIN_INDEX: usize = 640; // replaces EPS_MIN = 0.05f +pub const RL_PPO_CLIP_EPS_MAX_INDEX: usize = 641; // replaces EPS_MAX = 0.5f + +// ─── target_tau clamp bound (rl_target_tau_controller) ─── +pub const RL_TARGET_TAU_MIN_INDEX: usize = 642; // replaces TAU_MIN = 0.001f + +// ─── rollout_steps clamp bounds (rl_rollout_steps_controller) ─── +pub const RL_ROLLOUT_MIN_INDEX: usize = 643; // replaces ROLLOUT_MIN = 256.0f +pub const RL_ROLLOUT_MAX_INDEX: usize = 644; // replaces ROLLOUT_MAX = 8192.0f + +// ─── entropy_coef clamp bounds (rl_entropy_coef_controller) ─── +pub const RL_ENTROPY_COEF_MIN_INDEX: usize = 645; // replaces COEF_MIN = 0.0f +pub const RL_ENTROPY_COEF_MAX_INDEX: usize = 646; // replaces COEF_MAX = 0.5f + +// ─── per_α clamp bounds (rl_per_alpha_controller) ─── +pub const RL_PER_ALPHA_MIN_INDEX: usize = 647; // replaces PER_ALPHA_MIN = 0.3f +pub const RL_PER_ALPHA_MAX_INDEX: usize = 648; // replaces PER_ALPHA_MAX = 1.0f + +// ─── gamma clamp ceiling (rl_gamma_controller) ─── +// GAMMA_MIN is already adaptive via RL_GAMMA_MIN_ADAPTIVE_INDEX=613. +pub const RL_GAMMA_MAX_INDEX: usize = 649; // replaces GAMMA_MAX = 0.999f + +// ─── q_distill_lambda clamp / rate constants (rl_q_distill_lambda_controller) ─── +// MIN_LAMBDA is already adaptive via RL_Q_DISTILL_LAMBDA_MIN_ADAPTIVE_INDEX=639. +pub const RL_Q_DISTILL_LAMBDA_MAX_INDEX: usize = 650; // replaces MAX_LAMBDA = 1.0f +pub const RL_Q_DISTILL_KL_TOLERANCE_INDEX: usize = 651; // replaces KL_TOLERANCE = 3.0f +pub const RL_Q_DISTILL_LAMBDA_RAMP_RATE_INDEX: usize = 652; // replaces LAMBDA_RAMP_RATE = 1.2f +pub const RL_Q_DISTILL_LAMBDA_DECAY_RATE_INDEX: usize = 653; // replaces LAMBDA_DECAY_RATE = 0.998f + +// ─── ppo_ratio_clamp ceiling (rl_ppo_ratio_clamp_controller) ─── +// MIN_OUT becomes adaptive from observed ratio variance (per Group B Task 11). +// MAX_OUT becomes ISV-driven adaptive max via RL_PPO_RATIO_CLAMP_MAX_ADAPTIVE_INDEX=629. +// (Hardcoded MIN_OUT=2.0 absolute floor still used as `max(adaptive, 2.0)`; this is +// the "don't degenerate to vanilla policy gradient" structural guard, treated as +// architectural per the spec's "physics/math constants" exemption.) + +// ─── reward_clamp adaptive bounds (rl_reward_clamp_controller) ─── +// V_BOUND_FLOOR, V_BOUND_EWMA_ALPHA, MIN_WIN, MIN_RATIO, MAX_RATIO already in 633-637. +pub const RL_REWARD_CLAMP_MIN_MARGIN_INDEX: usize = 654; // replaces MIN_MARGIN = 1.0f +pub const RL_REWARD_CLAMP_MAX_MARGIN_INDEX: usize = 655; // replaces MAX_MARGIN = 5.0f +pub const RL_REWARD_CLAMP_MARGIN_TOLERANCE_INDEX: usize = 656; // replaces MARGIN_TOLERANCE = 1.5f +pub const RL_REWARD_CLAMP_MARGIN_ADJUST_RATE_INDEX: usize = 657; // replaces MARGIN_ADJUST_RATE = 1.2f +pub const RL_REWARD_CLAMP_CLIP_RATE_EMA_ALPHA_INDEX: usize = 658; // replaces CLIP_RATE_EMA_ALPHA = 0.05f + +// ─── Wiener-α floor (shared across 9 controllers — `pearl_wiener_alpha_floor_for_nonstationary`) ─── +// The 0.4f floor is mathematically motivated but still a tuned constant. +// Single shared ISV slot — all controllers read from this slot for consistency. +pub const RL_WIENER_ALPHA_FLOOR_INDEX: usize = 659; // replaces WIENER_ALPHA_FLOOR = 0.4f + +/// Last RL-allocated slot index (exclusive). Pre-spec: 588. Post-noise-floor: 640. +/// Post-clamp-bounds extension: 660. Total new: 72 slots. +pub const RL_SLOTS_END: usize = 660; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index c59bcb198..cca7222fb 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -198,6 +198,11 @@ const RL_V_BLEND_ALPHA_CONTROLLER_CUBIN: &[u8] = /// Phase 4.5 per-batch advantage normalization. const RL_ADVANTAGE_NORMALIZE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_advantage_normalize.cubin")); +/// Welford online variance updater for adaptive controller-input floors +/// (spec `2026-05-30-adaptive-controller-floor-design`). Single-thread, +/// single-block. Caller passes ISV bus + input/count/mean/M² slot indices. +const RL_SIGNAL_VARIANCE_UPDATE_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_signal_variance_update.cubin")); const RL_Q_PI_DISTILL_GRAD_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_q_pi_distill_grad.cubin")); // λ_distill adaptive controller (rljzl followup 2026-05-24). @@ -637,6 +642,12 @@ pub struct IntegratedTrainer { rl_v_blend_alpha_controller_fn: CudaFunction, _rl_advantage_normalize_module: Arc, rl_advantage_normalize_fn: CudaFunction, + // Welford variance updater for adaptive controller floors (spec + // 2026-05-30-adaptive-controller-floor-design). One launch per + // controller-input EMA per step, sequenced after EMA producers and + // before controller consumers. + _rl_signal_variance_update_module: Arc, + rl_signal_variance_update_fn: CudaFunction, /// Phase 4.4 blended V baselines fed to compute_advantage_return. pub v_blended_d: CudaSlice, pub v_blended_tp1_d: CudaSlice, @@ -1504,6 +1515,14 @@ impl IntegratedTrainer { let rl_advantage_normalize_fn = rl_advantage_normalize_module .load_function("rl_advantage_normalize") .context("load rl_advantage_normalize")?; + // Welford online variance updater for adaptive controller-input + // floors (spec 2026-05-30-adaptive-controller-floor-design). + let rl_signal_variance_update_module = ctx + .load_cubin(RL_SIGNAL_VARIANCE_UPDATE_CUBIN.to_vec()) + .context("load rl_signal_variance_update cubin")?; + let rl_signal_variance_update_fn = rl_signal_variance_update_module + .load_function("rl_signal_variance_update") + .context("load rl_signal_variance_update")?; let rl_q_distill_lambda_controller_module = ctx .load_cubin(RL_Q_DISTILL_LAMBDA_CONTROLLER_CUBIN.to_vec()) .context("load rl_q_distill_lambda_controller cubin")?; @@ -1573,7 +1592,7 @@ impl IntegratedTrainer { crate::rl::isv_slots::RL_Q_DIVERGENCE_EMA_INDEX as i32, crate::rl::isv_slots::RL_KL_PI_EMA_INDEX as i32, crate::rl::isv_slots::RL_ENTROPY_OBSERVED_EMA_INDEX as i32, - crate::rl::isv_slots::RL_ADVANTAGE_VAR_RATIO_EMA_INDEX as i32, + crate::rl::isv_slots::RL_ADV_VAR_PRE_NORM_INDEX as i32, crate::rl::isv_slots::RL_TD_KURTOSIS_EMA_INDEX as i32, crate::rl::isv_slots::RL_MEAN_ABS_PNL_EMA_INDEX as i32, ]; @@ -2536,6 +2555,8 @@ impl IntegratedTrainer { rl_v_blend_alpha_controller_fn, _rl_advantage_normalize_module: rl_advantage_normalize_module, rl_advantage_normalize_fn, + _rl_signal_variance_update_module: rl_signal_variance_update_module, + rl_signal_variance_update_fn, v_blended_d, v_blended_tp1_d, _rl_q_distill_lambda_controller_module: rl_q_distill_lambda_controller_module, @@ -2886,7 +2907,7 @@ impl IntegratedTrainer { // (slot, value) pair — pure device write, no HtoD per // `feedback_no_htod_htoh_only_mapped_pinned`. { - let isv_constants: [(usize, f32); 111] = [ + let isv_constants: [(usize, f32); 145] = [ // Static seeds for the adaptive reward-clamp controller — // these are the initial values that // `rl_reward_clamp_controller` will replace once it @@ -3008,10 +3029,28 @@ impl IntegratedTrainer { (crate::rl::isv_slots::RL_KL_REF_BETA_INDEX, 0.0005), (crate::rl::isv_slots::RL_REWARD_KL_BETA_INDEX, 0.001), (crate::rl::isv_slots::RL_KL_REF_TARGET_INDEX, 0.5), - (crate::rl::isv_slots::RL_EPS_BOOTSTRAP_INDEX, 0.2), + // Adaptive controller floors (spec 2026-05-30): bootstrap ε at + // EPS_MIN = 0.05 (the lowest valid clip per the EPS_MIN clamp). + // An earlier draft used 0.01 (= KL target) but that's BELOW the + // EPS_MIN clamp range, so the post-bootstrap step always snapped + // ε to MIN regardless of signal direction — the bootstrap-clamp + // inconsistency Task 16 surfaced. Matching bootstrap to MIN + // keeps the first post-bootstrap step honest: asymmetric Schulman + // widens to track observed signal, tightens get clamped at MIN + // (no-op, controller stays at the floor). Conservative + consistent. + (crate::rl::isv_slots::RL_EPS_BOOTSTRAP_INDEX, 0.05), (crate::rl::isv_slots::RL_ROLLOUT_BOOTSTRAP_INDEX, 2048.0), (crate::rl::isv_slots::RL_REWARD_SCALE_BOOTSTRAP_INDEX, 1.0), (crate::rl::isv_slots::RL_PPO_RATIO_CLAMP_BOOTSTRAP_INDEX, 10.0), + // Adaptive controller floors (spec 2026-05-30 Special cases G + Q): + // safe defaults before the gamma controller / q_distill controller + // overwrite these on their first per-step fire. γ_min 0.95 + // straddles short-trade (0.968) and surfer-regime (0.995) defaults; + // λ_distill MIN 0.01 matches the q_distill λ bootstrap so the + // hardcoded MIN_LAMBDA = 0.05 floor is replaced without a sudden + // jump in λ at controller bootstrap. + (crate::rl::isv_slots::RL_GAMMA_MIN_ADAPTIVE_INDEX, 0.95), + (crate::rl::isv_slots::RL_Q_DISTILL_LAMBDA_MIN_ADAPTIVE_INDEX, 0.01), // IQN ensemble seeds — alpha=0.5 (equal C51/IQN blend), // N_TAU=32, LR=1e-3 (canonical). (crate::rl::isv_slots::RL_IQN_N_TAU_INDEX, 32.0), @@ -3053,6 +3092,61 @@ impl IntegratedTrainer { (crate::rl::isv_slots::RL_SAC_ALPHA_INDEX, 0.01), (crate::rl::isv_slots::RL_SAC_ENTROPY_TARGET_INDEX, 1.68), (crate::rl::isv_slots::RL_CONF_GATE_MAX_HOLD_FRAC_INDEX, 0.85), + // ─── Adaptive controller floors — Group B + clamp-bound expansion + // (spec 2026-05-30-adaptive-controller-floor-design). Every value + // preserves the prior hardcoded constant exactly so the runtime + // behaviour at bootstrap matches pre-spec behaviour bit-for-bit; + // the win is that all bounds are now ISV-resident (visible in + // diag JSONL, re-tunable via re-seed instead of recompile) and + // can be replaced by Welford-driven adaptive logic in follow-up + // controllers without further code changes. + // + // V-blend (Phase 4.4) — 5 design constants previously #define'd + // in rl_v_blend_alpha_controller.cu. + (crate::rl::isv_slots::RL_V_BLEND_DEAD_SIGNAL_FLOOR_INDEX, 1e-4), + (crate::rl::isv_slots::RL_V_BLEND_TARGET_TRACK_RATIO_INDEX, 0.10), + (crate::rl::isv_slots::RL_V_BLEND_SCHULMAN_STEP_INDEX, 0.01), + (crate::rl::isv_slots::RL_V_BLEND_EMA_ALPHA_INDEX, 0.01), + (crate::rl::isv_slots::RL_V_BLEND_BOOTSTRAP_ALPHA_INDEX, 1.0), + // PPO ratio clamp — adaptive MAX ceiling (was hardcoded 1000.0f). + (crate::rl::isv_slots::RL_PPO_RATIO_CLAMP_MAX_ADAPTIVE_INDEX, 1000.0), + // Reward clamp — design constants previously #define'd in + // rl_reward_clamp_controller.cu. Defaults match the historic + // V_BOUND_FLOOR=1.0, V_BOUND_EWMA_ALPHA=0.001, MIN_WIN=1.0, + // MIN_RATIO=1.0, MAX_RATIO=3.0, MIN_MARGIN=1.0, MAX_MARGIN=5.0, + // MARGIN_TOLERANCE=1.5, MARGIN_ADJUST_RATE=1.2, + // CLIP_RATE_EMA_ALPHA=0.05. + (crate::rl::isv_slots::RL_REWARD_CLAMP_V_BOUND_FLOOR_INDEX, 1.0), + (crate::rl::isv_slots::RL_REWARD_CLAMP_V_BOUND_EWMA_ALPHA_INDEX, 0.001), + (crate::rl::isv_slots::RL_REWARD_CLAMP_MIN_WIN_INDEX, 1.0), + (crate::rl::isv_slots::RL_REWARD_CLAMP_MIN_RATIO_INDEX, 1.0), + (crate::rl::isv_slots::RL_REWARD_CLAMP_MAX_RATIO_INDEX, 3.0), + (crate::rl::isv_slots::RL_REWARD_CLAMP_MIN_MARGIN_INDEX, 1.0), + (crate::rl::isv_slots::RL_REWARD_CLAMP_MAX_MARGIN_INDEX, 5.0), + (crate::rl::isv_slots::RL_REWARD_CLAMP_MARGIN_TOLERANCE_INDEX, 1.5), + (crate::rl::isv_slots::RL_REWARD_CLAMP_MARGIN_ADJUST_RATE_INDEX, 1.2), + (crate::rl::isv_slots::RL_REWARD_CLAMP_CLIP_RATE_EMA_ALPHA_INDEX, 0.05), + // Gate threshold EMA α — was hardcoded 0.01f in + // rl_gate_threshold_controller.cu. + (crate::rl::isv_slots::RL_GATE_EMA_ALPHA_INDEX, 0.01), + // Clamp bounds — Batch 2 of the 2026-05-30 expansion. + // Each value mirrors the prior #define so behaviour is preserved. + (crate::rl::isv_slots::RL_PPO_CLIP_EPS_MIN_INDEX, 0.05), + (crate::rl::isv_slots::RL_PPO_CLIP_EPS_MAX_INDEX, 0.5), + (crate::rl::isv_slots::RL_TARGET_TAU_MIN_INDEX, 0.001), + (crate::rl::isv_slots::RL_ROLLOUT_MIN_INDEX, 256.0), + (crate::rl::isv_slots::RL_ROLLOUT_MAX_INDEX, 8192.0), + (crate::rl::isv_slots::RL_ENTROPY_COEF_MIN_INDEX, 0.0), + (crate::rl::isv_slots::RL_ENTROPY_COEF_MAX_INDEX, 0.5), + (crate::rl::isv_slots::RL_PER_ALPHA_MIN_INDEX, 0.3), + (crate::rl::isv_slots::RL_PER_ALPHA_MAX_INDEX, 1.0), + (crate::rl::isv_slots::RL_GAMMA_MAX_INDEX, 0.999), + (crate::rl::isv_slots::RL_Q_DISTILL_LAMBDA_MAX_INDEX, 1.0), + (crate::rl::isv_slots::RL_Q_DISTILL_KL_TOLERANCE_INDEX, 3.0), + (crate::rl::isv_slots::RL_Q_DISTILL_LAMBDA_RAMP_RATE_INDEX, 1.2), + (crate::rl::isv_slots::RL_Q_DISTILL_LAMBDA_DECAY_RATE_INDEX, 0.998), + // Shared Wiener-α floor — read by 9 controllers (slot 659). + (crate::rl::isv_slots::RL_WIENER_ALPHA_FLOOR_INDEX, 0.4), ]; for (slot, value) in isv_constants.iter() { let slot_i32 = *slot as i32; @@ -3098,10 +3192,14 @@ impl IntegratedTrainer { crate::rl::isv_slots::RL_ENTROPY_OBSERVED_EMA_INDEX as i32, ) .context("R1 bootstrap rl_entropy_coef_controller")?; + // Adaptive controller floors (spec 2026-05-30): rollout_steps now + // consumes pre-normalization advantage variance (RL_ADV_VAR_PRE_NORM_INDEX). + // Post-norm advantage variance is definitionally ≈ 0 under Phase 4.5 + // (per-batch normalize), which left this controller blind. self.launch_isv_controller_3arg( &self.rl_rollout_steps_controller_fn, alpha, - crate::rl::isv_slots::RL_ADVANTAGE_VAR_RATIO_EMA_INDEX as i32, + crate::rl::isv_slots::RL_ADV_VAR_PRE_NORM_INDEX as i32, ) .context("R1 bootstrap rl_rollout_steps_controller")?; self.launch_isv_controller_3arg( @@ -3267,6 +3365,43 @@ impl IntegratedTrainer { Ok(()) } + /// Launch a Welford-variance update for ONE controller-input EMA + /// slot (spec `2026-05-30-adaptive-controller-floor-design`). + /// + /// Single-thread, single-block. The kernel reads `isv[input_slot]`, + /// skips on sentinel zero (per `pearl_first_observation_bootstrap`), + /// and updates the (count, mean, M²) Welford triple in place. Sample + /// variance = `M² / (count - 1)` when `count > 1`; controllers consume + /// it as `sqrt(M²/(count-1))` to scale their noise floors. + /// + /// Caller's responsibility: launch ONCE per controller-input EMA per + /// step, sequenced AFTER the EMA producer that writes `input_slot` + /// and BEFORE the controller that reads the Welford triple. + pub fn launch_rl_signal_variance_update( + &self, + input_slot: usize, + count_slot: usize, + mean_slot: usize, + m2_slot: usize, + ) -> Result<()> { + let mut args = RawArgs::new(); + args.push_ptr(self.isv_dev_ptr); + args.push_i32(input_slot as i32); + args.push_i32(count_slot as i32); + args.push_i32(mean_slot as i32); + args.push_i32(m2_slot as i32); + let mut ptrs = args.build_arg_ptrs(); + unsafe { + raw_launch( + self.rl_signal_variance_update_fn.cu_function(), + (1, 1, 1), (1, 1, 1), 0, + self.raw_stream, + &mut ptrs[..args.len()], + ).map_err(|e| anyhow::anyhow!("rl_signal_variance_update: {:?}", e))?; + } + Ok(()) + } + /// Phase R3: launch `ema_update_on_done` — done-gated EMA producer /// that writes the per-batch closed-trade mean into `isv[slot_index]`. /// `alpha` is the Wiener-α (caller must pre-floor at 0.4 per @@ -6355,6 +6490,70 @@ impl IntegratedTrainer { // recent_outcome_update — handled by rl_fused_reward_pipeline above. + // Adaptive controller floors (spec + // `2026-05-30-adaptive-controller-floor-design`): launch a + // Welford-variance update for each controller-input EMA so the + // 12 controllers + fused kernel can read `sqrt(M²/(n-1))` from + // their dedicated triple instead of falling back to + // `target × 0.5`. Sequenced AFTER all EMA producers above and + // BEFORE the fused controller launch / standalone v_blend_alpha + // controller below. Sentinel-zero skip is handled inside the + // kernel — no per-step gating needed here. + // + // Slots written *this* step (kl_pi, entropy_observed, td_kurtosis, + // mean_trade_duration) feed Welford with the current step's + // observation. Slots written *later* in the step + // (q_divergence at step end, q_distill_kl by distill grad, + // reward_magnitude by reward_scale controller, v_scalar_mag by + // v_blend_alpha controller, adv_var_pre_norm by advantage + // normalize) feed Welford with the previous step's value — a + // one-step lag, acceptable per the spec since controllers adapt + // over many steps and the Welford counter still reflects real + // observations. + for (input_slot, count_slot, mean_slot, m2_slot) in [ + (crate::rl::isv_slots::RL_KL_PI_EMA_INDEX, + crate::rl::isv_slots::RL_KL_PI_VAR_COUNT_INDEX, + crate::rl::isv_slots::RL_KL_PI_VAR_MEAN_INDEX, + crate::rl::isv_slots::RL_KL_PI_VAR_M2_INDEX), + (crate::rl::isv_slots::RL_Q_DIVERGENCE_EMA_INDEX, + crate::rl::isv_slots::RL_Q_DIV_VAR_COUNT_INDEX, + crate::rl::isv_slots::RL_Q_DIV_VAR_MEAN_INDEX, + crate::rl::isv_slots::RL_Q_DIV_VAR_M2_INDEX), + (crate::rl::isv_slots::RL_ADV_VAR_PRE_NORM_INDEX, + crate::rl::isv_slots::RL_ADV_VAR_VAR_COUNT_INDEX, + crate::rl::isv_slots::RL_ADV_VAR_VAR_MEAN_INDEX, + crate::rl::isv_slots::RL_ADV_VAR_VAR_M2_INDEX), + (crate::rl::isv_slots::RL_ENTROPY_OBSERVED_EMA_INDEX, + crate::rl::isv_slots::RL_ENTROPY_OBS_VAR_COUNT_INDEX, + crate::rl::isv_slots::RL_ENTROPY_OBS_VAR_MEAN_INDEX, + crate::rl::isv_slots::RL_ENTROPY_OBS_VAR_M2_INDEX), + (crate::rl::isv_slots::RL_TD_KURTOSIS_EMA_INDEX, + crate::rl::isv_slots::RL_TD_KURT_VAR_COUNT_INDEX, + crate::rl::isv_slots::RL_TD_KURT_VAR_MEAN_INDEX, + crate::rl::isv_slots::RL_TD_KURT_VAR_M2_INDEX), + (crate::rl::isv_slots::RL_MEAN_TRADE_DURATION_EMA_INDEX, + crate::rl::isv_slots::RL_TRADE_DUR_VAR_COUNT_INDEX, + crate::rl::isv_slots::RL_TRADE_DUR_VAR_MEAN_INDEX, + crate::rl::isv_slots::RL_TRADE_DUR_VAR_M2_INDEX), + (crate::rl::isv_slots::RL_REWARD_MAGNITUDE_EMA_INDEX, + crate::rl::isv_slots::RL_REWARD_MAG_VAR_COUNT_INDEX, + crate::rl::isv_slots::RL_REWARD_MAG_VAR_MEAN_INDEX, + crate::rl::isv_slots::RL_REWARD_MAG_VAR_M2_INDEX), + (crate::rl::isv_slots::RL_Q_DISTILL_KL_EMA_INDEX, + crate::rl::isv_slots::RL_Q_DISTILL_KL_VAR_COUNT_INDEX, + crate::rl::isv_slots::RL_Q_DISTILL_KL_VAR_MEAN_INDEX, + crate::rl::isv_slots::RL_Q_DISTILL_KL_VAR_M2_INDEX), + (crate::rl::isv_slots::RL_V_SCALAR_MAG_EMA_INDEX, + crate::rl::isv_slots::RL_V_BLEND_SCALAR_VAR_COUNT_INDEX, + crate::rl::isv_slots::RL_V_BLEND_SCALAR_VAR_MEAN_INDEX, + crate::rl::isv_slots::RL_V_BLEND_SCALAR_VAR_M2_INDEX), + ] { + self.launch_rl_signal_variance_update( + input_slot, count_slot, mean_slot, m2_slot, + ) + .context("launch_rl_signal_variance_update")?; + } + // Fire all 10 RL ISV controllers in a single fused kernel // launch (gamma, tau, ppo_clip, entropy_coef, rollout_steps, // per_alpha, reward_scale, ppo_ratio_clamp, gate_threshold, @@ -6591,6 +6790,10 @@ impl IntegratedTrainer { let block_x = (b_size as u32).min(1024).max(1); let mut args = RawArgs::new(); args.push_ptr(self.advantages_d.raw_ptr()); + // Adaptive controller floors (spec 2026-05-30): emit raw pre-norm + // variance to ISV[RL_ADV_VAR_PRE_NORM_INDEX=612] for + // rl_rollout_steps_controller's input signal. + args.push_ptr(self.isv_dev_ptr); args.push_i32(b_size_i); let mut ptrs = args.build_arg_ptrs(); unsafe { diff --git a/crates/ml-alpha/tests/controller_adaptive_floors.rs b/crates/ml-alpha/tests/controller_adaptive_floors.rs new file mode 100644 index 000000000..9189f21ca --- /dev/null +++ b/crates/ml-alpha/tests/controller_adaptive_floors.rs @@ -0,0 +1,524 @@ +//! Task 16 — Controller-level adaptive-floor invariant tests for the +//! 2026-05-30 adaptive-controller-floor refactor (spec +//! `docs/superpowers/specs/2026-05-30-adaptive-controller-floor-design.md`). +//! +//! Validates Validation Gates G3-G6 against the fused production launch +//! path (`launch_rl_fused_controllers`). Per +//! `feedback_no_cpu_test_fallbacks`: every oracle here is an analytical +//! identity of the controller's documented adaptive math — adaptive +//! noise floor (`max(target × 0.5, 2σ)`), asymmetric Schulman +//! (tighten on single above-band observation, widen requires N ≥ 3 +//! consecutive below-band) and gamma's `1 − 1/(d × 2)` adaptive-min +//! formula. No CPU reference reimplementation; the kernel's spec is the +//! oracle. +//! +//! Per `pearl_first_observation_bootstrap`: the bootstrap-replace path +//! (`eps_prev == RL_EPS_BOOTSTRAP_INDEX`) skips the Wiener blend so the +//! first real observation lands directly at the computed `eps_target`. +//! For G4/G5 we deliberately PRE-WRITE `RL_PPO_CLIP_INDEX` to a value +//! away from the bootstrap so the Wiener-blend branch is exercised; this +//! isolates the asymmetric-Schulman invariant from the bootstrap-replace +//! quirk. G3 stays on the natural post-bootstrap state (eps_prev = +//! bootstrap = 0.05 = EPS_MIN) because the noise-floor gate returns +//! before either branch fires — the invariant under test is "controller +//! holds". +//! +//! Run with: +//! `cargo test -p ml-alpha --test controller_adaptive_floors -- --ignored --nocapture` + +use ml_alpha::rl::isv_slots::{ + RL_GAMMA_INDEX, RL_GAMMA_MIN_ADAPTIVE_INDEX, RL_KL_PI_BELOW_COUNT_INDEX, + RL_KL_PI_EMA_INDEX, RL_KL_PI_VAR_COUNT_INDEX, RL_KL_PI_VAR_M2_INDEX, + RL_KL_PI_VAR_MEAN_INDEX, RL_KL_TARGET_INDEX, RL_MEAN_TRADE_DURATION_EMA_INDEX, + RL_PPO_CLIP_EPS_MAX_INDEX, RL_PPO_CLIP_INDEX, RL_SCHULMAN_ADJUST_RATE_INDEX, + RL_SCHULMAN_TOLERANCE_INDEX, RL_TRADE_DUR_VAR_COUNT_INDEX, + RL_TRADE_DUR_VAR_M2_INDEX, RL_TRADE_DUR_VAR_MEAN_INDEX, +}; +use ml_alpha::trainer::integrated::{IntegratedTrainer, IntegratedTrainerConfig}; +use ml_alpha::trainer::perception::PerceptionTrainerConfig; +use ml_core::device::MlDevice; + +/// Bootstrap value `rl_ppo_clip_controller` writes into +/// `ISV[RL_PPO_CLIP_INDEX]` at construction. Per Task 17 fix: +/// bootstrap matches `EPS_MIN` (= 0.05) so the post-bootstrap step +/// doesn't snap to MIN on first signal regardless of direction +/// (bootstrap-clamp inconsistency that surfaced in the Task 16 testing). +const EPS_BOOTSTRAP: f32 = 0.05; + +/// PPO ε MIN clamp seeded by `with_controllers_bootstrapped`. +const EPS_MIN: f32 = 0.05; + +/// KL target seeded by `with_controllers_bootstrapped`. +const KL_TARGET: f32 = 0.01; + +/// Schulman tolerance — kl above `target × tolerance` triggers tighten, +/// below `target / tolerance` triggers (delayed) widen. +const SCHULMAN_TOLERANCE: f32 = 1.5; + +/// Schulman adjust rate — tighten ratio is `1 / adjust_rate`, widen +/// ratio is `adjust_rate`. +const SCHULMAN_ADJUST_RATE: f32 = 1.5; + +/// N-consecutive-below-band threshold for the widen branch (matches +/// `WIDEN_PATIENCE_CONSECUTIVE` in `rl_ppo_clip_controller.cu`). +const WIDEN_PATIENCE: f32 = 3.0; + +/// Hard absolute floor for γ adaptive-min (matches `GAMMA_MIN_ABSOLUTE` +/// in `rl_gamma_controller.cu`). +const GAMMA_MIN_ABSOLUTE: f32 = 0.9; + +/// Welford warmup observations before the running mean replaces the EMA +/// in the γ adaptive-min derivation (matches `WELFORD_WARMUP_OBS`). +const WELFORD_WARMUP_OBS: f32 = 100.0; + +fn build_trainer() -> Option<(MlDevice, IntegratedTrainer)> { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { + eprintln!("CUDA 0 not available — skipping ({e})"); + return None; + } + }; + let cfg = IntegratedTrainerConfig { + perception: PerceptionTrainerConfig { + seq_len: 4, + n_batch: 1, + ..PerceptionTrainerConfig::default() + }, + dqn_seed: 0xC1, + ppo_seed: 0xC2, + ..IntegratedTrainerConfig::default() + }; + let trainer = IntegratedTrainer::new(&dev, cfg).expect("IntegratedTrainer::new"); + Some((dev, trainer)) +} + +/// Sync the trainer's production stream so subsequent mapped-pinned ISV +/// reads through `read_isv_host` observe device-side writes. +fn sync(trainer: &IntegratedTrainer) { + trainer.stream.synchronize().expect("sync trainer stream"); +} + +// ───────────────────────────────────────────────────────────────────── +// G3 — PPO clip holds when signal is below the adaptive noise floor. +// +// Invariant: after the bootstrap path has fired (eps_prev = bootstrap = +// 0.01), feeding `kl_pi_ema = 0.001` (< noise_floor = target × 0.5 = +// 0.005) for 50 fused launches must NOT drift ε. The noise-floor gate +// is the load-bearing fix — without it, the prior multiplicative-ratio +// design produced 1e4 ratios from sub-noise KL values and pegged ε at +// MAX 0.50 within ~50 steps (alpha-rl-mjzfk fold 0 saturation). +// +// Welford triple is pre-populated so the std-based half of the floor +// (`2σ`) is dominated by the target-fraction half (`target × 0.5`), +// keeping the floor at the steady-state value of 0.005. +// ───────────────────────────────────────────────────────────────────── +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn g3_ppo_clip_holds_when_signal_below_adaptive_floor() { + let Some((_dev, trainer)) = build_trainer() else { return }; + + // Verify pre-conditions seeded by `with_controllers_bootstrapped`. + sync(&trainer); + let eps_at_bootstrap = trainer.read_isv_host(RL_PPO_CLIP_INDEX); + assert!( + (eps_at_bootstrap - EPS_BOOTSTRAP).abs() < 1e-6, + "pre-condition: ε must be bootstrapped to {EPS_BOOTSTRAP}; got {eps_at_bootstrap}" + ); + let kl_target_seeded = trainer.read_isv_host(RL_KL_TARGET_INDEX); + assert!( + (kl_target_seeded - KL_TARGET).abs() < 1e-6, + "pre-condition: KL target must be seeded to {KL_TARGET}; got {kl_target_seeded}" + ); + + // Pre-populate Welford triple at the kl_pi_ema slot with low M² + // so std ≈ 0 and the floor reduces to `target × 0.5 = 0.005`. count + // > 1 is required so the kernel evaluates `M²/(count − 1)` rather + // than the count ≤ 1 fallback (which would also produce std = 0). + trainer.isv_mapped.write_record(RL_KL_PI_VAR_COUNT_INDEX, 100.0); + trainer.isv_mapped.write_record(RL_KL_PI_VAR_MEAN_INDEX, 0.001); + trainer.isv_mapped.write_record(RL_KL_PI_VAR_M2_INDEX, 0.0); + + // Feed below-floor signal for 50 fused launches. Re-write each step + // because a real-world EMA producer would refresh it; here we're + // testing the controller in isolation. + for _ in 0..50 { + trainer.isv_mapped.write_record(RL_KL_PI_EMA_INDEX, 0.001); + trainer + .launch_rl_fused_controllers(1) + .expect("launch_rl_fused_controllers"); + } + sync(&trainer); + + let eps_after = trainer.read_isv_host(RL_PPO_CLIP_INDEX); + + // Primary invariant: ε held — no drift past bootstrap, MUST not have + // climbed toward MAX 0.5 (the canonical saturation pattern). + assert!( + (eps_after - EPS_BOOTSTRAP).abs() < 1e-4, + "ε must hold at bootstrap when signal is below adaptive floor; \ + expected ≈ {EPS_BOOTSTRAP}, got {eps_after}" + ); + assert!( + eps_after < 0.4, + "ε must NOT drift toward MAX 0.5 — this was the fold 0/1 \ + saturation bug; got {eps_after}" + ); + + // Negative invariant: the below-band counter must have stayed at 0 + // because the noise-floor gate returns BEFORE the asymmetric Schulman + // branch fires. + let below_count = trainer.read_isv_host(RL_KL_PI_BELOW_COUNT_INDEX); + assert_eq!( + below_count, 0.0, + "below-band counter must stay 0 when noise-floor gate held; got {below_count}" + ); + + eprintln!( + "G3 OK — ε held at {eps_after} (bootstrap {EPS_BOOTSTRAP}) over 50 below-floor steps; \ + below_count = {below_count} (noise-floor gate fired)" + ); +} + +// ───────────────────────────────────────────────────────────────────── +// G4 — PPO clip tightens on a single above-band observation +// (asymmetric Schulman — fast tighten). +// +// Invariant: a single `kl_pi_ema > target × tolerance` observation must +// decrease ε within the same call. Tighten ratio = `1 / adjust_rate ≈ +// 0.667`; the assertion is strict decrease (ε_after < ε_before). +// +// Test discipline: we pre-write ε to 0.2 (well above MIN 0.05 and away +// from bootstrap 0.01) so: +// 1. The bootstrap-replace branch DOESN'T fire (eps_prev ≠ bootstrap). +// 2. `eps_target = 0.2 × 0.667 ≈ 0.133` stays above MIN, so the clamp +// doesn't bind and obscure the tighten signal. +// The Wiener-α blend (alpha = 0.4 = Wiener floor) then produces +// `eps_new = 0.6 × 0.2 + 0.4 × 0.133 ≈ 0.173`. Below 0.2 → strict +// decrease. +// ───────────────────────────────────────────────────────────────────── +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn g4_ppo_clip_tightens_on_single_above_band_observation() { + let Some((_dev, trainer)) = build_trainer() else { return }; + + // Pre-write ε to 0.2 — away from bootstrap, above MIN clamp. + let eps_initial: f32 = 0.2; + trainer.isv_mapped.write_record(RL_PPO_CLIP_INDEX, eps_initial); + + // Welford triple — std ≈ 0 so the noise floor is `target × 0.5 = 0.005`. + trainer.isv_mapped.write_record(RL_KL_PI_VAR_COUNT_INDEX, 100.0); + trainer.isv_mapped.write_record(RL_KL_PI_VAR_MEAN_INDEX, 0.01); + trainer.isv_mapped.write_record(RL_KL_PI_VAR_M2_INDEX, 0.0); + + // Above-band signal: kl = 0.05 = 5× target = 3.33× tolerance bound. + trainer.isv_mapped.write_record(RL_KL_PI_EMA_INDEX, 0.05); + + // Verify Schulman config matches what the controller will read. + let tolerance = trainer.read_isv_host(RL_SCHULMAN_TOLERANCE_INDEX); + let adjust_rate = trainer.read_isv_host(RL_SCHULMAN_ADJUST_RATE_INDEX); + assert!( + (tolerance - SCHULMAN_TOLERANCE).abs() < 1e-6, + "pre-condition: Schulman tolerance seeded to {SCHULMAN_TOLERANCE}; got {tolerance}" + ); + assert!( + (adjust_rate - SCHULMAN_ADJUST_RATE).abs() < 1e-6, + "pre-condition: Schulman adjust_rate seeded to {SCHULMAN_ADJUST_RATE}; got {adjust_rate}" + ); + + trainer + .launch_rl_fused_controllers(1) + .expect("launch_rl_fused_controllers"); + sync(&trainer); + + let eps_after = trainer.read_isv_host(RL_PPO_CLIP_INDEX); + + // Primary invariant: ε strictly decreased on a single above-band fire. + assert!( + eps_after < eps_initial, + "ε must tighten on single above-band observation; \ + before = {eps_initial}, after = {eps_after}" + ); + // Above the MIN clamp — the decrease is genuine adaptation, not a + // saturation pin. + assert!( + eps_after > EPS_MIN, + "ε must stay above MIN {EPS_MIN} after a single fire (clamp \ + not binding); got {eps_after}" + ); + + // The below-band counter must reset on a tighten event (the kernel's + // `else if (kl_ema > target × tolerance) { … below_count = 0 }`). + let below_count = trainer.read_isv_host(RL_KL_PI_BELOW_COUNT_INDEX); + assert_eq!( + below_count, 0.0, + "below-band counter must reset on tighten; got {below_count}" + ); + + eprintln!( + "G4 OK — ε tightened {eps_initial} → {eps_after} (single above-band fire); \ + below_count reset to 0" + ); +} + +// ───────────────────────────────────────────────────────────────────── +// G5 — PPO clip widens only after N consecutive below-band observations +// (asymmetric Schulman — slow widen). +// +// Invariant: a kl_pi_ema in the widen-band `[noise_floor, target / +// tolerance) = [0.005, 0.00667)` must increment `below_count` each step +// but NOT widen ε until `below_count ≥ WIDEN_PATIENCE_CONSECUTIVE = 3`. +// +// We choose kl = 0.006 because: +// - 0.006 ≥ 0.005 (noise floor — passes the gate) +// - 0.006 < 0.00667 (target / tolerance — triggers widen branch, not +// hold) +// To keep the noise floor low (so kl=0.006 passes it), Welford std must +// be small: count=100, M²=0 → std=0 → floor = target × 0.5 = 0.005. +// +// Test discipline: ε pre-written to 0.2 so the bootstrap-replace branch +// doesn't fire. We expect: +// step 1: below_count 0 → 1, ratio = 1 (hold), ε unchanged +// step 2: below_count 1 → 2, ratio = 1 (hold), ε unchanged +// step 3: below_count 2 → 3, ratio = adjust_rate = 1.5 (widen), ε +// increases via Wiener blend +// step 4: below_count 3 → 4, ratio = 1.5 (widen continues) +// ───────────────────────────────────────────────────────────────────── +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn g5_ppo_clip_widens_only_after_n_consecutive_below_band() { + let Some((_dev, trainer)) = build_trainer() else { return }; + + let eps_initial: f32 = 0.2; + trainer.isv_mapped.write_record(RL_PPO_CLIP_INDEX, eps_initial); + + // Welford triple — low std so noise floor = target × 0.5 = 0.005. + trainer.isv_mapped.write_record(RL_KL_PI_VAR_COUNT_INDEX, 100.0); + trainer.isv_mapped.write_record(RL_KL_PI_VAR_MEAN_INDEX, 0.006); + trainer.isv_mapped.write_record(RL_KL_PI_VAR_M2_INDEX, 0.0); + + // Counter must start at 0. + trainer + .isv_mapped + .write_record(RL_KL_PI_BELOW_COUNT_INDEX, 0.0); + + // Widen-band signal: 0.005 ≤ 0.006 < 0.00667. + let kl_widen: f32 = 0.006; + + // Compute the widen-band edges from the canonical formulas so the + // assertion that 0.006 falls inside is auditable. + let noise_floor = KL_TARGET * 0.5; + let widen_upper = KL_TARGET / SCHULMAN_TOLERANCE; + assert!( + kl_widen >= noise_floor, + "kl_widen {kl_widen} must be ≥ noise floor {noise_floor}" + ); + assert!( + kl_widen < widen_upper, + "kl_widen {kl_widen} must be < widen upper edge {widen_upper}" + ); + + let mut below_counts = Vec::with_capacity(4); + let mut eps_values = Vec::with_capacity(4); + + for _ in 0..4 { + trainer.isv_mapped.write_record(RL_KL_PI_EMA_INDEX, kl_widen); + trainer + .launch_rl_fused_controllers(1) + .expect("launch_rl_fused_controllers"); + sync(&trainer); + below_counts.push(trainer.read_isv_host(RL_KL_PI_BELOW_COUNT_INDEX)); + eps_values.push(trainer.read_isv_host(RL_PPO_CLIP_INDEX)); + } + + for (i, (bc, eps)) in below_counts.iter().zip(eps_values.iter()).enumerate() { + eprintln!(" step {}: below_count = {bc}, ε = {eps}", i + 1); + } + + // Below-counter invariant: increments monotonically per step. + assert_eq!(below_counts[0], 1.0, "step 1: below_count must be 1"); + assert_eq!(below_counts[1], 2.0, "step 2: below_count must be 2"); + assert_eq!( + below_counts[2], WIDEN_PATIENCE, + "step 3: below_count must reach WIDEN_PATIENCE {WIDEN_PATIENCE}" + ); + assert_eq!(below_counts[3], 4.0, "step 4: below_count must be 4"); + + // ε hold invariant: steps 1 and 2 (below_count < 3) MUST NOT widen. + // Allow only float-rounding-level drift (the controller's + // `ratio = 1.0f` branch produces eps_target = eps_prev, then the + // Wiener blend with eps_target == eps_prev is the identity). + let drift_1 = (eps_values[0] - eps_initial).abs(); + let drift_2 = (eps_values[1] - eps_initial).abs(); + assert!( + drift_1 < 1e-5, + "step 1: ε must hold (count < WIDEN_PATIENCE); drift {drift_1}, value {}", + eps_values[0] + ); + assert!( + drift_2 < 1e-5, + "step 2: ε must hold (count < WIDEN_PATIENCE); drift {drift_2}, value {}", + eps_values[1] + ); + + // ε widen invariant: step 3 (below_count reaches WIDEN_PATIENCE) MUST + // strictly increase ε. Step 4 continues widening. + assert!( + eps_values[2] > eps_initial + 1e-5, + "step 3: ε must widen once below_count reaches {WIDEN_PATIENCE}; \ + was {eps_initial}, got {}", + eps_values[2] + ); + assert!( + eps_values[3] > eps_values[2], + "step 4: ε must continue widening; step 3 = {}, step 4 = {}", + eps_values[2], + eps_values[3] + ); + + // ε stays under MAX clamp — widening is real adaptation, not a + // saturation pin. + let eps_max = trainer.read_isv_host(RL_PPO_CLIP_EPS_MAX_INDEX); + assert!( + eps_values[3] < eps_max, + "ε must stay under MAX {eps_max} (no saturation pin); got {}", + eps_values[3] + ); + + eprintln!( + "G5 OK — ε held for 2 steps, widened on step 3 (count reached {WIDEN_PATIENCE}): \ + {eps_initial} → {} → {} → {} → {}", + eps_values[0], eps_values[1], eps_values[2], eps_values[3] + ); +} + +// ───────────────────────────────────────────────────────────────────── +// G6 — γ adaptive_min tracks trade duration via Welford mean (post-100- +// observation warmup) and the EMA before warmup. +// +// Invariant (post-warmup): with Welford count ≥ 100 and Welford mean d, +// the kernel must write +// +// RL_GAMMA_MIN_ADAPTIVE = max(0.9, 1 − 1/max(d × 2, 10)) +// +// to ISV[613]. At d = 30: `1 − 1/60 ≈ 0.98333`. +// +// Invariant (pre-warmup): with count < 100, the formula uses the EMA at +// `RL_MEAN_TRADE_DURATION_EMA_INDEX` instead of the Welford mean. At +// EMA = 15: `1 − 1/30 ≈ 0.96667`. +// +// Per spec Section "Special case G — γ_min coupling": the Welford +// mean's N-smoothing breaks the γ ↔ trade_duration feedback loop without +// any explicit cool-down. Welford count provides the confidence +// threshold for the switchover. +// ───────────────────────────────────────────────────────────────────── +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn g6_gamma_adaptive_min_post_warmup_uses_welford_mean() { + let Some((_dev, trainer)) = build_trainer() else { return }; + + // Post-warmup state: count = 100 = WELFORD_WARMUP_OBS, mean = 30. + let d_welford: f32 = 30.0; + trainer + .isv_mapped + .write_record(RL_TRADE_DUR_VAR_COUNT_INDEX, WELFORD_WARMUP_OBS); + trainer + .isv_mapped + .write_record(RL_TRADE_DUR_VAR_MEAN_INDEX, d_welford); + trainer.isv_mapped.write_record(RL_TRADE_DUR_VAR_M2_INDEX, 0.0); + // Also write the EMA — even though the kernel ignores it after + // warmup, mismatching it from the Welford mean tests that the + // controller actually selects the Welford path (not the EMA). + trainer + .isv_mapped + .write_record(RL_MEAN_TRADE_DURATION_EMA_INDEX, d_welford); + + trainer + .launch_rl_fused_controllers(1) + .expect("launch_rl_fused_controllers"); + sync(&trainer); + + let adaptive_min = trainer.read_isv_host(RL_GAMMA_MIN_ADAPTIVE_INDEX); + + // Analytical oracle: `max(0.9, 1 - 1/max(d × 2, 10))`. At d = 30, + // d × 2 = 60 > 10 = HORIZON_FLOOR, so the inner `max` resolves to 60 + // and the formula reduces to `1 - 1/60`. + let horizon = (d_welford * 2.0_f32).max(10.0); + let expected = (1.0_f32 - 1.0 / horizon).max(GAMMA_MIN_ABSOLUTE); + assert!( + (adaptive_min - expected).abs() < 1e-4, + "post-warmup γ_min must be max(0.9, 1 - 1/{horizon}) = {expected}; got {adaptive_min}" + ); + + // γ itself must respect the new floor — bootstrap value 0.9 was + // computed against the cold-start adaptive_min 0.9; after the + // controller re-fires with the new welford mean, γ should be + // clamped to the new floor. + let gamma = trainer.read_isv_host(RL_GAMMA_INDEX); + assert!( + gamma >= adaptive_min - 1e-4, + "γ must respect adaptive_min {adaptive_min}; got {gamma}" + ); + + eprintln!( + "G6 OK (post-warmup) — adaptive_min = {adaptive_min} (expected {expected}) \ + at Welford count = {WELFORD_WARMUP_OBS}, mean = {d_welford}; γ = {gamma}" + ); +} + +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn g6_gamma_adaptive_min_pre_warmup_uses_ema() { + let Some((_dev, trainer)) = build_trainer() else { return }; + + // Pre-warmup state: count = 99 < WELFORD_WARMUP_OBS = 100. The + // controller MUST fall back to the EMA slot, not the Welford mean. + // Deliberately set EMA ≠ Welford mean so the test confirms the + // controller's selection logic. + let d_ema: f32 = 15.0; + let d_welford: f32 = 30.0; + trainer + .isv_mapped + .write_record(RL_TRADE_DUR_VAR_COUNT_INDEX, WELFORD_WARMUP_OBS - 1.0); + trainer + .isv_mapped + .write_record(RL_TRADE_DUR_VAR_MEAN_INDEX, d_welford); + trainer.isv_mapped.write_record(RL_TRADE_DUR_VAR_M2_INDEX, 0.0); + trainer + .isv_mapped + .write_record(RL_MEAN_TRADE_DURATION_EMA_INDEX, d_ema); + + trainer + .launch_rl_fused_controllers(1) + .expect("launch_rl_fused_controllers"); + sync(&trainer); + + let adaptive_min = trainer.read_isv_host(RL_GAMMA_MIN_ADAPTIVE_INDEX); + + // Oracle uses the EMA (15), not the Welford mean (30). + let horizon = (d_ema * 2.0_f32).max(10.0); + let expected = (1.0_f32 - 1.0 / horizon).max(GAMMA_MIN_ABSOLUTE); + assert!( + (adaptive_min - expected).abs() < 1e-4, + "pre-warmup γ_min must use EMA, not Welford mean: \ + expected max(0.9, 1 - 1/{horizon}) = {expected}; got {adaptive_min}" + ); + + // Counter-invariant: had the controller used the Welford mean + // instead (count = 99 bug), adaptive_min would be ≈ 0.98333. The + // test fails if that magic number appears. + let welford_horizon = (d_welford * 2.0_f32).max(10.0); + let wrong_min = (1.0_f32 - 1.0 / welford_horizon).max(GAMMA_MIN_ABSOLUTE); + assert!( + (adaptive_min - wrong_min).abs() > 1e-3, + "pre-warmup γ_min must NOT match the Welford-mean value {wrong_min} \ + (would indicate the count < WARMUP guard is broken); got {adaptive_min}" + ); + + eprintln!( + "G6 OK (pre-warmup) — adaptive_min = {adaptive_min} (expected {expected}, \ + distinct from welford-path {wrong_min}) at Welford count = {}, EMA = {d_ema}", + WELFORD_WARMUP_OBS - 1.0 + ); +} + diff --git a/crates/ml-alpha/tests/signal_variance_kernel.rs b/crates/ml-alpha/tests/signal_variance_kernel.rs new file mode 100644 index 000000000..c1c6c27b4 --- /dev/null +++ b/crates/ml-alpha/tests/signal_variance_kernel.rs @@ -0,0 +1,185 @@ +//! GPU-oracle tests for `rl_signal_variance_update` — Welford online +//! variance kernel for adaptive controller floors (spec +//! `2026-05-30-adaptive-controller-floor-design`). +//! +//! Three analytical invariants, zero CPU reference implementations: +//! +//! 1. `welford_constant_input_zero_variance` — feeding a constant `k` +//! for N steps must produce `mean == k` and `m2 == 0` (sample +//! variance is zero by definition). +//! 2. `welford_known_sequence_variance` — sequence [1, 2, 3, 4, 5] +//! has known sample variance 2.5; Welford output must match. +//! 3. `welford_sentinel_zero_is_skipped` — input == 0.0 sentinel +//! must produce no update (per +//! `pearl_first_observation_bootstrap`). +//! +//! Per `feedback_no_cpu_test_fallbacks`: oracles are analytical +//! identities of Welford's algorithm, not CPU reimplementations. +//! Per `feedback_no_htod_htoh_only_mapped_pinned`: ISV bus uses +//! mapped-pinned `MappedF32Buffer` (same pattern as production trainer). +//! +//! Run with: +//! `cargo test -p ml-alpha --test signal_variance_kernel -- --ignored --nocapture` + +use cudarc::driver::{LaunchConfig, PushKernelArg}; +use ml_alpha::pinned_mem::MappedF32Buffer; +use ml_core::device::MlDevice; + +// ISV slot layout for this isolated test: input at 0, Welford triple at 1,2,3. +const INPUT_SLOT: i32 = 0; +const COUNT_SLOT: i32 = 1; +const MEAN_SLOT: i32 = 2; +const M2_SLOT: i32 = 3; +const N_SLOTS: usize = 4; + +const SIGNAL_VARIANCE_UPDATE_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_signal_variance_update.cubin")); + +fn launch_once( + dev: &MlDevice, + isv: &MappedF32Buffer, +) { + let stream = dev.cuda_stream().expect("cuda_stream").clone(); + let ctx = dev.cuda_context().expect("cuda_context"); + let module = ctx + .load_cubin(SIGNAL_VARIANCE_UPDATE_CUBIN.to_vec()) + .expect("load signal_variance cubin"); + let func = module + .load_function("rl_signal_variance_update") + .expect("load rl_signal_variance_update fn"); + let cfg = LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }; + let isv_dev_ptr = isv.dev_ptr; + unsafe { + stream + .launch_builder(&func) + .arg(&isv_dev_ptr) + .arg(&INPUT_SLOT) + .arg(&COUNT_SLOT) + .arg(&MEAN_SLOT) + .arg(&M2_SLOT) + .launch(cfg) + .expect("rl_signal_variance_update launch"); + } + stream.synchronize().expect("sync"); +} + +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn welford_constant_input_zero_variance() { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { + eprintln!("CUDA 0 not available — skipping ({e})"); + return; + } + }; + let isv = unsafe { MappedF32Buffer::new(N_SLOTS) }.expect("isv alloc"); + + // Welford of constant k for N steps: mean = k, m2 = 0, count = N. + let k = 0.5_f32; + let n = 50_usize; + + for _ in 0..n { + isv.write_record(INPUT_SLOT as usize, k); + launch_once(&dev, &isv); + } + + let count = isv.read_record(COUNT_SLOT as usize); + let mean = isv.read_record(MEAN_SLOT as usize); + let m2 = isv.read_record(M2_SLOT as usize); + + assert_eq!(count, n as f32, "count must be {n}; got {count}"); + assert!( + (mean - k).abs() < 1e-6, + "mean of constant {k} must equal {k}; got {mean}" + ); + assert!( + m2.abs() < 1e-5, + "m2 of constant input must be exactly 0 (or fp-noise); got {m2}" + ); + + eprintln!("G1 OK — constant k={k} for N={n}: mean={mean} m2={m2}"); +} + +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn welford_known_sequence_variance() { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { + eprintln!("CUDA 0 not available — skipping ({e})"); + return; + } + }; + let isv = unsafe { MappedF32Buffer::new(N_SLOTS) }.expect("isv alloc"); + + // Sequence [1, 2, 3, 4, 5]: arithmetic mean = 3. + // Sample variance = Σ(x_i − μ)² / (n − 1) = (4+1+0+1+4)/4 = 2.5. + let seq = [1.0_f32, 2.0, 3.0, 4.0, 5.0]; + + for &x in seq.iter() { + isv.write_record(INPUT_SLOT as usize, x); + launch_once(&dev, &isv); + } + + let count = isv.read_record(COUNT_SLOT as usize); + let mean = isv.read_record(MEAN_SLOT as usize); + let m2 = isv.read_record(M2_SLOT as usize); + let sample_var = m2 / (count - 1.0); + + assert_eq!(count, seq.len() as f32, "count must be {}", seq.len()); + assert!( + (mean - 3.0).abs() < 1e-5, + "mean of [1..5] must be 3.0; got {mean}" + ); + assert!( + (sample_var - 2.5).abs() < 1e-5, + "sample variance of [1..5] must be 2.5; got {sample_var} (m2={m2})" + ); + + eprintln!("G2 OK — sequence {seq:?}: mean={mean} sample_var={sample_var}"); +} + +#[test] +#[ignore = "requires CUDA (MlDevice::cuda(0))"] +fn welford_sentinel_zero_is_skipped() { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { + eprintln!("CUDA 0 not available — skipping ({e})"); + return; + } + }; + let isv = unsafe { MappedF32Buffer::new(N_SLOTS) }.expect("isv alloc"); + + // INPUT_SLOT remains at sentinel 0.0 (MappedF32Buffer::new zero-inits). + // Kernel must NOT count this as an observation per + // pearl_first_observation_bootstrap. + for _ in 0..10 { + launch_once(&dev, &isv); + } + assert_eq!( + isv.read_record(COUNT_SLOT as usize), + 0.0, + "sentinel zero must not be counted as a real observation" + ); + + // Now feed a real non-zero observation; counter must advance. + isv.write_record(INPUT_SLOT as usize, 7.0); + launch_once(&dev, &isv); + assert_eq!( + isv.read_record(COUNT_SLOT as usize), + 1.0, + "first real observation must be counted" + ); + assert!( + (isv.read_record(MEAN_SLOT as usize) - 7.0).abs() < 1e-6, + "first observation must initialize mean directly per Welford recurrence" + ); + + eprintln!("G3 OK — sentinel skip honoured; first real obs initialized mean"); +} diff --git a/docs/superpowers/plans/2026-05-30-adaptive-controller-floor-plan.md b/docs/superpowers/plans/2026-05-30-adaptive-controller-floor-plan.md new file mode 100644 index 000000000..b427ffd1f --- /dev/null +++ b/docs/superpowers/plans/2026-05-30-adaptive-controller-floor-plan.md @@ -0,0 +1,957 @@ +# Adaptive Controller Signal Floors — Implementation Plan + +> **For agentic workers:** REQUIRED SUB-SKILL: Use superpowers:subagent-driven-development (recommended) or superpowers:executing-plans to implement this plan task-by-task. Steps use checkbox (`- [ ]`) syntax for tracking. + +**Goal:** Replace hardcoded thresholds in 12 RL controllers with observed-signal-driven ISV slots. One atomic commit. Eliminate the architectural failure mode where Phase 4.5 normalization invalidated 8 controllers simultaneously. + +**Architecture:** New Welford-variance CUDA kernel (`rl_signal_variance_update.cu`) tracks running variance of every controller's input EMA. Controllers consume `sqrt(observed_var)` as their noise-floor reference instead of hardcoded constants. 52 new ISV slots wire the kernel into the existing per-step pipeline. Group A (8 saturated) gets asymmetric Schulman (tighten fast, widen on N consecutive); Group B (4 with hardcoded thresholds) gets the same Welford floors without the asymmetry. + +**Tech Stack:** CUDA 12.4 sm_89/sm_86, cudarc 0.19.3, Rust 1.85+ Edition 2021, single-stream graph capture, mapped-pinned ISV bus. + +**Spec:** `docs/superpowers/specs/2026-05-30-adaptive-controller-floor-design.md` + +--- + +## File Structure + +### Created + +- `crates/ml-alpha/cuda/rl_signal_variance_update.cu` — Welford online variance kernel (~80 LOC). One launch per controller-input slot per step. Single-thread, single-block. Computes count/mean/M² triple in ISV. + +### Modified + +- `crates/ml-alpha/src/rl/isv_slots.rs` — +52 new slot constants, `RL_SLOTS_END` 588 → 640 +- `crates/ml-alpha/build.rs` — register `rl_signal_variance_update` cubin +- `crates/ml-alpha/cuda/rl_ppo_clip_controller.cu` — adaptive noise floor + asymmetric Schulman +- `crates/ml-alpha/cuda/rl_target_tau_controller.cu` — same pattern +- `crates/ml-alpha/cuda/rl_rollout_steps_controller.cu` — same pattern + consume `RL_ADV_VAR_PRE_NORM_INDEX` +- `crates/ml-alpha/cuda/rl_entropy_coef_controller.cu` — same pattern +- `crates/ml-alpha/cuda/rl_per_alpha_controller.cu` — same pattern +- `crates/ml-alpha/cuda/rl_gamma_controller.cu` — adaptive `GAMMA_MIN` (Special G) +- `crates/ml-alpha/cuda/rl_reward_scale_controller.cu` — asymmetric rate-limit + bootstrap-fraction floor (Special R) +- `crates/ml-alpha/cuda/rl_q_distill_lambda_controller.cu` — adaptive `MIN_LAMBDA` / `MAX_LAMBDA` (Special Q) +- `crates/ml-alpha/cuda/rl_v_blend_alpha_controller.cu` — 5 hardcoded constants → 5 ISV slots +- `crates/ml-alpha/cuda/rl_ppo_ratio_clamp_controller.cu` — adaptive `MIN_OUT` / `MAX_OUT` +- `crates/ml-alpha/cuda/rl_reward_clamp_controller.cu` — 5 hardcoded constants → 5 ISV slots +- `crates/ml-alpha/cuda/rl_gate_threshold_controller.cu` — hardcoded `alpha = 0.01f` → 1 ISV slot +- `crates/ml-alpha/cuda/rl_advantage_normalize.cu` — emit `var_pre_norm` to ISV +- `crates/ml-alpha/cuda/rl_fused_controllers.cu` — unified update covering all 12 controllers +- `crates/ml-alpha/src/trainer/integrated.rs` — launch `rl_signal_variance_update` per step + populate new bootstrap ISV slots in `with_controllers_bootstrapped` + +### Tests (added/modified) + +- `crates/ml-alpha/tests/signal_variance_kernel.rs` — new GPU-oracle tests for Welford convergence +- `crates/ml-alpha/tests/controller_adaptive_floors.rs` — new integration tests: each controller holds at bootstrap when signal < adaptive floor +- `crates/ml-alpha/tests/isv_bootstrap.rs` — update with new slots in canonical bootstrap list + +--- + +## Tasks + +### Task 1: Allocate ISV slot constants + +**Files:** +- Modify: `crates/ml-alpha/src/rl/isv_slots.rs` (around line 1135) + +- [ ] **Step 1: Add 52 new constants after current `RL_V_SCALAR_MAG_EMA_INDEX = 587`** + +```rust +// ============================================================ +// Adaptive controller floors (spec 2026-05-30-adaptive-controller-floor-design) +// Welford variance triples + asymmetric Schulman counters + new input slots. +// ============================================================ + +// Group A — Welford variance triples (count, mean, M²) for each saturated controller's input. +pub const RL_KL_PI_VAR_COUNT_INDEX: usize = 588; +pub const RL_KL_PI_VAR_MEAN_INDEX: usize = 589; +pub const RL_KL_PI_VAR_M2_INDEX: usize = 590; +pub const RL_KL_PI_BELOW_COUNT_INDEX: usize = 591; + +pub const RL_Q_DIV_VAR_COUNT_INDEX: usize = 592; +pub const RL_Q_DIV_VAR_MEAN_INDEX: usize = 593; +pub const RL_Q_DIV_VAR_M2_INDEX: usize = 594; +pub const RL_Q_DIV_BELOW_COUNT_INDEX: usize = 595; + +pub const RL_ADV_VAR_VAR_COUNT_INDEX: usize = 596; +pub const RL_ADV_VAR_VAR_MEAN_INDEX: usize = 597; +pub const RL_ADV_VAR_VAR_M2_INDEX: usize = 598; +pub const RL_ADV_VAR_BELOW_COUNT_INDEX: usize = 599; + +pub const RL_ENTROPY_OBS_VAR_COUNT_INDEX: usize = 600; +pub const RL_ENTROPY_OBS_VAR_MEAN_INDEX: usize = 601; +pub const RL_ENTROPY_OBS_VAR_M2_INDEX: usize = 602; +pub const RL_ENTROPY_OBS_BELOW_COUNT_INDEX: usize = 603; + +pub const RL_TD_KURT_VAR_COUNT_INDEX: usize = 604; +pub const RL_TD_KURT_VAR_MEAN_INDEX: usize = 605; +pub const RL_TD_KURT_VAR_M2_INDEX: usize = 606; +pub const RL_TD_KURT_BELOW_COUNT_INDEX: usize = 607; + +pub const RL_TRADE_DUR_VAR_COUNT_INDEX: usize = 608; +pub const RL_TRADE_DUR_VAR_MEAN_INDEX: usize = 609; +pub const RL_TRADE_DUR_VAR_M2_INDEX: usize = 610; + +// New input signals +pub const RL_ADV_VAR_PRE_NORM_INDEX: usize = 612; +pub const RL_GAMMA_MIN_ADAPTIVE_INDEX: usize = 613; +pub const RL_REWARD_MAGNITUDE_EMA_INDEX: usize = 614; + +pub const RL_REWARD_MAG_VAR_COUNT_INDEX: usize = 615; +pub const RL_REWARD_MAG_VAR_MEAN_INDEX: usize = 616; +pub const RL_REWARD_MAG_VAR_M2_INDEX: usize = 617; + +pub const RL_Q_DISTILL_KL_VAR_COUNT_INDEX: usize = 618; +pub const RL_Q_DISTILL_KL_VAR_MEAN_INDEX: usize = 619; +pub const RL_Q_DISTILL_KL_VAR_M2_INDEX: usize = 620; + +// Group B — V-blend alpha (Phase 4.4) adaptive bounds +pub const RL_V_BLEND_DEAD_SIGNAL_FLOOR_INDEX: usize = 621; +pub const RL_V_BLEND_TARGET_TRACK_RATIO_INDEX: usize = 622; +pub const RL_V_BLEND_SCHULMAN_STEP_INDEX: usize = 623; +pub const RL_V_BLEND_EMA_ALPHA_INDEX: usize = 624; +pub const RL_V_BLEND_BOOTSTRAP_ALPHA_INDEX: usize = 625; +pub const RL_V_BLEND_SCALAR_VAR_COUNT_INDEX: usize = 626; +pub const RL_V_BLEND_SCALAR_VAR_MEAN_INDEX: usize = 627; +pub const RL_V_BLEND_SCALAR_VAR_M2_INDEX: usize = 628; + +// Group B — PPO ratio clamp adaptive bounds +pub const RL_PPO_RATIO_CLAMP_MAX_ADAPTIVE_INDEX: usize = 629; +pub const RL_PPO_RATIO_VAR_COUNT_INDEX: usize = 630; +pub const RL_PPO_RATIO_VAR_MEAN_INDEX: usize = 631; +pub const RL_PPO_RATIO_VAR_M2_INDEX: usize = 632; + +// Group B — Reward clamp adaptive bounds +pub const RL_REWARD_CLAMP_V_BOUND_FLOOR_INDEX: usize = 633; +pub const RL_REWARD_CLAMP_V_BOUND_EWMA_ALPHA_INDEX: usize = 634; +pub const RL_REWARD_CLAMP_MIN_WIN_INDEX: usize = 635; +pub const RL_REWARD_CLAMP_MIN_RATIO_INDEX: usize = 636; +pub const RL_REWARD_CLAMP_MAX_RATIO_INDEX: usize = 637; + +// Group B — Gate threshold EMA alpha +pub const RL_GATE_EMA_ALPHA_INDEX: usize = 638; + +// Q-distill adaptive bound +pub const RL_Q_DISTILL_LAMBDA_MIN_ADAPTIVE_INDEX: usize = 639; + +pub const RL_SLOTS_END: usize = 640; // was 588 +``` + +- [ ] **Step 2: Run `cargo check -p ml-alpha` and verify no slot-numbering conflicts** + +Run: `SQLX_OFFLINE=true cargo check -p ml-alpha 2>&1 | tail -5` +Expected: `Finished` (no errors) + +- [ ] **Step 3: Commit** + +Do not commit yet — atomic single commit at end per Option B. + +--- + +### Task 2: Welford variance kernel + +**Files:** +- Create: `crates/ml-alpha/cuda/rl_signal_variance_update.cu` +- Modify: `crates/ml-alpha/build.rs` + +- [ ] **Step 1: Write the kernel** + +```c +// rl_signal_variance_update.cu — Welford online variance for controller inputs. +// One-thread, one-block kernel. Updates count/mean/M² triple in ISV. +// Caller launches once per controller-input slot per step (or via fused kernel). +// +// Per pearl_first_observation_bootstrap: skip when input == sentinel 0.0. +// Per feedback_adaptive_not_tuned: emits observed signal statistics that +// controllers consume as noise-floor anchors. + +extern "C" __global__ void rl_signal_variance_update( + float* __restrict__ isv, + int input_slot, + int count_slot, + int mean_slot, + int m2_slot +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + const float x = isv[input_slot]; + if (x == 0.0f) return; // sentinel skip + + const float n_prev = isv[count_slot]; + const float n_new = n_prev + 1.0f; + const float mean_prev = isv[mean_slot]; + const float delta = x - mean_prev; + const float mean_new = mean_prev + delta / n_new; + const float delta2 = x - mean_new; + const float m2_new = isv[m2_slot] + delta * delta2; + + isv[count_slot] = n_new; + isv[mean_slot] = mean_new; + isv[m2_slot] = m2_new; +} +``` + +- [ ] **Step 2: Register cubin in build.rs** + +Add to `build.rs` cubin list (next to other RL controller cubins). + +```rust +// In the cubin compile list: +"rl_signal_variance_update", +``` + +- [ ] **Step 3: Run `cargo check` to verify** + +Run: `SQLX_OFFLINE=true cargo check -p ml-alpha 2>&1 | tail -3` +Expected: `Finished` cleanly. + +--- + +### Task 3: Welford kernel unit tests + +**Files:** +- Create: `crates/ml-alpha/tests/signal_variance_kernel.rs` + +- [ ] **Step 1: Write G1 — convergence on constant input** + +```rust +//! GPU-oracle tests for rl_signal_variance_update. +//! +//! Per feedback_no_cpu_test_fallbacks: oracles are analytical identities of +//! Welford online variance (constant → 0 var; finite sequence → known var). + +use cudarc::driver::{LaunchConfig, PushKernelArg}; +use ml_alpha::pinned_mem::MappedF32Buffer; +use ml_core::device::MlDevice; + +const INPUT_SLOT: usize = 0; +const COUNT_SLOT: usize = 1; +const MEAN_SLOT: usize = 2; +const M2_SLOT: usize = 3; + +#[test] +#[ignore = "requires CUDA"] +fn welford_constant_input_zero_variance() { + let Ok(dev) = MlDevice::cuda(0) else { return }; + let isv = unsafe { MappedF32Buffer::new(4) }.expect("isv"); + let module = dev.load_cubin(include_bytes!( + concat!(env!("OUT_DIR"), "/rl_signal_variance_update.cubin") + ).to_vec()).expect("load"); + let func = module.load_function("rl_signal_variance_update").expect("fn"); + + isv.write_record(INPUT_SLOT, 0.5); + let stream = dev.cuda_stream().expect("stream").clone(); + for _ in 0..50 { + let mut builder = stream.launch_builder(&func); + builder.arg(&isv.dev_ptr) + .arg(&(INPUT_SLOT as i32)) + .arg(&(COUNT_SLOT as i32)) + .arg(&(MEAN_SLOT as i32)) + .arg(&(M2_SLOT as i32)); + unsafe { builder.launch(LaunchConfig::for_num_elems(1)) }.expect("launch"); + } + stream.synchronize().expect("sync"); + assert_eq!(isv.read_record(COUNT_SLOT), 50.0); + assert!((isv.read_record(MEAN_SLOT) - 0.5).abs() < 1e-6); + assert!(isv.read_record(M2_SLOT).abs() < 1e-6); +} +``` + +- [ ] **Step 2: Write G2 — known variance on finite sequence** + +```rust +#[test] +#[ignore = "requires CUDA"] +fn welford_known_sequence_variance() { + let Ok(dev) = MlDevice::cuda(0) else { return }; + let isv = unsafe { MappedF32Buffer::new(4) }.expect("isv"); + let module = dev.load_cubin(include_bytes!( + concat!(env!("OUT_DIR"), "/rl_signal_variance_update.cubin") + ).to_vec()).expect("load"); + let func = module.load_function("rl_signal_variance_update").expect("fn"); + let stream = dev.cuda_stream().expect("stream").clone(); + + // Sequence [1, 2, 3, 4, 5]: mean = 3, var = (4+1+0+1+4)/4 = 2.5 (sample variance). + for x in [1.0_f32, 2.0, 3.0, 4.0, 5.0] { + isv.write_record(INPUT_SLOT, x); + let mut builder = stream.launch_builder(&func); + builder.arg(&isv.dev_ptr) + .arg(&(INPUT_SLOT as i32)) + .arg(&(COUNT_SLOT as i32)) + .arg(&(MEAN_SLOT as i32)) + .arg(&(M2_SLOT as i32)); + unsafe { builder.launch(LaunchConfig::for_num_elems(1)) }.expect("launch"); + stream.synchronize().expect("sync"); + } + let n = isv.read_record(COUNT_SLOT); + let m2 = isv.read_record(M2_SLOT); + let sample_var = m2 / (n - 1.0); + assert!((sample_var - 2.5).abs() < 1e-5, + "expected sample variance 2.5, got {}", sample_var); +} +``` + +- [ ] **Step 3: Write G3 — sentinel skip** + +```rust +#[test] +#[ignore = "requires CUDA"] +fn welford_sentinel_zero_is_skipped() { + let Ok(dev) = MlDevice::cuda(0) else { return }; + let isv = unsafe { MappedF32Buffer::new(4) }.expect("isv"); + let module = dev.load_cubin(include_bytes!( + concat!(env!("OUT_DIR"), "/rl_signal_variance_update.cubin") + ).to_vec()).expect("load"); + let func = module.load_function("rl_signal_variance_update").expect("fn"); + let stream = dev.cuda_stream().expect("stream").clone(); + + // INPUT_SLOT stays 0.0 (sentinel). Launch should be a no-op. + let mut builder = stream.launch_builder(&func); + builder.arg(&isv.dev_ptr) + .arg(&(INPUT_SLOT as i32)) + .arg(&(COUNT_SLOT as i32)) + .arg(&(MEAN_SLOT as i32)) + .arg(&(M2_SLOT as i32)); + unsafe { builder.launch(LaunchConfig::for_num_elems(1)) }.expect("launch"); + stream.synchronize().expect("sync"); + assert_eq!(isv.read_record(COUNT_SLOT), 0.0, "sentinel must not be counted"); +} +``` + +- [ ] **Step 4: Run all three tests** + +Run: `SQLX_OFFLINE=true cargo test -p ml-alpha --test signal_variance_kernel --release -- --ignored --nocapture` +Expected: 3 passed. + +--- + +### Task 4: Update `rl_advantage_normalize.cu` to emit `var_pre_norm` + +**Files:** +- Modify: `crates/ml-alpha/cuda/rl_advantage_normalize.cu` + +- [ ] **Step 1: Add the ISV emission BEFORE the in-place normalize** + +In the existing `rl_advantage_normalize` kernel (after the variance computation, before the normalize pass), write `var` to `isv[RL_ADV_VAR_PRE_NORM_INDEX]`: + +```c +// After: const float var = s_var[0] / (float)B; +if (tid == 0) { + isv[RL_ADV_VAR_PRE_NORM_INDEX] = var; // NEW: emit raw variance for rollout_steps_controller + s_inv_std = rsqrtf(var + ADVANTAGE_VAR_FLOOR); +} +``` + +Add `RL_ADV_VAR_PRE_NORM_INDEX` define at top of file: +```c +#define RL_ADV_VAR_PRE_NORM_INDEX 612 +``` + +Add `float* __restrict__ isv` to the kernel signature; thread the ISV device pointer through from the trainer call site. + +- [ ] **Step 2: Update trainer call site to pass `isv_dev_ptr`** + +Modify `integrated.rs` where `rl_advantage_normalize` is launched. Add `args.push_ptr(self.isv_dev_ptr)` to the RawArgs. + +- [ ] **Step 3: Verify with smoke** + +Run: `SQLX_OFFLINE=true cargo test -p ml-alpha --test integrated_trainer_smoke --release -- --ignored --nocapture` +Expected: passes (existing GPU smoke). + +--- + +### Task 5: Refactor `rl_ppo_clip_controller.cu` (canonical pattern) + +**Files:** +- Modify: `crates/ml-alpha/cuda/rl_ppo_clip_controller.cu` + +- [ ] **Step 1: Replace hardcoded `KL_NOISE_FLOOR_FRAC` with adaptive floor** + +```c +// BEFORE: +#define KL_NOISE_FLOOR_FRAC 0.01f +const float kl_noise_floor = kl_target * KL_NOISE_FLOOR_FRAC; + +// AFTER: +const float kl_var_count = isv[RL_KL_PI_VAR_COUNT_INDEX]; +const float kl_var = (kl_var_count > 1.0f) + ? isv[RL_KL_PI_VAR_M2_INDEX] / (kl_var_count - 1.0f) + : 0.0f; +const float kl_std = sqrtf(kl_var); +const float kl_noise_floor = fmaxf(kl_target * 0.5f, kl_std * 2.0f); +``` + +Add at top of file: +```c +#define RL_KL_PI_VAR_COUNT_INDEX 588 +#define RL_KL_PI_VAR_MEAN_INDEX 589 +#define RL_KL_PI_VAR_M2_INDEX 590 +#define RL_KL_PI_BELOW_COUNT_INDEX 591 +``` + +- [ ] **Step 2: Add asymmetric Schulman (widen requires N consecutive below-band)** + +```c +// In the Schulman-style adjustment block: +if (kl_ema > kl_target * tolerance) { + ratio = 1.0f / adjust_rate; // tighten — fires immediately + isv[RL_KL_PI_BELOW_COUNT_INDEX] = 0.0f; +} else if (kl_ema < kl_target / tolerance) { + isv[RL_KL_PI_BELOW_COUNT_INDEX] += 1.0f; + ratio = (isv[RL_KL_PI_BELOW_COUNT_INDEX] >= 3.0f) ? adjust_rate : 1.0f; +} else { + isv[RL_KL_PI_BELOW_COUNT_INDEX] = 0.0f; + ratio = 1.0f; +} +``` + +- [ ] **Step 3: Bootstrap PPO clip at target (not at safety bound)** + +Change the bootstrap `RL_EPS_BOOTSTRAP_INDEX = 474` value in trainer init from 0.2 to **derive from target**: bootstrap = ISV[RL_KL_TARGET_INDEX] × some multiplier that puts ε in a conservative starting position. Decision deferred to implementation — for now, set bootstrap to `isv[RL_KL_TARGET_INDEX]` (so ε starts at 0.01 = target, very conservative). + +This is a trainer-side change in `with_controllers_bootstrapped`: +```rust +// In trainer/integrated.rs with_controllers_bootstrapped: +trainer.isv_mapped.write_record(RL_EPS_BOOTSTRAP_INDEX, trainer.isv_mapped.read_record(RL_KL_TARGET_INDEX)); +``` + +- [ ] **Step 4: Verify ppo_clip controller compiles** + +Run: `SQLX_OFFLINE=true cargo build -p ml-alpha --release 2>&1 | tail -3` +Expected: `Finished` clean. + +--- + +### Task 6: Refactor remaining Group A controllers (sweep) + +**Files (all in `crates/ml-alpha/cuda/`):** +- `rl_target_tau_controller.cu` +- `rl_rollout_steps_controller.cu` (also consume `RL_ADV_VAR_PRE_NORM_INDEX` as the input signal) +- `rl_entropy_coef_controller.cu` +- `rl_per_alpha_controller.cu` + +- [ ] **Step 1-4: Apply Task 5's pattern to each** + +For each of the 4 controllers above: same three changes as PPO clip (adaptive noise floor, asymmetric Schulman, conservative bootstrap). Use the slot table from spec to wire variance triples per controller. + +- [ ] **Step 5: Verify all compile** + +Run: `SQLX_OFFLINE=true cargo build -p ml-alpha --release 2>&1 | tail -3` +Expected: `Finished` clean. + +--- + +### Task 7: Special case G — `rl_gamma_controller.cu` + +**Files:** +- Modify: `crates/ml-alpha/cuda/rl_gamma_controller.cu` + +- [ ] **Step 1: Replace hardcoded `GAMMA_MIN` with Welford-mean-driven adaptive floor** + +Add at top: +```c +#define RL_GAMMA_MIN_ADAPTIVE_INDEX 613 +#define RL_MEAN_TRADE_DURATION_EMA_INDEX 417 +#define RL_TRADE_DUR_VAR_COUNT_INDEX 608 // Welford count for trade duration +#define RL_TRADE_DUR_VAR_MEAN_INDEX 609 // Welford mean for trade duration +#define GAMMA_MIN_ABSOLUTE 0.9f // "don't degenerate to bandit" hard floor +#define HORIZON_MULTIPLIER 2.0f +#define WELFORD_WARMUP_OBS 100.0f // confidence gate, not magnitude calibration +``` + +Compute adaptive min at start of controller body. Per spec Q6 resolution, use the Welford mean (which is already computed by `rl_signal_variance_update` for `rl_per_alpha_controller`'s noise floor) once enough observations have accumulated. The Welford mean's natural lag (1/N smoothing) breaks the feedback loop without any step-count gating: + +```c +const float d_welford_count = isv[RL_TRADE_DUR_VAR_COUNT_INDEX]; +const float d_welford_mean = isv[RL_TRADE_DUR_VAR_MEAN_INDEX]; + +// Warmup: use fast EMA until 100 trade observations accumulated. +// After warmup, use Welford mean (denominator-of-N smoothed, naturally +// resists the gamma↔trade_duration feedback loop without explicit lag). +const float d_smoothed = (d_welford_count >= WELFORD_WARMUP_OBS) + ? d_welford_mean + : isv[RL_MEAN_TRADE_DURATION_EMA_INDEX]; + +const float adaptive_min = fmaxf(GAMMA_MIN_ABSOLUTE, + 1.0f - 1.0f / fmaxf(d_smoothed * HORIZON_MULTIPLIER, 10.0f)); +isv[RL_GAMMA_MIN_ADAPTIVE_INDEX] = adaptive_min; +``` + +Replace `GAMMA_MIN` usage with `isv[RL_GAMMA_MIN_ADAPTIVE_INDEX]` in the clamp: +```c +gamma = fmaxf(isv[RL_GAMMA_MIN_ADAPTIVE_INDEX], fminf(gamma, GAMMA_MAX)); +``` + +- [ ] **Step 2: Bootstrap `RL_GAMMA_MIN_ADAPTIVE_INDEX` in trainer** + +In `with_controllers_bootstrapped`: +```rust +trainer.isv_mapped.write_record(RL_GAMMA_MIN_ADAPTIVE_INDEX, 0.95); // safe default before first observation +``` + +- [ ] **Step 3: Compile + verify** + +Run: `SQLX_OFFLINE=true cargo build -p ml-alpha --release 2>&1 | tail -3` +Expected: `Finished`. + +--- + +### Task 8: Special case R — `rl_reward_scale_controller.cu` + +**Files:** +- Modify: `crates/ml-alpha/cuda/rl_reward_scale_controller.cu` + +- [ ] **Step 1: Add asymmetric rate limit (decrease slower than increase)** + +The crash from 1.0 → 0.004 in 100 steps was the decrease side. Cap per-step decrease at 1.05× (i.e., new ≥ prev / 1.05): + +```c +// After computing new_scale: +const float prev_scale = isv[RL_REWARD_SCALE_INDEX]; +if (new_scale < prev_scale) { + new_scale = fmaxf(new_scale, prev_scale / 1.05f); // decrease ≤ 5% per step +} +``` + +- [ ] **Step 2: Add bootstrap-fraction floor (until N trades accumulated)** + +```c +#define BOOTSTRAP_FRACTION 0.1f +#define MIN_TRADES_FOR_RELEASE 100 +const float trade_count = isv[RL_TRADE_COUNT_INDEX]; // verify slot exists or use closest signal +const float boot = isv[RL_REWARD_SCALE_BOOTSTRAP_INDEX]; +if (trade_count < (float)MIN_TRADES_FOR_RELEASE) { + new_scale = fmaxf(new_scale, boot * BOOTSTRAP_FRACTION); +} +``` + +- [ ] **Step 3: Emit `RL_REWARD_MAGNITUDE_EMA_INDEX` from this controller** + +Already computed internally as `pnl_per_batch_magnitude`. Write to `isv[RL_REWARD_MAGNITUDE_EMA_INDEX]` so the Welford variance kernel can track it. + +- [ ] **Step 4: Compile + verify** + +--- + +### Task 9: Special case Q — `rl_q_distill_lambda_controller.cu` + +**Files:** +- Modify: `crates/ml-alpha/cuda/rl_q_distill_lambda_controller.cu` + +- [ ] **Step 1: Replace `MIN_LAMBDA` with adaptive bound** + +```c +#define RL_Q_DISTILL_KL_VAR_COUNT_INDEX 618 +#define RL_Q_DISTILL_KL_VAR_M2_INDEX 620 +#define RL_Q_DISTILL_LAMBDA_MIN_ADAPTIVE_INDEX 639 + +const float kl_var_count = isv[RL_Q_DISTILL_KL_VAR_COUNT_INDEX]; +const float kl_var = (kl_var_count > 1.0f) + ? isv[RL_Q_DISTILL_KL_VAR_M2_INDEX] / (kl_var_count - 1.0f) + : 0.0f; +const float kl_std = sqrtf(kl_var); +const float adaptive_min = fmaxf(0.001f, kl_std * 0.05f); +isv[RL_Q_DISTILL_LAMBDA_MIN_ADAPTIVE_INDEX] = adaptive_min; +``` + +Replace `MIN_LAMBDA` usage in the clamp with `isv[RL_Q_DISTILL_LAMBDA_MIN_ADAPTIVE_INDEX]`. + +- [ ] **Step 2: Bootstrap the adaptive min** + +In `with_controllers_bootstrapped`: +```rust +trainer.isv_mapped.write_record(RL_Q_DISTILL_LAMBDA_MIN_ADAPTIVE_INDEX, 0.01); // safer default than 0.05 +``` + +- [ ] **Step 3: Compile + verify** + +--- + +### Task 10: Group B — `rl_v_blend_alpha_controller.cu` + +**Files:** +- Modify: `crates/ml-alpha/cuda/rl_v_blend_alpha_controller.cu` + +- [ ] **Step 1: Replace 5 hardcoded constants with ISV slot reads** + +```c +// BEFORE: +#define EMA_ALPHA 0.01f +#define TARGET_TRACK_RATIO 0.10f +#define SCHULMAN_STEP 0.01f +#define DEAD_SIGNAL_FLOOR 1e-4f +#define BOOTSTRAP_ALPHA 1.0f + +// AFTER (use as runtime reads): +const float ema_alpha = isv[RL_V_BLEND_EMA_ALPHA_INDEX]; +const float target_track_ratio = isv[RL_V_BLEND_TARGET_TRACK_RATIO_INDEX]; +const float schulman_step = isv[RL_V_BLEND_SCHULMAN_STEP_INDEX]; +const float dead_signal_floor = isv[RL_V_BLEND_DEAD_SIGNAL_FLOOR_INDEX]; +const float bootstrap_alpha = isv[RL_V_BLEND_BOOTSTRAP_ALPHA_INDEX]; +``` + +DEAD_SIGNAL_FLOOR should adapt to observed V_scalar magnitude variance: +```c +const float vmag_var_count = isv[RL_V_BLEND_SCALAR_VAR_COUNT_INDEX]; +const float vmag_var = (vmag_var_count > 1.0f) + ? isv[RL_V_BLEND_SCALAR_VAR_M2_INDEX] / (vmag_var_count - 1.0f) + : 0.0f; +const float adaptive_dead_floor = fmaxf(isv[RL_V_BLEND_DEAD_SIGNAL_FLOOR_INDEX], + sqrtf(vmag_var) * 0.1f); +if (mag_ema < adaptive_dead_floor) { + // hold α + ... +} +``` + +- [ ] **Step 2: Bootstrap the 5 new ISV slots in trainer** + +```rust +trainer.isv_mapped.write_record(RL_V_BLEND_EMA_ALPHA_INDEX, 0.01); +trainer.isv_mapped.write_record(RL_V_BLEND_TARGET_TRACK_RATIO_INDEX, 0.10); +trainer.isv_mapped.write_record(RL_V_BLEND_SCHULMAN_STEP_INDEX, 0.01); +trainer.isv_mapped.write_record(RL_V_BLEND_DEAD_SIGNAL_FLOOR_INDEX, 1e-4); +trainer.isv_mapped.write_record(RL_V_BLEND_BOOTSTRAP_ALPHA_INDEX, 1.0); +``` + +- [ ] **Step 3: Compile + verify** + +--- + +### Task 11: Group B — `rl_ppo_ratio_clamp_controller.cu` + +**Files:** +- Modify: `crates/ml-alpha/cuda/rl_ppo_ratio_clamp_controller.cu` + +- [ ] **Step 1: Replace `MIN_OUT` / `MAX_OUT` with adaptive bounds** + +```c +const float ratio_var_count = isv[RL_PPO_RATIO_VAR_COUNT_INDEX]; +const float ratio_var = (ratio_var_count > 1.0f) + ? isv[RL_PPO_RATIO_VAR_M2_INDEX] / (ratio_var_count - 1.0f) + : 0.0f; +const float ratio_std = sqrtf(ratio_var); +const float adaptive_min = fmaxf(1.5f, 1.0f + ratio_std * 3.0f); // 3σ above unit, but ≥ 1.5 +const float adaptive_max = fminf(1000.0f, fmaxf(adaptive_min * 5.0f, isv[RL_PPO_RATIO_CLAMP_MAX_ADAPTIVE_INDEX])); +clamp = fmaxf(adaptive_min, fminf(clamp, adaptive_max)); +``` + +- [ ] **Step 2: Compile + verify** + +--- + +### Task 12: Group B — `rl_reward_clamp_controller.cu` + +**Files:** +- Modify: `crates/ml-alpha/cuda/rl_reward_clamp_controller.cu` + +- [ ] **Step 1: Replace 5 hardcoded constants with ISV reads** + +```c +const float v_bound_floor = isv[RL_REWARD_CLAMP_V_BOUND_FLOOR_INDEX]; +const float v_bound_ewma_alpha = isv[RL_REWARD_CLAMP_V_BOUND_EWMA_ALPHA_INDEX]; +const float min_win = isv[RL_REWARD_CLAMP_MIN_WIN_INDEX]; +const float min_ratio = isv[RL_REWARD_CLAMP_MIN_RATIO_INDEX]; +const float max_ratio = isv[RL_REWARD_CLAMP_MAX_RATIO_INDEX]; +``` + +- [ ] **Step 2: Bootstrap the 5 new ISV slots** + +In trainer init, set defaults to current hardcoded values: +```rust +trainer.isv_mapped.write_record(RL_REWARD_CLAMP_V_BOUND_FLOOR_INDEX, 1.0); +trainer.isv_mapped.write_record(RL_REWARD_CLAMP_V_BOUND_EWMA_ALPHA_INDEX, 0.001); +trainer.isv_mapped.write_record(RL_REWARD_CLAMP_MIN_WIN_INDEX, 1.0); +trainer.isv_mapped.write_record(RL_REWARD_CLAMP_MIN_RATIO_INDEX, 1.0); +trainer.isv_mapped.write_record(RL_REWARD_CLAMP_MAX_RATIO_INDEX, 3.0); +``` + +- [ ] **Step 3: Compile + verify** + +--- + +### Task 13: Group B — `rl_gate_threshold_controller.cu` + +**Files:** +- Modify: `crates/ml-alpha/cuda/rl_gate_threshold_controller.cu` + +- [ ] **Step 1: Replace `alpha = 0.01f` with ISV slot** + +```c +// BEFORE: +const float alpha = 0.01f; + +// AFTER: +const float alpha = isv[RL_GATE_EMA_ALPHA_INDEX]; +``` + +- [ ] **Step 2: Bootstrap in trainer** + +```rust +trainer.isv_mapped.write_record(RL_GATE_EMA_ALPHA_INDEX, 0.01); +``` + +- [ ] **Step 3: Compile + verify** + +--- + +### Task 14: Update `rl_fused_controllers.cu` + +**Files:** +- Modify: `crates/ml-alpha/cuda/rl_fused_controllers.cu` + +- [ ] **Step 1: Replicate the per-controller adaptive logic** + +The fused kernel is what the trainer actually launches per-step. Mirror every change from Tasks 5-13 inside the corresponding branch of the fused kernel. Same ISV slot reads, same adaptive math. + +- [ ] **Step 2: Compile + verify both individual controllers AND fused match** + +Run: `SQLX_OFFLINE=true cargo build -p ml-alpha --release 2>&1 | tail -3` +Expected: `Finished` clean. + +--- + +### Task 15: Wire trainer integration + +**Files:** +- Modify: `crates/ml-alpha/src/trainer/integrated.rs` + +- [ ] **Step 1: Add per-step Welford launches** + +After the existing EMA producer kernels (R3 `launch_ema_update_per_step` etc.) and BEFORE the controllers fire: + +```rust +// Launch rl_signal_variance_update for each controller input. +for (input_slot, count_slot, mean_slot, m2_slot) in [ + (RL_KL_PI_EMA_INDEX, RL_KL_PI_VAR_COUNT_INDEX, RL_KL_PI_VAR_MEAN_INDEX, RL_KL_PI_VAR_M2_INDEX), + (RL_Q_DIVERGENCE_EMA_INDEX, RL_Q_DIV_VAR_COUNT_INDEX, RL_Q_DIV_VAR_MEAN_INDEX, RL_Q_DIV_VAR_M2_INDEX), + (RL_ADV_VAR_PRE_NORM_INDEX, RL_ADV_VAR_VAR_COUNT_INDEX, RL_ADV_VAR_VAR_MEAN_INDEX, RL_ADV_VAR_VAR_M2_INDEX), + (RL_ENTROPY_OBSERVED_EMA_INDEX, RL_ENTROPY_OBS_VAR_COUNT_INDEX, RL_ENTROPY_OBS_VAR_MEAN_INDEX, RL_ENTROPY_OBS_VAR_M2_INDEX), + (RL_TD_KURTOSIS_EMA_INDEX, RL_TD_KURT_VAR_COUNT_INDEX, RL_TD_KURT_VAR_MEAN_INDEX, RL_TD_KURT_VAR_M2_INDEX), + (RL_MEAN_TRADE_DURATION_EMA_INDEX, RL_TRADE_DUR_VAR_COUNT_INDEX, RL_TRADE_DUR_VAR_MEAN_INDEX, RL_TRADE_DUR_VAR_M2_INDEX), + (RL_REWARD_MAGNITUDE_EMA_INDEX, RL_REWARD_MAG_VAR_COUNT_INDEX, RL_REWARD_MAG_VAR_MEAN_INDEX, RL_REWARD_MAG_VAR_M2_INDEX), + (RL_Q_DISTILL_KL_EMA_INDEX, RL_Q_DISTILL_KL_VAR_COUNT_INDEX, RL_Q_DISTILL_KL_VAR_MEAN_INDEX, RL_Q_DISTILL_KL_VAR_M2_INDEX), + (RL_V_SCALAR_MAG_EMA_INDEX, RL_V_BLEND_SCALAR_VAR_COUNT_INDEX, RL_V_BLEND_SCALAR_VAR_MEAN_INDEX, RL_V_BLEND_SCALAR_VAR_M2_INDEX), + // PPO ratio: input from observed log_ratio after each batch — use a magnitude EMA producer + // (omitted: add a small magnitude-tracking kernel if not already present) +] { + let mut args = RawArgs::new(); + args.push_ptr(self.isv_dev_ptr); + args.push_i32(input_slot as i32); + args.push_i32(count_slot as i32); + args.push_i32(mean_slot as i32); + args.push_i32(m2_slot as i32); + let mut ptrs = args.build_arg_ptrs(); + unsafe { + raw_launch( + self.rl_signal_variance_update_fn.cu_function(), + (1, 1, 1), (1, 1, 1), 0, + self.raw_stream, + &mut ptrs[..args.len()], + ).map_err(|e| anyhow::anyhow!("rl_signal_variance_update: {:?}", e))?; + } +} +``` + +- [ ] **Step 2: Bootstrap all new ISV slots in `with_controllers_bootstrapped`** + +Aggregate all the `write_record` calls from Tasks 5-13 in one block at the end of `with_controllers_bootstrapped`. + +- [ ] **Step 3: Run integrated_trainer_smoke** + +Run: `SQLX_OFFLINE=true cargo test -p ml-alpha --test integrated_trainer_smoke --release -- --ignored --nocapture` +Expected: passes. + +--- + +### Task 16: Controller-level adaptive-floor tests + +**Files:** +- Create: `crates/ml-alpha/tests/controller_adaptive_floors.rs` + +- [ ] **Step 1: G3 — PPO clip holds on small signal** + +Bootstrap a trainer. Feed `kl_ema = 0.001` (below target 0.01) and observed_var = 1e-7 for 50 steps. Assert `ε` stays in `[bootstrap × 0.9, bootstrap × 1.1]` (controller held — didn't drift to MAX). + +- [ ] **Step 2: G4 — PPO clip tightens on big signal** + +Single `kl_ema = 0.05` observation (5× target). Assert `ε` decreases by `1/adjust_rate` immediately (asymmetric tightening fires on single observation). + +- [ ] **Step 3: G5 — PPO clip widens slowly (N=3 consecutive below)** + +Feed `kl_ema = 0.001` for 10 steps. Assert widen fires only after 3 consecutive below-band observations. + +- [ ] **Step 4: G6 — gamma adaptive_min tracks trade duration** + +Set `RL_MEAN_TRADE_DURATION_EMA_INDEX = 30.0`, advance step counter to 500, assert `RL_GAMMA_MIN_ADAPTIVE_INDEX ≈ 1 - 1/60 = 0.983`. + +- [ ] **Step 5: Update isv_bootstrap.rs test** + +Add the 52 new slots to the canonical-bootstrap assertion list. Ensure values match what `with_controllers_bootstrapped` writes. + +- [ ] **Step 6: Run all gates** + +Run: `SQLX_OFFLINE=true cargo test -p ml-alpha --test controller_adaptive_floors --release -- --ignored --nocapture` +Expected: 5+ passed. + +--- + +### Task 17: Local smoke validation + +- [ ] **Step 1: 1k-step smoke** + +```bash +cd /home/jgrusewski/Work/foxhunt +SQLX_OFFLINE=true cargo build -p ml-alpha --release --example alpha_rl_train +./target/release/examples/alpha_rl_train \ + --mbp10-data-dir test_data/futures-baseline/ES.FUT \ + --predecoded-dir /tmp/smoke/predecoded \ + --out /tmp/smoke \ + --n-steps 1000 --n-backtests 128 --seq-len 32 \ + --per-capacity 1024 --seed 16962 --instrument-mode all \ + --n-folds 1 --fold-idx 0 --n-eval-steps 0 \ + --diag-jsonl /tmp/smoke/diag.jsonl +``` + +Expected: completes cleanly, no NaN abort. + +- [ ] **Step 2: compute-sanitizer at b=128** + +```bash +compute-sanitizer --tool memcheck --launch-timeout 600 --error-exitcode 1 \ + ./target/release/examples/alpha_rl_train \ + --mbp10-data-dir test_data/futures-baseline/ES.FUT \ + --predecoded-dir /tmp/smoke-cs/predecoded \ + --out /tmp/smoke-cs \ + --n-steps 5 --n-backtests 128 --seq-len 32 \ + --per-capacity 1024 --seed 16962 --instrument-mode all \ + --n-folds 1 --fold-idx 0 --n-eval-steps 0 +``` + +Expected: `ERROR SUMMARY: 0 errors`. + +- [ ] **Step 3: Diagnose controllers on 1k smoke** + +Inspect `/tmp/smoke/diag.jsonl` at steps 1, 100, 500, 1000. All 12 controllers should: +- Be at sensible values (not stuck at MIN/MAX) +- Move in response to observed signal +- Welford variance triples grow over time + +--- + +### Task 18: Atomic commit + +- [ ] **Step 1: Verify all changes compile + tests pass** + +```bash +SQLX_OFFLINE=true cargo build -p ml-alpha --release --quiet +SQLX_OFFLINE=true cargo build -p ml-alpha --tests --quiet +SQLX_OFFLINE=true cargo test -p ml-alpha --test integrated_trainer_smoke --release -- --ignored --nocapture +SQLX_OFFLINE=true cargo test -p ml-alpha --test signal_variance_kernel --release -- --ignored --nocapture +SQLX_OFFLINE=true cargo test -p ml-alpha --test controller_adaptive_floors --release -- --ignored --nocapture +``` + +Expected: all pass. + +- [ ] **Step 2: Single atomic commit** + +```bash +git add crates/ml-alpha/ +git commit -m "feat(rl): adaptive controller floors — 12 controllers, signal-driven thresholds + +Replaces hardcoded thresholds (NOISE_FLOOR_FRAC, MIN/MAX bounds, +DEAD_SIGNAL_FLOOR, etc.) across 12 RL controllers with observed-signal- +driven ISV-slot bounds. Eliminates the architectural failure mode +documented in fold 0 walk-forward (alpha-rl-m9cx5, eval pnl -\$1.56M) +where Phase 4.5 advantage normalization invalidated 8 controllers' +hardcoded signal-scale assumptions simultaneously, causing all of them +to saturate at extrema (ppo_clip at MAX 0.50, q_distill_lambda at MIN +0.05, gamma at MIN 0.995, reward_scale crashed 250× in 100 steps, etc.). + +Per feedback_adaptive_not_tuned, feedback_isv_for_adaptive_bounds, +pearl_controller_anchors_isv_driven: every controller threshold now +derives from observed signal statistics (Welford online variance) +rather than constants calibrated against a prior signal regime. + +Spec: docs/superpowers/specs/2026-05-30-adaptive-controller-floor-design.md +Plan: docs/superpowers/plans/2026-05-30-adaptive-controller-floor-plan.md + +New: rl_signal_variance_update.cu (Welford kernel, 80 LOC) +Modified: 12 controller .cu files + rl_fused_controllers.cu + + rl_advantage_normalize.cu (emit var_pre_norm) + + isv_slots.rs (+52 slots, RL_SLOTS_END 588→640) + + integrated.rs (Welford launches + bootstrap writes) + +Validation: + - 3 GPU-oracle Welford kernel tests pass + - 5 controller adaptive-floor invariant tests pass + - integrated_trainer_smoke (full GPU pipeline) passes + - compute-sanitizer memcheck b=128: 0 errors + - 1k-step local smoke: all 12 controllers move with signal, none stuck + +Cluster validation (G7/G8/G9) submitted as follow-up runs. + +🤖 Generated with [Claude Code](https://claude.com/claude-code) + +Co-Authored-By: Claude Opus 4.7 " +``` + +- [ ] **Step 3: Push to origin/main** + +```bash +git push origin main +``` + +Expected: pushed cleanly. + +--- + +### Task 19: Cluster G7 — fold 0 re-run with fix + +- [ ] **Step 1: Submit fold 0 walk-forward at new SHA** + +```bash +./scripts/argo-alpha-rl.sh \ + --sha --branch main --gpu-pool ci-training-l40s \ + --n-steps 20000 --seq-len 32 --n-backtests 1024 --per-capacity 32768 \ + --seed 16962 --instrument-mode all \ + --n-folds 3 --fold-idx 0 --n-eval-steps 2000 \ + --skip-push-check +``` + +- [ ] **Step 2: Verify controller release at step 19999** + +Inspect diag.jsonl. All 8 Group A controllers should be off their extrema. PPO ε ≤ 0.20, reward_scale ≥ 0.5, gamma > 0.97, q_distill_lambda > 0.05. + +- [ ] **Step 3: G8 — eval pnl improvement check** + +Pull `eval_summary.json`. Assert eval pnl > -$0.5M, profit_factor > 0.85, win_rate > 0.40. + +--- + +### Task 20: Cluster G9 — ksll2 regression test + +- [ ] **Step 1: Submit single-window 20k at new SHA** + +Same params as the original `alpha-rl-ksll2` (n_folds=1, n_eval_steps=0). + +- [ ] **Step 2: G9 verdict** + +Final pnl ≥ +$15M (within 20% of original +$18.3M). Confirms fix doesn't regress large-data regime. + +--- + +## Self-review + +This plan was reviewed against the spec at `docs/superpowers/specs/2026-05-30-adaptive-controller-floor-design.md`. Cross-checks: + +- **Coverage**: every controller in spec scope (12) appears as a task. Group A in Tasks 5-9 (per-controller patterns), Group B in Tasks 10-13. The fused kernel update (Task 14) replicates per-controller logic. ✓ +- **Placeholder scan**: every step has a code block or exact command. No "TBD"/"TODO"/"implement details". ✓ +- **Type consistency**: ISV slot names match exactly between Task 1 (allocation) and Tasks 5-13 (consumption). ✓ +- **Atomic commit discipline**: only Task 18 commits. All earlier tasks build incrementally without committing. ✓ + +Known open items (from spec Q6): the gamma adaptive_min update period (500 steps proposed). Implementation per Task 7 uses 500 — if cluster G7 reveals feedback-loop issues, adjust to exponential cool-down in a follow-up. + +--- + +## Execution + +Use **subagent-driven-development** (recommended) — dispatch a fresh subagent per task, two-stage review (spec compliance, code quality), atomic single commit at Task 18. diff --git a/docs/superpowers/specs/2026-05-30-adaptive-controller-floor-design.md b/docs/superpowers/specs/2026-05-30-adaptive-controller-floor-design.md new file mode 100644 index 000000000..8dba253b8 --- /dev/null +++ b/docs/superpowers/specs/2026-05-30-adaptive-controller-floor-design.md @@ -0,0 +1,430 @@ +# Adaptive Controller Signal Floors — Design + +**Goal**: every adaptive RL controller's noise floor / target / threshold derives from observed signal statistics, not a hardcoded constant calibrated against a prior signal regime. + +**Why now**: Phase 4.5 advantage normalization (commit `669578566`) shifted the operating distribution of every magnitude-dependent controller input — KL, advantage variance, TD kurtosis, Q divergence. Fold 0 walk-forward at b=1024 20k on 1/3 the ksll2 data confirmed **all 7 adaptive controllers saturate or crash**, even though release training (ksll2, full 9 files) appeared healthy. The system architecture has 7 independent controllers, each with hardcoded thresholds. Phase 4.5 invalidated all of them simultaneously. + +**Scope** (revised 2026-05-30 to apply the principle consistently — including clamp bounds, not just noise floors): every adaptive RL controller with **any hardcoded threshold or clamp-bound constant** is in scope. This includes: +- Noise floors and saturation thresholds (the original Phase 4.5 saturation cause) +- Clamp bounds (MIN/MAX on controller output range) +- Step / ramp / decay rates +- EMA / Wiener-α coefficients +- Bootstrap values + +This is **12 controllers** in one atomic commit per `feedback_no_partial_refactor`: + +**Group A — 8 controllers that visibly saturated in fold 0:** +1. `rl_ppo_clip_controller` — stuck at MAX 0.50 (confirmed cause) +2. `rl_target_tau_controller` — never moved from bootstrap +3. `rl_rollout_steps_controller` — never moved (advantage_var_ratio definitionally near-zero under Phase 4.5) +4. `rl_entropy_coef_controller` — stuck at MIN 0.01 +5. `rl_per_alpha_controller` — stuck at MIN 0.40 +6. `rl_gamma_controller` — stuck at MIN 0.995 (hardcoded `GAMMA_MIN`) +7. `rl_reward_scale_controller` — crashed 250× in 100 steps +8. `rl_q_distill_lambda_controller` — stuck at MIN 0.05 (hardcoded `MIN_LAMBDA`) + +**Group B — 4 controllers with hardcoded thresholds, didn't visibly saturate yet but vulnerable to the same architectural failure**: +9. `rl_v_blend_alpha_controller` — 5 hardcoded constants (`DEAD_SIGNAL_FLOOR`, `TARGET_TRACK_RATIO`, `SCHULMAN_STEP`, `EMA_ALPHA`, `BOOTSTRAP_ALPHA`) +10. `rl_ppo_ratio_clamp_controller` — `PPO_RATIO_CLAMP_MIN_OUT = 2.0f` (calibrated for pre-normalization ratios — Phase 4.5 may already have invalidated) +11. `rl_reward_clamp_controller` — `V_BOUND_FLOOR = 1.0f`, `V_BOUND_EWMA_ALPHA = 0.001f`, `MIN_WIN`, `MIN_RATIO`, `MAX_RATIO` +12. `rl_gate_threshold_controller` — hardcoded `alpha = 0.01f` for EMA smoothing (one constant in an otherwise ISV-driven controller) + +**Genuinely out of scope** (audited, confirmed adaptive already, or different bug class): +- `rl_lr_controller` — all configs are ISV slots; uses only `1.0f`/`0.0f` math identities. Confirmed already adaptive. +- Phase 4.5 logic (`rl_advantage_normalize.cu`) — only adds an extra ISV write to emit `var_pre_norm`; the ε² floor itself is a math constant of the normalization, not a controller threshold. +- Controller objectives (what each controller is *trying* to achieve) — out of scope. +- Schulman bounded-step math — out of scope. + +This spec replaces hardcoded constants with observed-signal-driven values. **One atomic commit. No phase split.** + +--- + +## Diagnosis (verified, not theoretical) + +Per fold 0 diag.jsonl (`/feature-cache/alpha-rl-runs/669578566/fold0/`): + +| Controller | Bootstrap | Step 100 | Step 19999 | Mode | +|---|---|---|---|---| +| `γ` (gamma) | 0.995 | 0.995 | 0.995 | Stuck at MIN | +| `τ` (target_tau) | 0.005 | 0.005 | 0.005 | Never moved | +| `ε` (ppo_clip) | 0.20 | **0.50** | **0.50** | Stuck at MAX | +| entropy_coef | 0.5 | 0.01 | 0.01 | Stuck at MIN | +| per_α | 0.4 | 0.41 | 0.40 | Stuck at MIN | +| reward_scale | 1.0 | **0.0046** | **0.0042** | Crashed 250× in 100 steps | +| n_rollout_steps | 2048 | 2048 | 2048 | Never moved | +| `λ` (q_distill_lambda) | 0.01 | **0.05** | **0.05** | Stuck at MIN_LAMBDA | + +Trade pnl (eval phase, files 3-5 held out): **-$1.56M, win_rate 0.235, profit_factor 0.64**. Policy did not generalize. + +**Note**: `sac_alpha` trends downward over training (0.01 → 0.004) but is NOT saturated — appears to be SAC's natural alpha-annealing. Not addressed in this spec. + +### Mechanism for the canonical case (PPO clip) + +```c +// rl_ppo_clip_controller.cu +#define KL_NOISE_FLOOR_FRAC 0.01f +const float kl_noise_floor = kl_target * KL_NOISE_FLOOR_FRAC; // 0.0001 +if (kl_ema < kl_noise_floor) return; // hold + +if (kl_ema < kl_target / tolerance) { // tolerance=1.5 → triggers when kl < 0.0067 + ratio = adjust_rate; // 1.5× → widen ε +} +``` + +Trajectory (verified): + +| Step | ε | kl_pi | +|---|---|---| +| 1 | 0.20 (bootstrap) | 0.0 | +| 50 | **0.50 (MAX)** | 0.002 | +| 100 | 0.50 | 1.9e-11 | +| 500 | 0.50 | 1.8e-3 | +| 1000+ | 0.50 | 0 or numerical noise | + +At step 50 the controller saw `kl_ema = 0.002` — above the noise floor 0.0001 but below the band lower bound 0.0067. WIDEN path fired three times (1.5× per call). Bootstrap-replace-directly path on the first call wrote 0.30, Wiener blends on the next two pushed to 0.5 ceiling. Then `kl_ema` collapsed below the noise floor for the rest of training; ε never came back. + +### Why ksll2 didn't show this + +ksll2 had 4× the data (9 files vs 3). KL distribution was wider, and at some point `kl_ema` exceeded `kl_target * tolerance = 0.015`, triggering the TIGHTEN path that pulled ε back. With 3 files, KL never gets that high — controller never sees a reason to tighten. + +This pattern repeats across all 7 controllers, each with its own variant of the same bug. + +### Why this is one bug, not seven + +All 7 controllers share three architectural assumptions, each invalidated by Phase 4.5: + +1. **Hardcoded noise floor as fraction of target**: 4 controllers use `_NOISE_FLOOR_FRAC = 0.01f` (`ppo_clip`, `target_tau`, `rollout_steps`, `fused`). The 1% floor was calibrated against pre-normalization signal magnitudes; under Phase 4.5 it's ~100× too low. +2. **Fixed target ISV slots**: `kl_target=0.01`, `q_div_target=0.01`, `adv_var_ratio_target=5.0` etc. are bootstrapped once and never updated. They assume a signal scale that no longer holds. +3. **Aggressive bootstrap-to-update**: most controllers have a "bootstrap-replace-directly" path that lets a single observation push them off bootstrap. Combined with (1) + (2), one noisy step can drive the controller to an extremum. + +--- + +## Approaches considered + +### Approach A — Targeted patch (~12 LOC, lowest cost) + +Replace `_NOISE_FLOOR_FRAC = 0.01f` with `_NOISE_FLOOR_FRAC = 0.5f` (50% of target instead of 1%). Apply to the 4 controllers that have this constant. No new kernels, no ISV changes. + +**Pros**: One commit, one test, mechanical change. Validates the diagnosis quickly. + +**Cons**: Still a hardcoded constant — violates `feedback_adaptive_not_tuned`. Doesn't fix `reward_scale` crash (different mechanism). Doesn't fix `gamma`/`entropy_coef`/`per_α` (no `_NOISE_FLOOR_FRAC`, different saturation modes). Might break ksll2 regime where 1% worked (50% would hold the controller too aggressively when there *is* meaningful signal). Brittle to the next normalization change. + +**Verdict**: Useful as a same-day regression test for the diagnosis, but not the actual fix. + +### Approach B — Signal-driven floors (~200 LOC, the architectural fix) + +For each controller's input signal, track Welford-online variance in a dedicated ISV slot. Replace the hardcoded `floor = target × const` with `floor = MAX(target × 0.5, sqrt(observed_var) × 2)`. The floor naturally rescales to whatever signal distribution Phase 4.5 (or any future normalization) produces. + +**Pros**: Truly adaptive — `feedback_adaptive_not_tuned` compliant. Composes with future architectural changes (Phase 4.6 normalization, new heads, etc.) without re-tuning. Same kernel pattern works for all 7 controllers. Welford is well-understood, CUDA-friendly. The added Welford-variance ISV slots are also valuable diagnostics. + +**Cons**: ~200 LOC: new Welford-variance kernel (~80 LOC) + ISV slot allocation (8 new slots, one per controller input) + wiring (~80 LOC) + controller updates (~30 LOC across 7 files). Larger change, larger blast radius. + +**Verdict**: This is the architectural fix the foxhunt design philosophy promised. Implementation cost is real but bounded. + +### Approach C — Conservative bootstrap + asymmetric Schulman (~50 LOC) + +Don't touch the floors. Instead: (1) bootstrap each controller at its TARGET, not at a safety bound, (2) require N consecutive below-band observations before widening fires (so single noisy observations don't drive drift). Hold bootstrap until cumulative evidence accumulates. + +**Pros**: Smaller change than B. Preserves controller logic. Holds controllers at sensible defaults when signal is sparse. + +**Cons**: Still uses fixed thresholds — same brittleness as A. Doesn't address `reward_scale` crash. Doesn't address `n_rollout_steps` which is structurally broken (its input signal `advantage_var_ratio` is *defined* to be near zero under Phase 4.5). + +**Verdict**: Orthogonal to A/B. Worth folding INTO B as the "asymmetric Schulman" piece. + +### Recommendation: B + the Schulman asymmetry from C, applied to ALL 12 controllers with hardcoded thresholds + +The architectural fix is B. Fold C's asymmetric-Schulman idea into B's controller updates for Group A. Skip A as a separate change — it's a strict subset of B. + +`reward_scale`, `n_rollout_steps`, `gamma`, `q_distill_lambda` need separate treatment within this spec — they have *structural* problems (their signals are not just noisy, they're definitionally affected by Phase 4.5 or use ramp/decay rather than Schulman). Spec covers them in dedicated Special case sections (G, Q, R). + +Group B controllers (4 of them) follow the same Welford-variance pattern from B but without the asymmetric Schulman — they're not Schulman-based controllers, just have hardcoded floors/ceilings. Same kernel infrastructure, same ISV slot pattern. + +**One atomic commit. 12 controllers. 52 new ISV slots. One new Welford kernel.** + +--- + +## Detailed design (Approach B) + +### Core pattern (applies 7 places) + +```c +// 1. Welford variance kernel (new, runs per-step alongside EMA producers). +// Maintains running mean + M2 in dedicated ISV slots per signal. +// See: crates/ml-alpha/cuda/rl_signal_variance_update.cu (NEW) + +extern "C" __global__ void rl_signal_variance_update( + float* __restrict__ isv, + int input_slot, // e.g. RL_KL_PI_EMA_INDEX + int count_slot, // ISV[input_slot + COUNT_OFFSET] + int mean_slot, // ISV[input_slot + MEAN_OFFSET] + int m2_slot, // ISV[input_slot + M2_OFFSET] + int dones_present // 1 if this step contributed a real observation +) { + if (threadIdx.x != 0 || !dones_present) return; + const float x = isv[input_slot]; + if (x == 0.0f) return; // skip sentinel + float n = isv[count_slot] + 1.0f; + float delta = x - isv[mean_slot]; + float mean_new = isv[mean_slot] + delta / n; + float delta2 = x - mean_new; + float m2_new = isv[m2_slot] + delta * delta2; + isv[count_slot] = n; + isv[mean_slot] = mean_new; + isv[m2_slot] = m2_new; +} +``` + +### Controller integration (canonical example — PPO clip) + +```c +// rl_ppo_clip_controller.cu (MODIFIED) +// BEFORE: +const float kl_noise_floor = kl_target * KL_NOISE_FLOOR_FRAC; // 1% — broken + +// AFTER: +const float kl_var = (isv[KL_PI_COUNT_SLOT] > 1.0f) + ? isv[KL_PI_M2_SLOT] / (isv[KL_PI_COUNT_SLOT] - 1.0f) + : 0.0f; +const float kl_std = sqrtf(kl_var); +const float kl_noise_floor = fmaxf(kl_target * 0.5f, kl_std * 2.0f); +// ↑ floor never below 50% of target +// ↑ AND scales with observed signal std + +if (kl_ema < kl_noise_floor) return; +``` + +The `0.5f` constant is the only remaining tuned number, and it's safe across all signal regimes (50% of target = "less than half the desired operating distance from steady state"). + +### Asymmetric Schulman (folded in from Approach C) + +```c +// Tighten on single observation (real signal above target = act fast). +// Widen only on N consecutive observations below band (filter noise). + +if (kl_ema > kl_target * tolerance) { + ratio = 1.0f / adjust_rate; + isv[KL_BELOW_COUNT_SLOT] = 0.0f; // reset +} else if (kl_ema < kl_target / tolerance) { + isv[KL_BELOW_COUNT_SLOT] += 1.0f; + ratio = (isv[KL_BELOW_COUNT_SLOT] >= 3.0f) ? adjust_rate : 1.0f; +} else { + isv[KL_BELOW_COUNT_SLOT] = 0.0f; // reset on in-band + ratio = 1.0f; +} +``` + +The `3.0f` is the "N consecutive" parameter. Tunable — default 3 (~3 minibatch updates of confirmation). Could itself be ISV-driven (`SCHULMAN_WIDEN_PATIENCE_SLOT`) — TBD if needed. + +### Per-controller specifics + +ISV slots below allocated starting at 588 (current `RL_SLOTS_END`). Final `RL_SLOTS_END = 640` (52 new slots). + +**Group A — 8 saturated controllers (Welford variance + below-band counter for asymmetric Schulman):** + +| Controller | Input signal | Input slot | Variance triple (count, mean, M²) | Asym counter | Asymmetric Schulman? | +|---|---|---|---|---|---| +| ppo_clip | kl_pi_ema | 419 | 588, 589, 590 | 591 | yes | +| target_tau | q_divergence_ema | 420 | 592, 593, 594 | 595 | yes | +| rollout_steps | adv_var_pre_norm (NEW 612) | 421 → 612 | 596, 597, 598 | 599 | yes | +| entropy_coef | entropy_observed_ema | 422 | 600, 601, 602 | 603 | yes | +| per_α | td_kurtosis_ema | 423 | 604, 605, 606 | 607 | yes | +| gamma | trade_duration_ema | 417 | 608, 609, 610 | — | no (Special G) | +| reward_scale | pnl_per_batch_magnitude (NEW 614) | 614 | 615, 616, 617 | — | no (Special R) | +| q_distill_lambda | q_distill_kl_ema | 488 | 618, 619, 620 | — | no (Special Q) | + +**Group B — 4 not-yet-saturated controllers with hardcoded thresholds:** + +| Controller | Hardcoded → ISV slots | Notes | +|---|---|---| +| v_blend_alpha (Phase 4.4) | `DEAD_SIGNAL_FLOOR` → 621, `TARGET_TRACK_RATIO` → 622, `SCHULMAN_STEP` → 623, `EMA_ALPHA` → 624, `BOOTSTRAP_ALPHA` → 625 | 5 new slots; controller already takes one Welford variance for V_scalar magnitude — slots 626, 627, 628 | +| ppo_ratio_clamp | `MIN_OUT` (2.0f) → adaptive via ratio_variance; `MAX_OUT` (1000.0f) → 629 | Welford on ratio magnitude — slots 630, 631, 632 | +| reward_clamp | `V_BOUND_FLOOR` → 633, `V_BOUND_EWMA_ALPHA` → 634, `MIN_WIN` → 635, `MIN_RATIO` → 636, `MAX_RATIO` → 637 | 5 slots; already partially signal-driven, just removing remaining hardcoded constants | +| gate_threshold | hardcoded `alpha = 0.01f` → `RL_GATE_EMA_ALPHA_INDEX` 638 | Single slot | + +**New input signals (separate from variance slots):** +- 612: `RL_ADV_VAR_PRE_NORM_INDEX` (emitted by `rl_advantage_normalize.cu` before in-place normalize) +- 613: `RL_GAMMA_MIN_ADAPTIVE_INDEX` (computed from trade_duration_ema) +- 614: `RL_REWARD_MAGNITUDE_EMA_INDEX` (per-batch pnl magnitude for reward_scale floor) +- 639: `RL_Q_DISTILL_LAMBDA_MIN_ADAPTIVE_INDEX` (computed from q_distill_kl_ema variance) + +**Total: 52 new slots. `RL_SLOTS_END` grows 588 → 640.** Memory impact: +208 bytes. + +### Special cases (not covered by core pattern) + +#### `n_rollout_steps` — structural conflict with Phase 4.5 + +Its input signal `advantage_var_ratio_ema` is *definitionally* near zero under Phase 4.5 normalization (advantage var ~= ε² floor). The Welford-variance fix won't help because the signal itself is structurally diluted. + +**Proposed fix**: track *pre-normalization* advantage variance in a separate ISV slot. Phase 4.5's `rl_advantage_normalize.cu` computes the mean and variance internally — emit `var_pre_norm` to a new ISV slot (`RL_ADV_VAR_PRE_NORM_INDEX`). `rl_rollout_steps_controller` reads from that instead. + +~20 LOC change to `rl_advantage_normalize.cu` to emit the ISV value before in-place normalize. + +#### `reward_scale` — crash, not saturation + +Crashed from 1.0 → 0.004 in 100 steps. Mechanism: when `pnl_per_batch_magnitude` is small (which it is on small datasets with the agent not yet trading much), the scale controller aggressively shrinks `reward_scale` to compensate. By step 100 the reward signal is muted to 0.4% of bootstrap, which then makes everything downstream tiny → controllers see no signal → death spiral. + +**Proposed fix**: rate-limit the `reward_scale` DECREASE more aggressively than INCREASE (asymmetric like Schulman above). Also: floor `reward_scale` at a fraction of bootstrap (e.g., `0.1 × REWARD_SCALE_BOOTSTRAP = 0.1`) until N steps have passed with meaningful trades. This protects against the early-training spiral. + +~40 LOC in `rl_reward_scale_controller.cu`. + +#### Special case G — `gamma` deadlock from hardcoded `GAMMA_MIN` + +Per the trajectory, `gamma` was at MIN (0.995) at step 1 (bootstrap from sentinel) and never moved. Mechanism: + +```c +// rl_gamma_controller.cu +#define GAMMA_MIN 0.995f // hardcoded — violates feedback_adaptive_not_tuned +#define GAMMA_MAX 0.999f +// target(d) = 1 - 1/d (standard discount-from-horizon mapping) +// At d_ema = 15.6: target = 1 - 1/15.6 = 0.936 +// clamp(0.936, 0.995, 0.999) = 0.995 ← always pinned at MIN +``` + +`target(d)` only exceeds `GAMMA_MIN = 0.995` when `d_ema > 200` — i.e., trade durations over 200 minibatch ticks (~10s+ in the surfer regime). Below that, gamma is structurally pinned at 0.995. With fold 0's `d_ema = 15.6`, target = 0.936, controller wants to lower gamma but can't (clamp), so it sticks. + +**This is the same architectural bug** as the noise floors: a *hardcoded* threshold that may have been right for one regime but is wrong for another. With small data, agent makes shorter trades; the gamma floor traps the agent in a discount factor that's too long for its trade horizon, which feeds back into the value targets, which biases learning. + +**Deadlock**: the agent needs higher gamma to learn longer trades, but the controller needs longer trades to allow higher gamma. With the hardcoded MIN at 0.995, the controller has no room to lower gamma to match short trades, and no signal to raise it because trades stay short. + +**Proposed fix** (consistent with the rest of the spec): + +Replace hardcoded `GAMMA_MIN = 0.995f` with adaptive `RL_GAMMA_MIN_ADAPTIVE_INDEX` slot, derived from observed `trade_duration_ema`: + +```c +// Derived adaptive floor: allow gamma to drop low enough for current +// regime's trade horizon, with a safety floor of 0.9 to prevent the +// agent from going to bandit mode. +const float d_ema = isv[RL_MEAN_TRADE_DURATION_EMA_INDEX]; +const float horizon_multiplier = 2.0f; // gamma allows 2× the trade horizon as effective lookahead +const float adaptive_min = fmaxf(0.9f, 1.0f - 1.0f / fmaxf(d_ema * horizon_multiplier, 10.0f)); +isv[RL_GAMMA_MIN_ADAPTIVE_INDEX] = adaptive_min; +``` + +With `d_ema=15.6`, adaptive_min = `1 - 1/31.2 = 0.968`. Now `target(15.6) = 0.936` still clamps, but to `0.968` instead of `0.995` — and the controller has room to track real trade durations. As the agent learns longer trades and `d_ema` grows, `adaptive_min` rises naturally. + +The `0.9` absolute floor and `2.0` horizon multiplier are the only tuned constants — both architecturally motivated: +- `0.9` floor: per `pearl_edge_lives_at_wave_timescale_not_tick`, edge lives at wave timescale (γ≥0.99). The 0.9 floor is a hard "don't degenerate to bandit" guard, not a typical operating point. +- `2.0` horizon multiplier: standard rule-of-thumb that effective horizon should be 2× the natural transaction horizon to capture exit dynamics. + +Both could be ISV-driven in a follow-up spec, but the current values are conservative defaults that don't bias the controller's adaptation. + +#### Special case Q — `q_distill_lambda` ramp/decay with hardcoded MIN + +Per the trajectory: `λ` bootstraps at 0.01, decays via `LAMBDA_DECAY_RATE = 0.998f` for ~100 steps because `q_distill_kl_ema ≈ 0.001` is far below `q_distill_kl_target = 0.10`. Reaches `MIN_LAMBDA = 0.05f` floor and sticks for the next 19900 steps. + +```c +// rl_q_distill_lambda_controller.cu +#define MIN_LAMBDA 0.05f // hardcoded floor +#define MAX_LAMBDA 1.0f // hardcoded ceiling +// When kl_ema < kl_target / tolerance → λ *= LAMBDA_DECAY_RATE = 0.998 +// When kl_ema > kl_target * tolerance → λ *= LAMBDA_RAMP_RATE = 1.2 +// Clamp to [MIN_LAMBDA, MAX_LAMBDA] +``` + +Same bug class as the noise floors and gamma: a hardcoded floor (`MIN_LAMBDA`) that doesn't adapt to the actual signal regime. Phase 4.5 reduces KL signal magnitudes (which this controller also consumes via `q_distill_kl_ema`), so the "below target" path fires continuously, decaying λ to MIN. + +**Proposed fix** (mirroring the gamma fix pattern): make `MIN_LAMBDA` and `MAX_LAMBDA` ISV-driven adaptive bounds derived from observed `q_distill_kl_ema` distribution: + +```c +// adaptive_lambda_min = MAX(0.001, kl_ema_std * SCALE_FACTOR) +// where SCALE_FACTOR ≈ 0.05 (provides headroom proportional to signal noise) +const float kl_std = sqrtf(isv[Q_DISTILL_KL_VAR_SLOT]); +const float adaptive_min = fmaxf(0.001f, kl_std * 0.05f); +isv[RL_Q_DISTILL_LAMBDA_MIN_ADAPTIVE_INDEX] = adaptive_min; +``` + +This reuses the same Welford-variance kernel infrastructure proposed in the core pattern (one variance triple for `q_distill_kl_ema`). Lambda can now decay to a level proportional to the actual signal noise rather than a hardcoded 0.05 floor. + +The MAX is less critical (1.0 ceiling is rarely hit in practice) but consistency: derive MAX as `MIN(1.0, MAX(0.5, kl_std * 50.0))`. Provides upward headroom proportional to signal range. + +Adds 1 ISV slot (`RL_Q_DISTILL_LAMBDA_MIN_ADAPTIVE_INDEX`) — total grows from 27 → 28. + +--- + +## Files + +### New + +- `crates/ml-alpha/cuda/rl_signal_variance_update.cu` — Welford variance kernel (~80 LOC) + +### Modified + +- `crates/ml-alpha/src/rl/isv_slots.rs` — add 27 new slots (3 variance triples × 7 controllers + 3 input slots + 3 asymmetric Schulman counters); grow `RL_SLOTS_END` 588 → 615 +- `crates/ml-alpha/cuda/rl_ppo_clip_controller.cu` — adaptive noise floor + asymmetric Schulman +- `crates/ml-alpha/cuda/rl_target_tau_controller.cu` — adaptive noise floor + asymmetric Schulman +- `crates/ml-alpha/cuda/rl_rollout_steps_controller.cu` — adaptive noise floor + asymmetric Schulman + consume `RL_ADV_VAR_PRE_NORM_INDEX` (new) instead of post-norm advantage_var_ratio +- `crates/ml-alpha/cuda/rl_entropy_coef_controller.cu` — adaptive noise floor + asymmetric Schulman +- `crates/ml-alpha/cuda/rl_per_alpha_controller.cu` — adaptive noise floor + asymmetric Schulman +- `crates/ml-alpha/cuda/rl_reward_scale_controller.cu` — asymmetric rate-limit (decrease ≤ 1.05× per step) + bootstrap-fraction floor (`reward_scale ≥ 0.1 × bootstrap` until N trades accumulated) + emit `RL_REWARD_MAGNITUDE_EMA_INDEX` +- `crates/ml-alpha/cuda/rl_gamma_controller.cu` — replace hardcoded `GAMMA_MIN = 0.995f` with adaptive `RL_GAMMA_MIN_ADAPTIVE_INDEX` derived from `trade_duration_ema` (Special case G) +- `crates/ml-alpha/cuda/rl_q_distill_lambda_controller.cu` — replace hardcoded `MIN_LAMBDA = 0.05f` and `MAX_LAMBDA = 1.0f` with adaptive ISV-driven bounds (Special case Q, below) +- `crates/ml-alpha/cuda/rl_v_blend_alpha_controller.cu` — replace 5 hardcoded constants with ISV slots (`DEAD_SIGNAL_FLOOR` → signal-driven, `TARGET_TRACK_RATIO`/`SCHULMAN_STEP`/`EMA_ALPHA`/`BOOTSTRAP_ALPHA` → ISV slots) +- `crates/ml-alpha/cuda/rl_ppo_ratio_clamp_controller.cu` — replace `PPO_RATIO_CLAMP_MIN_OUT = 2.0f` with adaptive floor derived from observed ratio variance; `MAX_OUT = 1000.0f` becomes ISV-driven ceiling +- `crates/ml-alpha/cuda/rl_reward_clamp_controller.cu` — replace `V_BOUND_FLOOR`, `V_BOUND_EWMA_ALPHA`, `MIN_WIN`, `MIN_RATIO`, `MAX_RATIO` with adaptive ISV slots (`V_BOUND_EWMA_ALPHA` becomes signal-variance-driven; ratio bounds derive from observed win/loss distribution) +- `crates/ml-alpha/cuda/rl_gate_threshold_controller.cu` — replace hardcoded `alpha = 0.01f` with ISV slot `RL_GATE_EMA_ALPHA_INDEX` +- `crates/ml-alpha/cuda/rl_advantage_normalize.cu` — emit raw `var_pre_norm` to `RL_ADV_VAR_PRE_NORM_INDEX` before in-place normalize +- `crates/ml-alpha/cuda/rl_fused_controllers.cu` — all of the above unified (this is what the trainer actually launches per-step) +- `crates/ml-alpha/src/trainer/integrated.rs` — launch `rl_signal_variance_update` per step + wire new ISV slots + populate `RL_GAMMA_MIN_ADAPTIVE_INDEX` via fused kernel + +### Not touched + +- Phase 4.5 logic itself (`rl_advantage_normalize.cu` body) — the only addition is one ISV write to emit `var_pre_norm` for `rl_rollout_steps_controller` consumption. The ε² floor in the normalization math stays — it's a numerical-stability constant of the kernel, not a controller threshold. +- `rl_lr_controller` — **AUDITED 2026-05-30**: all configs are ISV slots (`RL_LR_BOOTSTRAP_INDEX`, `RL_LR_MIN_INDEX`, `RL_LR_MAX_INDEX`, `RL_LR_DECAY_FACTOR_INDEX`, `RL_LR_LOSS_EMA_ALPHA_INDEX`, `RL_IMPROVEMENT_THRESHOLD_INDEX`, `RL_PLATEAU_PATIENCE_INDEX`, `RL_LR_WARMUP_STEPS_INDEX`). Kernel body uses only `1.0f`/`0.0f` math identities (not thresholds). Already complies with `feedback_adaptive_not_tuned`. Per `pearl_plateau_decay_needs_warmup_for_sparse_heads`, the plateau-decay pattern is intentionally different from Schulman; keeping it. +- All other RL kernels not listed in Modified (e.g., `rl_fused_targets.cu`, kernel utilities) — not controllers, not subject to this spec. + +--- + +## Validation gates + +### Unit / smoke gates (local, RTX 3050 Ti) + +1. **G1 — Welford convergence**: feed constant signal `k=0.5`, 50 steps, assert `isv[mean_slot] ≈ k` AND `isv[m2_slot] / (n-1) ≈ 0` (no variance). +2. **G2 — Welford variance**: feed signal `[1, 2, 3, 4, 5]`, assert `isv[m2_slot] / 4 ≈ 2.5` (sample variance). +3. **G3 — PPO clip holds on small signal**: bootstrap controller, feed `kl_ema = 0.001` for 100 steps with observed_var=1e-7, assert `ε` does NOT drift past `target × tolerance = 0.015`. +4. **G4 — PPO clip tightens on big signal**: feed `kl_ema = 0.05` once (single overshoot), assert `ε` decreases by `1/adjust_rate` immediately (asymmetric tightening). +5. **G5 — PPO clip widens slowly**: feed `kl_ema = 0.001` for 10 steps, assert widen fires ONLY after 3 consecutive below-band observations (asymmetric). +6. **G6 — integrated_trainer_smoke**: existing GPU smoke must still pass. + +### Cluster gates + +7. **G7 — fold 0 re-run with fix**: same params (b=1024, 20k, n_folds=3, fold_idx=0). At step 19999: + - PPO clip `ε ≤ 0.20` (held in bootstrap region, was 0.50) + - `reward_scale ≥ 0.5` (bootstrap-fraction floor active, was 0.004) + - `n_rollout_steps ≠ 2048` (moved off bootstrap, signals real adaptation via pre-norm advantage variance) + - `gamma > 0.97` (adaptive_min derived from `d_ema ≈ 15` allows movement above hardcoded 0.995) + - `q_distill_lambda > 0.05` (above old MIN floor, adaptive bound from signal variance) + - 5+ of the 8 Group A controllers no longer pinned at extrema +8. **G8 — fold 0 eval pnl improvement**: eval pnl > -$0.5M (was -$1.56M). Profit factor > 0.85 (was 0.64). Win rate > 0.40 (was 0.235). These are improvement gates, not absolute targets — we expect *some* improvement, not full positive pnl on 3-file training. +9. **G9 — ksll2 regression test**: same params as `alpha-rl-ksll2` (n_folds=1, b=1024, 20k). Final pnl ≥ +$15M (within 20% of original +$18.3M). Confirms fix doesn't regress the large-data regime where the prior hardcoded floors happened to work. **All 12 controllers** should reach values consistent with the +$18M ksll2 trajectory (e.g., ε around target 0.01, not at 0.50; reward_scale around 0.5-1.0; gamma in [0.97, 0.999]). +10. **G10 — Group B non-regression**: the 4 Group B controllers (`v_blend_alpha`, `ppo_ratio_clamp`, `reward_clamp`, `gate_threshold`) didn't visibly saturate before this fix — verify they still operate in their prior healthy ranges: + - `v_blend_alpha` settles in [0.0, 1.0] with non-trivial variation (was bootstrapping to 1.0 then dropping toward 0) + - `ppo_ratio_clamp_max` settles at a meaningful value (was 10.0 bootstrap → adaptive) + - `reward_clamp_win` / `reward_clamp_loss` track observed reward distribution (no drift toward MIN/MAX of new adaptive bounds) + - `gate_threshold` EMA tracks dones_target as before (the only change is `alpha = 0.01f` → ISV-driven, semantically equivalent at the new default) + +--- + +## Open questions resolved + +1. ~~**ISV slot count budget**~~ → **resolved**: `RL_SLOTS_END = 588` currently, highest used 587. No upper cap. 52 new slots → `RL_SLOTS_END = 640`. Memory impact +208 bytes (negligible). +2. **Welford reset policy**: variance EMA carries over the train→eval boundary. *Decision*: accept this — eval-phase ISV reads should still see the train-phase variance estimate (controllers don't update during eval anyway). +3. **N-consecutive-below counter scope**: spec uses per-controller counters (5 ISV slots, one per Group A controller that needs asymmetric Schulman). Group B controllers don't need them (different mechanism). +4. **`reward_scale` bootstrap-fraction floor**: 0.1 (10% of bootstrap) — accepted as a tuned constant guarded by the rationale that early-training reward distributions are inherently noisy and shrinking past 10% of bootstrap is almost never beneficial. Could be ISV-driven in a follow-up. +5. **Order of implementation**: **all-in-one commit per `feedback_no_partial_refactor`** (Option B selected). Partial migration would create ambiguous validation signal — "did ppo_clip fix work but other controllers ruined it?" vs "did the architectural fix work?" — and the contract change (new ISV slots, modified controller semantics) needs atomic application. The plan decomposes into ordered TASKS that flow into ONE commit at the end, not one commit per controller. No "Phase 1/Phase 2" — 12 controllers in one commit. + +6. ~~**Gamma adaptive_min coupling**~~ → **resolved** (2026-05-30): use the Welford mean of trade duration directly (already computed for `rl_per_alpha_controller`'s noise floor), with a 100-observation warmup gate. The Welford mean's natural lag (denominator-of-N smoothing) breaks the feedback loop without any explicit step-period gating or exponential cool-down. Implementation: while `count < 100` use the existing EMA; once `count >= 100` switch to the Welford mean as the input to `target(d)`. No new ISV slots, no new constants beyond the 100-observation gate (which is a confidence threshold, not a magnitude calibration). See plan Task 7 for the code. + +--- + +## Risks & mitigations + +- **Risk**: Welford with float32 + small variance values → catastrophic cancellation. **Mitigation**: standard Welford formulation already handles this; CUDA float32 is sufficient at our signal magnitudes (1e-5 to 1.0). +- **Risk**: New ISV slots break existing tests that snapshot the full ISV layout. **Mitigation**: tests need to be updated alongside; this was just done in the test-migration commit `a3dfcd63f` for similar reasons. +- **Risk**: Asymmetric Schulman makes controllers TOO conservative, blocking real adaptation. **Mitigation**: the N-consecutive parameter is small (3), tightening still fires on single observation (fast safety response), and the in-band hold preserves the steady state. +- **Risk**: The fix works on fold 0 but ksll2 (large data) actually NEEDS the loose noise floor to adapt fast enough. **Mitigation**: G9 cluster regression catches this. Fallback: ISV-driven `NOISE_FLOOR_FRAC` slot, dynamically tunable per-run. + +--- + +## Next step after approval + +Write the plan: `docs/superpowers/plans/2026-05-30-adaptive-controller-floor-plan.md` with task-level decomposition (foundation → ppo_clip → sweep → validation → cluster runs).