gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

feat(ml-backtesting): seed_inflight_limits_batched kernel (P2)

Browse Source 
Replaces the host roundtrip loop at the latency path of step_decision_
with_latency with one GPU kernel launch. At n_parallel=140 the host
loop did up to 140 memcpy_dtoh + 140 seed_limit_order calls per
decision (~210M roundtrips per quarter at the threshold-tuning load).
The new kernel does the same work in one launch.

Per-backtest single-writer (threadIdx.x==0). Each backtest scans its
own MAX_LIMITS=32 slot range for an `active==0` slot. Slot allocation
is per-backtest (no cross-backtest atomics needed). Overflow path
increments pos.submission_overflow.

dispatch_latent_market_orders now takes &BatchedSimConfig (unused
inside — the latency_ns_d device buffer is already populated by
step_decision_with_latency's upload block from P1).

All 3 decision_floor_coldstart tests still pass via the new GPU path
(at latency_ns=0 the in-flight slot's arrival_ts==current_ts and gets
promoted on the next step_resting_orders, functionally equivalent to
the legacy submit_market_immediate kernel).

Independence test (different per-backtest latencies → different
arrival_ts) deferred to a later step where a read_first_inflight_arrival_ts
helper is added.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-19 17:01:25 +02:00
parent b8966fb1a6
commit b741f0c5ce
2 changed files with 77 additions and 44 deletions
Download Patch File Download Diff File Expand all files Collapse all files

47
crates/ml-backtesting/cuda/resting_orders.cu
View File

@@ -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*>(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;
}

74
crates/ml-backtesting/src/sim/mod.rs
View File

@@ -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<f32>,
@@ -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(&current_ts_ns)
.arg(&n)
.launch(cfg_launch)?;
}
self.stream.synchronize()?;
Ok(())
}