fix(sp11): B1b bug-hunt fix-up — stale rc[] init + cf_flip ordering

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_<component>
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) <noreply@anthropic.com>

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;
 

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.

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

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_<component>` 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.