gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

fix(rl): ISV-driven K-loop divisor + max ceiling (slot 450, 451)

Browse Source 
f2ggr confirmed K-loop wiring works mechanically but K=8 firing on
22 % of steps over-trained at b_size=1: KL excursions to 12.44
(vs prior 3.4e-4), policy overshoot, reward/trade -$0.585 → -$0.723.

Per `feedback_isv_for_adaptive_bounds` the K-loop config must live
in ISV, not as hardcoded values in the trainer. Two new slots:

  RL_K_LOOP_DIVISOR_INDEX (450) — divides n_rollout_steps to get K
                                  Default 2048 (matches ROLLOUT_BOOTSTRAP
                                  so K=1 at controller bootstrap)
  RL_K_LOOP_MAX_INDEX     (451) — clamp ceiling on K
                                  Default 4 (was hardcoded 8; halved
                                  to prevent gradient overtraining)

K computation in step_with_lobsim now reads both from ISV:
  K = clamp(isv[404] / isv[450], 1, isv[451])

Halves worst-case overtraining while preserving the controller
cascade activation (KL above noise floor, ε actively adapting,
ratio_clamp firing). Distribution shifts from K=8 @ 22% → K=4 @ 22%
(half the gradient updates in the high-noise case).

## Wiring

`rl_streaming_clamp_init.cu` extended to seed 5 ISV-resident design
constants (was 3): adv_var_clamp, td_kurt_clamp, adv_var_target,
k_loop_divisor, k_loop_max. Still one kernel call, no HtoD.

## Diag bake-in

JSONL `k_updates` field replaced with `k_loop` block:
  k_loop.k_updates  — actual K used this step
  k_loop.divisor    — current divisor (reads isv[450])
  k_loop.max        — current max (reads isv[451])

Post-hoc analysis can verify the K-computation by independently
recomputing K from isv[404] / k_loop.divisor.

## Slot allocation

RL_SLOTS_END: 450 → 452 (+2 new config slots).

## Test updates

G1 + G3 skip slots 450, 451 in sentinel-zero loop and assert seeded
values (2048.0 + 4.0) separately.

## Verified gates (local sm_86)

  G1 isv_bootstrap   ✅
  G3 controllers     ✅
  G4 target_update   ✅
  integrated_smoke   ✅

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-24 00:23:38 +02:00
parent 1d8ef94848
commit 644fbe0348
6 changed files with 113 additions and 22 deletions
Download Patch File Download Diff File Expand all files Collapse all files

12
crates/ml-alpha/cuda/rl_streaming_clamp_init.cu
View File

@@ -1,9 +1,11 @@
// rl_streaming_clamp_init.cu — single-thread device-side seeder for
// the streaming-kernel output clamp ceilings AND the
// rl_rollout_steps_controller regression target:
// the streaming-kernel output clamp ceilings, the rollout-controller
// regression target, AND the K-loop training-intensification config:
// ISV[RL_ADV_VAR_RATIO_CLAMP_INDEX = 447] — var/|mean| clamp
// ISV[RL_TD_KURTOSIS_CLAMP_INDEX = 448] — kurtosis clamp
// ISV[RL_ADV_VAR_RATIO_TARGET_INDEX = 449] — rollout-controller target
// ISV[RL_K_LOOP_DIVISOR_INDEX = 450] — K = n_rollout / divisor
// ISV[RL_K_LOOP_MAX_INDEX = 451] — K clamped to this ceiling
//
// Launched once from `with_controllers_bootstrapped` after ISV is
// alloc_zeros'd. Writes design-constant values directly on device —
@@ -32,10 +34,14 @@ 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,
int adv_var_target_slot, float adv_var_target_val
int adv_var_target_slot, float adv_var_target_val,
int k_loop_divisor_slot, float k_loop_divisor_val,
int k_loop_max_slot, float k_loop_max_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;
isv[adv_var_target_slot] = adv_var_target_val;
isv[k_loop_divisor_slot] = k_loop_divisor_val;
isv[k_loop_max_slot] = k_loop_max_val;
}

19
crates/ml-alpha/examples/alpha_rl_train.rs
View File

