fix(cuda): confidence gate exploration floor + symmetric threshold decay

The gate created a self-reinforcing Hold trap: gate forces Hold → Q
learns Hold is best → conf=0 for trading actions → gate blocks
everything → Hold=100%. Even max SAC (α=2.0, τ=50) couldn't break it.

Two fixes:

1. Exploration slots: the first `b_size × (1 - max_hold_frac)` batch
   elements are NEVER gated. At max_hold_frac=0.85 with b=1024, 154
   slots always use the Thompson-sampled action. This guarantees Q
   sees trading outcomes and can learn their value.

2. Symmetric threshold decay: the adaptive threshold had 100× asymmetry
   (10% raise vs 0.1% lower). When Hold exceeded target, the threshold
   barely decreased. Now both directions use THRESHOLD_ADJUST_RATE=1.1.

New ISV slot: RL_CONF_GATE_MAX_HOLD_FRAC_INDEX=584 (bootstrap 0.85).

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

crates/ml-alpha/cuda/rl_confidence_gate.cu

@@ -33,6 +33,7 @@
 #define RL_STEP_COUNTER_ISV_INDEX      548
 #define RL_HOLD_TARGET_FRAC_INDEX      575
 #define RL_HOLD_FRAC_EMA_INDEX         576
+#define RL_CONF_GATE_MAX_HOLD_FRAC_INDEX 584
 #define THRESHOLD_MIN                  0.05f
 #define THRESHOLD_MAX                  0.95f
 #define HOLD_EMA_ALPHA                 0.1f
@@ -54,6 +55,14 @@ extern "C" __global__ void rl_confidence_gate(
     const int warmup = (int)isv[RL_GATE_WARMUP_STEPS_INDEX];
     if (current_step < warmup) return;
 
+    // Exploration slots: the first `min_explore` batch elements are
+    // NEVER gated. This guarantees a minimum fraction of trading
+    // actions for Q to learn from, breaking the self-reinforcing
+    // Hold trap where gate→Hold→Q learns Hold→conf=0→gate.
+    const float max_hold = isv[RL_CONF_GATE_MAX_HOLD_FRAC_INDEX];
+    const int min_explore = (int)((float)b_size * (1.0f - max_hold));
+    if (b < min_explore) return;
+
     const int action = actions[b];
     if (action == ACTION_HOLD) return;
 
@@ -129,12 +138,14 @@ extern "C" __global__ void rl_confidence_gate(
         // Schulman-style bounded adjustment: if Hold fraction is below
         // target, raise threshold (gate more aggressively). If above
         // target, lower threshold (allow more trading).
+        // Symmetric rates — the old 100× asymmetry caused the threshold
+        // to ratchet up and never recover.
         const float hold_target = isv[RL_HOLD_TARGET_FRAC_INDEX];
         float new_threshold = threshold;
         if (new_ema < hold_target * 0.9f) {
             new_threshold = threshold * THRESHOLD_ADJUST_RATE;
         } else if (new_ema > hold_target * 1.1f) {
-            new_threshold = threshold / (1.0f + (THRESHOLD_ADJUST_RATE - 1.0f) * 0.01f);
+            new_threshold = threshold / THRESHOLD_ADJUST_RATE;
         }
         new_threshold = fmaxf(THRESHOLD_MIN, fminf(new_threshold, THRESHOLD_MAX));
         isv[RL_CONF_GATE_THRESHOLD_INDEX] = new_threshold;

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

@@ -1104,5 +1104,12 @@ pub const RL_SAC_ENTROPY_TARGET_INDEX: usize = 582;
 /// misleading when the confidence gate overrides actions to Hold.
 pub const RL_ACTION_ENTROPY_EMA_INDEX: usize = 583;
 
+/// Maximum Hold fraction the confidence gate may produce. Batch
+/// elements with `b < b_size × (1 - max_hold_frac)` are exploration
+/// slots that the gate never overrides. Prevents the gate from
+/// creating a 100% Hold trap where Q never sees trading actions.
+/// Bootstrap: 0.85 (15% exploration floor).
+pub const RL_CONF_GATE_MAX_HOLD_FRAC_INDEX: usize = 584;
+
 /// Last RL-allocated slot index (exclusive).
-pub const RL_SLOTS_END: usize = 584;
+pub const RL_SLOTS_END: usize = 585;

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); 110] = [
+            let isv_constants: [(usize, f32); 111] = [
                 // Static seeds for the adaptive reward-clamp controller —
                 // these are the initial values that
                 // `rl_reward_clamp_controller` will replace once it
@@ -2839,6 +2839,7 @@ impl IntegratedTrainer {
                 // maintain target entropy. H_target = 0.7 × ln(11) ≈ 1.68.
                 (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),
             ];
             for (slot, value) in isv_constants.iter() {
                 let slot_i32 = *slot as i32;