fix(sp11): symmetric reward cap — losses were unbounded

experience_kernels.cu:2788:
  float capped_pnl = fminf(base_reward, 10.0f);
                          ^^^^^^^^^^^^^ caps profits, NOT losses

Diagnostic instrumentation in smoke-test-k9drh on commit 774d7552a
captured the asymmetry empirically:
  ep1: r_popart min=-9186, max=+10
  ep2: r_popart min=-79089, max=+10 (growing)

`base_reward = 2.0f * vol_normalized_return` and
`vol_normalized_return = segment_return / vol_norm` where
`segment_return` has no structural lower bound (signed P&L). The
unilateral `fminf(base_reward, 10.0f)` capped the upper tail only,
so a single large adverse segment_return produced an arbitrarily
negative `capped_pnl` → r_popart → r_weighted →
reward_components[+0] → slot 63 (PopArt input EMA), inflating
C51/IQN/Bellman normalization scale and breaking Q-target
consistency across epochs. Empirical fingerprint matched the
within-fold sharpe degradation observed in smoke-test-gwfn8 on
commit fd24b5383 (10→4 within F0, 9→2.5 within F1).

Spec semantic: reward bounded in [-10, +10]. Fix:
  float capped_pnl = fmaxf(-10.0f, fminf(base_reward, 10.0f));

Audit findings (per task §2): all related fminf/fmaxf clamps in
experience_kernels.cu reviewed. Reward modifier chain (3260-3500)
verified bounded once r_popart is bilateral. Other bilateral
clamps already correct (515-517, 2174, 2191, 2227, 3300, 3452,
3729, 3763, 5076, 5210, 5213, 5461, 5535, 6353, 6693, 6694).
Intentional asymmetries verified at 1078, 1082-1085, 2564-2565,
2873, 2949, 4574, 5896 (each documented in the audit doc with
the structural reason the lower side is unbounded).

Latent finding flagged separately (NOT fixed here — feature-side
requires consumer audit per feedback_no_partial_refactor):
plan_isv[PNL_VS_TARGET] at line 850 and plan_isv[PNL_VS_STOP]
at line 853 are upper-clamped at 2.0 but lower-unbounded. These
feed assemble_state as policy features (not reward components),
so out of scope of this reward-chain fix. Mirrored in
backtest_plan_kernel.cu:164,168 (same pattern). Tracking as
follow-up.

Per pearl_bounded_modifier_outputs_require_structural_activation:
spec-bounded values require BILATERAL structural enforcement.
This bug was the asymmetric counterpart to the conviction sigmoid
(which IS correctly bounded structurally).

Diagnostic instrumentation (commit 774d7552a) NOT removed in this
commit — will be removed in a follow-up after the symmetric-cap
smoke validates the fix on L40S.

docs/dqn-wire-up-audit.md updated with Resolution (2026-05-04)
section reflecting root cause, fix, audit follow-through, and the
latent plan_isv finding (per Invariant 7).

Build: SQLX_OFFLINE=true cargo check -p ml --lib — clean.
Test: trainers::dqn::trainer::tests::test_reward_function_price_changes
      passes. (PPO test_reward_computation pre-existing failure on
      HEAD 774d7552a, unrelated — verified via stash+rerun.)

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

crates/ml/src/cuda_pipeline/experience_kernels.cu

