From 29640b6e6d9d5ec8cf294ffac56736e1a04e3cfe Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Mon, 25 May 2026 19:53:54 +0200 Subject: [PATCH] =?UTF-8?q?feat(rl):=20asymmetric=20trail=20+=20session=20?= =?UTF-8?q?risk=20=E2=80=94=20structural=20P&L=20edge?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Two structural mechanics that produce asymmetric win/loss ratio without the agent needing to learn the behavior: 1. Asymmetric trail decay (rl_asymmetric_trail_decay.cu): - LOSING: trail *= 0.995/step (halves in 139 steps = 35s) - WINNING beyond initial_r: trail = max(trail, profit × 0.5) - NEUTRAL: unchanged (proving zone) Produces: small/quick losses, large/extended wins. 2. Session risk circuit breaker (rl_session_risk_check.cu): - Tracks EMA of realized PnL (α=0.02, slow) - When EMA < -$50: blocks ALL opening actions - Prevents tilt — a losing streak stops the agent from digging deeper Pipeline order: action selection → confidence gate → FRD gate → session risk → min_hold → asymmetric trail → trail_mutate → trail_stop → heat_cap → actions_to_market_targets ISV slots 537-541. RL_SLOTS_END → 542. Co-Authored-By: Claude Opus 4.7 --- crates/ml-alpha/build.rs | 2 + .../cuda/rl_asymmetric_trail_decay.cu | 94 +++++++++++++++++++ crates/ml-alpha/cuda/rl_session_risk_check.cu | 64 +++++++++++++ crates/ml-alpha/src/rl/isv_slots.rs | 22 ++++- crates/ml-alpha/src/trainer/integrated.rs | 83 +++++++++++++++- 5 files changed, 263 insertions(+), 2 deletions(-) create mode 100644 crates/ml-alpha/cuda/rl_asymmetric_trail_decay.cu create mode 100644 crates/ml-alpha/cuda/rl_session_risk_check.cu diff --git a/crates/ml-alpha/build.rs b/crates/ml-alpha/build.rs index b288ef90e..ea7164b72 100644 --- a/crates/ml-alpha/build.rs +++ b/crates/ml-alpha/build.rs @@ -89,6 +89,8 @@ const KERNELS: &[&str] = &[ "rl_gate_threshold_controller", // adaptive gate thresholds from trade frequency (dones EMA) "rl_reward_shaping", // surfer-philosophy: entry cost + short-hold penalty + hold bonus "rl_min_hold_check", // hard minimum hold time — override close actions to Hold before N steps + "rl_asymmetric_trail_decay", // auto-tighten losers, auto-widen winners — structural P&L asymmetry + "rl_session_risk_check", // session-level loss limit circuit breaker ]; // Cache bust v31 — five new reduce / derive kernels populate the input diff --git a/crates/ml-alpha/cuda/rl_asymmetric_trail_decay.cu b/crates/ml-alpha/cuda/rl_asymmetric_trail_decay.cu new file mode 100644 index 000000000..27828f7f8 --- /dev/null +++ b/crates/ml-alpha/cuda/rl_asymmetric_trail_decay.cu @@ -0,0 +1,94 @@ +// rl_asymmetric_trail_decay.cu — structural asymmetric trail management. +// +// Runs EVERY STEP for each active unit. Automatically adjusts trail +// distance based on unrealized P&L direction: +// +// LOSING (mid moved against entry): +// trail *= loss_decay_rate per step (default 0.995) +// → halves in ~139 steps (~35 seconds) +// "Cut losers fast" — the longer a trade stays underwater, +// the tighter the stop gets, accelerating the exit. +// +// WINNING (unrealized > initial_r): +// trail = max(trail, unrealized_profit × win_trail_factor) +// → trail ratchets up with profit, never below current level +// "Let winners run" — the stop tracks 50% of open profit, +// locking in gains while giving room for continuation. +// +// NEUTRAL (between 0 and initial_r): +// trail unchanged — in the "proving zone" where the trade +// hasn't yet earned the right to a wider stop. +// +// This produces asymmetric P&L without the agent needing to LEARN +// when to tighten vs loosen — it's a structural edge baked into +// the mechanics. The agent's a7/a8 trail actions provide additional +// fine-tuning on top. +// +// Runs BEFORE rl_trail_stop_check so the updated trail distances +// are immediately used for breach detection. +// +// Per `feedback_no_atomicadd`: per-batch per-unit element-wise. +// Per `feedback_cpu_is_read_only`: pure device-side. + +#include + +#define MAX_UNITS 4 +#define RL_ASYM_LOSS_DECAY_RATE_INDEX 537 +#define RL_ASYM_WIN_TRAIL_FACTOR_INDEX 538 +#define RL_TRAIL_MIN_INDEX 494 + +extern "C" __global__ void rl_asymmetric_trail_decay( + float* __restrict__ unit_trail_distance, // [B × MAX_UNITS] IN/OUT + const unsigned char* __restrict__ unit_active, // [B × MAX_UNITS] + const float* __restrict__ unit_entry_price, // [B × MAX_UNITS] + const float* __restrict__ unit_initial_r, // [B × MAX_UNITS] + const int* __restrict__ unit_lots, // [B × MAX_UNITS] + const float* __restrict__ bid_px, // [BOOK_LEVELS] + const float* __restrict__ ask_px, // [BOOK_LEVELS] + const float* __restrict__ isv, + int b_size +) { + const int b = blockIdx.x; + const int u = threadIdx.x; + if (b >= b_size || u >= MAX_UNITS) return; + + const int idx = b * MAX_UNITS + u; + if (unit_active[idx] == 0) return; + + const int lots = unit_lots[idx]; + if (lots == 0) return; + + const float entry = unit_entry_price[idx]; + const float trail = unit_trail_distance[idx]; + const float init_r = unit_initial_r[idx]; + const float mid = 0.5f * (bid_px[0] + ask_px[0]); + + const float loss_decay = isv[RL_ASYM_LOSS_DECAY_RATE_INDEX]; + const float win_factor = isv[RL_ASYM_WIN_TRAIL_FACTOR_INDEX]; + const float trail_min = isv[RL_TRAIL_MIN_INDEX]; + + // Unrealized P&L in price units (positive = profitable). + const float direction = (lots > 0) ? 1.0f : -1.0f; + const float unrealized = direction * (mid - entry); + + float new_trail = trail; + + if (unrealized < 0.0f) { + // LOSING: tighten every step. + new_trail = trail * loss_decay; + } else if (unrealized > init_r) { + // WINNING beyond initial R: ratchet trail to track profit. + // Trail = max(current_trail, profit × win_factor). + // Never shrinks on winners — only grows. + const float profit_trail = unrealized * win_factor; + if (profit_trail > new_trail) { + new_trail = profit_trail; + } + } + // NEUTRAL (0 to init_r): trail unchanged — proving zone. + + // Floor at trail_min. + new_trail = fmaxf(new_trail, trail_min); + + unit_trail_distance[idx] = new_trail; +} diff --git a/crates/ml-alpha/cuda/rl_session_risk_check.cu b/crates/ml-alpha/cuda/rl_session_risk_check.cu new file mode 100644 index 000000000..19c107bfb --- /dev/null +++ b/crates/ml-alpha/cuda/rl_session_risk_check.cu @@ -0,0 +1,64 @@ +// rl_session_risk_check.cu — session-level risk circuit breaker. +// +// Tracks cumulative raw PnL over a rolling window. When cumulative +// loss exceeds the ISV-driven threshold, overrides ALL opening +// actions to Hold until the window advances enough to recover. +// +// "A pro trader who's down their daily limit STOPS trading." +// +// The window is implemented as an EMA of realized PnL with slow +// decay — effectively a running average of recent performance. +// When the EMA drops below the threshold, the breaker fires. +// +// Per `feedback_no_atomicadd`: single-thread kernel. +// Per `feedback_cpu_is_read_only`: pure device-side. + +#include + +#define ACTION_HOLD 2 +#define RL_SESSION_PNL_EMA_INDEX 539 +#define RL_SESSION_LOSS_LIMIT_INDEX 540 +#define RL_SESSION_EMA_ALPHA_INDEX 541 + +extern "C" __global__ void rl_session_risk_check( + int* __restrict__ actions, // [B] IN/OUT + const float* __restrict__ rewards, // [B] raw rewards this step + const float* __restrict__ dones, // [B] + float* __restrict__ isv, + int b_size +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + + const float alpha = isv[RL_SESSION_EMA_ALPHA_INDEX]; + const float limit = isv[RL_SESSION_LOSS_LIMIT_INDEX]; + float ema = isv[RL_SESSION_PNL_EMA_INDEX]; + + // Update EMA from this step's realized PnL (done events only). + float step_pnl = 0.0f; + int done_count = 0; + for (int b = 0; b < b_size; ++b) { + if (dones[b] > 0.5f) { + step_pnl += rewards[b]; + done_count++; + } + } + if (done_count > 0) { + const float avg_pnl = step_pnl / (float)done_count; + if (ema == 0.0f) { + ema = avg_pnl; + } else { + ema = (1.0f - alpha) * ema + alpha * avg_pnl; + } + isv[RL_SESSION_PNL_EMA_INDEX] = ema; + } + + // If session PnL EMA below limit, block ALL opening actions. + if (ema < limit) { + for (int b = 0; b < b_size; ++b) { + const int a = actions[b]; + if (a == 0 || a == 1 || a == 5 || a == 6) { + actions[b] = ACTION_HOLD; + } + } + } +} diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index c353ed8bc..91becdfb9 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -931,6 +931,26 @@ pub const RL_HOLD_BONUS_INDEX: usize = 535; /// regardless (safety overrides patience). Default: 20. pub const RL_MIN_HOLD_STEPS_INDEX: usize = 536; +/// Asymmetric trail: per-step decay on losing positions. Default 0.995 +/// (halves trail in ~139 steps = ~35 seconds). +pub const RL_ASYM_LOSS_DECAY_RATE_INDEX: usize = 537; + +/// Asymmetric trail: profit-tracking factor for winners. Trail ratchets +/// to max(current, unrealized × factor). Default 0.5 (trail = 50% of +/// open profit). Only fires when unrealized > initial_r. +pub const RL_ASYM_WIN_TRAIL_FACTOR_INDEX: usize = 538; + +/// Session risk: EMA of realized PnL (rolling performance tracker). +pub const RL_SESSION_PNL_EMA_INDEX: usize = 539; + +/// Session risk: loss limit. When PnL EMA drops below this value, +/// all opening actions are blocked. Default -50.0 (stop trading +/// after losing $50 on average per recent close). +pub const RL_SESSION_LOSS_LIMIT_INDEX: usize = 540; + +/// Session risk: EMA alpha for the PnL tracker. Default 0.02. +pub const RL_SESSION_EMA_ALPHA_INDEX: usize = 541; + /// 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 = 537; +pub const RL_SLOTS_END: usize = 542; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index 1a3bb2b5d..dacc5feaa 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -204,6 +204,10 @@ const RL_GATE_THRESHOLD_CONTROLLER_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_gate_threshold_controller.cubin")); const RL_REWARD_SHAPING_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_reward_shaping.cubin")); +const RL_ASYMMETRIC_TRAIL_DECAY_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_asymmetric_trail_decay.cubin")); +const RL_SESSION_RISK_CHECK_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_session_risk_check.cubin")); const RL_MIN_HOLD_CHECK_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/rl_min_hold_check.cubin")); const RL_TRADE_CONTEXT_UPDATE_CUBIN: &[u8] = @@ -469,6 +473,10 @@ pub struct IntegratedTrainer { rl_gate_threshold_controller_fn: CudaFunction, _rl_reward_shaping_module: Arc, rl_reward_shaping_fn: CudaFunction, + _rl_asymmetric_trail_decay_module: Arc, + rl_asymmetric_trail_decay_fn: CudaFunction, + _rl_session_risk_check_module: Arc, + rl_session_risk_check_fn: CudaFunction, _rl_min_hold_check_module: Arc, rl_min_hold_check_fn: CudaFunction, _rl_trade_context_update_module: Arc, @@ -994,6 +1002,18 @@ impl IntegratedTrainer { let rl_reward_shaping_fn = rl_reward_shaping_module .load_function("rl_reward_shaping") .context("load rl_reward_shaping")?; + let rl_asymmetric_trail_decay_module = ctx + .load_cubin(RL_ASYMMETRIC_TRAIL_DECAY_CUBIN.to_vec()) + .context("load rl_asymmetric_trail_decay cubin")?; + let rl_asymmetric_trail_decay_fn = rl_asymmetric_trail_decay_module + .load_function("rl_asymmetric_trail_decay") + .context("load rl_asymmetric_trail_decay")?; + let rl_session_risk_check_module = ctx + .load_cubin(RL_SESSION_RISK_CHECK_CUBIN.to_vec()) + .context("load rl_session_risk_check cubin")?; + let rl_session_risk_check_fn = rl_session_risk_check_module + .load_function("rl_session_risk_check") + .context("load rl_session_risk_check")?; let rl_min_hold_check_module = ctx .load_cubin(RL_MIN_HOLD_CHECK_CUBIN.to_vec()) .context("load rl_min_hold_check cubin")?; @@ -1429,6 +1449,10 @@ impl IntegratedTrainer { rl_gate_threshold_controller_fn, _rl_reward_shaping_module: rl_reward_shaping_module, rl_reward_shaping_fn, + _rl_asymmetric_trail_decay_module: rl_asymmetric_trail_decay_module, + rl_asymmetric_trail_decay_fn, + _rl_session_risk_check_module: rl_session_risk_check_module, + rl_session_risk_check_fn, _rl_min_hold_check_module: rl_min_hold_check_module, rl_min_hold_check_fn, _rl_trade_context_update_module: rl_trade_context_update_module, @@ -1576,7 +1600,7 @@ impl IntegratedTrainer { // (slot, value) pair — pure device write, no HtoD per // `feedback_no_htod_htoh_only_mapped_pinned`. { - let isv_constants: [(usize, f32); 73] = [ + let isv_constants: [(usize, f32); 77] = [ // Static seeds for the adaptive reward-clamp controller — // these are the initial values that // `rl_reward_clamp_controller` will replace once it @@ -1656,6 +1680,10 @@ impl IntegratedTrainer { (crate::rl::isv_slots::RL_SHORT_HOLD_PENALTY_INDEX, 0.5), (crate::rl::isv_slots::RL_HOLD_BONUS_INDEX, 2.0), (crate::rl::isv_slots::RL_MIN_HOLD_STEPS_INDEX, 100.0), + (crate::rl::isv_slots::RL_ASYM_LOSS_DECAY_RATE_INDEX, 0.995), + (crate::rl::isv_slots::RL_ASYM_WIN_TRAIL_FACTOR_INDEX, 0.5), + (crate::rl::isv_slots::RL_SESSION_LOSS_LIMIT_INDEX, -50.0), + (crate::rl::isv_slots::RL_SESSION_EMA_ALPHA_INDEX, 0.02), (crate::rl::isv_slots::RL_MULTIRES_HORIZON_1_INDEX, 1.0), (crate::rl::isv_slots::RL_MULTIRES_HORIZON_2_INDEX, 10.0), (crate::rl::isv_slots::RL_MULTIRES_HORIZON_3_INDEX, 600.0), @@ -3888,6 +3916,29 @@ impl IntegratedTrainer { let pos_bytes_i = lobsim.pos_bytes() as i32; let b_size_i = b_size as i32; + // Session risk circuit breaker — block opening actions when + // cumulative PnL EMA drops below threshold. + { + let cfg = LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }; + let b_size_i = b_size as i32; + let mut launch = self.stream.launch_builder(&self.rl_session_risk_check_fn); + launch + .arg(&mut self.actions_d) + .arg(&self.rewards_d) + .arg(&self.dones_d) + .arg(&self.isv_d) + .arg(&b_size_i); + unsafe { + launch + .launch(cfg) + .context("rl_session_risk_check launch")?; + } + } + // Min hold check — override closing actions to Hold when // hold time < ISV-driven minimum. Trail stops still fire // (they run AFTER this, safety overrides patience). @@ -3913,6 +3964,36 @@ impl IntegratedTrainer { } } + // Asymmetric trail decay — auto-tighten losers, widen winners. + // Runs BEFORE trail_mutate so the structural decay is applied + // first, then agent's a7/a8 fine-tuning on top. + { + let bid_px_d = lobsim.bid_px_d(); + let ask_px_d = lobsim.ask_px_d(); + let cfg = LaunchConfig { + grid_dim: (b_size as u32, 1, 1), + block_dim: (4, 1, 1), + shared_mem_bytes: 0, + }; + let mut launch = + self.stream.launch_builder(&self.rl_asymmetric_trail_decay_fn); + launch + .arg(&mut self.unit_trail_distance_d) + .arg(&self.unit_active_d) + .arg(&self.unit_entry_price_d) + .arg(&self.unit_initial_r_d) + .arg(&self.unit_lots_d) + .arg(bid_px_d) + .arg(ask_px_d) + .arg(&self.isv_d) + .arg(&b_size_i); + unsafe { + launch + .launch(cfg) + .context("rl_asymmetric_trail_decay launch")?; + } + } + // Trail mutate (a7/a8) — BEFORE actions_to_market_targets. // Mutates unit_trail_distance for active units if the chosen action // is TrailTighten/Loosen. Non-trail actions pass through.