feat(rl): adaptive C51 atom span ratchet to lift Q learning ceiling

rdgzl follow-up — chain hypothesis layer 2:
  reward clamp lift unlocked V regression + PPO advantage (R/done
  -$1.39 → -$0.48), but Q's distributional learning was structurally
  capped at hardcoded V_MAX=1.0 in bellman_target_projection.cu —
  any Bellman target > 1.0 categorically projected to atom 20 (top)
  regardless of clamp. Even with WIN=3.8 clamp, Q never saw a +3.8
  reward signal as distinct from a +1.0 reward signal.

This commit makes V_MIN/V_MAX ISV-driven with monotone-grow ratchet
coupled to the reward clamp. The C51 distribution support adapts
WITHOUT destabilising Q's learned values — atom 20 always represents
at least the widest WIN we've ever admitted (only grows, never shrinks).

Implementation:
  - 2 new ISV slots (484 V_MAX, 485 V_MIN) with [-1, +1] floors
    seeded by rl_isv_write
  - rl_reward_clamp_controller.cu also ratchets these slots:
    V_MAX_new = max(V_MAX_prev, max(1.0, WIN_clamp))
    V_MIN_new = min(V_MIN_prev, min(-1.0, -LOSS_clamp))
  - bellman_target_projection.cu reads V_MIN/V_MAX from ISV, derives
    DELTA_Z inline (was #define)
  - New rl_atom_support_update.cu (21-thread block) refreshes
    atom_supports_d = linspace(V_MIN, V_MAX, 21) per step so
    downstream C51 kernels (argmax_expected_q, rl_action_kernel,
    dqn_distributional_q) see the current span
  - Trainer launches atom-support updater after each reward-clamp
    controller launch (both helper + step_with_lobsim inline paths)
  - Diag exposes c51_v_max + c51_v_min for adaptation visibility

Floors at [-1, +1] preserve original C51 design as hard minimum —
the atom support can only become wider, never narrower than the
baseline.

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

crates/ml-alpha/cuda/rl_reward_clamp_controller.cu

@@ -68,6 +68,13 @@
 // MARGIN adaptation slots (audit 2026-05-24 follow-up).
 #define RL_REWARD_CLAMP_CLIP_RATE_EMA_INDEX       482
 #define RL_REWARD_CLAMP_CLIP_RATE_TARGET_INDEX    483
+// C51 atom span ratchet slots (audit 2026-05-24 second follow-up).
+// Coupling: V_MAX tracks max(V_MAX_prev, max(1.0, WIN_clamp));
+// V_MIN tracks min(V_MIN_prev, min(-1.0, -LOSS_clamp)). Ratchet
+// (monotone-grow) prevents destabilising Q's learned distribution
+// when WIN_clamp temporarily shrinks.
+#define RL_C51_V_MAX_INDEX                        484
+#define RL_C51_V_MIN_INDEX                        485
 
 #define MIN_WIN                  1.0f
 #define WIENER_ALPHA_FLOOR       0.4f
@@ -172,4 +179,31 @@ extern "C" __global__ void rl_reward_clamp_controller(
 
     isv[RL_REWARD_CLAMP_WIN_INDEX]  = win_eff;
     isv[RL_REWARD_CLAMP_LOSS_INDEX] = loss_eff;
+
+    // ── C51 atom span ratchet (audit follow-up). ──────────────────
+    //
+    // V_MAX/V_MIN track the reward clamp with monotone-grow semantics
+    // — once we've admitted a winning trade of magnitude X, the C51
+    // atom mapping permanently includes X in its support. This lets
+    // bellman_target_projection's clip pass through wins > 1.0 to
+    // distinct atoms instead of saturating the top atom.
+    //
+    // Floors: V_MAX ≥ 1.0, V_MIN ≤ -1.0 preserve the original C51
+    // [-1, +1] support as a hard minimum (Q's distribution can't
+    // become coarser than the design baseline).
+    const float v_max_prev = isv[RL_C51_V_MAX_INDEX];
+    const float v_min_prev = isv[RL_C51_V_MIN_INDEX];
+    const float v_max_floor = 1.0f;
+    const float v_min_floor = -1.0f;
+    // Ratchet: target = max(floor, win); kept = max(prev, target).
+    const float v_max_target = fmaxf(v_max_floor, win_eff);
+    const float v_min_target = fminf(v_min_floor, -loss_eff);
+    // Bootstrap on sentinel 0 — first emit replaces directly with
+    // target (per pearl_first_observation_bootstrap).
+    const float v_max_new = (v_max_prev == 0.0f) ? v_max_target
+                                                 : fmaxf(v_max_prev, v_max_target);
+    const float v_min_new = (v_min_prev == 0.0f) ? v_min_target
+                                                 : fminf(v_min_prev, v_min_target);
+    isv[RL_C51_V_MAX_INDEX] = v_max_new;
+    isv[RL_C51_V_MIN_INDEX] = v_min_new;
 }