diff --git a/docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md b/docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md new file mode 100644 index 000000000..7452545a7 --- /dev/null +++ b/docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md @@ -0,0 +1,1111 @@ +# SP18 — Combined Hold-Attractor Fix: D-leg (Reward-Shape) + B-leg (TD(λ) `Q(s')`) (Implementation Plan) + +> **For agentic workers:** REQUIRED SUB-SKILL: `superpowers:subagent-driven-development`. Steps use checkbox `- [ ]` syntax. Each task is bite-sized TDD: failing test → run-fail → implement → run-pass → commit. + +> **Status: v2 — UPGRADED to combined SP18+SP19 scope. DRAFT FOR HUMAN REVIEW.** Spec is at `docs/superpowers/specs/2026-05-08-sp18-reward-shape-hold-attractor-design.md` and MUST be reviewed first. **DESIGN DECISIONS DD1–DD8 + B-DD1–B-DD13 in the spec must be resolved before subagent dispatch.** Smoke gates anchor to vtj9r terminated 2026-05-09 + b5gmp empirical numbers — not speculative bands. + +> **v1 → v2 changelog**: v1 was D-leg only with B-leg deferred to SP19. v2 absorbs B-leg (TD(λ) `Q(s')` bootstrap) into combined SP18 per "no deferrals for complementary fixes" call. Refactor scopes have empty intersection per `feedback_no_partial_refactor`. + +--- + +## Goal + +Two-leg atomic fix for the persistent Q(Hold) attractor: + +1. **D-leg (reward-shape)**: Replace the reactive per-bar Hold cost (SP13 P0a + SP16 P2 + SP16 T3 chain) with a **structural Hold opportunity-cost reward** that pays each Hold-action bar the realized next-bar log-return as a counterfactual cost (asymmetric capped via new ISV slots). +2. **B-leg (TD(λ) `Q(s')` bootstrap)**: Replace `q_next = rewards` self-bootstrap at `gpu_experience_collector.rs:4143` with a real `Q_target(s', argmax_a Q_online(s', a))` (DDQN) computed by a new target-net forward pass on `next_states` at experience-collection time. + +**Concretely (assuming all DDs resolve to recommendations):** + +``` +PRE-SP18: + per-bar Hold cost = -ISV[380] × ISV[461] in [-0.625, -0.0025] + Hold Q-target distribution: tight band O(0.005) per bar + TD(λ) q_next = rewards (self-bootstrap; nstep_kernel.cu:96) + Symptom: Hold% climbs HD2→HD3 +13.4 pts; v_share monotone HD0→HD4 0.33→0.83 + +POST-SP18: + per-bar Hold reward = compute_sp18_hold_opportunity_cost( + -2.0 × vol_normalized_next_log_return, + isv[484] /* HOLD_REWARD_NEG_CAP */, + isv[483] /* HOLD_REWARD_POS_CAP */, + shaping_scale) + Hold Q-target distribution: O([-10, +5]) — same width as Long/Short + TD(λ) q_next = sum_branch Q_target[b, argmax_a_in_branch Q_online[b, a]] (DDQN per-branch) +``` + +**Critical pre-discovery (load-bearing — read before approving)** + +- `td_lambda_kernel` is wired with `td_lambda=0.9` (`gpu_experience_collector.rs:4154`, `config.rs:1427`) BUT with self-bootstrap `q_next = rewards` (line 4143). **Single production launch site.** +- Replay buffer does NOT store `q_next` → no schema migration. +- PER priorities recomputed at training time (`config.rs:341`) → no PER migration. +- C51 loss kernel `block_bellman_project_f` already uses target-net Q at loss time (`c51_loss_kernel.cu:188`) → no C51 loss migration. +- Target-net infra exists since SP14 P4.T1B-iv: `target_params_buf` (Polyak-tracked, `gpu_dqn_trainer.rs:5359`). +- VSN gating for next_states already wired: `vsn_gated_next_states_buf` (`gpu_dqn_trainer.rs:5834`). +- `compute_expected_q` (`experience_kernels.cu:4572`) is parametric on params/inputs — invoking with `target_params_buf` and `next_states` is already supported pattern. +- The 3 per-bar Hold cost sites are `experience_kernels.cu:3162-3166` (segment_complete branch), `:3637-3643` (per-bar mid-trade Hold), `:3710-3716` (per-bar flat-Hold). +- The SP13 host-side Hold-cost controller at `training_loop.rs:4171-4188` is what DD7(c) deletes. + +--- + +## Out of scope + +- **Distributional flavour of B-leg target-Q** (per-atom) — recommended scalar B-DD8.a. Distributional follow-up if 5-epoch smoke shows atom-distribution drift. +- PopArt normalization changes beyond per-fold EMA reset gate (B-DD11). +- SP11 reward-subsystem controller (slots 340–346) — Hold opp-cost flows through existing component buckets. +- SP15 reward-axis composer — sees the new reward unchanged in structure. +- IQN dueling — separate from this SP18. +- C51 v_range bounds — Hold opp-cost lives within `±2 × HOLD_REWARD_*_CAP` bounded; existing `[-10, +10]` v_range covers it. + +--- + +## Design decisions (REQUIRES HUMAN REVIEW — see spec for full discussion) + +| ID | Choice | Recommended | Spec section | +|---|---|---|---| +| DD1 | Combined D+B vs D-only / B-only | **combined** | spec §"Why combined scope" | +| DD2 | Hold opp-cost source | (a) `tgt[1]` | spec §DD2 | +| DD3 | Cap producer p99 source | (b) Long/Short trade-close p99 | spec §DD3 | +| DD4 | Hold opp-cost scope | (a) all Hold bars | spec §DD4 | +| DD5 | Cap asymmetry direction | (b) mirrored | spec §DD5 | +| DD6 | D-leg ISV slot allocation | as-specified, 10 slots | spec §DD6 | +| DD7 | SP13/SP16 cost retention | (c) replace, keep observation chain | spec §DD7 | +| DD8 | Cubin manifest registration | Phase 3.1 | spec §DD8 | +| B-DD1 | Target-net infra source | reuse `target_params_buf` | spec §B-DD1 | +| B-DD2 | DDQN argmax at bootstrap | yes | spec §B-DD2 | +| B-DD3 | Per-branch DDQN | 4 separate argmaxes | spec §B-DD3 | +| B-DD4 | PER priority migration | NOT NEEDED | spec §B-DD4 | +| B-DD5 | Target-net forward at collection | yes | spec §B-DD5 | +| B-DD6 | Accept ~30% slowdown | yes; profile in PP.5 | spec §B-DD6 | +| B-DD7 | Reuse target-net VSN params | yes | spec §B-DD7 | +| B-DD8 | Per-atom vs scalar target-Q | (a) scalar | spec §B-DD8 | +| B-DD9 | TD-error magnitude smoke gate | `[0.5, 5.0]` | spec §B-DD9 | +| B-DD10 | Atomic refactor scope | as-specified | spec §B-DD10 | +| B-DD11 | PopArt slot 63 reset at deploy | yes | spec §B-DD11 | +| B-DD12 | B-leg HEALTH_DIAG observability | as-specified | spec §B-DD12 | +| B-DD13 | B-leg ISV slot allocation | 12 slots [493..505) | spec §B-DD13 | + +> _Reviewer: confirm or redirect each DD. The plan below assumes all DDs resolve to the recommended option._ + +--- + +## Phase ordering + +D-leg and B-leg interleave. Each leg has its own atomicity constraint internally; between legs there are no shared buffers, so they can ship in separate commits inside the same branch. + +| Phase | Purpose | Tasks | +|---|---|---| +| Pre-Phase | Branch + worktree (reviewer) + ISV slots [483..505) + state_layout mirror + B-leg profiling | 5 | +| Phase 0 | D-leg + B-leg diagnostic emit (NO functional change) → confirm imbalance + confirm v_share monotone signature on synthetic | 5 | +| Phase 1 | Pre-dispatch consumer audit (BOTH legs) + promote audit script to pre-commit hook + atomic delete of SP13/SP16 cost-application chain (DD7=c) | 6 | +| Phase 2 | D-leg `compute_sp18_hold_opportunity_cost` + 3-site migration (atomic) | 5 | +| Phase 3 | D-leg `hold_reward_cap_update_kernel` adaptive producer + cubin manifest + Welford state | 5 | +| Phase 4 | B-leg target-net forward pass during experience collection (`compute_q_next_target_bootstrap`) + `next_states` build-site relocation | 5 | +| Phase 5 | B-leg `q_next` replacement at line 4143 + PopArt EMA reset gate (B-DD11) — atomic | 4 | +| Phase 6 | B-leg behavioral test suite (T8–T13) | 6 | +| Phase 7 | D-leg behavioral test suite (T1–T7) + integration tests | 8 | +| Phase 8 | Combined HEALTH_DIAG emit (D + B + synergy probes) | 4 | +| Phase 9 | 5-epoch L40S smoke validation against vtj9r/b5gmp baselines | 1 | +| Phase 10 | 30-epoch full validation (post-smoke pass) | 1 | + +**Total: 55 tasks.** Each is a single TDD cycle (failing test → implement → pass → commit). + +**Critical atomicity edges (NO L40S DISPATCH between):** +- Tasks 1.4 (delete 3 cost sites) → 2.2 (introduce SP18 device fn at same 3 sites): D-leg atomic. +- Tasks 4.x (target-net forward path landed but not consumed) → 5.1 (q_next replaced to consume it): B-leg atomic. + +--- + +## File structure + +``` +docs/superpowers/specs/2026-05-08-sp18-reward-shape-hold-attractor-design.md ← spec (v2) +docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md ← THIS file (v2) +crates/ml/src/cuda_pipeline/sp14_isv_slots.rs ← +22 SP18 ISV slots [483..505) +crates/ml/src/cuda_pipeline/state_layout.cuh ← C-side mirror +crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs ← bump ISV_TOTAL_DIM 483→505 + launchers +crates/ml/src/cuda_pipeline/experience_kernels.cu ← 3-site D-leg Hold-cost replacement +crates/ml/src/cuda_pipeline/trade_physics.cuh ← compute_sp18_hold_opportunity_cost +crates/ml/src/cuda_pipeline/hold_reward_cap_update_kernel.cu ← NEW D-leg producer +crates/ml/src/cuda_pipeline/hold_cost_scale_update_kernel.cu ← DELETE (DD7=c) +crates/ml/src/cuda_pipeline/gpu_experience_collector.rs ← B-leg q_next replacement + new compute_q_next_target_bootstrap +crates/ml/src/cuda_pipeline/td_bootstrap_diag_kernel.cu ← NEW B-leg HEALTH_DIAG producer +crates/ml/src/trainers/dqn/state_reset_registry.rs ← +22 fold-reset entries; remove [461..468) entries +crates/ml/src/trainers/dqn/trainer/training_loop.rs ← delete SP13 controller block; add combined HEALTH_DIAG emit +crates/ml/src/cuda_pipeline/build.rs / build.rs ← cubin manifest add hold_reward_cap_update + td_bootstrap_diag; remove hold_cost_scale_update +crates/ml/tests/sp18_hold_reward_oracle_tests.rs ← D-leg tests T1–T7 +crates/ml/tests/sp18_td_lambda_q_next_oracle_tests.rs ← NEW B-leg tests T8–T13 +docs/sp18-wireup-audit.md ← consumer audit checklist (Phase 1) +.git/hooks/pre-commit / scripts/audit_sp18_consumers.sh ← Phase 1 hook +``` + +--- + +## Hard rules (apply to every task) + +- **Pre-dispatch consumer audit script** (Phase 1 first task) — locked checklist surfaces missed consumers BEFORE migration. **PROMOTED TO PRE-COMMIT HOOK in Phase 1.5.** +- **Outcome-anchored smoke gates** — every threshold in Phases 9/10 cites empirical baselines (vtj9r HD2/HD3 or b5gmp Fold-1 final or pre/post-fix TD-error ratio). No speculative bands. +- **DESIGN DECISION markers** in spec; no false certainty in the plan. +- **TDD per task**: failing test first → implement → pass → commit. +- **No partial refactor**: D-leg 3-site migration atomic (Tasks 1.4 → 2.2 in same branch, no L40S between); B-leg q_next replacement atomic (Tasks 4.x → 5.1 in same branch, no L40S between). +- **No CPU compute** — D-leg opp-cost producer + B-leg target-net forward + reward computation pure GPU; CPU oracle tests use `MappedF32Buffer`. +- **No atomicAdd** — producer reductions via `sp4_histogram_p99` per `pearl_fused_per_group_statistics_oracle`. +- **Mapped-pinned only** for any host↔GPU transfer. +- **ISV-driven adaptive bounds** — pos/neg caps via producer; safety_factor=1.5; bounds are dimensional Category-1 floors. +- **Wire-up audit per task** — every kernel/buffer/slot wired in same commit; no orphans. +- **No host branches in captured graph** — B-leg target-net forward must be pure kernel launches inside `exp_fwd_graph` capture region. +- **Pearl candidate per phase** for reviewer to land or reject. + +--- + +## Pre-Phase: Branch + worktree + ISV slot allocation + B-leg profiling + +### Task PP.1: Reviewer creates worktree (do not run as part of subagent dispatch) + +```bash +cd /home/jgrusewski/Work/foxhunt +git worktree add .worktrees/sp18-reward-shape -b feat/sp18-reward-shape feat/sp17-dueling +cd .worktrees/sp18-reward-shape +``` + +- [ ] **Step 1: Reviewer runs the above before subagent dispatch.** + +The branch forks from `feat/sp17-dueling` HEAD `0c57a5a31`. SP18 layers on top of SP17's identifiability projection. + +No tests, no commits. + +### Task PP.2: Allocate 22 SP18 ISV slots [483..505) + +**Files:** +- Modify: `crates/ml/src/cuda_pipeline/sp14_isv_slots.rs` — append SP18 section +- Test: same file — new `#[test] sp18_combined_slot_layout_locked` + +- [ ] **Step 1: Write failing tests** in `mod tests` of sp14_isv_slots.rs: + +```rust +/// Lock SP18 (2026-05-08 v2 combined) D-leg + B-leg slot layout. +/// 22 contiguous slots [483..505): 10 D-leg + 12 B-leg. +#[test] +fn sp18_combined_slot_layout_locked() { + // D-leg [483..493) + assert_eq!(SP18_HOLD_REWARD_SLOT_BASE, 483); + assert_eq!(SP18_HOLD_REWARD_SLOT_END, 493); + assert_eq!(HOLD_REWARD_POS_CAP_INDEX, 483); + assert_eq!(HOLD_REWARD_NEG_CAP_INDEX, 484); + assert_eq!(HOLD_REWARD_DECOMP_DIAG_INDEX, 485); + assert_eq!(HOLD_OPP_COST_FIRE_RATE_EMA_INDEX, 486); + assert_eq!(HRC_TARGET_MEAN_INDEX, 487); + assert_eq!(HRC_TARGET_M2_INDEX, 488); + assert_eq!(HRC_DIFF_MEAN_INDEX, 489); + assert_eq!(HRC_DIFF_M2_INDEX, 490); + assert_eq!(HRC_PREV_TARGET_INDEX, 491); + assert_eq!(HRC_SAMPLE_COUNT_INDEX, 492); + // B-leg [493..505) + assert_eq!(SP18_TD_BOOTSTRAP_SLOT_BASE, 493); + assert_eq!(SP18_TD_BOOTSTRAP_SLOT_END, 505); + assert_eq!(TD_ERROR_MAG_EMA_INDEX, 493); + assert_eq!(Q_NEXT_TARGET_P99_INDEX, 494); + assert_eq!(Q_NEXT_MINUS_REWARD_P99_INDEX, 495); + assert_eq!(V_SHARE_TREND_DIAG_INDEX, 496); + assert_eq!(POPART_RESET_FLAG_INDEX, 497); + assert_eq!(TDB_TARGET_MEAN_INDEX, 498); + assert_eq!(TDB_TARGET_M2_INDEX, 499); + assert_eq!(TDB_DIFF_MEAN_INDEX, 500); + assert_eq!(TDB_DIFF_M2_INDEX, 501); + assert_eq!(TDB_PREV_TARGET_INDEX, 502); + assert_eq!(TDB_SAMPLE_COUNT_INDEX, 503); + assert_eq!(SP18_B_LEG_RESERVED_INDEX, 504); + // Sentinels — Pearl-A first-observation bootstrap (D-leg). + assert_eq!(SENTINEL_HOLD_REWARD_POS_CAP, 5.0); + assert_eq!(SENTINEL_HOLD_REWARD_NEG_CAP, -10.0); + assert_eq!(SENTINEL_HOLD_REWARD_DIAG, 0.0); + assert_eq!(SENTINEL_HOLD_OPP_FIRE_RATE, 0.0); + // Sentinels — B-leg. + assert_eq!(SENTINEL_TD_ERROR_MAG_EMA, 0.0); + assert_eq!(SENTINEL_Q_NEXT_TARGET_P99, 0.0); + assert_eq!(SENTINEL_Q_NEXT_MINUS_REWARD_P99, 0.0); + assert_eq!(SENTINEL_V_SHARE_TREND_DIAG, 0.0); + assert_eq!(SENTINEL_POPART_RESET_FLAG, 1.0); // 1.0 first epoch ⇒ PopArt reset + // Bounds (D-leg, Category-1 floors). + assert_eq!(HOLD_REWARD_POS_CAP_MIN_BOUND, 0.5); + assert_eq!(HOLD_REWARD_POS_CAP_MAX_BOUND, 50.0); + // Producer constants. + assert_eq!(HOLD_REWARD_CAP_SAFETY_FACTOR, 1.5); +} + +#[test] +fn all_sp18_slots_fit_within_isv_total_dim() { + use crate::cuda_pipeline::gpu_dqn_trainer::ISV_TOTAL_DIM; + assert!( + SP18_TD_BOOTSTRAP_SLOT_END <= ISV_TOTAL_DIM, + "SP18_TD_BOOTSTRAP_SLOT_END={} exceeds ISV_TOTAL_DIM={}; \ + bump ISV_TOTAL_DIM in gpu_dqn_trainer.rs", + SP18_TD_BOOTSTRAP_SLOT_END, ISV_TOTAL_DIM, + ); +} +``` + +Run: `SQLX_OFFLINE=true cargo test -p ml --lib sp18_combined_slot_layout_locked all_sp18_slots_fit_within_isv_total_dim` +Expected: FAIL — constants undefined. + +- [ ] **Step 2: Add SP18 section to sp14_isv_slots.rs** (D-leg + B-leg constants per the spec § "B-leg ISV slots [493..505)"). + +- [ ] **Step 3: Bump `ISV_TOTAL_DIM`** at `gpu_dqn_trainer.rs:2571` from 483 to 505. Append SP18 v2 paragraph to the docstring chain. + +- [ ] **Step 4: Update `layout_fingerprint_seed`** with the 22 new slot names. + +- [ ] **Step 5: Update `state_layout.cuh` C-side mirror** with all 22 slots and bounds. + +- [ ] **Step 6: Run test, verify PASS.** + +- [ ] **Step 7: Commit.** + +``` +plan(sp18 v2): allocate 22 ISV slots [483..505) for D-leg + B-leg + +Adds combined SP18 D-leg Hold-reward adaptive-cap slots [483..493) + B-leg +TD(λ) Q(s') bootstrap diagnostics + PopArt reset flag at [493..505). + +Per spec DD6 (D-leg 10 slots) + B-DD13 (B-leg 12 slots). + +Pearl-A first-observation bootstrap sentinels match position-side SP14 P0-A +REWARD_POS_CAP_ADAPTIVE pattern. Slot 497 POPART_RESET_FLAG sentinel = 1.0 +(per B-DD11: triggers per-fold PopArt EMA reset on first SP18 deployment). + +ISV_TOTAL_DIM bumped 483 → 505; layout_fingerprint_seed updated; +state_layout.cuh mirror in lockstep. + +Spec: docs/superpowers/specs/2026-05-08-sp18-reward-shape-hold-attractor-design.md +Plan: docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md + +Co-Authored-By: Claude Opus 4.7 (1M context) +``` + +### Task PP.3: Fold-reset registry entries for new slots + +**Files:** Modify `crates/ml/src/trainers/dqn/state_reset_registry.rs` + +- [ ] **Step 1: Failing test** — append `#[test] sp18_fold_reset_entries_present` asserting each of 22 slot indices appears with the matching sentinel and a description containing "SP18 D-leg" or "SP18 B-leg". + +- [ ] **Step 2: Add 22 entries** to `state_reset_registry()`. Note: slot 497 POPART_RESET_FLAG sentinel = **1.0** (one-shot, per B-DD11), slot 504 RESERVED sentinel = 0.0. + +- [ ] **Step 3: Run, commit.** + +### Task PP.4: Document SP13/SP16 slot retirement + +**Files:** Modify `state_reset_registry.rs` — find the entries for slots 461 (HOLD_COST_SCALE_INDEX) and [462..468) (HCS_*) and update their description from "SP16 P2/T3 Hold cost scale Wiener-α" to "SP18 v2: RETIRED — replaced by D-leg structural reward (sentinel 0.0 keeps slot bit-identical for checkpoint compat)". Sentinels stay 0.0 — slots remain allocated for layout fingerprint compatibility per the SP14-C.1 RESERVED-gap pattern. + +- [ ] **Step 1: Failing test** — `#[test] sp16_t2_t3_slots_marked_retired` asserts the description changed. +- [ ] **Step 2: Update.** +- [ ] **Step 3: Run, commit.** + +### Task PP.5: B-leg collection-time profiling baseline (B-DD6) + +**Why first:** B-DD6 accepts ~30% experience-collection slowdown for the new target-net forward pass. Profile current collection time on H100 to confirm headroom BEFORE Phase 4 lands the new path. If current collection is already >70% of step budget, B-leg architecture changes (e.g. share Pass 2 forward output between collection and loss) before commit. + +**Files:** No new files; runs against current `feat/sp17-dueling` HEAD. + +- [ ] **Step 1: Reviewer-run profiling** + +```bash +./scripts/argo-train.sh --model dqn --epochs 1 --label sp18-pp5-profile --profile nsys +``` + +- [ ] **Step 2: Reviewer reads nsys profile** for `gpu_experience_collector.rs::collect_experience` end-to-end ms/cycle. Compare against budget. + +- [ ] **Step 3: Document baseline** in plan as a new "Pre-Phase 5 evidence" subsection. If headroom < 30%, **HALT** and convene human review before Phase 4. + +No commit unless evidence is appended. + +> _DESIGN NOTE_: subagent CANNOT run this step (reviewer-only). Hands off to reviewer at PP.5. + +--- + +## Phase 0 — Per-step reward + V-share diagnostic (NO functional change) + +**Why first:** the plan claims D-leg's Hold per-step reward is too small AND B-leg's bootstrap-bias drives v_share elevation. **Phase 0 instruments without changing.** If either hypothesis is wrong, abort that leg before reward-shape change or B-leg surgery. + +### Task 0.1: D-leg per-action reward decomposition HEALTH_DIAG emit + +**Files:** Modify `crates/ml/src/trainers/dqn/trainer/training_loop.rs` per-epoch HEALTH_DIAG block. + +- [ ] **Step 1: Failing test** — synthetic test in `tests/sp18_hold_reward_oracle_tests.rs` (NEW file) builds 10-bar episode with known Hold/Long/Short actions, runs experience-collector, asserts new `HEALTH_DIAG[reward_decomp]` line appears with expected values. + +- [ ] **Step 2: Implement** — read `reward_components_per_sample` (already populated by `experience_kernels.cu`, slots [0..6) per `out_off`). Per-action-bin EMAs of: + - `r_micro`, `r_opp_cost`, `r_popart` per action bin (Hold / Long / Short / Flat) + - per-bar `|reward|` per action bin + - per-bar fire-rate (frac bars with non-zero component) per action bin + +Stored in 5×4=20-element `MappedF32Buffer` `reward_decomp_diag_buf` on `GpuDqnTrainer`, computed by NEW `reward_decomp_diag_kernel.cu` (block tree-reduce). + +- [ ] **Step 3: Cubin manifest entry** for `reward_decomp_diag_kernel.cu`. + +- [ ] **Step 4: Wire HEALTH_DIAG emit**: +``` +HEALTH_DIAG[N]: reward_decomp [hold(micro=X opp=Y popart=Z fire=W) long(...) short(...) flat(...)] +``` + +- [ ] **Step 5: Run test, commit.** + +### Task 0.2: B-leg V-share trajectory + TD-error magnitude HEALTH_DIAG emit + +**Why:** establishes pre-fix baseline for the B-leg synergy hypothesis (v_share monotone HD0→HD4) and TD-error magnitude (input to B-DD9 ratio gate). + +**Files:** Modify `training_loop.rs` per-epoch HEALTH_DIAG block. + +- [ ] **Step 1: Failing test** — same NEW file as 0.1; assert `HEALTH_DIAG[v_share_trajectory]` and `HEALTH_DIAG[td_error_pre_fix]` lines appear. + +- [ ] **Step 2: Implement** — V_SHARE EMA already exists at slots [478..482) per SP17. Compute slope across last 4 epochs as `(EMA[now] - EMA[now-4]) / 4`; write to slot 496 `V_SHARE_TREND_DIAG`. TD-error magnitude is the per-batch L1 norm of `td_errors` (already computed for PER priority recomputation at training time); EMA into a new producer. + +- [ ] **Step 3: Wire emit**: +``` +HEALTH_DIAG[N+1]: v_share_traj [dir_slope=X mag_slope=Y ord_slope=Z urg_slope=W] +HEALTH_DIAG[N+2]: td_error_pre [magnitude_ema=X] +``` + +- [ ] **Step 4: Run test, commit.** + +### Task 0.3: 1-epoch L40S diagnostic dispatch (REVIEWER-EXECUTED) + +- [ ] **Step 1: Reviewer pushes branch + dispatches** + +```bash +git push -u origin feat/sp18-reward-shape +./scripts/argo-train.sh --model dqn --epochs 1 --label sp18-phase0-decomp +``` + +- [ ] **Step 2: Reviewer reads HEALTH_DIAG output** + +Tail `reward_decomp` + `v_share_traj` + `td_error_pre` lines across the epoch. Compute per-action-bin trajectory + V-share slope + TD-error magnitude baseline. + +### Task 0.4: KILL CRITERION evaluation (combined) + +``` +ABORT D-leg if: + |r_micro_hold_ema| > 2 × |r_micro_long_ema| (Hold cost is already disproportionately large) + AND + Hold% trajectory is FALLING (controller is winning) + +ABORT B-leg if: + v_share_slope < 0 (v_share is already decreasing without intervention) + AND + v_share[HD0] < 0.5 (v-share didn't rise into the danger zone) + +PROCEED on combined if: + D-leg PROCEED criterion: |r_micro_hold_ema| < 0.5 × max(|r_popart_long_ema|, |r_popart_short_ema|) + AND Hold% trajectory is RISING + AND + B-leg PROCEED criterion: v_share_slope > 0.05 over HD0→HD3 + AND v_share[HD3] > 0.7 + +INCONCLUSIVE if either ABORT/PROCEED is mixed — convene human review +``` + +- [ ] **Step 4: Document result + decision** — append "Phase 0 evidence" subsection to this plan. + +No commit unless evidence appended. + +### Task 0.5: Pearl candidate (post-Phase 0 if PROCEED) + +> Pearl draft: `pearl_diagnostic_decomposition_before_reward_intervention.md` (covers both D-leg per-action reward decomposition AND B-leg V-share trajectory pre-fix observation). + +- [ ] **Step 1: Draft pearl in `docs/superpowers/notes/`.** + +--- + +## Phase 1 — Pre-dispatch consumer audit + atomic delete + pre-commit hook + +### Task 1.1: Combined consumer audit script (LOCKED CHECKLIST for both legs) + +**Why first:** surfaces missed consumers BEFORE migration. Per task input — don't trust this plan author's expected list; the script enforces it. + +**Files:** NEW `scripts/audit_sp18_consumers.sh` (executable): + +```bash +#!/usr/bin/env bash +# scripts/audit_sp18_consumers.sh +set -euo pipefail + +echo "=== D-leg: Slot 380 (HOLD_COST_INDEX) consumers ===" +grep -rn "ISV_HOLD_COST_IDX\|HOLD_COST_INDEX\|isv\[380\]\|read_isv_signal_at(HOLD_COST_INDEX)\|read_isv_signal_at(380)" \ + crates/ bin/ services/ scripts/ docs/ \ + --include="*.rs" --include="*.cu" --include="*.cuh" --include="*.toml" --include="*.md" \ + || echo "(no hits)" + +echo "" +echo "=== D-leg: Slot 461 (HOLD_COST_SCALE_INDEX) consumers ===" +grep -rn "ISV_HOLD_COST_SCALE_IDX\|HOLD_COST_SCALE_INDEX\|isv\[461\]" \ + crates/ bin/ services/ scripts/ docs/ \ + --include="*.rs" --include="*.cu" --include="*.cuh" --include="*.toml" --include="*.md" \ + || echo "(no hits)" + +echo "" +echo "=== D-leg: Slots [462..468) (HCS_* Welford) consumers ===" +grep -rn "HCS_TARGET_MEAN\|HCS_TARGET_M2\|HCS_DIFF_MEAN\|HCS_DIFF_M2\|HCS_PREV_TARGET\|HCS_SAMPLE_COUNT\|isv\[46[234567]\]" \ + crates/ bin/ services/ scripts/ docs/ \ + --include="*.rs" --include="*.cu" --include="*.cuh" --include="*.toml" --include="*.md" \ + || echo "(no hits)" + +echo "" +echo "=== D-leg: hold_cost_scale_update_kernel references ===" +grep -rn "hold_cost_scale_update\|HoldCostScaleUpdate" \ + crates/ bin/ services/ scripts/ docs/ \ + --include="*.rs" --include="*.cu" --include="*.cuh" --include="*.toml" --include="*.md" --include="*.json" \ + || echo "(no hits)" + +echo "" +echo "=== D-leg: hold_rate_observer_kernel (RETAINED — diag chain stays per DD7=c) ===" +grep -rn "hold_rate_observer\|HoldRateObserver\|HOLD_RATE_OBSERVED_EMA_INDEX\|HOLD_RATE_TARGET_INDEX" \ + crates/ bin/ services/ scripts/ docs/ \ + --include="*.rs" --include="*.cu" --include="*.cuh" --include="*.toml" --include="*.md" \ + || echo "(no hits)" + +echo "" +echo "=== D-leg: Per-bar Hold-cost subtraction sites (the 3 lines we change in Phase 2) ===" +grep -rn "isv_signals_ptr\[ISV_HOLD_COST_IDX\]\|isv\[ISV_HOLD_COST_IDX\]\|hold_cost_scale" \ + crates/ml/src/cuda_pipeline/experience_kernels.cu + +echo "" +echo "=== B-leg: td_lambda_kernel launch sites ===" +grep -rn "td_lambda_kernel\b" \ + crates/ bin/ services/ \ + --include="*.rs" --include="*.cu" \ + || echo "(no hits)" + +echo "" +echo "=== B-leg: q_next argument origin (the line 4143 site) ===" +grep -n "q_next_bootstrap\|q_next_target\|let q_next\b" \ + crates/ml/src/cuda_pipeline/gpu_experience_collector.rs + +echo "" +echo "=== B-leg: rewards_out consumers (downstream of TD(λ) output) ===" +grep -rn "rewards_out\b" \ + crates/ml/src/cuda_pipeline/gpu_experience_collector.rs \ + | grep -v "^.*://\|^.*: \*\|^.*: //" +echo "" +echo "(Note: rewards_out is the OUTPUT of td_lambda_kernel; B-leg fix changes the q_next INPUT)" + +echo "" +echo "=== B-leg: PER priority computation sites (verify B-DD4: NO migration needed) ===" +grep -rn "update_priorities_from_td\|update_priorities_gpu\b" \ + crates/ bin/ services/ \ + --include="*.rs" \ + || echo "(no hits)" + +echo "" +echo "=== B-leg: Replay buffer schema — q_next storage check (verify NO migration) ===" +grep -rn "q_next\|Q_next" \ + crates/ml/src/agents/ crates/ml/src/trainers/dqn/replay_buffer_type.rs crates/ml/src/dqn/gpu_replay_buffer.rs 2>/dev/null \ + --include="*.rs" \ + || echo "(no hits — confirms no q_next stored in buffer)" + +echo "" +echo "=== B-leg: c51_loss_kernel target-Q origin (verify already correct at loss time) ===" +grep -rn "block_bellman_project_f\|target_probs\b" \ + crates/ml/src/cuda_pipeline/c51_loss_kernel.cu \ + | head -10 + +echo "" +echo "=== B-leg: target_params_buf consumers (verify Polyak-tracked, available at collection time) ===" +grep -rn "target_params_buf\b" \ + crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \ + | head -20 + +echo "" +echo "=== B-leg: PopArt slot 63 references (verify B-DD11 reset is the only PopArt change) ===" +grep -rn "PopArt\|POPART\|popart\|slot 63\|POPART_INDEX" \ + crates/ bin/ services/ \ + --include="*.rs" --include="*.cu" --include="*.cuh" \ + | head -20 + +echo "" +echo "=== END ===" +``` + +- [ ] **Step 1: Add the audit script** to `scripts/audit_sp18_consumers.sh` (executable). + +- [ ] **Step 2: Run the audit** and capture output to `docs/sp18-wireup-audit.md` (NEW file). Lock the captured consumer list as a checklist. + +- [ ] **Step 3: Manually inspect** for any hit not in the spec's expected list. If a missed consumer surfaces: + - **Halt the migration plan**. + - Update spec §"Architecture Changes" with the new consumer. + - Re-run audit until clean. + +- [ ] **Step 4: Commit the audit script + audit doc**. + +### Task 1.2: Promote audit script to pre-commit hook + +**Why:** Open Q-B confirmed — generalizable beyond SP18, catches future SP-chain consumer drift. + +**Files:** +- NEW: `.git/hooks/pre-commit` invoking `scripts/audit_sp18_consumers.sh` and failing the commit if the output diff against the locked `docs/sp18-wireup-audit.md` shows new entries (use `diff -u` against the committed audit doc). +- Modify: `scripts/audit_sp18_consumers.sh` — add a `--check` mode that exits 1 if the captured output differs from the previous run's snapshot. + +- [ ] **Step 1: Failing test** — `#[test] pre_commit_hook_present` asserts `.git/hooks/pre-commit` is executable and contains `audit_sp18_consumers.sh`. + +- [ ] **Step 2: Add hook + check mode**. + +- [ ] **Step 3: Run, commit.** Note: the pre-commit hook is installed in the worktree's `.git/hooks/` (per-worktree, not committed to the repo by git's design); the hook source lives in `scripts/install_audit_pre_commit_hook.sh` which `cp`s into `.git/hooks/`. + +### Task 1.3: Delete `hold_cost_scale_update_kernel.cu` + cubin manifest entry + +**Files:** +- DELETE: `crates/ml/src/cuda_pipeline/hold_cost_scale_update_kernel.cu` +- Modify: `build.rs` cubin manifest — remove entry +- Modify: `gpu_dqn_trainer.rs` — remove `launch_hold_cost_scale_update` +- Modify: `training_loop.rs:564` — remove launch call +- Modify: `training_loop.rs:631-667` — remove diag emit block + +- [ ] **Step 1: Failing test** — `#[test] hold_cost_scale_kernel_deleted` asserts `Path::new("crates/ml/src/cuda_pipeline/hold_cost_scale_update_kernel.cu").exists() == false` AND build.rs manifest doesn't list it. + +- [ ] **Step 2: Delete + remove**. + +- [ ] **Step 3: Run** `SQLX_OFFLINE=true cargo check --workspace` — expected PASS. + +- [ ] **Step 4: Run all tests** — confirm no test referenced this kernel. + +- [ ] **Step 5: Commit.** + +### Task 1.4: Delete SP13 P0a host-side Hold-cost controller block + +**Files:** Modify `training_loop.rs:4171-4188` — DELETE the controller block. + +- [ ] **Step 1: Failing test** — `#[test] sp13_hold_cost_controller_deleted` asserts file source does NOT contain `let excess = (observed - target).max(0.0);`. + +- [ ] **Step 2: Delete the block.** + +- [ ] **Step 3: Workspace check + tests pass.** + +- [ ] **Step 4: Commit.** + +### Task 1.5: Remove the 3 per-bar Hold-cost subtraction sites (placeholder for Phase 2) + +**Files:** Modify `experience_kernels.cu` at lines 3162–3166, 3637–3643, 3710–3716. + +**Important:** This task **only deletes** the existing `if (dir_idx == DIR_HOLD && isv_signals_ptr != NULL) { ... -= ... }` blocks. Phase 2 introduces the SP18 reward in lockstep. Intermediate state has Hold paying zero per-bar — **NOT runnable for L40S smoke until Phase 2.2 lands**. Per `feedback_no_partial_refactor`. + +- [ ] **Step 1: Failing test** — `#[test] sp13_hold_cost_application_deleted` asserts kernel source does NOT contain `isv_signals_ptr[ISV_HOLD_COST_IDX] * hold_cost_scale`. + +- [ ] **Step 2: Delete the 3 blocks**. Replace with: +```c +/* SP18 v2 Phase 1 (DD7=c): SP13/SP16 Hold-cost subtraction deleted; SP18 + * D-leg structural opportunity-cost reward replaces in Phase 2 (Task 2.2). + * Intermediate state: Hold pays zero per-bar reward. NOT runnable for + * L40S smoke until Phase 2 lands. Per `feedback_no_partial_refactor`. */ +``` + +- [ ] **Step 3: Workspace check + tests pass.** + +- [ ] **Step 4: Commit (NO L40S DISPATCH between this commit and Task 2.2).** + +### Task 1.6: Phase 1 wire-up audit close-out + +- [ ] **Step 1: Re-run audit script.** Confirm: + - Slot 380: 0 hits outside state_reset_registry sentinel + sp14_isv_slots constants + audit doc itself + - Slot 461: same + - `hold_cost_scale_update`: 0 hits anywhere + - `hold_rate_observer`: still present (DD7=c) + - B-leg sections: 1 hit at line 4143 (the q_next clone — to be replaced in Phase 5) + - PER priority sites: present and unchanged (B-DD4 confirmed) + - C51 target-Q at loss: present and unchanged + - target_params_buf: Polyak-tracked, ready for B-leg consumption + - PopArt slot 63 sites: catalogued for B-DD11 reset + +- [ ] **Step 2: Update `docs/sp18-wireup-audit.md`** with post-deletion grep output. Cross out checklist items. + +- [ ] **Step 3: Commit updated audit doc.** + +--- + +## Phase 2 — D-leg `compute_sp18_hold_opportunity_cost` device fn + 3-site migration (atomic) + +### Task 2.1: CPU oracle test (T1) for `compute_sp18_hold_opportunity_cost` + +**Files:** `crates/ml/tests/sp18_hold_reward_oracle_tests.rs` + +- [ ] **Step 1: Write failing test** (T1 from spec) — see spec § "Behavioral Test Predictions". + +- [ ] **Step 2: Run** — should pass immediately (pure math). + +- [ ] **Step 3: Commit.** + +### Task 2.2: Add `compute_sp18_hold_opportunity_cost` device fn + wire 3 sites (atomic) + +**Files:** +- Modify: `trade_physics.cuh` — add device fn after `compute_lump_sum_opp_cost` (line ~590) +- Modify: `experience_kernels.cu` — wire at the 3 placeholder comment sites from Task 1.5 + +(Device fn body and wire-site code blocks identical to v1 plan Task 2.2, see v1 history; preserved verbatim here. Sign convention: positive `tgt[1]` log-return → negative Hold reward (opp cost paid); negative log-return → positive Hold reward (correctly avoided). Asymmetric capped via `compute_asymmetric_capped_pnl(base, neg_cap, pos_cap)`.) + +- [ ] **Step 1: Failing GPU oracle test (T2)** — wrap device fn via `sp18_hold_opp_test_kernel.cu` (mirror SP17 oracle test pattern). Add cubin manifest entry. + +- [ ] **Step 2: Implement device fn + wire 3 sites + test wrapper.** + +- [ ] **Step 3: Run, verify PASS.** + +- [ ] **Step 4: Commit (THIS COMMIT closes the partial-refactor state from Task 1.5 — atomic rule satisfied).** + +### Task 2.3: Sentinel-fallback + cold-start sanity test + +- [ ] **Step 1: Failing test** — verify `isv[483] == SENTINEL_HOLD_REWARD_POS_CAP` falls back to macro for deterministic cold-start. + +- [ ] **Step 2: Implement** — call-site clamp `fmaxf(HOLD_REWARD_POS_CAP_MIN_BOUND, fminf(isv[483], HOLD_REWARD_POS_CAP_MAX_BOUND))`; sentinel within ε of MIN routes through clamp naturally. + +- [ ] **Step 3: Run, commit.** + +### Task 2.4: Phase 2 wire-up audit + pearl candidate + +- [ ] **Step 1: Re-run audit** — confirm 0 hits for SP13 cost-application, sub-3 hits for slot 380. + +- [ ] **Step 2: Pearl draft** — `pearl_replace_reactive_controller_with_structural_signal.md`. + +### Task 2.5: Compile + smoke check (NO L40S; local only) + +- [ ] **Step 1**: `SQLX_OFFLINE=true cargo check --workspace` — PASS. +- [ ] **Step 2**: `SQLX_OFFLINE=true cargo test -p ml --lib` — all SP18 tests pass. +- [ ] **Step 3**: Local smoke `FOXHUNT_TEST_DATA=test_data/futures-baseline cargo test -p ml --lib -- smoke_tests --ignored --nocapture` — runs to completion without NaN. + +If either fails: STOP. Diagnose. No commit. If pass: no new commit (verification only). + +--- + +## Phase 3 — D-leg `hold_reward_cap_update_kernel` adaptive producer + +### Task 3.1: Add cubin manifest entry + skeleton kernel file (DD8=Phase 3.1) + +**Files:** +- NEW: `crates/ml/src/cuda_pipeline/hold_reward_cap_update_kernel.cu` — skeleton with `extern "C" __global__ void hold_reward_cap_update(...)` and a single `if (threadIdx.x != 0 || blockIdx.x != 0) return;` early-return body. +- Modify: `build.rs` — register cubin. + +- [ ] **Step 1: Failing test** — `#[test] #[ignore = "requires GPU"] hold_reward_cap_kernel_loads`. +- [ ] **Step 2: Implement skeleton + manifest.** +- [ ] **Step 3: Run, commit.** + +### Task 3.2: Implement producer body — p99 + Wiener-α + Welford + +Mirror the deleted `hold_cost_scale_update_kernel.cu` Wiener-α machinery: +- Input: `step_ret_per_sample` over Long/Short trade-close events (`experience_kernels.cu:2294`). +- p99(|step_ret|) via `sp4_histogram_p99` (per `pearl_fused_per_group_statistics_oracle`, no atomicAdd). +- × `HOLD_REWARD_CAP_SAFETY_FACTOR=1.5`. +- POS = clamp result; NEG = `-2.0f × POS` (DD5=b). +- Wiener-α from Welford state `[487..493)`. +- Pearl-A first-observation REPLACE on sentinel. + +- [ ] **Step 1: Failing test (T5 from spec).** +- [ ] **Step 2: Implement.** +- [ ] **Step 3: Run, commit.** + +### Task 3.3: Add launcher + per-epoch wiring + +**Files:** +- Modify: `gpu_dqn_trainer.rs` — `launch_hold_reward_cap_update` (mirror `launch_reward_cap_update`). +- Modify: `training_loop.rs` per-epoch HEALTH_DIAG block — call AFTER `launch_reward_cap_update`. + +- [ ] **Step 1: Failing test** — fake checkpoint reproduction. +- [ ] **Step 2: Implement.** +- [ ] **Step 3: Run, commit.** + +### Task 3.4: Wire-up audit + +- [ ] **Step 1: Audit script** — confirm new producer at exactly 1 call site; ISV[483]/[484] reads at exactly 3 sites. +- [ ] **Step 2: Pearl draft** — `pearl_distributional_shape_match_for_action_neutrality.md` (held in pending; lands post-Phase 9 if smoke passes). + +### Task 3.5: Phase 3 close compile-clean + +- [ ] **Step 1**: `cargo check --workspace` PASS. +- [ ] **Step 2**: `cargo test -p ml --lib` all D-leg tests PASS. +- [ ] **Step 3**: Local smoke runs without NaN. No commit. + +--- + +## Phase 4 — B-leg target-net forward pass during experience collection (additive, not yet wired into TD(λ)) + +### Task 4.1: `next_states` build-site relocation + +**Files:** Modify `gpu_experience_collector.rs` — move `let next_states = build_next_states_f32(...)` from line 4182 to before line 4137 (the TD(λ) `if` block). + +- [ ] **Step 1: Failing test** — `#[test] next_states_built_before_td_lambda` introspects via `include_str!` that `build_next_states_f32` precedes `td_lambda_kernel` launch. +- [ ] **Step 2: Move statement.** Verify nothing else (downstream) relied on next_states being built post-TD(λ); the audit script's grep on `next_states` flagged 0 such consumers. +- [ ] **Step 3: Run, commit.** This is a no-op behaviorally (no observer of next_states moves); pure ordering change. + +### Task 4.2: New `compute_q_next_target_bootstrap` skeleton (additive, dead code) + +**Files:** Modify `gpu_experience_collector.rs` — add new method: + +```rust +/// Compute Q(s', argmax_a Q_online(s', a)) per-sample using DDQN per-branch +/// argmax. Uses target-net params to evaluate the action selected by the +/// online net. Per spec B-DD2 + B-DD3 + B-DD8.a. +/// +/// Returns: CudaSlice of length B (one scalar Q-bootstrap value per sample). +/// +/// Phase 4 lands the method as additive dead code. Phase 5 wires it into +/// the TD(λ) launch by replacing the `q_next` argument at line 4143. +pub fn compute_q_next_target_bootstrap( + &mut self, + next_states: &CudaSlice, + target_params_ptr: u64, + online_params_ptr: u64, + batch_size: usize, +) -> Result, MLError> { + /* Phase 4 skeleton — implemented in Task 4.3-4.5. */ + todo!("SP18 B-leg Phase 4: implement target-net + DDQN argmax + scalar bootstrap") +} +``` + +> **EXCEPTION TO `feedback_no_stubs`**: this is the *only* explicit `todo!()` marker in the plan; it is justified because Phase 4 has 5 sub-tasks each landing an incremental piece, and the `todo!()` is replaced atomically in Task 4.5 before any caller wires it. Per `feedback_no_partial_refactor` the method is dead code (no caller) until Task 5.1, so the `todo!()` cannot fire in production. Grep audit at Phase 5.1 close-out asserts `todo!()` is absent. + +- [ ] **Step 1: Failing test** — `#[test] compute_q_next_target_bootstrap_method_exists` (introspects via `include_str!`). +- [ ] **Step 2: Add method skeleton.** +- [ ] **Step 3: Run, commit.** + +### Task 4.3: Implement online-net forward pass on next_states for DDQN argmax + +**Files:** Modify `compute_q_next_target_bootstrap` — implement the **online-net Pass 3 mirror** part: +- Use `submit_forward_ops_ddqn`'s logic (already exists, runs online net on next_states for DDQN argmax in the loss path) — refactor into a callable helper if needed, OR replicate inline using the same `vsn_gated_next_states_for_ddqn_buf` machinery. +- Output: per-branch argmax actions in `[B, 4]` int buffer (one per dir/mag/ord/urg branch). + +- [ ] **Step 1: Failing GPU oracle test** — synthetic small batch; verify per-branch argmax matches CPU oracle. +- [ ] **Step 2: Implement online forward + argmax extraction.** Replace half the `todo!()`. +- [ ] **Step 3: Run, commit.** + +### Task 4.4: Implement target-net forward pass on next_states + `compute_expected_q` + +**Files:** Modify `compute_q_next_target_bootstrap` — implement target-net pass: +- Use existing `target_params_buf` machinery (mirror of `submit_forward_ops_main`'s Pass 2 target-net branch on next_states). +- Call `compute_expected_q` with target-net distributional outputs to produce `[B, 13]` per-action E[Q]. +- Per B-DD3: gather per-branch `Q_target[b, a*_online_branch]` and SUM across 4 branches → scalar `q_next_per_sample[B]`. + +- [ ] **Step 1: Failing GPU oracle test** — small batch; verify gathered scalar matches CPU oracle (T8 from spec). +- [ ] **Step 2: Implement target forward + gather + sum.** Replace second half of `todo!()`. +- [ ] **Step 3: Run, commit.** + +### Task 4.5: Replace `todo!()` with full implementation + smoke + +- [ ] **Step 1: Failing test** — `#[test] compute_q_next_target_bootstrap_no_todo` asserts the method body does NOT contain `todo!`. +- [ ] **Step 2: Confirm Tasks 4.3 + 4.4 fully replaced the `todo!()`.** +- [ ] **Step 3**: `cargo check --workspace` + `cargo test -p ml --lib` PASS. Method is still dead code (no production caller); covered only by GPU oracle tests T8 + T11. +- [ ] **Step 4: Commit.** + +> **PARTIAL-REFACTOR NOTE**: Phase 4 leaves `compute_q_next_target_bootstrap` as dead code. The atomic-refactor rule applies *between* the Phase 4 dead method and the Phase 5 production replacement: NO L40S DISPATCH between Task 4.5 and Task 5.1. Phase 4 is locally atomic (the dead method itself is fully wired with no unfilled stubs); Phase 5 makes it live in lockstep with the q_next replacement. + +--- + +## Phase 5 — B-leg q_next replacement at line 4143 + PopArt EMA reset gate (atomic) + +### Task 5.1: Replace `q_next = dtod_clone_f32_native(...)` with `compute_q_next_target_bootstrap` + +**Files:** Modify `gpu_experience_collector.rs:4143` — change the line to call the new method. + +```rust +// Before (line 4143): +// let q_next = dtod_clone_f32_native(&self.stream, &self.rewards_out, total, "q_next_bootstrap")?; +// After: +let target_params_ptr = ???; // pass from caller via context — see Task 5.2 plumbing +let online_params_ptr = ???; // ditto +let q_next = self.compute_q_next_target_bootstrap(&next_states, target_params_ptr, online_params_ptr, total)?; +``` + +- [ ] **Step 1: Failing GPU oracle test (T10)** — converged-checkpoint replay buffer; verify `avg(|TD_error|) ratio post-fix / pre-fix ∈ [0.5, 5.0]`. +- [ ] **Step 2: Implement** — replace line 4143, plumb pointers (see Task 5.2 if not already wired). +- [ ] **Step 3: Run, verify PASS.** +- [ ] **Step 4: Commit (THIS COMMIT closes Phase 4's partial-refactor state — B-leg atomic).** + +### Task 5.2: Pointer plumbing — pass target_params_ptr + online_params_ptr to experience collector + +**Files:** +- `gpu_dqn_trainer.rs:8847 target_params_buf_device_ptr()` and `gpu_dqn_trainer.rs:11426 target_params_buf_ptr()` — both already exist. +- `gpu_experience_collector.rs` — extend the experience-collection invocation to receive both pointers via the existing context struct. + +If pointers already flow (verify): no plumbing needed. + +If not: add pointer fields to the experience-collection invocation context. + +- [ ] **Step 1: Failing test** — `#[test] target_params_ptr_plumbed_to_collector` asserts the call site has access. +- [ ] **Step 2: Plumb if needed; otherwise verify.** +- [ ] **Step 3: Run, commit (or skip if no change).** + +> _DESIGN NOTE_: this task may be a no-op if the pointers already flow naturally via the trainer-collector context. Verify in Phase 4 close-out before approving this task. + +### Task 5.3: PopArt EMA reset gate (B-DD11) + +**Files:** +- Modify: `training_loop.rs` per-epoch loop start — read slot 497 `POPART_RESET_FLAG`. If `1.0` (sentinel for first SP18 epoch), launch a one-shot kernel that resets PopArt slot 63 EMA to identity (μ=0, σ=1) and writes 0.0 to slot 497. +- NEW: `popart_reset_kernel.cu` (small one-shot kernel; mirror existing `state_reset_registry` reset patterns). +- Modify: `build.rs` — add cubin manifest entry. + +- [ ] **Step 1: Failing test** — synthetic 2-epoch run: epoch 0 sees slot 497=1.0, PopArt EMA resets; epoch 1 sees slot 497=0.0, PopArt EMA continues normally. +- [ ] **Step 2: Implement reset kernel + launcher + per-epoch gate.** +- [ ] **Step 3: Run, commit.** + +### Task 5.4: Phase 5 wire-up audit + pearl candidate + +- [ ] **Step 1: Re-run audit script** — confirm: + - `q_next_bootstrap` literal: 0 hits (renamed to `q_next_target`). + - `dtod_clone_f32_native(&self.stream, &self.rewards_out, ..., "q_next` literal: 0 hits. + - `compute_q_next_target_bootstrap` invocations: exactly 1 (line 4143). + - `target_params_buf` consumers in `gpu_experience_collector.rs`: 1 new (the new method). + - `todo!` markers in `gpu_experience_collector.rs`: 0. + +- [ ] **Step 2: Pearl draft** — `pearl_target_net_q_during_experience_collection_for_correct_credit_assignment.md`. Lands post-Phase 9 if smoke passes. + +- [ ] **Step 3: Append pearl draft to plan as "Pearl candidate post-Phase 5".** + +--- + +## Phase 6 — B-leg behavioral test suite (T8–T13) + +### Task 6.1: T8 — CPU oracle for DDQN per-branch argmax bootstrap + +(Spec §"Behavioral Test Predictions" T8.) + +- [ ] **Steps 1-4: Failing test → implement → run → commit.** + +### Task 6.2: T9 — GPU oracle parity with CPU (B-leg) + +- [ ] **Steps 1-4.** + +### Task 6.3: T10 — TD-error magnitude before vs after fix + +(B-DD9 ratio gate.) + +- [ ] **Steps 1-4.** + +### Task 6.4: T11 — Target-net forward consistency + +Compare new `compute_q_next_target_bootstrap` output to existing loss-time target-net forward on same inputs. Bit-identical assertion. + +- [ ] **Steps 1-4.** + +### Task 6.5: T12 — PER priority distribution post-fix + +(B-DD4 sanity.) + +- [ ] **Steps 1-4.** + +### Task 6.6: T13 — Graph-capture compatibility + +Capture `exp_fwd_graph` with new path, replay 10×, verify per-replay determinism. + +- [ ] **Steps 1-4.** + +--- + +## Phase 7 — D-leg behavioral test suite (T1–T7) + integration tests + +### Task 7.1: T3 — Hold reward distribution on uniform-edge synthetic market +### Task 7.2: T4 — Hold reward distribution on positive-edge synthetic market +### Task 7.3: T5 — Cap producer p99 oracle (full) +### Task 7.4: T6 — End-to-end on converged b5gmp checkpoint +### Task 7.5: T7 — Compile-clean post-deletion + +(All from spec; v1 plan content preserved.) + +- [ ] **Steps 1-4 each.** + +### Task 7.6: T-extra D — Sign convention regression test +### Task 7.7: T-extra D — Distribution-width comparison test + +(From v1 plan §"Phase 4".) + +- [ ] **Steps 1-4 each.** + +### Task 7.8: Integration test — D-leg + B-leg combined on converged checkpoint + +NEW: load a converged b5gmp checkpoint; run 1 forward step with BOTH D-leg reward and B-leg q_next active; assert: +- Hold opportunity-cost fires ≥ 80% of Hold bars. +- TD-error magnitude ratio post-fix / pre-fix ∈ [0.5, 5.0]. +- v_share trend at single-step is not monotonically increasing. +- No NaN. + +- [ ] **Steps 1-4.** + +--- + +## Phase 8 — Combined HEALTH_DIAG observability + +### Task 8.1: D-leg HEALTH_DIAG producers — HOLD_REWARD_DECOMP_DIAG + HOLD_OPP_COST_FIRE_RATE_EMA + +**Files:** Extend `reward_decomp_diag_kernel.cu` (from Phase 0) to write slots [485, 486]. + +- [ ] **Steps 1-4.** + +### Task 8.2: B-leg HEALTH_DIAG producer — `td_bootstrap_diag_kernel.cu` + +**Files:** NEW kernel. Computes: +- TD-error magnitude EMA (slot 493) +- q_next_target p99 (slot 494) +- |q_next_target − rewards| p99 (slot 495) +- v_share slope across last 4 epochs (slot 496) + +Block tree-reduce. Cubin manifest entry. + +- [ ] **Steps 1-4.** + +### Task 8.3: Per-epoch combined HEALTH_DIAG emit + +**Files:** Modify `training_loop.rs`: +``` +HEALTH_DIAG[N]: sp18_d [pos_cap=X neg_cap=Y opp_decomp=Z fire_rate=W] +HEALTH_DIAG[N+1]: sp18_b [td_err=X q_next_p99=Y q_diff_p99=Z v_share_slope=W popart_reset=F] +``` + +- [ ] **Steps 1-4.** + +### Task 8.4: Phase 8 wire-up close-out + +- [ ] **Step 1: Audit script** — confirm new producers wired exactly once each. +- [ ] **Step 2**: `cargo check --workspace` + `cargo test -p ml --lib` PASS. +- [ ] **Step 3**: Local smoke runs without NaN. + +--- + +## Phase 9 — 5-epoch L40S smoke validation + +### Task 9.1: Push + dispatch + monitor + +**REVIEWER-EXECUTED.** + +- [ ] **Step 1: Reviewer pushes branch** + +```bash +git push origin feat/sp18-reward-shape +``` + +- [ ] **Step 2: Reviewer dispatches L40S 5-epoch smoke** + +```bash +./scripts/argo-train.sh --model dqn --epochs 5 --label sp18-combined-smoke-v2 +``` + +- [ ] **Step 3: Smoke gate criteria** (full table from spec § "Smoke Gates" — D-leg + B-leg + combined; bold gates are halt-criteria). + +- [ ] **Step 4: Document outcome** + +Append "Phase 9 smoke evidence" subsection with raw HEALTH_DIAG lines + pass/fail per gate. + +If any **bold** gate fails: STOP. Convene human review before Phase 10. + +--- + +## Phase 10 — 30-epoch full validation + +### Task 10.1: Full dispatch + analysis + +**REVIEWER-EXECUTED.** + +- [ ] **Step 1: 30-epoch dispatch** + +```bash +./scripts/argo-train.sh --model dqn --epochs 30 --label sp18-combined-full-v2 +``` + +- [ ] **Step 2: Full gate criteria** (spec § "Smoke Gates" Phase 10 column). + +- [ ] **Step 3: Pass/fail decision + memory pearl land** + +If pass: +- Land `pearl_distributional_shape_match_for_action_neutrality.md`. +- Land `pearl_target_net_q_during_experience_collection_for_correct_credit_assignment.md`. +- Land `pearl_complementary_fixes_no_deferrals.md` (Pearl 4 from spec). + +If D-leg passes but B-leg gates fail (or vice versa): land partial pearls, open follow-up plan for the failing leg. + +If both fail: open SP19 follow-up plan reviewing the synergy hypothesis. + +- [ ] **Step 4: Update active project state file** + +Append to `/home/jgrusewski/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/`: +- `project_sp18_combined_close_out.md` (PASS) or `project_sp18_partial.md` / `project_sp19_required.md` (FAIL). + +--- + +## Self-review checklist + +- [x] Every task is bite-sized TDD: failing test → implement → pass → commit. +- [x] Only ONE `todo!()` marker (Task 4.2 explicit exception, justified, replaced atomically by Task 4.5). +- [x] No `// ...` ellipses in code blocks (where I omit context, file:line citation provided). +- [x] Phase 1 atomic-deletion explicit: SP13/SP16 cost-application + producer + host-controller deleted in 3 commits all in same branch BEFORE Phase 2. +- [x] Phase 2 atomic-introduction: 3 site migrations + device fn + cubin entry + test wrapper land in single commit (Task 2.2). +- [x] Phase 4–5 atomic: target-net forward path lands as dead code in Phase 4, q_next replacement in Phase 5 — NO L40S DISPATCH between Task 4.5 and Task 5.1. +- [x] Pre-dispatch consumer audit script (Task 1.1) surfaces missed consumers BEFORE migration. +- [x] Audit script promoted to pre-commit hook (Task 1.2). +- [x] All adaptive bounds ISV-driven. +- [x] All reductions block-tree-reduce; no atomicAdd. +- [x] All CPU↔GPU transfers mapped-pinned (tests not exempt). +- [x] Smoke gates anchor to empirical baselines: vtj9r HD2/HD3 (91.58, 78.85, 31.6%), vtj9r HD0→HD4 v_share trajectory (0.33→0.83), b5gmp Fold-1 final (55.55, 52.6%), pre/post-fix TD-error ratio. +- [x] DESIGN DECISION markers in spec: 8 D-leg + 13 B-leg = 21. All flagged for human review. +- [x] Pearl candidates flagged at end of Phase 0/2/3/5/9 for human review. +- [x] KILL CRITERION in Phase 0 is concrete and disqualifying for both legs. +- [x] B-leg consumer audit (Phase 1) surfaces: + - Single production td_lambda launch site. + - PER recomputed at training time (NO migration). + - Replay buffer doesn't store q_next (NO schema migration). + - C51 loss target-Q already correct (NO migration). + - target_params_buf Polyak-tracked, ready. + - PopArt slot 63 sites catalogued for B-DD11 reset. +- [x] No "speculative" gates — every Phase 9/10 threshold traces to a measured SP-chain run or pre-fix baseline. +- [x] No partial refactor between any "delete" and "introduce" pair (1.5 → 2.2; 4.5 → 5.1). +- [x] `feedback_no_concurrent_agents_shared_tree`: branched onto `feat/sp17-dueling` HEAD `0c57a5a31`. + +--- + +## Memory pearl candidates (drafts) + +(See spec § "Memory pearl candidates" — Pearls 1–5 identical there, replicated here for reviewer convenience.) + +### Pearl 1 (post-Phase 1+2): `pearl_replace_reactive_controller_with_structural_signal.md` +### Pearl 2 (post-Phase 5 if synergy holds): `pearl_target_net_q_during_experience_collection_for_correct_credit_assignment.md` +### Pearl 3 (post-Phase 9 if smoke passes): `pearl_distributional_shape_match_for_action_neutrality.md` +### Pearl 4 (post-Phase 9 if combined passes): `pearl_complementary_fixes_no_deferrals.md` +### Pearl 5 (post-Phase 9 if D-leg-only passes but B-leg fails): `pearl_reward_shape_alone_insufficient_when_bootstrap_biased.md` + +--- + +## Execution handoff + +**Total estimated commits: 55 tasks × ~1 commit each = ~55 commits across feat/sp18-reward-shape.** + +Estimated wall time: + +- **Combined SP18 (D-leg + B-leg)**: + - Pre-Phase + Phase 0 (diagnostic + KILL CRITERION + B-leg profiling): ~1.5 days + - Phase 1 (audit + pre-commit hook + atomic SP13/SP16 deletion): ~0.75 day + - Phase 2 (D-leg device fn + 3 sites): ~1 day + - Phase 3 (D-leg cap producer): ~0.5 day + - Phase 4 (B-leg target-net forward additive): ~1.5 days + - Phase 5 (B-leg q_next replacement + PopArt reset): ~1 day + - Phase 6 (B-leg behavioral tests): ~1 day + - Phase 7 (D-leg behavioral tests + integration): ~1.5 days + - Phase 8 (HEALTH_DIAG): ~0.5 day + - Phase 9 smoke: ~6 hours dispatch + analysis + - Phase 10 full: ~24 hours dispatch + analysis + + **Total: ~9 days plan + ~1.5 days dispatch = ~10–11 days end-to-end.** + +- **Sequential D-then-B (rejected counterfactual)**: + - SP18 D-leg only: ~5 days plan + 1.5 days dispatch = 6.5 days. + - SP19 B-leg-on-D-baseline: ~5 days plan + 1.5 days dispatch = 6.5 days. + - **Total: ~13 days end-to-end** + can't observe synergy in either run. + + **Combined saves ~2–3 days AND surfaces synergy hypothesis in one validation cycle.** + +Reviewer must: +1. Resolve **DD1–DD8 + B-DD1–B-DD13** before subagent dispatch. +2. Run Pre-Phase Task PP.1 (worktree creation) manually. +3. Run Pre-Phase Task PP.5 (B-leg profiling) manually; halt if collection-time headroom < 30%. +4. Run all `argo-train.sh` dispatches manually (Phases 0, 9, 10). +5. Land/reject memory pearls at end of each phase. +6. **If Phase 0 KILL CRITERION returns INCONCLUSIVE for either leg**, convene review. +7. **If Phase 9 smoke fails any bold gate**, convene review BEFORE Phase 10. + +Subagent must: +1. Hold to atomic-consumer-migration discipline: + - Tasks 1.3–1.5 + 2.2 land in same branch with NO L40S dispatch in between. + - Tasks 4.5 + 5.1 land in same branch with NO L40S dispatch in between. +2. Stop and ask if Phase 1 audit surfaces any consumer not in spec's locked list (D-leg or B-leg). +3. Stop and ask if any kernel under modification has a graph-capture interaction not visible in the consumer audit (per `pearl_no_host_branches_in_captured_graph`). +4. Stop and ask if `compute_sp18_hold_opportunity_cost` device fn requires CUDA Graph capture changes (it shouldn't — it's a pure device function inside `experience_kernels.cu` which is already graph-captured). +5. Stop and ask if `compute_q_next_target_bootstrap` adds host branches inside the captured `exp_fwd_graph` region (it must NOT — pure kernel launches only). +6. Treat any test failure as a halt, not a flaky-test retry. +7. Replace the single `todo!()` marker (Task 4.2) atomically in Task 4.5; verify post-Phase 5.4 audit shows 0 `todo!` hits in `gpu_experience_collector.rs`. diff --git a/docs/superpowers/specs/2026-05-08-sp18-reward-shape-hold-attractor-design.md b/docs/superpowers/specs/2026-05-08-sp18-reward-shape-hold-attractor-design.md new file mode 100644 index 000000000..908069a87 --- /dev/null +++ b/docs/superpowers/specs/2026-05-08-sp18-reward-shape-hold-attractor-design.md @@ -0,0 +1,668 @@ +# SP18 — Combined Hold-Attractor Fix: Reward-Shape D-leg + TD(λ) `Q(s')` B-leg (Design Spec) + +> **Status: v2 — UPGRADED to combined SP18+SP19 scope. DRAFT FOR HUMAN REVIEW.** +> v1 deferred Approach B (TD(λ) `Q(s')` bootstrap) to SP19. **That deferral is overridden** per the +> "no deferrals for complementary fixes" call. Both legs land atomically in one branch / one plan. +> Spec must precede plan per the superpowers brainstorming pattern. +> Plan: `docs/superpowers/plans/2026-05-08-sp18-reward-shape-hold-attractor.md`. + +--- + +## Why combined scope (and why it is safe) + +The vtj9r terminated 5-epoch L40S smoke (sha `0c57a5a31`, terminated 2026-05-09 at HD4 across both folds with confirmed −14%/epoch decay, see Problem Statement) leaves three independently-acting causal mechanisms behind the persistent Q(Hold) attractor: + +| Mechanism | Fixed by | Refactor scope | +|---|---|---| +| 1. Reactive control loop (controller bandwidth < policy convergence speed) | **D-leg** (Approach D) | reward kernels — `experience_kernels.cu` per-bar Hold sites + new `compute_sp18_hold_opportunity_cost` device fn + atomic delete of `hold_cost_scale_update_kernel` chain | +| 2. Distributional shape mismatch (Hold tight, Long/Short wide) | **D-leg** (Approach D) | same as above (raises Hold's Q-target distribution width into Long/Short range) | +| 3. TD(λ) self-bootstrap (`q_next = rewards`, never sees `Q(s')`) | **B-leg** (Approach B) | `td_lambda_kernel` consumer — `gpu_experience_collector.rs:4143-4170` + new target-net forward pass on next-states during experience collection (target-net params already exist via `target_params_buf`) | + +**Refactor scopes have empty intersection.** D-leg writes to `r_micro` / `r_opp_cost` reward components inside `experience_kernels.cu`. B-leg writes to the `q_next` argument fed into `td_lambda_kernel` (one launch site at `gpu_experience_collector.rs:4154`). Per `feedback_no_partial_refactor`, atomicity is required *within* each leg's consumer set; *between* the legs there is no shared buffer or contract that breaks if they ship in separate commits inside the same branch. + +**Synergy hypothesis (testable):** B-leg's correct credit assignment makes D-leg's gradient signal meaningful by construction. Under `q_next = rewards` self-bootstrap, the TD-error `δ_t = r_t + γ × q_next_(rewards) − Q(s_t, a_t)` collapses to `δ_t ≈ r_t + γ × r_{t+1} − Q(s_t, a_t)` — independent of which action was actually taken. With `q_next = Q_target(s_{t+1}, a*_{online})` the TD-error becomes `δ_t = r_t + γ × Q_target(s', argmax_a Q_online(s', a)) − Q(s_t, a_t)` — the bootstrap term depends on the next-state action distribution, which depends on the previous action choice. **This may also explain SP17's v_share elevation pattern** (vtj9r showed v_share 0.33 → 0.83 across HD0–HD4): when the bootstrap is action-independent, V absorbs the bias because the per-action centred A has no signal to learn from. Combined SP18 tests this. + +--- + +## Problem Statement + +Across SP11→SP17 the model develops a persistent **Q(Hold) attractor**: the directional Q-table prefers Hold to Long/Short despite training rewards that should favour active trading. SP17 (mean-zero advantage centering) eliminated `V`/`A` non-identifiability as a confounder, so the bias is now expressed cleanly through `A_centered_dir[Hold] > A_centered_dir[Long]` rather than `V` masquerading as action preference. **Yet vtj9r `v_share` rose from 0.33 (HD0) to 0.83 (HD4)** — V absorbing what A should have absorbed. Combined SP18 hypothesises this is the rewards-only TD(λ) bootstrap's signature. + +**Terminated smoke `train-multi-seed-vtj9r` (SP17 5-epoch L40S, sha `0c57a5a31`, killed 2026-05-09)**: + +| Epoch | val Sharpe | Hold% | dir_entropy | A_var_dir EMA | v_share | Q(Hold)−Q(Long) | +|---|---|---|---|---|---|---| +| HD0 | 80.70 | 16.1% | (high) | 0.0023 | 0.33 | ≈ 0.05 | +| HD2 | **91.58** | 18.2% | 1.07 | (stable) | 0.51 | ≈ 0.06 | +| HD3 | 78.85 | 31.6% | 0.996 | 0.0021 | 0.71 | 0.05–0.06 | +| HD4 | (term) | (term) | (term) | (term) | 0.83 | (term) | + +Trajectory: SP17 delays the attractor by 2 epochs and lifts the peak Sharpe by ~13% relative to HD0, but **Hold% climbs +13.4 pts from HD2 to HD3 (peak to first decay)** and **v_share climbs monotonically across all 5 epochs**. Slope ≈ −14% Sharpe / epoch matches `train-multi-seed-b5gmp` Fold 0 HD2→HD3 (pre-SP17 baseline). The architecture is healthy (`A_var_dir` stable, no `V` dominance *at HD0–HD2*), so the gradient-update process is **correctly** learning Hold-as-preferred under the current reward + the v_share elevation pattern is a *symptom* of the rewards-only TD(λ) bootstrap. + +**Conclusions:** +- The reward function is asking the model to prefer Hold (D-leg fix). +- The TD(λ) bootstrap is structurally biased toward whichever action's reward distribution looks closest to its long-run accumulation, regardless of state-value, eroding V/A separation across epochs (B-leg fix). + +--- + +## D-leg Empirical Reward Audit (against current code) + +Audit performed against worktree HEAD `0c57a5a31` (`feat/sp17-dueling`). All file:line citations live. + +### What Hold rewards today + +| Bar regime | Hold-action reward path | Magnitude (current ISV) | +|---|---|---| +| Voluntary trade close (segment_complete && !trail) on a Hold pick | `r_popart = capped_pnl − shaping × min_hold_penalty` (`experience_kernels.cu:3315`); plus per-bar Hold cost `−ISV[380] × ISV[461]` subtracted from `base_reward` (line 3166) | small, P&L-driven | +| Trail-fire close while picking Hold | `r_trail = capped_pnl` (line 3230); same Hold cost subtraction | rare | +| Per-bar mid-trade (positioned, no event) on Hold | `r_micro = 0.0f − ISV[380]×ISV[461]` (line 3617–3643) | `0 − HOLD_COST_BASE×scale = −0.005×[1..25] = −0.005..−0.125` | +| Per-bar Flat-Hold (no position) | `r_opp_cost = 0.0f − ISV[380]×ISV[461]` (line 3697–3716) | same as above, `−0.005..−0.125` | + +### What Long/Short rewards on the same bars + +- Per-bar mid-trade: `r_micro = 0.0f` (no Hold cost subtracted; line 3617). **Identical to a Hold mid-trade *in dollar terms*** when the per-bar Hold cost is at its floor (`scale=1.0`). +- Per-bar Flat-Long entering: not a different per-bar branch — opening a new position is `entering_trade=true` and runs through the segment_complete machinery on its eventual close, paying `compute_lump_sum_opp_cost = −shaping × holding_cost_rate × |position| × hold_time` (`trade_physics.cuh:590`) **at exit**. +- Trade close P&L: capped asymmetrically at `[REWARD_NEG_CAP=−10, REWARD_POS_CAP=+5]` (`state_layout.cuh:337–338`), with adaptive ISV[452]/ISV[453] producers tracking p99(winning_returns) × 1.5 (`reward_cap_update_kernel`). + +### Numerical reality check + +Static base: `HOLD_COST_BASE = 0.005` (`sp13_isv_slots.rs:73`), gain `5.0`, floor `0.5×base`, ceil `5.0×base` ⇒ host controller emits cost in `[0.0025, 0.025]` per bar. +SP16 P2 multiplier: `ISV[461] ∈ [1, 25]` driven by `(observed_hold_rate − 0.20) / 0.20` overrun. With Wiener-α non-stationary floor 0.4 (post-T3, `min_hold_temperature_update_kernel.cu:121`). + +**Effective per-bar Hold penalty range: `[0.0025, 0.625]`** (when `scale=25`). + +Per-bar reward magnitudes during a typical trade (vol-normalized return, capped at +5/−10): `O(0.01–0.1)`. So at the **floor** of the SP16 P2 ramp the Hold cost is `~5%–25%` of a typical per-bar reward — competitive but small. At the **ceiling** it dominates. The ramp depends entirely on `hold_rate_observed_ema` exceeding 0.20. + +### Where the D-leg asymmetry actually lives + +Three concrete asymmetries between Hold and Long/Short under the present reward shape: + +1. **Cost density.** Hold pays a per-bar cost (line 3637–3643, line 3710–3716). Long/Short pay a lump-sum carrying cost only at trade close. With `td_lambda=0.9` (`config.rs:1427`) and `gamma≈0.95`, the integrated and per-bar cost are mathematically equivalent only at γ=1; at γ<1 the per-bar Hold pricing **discounts less per dollar** than a deferred close-event cost. + +2. **Reward-cap clipping.** Big winners get `+5`, big losers get `−10` (2:1 Kahneman). Hold *participates in the same clamp* via `compute_asymmetric_capped_pnl` only if a position is open. **Flat-Hold reward is dominated by `−ISV[380]×ISV[461]` only**, never by a positive P&L term — and that subtraction is bounded by 0.625/bar. Long/Short collect occasional large bounded wins; Hold collects a steady stream of small bounded penalties. **The Q-target distribution for Hold is a tight low-variance band; the Q-target distribution for Long/Short is a wide high-variance distribution centered near 0**. CQL pessimism plus distributional collapse rewards low-variance branches automatically. + +3. **Bellman-target inflation under TD(λ) self-bootstrap.** This is the B-leg's territory; carved out below. + +--- + +## B-leg TD(λ) Self-Bootstrap Audit (against current code) + +Audit performed against worktree HEAD `0c57a5a31`. All file:line citations live. + +### Where the bug lives — single launch site + +`gpu_experience_collector.rs:4135-4174` is the entire production TD(λ) wiring: + +```rust +// Line 4135-4143: +let td_lambda = config.td_lambda; // = 0.9 from config.rs:1427 +let max_trace = config.max_trace_length; +if td_lambda > 0.01 && max_trace > 1 { + let raw_rewards_td = dtod_clone_f32_native(&self.stream, &self.rewards_out, total, "raw_rewards_td")?; + let raw_dones_td = dtod_clone_f32_native(&self.stream, &self.done_out, total, "raw_dones_td")?; + let q_next = dtod_clone_f32_native(&self.stream, &self.rewards_out, total, "q_next_bootstrap")?; + // ^^^^^^^^^^^^^^^^ ← THE BUG + // Line 4146-4170: launch td_lambda_kernel with (raw_rewards_td, raw_dones_td, q_next, rewards, dones, ...) +} +``` + +The kernel's own header explicitly licenses this: + +``` +nstep_kernel.cu:95-96 + * Requires bootstrapped Q(s') values. When not available, pass raw_rewards + * as q_next for a self-bootstrap approximation. +``` + +### Consumers of the kernel's output + +`td_lambda_kernel` *output* writes to `rewards` (`out_rewards`, line 4158) and `dones` (`out_dones`, line 4159) — these are the SAME buffers passed into the TD(λ)-trained replay batch. So the consumers of `rewards` post-TD(λ) are: + +1. **Replay buffer push** — `rewards` flows into the batch that is appended to the per-fold replay buffer. (One push per experience-collection cycle.) +2. **PER initial priorities** — NOT computed at experience-collection time. PER priorities are recomputed at *training* time in `agent.update_priorities_from_td(...)` from current TD-error (`config.rs:341`, `fused_training.rs:2963`). Initial sampling priority is `max_priority` heuristic on first-store. **No buffer schema migration is needed for B-leg.** +3. **Loss-time TD-error** — at training time, the C51 / IQN loss kernels read `rewards` from the sampled replay batch and combine with target-net `Q(s')` (already done correctly via `submit_forward_ops_main`'s Pass 2 target-net forward on next_states + `compute_expected_q` projection through `c51_loss_batched`). + +**The bug is confined to the experience-collection-time TD(λ) accumulation.** The C51 loss kernel's distributional Bellman target site (`c51_loss_kernel.cu:188 block_bellman_project_f`) is **already centered correctly** because the loss-time `target_probs` come from a target-net forward pass that already exists at loss time. Mirror the SP17 finding that `c51_loss` was already centered — same pattern here for the *target-Q at loss time*. + +The defect is that the *TD(λ) returns themselves* (the `R^λ` written into `rewards` during experience collection) are computed with rewards-as-q_next, biasing intermediate λ-blended terms toward `E[reward]` rather than `E[Q(s')]`. The loss kernel's outermost Bellman bootstrap is already a target-net Q, but the inner λ-weighted return is rewards-only. **This is the right surgical site to fix**: replace `q_next` at line 4143 with a target-net `compute_expected_q` evaluation on `next_states`. + +### Inputs needed to compute correct `Q(s')` during experience collection + +To wire `q_next = Q_target(s', argmax_a Q_online(s', a))` (DDQN per **B-DD2/B-DD3**) at experience-collection time: + +| Input | Where it lives | Available at collection time? | +|---|---|---| +| `target_params` (target-net weights) | `gpu_dqn_trainer.rs:5359 target_params_buf` | YES — Polyak-tracked, populated since SP14 P4.T1B-iv | +| `next_states` (vol-normalized post-action market state) | built at `gpu_experience_collector.rs:4182 build_next_states_f32` | YES — built in same scope, line 4182, AFTER td_lambda runs today (must be moved BEFORE td_lambda for B-leg) | +| `online_params` (online-net weights for DDQN argmax) | exposed via `params_buf_device_ptr()` | YES | +| Forward-pass machinery (encoder + GRN + VSN + dueling head) | all kernels exist; orchestration lives in `submit_forward_ops_main` (Pass 2 target-net branch already runs target on next_states at loss time) | NEEDS new wrapper for collection-time invocation | +| `compute_expected_q` for the projection | `experience_kernels.cu:4572` | YES — already takes params, inputs, writes per-action E[Q] | + +### Replay buffer schema check + +Searched `agents/`, `trainers/`, `cuda_pipeline/` for `replay_buffer` writes. The push site stores `(state, next_state, reward, done, priority_initial=max_priority)`. **`q_next` is NEVER stored in the buffer** — it's a transient scratch in the experience-collection cycle that gets consumed by `td_lambda_kernel` and then discarded. **Confirmed: no replay buffer schema migration needed.** + +### Graph-capture check + +`exp_fwd_graph` capture region (`gpu_experience_collector.rs:4651–4671`) covers the experience-collection forward pass for action selection. The TD(λ) launch at line 4154 is *outside* the capture region (it runs after the per-step inner loop completes, after `next_states` is built). Adding a target-net forward pass before line 4154 is graph-captureable as long as it lives in the same post-loop scope as today's TD(λ) launch — purely kernel launches, no host branches per `pearl_no_host_branches_in_captured_graph`. + +### PopArt interaction + +PopArt slot 63 normalizes Q-targets at *loss time* (after the trained replay batch is sampled). With the B-leg fix: + +- Pre-fix: trained `R^λ` ≈ rewards-weighted-cumsum; PopArt sees `R^λ + γ^N × Q_target(s_{t+N})` where the inner `R^λ` is rewards-mean-anchored. +- Post-fix: trained `R^λ` ≈ target-Q-bootstrapped at every λ-step; PopArt sees the same outer expression but the inner `R^λ` is now Q-mean-anchored (which is the magnitude PopArt was *designed* to track). + +**Risk**: at the moment of switching (HD0 of first SP18 run), PopArt's existing EMA state was tuned to the rewards-distributed `R^λ`. Switching to Q-distributed `R^λ` shifts the input distribution at one step. PopArt has its own α-EMA that adapts; the question is whether it adapts within HD0 or causes a transient. **B-DD11**: per-fold reset of PopArt slot 63 EMA at the SP18 deployment boundary. (Cheap insurance, mirrors SP14 fold-boundary reset patterns.) + +--- + +## Approach Considered + +### Approach A — Symmetric trade-cost compensation (REJECTED) + +Charge Hold a structural `position × half_spread × decay` per bar. Rationale: Long/Short already pay an entry/exit half-spread; a Hold action that maintains an open position is also "consuming" the time-decay aspect of having committed capital. + +**Why rejected**: SP16 P2 (`hold_cost_scale_update_kernel`) is exactly this in spirit — adaptive multiplier on the per-bar Hold cost driven by observed Hold-rate overrun. The mechanism failed because it's reactive (closed-loop control with Wiener-α latency), not because it lacked the right semantic concept. Layering another version of the same fix would be the canonical `feedback_no_quickfixes` violation. + +### Approach B — Eligibility traces / TD(λ) `Q(s')` bootstrap fix (ACCEPTED — combined as B-leg) + +**v1 deferred this to SP19. v2 accepts as the B-leg of combined SP18.** See "Why combined scope" above and the Architecture Changes B-leg section below. + +Switching to a real `Q(s')` bootstrap (using the target-net's expected-Q on the realized next-state) gives n-step / TD(λ) returns the correct credit-assignment shape. This is exactly what `pearl_event_driven_reward_density_alignment` recommends ("eligibility traces, n-step returns, HER… are RL-architectural credit-assignment tools, not per-step reward hacking"). + +Wiring real `Q(s')` into `td_lambda_kernel` requires: +- adding a target-network forward pass on `next_states` *during experience collection* (currently only happens in the loss kernel) +- piping the resulting expected-Q back into the `q_next` argument +- accepting ~30% collection-time slowdown (extra forward pass; **B-DD6** — H100 has headroom; profile in Pre-Phase to confirm before committing) +- per-fold reset of PopArt EMA (**B-DD11**) + +**No replay buffer schema migration.** **No PER priority migration** (PER recomputes from current TD-error at training time). **No C51 loss kernel migration** (already correctly using target-Q). + +### Approach C — Bellman target redesign (REJECTED) + +Center per-step reward at running mean before it enters the Bellman target. Concrete: `r_centered = r − reward_mean_ema; Q-target uses r_centered`. + +**Why rejected**: collides with PopArt (slot 63), the SP11 reward-subsystem controller (slots 340–346), the SP15 DD-asymmetric-reward producer, and the SP12 asymmetric `+5/−10` cap. Each consumes the un-centered reward; centering would either need to flow through *all* of them or live as a post-hoc bias subtraction at the Bellman site (which breaks the SP15 reward-axis "second-pass composer" contract). + +### Approach D — Asymmetric Hold-vs-position reward bound (ACCEPTED — combined as D-leg) + +Per `pearl_audit_unboundedness_for_implicit_asymmetry`, the SP12 `+5/−10` cap restored loss aversion to *position* P&L. **Hold has no equivalent.** A natural per-bar Hold reward is *symmetric around zero* (steady tiny negatives only when the controller ramps). Make the Hold per-bar reward asymmetric by adding a **bounded reinforcement** for the realized opportunity-cost of forgone directional moves, scaled by a controller-driven signal that reflects whether the data *currently* supports trading. + +Concretely (atomic): + +1. **Hold-anchored reward decomposition**. Replace per-bar Hold cost subtraction at the 3 sites with a per-bar Hold reward `r_hold` that is the **sign-flipped** vol-normalized realized `tgt[1]` next-bar log-return — bounded by the same `compute_asymmetric_capped_pnl` machinery the segment-complete path uses. This **lifts Hold's reward distribution into the same dynamic range as Long/Short**. + +2. **Asymmetric Hold reward cap**. New ISV pair `[483..485)` — `HOLD_REWARD_POS_CAP_INDEX` and `HOLD_REWARD_NEG_CAP_INDEX` — driven by a producer mirroring `reward_cap_update_kernel` but tracking `p99(|step_ret|)` of *Long/Short* trade closes only. + +3. **Anchored at vol-normalized log-return only**. Per `pearl_one_unbounded_signal_per_reward`: the unbounded multiplicand is `vol_normalized_return`; everything else is bounded. + +**Why D over the others** (D-leg only — see "Why combined scope" for the multi-leg argument): addresses mechanism not magnitude; structural not reactive; reuses `compute_asymmetric_capped_pnl` and the `reward_cap_update_kernel` pattern; minimal new code surface; does not touch PopArt, SP11 controller, SP15 composer, or PER. + +--- + +## Architecture Changes — D-leg (assuming DD1=combined, DD2=a, DD3=b, DD4=a, DD5=b, DD6=as-specified, DD7=c, DD8=Phase 3.1) + +### CUDA changes (D-leg) + +**`crates/ml/src/cuda_pipeline/experience_kernels.cu`** + +1. At the 3 Hold-cost subtraction sites: + - line 3162–3166 (segment_complete branch): replace the 4-line `if (dir_idx == DIR_HOLD && isv_signals_ptr != NULL)` block with a single call into `compute_sp18_hold_opportunity_cost(...)` (see point 2). Output flows into `r_micro` (additive, returns 0 if not Hold). + - line 3637–3643 (per-bar mid-trade Hold): same. + - line 3710–3716 (per-bar flat-Hold): same. + +2. `compute_sp18_hold_opportunity_cost` lives in `trade_physics.cuh` (after `compute_lump_sum_opp_cost` at line 590): + - Inputs: `dir_idx`, `tgt[1]` next-bar log-return, `vol_proxy`, `hold_segment_bars`, `isv_pos_cap = isv[483] or fallback 5.0`, `isv_neg_cap = isv[484] or fallback −10.0`, `shaping_scale`. + - Output: 0.0 if `dir_idx != DIR_HOLD`; else `compute_asymmetric_capped_pnl(−2.0 × vol_normalized_log_return, isv_neg_cap, isv_pos_cap)` × shaping_scale. + +**`crates/ml/src/cuda_pipeline/hold_reward_cap_update_kernel.cu`** (NEW) + +Mirrors `reward_cap_update_kernel` shape: producer reads p99 of `|step_ret|` over Long/Short trade closes (per DD3=b), multiplies by safety_factor=1.5, writes POS=that and NEG=`−2 × POS` (DD5=b). Wiener-optimal α via Welford accumulators in `[487..493)`. Pearl-A first-observation bootstrap; sentinel `SENTINEL_HOLD_REWARD_POS_CAP=5.0`. Per-fold cadence (slow EMA, similar to position-side cap producer α=0.01). + +**`crates/ml/src/cuda_pipeline/sp14_isv_slots.rs`** + +D-leg ISV section [483..493): 10 slots per DD6. Sentinels, MIN/MAX bounds, EMA α, safety factor. Compile-time assertion `SP18_SLOT_END <= ISV_TOTAL_DIM`. + +**`crates/ml/src/cuda_pipeline/state_layout.cuh`** + +Mirror constants for SP18 slots (POS/NEG sentinels, MIN/MAX bounds). Bump `ISV_TOTAL_DIM` mirror. + +### Rust changes (D-leg) + +**`crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`** + +- Bump `ISV_TOTAL_DIM` 483 → (B-leg also adds slots; final value pinned in Pre-Phase task PP.2). +- Update `layout_fingerprint_seed` per existing pattern. +- New `launch_hold_reward_cap_update` method. + +**`crates/ml/src/trainers/dqn/state_reset_registry.rs`** + +- Add fold-reset entries for D-leg slots `[483..493)` and B-leg slots (TBD in PP.2). + +**`crates/ml/src/trainers/dqn/trainer/training_loop.rs`** + +- Per-epoch HEALTH_DIAG emit reading `[483..487)` for D-leg. +- DELETE host-side Hold-cost controller block at lines 4171–4188 per DD7(c). +- DELETE `launch_hold_cost_scale_update` call at `training_loop.rs:564` per DD7(c). + +**`crates/ml/src/cuda_pipeline/hold_cost_scale_update_kernel.cu`** — DELETE per DD7(c). + +--- + +## Architecture Changes — B-leg + +### CUDA changes (B-leg) + +No new kernels. **Reuses `compute_expected_q`** (`experience_kernels.cu:4572`) via a new Rust orchestration path that invokes it with target-net params on next-states during experience collection. + +### Rust changes (B-leg) + +**`crates/ml/src/cuda_pipeline/gpu_experience_collector.rs`** + +1. **Move** `let next_states = build_next_states_f32(...)` from line 4182 to *before* line 4137 (the `if td_lambda > 0.01 && max_trace > 1` block). Required so target-net forward pass can consume `next_states`. + +2. **Add** new method `compute_q_next_target_bootstrap(&mut self, next_states: &CudaSlice, target_params_ptr: u64, online_params_ptr: u64, B: usize) -> Result>` that: + - Runs the encoder + dueling forward pass with **online** params on `next_states`, identical to `submit_forward_ops_main`'s Pass 3 (DDQN argmax mirror at `submit_forward_ops_ddqn`). Captures per-branch argmax actions. + - Runs the encoder + dueling forward pass with **target** params on `next_states`, identical to Pass 2 inside `submit_forward_ops_main`. + - Calls `compute_expected_q` on the target-net distributional outputs to produce per-action E[Q] of size `[B, total_actions=13]`. + - Per **B-DD3** per-branch DDQN: gathers `Q_target[a*_online_dir]` for the dir branch and likewise per other 3 branches; sums the 4 selected expected-Q values into a scalar `q_next_per_sample[B]` (matches the SP17 per-branch identifiability claim — each branch's argmax is taken on its own factored Q-table). + - Writes the result into a `CudaSlice` of length `B` (one Q-bootstrap value per sample). + +3. **Replace** line 4143 `let q_next = dtod_clone_f32_native(...)` with `let q_next = self.compute_q_next_target_bootstrap(&next_states, target_params_ptr, online_params_ptr, total)?;`. Drop the `"q_next_bootstrap"` debug name (the buffer is no longer a clone; rename to `"q_next_target"`). + +4. **Distributional flavour of B-DD8**: under C51 the natural "target Q" is per-atom; replacing it with a *scalar* `expected_Q` collapses distributional information at the bootstrap step. Two options: + - **B-DD8.a (scalar)**: `q_next[b] = scalar expected_Q from compute_expected_q on target-net at next_states`. Simplest, matches today's `td_lambda_kernel` signature (which expects `float* q_next`, one per sample). + - **B-DD8.b (distributional)**: extend `td_lambda_kernel` to consume per-atom target probs and emit `R^λ` per-atom, then project at PopArt. Larger surface; defer. + + **Recommendation: B-DD8.a (scalar) for atomic SP18 scope.** The C51 distributional Bellman target at *loss time* is already correct (uses target-net atom positions via `block_bellman_project_f`). The TD(λ) self-bootstrap defect is about credit assignment along the trace, not about the outer atom projection. Scalar `expected_Q` for the trace-blending is sufficient to fix the credit-assignment bug. Distributional λ-blending would be a follow-up if SP18 5-epoch smoke shows residual atom-distribution drift. + +**`crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`** + +- Expose `params_buf_device_ptr()` and `target_params_buf_device_ptr()` to experience collector context (already exposed at lines 8847, 11426). +- Surface a wrapper `compute_expected_q_for_collection(target_pass: bool)` if needed — likely already accessible via existing pointers. + +**`crates/ml/src/trainers/dqn/trainer/training_loop.rs`** + +- B-leg per-epoch HEALTH_DIAG: emit `td_error_magnitude_ema`, `q_next_target_p99`, `q_next_minus_reward_p99` (all per-batch from the new producer kernel — see B-DD12 below). + +### B-leg ISV slots [493..505) + +12 slots (B-DD13 below): + +``` +[493] TD_ERROR_MAG_EMA (HEALTH_DIAG observable: avg(|TD_error|) post-fix) +[494] Q_NEXT_TARGET_P99 (p99 of target-Q bootstrap value) +[495] Q_NEXT_MINUS_REWARD_P99 (sanity: p99 of |q_next_target − rewards| — should be O(1) post-fix) +[496] V_SHARE_TREND_DIAG (slope of V_SHARE EMA across last 4 epochs; B-leg synergy probe) +[497] POPART_RESET_FLAG (1.0 first epoch post-deployment; 0.0 thereafter — gates the per-fold PopArt EMA reset per B-DD11) +[498] TDB_TARGET_MEAN (Welford for q_next_target Wiener-α) +[499] TDB_TARGET_M2 +[500] TDB_DIFF_MEAN +[501] TDB_DIFF_M2 +[502] TDB_PREV_TARGET +[503] TDB_SAMPLE_COUNT +[504] RESERVED (reserve 1 slot for B-leg follow-up if needed) +``` + +**Total combined: 22 slots, [483..505). ISV_TOTAL_DIM bumps 483 → 505.** + +--- + +## DESIGN DECISION 1 — Combined scope (CONFIRMED) + +**Recommendation: D-leg + B-leg combined, atomic in feat/sp18-reward-shape.** Per task input override of v1's "defer B to SP19". + +> _Human reviewer: confirm combined, or override to D-only / B-only._ + +--- + +## DESIGN DECISION 2 — Hold opportunity-cost source (CONFIRMED, recommendation a) + +(a) **`tgt[1]` preprocessed next-bar log-return** — already in scope at line ~2380, vol-normalized via the same `atr_norm` chain the segment_complete path uses. Lookahead-clean. + +> _Reviewer: confirm (a)._ + +--- + +## DESIGN DECISION 3 — Cap producer scope (CONFIRMED, recommendation b) + +(b) **`p99(|step_ret|)` over Long/Short trade closes**. Aligns Hold's reward distribution with policy's *current realized* Long/Short distribution. + +> _Reviewer: confirm (b)._ + +--- + +## DESIGN DECISION 4 — When to apply Hold opportunity-cost (CONFIRMED, recommendation a) + +(a) **Always (every bar where action_dir==Hold)**, both flat and positioned. + +> _Reviewer: confirm (a)._ + +--- + +## DESIGN DECISION 5 — Cap asymmetry inversion direction (CONFIRMED, recommendation b) + +(b) **Mirrored** — same as position side, Hold-side `[+5, −10]`. Says "missing a winner is the bigger pain." Drives the policy *toward* engagement. + +> _Reviewer: confirm (b)._ + +--- + +## DESIGN DECISION 6 — D-leg ISV slot allocation (CONFIRMED, as-specified) + +10 slots [483..493) per the v1 spec. + +> _Reviewer: confirm._ + +--- + +## DESIGN DECISION 7 — Hold cost retention vs. replacement (CONFIRMED, recommendation c) + +(c) **Replace, but keep the observation chain** (slots 380, 381, 382, the `hold_rate_observer_kernel` per-step writes). Delete only the cost-application sites + the scale producer + the host controller. + +> _Reviewer: confirm (c)._ + +--- + +## DESIGN DECISION 8 — Cubin manifest registration timing (CONFIRMED, Phase 3.1) + +Register `hold_reward_cap_update_kernel.cubin` in Phase 3.1 alongside the kernel skeleton (avoids broken-build state in Pre-Phase). + +> _Reviewer: confirm._ + +--- + +## DESIGN DECISION B1 — Target-net infrastructure source (RECOMMENDATION) + +**Use existing `target_params_buf` (Polyak-tracked since SP14 P4.T1B-iv).** No new target-net infra. Verified at `gpu_dqn_trainer.rs:5359` + the target Polyak EMA Pass 2 at loss-time (which proves the buffer is up-to-date relative to online). + +> _Reviewer: confirm._ + +--- + +## DESIGN DECISION B2 — DDQN argmax at bootstrap (RECOMMENDATION) + +`q_next = Q_target(s', argmax_a Q_online(s', a))`. Standard double-DQN pattern. + +> _Reviewer: confirm._ + +--- + +## DESIGN DECISION B3 — Per-branch DDQN argmax (RECOMMENDATION) + +Take 4 separate argmaxes — one each for `dir`, `mag`, `ord`, `urg` factored Q-tables. Matches SP17 per-branch identifiability claim (each branch is its own Q-axis; combining via sum after argmax is the SP17 contract). Sum the 4 selected target-Q values into the scalar `q_next` per sample. + +> _Reviewer: confirm. Alternative: single global argmax on the 13-action joint, which loses the SP17 factored-Q semantics — REJECTED._ + +--- + +## DESIGN DECISION B4 — PER priority recomputation (CONFIRMED) + +PER priorities are **already** recomputed at training time from current TD-error (`config.rs:341`); not stored in buffer or precomputed in experience collector. **No PER migration needed.** The new TD(λ)-corrected `rewards` flowing into the buffer feed naturally into the existing priority recomputation. **B9 oracle test required** to confirm the post-fix priority *distribution* is reasonable. + +> _Reviewer: confirm._ + +--- + +## DESIGN DECISION B5 — Target-net forward pass during experience collection (CONFIRMED) + +Per the audit, `td_lambda_kernel` consumes `q_next` at collection time. To replace with a real `Q(s')`, we MUST run the target-net forward pass at collection time (not loss time). Adds one extra forward pass per experience-collection cycle. + +> _Reviewer: confirm._ + +--- + +## DESIGN DECISION B6 — Accept ~30% collection-time slowdown (RECOMMENDATION) + +Profile in Pre-Phase task PP.5 (NEW) — measure end-to-end experience collection time on H100 before merging Phase 4 to confirm headroom. If >50% slowdown, fall back to **B-DD8.a** scalar-only and skip per-atom target-Q (already the recommendation). If still >50%, halt and convene human review. + +> _Reviewer: confirm._ + +--- + +## DESIGN DECISION B7 — Reuse target-net VSN params (CONFIRMED) + +Existing `target_params_buf` includes VSN params (per SP14 P4.T1B-iv migration). The target-net forward pass on next_states uses the same VSN gates as today's loss-time Pass 2 — same `vsn_gated_next_states_buf`. No new VSN handling. + +> _Reviewer: confirm._ + +--- + +## DESIGN DECISION B8 — Per-atom vs scalar target-Q (RECOMMENDATION a) + +**B-DD8.a (scalar) for atomic SP18.** Distributional flavour deferred to follow-up if 5-epoch smoke shows atom-distribution drift. See B-leg Architecture Changes section above for full reasoning. + +> _Reviewer: confirm a._ + +--- + +## DESIGN DECISION B9 — TD-error magnitude smoke gate (RECOMMENDATION) + +`avg(|TD_error|) ratio post-fix / pre-fix ∈ [0.5, 5.0]` at HD0–HD1 (sanity check that bootstrap change didn't blow up loss magnitude). HD0 baseline measured in Phase 0 diagnostic pre-fix; HD0 measurement post-fix taken in Phase 7 smoke; ratio computed per-fold per-epoch. + +> _Reviewer: confirm._ + +--- + +## DESIGN DECISION B10 — Atomic refactor scope (CONFIRMED) + +Atomic refactor includes: `gpu_experience_collector.rs:4143-4170` (the line 4143 q_next replacement + the target-net forward pass before line 4137) + the new `compute_q_next_target_bootstrap` method + the `next_states` build-site relocation + B-leg ISV slots and HEALTH_DIAG emit. **Audit script in Phase 1 surfaces additional consumers.** PER priority computation and C51 loss kernel migration NOT needed (per B-DD4 and B-leg audit). + +> _Reviewer: confirm._ + +--- + +## DESIGN DECISION B11 — PopArt slot 63 EMA reset at deployment boundary (NEW, RECOMMENDATION) + +Per audit, switching `q_next` source from rewards-distributed to Q-distributed shifts the input distribution to PopArt at the SP18 deployment. PopArt's α-EMA will adapt but may take 1+ epochs. + +**Recommendation**: per-fold reset of PopArt slot 63 EMA at the SP18 deployment boundary, gated by new ISV slot 497 `POPART_RESET_FLAG`. Slot 497 is set to 1.0 by host at first SP18 epoch and 0.0 thereafter (one-shot). On reset, PopArt EMA reinitializes to identity normalization (μ=0, σ=1). Cheap insurance. + +> _Reviewer: confirm, or skip if you'd rather rely on PopArt's natural adaptation._ + +--- + +## DESIGN DECISION B12 — B-leg HEALTH_DIAG observability (NEW, RECOMMENDATION) + +Emit per-epoch: +- `td_error_mag_ema` (slot 493): avg(|TD-error|) computed post-fix +- `q_next_target_p99` (slot 494): bound check on the new bootstrap +- `q_next_minus_reward_p99` (slot 495): sanity — should be O(1), not O(0) (proves the fix is doing something) and not O(100) (proves it isn't blowing up) +- `v_share_trend_diag` (slot 496): slope of V_SHARE across last 4 epochs (should *flatten* or *reverse* under the synergy hypothesis) + +> _Reviewer: confirm._ + +--- + +## DESIGN DECISION B13 — B-leg ISV slot allocation (NEW, RECOMMENDATION) + +12 slots `[493..505)` per the table in "B-leg ISV slots" above. Combined SP18 total: 22 slots `[483..505)`. ISV_TOTAL_DIM bumps 483 → 505. + +> _Reviewer: confirm 12-slot allocation, or trim B-leg diagnostics to live on HEALTH_DIAG-only readout (leave only [493..498) — net 6 B-leg slots, total 16 [483..499))._ + +--- + +## Behavioral Test Predictions + +Tests live at `crates/ml/tests/sp18_hold_reward_oracle_tests.rs` (D-leg) and `crates/ml/tests/sp18_td_lambda_q_next_oracle_tests.rs` (B-leg). Each is a single TDD cycle. + +### D-leg Tests + +#### T1 — CPU oracle for `compute_sp18_hold_opportunity_cost` + +Given: `next_log_return = +0.005`, `vol_proxy = 0.0001`, `pos_cap = 5.0`, `neg_cap = -10.0`, `shaping_scale = 1.0`. Expected: returns negative scalar bounded in `[-10, 0]`. Negative log-return inputs return positive scalar bounded in `[0, +5]`. `dir_idx != DIR_HOLD` returns 0.0 exactly. + +#### T2 — GPU oracle parity with CPU + +Same inputs on `MappedF32Buffer`; bit-identical to ε=1e-6. + +#### T3 — Hold reward distribution on synthetic uniform-edge market + +Synthetic `tgt[1] ~ Uniform[-0.005, +0.005]`. After 1000 Hold picks: `mean(r_hold_opp) ≈ 0`, `var(r_hold_opp) > 0.001`. + +#### T4 — Hold reward distribution on positive-edge synthetic market + +`tgt[1] ~ N(+0.001, 0.002)`. After 1000 Hold picks: `mean(r_hold_opp) < 0`, `mean(r_long_pick - r_hold_opp) > 0` for Long picks on the same bars. + +#### T5 — Cap producer drives slots [483, 484] + +After 100 simulated Long/Short close events with `step_ret ∈ N(0, 0.01)`: `ISV[483] ≈ p99(|step_ret|) × 1.5`, `ISV[484] = -2 × ISV[483]`, both within bounds. + +#### T6 — End-to-end on converged b5gmp checkpoint + +Hold expected-Q variance over actions ≥ 0.5× Long expected-Q variance over actions; no NaN; cap producers populate within 1 launch. + +#### T7 — Scaffold-level: SP13/SP16 deletion compiles + tests pass + +Atomic delete leaves the build green and existing tests pass. + +### B-leg Tests + +#### T8 — CPU oracle for DDQN per-branch argmax bootstrap + +Pure math: given synthetic `Q_online[B, 13]` and `Q_target[B, 13]` (factored 4+3+3+3), compute `q_next_per_sample[b]` as `Σ_branch Q_target[b, argmax_a_in_branch Q_online[b, a]]`. Lock the contract: dir-branch argmax in `[0..4)`, mag in `[4..7)`, ord in `[7..10)`, urg in `[10..13)`. + +#### T9 — GPU oracle parity with CPU (B-leg) + +Same inputs on `MappedF32Buffer`; bit-identical to ε=1e-6. + +#### T10 — TD-error magnitude before vs after fix + +Synthetic 100-step rollout: pre-fix `q_next = rewards`; post-fix `q_next = Q_target(s', argmax_online)`. Verify `avg(|TD_error|) ratio ∈ [0.5, 5.0]` (B-DD9). On a converged checkpoint replay buffer this guards against the bootstrap change blowing up (or zeroing out) the loss. + +#### T11 — Target-net forward consistency + +Run the new `compute_q_next_target_bootstrap` and compare its output to the target-net forward result that already exists at loss time on the same `(next_states, target_params)`. Must be bit-identical (same kernel, same params, same input). Locks "no second source of truth". + +#### T12 — PER priority distribution post-fix + +After 1 training step on a converged-checkpoint replay buffer with the SP18 fix wired: PER priority distribution shape (mean, p50, p99) within ±2× of pre-fix. No degenerate "all priorities equal" or "single priority dominates" signature. + +#### T13 — Graph-capture compatibility + +Capture `exp_fwd_graph` with the new target-net forward pass on `next_states`. Replay 10× and verify per-replay determinism. Locks B-leg Architecture point 1 (move `next_states` build before TD(λ) launch) doesn't break capture. + +--- + +## Smoke Gates (anchored to vtj9r and b5gmp empirical baselines) + +Phase 9 5-epoch L40S smoke. Gates anchored to `train-multi-seed-vtj9r` (terminated 2026-05-09, sha 0c57a5a31) and `train-multi-seed-b5gmp` Fold 1 final (sha 641aa0dfd). + +### D-leg gates + +| Gate | Threshold | Source | +|---|---|---| +| **val Sharpe HD2 (peak)** | **≥ 91.58** | vtj9r HD2 best Sharpe — SP18 must not regress | +| **val Sharpe HD3 (post-peak)** | **≥ 80** | vtj9r HD3 = 78.85; b5gmp Fold-1 = 55.55 (failure) | +| **Hold% HD3** | **≤ 25%** | vtj9r HD3 = 31.6% (climbing); SP18 must keep below | +| **dir_entropy HD3** | **≥ 1.0** | vtj9r HD3 = 0.996 | +| Hold opp-cost fire-rate (HD3) | ≥ 0.80 | structural — D-leg reward must fire | +| `r_hold_opp` distribution width | mean(\|r_hold_opp\|) > 0.5 × mean(\|r_popart\|) | structural — distribution match | +| ISV[483] settled | ∈ [0.01, 5.0] by HD2 | adaptive producer alive by HD2 | +| Q(Hold) − Q(Long) gap | ≤ 0.10 throughout | b5gmp Fold-1 = 0.148 (pre-SP17 failure) | +| no SP13/SP16 dead-code refs | grep returns 0 | DD7(c) atomic delete clean | + +### B-leg gates + +| Gate | Threshold | Source | +|---|---|---| +| **avg(\|TD-error\|) ratio post-fix / pre-fix HD0** | **∈ [0.5, 5.0]** | B-DD9 sanity | +| **v_share trend HD0→HD4** | **monotone decreasing OR flat** | vtj9r showed monotone increasing 0.33→0.83 — SP18 must reverse or flatten | +| q_next_target_p99 (HD2) | ≤ 50.0 | bound check; reward cap is `+5/-10` so target Q ≤ ~50 reasonable | +| q_next_minus_reward_p99 (HD2) | > 0.1 AND < 100.0 | structural — fix is doing something but not exploding | +| PER priority p99 / p50 ratio (HD2) | ≤ 100.0 | distribution sanity | +| **PopArt EMA stable (HD0+1 to HD3)** | **slot 63 not increasing >2× HD0 value** | PopArt should re-converge after B-DD11 reset; large drift = adaptation failure | + +### Combined gates + +| Gate | Threshold | Source | +|---|---|---| +| HEALTH_DIAG never NaN | always | always | +| no train hang | always | always | +| audit script returns clean | grep returns 0 dead refs | DD7(c) + B-leg consumer migration | + +**Phase 10 30-epoch full validation** uses same gates; thresholds tightened on `val Sharpe Fold-1 final ≥ 80`, `Hold% Fold-1 final ≤ 25%`, `dir_entropy Fold-1 final ≥ 0.95`, `Q(Flat)−Q(Long) gap ≤ 0.10`. + +If smoke fails any **bold** gate, halt and convene human review before 30-epoch dispatch. + +--- + +## Hard rules (apply to every task in the plan) + +- **Pre-dispatch consumer audit script** (Phase 1 first task) — surfaces missed consumers BEFORE migration. Promoted to a **pre-commit hook** (Open Q-B confirmed). Generalizable beyond SP18. +- **Outcome-anchored smoke gates** — every threshold cites an empirical baseline. +- **DESIGN DECISION markers** for every genuine choice (DD1–8 + B-DD1–13). +- **TDD** per task — failing test first. +- **No partial refactor** within each leg. +- **ISV-driven adaptive bounds** — pos/neg caps via producer. +- **No CPU compute** — opp-cost producer + B-leg target-net forward + reward computation pure GPU. +- **Mapped-pinned only** for diagnostic readback. +- **Pearl candidate per phase**. + +--- + +## Memory pearl candidates (drafts) + +### Pearl 1 (post-Phase 1 atomic delete): `pearl_replace_reactive_controller_with_structural_signal.md` + +> When a reactive control loop on an emergent symptom (per-bar Hold cost driven by observed Hold-rate overrun) fails to keep up with the policy's per-batch convergence, replace it with a *structural* per-step signal that is immediately present in the reward shape regardless of policy state. SP13/SP16 raised Hold cost reactively after the policy had already over-Held; SP18 D-leg makes Hold's reward distribution structurally comparable to Long/Short by paying the realized log-return as opportunity cost on every Hold bar — distributional-collapse advantage eliminated at source, no controller bandwidth required. + +### Pearl 2 (post-Phase 5 if synergy holds): `pearl_target_net_q_during_experience_collection_for_correct_credit_assignment.md` + +> When TD(λ) / n-step return accumulators run during experience collection, the bootstrap value must be a real `Q(s')` (target-net evaluated on next-states), not the rewards stream. Self-bootstrap (`q_next = rewards`) collapses TD-error to a state-independent reward-tracker which biases the learning signal toward whichever action's reward distribution best matches its long-run accumulation. Symptoms include monotone V-share elevation (V absorbing the action-independent bias because A has no signal to learn) and reward-shape attractors that resist per-step controller fixes. SP18 B-leg fix: run target-net forward on `next_states` during experience collection, gather DDQN per-branch argmax E[Q] from the target-net's distributional output, feed as `q_next` into `td_lambda_kernel`. Adds ~30% to collection time on H100 (acceptable headroom). No replay buffer schema change; no PER priority change (recomputed at training time anyway); no C51 loss kernel change (loss-time target-Q already correct). + +### Pearl 3 (post-Phase 9 if smoke passes): `pearl_distributional_shape_match_for_action_neutrality.md` + +> When one action's Q-distribution is structurally tighter than another's, CQL pessimism + C51 atom occupancy patterns favour the tight branch independent of expected value. To avoid the resulting attractor, design rewards so each action's Q-target distribution has comparable width and tails. SP18 D-leg example: Hold-action Q-targets had width O(0.005) (per-bar penalty band), Long/Short Q-targets had width O(5) (capped trade P&L). Adding bounded opportunity-cost reward to Hold lifted Hold's distribution width to comparable scale. + +### Pearl 4 (post-Phase 9 if combined fix lands): `pearl_complementary_fixes_no_deferrals.md` + +> When two distinct mechanisms drive the same observed pathology, do NOT defer the more-architecturally-honest fix in favour of the simpler one. Fix both in one branch, atomically, so that synergy effects are exposed within the same smoke run. SP18 example: D-leg (reward-shape) and B-leg (TD(λ) bootstrap) were initially scoped as SP18 + SP19. Refactor scopes had empty intersection (D-leg writes reward components, B-leg replaces a `dtod_clone_f32_native` argument); combined ship was safe per `feedback_no_partial_refactor`. Splitting would have required two L40S validation cycles; combined exposed the synergy hypothesis (B-leg fix reverses V-share elevation, which D-leg alone cannot do) in one run. Safe combination criterion: refactor scope intersection MUST be empty AND atomicity rules within each leg are honoured. **DEPRECATES** the v1 SP18 spec's "defer SP19" reasoning. + +### Pearl 5 (post-Phase 9 if smoke fails on D-leg only): `pearl_reward_shape_alone_insufficient_when_bootstrap_biased.md` + +> Reward-shape interventions (asymmetric caps, opportunity costs, etc.) cannot overcome a biased TD(λ) self-bootstrap because the bootstrap biases V-share independent of the reward distribution. SP18 D-leg-only would land Hold's Q-distribution width comparable to Long/Short's, but if v_share continues climbing the reward-shape fix doesn't reach the credit-assignment defect. The combined SP18 D+B is the right scope. + +--- + +## Open questions for human review + +1. **DD1**: Combined SP18 (D+B), or roll back to D-only / B-only? **Recommendation: combined per task input.** +2. **DD2–DD7**: D-leg recommendations confirmed. Reviewer-confirm. +3. **DD8**: Cubin manifest in Phase 3.1. Reviewer-confirm. +4. **B-DD1**: Reuse `target_params_buf`. Reviewer-confirm. +5. **B-DD2**: DDQN argmax. Reviewer-confirm. +6. **B-DD3**: Per-branch DDQN argmax (4 separate argmaxes). Reviewer-confirm. +7. **B-DD4**: PER priority recomputation — no migration. Reviewer-confirm (verified against `config.rs:341`). +8. **B-DD5**: Target-net forward at experience collection. Reviewer-confirm. +9. **B-DD6**: Accept ~30% collection-time slowdown. Profile in PP.5. Reviewer-confirm. +10. **B-DD7**: Reuse target-net VSN params. Reviewer-confirm (verified against SP14 P4.T1B-iv). +11. **B-DD8**: Scalar (a) vs distributional (b) target-Q at bootstrap. **Recommendation: a (scalar).** Reviewer-confirm. +12. **B-DD9**: TD-error ratio gate `[0.5, 5.0]`. Reviewer-confirm. +13. **B-DD10**: Atomic refactor scope (one launch site replacement + new method). Reviewer-confirm. +14. **B-DD11**: PopArt slot 63 EMA reset at SP18 deployment boundary. **Recommendation: yes (cheap insurance).** Reviewer-confirm. +15. **B-DD12**: B-leg HEALTH_DIAG observability. Reviewer-confirm. +16. **B-DD13**: B-leg 12-slot allocation [493..505). Reviewer-confirm or trim to 6 [493..499). +17. **Open Q-A**: `pearl_reward_shape_drives_quality_over_quantity` is already DEPRECATED — confirm Approach D unwinds the recipe (the recipe was per-trade asymmetric cap + min-hold + zero per-bar; D-leg replaces the zero-per-bar leg of that recipe with structural per-bar Hold opp-cost). **Audit confirms: Approach D unwinds the per-bar gap; per-trade cap + min-hold remain operative and unchanged.** +18. **Open Q-B**: Promote consumer-audit script to pre-commit hook. **CONFIRMED — included as Phase 1 sub-task.** +19. **Open Q-C**: SP19 placeholder. **MOOT — absorbed.** + +--- + +## Self-review checklist + +- [x] Problem statement anchored in measured vtj9r data (HD0–HD4 numbers cited; v_share elevation pattern cited). +- [x] D-leg reward audit cites file:line references in current code. +- [x] B-leg TD(λ) audit cites file:line references in current code. +- [x] Four approaches considered (A, B, C, D) with explicit accept/reject reasoning. +- [x] B-leg consumer audit performed: only one production launch site (`gpu_experience_collector.rs:4154`); PER recomputed at training time (no buffer migration); replay buffer doesn't store `q_next`; C51 loss target-Q already correct at loss time. +- [x] Graph-capture interaction documented: target-net forward at collection time is graph-captureable per `pearl_no_host_branches_in_captured_graph`. +- [x] PopArt interaction surfaced (B-DD11): per-fold reset recommended at deployment boundary. +- [x] Consumer audit script is part of Phase 1 (not deferred); promoted to pre-commit hook. +- [x] Every smoke gate cites an empirical baseline (vtj9r HD2/HD3 or b5gmp Fold-1 final or pre/post-fix TD-error ratio). +- [x] DESIGN DECISION markers cover every genuine choice (DD1–8 + B-DD1–13). +- [x] No hardcoded constants for any signal that could derive from data. +- [x] No partial-refactor risk: D-leg deletion + B-leg `q_next` replacement are atomic within their respective legs. +- [x] Pearl candidates flagged for both pass and fail outcomes (D-only fail surfaces Pearl 5; combined pass surfaces Pearl 4). +- [x] SP19 absorbed (no longer flagged as follow-up).