fix: spread_cost wired to training kernel — eliminates last train/eval gap

- Add spread_cost as 28th param to experience_env_step kernel
- Wire from DQNHyperparameters.fill_spread_cost_frac through
  ExperienceCollectorConfig to the kernel launch
- Training execute_trade() and compute_tx_cost() now use spread_cost
  (was hardcoded 0.0, backtest used real spread_cost)
- All compute_tx_cost calls in training use spread_cost consistently

This was the LAST remaining train/eval transaction cost mismatch.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

crates/ml/src/cuda_pipeline/experience_kernels.cu

@@ -532,7 +532,8 @@ 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 (unshapen) */
-    int min_hold_bars  /* minimum bars to hold before exiting or reversing */
+    int min_hold_bars,  /* minimum bars to hold before exiting or reversing */
+    float spread_cost   /* bid-ask spread cost per unit (matches backtest) */
 ) {
     int i = blockIdx.x * blockDim.x + threadIdx.x;
     if (i >= N) return;
@@ -753,7 +754,7 @@ extern "C" __global__ void experience_env_step(
     /* ---- Execute trade (shared: trade_physics.cuh) ---- */
     int order_type_idx = decode_order_type(action_idx, b1_size, b2_size);
     float tx_cost = execute_trade(&position, &cash, target_position, raw_close,
-                                  tx_cost_multiplier, 0.0f, max_position,
+                                  tx_cost_multiplier, spread_cost, max_position,
                                   order_type_idx, spread_scale);
 
     /* ---- Mark-to-market PnL for this timestep ---- */
@@ -930,7 +931,7 @@ extern "C" __global__ void experience_env_step(
         /* Charge tx cost on forced liquidation — production exits aren't free */
         if (fabsf(position) > 0.001f) {
             int order_type_idx_liq = decode_order_type(action_idx, b1_size, b2_size);
-            float exit_cost = compute_tx_cost(position, raw_close, tx_cost_multiplier, 0.0f,
+            float exit_cost = compute_tx_cost(position, raw_close, tx_cost_multiplier, spread_cost,
                                               max_position, order_type_idx_liq, -1.0f);
             new_portfolio_value -= exit_cost;
         }

crates/ml/src/cuda_pipeline/gpu_experience_collector.rs

@@ -239,6 +239,8 @@ pub struct ExperienceCollectorConfig {
     pub enable_action_masking: bool,
     /// Minimum bars to hold a position before exiting or reversing (Layer 2 hold enforcement).
     pub min_hold_bars: i32,
+    /// Bid-ask spread cost per unit position change. Matches backtest_env_kernel's spread_cost.
+    pub spread_cost: f32,
 }
 
 impl Default for ExperienceCollectorConfig {
@@ -287,6 +289,7 @@ impl Default for ExperienceCollectorConfig {
             n_steps: 1,
             enable_action_masking: false,
             min_hold_bars: 5,
+            spread_cost: 0.0,  // default: no spread cost (overridden by hyperparams)
         }
     }
 }
@@ -1334,6 +1337,7 @@ impl GpuExperienceCollector {
                     .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.spread_cost)         // bid-ask spread cost (matches backtest)
                     .launch(launch_cfg)
                     .map_err(|e| MLError::ModelError(format!(
                         "experience_env_step t={t}: {e}"

crates/ml/src/trainers/dqn/trainer/training_loop.rs

@@ -1050,6 +1050,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,
+            // ES futures: tick_size=0.25, typical spread=1 tick, cost = tick * multiplier * fraction
+            // Matches backtest_env_kernel's spread_cost from GpuBacktestConfig
+            spread_cost: (0.25 * 50.0 * self.hyperparams.fill_spread_cost_frac) as f32,
             ..Default::default()
         };