diff --git a/crates/ml-backtesting/cuda/decision_policy.cu b/crates/ml-backtesting/cuda/decision_policy.cu index 38fb3fea1..aa0834465 100644 --- a/crates/ml-backtesting/cuda/decision_policy.cu +++ b/crates/ml-backtesting/cuda/decision_policy.cu @@ -50,24 +50,50 @@ __device__ static int stop_check_isv( const unsigned long long* __restrict__ max_hold_ns_per_b, unsigned long long current_ts_ns, const unsigned char* __restrict__ open_trade_state, - int* __restrict__ market_targets + int* __restrict__ market_targets, + // S2.1: max_hold enforcement diagnostic counters (7 per-backtest u32 arrays). + unsigned int* __restrict__ mh_kernel_calls, + unsigned int* __restrict__ mh_seen_zero, + unsigned int* __restrict__ mh_entry_ts_zero, + unsigned int* __restrict__ mh_ts_underflow, + unsigned int* __restrict__ mh_elapsed_below_cap, + unsigned int* __restrict__ mh_would_fire, + unsigned int* __restrict__ mh_force_flat_written ) { const Pos& pos = positions[b]; if (pos.position_lots == 0) return 0; + // S2.1: count every entry with a non-zero position (max_hold check relevant). + mh_kernel_calls[b] += 1u; + // Per spec §11 follow-up (cluster smoke gp74n max_hold=263985s): // max_hold force-close prevents pathological session-spanning positions. // max_hold_ns = 0 means disabled; max_hold_ns > 0 fires force-flat when // (current_ts - entry_ts) >= max_hold_ns. entry_ts comes from // open_trade_state[b][0..8] maintained by pnl_track_step. const unsigned long long max_hold = max_hold_ns_per_b[b]; - if (max_hold > 0ull) { + if (max_hold == 0ull) { + // S2.1: mode A — max_hold was zero (upload didn't land or disabled). + mh_seen_zero[b] += 1u; + } else { const unsigned long long entry_ts = *reinterpret_cast( open_trade_state + (size_t)b * 24); // OPEN_TRADE_STATE_BYTES = 24 - if (entry_ts > 0ull && current_ts_ns >= entry_ts && (current_ts_ns - entry_ts) >= max_hold) { + if (entry_ts == 0ull) { + // S2.1: mode B — pnl_track open branch never wrote entry_ts. + mh_entry_ts_zero[b] += 1u; + } else if (current_ts_ns < entry_ts) { + // S2.1: mode C — ts underflow (should never happen). + mh_ts_underflow[b] += 1u; + } else if ((current_ts_ns - entry_ts) < max_hold) { + // S2.1: mode D — elapsed < cap; check not yet due. + mh_elapsed_below_cap[b] += 1u; + } else { + // S2.1: elapsed >= max_hold; force-flat fires. + mh_would_fire[b] += 1u; market_targets[b * 2 + 0] = 3; // force-flat market_targets[b * 2 + 1] = 0; + mh_force_flat_written[b] += 1u; return 1; } } @@ -166,6 +192,14 @@ extern "C" __global__ void decision_policy_default( const unsigned long long* __restrict__ max_hold_ns_per_b, // [n_backtests] unsigned long long current_ts_ns, // scalar const unsigned char* __restrict__ open_trade_state_per_b, // [n_backtests * 24] + // S2.1: max_hold enforcement diagnostic counters. + unsigned int* __restrict__ mh_kernel_calls_per_b, // [n_backtests] + unsigned int* __restrict__ mh_seen_zero_per_b, // [n_backtests] + unsigned int* __restrict__ mh_entry_ts_zero_per_b, // [n_backtests] + unsigned int* __restrict__ mh_ts_underflow_per_b, // [n_backtests] + unsigned int* __restrict__ mh_elapsed_below_cap_per_b, // [n_backtests] + unsigned int* __restrict__ mh_would_fire_per_b, // [n_backtests] + unsigned int* __restrict__ mh_force_flat_written_per_b, // [n_backtests] int n_backtests ) { int b = blockIdx.x; @@ -177,7 +211,11 @@ extern "C" __global__ void decision_policy_default( atr_mid_ema_per_b, trail_hwm_per_b, books_per_b, cost_per_lot_per_side_per_b, max_hold_ns_per_b, current_ts_ns, open_trade_state_per_b, - market_targets)) { + market_targets, + mh_kernel_calls_per_b, mh_seen_zero_per_b, + mh_entry_ts_zero_per_b, mh_ts_underflow_per_b, + mh_elapsed_below_cap_per_b, mh_would_fire_per_b, + mh_force_flat_written_per_b)) { return; } @@ -356,6 +394,14 @@ extern "C" __global__ void decision_policy_program( const unsigned long long* __restrict__ max_hold_ns_per_b, // [n_backtests] unsigned long long current_ts_ns, // scalar const unsigned char* __restrict__ open_trade_state_per_b, // [n_backtests * 24] + // S2.1: max_hold enforcement diagnostic counters. + unsigned int* __restrict__ mh_kernel_calls_per_b, // [n_backtests] + unsigned int* __restrict__ mh_seen_zero_per_b, // [n_backtests] + unsigned int* __restrict__ mh_entry_ts_zero_per_b, // [n_backtests] + unsigned int* __restrict__ mh_ts_underflow_per_b, // [n_backtests] + unsigned int* __restrict__ mh_elapsed_below_cap_per_b, // [n_backtests] + unsigned int* __restrict__ mh_would_fire_per_b, // [n_backtests] + unsigned int* __restrict__ mh_force_flat_written_per_b, // [n_backtests] int n_backtests ) { int b = blockIdx.x; @@ -372,7 +418,11 @@ extern "C" __global__ void decision_policy_program( atr_mid_ema_per_b, trail_hwm_per_b, books_per_b, cost_per_lot_per_side_per_b, max_hold_ns_per_b, current_ts_ns, open_trade_state_per_b, - market_targets)) { + market_targets, + mh_kernel_calls_per_b, mh_seen_zero_per_b, + mh_entry_ts_zero_per_b, mh_ts_underflow_per_b, + mh_elapsed_below_cap_per_b, mh_would_fire_per_b, + mh_force_flat_written_per_b)) { return; } diff --git a/crates/ml-backtesting/cuda/resting_orders.cu b/crates/ml-backtesting/cuda/resting_orders.cu index 94c2bee9f..b224d2edb 100644 --- a/crates/ml-backtesting/cuda/resting_orders.cu +++ b/crates/ml-backtesting/cuda/resting_orders.cu @@ -606,12 +606,15 @@ extern "C" __global__ void submit_limit_alloc( // appropriate side. If all 32 slots are occupied, increments // pos.submission_overflow. extern "C" __global__ void seed_inflight_limits_batched( - unsigned char* __restrict__ orders_base, // [n_backtests * orders_bytes] + unsigned char* __restrict__ orders_base, // [n_backtests * orders_bytes] int orders_bytes, - Pos* __restrict__ positions, // [n_backtests] - const int* __restrict__ market_targets, // [n_backtests * 2] - const unsigned int* __restrict__ latency_ns_per_b, // [n_backtests] + Pos* __restrict__ positions, // [n_backtests] + const int* __restrict__ market_targets, // [n_backtests * 2] + const unsigned int* __restrict__ latency_ns_per_b, // [n_backtests] unsigned long long current_ts_ns, + // S2.1: max_hold enforcement diagnostic counter — incremented when this + // kernel reads market_targets[b*2+0] == 3 (force-flat from stop_check_isv). + unsigned int* __restrict__ mh_force_flat_seen_by_seed, // [n_backtests] int n_backtests ) { int b = blockIdx.x; @@ -620,6 +623,13 @@ extern "C" __global__ void seed_inflight_limits_batched( const int side = market_targets[b * 2 + 0]; const int abs_sz = market_targets[b * 2 + 1]; + // S2.1: count every time this kernel sees a force-flat target, regardless + // of whether the delta calculation produces a non-zero order (delta==0 means + // position is already flat, but the force-flat signal still arrived here). + if (side == 3) { + mh_force_flat_seen_by_seed[b] += 1u; + } + if (side == 2) return; // §7: no-op = preserve position int target_signed; diff --git a/crates/ml-backtesting/src/harness.rs b/crates/ml-backtesting/src/harness.rs index 179fe26f4..a0b82d113 100644 --- a/crates/ml-backtesting/src/harness.rs +++ b/crates/ml-backtesting/src/harness.rs @@ -328,6 +328,26 @@ impl BacktestHarness { Err(e) => eprintln!("nan_counters_v2 read failed: {e}"), } + // S2.1: max_hold enforcement diagnostic counters. Identifies which step + // of the max_hold chain fails when trades exceed the configured cap. + match self.sim.read_max_hold_counters() { + Ok(c) => { + let sum_u32 = |v: &[u32]| -> u64 { v.iter().map(|&x| x as u64).sum() }; + eprintln!( + "max_hold_counters: kernel_calls={} seen_zero={} entry_ts_zero={} ts_underflow={} below_cap={} would_fire={} force_flat_written={} seed_saw_force_flat={}", + sum_u32(&c.mh_kernel_calls), + sum_u32(&c.mh_seen_zero), + sum_u32(&c.mh_entry_ts_zero), + sum_u32(&c.mh_ts_underflow), + sum_u32(&c.mh_elapsed_below_cap), + sum_u32(&c.mh_would_fire), + sum_u32(&c.mh_force_flat_written), + sum_u32(&c.mh_force_flat_seen_by_seed), + ); + } + Err(e) => eprintln!("max_hold_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 f4ca95e4b..9009055fc 100644 --- a/crates/ml-backtesting/src/sim/mod.rs +++ b/crates/ml-backtesting/src/sim/mod.rs @@ -94,6 +94,32 @@ pub struct NanCountersV2 { pub last_bad_path: Vec, } +/// S2.1: max_hold enforcement diagnostic counters. +/// +/// Incremented inside `stop_check_isv` (7 counters) and +/// `seed_inflight_limits_batched` (1 counter). Read via +/// `LobSimCuda::read_max_hold_counters` at end-of-smoke to identify +/// which step of the max_hold chain fails when 86.5% of trades exceed +/// the configured 60s cap. +pub struct NanCountersMaxHold { + /// Every stop_check_isv entry with position_lots != 0. + pub mh_kernel_calls: Vec, + /// max_hold_ns_per_b[b] == 0 (mode A: upload didn't land). + pub mh_seen_zero: Vec, + /// entry_ts == 0 despite position open (mode B: pnl_track didn't write). + pub mh_entry_ts_zero: Vec, + /// current_ts_ns < entry_ts — ts underflow (mode C: should never happen). + pub mh_ts_underflow: Vec, + /// elapsed < max_hold — check not yet due (mode D: normal in-flight). + pub mh_elapsed_below_cap: Vec, + /// elapsed >= max_hold — force-flat would fire (mode D-now or success path). + pub mh_would_fire: Vec, + /// market_targets[b*2+0] = 3 was actually written. + pub mh_force_flat_written: Vec, + /// seed_inflight_limits_batched read market_targets[b*2+0] == 3. + pub mh_force_flat_seen_by_seed: Vec, +} + pub struct LobSimCuda { n_backtests: usize, stream: Arc, @@ -206,6 +232,16 @@ pub struct LobSimCuda { vwap_session_gap_was_bad_d: CudaSlice, // [n_backtests] last_bad_vwap_d: CudaSlice, // [n_backtests] last_bad_path_d: CudaSlice, // [n_backtests] + + // S2.1: max_hold enforcement diagnostic counters (7 in stop_check_isv + 1 in seed_inflight). + mh_kernel_calls_d: CudaSlice, // [n_backtests] + mh_seen_zero_d: CudaSlice, // [n_backtests] + mh_entry_ts_zero_d: CudaSlice, // [n_backtests] + mh_ts_underflow_d: CudaSlice, // [n_backtests] + mh_elapsed_below_cap_d: CudaSlice, // [n_backtests] + mh_would_fire_d: CudaSlice, // [n_backtests] + mh_force_flat_written_d: CudaSlice, // [n_backtests] + mh_force_flat_seen_by_seed_d: CudaSlice, // [n_backtests] } impl LobSimCuda { @@ -439,6 +475,32 @@ impl LobSimCuda { .alloc_zeros::(n_backtests) .context("alloc last_bad_path_d")?; + // S2.1: max_hold enforcement diagnostic counters, all zero-initialized. + let mh_kernel_calls_d = stream + .alloc_zeros::(n_backtests) + .context("alloc mh_kernel_calls_d")?; + let mh_seen_zero_d = stream + .alloc_zeros::(n_backtests) + .context("alloc mh_seen_zero_d")?; + let mh_entry_ts_zero_d = stream + .alloc_zeros::(n_backtests) + .context("alloc mh_entry_ts_zero_d")?; + let mh_ts_underflow_d = stream + .alloc_zeros::(n_backtests) + .context("alloc mh_ts_underflow_d")?; + let mh_elapsed_below_cap_d = stream + .alloc_zeros::(n_backtests) + .context("alloc mh_elapsed_below_cap_d")?; + let mh_would_fire_d = stream + .alloc_zeros::(n_backtests) + .context("alloc mh_would_fire_d")?; + let mh_force_flat_written_d = stream + .alloc_zeros::(n_backtests) + .context("alloc mh_force_flat_written_d")?; + let mh_force_flat_seen_by_seed_d = stream + .alloc_zeros::(n_backtests) + .context("alloc mh_force_flat_seen_by_seed_d")?; + Ok(Self { n_backtests, stream, @@ -511,6 +573,14 @@ impl LobSimCuda { vwap_session_gap_was_bad_d, last_bad_vwap_d, last_bad_path_d, + mh_kernel_calls_d, + mh_seen_zero_d, + mh_entry_ts_zero_d, + mh_ts_underflow_d, + mh_elapsed_below_cap_d, + mh_would_fire_d, + mh_force_flat_written_d, + mh_force_flat_seen_by_seed_d, }) } @@ -705,6 +775,30 @@ impl LobSimCuda { }) } + /// S2.1: max_hold enforcement diagnostic counters. + /// + /// All counters are zero at construction. Incremented in `stop_check_isv` + /// (7 counters) and `seed_inflight_limits_batched` (1 counter). Read at + /// end-of-smoke to localize which step of the max_hold chain is failing + /// when 86.5% of trades exceed the configured 60s cap. + pub fn read_max_hold_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(NanCountersMaxHold { + mh_kernel_calls: read(&self.mh_kernel_calls_d)?, + mh_seen_zero: read(&self.mh_seen_zero_d)?, + mh_entry_ts_zero: read(&self.mh_entry_ts_zero_d)?, + mh_ts_underflow: read(&self.mh_ts_underflow_d)?, + mh_elapsed_below_cap: read(&self.mh_elapsed_below_cap_d)?, + mh_would_fire: read(&self.mh_would_fire_d)?, + mh_force_flat_written: read(&self.mh_force_flat_written_d)?, + mh_force_flat_seen_by_seed: read(&self.mh_force_flat_seen_by_seed_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). @@ -1016,6 +1110,14 @@ impl LobSimCuda { .arg(&self.max_hold_ns_d) .arg(¤t_ts_ns) .arg(&self.open_trade_state_d) + // S2.1: max_hold diagnostic counters. + .arg(&mut self.mh_kernel_calls_d) + .arg(&mut self.mh_seen_zero_d) + .arg(&mut self.mh_entry_ts_zero_d) + .arg(&mut self.mh_ts_underflow_d) + .arg(&mut self.mh_elapsed_below_cap_d) + .arg(&mut self.mh_would_fire_d) + .arg(&mut self.mh_force_flat_written_d) .arg(&n) .launch(cfg)?; } @@ -1042,6 +1144,14 @@ impl LobSimCuda { .arg(&self.max_hold_ns_d) .arg(¤t_ts_ns) .arg(&self.open_trade_state_d) + // S2.1: max_hold diagnostic counters. + .arg(&mut self.mh_kernel_calls_d) + .arg(&mut self.mh_seen_zero_d) + .arg(&mut self.mh_entry_ts_zero_d) + .arg(&mut self.mh_ts_underflow_d) + .arg(&mut self.mh_elapsed_below_cap_d) + .arg(&mut self.mh_would_fire_d) + .arg(&mut self.mh_force_flat_written_d) .arg(&n) .launch(cfg)?; } @@ -1163,6 +1273,8 @@ impl LobSimCuda { .arg(&self.market_targets_d) .arg(&self.latency_ns_d) .arg(¤t_ts_ns) + // S2.1: counter 8 — seed_inflight saw force-flat target. + .arg(&mut self.mh_force_flat_seen_by_seed_d) .arg(&n) .launch(cfg_launch)?; } diff --git a/crates/ml-backtesting/tests/stop_controller.rs b/crates/ml-backtesting/tests/stop_controller.rs index 6f8bd1887..1c52dbf4f 100644 --- a/crates/ml-backtesting/tests/stop_controller.rs +++ b/crates/ml-backtesting/tests/stop_controller.rs @@ -1214,3 +1214,23 @@ fn nan_counters_v2_initialize_to_zero() -> Result<()> { assert_eq!(c.last_bad_path[0], 0, "last_bad_path must initialize to 0 (no bad write)"); Ok(()) } + +#[test] +#[ignore = "requires CUDA"] +fn max_hold_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 c = sim.read_max_hold_counters()?; + assert_eq!(c.mh_kernel_calls[0], 0, "mh_kernel_calls must initialize to 0"); + assert_eq!(c.mh_seen_zero[0], 0, "mh_seen_zero must initialize to 0"); + assert_eq!(c.mh_entry_ts_zero[0], 0, "mh_entry_ts_zero must initialize to 0"); + assert_eq!(c.mh_ts_underflow[0], 0, "mh_ts_underflow must initialize to 0"); + assert_eq!(c.mh_elapsed_below_cap[0], 0, "mh_elapsed_below_cap must initialize to 0"); + assert_eq!(c.mh_would_fire[0], 0, "mh_would_fire must initialize to 0"); + assert_eq!(c.mh_force_flat_written[0], 0, "mh_force_flat_written must initialize to 0"); + assert_eq!(c.mh_force_flat_seen_by_seed[0], 0, "mh_force_flat_seen_by_seed must initialize to 0"); + Ok(()) +}