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

feat(cuda): done-gated trade-level EMAs drive adaptive LOSS clamp

Browse Source 
The old approach used pre-clamp reward tail maxes (neg_ema/pos_ema ≈
0.07) → LOSS hit floor at 1.0 → no loss aversion → wr dropped to 0.44.

New approach: the clamp controller loops over done-step rewards,
separates wins from losses, maintains two ISV EMAs:
- RL_DONE_WIN_MAGNITUDE_EMA_INDEX (585) — avg winning trade magnitude
- RL_DONE_LOSS_MAGNITUDE_EMA_INDEX (586) — avg losing trade magnitude

LOSS = clamp(1.0, loss_ema/win_ema × 1.1, 3.0). With observed L/W
ratio of 1.32, LOSS should settle at ~1.45 — preserves loss aversion
while preventing the 3× tolerance that caused massive losses.

On top of dd049d9a4 baseline (wr=0.567). Only this one change.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-28 01:35:44 +02:00
parent c0143d7e34
commit 9355118e8b
4 changed files with 134 additions and 35 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -14,9 +14,14 @@
// isv[RL_POS_SCALED_REWARD_MAX_EMA_INDEX = 479] — controller's EMA state
// isv[RL_REWARD_CLAMP_MARGIN_INDEX = 480] — k multiplier (default 1.5)
// isv[RL_REWARD_CLAMP_RATIO_INDEX = 481] — loss/win ratio (default 3.0)
// rewards[b_size] — post-scale reward buffer
// dones[b_size] — done flags (0/1)
//
// Writes:
// isv[RL_POS_SCALED_REWARD_MAX_EMA_INDEX = 479] — updated EMA
// isv[RL_DONE_WIN_MAGNITUDE_EMA_INDEX = 585] — done-gated mean(+reward)
// isv[RL_DONE_LOSS_MAGNITUDE_EMA_INDEX = 586]— done-gated mean(|−reward|)
// isv[RL_REWARD_CLAMP_RATIO_INDEX = 481] — adaptive RATIO from trade EMAs
// isv[RL_REWARD_CLAMP_WIN_INDEX = 452] — adaptive WIN bound
// isv[RL_REWARD_CLAMP_LOSS_INDEX = 453] — adaptive LOSS = ratio × WIN
//
@@ -88,6 +93,12 @@
#define RL_NEG_SCALED_REWARD_MAX_INDEX 489
#define RL_NEG_SCALED_REWARD_MAX_EMA_INDEX 490
// Done-gated trade-level magnitude EMAs — mean(reward | done & sign).
// Drives adaptive RATIO from actual trade win/loss magnitudes instead
// of per-step tail maxes which conflate hold-step noise with trades.
#define RL_DONE_WIN_MAGNITUDE_EMA_INDEX 585
#define RL_DONE_LOSS_MAGNITUDE_EMA_INDEX 586
#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
@@ -110,7 +121,10 @@
extern "C" __global__ void rl_reward_clamp_controller(
float* __restrict__ isv,
float alpha
float alpha,
const float* __restrict__ rewards, // [b_size] post-scale rewards
const float* __restrict__ dones, // [b_size] 0/1 done flags
int b_size
) {
if (threadIdx.x != 0 || blockIdx.x != 0) return;
@@ -194,45 +208,97 @@ extern "C" __global__ void rl_reward_clamp_controller(
isv[RL_REWARD_CLAMP_MARGIN_INDEX] = margin;
}
// ── Adaptive RATIO from observed neg/pos EMAs (wwcsz followup). ──
// ── Done-gated trade-level EMAs for adaptive LOSS ratio. ─────────
//
// 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;
// The prior RATIO controller used per-step max(pos/neg scaled
// reward) EMAs. That conflated hold-step noise (pos_max from tiny
// shaping rewards) with actual closed-trade magnitudes. With deep
// replay, stale-scale drift collapsed the ratio to the MIN_RATIO
// floor (per `pearl_replay_reward_renormalization_at_sample_time`).
//
// Fix: loop over rewards[b] × dones[b] to accumulate ONLY
// closed-trade magnitudes. At b_size=1024, ~40 dones/step — the
// loop is trivially cheap for a single-thread kernel. No atomics
// needed (per `feedback_no_atomicadd`).
//
// Sparse-aware: only update each EMA when there are actual
// winning/losing trades this step. Sentinel-bootstrap per
// `pearl_first_observation_bootstrap`.
float win_sum = 0.0f, loss_sum = 0.0f;
int win_count = 0, loss_count = 0;
for (int b = 0; b < b_size; b++) {
if (dones[b] > 0.5f) {
if (rewards[b] > 0.0f) {
win_sum += rewards[b];
win_count++;
} else if (rewards[b] < 0.0f) {
loss_sum += -rewards[b];
loss_count++;
}
}
}
if (win_count > 0) {
const float win_avg = win_sum / (float)win_count;
const float prev = isv[RL_DONE_WIN_MAGNITUDE_EMA_INDEX];
if (prev == 0.0f) {
isv[RL_DONE_WIN_MAGNITUDE_EMA_INDEX] = win_avg;
} else {
const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR);
neg_new = (1.0f - a) * neg_prev + a * neg_max;
isv[RL_DONE_WIN_MAGNITUDE_EMA_INDEX] =
(1.0f - a) * prev + a * win_avg;
}
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;
if (loss_count > 0) {
const float loss_avg = loss_sum / (float)loss_count;
const float prev = isv[RL_DONE_LOSS_MAGNITUDE_EMA_INDEX];
if (prev == 0.0f) {
isv[RL_DONE_LOSS_MAGNITUDE_EMA_INDEX] = loss_avg;
} else {
const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR);
isv[RL_DONE_LOSS_MAGNITUDE_EMA_INDEX] =
(1.0f - a) * prev + a * loss_avg;
}
}
// Adaptive RATIO from trade-level win/loss magnitudes. The 1.1×
// multiplier preserves mild loss-aversion beyond break-even —
// without it, symmetric win/loss gives RATIO=1.0 and Q treats
// losses and wins equally (wr drops to 0.44 per the clean-clamp
// experiment). Defer if EMAs not warmed up — keep statically-seeded
// RATIO=3.0 until both EMAs are non-zero.
const float win_ema = isv[RL_DONE_WIN_MAGNITUDE_EMA_INDEX];
const float loss_ema = isv[RL_DONE_LOSS_MAGNITUDE_EMA_INDEX];
if (win_ema > 0.0f && loss_ema > 0.0f) {
const float trade_ratio = loss_ema / win_ema * 1.1f;
const float ratio_clamped = fmaxf(MIN_RATIO,
fminf(MAX_RATIO, ratio_target));
fminf(MAX_RATIO, trade_ratio));
isv[RL_REWARD_CLAMP_RATIO_INDEX] = ratio_clamped;
}
// Step 4 (G.2): WIN stays STRUCTURAL. LOSS adapts from observed
// neg/pos ratio to give Q accurate loss magnitude perception.
(void) margin;
if (ema_new > 0.0f && neg_new > 0.0f) {
const float observed_ratio = neg_new / ema_new;
const float adaptive_loss = fmaxf(1.0f, fminf(observed_ratio * 1.1f, 3.0f));
isv[RL_REWARD_CLAMP_LOSS_INDEX] = adaptive_loss;
// Legacy per-step neg-tail EMA — still maintained for diagnostic
// continuity (JSONL logs neg_scaled_max_ema). No longer drives
// RATIO.
{
const float neg_max = isv[RL_NEG_SCALED_REWARD_MAX_INDEX];
const float neg_prev = isv[RL_NEG_SCALED_REWARD_MAX_EMA_INDEX];
if (neg_max > 0.0f) {
if (neg_prev == 0.0f) {
isv[RL_NEG_SCALED_REWARD_MAX_EMA_INDEX] = neg_max;
} else {
const float a = fmaxf(alpha, WIENER_ALPHA_FLOOR);
isv[RL_NEG_SCALED_REWARD_MAX_EMA_INDEX] =
(1.0f - a) * neg_prev + a * neg_max;
}
}
}
// Step 4 (G.2): WIN / LOSS clamp bounds remain STRUCTURAL — the
// G.2 feedback loop required BOTH clamp widening AND span widening.
// With clamp frozen here, only the atom span adapts (Step 5 below).
// Suppress unused warnings on diagnostic-only state.
(void) margin;
// ── Step 5: C51 atom span adaptation from observed reward EMAs. ──
//
// G.2 disabled this because atom-span growth + clamp-bound growth

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

@@ -79,6 +79,7 @@ use ml_alpha::rl::isv_slots::{
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_DONE_WIN_MAGNITUDE_EMA_INDEX, RL_DONE_LOSS_MAGNITUDE_EMA_INDEX,
RL_Q_DISTILL_KL_TARGET_INDEX, RL_REWARD_SCALE_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,
@@ -984,6 +985,8 @@ fn main() -> Result<()> {
"sac_entropy_target": isv[RL_SAC_ENTROPY_TARGET_INDEX],
"action_entropy_ema": isv[RL_ACTION_ENTROPY_EMA_INDEX],
"reward_scale_min": isv[RL_REWARD_SCALE_MIN_INDEX],
"done_win_magnitude_ema": isv[RL_DONE_WIN_MAGNITUDE_EMA_INDEX],
"done_loss_magnitude_ema": isv[RL_DONE_LOSS_MAGNITUDE_EMA_INDEX],
},
// audit — Q vs π action agreement EMA at slot 407
// (previously dead). 1.0 = perfect ranking consistency,

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

@@ -1111,5 +1111,21 @@ pub const RL_ACTION_ENTROPY_EMA_INDEX: usize = 583;
/// Bootstrap: 0.85 (15% exploration floor).
pub const RL_CONF_GATE_MAX_HOLD_FRAC_INDEX: usize = 584;
/// Done-gated EMA of winning trade magnitudes — mean(reward | done,
/// reward > 0). Maintained by `rl_reward_clamp_controller` from the
/// post-scale `rewards_d` buffer. Sparse-aware: only updates on steps
/// with at least one positive done-step reward. Sentinel 0 →
/// bootstrap to first non-zero observation per
/// `pearl_first_observation_bootstrap`.
pub const RL_DONE_WIN_MAGNITUDE_EMA_INDEX: usize = 585;
/// Done-gated EMA of losing trade magnitudes — mean(|reward| | done,
/// reward < 0). Maintained by `rl_reward_clamp_controller` alongside
/// `RL_DONE_WIN_MAGNITUDE_EMA_INDEX`. The LOSS/WIN ratio from these
/// two EMAs drives the adaptive RATIO in slot 481 — trade-level
/// loss-aversion grounded in actual win/loss magnitudes rather than
/// per-step tail maxes.
pub const RL_DONE_LOSS_MAGNITUDE_EMA_INDEX: usize = 586;
/// Last RL-allocated slot index (exclusive).
pub const RL_SLOTS_END: usize = 585;
pub const RL_SLOTS_END: usize = 587;

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

@@ -3587,6 +3587,7 @@ impl IntegratedTrainer {
pub fn launch_apply_reward_scale(
&self,
rewards_d: &mut CudaSlice<f32>,
dones_d: &CudaSlice<f32>,
b_size: usize,
) -> Result<()> {
debug_assert_eq!(rewards_d.len(), b_size);
@@ -3614,13 +3615,18 @@ impl IntegratedTrainer {
// positive-tail max just published into slot 478, maintains
// an EMA in slot 479, writes adaptive WIN/LOSS bounds into
// slots 452/453 for the NEXT step's apply_reward_scale to
// consume. Single-thread kernel; α driven by the same
// controller-cohort floor as the other R5 EMAs.
// consume. Also loops over rewards/dones to compute done-gated
// trade-level EMAs driving RATIO. Single-thread kernel; α
// driven by the same controller-cohort floor as other R5 EMAs.
{
let alpha = RL_LR_CONTROLLER_ALPHA;
let b_size_i = b_size as i32;
let mut args = RawArgs::new();
args.push_ptr(self.isv_dev_ptr);
args.push_f32(alpha);
args.push_ptr(rewards_d.raw_ptr());
args.push_ptr(dones_d.raw_ptr());
args.push_i32(b_size_i);
let mut ptrs = args.build_arg_ptrs();
unsafe {
raw_launch(
@@ -5946,12 +5952,16 @@ impl IntegratedTrainer {
}
// Adaptive reward-clamp controller — reads per-step max from
// apply_reward_scale, maintains EMAs, adapts C51 atom span.
// apply_reward_scale, loops over rewards/dones for done-gated
// trade-level EMAs, maintains RATIO, adapts C51 atom span.
{
let alpha = RL_LR_CONTROLLER_ALPHA;
let mut args = RawArgs::new();
args.push_ptr(self.isv_dev_ptr);
args.push_f32(alpha);
args.push_ptr(self.rewards_d.raw_ptr());
args.push_ptr(self.dones_d.raw_ptr());
args.push_i32(b_size_i);
let mut ptrs = args.build_arg_ptrs();
unsafe {
raw_launch(
@@ -6219,13 +6229,17 @@ impl IntegratedTrainer {
).context("outcome_head.assign_labels")?;
// Adaptive reward-clamp controller — refresh slots 452/453 from
// the just-published positive-tail max in slot 478. The new
// bounds take effect on the NEXT step's apply_reward_scale.
// the just-published positive-tail max in slot 478 and done-gated
// trade-level EMAs from rewards/dones. The new bounds take effect
// on the NEXT step's apply_reward_scale.
{
let alpha = RL_LR_CONTROLLER_ALPHA;
let mut args = RawArgs::new();
args.push_ptr(self.isv_dev_ptr);
args.push_f32(alpha);
args.push_ptr(self.rewards_d.raw_ptr());
args.push_ptr(self.dones_d.raw_ptr());
args.push_i32(b_size_i);
let mut ptrs = args.build_arg_ptrs();
unsafe {
raw_launch(