diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu
index eeb0d9a6c..a5ed8af16 100644
--- a/crates/ml/src/cuda_pipeline/experience_kernels.cu
+++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu
@@ -722,59 +722,22 @@ extern "C" __global__ void experience_env_step(
         }
     }
 
-    /* ---- Execute trade (shared: trade_physics.cuh) ---- */
-    int order_type_idx = decode_order_type(action_idx, b1_size, b2_size);
-    float tx_cost = execute_trade(&position, &cash, target_position, raw_close,
-                                  tx_cost_multiplier, 0.0f, max_position,
-                                  order_type_idx, spread_scale);
-
-    /* ---- Mark-to-market PnL for this timestep ---- */
-    float raw_pnl = position * (raw_next - raw_close);
-    /* Old position's PnL — needed for reversal trade tracking. */
-    float old_pos_pnl = pre_trade_position * (raw_next - raw_close);
-
-    /* ---- Mark-to-market equity (needed for peak_equity and reward) ---- */
-    float equity = cash + position * raw_next;
-    peak_equity = (equity > peak_equity) ? equity : peak_equity;
-
-    /* ==== Trade lifecycle tracking (v4: segment-based, reversal-aware) ==== */
-    float is_flat = (fabsf(position) < 0.001f) ? 1.0f : 0.0f;
-    int curr_sign = (position > 0.001f) ? 1 : ((position < -0.001f) ? -1 : 0);
-
-    int entering_trade  = (prev_sign == 0 && curr_sign != 0);
-    int exiting_trade = trail_triggered ? 1 : 0;
-    int reversing_trade = (prev_sign != 0 && curr_sign != 0 && prev_sign != curr_sign);
-
-    /* NOTE: reversal P&L booking moved AFTER hold enforcement (below). */
-    float reversal_return = 0.0f;
-
-    /* On entry from flat */
-    if (entering_trade) {
-        entry_price = raw_close;
-        trade_start_pnl = ps[11];
-    }
-
     /* ════════════════════════════════════════════════════════════════════
-     * HOLD ENFORCEMENT: block ALL position changes during minimum hold.
-     * Applies to BOTH flat exits AND reversals.
+     * HOLD ENFORCEMENT: block exits and reversals during minimum hold.
+     *
+     * Runs BEFORE execute_trade (same order as backtest kernel).
+     * If hold is violated, override target_position to keep old position.
+     * execute_trade then sees delta≈0 and does nothing.
      *
      * Layer 2 safety net — Layer 1 (action masking) prevents most hold
      * violations, but epsilon exploration can still slip through.
-     * Uses enforce_hold() from trade_physics.cuh for the decision.
      * ════════════════════════════════════════════════════════════════════ */
     int is_last_bar = (bar_idx >= total_bars - 1) ? 1 : 0;
-    float enforced_position = enforce_hold(ps[0], position, hold_time, min_hold_bars, is_last_bar);
-    int hold_violation = (fabsf(enforced_position - position) > 0.001f);
-
-    /* wants_exit / wants_reversal needed downstream for trade lifecycle */
-    int wants_exit = (prev_sign != 0 && curr_sign == 0);
-    int wants_reversal = (prev_sign != 0 && curr_sign != 0 && prev_sign != curr_sign);
+    float enforced_position = enforce_hold(ps[0], target_position, hold_time, min_hold_bars, is_last_bar);
+    int hold_violation = (fabsf(enforced_position - target_position) > 0.001f);
 
     if (hold_violation) {
-        /* Override: keep previous position, cancel the action */
-        position = ps[0];
-        cash = ps[1];
-        is_flat = (fabsf(position) < 0.001f) ? 1.0f : 0.0f;
+        target_position = enforced_position;  /* keep old position */
 
         /* Fix the stored action to match held exposure (prevent action aliasing) */
         float prev_exposure_frac = ps[0] / (max_position > 0.0f ? max_position : 1.0f);
@@ -785,18 +748,40 @@ extern "C" __global__ void experience_env_step(
         int original_urgency = (b2_size > 0) ? action_idx % b2_size : 0;
         action_idx = held_exposure * b1_size * b2_size + original_order * b2_size + original_urgency;
         out_actions[out_off] = action_idx;
-
-        /* Recompute signs from overridden position */
-        curr_sign = prev_sign;
-        wants_exit = 0;
-        wants_reversal = 0;
     }
 
-    /* Compute final trade lifecycle from (possibly overridden) state */
-    exiting_trade = wants_exit;
-    reversing_trade = wants_reversal;
+    /* ---- Execute trade (shared: trade_physics.cuh) ---- */
+    int order_type_idx = decode_order_type(action_idx, b1_size, b2_size);
+    float tx_cost = execute_trade(&position, &cash, target_position, raw_close,
+                                  tx_cost_multiplier, 0.0f, max_position,
+                                  order_type_idx, spread_scale);
+
+    /* ---- Mark-to-market PnL for this timestep ---- */
+    float raw_pnl = position * (raw_next - raw_close);
+    float old_pos_pnl = pre_trade_position * (raw_next - raw_close);
+
+    /* ---- Mark-to-market equity ---- */
+    float equity = cash + position * raw_next;
+    peak_equity = (equity > peak_equity) ? equity : peak_equity;
+
+    /* ==== Trade lifecycle tracking ==== */
+    float is_flat = (fabsf(position) < 0.001f) ? 1.0f : 0.0f;
+    int curr_sign = (position > 0.001f) ? 1 : ((position < -0.001f) ? -1 : 0);
+
+    int entering_trade = (prev_sign == 0 && curr_sign != 0);
+    int wants_exit = hold_violation ? 0 : (prev_sign != 0 && curr_sign == 0);
+    int wants_reversal = hold_violation ? 0 : (prev_sign != 0 && curr_sign != 0 && prev_sign != curr_sign);
+    int exiting_trade = trail_triggered ? 1 : wants_exit;
+    int reversing_trade = wants_reversal;
     int segment_complete = exiting_trade || reversing_trade;
 
+    float reversal_return = 0.0f;
+
+    if (entering_trade) {
+        entry_price = raw_close;
+        trade_start_pnl = ps[11];
+    }
+
     /* On reversal: close old segment, open new one.
      * MUST run AFTER hold enforcement — if hold guard cancelled the reversal,
      * reversing_trade is 0 and this block correctly does nothing. */