Commit Graph

24 Commits

Author SHA1 Message Date
jgrusewski
95a77c4ac6 diag(ml-backtesting): max_hold enforcement counters localize why 86.5% of trades exceed the 60s cap
S2 cluster smoke ppcfk (29f7e923c) showed mean trade hold = 432s with
max_hold_ns=60s — 886/1024 trades exceed the cap that should hard-flat
them. Upload chain and pnl_track entry_ts write look structurally
correct on inspection. Need per-step counters to localize WHICH step
of the chain fails.

Adds 7 counters in stop_check_isv (kernel_calls, seen_zero, entry_ts_zero,
ts_underflow, elapsed_below_cap, would_fire, force_flat_written) plus 1
counter in seed_inflight_limits_batched (seed_saw_force_flat). Logged at
smoke end as a third nan_counters-like line.

Diagnostic only — no behavioral change.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-20 15:35:20 +02:00
jgrusewski
76ec68c1c9 fix(ml-backtesting): per-level price-range validation in walk_ask_for_buy/walk_bid_for_sell — eliminates sized-but-bad-priced sentinel propagation
v2 NaN instrumentation (S1.21) localized the bug to apply_fill_to_pos's
open-from-flat branch writing pos.vwap_entry = avg_px where avg_px was 0
(194 cases) or > 21M finite (216 cases). The arithmetic was clean — root
cause is walk_ask_for_buy/walk_bid_for_sell consuming size from deep
levels (k=1..9) that have lvl_sz > 0 but lvl_px = 0 (or huge sentinel).
MBP-10 fills empty depth slots beyond available levels with these
sentinels.

Existing per-level filter checked lvl_sz <= 0, !isfinite(lvl_sz),
!isfinite(lvl_px) — but allowed lvl_px = 0 and lvl_px > max range.

Fix: thread per-backtest min_reasonable_px / max_reasonable_px (already
uploaded for the top-of-book skip in book_update_apply_snapshot via
S1.19) through to walk_*, and reject any level whose price falls
outside [min_px, max_px]. Same fix shape as Bug C-b (top-of-book skip),
now applied at all 10 depths.

Changes: resting_orders.cu (walk_* signatures + kernel param + 2 call
sites), order_match.cu (walk_* signatures + kernel param + 1 call site),
sim/mod.rs (submit_market + step_resting_orders launch args). 19 CUDA
tests pass.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-20 13:49:28 +02:00
jgrusewski
fc41440dc9 diag(ml-backtesting): finer NaN instrumentation — per-vwap-write-site + last-bad-vwap capture
v1 instrumentation (S1.20) eliminated 3 of 6 hypotheses: apply_fill_to_pos
arithmetic is fully clean (avg_px=0 realised=0 realized_pnl=0). But
pnl_track open branch still saw vwap_entry=0 (194 times) or > 21M (216
times) at the open transition.

v2 adds per-vwap-write-site counters so we can pinpoint which apply_fill
branch produced the bad vwap, plus captures the actual last-bad-vwap
value and the path id (1..6 = open-flat, scale-in zero/huge, flip-beyond
zero/huge, session-gap-saw-stale).

The 7th counter (vwap_session_gap_was_bad) fires when the session-gap
force-close path sees pos.vwap_entry already in a corrupt state pre-gap.

Logged at end of smoke as `nan_counters_v2: zero_flat=N huge_flat=N
zero_scale=N huge_scale=N zero_flip=N huge_flip=N session_gap_was_bad=N
| b0_last_bad_vwap=X b0_last_bad_path=N`.

19/19 stop_controller CUDA tests pass; 5/5 decision_floor_coldstart pass.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-20 13:27:03 +02:00
jgrusewski
025554afcd diag(ml-backtesting): kernel-side NaN instrumentation for residual sentinel root-cause
Cluster smoke 88dbp (post-Fix-S1.19 parameterized price range) still
produces 97 zero + 85 i32::MAX sentinel entry_px values despite source
data being fully filtered. Sentinels originate in kernel arithmetic
paths post-sanitization, not from book input.

