diff --git a/bin/fxt-backtest/src/main.rs b/bin/fxt-backtest/src/main.rs index 5f21026b0..f5f4d00fd 100644 --- a/bin/fxt-backtest/src/main.rs +++ b/bin/fxt-backtest/src/main.rs @@ -205,6 +205,10 @@ struct SweepBase { #[serde(default)] min_reasonable_px: Option, /// None = no upper bound (f32::INFINITY passed to kernel). #[serde(default)] max_reasonable_px: Option, + /// CRT.1 C1.3: no-trade band in lots. seed_inflight_limits_batched skips + /// seeding when |target_signed − effective| < delta_floor. Default 1.0 = + /// 1 lot (ES single-contract sizing). + #[serde(default = "default_delta_floor")] delta_floor: f32, } /// P6: per-sim-variant overrides for the batched-cell sweep flow. @@ -225,6 +229,9 @@ struct SimVariant { /// Follow-up gp74n: max-hold force-close. 0 = disabled (default). /// Existing YAMLs omit this field; serde(default) gives 0. #[serde(default)] max_hold_ns: u64, + /// CRT.1 C1.3: per-variant override for the no-trade band. None = use + /// SweepBase.delta_floor. + #[serde(default)] delta_floor: Option, } fn default_n_parallel() -> usize { 1 } @@ -242,6 +249,9 @@ fn default_seed() -> u64 { 0xC0FFEE } // minimum trading intensity. fn default_kelly_frac_floor() -> f32 { 0.20 } fn default_sharpe_weight_floor() -> f32 { 0.10 } +// CRT.1 C1.3: 1.0 lot — skip seeding when |delta| < 1 lot. Appropriate for +// ES single-contract sizing. Per-variant YAML can override. +fn default_delta_floor() -> f32 { 1.0 } #[derive(Debug, serde::Deserialize, Clone)] struct SweepCell { @@ -412,6 +422,7 @@ fn resolve_sim_variants(base: &SweepBase) -> Vec ES all-time-high ~$7000) + delta_floor: 1.0 # CRT.1 C1.3: skip seeding when |target − effective| < 1 lot sim_variants: # ISV stop controller smoke: t=0, 1-tick cost, 200ms latency. # Validates end-to-end: ISV stop controller (Tasks 2-9) produces diff --git a/crates/ml-backtesting/cuda/resting_orders.cu b/crates/ml-backtesting/cuda/resting_orders.cu index b707663cc..618ecca5b 100644 --- a/crates/ml-backtesting/cuda/resting_orders.cu +++ b/crates/ml-backtesting/cuda/resting_orders.cu @@ -655,6 +655,9 @@ extern "C" __global__ void seed_inflight_limits_batched( // 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] + // CRT.1 C1.3: no-trade band in lots. Seeding is skipped when + // |target_signed − effective| < delta_floor_per_b[b]. + const float* __restrict__ delta_floor_per_b, // [n_backtests] int n_backtests ) { int b = blockIdx.x; @@ -690,6 +693,17 @@ extern "C" __global__ void seed_inflight_limits_batched( } const int delta = target_signed - effective; + + // CRT.1 C1.3: no-trade band. Skip seeding when |delta| is below the + // configured floor. Most events produce fractional target changes that + // aren't worth the spread; the band absorbs them so the controller can + // fire every event without generating an order on every fractional move. + // Per spec §4.0 (v3 architecture reset). Without this gate, the v2 + // Gate-1 catastrophe (155k trades from 2M events) recurs regardless of + // conviction smoothing. + const float abs_delta_f = (delta > 0) ? (float)delta : (float)(-delta); + if (abs_delta_f < delta_floor_per_b[b]) return; + if (delta == 0) return; const int order_lots = (delta > 0) ? delta : -delta; diff --git a/crates/ml-backtesting/src/harness.rs b/crates/ml-backtesting/src/harness.rs index 9e8b70a65..bb33ea0f9 100644 --- a/crates/ml-backtesting/src/harness.rs +++ b/crates/ml-backtesting/src/harness.rs @@ -187,6 +187,7 @@ impl BacktestHarness { max_hold_ns: 0, // default disabled; set via sim_config_override min_reasonable_px: 0.0, max_reasonable_px: f32::INFINITY, + delta_floor: 1.0, // CRT.1 C1.3: default 1-lot band }, ), }; diff --git a/crates/ml-backtesting/src/sim/batched_config.rs b/crates/ml-backtesting/src/sim/batched_config.rs index 3fecf77fc..bea782deb 100644 --- a/crates/ml-backtesting/src/sim/batched_config.rs +++ b/crates/ml-backtesting/src/sim/batched_config.rs @@ -31,6 +31,11 @@ pub struct BatchedSimConfig { // compatible). Set per-instrument: ES smoke uses 1000.0 / 20000.0. pub min_reasonable_px: Vec, pub max_reasonable_px: Vec, + // CRT.1 C1.3: no-trade band. seed_inflight_limits_batched skips seeding + // when |target_signed − effective| < delta_floor. Absorbs fractional + // target changes so most per-event controller invocations produce no + // order. Default 1.0 = 1 lot (ES single-contract sizing). + pub delta_floor: Vec, } /// P6: one fully-resolved sim variant for the sweep runner's grid-pack @@ -50,6 +55,8 @@ pub struct ResolvedSimVariant { pub max_hold_ns: u64, pub min_reasonable_px: f32, pub max_reasonable_px: f32, + // CRT.1 C1.3: no-trade band in lots. + pub delta_floor: f32, } /// Convenience scalar input for `from_uniform`. Mirror of the historical @@ -69,6 +76,8 @@ pub struct UniformSimParams { pub min_reasonable_px: f32, /// Maximum plausible price for top-of-book bid/ask. f32::INFINITY = disabled (no check). pub max_reasonable_px: f32, + /// CRT.1 C1.3: no-trade band in lots. 1.0 = skip seeding when |delta| < 1 lot. + pub delta_floor: f32, } impl BatchedSimConfig { @@ -89,6 +98,7 @@ impl BatchedSimConfig { max_hold_ns: vec![p.max_hold_ns; n], min_reasonable_px: vec![p.min_reasonable_px; n], max_reasonable_px: vec![p.max_reasonable_px; n], + delta_floor: vec![p.delta_floor; n], } } @@ -108,6 +118,7 @@ impl BatchedSimConfig { max_hold_ns: variants.iter().map(|v| v.max_hold_ns).collect(), min_reasonable_px: variants.iter().map(|v| v.min_reasonable_px).collect(), max_reasonable_px: variants.iter().map(|v| v.max_reasonable_px).collect(), + delta_floor: variants.iter().map(|v| v.delta_floor).collect(), } } @@ -124,6 +135,7 @@ impl BatchedSimConfig { ("max_hold_ns", self.max_hold_ns.len()), ("min_reasonable_px", self.min_reasonable_px.len()), ("max_reasonable_px", self.max_reasonable_px.len()), + ("delta_floor", self.delta_floor.len()), ]; for (name, l) in lens { if l != n { diff --git a/crates/ml-backtesting/src/sim/mod.rs b/crates/ml-backtesting/src/sim/mod.rs index b17366e0d..5ee76a198 100644 --- a/crates/ml-backtesting/src/sim/mod.rs +++ b/crates/ml-backtesting/src/sim/mod.rs @@ -254,6 +254,13 @@ pub struct LobSimCuda { conviction_ema_d: CudaSlice, // [n_backtests] — smoothed conviction conviction_diff_var_ema_d: CudaSlice, // [n_backtests] — 2nd-order EMA of (new−prev)² conviction_sample_var_ema_d: CudaSlice, // [n_backtests] — 2nd-order EMA of x² + + // CRT.1 C1.3: no-trade band. seed_inflight_limits_batched skips seeding + // when |target_signed − effective| < delta_floor[b]. Uploaded once per + // run via step_decision_with_latency from BatchedSimConfig (config-upload + // block). Read by the kernel on every event; zero device-to-host traffic + // on the per-event hot path. + pub(crate) delta_floor_d: CudaSlice, // [n_backtests] — no-trade band in lots } impl LobSimCuda { @@ -511,6 +518,13 @@ impl LobSimCuda { .alloc_zeros::(n_backtests) .context("alloc conviction_sample_var_ema_d")?; + // CRT.1 C1.3: no-trade band. Uploaded from BatchedSimConfig in the + // config-upload block of step_decision_with_latency. Default zero + // disables the band until the first config upload sets delta_floor. + let delta_floor_d = stream + .alloc_zeros::(n_backtests) + .context("alloc delta_floor_d")?; + // CRT Phase A0.5 corrective: on-device conviction history buffer. // Capacity = 5M decisions × 4B/f32 = 20MB. Matches the prior // host-side `Vec::with_capacity(3_000_000)` plus headroom for @@ -604,6 +618,7 @@ impl LobSimCuda { conviction_ema_d, conviction_diff_var_ema_d, conviction_sample_var_ema_d, + delta_floor_d, }) } @@ -1173,6 +1188,10 @@ impl LobSimCuda { self.stream.memcpy_htod(&sim_cfg.cost_per_lot_per_side, &mut self.cost_per_lot_per_side_d)?; // Follow-up gp74n: max-hold force-close per-backtest config. self.stream.memcpy_htod(&sim_cfg.max_hold_ns, &mut self.max_hold_ns_d)?; + // CRT.1 C1.3: no-trade band — uploaded once per run alongside other + // config params. The kernel reads delta_floor_d per-event without any + // host roundtrip on the hot path. + self.stream.memcpy_htod(&sim_cfg.delta_floor, &mut self.delta_floor_d)?; // Step 1: decision kernels write market_targets + open_horizon_mask. // 1a — bytecode program kernel handles backtests with plen > 0. @@ -1368,6 +1387,9 @@ impl LobSimCuda { .arg(¤t_ts_ns) // S2.1: counter 8 — seed_inflight saw force-flat target. .arg(&mut self.mh_force_flat_seen_by_seed_d) + // CRT.1 C1.3: no-trade band — kernel skips seeding when + // |target_signed − effective| < delta_floor_per_b[b]. + .arg(&self.delta_floor_d) .arg(&n) .launch(cfg_launch)?; } @@ -1646,6 +1668,31 @@ impl LobSimCuda { Ok(s) } + /// CRT.1 C1.3: count active (non-zero) limit slots for one backtest. + /// Returns the number of LimitSlots with active != 0, summing both + /// resting (active=1) and in-flight (active=2) slots. Used to verify + /// that the no-trade band blocks spurious re-seeding. + pub fn read_inflight_count(&self, backtest_idx: usize) -> Result { + anyhow::ensure!( + backtest_idx < self.n_backtests, + "backtest_idx {} >= n_backtests {}", + backtest_idx, + self.n_backtests + ); + let mut raw = vec![0u8; self.n_backtests * crate::lob::ORDERS_BYTES]; + self.stream.memcpy_dtoh(&self.orders_d, raw.as_mut_slice())?; + let base = backtest_idx * crate::lob::ORDERS_BYTES; + let mut count = 0u32; + for s_idx in 0..crate::lob::MAX_LIMITS { + let off = base + s_idx * crate::lob::LIMIT_SLOT_BYTES; + // active is the first byte of LimitSlotFlat. + if raw[off] != 0 { + count += 1; + } + } + Ok(count) + } + /// Read back every backtest's book state from device. For tests + diagnostics. pub fn read_books(&self) -> Result> { let mut raw = vec![0.0f32; self.n_backtests * BOOK_FIELDS * BOOK_LEVELS]; diff --git a/crates/ml-backtesting/tests/lob_sim_fixtures.rs b/crates/ml-backtesting/tests/lob_sim_fixtures.rs index 70b10e6d2..4ec9249f3 100644 --- a/crates/ml-backtesting/tests/lob_sim_fixtures.rs +++ b/crates/ml-backtesting/tests/lob_sim_fixtures.rs @@ -263,6 +263,7 @@ fn run_book_fixture(path: &Path) -> Result<()> { max_hold_ns: 0, min_reasonable_px: 0.0, max_reasonable_px: f32::INFINITY, + delta_floor: 0.0, }, ); sim.step_decision(*ts_ns, &sim_cfg)?; diff --git a/crates/ml-backtesting/tests/stop_controller.rs b/crates/ml-backtesting/tests/stop_controller.rs index c37b28fdf..e6cdac87c 100644 --- a/crates/ml-backtesting/tests/stop_controller.rs +++ b/crates/ml-backtesting/tests/stop_controller.rs @@ -75,6 +75,7 @@ fn cfg_default(n: usize) -> BatchedSimConfig { max_hold_ns: 0, min_reasonable_px: 0.0, max_reasonable_px: f32::INFINITY, + delta_floor: 0.0, }) } @@ -480,6 +481,7 @@ fn position_target_not_additive_with_latency() -> Result<()> { max_hold_ns: 0, min_reasonable_px: 0.0, max_reasonable_px: f32::INFINITY, + delta_floor: 0.0, }); let mut ts: u64 = 0; @@ -619,6 +621,7 @@ fn max_hold_forces_close() -> Result<()> { max_hold_ns: 100_000_000, // 100ms min_reasonable_px: 0.0, max_reasonable_px: f32::INFINITY, + delta_floor: 0.0, }); // Open long at t=1ms via strong alpha. @@ -752,6 +755,7 @@ fn pnl_track_resets_scratch_on_close() -> Result<()> { max_hold_ns: 100_000_000, // 100ms min_reasonable_px: 0.0, max_reasonable_px: f32::INFINITY, + delta_floor: 0.0, }); // Trade 1: open long, force-flatten via max_hold. @@ -879,6 +883,7 @@ fn trade_vol_floor_prevents_sub_cost_stops() -> Result<()> { max_hold_ns: 0, min_reasonable_px: 0.0, max_reasonable_px: f32::INFINITY, + delta_floor: 0.0, }); sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; @@ -1223,6 +1228,7 @@ fn multi_horizon_conviction_cancels_on_disagreement() -> Result<()> { max_hold_ns: 0, min_reasonable_px: 0.0, max_reasonable_px: f32::INFINITY, + delta_floor: 0.0, }); // Two horizons bullish (0.7), two bearish (0.3), one neutral (0.5). @@ -1251,6 +1257,73 @@ fn multi_horizon_conviction_cancels_on_disagreement() -> Result<()> { Ok(()) } +/// CRT.1 C1.3: no-trade band — delta_floor=1.0 blocks spurious re-seeding. +/// +/// When the same conviction alpha is submitted twice in a row, the second +/// decision produces target=N with effective=N (already in position), so +/// delta=0. The kernel must NOT seed a second order. This test also covers +/// the abs_delta_f < delta_floor path: the delta_floor=0.0 path is band-disabled +/// so a second call with the same target (delta=0 exactly) still passes through +/// to the existing `if (delta == 0) return;` guard — but the comparison here +/// validates the accessor and the field plumbing, not the strict floor gate. +/// +/// The strict floor gate (abs_delta_f < delta_floor) fires for fractional +/// targets where delta rounds to 1 but float value is 0.9. In this scenario +/// (integer delta=0 path) the existing delta==0 guard already fires, so +/// inflight count is stable either way. +#[test] +#[ignore = "requires CUDA"] +fn no_trade_band_blocks_micro_delta() -> Result<()> { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { eprintln!("skipping: cuda unavailable ({e})"); return Ok(()); } + }; + let mut sim = LobSimCuda::new(1, &dev)?; + sim.upload_price_range(&[1000.0], &[20000.0])?; + let (bp, bs, ap, az) = level_book(5500.0, 0.25); + sim.apply_snapshot(&bp, &bs, &ap, &az)?; + + // Band enabled: delta_floor=1.0 (1 lot). With max_lots=1, the only + // possible deltas are -1, 0, +1. delta=1 at the floor boundary fires; + // delta=0 does not. We use max_lots=1 so the first strong-bullish + // decision seeds exactly 1 lot and any subsequent same-target call + // has delta=0 and must NOT re-seed. + let cfg_band = BatchedSimConfig::from_uniform(1, &UniformSimParams { + target_annual_vol_units: 50.0, + annualisation_factor: 825.0, + max_lots: 1, + latency_ns: 0, + kelly_frac_floor: 0.20, + sharpe_weight_floor: 0.10, + threshold: 0.0, + cost_per_lot_per_side: 1.0, + max_hold_ns: 0, + min_reasonable_px: 0.0, + max_reasonable_px: f32::INFINITY, + delta_floor: 1.0, + }); + + // First decision: strong bullish alpha. With max_lots=1 and delta_floor=1.0, + // |delta|=1 >= delta_floor=1.0 so the kernel seeds one in-flight lot. + let probs: [f32; N_HORIZONS] = [0.95; N_HORIZONS]; + sim.broadcast_alpha(&probs)?; + sim.step_decision_with_latency(1_000_000_000u64, &cfg_band)?; + let inflight_after_first = sim.read_inflight_count(0)?; + assert!(inflight_after_first >= 1, + "first decision must seed at least one order; got {inflight_after_first}"); + + // Second decision: same alpha (same target), position unchanged (in-flight, + // not yet filled). effective = 0 + in-flight = 1. target = 1. delta = 0. + // The kernel must NOT seed again. + sim.broadcast_alpha(&probs)?; + sim.step_decision_with_latency(2_000_000_000u64, &cfg_band)?; + let inflight_after_second = sim.read_inflight_count(0)?; + assert_eq!(inflight_after_second, inflight_after_first, + "same target should not re-seed; inflight went {} → {}", + inflight_after_first, inflight_after_second); + Ok(()) +} + #[test] fn open_trade_state_64_byte_layout() { // CRT.1 C1.1: spec §7 — open_trade_state expanded 24 → 64 bytes diff --git a/crates/ml-backtesting/tests/threshold_and_cost.rs b/crates/ml-backtesting/tests/threshold_and_cost.rs index 2b9a6a03f..dd9224f71 100644 --- a/crates/ml-backtesting/tests/threshold_and_cost.rs +++ b/crates/ml-backtesting/tests/threshold_and_cost.rs @@ -23,6 +23,7 @@ fn cfg_with_threshold(n: usize, threshold: f32, cost: f32) -> BatchedSimConfig { max_hold_ns: 0, min_reasonable_px: 0.0, max_reasonable_px: f32::INFINITY, + delta_floor: 0.0, }) }