From b435d25becb9cf2cf26b6f7e22b880b4faad9848 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Mon, 4 May 2026 10:51:55 +0200 Subject: [PATCH] =?UTF-8?q?fix(sp11):=20B1b=20bug-hunt=20fix-up=20?= =?UTF-8?q?=E2=80=94=20stale=20rc[]=20init=20+=20cf=5Fflip=20ordering?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Bug-hunt review on 5e16b67ca found two real bugs in experience_env_step post-B1b structural refactor — same class as the slot 63 overload (pre-SP11 invariants exposed by post-decomposition semantic). Bug 1 (Critical): reward_components_per_sample[+1..+5] not zero-init at per-bar entry. Slots written only on execution paths that fire (rc[2] inside segment_complete trail-trigger, rc[3] inside positioned-non- complete, rc[4] inside flat-with-features, rc[5] inside multiple bonus paths, rc[1] inside CF block). Non-firing paths leave stale values from previous batch's same out_off, contaminating SP4 reward_component_ema_kernel -> slots 64..68 -> SP11 mag-ratio canary -> wrong controller weights. Fix: zero-init rc[0..6) at the same location where r_ locals are zero-initialized, so locals and buffer reset together. Eliminates the entire stale-rc class of bugs. Bug 2 (Important): cf_flip applied BEFORE inventory/churn/conviction modifiers. Spec section 3.4.4 requires cf_flip LAST so subtractive penalties operate on the right sign and multiplicative scales attenuate before direction is flipped. Pre-fix on flipped samples, inventory/churn ADDED to negated reward (penalty became bonus) and conviction scaled the wrong-signed value. Asymmetric gradient signals between flipped/non-flipped -> contaminated SP11 mag-ratio canary at the cf axis. Fix: move cf_flip to after conviction_scale, making it the last modifier before out_rewards write. CF-block contract preserved (reward at cf_off is still the post-flipped final value; CF block already does `do_flip ? -reward : reward` to recover unflipped base). Stale doc-strings: state_reset_registry.rs (2 sites) and training_loop.rs (1 site) referenced pre-fix-up SP5_PRODUCER_COUNT=186 / wiener_buffer=771. Updated to formula form `(71 + SP5_PRODUCER_COUNT) × 3` so they don't drift on the next SP; current value (post-B1b fix-up slot 360) is SP5_PRODUCER_COUNT=187, buffer=774 floats. dqn-wire-up-audit.md: appended SP11 B1b bug-hunt fix-up section documenting both bugs, the structural fix, the doc-string sweep, and verification (cargo check/build/tests). cargo check + build clean; 6/6 SP11 GPU tests + 10/10 sp5_isv_slots contract tests + 4/4 state_reset_registry tests still pass. The two bugs were latent in the local smoke (sharpe 3 vs B1a 30 likely traceable to either or both); structural fix sound. L40S smoke on this commit will validate the empirical recovery. Co-Authored-By: Claude Opus 4.7 (1M context) --- .../src/cuda_pipeline/experience_kernels.cu | 64 ++++++++++++++--- .../src/trainers/dqn/state_reset_registry.rs | 13 ++-- .../src/trainers/dqn/trainer/training_loop.rs | 10 ++- docs/dqn-wire-up-audit.md | 69 +++++++++++++++++++ 4 files changed, 141 insertions(+), 15 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 0502a6d53..4132ae13c 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -2629,6 +2629,41 @@ extern "C" __global__ void experience_env_step( float r_opp_cost = 0.0f; float r_bonus = 0.0f; + /* Zero-init per-component slots so non-firing paths don't leave stale + * data from the previous batch's `out_off` contaminating the SP4 reward- + * component EMAs (slots 64..68) and the SP11 mag-ratio canary that reads + * them. The local `r_*` variables are zero-initialized above; the backing + * buffer needs the same treatment so non-firing paths produce a true zero + * observation rather than stale prior-bar values. + * + * Slot population is conditional and sparse: + * rc[1] = cf_reward_weighted — written ONLY inside the CF block (~L3648) + * rc[2] = r_trail — written ONLY inside the segment_complete + * trail-trigger path (~L2730) + * rc[3] = r_micro — written ONLY inside positioned-non- + * complete paths (~L3006, L3016) + * rc[4] = r_opp_cost — written ONLY inside flat-with-features + * path (~L3121) + * rc[5] = bonus_* — overwrite at ~L3169, multiple `+=` / + * `-=` sites elsewhere + * + * Without this zero-init, a bar that doesn't fire the relevant path + * retains values from the previous batch's same `out_off` → reward- + * component EMAs (slots 64..68) accumulate stale magnitudes → SP11 mag- + * ratio canary reads contaminated ratios → controller emits wrong + * weights. Same class as the slot 63 overload bug fixed via B1b fix-up: + * pre-SP11 invariants exposed by post-decomposition semantic. + * + * Note: rc[+0] (PopArt input = total reward) is unconditionally + * overwritten at ~L3447, so it doesn't strictly need zero-init here, but + * including it removes ordering dependency for future readers. */ + reward_components_per_sample[out_off * 6 + 0] = 0.0f; + reward_components_per_sample[out_off * 6 + 1] = 0.0f; + reward_components_per_sample[out_off * 6 + 2] = 0.0f; + reward_components_per_sample[out_off * 6 + 3] = 0.0f; + reward_components_per_sample[out_off * 6 + 4] = 0.0f; + reward_components_per_sample[out_off * 6 + 5] = 0.0f; + /* Trade-cumulative running scalar — tracks the in-progress reward * across the if/else cascade. C.4/D.4b bonus blocks read this to * Q-cap their `|reward|` multiplicand (the existing Q-bounded @@ -3395,16 +3430,20 @@ extern "C" __global__ void experience_env_step( * the replay buffer. This prevents gradient explosion from propagating. */ if (isnan(reward) || isinf(reward)) reward = 0.0f; - /* G7: Counterfactual reward flip — negate reward when features were flipped. - * The model learns that flipped features -> flipped reward, enforcing - * directional symmetry. */ - if (do_flip) { - reward *= -1.0f; - } - /* ── Cost-driven hold timing (replaces min_hold_bars) ────────────── * The temporal attention learns optimal hold timing from these signals. - * No hardcoded gates — the model sees costs and learns to minimize them. */ + * No hardcoded gates — the model sees costs and learns to minimize them. + * + * Modifier ordering (SP11 spec §3.4.4): + * clamp_capital → -drawdown → -inventory → -churn → ×conviction → ×cf_flip + * + * cf_flip MUST be the LAST modifier so subtractive penalties operate on + * the correct sign and the multiplicative conviction scale attenuates + * before the direction is flipped. Pre-fix, on flipped samples + * inventory/churn ADDED to the negated reward (penalty became bonus) + * and conviction scaled the wrong-signed value — asymmetric gradient + * signals between flipped and non-flipped samples contaminated the + * SP11 mag-ratio canary at the cf axis. */ /* Inventory penalty: continuous cost proportional to position size. * Teaches: small positions cheap to hold, large positions expensive. @@ -3433,6 +3472,15 @@ extern "C" __global__ void experience_env_step( reward *= conviction; } + /* G7: Counterfactual reward flip — negate reward when features were flipped. + * The model learns that flipped features -> flipped reward, enforcing + * directional symmetry. Per spec §3.4.4 this is the LAST modifier so all + * additive penalties and the multiplicative conviction operate on the + * correctly-signed reward before the direction is inverted. */ + if (do_flip) { + reward *= -1.0f; + } + out_rewards[out_off] = reward; out_dones[out_off] = (float)done; diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 05f967423..ca518095b 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -534,7 +534,7 @@ impl StateResetRegistry { RegistryEntry { name: "sp4_wiener_state", category: ResetCategory::FoldReset, - description: "GpuDqnTrainer.wiener_state_buf [(71 + SP5_PRODUCER_COUNT) × 3 floats; post-SP11: (71 + 186) × 3 = 771] = (SP4 + SP5/SP6/SP7/SP8/SP9/SP10/SP11 producers) × {sample_var, diff_var, x_lag}] — SP4 Pearl D Wiener-EMA state for SP4 block [0..213) (40 SP4 + 29 Task-A13 retrofit + 2 Layer-C-#260 producers). Mapped-pinned f32; bulk write_bytes(0) at fold boundary covers the full buffer [0..(71 + SP5_PRODUCER_COUNT) × 3) so Pearl A's sentinel branch fires on the new fold's first producer launch for ALL SP4 and SP5+ producers (`prev_x_mean == 0 AND state.x_lag == 0` → first-observation replacement). Companion to the 11 SP4 ISV bound entries above — both halves of the sentinel contract reset together per `feedback_no_partial_refactor.md`. Buffer size derives from SP5_PRODUCER_COUNT to avoid drift on each subsequent SP.", + description: "GpuDqnTrainer.wiener_state_buf [(71 + SP5_PRODUCER_COUNT) × 3 floats; total = (SP4 + SP5/SP6/SP7/SP8/SP9/SP10/SP11 producers) × {sample_var, diff_var, x_lag}] — SP4 Pearl D Wiener-EMA state for SP4 block [0..213) (40 SP4 + 29 Task-A13 retrofit + 2 Layer-C-#260 producers). Mapped-pinned f32; bulk write_bytes(0) at fold boundary covers the full buffer [0..(71 + SP5_PRODUCER_COUNT) × 3) so Pearl A's sentinel branch fires on the new fold's first producer launch for ALL SP4 and SP5+ producers (`prev_x_mean == 0 AND state.x_lag == 0` → first-observation replacement). Companion to the 11 SP4 ISV bound entries above — both halves of the sentinel contract reset together per `feedback_no_partial_refactor.md`. Buffer size derives from SP5_PRODUCER_COUNT (formula form, not a hardcoded literal) to avoid drift on each subsequent SP.", }, RegistryEntry { name: "sp4_clamp_engage_counters", @@ -903,14 +903,19 @@ impl StateResetRegistry { // SP5 Task A1: Wiener-state companion reset. The wiener_state_buf // covers SP4 (SP4_PRODUCER_COUNT=71 producers × 3 = 213 floats) + // SP5 (SP5_PRODUCER_COUNT × 3 floats) = (71 + SP5_PRODUCER_COUNT) × 3 - // floats total. With SP5_PRODUCER_COUNT=186 (post-SP11 Fix 39) the - // buffer is 771 floats; the SP5 triples start at offset 213. + // floats total. The SP5 triples start at offset 213 and span + // SP5_PRODUCER_COUNT × 3 floats; the runtime constant lives in + // `cuda_pipeline::sp5_isv_slots::SP5_PRODUCER_COUNT` and grows as + // each subsequent SP adds producers (B1b fix-up landed slot 360 → + // SP5_PRODUCER_COUNT=187 → buffer = 774 floats, SP5 block + // [213..774)). Always derive sizes via the formula + // `(71 + SP5_PRODUCER_COUNT) × 3` rather than committing the literal. // Both halves of the Pearl A sentinel contract must reset together per // `feedback_no_partial_refactor.md`. RegistryEntry { name: "sp5_wiener_state", category: ResetCategory::FoldReset, - description: "GpuDqnTrainer.wiener_state_buf SP5 block at offsets [SP4_PRODUCER_COUNT × 3 .. (SP4_PRODUCER_COUNT + SP5_PRODUCER_COUNT) × 3) = [213..(213 + SP5_PRODUCER_COUNT × 3)). The total buffer is sized via `SP5_WIENER_TOTAL_FLOATS = (SP4_PRODUCER_COUNT=71 + SP5_PRODUCER_COUNT) × 3` floats; with SP5_PRODUCER_COUNT=186 (current value, post-SP11 Fix 39) the buffer is 771 floats and the SP5 block is [213..771). Each producer triple = {sample_var, diff_var, x_lag}. SP5 Pearl D Wiener-EMA state for ALL SP5 producers (Pearls 1, 2, 3, 4, 5, Layer D D1/D2/D3, SP7 T1, SP7 activation-flag fix, SP8 Fix 36, SP9 Kelly cold-start, SP10 eval Thompson temperature, SP11 reward-subsystem controller). Bulk write_bytes(0) at fold boundary covers the full wiener_state_buf [0..SP5_WIENER_TOTAL_FLOATS) so Pearl A's sentinel branch fires on the new fold's first producer launch. Pearl 6 (slots 280..286) does NOT use this buffer — Pearl 6 is cross-fold-persistent and uses in-kernel EWMA + cumulative-via-max(); the wiener slots [525..543) that Pearl 6 would map to under naive offset arithmetic (213 + (280-174)*3 = 525) are intentionally unused. Both SP4 and SP5 [non-Pearl-6] blocks reset atomically.", + description: "GpuDqnTrainer.wiener_state_buf SP5 block at offsets [SP4_PRODUCER_COUNT × 3 .. (SP4_PRODUCER_COUNT + SP5_PRODUCER_COUNT) × 3) = [213..(213 + SP5_PRODUCER_COUNT × 3)). The total buffer is sized via `SP5_WIENER_TOTAL_FLOATS = (SP4_PRODUCER_COUNT=71 + SP5_PRODUCER_COUNT) × 3` floats; refer to the runtime `SP5_PRODUCER_COUNT` (`cuda_pipeline::sp5_isv_slots`) for the current value rather than hardcoding a literal here — it grows on each SP follow-up (e.g. B1b fix-up landed slot 360 → SP5_PRODUCER_COUNT=187 → buffer=774 floats, SP5 block [213..774)). Each producer triple = {sample_var, diff_var, x_lag}. SP5 Pearl D Wiener-EMA state for ALL SP5 producers (Pearls 1, 2, 3, 4, 5, Layer D D1/D2/D3, SP7 T1, SP7 activation-flag fix, SP8 Fix 36, SP9 Kelly cold-start, SP10 eval Thompson temperature, SP11 reward-subsystem controller). Bulk write_bytes(0) at fold boundary covers the full wiener_state_buf [0..SP5_WIENER_TOTAL_FLOATS) so Pearl A's sentinel branch fires on the new fold's first producer launch. Pearl 6 (slots 280..286) does NOT use this buffer — Pearl 6 is cross-fold-persistent and uses in-kernel EWMA + cumulative-via-max(); the wiener slots [525..543) that Pearl 6 would map to under naive offset arithmetic (213 + (280-174)*3 = 525) are intentionally unused. Both SP4 and SP5 [non-Pearl-6] blocks reset atomically.", }, // SP5 Task A6: Pearl 6 cross-fold-persistent Kelly cap signals. // ISV[280..286) are EXPLICITLY EXEMPT from this registry. diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index f0fe81b0c..f590e36b6 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -6906,9 +6906,13 @@ impl DQNTrainer { // No-op: `reset_sp4_wiener_state` already bulk-zeros the // FULL `wiener_state_buf [0..(71 + SP5_PRODUCER_COUNT) × 3)` // which includes the SP5 producer block at offsets - // [71 × 3 .. (71 + SP5_PRODUCER_COUNT) × 3). Post-SP11 - // (SP5_PRODUCER_COUNT = 186): buffer is 771 floats, SP5 - // block at [213..771). The registry entry exists for + // [71 × 3 .. (71 + SP5_PRODUCER_COUNT) × 3). Buffer size + // is parameterised on the runtime `SP5_PRODUCER_COUNT` + // constant (`cuda_pipeline::sp5_isv_slots`) and grows as + // each SP follow-up adds producers — do NOT hardcode a + // float count here, it will go stale (B1b fix-up landed + // slot 360 → SP5_PRODUCER_COUNT=187 → buffer = 774 floats, + // SP5 block [213..774)). The registry entry exists for // documentation symmetry with the SP4 entry; the actual // byte-zero is performed by the SP4 arm. Pearl 6 (slots // 280..286) is intentionally exempt — its wiener-state diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 97d3e3677..3100dc615 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -5232,3 +5232,72 @@ update needed. B1c will pass `STATE_DIM_PADDED as i32` for the Pre-requisite for B1b (structural reward composition) and B1c (replay-time curiosity wiring against `trainer.states_buf`). + +### SP11 B1b bug-hunt fix-up (2026-05-04, on top of `5e16b67ca`) + +Bug-hunt review on the post-B1b kernel found two real bugs in +`experience_env_step` reward composition. Both are the same class +as the slot 63 overload bug fixed via the slot-360 follow-up: +pre-SP11 invariants exposed by post-decomposition semantic. Both +fixed atomically per `feedback_no_partial_refactor.md`. + +Bug 1 (Critical) — stale `reward_components_per_sample[+1..+5]`: +the per-component slots `rc[1..5]` are written only on execution +paths that fire (`rc[1]` inside the CF block, `rc[2]` inside the +`segment_complete` trail-trigger, `rc[3]` inside positioned-non- +complete, `rc[4]` inside flat-with-features, `rc[5]` inside multiple +bonus paths). Bars whose execution path doesn't fire retained values +from the previous batch's same `out_off`, contaminating the SP4 +`reward_component_ema_kernel` writes at ISV[64..68) and the SP11 +mag-ratio canary that reads them — the controller emitted the wrong +weights as a result. Fix: zero-init `rc[0..6)` at per-bar entry +(line ~2660), at the same site where the `r_` locals +are zero-initialized so locals + buffer reset together. Component +0 is unconditionally overwritten downstream so it doesn't strictly +need the zero-init, but including it eliminates ordering dependency +for future readers. + +Bug 2 (Important) — cf_flip ordering: per spec section 3.4.4 the +modifier order is +`clamp_capital -> -drawdown -> -inventory -> -churn -> ×conviction +-> ×cf_flip`. Pre-fix-up the `if (do_flip) reward *= -1.0f;` block +ran at the NaN-guard site BEFORE inventory/churn/conviction. On +flipped samples this caused inventory/churn (subtractive penalties) +to ADD to the negated reward (penalty became bonus) and conviction +to scale the wrong-signed value — asymmetric gradient signal between +flipped and non-flipped samples in the same batch, contaminating the +SP11 mag-ratio canary at the cf axis. Fix: move the cf_flip block +to AFTER the conviction multiplicand, making it the last modifier +before the `out_rewards` write. The CF block downstream still reads +the post-flipped final `reward` and uses `do_flip ? -reward : reward` +to recover the unflipped base — the CF contract is preserved, only +the underlying composition sequence changed (and that's the intended +fix). + +Stale doc-strings updated to formula form to prevent next-SP drift: +- `crates/ml/src/trainers/dqn/state_reset_registry.rs` lines 537 + and 905-915 (sp4_wiener_state, sp5_wiener_state SP5 block comment, + sp5_wiener_state RegistryEntry description). +- `crates/ml/src/trainers/dqn/trainer/training_loop.rs` lines + 6905-6917 (sp5_wiener_state dispatch arm comment). + +All three sites previously hardcoded `SP5_PRODUCER_COUNT=186 / buffer +=771`; post-B1b fix-up the values are 187/774, but the doc-strings +should never have committed the literal — they now reference the +runtime constant via the formula `(71 + SP5_PRODUCER_COUNT) × 3`. + +Verification: +- `SQLX_OFFLINE=true cargo check -p ml --lib` — clean +- `SQLX_OFFLINE=true cargo build -p ml --release` — clean +- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test + sp11_producer_unit_tests --features cuda -- --ignored` — 6/6 GPU + oracle tests pass on RTX 3050 Ti +- `SQLX_OFFLINE=true cargo test -p ml --lib sp5_isv_slots` — + 10/10 contract tests pass +- `SQLX_OFFLINE=true cargo test -p ml --lib state_reset_registry` — + 4/4 contract tests pass + +L40S smoke on this commit will validate the empirical sharpe-recovery +(local B1b smoke regressed from B1a sharpe ~30 to ~3; both bugs above +are plausible root causes). After this fix-up + L40S verification, +B1c (replay-time curiosity wiring) is unblocked.