Existing lob_sim_fuzz only exercised book_update + market orders.
This new test suite drives the FULL integrated pipeline under random
adversarial conditions:
apply_snapshot (random-walk book)
→ step_resting_orders (random signed trade-flow signal)
→ broadcast_alpha (random per-horizon probs every 4th event)
→ step_decision_with_latency (mixed latency=0 + latency=50ms cells)
→ submit_market_immediate (immediate path)
→ pnl_track_step + isv_kelly_update_on_close
Warm-starts every backtest with random-but-plausible Kelly state
(at least h4 set credibly profitable so decisions actually open
positions). Random target_annual_vol + annualisation_factor + max_lots
per decision to vary the Kelly cap.
Per-50-event invariants:
• book monotonicity (bid_px[k] ≤ bid_px[k-1], ask_px[k] ≥ ask_px[k-1])
• no-crossed-book (ask[0] > bid[0])
• Pos.realized_pnl + Pos.vwap_entry finite (no NaN/Inf leaks)
• Pos.position_lots in plausible range (|≤100|)
• All 5 per-horizon IsvKellyState fields finite
Three test sizes:
integrated_fuzz_n1_short — N=1, 200 events
integrated_fuzz_n8_medium — N=8, 500 events
integrated_fuzz_n64_long — N=64, 1000 events
All three pass on RTX 3050. The N=64 × 1000 case exercises 64,000
event-snapshots × 16,000 decisions × ~12,800 trade attempts without
any invariant violation or NaN propagation.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Closes the spec §8 Ring 3 deferral. Two integration tests against
FOXHUNT_TEST_DATA real MBP-10 day files:
buy_and_hold_full_day
Walk the first .dbn.zst from open to close; submit a market buy
of 1 lot at the first event, run the full step-loop
(apply_snapshot + step_resting_orders) over every subsequent
event, then close with a market sell at the last event. Assert
realized P&L matches (close_mid - open_mid) within ±2 ticks per
spec slippage budget. Proves the book-walking + position
accounting + step-loop orchestration are wire-level correct
against real data, not just hand-crafted JSON fixtures.
walk_book_one_full_day
Same fixture, no trades — just iterates every event through
apply_snapshot + step_resting_orders and asserts the book
invariants (bid/ask monotonicity, no-crossed-book) hold at every
10_000-event checkpoint across the full day. Stress test of
the kernel against real market microstructure (gaps, locked
markets, regime shifts).
Both tests skip gracefully when FOXHUNT_TEST_DATA is absent OR the
predecoded sidecars are empty (the current 40-byte placeholder
fixtures in test_data/futures-baseline/ES.FUT/*.predecoded.bin will
trigger the skip path; populating those with a real Databento decode
makes both tests fire end-to-end). Ring 3 is operational, not
CI-mandatory.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Replaces the placeholder `trade_signed_vol = trade_count_delta` (always
non-negative, sign-neutral) with a proper Databento-standard tick-rule
inference applied to L1 size + price deltas across consecutive MBP-10
snapshots:
• ask_px[0] unchanged AND ask_sz[0] decreased → aggressive buys
consumed ask depth; add (prev.ask_sz − cur.ask_sz).
• bid_px[0] unchanged AND bid_sz[0] decreased → aggressive sells hit
the bid; subtract (prev.bid_sz − cur.bid_sz).
• ask_px[0] moved up → previous best ask cleared; add prev.ask_sz.
• bid_px[0] moved down → previous best bid hit; subtract prev.bid_sz.
• Pure cancellations (size shrank but price moved AWAY from us) =
ambiguous; ignore.
Convention matches `Mbp10RawInput::trade_signed_vol`: positive =
buyer-initiated, negative = seller-initiated. This is a LOWER-BOUND
estimator — won't catch trades that cleared multiple levels (those
manifest only via deeper-level deltas) or trades against hidden /
off-book liquidity. Acceptable for v1 queue-decay signal; production
deployments can layer the trades-stream loader for ground-truth flow.
Wired into BacktestHarness::run() → sim.step_resting_orders(ts, vol)
so the queue-decay branch of resting_orders.cu finally fires with
non-zero input. Previously the harness passed 0.0 unconditionally,
which meant resting limits could only fill via the price-cross
marketability branch — same-price queue-decay was inert.
Six new unit tests cover each branch of the inference (pure cancel,
ask-shrank, bid-shrank, ask-px-up, bid-px-down, mixed both-sides).
34 ml-alpha lib tests + 33 ml-backtesting lib + 12 GPU fixtures + 3
fuzz still green.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The hardcoded WeightedByRealizedSharpe path from C7 is now joined by
a stack-based bytecode interpreter that consumes Strategy::flatten()
output, unlocking RegimeSwitch / Portfolio / non-default Ensemble /
single-Leaf compositions specified via policy-grid YAML.
cuda/decision_policy.cu — new kernel `decision_policy_program`:
Stack-based VM with parallel (value, attribution_mask) stacks.
Opcodes mirror src/policy/mod.rs::OpCode exactly:
NoOp / PushScalar / EvalRegime / BranchIfRegime
EmitPerHorizonSize (computes sized intent from alpha[h] +
IsvKellyState[h] using same Kelly + ISV-cap formula as the
hardcoded default)
AggMean / AggWeightedSharpe / AggMaxConfidence (pop n values,
push aggregated, OR attribution masks)
ApplyConflict (v1 no-op, reserved for Portfolio)
WriteOrder (terminal — converts top-of-stack to market_target +
open_horizon_masks attribution if currently flat)
AggWeightedSharpe recovers the source horizon from a single-bit
attribution mask to look up recent_sharpe; multi-bit masks (nested
aggregators that collapsed horizons) fall back to uniform weight.
decision_policy_default extended with a program_lens param: skips any
backtest whose plen > 0 (the program kernel handled it). The two
kernels run sequentially in step_decision_with_latency with mutual
exclusivity on each backtest slot.
LobSimCuda gains:
upload_program(b, &Program) — uploads a Strategy::flatten() output
to backtest b's slot in program_table_d, updates program_lens_d.
set_regime(b, regime_id) — writes regimes_d for OP_EVAL_REGIME /
OP_BRANCH_IF_REGIME consumption.
BacktestHarnessConfig.strategies (Vec<Strategy>) — empty means every
cell uses the hardcoded default; non-empty len must equal n_parallel
and each strategy is flattened + uploaded at construction.
bin/fxt-backtest --policy-grid <yaml> path now actually plumbs through
to the kernel (was parsed-but-ignored in C9). Empty grid keeps the
default behaviour.
New fixture decision_program_h4_only: uploads Strategy::Leaf(h4_only)
flattened to (EmitPerHorizonSize, WriteOrder) — 2 instructions, 16
bytes — and verifies the bytecode kernel produces equivalent end-state
to the hardcoded default (3 lot buy at vwap=5500). Proves the VM
dispatcher works end-to-end.
12/12 GPU fixtures green. 33 lib unit tests + 3 Ring 2 fuzz tests
still green.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The in-flight machinery from C11 is now reachable from the harness +
CLI. New step_decision_with_latency on LobSimCuda:
latency_ns == 0 → existing immediate-match path (submit_market_immediate
kernel fills against current book).
latency_ns > 0 → host reads market_targets, converts each non-noop
into seed_limit_order(active=2, price=very-aggressive,
arrival_ts_ns = current + latency_ns). The
resting_orders kernel promotes + fills at arrival
time, so the order sees whatever book exists then —
not the book at decision time.
Aggressive-price heuristic: buy at 1e9 / sell at 0 — the marketability
check (price ≥ best ask for buy, ≤ best bid for sell) unconditionally
crosses at arrival; the kernel then walks book levels for the actual
fill price. This models "market order with latency" correctly because
slippage emerges from the book-walking at arrival, not from a limit
price boundary.
BacktestHarness gains a latency_ns field; harness::run() always calls
step_decision_with_latency now. Also calls sim.step_resting_orders
per event (with trade_signed_vol=0.0 — the Mbp10RawInput trade-flow
hookup is deferred to a trades-feed integration commit) so in-flight
orders get a chance to promote on every snapshot.
bin/fxt-backtest no longer ignores --latency-ns; the flag value
propagates through BacktestHarnessConfig.latency_ns into the kernel
path. Default stays at 100_000_000 (IBKR + Scaleway baseline per
spec §4). Setting --latency-ns 0 selects the legacy immediate path.
Adds latency_in_flight_miss GPU fixture: limit at 5510 submitted
active=2 at T=1s with arrival_ts=T+100ms. step_resting at T+50ms
sees no promotion (still in-flight). Book moves +5 against buyer
during the 100ms window. step_resting at T+150ms promotes the
limit and fills it at the WORSE post-move book (vwap=5505 vs the
original 5500 it would have hit without latency). Slot ends active=0.
11/11 GPU fixtures green on RTX 3050. 33 lib tests still green.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
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 <noreply@anthropic.com>
Property-based fuzz tests with random-walk MBP-10 sequences applied to
the LOB simulator at three backtest counts; assert per-snapshot
invariants that must hold regardless of input or block scheduling:
fuzz_n1_book_only (200 events, no orders)
Pure book-update kernel — verifies mid-drift random walks preserve
book monotonicity and never produce a crossed book.
fuzz_n16_with_orders (200 events, market orders every 8th event)
16 backtests in parallel, each submitting random buy/sell market
orders 1-3 lots at every 8th event. Asserts book invariants on
each backtest's state PLUS:
- position_lots stays within ±30 (plausible given fixture book depth)
- realized_pnl + vwap_entry finite (no NaN/Inf leaks)
fuzz_n256_with_orders (100 events, market orders)
Production-scale parallelism. Same invariants as N=16. Each block
has its own per-backtest Pos + OpenTradeState + TradeLog, exercising
the per-block isolation discipline established in C5-C7.
All 3 pass on RTX 3050. Spec §8 Ring 2 confidence gate hit.
Adds rand + rand_chacha to ml-backtesting dev-dependencies.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
artifacts.rs:
- Summary struct (total_pnl_usd, sharpe_ann, sortino_ann,
max_drawdown_usd, calmar, n_trades, win_rate, avg_win/avg_loss,
profit_factor, total_fees_usd, exposure_pct,
kelly_cap_history_sample).
- compute_summary(records, pnl_curve_usd) — non-overlapping
annualisation × √825 per pearl_phase1d4_backtest_cost_edge_frontier
(K=6000 holding × 250 trading days ≈ 825 trades/year).
Sharpe + Sortino + max drawdown + Calmar.
- write_summary (JSON pretty-printed), write_trades_csv (with USD
conversion from fp ×100), write_pnl_curve_bin (bytemuck-cast f32
slice). 5 unit tests with tempdir.
aggregate.rs:
- aggregate_sweep_dir walks <root>/<cell>/summary.json, builds an
arrow RecordBatch (cell name + 9 stats columns), writes
SNAPPY-compressed aggregate.parquet.
- pareto_frontier: cells are kept unless another cell weakly
dominates on all three of (sharpe_ann maxed, max_drawdown_usd
minimised, total_fees_usd minimised) AND strictly improves on one.
Written to pareto_frontier.json (Vec<cell-name>).
- 2 unit tests (3-cell mutual-non-dominance; B-dominates-A).
harness.rs:
- run() now samples Pos.realized_pnl × $50/index-pt per event into
self.pnl_curves[b], so the per-cell P&L curve is ready for
write_artifacts() without an extra sim pass.
- write_artifacts(out_dir) — per-cell <out>/cell_NNNN/{summary.json,
trades.csv, pnl_curve.bin}.
bin/fxt-backtest:
- clap-derive CLI with two subcommands:
run --data <dir> [--predecoded-dir <dir>] [--policy-grid <yaml>]
[--n-parallel N] [--decision-stride S] [--latency-ns N]
[--target-annual-vol-units F] [--annualisation-factor F]
[--max-lots N] [--max-events N] [--seed N] --out <dir>
aggregate <sweep_dir>
- Constructs MlDevice::cuda(0) + CfcTrunk::new_random for the trunk
(v1 — ml-alpha has no checkpoint format yet; --seed gates init).
- Parses --policy-grid YAML if given but doesn't yet plumb to the
LobSimCuda decision kernel (the v1 kernel hardcodes the
Strategy::default_for path; bytecode VM is C7's deferred follow-up).
Parse step kept end-to-end so the YAML format is validated now.
- --latency-ns parsed but not consumed — reserved for follow-up
resting-order in-flight promotion (deferred from C5).
Adds parquet + arrow + arrow-array + arrow-schema + serde_yaml to
ml-backtesting deps; bin/fxt-backtest added to workspace members.
All 33 lib tests + 6 GPU fixture tests green. CLI --help renders both
subcommands correctly.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
BacktestHarness::new(cfg, &dev, trunk) constructs a MultiHorizonLoader
(inference_only=true), captures the trunk's perception Graph A using
loader.peek_first() as the template, then allocates a LobSimCuda.
run() walks the chronological snapshot stream, calls apply_snapshot on
every event, and at decision-stride boundaries:
trunk.update_input_buffers(raw)
→ trunk.perception_forward_captured() → (probs[N_HORIZONS], proj)
→ sim.broadcast_alpha(&probs)
→ sim.step_decision(ts, target_vol, ann_factor, max_lots)
Returns RunStats { events_processed, decisions_taken }. The harness
deliberately accepts an externally-built CfcTrunk (random-init in v1)
because a checkpoint format isn't pinned in ml-alpha yet; when one
lands, the binary CLI (C9) can switch from new_random to load_checkpoint.
trainer_parity.rs Ring 1b — two ignored tests:
peek_first_byte_equal_across_modes
Verifies the Mbp10RawInput produced by the loader path used by the
backtest harness (inference_only=true) is BYTE-EQUAL to what the
trainer's loader (inference_only=false) produces from the same
source. Guards against any future refactor accidentally diverging
the two paths (e.g. someone special-casing inference path to skip
regime feature computation). All 50 f32 fields + scalars compared
via .to_bits() equality.
inference_iteration_matches_chronological_snapshots
Verifies next_inference_input() yields monotone-ts snapshots with
correct cur==prev semantics on the first read and prev_ts==prior_ts
afterwards.
Both tests skip gracefully when FOXHUNT_TEST_DATA fixtures lack
populated sidecars (the placeholder-empty bins committed in the
test_data/ tree).
GPU-side inference parity (probs[N_HORIZONS] bit-equal across loader
modes) is deferred until ml-alpha pins a checkpoint format and we have
a small checkpoint fixture to gate it on FOXHUNT_TEST_CKPT.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Two new kernels in cuda/decision_policy.cu:
decision_policy_default — alpha[N_HORIZONS] × per-horizon IsvKellyState
→ market target. Per-horizon Kelly fraction × signal magnitude × ISV-cap
(target_annual_vol / sqrt(realised_return_var × annualisation_factor)),
aggregated via WeightedByRealizedSharpe (weights = max(0, recent_sharpe)
/ Σ; auto-shifts capital toward empirically winning horizons per spec §5
+ §6). No floor — sub-1-lot intents become no-op. Round-to-nearest on
the final lot count to dodge an f32-truncation off-by-one where
0.8(f32) * 0.75 * 5.0 ≈ 2.99999976 → trunc-int 2.
isv_kelly_update_on_close — for every horizon flagged in
closed_horizon_mask[b], updates pnl_ema_{win,loss} via Wiener-α (floor
0.4 per pearl_wiener_alpha_floor_for_nonstationary), win_rate_ema,
Welford-ish realised_return_var, and the recent_sharpe composite.
First-observation bootstrap (per pearl_first_observation_bootstrap):
sentinel n_trades_seen=0 → direct EMA replacement, no zero-bias warmup.
The full bytecode VM from spec §6 is NOT in this commit — the default
policy is hardcoded in the kernel as the path Strategy::default_for()
already produces. Bytecode plumbing in src/policy/mod.rs stays put for
v2 expansion (custom RegimeSwitch / Portfolio compositions).
IsvKellyState struct added to lob_state.cuh (24 bytes per horizon × 5
per backtest); host mirror IsvKellyStateHost from C3 cast-compatible
via bytemuck::Pod. LobSimCuda gains broadcast_alpha + step_decision +
read_isv_kelly + write_isv_kelly (warm-start). step_decision chains:
decision_policy → merge_open_mask → submit_market_immediate
→ pnl_track_step → host close-detect → isv_kelly_update_on_close.
PRE-submit pos/pnl/mask snapshots feed the host close-detection;
captured via three small DtoH copies (cold path, 24 bytes × n_backtests).
decision_alpha_buy_close fixture: warm-start h4 with positive Kelly
state (n=50, recent_sharpe=1.0), broadcast alpha[4]=0.9 → buy 3 lots
@ ask top 5500.00. Snapshot moves to bid 5505.00, broadcast alpha[4]=0.1
→ sell 3 lots → close at +15 P&L. Verify ISV-Kelly h4: n_trades 50→51
exactly, others unchanged. PASS — all 6 Ring 1 fixtures green on RTX 3050.
Out-of-scope for this commit (defer to follow-ups, per plan trim notes):
- stop_trigger / oco_one_cancels_other / submission_overflow fixtures
(need resting-order LimitSlot[32] machinery deferred from C5)
- Bytecode VM dispatch (RegimeSwitch / Portfolio compositions)
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
pnl_track_step runs after each matching pass, compares per-block
position-state-now against persisted OpenTradeState (24 B scratch)
and either:
- records entry context (entry_ts_ns, entry_px_x100, entry_size,
realised_at_open) on open transition (prev==0, now!=0); or
- emits a 40-byte TradeRecord into the per-block trade-log buffer
on close transition (prev!=0, now==0), reconstructing implied
exit_px from the realized P&L delta and converting to USD ×100
fixed-point ($50/index-point × 100 = ×5000 multiplier).
Multi-fill averaging (scale-in then partial close) deferred to v2 —
v1 covers the clean open→close case the spec calls out as primary.
LobSimCuda owns three new buffers: open_trade_state_d (n × 24),
trade_log_d (n × TRADE_LOG_CAP × 40), trade_log_head_d (n × u32).
submit_market now takes current_ts_ns and chains pnl_track_step
internally; step_pnl_track() exposed for caller-driven orchestration.
read_trade_records(backtest_idx) drains the per-block ring as
Vec<TradeRecord>; LSP-pinned 40-byte Pod struct from C2 lines up
1:1 with the kernel's hand-rolled byte writes.
pnl_accounting_buy_close fixture: buy 4 lots @ ask top (5500.00),
book moves +5 to bid top 5505.00, sell 4 to close. Expected
realized_pnl = (5505 − 5500.00) × 4 = $20 in price-units, which is
$20 × $50/contract × 100 = 100000 USD ×100 fixed-point. PASS within
$1 fixed-point tolerance.
All 5 Ring 1 fixtures green on RTX 3050.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
submit_market_immediate walks the ask/bid book for a buy/sell market
order, computes notional cost + filled lots across up to 10 levels,
and updates per-backtest Pos { position_lots, vwap_entry, realized_pnl }
with correct VWAP scale-in + counter-direction realized-P&L accounting.
Single-writer (thread 0) per block — no atomics per
feedback_no_atomicadd.md. Each backtest gets its own target slot
{ side, size } in the device-global targets array (side=2 = no-op).
Pos struct added to lob_state.cuh (24 bytes); host mirror PosFlat in
src/lob/mod.rs is repr(C) bytemuck::Pod for direct host↔device cast.
LobSimCuda gains submit_market(backtest_idx, side, size) +
read_pos(backtest_idx) -> PosFlat. Fixture harness extended with the
SubmitMarket event variant + ExpectedPos check (vwap_tol + realized_pnl_tol
for FP tolerance).
market_order_consumes_top fixture verifies a buy of 5 lots into
ask[3@5500.00, 10@5500.25] fills 3+2 → position_lots=5,
vwap_entry=5500.10 within 1e-3 tolerance. All 4 Ring 1 fixtures
(book_update × 3 + market_order × 1) green on RTX 3050.
Scope deliberately trimmed from plan C5: this commit lands market-fill
matching only. Queue decay on resting limits, in-flight latency
promotion, stop-trigger detection, and OCO leg-cancellation will land
in follow-up commits — each in its own kernel addition. This keeps the
incremental delta reviewable.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
build.rs compiles every .cu under crates/ml-backtesting/cuda/ to
\$OUT_DIR/<name>.cubin via nvcc (no nvrtc per feedback_no_nvrtc.md);
pairs every env::var with rerun-if-env-changed per
pearl_build_rs_rerun_if_env_changed.md. Default arch sm_86 (RTX 3050
local); production sets FOXHUNT_CUDA_ARCH=sm_89 (L40S) or sm_90 (H100).
First kernel: book_update_apply_snapshot — block-per-backtest replace
of the 10-level book from a broadcast MBP-10 snapshot. Single-writer
per level (thread tid = level idx), no atomics per feedback_no_atomicadd.md.
lob_state.cuh defines the canonical per-block shared-mem layout that
all subsequent kernels share (Book struct in this commit; Orders, Pos,
IsvKellyState[5] added in C5-C7).
LobSimCuda host struct (in src/sim.rs) constructed from an &MlDevice;
owns per-backtest device book state + mapped-pinned input snapshots.
apply_snapshot launches the kernel and synchronises; read_books drains
device state for tests.
Three Ring 1 fixture tests (book_update_replace, _two_step,
_n_backtests=4) — N≤4 bit-exact GPU oracle assertions loaded from JSON.
All three pass on RTX 3050 locally. Verifies the build-script → cubin →
cudarc.load_cubin → kernel launch → DtoH read pipeline end-to-end.
cudarc added as direct path dep (../../vendor/cudarc) since it's not in
[workspace.dependencies] — same vendored fork ml-alpha uses.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Strategy (recursive Leaf | Ensemble | Portfolio | RegimeSwitch) +
StrategyConfig (horizon_idx + sizing_policy + sl_tp_rules +
max_concurrent_lots). Strategy::default_for(max_lots) returns the v1
default: per-horizon adaptive Ensemble with WeightedByRealizedSharpe
aggregator. No static horizon mask — capital auto-shifts via per-horizon
ISV-Kelly recent_sharpe weights (see spec §5 + §6).
Strategy::flatten() walks the tree → linear bytecode Program
(repr(C) Pod Instruction { op, arg0, arg1, arg2 }, 8 bytes/instruction)
for upload to device-global memory at slot blockIdx.x. RegimeSwitch
emits BranchIfRegime chains with arg2 patched to the post-child jump
offset.
Tests cover: 5-horizon default tree shape, flatten output for default
+ single-leaf + 2-leaf Portfolio (verifying ApplyConflict op emission),
Instruction layout size pin (8 bytes), LatencyConfig default 100ms
(IBKR+Scaleway baseline), Empirical wrap-around + empty case,
Decomposed summation.
IsvKellyStateHost host-mirror of cuda IsvKellyState (24 bytes, Pod)
defined in policy::sizing for warm-start file I/O in later commits.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
OrderIntent (ergonomic host-side enum for policy YAML config + audit
reconstruction), OrderEvent (24-byte repr(C) Pod for the device-side
audit ring), TradeRecord (40-byte repr(C) Pod for the device-side
trade-log buffer), SlotTag (u32-packed {SlotKind, index, generation}
for cancel/modify against slot reuse), plus Side / LimitKind /
SlotKind / LimitParams.
Size pinning tests (24/40 bytes) ensure any future field-add forces
the kernel-side struct to migrate in lockstep. Serde roundtrips cover
OCO + Cancel variants. Renamed `.gen()` accessor to `.generation()`
to dodge the Rust 2024 reserved-keyword clash.
OCO is two limit legs linked by oco_pair byte on-device (not recursive
in OrderIntent) per spec §3 — keeps device-side layout flat for the
matching kernel coming in C5.
Adds ml-alpha + bytemuck + serde_json + tracing + thiserror to deps.
cudarc + parquet + arrow added later when CUDA + sweep aggregator
land (C4 / C9). No CUDA dependency in this commit.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
- Add backticks to type names in doc comments (doc_markdown)
- Mark eligible functions as const fn (missing_const_for_fn)
No behavior changes.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
module_name_repetitions: PpoConfig in ppo module is conventional ML naming.
Renaming breaks every import across the workspace.
integer_division: Basis point calculations, batch size math, combinatorial
formulas — truncation is intentional. Float conversion would introduce bugs.
Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
