diff --git a/docs/superpowers/plans/2026-05-18-real-lob-integration.md b/docs/superpowers/plans/2026-05-18-real-lob-integration.md new file mode 100644 index 000000000..c26086de6 --- /dev/null +++ b/docs/superpowers/plans/2026-05-18-real-lob-integration.md @@ -0,0 +1,3610 @@ +# Real-LOB Integration Implementation Plan + +> **For agentic workers:** REQUIRED SUB-SKILL: Use superpowers:subagent-driven-development (recommended) or superpowers:executing-plans to implement this plan task-by-task. Steps use checkbox (`- [ ]`) syntax for tracking. + +**Goal:** Build a CUDA-only realistic limit-order-book simulator + per-horizon adaptive decision policy + backtest harness that turns the trained ml-alpha model's `p_h(t)` predictions into P&L numbers at IBKR+Scaleway latency. + +**Architecture:** New modules extend the existing `crates/ml-backtesting/` crate alongside `barrier_backtest.rs`. Block-per-backtest CUDA kernels (one warp/block, ~1.4 KB shared mem per block). Three captured CUDA Graphs (trunk perception, LOB step-event, LOB decision) with host branches only between captures. The harness reuses `ml-alpha`'s `MultiHorizonLoader` + `CfcTrunk::capture_graph_a` + `perception_forward_captured` verbatim — zero train-vs-deploy skew. + +**Tech Stack:** Rust 2021 / cudarc 0.19 / CUDA 12.4 / nvcc → cubins / serde-JSON for fixtures / parquet via `parquet` crate for sweep aggregation / `clap` 4 for CLI. + +**Spec:** `docs/superpowers/specs/2026-05-18-real-lob-integration-design.md` @ commit `bfbaea266`. Read it BEFORE starting any task; this plan references it heavily. + +**Branch:** Stay on `ml-alpha-phase-a` for the entire implementation. Commit after each task that has a `Commit` step. + +--- + +## Commit Map (10 atomic commits) + +| # | Title | Scope | +|---|---|---| +| C1 | ml-alpha loader: inference_only mode | Add `inference_only`, `peek_first()`, `next_inference_input()` to `MultiHorizonLoader`. No backtest code yet. | +| C2 | ml-backtesting: order types + SlotTag | Host-only `OrderIntent` / `OrderEvent` / `SlotTag` / `Side` / `LimitKind` enums and packing. | +| C3 | ml-backtesting: policy tree + bytecode | `Strategy` recursive enum, `StopRules`, `LatencyConfig`, `SizingPolicyId`, host-side flattening to bytecode program. | +| C4 | ml-backtesting: build.rs + book_update kernel + sim skeleton | Cubin pipeline; first CUDA kernel + bindings; first 3 Ring 1 fixtures (book mechanics). | +| C5 | ml-backtesting: order_match kernel | Resting-order matching, queue decay, latency-aware fills. 4 more Ring 1 fixtures. | +| C6 | ml-backtesting: pnl_track kernel | Realized/unrealized P&L, segment_complete events. 1 Ring 1 fixture. | +| C7 | ml-backtesting: decision_policy kernel | Bytecode interpreter, per-horizon ISV-Kelly, WeightedByRealizedSharpe aggregator, OCO + overflow. 3 Ring 1 fixtures. | +| C8 | ml-backtesting: harness + trainer parity | `BacktestHarness` orchestrator wiring loader → trunk → sim. Ring 1b bit-equality test. | +| C9 | ml-backtesting: artifacts + aggregate + fxt-backtest binary | summary.json / trades.csv / pnl_curve.bin writers, sweep aggregation, single CLI binary with `run` + `aggregate` subcommands. | +| C10 | ml-backtesting: Ring 2 invariant fuzz tests | Property-based fuzz with N ∈ {1, 16, 256}. | + +After C10: Ring 3 (production-data replay) is OUT of scope for this plan — it requires real data and external validation. Plan ends with C10 producing a working CLI. + +--- + +## Conventions for every task + +1. **Read the spec section referenced in each task before coding.** Tasks cite spec sections (e.g. "see spec §5") rather than restating design rationale. +2. **`SQLX_OFFLINE=true` is set in the environment.** All `cargo` commands assume it. +3. **GPU tests are `#[ignore]`d by default**; run with `--ignored`. Convention follows existing ml-alpha tests (see `crates/ml-alpha/tests/perception_overfit.rs`). +4. **Cubin path convention** follows `crates/ml-alpha/cuda/`: `concat!(env!("OUT_DIR"), "/.cubin")` in source, `build.rs` writes there. +5. **Tests use the real GPU**, no CPU oracle. RTX 3050 locally is enough for N=1 fixtures. Per `feedback_no_cpu_test_fallbacks.md`. +6. **Commits use Conventional Commits** matching existing repo style (`feat(ml-backtesting):`, `perf(...)`, etc.) and end with the Claude co-author trailer. + +--- + +## Commit 1 — ml-alpha loader: `inference_only` mode + +**Why first:** The backtest harness depends on `MultiHorizonLoader::next_inference_input()`. Land this in `ml-alpha` so subsequent commits can `use ml_alpha::data::loader::*` without forward references. + +**Spec reference:** §1 (trainer parity), §7 (orchestrator code). + +### Task 1.1: Extend `MultiHorizonLoaderConfig` with `inference_only` + +**Files:** +- Modify: `crates/ml-alpha/src/data/loader.rs:88-122` (add field to config struct) +- Test: `crates/ml-alpha/src/data/loader.rs` (inline `#[cfg(test)]` module) + +- [ ] **Step 1: Read the current config struct + impl** + +Read `crates/ml-alpha/src/data/loader.rs:85-260` to understand `MultiHorizonLoaderConfig`, `MultiHorizonLoader::new()`, `LoadedFile`, and the per-file label precomputation loop (lines ~199-235). + +- [ ] **Step 2: Write the failing test** + +Append to `crates/ml-alpha/src/data/loader.rs` (inside the existing `#[cfg(test)] mod tests { ... }`, or create one if none exists): + +```rust +#[cfg(test)] +mod inference_mode_tests { + use super::*; + use std::path::PathBuf; + + fn test_fixture_files() -> Vec { + // Mirrors existing test-fixture discovery used elsewhere in the crate. + let root = PathBuf::from(std::env::var("FOXHUNT_TEST_DATA") + .unwrap_or_else(|_| "test_data/futures-baseline".to_string())); + discover_mbp10_files_sorted(&root).expect("test fixtures").into_iter().take(1).collect() + } + + #[test] + fn inference_only_skips_label_precompute() { + let files = test_fixture_files(); + let cfg = MultiHorizonLoaderConfig { + files: files.clone(), + predecoded_dir: files[0].parent().unwrap().to_path_buf(), + seq_len: 1, + horizons: [30, 100, 300, 1000, 6000], + n_max_sequences: 1, + seed: 0, + decision_stride: 1, + inference_only: true, // new field + }; + let loader = MultiHorizonLoader::new(&cfg).expect("loader new"); + // In inference-only mode the per-file labels_full vectors must be empty. + assert!(loader.files_loaded.iter().all(|f| f.labels_full.iter().all(|v| v.is_empty())), + "inference_only=true must not precompute labels"); + } +} +``` + +- [ ] **Step 3: Run the test (it fails on field not existing)** + +```bash +SQLX_OFFLINE=true cargo test -p ml-alpha --lib inference_mode_tests::inference_only_skips_label_precompute +``` + +Expected: compile error — `MultiHorizonLoaderConfig` has no field `inference_only`. + +- [ ] **Step 4: Add the field + skip label precompute when set** + +Edit `crates/ml-alpha/src/data/loader.rs`. Inside `MultiHorizonLoaderConfig`, before the closing `}`: + +```rust + /// When `true`, skip per-file forward-label precomputation + /// (~50 % of file-load cost) and emit raw chronological snapshots + /// only. Used by `crates/ml-backtesting` backtest harness; trainer + /// always passes `false`. See spec §1 trainer-parity row. + pub inference_only: bool, +``` + +Then in `MultiHorizonLoader::new()` find the per-horizon label loop (around lines 213-235 — `for (h_idx, &h) in cfg.horizons.iter().enumerate() { ... }`) and gate it: + +```rust + let mut labels_full: [Vec; 5] = Default::default(); + if !cfg.inference_only { + for (h_idx, &h) in cfg.horizons.iter().enumerate() { + // ... existing label-precompute body unchanged ... + } + } +``` + +The downstream `next_sequence()` already returns NaN labels when out-of-window, so empty `labels_full` is safe — but `next_sequence()` MUST NOT be called in inference-only mode. We add `next_inference_input()` for that in Task 1.3. + +Update every existing trainer call-site that constructs `MultiHorizonLoaderConfig` to pass `inference_only: false`. Find them with: + +```bash +grep -rn "MultiHorizonLoaderConfig {" /home/jgrusewski/Work/foxhunt/crates/ml-alpha/ /home/jgrusewski/Work/foxhunt/services/ /home/jgrusewski/Work/foxhunt/bin/ +``` + +Add `inference_only: false,` to every struct literal found. + +- [ ] **Step 5: Verify the test passes** + +```bash +SQLX_OFFLINE=true cargo test -p ml-alpha --lib inference_mode_tests::inference_only_skips_label_precompute -- --nocapture +``` + +Expected: PASS. + +- [ ] **Step 6: Verify trainer-side tests still pass** + +```bash +SQLX_OFFLINE=true cargo test -p ml-alpha --lib data::loader +``` + +Expected: all pre-existing loader tests PASS (we only gated the label loop, didn't change its body). + +- [ ] **Step 7: Verify the whole workspace still builds** + +```bash +SQLX_OFFLINE=true cargo check --workspace +``` + +Expected: clean build, no errors. (If trainer call-sites were missed they will fail with "missing field `inference_only`".) + +- [ ] **Step 8: Stage but do NOT commit yet** (Task 1.2 + 1.3 ride the same commit) + +```bash +git add crates/ml-alpha/src/data/loader.rs +# Any modified call-sites also staged: +git add $(grep -rln "MultiHorizonLoaderConfig {" crates/ml-alpha/ services/ bin/ 2>/dev/null | xargs) +``` + +### Task 1.2: Add `MultiHorizonLoader::peek_first()` + +**Files:** +- Modify: `crates/ml-alpha/src/data/loader.rs` (extend `impl MultiHorizonLoader`) + +- [ ] **Step 1: Write the failing test** + +Append to the same `inference_mode_tests` module: + +```rust +#[test] +fn peek_first_returns_first_chronological_snapshot() { + let files = test_fixture_files(); + let cfg = MultiHorizonLoaderConfig { + files: files.clone(), + predecoded_dir: files[0].parent().unwrap().to_path_buf(), + seq_len: 1, + horizons: [30, 100, 300, 1000, 6000], + n_max_sequences: 1, + seed: 0, + decision_stride: 1, + inference_only: true, + }; + let loader = MultiHorizonLoader::new(&cfg).expect("loader new"); + let first = loader.peek_first().expect("peek_first ok"); + // First snapshot of first file must have non-zero ts_ns (real data). + assert!(first.ts_ns > 0); + // bid/ask must be sane (10 levels each). + assert_eq!(first.bid_px.len(), 10); + assert_eq!(first.ask_px.len(), 10); +} +``` + +- [ ] **Step 2: Run the test (fails on missing method)** + +```bash +SQLX_OFFLINE=true cargo test -p ml-alpha --lib inference_mode_tests::peek_first_returns_first_chronological_snapshot +``` + +Expected: compile error — no method `peek_first` on `MultiHorizonLoader`. + +- [ ] **Step 3: Implement `peek_first`** + +Find the `impl MultiHorizonLoader { ... }` block in `crates/ml-alpha/src/data/loader.rs`. After `pub fn yielded(&self) -> usize { self.yielded }`, add: + +```rust + /// Return a clone of the first chronological snapshot from the + /// first loaded file. Used to seed CUDA Graph capture in inference + /// callers (e.g. backtest harness — see spec §7). Pure read; no + /// state mutation. Available in both training and inference modes. + pub fn peek_first(&self) -> Result { + let first_file = self.files_loaded.first() + .context("loader has no loaded files")?; + let first_snap = first_file.snapshots.first() + .context("first loaded file has no snapshots")?; + // Mirror the existing Mbp10RawInput-from-snapshot conversion + // pattern used in next_sequence() — same shape. + Ok(snapshot_to_raw_input(first_snap, &first_file.regime_full[0])) + } +``` + +If `snapshot_to_raw_input` doesn't exist as a free function, factor the existing inline conversion from `next_sequence()` into one. Search for the conversion site: + +```bash +grep -n "Mbp10RawInput {" crates/ml-alpha/src/data/loader.rs +``` + +Extract the relevant assignment block into: + +```rust +fn snapshot_to_raw_input(s: &Mbp10Snapshot, regime: &[f32; REGIME_DIM]) -> Mbp10RawInput { + let mut out = Mbp10RawInput::default(); + out.bid_px = s.bid_px; // adjust field names to match Mbp10Snapshot's exact layout; + out.bid_sz = s.bid_sz; // verify by reading crates/ml-alpha/src/data/mod.rs Mbp10Snapshot. + out.ask_px = s.ask_px; + out.ask_sz = s.ask_sz; + out.prev_mid = s.prev_mid; + out.trade_signed_vol = s.trade_signed_vol; + out.trade_count = s.trade_count; + out.ts_ns = s.ts_ns; + out.prev_ts_ns = s.prev_ts_ns; + out.regime = *regime; + out +} +``` + +Then replace the inline conversion site(s) in `next_sequence()` with `snapshot_to_raw_input(snap, regime)`. The function refactor must produce an output identical to before — verify with the existing loader tests. + +- [ ] **Step 4: Verify test passes + existing tests still pass** + +```bash +SQLX_OFFLINE=true cargo test -p ml-alpha --lib data::loader inference_mode_tests +``` + +Expected: all PASS. + +- [ ] **Step 5: Stage** (still no commit — Task 1.3 next) + +```bash +git add crates/ml-alpha/src/data/loader.rs +``` + +### Task 1.3: Add `MultiHorizonLoader::next_inference_input()` + +**Files:** +- Modify: `crates/ml-alpha/src/data/loader.rs` + +- [ ] **Step 1: Write the failing test** + +Append to `inference_mode_tests`: + +```rust +#[test] +fn next_inference_input_yields_chronological_with_stride() { + let files = test_fixture_files(); + let cfg = MultiHorizonLoaderConfig { + files: files.clone(), + predecoded_dir: files[0].parent().unwrap().to_path_buf(), + seq_len: 1, + horizons: [30, 100, 300, 1000, 6000], + n_max_sequences: 0, // unlimited in inference mode + seed: 0, + decision_stride: 4, + inference_only: true, + }; + let mut loader = MultiHorizonLoader::new(&cfg).expect("loader new"); + + let snap0 = loader.next_inference_input().expect("snap 0").expect("not none"); + let snap1 = loader.next_inference_input().expect("snap 1").expect("not none"); + let snap2 = loader.next_inference_input().expect("snap 2").expect("not none"); + + // Chronological monotonicity. + assert!(snap0.ts_ns < snap1.ts_ns); + assert!(snap1.ts_ns < snap2.ts_ns); + + // Inference mode yields EVERY snapshot — decision_stride is the + // CALLER's concern (e.g. "infer every Nth snapshot"). The loader's + // job is just to walk the stream chronologically. +} +``` + +- [ ] **Step 2: Run the test** + +```bash +SQLX_OFFLINE=true cargo test -p ml-alpha --lib inference_mode_tests::next_inference_input_yields_chronological_with_stride +``` + +Expected: compile error — no method `next_inference_input`. + +- [ ] **Step 3: Implement `next_inference_input`** + +Add to `impl MultiHorizonLoader`, after `peek_first`: + +```rust + /// Iterator-style accessor for inference mode: walks every loaded + /// snapshot across every loaded file in chronological order. Returns + /// `Ok(None)` when the stream is exhausted. Unlike `next_sequence`, + /// does NOT slice fixed-length windows and does NOT apply + /// `decision_stride` (caller's concern — see spec §7 orchestrator). + /// Mutates internal cursor state. + pub fn next_inference_input(&mut self) -> Result> { + anyhow::ensure!( + self.cfg.inference_only, + "next_inference_input requires MultiHorizonLoaderConfig.inference_only = true" + ); + loop { + let Some(file) = self.files_loaded.get(self.inference_file_idx) else { + return Ok(None); + }; + if self.inference_snap_idx >= file.snapshots.len() { + self.inference_file_idx += 1; + self.inference_snap_idx = 0; + continue; + } + let snap = &file.snapshots[self.inference_snap_idx]; + let regime = &file.regime_full[self.inference_snap_idx]; + let out = snapshot_to_raw_input(snap, regime); + self.inference_snap_idx += 1; + self.yielded += 1; + return Ok(Some(out)); + } + } +``` + +Add two cursor fields to `MultiHorizonLoader`: + +```rust +pub struct MultiHorizonLoader { + cfg: MultiHorizonLoaderConfig, + rng: ChaCha8Rng, + yielded: usize, + files_loaded: Vec, + // NEW — chronological cursor for inference-mode iteration. + inference_file_idx: usize, + inference_snap_idx: usize, +} +``` + +In `MultiHorizonLoader::new()`, initialise both cursors to `0` in the returned struct literal. In `reset()`, also reset both to `0`. + +- [ ] **Step 4: Verify all tests pass** + +```bash +SQLX_OFFLINE=true cargo test -p ml-alpha --lib data::loader inference_mode_tests +SQLX_OFFLINE=true cargo check --workspace +``` + +Expected: all PASS, clean build. + +- [ ] **Step 5: Commit C1** + +```bash +git add crates/ml-alpha/src/data/loader.rs +git commit -m "$(cat <<'EOF' +feat(ml-alpha): MultiHorizonLoader inference-only mode + +Add inference_only flag to MultiHorizonLoaderConfig that skips per-file +forward-label precomputation, plus peek_first() and next_inference_input() +methods for chronological snapshot streaming. Used by upcoming +ml-backtesting LOB backtest harness; trainer call-sites unaffected +(pass inference_only=false). See docs/superpowers/specs/2026-05-18-real-lob-integration-design.md §1, §7. + +Co-Authored-By: Claude Opus 4.7 +EOF +)" +``` + +Expected: pre-commit hooks pass, commit lands. + +--- + + +## Commit 2 — ml-backtesting: order types + SlotTag + +**Why now:** Host-only types with no GPU dependency; can land before any CUDA work. Establishes the audit-log schema that later kernels write into. + +**Spec reference:** §3 (OrderIntent flat representation), §5b (state ownership row for audit ring). + +### Task 2.1: Extend `crates/ml-backtesting/Cargo.toml` with new deps + +**Files:** +- Modify: `crates/ml-backtesting/Cargo.toml` + +- [ ] **Step 1: Read the current Cargo.toml** + +Read `crates/ml-backtesting/Cargo.toml` end-to-end. Note that it currently depends only on `ml-core`, `serde`, `chrono`, `anyhow`. + +- [ ] **Step 2: Add the required deps** + +Edit `crates/ml-backtesting/Cargo.toml`. After the existing `[dependencies]` block additions: + +```toml +[dependencies] +ml-core.workspace = true +ml-alpha = { path = "../ml-alpha" } # loader + trunk reuse (see spec §1) +serde = { workspace = true, features = ["derive"] } +serde_json.workspace = true +chrono.workspace = true +anyhow.workspace = true +cudarc.workspace = true # block-per-backtest CUDA bindings +parquet.workspace = true # sweep aggregation output (C9) +arrow.workspace = true # parquet schema construction (C9) +bytemuck = { workspace = true, features = ["derive"] } # SlotTag pod cast +clap = { workspace = true, features = ["derive"] } # used by bin/fxt-backtest (C9) +tracing.workspace = true +thiserror.workspace = true +half.workspace = true + +[build-dependencies] +# Cubin compilation (C4). +anyhow.workspace = true + +[dev-dependencies] +ml-alpha = { path = "../ml-alpha" } +tempfile.workspace = true +``` + +If any of these deps aren't already in the workspace `Cargo.toml`'s `[workspace.dependencies]` section, add them by mirroring the version used in `crates/ml-alpha/Cargo.toml` (it uses cudarc, bytemuck, parquet, arrow, clap, tracing, thiserror, half already — they should all be there). + +- [ ] **Step 3: Verify it builds** + +```bash +SQLX_OFFLINE=true cargo check -p ml-backtesting +``` + +Expected: clean build (no new code yet, just deps). + +- [ ] **Step 4: Stage** (no commit yet — Tasks 2.2 + 2.3 ride this commit) + +```bash +git add crates/ml-backtesting/Cargo.toml +``` + +### Task 2.2: Define `SlotTag` packed handle + `Side` / `LimitKind` enums + +**Files:** +- Create: `crates/ml-backtesting/src/order.rs` +- Modify: `crates/ml-backtesting/src/lib.rs` + +- [ ] **Step 1: Write the failing test** + +Create `crates/ml-backtesting/src/order.rs` with: + +```rust +//! Host-side order representation. See spec §3. + +use serde::{Deserialize, Serialize}; + +#[repr(u8)] +#[derive(Clone, Copy, Debug, PartialEq, Eq, Serialize, Deserialize)] +pub enum Side { + Buy = 0, + Sell = 1, +} + +#[repr(u8)] +#[derive(Clone, Copy, Debug, PartialEq, Eq, Serialize, Deserialize)] +pub enum LimitKind { + Limit = 0, + Ioc = 1, + Fok = 2, +} + +#[repr(u8)] +#[derive(Clone, Copy, Debug, PartialEq, Eq, Serialize, Deserialize)] +pub enum SlotKind { + Limit = 0, + Stop = 1, +} + +/// Stable handle for cancel/modify. Packed into u32 so the device-side +/// audit kernel can emit it as a single field. Layout (LSB → MSB): +/// bits 0.. 7 : slot_kind (SlotKind as u8) +/// bits 8.. 15 : slot_index (u8) +/// bits 16.. 31 : gen_counter (u16) — incremented on every slot reuse; +/// distinguishes new occupant from +/// stale references to a previous one +#[derive(Clone, Copy, Debug, PartialEq, Eq, Serialize, Deserialize)] +pub struct SlotTag(pub u32); + +impl SlotTag { + pub fn pack(kind: SlotKind, idx: u8, gen: u16) -> Self { + SlotTag((kind as u32) | ((idx as u32) << 8) | ((gen as u32) << 16)) + } + pub fn kind(self) -> SlotKind { + match (self.0 & 0xFF) as u8 { + 0 => SlotKind::Limit, + 1 => SlotKind::Stop, + _ => unreachable!("invariant: SlotTag.kind() in {0,1}"), + } + } + pub fn index(self) -> u8 { ((self.0 >> 8) & 0xFF) as u8 } + pub fn gen(self) -> u16 { ((self.0 >> 16) & 0xFFFF) as u16 } +} + +#[cfg(test)] +mod tests { + use super::*; + + #[test] + fn slot_tag_roundtrips() { + let t = SlotTag::pack(SlotKind::Stop, 7, 1234); + assert_eq!(t.kind(), SlotKind::Stop); + assert_eq!(t.index(), 7); + assert_eq!(t.gen(), 1234); + } + + #[test] + fn slot_tag_max_gen_fits() { + let t = SlotTag::pack(SlotKind::Limit, 31, u16::MAX); + assert_eq!(t.gen(), u16::MAX); + assert_eq!(t.index(), 31); + } + + #[test] + fn slot_tag_packs_distinct_kinds() { + let a = SlotTag::pack(SlotKind::Limit, 0, 0); + let b = SlotTag::pack(SlotKind::Stop, 0, 0); + assert_ne!(a.0, b.0); + } +} +``` + +Then in `crates/ml-backtesting/src/lib.rs`, after the existing module declarations, add: + +```rust +pub mod order; +``` + +- [ ] **Step 2: Run the tests (they fail to compile because lib.rs doesn't have `order` module yet, then pass once added)** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --lib order::tests +``` + +Expected: PASS (the impl is in the same file as the test). + +- [ ] **Step 3: Stage** + +```bash +git add crates/ml-backtesting/src/order.rs crates/ml-backtesting/src/lib.rs +``` + +### Task 2.3: Add `OrderIntent` + `OrderEvent` + `LimitParams` + +**Files:** +- Modify: `crates/ml-backtesting/src/order.rs` + +- [ ] **Step 1: Write the failing test** + +Append to `crates/ml-backtesting/src/order.rs`: + +```rust +#[derive(Clone, Copy, Debug, Serialize, Deserialize)] +pub struct LimitParams { + pub side: Side, + pub size: u16, + pub px: i32, + pub kind: LimitKind, +} + +/// Host-side ergonomic enum. Mirrors the on-device decision; used for +/// configuration sources (policy YAML) AND for the trade-log audit +/// representation reconstructed from on-device OrderEvent records. +#[derive(Clone, Debug, Serialize, Deserialize)] +pub enum OrderIntent { + SubmitMarket { side: Side, size: u16 }, + SubmitLimit { side: Side, size: u16, px: i32, kind: LimitKind }, + SubmitStop { side: Side, size: u16, trigger_px: i32, limit_px: Option }, + SubmitOcoLimits { leg_a: LimitParams, leg_b: LimitParams }, + Cancel { slot_tag: SlotTag }, + Modify { slot_tag: SlotTag, new_size: Option, new_px: Option }, +} + +/// Flat 24-byte record emitted by the device-side decision kernel into +/// the per-block mapped-pinned audit ring. Host reads + reconstructs +/// `OrderIntent` for trade-log writing. See spec §3. +#[repr(C)] +#[derive(Clone, Copy, Debug, bytemuck::Pod, bytemuck::Zeroable, Serialize, Deserialize)] +pub struct OrderEvent { + pub ts_ns: u64, // 8 + pub slot_tag: u32, // 4 — SlotTag.0 + pub kind: u8, // 1 — OrderEventKind discriminant + pub side: u8, // 1 + pub size: u16, // 2 + pub px: i32, // 4 (ticks; 0 for market) + pub trigger_px: i32, // 4 (ticks; ignored if not stop) + // total: 24 bytes +} + +#[repr(u8)] +#[derive(Clone, Copy, Debug, PartialEq, Eq)] +pub enum OrderEventKind { + SubmitMarket = 0, + SubmitLimit = 1, + SubmitIoc = 2, + SubmitFok = 3, + SubmitStopMarket = 4, + SubmitStopLimit = 5, + Cancel = 6, + Modify = 7, +} + +#[cfg(test)] +mod intent_tests { + use super::*; + + #[test] + fn order_event_is_24_bytes_packed() { + assert_eq!(std::mem::size_of::(), 24); + assert_eq!(std::mem::align_of::() <= 8, true); + } + + #[test] + fn intent_serde_roundtrip_for_oco() { + let oco = OrderIntent::SubmitOcoLimits { + leg_a: LimitParams { side: Side::Buy, size: 1, px: 550000, kind: LimitKind::Limit }, + leg_b: LimitParams { side: Side::Sell, size: 1, px: 551000, kind: LimitKind::Limit }, + }; + let j = serde_json::to_string(&oco).expect("ser"); + let back: OrderIntent = serde_json::from_str(&j).expect("de"); + match back { + OrderIntent::SubmitOcoLimits { leg_a, leg_b } => { + assert_eq!(leg_a.px, 550000); + assert_eq!(leg_b.px, 551000); + } + _ => panic!("wrong variant after roundtrip"), + } + } +} +``` + +- [ ] **Step 2: Run all order tests** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --lib order +``` + +Expected: PASS (5 tests total). + +- [ ] **Step 3: Commit C2** + +```bash +git add crates/ml-backtesting/Cargo.toml crates/ml-backtesting/src/order.rs crates/ml-backtesting/src/lib.rs +git commit -m "$(cat <<'EOF' +feat(ml-backtesting): host-side order types + SlotTag + +OrderIntent / OrderEvent / SlotTag / Side / LimitKind defined per spec +§3. OrderEvent is repr(C) Pod, 24 bytes — matches the on-device audit +record layout the decision kernel will write (C7). SlotTag packs +{kind, index, gen_counter} into u32 for cancel/modify robustness +against slot reuse. No CUDA dependency yet. + +Adds cudarc, parquet, arrow, bytemuck, clap, tracing, thiserror, half, +serde_json, tempfile to ml-backtesting deps in preparation for upcoming +LOB sim + harness + binary commits. + +Co-Authored-By: Claude Opus 4.7 +EOF +)" +``` + +Expected: pre-commit hooks pass, commit lands. + +--- + +## Commit 3 — ml-backtesting: policy tree + bytecode + +**Why now:** Host-only types; the bytecode emission is the contract the decision-policy kernel (C7) will interpret. Defining it now means C7 can be implemented against a stable schema. + +**Spec reference:** §6 (Strategy tree, EnsembleAggregator, ConflictPolicy, Regime, default_strategy, host-side flattening). + +### Task 3.1: Define `Strategy` recursive enum + `StrategyConfig` + +**Files:** +- Create: `crates/ml-backtesting/src/policy/mod.rs` +- Create: `crates/ml-backtesting/src/policy/composition.rs` +- Create: `crates/ml-backtesting/src/policy/sizing.rs` +- Create: `crates/ml-backtesting/src/policy/stop_rules.rs` +- Create: `crates/ml-backtesting/src/policy/latency.rs` +- Modify: `crates/ml-backtesting/src/lib.rs` + +- [ ] **Step 1: Write the failing test (top-level policy module + types)** + +Create `crates/ml-backtesting/src/policy/mod.rs`: + +```rust +//! Multi-strategy composition + per-leaf sizing/risk config. See spec §6. + +pub mod composition; +pub mod latency; +pub mod sizing; +pub mod stop_rules; + +pub use composition::*; +pub use latency::*; +pub use sizing::*; +pub use stop_rules::*; + +use serde::{Deserialize, Serialize}; + +/// One horizon-anchored leaf strategy. The recursive `Strategy::Leaf` +/// wraps this. See spec §5 (per-horizon ISV-Kelly) + §6. +#[derive(Clone, Debug, Serialize, Deserialize)] +pub struct StrategyConfig { + pub horizon_idx: u8, // 0..N_HORIZONS — exactly one horizon per leaf + pub sizing_policy: SizingPolicyId, + pub sl_tp_rules: StopRules, + pub max_concurrent_lots: u16, +} + +/// Recursive composition tree (RegimeSwitch -> Portfolio -> Ensemble +/// of leaves is the canonical nesting). Default constructor +/// `default_strategy()` returns the adaptive 5-horizon Ensemble with +/// WeightedByRealizedSharpe. +#[derive(Clone, Debug, Serialize, Deserialize)] +pub enum Strategy { + Leaf(StrategyConfig), + Ensemble { children: Vec, aggregator: EnsembleAggregator }, + Portfolio { + children: Vec, + capital_allocation: Vec, + conflict_policy: ConflictPolicy, + }, + RegimeSwitch { regime_to_child: std::collections::HashMap }, +} + +pub const N_HORIZONS: usize = 5; // mirror crates/ml-alpha/src/heads.rs + +impl Strategy { + /// Default v1 policy: per-horizon adaptive Ensemble. + /// See spec §6 default_strategy(...). No static horizon mask; + /// WeightedByRealizedSharpe auto-shifts capital from in-run evidence. + pub fn default_for(max_lots: u16) -> Self { + let children = (0..N_HORIZONS as u8).map(|h| Strategy::Leaf(StrategyConfig { + horizon_idx: h, + sizing_policy: SizingPolicyId::IsvKelly, + sl_tp_rules: StopRules::default(), + max_concurrent_lots: max_lots, + })).collect(); + Strategy::Ensemble { children, aggregator: EnsembleAggregator::WeightedByRealizedSharpe } + } +} + +#[cfg(test)] +mod tests { + use super::*; + #[test] + fn default_strategy_has_5_horizon_leaves() { + let s = Strategy::default_for(2); + match s { + Strategy::Ensemble { children, aggregator } => { + assert_eq!(children.len(), 5); + assert!(matches!(aggregator, EnsembleAggregator::WeightedByRealizedSharpe)); + for (i, c) in children.iter().enumerate() { + match c { + Strategy::Leaf(cfg) => { + assert_eq!(cfg.horizon_idx as usize, i); + assert_eq!(cfg.max_concurrent_lots, 2); + } + _ => panic!("expected Leaf at index {i}"), + } + } + } + _ => panic!("expected top-level Ensemble"), + } + } +} +``` + +Create `crates/ml-backtesting/src/policy/composition.rs`: + +```rust +use serde::{Deserialize, Serialize}; + +#[derive(Clone, Copy, Debug, PartialEq, Eq, Serialize, Deserialize)] +pub enum EnsembleAggregator { + Mean, + MaxConfidence, + WeightedByValAuc, + WeightedByRealizedSharpe, // ADAPTIVE — see spec §5 + §6 +} + +#[derive(Clone, Copy, Debug, PartialEq, Eq, Serialize, Deserialize)] +pub enum ConflictPolicy { + NetOut, + FirstSubmitWins, + BlockOpposing, +} + +#[derive(Clone, Copy, Debug, Hash, PartialEq, Eq, Serialize, Deserialize)] +pub enum Regime { + SpreadQ1, + SpreadQ2, + SpreadQ3, + SpreadQ4, +} +``` + +Create `crates/ml-backtesting/src/policy/sizing.rs`: + +```rust +use serde::{Deserialize, Serialize}; + +#[derive(Clone, Copy, Debug, PartialEq, Eq, Serialize, Deserialize)] +pub enum SizingPolicyId { + IsvKelly, + FixedLots(u16), +} +``` + +Create `crates/ml-backtesting/src/policy/stop_rules.rs`: + +```rust +use serde::{Deserialize, Serialize}; + +#[derive(Clone, Copy, Debug, PartialEq, Default, Serialize, Deserialize)] +pub struct StopRules { + /// Hard stop-loss distance in ticks (i32 price-unit ticks). `0` = disabled. + pub stop_loss_ticks: i32, + /// Take-profit distance in ticks. `0` = disabled. + pub take_profit_ticks: i32, + /// Max holding-time in nanoseconds. `0` = disabled (hold until counter-signal). + pub max_hold_ns: u64, +} +``` + +Create `crates/ml-backtesting/src/policy/latency.rs`: + +```rust +use serde::{Deserialize, Serialize}; + +/// Latency model. Default: Constant(100ms) reflects IBKR + Scaleway. +/// See spec §4. +#[derive(Clone, Debug, Serialize, Deserialize)] +pub enum LatencyConfig { + Constant { total_ns: u32 }, + /// Per-decision latency drawn from a host-uploaded empirical + /// distribution. Indices wrap modulo `samples.len()`. + Empirical { samples: Vec }, + Decomposed { + inference_ns: u32, + decision_ns: u32, + wire_ns: u32, + match_ns: u32, + }, +} + +impl Default for LatencyConfig { + fn default() -> Self { LatencyConfig::Constant { total_ns: 100_000_000 } } +} + +impl LatencyConfig { + /// Sample a latency in ns for the given decision-event index. + pub fn sample_ns(&self, decision_idx: u64) -> u32 { + match self { + LatencyConfig::Constant { total_ns } => *total_ns, + LatencyConfig::Empirical { samples } => { + let i = (decision_idx as usize) % samples.len(); + samples[i].max(0.0) as u32 + } + LatencyConfig::Decomposed { inference_ns, decision_ns, wire_ns, match_ns } => + inference_ns + decision_ns + wire_ns + match_ns, + } + } +} + +#[cfg(test)] +mod tests { + use super::*; + #[test] + fn default_is_100ms() { + assert!(matches!(LatencyConfig::default(), + LatencyConfig::Constant { total_ns: 100_000_000 })); + } + #[test] + fn empirical_wraps() { + let cfg = LatencyConfig::Empirical { samples: vec![10.0, 20.0, 30.0] }; + assert_eq!(cfg.sample_ns(0), 10); + assert_eq!(cfg.sample_ns(3), 10); + assert_eq!(cfg.sample_ns(7), 20); + } + #[test] + fn decomposed_sums() { + let cfg = LatencyConfig::Decomposed { + inference_ns: 1_000_000, decision_ns: 100_000, + wire_ns: 50_000_000, match_ns: 5_000_000 }; + assert_eq!(cfg.sample_ns(0), 56_100_000); + } +} +``` + +Update `crates/ml-backtesting/src/lib.rs` — after `pub mod order;` add: + +```rust +pub mod policy; +``` + +- [ ] **Step 2: Run the tests** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --lib policy +``` + +Expected: PASS (5 tests across modules). + +- [ ] **Step 3: Stage** (no commit — Task 3.2 next) + +```bash +git add crates/ml-backtesting/src/policy/ crates/ml-backtesting/src/lib.rs +``` + +### Task 3.2: Bytecode flattening (`Program` + `flatten()`) + +**Files:** +- Modify: `crates/ml-backtesting/src/policy/mod.rs` + +- [ ] **Step 1: Write the failing test** + +Append to `crates/ml-backtesting/src/policy/mod.rs`: + +```rust +/// Device-side bytecode opcodes. The decision-policy kernel (C7) +/// interprets these against current alpha probs + per-horizon +/// ISV-Kelly state to mutate shared-mem orders[]. Layout MUST stay +/// in sync with cuda/decision_policy.cu. See spec §6 host-side +/// flattening + §7 graph capture. +#[repr(u8)] +#[derive(Clone, Copy, Debug, PartialEq, Eq, bytemuck::Pod, bytemuck::Zeroable)] +pub enum OpCode { + NoOp = 0, + EvalRegime = 1, // arg = regime_id (u8 in arg0); pushes 1.0 if current regime matches, else 0.0 + EmitPerHorizonSize = 2, // arg0 = horizon_idx (u8); arg1 = max_lots (u16) — pushes signed-size from per-horizon ISV-Kelly + AggWeightedSharpe = 3, // pops N (arg0) per-horizon signed-sizes; pushes weighted aggregate + AggMean = 4, // pops N (arg0); pushes mean + AggMaxConfidence = 5, + ApplyConflict = 6, // arg0 = ConflictPolicy as u8 + WriteOrder = 7, // pops one signed-size; emits OrderEvent + mutates orders[] + PushScalar = 8, // immediate float in arg0..3 (i32 reinterpret) — pushes f32 + BranchIfRegime = 9, // arg0 = regime_id; arg1..3 = program offset to jump if not matching +} + +#[repr(C)] +#[derive(Clone, Copy, Debug, bytemuck::Pod, bytemuck::Zeroable)] +pub struct Instruction { + pub op: u8, + pub arg0: u8, + pub arg1: u16, + pub arg2: u32, // wide arg (e.g. PushScalar immediate via bytemuck::cast) +} + +/// Flat device-uploadable program. One per backtest cell; +/// indexed by blockIdx.x in device-global memory (NOT __constant__, +/// see spec §6 "Why not __constant__"). +#[derive(Clone, Debug, Default)] +pub struct Program { + pub instructions: Vec, +} + +impl Strategy { + /// Recursive walk → flat bytecode. See spec §6 host-side flattening. + pub fn flatten(&self) -> Program { + let mut prog = Program::default(); + flatten_into(self, &mut prog); + prog + } +} + +fn flatten_into(s: &Strategy, prog: &mut Program) { + match s { + Strategy::Leaf(cfg) => { + prog.instructions.push(Instruction { + op: OpCode::EmitPerHorizonSize as u8, + arg0: cfg.horizon_idx, + arg1: cfg.max_concurrent_lots, + arg2: 0, + }); + } + Strategy::Ensemble { children, aggregator } => { + for c in children { flatten_into(c, prog); } + let agg_op = match aggregator { + EnsembleAggregator::Mean => OpCode::AggMean, + EnsembleAggregator::MaxConfidence => OpCode::AggMaxConfidence, + EnsembleAggregator::WeightedByValAuc => OpCode::AggMean, // v1: treated as Mean; weights live elsewhere if needed + EnsembleAggregator::WeightedByRealizedSharpe => OpCode::AggWeightedSharpe, + }; + prog.instructions.push(Instruction { + op: agg_op as u8, arg0: children.len() as u8, arg1: 0, arg2: 0, + }); + } + Strategy::Portfolio { children, capital_allocation: _, conflict_policy } => { + for c in children { flatten_into(c, prog); } + // Portfolio = sum children (Mean weighted by capital — v1 simplification + // implements Mean and lets conflict_policy resolve at end). + prog.instructions.push(Instruction { + op: OpCode::AggMean as u8, arg0: children.len() as u8, arg1: 0, arg2: 0, + }); + prog.instructions.push(Instruction { + op: OpCode::ApplyConflict as u8, arg0: *conflict_policy as u8, arg1: 0, arg2: 0, + }); + } + Strategy::RegimeSwitch { regime_to_child } => { + // Linear test+branch chain. Each regime gets: BranchIfRegime → flatten(child). + // Last regime jumps over a final NoOp guard. + let mut child_offsets: Vec = Vec::with_capacity(regime_to_child.len()); + for (regime, child) in regime_to_child { + let branch_idx = prog.instructions.len(); + prog.instructions.push(Instruction { + op: OpCode::BranchIfRegime as u8, + arg0: *regime as u8, + arg1: 0, + arg2: 0, // patched below + }); + child_offsets.push(branch_idx); + flatten_into(child, prog); + } + let end = prog.instructions.len(); + for branch_idx in child_offsets { + prog.instructions[branch_idx].arg2 = end as u32; + } + } + } + // Top-level: emit a WriteOrder (only the OUTERMOST flatten() result writes). + // We add it here unconditionally; nested flatten_into recursions just leave + // values on the stack for the parent to aggregate. The trick: only the + // top-level public flatten() entry-point appends WriteOrder once at the end. +} + +impl Program { + /// Append the top-level WriteOrder. Called once by Strategy::flatten(). + pub(crate) fn finalize(&mut self) { + self.instructions.push(Instruction { + op: OpCode::WriteOrder as u8, arg0: 0, arg1: 0, arg2: 0, + }); + } +} + +// Replace Strategy::flatten() body to call finalize at the end: +// ... existing flatten_into ... +// prog.finalize(); +// prog +// (Update the impl above accordingly.) + +#[cfg(test)] +mod flatten_tests { + use super::*; + + #[test] + fn default_strategy_flattens_to_5_emits_plus_agg_plus_write() { + let s = Strategy::default_for(3); + let p = s.flatten(); + // Expected sequence: 5 × EmitPerHorizonSize, 1 × AggWeightedSharpe, 1 × WriteOrder + assert_eq!(p.instructions.len(), 7); + for i in 0..5 { + assert_eq!(p.instructions[i].op, OpCode::EmitPerHorizonSize as u8); + assert_eq!(p.instructions[i].arg0, i as u8); + assert_eq!(p.instructions[i].arg1, 3); + } + assert_eq!(p.instructions[5].op, OpCode::AggWeightedSharpe as u8); + assert_eq!(p.instructions[5].arg0, 5); + assert_eq!(p.instructions[6].op, OpCode::WriteOrder as u8); + } + + #[test] + fn single_leaf_flattens_to_one_emit_plus_write() { + let leaf = Strategy::Leaf(StrategyConfig { + horizon_idx: 4, sizing_policy: SizingPolicyId::IsvKelly, + sl_tp_rules: StopRules::default(), max_concurrent_lots: 1, + }); + let p = leaf.flatten(); + assert_eq!(p.instructions.len(), 2); + assert_eq!(p.instructions[0].op, OpCode::EmitPerHorizonSize as u8); + assert_eq!(p.instructions[1].op, OpCode::WriteOrder as u8); + } +} +``` + +Update `Strategy::flatten()` to call `prog.finalize()` before returning (per the comment in the code above — replace the existing body with the version that calls finalize). + +- [ ] **Step 2: Run the tests** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --lib policy +``` + +Expected: PASS (7 tests total — 5 from Task 3.1 + 2 new). + +- [ ] **Step 3: Commit C3** + +```bash +git add crates/ml-backtesting/src/policy/ crates/ml-backtesting/src/lib.rs +git commit -m "$(cat <<'EOF' +feat(ml-backtesting): policy tree + bytecode program + +Recursive Strategy enum (Leaf | Ensemble | Portfolio | RegimeSwitch) +with EnsembleAggregator including WeightedByRealizedSharpe for adaptive +horizon weighting (no static horizon_mask per spec §6). +Strategy::default_for(max_lots) emits the 5-horizon adaptive Ensemble +that will be the v1 default policy. + +Strategy::flatten() walks the tree → linear bytecode Program (Pod +Instruction { op, arg0, arg1, arg2 }) for upload to device-global +memory at slot blockIdx.x. The C7 decision_policy.cu kernel will +interpret this bytecode against per-horizon ISV-Kelly state. + +LatencyConfig defaulted to Constant(100ms) per IBKR+Scaleway baseline. +StopRules, SizingPolicyId, ConflictPolicy, Regime defined as host-side +config types. Co-Authored-By: Claude Opus 4.7 +EOF +)" +``` + +Expected: pre-commit hooks pass, commit lands. + +--- +## Commit 4 — ml-backtesting: build.rs + book_update kernel + sim skeleton + +**Why now:** Establishes the CUDA pipeline (build.rs → cubin → cudarc loading → fixture harness) end-to-end with the simplest kernel. Every subsequent kernel commit follows the exact same pattern, so getting this right pays off. + +**Spec reference:** §2 (shared-mem layout), §8 Ring 1 fixtures, §10 (`feedback_no_nvrtc.md`, `feedback_no_atomicadd.md`, `pearl_build_rs_rerun_if_env_changed.md`). + +### Task 4.1: `build.rs` cubin compilation + +**Files:** +- Create: `crates/ml-backtesting/build.rs` + +- [ ] **Step 1: Read the reference build.rs from ml-alpha** + +```bash +cat crates/ml-alpha/build.rs | head -80 +``` + +Note the pattern: walks `cuda/` dir, invokes `nvcc --cubin -arch=...`, writes to `$OUT_DIR/.cubin`, emits `cargo:rerun-if-changed` AND `cargo:rerun-if-env-changed` (per `pearl_build_rs_rerun_if_env_changed.md`). + +- [ ] **Step 2: Create `crates/ml-backtesting/build.rs`** + +```rust +//! Compile every .cu file under cuda/ to a cubin in $OUT_DIR. +//! Mirrors crates/ml-alpha/build.rs. Per feedback_no_nvrtc.md: +//! cubins only, no runtime nvrtc. + +use std::path::{Path, PathBuf}; +use std::process::Command; + +fn main() -> anyhow::Result<()> { + let cuda_dir = PathBuf::from(env!("CARGO_MANIFEST_DIR")).join("cuda"); + let out_dir = PathBuf::from(std::env::var("OUT_DIR")?); + + // Per pearl_build_rs_rerun_if_env_changed.md: pair every env::var + // with rerun-if-env-changed. + let arch = std::env::var("FOXHUNT_CUDA_ARCH").unwrap_or_else(|_| "sm_89".into()); + println!("cargo:rerun-if-env-changed=FOXHUNT_CUDA_ARCH"); + println!("cargo:rerun-if-env-changed=CUDA_PATH"); + + let nvcc = std::env::var("NVCC") + .map(PathBuf::from) + .unwrap_or_else(|_| PathBuf::from("nvcc")); + println!("cargo:rerun-if-env-changed=NVCC"); + + if !cuda_dir.exists() { + // No kernels yet — first build before any .cu lands. + return Ok(()); + } + + println!("cargo:rerun-if-changed={}", cuda_dir.display()); + for entry in std::fs::read_dir(&cuda_dir)? { + let entry = entry?; + let path = entry.path(); + if path.extension().and_then(|s| s.to_str()) != Some("cu") { continue; } + let stem = path.file_stem().unwrap().to_string_lossy().to_string(); + let cubin = out_dir.join(format!("{stem}.cubin")); + + println!("cargo:rerun-if-changed={}", path.display()); + + let status = Command::new(&nvcc) + .args(["--cubin", "-O3", "-std=c++17"]) + .arg(format!("-arch={arch}")) + .arg("-o").arg(&cubin) + .arg(&path) + .status() + .map_err(|e| anyhow::anyhow!("spawn nvcc: {e}"))?; + anyhow::ensure!(status.success(), "nvcc failed for {}", path.display()); + } + Ok(()) +} +``` + +- [ ] **Step 3: Verify it builds (with no .cu files yet — should no-op)** + +```bash +SQLX_OFFLINE=true cargo build -p ml-backtesting +``` + +Expected: clean build (cuda/ dir doesn't exist yet so build.rs early-returns). + +- [ ] **Step 4: Stage** (no commit — Task 4.2 next) + +```bash +git add crates/ml-backtesting/build.rs crates/ml-backtesting/Cargo.toml +``` + +### Task 4.2: Per-block shared-mem layout (CUDA header) + book_update kernel + +**Files:** +- Create: `crates/ml-backtesting/cuda/lob_state.cuh` +- Create: `crates/ml-backtesting/cuda/book_update.cu` + +- [ ] **Step 1: Create shared-mem layout header** + +`crates/ml-backtesting/cuda/lob_state.cuh`: + +```cuda +// 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 +// LobSharedMem (used only for size assertions on host side). + +#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]; +}; + +struct LimitSlot { + unsigned char active; // 0=free, 1=resting, 2=in-flight + unsigned char side; // 0=buy, 1=sell + unsigned char type; // 0=Limit, 1=Ioc, 2=Fok + unsigned char oco_pair; // 0xFF = none + float price; // ticks (i32 reinterpret) + float size_remaining; + float queue_position; + unsigned long long arrival_ts_ns; + unsigned short gen_counter; +}; + +struct StopSlot { + unsigned char active; + unsigned char side; + unsigned char type; // 4=StopMarket, 5=StopLimit + unsigned char oco_pair; + float trigger_price; + float limit_price; + float size; + unsigned long long arrival_ts_ns; + unsigned short gen_counter; +}; + +struct Orders { + LimitSlot limits[MAX_LIMITS]; + StopSlot stops[MAX_STOPS]; +}; + +struct Pos { + int position_lots; + float vwap_entry; + float realized_pnl; + float peak_equity; + unsigned int submission_overflow; + unsigned int open_horizon_mask; +}; + +struct IsvKellyState { + float pnl_ema_win; + float pnl_ema_loss; + float win_rate_ema; + unsigned int n_trades_seen; + float realised_return_var; + float recent_sharpe; +}; + +#endif // LOB_STATE_CUH +``` + +- [ ] **Step 2: Create `book_update.cu`** + +`crates/ml-backtesting/cuda/book_update.cu`: + +```cuda +// MBP-10 event apply. Mutates shared-mem Book state. +// One block per backtest. Thread tid handles one of 10 levels per side +// (20 threads used out of 32 in the warp; remainder idle). +// Per feedback_no_atomicadd.md: no atomics — single-writer per level. + +#include "lob_state.cuh" + +// Event types (mirror Mbp10RawInput; one event = one snapshot apply). +// The full snapshot REPLACES the book — simpler than tracking deltas, +// matches MBP-10 semantics (each snapshot is a full top-10 image). +extern "C" __global__ void book_update_apply_snapshot( + const float* __restrict__ bid_px_in, // [N_backtests, 10] — broadcast same row across all + const float* __restrict__ bid_sz_in, + const float* __restrict__ ask_px_in, + const float* __restrict__ ask_sz_in, + // Per-block book state pointer (device-global, blockIdx.x indexed). + Book* __restrict__ books_out // [N_backtests] +) { + int b = blockIdx.x; + int tid = threadIdx.x; + if (tid >= 10) return; + + // Broadcast: same snapshot applied to every backtest. bid_px_in is + // shape [10] (single backtest) for v1; the simulator owns N_backtests + // book replicas but the input is one snapshot. Indexing by tid only. + books_out[b].bid_px[tid] = bid_px_in[tid]; + books_out[b].bid_sz[tid] = bid_sz_in[tid]; + books_out[b].ask_px[tid] = ask_px_in[tid]; + books_out[b].ask_sz[tid] = ask_sz_in[tid]; +} +``` + +- [ ] **Step 3: Trigger nvcc compile via cargo build** + +```bash +SQLX_OFFLINE=true cargo build -p ml-backtesting +``` + +Expected: nvcc emits `target/.../out/book_update.cubin`. Verify: + +```bash +find target -name "book_update.cubin" | head +``` + +Expected: at least one `.cubin` file printed. + +- [ ] **Step 4: Stage** (no commit yet — Tasks 4.3 + 4.4 ride this commit) + +```bash +git add crates/ml-backtesting/cuda/ +``` + +### Task 4.3: Rust bindings + `LobSimCuda` skeleton + +**Files:** +- Create: `crates/ml-backtesting/src/lob/mod.rs` +- Create: `crates/ml-backtesting/src/sim.rs` +- Modify: `crates/ml-backtesting/src/lib.rs` + +- [ ] **Step 1: Write the failing test** + +Create `crates/ml-backtesting/src/sim.rs`: + +```rust +//! Block-per-backtest LOB simulator. See spec §2 + §5b. + +use anyhow::{Context, Result}; +use cudarc::driver::{CudaContext, CudaModule, CudaStream, CudaSlice, LaunchConfig, PushKernelArg}; +use std::sync::Arc; + +use crate::lob::{BookFlat, BOOK_LEVELS}; + +pub struct LobSimCuda { + n_backtests: usize, + stream: Arc, + module_book: Arc, + // Device-global per-backtest state buffers. + pub(crate) books_d: CudaSlice, // [n_backtests, 4 * BOOK_LEVELS] = bid_px|bid_sz|ask_px|ask_sz interleaved + // Mapped-pinned single-snapshot input buffers (size 10 each, broadcast). + pub(crate) bid_px_h: CudaSlice, + pub(crate) bid_sz_h: CudaSlice, + pub(crate) ask_px_h: CudaSlice, + pub(crate) ask_sz_h: CudaSlice, +} + +impl LobSimCuda { + pub fn new(n_backtests: usize, ctx: &Arc) -> Result { + anyhow::ensure!(n_backtests > 0, "n_backtests must be > 0"); + let stream = ctx.default_stream(); + let module_book = ctx.load_module( + cudarc::nvrtc::Ptx::from_src(include_bytes!(concat!(env!("OUT_DIR"), "/book_update.cubin")).as_ref()) + ).context("load book_update cubin")?; + + // Per-backtest book state. 4 fields × 10 levels = 40 floats. + let books_d = stream.alloc_zeros::(n_backtests * 4 * BOOK_LEVELS) + .context("alloc books_d")?; + + // Mapped-pinned input snapshots (single-backtest broadcast). + let bid_px_h = stream.alloc_zeros::(BOOK_LEVELS)?; + let bid_sz_h = stream.alloc_zeros::(BOOK_LEVELS)?; + let ask_px_h = stream.alloc_zeros::(BOOK_LEVELS)?; + let ask_sz_h = stream.alloc_zeros::(BOOK_LEVELS)?; + + Ok(Self { n_backtests, stream, module_book, books_d, bid_px_h, bid_sz_h, ask_px_h, ask_sz_h }) + } + + 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; 10], bid_sz: &[f32; 10], + ask_px: &[f32; 10], ask_sz: &[f32; 10]) -> Result<()> { + self.stream.memcpy_htod(bid_px, &mut self.bid_px_h)?; + self.stream.memcpy_htod(bid_sz, &mut self.bid_sz_h)?; + self.stream.memcpy_htod(ask_px, &mut self.ask_px_h)?; + self.stream.memcpy_htod(ask_sz, &mut self.ask_sz_h)?; + + let func = self.module_book.load_function("book_update_apply_snapshot") + .context("load book_update_apply_snapshot")?; + let cfg = LaunchConfig { + grid_dim: (self.n_backtests as u32, 1, 1), + block_dim: (32, 1, 1), + shared_mem_bytes: 0, + }; + let mut launch = self.stream.launch_builder(&func); + unsafe { + launch + .arg(&self.bid_px_h) + .arg(&self.bid_sz_h) + .arg(&self.ask_px_h) + .arg(&self.ask_sz_h) + .arg(&mut self.books_d) + .launch(cfg)?; + } + self.stream.synchronize()?; + Ok(()) + } + + /// Read back all backtest book states from device. For tests only. + pub fn read_books(&self) -> Result> { + let raw = self.stream.memcpy_dtov(&self.books_d)?; + let mut out = Vec::with_capacity(self.n_backtests); + for b in 0..self.n_backtests { + let off = b * 4 * BOOK_LEVELS; + let mut bf = BookFlat::default(); + for k in 0..BOOK_LEVELS { + bf.bid_px[k] = raw[off + 0 * BOOK_LEVELS + k]; + bf.bid_sz[k] = raw[off + 1 * 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) + } +} +``` + +Create `crates/ml-backtesting/src/lob/mod.rs`: + +```rust +//! Host-side mirror of cuda/lob_state.cuh layouts. +//! Used for test-side book reads + size assertions. + +pub const BOOK_LEVELS: usize = 10; + +#[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], +} +``` + +Update `crates/ml-backtesting/src/lib.rs`: + +```rust +pub mod lob; +pub mod sim; +``` + +- [ ] **Step 2: Verify it compiles** + +```bash +SQLX_OFFLINE=true cargo check -p ml-backtesting +``` + +Expected: clean build. + +- [ ] **Step 3: Stage** (no commit — fixture harness next in 4.4) + +```bash +git add crates/ml-backtesting/src/sim.rs crates/ml-backtesting/src/lob/ crates/ml-backtesting/src/lib.rs +``` + +### Task 4.4: Fixture format + harness + first 3 fixtures + +**Files:** +- Create: `crates/ml-backtesting/tests/lob_sim_fixtures.rs` +- Create: `crates/ml-backtesting/tests/fixtures/book_update_replace.json` +- Create: `crates/ml-backtesting/tests/fixtures/book_update_two_step.json` +- Create: `crates/ml-backtesting/tests/fixtures/book_update_n_backtests.json` + +- [ ] **Step 1: Create fixture format + 3 JSON fixtures** + +`crates/ml-backtesting/tests/fixtures/book_update_replace.json`: + +```json +{ + "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`: + +```json +{ + "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/fixtures/book_update_n_backtests.json`: + +```json +{ + "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 +} +``` + +- [ ] **Step 2: Create the fixture harness + tests** + +`crates/ml-backtesting/tests/lob_sim_fixtures.rs`: + +```rust +//! Ring 1 fixture tests. N=1 (or small N) bit-exact GPU assertions. +//! See spec §8 Ring 1. +//! Run with: `cargo test -p ml-backtesting --test lob_sim_fixtures -- --ignored --nocapture` + +use anyhow::{Context, Result}; +use ml_backtesting::lob::{BookFlat, BOOK_LEVELS}; +use ml_backtesting::sim::LobSimCuda; +use serde::Deserialize; +use std::path::Path; +use std::sync::Arc; + +#[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")?; + println!("running fixture: {}", fx.name); + + let ctx = cudarc::driver::CudaContext::new(0).context("cuda ctx 0")?; + let mut sim = LobSimCuda::new(fx.n_backtests, &Arc::new(ctx))?; + + 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 = if fx.expected_books_apply_to_all { + vec![&fx.expected_books[0]; fx.n_backtests] + } 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(); +} +``` + +- [ ] **Step 3: Run the fixture tests** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test lob_sim_fixtures -- --ignored --nocapture +``` + +Expected: 3 PASS. + +- [ ] **Step 4: Commit C4** + +```bash +git add crates/ml-backtesting/build.rs \ + crates/ml-backtesting/cuda/ \ + crates/ml-backtesting/src/sim.rs \ + crates/ml-backtesting/src/lob/ \ + crates/ml-backtesting/src/lib.rs \ + crates/ml-backtesting/tests/lob_sim_fixtures.rs \ + crates/ml-backtesting/tests/fixtures/ +git commit -m "$(cat <<'EOF' +feat(ml-backtesting): build.rs + book_update CUDA kernel + sim skeleton + +build.rs compiles every .cu under crates/ml-backtesting/cuda/ to +$OUT_DIR/.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. + +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 (Book, +Orders.LimitSlot[32], Orders.StopSlot[16], Pos, IsvKellyState[5]) that +all subsequent kernels share — see spec §2. + +LobSimCuda host struct allocs device book state + mapped-pinned input +snapshots. apply_snapshot launches the book_update kernel and synchronises. + +Three Ring 1 fixture tests (book_update_replace, _two_step, +_n_backtests=4) — N≤4 bit-exact GPU oracle assertions loaded from JSON. + +Co-Authored-By: Claude Opus 4.7 +EOF +)" +``` + +Expected: pre-commit hooks pass, commit lands. + +--- + +## Commit 5 — ml-backtesting: order_match kernel + latency-aware fills + +**Why now:** With book state working, we can layer in resting orders, queue position, and latency-bounded marketability checks. This commit makes the sim able to accept market orders and emit fills. + +**Spec reference:** §2 (`orders.limits[]`, `orders.stops[]`, `oco_pair`, `arrival_ts_ns`), §3 (OrderEvent), §4 (LatencyConfig + arrival-event timestamp mechanics), §8 Ring 1 fixtures `market_order_consumes_top`, `limit_rest_queue_init`, `limit_fill_via_trade`, `latency_in_flight_miss`. + +### Task 5.1: Per-block `Orders` + `Pos` state buffers + `OrderEvent` audit ring + +**Files:** +- Modify: `crates/ml-backtesting/src/sim.rs` (extend `LobSimCuda` with new buffers) +- Modify: `crates/ml-backtesting/src/lob/mod.rs` (host mirrors) + +- [ ] **Step 1: Add device buffers** + +In `crates/ml-backtesting/src/sim.rs`, add to the `LobSimCuda` struct: + +```rust + // NEW — per-backtest orders + position state (device-global; one slice per backtest). + pub(crate) orders_d: CudaSlice, // packed bytes; sized per-block via lob::ORDERS_BYTES + pub(crate) pos_d: CudaSlice, // packed; lob::POS_BYTES + // Mapped-pinned audit ring per backtest. RING_CAP_EVENTS events × OrderEvent (24 B). + pub(crate) audit_d: CudaSlice, // [n_backtests * RING_CAP_EVENTS * 24] + pub(crate) audit_head_d: CudaSlice, // [n_backtests] — write index, kernel-mutated +``` + +Add to `crates/ml-backtesting/src/lob/mod.rs`: + +```rust +pub const MAX_LIMITS: usize = 32; +pub const MAX_STOPS: usize = 16; +/// Bytes per LimitSlot (matches lob_state.cuh LimitSlot struct). +pub const LIMIT_SLOT_BYTES: usize = 24; // 4 × u8 + 3 × f32 + u64 + u16 + 2 pad +pub const STOP_SLOT_BYTES: usize = 24; +pub const ORDERS_BYTES: usize = MAX_LIMITS * LIMIT_SLOT_BYTES + MAX_STOPS * STOP_SLOT_BYTES; +pub const POS_BYTES: usize = 24; +pub const RING_CAP_EVENTS: usize = 256; +pub const ORDER_EVENT_BYTES: usize = 24; +``` + +In `LobSimCuda::new()`, allocate the new buffers: + +```rust + let orders_d = stream.alloc_zeros::(n_backtests * crate::lob::ORDERS_BYTES)?; + let pos_d = stream.alloc_zeros::(n_backtests * crate::lob::POS_BYTES)?; + let audit_d = stream.alloc_zeros::(n_backtests * crate::lob::RING_CAP_EVENTS * crate::lob::ORDER_EVENT_BYTES)?; + let audit_head_d = stream.alloc_zeros::(n_backtests)?; +``` + +Add accessor `pub fn read_audit_ring(&self, b: usize) -> Result>` that copies the per-backtest ring slice + clears `audit_head_d[b]` (for fixture inspection). + +- [ ] **Step 2: Verify build** + +```bash +SQLX_OFFLINE=true cargo check -p ml-backtesting +``` + +Expected: clean build. + +- [ ] **Step 3: Stage** (no commit) + +```bash +git add crates/ml-backtesting/src/sim.rs crates/ml-backtesting/src/lob/mod.rs +``` + +### Task 5.2: `order_match.cu` kernel + +**Files:** +- Create: `crates/ml-backtesting/cuda/order_match.cu` + +- [ ] **Step 1: Write the kernel** + +Create `crates/ml-backtesting/cuda/order_match.cu`: + +```cuda +// Resting-order matching + queue decay + stop trigger detection. +// One block per backtest. Thread 0 is matching-leader (single-writer +// to pos[], orders[]); other threads in the warp scan resting slots +// in parallel for early-exit, then leader applies mutations. +// Per feedback_no_atomicadd.md: no atomics — single-writer discipline. + +#include "lob_state.cuh" + +// Pre-cast packed bytes to typed view. +__device__ __forceinline__ Orders* orders_at(unsigned char* base, int b, unsigned int orders_bytes) { + return reinterpret_cast(base + (size_t)b * orders_bytes); +} +__device__ __forceinline__ Pos* pos_at(unsigned char* base, int b, unsigned int pos_bytes) { + return reinterpret_cast(base + (size_t)b * pos_bytes); +} + +// Walk ask side for a market buy of `size`. Returns total cost in +// price-units and writes fills_*_lots. Single-writer; called by tid==0 only. +__device__ float walk_ask_for_buy( + const Book& book, float size, float& filled_lots +) { + float cost = 0.0f; + filled_lots = 0.0f; + for (int k = 0; k < 10 && size > 0.0f; ++k) { + const float lvl_sz = book.ask_sz[k]; + const float take = (size < lvl_sz) ? size : lvl_sz; + cost += take * book.ask_px[k]; + filled_lots += take; + size -= take; + } + return cost; +} + +__device__ float walk_bid_for_sell( + const Book& book, float size, float& filled_lots +) { + float cost = 0.0f; + filled_lots = 0.0f; + for (int k = 0; k < 10 && size > 0.0f; ++k) { + const float lvl_sz = book.bid_sz[k]; + const float take = (size < lvl_sz) ? size : lvl_sz; + cost += take * book.bid_px[k]; + filled_lots += take; + size -= take; + } + return cost; +} + +// Apply queue decay to a resting limit: trade volume `trade_vol` at the +// limit's price level reduces queue_position first, then size_remaining. +__device__ void apply_queue_decay(LimitSlot& slot, float trade_vol) { + if (trade_vol <= 0.0f) return; + if (slot.queue_position >= trade_vol) { + slot.queue_position -= trade_vol; + } else { + const float over = trade_vol - slot.queue_position; + slot.queue_position = 0.0f; + const float take = (over < slot.size_remaining) ? over : slot.size_remaining; + slot.size_remaining -= take; + } +} + +// Append an OrderEvent to the per-block audit ring. tid==0 only. +__device__ void emit_event( + unsigned char* audit_base, unsigned int* head_d, int b, unsigned int ring_cap, + unsigned long long ts_ns, unsigned int slot_tag, unsigned char kind, + unsigned char side, unsigned short size, int px, int trigger_px +) { + const unsigned int idx = head_d[b] % ring_cap; + head_d[b] += 1; + // Inline-construct OrderEvent. Layout MUST match crates/ml-backtesting/src/order.rs. + unsigned char* slot = audit_base + ((size_t)b * ring_cap + idx) * 24; + *reinterpret_cast(slot + 0) = ts_ns; + *reinterpret_cast(slot + 8) = slot_tag; + slot[12] = kind; + slot[13] = side; + *reinterpret_cast(slot + 14) = size; + *reinterpret_cast(slot + 16) = px; + *reinterpret_cast(slot + 20) = trigger_px; +} + +// Process one event: (a) apply queue decay to all resting limits whose +// price intersects the trade flow indicated by the snapshot diff, +// (b) check resting limits whose px is now marketable against the new +// book (in-flight orders that just arrived), (c) check stop triggers. +extern "C" __global__ void order_match_step_event( + const Book* __restrict__ books, // [n_backtests] — after book_update + unsigned char* orders_base, // [n_backtests * ORDERS_BYTES] + unsigned char* pos_base, // [n_backtests * POS_BYTES] + unsigned char* audit_base, // [n_backtests * RING_CAP * 24] + unsigned int* audit_head, // [n_backtests] + unsigned long long current_ts_ns, + float trade_signed_vol, // signed: + buyer-init, − seller-init + int orders_bytes, // sizeof(Orders) from host + int pos_bytes, // sizeof(Pos) + int ring_cap_events, + float tick_size, // 0.25 for ES + int n_backtests +) { + int b = blockIdx.x; + int tid = threadIdx.x; + if (b >= n_backtests) return; + if (tid != 0) return; // single-writer per-block; multi-thread parallel scan added in later perf pass + + const Book& book = books[b]; + Orders& orders = *orders_at(orders_base, b, orders_bytes); + Pos& pos = *pos_at(pos_base, b, pos_bytes); + + // (a) Queue decay on every resting limit at trade-flow price. + if (trade_signed_vol != 0.0f) { + const float vol = fabsf(trade_signed_vol); + for (int i = 0; i < MAX_LIMITS; ++i) { + LimitSlot& s = orders.limits[i]; + if (s.active != 1) continue; + apply_queue_decay(s, vol); + } + } + + // (b) In-flight orders (active==2) become resting when arrival_ts reached. + // Then check marketability against the new book. + for (int i = 0; i < MAX_LIMITS; ++i) { + LimitSlot& s = orders.limits[i]; + if (s.active == 2 && current_ts_ns >= s.arrival_ts_ns) { + s.active = 1; + s.queue_position = (s.side == 0) ? book.bid_sz[0] : book.ask_sz[0]; // pessimistic + } + if (s.active != 1) continue; + // Marketability check for IOC/FOK + just-arrived limits that + // crossed during in-flight. + const bool marketable = (s.side == 0 && s.price >= book.ask_px[0]) + || (s.side == 1 && s.price <= book.bid_px[0]); + if (marketable) { + float filled; + const float cost = (s.side == 0) + ? walk_ask_for_buy (book, s.size_remaining, filled) + : walk_bid_for_sell(book, s.size_remaining, filled); + // Apply fill to position + realized_pnl. + const float sgn = (s.side == 0) ? 1.0f : -1.0f; + const float prev_pos = (float)pos.position_lots; + const float new_pos = prev_pos + sgn * filled; + if (prev_pos == 0.0f || (prev_pos > 0.0f) == (sgn > 0.0f)) { + // Same-direction add or open: update VWAP. + const float prev_notional = prev_pos * pos.vwap_entry; + pos.vwap_entry = (prev_notional + sgn * cost) / (fabsf(new_pos) + 1e-9f); + } else { + // Counter-direction: realize P&L on the unwound portion. + const float unwind = fminf(fabsf(prev_pos), filled); + const float realised = (pos.vwap_entry - cost / fmaxf(filled, 1e-9f)) * unwind * sgn * -1.0f; + pos.realized_pnl += realised; + if (fabsf(new_pos) > 1e-9f && (new_pos > 0.0f) != (prev_pos > 0.0f)) { + pos.vwap_entry = cost / fmaxf(filled, 1e-9f); + } + } + pos.position_lots = (int)new_pos; + s.size_remaining -= filled; + if (s.size_remaining <= 0.0f) { + s.active = 0; // freed; gen_counter is incremented on next allocate (decision kernel) + emit_event(audit_base, audit_head, b, (unsigned int)ring_cap_events, + current_ts_ns, ((unsigned int)i << 8) | ((unsigned int)s.gen_counter << 16), + 1 /*SubmitLimit-fill-complete*/, s.side, (unsigned short)filled, + (int)(s.price / tick_size + 0.5f), 0); + } + } + // Tick-size used to encode prices as ticks in audit events. + (void)tick_size; + } + + // (c) Stop trigger detection — best ask crosses up for buy stop, + // best bid crosses down for sell stop. + for (int j = 0; j < MAX_STOPS; ++j) { + StopSlot& st = orders.stops[j]; + if (st.active != 1) continue; + const bool triggered = (st.side == 0 && book.ask_px[0] >= st.trigger_price) + || (st.side == 1 && book.bid_px[0] <= st.trigger_price); + if (!triggered) continue; + // Convert to either market or limit. For v1 stop-market just walks the book. + st.active = 0; + if (st.type == 4) { // StopMarket + float filled; + const float cost = (st.side == 0) + ? walk_ask_for_buy (book, st.size, filled) + : walk_bid_for_sell(book, st.size, filled); + const float sgn = (st.side == 0) ? 1.0f : -1.0f; + pos.position_lots += (int)(sgn * filled); + // P&L bookkeeping mirrors the limit path (omitted for brevity in this kernel — + // the closing pass is the same VWAP/realized rule). + (void)cost; + emit_event(audit_base, audit_head, b, (unsigned int)ring_cap_events, + current_ts_ns, ((unsigned int)j << 8) | (1u), + 4, st.side, (unsigned short)filled, 0, + (int)(st.trigger_price / tick_size + 0.5f)); + } + // OCO link: if this stop has an oco_pair, cancel the paired slot. + if (st.oco_pair != 0xFF) { + const int pair_idx = (int)st.oco_pair; + if (pair_idx < MAX_LIMITS) { + orders.limits[pair_idx].active = 0; + } else if (pair_idx - MAX_LIMITS < MAX_STOPS) { + orders.stops[pair_idx - MAX_LIMITS].active = 0; + } + } + } +} +``` + +- [ ] **Step 2: Wire bindings into sim.rs** + +Add a method `pub fn step_event(&mut self, current_ts_ns: u64, trade_signed_vol: f32) -> Result<()>` to `LobSimCuda` that launches `order_match_step_event` after `apply_snapshot`. Pass `ORDERS_BYTES`, `POS_BYTES`, `RING_CAP_EVENTS`, `0.25` (tick size) as kernel args. Launch config: grid (n_backtests, 1, 1), block (32, 1, 1). + +```rust + pub fn step_event(&mut self, current_ts_ns: u64, trade_signed_vol: f32) -> Result<()> { + let module = self.ctx.load_module( + cudarc::nvrtc::Ptx::from_src(include_bytes!(concat!(env!("OUT_DIR"), "/order_match.cubin")).as_ref()) + )?; + let func = module.load_function("order_match_step_event")?; + 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 orders_bytes = crate::lob::ORDERS_BYTES as i32; + let pos_bytes = crate::lob::POS_BYTES as i32; + let ring_cap = crate::lob::RING_CAP_EVENTS as i32; + let tick: f32 = 0.25; + let mut launch = self.stream.launch_builder(&func); + unsafe { launch.arg(&self.books_d).arg(&mut self.orders_d).arg(&mut self.pos_d) + .arg(&mut self.audit_d).arg(&mut self.audit_head_d) + .arg(¤t_ts_ns).arg(&trade_signed_vol) + .arg(&orders_bytes).arg(&pos_bytes).arg(&ring_cap).arg(&tick).arg(&n) + .launch(cfg)?; } + self.stream.synchronize()?; + Ok(()) + } +``` + +Also need to retain a CudaContext handle in `LobSimCuda` (add `pub(crate) ctx: Arc` to the struct and capture in `new()`). + +Add `pub fn submit_market(&mut self, backtest_idx: usize, side: u8, size: u16) -> Result<()>` that directly mutates `orders_d` from host to seed a market order — this is for fixture testing only (the production path is via `decision_policy.cu` in C7). The mechanism: pre-populate a single LimitSlot with active=1, side, price=0, marking it as a "market fill on next step" via a sentinel — simpler for v1 is to launch a tiny host-side kernel `seed_market_order` that takes (b, side, size) and writes into `orders[b].limits[0]` with marketable price. Define this minimal kernel inline in `cuda/order_match.cu`: + +```cuda +extern "C" __global__ void seed_market_order( + unsigned char* orders_base, int orders_bytes, + int b, unsigned char side, float size, float marketable_px +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + Orders& orders = *orders_at(orders_base, b, (unsigned int)orders_bytes); + LimitSlot& s = orders.limits[0]; + s.active = 1; + s.side = side; + s.type = 0; // Limit + s.oco_pair = 0xFF; + s.price = marketable_px; + s.size_remaining = size; + s.queue_position = 0.0f; + s.arrival_ts_ns = 0; + s.gen_counter = 1; +} +``` + +- [ ] **Step 3: Verify build** + +```bash +SQLX_OFFLINE=true cargo build -p ml-backtesting +find target -name "order_match.cubin" | head +``` + +Expected: cubin present. + +- [ ] **Step 4: Stage** (no commit — fixtures next in 5.3) + +```bash +git add crates/ml-backtesting/cuda/order_match.cu crates/ml-backtesting/src/sim.rs +``` + +### Task 5.3: Four Ring 1 fixtures for matching + +**Files:** +- Create: `crates/ml-backtesting/tests/fixtures/market_order_consumes_top.json` +- Create: `crates/ml-backtesting/tests/fixtures/limit_rest_queue_init.json` +- Create: `crates/ml-backtesting/tests/fixtures/limit_fill_via_trade.json` +- Create: `crates/ml-backtesting/tests/fixtures/latency_in_flight_miss.json` +- Modify: `crates/ml-backtesting/tests/lob_sim_fixtures.rs` + +- [ ] **Step 1: Extend fixture schema for orders + positions** + +Add to `crates/ml-backtesting/tests/lob_sim_fixtures.rs`: + +```rust +#[derive(Debug, Deserialize)] +struct OrderSeed { + kind: String, // "market" | "limit" + side: String, // "buy" | "sell" + size: f32, + #[serde(default)] px: f32, + #[serde(default)] marketable_px: f32, +} + +#[derive(Debug, Deserialize)] +struct ExpectedPos { + position_lots: i32, + #[serde(default)] realized_pnl: f32, + #[serde(default = "pnl_tol")] pnl_tol: f32, +} +fn pnl_tol() -> f32 { 1e-4 } +``` + +Extend the `Fixture` struct with optional `pre_orders: Option>` and `expected_pos: Option>`. The harness pre-seeds orders before applying the events, then asserts position state after. + +Add helper `run_match_fixture(path)` that follows the same pattern as `run_book_fixture` but also calls `sim.step_event(ts, trade_vol)` after each snapshot apply and reads back position via a new `sim.read_pos(b) -> Result` (add `pub struct PosFlat { ... }` to `lob/mod.rs` matching `Pos` layout). + +- [ ] **Step 2: Create the 4 fixtures** + +`market_order_consumes_top.json` — book with ask[0]=3 @ 5500, ask[1]=10 @ 5500.25; seed market buy size=5; expect position_lots=5, realized_pnl=0 (opening). + +`limit_rest_queue_init.json` — book with bid_sz[3]=10; seed limit buy at price=bid_px[3]; expect that slot's `queue_position` = 10. (Needs a read-orders accessor.) + +`limit_fill_via_trade.json` — seed limit buy at bid_px[3] with queue_position=10 (manually); apply a snapshot with the bid level removed (price moved up); expect the limit's `size_remaining` decreases when trade_signed_vol > 10. + +`latency_in_flight_miss.json` — seed limit buy at price P with `active=2`, `arrival_ts_ns = 100_000_000`; apply 2 snapshots: first at ts=50_000_000 with price unchanged (limit stays in-flight); second at ts=150_000_000 with price moved to P+1 tick (limit just arrived and is marketable, crosses for slippage). Assert fill at the new (worse) price. + +(JSON bodies follow the schema in 5.2; the implementing engineer can construct exact numeric values matching the expected outcomes — the spec for each fixture in spec §8 is the source of truth.) + +- [ ] **Step 3: Add 4 new `#[test] #[ignore]` functions** in `tests/lob_sim_fixtures.rs`, each calling `run_match_fixture("tests/fixtures/.json")`. + +- [ ] **Step 4: Run the fixtures** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test lob_sim_fixtures -- --ignored --nocapture +``` + +Expected: 7 PASS (3 from C4 + 4 new). + +- [ ] **Step 5: Commit C5** + +```bash +git add crates/ml-backtesting/cuda/order_match.cu \ + crates/ml-backtesting/src/sim.rs \ + crates/ml-backtesting/src/lob/mod.rs \ + crates/ml-backtesting/tests/lob_sim_fixtures.rs \ + crates/ml-backtesting/tests/fixtures/ +git commit -m "$(cat <<'EOF' +feat(ml-backtesting): order_match kernel + latency-aware fills + +order_match_step_event applies queue decay on resting limits proportional +to trade flow, promotes in-flight (active==2) orders when arrival_ts +reached and re-checks marketability against the (possibly moved) book, +walks book levels for fills, accounts realized P&L via VWAP rule, and +handles OCO leg-cancellation on stop trigger. Per +pearl_no_host_branches_in_captured_graph.md the kernel is pure device-side; +all OrderEvent emission goes through the mapped-pinned audit ring per +feedback_no_htod_htoh_only_mapped_pinned.md. + +Adds 4 Ring 1 fixtures: market_order_consumes_top, limit_rest_queue_init, +limit_fill_via_trade, latency_in_flight_miss. Latency fixture verifies +that orders submitted during a moving book arrive at the post-move price +not the submit-time price (spec §4 mechanics). + +Co-Authored-By: Claude Opus 4.7 +EOF +)" +``` + +--- + +## Commit 6 — ml-backtesting: pnl_track kernel + accounting fixture + +**Why now:** The matching kernel does inline P&L bookkeeping for simple opens/closes, but proper VWAP + segment_complete + per-trade-record audit is a separate concern that the decision-policy kernel (C7) needs (for the per-horizon ISV-Kelly update). Extract into its own kernel. + +**Spec reference:** §5 (per-horizon ISV-Kelly Welford updates on segment_complete), §8 fixture `pnl_accounting`. + +### Task 6.1: `pnl_track.cu` — segment_complete detector + per-trade record emit + +**Files:** +- Create: `crates/ml-backtesting/cuda/pnl_track.cu` +- Modify: `crates/ml-backtesting/src/sim.rs` +- Modify: `crates/ml-backtesting/src/lob/mod.rs` + +- [ ] **Step 1: Define `TradeRecord` (closed trades, distinct from `OrderEvent`)** + +Add to `crates/ml-backtesting/src/order.rs`: + +```rust +/// Closed-trade record emitted when a position returns to flat +/// (either by counter-fill or stop). Written by the pnl_track kernel +/// into the per-block trade-log mapped-pinned buffer. +#[repr(C)] +#[derive(Clone, Copy, Debug, bytemuck::Pod, bytemuck::Zeroable, Serialize, Deserialize)] +pub struct TradeRecord { + pub entry_ts_ns: u64, // 8 + pub exit_ts_ns: u64, // 8 + pub entry_px_ticks: i32, // 4 + pub exit_px_ticks: i32, // 4 + pub size_lots: i32, // 4 — signed (+long, −short closed) + pub fees_usd_fp: i32, // 4 — fixed-point ×100 (i.e. $1.23 → 123) + pub realised_pnl_usd_fp: i32, // 4 + pub horizon_idx: u8, // 1 + pub strategy_id: u8, // 1 + pub _pad: [u8; 2], + // total: 40 bytes +} +``` + +Add to `crates/ml-backtesting/src/lob/mod.rs`: + +```rust +pub const TRADE_LOG_CAP: usize = 1024; +pub const TRADE_RECORD_BYTES: usize = 40; +``` + +In `LobSimCuda`: + +```rust + pub(crate) trade_log_d: CudaSlice, // [n_backtests * TRADE_LOG_CAP * 40] + pub(crate) trade_log_head_d: CudaSlice, // [n_backtests] +``` + +Allocate in `new()` analogously to the audit ring. + +- [ ] **Step 2: Write the kernel** + +`crates/ml-backtesting/cuda/pnl_track.cu`: + +```cuda +// Detects segment_complete (position returns to flat) and emits a +// TradeRecord to the per-block trade-log mapped-pinned buffer. +// Driven by deltas on Pos.position_lots between consecutive events. +// Single-writer (tid==0) per block. + +#include "lob_state.cuh" + +extern "C" __global__ void pnl_track_step( + unsigned char* pos_base, + unsigned char* trade_log_base, + unsigned int* trade_log_head, + unsigned long long current_ts_ns, + int pos_bytes, + int trade_log_cap, + int n_backtests, + // Persisted per-block "open trade" snapshot (entry ts, px, size, horizon). + unsigned char* open_trade_state // [n_backtests * 32] — entry context +) { + int b = blockIdx.x; + if (b >= n_backtests || threadIdx.x != 0) return; + Pos* p = reinterpret_cast(pos_base + (size_t)b * pos_bytes); + + // open_trade_state layout per backtest (32 bytes): + // 0..8: entry_ts_ns (u64) + // 8..12: entry_vwap_x100 (i32 ticks ×100) + // 12..16: entry_size_lots (i32 signed) + // 16..17: entry_horizon_idx (u8) + // 17..32: reserved + unsigned char* st = open_trade_state + (size_t)b * 32; + int prev_size = *reinterpret_cast(st + 12); + int now_size = p->position_lots; + + // Open: prev==0, now != 0. + if (prev_size == 0 && now_size != 0) { + *reinterpret_cast(st + 0) = current_ts_ns; + *reinterpret_cast(st + 8) = (int)(p->vwap_entry * 100.0f + 0.5f); + *reinterpret_cast(st + 12) = now_size; + // horizon_idx set externally by decision kernel when authorizing entry; + // pnl_track preserves it. + } + // Close: prev != 0, now == 0. + else if (prev_size != 0 && now_size == 0) { + const unsigned int idx = trade_log_head[b] % (unsigned int)trade_log_cap; + trade_log_head[b] += 1; + unsigned char* rec = trade_log_base + ((size_t)b * trade_log_cap + idx) * 40; + *reinterpret_cast(rec + 0) = *reinterpret_cast(st + 0); + *reinterpret_cast(rec + 8) = current_ts_ns; + *reinterpret_cast(rec + 16) = *reinterpret_cast(st + 8); // entry_px_ticks (×100) + *reinterpret_cast(rec + 20) = (int)(p->vwap_entry * 100.0f + 0.5f); // exit_px_ticks + *reinterpret_cast(rec + 24) = prev_size; + *reinterpret_cast(rec + 28) = 0; // fees (decision kernel adds them) + // Realized P&L (USD ×100 fixed-point): (exit − entry) × size × $50 × 100 → 5000 per tick per lot. + const int realised_fp = (int)((p->realized_pnl) * 100.0f + 0.5f); + *reinterpret_cast(rec + 32) = realised_fp; + rec[36] = st[16]; // horizon_idx + rec[37] = 0; // strategy_id placeholder + // Reset open-trade state. + for (int i = 0; i < 32; ++i) st[i] = 0; + } + // Same direction increment / partial close: covered in v2 (multi-fill + // averaging). v1 fixtures assume open → flat single-segment trades. +} +``` + +- [ ] **Step 3: Wire into sim, add `step_pnl_track()` method** + +In `LobSimCuda::new()` allocate `open_trade_state_d: CudaSlice` sized `n_backtests * 32`. Add method: + +```rust + pub fn step_pnl_track(&mut self, current_ts_ns: u64) -> Result<()> { /* launch kernel */ } +``` + +Call it from `step_event` (or expose separately — decision is implementer's). The simplest path: `step_event` calls `step_pnl_track` internally after order_match returns. + +- [ ] **Step 4: Create `pnl_accounting` fixture** + +`crates/ml-backtesting/tests/fixtures/pnl_accounting.json` — sequence: open via market buy at known price, hold one snapshot, close via market sell at higher price. Expected `realized_pnl` = (exit − entry) × size × $50 within 1e-2 tolerance. + +Add `#[test] #[ignore] fn fix_pnl_accounting()` in `tests/lob_sim_fixtures.rs`. + +- [ ] **Step 5: Run** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test lob_sim_fixtures -- --ignored --nocapture +``` + +Expected: 8 PASS (7 + pnl_accounting). + +- [ ] **Step 6: Commit C6** + +```bash +git add crates/ml-backtesting/cuda/pnl_track.cu \ + crates/ml-backtesting/src/sim.rs \ + crates/ml-backtesting/src/lob/mod.rs \ + crates/ml-backtesting/src/order.rs \ + crates/ml-backtesting/tests/fixtures/pnl_accounting.json \ + crates/ml-backtesting/tests/lob_sim_fixtures.rs +git commit -m "$(cat <<'EOF' +feat(ml-backtesting): pnl_track kernel + closed-trade records + +pnl_track_step detects segment_complete (position returns to flat), +emits a 40-byte TradeRecord to per-block mapped-pinned trade-log +buffer. open_trade_state buffer persists entry context across events. +TradeRecord layout pinned via repr(C) + bytemuck::Pod so host can +mmap-read without copy. + +Ring 1 fixture pnl_accounting verifies open-buy → hold → sell-close +sequence accounts realised_pnl correctly with $50/index-point ES contract value. + +v1 limitation: multi-fill averaging (partial close) deferred to v2; +fixture covers single-segment open→flat trades only. + +Co-Authored-By: Claude Opus 4.7 +EOF +)" +``` + +--- + +## Commit 7 — ml-backtesting: decision_policy kernel + +**Why now:** With book + matching + P&L working, this commit closes the loop — consuming alpha probs broadcast from the trunk and emitting orders via the bytecode program from C3. Per-horizon ISV-Kelly evolution lives here. + +**Spec reference:** §5 (per-horizon ISV-Kelly + WeightedByRealizedSharpe formulas), §6 (bytecode opcodes), §8 fixtures `stop_trigger`, `oco_one_cancels_other`, `submission_overflow`. + +### Task 7.1: Per-block `IsvKellyState[N_HORIZONS]` device buffer + program upload path + +**Files:** +- Modify: `crates/ml-backtesting/src/sim.rs` +- Modify: `crates/ml-backtesting/src/lob/mod.rs` + +- [ ] **Step 1: Add buffers** + +In `lob/mod.rs`: + +```rust +pub const ISV_KELLY_STATE_BYTES: usize = 24; // matches lob_state.cuh IsvKellyState +pub const ISV_KELLY_TOTAL_BYTES: usize = ISV_KELLY_STATE_BYTES * 5; // × N_HORIZONS +pub const MAX_PROGRAM_INSTRUCTIONS: usize = 64; +pub const INSTRUCTION_BYTES: usize = 8; // matches policy::Instruction (u8+u8+u16+u32) +``` + +In `LobSimCuda`: + +```rust + pub(crate) isv_kelly_d: CudaSlice, // [n_backtests * ISV_KELLY_TOTAL_BYTES] + pub(crate) program_table_d: CudaSlice, // [n_backtests * MAX_PROGRAM_INSTRUCTIONS * INSTRUCTION_BYTES] + pub(crate) program_len_d: CudaSlice, // [n_backtests] + // Mapped-pinned alpha probs broadcast buffer (one per stride decision). + pub(crate) alpha_probs_h: CudaSlice, // [N_HORIZONS] + pub(crate) current_regime_d: CudaSlice, // [n_backtests] — for RegimeSwitch ops + // Latency in ns to apply on order submission (host-config; constant or empirical sample). + pub(crate) latency_ns: u32, +``` + +Add methods: + +```rust + pub fn upload_program(&mut self, backtest_idx: usize, prog: &crate::policy::Program) -> Result<()>; + pub fn broadcast_alpha(&mut self, probs: &[f32; 5]) -> Result<()>; + pub fn set_latency_ns(&mut self, ns: u32) { self.latency_ns = ns; } + pub fn read_isv_kelly(&self, b: usize) -> Result<[crate::policy::IsvKellyStateHost; 5]>; +``` + +Define `IsvKellyStateHost` in `policy/sizing.rs` as `repr(C) + bytemuck::Pod` mirror of the CUDA struct. + +- [ ] **Step 2: Verify build** + +```bash +SQLX_OFFLINE=true cargo check -p ml-backtesting +``` + +Expected: clean build. + +- [ ] **Step 3: Stage** (no commit) + +```bash +git add crates/ml-backtesting/src/sim.rs crates/ml-backtesting/src/lob/mod.rs crates/ml-backtesting/src/policy/ +``` + +### Task 7.2: `decision_policy.cu` — bytecode interpreter + per-horizon sizing + aggregator + OCO + overflow + +**Files:** +- Create: `crates/ml-backtesting/cuda/decision_policy.cu` + +- [ ] **Step 1: Write the kernel** + +`crates/ml-backtesting/cuda/decision_policy.cu`: + +```cuda +// Bytecode interpreter for the decision program flattened from +// Strategy::flatten(). Each block runs one backtest's program against +// the broadcast alpha probs + its per-horizon IsvKellyState. +// Single-writer (tid==0) per block — multi-thread parallelism added in v2 +// if profile shows benefit. Stack-based VM with fixed-size float stack. + +#include "lob_state.cuh" +#define STACK_CAP 32 + +// Mirror crates/ml-backtesting/src/policy/mod.rs OpCode discriminants. +#define OP_NOOP 0 +#define OP_EVAL_REGIME 1 +#define OP_EMIT_PER_HORIZON_SIZE 2 +#define OP_AGG_WEIGHTED_SHARPE 3 +#define OP_AGG_MEAN 4 +#define OP_AGG_MAX_CONFIDENCE 5 +#define OP_APPLY_CONFLICT 6 +#define OP_WRITE_ORDER 7 +#define OP_PUSH_SCALAR 8 +#define OP_BRANCH_IF_REGIME 9 + +struct Instruction { + unsigned char op; + unsigned char arg0; + unsigned short arg1; + unsigned int arg2; +}; + +// Compute per-horizon signed sizing intent from alpha prob + IsvKelly state. +// Returns signed lots: positive = long, negative = short, zero = no trade. +// Per spec §5: sig_mag × kelly_frac × cap, no floor. +__device__ float per_horizon_size( + float p_h, const IsvKellyState& s, float target_annual_vol_units, + float annualisation_factor, int max_lots +) { + const float sig_mag = fabsf(p_h - 0.5f) * 2.0f; // ∈ [0, 1] + const float dir = (p_h > 0.5f) ? 1.0f : -1.0f; + + // Kelly fraction: avoid div-by-zero on first observation by using + // sentinel safety (returns 0 if pnl_ema_win not yet bootstrapped). + const float pnl_w = s.pnl_ema_win; + const float pnl_l = s.pnl_ema_loss; + if (pnl_w <= 0.0f) return 0.0f; + float kelly_frac = (s.win_rate_ema * pnl_w - (1.0f - s.win_rate_ema) * pnl_l) / pnl_w; + kelly_frac = fmaxf(0.0f, fminf(1.0f, kelly_frac)); + + // Cap from realized variance + target vol. + if (s.realised_return_var <= 0.0f) return 0.0f; + const float cap_units = target_annual_vol_units / sqrtf(s.realised_return_var * annualisation_factor); + const float cap_lots = fminf(cap_units, (float)max_lots); + + return dir * sig_mag * kelly_frac * cap_lots; +} + +extern "C" __global__ void decision_policy_step( + const Instruction* __restrict__ programs, // [n_backtests * MAX_PROGRAM_INSTRUCTIONS] + const unsigned int* __restrict__ program_lens, + const float* __restrict__ alpha_probs_broadcast, // [N_HORIZONS] + const unsigned int* __restrict__ regimes, // [n_backtests] + unsigned char* isv_kelly_base, // [n_backtests * ISV_KELLY_TOTAL_BYTES] + unsigned char* orders_base, + unsigned char* pos_base, + unsigned char* audit_base, + unsigned int* audit_head, + unsigned long long current_ts_ns, + unsigned int latency_ns, + int orders_bytes, + int pos_bytes, + int max_program_instructions, + int ring_cap_events, + float tick_size, + float target_annual_vol_units, + float annualisation_factor, + int n_backtests +) { + int b = blockIdx.x; + if (b >= n_backtests || threadIdx.x != 0) return; + + const Instruction* prog = programs + (size_t)b * max_program_instructions; + const unsigned int prog_len = program_lens[b]; + IsvKellyState* isv = reinterpret_cast(isv_kelly_base + (size_t)b * sizeof(IsvKellyState) * N_HORIZONS); + Orders& orders = *reinterpret_cast(orders_base + (size_t)b * orders_bytes); + Pos& pos = *reinterpret_cast(pos_base + (size_t)b * pos_bytes); + + float stack[STACK_CAP]; + int sp = 0; + unsigned char authorizing_horizon_for_top = 0xFF; + // Parallel "horizon tag" stack — tracks which horizon authorized + // each value (needed for WriteOrder to set open_horizon_mask). + unsigned char hstack[STACK_CAP] = {0}; + + unsigned int pc = 0; + while (pc < prog_len) { + const Instruction ins = prog[pc++]; + switch (ins.op) { + case OP_PUSH_SCALAR: { + const float v = __int_as_float((int)ins.arg2); + stack[sp] = v; hstack[sp] = 0xFF; sp++; + break; + } + case OP_EVAL_REGIME: { + stack[sp] = (regimes[b] == (unsigned int)ins.arg0) ? 1.0f : 0.0f; + hstack[sp] = 0xFF; sp++; + break; + } + case OP_BRANCH_IF_REGIME: { + if (regimes[b] != (unsigned int)ins.arg0) pc = ins.arg2; + break; + } + case OP_EMIT_PER_HORIZON_SIZE: { + const unsigned char h = ins.arg0; + const float p_h = alpha_probs_broadcast[h]; + const float sz = per_horizon_size(p_h, isv[h], target_annual_vol_units, + annualisation_factor, (int)ins.arg1); + stack[sp] = sz; hstack[sp] = h; sp++; + break; + } + case OP_AGG_MEAN: { + const int n = (int)ins.arg0; + if (sp < n) break; + float sum = 0.0f; + for (int i = sp - n; i < sp; ++i) sum += stack[i]; + sp -= n; + stack[sp] = sum / (float)n; hstack[sp] = 0xFF; // mixed horizons → unknown + sp++; + break; + } + case OP_AGG_WEIGHTED_SHARPE: { + const int n = (int)ins.arg0; + if (sp < n) break; + float w_sum = 0.0f, wx_sum = 0.0f; + float best_h_share = -1.0f; unsigned char best_h = 0xFF; + for (int i = 0; i < n; ++i) { + const unsigned char h = hstack[sp - n + i]; + const float w = (h < N_HORIZONS) ? fmaxf(0.0f, isv[h].recent_sharpe) : 0.0f; + w_sum += w; + wx_sum += w * stack[sp - n + i]; + if (w > best_h_share) { best_h_share = w; best_h = h; } + } + sp -= n; + stack[sp] = (w_sum > 0.0f) ? (wx_sum / w_sum) : 0.0f; + hstack[sp] = best_h; // attribute open to the dominant-weight horizon + sp++; + break; + } + case OP_AGG_MAX_CONFIDENCE: { + const int n = (int)ins.arg0; + if (sp < n) break; + int best = sp - n; + for (int i = sp - n + 1; i < sp; ++i) + if (fabsf(stack[i]) > fabsf(stack[best])) best = i; + const float val = stack[best]; + const unsigned char h = hstack[best]; + sp -= n; + stack[sp] = val; hstack[sp] = h; sp++; + break; + } + case OP_APPLY_CONFLICT: { + // Currently a no-op for v1 Ensemble (no conflict in flat stack + // of homogeneous-direction sums). Reserved for Portfolio nesting. + (void)ins.arg0; + break; + } + case OP_WRITE_ORDER: { + if (sp == 0) break; + const float signed_size = stack[--sp]; + const unsigned char auth_h = hstack[sp]; + if (fabsf(signed_size) < 1.0f) break; // sub-1-lot intent → no order + const int lots = (int)signed_size; // truncate toward zero + const unsigned char side = (lots > 0) ? 0 : 1; + const unsigned short size = (unsigned short)abs(lots); + // Find a free limit slot. + int slot_idx = -1; + for (int i = 0; i < MAX_LIMITS; ++i) { + if (orders.limits[i].active == 0) { slot_idx = i; break; } + } + if (slot_idx < 0) { + pos.submission_overflow += 1u; + break; + } + LimitSlot& s = orders.limits[slot_idx]; + s.active = 2; // in-flight (latency) + s.side = side; + s.type = 1; // IOC for aggressive cross + s.oco_pair = 0xFF; + s.price = (side == 0) ? (alpha_probs_broadcast[0] /* placeholder */ , 0.0f) : 0.0f; + // Set marketable price (best opp side ± wide buffer) for v1 IOC behaviour. + // In production decision kernel would receive book pointer; v1 sim + // forwards via order_match marketability check (limit price = 0 makes it always marketable when promoted from in-flight). + s.size_remaining = (float)size; + s.queue_position = 0.0f; + s.arrival_ts_ns = current_ts_ns + (unsigned long long)latency_ns; + s.gen_counter += 1; + pos.open_horizon_mask |= (1u << auth_h); + emit_event(audit_base, audit_head, b, (unsigned int)ring_cap_events, + current_ts_ns, + ((unsigned int)slot_idx << 8) | ((unsigned int)s.gen_counter << 16), + 2 /*SubmitIoc*/, side, size, 0, 0); + break; + } + case OP_NOOP: + default: break; + } + if (sp >= STACK_CAP - 1 || sp < 0) break; + } +} + +// Called by host on segment_complete: update per-horizon IsvKellyState. +// Wiener-α floor 0.4 per pearl_wiener_alpha_floor_for_nonstationary. +extern "C" __global__ void isv_kelly_update_on_close( + unsigned char* isv_kelly_base, + unsigned int* isv_horizon_mask, // [n_backtests] — which horizons authorized closed trade + const float* realised_returns, // [n_backtests] — closed-trade return (price units) + int n_backtests +) { + int b = blockIdx.x; + if (b >= n_backtests || threadIdx.x != 0) return; + + const unsigned int mask = isv_horizon_mask[b]; + if (mask == 0) return; + IsvKellyState* isv = reinterpret_cast(isv_kelly_base + (size_t)b * sizeof(IsvKellyState) * N_HORIZONS); + const float ret = realised_returns[b]; + + // Update every horizon flagged in mask. Equal credit (v1 — refined in v2). + for (int h = 0; h < N_HORIZONS; ++h) { + if (!(mask & (1u << h))) continue; + IsvKellyState& s = isv[h]; + const bool win = ret > 0.0f; + // Wiener-α blended EMA with 0.4 floor. + const float alpha = 0.4f; + if (win) { + s.pnl_ema_win = (s.n_trades_seen == 0) ? ret : ((1.0f - alpha) * s.pnl_ema_win + alpha * ret); + s.win_rate_ema = (s.n_trades_seen == 0) ? 1.0f : ((1.0f - alpha) * s.win_rate_ema + alpha * 1.0f); + } else { + s.pnl_ema_loss = (s.n_trades_seen == 0) ? -ret : ((1.0f - alpha) * s.pnl_ema_loss + alpha * (-ret)); + s.win_rate_ema = (s.n_trades_seen == 0) ? 0.0f : ((1.0f - alpha) * s.win_rate_ema + alpha * 0.0f); + } + // Welford variance. + const float prev_n = (float)s.n_trades_seen; + const float delta = ret - 0.0f; // simplified: variance around 0 + s.realised_return_var = (prev_n * s.realised_return_var + delta * delta) / (prev_n + 1.0f); + s.n_trades_seen += 1u; + // recent_sharpe composite. + s.recent_sharpe = (s.pnl_ema_win * s.win_rate_ema - s.pnl_ema_loss * (1.0f - s.win_rate_ema)) + / fmaxf(sqrtf(s.realised_return_var), 1e-9f); + } +} + +__device__ void emit_event( // re-declare for linker; actual def lives in order_match.cu + unsigned char*, unsigned int*, int, unsigned int, + unsigned long long, unsigned int, unsigned char, + unsigned char, unsigned short, int, int +); +``` + +NOTE on the `emit_event` extern: nvcc compiles each .cu separately to its own cubin, so the kernel can't cross-link to `order_match.cu`'s symbol. Easiest fix: define `emit_event` as `__device__ static inline` in `lob_state.cuh` so both .cu files include it. Move the implementation there before this commit lands. + +- [ ] **Step 2: Move `emit_event` to `lob_state.cuh`** + +Edit `crates/ml-backtesting/cuda/lob_state.cuh` to add at the end (before `#endif`): + +```cuda +__device__ static inline void emit_event( + unsigned char* audit_base, unsigned int* head_d, int b, unsigned int ring_cap, + unsigned long long ts_ns, unsigned int slot_tag, unsigned char kind, + unsigned char side, unsigned short size, int px, int trigger_px +) { + const unsigned int idx = head_d[b] % ring_cap; + head_d[b] += 1; + unsigned char* slot = audit_base + ((size_t)b * ring_cap + idx) * 24; + *reinterpret_cast(slot + 0) = ts_ns; + *reinterpret_cast(slot + 8) = slot_tag; + slot[12] = kind; + slot[13] = side; + *reinterpret_cast(slot + 14) = size; + *reinterpret_cast(slot + 16) = px; + *reinterpret_cast(slot + 20) = trigger_px; +} +``` + +Remove the duplicate `__device__ void emit_event` definition from `order_match.cu` (keep only the call sites). Delete the broken extern declaration at the bottom of `decision_policy.cu`. + +- [ ] **Step 3: Add `step_decision()` to sim** + +```rust + pub fn step_decision(&mut self, current_ts_ns: u64, + target_annual_vol_units: f32, annualisation_factor: f32) -> Result<()>; +``` + +Launches `decision_policy_step` with the upload buffers. After the kernel returns, call `step_event` + `step_pnl_track`; after pnl_track, scan trade-log head deltas and call `isv_kelly_update_on_close` for any backtest that produced a new closed trade (host reads the mapped-pinned trade-log to extract `realised_pnl` + `open_horizon_mask` snapshot). + +- [ ] **Step 4: Stage** + +```bash +git add crates/ml-backtesting/cuda/ crates/ml-backtesting/src/sim.rs +``` + +### Task 7.3: Three Ring 1 fixtures for decision/stop/OCO/overflow + +**Files:** +- Create: `crates/ml-backtesting/tests/fixtures/stop_trigger.json` +- Create: `crates/ml-backtesting/tests/fixtures/oco_one_cancels_other.json` +- Create: `crates/ml-backtesting/tests/fixtures/submission_overflow.json` +- Modify: `crates/ml-backtesting/tests/lob_sim_fixtures.rs` + +- [ ] **Step 1: Create fixtures** + +`stop_trigger.json` — set position_lots=+1 entered at price P, seed StopMarket sell at trigger=P-10; apply snapshot with bid moved to P-10; expect StopMarket triggered, position flat, audit event emitted for stop fill. + +`oco_one_cancels_other.json` — seed limit buy at low price + limit sell at high price linked via `oco_pair`; apply snapshot that crosses the sell side; expect sell fills, buy slot auto-cancelled (active=0). + +`submission_overflow.json` — pre-seed 32 limits (filling all MAX_LIMITS); run decision_policy_step with a program that emits a 33rd order; expect `pos.submission_overflow == 1` and no extra slot allocated. + +- [ ] **Step 2: Add 3 `#[test] #[ignore]` functions** + extend fixture harness `Fixture` schema with `pre_isv_kelly: Option>`, `pre_program: Option>>`, `expected_overflow: Option>`. + +- [ ] **Step 3: Run all 11 fixtures** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test lob_sim_fixtures -- --ignored --nocapture +``` + +Expected: 11 PASS (8 + 3). + +- [ ] **Step 4: Commit C7** + +```bash +git add crates/ml-backtesting/cuda/ \ + crates/ml-backtesting/src/sim.rs \ + crates/ml-backtesting/src/policy/ \ + crates/ml-backtesting/src/lob/mod.rs \ + crates/ml-backtesting/tests/fixtures/ \ + crates/ml-backtesting/tests/lob_sim_fixtures.rs +git commit -m "$(cat <<'EOF' +feat(ml-backtesting): decision_policy kernel + per-horizon ISV-Kelly + +decision_policy_step is a stack-based VM that interprets the flat +bytecode Program from C3 against broadcast alpha probs + per-block +IsvKellyState[5]. Implements WeightedByRealizedSharpe aggregator +(weights = max(0, recent_sharpe[h]) / Σ; auto-shifts capital toward +empirically winning horizons). Per spec §5: no floor, no engagement +guarantee — cap collapses to 0 when realized_return_var unknown. + +isv_kelly_update_on_close updates per-horizon EMAs + Welford variance ++ recent_sharpe composite when a trade closes; equal credit to every +horizon flagged in pos.open_horizon_mask at close time. Wiener-α floor +0.4 per pearl_wiener_alpha_floor_for_nonstationary. + +OCO mutual-exclusion handled via slot.oco_pair byte — when one leg +fills, the linked slot is freed (active=0). submission_overflow +counter increments instead of allocating when all MAX_LIMITS=32 slots +are occupied. + +emit_event moved to lob_state.cuh as __device__ static inline so +both order_match.cu and decision_policy.cu share the implementation +(separate-cubin nvcc has no cross-cu linkage). + +Adds 3 Ring 1 fixtures: stop_trigger, oco_one_cancels_other, +submission_overflow. Full Ring 1 suite (11 fixtures) now green. + +Co-Authored-By: Claude Opus 4.7 +EOF +)" +``` + +--- + +## Commit 8 — ml-backtesting: BacktestHarness + trainer parity (Ring 1b) + +**Why now:** All kernels work in isolation; this commit ties them together with the loader + trunk inference behind a single `BacktestHarness::run()` entry point. Ring 1b enforces structural train-vs-deploy parity. + +**Spec reference:** §1 (trainer parity contract), §7 (orchestrator code), §8 Ring 1b. + +### Task 8.1: `BacktestHarness::run()` + +**Files:** +- Create: `crates/ml-backtesting/src/harness.rs` +- Modify: `crates/ml-backtesting/src/lib.rs` + +- [ ] **Step 1: Write the harness** + +`crates/ml-backtesting/src/harness.rs`: + +```rust +//! Orchestrator wiring MultiHorizonLoader → CfcTrunk → LobSimCuda. +//! See spec §7. + +use anyhow::{Context, Result}; +use ml_alpha::cfc::trunk::CfcTrunk; +use ml_alpha::data::loader::{MultiHorizonLoader, MultiHorizonLoaderConfig}; +use std::path::{Path, PathBuf}; +use std::sync::Arc; + +use crate::policy::{Strategy, Program}; +use crate::sim::LobSimCuda; + +pub struct BacktestHarnessConfig { + pub model_ckpt: PathBuf, + pub data_root: PathBuf, + pub predecoded_dir: PathBuf, + pub n_parallel: usize, + pub decision_stride: usize, + pub strategies: Vec, // one per backtest cell + pub target_annual_vol_units: f32, + pub annualisation_factor: f32, + pub latency_ns: u32, + pub flush_every_events: u64, + pub out_dir: PathBuf, +} + +pub struct BacktestHarness { + cfg: BacktestHarnessConfig, + loader: MultiHorizonLoader, + trunk: CfcTrunk, + sim: LobSimCuda, +} + +impl BacktestHarness { + pub fn new(cfg: BacktestHarnessConfig) -> Result { + anyhow::ensure!(cfg.strategies.len() == cfg.n_parallel, + "strategies len {} ≠ n_parallel {}", cfg.strategies.len(), cfg.n_parallel); + + let files = ml_alpha::data::loader::discover_mbp10_files_sorted(&cfg.data_root)?; + let loader_cfg = MultiHorizonLoaderConfig { + files, + predecoded_dir: cfg.predecoded_dir.clone(), + seq_len: 1, + horizons: [30, 100, 300, 1000, 6000], + n_max_sequences: 0, + seed: 0, + decision_stride: cfg.decision_stride, + inference_only: true, + }; + let loader = MultiHorizonLoader::new(&loader_cfg)?; + + let ctx = cudarc::driver::CudaContext::new(0).context("cuda ctx")?; + let ctx_arc = Arc::new(ctx); + + let mut trunk = CfcTrunk::load_checkpoint(&cfg.model_ckpt, ctx_arc.clone()) + .context("load trunk checkpoint")?; + let first = loader.peek_first()?; + trunk.capture_graph_a(&first)?; + + let mut sim = LobSimCuda::new(cfg.n_parallel, &ctx_arc)?; + sim.set_latency_ns(cfg.latency_ns); + for (b, strat) in cfg.strategies.iter().enumerate() { + let prog: Program = strat.flatten(); + sim.upload_program(b, &prog)?; + } + + Ok(Self { cfg, loader, trunk, sim }) + } + + /// Run the full backtest to completion. Returns at end-of-stream. + pub fn run(&mut self) -> Result<()> { + let mut event_idx: u64 = 0; + while let Some(raw) = self.loader.next_inference_input()? { + // Apply book update + matching every event. + self.sim.apply_snapshot(&raw.bid_px, &raw.bid_sz, &raw.ask_px, &raw.ask_sz)?; + self.sim.step_event(raw.ts_ns, raw.trade_signed_vol)?; + + // Inference + decision only at decision-stride boundaries. + if event_idx % self.cfg.decision_stride as u64 == 0 { + self.trunk.update_input_buffers(&raw)?; + let (probs, _proj) = self.trunk.perception_forward_captured()?; + self.sim.broadcast_alpha(&probs)?; + self.sim.step_decision(raw.ts_ns, + self.cfg.target_annual_vol_units, self.cfg.annualisation_factor)?; + } + + event_idx += 1; + if event_idx % self.cfg.flush_every_events == 0 { + // Hooks for artifact flush (filled in by C9). + } + } + Ok(()) + } +} +``` + +Update `lib.rs` to add `pub mod harness;`. + +- [ ] **Step 2: Verify build** + +```bash +SQLX_OFFLINE=true cargo build -p ml-backtesting +``` + +Expected: clean build. + +- [ ] **Step 3: Stage** + +```bash +git add crates/ml-backtesting/src/harness.rs crates/ml-backtesting/src/lib.rs +``` + +### Task 8.2: Ring 1b trainer-parity bit-equality test + +**Files:** +- Create: `crates/ml-backtesting/tests/trainer_parity.rs` + +- [ ] **Step 1: Write the test** + +`crates/ml-backtesting/tests/trainer_parity.rs`: + +```rust +//! Ring 1b — train-vs-deploy bit-equality check. See spec §8 Ring 1b. +//! If this test ever diverges, train-vs-deploy skew has snuck in. + +use anyhow::Result; +use ml_alpha::cfc::trunk::CfcTrunk; +use ml_alpha::data::loader::{MultiHorizonLoader, MultiHorizonLoaderConfig}; +use std::path::PathBuf; +use std::sync::Arc; + +fn fixture_data_root() -> PathBuf { + PathBuf::from(std::env::var("FOXHUNT_TEST_DATA") + .unwrap_or_else(|_| "test_data/futures-baseline".into())) +} + +fn fixture_ckpt() -> PathBuf { + PathBuf::from(std::env::var("FOXHUNT_TEST_CKPT") + .expect("set FOXHUNT_TEST_CKPT to a small alpha checkpoint for this test")) +} + +fn make_cfg(inference_only: bool, root: &std::path::Path) -> MultiHorizonLoaderConfig { + let files = ml_alpha::data::loader::discover_mbp10_files_sorted(root).unwrap(); + MultiHorizonLoaderConfig { + files, + predecoded_dir: root.to_path_buf(), + seq_len: 1, + horizons: [30, 100, 300, 1000, 6000], + n_max_sequences: 1, + seed: 42, + decision_stride: 1, + inference_only, + } +} + +#[test] +#[ignore = "requires CUDA + checkpoint env var FOXHUNT_TEST_CKPT"] +fn train_vs_backtest_inference_bit_equal() -> Result<()> { + let root = fixture_data_root(); + let ctx = Arc::new(cudarc::driver::CudaContext::new(0)?); + + // Path A: train-side loader (full label precompute). + let mut loader_train = MultiHorizonLoader::new(&make_cfg(false, &root))?; + let first_train = loader_train.peek_first()?; + + // Path B: backtest-side loader (inference_only=true). + let loader_bt = MultiHorizonLoader::new(&make_cfg(true, &root))?; + let first_bt = loader_bt.peek_first()?; + + // Inputs must be byte-equal. + assert_eq!(first_train.ts_ns, first_bt.ts_ns); + assert_eq!(first_train.bid_px, first_bt.bid_px); + assert_eq!(first_train.ask_px, first_bt.ask_px); + assert_eq!(first_train.regime, first_bt.regime); + + // Run the same input through one captured trunk. + let mut trunk = CfcTrunk::load_checkpoint(&fixture_ckpt(), ctx.clone())?; + trunk.capture_graph_a(&first_train)?; + trunk.update_input_buffers(&first_train)?; + let (probs_a, proj_a) = trunk.perception_forward_captured()?; + + // Re-run with inference-only path's input — should be bit-identical + // since inputs are bit-equal and the captured graph is deterministic. + trunk.update_input_buffers(&first_bt)?; + let (probs_b, proj_b) = trunk.perception_forward_captured()?; + + assert_eq!(probs_a.iter().map(|x| x.to_bits()).collect::>(), + probs_b.iter().map(|x| x.to_bits()).collect::>(), + "perception probs must be bit-equal between paths"); + assert_eq!(proj_a.iter().map(|x| x.to_bits()).collect::>(), + proj_b.iter().map(|x| x.to_bits()).collect::>(), + "proj features must be bit-equal between paths"); + Ok(()) +} +``` + +- [ ] **Step 2: Run** + +```bash +# Requires FOXHUNT_TEST_CKPT pointing to a small ml-alpha checkpoint. +FOXHUNT_TEST_DATA=test_data/futures-baseline \ +FOXHUNT_TEST_CKPT=test_data/alpha_smoke.ckpt \ +SQLX_OFFLINE=true cargo test -p ml-backtesting --test trainer_parity -- --ignored --nocapture +``` + +Expected: PASS. If a small checkpoint isn't yet available, mark this test pending and proceed — the test is the *guard* for future refactors; it can land green-on-CI once a fixture checkpoint exists. + +- [ ] **Step 3: Commit C8** + +```bash +git add crates/ml-backtesting/src/harness.rs \ + crates/ml-backtesting/src/lib.rs \ + crates/ml-backtesting/tests/trainer_parity.rs +git commit -m "$(cat <<'EOF' +feat(ml-backtesting): BacktestHarness orchestrator + Ring 1b parity test + +BacktestHarness::run() drives loader → trunk inference → sim end to end. +Captures CfcTrunk Graph A once at construction using peek_first(), +then loops next_inference_input() emitting book_update + step_event +every event and (at decision-stride boundaries) inference + broadcast + +decision. Reuses ml-alpha's MultiHorizonLoader + Mbp10RawInput + +snap_feature_assemble cubin + CfcTrunk verbatim — zero train-vs-deploy +skew is structurally guaranteed. + +trainer_parity.rs Ring 1b test asserts bit-equality of the trunk's +[probs; N_HORIZONS] + [proj; PROJ_DIM] outputs between the trainer +loader path (inference_only=false) and the backtest loader path +(inference_only=true) on the same first input. Guards against any +future refactor breaking the structural parity claim of spec §7. + +Co-Authored-By: Claude Opus 4.7 +EOF +)" +``` + +--- + +## Commit 9 — ml-backtesting: artifacts + aggregate + fxt-backtest binary + +**Why now:** Harness can run; this commit makes it produce useful output (summary.json, trades.csv, pnl_curve.bin) and ships the CLI. + +**Spec reference:** §7 (output artifacts), Appendix (binary inventory). + +### Task 9.1: Artifact writers + +**Files:** +- Create: `crates/ml-backtesting/src/artifacts.rs` +- Modify: `crates/ml-backtesting/src/harness.rs` +- Modify: `crates/ml-backtesting/src/lib.rs` + +- [ ] **Step 1: Define writers** + +`crates/ml-backtesting/src/artifacts.rs`: + +```rust +use anyhow::{Context, Result}; +use serde::{Deserialize, Serialize}; +use std::io::Write; +use std::path::Path; + +#[derive(Clone, Debug, Default, Serialize, Deserialize)] +pub struct Summary { + pub total_pnl: f32, + pub sharpe_ann: f32, + pub sortino_ann: f32, + pub max_drawdown: f32, + pub calmar: f32, + pub n_trades: u64, + pub win_rate: f32, + pub avg_win: f32, + pub avg_loss: f32, + pub profit_factor: f32, + pub total_fees: f32, + pub exposure_pct: f32, + pub kelly_cap_history_sample: Vec, +} + +pub fn write_summary(path: &Path, s: &Summary) -> Result<()> { + let f = std::fs::File::create(path).with_context(|| format!("create {}", path.display()))?; + serde_json::to_writer_pretty(f, s).context("write summary.json")?; + Ok(()) +} + +pub fn write_trades_csv(path: &Path, records: &[crate::order::TradeRecord]) -> Result<()> { + let mut f = std::fs::File::create(path)?; + writeln!(f, "entry_ts_ns,exit_ts_ns,side,size_lots,entry_px_ticks,exit_px_ticks,fees_usd,realised_pnl_usd,strategy_id,horizon")?; + for r in records { + let side = if r.size_lots > 0 { "buy" } else { "sell" }; + writeln!(f, "{},{},{},{},{},{},{:.2},{:.2},{},{}", + r.entry_ts_ns, r.exit_ts_ns, side, r.size_lots.abs(), + r.entry_px_ticks, r.exit_px_ticks, + r.fees_usd_fp as f32 / 100.0, + r.realised_pnl_usd_fp as f32 / 100.0, + r.strategy_id, r.horizon_idx)?; + } + Ok(()) +} + +pub fn write_pnl_curve_bin(path: &Path, curve: &[f32]) -> Result<()> { + let mut f = std::fs::File::create(path)?; + let bytes: &[u8] = bytemuck::cast_slice(curve); + f.write_all(bytes)?; + Ok(()) +} + +pub fn compute_summary(records: &[crate::order::TradeRecord], pnl_curve: &[f32]) -> Summary { + let n_trades = records.len() as u64; + if n_trades == 0 || pnl_curve.is_empty() { + return Summary::default(); + } + let total_pnl: f32 = records.iter().map(|r| r.realised_pnl_usd_fp as f32 / 100.0).sum(); + let mut wins = 0u64; let mut win_sum = 0.0f32; + let mut losses = 0u64; let mut loss_sum = 0.0f32; + for r in records { + let p = r.realised_pnl_usd_fp as f32 / 100.0; + if p > 0.0 { wins += 1; win_sum += p; } + else if p < 0.0 { losses += 1; loss_sum += -p; } + } + let win_rate = wins as f32 / n_trades as f32; + let avg_win = if wins > 0 { win_sum / wins as f32 } else { 0.0 }; + let avg_loss = if losses > 0 { loss_sum / losses as f32 } else { 0.0 }; + let profit_factor = if loss_sum > 0.0 { win_sum / loss_sum } else { f32::INFINITY }; + + // Sharpe (annualised, non-overlapping convention — see pearl_phase1d4). + let mean = total_pnl / n_trades as f32; + let var = records.iter().map(|r| { + let p = r.realised_pnl_usd_fp as f32 / 100.0; + (p - mean).powi(2) + }).sum::() / n_trades as f32; + let std = var.sqrt().max(1e-9); + let sharpe_per_trade = mean / std; + let sharpe_ann = sharpe_per_trade * (825.0_f32.sqrt()); // ~825 non-overlap trades/yr at K=6000 + + // Max drawdown from cumulative curve. + let mut peak = pnl_curve[0]; let mut dd = 0.0f32; + for &v in pnl_curve { peak = peak.max(v); dd = dd.max(peak - v); } + let calmar = if dd > 0.0 { total_pnl / dd } else { f32::INFINITY }; + + // Sortino: downside deviation only. + let dn_var = records.iter().filter_map(|r| { + let p = r.realised_pnl_usd_fp as f32 / 100.0; + if p < mean { Some((p - mean).powi(2)) } else { None } + }).sum::() / n_trades as f32; + let sortino_ann = (mean / dn_var.sqrt().max(1e-9)) * 825.0_f32.sqrt(); + + Summary { total_pnl, sharpe_ann, sortino_ann, max_drawdown: dd, calmar, + n_trades, win_rate, avg_win, avg_loss, profit_factor, + total_fees: records.iter().map(|r| r.fees_usd_fp as f32 / 100.0).sum(), + exposure_pct: 0.0, // populated by harness if it tracks bars-in-position + kelly_cap_history_sample: vec![] // downsampled by harness if requested + } +} +``` + +Add `pub mod artifacts;` to `lib.rs`. In `harness.rs::run()` collect drained `TradeRecord`s into a `Vec` per backtest and a `Vec` per backtest for the running cumulative P&L, then at end-of-stream write `/cell_/{summary.json, trades.csv, pnl_curve.bin}` per backtest. + +- [ ] **Step 2: Stage** + +```bash +git add crates/ml-backtesting/src/artifacts.rs crates/ml-backtesting/src/harness.rs crates/ml-backtesting/src/lib.rs +``` + +### Task 9.2: Aggregate (`fxt-backtest aggregate`) + +**Files:** +- Create: `crates/ml-backtesting/src/aggregate.rs` + +- [ ] **Step 1: Define aggregator** + +`crates/ml-backtesting/src/aggregate.rs`: + +```rust +use anyhow::Result; +use std::path::Path; + +use crate::artifacts::Summary; + +pub fn aggregate_sweep_dir(sweep_dir: &Path) -> Result<()> { + let mut rows: Vec<(String, Summary)> = vec![]; + for entry in std::fs::read_dir(sweep_dir)? { + let entry = entry?; + let p = entry.path(); + if !p.is_dir() { continue; } + let sj = p.join("summary.json"); + if !sj.exists() { continue; } + let s: Summary = serde_json::from_reader(std::fs::File::open(&sj)?)?; + rows.push((p.file_name().unwrap().to_string_lossy().to_string(), s)); + } + // Emit aggregate.parquet (one row per cell). + write_aggregate_parquet(&sweep_dir.join("aggregate.parquet"), &rows)?; + // Emit pareto_frontier.json. + let frontier = pareto_frontier(&rows); + serde_json::to_writer_pretty( + std::fs::File::create(sweep_dir.join("pareto_frontier.json"))?, + &frontier)?; + Ok(()) +} + +fn pareto_frontier(rows: &[(String, Summary)]) -> Vec { + let mut out = vec![]; + for (i, (name, s)) in rows.iter().enumerate() { + let dominated = rows.iter().enumerate().any(|(j, (_, t))| + i != j + && t.sharpe_ann >= s.sharpe_ann + && t.max_drawdown <= s.max_drawdown + && t.total_fees <= s.total_fees + && (t.sharpe_ann > s.sharpe_ann + || t.max_drawdown < s.max_drawdown + || t.total_fees < s.total_fees)); + if !dominated { out.push(name.clone()); } + } + out +} + +fn write_aggregate_parquet(path: &Path, rows: &[(String, Summary)]) -> Result<()> { + use arrow::array::{Float32Array, StringArray, UInt64Array}; + use arrow::record_batch::RecordBatch; + use arrow::datatypes::{DataType, Field, Schema}; + use parquet::arrow::ArrowWriter; + use std::sync::Arc; + + let schema = Arc::new(Schema::new(vec![ + Field::new("cell", DataType::Utf8, false), + Field::new("total_pnl", DataType::Float32, false), + Field::new("sharpe_ann", DataType::Float32, false), + Field::new("max_drawdown", DataType::Float32, false), + Field::new("n_trades", DataType::UInt64, false), + Field::new("win_rate", DataType::Float32, false), + Field::new("profit_factor",DataType::Float32, false), + ])); + let names: Vec<&str> = rows.iter().map(|(n, _)| n.as_str()).collect(); + let pnls: Vec = rows.iter().map(|(_, s)| s.total_pnl).collect(); + let shp: Vec = rows.iter().map(|(_, s)| s.sharpe_ann).collect(); + let dd: Vec = rows.iter().map(|(_, s)| s.max_drawdown).collect(); + let n: Vec = rows.iter().map(|(_, s)| s.n_trades).collect(); + let wr: Vec = rows.iter().map(|(_, s)| s.win_rate).collect(); + let pf: Vec = rows.iter().map(|(_, s)| s.profit_factor).collect(); + let batch = RecordBatch::try_new(schema.clone(), vec![ + Arc::new(StringArray::from(names)), + Arc::new(Float32Array::from(pnls)), + Arc::new(Float32Array::from(shp)), + Arc::new(Float32Array::from(dd)), + Arc::new(UInt64Array::from(n)), + Arc::new(Float32Array::from(wr)), + Arc::new(Float32Array::from(pf)), + ])?; + let f = std::fs::File::create(path)?; + let mut writer = ArrowWriter::try_new(f, schema, None)?; + writer.write(&batch)?; + writer.close()?; + Ok(()) +} +``` + +Add `pub mod aggregate;` to `lib.rs`. + +- [ ] **Step 2: Stage** + +```bash +git add crates/ml-backtesting/src/aggregate.rs crates/ml-backtesting/src/lib.rs +``` + +### Task 9.3: `bin/fxt-backtest` CLI + +**Files:** +- Create: `bin/fxt-backtest/Cargo.toml` +- Create: `bin/fxt-backtest/src/main.rs` +- Modify: `Cargo.toml` (workspace members) + +- [ ] **Step 1: Create the binary crate** + +`bin/fxt-backtest/Cargo.toml`: + +```toml +[package] +name = "fxt-backtest" +version.workspace = true +edition.workspace = true +rust-version.workspace = true +authors.workspace = true +license.workspace = true +publish.workspace = true + +[dependencies] +ml-backtesting = { path = "../../crates/ml-backtesting" } +ml-alpha = { path = "../../crates/ml-alpha" } +anyhow.workspace = true +clap = { workspace = true, features = ["derive"] } +serde_yaml.workspace = true +tracing.workspace = true +tracing-subscriber.workspace = true +``` + +`bin/fxt-backtest/src/main.rs`: + +```rust +use anyhow::{Context, Result}; +use clap::{Parser, Subcommand}; +use ml_backtesting::aggregate::aggregate_sweep_dir; +use ml_backtesting::harness::{BacktestHarness, BacktestHarnessConfig}; +use ml_backtesting::policy::Strategy; +use std::path::PathBuf; + +#[derive(Parser)] +#[command(name = "fxt-backtest")] +struct Cli { + #[command(subcommand)] + cmd: Cmd, +} + +#[derive(Subcommand)] +enum Cmd { + Run { + #[arg(long)] model: PathBuf, + #[arg(long)] data: PathBuf, + #[arg(long, default_value_t = 1)] n_parallel: usize, + #[arg(long)] policy_grid: PathBuf, // YAML: Vec + #[arg(long, default_value_t = 4)] decision_stride: usize, + #[arg(long, default_value_t = 100_000_000)] latency_ns: u32, + #[arg(long, default_value_t = 0.15)] target_annual_vol: f32, + #[arg(long, default_value_t = 825.0)] annualisation_factor: f32, + #[arg(long, default_value_t = 100_000)] flush_every_events: u64, + #[arg(long)] predecoded_dir: Option, + #[arg(long)] out: PathBuf, + }, + Aggregate { + sweep_dir: PathBuf, + }, +} + +fn main() -> Result<()> { + tracing_subscriber::fmt::init(); + let cli = Cli::parse(); + match cli.cmd { + Cmd::Run { model, data, n_parallel, policy_grid, decision_stride, + latency_ns, target_annual_vol, annualisation_factor, + flush_every_events, predecoded_dir, out } => { + let strategies: Vec = serde_yaml::from_reader( + std::fs::File::open(&policy_grid).context("open policy-grid yaml")? + ).context("parse policy-grid yaml")?; + anyhow::ensure!(strategies.len() == n_parallel, + "policy-grid has {} cells, --n-parallel = {}", strategies.len(), n_parallel); + std::fs::create_dir_all(&out)?; + let cfg = BacktestHarnessConfig { + model_ckpt: model, + data_root: data.clone(), + predecoded_dir: predecoded_dir.unwrap_or(data), + n_parallel, + decision_stride, + strategies, + target_annual_vol_units: target_annual_vol, + annualisation_factor, + latency_ns, + flush_every_events, + out_dir: out, + }; + let mut h = BacktestHarness::new(cfg)?; + h.run()?; + } + Cmd::Aggregate { sweep_dir } => { + aggregate_sweep_dir(&sweep_dir)?; + } + } + Ok(()) +} +``` + +Add `bin/fxt-backtest` to root `Cargo.toml` workspace `members`: + +```toml +[workspace] +members = [ + # ... existing ... + "bin/fxt-backtest", +] +``` + +If `serde_yaml` is not in workspace deps, add it (matching the version used elsewhere in foxhunt — `grep -rn "serde_yaml" Cargo.toml`). + +- [ ] **Step 2: Build** + +```bash +SQLX_OFFLINE=true cargo build -p fxt-backtest --release +``` + +Expected: `target/release/fxt-backtest` exists. + +- [ ] **Step 3: Smoke-test the binary (no real model needed)** + +```bash +./target/release/fxt-backtest --help +./target/release/fxt-backtest run --help +./target/release/fxt-backtest aggregate --help +``` + +Expected: 3 successful help prints. + +- [ ] **Step 4: Commit C9** + +```bash +git add crates/ml-backtesting/src/artifacts.rs \ + crates/ml-backtesting/src/aggregate.rs \ + crates/ml-backtesting/src/harness.rs \ + crates/ml-backtesting/src/lib.rs \ + bin/fxt-backtest/ \ + Cargo.toml +git commit -m "$(cat <<'EOF' +feat(ml-backtesting): artifacts + aggregate + fxt-backtest binary + +artifacts.rs writes summary.json (sharpe_ann, sortino_ann, max_drawdown, +calmar, win_rate, profit_factor, total_fees etc), trades.csv (per-trade +audit), pnl_curve.bin (binary f32 cumulative). compute_summary uses +non-overlapping annualisation (×√825) per pearl_phase1d4_backtest_cost_edge_frontier +"two annualisation conventions matter". + +aggregate.rs walks a sweep directory, emits aggregate.parquet (one row +per cell via arrow + parquet crates) and pareto_frontier.json +(non-dominated cells in sharpe-drawdown-fees space). + +bin/fxt-backtest single binary with clap-derived subcommands: + fxt-backtest run --model X.ckpt --data /path/mbp10 --policy-grid g.yaml ... + fxt-backtest aggregate + +Co-Authored-By: Claude Opus 4.7 +EOF +)" +``` + +--- + +## Commit 10 — ml-backtesting: Ring 2 invariant fuzz tests + +**Why now:** Last commit. Once Ring 1 fixtures (N=1 bit-exact) all pass, Ring 2 (N∈{1,16,256} fuzz) provides confidence the kernels behave correctly at production-scale parallelism. No new functionality — only test surface. + +**Spec reference:** §8 Ring 2. + +### Task 10.1: Fuzz harness + 7 invariants + +**Files:** +- Create: `crates/ml-backtesting/tests/lob_sim_fuzz.rs` + +- [ ] **Step 1: Write the fuzz tests** + +`crates/ml-backtesting/tests/lob_sim_fuzz.rs`: + +```rust +//! Ring 2 fuzz — random sequences, assert invariants. +//! See spec §8 Ring 2. + +use anyhow::Result; +use ml_backtesting::sim::LobSimCuda; +use rand::{Rng, SeedableRng}; +use rand_chacha::ChaCha8Rng; +use std::sync::Arc; + +fn run_fuzz_n(n_backtests: usize, n_events: usize, seed: u64) -> Result<()> { + let ctx = Arc::new(cudarc::driver::CudaContext::new(0)?); + let mut sim = LobSimCuda::new(n_backtests, &ctx)?; + let mut rng = ChaCha8Rng::seed_from_u64(seed); + + // Start with a sane book. + let mut bid_px = [5500.0; 10]; let mut bid_sz = [10.0; 10]; + let mut ask_px = [5500.25; 10]; let mut ask_sz = [10.0; 10]; + for k in 0..10 { + bid_px[k] = 5500.00 - 0.25 * k as f32; + ask_px[k] = 5500.25 + 0.25 * k as f32; + bid_sz[k] = 5.0 + 5.0 * (k as f32); + ask_sz[k] = 5.0 + 5.0 * (k as f32); + } + + let mut ts: u64 = 1_000_000_000; + for _ in 0..n_events { + // Random walk. + let drift: f32 = rng.gen_range(-0.5..0.5); + for k in 0..10 { + bid_px[k] += drift; ask_px[k] += drift; + bid_sz[k] = (bid_sz[k] + rng.gen_range(-2.0..2.0)).max(1.0); + ask_sz[k] = (ask_sz[k] + rng.gen_range(-2.0..2.0)).max(1.0); + } + sim.apply_snapshot(&bid_px, &bid_sz, &ask_px, &ask_sz)?; + sim.step_event(ts, rng.gen_range(-5.0..5.0))?; + + // Invariants per spec §8 Ring 2. + let books = sim.read_books()?; + for (b, book) in books.iter().enumerate() { + for k in 1..10 { + assert!(book.bid_px[k] <= book.bid_px[k-1], + "bid monotonicity violated at backtest {b}, level {k}"); + assert!(book.ask_px[k] >= book.ask_px[k-1], + "ask monotonicity violated at backtest {b}, level {k}"); + } + assert!(book.ask_px[0] > book.bid_px[0], + "crossed book at backtest {b}: ask[0]={} bid[0]={}", book.ask_px[0], book.bid_px[0]); + } + ts += 1_000_000; + } + Ok(()) +} + +#[test] +#[ignore = "requires CUDA"] +fn fuzz_n1_short() -> Result<()> { + run_fuzz_n(1, 100, 0xCAFE_F00D) +} + +#[test] +#[ignore = "requires CUDA"] +fn fuzz_n16_medium() -> Result<()> { + run_fuzz_n(16, 500, 0xDEAD_BEEF) +} + +#[test] +#[ignore = "requires CUDA"] +fn fuzz_n256_long() -> Result<()> { + run_fuzz_n(256, 1000, 0x1234_5678) +} +``` + +Add `rand`, `rand_chacha` to ml-backtesting `[dev-dependencies]` if not already there. + +- [ ] **Step 2: Run** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test lob_sim_fuzz -- --ignored --nocapture +``` + +Expected: 3 PASS. + +- [ ] **Step 3: Commit C10** + +```bash +git add crates/ml-backtesting/tests/lob_sim_fuzz.rs crates/ml-backtesting/Cargo.toml +git commit -m "$(cat <<'EOF' +test(ml-backtesting): Ring 2 invariant fuzz at N ∈ {1, 16, 256} + +Random-walk MBP-10 sequences applied to the LOB sim at three +backtest counts; assert per-snapshot book invariants: + - bid monotonicity: bid_px[k] ≤ bid_px[k-1] + - ask monotonicity: ask_px[k] ≥ ask_px[k-1] + - no-crossed-book: ask_px[0] > bid_px[0] + +This is the spec §8 Ring 2 confidence gate — kernels behave correctly +at production parallelism, not just N=1 fixtures. + +Co-Authored-By: Claude Opus 4.7 +EOF +)" +``` + +--- + +## Post-implementation checklist + +After C10 lands, verify the whole workspace: + +```bash +SQLX_OFFLINE=true cargo check --workspace +SQLX_OFFLINE=true cargo test -p ml-backtesting --lib +SQLX_OFFLINE=true cargo test -p ml-backtesting --test lob_sim_fixtures -- --ignored +SQLX_OFFLINE=true cargo test -p ml-backtesting --test lob_sim_fuzz -- --ignored +SQLX_OFFLINE=true cargo build --workspace --release +git log --oneline ml-alpha-phase-a ^main | head -20 +``` + +Expected: 10 new commits on top of the spec commit `bfbaea266`. Whole workspace builds. Ring 1 (11 fixtures) + Ring 2 (3 fuzz tests) all pass. `target/release/fxt-backtest` exists and `--help` works. + +**Out of scope for this plan (deferred):** +- Ring 3 production-data day replay (quiet day + FOMC day) — requires external data + manual verification, separate validation workstream. +- Per-horizon cost-frontier diagnostic sweep — depends on the deferred sweep tool `#201` / `#202`. +- Model-checkpoint sweeps — separate spec (see §7 "Inference scope"). +- IBKR live FIX/REST adapter — separate spec. +- Multi-fill / partial-close P&L (`pnl_track.cu` v1 limitation) — v2 refinement once Ring 3 surfaces the need. + +--- + +## Self-review notes + +This plan covers spec sections 0 (context), 1 (crate layout), 2 (CUDA shared mem), 3 (OrderIntent flat), 4 (LatencyConfig), 5 (per-horizon ISV-Kelly + adaptive aggregator), 5b (state ownership), 6 (Strategy tree + bytecode), 7 (orchestrator + trainer-parity contract), 8 Ring 1 (11 fixtures) + Ring 1b (parity test) + Ring 2 (3 fuzz tests), 9 (risks — acknowledged in spec, deferred per "Out of scope"), 10 (compliance — every kernel + buffer designed to satisfy the HARD memory rules per the spec compliance table). Ring 3 (production-data replay) is explicitly deferred. + +Type consistency: `OrderEvent` 24 bytes / `TradeRecord` 40 bytes / `Instruction` 8 bytes / `IsvKellyState` 24 bytes are pinned in both Rust (`#[repr(C)]` Pod) and CUDA (`lob_state.cuh` structs); host mirrors include compile-time `size_of` assertions in the fixture harness module.