Replaces the flawed Phase F+G LobEnv trait + LobSimEnvAdapter +
MockLobEnv fixture with a GPU-pure env-interaction path.
DELETED:
- LobEnv trait (host-method-per-batch surface)
- MockLobEnv fixture (the toy-bandit pattern that hid the production defects)
ADDED:
- RlLobBackend trait (src/rl/reward.rs): narrow device-oriented
surface (apply_snapshot, pos_and_market_targets_mut, pos_d,
pos_bytes, step_fill_from_market_targets, n_backtests). Single
purpose: break the ml-alpha ↔ ml-backtesting dep cycle. No
mocking layer.
- 3 new GPU-pure kernels (cuda/):
* extract_realized_pnl_delta.cu — reads pos.realized_pnl +
position_lots from device Pos array, writes per-batch
(reward delta, done flag), updates trainer's prev_* buffers.
* apply_reward_scale.cu — element-wise rewards *= ISV[406].
* actions_to_market_targets.cu — 9-action grid → market_targets
{side, size} on device, reading current position_lots for
conditional Flat-from-Long/Short.
- LobSimCuda impls RlLobBackend (ml-backtesting/src/sim/mod.rs)
plus a new step_fill_from_market_targets entry that runs
submit_market_immediate + step_pnl_track without the host-side
targets-vec build. pos_and_market_targets_mut returns disjoint
field borrows (&pos_d, &mut market_targets_d).
- IntegratedTrainer: 3 cubin includes, 3 module/function fields,
launch_apply_reward_scale launcher, prev_realized_pnl_d /
prev_position_lots_d / rewards_d / dones_d buffers,
step_with_lobsim signature switched to RlLobBackend, body's
Step 5 rewritten as kernel-driven (no per-batch host loop, no
individual submit_action calls).
R6 PARTIAL — work R7 lifts:
- Thompson + log_pi stay host (R7 uses R4 kernels)
- mean_abs_pnl EMA stays host (R7 uses R3 ema_update_on_done)
- Advantage/return stays host (R7 uses R3 compute_advantage_return)
- 4 final DtoH copies for step_synthetic's host-slice signature
(R7 lifts step_after_encoder_forward to device-buffer args)
The "DtoH the device rewards/dones at end" comment in
step_with_lobsim documents the boundary. After R7, hot path has
zero host loops other than the now-unused Thompson host loop
(which R7 retires in favour of rl_action_kernel from R4).
Tests:
- integrated_trainer_smoke.rs: rewritten — real LobSimCuda with
synthetic book, one step_with_lobsim call, asserts finite losses
+ λs sum to 1. Smoke gate, not a convergence test.
- dqn_toy.rs, ppo_toy.rs: retired (empty stubs documenting
rationale). Files preserved so rename history survives. The
MockLobEnv toy-bandit pattern intrinsically couldn't catch the
production defects #1, #2, #5 that motivated this rebuild.
ml-backtesting added as ml-alpha dev-dep (cycle-safe: production
direction is ml-backtesting → ml-alpha; dev edge only loads when
building ml-alpha's own tests/examples).
Build cache-bust v29. cargo check + cargo build for all R-phase
tests green (pre-existing heads_bit_equiv index-OOB persists).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Ports two scope-complete fixes from the ml-alpha-phase-f-g-flawed
reference branch:
A7 fee model:
- order_match.cu::submit_market_immediate gains 2 new kernel args
(cost_per_lot_per_side, total_fees_per_b) mirroring the per-fill
fee deduction in resting_orders.cu::apply_fill_to_pos:213-219.
Fee deducts from pos.realized_pnl on EVERY fill (open, scale-in,
counter); total_fees_per_b accumulates for telemetry.
Single-writer-per-block plain += is safe — line 79 has
`threadIdx.x != 0 → return` (feedback_no_atomicadd).
- LobSimCuda::submit_market launch updated to pass the 2 new args.
- LobSimCuda::upload_cost_per_lot_per_side host API lets callers
configure ES-realistic fees (≈$1.25/contract/side). Default
alloc_zeros = $0; production decision-policy path is unchanged
(uploads its own cost via step_decision_with_latency).
A8 loader pair API:
- MultiHorizonLoader::next_sequence_pair returns
(LabeledSequence, LabeledSequence) at adjacent anchors in the
same source file. anchor_t sampled from [min_anchor, max_anchor−1)
so anchor+1 also fits the upper-bound. Counts as ONE yielded
sequence against n_max_sequences.
- next_sequence_random and next_sequence_pair share a new private
helper build_sequence_at(lf, anchor) -> LabeledSequence that
contains the multi-resolution windowing logic. Single source of
truth for the build (feedback_single_source_of_truth_no_duplicates).
- next_sequence's caller-facing contract (random anchor, one
sequence per call) is unchanged — alpha_train.rs supervised
pipeline keeps working as-is.
No new local tests this phase per the rebuild plan (R2): the fee
deduction with default cost=0 is a no-op for existing callers, and
G7 in R6 covers the with-fees path via a rebuilt reward_calibration
test driving LobSimCuda directly (no LobEnv adapter).
cargo check -p ml-alpha -p ml-backtesting + cargo build --tests on
ml-backtesting both green; baseline test suite unaffected.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
First commit of the horizon-rebase refactor (full plan at
docs/superpowers/plans/2026-05-22-horizon-rebase-n3.md). Motivated by
empirical evidence from this session's investigations:
- C1: ES MBP-10 input ACF dies by L=300; K=1000/6000 are below noise floor
- D3: binarized labels ρ(30,6000)=0.06; K=6000 carries no signal
- Heads_w1 mask experiment: h6000 AUC collapsed 0.73→0.54, confirming
long-horizon prediction needs cross-channel integration that 5-horizon
bucket structure couldn't provide
New horizon set [10, 100, 1000] matches C1's empirical autocorrelation
elbows (~50ms, ~3.5s, ~45s wall-clock at 74µs median Δt).
Scope of this commit:
- crates/ml-alpha/src/heads.rs: N_HORIZONS 5→3, HORIZONS=[10,100,1000]
- crates/ml-alpha/src/cfc/bucket_routing.rs: MAX_BUCKET_DIM 28→96,
BUCKET_DIM_K=[43,43,42], BUCKET_CHANNEL_OFFSET=[0,43,86,128]
- crates/ml-alpha/src/data/loader.rs: 5 hardcoded [T;5] types migrated
- crates/ml-backtesting/src/sim/mod.rs: broadcast_alpha signature,
local N_HORIZONS re-exports ml_alpha::heads::N_HORIZONS
- crates/ml-backtesting/src/harness.rs: local N_HORIZONS re-export,
per-horizon log labels, MultiHorizonLoaderConfig.horizons
Tests + examples + CUDA kernels still need migration (subsequent commits
per plan tasks 2-8).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Anti-calibration evidence (smoke ckpts across architecture variants
showed higher reported conviction → MORE negative PnL, -15.6 → -145.2)
traced to a magnitude/direction mismatch in the sizing formula:
conv_ema = EMA(|conviction_signed|) # magnitude smoothing
target_lots = sign(conviction_signed) # INSTANTANEOUS sign
× conv_ema × max_lots # × SMOOTHED magnitude
When per-horizon directions disagree event-to-event, the magnitude EMA
stays high (|x| EMA can't cancel) but the sign flips on noise. Result:
big trades in random direction whenever the 5 horizons disagree.
Replace with a single signed-conviction EMA:
conv_signed_ema = EMA(conviction_signed)
target_lots = sign(conv_signed_ema)
× |conv_signed_ema| × max_lots
Now BOTH sign AND magnitude come from the same smoothed signal. When
directions disagree, signed EMA → 0 collapses size to zero naturally.
Atomic migration across decision_policy.cu (2 kernels), pnl_track.cu
(open-branch reads the SIGNED buffer and takes fabsf for the magnitude-
semantics open_trade_state offsets that composite_exit_check forms ratios
from), and src/sim/mod.rs (renamed device buffers + accessor + 4 launch
sites). No legacy aliases per feedback_no_legacy_aliases; no parallel
buffers per feedback_single_source_of_truth_no_duplicates; α floor 0.4
preserved per pearl_wiener_alpha_floor_for_nonstationary; first-
observation bootstrap preserved per pearl_first_observation_bootstrap.
ml-backtesting lib: 33 passed.
GPU oracle: signed_conviction_ema_collapses_on_sign_flips passes —
observed steady-state shows |signed_ema| ≈ 0.205 (post-fix) producing
|target_lots| = 2 every tail event, vs the pre-fix bug's predicted
|target_lots| = 7 every tail event. Existing
multi_horizon_conviction_cancels_on_disagreement still passes (zero-
cancellation regime is even cleaner under signed EMA).
ml-alpha lib: 33 passed (no regression).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Adds --per-horizon-logit-sampling-stride CLI flag to fxt-backtest run/sweep
+ harness instrumentation that DtoHs alpha_probs every N events and
accumulates per-horizon mean/stddev/min/max + 10-bin histogram.
Emitted at end of CRT.diag block as:
per_horizon_logit_diag h{N}: count=... mean=... std=... min=... max=... hist=[...]
Goal: distinguish whether Smoke 2's uniform mean_run_len 1.04x across
horizons is caused by uniform per-horizon LOGITS (-> Task 11 scope
incomplete; heads_w1/w2/gate/main need block-diagonal too) or by
DIFFERENTIATED logits with collapsed downstream (-> different fix).
Default stride=None (disabled, zero overhead). With stride=1000 on
2M events, ~2000 DtoD stops x ~50us = ~100ms total overhead per run.
Per feedback_no_htod_htoh_only_mapped_pinned: uses mapped-pinned buffer
(single 5xf32 allocation reused on every sampling tick).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Driven by ffr59 (commit b44a97ff9) findings: all 5 horizons flip
direction every 2.5 events; win rate flat 24% across conviction; mean
PnL anti-correlated with conviction. Hypothesis: per-event alpha output
is high-frequency noise on top of a slower signal. If true, smoothing
input alpha BEFORE the conviction formula should reduce direction flips
and recover signal.
Adds Wiener-α adaptive EMA on raw alpha_probs[h] for each horizon,
applied BEFORE the multi-horizon conviction formula. Floor at 0.1
(stronger than the 0.4 floor on the output-side conviction EMA — this
tests whether INPUT smoothing has different impact than OUTPUT smoothing).
Three new device slots:
- alpha_ema_per_b_per_h (per-horizon EMA state)
- alpha_diff_var_per_b_per_h (variance of changes)
- alpha_sample_var_per_b_per_h (variance of value)
The CRT.diag Group A direction-flip counter still reads RAW alpha_probs
so we have a head-to-head comparison: raw flip rate vs smoothed flip rate.
Group E adds the smoothed-direction counter + mean run length.
End-of-run log adds one line per horizon:
crt_diag h<X> smoothed: flips=Y mean_run_len=Z events (vs raw F / M)
If smoothed mean_run_len >> raw mean_run_len: hypothesis is RIGHT, the
input had signal under the noise. Next step would be to make this an
operational EMA in the controller.
If smoothed and raw are similar: hypothesis is WRONG, per-event output
is genuinely noisy. Next step would be to investigate model training
(horizon collapse) OR the AUC=0.66 measurement definition.
Either way, definitive result from one smoke run.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Pure-instrumentation commit. Zero behavior change. The Gate CRT.1 smoke
p9cnk (commit 1e656948b) produced 222k trades and 29688% drawdown,
worse than the failed CRT.A. The structural fixes (multi-horizon
conviction, no-trade band, composite exit) all work as designed — but
the trade cadence is fundamentally too fast for the signal/cost ratio.
Before another round of threshold tuning, measure the structural truth.
Adds device-side counters dumped at smoke close (single end-of-run
memcpy_dtoh batch in read_diagnostics(), NOT per-event):
Group A — per-horizon signal persistence:
- flip_count[h]: total direction-sign flips
- sum_run_length[h]: cumulative events between flips (u64)
- run_length_hist[h]: 5-bucket log-scale histogram (1-9, 10-99,
100-999, 1k-9.9k, 10k+)
Output: mean_run_len per horizon = signal half-life proxy
Group B — conviction smoothed-EMA distribution:
- conv_hist[10]: 10-bucket histogram over [0, 1)
Output: shape of conviction distribution at decision time
Group C — hold-time distribution:
- hold_hist[6]: <1s, 1-10s, 10-60s, 1-10m, 10-60m, >1h
Output: empirical distribution of trade hold times
Group D — outcome by entry-conviction:
- outcome_n[10]: trades per conviction bucket
- outcome_sum_pnl[10]: cumulative realized PnL per bucket (price-units)
- outcome_n_wins[10]: wins per bucket
Output: validates "high conviction = better trades" hypothesis OR
surfaces signal miscalibration
Per-backtest memory: ~330 bytes. Negligible vs existing state.
All counters single-writer-per-block (threadIdx.x == 0 convention; no
atomicAdd per feedback_no_atomicadd). All updates kernel-internal (no
host roundtrip per event). The end-of-run dump uses memcpy_dtoh in
read_diagnostics() — called once at harness termination from the
post-stop_ctrl_counters block, never per-event. Per
feedback_no_htod_htoh_only_mapped_pinned, the hot path is untouched.
N_HORIZONS = 5 (heads.rs canonical {30, 100, 300, 1000, 6000}); the
plan text says 4 but the workspace constant is 5 — used the code value.
Output at end of cluster smoke as a `crt_diag` log block:
crt_diag h30: flips=X mean_run_len=Y events
crt_diag h30 run_length_hist: 1-9:N 10-99:N 100-999:N ...
...
crt_diag conv_ema_hist: [0.0-0.1]:N [0.1-0.2]:N ...
crt_diag hold_time_hist: <1s:N 1-10s:N ...
crt_diag outcome_by_entry_conv[0.0-0.1]: n=N win_rate=X.X% mean_pnl_pu=Y.YYY
...
The output drives the next round of CRT.1 threshold design (delta_floor,
composite exit factors, min-hold cooldown) from data instead of guesses.
Per pearl_adaptive_not_tuned and pearl_controller_anchors_isv_driven:
thresholds will become ISV-derived in CRT.2; CRT.1 must first establish
defensible defaults from the empirical signal characteristics.
Verified: stop_controller (22/22 pass), decision_floor_coldstart (3/3
pass), threshold_and_cost (3/3 pass), parallel_sim_correctness (1/1
pass), lob_sim_integrated_fuzz (3/3 pass). Two lob_sim_fixtures
failures (fix_decision_alpha_buy_close, fix_decision_program_h4_only)
were pre-existing at HEAD 1e656948b — not caused by this commit.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per spec §4.3 + plan C1.4. Safety backup for the primary exit path
(target → 0 from collapsed multi-horizon conviction, C1.2). Catches
the edge case where conviction stays bounded above zero but the trade
is deep in unrealized loss, or short/long horizons disagree for many
events while threshold-gate logic still permits sizing.
Three terms, max-of-three composite:
degradation_term = (1 − conv_ema_now / conviction_at_entry) / 2.0
disagreement_term = disagree_consecutive_events / 32.0
pnl_decay_term = max(−pnl_adjusted_conv_ema, 0) / 1.0
exit_signal = max(deg, dis, pnl_decay)
Fires force-flat (market_targets = (3, 0)) when exit_signal >= 1.0.
State-update flow:
- pnl_track open branch: write conviction_at_entry (offset 20),
conviction_ema (offset 24, initialised to entry value),
pnl_adjusted_conviction_ema (offset 44, = 0),
peak_unrealized_pnl (offset 48, = 0). Takes new conviction_ema_per_b
read-only arg to snapshot at open.
- decision kernel intra-event (when pos open): update offsets
24/44/48/56 (disagreement counter) via update_open_trade_trajectory.
- composite_exit_check device helper: reads offsets 20/24/44/56,
writes (3, 0) on fire. Sibling of stop_check_isv with same return
semantics; called after trajectory update so all four reads see
the freshly-written values.
Fields not read in CRT.1 (offsets 28-43 per-horizon EMA, 52 degradation
counter, 60 horizon_idx_dominant) are NOT written by the open branch —
alloc_zeros covers init. Per feedback_no_hiding: don't populate unused
fields. Also fixes a latent C1.1 bug: pnl_track close branch read
horizon_idx from offset 20 (now conviction_at_entry, f32); moved the
read to offset 60 per the documented layout.
Threshold-bootstrap defaults (will be ISV-derived in CRT.2 alongside
the rest of the controller anchors per pearl_controller_anchors_isv_driven):
degradation_factor = 2.0 (50% conviction decay vs entry → ≤0.5 term)
disagreement_factor = 32.0 (32 events of short/long disagreement)
pnl_decay_factor = 1.0 (pnl_adj_conv structurally ∈ [-1, +1])
Under these defaults disagreement_term is the only term that can fire
on its own — degradation_term caps at 0.5 (conv_ema/conviction_at_entry
≥ 0) and pnl_decay_term caps at 1.0 in the limit. The three terms still
add evidence in superposition; CRT.2 will retune factors against ISV
percentiles. Unit-test coverage validates the disagreement path
(composite_exit_signal_fires_on_disagreement).
Per pearl_no_host_branches_in_captured_graph — all composite logic
lives inside the decision kernel; no host roundtrip in the per-event
hot path. Per feedback_no_htod_htoh_only_mapped_pinned — no new
memcpy or stream.synchronize() in step_pnl_track or
dispatch_latent_market_orders.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per spec §4.0 + plan C1.3. Without this gate, every fractional change
in target_lots seeds an order via target-delta semantics; combined with
the continuous controller invocation (every event after A1) this
generates hyperactivity even with the multi-horizon §4.4 conviction
formula (C1.2). v2 Gate-1 failures confirmed: scalar EMA + amplification
clamp + per-event controller = 155k trades on a 2M-event smoke.
The band absorbs fractional target adjustments so most events produce
no order. When the signal genuinely shifts and target moves enough to
cross the band, the kernel seeds. Continuous evaluation, sparse action.
Greenfields atomic refactor — single config knob threaded through every
layer in one commit:
SweepBase.delta_floor: f32 (YAML, default 1.0 = 1 lot)
SimVariant.delta_floor: Option<f32> (per-variant override)
ResolvedSimVariant.delta_floor: f32
UniformSimParams.delta_floor: f32
BatchedSimConfig.delta_floor: Vec<f32>
LobSimCuda.delta_floor_d: CudaSlice<f32>
seed_inflight_limits_batched kernel param + skip-if-below-floor logic
New accessor: LobSimCuda::read_inflight_count(b) — counts active != 0
limit slots for backtest b. Not cfg(test); follows existing read_limit_slot
pattern.
Test: no_trade_band_blocks_micro_delta verifies that a second decision
with the same strong-bullish alpha (same target, effective = in-flight
lots) produces delta=0 and does not re-seed. max_lots=1 keeps the
arithmetic unambiguous.
Existing tests: all UniformSimParams struct literals updated with
delta_floor=0.0 (band disabled) so existing behaviour is preserved.
harness.rs from_uniform path uses delta_floor=1.0 (production default).
Hot-path discipline: the delta_floor_d upload happens once per run in
the existing config-upload block (alongside threshold, cost, max_hold_ns);
the kernel reads the device slot per-event without any host roundtrip. No
memcpy_htod / dtoh / dtov / synchronize introduced on the per-event path.
Per pearl_controller_anchors_isv_driven, this floor is currently a
config constant; CRT.2 (Phase 2) makes it ISV-derived from rolling
spread cost / signal volatility.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per spec §4.4 + plan C1.2. Replaces v2 A2 (1d889d2de) + A2.1
(fe2498769) scalar approach that failed Gate 1 catastrophically on
smokes vjmwc and lkrdf (155k trades, 9100% drawdown, Sharpe -15.7).
The structural fix: direction comes from per-horizon SIGNED sum, so
disagreeing horizons cancel. The scalar EMA on max(|p-0.5|) could only
smooth magnitude — it could not smooth direction jitter. Only this
multi-horizon weighted formula can.
Formula (spec §4.4):
weight_h = max(isv[h].net_edge, ε) / (isv[h].var + cost²)
weighted_h = magnitude_h × weight_h × direction_h
conviction_signed = sum(weighted_h) / sum(|weight_h|)
conviction_ema = Wiener-α EMA(|conviction_signed|) (reuses A2 slots)
target_lots = round(sign(conviction_signed) × conviction_ema × max_lots)
Pearl conformance:
- pearl_controller_anchors_isv_driven: weights from ISV state
- pearl_one_unbounded_signal_per_reward: net_edge/(var+cost²) is the
single unbounded factor; magnitude × direction is bounded in [-1,1]
- pearl_zscore_normalization_for_magnitude_asymmetric_signals: divide
by total_abs_weight for scale invariance
- pearl_trade_level_vol_for_stop_distance: cost² is variance bootstrap
- pearl_blend_formulas_must_have_permanent_floor: ε = cost · 0.01 on
net_edge — no cold-start regime switch
- pearl_audit_unboundedness_for_implicit_asymmetry: net_edge is
one-sided positive (losing-edge horizons drop out, not flip sign)
- pearl_wiener_alpha_floor_for_nonstationary: α floor at 0.4 (unchanged
EMA mechanism, repointed at the multi-horizon scalar)
DELETED:
- The `final_size *= (conv_ema/raw_max_conv)` aggregate rescale step
(the v2 bug that amplified weak signals)
- The per-horizon `signed_sizes[h]` Kelly + Sharpe-weight aggregation
loop in decision_policy_default (replaced by direct §4.4 formula)
- Three obsolete tests asserting on the old scalar path:
* conviction_ema_smooths_micro_oscillations
* conviction_ema_does_not_lag_reversals
* conviction_ema_rescale_never_amplifies_weak_signal
* cfg_high_kelly_floor helper (only used by deleted tests)
- cold_start_with_zero_floor_reproduces_old_bug (tested OLD kelly_floor
semantics that v3 doesn't expose)
- cold_start_persistent_bullish_now_closes (was passing under v2 via
target oscillation from integer rounding; v3 produces stable target
on stable signal — closes need price movement)
- alpha_noop_side_2_preserved (anchored on the v2 kelly_frac_floor
integer-truncation path that doesn't exist under v3)
Applied to BOTH decision_policy_default and decision_policy_program
(bytecode VM at OP_WRITE_ORDER). Per feedback_no_partial_refactor —
both consumers migrate atomically. The VM's per-horizon emit + aggregate
opcodes still execute but their stack `final_size` is unused at
OP_WRITE_ORDER (the §4.4 conviction-driven target replaces it). The
stack's attribution mask is still consumed for entry crediting.
The three conviction-EMA device slots (conviction_ema_d, _diff_var,
_sample_var) STAY in LobSimCuda. The Wiener-α EMA mechanism is reused
on the multi-horizon conviction scalar; the device-helper signature
changed from taking alpha_probs[] to taking the precomputed scalar
|conviction_signed|.
Kernel ABI change: decision_policy_default no longer takes
target_annual_vol_units / annualisation_factor / kelly_frac_floor /
sharpe_weight_floor (the §4.4 formula doesn't use them). Removed from
both the CUDA signature and the launch builder in sim/mod.rs. The
bytecode VM (decision_policy_program) still consumes those params for
its OP_EMIT_PER_HORIZON_SIZE / OP_AGG_WEIGHTED_SHARPE opcodes — their
device slots remain allocated.
Test: multi_horizon_conviction_cancels_on_disagreement verifies the
structural fix — bullish [0.7,0.3,0.7,0.3,0.5] horizons cancel to
no-trade output (side=2/3, size=0, conv_ema≈0).
Existing tests rebased: cfgs that used cost_per_lot_per_side=0.0 now
use 1.0 (the §4.4 formula's eps_edge floor requires cost > 0). Each
affected test file documents the rebase inline. Tests that were
anchored on observed v2-kernel values (per
pearl_tests_must_prove_not_lock_observations) are deleted; tests with
genuine invariants (boundedness, direction-sanity, stop-controller
behavior) are preserved.
Hot-path discipline: no memcpy_htod/dtoh/dtov/synchronize introduced.
Only kernel-internal computation; one launch arg removed from the
default decision kernel.
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>
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>
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>
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>
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>
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>
Both decision kernels now call the same __device__ helper at the top
of per-backtest dispatch (after the program_lens early-return). Single
source of truth for stop logic across default and bytecode-VM paths.
Parity test ensures they produce identical market_target output for
matched scenarios.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Skeleton helper called from decision_policy_default at top of
per-backtest dispatch. Returns 0 (no-op) on flat positions; trigger
logic for open positions added in Tasks 4-6. Single source of truth
for stop logic across decision_policy_{default,program}.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
cudarc kernel launches require `unsafe` blocks — the driver API has no
way to type-check kernel args against the cubin signature. The
workspace-wide `-W unsafe-code` lint produces noise on every launch
site (10+ blocks); suppressing at the module level keeps the lint
useful elsewhere in the crate.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per-event Wiener-α=0.4 EMA on |Δmid| with first-observation bootstrap.
Floor source for the SL/trail-distance controller (spec §6). Thread 0
of the broadcast-snapshot kernel handles the update per backtest;
threads 1..9 still handle the 10 level writes.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Three CudaSlice<f32> per-backtest slots (prev_mid_d, atr_mid_ema_d,
trail_hwm_d) plus test-only accessors. Foundation for the
ISV-driven stop controller; no behavioral change until subsequent
tasks wire them into kernels.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Adds LobSimCuda::read_total_trade_count (cheap n_backtests*4 byte DtoH of
the trade-log head counters) and emits the sum on each PROGRESS_EVERY
eprintln. Lets future smokes (Q2 CBSW validation, P7 sweep) see trade
firing mid-run instead of waiting for write_artifacts at harness end.
Heads count kernel writes (clamped to ring cap), so the sum is a
monotone trade counter across the sweep — drops back only on harness
reset which doesn't happen mid-run.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The threshold-tuning smoke at 81decf40f produced n_trades=0 despite
74.6% of decisions having max_conv ≥ 0.30 — the linear-weighted-mean
aggregator in decision_policy_default is structurally dilution-bound
at cold-start (per spec §1).
Q1 stopgap: when sim_variants[i].use_cold_start_stopgap = true, the
harness uploads a max-confidence Strategy bytecode program for that
backtest, routing decisions through decision_policy_program with
OP_AGG_MAX_CONFIDENCE. Existing kernel; zero CUDA changes.
Field additions (atomically across BatchedSimConfig + UniformSimParams
+ ResolvedSimVariant + SweepBase.SimVariant) — every UniformSimParams
literal migrated to include use_cold_start_stopgap: false (default).
The sweep YAML's sim_variants entry sets it to true only for the
validation run; production deployability uses Q2's kernel fix instead.
Sweep YAML (config/ml/sweep_smoke.yaml) flipped to use_cold_start_stopgap=true
at threshold=0.0, cost=0.125 — same anchor as the threshold-tuning
smoke that produced n_trades=0, for direct comparison.
This is a VALIDATION step. Cluster smoke at this commit MUST produce
n_trades > 100 + finite metrics. Q2's kernel CBSW immediately follows
and deletes this entire stopgap atomically (field, harness branch,
YAML setting, every literal).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
snapshot_realized_pnl / snapshot_position_lots / snapshot_open_horizon_mask
were the host-side read paths for the close-detection loop that P3
(3836e2578) replaced with snapshot_pos_state + detect_close_transitions_
batched kernels. The helpers stayed dead-code-warned after P3; per
feedback_no_legacy_aliases + feedback_no_hiding, delete them.
submit_market_fn is NOT dead — pub fn submit_market wraps it and has 4
test callers (lob_sim_fixtures, lob_sim_fuzz, ring3_replay × 2). Kept.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Sweep YAML now supports the batched flow per spec §3.3 + Task 6:
- SweepBase.sim_variants: Vec<SimVariant> — list of (cost, latency,
threshold, ...) variants. When non-empty, each cell runs ONE harness
at n_parallel=variants.len() with BatchedSimConfig::from_grid instead
of the legacy one-harness-per-cell fan-out.
- SweepBase.data_template: Option<String> — when set with `{window}`
placeholder, each cell's `window` field interpolates the per-cell
data path. Replaces single scalar `data` for the windowed flow.
- SweepCell.window: Option<String> — window identifier (e.g., "2025-Q2").
- SimVariant: threshold + cost_per_lot_per_side required (the spec's
primary axes); other fields optional overrides on top of SweepBase
scalars.
New runner pieces:
- BatchedSimConfig::from_grid(&[ResolvedSimVariant]) in
crates/ml-backtesting/src/sim/batched_config.rs.
- ResolvedSimVariant — per-variant fully-resolved sim params.
- resolve_sim_variants(&SweepBase) in main.rs — layers per-variant
overrides over base scalars.
- run_batched_cell() in main.rs — builds the harness with
sim_config_override + variant_names plumbed through. Writes per-
backtest artifacts to sim_<variant_name>/ subdirs (spec §3.3).
Harness side:
- BacktestHarnessConfig gains variant_names + sim_config_override
Option fields. When sim_config_override is Some, harness uses that
directly instead of building from_uniform off scalar cfg. When
variant_names is Some, write_artifacts uses sim_<name>/ instead of
cell_NNNN/ subdirs. Both None preserve legacy single-cell behaviour
(smoke, fixtures unchanged).
YAML configs:
- config/ml/sweep_threshold_tuning.yaml: 1 cell (W0) × 8 sim_variants
(p60-p95 in 5pt steps) with cost=0.125 (1-tick anchor). Threshold
pre-registration pass.
- config/ml/sweep_deployability.yaml: 4 cells (W1-W4) × 140 variants
each (7 costs × 4 latencies × 5 thresholds). Generated by
scripts/generate_sweep_variants.py — placeholder threshold values
(p60-p95) until threshold-tuning publishes calibrated absolutes to
config/ml/v2_prod_thresholds.json.
Deferred to P7 (operational glue):
- argo-lob-sweep.sh adaptation for the batched flow (cells = windows,
not sim-variants; one Argo task per window invokes `fxt-backtest sweep`
end-to-end inside the pod rather than `fxt-backtest run`).
Regression: all 7 existing CUDA tests pass through the new harness
construction path (sim_config_override = None → from_uniform fallback).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Adds the two sweep axes that the spec's deployability grid needs but
were missing from the kernels:
Threshold gate (decision_policy.cu, both kernels):
- New per-backtest `threshold_per_b` array kernel arg.
- Pre-Kelly prelude: if max_h |alpha[h] - 0.5| * 2 < threshold[b],
emit noop and return. Kept deterministic from alpha alone so the
threshold pre-registration step (p60-p95 absolute calibration on a
validation window, future P6) reflects exactly what gets gated in
deployment.
Per-fill cost integration (resting_orders.cu / apply_fill_to_pos):
- apply_fill_to_pos signature grows three args: b, cost_per_lot_per_side_per_b,
total_fees_per_b. Single insertion point at line 90.
- After the close-leg realized_pnl math runs (so the gross unwind P&L
is preserved), deduct fill_cost = filled_lots * cost_per_lot_per_side[b]
from pos.realized_pnl AND accumulate into total_fees_per_b[b].
- Net-of-cost semantics: isv_kelly_update_on_close reads realized_pnl
delta which is now net of cost — Kelly state learns from realistic
return distribution.
- All 3 apply_fill_to_pos call sites in step_resting_orders updated.
order_match.cu's submit_market_immediate path is dead code in the
post-P1 flow (everything routes through seed_inflight_limits_batched
→ step_resting_orders → apply_fill_to_pos) so not touched here.
BatchedSimConfig + UniformSimParams + BacktestHarnessConfig gain
threshold + cost_per_lot_per_side fields. All UniformSimParams
constructors in tests and main.rs updated with defaults (0.0, 0.0 =
gate disabled, frictionless).
Regression:
- threshold_gate_skips_low_conviction (p=0.51 + threshold=0.10 → noop)
- threshold_gate_allows_high_conviction (p=0.8 + threshold=0.10 → buy 1+)
- threshold_zero_is_passthrough (sanity)
- All P1+P2+P3 tests continue to pass via the new ABI.
cost_deducted_at_each_fill + kelly_state_sees_net_return end-to-end
tests deferred — they require a full submit_market → fill → close
sequence, which the production smoke exercises.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Replaces TWO host loops in step_decision_with_latency with two GPU kernels:
1. snapshot_pos_state — replaces the host-side snapshot_realized_pnl +
snapshot_position_lots + snapshot_open_horizon_mask trio (3 separate
memcpy_dtoh per decision). Now one kernel launch writes
prev_pos_lots_d / prev_realized_pnl_d / prev_open_horizon_mask_d
directly on the device.
2. detect_close_transitions_batched — replaces the host close-detect
loop that called read_pos per close-eligible backtest (up to
n_backtests memcpy_dtoh per decision). Now one kernel writes
closed_horizon_mask_d + realised_return_d on the device, and
isv_kelly_update_on_close consumes them with no host roundtrip.
At n_parallel=140 these two loops together accounted for ~350M+ small
host roundtrips per quarter. Combined with P2 the latency-path of
step_decision_with_latency is now fully GPU-resident.
isv_kelly_update_on_close kernel always launches (skips backtests with
mask=0 internally) rather than gating via a host any_close check.
All P1+P2 regression tests pass through the new GPU close-detect path
(at n=1 with uniform config + immediate-fill latency, no close happens
in the cold-start tests since they only check market_target; the close
path is exercised indirectly).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Replaces the host roundtrip loop at the latency path of step_decision_
with_latency with one GPU kernel launch. At n_parallel=140 the host
loop did up to 140 memcpy_dtoh + 140 seed_limit_order calls per
decision (~210M roundtrips per quarter at the threshold-tuning load).
The new kernel does the same work in one launch.
Per-backtest single-writer (threadIdx.x==0). Each backtest scans its
own MAX_LIMITS=32 slot range for an `active==0` slot. Slot allocation
is per-backtest (no cross-backtest atomics needed). Overflow path
increments pos.submission_overflow.
dispatch_latent_market_orders now takes &BatchedSimConfig (unused
inside — the latency_ns_d device buffer is already populated by
step_decision_with_latency's upload block from P1).
All 3 decision_floor_coldstart tests still pass via the new GPU path
(at latency_ns=0 the in-flight slot's arrival_ts==current_ts and gets
promoted on the next step_resting_orders, functionally equivalent to
the legacy submit_market_immediate kernel).
Independence test (different per-backtest latencies → different
arrival_ts) deferred to a later step where a read_first_inflight_arrival_ts
helper is added.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Migrates target_annual_vol_units, annualisation_factor, max_lots,
latency_ns, kelly_frac_floor, sharpe_weight_floor from scalar-broadcast
kernel args to per-backtest device arrays via BatchedSimConfig. Atomic
contract change per feedback_no_partial_refactor — kernel + sim + harness
+ all 3 existing test files migrate in this commit.
- crates/ml-backtesting/src/sim.rs → sim/mod.rs (directory module)
- crates/ml-backtesting/src/sim/batched_config.rs (NEW): BatchedSimConfig
+ UniformSimParams + validate(). from_uniform rebuilds the legacy
uniform-broadcast behaviour at n=1 (smoke/fixtures). from_grid lands
in P6 for the 140-variant sweep packing.
- LobSimCuda gains 6 per-backtest device buffers (target_annual_vol_units_d,
annualisation_factor_d, max_lots_d, latency_ns_d, kelly_frac_floor_d,
sharpe_weight_floor_d). step_decision_with_latency uploads from
BatchedSimConfig each call; both decision kernel launches now pass
per-backtest array pointers.
- decision_policy_default + decision_policy_program: scalar args become
const float* / const int* per_b arrays; first lines of each kernel
index by `b` into the arrays. Behaviour preserved at n=1 uniform.
- dispatch_latent_market_orders: reads latency per-backtest from
&BatchedSimConfig (host loop stays for P1; P2 replaces with kernel).
- step_decision_with_latency now ALWAYS dispatches through the latency
path; when cfg.latency_ns[b]=0 the in-flight slot's arrival_ts equals
current_ts and gets promoted immediately on next snapshot. Eliminates
the if/else branch and consolidates the launch path.
- harness.rs: BacktestHarness gains a sim_config field, built via
BatchedSimConfig::from_uniform at new() from the harness cfg's scalar
fields. The run loop passes &self.sim_config to step_decision_with_latency.
Regression coverage:
- parallel_sim_correctness::parallel_sim_equivalence_with_uniform_config
— n=8 with uniform config produces 8 bit-identical market_targets
(proves per-backtest indexing reduces correctly).
- Existing decision_floor_coldstart tests (3) all pass through the new
ABI — proves cold-start floor + variance-cap gate behaviour preserved.
- parallel_sim_independence_per_backtest deferred to P2 (needs
read_first_inflight_arrival_ts helper that depends on LimitSlot layout).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
