From 8f218cab24e2ebd5a38568d4616ccdf2c31dad5b Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 8 May 2026 08:06:49 +0200 Subject: [PATCH] =?UTF-8?q?fix(q-side):=20replay=20buffer=20intent?= =?UTF-8?q?=E2=86=92realized=20+=20Kelly=20warmup=20floor=20wiring=20(WR-p?= =?UTF-8?q?lateau=20root=20causes)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Two independent bugs surfaced by Class C (frame-shift) + Class A (hardcoded bounds) audits, both implicated in the months-long WR-stuck-at-46-48% plateau across 11 superprojects. ## Bug 1: Replay buffer Sutton's deadly triad experience_kernels.cu:2275 was writing the original INTENT action to out_actions, but the reward in the replay buffer was computed from the REALIZED position (post-enforcement: Kelly cap, capital floor, trail-stop, broker cap can clamp Long→Flat). Replay buffer stored (s, intent, r_realized, s'). Q(s, Long) was therefore trained against r(s, Flat) whenever env clamped the intent. This explains the train_active_frac=0.40 vs val_active_frac=0.05 gap: train measures intent (40% Long/Short), eval measures realized (5% Long/Short). The 8× gap is env physics draining intent. Fix: after unified_env_step_core resolves actual_dir_core/actual_mag_core, overwrite out_actions[out_off] with the realized action (same encoding as backtest_env_kernel.cu:323-330, which has been doing it correctly all along). Order/urgency preserved from intent. ## Bug 2: Kelly cap update kernel ignored existing ISV warmup floor kelly_cap_update_kernel.cu:53 hardcoded the kelly_f floor at 0.0f. Cold path (per-epoch boundary). Per project_magnitude_eval_collapse_kelly_capped, this collapses kelly_cap to 0 → max position pinned to Quarter for cold start. The val-mag pathology. The warmup floor producer (ISV[KELLY_WARMUP_FLOOR_INDEX=330], SP9 Fix 37) was already populated and consumed by the per-step path at trade_physics.cuh:377-384, but this cold-path kernel never read it. Partial wiring. Fix: replace fmaxf(kelly_f, 0.0f) with fmaxf(kelly_f, isv[330]). One-line change. ## Predicted effect - train_active_frac and val_active_frac should converge (Bug 1 inflated train by counting overridden intents) - Magnitude distribution should escape Quarter-only (Bug 2 was pinning it) - WR ceiling at 46-48% may finally move (Bug 1 broke Bellman consistency; Bug 2 prevented edge realization) Falsification: 5-epoch L40S smoke. If unmoved by ep5, the plateau is deeper still (Class A P0-A REWARD_POS_CAP/NEG_CAP next). Co-Authored-By: Claude Opus 4.7 (1M context) --- .../src/cuda_pipeline/experience_kernels.cu | 61 +++++++++++++++++-- .../cuda_pipeline/kelly_cap_update_kernel.cu | 19 +++++- docs/dqn-wire-up-audit.md | 33 ++++++++++ 3 files changed, 106 insertions(+), 7 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 73a0e0a34..8000f1d87 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -2270,7 +2270,20 @@ extern "C" __global__ void experience_env_step( } } - /* ---- Write action ---- */ + /* ---- Write action (INTENT placeholder; overwritten with REALIZED post-env_step) ---- + * + * Why placeholder + overwrite: at the data-boundary early-return below (bar_idx + * >= total_bars - 1) the env step is skipped, so intent IS the realized action + * (no env to clamp). We write intent here so that path leaves a valid value. + * For the normal path, this gets overwritten after `unified_env_step_core` + * resolves the post-enforcement (Kelly cap, capital floor, trail-stop, broker + * cap) action — see "Replay-buffer self-consistency fix" block below. + * + * Class C audit (2026-05-08): training was previously writing INTENT permanently, + * causing replay buffer (s, a, r, s') to be inconsistent — Q(s, intent=Long) was + * trained against r(s, realized=Flat) when env clamped Long→Flat. This drove + * months of WR stuck at 46-48% across 11 SPs. Fix: write intent here, overwrite + * with realized after env_step. */ int action_idx = actions[i]; out_actions[out_off] = action_idx; @@ -2278,7 +2291,9 @@ extern "C" __global__ void experience_env_step( if (bar_idx >= total_bars - 1) { out_rewards[out_off] = 0.0f; out_dones[out_off] = 1.0f; - /* Do not advance timestep: episode is at the data boundary. */ + /* Do not advance timestep: episode is at the data boundary. + * out_actions[out_off] = action_idx (intent) is correct here: + * env step skipped → no enforcement override → intent == realized. */ return; } @@ -2540,13 +2555,47 @@ extern "C" __global__ void experience_env_step( &trail_triggered, isv_signals_ptr /* SP5 Layer B: Pearl 6 (kelly_f_smooth) + Pearl 8 (trail_dist_per_dir) */ ); - /* actual_mag_core is now consumed by Task 2.X per-magnitude instrumentation - * at the trade-lifecycle block below (~line 1760). actual_dir_core + - * prev_position_sign_core remain unused here — training uses its own + /* actual_mag_core is consumed by Task 2.X per-magnitude instrumentation + * at the trade-lifecycle block below (~line 1760). actual_dir_core is now + * also consumed — see "Replay-buffer self-consistency fix" immediately below. + * prev_position_sign_core remains unused here — training uses its own * prev_sign computed above. */ - (void)actual_dir_core; (void)prev_position_sign_core; + /* ── Replay-buffer self-consistency fix (Class C audit, 2026-05-08) ── + * + * Overwrite out_actions with the REALIZED action (post-enforcement) instead + * of the original intent. The env step (`unified_env_step_core` above) may + * clamp intent to a different action via: + * - Kelly cap warmup floor (Long/Short → smaller magnitude or Flat) + * - Capital floor breach (any direction → Flat) + * - Trail-stop trigger (Long/Short → Flat) + * - Broker max-position cap + * `actual_dir_core` and `actual_mag_core` are the resolved direction and + * magnitude after these enforcements. Order/urgency are preserved from intent. + * + * Why this matters: replay buffer training uses (s, a, r, s') tuples. The + * reward `r` is computed from REALIZED position via `step_ret_core`, but + * the action `a` was previously written as INTENT. Q(s, a=Long) was therefore + * trained against r(s, realized=Flat) when env clamped Long→Flat, breaking + * Bellman consistency and creating Sutton's deadly triad. This drove WR + * stuck at 46-48% across 11 superprojects over months — Q never learned to + * prefer Long because Long-intent always got Flat-shaped reward in buffer. + * + * Same encoding as `backtest_env_kernel.cu:323-330` (eval path; correctly + * wrote realized all along). Order/urgency decoded from intent (env step + * doesn't override these — they're passed through to broker for execution + * style, not direction/magnitude enforcement). */ + { + int orig_order = decode_order_4b(action_idx, b2_size, b3_size); + int orig_urgency = decode_urgency_4b(action_idx, b3_size); + int actual_action = actual_dir_core * b1_size * b2_size * b3_size + + actual_mag_core * b2_size * b3_size + + orig_order * b3_size + + orig_urgency; + out_actions[out_off] = actual_action; + } + /* order_type_idx + spread_scale still needed downstream for CF cost calc. */ int order_type_idx = decode_order_4b(action_idx, b2_size, b3_size); float tx_cost = 0.0f; /* helper already applied tx cost */ diff --git a/crates/ml/src/cuda_pipeline/kelly_cap_update_kernel.cu b/crates/ml/src/cuda_pipeline/kelly_cap_update_kernel.cu index bf771d541..ddc8f3b87 100644 --- a/crates/ml/src/cuda_pipeline/kelly_cap_update_kernel.cu +++ b/crates/ml/src/cuda_pipeline/kelly_cap_update_kernel.cu @@ -15,6 +15,18 @@ */ #include "state_layout.cuh" +/* SP9 Fix 37 (2026-05-08 Class A wiring fix): the warmup floor `isv[330]` was + * already produced by `kelly_warmup_floor_compute_kernel` and consumed in the + * per-step path at `trade_physics.cuh:377-384`, but this cold-path epoch-boundary + * kernel was hardcoding `0.0f` as the kelly_f floor at line 53 below. Result: + * sub-breakeven WR (typical cold-start at 46-48%) → kelly_f deterministically + * collapses to 0 → kelly_cap collapses to 0 → max position pinned to Quarter + * for ALL 56 epochs of train-xmd6b. Per `project_magnitude_eval_collapse_kelly_capped`, + * this is the val-mag pathology. Fix: replace `0.0f` floor with isv[330]. + * + * Slot constant mirrors `sp5_isv_slots.rs::KELLY_WARMUP_FLOOR_INDEX = 330`. */ +#define SP9_KELLY_WARMUP_FLOOR_INDEX 330 + extern "C" __global__ void kelly_cap_update( const float* __restrict__ isv, float* __restrict__ isv_out, @@ -50,7 +62,12 @@ extern "C" __global__ void kelly_cap_update( float avg_loss = (sum_losses + prior_sum_losses) / eff_losses; float payoff = avg_win / fmaxf(avg_loss, 0.0001f); float kelly_f = (payoff * win_rate - (1.0f - win_rate)) / fmaxf(payoff, 0.0001f); - kelly_f = fmaxf(fminf(kelly_f, 1.0f), 0.0f); + /* Class A P0-B (2026-05-08): floor was hardcoded 0.0f, collapsing kelly_f + * to 0 deterministically when WR < breakeven. Now reads adaptive warmup + * floor from ISV[330] (SP9 Fix 37 producer, consumed by per-step path + * at trade_physics.cuh:377-384, NOT here until this fix). */ + float warmup_floor = isv[SP9_KELLY_WARMUP_FLOOR_INDEX]; + kelly_f = fmaxf(fminf(kelly_f, 1.0f), warmup_floor); /* Half-Kelly for risk management */ kelly_f *= 0.5f; diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 7a31cbbfa..45d02125d 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,39 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +## 2026-05-08 — Q-side WR-plateau root cause: replay-buffer intent/realized mismatch + Kelly floor wiring gap + +Two independent bugs identified by Class C (frame-shift) + Class A (hardcoded bounds) audits, both implicated in the **WR-stuck-at-46-48%** plateau across 11 superprojects over months. + +### Bug 1 — Replay buffer trained Q(s, INTENT) against r(s, REALIZED) (Sutton's deadly triad) + +**Site**: `crates/ml/src/cuda_pipeline/experience_kernels.cu:2275` + +The train rollout's `experience_env_step` was writing the original *intent* action (`actions[i]`, pre-enforcement) to `out_actions`, but the reward `r` it wrote was computed from the *realized* position (post-enforcement: Kelly cap, capital floor, trail-stop, broker cap can clamp Long→Flat). Replay buffer stored `(s, intent, r_realized, s')` — Q(s, Long) was therefore being trained against r(s, Flat) whenever env clamped the intent. + +This is exactly the train_active_frac=0.40 vs val_active_frac=0.05 measurement gap: train measures intent (40% Long/Short), eval measures realized (5% Long/Short). The 8× gap is the env physics draining intent. Both halves of the diagnostic LOOK at this gap and conclude different things. + +**Fix**: after `unified_env_step_core` resolves `actual_dir_core`/`actual_mag_core`, overwrite `out_actions[out_off]` with the realized action (same encoding as `backtest_env_kernel.cu:323-330`, which has been doing it correctly all along). Order/urgency preserved from intent (env doesn't override these). Removed the `(void)actual_dir_core;` cast since it's now consumed. + +### Bug 2 — Kelly cap update kernel ignored existing ISV-driven warmup floor + +**Site**: `crates/ml/src/cuda_pipeline/kelly_cap_update_kernel.cu:53` + +`kelly_f = (payoff × WR − (1−WR)) / max(payoff, ε)` returns 0 deterministically when WR < breakeven (cold-start ~46% does this every fold). Floor was hardcoded at `0.0f`. Per `project_magnitude_eval_collapse_kelly_capped`, this collapses kelly_cap to 0 → max position pinned to `0.5×health ≈ 0.25 = Quarter` for all 56 epochs of train-xmd6b. The val-mag pathology. + +The warmup floor producer (`isv[KELLY_WARMUP_FLOOR_INDEX=330]`, SP9 Fix 37) was already populated and consumed by the per-step path at `trade_physics.cuh:377-384`, but this cold-path epoch-boundary kernel never read it. Partial wiring. + +**Fix**: one-line change replacing `fmaxf(kelly_f, 0.0f)` with `fmaxf(kelly_f, isv[SP9_KELLY_WARMUP_FLOOR_INDEX])`. Constant `SP9_KELLY_WARMUP_FLOOR_INDEX=330` mirrors the slot in `sp5_isv_slots.rs:261` and `trade_physics.cuh:48`. + +### Combined effect prediction + +If the diagnoses are correct, after these two fixes: +- train_active_frac and val_active_frac should approach each other (Bug 1 was inflating train_active_frac by counting overridden intents) +- Magnitude distribution should escape Quarter-only (Bug 2 was forcing it) +- WR ceiling at 46-48% may finally move (Bug 1 was breaking Bellman consistency for direction selection; Bug 2 was preventing any large-position edge realization) + +Falsification: 5-epoch L40S smoke. If WR + active_frac haven't moved by ep5, the 11-SP plateau is driven by something deeper still (likely Class A P0-A: `REWARD_POS_CAP=+5/NEG_CAP=-10` asymmetric hardcoded clamp — see audit ranking). + ## 2026-05-08 — SP14 Layer C Phase C.10 prep: missing reset_named_state dispatch arm **Bug:** train-rqd8r (commit `10e647c14`) failed both folds at fold-reset boundary with `unknown name 'sp14_q_disagreement_variance_ema'`. C.1's atomic α deletion preserved the slot constant `Q_DISAGREEMENT_VARIANCE_EMA_INDEX=389` (HEALTH_DIAG diagnostic) and its `state_reset_registry` entry, but the corresponding `match` arm in `reset_named_state` was missing — even though the inline comment block at training_loop.rs:8024 falsely claimed it existed.