diff --git a/docs/superpowers/plans/2026-05-19-isv-driven-stop-controller.md b/docs/superpowers/plans/2026-05-19-isv-driven-stop-controller.md new file mode 100644 index 000000000..05fbf200e --- /dev/null +++ b/docs/superpowers/plans/2026-05-19-isv-driven-stop-controller.md @@ -0,0 +1,1735 @@ +# ISV-Driven Stop Controller Implementation Plan + +> **For agentic workers:** REQUIRED SUB-SKILL: Use superpowers:subagent-driven-development (recommended) or superpowers:executing-plans to implement this plan task-by-task. Steps use checkbox (`- [ ]`) syntax for tracking. + +**Goal:** Implement the ISV-driven stop controller specified in `docs/superpowers/specs/2026-05-19-isv-driven-stop-controller-design.md` so the cluster smoke produces `n_trades ≥ 100`, `mean_trade_hold_time_ns ∈ (0, 30s)`, and `|pos.position_lots| ≤ max_lots`. + +**Architecture:** Stop check lives in a shared `__device__` helper (`stop_check_isv`) called at the top of both `decision_policy_default` and `decision_policy_program`. Distances derive from per-horizon `pnl_ema_win / pnl_ema_loss` with a per-backtest ATR-on-mid EMA floor (`max(real, floor)` per `pearl_blend_formulas_must_have_permanent_floor.md`). The `seed_inflight_limits_batched` kernel converts market_targets to delta orders against `effective_position = pos.position_lots + Σ unfilled slot signed-sizes`, ensuring `|pos.position_lots| ≤ max_lots` invariant. Stop-fire uses a new `side=3` encoding distinct from `side=2` (alpha no-op). + +**Tech Stack:** Rust 1.85+, Tokio 1.40, cudarc 0.19, CUDA 12.4 (sm_86 local / sm_89 cluster). + +**Branch:** `ml-alpha-phase-a` (current working branch). + +--- + +## File Map + +### New device code +- `crates/ml-backtesting/cuda/decision_policy.cu` — append `__device__ stop_check_isv()` helper; add helper calls in `decision_policy_default` and `decision_policy_program`; add `(3, 0)` write path; extend kernel signatures. + +### Modified device code +- `crates/ml-backtesting/cuda/book_update.cu` — extend `book_update_apply_snapshot` with ATR-EMA update block. +- `crates/ml-backtesting/cuda/pnl_track.cu` — `pnl_track_step` adds `trail_hwm[b] = 0` in close branch. +- `crates/ml-backtesting/cuda/resting_orders.cu` — `seed_inflight_limits_batched` reworked for target-delta + in-flight summation + `side=3` handling. + +### Modified host code +- `crates/ml-backtesting/src/sim/mod.rs` — three new `CudaSlice` fields, allocation, kernel-arg threading, three new accessor methods. +- `crates/ml-backtesting/src/policy/mod.rs` — drop `StopRules` re-export, drop `sl_tp_rules` field, update `StrategyConfig` literals. +- `crates/ml-backtesting/src/harness.rs` — delete Q1 stopgap branch + `use_cold_start_stopgap` propagation. +- `crates/ml-backtesting/src/sim/batched_config.rs` — drop `use_cold_start_stopgap` field from three structs. +- `bin/fxt-backtest/src/main.rs` — drop `use_cold_start_stopgap` from `SweepBase.SimVariant`. +- `config/ml/sweep_smoke.yaml` — rename variant `t0c1l200_stopgap` → `t0c1l200_isv_stops`, drop the flag. + +### Deleted files +- `crates/ml-backtesting/src/policy/stop_rules.rs` — type no longer used. + +### New tests +- `crates/ml-backtesting/tests/stop_controller.rs` — 10 of the 11 spec §9.1 tests. + +### Modified tests +- `crates/ml-backtesting/tests/decision_floor_coldstart.rs` — retarget `cold_start_persistent_bullish_with_default_stops_never_closes` to `cold_start_persistent_bullish_now_closes` with inverted assertion. +- `crates/ml-backtesting/tests/{threshold_and_cost, lob_sim_fixtures, lob_sim_integrated_fuzz, parallel_sim_correctness}.rs` — drop `use_cold_start_stopgap: false` from `UniformSimParams` literals. + +### Modified docs +- `docs/superpowers/specs/2026-05-19-cbsw-cold-start-aggregator-design.md` — prepend SUPERSEDED header. +- `docs/superpowers/plans/2026-05-19-cbsw-cold-start-aggregator.md` — prepend SUPERSEDED header. + +--- + +## Task 1: Sim state slots + accessors (foundation) + +**Files:** +- Modify: `crates/ml-backtesting/src/sim/mod.rs` (field declarations near line 85, allocation near line 222, accessor methods after `n_backtests()` at line 389) + +- [ ] **Step 1: Add the three slot fields to `LobSimCuda` struct** + +Find the struct field declarations (around line 60-110 in `sim/mod.rs`) and append: + +```rust + /// Per-event ATR EMA on mid-price; permanent floor source for §5 controller. + prev_mid_d: CudaSlice, // [n_backtests] + atr_mid_ema_d: CudaSlice, // [n_backtests] + /// Per-position trail-stop high-water-mark of unrealized-P&L-per-lot. + /// Reset to 0 by `pnl_track_step` on close transition. + trail_hwm_d: CudaSlice, // [n_backtests] +``` + +- [ ] **Step 2: Allocate the three slots in `LobSimCuda::new`** + +In the allocation block around line 222 (after `program_lens_d`), append: + +```rust + let prev_mid_d = stream + .alloc_zeros::(n_backtests) + .context("alloc prev_mid_d")?; + let atr_mid_ema_d = stream + .alloc_zeros::(n_backtests) + .context("alloc atr_mid_ema_d")?; + let trail_hwm_d = stream + .alloc_zeros::(n_backtests) + .context("alloc trail_hwm_d")?; +``` + +Add them to the `Ok(Self { ... })` struct literal near line 270: + +```rust + prev_mid_d, + atr_mid_ema_d, + trail_hwm_d, +``` + +- [ ] **Step 3: Add three accessor methods after `n_backtests()`** + +After line 391 (end of `pub fn n_backtests`), append: + +```rust + /// Read `atr_mid_ema_d[backtest_idx]`. Test-only accessor. + pub fn read_atr_mid_ema(&self, backtest_idx: usize) -> Result { + anyhow::ensure!(backtest_idx < self.n_backtests, "atr idx oob"); + let mut buf = vec![0.0f32; self.n_backtests]; + self.stream.memcpy_dtoh(&self.atr_mid_ema_d, buf.as_mut_slice())?; + Ok(buf[backtest_idx]) + } + + /// Read `prev_mid_d[backtest_idx]`. Test-only accessor. + pub fn read_prev_mid(&self, backtest_idx: usize) -> Result { + anyhow::ensure!(backtest_idx < self.n_backtests, "prev_mid idx oob"); + let mut buf = vec![0.0f32; self.n_backtests]; + self.stream.memcpy_dtoh(&self.prev_mid_d, buf.as_mut_slice())?; + Ok(buf[backtest_idx]) + } + + /// Read `trail_hwm_d[backtest_idx]`. Test-only accessor. + pub fn read_trail_hwm(&self, backtest_idx: usize) -> Result { + anyhow::ensure!(backtest_idx < self.n_backtests, "trail_hwm idx oob"); + let mut buf = vec![0.0f32; self.n_backtests]; + self.stream.memcpy_dtoh(&self.trail_hwm_d, buf.as_mut_slice())?; + Ok(buf[backtest_idx]) + } +``` + +- [ ] **Step 4: Verify compilation** + +```bash +SQLX_OFFLINE=true cargo check -p ml-backtesting +``` + +Expected: compiles cleanly (the new fields are unused for now — Rust does not warn on private struct fields). + +- [ ] **Step 5: Commit** + +```bash +git add crates/ml-backtesting/src/sim/mod.rs +git commit -m "$(cat <<'EOF' +feat(ml-backtesting): add stop-controller state slots + accessors + +Three CudaSlice 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) +EOF +)" +``` + +--- + +## Task 2: ATR EMA update in `book_update_apply_snapshot` + +**Files:** +- Test: `crates/ml-backtesting/tests/stop_controller.rs` (NEW) +- Modify: `crates/ml-backtesting/cuda/book_update.cu` +- Modify: `crates/ml-backtesting/src/sim/mod.rs` (kernel-launch site for `book_update_apply_snapshot`) + +- [ ] **Step 1: Create the new test file with the ATR-bootstrap test** + +Create `crates/ml-backtesting/tests/stop_controller.rs`: + +```rust +//! Stop-controller spec §9.1 unit tests. +//! See docs/superpowers/specs/2026-05-19-isv-driven-stop-controller-design.md. + +use anyhow::Result; +use ml_backtesting::sim::LobSimCuda; +use ml_core::device::MlDevice; + +fn level_book(mid: f32, tick: f32) -> ([f32; 10], [f32; 10], [f32; 10], [f32; 10]) { + let mut bid_px = [0.0; 10]; + let mut bid_sz = [0.0; 10]; + let mut ask_px = [0.0; 10]; + let mut ask_sz = [0.0; 10]; + for k in 0..10 { + bid_px[k] = mid - tick * (k as f32 + 1.0); + ask_px[k] = mid + tick * (k as f32 + 1.0); + bid_sz[k] = 20.0; + ask_sz[k] = 20.0; + } + (bid_px, bid_sz, ask_px, ask_sz) +} + +#[test] +#[ignore = "requires CUDA"] +fn atr_ema_bootstrap_first_observation() -> Result<()> { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { eprintln!("skipping: cuda unavailable ({e})"); return Ok(()); } + }; + let mut sim = LobSimCuda::new(1, &dev)?; + + // Event 0: snapshot at mid=5500. After: prev_mid=5500, atr_mid_ema=0. + let (bp, bs, ap, az) = level_book(5500.0, 0.25); + sim.apply_snapshot(&bp, &bs, &ap, &az)?; + assert_eq!(sim.read_atr_mid_ema(0)?, 0.0, "ATR is 0 after first snapshot"); + let prev0 = sim.read_prev_mid(0)?; + assert!((prev0 - 5500.0).abs() < 1e-4, "prev_mid bootstrapped to 5500, got {prev0}"); + + // Event 1: snapshot at mid=5500.5. Δ=0.5. After: atr_mid_ema=0.5 (replace, not blend). + let (bp, bs, ap, az) = level_book(5500.5, 0.25); + sim.apply_snapshot(&bp, &bs, &ap, &az)?; + let atr1 = sim.read_atr_mid_ema(0)?; + assert!((atr1 - 0.5).abs() < 1e-4, "ATR after 2nd snapshot equals |Δmid|=0.5, got {atr1}"); + + // Events 2..200: jitter book. ATR stays positive and finite, bounded by max Δ. + let mut mid = 5500.5_f32; + let mut max_delta: f32 = 0.5; + for i in 0..198 { + let drift = if i % 3 == 0 { 0.25 } else if i % 3 == 1 { -0.25 } else { 0.0 }; + mid += drift; + max_delta = max_delta.max(drift.abs()); + let (bp, bs, ap, az) = level_book(mid, 0.25); + sim.apply_snapshot(&bp, &bs, &ap, &az)?; + } + let atr_final = sim.read_atr_mid_ema(0)?; + assert!(atr_final > 0.0 && atr_final.is_finite(), + "ATR positive + finite after 200 snapshots, got {atr_final}"); + assert!(atr_final <= max_delta + 1e-3, + "ATR bounded by max observed Δ={max_delta}, got {atr_final}"); + Ok(()) +} +``` + +- [ ] **Step 2: Run test to verify it fails (functional + API)** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test stop_controller -- atr_ema_bootstrap_first_observation --ignored --nocapture +``` + +Expected: test runs but **fails on `assert!((atr1 - 0.5).abs() < 1e-4`** — ATR remains 0 on event 1 because `book_update_apply_snapshot` doesn't update it yet. (read_atr_mid_ema compiles and returns 0.0 from the alloc_zeros slot.) + +- [ ] **Step 3: Extend the CUDA kernel with the ATR update block** + +In `crates/ml-backtesting/cuda/book_update.cu`, replace the entire kernel with: + +```c +extern "C" __global__ void book_update_apply_snapshot( + const float* __restrict__ bid_px_in, // [10] — broadcast + const float* __restrict__ bid_sz_in, // [10] + const float* __restrict__ ask_px_in, // [10] + const float* __restrict__ ask_sz_in, // [10] + Book* __restrict__ books_out, // [n_backtests] + float* __restrict__ prev_mid, // [n_backtests] + float* __restrict__ atr_mid_ema, // [n_backtests] + int n_backtests +) { + int b = blockIdx.x; + int tid = threadIdx.x; + if (b >= n_backtests) return; + + if (tid < 10) { + Book& bk = books_out[b]; + bk.bid_px[tid] = bid_px_in[tid]; + bk.bid_sz[tid] = bid_sz_in[tid]; + bk.ask_px[tid] = ask_px_in[tid]; + bk.ask_sz[tid] = ask_sz_in[tid]; + } + + // Thread 0 handles per-backtest ATR-EMA update against the broadcast mid. + if (tid == 0) { + const float mid = 0.5f * (bid_px_in[0] + ask_px_in[0]); + const float prev = prev_mid[b]; + if (prev == 0.0f) { + // First-observation bootstrap (pearl_first_observation_bootstrap). + prev_mid[b] = mid; + // atr_mid_ema stays 0 until event 2. + } else { + const float delta = fabsf(mid - prev); + const float ema = atr_mid_ema[b]; + if (ema == 0.0f) { + // First non-zero delta: replace directly. + atr_mid_ema[b] = delta; + } else { + // Wiener-α=0.4 EMA (pearl_wiener_alpha_floor_for_nonstationary). + atr_mid_ema[b] = 0.4f * delta + 0.6f * ema; + } + prev_mid[b] = mid; + } + } +} +``` + +- [ ] **Step 4: Thread the new args through the kernel launch in `sim/mod.rs`** + +Find the `apply_snapshot` method (around line 407 in `sim/mod.rs`). Locate the kernel launch (look for `book_update_apply_snapshot` function load + `LaunchConfig` usage). Append the two new args to the `.arg(...)` chain: + +```rust + .arg(&mut self.books_d) + .arg(&mut self.prev_mid_d) // NEW + .arg(&mut self.atr_mid_ema_d) // NEW + .arg(&n_backtests) +``` + +(Keep the existing arg ordering for everything else.) + +- [ ] **Step 5: Run test to verify it passes** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test stop_controller -- atr_ema_bootstrap_first_observation --ignored --nocapture +``` + +Expected: `test result: ok. 1 passed`. + +- [ ] **Step 6: Commit** + +```bash +git add crates/ml-backtesting/cuda/book_update.cu crates/ml-backtesting/src/sim/mod.rs crates/ml-backtesting/tests/stop_controller.rs +git commit -m "$(cat <<'EOF' +feat(ml-backtesting): ATR-EMA on mid-price in book_update_apply_snapshot + +Per-event Wiener-α=0.4 EMA on |Δmid| with first-observation bootstrap. +Floor source for the SL/trail-distance controller (spec §6). Thread 0 +of the broadcast-snapshot kernel handles the update per backtest; +threads 1..9 still handle the 10 level writes. + +Co-Authored-By: Claude Opus 4.7 (1M context) +EOF +)" +``` + +--- + +## Task 3: `stop_check_isv` helper + gate-on-flat in `decision_policy_default` + +**Files:** +- Modify: `crates/ml-backtesting/tests/stop_controller.rs` +- Modify: `crates/ml-backtesting/cuda/decision_policy.cu` +- Modify: `crates/ml-backtesting/src/sim/mod.rs` (kernel launch args for `decision_policy_default`) + +- [ ] **Step 1: Add the `stop_check_skipped_when_flat` test** + +Append to `crates/ml-backtesting/tests/stop_controller.rs`: + +```rust +use ml_backtesting::sim::{BatchedSimConfig, UniformSimParams}; + +fn cfg_default(n: usize) -> BatchedSimConfig { + BatchedSimConfig::from_uniform(n, &UniformSimParams { + target_annual_vol_units: 50.0, + annualisation_factor: 825.0, + max_lots: 5, + latency_ns: 0, + kelly_frac_floor: 0.20, + sharpe_weight_floor: 0.10, + threshold: 0.0, + cost_per_lot_per_side: 0.0, + use_cold_start_stopgap: false, + }) +} + +#[test] +#[ignore = "requires CUDA"] +fn stop_check_skipped_when_flat() -> Result<()> { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { eprintln!("skipping: cuda unavailable ({e})"); return Ok(()); } + }; + let mut sim = LobSimCuda::new(1, &dev)?; + let (bp, bs, ap, az) = level_book(5500.0, 0.25); + sim.apply_snapshot(&bp, &bs, &ap, &az)?; + + // Position is flat (just allocated, zeros). Stop check must not fire; + // entry path runs and broadcasts alpha decides the action. + sim.broadcast_alpha(&[0.8, 0.8, 0.8, 0.8, 0.8])?; + sim.step_decision_with_latency(0, &cfg_default(1))?; + + let (side, size) = sim.read_market_target(0)?; + // Strong bullish alpha + flat position → entry side=0 size>=1. NOT (3, 0). + assert_eq!(side, 0, "with flat position, stop must not write side=3; got side={side}"); + assert!(size >= 1, "entry should fire under p=0.8 + flat (size={size})"); + Ok(()) +} +``` + +- [ ] **Step 2: Run test to verify it fails (API: side=3 doesn't exist yet, but the test asserts side != 3 which would pass on main)** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test stop_controller -- stop_check_skipped_when_flat --ignored --nocapture +``` + +Expected: this test actually PASSES on main (no helper, flat-position alpha entry works). The failure mode kicks in once the helper is added — if the helper incorrectly fires on flat positions, the test then catches it. Document this in the commit: this is a regression-prevention test that protects against an over-eager helper. + +For TDD purposes we accept this — proceed to Step 3. + +- [ ] **Step 3: Add the `stop_check_isv` helper declaration in `decision_policy.cu`** + +In `crates/ml-backtesting/cuda/decision_policy.cu`, before `extern "C" __global__ void decision_policy_default` (around line 30), add: + +```c +// Per spec §3 + §5: shared stop-check helper called from both +// decision_policy_default and decision_policy_program. Pure-device, +// no host branches (CUDA Graph capture safe). +// +// Returns 1 if a stop fired (and market_targets[b] was written to +// the (3, 0) force-flat encoding); 0 otherwise. Caller checks return +// value to decide whether to skip the VM / sizing logic. +__device__ static int stop_check_isv( + int b, + const Pos* __restrict__ positions, + const unsigned char* __restrict__ isv_kelly_base, + const unsigned int* __restrict__ open_horizon_masks, + const float* __restrict__ atr_mid_ema, + float* __restrict__ trail_hwm, + const Book* __restrict__ books, + int* __restrict__ market_targets +) { + if (positions[b].position_lots == 0) { + return 0; // flat — entry path handles this + } + // Stop trigger logic added in Tasks 4-6. + return 0; +} +``` + +(Skeleton only — the SL / trail / averaging logic lands in Tasks 4-6. Tests for those will fail until then.) + +- [ ] **Step 4: Add helper call at top of `decision_policy_default`** + +In the same file, find the existing kernel body for `decision_policy_default` (around line 36-170). The existing kernel signature takes `Book* books` (or similar) as one of its args — verify by reading line 36-50. After the existing `if (program_lens[b] != 0u) return;` line (~line 56), add the helper call: + +```c + // Stop-check (spec §3) — pre-VM. Returns 1 if force-flat was written. + if (stop_check_isv(b, positions, isv_kelly_base, open_horizon_masks, + atr_mid_ema, trail_hwm, books, market_targets)) { + return; + } +``` + +If `decision_policy_default` does not currently take a `Book*` arg, add it as a new arg to the kernel signature now (spec §4.1 requires `const float* bid_px, const float* ask_px (latter two if not already present)` — but unifying as `const Book*` matches the existing book-update kernel and is cleaner). + +- [ ] **Step 5: Thread new args through `step_decision_with_latency` in `sim/mod.rs`** + +Find the launch of `decision_policy_default` (in `step_decision_with_latency`, around line 660). Append after the existing args: + +```rust + .arg(&self.atr_mid_ema_d) + .arg(&mut self.trail_hwm_d) + .arg(&self.books_d) // if books not already passed +``` + +Order must match the kernel signature. + +- [ ] **Step 6: Run the test to verify it still passes (regression-prevention)** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test stop_controller -- stop_check_skipped_when_flat --ignored --nocapture +``` + +Expected: PASS (helper returns 0 on flat position; entry path runs unchanged). + +- [ ] **Step 7: Run the previously-passing ATR test to confirm no regression** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test stop_controller -- atr_ema_bootstrap_first_observation --ignored --nocapture +``` + +Expected: PASS. + +- [ ] **Step 8: Commit** + +```bash +git add crates/ml-backtesting/cuda/decision_policy.cu crates/ml-backtesting/src/sim/mod.rs crates/ml-backtesting/tests/stop_controller.rs +git commit -m "$(cat <<'EOF' +feat(ml-backtesting): stop_check_isv __device__ helper + gate-on-flat + +Skeleton helper called from decision_policy_default at top of +per-backtest dispatch. Returns 0 (no-op) on flat positions; trigger +logic for open positions added in Tasks 4-6. Single source of truth +for stop logic across decision_policy_{default,program}. + +Co-Authored-By: Claude Opus 4.7 (1M context) +EOF +)" +``` + +--- + +## Task 4: Hard SL trigger + +**Files:** +- Modify: `crates/ml-backtesting/tests/stop_controller.rs` +- Modify: `crates/ml-backtesting/cuda/decision_policy.cu` (`stop_check_isv` body) + +- [ ] **Step 1: Add the `sl_fires_when_unrealized_breaks_distance` test** + +Append to `crates/ml-backtesting/tests/stop_controller.rs`: + +```rust +use ml_backtesting::policy::IsvKellyStateHost; + +#[test] +#[ignore = "requires CUDA"] +fn sl_fires_when_unrealized_breaks_distance() -> Result<()> { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { eprintln!("skipping: cuda unavailable ({e})"); return Ok(()); } + }; + let mut sim = LobSimCuda::new(1, &dev)?; + + // Seed isv: 5 horizons with pnl_ema_loss=2.0 (so sl_distance=2.0). + let seeded: [IsvKellyStateHost; 5] = std::array::from_fn(|_| IsvKellyStateHost { + pnl_ema_win: 1.0, + pnl_ema_loss: 2.0, + win_rate_ema: 0.5, + n_trades_seen: 20, // past MIN_TRADES_FOR_VAR_CAP + realised_return_var: 0.5, + recent_sharpe: 0.5, + }); + sim.write_isv_kelly(0, &seeded)?; + + // Bootstrap ATR + book at mid=5500. + let (bp, bs, ap, az) = level_book(5500.0, 0.25); + sim.apply_snapshot(&bp, &bs, &ap, &az)?; + // Event 2 so atr_mid_ema becomes nonzero (the test of stop firing + // independent of bootstrap-zero floor). + let (bp2, bs2, ap2, az2) = level_book(5500.1, 0.25); + sim.apply_snapshot(&bp2, &bs2, &ap2, &az2)?; + + // Open a long position via strong-alpha + step_decision + manual fill. + sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + sim.step_decision_with_latency(0, &cfg_default(1))?; + // For latency=0, the inflight-seed → fill cycle runs same-step on a + // matching market price. Walk the simulation a couple events to fill. + let mut ts: u64 = 1_000_000; + for _ in 0..3 { + sim.step_resting_orders(ts, 0.0)?; + ts += 1_000_000; + } + let pos_open = sim.read_pos(0)?; + assert!(pos_open.position_lots > 0, + "test setup: long position should open, got {}", pos_open.position_lots); + + // Drive mid down by 3.0 (>= sl_distance=2.0). Stop must fire. + let (bp3, bs3, ap3, az3) = level_book(5497.0, 0.25); + sim.apply_snapshot(&bp3, &bs3, &ap3, &az3)?; + sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + sim.step_decision_with_latency(ts, &cfg_default(1))?; + + let (side, size) = sim.read_market_target(0)?; + assert_eq!(side, 3, "SL must fire with force-flat (3, 0); got side={side}"); + assert_eq!(size, 0, "force-flat encoding has abs_sz=0; got size={size}"); + Ok(()) +} +``` + +- [ ] **Step 2: Run test to verify it fails** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test stop_controller -- sl_fires_when_unrealized_breaks_distance --ignored --nocapture +``` + +Expected: FAIL on `assert_eq!(side, 3, ...)` — currently the helper returns 0 and the entry path writes whatever alpha says. + +- [ ] **Step 3: Implement SL trigger in `stop_check_isv`** + +Replace the helper body in `decision_policy.cu` with: + +```c +__device__ static int stop_check_isv( + int b, + const Pos* __restrict__ positions, + const unsigned char* __restrict__ isv_kelly_base, + const unsigned int* __restrict__ open_horizon_masks, + const float* __restrict__ atr_mid_ema, + float* __restrict__ trail_hwm, + const Book* __restrict__ books, + int* __restrict__ market_targets +) { + const Pos& pos = positions[b]; + if (pos.position_lots == 0) return 0; + + // §5 controller math. + const float atr = atr_mid_ema[b]; + const unsigned int mask = open_horizon_masks[b]; + float ema_loss = 0.0f; + int n_open = 0; + #pragma unroll + for (int h = 0; h < N_HORIZONS; ++h) { + if (mask & (1u << h)) { + const IsvKellyState* isv = reinterpret_cast( + isv_kelly_base + (size_t)b * N_HORIZONS * sizeof(IsvKellyState) + ); + ema_loss += isv[h].pnl_ema_loss; + n_open += 1; + } + } + if (n_open > 0) ema_loss /= (float)n_open; + + const float sl_distance = fmaxf(ema_loss, atr); + + const bool is_long = pos.position_lots > 0; + const Book& bk = books[b]; + const float close_px = is_long ? bk.bid_px[0] : bk.ask_px[0]; + const float entry_px = pos.vwap_entry; + const float unrealized_pl_per_lot = + is_long ? (close_px - entry_px) : (entry_px - close_px); + + const bool sl_fired = unrealized_pl_per_lot <= -sl_distance; + + if (sl_fired) { + market_targets[b * 2 + 0] = 3; // §7 force-flat encoding + market_targets[b * 2 + 1] = 0; + return 1; + } + return 0; +} +``` + +(The trail logic + multi-horizon ema_win averaging join in Tasks 5-6.) + +- [ ] **Step 4: Run test to verify it passes** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test stop_controller -- sl_fires_when_unrealized_breaks_distance --ignored --nocapture +``` + +Expected: PASS. + +- [ ] **Step 5: Run all stop_controller tests to confirm no regression** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test stop_controller -- --ignored --nocapture +``` + +Expected: 3 passed (`atr_ema_bootstrap_first_observation`, `stop_check_skipped_when_flat`, `sl_fires_when_unrealized_breaks_distance`). + +- [ ] **Step 6: Commit** + +```bash +git add crates/ml-backtesting/cuda/decision_policy.cu crates/ml-backtesting/tests/stop_controller.rs +git commit -m "$(cat <<'EOF' +feat(ml-backtesting): hard-SL trigger in stop_check_isv + +sl_distance = max(pnl_ema_loss, atr_mid_ema) per pearl_blend_formulas +permanent floor. Fires force-flat (3, 0) when unrealized_pl_per_lot +drops below -sl_distance. Multi-horizon mask averaging + trail trigger +added in Tasks 5-6. + +Co-Authored-By: Claude Opus 4.7 (1M context) +EOF +)" +``` + +--- + +## Task 5: Trail trigger + HWM ratchet + +**Files:** +- Modify: `crates/ml-backtesting/tests/stop_controller.rs` +- Modify: `crates/ml-backtesting/cuda/decision_policy.cu` (extend `stop_check_isv`) + +- [ ] **Step 1: Add the `trail_arms_then_fires` test** + +Append: + +```rust +#[test] +#[ignore = "requires CUDA"] +fn trail_arms_then_fires() -> Result<()> { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { eprintln!("skipping: cuda unavailable ({e})"); return Ok(()); } + }; + let mut sim = LobSimCuda::new(1, &dev)?; + + // Seeded isv: pnl_ema_win=1.5 (so trail_distance=1.5 with atr floor below). + let seeded: [IsvKellyStateHost; 5] = std::array::from_fn(|_| IsvKellyStateHost { + pnl_ema_win: 1.5, + pnl_ema_loss: 10.0, // big SL so trail fires first + win_rate_ema: 0.5, + n_trades_seen: 20, + realised_return_var: 0.5, + recent_sharpe: 0.5, + }); + sim.write_isv_kelly(0, &seeded)?; + + let (bp, bs, ap, az) = level_book(5500.0, 0.25); + sim.apply_snapshot(&bp, &bs, &ap, &az)?; + let (bp2, bs2, ap2, az2) = level_book(5500.1, 0.25); + sim.apply_snapshot(&bp2, &bs2, &ap2, &az2)?; + + // Open long. + sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + sim.step_decision_with_latency(0, &cfg_default(1))?; + let mut ts: u64 = 1_000_000; + for _ in 0..3 { + sim.step_resting_orders(ts, 0.0)?; + ts += 1_000_000; + } + assert!(sim.read_pos(0)?.position_lots > 0, "setup: long opened"); + + // Walk mid up to ratchet trail HWM past trail_distance=1.5. + let mid_up = [5500.5, 5501.0, 5502.0, 5502.5]; + for &m in &mid_up { + let (b, bs, a, asz) = level_book(m, 0.25); + sim.apply_snapshot(&b, &bs, &a, &asz)?; + sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + sim.step_decision_with_latency(ts, &cfg_default(1))?; + ts += 1_000_000; + } + // HWM should now be ~2.5 (5502.5 - 5500 entry). Trail not fired (still rising). + let hwm_armed = sim.read_trail_hwm(0)?; + assert!(hwm_armed > 1.5, "HWM should ratchet past trail_distance=1.5, got {hwm_armed}"); + + // Drop mid by trail_distance+ε to fire trail. + let (b, bs, a, asz) = level_book(5500.5, 0.25); // drop from 5502.5 to 5500.5 = 2.0 + sim.apply_snapshot(&b, &bs, &a, &asz)?; + sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + sim.step_decision_with_latency(ts, &cfg_default(1))?; + + let (side, size) = sim.read_market_target(0)?; + assert_eq!(side, 3, "trail must fire force-flat; got side={side}"); + assert_eq!(size, 0); + Ok(()) +} +``` + +- [ ] **Step 2: Run test to verify it fails** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test stop_controller -- trail_arms_then_fires --ignored --nocapture +``` + +Expected: FAIL — helper currently only does SL. + +- [ ] **Step 3: Extend `stop_check_isv` with trail logic** + +In `decision_policy.cu`, modify the helper to add trail logic. Insert after the `n_open` averaging block and before the `sl_distance` line: + +```c + float ema_win = 0.0f; + if (n_open > 0) { + // Recompute mean of pnl_ema_win across the same open-mask bits. + ema_win = 0.0f; + const IsvKellyState* isv_for_win = reinterpret_cast( + isv_kelly_base + (size_t)b * N_HORIZONS * sizeof(IsvKellyState) + ); + #pragma unroll + for (int h = 0; h < N_HORIZONS; ++h) { + if (mask & (1u << h)) { + ema_win += isv_for_win[h].pnl_ema_win; + } + } + ema_win /= (float)n_open; + } +``` + +After the `sl_distance` line, add: + +```c + const float trail_distance = fmaxf(ema_win, atr); + + // Trail ratchet (single thread per backtest, no atomics). + const float new_hwm = fmaxf(trail_hwm[b], unrealized_pl_per_lot); + trail_hwm[b] = new_hwm; + const bool trail_fired = + (new_hwm > trail_distance) + && (unrealized_pl_per_lot <= new_hwm - trail_distance); +``` + +Change the fire condition from `if (sl_fired)` to `if (sl_fired || trail_fired)`. + +- [ ] **Step 4: Run test to verify it passes** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test stop_controller -- trail_arms_then_fires --ignored --nocapture +``` + +Expected: PASS. + +- [ ] **Step 5: Re-run prior tests for regression** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test stop_controller -- --ignored --nocapture +``` + +Expected: 4 passed. + +- [ ] **Step 6: Commit** + +```bash +git add crates/ml-backtesting/cuda/decision_policy.cu crates/ml-backtesting/tests/stop_controller.rs +git commit -m "$(cat <<'EOF' +feat(ml-backtesting): trail-TP trigger with HWM ratchet in stop_check_isv + +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) +EOF +)" +``` + +--- + +## Task 6: Multi-horizon mask averaging invariant test + +**Files:** +- Modify: `crates/ml-backtesting/tests/stop_controller.rs` + +(The averaging logic was already implemented in Tasks 4–5. This task adds the explicit invariant test that catches bit-pick-first / bit-pick-max regressions in future edits.) + +- [ ] **Step 1: Add the `multi_horizon_mask_averages_emas` test** + +Append: + +```rust +#[test] +#[ignore = "requires CUDA"] +fn multi_horizon_mask_averages_emas() -> Result<()> { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { eprintln!("skipping: cuda unavailable ({e})"); return Ok(()); } + }; + let mut sim = LobSimCuda::new(1, &dev)?; + + // Seed isv: h0.pnl_ema_loss=2.0, h1.pnl_ema_loss=4.0, others=10.0. + // With open_horizon_masks=0b00011 the mean = (2+4)/2 = 3.0. + let mut seeded: [IsvKellyStateHost; 5] = std::array::from_fn(|_| IsvKellyStateHost { + pnl_ema_win: 0.5, + pnl_ema_loss: 10.0, + win_rate_ema: 0.5, + n_trades_seen: 20, + realised_return_var: 0.5, + recent_sharpe: 0.5, + }); + seeded[0].pnl_ema_loss = 2.0; + seeded[1].pnl_ema_loss = 4.0; + sim.write_isv_kelly(0, &seeded)?; + + let (bp, bs, ap, az) = level_book(5500.0, 0.25); + sim.apply_snapshot(&bp, &bs, &ap, &az)?; + let (bp2, bs2, ap2, az2) = level_book(5500.1, 0.25); + sim.apply_snapshot(&bp2, &bs2, &ap2, &az2)?; + + sim.broadcast_alpha(&[0.95, 0.95, 0.5, 0.5, 0.5])?; // h0+h1 strongest + sim.step_decision_with_latency(0, &cfg_default(1))?; + let mut ts: u64 = 1_000_000; + for _ in 0..3 { + sim.step_resting_orders(ts, 0.0)?; + ts += 1_000_000; + } + let pos_open = sim.read_pos(0)?; + assert!(pos_open.position_lots > 0, "setup: long opens, got {}", pos_open.position_lots); + + // Drive mid down by 2.5 — between the per-h0 SL (2.0) and per-h1 SL (4.0). + // With mean=3.0 SL, 2.5 < 3.0 → trail/SL does NOT fire. + let (bp3, bs3, ap3, az3) = level_book(5497.5, 0.25); + sim.apply_snapshot(&bp3, &bs3, &ap3, &az3)?; + sim.broadcast_alpha(&[0.95, 0.95, 0.5, 0.5, 0.5])?; + sim.step_decision_with_latency(ts, &cfg_default(1))?; + let (side_at_25, _) = sim.read_market_target(0)?; + assert_ne!(side_at_25, 3, "Δ=2.5 < mean_sl=3.0 must not fire; got side={side_at_25}"); + ts += 1_000_000; + + // Drive mid down further to 3.5 total → > mean_sl=3.0 → fires. + let (bp4, bs4, ap4, az4) = level_book(5496.5, 0.25); + sim.apply_snapshot(&bp4, &bs4, &ap4, &az4)?; + sim.broadcast_alpha(&[0.95, 0.95, 0.5, 0.5, 0.5])?; + sim.step_decision_with_latency(ts, &cfg_default(1))?; + let (side_at_35, size_at_35) = sim.read_market_target(0)?; + assert_eq!(side_at_35, 3, "Δ=3.5 > mean_sl=3.0 must fire; got side={side_at_35}"); + assert_eq!(size_at_35, 0); + Ok(()) +} +``` + +- [ ] **Step 2: Run test to verify it passes (averaging already implemented in Task 4-5)** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test stop_controller -- multi_horizon_mask_averages_emas --ignored --nocapture +``` + +Expected: PASS — verifies the averaging logic by exercising the boundary between per-h0 SL=2.0 and per-h1 SL=4.0. If a future edit collapses to bit-pick-first (would use h0 SL=2.0) the Δ=2.5 case would fire (FAIL on `assert_ne!`). Bit-pick-max (h1 SL=4.0) would not fire at Δ=3.5 (FAIL on `assert_eq!`). + +- [ ] **Step 3: Commit** + +```bash +git add crates/ml-backtesting/tests/stop_controller.rs +git commit -m "$(cat <<'EOF' +test(ml-backtesting): multi-horizon mask averaging invariant test + +Boundary test that fails on bit-pick-first or bit-pick-max regressions +of the open_horizon_masks averaging in stop_check_isv. + +Co-Authored-By: Claude Opus 4.7 (1M context) +EOF +)" +``` + +--- + +## Task 7: Bytecode VM parity — call `stop_check_isv` from `decision_policy_program` + +**Files:** +- Modify: `crates/ml-backtesting/tests/stop_controller.rs` +- Modify: `crates/ml-backtesting/cuda/decision_policy.cu` (`decision_policy_program`) +- Modify: `crates/ml-backtesting/src/sim/mod.rs` (kernel launch args) + +- [ ] **Step 1: Add the `bytecode_and_default_kernel_agree_on_stops` test** + +Append: + +```rust +use ml_backtesting::policy::{Strategy, StrategyConfig, EnsembleAggregator, SizingPolicyId, StopRules, N_HORIZONS}; + +#[test] +#[ignore = "requires CUDA"] +fn bytecode_and_default_kernel_agree_on_stops() -> Result<()> { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { eprintln!("skipping: cuda unavailable ({e})"); return Ok(()); } + }; + + // Run scenario A: default kernel, p=0.95 with SL setup → expect (3, 0). + let mut sim_a = LobSimCuda::new(1, &dev)?; + let seeded: [IsvKellyStateHost; 5] = std::array::from_fn(|_| IsvKellyStateHost { + pnl_ema_win: 1.0, pnl_ema_loss: 2.0, win_rate_ema: 0.5, + n_trades_seen: 20, realised_return_var: 0.5, recent_sharpe: 0.5, + }); + sim_a.write_isv_kelly(0, &seeded)?; + let (bp, bs, ap, az) = level_book(5500.0, 0.25); + sim_a.apply_snapshot(&bp, &bs, &ap, &az)?; + let (bp2, bs2, ap2, az2) = level_book(5500.1, 0.25); + sim_a.apply_snapshot(&bp2, &bs2, &ap2, &az2)?; + sim_a.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + sim_a.step_decision_with_latency(0, &cfg_default(1))?; + let mut ts: u64 = 1_000_000; + for _ in 0..3 { sim_a.step_resting_orders(ts, 0.0)?; ts += 1_000_000; } + let (bp3, bs3, ap3, az3) = level_book(5497.0, 0.25); // Δ = 3.0 > sl=2.0 + sim_a.apply_snapshot(&bp3, &bs3, &ap3, &az3)?; + sim_a.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + sim_a.step_decision_with_latency(ts, &cfg_default(1))?; + let mt_a = sim_a.read_market_target(0)?; + + // Run scenario B: same scenario through decision_policy_program via upload. + let mut sim_b = LobSimCuda::new(1, &dev)?; + let prog = Strategy::Ensemble { + children: (0..N_HORIZONS as u8) + .map(|h| Strategy::Leaf(StrategyConfig { + horizon_idx: h, + sizing_policy: SizingPolicyId::IsvKelly, + sl_tp_rules: StopRules::default(), + max_concurrent_lots: 5, + })) + .collect(), + aggregator: EnsembleAggregator::MaxConfidence, + }; + sim_b.upload_program(0, &prog.flatten())?; + sim_b.write_isv_kelly(0, &seeded)?; + let (bp, bs, ap, az) = level_book(5500.0, 0.25); + sim_b.apply_snapshot(&bp, &bs, &ap, &az)?; + let (bp2, bs2, ap2, az2) = level_book(5500.1, 0.25); + sim_b.apply_snapshot(&bp2, &bs2, &ap2, &az2)?; + sim_b.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + sim_b.step_decision_with_latency(0, &cfg_default(1))?; + let mut ts: u64 = 1_000_000; + for _ in 0..3 { sim_b.step_resting_orders(ts, 0.0)?; ts += 1_000_000; } + let (bp3, bs3, ap3, az3) = level_book(5497.0, 0.25); + sim_b.apply_snapshot(&bp3, &bs3, &ap3, &az3)?; + sim_b.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + sim_b.step_decision_with_latency(ts, &cfg_default(1))?; + let mt_b = sim_b.read_market_target(0)?; + + assert_eq!(mt_a, mt_b, + "default kernel vs program kernel produce different market_target: a={mt_a:?} b={mt_b:?}"); + assert_eq!(mt_a.0, 3, "both kernels must fire stop force-flat; got {:?}", mt_a); + Ok(()) +} +``` + +- [ ] **Step 2: Run test to verify it fails** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test stop_controller -- bytecode_and_default_kernel_agree_on_stops --ignored --nocapture +``` + +Expected: FAIL — `decision_policy_program` does not yet call `stop_check_isv`, so it produces whatever the VM decides (likely (0, n) for strong bullish alpha), not (3, 0). + +- [ ] **Step 3: Add helper call at top of `decision_policy_program`** + +In `decision_policy.cu`, find `decision_policy_program` (around line 206). After its early-return on `program_lens[b] == 0` (around line 230), and BEFORE the threshold gate (line 235), add: + +```c + // Stop-check (spec §3) — same helper as decision_policy_default. + if (stop_check_isv(b, positions, isv_kelly_base, open_horizon_masks, + atr_mid_ema, trail_hwm, books, market_targets)) { + return; + } +``` + +- [ ] **Step 4: Thread new args through `decision_policy_program` launch in `sim/mod.rs`** + +Find the second launch site (the one for `decision_policy_program`, around line 681 in `sim/mod.rs`). Append the same new args: + +```rust + .arg(&self.atr_mid_ema_d) + .arg(&mut self.trail_hwm_d) + .arg(&self.books_d) +``` + +Match the order to the kernel signature. + +- [ ] **Step 5: Run test to verify it passes** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test stop_controller -- bytecode_and_default_kernel_agree_on_stops --ignored --nocapture +``` + +Expected: PASS. + +- [ ] **Step 6: Re-run all stop_controller tests for regression** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test stop_controller -- --ignored --nocapture +``` + +Expected: 6 passed. + +- [ ] **Step 7: Commit** + +```bash +git add crates/ml-backtesting/cuda/decision_policy.cu crates/ml-backtesting/src/sim/mod.rs crates/ml-backtesting/tests/stop_controller.rs +git commit -m "$(cat <<'EOF' +feat(ml-backtesting): wire stop_check_isv into decision_policy_program + +Both decision kernels now call the same __device__ helper at the top +of per-backtest dispatch. 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) +EOF +)" +``` + +--- + +## Task 8: Target-delta in `seed_inflight_limits_batched` + in-flight summation + `side=3` + +**Files:** +- Modify: `crates/ml-backtesting/tests/stop_controller.rs` +- Modify: `crates/ml-backtesting/cuda/resting_orders.cu` (`seed_inflight_limits_batched`) + +- [ ] **Step 1: Add the three target-semantics tests** + +Append: + +```rust +#[test] +#[ignore = "requires CUDA"] +fn position_target_not_additive() -> Result<()> { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { eprintln!("skipping: cuda unavailable ({e})"); return Ok(()); } + }; + let mut sim = LobSimCuda::new(1, &dev)?; + let (bp, bs, ap, az) = level_book(5500.0, 0.25); + sim.apply_snapshot(&bp, &bs, &ap, &az)?; + // Force target = (0, 3) via strong-alpha that produces lots=3 under + // cfg_default's max_lots=5 + kelly_frac_floor=0.20: + // sig_mag = |0.95 - 0.5| * 2 = 0.9 + // ss = 1 * 0.9 * 0.20 * 5 = 0.9 → truncf(0.9 + 0.5) = 1. + // So we need an aggressive isv warm-start to get to 3 lots... simpler: + // use the smoke path's natural output and verify pos cap-limited. + sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + let mut ts: u64 = 0; + for _ in 0..10 { + sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + sim.step_decision_with_latency(ts, &cfg_default(1))?; + sim.step_resting_orders(ts, 0.0)?; + ts += 1_000_000; + } + let pos = sim.read_pos(0)?; + assert!(pos.position_lots.abs() <= 5, + "|position_lots| must stay <= max_lots=5 with target semantics; got {}", + pos.position_lots); + Ok(()) +} + +#[test] +#[ignore = "requires CUDA"] +fn position_target_not_additive_with_latency() -> Result<()> { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { eprintln!("skipping: cuda unavailable ({e})"); return Ok(()); } + }; + let mut sim = LobSimCuda::new(1, &dev)?; + let (bp, bs, ap, az) = level_book(5500.0, 0.25); + sim.apply_snapshot(&bp, &bs, &ap, &az)?; + + let cfg_lat = BatchedSimConfig::from_uniform(1, &UniformSimParams { + target_annual_vol_units: 50.0, annualisation_factor: 825.0, + max_lots: 5, latency_ns: 200_000_000, // 200ms latency + kelly_frac_floor: 0.20, sharpe_weight_floor: 0.10, + threshold: 0.0, cost_per_lot_per_side: 0.0, + use_cold_start_stopgap: false, + }); + + // 200 events of persistent target = (0, n) with n in (0..5] from alpha. + let mut ts: u64 = 0; + for _ in 0..200 { + sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + sim.step_decision_with_latency(ts, &cfg_lat)?; + sim.step_resting_orders(ts, 0.0)?; + ts += 1_000_000; + } + // Drain: 300 more events with no new alpha-change so all in-flight fills. + for _ in 0..300 { + sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + sim.step_decision_with_latency(ts, &cfg_lat)?; + sim.step_resting_orders(ts, 0.0)?; + ts += 1_000_000; + } + let pos = sim.read_pos(0)?; + assert!(pos.position_lots.abs() <= 5, + "|position_lots| must stay <= max_lots=5 even under 200ms latency; got {}", + pos.position_lots); + Ok(()) +} + +#[test] +#[ignore = "requires CUDA"] +fn alpha_noop_side_2_preserved() -> Result<()> { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { eprintln!("skipping: cuda unavailable ({e})"); return Ok(()); } + }; + let mut sim = LobSimCuda::new(1, &dev)?; + + // Get to long-3 first with strong alpha. + let (bp, bs, ap, az) = level_book(5500.0, 0.25); + sim.apply_snapshot(&bp, &bs, &ap, &az)?; + let mut ts: u64 = 0; + for _ in 0..20 { + sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + sim.step_decision_with_latency(ts, &cfg_default(1))?; + sim.step_resting_orders(ts, 0.0)?; + ts += 1_000_000; + } + let pos_open = sim.read_pos(0)?.position_lots; + assert!(pos_open > 0, "setup: long open, got {}", pos_open); + + // Force weak alpha that produces lots=0 (sig_mag * 0.20 * 5 < 0.5). + // With p=0.55: sig_mag=0.1, ss=0.1, truncf(0.1+0.5)=0 → (2, 0). + // Verify position unchanged across many weak-alpha events. + for _ in 0..20 { + sim.broadcast_alpha(&[0.55, 0.55, 0.55, 0.55, 0.55])?; + sim.step_decision_with_latency(ts, &cfg_default(1))?; + sim.step_resting_orders(ts, 0.0)?; + ts += 1_000_000; + } + let pos_after = sim.read_pos(0)?.position_lots; + assert_eq!(pos_after, pos_open, + "side=2 must preserve position; opened={} after 20 weak-alpha events={}", + pos_open, pos_after); + Ok(()) +} +``` + +- [ ] **Step 2: Run tests to verify they fail** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test stop_controller -- position_target_not_additive --ignored --nocapture +``` + +Expected: FAIL — position_lots accumulates past 5. + +- [ ] **Step 3: Modify `seed_inflight_limits_batched` to compute target-delta with in-flight summation** + +Replace the kernel body in `crates/ml-backtesting/cuda/resting_orders.cu` (line 439–475): + +```c +extern "C" __global__ void seed_inflight_limits_batched( + unsigned char* __restrict__ orders_base, // [n_backtests * orders_bytes] + int orders_bytes, + Pos* __restrict__ positions, // [n_backtests] + const int* __restrict__ market_targets, // [n_backtests * 2] + const unsigned int* __restrict__ latency_ns_per_b, // [n_backtests] + unsigned long long current_ts_ns, + int n_backtests +) { + int b = blockIdx.x; + if (b >= n_backtests || threadIdx.x != 0) return; + + const int side = market_targets[b * 2 + 0]; + const int abs_sz = market_targets[b * 2 + 1]; + + if (side == 2) return; // §7: no-op = preserve position + + int target_signed; + if (side == 0) target_signed = +abs_sz; + else if (side == 1) target_signed = -abs_sz; + else target_signed = 0; // side == 3: force flat + + Orders* orders = reinterpret_cast(orders_base + (size_t)b * orders_bytes); + + // Effective position = filled position + all unfilled signed lots. + int effective = positions[b].position_lots; + #pragma unroll + for (int s_idx = 0; s_idx < MAX_LIMITS; ++s_idx) { + const LimitSlot& s = orders->limits[s_idx]; + if (s.active == 0) continue; + const int slot_lots = (int)s.size_remaining; + effective += (s.side == 0) ? slot_lots : -slot_lots; + } + + const int delta = target_signed - effective; + if (delta == 0) return; + + const int order_lots = (delta > 0) ? delta : -delta; + const int order_side = (delta > 0) ? 0 : 1; + + // Find a free slot and seed an in-flight order. + for (int s_idx = 0; s_idx < MAX_LIMITS; ++s_idx) { + LimitSlot& s = orders->limits[s_idx]; + if (s.active == 0) { + s.active = 2; // in-flight + s.side = (unsigned char)order_side; + s.kind = 1; // Limit + s.oco_pair = 0xFF; + s.price = (order_side == 0) ? 1.0e9f : 0.0f; // always crosses + s.size_remaining = (float)order_lots; + s.queue_position = 0.0f; + s.arrival_ts_ns = current_ts_ns + (unsigned long long)latency_ns_per_b[b]; + s.gen_counter = (unsigned short)(s.gen_counter + 1); + return; + } + } + positions[b].submission_overflow += 1u; +} +``` + +- [ ] **Step 4: Run all three tests to verify they pass** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test stop_controller -- position_target --ignored --nocapture +SQLX_OFFLINE=true cargo test -p ml-backtesting --test stop_controller -- alpha_noop_side_2_preserved --ignored --nocapture +``` + +Expected: all 3 pass. + +- [ ] **Step 5: Re-run all stop_controller tests** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test stop_controller -- --ignored --nocapture +``` + +Expected: 9 passed. + +- [ ] **Step 6: Commit** + +```bash +git add crates/ml-backtesting/cuda/resting_orders.cu crates/ml-backtesting/tests/stop_controller.rs +git commit -m "$(cat <<'EOF' +feat(ml-backtesting): target-delta semantics in seed_inflight_limits_batched + +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). + +Co-Authored-By: Claude Opus 4.7 (1M context) +EOF +)" +``` + +--- + +## Task 9: `pnl_track_step` trail_hwm reset on close + +**Files:** +- Modify: `crates/ml-backtesting/tests/stop_controller.rs` +- Modify: `crates/ml-backtesting/cuda/pnl_track.cu` +- Modify: `crates/ml-backtesting/src/sim/mod.rs` (kernel launch args for `pnl_track_step`) + +- [ ] **Step 1: Add the `trail_hwm_reset_on_close` test** + +Append: + +```rust +#[test] +#[ignore = "requires CUDA"] +fn trail_hwm_reset_on_close() -> Result<()> { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { eprintln!("skipping: cuda unavailable ({e})"); return Ok(()); } + }; + let mut sim = LobSimCuda::new(1, &dev)?; + + let seeded: [IsvKellyStateHost; 5] = std::array::from_fn(|_| IsvKellyStateHost { + pnl_ema_win: 1.5, pnl_ema_loss: 10.0, win_rate_ema: 0.5, + n_trades_seen: 20, realised_return_var: 0.5, recent_sharpe: 0.5, + }); + sim.write_isv_kelly(0, &seeded)?; + + let (bp, bs, ap, az) = level_book(5500.0, 0.25); + sim.apply_snapshot(&bp, &bs, &ap, &az)?; + let (bp2, bs2, ap2, az2) = level_book(5500.1, 0.25); + sim.apply_snapshot(&bp2, &bs2, &ap2, &az2)?; + + // Open long. + sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + let mut ts: u64 = 0; + sim.step_decision_with_latency(ts, &cfg_default(1))?; + for _ in 0..3 { sim.step_resting_orders(ts, 0.0)?; ts += 1_000_000; } + assert!(sim.read_pos(0)?.position_lots > 0, "setup: long opens"); + + // Ratchet HWM up to high value. + for m in [5501.0, 5502.0, 5503.0, 5504.0, 5505.0] { + let (b, bs, a, asz) = level_book(m, 0.25); + sim.apply_snapshot(&b, &bs, &a, &asz)?; + sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + sim.step_decision_with_latency(ts, &cfg_default(1))?; + sim.step_resting_orders(ts, 0.0)?; + ts += 1_000_000; + } + let hwm_armed = sim.read_trail_hwm(0)?; + assert!(hwm_armed > 1.5, "HWM ratcheted up, got {hwm_armed}"); + + // Trigger trail-close via price drop. + let (b, bs, a, asz) = level_book(5502.0, 0.25); // drop 3.0 from peak 5505 + sim.apply_snapshot(&b, &bs, &a, &asz)?; + sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?; + sim.step_decision_with_latency(ts, &cfg_default(1))?; + let (s, _) = sim.read_market_target(0)?; + assert_eq!(s, 3, "trail must force-flat on drop"); + // Drain to actually close. + for _ in 0..5 { sim.step_resting_orders(ts, 0.0)?; sim.step_pnl_track(ts)?; ts += 1_000_000; } + let pos_closed = sim.read_pos(0)?; + assert_eq!(pos_closed.position_lots, 0, "position must flatten"); + let hwm_after_close = sim.read_trail_hwm(0)?; + assert_eq!(hwm_after_close, 0.0, + "pnl_track_step must reset trail_hwm to 0 on close; got {hwm_after_close}"); + Ok(()) +} +``` + +- [ ] **Step 2: Run test to verify it fails** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test stop_controller -- trail_hwm_reset_on_close --ignored --nocapture +``` + +Expected: FAIL on `assert_eq!(hwm_after_close, 0.0, ...)` — pnl_track doesn't yet reset. + +- [ ] **Step 3: Modify `pnl_track_step` in `crates/ml-backtesting/cuda/pnl_track.cu`** + +Add `trail_hwm` to the kernel signature and reset it inside the close branch. + +```c +extern "C" __global__ void pnl_track_step( + unsigned char* pos_base, + unsigned char* open_trade_state, + unsigned char* trade_log_base, + unsigned int* trade_log_head, + float* trail_hwm, // NEW [n_backtests] + unsigned long long current_ts_ns, + int trade_log_cap, + int n_backtests +) { + int b = blockIdx.x; + if (b >= n_backtests || threadIdx.x != 0) return; + + Pos* pos = reinterpret_cast(pos_base + (size_t)b * sizeof(Pos)); + unsigned char* st = open_trade_state + (size_t)b * OPEN_TRADE_STATE_BYTES; + int prev_size = *reinterpret_cast(st + 12); + int now_size = pos->position_lots; + + if (prev_size == 0 && now_size != 0) { + // ... existing entry-snapshot code unchanged ... + } else if (prev_size != 0 && now_size == 0) { + // Close — emit TradeRecord (existing code). + // ... existing TradeRecord emission unchanged ... + + // NEW: reset trail_hwm so next entry starts fresh. + trail_hwm[b] = 0.0f; + } +} +``` + +(Keep all existing fields and code blocks. Add only the new arg + the one-line reset in the close branch.) + +- [ ] **Step 4: Thread the new arg through `step_pnl_track` in `sim/mod.rs`** + +Find the `step_pnl_track` method (around line 493) and append `&mut self.trail_hwm_d` to the kernel launch arg chain. + +- [ ] **Step 5: Run test to verify it passes** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test stop_controller -- trail_hwm_reset_on_close --ignored --nocapture +``` + +Expected: PASS. + +- [ ] **Step 6: Run all stop_controller tests for regression** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test stop_controller -- --ignored --nocapture +``` + +Expected: 10 passed. + +- [ ] **Step 7: Commit** + +```bash +git add crates/ml-backtesting/cuda/pnl_track.cu crates/ml-backtesting/src/sim/mod.rs crates/ml-backtesting/tests/stop_controller.rs +git commit -m "$(cat <<'EOF' +feat(ml-backtesting): pnl_track_step resets trail_hwm on close transition + +Two-instruction addition to the existing close-emission branch +(prev != 0 && now == 0). Without the reset, the next entry inherits +a stale HWM that could arm the trail spuriously on event 0 of the +new trade. + +Co-Authored-By: Claude Opus 4.7 (1M context) +EOF +)" +``` + +--- + +## Task 10: Atomic cleanup — delete `StopRules` + `use_cold_start_stopgap` + retarget integration test + +**Files:** +- Delete: `crates/ml-backtesting/src/policy/stop_rules.rs` +- Modify: `crates/ml-backtesting/src/policy/mod.rs` +- Modify: `crates/ml-backtesting/src/harness.rs` +- Modify: `crates/ml-backtesting/src/sim/batched_config.rs` +- Modify: `bin/fxt-backtest/src/main.rs` +- Modify: `config/ml/sweep_smoke.yaml` +- Modify: `crates/ml-backtesting/tests/decision_floor_coldstart.rs` +- Modify: `crates/ml-backtesting/tests/threshold_and_cost.rs` +- Modify: `crates/ml-backtesting/tests/lob_sim_fixtures.rs` +- Modify: `crates/ml-backtesting/tests/lob_sim_integrated_fuzz.rs` +- Modify: `crates/ml-backtesting/tests/parallel_sim_correctness.rs` +- Modify: `crates/ml-backtesting/tests/stop_controller.rs` (drop `StopRules` import) +- Modify: `docs/superpowers/specs/2026-05-19-cbsw-cold-start-aggregator-design.md` +- Modify: `docs/superpowers/plans/2026-05-19-cbsw-cold-start-aggregator.md` + +- [ ] **Step 1: Retarget the existing integration test (decision_floor_coldstart.rs)** + +Replace the test body of `cold_start_persistent_bullish_with_default_stops_never_closes` (around line 130-200) with: + +```rust +/// Integration retarget from the falsified-spec sibling test. Asserts the +/// stop controller closes positions and respects the |pos| <= max_lots cap. +#[test] +#[ignore = "requires CUDA"] +fn cold_start_persistent_bullish_now_closes() -> Result<()> { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { eprintln!("skipping: cuda unavailable ({e})"); return Ok(()); } + }; + let mut sim = LobSimCuda::new(1, &dev)?; + + let mid: f32 = 5500.00; + let tick: f32 = 0.25; + let mut bid_px = [0.0; 10]; let mut bid_sz = [0.0; 10]; + let mut ask_px = [0.0; 10]; let mut ask_sz = [0.0; 10]; + for k in 0..10 { + bid_px[k] = mid - tick * (k as f32 + 1.0); + ask_px[k] = mid + tick * (k as f32 + 1.0); + bid_sz[k] = 20.0 + 10.0 * k as f32; + ask_sz[k] = 20.0 + 10.0 * k as f32; + } + + let cfg = cfg_uniform(1, 0.20, 0.10); + let mut ts: u64 = 1_000_000_000; + for _ in 0..200 { + sim.apply_snapshot(&bid_px, &bid_sz, &ask_px, &ask_sz)?; + sim.step_resting_orders(ts, 0.0)?; + sim.step_pnl_track(ts)?; + sim.broadcast_alpha(&[0.8, 0.8, 0.8, 0.8, 0.8])?; + sim.step_decision_with_latency(ts, &cfg)?; + ts += 1_000_000; + } + let pos = sim.read_pos(0)?; + let trades = sim.read_total_trade_count()?; + eprintln!("after 200 steps: position_lots={} trades_closed={}", + pos.position_lots, trades); + assert!(pos.position_lots.abs() as i32 <= 5, + "|position_lots| must respect max_lots=5, got {}", pos.position_lots); + assert!(trades > 0, + "stop controller must produce closed trades over 200 events of bullish alpha; got {}", + trades); + Ok(()) +} +``` + +Also drop the now-unused `_` prefix or whatever framing the original test had. Delete `cfg_uniform` `use_cold_start_stopgap` field (next step covers). + +- [ ] **Step 2: Run retargeted test — must FAIL because StopRules is still wired but the kernel ignores it** + +Wait, no — the kernel was extended in Tasks 4-9 to do real stops. Run it. + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --test decision_floor_coldstart -- cold_start_persistent_bullish_now_closes --ignored --nocapture +``` + +Expected: PASS (stop controller now produces closed trades). + +If it fails, debug — likely a kernel-arg or wiring issue from Tasks 2-9. + +- [ ] **Step 3: Delete `StopRules` struct + uses** + +```bash +rm crates/ml-backtesting/src/policy/stop_rules.rs +``` + +In `crates/ml-backtesting/src/policy/mod.rs`, remove these lines: + +```rust +pub use stop_rules::StopRules; +``` + +and the entire `mod stop_rules;` declaration. + +Remove `sl_tp_rules: StopRules` field from `StrategyConfig` and every literal that uses it: + +```bash +grep -rln "StopRules\|sl_tp_rules" crates/ml-backtesting/src crates/ml-backtesting/tests bin/fxt-backtest/src +``` + +For each hit, delete the field reference and the import. + +Specifically the helper builders in `src/policy/mod.rs` `default_for` and similar use `sl_tp_rules: StopRules::default()` — delete that field from the struct-literal. + +- [ ] **Step 4: Delete `use_cold_start_stopgap` propagation** + +```bash +grep -rln "use_cold_start_stopgap" crates/ml-backtesting/src bin/fxt-backtest/src config/ml crates/ml-backtesting/tests +``` + +For each hit, delete the field, the YAML key, and the runtime branches that read it. + +In `crates/ml-backtesting/src/harness.rs` (around lines 188-216), delete the entire `if any_stopgap { ... }` block (introduced in Q1) and replace with the original simple loop: + +```rust + for (b, strat) in cfg.strategies.iter().enumerate() { + let prog = strat.flatten(); + sim.upload_program(b, &prog)?; + } +``` + +In `crates/ml-backtesting/src/sim/batched_config.rs`, delete `use_cold_start_stopgap: bool` from `UniformSimParams` and `ResolvedSimVariant`, and `use_cold_start_stopgap: Vec` from `BatchedSimConfig`, plus all `_per_b` propagation in `from_uniform`, `from_grid`, `validate`. + +In `bin/fxt-backtest/src/main.rs`, delete `use_cold_start_stopgap: Option` from `SweepBase.SimVariant` and the `.unwrap_or(false)` propagation. + +In `config/ml/sweep_smoke.yaml`, rename variant and drop the flag: + +```yaml + - { name: t0c1l200_isv_stops, threshold: 0.0, cost_per_lot_per_side: 0.125, + latency_ns: 200000000 } +``` + +In each test file, delete `use_cold_start_stopgap: false,` from the `UniformSimParams { ... }` literal. + +- [ ] **Step 5: Verify compilation across the whole workspace** + +```bash +SQLX_OFFLINE=true cargo check -p ml-backtesting -p fxt-backtest +``` + +Expected: clean compile. If anything errors, fix the missed-reference and continue. + +- [ ] **Step 6: Run the entire ml-backtesting test suite to confirm no regression** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --lib +SQLX_OFFLINE=true cargo test -p ml-backtesting --test stop_controller -- --ignored --nocapture +SQLX_OFFLINE=true cargo test -p ml-backtesting --test decision_floor_coldstart -- --ignored --nocapture +SQLX_OFFLINE=true cargo test -p ml-backtesting --test parallel_sim_correctness -- --ignored --nocapture +SQLX_OFFLINE=true cargo test -p ml-backtesting --test threshold_and_cost -- --ignored --nocapture +SQLX_OFFLINE=true cargo test -p ml-backtesting --test lob_sim_fixtures -- --ignored --nocapture +SQLX_OFFLINE=true cargo test -p ml-backtesting --test lob_sim_integrated_fuzz -- --ignored --nocapture +``` + +Expected: all pass. Any failures are bugs introduced by the cleanup and must be fixed before commit. + +- [ ] **Step 7: Add SUPERSEDED headers to the falsified CBSW docs** + +In `docs/superpowers/specs/2026-05-19-cbsw-cold-start-aggregator-design.md`, prepend at line 1: + +```markdown +> **STATUS: SUPERSEDED** by [2026-05-19-isv-driven-stop-controller-design.md](2026-05-19-isv-driven-stop-controller-design.md). +> Empirically falsified — the convictions-log diagnosis was orthogonal to the +> actual n_trades=0 cause (StopRules::default() + additive seed_inflight bug). +> Kept for audit trail. +``` + +In `docs/superpowers/plans/2026-05-19-cbsw-cold-start-aggregator.md`, prepend the same header pointing to `../specs/2026-05-19-isv-driven-stop-controller-design.md` and `2026-05-19-isv-driven-stop-controller.md`. + +- [ ] **Step 8: Commit the atomic cleanup** + +```bash +git add -A +git commit -m "$(cat <<'EOF' +refactor(ml-backtesting): delete StopRules + use_cold_start_stopgap atomically + +- StopRules struct + sl_tp_rules field + all literals deleted; the ISV + stop controller (Tasks 2-9) replaces this dead data path. +- use_cold_start_stopgap propagation deleted across harness, + batched_config, fxt-backtest, sweep_smoke.yaml, and 5 test files. +- Q1 stopgap branch in harness.rs deleted; replaced with the original + simple strategy-upload loop. +- decision_floor_coldstart test retargeted: cold_start_persistent_ + bullish_with_default_stops_never_closes → ..._now_closes, asserts + trades > 0 AND |pos| <= max_lots. +- 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) +EOF +)" +``` + +--- + +## Task 11: Push + cluster smoke validation + +**Files:** none (deployment + monitoring) + +- [ ] **Step 1: Run the full local test suite one final time** + +```bash +SQLX_OFFLINE=true cargo test -p ml-backtesting --lib && \ + SQLX_OFFLINE=true cargo test -p ml-backtesting --test stop_controller -- --ignored --nocapture && \ + SQLX_OFFLINE=true cargo test -p ml-backtesting --test decision_floor_coldstart -- --ignored --nocapture +``` + +Expected: all pass. If anything fails, fix before pushing. + +- [ ] **Step 2: Push the branch** + +```bash +git push origin ml-alpha-phase-a +``` + +Expected: push succeeds. + +- [ ] **Step 3: Deploy the cluster smoke** + +```bash +./scripts/argo-lob-sweep.sh --grid config/ml/sweep_smoke.yaml --tag isv-stops-smoke +``` + +Expected: workflow submitted. Capture the workflow name (e.g., `lob-backtest-sweep-XXXXX`). + +- [ ] **Step 4: Monitor smoke to completion (~16 minutes)** + +```bash +# Monitor pod state until run-sweep + aggregate + notify-result complete. +mc ls fh/argo-logs// # replace with actual name +``` + +Wait for `notify-result` pod to reach Completed. Then fetch logs: + +```bash +mc cat fh/argo-logs//-run-sweep-*/main.log | tail -10 +mc cat fh/argo-logs//-aggregate-*/main.log +``` + +- [ ] **Step 5: Fetch summary.json from training-data PVC** + +Apply the reader pod template (adjust workflow name in the path): + +```bash +cat > /tmp/q-summary-read.yaml < + find "\$ROOT/smoke" -name summary.json -print -exec sh -c 'echo "--- \$1 ---"; cat "\$1"' _ {} \; + volumeMounts: + - { name: training-data, mountPath: /mnt/training-data } + volumes: + - { name: training-data, persistentVolumeClaim: { claimName: training-data-pvc } } +EOF + +kubectl -n foxhunt apply -f /tmp/q-summary-read.yaml +# wait for pod to complete, then: +kubectl -n foxhunt logs isv-stops-summary-read +kubectl -n foxhunt delete pod isv-stops-summary-read +``` + +- [ ] **Step 6: Validate pass criteria** + +The summary.json must satisfy ALL three: + +| Metric | Threshold | Source field | +|---|---|---| +| `n_trades` | `>= 100` | `n_trades` | +| `mean_trade_hold_time_ns ∈ (0, 30·10⁹)` | yes | derived from trade records — fetch with the trade-records read path if not in summary.json | +| `max(|pos.position_lots|)` at smoke end | `<= 5` | requires reading final positions from the PVC; or inferred from absence of overflow flags in the smoke log | + +If `n_trades >= 100` AND `mean_trade_hold_time_ns ∈ (0, 30·10⁹)` AND all positions respect the cap: **SUCCESS**, the spec is validated. + +If any criterion fails: investigate the smoke log for clues, then update the spec and roll a new task plan. Do NOT silently widen the gates. + +- [ ] **Step 7: Commit final validation result memo (if successful)** + +```bash +# Add a project memory entry summarizing the validation result. +# Use the auto-memory system: write to /home/jgrusewski/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/project_isv_stops_smoke_validation.md +# Then update MEMORY.md index entry. +# This memory captures the smoke outcome for future-conversation reference. +``` + +(Memory write is per the auto-memory rules; the engineer fills in the actual smoke numbers.) + +--- + +## Notes for the Implementer + +- **TDD discipline**: every task starts with a failing test (functional, API, or setup failure mode), then implementation, then verifying pass. Do not skip the "verify it fails" step — it is the proof the test is invariant-asserting rather than observation-locked. +- **Each task is one commit**. No mid-task commits. If a task hits a blocker, stop and ask rather than half-committing. +- **Memory-pearl conformance**: Tasks 2, 4, 5 use the `max(real, floor)` pattern from `pearl_blend_formulas_must_have_permanent_floor.md`. Task 2 uses first-observation bootstrap per `pearl_first_observation_bootstrap.md`. Task 5's trail uses `pos.position_lots` pre-action mag per `pearl_trail_fire_pre_vs_action_mag.md`. Task 7 enforces single-source-of-truth per `feedback_single_source_of_truth_no_duplicates.md` via the shared `__device__` helper. Task 10 atomically migrates every consumer per `feedback_no_partial_refactor.md`. +- **GPU contract**: `stop_check_isv` is pure-device with no host branches. The new args to `book_update_apply_snapshot` and `decision_policy_*` are pointers, not scalars-changing-mid-kernel, so CUDA Graph capture (P5) remains valid. +- **If Task 11 cluster smoke fails**: do not push fixes blindly. Run a local 200-event reproduction first (the `cold_start_persistent_bullish_now_closes` test scaffold is the closest local proxy). Iterate locally; only re-deploy when the local repro produces the expected smoke metrics.