diff --git a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu
index c858f76dd..7fa24cab7 100644
--- a/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu
+++ b/crates/ml/src/cuda_pipeline/backtest_env_kernel.cu
@@ -294,33 +294,29 @@ extern "C" __global__ void backtest_env_step(
         return;
     }
 
-    // Step return
+    // Step return — PURE per-bar P&L (Phase 3 unified-env policy).
+    //
+    // Validation reports what *deployment* would experience: the actual
+    // change in portfolio equity. Behavioral shaping (inventory penalty,
+    // churn penalty, opportunity cost) belongs in TRAINING reward only —
+    // those terms steer the policy but should not be subtracted from the
+    // measurement we use to judge whether the policy is profitable in
+    // production. Equivalent to running experience_env_step with
+    // shaping_scale = 0.
+    //
+    // Removed (2026-04-21):
+    //   inventory_penalty = holding_cost_rate * |position|     // shaping
+    //   churn_penalty     = scale * (threshold - hold_time)/.. // shaping
+    //   opp_cost penalty  (was opp_cost_scale * is_flat * 0)    // dead, scale=0
+    // These are still suppressed via the unused-args contract — the kernel
+    // signature retains the parameters for ABI stability but ignores them.
     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;
-        }
-    }
-
-    // Opportunity cost: penalize flat when model predicts edge (matches training kernel).
-    // opp_cost_scale = 0.0 during backtest evaluation (no Q-gap available post-hoc).
-    {
-        float is_flat = (fabsf(position) < 0.001f) ? 1.0f : 0.0f;
-        step_ret -= 0.0f * opp_cost_scale * is_flat;  // q_gap not available in backtest
-    }
+    (void)holding_cost_rate;
+    (void)churn_threshold;
+    (void)churn_penalty_scale;
+    (void)opp_cost_scale;
 
     float new_cum_return = cum_return + step_ret;
     float new_max = fmaxf(max_equity, new_value);
@@ -416,6 +412,13 @@ extern "C" __global__ void backtest_env_step_batch(
     float*       kelly_stats,          /* [n_windows * KELLY_STATS_SIZE=4] win_count, loss_count, sum_wins, sum_losses */
     const float* __restrict__ isv_signals  /* [13] pinned; [12]=health. NULL → 0.5 fallback. */
 ) {
+    /* Phase 3 unified-env: shaping params retained for ABI but no longer
+     * applied to step_returns. See single-step variant header for rationale. */
+    (void)holding_cost_rate;
+    (void)churn_threshold;
+    (void)churn_penalty_scale;
+    (void)opp_cost_scale;
+
     int w = blockIdx.x * blockDim.x + threadIdx.x;
     if (w >= n_windows) return;
 
@@ -591,33 +594,17 @@ extern "C" __global__ void backtest_env_step_batch(
             return; /* Exit kernel — episode over */
         }
 
-        /* Step return */
+        /* Step return — PURE per-bar P&L (Phase 3 unified-env policy).
+         *
+         * Validation reports deployment-equivalent P&L. Behavioral shaping
+         * (inventory, churn, opportunity cost) belongs in TRAINING reward
+         * only. The single-step variant above documents the same policy.
+         *
+         * Removed (2026-04-21): inv_pen, churn_pen, opp_cost penalty.
+         * Args retained for ABI stability; suppressed via (void) below. */
         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;
-            }
-        }
-
-        /* Opportunity cost: penalize flat when model predicts edge (matches training kernel).
-         * opp_cost_scale = 0.0 during backtest evaluation (no Q-gap available post-hoc). */
-        {
-            float is_flat = (fabsf(position) < 0.001f) ? 1.0f : 0.0f;
-            step_ret -= 0.0f * opp_cost_scale * is_flat;  /* q_gap not available in backtest */
-        }
-
         cum_return = cum_return + step_ret;
         float new_max = fmaxf(max_equity, new_value);
         max_equity = new_max;