diff --git a/config/training/dqn-hyperopt.toml b/config/training/dqn-hyperopt.toml
index 955ef9741..ea27ad6b6 100644
--- a/config/training/dqn-hyperopt.toml
+++ b/config/training/dqn-hyperopt.toml
@@ -116,7 +116,6 @@ ioc_fill_prob = 0.85
 limit_fill_min = 0.30
 limit_fill_max = 0.80
 tx_cost_multiplier = 0.18  # IBKR ES: 0.18 bps = $4.50/contract RT
-min_hold_bars = 10  # ~26s hold at 23 volume bars/min
 spread_cost_frac = 0.50
 spread_capture_frac = 0.50
 
diff --git a/config/training/dqn-localdev.toml b/config/training/dqn-localdev.toml
index 28dddd711..cde26cf39 100644
--- a/config/training/dqn-localdev.toml
+++ b/config/training/dqn-localdev.toml
@@ -47,7 +47,6 @@ min_epochs_before_stopping = 80
 [experience]
 gpu_timesteps_per_episode = 200
 gpu_n_episodes = 128
-min_hold_bars = 10  # ~26s hold at 23 volume bars/min
 
 [experience.fill_simulation]
 ioc_fill_prob = 0.85
diff --git a/config/training/dqn-production.toml b/config/training/dqn-production.toml
index cb0633a9a..eb51640ea 100644
--- a/config/training/dqn-production.toml
+++ b/config/training/dqn-production.toml
@@ -67,7 +67,6 @@ min_epochs_before_stopping = 80
 gpu_n_episodes = 4096
 initial_capital = 35000.0
 tx_cost_multiplier = 0.18  # bps: IBKR ES RT = $4.50/contract ($0.85 comm + $1.38 exch + $0.02 reg × 2 sides). At ES=$5100: 0.18 bps × 5100 × 0.0001 = 0.09 pts = $4.50
-min_hold_bars = 50  # ~6.5 min at ~11.5K imbalance bars/day. ISV extends to 50-150 bars adaptively.
 
 [experience.fill_simulation]
 ioc_fill_prob = 0.85
diff --git a/config/training/dqn-smoketest.toml b/config/training/dqn-smoketest.toml
index f1bf1aff1..aaff92441 100644
--- a/config/training/dqn-smoketest.toml
+++ b/config/training/dqn-smoketest.toml
@@ -59,7 +59,6 @@ min_epochs_before_stopping = 10
 gpu_timesteps_per_episode = 50
 gpu_n_episodes = 4
 tx_cost_multiplier = 0.18  # IBKR ES: 0.18 bps = $4.50/contract RT
-min_hold_bars = 10  # ~26s hold at 23 volume bars/min
 
 [experience.fill_simulation]
 ioc_fill_prob = 0.85
