diff --git a/crates/ml-backtesting/cuda/resting_orders.cu b/crates/ml-backtesting/cuda/resting_orders.cu index f464779f9..74ae70a82 100644 --- a/crates/ml-backtesting/cuda/resting_orders.cu +++ b/crates/ml-backtesting/cuda/resting_orders.cu @@ -397,3 +397,50 @@ extern "C" __global__ void submit_limit_alloc( s.gen_counter = (unsigned short)(s.gen_counter + 1); *allocated_slot_out = idx; } + +// 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. +// +// noop targets (side == 2) and zero-size targets are skipped — the prior +// host loop had the same behaviour. +extern "C" __global__ void seed_inflight_limits_batched( + unsigned char* __restrict__ orders_base, // [n_backtests * orders_bytes] + int orders_bytes, + Pos* __restrict__ positions, // [n_backtests] + const int* __restrict__ market_targets, // [n_backtests * 2] + const unsigned int* __restrict__ latency_ns_per_b, // [n_backtests] + unsigned long long current_ts_ns, + int n_backtests +) { + 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 + + Orders* orders = reinterpret_cast(orders_base + (size_t)b * orders_bytes); + + // Scan per-backtest slots for a free one. + 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.kind = 1; // Limit + s.oco_pair = 0xFF; + s.price = (side == 0) ? 1.0e9f : 0.0f; // always crosses + s.size_remaining = (float)size; + 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/src/sim/mod.rs b/crates/ml-backtesting/src/sim/mod.rs index a459288a3..e057d8bb2 100644 --- a/crates/ml-backtesting/src/sim/mod.rs +++ b/crates/ml-backtesting/src/sim/mod.rs @@ -44,6 +44,9 @@ pub struct LobSimCuda { isv_kelly_update_fn: cudarc::driver::CudaFunction, resting_orders_fn: cudarc::driver::CudaFunction, seed_limit_fn: cudarc::driver::CudaFunction, + /// P2: GPU-side seed of in-flight Limit slots from market_targets[b]. + /// Replaces the host roundtrip loop in dispatch_latent_market_orders. + seed_inflight_limits_fn: cudarc::driver::CudaFunction, seed_stop_fn: cudarc::driver::CudaFunction, books_d: CudaSlice, @@ -133,6 +136,9 @@ impl LobSimCuda { let seed_limit_fn = resting_module .load_function("seed_limit_slot") .context("load seed_limit_slot")?; + let seed_inflight_limits_fn = resting_module + .load_function("seed_inflight_limits_batched") + .context("load seed_inflight_limits_batched")?; let seed_stop_fn = resting_module .load_function("seed_stop_slot") .context("load seed_stop_slot")?; @@ -226,6 +232,7 @@ impl LobSimCuda { isv_kelly_update_fn, resting_orders_fn, seed_limit_fn, + seed_inflight_limits_fn, seed_stop_fn, books_d, pos_d, @@ -718,52 +725,31 @@ impl LobSimCuda { fn dispatch_latent_market_orders( &mut self, current_ts_ns: u64, - sim_cfg: &BatchedSimConfig, + _sim_cfg: &BatchedSimConfig, ) -> Result<()> { - let n = self.n_backtests; - let mut targets = vec![0i32; n * 2]; - self.stream.memcpy_dtoh(&self.market_targets_d, targets.as_mut_slice())?; - for b in 0..n { - let side = targets[b * 2]; - let size = targets[b * 2 + 1]; - if side == 2 || size <= 0 { - continue; // no-op - } - let latency_ns = sim_cfg.latency_ns[b]; - // Find a free slot via the existing seed path (which sets - // overflow_flag if all slots are in use, surfaced as Err). - let mut placed = false; - for slot in 0..crate::lob::MAX_LIMITS { - if self.seed_limit_order( - b, - slot, - side as u8, - 1, // Limit - 2, // in-flight - 0xFF, - if side == 0 { 1.0e9 } else { 0.0 }, - size as f32, - 0.0, - current_ts_ns + latency_ns as u64, - ).is_ok() { - placed = true; - break; - } - } - if !placed { - // All 32 slots in use; record overflow on the Pos counter. - // (seed_limit_order errored on each attempt; here we'd want - // pos.submission_overflow++ but doing it via host is wasteful - // when the device-side path already increments naturally for - // in-kernel allocations. v1 simply drops the order silently - // and the test fixture for submission_overflow already - // exercises the device-side counter.) - tracing::warn!( - backtest_idx = b, - "dispatch_latent_market_orders: all limit slots in use; dropping" - ); - } + // P2: GPU-side seed of in-flight LimitSlots. Replaces the host + // roundtrip loop. Per-backtest latency reads from self.latency_ns_d + // (already uploaded by step_decision_with_latency from sim_cfg). + let cfg_launch = LaunchConfig { + grid_dim: (self.n_backtests as u32, 1, 1), + block_dim: (32, 1, 1), + shared_mem_bytes: 0, + }; + let n = self.n_backtests as i32; + let orders_bytes = crate::lob::ORDERS_BYTES as i32; + let mut launch = self.stream.launch_builder(&self.seed_inflight_limits_fn); + unsafe { + launch + .arg(&mut self.orders_d) + .arg(&orders_bytes) + .arg(&mut self.pos_d) + .arg(&self.market_targets_d) + .arg(&self.latency_ns_d) + .arg(¤t_ts_ns) + .arg(&n) + .launch(cfg_launch)?; } + self.stream.synchronize()?; Ok(()) }