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

Deployability-Sweep Parallelism — Design

Author: brainstormed live 2026-05-19 with claude-opus-4-7 Status: Draft (awaiting review) Supersedes: none — this is a follow-up to docs/superpowers/specs/2026-05-19-ml-alpha-v2-trunk-grows-and-deployability-design.md (commit da1dd92bf) Depends on: c7fdc617d and earlier (cold-start floor + variance-cap gate landed; smoke validated end-to-end with kernel fixes)

---

1. Problem

The deployability sweep grid (per the v2 design spec §3.3, Pass 2) is 7 cost × 4 latency × 5 threshold × 4 window = 560 cells. At the smoke's measured rate of 1988 events/s ≈ 2 ms per forward on L40S with n_parallel = 1, one cell takes ~83 min per quarter. Naive Argo fan-out at 1 cell per pod = ~750 GPU-hours per sweep.

We have 1 L40S GPU pod, scaling to 3 if a 1-pod proof-of-life passes. The budget is ≤ 1 hour wall-clock per sweep because the development loop iterates on multiple sweeps per day (hyperparam search, A/B candidates). At today's per-cell rate this is impossible.

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.
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.

Spec covers the refactor + new functionality + rate-side improvements to land a sweep under the 1-hour budget.

---

2. Goal and verdict criteria

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.
• 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).

---

3. Architecture — three multiplicative levers

3.1 Per-backtest matrix (the dominant lever, ~140×)

Every sim parameter becomes a per-backtest f32/u32 array. Single BatchedSimConfig is the new API:

pub struct BatchedSimConfig {
    pub target_annual_vol_units: Vec<f32>,  // len = n_backtests
    pub annualisation_factor:    Vec<f32>,
    pub max_lots:                Vec<u16>,
    pub latency_ns:              Vec<u32>,
    pub kelly_frac_floor:        Vec<f32>,
    pub sharpe_weight_floor:     Vec<f32>,
    pub threshold:               Vec<f32>,   // NEW
    pub cost_per_lot_per_side:   Vec<f32>,   // NEW
}

impl BatchedSimConfig {
    /// Used by smoke / fixture tests where every cell shares the same config.
    pub fn from_uniform(n: usize, cfg: &UniformSimParams) -> Self;
    /// Used by the sweep runner to pack a 140-variant grid into one pod.
    pub fn from_grid(cells: &[SweepCellResolved]) -> Self;
}

The previous scalar BacktestHarnessConfig is the legacy API; dropped per feedback_no_legacy_aliases and feedback_single_source_of_truth_no_duplicates. BacktestHarness constructs a BatchedSimConfig from its inputs (uniform broadcast at n=1 for smoke, grid pack at n=140 for sweep).

Kernel ABI changes (all in one atomic commit per feedback_no_partial_refactor):

// decision_policy_default + decision_policy_program:
//   was: float target_annual_vol_units, float annualisation_factor, int max_lots,
//        float kelly_frac_floor, float sharpe_weight_floor
//   new: const float* target_annual_vol_units_per_b,
//        const float* annualisation_factor_per_b,
//        const int*   max_lots_per_b,
//        const float* kelly_frac_floor_per_b,
//        const float* sharpe_weight_floor_per_b,
//        const float* threshold_per_b           // NEW

// pnl_track_step + apply_fill_to_pos:
//   const float* cost_per_lot_per_side_per_b   // NEW

// step_resting_orders + the latency-path dispatch:
//   const u32* latency_ns_per_b

Host roundtrip removal — TWO new kernels (verified during scoping; see §6 for the verified host loops):

Kernel	Replaces	LOC	Notes
seed_inflight_limits_batched	dispatch_latent_market_orders host loop (sim.rs:706) — reads market_targets to host, loops over backtests, calls seed_limit_order per non-noop	~80	One kernel launch per decision; per-backtest scans MAX_LIMITS slots for a free one. Slot allocation is per-backtest so no cross-backtest atomics needed. Overflow path writes to a per-backtest counter on Pos.
detect_close_transitions_batched	host loop at sim.rs:658-674 — compares prev_pos vs current pos per backtest, calls read_pos per close-eligible backtest, builds closed_mask + realised_return vectors host-side	~70	Writes closed_horizon_mask_d and realised_return_d directly on GPU. isv_kelly_update_on_close then dispatched without host round-trip.

