diff --git a/crates/ml-backtesting/cuda/decision_policy.cu b/crates/ml-backtesting/cuda/decision_policy.cu index aa0834465..09caef7bc 100644 --- a/crates/ml-backtesting/cuda/decision_policy.cu +++ b/crates/ml-backtesting/cuda/decision_policy.cu @@ -47,57 +47,19 @@ __device__ static int stop_check_isv( float* __restrict__ trail_hwm, const Book* __restrict__ books, const float* __restrict__ cost_per_lot_per_side, - 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, - // 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 + // S2.1: mh_kernel_calls counts stop_check_isv entries with non-zero position. + // Kept as an ongoing generic diagnostic (always-non-decreasing means SL/trail + // are being evaluated). The 6 max_hold-specific counters were removed when + // max_hold enforcement moved to resting_orders_step (event-rate). + unsigned int* __restrict__ mh_kernel_calls ) { 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). + // S2.1: count every entry with a non-zero position. 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) { - // 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) { - // 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; - } - } - // Spec §5 controller math — SL only in this task; trail in Task 5. const float atr = atr_mid_ema[b]; const unsigned int mask = open_horizon_masks[b]; @@ -188,18 +150,8 @@ extern "C" __global__ void decision_policy_default( const float* __restrict__ atr_mid_ema_per_b, // [n_backtests] float* __restrict__ trail_hwm_per_b, // [n_backtests] — read + write const Book* __restrict__ books_per_b, // [n_backtests * BOOK_LEVELS] (via lob_state.cuh) - // Follow-up gp74n: max-hold force-close. - 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. + // S2.1: mh_kernel_calls diagnostic counter (non-zero-position entries). 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; @@ -210,12 +162,8 @@ extern "C" __global__ void decision_policy_default( if (stop_check_isv(b, positions, isv_kelly_base, open_horizon_masks, 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, - 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)) { + mh_kernel_calls_per_b)) { return; } @@ -390,18 +338,8 @@ extern "C" __global__ void decision_policy_program( const float* __restrict__ atr_mid_ema_per_b, // [n_backtests] float* __restrict__ trail_hwm_per_b, // [n_backtests] — read + write const Book* __restrict__ books_per_b, // [n_backtests * BOOK_LEVELS] - // Follow-up gp74n: max-hold force-close. - 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. + // S2.1: mh_kernel_calls diagnostic counter (non-zero-position entries). 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; @@ -417,12 +355,8 @@ extern "C" __global__ void decision_policy_program( if (stop_check_isv(b, positions, isv_kelly_base, open_horizon_masks, 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, - 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)) { + mh_kernel_calls_per_b)) { return; } diff --git a/crates/ml-backtesting/cuda/resting_orders.cu b/crates/ml-backtesting/cuda/resting_orders.cu index b224d2edb..501fe1e16 100644 --- a/crates/ml-backtesting/cuda/resting_orders.cu +++ b/crates/ml-backtesting/cuda/resting_orders.cu @@ -248,6 +248,9 @@ 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] + // S2.2: max_hold force-close at event rate. + const unsigned long long* __restrict__ max_hold_ns_per_b, // [n_backtests] + const unsigned char* __restrict__ open_trade_state, // [n_backtests * 24] // S1.20: NaN instrumentation counters (per-backtest). unsigned int* __restrict__ nan_avg_px, unsigned int* __restrict__ nan_realised, @@ -305,6 +308,23 @@ extern "C" __global__ void resting_orders_step( } last_event_ts[b] = current_ts_ns; + // S2.2: max_hold force-close at event rate. + // Mirrors the decision-rate check that used to live in stop_check_isv (removed), + // but runs every event so the 60s cap is actually enforceable. The previous + // decision-rate check fired every decision_stride events (~13 min sim-time on ES + // data at stride=200), so it was ~13 min LATE on average. Event-rate enforces + // within one event of the cap. + const unsigned long long max_hold = max_hold_ns_per_b[b]; + if (max_hold > 0ull && pos.position_lots != 0) { + 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) { + pos.position_lots = 0; // force-flat; pnl_track close fires next step + } + } + // (1) Promote in-flight limits to resting; (2) queue-decay; (3) marketability check. const float trade_abs = (trade_signed_vol < 0.0f) ? -trade_signed_vol : trade_signed_vol; for (int i = 0; i < MAX_LIMITS; ++i) { diff --git a/crates/ml-backtesting/src/harness.rs b/crates/ml-backtesting/src/harness.rs index a0b82d113..fdce15155 100644 --- a/crates/ml-backtesting/src/harness.rs +++ b/crates/ml-backtesting/src/harness.rs @@ -328,24 +328,21 @@ 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. + // S2.1/S2.2: generic stop-controller diagnostic counters. The 6 + // max_hold-specific decision-rate counters were removed when max_hold + // enforcement moved to resting_orders_step (event-rate). Two remain: + // kernel_calls — stop_check_isv entries with open position (SL/trail active) + // seed_saw_force_flat — force-flat signals seen by seed_inflight (SL/trail) 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={}", + "stop_ctrl_counters: kernel_calls={} 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}"), + Err(e) => eprintln!("stop_ctrl_counters read failed: {e}"), } Ok(RunStats { diff --git a/crates/ml-backtesting/src/sim/mod.rs b/crates/ml-backtesting/src/sim/mod.rs index 9009055fc..844fe254c 100644 --- a/crates/ml-backtesting/src/sim/mod.rs +++ b/crates/ml-backtesting/src/sim/mod.rs @@ -94,28 +94,19 @@ pub struct NanCountersV2 { pub last_bad_path: Vec, } -/// S2.1: max_hold enforcement diagnostic counters. +/// S2.1/S2.2: 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. +/// After S2.2, max_hold enforcement lives in `resting_orders_step` (event-rate). +/// The decision-rate max_hold code in `stop_check_isv` was removed. Two counters +/// remain as ongoing generics: +/// - `mh_kernel_calls`: every `stop_check_isv` entry with position_lots != 0 +/// (SL/trail are being evaluated). Should be non-zero whenever a position is +/// open and a decision fires. +/// - `mh_force_flat_seen_by_seed`: every `seed_inflight_limits_batched` entry +/// that reads market_targets[b*2+0] == 3 (force-flat from SL or trail). 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, } @@ -233,14 +224,8 @@ pub struct LobSimCuda { 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). + // S2.1/S2.2: max_hold diagnostic counters (2 remain after S2.2 prune). 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] } @@ -475,28 +460,10 @@ impl LobSimCuda { .alloc_zeros::(n_backtests) .context("alloc last_bad_path_d")?; - // S2.1: max_hold enforcement diagnostic counters, all zero-initialized. + // S2.1/S2.2: max_hold diagnostic counters (2 remain after S2.2 prune). 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")?; @@ -574,12 +541,6 @@ impl LobSimCuda { 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, }) } @@ -775,12 +736,14 @@ impl LobSimCuda { }) } - /// S2.1: max_hold enforcement diagnostic counters. + /// S2.1/S2.2: max_hold 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. + /// After S2.2 the 6 max_hold-specific decision-rate counters were removed. + /// Two generic counters remain: + /// - `mh_kernel_calls`: stop_check_isv entries with position_lots != 0 + /// (SL/trail evaluation is running; non-zero when open + decision fires). + /// - `mh_force_flat_seen_by_seed`: force-flat signals (SL or trail) seen + /// by seed_inflight_limits_batched. pub fn read_max_hold_counters(&self) -> Result { let read = |slot: &CudaSlice| -> Result> { let mut buf = vec![0u32; self.n_backtests]; @@ -789,12 +752,6 @@ impl LobSimCuda { }; 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)?, }) } @@ -1107,17 +1064,8 @@ impl LobSimCuda { .arg(&self.atr_mid_ema_d) .arg(&mut self.trail_hwm_d) .arg(&self.books_d) - .arg(&self.max_hold_ns_d) - .arg(¤t_ts_ns) - .arg(&self.open_trade_state_d) - // S2.1: max_hold diagnostic counters. + // S2.1: mh_kernel_calls diagnostic counter. .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)?; } @@ -1141,17 +1089,8 @@ impl LobSimCuda { .arg(&self.atr_mid_ema_d) .arg(&mut self.trail_hwm_d) .arg(&self.books_d) - .arg(&self.max_hold_ns_d) - .arg(¤t_ts_ns) - .arg(&self.open_trade_state_d) - // S2.1: max_hold diagnostic counters. + // S2.1: mh_kernel_calls diagnostic counter. .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)?; } @@ -1492,6 +1431,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) + // S2.2: max_hold force-close at event rate. + .arg(&self.max_hold_ns_d) + .arg(&self.open_trade_state_d) .arg(&mut self.nan_avg_px_d) .arg(&mut self.nan_realised_d) .arg(&mut self.nan_realized_pnl_d) diff --git a/crates/ml-backtesting/tests/stop_controller.rs b/crates/ml-backtesting/tests/stop_controller.rs index 1c52dbf4f..c138acee5 100644 --- a/crates/ml-backtesting/tests/stop_controller.rs +++ b/crates/ml-backtesting/tests/stop_controller.rs @@ -626,11 +626,15 @@ fn trail_hwm_reset_on_close() -> Result<()> { /// Validates that max_hold_ns fires a force-flat when the elapsed hold /// time exceeds the configured limit, and does NOT fire before that limit. /// +/// S2.2: max_hold enforcement moved to resting_orders_step (event-rate). +/// The test now exercises the event-rate path via step_resting_orders +/// rather than checking the market_target written by stop_check_isv. +/// /// Test geometry: /// - max_hold_ns = 100ms = 100_000_000ns /// - Open long at t=1ms via strong alpha -/// - At entry_ts + 50ms: still under limit → must NOT fire (side != 3) -/// - At entry_ts + 150ms: past limit → must fire force-flat (side=3, size=0) +/// - At entry_ts + 50ms: step_resting_orders at 50ms must NOT close (elapsed < cap) +/// - At entry_ts + 150ms: step_resting_orders at 150ms must close (elapsed >= cap) #[test] #[ignore = "requires CUDA"] fn max_hold_forces_close() -> Result<()> { @@ -679,24 +683,25 @@ fn max_hold_forces_close() -> Result<()> { assert!(sim.read_pos(0)?.position_lots > 0, "setup: long opens"); let entry_ts = ts; - // At entry_ts + 50ms — still under 100ms hold. Book stable so SL/trail don't fire. + // At entry_ts + 50ms — still under 100ms hold. step_resting_orders must NOT close. ts = entry_ts + 50_000_000; let (bp3, bs3, ap3, az3) = level_book(5500.1, 0.25); sim.apply_snapshot(&bp3, &bs3, &ap3, &az3)?; - sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; - sim.step_decision_with_latency(ts, &cfg_max_hold)?; - let (side_pre, _) = sim.read_market_target(0)?; - assert_ne!(side_pre, 3, "elapsed 50ms < max_hold=100ms: must NOT fire; got side={side_pre}"); + sim.step_resting_orders(ts, 0.0)?; + let pos_pre = sim.read_pos(0)?; + assert_ne!(pos_pre.position_lots, 0, + "elapsed 50ms < max_hold=100ms: resting_orders must NOT close; position_lots={}", + pos_pre.position_lots); - // At entry_ts + 150ms — past max_hold. + // At entry_ts + 150ms — past max_hold. step_resting_orders must force-close. ts = entry_ts + 150_000_000; let (bp4, bs4, ap4, az4) = level_book(5500.1, 0.25); sim.apply_snapshot(&bp4, &bs4, &ap4, &az4)?; - sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; - sim.step_decision_with_latency(ts, &cfg_max_hold)?; - let (side_post, size_post) = sim.read_market_target(0)?; - assert_eq!(side_post, 3, "elapsed 150ms > max_hold=100ms: must fire force-flat; got side={side_post}"); - assert_eq!(size_post, 0); + sim.step_resting_orders(ts, 0.0)?; + let pos_post = sim.read_pos(0)?; + assert_eq!(pos_post.position_lots, 0, + "elapsed 150ms > max_hold=100ms: resting_orders must force-close; position_lots={}", + pos_post.position_lots); Ok(()) } @@ -1224,13 +1229,8 @@ fn max_hold_counters_initialize_to_zero() -> Result<()> { }; let sim = LobSimCuda::new(1, &dev)?; let c = sim.read_max_hold_counters()?; + // S2.2: 6 decision-rate max_hold counters removed; 2 generic ones remain. 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(()) }