diff --git a/Cargo.lock b/Cargo.lock index 7fe0ee372..ea912f7f1 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -6090,6 +6090,7 @@ dependencies = [ "approx", "arrow 56.2.0", "bincode", + "bytemuck", "clap", "cudarc", "data", diff --git a/crates/ml-alpha/build.rs b/crates/ml-alpha/build.rs index 9695c3f65..931568813 100644 --- a/crates/ml-alpha/build.rs +++ b/crates/ml-alpha/build.rs @@ -73,6 +73,7 @@ const KERNELS: &[&str] = &[ "rl_kl_reference_grad", "rl_q_pi_distill_grad", // audit 2026-05-24 vj5f6 followup: KL(softmax(E_Q/τ) || π_new) gradient ADDED to pi_grad_logits — couples Q's improved C51 calibration to action selection (was decoupled per Option B) "action_entropy_per_step", // POST-gate action entropy EMA for SAC α/τ co-tuning + "rl_drawdown_stop", // per-step drawdown penalty + hard stop-loss "rl_q_distill_lambda_controller",// audit 2026-05-24 rljzl followup: adaptive λ_distill via Schulman bounded step on KL_EMA vs target "rl_unit_state_update", // SP20 P1+P5 audit fix: per-unit trade state machine — detects open/close/reverse transitions, sets up unit slot 0 entry+trail "rl_trail_mutate", // SP20 P1+P5 audit fix (a7/a8 dead): TrailTighten/TrailLoosen mutate unit_trail_distance bounded MIN/MAX, symmetric reciprocal adjust diff --git a/crates/ml-alpha/cuda/rl_confidence_gate.cu b/crates/ml-alpha/cuda/rl_confidence_gate.cu index 645d7d32c..158b17593 100644 --- a/crates/ml-alpha/cuda/rl_confidence_gate.cu +++ b/crates/ml-alpha/cuda/rl_confidence_gate.cu @@ -65,6 +65,7 @@ extern "C" __global__ void rl_confidence_gate( const int action = actions[b]; if (action == ACTION_HOLD) return; + if (action == 3 || action == 4) return; // FlatFromLong/Short: never gate exits const float threshold = isv[RL_CONF_GATE_THRESHOLD_INDEX]; const float lambda = isv[RL_CONF_GATE_LAMBDA_INDEX]; diff --git a/crates/ml-alpha/cuda/rl_drawdown_stop.cu b/crates/ml-alpha/cuda/rl_drawdown_stop.cu new file mode 100644 index 000000000..c3e322c9d --- /dev/null +++ b/crates/ml-alpha/cuda/rl_drawdown_stop.cu @@ -0,0 +1,48 @@ +// rl_drawdown_stop.cu — per-step drawdown penalty + hard stop-loss. +// +// Runs AFTER rl_trade_context_update (which computes unrealized_r). +// Reads unrealized_r from trade_context_d and applies: +// +// 1. Per-step drawdown penalty: adds min(0, unrealized_r) × PENALTY_RATE +// to rewards[b]. Creates continuous exit gradient on losing positions. +// Penalty-only (never positive) = no exposure incentive. +// +// 2. Hard stop-loss: when unrealized_r < -STOP_THRESHOLD, force-close +// by setting dones[b] = 1. The realized PnL at this point becomes +// the final trade reward. Caps max loss per trade. +// +// Both thresholds are ISV-driven per feedback_isv_for_adaptive_bounds. +// Grid: ceil(B/32), Block: 32. + +#define RL_DRAWDOWN_PENALTY_RATE_INDEX 586 +#define RL_STOP_LOSS_THRESHOLD_INDEX 587 +#define TRADE_CONTEXT_STRIDE 5 +#define UNREALIZED_R_OFFSET 1 + +extern "C" __global__ void rl_drawdown_stop( + const float* __restrict__ trade_context, // [B, TRADE_CONTEXT_STRIDE] + float* __restrict__ rewards, // [B] IN/OUT + float* __restrict__ dones, // [B] IN/OUT + float* __restrict__ isv, + int b_size +) { + const int b = blockIdx.x * blockDim.x + threadIdx.x; + if (b >= b_size) return; + + const float unrealized_r = trade_context[b * TRADE_CONTEXT_STRIDE + UNREALIZED_R_OFFSET]; + const float penalty_rate = isv[RL_DRAWDOWN_PENALTY_RATE_INDEX]; + const float stop_thresh = isv[RL_STOP_LOSS_THRESHOLD_INDEX]; + + // Per-step drawdown penalty: only fires when losing (unrealized_r < 0). + // min(0, x) * rate is always <= 0. + if (unrealized_r < 0.0f) { + rewards[b] += unrealized_r * penalty_rate; + } + + // Hard stop-loss: force-close when drawdown exceeds threshold. + // unrealized_r is normalized by initial_r (trade-level vol), so + // threshold is in units of "initial risk multiples." + if (unrealized_r < -stop_thresh && dones[b] < 0.5f) { + dones[b] = 1.0f; + } +} diff --git a/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu b/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu index d9bfe7daa..adecc96a4 100644 --- a/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu +++ b/crates/ml-alpha/cuda/rl_reward_clamp_controller.cu @@ -230,6 +230,15 @@ extern "C" __global__ void rl_reward_clamp_controller( // Suppress unused warnings on diagnostic-only state. (void) margin; + // Adaptive LOSS clamp from observed neg/pos ratio. WIN stays + // structural at 1.0; LOSS adapts to actual loss magnitude to + // give Q accurate resolution without over-allocating to rare tails. + 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; + } + // ── 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/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs index 48eb976fa..c035e2c94 100644 --- a/crates/ml-alpha/src/rl/isv_slots.rs +++ b/crates/ml-alpha/src/rl/isv_slots.rs @@ -1111,5 +1111,16 @@ 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; +/// Per-step drawdown penalty rate. Applied as min(0, unrealized_r) × rate +/// to the reward each step a position is losing. Penalty-only (never +/// positive) creates continuous exit gradient without exposure incentive. +/// Bootstrap: 0.01 (1% of unrealized_r magnitude per step). +pub const RL_DRAWDOWN_PENALTY_RATE_INDEX: usize = 586; + +/// Hard stop-loss threshold in initial-risk multiples. When unrealized_r +/// drops below -threshold, the position force-closes (done=1). Caps max +/// loss per trade. Bootstrap: 2.0 (2× initial risk). +pub const RL_STOP_LOSS_THRESHOLD_INDEX: usize = 587; + /// Last RL-allocated slot index (exclusive). -pub const RL_SLOTS_END: usize = 585; +pub const RL_SLOTS_END: usize = 588; diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index a04b284f2..3cbf24ab5 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -150,6 +150,8 @@ const COMPUTE_ADVANTAGE_RETURN_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/compute_advantage_return.cubin")); const ACTION_ENTROPY_PER_STEP_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/action_entropy_per_step.cubin")); +const RL_DRAWDOWN_STOP_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/rl_drawdown_stop.cubin")); // Phase R4: GPU-resident action sampling kernels. Replace the host // Thompson + host argmax + host log-softmax loops the flawed Phase F @@ -566,6 +568,8 @@ pub struct IntegratedTrainer { compute_advantage_return_fn: CudaFunction, _action_entropy_per_step_module: Arc, action_entropy_per_step_fn: CudaFunction, + _rl_drawdown_stop_module: Arc, + rl_drawdown_stop_fn: CudaFunction, // ── Phase R4: GPU action sampling ───────────────────────────────── _rl_action_kernel_module: Arc, @@ -1327,6 +1331,12 @@ impl IntegratedTrainer { let action_entropy_per_step_module = ctx .load_cubin(ACTION_ENTROPY_PER_STEP_CUBIN.to_vec()) .context("load action_entropy_per_step cubin")?; + let rl_drawdown_stop_module = ctx + .load_cubin(RL_DRAWDOWN_STOP_CUBIN.to_vec()) + .context("load rl_drawdown_stop cubin")?; + let rl_drawdown_stop_fn = rl_drawdown_stop_module + .load_function("rl_drawdown_stop") + .context("load rl_drawdown_stop")?; let action_entropy_per_step_fn = action_entropy_per_step_module .load_function("action_entropy_per_step") .context("load action_entropy_per_step")?; @@ -2323,6 +2333,8 @@ impl IntegratedTrainer { ema_update_per_step_fn, _action_entropy_per_step_module: action_entropy_per_step_module, action_entropy_per_step_fn, + _rl_drawdown_stop_module: rl_drawdown_stop_module, + rl_drawdown_stop_fn, _compute_advantage_return_module: compute_advantage_return_module, compute_advantage_return_fn, _rl_action_kernel_module: rl_action_kernel_module, @@ -2673,7 +2685,7 @@ impl IntegratedTrainer { // (slot, value) pair — pure device write, no HtoD per // `feedback_no_htod_htoh_only_mapped_pinned`. { - let isv_constants: [(usize, f32); 111] = [ + let isv_constants: [(usize, f32); 113] = [ // Static seeds for the adaptive reward-clamp controller — // these are the initial values that // `rl_reward_clamp_controller` will replace once it @@ -2840,6 +2852,8 @@ impl IntegratedTrainer { (crate::rl::isv_slots::RL_SAC_ALPHA_INDEX, 0.01), (crate::rl::isv_slots::RL_SAC_ENTROPY_TARGET_INDEX, 1.68), (crate::rl::isv_slots::RL_CONF_GATE_MAX_HOLD_FRAC_INDEX, 0.85), + (crate::rl::isv_slots::RL_DRAWDOWN_PENALTY_RATE_INDEX, 0.01), + (crate::rl::isv_slots::RL_STOP_LOSS_THRESHOLD_INDEX, 2.0), ]; for (slot, value) in isv_constants.iter() { let slot_i32 = *slot as i32; @@ -6011,6 +6025,29 @@ impl IntegratedTrainer { } } + // Per-step drawdown penalty + hard stop-loss. Reads unrealized_r + // from trade_context_d (just written above). Modifies rewards_d + // and dones_d in place before apply_reward_scale. + { + let grid_x = ((b_size as u32) + 31) / 32; + let b_size_i = b_size as i32; + let mut args = RawArgs::new(); + args.push_ptr(self.trade_context_d.raw_ptr()); + args.push_ptr(self.rewards_d.raw_ptr()); + args.push_ptr(self.dones_d.raw_ptr()); + args.push_ptr(self.isv_dev_ptr); + args.push_i32(b_size_i); + let mut ptrs = args.build_arg_ptrs(); + unsafe { + raw_launch( + self.rl_drawdown_stop_fn.cu_function(), + (grid_x.max(1), 1, 1), (32, 1, 1), 0, + self.raw_stream, + &mut ptrs[..args.len()], + ).map_err(|e| anyhow::anyhow!("rl_drawdown_stop: {:?}", e))?; + } + } + // Multi-resolution streaming features — time-weighted EMA at 3 horizons. { let bid_px_d = lobsim.bid_px_d(); diff --git a/crates/ml-backtesting/src/sim/mod.rs b/crates/ml-backtesting/src/sim/mod.rs index 38c0dedbb..36bb307bf 100644 --- a/crates/ml-backtesting/src/sim/mod.rs +++ b/crates/ml-backtesting/src/sim/mod.rs @@ -1143,7 +1143,6 @@ impl LobSimCuda { .arg(&n) .launch(cfg)?; } - self.stream.synchronize()?; Ok(()) }