feat(rl): adaptive surfer-scaffold reward shaping (spec v5 A+B combined)

Diagnosis from alpha-rl-8fb55 (fa347e481 b=1024 fold-1, 4500 steps):
the v4 binary `pure_pnl_mode=1` correctly aligned the gradient (Pearson
0.92) but left the agent unable to LEARN. action_entropy stayed at
1.9-2.1 (near uniform log(11)=2.40) — the policy never converged on a
strategy, just oscillated between random trading and all-flat. pnl
bled −$3.9M, wr stuck at 0.30 (random-baseline). The Phase 5 hold-bonus
was NOT pure contamination — it was the INDUCTIVE BIAS SCAFFOLD that
taught a random-init network what to optimize. Pure pnl alone is too
sparse to bootstrap learning in foxhunt's reward-density regime.

REDESIGN — replace binary pure_pnl_mode with continuous adaptive
surfer-scaffold weight w ∈ [0,1] multiplying Phase 5 shaping:
  w=1 → identical to dd049d9a4 surfer-baseline shaping (full scaffold)
  w=0 → identical to v4 pure pnl reward (no shaping)
  0<w<1 → smooth lerp

Controller `rl_surfer_scaffold_controller.cu` writes w each step from:
  - RL_WIN_RATE_EMA_INDEX (677): agent's competence signal
  - RL_CUMULATIVE_DONES_INDEX (660): trade-count confidence gate
  - RL_EDGE_PH_FRAC_ALERTED_INDEX (751): edge-decay re-engagement

Math:
  confidence  = sigmoid((n_trades - warmup) / (warmup*0.3))
  competence  = confidence × sigmoid(k_sharp × max(0, wr_ema - break_even))
  w_competence = 1 - competence
  w_decay     = min(1, 2 × frac_alerted)
  w = clamp(max(w_competence, w_decay), 0, 1)

At random init (n_trades≈0, wr≈0): w → 1 (full scaffold, agent learns
to hold via Phase 5 like dd049d9a4 surfer baseline). After warmup with
wr crossing break-even: w → 0 (pure pnl reward, no Phase 5
contamination for eval-relevant strategies). Edge-decay PH alerts
force w back up if the policy degrades — Option A "rethink strategy"
mechanism without an explicit reset.

Slot reuse: 753 was `RL_REWARD_PURE_PNL_MODE_INDEX` (v4 binary, never
deployed beyond fa347e481 cluster smoke alpha-rl-kwppb/8fb55, both
verdict=fail). v5 renames to `RL_SURFER_SCAFFOLD_WEIGHT_INDEX` with
semantic flip (v5 w=1 ⇔ v4 mode=0; bootstrap 1.0 invariant). 3 new
config slots (754/755/756) for break_even_wr / k_sharpness /
warmup_trades. RL_SLOTS_END 754 → 757.

Phase 5 kernel modifications — multiplicative lerp form preserves
w=1 ⇔ legacy shaping:
  additive a:        r += w × a
  multiplicative m:  r *= (1 + w × (m - 1))

Diag: `rewards.surfer_scaffold_weight` replaces `rewards.pure_pnl_mode`
(1 leaf in, 1 leaf out → EXPECTED_LEAVES still 679).

Invariant test:
  - I1+I2: scaffold_weight ∈ [0, 1] across train + eval (250 rows)
  - I3: scaffold_weight ≥ 0.9 at step 0 (novice agent ⇒ full scaffold)

Local validation:
  reward_alignment_surfer_scaffold_invariants OK: 251 rows validated;
    train[0] bootstrap = 0.966 (matches math: 1 - sigmoid(-3.33)×0.5 ≈ 0.98)
  eval_diag_emission OK: 679 leaves train + eval (schema parity)

Cluster smoke (next): fold-1 b=1024 20k+5k. Falsification:
  PASS: action_entropy drops below 1.7 by step 5000 (policy concentrating)
        AND wr_ema crosses 0.35 by step 10000
        AND surfer_scaffold_weight fades to < 0.5 by step 15000
  FAIL: any criterion below → revisit signal weights or anneal schedule

Spec: docs/superpowers/specs/2026-06-01-reward-policy-alignment-investigation.md §6
Pearls:
  - pearl_reward_signal_anti_aligned_with_pnl (the v4 diagnosis)
  - pearl_pure_pnl_mode_starves_b16_controllers (the v4 cluster verdict)

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

crates/ml-alpha/build.rs

