diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu
index 6fc9df77d..e043d8d13 100644
--- a/crates/ml/src/cuda_pipeline/experience_kernels.cu
+++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu
@@ -440,7 +440,8 @@ extern "C" __global__ void experience_env_step(
     int b2_size,
     int current_t,
     const float* __restrict__ cvar_scales,  /* [N] or NULL — CVaR position scaling */
-    const float* __restrict__ q_gaps        /* [N] or NULL — Q-gap conviction scaling */
+    const float* __restrict__ q_gaps,       /* [N] or NULL — Q-gap conviction scaling */
+    float kelly_scale                        /* Kelly optimal fraction (0.0-1.0, from CPU per-epoch) */
 ) {
     int i = blockIdx.x * blockDim.x + threadIdx.x;
     if (i >= N) return;
@@ -528,6 +529,15 @@ extern "C" __global__ void experience_env_step(
         target_position *= conviction;
     }
 
+    /* Kelly optimal fraction: computed on CPU per-epoch from trade statistics.
+     * f* = (b*p - q) / b where p=win_rate, b=avg_win/avg_loss, q=1-p.
+     * Half-Kelly (×0.5) is standard conservative sizing.
+     * Applied AFTER CVaR and conviction — all three stack multiplicatively:
+     * final_size = exposure × CVaR_scale × conviction × kelly_scale */
+    if (kelly_scale > 0.0f && kelly_scale < 1.0f) {
+        target_position *= kelly_scale;
+    }
+
     /* ---- Position adjustment with volatility-scaled transaction cost ---- */
     /* Real spread widens in volatile markets. vol_scale (from CUSUM) is
      * computed later, but we can read CUSUM here for the cost calculation.
@@ -543,13 +553,23 @@ extern "C" __global__ void experience_env_step(
     float tx_cost = 0.0f;
 
     if (delta != 0.0f) {
-        /* Market impact: larger trades cost more (quadratic in size).
-         * A 1-lot order fills at bid/ask. A 4-lot order moves the market.
-         * impact_scale = 1 + |delta/max_position|^2 — ranges from 1.0 to 2.0.
-         * This teaches the model that scaling to max position is expensive. */
+        /* Square-root market impact (Almgren & Chriss, 2000).
+         * impact_scale = 1 + sqrt(|delta|/max_position) — standard academic model.
+         * √x grows slower than x² → less penalty for moderate sizes,
+         * more realistic than quadratic for futures markets. */
         float size_ratio = fabsf(delta) / (max_position > 0.0f ? max_position : 1.0f);
-        float impact_scale = 1.0f + size_ratio * size_ratio;
-        tx_cost   = fabsf(delta) * raw_close * tx_cost_multiplier * spread_scale * impact_scale;
+        float impact_scale = 1.0f + sqrtf(size_ratio);
+
+        /* Order-type cost differentiation from branching action:
+         * Factored action = exposure * 9 + order_type * 3 + urgency
+         * order_type: 0=Market (+0 bps), 1=IoC (+2 bps), 2=LimitMaker (-5 bps rebate)
+         * LimitMaker gets a REBATE — teaches the model to prefer limit orders. */
+        int order_type_idx = (action_idx / b2_size) % b1_size;
+        float order_premium = (order_type_idx == 0) ? 0.0f
+                            : (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);
         cash     -= delta * raw_close;
         cash     -= tx_cost;
         position  = target_position;
diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
index 10116cff4..1862ada63 100644
--- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
@@ -205,6 +205,10 @@ pub struct ExperienceCollectorConfig {
     /// When greedy Q(best_exposure) - Q(flat) < threshold, default to flat.
     /// 0.0 = disabled (trade on every bar). Typical range [0.0, 0.5].
     pub q_gap_threshold: f32,
+    /// Kelly optimal fraction for position sizing (0.0-1.0).
+    /// Computed per-epoch on CPU from trade statistics, uploaded as scalar.
+    /// 1.0 = full Kelly (aggressive), 0.5 = half-Kelly (conservative).
+    pub kelly_scale: f32,
     /// DSR EMA decay rate (e.g. 0.01 = ~100-step window). DSR is always enabled.
     pub dsr_eta: f32,
     /// N-step returns lookahead (1 = standard TD, 3-5 typical for Rainbow DQN)
@@ -259,6 +263,7 @@ impl Default for ExperienceCollectorConfig {
             fill_spread_capture_frac: 0.50,
             fill_simulation_enabled: false,
             q_gap_threshold: 0.0,
+            kelly_scale: 1.0, // full Kelly by default (no reduction)
             dsr_eta: 0.01,
             n_steps: 1,
             enable_action_masking: false,
@@ -369,6 +374,8 @@ pub struct GpuExperienceCollector {
     batch_actions: CudaSlice<i32>,
     /// Q-gap conviction per episode [N]: output by action_select, input to env_step.
     q_gaps_buf: CudaSlice<f32>,
+    /// Kelly optimal fraction (per-epoch scalar from CPU).
+    kelly_scale: f32,
     /// Current timestep counter per episode: [N].
     current_timesteps: CudaSlice<i32>,
     /// Q-values output from expected Q or direct logits: [N, total_branch_actions].
@@ -788,6 +795,7 @@ impl GpuExperienceCollector {
             batch_states,
             batch_actions,
             q_gaps_buf,
+            kelly_scale: 1.0,
             current_timesteps,
             q_values,
             exp_h_s1,
@@ -1094,6 +1102,7 @@ impl GpuExperienceCollector {
                     .arg(&t_i32)
                     .arg(&self.cvar_scales_ptr) // CVaR position scaling (0 = NULL = no scaling)
                     .arg(&self.q_gaps_buf)     // Q-gap conviction scaling
+                    .arg(&config.kelly_scale)  // Kelly optimal fraction (per-epoch from CPU)
                     .launch(launch_cfg)
                     .map_err(|e| MLError::ModelError(format!(
                         "experience_env_step t={t}: {e}"