feat(rl): asymmetric trail + session risk — structural P&L edge
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 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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
|
||||
|
||||
94
crates/ml-alpha/cuda/rl_asymmetric_trail_decay.cu
Normal file
94
crates/ml-alpha/cuda/rl_asymmetric_trail_decay.cu
Normal file
@@ -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 <stdint.h>
|
||||
|
||||
#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;
|
||||
}
|
||||
64
crates/ml-alpha/cuda/rl_session_risk_check.cu
Normal file
64
crates/ml-alpha/cuda/rl_session_risk_check.cu
Normal file
@@ -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 <stdint.h>
|
||||
|
||||
#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;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -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;
|
||||
|
||||
@@ -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<CudaModule>,
|
||||
rl_reward_shaping_fn: CudaFunction,
|
||||
_rl_asymmetric_trail_decay_module: Arc<CudaModule>,
|
||||
rl_asymmetric_trail_decay_fn: CudaFunction,
|
||||
_rl_session_risk_check_module: Arc<CudaModule>,
|
||||
rl_session_risk_check_fn: CudaFunction,
|
||||
_rl_min_hold_check_module: Arc<CudaModule>,
|
||||
rl_min_hold_check_fn: CudaFunction,
|
||||
_rl_trade_context_update_module: Arc<CudaModule>,
|
||||
@@ -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.
|
||||
|
||||
Reference in New Issue
Block a user