diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu
index 30f453b65..c47939a37 100644
--- a/crates/ml/src/cuda_pipeline/experience_kernels.cu
+++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu
@@ -1421,7 +1421,8 @@ extern "C" __global__ void experience_env_step(
         float segment_pnl = __bfloat162float(ps[11]) + raw_pnl - trade_start_pnl;
         float segment_return = segment_pnl / (prev_equity > 1.0f ? prev_equity : 1.0f);
 
-        /* Kelly statistics (unchanged from v6) */
+        /* Kelly statistics use RAW equity-normalized return (not magnitude-normalized).
+         * Kelly measures actual returns for position sizing — needs true scale. */
         if (exiting_trade) {
             if (segment_return > 0.0f) {
                 win_count += 1.0f;
@@ -1437,6 +1438,30 @@ extern "C" __global__ void experience_env_step(
             segment_return = reversal_return;
         }
 
+        /* ── Magnitude-neutral reward normalization ──────────────────────
+         * Divide return by position fraction (|pos|/max_pos) so the reward
+         * measures DIRECTIONAL SKILL PER UNIT OF RISK, not absolute PnL.
+         *
+         * Without this, Quarter (0.25) produces 4× lower PnL variance than
+         * Full (1.00) for the same price move. The MSE Bellman target for
+         * Quarter has 16× lower variance (squared) → gradient is 16× more
+         * consistent → Quarter Q converges first → Boltzmann locks in →
+         * model never learns Full's true Q. Same mechanism as C51 collapse
+         * but through reward variance instead of distributional shape.
+         *
+         * After normalization, all magnitudes produce the SAME reward
+         * variance for the same directional skill. The model then selects
+         * magnitude based on risk-adjusted return, not convergence speed.
+         *
+         * Applied AFTER Kelly (which needs raw returns) and AFTER reversal
+         * override (which also needs magnitude normalization). */
+        {
+            float pos_frac = fabsf(position) / fmaxf(max_position, 0.001f);
+            if (pos_frac > 0.01f) {
+                segment_return /= pos_frac;
+            }
+        }
+
         /* Vol normalization (unchanged from v6) */
         float atr_norm = 0.0f;
         if (features != NULL && bar_idx < total_bars && market_dim > 9) {