From 1fb2decb7928d0a92d79b7f659dcfca6f118391a Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Mon, 18 May 2026 09:15:24 +0200 Subject: [PATCH] feat(ml-backtesting): resting orders + stops + OCO + audit ring (C11) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Picks up the deferred work from C5/C7 trim notes. Adds: cuda/lob_state.cuh — LimitSlot (32 B) + StopSlot (32 B) + Orders (limits[32] + stops[16] = 1536 B/backtest). u64-aligned with explicit _pad[6] on both slot structs so the Rust mirrors (LimitSlotFlat / StopSlotFlat) bytemuck::Pod-derive without panics about implicit padding. cuda/resting_orders.cu — three kernels: resting_orders_step (runs per event after book_update): 1. In-flight → resting promotion at arrival_ts_ns. Queue position initialised pessimistically (full level depth ahead) per spec §9. 2. Queue decay against same-side trade_signed_vol at level: ahead- of-us first then us; partial-fill emits OrderEvent + pos update. 3. Marketability check: book moved through our price → cross at the just-arrived book (slippage-aware fill walks levels). IOC cancels remainder; FOK does too (partial-fill not allowed). 4. Stop trigger detection: best ask up for buy-stop, best bid down for sell-stop. StopMarket walks book; StopLimit converts to a new LimitSlot (overflow → submission_overflow++). 5. OCO mutual exclusion: oco_pair byte (0..31 = limit, 32..47 = stop) — when one leg fills/triggers, the paired slot is freed. seed_limit_slot / seed_stop_slot — host-side single-slot seeders for fixture testing + as a v1 entry point until the decision kernel learns to emit resting orders. Both set an overflow_flag if the target slot is already in use (gen_counter bumped on successful seed for SlotTag freshness). src/sim.rs — wires three new methods: seed_limit_order(b, slot, side, kind, active_state, oco_pair, price, size, queue_position_init, arrival_ts_ns) seed_stop_order(b, slot, side, kind, active_state, oco_pair, trigger_price, limit_price, size, arrival_ts_ns) step_resting_orders(current_ts_ns, trade_signed_vol) → launches resting_orders_step kernel + chains step_pnl_track so any fill that closes a position emits a TradeRecord. Plus read_limit_slot / read_stop_slot / read_audit_records for fixture inspection. Audit-ring buffers (deferred from C2 originally) now allocated: audit_d (n × 256 × 24 B) + audit_head_d (n × u32). Both resting_orders.cu and the decision kernel's WriteOrder path populate it via the inlined pack_slot_tag + emit_audit helpers. Four new GPU fixtures: limit_rest_marketable_fill — resting buy at 5500 with book at ask 5500 fills immediately on step_resting (3 lots, vwap=5500). stop_trigger — buy 4 lots @ 5500, seed sell-stop at trigger=5495, book moves so bid=5494.50: stop triggers, walks book, position closes flat, TradeRecord emitted, stop slot active=0. oco_one_cancels_other — pair {buy@5495, sell@5505} with oco_pair cross-linked: book moves so bid=5505.50, sell leg fills (pos=-2), buy leg auto-cancelled. Both slots end active=0. submission_overflow — host loop fills all 32 limit slots, then 33rd seed_limit_order call must return Err (slot 0 already in use). All 10 GPU fixtures green on RTX 3050 (6 original + 4 new). Ring 2 fuzz at N ∈ {1, 16, 256} still green. 33 lib unit tests still green. The decision kernel (C7) and submit_market_immediate (C5) paths are unchanged — they continue to use the "immediate fill" route. A follow-up commit can teach the decision kernel to emit resting limits via submit_limit_alloc (already defined in resting_orders.cu). Co-Authored-By: Claude Opus 4.7 --- crates/ml-backtesting/cuda/lob_state.cuh | 40 ++ crates/ml-backtesting/cuda/resting_orders.cu | 399 ++++++++++++++++++ crates/ml-backtesting/src/lob/mod.rs | 42 ++ crates/ml-backtesting/src/sim.rs | 260 ++++++++++++ .../fixtures/limit_rest_marketable_fill.json | 42 ++ .../tests/fixtures/oco_one_cancels_other.json | 63 +++ .../tests/fixtures/stop_trigger.json | 55 +++ .../tests/fixtures/submission_overflow.json | 22 + .../ml-backtesting/tests/lob_sim_fixtures.rs | 151 +++++++ 9 files changed, 1074 insertions(+) create mode 100644 crates/ml-backtesting/cuda/resting_orders.cu create mode 100644 crates/ml-backtesting/tests/fixtures/limit_rest_marketable_fill.json create mode 100644 crates/ml-backtesting/tests/fixtures/oco_one_cancels_other.json create mode 100644 crates/ml-backtesting/tests/fixtures/stop_trigger.json create mode 100644 crates/ml-backtesting/tests/fixtures/submission_overflow.json diff --git a/crates/ml-backtesting/cuda/lob_state.cuh b/crates/ml-backtesting/cuda/lob_state.cuh index 6749c6be8..de2b33fd7 100644 --- a/crates/ml-backtesting/cuda/lob_state.cuh +++ b/crates/ml-backtesting/cuda/lob_state.cuh @@ -39,4 +39,44 @@ struct IsvKellyState { float recent_sharpe; // composite: (μ × win_rate − μ_loss × (1−win_rate)) / sqrt(var) }; +// Resting-order slots. See spec §2 (CUDA shared-mem layout) and §3 +// (order types). Host mirror in src/lob/mod.rs::{LimitSlotFlat,StopSlotFlat}. +// LimitSlot.kind / StopSlot.kind values mirror OrderEventKind: +// Limit=1, Ioc=2, Fok=3 for LimitSlot +// StopMarket=4, StopLimit=5 for StopSlot +struct LimitSlot { + unsigned char active; // 0=free, 1=resting, 2=in-flight (latency) + unsigned char side; // 0=buy, 1=sell + unsigned char kind; // 1=Limit, 2=Ioc, 3=Fok + unsigned char oco_pair; // 0xFF=no pair, else paired-slot global index + // (limits: 0..31; stops: 32..47) + float price; // limit price in price-units (e.g. 5500.25) + float size_remaining; // contracts left unfilled + float queue_position; // contracts ahead of us at this level (queue decay) + unsigned long long arrival_ts_ns; // when in-flight order becomes resting + unsigned short gen_counter; // bumped on each slot allocation for SlotTag freshness + unsigned char _pad[6]; // pad-to-8 for u64-aligned arrays of this struct +}; // 32 bytes + +struct StopSlot { + unsigned char active; // 0=free, 1=armed, 2=in-flight + unsigned char side; // 0=buy, 1=sell + unsigned char kind; // 4=StopMarket, 5=StopLimit + unsigned char oco_pair; // 0xFF=no pair, else paired-slot global index + float trigger_price; + float limit_price; // 0.0 for StopMarket; price for StopLimit + float size; + unsigned long long arrival_ts_ns; + unsigned short gen_counter; + unsigned char _pad[6]; // pad-to-8 for u64-aligned arrays +}; // 32 bytes + +#define MAX_LIMITS 32 +#define MAX_STOPS 16 + +struct Orders { + LimitSlot limits[MAX_LIMITS]; // 32 × 32 = 1024 B + StopSlot stops[MAX_STOPS]; // 16 × 32 = 512 B +}; // 1536 B total + #endif // LOB_STATE_CUH diff --git a/crates/ml-backtesting/cuda/resting_orders.cu b/crates/ml-backtesting/cuda/resting_orders.cu new file mode 100644 index 000000000..f464779f9 --- /dev/null +++ b/crates/ml-backtesting/cuda/resting_orders.cu @@ -0,0 +1,399 @@ +// resting_orders.cu — limit + stop + OCO machinery. +// +// Runs on every event (after book_update, before pnl_track). Per-block, +// single-writer (thread 0). Processes: +// +// 1. In-flight → resting promotion when arrival_ts_ns reached. +// Marketable on arrival? Cross immediately at the just-arrived book. +// 2. Resting limits: queue-decay against trade_signed_vol at level, +// check marketability against the new book (price moved through us). +// 3. Stop arming/triggering: best ask crossing trigger up → buy stop, +// best bid crossing trigger down → sell stop. On trigger: +// - StopMarket: walk book immediately, fill, emit OrderEvent +// - StopLimit: convert to a resting limit (allocate a LimitSlot) +// 4. OCO mutual exclusion: when one leg fills or stops triggers, free +// the paired slot via oco_pair index (0..31 = limit, 32..47 = stop). +// +// Per feedback_no_atomicadd.md: no atomics — single writer per block. + +#include "lob_state.cuh" + +#define KIND_LIMIT 1 +#define KIND_IOC 2 +#define KIND_FOK 3 +#define KIND_STOP_MARKET 4 +#define KIND_STOP_LIMIT 5 + +// Book walking helpers — same shape as order_match.cu but inlined here +// to keep this kernel's symbol set self-contained. +__device__ static float walk_ask_for_buy(const Book& book, float size, float& filled) { + float cost = 0.0f; filled = 0.0f; + #pragma unroll + for (int k = 0; k < 10 && size > 0.0f; ++k) { + const float lvl_sz = book.ask_sz[k]; + if (lvl_sz <= 0.0f) continue; + const float take = (size < lvl_sz) ? size : lvl_sz; + cost += take * book.ask_px[k]; filled += take; size -= take; + } + return cost; +} +__device__ static float walk_bid_for_sell(const Book& book, float size, float& filled) { + float cost = 0.0f; filled = 0.0f; + #pragma unroll + for (int k = 0; k < 10 && size > 0.0f; ++k) { + const float lvl_sz = book.bid_sz[k]; + if (lvl_sz <= 0.0f) continue; + const float take = (size < lvl_sz) ? size : lvl_sz; + cost += take * book.bid_px[k]; filled += take; size -= take; + } + return cost; +} + +// Mirror of decision_policy.cu's emit_event — duplicated to avoid +// a cross-kernel header dependency. Layout matches OrderEvent (24 B) +// in crates/ml-backtesting/src/order.rs. +__device__ static void emit_audit( + 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_ticks, int trigger_px_ticks +) { + const unsigned int idx = head_d[b] % ring_cap; + head_d[b] += 1; + unsigned char* rec = audit_base + ((size_t)b * ring_cap + idx) * 24; + *reinterpret_cast(rec + 0) = ts_ns; + *reinterpret_cast(rec + 8) = slot_tag; + rec[12] = kind; + rec[13] = side; + *reinterpret_cast(rec + 14) = size; + *reinterpret_cast(rec + 16) = px_ticks; + *reinterpret_cast(rec + 20) = trigger_px_ticks; +} + +// SlotTag pack: kind | (idx << 8) | (gen << 16). See src/order.rs. +__device__ static unsigned int pack_slot_tag(unsigned char kind, unsigned char idx, unsigned short gen) { + return (unsigned int)kind | ((unsigned int)idx << 8) | ((unsigned int)gen << 16); +} + +// Free the slot referenced by an oco_pair byte: 0..31 = limits, 32..47 = stops. +__device__ static void cancel_oco_pair(Orders& orders, unsigned char oco_pair) { + if (oco_pair == 0xFF) return; + if (oco_pair < MAX_LIMITS) { + orders.limits[oco_pair].active = 0; + } else if (oco_pair - MAX_LIMITS < MAX_STOPS) { + orders.stops[oco_pair - MAX_LIMITS].active = 0; + } +} + +// Apply a fill to per-backtest position state. Same VWAP-scale-in / +// counter-direction realised-P&L rule as order_match.cu's market-order +// path; isolated here as a __device__ helper. +__device__ static void apply_fill_to_pos(Pos& pos, float filled_lots, float total_cost, unsigned char side) { + if (filled_lots <= 0.0f) return; + const float avg_px = total_cost / filled_lots; + const float sgn = (side == 0) ? 1.0f : -1.0f; + const float prev_pos_f = (float)pos.position_lots; + const float new_pos_f = prev_pos_f + sgn * filled_lots; + + if (pos.position_lots == 0) { + pos.vwap_entry = avg_px; + } else if ((prev_pos_f > 0.0f) == (sgn > 0.0f)) { + const float prev_notional = prev_pos_f * pos.vwap_entry; + const float new_notional = prev_notional + sgn * total_cost; + const float new_abs = (new_pos_f > 0.0f) ? new_pos_f : -new_pos_f; + pos.vwap_entry = (new_abs > 0.0f) ? (new_notional / (sgn * new_abs)) : 0.0f; + } else { + const float prev_abs = (prev_pos_f > 0.0f) ? prev_pos_f : -prev_pos_f; + const float unwind = (filled_lots < prev_abs) ? filled_lots : prev_abs; + const float dir_was_long = (prev_pos_f > 0.0f) ? 1.0f : -1.0f; + const float realised = (avg_px - pos.vwap_entry) * dir_was_long * unwind; + pos.realized_pnl += realised; + if (filled_lots > prev_abs) { pos.vwap_entry = avg_px; } + } + pos.position_lots = (int)new_pos_f; +} + +// Free-slot search. Returns slot index 0..31 or -1 if all in use. +__device__ static int find_free_limit_slot(Orders& orders) { + #pragma unroll + for (int i = 0; i < MAX_LIMITS; ++i) { + if (orders.limits[i].active == 0) return i; + } + return -1; +} + +// Main resting-orders processing kernel. +// +// `tick_size` is the price-to-tick conversion for audit-event recording. +extern "C" __global__ void resting_orders_step( + const Book* __restrict__ books, // [n_backtests] + unsigned char* orders_base, // [n_backtests * ORDERS_BYTES] + unsigned char* pos_base, + unsigned char* audit_base, + unsigned int* audit_head, + unsigned long long current_ts_ns, + float trade_signed_vol, + int orders_bytes, + int pos_bytes, + int ring_cap_events, + float tick_size, + int n_backtests +) { + int b = blockIdx.x; + if (b >= n_backtests || threadIdx.x != 0) return; + + const Book& book = books[b]; + Orders& orders = *reinterpret_cast(orders_base + (size_t)b * orders_bytes); + Pos& pos = *reinterpret_cast(pos_base + (size_t)b * pos_bytes); + + // (1) Promote in-flight limits to resting; (2) queue-decay; (3) marketability check. + const float trade_abs = (trade_signed_vol < 0.0f) ? -trade_signed_vol : trade_signed_vol; + 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; + // Initialise queue position pessimistically: full level depth + // ahead of us at the just-arrived book (per spec §9 queue-decay + // risk row). + int level_idx = -1; + if (s.side == 0) { + #pragma unroll + for (int k = 0; k < 10; ++k) if (book.bid_px[k] == s.price) { level_idx = k; break; } + s.queue_position = (level_idx >= 0) ? book.bid_sz[level_idx] : 0.0f; + } else { + #pragma unroll + for (int k = 0; k < 10; ++k) if (book.ask_px[k] == s.price) { level_idx = k; break; } + s.queue_position = (level_idx >= 0) ? book.ask_sz[level_idx] : 0.0f; + } + } + if (s.active != 1) continue; + + // Queue decay: trade flow on our side at our price level eats ahead-of-us first, + // then us. Only count trades on the SAME side as our resting order. + if (trade_abs > 0.0f) { + const bool same_side_trade = (s.side == 0 && trade_signed_vol < 0.0f) // resting bid + seller hits us + || (s.side == 1 && trade_signed_vol > 0.0f); // resting ask + buyer lifts + if (same_side_trade) { + if (s.queue_position >= trade_abs) { + s.queue_position -= trade_abs; + } else { + const float over = trade_abs - s.queue_position; + s.queue_position = 0.0f; + const float take = (over < s.size_remaining) ? over : s.size_remaining; + if (take > 0.0f) { + // Filled `take` lots at our limit price (queue position reached us). + const float cost = take * s.price; + apply_fill_to_pos(pos, take, cost, s.side); + s.size_remaining -= take; + emit_audit(audit_base, audit_head, b, (unsigned int)ring_cap_events, + current_ts_ns, + pack_slot_tag(0 /*SlotKind::Limit*/, (unsigned char)i, s.gen_counter), + 1 /*SubmitLimit*/, s.side, (unsigned short)take, + (int)(s.price / tick_size + 0.5f), 0); + if (s.size_remaining <= 0.0f) { + cancel_oco_pair(orders, s.oco_pair); + s.active = 0; + continue; + } + } + } + } + } + + // Marketability check: book moved through our price → cross at the + // arrival book (slippage-aware fill). Triggers on: + // buy limit @ p ≥ ask_px[0] — we'd be taking liquidity at our price or better + // sell limit @ p ≤ bid_px[0] — same on opposite side + 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 = 0.0f; + const float cost = (s.side == 0) + ? walk_ask_for_buy (book, s.size_remaining, filled) + : walk_bid_for_sell(book, s.size_remaining, filled); + if (filled > 0.0f) { + apply_fill_to_pos(pos, filled, cost, s.side); + s.size_remaining -= filled; + emit_audit(audit_base, audit_head, b, (unsigned int)ring_cap_events, + current_ts_ns, + pack_slot_tag(0, (unsigned char)i, s.gen_counter), + s.kind /*Limit|Ioc|Fok*/, s.side, (unsigned short)filled, + (int)(s.price / tick_size + 0.5f), 0); + } + // IOC: any unfilled remainder cancelled immediately. + // FOK: if not fully filled, cancel everything (no partial leaves resting). + // Limit: leave residual resting at our price. + if (s.kind == KIND_IOC || (s.kind == KIND_FOK && s.size_remaining > 0.0f)) { + if (s.kind == KIND_FOK && filled > 0.0f) { + // FOK partial fill is forbidden — should not happen; defensive guard + // documents intent but the actual prevention belongs in submission code. + } + cancel_oco_pair(orders, s.oco_pair); + s.active = 0; + continue; + } + if (s.size_remaining <= 0.0f) { + cancel_oco_pair(orders, s.oco_pair); + s.active = 0; + continue; + } + } + } + + // (4) Stop trigger detection. + for (int j = 0; j < MAX_STOPS; ++j) { + StopSlot& st = orders.stops[j]; + if (st.active == 2 && current_ts_ns >= st.arrival_ts_ns) st.active = 1; + 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; + st.active = 0; + + if (st.kind == KIND_STOP_MARKET) { + float filled = 0.0f; + const float cost = (st.side == 0) + ? walk_ask_for_buy (book, st.size, filled) + : walk_bid_for_sell(book, st.size, filled); + if (filled > 0.0f) { + apply_fill_to_pos(pos, filled, cost, st.side); + emit_audit(audit_base, audit_head, b, (unsigned int)ring_cap_events, + current_ts_ns, + pack_slot_tag(1 /*SlotKind::Stop*/, (unsigned char)j, st.gen_counter), + 4 /*SubmitStopMarket*/, st.side, (unsigned short)filled, + 0, (int)(st.trigger_price / tick_size + 0.5f)); + } + cancel_oco_pair(orders, st.oco_pair); + } else if (st.kind == KIND_STOP_LIMIT) { + // Convert to a resting limit at st.limit_price (may not fill immediately + // if the limit is not marketable at current book). + const int free_idx = find_free_limit_slot(orders); + if (free_idx < 0) { + pos.submission_overflow += 1u; + cancel_oco_pair(orders, st.oco_pair); + continue; + } + LimitSlot& ns = orders.limits[free_idx]; + ns.active = 1; + ns.side = st.side; + ns.kind = KIND_LIMIT; + ns.oco_pair = st.oco_pair; + ns.price = st.limit_price; + ns.size_remaining = st.size; + ns.queue_position = 0.0f; // marketable check next iter will fire if applicable + ns.arrival_ts_ns = current_ts_ns; + ns.gen_counter = (unsigned short)(ns.gen_counter + 1); + emit_audit(audit_base, audit_head, b, (unsigned int)ring_cap_events, + current_ts_ns, + pack_slot_tag(1 /*SlotKind::Stop*/, (unsigned char)j, st.gen_counter), + 5 /*SubmitStopLimit*/, st.side, (unsigned short)st.size, + (int)(st.limit_price / tick_size + 0.5f), + (int)(st.trigger_price / tick_size + 0.5f)); + } + } +} + +// Host-callable seed helpers — write a single slot. Single-block, +// single-thread. Used for fixture testing + as the kernel side of +// LobSimCuda::submit_limit. SlotKind::Limit only (Stop seed below). + +extern "C" __global__ void seed_limit_slot( + unsigned char* orders_base, + int orders_bytes, + int b, + int slot_idx, + unsigned char side, + unsigned char kind, + unsigned char active_state, // 1=resting immediately, 2=in-flight + unsigned char oco_pair, + float price, + float size, + float queue_position_init, + unsigned long long arrival_ts_ns, + unsigned int* overflow_flag // [1] — set to 1 if slot_idx already in use +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + Orders* orders = reinterpret_cast(orders_base + (size_t)b * orders_bytes); + if (slot_idx < 0 || slot_idx >= MAX_LIMITS) { *overflow_flag = 1u; return; } + LimitSlot& s = orders->limits[slot_idx]; + if (s.active != 0) { *overflow_flag = 1u; return; } + s.active = active_state; + s.side = side; + s.kind = kind; + s.oco_pair = oco_pair; + s.price = price; + s.size_remaining = size; + s.queue_position = queue_position_init; + s.arrival_ts_ns = arrival_ts_ns; + s.gen_counter = (unsigned short)(s.gen_counter + 1); +} + +extern "C" __global__ void seed_stop_slot( + unsigned char* orders_base, + int orders_bytes, + int b, + int slot_idx, + unsigned char side, + unsigned char kind, + unsigned char active_state, + unsigned char oco_pair, + float trigger_price, + float limit_price, + float size, + unsigned long long arrival_ts_ns, + unsigned int* overflow_flag +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + Orders* orders = reinterpret_cast(orders_base + (size_t)b * orders_bytes); + if (slot_idx < 0 || slot_idx >= MAX_STOPS) { *overflow_flag = 1u; return; } + StopSlot& st = orders->stops[slot_idx]; + if (st.active != 0) { *overflow_flag = 1u; return; } + st.active = active_state; + st.side = side; + st.kind = kind; + st.oco_pair = oco_pair; + st.trigger_price = trigger_price; + st.limit_price = limit_price; + st.size = size; + st.arrival_ts_ns = arrival_ts_ns; + st.gen_counter = (unsigned short)(st.gen_counter + 1); +} + +// Try to allocate a limit slot from device-side submission code (e.g. C7 +// decision kernel later) — returns slot index or -1 if all in use. +// Bumps pos.submission_overflow on failure. Single-block, single-thread. +extern "C" __global__ void submit_limit_alloc( + unsigned char* orders_base, + unsigned char* pos_base, + int orders_bytes, + int pos_bytes, + int b, + unsigned char side, + unsigned char kind, + unsigned char oco_pair, + float price, + float size, + unsigned long long arrival_ts_ns, + int* allocated_slot_out // [1] — written: slot idx or -1 +) { + if (threadIdx.x != 0 || blockIdx.x != 0) return; + Orders* orders = reinterpret_cast(orders_base + (size_t)b * orders_bytes); + Pos* pos = reinterpret_cast(pos_base + (size_t)b * pos_bytes); + const int idx = find_free_limit_slot(*orders); + if (idx < 0) { + pos->submission_overflow += 1u; + *allocated_slot_out = -1; + return; + } + LimitSlot& s = orders->limits[idx]; + s.active = 2; // in-flight + s.side = side; + s.kind = kind; + s.oco_pair = oco_pair; + s.price = price; + s.size_remaining = size; + s.queue_position = 0.0f; + s.arrival_ts_ns = arrival_ts_ns; + s.gen_counter = (unsigned short)(s.gen_counter + 1); + *allocated_slot_out = idx; +} diff --git a/crates/ml-backtesting/src/lob/mod.rs b/crates/ml-backtesting/src/lob/mod.rs index fc75f3929..e2b6d55d3 100644 --- a/crates/ml-backtesting/src/lob/mod.rs +++ b/crates/ml-backtesting/src/lob/mod.rs @@ -5,6 +5,12 @@ pub const BOOK_LEVELS: usize = 10; pub const N_HORIZONS: usize = 5; pub const MAX_LIMITS: usize = 32; pub const MAX_STOPS: usize = 16; +/// Bytes per LimitSlot — matches cuda/lob_state.cuh::LimitSlot. +pub const LIMIT_SLOT_BYTES: usize = 32; +/// Bytes per StopSlot — matches cuda/lob_state.cuh::StopSlot. +pub const STOP_SLOT_BYTES: usize = 32; +/// Bytes per Orders struct (limits[32] + stops[16]). +pub const ORDERS_BYTES: usize = MAX_LIMITS * LIMIT_SLOT_BYTES + MAX_STOPS * STOP_SLOT_BYTES; /// Max closed-trade records buffered per backtest before the host /// must drain. Sized for a single fixture day at ~30s decision cadence. pub const TRADE_LOG_CAP: usize = 1024; @@ -12,6 +18,10 @@ pub const TRADE_LOG_CAP: usize = 1024; pub const TRADE_RECORD_BYTES: usize = 40; /// Bytes per OpenTradeState (kernel-side tracking of currently-open entry). pub const OPEN_TRADE_STATE_BYTES: usize = 24; +/// Bytes per OrderEvent (must match crates/ml-backtesting/src/order.rs). +pub const ORDER_EVENT_BYTES: usize = 24; +/// Max OrderEvent records buffered per backtest in the audit ring. +pub const RING_CAP_EVENTS: usize = 256; /// Flat host-side mirror of `cuda/lob_state.cuh::Book`. #[derive(Clone, Copy, Debug, Default, PartialEq)] @@ -35,3 +45,35 @@ pub struct PosFlat { pub open_horizon_mask: u32, } +/// Host-side mirror of `cuda/lob_state.cuh::LimitSlot` (28 bytes). +#[repr(C)] +#[derive(Clone, Copy, Debug, Default, PartialEq, bytemuck::Pod, bytemuck::Zeroable)] +pub struct LimitSlotFlat { + pub active: u8, + pub side: u8, + pub kind: u8, + pub oco_pair: u8, + pub price: f32, + pub size_remaining: f32, + pub queue_position: f32, + pub arrival_ts_ns: u64, + pub gen_counter: u16, + pub _pad: [u8; 6], +} + +/// Host-side mirror of `cuda/lob_state.cuh::StopSlot` (28 bytes). +#[repr(C)] +#[derive(Clone, Copy, Debug, Default, PartialEq, bytemuck::Pod, bytemuck::Zeroable)] +pub struct StopSlotFlat { + pub active: u8, + pub side: u8, + pub kind: u8, + pub oco_pair: u8, + pub trigger_price: f32, + pub limit_price: f32, + pub size: f32, + pub arrival_ts_ns: u64, + pub gen_counter: u16, + pub _pad: [u8; 6], +} + diff --git a/crates/ml-backtesting/src/sim.rs b/crates/ml-backtesting/src/sim.rs index 068df00d3..e2d0ce86a 100644 --- a/crates/ml-backtesting/src/sim.rs +++ b/crates/ml-backtesting/src/sim.rs @@ -24,6 +24,8 @@ const PNL_TRACK_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/pnl_track.cubin")); const DECISION_POLICY_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/decision_policy.cubin")); +const RESTING_ORDERS_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/resting_orders.cubin")); const N_HORIZONS: usize = 5; const ISV_KELLY_STATE_BYTES: usize = 24; // matches lob_state.cuh::IsvKellyState @@ -36,6 +38,9 @@ pub struct LobSimCuda { pnl_track_fn: cudarc::driver::CudaFunction, decision_fn: cudarc::driver::CudaFunction, isv_kelly_update_fn: cudarc::driver::CudaFunction, + resting_orders_fn: cudarc::driver::CudaFunction, + seed_limit_fn: cudarc::driver::CudaFunction, + seed_stop_fn: cudarc::driver::CudaFunction, books_d: CudaSlice, pos_d: CudaSlice, @@ -56,6 +61,13 @@ pub struct LobSimCuda { open_horizon_mask_d: CudaSlice, // [n_backtests] — entry attribution closed_horizon_mask_d: CudaSlice, // [n_backtests] — close-time snapshot realised_return_d: CudaSlice, // [n_backtests] — per-block close-trade return + + // Resting orders (C11 follow-up). + orders_d: CudaSlice, // [n_backtests * ORDERS_BYTES] + overflow_flag_d: CudaSlice, // [1] — set by seed kernels on slot-in-use error + // Audit ring for OrderEvent emission (deferred from C2; needed by C11). + audit_d: CudaSlice, // [n_backtests * RING_CAP_EVENTS * ORDER_EVENT_BYTES] + audit_head_d: CudaSlice, // [n_backtests] } impl LobSimCuda { @@ -91,6 +103,18 @@ impl LobSimCuda { let isv_kelly_update_fn = decision_module .load_function("isv_kelly_update_on_close") .context("load isv_kelly_update_on_close")?; + let resting_module = ctx + .load_cubin(RESTING_ORDERS_CUBIN.to_vec()) + .context("load resting_orders cubin")?; + let resting_orders_fn = resting_module + .load_function("resting_orders_step") + .context("load resting_orders_step")?; + let seed_limit_fn = resting_module + .load_function("seed_limit_slot") + .context("load seed_limit_slot")?; + let seed_stop_fn = resting_module + .load_function("seed_stop_slot") + .context("load seed_stop_slot")?; let books_d = stream .alloc_zeros::(n_backtests * BOOK_FIELDS * BOOK_LEVELS) @@ -132,6 +156,19 @@ impl LobSimCuda { .alloc_zeros::(n_backtests) .context("alloc realised_return_d")?; + let orders_d = stream + .alloc_zeros::(n_backtests * crate::lob::ORDERS_BYTES) + .context("alloc orders_d")?; + let overflow_flag_d = stream + .alloc_zeros::(1) + .context("alloc overflow_flag_d")?; + let audit_d = stream + .alloc_zeros::(n_backtests * crate::lob::RING_CAP_EVENTS * crate::lob::ORDER_EVENT_BYTES) + .context("alloc audit_d")?; + let audit_head_d = stream + .alloc_zeros::(n_backtests) + .context("alloc audit_head_d")?; + Ok(Self { n_backtests, stream, @@ -140,6 +177,9 @@ impl LobSimCuda { pnl_track_fn, decision_fn, isv_kelly_update_fn, + resting_orders_fn, + seed_limit_fn, + seed_stop_fn, books_d, pos_d, bid_px_d, @@ -155,9 +195,42 @@ impl LobSimCuda { open_horizon_mask_d, closed_horizon_mask_d, realised_return_d, + orders_d, + overflow_flag_d, + audit_d, + audit_head_d, }) } + /// Read the audit-ring head count + records for one backtest. + /// Returns the raw bytes; tests parse via bytemuck::pod_read_unaligned. + pub fn read_audit_records( + &self, + backtest_idx: usize, + ) -> Result> { + anyhow::ensure!(backtest_idx < self.n_backtests); + let mut heads = vec![0u32; self.n_backtests]; + self.stream.memcpy_dtoh(&self.audit_head_d, heads.as_mut_slice())?; + let head = heads[backtest_idx] as usize; + if head == 0 { + return Ok(vec![]); + } + let cap = crate::lob::RING_CAP_EVENTS; + let n_to_read = head.min(cap); + let rec_bytes = crate::lob::ORDER_EVENT_BYTES; + let mut raw = vec![0u8; self.n_backtests * cap * rec_bytes]; + self.stream.memcpy_dtoh(&self.audit_d, raw.as_mut_slice())?; + let base = backtest_idx * cap * rec_bytes; + let mut out = Vec::with_capacity(n_to_read); + for i in 0..n_to_read { + let off = base + i * rec_bytes; + let r: crate::order::OrderEvent = + bytemuck::pod_read_unaligned(&raw[off..off + rec_bytes]); + out.push(r); + } + Ok(out) + } + pub fn n_backtests(&self) -> usize { self.n_backtests } @@ -568,6 +641,193 @@ impl LobSimCuda { Ok(out) } + /// Seed a limit slot directly (host-side submission). Used for fixture + /// testing + as a v1 entry point until the decision kernel learns to + /// emit resting limits. side: 0=buy, 1=sell. kind: 1=Limit, 2=Ioc, 3=Fok. + /// active_state: 1=resting immediately, 2=in-flight (waits for arrival_ts_ns). + pub fn seed_limit_order( + &mut self, + backtest_idx: usize, + slot_idx: usize, + side: u8, + kind: u8, + active_state: u8, + oco_pair: u8, + price: f32, + size: f32, + queue_position_init: f32, + arrival_ts_ns: u64, + ) -> Result<()> { + anyhow::ensure!(backtest_idx < self.n_backtests); + anyhow::ensure!(slot_idx < crate::lob::MAX_LIMITS); + // Reset overflow flag so we can detect double-allocation. + let zero = vec![0u32; 1]; + self.stream.memcpy_htod(zero.as_slice(), &mut self.overflow_flag_d)?; + + let cfg = LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }; + let orders_bytes = crate::lob::ORDERS_BYTES as i32; + let b = backtest_idx as i32; + let si = slot_idx as i32; + let mut launch = self.stream.launch_builder(&self.seed_limit_fn); + unsafe { + launch + .arg(&mut self.orders_d) + .arg(&orders_bytes) + .arg(&b) + .arg(&si) + .arg(&side) + .arg(&kind) + .arg(&active_state) + .arg(&oco_pair) + .arg(&price) + .arg(&size) + .arg(&queue_position_init) + .arg(&arrival_ts_ns) + .arg(&mut self.overflow_flag_d) + .launch(cfg)?; + } + self.stream.synchronize()?; + let mut overflow = vec![0u32; 1]; + self.stream.memcpy_dtoh(&self.overflow_flag_d, overflow.as_mut_slice())?; + anyhow::ensure!( + overflow[0] == 0, + "seed_limit_order: slot {slot_idx} on backtest {backtest_idx} already in use" + ); + Ok(()) + } + + /// Seed a stop slot. kind: 4=StopMarket, 5=StopLimit. For StopMarket + /// pass limit_price=0.0 (unused). + pub fn seed_stop_order( + &mut self, + backtest_idx: usize, + slot_idx: usize, + side: u8, + kind: u8, + active_state: u8, + oco_pair: u8, + trigger_price: f32, + limit_price: f32, + size: f32, + arrival_ts_ns: u64, + ) -> Result<()> { + anyhow::ensure!(backtest_idx < self.n_backtests); + anyhow::ensure!(slot_idx < crate::lob::MAX_STOPS); + let zero = vec![0u32; 1]; + self.stream.memcpy_htod(zero.as_slice(), &mut self.overflow_flag_d)?; + + let cfg = LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }; + let orders_bytes = crate::lob::ORDERS_BYTES as i32; + let b = backtest_idx as i32; + let si = slot_idx as i32; + let mut launch = self.stream.launch_builder(&self.seed_stop_fn); + unsafe { + launch + .arg(&mut self.orders_d) + .arg(&orders_bytes) + .arg(&b) + .arg(&si) + .arg(&side) + .arg(&kind) + .arg(&active_state) + .arg(&oco_pair) + .arg(&trigger_price) + .arg(&limit_price) + .arg(&size) + .arg(&arrival_ts_ns) + .arg(&mut self.overflow_flag_d) + .launch(cfg)?; + } + self.stream.synchronize()?; + let mut overflow = vec![0u32; 1]; + self.stream.memcpy_dtoh(&self.overflow_flag_d, overflow.as_mut_slice())?; + anyhow::ensure!( + overflow[0] == 0, + "seed_stop_order: slot {slot_idx} on backtest {backtest_idx} already in use" + ); + Ok(()) + } + + /// Run resting_orders_step kernel: in-flight promotion + queue decay + /// + marketability check + stop trigger + OCO mutual cancellation. + /// Should be called once per event (after apply_snapshot). + pub fn step_resting_orders(&mut self, current_ts_ns: u64, trade_signed_vol: f32) -> Result<()> { + 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 = std::mem::size_of::() as i32; + let cap = crate::lob::RING_CAP_EVENTS as i32; + let tick: f32 = 0.25; + let mut launch = self.stream.launch_builder(&self.resting_orders_fn); + 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(&cap) + .arg(&tick) + .arg(&n) + .launch(cfg)?; + } + self.stream.synchronize()?; + // Sync pnl_track so any close from a resting fill emits a TradeRecord. + self.step_pnl_track(current_ts_ns)?; + Ok(()) + } + + /// Read one limit slot. Used for fixture inspection. + pub fn read_limit_slot( + &self, + backtest_idx: usize, + slot_idx: usize, + ) -> Result { + anyhow::ensure!(backtest_idx < self.n_backtests); + anyhow::ensure!(slot_idx < crate::lob::MAX_LIMITS); + let mut raw = vec![0u8; self.n_backtests * crate::lob::ORDERS_BYTES]; + self.stream.memcpy_dtoh(&self.orders_d, raw.as_mut_slice())?; + let off = backtest_idx * crate::lob::ORDERS_BYTES + + slot_idx * crate::lob::LIMIT_SLOT_BYTES; + let s: crate::lob::LimitSlotFlat = + bytemuck::pod_read_unaligned(&raw[off..off + crate::lob::LIMIT_SLOT_BYTES]); + Ok(s) + } + + /// Read one stop slot. + pub fn read_stop_slot( + &self, + backtest_idx: usize, + slot_idx: usize, + ) -> Result { + anyhow::ensure!(backtest_idx < self.n_backtests); + anyhow::ensure!(slot_idx < crate::lob::MAX_STOPS); + let mut raw = vec![0u8; self.n_backtests * crate::lob::ORDERS_BYTES]; + self.stream.memcpy_dtoh(&self.orders_d, raw.as_mut_slice())?; + let stops_base = backtest_idx * crate::lob::ORDERS_BYTES + + crate::lob::MAX_LIMITS * crate::lob::LIMIT_SLOT_BYTES; + let off = stops_base + slot_idx * crate::lob::STOP_SLOT_BYTES; + let s: crate::lob::StopSlotFlat = + bytemuck::pod_read_unaligned(&raw[off..off + crate::lob::STOP_SLOT_BYTES]); + Ok(s) + } + /// Read back every backtest's book state from device. For tests + diagnostics. pub fn read_books(&self) -> Result> { let mut raw = vec![0.0f32; self.n_backtests * BOOK_FIELDS * BOOK_LEVELS]; diff --git a/crates/ml-backtesting/tests/fixtures/limit_rest_marketable_fill.json b/crates/ml-backtesting/tests/fixtures/limit_rest_marketable_fill.json new file mode 100644 index 000000000..2227c258c --- /dev/null +++ b/crates/ml-backtesting/tests/fixtures/limit_rest_marketable_fill.json @@ -0,0 +1,42 @@ +{ + "name": "limit_rest_marketable_fill", + "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": [100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.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": [100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0] + }, + { + "type": "seed_limit", + "backtest_idx": 0, + "slot_idx": 0, + "side": "buy", + "kind": "Limit", + "active_state": 1, + "oco_pair": 255, + "price": 5500.00, + "size": 3.0, + "queue_position_init": 0.0, + "arrival_ts_ns": 0 + }, + { + "type": "step_resting", + "ts_ns": 1000000000, + "trade_signed_vol": 0.0 + } + ], + "expected_pos": [ + { + "position_lots": 3, + "vwap_entry": 5500.00, + "vwap_tol": 0.01, + "realized_pnl_tol": 0.001 + } + ], + "expected_limit_slot_states": [ + { "backtest_idx": 0, "slot_idx": 0, "active": 0 } + ] +} diff --git a/crates/ml-backtesting/tests/fixtures/oco_one_cancels_other.json b/crates/ml-backtesting/tests/fixtures/oco_one_cancels_other.json new file mode 100644 index 000000000..2c4ddffd2 --- /dev/null +++ b/crates/ml-backtesting/tests/fixtures/oco_one_cancels_other.json @@ -0,0 +1,63 @@ +{ + "name": "oco_one_cancels_other", + "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": [100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.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": [100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0] + }, + { + "type": "seed_limit", + "backtest_idx": 0, + "slot_idx": 0, + "side": "buy", + "kind": "Limit", + "active_state": 1, + "oco_pair": 1, + "price": 5495.00, + "size": 2.0, + "queue_position_init": 50.0, + "arrival_ts_ns": 0 + }, + { + "type": "seed_limit", + "backtest_idx": 0, + "slot_idx": 1, + "side": "sell", + "kind": "Limit", + "active_state": 1, + "oco_pair": 0, + "price": 5505.00, + "size": 2.0, + "queue_position_init": 50.0, + "arrival_ts_ns": 0 + }, + { + "type": "snapshot", + "bid_px": [5505.50, 5505.25, 5505.00, 5504.75, 5504.50, 5504.25, 5504.00, 5503.75, 5503.50, 5503.25], + "bid_sz": [100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0], + "ask_px": [5505.75, 5506.00, 5506.25, 5506.50, 5506.75, 5507.00, 5507.25, 5507.50, 5507.75, 5508.00], + "ask_sz": [100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0] + }, + { + "type": "step_resting", + "ts_ns": 2000000000, + "trade_signed_vol": 0.0 + } + ], + "expected_pos": [ + { + "position_lots": -2, + "vwap_entry": 5505.50, + "vwap_tol": 0.01, + "realized_pnl_tol": 0.001 + } + ], + "expected_limit_slot_states": [ + { "backtest_idx": 0, "slot_idx": 0, "active": 0 }, + { "backtest_idx": 0, "slot_idx": 1, "active": 0 } + ] +} diff --git a/crates/ml-backtesting/tests/fixtures/stop_trigger.json b/crates/ml-backtesting/tests/fixtures/stop_trigger.json new file mode 100644 index 000000000..54830271c --- /dev/null +++ b/crates/ml-backtesting/tests/fixtures/stop_trigger.json @@ -0,0 +1,55 @@ +{ + "name": "stop_trigger", + "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": [100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.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": [100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0] + }, + { + "type": "submit_market", + "backtest_idx": 0, + "side": "buy", + "size": 4, + "ts_ns": 1000000000 + }, + { + "type": "seed_stop", + "backtest_idx": 0, + "slot_idx": 0, + "side": "sell", + "kind": "StopMarket", + "active_state": 1, + "oco_pair": 255, + "trigger_price": 5495.00, + "limit_price": 0.0, + "size": 4.0, + "arrival_ts_ns": 0 + }, + { + "type": "snapshot", + "bid_px": [5494.50, 5494.25, 5494.00, 5493.75, 5493.50, 5493.25, 5493.00, 5492.75, 5492.50, 5492.25], + "bid_sz": [100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0], + "ask_px": [5494.75, 5495.00, 5495.25, 5495.50, 5495.75, 5496.00, 5496.25, 5496.50, 5496.75, 5497.00], + "ask_sz": [100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0] + }, + { + "type": "step_resting", + "ts_ns": 2000000000, + "trade_signed_vol": 0.0 + } + ], + "expected_pos": [ + { + "position_lots": 0, + "realized_pnl_tol": 50.0 + } + ], + "expected_stop_slot_states": [ + { "backtest_idx": 0, "slot_idx": 0, "active": 0 } + ], + "expected_min_trade_records": 1 +} diff --git a/crates/ml-backtesting/tests/fixtures/submission_overflow.json b/crates/ml-backtesting/tests/fixtures/submission_overflow.json new file mode 100644 index 000000000..f609b3924 --- /dev/null +++ b/crates/ml-backtesting/tests/fixtures/submission_overflow.json @@ -0,0 +1,22 @@ +{ + "name": "submission_overflow", + "n_backtests": 1, + "comment": "Fills 32 limit slots (the MAX_LIMITS cap), then verifies a 33rd seed attempt fails with overflow=true.", + "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": [100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.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": [100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0, 100.0] + } + ], + "fill_all_limit_slots": true, + "expect_33rd_overflow": true, + "expected_pos": [ + { + "position_lots": 0, + "realized_pnl_tol": 0.001 + } + ] +} diff --git a/crates/ml-backtesting/tests/lob_sim_fixtures.rs b/crates/ml-backtesting/tests/lob_sim_fixtures.rs index fdae6074a..95a724777 100644 --- a/crates/ml-backtesting/tests/lob_sim_fixtures.rs +++ b/crates/ml-backtesting/tests/lob_sim_fixtures.rs @@ -35,6 +35,41 @@ enum FixtureEvent { #[serde(default = "default_ts_ns")] ts_ns: u64, }, + SeedLimit { + backtest_idx: usize, + slot_idx: usize, + side: String, + kind: String, + active_state: u8, + oco_pair: u8, + price: f32, + size: f32, + #[serde(default)] queue_position_init: f32, + #[serde(default)] arrival_ts_ns: u64, + }, + SeedStop { + backtest_idx: usize, + slot_idx: usize, + side: String, + kind: String, + active_state: u8, + oco_pair: u8, + trigger_price: f32, + #[serde(default)] limit_price: f32, + size: f32, + #[serde(default)] arrival_ts_ns: u64, + }, + StepResting { + ts_ns: u64, + #[serde(default)] trade_signed_vol: f32, + }, +} + +#[derive(Debug, Deserialize)] +struct ExpectedSlotState { + backtest_idx: usize, + slot_idx: usize, + active: u8, } fn default_ts_ns() -> u64 { 1 } @@ -122,6 +157,37 @@ struct Fixture { expected_min_trade_records: usize, #[serde(default)] expected_isv_kelly_after: Vec<[ExpectedIsvKelly; 5]>, + #[serde(default)] + expected_limit_slot_states: Vec, + #[serde(default)] + expected_stop_slot_states: Vec, + #[serde(default)] + fill_all_limit_slots: bool, + #[serde(default)] + expect_33rd_overflow: bool, +} + +fn parse_kind_limit(s: &str) -> u8 { + match s { + "Limit" => 1, + "Ioc" => 2, + "Fok" => 3, + other => panic!("unknown limit kind: {other}"), + } +} +fn parse_kind_stop(s: &str) -> u8 { + match s { + "StopMarket" => 4, + "StopLimit" => 5, + other => panic!("unknown stop kind: {other}"), + } +} +fn parse_side(s: &str) -> u8 { + match s { + "buy" => 0, + "sell" => 1, + other => panic!("unknown side: {other}"), + } } fn run_book_fixture(path: &Path) -> Result<()> { @@ -169,6 +235,29 @@ fn run_book_fixture(path: &Path) -> Result<()> { sim.broadcast_alpha(alpha_probs)?; sim.step_decision(*ts_ns, *target_annual_vol_units, *annualisation_factor, *max_lots)?; } + FixtureEvent::SeedLimit { + backtest_idx, slot_idx, side, kind, active_state, + oco_pair, price, size, queue_position_init, arrival_ts_ns, + } => { + sim.seed_limit_order( + *backtest_idx, *slot_idx, + parse_side(side), parse_kind_limit(kind), *active_state, + *oco_pair, *price, *size, *queue_position_init, *arrival_ts_ns, + )?; + } + FixtureEvent::SeedStop { + backtest_idx, slot_idx, side, kind, active_state, + oco_pair, trigger_price, limit_price, size, arrival_ts_ns, + } => { + sim.seed_stop_order( + *backtest_idx, *slot_idx, + parse_side(side), parse_kind_stop(kind), *active_state, + *oco_pair, *trigger_price, *limit_price, *size, *arrival_ts_ns, + )?; + } + FixtureEvent::StepResting { ts_ns, trade_signed_vol } => { + sim.step_resting_orders(*ts_ns, *trade_signed_vol)?; + } } } @@ -182,6 +271,44 @@ fn run_book_fixture(path: &Path) -> Result<()> { ); } + if fx.fill_all_limit_slots { + // Seed all 32 limit slots successfully; expect each call to return Ok. + for slot in 0..ml_backtesting::lob::MAX_LIMITS { + sim.seed_limit_order( + 0, slot, 0 /*buy*/, 1 /*Limit*/, 1 /*resting*/, 0xFF, + 5400.00 /* deep below the book — never marketable */, + 1.0, 0.0, 0, + )?; + } + // The 33rd seed (re-seeding slot 0 which is in use) must fail. + if fx.expect_33rd_overflow { + let res = sim.seed_limit_order( + 0, 0, 0, 1, 1, 0xFF, 5400.00, 1.0, 0.0, 0, + ); + assert!( + res.is_err(), + "expected slot-already-in-use error, got Ok" + ); + } + } + + for want in &fx.expected_limit_slot_states { + let got = sim.read_limit_slot(want.backtest_idx, want.slot_idx)?; + assert_eq!( + got.active, want.active, + "limit_slot[b={},s={}].active got {} want {}", + want.backtest_idx, want.slot_idx, got.active, want.active + ); + } + for want in &fx.expected_stop_slot_states { + let got = sim.read_stop_slot(want.backtest_idx, want.slot_idx)?; + assert_eq!( + got.active, want.active, + "stop_slot[b={},s={}].active got {} want {}", + want.backtest_idx, want.slot_idx, got.active, want.active + ); + } + if !fx.expected_isv_kelly_after.is_empty() { for (b, expected) in fx.expected_isv_kelly_after.iter().enumerate() { let got = sim.read_isv_kelly(b)?; @@ -312,3 +439,27 @@ fn fix_pnl_accounting_buy_close() { fn fix_decision_alpha_buy_close() { run_book_fixture(Path::new("tests/fixtures/decision_alpha_buy_close.json")).unwrap(); } + +#[test] +#[ignore = "requires CUDA"] +fn fix_limit_rest_marketable_fill() { + run_book_fixture(Path::new("tests/fixtures/limit_rest_marketable_fill.json")).unwrap(); +} + +#[test] +#[ignore = "requires CUDA"] +fn fix_stop_trigger() { + run_book_fixture(Path::new("tests/fixtures/stop_trigger.json")).unwrap(); +} + +#[test] +#[ignore = "requires CUDA"] +fn fix_oco_one_cancels_other() { + run_book_fixture(Path::new("tests/fixtures/oco_one_cancels_other.json")).unwrap(); +} + +#[test] +#[ignore = "requires CUDA"] +fn fix_submission_overflow() { + run_book_fixture(Path::new("tests/fixtures/submission_overflow.json")).unwrap(); +}