feat(rl): ISV-driven PPO importance-ratio clamp + log-ratio diagnostic
pt67l confirmed reward-scale + V-target clamp eliminate V regression
spikes — but exposed a residual: |l_pi| max=586 with mean 0.22. Root
cause: PPO's clip(r, 1-ε, 1+ε) bounds the loss only when surr2 is
the active min. The unclipped branch IS active when A<0,r>1+ε
(surr1=A·r is then more negative than surr2=A·(1+ε), so min selects
surr1) and when A>0,r<1-ε. In the first case `r` can blow up: we've
seen r reach 1e10 from policy drift over a multi-step rollout
producing l_pi=O(1e10) spikes that contaminate the loss-balance
controller and the LR controller's plateau detection.
## Fix: ISV-driven ratio clamp
Per `feedback_isv_for_adaptive_bounds` and
`pearl_controller_anchors_isv_driven`: the clamp ceiling lives in
ISV[RL_PPO_RATIO_CLAMP_MAX_INDEX = 440], not as a hardcoded #define.
New controller `rl_ppo_ratio_clamp_controller.cu`:
* Anchors on the (already KL-adaptive) PPO clip ε at ISV[402]
* target = (1 + ε) × PPO_CLAMP_MARGIN (MARGIN = 10.0)
* Wiener-α blend with floor 0.4 per
pearl_wiener_alpha_floor_for_nonstationary (ε is non-stationary)
* Permanent floor 2.0 / ceiling 1000 per
pearl_blend_formulas_must_have_permanent_floor
* Bootstrap 10.0, replace-directly on first non-bootstrap ε
observation per pearl_first_observation_bootstrap
When ε is small (rl_ppo_clip_controller seeing low KL → tight clip
band), the ratio clamp tightens — outliers should be rare anomalies.
When ε widens (large KL → wide clip band), the clamp widens
proportionally — outliers are expected so we permit more
magnitude before bounding.
## Wiring
ppo_clipped_surrogate_fwd and _bwd both read
isv[RL_PPO_RATIO_CLAMP_MAX_INDEX] and clamp ratio to
[1/ratio_max, ratio_max] before forming surr1/surr2. The clamp is
forward-only in effect (bwd gates pg_grad inside [1-ε, 1+ε] anyway
so gradients were already bounded), but bounding the FORWARD ratio
keeps l_pi sane for the controllers downstream.
The new controller is wired into both:
* `with_controllers_bootstrapped` — bootstrap launch alongside
the other 7 R1 controllers
* `launch_rl_controllers_per_step` — per-step refresh alongside
the other 7 R5 controllers
## Diagnostic: per-step max |log_ratio|
New kernel `ppo_log_ratio_abs_max_b.cu` (same tree-reduce shape as
rl_kl_approx_b) writes per-batch max(|log π_new − log π_old|) to
ISV[RL_PPO_LOG_RATIO_ABS_MAX_INDEX = 441]. Launched right after
rl_kl_approx_b (uses the same log_pi_old_d + pi_log_prob_d inputs).
Surfaces in diag JSONL as:
"ppo": {
"ratio_clamp_max": isv[440], # adaptive ceiling
"log_ratio_abs_max": isv[441] # per-step observed max
}
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.
## Slot allocation
RL_PPO_RATIO_CLAMP_MAX_INDEX = 440 (controller output)
RL_PPO_LOG_RATIO_ABS_MAX_INDEX = 441 (per-step diag)
RL_SLOTS_END = 442 (was 440)
## Test updates
G1 (isv_bootstrap) + G3 (r5_controllers) blanket-assert ISV[417..END]
== 0.0 to catch slot-wiring bugs. Slot 440 is now a controller
OUTPUT bootstrapped to 10.0, so both tests skip it in the loop and
assert == 10.0 separately.
## Verified gates (local sm_86)
G1 isv_bootstrap ✅ (with new slot-440 assertion)
G3 controllers ✅
G4 target_update ✅
G6 r7d_per_wiring ✅
integrated_smoke ✅
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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
|
||||
|
||||
@@ -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];
|
||||
|
||||
|
||||
58
crates/ml-alpha/cuda/ppo_log_ratio_abs_max_b.cu
Normal file
58
crates/ml-alpha/cuda/ppo_log_ratio_abs_max_b.cu
Normal file
@@ -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];
|
||||
}
|
||||
}
|
||||
116
crates/ml-alpha/cuda/rl_ppo_ratio_clamp_controller.cu
Normal file
116
crates/ml-alpha/cuda/rl_ppo_ratio_clamp_controller.cu
Normal file
@@ -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;
|
||||
}
|
||||
@@ -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,
|
||||
|
||||
@@ -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;
|
||||
|
||||
@@ -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<CudaModule>,
|
||||
rl_l2_norm_fn: CudaFunction,
|
||||
// R9 PPO ratio clamp + log-ratio diag.
|
||||
_rl_ppo_ratio_clamp_controller_module: Arc<CudaModule>,
|
||||
rl_ppo_ratio_clamp_controller_fn: CudaFunction,
|
||||
_ppo_log_ratio_abs_max_b_module: Arc<CudaModule>,
|
||||
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<i32>,
|
||||
@@ -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::<f32>()) 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,
|
||||
|
||||
@@ -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} \
|
||||
|
||||
@@ -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
|
||||
|
||||
Reference in New Issue
Block a user