diff --git a/crates/ml-alpha/cuda/rl_confidence_gate.cu b/crates/ml-alpha/cuda/rl_confidence_gate.cu
index 158b17593..c97b36c17 100644
--- a/crates/ml-alpha/cuda/rl_confidence_gate.cu
+++ b/crates/ml-alpha/cuda/rl_confidence_gate.cu
@@ -34,6 +34,8 @@
 #define RL_HOLD_TARGET_FRAC_INDEX      575
 #define RL_HOLD_FRAC_EMA_INDEX         576
 #define RL_CONF_GATE_MAX_HOLD_FRAC_INDEX 584
+#define RL_STOP_LOSS_THRESHOLD_INDEX    587
+#define RL_MEAN_ABS_PNL_EMA_INDEX      423
 #define THRESHOLD_MIN                  0.05f
 #define THRESHOLD_MAX                  0.95f
 #define HOLD_EMA_ALPHA                 0.1f
@@ -51,6 +53,35 @@ extern "C" __global__ void rl_confidence_gate(
     const int b = blockIdx.x;
     if (b >= b_size) return;
 
+    // ── Hard stop-loss: override action to Flat when unrealized loss
+    // exceeds ISV-driven threshold. Fires BEFORE gate logic so the
+    // lobsim actually closes the position (no state mismatch).
+    // unrealized_loss = |vwap - mid| × |lots|; normalized by mean_abs_pnl.
+    {
+        const unsigned char* p = pos_state + b * pos_bytes;
+        const int lots = *(const int*)(p + 0);
+        if (lots != 0) {
+            const float stop_thresh = isv[RL_STOP_LOSS_THRESHOLD_INDEX];
+            const float mean_abs = fmaxf(isv[RL_MEAN_ABS_PNL_EMA_INDEX], 1e-6f);
+            // Approximate mid from the first bid/ask level would need
+            // book data; instead use the realized_pnl delta as a proxy
+            // for position health. Simpler: use vwap vs realized_pnl trend.
+            // For now, use the unit_entry_price-based unrealized_r from
+            // trade_context — but that runs AFTER this kernel.
+            //
+            // Practical approach: read peak_equity (offset 12) and
+            // realized_pnl (offset 8). Drawdown from peak = peak - realized.
+            const float realized = *(const float*)(p + 8);
+            const float peak = *(const float*)(p + 12);
+            const float drawdown = peak - realized;
+            const float normalized_dd = drawdown / mean_abs;
+            if (normalized_dd > stop_thresh) {
+                actions[b] = (lots > 0) ? 3 : 4;  // FlatFromLong / FlatFromShort
+                return;
+            }
+        }
+    }
+
     const int current_step = (int)isv[RL_STEP_COUNTER_ISV_INDEX];
     const int warmup = (int)isv[RL_GATE_WARMUP_STEPS_INDEX];
     if (current_step < warmup) return;