From 66115007abe108cbed8ef76d09ec69729fc73999 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sat, 23 May 2026 22:26:12 +0200 Subject: [PATCH] fix(rl): ISV-driven output clamp on streaming var/kurtosis kernels MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit gxhr8 confirmed the streaming kernels work — both formerly-dead controllers (rl_rollout_steps, rl_per_alpha) now adapt instead of pegging at MIN. But the unclamped streaming outputs reached advantage_var_ratio = 3e5 (when streaming-mean passed through zero and `var/|mean|` blew up under the 1e-6 denominator floor) and td_kurtosis = 50.6, pegging both downstream controllers at MAX instead. Per_α at MAX over-concentrates PER sampling on outliers, which hurts distributional Q learning (best l_q window regressed from 2.41 → 2.69 between pdgxn and gxhr8). ## Fix: ISV-resident output clamp ceilings Two new ISV slots hold the streaming-kernel output ceilings: RL_ADV_VAR_RATIO_CLAMP_INDEX = 447 (default 100.0) RL_TD_KURTOSIS_CLAMP_INDEX = 448 (default 30.0) * 100.0 for var_ratio = 1000× ADV_VAR_RATIO_TARGET (= 0.1) — wide enough that healthy signal (typical 1-10) passes through, tight enough that 3e5 outliers don't peg rollout_steps. * 30.0 for kurtosis = 3× (KURT_GAUSSIAN + KURT_LIFT_SCALE) — lets the full per_α response range engage on heavy-tailed signal (≤ 10), bounds runaway above that. Per `feedback_isv_for_adaptive_bounds`: the clamps live in ISV (visible in diag, modifiable at runtime via re-launching the init kernel or a future adaptive controller) rather than as kernel-side `#define`s. ## Seeding (no HtoD per feedback_no_htod_htoh_only_mapped_pinned) New device kernel `rl_streaming_clamp_init.cu` — single thread, writes both clamp ceilings directly to ISV. Launched once at the end of `with_controllers_bootstrapped` alongside the 8 existing controller-bootstrap launches. Zero host→device transfer. ## Diag bake-in (per user request "ensure to bake in diags") JSONL gains a new `streaming` block exposing: * `streaming.adv_var.{mean, m2, clamp}` * `streaming.td_kurt.{mean, m2, m4, clamp}` Cross-check: when consumer-input slot (RL_ADVANTAGE_VAR_RATIO_EMA_INDEX or RL_TD_KURTOSIS_EMA_INDEX) reads exactly the same value as `streaming.*.clamp`, the clamp fired this step. ## Test updates G1 (isv_bootstrap) + G3 (r5_controllers) blanket-assert that ISV[417..END] is sentinel-zero at bootstrap. Both new slots are seeded to non-zero values by rl_streaming_clamp_init during bootstrap, so both tests skip these slots in the loop and assert the seeded values separately. ## Verified gates (local sm_86) G1 isv_bootstrap ✅ G3 controllers ✅ G4 target_update ✅ integrated_smoke ✅ Co-Authored-By: Claude Opus 4.7 --- crates/ml-alpha/build.rs | 1 + crates/ml-alpha/cuda/rl_kurtosis_streaming.cu | 19 +++++- .../ml-alpha/cuda/rl_streaming_clamp_init.cu | 35 +++++++++++ .../cuda/rl_var_over_abs_mean_streaming.cu | 28 ++++++--- crates/ml-alpha/examples/alpha_rl_train.rs | 29 ++++++++- crates/ml-alpha/src/rl/isv_slots.rs | 30 +++++++++- crates/ml-alpha/src/trainer/integrated.rs | 59 ++++++++++++++++++- crates/ml-alpha/tests/isv_bootstrap.rs | 25 ++++++-- .../tests/r5_controllers_and_soft_update.rs | 18 ++++-- 9 files changed, 220 insertions(+), 24 deletions(-) create mode 100644 crates/ml-alpha/cuda/rl_streaming_clamp_init.cu diff --git a/crates/ml-alpha/build.rs b/crates/ml-alpha/build.rs index d19648333..5e3898ec6 100644 --- a/crates/ml-alpha/build.rs +++ b/crates/ml-alpha/build.rs @@ -64,6 +64,7 @@ const KERNELS: &[&str] = &[ // (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 + "rl_streaming_clamp_init", // RL R9: device-side seeder for streaming-kernel output clamp ceilings (ISV[447], ISV[448]) — no HtoD ]; // Cache bust v31 — five new reduce / derive kernels populate the input diff --git a/crates/ml-alpha/cuda/rl_kurtosis_streaming.cu b/crates/ml-alpha/cuda/rl_kurtosis_streaming.cu index 537450fb2..26ebbf7d1 100644 --- a/crates/ml-alpha/cuda/rl_kurtosis_streaming.cu +++ b/crates/ml-alpha/cuda/rl_kurtosis_streaming.cu @@ -46,6 +46,14 @@ #define STREAM_ALPHA 0.05f #define M2_SQ_FLOOR 1e-12f +// clamp_slot: ISV slot holding the output ceiling +// (RL_TD_KURTOSIS_CLAMP_INDEX = 448). Seeded once at +// trainer init by rl_streaming_clamp_init. Bounds the +// kurtosis output before it feeds the per_α controller — +// without this, kurtosis can reach 50+ (seen in +// alpha-rl-gxhr8 fold0) and pegs per_α at MAX, which +// over-concentrates PER sampling and hurts distributional +// Q learning. extern "C" __global__ void rl_kurtosis_streaming( const float* __restrict__ x, float* __restrict__ isv, @@ -53,7 +61,8 @@ extern "C" __global__ void rl_kurtosis_streaming( int out_slot, int mean_slot, int m2_slot, - int m4_slot + int m4_slot, + int clamp_slot ) { if (threadIdx.x != 0 || blockIdx.x != 0) return; @@ -100,6 +109,12 @@ extern "C" __global__ void rl_kurtosis_streaming( // Kurtosis = M4 / M2² (Pearson). Permanent-floor on M2² prevents // div-by-zero before the streaming variance has accumulated. + // ISV-driven clamp at isv[clamp_slot] bounds outlier kurtosis + // (typical 3-15 for heavy-tailed real signal; >30 is numerical + // artefact from tiny M2 denominator — clamp default 30.0 lets + // healthy signal through unchanged). const float m2_sq = m2_new * m2_new; - isv[out_slot] = m4_new / (m2_sq + M2_SQ_FLOOR); + const float raw_kurt = m4_new / (m2_sq + M2_SQ_FLOOR); + const float clamp_max = isv[clamp_slot]; + isv[out_slot] = fminf(raw_kurt, clamp_max); } diff --git a/crates/ml-alpha/cuda/rl_streaming_clamp_init.cu b/crates/ml-alpha/cuda/rl_streaming_clamp_init.cu new file mode 100644 index 000000000..98fcf948f --- /dev/null +++ b/crates/ml-alpha/cuda/rl_streaming_clamp_init.cu @@ -0,0 +1,35 @@ +// rl_streaming_clamp_init.cu — single-thread device-side seeder for +// the streaming-kernel output clamp ceilings at +// ISV[RL_ADV_VAR_RATIO_CLAMP_INDEX = 447] and +// ISV[RL_TD_KURTOSIS_CLAMP_INDEX = 448]. +// +// Launched once from `with_controllers_bootstrapped` after ISV is +// alloc_zeros'd. Writes design-constant clamp ceilings directly on +// device — no host→device transfer, in line with +// `feedback_no_htod_htoh_only_mapped_pinned`. +// +// The streaming variance and kurtosis kernels then read these slots +// each step and clamp their output before writing to the consumer +// controllers' EMA-input slots. Without the clamp, streaming +// `var/|mean|` reached 3e5 in alpha-rl-gxhr8 fold0 (when streaming +// mean passed through zero — `var / max(|μ|, 1e-6)` blows up), +// pegging the rollout_steps controller at MAX; streaming kurtosis +// reached 50.6, pegging per_α at MAX which over-concentrated PER +// sampling and hurt distributional Q learning. +// +// Per `feedback_isv_for_adaptive_bounds`: the ceilings live in ISV +// (visible in diag, modifiable at runtime by re-launching this +// kernel or by a future adaptive controller) rather than as kernel- +// side `#define`s. + +extern "C" __global__ void rl_streaming_clamp_init( + float* __restrict__ isv, + int adv_var_clamp_slot, + float adv_var_clamp_val, + int td_kurt_clamp_slot, + float td_kurt_clamp_val +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + isv[adv_var_clamp_slot] = adv_var_clamp_val; + isv[td_kurt_clamp_slot] = td_kurt_clamp_val; +} diff --git a/crates/ml-alpha/cuda/rl_var_over_abs_mean_streaming.cu b/crates/ml-alpha/cuda/rl_var_over_abs_mean_streaming.cu index 9f9561197..2e1430a20 100644 --- a/crates/ml-alpha/cuda/rl_var_over_abs_mean_streaming.cu +++ b/crates/ml-alpha/cuda/rl_var_over_abs_mean_streaming.cu @@ -41,18 +41,25 @@ #define STREAM_ALPHA 0.05f #define ABS_MEAN_FLOOR 1e-6f -// out_slot: ISV slot to write var/|mean| to (e.g. -// RL_ADVANTAGE_VAR_RATIO_EMA_INDEX). -// mean_slot: streaming mean state (e.g. -// RL_ADV_VAR_STREAM_MEAN_INDEX). -// m2_slot: streaming M2 state. +// out_slot: ISV slot to write var/|mean| to (e.g. +// RL_ADVANTAGE_VAR_RATIO_EMA_INDEX). +// mean_slot: streaming mean state (e.g. RL_ADV_VAR_STREAM_MEAN_INDEX). +// m2_slot: streaming M2 state. +// clamp_slot: ISV slot holding the output ceiling +// (RL_ADV_VAR_RATIO_CLAMP_INDEX = 447). Seeded once at +// trainer init by rl_streaming_clamp_init. Bounds the +// `var/|mean|` output before it feeds the rollout_steps +// controller — without this, when streaming mean passes +// through zero the ratio blows up (3e5 outliers in +// alpha-rl-gxhr8 fold0) and pegs rollout_steps at MAX. extern "C" __global__ void rl_var_over_abs_mean_streaming( const float* __restrict__ x, float* __restrict__ isv, int b_size, int out_slot, int mean_slot, - int m2_slot + int m2_slot, + int clamp_slot ) { if (threadIdx.x != 0 || blockIdx.x != 0) return; @@ -90,7 +97,12 @@ extern "C" __global__ void rl_var_over_abs_mean_streaming( isv[m2_slot] = m2_new; // var/|mean| → controller input slot directly (no separate EMA - // update — the streaming kernel IS the EMA). + // update — the streaming kernel IS the EMA). ISV-driven clamp at + // isv[clamp_slot] bounds outlier magnitudes (typical: 1–10; + // outliers when streaming-mean ≈ 0 can reach 1e5+; clamp default + // 100.0 lets healthy signal through unchanged). const float abs_mean = fmaxf(fabsf(mean_new), ABS_MEAN_FLOOR); - isv[out_slot] = m2_new / abs_mean; + const float raw_out = m2_new / abs_mean; + const float clamp_max = isv[clamp_slot]; + isv[out_slot] = fminf(raw_out, clamp_max); } diff --git a/crates/ml-alpha/examples/alpha_rl_train.rs b/crates/ml-alpha/examples/alpha_rl_train.rs index 271cf9307..a48e28542 100644 --- a/crates/ml-alpha/examples/alpha_rl_train.rs +++ b/crates/ml-alpha/examples/alpha_rl_train.rs @@ -45,11 +45,14 @@ use ml_alpha::rl::isv_slots::{ RL_LR_Q_LOSS_EMA_INDEX, RL_LR_Q_STEPS_SINCE_BEST_INDEX, RL_LR_Q_WARMUP_COUNTER_INDEX, RL_LR_V_BEST_LOSS_INDEX, RL_LR_V_INDEX, RL_LR_V_LOSS_EMA_INDEX, RL_LR_V_STEPS_SINCE_BEST_INDEX, RL_LR_V_WARMUP_COUNTER_INDEX, + RL_ADV_VAR_RATIO_CLAMP_INDEX, RL_ADV_VAR_STREAM_M2_INDEX, RL_ADV_VAR_STREAM_MEAN_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_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, + RL_REWARD_SCALE_INDEX, RL_TARGET_TAU_INDEX, RL_TD_KURTOSIS_CLAMP_INDEX, + RL_TD_KURTOSIS_EMA_INDEX, RL_TD_KURT_STREAM_M2_INDEX, RL_TD_KURT_STREAM_M4_INDEX, + RL_TD_KURT_STREAM_MEAN_INDEX, RL_V_GRAD_NORM_EMA_INDEX, }; use ml_alpha::trainer::integrated::{ read_slice_d_pub, read_slice_i32_d_pub, IntegratedStepStats, IntegratedTrainer, @@ -576,6 +579,30 @@ fn main() -> Result<()> { "ratio_clamp_max": isv[RL_PPO_RATIO_CLAMP_MAX_INDEX], "log_ratio_abs_max": isv[RL_PPO_LOG_RATIO_ABS_MAX_INDEX], }, + // R9 — streaming-kernel state for the variance-ratio + + // kurtosis estimators that feed rollout_steps and per_α. + // * `*_stream.{mean, m2, m4}` are the Welford-EMA + // internal state slots (442-446) — confirms the + // kernel is folding observations across STEPS. + // * `*_clamp` is the ISV-resident output ceiling + // (447, 448) seeded once at trainer init by + // rl_streaming_clamp_init — visible here so the diag + // can spot when the clamp is firing + // (post-clamp EMA at the consumer-input slot equals + // `clamp` ↔ clamp active). + "streaming": { + "adv_var": { + "mean": isv[RL_ADV_VAR_STREAM_MEAN_INDEX], + "m2": isv[RL_ADV_VAR_STREAM_M2_INDEX], + "clamp": isv[RL_ADV_VAR_RATIO_CLAMP_INDEX], + }, + "td_kurt": { + "mean": isv[RL_TD_KURT_STREAM_MEAN_INDEX], + "m2": isv[RL_TD_KURT_STREAM_M2_INDEX], + "m4": isv[RL_TD_KURT_STREAM_M4_INDEX], + "clamp": isv[RL_TD_KURTOSIS_CLAMP_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 3960b2ec3..a68967278 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -283,6 +283,34 @@ pub const RL_TD_KURT_STREAM_M2_INDEX: usize = 445; /// Kurtosis = M4 / (M2² + ε). pub const RL_TD_KURT_STREAM_M4_INDEX: usize = 446; +/// Output clamp ceiling for the streaming advantage variance ratio. +/// Bounds the EMA value written to `RL_ADVANTAGE_VAR_RATIO_EMA_INDEX` +/// so the downstream `rl_rollout_steps_controller` isn't driven by +/// outliers when the streaming-mean passes through zero +/// (`var/|mean|` blows up under tiny denominator). Default 100.0 — +/// 1000× the consumer controller's `ADV_VAR_RATIO_TARGET = 0.1` — +/// generous enough that the clamp is a no-op for healthy signal +/// (typical values 1–10) but tight enough that the 3e5 outliers seen +/// in alpha-rl-gxhr8 fold0 (commit 39f90f372) don't peg the +/// rollout_steps controller at MAX. +/// +/// Per `feedback_isv_for_adaptive_bounds`: the clamp lives in ISV +/// (visible in diag, modifiable at runtime), not as a kernel-side +/// `#define`. Per `pearl_first_observation_bootstrap`: trainer seeds +/// the slot at init; streaming kernel does NOT bootstrap because +/// sentinel-zero would mean "no clamp" which defeats the purpose. +pub const RL_ADV_VAR_RATIO_CLAMP_INDEX: usize = 447; + +/// Output clamp ceiling for the streaming TD kurtosis. Bounds the +/// EMA value written to `RL_TD_KURTOSIS_EMA_INDEX` so the +/// downstream `rl_per_alpha_controller` isn't driven to MAX by +/// outlier kurtosis (which over-concentrates PER sampling and hurts +/// distributional Q learning). Default 30.0 — 3× the consumer's +/// `(KURT_GAUSSIAN + KURT_LIFT_SCALE) = 10` saturation point — +/// preserves the controller's full responsive range while bounding +/// runaway. +pub const RL_TD_KURTOSIS_CLAMP_INDEX: usize = 448; + /// 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 = 447; +pub const RL_SLOTS_END: usize = 449; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index 78d7a2350..c30b948bc 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -193,6 +193,8 @@ 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")); +const RL_STREAMING_CLAMP_INIT_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_streaming_clamp_init.cubin")); /// Per-head LR controller Wiener-α target. The controller kernel floors /// this at 0.4 per `pearl_wiener_alpha_floor_for_nonstationary` so the @@ -399,6 +401,8 @@ pub struct IntegratedTrainer { rl_ppo_ratio_clamp_controller_fn: CudaFunction, _ppo_log_ratio_abs_max_b_module: Arc, ppo_log_ratio_abs_max_b_fn: CudaFunction, + _rl_streaming_clamp_init_module: Arc, + rl_streaming_clamp_init_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, @@ -605,6 +609,11 @@ impl IntegratedTrainer { .alloc_zeros::(RL_SLOTS_END) .context("alloc isv_d")?; let isv_host = vec![0.0_f32; RL_SLOTS_END]; + // R9 clamp-ceiling seeding happens via the + // rl_streaming_clamp_init kernel launched alongside the + // controller bootstraps below — per + // `feedback_no_htod_htoh_only_mapped_pinned` we never + // host-seed device-resident state. // Combined-encoder-grad slot for item 1 (perception backward). // Sized [B × HIDDEN_DIM]; the perception config holds n_batch. @@ -808,6 +817,12 @@ impl IntegratedTrainer { 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")?; + let rl_streaming_clamp_init_module = ctx + .load_cubin(RL_STREAMING_CLAMP_INIT_CUBIN.to_vec()) + .context("load rl_streaming_clamp_init cubin")?; + let rl_streaming_clamp_init_fn = rl_streaming_clamp_init_module + .load_function("rl_streaming_clamp_init") + .context("load rl_streaming_clamp_init")?; // Per-batch PRNG state for the Thompson sampler. Seeded // deterministically from cfg.dqn_seed via ChaCha8 host RNG so @@ -1033,6 +1048,8 @@ impl IntegratedTrainer { 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, + _rl_streaming_clamp_init_module: rl_streaming_clamp_init_module, + rl_streaming_clamp_init_fn, steps_since_done_d, trade_duration_emit_d, prev_realized_pnl_d, @@ -1163,6 +1180,38 @@ impl IntegratedTrainer { crate::rl::isv_slots::RL_PPO_CLIP_INDEX as i32, ) .context("R9 bootstrap rl_ppo_ratio_clamp_controller")?; + // R9 — seed streaming-kernel output clamp ceilings device-side + // (no HtoD per feedback_no_htod_htoh_only_mapped_pinned). + // Single-thread kernel writes design constants to ISV[447] + // and ISV[448] which the streaming variance / kurtosis kernels + // read each step to clamp their outputs. + { + let adv_var_slot: i32 = + crate::rl::isv_slots::RL_ADV_VAR_RATIO_CLAMP_INDEX as i32; + let adv_var_val: f32 = 100.0; + let td_kurt_slot: i32 = + crate::rl::isv_slots::RL_TD_KURTOSIS_CLAMP_INDEX as i32; + let td_kurt_val: f32 = 30.0; + let cfg = LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }; + let mut launch = self + .stream + .launch_builder(&self.rl_streaming_clamp_init_fn); + launch + .arg(&self.isv_d) + .arg(&adv_var_slot) + .arg(&adv_var_val) + .arg(&td_kurt_slot) + .arg(&td_kurt_val); + unsafe { + launch + .launch(cfg) + .context("R9 launch rl_streaming_clamp_init")?; + } + } self.stream .synchronize() .context("R1 controller bootstrap sync")?; @@ -2937,6 +2986,8 @@ impl IntegratedTrainer { crate::rl::isv_slots::RL_ADV_VAR_STREAM_MEAN_INDEX as i32; let m2_slot: i32 = crate::rl::isv_slots::RL_ADV_VAR_STREAM_M2_INDEX as i32; + let clamp_slot: i32 = + crate::rl::isv_slots::RL_ADV_VAR_RATIO_CLAMP_INDEX as i32; let mut launch = self.stream.launch_builder( &self.rl_var_over_abs_mean_streaming_fn, ); @@ -2946,7 +2997,8 @@ impl IntegratedTrainer { .arg(&b_size_i) .arg(&out_slot) .arg(&mean_slot) - .arg(&m2_slot); + .arg(&m2_slot) + .arg(&clamp_slot); unsafe { launch .launch(cfg) @@ -3014,6 +3066,8 @@ impl IntegratedTrainer { crate::rl::isv_slots::RL_TD_KURT_STREAM_M2_INDEX as i32; let m4_slot: i32 = crate::rl::isv_slots::RL_TD_KURT_STREAM_M4_INDEX as i32; + let clamp_slot: i32 = + crate::rl::isv_slots::RL_TD_KURTOSIS_CLAMP_INDEX as i32; let mut launch = self.stream.launch_builder( &self.rl_kurtosis_streaming_fn, ); @@ -3024,7 +3078,8 @@ impl IntegratedTrainer { .arg(&out_slot) .arg(&mean_slot) .arg(&m2_slot) - .arg(&m4_slot); + .arg(&m4_slot) + .arg(&clamp_slot); unsafe { launch .launch(cfg) diff --git a/crates/ml-alpha/tests/isv_bootstrap.rs b/crates/ml-alpha/tests/isv_bootstrap.rs index a6e204a53..0b3c9006f 100644 --- a/crates/ml-alpha/tests/isv_bootstrap.rs +++ b/crates/ml-alpha/tests/isv_bootstrap.rs @@ -23,9 +23,10 @@ //! `cargo test -p ml-alpha --test isv_bootstrap -- --ignored --nocapture` 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_PPO_RATIO_CLAMP_MAX_INDEX, RL_REWARD_SCALE_INDEX, RL_SLOTS_END, RL_TARGET_TAU_INDEX, + RL_ADV_VAR_RATIO_CLAMP_INDEX, 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_PPO_RATIO_CLAMP_MAX_INDEX, RL_REWARD_SCALE_INDEX, RL_SLOTS_END, + RL_TARGET_TAU_INDEX, RL_TD_KURTOSIS_CLAMP_INDEX, }; use ml_alpha::trainer::integrated::{IntegratedTrainer, IntegratedTrainerConfig}; use ml_alpha::trainer::perception::PerceptionTrainerConfig; @@ -143,7 +144,11 @@ fn g1_isv_bootstrap_writes_canonical_values() { // 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 { + // R9 controller-OUTPUT slots that bootstrap to non-zero values. + if slot == RL_PPO_RATIO_CLAMP_MAX_INDEX + || slot == RL_ADV_VAR_RATIO_CLAMP_INDEX + || slot == RL_TD_KURTOSIS_CLAMP_INDEX + { continue; } assert_eq!( @@ -158,6 +163,18 @@ fn g1_isv_bootstrap_writes_canonical_values() { "ISV[ppo_ratio_clamp_max={RL_PPO_RATIO_CLAMP_MAX_INDEX}] expected 10.0, got {}", isv[RL_PPO_RATIO_CLAMP_MAX_INDEX] ); + // R9 — streaming-kernel output clamp ceilings (seeded device-side + // by rl_streaming_clamp_init during bootstrap). + assert!( + (isv[RL_ADV_VAR_RATIO_CLAMP_INDEX] - 100.0).abs() < EPS, + "ISV[adv_var_ratio_clamp={RL_ADV_VAR_RATIO_CLAMP_INDEX}] expected 100.0, got {}", + isv[RL_ADV_VAR_RATIO_CLAMP_INDEX] + ); + assert!( + (isv[RL_TD_KURTOSIS_CLAMP_INDEX] - 30.0).abs() < EPS, + "ISV[td_kurtosis_clamp={RL_TD_KURTOSIS_CLAMP_INDEX}] expected 30.0, got {}", + isv[RL_TD_KURTOSIS_CLAMP_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 1b703c117..922591376 100644 --- a/crates/ml-alpha/tests/r5_controllers_and_soft_update.rs +++ b/crates/ml-alpha/tests/r5_controllers_and_soft_update.rs @@ -26,11 +26,12 @@ use cudarc::driver::CudaStream; 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_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, + RL_ADVANTAGE_VAR_RATIO_EMA_INDEX, RL_ADV_VAR_RATIO_CLAMP_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_PPO_RATIO_CLAMP_MAX_INDEX, + RL_Q_DIVERGENCE_EMA_INDEX, RL_REWARD_SCALE_INDEX, RL_SLOTS_END, RL_TARGET_TAU_INDEX, + RL_TD_KURTOSIS_CLAMP_INDEX, RL_TD_KURTOSIS_EMA_INDEX, }; use ml_alpha::trainer::integrated::{IntegratedTrainer, IntegratedTrainerConfig}; use ml_alpha::trainer::perception::PerceptionTrainerConfig; @@ -115,12 +116,17 @@ fn g3_per_step_controllers_move_isv_outputs_when_fed_real_emas() { // 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 { + if slot == RL_PPO_RATIO_CLAMP_MAX_INDEX + || slot == RL_ADV_VAR_RATIO_CLAMP_INDEX + || slot == RL_TD_KURTOSIS_CLAMP_INDEX + { continue; } assert_eq!(isv_before[slot], 0.0); } assert_eq!(isv_before[RL_PPO_RATIO_CLAMP_MAX_INDEX], 10.0); + assert_eq!(isv_before[RL_ADV_VAR_RATIO_CLAMP_INDEX], 100.0); + assert_eq!(isv_before[RL_TD_KURTOSIS_CLAMP_INDEX], 30.0); // Populate each EMA-input slot with a non-zero value via the // R3 ema_update_per_step bootstrap path (sentinel-zero → first