Adds 6 per-backtest u32 counters incremented at NaN-producing sites:
- nan_avg_px: avg_px = total_cost/filled_lots -> NaN/Inf
- nan_realised: (avg_px - vwap_entry) * dir * unwind -> NaN/Inf
- nan_realized_pnl: pos.realized_pnl becomes non-finite after += or -=
- zero_vwap_at_open: pnl_track open branch saw vwap_entry == 0
- saturated_vwap_at_open: pnl_track open saw |vwap_entry| > 21M or NaN/Inf
- defensive_exit_clamp: pnl_track close defensive clamp fired

Each counter printed at end of smoke via nan_counters: log line.

apply_fill_to_pos (resting_orders.cu) gains 3 counter args; NaN-producing
paths return early WITHOUT propagating into pos state. pnl_track_step
gains 3 counter args. All call sites threaded through sim/mod.rs launches.

NanCounters struct + read_nan_counters() accessor added to LobSimCuda.
Unit test nan_counters_initialize_to_zero confirms zero-init and accessor
compile. 18/18 stop_controller, 5/5 decision_floor_coldstart pass.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-20 13:00:51 +02:00
jgrusewski
c03cf9aa38 fix(ml-backtesting): parameterized price-range sanitization replaces hardcoded 1e8
Audit (fxt-data-audit on ES.FUT_2024-Q1) revealed real source-data
outliers that the hardcoded < 1e8 threshold didn't catch:
- bid_min = -$4.85 (negative bid)
- bid_p1 = $64.15 (1% of bids in sub-$100 range, far below ES)
- ask_max = $53,012 (10x above any plausible ES price)

The < 1e8 threshold = $100M was useless: never triggered on real ES.

Fix: parameterize the range. min_reasonable_px / max_reasonable_px as
per-backtest fields in UniformSimParams, ResolvedSimVariant, and
BatchedSimConfig (defaults 0.0 / f32::INFINITY = no-check, preserves
existing test fixtures). Sweep_smoke.yaml sets 1000/20000 for ES
futures — catches all observed outliers without rejecting any plausible
price. BacktestHarness calls upload_price_range() once after creating
the sim so the bounds are active before the first apply_snapshot.

CUDA kernel book_update_apply_snapshot gains two new args
(min_reasonable_px[n_backtests], max_reasonable_px[n_backtests]) that
replace the hardcoded > 0.0f && < 1.0e8f checks at both the top-of-book
gate and the per-level sanitization pass.

Test price_range_rejection_skips_snapshot validates: sub-$1000
snapshot, super-$20000 snapshot, and negative bid all skip with
snapshots_skipped counter increment; valid ES snapshot passes through.
17/17 stop_controller tests pass.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-20 12:31:43 +02:00
jgrusewski
a85f38e97a fix(ml-backtesting): three residual sentinel + session-gap fixes
Three orthogonal fixes for residual issues in smoke stx9p (97 zero
sentinels, 85 i32::MAX sentinels, 840/1024 over-60s holds):

Fix 1 — zero-sentinel residue (px > 0):
- Per-level sanitization in apply_snapshot_kernel only checked sz > 0.
- A book level with px=0, sz>0 passed → walk_* computed total_cost=0
  → avg_px=0 → vwap_entry=0. Trade record reported entry_px=0.
- Add px > 0.0f to bid_ok/ask_ok conditions.

Fix 2 — i32::MAX upper bound (px < 1e8):
- Post-Bug-D (1e9 nanoprice scaling) some boundary events carried prices
  larger than any plausible instrument. Saturating cast to i32 produced
  the 21474836 sentinel even when isfinite() passed.
- Add px < 1.0e8f upper bound to per-level AND top-of-book validation.

Fix 3 — session-gap force-close:
- max_hold check fires at decision_stride frequency, but during weekend
  halts no events advance current_ts → max_hold never fires until next
  session. Result: 49h holds in stx9p (840/1024 over 60s threshold).
- Detect ts gap > 1 hour in resting_orders_step. If position is open,
  zero position_lots directly. pnl_track_step's existing close branch
  emits the TradeRecord on the next call. No synthetic P&L added —
  records show realised_pnl from whatever was accumulated before the gap
  (honest: cannot fill across a halt).
