docs(sp22): H6 Phase 3 implementation runbook — 36 tasks across 6 phases

Bite-sized task plan derived from the Phase 3 spec
(docs/plans/2026-05-12-sp22-h6-phase3-alpha-beta.md). Tasks grouped
into six phases; all EDIT tasks land in a single atomic commit at
Phase F per `feedback_no_partial_refactor`. There are NO intermediate
commits across Phase A–E; each task's cargo-check step is the
in-flight gate.

Phase A (5 tasks): foundation — new ISV slots
(REWARD_AUX_ALIGN_EMA_INDEX + SP22_AUX_ALIGN_SCALE_INDEX),
state_reset_registry entries, new alpha forward + backward kernels,
build.rs registration.

Phase B (11 tasks): 7-component contract migration —
experience_kernels.cu (producer + beta), backtest_env_kernel.cu
(eval producer + beta), reward_component_ema_kernel.cu (0..7
iteration), reward_decomp_diag_kernel.cu (column 6),
reward_component_mag_ratio_compute_kernel.cu (axis 6),
reward_subsystem_controller_kernel.cu (7-weight emit + w_aux_align
anchor), health_diag_kernel.cu (snap slot),
gpu_experience_collector.rs (7-slot buffer + HEALTH_DIAG print),
gpu_dqn_trainer.rs (7-slot buffer + HEALTH_DIAG + alpha param +
Adam + captured graph wiring), reward_component_monitor.rs (reader),
training_loop.rs (sp11_reward + reward_split print lines).

Phase C (1 task): alpha plumbing collector-side / rollout-time
captured graph (Phase C is otherwise folded into Phase B Task B9).

Phase D (7 tasks): A2 eval-side — aux trunk weight loading +
per-window prev_aux_dir_prob buffer + per-step aux trunk forward +
softmax-to-per-env launch + alpha forward + non-NULL state-gather
wiring across 3 launchers + W ptr sharing trainer ↔ evaluator.

Phase E (7 tasks): verification gates — cargo check,
gpu_backtest_validation, compute-sanitizer, captured-graph integrity
(deferred to smoke), alpha weight movement (deferred to smoke),
7-component contract sanity (deferred to smoke), A2 eval-side aux
activity (deferred to validation eval).

Phase F (5 tasks): audit doc append (Invariant 7), single atomic
commit covering all changes, push, smoke dispatch, verdict
monitoring + audit doc verdict.

Total: 36 tasks. Each task has exact file paths, exact commands with
expected output, and code-block-level specificity for kernel
signatures, struct field additions, and HEALTH_DIAG format-string
extensions. Self-review section maps every spec requirement to a
specific task.

Refs
────
- docs/plans/2026-05-12-sp22-h6-phase3-alpha-beta.md (spec)
- pearl_no_partial_refactor (single atomic commit Phase F)
- feedback_push_before_deploy (Phase F3)
- feedback_no_redundant_monitor (Phase F5 single-channel waiter)
- feedback_kill_runs_on_anomaly_quickly (verdict criteria)

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

docs/plans/2026-05-13-sp22-h6-phase3-alpha-beta-runbook.md

