gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

fix(ml-backtesting): move max_hold force-close from decision-rate to event-rate — actual cap enforcement

Browse Source 
S2.1 instrumentation (smoke cqpph @ 95a77c4ac) revealed the chain
worked end-to-end: max_hold check fired 1661 times, force-flat target
was written and seen by seed_inflight. Yet 86.5% of trades exceeded
the 60s cap with mean hold = 432s. Root cause: decision_policy's
stop_check_isv only runs every decision_stride events. At
decision_stride=200 on ES Q1 (2M events / ~91 days = ~3.9s sim-time
per event), decisions are ~13 min apart in sim-time, so the cap is
enforced 13 min late on average.

Fix: move the max_hold check to resting_orders_step (event-rate, runs
every iteration), same pattern as the session-gap force-close at
resting_orders.cu:282. Thread max_hold_ns_per_b and open_trade_state
into resting_orders_step. SL/trail stay in stop_check_isv because
they depend on ISV controller state at decision-rate.

Remove the dead decision-rate max_hold code and its 6 diagnostic
counter params from stop_check_isv per single-source-of-truth and
no-hiding rules. Keep mh_kernel_calls and mh_force_flat_seen_by_seed
counters as ongoing generic diagnostics. Update harness.rs log line
and stop_controller tests to exercise the event-rate path.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-20 16:07:08 +02:00
parent 95a77c4ac6
commit b92bd72c72
5 changed files with 79 additions and 186 deletions
Download Patch File Download Diff File Expand all files Collapse all files

86
crates/ml-backtesting/cuda/decision_policy.cu
View File

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

20
crates/ml-backtesting/cuda/resting_orders.cu
View File

@@ -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<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) {
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) {

17
crates/ml-backtesting/src/harness.rs
View File

@@ -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 {

104
crates/ml-backtesting/src/sim/mod.rs
View File

@@ -94,28 +94,19 @@ pub struct NanCountersV2 {
pub last_bad_path: Vec<u32>,
}
/// 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<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>,
}
@@ -233,14 +224,8 @@ pub struct LobSimCuda {
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).
// S2.1/S2.2: max_hold diagnostic counters (2 remain after S2.2 prune).
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]
}
@@ -475,28 +460,10 @@ impl LobSimCuda {
.alloc_zeros::<u32>(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::<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")?;
@@ -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<NanCountersMaxHold> {
let read = |slot: &CudaSlice<u32>| -> Result<Vec<u32>> {
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(&current_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(&current_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)

38
crates/ml-backtesting/tests/stop_controller.rs
View File

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