fix(rl): v9 defensive eval-boundary calibration
Implements docs/superpowers/specs/2026-05-31-v9-defensive-eval-boundary-calibration.md.
Closes the [[pearl_adaptive_carryover_discipline]] gap: every adaptive
EMA must reset OR re-bootstrap at regime boundaries, never selectively
preserve train-acquired statistics.
Layer 1 — full EMA reset in `reset_session_state`:
- win_rate, avg_win/loss EMAs → neutral sentinels
- cumulative_dones → 0 (re-enter Kelly bootstrap)
- inventory_beta + inventory_variance EMAs → 0
- reward_clamp pos/neg max EMAs + clip_rate EMA → 0
- Kelly fraction → 1.0 (bootstrap)
Lifts Fix D's deliberate carryover (commit 7064c9269), which was
diagnosed in v8 fold-1 eval as the primary -$507k driver: agent
entered eval with train wr=0.34 EMA and took aggressively-sized
losing trades while EMA decayed to true eval wr=0.25.
Layer 2 — defensive warmup window:
- New `rl_eval_warmup_decay` CUDA kernel: single-thread single-block,
no atomics, mapped-pinned only. Runs after fused controllers each
step. Overrides 4 risk-sizing floors (Kelly safety_frac, IQN τ_min,
entropy_coef_min, PPO clip ε_min) with conservative defensive
values for the first 500 steps post-boundary, with linear decay
over the final 200 steps back to normal targets.
- 11 new ISV slots (685-695): remaining counter, configured
durations, defensive overrides (0.25/0.30/0.05/0.10), normal
targets (0.50/0.10/0.01/0.05).
- Sentinel counter = -1 at boot → kernel is no-op until reset arms it.
All boundary semantics live in ISV; no hardcoded constants in the
kernel. Build verified, all 13 risk-stack invariants + 5 controller
adaptive-floor tests + integrated-trainer smoke pass on RTX 3050.
This commit is contained in:
jgrusewski
@@ -111,6 +111,7 @@ const KERNELS: &[&str] = &[
|
||||
"rl_iqn_action_tau_controller", // Adaptive risk management (2026-05-30): Layer 2 — adapts IQN action-selection quantile τ from session drawdown signal
|
||||
"rl_inventory_beta_controller", // Adaptive risk management (2026-05-30): Layer 3 — adapts inventory penalty β from observed inventory variance vs reward magnitude
|
||||
"rl_kelly_fraction_controller", // Adaptive risk management (2026-05-30): Layer 4 — half-Kelly position-size multiplier from observed win_rate × R-multiple; warmup-gated
|
||||
"rl_eval_warmup_decay", // v9 (2026-05-31): defensive eval-boundary calibration — overrides Kelly/τ_min/entropy_min/ε_min during warmup, linearly decays back to normal; runs every step
|
||||
"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
|
||||
|
||||
105
crates/ml-alpha/cuda/rl_eval_warmup_decay.cu
Normal file
105
crates/ml-alpha/cuda/rl_eval_warmup_decay.cu
Normal file
@@ -0,0 +1,105 @@
|
||||
// rl_eval_warmup_decay.cu — defensive eval-boundary calibration (v9).
|
||||
//
|
||||
// Spec: docs/superpowers/specs/2026-05-31-v9-defensive-eval-boundary-calibration.md
|
||||
// Pearl: pearl_adaptive_carryover_discipline
|
||||
//
|
||||
// At every regime boundary (train→eval, fold transition), the trainer
|
||||
// calls `reset_session_state` which sets `RL_EVAL_WARMUP_REMAINING_INDEX`
|
||||
// to the configured warmup duration. This kernel runs once per step
|
||||
// thereafter; while the counter is positive, it overrides four
|
||||
// risk-sizing controller floors with conservative defensive values.
|
||||
//
|
||||
// Schedule (let R = warmup_remaining, W = warmup_steps, D = decay_steps):
|
||||
//
|
||||
// R > D full defensive overrides written
|
||||
// 0 < R ≤ D linear interpolation: defensive (1) → normal (0)
|
||||
// R = 0 normal values written once, counter goes to -1
|
||||
// R < 0 no-op (warmup completed, controllers run unimpeded)
|
||||
//
|
||||
// The override slots affected are read-only by their consuming controllers
|
||||
// (Kelly safety_frac, IQN τ_min, entropy_coef_min, PPO clip ε_min), so
|
||||
// last-writer-wins applies cleanly when this kernel runs AFTER the
|
||||
// controller batch.
|
||||
//
|
||||
// Per `feedback_no_atomicadd`: single-thread single-block, no atomics.
|
||||
// Per `feedback_cpu_is_read_only`: pure device kernel.
|
||||
// Per `feedback_isv_for_adaptive_bounds`: defensive override values, normal
|
||||
// target values, warmup/decay durations are all ISV-driven.
|
||||
|
||||
#include <stdint.h>
|
||||
|
||||
#define RL_EVAL_WARMUP_REMAINING_INDEX 685
|
||||
#define RL_EVAL_WARMUP_STEPS_CONFIG_INDEX 686
|
||||
#define RL_EVAL_WARMUP_DECAY_STEPS_CONFIG_INDEX 687
|
||||
#define RL_EVAL_KELLY_SAFETY_DEFENSIVE_INDEX 688
|
||||
#define RL_EVAL_IQN_TAU_MIN_DEFENSIVE_INDEX 689
|
||||
#define RL_EVAL_ENTROPY_COEF_MIN_DEFENSIVE_INDEX 690
|
||||
#define RL_EVAL_PPO_CLIP_EPS_MIN_DEFENSIVE_INDEX 691
|
||||
#define RL_EVAL_KELLY_SAFETY_NORMAL_INDEX 692
|
||||
#define RL_EVAL_IQN_TAU_MIN_NORMAL_INDEX 693
|
||||
#define RL_EVAL_ENTROPY_COEF_MIN_NORMAL_INDEX 694
|
||||
#define RL_EVAL_PPO_CLIP_EPS_MIN_NORMAL_INDEX 695
|
||||
|
||||
// Consumer slots (controller floors that this kernel overrides during warmup).
|
||||
#define RL_KELLY_SAFETY_FRAC_INDEX 680
|
||||
#define RL_IQN_ACTION_TAU_MIN_INDEX 672
|
||||
#define RL_ENTROPY_COEF_MIN_INDEX 645
|
||||
#define RL_PPO_CLIP_EPS_MIN_INDEX 640
|
||||
|
||||
extern "C" __global__ void rl_eval_warmup_decay(
|
||||
float* __restrict__ isv
|
||||
) {
|
||||
if (threadIdx.x != 0 || threadIdx.y != 0 || threadIdx.z != 0) return;
|
||||
if (blockIdx.x != 0 || blockIdx.y != 0 || blockIdx.z != 0) return;
|
||||
|
||||
const float remaining = isv[RL_EVAL_WARMUP_REMAINING_INDEX];
|
||||
|
||||
// Post-warmup steady state: do nothing. The normal floors stay at
|
||||
// their bootstrap values (or whatever the last-warmup-step wrote).
|
||||
if (remaining < 0.0f) {
|
||||
return;
|
||||
}
|
||||
|
||||
// Read configured durations + override / normal values.
|
||||
const float warmup_steps = isv[RL_EVAL_WARMUP_STEPS_CONFIG_INDEX];
|
||||
const float decay_steps = isv[RL_EVAL_WARMUP_DECAY_STEPS_CONFIG_INDEX];
|
||||
|
||||
const float kelly_def = isv[RL_EVAL_KELLY_SAFETY_DEFENSIVE_INDEX];
|
||||
const float tau_def = isv[RL_EVAL_IQN_TAU_MIN_DEFENSIVE_INDEX];
|
||||
const float entropy_def = isv[RL_EVAL_ENTROPY_COEF_MIN_DEFENSIVE_INDEX];
|
||||
const float ppo_def = isv[RL_EVAL_PPO_CLIP_EPS_MIN_DEFENSIVE_INDEX];
|
||||
|
||||
const float kelly_norm = isv[RL_EVAL_KELLY_SAFETY_NORMAL_INDEX];
|
||||
const float tau_norm = isv[RL_EVAL_IQN_TAU_MIN_NORMAL_INDEX];
|
||||
const float entropy_norm = isv[RL_EVAL_ENTROPY_COEF_MIN_NORMAL_INDEX];
|
||||
const float ppo_norm = isv[RL_EVAL_PPO_CLIP_EPS_MIN_NORMAL_INDEX];
|
||||
|
||||
// Phase selection:
|
||||
// pure warmup remaining > decay_steps → blend = 1.0 (full defensive)
|
||||
// decay phase 0 < remaining ≤ decay_steps → blend = remaining/decay_steps
|
||||
// final boundary remaining == 0 → blend = 0.0 (full normal)
|
||||
float blend;
|
||||
if (remaining > decay_steps) {
|
||||
blend = 1.0f;
|
||||
} else if (decay_steps > 0.0f) {
|
||||
blend = remaining / decay_steps; // 1.0 → 0.0 as remaining → 0
|
||||
} else {
|
||||
// Defensive: decay_steps configured to 0 means no decay phase.
|
||||
blend = (remaining > 0.0f) ? 1.0f : 0.0f;
|
||||
}
|
||||
|
||||
// Linear interpolation: blend × defensive + (1−blend) × normal.
|
||||
isv[RL_KELLY_SAFETY_FRAC_INDEX] = blend * kelly_def + (1.0f - blend) * kelly_norm;
|
||||
isv[RL_IQN_ACTION_TAU_MIN_INDEX] = blend * tau_def + (1.0f - blend) * tau_norm;
|
||||
isv[RL_ENTROPY_COEF_MIN_INDEX] = blend * entropy_def + (1.0f - blend) * entropy_norm;
|
||||
isv[RL_PPO_CLIP_EPS_MIN_INDEX] = blend * ppo_def + (1.0f - blend) * ppo_norm;
|
||||
|
||||
// Decrement counter for next step. When it crosses 0, the next step
|
||||
// sees remaining = -1 and returns early (controllers run unimpeded).
|
||||
isv[RL_EVAL_WARMUP_REMAINING_INDEX] = remaining - 1.0f;
|
||||
// Spec note: warmup_steps is used only by the trainer to set the
|
||||
// initial counter value; this kernel reads warmup_steps only to make
|
||||
// the `remaining > decay_steps` comparison meaningful for non-default
|
||||
// configurations (e.g., shorter warmup with same decay).
|
||||
(void) warmup_steps;
|
||||
}
|
||||
@@ -1460,6 +1460,48 @@ pub const RL_TRAIL_LOOSEN_FACTOR_INDEX: usize = 683;
|
||||
/// the [[feedback_no_partial_refactor]] per-batch fix.
|
||||
pub const RL_SESSION_PNL_WORST_INDEX: usize = 684;
|
||||
|
||||
/// Last RL-allocated slot index (exclusive). Pre-risk-stack: 662.
|
||||
/// Post-risk-stack (Fix B IQN-τ split): 685.
|
||||
pub const RL_SLOTS_END: usize = 685;
|
||||
// ============================================================
|
||||
// v9 — Defensive eval-boundary calibration (spec 2026-05-31-v9)
|
||||
//
|
||||
// Implements [[pearl_adaptive_carryover_discipline]]: at every regime
|
||||
// boundary (fold/eval transition), defensive overrides force conservative
|
||||
// behavior for the first N steps while EMAs re-bootstrap from new-regime
|
||||
// data. After N steps, override values linearly decay back to the normal
|
||||
// (bootstrap-configured) controller floors over `decay_steps` steps.
|
||||
// ============================================================
|
||||
|
||||
/// Counter — set to WARMUP_STEPS at every regime boundary, decremented each
|
||||
/// step. >0 = warmup/decay active; ≤0 = inactive (normal controller floors).
|
||||
pub const RL_EVAL_WARMUP_REMAINING_INDEX: usize = 685;
|
||||
/// Configured warmup duration (steps with full defensive overrides).
|
||||
/// Bootstrap 500 (~10% of a 5000-step eval phase).
|
||||
pub const RL_EVAL_WARMUP_STEPS_CONFIG_INDEX: usize = 686;
|
||||
/// Configured decay duration (steps over which override → normal interpolation).
|
||||
/// Bootstrap 200. Total transition: warmup_steps + decay_steps = 700 steps.
|
||||
pub const RL_EVAL_WARMUP_DECAY_STEPS_CONFIG_INDEX: usize = 687;
|
||||
|
||||
/// Defensive override: Kelly safety_frac. Bootstrap 0.25 (quarter-Kelly).
|
||||
/// Active during warmup; decays to NORMAL (slot 692) over decay phase.
|
||||
pub const RL_EVAL_KELLY_SAFETY_DEFENSIVE_INDEX: usize = 688;
|
||||
/// Defensive override: IQN τ_min. Bootstrap 0.30 (pessimistic action floor).
|
||||
/// Active during warmup; decays to NORMAL (slot 693) over decay phase.
|
||||
pub const RL_EVAL_IQN_TAU_MIN_DEFENSIVE_INDEX: usize = 689;
|
||||
/// Defensive override: entropy_coef_min. Bootstrap 0.05 (more exploration).
|
||||
/// Active during warmup; decays to NORMAL (slot 694) over decay phase.
|
||||
pub const RL_EVAL_ENTROPY_COEF_MIN_DEFENSIVE_INDEX: usize = 690;
|
||||
/// Defensive override: PPO clip ε_min. Bootstrap 0.10 (larger policy steps).
|
||||
/// Active during warmup; decays to NORMAL (slot 695) over decay phase.
|
||||
pub const RL_EVAL_PPO_CLIP_EPS_MIN_DEFENSIVE_INDEX: usize = 691;
|
||||
|
||||
/// Normal-value saved slots — read by warmup_decay kernel as the linear-
|
||||
/// interpolation target during the decay phase, and as the final restored
|
||||
/// value after warmup completes. Bootstrap to match the existing controller
|
||||
/// floor constants. These are NOT modified after bootstrap (read-only refs).
|
||||
pub const RL_EVAL_KELLY_SAFETY_NORMAL_INDEX: usize = 692; // 0.50
|
||||
pub const RL_EVAL_IQN_TAU_MIN_NORMAL_INDEX: usize = 693; // 0.10
|
||||
pub const RL_EVAL_ENTROPY_COEF_MIN_NORMAL_INDEX: usize = 694; // 0.01
|
||||
pub const RL_EVAL_PPO_CLIP_EPS_MIN_NORMAL_INDEX: usize = 695; // 0.05
|
||||
|
||||
/// Last RL-allocated slot index (exclusive).
|
||||
/// Pre-risk-stack: 662. Post-Fix-B: 685. Post-v9 (warmup boundary): 696.
|
||||
pub const RL_SLOTS_END: usize = 696;
|
||||
|
||||
@@ -214,6 +214,8 @@ const RL_INVENTORY_BETA_CONTROLLER_CUBIN: &[u8] =
|
||||
include_bytes!(concat!(env!("OUT_DIR"), "/rl_inventory_beta_controller.cubin"));
|
||||
const RL_KELLY_FRACTION_CONTROLLER_CUBIN: &[u8] =
|
||||
include_bytes!(concat!(env!("OUT_DIR"), "/rl_kelly_fraction_controller.cubin"));
|
||||
const RL_EVAL_WARMUP_DECAY_CUBIN: &[u8] =
|
||||
include_bytes!(concat!(env!("OUT_DIR"), "/rl_eval_warmup_decay.cubin"));
|
||||
const RL_WIN_RATE_EMA_UPDATE_CUBIN: &[u8] =
|
||||
include_bytes!(concat!(env!("OUT_DIR"), "/rl_win_rate_ema_update.cubin"));
|
||||
const RL_AVG_WIN_LOSS_EMA_UPDATE_CUBIN: &[u8] =
|
||||
@@ -678,6 +680,9 @@ pub struct IntegratedTrainer {
|
||||
rl_inventory_beta_controller_fn: CudaFunction,
|
||||
_rl_kelly_fraction_controller_module: Arc<CudaModule>,
|
||||
rl_kelly_fraction_controller_fn: CudaFunction,
|
||||
// v9 (2026-05-31): defensive eval-boundary calibration kernel.
|
||||
_rl_eval_warmup_decay_module: Arc<CudaModule>,
|
||||
rl_eval_warmup_decay_fn: CudaFunction,
|
||||
// EMA producers feeding the new controllers.
|
||||
_rl_win_rate_ema_update_module: Arc<CudaModule>,
|
||||
rl_win_rate_ema_update_fn: CudaFunction,
|
||||
@@ -1595,6 +1600,13 @@ impl IntegratedTrainer {
|
||||
let rl_kelly_fraction_controller_fn = rl_kelly_fraction_controller_module
|
||||
.load_function("rl_kelly_fraction_controller")
|
||||
.context("load rl_kelly_fraction_controller")?;
|
||||
// v9 — defensive eval-boundary calibration kernel.
|
||||
let rl_eval_warmup_decay_module = ctx
|
||||
.load_cubin(RL_EVAL_WARMUP_DECAY_CUBIN.to_vec())
|
||||
.context("load rl_eval_warmup_decay cubin")?;
|
||||
let rl_eval_warmup_decay_fn = rl_eval_warmup_decay_module
|
||||
.load_function("rl_eval_warmup_decay")
|
||||
.context("load rl_eval_warmup_decay")?;
|
||||
let rl_win_rate_ema_update_module = ctx
|
||||
.load_cubin(RL_WIN_RATE_EMA_UPDATE_CUBIN.to_vec())
|
||||
.context("load rl_win_rate_ema_update cubin")?;
|
||||
@@ -2671,6 +2683,8 @@ impl IntegratedTrainer {
|
||||
rl_inventory_beta_controller_fn,
|
||||
_rl_kelly_fraction_controller_module: rl_kelly_fraction_controller_module,
|
||||
rl_kelly_fraction_controller_fn,
|
||||
_rl_eval_warmup_decay_module: rl_eval_warmup_decay_module,
|
||||
rl_eval_warmup_decay_fn,
|
||||
_rl_win_rate_ema_update_module: rl_win_rate_ema_update_module,
|
||||
rl_win_rate_ema_update_fn,
|
||||
_rl_avg_win_loss_ema_update_module: rl_avg_win_loss_ema_update_module,
|
||||
@@ -2970,12 +2984,43 @@ impl IntegratedTrainer {
|
||||
/// inheriting all the dead/cooldown accounts from training.
|
||||
pub fn reset_session_state(&self) -> Result<()> {
|
||||
// Summary slots — direct ISV writes.
|
||||
//
|
||||
// v9 (2026-05-31, [[pearl_adaptive_carryover_discipline]]): reset
|
||||
// EVERY adaptive EMA driving downstream behavior, including the
|
||||
// ones Fix D previously preserved. The rationale for the prior
|
||||
// selective preservation ("training-acquired statistics should
|
||||
// carry into eval") proved wrong at the v8 fold-1 boundary:
|
||||
// Kelly entered eval with train wr=0.34 vs eval wr=0.25, took
|
||||
// overconfident positions, and lost $500k before EMAs could
|
||||
// re-converge. Pure adaptive principle now: every signal that
|
||||
// drives behavior must be re-bootstrapped from new-regime data.
|
||||
// Defensive warmup duration. Matches RL_EVAL_WARMUP_STEPS_CONFIG_INDEX
|
||||
// bootstrap (see ISV bootstrap block in `with_controllers_bootstrapped`).
|
||||
// Hardcoded here because reset_session_state runs before any ISV read
|
||||
// would observe a host-visible value; the kernel reads its own config.
|
||||
let warmup_steps = 500.0_f32;
|
||||
for (slot, value) in [
|
||||
// Layer 1 CMDP summary slots (kept from Fix D).
|
||||
(crate::rl::isv_slots::RL_SESSION_PNL_USD_INDEX, 0.0_f32),
|
||||
(crate::rl::isv_slots::RL_SESSION_PNL_WORST_INDEX, 0.0_f32),
|
||||
(crate::rl::isv_slots::RL_SESSION_DD_TRIGGERED_INDEX, 0.0_f32),
|
||||
(crate::rl::isv_slots::RL_CONSEC_LOSS_COUNT_INDEX, 0.0_f32),
|
||||
(crate::rl::isv_slots::RL_COOLDOWN_REMAINING_STEPS_INDEX, 0.0_f32),
|
||||
// Layer 4 (Kelly) EMAs — NEW in v9, was preserved by Fix D.
|
||||
(crate::rl::isv_slots::RL_WIN_RATE_EMA_INDEX, 0.5_f32), // neutral sentinel
|
||||
(crate::rl::isv_slots::RL_AVG_WIN_USD_EMA_INDEX, 0.0_f32), // sentinel
|
||||
(crate::rl::isv_slots::RL_AVG_LOSS_USD_EMA_INDEX, 0.0_f32), // sentinel
|
||||
(crate::rl::isv_slots::RL_CUMULATIVE_DONES_INDEX, 0.0_f32), // re-enter warmup
|
||||
(crate::rl::isv_slots::RL_KELLY_FRACTION_INDEX, 1.0_f32), // bootstrap
|
||||
// Layer 3 (Inventory) EMAs — NEW in v9.
|
||||
(crate::rl::isv_slots::RL_INVENTORY_PENALTY_BETA_INDEX, 0.0_f32), // sentinel
|
||||
(crate::rl::isv_slots::RL_INVENTORY_VARIANCE_EMA_INDEX, 0.0_f32), // sentinel
|
||||
// Reward-clamp EMAs — NEW in v9.
|
||||
(crate::rl::isv_slots::RL_POS_SCALED_REWARD_MAX_EMA_INDEX, 0.0_f32),
|
||||
(crate::rl::isv_slots::RL_NEG_SCALED_REWARD_MAX_EMA_INDEX, 0.0_f32),
|
||||
(crate::rl::isv_slots::RL_REWARD_CLAMP_CLIP_RATE_EMA_INDEX, 0.0_f32),
|
||||
// v9 — arm the defensive warmup window.
|
||||
(crate::rl::isv_slots::RL_EVAL_WARMUP_REMAINING_INDEX, warmup_steps),
|
||||
] {
|
||||
let mut args = RawArgs::new();
|
||||
args.push_ptr(self.isv_dev_ptr);
|
||||
@@ -3091,7 +3136,7 @@ impl IntegratedTrainer {
|
||||
// (slot, value) pair — pure device write, no HtoD per
|
||||
// `feedback_no_htod_htoh_only_mapped_pinned`.
|
||||
{
|
||||
let isv_constants: [(usize, f32); 168] = [
|
||||
let isv_constants: [(usize, f32); 179] = [
|
||||
// Static seeds for the adaptive reward-clamp controller —
|
||||
// these are the initial values that
|
||||
// `rl_reward_clamp_controller` will replace once it
|
||||
@@ -3369,6 +3414,23 @@ impl IntegratedTrainer {
|
||||
// Layer D (TrailTighten/Loosen factors).
|
||||
(crate::rl::isv_slots::RL_TRAIL_TIGHTEN_FACTOR_INDEX, 0.9),
|
||||
(crate::rl::isv_slots::RL_TRAIL_LOOSEN_FACTOR_INDEX, 1.1),
|
||||
// v9 (defensive eval-boundary calibration, [[pearl_adaptive_carryover_discipline]]):
|
||||
// 500-step warmup with 200-step linear decay phase. During warmup,
|
||||
// risk-sizing floors are temporarily overridden with defensive
|
||||
// values; reset_session_state arms the counter at each boundary.
|
||||
(crate::rl::isv_slots::RL_EVAL_WARMUP_REMAINING_INDEX, -1.0), // sentinel: no active warmup at boot
|
||||
(crate::rl::isv_slots::RL_EVAL_WARMUP_STEPS_CONFIG_INDEX, 500.0), // total warmup duration
|
||||
(crate::rl::isv_slots::RL_EVAL_WARMUP_DECAY_STEPS_CONFIG_INDEX, 200.0), // linear decay window
|
||||
// Defensive overrides (active for ~first 300 steps post-boundary).
|
||||
(crate::rl::isv_slots::RL_EVAL_KELLY_SAFETY_DEFENSIVE_INDEX, 0.25), // quarter-Kelly
|
||||
(crate::rl::isv_slots::RL_EVAL_IQN_TAU_MIN_DEFENSIVE_INDEX, 0.30), // raise pessimism floor
|
||||
(crate::rl::isv_slots::RL_EVAL_ENTROPY_COEF_MIN_DEFENSIVE_INDEX, 0.05), // more exploration
|
||||
(crate::rl::isv_slots::RL_EVAL_PPO_CLIP_EPS_MIN_DEFENSIVE_INDEX, 0.10), // wider trust region
|
||||
// Normal target values (restored post-decay; mirror Phase-3/4 bootstrap floors).
|
||||
(crate::rl::isv_slots::RL_EVAL_KELLY_SAFETY_NORMAL_INDEX, 0.50), // half-Kelly (Thorp)
|
||||
(crate::rl::isv_slots::RL_EVAL_IQN_TAU_MIN_NORMAL_INDEX, 0.10), // existing IQN floor
|
||||
(crate::rl::isv_slots::RL_EVAL_ENTROPY_COEF_MIN_NORMAL_INDEX, 0.01), // existing entropy floor
|
||||
(crate::rl::isv_slots::RL_EVAL_PPO_CLIP_EPS_MIN_NORMAL_INDEX, 0.05), // existing PPO ε floor
|
||||
];
|
||||
for (slot, value) in isv_constants.iter() {
|
||||
let slot_i32 = *slot as i32;
|
||||
@@ -3710,6 +3772,27 @@ impl IntegratedTrainer {
|
||||
Ok(())
|
||||
}
|
||||
|
||||
/// v9 — defensive eval-boundary calibration. Launch once per step
|
||||
/// AFTER all risk-stack controllers run (so this kernel's overrides
|
||||
/// of Kelly_safety_frac / IQN τ_min / entropy_coef_min / PPO ε_min
|
||||
/// take effect for the next step's consumers). Returns early if
|
||||
/// `RL_EVAL_WARMUP_REMAINING_INDEX` ≤ 0, so the runtime cost when
|
||||
/// warmup is inactive is just a single mapped-pinned ISV read.
|
||||
pub fn launch_rl_eval_warmup_decay(&self) -> Result<()> {
|
||||
let mut args = RawArgs::new();
|
||||
args.push_ptr(self.isv_dev_ptr);
|
||||
let mut ptrs = args.build_arg_ptrs();
|
||||
unsafe {
|
||||
raw_launch(
|
||||
self.rl_eval_warmup_decay_fn.cu_function(),
|
||||
(1, 1, 1), (1, 1, 1), 0,
|
||||
self.raw_stream,
|
||||
&mut ptrs[..args.len()],
|
||||
).map_err(|e| anyhow::anyhow!("rl_eval_warmup_decay: {:?}", e))?;
|
||||
}
|
||||
Ok(())
|
||||
}
|
||||
|
||||
/// Adaptive risk management — Layer 4 input EMA producer (win_rate).
|
||||
/// Reads `rewards_d`, `dones_d` [b_size] f32; outcome derived inline.
|
||||
/// Writes ISV[RL_WIN_RATE_EMA_INDEX = 677].
|
||||
@@ -6994,6 +7077,15 @@ impl IntegratedTrainer {
|
||||
self.launch_rl_fused_controllers(b_size)
|
||||
.context("launch_rl_fused_controllers")?;
|
||||
|
||||
// v9 defensive eval-boundary calibration ([[pearl_adaptive_carryover_discipline]]).
|
||||
// MUST run AFTER fused controllers: last-writer-wins on the four risk-sizing
|
||||
// floors (Kelly safety_frac, IQN τ_min, entropy_coef_min, PPO clip ε_min).
|
||||
// Steady-state (counter < 0): kernel returns early — zero behavioral effect.
|
||||
// After `reset_session_state` arms the counter, kernel overrides floors with
|
||||
// defensive defaults for warmup_steps, with a linear decay over decay_steps.
|
||||
self.launch_rl_eval_warmup_decay()
|
||||
.context("launch_rl_eval_warmup_decay")?;
|
||||
|
||||
// PopArt: normalize rewards in-place (replaces apply_reward_scale).
|
||||
// Welford-EMA on batch mean+variance → whitened rewards → ISV
|
||||
// stats for V-correct. Grid=(1), Block=(min(B, 256)). Shared mem
|
||||
|
||||
Reference in New Issue
Block a user