Boundary test that fails on bit-pick-first or bit-pick-max regressions
of the open_horizon_masks averaging in stop_check_isv. Per-horizon
pnl_ema_loss values (2, 4, 3, 3, 3) yield mean=3.0 across the 0x1F
mask; snapshots placed relative to vwap_entry for ask-spread safety.
Tests Δ_entry=2.5 (no-fire) and Δ_entry=4.5 (fire).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
trail_distance = max(pnl_ema_win, atr_mid_ema). HWM ratchets up via
fmaxf each event while position open; trail fires when HWM clears
the arming threshold AND unrealized drops by trail_distance from
HWM. Same force-flat (3, 0) write as SL.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Skeleton helper called from decision_policy_default at top of
per-backtest dispatch. Returns 0 (no-op) on flat positions; trigger
logic for open positions added in Tasks 4-6. Single source of truth
for stop logic across decision_policy_{default,program}.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
cudarc kernel launches require `unsafe` blocks — the driver API has no
way to type-check kernel args against the cubin signature. The
workspace-wide `-W unsafe-code` lint produces noise on every launch
site (10+ blocks); suppressing at the module level keeps the lint
useful elsewhere in the crate.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per-event Wiener-α=0.4 EMA on |Δmid| with first-observation bootstrap.
Floor source for the SL/trail-distance controller (spec §6). Thread 0
of the broadcast-snapshot kernel handles the update per backtest;
threads 1..9 still handle the 10 level writes.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Three CudaSlice<f32> per-backtest slots (prev_mid_d, atr_mid_ema_d,
trail_hwm_d) plus test-only accessors. Foundation for the
ISV-driven stop controller; no behavioral change until subsequent
tasks wire them into kernels.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Adds LobSimCuda::read_total_trade_count (cheap n_backtests*4 byte DtoH of
the trade-log head counters) and emits the sum on each PROGRESS_EVERY
eprintln. Lets future smokes (Q2 CBSW validation, P7 sweep) see trade
firing mid-run instead of waiting for write_artifacts at harness end.
Heads count kernel writes (clamped to ring cap), so the sum is a
monotone trade counter across the sweep — drops back only on harness
reset which doesn't happen mid-run.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The threshold-tuning smoke at 81decf40f produced n_trades=0 despite
74.6% of decisions having max_conv ≥ 0.30 — the linear-weighted-mean
aggregator in decision_policy_default is structurally dilution-bound
at cold-start (per spec §1).
Q1 stopgap: when sim_variants[i].use_cold_start_stopgap = true, the
harness uploads a max-confidence Strategy bytecode program for that
backtest, routing decisions through decision_policy_program with
OP_AGG_MAX_CONFIDENCE. Existing kernel; zero CUDA changes.
Field additions (atomically across BatchedSimConfig + UniformSimParams
+ ResolvedSimVariant + SweepBase.SimVariant) — every UniformSimParams
literal migrated to include use_cold_start_stopgap: false (default).
The sweep YAML's sim_variants entry sets it to true only for the
validation run; production deployability uses Q2's kernel fix instead.
Sweep YAML (config/ml/sweep_smoke.yaml) flipped to use_cold_start_stopgap=true
at threshold=0.0, cost=0.125 — same anchor as the threshold-tuning
smoke that produced n_trades=0, for direct comparison.
This is a VALIDATION step. Cluster smoke at this commit MUST produce
n_trades > 100 + finite metrics. Q2's kernel CBSW immediately follows
and deletes this entire stopgap atomically (field, harness branch,
YAML setting, every literal).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
P4 plumbed the threshold-gate kernel side but deferred the side-channel
that captures observed max_conviction per decision. Wire it now so the
threshold pre-registration step (spec §3.4) can compute the calibrated
p60-p95 absolute threshold values from a real model-on-data run.
Harness changes:
- BacktestHarness gains conviction_log: Vec<f32>. Per decision, computes
max_h |alpha[h] - 0.5| * 2 from the SAME probs that go into broadcast_
alpha (same value the threshold gate would compare against), pushes
to the log. One shared vec — batched cells broadcast the same probs
to every backtest, so per-backtest is redundant.
- write_artifacts emits convictions.bin (raw little-endian f32) +
conviction_percentiles.json with pre-computed p10/p25/p50/p60/p70/
p80/p90/p95/p99 + mean/min/max. Also eprintln-prints the summary
line for at-a-glance log inspection.
Smoke YAML switched to the threshold-tuning configuration: threshold=0
(no gate, full distribution captured), cost=0.125 (1-tick realistic
anchor so the observed Sharpe is the no-gate net-of-cost floor for
the sweep's deployability story).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Legacy fan-out writes <sweep-dir>/<cell>/summary.json (one cell per
Argo task). P6 batched flow writes <sweep-dir>/<cell>/sim_<variant>/
summary.json (one Argo task → run_batched_cell → harness with
variant_names → sim_<name>/ subdirs per spec §3.3).
The aggregator was looking only at <sweep-dir>/<cell>/summary.json,
so the realistic batched smoke completed the actual backtest fine
(2M events, 500k decisions, real artifacts written) but the
end-of-sweep aggregate step errored with "no cell directories with
summary.json".
Walk both layouts: directories containing summary.json directly are
flat cells (legacy); directories one level deeper that contain
summary.json are batched-cell variants. Cell labels become
"<cell>/<variant>" so the aggregate.parquet rows distinguish.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
snapshot_realized_pnl / snapshot_position_lots / snapshot_open_horizon_mask
were the host-side read paths for the close-detection loop that P3
(3836e2578) replaced with snapshot_pos_state + detect_close_transitions_
batched kernels. The helpers stayed dead-code-warned after P3; per
feedback_no_legacy_aliases + feedback_no_hiding, delete them.
submit_market_fn is NOT dead — pub fn submit_market wraps it and has 4
test callers (lob_sim_fixtures, lob_sim_fuzz, ring3_replay × 2). Kept.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Sweep YAML now supports the batched flow per spec §3.3 + Task 6:
- SweepBase.sim_variants: Vec<SimVariant> — list of (cost, latency,
threshold, ...) variants. When non-empty, each cell runs ONE harness
at n_parallel=variants.len() with BatchedSimConfig::from_grid instead
of the legacy one-harness-per-cell fan-out.
- SweepBase.data_template: Option<String> — when set with `{window}`
placeholder, each cell's `window` field interpolates the per-cell
data path. Replaces single scalar `data` for the windowed flow.
- SweepCell.window: Option<String> — window identifier (e.g., "2025-Q2").
- SimVariant: threshold + cost_per_lot_per_side required (the spec's
primary axes); other fields optional overrides on top of SweepBase
scalars.
New runner pieces:
- BatchedSimConfig::from_grid(&[ResolvedSimVariant]) in
crates/ml-backtesting/src/sim/batched_config.rs.
- ResolvedSimVariant — per-variant fully-resolved sim params.
- resolve_sim_variants(&SweepBase) in main.rs — layers per-variant
overrides over base scalars.
- run_batched_cell() in main.rs — builds the harness with
sim_config_override + variant_names plumbed through. Writes per-
backtest artifacts to sim_<variant_name>/ subdirs (spec §3.3).
Harness side:
- BacktestHarnessConfig gains variant_names + sim_config_override
Option fields. When sim_config_override is Some, harness uses that
directly instead of building from_uniform off scalar cfg. When
variant_names is Some, write_artifacts uses sim_<name>/ instead of
cell_NNNN/ subdirs. Both None preserve legacy single-cell behaviour
(smoke, fixtures unchanged).
YAML configs:
- config/ml/sweep_threshold_tuning.yaml: 1 cell (W0) × 8 sim_variants
(p60-p95 in 5pt steps) with cost=0.125 (1-tick anchor). Threshold
pre-registration pass.
- config/ml/sweep_deployability.yaml: 4 cells (W1-W4) × 140 variants
each (7 costs × 4 latencies × 5 thresholds). Generated by
scripts/generate_sweep_variants.py — placeholder threshold values
(p60-p95) until threshold-tuning publishes calibrated absolutes to
config/ml/v2_prod_thresholds.json.
Deferred to P7 (operational glue):
- argo-lob-sweep.sh adaptation for the batched flow (cells = windows,
not sim-variants; one Argo task per window invokes `fxt-backtest sweep`
end-to-end inside the pod rather than `fxt-backtest run`).
Regression: all 7 existing CUDA tests pass through the new harness
construction path (sim_config_override = None → from_uniform fallback).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Adds the two sweep axes that the spec's deployability grid needs but
were missing from the kernels:
Threshold gate (decision_policy.cu, both kernels):
- New per-backtest `threshold_per_b` array kernel arg.
- Pre-Kelly prelude: if max_h |alpha[h] - 0.5| * 2 < threshold[b],
emit noop and return. Kept deterministic from alpha alone so the
threshold pre-registration step (p60-p95 absolute calibration on a
validation window, future P6) reflects exactly what gets gated in
deployment.
Per-fill cost integration (resting_orders.cu / apply_fill_to_pos):
- apply_fill_to_pos signature grows three args: b, cost_per_lot_per_side_per_b,
total_fees_per_b. Single insertion point at line 90.
- After the close-leg realized_pnl math runs (so the gross unwind P&L
is preserved), deduct fill_cost = filled_lots * cost_per_lot_per_side[b]
from pos.realized_pnl AND accumulate into total_fees_per_b[b].
- Net-of-cost semantics: isv_kelly_update_on_close reads realized_pnl
delta which is now net of cost — Kelly state learns from realistic
return distribution.
- All 3 apply_fill_to_pos call sites in step_resting_orders updated.
order_match.cu's submit_market_immediate path is dead code in the
post-P1 flow (everything routes through seed_inflight_limits_batched
→ step_resting_orders → apply_fill_to_pos) so not touched here.
BatchedSimConfig + UniformSimParams + BacktestHarnessConfig gain
threshold + cost_per_lot_per_side fields. All UniformSimParams
constructors in tests and main.rs updated with defaults (0.0, 0.0 =
gate disabled, frictionless).
Regression:
- threshold_gate_skips_low_conviction (p=0.51 + threshold=0.10 → noop)
- threshold_gate_allows_high_conviction (p=0.8 + threshold=0.10 → buy 1+)
- threshold_zero_is_passthrough (sanity)
- All P1+P2+P3 tests continue to pass via the new ABI.
cost_deducted_at_each_fill + kelly_state_sees_net_return end-to-end
tests deferred — they require a full submit_market → fill → close
sequence, which the production smoke exercises.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Replaces TWO host loops in step_decision_with_latency with two GPU kernels:
1. snapshot_pos_state — replaces the host-side snapshot_realized_pnl +
snapshot_position_lots + snapshot_open_horizon_mask trio (3 separate
memcpy_dtoh per decision). Now one kernel launch writes
prev_pos_lots_d / prev_realized_pnl_d / prev_open_horizon_mask_d
directly on the device.
2. detect_close_transitions_batched — replaces the host close-detect
loop that called read_pos per close-eligible backtest (up to
n_backtests memcpy_dtoh per decision). Now one kernel writes
closed_horizon_mask_d + realised_return_d on the device, and
isv_kelly_update_on_close consumes them with no host roundtrip.
At n_parallel=140 these two loops together accounted for ~350M+ small
host roundtrips per quarter. Combined with P2 the latency-path of
step_decision_with_latency is now fully GPU-resident.
isv_kelly_update_on_close kernel always launches (skips backtests with
mask=0 internally) rather than gating via a host any_close check.
All P1+P2 regression tests pass through the new GPU close-detect path
(at n=1 with uniform config + immediate-fill latency, no close happens
in the cold-start tests since they only check market_target; the close
path is exercised indirectly).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Replaces the host roundtrip loop at the latency path of step_decision_
with_latency with one GPU kernel launch. At n_parallel=140 the host
loop did up to 140 memcpy_dtoh + 140 seed_limit_order calls per
decision (~210M roundtrips per quarter at the threshold-tuning load).
The new kernel does the same work in one launch.
Per-backtest single-writer (threadIdx.x==0). Each backtest scans its
own MAX_LIMITS=32 slot range for an `active==0` slot. Slot allocation
is per-backtest (no cross-backtest atomics needed). Overflow path
increments pos.submission_overflow.
dispatch_latent_market_orders now takes &BatchedSimConfig (unused
inside — the latency_ns_d device buffer is already populated by
step_decision_with_latency's upload block from P1).
All 3 decision_floor_coldstart tests still pass via the new GPU path
(at latency_ns=0 the in-flight slot's arrival_ts==current_ts and gets
promoted on the next step_resting_orders, functionally equivalent to
the legacy submit_market_immediate kernel).
Independence test (different per-backtest latencies → different
arrival_ts) deferred to a later step where a read_first_inflight_arrival_ts
helper is added.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Migrates target_annual_vol_units, annualisation_factor, max_lots,
latency_ns, kelly_frac_floor, sharpe_weight_floor from scalar-broadcast
kernel args to per-backtest device arrays via BatchedSimConfig. Atomic
contract change per feedback_no_partial_refactor — kernel + sim + harness
+ all 3 existing test files migrate in this commit.
- crates/ml-backtesting/src/sim.rs → sim/mod.rs (directory module)
- crates/ml-backtesting/src/sim/batched_config.rs (NEW): BatchedSimConfig
+ UniformSimParams + validate(). from_uniform rebuilds the legacy
uniform-broadcast behaviour at n=1 (smoke/fixtures). from_grid lands
in P6 for the 140-variant sweep packing.
- LobSimCuda gains 6 per-backtest device buffers (target_annual_vol_units_d,
annualisation_factor_d, max_lots_d, latency_ns_d, kelly_frac_floor_d,
sharpe_weight_floor_d). step_decision_with_latency uploads from
BatchedSimConfig each call; both decision kernel launches now pass
per-backtest array pointers.
- decision_policy_default + decision_policy_program: scalar args become
const float* / const int* per_b arrays; first lines of each kernel
index by `b` into the arrays. Behaviour preserved at n=1 uniform.
- dispatch_latent_market_orders: reads latency per-backtest from
&BatchedSimConfig (host loop stays for P1; P2 replaces with kernel).
- step_decision_with_latency now ALWAYS dispatches through the latency
path; when cfg.latency_ns[b]=0 the in-flight slot's arrival_ts equals
current_ts and gets promoted immediately on next snapshot. Eliminates
the if/else branch and consolidates the launch path.
- harness.rs: BacktestHarness gains a sim_config field, built via
BatchedSimConfig::from_uniform at new() from the harness cfg's scalar
fields. The run loop passes &self.sim_config to step_decision_with_latency.
Regression coverage:
- parallel_sim_correctness::parallel_sim_equivalence_with_uniform_config
— n=8 with uniform config produces 8 bit-identical market_targets
(proves per-backtest indexing reduces correctly).
- Existing decision_floor_coldstart tests (3) all pass through the new
ABI — proves cold-start floor + variance-cap gate behaviour preserved.
- parallel_sim_independence_per_backtest deferred to P2 (needs
read_first_inflight_arrival_ts helper that depends on LimitSlot layout).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
After widening the smoke to max_events=0 (full quarter), the inherited
decision_stride=1 became 10M+ forward passes — hours of GPU per cell,
indistinguishable from a deadlock in pod logs.
Two operational fixes:
- sweep_smoke.yaml: decision_stride 1 -> 4 (matches the sweep
template's own default).
- harness.rs: emit a `progress: events=... decisions=... elapsed=...
rate=...ev/s` line every 1M events on stderr so multi-million-event
cells are observable from kubectl logs.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Three follow-ups to the cold-start floor fix:
1. Kernel: MIN_TRADES_FOR_VAR_CAP gate. After the first trade closed,
`isv_kelly_update_on_close` set `realised_return_var = ret²` — a
single-sample variance proxy that systematically collapses
`cap_units = target_vol / sqrt(var × ann_factor)` near zero for
any biased return. cap_lots → 0 → no further trades despite strong
alpha. Gate the variance-derived cap behind `n_trades_seen >= 10`;
below the threshold cap falls back to host-supplied `max_lots`,
same as the pre-first-trade path. Same gate applied to both
`decision_policy_default` and `decision_policy_program`.
Regression: `post_first_loss_state_does_not_lock_out_further_trades`
reproduces the exact pre-fix state from the smoke (n_trades_seen=1,
var=103.6) and asserts the kernel still fires a long with p=0.8.
2. Aggregate: add `nvidia.com/gpu: 1` resource request. Scaleway's
L40S device plugin mounts libcuda.so.1 into the container only on
GPU-requesting pods; the aggregate logic is CPU-only but the
binary's dynamic loader needs the driver libs. Cheapest correct
fix until a separate CPU-only aggregator binary exists.
3. Smoke YAML: `max_events: 0` (exhaust loader). 100k events is
minutes of ES.FUT, far shorter than the h6000 holding horizon the
model was trained on. Full quarter exercises sustained trading +
variance estimate ramp-up.
All three regression tests pass locally on RTX 3050 Ti.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Production smoke completed end-to-end but produced n_trades=0 across 99,969
decisions — `decision_policy_default` and `decision_policy_program` both
applied a sentinel-skip pattern: if `isv_kelly_d` had not been seeded
(pnl_ema_win == 0), each horizon's signed-size stayed zero, AND each
horizon's aggregation weight (= recent_sharpe) also stayed zero. The
cross-horizon w_sum was therefore 0, final_size was 0, every market_target
was noop. State only updates on trade close → no trade ever fires →
infinite cold-start.
Per pearl_blend_formulas_must_have_permanent_floor (`max(real, floor)`,
not blend) and pearl_kelly_cap_signal_driven_floors, replace the sentinel-
skip with a two-layer floor on each kernel:
1. Kelly fraction: `max(kelly_frac_floor, computed_kelly)` — when state
is sentinel, falls back to the floor directly. Cap_lots falls back
to `max_lots` when realised_return_var is sentinel.
2. Aggregation weight: `max(sharpe_weight_floor, recent_sharpe)` — lets
cross-horizon sum produce a non-zero size before recent_sharpe is
populated. Once a horizon shows positive sharpe it dominates.
Plumbed through `step_decision_with_latency` / `step_decision` as two
new f32 args (atomic contract change, every caller migrated). Defaults
0.20 / 0.10 chosen so a strong-conviction signal (sig_mag ≥ 0.5) fires
1 lot at cold-start under max_lots=5 while weaker signals stay flat
(see `default_kelly_frac_floor` comment for the arithmetic). Exposed
as CLI flags + sweep-grid base/cell overrides.
Regression test `decision_floor_coldstart` proves:
- default floors (0.20/0.10) fire a 1-lot buy with p_h=0.8 and zero state
- zero floors reproduce the original noop bug
Also moves `aggregate` step to the GPU pool because fxt-backtest is
dynamically linked against libcuda.so.1 (the ci-compile-cpu hosts
don't expose CUDA driver libs).
Verified locally on RTX 3050 Ti — workspace cargo check passes, both
regression tests pass, trunk save/load roundtrip still passes.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
BacktestHarness now owns a PerceptionTrainer (in inference role) instead
of a raw CfcTrunk. The sliding K-window of recent snapshots accumulates
in the harness; at each decision-stride boundary (and only once the
window has reached cfg.seq_len), the harness calls
trainer.forward_only(&window) and broadcasts the last K position's
per-horizon probs to the LobSim.
fxt-backtest's main.rs constructs the trainer via
PerceptionTrainer::from_checkpoint when --checkpoint is supplied (else
random init for noise baseline).
Why this shape: PerceptionTrainer's evaluate_batched already runs the
full inference chain (snap → vsn → mamba2 → ln → mamba2 → ln →
attn_pool → cfc K-loop → grn heads) correctly. Duplicating that 400-line
forward chain on CfcTrunk would double the surface area for the same
result — the trunk's role is weight-source-of-truth (achieved in X1-X9),
not kernel-launch orchestration.
End-to-end status: alpha_train emits Checkpoint files via X14 wiring;
fxt-backtest now loads those Checkpoints via from_checkpoint and drives
forward via forward_only. Phase 2 (Argo runtime: training → smoke →
threshold pre-reg → 560-cell deployability sweep → verdict) is unblocked.
Adds PerceptionTrainer::config() accessor so the harness can read seq_len.
Verification: ml-alpha + ml-backtesting + fxt-backtest all build clean.
Per project_ml_alpha_starting_capital greenfield posture: there is no V1
to differentiate from (the V1 trunk forward was dead code, no V1
checkpoint files exist in the wild). The 'v2' prefix on every identifier
was historical baggage from the migration period.
Renames:
- CheckpointV2 -> Checkpoint (also drops the version: u32 field —
bincode either deserialises a current envelope or errors; no migration
path needed)
- CheckpointVersionProbe removed (was only for V1 rejection)
- LAYER_NORM_CUBIN_V2 / VARIABLE_SELECTION_CUBIN_V2 / ATTENTION_POOL_CUBIN_V2
-> LAYER_NORM_CUBIN / VARIABLE_SELECTION_CUBIN / ATTENTION_POOL_CUBIN
- _ln_module_v2 / _vsn_module_v2 / _attn_module_v2 -> drop _v2 suffix
- smoke_load_v2_checkpoint test -> smoke_load_checkpoint
- config/ml/sweep_v2_*.yaml -> config/ml/sweep_*.yaml
- migration-era 'V2 weight skeleton' / 'V2 fields' / etc. comments
cleaned to remove the v2 prefix
Pre-existing 'v2' references in ml-backtesting CUDA files
(decision_policy.cu, pnl_track.cu) are NOT touched — those refer to
future planned 'v2' refinements (Portfolio mode, multi-fill averaging)
from the C1-C19 commits and reflect aspirational features unrelated to
this session's trunk-grows work.
Verification: ml-alpha + ml-backtesting + fxt-backtest all build clean.
perception_forward_golden bit-exact (max_diff = 0.000000).
Adds #[ignore]d integration test that loads a CheckpointV2 file produced
by alpha_train and verifies CfcTrunk::load_checkpoint accepts it.
Activated via FOXHUNT_SMOKE_CKPT env var on a CUDA-capable host.
Also adds a compile-time witness test confirming Summary.max_drawdown_pct
field exists post-X16 (required by emit_deployability_verdict).
Note: full BacktestHarness end-to-end smoke (running one cell against
fixture MBP-10) requires the v2 forward path to live on the trunk
(X11 — deferred). This commit verifies the producer/consumer wire-up.
Per spec §3.4, §4.1 (X18).
X16: Adds Summary.max_drawdown_pct field as |max_drawdown_usd| /
STARTING_CAPITAL_USD (pinned at $35k per project_ml_alpha_starting_capital
memory — realistic ES single-contract small-account anchor). Used by
the verdict emitter as the capital-deployability hard gate (median
across windows < 20%).
X17: Adds VerdictTier enum (Pass-robust / Pass-nominal / Fail-inconclusive
/ Fail / Fail-degenerate), AnchorSpec / AnchorReport / DeployabilityVerdict
types, classify_verdict (tiered logic per spec §3.5), and
emit_deployability_verdict that reads per-cell summary.json files at
both realistic (1.0 tick, 200 ms) and stress (1.5 tick, 400 ms) anchors,
computes median Sharpe / Sortino / max_dd_pct / profit_factor across
walk-forward windows, and applies the gates.
Per spec §3.2 (X16), §3.5 (X17).
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>
