From fc41440dc9d8eb2177ecfb80f0886c67b4facf71 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Wed, 20 May 2026 13:27:03 +0200 Subject: [PATCH] =?UTF-8?q?diag(ml-backtesting):=20finer=20NaN=20instrumen?= =?UTF-8?q?tation=20=E2=80=94=20per-vwap-write-site=20+=20last-bad-vwap=20?= =?UTF-8?q?capture?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit v1 instrumentation (S1.20) eliminated 3 of 6 hypotheses: apply_fill_to_pos arithmetic is fully clean (avg_px=0 realised=0 realized_pnl=0). But pnl_track open branch still saw vwap_entry=0 (194 times) or > 21M (216 times) at the open transition. v2 adds per-vwap-write-site counters so we can pinpoint which apply_fill branch produced the bad vwap, plus captures the actual last-bad-vwap value and the path id (1..6 = open-flat, scale-in zero/huge, flip-beyond zero/huge, session-gap-saw-stale). The 7th counter (vwap_session_gap_was_bad) fires when the session-gap force-close path sees pos.vwap_entry already in a corrupt state pre-gap. Logged at end of smoke as `nan_counters_v2: zero_flat=N huge_flat=N zero_scale=N huge_scale=N zero_flip=N huge_flip=N session_gap_was_bad=N | b0_last_bad_vwap=X b0_last_bad_path=N`. 19/19 stop_controller CUDA tests pass; 5/5 decision_floor_coldstart pass. Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml-backtesting/cuda/resting_orders.cu | 86 ++++++++++++- crates/ml-backtesting/src/harness.rs | 23 ++++ crates/ml-backtesting/src/sim/mod.rs | 119 ++++++++++++++++++ .../ml-backtesting/tests/stop_controller.rs | 23 ++++ 4 files changed, 246 insertions(+), 5 deletions(-) diff --git a/crates/ml-backtesting/cuda/resting_orders.cu b/crates/ml-backtesting/cuda/resting_orders.cu index e934c0bbb..b99539a50 100644 --- a/crates/ml-backtesting/cuda/resting_orders.cu +++ b/crates/ml-backtesting/cuda/resting_orders.cu @@ -105,6 +105,12 @@ __device__ static void cancel_oco_pair(Orders& orders, unsigned char oco_pair) { // 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. +// +// S1.21: per-vwap-write-site counters. +// vwap_zero_open_flat / vwap_huge_open_flat — open-from-flat branch (line 131). +// vwap_zero_scale_in / vwap_huge_scale_in — scale-in branch (line 136). +// vwap_zero_flip / vwap_huge_flip — flip-beyond branch (line 150). +// last_bad_vwap / last_bad_path — capture site + value. __device__ static void apply_fill_to_pos( Pos& pos, float filled_lots, @@ -115,7 +121,15 @@ __device__ static void apply_fill_to_pos( float* __restrict__ total_fees_per_b, unsigned int* __restrict__ nan_avg_px, unsigned int* __restrict__ nan_realised, - unsigned int* __restrict__ nan_realized_pnl + unsigned int* __restrict__ nan_realized_pnl, + unsigned int* __restrict__ vwap_zero_open_flat, + unsigned int* __restrict__ vwap_huge_open_flat, + unsigned int* __restrict__ vwap_zero_scale_in, + unsigned int* __restrict__ vwap_huge_scale_in, + unsigned int* __restrict__ vwap_zero_flip, + unsigned int* __restrict__ vwap_huge_flip, + float* __restrict__ last_bad_vwap, + unsigned int* __restrict__ last_bad_path ) { if (filled_lots <= 0.0f) return; const float avg_px = total_cost / filled_lots; @@ -129,11 +143,31 @@ __device__ static void apply_fill_to_pos( if (pos.position_lots == 0) { pos.vwap_entry = avg_px; + // S1.21: open-from-flat write site. + if (pos.vwap_entry == 0.0f) { + vwap_zero_open_flat[b] += 1u; + last_bad_vwap[b] = 0.0f; + last_bad_path[b] = 1u; + } else if (!isfinite(pos.vwap_entry) || fabsf(pos.vwap_entry) > 21000000.0f) { + vwap_huge_open_flat[b] += 1u; + last_bad_vwap[b] = pos.vwap_entry; + last_bad_path[b] = 1u; + } } else if ((prev_pos_f > 0.0f) == (sgn > 0.0f)) { const float prev_notional = prev_pos_f * pos.vwap_entry; const float new_notional = prev_notional + sgn * total_cost; const float new_abs = (new_pos_f > 0.0f) ? new_pos_f : -new_pos_f; pos.vwap_entry = (new_abs > 0.0f) ? (new_notional / (sgn * new_abs)) : 0.0f; + // S1.21: scale-in write site. + if (pos.vwap_entry == 0.0f) { + vwap_zero_scale_in[b] += 1u; + last_bad_vwap[b] = 0.0f; + last_bad_path[b] = 2u; + } else if (!isfinite(pos.vwap_entry) || fabsf(pos.vwap_entry) > 21000000.0f) { + vwap_huge_scale_in[b] += 1u; + last_bad_vwap[b] = pos.vwap_entry; + last_bad_path[b] = 3u; + } } else { const float prev_abs = (prev_pos_f > 0.0f) ? prev_pos_f : -prev_pos_f; const float unwind = (filled_lots < prev_abs) ? filled_lots : prev_abs; @@ -147,7 +181,19 @@ __device__ static void apply_fill_to_pos( if (!isfinite(pos.realized_pnl)) { nan_realized_pnl[b] += 1u; } - if (filled_lots > prev_abs) { pos.vwap_entry = avg_px; } + if (filled_lots > prev_abs) { + pos.vwap_entry = avg_px; + // S1.21: flip-beyond write site. + if (pos.vwap_entry == 0.0f) { + vwap_zero_flip[b] += 1u; + last_bad_vwap[b] = 0.0f; + last_bad_path[b] = 4u; + } else if (!isfinite(pos.vwap_entry) || fabsf(pos.vwap_entry) > 21000000.0f) { + vwap_huge_flip[b] += 1u; + last_bad_vwap[b] = pos.vwap_entry; + last_bad_path[b] = 5u; + } + } } pos.position_lots = (int)new_pos_f; @@ -193,6 +239,16 @@ extern "C" __global__ void resting_orders_step( unsigned int* __restrict__ nan_avg_px, unsigned int* __restrict__ nan_realised, unsigned int* __restrict__ nan_realized_pnl, + // S1.21: per-vwap-write-site counters + last-bad-vwap capture. + unsigned int* __restrict__ vwap_zero_open_flat, + unsigned int* __restrict__ vwap_huge_open_flat, + unsigned int* __restrict__ vwap_zero_scale_in, + unsigned int* __restrict__ vwap_huge_scale_in, + unsigned int* __restrict__ vwap_zero_flip, + unsigned int* __restrict__ vwap_huge_flip, + unsigned int* __restrict__ vwap_session_gap_was_bad, + float* __restrict__ last_bad_vwap, + unsigned int* __restrict__ last_bad_path, int n_backtests ) { int b = blockIdx.x; @@ -217,6 +273,14 @@ extern "C" __global__ void resting_orders_step( // 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). + // S1.21: capture pre-close vwap_entry corruption if already bad. + if (pos.vwap_entry == 0.0f || !isfinite(pos.vwap_entry) + || fabsf(pos.vwap_entry) > 21000000.0f) + { + vwap_session_gap_was_bad[b] += 1u; + last_bad_vwap[b] = pos.vwap_entry; + last_bad_path[b] = 6u; + } pos.position_lots = 0; } last_event_ts[b] = current_ts_ns; @@ -260,7 +324,11 @@ extern "C" __global__ void resting_orders_step( 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, - nan_avg_px, nan_realised, nan_realized_pnl); + nan_avg_px, nan_realised, nan_realized_pnl, + vwap_zero_open_flat, vwap_huge_open_flat, + vwap_zero_scale_in, vwap_huge_scale_in, + vwap_zero_flip, vwap_huge_flip, + last_bad_vwap, last_bad_path); s.size_remaining -= take; emit_audit(audit_base, audit_head, b, (unsigned int)ring_cap_events, current_ts_ns, @@ -291,7 +359,11 @@ extern "C" __global__ void resting_orders_step( 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, - nan_avg_px, nan_realised, nan_realized_pnl); + nan_avg_px, nan_realised, nan_realized_pnl, + vwap_zero_open_flat, vwap_huge_open_flat, + vwap_zero_scale_in, vwap_huge_scale_in, + vwap_zero_flip, vwap_huge_flip, + last_bad_vwap, last_bad_path); s.size_remaining -= filled; emit_audit(audit_base, audit_head, b, (unsigned int)ring_cap_events, current_ts_ns, @@ -338,7 +410,11 @@ extern "C" __global__ void resting_orders_step( 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, - nan_avg_px, nan_realised, nan_realized_pnl); + nan_avg_px, nan_realised, nan_realized_pnl, + vwap_zero_open_flat, vwap_huge_open_flat, + vwap_zero_scale_in, vwap_huge_scale_in, + vwap_zero_flip, vwap_huge_flip, + last_bad_vwap, last_bad_path); 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 0ef84d3ac..179fe26f4 100644 --- a/crates/ml-backtesting/src/harness.rs +++ b/crates/ml-backtesting/src/harness.rs @@ -305,6 +305,29 @@ impl BacktestHarness { Err(e) => eprintln!("nan_counters read failed: {e}"), } + // S1.21: per-vwap-write-site counters. Printed after v1 line so both + // appear together in the smoke log. + match self.sim.read_nan_counters_v2() { + Ok(c) => { + let sum_u32 = |v: &[u32]| -> u64 { v.iter().map(|&x| x as u64).sum() }; + let last_bad_vwap_b0 = c.last_bad_vwap.first().copied().unwrap_or(0.0); + let last_bad_path_b0 = c.last_bad_path.first().copied().unwrap_or(0); + eprintln!( + "nan_counters_v2: zero_flat={} huge_flat={} zero_scale={} huge_scale={} zero_flip={} huge_flip={} session_gap_was_bad={} | b0_last_bad_vwap={} b0_last_bad_path={}", + sum_u32(&c.vwap_zero_open_flat), + sum_u32(&c.vwap_huge_open_flat), + sum_u32(&c.vwap_zero_scale_in), + sum_u32(&c.vwap_huge_scale_in), + sum_u32(&c.vwap_zero_flip), + sum_u32(&c.vwap_huge_flip), + sum_u32(&c.vwap_session_gap_was_bad), + last_bad_vwap_b0, + last_bad_path_b0, + ); + } + Err(e) => eprintln!("nan_counters_v2 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 c02509885..a432c4a80 100644 --- a/crates/ml-backtesting/src/sim/mod.rs +++ b/crates/ml-backtesting/src/sim/mod.rs @@ -60,6 +60,40 @@ pub struct NanCounters { pub defensive_exit_clamp: Vec, } +/// Per-backtest per-vwap-write-site NaN instrumentation counters (S1.21). +/// +/// Finer-grained than NanCounters: each counter corresponds to exactly one +/// vwap_entry write site inside apply_fill_to_pos (or the session-gap path). +/// Read via `LobSimCuda::read_nan_counters_v2`. +/// +/// Path ids stored in last_bad_path: +/// 1 = open-from-flat (line 131) +/// 2 = scale-in zero (line 136) +/// 3 = scale-in huge (line 136) +/// 4 = flip-beyond zero (line 150) +/// 5 = flip-beyond huge (line 150) +/// 6 = session-gap saw already-corrupt vwap_entry +pub struct NanCountersV2 { + /// apply_fill_to_pos open-from-flat branch wrote vwap_entry == 0. + pub vwap_zero_open_flat: Vec, + /// apply_fill_to_pos open-from-flat branch wrote |vwap_entry| > 21M or non-finite. + pub vwap_huge_open_flat: Vec, + /// apply_fill_to_pos scale-in branch wrote vwap_entry == 0. + pub vwap_zero_scale_in: Vec, + /// apply_fill_to_pos scale-in branch wrote |vwap_entry| > 21M or non-finite. + pub vwap_huge_scale_in: Vec, + /// apply_fill_to_pos flip-beyond branch wrote vwap_entry == 0. + pub vwap_zero_flip: Vec, + /// apply_fill_to_pos flip-beyond branch wrote |vwap_entry| > 21M or non-finite. + pub vwap_huge_flip: Vec, + /// session-gap force-close saw pos.vwap_entry already corrupt (zero / non-finite / > 21M). + pub vwap_session_gap_was_bad: Vec, + /// Last bad vwap_entry value observed (per backtest; 0.0 if none). + pub last_bad_vwap: Vec, + /// Path id of the last bad vwap_entry write (per backtest; 0 if none). + pub last_bad_path: Vec, +} + pub struct LobSimCuda { n_backtests: usize, stream: Arc, @@ -162,6 +196,16 @@ pub struct LobSimCuda { zero_vwap_at_open_d: CudaSlice, // [n_backtests] saturated_vwap_at_open_d: CudaSlice, // [n_backtests] defensive_exit_clamp_d: CudaSlice, // [n_backtests] + // S1.21: per-vwap-write-site counters + last-bad-vwap capture. + vwap_zero_open_flat_d: CudaSlice, // [n_backtests] + vwap_huge_open_flat_d: CudaSlice, // [n_backtests] + vwap_zero_scale_in_d: CudaSlice, // [n_backtests] + vwap_huge_scale_in_d: CudaSlice, // [n_backtests] + vwap_zero_flip_d: CudaSlice, // [n_backtests] + vwap_huge_flip_d: CudaSlice, // [n_backtests] + vwap_session_gap_was_bad_d: CudaSlice, // [n_backtests] + last_bad_vwap_d: CudaSlice, // [n_backtests] + last_bad_path_d: CudaSlice, // [n_backtests] } impl LobSimCuda { @@ -366,6 +410,34 @@ impl LobSimCuda { let defensive_exit_clamp_d = stream .alloc_zeros::(n_backtests) .context("alloc defensive_exit_clamp_d")?; + // S1.21: per-vwap-write-site counters + last-bad-vwap capture. + let vwap_zero_open_flat_d = stream + .alloc_zeros::(n_backtests) + .context("alloc vwap_zero_open_flat_d")?; + let vwap_huge_open_flat_d = stream + .alloc_zeros::(n_backtests) + .context("alloc vwap_huge_open_flat_d")?; + let vwap_zero_scale_in_d = stream + .alloc_zeros::(n_backtests) + .context("alloc vwap_zero_scale_in_d")?; + let vwap_huge_scale_in_d = stream + .alloc_zeros::(n_backtests) + .context("alloc vwap_huge_scale_in_d")?; + let vwap_zero_flip_d = stream + .alloc_zeros::(n_backtests) + .context("alloc vwap_zero_flip_d")?; + let vwap_huge_flip_d = stream + .alloc_zeros::(n_backtests) + .context("alloc vwap_huge_flip_d")?; + let vwap_session_gap_was_bad_d = stream + .alloc_zeros::(n_backtests) + .context("alloc vwap_session_gap_was_bad_d")?; + let last_bad_vwap_d = stream + .alloc_zeros::(n_backtests) + .context("alloc last_bad_vwap_d")?; + let last_bad_path_d = stream + .alloc_zeros::(n_backtests) + .context("alloc last_bad_path_d")?; Ok(Self { n_backtests, @@ -430,6 +502,15 @@ impl LobSimCuda { zero_vwap_at_open_d, saturated_vwap_at_open_d, defensive_exit_clamp_d, + vwap_zero_open_flat_d, + vwap_huge_open_flat_d, + vwap_zero_scale_in_d, + vwap_huge_scale_in_d, + vwap_zero_flip_d, + vwap_huge_flip_d, + vwap_session_gap_was_bad_d, + last_bad_vwap_d, + last_bad_path_d, }) } @@ -595,6 +676,35 @@ impl LobSimCuda { }) } + /// S1.21: per-vwap-write-site counters. + /// + /// All counters are zero at construction. Incremented in + /// `apply_fill_to_pos` at each vwap_entry write site and in the + /// session-gap force-close path in `resting_orders_step`. + pub fn read_nan_counters_v2(&self) -> Result { + let read_u32 = |slot: &CudaSlice| -> Result> { + let mut buf = vec![0u32; self.n_backtests]; + self.stream.memcpy_dtoh(slot, buf.as_mut_slice())?; + Ok(buf) + }; + let read_f32 = |slot: &CudaSlice| -> Result> { + let mut buf = vec![0f32; self.n_backtests]; + self.stream.memcpy_dtoh(slot, buf.as_mut_slice())?; + Ok(buf) + }; + Ok(NanCountersV2 { + vwap_zero_open_flat: read_u32(&self.vwap_zero_open_flat_d)?, + vwap_huge_open_flat: read_u32(&self.vwap_huge_open_flat_d)?, + vwap_zero_scale_in: read_u32(&self.vwap_zero_scale_in_d)?, + vwap_huge_scale_in: read_u32(&self.vwap_huge_scale_in_d)?, + vwap_zero_flip: read_u32(&self.vwap_zero_flip_d)?, + vwap_huge_flip: read_u32(&self.vwap_huge_flip_d)?, + vwap_session_gap_was_bad: read_u32(&self.vwap_session_gap_was_bad_d)?, + last_bad_vwap: read_f32(&self.last_bad_vwap_d)?, + last_bad_path: read_u32(&self.last_bad_path_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). @@ -1271,6 +1381,15 @@ impl LobSimCuda { .arg(&mut self.nan_avg_px_d) .arg(&mut self.nan_realised_d) .arg(&mut self.nan_realized_pnl_d) + .arg(&mut self.vwap_zero_open_flat_d) + .arg(&mut self.vwap_huge_open_flat_d) + .arg(&mut self.vwap_zero_scale_in_d) + .arg(&mut self.vwap_huge_scale_in_d) + .arg(&mut self.vwap_zero_flip_d) + .arg(&mut self.vwap_huge_flip_d) + .arg(&mut self.vwap_session_gap_was_bad_d) + .arg(&mut self.last_bad_vwap_d) + .arg(&mut self.last_bad_path_d) .arg(&n) .launch(cfg)?; } diff --git a/crates/ml-backtesting/tests/stop_controller.rs b/crates/ml-backtesting/tests/stop_controller.rs index 3483774ea..6f8bd1887 100644 --- a/crates/ml-backtesting/tests/stop_controller.rs +++ b/crates/ml-backtesting/tests/stop_controller.rs @@ -1191,3 +1191,26 @@ fn price_range_rejection_skips_snapshot() -> Result<()> { Ok(()) } + +/// S1.21: All per-vwap-write-site v2 counters initialize to zero at +/// LobSimCuda construction; last_bad_path initializes to zero (no bad +/// write observed yet). +#[test] +#[ignore = "requires CUDA"] +fn nan_counters_v2_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 c = sim.read_nan_counters_v2()?; + assert_eq!(c.vwap_zero_open_flat[0], 0, "vwap_zero_open_flat must initialize to 0"); + assert_eq!(c.vwap_huge_open_flat[0], 0, "vwap_huge_open_flat must initialize to 0"); + assert_eq!(c.vwap_zero_scale_in[0], 0, "vwap_zero_scale_in must initialize to 0"); + assert_eq!(c.vwap_huge_scale_in[0], 0, "vwap_huge_scale_in must initialize to 0"); + assert_eq!(c.vwap_zero_flip[0], 0, "vwap_zero_flip must initialize to 0"); + assert_eq!(c.vwap_huge_flip[0], 0, "vwap_huge_flip must initialize to 0"); + assert_eq!(c.vwap_session_gap_was_bad[0], 0, "vwap_session_gap_was_bad must initialize to 0"); + assert_eq!(c.last_bad_path[0], 0, "last_bad_path must initialize to 0 (no bad write)"); + Ok(()) +}