From 7e1995d0adf51a5bdd7fad8a9a44f60b00be95c5 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Wed, 20 May 2026 00:06:23 +0200 Subject: [PATCH] feat(ml-backtesting): trail-TP trigger with HWM ratchet in stop_check_isv trail_distance = max(pnl_ema_win, atr_mid_ema). HWM ratchets up via fmaxf each event while position open; trail fires when HWM clears the arming threshold AND unrealized drops by trail_distance from HWM. Same force-flat (3, 0) write as SL. Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml-backtesting/cuda/decision_policy.cu | 19 ++++-- .../ml-backtesting/tests/stop_controller.rs | 68 +++++++++++++++++++ 2 files changed, 82 insertions(+), 5 deletions(-) diff --git a/crates/ml-backtesting/cuda/decision_policy.cu b/crates/ml-backtesting/cuda/decision_policy.cu index 2e8f460eb..7f0cc6036 100644 --- a/crates/ml-backtesting/cuda/decision_policy.cu +++ b/crates/ml-backtesting/cuda/decision_policy.cu @@ -59,18 +59,24 @@ __device__ static int stop_check_isv( isv_kelly_base + (size_t)b * N_HORIZONS * sizeof(IsvKellyState) ); float ema_loss = 0.0f; + float ema_win = 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; + ema_win += isv[h].pnl_ema_win; n_open += 1; } } - if (n_open > 0) ema_loss /= (float)n_open; + if (n_open > 0) { + ema_loss /= (float)n_open; + ema_win /= (float)n_open; + } // pearl_blend_formulas_must_have_permanent_floor — max, not blend. - const float sl_distance = fmaxf(ema_loss, atr); + const float sl_distance = fmaxf(ema_loss, atr); + const float trail_distance = fmaxf(ema_win, atr); const bool is_long = pos.position_lots > 0; const Book& bk = books[b]; @@ -81,10 +87,13 @@ __device__ static int stop_check_isv( 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; + const float new_hwm = fmaxf(trail_hwm[b], unrealized_pl_per_lot); + trail_hwm[b] = new_hwm; + const bool trail_fired = + (new_hwm > trail_distance) + && (unrealized_pl_per_lot <= new_hwm - trail_distance); - if (sl_fired) { + if (sl_fired || trail_fired) { // §7 force-flat encoding. market_targets[b * 2 + 0] = 3; market_targets[b * 2 + 1] = 0; diff --git a/crates/ml-backtesting/tests/stop_controller.rs b/crates/ml-backtesting/tests/stop_controller.rs index 60aee70d7..d305a9673 100644 --- a/crates/ml-backtesting/tests/stop_controller.rs +++ b/crates/ml-backtesting/tests/stop_controller.rs @@ -151,3 +151,71 @@ fn sl_fires_when_unrealized_breaks_distance() -> Result<()> { assert_eq!(size, 0, "force-flat encoding has abs_sz=0; got size={size}"); Ok(()) } + +#[test] +#[ignore = "requires CUDA"] +fn trail_arms_then_fires() -> 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)?; + + // Cold-start ISV (forces kelly_frac_floor + max_lots cap) so position opens. + // pnl_ema_loss large enough that SL won't fire during the +5pt walk. + let cold_start: [IsvKellyStateHost; 5] = std::array::from_fn(|_| IsvKellyStateHost { + pnl_ema_win: 0.0, + pnl_ema_loss: 10.0, // big SL so trail fires first + win_rate_ema: 0.0, + n_trades_seen: 0, + realised_return_var: 0.0, + recent_sharpe: 0.0, + }); + sim.write_isv_kelly(0, &cold_start)?; + + 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 long. + 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; } + assert!(sim.read_pos(0)?.position_lots > 0, "setup: long opened"); + + // After opening, switch ISV to one with pnl_ema_win=1.5 (trail_distance target). + // The trail check reads this fresh state on subsequent decision steps. + let trail_state: [IsvKellyStateHost; 5] = std::array::from_fn(|_| IsvKellyStateHost { + pnl_ema_win: 1.5, + pnl_ema_loss: 10.0, + win_rate_ema: 0.5, + n_trades_seen: 20, + realised_return_var: 0.5, + recent_sharpe: 0.5, + }); + sim.write_isv_kelly(0, &trail_state)?; + + // Walk mid up to ratchet trail HWM past trail_distance=1.5. + for &m in &[5501.0_f32, 5502.0, 5503.0, 5504.0, 5505.0] { + let (b, bs, a, asz) = level_book(m, 0.25); + sim.apply_snapshot(&b, &bs, &a, &asz)?; + sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + sim.step_decision_with_latency(ts, &cfg_default(1))?; + ts += 1_000_000; + } + let hwm_armed = sim.read_trail_hwm(0)?; + assert!(hwm_armed > 1.5, "HWM should ratchet past trail_distance=1.5, got {hwm_armed}"); + + // Drop mid by trail_distance+ε from peak to fire trail. + let (b, bs, a, asz) = level_book(5502.0, 0.25); // drop from peak 5505 to 5502 = 3.0 + sim.apply_snapshot(&b, &bs, &a, &asz)?; + 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, "trail must fire force-flat; got side={side}"); + assert_eq!(size, 0); + Ok(()) +}