Self-critical review of v1 (commit 0f3484325) found 12 issues. All
fixed inline. Spec marked Status: Design v2 with explicit changelog.
LOAD-BEARING FIXES
1. §4.4 conviction cold-start: removed backwards "uniform fallback"
branch that maintained trading authority when net_edge=0 across
all horizons (exactly when model has no edge). Replaced with ε
floor on net_edge_h. Single formula, smooth bootstrap, no regime
switch.
3. §4.3 exit logic: 5 OR conditions → 1 composite exit_signal scalar
(max of degradation, disagreement, pnl_decay terms). SL/trail and
8h circuit-breaker stay orthogonal as safety, not exit policy.
Attribution per term recorded in §7 state for observability.
6. §9 Gate 2: PF > 1.0 (2.5× improvement, unrealistic) → tiered
MUST (no-regression) + WIN (real lift, four candidate criteria) +
explicit STOP condition if MUST passes without WIN.
7. §9.1.1: alpha-realization metric defined explicitly as
realized_pnl / max(unrealized during signal-validity window).
Was hand-waved in v1; now formulated.
TRACTABLE INLINE FIXES
2. §7 open_trade_state: 128 → 64 bytes. Diagnostic-only fields
(scale_up/down counts, max/min target reached) moved to a
separate audit buffer (single source of truth keeps state lean).
4. §5.1 envelope: multiplicative formula → additive bounded
contributions. tanh(sharpe), min(dd/cap), clamp(vol). No single
factor can crash envelope to zero. Floor (envelope_floor lots)
and hard-cap (envelope_hard_cap) remain.
5. §5.2 threshold: hand-waved "self-tune toward PF" → explicit
3-arm bandit (lower, current, higher percentile) every K_eval
closed trades, Wiener-α EMA target with floor, bounded
[0.5, 0.95].
8. §3.4 + §3.6: Layer A wall-time ≤ 5× → ≤ 2× (controller is
light, 5× would indicate structural bug, not config tuning).
9. §4.2: explicit Wiener-α formula for conviction EMA with floor
at 0.4 per pearl_wiener_alpha_floor_for_nonstationary.
10. §6.2.1: shadow eval defined explicitly with backtest-mode (last
15% of fold) and live-mode (parallel-prediction-only) variants.
Promotion criteria spelled out (PF, Sharpe, tail loss).
11. §6.2: Layer D loss explicit (primary = value regression on
realized PnL conditional on state; secondary = EWC++ regularizer
toward base). Was three losses with no specified combination.
12. §10.5: adapter regime over-fit risk added with sample-count
down-weighting, cross-regime shadow eval, periodic adapter reset
as hard backstop.
GREENFIELDS (per user direction)
- §3.3: decision_stride field DELETED from YAML + Rust (not
deprecated). No backwards-compat shim. Per
feedback_no_legacy_aliases.
- §7.1: open_trade_state 24→64 byte migration is atomic, no old
layout shim.
- §8: boundary matrix decision_stride row marked REMOVED, not
DEPRECATED.
§13 pearl conformance: per-section pearl application (not bulk
citation). Added pearl_audit_unboundedness_for_implicit_asymmetry,
pearl_no_deferrals_for_complementary_fixes acknowledgement (layers
are sequenced amplifications, not parallel fixes), and the new
pearl_stop_checks_run_at_deadline_cadence from S2.
Status: Design v2 — awaiting user review before writing-plans.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
S2.2 moved max_hold to event-rate at step_resting_orders (cap is now
enforced — mean hold 432s → 58.43s, max 173893s → 96s). But it mirrored
the session-gap pattern `pos.position_lots = 0;` which intentionally
skips PnL ("cannot fill across halt"). Max_hold differs: the market is
live so the close SHOULD realize PnL via the current top-of-book.
Smoke t9msj (b92bd72c7) showed the consequence: 84% of trades record
$0 realised_pnl, win_rate dropped to 3.3%, total_pnl looks artificially
better (-$225k) only because losses aren't recorded.
Fix: route the force-close through apply_fill_to_pos by synthesizing a
closing fill at book.bid_px[0] (long) or book.ask_px[0] (short). The
existing counter-direction branch (realised = (avg_px − vwap) × dir ×
unwind) correctly realizes PnL on the unwound position. Top-of-book is
already validated by book_update_apply_snapshot's skip, so close_px is
guaranteed in range.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
S2.1 instrumentation (smoke cqpph @ 95a77c4ac) revealed the chain
worked end-to-end: max_hold check fired 1661 times, force-flat target
was written and seen by seed_inflight. Yet 86.5% of trades exceeded
the 60s cap with mean hold = 432s. Root cause: decision_policy's
stop_check_isv only runs every decision_stride events. At
decision_stride=200 on ES Q1 (2M events / ~91 days = ~3.9s sim-time
per event), decisions are ~13 min apart in sim-time, so the cap is
enforced 13 min late on average.
Fix: move the max_hold check to resting_orders_step (event-rate, runs
every iteration), same pattern as the session-gap force-close at
resting_orders.cu:282. Thread max_hold_ns_per_b and open_trade_state
into resting_orders_step. SL/trail stay in stop_check_isv because
they depend on ISV controller state at decision-rate.
Remove the dead decision-rate max_hold code and its 6 diagnostic
counter params from stop_check_isv per single-source-of-truth and
no-hiding rules. Keep mh_kernel_calls and mh_force_flat_seen_by_seed
counters as ongoing generic diagnostics. Update harness.rs log line
and stop_controller tests to exercise the event-rate path.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
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>
S1.20-S1.22 chased a wrong hypothesis (walk_* deep-level filter) for
3 rounds. The actual bug: resting_orders.cu:344 fills at cost = take *
s.price in the queue-decay maturation arm. Market-like orders seed at
line 644 with sentinel s.price = 1e9 (buy) / 0 (sell), counting on the
marketability check at lines 372-... to fire first via walk_*. But
queue_position initializes to 0 (no book level matches the sentinel),
so on the first same-side trade volume the queue-decay arm fires and
injects the sentinel into cost: buy cost = take * 1e9 -> avg_px =
1e9 -> huge_flat=216, sell cost = take * 0 -> avg_px = 0 ->
zero_flat=194. Same pattern explains flip=9+6.
Fix: detect s.price outside [min_reasonable_px, max_reasonable_px] in
the queue-decay arm; absorb queue_position depletion without filling.
Marketability check fires in the same iteration via walk_* with proper
book prices.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
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>
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>
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>
One-shot diagnostic that loads a .dbn.zst via
load_or_predecode_mbp10 and reports symbol distribution, per-symbol
price stats (min/p1/p50/p99/max for bid_px[0]/ask_px[0]), outlier
counts (zero-price, huge-price >$100k, zero-price-with-size),
and first-record samples per symbol.
Built to investigate the 18% sentinel-laden trade residue in the
ISV stop-controller cluster smokes (97 zero, 85 i32::MAX, 14 weird-
other). The controller path is structurally correct; remaining bad
records hypothesized to originate from source MBP-10 data (mixed
symbols, corrupt records, or unreported instrument families).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
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>
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>
Smoke n59t4 (commit 691dec144, root-cause NaN/Inf fix) revealed 96
trades/sec at decision_stride=4 — decisions fire every ~4ms while
orders take 200ms to fill. Controller makes decisions on phantom
in-flight positions; stops fire against positions that haven't
materialized yet. At $0.25 round-trip cost × 96 trades/sec, fee burn
is $86k/hour.
Architectural mismatch: decision_stride should be >= latency_ns /
event_dt_ns. At 200ms latency and ~1ms/event, stride=200 means
decisions fire every ~200ms — once per latency cycle. Resulting
~10k decisions over 2M events is the natural cadence for the
deployment anchor.
Not a controller bug. The controller logic (NaN-clean per Task 15)
is correct; the smoke config was tuned for sub-millisecond latency
that doesn't match the 200ms Scaleway PAR → IBKR anchor.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Two orthogonal bugs that compounded to produce the exit_px=±21474836
sentinel in cluster smokes (315 trades baseline, 500 trades pearl):
Bug A — NaN/Inf book propagation:
- apply_snapshot_kernel passes through NaN/Inf prices from MBP-10
predecoded data (session boundaries, gap-fills).
- walk_ask_for_buy / walk_bid_for_sell accumulate cost += take * NaN.
- apply_fill_to_pos: avg_px = NaN; realized_pnl += NaN → permanent NaN.
- Subsequent close: (int)(NaN-derived * 5000) → i32::MAX sentinel.
- Fix: zero-out non-finite levels in apply_snapshot_kernel; add
isfinite() guards in walk helpers as defense-in-depth; ATR update
guards against non-finite mid.
Bug B — pnl_track close branch doesn't reset scratch:
- Scratch reset (full 24-byte memset to 0) was already present in the
current code; no source change required for Bug B.
Tests:
- book_nan_inf_prices_dont_corrupt_realized_pnl: NaN at ask[3] + Inf
at bid[5] survives the fill pipeline without making realized_pnl
non-finite.
- pnl_track_resets_scratch_on_close: open+close+5 flat events emits
exactly 1 record; second open+close emits exactly 1 more (=2 total).
14/14 stop_controller tests pass; all other ml-backtesting test suites
unaffected.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
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>
Replaces the Task 12 `2.0f * cost` floor (hardcoded multiplier) with
trade_vol = sqrt(realised_return_var) bootstrapped from cost². Per
pearl_trade_level_vol_for_stop_distance.md: microstructure ATR is the
wrong time scale for trade-level stop decisions; per-horizon
realised_return_var is the right one, with cost² as a structural cold-
start sentinel.
cost now appears exactly once — inside the sqrt as a bootstrap sentinel,
never as a distance multiplier. The 2.0f literal is eliminated;
controller is fully ISV-driven.
var_avg accumulates realised_return_var in the same single-pass horizon
loop as ema_loss/ema_win. Cold-start (var_avg=0): trade_vol = cost.
Post-bootstrap: sqrt(var_avg) dominates.
Test retargeted: cost_floor_prevents_sub_cost_stops →
trade_vol_floor_prevents_sub_cost_stops, with boundaries straddling
trade_vol=cost=0.125 instead of the prior 2*cost=0.25 (no-fire Δ=0.08,
fire Δ=0.20).
Spec §5, §10, §11, §12 amended.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
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>
- 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>
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>
market_targets[b] now interpreted as target absolute position (side=0
long, side=1 short, side=2 no-op preserved, side=3 force-flat).
seed_inflight_limits_batched computes order_lots = target_signed -
effective_position where effective_position sums pos.position_lots
plus all unfilled signed slot sizes (active in {1, 2}). Fixes both
the additive accumulation bug AND the worse latency-bypass variant
(naive delta would have made things 200x worse under 200ms latency).
3 new tests added (all passing):
- position_target_not_additive: latency=0, 10 repeated target events stay <= max_lots=5
- position_target_not_additive_with_latency: 200ms latency, 500 events, in-flight summation prevents runaway
- alpha_noop_side_2_preserved: side=2 events leave filled position unchanged
All 9 stop_controller tests pass; parallel_sim and fuzz regression clean.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
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>
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>
11 tasks, each one commit, TDD-disciplined (test → fail → impl →
pass → commit). Covers all 11 sections of the spec at
docs/superpowers/specs/2026-05-19-isv-driven-stop-controller-design.md:
- Tasks 1-2: state slots + ATR EMA in book_update_apply_snapshot
- Task 3: stop_check_isv shared __device__ helper skeleton
- Tasks 4-6: SL trigger, trail-TP + HWM ratchet, multi-horizon avg
- Task 7: bytecode VM parity (wire helper into decision_policy_program)
- Task 8: seed_inflight_limits_batched target-delta + in-flight summation
- Task 9: pnl_track_step trail_hwm reset on close
- Task 10: atomic deletion of StopRules + use_cold_start_stopgap +
integration-test retarget + CBSW SUPERSEDED markers
- Task 11: cluster smoke validation via argo-lob-sweep.sh
All 10 §9.1 unit tests have explicit failing-test scaffolds before
their implementations, and the retargeted integration test
(cold_start_persistent_bullish_now_closes) flips the assertion to
prove the fix.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Issues caught in self-review and fixed:
1. Encoding ambiguity: stop-fire writing (2, 0) would have silently
collapsed weak-alpha no-op semantics with force-flat. Introduced
distinct side=3 = force-flat; preserved side=2 = no-op for entry
path. Updated §3 + §5 + §7 truth table.
2. trail_hwm reset on close was implied in §4 but missing from §8's
touched-files list. Added pnl_track.cu modification (§7.1) +
added to §8 Added list.
3. __device__ helper enforcement: §3 now mandates stop_check_isv()
as a shared __device__ function called from both decision kernels,
not duplicated code.
4. Kernel arg additions enumerated explicitly in new §4.1 table —
every kernel-launch site gets named, no "thread through every kernel"
hand-wave.
5. seed_inflight_limits_batched (resting_orders.cu:439–475) named
explicitly in §1 + §7 — the actual kernel that needs the position-
target-semantics fix.
6. In-flight order summation: naive delta = target_signed - pos was
strictly worse than the original additive bug under non-zero
latency (would produce pos = pre + K·delta after fills). Fixed
§7 algorithm to compute effective_position by scanning all
active∈{1, 2} slots and summing their signed sizes. Added
position_target_not_additive_with_latency test.
7. ATR α=0.4 honest justification in §10: not strictly derived from
pearl_wiener_alpha_floor_for_nonstationary (which is about co-
adapting control loops, not passive observation); chosen as a
pragmatic project-wide constant matching isv_kelly_update_on_close.
8. CUDA Graph host-branch-free affirmation added to §3 — explicit
compatibility with P5 graph capture.
9. Cold-start symmetry caveat (§5): at n_trades_seen=0 both EMAs are
0 so sl_distance == trail_distance == atr — asymmetry only
emerges post-bootstrap. Acknowledged honestly.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Replaces StopRules::default()-all-zeros (the actual cause of the cluster
smoke's n_trades=0, falsifying the CBSW dilution diagnosis) with an
ISV-driven SL + trail-stop controller folded into the existing decision
kernel(s). Per-backtest ATR EMA on mid-price + per-horizon pnl_ema_*
drive distances under max(real, floor) per
pearl_blend_formulas_must_have_permanent_floor. No new kernel; single
source of truth in step_decision*. Folds in the position-target
semantics fix (200-event local repro showed position_lots=199 from
additive-vs-target order semantics) — same commit.
Removes: StopRules, sl_tp_rules, Q1 stopgap field/branch,
use_cold_start_stopgap propagation. Marks CBSW spec + plan SUPERSEDED.
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>
4-task atomic ladder (Q1-Q4) implementing spec 566e8bcb0. Each task
maps to one commit per feedback_no_partial_refactor:
Q1: Cold-start stopgap (bytecode max-confidence policy upload)
Q2: CBSW kernel aggregator + Tier 1 revert (atomic)
Q3: Memory pearl pearl_conviction_bootstrap_for_kelly_aggregation
Q4: Parallelism spec §3.4 cross-reference
Kernel ABI unchanged across all 4 commits — Q2 modifies only
decision_policy_default's body (decision_policy_program left as a
pluggable bytecode-VM experiment surface). Q2 atomically deletes
the Q1 stopgap field/CLI/YAML/harness branch in the same commit
that lands the kernel fix (feedback_no_legacy_aliases compliance:
grep returns zero hits for use_cold_start_stopgap post-Q2).
Validation gates per spec §9:
- Q1: cluster smoke n_trades > 100, total_pnl != 0 (diagnosis
confirmation). If n_trades = 0 still, STOP — downstream bug.
- Q2: pre-existing 11 CUDA tests still pass; 7 new cbsw_*
regression tests pass; cluster smoke n_trades > 100; ev/s
rate ratio Q2/Q1 ≥ 0.95.
Self-review confirms all 11 spec sections covered, no placeholders,
type/signature consistency across tasks.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Critical review surfaced 11 actionable issues in the v1 spec; 10 fixed
inline (#11 — legacy-comparison test — dropped per user decision since
empirical legacy behavior is already known: n_trades=0).
Performance:
1. expf → piecewise-linear ramp (~5× cheaper on GPU, no transcendental
in the hot path). Operationally equivalent: monotonic, saturating,
midpoint at K. Side-benefit: pure max-confidence at cold-start
(sq=0 instead of sigmoid's 11.9% leakage).
2. Single fused per-horizon pass replaces the two-loop sketch.
3. Tier 2 scope narrowed to decision_policy_default ONLY. The bytecode
VM (decision_policy_program) stays unchanged — runs only for custom
strategy experiments, never in production policy. Halves Tier 2's
surface area + test cost.
4. __device__ helpers cbsw_signal_quality + cbsw_weight introduced for
single source of truth.
Correctness bugs (silently present in v1 sketch, would have shipped):
5. strong_h and sq_min_active initializers added (were referenced
before init in the per-horizon loop).
6. Attribution mask at cold-start was setting all 5 bits because every
w[h] >= floor > 1e-9; trades would pollute ALL horizons'
recent_sharpe. Fix: binary split — at sq_min_active < 0.5 attribute
only to strong_h; at mature, attribute per weights > floor + epsilon.
7. sq_min was "min over h", which permanently locked the aggregator
into cold-start mode if any horizon never gets attributed (e.g.,
h6000-only-trading regime). Fix: sq_min_active = min over h with
n_trades_seen > 0; cold horizons don't gate maturity.
8. Opposing-horizons case now correctly fires on the strongest single
horizon at cold-start (piecewise-linear sq=0 → pure max-conf), with
explicit design-choice note explaining why this conservatism trade-
off favors firing.
Spec hygiene:
9. Rate validation committed to a measurement gate (Q2 ev/s ≥ 95% of
Q1 ev/s) instead of back-of-envelope estimate.
10. Kernel ABI explicitly stated as unchanged → feedback_no_partial_
refactor compliance is trivial at the kernel boundary.
12. Tier 1's use_cold_start_stopgap field is DROPPED in Q2 atomically
(not orphaned), and DoD checklist includes a grep-zero gate for
feedback_no_legacy_aliases compliance.
Risks updated: removed the sigmoid-narrowing risk (irrelevant now);
added register-pressure risk with the rate-gate mitigation; added
explicit "cold-start may lose on noisy trades" risk with empirical
escalation path (bump kelly_floor 0.20 → 0.40 if Q2 smoke shows large
negative PnL).
Math walk-throughs rewritten for piecewise-linear (different numbers
at cold-start): cold = pure max-conf, mature = pure weighted-sharpe,
clean separation at sq_min_active threshold.
Awaiting review of the revised spec before writing-plans dispatch.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The post-trunk-grows threshold-tuning smoke (81decf40f) produced
n_trades=0 despite the model having a HEALTHY max-conviction
distribution (74.6% of decisions ≥ 0.30, 26.7% ≥ 0.70, full spread
across [0,1]). Diagnosis: the linear-weighted-mean aggregator in
decision_policy_default is structurally dilution-bound at cold-start
— single-horizon strong signals get washed out when uniform-floor
weights produce mean-over-horizons aggregation.
Solution: Conviction-Bootstrapped Sharpe Weighting (CBSW). Hybrid
max-confidence × weighted-sharpe with a per-horizon sigmoid transition
keyed on n_trades_seen vs MIN_TRADES_FOR_VAR_CAP. Cold-start: sig_mag
(decision-time conviction) drives weights AND max-confidence aggregator
fires single-horizon trades. Mature: recent_sharpe (historical) drives
weights AND linear-mean aggregator emphasizes strong-Sharpe horizons.
Permanent floor preserved per pearl_blend_formulas_must_have_permanent_floor.
Mirrors DQN bootstrap pattern (pearl_thompson_for_distributional_action_
selection): when historical estimates are uncertain, use available
signal as the bootstrap. Sig_mag is the ISV signal at decision time;
recent_sharpe is the ISV signal at trade-close time. Transition
self-terminates based on data accumulation, not time constants.
3-tier delivery (one spec, atomic commits):
Q1: Bytecode VM stopgap — upload max-confidence 7-instruction
program per backtest. Validates diagnosis; zero kernel work.
Q2: Kernel CBSW — replace weight + aggregator in both decision
kernels. 5 new regression tests covering cold/mature/transition.
Q3: New memory pearl pearl_conviction_bootstrap_for_kelly_aggregation
capturing the lesson.
Q4: Cross-reference from parallelism spec (deployability sweep
depends on CBSW being live to produce meaningful verdict).
The parallelism work (P1-P6) is fully working — confirmed by the
threshold-tuning smoke completing end-to-end (Succeeded status, 500k
decisions, artifacts written, aggregator parquet emitted). What's
blocked is the deployability VERDICT, because the dilution bug means
all variants would show n_trades=0 regardless of cost/latency/threshold.
CBSW unblocks the verdict.
Awaiting review before transitioning to writing-plans for Q1-Q4
implementation plan.
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>
data_template + cell.window interpolation produces a single-file path,
but MultiHorizonLoader expects a directory (iterates all files inside).
The data_template path is reserved for a future per-quarter loader
filter; for the smoke, fall back to scalar `data` pointing at the
directory + cap max_events at 2M to keep wall under ~20 min.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Frictionless + threshold=0 produces an uninformative number (every
signal trades for free). Swap to realistic anchor:
- latency_ns: 200ms (Scaleway PAR → IBKR baseline)
- cost_per_lot_per_side: 0.125 (1 tick on ES)
- threshold: 0.30 (moderate gate, ~p70-80 of conviction)
Uses the P6 batched flow (sim_variants list with 1 entry) so this run
also exercises the run-sweep Argo template + BatchedSimConfig::from_grid
end-to-end. Output: cell_W0/sim_t30c1l200/{summary.json, trades.csv,
pnl_curve.bin} with REAL net-of-cost Sharpe + max-drawdown.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The cargo-target-cuda PVC persists `$BUILD` across runs. cargo build
mutates Cargo.lock, so the next run's `git checkout $SHA` fails with
"Your local changes to the following files would be overwritten by
checkout: Cargo.lock". Same pattern alpha-perception-template uses
(`git checkout --force $SHA; git clean -fd`) handles this cleanly.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Closes the operational gap from P6: when a sweep grid YAML carries
sim_variants, argo-lob-sweep.sh now emits ONE run-sweep-batched task
instead of N fan-out run-cell tasks. The new run-sweep template
invokes `fxt-backtest sweep` against the full grid (base64-encoded
inline as Argo parameter to avoid YAML special-char encoding traps).
The binary's P6 sweep() function handles cell × variant fan-out
internally via BatchedSimConfig::from_grid, so one pod processes all
4 windows × 140 variants sequentially. Trade-off: no inter-window
parallelism in this rev (4 quarters sequential in one pod ≈ firm-bound
2h wall per spec §9). 3-pod scale-out is P7 future work — needs the
script to split the YAML into per-window sub-grids and emit one
run-sweep task per sub-grid.
Backward-compat: legacy (no sim_variants) flow unchanged. The dry-run
of sweep_smoke.yaml continues to emit run-cell-* fan-out tasks; the
dry-run of sweep_deployability.yaml emits one run-sweep-batched task.
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>
X11's original plan said "capture_graph_a covers full v2 forward" but
the X11 commit (4f888abbf) only shipped forward_only + from_checkpoint.
Graph capture is now actually implemented for the inference path.
Mirrors the pattern already in step_batched (perception.rs:1221-1257):
- First call: eager dispatch + set forward_warmed flag.
- Second call: begin_capture -> dispatch_forward_kernels -> end_capture
-> store CudaGraph.
- Subsequent calls: graph.launch() — captured replay.
forward_only now performs its own staging-fill of the mapped-pinned
host buffers (input data varies per call), then dispatches through
the three-state machine. The captured region is the new private
dispatch_forward_kernels helper: a copy of evaluate_batched's
forward chain (VSN -> Mamba2 x2 -> LN x2 -> attn-pool -> CfC K-loop
-> heads) that omits labels, BCE, and any stream syncs. The final
sync + dtoh of probs_per_k_d happens OUTSIDE capture in forward_only.
Per pearl_no_host_branches_in_captured_graph: no host branches /
scalar-arg-changes / host-mallocs inside the captured region; all
kernel launches use pre-bound device pointers stable across replays.
Per pearl_cudarc_disable_event_tracking_for_graph_capture: event
tracking is already disabled for the trainer's lifetime at
construction (see PerceptionTrainer::new ~line 529), so the captured
region is free of cuStreamWaitEvent / event.record() insertions.
Vestigial loader.rs:272 doc comment referencing the never-shipped
CfcTrunk::capture_graph_a updated to point at the now-real
PerceptionTrainer::forward_only warmup path.
Regression: forward_captured_matches_uncaptured — eager (call 1) vs
captured replay (call 3) agree within 1e-5 relative tolerance per
element. NOT strict bit-identity because CUDA Graph capture can
reorder kernel launches and flip f32 reductions by 1 ULP harmlessly.
Local RTX 3050 Ti run: 160 elements, max rel_err = 0e0.
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>
