diff --git a/crates/ml-alpha/build.rs b/crates/ml-alpha/build.rs index 4a239f6b2..fb1653184 100644 --- a/crates/ml-alpha/build.rs +++ b/crates/ml-alpha/build.rs @@ -69,6 +69,7 @@ const KERNELS: &[&str] = &[ "rl_q_pi_agree_b", // RL R9 audit: per-batch fraction (argmax Q == argmax π) → ISV[407] EMA; wires previously-dead slot "rl_pi_action_kernel", // audit Option B: π drives action selection via multinomial sampling from softmax(pi_logits); Q becomes pure critic "rl_reward_clamp_controller", // audit 2026-05-24: adaptive [-LOSS, +WIN] clamp from positive-tail EMA; replaces static [-3, +1] that crushed winning-trade signal in rmgm5 + "rl_atom_support_update", // audit 2026-05-24 followup: refreshes atom_supports_d from ISV V_MIN/V_MAX so C51 atom span adapts with reward clamp (Q learning was capped at V_MAX=1.0) ]; // Cache bust v31 — five new reduce / derive kernels populate the input diff --git a/crates/ml-alpha/cuda/bellman_target_projection.cu b/crates/ml-alpha/cuda/bellman_target_projection.cu index d8a3b2433..497ada90f 100644 --- a/crates/ml-alpha/cuda/bellman_target_projection.cu +++ b/crates/ml-alpha/cuda/bellman_target_projection.cu @@ -42,9 +42,13 @@ #define Q_N_ATOMS 21 #define N_ACTIONS 9 -#define V_MIN -1.0f -#define V_MAX 1.0f -#define DELTA_Z ((V_MAX - V_MIN) / (Q_N_ATOMS - 1)) +// V_MIN/V_MAX/DELTA_Z are now ISV-driven per audit 2026-05-24 followup +// so adaptive reward clamps also lift Q's distributional support. +// Slots RL_C51_V_MIN_INDEX/RL_C51_V_MAX_INDEX are ratchet (monotone- +// grow), so the C51 atom mapping never coarsens — only expands as +// wider trades are observed. +#define RL_C51_V_MAX_INDEX 484 +#define RL_C51_V_MIN_INDEX 485 #define RL_GAMMA_INDEX 400 @@ -134,11 +138,19 @@ extern "C" __global__ void bellman_target_projection( const float r = rewards[batch]; const float gamma_eff = gamma * (1.0f - dones[batch]); + // ── Adaptive atom span from ISV (audit follow-up). DELTA_Z is + // recomputed every kernel launch — ratchet semantics in the + // controller guarantee V_MAX > V_MIN strictly (floored at + // [-1, +1] baseline) so this is always well-defined. + const float V_MIN_eff = isv[RL_C51_V_MIN_INDEX]; + const float V_MAX_eff = isv[RL_C51_V_MAX_INDEX]; + const float DELTA_Z = (V_MAX_eff - V_MIN_eff) / (float)(Q_N_ATOMS - 1); + // ── Per-source-atom target value + clamp + interpolation indices ─ - const float atom_value = V_MIN + (float)atom * DELTA_Z; + const float atom_value = V_MIN_eff + (float)atom * DELTA_Z; const float t_z = r + gamma_eff * atom_value; - const float t_z_clamp = fmaxf(V_MIN, fminf(V_MAX, t_z)); - const float b_frac = (t_z_clamp - V_MIN) / DELTA_Z; + const float t_z_clamp = fmaxf(V_MIN_eff, fminf(V_MAX_eff, t_z)); + const float b_frac = (t_z_clamp - V_MIN_eff) / DELTA_Z; const int l = max(0, min(Q_N_ATOMS - 1, (int)floorf(b_frac))); const int u = max(0, min(Q_N_ATOMS - 1, (int)ceilf(b_frac))); const float frac = b_frac - (float)l; diff --git a/crates/ml-alpha/cuda/rl_atom_support_update.cu b/crates/ml-alpha/cuda/rl_atom_support_update.cu new file mode 100644 index 000000000..fb7b03e22 --- /dev/null +++ b/crates/ml-alpha/cuda/rl_atom_support_update.cu @@ -0,0 +1,36 @@ +// rl_atom_support_update.cu — refresh `atom_supports_d` from the +// ISV-driven [V_MIN, V_MAX] span (audit 2026-05-24 second follow-up). +// +// Companion to the C51 atom-span ratchet in rl_reward_clamp_controller. +// `atom_supports_d` is the device buffer of 21 float values that the +// non-projection C51 kernels (`argmax_expected_q`, `rl_action_kernel`, +// `dqn_distributional_q`) read instead of recomputing the per-atom +// values themselves. When V_MIN/V_MAX adapt, this buffer must be +// rewritten or those kernels see a stale span. +// +// One block, Q_N_ATOMS threads. Each thread writes one atom value: +// atom_supports[i] = V_MIN + i * (V_MAX - V_MIN) / (N_ATOMS - 1) +// +// Per `feedback_no_atomicadd`: no atomics. Per `feedback_cpu_is_read_only`: +// pure device write. Per `feedback_no_htod_htoh_only_mapped_pinned`: no +// host transfer — values come from ISV slots, written from the trainer +// per-step launch right after the reward clamp controller refreshes +// V_MIN/V_MAX. + +#define Q_N_ATOMS 21 +#define RL_C51_V_MAX_INDEX 484 +#define RL_C51_V_MIN_INDEX 485 + +extern "C" __global__ void rl_atom_support_update( + const float* __restrict__ isv, // ≥ RL_C51_V_MAX_INDEX + 1 + float* __restrict__ atom_supports // [Q_N_ATOMS] +) { + const int tid = threadIdx.x; + if (tid >= Q_N_ATOMS) return; + + const float v_min = isv[RL_C51_V_MIN_INDEX]; + const float v_max = isv[RL_C51_V_MAX_INDEX]; + const float delta = (v_max - v_min) / (float)(Q_N_ATOMS - 1); + + atom_supports[tid] = v_min + (float)tid * delta; +} diff --git a/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu b/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu index 2a5cf0beb..6e466f656 100644 --- a/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu +++ b/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu @@ -68,6 +68,13 @@ // MARGIN adaptation slots (audit 2026-05-24 follow-up). #define RL_REWARD_CLAMP_CLIP_RATE_EMA_INDEX 482 #define RL_REWARD_CLAMP_CLIP_RATE_TARGET_INDEX 483 +// C51 atom span ratchet slots (audit 2026-05-24 second follow-up). +// Coupling: V_MAX tracks max(V_MAX_prev, max(1.0, WIN_clamp)); +// V_MIN tracks min(V_MIN_prev, min(-1.0, -LOSS_clamp)). Ratchet +// (monotone-grow) prevents destabilising Q's learned distribution +// when WIN_clamp temporarily shrinks. +#define RL_C51_V_MAX_INDEX 484 +#define RL_C51_V_MIN_INDEX 485 #define MIN_WIN 1.0f #define WIENER_ALPHA_FLOOR 0.4f @@ -172,4 +179,31 @@ extern "C" __global__ void rl_reward_clamp_controller( isv[RL_REWARD_CLAMP_WIN_INDEX] = win_eff; isv[RL_REWARD_CLAMP_LOSS_INDEX] = loss_eff; + + // ── C51 atom span ratchet (audit follow-up). ────────────────── + // + // V_MAX/V_MIN track the reward clamp with monotone-grow semantics + // — once we've admitted a winning trade of magnitude X, the C51 + // atom mapping permanently includes X in its support. This lets + // bellman_target_projection's clip pass through wins > 1.0 to + // distinct atoms instead of saturating the top atom. + // + // Floors: V_MAX ≥ 1.0, V_MIN ≤ -1.0 preserve the original C51 + // [-1, +1] support as a hard minimum (Q's distribution can't + // become coarser than the design baseline). + const float v_max_prev = isv[RL_C51_V_MAX_INDEX]; + const float v_min_prev = isv[RL_C51_V_MIN_INDEX]; + const float v_max_floor = 1.0f; + const float v_min_floor = -1.0f; + // Ratchet: target = max(floor, win); kept = max(prev, target). + const float v_max_target = fmaxf(v_max_floor, win_eff); + const float v_min_target = fminf(v_min_floor, -loss_eff); + // Bootstrap on sentinel 0 — first emit replaces directly with + // target (per pearl_first_observation_bootstrap). + const float v_max_new = (v_max_prev == 0.0f) ? v_max_target + : fmaxf(v_max_prev, v_max_target); + const float v_min_new = (v_min_prev == 0.0f) ? v_min_target + : fminf(v_min_prev, v_min_target); + isv[RL_C51_V_MAX_INDEX] = v_max_new; + isv[RL_C51_V_MIN_INDEX] = v_min_new; } diff --git a/crates/ml-alpha/examples/alpha_rl_train.rs b/crates/ml-alpha/examples/alpha_rl_train.rs index 2f427e752..4e3f65b11 100644 --- a/crates/ml-alpha/examples/alpha_rl_train.rs +++ b/crates/ml-alpha/examples/alpha_rl_train.rs @@ -67,6 +67,7 @@ use ml_alpha::rl::isv_slots::{ RL_POS_SCALED_REWARD_MAX_INDEX, RL_POS_SCALED_REWARD_MAX_EMA_INDEX, RL_REWARD_CLAMP_MARGIN_INDEX, RL_REWARD_CLAMP_RATIO_INDEX, RL_REWARD_CLAMP_CLIP_RATE_EMA_INDEX, RL_REWARD_CLAMP_CLIP_RATE_TARGET_INDEX, + RL_C51_V_MAX_INDEX, RL_C51_V_MIN_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, @@ -676,6 +677,13 @@ fn main() -> Result<()> { // or fighting the tail (rate > target × tolerance). "clip_rate_ema": isv[RL_REWARD_CLAMP_CLIP_RATE_EMA_INDEX], + // C51 atom span — ratcheted from WIN/LOSS clamp by the + // same controller. Watching these expand confirms Q's + // distributional learning ceiling is being lifted. + "c51_v_max": + isv[RL_C51_V_MAX_INDEX], + "c51_v_min": + isv[RL_C51_V_MIN_INDEX], }, // audit — PPO importance-ratio clamp diagnostics. // diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index f02c8d769..a492d60dd 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -40,8 +40,9 @@ //! | 468-477 | 10 design constants (Schulman + bootstraps + streaming α + kurt refs) | rl_isv_write | //! | 478-481 | 4 adaptive reward-clamp slots (raw + EMA + margin + ratio) | apply_reward_scale + rl_reward_clamp_controller | //! | 482-483 | 2 MARGIN-adaptation slots (clip-rate EMA + target) | rl_reward_clamp_controller (Schulman step on MARGIN) | +//! | 484-485 | 2 C51 atom-span ratchet slots (V_MAX + V_MIN) | rl_reward_clamp_controller (ratchet) + rl_atom_support_update (writes atom_supports_d) | //! -//! Total: 84 slots; `RL_SLOTS_END = 484`. +//! Total: 86 slots; `RL_SLOTS_END = 486`. /// Discount factor γ. Controller input: mean trade duration / anchor. /// Bootstrap 0.99. @@ -619,6 +620,32 @@ pub const RL_REWARD_CLAMP_CLIP_RATE_EMA_INDEX: usize = 482; /// (e.g. 0.10) accepts more clipping for tighter regression scale. pub const RL_REWARD_CLAMP_CLIP_RATE_TARGET_INDEX: usize = 483; +// ─── Adaptive C51 atom span (audit 2026-05-24 follow-up) ─────────── +// +// The reward clamp lift only unlocked V regression and PPO advantage; +// Q's distributional learning remained capped at the hardcoded +// bellman_target_projection's V_MIN=-1.0 / V_MAX=+1.0 — any Bellman +// target outside that range was categorically projected to atom 0 or +// atom 20 regardless of clamp. These slots make V_MIN/V_MAX +// per-step adaptive AND coupled to the reward clamp (V_MAX tracks +// WIN_clamp, V_MIN tracks -LOSS_clamp). Ratchet (monotone-grow) +// semantics: once a wider span is seen, the C51 atom mapping doesn't +// shrink — Q's learned distribution stays interpretable across +// training as "atom 20 represents at least the widest WIN we ever +// admitted." Floors: V_MAX ≥ 1.0 and V_MIN ≤ -1.0 preserve the +// original C51 [-1, +1] support as a hard minimum. + +/// C51 upper bound on the atom support. Ratchet — never shrinks below +/// max(1.0, current WIN_clamp). Written by `rl_reward_clamp_controller`; +/// consumed by `bellman_target_projection` (Bellman target clip) and +/// `rl_atom_support_update` (rebuilds `atom_supports_d` per step). +pub const RL_C51_V_MAX_INDEX: usize = 484; + +/// C51 lower bound on the atom support. Ratchet — never rises above +/// min(-1.0, -LOSS_clamp). Pair to `RL_C51_V_MAX_INDEX`; same +/// producers / consumers. +pub const RL_C51_V_MIN_INDEX: usize = 485; + /// 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 = 484; +pub const RL_SLOTS_END: usize = 486; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index 4880ff8eb..0592592ba 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -165,6 +165,14 @@ const APPLY_REWARD_SCALE_CUBIN: &[u8] = // kernel header for the design rationale (rmgm5 clamp-firing audit). const RL_REWARD_CLAMP_CONTROLLER_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_reward_clamp_controller.cubin")); +// C51 atom-support updater — audit 2026-05-24 followup. Refreshes +// `atom_supports_d` from the ISV V_MIN/V_MAX slots that the reward +// clamp controller writes (with ratchet semantics). Launched right +// after the reward clamp controller so atom_supports_d reflects the +// just-updated V_MIN/V_MAX before any C51 forward / Bellman target +// projection in the next step. 21-thread one-block kernel. +const RL_ATOM_SUPPORT_UPDATE_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_atom_support_update.cubin")); const ACTIONS_TO_MARKET_TARGETS_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/actions_to_market_targets.cubin")); @@ -383,6 +391,9 @@ pub struct IntegratedTrainer { // Adaptive reward-clamp controller (audit 2026-05-24). _rl_reward_clamp_controller_module: Arc, rl_reward_clamp_controller_fn: CudaFunction, + // C51 atom-support updater (audit 2026-05-24 followup). + _rl_atom_support_update_module: Arc, + rl_atom_support_update_fn: CudaFunction, _actions_to_market_targets_module: Arc, actions_to_market_targets_fn: CudaFunction, @@ -774,6 +785,12 @@ impl IntegratedTrainer { let rl_reward_clamp_controller_fn = rl_reward_clamp_controller_module .load_function("rl_reward_clamp_controller") .context("load rl_reward_clamp_controller")?; + let rl_atom_support_update_module = ctx + .load_cubin(RL_ATOM_SUPPORT_UPDATE_CUBIN.to_vec()) + .context("load rl_atom_support_update cubin")?; + let rl_atom_support_update_fn = rl_atom_support_update_module + .load_function("rl_atom_support_update") + .context("load rl_atom_support_update")?; let actions_to_market_targets_module = ctx .load_cubin(ACTIONS_TO_MARKET_TARGETS_CUBIN.to_vec()) .context("load actions_to_market_targets cubin")?; @@ -1082,6 +1099,8 @@ impl IntegratedTrainer { apply_reward_scale_fn, _rl_reward_clamp_controller_module: rl_reward_clamp_controller_module, rl_reward_clamp_controller_fn, + _rl_atom_support_update_module: rl_atom_support_update_module, + rl_atom_support_update_fn, _actions_to_market_targets_module: actions_to_market_targets_module, actions_to_market_targets_fn, _abs_copy_module: abs_copy_module, @@ -1199,7 +1218,7 @@ impl IntegratedTrainer { // (slot, value) pair — pure device write, no HtoD per // `feedback_no_htod_htoh_only_mapped_pinned`. { - let isv_constants: [(usize, f32); 29] = [ + let isv_constants: [(usize, f32); 31] = [ // Static seeds for the adaptive reward-clamp controller — // these are the initial values that // `rl_reward_clamp_controller` will replace once it @@ -1209,12 +1228,16 @@ impl IntegratedTrainer { // the controller's Schulman bounded-step from clip-rate // EMA (482) vs target (483). CLIP_RATE_TARGET seeded to // 0.05 = "let 5% of wins ride the clamp as tail - // outliers." + // outliers." V_MAX/V_MIN seeded to the original C51 + // [-1, +1] span as the canonical floor — the ratchet + // in `rl_reward_clamp_controller` only grows magnitude. (crate::rl::isv_slots::RL_REWARD_CLAMP_WIN_INDEX, 1.0), (crate::rl::isv_slots::RL_REWARD_CLAMP_LOSS_INDEX, 3.0), (crate::rl::isv_slots::RL_REWARD_CLAMP_MARGIN_INDEX, 1.5), (crate::rl::isv_slots::RL_REWARD_CLAMP_RATIO_INDEX, 3.0), (crate::rl::isv_slots::RL_REWARD_CLAMP_CLIP_RATE_TARGET_INDEX, 0.05), + (crate::rl::isv_slots::RL_C51_V_MAX_INDEX, 1.0), + (crate::rl::isv_slots::RL_C51_V_MIN_INDEX, -1.0), (crate::rl::isv_slots::RL_KL_TARGET_INDEX, 0.01), (crate::rl::isv_slots::RL_IMPROVEMENT_THRESHOLD_INDEX, 0.99), (crate::rl::isv_slots::RL_PLATEAU_PATIENCE_INDEX, 1000.0), @@ -1740,6 +1763,28 @@ impl IntegratedTrainer { .context("rl_reward_clamp_controller launch")?; } } + + // C51 atom support refresh — the reward clamp controller just + // ratcheted V_MIN/V_MAX; rewrite atom_supports_d so downstream + // C51 kernels (argmax_expected_q, rl_action_kernel, + // dqn_distributional_q) see the updated span. 21-thread + // single-block kernel; trivially cheap. + { + let cfg_atom = LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (21, 1, 1), + shared_mem_bytes: 0, + }; + let mut launch = self + .stream + .launch_builder(&self.rl_atom_support_update_fn); + launch.arg(&self.isv_d).arg(&self.atom_supports_d); + unsafe { + launch + .launch(cfg_atom) + .context("rl_atom_support_update launch")?; + } + } Ok(()) } @@ -3068,6 +3113,22 @@ impl IntegratedTrainer { .launch(cfg_ctrl) .context("rl_reward_clamp_controller launch (step_with_lobsim path)")?; } + + // C51 atom support refresh from updated V_MIN/V_MAX. + let cfg_atom = LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (21, 1, 1), + shared_mem_bytes: 0, + }; + let mut launch = self + .stream + .launch_builder(&self.rl_atom_support_update_fn); + launch.arg(&self.isv_d).arg(&self.atom_supports_d); + unsafe { + launch + .launch(cfg_atom) + .context("rl_atom_support_update launch (step_with_lobsim path)")?; + } } // GPU compute_advantage_return: returns_d, advantages_d