feat: dynamic trade management — stop-loss, take-profit, regime-adaptive

Three critical fixes from H100 Trial 0 analysis:

1. REMOVE IDLE PENALTY — was making Flat the worst Q-value (Q=-19.4),
   forcing the model to trade on 95% of bars. Now Flat = reward 0.0.
   DSR already penalizes inaction through the Sharpe denominator.

2. TRADE COMPLETION SCALING — ×1000 clamped to ±1 made ALL trades look
   the same (can't distinguish 1-tick from 10-tick return). Now ×200
   with clamp ±2, plus sqrt(hold_time) bonus for holding longer.

3. DYNAMIC TRADE MANAGEMENT — model learns ENTRY quality, exits are
   managed by the system:
   - Stop loss: -0.3% × vol_mult (CUSUM-adaptive, up to -0.75%)
   - Take profit: +0.5% × vol_mult × trend_mult (up to +2.5% in trends)
   - Time stop: 100 bars × vol_mult (up to 250 in volatile markets)
   - Min hold: 5 bars before model can exit voluntarily
   - Auto-exit forces flat and computes trade_completion_reward

   Regime-adaptive via ADX (trend strength) and CUSUM (volatility)
   features already in the GPU memory.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

crates/ml/src/cuda_pipeline/experience_kernels.cu

@@ -629,21 +629,82 @@ extern "C" __global__ void experience_env_step(
     if (exiting_trade && hold_time > 0.0f) {
         float trade_pnl = ps[11] + raw_pnl - trade_start_pnl;
         float trade_return = trade_pnl / (prev_equity > 1.0f ? prev_equity : 1.0f);
-        /* Scale: 1 ES tick ($12.50 on $35K) = 0.036% return × 1000 = 0.36 reward.
-         * Must be stronger than idle penalty (max 0.05) to incentivize quality trades.
-         * Clamp to [-1, +1] so exceptional trades don't dominate. */
-        trade_completion_reward = fmaxf(-1.0f, fminf(1.0f, trade_return * 1000.0f));
+        /* Scale ×200: 1 ES tick = 0.036% × 200 = 0.072 reward.
+         * 10 ticks (a meaningful move) = 0.72 reward.
+         * Clamp [-2, +2] so the model sees the MAGNITUDE of returns,
+         * not just the sign. Old ×1000 clamped to ±1 made all trades equal. */
+        trade_completion_reward = fmaxf(-2.0f, fminf(2.0f, trade_return * 200.0f));
+        /* Bonus for holding longer — reward scales with sqrt(hold_time).
+         * A 50-bar trade gets 7x bonus vs a 1-bar trade.
+         * This teaches the model to hold for meaningful periods. */
+        float hold_bonus = sqrtf(hold_time) * 0.02f;
+        trade_completion_reward += (trade_completion_reward > 0.0f) ? hold_bonus : -hold_bonus;
     }
 
-    /* Idle counter (flat bars) */
+    /* Dynamic trade management: once entered, the trade runs until a stop
+     * condition is met. The model learns ENTRY quality, not exit timing.
+     *
+     * Stop conditions (checked when model tries to exit):
+     * 1. Stop loss: unrealized P&L < -0.3% of equity → force exit (bad entry)
+     * 2. Take profit: unrealized P&L > +0.5% of equity → force exit (good entry)
+     * 3. Time stop: held > 100 bars → force exit (stale position)
+     * 4. Model exit: model chooses Flat AND held >= 5 bars (learned exit)
+     *
+     * If none met and model tries to exit early (<5 bars), override to hold. */
+    float unrealized_trade_pnl = 0.0f;
+    if (!was_flat && hold_time > 0.0f) {
+        unrealized_trade_pnl = (ps[11] + raw_pnl - trade_start_pnl) / (prev_equity > 1.0f ? prev_equity : 1.0f);
+    }
+
+    /* Regime-adaptive stop levels using ADX and CUSUM (already in feature vector).
+     * High volatility (CUSUM > 0.5) → wider stops (more room for noise).
+     * Strong trend (ADX > 25) → wider take-profit (let trends run).
+     * Read from features buffer (same as used for regime scaling earlier). */
+    float stop_adx = 0.0f, stop_cusum = 0.0f;
+    if (bar_idx < total_bars && market_dim >= 42) {
+        long long feat_off = (long long)bar_idx * market_dim;
+        stop_adx   = features[feat_off + 40];
+        stop_cusum = features[feat_off + 41];
+    }
+    /* Vol scale: [1.0, 2.5] — wider stops in volatile markets */
+    float vol_mult = 1.0f + (stop_cusum > 0.5f ? stop_cusum : 0.0f);
+    vol_mult = (vol_mult > 2.5f) ? 2.5f : vol_mult;
+    /* Trend scale: [1.0, 2.0] — wider TP in strong trends */
+    float trend_mult = 1.0f + (stop_adx > 25.0f ? (stop_adx - 25.0f) / 50.0f : 0.0f);
+    trend_mult = (trend_mult > 2.0f) ? 2.0f : trend_mult;
+
+    float sl_level = -0.003f * vol_mult;           /* base -0.3%, up to -0.75% in vol */
+    float tp_level =  0.005f * vol_mult * trend_mult; /* base +0.5%, up to +2.5% in trending+vol */
+    float time_stop_bars = 100.0f * vol_mult;       /* base 100 bars, up to 250 in vol */
+
+    /* Auto-exit triggers (force flat regardless of model's action) */
+    int auto_stop_loss  = (!was_flat && unrealized_trade_pnl < sl_level) ? 1 : 0;
+    int auto_take_profit = (!was_flat && unrealized_trade_pnl > tp_level) ? 1 : 0;
+    int auto_time_stop  = (!was_flat && hold_time >= time_stop_bars) ? 1 : 0;
+
+    if (auto_stop_loss || auto_take_profit || auto_time_stop) {
+        /* Force exit — override model's action to flat */
+        position = 0.0f;
+        float exit_cost = fabsf(ps[0]) * raw_close * tx_cost_multiplier;
+        cash = ps[1] + ps[0] * raw_close - exit_cost;
+        is_flat = 1.0f;
+        exiting_trade = 1;
+    } else if (!was_flat && is_flat > 0.5f && hold_time < 5.0f && hold_time > 0.0f) {
+        /* Model tries to exit too early — override to hold */
+        position = ps[0];
+        cash = ps[1];
+        is_flat = 0.0f;
+        exiting_trade = 0;
+    }
+
+    /* Idle counter (flat bars) — NO penalty. DSR already penalizes inaction
+     * (zero returns reduce Sharpe). Idle penalty was making Flat the worst
+     * Q-value, forcing the model to trade on every bar. */
     if (is_flat > 0.5f) {
         flat_counter += 1.0f;
     } else {
         flat_counter = 0.0f;
     }
-    /* Idle penalty: tiny nudge (max 0.01), must be << trade completion reward.
-     * Old max 0.05 competed with 1-tick trade reward 0.036 — wrong incentive. */
-    float idle_penalty_t = fminf(0.01f, flat_counter * 0.0002f);
 
     /* ==== Component 5: Position-time decay (losing positions only) ==== */
     float hold_time_penalty = 0.0f;
@@ -695,8 +756,9 @@ extern "C" __global__ void experience_env_step(
                  + w_dsr * dsr_t
                  - tx_cost;
     } else if (is_flat > 0.5f) {
-        /* NEAR-ZERO: flat is almost free */
-        reward = -w_idle * idle_penalty_t;
+        /* ZERO: flat costs nothing. No idle penalty — DSR handles it.
+         * The model must learn that flat is a VALID action, not a penalty. */
+        reward = 0.0f;
     } else {
         /* WEAK: 0.1x dense shaping while holding */
         reward = 0.1f * (w_dsr * dsr_t + w_pnl * asym_pnl)