diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index d1e91cf09..ca1c448e8 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/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: - * [0] popart — final on-policy reward (denominator for PopArt drift) - * [1] cf — counterfactual reward written to CF slot - * [2] trail — 0.0 (placeholder; future term) - * [3] micro — dense OFI micro-reward (R5) - * [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) + * Layout [N*L, 7] row-major (SP22 H6 Phase 3 contract migration); + * component indices: + * [0] popart — final on-policy reward (denominator for PopArt drift) + * [1] cf — counterfactual reward written to CF slot + * [2] trail — 0.0 (placeholder; future term) + * [3] micro — dense OFI micro-reward (R5) + * [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). */ - float* __restrict__ reward_components_per_sample, /* [N*L * 6] */ + * Consumer: reward_component_ema GPU kernel → ISV[63..69) and + * 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 * 6 + 1] = 0.0f; - reward_components_per_sample[out_off * 6 + 2] = 0.0f; - reward_components_per_sample[out_off * 6 + 3] = 0.0f; - reward_components_per_sample[out_off * 6 + 4] = 0.0f; - reward_components_per_sample[out_off * 6 + 5] = 0.0f; + reward_components_per_sample[out_off * 7 + 0] = 0.0f; + reward_components_per_sample[out_off * 7 + 1] = 0.0f; + reward_components_per_sample[out_off * 7 + 2] = 0.0f; + reward_components_per_sample[out_off * 7 + 3] = 0.0f; + reward_components_per_sample[out_off * 7 + 4] = 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 ---- */ diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 246979e89..e0a726b04 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/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::(total_output * 6) + .alloc_zeros::(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}" diff --git a/crates/ml/src/cuda_pipeline/reward_component_ema_kernel.cu b/crates/ml/src/cuda_pipeline/reward_component_ema_kernel.cu index 0c6bdb78c..5adb57074 100644 --- a/crates/ml/src/cuda_pipeline/reward_component_ema_kernel.cu +++ b/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); diff --git a/crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu b/crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu index 332729b17..347129f36 100644 --- a/crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu +++ b/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 diff --git a/crates/ml/src/cuda_pipeline/reward_decomp_diag_kernel.cu b/crates/ml/src/cuda_pipeline/reward_decomp_diag_kernel.cu index 8c074d2d4..d999d50b8 100644 --- a/crates/ml/src/cuda_pipeline/reward_decomp_diag_kernel.cu +++ b/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 diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index f50002aeb..b873f4a36 100644 --- a/docs/dqn-wire-up-audit.md +++ b/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`.