diff --git a/crates/ml-alpha/cuda/rl_fused_reward_pipeline.cu b/crates/ml-alpha/cuda/rl_fused_reward_pipeline.cu
index 2ef0dfb1a..ff2b5e233 100644
--- a/crates/ml-alpha/cuda/rl_fused_reward_pipeline.cu
+++ b/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;
 
diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs
index 8cfbe100d..b43507a01 100644
--- a/crates/ml-alpha/src/rl/isv_slots.rs
+++ b/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;
diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs
index 9e0fc739a..bad1c4d77 100644
--- a/crates/ml-alpha/src/trainer/integrated.rs
+++ b/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;