fix(rl): replace-directly on first warm observation (4 controllers)

R9 cluster smoke alpha-rl-qzstj step 7 caught the second half of
the cold-start fix: even with the input==0 gate holding controllers
at bootstrap until the first observation, the Wiener α-floor=0.4
blend then produced `0.6 × bootstrap + 0.4 × target` on the FIRST
warm step — 60% bootstrap contamination distorting the controller's
first emit.

For `rl_reward_scale` this was the load-bearing failure: with
prev=1.0 (bootstrap) and target=1/832=0.0012 on the first closed
trade, blend gave scale=0.600 → real $832 × 0.6 = $499 fed to V
regression → l_v = 249,782. Replace-directly: scale = 0.0012
immediately → V target = 1.0 → l_v ≈ 1. Three orders of magnitude
reduction in cold-start contamination.

## Fix

For each of the 4 cold-start-gated controllers (τ, ε, n_roll, scale),
detect "first warm observation" via `prev == BOOTSTRAP_VALUE` and
write target directly instead of Wiener blending. Subsequent steps
(where prev has drifted via earlier blends) take the Wiener path
unchanged.

```cuda
// (cold-start gate, then target computation already done)
if (prev == HARDCODED_BOOTSTRAP_VALUE) {
    isv[OUTPUT_INDEX] = target;
    return;
}
// ... Wiener blend
```

The `prev == HARDCODED_BOOTSTRAP` check uses float equality but is
safe: the sentinel-bootstrap path WROTE that exact value, and the
cold-start gate prevents any arithmetic from touching it until input
becomes non-zero. The first non-zero input triggers this branch
exactly once.

This is `pearl_first_observation_bootstrap` ("sentinel = 0; first
observation replaces directly") applied at the controller's bootstrap
→ warm transition. The pearl was originally framed for EMA producers;
the R9 audit shows it applies equally to adaptive controllers whose
hardcoded bootstrap doubles as a "no data yet" sentinel.

## Test impact

`g3_per_step_controllers_move_isv_outputs_when_fed_real_emas` now
shows stronger first-observation moves (replace-directly hits target
cleanly):

  Before R9 fixes:   τ 0.005 → 0.023   ε 0.2 → 0.14   scale 1 → 0.608
  After cold-gate:   τ 0.005 → 0.023   ε 0.2 → 0.14   scale 1 → 0.608
  After this fix:    τ 0.005 → 0.05    ε 0.2 → 0.05   scale 1 → 0.02

All gates still green:
  G1  isv_bootstrap            ✅
  G3  controllers_emit         ✅ (stronger first-emit movements)
  G4  target_soft_update       ✅
  G6  r7d_per_wiring           ✅
  R3, R4, smoke                ✅

## What's NOT in this commit

The 6 missing EMA input wirings (kl_pi, q_divergence,
entropy_observed, advantage_var_ratio, td_kurtosis, trade_duration)
remain. Six of seven controllers will still hold at bootstrap during
the cluster smoke because their input EMAs receive no signal. That
fix is the next commit — it requires new reduce kernels (var-over-
abs-mean, kurtosis, KL-approx, L2-diff-norm) and a per-batch
trade-duration counter.

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

crates/ml-alpha/cuda/rl_ppo_clip_controller.cu

@@ -83,6 +83,16 @@ extern "C" __global__ void rl_ppo_clip_controller(
     float       eps_target = eps_prev * ratio;
     eps_target = fmaxf(EPS_MIN, fminf(eps_target, EPS_MAX));
 
+    // First-observation replace-directly per
+    // `pearl_first_observation_bootstrap`. See rl_target_tau_controller
+    // for the rationale — the Wiener blend's bootstrap contamination
+    // is what produced the step-7 l_v=249,782 spike in alpha-rl-qzstj
+    // (R9 audit 2026-05-23).
+    if (eps_prev == EPS_BOOTSTRAP) {
+        isv[RL_PPO_CLIP_INDEX] = eps_target;
+        return;
+    }
+
     // Wiener-α blend with floor per pearl_wiener_alpha_floor_for_nonstationary.
     const float a       = fmaxf(alpha, WIENER_ALPHA_FLOOR);
     float       eps_new = (1.0f - a) * eps_prev + a * eps_target;

crates/ml-alpha/cuda/rl_reward_scale_controller.cu

@@ -107,6 +107,19 @@ extern "C" __global__ void rl_reward_scale_controller(
     float       target = 1.0f / denom;
     target = fmaxf(REWARD_SCALE_MIN, fminf(target, REWARD_SCALE_MAX));
 
+    // First-observation replace-directly per
+    // `pearl_first_observation_bootstrap`. This is the load-bearing
+    // fix for the step-7 contamination cascade in alpha-rl-qzstj
+    // (R9 audit 2026-05-23): with prev=1.0 (bootstrap) and target=
+    // 1/832=0.0012, the Wiener blend at α=0.4 produced scale=0.600
+    // — applying 0.6 × $832 = $499 to V regression → MSE = 249,000.
+    // Replace-directly: scale = 0.0012 immediately → V target = 1.0,
+    // MSE = 1.0. Three orders of magnitude reduction.
+    if (prev == REWARD_SCALE_BOOTSTRAP) {
+        isv[RL_REWARD_SCALE_INDEX] = target;
+        return;
+    }
+
     // Wiener-α blend with floor per pearl_wiener_alpha_floor_for_nonstationary.
     const float a   = fmaxf(alpha_step, WIENER_ALPHA_FLOOR);
     float       out = (1.0f - a) * prev + a * target;

crates/ml-alpha/cuda/rl_rollout_steps_controller.cu

@@ -92,6 +92,13 @@ extern "C" __global__ void rl_rollout_steps_controller(
     float       target    = prev * scale;
     target = fmaxf(ROLLOUT_MIN, fminf(target, ROLLOUT_MAX));
 
+    // First-observation replace-directly per
+    // `pearl_first_observation_bootstrap`.
+    if (prev == ROLLOUT_BOOTSTRAP) {
+        isv[RL_N_ROLLOUT_STEPS_INDEX] = target;
+        return;
+    }
+
     // Wiener-α blend with floor per pearl_wiener_alpha_floor_for_nonstationary.
     const float a   = fmaxf(alpha, WIENER_ALPHA_FLOOR);
     float       out = (1.0f - a) * prev + a * target;

crates/ml-alpha/cuda/rl_target_tau_controller.cu

@@ -85,6 +85,20 @@ extern "C" __global__ void rl_target_tau_controller(
     float       tau_target = tau_prev * ratio;
     tau_target = fmaxf(TAU_MIN, fminf(tau_target, TAU_MAX));
 
+    // First-observation replace-directly (R9 audit 2026-05-23): if
+    // prev is still exactly the hardcoded bootstrap value, this is
+    // the FIRST per-step fire that has real input signal. Per
+    // `pearl_first_observation_bootstrap`, "first observation
+    // replaces directly" — the Wiener blend `(1-α)·bootstrap +
+    // α·target` would leave 60% bootstrap contamination, distorting
+    // the controller's first emit. Write target directly instead.
+    // Subsequent steps (where prev has drifted off bootstrap via
+    // earlier blends) take the Wiener path below.
+    if (tau_prev == TAU_BOOTSTRAP) {
+        isv[RL_TARGET_TAU_INDEX] = tau_target;
+        return;
+    }
+
     // Wiener-α blend with floor per pearl_wiener_alpha_floor_for_nonstationary.
     const float a       = fmaxf(alpha, WIENER_ALPHA_FLOOR);
     float       tau_new = (1.0f - a) * tau_prev + a * tau_target;