revert: remove reward normalization + magnitude gradient restore

Both destabilized H100 training:
- Reward normalization (÷pos_frac): collapsed gradients at epoch 2-3
- Magnitude gradient restore (beta*MSE): Q-value explosion at epoch 25
  (tested beta=0.10, 0.05, 0.02 — ALL cause Q explosion >50)

The magnitude MSE gradient is structurally unstable post-warmup.
ANY non-zero beta feeds a Q-overestimation loop that CQL can't counter.
The magnitude branch head learns during MSE warmup (epochs 1-5), then
the trunk continues improving via IQN (60% budget). This is the ONLY
stable configuration proven on H100 (RUN 4: stable through 9 epochs,
grad_norm=0.44, Q=2.03).

Reverts 76478559b (reward norm) and 22b7bc083 (gradient restore).
Keeps 1540c0287 (CUDA Graph fixes — the critical root cause).

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

crates/ml/src/cuda_pipeline/experience_kernels.cu

@@ -1421,8 +1421,7 @@ 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 use RAW equity-normalized return (not magnitude-normalized).
-         * Kelly measures actual returns for position sizing — needs true scale. */
+        /* Kelly statistics (unchanged from v6) */
         if (exiting_trade) {
             if (segment_return > 0.0f) {
                 win_count += 1.0f;
@@ -1438,30 +1437,6 @@ 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) {

crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs

@@ -4573,26 +4573,6 @@ impl GpuDqnTrainer {
                     .arg(&adv_ptr).arg(&adv_mse_ptr)
                     .arg(&scale_mse).arg(&n_adv)
                     .launch(cfg_adv).map_err(|e| MLError::ModelError(format!("blend mse adv: {e}")))?;
-
-                // ── Restore full MSE gradient for magnitude branch (d==1) ────
-                // The alpha blend above gave magnitude: (1-alpha)*MSE_grad
-                // (because C51_grad is zeroed for d==1). At alpha=1.0, this is ZERO.
-                // Magnitude needs full MSE gradient always — it's the ONLY training
-                // signal for the magnitude branch head (C51 zeroed, IQN trunk-only).
-                // Fix: SAXPY d_adv[d==1] += alpha * d_adv_mse[d==1] to restore
-                // the alpha*MSE portion that was removed by the (1-alpha) scale.
-                if alpha > 1e-6 {
-                    let mag_offset = (b * b0 * na) as u64 * 4; // bytes: b0 branch comes first
-                    let n_mag = (b * b1 * na) as i32;
-                    let mag_blocks = ((b * b1 * na + 255) / 256) as u32;
-                    let cfg_mag = LaunchConfig { grid_dim: (mag_blocks, 1, 1), block_dim: (256, 1, 1), shared_mem_bytes: 0 };
-                    let mag_adv_ptr = adv_ptr + mag_offset;
-                    let mag_mse_ptr = adv_mse_ptr + mag_offset;
-                    self.stream.launch_builder(&self.saxpy_f32_kernel)
-                        .arg(&mag_adv_ptr).arg(&mag_mse_ptr)
-                        .arg(&alpha).arg(&n_mag)
-                        .launch(cfg_mag).map_err(|e| MLError::ModelError(format!("restore mag mse: {e}")))?;
-                }
             }
         }
 

test_data/futures-baseline/00b346049e49a136b0fe649c3a0f4e49db4a695b15a76f9a28ba032cac3e93a0.norm_stats.json

@@ -0,0 +1,90 @@
+{
+  "mean": [
+    -1.5753783410475906e-6,
+    0.00015586458233923567,
+    -0.0001602226610409705,
+    2.6252690595741797e-7,
+    0.23841375572607446,
+    0.5006594472619454,
+    -0.14025914465123132,
+    0.004148928364065063,
+    0.7734723809594998,
+    0.8080563347546004,
+    2.6252690595741797e-7,
+    1.8379052470046046e-6,
+    -1.5753783410475906e-6,
+    0.06705760455136711,
+    0.07124305868450649,
+    -0.00009607003103403483,
+    -0.058209735129103,
+    0.13092338921053032,
+    0.0021655944675779393,
+    -0.0019227122071590423,
+    0.09566432933362572,
+    0.1528548103323881,
+    0.5068672252326588,
+    0.3432918637676055,
+    0.3300639323722491,
+    0.5068513131249612,
+    0.4931486868750401,
+    0.011071998668295478,
+    0.01219140812646202,
+    0.6216440027638934,
+    0.6673510439809522,
+    -0.5260726329047711,
+    0.0013139637674245343,
+    -0.0004938218591171064,
+    -0.33302277745951275,
+    -0.6539613943795362,
+    -0.9232688159360032,
+    -0.7168433678765663,
+    0.31642407821366536,
+    0.8185449985393084,
+    0.06258714301096246,
+    -0.00019736072450139626
+  ],
+  "std": [
+    2.435468244032963,
+    2.435424311620899,
+    2.4355093178271687,
+    2.4354647373410425,
+    0.1391085693432111,
+    0.10677570816817682,
+    0.7786624019131647,
+    0.5104368912046735,
+    0.40732640106624063,
+    0.17846192346988113,
+    2.4354647373410425,
+    0.0005818855392673685,
+    2.435468244032963,
+    0.2904432575934772,
+    0.2903293458531898,
+    0.006964157057777831,
+    0.8903249601344784,
+    0.337316550684594,
+    0.0501037659817767,
+    0.050409221333257176,
+    0.6042820535545856,
+    0.3605113712831862,
+    0.2845791873762313,
+    0.24413043338060986,
+    0.4702358269148673,
+    0.4420676129790175,
+    0.4420676129790175,
+    1.019288177963432,
+    0.9710412756563976,
+    0.46074711703678956,
+    0.44981448297335824,
+    0.05306725197147951,
+    0.08276483960102538,
+    0.07766493953874622,
+    1.0142738579878101,
+    1.5320938155185901,
+    1.7245047881189186,
+    0.9469493748572908,
+    1.8734588386298756,
+    1.9146394511311065,
+    0.02358036785123252,
+    0.4524640966211648
+  ]
+}