feat(ml-backtesting): build.rs + book_update CUDA kernel + sim skeleton

build.rs compiles every .cu under crates/ml-backtesting/cuda/ to
\$OUT_DIR/<name>.cubin via nvcc (no nvrtc per feedback_no_nvrtc.md);
pairs every env::var with rerun-if-env-changed per
pearl_build_rs_rerun_if_env_changed.md. Default arch sm_86 (RTX 3050
local); production sets FOXHUNT_CUDA_ARCH=sm_89 (L40S) or sm_90 (H100).

First kernel: book_update_apply_snapshot — block-per-backtest replace
of the 10-level book from a broadcast MBP-10 snapshot. Single-writer
per level (thread tid = level idx), no atomics per feedback_no_atomicadd.md.

lob_state.cuh defines the canonical per-block shared-mem layout that
all subsequent kernels share (Book struct in this commit; Orders, Pos,
IsvKellyState[5] added in C5-C7).

LobSimCuda host struct (in src/sim.rs) constructed from an &MlDevice;
owns per-backtest device book state + mapped-pinned input snapshots.
apply_snapshot launches the kernel and synchronises; read_books drains
device state for tests.

Three Ring 1 fixture tests (book_update_replace, _two_step,
_n_backtests=4) — N≤4 bit-exact GPU oracle assertions loaded from JSON.
All three pass on RTX 3050 locally. Verifies the build-script → cubin →
cudarc.load_cubin → kernel launch → DtoH read pipeline end-to-end.

cudarc added as direct path dep (../../vendor/cudarc) since it's not in
[workspace.dependencies] — same vendored fork ml-alpha uses.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>

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
 

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(())
+}

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];
+}

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

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};

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],
+}

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<CudaStream>,
+    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<f32>,
+
+    // Mapped-pinned input snapshot buffers (size 10 each) — broadcast across
+    // all N backtests in v1 (single-instrument, single-data-stream).
+    bid_px_d: CudaSlice<f32>,
+    bid_sz_d: CudaSlice<f32>,
+    ask_px_d: CudaSlice<f32>,
+    ask_sz_d: CudaSlice<f32>,
+}
+
+impl LobSimCuda {
+    pub fn new(n_backtests: usize, dev: &MlDevice) -> Result<Self> {
+        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::<f32>(n_backtests * BOOK_FIELDS * BOOK_LEVELS)
+            .context("alloc books_d")?;
+        let bid_px_d = stream.alloc_zeros::<f32>(BOOK_LEVELS).context("bid_px alloc")?;
+        let bid_sz_d = stream.alloc_zeros::<f32>(BOOK_LEVELS).context("bid_sz alloc")?;
+        let ask_px_d = stream.alloc_zeros::<f32>(BOOK_LEVELS).context("ask_px alloc")?;
+        let ask_sz_d = stream.alloc_zeros::<f32>(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<Vec<BookFlat>> {
+        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)
+    }
+}

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
+}

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]
+    }
+  ]
+}

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]
+    }
+  ]
+}

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<FixtureEvent>,
+    expected_books: Vec<FixtureBook>,
+    #[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();
+}