- New per-backtest last_event_ts_d slot tracks the previous event ts.
- Test session_gap_force_closes_open_positions: 2-hour ts jump after
  open verifies force-flat fires and exactly 1 TradeRecord is emitted.

All 16 stop_controller tests pass. All 5 decision_floor_coldstart pass.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-20 10:14:18 +02:00
jgrusewski
5235b4515b fix(data,ml-backtesting): DBN nanoprice scaling + skip-on-corrupt-top-of-book
Two root-cause fixes surfaced by cluster smoke v74v4 (sweep_smoke-a2dfc6d99):

Bug D — DBN parser price scaling:
- BidAskPair::price_to_f64/price_from_f64 used /1e12 / *1e12 from test-data
  calibration. DBN production uses 1e-9 nanoprice (the DBN standard). ES at
  5500 raw 5_500_000_000_000 → 5.5 instead of 5500. Smoke trade records
  showed entry_px=5.24 instead of expected ~5240 ES index points (1000×
  too small). Fix: 1e12 → 1e9 in both functions. Round-trip symmetric;
  tests updated.

Bug C-b — corrupt top-of-book sentinel:
- Per-level sanitization (Task 15) zeros each unhealthy MBP-10 level
  individually. At session-boundary events with all 10 levels invalid,
  the book becomes uniformly zero. apply_fill_to_pos then reads
  bid_px[0]=0 / ask_px[0]=0 → vwap_entry=0 → trade record entry_px=0
  (zero sentinel in v74v4 CSV, 162/1024 trades in n59t4).
- Fix: pre-validate top-of-book in apply_snapshot_kernel. If
  bid_px[0]/ask_px[0]/bid_sz[0]/ask_sz[0] are non-finite or
  bid_px[0]<=0/ask_px[0]<=0/bid_sz[0]<=0/ask_sz[0]<=0, atomically skip
  the entire snapshot (book/prev_mid/atr_mid_ema unchanged). Add
  per-backtest snapshots_skipped_d counter for observability.

Test corrupt_top_of_book_skips_snapshot_and_increments_counter validates
NaN top-price + zero top-size cases both increment the counter without
mutating state, and that a subsequent valid snapshot updates normally.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-20 09:37:07 +02:00
jgrusewski
bf619a2e7b feat(ml-backtesting): max-hold + exit_px defensive + spec criteria revision
Three follow-ups from cluster smoke gp74n (trade_vol pearl validation):

1. Max-hold force-close: max_hold_ns added as per-backtest config
   (default 0 = disabled). Fires force-flat (3, 0) when current_ts -
   entry_ts >= max_hold_ns, BEFORE SL/trail check. Tested via
   max_hold_forces_close. Sweep YAML sets 60s cap to bound the long
   tail observed in gp74n (263985s pathological hold).

2. exit_px defensive sanity check: 500/1024 gp74n trades reported
   exit_px = ±i32::MAX/100 (float→int saturation sentinel from likely
   NaN segment_realized). Defensive fix in pnl_track_step: if exit_px
   is non-finite or diverges from entry_px by >10%, fall back to
   entry_px with zero realised_pnl. Root cause to be traced separately.

3. Spec §9.2 criteria revision: mean_hold<30s replaced with p95<600s
   + max<=max_hold_ns. The 30s threshold reflected the pre-pearl
   sub-cost churning bug, not a real feature criterion. p95 catches
   long-tail pathology while letting alpha-driven exit timing breathe.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-20 08:10:00 +02:00
jgrusewski
da6e887174 feat(ml-backtesting): cost-floor in stop_check_isv prevents sub-cost churning
Cluster smoke w7b4p showed 44864 closed trades in 1M events at 200ms
latency — physically impossible without sub-cost churning. Root cause:
sl_distance = max(pnl_ema_loss, atr) ≈ 0.05 at cold-start, but round-trip
cost = 2 × 0.125 = 0.25. Every stop-out guaranteed loss > 5× distance;
EMA converges sub-cost.