Total: ~150 LOC of CUDA, plus existing kernel signature updates (~50 LOC) and Rust side (~200 LOC).

3.2 CUDA Graph capture for the forward pass (~2-3× per-forward speedup)

Verified: forward_only → evaluate_batched uses launch_builder, NOT launch_captured. The X11 plan said "capture_graph_a covers full v2 forward" but the actual X11 commit (4f888abbf) shipped only forward_only + from_checkpoint. Graph capture is new work, not plumbing.

Per pearl_no_host_branches_in_captured_graph + pearl_cudarc_disable_event_tracking_for_graph_capture:

• Disable cudarc 0.19 event tracking around the capture (bracket with disable_event_tracking() + enable_event_tracking() per the pearl)
• No host branches inside the capture region
• No host-malloc / scalar-arg changes / pointer swaps inside capture
• No per-call kernel parameter changes (use pre-bound device pointers; bind once at trunk construction)

Capture scope. Wrap the entire forward chain into one captured graph that takes (window_d, h_init) device pointers and writes probs_d. Components:

VSN → Mamba2 stack 1 → LN_a → Mamba2 stack 2 → LN_b → attn_pool → CfC step → heads (×N_HORIZONS)

Snapshot-assemble (snap_feature_assemble_batched) stays OUTSIDE the captured graph because it consumes a fresh Mbp10RawInput per call (its inputs change per decision; the captured forward consumes the assembled tensor instead).

Capture is built ONCE in PerceptionTrainer::new (or from_checkpoint), replayed per decision via launch_captured. Verify with a forward_captured_matches_uncaptured bit-equivalence test that doesn't ship to production but blocks merge.

Realistic estimate: ~150 LOC of plumbing + 1 careful test. Expected impact: ~2× on the forward pass on L40S (the kernel-launch latency overhead is a real chunk of the 2 ms/forward; on Ampere/Ada a captured graph drops it by 60-70% per official benchmarks).

3.3 Pod-level parallelism (1 → 3 pods, ~3× wall reduction)

argo-lob-sweep.sh already fans cells out as separate WorkflowTasks. Change WHAT a cell is:

• Old: cell = (cost, latency, threshold), 140 cells per window × 4 windows = 560 cells
• New: cell = (window), 4 cells per sweep. Each cell's BatchedSimConfig packs the 140 sim variants for that window.

Sweep YAML schema:

base:
  data_template: /mnt/training-data/futures-baseline-mbp10/ES.FUT/ES.FUT_{window}.dbn.zst
  predecoded_dir: /feature-cache/predecoded
  decision_stride: 4
  n_parallel: 140                    # = len(sim_variants)
  checkpoint: /feature-cache/alpha-perception-runs/<sha>/trunk_best_h6000.bin
  sim_variants:                      # 140 entries
    - { name: c0_l0_t0, cost: 0.0,    latency_ns: 100000000, threshold: 0.30 }
    - { name: c0_l0_t1, cost: 0.0,    latency_ns: 100000000, threshold: 0.35 }
    # ... 140 total
cells:
  - { name: W1, window: 2025-Q2 }
  - { name: W2, window: 2025-Q3 }
  - { name: W3, window: 2025-Q4 }
  - { name: W4, window: 2026-Q1 }

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.

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
• Per-backtest summary.json files have distinct metrics (no indexing copy-paste)
• Wall-clock ≤ 30 min

Only then scale to 3-pod parallel. Per cluster autoscaler cap (3 GPU nodes), the 4-cell sweep gets 3 cells parallel + 1 queued. Wall = max(2 × per-cell, 1 × per-cell) ≈ 2× per-cell = ≤ 60 min total.

---

4. Threshold gate (kernel-side absolute cutoff)

Insertion at the top of decision_policy_default and decision_policy_program, before the per-horizon ss[] loop:

