diff --git a/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu b/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu index 23b39a3f0..f89479026 100644 --- a/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu +++ b/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu @@ -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 diff --git a/crates/ml-alpha/examples/alpha_rl_train.rs b/crates/ml-alpha/examples/alpha_rl_train.rs index b81590201..b62432198 100644 --- a/crates/ml-alpha/examples/alpha_rl_train.rs +++ b/crates/ml-alpha/examples/alpha_rl_train.rs @@ -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, diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index 48eb976fa..1c52b7b6f 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -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; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index 003420968..2327f4b69 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -3587,6 +3587,7 @@ impl IntegratedTrainer { pub fn launch_apply_reward_scale( &self, rewards_d: &mut CudaSlice, + dones_d: &CudaSlice, 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(