@@ -2785,7 +2785,18 @@ extern "C" __global__ void experience_env_step(
          * `exiting_trade=1` upstream (line ~2371), so this branch fires
          * for both — the segregation matters only for component attribution. */
         float base_reward = 2.0f * vol_normalized_return;
-        float capped_pnl = fminf(base_reward, 10.0f);
+        /* SP11 fix (2026-05-04): symmetric cap. The previous unilateral
+         * `fminf(base_reward, 10.0f)` capped profits at +10 but left losses
+         * unbounded — vol_normalized_return has no lower bound, so a single
+         * large adverse segment_return propagated as a large negative reward
+         * straight through the modifier chain into reward_components[+0],
+         * inflating PopArt input EMA (slot 63) and breaking C51/IQN/Bellman
+         * Q-target normalization. Empirical fingerprint from
+         * smoke-test-k9drh @ 774d7552a: ep1 r_popart min=-9186, ep2
+         * min=-79089 (growing), max stuck at +10. Spec semantic is bilateral
+         * [-10, +10]; enforce bilaterally per pearl_bounded_modifier_outputs_
+         * require_structural_activation. */
+        float capped_pnl = fmaxf(-10.0f, fminf(base_reward, 10.0f));
         if (trail_triggered) {
             r_trail = capped_pnl;
         } else {

docs/dqn-wire-up-audit.md

@@ -5585,7 +5585,11 @@ Sharpe degrades within fold despite the SP11 controller working
 correctly.
 
 All known multiplicative modifiers are structurally bounded:
-- `r_popart = capped_pnl ≤ 10` (experience_kernels.cu:2744)
+- `r_popart = capped_pnl ∈ [-10, +10]` (experience_kernels.cu:2799,
+  bilateral as of the SP11 symmetric-cap fix — see "Resolution
+  (2026-05-04)" below). Pre-fix: `fminf(base_reward, 10.0f)` at
+  line 2788 only capped the upper tail; losses were unbounded —
+  that WAS the inflater this diagnostic was searching for.
 - `× conviction ∈ (0, 1)` via sigmoid at experience_kernels.cu:7579
 - `× cf_flip ∈ {-1, +1}`
 - `× shaping_scale ∈ [0, 1]`
@@ -5597,6 +5601,58 @@ inflater, this commit adds GPU-side instrumentation (data-gathering
 only, no fix) capturing the running `reward` cascade variable + per-
 component locals at every checkpoint.
 
+### Resolution (2026-05-04)
+
+The diagnostic instrumentation in `smoke-test-k9drh @ 774d7552a`
+captured the asymmetry directly: `ep1 r_popart min=-9186, max=+10`
+and `ep2 min=-79089, max=+10` (max stuck, min growing each epoch).
+Root cause: line 2788 `float capped_pnl = fminf(base_reward, 10.0f)`
+— unilateral upper cap; `base_reward = 2.0f * vol_normalized_return`
+where `vol_normalized_return = segment_return / vol_norm` has no
+structural lower bound (signed P&L). A single large adverse
+`segment_return` produced an arbitrarily negative `capped_pnl` →
+`r_popart` → `r_weighted` → `reward_components[+0]` → slot 63
+(PopArt input EMA), inflating C51/IQN/Bellman normalization scale
+and breaking Q-target consistency across epochs.
+
+Fix:
+```c
+float capped_pnl = fmaxf(-10.0f, fminf(base_reward, 10.0f));
+```
+Spec semantic was always bilateral `[-10, +10]`; the implementation
+was unilateral. Per
+`pearl_bounded_modifier_outputs_require_structural_activation`:
+spec-bounded values require BILATERAL structural enforcement.
+
+Audit follow-through (per `feedback_no_partial_refactor`): all
+other `fminf(.,K)` and `fmaxf(.,K)` clamps in
+`experience_kernels.cu` reviewed for the same asymmetric-clamp
+pattern. Findings:
+- Reward modifier chain (3260-3500): all bounded modifiers operate
+  on the post-cap `r_weighted`; once `capped_pnl` is bilateral, the
+  whole downstream chain is bounded.
+- Bilateral clamps verified correct at: 515-517 (cost bands),
+  2174/2191/2227 (conviction/Kelly), 3300 (stability),
+  3729/3763 (cf_reward), 3452 (capital-floor), 5076/5210/5213/
+  5461/5535 ([0,1] gates), 6353 (d_h0), 6693/6694 (ADX/CUSUM).
+- Intentional asymmetries verified at: 1078 (cosine_progress ≥ 0
+  structurally), 1082-1085 (eps×mult ≥ 0 structurally + EPS_FLOOR
+  at 1090-1094), 2564-2565 (max_dd lower-unbounded by drawdown
+  semantic), 2873/2949 (fabsf-based unit ≥ 0 structurally), 4574
+  (guarded by `if best_pnl > 0.0f`), 5896 (input-gate underflow
+  toward 0 is correct sigmoid-like semantic).
+- Latent finding flagged separately (NOT fixed in this commit —
+  feature-side, requires consumer audit): `plan_isv[PNL_VS_TARGET]`
+  at 850 and `plan_isv[PNL_VS_STOP]` at 853 are upper-clamped at
+  2.0 but lower-unbounded; mirrored in `backtest_plan_kernel.cu`.
+  These feed `assemble_state` (line 956) as policy features, not
+  reward components, so out of scope for this reward-chain fix.
+
+Diagnostic instrumentation (commit 774d7552a) NOT removed in this
+commit — kept for the symmetric-cap validation smoke on L40S.
+Removal lands in the follow-up commit after smoke validates the
+fix end-to-end.
+
 ### What it adds
 
 | Component | Detail |