feat(rl): four-layer loss defense + perf sync removal

Loss minimization (wr=0.561 but PnL negative — losses 32% > wins):

A1: Exempt FlatL/FlatS from confidence gate — exit actions never
    blocked, model can always close losing positions.

A2+A3: New rl_drawdown_stop kernel — per-step drawdown penalty
    (min(0, unrealized_r) × rate) creates continuous exit gradient.
    Hard stop-loss force-closes when unrealized_r < -threshold.
    Both ISV-driven (slots 586, 587).

A4: Adaptive LOSS clamp — tracks observed neg/pos EMA ratio instead
    of static 3.0. LOSS = clamp(1.0, ratio×1.1, 3.0). Q sees
    accurate loss magnitudes.

Performance:

B0: Remove gratuitous stream.synchronize() in apply_snapshot
    (sim/mod.rs) — same-stream ordering makes it unnecessary.
    Expected: -12-49ms/step.

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

Cargo.lock

@@ -6090,6 +6090,7 @@ dependencies = [
  "approx",
  "arrow 56.2.0",
  "bincode",
+ "bytemuck",
  "clap",
  "cudarc",
  "data",

crates/ml-alpha/build.rs

@@ -73,6 +73,7 @@ const KERNELS: &[&str] = &[
     "rl_kl_reference_grad",
     "rl_q_pi_distill_grad",          // audit 2026-05-24 vj5f6 followup: KL(softmax(E_Q/τ) || π_new) gradient ADDED to pi_grad_logits — couples Q's improved C51 calibration to action selection (was decoupled per Option B)
     "action_entropy_per_step",       // POST-gate action entropy EMA for SAC α/τ co-tuning
+    "rl_drawdown_stop",              // per-step drawdown penalty + hard stop-loss
     "rl_q_distill_lambda_controller",// audit 2026-05-24 rljzl followup: adaptive λ_distill via Schulman bounded step on KL_EMA vs target
     "rl_unit_state_update",          // SP20 P1+P5 audit fix: per-unit trade state machine — detects open/close/reverse transitions, sets up unit slot 0 entry+trail
     "rl_trail_mutate",               // SP20 P1+P5 audit fix (a7/a8 dead): TrailTighten/TrailLoosen mutate unit_trail_distance bounded MIN/MAX, symmetric reciprocal adjust

crates/ml-alpha/cuda/rl_confidence_gate.cu

@@ -65,6 +65,7 @@ extern "C" __global__ void rl_confidence_gate(
 
     const int action = actions[b];
     if (action == ACTION_HOLD) return;
+    if (action == 3 || action == 4) return;  // FlatFromLong/Short: never gate exits
 
     const float threshold  = isv[RL_CONF_GATE_THRESHOLD_INDEX];
     const float lambda     = isv[RL_CONF_GATE_LAMBDA_INDEX];

crates/ml-alpha/cuda/rl_drawdown_stop.cu

@@ -0,0 +1,48 @@
+// rl_drawdown_stop.cu — per-step drawdown penalty + hard stop-loss.
+//
+// Runs AFTER rl_trade_context_update (which computes unrealized_r).
+// Reads unrealized_r from trade_context_d and applies:
+//
+//   1. Per-step drawdown penalty: adds min(0, unrealized_r) × PENALTY_RATE
+//      to rewards[b]. Creates continuous exit gradient on losing positions.
+//      Penalty-only (never positive) = no exposure incentive.
+//
+//   2. Hard stop-loss: when unrealized_r < -STOP_THRESHOLD, force-close
+//      by setting dones[b] = 1. The realized PnL at this point becomes
+//      the final trade reward. Caps max loss per trade.
+//
+// Both thresholds are ISV-driven per feedback_isv_for_adaptive_bounds.
+// Grid: ceil(B/32), Block: 32.
+
+#define RL_DRAWDOWN_PENALTY_RATE_INDEX  586
+#define RL_STOP_LOSS_THRESHOLD_INDEX    587
+#define TRADE_CONTEXT_STRIDE            5
+#define UNREALIZED_R_OFFSET             1
+
+extern "C" __global__ void rl_drawdown_stop(
+    const float* __restrict__ trade_context,  // [B, TRADE_CONTEXT_STRIDE]
+    float*       __restrict__ rewards,         // [B] IN/OUT
+    float*       __restrict__ dones,           // [B] IN/OUT
+    float*       __restrict__ isv,
+    int          b_size
+) {
+    const int b = blockIdx.x * blockDim.x + threadIdx.x;
+    if (b >= b_size) return;
+
+    const float unrealized_r = trade_context[b * TRADE_CONTEXT_STRIDE + UNREALIZED_R_OFFSET];
+    const float penalty_rate = isv[RL_DRAWDOWN_PENALTY_RATE_INDEX];
+    const float stop_thresh  = isv[RL_STOP_LOSS_THRESHOLD_INDEX];
+
+    // Per-step drawdown penalty: only fires when losing (unrealized_r < 0).
+    // min(0, x) * rate is always <= 0.
+    if (unrealized_r < 0.0f) {
+        rewards[b] += unrealized_r * penalty_rate;
+    }
+
+    // Hard stop-loss: force-close when drawdown exceeds threshold.
+    // unrealized_r is normalized by initial_r (trade-level vol), so
+    // threshold is in units of "initial risk multiples."
+    if (unrealized_r < -stop_thresh && dones[b] < 0.5f) {
+        dones[b] = 1.0f;
+    }
+}

crates/ml-alpha/cuda/rl_reward_clamp_controller.cu

@@ -230,6 +230,15 @@ extern "C" __global__ void rl_reward_clamp_controller(
     // Suppress unused warnings on diagnostic-only state.
     (void) margin;
 
+    // Adaptive LOSS clamp from observed neg/pos ratio. WIN stays
+    // structural at 1.0; LOSS adapts to actual loss magnitude to
+    // give Q accurate resolution without over-allocating to rare tails.
+    if (ema_new > 0.0f && neg_new > 0.0f) {
+        const float observed_ratio = neg_new / ema_new;
+        const float adaptive_loss = fmaxf(1.0f, fminf(observed_ratio * 1.1f, 3.0f));
+        isv[RL_REWARD_CLAMP_LOSS_INDEX] = adaptive_loss;
+    }
+
     // ── Step 5: C51 atom span adaptation from observed reward EMAs. ──
     //
     // G.2 disabled this because atom-span growth + clamp-bound growth

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

@@ -1111,5 +1111,16 @@ pub const RL_ACTION_ENTROPY_EMA_INDEX: usize = 583;
 /// Bootstrap: 0.85 (15% exploration floor).
 pub const RL_CONF_GATE_MAX_HOLD_FRAC_INDEX: usize = 584;
 
+/// Per-step drawdown penalty rate. Applied as min(0, unrealized_r) × rate
+/// to the reward each step a position is losing. Penalty-only (never
+/// positive) creates continuous exit gradient without exposure incentive.
+/// Bootstrap: 0.01 (1% of unrealized_r magnitude per step).
+pub const RL_DRAWDOWN_PENALTY_RATE_INDEX: usize = 586;
+
+/// Hard stop-loss threshold in initial-risk multiples. When unrealized_r
+/// drops below -threshold, the position force-closes (done=1). Caps max
+/// loss per trade. Bootstrap: 2.0 (2× initial risk).
+pub const RL_STOP_LOSS_THRESHOLD_INDEX: usize = 587;
+
 /// Last RL-allocated slot index (exclusive).
-pub const RL_SLOTS_END: usize = 585;
+pub const RL_SLOTS_END: usize = 588;

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

@@ -150,6 +150,8 @@ const COMPUTE_ADVANTAGE_RETURN_CUBIN: &[u8] =
     include_bytes!(concat!(env!("OUT_DIR"), "/compute_advantage_return.cubin"));
 const ACTION_ENTROPY_PER_STEP_CUBIN: &[u8] =
     include_bytes!(concat!(env!("OUT_DIR"), "/action_entropy_per_step.cubin"));
+const RL_DRAWDOWN_STOP_CUBIN: &[u8] =
+    include_bytes!(concat!(env!("OUT_DIR"), "/rl_drawdown_stop.cubin"));
 
 // Phase R4: GPU-resident action sampling kernels. Replace the host
 // Thompson + host argmax + host log-softmax loops the flawed Phase F
@@ -566,6 +568,8 @@ pub struct IntegratedTrainer {
     compute_advantage_return_fn: CudaFunction,
     _action_entropy_per_step_module: Arc<CudaModule>,
     action_entropy_per_step_fn: CudaFunction,
+    _rl_drawdown_stop_module: Arc<CudaModule>,
+    rl_drawdown_stop_fn: CudaFunction,
 
     // ── Phase R4: GPU action sampling ─────────────────────────────────
     _rl_action_kernel_module: Arc<CudaModule>,
@@ -1327,6 +1331,12 @@ impl IntegratedTrainer {
         let action_entropy_per_step_module = ctx
             .load_cubin(ACTION_ENTROPY_PER_STEP_CUBIN.to_vec())
             .context("load action_entropy_per_step cubin")?;
+        let rl_drawdown_stop_module = ctx
+            .load_cubin(RL_DRAWDOWN_STOP_CUBIN.to_vec())
+            .context("load rl_drawdown_stop cubin")?;
+        let rl_drawdown_stop_fn = rl_drawdown_stop_module
+            .load_function("rl_drawdown_stop")
+            .context("load rl_drawdown_stop")?;
         let action_entropy_per_step_fn = action_entropy_per_step_module
             .load_function("action_entropy_per_step")
             .context("load action_entropy_per_step")?;
@@ -2323,6 +2333,8 @@ impl IntegratedTrainer {
             ema_update_per_step_fn,
             _action_entropy_per_step_module: action_entropy_per_step_module,
             action_entropy_per_step_fn,
+            _rl_drawdown_stop_module: rl_drawdown_stop_module,
+            rl_drawdown_stop_fn,
             _compute_advantage_return_module: compute_advantage_return_module,
             compute_advantage_return_fn,
             _rl_action_kernel_module: rl_action_kernel_module,
@@ -2673,7 +2685,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); 113] = [
                 // Static seeds for the adaptive reward-clamp controller —
                 // these are the initial values that
                 // `rl_reward_clamp_controller` will replace once it
@@ -2840,6 +2852,8 @@ 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),
+                (crate::rl::isv_slots::RL_DRAWDOWN_PENALTY_RATE_INDEX,    0.01),
+                (crate::rl::isv_slots::RL_STOP_LOSS_THRESHOLD_INDEX,      2.0),
             ];
             for (slot, value) in isv_constants.iter() {
                 let slot_i32 = *slot as i32;
@@ -6011,6 +6025,29 @@ impl IntegratedTrainer {
             }
         }
 
+        // Per-step drawdown penalty + hard stop-loss. Reads unrealized_r
+        // from trade_context_d (just written above). Modifies rewards_d
+        // and dones_d in place before apply_reward_scale.
+        {
+            let grid_x = ((b_size as u32) + 31) / 32;
+            let b_size_i = b_size as i32;
+            let mut args = RawArgs::new();
+            args.push_ptr(self.trade_context_d.raw_ptr());
+            args.push_ptr(self.rewards_d.raw_ptr());
+            args.push_ptr(self.dones_d.raw_ptr());
+            args.push_ptr(self.isv_dev_ptr);
+            args.push_i32(b_size_i);
+            let mut ptrs = args.build_arg_ptrs();
+            unsafe {
+                raw_launch(
+                    self.rl_drawdown_stop_fn.cu_function(),
+                    (grid_x.max(1), 1, 1), (32, 1, 1), 0,
+                    self.raw_stream,
+                    &mut ptrs[..args.len()],
+                ).map_err(|e| anyhow::anyhow!("rl_drawdown_stop: {:?}", e))?;
+            }
+        }
+
         // Multi-resolution streaming features — time-weighted EMA at 3 horizons.
         {
             let bid_px_d = lobsim.bid_px_d();

crates/ml-backtesting/src/sim/mod.rs

@@ -1143,7 +1143,6 @@ impl LobSimCuda {
                 .arg(&n)
                 .launch(cfg)?;
         }
-        self.stream.synchronize()?;
         Ok(())
     }