Both decision kernels now call the same __device__ helper at the top
of per-backtest dispatch (after the program_lens early-return). Single
source of truth for stop logic across default and bytecode-VM paths.
Parity test ensures they produce identical market_target output for
matched scenarios.
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>
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.
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).
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>
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>
