diag(ml-backtesting): max_hold enforcement counters localize why 86.5% of trades exceed the 60s cap

S2 cluster smoke ppcfk (29f7e923c) showed mean trade hold = 432s with
max_hold_ns=60s — 886/1024 trades exceed the cap that should hard-flat
them. Upload chain and pnl_track entry_ts write look structurally
correct on inspection. Need per-step counters to localize WHICH step
of the chain fails.

Adds 7 counters in stop_check_isv (kernel_calls, seen_zero, entry_ts_zero,
ts_underflow, elapsed_below_cap, would_fire, force_flat_written) plus 1
counter in seed_inflight_limits_batched (seed_saw_force_flat). Logged at
smoke end as a third nan_counters-like line.

Diagnostic only — no behavioral change.

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

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<const unsigned long long*>(
                 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;
     }
 

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;

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,

crates/ml-backtesting/src/sim/mod.rs

@@ -94,6 +94,32 @@ pub struct NanCountersV2 {
     pub last_bad_path: Vec<u32>,
 }
 
+/// 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<u32>,
+    /// max_hold_ns_per_b[b] == 0 (mode A: upload didn't land).
+    pub mh_seen_zero: Vec<u32>,
+    /// entry_ts == 0 despite position open (mode B: pnl_track didn't write).
+    pub mh_entry_ts_zero: Vec<u32>,
+    /// current_ts_ns < entry_ts — ts underflow (mode C: should never happen).
+    pub mh_ts_underflow: Vec<u32>,
+    /// elapsed < max_hold — check not yet due (mode D: normal in-flight).
+    pub mh_elapsed_below_cap: Vec<u32>,
+    /// elapsed >= max_hold — force-flat would fire (mode D-now or success path).
+    pub mh_would_fire: Vec<u32>,
+    /// market_targets[b*2+0] = 3 was actually written.
+    pub mh_force_flat_written: Vec<u32>,
+    /// seed_inflight_limits_batched read market_targets[b*2+0] == 3.
+    pub mh_force_flat_seen_by_seed: Vec<u32>,
+}
+
 pub struct LobSimCuda {
     n_backtests: usize,
     stream: Arc<CudaStream>,
@@ -206,6 +232,16 @@ pub struct LobSimCuda {
     vwap_session_gap_was_bad_d: CudaSlice<u32>,    // [n_backtests]
     last_bad_vwap_d:            CudaSlice<f32>,    // [n_backtests]
     last_bad_path_d:            CudaSlice<u32>,    // [n_backtests]
+
+    // S2.1: max_hold enforcement diagnostic counters (7 in stop_check_isv + 1 in seed_inflight).
+    mh_kernel_calls_d:            CudaSlice<u32>,  // [n_backtests]
+    mh_seen_zero_d:               CudaSlice<u32>,  // [n_backtests]
+    mh_entry_ts_zero_d:           CudaSlice<u32>,  // [n_backtests]
+    mh_ts_underflow_d:            CudaSlice<u32>,  // [n_backtests]
+    mh_elapsed_below_cap_d:       CudaSlice<u32>,  // [n_backtests]
+    mh_would_fire_d:              CudaSlice<u32>,  // [n_backtests]
+    mh_force_flat_written_d:      CudaSlice<u32>,  // [n_backtests]
+    mh_force_flat_seen_by_seed_d: CudaSlice<u32>,  // [n_backtests]
 }
 
 impl LobSimCuda {
@@ -439,6 +475,32 @@ impl LobSimCuda {
             .alloc_zeros::<u32>(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::<u32>(n_backtests)
+            .context("alloc mh_kernel_calls_d")?;
+        let mh_seen_zero_d = stream
+            .alloc_zeros::<u32>(n_backtests)
+            .context("alloc mh_seen_zero_d")?;
+        let mh_entry_ts_zero_d = stream
+            .alloc_zeros::<u32>(n_backtests)
+            .context("alloc mh_entry_ts_zero_d")?;
+        let mh_ts_underflow_d = stream
+            .alloc_zeros::<u32>(n_backtests)
+            .context("alloc mh_ts_underflow_d")?;
+        let mh_elapsed_below_cap_d = stream
+            .alloc_zeros::<u32>(n_backtests)
+            .context("alloc mh_elapsed_below_cap_d")?;
+        let mh_would_fire_d = stream
+            .alloc_zeros::<u32>(n_backtests)
+            .context("alloc mh_would_fire_d")?;
+        let mh_force_flat_written_d = stream
+            .alloc_zeros::<u32>(n_backtests)
+            .context("alloc mh_force_flat_written_d")?;
+        let mh_force_flat_seen_by_seed_d = stream
+            .alloc_zeros::<u32>(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<NanCountersMaxHold> {
+        let read = |slot: &CudaSlice<u32>| -> Result<Vec<u32>> {
+            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(&current_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(&current_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(&current_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)?;
         }

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(())
+}