spec(ml-backtesting): apply 9 critical-review fixes to parallelism spec

Critical review surfaced 9 issues; all fixed inline:

1. §1 — Reframed "140× amortization" as "amortization on forward; sim
   runs parallel". True win is on the ~2ms forward shared across 140
   cells, not on sim work which scales linearly with n_backtests.
2. §2 — Made the 1h target vs firm bound explicit (≤1h target, ≤2h
   firm). Acknowledges Graph capture realistic speedup is 1.2-1.5×,
   not 2×.
3. §5 — Dropped atomicAdd. Plain `+=` under the single-writer-per-block
   convention (existing pattern across sim kernels). No race.
4. §4 — Documented max-over-horizons threshold rationale (vs per-
   horizon or aggregate-conviction). Flagged per-horizon as a follow-up
   tweak if dilution pathway matters in the verdict.
5. §7 P5 + §10 risk — Captured-vs-uncaptured tolerance is 1e-5 relative,
   NOT strict bit-identity. CUDA Graph capture can reorder reductions
   harmlessly by 1 ULP; strict bit-identity would be a false-positive.
6. §8 — Made explicit that parallel_sim_equivalence + independence
   tests are BOTH required. Equivalence alone is necessary but not
   sufficient (a shadow-backtest[0] bug still passes equivalence).
7. §3.3 — Specified output schema: `cell_W{n}/sim_<variant_name>/`
   with summary.json carrying a resolved `sim_config` block (verdict
   emitter reads that, not the directory name).
8. §7 P4 — Enumerated apply_fill_to_pos call sites + added grep-verify
   step before commit, so no fill path silently loses cost.
9. §9 — Tightened rate-validation gates with hard targets (P2 ≤90s,
   P5 ≤60s) instead of generous minute envelopes. Added §9.1 stride=8
   fallback as explicit Plan-B if Graph capture under-delivers.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>

docs/superpowers/specs/2026-05-19-deployability-sweep-parallelism-design.md

@@ -15,7 +15,7 @@ We have 1 L40S GPU pod, scaling to 3 if a 1-pod proof-of-life passes. The budget
 
 Three architectural facts make the goal achievable:
 
-1. **Forward sharing across cells.** Of the 560-cell grid, only the `data window` (4 quarters) varies the forward pass. The other three axes — `cost`, `latency`, `threshold` — are sim parameters downstream of the forward. All `7 × 4 × 5 = 140` cells per window can share ONE forward pass and differ only in sim state. This is a ~140× amortization lever.
+1. **Forward sharing across cells, sim runs parallel.** Of the 560-cell grid, only the `data window` (4 quarters) varies the forward pass. The other three axes — `cost`, `latency`, `threshold` — are sim parameters downstream of the forward. All `7 × 4 × 5 = 140` cells per window can share ONE forward pass; the *forward* work amortizes across 140 cells. Sim work itself does NOT amortize — it scales linearly with `n_backtests`, but runs in parallel across thread blocks (L40S has ~142 SMs so n=140 fits in one wave with bandwidth contention as the only cost). The win comes from sharing the ~2 ms/forward across 140 cells; not from somehow making sim work disappear.
 2. **Existing sim kernels already loop over `n_backtests`.** The decision/pnl/order kernels were written for `n_backtests > 1`. What's missing is per-backtest parameter arrays — current kernels take scalar params and broadcast.
 3. **Three new pieces of functionality** (not just refactor): a threshold gate (currently no kernel-level threshold), per-fill cost deduction (currently `fees_usd_fp` is a literal-zero placeholder in `pnl_track.cu:92`), and per-backtest variants of every existing sim parameter.
 
@@ -27,12 +27,12 @@ Spec covers the refactor + new functionality + rate-side improvements to land a
 
 **Falsifiable goal.** After this spec lands:
 
