feat(sp22): H6 Phase 3a — 7-component contract migration + β producer (WIP)

Phase 3a builds on the Phase A foundation (464bc5f7a). Migrates the
7-component reward_components_per_sample contract atomically across
producer + readers + buffer alloc, and installs the β producer at
training-side trade-close. α kernels exist (compiled in Phase A) but
are not yet launched in captured graphs — that's Phase 3b alongside
the SP11 controller extension and A2 eval-side aux infrastructure.

Why split into 3a/3b
────────────────────
Full Phase 3 (α + β + SP11 controller + A2) is ~25-35 hr engineering
spanning ~19 files. Phase 3a is the SAFE atomic contract migration
(7-stride buffer + β producer; no α captured-graph integration yet)
— runtime-equivalent to Phase 2 (β no-op at scale_β=0 sentinel; α
kernels loaded but never launched). This commits the foundation +
contract change as a clean checkpoint per
`feedback_no_partial_refactor` (the 7-component contract spans every
consumer; partial migration would produce stride mismatches; this
commit migrates ALL consumers that read the buffer).

Files
─────
- crates/ml/src/cuda_pipeline/experience_kernels.cu:
    Preamble doc → 7-component layout.
    Buffer stride `* 6 +` → `* 7 +` (~14 sites, atomic).
    7th-slot init at the per-step zero block (`rc[6] = 0.0f`).
    β producer at the segment_complete branch:
      r_aux_align = scale_β * max(0, aux × pos_sign) * max(0, pnl)
      with NULL-safe fallbacks (aux_dir_prob_per_env NULL OR
      isv_signals_ptr NULL → β no-op).
    Two new kernel args: aux_dir_prob_per_env + aux_align_scale_idx
    (slot index for scale_β, decoupled per the loss_cap_idx pattern).

- crates/ml/src/cuda_pipeline/reward_component_ema_kernel.cu:
    Stride `idx * 6` → `idx * 7` (3 sites). Iteration stays c=0..5;
    the 7th component (aux_align) is intentionally NOT EMA'd here.
    A dedicated reward_aux_align_ema_kernel writing directly to
    ISV[REWARD_AUX_ALIGN_EMA_INDEX=536] is Phase 3b scope (avoids
    extending the apply_pearls_ad chain).
    Preamble doc updated.

- crates/ml/src/cuda_pipeline/reward_decomp_diag_kernel.cu:
    #define RCP_NUM_COMPS 6 → 7. The kernel's per-bin abs-sum
    (col 3) now naturally includes r_aux_align; popart/micro/
    opp_cost per-bin means unchanged.

- crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu:
    Documentation only: aux_align excluded from the 6-axis
    cf_others ratio (non-contiguous with cf_others_base_slot at
    64..68; aux_align EMA lives at ISV[536]).

- crates/ml/src/cuda_pipeline/gpu_experience_collector.rs:
    Buffer alloc `total_output * 6` → `* 7` (critical for runtime
    safety — partial migration would produce OOB writes since
    experience_env_step writes to `out_off * 7 + N`).
    experience_env_step launcher gains 2 new `.arg(...)` calls
    passing `self.prev_aux_dir_prob.raw_ptr()` and
    `SP22_AUX_ALIGN_SCALE_INDEX as i32`.

- docs/dqn-wire-up-audit.md:
    Phase 3a entry documenting the partial commit + Phase 3b
    remaining-work breakdown.

Verification
────────────
- cargo check -p ml --features cuda: 0 errors, 21 pre-existing
  warnings (Phase 2 baseline parity).