float max_conviction = 0.0f;
#pragma unroll
for (int h = 0; h < N_HORIZONS; ++h) {
    max_conviction = fmaxf(max_conviction,
                           fabsf(alpha_probs[h] - 0.5f) * 2.0f);
}
if (max_conviction < threshold_per_b[b]) {
    market_targets[b * 2 + 0] = 2;   // noop
    market_targets[b * 2 + 1] = 0;
    open_horizon_masks[b]    = 0u;
    return;
}

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.

Pre-registration step (in sweep_threshold_tuning.yaml):

• Run 1 cell with threshold = 0.0 (never gates) on the validation window
• Harness logs every observed max_conviction per decision into a side-channel Vec<f32> (host-side, computed from the SAME probs that get broadcast — no extra GPU readback)
• After run: compute empirical p60/p70/p80/p90/p95 → write to config/ml/v2_prod_thresholds.json
• Full deployability sweep then uses those 5 absolute values

Side-channel logging is ~20 LOC in the harness.

Regression tests:

• threshold_gate_skips_low_conviction: alpha p_h = 0.51 (max_conviction = 0.02), threshold = 0.10 → market_target = noop
• threshold_gate_allows_high_conviction: alpha p_h = 0.7 (max_conviction = 0.4), threshold = 0.10 → market_target buy, size > 0

---

5. Per-fill cost (per-backtest f32 array, kernel-side deduction)

Verified: pnl_track.cu:92 writes a literal 0 to fees_usd_fp as a placeholder. The cost system is unimplemented.

Insertion point (single location): apply_fill_to_pos in resting_orders.cu:90. This function is called from both the in-flight fill path (step_resting_orders) and the immediate-fill path (submit_market_immediate).

// existing:
//   pos.realized_pnl += sgn * unwind * (vwap_entry - avg_px);  // close-leg P&L
//
// 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
// 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

(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.)

Net-of-cost semantics. Both entry-fill and exit-fill incur cost. A round-trip 1-lot trade with cost_per_lot_per_side = 2.0 produces total_fees = 4.0 and realized_pnl_net = realized_pnl_gross - 4.0. isv_kelly_update_on_close reads realised_return = (pos.realized_pnl_now - pos.realized_pnl_at_open) / abs_size, which is net of cost — Kelly state learns from realistic return distribution. This is the correct architecture per the user's preference for "per-fill in pnl_track kernel" (over "per-close in summary stats" which would mismatch Kelly's view from reported metrics).

The trade-log record's fees_usd_fp field (currently placeholder zero at pnl_track.cu:92) gets the running total at close-time, written by pnl_track_step when it emits the close record.

Regression tests:

• cost_deducted_at_each_fill: round-trip 1-lot trade with cost=2.0 should produce net_pnl = gross_pnl - 4.0
• kelly_state_sees_net_return: assert isv_kelly_update_on_close's realised_return arg = (gross - 2*cost) / abs_size

---

6. Verified host-roundtrip locations (the GPU-ization targets)

This section documents the two host loops verified by code reading during spec scoping. Both must be replaced for the per-backtest matrix to actually pay off — without them, host roundtrip overhead at n_parallel=140 cancels the speedup.

Loop 1 — dispatch_latent_market_orders at crates/ml-backtesting/src/sim.rs:706:

let mut targets = vec![0i32; n * 2];
self.stream.memcpy_dtoh(&self.market_targets_d, targets.as_mut_slice())?;
for b in 0..n {
    let side = targets[b * 2];
    let size = targets[b * 2 + 1];
    if side == 2 || size <= 0 { continue; }
    for slot in 0..MAX_LIMITS {
        if self.seed_limit_order(b, slot, side as u8, 1, 2, 0xFF,
            if side == 0 { 1.0e9 } else { 0.0 }, size as f32, 0.0,
            current_ts_ns + latency_ns as u64).is_ok() { break; }
    }
}

Cost at n=140 × 2.5M decisions × 60% non-noop = ~210M host roundtrips per quarter (assuming worst case: every backtest non-noop on most decisions). Real cost is lower if many backtests stay noop, but at threshold-tuning-time when threshold=0 it's near the worst case.

Replaces with: seed_inflight_limits_batched kernel. Inputs: market_targets_d, limit_slots_d (existing per-backtest array), per-backtest latency_ns_per_b. For each backtest with non-noop target, atomically scans its own MAX_LIMITS slots for active=0, writes an in-flight slot record. Slot allocation is per-backtest (each backtest has its own MAX_LIMITS slots) so no cross-backtest atomics. Overflow path increments a per-Pos overflow counter (already exists).

