diff --git a/crates/ml-alpha/build.rs b/crates/ml-alpha/build.rs index 04a40d4af..060030251 100644 --- a/crates/ml-alpha/build.rs +++ b/crates/ml-alpha/build.rs @@ -62,6 +62,8 @@ const KERNELS: &[&str] = &[ "rl_step_counter_update", // per-batch trade-duration counter + done-gated emit → mean_trade_duration_ema (ISV[417]) feeding rl_gamma "rl_l2_norm", // ‖x‖₂ single-buffer reduction → q/pi/v grad-norm EMAs (ISV[424..427]) feeding rl_lr_controller // (entropy_observed_ema, ISV[420], feeds rl_entropy_coef directly via ema_update_per_step's internal mean reduce on entropy_d — no separate kernel needed.) + "rl_ppo_ratio_clamp_controller", // RL R9: PPO ratio clamp ceiling at ISV[440], anchored on ε at ISV[402] — bounds catastrophic unclipped-branch surrogate + "ppo_log_ratio_abs_max_b", // RL R9 diag: per-batch max|log π_new − log π_old| → ISV[441]; surfaces ratio-clamp activity in diag JSONL ]; // Cache bust v31 — five new reduce / derive kernels populate the input diff --git a/crates/ml-alpha/cuda/ppo_clipped_surrogate.cu b/crates/ml-alpha/cuda/ppo_clipped_surrogate.cu index 6b01c9c69..e0322d386 100644 --- a/crates/ml-alpha/cuda/ppo_clipped_surrogate.cu +++ b/crates/ml-alpha/cuda/ppo_clipped_surrogate.cu @@ -58,10 +58,14 @@ // action, so no parallelism is wasted there). // Backward: same layout. Each thread writes its own logit's gradient. -#define HIDDEN_DIM 128 -#define N_ACTIONS 9 -#define RL_PPO_CLIP_INDEX 402 -#define RL_ENTROPY_COEF_INDEX 403 +#define HIDDEN_DIM 128 +#define N_ACTIONS 9 +#define RL_PPO_CLIP_INDEX 402 +#define RL_ENTROPY_COEF_INDEX 403 +// ISV slot carrying the adaptive importance-ratio clamp ceiling +// produced by rl_ppo_ratio_clamp_controller. Floor is reciprocal +// (clamp window is symmetric in log space: [1/ratio_max, ratio_max]). +#define RL_PPO_RATIO_CLAMP_MAX_INDEX 440 // ───────────────────────────────────────────────────────────────────── @@ -200,12 +204,42 @@ extern "C" __global__ void ppo_clipped_surrogate_fwd( const float log_p = logf(fmaxf(s_softmax[a_taken] / s_sumexp, 1e-7f)); pi_log_prob[batch] = log_p; - const float ratio = expf(log_p - log_pi_old[batch]); - const float A = advantages[batch]; - const float eps = isv[RL_PPO_CLIP_INDEX]; - const float surr1 = ratio * A; - const float surr2 = fminf(fmaxf(ratio, 1.0f - eps), 1.0f + eps) * A; - const float l_pi = -fminf(surr1, surr2); + // Importance ratio with ISV-driven magnitude clamp. + // + // PPO's clip(r, 1-ε, 1+ε) bounds the LOSS when surr2 is the + // active min — i.e. A>0,r>1+ε and A<0,r<1-ε. The + // unclipped branch IS active when A<0,r>1+ε (surr1=A·r is + // more negative than surr2=A·(1+ε) so min=surr1) and when + // A>0,r<1-ε. In the first case `r` can blow up — we've + // seen r reach 1e10 within a few steps of policy drift, + // producing l_pi=-A·r=O(1e10) spikes that contaminate + // the loss-balance controller and the LR controller's + // plateau detection. + // + // Per `feedback_isv_for_adaptive_bounds` and + // `pearl_controller_anchors_isv_driven`: the clamp bound + // lives in ISV[RL_PPO_RATIO_CLAMP_MAX_INDEX = 440], emitted + // by `rl_ppo_ratio_clamp_controller` which derives it from + // the (already KL-adaptive) PPO clip ε at ISV[402]. The + // controller widens the clamp when ε widens (rare but + // larger excursions are expected) and tightens when ε + // tightens (small excursions should be rare anomalies). + // + // Backward (ppo_clipped_surrogate_bwd) gates pg_grad + // inside [1-ε, 1+ε] anyway so this clamp is forward-only — + // gradients are bounded already; this bounds the reported + // loss value to keep downstream controllers sane. + const float ratio_raw = expf(log_p - log_pi_old[batch]); + const float ratio_max = isv[RL_PPO_RATIO_CLAMP_MAX_INDEX]; + const float ratio_min = 1.0f / ratio_max; + const float ratio = fmaxf(ratio_min, + fminf(ratio_raw, ratio_max)); + const float A = advantages[batch]; + const float eps = isv[RL_PPO_CLIP_INDEX]; + const float surr1 = ratio * A; + const float surr2 = fminf(fmaxf(ratio, 1.0f - eps), + 1.0f + eps) * A; + const float l_pi = -fminf(surr1, surr2); const float coef = isv[RL_ENTROPY_COEF_INDEX]; const float l_ent = -coef * h; @@ -300,7 +334,15 @@ extern "C" __global__ void ppo_clipped_surrogate_bwd( } const float log_p = logf(fmaxf(s_softmax[a_taken] / s_sumexp, 1e-7f)); - const float ratio = expf(log_p - log_pi_old[batch]); + // Match the forward's ratio clamp from ISV[RL_PPO_RATIO_CLAMP_MAX_INDEX]. + // For inside_clip the clamp is a no-op (since the controller floors + // at well above 1+ε); for outside_clip the gradient is zero anyway + // — but keeping fwd/bwd algebraically consistent prevents + // future-edit drift. + const float ratio_raw = expf(log_p - log_pi_old[batch]); + const float ratio_max = isv[RL_PPO_RATIO_CLAMP_MAX_INDEX]; + const float ratio_min = 1.0f / ratio_max; + const float ratio = fmaxf(ratio_min, fminf(ratio_raw, ratio_max)); const float A = advantages[batch]; const float eps = isv[RL_PPO_CLIP_INDEX]; diff --git a/crates/ml-alpha/cuda/ppo_log_ratio_abs_max_b.cu b/crates/ml-alpha/cuda/ppo_log_ratio_abs_max_b.cu new file mode 100644 index 000000000..d10af60c2 --- /dev/null +++ b/crates/ml-alpha/cuda/ppo_log_ratio_abs_max_b.cu @@ -0,0 +1,58 @@ +// ppo_log_ratio_abs_max_b.cu — per-batch max |log π_new − log π_old| +// reduced to a single scalar written to +// ISV[RL_PPO_LOG_RATIO_ABS_MAX_INDEX = 441]. +// +// Diagnostic-only kernel. Tells the diag JSONL what the largest raw +// log-ratio in this batch was BEFORE the kernel-side clamp in +// ppo_clipped_surrogate.cu applied. Useful for: +// +// * Confirming the ratio clamp is firing only when log_ratio_abs_max +// exceeds log(isv[RL_PPO_RATIO_CLAMP_MAX_INDEX]) — same comparison +// as scaled_pre_clamp_max vs reward clamp bound. +// * Surfacing the actual size of policy excursions step-to-step so +// we can spot whether the rl_ppo_clip_controller is keeping KL +// reasonable. +// * Cross-checking the rl_kl_approx_b mean-KL estimator against the +// PER-BATCH MAX |log ratio| — KL EMA can stay small while one +// outlier batch entry has a huge log ratio (the outlier the clamp +// catches). +// +// Same tree-reduce shape as rl_kl_approx_b and rl_l2_norm: single +// block, grid-stride loop into per-thread local max, then in-block +// reduction. No atomics per `pearl_no_atomicadd`. + +#define RL_PPO_LOG_RATIO_ABS_MAX_INDEX 441 + +extern "C" __global__ void ppo_log_ratio_abs_max_b( + const float* __restrict__ log_pi_old, // [b_size] + const float* __restrict__ log_pi_new, // [b_size] + float* __restrict__ isv, // [≥ 442] + int b_size +) { + extern __shared__ float s[]; + const int tid = threadIdx.x; + + // Per-thread: scan strided over b_size, accumulate local max |diff|. + float local_max = 0.0f; + for (int b = tid; b < b_size; b += blockDim.x) { + const float diff = log_pi_new[b] - log_pi_old[b]; + const float a = fabsf(diff); + if (a > local_max) local_max = a; + } + s[tid] = local_max; + __syncthreads(); + + // Tree reduce max — block dim must be power of 2; caller enforces. + for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) { + if (tid < stride) { + const float a = s[tid]; + const float b = s[tid + stride]; + s[tid] = a > b ? a : b; + } + __syncthreads(); + } + + if (tid == 0) { + isv[RL_PPO_LOG_RATIO_ABS_MAX_INDEX] = s[0]; + } +} diff --git a/crates/ml-alpha/cuda/rl_ppo_ratio_clamp_controller.cu b/crates/ml-alpha/cuda/rl_ppo_ratio_clamp_controller.cu new file mode 100644 index 000000000..0e3727ae8 --- /dev/null +++ b/crates/ml-alpha/cuda/rl_ppo_ratio_clamp_controller.cu @@ -0,0 +1,116 @@ +// rl_ppo_ratio_clamp_controller.cu — emits adaptive PPO importance-ratio +// clamp ceiling to ISV[RL_PPO_RATIO_CLAMP_MAX_INDEX = 440]. +// +// Per `feedback_isv_for_adaptive_bounds` + `pearl_controller_anchors_isv_driven`: +// the importance-ratio clamp bound used by `ppo_clipped_surrogate.cu` +// can't be a hardcoded #define — it has to track the controller chain. +// The natural anchor is the PPO clip ε at ISV[RL_PPO_CLIP_INDEX = 402] +// which is itself KL-driven by rl_ppo_clip_controller (small KL EMA → +// small ε → tight clip band; large KL EMA → wide ε → wide clip band). +// +// Mapping: target = (1 + ε) × PPO_CLAMP_MARGIN. +// +// * ε = 0.05 → target = 1.05 × 10 = 10.5 +// * ε = 0.20 → target = 1.20 × 10 = 12.0 +// * ε = 0.50 → target = 1.50 × 10 = 15.0 +// * ε = 1.00 → target = 2.00 × 10 = 20.0 +// +// The clamp window is symmetric in log space: [1/target, target]. With +// PPO_CLAMP_MARGIN = 10 the clamp sits ~10× the controller's stated +// clip half-width — wide enough that the clamp is a no-op for +// typical updates (which sit inside [1-ε, 1+ε]) but tight enough that +// catastrophic excursions (r = 1e10 from policy drift over a multi-step +// rollout) are bounded into l_pi = O(|A| · 10) instead of O(|A| · 1e10). +// +// Per `pearl_one_unbounded_signal_per_reward`: PPO surrogate has the +// importance ratio as its single unbounded multiplicand. PPO clip(r) +// bounds the loss when its branch is the active min; this clamp bounds +// the OTHER branch (A<0,r>1+ε and A>0,r<1-ε where surr1=A·r is the +// active min). +// +// Per `pearl_first_observation_bootstrap`: ISV slot starts at 0.0 +// sentinel. First emit writes PPO_RATIO_CLAMP_BOOTSTRAP = 10.0 directly +// (same magnitude as the would-have-been hardcoded constant; the +// controller then refines from this baseline as ε adapts). +// +// Per `pearl_wiener_alpha_floor_for_nonstationary`: ε itself is +// non-stationary (KL EMA drifts as the policy co-adapts), so the +// Wiener α is floored at 0.4 to keep the clamp responsive enough to +// follow ε changes without oscillation. +// +// Per `pearl_blend_formulas_must_have_permanent_floor`: output clamped +// to [PPO_RATIO_CLAMP_MIN_OUT, PPO_RATIO_CLAMP_MAX_OUT] — a permanent +// floor at 2.0 ensures the clamp NEVER collapses to the PPO clip band +// itself (which would gate even routine policy updates), and a +// permanent ceiling at 1000 prevents runaway when ε is large. + +#define RL_PPO_CLIP_INDEX 402 +#define RL_PPO_RATIO_CLAMP_MAX_INDEX 440 + +#define PPO_RATIO_CLAMP_BOOTSTRAP 10.0f +#define PPO_CLAMP_MARGIN 10.0f +#define PPO_RATIO_CLAMP_MIN_OUT 2.0f +#define PPO_RATIO_CLAMP_MAX_OUT 1000.0f +#define WIENER_ALPHA_FLOOR 0.4f + +// ───────────────────────────────────────────────────────────────────── +// rl_ppo_ratio_clamp_controller: +// Single-thread controller — writes ONE float to +// isv[RL_PPO_RATIO_CLAMP_MAX_INDEX]. +// +// Signature matches the other R5 controllers (isv, alpha, input_slot) +// so the trainer can launch via `launch_isv_controller_3arg`. +// +// Inputs: +// isv [≥ 442] — ISV bus. +// alpha_step — Wiener α from the controller's signal stats; +// floored at WIENER_ALPHA_FLOOR before blend. +// input_slot — ISV index of the ε signal (caller passes +// RL_PPO_CLIP_INDEX = 402). +// +// Outputs: +// isv[RL_PPO_RATIO_CLAMP_MAX_INDEX] — clamp ceiling ∈ +// [PPO_RATIO_CLAMP_MIN_OUT, +// PPO_RATIO_CLAMP_MAX_OUT] +// ───────────────────────────────────────────────────────────────────── +extern "C" __global__ void rl_ppo_ratio_clamp_controller( + float* __restrict__ isv, + float alpha_step, + int input_slot +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + const float prev = isv[RL_PPO_RATIO_CLAMP_MAX_INDEX]; + + // Bootstrap on sentinel 0.0 per pearl_first_observation_bootstrap. + if (prev == 0.0f) { + isv[RL_PPO_RATIO_CLAMP_MAX_INDEX] = PPO_RATIO_CLAMP_BOOTSTRAP; + return; + } + + // ── Derive target from current PPO clip ε at input_slot. ──────── + const float eps = isv[input_slot]; + float target = (1.0f + eps) * PPO_CLAMP_MARGIN; + + // 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)); + + // First-observation replace-directly per pearl_first_observation_bootstrap. + // prev == PPO_RATIO_CLAMP_BOOTSTRAP means "we just bootstrapped last + // step and this is the first real ε observation"; blending with the + // bootstrap would leave 60% of the 10.0 constant in the output even + // when ε says target should be much smaller (or larger). + if (prev == PPO_RATIO_CLAMP_BOOTSTRAP) { + isv[RL_PPO_RATIO_CLAMP_MAX_INDEX] = target; + return; + } + + // Wiener-α blend with floor. + const float a = fmaxf(alpha_step, WIENER_ALPHA_FLOOR); + float out = (1.0f - a) * prev + a * target; + + out = fmaxf(PPO_RATIO_CLAMP_MIN_OUT, + fminf(out, PPO_RATIO_CLAMP_MAX_OUT)); + isv[RL_PPO_RATIO_CLAMP_MAX_INDEX] = out; +} diff --git a/crates/ml-alpha/examples/alpha_rl_train.rs b/crates/ml-alpha/examples/alpha_rl_train.rs index 85303d6d9..271cf9307 100644 --- a/crates/ml-alpha/examples/alpha_rl_train.rs +++ b/crates/ml-alpha/examples/alpha_rl_train.rs @@ -47,9 +47,9 @@ use ml_alpha::rl::isv_slots::{ RL_LR_V_STEPS_SINCE_BEST_INDEX, RL_LR_V_WARMUP_COUNTER_INDEX, RL_MAX_ABS_SCALED_REWARD_PRE_CLAMP_INDEX, RL_MEAN_ABS_PNL_EMA_INDEX, RL_MEAN_TRADE_DURATION_EMA_INDEX, RL_N_ROLLOUT_STEPS_INDEX, RL_PER_ALPHA_INDEX, - RL_PI_GRAD_NORM_EMA_INDEX, RL_PPO_CLIP_INDEX, RL_Q_DIVERGENCE_EMA_INDEX, - RL_Q_GRAD_NORM_EMA_INDEX, RL_REWARD_SCALE_INDEX, RL_TARGET_TAU_INDEX, RL_TD_KURTOSIS_EMA_INDEX, - RL_V_GRAD_NORM_EMA_INDEX, + RL_PI_GRAD_NORM_EMA_INDEX, RL_PPO_CLIP_INDEX, RL_PPO_LOG_RATIO_ABS_MAX_INDEX, + RL_PPO_RATIO_CLAMP_MAX_INDEX, RL_Q_DIVERGENCE_EMA_INDEX, RL_Q_GRAD_NORM_EMA_INDEX, + RL_REWARD_SCALE_INDEX, RL_TARGET_TAU_INDEX, RL_TD_KURTOSIS_EMA_INDEX, RL_V_GRAD_NORM_EMA_INDEX, }; use ml_alpha::trainer::integrated::{ read_slice_d_pub, read_slice_i32_d_pub, IntegratedStepStats, IntegratedTrainer, @@ -558,6 +558,24 @@ fn main() -> Result<()> { "scaled_pre_clamp_max": isv[RL_MAX_ABS_SCALED_REWARD_PRE_CLAMP_INDEX], }, + // R9 — PPO importance-ratio clamp diagnostics. + // + // ratio_clamp_max : the adaptive ceiling at ISV[440], + // emitted by rl_ppo_ratio_clamp_controller + // from (1+ε) × PPO_CLAMP_MARGIN. + // log_ratio_abs_max : per-step max(|log π_new − log π_old|) + // over the batch at ISV[441], written + // by ppo_log_ratio_abs_max_b. + // + // The clamp fires when log_ratio_abs_max > ln(ratio_clamp_max). + // For ratio_clamp_max = 10, ln = 2.30. Healthy training has + // log_ratio_abs_max well below this most steps; outliers + // touch or exceed it on rare excursions which the clamp + // bounds before they pollute l_pi. + "ppo": { + "ratio_clamp_max": isv[RL_PPO_RATIO_CLAMP_MAX_INDEX], + "log_ratio_abs_max": isv[RL_PPO_LOG_RATIO_ABS_MAX_INDEX], + }, "done_count": done_count, // Action histogram: indices match `Action` enum // (0=ShortLarge, 1=ShortSmall, 2=Hold, 3=FlatFromLong, diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index 8e848710a..9ac02bf0b 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -211,6 +211,45 @@ pub const RL_LR_V_WARMUP_COUNTER_INDEX: usize = 438; /// magnitude and the clamp is materially shaping training signal. pub const RL_MAX_ABS_SCALED_REWARD_PRE_CLAMP_INDEX: usize = 439; +/// Adaptive importance-ratio clamp ceiling for the PPO surrogate. +/// Emitted by `rl_ppo_ratio_clamp_controller` from the (KL-adaptive) +/// PPO clip ε at `RL_PPO_CLIP_INDEX`. The clamp window is symmetric +/// in log space: `[1/this, this]`. Read directly by +/// `ppo_clipped_surrogate_fwd` and `_bwd` to bound the importance +/// ratio before it forms the surrogate or backward gate. +/// +/// Per `pearl_one_unbounded_signal_per_reward` + `pearl_audit_unboundedness_for_implicit_asymmetry`: +/// the ratio is the unbounded multiplicand in the PPO surrogate. PPO +/// clip(r, 1-ε, 1+ε) bounds the LOSS only when surr2 is the active +/// min — A<0,r>1+ε and A>0,r<1-ε leave the unclipped branch active +/// and the surrogate can blow up (canonical r=1e10 → l_pi=O(1e10) in +/// alpha-rl-pt67l, max|l_pi|=586). This clamp bounds the unclipped +/// branches. +/// +/// Bootstrap 10.0 (matches the safe-default constant we'd otherwise +/// hardcode). Wiener-α blend with floor 0.4 per `pearl_wiener_alpha_floor_for_nonstationary` +/// — the underlying signal (ε) is itself adapting via KL EMA so the +/// target drifts. +pub const RL_PPO_RATIO_CLAMP_MAX_INDEX: usize = 440; + +/// Per-step diagnostic: `max(|log_pi_new − log_pi_old|)` over the +/// current batch. Written by `ppo_log_ratio_abs_max_b` immediately +/// after `ppo_clipped_surrogate_fwd` consumes the new policy +/// log-probs. POINT measurement (overwritten each step). Surfaces in +/// diag JSONL as `ppo.log_ratio_abs_max`. +/// +/// Interpretation: +/// * < 0.5 — typical policy update, well inside the PPO clip band +/// * 0.5–2 — large but expected drift after several steps +/// * > 2 — outlier excursion; ratio_clamp ceiling activates +/// * > 5 — extreme excursion (raw ratio > 150); controller +/// widening ε but clamp still firing +/// +/// Cross-check: `log(isv[RL_PPO_RATIO_CLAMP_MAX_INDEX])` is the +/// ceiling — when this slot exceeds that, the kernel-side clamp +/// fired this step. +pub const RL_PPO_LOG_RATIO_ABS_MAX_INDEX: usize = 441; + /// Last RL-allocated slot index (exclusive). The integrated trainer /// extends `ISV_TOTAL_DIM` to at least this value at trainer init time. -pub const RL_SLOTS_END: usize = 440; +pub const RL_SLOTS_END: usize = 442; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index a0489af29..6360cc781 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -187,6 +187,12 @@ const RL_STEP_COUNTER_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_step_counter_update.cubin")); const RL_L2_NORM_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_l2_norm.cubin")); +// R9 — PPO ratio clamp controller (emits ISV[440]) + per-step +// log-ratio max diagnostic kernel (writes ISV[441]). +const RL_PPO_RATIO_CLAMP_CONTROLLER_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_ppo_ratio_clamp_controller.cubin")); +const PPO_LOG_RATIO_ABS_MAX_B_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/ppo_log_ratio_abs_max_b.cubin")); /// Per-head LR controller Wiener-α target. The controller kernel floors /// this at 0.4 per `pearl_wiener_alpha_floor_for_nonstationary` so the @@ -388,6 +394,11 @@ pub struct IntegratedTrainer { rl_step_counter_update_fn: CudaFunction, _rl_l2_norm_module: Arc, rl_l2_norm_fn: CudaFunction, + // R9 PPO ratio clamp + log-ratio diag. + _rl_ppo_ratio_clamp_controller_module: Arc, + rl_ppo_ratio_clamp_controller_fn: CudaFunction, + _ppo_log_ratio_abs_max_b_module: Arc, + ppo_log_ratio_abs_max_b_fn: CudaFunction, /// Per-batch step counter for `trade_duration_ema` — increments /// every step, resets on done. Zero-init at construction. pub steps_since_done_d: CudaSlice, @@ -778,6 +789,20 @@ impl IntegratedTrainer { .load_function("rl_l2_norm") .context("load rl_l2_norm")?; + // R9 — PPO ratio clamp controller + per-step log-ratio diag. + let rl_ppo_ratio_clamp_controller_module = ctx + .load_cubin(RL_PPO_RATIO_CLAMP_CONTROLLER_CUBIN.to_vec()) + .context("load rl_ppo_ratio_clamp_controller cubin")?; + let rl_ppo_ratio_clamp_controller_fn = rl_ppo_ratio_clamp_controller_module + .load_function("rl_ppo_ratio_clamp_controller") + .context("load rl_ppo_ratio_clamp_controller")?; + let ppo_log_ratio_abs_max_b_module = ctx + .load_cubin(PPO_LOG_RATIO_ABS_MAX_B_CUBIN.to_vec()) + .context("load ppo_log_ratio_abs_max_b cubin")?; + let ppo_log_ratio_abs_max_b_fn = ppo_log_ratio_abs_max_b_module + .load_function("ppo_log_ratio_abs_max_b") + .context("load ppo_log_ratio_abs_max_b")?; + // Per-batch PRNG state for the Thompson sampler. Seeded // deterministically from cfg.dqn_seed via ChaCha8 host RNG so // (cfg.dqn_seed, b_size) → identical Thompson draws across @@ -998,6 +1023,10 @@ impl IntegratedTrainer { rl_step_counter_update_fn, _rl_l2_norm_module: rl_l2_norm_module, rl_l2_norm_fn, + _rl_ppo_ratio_clamp_controller_module: rl_ppo_ratio_clamp_controller_module, + rl_ppo_ratio_clamp_controller_fn, + _ppo_log_ratio_abs_max_b_module: ppo_log_ratio_abs_max_b_module, + ppo_log_ratio_abs_max_b_fn, steps_since_done_d, trade_duration_emit_d, prev_realized_pnl_d, @@ -1119,6 +1148,15 @@ impl IntegratedTrainer { crate::rl::isv_slots::RL_MEAN_ABS_PNL_EMA_INDEX as i32, ) .context("R1 bootstrap rl_reward_scale_controller")?; + // R9 — PPO ratio clamp ceiling. At bootstrap RL_PPO_CLIP_INDEX + // is still at sentinel zero; the kernel writes + // PPO_RATIO_CLAMP_BOOTSTRAP=10.0 and returns before reading. + self.launch_isv_controller_3arg( + &self.rl_ppo_ratio_clamp_controller_fn, + alpha, + crate::rl::isv_slots::RL_PPO_CLIP_INDEX as i32, + ) + .context("R9 bootstrap rl_ppo_ratio_clamp_controller")?; self.stream .synchronize() .context("R1 controller bootstrap sync")?; @@ -1221,6 +1259,14 @@ impl IntegratedTrainer { crate::rl::isv_slots::RL_MEAN_ABS_PNL_EMA_INDEX as i32, ) .context("R5 launch rl_reward_scale_controller")?; + // R9 — PPO ratio clamp ceiling, anchored on ε at RL_PPO_CLIP_INDEX + // (which the rl_ppo_clip_controller just refreshed two calls up). + self.launch_isv_controller_3arg( + &self.rl_ppo_ratio_clamp_controller_fn, + alpha, + crate::rl::isv_slots::RL_PPO_CLIP_INDEX as i32, + ) + .context("R9 launch rl_ppo_ratio_clamp_controller")?; Ok(()) } @@ -1986,6 +2032,34 @@ impl IntegratedTrainer { .context("rl_kl_approx_b launch")?; } } + + // R9 diag — per-batch max |log π_new − log π_old| → ISV[441]. + // Same input buffers as rl_kl_approx_b but max-reduce |diff| + // instead of mean signed-diff. Surfaces the largest single + // policy excursion in this step's batch for the diag JSONL. + // Cross-check: log(isv[RL_PPO_RATIO_CLAMP_MAX_INDEX]) is the + // ceiling — when this slot exceeds that, the ratio clamp fired + // this step. + { + let block_dim = (b_size.next_power_of_two() as u32).max(1); + let cfg = LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (block_dim, 1, 1), + shared_mem_bytes: (block_dim as usize * std::mem::size_of::()) as u32, + }; + let b_size_i = b_size as i32; + let mut launch = self.stream.launch_builder(&self.ppo_log_ratio_abs_max_b_fn); + launch + .arg(&self.log_pi_old_d) + .arg(&pi_log_prob_d) + .arg(&mut self.isv_d) + .arg(&b_size_i); + unsafe { + launch + .launch(cfg) + .context("ppo_log_ratio_abs_max_b launch")?; + } + } self.launch_ema_update_per_step( crate::rl::isv_slots::RL_KL_PI_EMA_INDEX, RL_LR_CONTROLLER_ALPHA, diff --git a/crates/ml-alpha/tests/isv_bootstrap.rs b/crates/ml-alpha/tests/isv_bootstrap.rs index f74eb7b92..a6e204a53 100644 --- a/crates/ml-alpha/tests/isv_bootstrap.rs +++ b/crates/ml-alpha/tests/isv_bootstrap.rs @@ -24,8 +24,8 @@ use ml_alpha::rl::isv_slots::{ RL_ENTROPY_COEF_INDEX, RL_GAMMA_INDEX, RL_MEAN_TRADE_DURATION_EMA_INDEX, - RL_N_ROLLOUT_STEPS_INDEX, RL_PER_ALPHA_INDEX, RL_PPO_CLIP_INDEX, RL_REWARD_SCALE_INDEX, - RL_SLOTS_END, RL_TARGET_TAU_INDEX, + RL_N_ROLLOUT_STEPS_INDEX, RL_PER_ALPHA_INDEX, RL_PPO_CLIP_INDEX, + RL_PPO_RATIO_CLAMP_MAX_INDEX, RL_REWARD_SCALE_INDEX, RL_SLOTS_END, RL_TARGET_TAU_INDEX, }; use ml_alpha::trainer::integrated::{IntegratedTrainer, IntegratedTrainerConfig}; use ml_alpha::trainer::perception::PerceptionTrainerConfig; @@ -131,19 +131,33 @@ fn g1_isv_bootstrap_writes_canonical_values() { isv[RL_REWARD_SCALE_INDEX] ); - // Invariant: the 7 EMA-input slots (417..424) are NOT yet - // bootstrapped. Phase R3 wires the EMA producer kernels that fill - // these; until then they MUST remain at `alloc_zeros` sentinel 0. - // Asserting this here protects against accidental controller-side - // writes to the wrong slot (a defect a one-character bug could - // introduce, given the 7 slots are sequential). + // Invariant: the EMA-input slots are NOT yet bootstrapped. Phase + // R3 wires the EMA producer kernels that fill these; until then + // they MUST remain at `alloc_zeros` sentinel 0. Asserting this + // here protects against accidental controller-side writes to the + // wrong slot (a defect a one-character bug could introduce, given + // the slots are sequential). + // + // Exception: RL_PPO_RATIO_CLAMP_MAX_INDEX is a controller-OUTPUT + // slot (the R9 ppo ratio clamp ceiling). Its bootstrap path fires + // during `with_controllers_bootstrapped`, leaving the slot at + // PPO_RATIO_CLAMP_BOOTSTRAP = 10.0 — checked separately below. for slot in RL_MEAN_TRADE_DURATION_EMA_INDEX..RL_SLOTS_END { + if slot == RL_PPO_RATIO_CLAMP_MAX_INDEX { + continue; + } assert_eq!( isv[slot], 0.0, "ISV[{slot}] expected sentinel 0.0 (R3 wires EMA producers), got {}", isv[slot] ); } + // R9 — PPO ratio clamp ceiling bootstrap. + assert!( + (isv[RL_PPO_RATIO_CLAMP_MAX_INDEX] - 10.0).abs() < EPS, + "ISV[ppo_ratio_clamp_max={RL_PPO_RATIO_CLAMP_MAX_INDEX}] expected 10.0, got {}", + isv[RL_PPO_RATIO_CLAMP_MAX_INDEX] + ); eprintln!( "G1 OK — bootstraps: γ={:.4} τ={:.4} ε={:.4} entropy_coef={:.4} \ diff --git a/crates/ml-alpha/tests/r5_controllers_and_soft_update.rs b/crates/ml-alpha/tests/r5_controllers_and_soft_update.rs index 58ddc2b5e..1b703c117 100644 --- a/crates/ml-alpha/tests/r5_controllers_and_soft_update.rs +++ b/crates/ml-alpha/tests/r5_controllers_and_soft_update.rs @@ -29,8 +29,8 @@ use ml_alpha::rl::isv_slots::{ RL_ADVANTAGE_VAR_RATIO_EMA_INDEX, RL_ENTROPY_COEF_INDEX, RL_ENTROPY_OBSERVED_EMA_INDEX, RL_GAMMA_INDEX, RL_KL_PI_EMA_INDEX, RL_MEAN_ABS_PNL_EMA_INDEX, RL_MEAN_TRADE_DURATION_EMA_INDEX, RL_N_ROLLOUT_STEPS_INDEX, RL_PER_ALPHA_INDEX, - RL_PPO_CLIP_INDEX, RL_Q_DIVERGENCE_EMA_INDEX, RL_REWARD_SCALE_INDEX, RL_SLOTS_END, - RL_TARGET_TAU_INDEX, RL_TD_KURTOSIS_EMA_INDEX, + RL_PPO_CLIP_INDEX, RL_PPO_RATIO_CLAMP_MAX_INDEX, RL_Q_DIVERGENCE_EMA_INDEX, + RL_REWARD_SCALE_INDEX, RL_SLOTS_END, RL_TARGET_TAU_INDEX, RL_TD_KURTOSIS_EMA_INDEX, }; use ml_alpha::trainer::integrated::{IntegratedTrainer, IntegratedTrainerConfig}; use ml_alpha::trainer::perception::PerceptionTrainerConfig; @@ -111,10 +111,16 @@ fn g3_per_step_controllers_move_isv_outputs_when_fed_real_emas() { assert_eq!(isv_before[RL_N_ROLLOUT_STEPS_INDEX], ROLLOUT_BOOTSTRAP); assert_eq!(isv_before[RL_PER_ALPHA_INDEX], PER_ALPHA_BOOTSTRAP); assert_eq!(isv_before[RL_REWARD_SCALE_INDEX], REWARD_SCALE_BOOTSTRAP); - // And all 7 EMA-input slots are still at sentinel zero. + // And all EMA-input slots are still at sentinel zero. Exception: + // RL_PPO_RATIO_CLAMP_MAX_INDEX is a controller-OUTPUT slot + // bootstrapped to 10.0 by `with_controllers_bootstrapped`. for slot in RL_MEAN_TRADE_DURATION_EMA_INDEX..RL_SLOTS_END { + if slot == RL_PPO_RATIO_CLAMP_MAX_INDEX { + continue; + } assert_eq!(isv_before[slot], 0.0); } + assert_eq!(isv_before[RL_PPO_RATIO_CLAMP_MAX_INDEX], 10.0); // Populate each EMA-input slot with a non-zero value via the // R3 ema_update_per_step bootstrap path (sentinel-zero → first