- All nvcc cubins recompile (experience_kernels, reward_component_ema,
  reward_decomp_diag, reward_component_mag_ratio_compute, plus
  Phase A's aux_to_q_dir_bias_kernel + backward).
- Runtime equivalent to Phase 2: β no-op (scale_β=0 sentinel since
  SP11 controller not yet extended), α no-op (kernels dead-code
  until Phase 3b wires them into captured graphs).

Phase 3b scope (resume in fresh session)
────────────────────────────────────────
- B6: SP11 controller extension (w_aux_align emit at ISV[537])
- B7: HEALTH_DIAG snap layout extension
- B9: α plumbing — W_aux_to_Q_dir param + Adam + captured-graph
       forward + backward integration in gpu_dqn_trainer.rs
- B10/B11: HEALTH_DIAG print-line extensions
- C1: α forward in collector's rollout-time captured graph
- D1-D7: A2 eval-side aux trunk + α + state-gather wiring
- E + F: verification gates + atomic Phase F commit + smoke + verdict

Refs
────
- docs/plans/2026-05-12-sp22-h6-phase3-alpha-beta.md (spec)
- docs/plans/2026-05-13-sp22-h6-phase3-alpha-beta-runbook.md (runbook)
- 464bc5f7a (Phase A foundation)
- pearl_no_partial_refactor (atomic 7-component contract migration)
- pearl_event_driven_reward_density_alignment (β at segment_complete)
- pearl_one_unbounded_signal_per_reward (β bounded by scale_β + alignment caps)

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

crates/ml/src/cuda_pipeline/experience_kernels.cu

@@ -2051,16 +2051,27 @@ extern "C" __global__ void experience_env_step(
     const float* __restrict__ magnitude_conviction_ptr,
     /* C.2 Reward-component attribution (Plan 3 Task 1, spec §4.C.2).
      * Per-sample per-component reward magnitudes written at every (i,t).
-     * Layout [N*L, 6] row-major; component indices:
+     * Layout [N*L, 7] row-major (SP22 H6 Phase 3 contract migration);
-     *   [0] popart   — final on-policy reward (denominator for PopArt drift)
+     * component indices:
-     *   [1] cf       — counterfactual reward written to CF slot
+     *   [0] popart    — final on-policy reward (denominator for PopArt drift)
-     *   [2] trail    — 0.0 (placeholder; future term)
+     *   [1] cf        — counterfactual reward written to CF slot
-     *   [3] micro    — dense OFI micro-reward (R5)
+     *   [2] trail     — 0.0 (placeholder; future term)
-     *   [4] opp_cost — Flat opportunity-cost reward (Plan 3 Task 2 B.1)
+     *   [3] micro     — dense OFI micro-reward (R5)
-     *   [5] bonus    — novelty-scaled trade-attempt bonus (Plan 3 Task 3 B.2)
+     *   [4] opp_cost  — Flat opportunity-cost reward (Plan 3 Task 2 B.1)
+     *   [5] bonus     — novelty-scaled trade-attempt bonus (Plan 3 Task 3 B.2)
+     *   [6] aux_align — SP22 H6 Phase 3 β: event-driven aux-aligned trade-close
+     *                   bonus. Fires at `segment_complete` when policy direction
+     *                   aligns with aux conviction AND trade was profitable.
+     *                   `r_aux_align = scale_β × alignment × profit_pos` where
+     *                   `alignment = max(0, aux_at_close × position_sign)` and
+     *                   `profit_pos = max(0, realized_pnl)`. scale_β read from
+     *                   ISV[SP22_AUX_ALIGN_SCALE_INDEX] (controller-emitted).
+     *                   NULL-safe via two fallbacks: aux_dir_prob_per_env NULL
+     *                   OR isv_signals_ptr NULL → β no-op.
      * Producer: this kernel (experience_env_step).
-     * Consumer: reward_component_ema GPU kernel → ISV[63..69). */
+     * Consumer: reward_component_ema GPU kernel → ISV[63..69) and
-    float* __restrict__ reward_components_per_sample, /* [N*L * 6] */
+     * ISV[REWARD_AUX_ALIGN_EMA_INDEX=536]. */
+    float* __restrict__ reward_components_per_sample, /* [N*L * 7] */
     /* B.2 Plan 3 Task 3: Flat→Positioned transition flag per (i,t).
      * Written unconditionally (0/1) at every reached slot. Consumer:
      * trade_attempt_rate_ema_update → ISV[TRADE_ATTEMPT_RATE_EMA_INDEX]. */
@@ -2310,7 +2321,26 @@ extern "C" __global__ void experience_env_step(
      * gather (sample_proportional → GpuBatchPtrs.aux_conf_ptr), and
      * eventual Phase 5 Aux→Q gate read in the Bellman-target loss
      * computation. */
-    float* __restrict__ aux_conf_per_sample
+    float* __restrict__ aux_conf_per_sample,
+    /* SP22 H6 Phase 3 β (2026-05-13) — per-env recentered aux directional
+     * probability from the PREVIOUS step's aux trunk forward. Layout [N]
+     * f32. Same buffer as `experience_state_gather`'s arg of the same name
+     * — both are populated by `aux_softmax_to_per_env_kernel`. Read here
+     * at the `segment_complete` branch to compute the alignment factor
+     * for `r_aux_align = scale_β × max(0, aux_at_close × position_sign)
+     * × max(0, realized_pnl)`. NULL-safe: when NULL (test scaffolds, or
+     * pre-aux-wireup), `aux_at_close = 0` → `alignment = 0` → β no-op.
+     * Identical to Phase 2 recentered encoding: 0 = neutral, +1 = up
+     * with full conviction, -1 = down with full conviction. */
+    const float* __restrict__ aux_dir_prob_per_env,
+    /* SP22 H6 Phase 3 β — ISV slot index for the SP11-controller-emitted
+     * adaptive scale_β (canonical SP22_AUX_ALIGN_SCALE_INDEX = 537).
+     * Passed as a kernel arg per the existing pattern (mirrors
+     * `loss_cap_idx`, `alpha_ema_idx`, `hold_cost_scale_idx`,
+     * `hold_reward_ema_idx`) so the kernel stays decoupled from SP slot
+     * drift. The kernel reads `isv_signals_ptr[aux_align_scale_idx]`
+     * NULL-safely (returns 0 if isv_signals_ptr == NULL → β no-op). */
+    int aux_align_scale_idx
 ) {
     int i = blockIdx.x * blockDim.x + threadIdx.x;
     if (i >= N) return;
@@ -2446,7 +2476,7 @@ extern "C" __global__ void experience_env_step(
      * write pattern self-documenting and forward-safe. Components [4, 5] are
      * overwritten later by the Flat opp-cost and Flat→Positioned bonus paths. */
     {
-        float* rc = reward_components_per_sample + out_off * 6;
+        float* rc = reward_components_per_sample + out_off * 7;
         rc[0] = 0.0f;  /* popart (final reward) — overwritten below */
         rc[1] = 0.0f;  /* cf reward — overwritten below at cf_off */
         rc[2] = 0.0f;  /* trail penalty (future term) */
@@ -3296,12 +3326,17 @@ extern "C" __global__ void experience_env_step(
      * Note: rc[+0] (PopArt input = total reward) is unconditionally
      * overwritten at ~L3447, so it doesn't strictly need zero-init here, but
      * including it removes ordering dependency for future readers. */
-    reward_components_per_sample[out_off * 6 + 0] = 0.0f;
+    reward_components_per_sample[out_off * 7 + 0] = 0.0f;
-    reward_components_per_sample[out_off * 6 + 1] = 0.0f;
+    reward_components_per_sample[out_off * 7 + 1] = 0.0f;
-    reward_components_per_sample[out_off * 6 + 2] = 0.0f;
+    reward_components_per_sample[out_off * 7 + 2] = 0.0f;
-    reward_components_per_sample[out_off * 6 + 3] = 0.0f;
+    reward_components_per_sample[out_off * 7 + 3] = 0.0f;
-    reward_components_per_sample[out_off * 6 + 4] = 0.0f;
+    reward_components_per_sample[out_off * 7 + 4] = 0.0f;
-    reward_components_per_sample[out_off * 6 + 5] = 0.0f;
+    reward_components_per_sample[out_off * 7 + 5] = 0.0f;
+    /* SP22 H6 Phase 3 β (2026-05-13): 7th reward component slot.
+     * Overwritten below at segment_complete with `scale_β × alignment
+     * × profit_pos` when policy direction aligns with aux conviction
+     * AND trade was profitable; otherwise stays at 0 (β no-op). */
+    reward_components_per_sample[out_off * 7 + 6] = 0.0f;
 
     /* Trade-cumulative running scalar — tracks the in-progress reward
      * across the if/else cascade. C.4/D.4b bonus blocks read this to
@@ -3504,7 +3539,7 @@ extern "C" __global__ void experience_env_step(
          * trail-fire branch fires (was a structural placeholder pre-B1b).
          * The reward_component_ema kernel will pick up real trail magnitude
          * for ISV-driven component-mag-ratio canaries. */
-        reward_components_per_sample[out_off * 6 + 2] = r_trail;
+        reward_components_per_sample[out_off * 7 +2] = r_trail;
 
         /* SP12 v3 Change 3: lump-sum opportunity cost on trade close.
          *
@@ -3545,7 +3580,52 @@ extern "C" __global__ void experience_env_step(
         r_opp_cost = compute_lump_sum_opp_cost(
             shaping_scale, holding_cost_rate, pre_trade_position, segment_hold_time
         );
-        reward_components_per_sample[out_off * 6 + 4] = r_opp_cost;
+        reward_components_per_sample[out_off * 7 +4] = r_opp_cost;
+
+        /* ── SP22 H6 Phase 3 β (2026-05-13): aux-aligned trade-close bonus ──
+         *
+         * Fires at segment_complete when policy direction aligns with aux
+         * conviction AND trade was profitable. Non-negative-only (no
+         * anti-alignment penalty; let r_trail / step_ret_core's negative
+         * carry the loss signal — let r_opp_cost's negative carry
+         * cost-dominance). Bounded by `pearl_one_unbounded_signal_per_reward`:
+         * `alignment ∈ [0, 1]`, `scale_β` controller-bound to [0.05, 2.0],
+         * `step_ret_core` is the natural per-trade scale (single unbounded
+         * multiplicand per the pearl).
+         *
+         * NULL-safe via two fallbacks: `aux_dir_prob_per_env == NULL` (test
+         * scaffolds) → `aux_at_close = 0` → `alignment = 0` → β no-op;
+         * `isv_signals_ptr == NULL` (same path) → `scale_β = 0` → β no-op.
+         *
+         * `pre_trade_position` is the position held BEFORE the exit
+         * (`unified_env_step_core` mutates `position` in-place; the saved
+         * pre-step value is what defines the direction of the trade that
+         * just closed). `step_ret_core` is the realized P&L of this
+         * trade-close (the segment return).
+         *
+         * Pearl notes:
+         *   - `pearl_event_driven_reward_density_alignment` — β fires
+         *     ONLY at segment_complete (event-driven density matches the
+         *     trade-close objective; not per-step shaping).
+         *   - `pearl_one_unbounded_signal_per_reward` — only step_ret_core
+         *     is unbounded; alignment + scale_β are bounded.
+         *   - `pearl_first_observation_bootstrap` — scale_β cold-starts at
+         *     0 (β no-op until SP11 controller's first emit). */
+        {
+            float position_sign = (pre_trade_position > 0.001f) ? 1.0f
+                                : (pre_trade_position < -0.001f) ? -1.0f
+                                : 0.0f;
+            float aux_at_close = (aux_dir_prob_per_env != NULL)
+                               ? aux_dir_prob_per_env[i]
+                               : 0.0f;
+            float alignment = fmaxf(0.0f, aux_at_close * position_sign);
+            float profit_pos = fmaxf(0.0f, step_ret_core);
+            float scale_beta_isv = (isv_signals_ptr != NULL)
+                                 ? isv_signals_ptr[aux_align_scale_idx]
+                                 : 0.0f;
+            float r_aux_align = scale_beta_isv * alignment * profit_pos;
+            reward_components_per_sample[out_off * 7 + 6] = r_aux_align;
+        }
 
         /* ── Layer 3: CEA counterfactual loop REMOVED (4-branch: direction(4) replaces exposure(9)).
          * Dense micro-reward (Gem 1) below provides directional gradient signal. ── */
@@ -3606,7 +3686,7 @@ extern "C" __global__ void experience_env_step(
             r_bonus += timing_bonus;
             reward  += timing_bonus;  /* keep cascade scalar in lockstep for Q-cap reads below */
             /* C.2 attribution — accumulate into rc[5] bonus slot (see above). */
-            reward_components_per_sample[out_off * 6 + 5] += timing_bonus;
+            reward_components_per_sample[out_off * 7 +5] += timing_bonus;
         }
 
         /* ── D.4a (Plan 3 Task 6a): Persistence credit ──
@@ -3627,7 +3707,7 @@ extern "C" __global__ void experience_env_step(
                                           * tanhf(reward / fmaxf(1e-4f, drawdown_depth));
                 r_bonus += persist_bonus;
                 reward  += persist_bonus;  /* keep cascade scalar in lockstep for Q-cap reads below */
-                reward_components_per_sample[out_off * 6 + 5]  += persist_bonus;
+                reward_components_per_sample[out_off * 7 +5]  += persist_bonus;
             }
         }
 
@@ -3679,7 +3759,7 @@ extern "C" __global__ void experience_env_step(
                                 * bars_late_frac * reward_capped;
             r_bonus -= penalty;
             reward  -= penalty;  /* keep cascade scalar in lockstep for Q-cap reads below */
-            reward_components_per_sample[out_off * 6 + 5]  -= penalty;
+            reward_components_per_sample[out_off * 7 +5]  -= penalty;
         }
         /* Reset shift-bar for the NEXT trade (in addition to Step 3's lifecycle resets,
          * this handles the exit → next-entry transition cleanly). */
@@ -3881,7 +3961,7 @@ extern "C" __global__ void experience_env_step(
         }
         reward  = r_micro;
         micro_reward_per_sample[out_off] = r_micro;
-        reward_components_per_sample[out_off * 6 + 3] = r_micro;
+        reward_components_per_sample[out_off * 7 +3] = r_micro;
     } else if (!segment_complete) {
         /* Plan 3 D.4c: Welford-style EMA of conviction_core during Flat bars.
          * Updates running mean and squared-deviation EMAs that the consumer
@@ -3986,7 +4066,7 @@ extern "C" __global__ void experience_env_step(
             }
         }
         reward = r_opp_cost;
-        reward_components_per_sample[out_off * 6 + 4] = r_opp_cost;
+        reward_components_per_sample[out_off * 7 +4] = r_opp_cost;
     }
 
     /* ── B.2 (Plan 3 Task 3): ISV-driven trade-attempt bonus ──
@@ -4033,7 +4113,7 @@ extern "C" __global__ void experience_env_step(
             r_bonus            += bonus_scaled;
             reward             += bonus_scaled;  /* keep cascade scalar in lockstep */
             /* C.2 attribution: record bonus component magnitude at rc[5]. */
-            reward_components_per_sample[out_off * 6 + 5] = bonus_scaled;
+            reward_components_per_sample[out_off * 7 +5] = bonus_scaled;
         }
     }
 
@@ -4083,7 +4163,7 @@ extern "C" __global__ void experience_env_step(
             const float bonus        = shaping_scale * vol_proxy_c * stability * conviction_core;
             r_bonus                                        += bonus;
             reward                                         += bonus;  /* keep cascade scalar in lockstep */
-            reward_components_per_sample[out_off * 6 + 5]  += bonus;
+            reward_components_per_sample[out_off * 7 +5]  += bonus;
         }
     }
     /* Defensive reset of D.4c EMA slots at entry — redundant with the
@@ -4355,7 +4435,7 @@ extern "C" __global__ void experience_env_step(
      * to c51_loss_kernel.cu::block_bellman_project_f. */
     total_reward_per_sample[out_off] = reward;
     /* C.2: component [0] = final on-policy reward (PopArt input denominator). */
-    reward_components_per_sample[out_off * 6 + 0] = reward;
+    reward_components_per_sample[out_off * 7 +0] = reward;
 
     /* SP11 Fix 39 (A1.2): per-bar saboteur engagement signal — see param
      * docstring above. Computed after `total_reward_per_sample` is
@@ -4580,7 +4660,7 @@ extern "C" __global__ void experience_env_step(
          * (popart-component) is fed by raw r_popart via
          * popart_component_per_sample. Only rc[+0] is intentionally
          * post-composition (PopArt input). */
-        reward_components_per_sample[out_off * 6 + 1] = cf_reward;
+        reward_components_per_sample[out_off * 7 +1] = cf_reward;
     }
 
     /* ---- Advance timestep counter ---- */