Amendment: triple-max sl_distance = max(pnl_ema_loss, atr, 2*cost);
trail_distance = max(pnl_ema_win, atr, 2*cost). cost is an ISV-discipline
strategy anchor (already per-backtest from P4 — no new state slot, no
tuned constants). Added cost_per_lot_per_side_per_b param to both decision
kernel signatures and threaded cost_per_lot_per_side_d through both launch
sites in step_decision_with_latency.

Test: cost_floor_prevents_sub_cost_stops validates unrealized in (ATR, 0.25)
does NOT fire SL; unrealized > 0.25 DOES fire. Spec amended (§10, §12).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-20 01:27:28 +02:00
jgrusewski
4101796ad9 refactor(ml-backtesting): delete StopRules + use_cold_start_stopgap atomically
- StopRules struct + sl_tp_rules field + all literals deleted; the ISV
  stop controller (Tasks 2-9) replaces this dead data path.
- use_cold_start_stopgap propagation deleted across harness,
  batched_config, fxt-backtest, sweep_smoke.yaml, and 5 test files.
- Q1 stopgap branch in harness.rs deleted; replaced with the original
  simple strategy-upload loop.
- decision_floor_coldstart test retargeted: cold_start_persistent_
  bullish_with_default_stops_never_closes -> ..._now_closes, asserts
  trades > 0 AND |pos| <= max_lots (99 trades, pos=1 on local GPU).
- CBSW spec + plan prefixed with SUPERSEDED headers.

Per feedback_single_source_of_truth_no_duplicates +
feedback_no_partial_refactor: contract change migrates every consumer
in one commit. No legacy wrappers; no version suffixes.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-20 00:40:14 +02:00
jgrusewski
0c4b106394 feat(ml-backtesting): pnl_track_step resets trail_hwm on close transition
Two-instruction addition to the existing close-emission branch
(prev != 0 && now == 0). Without the reset, the next entry inherits
a stale HWM that could arm the trail spuriously on event 0 of the
new trade.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-05-20 00:31:05 +02:00
jgrusewski
1d852a994d feat(ml-backtesting): wire stop_check_isv into decision_policy_program
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>
2026-05-20 00:20:40 +02:00
jgrusewski
2cb4dfed28 feat(ml-backtesting): stop_check_isv __device__ helper + gate-on-flat
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>
2026-05-19 23:54:48 +02:00
jgrusewski
24a8b4c621 chore(ml-backtesting): allow(unsafe_code) on sim/mod.rs for cudarc launches
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>
2026-05-19 23:50:31 +02:00
jgrusewski
dc8c37e11e feat(ml-backtesting): ATR-EMA on mid-price in book_update_apply_snapshot
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>
2026-05-19 23:49:02 +02:00
jgrusewski
5b5292aecd feat(ml-backtesting): add stop-controller state slots + accessors
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>
2026-05-19 23:45:39 +02:00
jgrusewski
ef831ba598 obs(ml-backtesting): trades=N on progress line for intra-run firing visibility
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>
2026-05-19 22:29:52 +02:00
jgrusewski
fef5939556 feat(ml-backtesting): cold-start stopgap — max-confidence bytecode policy (Q1/Tier1)
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>
2026-05-19 22:00:52 +02:00
jgrusewski
b1f8cd4389 chore(ml-backtesting): drop dead snapshot helpers from P3 migration
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>
2026-05-19 18:11:23 +02:00
jgrusewski
9d4fda36ab feat(ml-backtesting): batched-cell sweep schema + 140-variant runner (P6)
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>
2026-05-19 17:41:20 +02:00
jgrusewski
cd82f9a4a0 feat(ml-backtesting): threshold gate + per-fill cost integration (P4)
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>
2026-05-19 17:13:29 +02:00
jgrusewski
3836e25783 feat(ml-backtesting): detect_close_transitions_batched kernel (P3)
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>
2026-05-19 17:04:42 +02:00
jgrusewski
b741f0c5ce feat(ml-backtesting): seed_inflight_limits_batched kernel (P2)
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>
2026-05-19 17:01:25 +02:00
jgrusewski
b8966fb1a6 feat(ml-backtesting): per-backtest sim parameter arrays (P1)
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>
2026-05-19 16:58:10 +02:00