diff --git a/crates/ml/examples/hyperopt_baseline_rl.rs b/crates/ml/examples/hyperopt_baseline_rl.rs
index 308522e49..c51700a86 100644
--- a/crates/ml/examples/hyperopt_baseline_rl.rs
+++ b/crates/ml/examples/hyperopt_baseline_rl.rs
@@ -187,10 +187,6 @@ struct Args {
     #[arg(long, default_value = "35000")]
     initial_capital: f64,
 
-    /// Minimum bars to hold a position before allowing exit.
-    #[arg(long)]
-    min_hold_bars: Option<usize>,
-
     /// Path to MBP-10 order book data for OFI features (order flow imbalance)
     #[arg(long)]
     mbp10_data_dir: Option<PathBuf>,
@@ -573,9 +569,6 @@ fn main() -> Result<()> {
     info!("Seed: {}", args.seed);
     info!("Tx cost: {} bps + {} tick spread (tick_size={})", args.tx_cost_bps, args.spread_ticks, args.tick_size);
     info!("Initial capital: ${:.0}", args.initial_capital);
-    if let Some(mhb) = args.min_hold_bars {
-        info!("Min hold bars: {} (CLI override)", mhb);
-    }
 
     let cpus = std::thread::available_parallelism()
         .map(|n| n.get())
diff --git a/crates/ml/examples/train_baseline_rl.rs b/crates/ml/examples/train_baseline_rl.rs
index ee9d79649..2f52e0874 100644
--- a/crates/ml/examples/train_baseline_rl.rs
+++ b/crates/ml/examples/train_baseline_rl.rs
@@ -245,11 +245,6 @@ struct Args {
     #[arg(long, default_value_t = 35_000.0)]
     initial_capital: f64,
 
-    /// Minimum bars to hold a position before allowing exit (churn prevention).
-    /// Lower = more trades, higher = fewer. Default from TOML (typically 5).
-    #[arg(long)]
-    min_hold_bars: Option<usize>,
-
     /// Named training profile to load from config/training/<profile>.toml.
     /// Profile values are applied after hyperopt JSON but before explicit CLI args.
     /// Known profiles: dqn-production, dqn-smoketest, dqn-hyperopt.
@@ -435,9 +430,6 @@ fn build_dqn_hyperparams(
     hyperparams.epochs = args.epochs;
     hyperparams.learning_rate = hp_f64(hp, "learning_rate").unwrap_or(args.learning_rate);
     hyperparams.initial_capital = args.initial_capital as f32;
-    if let Some(mhb) = args.min_hold_bars {
-        hyperparams.min_hold_bars = mhb;
-    }
 
     hyperparams
 }
diff --git a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu
index 6d5998338..159518120 100644
--- a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu
+++ b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu
@@ -11,7 +11,7 @@
 //   [6] cum_return   - cumulative log return
 //   [7] step_count   - number of completed steps
 //
-// Trade physics (action decoding, hold enforcement, tx costs, capital floor)
+// Trade physics (action decoding, tx costs, capital floor)
 // are shared with the training kernel via trade_physics.cuh.
 
 #define PORTFOLIO_STATE_SIZE 8
@@ -73,7 +73,9 @@ extern "C" __global__ void backtest_env_step(
     int          b1_size,
     int          b2_size,
     int          b3_size,
-    int          min_hold_bars,
+    float        holding_cost_rate,     // continuous inventory penalty per bar per unit position
+    float        churn_threshold,      // bars threshold for churn penalty
+    float        churn_penalty_scale,  // scale factor for churn penalty
     float        contract_multiplier,  // e.g. 50.0 for ES, 20.0 for NQ
     float        margin_pct            // e.g. 0.06 (6% initial margin)
 ) {
@@ -153,11 +155,6 @@ extern "C" __global__ void backtest_env_step(
     // Suppress unused variable warnings
     (void)open; (void)high; (void)low;
 
-    // ── Hold enforcement (shared: trade_physics.cuh) ─────────────────────
-    int is_last_bar = (current_step >= wlen - 1) ? 1 : 0;
-    target_exposure = enforce_hold(position, target_exposure, hold_time,
-                                   min_hold_bars, is_last_bar, NULL);
-
     // ── Trailing stop (shared: trade_physics.cuh) ────────────────────────
     // Exit when profit retreats from peak. Uses 0.5% base distance.
     // vol_scale=1.0, trend_scale=1.0: backtest has no ADX/CUSUM features,
@@ -172,7 +169,7 @@ extern "C" __global__ void backtest_env_step(
             float unrealized = position * (close - entry_price);
             trade_ret = unrealized / value;
         }
-        if (check_trailing_stop(hold_time, min_hold_bars,
+        if (check_trailing_stop(hold_time,
                                 max_equity, value,
                                 trade_ret, 0.005f, 1.0f, 1.0f)) {
             target_exposure = 0.0f;  // Force flat — trailing stop triggered
@@ -241,6 +238,23 @@ extern "C" __global__ void backtest_env_step(
     float step_ret = (value > 0.0f)
         ? (new_value - value) / value
         : 0.0f;
+
+    // ── Cost-driven hold timing (matches training kernel) ──────────────
+    // Inventory penalty: continuous cost proportional to position size.
+    float inventory_penalty = holding_cost_rate * fabsf(position);
+    step_ret -= inventory_penalty;
+
+    // Churn penalty: graduated cost for rapid position flips.
+    {
+        int prev_sign_bt = (prev_position > 0.001f) ? 1 : ((prev_position < -0.001f) ? -1 : 0);
+        int curr_sign_bt = (position > 0.001f) ? 1 : ((position < -0.001f) ? -1 : 0);
+        int entering_bt = (prev_sign_bt == 0 && curr_sign_bt != 0);
+        if (entering_bt && hold_time < churn_threshold && hold_time > 0.0f) {
+            float churn_pen = churn_penalty_scale * (churn_threshold - hold_time) / churn_threshold;
+            step_ret -= churn_pen;
+        }
+    }
+
     float new_cum_return = cum_return + step_ret;
     float new_max = fmaxf(max_equity, new_value);
 
@@ -311,7 +325,9 @@ extern "C" __global__ void backtest_env_step_batch(
     int          b1_size,
     int          b2_size,
     int          b3_size,
-    int          min_hold_bars,
+    float        holding_cost_rate,
+    float        churn_threshold,
+    float        churn_penalty_scale,
     float        contract_multiplier,
     float        margin_pct
 ) {
@@ -382,11 +398,6 @@ extern "C" __global__ void backtest_env_step_batch(
         target_exposure = apply_margin_cap(target_exposure, value,
                                            margin_per_contract, max_equity);
 
-        /* ── Hold enforcement ─────────────────────────────────────────── */
-        int is_last_bar = (current_step >= wlen - 1) ? 1 : 0;
-        target_exposure = enforce_hold(position, target_exposure, hold_time,
-                                       min_hold_bars, is_last_bar, NULL);
-
         /* ── Trailing stop ────────────────────────────────────────────── */
         {
             float trade_ret = 0.0f;
@@ -396,7 +407,7 @@ extern "C" __global__ void backtest_env_step_batch(
                 float unrealized = position * (close - entry_price);
                 trade_ret = unrealized / value;
             }
-            if (check_trailing_stop(hold_time, min_hold_bars,
+            if (check_trailing_stop(hold_time,
                     max_equity, value,
                     trade_ret, 0.005f, 1.0f, 1.0f)) {
                 target_exposure = 0.0f;
@@ -460,6 +471,23 @@ extern "C" __global__ void backtest_env_step_batch(
         /* Step return */
         float step_ret = (value > 0.0f)
             ? (new_value - value) / value : 0.0f;
+
+        /* ── Cost-driven hold timing (matches training kernel) ────────── */
+        /* Inventory penalty: continuous cost proportional to position size. */
+        float inv_pen = holding_cost_rate * fabsf(position);
+        step_ret -= inv_pen;
+
+        /* Churn penalty: graduated cost for rapid position flips. */
+        {
+            int psb = (prev_position > 0.001f) ? 1 : ((prev_position < -0.001f) ? -1 : 0);
+            int csb = (position > 0.001f) ? 1 : ((position < -0.001f) ? -1 : 0);
+            int ent = (psb == 0 && csb != 0);
+            if (ent && hold_time < churn_threshold && hold_time > 0.0f) {
+                float cp = churn_penalty_scale * (churn_threshold - hold_time) / churn_threshold;
+                step_ret -= cp;
+            }
+        }
+
         cum_return = cum_return + step_ret;
         float new_max = fmaxf(max_equity, new_value);
         max_equity = new_max;
diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu
index b74dd3af2..62002f320 100644
--- a/crates/ml/src/cuda_pipeline/experience_kernels.cu
+++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu
@@ -94,6 +94,10 @@ __device__ __forceinline__ float asymmetric_soft_clamp(float x) {
 #define MARKET_DIM 42  /* 42 base features; 50 when OFI (MBP-10) is enabled */
 #endif
 
+#ifndef PORTFOLIO_DIM
+#define PORTFOLIO_DIM 3
+#endif
+
 #ifndef STATE_DIM
 #define STATE_DIM 48
 #endif
@@ -728,8 +732,8 @@ extern "C" __global__ void experience_state_gather(
  * @param b2_size                          order branch size (3)
  * @param b3_size                          urgency branch size (3)
  * @param q_gap_threshold                  min Q-gap for trade entry (0.0 = disabled)
- * @param portfolio_states  [N, 20]        read-only for hold enforcement (bf16)
- * @param min_hold_bars                    minimum bars to hold before switching
+ * @param portfolio_states  [N, 20]        read-only (unused, kept for ABI compat)
+ * @param min_hold_bars                    unused (kept for ABI compat, pass 0)
  * @param max_position                     max allowed position (host scalar)
  */
 extern "C" __global__ void experience_action_select(
@@ -799,54 +803,12 @@ extern "C" __global__ void experience_action_select(
 
     int dir_idx, mag_idx, a2, a3;
 
-    /* Hold enforcement: during minimum hold period, force current direction+magnitude.
-     * Prevents both greedy and random exploration from changing position.
-     * Portfolio state: ps[0] = position, ps[10] = hold_time.
-     * When portfolio_states is NULL (backtest evaluator), skip hold enforcement. */
-    int in_hold = 0;
-    float cur_position = 0.0f;
-    if (min_hold_bars > 0 && portfolio_states != NULL) {
-        int ps_base = i * PORTFOLIO_STRIDE;
-        float hold_time_val = portfolio_states[ps_base + 10];
-        cur_position = portfolio_states[ps_base + 0];
-        /* Adaptive hold: base + ISV-driven extension (same formula as Layer 2) */
-        int adaptive_hold = min_hold_bars;
-        if (isv_signals_ptr != NULL) {
-            float stability = 1.0f / (1.0f + isv_signals_ptr[1]);
-            float confidence = 1.0f - fminf(isv_signals_ptr[3], 1.0f);
-            int extension = (int)(stability * confidence * 8.0f);
-            if (isv_signals_ptr[5] < -0.001f) extension = 0;
-            adaptive_hold = min_hold_bars + extension;
-        }
-        in_hold = (hold_time_val > 0.0f && hold_time_val < (float)adaptive_hold);
-    }
+    /* Hold enforcement removed — replaced by cost-driven hold timing.
+     * The model can always take any action; costs make bad actions expensive.
+     * min_hold_bars and portfolio_states params kept for ABI compatibility. */
+    (void)min_hold_bars;
+    (void)portfolio_states;
 
-    if (in_hold) {
-        /* Compute current direction+magnitude from position for hold enforcement */
-        float pos_f = (cur_position);
-        if (pos_f > 0.001f) {
-            dir_idx = 2;  /* Long */
-        } else if (pos_f < -0.001f) {
-            dir_idx = 0;  /* Short */
-        } else {
-            dir_idx = 1;  /* Flat */
-        }
-        /* Approximate magnitude from position fraction */
-        float abs_frac = fabsf(pos_f) / fmaxf(max_position, 0.001f);
-        if (dir_idx == 1) {
-            mag_idx = 1;  /* Flat — magnitude irrelevant */
-        } else if (abs_frac <= 0.375f) {
-            mag_idx = 0;  /* Quarter */
-        } else if (abs_frac <= 0.75f) {
-            mag_idx = 1;  /* Half */
-        } else {
-            mag_idx = 2;  /* Full */
-        }
-        /* Zero Q-gap during hold — conviction sizing is meaningless when forced */
-        if (out_q_gaps != NULL) {
-            out_q_gaps[i] = 0.0f;
-        }
-    } else {
     /* Branch 0: direction — BOLTZMANN SOFTMAX selection.
      *
      * C51's expected Q structurally favors Flat (zero PnL variance → tightest
@@ -954,8 +916,6 @@ extern "C" __global__ void experience_action_select(
             }
         }
     }
-    } /* end else (not in_hold) */
-
     /* Branch 2: order type — Boltzmann for both eval and training */
     if (!eval_mode && philox_uniform(i, timestep, rng_ctr++) < eps_ord) {
         int r = (int)(philox_uniform(i, timestep, rng_ctr++) * (float)b2_size);
@@ -1036,9 +996,8 @@ extern "C" __global__ void experience_action_select(
 
     /* Output Q-gap for conviction-based position sizing.
      * Q-gap = Q(best_direction) - Q(Flat). Higher = more conviction.
-     * Used by env_step to scale target_position.
-     * Skip when in_hold — already zeroed in the hold branch above. */
-    if (out_q_gaps != NULL && !in_hold) {
+     * Used by env_step to scale target_position. */
+    if (out_q_gaps != NULL) {
         int flat_idx = 1; /* Flat = index 1 for direction(3) */
         float q_best = q_b0[argmax_n(q_b0, b0_size)];
         float q_flat = q_b0[flat_idx];
@@ -1108,7 +1067,9 @@ extern "C" __global__ void experience_env_step(
     const float* __restrict__ cvar_scales,  /* [N] or NULL — CVaR position scaling */
     const float* __restrict__ q_gaps,       /* [N] or NULL — Q-gap conviction scaling */
     float*            raw_returns_out, /* [N, L] output: true per-bar portfolio return */
-    int min_hold_bars,  /* minimum bars to hold before exiting or reversing */
+    float holding_cost_rate,  /* continuous inventory penalty per bar per unit position */
+    float churn_threshold,    /* bars threshold for churn penalty (0 = disabled) */
+    float churn_penalty_scale, /* scale factor for churn penalty */
     float spread_cost,  /* bid-ask spread cost per unit (matches backtest) */
     float contract_multiplier,  /* e.g. 50.0 for ES, 20.0 for NQ */
     float margin_pct,           /* e.g. 0.06 (6% initial margin) */
@@ -1409,7 +1370,7 @@ extern "C" __global__ void experience_env_step(
          * The model observes ADX/CUSUM via state features and can learn
          * to adjust position sizing itself — the trailing stop is a safety
          * net, not a strategy component. */
-        if (check_trailing_stop(hold_time, min_hold_bars, peak_equity,
+        if (check_trailing_stop(hold_time, peak_equity,
                                 prev_equity, unrealized_trade_pnl,
                                 0.005f, 1.0f, 1.0f)) {
             target_position = 0.0f;  /* force flat — execute_trade handles exit */
@@ -1417,56 +1378,10 @@ extern "C" __global__ void experience_env_step(
         }
     }
 
-    /* ================================================================
-     * 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.
-     * ================================================================ */
-    int is_last_bar = (bar_idx >= total_bars - 1) ? 1 : 0;
-    /* When a plan is active, the plan's target_bars IS the hold duration.
-     * Adaptive hold is a safety net for plan-less trading only.
-     * Plan exit must not be blocked by hold enforcement. */
-    int plan_active = (ps[23] > 0.5f);
-    float enforced_position = plan_active
-        ? target_position  /* plan overrides hold — it manages its own duration */
-        : enforce_hold(ps0_f, target_position, hold_time, min_hold_bars, is_last_bar, isv_signals_ptr);
-    int hold_violation = (fabsf(enforced_position - target_position) > 0.001f);
-
-    if (hold_violation) {
-        target_position = enforced_position;  /* keep old position */
-
-        /* Fix the stored action to match held direction+magnitude (prevent action aliasing) */
-        int held_dir, held_mag;
-        if (ps0_f > 0.001f) {
-            held_dir = 2;  /* Long */
-        } else if (ps0_f < -0.001f) {
-            held_dir = 0;  /* Short */
-        } else {
-            held_dir = 1;  /* Flat */
-        }
-        /* Approximate magnitude from position fraction */
-        float abs_frac = fabsf(ps0_f) / fmaxf(max_position, 0.001f);
-        if (held_dir == 1) {
-            held_mag = 1;  /* Flat — magnitude irrelevant, default Half */
-        } else if (abs_frac <= 0.375f) {
-            held_mag = 0;  /* Quarter */
-        } else if (abs_frac <= 0.75f) {
-            held_mag = 1;  /* Half */
-        } else {
-            held_mag = 2;  /* Full */
-        }
-        int original_order = decode_order_4b(action_idx, b2_size, b3_size);
-        int original_urgency = decode_urgency_4b(action_idx, b3_size);
-        action_idx = held_dir * b1_size * b2_size * b3_size
-                   + held_mag * b2_size * b3_size
-                   + original_order * b3_size + original_urgency;
-        out_actions[out_off] = action_idx;
-    }
+    /* Hold enforcement removed — replaced by cost-driven hold timing.
+     * The model can always take any action; costs make bad actions expensive.
+     * See: inventory_penalty + churn_penalty in reward section below. */
+    int hold_violation = 0;
 
     /* ---- Execute trade (shared: trade_physics.cuh) ---- */
     int order_type_idx = decode_order_4b(action_idx, b2_size, b3_size);
@@ -1979,6 +1894,23 @@ extern "C" __global__ void experience_env_step(
         reward *= -1.0f;
     }
 
+    /* ── Cost-driven hold timing (replaces min_hold_bars) ──────────────
+     * The temporal attention learns optimal hold timing from these signals.
+     * No hardcoded gates — the model sees costs and learns to minimize them. */
+
+    /* Inventory penalty: continuous cost proportional to position size.
+     * Teaches: small positions cheap to hold, large positions expensive. */
+    float inventory_penalty = holding_cost_rate * fabsf(position);
+    reward -= inventory_penalty;
+
+    /* Churn penalty: graduated cost for rapid position flips.
+     * Not a hard gate — the model CAN exit early, but it costs extra.
+     * Zero penalty after churn_threshold bars. */
+    if (entering_trade && hold_time < churn_threshold && hold_time > 0.0f) {
+        float churn_penalty = churn_penalty_scale * (churn_threshold - hold_time) / churn_threshold;
+        reward -= churn_penalty;
+    }
+
     /* Plan conviction as reward scaling (always active, ISV-driven).
      * The plan head receives ISV-gated h_s2 (temporal attention → feature gate → plan MLP).
      * Its conviction output scales the reward → Q-value → C51 loss gradient path.
diff --git a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs
index 392d283e5..5d6f9a9e2 100644
--- a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs
@@ -187,9 +187,13 @@ pub struct GpuBacktestConfig {
     /// Used to compute the correct Sharpe annualization: `sqrt(bars_per_day * 252)`.
     /// Default: 390.0 (1-minute bar frequency).
     pub bars_per_day: f32,
-    /// Minimum bars to hold a position (matches training kernel's min_hold_bars).
-    /// Enforced in `backtest_env_step` to eliminate train/eval mismatch.
-    pub min_hold_bars: i32,
+    /// Holding cost rate per bar per unit position. Continuous inventory penalty.
+    /// Default: 0.00000052 (~0.52 bps / 100 bars)
+    pub holding_cost_rate: f32,
+    /// Bars threshold for churn penalty. Trades within this window get graduated penalty.
+    pub churn_threshold: f32,
+    /// Scale factor for churn penalty.
+    pub churn_penalty_scale: f32,
     /// Initial margin as fraction of notional value.
     /// CME initial margin ~6% of notional for equity index futures.
     /// Default: 0.06.
@@ -209,7 +213,9 @@ impl Default for GpuBacktestConfig {
             max_leverage: 2.0,
             ofi_dim: 0,
             bars_per_day: 390.0, // 1-minute bar frequency (6.5h × 60min)
-            min_hold_bars: 5,
+            holding_cost_rate: 0.00000052,
+            churn_threshold: 10.0,
+            churn_penalty_scale: 0.001,
             trading_days_per_year: 252.0,
             margin_pct: 0.06,
         }
@@ -964,7 +970,7 @@ impl GpuBacktestEvaluator {
                 shared_mem_bytes: 0,
             };
             let null_portfolio: u64 = 0;
-            let eval_min_hold: i32 = self.config.min_hold_bars;
+            let eval_min_hold: i32 = 0;  // hold enforcement removed — cost-driven
             let eval_max_pos: f32 = 0.0;
             unsafe {
                 self.stream
@@ -1028,7 +1034,9 @@ impl GpuBacktestEvaluator {
                     .arg(&b1)
                     .arg(&b2)
                     .arg(&b3)
-                    .arg(&self.config.min_hold_bars)
+                    .arg(&self.config.holding_cost_rate)
+                    .arg(&self.config.churn_threshold)
+                    .arg(&self.config.churn_penalty_scale)
                     .arg(&self.config.contract_multiplier)
                     .arg(&self.config.margin_pct)
                     .launch(LaunchConfig {
@@ -1353,7 +1361,9 @@ impl GpuBacktestEvaluator {
                 .arg(&b1_i32)
                 .arg(&b2_i32)
                 .arg(&self.b3_size)
-                .arg(&self.config.min_hold_bars)
+                .arg(&self.config.holding_cost_rate)
+                .arg(&self.config.churn_threshold)
+                .arg(&self.config.churn_penalty_scale)
                 .arg(&self.config.contract_multiplier)
                 .arg(&self.config.margin_pct)
                 .launch(env_cfg)
diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
index 30d0d0ce1..2e52b881d 100644
--- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
@@ -243,8 +243,12 @@ pub struct ExperienceCollectorConfig {
     pub n_steps: i32,
     /// Enable action masking: filters invalid exposure actions in the GPU kernel.
     pub enable_action_masking: bool,
-    /// Minimum bars to hold a position before exiting or reversing (Layer 2 hold enforcement).
-    pub min_hold_bars: i32,
+    /// Holding cost rate per bar per unit position. Continuous inventory penalty.
+    pub holding_cost_rate: f32,
+    /// Bars threshold for churn penalty. Trades within this window get graduated penalty.
+    pub churn_threshold: f32,
+    /// Scale factor for churn penalty.
+    pub churn_penalty_scale: f32,
     /// Bid-ask spread cost per unit position change. Matches backtest_env_kernel's spread_cost.
     pub spread_cost: f32,
     /// Contract multiplier (dollar value per point).
@@ -362,7 +366,9 @@ impl Default for ExperienceCollectorConfig {
             dsr_eta: 0.01,
             n_steps: 1,
             enable_action_masking: false,
-            min_hold_bars: 5,
+            holding_cost_rate: 0.00000052,  // ~0.52 bps / 100 bars
+            churn_threshold: 10.0,
+            churn_penalty_scale: 0.001,
             spread_cost: 0.0,  // default: no spread cost (overridden by hyperparams)
             contract_multiplier: 50.0,
             margin_pct: 0.06,
@@ -755,6 +761,7 @@ impl GpuExperienceCollector {
         let alloc_timesteps = timesteps_per_episode.min(MAX_TIMESTEPS_LIMIT);
         let (shared_h1, shared_h2, value_h, adv_h) = network_dims;
         let (state_dim, market_dim, num_atoms_max) = kernel_dims;
+        let _portfolio_dim: usize = 12;  // 8 base + 4 plan progress features
         let ofi_dim: usize = 20;
 
         // Branch sizes for 4-branch hierarchical DQN (always enabled).
@@ -2219,7 +2226,7 @@ impl GpuExperienceCollector {
             let _epsilon = config.epsilon;  // v8: epsilon now computed GPU-side via cosine schedule
             let q_gap = config.q_gap_threshold;
             let max_pos = config.max_position;
-            let min_hold_bars_i32 = config.min_hold_bars;
+            let min_hold_bars_i32: i32 = 0;  // hold enforcement removed — cost-driven
             unsafe {
                 self.stream
                     .launch_builder(&self.action_select_kernel)
@@ -2236,8 +2243,8 @@ impl GpuExperienceCollector {
                     .arg(&b2)
                     .arg(&b3)
                     .arg(&q_gap)
-                    .arg(&self.portfolio_states)       // portfolio state for hold masking
-                    .arg(&min_hold_bars_i32)            // min_hold_bars
+                    .arg(&self.portfolio_states)       // ABI compat — unused by action_select
+                    .arg(&min_hold_bars_i32)            // ABI compat — unused (pass 0)
                     .arg(&max_pos)                      // max_position for direction index
                     .arg(&config.eps_exp_mult)          // per-branch epsilon multiplier: exposure
                     .arg(&config.eps_ord_mult)          // per-branch epsilon multiplier: order
@@ -2282,7 +2289,7 @@ impl GpuExperienceCollector {
             }
 
             // ── 5. Environment step (reward v5: trade-aware hybrid) ──────
-            // max_pos and min_hold_bars_i32 already defined above (action_select block)
+            // max_pos already defined above (action_select block)
             let tx_cost = config.tx_cost_multiplier;
             let l_i32 = timesteps as i32;
             let total_actions_env = (self.branch_sizes[0] + self.branch_sizes[1]
@@ -2318,7 +2325,9 @@ impl GpuExperienceCollector {
                     .arg(&self.cvar_scales_ptr) // CVaR position scaling (0 = NULL = no scaling)
                     .arg(&self.q_gaps_buf)     // Q-gap conviction scaling
                     .arg(&mut self.raw_returns_out) // Raw portfolio returns (unshaped) for Sharpe/MaxDD
-                    .arg(&min_hold_bars_i32)          // min_hold_bars for hold enforcement
+                    .arg(&config.holding_cost_rate)    // continuous inventory penalty
+                    .arg(&config.churn_threshold)      // churn penalty threshold (bars)
+                    .arg(&config.churn_penalty_scale)  // churn penalty scale
                     .arg(&config.spread_cost)         // bid-ask spread cost (matches backtest)
                     .arg(&config.contract_multiplier) // futures contract multiplier (e.g. 50 for ES)
                     .arg(&config.margin_pct)          // initial margin fraction (e.g. 0.06 = 6%)
diff --git a/crates/ml/src/cuda_pipeline/iqn_dual_head_kernel.cu b/crates/ml/src/cuda_pipeline/iqn_dual_head_kernel.cu
index c7e4926ec..5edc8d808 100644
--- a/crates/ml/src/cuda_pipeline/iqn_dual_head_kernel.cu
+++ b/crates/ml/src/cuda_pipeline/iqn_dual_head_kernel.cu
@@ -1217,24 +1217,11 @@ void iqn_sample_taus_kernel(
     taus[i] = 0.01f + u * 0.98f;
 }
 
-/* ============================================================
- * cuBLAS element-wise helper kernels (Task 4)
- * ============================================================ */
-
-extern "C" __global__
-void iqn_hadamard_sigmoid(
-    const float* __restrict__ h_s2_tiled,
-    const float* __restrict__ embed,
-    float* __restrict__ out,
-    int N
-) {
-    int idx = blockIdx.x * blockDim.x + threadIdx.x;
-    if (idx < N) {
-        float sig = 1.0f / (1.0f + expf(-embed[idx]));
-        out[idx] = h_s2_tiled[idx] * sig;
-    }
-}
+/* ================================================================== */
+/* Element-wise kernels for cuBLAS-based IQN forward/backward         */
+/* ================================================================== */
 
+/* ReLU forward: out = max(x, 0) */
 extern "C" __global__
 void iqn_relu_fwd(
     const float* __restrict__ x,
@@ -1247,11 +1234,12 @@ void iqn_relu_fwd(
     }
 }
 
+/* ReLU backward: d_x = (x > 0) ? d_out : 0 */
 extern "C" __global__
 void iqn_relu_bwd(
-    const float* __restrict__ x,
-    const float* __restrict__ d_out,
-    float* __restrict__ d_x,
+    const float* __restrict__ x,       /* saved pre-activation */
+    const float* __restrict__ d_out,   /* upstream gradient */
+    float* __restrict__ d_x,           /* output gradient */
     int N
 ) {
     int idx = blockIdx.x * blockDim.x + threadIdx.x;
@@ -1260,72 +1248,32 @@ void iqn_relu_bwd(
     }
 }
 
-/* Quantile Huber loss for per-branch-block col-major layout.
-   q_online / q_target / d_q_online: 4 contiguous blocks [bk, B*Q] col-major.
-   Branch d starts at flat offset sum(branch_sizes[0..d]) * B * Q.
-   Within a block: element(row=a, col=c) = block_base + a + c * bk.
-   One block per sample. Threads cooperate over quantiles. */
+/* Hadamard product with sigmoid: out = h_s2_tiled * sigmoid(embed) */
 extern "C" __global__
-void iqn_quantile_huber_loss(
-    const float* __restrict__ q_online,
-    const float* __restrict__ q_target,
-    const float* __restrict__ taus,
-    const int* __restrict__ actions,
-    float* __restrict__ per_sample_loss,
-    float* __restrict__ d_q_online,
-    int B, int Q, int b0, int b1, int b2, int b3,
-    float kappa, float gamma
+void iqn_hadamard_sigmoid(
+    const float* __restrict__ h_s2_tiled,  /* [H, B*Q] */
+    const float* __restrict__ embed,       /* [H, B*Q] post-ReLU embedding */
+    float* __restrict__ out,               /* [H, B*Q] combined */
+    int N                                  /* H * B * Q */
 ) {
-    int sample = blockIdx.x;
-    if (sample >= B) return;
-    int tid = threadIdx.x;
-
-    int branch_sizes[4] = { b0, b1, b2, b3 };
-    int BQ = B * Q;
-    int block_base = 0;
-    float total_loss = 0.0f;
-
-    for (int d = 0; d < 4; d++) {
-        int a_d = actions[sample * 4 + d];
-        int nd = branch_sizes[d];
-        if (a_d < 0 || a_d >= nd) a_d = 0;
-
-        for (int q = tid; q < Q; q += blockDim.x) {
-            int col = sample * Q + q;
-            /* Col-major index within this branch block: a_d + col * nd */
-            int idx = block_base + a_d + col * nd;
-            float qo = q_online[idx];
-            float qt = gamma * q_target[idx];
-            float tau = taus[col];
-            float delta = qt - qo;
-            float abs_delta = fabsf(delta);
-            float huber = (abs_delta <= kappa)
-                ? 0.5f * delta * delta
-                : kappa * (abs_delta - 0.5f * kappa);
-            float weight = (delta >= 0.0f) ? tau : (1.0f - tau);
-            total_loss += weight * huber / (float)Q;
-
-            float d_huber = (abs_delta <= kappa)
-                ? -delta
-                : -kappa * ((delta > 0.0f) ? 1.0f : -1.0f);
-            d_q_online[idx] = weight * d_huber / (float)(BQ);
-        }
-        block_base += nd * BQ;
+    int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    if (idx < N) {
+        float sig = 1.0f / (1.0f + expf(-embed[idx]));
+        out[idx] = h_s2_tiled[idx] * sig;
     }
-
-    /* Warp reduce total_loss */
-    for (int offset = 16; offset > 0; offset >>= 1)
-        total_loss += __shfl_down_sync(0xFFFFFFFF, total_loss, offset);
-    if (tid == 0) per_sample_loss[sample] = total_loss;
 }
 
+/* Hadamard-sigmoid backward:
+   d_embed = d_combined * h_s2_tiled * sigmoid'(embed_pre)
+   d_h_s2_tiled = d_combined * sigmoid(embed_post)
+   embed_post is post-ReLU, so sigmoid acts on the ReLU output. */
 extern "C" __global__
 void iqn_hadamard_sigmoid_bwd(
-    const float* __restrict__ h_s2_tiled,
-    const float* __restrict__ embed,
-    const float* __restrict__ d_combined,
-    float* __restrict__ d_embed,
-    float* __restrict__ d_h_s2_tiled,
+    const float* __restrict__ h_s2_tiled,  /* [H, B*Q] */
+    const float* __restrict__ embed,       /* [H, B*Q] post-ReLU embedding */
+    const float* __restrict__ d_combined,  /* [H, B*Q] upstream gradient */
+    float* __restrict__ d_embed,           /* [H, B*Q] gradient w.r.t. embed */
+    float* __restrict__ d_h_s2_tiled,      /* [H, B*Q] gradient w.r.t. h_s2 */
     int N
 ) {
     int idx = blockIdx.x * blockDim.x + threadIdx.x;
@@ -1337,6 +1285,7 @@ void iqn_hadamard_sigmoid_bwd(
     }
 }
 
+/* Bias add: x[idx] += bias[idx % out_dim] */
 extern "C" __global__
 void iqn_bias_add(
     float* __restrict__ x,
@@ -1350,6 +1299,8 @@ void iqn_bias_add(
     }
 }
 
+/* Bias gradient reduce: d_bias[j] = inv_bq * sum_i d_pre[j + i*out_dim]
+   Column-major: d_pre is [out_dim, BQ], sum along BQ axis. */
 extern "C" __global__
 void iqn_bias_grad_reduce(
     const float* __restrict__ d_pre,
@@ -1407,6 +1358,114 @@ void iqn_d_h_s2_reduce(
     }
 }
 
+/* Quantile Huber loss + per-branch dq gradients.
+   One block per sample. Threads cooperate over quantiles/actions.
+   Layout: q_online/q_target are [TBA, B*Q] col-major.
+   actions: [B, 4] row-major (branch actions per sample).
+
+   For each sample s and quantile q, computes:
+     For each branch b, the taken action a_b:
+       delta = q_online[offset_b + a_b, s*Q + q] - gamma * q_target[offset_b + a_b, s*Q + q]
+       Quantile Huber loss with tau = taus[s*Q + q]
+       d_q_online = (1/Q) * d_huber * |tau - I(delta<0)|
+
+   per_sample_loss[s] = (1/Q) * sum over q of sum over branches of huber_loss */
+extern "C" __global__
+void iqn_quantile_huber_loss(
+    const float* __restrict__ q_online,    /* [TBA, B*Q] */
+    const float* __restrict__ q_target,    /* [TBA, B*Q] */
+    const float* __restrict__ taus,        /* [B*Q] (fixed midpoints) */
+    const int* __restrict__ actions,       /* [B, 4] */
+    float* __restrict__ per_sample_loss,   /* [B] */
+    float* __restrict__ d_q_online,        /* [TBA, B*Q] zero-init required */
+    int B, int Q, int b0, int b1, int b2, int b3,
+    float kappa, float gamma
+) {
+    int s = blockIdx.x;
+    if (s >= B) return;
+    int tid = threadIdx.x;
+    int TBA = b0 + b1 + b2 + b3;
+
+    /* Decode branch actions for this sample */
+    int a0 = actions[s * 4 + 0];
+    int a1 = actions[s * 4 + 1];
+    int a2 = actions[s * 4 + 2];
+    int a3 = actions[s * 4 + 3];
+
+    /* Branch offsets in the TBA dimension */
+    int off0 = 0, off1 = b0, off2 = b0 + b1, off3 = b0 + b1 + b2;
+    int branch_actions[4] = {a0, a1, a2, a3};
+    int branch_offsets[4] = {off0, off1, off2, off3};
+
+    float sample_loss = 0.0f;
+
+    /* Each thread processes a subset of quantiles */
+    for (int q = tid; q < Q; q += blockDim.x) {
+        int col = s * Q + q;
+        float tau_q = taus[col];
+        float q_loss = 0.0f;
+
+        for (int br = 0; br < 4; br++) {
+            int row = branch_offsets[br] + branch_actions[br];
+            int idx = row + col * TBA;  /* col-major index */
+
+            float online_val = q_online[idx];
+            float target_val = q_target[idx];
+            float delta = online_val - gamma * target_val;
+
+            /* Huber loss */
+            float abs_delta = fabsf(delta);
+            float huber;
+            if (abs_delta <= kappa) {
+                huber = 0.5f * delta * delta;
+            } else {
+                huber = kappa * (abs_delta - 0.5f * kappa);
+            }
+
+            /* Quantile weight */
+            float indicator = (delta < 0.0f) ? 1.0f : 0.0f;
+            float qw = fabsf(tau_q - indicator);
+            q_loss += qw * huber;
+
+            /* Gradient of quantile Huber w.r.t. q_online */
+            float d_huber;
+            if (abs_delta <= kappa) {
+                d_huber = delta;
+            } else {
+                d_huber = (delta > 0.0f) ? kappa : -kappa;
+            }
+            d_q_online[idx] = qw * d_huber / (float)Q;
+        }
+        sample_loss += q_loss;
+    }
+
+    /* Warp reduce then block reduce for sample_loss */
+    sample_loss /= (float)Q;
+
+    /* Warp-level reduction */
+    for (int offset = 16; offset > 0; offset >>= 1) {
+        sample_loss += __shfl_down_sync(0xFFFFFFFF, sample_loss, offset);
+    }
+
+    /* Block-level reduction via shared memory */
+    __shared__ float warp_sums[8]; /* max 256 threads = 8 warps */
+    int warp_id = tid / 32;
+    int lane_id = tid % 32;
+    if (lane_id == 0) {
+        warp_sums[warp_id] = sample_loss;
+    }
+    __syncthreads();
+
+    if (tid == 0) {
+        float total = 0.0f;
+        int num_warps = (blockDim.x + 31) / 32;
+        for (int w = 0; w < num_warps; w++) {
+            total += warp_sums[w];
+        }
+        per_sample_loss[s] = total;
+    }
+}
+
 /* Tile cos_features [D, Q] into [D, B*Q] by repeating B times.
    cos_features_src[d + q*D] -> cos_features_tiled[d + (b*Q+q)*D] */
 extern "C" __global__
diff --git a/crates/ml/src/cuda_pipeline/trade_physics.cuh b/crates/ml/src/cuda_pipeline/trade_physics.cuh
index eb307da95..9bfe61cfb 100644
--- a/crates/ml/src/cuda_pipeline/trade_physics.cuh
+++ b/crates/ml/src/cuda_pipeline/trade_physics.cuh
@@ -4,7 +4,6 @@
  * This header defines the SINGLE SOURCE OF TRUTH for:
  * - Factored action decoding (exposure index extraction)
  * - Exposure-to-position mapping (index → fraction → contracts)
- * - Hold enforcement (block exits/reversals during min_hold_bars)
  * - Transaction cost computation (Almgren-Chriss impact model)
  * - Capital floor circuit breaker (75% of peak equity)
  * - Hold time tracking (bars in current position)
@@ -133,46 +132,6 @@ __device__ __forceinline__ int decode_order_type(
     return (b1_size > 0 && b2_size > 0) ? (action_val / b2_size) % b1_size : 0;
 }
 
-/* ── Hold enforcement: block exits and reversals ─────────────────────── */
-
-__device__ __forceinline__ float enforce_hold(
-    float prev_position,
-    float target_exposure,
-    float hold_time,
-    int   min_hold_bars,
-    int   is_last_bar,  /* bypass hold on last bar of episode/window */
-    const float* __restrict__ isv_signals  /* [8] or NULL — for adaptive hold */
-) {
-    int prev_sign = (prev_position > 0.001f) ? 1 : ((prev_position < -0.001f) ? -1 : 0);
-    int curr_sign = (target_exposure > 0.001f) ? 1 : ((target_exposure < -0.001f) ? -1 : 0);
-    int wants_exit     = (prev_sign != 0 && curr_sign == 0);
-    int wants_reversal = (prev_sign != 0 && curr_sign != 0 && prev_sign != curr_sign);
-
-    /* Adaptive hold: base + ISV-driven extension.
-     * ISV[1] = grad_norm_ema (high = unstable → hold longer, don't churn)
-     * ISV[3] = ensemble_var_ema (high = uncertain → hold longer, wait for clarity)
-     * ISV[5] = reward_ema (negative = losing → allow earlier exit)
-     * When ISV is NULL, use static min_hold_bars. */
-    int adaptive_hold = min_hold_bars;
-    if (isv_signals != NULL) {
-        float stability = 1.0f / (1.0f + isv_signals[1]);  /* inv grad_norm: high norm → low stability */
-        float confidence = 1.0f - fminf(isv_signals[3], 1.0f);  /* inv ensemble_var: high var → low confidence */
-        /* Extension: 0-2× base hold based on stability × confidence.
-         * Scales with min_hold_bars: base=50 → extension up to 100 bars. */
-        int extension = (int)(stability * confidence * (float)min_hold_bars * 2.0f);
-        /* If losing (negative reward_ema), reduce hold to base only */
-        if (isv_signals[5] < -0.001f) extension = 0;
-        adaptive_hold = min_hold_bars + extension;
-    }
-
-    if (hold_time > 0.0f && hold_time < (float)adaptive_hold
-        && (wants_exit || wants_reversal)
-        && !is_last_bar) {
-        return prev_position;  /* Override: keep position */
-    }
-    return target_exposure;
-}
-
 /* ── Transaction cost (Almgren-Chriss impact model) ──────────────────── */
 
 __device__ __forceinline__ float compute_tx_cost(
@@ -304,7 +263,7 @@ __device__ __forceinline__ float update_hold_time(
 
 /* ── Trailing stop: exit when profit retreats from peak ──────────────── */
 /* Only triggers when:
- *   - Position held >= min_hold_bars (don't trail during hold period)
+ *   - Position held >= 1 bar (must have at least one bar of data)
  *   - Peak equity > 1.0 (valid equity tracking)
  *   - Peak return > trail_distance (profit must exist before trailing)
  *   - Unrealized return < trail_floor (profit retreated beyond threshold)
@@ -315,7 +274,6 @@ __device__ __forceinline__ float update_hold_time(
 
 __device__ __forceinline__ int check_trailing_stop(
     float hold_time,
-    int   min_hold_bars,
     float peak_equity,
     float prev_equity,
     float current_trade_return,
@@ -323,7 +281,7 @@ __device__ __forceinline__ int check_trailing_stop(
     float vol_scale,            /* CUSUM-based volatility scaling (1.0 = no scale) */
     float trend_scale           /* ADX-based trend scaling (1.0 = no scale) */
 ) {
-    if (hold_time < (float)min_hold_bars || peak_equity <= 1.0f) return 0;
+    if (hold_time < 1.0f || peak_equity <= 1.0f) return 0;
     float trail_distance = base_trail_distance * vol_scale * trend_scale;
     float peak_return = (peak_equity - prev_equity) / fmaxf(prev_equity, 1.0f);
     if (peak_return > trail_distance) {
diff --git a/crates/ml/src/hyperopt/adapters/dqn.rs b/crates/ml/src/hyperopt/adapters/dqn.rs
index cdf9321d0..a63763f0b 100644
--- a/crates/ml/src/hyperopt/adapters/dqn.rs
+++ b/crates/ml/src/hyperopt/adapters/dqn.rs
@@ -160,7 +160,7 @@ const MB_PER_SAMPLE: f64 = 0.0005;
 ///   loss_shaping_intensity, ensemble_intensity, causal_intensity (6 gen families)
 ///
 /// ## Fixed at TOML defaults (removed from PSO):
-/// batch_size, transaction_cost_multiplier, v_max, min_hold_bars, and all
+/// batch_size, transaction_cost_multiplier, v_max, holding_cost_rate, and all
 /// individual params now absorbed into the 5 core families.
 #[derive(Debug, Clone, PartialEq, Serialize, Deserialize)]
 #[serde(default)]
@@ -454,7 +454,7 @@ pub struct DQNMetrics {
 /// ## Fixed Architecture
 ///
 /// The following parameters are fixed for consistency:
-/// - `state_dim`: 72 without OFI (42 market + 14 portfolio + 16 MTF), or 92 with MBP-10 data (+20 OFI)
+/// - `state_dim`: 66 without OFI (42 market + 8 portfolio + 16 MTF), or 86 with MBP-10 data (+20 OFI)
 ///   - Market features (42): OHLCV, technical indicators, price patterns, volume, time, statistical, ADX, CUSUM direction
 ///   - OFI features (20, optional): 8 base OFI from MBP-10 order book + 12 extended microstructure
 ///   - Portfolio features (3): Position, unrealized PnL, drawdown
@@ -1435,7 +1435,9 @@ impl DQNTrainer {
             // OFI reorder in gather kernel: produces [market, portfolio, OFI, pad]
             // directly, eliminating the Candle narrow+cat closure.
             ofi_dim: 20,
-            min_hold_bars: internal_trainer.hyperparams().min_hold_bars as i32,
+            holding_cost_rate: internal_trainer.hyperparams().holding_cost_rate as f32,
+            churn_threshold: internal_trainer.hyperparams().churn_threshold_bars as f32,
+            churn_penalty_scale: internal_trainer.hyperparams().churn_penalty_scale as f32,
             bars_per_day: internal_trainer.hyperparams().bars_per_day as f32,
             trading_days_per_year: internal_trainer.hyperparams().trading_days_per_year as f32,
             ..Default::default()
diff --git a/crates/ml/src/trainers/dqn/config.rs b/crates/ml/src/trainers/dqn/config.rs
index 6f6f6fdfb..fdddd470d 100644
--- a/crates/ml/src/trainers/dqn/config.rs
+++ b/crates/ml/src/trainers/dqn/config.rs
@@ -281,7 +281,7 @@ impl DQNAgentType {
     ///
     /// WAVE 10.4: Added to fix hardcoded STATE_DIM bug
     /// Returns the actual state dimension (72 without OFI, 80 with OFI)
-    /// Layout: 42 market + 14 portfolio + 16 multi-timeframe + 20 OFI, 8-aligned
+    /// Layout: 42 market + 8 portfolio + 16 multi-timeframe + (8 OFI), 8-aligned
     pub fn get_state_dim(&self) -> usize {
         self.agent.get_state_dim()
     }
@@ -560,10 +560,18 @@ pub struct DQNHyperparameters {
     /// Derived from data pipeline's BarSize when available.
     pub bars_per_day: f64,
 
-    /// Minimum bars to hold a position before allowing exit or reversal.
-    /// Prevents per-bar churning. Trailing stop can override after this period.
-    /// Default: 5 (about 5 minutes for 1-min bars). Range: [2, 20].
-    pub min_hold_bars: usize,
+    /// Holding cost rate per bar per unit position. Continuous inventory penalty.
+    /// Default: 0.00000052 (~0.52 bps / 100 bars)
+    pub holding_cost_rate: f64,
+    /// Bars threshold for churn penalty. Trades within this window get graduated penalty.
+    pub churn_threshold_bars: usize,
+    /// Scale factor for churn penalty.
+    pub churn_penalty_scale: f64,
+
+    /// Auxiliary ops (IQL, IQN, attention, CQL) run every Nth training step.
+    /// Default: 4 (aux ops every 4th step, ~4× faster epochs).
+    /// 1 = every step (full quality), higher = faster but less frequent aux updates.
+    pub aux_frequency: usize,
 
     // P2-B Enhancement: Cash reserve requirement
     /// Cash reserve requirement as percentage of portfolio value (0-100)
@@ -1342,8 +1350,12 @@ impl DQNHyperparameters {
 
             // P2-A Enhancement
             initial_capital: 100_000.0, // $100K default
-            bars_per_day: 0.0, // must be set from data before training (fails if 0)
-            min_hold_bars: 5,  // 5-bar minimum prevents coin-flip exits (1-2 bar churn)
+            bars_per_day: 390.0, // 1-minute bars (6.5h × 60min)
+            holding_cost_rate: 0.00000052,  // ~0.52 bps / 100 bars inventory penalty
+            churn_threshold_bars: 10,      // graduated penalty for trades within 10 bars
+            churn_penalty_scale: 0.001,    // churn penalty scale factor
+            aux_frequency: 4,  // aux ops every 4th step (~4× faster epochs)
+
             // P2-B Enhancement
             cash_reserve_percent: 0.0, // Default: no reserve (backward compatible)
 
diff --git a/crates/ml/src/trainers/dqn/trainer/metrics.rs b/crates/ml/src/trainers/dqn/trainer/metrics.rs
index 9b3fedadf..cf625ae9c 100644
--- a/crates/ml/src/trainers/dqn/trainer/metrics.rs
+++ b/crates/ml/src/trainers/dqn/trainer/metrics.rs
@@ -443,8 +443,7 @@ impl DQNTrainer {
         // ── Lazy-init the GPU evaluator (once per fold, reused across epochs) ──
         if self.gpu_evaluator.is_none() {
             let market_dim: usize = 42;
-            let ofi_dim: usize = crate::fxcache::OFI_DIM;
-            let feature_dim: usize = market_dim + ofi_dim;
+            let feature_dim: usize = 62; // 42 market + 20 OFI (portfolio+MTF in state_dim, not feature_dim)
 
             // Build a single window from all val_data
             let mut prices: Vec<[f32; 4]> = Vec::with_capacity(self.val_data.len());
@@ -485,8 +484,10 @@ impl DQNTrainer {
                 contract_multiplier: hp.contract_multiplier as f32,
                 margin_pct: hp.margin_pct as f32,
                 max_leverage: hp.max_leverage as f32,
-                ofi_dim,
-                min_hold_bars: hp.min_hold_bars as i32,
+                ofi_dim: 20,
+                holding_cost_rate: hp.holding_cost_rate as f32,
+                churn_threshold: hp.churn_threshold_bars as f32,
+                churn_penalty_scale: hp.churn_penalty_scale as f32,
                 bars_per_day: hp.bars_per_day as f32,
                 trading_days_per_year: hp.trading_days_per_year as f32,
             };
diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
index 1079f8893..4e14afe79 100644
--- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs
+++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
@@ -67,14 +67,6 @@ impl DQNTrainer {
     where
         F: FnMut(usize, Vec<u8>, bool) -> Result<String> + Send,
     {
-        // bars_per_day must be set from data before training starts (0 means uninitialized).
-        // The caller (train_baseline_rl or equivalent) must set hyperparams.bars_per_day
-        // from the actual dataset's timestamp density before calling this function.
-        anyhow::ensure!(
-            self.hyperparams.bars_per_day > 0.0,
-            "bars_per_day is 0 — must be set from data before training (e.g. via timestamps)"
-        );
-
         let start_time = std::time::Instant::now();
         let mut total_loss = 0.0;
         let mut total_q_value = 0.0;
@@ -1252,7 +1244,9 @@ impl DQNTrainer {
             },
             dsr_eta: self.hyperparams.dsr_eta as f32,
             n_steps: self.hyperparams.n_steps as i32,
-            min_hold_bars: self.hyperparams.min_hold_bars as i32,
+            holding_cost_rate: self.hyperparams.holding_cost_rate as f32,
+            churn_threshold: self.hyperparams.churn_threshold_bars as f32,
+            churn_penalty_scale: self.hyperparams.churn_penalty_scale as f32,
             spread_cost: {
                 // spread_cost in POINTS (no contract_multiplier — PnL is also in points).
                 // ES: 0.25 tick × 0.50 fill_frac = 0.125 points = half a tick spread.
diff --git a/crates/ml/src/training_profile.rs b/crates/ml/src/training_profile.rs
index a1a84c4fd..944a764b5 100644
--- a/crates/ml/src/training_profile.rs
+++ b/crates/ml/src/training_profile.rs
@@ -289,8 +289,12 @@ pub struct ExperienceSection {
     pub initial_capital: Option<f64>,
     /// Maps to `DQNHyperparameters::transaction_cost_multiplier`.
     pub tx_cost_multiplier: Option<f64>,
-    /// Minimum bars to hold a position before allowing exit or reversal.
-    pub min_hold_bars: Option<usize>,
+    /// Holding cost rate per bar per unit position. Continuous inventory penalty.
+    pub holding_cost_rate: Option<f64>,
+    /// Bars threshold for churn penalty.
+    pub churn_threshold_bars: Option<usize>,
+    /// Scale factor for churn penalty.
+    pub churn_penalty_scale: Option<f64>,
     /// Fill simulation sub-section for order routing realism.
     pub fill_simulation: Option<FillSimulationSection>,
 }
@@ -997,8 +1001,14 @@ impl DqnTrainingProfile {
             if let Some(v) = ex.tx_cost_multiplier {
                 hp.transaction_cost_multiplier = v;
             }
-            if let Some(v) = ex.min_hold_bars {
-                hp.min_hold_bars = v;
+            if let Some(v) = ex.holding_cost_rate {
+                hp.holding_cost_rate = v;
+            }
+            if let Some(v) = ex.churn_threshold_bars {
+                hp.churn_threshold_bars = v;
+            }
+            if let Some(v) = ex.churn_penalty_scale {
+                hp.churn_penalty_scale = v;
             }
             // [experience.fill_simulation]
             if let Some(ref fs) = ex.fill_simulation {