diff --git a/crates/ml-backtesting/cuda/pnl_track.cu b/crates/ml-backtesting/cuda/pnl_track.cu index 7be99fea3..c6e6675c2 100644 --- a/crates/ml-backtesting/cuda/pnl_track.cu +++ b/crates/ml-backtesting/cuda/pnl_track.cu @@ -42,7 +42,11 @@ extern "C" __global__ void pnl_track_step( unsigned long long current_ts_ns, int trade_log_cap, int n_backtests, - float* trail_hwm // [n_backtests] — reset to 0 on close + float* trail_hwm, // [n_backtests] — reset to 0 on close + // S1.20: NaN instrumentation counters (per-backtest). + unsigned int* __restrict__ zero_vwap_at_open, // open branch: vwap_entry == 0 + unsigned int* __restrict__ saturated_vwap_at_open, // open branch: |vwap_entry| > 21M or NaN/Inf + unsigned int* __restrict__ defensive_exit_clamp // close branch: defensive clamp fired ) { int b = blockIdx.x; if (b >= n_backtests || threadIdx.x != 0) return; @@ -54,8 +58,15 @@ extern "C" __global__ void pnl_track_step( if (prev_size == 0 && now_size != 0) { // Open entry — snapshot context. + // S1.20: instrument vwap_entry at open to catch zero/saturated sentinels. + const float vwap_at_open = pos->vwap_entry; + if (vwap_at_open == 0.0f) { + zero_vwap_at_open[b] += 1u; + } else if (!isfinite(vwap_at_open) || fabsf(vwap_at_open) > 21000000.0f) { + saturated_vwap_at_open[b] += 1u; + } *reinterpret_cast(st + 0) = current_ts_ns; - *reinterpret_cast(st + 8) = (int)(pos->vwap_entry * 100.0f + 0.5f); + *reinterpret_cast(st + 8) = (int)(vwap_at_open * 100.0f + 0.5f); *reinterpret_cast(st + 12) = now_size; *reinterpret_cast(st + 16) = pos->realized_pnl; st[20] = 0; // horizon_idx; set by decision kernel in C7 @@ -97,6 +108,7 @@ extern "C" __global__ void pnl_track_step( if (!isfinite(exit_px_safe) || fabsf(exit_px_safe - entry_px) > 0.1f * fabsf(entry_px)) { + defensive_exit_clamp[b] += 1u; // S1.20: defensive clamp fired exit_px_safe = entry_px; realised_usd_to_write = 0; } diff --git a/crates/ml-backtesting/cuda/resting_orders.cu b/crates/ml-backtesting/cuda/resting_orders.cu index 9ddd1b49f..e934c0bbb 100644 --- a/crates/ml-backtesting/cuda/resting_orders.cu +++ b/crates/ml-backtesting/cuda/resting_orders.cu @@ -99,6 +99,12 @@ __device__ static void cancel_oco_pair(Orders& orders, unsigned char oco_pair) { // Plain `+=` for total_fees_per_b is safe under the single-writer-per-block // convention enforced by `if (threadIdx.x != 0) return;` at the top of // every block-per-backtest kernel. +// +// S1.20: NaN instrumentation counters. +// nan_avg_px — incremented when avg_px = total_cost/filled_lots is NaN/Inf. +// nan_realised — incremented when realised P&L in counter-direction branch is NaN/Inf. +// nan_realized_pnl — incremented when pos.realized_pnl becomes non-finite after += or -=. +// On NaN-producing events the function returns early WITHOUT poisoning pos state. __device__ static void apply_fill_to_pos( Pos& pos, float filled_lots, @@ -106,10 +112,17 @@ __device__ static void apply_fill_to_pos( unsigned char side, int b, const float* __restrict__ cost_per_lot_per_side_per_b, - float* __restrict__ total_fees_per_b + float* __restrict__ total_fees_per_b, + unsigned int* __restrict__ nan_avg_px, + unsigned int* __restrict__ nan_realised, + unsigned int* __restrict__ nan_realized_pnl ) { if (filled_lots <= 0.0f) return; const float avg_px = total_cost / filled_lots; + if (!isfinite(avg_px)) { + nan_avg_px[b] += 1u; + return; // do NOT propagate NaN into pos state + } const float sgn = (side == 0) ? 1.0f : -1.0f; const float prev_pos_f = (float)pos.position_lots; const float new_pos_f = prev_pos_f + sgn * filled_lots; @@ -126,7 +139,14 @@ __device__ static void apply_fill_to_pos( const float unwind = (filled_lots < prev_abs) ? filled_lots : prev_abs; const float dir_was_long = (prev_pos_f > 0.0f) ? 1.0f : -1.0f; const float realised = (avg_px - pos.vwap_entry) * dir_was_long * unwind; + if (!isfinite(realised)) { + nan_realised[b] += 1u; + return; // do not poison realized_pnl + } pos.realized_pnl += realised; + if (!isfinite(pos.realized_pnl)) { + nan_realized_pnl[b] += 1u; + } if (filled_lots > prev_abs) { pos.vwap_entry = avg_px; } } pos.position_lots = (int)new_pos_f; @@ -134,6 +154,9 @@ __device__ static void apply_fill_to_pos( // P4: per-fill cost deduction (single-writer-per-block; plain += safe). const float fill_cost = filled_lots * cost_per_lot_per_side_per_b[b]; pos.realized_pnl -= fill_cost; + if (!isfinite(pos.realized_pnl)) { + nan_realized_pnl[b] += 1u; + } total_fees_per_b[b] += fill_cost; } @@ -166,6 +189,10 @@ extern "C" __global__ void resting_orders_step( 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] + // S1.20: NaN instrumentation counters (per-backtest). + unsigned int* __restrict__ nan_avg_px, + unsigned int* __restrict__ nan_realised, + unsigned int* __restrict__ nan_realized_pnl, int n_backtests ) { int b = blockIdx.x; @@ -232,7 +259,8 @@ extern "C" __global__ void resting_orders_step( // Filled `take` lots at our limit price (queue position reached us). const float cost = take * s.price; apply_fill_to_pos(pos, take, cost, s.side, - b, cost_per_lot_per_side_per_b, total_fees_per_b); + b, cost_per_lot_per_side_per_b, total_fees_per_b, + nan_avg_px, nan_realised, nan_realized_pnl); s.size_remaining -= take; emit_audit(audit_base, audit_head, b, (unsigned int)ring_cap_events, current_ts_ns, @@ -262,7 +290,8 @@ extern "C" __global__ void resting_orders_step( : walk_bid_for_sell(book, s.size_remaining, filled); if (filled > 0.0f) { apply_fill_to_pos(pos, filled, cost, s.side, - b, cost_per_lot_per_side_per_b, total_fees_per_b); + b, cost_per_lot_per_side_per_b, total_fees_per_b, + nan_avg_px, nan_realised, nan_realized_pnl); s.size_remaining -= filled; emit_audit(audit_base, audit_head, b, (unsigned int)ring_cap_events, current_ts_ns, @@ -308,7 +337,8 @@ extern "C" __global__ void resting_orders_step( : walk_bid_for_sell(book, st.size, filled); if (filled > 0.0f) { apply_fill_to_pos(pos, filled, cost, st.side, - b, cost_per_lot_per_side_per_b, total_fees_per_b); + b, cost_per_lot_per_side_per_b, total_fees_per_b, + nan_avg_px, nan_realised, nan_realized_pnl); emit_audit(audit_base, audit_head, b, (unsigned int)ring_cap_events, current_ts_ns, pack_slot_tag(1 /*SlotKind::Stop*/, (unsigned char)j, st.gen_counter), diff --git a/crates/ml-backtesting/src/harness.rs b/crates/ml-backtesting/src/harness.rs index d998bcb6e..0ef84d3ac 100644 --- a/crates/ml-backtesting/src/harness.rs +++ b/crates/ml-backtesting/src/harness.rs @@ -286,6 +286,25 @@ impl BacktestHarness { } } + // S1.20: NaN instrumentation summary. Print ONCE per smoke so the + // cluster log shows which kernel-side arithmetic site produced the + // residual 97 zero + 85 i32::MAX sentinel trade records. + match self.sim.read_nan_counters() { + Ok(counters) => { + let sum_u32 = |v: &[u32]| -> u64 { v.iter().map(|&x| x as u64).sum() }; + eprintln!( + "nan_counters: avg_px={} realised={} realized_pnl={} zero_vwap_open={} saturated_vwap_open={} defensive_clamp={}", + sum_u32(&counters.nan_avg_px), + sum_u32(&counters.nan_realised), + sum_u32(&counters.nan_realized_pnl), + sum_u32(&counters.zero_vwap_at_open), + sum_u32(&counters.saturated_vwap_at_open), + sum_u32(&counters.defensive_exit_clamp), + ); + } + Err(e) => eprintln!("nan_counters read failed: {e}"), + } + Ok(RunStats { events_processed: self.event_count, decisions_taken: total_decisions, diff --git a/crates/ml-backtesting/src/sim/mod.rs b/crates/ml-backtesting/src/sim/mod.rs index 4210e00d7..c02509885 100644 --- a/crates/ml-backtesting/src/sim/mod.rs +++ b/crates/ml-backtesting/src/sim/mod.rs @@ -40,6 +40,26 @@ const RESTING_ORDERS_CUBIN: &[u8] = const N_HORIZONS: usize = 5; const ISV_KELLY_STATE_BYTES: usize = 24; // matches lob_state.cuh::IsvKellyState +/// Per-backtest NaN instrumentation counters (S1.20). +/// +/// Incremented by the CUDA kernels at specific NaN-producing arithmetic sites. +/// Read via `LobSimCuda::read_nan_counters` at end-of-smoke to identify which +/// code path produces the residual 97 zero + 85 i32::MAX sentinel trade records. +pub struct NanCounters { + /// avg_px = total_cost/filled_lots was NaN/Inf in apply_fill_to_pos. + pub nan_avg_px: Vec, + /// realised = (avg_px - vwap_entry) * dir * unwind was NaN/Inf. + pub nan_realised: Vec, + /// pos.realized_pnl became non-finite after += or -= in apply_fill_to_pos. + pub nan_realized_pnl: Vec, + /// pnl_track open branch saw vwap_entry == 0. + pub zero_vwap_at_open: Vec, + /// pnl_track open branch saw |vwap_entry| > 21M or NaN/Inf. + pub saturated_vwap_at_open: Vec, + /// pnl_track close branch defensive clamp fired. + pub defensive_exit_clamp: Vec, +} + pub struct LobSimCuda { n_backtests: usize, stream: Arc, @@ -131,6 +151,17 @@ pub struct LobSimCuda { // 0.0 / f32::INFINITY mean no-check (backward compatible). min_reasonable_px_d: CudaSlice, // [n_backtests] max_reasonable_px_d: CudaSlice, // [n_backtests] + + // S1.20: NaN instrumentation slots. Per-backtest counters incremented + // at kernel-internal arithmetic sites that can produce NaN/Inf/sentinel + // values in trade records. Read at end-of-smoke to identify which path + // produces the 97 zero + 85 i32::MAX sentinels. + nan_avg_px_d: CudaSlice, // [n_backtests] + nan_realised_d: CudaSlice, // [n_backtests] + nan_realized_pnl_d: CudaSlice, // [n_backtests] + zero_vwap_at_open_d: CudaSlice, // [n_backtests] + saturated_vwap_at_open_d: CudaSlice, // [n_backtests] + defensive_exit_clamp_d: CudaSlice, // [n_backtests] } impl LobSimCuda { @@ -316,6 +347,26 @@ impl LobSimCuda { stream.memcpy_htod(&max_reasonable_px_init, &mut max_reasonable_px_d) .context("init max_reasonable_px_d")?; + // S1.20: NaN instrumentation counters, all zero-initialized. + let nan_avg_px_d = stream + .alloc_zeros::(n_backtests) + .context("alloc nan_avg_px_d")?; + let nan_realised_d = stream + .alloc_zeros::(n_backtests) + .context("alloc nan_realised_d")?; + let nan_realized_pnl_d = stream + .alloc_zeros::(n_backtests) + .context("alloc nan_realized_pnl_d")?; + let zero_vwap_at_open_d = stream + .alloc_zeros::(n_backtests) + .context("alloc zero_vwap_at_open_d")?; + let saturated_vwap_at_open_d = stream + .alloc_zeros::(n_backtests) + .context("alloc saturated_vwap_at_open_d")?; + let defensive_exit_clamp_d = stream + .alloc_zeros::(n_backtests) + .context("alloc defensive_exit_clamp_d")?; + Ok(Self { n_backtests, stream, @@ -373,6 +424,12 @@ impl LobSimCuda { last_event_ts_d, min_reasonable_px_d, max_reasonable_px_d, + nan_avg_px_d, + nan_realised_d, + nan_realized_pnl_d, + zero_vwap_at_open_d, + saturated_vwap_at_open_d, + defensive_exit_clamp_d, }) } @@ -518,6 +575,26 @@ impl LobSimCuda { Ok(buf) } + /// Read the 6 per-backtest NaN instrumentation counters (S1.20). + /// Returns a `NanCounters` with one `Vec` per counter, length = n_backtests. + /// All counters are zero at construction and increment on NaN-producing events + /// in the fill + pnl_track kernels. + pub fn read_nan_counters(&self) -> Result { + let read = |slot: &CudaSlice| -> Result> { + let mut buf = vec![0u32; self.n_backtests]; + self.stream.memcpy_dtoh(slot, buf.as_mut_slice())?; + Ok(buf) + }; + Ok(NanCounters { + nan_avg_px: read(&self.nan_avg_px_d)?, + nan_realised: read(&self.nan_realised_d)?, + nan_realized_pnl: read(&self.nan_realized_pnl_d)?, + zero_vwap_at_open: read(&self.zero_vwap_at_open_d)?, + saturated_vwap_at_open: read(&self.saturated_vwap_at_open_d)?, + defensive_exit_clamp: read(&self.defensive_exit_clamp_d)?, + }) + } + /// Upload per-backtest price-range bounds. Call ONCE before the first /// apply_snapshot. min=0.0 / max=f32::INFINITY disables the check for /// that backtest (default after LobSimCuda::new). @@ -655,6 +732,9 @@ impl LobSimCuda { .arg(&cap) .arg(&n) .arg(&mut self.trail_hwm_d) + .arg(&mut self.zero_vwap_at_open_d) + .arg(&mut self.saturated_vwap_at_open_d) + .arg(&mut self.defensive_exit_clamp_d) .launch(cfg)?; } self.stream.synchronize()?; @@ -901,6 +981,9 @@ impl LobSimCuda { .arg(&cap) .arg(&n) .arg(&mut self.trail_hwm_d) + .arg(&mut self.zero_vwap_at_open_d) + .arg(&mut self.saturated_vwap_at_open_d) + .arg(&mut self.defensive_exit_clamp_d) .launch(cfg)?; } self.stream.synchronize()?; @@ -1185,6 +1268,9 @@ impl LobSimCuda { .arg(&self.cost_per_lot_per_side_d) .arg(&mut self.total_fees_per_b_d) .arg(&mut self.last_event_ts_d) + .arg(&mut self.nan_avg_px_d) + .arg(&mut self.nan_realised_d) + .arg(&mut self.nan_realized_pnl_d) .arg(&n) .launch(cfg)?; } diff --git a/crates/ml-backtesting/tests/stop_controller.rs b/crates/ml-backtesting/tests/stop_controller.rs index a4726ee05..3483774ea 100644 --- a/crates/ml-backtesting/tests/stop_controller.rs +++ b/crates/ml-backtesting/tests/stop_controller.rs @@ -1110,6 +1110,27 @@ fn corrupt_top_of_book_skips_snapshot_and_increments_counter() -> Result<()> { Ok(()) } +/// Validates that all 6 NaN instrumentation counters (S1.20) are zero after +/// LobSimCuda::new() with no kernel launches. This confirms the device buffers +/// are correctly zero-initialized and the accessor compiles + works. +#[test] +#[ignore = "requires CUDA"] +fn nan_counters_initialize_to_zero() -> Result<()> { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { eprintln!("skipping: cuda unavailable ({e})"); return Ok(()); } + }; + let sim = LobSimCuda::new(1, &dev)?; + let counters = sim.read_nan_counters()?; + assert_eq!(counters.nan_avg_px[0], 0, "nan_avg_px must be zero at init"); + assert_eq!(counters.nan_realised[0], 0, "nan_realised must be zero at init"); + assert_eq!(counters.nan_realized_pnl[0], 0, "nan_realized_pnl must be zero at init"); + assert_eq!(counters.zero_vwap_at_open[0], 0, "zero_vwap_at_open must be zero at init"); + assert_eq!(counters.saturated_vwap_at_open[0], 0, "saturated_vwap_at_open must be zero at init"); + assert_eq!(counters.defensive_exit_clamp[0], 0, "defensive_exit_clamp must be zero at init"); + Ok(()) +} + /// Validates that the parameterized price-range filter rejects real-data outliers. /// /// Audit (fxt-data-audit on ES.FUT_2024-Q1) found: