diff --git a/crates/ml-backtesting/Cargo.toml b/crates/ml-backtesting/Cargo.toml index 701ac1da7..6c46beba7 100644 --- a/crates/ml-backtesting/Cargo.toml +++ b/crates/ml-backtesting/Cargo.toml @@ -16,12 +16,13 @@ description = "ML backtesting framework + LOB simulator + per-horizon execution [dependencies] ml-core.workspace = true ml-alpha = { path = "../ml-alpha" } # loader + trunk reuse — see real-LOB spec §1 +cudarc = { path = "../../vendor/cudarc" } # CUDA bindings — same vendored fork as ml-alpha serde = { workspace = true, features = ["derive"] } serde_json.workspace = true chrono.workspace = true anyhow.workspace = true csv.workspace = true -bytemuck = { workspace = true, features = ["derive"] } # SlotTag / OrderEvent / TradeRecord Pod casts +bytemuck = { workspace = true, features = ["derive"] } tracing.workspace = true thiserror.workspace = true diff --git a/crates/ml-backtesting/build.rs b/crates/ml-backtesting/build.rs new file mode 100644 index 000000000..343f076d5 --- /dev/null +++ b/crates/ml-backtesting/build.rs @@ -0,0 +1,62 @@ +//! Cubin compilation for the LOB simulator CUDA kernels. +//! Mirrors crates/ml-alpha/build.rs. +//! Per feedback_no_nvrtc.md: cubins only, no runtime nvrtc. + +use std::path::PathBuf; +use std::process::Command; + +fn main() -> Result<(), String> { + let cuda_dir = PathBuf::from(env!("CARGO_MANIFEST_DIR")).join("cuda"); + let out_dir = PathBuf::from(std::env::var("OUT_DIR").map_err(|e| e.to_string())?); + + // Per pearl_build_rs_rerun_if_env_changed.md: pair every env::var + // with rerun-if-env-changed. + // Default sm_86 covers RTX 3050 (local dev). Production Argo workflows + // override via FOXHUNT_CUDA_ARCH=sm_89 (L40S Ada) or sm_90 (H100). + let arch = std::env::var("FOXHUNT_CUDA_ARCH").unwrap_or_else(|_| "sm_86".into()); + println!("cargo:rerun-if-env-changed=FOXHUNT_CUDA_ARCH"); + println!("cargo:rerun-if-env-changed=CUDA_PATH"); + println!("cargo:rerun-if-env-changed=NVCC"); + + let nvcc = std::env::var("NVCC") + .map(PathBuf::from) + .unwrap_or_else(|_| PathBuf::from("nvcc")); + + if !cuda_dir.exists() { + // No kernels yet (e.g. before C4 lands them) — quietly noop. + return Ok(()); + } + + println!("cargo:rerun-if-changed={}", cuda_dir.display()); + let entries = std::fs::read_dir(&cuda_dir).map_err(|e| format!("read cuda dir: {e}"))?; + for entry in entries { + let entry = entry.map_err(|e| e.to_string())?; + let path = entry.path(); + if path.extension().and_then(|s| s.to_str()) != Some("cu") { + continue; + } + let stem = path + .file_stem() + .ok_or("cu file has no stem")? + .to_string_lossy() + .to_string(); + let cubin = out_dir.join(format!("{stem}.cubin")); + println!("cargo:rerun-if-changed={}", path.display()); + + let mut cmd = Command::new(&nvcc); + cmd.args(["--cubin", "-O3", "-std=c++17"]) + .arg(format!("-arch={arch}")) + .arg("-I") + .arg(&cuda_dir) + .arg("-o") + .arg(&cubin) + .arg(&path); + let status = cmd + .status() + .map_err(|e| format!("spawn nvcc for {}: {e}", path.display()))?; + if !status.success() { + return Err(format!("nvcc failed for {}", path.display())); + } + } + Ok(()) +} diff --git a/crates/ml-backtesting/cuda/book_update.cu b/crates/ml-backtesting/cuda/book_update.cu new file mode 100644 index 000000000..a598cb47d --- /dev/null +++ b/crates/ml-backtesting/cuda/book_update.cu @@ -0,0 +1,32 @@ +// book_update.cu — MBP-10 snapshot apply. +// +// Replaces the per-block book state with the broadcast snapshot. The +// full top-10 image overwrites the previous state (MBP-10 semantics). +// One block per backtest; thread tid handles one of 10 levels per side +// (10 threads used out of 32 in the warp; remainder idle). +// +// Per feedback_no_atomicadd.md: no atomics — single-writer per level. +// Per spec §2: same snapshot broadcast to every backtest in v1. + +#include "lob_state.cuh" + +// Books are laid out [n_backtests, 4 fields × 10 levels]. +// Field order: bid_px, bid_sz, ask_px, ask_sz (one contiguous Book per backtest). +extern "C" __global__ void book_update_apply_snapshot( + const float* __restrict__ bid_px_in, // [10] — broadcast + const float* __restrict__ bid_sz_in, // [10] + const float* __restrict__ ask_px_in, // [10] + const float* __restrict__ ask_sz_in, // [10] + Book* __restrict__ books_out, // [n_backtests] + int n_backtests +) { + int b = blockIdx.x; + int tid = threadIdx.x; + if (b >= n_backtests || tid >= 10) return; + + Book& bk = books_out[b]; + bk.bid_px[tid] = bid_px_in[tid]; + bk.bid_sz[tid] = bid_sz_in[tid]; + bk.ask_px[tid] = ask_px_in[tid]; + bk.ask_sz[tid] = ask_sz_in[tid]; +} diff --git a/crates/ml-backtesting/cuda/lob_state.cuh b/crates/ml-backtesting/cuda/lob_state.cuh new file mode 100644 index 000000000..994dbc390 --- /dev/null +++ b/crates/ml-backtesting/cuda/lob_state.cuh @@ -0,0 +1,20 @@ +// Per-block shared-mem layout for the LOB simulator. +// One block = one parallel backtest. See spec §2 + §5b. +// +// MUST stay in sync with crates/ml-backtesting/src/lob/mod.rs's host mirrors. + +#ifndef LOB_STATE_CUH +#define LOB_STATE_CUH + +#define N_HORIZONS 5 +#define MAX_LIMITS 32 +#define MAX_STOPS 16 + +struct Book { + float bid_px[10]; + float bid_sz[10]; + float ask_px[10]; + float ask_sz[10]; +}; + +#endif // LOB_STATE_CUH diff --git a/crates/ml-backtesting/src/lib.rs b/crates/ml-backtesting/src/lib.rs index 9bec3ff3e..f99366986 100644 --- a/crates/ml-backtesting/src/lib.rs +++ b/crates/ml-backtesting/src/lib.rs @@ -17,9 +17,11 @@ pub use ml_core::MLError; pub mod action_loader; pub mod barrier_backtest; +pub mod lob; pub mod order; pub mod policy; pub mod report; +pub mod sim; pub use action_loader::{load_actions_from_csv, DQNActionRecord}; pub use barrier_backtest::{BacktestResults, BarrierBacktester, BarrierParams}; diff --git a/crates/ml-backtesting/src/lob/mod.rs b/crates/ml-backtesting/src/lob/mod.rs new file mode 100644 index 000000000..48d1b998b --- /dev/null +++ b/crates/ml-backtesting/src/lob/mod.rs @@ -0,0 +1,19 @@ +//! Host-side mirrors of cuda/lob_state.cuh layouts. +//! See spec §2 (CUDA data layout) and §5b (state ownership). + +pub const BOOK_LEVELS: usize = 10; +pub const N_HORIZONS: usize = 5; +pub const MAX_LIMITS: usize = 32; +pub const MAX_STOPS: usize = 16; + +/// Flat host-side mirror of `cuda/lob_state.cuh::Book`. +/// Fields stored field-major (bid_px[10] then bid_sz[10] then ask_px[10] +/// then ask_sz[10]) to match the cuda struct layout when interpreted as +/// 40 contiguous floats. See `LobSimCuda::read_books`. +#[derive(Clone, Copy, Debug, Default, PartialEq)] +pub struct BookFlat { + pub bid_px: [f32; BOOK_LEVELS], + pub bid_sz: [f32; BOOK_LEVELS], + pub ask_px: [f32; BOOK_LEVELS], + pub ask_sz: [f32; BOOK_LEVELS], +} diff --git a/crates/ml-backtesting/src/sim.rs b/crates/ml-backtesting/src/sim.rs new file mode 100644 index 000000000..179ba817a --- /dev/null +++ b/crates/ml-backtesting/src/sim.rs @@ -0,0 +1,130 @@ +//! Block-per-backtest LOB simulator. +//! +//! One CUDA block = one parallel backtest. The simulator owns per-backtest +//! book state (and, in later commits, orders / pos / ISV-Kelly state + +//! mapped-pinned audit & trade-log rings). +//! +//! See docs/superpowers/specs/2026-05-18-real-lob-integration-design.md §2 + §5b. + +use anyhow::{Context, Result}; +use cudarc::driver::{CudaSlice, CudaStream, LaunchConfig, PushKernelArg}; +use ml_core::device::MlDevice; +use std::sync::Arc; + +use crate::lob::{BookFlat, BOOK_LEVELS}; + +const BOOK_FIELDS: usize = 4; // bid_px, bid_sz, ask_px, ask_sz + +/// Precompiled book_update cubin embedded at compile time. +const BOOK_UPDATE_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/book_update.cubin")); + +pub struct LobSimCuda { + n_backtests: usize, + stream: Arc, + book_update_fn: cudarc::driver::CudaFunction, + + // Per-backtest book state on device. Laid out + // [n_backtests, 4 fields × 10 levels] = N × 40 floats. + books_d: CudaSlice, + + // Mapped-pinned input snapshot buffers (size 10 each) — broadcast across + // all N backtests in v1 (single-instrument, single-data-stream). + bid_px_d: CudaSlice, + bid_sz_d: CudaSlice, + ask_px_d: CudaSlice, + ask_sz_d: CudaSlice, +} + +impl LobSimCuda { + pub fn new(n_backtests: usize, dev: &MlDevice) -> Result { + anyhow::ensure!(n_backtests > 0, "n_backtests must be > 0"); + let stream = dev.cuda_stream().context("cuda_stream")?.clone(); + let ctx = dev.cuda_context().context("cuda_context")?; + + let module = ctx + .load_cubin(BOOK_UPDATE_CUBIN.to_vec()) + .context("load book_update cubin")?; + let book_update_fn = module + .load_function("book_update_apply_snapshot") + .context("load book_update_apply_snapshot")?; + + let books_d = stream + .alloc_zeros::(n_backtests * BOOK_FIELDS * BOOK_LEVELS) + .context("alloc books_d")?; + let bid_px_d = stream.alloc_zeros::(BOOK_LEVELS).context("bid_px alloc")?; + let bid_sz_d = stream.alloc_zeros::(BOOK_LEVELS).context("bid_sz alloc")?; + let ask_px_d = stream.alloc_zeros::(BOOK_LEVELS).context("ask_px alloc")?; + let ask_sz_d = stream.alloc_zeros::(BOOK_LEVELS).context("ask_sz alloc")?; + + Ok(Self { + n_backtests, + stream, + book_update_fn, + books_d, + bid_px_d, + bid_sz_d, + ask_px_d, + ask_sz_d, + }) + } + + pub fn n_backtests(&self) -> usize { + self.n_backtests + } + + /// Apply one MBP-10 snapshot's top-10 levels to every backtest's book. + /// v1 inputs are single-snapshot broadcast (same market data to all + /// parallel backtests — see spec §7 inference scope, policy sweeps only). + pub fn apply_snapshot( + &mut self, + bid_px: &[f32; BOOK_LEVELS], + bid_sz: &[f32; BOOK_LEVELS], + ask_px: &[f32; BOOK_LEVELS], + ask_sz: &[f32; BOOK_LEVELS], + ) -> Result<()> { + self.stream.memcpy_htod(bid_px.as_slice(), &mut self.bid_px_d)?; + self.stream.memcpy_htod(bid_sz.as_slice(), &mut self.bid_sz_d)?; + self.stream.memcpy_htod(ask_px.as_slice(), &mut self.ask_px_d)?; + self.stream.memcpy_htod(ask_sz.as_slice(), &mut self.ask_sz_d)?; + + let cfg = 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 mut launch = self.stream.launch_builder(&self.book_update_fn); + unsafe { + launch + .arg(&self.bid_px_d) + .arg(&self.bid_sz_d) + .arg(&self.ask_px_d) + .arg(&self.ask_sz_d) + .arg(&mut self.books_d) + .arg(&n) + .launch(cfg)?; + } + self.stream.synchronize()?; + Ok(()) + } + + /// Read back every backtest's book state from device. For tests + diagnostics. + pub fn read_books(&self) -> Result> { + let mut raw = vec![0.0f32; self.n_backtests * BOOK_FIELDS * BOOK_LEVELS]; + self.stream.memcpy_dtoh(&self.books_d, raw.as_mut_slice())?; + let mut out = Vec::with_capacity(self.n_backtests); + for b in 0..self.n_backtests { + let off = b * BOOK_FIELDS * BOOK_LEVELS; + let mut bf = BookFlat::default(); + for k in 0..BOOK_LEVELS { + bf.bid_px[k] = raw[off + k]; + bf.bid_sz[k] = raw[off + BOOK_LEVELS + k]; + bf.ask_px[k] = raw[off + 2 * BOOK_LEVELS + k]; + bf.ask_sz[k] = raw[off + 3 * BOOK_LEVELS + k]; + } + out.push(bf); + } + Ok(out) + } +} diff --git a/crates/ml-backtesting/tests/fixtures/book_update_n_backtests.json b/crates/ml-backtesting/tests/fixtures/book_update_n_backtests.json new file mode 100644 index 000000000..dedaa3e35 --- /dev/null +++ b/crates/ml-backtesting/tests/fixtures/book_update_n_backtests.json @@ -0,0 +1,22 @@ +{ + "name": "book_update_n_backtests", + "n_backtests": 4, + "events": [ + { + "type": "snapshot", + "bid_px": [5499.75, 5499.50, 5499.25, 5499.00, 5498.75, 5498.50, 5498.25, 5498.00, 5497.75, 5497.50], + "bid_sz": [12.0, 18.0, 24.0, 30.0, 36.0, 42.0, 48.0, 54.0, 60.0, 66.0], + "ask_px": [5500.00, 5500.25, 5500.50, 5500.75, 5501.00, 5501.25, 5501.50, 5501.75, 5502.00, 5502.25], + "ask_sz": [10.0, 16.0, 22.0, 28.0, 34.0, 40.0, 46.0, 52.0, 58.0, 64.0] + } + ], + "expected_books": [ + { + "bid_px": [5499.75, 5499.50, 5499.25, 5499.00, 5498.75, 5498.50, 5498.25, 5498.00, 5497.75, 5497.50], + "bid_sz": [12.0, 18.0, 24.0, 30.0, 36.0, 42.0, 48.0, 54.0, 60.0, 66.0], + "ask_px": [5500.00, 5500.25, 5500.50, 5500.75, 5501.00, 5501.25, 5501.50, 5501.75, 5502.00, 5502.25], + "ask_sz": [10.0, 16.0, 22.0, 28.0, 34.0, 40.0, 46.0, 52.0, 58.0, 64.0] + } + ], + "expected_books_apply_to_all": true +} diff --git a/crates/ml-backtesting/tests/fixtures/book_update_replace.json b/crates/ml-backtesting/tests/fixtures/book_update_replace.json new file mode 100644 index 000000000..3fdbec0f6 --- /dev/null +++ b/crates/ml-backtesting/tests/fixtures/book_update_replace.json @@ -0,0 +1,21 @@ +{ + "name": "book_update_replace", + "n_backtests": 1, + "events": [ + { + "type": "snapshot", + "bid_px": [5499.75, 5499.50, 5499.25, 5499.00, 5498.75, 5498.50, 5498.25, 5498.00, 5497.75, 5497.50], + "bid_sz": [12.0, 18.0, 24.0, 30.0, 36.0, 42.0, 48.0, 54.0, 60.0, 66.0], + "ask_px": [5500.00, 5500.25, 5500.50, 5500.75, 5501.00, 5501.25, 5501.50, 5501.75, 5502.00, 5502.25], + "ask_sz": [10.0, 16.0, 22.0, 28.0, 34.0, 40.0, 46.0, 52.0, 58.0, 64.0] + } + ], + "expected_books": [ + { + "bid_px": [5499.75, 5499.50, 5499.25, 5499.00, 5498.75, 5498.50, 5498.25, 5498.00, 5497.75, 5497.50], + "bid_sz": [12.0, 18.0, 24.0, 30.0, 36.0, 42.0, 48.0, 54.0, 60.0, 66.0], + "ask_px": [5500.00, 5500.25, 5500.50, 5500.75, 5501.00, 5501.25, 5501.50, 5501.75, 5502.00, 5502.25], + "ask_sz": [10.0, 16.0, 22.0, 28.0, 34.0, 40.0, 46.0, 52.0, 58.0, 64.0] + } + ] +} diff --git a/crates/ml-backtesting/tests/fixtures/book_update_two_step.json b/crates/ml-backtesting/tests/fixtures/book_update_two_step.json new file mode 100644 index 000000000..5af53533d --- /dev/null +++ b/crates/ml-backtesting/tests/fixtures/book_update_two_step.json @@ -0,0 +1,28 @@ +{ + "name": "book_update_two_step", + "n_backtests": 1, + "events": [ + { + "type": "snapshot", + "bid_px": [5499.75, 5499.50, 5499.25, 5499.00, 5498.75, 5498.50, 5498.25, 5498.00, 5497.75, 5497.50], + "bid_sz": [12.0, 18.0, 24.0, 30.0, 36.0, 42.0, 48.0, 54.0, 60.0, 66.0], + "ask_px": [5500.00, 5500.25, 5500.50, 5500.75, 5501.00, 5501.25, 5501.50, 5501.75, 5502.00, 5502.25], + "ask_sz": [10.0, 16.0, 22.0, 28.0, 34.0, 40.0, 46.0, 52.0, 58.0, 64.0] + }, + { + "type": "snapshot", + "bid_px": [5500.00, 5499.75, 5499.50, 5499.25, 5499.00, 5498.75, 5498.50, 5498.25, 5498.00, 5497.75], + "bid_sz": [5.0, 15.0, 25.0, 35.0, 45.0, 55.0, 65.0, 75.0, 85.0, 95.0], + "ask_px": [5500.25, 5500.50, 5500.75, 5501.00, 5501.25, 5501.50, 5501.75, 5502.00, 5502.25, 5502.50], + "ask_sz": [8.0, 20.0, 32.0, 44.0, 56.0, 68.0, 80.0, 92.0, 104.0, 116.0] + } + ], + "expected_books": [ + { + "bid_px": [5500.00, 5499.75, 5499.50, 5499.25, 5499.00, 5498.75, 5498.50, 5498.25, 5498.00, 5497.75], + "bid_sz": [5.0, 15.0, 25.0, 35.0, 45.0, 55.0, 65.0, 75.0, 85.0, 95.0], + "ask_px": [5500.25, 5500.50, 5500.75, 5501.00, 5501.25, 5501.50, 5501.75, 5502.00, 5502.25, 5502.50], + "ask_sz": [8.0, 20.0, 32.0, 44.0, 56.0, 68.0, 80.0, 92.0, 104.0, 116.0] + } + ] +} diff --git a/crates/ml-backtesting/tests/lob_sim_fixtures.rs b/crates/ml-backtesting/tests/lob_sim_fixtures.rs new file mode 100644 index 000000000..72db6881c --- /dev/null +++ b/crates/ml-backtesting/tests/lob_sim_fixtures.rs @@ -0,0 +1,100 @@ +//! Ring 1 fixture tests: N=1 (or small N) bit-exact GPU oracle assertions. +//! See docs/superpowers/specs/2026-05-18-real-lob-integration-design.md §8. +//! +//! Run with: `cargo test -p ml-backtesting --test lob_sim_fixtures -- --ignored --nocapture` + +use anyhow::{Context, Result}; +use ml_backtesting::lob::BOOK_LEVELS; +use ml_backtesting::sim::LobSimCuda; +use ml_core::device::MlDevice; +use serde::Deserialize; +use std::path::Path; + +#[derive(Debug, Deserialize)] +#[serde(tag = "type", rename_all = "snake_case")] +enum FixtureEvent { + Snapshot { + bid_px: [f32; BOOK_LEVELS], + bid_sz: [f32; BOOK_LEVELS], + ask_px: [f32; BOOK_LEVELS], + ask_sz: [f32; BOOK_LEVELS], + }, +} + +#[derive(Debug, Deserialize)] +struct FixtureBook { + bid_px: [f32; BOOK_LEVELS], + bid_sz: [f32; BOOK_LEVELS], + ask_px: [f32; BOOK_LEVELS], + ask_sz: [f32; BOOK_LEVELS], +} + +#[derive(Debug, Deserialize)] +struct Fixture { + name: String, + n_backtests: usize, + events: Vec, + expected_books: Vec, + #[serde(default)] + expected_books_apply_to_all: bool, +} + +fn run_book_fixture(path: &Path) -> Result<()> { + let s = std::fs::read_to_string(path) + .with_context(|| format!("read {}", path.display()))?; + let fx: Fixture = serde_json::from_str(&s).context("parse fixture")?; + eprintln!("running fixture: {}", fx.name); + + let dev = MlDevice::cuda(0) + .map_err(|e| anyhow::anyhow!("cuda device: {e}"))?; + let mut sim = LobSimCuda::new(fx.n_backtests, &dev)?; + + for ev in &fx.events { + match ev { + FixtureEvent::Snapshot { bid_px, bid_sz, ask_px, ask_sz } => { + sim.apply_snapshot(bid_px, bid_sz, ask_px, ask_sz)?; + } + } + } + + let books = sim.read_books()?; + let expected: Vec<&FixtureBook> = if fx.expected_books_apply_to_all { + (0..fx.n_backtests).map(|_| &fx.expected_books[0]).collect() + } else { + anyhow::ensure!( + fx.expected_books.len() == fx.n_backtests, + "expected_books len {} ≠ n_backtests {}", + fx.expected_books.len(), + fx.n_backtests + ); + fx.expected_books.iter().collect() + }; + + for (b, (got, want)) in books.iter().zip(expected.iter()).enumerate() { + for k in 0..BOOK_LEVELS { + assert_eq!(got.bid_px[k], want.bid_px[k], "backtest {b} bid_px[{k}]"); + assert_eq!(got.bid_sz[k], want.bid_sz[k], "backtest {b} bid_sz[{k}]"); + assert_eq!(got.ask_px[k], want.ask_px[k], "backtest {b} ask_px[{k}]"); + assert_eq!(got.ask_sz[k], want.ask_sz[k], "backtest {b} ask_sz[{k}]"); + } + } + Ok(()) +} + +#[test] +#[ignore = "requires CUDA"] +fn fix_book_update_replace() { + run_book_fixture(Path::new("tests/fixtures/book_update_replace.json")).unwrap(); +} + +#[test] +#[ignore = "requires CUDA"] +fn fix_book_update_two_step() { + run_book_fixture(Path::new("tests/fixtures/book_update_two_step.json")).unwrap(); +} + +#[test] +#[ignore = "requires CUDA"] +fn fix_book_update_n_backtests() { + run_book_fixture(Path::new("tests/fixtures/book_update_n_backtests.json")).unwrap(); +}