@@ -0,0 +1,505 @@
+# SP22 H6 Phase 3 — alpha + beta + 7-component + SP11 + A2 eval-side Implementation Runbook
+
+> **For agentic workers:** REQUIRED SUB-SKILL: Use superpowers:executing-plans (or superpowers:subagent-driven-development) to implement this plan task-by-task. Steps use checkbox (`- [ ]`) syntax for tracking.
+
+**Goal:** Wire the aux to-policy bypass routing end-to-end — alpha (post-encoder Q_dir bias) + beta (event-driven aux-aligned trade-close bonus, 7th reward component) + SP11 controller for scale_beta + A2 eval-side parity — so the policy can actually use the aux signal it has been ignoring across Phase 1 / Phase 2.
+
+**Architecture:** Two parallel routing paths bypass the cold-encoder bottleneck (state[121] encoder weight stuck near zero). **alpha** adds a 4-weight learnable bias to Q_dir post-encoder, fed by `state_121 in [-1, +1]` and trained via standard Adam. **beta** fires at trade-close as a 7th reward component (`r_aux_align = scale_beta * alignment * profit_pos`) with adaptive scale_beta emitted by the existing SP11 controller. The 7-component contract migration cascades through every reward-decomposition consumer atomically per `feedback_no_partial_refactor`. A2 brings the eval pipeline to production parity: aux trunk forward at eval, per-window p_up buffer, alpha launch at eval, beta producer in `backtest_env_kernel`.
+
+**Tech Stack:** Rust 1.85+ Edition 2021; CUDA 12.4 sm_89 (L40S target); cudarc 0.19; SQLX_OFFLINE for cargo check; argo-workflows for smoke dispatch.
+
+**Working tree:** `/home/jgrusewski/Work/foxhunt/.worktrees/sp19-20-wr-first/` on branch `sp20-aux-h-fixed`. All commands assume cwd is this worktree.
+
+**Spec:** `docs/plans/2026-05-12-sp22-h6-phase3-alpha-beta.md` (canonical design + decision rationale + verdict criteria).
+
+---
+
+## Phase-level structure
+
+Tasks are grouped into six phases. **All EDIT tasks land in a single atomic commit at Phase F** per `feedback_no_partial_refactor` — the 7-component contract spans many files and partial migration produces inconsistent reward stride across consumers. There are NO intermediate commits across Phase A–E; each task's `cargo check` step is the in-flight gate to ensure no consumer is left behind.
+
+| Phase | Purpose | # tasks |
+|---|---|---|
+| A | Foundation: new ISV slots + state_reset_registry + new alpha kernels + build.rs registration | 5 |
+| B | 7-component contract migration (cascade across all reward-decomp consumers) | 11 |
+| C | alpha plumbing (trainer-side param, Adam, captured graph wiring) | 1 |
+| D | A2 eval-side aux trunk + alpha + beta + state-gather wiring | 7 |
+| E | Verification gates (cargo check, gpu_backtest_validation, compute-sanitizer, capture-pulse, post-smoke checks) | 7 |
+| F | Audit doc + atomic commit + push + smoke dispatch + monitor verdict | 5 |
+
+Total: **36 tasks**. (Phase C is folded inside Phase B Task B9 — alpha plumbing in `gpu_dqn_trainer.rs` lands in the same edit as the 7-component buffer alloc; Task C1 covers only the collector-side rollout-time wiring.)
+
+---
+
+## File map (recap from spec, with task ownership)
+
+| File | First-touched in task | Owner-phase |
+|---|---|---|
+| `crates/ml/src/cuda_pipeline/aux_to_q_dir_bias_kernel.cu` (NEW) | A3 | alpha forward kernel |
+| `crates/ml/src/cuda_pipeline/aux_to_q_dir_bias_backward_kernel.cu` (NEW) | A4 | alpha backward kernel |
+| `crates/ml/build.rs` | A5 | Cubin registration |
+| `crates/ml/src/cuda_pipeline/sp22_isv_slots.rs` (NEW or EXTEND) | A1 | New ISV slot constants |
+| `crates/ml/src/trainers/dqn/state_reset_registry.rs` | A2 | Slot registry entries |
+| `crates/ml/src/cuda_pipeline/experience_kernels.cu` | B1 | 7-component producer + beta |
+| `crates/ml/src/cuda_pipeline/backtest_env_kernel.cu` | B2 | 7-component producer + beta (eval) |
+| `crates/ml/src/cuda_pipeline/reward_component_ema_kernel.cu` | B3 | EMA iteration 0..7 |
+| `crates/ml/src/cuda_pipeline/reward_decomp_diag_kernel.cu` | B4 | Diag emit column 6 |
+| `crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu` | B5 | Axis 6 |
+| `crates/ml/src/cuda_pipeline/reward_subsystem_controller_kernel.cu` | B6 | 7-weight output + w_aux_align anchor |
+| `crates/ml/src/cuda_pipeline/health_diag_kernel.cu` | B7 | Snap layout |
+| `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` | B8 / C1 | 7-slot alloc + HEALTH_DIAG + alpha rollout-side fwd |
+| `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` | B9 | 7-slot alloc + HEALTH_DIAG + alpha param + Adam + graphs |
+| `crates/ml/src/trainers/dqn/monitors/reward_component_monitor.rs` | B10 | 7-component EMA reader |
+| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | B11 | sp11_reward + reward_split print lines |
+| `crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs` | D1 | A2 eval-side aux trunk + alpha + state-gather |
+| `docs/dqn-wire-up-audit.md` | F1 | Audit doc append |
+
+---
+
+## Phase A — Foundation
+
+### Task A1: Define new ISV slot constants
+
+**Files:**
+- Modify (or Create): `crates/ml/src/cuda_pipeline/sp22_isv_slots.rs`
+
+- [ ] **Step 1: Check whether `sp22_isv_slots.rs` exists.** Run `ls crates/ml/src/cuda_pipeline/sp22_isv_slots.rs`. If exists, grep current `SP22_*_INDEX` constants to find the next available slot number. If does not exist, create as a new file.
+
+- [ ] **Step 2: Add the slot constants.** Both `REWARD_AUX_ALIGN_EMA_INDEX` (EMA producer output, consumed as SP11 controller anchor + HEALTH_DIAG) and `SP22_AUX_ALIGN_SCALE_INDEX` (controller-emitted adaptive scale_beta, consumed by beta producers at trade-close). Pick the lowest two contiguous unused indices. Document the contract in the file's module-level doc-comment.
+
+- [ ] **Step 3: Verify compile via `SQLX_OFFLINE=true cargo check -p ml --features cuda`.** Expected: 0 errors; unused-import warnings acceptable since consumers wire up in subsequent tasks.
+
+### Task A2: Register the new ISV slots in state_reset_registry
+
+**Files:**
+- Modify: `crates/ml/src/trainers/dqn/state_reset_registry.rs`
+
+- [ ] **Step 1: Find an existing FoldReset entry** to mirror via `grep -nB1 -A4 "category: ResetCategory::FoldReset" crates/ml/src/trainers/dqn/state_reset_registry.rs | head -30`. Note the entry format (name, index, category, description).
+
+- [ ] **Step 2: Add two FoldReset entries to the registry slice** — one for `REWARD_AUX_ALIGN_EMA_INDEX` ("sp22_reward_aux_align_ema") with description covering producer (`reward_component_ema_kernel`), consumer (SP11 controller anchor + HEALTH_DIAG), Pearl A bootstrap from sentinel 0. Mirror for `SP22_AUX_ALIGN_SCALE_INDEX` ("sp22_aux_align_scale") with description covering producer (extended `reward_subsystem_controller_kernel`), consumers (training-side beta + eval-side beta), cold-start sentinel 0 → first emit replaces. Add matching `use crate::cuda_pipeline::sp22_isv_slots::*` imports.
+
+- [ ] **Step 3: Verify compile.** If `training_loop.rs:9008` (the name-string match-expression on registry entries) needs handling for the new names, surface and fix.
+
+### Task A3: Create alpha forward kernel
+
+**Files:**
+- Create: `crates/ml/src/cuda_pipeline/aux_to_q_dir_bias_kernel.cu`
+
+- [ ] **Step 1: Write the kernel source.** Signature: `aux_to_q_dir_bias(float* q_dir [INOUT, B*b0_size], const float* w_aux [b0_size], const float* batch_states [B*state_dim], int aux_dir_prob_index, int state_dim, int batch_size, int b0_size)`. Kernel body: `int tid = blockIdx.x * blockDim.x + threadIdx.x; if (tid >= B*b0_size) return; int b = tid / b0_size; int a = tid - b * b0_size; float state_121 = batch_states[(long long)b * state_dim + aux_dir_prob_index]; q_dir[(long long)b * b0_size + a] += w_aux[a] * state_121;`. Document the architectural role (reads encoder INPUT, bypassing cold encoder), discipline (no atomicAdd, no host branches, capture-safe), and launch shape (grid = ceil(B*b0_size / 256), block = 256). Header preamble cites `feedback_no_atomicadd`, `pearl_no_host_branches_in_captured_graph`, `feedback_no_cpu_compute_strict`.
+
+- [ ] **Step 2: Verify file created** via `ls -la crates/ml/src/cuda_pipeline/aux_to_q_dir_bias_kernel.cu`.
+
+### Task A4: Create alpha backward kernel
+
+**Files:**
+- Create: `crates/ml/src/cuda_pipeline/aux_to_q_dir_bias_backward_kernel.cu`
+
+- [ ] **Step 1: Write the kernel source with TWO kernels in one .cu file.** Kernel `aux_to_q_dir_bias_backward_dw`: per-action block tree-reduce over batch; signature `(const float* dq_dir [B*b0_size], const float* batch_states [B*state_dim], float* dw_aux [b0_size, OUT], int aux_dir_prob_index, int state_dim, int batch_size, int b0_size)`. Launch one block per action, 256 threads, standard log2 tree reduce in shared mem; thread 0 writes `dw_aux[a]`. Kernel `aux_to_q_dir_bias_backward_dstate`: per-sample sum; signature `(const float* dq_dir, const float* w_aux, float* dbatch_states [B*state_dim, INOUT], int aux_dir_prob_index, int state_dim, int batch_size, int b0_size)`. Launch ceil(B/256), 256 threads; each thread computes `sum = Sum_a w_aux[a] * dq_dir[b*b0_size+a]` and ACCUMULATES (`+=`) into `dbatch_states[b*state_dim + aux_dir_prob_index]`. Both kernels: no atomicAdd, no host branches, capture-safe. Document option-(i) gradient routing (both alpha and encoder paths get gradient).
+
+- [ ] **Step 2: Verify file created.**
+
+### Task A5: Register both alpha kernels in build.rs
+
+**Files:**
+- Modify: `crates/ml/build.rs`
+
+- [ ] **Step 1: Find the existing `aux_softmax_to_per_env_kernel.cu` entry** via `grep -nE "aux_softmax_to_per_env_kernel.cu" crates/ml/build.rs`.
+
+- [ ] **Step 2: Add the two new kernel filenames** right after, in the `kernels_with_common` array: `"aux_to_q_dir_bias_kernel.cu",` and `"aux_to_q_dir_bias_backward_kernel.cu",`. Each gets a brief comment citing SP22 H6 Phase 3 alpha and the kernel role.
+
+- [ ] **Step 3: Verify the build picks up the new cubins** via `SQLX_OFFLINE=true cargo check -p ml --features cuda 2>&1 | grep -iE "aux_to_q_dir_bias|error" | head -10`. Expected: nvcc compiles both; 0 errors. If nvcc reports kernel-source errors, fix the kernel.
+
+---
+
+## Phase B — 7-component contract migration (atomic across consumers)
+
+### Task B1: experience_kernels.cu — 7-component producer + beta at training-side trade close
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/experience_kernels.cu`
+
+- [ ] **Step 1: Locate the current 6-component stride references** via `grep -nE "reward_components_per_sample\\[.*\\* 6 \\+|rc\\[5\\]|reward_components.*\\* 6" crates/ml/src/cuda_pipeline/experience_kernels.cu`. Note every `out_off * 6 + N` and `rc[N]` reference.
+
+- [ ] **Step 2: Update the kernel preamble doc** describing the reward-component layout: replace the 6-component documentation with 7-component (existing 0–5 unchanged; new `[6] aux_align — SP22 H6 Phase 3 beta: event-driven aux-aligned trade-close bonus, scale_beta read from ISV[SP22_AUX_ALIGN_SCALE_INDEX]`).
+
+- [ ] **Step 3: Migrate every `out_off * 6 +` to `out_off * 7 +`.** Use Edit with `replace_all` on the exact string `out_off * 6 +` only if grep confirms no false matches in non-reward contexts (which it should — that stride is specific to reward components). Otherwise edit one match at a time.
+
+- [ ] **Step 4: Add beta producer at the segment_complete branch.** Locate by grepping for the existing `r_opp_cost = compute_lump_sum_opp_cost` site (around line 3500 area, as discovered during spec investigation). Add code that computes `position_sign`, `aux_at_close` (NULL-safe from `aux_dir_prob_per_env`), `alignment = fmaxf(0.0f, aux_at_close * position_sign)`, `profit_pos = fmaxf(0.0f, step_ret_core)`, `scale_beta_isv` (NULL-safe from `isv_signals_ptr[SP22_AUX_ALIGN_SCALE_INDEX]`), `r_aux_align = scale_beta_isv * alignment * profit_pos`. Write to `reward_components_per_sample[out_off * 7 + 6] = r_aux_align`. Include the necessary header/define for `SP22_AUX_ALIGN_SCALE_INDEX`.
+
+- [ ] **Step 5: Verify compile.** Expected: 0 errors OR only errors that are propagating stride mismatches to other consumers (which Tasks B2–B11 fix). Proceed only if errors are exclusively the expected cascade.
+
+### Task B2: backtest_env_kernel.cu — 7-component producer + beta at eval-side trade close
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/backtest_env_kernel.cu`
+
+- [ ] **Step 1: Find analogous reward-component writes** via `grep -nE "reward_components_per_sample\\[.*\\* 6 \\+|rc\\[" crates/ml/src/cuda_pipeline/backtest_env_kernel.cu`.
+
+- [ ] **Step 2: Update preamble doc** to 7-component layout, same content as B1 Step 2.
+
+- [ ] **Step 3: Migrate `* 6 +` → `* 7 +`** in reward-component writes. Same pattern as B1 Step 3.
+
+- [ ] **Step 4: Add beta producer at segment_complete** mirroring B1 Step 4 with the eval-kernel's local variable names (`w` instead of `i` for window index; whatever local realized-PnL variable the kernel uses). NULL-safe: both `aux_dir_prob_per_env` and `isv_signals` NULL paths give 0.
+
+- [ ] **Step 5: Verify compile.** Same cascade-tolerance as B1 Step 5.
+
+### Task B3: reward_component_ema_kernel.cu — iterate 0..7
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/reward_component_ema_kernel.cu`
+
+- [ ] **Step 1: Locate the component iteration** via `grep -nE "for.*c <|c < 6|N_COMPONENTS|num_components" crates/ml/src/cuda_pipeline/reward_component_ema_kernel.cu`.
+
+- [ ] **Step 2: Update the bound from 6 to 7.** Whether the kernel uses a constant macro (`#define N_COMPONENTS 6`) or a literal (`c < 6`), change to 7. Also extend `var_scratch_first_index` slot count by 1 if applicable (per kernel preamble describing the var-scratch layout).
+
+- [ ] **Step 3: Verify compile.**
+
+### Task B4: reward_decomp_diag_kernel.cu — emit column 6
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/reward_decomp_diag_kernel.cu`
+
+- [ ] **Step 1: Locate column-emit loop** via `grep -nE "col [0-9]|component idx|c <|< 6" crates/ml/src/cuda_pipeline/reward_decomp_diag_kernel.cu`.
+
+- [ ] **Step 2: Update iteration / column count to 7** and add a `col 6: mean(r_aux_align) per action bin (component idx 6)` line to the kernel preamble doc.
+
+- [ ] **Step 3: Verify compile.**
+
+### Task B5: reward_component_mag_ratio_compute_kernel.cu — axis 6
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu`
+
+- [ ] **Step 1: Locate the cf_others axis map** via `grep -nE "axis [0-9]|cf_others_base_slot|isv\\[.*\\+ [0-9]+\\]" crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu`.
+
+- [ ] **Step 2: Add axis 5 (aux_align) to the map.** The cf_others axes are 0..4 covering cf, trail, micro, opp_cost, bonus; add axis 5 reading the ISV slot for aux_align EMA. If the slot layout for cf_others is contiguous, allocate ONE more slot at the end and update the base + count constants; otherwise allocate aux_align at a separate slot and document non-contiguity.
+
+- [ ] **Step 3: Verify compile.**
+
+### Task B6: reward_subsystem_controller_kernel.cu — 7-weight output + w_aux_align anchor
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/reward_subsystem_controller_kernel.cu`
+
+- [ ] **Step 1: Locate the existing 6-weight emit** via `grep -nE "w_pop|w_cf|w_tr|w_mi|w_oc|w_bn|6 components|0\\.\\.6" crates/ml/src/cuda_pipeline/reward_subsystem_controller_kernel.cu`.
+
+- [ ] **Step 2: Extend to 7-weight output.** Add `w_aux_align` to the emit list following the same per-weight controller pattern as the other 6 weights. Anchor signal: `isv[REWARD_AUX_ALIGN_EMA_INDEX]`. Target: `0.10 × Sum_c |reward_components_ema[c]|` so beta contributes ~10% of total reward signal at steady state. Bounds: floor 0.05, ceiling 2.0 (enforced inside the kernel). Cold-start handling per `pearl_first_observation_bootstrap`: if `isv[SP22_AUX_ALIGN_SCALE_INDEX] == 0` (sentinel) AND anchor is non-zero, write the new value directly. Write the emitted weight to `isv[SP22_AUX_ALIGN_SCALE_INDEX]`.
+
+- [ ] **Step 3: Verify compile.**
+
+### Task B7: health_diag_kernel.cu — snap layout slot for r_aux_align
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/health_diag_kernel.cu`
+
+- [ ] **Step 1: Locate snap layout doc** via `grep -nE "REWARD_.*_EMA_INDEX|reward_split|snap.reward" crates/ml/src/cuda_pipeline/health_diag_kernel.cu`.
+
+- [ ] **Step 2: Add `ISV[REWARD_AUX_ALIGN_EMA_INDEX] -> snap.reward_split_aux_align` mapping** in both the kernel preamble doc and the kernel body. If the snap struct is defined in a header, extend that struct by one field. Bump the snap word-count by 1 if fixed-size.
+
+- [ ] **Step 3: Verify compile.**
+
+### Task B8: gpu_experience_collector.rs — 7-slot buffer alloc + HEALTH_DIAG print update
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs`
+
+- [ ] **Step 1: Locate the reward_components_per_sample allocation** via `grep -nE "reward_components_per_sample|N_REWARD_COMPONENTS|6 components" crates/ml/src/cuda_pipeline/gpu_experience_collector.rs`.
+
+- [ ] **Step 2: Migrate allocation from 6-slot → 7-slot stride per sample.** Find any `alloc_episodes * 6` or `N_REWARD_COMPONENTS = 6` constants and change to 7.
+
+- [ ] **Step 3: Update HEALTH_DIAG print line for `reward_split`** — extend the format string and arg list to include `aux_align`. Read the 7th-component EMA from `trainer.read_isv_signal_at(REWARD_AUX_ALIGN_EMA_INDEX)` or the monitor's reader (updated in Task B10).
+
+- [ ] **Step 4: Verify compile.**
+
+### Task B9: gpu_dqn_trainer.rs — 7-slot buffer alloc + HEALTH_DIAG + alpha param + Adam + graphs
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`
+
+This is the largest single change. It covers both the 7-component contract migration AND the alpha plumbing. Lands in the same atomic Phase F commit.
+
+- [ ] **Step 1: Migrate reward_components_per_sample buffer alloc to 7 slots** (same pattern as B8).
+
+- [ ] **Step 2: Update HEALTH_DIAG print for `sp11_reward`** — extend the 6-weight format to 7 weights, adding `w_aux`. Read it from `read_isv_signal_at(SP22_AUX_ALIGN_SCALE_INDEX)`.
+
+- [ ] **Step 3: Update HEALTH_DIAG print for `reward_split`** to include `aux_align` (same as B8 Step 3 at the trainer's own print site).
+
+- [ ] **Step 4: Add `pub(crate) static SP22_AUX_TO_Q_DIR_BIAS_CUBIN` and `SP22_AUX_TO_Q_DIR_BIAS_BWD_CUBIN` statics** at the top of the file alongside the other SP22_* CUBIN statics, using `include_bytes!(concat!(env!("OUT_DIR"), "/aux_to_q_dir_bias_kernel.cubin"))` pattern (and similar for the backward cubin).
+
+- [ ] **Step 5: Add alpha-related struct fields**: `w_aux_to_q_dir: CudaSlice<f32>` [b0_size=4], `adam_m_w_aux: CudaSlice<f32>` [4], `adam_v_w_aux: CudaSlice<f32>` [4], `dw_aux_buf: CudaSlice<f32>` [4], `aux_to_q_dir_bias_kernel: CudaFunction`, `aux_to_q_dir_bias_backward_dw_kernel: CudaFunction`, `aux_to_q_dir_bias_backward_dstate_kernel: CudaFunction`. Document each with SP22 H6 Phase 3 alpha pointer.
+
+- [ ] **Step 6: Allocate buffers + load kernels in `new()`.** Allocate the 4 alpha-related `CudaSlice<f32>` buffers with `alloc_zeros::<f32>(b0_size)`. Load the alpha forward kernel from `SP22_AUX_TO_Q_DIR_BIAS_CUBIN`; load both backward kernels (`aux_to_q_dir_bias_backward_dw` + `aux_to_q_dir_bias_backward_dstate`) from `SP22_AUX_TO_Q_DIR_BIAS_BWD_CUBIN`. Add all 7 new fields to the struct construction at the end of `new()`.
+
+- [ ] **Step 7: Wire alpha forward into the training-time forward captured graph.** Find the Q-head forward site (grep for `q_dir`). After Q_dir is computed but before Q_dir is consumed by downstream branches (mag concat / action selection), add a launch of `aux_to_q_dir_bias_kernel` with args (q_dir_buf, w_aux_to_q_dir, batch_states, AUX_DIR_PROB_INDEX, STATE_DIM, batch_size, b0_size). Launch shape: grid = ceil(B*b0_size / 256), block = 256.
+
+- [ ] **Step 8: Wire alpha backward into the training-time backward captured graph.** Find the site where `dq_dir` is computed (after C51 / IQN heads emit gradients). Launch `aux_to_q_dir_bias_backward_dw_kernel` (grid = b0_size blocks, 256 threads each) and `aux_to_q_dir_bias_backward_dstate_kernel` (grid = ceil(B/256), 256 threads). Both consume `dq_dir`; the first writes `dw_aux_buf`, the second accumulates into `dbatch_states_buf[, AUX_DIR_PROB_INDEX]`.
+
+- [ ] **Step 9: Wire alpha Adam-step update.** Find the existing Adam-step site (where other parameter buffers get their per-step updates). Add an Adam call for `w_aux_to_q_dir` using the existing generic per-param Adam kernel. Launch shape: grid = (1, 1, 1), block = (4, 1, 1) since only 4 parameters.
+
+- [ ] **Step 10: Verify compile.** Expected: 0 errors. If `q_dir_buf` / `dq_dir_buf` / `dbatch_states_buf` / `adam_step_kernel` / `lr_current` symbol names differ in the trainer source, grep for the actual names and adapt.
+
+### Task B10: reward_component_monitor.rs — 7-component EMA reader
+
+**Files:**
+- Modify: `crates/ml/src/trainers/dqn/monitors/reward_component_monitor.rs`
+
+- [ ] **Step 1: Locate the existing 6-component reader** via `grep -nE "REWARD_.*_EMA_INDEX|REWARD_BONUS_EMA|isv.read" crates/ml/src/trainers/dqn/monitors/reward_component_monitor.rs`.
+
+- [ ] **Step 2: Add `aux_align` field + `isv.read(REWARD_AUX_ALIGN_EMA_INDEX)` reader.** Mirror the existing component-reader fields. Add the `use crate::cuda_pipeline::sp22_isv_slots::REWARD_AUX_ALIGN_EMA_INDEX` import.
+
+- [ ] **Step 3: Verify compile.**
+
+### Task B11: training_loop.rs — sp11_reward + reward_split print lines
+
+**Files:**
+- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs`
+
+- [ ] **Step 1: Locate HEALTH_DIAG print sites** via `grep -nE "sp11_reward \\[w_pop|reward_split \\[popart" crates/ml/src/trainers/dqn/trainer/training_loop.rs`.
+
+- [ ] **Step 2: Extend both format strings** — `sp11_reward` gains `w_aux={:.3}` after `w_bn={:.3}`; `reward_split` gains `aux_align={:.4}` after `bonus={:.4}`. New format args: `trainer.read_isv_signal_at(SP22_AUX_ALIGN_SCALE_INDEX)` and `trainer.read_isv_signal_at(REWARD_AUX_ALIGN_EMA_INDEX)`. Add necessary `use` imports.
+
+- [ ] **Step 3: Verify compile.** Expected: 0 errors. With Tasks B1–B11 complete, the 7-component contract is fully migrated across training-side consumers.
+
+---
+
+## Phase C — alpha plumbing (collector / rollout-side)
+
+### Task C1: gpu_experience_collector.rs — wire alpha forward into rollout-time captured graph
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs`
+
+- [ ] **Step 1: Locate the rollout-time Q-head forward site** via `grep -nE "q_dir|on_dir_logits|forward_online" crates/ml/src/cuda_pipeline/gpu_experience_collector.rs`.
+
+- [ ] **Step 2: Load the alpha forward kernel handle in `new()`** via `super::gpu_dqn_trainer::SP22_AUX_TO_Q_DIR_BIAS_CUBIN.to_vec()` → `load_function("aux_to_q_dir_bias")`. Store in `exp_aux_to_q_dir_bias_kernel: CudaFunction`. Add `w_aux_to_q_dir_dev_ptr: u64` field initialized to 0 (set via setter post-construction by the trainer).
+
+- [ ] **Step 3: Add `set_w_aux_ptr(&mut self, ptr: u64)` setter method.**
+
+- [ ] **Step 4: Wire alpha forward into the rollout-time forward graph.** Inside the rollout-step forward loop after Q_dir is computed and before action selection, add the same launch as B9 Step 7 with collector's local variables (n_episodes instead of B; w_aux_ptr from the setter).
+
+- [ ] **Step 5: Trigger the setter from the trainer**: in `gpu_dqn_trainer.rs`, after constructing the collector, call `collector.set_w_aux_ptr(self.w_aux_to_q_dir.raw_ptr())`.
+
+- [ ] **Step 6: Verify compile.**
+
+---
+
+## Phase D — A2 eval-side aux trunk + alpha + beta + state-gather wiring
+
+### Task D1: gpu_backtest_evaluator.rs — discovery of aux trunk weight references
+
+**Files:**
+- Modify (deferred to D2+): `crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs`
+- Inspect (read-only): `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`
+
+- [ ] **Step 1: Map the aux trunk weight buffers in the trainer** via `grep -nE "aux_trunk|aux_heads|aux_W|aux_in_proj|aux_out_proj|exp_h_s2_aux|aux_logits" crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`. Identify the input projection, hidden, and output projection buffers.
+
+- [ ] **Step 2: Determine if a reusable aux forward orchestrator exists** via `grep -nE "fn .*aux.*forward|AuxHeadsForward|aux_next_bar_forward|aux_trunk_forward" crates/ml/src/cuda_pipeline/`. If yes (e.g., `AuxHeadsForward::forward`), prefer reuse at eval-time. If only training-coupled code exists, plan a thin eval-only variant calling `cublasSgemm` directly.
+
+- [ ] **Step 3: Record findings as a code comment** in `gpu_backtest_evaluator.rs` module header documenting the chosen approach (reuse vs thin variant) and the specific weight-buffer field names. No verification step; D2+ does the actual wiring.
+
+### Task D2: gpu_backtest_evaluator.rs — per-window prev_aux_dir_prob buffer + reset
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs`
+
+- [ ] **Step 1: Add `prev_aux_dir_prob: CudaSlice<f32>` and `eval_fill_f32_kernel: CudaFunction` fields** to the evaluator struct. Document the Phase 2 sentinel = 0 semantics.
+
+- [ ] **Step 2: Allocate in `new()`** — `stream.alloc_zeros::<f32>(n_windows)` (already 0.0 sentinel; no fill_f32 needed at init). Load `fill_f32` from `super::gpu_action_selector::EPSILON_GREEDY_CUBIN`.
+
+- [ ] **Step 3: Add per-evaluator-call reset.** Find the entry point of `pub fn evaluate(...)` (the eval entry). At its start, launch `fill_f32` with value 0.0 over `n_windows`. Document cold-start semantics matching Phase 2.
+
+- [ ] **Step 4: Append both new fields to the struct construction at end of `new()`.**
+
+- [ ] **Step 5: Verify compile.**
+
+### Task D3: gpu_backtest_evaluator.rs — per-step aux trunk forward + softmax-to-per-env launch
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs`
+
+- [ ] **Step 1: Allocate per-window aux activation buffers**: `aux_h_s2_buf` [n_windows × shared_h2], `aux_logits_buf` [n_windows × K=2], `aux_softmax_buf` [n_windows × K=2]. All `alloc_zeros::<f32>` in `new()`.
+
+- [ ] **Step 2: Load `aux_softmax_to_per_env_kernel`** (existing Phase 1 cubin `SP22_AUX_SOFTMAX_TO_PER_ENV_CUBIN` from `super::gpu_dqn_trainer`). Store in `eval_aux_softmax_to_per_env_kernel: CudaFunction`.
+
+- [ ] **Step 3: Per-step aux trunk forward in the per-step loop of the evaluator.** Between encoder forward and action selection, invoke either the reused `AuxHeadsForward` orchestrator (option a) or a thin eval-only forward (option b) per D1's findings. The forward writes `aux_softmax_buf [n_windows × K=2]`.
+
+- [ ] **Step 4: Launch `aux_softmax_to_per_env_kernel`** immediately after the aux forward: args (aux_softmax_buf.raw_ptr, prev_aux_dir_prob.raw_ptr, n_windows_i32, aux_k_i32=2). Launch shape: grid = ceil(n_windows / 256), block = 256.
+
+- [ ] **Step 5: Verify compile.**
+
+### Task D4: gpu_backtest_evaluator.rs — alpha post-Q_dir bias launch at eval
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs`
+
+- [ ] **Step 1: Load alpha forward kernel handle** from `super::gpu_dqn_trainer::SP22_AUX_TO_Q_DIR_BIAS_CUBIN` in `new()`. Store in `eval_aux_to_q_dir_bias_kernel: CudaFunction`. Add `w_aux_to_q_dir_dev_ptr: u64` field initialized to 0 (set via setter).
+
+- [ ] **Step 2: Add `set_w_aux_ptr(&mut self, ptr: u64)` setter method.**
+
+- [ ] **Step 3: Wire alpha forward into eval per-step**: after Q_dir is computed by the eval forward, before action selection / model closure invocation, launch the alpha kernel mirroring B9 Step 7 with eval-side local variables (n_windows instead of B; w_aux_ptr from setter; states_buf as encoder input).
+
+- [ ] **Step 4: Verify compile.**
+
+### Task D5: gpu_backtest_evaluator.rs — wire non-NULL prev_aux_dir_prob to 3 state-gather launchers
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs`
+
+- [ ] **Step 1: Locate all 3 state-gather launchers** via `grep -nE "backtest_state_gather launch|launch_builder\\(&self.gather_kernel\\)|launch_builder\\(&self.gather_chunk_kernel\\)" crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs`. Expected: 3 sites.
+
+- [ ] **Step 2: At each launcher, replace the `aux_dir_prob_null: u64 = 0` arg** with `let aux_dir_prob_ptr = self.prev_aux_dir_prob.raw_ptr();` and pass `&aux_dir_prob_ptr` instead. Update the inline arg-order docstring to reflect non-NULL.
+
+- [ ] **Step 3: Verify compile.**
+
+### Task D6: gpu_dqn_trainer.rs — call setters on the evaluator post-construction
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`
+
+- [ ] **Step 1: Find where the trainer constructs or hands off the evaluator** via `grep -nE "GpuBacktestEvaluator|backtest_evaluator|evaluator\\." crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`.
+
+- [ ] **Step 2: After the evaluator is fully initialized, call both setters.** `evaluator.set_w_aux_ptr(self.w_aux_to_q_dir.raw_ptr());` and `evaluator.set_aux_weights(...)` (Task D7 adds the second setter).
+
+- [ ] **Step 3: Verify compile.**
+
+### Task D7: gpu_backtest_evaluator.rs — share trainer's aux trunk weight set
+
+**Files:**
+- Modify: `crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs`
+- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`
+
+- [ ] **Step 1: Add `set_aux_weights(...)` setter** on the evaluator accepting the trainer's aux trunk weight buffer dev_ptrs (one per buffer per Task D1's discovery: input projection, hidden, output projection at minimum).
+
+- [ ] **Step 2: Add matching `aux_*_dev_ptr: u64` fields** to the evaluator struct (one per weight buffer). Initialize to 0 in struct construction.
+
+- [ ] **Step 3: Call the setter from the trainer** after the evaluator is constructed (already framed in D6 Step 2).
+
+- [ ] **Step 4: Replace placeholder `&self.aux_weights_set` in Task D3 Step 3** with actual dispatching using the per-buffer dev_ptr fields. The aux trunk forward (option a or b per D1) reads these ptrs.
+
+- [ ] **Step 5: Verify compile.** Expected: 0 errors. With Tasks D1–D7 complete, the eval pipeline has aux trunk forward + alpha + non-NULL state-gather wiring + beta producer.
+
+---
+
+## Phase E — Verification gates
+
+All gates must pass before the Phase F commit.
+
+### Task E1: cargo check
+
+- [ ] **Step 1:** Run `SQLX_OFFLINE=true cargo check -p ml --features cuda`. Expected: 0 errors, warning count near Phase 2 baseline (21) + possibly a few new unused-import warnings. Resolve any new ERROR; warnings can defer.
+
+### Task E2: gpu_backtest_validation
+
+- [ ] **Step 1:** Run `SQLX_OFFLINE=true cargo test -p ml --test gpu_backtest_validation --features cuda --release`. Expected: 6/6 pass (post-test-fix baseline from `71eab9a25`). The 4 previously-passing should remain. Potential numerical shifts in extended-metrics assertions (CVaR/Omega) only if scale_beta is non-zero at test-time. At test-time scale_beta is sentinel 0 (no SP11 controller running), so beta is a no-op → metrics bit-identical expected.
+
+### Task E3: compute-sanitizer
+
+- [ ] **Step 1:** Run `/usr/local/cuda/bin/compute-sanitizer --tool=memcheck --error-exitcode=42 $(ls -t target/release/deps/gpu_backtest_validation-* | grep -v '\\.d$' | head -1) --test-threads=1 2>&1 | tail -10`. Expected: `ERROR SUMMARY: 0 errors`. Non-zero exit from the Rust test harness is fine (test failures don't trigger the exit-code-42 path).
+
+### Task E4: Captured-graph integrity (deferred to smoke)
+
+- [ ] **Step 1:** Document the post-smoke check. The smoke log should show `CAPTURE_PHASE_AUX_DONE`, `CAPTURE_PHASE_POST_AUX_DONE`, `CAPTURE_PHASE_FORWARD_DONE`, and the parent-graph compose line. Failure to capture = `pearl_no_host_branches_in_captured_graph` violation; debug the kernel for any host branch.
+
+### Task E5: alpha weight movement check (deferred to smoke)
+
+- [ ] **Step 1:** Document protocol. After smoke cycle 1, verify `W_aux_to_Q_dir` magnitudes are non-zero. If `[0, 0, 0, 0]`, backward gradient flow is broken.
+
+### Task E6: 7-component contract sanity (deferred to smoke)
+
+- [ ] **Step 1:** Document protocol. Smoke cycle 1 HEALTH_DIAG `sp11_reward` line shows 7 weights including `w_aux=…`; `reward_split` line shows 7 components including `aux_align=…`. Both `w_aux` and `aux_align` non-zero within epoch 1 after first SP11 controller emit. If still zero → either the controller emit is broken OR the beta producer isn't firing.
+
+### Task E7: A2 eval-side aux activity (deferred to validation eval after epoch 1)
+
+- [ ] **Step 1:** Document protocol. Post-epoch-1 per-cycle validation eval shows non-zero `r_aux_align` in eval-side `WindowMetrics` reward decomposition. If eval `r_aux_align ≈ 0` while training-side > 0 → A2 wiring failure.
+
+---
+
+## Phase F — Audit doc + atomic commit + push + smoke + verdict
+
+### Task F1: Append H6 Phase 3 entry to audit doc
+
+**Files:**
+- Modify: `docs/dqn-wire-up-audit.md`
+
+- [ ] **Step 1: Locate the audit doc tail** via `tail -5 docs/dqn-wire-up-audit.md`.
+
+- [ ] **Step 2: Append the H6 Phase 3 entry** using Edit tool with `replace_all=false` on the last paragraph as anchor. The entry covers: framing (refines Phase 2 verdict's mechanistic finding), the four in-scope additions (alpha + beta + 7-component + SP11 controller + A2), out-of-scope statement (aux trunk stop-grad property), verification-gate placeholders to be filled by the smoke run, verdict criteria.
+
+- [ ] **Step 3: Fill verification-gate placeholders** with actual results from Phase E.
+
+### Task F2: Single atomic commit
+
+- [ ] **Step 1: Verify changes via `git status -s`.** All Phase A–F1 file modifications should be unstaged.
+
+- [ ] **Step 2: Stage explicitly (no `git add -A`)** the ~19 modified or new files listed in the File Map plus the audit doc.
+
+- [ ] **Step 3: Commit with comprehensive message** covering all the items: alpha (new kernels + trainer plumbing + Adam), beta (7th component producer), 7-component contract migration cascade, SP11 controller extension, A2 eval-side aux trunk + alpha + beta, verification gate results, verdict criteria, pearl references.
+
+- [ ] **Step 4: Verify via `git log --oneline -3`.**
+
+### Task F3: Push
+
+- [ ] **Step 1:** Run `git push`. Per `feedback_push_before_deploy`, must succeed before smoke dispatch.
+
+### Task F4: Dispatch smoke
+
+- [ ] **Step 1:** Run `./scripts/argo-train.sh dqn --baseline --branch sp20-aux-h-fixed --epochs 3 --gpu-pool ci-training-l40s`.
+
+- [ ] **Step 2: Capture workflow name** `train-<id>` from output for monitoring + audit doc fill-in.
+
+### Task F5: Monitor for verdict
+
+- [ ] **Step 1: Start a background log tail** via `argo logs -f train-<id> -n foxhunt --no-color 2>&1 | tee /tmp/train-<id>.log` with `run_in_background: true`. Per `feedback_no_redundant_monitor`: one channel only.
+
+- [ ] **Step 2: Wait for cycle 1 + HEALTH_DIAG** via the canonical grep-anchored bg waiter pattern from prior phases.
+
+- [ ] **Step 3: Interpret verdict per spec criteria.** Compute WR = wins / (wins + losses) from epoch-1 trade stats. Check W_aux_to_Q_dir magnitudes, a_var per-branch, action distribution, w_aux + aux_align HEALTH_DIAG. Apply the spec's verdict tree: WR > 50.5% AND sub-branches gradient-coupled → confirmed; WR pinned + W non-zero + action distribution moved → bypass works but signal not actionable; WR pinned + W near-zero → debug gradient flow.
+
+- [ ] **Step 4: Update audit doc with verdict + close Phase 3 narrative.** Replace workflow-name placeholder. Add "Phase 3 smoke verdict" subsection. Commit + push.
+
+---
+
+## Self-review
+
+| Spec requirement | Plan task(s) |
+|---|---|
+| alpha 4-weight W_aux_to_Q_dir + Adam states | A3–A5, B9 Steps 4–9, C1 |
+| alpha forward kernel (post-Q_dir bias) | A3, B9 Step 7, C1 Step 4, D4 |
+| alpha backward kernel (dW + dstate) | A4, B9 Step 8 |
+| build.rs cubin registration | A5 |
+| W shared trainer ↔ collector | C1 Step 5, D6 |
+| 7-component contract migration training side | B1, B3–B11 |
+| 7-component contract migration eval side | B2, D5, F2 commit |
+| beta producer training-side | B1 Step 4 |
+| beta producer eval-side | B2 Step 4 |
+| New ISV slots | A1, A2 |
+| state_reset_registry entries | A2 |
+| SP11 controller w_aux_align emit | B6 |
+| HEALTH_DIAG sp11_reward + reward_split 7-entry print | B8, B9 Steps 2–3, B11 |
+| reward_component_ema 0..7 | B3 |
+| reward_decomp_diag column 6 | B4 |
+| reward_component_mag_ratio axis 6 | B5 |
+| health_diag_kernel snap slot | B7 |
+| reward_component_monitor 7-component reader | B10 |
+| A2.1 aux trunk weight loading at eval | D1, D7 |
+| A2.2 per-window prev_aux_dir_prob buffer | D2 |
+| A2.3 per-step aux trunk forward at eval + softmax-to-per-env | D3 |
+| A2.4 alpha at eval | D4 |
+| A2.5 beta at eval (covered by B2) | B2 |
+| A2.6 non-NULL state-gather wiring | D5 |
+| Encoder-state[121] gradient routing option (i) — both paths open | A4 dstate kernel accumulates |
+| Verification gate 1 — cargo check | E1 |
+| Verification gate 2 — gpu_backtest_validation | E2 |
+| Verification gate 3 — compute-sanitizer | E3 |
+| Verification gate 4 — captured-graph integrity (post-smoke) | E4 |
+| Verification gate 5 — alpha weight movement (post-smoke) | E5 |
+| Verification gate 6 — 7-component contract sanity (post-smoke) | E6 |
+| Verification gate 7 — A2 eval-side aux activity (post-smoke) | E7 |
+| Audit doc append (Invariant 7) | F1 |
+| Atomic commit | F2 |
+| Push | F3 |
+| Smoke dispatch | F4 |
+| Verdict monitoring + audit doc verdict | F5 |
+| Out of scope: aux trunk stop-grad property | N/A |
+
+All spec requirements have plan task assignments. No orphan tasks. Type consistency: `W_aux_to_Q_dir`, `aux_to_q_dir_bias_kernel`, `aux_to_q_dir_bias_backward_dw_kernel`, `aux_to_q_dir_bias_backward_dstate_kernel`, `SP22_AUX_ALIGN_SCALE_INDEX`, `REWARD_AUX_ALIGN_EMA_INDEX`, `r_aux_align`, `prev_aux_dir_prob` referenced consistently across tasks.