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

fix(rl): asymmetric clamp on scaled reward + pre-clamp |max| diag

Browse Source 
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 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-23 20:13:50 +02:00
parent d5c29fb4fa
commit 20c7852b66
4 changed files with 162 additions and 53 deletions
Download Patch File Download Diff File Expand all files Collapse all files

123
crates/ml-alpha/cuda/apply_reward_scale.cu
View File

@@ -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];
}
}

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

@@ -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

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

@@ -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;

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

@@ -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<f32>,
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::<f32>() 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::<f32>() as u32),
};
let mut launch = self.stream.launch_builder(&self.apply_reward_scale_fn);
launch