fix(sp11): B0 — controller renorm Σ=1 → mean=1 (post-A2 spec amendment)

Per spec §3.4.3 amended at 7ddaf9c51 on main: A2's Σweights=1
renormalization caused 6× reward-magnitude collapse on mutually-
exclusive components in experience_env_step (popart/micro/opp_cost
paths fire at most one per bar; weight 1/6 per fired path averages
1/6 of pre-SP11 reward magnitude).

Amended: weights normalize to mean(weights) = 1 (i.e., Σ = N = 6).
Each weight in [WEIGHT_HARD_FLOOR=0.01, MAX_WEIGHT=3.0]. Default
uniform = 1.0 each. Preserves pre-SP11 absolute scale on average.

Code change: reward_subsystem_controller_kernel.cu renormalization
step changes from `weights[c] = blends[c] / blend_sum` to
`weights[c] = min(MAX_WEIGHT, blends[c] × N / blend_sum)`. Anchors
N_COMPONENTS=6.0f and MAX_WEIGHT=3.0f added to the Invariant-1 const
float block at the top of the kernel.

3 A2 controller unit-test assertions updated:
- z_score_at_zero: weight_sum 1.0→6.0; per-component 1/6→1.0
- weights_renormalize_after_floor: assertion strengthened to
  weight_sum ≤ N (cap binds in this pathological test where pre-cap
  dominant weight ≈ 4.18 > MAX_WEIGHT=3.0); added per-component
  ≤ MAX_WEIGHT envelope check; added explicit cap-binding assertion
  on dominant weight.
- saboteur_post_clamp_holds_min: weight assertions unaffected (this
  test asserts only on s[6]/s[7], saboteur+curiosity are independent
  of mean-vs-Σ choice).

Audit doc updated: docs/dqn-wire-up-audit.md gets a new
"SP11 B0 — controller renorm Σ=1 → mean=1 (2026-05-04)" section.

cargo check + release build clean. 6/6 SP11 GPU oracle tests pass on
RTX 3050 Ti. sp5_isv_slots (10/10) + state_reset_registry (4/4)
contract tests still pass.

Pre-requisite for B1b structural reward-composition refactor (§3.5.3)
which depends on the mean=1 semantic.

crates/ml/src/cuda_pipeline/reward_subsystem_controller_kernel.cu

@@ -18,9 +18,15 @@
 //     saturation, collapsing the controller's effective range.
 //   - improving = sigmoid(dz);  stagnant_or_worse = 1 - improving.
 //   - Per-component blend (winner × improving + diversifier × stagnant)
-//     pre-floor, then floor-clamped, then RENORMALIZED to Σ=1. Renormalization
-//     is non-negotiable: PopArt's reward-magnitude EMA assumes a consistent
-//     reward-scale; sum(weights) drifting > 1 inflates running |reward|.
+//     pre-floor, then floor-clamped, then RENORMALIZED to mean(weights) = 1
+//     (i.e., Σ = N = 6). Renormalization is non-negotiable: PopArt's reward-
+//     magnitude EMA assumes a consistent reward-scale; sum(weights) drifting
+//     above N inflates running |reward|. Mean=1 (amended 2026-05-04 per spec
+//     §3.4.3) preserves pre-SP11 absolute reward magnitude on average across
+//     the mutually-exclusive component paths in `experience_env_step` (popart
+//     / micro / opp_cost fire at most one per bar), which the original Σ=1
+//     form scaled down 6× silently. Per-component upper cap MAX_WEIGHT=3.0
+//     prevents any single component from dominating after renormalization.
 //   - Curiosity uses the permanent-floor pattern (max(), NOT blend):
 //     `pressure = max(stagnant × bound, CURIOSITY_PERMANENT_FRACTION × bound)`.
 //     Per `pearl_blend_formulas_must_have_permanent_floor.md` — blend
