diff --git a/crates/ml-backtesting/cuda/resting_orders.cu b/crates/ml-backtesting/cuda/resting_orders.cu index 97560ad76..9bfe2b1ed 100644 --- a/crates/ml-backtesting/cuda/resting_orders.cu +++ b/crates/ml-backtesting/cuda/resting_orders.cu @@ -429,13 +429,21 @@ extern "C" __global__ void submit_limit_alloc( // P2: GPU-side seeder for in-flight LimitSlots driven by market_targets[b]. // Replaces the host-roundtrip loop in sim.rs::dispatch_latent_market_orders. -// Per-backtest single-writer (threadIdx.x == 0). Scans this backtest's own -// MAX_LIMITS slot range for an `active == 0` slot; if found, writes the -// in-flight Limit with arrival_ts = current_ts + latency_ns_per_b[b]. If -// all 32 slots are in use, increments pos.submission_overflow. +// Per-backtest single-writer (threadIdx.x == 0). // -// noop targets (side == 2) and zero-size targets are skipped — the prior -// host loop had the same behaviour. +// §7 truth table: +// side=0, abs_sz=N → target signed position = +N (long) +// side=1, abs_sz=N → target signed position = -N (short) +// side=2 → no-op: preserve current position +// side=3 → force flat: target signed position = 0 +// +// effective_position = pos.position_lots + Σ unfilled signed slot sizes +// where unfilled = active∈{1,2} (resting or in-flight). +// delta = target_signed − effective_position. +// If delta == 0: nothing to do (idempotent re-submission under latency). +// If delta != 0: seed an in-flight order for |delta| lots on the +// appropriate side. If all 32 slots are occupied, increments +// pos.submission_overflow. extern "C" __global__ void seed_inflight_limits_batched( unsigned char* __restrict__ orders_base, // [n_backtests * orders_bytes] int orders_bytes, @@ -448,28 +456,49 @@ extern "C" __global__ void seed_inflight_limits_batched( int b = blockIdx.x; if (b >= n_backtests || threadIdx.x != 0) return; - const int side = market_targets[b * 2 + 0]; - const int size = market_targets[b * 2 + 1]; - if (side == 2 || size <= 0) return; // noop / no order to seed + const int side = market_targets[b * 2 + 0]; + const int abs_sz = market_targets[b * 2 + 1]; + + if (side == 2) return; // §7: no-op = preserve position + + int target_signed; + if (side == 0) target_signed = +abs_sz; + else if (side == 1) target_signed = -abs_sz; + else target_signed = 0; // side == 3: force flat Orders* orders = reinterpret_cast(orders_base + (size_t)b * orders_bytes); - // Scan per-backtest slots for a free one. + // Effective position = filled position + all unfilled signed lots. + int effective = positions[b].position_lots; + #pragma unroll + for (int s_idx = 0; s_idx < MAX_LIMITS; ++s_idx) { + const LimitSlot& s = orders->limits[s_idx]; + if (s.active == 0) continue; + const int slot_lots = (int)lroundf(s.size_remaining); + effective += (s.side == 0) ? slot_lots : -slot_lots; + } + + const int delta = target_signed - effective; + if (delta == 0) return; + + const int order_lots = (delta > 0) ? delta : -delta; + const int order_side = (delta > 0) ? 0 : 1; + + // Find a free slot and seed an in-flight order. for (int s_idx = 0; s_idx < MAX_LIMITS; ++s_idx) { LimitSlot& s = orders->limits[s_idx]; if (s.active == 0) { s.active = 2; // in-flight - s.side = (unsigned char)side; + s.side = (unsigned char)order_side; s.kind = 1; // Limit s.oco_pair = 0xFF; - s.price = (side == 0) ? 1.0e9f : 0.0f; // always crosses - s.size_remaining = (float)size; + s.price = (order_side == 0) ? 1.0e9f : 0.0f; // always crosses + s.size_remaining = (float)order_lots; s.queue_position = 0.0f; s.arrival_ts_ns = current_ts_ns + (unsigned long long)latency_ns_per_b[b]; s.gen_counter = (unsigned short)(s.gen_counter + 1); return; } } - // No free slot — record overflow. positions[b].submission_overflow += 1u; } diff --git a/crates/ml-backtesting/tests/stop_controller.rs b/crates/ml-backtesting/tests/stop_controller.rs index 592edadb7..dda8bc312 100644 --- a/crates/ml-backtesting/tests/stop_controller.rs +++ b/crates/ml-backtesting/tests/stop_controller.rs @@ -425,3 +425,125 @@ fn bytecode_and_default_kernel_agree_on_stops() -> Result<()> { assert_eq!(mt_default.1, 0, "force-flat encoding has abs_sz=0; got size={}", mt_default.1); Ok(()) } + +#[test] +#[ignore = "requires CUDA"] +fn position_target_not_additive() -> Result<()> { + // latency=0: target (0, n) repeated must NOT accumulate the position + // past n. Cap is max_lots=5. + 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)?; + let (bp, bs, ap, az) = level_book(5500.0, 0.25); + sim.apply_snapshot(&bp, &bs, &ap, &az)?; + + let cfg = cfg_default(1); + let mut ts: u64 = 0; + for _ in 0..10 { + sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + sim.step_decision_with_latency(ts, &cfg)?; + sim.step_resting_orders(ts, 0.0)?; + ts += 1_000_000; + } + let pos = sim.read_pos(0)?; + assert!(pos.position_lots.abs() <= 5, + "|position_lots| must respect max_lots=5 with target semantics; got {}", + pos.position_lots); + Ok(()) +} + +#[test] +#[ignore = "requires CUDA"] +fn position_target_not_additive_with_latency() -> Result<()> { + // 200ms latency: in-flight orders must count toward effective_position + // so re-submission of the same target produces delta=0 and no new order. + 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)?; + let (bp, bs, ap, az) = level_book(5500.0, 0.25); + sim.apply_snapshot(&bp, &bs, &ap, &az)?; + + let cfg_lat = BatchedSimConfig::from_uniform(1, &UniformSimParams { + target_annual_vol_units: 50.0, + annualisation_factor: 825.0, + max_lots: 5, + latency_ns: 200_000_000, + kelly_frac_floor: 0.20, + sharpe_weight_floor: 0.10, + threshold: 0.0, + cost_per_lot_per_side: 0.0, + use_cold_start_stopgap: false, + }); + + let mut ts: u64 = 0; + // 500 events of persistent target — most should produce delta=0. + for _ in 0..500 { + sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + sim.step_decision_with_latency(ts, &cfg_lat)?; + sim.step_resting_orders(ts, 0.0)?; + ts += 1_000_000; + } + let pos = sim.read_pos(0)?; + assert!(pos.position_lots.abs() <= 5, + "|position_lots| must stay <= max_lots=5 under 200ms latency with in-flight summation; got {}", + pos.position_lots); + Ok(()) +} + +#[test] +#[ignore = "requires CUDA"] +fn alpha_noop_side_2_preserved() -> 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)?; + + // Seed ISV with large pnl_ema_loss (= 10.0) so sl_distance stays far enough + // that the spread (0.25) and small price moves never fire the SL during the + // open-position setup phase. Same pattern as other open-setup tests. + let cold_start: [IsvKellyStateHost; 5] = std::array::from_fn(|_| IsvKellyStateHost { + pnl_ema_win: 0.0, + pnl_ema_loss: 10.0, + win_rate_ema: 0.0, + n_trades_seen: 0, + realised_return_var: 0.0, + recent_sharpe: 0.0, + }); + sim.write_isv_kelly(0, &cold_start)?; + + let (bp, bs, ap, az) = level_book(5500.0, 0.25); + sim.apply_snapshot(&bp, &bs, &ap, &az)?; + let (bp2, bs2, ap2, az2) = level_book(5500.1, 0.25); + sim.apply_snapshot(&bp2, &bs2, &ap2, &az2)?; + let mut ts: u64 = 0; + + // Open long via strong alpha. + for _ in 0..10 { + sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + sim.step_decision_with_latency(ts, &cfg_default(1))?; + sim.step_resting_orders(ts, 0.0)?; + ts += 1_000_000; + } + let pos_open = sim.read_pos(0)?.position_lots; + assert!(pos_open > 0, "setup: long opens with strong alpha, got {}", pos_open); + + // Weak alpha that produces sig_mag * 0.20 * cap_lots < 0.5, rounding lots=0 + // → market_target = (2, 0). The pos must stay unchanged. + // Compute: sig_mag = |0.55-0.5|*2 = 0.1; ss = 0.1*0.20*5 = 0.1; truncf(0.1+0.5)=0. + for _ in 0..20 { + sim.broadcast_alpha(&[0.55, 0.55, 0.55, 0.55, 0.55])?; + sim.step_decision_with_latency(ts, &cfg_default(1))?; + sim.step_resting_orders(ts, 0.0)?; + ts += 1_000_000; + } + let pos_after = sim.read_pos(0)?.position_lots; + assert_eq!(pos_after, pos_open, + "side=2 must preserve position; opened={} after 20 weak-alpha events={}", + pos_open, pos_after); + Ok(()) +}