fix(ml-backtesting): three residual sentinel + session-gap fixes

Three orthogonal fixes for residual issues in smoke stx9p (97 zero
sentinels, 85 i32::MAX sentinels, 840/1024 over-60s holds):

Fix 1 — zero-sentinel residue (px > 0):
- Per-level sanitization in apply_snapshot_kernel only checked sz > 0.
- A book level with px=0, sz>0 passed → walk_* computed total_cost=0
  → avg_px=0 → vwap_entry=0. Trade record reported entry_px=0.
- Add px > 0.0f to bid_ok/ask_ok conditions.

Fix 2 — i32::MAX upper bound (px < 1e8):
- Post-Bug-D (1e9 nanoprice scaling) some boundary events carried prices
  larger than any plausible instrument. Saturating cast to i32 produced
  the 21474836 sentinel even when isfinite() passed.
- Add px < 1.0e8f upper bound to per-level AND top-of-book validation.

Fix 3 — session-gap force-close:
- max_hold check fires at decision_stride frequency, but during weekend
  halts no events advance current_ts → max_hold never fires until next
  session. Result: 49h holds in stx9p (840/1024 over 60s threshold).
- Detect ts gap > 1 hour in resting_orders_step. If position is open,
  zero position_lots directly. pnl_track_step's existing close branch
  emits the TradeRecord on the next call. No synthetic P&L added —
  records show realised_pnl from whatever was accumulated before the gap
  (honest: cannot fill across a halt).
- New per-backtest last_event_ts_d slot tracks the previous event ts.
- Test session_gap_force_closes_open_positions: 2-hour ts jump after
  open verifies force-flat fires and exactly 1 TradeRecord is emitted.

All 16 stop_controller tests pass. All 5 decision_floor_coldstart pass.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>

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;

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*>(orders_base + (size_t)b * orders_bytes);
     Pos&    pos    = *reinterpret_cast<Pos*>(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) {

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<u32>, // [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<u64>,    // [n_backtests] — 0 = sentinel (first event)
 }
 
 impl LobSimCuda {
@@ -289,6 +293,9 @@ impl LobSimCuda {
         let snapshots_skipped_d = stream
             .alloc_zeros::<u32>(n_backtests)
             .context("alloc snapshots_skipped_d")?;
+        let last_event_ts_d = stream
+            .alloc_zeros::<u64>(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)?;
         }

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.
 ///