fix(sp11): B1b bug 3 — cf-component feedback loop in mag-ratio canary

Deep audit on b435d25be found a third post-B1b bug in same class as
the slot 63 overload: experience_env_step writes cf_reward_weighted
(= w_cf × cf_reward, post-controller-weight) to reward_components_
per_sample[+1] at line ~3696. SP4's reward_component_ema_kernel EMAs
this into ISV slot 64 (REWARD_CF_EMA_INDEX) which the SP11 mag-ratio
canary reads as ratio[1] = slot[64] / Σ.

Self-reinforcing loop:
  high w_cf → high cf_reward_weighted → high slot 64 →
  high ratio[1] → controller raises w_cf → tighter loop

Until mean=1 normalization saturates other components to floor.

The other 5 component slots correctly write RAW pre-weight values:
- rc[+0] = total_reward (intentional, PopArt input — post-composition)
- rc[+2..+5] = r_trail / r_micro / r_opp_cost / r_bonus (all raw, pre-Σ)
- Slot 360 (popart-component) fed by `popart_component_per_sample`
  which receives `r_popart` raw

Only rc[+1] was wrong. Fixed: write raw cf_reward to rc[+1] so the
canary tracks intrinsic cf magnitude. The replay-buffer cf-tuple
reward (out_rewards[cf_off]) still uses cf_reward_weighted — that's
correct, loss kernels train on controller-weighted signal.

This was the third bug in a class — pre-SP11 invariants exposed by
post-decomposition semantic. Trio: (1) slot 63 overload (5e16b67ca),
(2) stale rc[] init + cf_flip ordering (b435d25be), (3) cf-component
feedback loop (this commit).

cargo check + build clean; 6/6 SP11 GPU tests + 14/14 contract tests
still pass. After this fix-up, all known SP11 reward-system bugs
identified by deep audit are resolved. L40S smoke validates empirical
sharpe recovery.

crates/ml/src/cuda_pipeline/experience_kernels.cu

@@ -3691,9 +3691,24 @@ extern "C" __global__ void experience_env_step(
         /* C.2: component [1] = counterfactual reward magnitude for this sample.
          * Written at out_off (on-policy slot), not cf_off, so the EMA kernel's
          * [N*L, 6] reduction over out_off slots captures it alongside the other
-         * on-policy components. Stores the post-weight value so the EMA
-         * tracks the same magnitude the loss kernels actually train on. */
-        reward_components_per_sample[out_off * 6 + 1] = cf_reward_weighted;
+         * on-policy components.
+         *
+         * Writes the RAW intrinsic cf_reward, NOT cf_reward_weighted. The
+         * weighted value is still used at out_rewards[cf_off] above so the
+         * loss kernels train on the controller-weighted signal — that's
+         * correct. But this slot feeds the SP11 mag-ratio canary's cf axis
+         * via SP4 reward_component_ema_kernel → ISV slot 64
+         * (REWARD_CF_EMA_INDEX), and the canary must observe the
+         * INTRINSIC component magnitude, not the controller's own output.
+         * Otherwise: high w_cf → high cf_reward_weighted → high slot 64 →
+         * high ratio[1] → controller raises w_cf → self-reinforcing loop
+         * until mean=1 normalization saturates the other 5 components to
+         * floor. Match the convention of the other component slots:
+         * rc[+2..+5] all write raw r_<component> pre-Σ, and slot 360
+         * (popart-component) is fed by raw r_popart via
+         * popart_component_per_sample. Only rc[+0] is intentionally
+         * post-composition (PopArt input). */
+        reward_components_per_sample[out_off * 6 + 1] = cf_reward;
     }
 
     /* ---- Advance timestep counter ---- */

docs/dqn-wire-up-audit.md

@@ -5301,3 +5301,53 @@ 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.
+
+### SP11 B1b bug-hunt fix-up #2 (2026-05-04, on top of `b435d25be`)
+
+Deep-audit pass on the post-B1b kernel found a third bug in the same
+class as the slot 63 overload (5e16b67ca) and the stale-rc[]+cf_flip
+pair (b435d25be): pre-SP11 invariants exposed by post-decomposition
+semantic. Triad complete with this fix.
+
+Bug 3 (Critical) — cf-component feedback loop in mag-ratio canary:
+`experience_env_step` wrote `cf_reward_weighted = w_cf × cf_reward`
+(POST-controller-weight) to `reward_components_per_sample[+1]` at
+the cf block (line ~3696 pre-fix). SP4's
+`reward_component_ema_kernel` EMAs this into ISV slot 64
+(`REWARD_CF_EMA_INDEX`), which the SP11 mag-ratio canary reads as
+`ratio[1] = slot[64] / Σ`. Self-reinforcing loop:
+
+  high w_cf → high cf_reward_weighted → high slot 64 →
+  high ratio[1] → controller raises w_cf → tighter loop
+
+…until the mean=1 normalization saturates the other 5 components to
+floor. The other 5 component slots correctly write RAW pre-weight
+values:
+- `rc[+0]` = `total_reward` (intentional — PopArt input,
+  post-composition)
+- `rc[+2..+5]` = `r_trail / r_micro / r_opp_cost / r_bonus`
+  (all raw, pre-Σ)
+- Slot 360 (popart-component) is fed by raw `r_popart` via
+  `popart_component_per_sample`.
+
+Only `rc[+1]` was wrong. Fix: write raw `cf_reward` to `rc[+1]` so
+the canary tracks intrinsic cf magnitude. The replay-buffer cf-tuple
+reward (`out_rewards[cf_off] = cf_reward_weighted`) is unchanged —
+loss kernels still train on the controller-weighted reward, which is
+correct; only the canary's input changes.
+
+Verification (RTX 3050 Ti):
+- `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
+- `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
+
+After this fix-up, all three known SP11 reward-system bugs surfaced
+by deep audit are resolved (slot 63 overload, stale rc[] init,
+cf_flip ordering, cf-component feedback loop). L40S smoke validates
+the empirical sharpe recovery.