Per spec §7 (v3 architecture reset, promoted from v2 Phase B into CRT.1).
Single-cache-line layout. Every reader and writer of open_trade_state
migrates in this commit per feedback_no_partial_refactor.
New 64-byte layout (existing 24-byte fields preserved at original offsets):
Offset Size Field
0 8 entry_ts_ns (u64)
8 4 entry_px_x100 (i32)
12 4 entry_size (i32)
16 4 realized_at_open (f32)
20 4 conviction_at_entry (f32) ← NEW
24 4 conviction_ema (f32) ← NEW
28 16 conviction_per_horizon_ema [4 × f32] ← NEW
44 4 pnl_adjusted_conviction_ema (f32) ← NEW
48 4 peak_unrealized_pnl (f32) ← NEW
52 4 degradation_consecutive_events (u32) ← NEW
56 4 disagreement_consecutive_events (u32) ← NEW
60 1 horizon_idx_dominant_at_entry (u8) ← NEW
61 3 pad
The 24-byte prefix is unchanged so downstream readers (stop_check_isv in
decision_policy.cu, max_hold event-rate check in resting_orders.cu) need
only update the stride constant. The new fields at offsets 20..63 are
populated by subsequent CRT.1 tasks (C1.2 multi-horizon conviction, C1.4
composite exit signal).
Open branch in pnl_track.cu zeros the new fields implicitly because
alloc_zeros on the device slot zero-initialises everything; subsequent
writes only touch the 0-23 byte range (existing fields). Close branch
reset-loop iterates OPEN_TRADE_STATE_BYTES which now zeros all 64.
Files changed:
crates/ml-backtesting/src/lob/mod.rs — pub const OPEN_TRADE_STATE_BYTES: usize = 64
crates/ml-backtesting/cuda/pnl_track.cu — #define OPEN_TRADE_STATE_BYTES 64
crates/ml-backtesting/cuda/resting_orders.cu — comment + stride literal 24→64
crates/ml-backtesting/tests/stop_controller.rs — open_trade_state_64_byte_layout test
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Empirical trigger: Phase A as designed in v2 failed Gate 1 catastrophically
on both smoke vjmwc (commit 3d8f12deb) and lkrdf (commit fe2498769 with
amplification clamp). Both produced ~155k trades (62× baseline), 9100%
max-drawdown (175× baseline), and Sharpe -15.7. The clamp had effectively
zero impact on the outcome, confirming the bug is structural not formulaic.
Architectural finding: the v2 phase split (A: continuous control / B:
signal-driven position management) is artificial. They are the same
mechanism. A scalar EMA on max(|p-0.5|) cannot smooth direction jitter
between horizons; only the multi-horizon ISV-weighted conviction formula
(v2 §4.4, deferred to Phase B in v2) is structurally coherent.
User-validated mental model: signal vector at any moment IS a forward
prediction of the trade trajectory; optimal position = inventory implied
by current beliefs; trade actions are sparse because most events
produce only fractional adjustments to a stable optimum. Continuous
evaluation, discrete action.
v3 phase structure (replaces v2 A/B/C/D):
CRT.1 Unified Inventory Controller (was A+B merged)
CRT.2 Adaptive Risk Envelope (was Layer C, unchanged)
CRT.3 Online Weight Adaptation (was Layer D, unchanged)
CRT.1 components:
- forward_step every event (already shipped: A0.5)
- decision_stride deleted (already shipped: A1)
- Multi-horizon ISV-weighted conviction (§4.4 promoted from Phase B)
- Wiener-α EMA on multi-horizon conviction (§4.2 promoted)
- target_lots = direction × |conviction_ema| × envelope_max
- No-trade band in seed_inflight: skip if |delta| < delta_floor (NEW)
- open_trade_state 24→64 expansion (§7 promoted from Phase B)
- Conviction-degradation exit emergent from target→0; composite-signal
safety layer (§4.3) preserved as circuit-breaker
What survives from v2 commits on the branch:
a0e81fbdf forward_step incremental SSM (A0.5)
92f8b10ed forward_step_into eliminates GPU↔CPU round-trip
045850e8f delete decision_stride field (A1)
3d8f12deb buffer-level seed bit-identity test
What gets superseded in CRT.1 implementation:
1d889d2de A2 scalar conviction-EMA with aggregate rescale — replaced
by multi-horizon §4.4 formula
fe2498769 A2.1 clamp on broken rescale — same reason
The three conviction-EMA device slots (conviction_ema_d et al.) STAY
in LobSimCuda — the EMA mechanism is reused on the multi-horizon
conviction scalar. The `final_size *= scale` rescale step is deleted.
Gate restructure:
v2 Gate 1 (Layer A standalone) — REMOVED
v2 Gate 2 (Layer B) — PROMOTED as Gate CRT.1
v2 Gate 3/4 — unchanged as CRT.2/CRT.3
Status: Design v3 — awaiting plan rewrite before implementation.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
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>
forward_only requires the full K=seq_len window on every call and resets
Mamba2 h_s2 to zero each invocation — no cross-call state carry. Additionally,
the call is inside the stride=200 gate in harness.rs (spec §3.2 claim "already
every event" is incorrect as of HEAD). Moving to stride=1 without a forward_step
implementation would be a ~200x GPU cost increase. A0.5 must implement
forward_step with persistent h_s2, a dedicated K=1 CUDA graph, and session-reset
hook. Memo documents exact file paths, line numbers, required struct/method
changes, and open questions for the A0.5 implementer.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per user direction: no fallbacks, greenfields, kernel updates acceptable.
Changes:
1. Task A0.5 ADDED — pre-planned task that fires if A0 investigation
finds forward_only is stateless K-window (Case 2) or has periodic
reset (Case 3). Refactors PerceptionTrainer to maintain persistent
SSM state and expose forward_step(snapshot, b) that advances by 1
event. Bit-identical to forward_only of equivalent window per a
new golden test in incremental_forward.rs.
This is NOT a fallback — it's a pre-planned task whose execution
is determined by actual investigation outcome. Case 1 → A0.5 is
a no-op. Case 2 or 3 → A0.5 fires in full. Either way the plan
proceeds without user-approval pause.
2. Task A2 reframed — Wiener-α conviction-EMA is a "load-bearing
component of continuous control," not a "minimal Phase B subset."
Without smoothing, event-rate trading is structurally incoherent.
3. Task A2 test fallback REMOVED — instead of "if helpers don't
exist, omit the unit test," the plan now says "add the helpers
properly to the public API." Step 1 audits the LobSimCuda public
API and adds missing accessors. Greenfields — accessors stay on
the API permanently if testing needs them.
4. Task A5 (conditional hyperactivity mitigation) DELETED. A2 is
designed correctly the first time with Wiener-α floor at 0.4. If
A4 cluster smoke shows hyperactivity, that's an A2 bug to fix
properly, not a tuning knob to nudge.
5. Scope contract updated — Phase A vs Phase B split is by code-path
responsibility (continuous control infrastructure vs multi-horizon
signal-driven policy), NOT by "shipping less to be safe."
6. Notes for the Implementer — explicit "No fallbacks" section
replaces the implicit-fallback language. STOP-and-notify-user for
Case 2 deleted (A0.5 handles it inline).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per docs/superpowers/specs/2026-05-20-continuous-reasoning-trader-design.md
(v2, commit 8ba8755b4). One plan per phase per writing-plans guidance —
Phase B/C/D plans written after each prior gate closes.
Phase A scope:
A0 - Investigate forward_only cost model (read-only research,
writes memo to docs/superpowers/memos/). STOP condition if
trunk forward is stateless K-window per call (Case 2) —
requires user approval to expand scope.
A1 - Atomic delete of decision_stride (greenfields, no compat).
Removes from SweepBase, ResolvedSimVariant, BatchedSimConfig,
BacktestHarness, all YAMLs.
A2 - Minimal Wiener-α conviction-EMA smoothing in decision_policy_*
kernels. Three new device slots: conviction_ema_d,
conviction_diff_var_ema_d, conviction_sample_var_ema_d. Floor
at 0.4 per pearl_wiener_alpha_floor_for_nonstationary.
First-observation bootstrap. Direction stays from raw alpha;
magnitude/sizing uses smoothed value.
A3 - Local compile + tests.
A4 - Cluster smoke + Gate 1 validation against spec §3.6 acceptance.
Memory entry on closure (green/red/yellow).
A5 - Conditional hyperactivity mitigation (raise floor 0.4→0.6 OR
add target-delta hysteresis). Fires only if Gate 1 reveals
trade-count balloon.
Out of scope (deferred): full §4.4 multi-horizon conviction formula,
§4.3 conviction-degradation exit, §5 envelope, §6 LoRA adapter, §7
open_trade_state expansion — all Phase B/C/D.
Pearl conformance:
- feedback_no_partial_refactor (atomic delete)
- feedback_no_legacy_aliases (no compat shim)
- feedback_push_before_deploy
- pearl_wiener_optimal_adaptive_alpha
- pearl_wiener_alpha_floor_for_nonstationary
- pearl_first_observation_bootstrap
- feedback_stop_on_anomaly (Gate 1 RED → diagnose, don't advance)
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Self-critical review of v1 (commit 0f3484325) found 12 issues. All
fixed inline. Spec marked Status: Design v2 with explicit changelog.
LOAD-BEARING FIXES
1. §4.4 conviction cold-start: removed backwards "uniform fallback"
branch that maintained trading authority when net_edge=0 across
all horizons (exactly when model has no edge). Replaced with ε
floor on net_edge_h. Single formula, smooth bootstrap, no regime
switch.
3. §4.3 exit logic: 5 OR conditions → 1 composite exit_signal scalar
(max of degradation, disagreement, pnl_decay terms). SL/trail and
8h circuit-breaker stay orthogonal as safety, not exit policy.
Attribution per term recorded in §7 state for observability.
6. §9 Gate 2: PF > 1.0 (2.5× improvement, unrealistic) → tiered
MUST (no-regression) + WIN (real lift, four candidate criteria) +
explicit STOP condition if MUST passes without WIN.
7. §9.1.1: alpha-realization metric defined explicitly as
realized_pnl / max(unrealized during signal-validity window).
Was hand-waved in v1; now formulated.
TRACTABLE INLINE FIXES
2. §7 open_trade_state: 128 → 64 bytes. Diagnostic-only fields
(scale_up/down counts, max/min target reached) moved to a
separate audit buffer (single source of truth keeps state lean).
4. §5.1 envelope: multiplicative formula → additive bounded
contributions. tanh(sharpe), min(dd/cap), clamp(vol). No single
factor can crash envelope to zero. Floor (envelope_floor lots)
and hard-cap (envelope_hard_cap) remain.
5. §5.2 threshold: hand-waved "self-tune toward PF" → explicit
3-arm bandit (lower, current, higher percentile) every K_eval
closed trades, Wiener-α EMA target with floor, bounded
[0.5, 0.95].
8. §3.4 + §3.6: Layer A wall-time ≤ 5× → ≤ 2× (controller is
light, 5× would indicate structural bug, not config tuning).
9. §4.2: explicit Wiener-α formula for conviction EMA with floor
at 0.4 per pearl_wiener_alpha_floor_for_nonstationary.
10. §6.2.1: shadow eval defined explicitly with backtest-mode (last
15% of fold) and live-mode (parallel-prediction-only) variants.
Promotion criteria spelled out (PF, Sharpe, tail loss).
11. §6.2: Layer D loss explicit (primary = value regression on
realized PnL conditional on state; secondary = EWC++ regularizer
toward base). Was three losses with no specified combination.
12. §10.5: adapter regime over-fit risk added with sample-count
down-weighting, cross-regime shadow eval, periodic adapter reset
as hard backstop.
GREENFIELDS (per user direction)
- §3.3: decision_stride field DELETED from YAML + Rust (not
deprecated). No backwards-compat shim. Per
feedback_no_legacy_aliases.
- §7.1: open_trade_state 24→64 byte migration is atomic, no old
layout shim.
- §8: boundary matrix decision_stride row marked REMOVED, not
DEPRECATED.
§13 pearl conformance: per-section pearl application (not bulk
citation). Added pearl_audit_unboundedness_for_implicit_asymmetry,
pearl_no_deferrals_for_complementary_fixes acknowledgement (layers
are sequenced amplifications, not parallel fixes), and the new
pearl_stop_checks_run_at_deadline_cadence from S2.
Status: Design v2 — awaiting user review before writing-plans.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
S2.2 moved max_hold to event-rate at step_resting_orders (cap is now
enforced — mean hold 432s → 58.43s, max 173893s → 96s). But it mirrored
the session-gap pattern `pos.position_lots = 0;` which intentionally
skips PnL ("cannot fill across halt"). Max_hold differs: the market is
live so the close SHOULD realize PnL via the current top-of-book.
Smoke t9msj (b92bd72c7) showed the consequence: 84% of trades record
$0 realised_pnl, win_rate dropped to 3.3%, total_pnl looks artificially
better (-$225k) only because losses aren't recorded.
Fix: route the force-close through apply_fill_to_pos by synthesizing a
closing fill at book.bid_px[0] (long) or book.ask_px[0] (short). The
existing counter-direction branch (realised = (avg_px − vwap) × dir ×
unwind) correctly realizes PnL on the unwound position. Top-of-book is
already validated by book_update_apply_snapshot's skip, so close_px is
guaranteed in range.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
S2.1 instrumentation (smoke cqpph @ 95a77c4ac) revealed the chain
worked end-to-end: max_hold check fired 1661 times, force-flat target
was written and seen by seed_inflight. Yet 86.5% of trades exceeded
the 60s cap with mean hold = 432s. Root cause: decision_policy's
stop_check_isv only runs every decision_stride events. At
decision_stride=200 on ES Q1 (2M events / ~91 days = ~3.9s sim-time
per event), decisions are ~13 min apart in sim-time, so the cap is
enforced 13 min late on average.
Fix: move the max_hold check to resting_orders_step (event-rate, runs
every iteration), same pattern as the session-gap force-close at
resting_orders.cu:282. Thread max_hold_ns_per_b and open_trade_state
into resting_orders_step. SL/trail stay in stop_check_isv because
they depend on ISV controller state at decision-rate.
Remove the dead decision-rate max_hold code and its 6 diagnostic
counter params from stop_check_isv per single-source-of-truth and
no-hiding rules. Keep mh_kernel_calls and mh_force_flat_seen_by_seed
counters as ongoing generic diagnostics. Update harness.rs log line
and stop_controller tests to exercise the event-rate path.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
S2 cluster smoke ppcfk (29f7e923c) showed mean trade hold = 432s with
max_hold_ns=60s — 886/1024 trades exceed the cap that should hard-flat
them. Upload chain and pnl_track entry_ts write look structurally
correct on inspection. Need per-step counters to localize WHICH step
of the chain fails.
Adds 7 counters in stop_check_isv (kernel_calls, seen_zero, entry_ts_zero,
ts_underflow, elapsed_below_cap, would_fire, force_flat_written) plus 1
counter in seed_inflight_limits_batched (seed_saw_force_flat). Logged at
smoke end as a third nan_counters-like line.
Diagnostic only — no behavioral change.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
S1.20-S1.22 chased a wrong hypothesis (walk_* deep-level filter) for
3 rounds. The actual bug: resting_orders.cu:344 fills at cost = take *
s.price in the queue-decay maturation arm. Market-like orders seed at
line 644 with sentinel s.price = 1e9 (buy) / 0 (sell), counting on the
marketability check at lines 372-... to fire first via walk_*. But
queue_position initializes to 0 (no book level matches the sentinel),
so on the first same-side trade volume the queue-decay arm fires and
injects the sentinel into cost: buy cost = take * 1e9 -> avg_px =
1e9 -> huge_flat=216, sell cost = take * 0 -> avg_px = 0 ->
zero_flat=194. Same pattern explains flip=9+6.
Fix: detect s.price outside [min_reasonable_px, max_reasonable_px] in
the queue-decay arm; absorb queue_position depletion without filling.
Marketability check fires in the same iteration via walk_* with proper
book prices.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
v2 NaN instrumentation (S1.21) localized the bug to apply_fill_to_pos's
open-from-flat branch writing pos.vwap_entry = avg_px where avg_px was 0
(194 cases) or > 21M finite (216 cases). The arithmetic was clean — root
cause is walk_ask_for_buy/walk_bid_for_sell consuming size from deep
levels (k=1..9) that have lvl_sz > 0 but lvl_px = 0 (or huge sentinel).
MBP-10 fills empty depth slots beyond available levels with these
sentinels.
Existing per-level filter checked lvl_sz <= 0, !isfinite(lvl_sz),
!isfinite(lvl_px) — but allowed lvl_px = 0 and lvl_px > max range.
Fix: thread per-backtest min_reasonable_px / max_reasonable_px (already
uploaded for the top-of-book skip in book_update_apply_snapshot via
S1.19) through to walk_*, and reject any level whose price falls
outside [min_px, max_px]. Same fix shape as Bug C-b (top-of-book skip),
now applied at all 10 depths.
Changes: resting_orders.cu (walk_* signatures + kernel param + 2 call
sites), order_match.cu (walk_* signatures + kernel param + 1 call site),
sim/mod.rs (submit_market + step_resting_orders launch args). 19 CUDA
tests pass.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
v1 instrumentation (S1.20) eliminated 3 of 6 hypotheses: apply_fill_to_pos
arithmetic is fully clean (avg_px=0 realised=0 realized_pnl=0). But
pnl_track open branch still saw vwap_entry=0 (194 times) or > 21M (216
times) at the open transition.
v2 adds per-vwap-write-site counters so we can pinpoint which apply_fill
branch produced the bad vwap, plus captures the actual last-bad-vwap
value and the path id (1..6 = open-flat, scale-in zero/huge, flip-beyond
zero/huge, session-gap-saw-stale).
The 7th counter (vwap_session_gap_was_bad) fires when the session-gap
force-close path sees pos.vwap_entry already in a corrupt state pre-gap.
Logged at end of smoke as `nan_counters_v2: zero_flat=N huge_flat=N
zero_scale=N huge_scale=N zero_flip=N huge_flip=N session_gap_was_bad=N
| b0_last_bad_vwap=X b0_last_bad_path=N`.
19/19 stop_controller CUDA tests pass; 5/5 decision_floor_coldstart pass.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Audit (fxt-data-audit on ES.FUT_2024-Q1) revealed real source-data
outliers that the hardcoded < 1e8 threshold didn't catch:
- bid_min = -$4.85 (negative bid)
- bid_p1 = $64.15 (1% of bids in sub-$100 range, far below ES)
- ask_max = $53,012 (10x above any plausible ES price)
The < 1e8 threshold = $100M was useless: never triggered on real ES.
Fix: parameterize the range. min_reasonable_px / max_reasonable_px as
per-backtest fields in UniformSimParams, ResolvedSimVariant, and
BatchedSimConfig (defaults 0.0 / f32::INFINITY = no-check, preserves
existing test fixtures). Sweep_smoke.yaml sets 1000/20000 for ES
futures — catches all observed outliers without rejecting any plausible
price. BacktestHarness calls upload_price_range() once after creating
the sim so the bounds are active before the first apply_snapshot.
CUDA kernel book_update_apply_snapshot gains two new args
(min_reasonable_px[n_backtests], max_reasonable_px[n_backtests]) that
replace the hardcoded > 0.0f && < 1.0e8f checks at both the top-of-book
gate and the per-level sanitization pass.
Test price_range_rejection_skips_snapshot validates: sub-$1000
snapshot, super-$20000 snapshot, and negative bid all skip with
snapshots_skipped counter increment; valid ES snapshot passes through.
17/17 stop_controller tests pass.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
One-shot diagnostic that loads a .dbn.zst via
load_or_predecode_mbp10 and reports symbol distribution, per-symbol
price stats (min/p1/p50/p99/max for bid_px[0]/ask_px[0]), outlier
counts (zero-price, huge-price >$100k, zero-price-with-size),
and first-record samples per symbol.
Built to investigate the 18% sentinel-laden trade residue in the
ISV stop-controller cluster smokes (97 zero, 85 i32::MAX, 14 weird-
other). The controller path is structurally correct; remaining bad
records hypothesized to originate from source MBP-10 data (mixed
symbols, corrupt records, or unreported instrument families).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Three orthogonal fixes for residual issues in smoke stx9p (97 zero
sentinels, 85 i32::MAX sentinels, 840/1024 over-60s holds):
Fix 1 — zero-sentinel residue (px > 0):
- Per-level sanitization in apply_snapshot_kernel only checked sz > 0.
- A book level with px=0, sz>0 passed → walk_* computed total_cost=0
→ avg_px=0 → vwap_entry=0. Trade record reported entry_px=0.
- Add px > 0.0f to bid_ok/ask_ok conditions.
Fix 2 — i32::MAX upper bound (px < 1e8):
- Post-Bug-D (1e9 nanoprice scaling) some boundary events carried prices
larger than any plausible instrument. Saturating cast to i32 produced
the 21474836 sentinel even when isfinite() passed.
- Add px < 1.0e8f upper bound to per-level AND top-of-book validation.
Fix 3 — session-gap force-close:
- max_hold check fires at decision_stride frequency, but during weekend
halts no events advance current_ts → max_hold never fires until next
session. Result: 49h holds in stx9p (840/1024 over 60s threshold).
- Detect ts gap > 1 hour in resting_orders_step. If position is open,
zero position_lots directly. pnl_track_step's existing close branch
emits the TradeRecord on the next call. No synthetic P&L added —
records show realised_pnl from whatever was accumulated before the gap
(honest: cannot fill across a halt).
- New per-backtest last_event_ts_d slot tracks the previous event ts.
- Test session_gap_force_closes_open_positions: 2-hour ts jump after
open verifies force-flat fires and exactly 1 TradeRecord is emitted.
All 16 stop_controller tests pass. All 5 decision_floor_coldstart pass.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Two root-cause fixes surfaced by cluster smoke v74v4 (sweep_smoke-a2dfc6d99):
Bug D — DBN parser price scaling:
- BidAskPair::price_to_f64/price_from_f64 used /1e12 / *1e12 from test-data
calibration. DBN production uses 1e-9 nanoprice (the DBN standard). ES at
5500 raw 5_500_000_000_000 → 5.5 instead of 5500. Smoke trade records
showed entry_px=5.24 instead of expected ~5240 ES index points (1000×
too small). Fix: 1e12 → 1e9 in both functions. Round-trip symmetric;
tests updated.
Bug C-b — corrupt top-of-book sentinel:
- Per-level sanitization (Task 15) zeros each unhealthy MBP-10 level
individually. At session-boundary events with all 10 levels invalid,
the book becomes uniformly zero. apply_fill_to_pos then reads
bid_px[0]=0 / ask_px[0]=0 → vwap_entry=0 → trade record entry_px=0
(zero sentinel in v74v4 CSV, 162/1024 trades in n59t4).
- Fix: pre-validate top-of-book in apply_snapshot_kernel. If
bid_px[0]/ask_px[0]/bid_sz[0]/ask_sz[0] are non-finite or
bid_px[0]<=0/ask_px[0]<=0/bid_sz[0]<=0/ask_sz[0]<=0, atomically skip
the entire snapshot (book/prev_mid/atr_mid_ema unchanged). Add
per-backtest snapshots_skipped_d counter for observability.
Test corrupt_top_of_book_skips_snapshot_and_increments_counter validates
NaN top-price + zero top-size cases both increment the counter without
mutating state, and that a subsequent valid snapshot updates normally.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Smoke n59t4 (commit 691dec144, root-cause NaN/Inf fix) revealed 96
trades/sec at decision_stride=4 — decisions fire every ~4ms while
orders take 200ms to fill. Controller makes decisions on phantom
in-flight positions; stops fire against positions that haven't
materialized yet. At $0.25 round-trip cost × 96 trades/sec, fee burn
is $86k/hour.
Architectural mismatch: decision_stride should be >= latency_ns /
event_dt_ns. At 200ms latency and ~1ms/event, stride=200 means
decisions fire every ~200ms — once per latency cycle. Resulting
~10k decisions over 2M events is the natural cadence for the
deployment anchor.
Not a controller bug. The controller logic (NaN-clean per Task 15)
is correct; the smoke config was tuned for sub-millisecond latency
that doesn't match the 200ms Scaleway PAR → IBKR anchor.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Two orthogonal bugs that compounded to produce the exit_px=±21474836
sentinel in cluster smokes (315 trades baseline, 500 trades pearl):
Bug A — NaN/Inf book propagation:
- apply_snapshot_kernel passes through NaN/Inf prices from MBP-10
predecoded data (session boundaries, gap-fills).
- walk_ask_for_buy / walk_bid_for_sell accumulate cost += take * NaN.
- apply_fill_to_pos: avg_px = NaN; realized_pnl += NaN → permanent NaN.
- Subsequent close: (int)(NaN-derived * 5000) → i32::MAX sentinel.
- Fix: zero-out non-finite levels in apply_snapshot_kernel; add
isfinite() guards in walk helpers as defense-in-depth; ATR update
guards against non-finite mid.
Bug B — pnl_track close branch doesn't reset scratch:
- Scratch reset (full 24-byte memset to 0) was already present in the
current code; no source change required for Bug B.
Tests:
- book_nan_inf_prices_dont_corrupt_realized_pnl: NaN at ask[3] + Inf
at bid[5] survives the fill pipeline without making realized_pnl
non-finite.
- pnl_track_resets_scratch_on_close: open+close+5 flat events emits
exactly 1 record; second open+close emits exactly 1 more (=2 total).
14/14 stop_controller tests pass; all other ml-backtesting test suites
unaffected.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Three follow-ups from cluster smoke gp74n (trade_vol pearl validation):
1. Max-hold force-close: max_hold_ns added as per-backtest config
(default 0 = disabled). Fires force-flat (3, 0) when current_ts -
entry_ts >= max_hold_ns, BEFORE SL/trail check. Tested via
max_hold_forces_close. Sweep YAML sets 60s cap to bound the long
tail observed in gp74n (263985s pathological hold).
2. exit_px defensive sanity check: 500/1024 gp74n trades reported
exit_px = ±i32::MAX/100 (float→int saturation sentinel from likely
NaN segment_realized). Defensive fix in pnl_track_step: if exit_px
is non-finite or diverges from entry_px by >10%, fall back to
entry_px with zero realised_pnl. Root cause to be traced separately.
3. Spec §9.2 criteria revision: mean_hold<30s replaced with p95<600s
+ max<=max_hold_ns. The 30s threshold reflected the pre-pearl
sub-cost churning bug, not a real feature criterion. p95 catches
long-tail pathology while letting alpha-driven exit timing breathe.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Replaces the Task 12 `2.0f * cost` floor (hardcoded multiplier) with
trade_vol = sqrt(realised_return_var) bootstrapped from cost². Per
pearl_trade_level_vol_for_stop_distance.md: microstructure ATR is the
wrong time scale for trade-level stop decisions; per-horizon
realised_return_var is the right one, with cost² as a structural cold-
start sentinel.
cost now appears exactly once — inside the sqrt as a bootstrap sentinel,
never as a distance multiplier. The 2.0f literal is eliminated;
controller is fully ISV-driven.
var_avg accumulates realised_return_var in the same single-pass horizon
loop as ema_loss/ema_win. Cold-start (var_avg=0): trade_vol = cost.
Post-bootstrap: sqrt(var_avg) dominates.
Test retargeted: cost_floor_prevents_sub_cost_stops →
trade_vol_floor_prevents_sub_cost_stops, with boundaries straddling
trade_vol=cost=0.125 instead of the prior 2*cost=0.25 (no-fire Δ=0.08,
fire Δ=0.20).
Spec §5, §10, §11, §12 amended.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Cluster smoke w7b4p showed 44864 closed trades in 1M events at 200ms
latency — physically impossible without sub-cost churning. Root cause:
sl_distance = max(pnl_ema_loss, atr) ≈ 0.05 at cold-start, but round-trip
cost = 2 × 0.125 = 0.25. Every stop-out guaranteed loss > 5× distance;
EMA converges sub-cost.
Amendment: triple-max sl_distance = max(pnl_ema_loss, atr, 2*cost);
trail_distance = max(pnl_ema_win, atr, 2*cost). cost is an ISV-discipline
strategy anchor (already per-backtest from P4 — no new state slot, no
tuned constants). Added cost_per_lot_per_side_per_b param to both decision
kernel signatures and threaded cost_per_lot_per_side_d through both launch
sites in step_decision_with_latency.
Test: cost_floor_prevents_sub_cost_stops validates unrealized in (ATR, 0.25)
does NOT fire SL; unrealized > 0.25 DOES fire. Spec amended (§10, §12).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
- StopRules struct + sl_tp_rules field + all literals deleted; the ISV
stop controller (Tasks 2-9) replaces this dead data path.
- use_cold_start_stopgap propagation deleted across harness,
batched_config, fxt-backtest, sweep_smoke.yaml, and 5 test files.
- Q1 stopgap branch in harness.rs deleted; replaced with the original
simple strategy-upload loop.
- decision_floor_coldstart test retargeted: cold_start_persistent_
bullish_with_default_stops_never_closes -> ..._now_closes, asserts
trades > 0 AND |pos| <= max_lots (99 trades, pos=1 on local GPU).
- CBSW spec + plan prefixed with SUPERSEDED headers.
Per feedback_single_source_of_truth_no_duplicates +
feedback_no_partial_refactor: contract change migrates every consumer
in one commit. No legacy wrappers; no version suffixes.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Two-instruction addition to the existing close-emission branch
(prev != 0 && now == 0). Without the reset, the next entry inherits
a stale HWM that could arm the trail spuriously on event 0 of the
new trade.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
market_targets[b] now interpreted as target absolute position (side=0
long, side=1 short, side=2 no-op preserved, side=3 force-flat).
seed_inflight_limits_batched computes order_lots = target_signed -
effective_position where effective_position sums pos.position_lots
plus all unfilled signed slot sizes (active in {1, 2}). Fixes both
the additive accumulation bug AND the worse latency-bypass variant
(naive delta would have made things 200x worse under 200ms latency).
3 new tests added (all passing):
- position_target_not_additive: latency=0, 10 repeated target events stay <= max_lots=5
- position_target_not_additive_with_latency: 200ms latency, 500 events, in-flight summation prevents runaway
- alpha_noop_side_2_preserved: side=2 events leave filled position unchanged
All 9 stop_controller tests pass; parallel_sim and fuzz regression clean.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Both decision kernels now call the same __device__ helper at the top
of per-backtest dispatch (after the program_lens early-return). Single
source of truth for stop logic across default and bytecode-VM paths.
Parity test ensures they produce identical market_target output for
matched scenarios.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Boundary test that fails on bit-pick-first or bit-pick-max regressions
of the open_horizon_masks averaging in stop_check_isv. Per-horizon
pnl_ema_loss values (2, 4, 3, 3, 3) yield mean=3.0 across the 0x1F
mask; snapshots placed relative to vwap_entry for ask-spread safety.
Tests Δ_entry=2.5 (no-fire) and Δ_entry=4.5 (fire).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
trail_distance = max(pnl_ema_win, atr_mid_ema). HWM ratchets up via
fmaxf each event while position open; trail fires when HWM clears
the arming threshold AND unrealized drops by trail_distance from
HWM. Same force-flat (3, 0) write as SL.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Skeleton helper called from decision_policy_default at top of
per-backtest dispatch. Returns 0 (no-op) on flat positions; trigger
logic for open positions added in Tasks 4-6. Single source of truth
for stop logic across decision_policy_{default,program}.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
cudarc kernel launches require `unsafe` blocks — the driver API has no
way to type-check kernel args against the cubin signature. The
workspace-wide `-W unsafe-code` lint produces noise on every launch
site (10+ blocks); suppressing at the module level keeps the lint
useful elsewhere in the crate.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per-event Wiener-α=0.4 EMA on |Δmid| with first-observation bootstrap.
Floor source for the SL/trail-distance controller (spec §6). Thread 0
of the broadcast-snapshot kernel handles the update per backtest;
threads 1..9 still handle the 10 level writes.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Three CudaSlice<f32> per-backtest slots (prev_mid_d, atr_mid_ema_d,
trail_hwm_d) plus test-only accessors. Foundation for the
ISV-driven stop controller; no behavioral change until subsequent
tasks wire them into kernels.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
11 tasks, each one commit, TDD-disciplined (test → fail → impl →
pass → commit). Covers all 11 sections of the spec at
docs/superpowers/specs/2026-05-19-isv-driven-stop-controller-design.md:
- Tasks 1-2: state slots + ATR EMA in book_update_apply_snapshot
- Task 3: stop_check_isv shared __device__ helper skeleton
- Tasks 4-6: SL trigger, trail-TP + HWM ratchet, multi-horizon avg
- Task 7: bytecode VM parity (wire helper into decision_policy_program)
- Task 8: seed_inflight_limits_batched target-delta + in-flight summation
- Task 9: pnl_track_step trail_hwm reset on close
- Task 10: atomic deletion of StopRules + use_cold_start_stopgap +
integration-test retarget + CBSW SUPERSEDED markers
- Task 11: cluster smoke validation via argo-lob-sweep.sh
All 10 §9.1 unit tests have explicit failing-test scaffolds before
their implementations, and the retargeted integration test
(cold_start_persistent_bullish_now_closes) flips the assertion to
prove the fix.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Issues caught in self-review and fixed:
1. Encoding ambiguity: stop-fire writing (2, 0) would have silently
collapsed weak-alpha no-op semantics with force-flat. Introduced
distinct side=3 = force-flat; preserved side=2 = no-op for entry
path. Updated §3 + §5 + §7 truth table.
2. trail_hwm reset on close was implied in §4 but missing from §8's
touched-files list. Added pnl_track.cu modification (§7.1) +
added to §8 Added list.
3. __device__ helper enforcement: §3 now mandates stop_check_isv()
as a shared __device__ function called from both decision kernels,
not duplicated code.
4. Kernel arg additions enumerated explicitly in new §4.1 table —
every kernel-launch site gets named, no "thread through every kernel"
hand-wave.
5. seed_inflight_limits_batched (resting_orders.cu:439–475) named
explicitly in §1 + §7 — the actual kernel that needs the position-
target-semantics fix.
6. In-flight order summation: naive delta = target_signed - pos was
strictly worse than the original additive bug under non-zero
latency (would produce pos = pre + K·delta after fills). Fixed
§7 algorithm to compute effective_position by scanning all
active∈{1, 2} slots and summing their signed sizes. Added
position_target_not_additive_with_latency test.
7. ATR α=0.4 honest justification in §10: not strictly derived from
pearl_wiener_alpha_floor_for_nonstationary (which is about co-
adapting control loops, not passive observation); chosen as a
pragmatic project-wide constant matching isv_kelly_update_on_close.
8. CUDA Graph host-branch-free affirmation added to §3 — explicit
compatibility with P5 graph capture.
9. Cold-start symmetry caveat (§5): at n_trades_seen=0 both EMAs are
0 so sl_distance == trail_distance == atr — asymmetry only
emerges post-bootstrap. Acknowledged honestly.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Replaces StopRules::default()-all-zeros (the actual cause of the cluster
smoke's n_trades=0, falsifying the CBSW dilution diagnosis) with an
ISV-driven SL + trail-stop controller folded into the existing decision
kernel(s). Per-backtest ATR EMA on mid-price + per-horizon pnl_ema_*
drive distances under max(real, floor) per
pearl_blend_formulas_must_have_permanent_floor. No new kernel; single
source of truth in step_decision*. Folds in the position-target
semantics fix (200-event local repro showed position_lots=199 from
additive-vs-target order semantics) — same commit.
Removes: StopRules, sl_tp_rules, Q1 stopgap field/branch,
use_cold_start_stopgap propagation. Marks CBSW spec + plan SUPERSEDED.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Adds LobSimCuda::read_total_trade_count (cheap n_backtests*4 byte DtoH of
the trade-log head counters) and emits the sum on each PROGRESS_EVERY
eprintln. Lets future smokes (Q2 CBSW validation, P7 sweep) see trade
firing mid-run instead of waiting for write_artifacts at harness end.
Heads count kernel writes (clamped to ring cap), so the sum is a
monotone trade counter across the sweep — drops back only on harness
reset which doesn't happen mid-run.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The threshold-tuning smoke at 81decf40f produced n_trades=0 despite
74.6% of decisions having max_conv ≥ 0.30 — the linear-weighted-mean
aggregator in decision_policy_default is structurally dilution-bound
at cold-start (per spec §1).
Q1 stopgap: when sim_variants[i].use_cold_start_stopgap = true, the
harness uploads a max-confidence Strategy bytecode program for that
backtest, routing decisions through decision_policy_program with
OP_AGG_MAX_CONFIDENCE. Existing kernel; zero CUDA changes.
Field additions (atomically across BatchedSimConfig + UniformSimParams
+ ResolvedSimVariant + SweepBase.SimVariant) — every UniformSimParams
literal migrated to include use_cold_start_stopgap: false (default).
The sweep YAML's sim_variants entry sets it to true only for the
validation run; production deployability uses Q2's kernel fix instead.
Sweep YAML (config/ml/sweep_smoke.yaml) flipped to use_cold_start_stopgap=true
at threshold=0.0, cost=0.125 — same anchor as the threshold-tuning
smoke that produced n_trades=0, for direct comparison.
This is a VALIDATION step. Cluster smoke at this commit MUST produce
n_trades > 100 + finite metrics. Q2's kernel CBSW immediately follows
and deletes this entire stopgap atomically (field, harness branch,
YAML setting, every literal).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
4-task atomic ladder (Q1-Q4) implementing spec 566e8bcb0. Each task
maps to one commit per feedback_no_partial_refactor:
Q1: Cold-start stopgap (bytecode max-confidence policy upload)
Q2: CBSW kernel aggregator + Tier 1 revert (atomic)
Q3: Memory pearl pearl_conviction_bootstrap_for_kelly_aggregation
Q4: Parallelism spec §3.4 cross-reference
Kernel ABI unchanged across all 4 commits — Q2 modifies only
decision_policy_default's body (decision_policy_program left as a
pluggable bytecode-VM experiment surface). Q2 atomically deletes
the Q1 stopgap field/CLI/YAML/harness branch in the same commit
that lands the kernel fix (feedback_no_legacy_aliases compliance:
grep returns zero hits for use_cold_start_stopgap post-Q2).
Validation gates per spec §9:
- Q1: cluster smoke n_trades > 100, total_pnl != 0 (diagnosis
confirmation). If n_trades = 0 still, STOP — downstream bug.
- Q2: pre-existing 11 CUDA tests still pass; 7 new cbsw_*
regression tests pass; cluster smoke n_trades > 100; ev/s
rate ratio Q2/Q1 ≥ 0.95.
Self-review confirms all 11 spec sections covered, no placeholders,
type/signature consistency across tasks.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Critical review surfaced 11 actionable issues in the v1 spec; 10 fixed
inline (#11 — legacy-comparison test — dropped per user decision since
empirical legacy behavior is already known: n_trades=0).
Performance:
1. expf → piecewise-linear ramp (~5× cheaper on GPU, no transcendental
in the hot path). Operationally equivalent: monotonic, saturating,
midpoint at K. Side-benefit: pure max-confidence at cold-start
(sq=0 instead of sigmoid's 11.9% leakage).
2. Single fused per-horizon pass replaces the two-loop sketch.
3. Tier 2 scope narrowed to decision_policy_default ONLY. The bytecode
VM (decision_policy_program) stays unchanged — runs only for custom
strategy experiments, never in production policy. Halves Tier 2's
surface area + test cost.
4. __device__ helpers cbsw_signal_quality + cbsw_weight introduced for
single source of truth.
Correctness bugs (silently present in v1 sketch, would have shipped):
5. strong_h and sq_min_active initializers added (were referenced
before init in the per-horizon loop).
6. Attribution mask at cold-start was setting all 5 bits because every
w[h] >= floor > 1e-9; trades would pollute ALL horizons'
recent_sharpe. Fix: binary split — at sq_min_active < 0.5 attribute
only to strong_h; at mature, attribute per weights > floor + epsilon.
7. sq_min was "min over h", which permanently locked the aggregator
into cold-start mode if any horizon never gets attributed (e.g.,
h6000-only-trading regime). Fix: sq_min_active = min over h with
n_trades_seen > 0; cold horizons don't gate maturity.
8. Opposing-horizons case now correctly fires on the strongest single
horizon at cold-start (piecewise-linear sq=0 → pure max-conf), with
explicit design-choice note explaining why this conservatism trade-
off favors firing.
Spec hygiene:
9. Rate validation committed to a measurement gate (Q2 ev/s ≥ 95% of
Q1 ev/s) instead of back-of-envelope estimate.
10. Kernel ABI explicitly stated as unchanged → feedback_no_partial_
refactor compliance is trivial at the kernel boundary.
12. Tier 1's use_cold_start_stopgap field is DROPPED in Q2 atomically
(not orphaned), and DoD checklist includes a grep-zero gate for
feedback_no_legacy_aliases compliance.
Risks updated: removed the sigmoid-narrowing risk (irrelevant now);
added register-pressure risk with the rate-gate mitigation; added
explicit "cold-start may lose on noisy trades" risk with empirical
escalation path (bump kelly_floor 0.20 → 0.40 if Q2 smoke shows large
negative PnL).
Math walk-throughs rewritten for piecewise-linear (different numbers
at cold-start): cold = pure max-conf, mature = pure weighted-sharpe,
clean separation at sq_min_active threshold.
Awaiting review of the revised spec before writing-plans dispatch.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The post-trunk-grows threshold-tuning smoke (81decf40f) produced
n_trades=0 despite the model having a HEALTHY max-conviction
distribution (74.6% of decisions ≥ 0.30, 26.7% ≥ 0.70, full spread
across [0,1]). Diagnosis: the linear-weighted-mean aggregator in
decision_policy_default is structurally dilution-bound at cold-start
— single-horizon strong signals get washed out when uniform-floor
weights produce mean-over-horizons aggregation.
Solution: Conviction-Bootstrapped Sharpe Weighting (CBSW). Hybrid
max-confidence × weighted-sharpe with a per-horizon sigmoid transition
keyed on n_trades_seen vs MIN_TRADES_FOR_VAR_CAP. Cold-start: sig_mag
(decision-time conviction) drives weights AND max-confidence aggregator
fires single-horizon trades. Mature: recent_sharpe (historical) drives
weights AND linear-mean aggregator emphasizes strong-Sharpe horizons.
Permanent floor preserved per pearl_blend_formulas_must_have_permanent_floor.
Mirrors DQN bootstrap pattern (pearl_thompson_for_distributional_action_
selection): when historical estimates are uncertain, use available
signal as the bootstrap. Sig_mag is the ISV signal at decision time;
recent_sharpe is the ISV signal at trade-close time. Transition
self-terminates based on data accumulation, not time constants.
3-tier delivery (one spec, atomic commits):
Q1: Bytecode VM stopgap — upload max-confidence 7-instruction
program per backtest. Validates diagnosis; zero kernel work.
Q2: Kernel CBSW — replace weight + aggregator in both decision
kernels. 5 new regression tests covering cold/mature/transition.
Q3: New memory pearl pearl_conviction_bootstrap_for_kelly_aggregation
capturing the lesson.
Q4: Cross-reference from parallelism spec (deployability sweep
depends on CBSW being live to produce meaningful verdict).
The parallelism work (P1-P6) is fully working — confirmed by the
threshold-tuning smoke completing end-to-end (Succeeded status, 500k
decisions, artifacts written, aggregator parquet emitted). What's
blocked is the deployability VERDICT, because the dilution bug means
all variants would show n_trades=0 regardless of cost/latency/threshold.
CBSW unblocks the verdict.
Awaiting review before transitioning to writing-plans for Q1-Q4
implementation plan.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
P4 plumbed the threshold-gate kernel side but deferred the side-channel
that captures observed max_conviction per decision. Wire it now so the
threshold pre-registration step (spec §3.4) can compute the calibrated
p60-p95 absolute threshold values from a real model-on-data run.
Harness changes:
- BacktestHarness gains conviction_log: Vec<f32>. Per decision, computes
max_h |alpha[h] - 0.5| * 2 from the SAME probs that go into broadcast_
alpha (same value the threshold gate would compare against), pushes
to the log. One shared vec — batched cells broadcast the same probs
to every backtest, so per-backtest is redundant.
- write_artifacts emits convictions.bin (raw little-endian f32) +
conviction_percentiles.json with pre-computed p10/p25/p50/p60/p70/
p80/p90/p95/p99 + mean/min/max. Also eprintln-prints the summary
line for at-a-glance log inspection.
Smoke YAML switched to the threshold-tuning configuration: threshold=0
(no gate, full distribution captured), cost=0.125 (1-tick realistic
anchor so the observed Sharpe is the no-gate net-of-cost floor for
the sweep's deployability story).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
