diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu
index 2699bd3cc..3badeaa65 100644
--- a/crates/ml/src/cuda_pipeline/experience_kernels.cu
+++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu
@@ -646,9 +646,10 @@ extern "C" __global__ void experience_env_step(
     }
 
     /* ---- 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.
-     * This teaches the model that trading in choppy markets is MORE expensive. */
+    /* Real spread widens in volatile markets. CUSUM direction (feature[41])
+     * is used here as a market-stress proxy for spread scaling — this is
+     * CORRECT (high |CUSUM| = trending = wider spreads). NOT used as vol
+     * proxy for reward normalization — ATR(14) at feature[9] handles that. */
     float cusum_raw = 0.0f;
     if (features != NULL && bar_idx < total_bars && market_dim > 41) {
         cusum_raw = features[(long long)bar_idx * market_dim + 41];
@@ -795,7 +796,8 @@ extern "C" __global__ void experience_env_step(
      * Only blocks flat exits, NOT reversals — reversals always allowed.
      * CRITICAL: fix the stored action to match what actually executed.
      * Without this, the replay buffer stores (state, action=Flat, reward_from_Hold)
-     * which corrupts Q-values for Flat — classic action aliasing bug. */
+     * which corrupts Q-values for Flat — classic action aliasing bug.
+     * TODO(hyperopt): min_hold_bars=5 is hardcoded — candidate for hyperopt parameter. */
     if (prev_sign != 0 && curr_sign == 0 && hold_time < 5.0f && hold_time > 0.0f && !exiting_trade) {
         position = ps[0];
         cash = ps[1];
@@ -891,35 +893,42 @@ extern "C" __global__ void experience_env_step(
             segment_return = reversal_return;
         }
 
-        /* Vol normalization: CUSUM proxy for realized volatility.
-         * Makes reward comparable across trending (wide moves) and
-         * ranging (tight moves) regimes. Target: unit-variance returns. */
-        float vol_proxy = fmaxf(cusum_raw * 0.01f, 0.0001f);
+        /* Vol normalization: ATR(14) proxy for realized volatility.
+         * Feature[9] stores safe_normalize(ln(atr), -7.0, 9.0) = (ln(atr)+7)/16.
+         * Reverse: ln(atr) = feature[9]*16 - 7, atr = exp(ln(atr)).
+         * atr_pct = atr / price gives per-bar percentage volatility.
+         * Makes reward comparable across trending and ranging regimes. */
+        float atr_norm = 0.0f;
+        if (features != NULL && bar_idx < total_bars && market_dim > 9) {
+            atr_norm = features[(long long)bar_idx * market_dim + 9];
+        }
+        float log_atr = atr_norm * 16.0f - 7.0f;
+        float atr_pct = expf(log_atr) / fmaxf(raw_close, 1.0f);
+        float vol_proxy = fmaxf(atr_pct, 0.0001f);
         float vol_norm = vol_proxy * sqrtf(fmaxf(segment_hold_time, 1.0f));
         float vol_normalized_return = segment_return / vol_norm;
 
-        /* Scale to learnable magnitude: target reward range [-3, +3].
+        /* Scale to learnable magnitude: target reward range ~[-10, +10].
          * With Xavier init weight noise ~0.001, a reward of 1.0 has SNR=1000.
          * ETDQN showed this range produces stable C51 learning. */
         reward = 10.0f * vol_normalized_return;
 
-        /* Clamp extreme outliers (prevent single-trade Q-value explosion) */
-        reward = fmaxf(-10.0f, fminf(10.0f, reward));
+        /* Loss aversion BEFORE squash (prospect theory): negative rewards
+         * weighted 1.5x. Applied pre-squash so the asymmetry is smooth —
+         * post-squash would create [-15, +10] hard asymmetry. */
+        if (reward < 0.0f) {
+            reward *= loss_aversion;
+        }
+
+        /* Soft squash: tanh maps [-inf,+inf] -> [-10,+10] smoothly.
+         * Preserves tail information (5% win > 1% win) unlike hard clamp
+         * which maps both to 10.0. Gradient never exactly zero. */
+        reward = 10.0f * tanhf(reward / 10.0f);
     }
 
-    /* ---- Turnover penalty: position changes are expensive ---- */
-    /* Teaches the model that CHANGING position costs money.
-     * This is IN ADDITION to the tx cost already deducted from cash.
-     * The model learns: "unless the edge justifies the cost, stay put." */
-    if (delta != 0.0f) {
-        float turnover = fabsf(delta) / fmaxf(max_position, 1.0f);
-        reward -= 0.05f * turnover;
-    }
-
-    /* Loss aversion: negative rewards weighted 1.5x (prospect theory) */
-    if (reward < 0.0f) {
-        reward *= loss_aversion;
-    }
+    /* Turnover penalty REMOVED (reward v6): tx costs already deducted from
+     * cash at line ~679 (Almgren-Chriss impact model). The old 0.05*|delta|
+     * penalty double-counted costs and over-penalized necessary rebalancing. */
 
     /* ---- Portfolio value for floor check ---- */
     float new_portfolio_value_pre_floor = cash + position * raw_next;