diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index dc49df0ee..f16a0691a 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -1696,7 +1696,19 @@ extern "C" __global__ void experience_env_step( * Kernel reads isv_signals_ptr[trade_attempt_rate_ema_idx] and * isv_signals_ptr[trade_target_rate_idx] at Flat→Positioned transitions. */ int trade_attempt_rate_ema_idx, - int trade_target_rate_idx + int trade_target_rate_idx, + /* C.4 / D.4b bonus optimism break (2026-04-29): direction-branch + * Q-scale EMA slot index (Q_DIR_ABS_REF_INDEX = 21). The C.4 timing + * bonus and D.4b regime penalty multiply by |final_pnl| / |reward| + * which is itself the trade-cumulative shaped reward — across trades + * this compounds (bonus inflates Q which inflates next-trade reward + * which inflates bonus, …). Replacing the unbounded multiplicand + * with a Q-scale-bounded variant breaks the loop. ISV[21] is a + * gradient-decoupled EMA of the direction-branch max(|Q_mean|), + * already produced by q_stats_kernel.cu — no new producer needed. + * Per pearl_one_unbounded_signal_per_reward.md + + * feedback_isv_for_adaptive_bounds.md. */ + int q_dir_abs_ref_idx ) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i >= N) return; @@ -2518,9 +2530,23 @@ extern "C" __global__ void experience_env_step( const float peak_bar = ps[PS_PEAK_PNL_BAR]; const float bars_early = fmaxf(0.0f, segment_hold_time - peak_bar); const float final_pnl = reward; + /* pearl_one_unbounded_signal_per_reward + feedback_isv_for_adaptive_bounds + * (2026-04-29): |final_pnl| compounds across trades through Layer 4-9 + * credits, creating a multi-trade optimism loop (bonus → Q → reward → + * bonus). Bound it to the direction-branch Q-scale EMA + * (ISV[Q_DIR_ABS_REF_INDEX=21]) so the term's unbounded multiplicand + * is the stable, gradient-decoupled Q-scale, not the trade-cumulative + * shaped reward. ISV[21] is already produced by q_stats_kernel.cu. + * NOTE: ISV slot index passed as kernel arg (no magic numbers per + * feedback_no_quickfixes / feedback_isv_for_adaptive_bounds). */ + const float pnl_norm = (isv_signals_ptr != NULL) + ? isv_signals_ptr[q_dir_abs_ref_idx] + : 0.0f; + const float pnl_unit = fabsf(final_pnl) / fmaxf(pnl_norm, 1e-6f); + const float pnl_capped = fminf(pnl_unit, 1.0f) * pnl_norm; /* ≤ Q-scale */ const float timing_bonus = shaping_scale * (bars_early / segment_hold_time) - * fabsf(final_pnl) + * pnl_capped * conviction_core; reward += timing_bonus; /* C.2 attribution — accumulate into rc[5] bonus slot (see above). */ @@ -2559,12 +2585,27 @@ extern "C" __global__ void experience_env_step( * absolute Q-magnitude) AND |reward| simultaneously — produced * penalties 5–50× larger than the reward, flipping trade outcomes * wholesale and destabilising training (smoke showed Return swings - * −400% / +170% across epochs). Fix: pick exactly ONE unbounded - * scaling signal per reward term — here, |reward| itself. + * −400% / +170% across epochs). The follow-on fix used |reward| as + * the single unbounded multiplicand — but |reward| at this point IS + * the trade-cumulative shaped reward (Layer 2 + Layer 4-9 credits + + * C.4 timing bonus + D.4a persistence bonus accumulated above). Across + * trades this compounds: penalty inflates with reward inflation, + * which the Bellman target then bootstraps off, which inflates the + * NEXT trade's reward, etc. (Plan C Phase 2 smoke diagnosis 2026-04-29: + * bonus EMA hit 235 by F2 ep2 in the a52d99613 baseline as well — + * pre-existing structural pathology.) + * + * Final fix: bound |reward| by ISV[Q_DIR_ABS_REF_INDEX=21], the + * direction-branch Q-scale EMA. The unbounded multiplicand becomes + * the stable, gradient-decoupled Q-scale rather than trade-cumulative + * shaping output. Same surgical pattern as C.4 above — pearl + * pearl_one_unbounded_signal_per_reward.md + + * feedback_isv_for_adaptive_bounds.md. * * bars_late ∈ [0, hold_time]. When shift fires late in the trade, penalty * is small; when shift fires early and we held until exit, penalty is - * up to |reward| — flipping a positive reward to zero in the worst case. + * up to ISV[21] — bounded by the Q-scale instead of the (compounding) + * trade-cumulative reward. * * Attribution: subtracts from rc[5] via -= (co-exists with the three positive * contributions B.2/C.4/D.4a at other (i,t) slots; cancellation within the @@ -2572,8 +2613,13 @@ extern "C" __global__ void experience_env_step( if (segment_complete && ps[PS_REGIME_SHIFT_BAR] > 0.0f && saved_hold_time > 0.0f) { const float bars_late_frac = fmaxf(0.0f, fminf(1.0f, ((float)saved_hold_time - ps[PS_REGIME_SHIFT_BAR]) / (float)saved_hold_time)); + const float reward_norm = (isv_signals_ptr != NULL) + ? isv_signals_ptr[q_dir_abs_ref_idx] + : 0.0f; + const float reward_unit = fabsf(reward) / fmaxf(reward_norm, 1e-6f); + const float reward_capped = fminf(reward_unit, 1.0f) * reward_norm; const float penalty = shaping_scale * conviction_core - * bars_late_frac * fabsf(reward); + * bars_late_frac * reward_capped; reward -= penalty; reward_components_per_sample[out_off * 6 + 5] -= penalty; } diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index b79722a64..4eddae72f 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -3797,6 +3797,18 @@ impl GpuExperienceCollector { use crate::cuda_pipeline::gpu_dqn_trainer::TRADE_TARGET_RATE_INDEX; TRADE_TARGET_RATE_INDEX as i32 }) + // C.4 / D.4b bonus optimism break (2026-04-29): + // direction-branch Q-scale EMA slot index. The C.4 timing + // bonus and D.4b regime penalty bound their |reward|/|pnl| + // multiplicands by ISV[Q_DIR_ABS_REF_INDEX=21] to break + // the cross-trade compounding optimism loop. ISV[21] is + // produced by q_stats_kernel.cu — no new producer needed. + // Per pearl_one_unbounded_signal_per_reward.md + + // feedback_isv_for_adaptive_bounds.md. + .arg(&{ + use crate::cuda_pipeline::gpu_dqn_trainer::Q_DIR_ABS_REF_INDEX; + Q_DIR_ABS_REF_INDEX as i32 + }) .launch(launch_cfg) .map_err(|e| MLError::ModelError(format!( "experience_env_step t={t}: {e}" diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index bf16e1874..4e7c3256f 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2192,3 +2192,5 @@ Predicted impact: Plan C fold 0 ep2 with `q_mean(t)=0.82, q_mean(t-1)=-0.018, IS Per `pearl_adaptive_moe_lambda.md` — canonical "EMA-tracked diagnostic drives a controller" pattern: kernel + ISV slot + bootstrap (0.0 cold-start) + reset (0.0 fold boundary) + observability. Per `pearl_cold_path_no_exception_to_gpu_drives.md` — even cold-path scalar arithmetic stays on GPU when its reduction inputs already live on GPU (ISV[16,21] both produced by `q_stats_kernel.cu`). Per `feedback_adaptive_not_tuned.md` — tau decay is ISV-driven, not a constant. Per `feedback_isv_for_adaptive_bounds.md` — the `4.0` upper clip and `1e-6` denom floor are structural bounds (drift saturation matching the kill-criterion's 3.0× ratio threshold; numerical-stability floor) carried in ISV slot semantics. Plan C A.2 wire-up (2026-04-29, second commit): `tau_update_kernel.cu`, `state_reset_registry.rs`, `training_loop.rs::reset_named_state` + per-epoch launch site, and `trainer/mod.rs::reset_for_fold` comment all updated in lockstep with the kernel+slot landed by the previous commit. Producer launch lives between the existing Q-drift kill criterion and the `prev_epoch_q_mean = q_mean` update — both prev and curr q_means are in scope, and the `prev_epoch_q_mean.abs() > 1e-6` cold-start gate matches the kill check's gate. Tau consumer applies `1/(1+clip(ISV[129],0,4))` AFTER the existing `[tau_final, tau_base]` clamp so dampening can drive tau below tau_final on runaway drift (intent: tau_final is the design floor for healthy runs, but a runaway requires further suppression). Defensive `fminf(...,4.0f)` in the consumer guards against producer-side bugs even though the producer already clips (Invariant 1 carve-out, belt-and-braces). + +Plan C Phase 2 follow-up (2026-04-29): bonus reward optimism-coupling break in `experience_kernels.cu`. The C.4 timing bonus (~line 2521) and D.4b regime penalty (~line 2581) both used the trade-cumulative `|final_pnl|` / `|reward|` as their unbounded multiplicand. By the time bonus shaping runs in `experience_env_step`, `reward` is the Layer 2 vol-normalised return PLUS Layer 4-9 credits PLUS the C.4 timing bonus PLUS the D.4a persistence bonus — i.e. it is itself a function of all prior shaping inflation. Across trades this compounds into a multi-trade optimism feedback loop: bonus inflates → Q-target inflates (Bellman bootstraps off shaped reward) → next trade's reward inflates → bonus inflates further. Researcher reproduction (a25f669e9df953174, 2026-04-29) found the a52d99613 baseline ALSO hits bonus EMA=235 by F2 ep2 — the pathology is structural and pre-dates Plan C Phase 2. Per `pearl_one_unbounded_signal_per_reward.md` the rule is exactly ONE unbounded multiplicand per reward term, and per `feedback_isv_for_adaptive_bounds.md` adaptive bounds live in the ISV signal bus. **Surgical fix**: replace `|final_pnl|` (C.4) and `|reward|` (D.4b) with `min(|x|/ISV[Q_DIR_ABS_REF_INDEX=21], 1) × ISV[21]` — the multiplicand is now the direction-branch Q-scale EMA (already produced gradient-decoupled by `q_stats_kernel.cu` since Plan 1) rather than the trade-cumulative shaping output. Cap is adaptive (matches Q magnitude as it evolves) and breaks cross-trade compounding; |reward| can still drive the cap to its max (= ISV[21]) so directional information is preserved, just bounded. **Kernel signature**: `experience_env_step` gains a final `int q_dir_abs_ref_idx` parameter (after the existing `trade_target_rate_idx`); single Rust call site `gpu_experience_collector.rs:3800` passes `Q_DIR_ABS_REF_INDEX as i32`. `experience_action_select` not affected (it never read `|reward|`/`|final_pnl|` as a multiplicand). **Other shaping sites NOT touched** because they were already bounded — D.4a persist (multiplicand `drawdown_depth ≤ ~0.1` gates the term, `tanh(reward/dd)` is bounded), B.2 novelty (vol_proxy ≤ 0.01 gates, kl_amp ≤ 1.0), D.4c stable (vol_proxy gates + bounded stability/conviction). Per `feedback_no_partial_refactor.md` consistency: both fixed sites use the SAME ISV[21] bound and the SAME `(unit, capped)` two-step formula — symmetric migration. Predicted impact: bonus EMA O(100-256) → O(0.05-2.0); rc[5] → Bellman-target optimism loop broken; Plan C smoke F0 ep2 Q-drift kill should no longer fire on this pathway; a52d99613 baseline gets the same fix automatically (validates pre-existing pathology, not Plan-C-specific). Layout fingerprint unchanged — kernel signature change but no ISV slot add and no param-tensor change.