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

feat(rl): adaptive RATIO + EWMA V_MIN/V_MAX + λ_distill bump

Browse Source 
wwcsz analysis identified atom-resolution starvation + asymmetric
RATIO mismatch as the empirical ceiling on win rate (38.56% vs
break-even 45.3%). Three coupled fixes shipped in one pass per
"no deferrals":

(a) Adaptive RATIO from observed |loss|/|win| EMAs:
  - apply_reward_scale tracks max(-scaled, 0) per step (slot 489)
  - rl_reward_clamp_controller maintains neg_max_ema (slot 490,
    sparse-aware like pos_max_ema)
  - RATIO = clamp(MIN=1.0, neg_ema/pos_ema, MAX=3.0); writes to
    slot 481
  - Removes built-in 3:1 loss-aversion bias when reality is
    symmetric (wwcsz: actual avg|loss|/avg(win) = 0.83). Floor
    1.0 prevents inverted asymmetry; ceiling 3.0 preserves
    original loss-aversion as the worst case.

(b) C51 V_MAX/V_MIN: ratchet → slow EWMA (α=0.001, half-life ~700
    steps):
  - Static ratchet wasted atom resolution on rare tails — wwcsz
    had V_MIN=-60, V_MAX=20 but realized rewards mostly in [-5, +5]
    (Δz=4 vs typical reward magnitude 1-5)
  - Slow EWMA lets atom span shrink toward active reward range,
    gaining resolution where data lives. Floors at [-1, +1]
    preserve original C51 baseline as the worst case.
  - Slow α gives Q's atom mapping time to be valid across
    encoder/head co-adaptation (vs aggressive EWMA which would
    invalidate Q's learned distribution every step)

(c) Q→π distillation λ bumped 0.01 → 0.05:
  - wwcsz showed KL dropped 2.10 → 0.30 with λ=0.01 — Q signal
    landing but conservatively. Bump tests whether stronger Q
    pull translates to better policy → better R/done.

Diag exposes neg_scaled_max + neg_scaled_max_ema so the RATIO
adaptation chain is observable.

apply_reward_scale shared_mem doubled from 2× to 3× block × f32
to fit the three parallel reductions (abs, pos, neg).

Companion to investigation (e) — n_rollout_steps controller was
suspected of misalignment (256-8192 vs trade_duration ≈ 6 steps)
but turned out to be a K-loop param, not used in Bellman target.
1-step Bellman with γ-bootstrap is the actual mechanism; closed
without code change.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-24 14:20:05 +02:00
parent 79756a2153
commit 185add7dc8
5 changed files with 152 additions and 54 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -70,13 +70,21 @@
// reach the C51 atom span properly.
#define RL_POS_SCALED_REWARD_MAX_INDEX 478
// wwcsz followup 2026-05-24: also track per-step max(-scaled, 0) — the
// loss-side tail — so the RATIO (LOSS/WIN) can become adaptive from
// observed loss/win magnitudes rather than being hardcoded 3:1. wwcsz
// data showed avg_loss=-$4.36 vs avg_win=+$5.27 (actual ratio 0.83);
// the 3:1 built-in bias was over-emphasising loss-aversion vs reality.
#define RL_NEG_SCALED_REWARD_MAX_INDEX 489
// 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).
// Shared-mem layout: caller passes `shared_mem_bytes = 2 * block_dim *
// sizeof(float)`. Lower half [0..block_dim) tracks max(|scaled|) for
// the diag slot; upper half [block_dim..2*block_dim) tracks max(positive
// scaled, 0) for the adaptive clamp controller's input signal.
// Shared-mem layout: caller passes `shared_mem_bytes = 3 * block_dim *
// sizeof(float)`. Three parallel reductions:
// [0..block_dim) s_abs — max(|scaled|) for diag slot
// [block_dim..2*block_dim) s_pos — max(positive scaled, 0)
// [2*block_dim..3*block_dim) s_neg — max(-scaled, 0) for adaptive RATIO
extern "C" __global__ void apply_reward_scale(
float* __restrict__ isv, // ISV bus IN/OUT
float* __restrict__ rewards, // [b_size] IN/OUT (scaled+clamped)
@@ -85,6 +93,7 @@ extern "C" __global__ void apply_reward_scale(
extern __shared__ float smem[];
float* s_abs = smem;
float* s_pos = smem + blockDim.x;
float* s_neg = smem + 2 * blockDim.x;
const int tid = threadIdx.x;
const float scale = isv[RL_REWARD_SCALE_INDEX];
// ISV-driven clamp bounds (was hardcoded 1.0 / 3.0). Reading per
@@ -93,11 +102,13 @@ extern "C" __global__ void apply_reward_scale(
const float clamp_win = isv[RL_REWARD_CLAMP_WIN_INDEX];
const float clamp_loss = isv[RL_REWARD_CLAMP_LOSS_INDEX];
// Pass 1 — scale, clamp, accumulate per-thread max |scaled| AND max
// max(positive scaled, 0). Two independent local maxes tracked in
// a single grid-stride loop so the kernel's HBM read pass is shared.
// Pass 1 — scale, clamp, accumulate per-thread max |scaled|, max
// max(positive scaled, 0), AND max(-scaled, 0). Three independent
// local maxes tracked in a single grid-stride loop so the kernel's
// HBM read pass is shared.
float local_abs = 0.0f;
float local_pos = 0.0f;
float local_neg = 0.0f;
for (int b = tid; b < b_size; b += blockDim.x) {
const float raw = rewards[b];
const float scaled = raw * scale;
@@ -109,14 +120,17 @@ extern "C" __global__ void apply_reward_scale(
// Track positive tail BEFORE clamping for the adaptive clamp
// controller — winning-trade signal that the current static
// WIN=1.0 bound clips. Negative-tail (losses) tracked
// implicitly via max(|scaled|) at the diag slot since losses
// dominate the abs-distribution; the controller assumes the
// documented loss-aversion ratio (LOSS = ratio × WIN) so we
// don't need a separate negative-tail tracker.
// WIN=1.0 bound clips.
const float p = fmaxf(scaled, 0.0f);
if (p > local_pos) local_pos = p;
// Track negative tail magnitude (max(-scaled, 0)) for adaptive
// RATIO calculation per wwcsz followup. Separate from |scaled|
// because we need positive-only and negative-only EMAs to
// compute observed loss/win magnitude ratio.
const float n = fmaxf(-scaled, 0.0f);
if (n > local_neg) local_neg = n;
// Asymmetric clamp: [-clamp_loss, +clamp_win] from ISV.
// Defaults preserve loss aversion (loss > win) while bounding
// V regression target.
@@ -126,10 +140,11 @@ extern "C" __global__ void apply_reward_scale(
}
s_abs[tid] = local_abs;
s_pos[tid] = local_pos;
s_neg[tid] = local_neg;
__syncthreads();
// Pass 2 — block tree reduce (no atomics per pearl_no_atomicadd).
// Both reductions advance together to share the __syncthreads barrier.
// Three reductions advance together to share the __syncthreads barrier.
for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
if (tid < stride) {
const float a_a = s_abs[tid];
@@ -138,15 +153,19 @@ extern "C" __global__ void apply_reward_scale(
const float p_a = s_pos[tid];
const float p_b = s_pos[tid + stride];
s_pos[tid] = p_a > p_b ? p_a : p_b;
const float n_a = s_neg[tid];
const float n_b = s_neg[tid + stride];
s_neg[tid] = n_a > n_b ? n_a : n_b;
}
__syncthreads();
}
// Thread 0 publishes per-step maxes to ISV. Both are POINT
// Thread 0 publishes per-step maxes to ISV. All are POINT
// measurements; the controller maintains the EMA over the positive
// signal in its own slot.
// and negative signals in its own slots.
if (tid == 0) {
isv[RL_MAX_ABS_SCALED_REWARD_PRE_CLAMP_INDEX] = s_abs[0];
isv[RL_POS_SCALED_REWARD_MAX_INDEX] = s_pos[0];
isv[RL_NEG_SCALED_REWARD_MAX_INDEX] = s_neg[0];
}
}

114
crates/ml-alpha/cuda/rl_reward_clamp_controller.cu
View File

@@ -68,15 +68,31 @@
// 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.
// C51 atom span slots (audit 2026-05-24 second follow-up).
// Initial design: ratchet (monotone-grow). wwcsz followup
// 2026-05-24: replaced with slow EWMA (α=0.001, half-life ~700
// steps) to focus atom resolution on the ACTIVE reward range.
// Static ratchet wasted resolution on rare tails — with avg
// rewards in [-5, +5] but atom span at [-60, +20], Δz=4 meant Q
// couldn't differentiate "slightly winning" from "slightly losing"
// actions. Slow EWMA lets the span shrink toward the current
// active range while floors at [-1, +1] preserve the original C51
// baseline. The slow α gives Q's atom mapping time to be valid
// across encoder/head co-adaptation.
#define RL_C51_V_MAX_INDEX 484
#define RL_C51_V_MIN_INDEX 485
// Adaptive RATIO slots (wwcsz followup) — negative tail signal that
// drives RATIO = LOSS/WIN bound asymmetry. The static RATIO=3.0
// over-emphasised loss-aversion vs observed |loss|/|win| ≈ 0.83.
#define RL_NEG_SCALED_REWARD_MAX_INDEX 489
#define RL_NEG_SCALED_REWARD_MAX_EMA_INDEX 490
#define MIN_WIN 1.0f
#define MIN_RATIO 1.0f // no inverted asymmetry (loss never < win)
#define MAX_RATIO 3.0f // no worse than original loss-aversion
#define V_BOUND_FLOOR 1.0f // |V_MIN| and V_MAX both ≥ 1.0
#define V_BOUND_EWMA_ALPHA 0.001f // slow — half-life ~700 steps
#define WIENER_ALPHA_FLOOR 0.4f
// MARGIN Schulman bounded-step bounds + adjust rate per
@@ -178,43 +194,77 @@ extern "C" __global__ void rl_reward_clamp_controller(
isv[RL_REWARD_CLAMP_MARGIN_INDEX] = margin;
}
// Step 4: compute adaptive WIN bound: MARGIN × EMA, floored.
// No upper cap per the header note — the MARGIN ∈ [1, 5] × EMA
// composition is the structural bound; downstream C51 projection
// saturates beyond V_MAX without needing this kernel to enforce.
const float win_eff = fmaxf(MIN_WIN, margin * ema_new);
// ── Adaptive RATIO from observed neg/pos EMAs (wwcsz followup). ──
//
// The static RATIO=3.0 baked 3:1 loss-aversion into Q's value
// representation. wwcsz showed observed |loss|/|win| ≈ 0.83,
// making Q over-cautious. Track per-step max(-scaled, 0) via a
// sparse-aware EMA (same discipline as pos_max_ema) and compute
// RATIO = clamp(MIN=1.0, neg_ema/pos_ema, MAX=3.0). Floor 1.0
// prevents inverted asymmetry; ceiling 3.0 preserves original
// loss-aversion as the worst case.
const float neg_max = isv[RL_NEG_SCALED_REWARD_MAX_INDEX];
const float neg_prev = isv[RL_NEG_SCALED_REWARD_MAX_EMA_INDEX];
float neg_new = neg_prev;
if (neg_max > 0.0f) {
if (neg_prev == 0.0f) {
neg_new = neg_max;
} else {
const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR);
neg_new = (1.0f - a) * neg_prev + a * neg_max;
}
isv[RL_NEG_SCALED_REWARD_MAX_EMA_INDEX] = neg_new;
}
// Adaptive RATIO. Defer if EMAs not warmed up yet — keep
// statically-seeded RATIO=3.0 until both EMAs are non-zero.
if (ema_new > 0.0f && neg_new > 0.0f) {
const float ratio_target = neg_new / ema_new;
const float ratio_clamped = fmaxf(MIN_RATIO,
fminf(MAX_RATIO, ratio_target));
isv[RL_REWARD_CLAMP_RATIO_INDEX] = ratio_clamped;
}
// LOSS = RATIO × WIN preserves asymmetry. RATIO seeded to 3.0.
// Step 4: compute adaptive WIN/LOSS bounds: MARGIN × EMA, floored.
// No upper cap on WIN per the header note — the MARGIN ∈ [1, 5]
// × EMA composition is the structural bound. LOSS = RATIO × WIN
// with RATIO now adaptive.
const float win_eff = fmaxf(MIN_WIN, margin * ema_new);
const float ratio = isv[RL_REWARD_CLAMP_RATIO_INDEX];
const float loss_eff = ratio * win_eff;
isv[RL_REWARD_CLAMP_WIN_INDEX] = win_eff;
isv[RL_REWARD_CLAMP_LOSS_INDEX] = loss_eff;
// ── C51 atom span ratchet (audit follow-up). ──────────────────
// ── C51 atom span — slow EWMA (was ratchet). ──────────────────
//
// 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).
// wwcsz followup: replace monotone-grow with slow EWMA so atom
// resolution focuses on the active reward range. Static ratchet
// wasted Δz on rare tails (V_MAX=20 with rewards mostly in
// [-5, +5] → Δz=2 wasted on extremes). Slow α=0.001
// (half-life ~700 steps) gives Q's atom mapping time to be valid
// across encoder/head co-adaptation. Floors at [-1, +1]
// preserve the original C51 baseline as the worst case.
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);
const float v_max_target = fmaxf(V_BOUND_FLOOR, win_eff);
const float v_min_target = fminf(-V_BOUND_FLOOR, -loss_eff);
// Bootstrap on sentinel 0 — first emit replaces directly (per
// pearl_first_observation_bootstrap).
float v_max_new, v_min_new;
if (v_max_prev == 0.0f) {
v_max_new = v_max_target;
} else {
v_max_new = (1.0f - V_BOUND_EWMA_ALPHA) * v_max_prev
+ V_BOUND_EWMA_ALPHA * v_max_target;
v_max_new = fmaxf(V_BOUND_FLOOR, v_max_new);
}
if (v_min_prev == 0.0f) {
v_min_new = v_min_target;
} else {
v_min_new = (1.0f - V_BOUND_EWMA_ALPHA) * v_min_prev
+ V_BOUND_EWMA_ALPHA * v_min_target;
v_min_new = fminf(-V_BOUND_FLOOR, v_min_new);
}
isv[RL_C51_V_MAX_INDEX] = v_max_new;
isv[RL_C51_V_MIN_INDEX] = v_min_new;
}

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

@@ -69,6 +69,7 @@ use ml_alpha::rl::isv_slots::{
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_Q_DISTILL_LAMBDA_INDEX, RL_Q_DISTILL_TEMPERATURE_INDEX, RL_Q_DISTILL_KL_EMA_INDEX,
RL_NEG_SCALED_REWARD_MAX_INDEX, RL_NEG_SCALED_REWARD_MAX_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,
RL_PPO_RATIO_CLAMP_MAX_INDEX, RL_Q_DIVERGENCE_EMA_INDEX, RL_Q_GRAD_NORM_EMA_INDEX,
@@ -689,6 +690,11 @@ fn main() -> Result<()> {
// successfully chasing Q's preferences.
"q_distill_kl_ema":
isv[RL_Q_DISTILL_KL_EMA_INDEX],
// Negative-tail tracking — input to adaptive RATIO.
"neg_scaled_max":
isv[RL_NEG_SCALED_REWARD_MAX_INDEX],
"neg_scaled_max_ema":
isv[RL_NEG_SCALED_REWARD_MAX_EMA_INDEX],
},
// audit — PPO importance-ratio clamp diagnostics.
//

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

@@ -42,8 +42,9 @@
//! | 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) |
//! | 486-488 | 3 Q→π distillation slots (λ + τ + KL_ema) | rl_isv_write (seed) + rl_q_pi_distill_grad (ADDS to pi_grad_logits + writes KL_ema) |
//! | 489-490 | 2 negative-tail tracking slots (raw + EMA) for adaptive RATIO | apply_reward_scale + rl_reward_clamp_controller |
//!
//! Total: 89 slots; `RL_SLOTS_END = 489`.
//! Total: 91 slots; `RL_SLOTS_END = 491`.
/// Discount factor γ. Controller input: mean trade duration / anchor.
/// Bootstrap 0.99.
@@ -684,6 +685,27 @@ pub const RL_Q_DISTILL_TEMPERATURE_INDEX: usize = 487;
/// (KL should drop over training).
pub const RL_Q_DISTILL_KL_EMA_INDEX: usize = 488;
// ─── Adaptive RATIO from observed loss/win magnitudes (wwcsz followup) ──
//
// wwcsz showed avg_loss=-$4.36 vs avg_win=+$5.27 (actual |loss|/|win|
// = 0.83). The static RATIO=3.0 baked into the reward clamp +
// C51 atom span was over-emphasising loss-aversion vs reality.
// Make RATIO adaptive from observed negative/positive tail EMAs,
// floored at 1.0 (no inverted asymmetry) and ceiled at 3.0
// (no worse than original loss-aversion).
/// Per-step max(-scaled, 0) — POINT measurement published by
/// `apply_reward_scale` for the adaptive RATIO controller.
/// Overwritten every step.
pub const RL_NEG_SCALED_REWARD_MAX_INDEX: usize = 489;
/// EMA of `RL_NEG_SCALED_REWARD_MAX_INDEX` — maintained by
/// `rl_reward_clamp_controller` via Wiener-α blend (sparse-aware:
/// only updates when neg_max > 0). The controller computes
/// RATIO = clamp(1.0, neg_ema / pos_ema, 3.0) from this and the
/// existing positive-tail EMA.
pub const RL_NEG_SCALED_REWARD_MAX_EMA_INDEX: usize = 490;
/// 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 = 489;
pub const RL_SLOTS_END: usize = 491;

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

@@ -1256,9 +1256,10 @@ impl IntegratedTrainer {
(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),
// Q→π distillation (vj5f6 followup) — λ small so PPO
// dominates; τ=1.0 for canonical softmax(E_Q).
(crate::rl::isv_slots::RL_Q_DISTILL_LAMBDA_INDEX, 0.01),
// Q→π distillation (vj5f6 followup) — λ raised to 0.05
// (wwcsz followup) so Q's now-calibrated signal more
// strongly drives π updates; τ=1.0 canonical Boltzmann.
(crate::rl::isv_slots::RL_Q_DISTILL_LAMBDA_INDEX, 0.05),
(crate::rl::isv_slots::RL_Q_DISTILL_TEMPERATURE_INDEX, 1.0),
(crate::rl::isv_slots::RL_KL_TARGET_INDEX, 0.01),
(crate::rl::isv_slots::RL_IMPROVEMENT_THRESHOLD_INDEX, 0.99),
@@ -1753,7 +1754,7 @@ impl IntegratedTrainer {
let cfg = LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (block_x, 1, 1),
shared_mem_bytes: 2 * block_x * (std::mem::size_of::<f32>() as u32),
shared_mem_bytes: 3 * 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);
@@ -3134,7 +3135,7 @@ impl IntegratedTrainer {
let cfg = LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (block_x, 1, 1),
shared_mem_bytes: 2 * block_x * (std::mem::size_of::<f32>() as u32),
shared_mem_bytes: 3 * block_x * (std::mem::size_of::<f32>() as u32),
};
let mut launch = self.stream.launch_builder(&self.apply_reward_scale_fn);
launch