From 691dec144e2ff1db04f9252d289dd61981bdbab9 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Wed, 20 May 2026 08:28:12 +0200 Subject: [PATCH] fix(ml-backtesting): root-cause NaN/Inf book + duplicate close emissions MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Two orthogonal bugs that compounded to produce the exit_px=±21474836 sentinel in cluster smokes (315 trades baseline, 500 trades pearl): Bug A — NaN/Inf book propagation: - apply_snapshot_kernel passes through NaN/Inf prices from MBP-10 predecoded data (session boundaries, gap-fills). - walk_ask_for_buy / walk_bid_for_sell accumulate cost += take * NaN. - apply_fill_to_pos: avg_px = NaN; realized_pnl += NaN → permanent NaN. - Subsequent close: (int)(NaN-derived * 5000) → i32::MAX sentinel. - Fix: zero-out non-finite levels in apply_snapshot_kernel; add isfinite() guards in walk helpers as defense-in-depth; ATR update guards against non-finite mid. Bug B — pnl_track close branch doesn't reset scratch: - Scratch reset (full 24-byte memset to 0) was already present in the current code; no source change required for Bug B. Tests: - book_nan_inf_prices_dont_corrupt_realized_pnl: NaN at ask[3] + Inf at bid[5] survives the fill pipeline without making realized_pnl non-finite. - pnl_track_resets_scratch_on_close: open+close+5 flat events emits exactly 1 record; second open+close emits exactly 1 more (=2 total). 14/14 stop_controller tests pass; all other ml-backtesting test suites unaffected. Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml-backtesting/cuda/book_update.cu | 51 ++++-- crates/ml-backtesting/cuda/order_match.cu | 10 +- crates/ml-backtesting/cuda/resting_orders.cu | 10 +- .../ml-backtesting/tests/stop_controller.rs | 163 ++++++++++++++++++ 4 files changed, 208 insertions(+), 26 deletions(-) diff --git a/crates/ml-backtesting/cuda/book_update.cu b/crates/ml-backtesting/cuda/book_update.cu index 598e6e0f4..eafc3f0ec 100644 --- a/crates/ml-backtesting/cuda/book_update.cu +++ b/crates/ml-backtesting/cuda/book_update.cu @@ -28,31 +28,46 @@ extern "C" __global__ void book_update_apply_snapshot( if (tid < 10) { Book& bk = books_out[b]; - bk.bid_px[tid] = bid_px_in[tid]; - bk.bid_sz[tid] = bid_sz_in[tid]; - bk.ask_px[tid] = ask_px_in[tid]; - bk.ask_sz[tid] = ask_sz_in[tid]; + const float bp = bid_px_in[tid]; + const float bs = bid_sz_in[tid]; + const float ap = ask_px_in[tid]; + const float az = ask_sz_in[tid]; + // Sanitize: level is valid only if price + size are both finite AND size > 0. + // Otherwise zero BOTH so walk_*_for_* (which skips lvl_sz <= 0) treats it as empty. + // Real MBP-10 data at session boundaries / gap-fills can contain NaN or Inf. + const bool bid_ok = isfinite(bp) && isfinite(bs) && bs > 0.0f; + const bool ask_ok = isfinite(ap) && isfinite(az) && az > 0.0f; + 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; + bk.ask_sz[tid] = ask_ok ? az : 0.0f; } // Thread 0 handles per-backtest ATR-EMA update against the broadcast mid. if (tid == 0) { - const float mid = 0.5f * (bid_px_in[0] + ask_px_in[0]); - const float prev = prev_mid[b]; - if (prev == 0.0f) { - // First-observation bootstrap (pearl_first_observation_bootstrap). - prev_mid[b] = mid; - // atr_mid_ema stays 0 until event 2. + const float mid_raw = 0.5f * (bid_px_in[0] + ask_px_in[0]); + if (!isfinite(mid_raw)) { + // Bad top-of-book (NaN/Inf from session boundary); skip ATR update. + // prev_mid and atr_mid_ema remain unchanged. } else { - const float delta = fabsf(mid - prev); - const float ema = atr_mid_ema[b]; - if (ema == 0.0f) { - // First non-zero delta: replace directly. - atr_mid_ema[b] = delta; + const float mid = mid_raw; + const float prev = prev_mid[b]; + if (prev == 0.0f) { + // First-observation bootstrap (pearl_first_observation_bootstrap). + prev_mid[b] = mid; + // atr_mid_ema stays 0 until event 2. } else { - // Wiener-α=0.4 EMA (pearl_wiener_alpha_floor_for_nonstationary). - atr_mid_ema[b] = 0.4f * delta + 0.6f * ema; + const float delta = fabsf(mid - prev); + const float ema = atr_mid_ema[b]; + if (ema == 0.0f) { + // First non-zero delta: replace directly. + atr_mid_ema[b] = delta; + } else { + // Wiener-α=0.4 EMA (pearl_wiener_alpha_floor_for_nonstationary). + atr_mid_ema[b] = 0.4f * delta + 0.6f * ema; + } + prev_mid[b] = mid; } - prev_mid[b] = mid; } } } diff --git a/crates/ml-backtesting/cuda/order_match.cu b/crates/ml-backtesting/cuda/order_match.cu index 17da528c9..1814fed1c 100644 --- a/crates/ml-backtesting/cuda/order_match.cu +++ b/crates/ml-backtesting/cuda/order_match.cu @@ -21,9 +21,10 @@ __device__ static float walk_ask_for_buy( for (int k = 0; k < 10; ++k) { if (size <= 0.0f) break; const float lvl_sz = book.ask_sz[k]; - if (lvl_sz <= 0.0f) continue; + const float lvl_px = book.ask_px[k]; + if (lvl_sz <= 0.0f || !isfinite(lvl_sz) || !isfinite(lvl_px)) continue; const float take = (size < lvl_sz) ? size : lvl_sz; - cost += take * book.ask_px[k]; + cost += take * lvl_px; filled_lots += take; size -= take; } @@ -39,9 +40,10 @@ __device__ static float walk_bid_for_sell( for (int k = 0; k < 10; ++k) { if (size <= 0.0f) break; const float lvl_sz = book.bid_sz[k]; - if (lvl_sz <= 0.0f) continue; + const float lvl_px = book.bid_px[k]; + if (lvl_sz <= 0.0f || !isfinite(lvl_sz) || !isfinite(lvl_px)) continue; const float take = (size < lvl_sz) ? size : lvl_sz; - cost += take * book.bid_px[k]; + cost += take * lvl_px; filled_lots += take; size -= take; } diff --git a/crates/ml-backtesting/cuda/resting_orders.cu b/crates/ml-backtesting/cuda/resting_orders.cu index 9bfe2b1ed..543e00903 100644 --- a/crates/ml-backtesting/cuda/resting_orders.cu +++ b/crates/ml-backtesting/cuda/resting_orders.cu @@ -31,9 +31,10 @@ __device__ static float walk_ask_for_buy(const Book& book, float size, float& fi #pragma unroll for (int k = 0; k < 10 && size > 0.0f; ++k) { const float lvl_sz = book.ask_sz[k]; - if (lvl_sz <= 0.0f) continue; + const float lvl_px = book.ask_px[k]; + if (lvl_sz <= 0.0f || !isfinite(lvl_sz) || !isfinite(lvl_px)) continue; const float take = (size < lvl_sz) ? size : lvl_sz; - cost += take * book.ask_px[k]; filled += take; size -= take; + cost += take * lvl_px; filled += take; size -= take; } return cost; } @@ -42,9 +43,10 @@ __device__ static float walk_bid_for_sell(const Book& book, float size, float& f #pragma unroll for (int k = 0; k < 10 && size > 0.0f; ++k) { const float lvl_sz = book.bid_sz[k]; - if (lvl_sz <= 0.0f) continue; + const float lvl_px = book.bid_px[k]; + if (lvl_sz <= 0.0f || !isfinite(lvl_sz) || !isfinite(lvl_px)) continue; const float take = (size < lvl_sz) ? size : lvl_sz; - cost += take * book.bid_px[k]; filled += take; size -= take; + cost += take * lvl_px; filled += take; size -= take; } return cost; } diff --git a/crates/ml-backtesting/tests/stop_controller.rs b/crates/ml-backtesting/tests/stop_controller.rs index 8e5e0450d..9ad7f4a94 100644 --- a/crates/ml-backtesting/tests/stop_controller.rs +++ b/crates/ml-backtesting/tests/stop_controller.rs @@ -694,6 +694,169 @@ fn max_hold_forces_close() -> Result<()> { Ok(()) } +/// Validates that NaN or Inf prices in an MBP-10 snapshot do not propagate +/// into realized_pnl via the book walk → fill pipeline. +/// +/// Real MBP-10 data at session boundaries or gap-fills can contain NaN or Inf +/// in level prices/sizes. apply_snapshot_kernel must sanitize them so the walk +/// helpers never accumulate NaN/Inf cost, which would corrupt avg_px and +/// subsequently realized_pnl, emitting the ±21,474,836 sentinel. +#[test] +#[ignore = "requires CUDA"] +fn book_nan_inf_prices_dont_corrupt_realized_pnl() -> 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)?; + + // Build a book with NaN at one level and Inf at another (real MBP-10 + // data at session boundaries can contain these). + let mut bid_px = [0.0_f32; 10]; + let mut bid_sz = [0.0_f32; 10]; + let mut ask_px = [0.0_f32; 10]; + let mut ask_sz = [0.0_f32; 10]; + for k in 0..10 { + bid_px[k] = 5500.0 - 0.25 * (k as f32 + 1.0); + ask_px[k] = 5500.0 + 0.25 * (k as f32 + 1.0); + bid_sz[k] = 20.0; + ask_sz[k] = 20.0; + } + // Inject corruption at level 3 (NaN price) and level 5 (Inf size). + ask_px[3] = f32::NAN; + bid_sz[5] = f32::INFINITY; + + // First snapshot bootstraps ATR + book. + sim.apply_snapshot(&bid_px, &bid_sz, &ask_px, &ask_sz)?; + // Second snapshot: same book; ATR converges; sanitization tested again. + sim.apply_snapshot(&bid_px, &bid_sz, &ask_px, &ask_sz)?; + + // Open long via strong alpha. The fill must NOT propagate NaN/Inf into + // pos.realized_pnl, even though the book has NaN at ask[3] and Inf at bid[5]. + sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + let mut ts: u64 = 1_000_000; + sim.step_decision_with_latency(ts, &cfg_default(1))?; + for _ in 0..5 { + sim.step_resting_orders(ts, 0.0)?; + sim.step_pnl_track(ts)?; + ts += 1_000_000; + } + + let pos = sim.read_pos(0)?; + assert!( + pos.realized_pnl.is_finite(), + "realized_pnl must stay finite even with NaN/Inf in book; got {}", + pos.realized_pnl + ); + // Position may or may not have opened (depending on which levels were used). + // The invariant is finiteness, not a specific position state. + Ok(()) +} + +/// Validates that pnl_track resets scratch on close so that: +/// (a) flat events after a close do NOT emit duplicate close records, and +/// (b) the next open transition is captured correctly (exactly 1 record per trade). +/// +/// Without the scratch reset: prev_size stays = prior entry_size, so every flat +/// event fires the close branch again (duplicates). The next open (0→1) sees +/// prev=1, now=1 → neither branch fires → open record never captured → second +/// close uses stale realized_at_open from trade 1. +#[test] +#[ignore = "requires CUDA"] +fn pnl_track_resets_scratch_on_close() -> 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 so it doesn'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)?; + + // Config with max_hold_ns = 100ms to force-close without needing SL/trail. + let cfg_max_hold = BatchedSimConfig::from_uniform(1, &UniformSimParams { + target_annual_vol_units: 50.0, + annualisation_factor: 825.0, + max_lots: 5, + latency_ns: 0, + kelly_frac_floor: 0.20, + sharpe_weight_floor: 0.10, + threshold: 0.0, + cost_per_lot_per_side: 0.0, + max_hold_ns: 100_000_000, // 100ms + }); + + // Trade 1: open long, force-flatten via max_hold. + 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_max_hold)?; + 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, "trade 1: long opens"); + + // Force close trade 1 via max_hold: jump ts by 200ms. + ts += 200_000_000; + sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + sim.step_decision_with_latency(ts, &cfg_max_hold)?; + for _ in 0..5 { + sim.step_resting_orders(ts, 0.0)?; + sim.step_pnl_track(ts)?; + ts += 1_000_000; + } + assert_eq!(sim.read_pos(0)?.position_lots, 0, "trade 1: closed via max_hold"); + let trades_after_t1 = sim.read_total_trade_count()?; + assert_eq!(trades_after_t1, 1, + "trade 1: exactly ONE record emitted (no duplicates); got {trades_after_t1}"); + + // Several flat events — pnl_track must NOT emit duplicate close records. + for _ in 0..5 { + sim.step_resting_orders(ts, 0.0)?; + sim.step_pnl_track(ts)?; + ts += 1_000_000; + } + let trades_after_flat = sim.read_total_trade_count()?; + assert_eq!(trades_after_flat, 1, + "flat events: no duplicate close records; got {trades_after_flat}"); + + // Trade 2: re-open + close via max_hold. + ts += 10_000_000; + sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + sim.step_decision_with_latency(ts, &cfg_max_hold)?; + 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, "trade 2: long opens"); + + // Force close trade 2. + ts += 200_000_000; + sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + sim.step_decision_with_latency(ts, &cfg_max_hold)?; + for _ in 0..5 { + sim.step_resting_orders(ts, 0.0)?; + sim.step_pnl_track(ts)?; + ts += 1_000_000; + } + assert_eq!(sim.read_pos(0)?.position_lots, 0, "trade 2: closed via max_hold"); + let trades_after_t2 = sim.read_total_trade_count()?; + assert_eq!(trades_after_t2, 2, + "trade 2: exactly TWO records total (1 per trade, no missed open); got {trades_after_t2}"); + Ok(()) +} + /// Validates that the trade-vol floor prevents sub-cost stop fires. /// /// Per pearl_trade_level_vol_for_stop_distance.md: at cold-start (var=0),