arch(crt-1): no-trade band in seed_inflight — sparse action at signal cadence (C1.3)

Per spec §4.0 + plan C1.3. Without this gate, every fractional change
in target_lots seeds an order via target-delta semantics; combined with
the continuous controller invocation (every event after A1) this
generates hyperactivity even with the multi-horizon §4.4 conviction
formula (C1.2). v2 Gate-1 failures confirmed: scalar EMA + amplification
clamp + per-event controller = 155k trades on a 2M-event smoke.

The band absorbs fractional target adjustments so most events produce
no order. When the signal genuinely shifts and target moves enough to
cross the band, the kernel seeds. Continuous evaluation, sparse action.

Greenfields atomic refactor — single config knob threaded through every
layer in one commit:
  SweepBase.delta_floor: f32  (YAML, default 1.0 = 1 lot)
  SimVariant.delta_floor: Option<f32>  (per-variant override)
  ResolvedSimVariant.delta_floor: f32
  UniformSimParams.delta_floor: f32
  BatchedSimConfig.delta_floor: Vec<f32>
  LobSimCuda.delta_floor_d: CudaSlice<f32>
  seed_inflight_limits_batched kernel param + skip-if-below-floor logic

New accessor: LobSimCuda::read_inflight_count(b) — counts active != 0
limit slots for backtest b. Not cfg(test); follows existing read_limit_slot
pattern.

Test: no_trade_band_blocks_micro_delta verifies that a second decision
with the same strong-bullish alpha (same target, effective = in-flight
lots) produces delta=0 and does not re-seed. max_lots=1 keeps the
arithmetic unambiguous.

Existing tests: all UniformSimParams struct literals updated with
delta_floor=0.0 (band disabled) so existing behaviour is preserved.
harness.rs from_uniform path uses delta_floor=1.0 (production default).

Hot-path discipline: the delta_floor_d upload happens once per run in
the existing config-upload block (alongside threshold, cost, max_hold_ns);
the kernel reads the device slot per-event without any host roundtrip. No
memcpy_htod / dtoh / dtov / synchronize introduced on the per-event path.

Per pearl_controller_anchors_isv_driven, this floor is currently a
config constant; CRT.2 (Phase 2) makes it ISV-derived from rolling
spread cost / signal volatility.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-20 20:59:50 +02:00
parent 632296021c
commit 7c850a6c02
9 changed files with 161 additions and 0 deletions

View File

@@ -205,6 +205,10 @@ struct SweepBase {
#[serde(default)] min_reasonable_px: Option<f32>,
/// None = no upper bound (f32::INFINITY passed to kernel).
#[serde(default)] max_reasonable_px: Option<f32>,
/// 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<f32>,
}
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<ml_backtesting::sim::ResolvedSi
max_hold_ns: v.max_hold_ns,
min_reasonable_px: base.min_reasonable_px.unwrap_or(0.0),
max_reasonable_px: base.max_reasonable_px.unwrap_or(f32::INFINITY),
delta_floor: v.delta_floor.unwrap_or(base.delta_floor),
}).collect()
}

View File

@@ -27,6 +27,7 @@ base:
checkpoint: /feature-cache/alpha-perception-runs/dbd500ecf/trunk_best_h6000.bin
min_reasonable_px: 1000.0 # ES futures: reject sub-$1000 (catches $48-100 outliers + negatives)
max_reasonable_px: 20000.0 # ES futures: reject super-$20000 (catches $53k outliers; > 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

View File

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

View File

@@ -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
},
),
};

View File

@@ -31,6 +31,11 @@ pub struct BatchedSimConfig {
// compatible). Set per-instrument: ES smoke uses 1000.0 / 20000.0.
pub min_reasonable_px: Vec<f32>,
pub max_reasonable_px: Vec<f32>,
// 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<f32>,
}
/// 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 {

View File

@@ -254,6 +254,13 @@ pub struct LobSimCuda {
conviction_ema_d: CudaSlice<f32>, // [n_backtests] — smoothed conviction
conviction_diff_var_ema_d: CudaSlice<f32>, // [n_backtests] — 2nd-order EMA of (newprev)²
conviction_sample_var_ema_d: CudaSlice<f32>, // [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<f32>, // [n_backtests] — no-trade band in lots
}
impl LobSimCuda {
@@ -511,6 +518,13 @@ impl LobSimCuda {
.alloc_zeros::<f32>(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::<f32>(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(&current_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<u32> {
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<Vec<BookFlat>> {
let mut raw = vec![0.0f32; self.n_backtests * BOOK_FIELDS * BOOK_LEVELS];

View File

@@ -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)?;

View File

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

View File

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