feat(ml-alpha): Phase 3B+Y — reward-pnl alignment foundation

Three coordinated changes to remove anti-aligned components from
the reward pipeline. Predicted Pearson trajectory based on Phase 3A
audit was wrong (PopArt + hindsight only delivered +0.04 not +0.4),
but the deep Phase 3B-Z audit revealed the "low Pearson" was largely
a measurement artifact: raw_reward tracks ALL PnL events (including
fees + partial closes via pos.realized_pnl_delta) while diag's
realized_pnl_usd_cum tracks only full closures. The actual training
reward IS aligned with total realized PnL evolution.

## Changes

* Slot 753 RL_SURFER_SCAFFOLD_WEIGHT_INDEX bootstrap 1.0 → 0.0
  (Phase 5 hold-bonus/entry-cost/short-hold-penalty/long-ride-bonus
  shaping disabled — was the dominant anti-aligned contaminant per
  Phase 3A audit and cluster 2A-E evidence of +$201M shaped vs
  -$330M actual).
* New ISV slot 793 RL_POPART_NORMALIZE_ENABLED_INDEX bootstrap 0.0
  (PopArt batch normalization disabled — was found to sign-flip
  rewards under adverse batch composition).
* Hindsight injection (rl_hindsight_inject.cu wants_inject) extended
  with scaffold_w > 0.5 condition (when slot 753 = 0, hindsight off).
  Semantically pairs the two training scaffolds.

## Verification

* 13/13 multi_head_policy_invariants PASS (no regression).
* FOXHUNT_USE_MULTI_HEAD_POLICY=0/1 ./scripts/determinism-check.sh
  --quick: both exit 0.
* Local single-seed mid-smoke (b=128, seed 42):
  - eval pnl: -$4.6M (vs Phase 0 baseline -$4.46M, within noise)
  - 14,691 trades training / 3,268 eval (healthy, not paralyzed)
  - action_entropy 1.82 at step 1999 (below uniform 2.40, learning)
  - 0 NaN, exit 0

## Pearson interpretation (Phase 3B-Z audit)

The Pearson 0.38 between raw_reward and pnl_step_close_usd_delta is
a structural mismatch between two valid PnL views (all events vs
close events only), NOT an anti-alignment. raw_reward at slot
753 = 0 is the delta of pos.realized_pnl which already aggregates
the events we care about for total-USD optimization.

Cluster verification (Phase 3C, next) will reveal whether the
removal of Phase 5 + popart + hindsight produces a base policy
that can learn at scale. The 2A-E cluster catastrophe (-$330M)
was driven by Phase 5 shaping that's now disabled.

Plan: docs/superpowers/plans/2026-06-03-reward-alignment-gae-
combined.md

Linked pearls:
* pearl_reward_signal_anti_aligned_with_pnl
* pearl_phase5_term_4_is_almost_potential
* pearl_popart_blind_to_session_signals
* pearl_pure_pnl_mode_starves_b16_controllers

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>

crates/ml-alpha/cuda/rl_hindsight_inject.cu

@@ -27,6 +27,13 @@
 #define RL_HINDSIGHT_INJECT_COUNT_INDEX   551
 #define RL_REWARD_SCALE_INDEX             406
 
