phase3(env-unification): val step_returns measure pure P&L (no shaping)

Root cause of the long-running "catastrophic train/val Sharpe gap":
backtest_env_kernel was subtracting behavioral shaping (inventory penalty,
churn penalty, opportunity cost) from step_returns BEFORE the metrics layer
computed Sharpe / Sortino / WinRate. Validation was reporting "P&L minus
shaping" as if it were realized P&L.

Both single-step and batched variants of backtest_env_step had the bug.

The shaping terms exist for a reason — they steer the training policy toward
risk-aware behavior. They belong in TRAINING reward, where they shape the
gradient. They do NOT belong in VALIDATION step_returns, which is the
measurement we use to judge whether the model would be profitable in
production. Production deployment doesn't pay an inventory penalty for
holding a position — it pays the actual market P&L of holding it.

Equivalent semantically to running experience_env_step with shaping_scale = 0
(the Phase 3 control scalar landed in commit 3f6eb006c).

Smoke-test verification (TD-propagation, RTX 3050 Ti, 20 epochs):

  metric                          before      after
  val_Sharpe range                -17 to -33   -1.24 to +2.34
  epochs val_Sharpe > 0           0 / 20       10 / 20
  Best (training) Sharpe          ~15-19       +21.04
  train Sharpe trajectory         unchanged    unchanged

The ~30-point Sharpe gap that motivated the entire env-unification effort
was ~80% measurement bug and ~20% legitimate train/val differences. The
remaining gap (val WinRate still anomalously low, 1.5–4.7% vs training
15–23%) suggests one more accounting issue in the val win-rate counter
but is non-blocking — Sharpe is now an honest production-equivalent
measurement.

Kernel signature kept stable (holding_cost_rate, churn_threshold,
churn_penalty_scale, opp_cost_scale args still present, suppressed via
(void) casts) so the Rust launch site does not need to change. Clean
deletion of those args is a follow-up after validation that no other
caller depends on the ABI.

Files touched:
  crates/ml/src/cuda_pipeline/backtest_env_kernel.cu  (-48 / +35)

Verified: SQLX_OFFLINE=true cargo check -p ml --lib --tests passes.
TD-propagation smoke test runs cleanly end-to-end (33s).

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

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;
-
+    (void)holding_cost_rate;
-    // ── Cost-driven hold timing (matches training kernel) ──────────────
+    (void)churn_threshold;
-    // Inventory penalty: continuous cost proportional to position size.
+    (void)churn_penalty_scale;
-    float inventory_penalty = holding_cost_rate * fabsf(position);
+    (void)opp_cost_scale;
-    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
-    }
 
     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;