From 5bdeb6c6d6ddecabba7ebd6f065055a8f78c2beb Mon Sep 17 00:00:00 2001
From: jgrusewski <jgrusewski@localhost>
Date: Thu, 9 Apr 2026 12:33:03 +0200
Subject: [PATCH] =?UTF-8?q?fix:=20restore=20reward=20normalization=20?=
 =?UTF-8?q?=E2=80=94=20was=20present=20in=20the=20only=20stable=20H100=20r?=
 =?UTF-8?q?un?=
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

RUN 4 (1540c0287) survived 9 epochs with stable grad_norm=0.44.
That commit INCLUDED reward normalization (÷pos_frac). The revert
in 09fb80baa removed it, and the resulting H100 run collapsed at
epoch 7 (grad_norm=0.000).

The reward normalization is part of the stable configuration.
experience_kernels.cu restored to exact RUN 4 state.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
---
 .../src/cuda_pipeline/experience_kernels.cu   | 27 ++++++++++++++++++-
 1 file changed, 26 insertions(+), 1 deletion(-)

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) {