Loop 2 — close-detection at crates/ml-backtesting/src/sim.rs:658-674:

let mut closed_masks = vec![0u32; self.n_backtests];
let mut realised_returns = vec![0.0f32; self.n_backtests];
for b in 0..self.n_backtests {
    if prev_pos_per_block[b] != 0 {
        let now_pos = self.read_pos(b)?;          // host roundtrip per b
        if now_pos.position_lots == 0 {
            let abs_size = prev_pos_per_block[b].unsigned_abs() as f32;
            let pnl_delta = now_pos.realized_pnl - prev_pnl_per_block[b];
            let ret_per_lot = if abs_size > 0.0 { pnl_delta / abs_size } else { 0.0 };
            closed_masks[b] = prev_mask_per_block[b];
            realised_returns[b] = ret_per_lot;
        }
    }
}

read_pos(b) issues a memcpy_dtoh for the full PosFlat array per call, indexed at offset b. At n=140, every close-eligible backtest triggers a full 140-PosFlat readback. Cost: 140 × 2.5M = ~350M small dtoh transfers per quarter.

Replaces with: detect_close_transitions_batched kernel. Inputs: pos_d, prev_pos_lots_d, prev_pnl_d, prev_open_mask_d. Per backtest: detect close transition (prev != 0, now == 0), compute realised_return_per_lot, write to closed_horizon_mask_d + realised_return_d. The subsequent isv_kelly_update_on_close then dispatches without host involvement.

---

7. Migration plan (atomic commit boundaries)

Per feedback_no_partial_refactor: every contract change with all consumers in one commit. Per feedback_wire_everything_up: every kernel/module wired in the same commit as introduced.

Commit	Subject	Files touched	Test gate
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).
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.

Each commit gates the next. If any test fails, fix in the same commit (no follow-up debt).

---

8. Test coverage

All CUDA tests stay #[ignore = "requires CUDA"] and run locally via cargo test -- --ignored on RTX 3050 Ti. No CI gating for CUDA tests; convention is documented in CLAUDE.md and developers run them before commits.

New tests landing alongside their commit:

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
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
threshold_gate_allows_high_conviction	max_conviction ≥ threshold → buy.	P4
cost_deducted_at_each_fill	Round-trip 1-lot with cost=2.0 → net_pnl = gross - 4.0.	P4
kelly_state_sees_net_return	isv_kelly_update_on_close receives net-of-cost realised_return.	P4
forward_captured_matches_uncaptured	Captured graph reproduces uncaptured forward bit-identically.	P5
lob_sweep_140_variants_smoke	1-pod 1-window with n_parallel=140 distinct configs, all summary.json files differentiate.	P6

Existing tests that must continue to pass unchanged:

• All decision_floor_coldstart::* tests (P1 migrates them to BatchedSimConfig::from_uniform; behaviour unchanged)
• All lob_sim_fuzz / lob_sim_fixtures / lob_sim_integrated_fuzz tests
• save_load_roundtrip_preserves_all_weights (unchanged)
• cold_start_with_zero_floor_reproduces_old_bug (unchanged; tests the kernel pre-gate behaviour with floors=0)

---

9. Rate target validation

Measure at each gate:

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

If any gate misses, stop and diagnose before proceeding. Per feedback_stop_on_anomaly.

---

10. Risk register

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.
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.
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.

---

11. Out-of-scope / explicit deferrals

• bf16 mixed precision — separate follow-up spec, only if 1-h target missed. Stability validation gate (bit-near-equivalence on a 100k window) must precede ship.
• n_batch > 1 in forward across windows — would need a multi-window batched data loader; revisit only if 3-pod budget misses.
• Per-cell variation of model checkpoint (A/B candidates) — requires per-backtest forward, which breaks the forward-sharing premise. Treated as a separate sweep mode in a future spec.
• decision_stride reduction (4 → 8) — easy escape hatch flagged in risk register but not pursued here.

---

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.
• The vestigial loader.rs:274 capture_graph_a comment is fixed in P5.