feat(rl): Layer 3 — per-step drawdown penalty for exit gradient

Layer 3 of the four-layer defense per pearl_wr_above_50_with_negative_pnl_loss_cutting.

The wr-targeting controller (Layer 4) successfully drives wr=0.55+ but
PnL spirals negative because Q can only learn exit timing from sparse
close-event rewards. Without per-step exit signal during open trades,
Q can't learn "close losing trade early".

Fix: add continuous drawdown penalty during open trades.
- rl_fused_reward_pipeline.cu PHASE 5 shaping:
  if (current_lots != 0 && unrealized < 0)
      r += unrealized * rate * reward_scale
- Penalty-only (no symmetric reward for unrealized gain) — avoids
  exposure-positive bias per pearl_event_driven_reward_density_alignment
- Reads unrealized_pnl from PosFlat offset 24 (added this session)

ISV slot 592 RL_DRAWDOWN_PENALTY_RATE_INDEX, bootstrap 0.001.
Conservative rate to avoid "Q learns trading is punished" pathology
(observed when rate was 0.01 in earlier session).

Local smoke 1000 steps (b=128):
- wr trajectory: 0.31 → 0.46 (Layer 4 controller still working)
- pnl: -$105k (vs -$5M without layer 3 at cluster scale)
- Completes clean, no NaN

The per-step penalty magnitude (~1e-5 to 1e-3 scaled units at rate=0.001)
provides Q with the exit gradient it was missing. Cluster validation
will show if surfer + positive PnL emerges at 5k+ steps.

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

crates/ml-alpha/cuda/rl_fused_reward_pipeline.cu

@@ -39,6 +39,7 @@
 #define RL_SHORT_HOLD_PENALTY_INDEX        534
 #define RL_HOLD_BONUS_INDEX                535
 #define RL_OUTCOME_ALPHA_INDEX             520
+#define RL_DRAWDOWN_PENALTY_RATE_INDEX     592
 
 #define MAX_UNITS 4
 
@@ -243,6 +244,31 @@ extern "C" __global__ void rl_fused_reward_pipeline(
         r += hold_bonus * sqrtf((float)hold_time);
     }
 
+    // 5. Layer 3 of the four-layer defense: per-step drawdown penalty.
+    //    Creates a CONTINUOUS exit gradient during open trades. Without
+    //    this, model learns entries but not exits — the close-event
+    //    reward is too sparse for Q to learn early loss-cutting (see
+    //    `pearl_wr_above_50_with_negative_pnl_loss_cutting`).
+    //
+    //    Reads unrealized_pnl from PosFlat (offset 24). Penalty fires
+    //    ONLY when position is open AND unrealized < 0. Penalty-only
+    //    (no symmetric reward for unrealized gain) avoids exposure-
+    //    positive bias per
+    //    `pearl_event_driven_reward_density_alignment`.
+    //
+    //    The penalty is applied in the SAME scaled reward space as the
+    //    other rewards (multiplied by reward_scale to match downstream
+    //    `apply_reward_scale` semantics).
+    {
+        const float drawdown_penalty_rate = isv[RL_DRAWDOWN_PENALTY_RATE_INDEX];
+        if (drawdown_penalty_rate > 0.0f && current_lots != 0) {
+            const float unrealized = *(const float*)(p + 24);
+            if (unrealized < 0.0f) {
+                r += unrealized * drawdown_penalty_rate * reward_scale;
+            }
+        }
+    }
+
     // Write shaped reward back.
     rewards[b] = r;
 

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

@@ -1175,5 +1175,24 @@ pub const RL_WR_EMA_INDEX: usize = 590;
 /// `rl_isv_write` at trainer init.
 pub const RL_WR_TARGET_INDEX: usize = 591;
 
+/// Per-step drawdown penalty rate — Layer 3 of the four-layer defense
+/// (wr-targeting LOSS controller is Layer 4). Provides Q with a
+/// CONTINUOUS exit gradient during open trades; without it, the close-
+/// event reward is sparse and Q can't learn to close losing trades
+/// early (canonical incident: wr climbed 0.34→0.45 while pnl
+/// deteriorated -$1.97M→-$5.02M because entries improved but exits
+/// did not — see `pearl_wr_above_50_with_negative_pnl_loss_cutting`).
+///
+/// The reward pipeline (rl_fused_reward_pipeline.cu) reads this slot
+/// and adds `unrealized × rate × reward_scale` to the per-step reward
+/// when (a) position is open and (b) unrealized PnL is negative.
+/// Penalty-only (no symmetric reward for unrealized gain) avoids
+/// exposure-positive bias per
+/// `pearl_event_driven_reward_density_alignment`.
+///
+/// Default: 0.001 — CONSERVATIVE. Higher rates (>0.01) cause Q to
+/// learn "any trade is bad" → never-trade pathology.
+pub const RL_DRAWDOWN_PENALTY_RATE_INDEX: usize = 592;
+
 /// Last RL-allocated slot index (exclusive).
-pub const RL_SLOTS_END: usize = 592;
+pub const RL_SLOTS_END: usize = 593;

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

@@ -2825,7 +2825,7 @@ impl IntegratedTrainer {
         // (slot, value) pair — pure device write, no HtoD per
         // `feedback_no_htod_htoh_only_mapped_pinned`.
         {
-            let isv_constants: [(usize, f32); 114] = [
+            let isv_constants: [(usize, f32); 115] = [
                 // Static seeds for the adaptive reward-clamp controller —
                 // these are the initial values that
                 // `rl_reward_clamp_controller` will replace once it
@@ -3009,6 +3009,14 @@ impl IntegratedTrainer {
                 // pick a uniform random action, preventing entropy collapse
                 // per `pearl_pi_actor_collapses_without_entropy_floor`.
                 (crate::rl::isv_slots::RL_THOMPSON_FLOOR_INDEX,              0.05),
+                // Layer 3 of the four-layer defense — per-step drawdown
+                // penalty rate. Provides Q a continuous exit gradient
+                // during open trades (close-event reward alone is too
+                // sparse for Q to learn to cut losers early). Conservative
+                // 0.001 — higher rates cause never-trade pathology. Read
+                // by `rl_fused_reward_pipeline` when position is open
+                // and unrealized PnL is negative.
+                (crate::rl::isv_slots::RL_DRAWDOWN_PENALTY_RATE_INDEX,     0.001),
             ];
             for (slot, value) in isv_constants.iter() {
                 let slot_i32 = *slot as i32;