diff --git a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu
new file mode 100644
index 000000000..2bef36b11
--- /dev/null
+++ b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu
@@ -0,0 +1,123 @@
+// Vectorized backtest environment step kernel.
+// One thread per walk-forward window. Each thread steps sequentially.
+//
+// Portfolio state layout per window [8 floats]:
+//   [0] value        - current portfolio value
+//   [1] position     - current position size (-1.0 to +1.0)
+//   [2] cash         - cash balance
+//   [3] entry_price  - entry price of current position (0 if flat)
+//   [4] max_equity   - peak equity for drawdown tracking
+//   [5] step_pnl     - PnL this step (for reward)
+//   [6] cum_return   - cumulative log return
+//   [7] step_count   - number of completed steps
+
+#define PORTFOLIO_STATE_SIZE 8
+
+extern "C" __global__ void backtest_env_step(
+    // Market data (read-only, uploaded once)
+    const float* __restrict__ prices,       // [n_windows * max_len * 4] (OHLC)
+    const int*   __restrict__ window_lens,  // [n_windows]
+
+    // Actions from model for current step
+    const int*   __restrict__ actions,      // [n_windows] (0-4: Short100..Long100)
+
+    // Portfolio state (read-write, persistent across steps)
+    float*       portfolio_state,           // [n_windows * PORTFOLIO_STATE_SIZE]
+
+    // Step outputs
+    float*       step_rewards,              // [n_windows]
+    float*       step_returns,              // [n_windows * max_len] (accumulated)
+    int*         done_flags,                // [n_windows]
+
+    // Config
+    int          n_windows,
+    int          max_len,
+    float        max_position,
+    float        tx_cost_bps,
+    float        spread_cost,
+    int          current_step
+) {
+    int w = blockIdx.x * blockDim.x + threadIdx.x;
+    if (w >= n_windows) return;
+    if (done_flags[w]) return;
+
+    int wlen = window_lens[w];
+    if (current_step >= wlen) {
+        done_flags[w] = 1;
+        return;
+    }
+
+    // Read current prices
+    int price_base = (w * max_len + current_step) * 4;
+    float open  = prices[price_base + 0];
+    float high  = prices[price_base + 1];
+    float low   = prices[price_base + 2];
+    float close = prices[price_base + 3];
+
+    // Read portfolio state
+    int ps = w * PORTFOLIO_STATE_SIZE;
+    float value      = portfolio_state[ps + 0];
+    float position   = portfolio_state[ps + 1];
+    float cash       = portfolio_state[ps + 2];
+    float entry_price = portfolio_state[ps + 3];
+    float max_equity = portfolio_state[ps + 4];
+    float cum_return = portfolio_state[ps + 6];
+
+    // Map action (0-4) to target exposure
+    float target_exposure;
+    switch (actions[w]) {
+        case 0: target_exposure = -1.0f; break;  // Short100
+        case 1: target_exposure = -0.5f; break;  // Short50
+        case 2: target_exposure =  0.0f; break;  // Flat
+        case 3: target_exposure =  0.5f; break;  // Long50
+        case 4: target_exposure =  1.0f; break;  // Long100
+        default: target_exposure = 0.0f; break;
+    }
+    target_exposure *= max_position;
+
+    // Execute trade if position changes
+    float delta = target_exposure - position;
+    float trade_cost = 0.0f;
+    if (fabsf(delta) > 0.001f && close > 0.0f) {
+        trade_cost = fabsf(delta) * close * tx_cost_bps * 0.0001f
+                   + fabsf(delta) * spread_cost * 0.5f;
+        cash -= trade_cost;
+
+        // Mark-to-market old position
+        if (fabsf(position) > 0.001f && entry_price > 0.0f) {
+            float pnl = position * (close - entry_price);
+            cash += pnl;
+        }
+
+        position = target_exposure;
+        entry_price = close;
+    }
+
+    // Mark-to-market current position
+    float unrealized = 0.0f;
+    if (fabsf(position) > 0.001f && entry_price > 0.0f) {
+        unrealized = position * (close - entry_price);
+    }
+    float new_value = cash + unrealized;
+
+    // Step return
+    float step_ret = (value > 0.0f) ? (new_value - value) / value : 0.0f;
+    float new_cum_return = cum_return + step_ret;
+
+    // Update max equity for drawdown
+    float new_max = fmaxf(max_equity, new_value);
+
+    // Write portfolio state
+    portfolio_state[ps + 0] = new_value;
+    portfolio_state[ps + 1] = position;
+    portfolio_state[ps + 2] = cash;
+    portfolio_state[ps + 3] = entry_price;
+    portfolio_state[ps + 4] = new_max;
+    portfolio_state[ps + 5] = step_ret;       // step PnL (for reward)
+    portfolio_state[ps + 6] = new_cum_return;
+    portfolio_state[ps + 7] += 1.0f;          // step count
+
+    // Outputs
+    step_rewards[w] = step_ret;
+    step_returns[w * max_len + current_step] = step_ret;
+}