crates/ml/src/cuda_pipeline/gpu_experience_collector.rs

@@ -2042,8 +2042,13 @@ impl GpuExperienceCollector {
         // [total_output * 6] row-major: [popart, cf, trail, micro, opp_cost, bonus] per (i,t).
         // Producer: experience_env_step (writes rc[0], rc[1], rc[3], rc[4], rc[5]; rc[2]=0 placeholder).
         // Consumer: reward_component_ema GPU kernel → ISV[63..69).
+        // SP22 H6 Phase 3 (2026-05-13): bumped 6 → 7 for the new aux_align
+        // reward component (β producer at trade-close in experience_env_step).
+        // Stride must match `experience_kernels.cu`'s `out_off * 7 +` writes;
+        // partial migration would produce OOB writes (kernel writes slot 6
+        // past row end of a 6-wide buffer). Per `feedback_no_partial_refactor`.
         let reward_components_per_sample = stream
-            .alloc_zeros::<f32>(total_output * 6)
+            .alloc_zeros::<f32>(total_output * 7)
             .map_err(|e| MLError::ModelError(format!("alloc reward_components_per_sample: {e}")))?;
         // B.2 Plan 3 Task 3: per-sample Flat→Positioned transition flag.
         // 1 iff (prev_sign == 0 && curr_sign != 0) for this (i,t), 0 otherwise.
@@ -6605,6 +6610,25 @@ impl GpuExperienceCollector {
                     // end of `collect_experiences_gpu`, mirroring the
                     // `aux_sign_labels` `[N × L × 2]` final shape).
                     .arg(&mut self.aux_conf_per_sample)
+                    // SP22 H6 Phase 3 β (2026-05-13): two new args for the
+                    // event-driven aux-aligned trade-close bonus.
+                    // - aux_dir_prob_per_env: same per-env recentered p_up
+                    //   buffer that experience_state_gather reads (Phase 1+2
+                    //   H6 wiring); β reads at segment_complete to compute
+                    //   alignment. Reuses the existing prev_aux_dir_prob
+                    //   buffer; no new allocation needed.
+                    // - aux_align_scale_idx: ISV slot index for the
+                    //   SP11-controller-emitted adaptive scale_β. In
+                    //   Phase 3a (this commit) the SP11 controller has not
+                    //   yet been extended, so the slot stays at sentinel 0
+                    //   and β is functionally a no-op (r_aux_align = 0).
+                    //   Phase 3b lands the controller extension to bring
+                    //   scale_β non-zero.
+                    .arg(&self.prev_aux_dir_prob.raw_ptr())
+                    .arg(&{
+                        use crate::cuda_pipeline::sp22_isv_slots::SP22_AUX_ALIGN_SCALE_INDEX;
+                        SP22_AUX_ALIGN_SCALE_INDEX as i32
+                    })
                     .launch(launch_cfg)
                     .map_err(|e| MLError::ModelError(format!(
                         "experience_env_step t={t}: {e}"

crates/ml/src/cuda_pipeline/reward_component_ema_kernel.cu

@@ -14,7 +14,11 @@
  * `popart_component_per_sample` GPU buffer (B1b fix-up).
  *
  * Reads per-sample per-component rewards from `reward_components_per_sample`
- * [n_samples * 6], where the 6 components are laid out as:
+ * [n_samples * 7] (SP22 H6 Phase 3 stride change; this kernel still
+ * processes components 0..5 only, the 7th component aux_align is
+ * handled by the separate `reward_aux_align_ema_kernel` writing
+ * directly to ISV[REWARD_AUX_ALIGN_EMA_INDEX=536]), where the
+ * 6 components processed here are laid out as:
  *
  *   [c=0] popart  — final on-policy reward (pre-PopArt; denominator signal)
  *   [c=1] cf      — counterfactual replay reward written to the CF slot
@@ -22,6 +26,9 @@
  *   [c=3] micro   — dense OFI-alignment micro-reward (R5)
  *   [c=4] opp_cost — Flat opportunity cost (Plan 3 Task 2 B.1; ISV[21]-scaled)
  *   [c=5] bonus    — 0.0 placeholder (future trade-attempt/timing bonus)
+ *   [c=6] aux_align — SP22 H6 Phase 3 (NOT processed by this kernel;
+ *                     see `reward_aux_align_ema_kernel` for the EMA
+ *                     into ISV[REWARD_AUX_ALIGN_EMA_INDEX=536])
  *
  * Reduces mean |r_c| over the batch per component, then writes the step
  * observation to `scratch_buf[scratch_first_index + c]` for c in 0..6.
@@ -54,7 +61,7 @@
  */
 
 extern "C" __global__ void reward_component_ema(
-    const float* __restrict__ reward_components, /* [n_samples * 6] row-major */
+    const float* __restrict__ reward_components, /* [n_samples * 7] row-major (SP22 H6 Phase 3 stride; kernel reads only c=0..5) */
     int          n_samples,
     float*       __restrict__ scratch_buf,        /* producer_step_scratch_buf */
     int          scratch_first_index,             /* slot 63 — first of 6 contiguous mag slots */
@@ -75,7 +82,7 @@ extern "C" __global__ void reward_component_ema(
     /* ── Pass 1: mean |r| over batch for this component. ────────────── */
     float acc = 0.0f;
     for (int idx = 0; idx < n_samples; ++idx) {
-        acc += fabsf(reward_components[idx * 6 + c]);
+        acc += fabsf(reward_components[idx * 7 + c]);
     }
     float mean_abs = acc / fmaxf(1.0f, (float)n_samples);
 
@@ -88,7 +95,7 @@ extern "C" __global__ void reward_component_ema(
     if (c >= 1 && c <= 5) {
         float v_acc = 0.0f;
         for (int idx = 0; idx < n_samples; ++idx) {
-            float d = fabsf(reward_components[idx * 6 + c]) - mean_abs;
+            float d = fabsf(reward_components[idx * 7 + c]) - mean_abs;
             v_acc += d * d;
         }
         float var_abs = v_acc / fmaxf(1.0f, (float)n_samples);

crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu

@@ -18,6 +18,17 @@
 //   axis 4 (opp_cost) = isv[cf_others_base_slot + 3]   ← slot 67
 //   axis 5 (bonus)    = isv[cf_others_base_slot + 4]   ← slot 68
 //
+// SP22 H6 Phase 3 (2026-05-13): aux_align (component idx 6 in the
+// 7-stride reward_components_per_sample buffer) is INTENTIONALLY
+// excluded from this kernel's 6-axis ratio computation. The aux_align
+// EMA lives at ISV[REWARD_AUX_ALIGN_EMA_INDEX=536] (non-contiguous
+// with cf_others_base_slot at 64..68), and the SP11 controller's
+// `w_aux_align` adaptive weight has its own anchor on that slot
+// directly (see reward_subsystem_controller_kernel.cu Phase 3 extension).
+// Keeping aux_align out of this kernel's ratio preserves the
+// pre-Phase-3 6-axis mag/var z-score machinery bit-for-bit; only the
+// SP11 controller's emit list gains the 7th weight.
+//
 // And per-component variance sources (B1b smoke-recovery):
 //
 //   axis 0 var = isv[popart_var_slot]                  ← slot 361

crates/ml/src/cuda_pipeline/reward_decomp_diag_kernel.cu

@@ -95,7 +95,11 @@
 #define RCP_POPART_IDX   0
 #define RCP_MICRO_IDX    3
 #define RCP_OPP_COST_IDX 4
-#define RCP_NUM_COMPS    6
+/* SP22 H6 Phase 3 (2026-05-13): bumped 6 → 7 for the new aux_align
+ * reward component (c=6). The kernel's abs-sum aggregate (col 3
+ * `mean(|reward|)`) now naturally includes r_aux_align in the per-bin
+ * total — popart/micro/opp_cost per-bin means remain unchanged. */
+#define RCP_NUM_COMPS    7
 
 /* Output layout: 4 bins × 5 cols = 20 floats. */
 #define SP18_DECOMP_NUM_BINS  4

docs/dqn-wire-up-audit.md

@@ -16743,3 +16743,92 @@ Phase B partial work (file diffs for `experience_kernels.cu` +
 `/tmp/sp22-h6-phase3-b-partial.patch` (300 lines, local-only). Next
 session can apply if useful as a starting point, or discard and
 re-derive directly from the runbook.
+
+### Phase 3a — partial Phase B + reward-component contract migration
+
+Built on the Phase A foundation (`464bc5f7a`). Phase 3a migrates the
+7-component reward_components_per_sample contract atomically and
+installs the β producer at training-side trade-close. α kernels
+exist (Phase A cubins) but are NOT yet wired into the captured
+training/rollout graphs — that lands in Phase 3b alongside the SP11
+controller extension and the A2 eval-side aux infrastructure.
+
+Phase 3a files (5):
+
+- `crates/ml/src/cuda_pipeline/experience_kernels.cu`:
+  - Preamble doc updated to 7-component layout
+  - `[N*L * 6]` → `[N*L * 7]` (kernel-arg comment + every
+    `reward_components_per_sample[out_off * 6 + N]` → `* 7 + N`,
+    ~14 sites, atomic via `replace_all`)
+  - New 7th slot init (`rc[6] = 0.0f` at the per-step zero block)
+  - β producer at the `segment_complete` branch:
+    `r_aux_align = scale_β × max(0, aux × position_sign) × max(0, realized_pnl)`
+    with NULL-safe fallbacks (aux_dir_prob_per_env NULL or
+    isv_signals_ptr NULL → β no-op).
+  - Two new kernel args: `const float* aux_dir_prob_per_env` and
+    `int aux_align_scale_idx` (slot index for scale_β, per the
+    `loss_cap_idx`/`hold_cost_scale_idx` decoupling pattern).
+
+- `crates/ml/src/cuda_pipeline/reward_component_ema_kernel.cu`:
+  - Stride `idx * 6` → `idx * 7` (3 sites). Kernel still iterates
+    c=0..5; the 7th component aux_align is intentionally NOT
+    EMA'd here (a dedicated `reward_aux_align_ema_kernel` writing
+    directly to ISV[REWARD_AUX_ALIGN_EMA_INDEX=536] is Phase 3b
+    scope to avoid extending the apply_pearls_ad chain).
+  - Preamble doc updated to document the 7-stride + aux_align
+    explicit non-processing.
+
+- `crates/ml/src/cuda_pipeline/reward_decomp_diag_kernel.cu`:
+  - `#define RCP_NUM_COMPS 6` → `7`. The kernel's abs-sum aggregate
+    (col 3 `mean(|reward|)`) now naturally includes r_aux_align in
+    the per-bin total; popart/micro/opp_cost per-bin means
+    unchanged.
+
+- `crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu`:
+  - Documentation-only: aux_align is intentionally excluded from
+    this kernel's 6-axis cf_others ratio computation (the
+    aux_align EMA lives at ISV[536], non-contiguous with the
+    cf_others_base_slot range at 64..68).
+
+- `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs`:
+  - Buffer alloc bump: `total_output * 6` → `* 7` for
+    `reward_components_per_sample`. Stride must match
+    experience_env_step's `out_off * 7 + N` writes per
+    `feedback_no_partial_refactor` (partial migration → OOB writes).
+  - experience_env_step launcher: two new `.arg(...)` calls passing
+    `self.prev_aux_dir_prob.raw_ptr()` (the existing Phase 1 H6
+    buffer) and `SP22_AUX_ALIGN_SCALE_INDEX as i32`.
+
+Functional state post-Phase-3a:
+
+- Compile + cubin clean (cargo check 0 errors, 21 pre-existing
+  warnings parity; nvcc compiles all kernels).
+- Runtime equivalent to Phase 2: β is a no-op at sentinel
+  (scale_β=0 since SP11 controller hasn't been extended to emit
+  the new w_aux_align), so r_aux_align stays at 0. α kernels are
+  "dead code" — loaded by the trainer's static CUBIN consts but
+  never launched (captured-graph integration is Phase 3b).
+
+Phase 3a NOT yet implemented (Phase 3b scope):
+
+- B6: SP11 controller kernel extension for w_aux_align emit
+  (block_dim 10→11, slot conflicts at ISV[350], apply_pearls_ad
+  chain refactor). Without this, scale_β stays at sentinel 0 →
+  β no-op.
+- B7: HEALTH_DIAG snap layout slot for r_aux_align EMA.
+- B9 (α plumbing): trainer-side α param `W_aux_to_Q_dir` + Adam
+  states + α forward kernel launch in training captured forward
+  graph + α backward kernels (dw + dstate) in captured backward
+  graph + Adam-step update for W.
+- B10: reward_component_monitor.rs 7-component EMA reader.
+- B11: training_loop.rs sp11_reward + reward_split print-line
+  extension to include `w_aux=...` and `aux_align=...`.
+- C1: α forward in collector's rollout-time captured graph.
+- D1-D7: A2 eval-side aux trunk forward + α + state-gather
+  wiring + W/aux-trunk-ptr setters trainer↔evaluator.
+- E + F: verification gates + atomic Phase F commit + smoke +
+  verdict.
+
+Estimated remaining work: ~25-35 hr engineering. Resume in a fresh
+session from Phase 3a commit per the runbook
+`docs/plans/2026-05-13-sp22-h6-phase3-alpha-beta-runbook.md`.