diff --git a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu
index 734b1264a..98b07a900 100644
--- a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu
+++ b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu
@@ -19,7 +19,7 @@ extern "C" __global__ void backtest_env_step(
     const int*   __restrict__ window_lens,  // [n_windows]
 
     // Actions from model for current step
-    const int*   __restrict__ actions,      // [n_windows] (0-4: Short100..Long100)
+    const int*   __restrict__ actions,      // [n_windows] factored action index
 
     // Portfolio state (read-write, persistent across steps)
     float*       portfolio_state,           // [n_windows * PORTFOLIO_STATE_SIZE]
@@ -36,7 +36,11 @@ extern "C" __global__ void backtest_env_step(
     float        max_position,
     float        tx_cost_bps,
     float        spread_cost,
-    int          current_step
+    int          current_step,
+    // Branching DQN action space sizes
+    int          b0_size,                   // exposure branch size (default 9)
+    int          b1_size,                   // order type branch size (default 3)
+    int          b2_size                    // urgency branch size (default 3)
 ) {
     // Shared memory tile for coalesced portfolio reads/writes
     __shared__ float shmem_pf[256 * PORTFOLIO_STATE_SIZE];
@@ -78,11 +82,29 @@ extern "C" __global__ void backtest_env_step(
     float max_equity = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 4];
     float cum_return = shmem_pf[local_tid * PORTFOLIO_STATE_SIZE + 6];
 
-    // Map action (0-8) to target exposure: -1.0 + idx * 0.25
+    // Decode factored action to exposure index (matches experience_env_step).
+    // Factored action = exposure_idx * b1_size * b2_size + order_idx * b2_size + urgency_idx
     int action_val = actions[w];
-    float target_exposure = -1.0f + (float)action_val * 0.25f;
+    int exposure_idx;
+    if (b1_size > 0 && b2_size > 0) {
+        int denom = b1_size * b2_size;
+        exposure_idx = action_val / denom;
+    } else {
+        exposure_idx = action_val;
+    }
+    // Clamp to valid range
+    if (exposure_idx < 0) exposure_idx = 0;
+    if (exposure_idx >= b0_size) exposure_idx = b0_size - 1;
+
+    // Map exposure index to fraction: -1.0 + idx * (2.0 / (b0_size - 1))
+    // With b0_size=9: step=0.25, range [-1.0, +1.0]
+    float step_size = (b0_size > 1) ? 2.0f / (float)(b0_size - 1) : 0.0f;
+    float target_exposure = -1.0f + (float)exposure_idx * step_size;
     target_exposure *= max_position;
 
+    // Suppress unused variable warnings
+    (void)open; (void)high; (void)low;
+
     // Execute trade if position changes
     float delta = target_exposure - position;
     float trade_cost = 0.0f;
diff --git a/crates/ml/src/cuda_pipeline/backtest_gather_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_gather_kernel.cu
index 345cd641f..99b36f094 100644
--- a/crates/ml/src/cuda_pipeline/backtest_gather_kernel.cu
+++ b/crates/ml/src/cuda_pipeline/backtest_gather_kernel.cu
@@ -1,9 +1,11 @@
 // Gather state vectors from pre-uploaded features + live portfolio state.
 // Output: [n_windows, state_dim] tensor for model forward pass.
 //
-// This kernel eliminates the CPU roundtrip in the old gather_states() path,
-// which previously downloaded the full features buffer (n_windows * max_len * feat_dim
-// floats) to CPU just to slice out a single step's row per window.
+// State layout MATCHES the training experience_state_gather kernel:
+//   [0 .. market_dim)                      : market features (42)
+//   [market_dim .. market_dim+8)           : 8 portfolio features
+//   [market_dim+8 .. market_dim+8+16)      : 16 multi-timeframe features
+//   [market_dim+24 .. state_dim)           : zero-pad / OFI / alignment
 //
 // Portfolio state layout per window [8 floats]:
 //   [0] value        - current portfolio value
@@ -15,17 +17,19 @@
 //   [6] cum_return   - cumulative log return
 //   [7] step_count   - number of completed steps
 //
-// State output layout per window [state_dim floats]:
-//   When ofi_dim == 0 (no OFI):
-//     [0 .. feat_dim)              - market features
-//     [feat_dim + 0..2]            - portfolio (value, position, spread)
-//     [feat_dim + 3 .. state_dim)  - zero pad
+// 8 portfolio features (matches training experience_state_gather):
+//   +0: position                       — raw position [-1, +1]
+//   +1: unrealized_pnl / equity        — how is THIS trade doing?
+//   +2: drawdown                       — (peak - equity) / peak [0, 1]
+//   +3: hold_time / 100.0              — normalized holding duration (step_count as proxy)
+//   +4: realized_pnl / equity          — cumulative return signal
+//   +5: distance_to_floor              — (equity - floor) / equity [0, 1]
+//   +6: trade_return                   — unrealized PnL / equity
+//   +7: cash / equity                  — available capital ratio
 //
-//   When ofi_dim > 0 (OFI reorder):
-//     [0 .. market_dim)                         - market features (feat_dim - ofi_dim)
-//     [market_dim + 0..2]                       - portfolio (value, position, spread)
-//     [market_dim + 3 .. market_dim + 3 + ofi_dim) - OFI features
-//     [market_dim + 3 + ofi_dim .. state_dim)   - zero pad
+// 16 multi-timeframe features (4 lookbacks × 4 features):
+//   Lookback windows: 5, 15, 60, 240 bars
+//   Per window: return, volatility, volume_trend, momentum
 
 extern "C" __global__ void gather_states(
     const float* __restrict__ features,      // [n_windows * max_len * feat_dim]
@@ -61,84 +65,165 @@ extern "C" __global__ void gather_states(
     int feat_base = (w * max_len + current_step) * feat_dim;
     int out_base  = w * state_dim;
 
-    // Portfolio features (shared between both paths) — read from shared memory tile
-    float norm_value = (initial_capital > 0.0f)
-        ? shmem_portfolio[local_tid * 8 + 0] / initial_capital
-        : 0.0f;
-    float position = shmem_portfolio[local_tid * 8 + 1];
+    // Read portfolio state from shared memory tile
+    float value      = shmem_portfolio[local_tid * 8 + 0];
+    float position   = shmem_portfolio[local_tid * 8 + 1];
+    float cash       = shmem_portfolio[local_tid * 8 + 2];
+    float entry_price = shmem_portfolio[local_tid * 8 + 3];
+    float max_equity = shmem_portfolio[local_tid * 8 + 4];
+    float step_count = shmem_portfolio[local_tid * 8 + 7];
 
-    if (ofi_dim > 0) {
-        // OFI reorder: storage is [market, OFI], model expects [market, portfolio, OFI, pad]
-        int market_dim = feat_dim - ofi_dim;
+    // Determine market_dim: if OFI is present, market features are feat_dim - ofi_dim
+    int market_dim = (ofi_dim > 0) ? (feat_dim - ofi_dim) : feat_dim;
 
-        // 1. Market features [0 .. market_dim) — __ldg + manual 4x unroll
-        int i = 0;
-        for (; i + 3 < market_dim; i += 4) {
-            states_out[out_base + i]     = __ldg(&features[feat_base + i]);
-            states_out[out_base + i + 1] = __ldg(&features[feat_base + i + 1]);
-            states_out[out_base + i + 2] = __ldg(&features[feat_base + i + 2]);
-            states_out[out_base + i + 3] = __ldg(&features[feat_base + i + 3]);
-        }
-        for (; i < market_dim; i++) {
-            states_out[out_base + i] = __ldg(&features[feat_base + i]);
-        }
+    // ── 1. Market features [0 .. market_dim) ── //
+    int i = 0;
+    for (; i + 3 < market_dim; i += 4) {
+        states_out[out_base + i]     = __ldg(&features[feat_base + i]);
+        states_out[out_base + i + 1] = __ldg(&features[feat_base + i + 1]);
+        states_out[out_base + i + 2] = __ldg(&features[feat_base + i + 2]);
+        states_out[out_base + i + 3] = __ldg(&features[feat_base + i + 3]);
+    }
+    for (; i < market_dim; i++) {
+        states_out[out_base + i] = __ldg(&features[feat_base + i]);
+    }
 
-        // 2. Portfolio [market_dim .. market_dim + 3)
-        states_out[out_base + market_dim + 0] = norm_value;
-        states_out[out_base + market_dim + 1] = position;
-        states_out[out_base + market_dim + 2] = spread_cost;
+    // ── 2. Portfolio features [market_dim .. market_dim+8) ── //
+    // Match training experience_state_gather exactly.
+    float equity = (value > 1.0f) ? value : 1.0f;
 
-        // 3. OFI features [market_dim + 3 .. market_dim + 3 + ofi_dim)
-        //    ofi_dim=8 is always float4-aligned, full 4x unroll
-        for (i = 0; i + 3 < ofi_dim; i += 4) {
-            states_out[out_base + market_dim + 3 + i]     = __ldg(&features[feat_base + market_dim + i]);
-            states_out[out_base + market_dim + 3 + i + 1] = __ldg(&features[feat_base + market_dim + i + 1]);
-            states_out[out_base + market_dim + 3 + i + 2] = __ldg(&features[feat_base + market_dim + i + 2]);
-            states_out[out_base + market_dim + 3 + i + 3] = __ldg(&features[feat_base + market_dim + i + 3]);
-        }
-        for (; i < ofi_dim; i++) {
-            states_out[out_base + market_dim + 3 + i] = __ldg(&features[feat_base + market_dim + i]);
-        }
+    // Drawdown from peak
+    float drawdown = (max_equity > 1.0f)
+                   ? (max_equity - value) / max_equity
+                   : 0.0f;
+    if (drawdown < 0.0f) drawdown = 0.0f;
 
-        // 4. Zero-pad [market_dim + 3 + ofi_dim .. state_dim) — 4x unroll
-        i = market_dim + 3 + ofi_dim;
-        for (; i + 3 < state_dim; i += 4) {
-            states_out[out_base + i]     = 0.0f;
-            states_out[out_base + i + 1] = 0.0f;
-            states_out[out_base + i + 2] = 0.0f;
-            states_out[out_base + i + 3] = 0.0f;
-        }
-        for (; i < state_dim; i++) {
-            states_out[out_base + i] = 0.0f;
-        }
-    } else {
-        // No OFI: straight copy [market, portfolio, pad]
-        // Market features with __ldg and manual 4x unroll
-        int i = 0;
-        for (; i + 3 < feat_dim; i += 4) {
-            states_out[out_base + i]     = __ldg(&features[feat_base + i]);
-            states_out[out_base + i + 1] = __ldg(&features[feat_base + i + 1]);
-            states_out[out_base + i + 2] = __ldg(&features[feat_base + i + 2]);
-            states_out[out_base + i + 3] = __ldg(&features[feat_base + i + 3]);
-        }
-        for (; i < feat_dim; i++) {
-            states_out[out_base + i] = __ldg(&features[feat_base + i]);
-        }
+    // Unrealized PnL: position * (current_close - entry_price)
+    // Use feature[0] as close price proxy (same as training kernel)
+    float close_now = __ldg(&features[feat_base]);
+    float unrealized_pnl = (entry_price > 0.0f && position != 0.0f)
+                         ? position * (close_now - entry_price)
+                         : 0.0f;
 
-        states_out[out_base + feat_dim + 0] = norm_value;
-        states_out[out_base + feat_dim + 1] = position;
-        states_out[out_base + feat_dim + 2] = spread_cost;
+    // Realized PnL: derive from equity state
+    // In backtest, realized PnL = value - initial_capital (approximate)
+    float realized_pnl = value - initial_capital;
 
-        // Zero-pad [feat_dim + 3 .. state_dim) — 4x unroll
-        i = feat_dim + 3;
-        for (; i + 3 < state_dim; i += 4) {
-            states_out[out_base + i]     = 0.0f;
-            states_out[out_base + i + 1] = 0.0f;
-            states_out[out_base + i + 2] = 0.0f;
-            states_out[out_base + i + 3] = 0.0f;
-        }
-        for (; i < state_dim; i++) {
-            states_out[out_base + i] = 0.0f;
+    // Trade return (unrealized / equity)
+    float trade_return = unrealized_pnl / equity;
+
+    // Capital floor distance (75% of peak)
+    float floor = max_equity * 0.75f;
+    float floor_dist = (equity > floor && equity > 1.0f)
+                     ? (equity - floor) / equity
+                     : 0.0f;
+
+    int pf_base = market_dim;
+    if (pf_base + 7 < state_dim) {
+        states_out[out_base + pf_base + 0] = position;                    // raw position
+        states_out[out_base + pf_base + 1] = unrealized_pnl / equity;     // trade P&L signal
+        states_out[out_base + pf_base + 2] = drawdown;                    // risk: how deep are we?
+        states_out[out_base + pf_base + 3] = step_count / 100.0f;         // how long in trade?
+        states_out[out_base + pf_base + 4] = realized_pnl / equity;       // session P&L
+        states_out[out_base + pf_base + 5] = floor_dist;                  // distance to game over
+        states_out[out_base + pf_base + 6] = trade_return;                // this trade's return
+        states_out[out_base + pf_base + 7] = cash / equity;               // available capital
+    }
+
+    // ── 3. Multi-timeframe features [market_dim+8 .. market_dim+8+16) ── //
+    // 4 lookback windows × 4 features = 16 GPU-native features.
+    // Matches training experience_state_gather exactly.
+    const int lookbacks[4] = {5, 15, 60, 240};
+    int mtf_base = market_dim + 8;
+
+    // Global bar index for this window at current step
+    int window_bar_base = w * max_len;
+
+    for (int lb = 0; lb < 4; lb++) {
+        int N_lb = lookbacks[lb];
+        int past_step = current_step - N_lb;
+        int slot = mtf_base + lb * 4;
+
+        if (past_step >= 0 && slot + 3 < state_dim) {
+            int now_feat_off  = (window_bar_base + current_step) * feat_dim;
+            int past_feat_off = (window_bar_base + past_step) * feat_dim;
+            float close_cur  = __ldg(&features[now_feat_off]);
+            float close_past = __ldg(&features[past_feat_off]);
+
+            // Return over N bars
+            float ret = (close_past > 0.0f) ? (close_cur - close_past) / close_past : 0.0f;
+            float scaled_ret = ret * 100.0f;
+            states_out[out_base + slot + 0] = fmaxf(-10.0f, fminf(10.0f, scaled_ret));
+
+            // Volatility: scan high/low over window (from close changes)
+            float max_val = close_cur;
+            float min_val = close_cur;
+            float vol_sum = 0.0f;
+            int vol_count = 0;
+            for (int j = past_step; j <= current_step; j++) {
+                int j_off = (window_bar_base + j) * feat_dim;
+                float v = __ldg(&features[j_off]);
+                if (v > max_val) max_val = v;
+                if (v < min_val) min_val = v;
+                // Volume proxy: feature index 4
+                if (feat_dim > 4) {
+                    vol_sum += __ldg(&features[j_off + 4]);
+                    vol_count++;
+                }
+            }
+            float range = (close_cur > 0.0f) ? (max_val - min_val) / close_cur : 0.0f;
+            float scaled_range = range * 100.0f;
+            states_out[out_base + slot + 1] = fmaxf(0.0f, fminf(10.0f, scaled_range));
+
+            // Volume trend: current vs average
+            float avg_vol = (vol_count > 0) ? vol_sum / (float)vol_count : 1.0f;
+            float cur_vol = (feat_dim > 4) ? __ldg(&features[now_feat_off + 4]) : 1.0f;
+            float vol_ratio = (avg_vol > 0.0f) ? cur_vol / avg_vol : 1.0f;
+            states_out[out_base + slot + 2] = fmaxf(0.0f, fminf(5.0f, vol_ratio));
+
+            // Momentum: position within range [0=bottom, 1=top]
+            float range_size = max_val - min_val;
+            states_out[out_base + slot + 3] = (range_size > 0.0f)
+                          ? (close_cur - min_val) / range_size
+                          : 0.5f;
+        } else {
+            // Not enough history — zero pad
+            if (slot + 3 < state_dim) {
+                states_out[out_base + slot + 0] = 0.0f;
+                states_out[out_base + slot + 1] = 0.0f;
+                states_out[out_base + slot + 2] = 0.0f;
+                states_out[out_base + slot + 3] = 0.0f;
+            }
         }
     }
+
+    // ── 4. OFI features (if present) ── //
+    if (ofi_dim > 0) {
+        int ofi_out_base = market_dim + 8 + 16;
+        // OFI features are stored after market features in the feature vector
+        for (i = 0; i + 3 < ofi_dim; i += 4) {
+            states_out[out_base + ofi_out_base + i]     = __ldg(&features[feat_base + market_dim + i]);
+            states_out[out_base + ofi_out_base + i + 1] = __ldg(&features[feat_base + market_dim + i + 1]);
+            states_out[out_base + ofi_out_base + i + 2] = __ldg(&features[feat_base + market_dim + i + 2]);
+            states_out[out_base + ofi_out_base + i + 3] = __ldg(&features[feat_base + market_dim + i + 3]);
+        }
+        for (; i < ofi_dim; i++) {
+            states_out[out_base + ofi_out_base + i] = __ldg(&features[feat_base + market_dim + i]);
+        }
+    }
+
+    // ── 5. Zero-pad remaining [filled .. state_dim) ── //
+    int filled = market_dim + 8 + 16 + ofi_dim;
+    for (i = filled; i + 3 < state_dim; i += 4) {
+        states_out[out_base + i]     = 0.0f;
+        states_out[out_base + i + 1] = 0.0f;
+        states_out[out_base + i + 2] = 0.0f;
+        states_out[out_base + i + 3] = 0.0f;
+    }
+    for (; i < state_dim; i++) {
+        states_out[out_base + i] = 0.0f;
+    }
+
+    // Suppress unused variable warning
+    (void)spread_cost;
 }
diff --git a/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu
index 193b70cc4..635bf933c 100644
--- a/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu
+++ b/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu
@@ -12,10 +12,10 @@
 //   [7] cvar_95 (mean of returns below VaR — Expected Shortfall)
 //   [8] calmar_ratio (annualized mean return / max drawdown)
 //   [9] omega_ratio (sum of gains / sum of losses)
-//  [10] buy_count  (actions 3,4 — Long50, Long100)
-//  [11] sell_count (actions 0,1 — Short100, Short50)
+//  [10] buy_count  (exposure_idx 5-8 — Long25..Long100)
+//  [11] sell_count (exposure_idx 0-3 — Short100..Short25)
 //  [12] hold_count (action 2 — Flat)
-//  [13] unique_action_count (popcount of 5-bit mask — how many of 0..4 were used)
+//  [13] unique_action_count (popcount of 9-bit mask — how many exposure levels 0..8 were used)
 
 extern "C" __global__ void compute_backtest_metrics(
     const float* __restrict__ step_returns,   // [n_windows * max_len]
@@ -65,7 +65,7 @@ extern "C" __global__ void compute_backtest_metrics(
     float local_cum = 0.0f, local_peak = 0.0f, local_max_dd = 0.0f;
     int local_wins = 0, local_trades = 0;
     int local_buys = 0, local_sells = 0, local_holds = 0;
-    int local_action_mask = 0;  // 5-bit bitmask: bit i set if action i was seen
+    int local_action_mask = 0;  // 9-bit bitmask: bit i set if exposure_idx i was seen
     int prev_action = -1;
 
     for (int i = tid; i < wlen; i += stride) {
@@ -89,13 +89,17 @@ extern "C" __global__ void compute_backtest_metrics(
         if (act != prev_action && i > 0) local_trades++;
         prev_action = act;
 
-        // Action distribution: 0,1=sell  2=hold  3,4=buy
-        if (act <= 1)      local_sells++;
-        else if (act == 2) local_holds++;
-        else               local_buys++;
+        // Action distribution: decode exposure from factored action.
+        // Factored action = exposure * (b1*b2) + order * b2 + urgency.
+        // exposure_idx: 0-3=sell, 4=flat, 5-8=buy (for 9-action space)
+        int exp_idx = act / (DQN_ORDER_ACTIONS * DQN_URGENCY_ACTIONS);
+        if (exp_idx < DQN_NUM_ACTIONS / 2)       local_sells++;
+        else if (exp_idx == DQN_NUM_ACTIONS / 2)  local_holds++;
+        else                                      local_buys++;
 
-        // Unique action tracking: set bit for each distinct action ID (0..4)
-        if (act >= 0 && act < 5) local_action_mask |= (1 << act);
+        // Unique action tracking: set bit per exposure index (0..8)
+        if (exp_idx >= 0 && exp_idx < DQN_NUM_ACTIONS)
+            local_action_mask |= (1 << exp_idx);
     }
 
     // Store ALL local values to shared memory
@@ -148,8 +152,8 @@ extern "C" __global__ void compute_backtest_metrics(
 
     // ── Unique action count via bitmask OR-reduction ─────────────────────────
     // Reuses s_sorted memory (safe: OR-reduction completes before bitonic sort).
-    // Each thread accumulated a 5-bit mask of observed action IDs. OR-reduce
-    // across the block, then popcount gives unique_action_count (1-5).
+    // Each thread accumulated a 9-bit mask of observed exposure indices. OR-reduce
+    // across the block, then popcount gives unique_action_count (1-9).
     {
         int* s_action_mask = (int*)s_sorted;  // alias — sequential use, no conflict
         s_action_mask[tid] = local_action_mask;
diff --git a/crates/ml/src/cuda_pipeline/curiosity_training_kernel.cu b/crates/ml/src/cuda_pipeline/curiosity_training_kernel.cu
index b7e31cf40..a2a2b3c3d 100644
--- a/crates/ml/src/cuda_pipeline/curiosity_training_kernel.cu
+++ b/crates/ml/src/cuda_pipeline/curiosity_training_kernel.cu
@@ -94,11 +94,15 @@ extern "C" __global__ void curiosity_forward_backward(
     input[MARKET_DIM + 1] = 0.0f;
     input[MARKET_DIM + 2] = 0.0f;
 
-    /* Action to category one-hot (5-action DQN mapping from common_device_functions.cuh) */
+    /* Action to category one-hot — decode exposure from factored action.
+     * Factored action = exposure_idx * (b1*b2) + order_idx * b2 + urgency_idx.
+     * With b0=9, b1=3, b2=3: exposure_idx = action / 9.
+     * Categories: Short (idx 0-3), Flat (idx 4), Long (idx 5-8). */
+    int exposure_idx = action_idx / (DQN_ORDER_ACTIONS * DQN_URGENCY_ACTIONS);
     int category;
-    if (action_idx <= 1)      category = 0;  /* Short100/Short50 */
-    else if (action_idx == 2) category = 1;  /* Flat */
-    else                      category = 2;  /* Long50/Long100 */
+    if (exposure_idx < DQN_NUM_ACTIONS / 2)       category = 0;  /* Short */
+    else if (exposure_idx == DQN_NUM_ACTIONS / 2)  category = 1;  /* Flat */
+    else                                           category = 2;  /* Long */
     input[MARKET_DIM + category] = 1.0f;
 
     /* Layer 1: pre_act = w1 * input + b1, hidden = LeakyReLU(pre_act, alpha=0.01) */
diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu
index aa767d2b8..b1f479944 100644
--- a/crates/ml/src/cuda_pipeline/experience_kernels.cu
+++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu
@@ -104,16 +104,21 @@ __device__ __forceinline__ int argmax_n(const float* arr, int n) {
 }
 
 /**
- * Map exposure branch index (0–4) to target exposure fraction.
- *   0 → -1.0  (Short100)
- *   1 → -0.5  (Short50)
- *   2 →  0.0  (Flat)
- *   3 →  0.5  (Long50)
- *   4 →  1.0  (Long100)
+ * Map exposure branch index (0–8) to target exposure fraction.
+ *   0 → -1.00  (S100)
+ *   1 → -0.75  (S75)
+ *   2 → -0.50  (S50)
+ *   3 → -0.25  (S25)
+ *   4 →  0.00  (Flat)
+ *   5 → +0.25  (L25)
+ *   6 → +0.50  (L50)
+ *   7 → +0.75  (L75)
+ *   8 → +1.00  (L100)
+ *
+ * Formula: -1.0 + idx * (2.0 / (b0_size - 1)).
+ * With b0_size=9: step = 0.25. Flat = index 4.
  */
 __device__ __forceinline__ float exposure_idx_to_fraction(int idx) {
-    /* 9 levels: -1.0, -0.75, -0.50, -0.25, 0.0, +0.25, +0.50, +0.75, +1.0
-     * Formula: -1.0 + idx * 0.25. Flat = index 4. */
     return -1.0f + (float)idx * 0.25f;
 }
 
@@ -597,84 +602,31 @@ extern "C" __global__ void experience_env_step(
     float target_exposure = exposure_idx_to_fraction(exposure_idx);
     float target_position = target_exposure * max_position;
 
-    /* Apply CVaR risk scaling: reduce position when tail risk is high */
-    if (cvar_scales != NULL) {
-        target_position *= cvar_scales[i];
-    }
-
-    /* Apply Q-gap conviction scaling: size UP with conviction.
-     * conviction = clamp(q_gap / 2.0, 0.25, 1.0)
-     * Low Q-gap (0.1-0.5) → small position (25-50% of target)
-     * High Q-gap (2.0+) → full position (100%)
-     * Combined with CVaR: size = exposure × CVaR_scale × conviction */
-    if (q_gaps != NULL) {
-        float conviction = q_gaps[i] / 2.0f;
-        conviction = (conviction < 0.25f) ? 0.25f : ((conviction > 1.0f) ? 1.0f : conviction);
-        target_position *= conviction;
-    }
-
-    /* ════════════════════════════════════════════════════════════════
-     * GPU-NATIVE ENHANCED KELLY CRITERION (full implementation)
+    /* CVaR, conviction, and Kelly scaling DISABLED during training.
      *
-     * Weighted blend of two Kelly estimators:
-     *   1. Continuous Kelly: f* = μ / σ² (Gaussian approximation)
-     *   2. Discrete Kelly:   f* = (b*p - q) / b (Bernoulli model)
-     *   Enhanced: f* = 0.7 × continuous + 0.3 × discrete
+     * These overrides destroy credit assignment: the agent selects an exposure
+     * level, but the actual position is silently rescaled by CVaR, Q-gap
+     * conviction, and Kelly fraction. The reward for action A then reflects
+     * the outcome of a DIFFERENT position size, corrupting Q-value learning.
      *
-     * With confidence scaling from sample size:
-     *   confidence = 1 - 1/√(n_trades)  (20 trades = 0.78, 100 = 0.90)
-     *   final_f = f* × confidence × fractional_kelly (0.5 = half-Kelly)
-     *
-     * All statistics accumulated in portfolio state ps[14]-ps[19].
-     * Zero CPU involvement. Pure GPU. Updated only at trade completions.
-     * ════════════════════════════════════════════════════════════════ */
-    {
-        float total_trades = win_count + loss_count;
-        if (total_trades >= 20.0f) {
-            /* ── Continuous Kelly: f* = μ / σ² ── */
-            float mu = sum_returns / total_trades;
-            float variance = (sum_sq_returns / total_trades) - (mu * mu);
-            float continuous_kelly = (variance > 1e-10f)
-                                   ? mu / variance
-                                   : 0.0f;
+     * Stop-loss/take-profit and capital floor are kept — they are environment
+     * rules (hard constraints), not action overrides. */
+    (void)cvar_scales;
+    (void)q_gaps;
 
-            /* ── Discrete Kelly: f* = (b*p - q) / b ── */
-            float p = win_count / total_trades;
-            float q = 1.0f - p;
-            float avg_win = sum_wins / fmaxf(win_count, 1.0f);
-            float avg_loss = sum_losses / fmaxf(loss_count, 1.0f);
-            float b = avg_win / fmaxf(avg_loss, 1e-6f);
-            float discrete_kelly = (b > 1e-6f)
-                                 ? (b * p - q) / b
-                                 : 0.0f;
+    /* Kelly criterion DISABLED during training (same rationale as CVaR/conviction above).
+     * Kelly statistics (ps[14]-ps[19]) are still accumulated for logging/diagnostics,
+     * but the position size is NOT rescaled. The agent must learn its own sizing
+     * through the reward signal, not have it overridden by a running Kelly estimate. */
+    (void)win_count; (void)loss_count;
+    (void)sum_wins; (void)sum_losses;
+    (void)sum_returns; (void)sum_sq_returns;
 
-            /* ── Enhanced Kelly: weighted blend ── */
-            /* NaN guard: if either Kelly produces NaN/Inf, use 0.0 (safe default) */
-            if (isnan(continuous_kelly) || isinf(continuous_kelly)) continuous_kelly = 0.0f;
-            if (isnan(discrete_kelly) || isinf(discrete_kelly)) discrete_kelly = 0.0f;
-            float enhanced_kelly = 0.7f * continuous_kelly + 0.3f * discrete_kelly;
-
-            /* ── Confidence scaling: penalize small samples ── */
-            float confidence = 1.0f - 1.0f / sqrtf(total_trades);
-
-            /* ── Fractional Kelly (half-Kelly = 0.5) ── */
-            float fractional = 0.5f;
-            float kelly_frac = enhanced_kelly * confidence * fractional;
-
-            /* ── Clamp to [0.05, 0.25] ── */
-            kelly_frac = fmaxf(0.05f, fminf(0.25f, kelly_frac));
-
-            /* ── Apply to position (normalize: 0.25 = full size) ── */
-            target_position *= kelly_frac / 0.25f;
-        }
-        /* < 20 trades: no Kelly applied (full position until we have statistics) */
-
-    /* Final NaN/Inf guard on target_position after ALL scaling.
-     * If any scaling (CVaR, conviction, Kelly) produced NaN, zero the position. */
+    /* Final NaN/Inf guard on target_position.
+     * If exposure_idx_to_fraction produced NaN, zero the position. */
     if (isnan(target_position) || isinf(target_position)) {
         target_position = 0.0f;
     }
-    }
 
     /* ---- Position adjustment with volatility-scaled transaction cost ---- */
     /* Real spread widens in volatile markets. vol_scale (from CUSUM) is
@@ -707,7 +659,7 @@ extern "C" __global__ void experience_env_step(
                             : (order_type_idx == 1) ? 0.0002f   /* IoC: +2 bps */
                             : -0.0005f;                          /* LimitMaker: -5 bps rebate */
 
-        tx_cost   = fabsf(delta) * raw_close * (tx_cost_multiplier * spread_scale * impact_scale + order_premium);
+        tx_cost   = fabsf(delta) * raw_close * (tx_cost_multiplier * 0.0001f * spread_scale * impact_scale + order_premium);
         cash     -= delta * raw_close;
         cash     -= tx_cost;
         position  = target_position;
@@ -905,7 +857,7 @@ extern "C" __global__ void experience_env_step(
     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;
+        float exit_cost = fabsf(ps[0]) * raw_close * tx_cost_multiplier * 0.0001f;
         cash = ps[1] + ps[0] * raw_close - exit_cost;
         is_flat = 1.0f;
         exiting_trade = 1;
diff --git a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs
index 176604537..f9898cc65 100644
--- a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs
@@ -79,13 +79,25 @@ fn compile_env_ptx(context: &CudaContext) -> Result<Ptx, String> {
 }
 
 fn compile_metrics_ptx(context: &CudaContext) -> Result<Ptx, String> {
+    let defines = "\
+        #define STATE_DIM 72\n\
+        #define MARKET_DIM 42\n\
+        #define PORTFOLIO_DIM 8\n";
+    let common_src = include_str!("common_device_functions.cuh");
     let src = include_str!("backtest_metrics_kernel.cu");
-    crate::cuda_pipeline::compile_ptx_for_device(src, context)
+    let full_source = format!("{defines}{common_src}\n{src}");
+    crate::cuda_pipeline::compile_ptx_for_device(&full_source, context)
 }
 
 fn compile_gather_ptx(context: &CudaContext) -> Result<Ptx, String> {
+    let defines = "\
+        #define STATE_DIM 72\n\
+        #define MARKET_DIM 42\n\
+        #define PORTFOLIO_DIM 8\n";
+    let common_src = include_str!("common_device_functions.cuh");
     let src = include_str!("backtest_gather_kernel.cu");
-    crate::cuda_pipeline::compile_ptx_for_device(src, context)
+    let full_source = format!("{defines}{common_src}\n{src}");
+    crate::cuda_pipeline::compile_ptx_for_device(&full_source, context)
 }
 
 /// Compile PPO backtest forward kernel with common device functions prepended.
@@ -570,11 +582,12 @@ impl GpuBacktestEvaluator {
             .alloc_zeros::<f32>(n_windows * 14)
             .map_err(|e| MLError::ModelError(format!("metrics alloc: {e}")))?;
 
-        // Portfolio dimension is always 3: (normalised value, position, spread_cost).
-        // state_dim = feature_dim + 3, then 8-aligned for H100 tensor core HMMA dispatch.
-        // The CUDA gather_states kernel zero-pads positions [feat_dim+3 .. state_dim).
-        const PORTFOLIO_DIM: usize = 8;
-        let state_dim = (feature_dim + PORTFOLIO_DIM + 7) & !7;
+        // State layout matches training experience_state_gather:
+        //   [0..42) market features + [42..50) 8 portfolio + [50..66) 16 multi-timeframe
+        //   Total filled = feat_dim + 8 portfolio + 16 multi-timeframe = feat_dim + 24
+        //   8-aligned for H100 tensor core HMMA dispatch.
+        const PORTFOLIO_AND_MTF_DIM: usize = 24; // 8 portfolio + 16 multi-timeframe
+        let state_dim = (feature_dim + PORTFOLIO_AND_MTF_DIM + 7) & !7;
         let states_buf = stream
             .alloc_zeros::<f32>(n_windows * state_dim)
             .map_err(|e| MLError::ModelError(format!("states_buf alloc: {e}")))?;
@@ -611,7 +624,7 @@ impl GpuBacktestEvaluator {
             n_windows,
             max_len,
             feature_dim,
-            portfolio_dim: PORTFOLIO_DIM,
+            portfolio_dim: PORTFOLIO_AND_MTF_DIM,
             state_dim,
             config,
             forward_actions_buf,
@@ -691,8 +704,8 @@ impl GpuBacktestEvaluator {
     /// CUDA storage — eliminating the GPU→CPU→GPU roundtrip entirely.
     ///
     /// # Panics
-    /// `portfolio_dim` must equal `self.portfolio_dim` (always 3). Callers using a
-    /// different value should be updated — the kernel signature is fixed.
+    /// `portfolio_dim` must equal `self.portfolio_dim` (24: 8 portfolio + 16 multi-timeframe).
+    /// This matches the training experience_state_gather layout.
     pub fn gather_states(
         &self,
         step: usize,
@@ -1315,7 +1328,7 @@ impl GpuBacktestEvaluator {
     /// * `forward_fn` — model forward: `&CudaSlice<f32>[N * state_dim]` → `CudaSlice<f32>[N]` predictions
     /// * `high_threshold_bps` — strong signal threshold (actions 0 and 4)
     /// * `low_threshold_bps` — mild signal threshold (actions 1 and 3)
-    /// * `portfolio_dim` — must be 3
+    /// * `portfolio_dim` — must be 24 (8 portfolio + 16 multi-timeframe)
     pub fn evaluate_supervised<F>(
         &mut self,
         forward_fn: &F,
@@ -1731,6 +1744,10 @@ impl GpuBacktestEvaluator {
         let n_win_i32 = self.n_windows as i32;
         let max_len_i32 = self.max_len as i32;
         let step_i32 = step as i32;
+        // Branch sizes for factored action decoding (default 9/3/3)
+        let b0_i32: i32 = 9;
+        let b1_i32: i32 = 3;
+        let b2_i32: i32 = 3;
 
         unsafe {
             self.stream
@@ -1749,6 +1766,9 @@ impl GpuBacktestEvaluator {
                 .arg(&self.config.tx_cost_bps)
                 .arg(&self.config.spread_cost)
                 .arg(&step_i32)
+                .arg(&b0_i32)
+                .arg(&b1_i32)
+                .arg(&b2_i32)
                 .launch(env_cfg)
                 .map_err(|e| {
                     MLError::ModelError(format!("backtest_env_step launch step {step}: {e}"))
@@ -1980,13 +2000,14 @@ mod tests {
         }
     }
 
-    /// Verify that gather_states() returns a ConfigError when portfolio_dim != 3.
+    /// Verify that gather_states() returns a ConfigError when portfolio_dim != 24.
     #[test]
     fn test_gather_states_portfolio_dim_mismatch() {
         // We don't have a CUDA device in unit tests, so test the validation path
-        // indirectly by checking the struct's portfolio_dim field is always 3.
-        // The real mismatch check is exercised at runtime when portfolio_dim != 3.
-        assert_eq!(3_usize, 3_usize, "PORTFOLIO_DIM constant is 3");
+        // indirectly by checking the struct's portfolio_dim field is always 24.
+        // The real mismatch check is exercised at runtime when portfolio_dim != 24.
+        // 24 = 8 portfolio features + 16 multi-timeframe features.
+        assert_eq!(24_usize, 24_usize, "PORTFOLIO_AND_MTF_DIM constant is 24");
     }
 
     /// Verify that the cuBLAS DQN helper kernels (compute_expected_q,
@@ -2031,16 +2052,19 @@ mod tests {
 
     #[test]
     fn test_gpu_backtest_evaluator_state_dim_calculation() {
-        // state_dim must equal feature_dim + PORTFOLIO_DIM (3)
+        // state_dim = (feature_dim + 24 + 7) & !7
+        // 24 = 8 portfolio + 16 multi-timeframe (matches training layout)
         let feature_dim: usize = 42;
-        let portfolio_dim: usize = 3;
-        let state_dim = feature_dim + portfolio_dim;
-        assert_eq!(state_dim, 45);
+        let portfolio_and_mtf: usize = 24;
+        let state_dim = (feature_dim + portfolio_and_mtf + 7) & !7;
+        // 42 + 24 = 66, aligned to 72
+        assert_eq!(state_dim, 72);
 
-        // With 8-alignment padding (56 features + 3)
-        let feature_dim_aligned: usize = 53; // 53-feature state
-        let state_dim_aligned = feature_dim_aligned + portfolio_dim;
-        assert_eq!(state_dim_aligned, 56);
+        // With OFI (feature_dim = 50 = 42 market + 8 OFI in feature vector)
+        let feature_dim_ofi: usize = 50;
+        let state_dim_ofi = (feature_dim_ofi + portfolio_and_mtf + 7) & !7;
+        // 50 + 24 = 74, aligned to 80
+        assert_eq!(state_dim_ofi, 80);
     }
 
     /// Verify that CUDA Graph infrastructure compiles and struct fields are present.
diff --git a/crates/ml/src/cuda_pipeline/gpu_monitoring.rs b/crates/ml/src/cuda_pipeline/gpu_monitoring.rs
index 7878e4026..39ac4821b 100644
--- a/crates/ml/src/cuda_pipeline/gpu_monitoring.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_monitoring.rs
@@ -34,8 +34,14 @@ pub struct GpuMonitoringReducer {
 static MONITORING_PTX: OnceLock<Result<Ptx, String>> = OnceLock::new();
 
 fn compile_monitoring_ptx(context: &CudaContext) -> Result<Ptx, String> {
+    let defines = "\
+        #define STATE_DIM 72\n\
+        #define MARKET_DIM 42\n\
+        #define PORTFOLIO_DIM 8\n";
+    let common_src = include_str!("common_device_functions.cuh");
     let kernel_src = include_str!("monitoring_kernel.cu");
-    crate::cuda_pipeline::compile_ptx_for_device(kernel_src, context)
+    let full_source = format!("{defines}{common_src}\n{kernel_src}");
+    crate::cuda_pipeline::compile_ptx_for_device(&full_source, context)
 }
 
 impl GpuMonitoringReducer {
diff --git a/crates/ml/src/cuda_pipeline/monitoring_kernel.cu b/crates/ml/src/cuda_pipeline/monitoring_kernel.cu
index cb05e2e57..018312fbb 100644
--- a/crates/ml/src/cuda_pipeline/monitoring_kernel.cu
+++ b/crates/ml/src/cuda_pipeline/monitoring_kernel.cu
@@ -57,8 +57,11 @@ extern "C" __global__ void monitoring_reduce(
         local_sq += r * r;
         local_min = fminf(local_min, r);
         local_max = fmaxf(local_max, r);
+        /* Decode exposure index from factored action for monitoring.
+         * Factored action = exposure * (b1*b2) + order * b2 + urgency. */
         int a = actions[i];
-        if (a >= 0 && a < 9) local_counts[a]++;
+        int exp_idx = a / (DQN_ORDER_ACTIONS * DQN_URGENCY_ACTIONS);
+        if (exp_idx >= 0 && exp_idx < DQN_NUM_ACTIONS) local_counts[exp_idx]++;
     }
 
     // Warp-level reduction before atomic to shared (reduces contention 32x)
diff --git a/crates/ml/src/hyperopt/adapters/ppo.rs b/crates/ml/src/hyperopt/adapters/ppo.rs
index a7fd78ef0..0f8251e92 100644
--- a/crates/ml/src/hyperopt/adapters/ppo.rs
+++ b/crates/ml/src/hyperopt/adapters/ppo.rs
@@ -1328,7 +1328,7 @@ impl PPOTrainer {
                     MLError::ModelError(format!("PPO backtest action upload: {e}"))
                 })
             },
-            3, // portfolio_dim
+            24, // portfolio_dim: 8 portfolio + 16 multi-timeframe (matches training state layout)
         )?;
 
         let first = metrics.first().ok_or_else(|| {
diff --git a/crates/ml/src/trainers/dqn/config.rs b/crates/ml/src/trainers/dqn/config.rs
index 1f0739843..9f859fef8 100644
--- a/crates/ml/src/trainers/dqn/config.rs
+++ b/crates/ml/src/trainers/dqn/config.rs
@@ -1566,7 +1566,7 @@ use crate::dqn::dqn::DQNConfig;
 pub(crate) fn dqn_default_config() -> DQNConfig {
     DQNConfig {
         // Network Architecture
-        state_dim: 48,  // 42 market + 3 portfolio = 45, aligned to 48 for tensor cores
+        state_dim: 48,  // Default placeholder; overridden by constructor to (42+8+16+7)&!7=72
         num_actions: 9,  // 9 exposure levels (25% steps)
         hidden_dims: vec![512, 256, 128],  // Sufficient capacity without overfitting (Trial 19)
         leaky_relu_alpha: 0.01,
diff --git a/crates/ml/src/trainers/dqn/trainer/metrics.rs b/crates/ml/src/trainers/dqn/trainer/metrics.rs
index 1cb08f0bd..0f0fbf3da 100644
--- a/crates/ml/src/trainers/dqn/trainer/metrics.rs
+++ b/crates/ml/src/trainers/dqn/trainer/metrics.rs
@@ -241,7 +241,7 @@ impl DQNTrainer {
             let row_offset = r * cols;
             let mut best_val = f32::NEG_INFINITY;
             let mut second_best = f32::NEG_INFINITY;
-            for c in 0..cols.min(5) {
+            for c in 0..cols.min(9) {
                 let v = host_q.get(row_offset + c).copied().unwrap_or(0.0);
                 if let Some(s) = col_sums.get_mut(c) {
                     *s += v as f64;
@@ -255,8 +255,9 @@ impl DQNTrainer {
                     // v <= second_best: no update needed
                 }
             }
-            // Also check remaining columns (6..total_actions) for best/second_best
-            for c in 5..cols {
+            // Also check remaining columns (9..total_actions) for best/second_best
+            // (order_type + urgency branches, if total_actions > 9)
+            for c in 9..cols {
                 let v = host_q.get(row_offset + c).copied().unwrap_or(0.0);
                 if v > best_val {
                     second_best = best_val;
diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
index 1ec9c6167..fb108668e 100644
--- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs
+++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
@@ -1884,7 +1884,7 @@ impl DQNTrainer {
                     gap_mean, gap_min, gap_max
                 );
 
-                let names = ["S100", "S50", "Flat", "L50", "L100"];
+                let names = ["S100", "S75", "S50", "S25", "Flat", "L25", "L50", "L75", "L100"];
                 let parts: Vec<String> = names.iter()
                     .zip(per_action_avgs.iter())
                     .map(|(n, q)| format!("{}={:.4}", n, q))