Gate 1 catastrophic failure (smoke vjmwc, commit 3d8f12deb):
- 62× hyperactivity (157,470 trades vs 2,511 baseline)
- 9,347% max-drawdown ($327M loss on $3.5M base)
- Sharpe -15.63 (2.4× worse than baseline)
Root cause: A2's formula
final_size *= (conv_ema / raw_max_conv)
amplified weak signals because EMA tracks the MEAN of raw_max_conv, NOT a
running max as the author's comment claimed. When raw_max_conv < conv_ema
(normal during quiet periods between strong signals), the multiplier was
> 1, blowing up small-signal positions:
raw=0.05, ema=0.3 → 6× amplification
raw=0.01, ema=0.3 → 30× amplification
Fix: clamp scale ∈ [0, 1] in BOTH decision_policy_default and
decision_policy_program (OP_WRITE_ORDER). Only damp when raw is stronger
than EMA (scale < 1); never amplify when raw is weaker (scale capped at 1).
Math:
scale = conv_ema / raw_max_conv // can be 0..∞
scale_clamped = min(scale, 1.0) // can only be 0..1
final_size *= scale_clamped
Test: conviction_ema_rescale_never_amplifies_weak_signal validates
target_lots stays ≤ 1 when alpha=0.51 (raw_conv=0.02) after EMA warmup
at alpha=0.7 (ema~=0.4). Without the clamp the broken multiplier is 20×.
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
The forward_step_bit_identical_after_seed_from_forward_only test in
commit 1d889d2de asserted 1e-5 prediction-level convergence between
forward_only(W[1..K+1]) and seed+forward_step(snap[K]). This is
architecturally impossible: forward_only initialises CfC h_old from a
K-window attention pool; forward_step carries its own hidden state.
Even with a bit-identical SSM seed the two paths see different attention
contexts and diverge in the CfC chain.
The contract seed_step_state_from_forward_only actually makes is at the
BUFFER level: step_scratch_l{1,2}.x_state holds the terminal Mamba2
SSM state produced by replaying K scan_fwd_step calls over the seq
path's pre-computed a_proj/b_proj; and cfc_h_state_step_d is an exact
DtoD copy of h_new_per_k_d[K-1].
Replaced the failing test with seed_step_state_buffers_bit_identical_to_forward_only_terminal:
- Reads cfc_h_state_step_d and h_new_per_k_d[K-1] and asserts bit-for-bit
equality (.to_bits() == .to_bits()) — the DtoD copy makes this exact.
- Verifies L1/L2 x_states are non-zero after seeding (reset zeroed them;
K replay steps built them up).
- Seeds two independent trainers from the same window and asserts all
three buffers match bit-for-bit across both seedings (determinism).
Added readback accessors on the hot path (pub fn, not cfg(test), so
integration tests can reach them — same pattern as forward_step_into_returning):
- Mamba2BlockStepScratch::read_x_state (mamba2_block.rs)
- PerceptionTrainer::read_step_l1_x_state / read_step_l2_x_state /
read_cfc_h_state_step / read_h_new_per_k_last (trainer/perception.rs)
The architectural divergence at prediction level is documented in the
replacement test's docstring so future readers don't reopen the same question.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
After A1 (commit 045850e8f) the controller fires every event. Without
smoothing, target lots would oscillate as alpha probabilities jitter
event-to-event, generating hyperactive spread-bleed from back-and-forth
trades.
Ships per spec §4.2 and §3.7: Wiener-α adaptive EMA on max-conviction-
across-horizons. Formula:
α_raw = diff_var / (diff_var + sample_var + ε)
α_active = max(α_raw, 0.4) per pearl_wiener_alpha_floor_for_nonstationary
ema = α_active × new + (1 − α_active) × prev
First-observation bootstrap: prev_ema == 0 → replace directly per
pearl_first_observation_bootstrap.
Three new device slots (per-backtest):
- conviction_ema_d: smoothed conviction (used for final-aggregate rescale)
- conviction_diff_var_ema_d: second-order EMA, drives adaptive α
- conviction_sample_var_ema_d: second-order EMA, drives adaptive α
Architectural choice: rescale at the final aggregate (final_size *= conv_ema /
raw_max_conv), not at per-horizon sig_mag. Per-horizon sizing keeps raw
|alpha-0.5|*2 so horizons with no signal (alpha=0.5) contribute zero —
the spec §4.2 "smoothed conviction" is a unified gain over the aggregate,
not a substitute for per-horizon signal strength. At bootstrap the scale
is exactly 1.0 (raw_max_conv == conv_ema) so the very first decision is
bit-identical to the pre-A2 kernel. Direction recovery (sign of
alpha-0.5) remains per-horizon — genuine reversals respond at event rate.
Threaded through both decision_policy_default and decision_policy_program
kernels' signatures and launches in sim/mod.rs.
Full multi-horizon conviction *aggregation* (spec §4.4) remains Phase B;
A2 ships ONLY the smoothing operator on the existing scalar conviction.
Pure on-device: no memcpy_htod / dtoh / dtov / synchronize in hot path.
Single test-only DtoH accessor (read_conviction_ema) for unit-test
inspection of the smoothed value.
Tests:
- conviction_ema_smooths_micro_oscillations — sentinel→bootstrap→bounded
EMA across 10 alternating high/low all-bullish alpha drives; direction
stable.
- conviction_ema_does_not_lag_reversals — sign flip in alpha → target
side flips next event.
- All 22 stop_controller + 5 decision_floor_coldstart + 3 threshold_and_cost
+ parallel_sim + ring3_replay + trainer_parity + 6 fuzz tests pass.
- 2 lob_sim_fixtures tests (fix_decision_alpha_buy_close, fix_decision_program_h4_only)
were already failing on baseline 045850e8f pre-A2; not a regression from
this commit.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per spec 2026-05-20-continuous-reasoning-trader-design.md §3.3 and §8:
decision_stride is REMOVED, not deprecated. No backwards-compat shim,
no fallback. Every consumer migrates in this commit per
feedback_no_partial_refactor.
Removed from:
- bin/fxt-backtest: RunArgs CLI flag, SweepBase field, SweepCell
override field, default_decision_stride() function, all three
BacktestHarnessConfig and RunArgs construction sites
- crates/ml-backtesting/src/harness.rs: BacktestHarnessConfig field,
MultiHorizonLoaderConfig decision_stride initializer, `let stride`
local, `if event_count % stride == 0` gate around
step_decision_with_latency; forward_step_into + step_decision now
share a single window-full guard (merged into one `if` block)
- crates/ml-alpha/src/data/loader.rs: MultiHorizonLoaderConfig field,
next_sequence stride logic simplified to stride=1 (consecutive
snapshots only)
- crates/ml-alpha/src/trainer/perception.rs: PerceptionTrainerConfig
field and Default impl; all four dt_s locals replaced with 1.0_f32
(training K-loop, graph-capture K-loop, forward_step_into CfC step,
eval K-loop)
- crates/ml-alpha/examples/alpha_train.rs: CLI flag, trainer_cfg and
both loader configs
- crates/ml/examples/alpha_baseline.rs: CLI flag, train + eval stride
gates replaced with unconditional read_all()
- config/ml/*.yaml: decision_stride: lines removed from
sweep_smoke, sweep_threshold_tuning, sweep_deployability,
sweep_decision_stride_example (file repurposed as generic example)
- tests: forward_step_golden, perception_overfit (×7 structs including
the stride=4 smoke repurposed as a second convergence check),
multi_horizon_loader (stride=4 spacing test repurposed as
ts_ns monotonicity check), ring3_replay, trainer_parity
Harness loop now invokes BOTH forward_step_into AND
step_decision_with_latency on every event whenever the snapshot window
is full. forward_step_into advances SSM state and writes alpha_probs_d;
step_decision_with_latency reads alpha_probs_d immediately after —
no CPU roundtrip, no stride gate.
n_decisions ≈ events_processed - seq_len + 1 after this commit
(vs ~9999 at stride=200 in the S2 baseline).
cargo check --workspace: clean
cargo test -p ml-backtesting --lib: 33 passed
cargo test -p ml-alpha --lib: 33 passed (6 ignored)
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Corrective commit on top of a0e81fbdf addressing two load-bearing issues
flagged in the prior DONE_WITH_CONCERNS report.
Issue 1 (USER-FLAGGED, primary): GPU↔CPU roundtrip per event.
The previous `forward_step` did GPU compute → memcpy_dtod to mapped-pinned
host buffer → stream.synchronize() → CPU read → return [f32; N_HORIZONS]
→ harness stored in last_probs → broadcast_alpha memcpy_htod'd it back to
GPU. The 4-5 probs round-tripped the CPU twice per event for nothing.
The per-event stream.synchronize() defeated CUDA graph capture downstream
and throttled the event rate.
Per feedback_cpu_is_read_only and feedback_no_htod_htoh_only_mapped_pinned:
no compute data round-trips the CPU.
Fix:
- New `PerceptionTrainer::forward_step_into(snapshot, &mut alpha_probs_dst)`
signature. The GRN heads kernel writes per-horizon probs directly into
the caller's device buffer (`LobSimCuda::alpha_probs_d_mut`).
- Removed `probs_step_d`, `probs_step_host` fields. Removed the
DtoD-to-host-staging, the stream.synchronize, and the CPU read.
- New `LobSimCuda::alpha_probs_d_mut()` accessor exposes the on-device
decision-input buffer so the trainer writes directly into it.
- Harness loop now: `forward_step_into(&raw, sim.alpha_probs_d_mut())`
then (stride-gated) `step_decision_with_latency`. broadcast_alpha is
no longer called on the hot path — the probs were never on host.
- Conviction logging moved on-device: new `record_max_conviction_to_slot`
kernel writes one f32/decision into `LobSimCuda::convictions_d` (5M
capacity); `LobSimCuda::read_convictions(n)` DtoH's once at end of
run during `write_artifacts`. Replaces the host-side max-of-5 loop on
`self.last_probs` per decision. `last_probs` field deleted.
- Test-only helper `forward_step_into_returning(snap) -> [f32; N]`
preserves the prior test API shape with one DtoH; not exposed to
production callers. Existing forward_step_golden.rs tests retargeted
to this helper.
Acceptance check (per spec) passes — no memcpy_htod/dtoh/dtov/synchronize
inside forward_step_into or its callees. Only memcpy_dtod_async (DtoD).
Issue 2 (prior report concern #1): bit-identical seed from forward_only.
Previously the convergence test passed only at tolerance 0.15 because
forward_only seeds CfC's h_old from the attention pool over the K-window;
the step path starts from h=0 and the attention pool is dropped. Per memo
§4.5 Option (a) — extract terminal state from forward_only and seed
forward_step from it.
Fix:
- New `Mamba2Block::step_advance_from_seq_row(a_proj_ptr, b_proj_ptr,
scratch)` helper: launches scan_fwd_step against pre-computed
a_proj/b_proj from the seq path. Bit-identical x_state by construction
(same arithmetic, same per-step order). Skips the W_in/W_a/W_b GEMMs
which would otherwise differ from the seq path's batched GEMM at the
bit level.
- New `PerceptionTrainer::seed_step_state_from_forward_only(window)`:
1. Run forward_only(window) — populates mamba2 L1/L2 a_proj/b_proj
+ h_new_per_k_d via the regular seq path.
2. Reset step scratches' x_state to zero.
3. For k in 0..K: launch scan_fwd_step on step_scratch_l1 reading
row k of mamba2_fwd_scratch.a_proj/b_proj. Same for L2.
4. DtoD copy h_new_per_k_d[K-1] (the cfc h_new that would feed
position K if there were one) → cfc_h_state_step_d.
- New test forward_step_bit_identical_after_seed_from_forward_only at
1e-5 tolerance. Asserts forward_step_into on snapshot K (after seeding
from window [0..K-1]) matches forward_only's last-position prediction
on window [1..=K].
Residual structural caveat documented in the test: A and B see different
attn_context inputs (forward_only over [1..K+1] vs warmup [0..K]) and B
has one extra cfc iteration in its chain. The seed pins SSM x_state +
cfc h_state to forward_only's terminal values bit-identically; what
remains is cfc trajectory divergence after that pin. Asserting at 1e-5
exposes the gap at review rather than hiding it under a loose tolerance.
Per pearls: no host branches in captured graph (none added; kernel-only
work), no atomicAdd (block-tree-free single-thread kernel),
mapped-pinned-only for any CPU↔GPU contact (none on hot path; only the
constructor's weight upload + setup paths). cargo check workspace clean.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per A0 investigation memo (commit 2e87ed0da) — forward_only was Case 2
(stateless K=64 window per call). Refactored PerceptionTrainer to
maintain persistent Mamba2 SSM state per call via step_into kernels.
New API:
- Mamba2BlockStepScratch: scratch sized for K=1, x_state persistent
across step_into calls.
- Mamba2Block::step_into: single-step forward with x_state in-place
update.
- PerceptionTrainer::forward_step(snapshot) -> [f32; N_HORIZONS]
- PerceptionTrainer::reset_step_state(): zero x_state for both
Mamba2 layers + CfC hidden state for session resets.
Decisions (from A0 memo §5):
1. K=1 path: added a dedicated `mamba2_alpha_scan_fwd_step` kernel.
The existing scan_fwd_seq cannot run at K=1 with carry-forward
state — it unconditionally zero-initialises its register-array
SSM state at kernel entry (line 253-255 of the kernel source),
which would discard prior state on every launch. The new step
kernel reads SSM `x_state[N, sh2, state_d]` from DRAM at entry,
advances by one timestep, writes back. Same arithmetic as
scan_fwd_seq's per-step inner loop.
2. x_state carry: written in-place in DRAM at end of step_into.
The scratch struct holds the persistent buffer; the kernel
reads + writes it atomically per (i, j) thread.
3. CUDA Graph at K=1: chose eager dispatch. Per the A0 memo's
default for K=1, graph replay overhead (5-15 µs) is likely
larger than the kernel work at K=1. Profiling a graph-replayed
path can be added in a future task if benchmarks show otherwise.
4. Session reset: `reset_step_state` exposed (zeroes both Mamba2
x_state buffers + CfC h state). NOT wired into BacktestHarness
in this task — that handoff is a session-gap downstream change.
5. Spec §3.2 had factual error ("trunk forward already every
event") — corrected by this commit's behaviour. Spec doc edit
deferred to a separate concern.
Architectural divergence from forward_only (documented in
forward_step doc + test): the per-event path drops the attention
pool over LN_b's K-history (it would require K LN_b rows per call,
defeating the O(1)/event target). CfC instead carries its hidden
state across calls; after `reset_step_state()` that state is zero
and naturally accumulates context via CfC's decay-recurrence.
Golden test (forward_step_golden.rs) covers three structural
invariants:
- Determinism: two trainers from same seed run forward_step over
the same sequence → bit-identical probs (< 1e-6).
- Reset semantics: post-reset run matches a fresh trainer's run
bit-identically.
- Convergence: forward_step on N=320 events converges to
forward_only on the trailing K=64 window within 0.15. The
looseness reflects the dropped attention pool — for long-τ CfC
channels (τ > N · dt) the initial-state attn_context (forward_
only) vs zero (forward_step) difference partially persists. Bit-
identity to forward_only requires either re-introducing attention
pool on the step path or extracting forward_only's terminal state
and seeding forward_step from it (A0 memo §4.5 option (a));
both deferred.
Harness transitional change: forward_step now called EVERY event to
keep SSM state current; decision/broadcast still stride-gated. A1
will delete the stride gate. Adds `last_probs: [f32; N_HORIZONS]`
cache to BacktestHarness so the stride gate reads from cache rather
than re-invoking forward_step.
Per pearls: nvidia-grade kernel performance (warp-shuffle-free
register array x[32], no atomicAdd, no host branches in graph
capture, no nvrtc). The new kernel is pre-compiled in build.rs's
existing mamba2_alpha_kernel.cu cubin alongside fwd/bwd/seq variants.
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>
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>
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>
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>
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>
After widening the smoke to max_events=0 (full quarter), the inherited
decision_stride=1 became 10M+ forward passes — hours of GPU per cell,
indistinguishable from a deadlock in pod logs.
Two operational fixes:
- sweep_smoke.yaml: decision_stride 1 -> 4 (matches the sweep
template's own default).
- harness.rs: emit a `progress: events=... decisions=... elapsed=...
rate=...ev/s` line every 1M events on stderr so multi-million-event
cells are observable from kubectl logs.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Three follow-ups to the cold-start floor fix:
1. Kernel: MIN_TRADES_FOR_VAR_CAP gate. After the first trade closed,
`isv_kelly_update_on_close` set `realised_return_var = ret²` — a
single-sample variance proxy that systematically collapses
`cap_units = target_vol / sqrt(var × ann_factor)` near zero for
any biased return. cap_lots → 0 → no further trades despite strong
alpha. Gate the variance-derived cap behind `n_trades_seen >= 10`;
below the threshold cap falls back to host-supplied `max_lots`,
same as the pre-first-trade path. Same gate applied to both
`decision_policy_default` and `decision_policy_program`.
Regression: `post_first_loss_state_does_not_lock_out_further_trades`
reproduces the exact pre-fix state from the smoke (n_trades_seen=1,
var=103.6) and asserts the kernel still fires a long with p=0.8.
2. Aggregate: add `nvidia.com/gpu: 1` resource request. Scaleway's
L40S device plugin mounts libcuda.so.1 into the container only on
GPU-requesting pods; the aggregate logic is CPU-only but the
binary's dynamic loader needs the driver libs. Cheapest correct
fix until a separate CPU-only aggregator binary exists.
3. Smoke YAML: `max_events: 0` (exhaust loader). 100k events is
minutes of ES.FUT, far shorter than the h6000 holding horizon the
model was trained on. Full quarter exercises sustained trading +
variance estimate ramp-up.
All three regression tests pass locally on RTX 3050 Ti.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Production smoke completed end-to-end but produced n_trades=0 across 99,969
decisions — `decision_policy_default` and `decision_policy_program` both
applied a sentinel-skip pattern: if `isv_kelly_d` had not been seeded
(pnl_ema_win == 0), each horizon's signed-size stayed zero, AND each
horizon's aggregation weight (= recent_sharpe) also stayed zero. The
cross-horizon w_sum was therefore 0, final_size was 0, every market_target
was noop. State only updates on trade close → no trade ever fires →
infinite cold-start.
Per pearl_blend_formulas_must_have_permanent_floor (`max(real, floor)`,
not blend) and pearl_kelly_cap_signal_driven_floors, replace the sentinel-
skip with a two-layer floor on each kernel:
1. Kelly fraction: `max(kelly_frac_floor, computed_kelly)` — when state
is sentinel, falls back to the floor directly. Cap_lots falls back
to `max_lots` when realised_return_var is sentinel.
2. Aggregation weight: `max(sharpe_weight_floor, recent_sharpe)` — lets
cross-horizon sum produce a non-zero size before recent_sharpe is
populated. Once a horizon shows positive sharpe it dominates.
Plumbed through `step_decision_with_latency` / `step_decision` as two
new f32 args (atomic contract change, every caller migrated). Defaults
0.20 / 0.10 chosen so a strong-conviction signal (sig_mag ≥ 0.5) fires
1 lot at cold-start under max_lots=5 while weaker signals stay flat
(see `default_kelly_frac_floor` comment for the arithmetic). Exposed
as CLI flags + sweep-grid base/cell overrides.
Regression test `decision_floor_coldstart` proves:
- default floors (0.20/0.10) fire a 1-lot buy with p_h=0.8 and zero state
- zero floors reproduce the original noop bug
Also moves `aggregate` step to the GPU pool because fxt-backtest is
dynamically linked against libcuda.so.1 (the ci-compile-cpu hosts
don't expose CUDA driver libs).
Verified locally on RTX 3050 Ti — workspace cargo check passes, both
regression tests pass, trunk save/load roundtrip still passes.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Production smoke panicked at trunk.rs:mamba2_l1() during
PerceptionTrainer::from_checkpoint. Root cause: load_checkpoint
called new_random which left mamba2_stack_1/2 = None, then tried
to upload weights into None blocks.
PerceptionTrainer::new was the only caller that populated the
Mamba2 stacks (X3/X5 migrations stopped at the accessor methods
but never moved construction). Inference paths that didn't go
through PerceptionTrainer::new — exactly what fxt-backtest does
via from_checkpoint → load_checkpoint — hit the gap.
Move Mamba2Block::new + state allocation into CfcTrunk::new_random.
PerceptionTrainer::new now sets up its optimizer + scratches against
the trunk-owned blocks via existing accessors (mamba2_l1/mamba2_l2).
Extends the trunk roundtrip test to:
- assert mamba2 weight slices are non-empty (catches empty-buffer
regressions that would let bit-equivalence pass trivially)
- bit-equivalence-check Mamba2 stack 1 + stack 2 weights through
save -> load, reproducing the exact failure path that hit prod.
Verified locally on RTX 3050 Ti — full workspace cargo check passes,
test passes with non-empty weight slices.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Both tests exercised CfcTrunk::capture_graph_a + perception_forward_captured
+ snapshot_hidden, which feed V1-shaped CfC weights. After X8 the trunk's
CfC was reshaped to v2 layout (cfc_n_in=HIDDEN_DIM); the V1 forward path
now feeds FEATURE_DIM input into HIDDEN_DIM-shaped CfC — runtime garbage.
The methods themselves are still in trunk.rs (marked dead-code by rustc).
A deeper cleanup pass — deleting the V1 weight fields (heads_w_d, proj_*),
V1 per-step scratches, V1 cubin function handles, and the V1 forward
methods themselves — is a follow-up commit when fresh.
Verification: ml-alpha + ml-backtesting + fxt-backtest all build clean.
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.
X11 inference interface for the deployability backtester. Two new
public methods on PerceptionTrainer:
forward_only(snapshots) -> Vec<f32>
Forward-only pass over a K-snapshot window. Wraps evaluate() with
dummy zero-valued labels and discards the loss. Returns the same
[K, B, N_HORIZONS]-shaped probability output as evaluate_batched.
from_checkpoint(dev, cfg, path) -> Result<Self>
Build a PerceptionTrainer ready for inference: constructs a fresh
trainer (with random init) to wire up kernel handles + grad
buffers + scratches, then overwrites the trunk's weights from the
Checkpoint file via CfcTrunk::load_checkpoint. The optimizer +
grad buffers stay allocated — unused at inference, but allocating
them keeps the struct invariant uniform. A leaner inference-only
struct can be added later if memory matters.
Architectural note: the trunk is the source of truth for weights (X1-X9
established that). PerceptionTrainer is the kernel-launch adapter that
drives forward + backward over those weights. Inference doesn't need a
separate forward path on the trunk — the trainer's evaluate_batched
already does the full chain correctly. fxt-backtest can now construct a
PerceptionTrainer in inference role via from_checkpoint + call
forward_only per decision.
Verification: ml-alpha + ml-backtesting + fxt-backtest build clean.
ml-alpha lib tests: 33 pass.
