diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 97a645b99..65fc38a2e 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/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 { diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 3085a6c36..a1107a5d1 100644 --- a/docs/dqn-wire-up-audit.md +++ b/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 |