From 302992f63a4df9da86e986e23b4062cbc7f7de6e Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Mon, 4 May 2026 08:40:33 +0200 Subject: [PATCH] =?UTF-8?q?fix(sp11):=20B0=20=E2=80=94=20controller=20reno?= =?UTF-8?q?rm=20=CE=A3=3D1=20=E2=86=92=20mean=3D1=20(post-A2=20spec=20amen?= =?UTF-8?q?dment)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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. --- .../reward_subsystem_controller_kernel.cu | 28 +++++-- crates/ml/tests/sp11_producer_unit_tests.rs | 77 ++++++++++++------- docs/dqn-wire-up-audit.md | 50 ++++++++++++ 3 files changed, 120 insertions(+), 35 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/reward_subsystem_controller_kernel.cu b/crates/ml/src/cuda_pipeline/reward_subsystem_controller_kernel.cu index 3dc49425c..afb94094c 100644 --- a/crates/ml/src/cuda_pipeline/reward_subsystem_controller_kernel.cu +++ b/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 diff --git a/crates/ml/tests/sp11_producer_unit_tests.rs b/crates/ml/tests/sp11_producer_unit_tests.rs index cce342bd1..22a28598a 100644 --- a/crates/ml/tests/sp11_producer_unit_tests.rs +++ b/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. diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 0a0ebc299..1938e67a7 100644 --- a/docs/dqn-wire-up-audit.md +++ b/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.