-- A single 1-pod `lob-backtest-sweep` invocation runs **1 quarter × 140 sim variants** end-to-end in ≤ 30 min wall-clock.
-- 4-quarter sweep across 3 pods completes in ≤ 1 h wall-clock.
+- A single 1-pod `lob-backtest-sweep` invocation runs **1 quarter × 140 sim variants** end-to-end in ≤ 30 min wall-clock at decision_stride=4.
+- 4-quarter sweep across 3 pods completes in **≤ 1 h wall-clock target** (firm bound: ≤ 2 h). The 1-h target depends on Graph capture delivering ≥ 1.5× on the forward; if it under-delivers, the explicit Plan-B is `decision_stride 4 → 8` (gate criterion in §9, automatically activated based on P5 measurement).
 - All 140 per-backtest summary.json files differentiate by their cost/latency/threshold configs (no copy-paste failures from indexing bugs).
 - The existing 100k-event smoke (n_parallel=1, uniform config) reproduces bit-identically to the pre-refactor result (proves the refactor doesn't change single-backtest behaviour).
 
-**Out of scope.** bf16 mixed precision (deferred to a follow-up spec, only if 1-h budget isn't reached). seq_len reduction (would require retraining). n_batch>1 in forward across windows (would require multi-window batched data loader; revisit if budget is missed).
+**Out of scope.** bf16 mixed precision (deferred — see §10 risk register; would be a separate follow-up spec only if both Graph capture AND stride=8 fallback miss the target). seq_len reduction (would require retraining). n_batch>1 in forward across windows (would require multi-window batched data loader; revisit if budget is still missed after both levers).
 
 ---
 
@@ -145,6 +145,24 @@ cells:
 
 The sweep runner expands `sim_variants` into a `BatchedSimConfig` for each cell. The aggregate step reads all 4 cells' per-backtest summary.json files (140 each = 560 total) and writes the standard aggregate.parquet + pareto_frontier.json.
 
+**Output schema for n=140.** Each cell writes one directory per sim-variant, named after the variant's `name` field (which has to be unique within `sim_variants`):
+
+```
+<sweep-root>/<sweep-tag>/
+  W1/
+    sim_c0_l0_t0/{summary.json, trades.csv, pnl_curve.bin}
+    sim_c0_l0_t1/{summary.json, trades.csv, pnl_curve.bin}
+    ...
+    sim_c6_l3_t4/{...}
+  W2/
+    sim_c0_l0_t0/{...}
+    ...
+  W3/
+  W4/
+```
+
+Aggregator + verdict emitter parse `sim_<variant_name>/` to extract (cost, latency, threshold) from the variant name (or, more robustly, write the resolved sim params into each `summary.json` as a `sim_config` block — verdict emitter prefers reading that to parsing the name). The verdict emitter then groups by `(threshold_value, anchor)` across windows for the median Sharpe / max-dd math from the original v2 spec §3.5.
+
 **3-pod scaling gate.** Run a single-cell `lob-backtest-sweep` (1 window, 140 sim variants) on 1 L40S. Verify:
 
 - All 140 backtests complete cleanly
@@ -176,6 +194,10 @@ if (max_conviction < threshold_per_b[b]) {
 
 Pre-Kelly placement keeps the gate deterministic from alpha alone, so the threshold pre-registration step (compute p60–p95 absolute values from observed convictions on a validation window) reflects exactly what gets gated in deployment.
 
+**Why max-over-horizons** (vs. per-horizon thresholding or aggregate-conviction gate after Kelly): the threshold's role is "is the model confident enough on ANY horizon to consider trading at all". A signal strong on h=4 (long-horizon) but neutral on h=0 (short-horizon) should still be allowed to trade — the downstream Kelly aggregation already weights horizons by `recent_sharpe`, which gives the model agency over which horizon's view actually drives sizing. Per-horizon thresholding would over-prune (gating out a horizon's contribution before Kelly even sees it) and would require a 5×n_backtests threshold array. Aggregate-conviction gating (after Kelly) couples threshold semantics to Kelly state, defeating the pre-registration step's purpose (the threshold value calibrated at p60 of CONVICTION wouldn't map to p60 of post-Kelly sizing).
+
+Max-over-horizons is the simplest interpretation of "p60–p95 of conviction" that keeps the gate decoupled from sim-side state. Future tweak: if the dilution pathway turns out to matter (verdict shows threshold is selecting trades dominated by a single horizon's strong signal that gets washed out in the aggregate), we revisit per-horizon thresholding as a follow-up.
+
 **Pre-registration step** (in `sweep_threshold_tuning.yaml`):
 
 - Run 1 cell with `threshold = 0.0` (never gates) on the validation window
@@ -204,13 +226,16 @@ Side-channel logging is ~20 LOC in the harness.
 //
 // add (BEFORE the existing realized_pnl math, to keep gross-vs-net clearly separated):
 const float fill_cost = filled_lots * cost_per_lot_per_side_per_b[b];
-pos.realized_pnl -= fill_cost;     // deducted from realized P&L
+pos.realized_pnl -= fill_cost;        // deducted from realized P&L
 // Plus: accumulate fill_cost into a per-backtest device array
 // `total_fees_per_b_d` (NEW). pnl_track_step's trade-log write then
 // reads the running total and writes `fees_usd_fp` per trade record
 // (currently a placeholder zero at pnl_track.cu:92).
-atomicAdd(&total_fees_per_b_d[b], fill_cost);   // single-writer per block; can be plain += under the
-                                                 // existing `if (threadIdx.x != 0) return;` convention
+//
+// Plain `+=` — single-writer-per-block convention is already enforced
+// across the sim kernels (`if (threadIdx.x != 0) return;` at the top
+// of every block). No race; no atomic needed.
+total_fees_per_b_d[b] += fill_cost;
 ```
 
 (`Pos` struct currently has no `total_fees` field — verified by reading `lob_state.cuh:22-29`. We add a separate per-backtest device array `total_fees_per_b_d` rather than growing `Pos`, to keep the existing Pos<->PosFlat Rust mirror layout-stable and to keep fee accumulation orthogonal to position tracking. `artifacts.rs::write_summary_json` reads `total_fees_per_b_d[b]` to populate `summary.json::total_fees_usd`.)
@@ -285,8 +310,8 @@ Per `feedback_no_partial_refactor`: every contract change with all consumers in
 | P1 | `feat(ml-backtesting): per-backtest sim parameter arrays (kernel ABI + Rust BatchedSimConfig)` | decision_policy.cu, pnl_track.cu, order_match.cu, resting_orders.cu, sim.rs, harness.rs, fixture/fuzz tests, main.rs CLI, sweep YAML schema | All existing decision_floor_coldstart + fuzz + fixture tests pass with `BatchedSimConfig::from_uniform`. NEW: `parallel_sim_equivalence_with_uniform_config` (n=8 with same config → 8 bit-identical results). NEW: `parallel_sim_independence_per_backtest` (different latencies → different fill timing). |
 | P2 | `feat(ml-backtesting): seed_inflight_limits_batched kernel (replaces host roundtrip loop)` | resting_orders.cu (or new cu), sim.rs (drop dispatch_latent_market_orders) | New `seed_inflight_limits_batched_smoke` test: 4 backtests, 2 non-noop with different latencies → 2 in-flight slots populated, 2 untouched. |
 | P3 | `feat(ml-backtesting): detect_close_transitions_batched kernel (replaces close-detection host loop)` | pnl_track.cu (or new cu), sim.rs (drop close-detection host loop) | New `detect_close_smoke` test: 4 backtests, 1 close, 3 unchanged → closed_mask[1]==prev_mask[1], rest = 0. |
-| P4 | `feat(ml-backtesting): threshold gate + cost integration` | decision_policy.cu (threshold), resting_orders.cu (cost in apply_fill_to_pos), pnl_track.cu (drop placeholder), harness side-channel for threshold tuning | `threshold_gate_*` tests, `cost_deducted_at_each_fill`, `kelly_state_sees_net_return`. |
-| P5 | `feat(ml-alpha): CUDA Graph capture of forward_only` | cfc/trunk.rs (capture_graph_a), trainer/perception.rs (forward_only uses launch_captured), data/loader.rs (drop misleading comment) | `forward_captured_matches_uncaptured` bit-equivalence test (already-trained checkpoint). |
+| P4 | `feat(ml-backtesting): threshold gate + cost integration` | **decision_policy.cu** (add threshold gate pre-Kelly in BOTH `decision_policy_default` and `decision_policy_program`). **resting_orders.cu** (`apply_fill_to_pos` gains `cost_per_lot_per_side_per_b` arg; both call sites updated — line ~185 in the StopMarket fill path and line ~214 in the in-flight arrival fill path). **order_match.cu / submit_market_immediate** (if it has its own fill path, also update — verify during P4 implementation). **pnl_track.cu** (drop `fees_usd_fp` literal-zero placeholder; read from `total_fees_per_b_d` at close). **harness.rs** (side-channel `max_conviction` logger for threshold tuning pre-registration step). | `threshold_gate_*` tests, `cost_deducted_at_each_fill`, `kelly_state_sees_net_return`. ALSO: search-and-verify before commit `grep -rn 'apply_fill_to_pos' crates/ml-backtesting/cuda/` — every call site must pass the new arg, otherwise some fill paths lose cost silently. |
+| P5 | `feat(ml-alpha): CUDA Graph capture of forward_only` | cfc/trunk.rs (new `capture_graph_a` + bracketing per `pearl_cudarc_disable_event_tracking_for_graph_capture`), trainer/perception.rs (forward_only uses launch_captured), data/loader.rs (drop misleading comment that references the not-yet-existing `capture_graph_a`). | `forward_captured_matches_uncaptured` — captured vs uncaptured forward must agree **within 1e-5 relative tolerance** (NOT strict bit-identity; CUDA Graph capture can reorder kernel launches in ways that flip f32 reduction sums by 1 ULP without changing semantics). Strict bit-identity would be a false-positive trigger from harmless ordering effects. Tolerance picked at 1e-5 because the heads' sigmoid output is the only test target, and 1e-5 is below probability-rounding precision (3 digits effective). |
 | P6 | `config(ml-alpha): sweep YAML schema for batched cells + sweep runner` | bin/fxt-backtest/src/main.rs (SweepCellResolved → BatchedSimConfig::from_grid), scripts/argo-lob-sweep.sh (cell = window, not sim-variant), config/ml/sweep_threshold_tuning.yaml + sweep_deployability.yaml (new schema) | 1-pod end-to-end smoke at n_parallel=140 against the trained checkpoint, validates rate target. |
 | P7 | `docs(ml-alpha): scale to 3-pod sweep` | argo-lob-sweep.sh (no code change; verify 3-pod fan-out behaves) | Run full 4-quarter sweep on 3 pods, verify ≤ 1h wall-clock and 560 differentiable summary.json files. |
 
@@ -300,10 +325,12 @@ All CUDA tests stay `#[ignore = "requires CUDA"]` and run locally via `cargo tes
 
 **New tests landing alongside their commit:**
 
+> **Important — equivalence + independence pair.** The two P1 tests below MUST both ship in the same commit. The equivalence test alone is necessary but NOT sufficient: a per-backtest indexing bug where `backtest[0]`'s result is copied to all slots (e.g., `b` index dropped, or stride math broken) still produces 8 identical results, so the equivalence test passes despite the bug. The independence test catches that pathway: with different latencies, two backtests MUST produce different fill timing, so any "shadow backtest[0] everywhere" bug shows up as identical timings where they should differ. Neither test substitutes for the other.
+
 | Test | Purpose | Commit |
 |---|---|---|
-| `parallel_sim_equivalence_with_uniform_config` | n=8, all-equal configs → 8 bit-identical results. Catches per-backtest indexing bugs. | P1 |
-| `parallel_sim_independence_per_backtest` | n=2, different latencies → different fill timing. Proves per-backtest path is real. | P1 |
+| `parallel_sim_equivalence_with_uniform_config` | n=8, all-equal configs → 8 bit-identical results. Necessary but NOT sufficient — also see independence test below. | P1 |
+| `parallel_sim_independence_per_backtest` | n=2, different latencies → different fill timing. Catches the "shadow backtest[0] everywhere" bug that equivalence alone misses. Required alongside the equivalence test. | P1 |
 | `seed_inflight_limits_batched_smoke` | 4 backtests, 2 non-noop → exactly those slots populated. | P2 |
 | `detect_close_smoke` | 4 backtests, 1 close → closed_mask written correctly. | P3 |
 | `threshold_gate_skips_low_conviction` | max_conviction < threshold → noop. | P4 |
@@ -324,16 +351,26 @@ All CUDA tests stay `#[ignore = "requires CUDA"]` and run locally via `cargo tes
 
 ## 9. Rate target validation
 
-Measure at each gate:
+Measurements are TIGHT (lower bound a known-passing config would hit, not a generous envelope). Misses trigger an investigation per `feedback_stop_on_anomaly` AND may activate the stride=8 fallback (see below).
 
 | Gate | Measurement | Threshold |
 |---|---|---|
-| P2 done | 100k-event smoke with n_parallel=140, all-uniform config, no graph capture | ≤ 5 min wall (proves roundtrip-loop GPU-ization paid off) |
-| P5 done | Same smoke + graph capture | ≤ 2.5 min wall (proves 2× forward speedup) |
-| P6 done | 1-quarter smoke with n_parallel=140, full sim-variant grid | ≤ 30 min wall (1-pod proof-of-life) |
-| P7 done | 4-quarter sweep on 3 pods | ≤ 60 min wall, ≤ 5 GPU-hours total |
+| P2 done | 100k-event smoke with n_parallel=140, all-uniform config, no graph capture, stride=4 | **≤ 90 s wall**. Reference: 100k events at n=1 takes ~30 s today (measured). At n=140 with fixed roundtrip-loop kernels, throughput should be approximately unchanged because sim runs parallel and forward is amortized. >90s = roundtrip-loop fix didn't pay off (investigate before proceeding to P3). |
+| P5 done | Same smoke + graph capture | **≤ 60 s wall** (≥ 1.3× speedup vs P2). If captured-vs-uncaptured matches but speedup < 1.3×, the captured graph isn't engaging properly (likely a host-branch or stale-pointer issue trapping launch into eager mode). |
+| P6 done | 1-quarter smoke with n_parallel=140, full sim-variant grid, stride=4 | **≤ 30 min wall** (1-pod proof-of-life). If miss, activate stride=8 fallback (see below) and re-run; new threshold ≤ 16 min. |
+| P7 done | 4-quarter sweep on 3 pods, stride=4 (or 8 if fallback activated) | **≤ 60 min wall target** (firm bound ≤ 120 min). |
 
-If any gate misses, stop and diagnose before proceeding. Per `feedback_stop_on_anomaly`.
+### 9.1 Stride=8 fallback (the Plan-B lever)
+
+Activation criterion: **if P6 measurement at stride=4 misses the 30-min gate by > 20% (i.e., > 36 min)**, the spec's primary plan has insufficient headroom for the 1-h sweep target. In that case, the implementation plan automatically:
+
+1. Reruns P6 measurement at `decision_stride = 8` (half the decisions, ~2× wins on the forward-dominated portion).
+2. If P6@stride=8 hits the ≤ 16 min gate, commits `config(ml-alpha): bump deployability sweep decision_stride 4 → 8` and proceeds to P7 with stride=8.
+3. Documents the decision in `config/ml/v2_prod_thresholds.json` notes so the verdict-emitter knows which stride generated the input.
+
+Rationale for stride=8 as the cleanest Plan-B (vs bf16): stride=8 is a 1-line YAML change, no kernel work, no precision-validation gate. It buys ~2× wall reduction at the cost of slightly coarser decision cadence — model was trained per-snapshot, but inference cadence is independent of the model's intrinsic timescale (h6000 ≈ 6000 snapshots forward). Down-sampling decisions by 2× has minimal effect on a multi-thousand-snapshot horizon.
+
+If P6@stride=8 ALSO misses ≤ 16 min, the spec is in unrecoverable territory and we open a follow-up spec for bf16. Don't paper over it by accepting > 1h.
 
 ---
 
@@ -341,11 +378,11 @@ If any gate misses, stop and diagnose before proceeding. Per `feedback_stop_on_a
 
 | Risk | Likelihood | Mitigation |
 |---|---|---|
-| Graph capture has a pearl violation (host branch / scalar arg change) we don't catch until P5 ships | Medium | Bit-equivalence test (`forward_captured_matches_uncaptured`) gates P5; if it diverges, the capture is wrong by definition. |
+| Graph capture has a pearl violation (host branch / scalar arg change) we don't catch until P5 ships | Medium | Captured-vs-uncaptured equivalence test (`forward_captured_matches_uncaptured`) gates P5 with **1e-5 relative tolerance** (not strict bit-identity — captured graphs can reorder kernel launches and flip f32 reductions by 1 ULP harmlessly). If divergence exceeds 1e-5, the capture is semantically broken (or carries a host-branch / stale-pointer / scalar-arg-change pearl violation). |
 | Per-backtest indexing bug at n=140 (off-by-one in stride math) | Medium | `parallel_sim_equivalence_with_uniform_config` test catches any per-backtest path that doesn't reduce to broadcast at uniform config. |
 | `seed_inflight_limits_batched` slot allocation has a race | Low | Each backtest's MAX_LIMITS slots are PER-backtest (different memory ranges), no cross-backtest atomics needed. Per-backtest scan is single-threaded inside `if (threadIdx.x != 0) return;` pattern (same convention as `decision_policy_default`). |
 | Cost deduction breaks Kelly behaviour vs trained distribution | Low | The model was trained without cost in the loss (forward inference doesn't see cost). Kelly state ON the backtest side learns from net returns, which adapts to the cost regime. If a specific cost regime trades degenerately at gate P6, that's information about the deployability question — NOT a refactor bug. |
-| 1-h budget still missed after all levers | Medium | bf16 follow-up spec is the next escape hatch (~1.5× additional). Beyond that: reduce decision_stride from 4 to 8 (2× decisions saved, slight signal-cadence cost). Out of scope for THIS spec; flagged here so we know what's left. |
+| 1-h target missed at P5 (Graph capture under-delivers) | **High** — Graph capture's realistic L40S speedup is 1.2-1.5× on a GEMM-dominated forward, not 2×. Plan numbers assumed 2× headroom. | **Explicit Plan-B in §9.1**: stride=4 → stride=8 if P6 misses by > 20%. 1-line YAML change. No kernel work needed. If stride=8 ALSO misses: open follow-up bf16 spec. Don't accept > 1h silently. |
 | Misleading X11 comment in loader.rs:274 references `capture_graph_a` that doesn't exist | Trivial | Drop the comment in commit P5 when capture is implemented. |
 
 ---
@@ -361,8 +398,10 @@ If any gate misses, stop and diagnose before proceeding. Per `feedback_stop_on_a
 
 ## 12. Definition of done
 
-- Commits P1–P7 landed with all gating tests green.
-- Full 4-quarter deployability sweep runs at ≤ 1 h wall-clock on 3 pods, produces `deployability_verdict.json` per the original v2 spec §3.5.
-- 1-pod proof-of-life passes before scaling to 3.
-- Memory pearls hold: no host branches in captured graph, max-with-floor pattern preserved for both cold-start floors and the new cost-net Kelly path, no version suffixes in code identifiers, no legacy aliases.
+- Commits P1–P7 landed with all gating tests green (per §8 and §9 measurements).
+- Full 4-quarter deployability sweep runs at **≤ 1 h wall-clock target** (firm bound ≤ 2 h) on 3 pods, produces `deployability_verdict.json` per the original v2 spec §3.5.
+- 1-pod proof-of-life (§9 P6) passes before scaling to 3 pods.
+- If stride=8 fallback (§9.1) was activated, the decision is documented in `config/ml/v2_prod_thresholds.json` notes.
+- Memory pearls hold: no host branches in captured graph, max-with-floor pattern preserved for both cold-start floors and the new cost-net Kelly path, no version suffixes in code identifiers, no legacy aliases, plain += in single-writer-per-block contexts (not atomicAdd).
 - The vestigial loader.rs:274 `capture_graph_a` comment is fixed in P5.
+- `apply_fill_to_pos` call-site verification (`grep -rn 'apply_fill_to_pos' crates/ml-backtesting/cuda/`) confirms every call passes the new `cost_per_lot_per_side_per_b` arg — no silent fee loss on any fill path.