@@ -67,10 +73,12 @@ extern "C" __global__ void reward_subsystem_controller_kernel(
     const float SABOTEUR_MIN                  = 0.5f;
     const float SABOTEUR_MAX                  = 2.0f;
     const float WEIGHT_HARD_FLOOR             = 0.01f;
+    const float MAX_WEIGHT                    = 3.0f;
     const float CURIOSITY_PERMANENT_FRACTION  = 0.2f;
     const float ENGAGEMENT_FLOOR              = 0.1f;
     const float CURIOSITY_BOUND_FRACTION      = 0.3f;
     const float WEIGHT_FLOOR_FRACTION         = 0.5f;
+    const float N_COMPONENTS                  = 6.0f;
 
     /* Block-shared scratch — block-loaded once by threads 0..5, then read by
      * all 10 threads after a single __syncthreads barrier. */
@@ -148,11 +156,17 @@ extern "C" __global__ void reward_subsystem_controller_kernel(
     }
     __syncthreads();
 
-    /* Threads 0..5: renormalize so Σweights = 1 exactly. The fmaxf with
-     * EPS_DIV guards the (algebraically impossible) all-floor-zero case. */
+    /* Threads 0..5: renormalize so mean(weights) = 1 exactly (i.e.,
+     * Σweights = N_COMPONENTS = 6). Amended 2026-05-04 per spec §3.4.3 —
+     * the original Σ=1 form caused 6× reward-magnitude collapse on the
+     * mutually-exclusive component paths in `experience_env_step`
+     * (popart / micro / opp_cost fire at most one per bar). Each weight
+     * is then capped at MAX_WEIGHT=3.0 to prevent any single component
+     * from dominating after renormalization. The fmaxf with EPS_DIV
+     * guards the (algebraically impossible) all-floor-zero case. */
     if (threadIdx.x < 6) {
-        scratch_out[threadIdx.x] = blends[threadIdx.x]
-                                 / fmaxf(blend_sum, EPS_DIV);
+        const float scale = N_COMPONENTS / fmaxf(blend_sum, EPS_DIV);
+        scratch_out[threadIdx.x] = fminf(MAX_WEIGHT, blends[threadIdx.x] * scale);
     }
 
     /* Thread 6: curiosity_pressure — permanent-floor pattern (max(), NOT

crates/ml/tests/sp11_producer_unit_tests.rs

@@ -365,7 +365,8 @@ fn launch_controller_kernel(
 ///
 /// Per-component winner_weight = 1/6, diversifier = (1 - 1/6)/5 = 1/6;
 /// blend = 0.5 × 1/6 + 0.5 × 1/6 = 1/6 each. After floor (0.0833 < 1/6) and
-/// renorm to Σ=1, each weight stays at 1/6.
+/// renorm to mean = 1 (Σ = N = 6), each weight becomes 1.0 (uniform pre-
+/// renorm × 6/Σ = 1/6 × 6/1 = 1.0). Per spec §3.4.3 amended 2026-05-04.
 #[test]
 #[ignore = "requires GPU"]
 fn controller_z_score_at_zero_yields_midpoint_outputs() {
@@ -394,19 +395,20 @@ fn controller_z_score_at_zero_yields_midpoint_outputs() {
 
     let s = scratch.read_all();
 
-    // Σweights = 1.0 ± 1e-5 — renormalization invariant per spec §3.4.
+    // mean(weights) = 1.0 ± 1e-5 (Σweights = N = 6) — renormalization
+    // invariant per spec §3.4.3 amended 2026-05-04.
     let weight_sum: f32 = s[0..6].iter().sum();
     assert!(
-        (weight_sum - 1.0).abs() < 1e-5,
-        "weights sum = {weight_sum}, expected 1.0"
+        (weight_sum - 6.0).abs() < 1e-5,
+        "weights sum = {weight_sum}, expected 6.0 (mean = 1.0)"
     );
 
-    // Each weight at 1/6 (uniform input → uniform output through the
-    // 50/50 blend → renorm preserves uniformity).
+    // Each weight at 1.0 (uniform input → uniform output through the
+    // 50/50 blend → renorm preserves uniformity at the mean=1 level).
     for c in 0..6 {
         assert!(
-            (s[c] - 1.0_f32 / 6.0).abs() < 1e-4,
-            "weight[{c}] = {}, expected 1/6 (≈ 0.16667)",
+            (s[c] - 1.0_f32).abs() < 1e-4,
+            "weight[{c}] = {}, expected 1.0 (uniform mean)",
             s[c]
         );
     }
@@ -444,13 +446,17 @@ fn controller_z_score_at_zero_yields_midpoint_outputs() {
 /// SP11 A2 — weight renormalization works after the per-component floor.
 ///
 /// Pathological mag_ratios: one large (0.95), five small (0.01). The
-/// pre-floor blend with improving=stagnant=0.5 yields heavily skewed
-/// weights; the floor (`0.5 × 0.01 = 0.005` ≤ WEIGHT_HARD_FLOOR=0.01,
-/// so adaptive_floor = 0.01) clamps the small weights to 0.01; renorm to
-/// Σ=1 then redistributes the residual mass.
+/// pre-floor blend with delta=1, var=1 (dz=1, improving≈0.731,
+/// stagnant≈0.269) yields heavily skewed weights; the floor
+/// (`0.5 × 0.01 = 0.005` ≤ WEIGHT_HARD_FLOOR=0.01, so adaptive_floor =
+/// 0.01) clamps the small weights to 0.01; renorm to mean = 1 (Σ = N = 6)
+/// then scales by 6/Σ_pre. With the dominant pre-renorm weight ≈ 0.697,
+/// post-renorm w[0] = 4.18 which is above MAX_WEIGHT=3.0 — the
+/// per-component cap binds, pulling Σ_post below the nominal 6.
 ///
-/// The test asserts Σweights = 1 ± 1e-5 (the renormalization invariant)
-/// and that no weight drops below the WEIGHT_HARD_FLOOR after renorm.
+/// The test asserts the structural envelope per spec §3.4.3 amended
+/// 2026-05-04: each weight in [WEIGHT_HARD_FLOOR, MAX_WEIGHT], dominant
+/// weight saturates at MAX_WEIGHT, and Σweights ≤ N + tolerance.
 #[test]
 #[ignore = "requires GPU"]
 fn controller_weights_renormalize_after_floor() {
@@ -479,31 +485,46 @@ fn controller_weights_renormalize_after_floor() {
 
     let s = scratch.read_all();
 
-    // Renormalization invariant: Σweights = 1.0 ± 1e-5. PopArt's reward-
-    // magnitude EMA assumes a consistent reward-scale; this is the
-    // non-negotiable guarantee per spec §3.4.
+    // Renormalization invariant: Σweights ≤ N (= 6.0). PopArt's reward-
+    // magnitude EMA assumes a consistent reward-scale; per spec §3.4.3
+    // amended 2026-05-04 the controller normalises to mean=1 (Σ=N), with
+    // a per-component MAX_WEIGHT=3.0 cap that may pull Σ_post below N
+    // when one component dominates pre-renorm — as it does here (w[0]
+    // would be ≈ 4.18 without the cap). The cap is structural, not a
+    // soft suggestion; clipping below 6 is the expected behaviour.
     let weight_sum: f32 = s[0..6].iter().sum();
     assert!(
-        (weight_sum - 1.0).abs() < 1e-5,
-        "weights sum = {weight_sum} after renormalization, expected 1.0"
+        weight_sum <= 6.0 + 1e-5,
+        "weights sum = {weight_sum} after renormalization, expected ≤ 6.0 (mean ≤ 1.0)"
     );
 
-    // No weight below WEIGHT_HARD_FLOOR=0.01 after renorm. Renorm scales
-    // by 1/Σ_pre_floor_blend; when the pre-floor blend already had each
-    // weight ≥ 0.01 (because the floor pass clamped at adaptive_floor =
-    // 0.01), the post-renorm weight is `0.01 / blend_sum` which is at
-    // least slightly above 0.01 only if blend_sum < 1. With one weight
-    // dominating (0.95-derived blend ≈ 0.5) and five at floor, blend_sum
-    // ≈ 0.5 + 5 × 0.01 = 0.55, so the floored weights post-renorm are
-    // 0.01 / 0.55 ≈ 0.018 — comfortably above 0.01. The lower bound check
-    // `>= 0.01 - tolerance` validates the floor's structural guarantee.
+    // No weight below WEIGHT_HARD_FLOOR=0.01 after renorm. With one
+    // weight dominating (0.95-derived blend ≈ 0.697) and five at floor
+    // (≈ 0.0606), blend_sum ≈ 1.0; scale = 6/1.0 = 6.0; floored weights
+    // post-renorm = 0.0606 × 6 ≈ 0.36 — comfortably above 0.01. The
+    // lower-bound check `>= 0.01 - tolerance` validates the floor's
+    // structural guarantee. The companion upper-bound check (≤
+    // MAX_WEIGHT) validates the new cap.
     for c in 0..6 {
         assert!(
             s[c] >= 0.01 - 1e-5,
             "weight[{c}] = {} below WEIGHT_HARD_FLOOR=0.01 after renorm",
             s[c]
         );
+        assert!(
+            s[c] <= 3.0 + 1e-5,
+            "weight[{c}] = {} above MAX_WEIGHT=3.0 after renorm",
+            s[c]
+        );
     }
+
+    // Dominant component saturates at MAX_WEIGHT — the cap is binding
+    // for this pathological input, demonstrating the cap mechanism.
+    assert!(
+        (s[0] - 3.0).abs() < 1e-4,
+        "dominant weight[0] = {}, expected MAX_WEIGHT=3.0 (cap should bind)",
+        s[0]
+    );
 }
 
 /// SP11 A2 — saboteur post-clamp holds SABOTEUR_MIN at extreme regression.

docs/dqn-wire-up-audit.md

@@ -5105,3 +5105,53 @@ pass under `--features cuda` (GPU-gated); 10/10 sp5_isv_slots +
 `#[allow(dead_code)]` count is exactly the 12 pre-existing
 attributes (audited via `git show 25eba79ad^:...gpu_dqn_trainer.rs |
 grep -c allow(dead_code)`), no new ones.
+
+## SP11 B0 — controller renorm Σ=1 → mean=1 (2026-05-04)
+
+Spec §3.4.3 amended on main at `7ddaf9c51`: A2's `Σweights = 1`
+renormalization was discovered (during the implementer's B1 audit)
+to cause 6× reward-magnitude collapse on the mutually-exclusive
+component paths in `experience_env_step` (popart / micro / opp_cost
+fire at most one per bar — each carries weight 1/6 under Σ=1, so the
+running |reward| averages 1/6 of the pre-SP11 absolute scale). B0
+amends `reward_subsystem_controller_kernel.cu` to normalize to
+`mean(weights) = 1` (i.e., Σ = N = 6) with a per-component
+`MAX_WEIGHT = 3.0` cap to prevent any single component from
+dominating after renormalization.
+
+Code change scope:
+- `reward_subsystem_controller_kernel.cu` renormalization step:
+  `weights[c] = blends[c] / blend_sum` →
+  `weights[c] = min(MAX_WEIGHT, blends[c] × N_COMPONENTS / blend_sum)`
+- `MAX_WEIGHT=3.0f` and `N_COMPONENTS=6.0f` added to the
+  Invariant-1 const float anchor block at the top of the kernel
+  (Invariant-1 carve-out: rate-limiters / structural envelopes,
+  not regime thresholds — every other adaptive bound lives on ISV
+  per `feedback_isv_for_adaptive_bounds`).
+- Header comment updated to describe the amended renorm semantic
+  and the historical rationale (6× collapse on mutually-exclusive
+  components).
+
+Test changes (`crates/ml/tests/sp11_producer_unit_tests.rs`):
+- `controller_z_score_at_zero_yields_midpoint_outputs`:
+  `weight_sum ≈ 1.0` → `weight_sum ≈ 6.0`; per-component uniform
+  `1/6` → `1.0`.
+- `controller_weights_renormalize_after_floor`: assertion
+  generalized from exact-equality `Σ=1` to envelope `Σ ≤ N`
+  (because in this pathological test the cap binds — pre-cap
+  dominant weight ≈ 4.18 > MAX_WEIGHT=3.0 — pulling Σ_post below
+  the nominal 6). Added per-component `≤ MAX_WEIGHT` envelope
+  check; added explicit `dominant weight == MAX_WEIGHT` assertion
+  to validate the cap mechanism.
+- `controller_saboteur_post_clamp_holds_min`: no changes — this
+  test asserts only on `s[6]` (curiosity) and `s[7]` (saboteur),
+  both independent of the mean-vs-Σ renormalization choice.
+
+Pre-requisite for B1b structural reward-composition refactor
+(spec §3.5.3) which depends on the mean=1 semantic. B0 lands as a
+contained ~30 LOC commit before B1a/B1b/B1c so the contract
+change is reviewable in isolation.
+
+Verified post-fix: cargo check + release build clean; 6/6 SP11 GPU
+oracle tests pass on RTX 3050 Ti; 10/10 sp5_isv_slots + 4/4
+state_reset_registry contract tests pass.