diff --git a/crates/ml-alpha/cuda/rl_confidence_gate.cu b/crates/ml-alpha/cuda/rl_confidence_gate.cu
index 645d7d32c..1c6603e56 100644
--- a/crates/ml-alpha/cuda/rl_confidence_gate.cu
+++ b/crates/ml-alpha/cuda/rl_confidence_gate.cu
@@ -34,6 +34,7 @@
 #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_MARGIN_FLOOR_USD_INDEX      587
 #define THRESHOLD_MIN                  0.05f
 #define THRESHOLD_MAX                  0.95f
 #define HOLD_EMA_ALPHA                 0.1f
@@ -51,6 +52,26 @@ extern "C" __global__ void rl_confidence_gate(
     const int b = blockIdx.x;
     if (b >= b_size) return;
 
+    // ── Margin floor: force Hold when cumulative loss exceeds limit.
+    // Reads realized_pnl from PosFlat (offset 8). If the sim has lost
+    // more than the margin floor, ALL non-Hold actions become Hold.
+    // This is a HARD constraint — broker will liquidate at margin call.
+    // Fires before warmup check, exploration slots, everything.
+    {
+        const unsigned char* p = pos_state + b * pos_bytes;
+        const float realized = *(const float*)(p + 8);
+        const float margin_floor = isv[RL_MARGIN_FLOOR_USD_INDEX];
+        if (realized < -margin_floor) {
+            const int lots = *(const int*)(p + 0);
+            if (lots != 0) {
+                actions[b] = (lots > 0) ? 3 : 4;  // force flat
+            } else {
+                actions[b] = ACTION_HOLD;  // stay flat, don't re-enter
+            }
+            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;
diff --git a/crates/ml-alpha/src/rl/isv_slots.rs b/crates/ml-alpha/src/rl/isv_slots.rs
index 1c52b7b6f..f026bd442 100644
--- a/crates/ml-alpha/src/rl/isv_slots.rs
+++ b/crates/ml-alpha/src/rl/isv_slots.rs
@@ -1127,5 +1127,11 @@ pub const RL_DONE_WIN_MAGNITUDE_EMA_INDEX: usize = 585;
 /// per-step tail maxes.
 pub const RL_DONE_LOSS_MAGNITUDE_EMA_INDEX: usize = 586;
 
+/// Maximum allowed drawdown in USD before force-Hold (margin floor).
+/// $35k capital, ES maintenance margin ~$14.4k → max drawdown ~$20k.
+/// The confidence gate reads realized_pnl from PosFlat — if loss
+/// exceeds this, ALL non-Hold actions are overridden to Hold.
+pub const RL_MARGIN_FLOOR_USD_INDEX: usize = 587;
+
 /// Last RL-allocated slot index (exclusive).
-pub const RL_SLOTS_END: usize = 587;
+pub const RL_SLOTS_END: usize = 588;
diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs
index 23523f4ec..2dede7d93 100644
--- a/crates/ml-alpha/src/trainer/integrated.rs
+++ b/crates/ml-alpha/src/trainer/integrated.rs
@@ -2673,7 +2673,7 @@ impl IntegratedTrainer {
         // (slot, value) pair — pure device write, no HtoD per
         // `feedback_no_htod_htoh_only_mapped_pinned`.
         {
-            let isv_constants: [(usize, f32); 111] = [
+            let isv_constants: [(usize, f32); 112] = [
                 // Static seeds for the adaptive reward-clamp controller —
                 // these are the initial values that
                 // `rl_reward_clamp_controller` will replace once it
@@ -2840,6 +2840,9 @@ impl IntegratedTrainer {
                 (crate::rl::isv_slots::RL_SAC_ALPHA_INDEX,                  0.01),
                 (crate::rl::isv_slots::RL_SAC_ENTROPY_TARGET_INDEX,         1.68),
                 (crate::rl::isv_slots::RL_CONF_GATE_MAX_HOLD_FRAC_INDEX,   0.85),
+                // Hard margin floor: $35k capital, ~$14.4k ES maintenance.
+                // Max drawdown ~$20k before broker liquidates.
+                (crate::rl::isv_slots::RL_MARGIN_FLOOR_USD_INDEX,         20000.0),
             ];
             for (slot, value) in isv_constants.iter() {
                 let slot_i32 = *slot as i32;