From 4aca6330f03de5a4bae467ca314a9038ee0dfa56 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Wed, 20 May 2026 00:01:06 +0200 Subject: [PATCH] feat(ml-backtesting): hard-SL trigger in stop_check_isv sl_distance = max(pnl_ema_loss, atr_mid_ema) per pearl_blend_formulas_must_have_permanent_floor. Fires force-flat (3, 0) when unrealized_pl_per_lot drops below -sl_distance. Multi-horizon mask averaging + trail trigger land in Tasks 5-6. Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml-backtesting/cuda/decision_policy.cu | 47 ++++++++++++++-- .../ml-backtesting/tests/stop_controller.rs | 55 +++++++++++++++++++ 2 files changed, 96 insertions(+), 6 deletions(-) diff --git a/crates/ml-backtesting/cuda/decision_policy.cu b/crates/ml-backtesting/cuda/decision_policy.cu index 1d7c71ec5..2e8f460eb 100644 --- a/crates/ml-backtesting/cuda/decision_policy.cu +++ b/crates/ml-backtesting/cuda/decision_policy.cu @@ -48,13 +48,48 @@ __device__ static int stop_check_isv( const Book* __restrict__ books, int* __restrict__ market_targets ) { - if (positions[b].position_lots == 0) { - return 0; // flat — entry path handles this + const Pos& pos = positions[b]; + if (pos.position_lots == 0) return 0; + + // Spec §5 controller math — SL only in this task; trail in Task 5. + const float atr = atr_mid_ema[b]; + const unsigned int mask = open_horizon_masks[b]; + + const IsvKellyState* isv = reinterpret_cast( + isv_kelly_base + (size_t)b * N_HORIZONS * sizeof(IsvKellyState) + ); + float ema_loss = 0.0f; + int n_open = 0; + #pragma unroll + for (int h = 0; h < N_HORIZONS; ++h) { + if (mask & (1u << h)) { + ema_loss += isv[h].pnl_ema_loss; + n_open += 1; + } + } + if (n_open > 0) ema_loss /= (float)n_open; + + // pearl_blend_formulas_must_have_permanent_floor — max, not blend. + const float sl_distance = fmaxf(ema_loss, atr); + + const bool is_long = pos.position_lots > 0; + const Book& bk = books[b]; + const float close_px = is_long ? bk.bid_px[0] : bk.ask_px[0]; + const float entry_px = pos.vwap_entry; + const float unrealized_pl_per_lot = + is_long ? (close_px - entry_px) : (entry_px - close_px); + + const bool sl_fired = unrealized_pl_per_lot <= -sl_distance; + + // trail_hwm is unused in this task — Task 5 will wire it. + (void)trail_hwm; + + if (sl_fired) { + // §7 force-flat encoding. + market_targets[b * 2 + 0] = 3; + market_targets[b * 2 + 1] = 0; + return 1; } - // Stop trigger logic added in Tasks 4-6. - // Mark args as referenced so -Wunused-parameter stays quiet across cubins: - (void)isv_kelly_base; (void)open_horizon_masks; (void)atr_mid_ema; - (void)trail_hwm; (void)books; (void)market_targets; return 0; } diff --git a/crates/ml-backtesting/tests/stop_controller.rs b/crates/ml-backtesting/tests/stop_controller.rs index 384bdb01a..60aee70d7 100644 --- a/crates/ml-backtesting/tests/stop_controller.rs +++ b/crates/ml-backtesting/tests/stop_controller.rs @@ -96,3 +96,58 @@ fn stop_check_skipped_when_flat() -> Result<()> { assert!(size >= 1, "entry should fire under p=0.8 + flat (size={size})"); Ok(()) } + +use ml_backtesting::policy::IsvKellyStateHost; + +#[test] +#[ignore = "requires CUDA"] +fn sl_fires_when_unrealized_breaks_distance() -> Result<()> { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { eprintln!("skipping: cuda unavailable ({e})"); return Ok(()); } + }; + let mut sim = LobSimCuda::new(1, &dev)?; + + // Seed isv: 5 horizons with pnl_ema_loss=2.0 → sl_distance=2.0. + // pnl_ema_win=0.0 forces cold-start branch in decision kernel (kelly_frac → + // floor=0.20, cap_lots → max_lots=5). With alpha=0.95 this gives + // lots=round(0.9*0.20*5)=round(0.9)=1 — position opens. + let seeded: [IsvKellyStateHost; 5] = std::array::from_fn(|_| IsvKellyStateHost { + pnl_ema_win: 0.0, + pnl_ema_loss: 2.0, + win_rate_ema: 0.5, + n_trades_seen: 0, + realised_return_var: 0.0, + recent_sharpe: 0.5, + }); + sim.write_isv_kelly(0, &seeded)?; + + // Two snapshots so atr_mid_ema becomes nonzero (small). + let (bp, bs, ap, az) = level_book(5500.0, 0.25); + sim.apply_snapshot(&bp, &bs, &ap, &az)?; + let (bp2, bs2, ap2, az2) = level_book(5500.1, 0.25); + sim.apply_snapshot(&bp2, &bs2, &ap2, &az2)?; + + // Open a long position via strong-alpha + step_decision + manual fill. + sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + sim.step_decision_with_latency(0, &cfg_default(1))?; + let mut ts: u64 = 1_000_000; + for _ in 0..3 { + sim.step_resting_orders(ts, 0.0)?; + ts += 1_000_000; + } + let pos_open = sim.read_pos(0)?; + assert!(pos_open.position_lots > 0, + "test setup: long position should open, got {}", pos_open.position_lots); + + // Drive mid down by 3.0 (>= sl_distance=2.0). Stop must fire force-flat. + let (bp3, bs3, ap3, az3) = level_book(5497.0, 0.25); + sim.apply_snapshot(&bp3, &bs3, &ap3, &az3)?; + sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + sim.step_decision_with_latency(ts, &cfg_default(1))?; + + let (side, size) = sim.read_market_target(0)?; + assert_eq!(side, 3, "SL must fire with force-flat (3, 0); got side={side}"); + assert_eq!(size, 0, "force-flat encoding has abs_sz=0; got size={size}"); + Ok(()) +}