@@ -115,6 +115,7 @@ const KERNELS: &[&str] = &[
     "rl_iqn_action_tau_controller",        // Adaptive risk management (2026-05-30): Layer 2 — adapts IQN action-selection quantile τ from session drawdown signal
     "rl_inventory_beta_controller",        // Adaptive risk management (2026-05-30): Layer 3 — adapts inventory penalty β from observed inventory variance vs reward magnitude
     "rl_kelly_fraction_controller",        // Adaptive risk management (2026-05-30): Layer 4 — half-Kelly position-size multiplier from observed win_rate × R-multiple; warmup-gated
+    "rl_surfer_scaffold_controller",       // Reward-policy realignment v5 (2026-06-01): adaptive Phase 5 shaping weight ∈ [0,1] — fades surfer-bias scaffold as agent crosses break-even, re-engages on edge-decay PH alerts
     "rl_eval_warmup_decay",                // v9 (2026-05-31): defensive eval-boundary calibration — overrides Kelly/τ_min/entropy_min/ε_min during warmup, linearly decays back to normal; runs every step
     "rl_regime_flat_count",                // F1.2 (2026-05-31): block-reduce per-account lots[b] → flat_count single int; prereq for regime_observer (F1.3)
     "rl_regime_observer",                  // F1.3 (2026-05-31): unified regime state-machine — dead-zone, Welford PnL variance, tail-event, recovery_factor/eps_live

crates/ml-alpha/cuda/rl_fused_reward_pipeline.cu

@@ -41,11 +41,14 @@
 #define RL_OUTCOME_ALPHA_INDEX             520
 // Layer 3 (inventory penalty — spec 2026-05-30-adaptive-risk-management).
 #define RL_INVENTORY_PENALTY_BETA_INDEX    674
-// Reward-policy realignment (spec 2026-06-01-reward-policy-alignment-investigation).
-// When > 0.5: skip ALL Phase 5 shaping (entry cost, short-hold penalty,
-// long-ride bonus, per-step hold bonus) AND Phase 5b inventory penalty.
-// Reward = Phase 1 realized_pnl_delta unchanged.
-#define RL_REWARD_PURE_PNL_MODE_INDEX      753
+// Adaptive surfer-scaffold shaping (spec v5 2026-06-01 §6).
+// w ∈ [0, 1] multiplies the Phase 5 shaping contribution:
+//   w = 1 → full shaping (surfer-baseline dd049d9a4 path)
+//   w = 0 → pure realized_pnl_delta reward
+//   0<w<1 → smooth interpolation
+// Driven by rl_surfer_scaffold_controller (reads wr_ema + cumulative
+// dones + edge_decay_frac_alerted; fades w as agent crosses break-even).
+#define RL_SURFER_SCAFFOLD_WEIGHT_INDEX    753
 
 #define MAX_UNITS 4
 
@@ -211,14 +214,18 @@ extern "C" __global__ void rl_fused_reward_pipeline(
     reward_abs[b] = fabsf(reward);
 
     // ================================================================
-    // PHASE 5: rl_reward_shaping (surfer philosophy)
+    // PHASE 5: rl_reward_shaping (adaptive surfer-scaffold, spec v5)
     // ================================================================
-    // Reward-policy alignment (spec 2026-06-01): when pure_pnl_mode > 0.5
-    // ALL of steps 1-4 below + Phase 5b inventory penalty are SKIPPED.
-    // r remains at the Phase 1 realized_pnl_delta — pure pnl alignment.
-    // Default is mode=1 (gating active); mode=0 preserves legacy shaping
-    // for one regression cluster smoke, then deleted in follow-up.
-    const float pure_pnl_mode = isv[RL_REWARD_PURE_PNL_MODE_INDEX];
+    // Adaptive scaffold weight w ∈ [0,1]: each shaping contribution is
+    // multiplied by w so the kernel smoothly interpolates between the
+    // pre-v4 surfer-baseline path (w=1) and pure pnl (w=0). Controller
+    // `rl_surfer_scaffold_controller` writes w per step from wr_ema +
+    // cumulative_dones + edge_decay_frac_alerted.
+    //
+    // LERP forms (preserve w=1 ↔ legacy shaping):
+    //   additive a:  r += w * a               (entry_cost neg, hold_bonus pos)
+    //   multiplicative m:  r *= (1 + w*(m-1)) (penalty, ride_mult)
+    const float w = fmaxf(0.0f, fminf(1.0f, isv[RL_SURFER_SCAFFOLD_WEIGHT_INDEX]));
 
     const float entry_cost = isv[RL_ENTRY_COST_INDEX];
     const float min_hold   = isv[RL_SHORT_HOLD_MIN_STEPS_INDEX];
@@ -229,42 +236,36 @@ extern "C" __global__ void rl_fused_reward_pipeline(
 
     float r = reward;
 
-    if (pure_pnl_mode <= 0.5f) {
-        // 1. Entry cost: flat → positioned transition.
-        if (prev_lots == 0 && current_lots != 0) {
-            r -= entry_cost;
-        }
+    // 1. Entry cost: flat → positioned transition.
+    if (prev_lots == 0 && current_lots != 0) {
+        r -= w * entry_cost;
+    }
 
-        // 2. Short-hold penalty: trade close with hold time below minimum.
-        if (done > 0.5f && (float)hold_time < min_hold) {
-            r *= penalty;
-        }
+    // 2. Short-hold penalty: trade close with hold time below minimum.
+    if (done > 0.5f && (float)hold_time < min_hold) {
+        r *= (1.0f + w * (penalty - 1.0f));
+    }
 
-        // 3. Long-ride bonus: profitable close amplified by hold time.
-        if (done > 0.5f && r > 0.0f && hold_time > 0) {
-            const float ride_mult = 1.0f + hold_bonus * sqrtf((float)hold_time);
-            r *= ride_mult;
-        }
+    // 3. Long-ride bonus: profitable close amplified by hold time.
+    if (done > 0.5f && r > 0.0f && hold_time > 0) {
+        const float ride_mult = 1.0f + hold_bonus * sqrtf((float)hold_time);
+        r *= (1.0f + w * (ride_mult - 1.0f));
+    }
 
-        // 4. Per-step hold bonus for staying in a profitable position.
-        if (prev_lots != 0 && current_lots != 0 && r > 0.0f && hold_time > 0) {
-            r += hold_bonus * sqrtf((float)hold_time);
-        }
+    // 4. Per-step hold bonus for staying in a profitable position.
+    if (prev_lots != 0 && current_lots != 0 && r > 0.0f && hold_time > 0) {
+        r += w * hold_bonus * sqrtf((float)hold_time);
     }
 
     // ================================================================
     // PHASE 5b: Layer 3 inventory penalty (spec 2026-05-30-adaptive-risk-management).
-    // Apply BEFORE raw_rewards snapshot so the penalty is part of the
-    // shaped reward seen by Q/V learning. β = 0 (sentinel) → no-op.
-    // Also gated by pure_pnl_mode — when mode>0.5, Layers 1/2/4 own
-    // inventory risk via action overrides + Kelly sizing.
+    // Apply BEFORE raw_rewards snapshot. β = 0 (sentinel) → no-op.
+    // Also weighted by w: Layers 1/2/4 own inventory risk when w → 0.
     // ================================================================
-    if (pure_pnl_mode <= 0.5f) {
-        const float inv_beta = isv[RL_INVENTORY_PENALTY_BETA_INDEX];
-        if (inv_beta > 0.0f) {
-            const float net_pos_mag = (float)((current_lots < 0) ? -current_lots : current_lots);
-            r -= inv_beta * net_pos_mag;
-        }
+    const float inv_beta = isv[RL_INVENTORY_PENALTY_BETA_INDEX];
+    if (inv_beta > 0.0f) {
+        const float net_pos_mag = (float)((current_lots < 0) ? -current_lots : current_lots);
+        r -= w * inv_beta * net_pos_mag;
     }
 
     // Write shaped reward back.

crates/ml-alpha/cuda/rl_surfer_scaffold_controller.cu

@@ -0,0 +1,86 @@
+// rl_surfer_scaffold_controller.cu — adaptive Phase 5 shaping weight.
+//
+// Spec: docs/superpowers/specs/2026-06-01-reward-policy-alignment-investigation.md §6
+//
+// Reads (from ISV):
+//   RL_WIN_RATE_EMA_INDEX (677)            — agent's running win rate
+//   RL_CUMULATIVE_DONES_INDEX (660)        — total closed-trade count (resets at
+//                                            fold/eval boundary; controller bootstraps
+//                                            scaffold back to 1.0 from there)
+//   RL_EDGE_PH_FRAC_ALERTED_INDEX (751)    — fraction of batches whose Page-Hinkley
+//                                            edge-decay detector has alerted; used as
+//                                            re-engagement trigger when policy edge
+//                                            degrades
+//   RL_SURFER_BREAK_EVEN_WR_INDEX (754)    — config, default 0.30
+//   RL_SURFER_K_SHARPNESS_INDEX (755)      — config, default 30.0
+//   RL_SURFER_WARMUP_TRADES_INDEX (756)    — config, default 200.0
+//
+// Writes (to ISV):
+//   RL_SURFER_SCAFFOLD_WEIGHT_INDEX (753)  — w ∈ [0,1] consumed by
+//                                            rl_fused_reward_pipeline.cu Phase 5
+//
+// Math:
+//   competence = confidence(n_trades) × sigmoid(k_sharp × max(0, wr_ema − break_even))
+//   w_competence = 1 − competence
+//   w_decay      = min(1, 2 × frac_alerted)
+//   w            = clamp(max(w_competence, w_decay), 0, 1)
+//
+// At training start (n_trades=0, wr_ema=0): competence ≈ 0, w_competence ≈ 1 → w=1 (full scaffold)
+// After warmup and wr_ema crossing break-even: competence → 1, w → 0 (pure pnl)
+// If edge-decay PH alerts on >50% of batches: w_decay forces scaffold back up
+//
+// Per `feedback_no_atomicadd`: single-thread launch (1×1×1), no atomic ops.
+// Per `feedback_cpu_is_read_only`: pure device-side.
+// Per `feedback_no_nvrtc`: pre-compiled cubin.
+
+#include <stdint.h>
+#include <math_constants.h>
+
+// ISV slot indices (must match crates/ml-alpha/src/rl/isv_slots.rs)
+#define RL_WIN_RATE_EMA_INDEX               677
+#define RL_CUMULATIVE_DONES_INDEX           660
+#define RL_EDGE_PH_FRAC_ALERTED_INDEX       751
+#define RL_SURFER_SCAFFOLD_WEIGHT_INDEX     753
+#define RL_SURFER_BREAK_EVEN_WR_INDEX       754
+#define RL_SURFER_K_SHARPNESS_INDEX         755
+#define RL_SURFER_WARMUP_TRADES_INDEX       756
+
+__device__ static inline float sigmoidf(float x) {
+    return 1.0f / (1.0f + expf(-x));
+}
+
+extern "C" __global__ void rl_surfer_scaffold_controller(float* isv) {
+    // Single-thread kernel — launched (1,1,1)/(1,1,1).
+    if (threadIdx.x != 0 || blockIdx.x != 0) return;
+
+    const float wr_ema       = isv[RL_WIN_RATE_EMA_INDEX];
+    const float n_trades     = isv[RL_CUMULATIVE_DONES_INDEX];
+    const float frac_decay   = isv[RL_EDGE_PH_FRAC_ALERTED_INDEX];
+    const float break_even   = isv[RL_SURFER_BREAK_EVEN_WR_INDEX];
+    const float k_sharp      = isv[RL_SURFER_K_SHARPNESS_INDEX];
+    const float warmup_n     = isv[RL_SURFER_WARMUP_TRADES_INDEX];
+
+    // Confidence: sigmoid centered on warmup_n, scale = 30% of warmup.
+    // n_trades=0 → confidence ≈ sigmoid(-warmup/(0.3*warmup)) = sigmoid(-3.33) ≈ 0.034
+    // n_trades=warmup → confidence = sigmoid(0) = 0.5
+    // n_trades=2*warmup → confidence ≈ sigmoid(3.33) ≈ 0.966
+    const float warmup_scale = fmaxf(warmup_n * 0.3f, 1.0f);  // floor to avoid div-by-zero
+    const float confidence   = sigmoidf((n_trades - warmup_n) / warmup_scale);
+
+    // Win-rate excess above break-even, gated nonneg.
+    const float wr_excess = fmaxf(0.0f, wr_ema - break_even);
+
+    // Competence: agent must have BOTH trade-count confidence AND wr above break-even.
+    const float competence = confidence * sigmoidf(k_sharp * wr_excess);
+
+    // Inverse: scaffold weight from competence.
+    const float w_competence = 1.0f - competence;
+
+    // Edge-decay re-engagement: if >50% of batches are PH-alerted, force w back up.
+    const float w_decay = fminf(1.0f, 2.0f * frac_decay);
+
+    const float w_raw = fmaxf(w_competence, w_decay);
+    const float w     = fmaxf(0.0f, fminf(1.0f, w_raw));
+
+    isv[RL_SURFER_SCAFFOLD_WEIGHT_INDEX] = w;
+}

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

@@ -1737,24 +1737,46 @@ pub const RL_EDGE_PH_MEAN_INDEX:                    usize = 750;  // mean ph_sta
 pub const RL_EDGE_PH_FRAC_ALERTED_INDEX:            usize = 751;  // n_alerted / n_active
 pub const RL_EDGE_PH_FRAC_WARMUP_INDEX:             usize = 752;  // n_in_warmup / b_size
 
-// Reward-policy realignment (2026-06-01). When > 0.5, rl_fused_reward_pipeline
-// Phase 5 skips ALL hold-time shaping (steps 2-4) AND entry-cost (step 1).
-// Reward = realized_pnl_delta (Phase 1 output). Pure pnl alignment.
+// ─── Adaptive Surfer-Scaffold Shaping (v5, spec 2026-06-01 §6) ──────
 //
-// Diagnosis: alpha-rl-8gtk2 at 87a8259c6 showed Pearson(rewards.sum,
-// Δrealized_pnl_cum_usd) = 0.16 train / 0.31 eval — gradient signal
-// only weakly correlated with the trade-level pnl reality. Phase 5
-// step 4 adds hold_bonus × √hold_time = 2.0 × 10 = +20/step for
-// 100-step holds, dominating realized pnl by ~2600× per held trade.
+// EMPIRICAL FINDING (alpha-rl-8fb55, fa347e481, b=1024 fold-1, 2026-06-01):
+// the binary `pure_pnl_mode=1` killed the agent. Pearson(rewards, Δpnl) =
+// 0.92 was achieved (gradient direction correctly aligned with pnl), BUT
+// action_entropy stayed at 1.9-2.1 (random ≈ uniform log(11)=2.40) across
+// 4500+ steps. The policy never learned a strategy: pnl bled -$3.9M, wr
+// stuck at 0.30 (random baseline), positions oscillated 50% flat / 25%
+// long / 25% short without directional bias. Phase 5's hold-bonus was
+// not contamination alone — it was the INDUCTIVE BIAS SCAFFOLD that
+// taught the random-init policy what to optimize. Pure pnl alone, in
+// foxhunt's reward-density regime, is too weak to bootstrap learning.
 //
-// Bootstrap = 1.0 (NEW path default). Legacy mode=0 path is preserved
-// for ONE regression cluster smoke, then DELETED in a follow-up commit
-// per feedback_no_feature_flags + feedback_single_source_of_truth_no_duplicates.
-// This is a permanent reward redesign, not a permanent toggle.
+// REDESIGN — adaptive scaffold weight w ∈ [0,1] multiplies Phase 5
+// shaping (entry cost + short-hold penalty + long-ride bonus + per-step
+// hold bonus + Phase 5b inventory penalty). Driven by a controller that
+// fades the scaffold as the agent becomes competent (wr_ema rises above
+// break-even) but re-engages it when competence degrades (edge_decay PH
+// alert rate climbs). Replaces v4 binary toggle.
 //
-// See spec docs/superpowers/specs/2026-06-01-reward-policy-alignment-investigation.md
-// and pearl pearl_reward_signal_anti_aligned_with_pnl.
-pub const RL_REWARD_PURE_PNL_MODE_INDEX:            usize = 753;
+//   w = 1: identical to dd049d9a4 surfer-baseline shaping path
+//   w = 0: identical to v4 pure-pnl mode (mode=1)
+//   0<w<1: smooth interpolation (entry_cost × w, hold_bonus × w, etc.)
+//
+// Bootstrap w = 1.0 (full scaffold at random init), fade as the
+// `rl_surfer_scaffold_controller` proves the policy has crossed
+// break-even with statistical confidence.
+//
+// This SLOT REUSE (was RL_REWARD_PURE_PNL_MODE_INDEX in v4 commit
+// fa347e481) is acceptable per feedback_single_source_of_truth_no_duplicates:
+// v4 binary was a one-cluster-smoke artifact, never deployed. The
+// semantic flips: v4 1=pure-pnl, v5 1=full-shaping. Bootstrap value
+// changes from 1.0 (pure pnl) to 1.0 (full shaping). Conceptually it's
+// the SAME knob in continuous form with the meaning inverted.
+pub const RL_SURFER_SCAFFOLD_WEIGHT_INDEX:          usize = 753;
+
+// Config slots — controller hyperparameters (set by trainer bootstrap):
+pub const RL_SURFER_BREAK_EVEN_WR_INDEX:            usize = 754;  // default 0.30
+pub const RL_SURFER_K_SHARPNESS_INDEX:              usize = 755;  // default 30.0 (sigmoid slope around break-even)
+pub const RL_SURFER_WARMUP_TRADES_INDEX:            usize = 756;  // default 200.0 (cumulative dones before confidence ≈ 1)
 
 /// Last RL-allocated slot index (exclusive).
 /// Pre-risk-stack: 662. Post-Fix-B: 685. Post-v9 (warmup boundary): 696.
@@ -1767,5 +1789,6 @@ pub const RL_REWARD_PURE_PNL_MODE_INDEX:            usize = 753;
 /// Post-B-10 policy-quality cascade diagnostic: 743.
 /// Post-CMDP-fleet-fraction diagnostic: 747.
 /// Post-edge-decay-detector Phase 1: 753.
-/// Post-reward-policy-realignment: 754.
-pub const RL_SLOTS_END: usize = 754;
+/// Post-reward-policy-realignment v4 binary: 754 (slot reused in v5).
+/// Post-surfer-scaffold v5 adaptive (3 config slots): 757.
+pub const RL_SLOTS_END: usize = 757;

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

@@ -296,6 +296,8 @@ const RL_INVENTORY_BETA_CONTROLLER_CUBIN: &[u8] =
     include_bytes!(concat!(env!("OUT_DIR"), "/rl_inventory_beta_controller.cubin"));
 const RL_KELLY_FRACTION_CONTROLLER_CUBIN: &[u8] =
     include_bytes!(concat!(env!("OUT_DIR"), "/rl_kelly_fraction_controller.cubin"));
+const RL_SURFER_SCAFFOLD_CONTROLLER_CUBIN: &[u8] =
+    include_bytes!(concat!(env!("OUT_DIR"), "/rl_surfer_scaffold_controller.cubin"));
 const RL_EVAL_WARMUP_DECAY_CUBIN: &[u8] =
     include_bytes!(concat!(env!("OUT_DIR"), "/rl_eval_warmup_decay.cubin"));
 // Regime observer (F1.2 + F1.3) — flat-count block-reduce + state-machine
@@ -834,6 +836,8 @@ pub struct IntegratedTrainer {
     rl_inventory_beta_controller_fn: CudaFunction,
     _rl_kelly_fraction_controller_module: Arc<CudaModule>,
     rl_kelly_fraction_controller_fn: CudaFunction,
+    _rl_surfer_scaffold_controller_module: Arc<CudaModule>,
+    rl_surfer_scaffold_controller_fn: CudaFunction,
     // v9 (2026-05-31): defensive eval-boundary calibration kernel.
     _rl_eval_warmup_decay_module: Arc<CudaModule>,
     rl_eval_warmup_decay_fn: CudaFunction,
@@ -1800,6 +1804,12 @@ impl IntegratedTrainer {
         let rl_kelly_fraction_controller_fn = rl_kelly_fraction_controller_module
             .load_function("rl_kelly_fraction_controller")
             .context("load rl_kelly_fraction_controller")?;
+        let rl_surfer_scaffold_controller_module = ctx
+            .load_cubin(RL_SURFER_SCAFFOLD_CONTROLLER_CUBIN.to_vec())
+            .context("load rl_surfer_scaffold_controller cubin")?;
+        let rl_surfer_scaffold_controller_fn = rl_surfer_scaffold_controller_module
+            .load_function("rl_surfer_scaffold_controller")
+            .context("load rl_surfer_scaffold_controller")?;
         // v9 — defensive eval-boundary calibration kernel.
         let rl_eval_warmup_decay_module = ctx
             .load_cubin(RL_EVAL_WARMUP_DECAY_CUBIN.to_vec())
@@ -2946,6 +2956,8 @@ impl IntegratedTrainer {
             rl_inventory_beta_controller_fn,
             _rl_kelly_fraction_controller_module: rl_kelly_fraction_controller_module,
             rl_kelly_fraction_controller_fn,
+            _rl_surfer_scaffold_controller_module: rl_surfer_scaffold_controller_module,
+            rl_surfer_scaffold_controller_fn,
             _rl_eval_warmup_decay_module: rl_eval_warmup_decay_module,
             rl_eval_warmup_decay_fn,
             _rl_regime_flat_count_module: rl_regime_flat_count_module,
@@ -3479,7 +3491,7 @@ impl IntegratedTrainer {
         // (slot, value) pair — pure device write, no HtoD per
         // `feedback_no_htod_htoh_only_mapped_pinned`.
         {
-            let isv_constants: [(usize, f32); 241] = [
+            let isv_constants: [(usize, f32); 244] = [
                 // Static seeds for the adaptive reward-clamp controller —
                 // these are the initial values that
                 // `rl_reward_clamp_controller` will replace once it
@@ -3899,11 +3911,19 @@ impl IntegratedTrainer {
                 (crate::rl::isv_slots::RL_EDGE_PH_MEAN_INDEX,                     0.0_f32),
                 (crate::rl::isv_slots::RL_EDGE_PH_FRAC_ALERTED_INDEX,             0.0_f32),
                 (crate::rl::isv_slots::RL_EDGE_PH_FRAC_WARMUP_INDEX,              0.0_f32),
-                // Reward-policy realignment (spec 2026-06-01).
-                // 1.0 = pure pnl reward (Phase 5/5b shaping SKIPPED).
-                // 0.0 = legacy shaping path, preserved for ONE regression
-                // cluster smoke then deleted per feedback_no_feature_flags.
-                (crate::rl::isv_slots::RL_REWARD_PURE_PNL_MODE_INDEX,             1.0_f32),
+                // Adaptive surfer-scaffold shaping (spec v5 2026-06-01).
+                // w = 1.0: full Phase 5 shaping (surfer-baseline path, like dd049d9a4)
+                // w = 0.0: pure pnl reward (no shaping)
+                // 0<w<1:   smooth interpolation
+                // Bootstrap = 1.0 (random init needs the inductive bias scaffold);
+                // 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),
+                // 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),
             ];
             for (slot, value) in isv_constants.iter() {
                 let slot_i32 = *slot as i32;
@@ -4236,6 +4256,24 @@ impl IntegratedTrainer {
         Ok(())
     }
 
+    /// Adaptive surfer-scaffold weight controller (spec v5 2026-06-01).
+    /// Single-thread kernel; reads wr_ema/cumulative_dones/edge_decay_frac;
+    /// writes `RL_SURFER_SCAFFOLD_WEIGHT_INDEX`.
+    pub fn launch_rl_surfer_scaffold_controller(&self) -> Result<()> {
+        let mut args = RawArgs::new();
+        args.push_ptr(self.isv_dev_ptr);
+        let mut ptrs = args.build_arg_ptrs();
+        unsafe {
+            raw_launch(
+                self.rl_surfer_scaffold_controller_fn.cu_function(),
+                (1, 1, 1), (1, 1, 1), 0,
+                self.raw_stream,
+                &mut ptrs[..args.len()],
+            ).map_err(|e| anyhow::anyhow!("rl_surfer_scaffold_controller: {:?}", e))?;
+        }
+        Ok(())
+    }
+
     /// Adaptive risk management — Layer 4 Kelly fraction controller (spec 2026-05-30).
     /// Standalone launcher; production calls go through `rl_fused_controllers`.
     pub fn launch_rl_kelly_fraction_controller(&self) -> Result<()> {
@@ -7714,6 +7752,15 @@ impl IntegratedTrainer {
         self.launch_rl_eval_warmup_decay()
             .context("launch_rl_eval_warmup_decay")?;
 
+        // Adaptive surfer-scaffold weight controller (spec v5 §6, 2026-06-01).
+        // Reads wr_ema + cumulative_dones + edge_decay_frac_alerted; writes
+        // `RL_SURFER_SCAFFOLD_WEIGHT_INDEX` ∈ [0,1] consumed by
+        // `rl_fused_reward_pipeline.cu` Phase 5 NEXT step. MUST run after
+        // any controller that updates wr_ema/cumulative_dones so the
+        // weight reflects this-step state.
+        self.launch_rl_surfer_scaffold_controller()
+            .context("launch_rl_surfer_scaffold_controller")?;
+
         // PopArt: normalize rewards in-place (replaces apply_reward_scale).
         // Welford-EMA on batch mean+variance → whitened rewards → ISV
         // stats for V-correct. Grid=(1), Block=(min(B, 256)). Shared mem
@@ -10003,8 +10050,8 @@ impl IntegratedTrainer {
                     isv[RL_NEG_SCALED_REWARD_MAX_INDEX],
                 "neg_scaled_max_ema":
                     isv[RL_NEG_SCALED_REWARD_MAX_EMA_INDEX],
-                "pure_pnl_mode":
-                    isv[crate::rl::isv_slots::RL_REWARD_PURE_PNL_MODE_INDEX],
+                "surfer_scaffold_weight":
+                    isv[crate::rl::isv_slots::RL_SURFER_SCAFFOLD_WEIGHT_INDEX],
             },
             "ppo": {
                 "ratio_clamp_max":   isv[RL_PPO_RATIO_CLAMP_MAX_INDEX],

crates/ml-alpha/tests/reward_alignment_invariants.rs

@@ -1,24 +1,21 @@
-//! Reward-policy realignment invariants (spec 2026-06-01).
+//! Reward-policy realignment invariants (spec v5 2026-06-01 §6).
 //!
-//! Validates that the new `RL_REWARD_PURE_PNL_MODE` ISV slot propagates
-//! correctly from trainer bootstrap → ISV → kernel `rl_fused_reward_pipeline`
-//! → diag.jsonl, and that the mode=1 default holds across train and eval
-//! phases.
+//! Validates that the `RL_SURFER_SCAFFOLD_WEIGHT` ISV slot propagates
+//! correctly from trainer bootstrap → ISV → kernel
+//! `rl_surfer_scaffold_controller` → `rl_fused_reward_pipeline` →
+//! diag.jsonl, and that the adaptive controller stays in bounds.
 //!
 //! Three invariants:
-//!   I1. `rewards.pure_pnl_mode == 1.0` on every train row (bootstrap
-//!       discipline; mode is a constant for the whole run).
-//!   I2. `rewards.pure_pnl_mode == 1.0` on every eval row (no slot rewrite
-//!       at reset_session_state — see [[pearl_adaptive_carryover_discipline]]
-//!       for the orthogonal "EMA boundary reset" rule; const-bootstrap slots
-//!       are not subject to it).
-//!   I3. `rewards.abs_max <= REWARD_INFLATION_FENCE` ($25k) on every row.
-//!       This excludes the Phase 5 step-4 per-step hold bonus +20×realized
-//!       inflation pattern that produced the alpha-rl-8gtk2 misalignment
-//!       (hold_bonus=2.0 × √100 = +20/step on hold_time=100 trades, ×~$1k
-//!       realized scale = effective $20k+ per-step bonus). With Phase 5
-//!       skipped, magnitudes track raw pnl deltas which stay well under the
-//!       fence at the test's b_size=16 + 200 steps.
+//!   I1. `rewards.surfer_scaffold_weight ∈ [0, 1]` on every train row
+//!       (controller output is clamped — kernel writes in [0,1]).
+//!   I2. `rewards.surfer_scaffold_weight ∈ [0, 1]` on every eval row
+//!       (same controller runs in eval; per-step launches preserved).
+//!   I3. At step 0, scaffold_weight is near 1.0 (≥ 0.9). The controller
+//!       fires before the per-step diag snapshot, so the bootstrap value
+//!       is not literally observable; what IS observable is that a novice
+//!       agent (n_trades ≈ 0, wr_ema ≈ 0) drives the controller toward
+//!       w_competence ≈ 1.0. If this falls, controller math broke OR the
+//!       trainer bootstrap didn't propagate (config slots wrong).
 //!
 //! Per `feedback_no_cpu_test_fallbacks.md`: GPU-oracle end-to-end (drives
 //! the alpha_rl_train release binary; parses the diag JSONL it emits).
@@ -75,7 +72,7 @@ fn dot_get(v: &Value, path: &str) -> Result<f64> {
 
 #[test]
 #[ignore = "requires CUDA + pre-built release binary + MBP-10 test data"]
-fn reward_alignment_pure_pnl_mode_invariants() -> Result<()> {
+fn reward_alignment_surfer_scaffold_invariants() -> Result<()> {
     let bin = binary_path();
     anyhow::ensure!(
         bin.exists(),
@@ -134,40 +131,35 @@ fn reward_alignment_pure_pnl_mode_invariants() -> Result<()> {
         "eval rows {} != n_eval_steps {n_eval_steps}", eval_rows.len()
     );
 
-    // I1 + I2: pure_pnl_mode must read exactly 1.0 on every row of every phase.
-    // I3: reward.abs_max must stay below the hold-bonus inflation fence.
-    //
-    // Fence rationale: hold_bonus=2.0 × √100=10 = +20/step additive bonus per
-    // held trade at b=16 produces rewards.abs_max in the 10^3+ range. With
-    // Phase 5 skipped, abs_max reflects realized pnl deltas alone which at
-    // b=16 + 200 steps + LobSim retail-fee structure rarely exceed $5k/step.
-    // $25k is a 5× safety margin that any return of Phase 5 step 4 would
-    // blow past.
-    const REWARD_INFLATION_FENCE: f64 = 25_000.0;
-
+    // I1 + I2: scaffold_weight clamped to [0,1] (kernel-enforced).
+    // I3: bootstrap value 1.0 visible in train[0] (controller hasn't fired yet).
     let mut checked = 0usize;
+    let train_t0 = dot_get(&train_rows[0], "rewards.surfer_scaffold_weight")?;
+    anyhow::ensure!(
+        train_t0 >= 0.9 && train_t0 <= 1.0 + 1e-6,
+        "I3 violated: train[0] rewards.surfer_scaffold_weight = {train_t0:.4} \
+         (expected ≥ 0.9 at step 0 with novice agent; the controller's first \
+         fire on n_trades=0 wr_ema≈0 should yield w_competence ≈ 1.0 minus a \
+         tiny confidence-bootstrap residual. If train_t0 < 0.9, the controller \
+         math is off or config-slot bootstrap is missing)"
+    );
+    checked += 1;
+
     for (phase, rows) in &[("train", &train_rows), ("eval", &eval_rows)] {
         for (i, row) in rows.iter().enumerate() {
-            let mode = dot_get(row, "rewards.pure_pnl_mode")?;
+            let w = dot_get(row, "rewards.surfer_scaffold_weight")?;
             anyhow::ensure!(
-                (mode - 1.0).abs() < 1e-6,
-                "{phase}[{i}] rewards.pure_pnl_mode = {mode} (expected 1.0; \
-                 bootstrap or kernel branch broken)"
-            );
-            let abs_max = dot_get(row, "rewards.abs_max")?;
-            anyhow::ensure!(
-                abs_max <= REWARD_INFLATION_FENCE,
-                "{phase}[{i}] rewards.abs_max = {abs_max:.2} exceeds fence \
-                 ${REWARD_INFLATION_FENCE:.0} — Phase 5 hold-bonus inflation \
-                 may have returned; check rl_fused_reward_pipeline.cu Phase 5 gate"
+                w >= -1e-6 && w <= 1.0 + 1e-6,
+                "{phase}[{i}] rewards.surfer_scaffold_weight = {w} out of [0,1] \
+                 — controller clamp broken or slot corrupted"
             );
             checked += 1;
         }
     }
 
     eprintln!(
-        "reward_alignment_pure_pnl_mode_invariants OK: {} rows validated \
-         (train[0..{}] + eval[0..{}])",
+        "reward_alignment_surfer_scaffold_invariants OK: {} rows validated \
+         (train[0..{}] + eval[0..{}]); train[0] bootstrap = {train_t0:.3}",
         checked, n_steps, n_eval_steps
     );
     Ok(())