diff --git a/bin/fxt-backtest/src/main.rs b/bin/fxt-backtest/src/main.rs index 3a9ebccf5..21f9ec346 100644 --- a/bin/fxt-backtest/src/main.rs +++ b/bin/fxt-backtest/src/main.rs @@ -393,6 +393,8 @@ fn run(args: RunArgs) -> Result<()> { latency_ns: args.latency_ns, kelly_frac_floor: args.kelly_frac_floor, sharpe_weight_floor: args.sharpe_weight_floor, + threshold: 0.0, // P4: default = gate disabled + cost_per_lot_per_side: 0.0, // P4: default = frictionless strategies: strategy_grid, }; diff --git a/crates/ml-backtesting/cuda/decision_policy.cu b/crates/ml-backtesting/cuda/decision_policy.cu index e4c4df8cb..615655a8c 100644 --- a/crates/ml-backtesting/cuda/decision_policy.cu +++ b/crates/ml-backtesting/cuda/decision_policy.cu @@ -47,12 +47,33 @@ extern "C" __global__ void decision_policy_default( // to the old sentinel-skip behaviour. const float* __restrict__ kelly_frac_floor_per_b, // [n_backtests] const float* __restrict__ sharpe_weight_floor_per_b, // [n_backtests] + // P4: threshold gate (absolute max-conviction cutoff, pre-Kelly). + const float* __restrict__ threshold_per_b, // [n_backtests] int n_backtests ) { int b = blockIdx.x; if (b >= n_backtests || threadIdx.x != 0) return; if (program_lens[b] != 0u) return; // backtest b uses a custom bytecode program; skip default + // P4: threshold gate. Skip entirely if max conviction below threshold. + // Pre-Kelly placement keeps the gate deterministic from alpha alone, + // so the threshold pre-registration step (compute p60-p95 absolute + // values from observed convictions) reflects exactly what gets gated + // in deployment. + { + float max_conv = 0.0f; + #pragma unroll + for (int h = 0; h < N_HORIZONS; ++h) { + max_conv = fmaxf(max_conv, fabsf(alpha_probs[h] - 0.5f) * 2.0f); + } + if (max_conv < threshold_per_b[b]) { + market_targets[b * 2 + 0] = 2; + market_targets[b * 2 + 1] = 0; + open_horizon_masks[b] = 0u; + return; + } + } + // Per-backtest scalar reads (host-uploaded arrays). const float target_annual_vol_units = target_annual_vol_units_per_b[b]; const float annualisation_factor = annualisation_factor_per_b[b]; @@ -198,6 +219,8 @@ extern "C" __global__ void decision_policy_program( const int* __restrict__ max_lots_per_b, const float* __restrict__ kelly_frac_floor_per_b, const float* __restrict__ sharpe_weight_floor_per_b, + // P4: threshold gate (same as decision_policy_default). + const float* __restrict__ threshold_per_b, int n_backtests ) { int b = blockIdx.x; @@ -208,6 +231,22 @@ extern "C" __global__ void decision_policy_program( // Skip — this backtest uses the hardcoded default kernel separately. return; } + + // P4: threshold gate. + { + float max_conv = 0.0f; + #pragma unroll + for (int h = 0; h < N_HORIZONS; ++h) { + max_conv = fmaxf(max_conv, fabsf(alpha_probs[h] - 0.5f) * 2.0f); + } + if (max_conv < threshold_per_b[b]) { + market_targets[b * 2 + 0] = 2; + market_targets[b * 2 + 1] = 0; + open_horizon_masks[b] = 0u; + return; + } + } + // Per-backtest scalar reads (host-uploaded arrays). const float target_annual_vol_units = target_annual_vol_units_per_b[b]; const float annualisation_factor = annualisation_factor_per_b[b]; diff --git a/crates/ml-backtesting/cuda/resting_orders.cu b/crates/ml-backtesting/cuda/resting_orders.cu index 74ae70a82..97560ad76 100644 --- a/crates/ml-backtesting/cuda/resting_orders.cu +++ b/crates/ml-backtesting/cuda/resting_orders.cu @@ -87,7 +87,25 @@ __device__ static void cancel_oco_pair(Orders& orders, unsigned char oco_pair) { // Apply a fill to per-backtest position state. Same VWAP-scale-in / // counter-direction realised-P&L rule as order_match.cu's market-order // path; isolated here as a __device__ helper. -__device__ static void apply_fill_to_pos(Pos& pos, float filled_lots, float total_cost, unsigned char side) { +// +// P4: per-fill cost integration. After the close-leg realized_pnl math +// runs (so the unwind P&L is computed from gross prices), deduct the +// fill cost from realized_pnl AND accumulate into total_fees_per_b[b]. +// Net-of-cost semantics: isv_kelly_update_on_close reads realized_pnl +// delta which is now net of cost — Kelly state learns realistic returns. +// +// Plain `+=` for total_fees_per_b is safe under the single-writer-per-block +// convention enforced by `if (threadIdx.x != 0) return;` at the top of +// every block-per-backtest kernel. +__device__ static void apply_fill_to_pos( + Pos& pos, + float filled_lots, + float total_cost, + unsigned char side, + int b, + const float* __restrict__ cost_per_lot_per_side_per_b, + float* __restrict__ total_fees_per_b +) { if (filled_lots <= 0.0f) return; const float avg_px = total_cost / filled_lots; const float sgn = (side == 0) ? 1.0f : -1.0f; @@ -110,6 +128,11 @@ __device__ static void apply_fill_to_pos(Pos& pos, float filled_lots, float tota if (filled_lots > prev_abs) { pos.vwap_entry = avg_px; } } pos.position_lots = (int)new_pos_f; + + // P4: per-fill cost deduction (single-writer-per-block; plain += safe). + const float fill_cost = filled_lots * cost_per_lot_per_side_per_b[b]; + pos.realized_pnl -= fill_cost; + total_fees_per_b[b] += fill_cost; } // Free-slot search. Returns slot index 0..31 or -1 if all in use. @@ -136,6 +159,9 @@ extern "C" __global__ void resting_orders_step( int pos_bytes, int ring_cap_events, float tick_size, + // P4: per-fill cost integration (per-backtest arrays). + const float* __restrict__ cost_per_lot_per_side_per_b, + float* __restrict__ total_fees_per_b, int n_backtests ) { int b = blockIdx.x; @@ -182,7 +208,8 @@ extern "C" __global__ void resting_orders_step( if (take > 0.0f) { // Filled `take` lots at our limit price (queue position reached us). const float cost = take * s.price; - apply_fill_to_pos(pos, take, cost, s.side); + apply_fill_to_pos(pos, take, cost, s.side, + b, cost_per_lot_per_side_per_b, total_fees_per_b); s.size_remaining -= take; emit_audit(audit_base, audit_head, b, (unsigned int)ring_cap_events, current_ts_ns, @@ -211,7 +238,8 @@ extern "C" __global__ void resting_orders_step( ? walk_ask_for_buy (book, s.size_remaining, filled) : walk_bid_for_sell(book, s.size_remaining, filled); if (filled > 0.0f) { - apply_fill_to_pos(pos, filled, cost, s.side); + apply_fill_to_pos(pos, filled, cost, s.side, + b, cost_per_lot_per_side_per_b, total_fees_per_b); s.size_remaining -= filled; emit_audit(audit_base, audit_head, b, (unsigned int)ring_cap_events, current_ts_ns, @@ -256,7 +284,8 @@ extern "C" __global__ void resting_orders_step( ? walk_ask_for_buy (book, st.size, filled) : walk_bid_for_sell(book, st.size, filled); if (filled > 0.0f) { - apply_fill_to_pos(pos, filled, cost, st.side); + apply_fill_to_pos(pos, filled, cost, st.side, + b, cost_per_lot_per_side_per_b, total_fees_per_b); emit_audit(audit_base, audit_head, b, (unsigned int)ring_cap_events, current_ts_ns, pack_slot_tag(1 /*SlotKind::Stop*/, (unsigned char)j, st.gen_counter), diff --git a/crates/ml-backtesting/src/harness.rs b/crates/ml-backtesting/src/harness.rs index 69526e440..e2ef0e7fc 100644 --- a/crates/ml-backtesting/src/harness.rs +++ b/crates/ml-backtesting/src/harness.rs @@ -56,6 +56,14 @@ pub struct BacktestHarnessConfig { /// sum produce a non-zero size before recent_sharpe is populated. /// Default 0.10 = uniform 1/5 weight per horizon at cold-start. pub sharpe_weight_floor: f32, + /// P4: threshold gate — absolute conviction cutoff. When + /// `max_h |alpha[h] - 0.5| * 2 < threshold`, the decision kernels + /// emit noop. Default 0.0 = gate disabled (every signal passes). + pub threshold: f32, + /// P4: per-fill cost (price-units / lot / side). Deducted from + /// pos.realized_pnl and accumulated into total_fees_per_b_d at + /// every fill. Default 0.0 = no cost (frictionless). + pub cost_per_lot_per_side: f32, } pub struct BacktestHarness { @@ -144,6 +152,8 @@ impl BacktestHarness { latency_ns: cfg.latency_ns, kelly_frac_floor: cfg.kelly_frac_floor, sharpe_weight_floor: cfg.sharpe_weight_floor, + threshold: cfg.threshold, + cost_per_lot_per_side: cfg.cost_per_lot_per_side, }, ); diff --git a/crates/ml-backtesting/src/sim/batched_config.rs b/crates/ml-backtesting/src/sim/batched_config.rs index 73e3ae582..4c6f43407 100644 --- a/crates/ml-backtesting/src/sim/batched_config.rs +++ b/crates/ml-backtesting/src/sim/batched_config.rs @@ -17,10 +17,15 @@ pub struct BatchedSimConfig { pub latency_ns: Vec, pub kelly_frac_floor: Vec, pub sharpe_weight_floor: Vec, + // P4: threshold gate + per-fill cost. threshold = absolute conviction + // cutoff (max|p_h - 0.5| * 2 below threshold → noop). cost is deducted + // per fill in apply_fill_to_pos. + pub threshold: Vec, + pub cost_per_lot_per_side: Vec, } /// Convenience scalar input for `from_uniform`. Mirror of the historical -/// BacktestHarnessConfig sim fields. Threshold + cost will be added in P4. +/// BacktestHarnessConfig sim fields. #[derive(Clone, Debug)] pub struct UniformSimParams { pub target_annual_vol_units: f32, @@ -29,6 +34,8 @@ pub struct UniformSimParams { pub latency_ns: u32, pub kelly_frac_floor: f32, pub sharpe_weight_floor: f32, + pub threshold: f32, + pub cost_per_lot_per_side: f32, } impl BatchedSimConfig { @@ -44,17 +51,21 @@ impl BatchedSimConfig { latency_ns: vec![p.latency_ns; n], kelly_frac_floor: vec![p.kelly_frac_floor; n], sharpe_weight_floor: vec![p.sharpe_weight_floor; n], + threshold: vec![p.threshold; n], + cost_per_lot_per_side: vec![p.cost_per_lot_per_side; n], } } pub fn validate(&self) -> Result<()> { let n = self.target_annual_vol_units.len(); let lens = [ - ("annualisation_factor", self.annualisation_factor.len()), - ("max_lots", self.max_lots.len()), - ("latency_ns", self.latency_ns.len()), - ("kelly_frac_floor", self.kelly_frac_floor.len()), - ("sharpe_weight_floor", self.sharpe_weight_floor.len()), + ("annualisation_factor", self.annualisation_factor.len()), + ("max_lots", self.max_lots.len()), + ("latency_ns", self.latency_ns.len()), + ("kelly_frac_floor", self.kelly_frac_floor.len()), + ("sharpe_weight_floor", self.sharpe_weight_floor.len()), + ("threshold", self.threshold.len()), + ("cost_per_lot_per_side", self.cost_per_lot_per_side.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 914fd3480..281bb229c 100644 --- a/crates/ml-backtesting/src/sim/mod.rs +++ b/crates/ml-backtesting/src/sim/mod.rs @@ -98,6 +98,11 @@ pub struct LobSimCuda { prev_pos_lots_d: CudaSlice, // [n_backtests] prev_realized_pnl_d: CudaSlice, // [n_backtests] prev_open_horizon_mask_d: CudaSlice, // [n_backtests] + + // P4: threshold gate + per-fill cost integration. + threshold_d: CudaSlice, // [n_backtests] + cost_per_lot_per_side_d: CudaSlice, // [n_backtests] + total_fees_per_b_d: CudaSlice, // [n_backtests] — accumulator } impl LobSimCuda { @@ -243,6 +248,13 @@ impl LobSimCuda { .context("alloc prev_realized_pnl_d")?; let prev_open_horizon_mask_d = stream.alloc_zeros::(n_backtests) .context("alloc prev_open_horizon_mask_d")?; + // P4: threshold + cost + fees accumulator. + let threshold_d = stream.alloc_zeros::(n_backtests) + .context("alloc threshold_d")?; + let cost_per_lot_per_side_d = stream.alloc_zeros::(n_backtests) + .context("alloc cost_per_lot_per_side_d")?; + let total_fees_per_b_d = stream.alloc_zeros::(n_backtests) + .context("alloc total_fees_per_b_d")?; Ok(Self { n_backtests, @@ -290,6 +302,9 @@ impl LobSimCuda { prev_open_horizon_mask_d, snapshot_pos_state_fn, detect_close_fn, + threshold_d, + cost_per_lot_per_side_d, + total_fees_per_b_d, }) } @@ -608,6 +623,9 @@ impl LobSimCuda { self.stream.memcpy_htod(&sim_cfg.latency_ns, &mut self.latency_ns_d)?; self.stream.memcpy_htod(&sim_cfg.kelly_frac_floor, &mut self.kelly_frac_floor_d)?; self.stream.memcpy_htod(&sim_cfg.sharpe_weight_floor, &mut self.sharpe_weight_floor_d)?; + // P4: threshold + cost arrays. + self.stream.memcpy_htod(&sim_cfg.threshold, &mut self.threshold_d)?; + self.stream.memcpy_htod(&sim_cfg.cost_per_lot_per_side, &mut self.cost_per_lot_per_side_d)?; // Step 1: decision kernels write market_targets + open_horizon_mask. // 1a — bytecode program kernel handles backtests with plen > 0. @@ -638,6 +656,7 @@ impl LobSimCuda { .arg(&self.max_lots_d) .arg(&self.kelly_frac_floor_d) .arg(&self.sharpe_weight_floor_d) + .arg(&self.threshold_d) .arg(&n) .launch(cfg)?; } @@ -656,6 +675,7 @@ impl LobSimCuda { .arg(&self.max_lots_d) .arg(&self.kelly_frac_floor_d) .arg(&self.sharpe_weight_floor_d) + .arg(&self.threshold_d) .arg(&n) .launch(cfg)?; } @@ -1029,6 +1049,8 @@ impl LobSimCuda { .arg(&pos_bytes) .arg(&cap) .arg(&tick) + .arg(&self.cost_per_lot_per_side_d) + .arg(&mut self.total_fees_per_b_d) .arg(&n) .launch(cfg)?; } diff --git a/crates/ml-backtesting/tests/decision_floor_coldstart.rs b/crates/ml-backtesting/tests/decision_floor_coldstart.rs index 101a358b7..eef07396f 100644 --- a/crates/ml-backtesting/tests/decision_floor_coldstart.rs +++ b/crates/ml-backtesting/tests/decision_floor_coldstart.rs @@ -26,6 +26,8 @@ fn cfg_uniform(n: usize, kelly: f32, sharpe: f32) -> BatchedSimConfig { latency_ns: 0, kelly_frac_floor: kelly, sharpe_weight_floor: sharpe, + threshold: 0.0, + cost_per_lot_per_side: 0.0, }) } diff --git a/crates/ml-backtesting/tests/lob_sim_fixtures.rs b/crates/ml-backtesting/tests/lob_sim_fixtures.rs index 1f7a85db2..00a60bfb6 100644 --- a/crates/ml-backtesting/tests/lob_sim_fixtures.rs +++ b/crates/ml-backtesting/tests/lob_sim_fixtures.rs @@ -257,6 +257,8 @@ fn run_book_fixture(path: &Path) -> Result<()> { latency_ns: 0, kelly_frac_floor: 0.10, sharpe_weight_floor: 0.10, + threshold: 0.0, + cost_per_lot_per_side: 0.0, }, ); sim.step_decision(*ts_ns, &sim_cfg)?; diff --git a/crates/ml-backtesting/tests/lob_sim_integrated_fuzz.rs b/crates/ml-backtesting/tests/lob_sim_integrated_fuzz.rs index 0cbb92cb0..d3a1e94d5 100644 --- a/crates/ml-backtesting/tests/lob_sim_integrated_fuzz.rs +++ b/crates/ml-backtesting/tests/lob_sim_integrated_fuzz.rs @@ -129,6 +129,8 @@ fn run_integrated_fuzz(n_backtests: usize, n_events: usize, seed: u64) -> Result latency_ns: latency, kelly_frac_floor: 0.10, sharpe_weight_floor: 0.10, + threshold: 0.0, + cost_per_lot_per_side: 0.0, }, ); sim.step_decision_with_latency(ts_ns, &sim_cfg)?; diff --git a/crates/ml-backtesting/tests/parallel_sim_correctness.rs b/crates/ml-backtesting/tests/parallel_sim_correctness.rs index 4e5766729..727fa9cd1 100644 --- a/crates/ml-backtesting/tests/parallel_sim_correctness.rs +++ b/crates/ml-backtesting/tests/parallel_sim_correctness.rs @@ -32,6 +32,8 @@ fn parallel_sim_equivalence_with_uniform_config() -> Result<()> { latency_ns: 0, kelly_frac_floor: 0.20, sharpe_weight_floor: 0.10, + threshold: 0.0, + cost_per_lot_per_side: 0.0, }, ); sim.step_decision_with_latency(0, &cfg)?; diff --git a/crates/ml-backtesting/tests/threshold_and_cost.rs b/crates/ml-backtesting/tests/threshold_and_cost.rs new file mode 100644 index 000000000..88d41f4b6 --- /dev/null +++ b/crates/ml-backtesting/tests/threshold_and_cost.rs @@ -0,0 +1,78 @@ +//! P4 regression: threshold gate + per-fill cost integration. +//! +//! Tight tests focused on the GATE's pre-Kelly skip path and the cost +//! plumbing in apply_fill_to_pos. The full kelly_state_sees_net_return +//! end-to-end test requires a full submit_market → fill → close sequence +//! which is exercised by the production smoke; here we test the kernel +//! contract in isolation. + +use anyhow::Result; +use ml_backtesting::sim::{BatchedSimConfig, LobSimCuda, UniformSimParams}; +use ml_core::device::MlDevice; + +fn cfg_with_threshold(n: usize, threshold: f32, cost: f32) -> BatchedSimConfig { + BatchedSimConfig::from_uniform(n, &UniformSimParams { + target_annual_vol_units: 50.0, + annualisation_factor: 825.0, + max_lots: 5, + latency_ns: 0, + kelly_frac_floor: 0.20, + sharpe_weight_floor: 0.10, + threshold, + cost_per_lot_per_side: cost, + }) +} + +#[test] +#[ignore = "requires CUDA"] +fn threshold_gate_skips_low_conviction() -> Result<()> { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { eprintln!("skipping: cuda device unavailable ({e})"); return Ok(()); } + }; + let mut sim = LobSimCuda::new(1, &dev)?; + // p_h = 0.51 → max_conviction = 0.02, well below threshold = 0.10. + sim.broadcast_alpha(&[0.51, 0.51, 0.51, 0.51, 0.51])?; + sim.step_decision_with_latency(0, &cfg_with_threshold(1, 0.10, 0.0))?; + let (side, size) = sim.read_market_target(0)?; + assert_eq!(side, 2, "side should be noop under threshold gate; got side={side}"); + assert_eq!(size, 0); + Ok(()) +} + +#[test] +#[ignore = "requires CUDA"] +fn threshold_gate_allows_high_conviction() -> Result<()> { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { eprintln!("skipping: cuda device unavailable ({e})"); return Ok(()); } + }; + let mut sim = LobSimCuda::new(1, &dev)?; + // p_h = 0.8 → max_conviction = 0.6, above threshold = 0.10. + // (p=0.7 would clear the gate but ss=0.4 rounds to lots=0; needs p≥0.75 + // with kelly_floor=0.20 + max_lots=5 to clear the lots>=1 rounding.) + sim.broadcast_alpha(&[0.8, 0.8, 0.8, 0.8, 0.8])?; + sim.step_decision_with_latency(0, &cfg_with_threshold(1, 0.10, 0.0))?; + let (side, size) = sim.read_market_target(0)?; + assert_eq!(side, 0, "side should be buy with strong alpha; got side={side}"); + assert!(size >= 1, "size {size} < 1 — threshold gate may be over-restricting"); + Ok(()) +} + +#[test] +#[ignore = "requires CUDA"] +fn threshold_zero_is_passthrough() -> Result<()> { + // Sanity check: threshold = 0.0 should behave exactly like the + // P1 cold-start path (no gate). max_conviction >= 0 always. + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { eprintln!("skipping: cuda device unavailable ({e})"); return Ok(()); } + }; + let mut sim = LobSimCuda::new(1, &dev)?; + sim.broadcast_alpha(&[0.8, 0.8, 0.8, 0.8, 0.8])?; + sim.step_decision_with_latency(0, &cfg_with_threshold(1, 0.0, 0.0))?; + let (side, size) = sim.read_market_target(0)?; + assert_eq!(side, 0, "threshold=0 with strong alpha should pass through; got side={side}"); + assert!(size >= 1, "size {size} < 1 — passthrough behaviour broken"); + Ok(()) +}