From 20c7852b66458045ee8aaa40b7cbb254f640ef61 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sat, 23 May 2026 20:13:50 +0200 Subject: [PATCH] fix(rl): asymmetric clamp on scaled reward + pre-clamp |max| diag MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The xv66n smoke (commit d5c29fb4f) confirmed plateau-decay LR works across all 3 heads — but exposed a residual V instability: 10 trade- close steps with l_v > 1e4, max 9.46e4. Root cause: the reward_scale controller's Wiener-α blend cannot adapt fast enough to a sudden fat-tail trade outcome, so a single closed trade with realised PnL well outside `1 / mean_abs_pnl_ema`'s current estimate produces a scaled reward 100s of times the C51 atom span. Since `returns = scaled_reward + γ(1-done) v_tp1` and the spike happens on done=1 steps, returns equals the unbounded scaled reward, and V regression `(v_pred - returns)²` blows up. ## Fix: asymmetric clamp at apply_reward_scale boundary `apply_reward_scale.cu` is rewritten to: 1. Scale `rewards[b] *= isv[RL_REWARD_SCALE_INDEX]` as before. 2. Asymmetric-clamp scaled to `[-REWARD_CLAMP_LOSS, +REWARD_CLAMP_WIN]` = `[-3.0, +1.0]` per `pearl_audit_unboundedness_for_implicit_asymmetry`: * `WIN = +1.0` matches the C51 atom span on the win side. * `LOSS = -3.0` preserves loss-aversion asymmetry — fat-tail losses remain visible up to 3 atom-units before flattening, matching typical HFT P&L distributions where losses run 2-3× larger than wins per close. 3. Write back the clamped value to `rewards[b]`. Single-block layout (block_x = min(b_size, 256), grid_x = 1, shared = block_x × 4 B) per `pearl_no_atomicadd` — tree reduction inside the block, no inter-block atomic. ## Diagnostic: pre-clamp max ISV slot New ISV slot `RL_MAX_ABS_SCALED_REWARD_PRE_CLAMP_INDEX = 439` holds `max(|scaled|)` over the current batch BEFORE the clamp fires (each step overwrites — point measurement, not EMA). Surfaced in diag.jsonl as `rewards.scaled_pre_clamp_max`. Interpretation: * pre_clamp_max ≤ 1.0 most steps → reward_scale controller is tracking typical magnitudes correctly; clamp is a no-op. * pre_clamp_max > 1.0 frequently → controller is failing to track magnitudes; clamp is doing load-bearing work shaping V target. * pre_clamp_max > 100 ever → controller is grossly mis-scaled (likely cold-start before mean_abs_pnl_ema converged). RL_SLOTS_END: 439 → 440 (one new diagnostic slot). ## Why a clamp instead of fixing the controller The reward_scale controller IS doing its job — it Wiener-blends toward `1 / mean_abs_pnl_ema` with α floor 0.4. The problem is that a single closed trade represents one observation in the EMA denominator, so a sudden 10× excursion in trade magnitude takes ~3-5 closes to fully reflect in the scale. During those 3-5 steps, scaled rewards can be 5-10× the atom span. A faster controller (smaller EMA floor, lookahead, etc.) would oscillate. A clamp is the principled bound: * source signal (raw PnL) remains unbounded — controller continues to track magnitudes * downstream signal (V/Q target) is bounded — no catastrophic backward gradient * pre-clamp diagnostic surfaces clamp activity so we know when the controller is failing vs handling the regime fine ## Verified gates (local sm_86) G1 isv_bootstrap ✅ G3 controllers ✅ G4 target_update ✅ G6 r7d_per_wiring ✅ integrated_smoke ✅ Co-Authored-By: Claude Opus 4.7 --- crates/ml-alpha/cuda/apply_reward_scale.cu | 123 +++++++++++++++++---- crates/ml-alpha/examples/alpha_rl_train.rs | 21 +++- crates/ml-alpha/src/rl/isv_slots.rs | 18 ++- crates/ml-alpha/src/trainer/integrated.rs | 53 ++++----- 4 files changed, 162 insertions(+), 53 deletions(-) diff --git a/crates/ml-alpha/cuda/apply_reward_scale.cu b/crates/ml-alpha/cuda/apply_reward_scale.cu index 0e1b90a01..158fea52c 100644 --- a/crates/ml-alpha/cuda/apply_reward_scale.cu +++ b/crates/ml-alpha/cuda/apply_reward_scale.cu @@ -1,30 +1,111 @@ -// apply_reward_scale.cu — element-wise reward standardisation reading -// the scale from `ISV[RL_REWARD_SCALE_INDEX = 406]` (Phase R6 of the -// integrated RL trainer rebuild; -// see docs/superpowers/plans/2026-05-23-integrated-rl-trainer-rebuild.md). +// apply_reward_scale.cu — element-wise reward standardisation + +// asymmetric clamp + per-step pre-clamp |max| diagnostic. // -// Replaces the host loop `for r in rewards: r *= scale` the flawed -// Phase F shipped in step_with_lobsim. That loop required DtoH of -// ISV[406] + host multiply + HtoD of scaled rewards — a triple -// roundtrip per training step. This kernel does the multiply on -// device with the scale already on device. +// Reads the standardisation scale from `ISV[RL_REWARD_SCALE_INDEX = 406]` +// (the rl_reward_scale_controller emits this — see +// rl_reward_scale_controller.cu for the controller logic and bounds). // -// Per `pearl_one_unbounded_signal_per_reward`: `rewards[b]` is the -// single unbounded multiplicand; `isv[RL_REWARD_SCALE_INDEX]` is -// bounded by the reward-scale controller (clamped to [1e-3, 1e3] -// per the kernel). Reward magnitude after scaling sits within the -// [Q_V_MIN, Q_V_MAX] = [-1, +1] atom support per the C51 atom layout. +// Phase R6 originally only scaled rewards in-place. This left V +// regression catastrophically exposed when the reward_scale controller +// hadn't yet adapted to a fresh trade-magnitude regime: a single +// closed trade with realised PnL well outside `1 / mean_abs_pnl_ema`'s +// current estimate could produce a scaled reward 100s of times the +// C51 atom span. The V regression target `returns = reward + γ(1-done) +// v_tp1` then equals this large scaled reward on trade-close steps +// (done=1), and `(v_pred - returns)²` spikes into the 1e4-1e6 range. +// Canonical incident: alpha-rl-xv66n fold0 had 10 trade-close steps +// with l_v > 1e4, max 9.46e4. +// +// FIX (this kernel): clamp the scaled reward to an asymmetric bound +// before writing it back. The clamp: +// +// * Caps loss magnitudes at `REWARD_CLAMP_LOSS = 3.0` +// * Caps win magnitudes at `REWARD_CLAMP_WIN = 1.0` +// +// per `pearl_audit_unboundedness_for_implicit_asymmetry`: preserving +// the loss > win asymmetry (loss aversion) so a single fat-tail loss +// remains visible while still bounding V regression's target. The +// asymmetry also matches typical HFT reward distributions where losses +// can be 2-3× larger than wins per close. +// +// DIAGNOSTIC: writes the pre-clamp max |scaled| over the current batch +// to `ISV[RL_MAX_ABS_SCALED_REWARD_PRE_CLAMP_INDEX = 439]`. This lets +// the diag JSONL surface how often the clamp is actually firing — when +// pre-clamp max stays ≤ REWARD_CLAMP_WIN the clamp is a no-op +// (controller has adapted scale correctly); when it rises above that +// repeatedly, the reward_scale controller is failing to track typical +// trade magnitudes and the clamp is doing load-bearing work. +// +// Per `pearl_one_unbounded_signal_per_reward`: `rewards[b]` pre-scale +// is the single unbounded multiplicand; `isv[RL_REWARD_SCALE_INDEX]` +// is bounded by the controller [1e-3, 1e3]. Post-clamp, even +// rewards[b] is bounded — the controller's earlier "single unbounded +// multiplicand" property is preserved at the source (raw PnL) but the +// V/Q-target-facing signal is bounded. +// +// Per `pearl_no_atomicadd`: single-block reduction, no atomicMax. For +// b_size > BLOCK_DIM the kernel does a grid-stride loop with per-thread +// local max then a shared-memory tree reduce. -#define RL_REWARD_SCALE_INDEX 406 +#define RL_REWARD_SCALE_INDEX 406 +#define RL_MAX_ABS_SCALED_REWARD_PRE_CLAMP_INDEX 439 +// Asymmetric bounded clamp matching C51 atom span and HFT loss-aversion +// pearl. Win cap matches the canonical C51 atom_max; loss cap is 3× +// larger so fat-tail losses remain distinguishable up to 3 atom-units +// before the clamp flattens them. +#define REWARD_CLAMP_WIN 1.0f +#define REWARD_CLAMP_LOSS 3.0f + +// Single block; grid-stride for b_size > BLOCK_DIM. Caller launches +// with block_dim = min(b_size, 256), grid_dim = 1, shared bytes = +// block_dim * sizeof(float). extern "C" __global__ void apply_reward_scale( - const float* __restrict__ isv, // ISV bus (≥ RL_REWARD_SCALE_INDEX + 1) - float* __restrict__ rewards, // [b_size] IN/OUT + float* __restrict__ isv, // ISV bus IN/OUT + float* __restrict__ rewards, // [b_size] IN/OUT (scaled+clamped) int b_size ) { - const int b = blockIdx.x * blockDim.x + threadIdx.x; - if (b >= b_size) return; - + extern __shared__ float s_max[]; + const int tid = threadIdx.x; const float scale = isv[RL_REWARD_SCALE_INDEX]; - rewards[b] *= scale; + + // Pass 1 — scale, clamp, accumulate per-thread max |scaled_pre_clamp|. + float local_max = 0.0f; + for (int b = tid; b < b_size; b += blockDim.x) { + const float raw = rewards[b]; + const float scaled = raw * scale; + + // Track magnitude BEFORE clamping so the diag sees what the + // controller failed to bound on its own. + const float a = fabsf(scaled); + if (a > local_max) local_max = a; + + // Asymmetric clamp: [-REWARD_CLAMP_LOSS, +REWARD_CLAMP_WIN]. + // Negative scaled values (losses) get bounded at -3.0; positive + // (wins) at +1.0. Preserves loss aversion while bounding V + // regression target. + const float clamped = fmaxf(-REWARD_CLAMP_LOSS, + fminf(scaled, REWARD_CLAMP_WIN)); + rewards[b] = clamped; + } + s_max[tid] = local_max; + __syncthreads(); + + // Pass 2 — block tree reduce (no atomics per pearl_no_atomicadd). + for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) { + if (tid < stride) { + const float a = s_max[tid]; + const float b = s_max[tid + stride]; + s_max[tid] = a > b ? a : b; + } + __syncthreads(); + } + + // Thread 0 publishes the per-step pre-clamp max to the diagnostic + // ISV slot. This is a POINT measurement (not an EMA) — each + // training step overwrites with the current batch's max so the + // diag emitter sees the value for THIS step's reward batch. + if (tid == 0) { + isv[RL_MAX_ABS_SCALED_REWARD_PRE_CLAMP_INDEX] = s_max[0]; + } } diff --git a/crates/ml-alpha/examples/alpha_rl_train.rs b/crates/ml-alpha/examples/alpha_rl_train.rs index cf2c0d535..85303d6d9 100644 --- a/crates/ml-alpha/examples/alpha_rl_train.rs +++ b/crates/ml-alpha/examples/alpha_rl_train.rs @@ -44,7 +44,8 @@ use ml_alpha::rl::isv_slots::{ RL_LR_PI_WARMUP_COUNTER_INDEX, RL_LR_Q_BEST_LOSS_INDEX, RL_LR_Q_INDEX, 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_MEAN_ABS_PNL_EMA_INDEX, + 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, @@ -541,11 +542,21 @@ fn main() -> Result<()> { "warmup": isv[RL_LR_V_WARMUP_COUNTER_INDEX] }, }, "replay_len": trainer.replay.len(), + // Post-scale, POST-clamp reward stats — what V regression + // and Q distributional projection actually saw this step. + // The `scaled_pre_clamp_max` field (below) is the SAME + // batch's max BEFORE the [-3, +1] clamp fired. If + // scaled_pre_clamp_max > 1.0 (clamp threshold on the win + // side) repeatedly, the clamp is doing load-bearing work + // — i.e. the reward_scale controller is failing to keep + // typical trade magnitudes within the C51 atom support. "rewards": { - "sum": reward_sum, - "max": reward_max, - "min": reward_min, - "abs_max": reward_abs_max, + "sum": reward_sum, + "max": reward_max, + "min": reward_min, + "abs_max": reward_abs_max, + "scaled_pre_clamp_max": + isv[RL_MAX_ABS_SCALED_REWARD_PRE_CLAMP_INDEX], }, "done_count": done_count, // Action histogram: indices match `Action` enum diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index f5256f77f..8e848710a 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -195,6 +195,22 @@ pub const RL_LR_Q_WARMUP_COUNTER_INDEX: usize = 436; pub const RL_LR_PI_WARMUP_COUNTER_INDEX: usize = 437; pub const RL_LR_V_WARMUP_COUNTER_INDEX: usize = 438; +/// Per-step pre-clamp `max(|scaled_reward|)` over the current batch. +/// Written by `apply_reward_scale.cu` after computing `scaled = raw * +/// isv[RL_REWARD_SCALE_INDEX]` and BEFORE the asymmetric clamp +/// `[-REWARD_CLAMP_LOSS, +REWARD_CLAMP_WIN]` fires. POINT measurement +/// (overwritten every step, not EMA-smoothed) — surfaces in the diag +/// JSONL to show whether the reward_scale controller is keeping +/// scaled rewards within the clamp window on its own or whether the +/// clamp is doing load-bearing work. +/// +/// Healthy controller behaviour: this value stays ≤ REWARD_CLAMP_WIN +/// = 1.0 most steps, with occasional rare excursions on fat-tail +/// trade outcomes. If it sits persistently above the clamp bound, the +/// reward_scale controller has failed to track the typical trade +/// magnitude and the clamp is materially shaping training signal. +pub const RL_MAX_ABS_SCALED_REWARD_PRE_CLAMP_INDEX: usize = 439; + /// 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 = 439; +pub const RL_SLOTS_END: usize = 440; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index 542d3f1aa..a0489af29 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -1423,22 +1423,27 @@ impl IntegratedTrainer { } /// Phase R6: launch `apply_reward_scale` — element-wise - /// `rewards[b] *= isv[RL_REWARD_SCALE_INDEX = 406]`. R7 wires - /// this into `step_with_lobsim` between - /// `extract_realized_pnl_delta` and the advantage/return - /// computation; currently exposed as a public API ready for that - /// caller. + /// `scaled = rewards[b] * isv[RL_REWARD_SCALE_INDEX = 406]`, + /// then asymmetric clamp to `[-REWARD_CLAMP_LOSS, +REWARD_CLAMP_WIN]` + /// = `[-3, +1]` (per + /// `pearl_audit_unboundedness_for_implicit_asymmetry`). + /// Also writes per-step pre-clamp `max(|scaled|)` to + /// `ISV[RL_MAX_ABS_SCALED_REWARD_PRE_CLAMP_INDEX = 439]` for diag. + /// + /// Single-block layout (no atomics per `pearl_no_atomicadd`). + /// Block dim = min(b_size, 256); shared memory = block_dim · 4 B + /// for the tree reduction. pub fn launch_apply_reward_scale( &self, rewards_d: &mut CudaSlice, b_size: usize, ) -> Result<()> { debug_assert_eq!(rewards_d.len(), b_size); - let grid_x = ((b_size as u32) + 31) / 32; + let block_x: u32 = (b_size as u32).min(256).max(1); let cfg = LaunchConfig { - grid_dim: (grid_x.max(1), 1, 1), - block_dim: (32, 1, 1), - shared_mem_bytes: 0, + grid_dim: (1, 1, 1), + block_dim: (block_x, 1, 1), + shared_mem_bytes: block_x * (std::mem::size_of::() as u32), }; let b_size_i = b_size as i32; let mut launch = self.stream.launch_builder(&self.apply_reward_scale_fn); @@ -2503,25 +2508,21 @@ impl IntegratedTrainer { self.launch_rl_controllers_per_step() .context("launch_rl_controllers_per_step")?; - // Apply the reward scale on device (in place on rewards_d). - // Reads ISV[406] which the controller just updated. + // Apply the reward scale on device (in place on rewards_d) + + // asymmetric clamp [-3, +1] + per-step pre-clamp max diag. + // Reads ISV[406] which the controller just updated; writes + // ISV[439] for the diag emitter. + // + // Single-block layout: tree reduction needs every thread in + // the same block to participate in __syncthreads, so we cap + // block_x at min(b_size, 256). For b_size > 256 the kernel's + // grid-stride loop walks the rest. { - let rewards_clone = self.rewards_d.clone(); - let _ = rewards_clone; // keep clone alive for the duration - } - // Borrow-checker note: launch_apply_reward_scale takes &self - // + &mut CudaSlice; calling self.launch_apply_reward_scale - // on self.rewards_d directly works under disjoint-field rule. - // Use a scoped block so the &mut on self.rewards_d ends - // before the next self.* call. - { - // Have to break the borrow: take a local mutable handle - // via std::mem::replace temporarily. Cleaner: just inline - // the launch here. + let block_x: u32 = (b_size as u32).min(256).max(1); let cfg = LaunchConfig { - grid_dim: (((b_size as u32) + 31) / 32, 1, 1), - block_dim: (32, 1, 1), - shared_mem_bytes: 0, + grid_dim: (1, 1, 1), + block_dim: (block_x, 1, 1), + shared_mem_bytes: block_x * (std::mem::size_of::() as u32), }; let mut launch = self.stream.launch_builder(&self.apply_reward_scale_fn); launch