diff --git a/crates/ml-backtesting/cuda/book_update.cu b/crates/ml-backtesting/cuda/book_update.cu index 8d0421974..5e12f373b 100644 --- a/crates/ml-backtesting/cuda/book_update.cu +++ b/crates/ml-backtesting/cuda/book_update.cu @@ -35,7 +35,8 @@ extern "C" __global__ void book_update_apply_snapshot( const bool top_ok = isfinite(bid_px_in[0]) && isfinite(ask_px_in[0]) && isfinite(bid_sz_in[0]) && isfinite(ask_sz_in[0]) && bid_sz_in[0] > 0.0f && ask_sz_in[0] > 0.0f - && bid_px_in[0] > 0.0f && ask_px_in[0] > 0.0f; + && bid_px_in[0] > 0.0f && ask_px_in[0] > 0.0f + && bid_px_in[0] < 1.0e8f && ask_px_in[0] < 1.0e8f; if (!top_ok) { if (tid == 0) snapshots_skipped[b] += 1u; return; // Leave book/prev_mid/atr_mid_ema unchanged. @@ -52,8 +53,8 @@ extern "C" __global__ void book_update_apply_snapshot( const float bs = bid_sz_in[tid]; const float ap = ask_px_in[tid]; const float az = ask_sz_in[tid]; - const bool bid_ok = isfinite(bp) && isfinite(bs) && bs > 0.0f; - const bool ask_ok = isfinite(ap) && isfinite(az) && az > 0.0f; + const bool bid_ok = isfinite(bp) && isfinite(bs) && bs > 0.0f && bp > 0.0f && bp < 1.0e8f; + const bool ask_ok = isfinite(ap) && isfinite(az) && az > 0.0f && ap > 0.0f && ap < 1.0e8f; bk.bid_px[tid] = bid_ok ? bp : 0.0f; bk.bid_sz[tid] = bid_ok ? bs : 0.0f; bk.ask_px[tid] = ask_ok ? ap : 0.0f; diff --git a/crates/ml-backtesting/cuda/resting_orders.cu b/crates/ml-backtesting/cuda/resting_orders.cu index 543e00903..9ddd1b49f 100644 --- a/crates/ml-backtesting/cuda/resting_orders.cu +++ b/crates/ml-backtesting/cuda/resting_orders.cu @@ -164,6 +164,8 @@ extern "C" __global__ void resting_orders_step( // P4: per-fill cost integration (per-backtest arrays). const float* __restrict__ cost_per_lot_per_side_per_b, float* __restrict__ total_fees_per_b, + // Fix 3: session-gap force-close. Tracks last event ts per backtest. + unsigned long long* __restrict__ last_event_ts, // [n_backtests] int n_backtests ) { int b = blockIdx.x; @@ -173,6 +175,25 @@ extern "C" __global__ void resting_orders_step( Orders& orders = *reinterpret_cast(orders_base + (size_t)b * orders_bytes); Pos& pos = *reinterpret_cast(pos_base + (size_t)b * pos_bytes); + // Session-gap force-close: when event timestamps jump > 1 hour (weekend + // halt, session boundary, gap-fill), max_hold can't fire because no + // intervening events advance current_ts. Force-close all open positions + // before processing the new event's fills. + const unsigned long long SESSION_GAP_NS = 3600000000000ull; // 1 hour + const unsigned long long last_ts = last_event_ts[b]; + if (last_ts > 0ull && current_ts_ns > last_ts + && (current_ts_ns - last_ts) >= SESSION_GAP_NS + && pos.position_lots != 0) + { + // Force-flat: zero position_lots directly. pnl_track_step (called by + // step_resting_orders after this kernel) will detect the prev!=0 && now==0 + // transition via open_trade_state scratch and emit a close TradeRecord. + // No synthetic P&L is added — records show realised_pnl from whatever + // was accumulated before the gap (honest: cannot fill across a halt). + pos.position_lots = 0; + } + last_event_ts[b] = current_ts_ns; + // (1) Promote in-flight limits to resting; (2) queue-decay; (3) marketability check. const float trade_abs = (trade_signed_vol < 0.0f) ? -trade_signed_vol : trade_signed_vol; for (int i = 0; i < MAX_LIMITS; ++i) { diff --git a/crates/ml-backtesting/src/sim/mod.rs b/crates/ml-backtesting/src/sim/mod.rs index 05556f7ee..cd8af6630 100644 --- a/crates/ml-backtesting/src/sim/mod.rs +++ b/crates/ml-backtesting/src/sim/mod.rs @@ -121,6 +121,10 @@ pub struct LobSimCuda { // Bug C-b (smoke v74v4 2026-05-20): corrupt-top-of-book skip counter. snapshots_skipped_d: CudaSlice, // [n_backtests] — incremented each time apply_snapshot skips a corrupt snapshot + + // Fix 3 (smoke stx9p 2026-05-20): session-gap force-close. Tracks the + // previous event ts per backtest so resting_orders_step can detect halts. + last_event_ts_d: CudaSlice, // [n_backtests] — 0 = sentinel (first event) } impl LobSimCuda { @@ -289,6 +293,9 @@ impl LobSimCuda { let snapshots_skipped_d = stream .alloc_zeros::(n_backtests) .context("alloc snapshots_skipped_d")?; + let last_event_ts_d = stream + .alloc_zeros::(n_backtests) + .context("alloc last_event_ts_d")?; Ok(Self { n_backtests, @@ -344,6 +351,7 @@ impl LobSimCuda { trail_hwm_d, max_hold_ns_d, snapshots_skipped_d, + last_event_ts_d, }) } @@ -1129,6 +1137,7 @@ impl LobSimCuda { .arg(&tick) .arg(&self.cost_per_lot_per_side_d) .arg(&mut self.total_fees_per_b_d) + .arg(&mut self.last_event_ts_d) .arg(&n) .launch(cfg)?; } diff --git a/crates/ml-backtesting/tests/stop_controller.rs b/crates/ml-backtesting/tests/stop_controller.rs index 24a583fbd..cd7cc606a 100644 --- a/crates/ml-backtesting/tests/stop_controller.rs +++ b/crates/ml-backtesting/tests/stop_controller.rs @@ -958,6 +958,66 @@ fn trade_vol_floor_prevents_sub_cost_stops() -> Result<()> { Ok(()) } +/// Validates that a timestamp gap > 1 hour (session boundary / weekend halt) +/// force-closes any open position before the next event is processed. +/// +/// max_hold_ns can't fire during a halt because no events advance current_ts. +/// The session-gap check in resting_orders_step detects the jump and zeroes +/// position_lots directly; pnl_track_step's existing close branch then emits +/// exactly one TradeRecord. +#[test] +#[ignore = "requires CUDA"] +fn session_gap_force_closes_open_positions() -> 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)?; + + let cold_start: [IsvKellyStateHost; 5] = std::array::from_fn(|_| IsvKellyStateHost { + pnl_ema_win: 0.0, pnl_ema_loss: 100.0, // huge SL — won't fire + 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 at t=1ms. + let mut ts: u64 = 1_000_000; + sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + sim.step_decision_with_latency(ts, &cfg_default(1))?; + for _ in 0..3 { + sim.step_resting_orders(ts, 0.0)?; + sim.step_pnl_track(ts)?; + ts += 1_000_000; + } + assert!(sim.read_pos(0)?.position_lots > 0, "setup: long opens"); + let trades_before_gap = sim.read_total_trade_count()?; + + // Advance ts by 2 hours (> 1 hour SESSION_GAP_NS). Call step_resting_orders + // with the new ts — session-gap force-close fires, then pnl_track detects + // the prev!=0 && now==0 close transition and emits a TradeRecord. + ts += 2 * 3_600_000_000_000u64; + sim.step_resting_orders(ts, 0.0)?; + sim.step_pnl_track(ts)?; + + assert_eq!( + sim.read_pos(0)?.position_lots, 0, + "session gap > 1h must force-close the position" + ); + let trades_after_gap = sim.read_total_trade_count()?; + assert_eq!( + trades_after_gap, trades_before_gap + 1, + "session-gap close emits exactly one TradeRecord; got {} -> {}", + trades_before_gap, trades_after_gap + ); + Ok(()) +} + /// Validates that a corrupt top-of-book snapshot is skipped entirely and the /// per-backtest `snapshots_skipped` counter is incremented. ///