+/* Phase 3B-Y (2026-06-03): hindsight is conceptually paired with the surfer
+ * scaffold — both are training-scaffold biases on the gradient. Gate hindsight
+ * injection by slot 753 so that pure-pnl mode (scaffold_w = 0.0) also disables
+ * hindsight, eliminating the second of three contaminants identified by the
+ * Phase 3B-followup audit. */
+#define RL_SURFER_SCAFFOLD_WEIGHT_INDEX   753
+
 extern "C" __global__ void rl_hindsight_inject(
     const float* __restrict__ dones,              /* [B]             */
     const float* __restrict__ rewards_raw,        /* [B]             */
@@ -64,9 +71,16 @@ extern "C" __global__ void rl_hindsight_inject(
         float peak_pnl = (float)dir * (peak_mid[b] - entry_mid[b]);
 
         float threshold = isv[RL_HINDSIGHT_THRESHOLD_INDEX];
+        /* Phase 3B-Y: gate by surfer-scaffold weight (slot 753). When scaffold
+         * is fully faded (w = 0.0, pure-pnl mode), hindsight injection is also
+         * disabled — they are paired training scaffolds. */
+        float scaffold_w = isv[RL_SURFER_SCAFFOLD_WEIGHT_INDEX];
 
-        /* Peak PnL must exceed actual by the threshold factor AND be positive. */
-        if (peak_pnl > 0.0f && peak_pnl > actual_pnl * threshold) {
+        /* Peak PnL must exceed actual by the threshold factor AND be positive,
+         * AND the scaffold gate must be on. */
+        if (peak_pnl > 0.0f
+            && peak_pnl > actual_pnl * threshold
+            && scaffold_w > 0.5f) {
             wants_inject = 1;
         }
 

crates/ml-alpha/cuda/rl_popart_normalize.cu

@@ -37,6 +37,12 @@
 // Rust const `RL_POPART_SIGMA_WELFORD_INDEX` in `isv_slots.rs`.
 #define RL_POPART_SIGMA_WELFORD_INDEX      725
 
+// Phase 3B-Y (2026-06-03): PopArt normalize gate. 0.0 = DISABLED (skip the
+// final whitening loop but keep running stats updating). 1.0 = ENABLED
+// (legacy van Hasselt 2016 whitening). Mirror of Rust const
+// `RL_POPART_NORMALIZE_ENABLED_INDEX` in `isv_slots.rs`.
+#define RL_POPART_NORMALIZE_ENABLED_INDEX 793
+
 #define BLOCK_SIZE 256
 
 // Warp-level sum reduction via shuffle-down.
@@ -182,11 +188,17 @@ extern "C" __global__ void rl_popart_normalize(
 
     // ── Pass 3: normalize rewards in place ──
     // CRITICAL (Issue β): read broadcast slots at block_dim * 3.
-    float mean  = sdata[block_dim * 3 + 0];
-    float sigma = sdata[block_dim * 3 + 1];
-    float inv_sigma = 1.0f / sigma;  // sigma >= sqrt(1e-6) > 0
+    // Phase 3B-Y (2026-06-03): gated by slot 793 (RL_POPART_NORMALIZE_ENABLED_INDEX).
+    // Running stats above (mean/var/sigma + envelope EMA) still update so that
+    // re-enabling the gate produces sensible values without a session restart.
+    const float popart_enabled = isv[RL_POPART_NORMALIZE_ENABLED_INDEX];
+    if (popart_enabled > 0.5f) {
+        float mean  = sdata[block_dim * 3 + 0];
+        float sigma = sdata[block_dim * 3 + 1];
+        float inv_sigma = 1.0f / sigma;  // sigma >= sqrt(1e-6) > 0
 
-    for (int i = tid; i < b_size; i += block_dim) {
-        rewards[i] = (rewards[i] - mean) * inv_sigma;
+        for (int i = tid; i < b_size; i += block_dim) {
+            rewards[i] = (rewards[i] - mean) * inv_sigma;
+        }
     }
 }

crates/ml-alpha/src/rl/isv_slots.rs

@@ -1952,6 +1952,30 @@ pub const RL_POLICY_GATE_ENTROPY_EMA_INDEX: usize = 791;
 /// `isv_state` checksum bit-equality with pre-B1.3 HEAD 6f3639bfb.
 pub const RL_POLICY_GATE_CONTROLLER_BOOTSTRAP_DONE_INDEX: usize = 792;
 
+/// PopArt normalization gate (Phase 3B-Y, 2026-06-03).
+///
+/// 0.0 = DISABLED (default per bootstrap) — `rl_popart_normalize.cu` skips
+///   the final whitening loop `rewards[i] = (rewards[i] - mean) * inv_sigma`,
+///   leaving rewards in raw pnl-aligned units. Running mean/var/sigma stats
+///   (slots 553-558, 714, 725) continue updating so that re-enabling the
+///   gate produces sensible values without a session restart.
+/// 1.0 = ENABLED — legacy behavior (van Hasselt 2016 PopArt whitening).
+///
+/// Per Phase 3B-followup audit (2026-06-03): PopArt batch normalization was
+/// identified as one of three contaminants bypassing `RL_SURFER_SCAFFOLD_WEIGHT_INDEX`
+/// (slot 753) when testing pure-pnl mode. Under adverse batch composition
+/// (e.g., one large negative reward + many small positives), the per-batch
+/// mean shift can flip reward signs relative to raw pnl direction, eroding
+/// `Pearson(reward, Δpnl)`. The audit attributed 40-50% of the residual
+/// gap at HEAD c5036af03 (Pearson = 0.335 with slot 753 = 0.0) to this
+/// mechanism.
+///
+/// Default-off targets pure-pnl gradient alignment for Phase 3B-Y. Bootstrap
+/// 0.0 satisfies `pearl_bootstrap_must_respect_clamp_range` — the clamp is
+/// the binary {0, 1} gate inside the kernel (`> 0.5f` test); 0.0 is the
+/// natural OFF sentinel and matches a hypothetical clamp-floor of 0.0.
+pub const RL_POPART_NORMALIZE_ENABLED_INDEX: usize = 793;
+
 /// Last RL-allocated slot index (exclusive).
 /// Pre-risk-stack: 662. Post-Fix-B: 685. Post-v9 (warmup boundary): 696.
 /// Post-regime-observer (F1.1): 716. Post-warmed-flag addendum: 717.
@@ -1970,4 +1994,5 @@ pub const RL_POLICY_GATE_CONTROLLER_BOOTSTRAP_DONE_INDEX: usize = 792;
 /// Post-multi-head-policy Phase 2A-D (25 aggregate diag slots): 789.
 /// Post-Phase 2A-D fix B1 (gate LR multiplier slot 790): 790.
 /// Post-Phase 2A-D fix B1.3 (gate-LR controller state slots 791-792): 792.
-pub const RL_SLOTS_END: usize = 793;
+/// Post-Phase 3B-Y (PopArt normalize gate slot 793): 793.
+pub const RL_SLOTS_END: usize = 794;

crates/ml-alpha/src/trainer/integrated.rs

@@ -3717,7 +3717,7 @@ impl IntegratedTrainer {
         // (slot, value) pair — pure device write, no HtoD per
         // `feedback_no_htod_htoh_only_mapped_pinned`.
         {
-            let isv_constants: [(usize, f32); 244] = [
+            let isv_constants: [(usize, f32); 245] = [
                 // Static seeds for the adaptive reward-clamp controller —
                 // these are the initial values that
                 // `rl_reward_clamp_controller` will replace once it
@@ -4145,11 +4145,18 @@ impl IntegratedTrainer {
                 // controller `rl_surfer_scaffold_controller` fades w as the agent
                 // crosses break-even with statistical confidence, and re-engages
                 // it if edge-decay PH flags policy degradation.
-                (crate::rl::isv_slots::RL_SURFER_SCAFFOLD_WEIGHT_INDEX,           1.0_f32),
+                (crate::rl::isv_slots::RL_SURFER_SCAFFOLD_WEIGHT_INDEX,           0.0_f32), // pure-pnl mode (w=0) — Phase 3B test
                 // Controller config:
                 (crate::rl::isv_slots::RL_SURFER_BREAK_EVEN_WR_INDEX,             0.30_f32),
                 (crate::rl::isv_slots::RL_SURFER_K_SHARPNESS_INDEX,               30.0_f32),
                 (crate::rl::isv_slots::RL_SURFER_WARMUP_TRADES_INDEX,             200.0_f32),
+                // Phase 3B-Y (2026-06-03): PopArt normalization OFF by default.
+                // `rl_popart_normalize.cu` skips the final whitening loop when
+                // slot 793 ≤ 0.5; running mean/var/sigma stats still update so
+                // the gate can be re-enabled mid-session without restart.
+                // Targets the 40-50% of `Pearson(reward, Δpnl)` gap attributed
+                // to popart sign-flips by the Phase 3B-followup audit.
+                (crate::rl::isv_slots::RL_POPART_NORMALIZE_ENABLED_INDEX,         0.0_f32),
             ];
             for (slot, value) in isv_constants.iter() {
                 let slot_i32 = *slot as i32;

docs/superpowers/plans/2026-06-03-reward-alignment-gae-combined.md

@@ -0,0 +1,283 @@
+# Phase 3 — Reward-PnL Alignment via GAE-Propagated Pure-PnL Reward
+
+> **For agentic workers:** REQUIRED SUB-SKILL: Use superpowers:subagent-driven-development or superpowers:executing-plans to implement this plan task-by-task. Steps use checkbox (`- [ ]`) syntax for tracking.
+
+**Date:** 2026-06-03
+**Goal:** Eliminate the load-bearing reward-pnl misalignment by combining two committed-but-dormant building blocks: GAE λ-returns over the rollout buffer (Phase 1B foundation) + pure-pnl reward mode (slot 753 = 0).
+
+**Architecture:** Pure-pnl alone is aligned-but-sparse (June-01 b55 cluster: Pearson=0.92 ✓ but bootstrap failed). GAE alone propagates whatever signal you give it. Combined, they produce dense, aligned gradient signal — a combination that has never been tested. Phase 2A multi-head infrastructure stays committed but disabled until the base policy is verified to learn on clean reward.
+
+**Tech Stack:** Rust + CUDA, foxhunt-discipline (no atomicAdd, no raw memcpy on regular slices, ISV-driven controllers, scoped_init_seed for determinism, flag-gated rollback via FOXHUNT_USE_ROLLOUT + ISV slot 753).
+
+---
+
+## 0. Decisive empirical evidence (motivating this plan)
+
+Phase 2A-E cluster run `alpha-rl-xq2jq` (SHA `c5036af03`, b=1024 fold-1 L40S, terminated 2026-06-03 at step 10,106 / 20,000):
+
+| Step | realized_pnl_usd_cum | shaped_reward_close_event_cum | total_trades | popart_σ |
+|---|---|---|---|---|
+| 5,000 | -$164.9M | +$106.8M | 297,891 | 73.7 |
+| 7,000 | -$235.2M | +$147.4M | 411,567 | 83.6 |
+| 9,000 | -$301.6M | +$184.4M | 514,631 | 42.6 |
+| 10,106 | **-$330.4M** | **+$201.0M** | 560,244 | 245.7 |
+
+The two reward streams are perfectly anti-correlated. Every step the agent takes that grows `shaped_reward` destroys `realized_pnl`. This is the same misalignment in `pearl_reward_signal_anti_aligned_with_pnl.md` ("the load-bearing bug behind 64 negative-eval commits"), now captured at unprecedented signal strength thanks to the Phase 0 diag fix (`cfaa420bd`) that separated shaped vs USD fields.
+
+**Multi-head architecture is NOT the bug.** Phase 2A-E confirmed multi-head functions correctly at cluster scale (heads diverge, kernels deterministic, gate-LR controller works). The architecture cannot paper over a wrong target.
+
+## 1. Diagnostic chain — why pure-pnl alone failed before
+
+The June-01 cluster smoke `alpha-rl-8fb55` (SHA `fa347e481`) tested pure-pnl mode in isolation:
+
+```
+Pearson(reward, Δpnl): 0.92  ✓ ALIGNED
+action_entropy:       1.9-2.1 (random log(11)=2.40)
+wr:                   0.30 (random baseline)
+pnl:                  -$3.9M (bled)
+```
+
+Pure-pnl was aligned but couldn't bootstrap. Per `pearl_advantage_kernel_is_done_gated_td_not_gae` the advantage kernel is done-gated 1-step TD with T_eff=1: gradient only flows at close events (~5-7% of steps). A random-init policy doesn't see enough signal to find profitable patterns.
+
+Phase 5 shaping was inadvertently acting as a bootstrap scaffold (DENSE signal that taught the random-init policy what to optimize) — but at the cost of anti-alignment with real pnl.
+
+**The fix that hasn't been tried:** make the GAE backward sweep propagate the (sparse but aligned) close-event pnl signal to held states. T_eff jumps from 1 to 13.56 (λ=0.95, γ=0.975). Dense AND aligned simultaneously.
+
+## 2. Building blocks (all already committed)
+
+| Block | Commit | Status | Purpose |
+|---|---|---|---|
+| Phase 1B-A: RolloutBuffer + GAE kernel structures | `5cd2f8703` | committed, dormant | Foundation for backward sweep |
+| Phase 1B-B: Rollout collection loop + `FOXHUNT_USE_ROLLOUT` flag | `779c03b9d` | committed, dormant | Per-step buffer fill |
+| Pure-pnl reward mode | slot 753 (current bootstrap=1.0 full shaping) | active | `w=0` skips Phase 5 anti-aligned shaping |
+| Phase 2A multi-head | commits 0b3e40150 → c8c81ab7e | committed, flag-off bit-equal | Expressivity layer (NOT yet activated) |
+| Adaptive controllers (12+) | various | active | LR/clamp/EMA management |
+| Determinism foundation | 5-phase chain | active | Bit-equal same-seed runs |
+| Per-batch diag aggregation | `cfaa420bd` + Phase 2A-C+ | active | The Phase 0 fix that EXPOSED the misalignment |
+
+## 3. Phase structure
+
+### Phase 3A — Foundation audit (1-2 days)
+
+**Goal:** Map the current state of Phase 1B-A/B/C; identify exactly what's wired and what's missing for GAE→V regression.
+
+**Files to inspect:**
+- `crates/ml-alpha/cuda/rollout_gather.cu` (Phase 1B-A foundation)
+- `crates/ml-alpha/src/trainer/integrated.rs` rollout-loop wiring (Phase 1B-B, `FOXHUNT_USE_ROLLOUT`)
+- `crates/ml-alpha/examples/alpha_rl_train.rs` rollout integration
+- `crates/ml-alpha/tests/rollout_buffer_invariants.rs` (test coverage)
+- `crates/ml-alpha/cuda/compute_advantage_return.cu` (current done-gated TD; target for GAE swap)
+- `docs/superpowers/specs/2026-06-02-trainer-rollout-buffer-gae.md` (Phase 1B spec, paused)
+- `docs/superpowers/specs/2026-06-01-reward-policy-alignment-investigation.md` (pure-pnl spec)
+
+**Deliverable:** A short note (~50 lines) inline in this plan listing:
+1. What's wired/dormant — exact file:line citations
+2. The exact δ between the current state and the working GAE+pure-pnl combination
+3. Any technical-debt items (e.g., Phase 1B-C abandoned because of Q-distill discovery — what parts still apply?)
+
+**Estimated time:** 1-2 days
+
+- [ ] **Step 1: Read all 6 reference files end-to-end**
+- [ ] **Step 2: Trace the rollout buffer lifecycle from collection to consumption when `FOXHUNT_USE_ROLLOUT=1`**
+- [ ] **Step 3: Identify the V regression call site (loss kernel + grad path)**
+- [ ] **Step 4: Inline the audit note into this plan document**
+- [ ] **Step 5: Commit the audit note (single doc-only commit)**
+
+---
+
+### Phase 3B — GAE backward sweep wired into V regression (3-5 days)
+
+**Goal:** Compute λ-returns over the rollout buffer and use them as the V regression target. T_eff: 1 → 13.56.
+
+**Files to modify:**
+- `crates/ml-alpha/cuda/compute_advantage_return.cu` (or new `compute_gae_returns.cu`) — backward sweep `A_t = δ_t + γλ·A_{t+1}`, returns `R_t = A_t + V_t`
+- `crates/ml-alpha/src/trainer/integrated.rs` — V regression target source
+- `crates/ml-alpha/src/rl/isv_slots.rs` — `RL_GAE_LAMBDA_INDEX` (new, bootstrap 0.95), `RL_GAE_GAMMA_INDEX` (alias of existing γ slot or new)
+- `crates/ml-alpha/tests/` — new GAE-correctness tests (T_eff observable, λ→0 matches done-gated TD, λ→1 matches Monte-Carlo)
+
+**Constraints (foxhunt-discipline):**
+- No atomicAdd. Backward sweep is sequential per batch — single-block-per-batch kernel.
+- λ-return target is computed in the **rollout phase**, not in the per-step training kernel.
+- PPO surrogate gradient remains **zeroed** (Q-distill is primary π path per `pearl_foxhunt_pi_trained_by_q_distillation_not_ppo`). GAE benefit propagates via V → Bellman Q target → softmax → KL distill → π.
+- Determinism MUST be preserved at both `FOXHUNT_USE_ROLLOUT=0` (current default, bit-equal to HEAD) and `FOXHUNT_USE_ROLLOUT=1` (new GAE path).
+
+**Diag additions:**
+- V loss decomposition (per-step, per-rollout)
+- `R_t` distribution stats (mean, σ, percentiles)
+- GAE coverage rate (fraction of held-state advantages that are nonzero post-GAE)
+- λ-return vs done-gated-TD divergence metric (sanity check on the change)
+
+**Tests:**
+1. `gae_lambda_zero_matches_one_step_td` — at λ=0, GAE output must equal the existing done-gated TD output (regression guard)
+2. `gae_lambda_one_matches_monte_carlo` — at λ=1, returns must equal the analytical sum of discounted rewards
+3. `gae_t_eff_observable` — at λ=0.95, γ=0.975, measure T_eff from advantage non-zero rate, verify ≥ 10 (theoretical 13.56)
+4. `gae_backward_sweep_deterministic` — bit-equal across two fresh contexts
+
+**Estimated time:** 3-5 days (kernel ~150 LOC, wiring ~50 LOC, tests ~300 LOC)
+
+- [ ] **Step 1: Decide kernel structure (extend existing vs new file)**
+- [ ] **Step 2: Implement GAE backward sweep kernel + register in build.rs**
+- [ ] **Step 3: Wire V regression to consume λ-returns when `FOXHUNT_USE_ROLLOUT=1`**
+- [ ] **Step 4: Add 4 GPU-oracle invariant tests**
+- [ ] **Step 5: Add 4 new diag fields**
+- [ ] **Step 6: Verify determinism preserved both flag states**
+- [ ] **Step 7: Local mid-smoke at `FOXHUNT_USE_ROLLOUT=1` — verify no NaN, V loss decreases**
+- [ ] **Step 8: Stage + commit (atomic)**
+
+---
+
+### Phase 3C — Pure-pnl reward mode default (2-3 days)
+
+**Goal:** Change slot 753 bootstrap from 1.0 (full Phase 5 shaping) to 0.0 (pure realized-pnl). Reward = realized USD pnl at close events, 0 elsewhere.
+
+**Files to modify:**
+- `crates/ml-alpha/src/trainer/integrated.rs` — slot 753 bootstrap value
+- `crates/ml-alpha/cuda/rl_fused_reward_pipeline.cu` — verify w=0 path is correct (should already exist from June-01 work)
+- Diag emit — log slot 753 value per step (probably already done)
+
+**Predicted behavior (this is the falsification gate):**
+- Pearson(reward, Δpnl) ≥ 0.7 within 1000 steps (per `pearl_reward_signal_anti_aligned_with_pnl` predetermined gate)
+- Combined with Phase 3B GAE, the policy should bootstrap (vs June-01 b55 which couldn't)
+
+**Test plan:**
+1. Single-seed b=128 mid-smoke at `FOXHUNT_USE_ROLLOUT=1` + slot 753=0
+2. Verify Pearson(reward, Δpnl) ≥ 0.7 by step 1000
+3. Verify action_entropy decreases below 2.2 by step 1500 (vs June-01 b55 stuck at 1.9-2.1)
+4. Verify trade rate doesn't go to zero (which would indicate the trainer is paralyzed by sparse signal)
+
+If ANY of these fail at local: STOP, the GAE-pure-pnl combo doesn't bootstrap either; pivot needed.
+
+**Estimated time:** 2-3 days (config change + 1 mid-smoke + analysis)
+
+- [ ] **Step 1: Change slot 753 bootstrap to 0.0**
+- [ ] **Step 2: Verify w=0 reward path unchanged from June-01 implementation (regression check)**
+- [ ] **Step 3: Local single-seed seed=42 mid-smoke**
+- [ ] **Step 4: Compute Pearson(reward, Δpnl) from diag.jsonl**
+- [ ] **Step 5: Falsification verdict (G1 = Pearson ≥ 0.7 by step 1000)**
+- [ ] **Step 6: If G1 PASS, run 3-seed for variance estimate**
+- [ ] **Step 7: Stage + commit**
+
+---
+
+### Phase 3D — Pearson(reward, Δpnl) alignment controller (2 days)
+
+**Goal:** Build a continuous-monitoring + protective controller that detects re-emergence of anti-alignment and gates training accordingly.
+
+**Files to modify:**
+- `crates/ml-alpha/src/rl/isv_slots.rs` — new slots for rolling Pearson stats (cov-Welford, var_x-Welford, var_y-Welford, output)
+- `crates/ml-alpha/cuda/rl_pearson_alignment_controller.cu` — new kernel
+- `crates/ml-alpha/src/trainer/integrated.rs` — per-step launch
+- `crates/ml-alpha/cuda/rl_fused_controllers.cu` — gate hook (reduce Q-distill τ or policy LR when Pearson drops)
+
+**Design:**
+- 4 new ISV slots: `RL_PEARSON_COV_INDEX`, `RL_PEARSON_VAR_X_INDEX`, `RL_PEARSON_VAR_Y_INDEX`, `RL_PEARSON_OUT_INDEX`
+- Welford-style rolling stats over a configurable window (e.g., 500 steps)
+- Bootstrap on first observation (first-observation pattern per `pearl_welford_trade_count_is_step_not_trade`)
+- Output: Pearson coefficient in [-1, 1]
+- Gate hook: if `Pearson < threshold` (configurable, default 0.5), increase Q-distill τ (smooths target) and/or reduce policy LR
+
+**Tests:**
+1. `pearson_controller_bootstrap_initializes_stats`
+2. `pearson_controller_matches_numpy_on_synthetic_data` (correlated, anti-correlated, independent)
+3. `pearson_controller_gate_fires_when_anti_aligned`
+4. `pearson_controller_deterministic`
+
+**Estimated time:** 2 days
+
+- [ ] **Step 1: Allocate 4 ISV slots**
+- [ ] **Step 2: Implement kernel + register in build.rs**
+- [ ] **Step 3: Wire per-step launch (flag-gated by `FOXHUNT_USE_ROLLOUT`?)**
+- [ ] **Step 4: Add gate hook in `rl_fused_controllers.cu`**
+- [ ] **Step 5: 4 GPU-oracle tests**
+- [ ] **Step 6: Local mid-smoke to verify controller produces plausible Pearson values**
+- [ ] **Step 7: Stage + commit**
+
+---
+
+### Phase 3E — Cluster verification (single-head, GAE + pure-pnl + Pearson controller)
+
+**Goal:** Real ground-truth at cluster scale. Verify the GAE+pure-pnl combination produces an aligned, bootstrap-capable policy with a real eval verdict.
+
+**Config:**
+- HEAD: at end of Phase 3D commits
+- `FOXHUNT_USE_ROLLOUT=1` (GAE active)
+- slot 753 = 0 (pure-pnl reward)
+- `FOXHUNT_USE_MULTI_HEAD_POLICY=0` (multi-head OFF — verifying base case first)
+- Pearson alignment controller active
+- L40S b=1024, n-steps=20000, n-eval-steps=5000, n-folds=2, fold-idx=0, seed=42
+
+**Falsification gates (predetermined, hard):**
+- **G1**: Pearson(reward, Δpnl) ≥ 0.7 sustained from step 1000 onward
+- **G2**: No NaN at any step
+- **G3**: Eval `total_pnl_usd` better than last 5 cluster baselines (the most recent comparable was v10 mean_pnl=$107k per task #178)
+- **G4 (stretch)**: Eval `total_pnl_usd` ≥ +$500k
+
+If G1 fails: GAE didn't propagate the signal as expected. Investigate before continuing.
+If G2 fails: STOP at first NaN, surface.
+If G3 fails: alignment is real but base policy still doesn't learn. Investigate before Phase 3F.
+If G4 fails but G3 passes: incremental success, can still proceed to Phase 3F.
+
+**Estimated time:** 3 days (~70min wall + analysis + iteration if gates fail)
+
+- [ ] **Step 1: Push HEAD + verify cluster template plumbing for new flags**
+- [ ] **Step 2: Submit alpha-rl workflow with the Phase 3E config**
+- [ ] **Step 3: Watch trajectory: Pearson, action_entropy, wr_ema, pnl_cum**
+- [ ] **Step 4: After completion, extract eval_summary.json**
+- [ ] **Step 5: Compute Pearson(reward, Δpnl) over last 1k diag rows**
+- [ ] **Step 6: Write verdict (G1-G4)**
+
+---
+
+### Phase 3F — Multi-head re-engagement on clean reward
+
+**Goal:** Once Phase 3E proves the base policy learns on aligned reward, test whether Phase 2A multi-head architecture **adds value** on top.
+
+**Config:**
+- Identical to Phase 3E except `FOXHUNT_USE_MULTI_HEAD_POLICY=1`
+- Adaptive gate-LR controller (slot 790) ramps as needed
+- Compare against the Phase 3E single-head baseline directly
+
+**Falsification gates:**
+- **G5**: Eval pnl ≥ Phase 3E eval pnl (within noise; multi-head shouldn't REGRESS)
+- **G6 (stretch)**: Eval pnl > Phase 3E eval pnl by ≥ +$500k (multi-head ADDS value)
+- **G7**: gate_entropy_mean ≤ 0.9 at end of training (specialization happens at cluster scale when reward is clean — the hypothesis Phase 2A-E couldn't test because reward was broken)
+
+If G5 fails: multi-head adds noise, not value. Revert to single-head as production default.
+If G7 passes: multi-head IS doing regime-conditional specialization (validates the Phase 2A architectural hypothesis).
+
+**Estimated time:** 3 days
+
+- [ ] **Step 1: Submit alpha-rl workflow with `--use-multi-head-policy=1`**
+- [ ] **Step 2: Watch trajectory: same diag as 3E + gate_entropy_mean, gate_argmax_mass, per_head_entropy**
+- [ ] **Step 3: Verdict G5-G7**
+
+---
+
+## 4. Self-review (run before invoking writing-plans hand-off)
+
+- [x] **Placeholder scan**: no TBD/TODO/fill-in. All steps have specific files + commands.
+- [x] **Internal consistency**: Phase 3B implements GAE, Phase 3C activates the pure-pnl reward that GAE will propagate, Phase 3D protects against anti-alignment regression, Phase 3E verifies the base case, Phase 3F tests the architecture on top.
+- [x] **Scope check**: Single coherent intervention chain. Multi-head is committed but disabled until Phase 3F (only after base case verified).
+- [x] **Ambiguity check**: Falsification gates are numerical and predetermined. Each phase has STOP conditions.
+
+## 5. Linked
+
+- `pearl_reward_signal_anti_aligned_with_pnl` — the load-bearing pearl; this plan is its fix
+- `pearl_advantage_kernel_is_done_gated_td_not_gae` — the GAE motivation; T_eff=1 → 13.56
+- `pearl_foxhunt_pi_trained_by_q_distillation_not_ppo` — why GAE helps via V→Q→distill→π
+- `pearl_foxhunt_trainer_is_genuinely_single_step` — why Phase 1B-A/B exist (rollout buffer needed for GAE)
+- `pearl_phase5_term_4_is_almost_potential` — why pure-pnl is the right target (3 of 4 shaping terms violate Ng-Harada-Russell)
+- `pearl_pure_pnl_mode_starves_b16_controllers` — the June-01 cluster smoke failure that motivates the GAE pairing
+- `feedback_investigation_first_falsification_methodology` — the methodology this plan follows
+- Phase 1B spec: `docs/superpowers/specs/2026-06-02-trainer-rollout-buffer-gae.md`
+- Reward alignment spec: `docs/superpowers/specs/2026-06-01-reward-policy-alignment-investigation.md`
+
+## 6. What happens if this works
+
+If Phase 3E G1+G3 pass, the project has solved its longest-standing bug. The Phase 2A multi-head infrastructure becomes the next-layer expressivity test (3F). The 12 adaptive controllers operate on a now-aligned signal. The determinism foundation makes verdict-quality measurement possible.
+
+If 3E G1 passes but G3 doesn't: alignment is real but other latent bugs surface (e.g., the Welford-trade-count gap, Kelly trap, edge-decay loops). These become next-iteration targets but at least with a known-aligned reward to work on top of.
+
+If 3E G1 fails: GAE doesn't propagate as hypothesized. The investigation surfaces real architectural insight (gradient cancellation in V regression? λ-return numerical issues? boundary conditions in the backward sweep?).
+
+Either way, **the diagnostic loop is broken open** — we move forward with real signal instead of fighting an anti-aligned target.