@@ -46,7 +46,7 @@ use ml_alpha::rl::isv_slots::{
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_RATIO_TARGET_INDEX, RL_ADV_VAR_STREAM_M2_INDEX,
RL_ADV_VAR_STREAM_MEAN_INDEX,
RL_ADV_VAR_STREAM_MEAN_INDEX, RL_K_LOOP_DIVISOR_INDEX, RL_K_LOOP_MAX_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,
@@ -626,12 +626,17 @@ fn main() -> Result<()> {
// design constants because they're fundamental to the
// controller's behaviour, not adaptive).
// R9 — K-loop training-intensification iterations performed
// this step. K = clamp(isv[404]/1024, 1, 8). Surfaces the
// wire-up of n_rollout_steps as the DQN-replay/PPO multi-
// update knob: when advantages are noisy, do K Q+π+V
// updates per env step (each with fresh PER sample) so
// gradient signal isn't starved at b_size=1.
"k_updates": trainer.last_k_updates,
// this step. K = clamp(isv[404] / isv[450], 1, isv[451]).
// Both divisor + max are ISV-driven (per
// feedback_isv_for_adaptive_bounds) — surfaced here so
// post-hoc analysis can compute K from any sample without
// recompile. Defaults: divisor=2048 (K=1 at controller
// bootstrap), max=4 (prevents gradient overtraining).
"k_loop": {
"k_updates": trainer.last_k_updates,
"divisor": isv[RL_K_LOOP_DIVISOR_INDEX],
"max": isv[RL_K_LOOP_MAX_INDEX],
},
"controller_branch": {
"rollout_steps_input": isv[RL_ADVANTAGE_VAR_RATIO_EMA_INDEX],
// ISV-driven target — seeded by rl_streaming_clamp_init,

34
crates/ml-alpha/src/rl/isv_slots.rs
View File

@@ -332,6 +332,38 @@ pub const RL_TD_KURTOSIS_CLAMP_INDEX: usize = 448;
/// target into a true feedback loop with the signal.
pub const RL_ADV_VAR_RATIO_TARGET_INDEX: usize = 449;
/// K-loop divisor for the n_rollout_steps training-intensification
/// wire-up. The trainer computes
/// `K = clamp(isv[RL_N_ROLLOUT_STEPS_INDEX] / isv[RL_K_LOOP_DIVISOR_INDEX],
/// 1, isv[RL_K_LOOP_MAX_INDEX])`
/// to determine how many full (PER sample + step_synthetic + PER
/// priority) iterations to run per env step.
///
/// Default 2048.0 — matches `ROLLOUT_BOOTSTRAP` so at controller
/// bootstrap (before any adaptation), `K = 2048/2048 = 1` and the
/// loop behaves like a single per-step update (current baseline).
/// As n_rollout grows toward MAX=8192, K rises to K_MAX.
///
/// Raised from the prior hardcoded `1024.0` after alpha-rl-f2ggr
/// fold0 (commit 1d8ef9484) confirmed K=8 firing on 22 % of steps
/// caused gradient overtraining: KL excursions to 12.44 (vs prior
/// 3.4e-4), policy overshoot, reward/trade -$0.585 → -$0.723.
/// 2048 halves the worst-case K to 4 while preserving the controller
/// cascade activation (KL above noise floor, ε actively adapting).
///
/// Per `feedback_isv_for_adaptive_bounds`: lives in ISV, discoverable
/// in diag, tunable at runtime via re-launching rl_streaming_clamp_init.
pub const RL_K_LOOP_DIVISOR_INDEX: usize = 450;
/// K-loop maximum ceiling. K is clamped to this value regardless of
/// how high n_rollout_steps grows. Default 4.0 — prevents
/// gradient overtraining at b_size=1 even if n_rollout_steps pegs
/// at MAX=8192 (would otherwise give K=4 with divisor=2048, but
/// kept explicit as defense-in-depth).
///
/// Per `feedback_isv_for_adaptive_bounds`: lives in ISV.
pub const RL_K_LOOP_MAX_INDEX: usize = 451;
/// 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 = 450;
pub const RL_SLOTS_END: usize = 452;

36
crates/ml-alpha/src/trainer/integrated.rs
View File

@@ -1209,6 +1209,19 @@ impl IntegratedTrainer {
let adv_var_target_slot: i32 =
crate::rl::isv_slots::RL_ADV_VAR_RATIO_TARGET_INDEX as i32;
let adv_var_target_val: f32 = 5.0;
// K-loop config — also ISV-driven per
// `feedback_isv_for_adaptive_bounds`. Divisor 2048 matches
// ROLLOUT_BOOTSTRAP (so K=1 at controller bootstrap, before
// any adaptation). K_MAX=4 prevents gradient overtraining
// at b_size=1; raised from default 8 after alpha-rl-f2ggr
// confirmed K=8 firing on 22% of steps over-trained
// (KL excursions to 12.44, reward/trade -$0.585 → -$0.723).
let k_loop_divisor_slot: i32 =
crate::rl::isv_slots::RL_K_LOOP_DIVISOR_INDEX as i32;
let k_loop_divisor_val: f32 = 2048.0;
let k_loop_max_slot: i32 =
crate::rl::isv_slots::RL_K_LOOP_MAX_INDEX as i32;
let k_loop_max_val: f32 = 4.0;
let cfg = LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (1, 1, 1),
@@ -1224,7 +1237,11 @@ impl IntegratedTrainer {
.arg(&td_kurt_slot)
.arg(&td_kurt_val)
.arg(&adv_var_target_slot)
.arg(&adv_var_target_val);
.arg(&adv_var_target_val)
.arg(&k_loop_divisor_slot)
.arg(&k_loop_divisor_val)
.arg(&k_loop_max_slot)
.arg(&k_loop_max_val);
unsafe {
launch
.launch(cfg)
@@ -2767,10 +2784,23 @@ impl IntegratedTrainer {
// addresses the per-step gradient signal starvation that left
// l_q stuck at 2.82 (uniform=3.04, only 7% information gain)
// in alpha-rl-frt7s.
// K-loop config — both divisor and max are ISV-driven per
// `feedback_isv_for_adaptive_bounds` so they're discoverable
// in diag and tunable at runtime by re-launching
// rl_streaming_clamp_init. Defaults seeded at trainer init:
// divisor=2048 (matches ROLLOUT_BOOTSTRAP so K=1 at controller
// bootstrap), max=4 (prevents gradient overtraining at
// b_size=1).
let n_rollout_steps =
self.isv_host[crate::rl::isv_slots::RL_N_ROLLOUT_STEPS_INDEX];
let k_updates = ((n_rollout_steps / 1024.0).round() as usize)
.clamp(1, 8);
let k_divisor = self
.isv_host[crate::rl::isv_slots::RL_K_LOOP_DIVISOR_INDEX]
.max(1.0);
let k_max = self
.isv_host[crate::rl::isv_slots::RL_K_LOOP_MAX_INDEX]
.max(1.0) as usize;
let k_updates = ((n_rollout_steps / k_divisor).round() as usize)
.clamp(1, k_max);
self.last_k_updates = k_updates;
// Track the LAST iteration's stats for caller reporting (per-

20
crates/ml-alpha/tests/isv_bootstrap.rs
View File

@@ -24,9 +24,10 @@
use ml_alpha::rl::isv_slots::{
RL_ADV_VAR_RATIO_CLAMP_INDEX, RL_ADV_VAR_RATIO_TARGET_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,
RL_GAMMA_INDEX, RL_K_LOOP_DIVISOR_INDEX, RL_K_LOOP_MAX_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;
@@ -150,6 +151,8 @@ fn g1_isv_bootstrap_writes_canonical_values() {
|| slot == RL_ADV_VAR_RATIO_CLAMP_INDEX
|| slot == RL_TD_KURTOSIS_CLAMP_INDEX
|| slot == RL_ADV_VAR_RATIO_TARGET_INDEX
|| slot == RL_K_LOOP_DIVISOR_INDEX
|| slot == RL_K_LOOP_MAX_INDEX
{
continue;
}
@@ -184,6 +187,17 @@ fn g1_isv_bootstrap_writes_canonical_values() {
"ISV[adv_var_ratio_target={RL_ADV_VAR_RATIO_TARGET_INDEX}] expected 5.0, got {}",
isv[RL_ADV_VAR_RATIO_TARGET_INDEX]
);
// R9 — K-loop config (divisor + max), seeded by rl_streaming_clamp_init.
assert!(
(isv[RL_K_LOOP_DIVISOR_INDEX] - 2048.0).abs() < EPS,
"ISV[k_loop_divisor={RL_K_LOOP_DIVISOR_INDEX}] expected 2048.0, got {}",
isv[RL_K_LOOP_DIVISOR_INDEX]
);
assert!(
(isv[RL_K_LOOP_MAX_INDEX] - 4.0).abs() < EPS,
"ISV[k_loop_max={RL_K_LOOP_MAX_INDEX}] expected 4.0, got {}",
isv[RL_K_LOOP_MAX_INDEX]
);
eprintln!(
"G1 OK — bootstraps: γ={:.4} τ={:.4} ε={:.4} entropy_coef={:.4} \

14
crates/ml-alpha/tests/r5_controllers_and_soft_update.rs
View File

@@ -28,11 +28,11 @@ use cudarc::driver::CudaStream;
use ml_alpha::rl::isv_slots::{
RL_ADVANTAGE_VAR_RATIO_EMA_INDEX, RL_ADV_VAR_RATIO_CLAMP_INDEX,
RL_ADV_VAR_RATIO_TARGET_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,
RL_GAMMA_INDEX, RL_KL_PI_EMA_INDEX, RL_K_LOOP_DIVISOR_INDEX, RL_K_LOOP_MAX_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;
@@ -121,6 +121,8 @@ fn g3_per_step_controllers_move_isv_outputs_when_fed_real_emas() {
|| slot == RL_ADV_VAR_RATIO_CLAMP_INDEX
|| slot == RL_TD_KURTOSIS_CLAMP_INDEX
|| slot == RL_ADV_VAR_RATIO_TARGET_INDEX
|| slot == RL_K_LOOP_DIVISOR_INDEX
|| slot == RL_K_LOOP_MAX_INDEX
{
continue;
}
@@ -130,6 +132,8 @@ fn g3_per_step_controllers_move_isv_outputs_when_fed_real_emas() {
assert_eq!(isv_before[RL_ADV_VAR_RATIO_CLAMP_INDEX], 100.0);
assert_eq!(isv_before[RL_TD_KURTOSIS_CLAMP_INDEX], 30.0);
assert_eq!(isv_before[RL_ADV_VAR_RATIO_TARGET_INDEX], 5.0);
assert_eq!(isv_before[RL_K_LOOP_DIVISOR_INDEX], 2048.0);
assert_eq!(isv_before[RL_K_LOOP_MAX_INDEX], 4.0);
// Populate each EMA-input slot with a non-zero value via the
// R3 ema_update_per_step bootstrap path (sentinel-zero → first