diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index ce5a8b80f..12aec6739 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -983,7 +983,7 @@ const ISV_NETWORK_DIM: usize = 23; /// (shifted 112→116 in Plan 4 Task 6 Commit A). /// Written by the constructor; checked at checkpoint load. Fail-fast only — no migration /// path exists. See spec §4.A.2 and `LAYOUT_FINGERPRINT_CURRENT` for structural-hash rationale. -pub(crate) const ISV_TOTAL_DIM: usize = 361; // SP5 + Layer D D1+D2+D3 + SP7 + SP8 + SP9 + SP10 + SP11: 173 + 188 (164 SP5 slots @ 174..278/280..340, with 2-slot gap before cross-fold-persistent Kelly block; Layer D D1 PnL outputs at [286..290); Layer D D2 health composition outputs at [290..294); Layer D D3 training metrics EMA at [294..297); SP7 T1 loss-balance Wiener stats at [297..313); SP7 activation-flag fix per-(head,branch) flags at [313..321); SP8 Fix 36 train_active_frac canary @ [321..322) + LB_MAX_BUDGET per-(head,branch) at [322..330); SP9 Fix 37 Kelly warmup floor at [330..331) + Q_VAR_MAG_EMA at [331..332) + INTENT_EVAL_DIVERGENCE at [332..333) + 3 EMA targets at [333..336) + 3 eval_dist mag bins at [336..339); SP10 Fix 38 EVAL_THOMPSON_TEMP at [339..340); SP11 Fix 39 reward-subsystem controller at [340..361) — 6 component weights [340..346) + 4 controller scalar outputs [346..350) + 2 val-sharpe canaries [350..352) + 6 mag-ratio canaries [352..358) + saboteur engagement [358] + PnL magnitude EMA [359] + popart-component magnitude EMA [360, B1b fix-up]) +pub(crate) const ISV_TOTAL_DIM: usize = 367; // SP5 + Layer D D1+D2+D3 + SP7 + SP8 + SP9 + SP10 + SP11: 173 + 194 (164 SP5 slots @ 174..278/280..340, with 2-slot gap before cross-fold-persistent Kelly block; Layer D D1 PnL outputs at [286..290); Layer D D2 health composition outputs at [290..294); Layer D D3 training metrics EMA at [294..297); SP7 T1 loss-balance Wiener stats at [297..313); SP7 activation-flag fix per-(head,branch) flags at [313..321); SP8 Fix 36 train_active_frac canary @ [321..322) + LB_MAX_BUDGET per-(head,branch) at [322..330); SP9 Fix 37 Kelly warmup floor at [330..331) + Q_VAR_MAG_EMA at [331..332) + INTENT_EVAL_DIVERGENCE at [332..333) + 3 EMA targets at [333..336) + 3 eval_dist mag bins at [336..339); SP10 Fix 38 EVAL_THOMPSON_TEMP at [339..340); SP11 Fix 39 reward-subsystem controller at [340..367) — 6 component weights [340..346) + 4 controller scalar outputs [346..350) + 2 val-sharpe canaries [350..352) + 6 mag-ratio canaries [352..358) + saboteur engagement [358] + PnL magnitude EMA [359] + popart-component magnitude EMA [360, B1b fix-up] + per-component variance EMAs [361..367, B1b smoke-recovery z-score normalization]) /// Legacy alias preserved for call sites that haven't been audited for the /// network-vs-total split. New code should pick `ISV_NETWORK_DIM` (for weight /// tensor sizing) or `ISV_TOTAL_DIM` (for the broadcast bus buffer). @@ -1304,8 +1304,25 @@ pub const SP5_WIENER_TOTAL_FLOATS: usize = /// SP11 (Fix 39, 2026-05-04, B1b fix-up) adds: /// 1 float for popart_component_ema (SCRATCH_SP11_POPART_COMPONENT_EMA=286, [286..287)) /// → ISV[POPART_COMPONENT_MAG_EMA_INDEX=360] via chained Pearls A+D singleton -/// Combined: 259 + 6 + 1 + 10 + 10 + 1 = 287 scratch slots [0..287). -pub const SP5_SCRATCH_TOTAL: usize = 287; +/// SP11 (Fix 39, 2026-05-04, B1b smoke-recovery) adds: +/// 1 float for popart_component_var (SCRATCH_SP11_POPART_COMPONENT_VAR=287, [287..288)) +/// → ISV[REWARD_COMPONENT_VAR_EMA_BASE=361] via chained Pearls A+D +/// (paired with the popart magnitude EMA to enable z-score normalisation +/// in the mag-ratio canary). Computed alongside the magnitude in the +/// extended `popart_component_ema_kernel` via Welford's online algorithm +/// (single-pass two-block-tree-reduce: first pass produces `mean(|x|)`, +/// second pass produces `var(|x|)` against that mean). +/// 5 floats for reward_component_var (SCRATCH_SP11_REWARD_COMPONENT_VAR_BASE=288, [288..293)) +/// → ISV[REWARD_COMPONENT_VAR_EMA_BASE+1..+6) = ISV[362..367) via 5 chained +/// Pearls A+D singleton launches. Layout: [288]=cf var, [289]=trail var, +/// [290]=micro var, [291]=opp_cost var, [292]=bonus var. Computed by the +/// extended `reward_component_ema_kernel` via Welford's online algorithm +/// per non-popart component (component 0 / popart variance is sourced +/// from the `popart_component_ema_kernel` chain because slot 63 is the +/// pre-SP11 invariant total-reward magnitude EMA, not the popart-component +/// magnitude). Spec §4 amendment "Why z-score" (lines 564-619). +/// Combined: 259 + 6 + 1 + 10 + 10 + 1 + 1 + 5 = 293 scratch slots [0..293). +pub const SP5_SCRATCH_TOTAL: usize = 293; /// SP5 Layer D Task D1 (rewrite, 2026-05-02): scratch index base for /// `pnl_aggregation_update`. @@ -1448,12 +1465,45 @@ pub const SCRATCH_SP11_MAG_RATIO_BASE: usize = 269; // [269..276) — f pub const SCRATCH_SP11_CONTROLLER_BASE: usize = 276; // [276..286) → ISV[340..350) /// SP11 (Fix 39, 2026-05-04, B1b fix-up): scratch slot for the -/// `popart_component_ema` producer. Single float at [286..287) → -/// ISV[POPART_COMPONENT_MAG_EMA_INDEX=360] via chained Pearls A+D -/// singleton launch. Resolves the slot-63 PopArt-input vs popart- -/// component-magnitude overload that B1b smoke surfaced. +/// `popart_component_ema` producer's mean(|x|) output. Single float at +/// [286..287) → ISV[POPART_COMPONENT_MAG_EMA_INDEX=360] via chained +/// Pearls A+D singleton launch. Resolves the slot-63 PopArt-input vs +/// popart-component-magnitude overload that B1b smoke surfaced. pub const SCRATCH_SP11_POPART_COMPONENT_EMA: usize = 286; // [286..287) → ISV[POPART_COMPONENT_MAG_EMA_INDEX=360] +/// SP11 (Fix 39, 2026-05-04, B1b smoke-recovery): scratch slot for the +/// `popart_component_ema` producer's var(|x|) Welford output. Single +/// float at [287..288) → ISV[REWARD_COMPONENT_VAR_EMA_BASE=361] via +/// chained Pearls A+D singleton launch. Computed in the same single- +/// block kernel as the magnitude: a two-pass tree-reduce (mean → var) +/// against the just-computed mean. Pairs with the magnitude EMA to +/// drive z-score normalisation in `reward_component_mag_ratio_compute_kernel`. +pub const SCRATCH_SP11_POPART_COMPONENT_VAR: usize = 287; // [287..288) → ISV[REWARD_COMPONENT_VAR_EMA_BASE=361] + +/// SP11 (Fix 39, 2026-05-04, B1b smoke-recovery): scratch base for the +/// 5 non-popart per-reward-component variance EMAs computed by the +/// extended `reward_component_ema_kernel`. Layout (matching component +/// indices c=1..5 in `reward_components_per_sample`): +/// [288] = cf variance → ISV[REWARD_COMPONENT_VAR_EMA_BASE+1=362] +/// [289] = trail variance → ISV[REWARD_COMPONENT_VAR_EMA_BASE+2=363] +/// [290] = micro variance → ISV[REWARD_COMPONENT_VAR_EMA_BASE+3=364] +/// [291] = opp_cost variance → ISV[REWARD_COMPONENT_VAR_EMA_BASE+4=365] +/// [292] = bonus variance → ISV[REWARD_COMPONENT_VAR_EMA_BASE+5=366] +/// +/// c=0 (popart) variance lives at SCRATCH_SP11_POPART_COMPONENT_VAR +/// (slot 287) and is sourced from `popart_component_ema_kernel` because +/// `reward_components_per_sample[+0]` feeds slot 63 = pre-SP11 +/// invariant total-reward magnitude EMA, NOT the popart-component +/// magnitude (see B1b fix-up rationale in `popart_component_ema_kernel.cu`). +/// Each variance is computed via two-pass Welford: first reduce produces +/// `mean(|r_c|)`, second reduce produces `var(|r_c|)` against the just- +/// computed mean. Smoothed per-slot by chained Pearls A+D singletons +/// (5 independent launches because the destination ISV slots 362..367 +/// are contiguous but the sources occupy a contiguous scratch block — +/// the actual chained launcher uses `n_slots=5` in one shot since both +/// scratch and ISV layouts are matched-stride contiguous). +pub const SCRATCH_SP11_REWARD_COMPONENT_VAR_BASE: usize = 288; // [288..293) → ISV[362..367) + /// SP5 Task A7: scratch index base for pearl_8_trail_update trail_dist[4] output block. /// Slots [199..203): per-direction trail-stop distance (Short=0, Hold=1, Long=2, Flat=3). /// Written by `pearl_8_trail_update`; consumed by apply_pearls_ad_kernel → @@ -1989,7 +2039,8 @@ const fn layout_fingerprint_seed() -> &'static [u8] { REWARD_COMPONENT_MAG_RATIO_BASE=352;\ SABOTEUR_ENGAGEMENT_RATE=358;PNL_REWARD_MAGNITUDE_EMA=359;\ POPART_COMPONENT_MAG_EMA=360;\ - ISV_TOTAL_DIM=361;\ + REWARD_COMPONENT_VAR_EMA_BASE=361;\ + ISV_TOTAL_DIM=367;\ PARAM_W_A_H_S1=0;PARAM_B_A_H_S1=1;PARAM_W_B_H_S1=2;PARAM_B_B_H_S1=3;\ PARAM_W_RESIDUAL_H_S1=4;PARAM_GAMMA_H_S1=5;PARAM_BETA_H_S1=6;\ PARAM_W_A_H_S2=7;PARAM_B_A_H_S2=8;PARAM_W_B_H_S2=9;PARAM_B_B_H_S2=10;\ @@ -12798,6 +12849,7 @@ impl GpuDqnTrainer { use crate::cuda_pipeline::sp11_isv_slots::{ REWARD_COMPONENT_MAG_RATIO_BASE, REWARD_COMPONENT_MAG_RATIO_COUNT, PNL_REWARD_MAGNITUDE_EMA_INDEX, POPART_COMPONENT_MAG_EMA_INDEX, + REWARD_COMPONENT_VAR_EMA_BASE, }; debug_assert!(self.isv_signals_dev_ptr != 0, @@ -12816,6 +12868,13 @@ impl GpuDqnTrainer { // contiguous slots starting at slot 63 — which was overloaded. let popart_specific_slot_i32: i32 = POPART_COMPONENT_MAG_EMA_INDEX as i32; let cf_others_base_slot_i32: i32 = (REWARD_POPART_EMA_INDEX + 1) as i32; + // B1b smoke-recovery (2026-05-04): per-component variance slots + // for z-score normalisation. Popart variance lives at slot 361 + // (paired with mag at slot 360 — separate producer per the + // slot-63 fix-up rationale); cf/trail/micro/opp_cost/bonus + // variances live at slots 362..366 (paired with mags at 64..68). + let popart_var_slot_i32: i32 = REWARD_COMPONENT_VAR_EMA_BASE as i32; + let cf_others_var_base_slot_i32: i32 = (REWARD_COMPONENT_VAR_EMA_BASE + 1) as i32; unsafe { self.stream @@ -12824,6 +12883,8 @@ impl GpuDqnTrainer { .arg(&scratch_out_dev) .arg(&popart_specific_slot_i32) .arg(&cf_others_base_slot_i32) + .arg(&popart_var_slot_i32) + .arg(&cf_others_var_base_slot_i32) .launch(LaunchConfig { grid_dim: (1, 1, 1), block_dim: (6, 1, 1), @@ -12893,7 +12954,9 @@ impl GpuDqnTrainer { pub(crate) fn launch_sp11_popart_component_ema(&self) -> Result<(), MLError> { use crate::cuda_pipeline::sp4_wiener_ema::{launch_apply_pearls, ALPHA_META}; use crate::cuda_pipeline::sp5_isv_slots::SP5_SLOT_BASE; - use crate::cuda_pipeline::sp11_isv_slots::POPART_COMPONENT_MAG_EMA_INDEX; + use crate::cuda_pipeline::sp11_isv_slots::{ + POPART_COMPONENT_MAG_EMA_INDEX, REWARD_COMPONENT_VAR_EMA_BASE, + }; debug_assert!(self.isv_signals_dev_ptr != 0, "launch_sp11_popart_component_ema: isv_signals_dev_ptr must be allocated"); @@ -12924,7 +12987,10 @@ impl GpuDqnTrainer { // Step 1: producer kernel — block tree-reduce, BLOCK_SIZE=256. // Shared memory = 256 floats = 1024 bytes (matches the kernel's - // `extern __shared__ float sdata[]` declaration). + // `extern __shared__ float sdata[]` declaration). Writes mean(|x|) + // to scratch_out[0] AND var(|x|) to scratch_out[1] via single-pass + // Welford (B1b smoke-recovery, 2026-05-04 — spec §4 amendment + // "Why z-score" lines 564-619). unsafe { self.stream .launch_builder(&self.sp11_popart_component_ema_kernel) @@ -12940,17 +13006,36 @@ impl GpuDqnTrainer { } // Step 2: chained apply_pearls_ad_kernel singleton → - // ISV[POPART_COMPONENT_MAG_EMA_INDEX=360]. + // ISV[POPART_COMPONENT_MAG_EMA_INDEX=360] (mag). let base_wiener_offset = SP4_PRODUCER_COUNT as i32 * 3; - let isv_idx_i32 = POPART_COMPONENT_MAG_EMA_INDEX as i32; - let wiener_off = base_wiener_offset + (isv_idx_i32 - SP5_SLOT_BASE as i32) * 3; + let mag_isv_idx_i32 = POPART_COMPONENT_MAG_EMA_INDEX as i32; + let mag_wiener_off = base_wiener_offset + (mag_isv_idx_i32 - SP5_SLOT_BASE as i32) * 3; unsafe { launch_apply_pearls( &self.stream, &self.apply_pearls_ad_kernel, scratch_dev, scratch_idx_i32, - isv_dev, isv_idx_i32, - wiener_dev, wiener_off, + isv_dev, mag_isv_idx_i32, + wiener_dev, mag_wiener_off, + 1, + ALPHA_META, + )?; + } + + // Step 3 (B1b smoke-recovery): chained singleton → + // ISV[REWARD_COMPONENT_VAR_EMA_BASE=361] (popart variance). + // Scratch source is the second float written by the kernel + // (scratch_idx + 1). + let var_isv_idx_i32 = REWARD_COMPONENT_VAR_EMA_BASE as i32; + let var_wiener_off = base_wiener_offset + (var_isv_idx_i32 - SP5_SLOT_BASE as i32) * 3; + let var_scratch_idx_i32 = scratch_idx_i32 + 1; + unsafe { + launch_apply_pearls( + &self.stream, + &self.apply_pearls_ad_kernel, + scratch_dev, var_scratch_idx_i32, + isv_dev, var_isv_idx_i32, + wiener_dev, var_wiener_off, 1, ALPHA_META, )?; diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index cb0240396..69e31c793 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs @@ -2413,9 +2413,12 @@ impl GpuExperienceCollector { n_samples: usize, ) -> Result<(), crate::MLError> { use crate::cuda_pipeline::gpu_dqn_trainer::{ - REWARD_COMPONENT_COUNT, REWARD_POPART_EMA_INDEX, + REWARD_COMPONENT_COUNT, REWARD_POPART_EMA_INDEX, SP4_PRODUCER_COUNT, + SCRATCH_SP11_REWARD_COMPONENT_VAR_BASE, }; use crate::cuda_pipeline::sp4_wiener_ema::launch_apply_pearls; + use crate::cuda_pipeline::sp5_isv_slots::SP5_SLOT_BASE; + use crate::cuda_pipeline::sp11_isv_slots::REWARD_COMPONENT_VAR_EMA_BASE; if n_samples == 0 || self.isv_signals_dev_ptr == 0 { return Ok(()); @@ -2450,6 +2453,11 @@ impl GpuExperienceCollector { let n_i32 = n_samples as i32; let scratch_first_i32 = SCRATCH_FIRST as i32; + // SP11 B1b smoke-recovery: kernel also writes 5 variance outputs + // for c=1..5 to scratch[SCRATCH_SP11_REWARD_COMPONENT_VAR_BASE..+5). + // Component 0 (popart) variance is sourced from the separate + // `popart_component_ema_kernel` (B1b fix-up rationale). + let var_scratch_first_i32 = SCRATCH_SP11_REWARD_COMPONENT_VAR_BASE as i32; let rc_ptr = self.reward_components_per_sample.raw_ptr(); let scratch_dev = self.reward_component_pearls_scratch_dev_ptr; unsafe { @@ -2458,6 +2466,7 @@ impl GpuExperienceCollector { .arg(&n_i32) .arg(&scratch_dev) .arg(&scratch_first_i32) + .arg(&var_scratch_first_i32) .launch(LaunchConfig { grid_dim: (1, 1, 1), block_dim: (REWARD_COMPONENT_COUNT as u32, 1, 1), @@ -2487,6 +2496,33 @@ impl GpuExperienceCollector { )?; } + // SP11 B1b smoke-recovery: chained Pearls A+D for the 5 non-popart + // variance EMAs. Source is scratch[SCRATCH_SP11_REWARD_COMPONENT_VAR_BASE..+5); + // destination is ISV[REWARD_COMPONENT_VAR_EMA_BASE+1..+6) = + // ISV[362..367). Wiener offset uses the SP5 base (SP4_PRODUCER_COUNT*3 + // + (isv_idx − SP5_SLOT_BASE) * 3) — matching the convention in + // `launch_sp11_popart_component_ema` etc., not the SP4 retrofit + // convention used for the magnitudes above (which use scratch_idx*3 + // offsets because they mapped to slot 63 in the SP4 retrofit block). + let base_wiener_offset = SP4_PRODUCER_COUNT as i32 * 3; + let var_isv_base_i32 = (REWARD_COMPONENT_VAR_EMA_BASE + 1) as i32; + let var_wiener_off = base_wiener_offset + + (var_isv_base_i32 - SP5_SLOT_BASE as i32) * 3; + unsafe { + launch_apply_pearls( + &self.stream, + &self.apply_pearls_ad_kernel, + scratch_dev, + var_scratch_first_i32, + self.reward_component_pearls_isv_dev_ptr, + var_isv_base_i32, + self.reward_component_pearls_wiener_dev_ptr, + var_wiener_off, + 5, + crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META, + )?; + } + self.stream.memset_zeros(&mut self.reward_components_per_sample) .map_err(|e| crate::MLError::ModelError(format!("memset reward_components_per_sample: {e}")))?; Ok(()) diff --git a/crates/ml/src/cuda_pipeline/popart_component_ema_kernel.cu b/crates/ml/src/cuda_pipeline/popart_component_ema_kernel.cu index 9517377c1..e495fe27c 100644 --- a/crates/ml/src/cuda_pipeline/popart_component_ema_kernel.cu +++ b/crates/ml/src/cuda_pipeline/popart_component_ema_kernel.cu @@ -1,12 +1,23 @@ // crates/ml/src/cuda_pipeline/popart_component_ema_kernel.cu // -// SP11 B1b fix-up (2026-05-04): popart-component-specific magnitude EMA. +// SP11 B1b fix-up (2026-05-04): popart-component-specific magnitude EMA +// AND variance EMA (B1b smoke-recovery). // // Reads `popart_component_per_sample[N]` (an N-bar mapped-pinned f32 // buffer populated by `experience_env_step` at the `r_popart` write site) -// and computes mean(|x|) into `scratch_out[0]`. The chained -// `apply_pearls_ad_kernel` launch on the same stream targets ISV slot -// `POPART_COMPONENT_MAG_EMA_INDEX = 360`. +// and computes: +// scratch_out[0] = mean(|x|) → ISV[POPART_COMPONENT_MAG_EMA_INDEX=360] +// scratch_out[1] = var(|x|) → ISV[REWARD_COMPONENT_VAR_EMA_BASE=361] +// +// The chained `apply_pearls_ad_kernel` launches on the same stream +// target ISV slots 360 (n_slots=1 mag) and 361 (n_slots=1 var) — see +// the two-singleton-chain in `launch_sp11_popart_component_ema`. +// +// Why both outputs in ONE kernel: a single-pass Welford via +// two-block-tree-reduces in the same kernel keeps the producer +// stream-ordered (no host sync between mag and var) and reuses the +// already-loaded |x| values across both reduces. The cost is a single +// extra `__syncthreads()` between the mean and variance passes. // // Why a separate kernel from `reward_component_ema_kernel`: // The pre-SP11 SP4 reward_component_ema kernel writes @@ -26,8 +37,17 @@ // preserved); the SP11 mag-ratio canary reads slot 360 for the popart // axis (this kernel's output) and slots 64..68 for the other 5 axes. // -// Single block (BLOCK_SIZE=256). Block tree-reduce — no atomicAdd per -// `feedback_no_atomicadd.md`. Pure GPU compute per +// Why variance (B1b smoke-recovery, 2026-05-04, spec §4 amendment +// "Why z-score" lines 564-619): the linear mag-ratio formula +// `winner_weight = mag_ratio` amplified popart's intrinsic O(100) +// magnitude over the other 5 components' O(0.1-2). w_pop saturated +// toward MAX_WEIGHT, curiosity_b exploded, sharpe collapsed. Z-score +// makes ratios scale-invariant — every component contributes "shares +// of own scale". The variance EMA at slot 361 is the σ² denominator +// in the canary's z[0] = mag[0] / max(sqrt(var[0]), EPS_DIV). +// +// Single block (BLOCK_SIZE=256). Two block tree-reduces — no atomicAdd +// per `feedback_no_atomicadd.md`. Pure GPU compute per // `feedback_no_cpu_compute_strict.md`. #include @@ -42,8 +62,9 @@ extern "C" __global__ void popart_component_ema_kernel( const float* __restrict__ popart_component_per_sample, /* Producer scratch buffer. The launcher passes a pointer offset to * `scratch[SCRATCH_SP11_POPART_COMPONENT_EMA]` so this kernel writes - * to `scratch_out[0]`. Chained `apply_pearls_ad_kernel` (n_slots=1) - * targets ISV[POPART_COMPONENT_MAG_EMA_INDEX=360]. */ + * to `scratch_out[0]` (mean) and `scratch_out[1]` (variance). + * Chained `apply_pearls_ad_kernel` launches target ISV[360] (mag) + * and ISV[361] (var) — see `launch_sp11_popart_component_ema`. */ float* __restrict__ scratch_out, const int n_samples) { @@ -51,7 +72,8 @@ extern "C" __global__ void popart_component_ema_kernel( if (blockIdx.x != 0) return; extern __shared__ float sdata[]; - /* Strided accumulation over the per-bar buffer. Each thread folds + /* ── Pass 1: mean(|x|) ─────────────────────────────────────────── + * Strided accumulation over the per-bar buffer. Each thread folds * |popart_component_per_sample[i]| into its private accumulator, * then we tree-reduce in shared memory. */ float acc = 0.0f; @@ -69,10 +91,37 @@ extern "C" __global__ void popart_component_ema_kernel( __syncthreads(); } + /* `fmaxf(1.0f, n_samples)` guards the impossible zero-bar edge case + * (no producer ever writes; debug-only). All threads need `mean` + * for the variance pass below — broadcast via shared mem. */ + __shared__ float mean_abs; if (threadIdx.x == 0) { - /* Mean(|x|) — `fmaxf(1.0f, n_samples)` guards the impossible - * zero-bar edge case (no producer ever writes; debug-only). */ - scratch_out[0] = sdata[0] / fmaxf(1.0f, (float)n_samples); + mean_abs = sdata[0] / fmaxf(1.0f, (float)n_samples); + scratch_out[0] = mean_abs; + } + __syncthreads(); + + /* ── Pass 2: var(|x|) via Welford against the just-computed mean ── + * var = (1/N) × Σ (|x_i| − mean)². Population variance (matches the + * pre-SP11 `mean_abs` semantics — both are batch-scope statistics, + * not sample-corrected estimators of an unknown population). */ + float v_acc = 0.0f; + for (int i = threadIdx.x; i < n_samples; i += blockDim.x) { + float d = fabsf(popart_component_per_sample[i]) - mean_abs; + v_acc += d * d; + } + sdata[threadIdx.x] = v_acc; + __syncthreads(); + + for (int s = blockDim.x / 2; s > 0; s >>= 1) { + if (threadIdx.x < s) { + sdata[threadIdx.x] += sdata[threadIdx.x + s]; + } + __syncthreads(); + } + + if (threadIdx.x == 0) { + scratch_out[1] = sdata[0] / fmaxf(1.0f, (float)n_samples); __threadfence_system(); } } diff --git a/crates/ml/src/cuda_pipeline/reward_component_ema_kernel.cu b/crates/ml/src/cuda_pipeline/reward_component_ema_kernel.cu index 95d3a8a97..0c6bdb78c 100644 --- a/crates/ml/src/cuda_pipeline/reward_component_ema_kernel.cu +++ b/crates/ml/src/cuda_pipeline/reward_component_ema_kernel.cu @@ -4,6 +4,15 @@ * Plan 3 Task 1, spec §4.C.2; SP4 Layer A Task A13.5 retrofit (2026-05-01). * GPU-drives-CPU-reads (spec §4.C.6). * + * SP11 B1b smoke-recovery (2026-05-04, spec §4 amendment lines 564-619): + * also computes per-component variance for c=1..5 via two-pass Welford + * against the just-computed mean. Variance for c=0 (popart) is NOT + * computed here — slot 63 is the pre-SP11 invariant total-reward + * magnitude EMA (PopArt's normalisation input), NOT popart-component + * magnitude. Popart-component variance lives at scratch slot 287 and is + * produced by `popart_component_ema_kernel` from the dedicated + * `popart_component_per_sample` GPU buffer (B1b fix-up). + * * Reads per-sample per-component rewards from `reward_components_per_sample` * [n_samples * 6], where the 6 components are laid out as: * @@ -22,6 +31,13 @@ * to ISV[REWARD_POPART_EMA_INDEX + c]. 2026-05-01 GPU-Pearls refactor — * graph-capture-compatible (no host sync between kernels). * + * SP11 B1b smoke-recovery: ALSO writes per-component variance for c=1..5 + * to `scratch_buf[var_scratch_first_index + (c - 1)]` (5 contiguous floats). + * The launcher chains a second `apply_pearls_ad_kernel` (n_slots=5) into + * ISV[REWARD_COMPONENT_VAR_EMA_BASE+1..+6) = ISV[362..367). Component 0 + * (popart) variance is sourced from the separate `popart_component_ema_kernel` + * (see B1b fix-up rationale) and lands at ISV[REWARD_COMPONENT_VAR_EMA_BASE=361]. + * * **α dropped per SP4 — Pearls A+D adapt α from per-slot signal-vs-noise * variance.** Cross-boundary wiring: the trainer's mapped-pinned * `producer_step_scratch_buf` + `wiener_state_buf` + `isv_signals_pinned` @@ -31,14 +47,21 @@ * collector construction (mirrors A14/A15 Pearl C wiring path). * * Single-block kernel (one thread per component). Cold-path — once per - * training step. Cost: 6 threads × n_samples additions + 6 scratch writes. + * training step. Cost: 6 threads × n_samples additions for the mean + * pass + 5 threads × n_samples deviations for the variance pass + 11 + * scratch writes total (6 mag + 5 var). No atomicAdd — each thread + * owns its component's accumulator independently. */ extern "C" __global__ void reward_component_ema( const float* __restrict__ reward_components, /* [n_samples * 6] row-major */ int n_samples, float* __restrict__ scratch_buf, /* producer_step_scratch_buf */ - int scratch_first_index /* slot 63 — first of 6 contiguous slots */ + int scratch_first_index, /* slot 63 — first of 6 contiguous mag slots */ + /* SP11 B1b smoke-recovery: scratch slot for variance outputs c=1..5 + * (5 contiguous floats). Component 0 (popart) variance lives at the + * separate popart_component_ema_kernel's output. */ + int var_scratch_first_index /* slot 288 — first of 5 contiguous var slots (cf, trail, micro, opp_cost, bonus) */ ) { /* Per-thread one component (REWARD_COMPONENT_COUNT=6 in production layout — * see `REWARD_COMPONENT_COUNT` in `gpu_dqn_trainer.rs`). The host launcher @@ -49,7 +72,7 @@ extern "C" __global__ void reward_component_ema( int c = (int)threadIdx.x; if (c >= 6 || blockIdx.x != 0) return; - /* Mean |r| over batch for this component. */ + /* ── Pass 1: mean |r| over batch for this component. ────────────── */ float acc = 0.0f; for (int idx = 0; idx < n_samples; ++idx) { acc += fabsf(reward_components[idx * 6 + c]); @@ -57,5 +80,22 @@ extern "C" __global__ void reward_component_ema( float mean_abs = acc / fmaxf(1.0f, (float)n_samples); scratch_buf[scratch_first_index + c] = mean_abs; + + /* ── Pass 2: var |r| via Welford against the just-computed mean. ── + * Component 0 (popart) variance is NOT written here — see file + * header for the B1b fix-up rationale (slot 63 vs slot 360 + * separation). Threads c=1..5 each own one variance-output slot. */ + if (c >= 1 && c <= 5) { + float v_acc = 0.0f; + for (int idx = 0; idx < n_samples; ++idx) { + float d = fabsf(reward_components[idx * 6 + c]) - mean_abs; + v_acc += d * d; + } + float var_abs = v_acc / fmaxf(1.0f, (float)n_samples); + /* Component 1 (cf) → var_scratch_first_index + 0, + * component 5 (bonus) → var_scratch_first_index + 4. */ + scratch_buf[var_scratch_first_index + (c - 1)] = var_abs; + } + __threadfence_system(); /* PCIe-visible to mapped-pinned host_ptr */ } diff --git a/crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu b/crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu index d488791dc..332729b17 100644 --- a/crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu +++ b/crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu @@ -1,11 +1,15 @@ // crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu // // SP11 (Fix 39, 2026-05-04): per-component reward-magnitude-ratio canary -// producer. B1b fix-up (2026-05-04) split the popart axis off slot 63 -// (which is PopArt's normalisation input = total reward magnitude EMA; -// pre-SP11 invariant preserved) onto a dedicated slot 360 = +// producer with z-score normalisation (B1b smoke-recovery 2026-05-04). +// +// B1b fix-up (2026-05-04) split the popart axis off slot 63 (which is +// PopArt's normalisation input = total reward magnitude EMA; pre-SP11 +// invariant preserved) onto a dedicated slot 360 = // `POPART_COMPONENT_MAG_EMA_INDEX` produced by `popart_component_ema_kernel`. -// This kernel reads non-contiguous sources after the fix: +// B1b smoke-recovery (2026-05-04) added per-component variance EMAs at +// slots 361..367 to drive z-score normalisation. This kernel reads +// non-contiguous magnitude sources after the fix-up: // // axis 0 (popart) = isv[popart_specific_slot] ← slot 360 (B1b fix-up) // axis 1 (cf) = isv[cf_others_base_slot + 0] ← slot 64 @@ -14,30 +18,55 @@ // axis 4 (opp_cost) = isv[cf_others_base_slot + 3] ← slot 67 // axis 5 (bonus) = isv[cf_others_base_slot + 4] ← slot 68 // -// Why two slot indices: pre-fix-up the kernel read 6 contiguous slots -// at `popart_ema_base_slot..+6` (= REWARD_POPART_EMA_INDEX=63..69). B1b -// smoke surfaced that slot 63 was overloaded — see spec §4 amendment -// "Why slot 360" + the popart_component_ema_kernel.cu header for the -// full diagnosis. The cf/trail/micro/opp_cost/bonus axes still live at -// slots 64..68 (their producer `reward_component_ema_kernel` is -// untouched). Only the popart axis is sourced from a different slot. +// And per-component variance sources (B1b smoke-recovery): +// +// axis 0 var = isv[popart_var_slot] ← slot 361 +// axis 1 var = isv[cf_others_var_base_slot + 0] ← slot 362 +// ... +// axis 5 var = isv[cf_others_var_base_slot + 4] ← slot 366 +// +// Why z-score (B1b smoke-recovery, spec §4 amendment lines 564-619): +// the linear ratio formula `winner_weight = mag_ratio` amplified +// whichever reward component had the largest intrinsic magnitude. +// Popart (trade P&L) is O(100) per fire; cf/trail/micro/opp_cost/bonus +// are O(0.1-2). Result: `w_pop` saturated toward MAX_WEIGHT in the +// controller, curiosity_b exploded, sharpe collapsed (B1b smoke +// `smoke-test-6wd2c` on commit `61b2fa962`). Z-score normalisation +// makes ratios scale-invariant — every component contributes "shares +// of own scale" rather than "shares of absolute magnitude". The +// controller's redistribution logic (`winner_weight = ratio`, +// unchanged) works correctly when it sees ratios ≈ 1/6 each at +// stability rather than ratios skewed by intrinsic O(100) vs O(1) +// scale differences. +// +// Why two slot-pair indices: the popart axis lives at non-contiguous +// magnitude (360) and variance (361) slots, while the other 5 axes +// have contiguous magnitude (64..68) and variance (362..366) slots. +// Pre-fix-up the kernel read 6 contiguous mag slots at slot 63..69 — +// see the popart_component_ema_kernel.cu header for the slot-63 +// overload diagnosis. The cf/trail/micro/opp_cost/bonus axes still +// share their producer (`reward_component_ema_kernel`); only popart +// is sourced from a different magnitude+variance slot pair. // // Writes 7 floats to `scratch_out`: -// scratch_out[0..6) = ratios (each / sum), Σ=1 modulo EPS_DIV floor -// scratch_out[6] = isv[popart_specific_slot] — popart-component -// magnitude mirrored to PNL_REWARD_MAGNITUDE_EMA_INDEX -// = 359 via the chained Pearls-A+D launch (n_slots=1). +// scratch_out[0..6) = ratios (each = z[c] / Σz, Σ=1 modulo EPS_DIV floor), +// where z[c] = mag[c] / max(sqrt(var[c]), EPS_DIV). +// scratch_out[6] = mag[0] (popart raw magnitude) — mirrored to +// PNL_REWARD_MAGNITUDE_EMA_INDEX = 359 via the +// chained Pearls A+D launch (n_slots=1). // -// Mirror semantic preserved post-fix-up: the popart-COMPONENT-specific -// magnitude (not slot 63's overloaded value) is the right "PnL signal -// scale" for curiosity bounding (spec §3.4.2) and saboteur-engagement -// thresholding (spec §3.3.1). Slot 359 stays semantically PnL signal -// scale; the source value is now correct. +// Mirror semantic preserved post-z-score: scratch_out[6] stays as raw +// `mag[0]` (popart magnitude EMA, slot 360 source) because slot 359 +// = `PNL_REWARD_MAGNITUDE_EMA_INDEX` is consumed by curiosity bounding +// (spec §3.4.2) and saboteur engagement thresholding (spec §3.3.1) — +// both need raw P&L scale, NOT z-score. Z-scoring is purely the +// normalisation step inside the ratio computation; the mirror feeds +// downstream consumers that need physical scale information. // // Pearls A+D applied chain-style by `apply_pearls_ad_kernel` (n_slots=6 // for the contiguous ratio block at ISV[REWARD_COMPONENT_MAG_RATIO_BASE -// =352..358), then n_slots=1 for the singleton mirror at ISV[359]) by the -// launcher `launch_sp11_mag_ratio_compute` in `gpu_dqn_trainer.rs`. +// =352..358), then n_slots=1 for the singleton mirror at ISV[359]) by +// the launcher `launch_sp11_mag_ratio_compute` in `gpu_dqn_trainer.rs`. // The kernel emits all 7 scratch values in one shot. // // Single block, 6 threads (one per component); thread 0 also computes @@ -58,48 +87,84 @@ extern "C" __global__ void reward_component_mag_ratio_compute_kernel( /* ISV slot index of the first non-popart per-component magnitude EMA * = REWARD_POPART_EMA_INDEX + 1 = 64 (cf at 64, trail 65, micro 66, * opp_cost 67, bonus 68 — produced by SP4 reward_component_ema_kernel). */ - int cf_others_base_slot + int cf_others_base_slot, + /* ISV slot index of popart-component-specific variance EMA + * = REWARD_COMPONENT_VAR_EMA_BASE = 361 (B1b smoke-recovery, 2026-05-04). */ + int popart_var_slot, + /* ISV slot index of the first non-popart per-component variance EMA + * = REWARD_COMPONENT_VAR_EMA_BASE + 1 = 362 (cf var at 362, trail 363, + * micro 364, opp_cost 365, bonus 366 — produced by SP4 reward_component_ema + * extended in B1b smoke-recovery). */ + int cf_others_var_base_slot ) { if (blockIdx.x != 0) return; if (threadIdx.x >= 6) return; - /* Block-load all 6 component magnitudes into shared mem. The popart - * axis (component 0) is sourced from a separate ISV slot than the - * rest — see the header for the B1b-fix-up rationale. */ - __shared__ float values[6]; - __shared__ float sum; + /* Block-load all 6 component magnitudes AND variances into shared + * mem. The popart axis (component 0) is sourced from a separate + * ISV slot than the rest — see the header for the B1b-fix-up + * rationale. The variance pair is identically split (slot 361 vs + * slots 362..366). */ + __shared__ float mag[6]; + __shared__ float var[6]; + __shared__ float sum_z; if (threadIdx.x == 0) { - values[0] = isv[popart_specific_slot]; + mag[0] = isv[popart_specific_slot]; + var[0] = isv[popart_var_slot]; } else { - /* Components 1..5 → ISV[cf_others_base_slot + (c - 1)] */ - values[threadIdx.x] = isv[cf_others_base_slot + (int)threadIdx.x - 1]; + /* Components 1..5 → ISV[cf_others_base_slot + (c - 1)] for mag, + * ISV[cf_others_var_base_slot + (c - 1)] for var. */ + mag[threadIdx.x] = isv[cf_others_base_slot + (int)threadIdx.x - 1]; + var[threadIdx.x] = isv[cf_others_var_base_slot + (int)threadIdx.x - 1]; } __syncthreads(); + /* Z-score per component. EPS_DIV (1e-6) is the Invariant-1 + * numerical-stability anchor matching SP4/SP5/SP9/SP10/SP11 + * conventions. Guards the impossible "all-zero variance" case + * during the first observations after a fold reset (Pearl A + * bootstraps each component variance EMA from its first + * observation, but timing of Pearls A+D fires across the 6 + * sources is per-slot, so a literal 0/0 is conceivable in the + * cold-start window — and on a constant input where var = 0 + * exactly). When var[c] = 0 the floor sets σ = EPS_DIV ⇒ + * z[c] = mag[c] × 1e6 (huge but well-defined); since all 6 + * z[c] scale linearly with the same EPS_DIV factor when all + * variances are zero, the ratios remain proportional and well- + * defined (each component still contributes its share of the + * total z-score envelope). */ + const float EPS_DIV = 1.0e-6f; + float z_local = mag[threadIdx.x] / fmaxf(sqrtf(var[threadIdx.x]), EPS_DIV); + + /* Block-publish per-thread z so thread 0 can compute the sum. We + * reuse the `mag` shared array as a temporary z-score store — + * `mag` is no longer needed past this point. */ + mag[threadIdx.x] = z_local; + __syncthreads(); + if (threadIdx.x == 0) { float s = 0.0f; #pragma unroll - for (int c = 0; c < 6; c++) s += values[c]; - /* EPS_DIV (1e-6) floor — Invariant-1 numerical-stability anchor - * matching SP4/SP5/SP9/SP10 conventions. Guards the impossible - * "all-zero magnitudes" case during the first observations after - * a fold reset (Pearl A bootstraps each component EMA from its - * first observation, but timing of Pearls A+D fires across the 6 - * sources is per-slot, so a literal 0/0 is conceivable in the - * cold-start window). */ - sum = fmaxf(s, 1e-6f); - /* Side-output mirror: popart-component-specific magnitude → - * ISV[359] via chained Pearls A+D singleton launch. Read directly - * from values[0] which thread 0 just loaded above (= isv[popart_ - * specific_slot]). Post-B1b-fix this is the popart-specific - * value, not the contaminated slot-63 total-reward value. */ - scratch_out[6] = values[0]; + for (int c = 0; c < 6; c++) s += mag[c]; + /* EPS_DIV (1e-6) floor on the sum — same Invariant-1 anchor. + * Guards the all-zero-z case (which can't happen given σ-floor + * above, but keeps the formula structurally robust). */ + sum_z = fmaxf(s, EPS_DIV); + /* Side-output mirror: popart-component RAW magnitude (not the + * z-score) → ISV[359] via chained Pearls A+D singleton launch. + * Slot 359 = `PNL_REWARD_MAGNITUDE_EMA_INDEX` is consumed by + * curiosity bounding + saboteur engagement thresholds — both + * need physical P&L scale, not z-score. Read directly from + * `var[0]`'s old neighbour `isv[popart_specific_slot]` — but + * we no longer have it as a shared (mag[0] is now z[0]). Re- + * read from ISV; thread 0 is the only writer to scratch_out[6]. */ + scratch_out[6] = isv[popart_specific_slot]; } __syncthreads(); /* Each of threads 0..5 writes its own ratio. */ - scratch_out[threadIdx.x] = values[threadIdx.x] / sum; + scratch_out[threadIdx.x] = mag[threadIdx.x] / sum_z; if (threadIdx.x == 0) { __threadfence_system(); diff --git a/crates/ml/src/cuda_pipeline/sp11_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp11_isv_slots.rs index 367dad5e2..5d22cd1ab 100644 --- a/crates/ml/src/cuda_pipeline/sp11_isv_slots.rs +++ b/crates/ml/src/cuda_pipeline/sp11_isv_slots.rs @@ -1,6 +1,6 @@ //! SP11 — Reward-subsystem controller ISV slot constants. //! -//! Slots [340..361) populated by the SP11 controller chain: +//! Slots [340..367) populated by the SP11 controller chain: //! - 6 component weights [340..346) //! - 4 controller scalar outputs [346..350) //! - 2 val-sharpe canaries [350..352) @@ -8,6 +8,8 @@ //! - 1 saboteur engagement (358) //! - 1 PnL magnitude EMA (359) — total reward magnitude (PopArt's input) //! - 1 PopArt-component-specific magnitude EMA (360) — B1b fix-up +//! - 6 per-reward-component variance EMAs [361..367) — B1b smoke-recovery +//! (z-score normalization for the mag-ratio canary) //! //! Spec: docs/superpowers/specs/2026-05-04-sp11-reward-as-controlled-subsystem.md //! Pearl: pearl_reward_as_controlled_subsystem.md @@ -64,6 +66,39 @@ pub const PNL_REWARD_MAGNITUDE_EMA_INDEX: usize = 359; // §4 amendment "Why slot 360". pub const POPART_COMPONENT_MAG_EMA_INDEX: usize = 360; +/// Per-reward-component variance EMA base. 6 slots [361..367) for +/// {popart, cf, trail, micro, opp_cost, bonus} variance EMAs. Computed +/// alongside their magnitude EMAs (slots 360, 64..68) via Welford's +/// online algorithm in popart_component_ema_kernel + reward_component_ema_kernel. +/// Read by reward_component_mag_ratio_compute_kernel for z-score +/// normalization (z[c] = mag[c] / max(sqrt(var[c]), EPS_DIV)). +/// +/// Why z-score (B1b smoke-recovery, 2026-05-04): the linear mag-ratio +/// formula `winner_weight = mag_ratio` amplified whichever reward +/// component had the largest intrinsic magnitude. Popart (trade P&L) is +/// O(100) per fire; cf/trail/micro/opp_cost/bonus are O(0.1-2). Result: +/// `w_pop` saturated toward MAX_WEIGHT, curiosity_b exploded, sharpe +/// collapsed. Z-score makes ratios scale-invariant — every component +/// contributes "shares of own scale" rather than "shares of absolute +/// magnitude". The controller's redistribution logic works correctly +/// when it sees ratios ≈ 1/6 each at stability. +/// +/// Layout: +/// [361] = popart variance EMA (paired with mag at slot 360) +/// [362] = cf variance EMA (paired with mag at slot 64) +/// [363] = trail variance EMA (paired with mag at slot 65) +/// [364] = micro variance EMA (paired with mag at slot 66) +/// [365] = opp_cost variance EMA (paired with mag at slot 67) +/// [366] = bonus variance EMA (paired with mag at slot 68) +/// +/// FoldReset sentinel 0; Pearl A bootstraps from the first observation +/// alongside the magnitude EMAs (the magnitudes already paired with +/// their per-fold sentinel-bootstrap path; the variance partners share +/// it). NOT cross-fold persistent — variances reset each fold so the +/// canary tracks the new fold's reward-component scale distribution. +pub const REWARD_COMPONENT_VAR_EMA_BASE: usize = 361; +pub const REWARD_COMPONENT_VAR_EMA_COUNT: usize = 6; + // Invariant-1 anchors (rate-limiters, NOT regime thresholds). // // Per `feedback_isv_for_adaptive_bounds`: every adaptive bound lives on @@ -82,7 +117,7 @@ pub const SP11_WEIGHT_FLOOR_FRACTION: f32 = 0.5; pub const SP11_ENGAGEMENT_FLOOR: f32 = 0.1; pub const SP11_SLOT_BASE: usize = 340; -pub const SP11_SLOT_END: usize = 361; +pub const SP11_SLOT_END: usize = 367; pub const SP11_PRODUCER_COUNT: usize = SP11_SLOT_END - SP11_SLOT_BASE; /// Salt added to `config.seed` when seeding the novelty SimHash projection diff --git a/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs index 319b0797b..ce1765ccf 100644 --- a/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs +++ b/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs @@ -18,10 +18,11 @@ //! 321..330 SP8 (Fix 36) train_active_frac canary + LB_MAX_BUDGET (9 slots, fold-reset) //! 330..339 SP9 (Fix 37) Kelly warmup floor + EMA targets + eval_dist (9 slots, fold-reset) //! 339..340 SP10 (Fix 38) eval Thompson temperature (1 slot, fold-reset) -//! 340..361 SP11 (Fix 39) reward-subsystem controller (21 slots, fold-reset) — see sp11_isv_slots -//! [includes B1b fix-up slot 360 = POPART_COMPONENT_MAG_EMA_INDEX] +//! 340..367 SP11 (Fix 39) reward-subsystem controller (27 slots, fold-reset) — see sp11_isv_slots +//! [includes B1b fix-up slot 360 = POPART_COMPONENT_MAG_EMA_INDEX +//! and B1b smoke-recovery 6 var EMAs at 361..367 for z-score normalization] //! -//! Total: 185 new SP5 ISV slots (52 + 24 + 20 + 4 + 4 + 6 + 4 + 4 + 3 + 16 + 8 + 9 + 9 + 1 + 21). +//! Total: 191 new SP5 ISV slots (52 + 24 + 20 + 4 + 4 + 6 + 4 + 4 + 3 + 16 + 8 + 9 + 9 + 1 + 27). // Re-export SP11 slot constants so existing consumers can import every // slot constant from a single location (`crate::cuda_pipeline::sp5_isv_slots::*`). @@ -303,7 +304,7 @@ pub const EVAL_THOMPSON_TEMP_INDEX: usize = 339; // [1] eval Thompson se // handled by the existing bulk memset of `wiener_state_buf` covered by the // `sp4_wiener_state` registry entry's dispatch arm. -pub const SP5_SLOT_END: usize = 361; +pub const SP5_SLOT_END: usize = 367; /// Wiener-buffer producer-count constant. Sizes `wiener_state_buf` via /// `(SP4_PRODUCER_COUNT + SP5_PRODUCER_COUNT) * SP4_WIENER_FLOATS_PER_SLOT`. @@ -382,23 +383,28 @@ pub const SP5_SLOT_END: usize = 361; /// two outputs); FoldReset (sentinel 0; Pearl A bootstrap on first /// observation). /// -/// SP11 (2026-05-04, Fix 39): allocates 21 new SP11 ISV slots at -/// ISV[340..361) — 6 reward component weights + 4 controller scalar +/// SP11 (2026-05-04, Fix 39): allocates 27 new SP11 ISV slots at +/// ISV[340..367) — 6 reward component weights + 4 controller scalar /// outputs + 11 canaries (val-sharpe Δ/var, 6 mag-ratios, saboteur -/// engagement, PnL EMA, popart-component magnitude EMA). Unique-slot -/// count grows 164 → 185; the linear span (and `SP5_PRODUCER_COUNT`) -/// grows 166 → 187. The 21st slot — `POPART_COMPONENT_MAG_EMA_INDEX` -/// at 360 — is the B1b fix-up (2026-05-04) for the PopArt-input vs -/// popart-component-magnitude overload (see spec §4 amendment "Why -/// slot 360"). Drives the reward subsystem controller per -/// `pearl_reward_as_controlled_subsystem.md` + -/// `pearl_controller_anchors_isv_driven.md`. All 21 slots are FoldReset -/// (sentinel 0; Pearl A bootstrap on first observation). Constants for -/// the SP11 block live in `crate::cuda_pipeline::sp11_isv_slots` and are -/// re-exported above. -pub const SP5_PRODUCER_COUNT: usize = 187; -// linear span = SP5_SLOT_END - SP5_SLOT_BASE = 361 - 174 = 187 wiener triples -// unique-slot count = 185 (52 per-branch + 24 Adam + 20 IQN τ + 4 trail +/// engagement, PnL EMA, popart-component magnitude EMA) + 6 per- +/// reward-component variance EMAs (B1b smoke-recovery 2026-05-04 for +/// the z-score normalised mag-ratio canary). Unique-slot count grows +/// 164 → 191; the linear span (and `SP5_PRODUCER_COUNT`) grows +/// 166 → 193. The 21st slot — `POPART_COMPONENT_MAG_EMA_INDEX` at 360 — +/// is the B1b fix-up (2026-05-04) for the PopArt-input vs popart- +/// component-magnitude overload (see spec §4 amendment "Why slot 360"). +/// Slots 22..27 are the B1b smoke-recovery (2026-05-04, spec §4 +/// amendment lines 564-619): per-component variance EMAs at +/// `REWARD_COMPONENT_VAR_EMA_BASE=361..367) feeding the z-score +/// normaliser in `reward_component_mag_ratio_compute_kernel`. Drives +/// the reward subsystem controller per `pearl_reward_as_controlled_subsystem.md` +/// + `pearl_controller_anchors_isv_driven.md`. All 27 slots are +/// FoldReset (sentinel 0; Pearl A bootstrap on first observation). +/// Constants for the SP11 block live in `crate::cuda_pipeline::sp11_isv_slots` +/// and are re-exported above. +pub const SP5_PRODUCER_COUNT: usize = 193; +// linear span = SP5_SLOT_END - SP5_SLOT_BASE = 367 - 174 = 193 wiener triples +// unique-slot count = 191 (52 per-branch + 24 Adam + 20 IQN τ + 4 trail // + 4 num_atoms + 6 Kelly + 4 Layer D D1 PnL aggregation // + 4 Layer D D2 health composition // + 3 Layer D D3 training metrics EMA @@ -408,7 +414,8 @@ pub const SP5_PRODUCER_COUNT: usize = 187; // + 8 SP8 LB_MAX_BUDGET per-(head,branch) // + 9 SP9 Kelly cold-start warmup floor + targets + eval_dist // + 1 SP10 eval Thompson temperature -// + 21 SP11 reward-subsystem controller [includes B1b fix-up slot 360]) +// + 27 SP11 reward-subsystem controller [includes B1b fix-up +// slot 360 + B1b smoke-recovery 6 per-component var EMAs 361..367]) // ── Convenience accessors ──────────────────────────────────────────── #[inline] pub const fn atom_v_center(b: usize) -> usize { ATOM_V_CENTER_BASE + b } @@ -472,7 +479,8 @@ pub const SP5_LAYOUT_FINGERPRINT_FRAGMENT: &str = REWARD_COMPONENT_MAG_RATIO_BASE=352;\ SABOTEUR_ENGAGEMENT_RATE=358;PNL_REWARD_MAGNITUDE_EMA=359;\ POPART_COMPONENT_MAG_EMA=360;\ - ISV_TOTAL_DIM=361"; + REWARD_COMPONENT_VAR_EMA_BASE=361;\ + ISV_TOTAL_DIM=367"; #[cfg(test)] mod tests { @@ -580,35 +588,38 @@ mod tests { // SP10 (Fix 38): eval Thompson selector temperature slots.insert(EVAL_THOMPSON_TEMP_INDEX); - // SP11 (Fix 39): include the 21 new reward-subsystem controller slots. + // SP11 (Fix 39): include the 27 new reward-subsystem controller slots. for s in SP11_SLOT_BASE..SP11_SLOT_END { slots.insert(s); } - // 1. Exactly 185 unique slots (164 pre-SP11 + 21 SP11). - assert_eq!(slots.len(), 185, "expected 185 unique slots, got {}", slots.len()); + // 1. Exactly 191 unique slots (164 pre-SP11 + 27 SP11). + assert_eq!(slots.len(), 191, "expected 191 unique slots, got {}", slots.len()); // 2. Min slot is SP5_SLOT_BASE = 174. assert_eq!(*slots.iter().min().unwrap(), 174); - // 3. Max slot is SP5_SLOT_END - 1 = 360. - assert_eq!(*slots.iter().max().unwrap(), 360); + // 3. Max slot is SP5_SLOT_END - 1 = 366. + assert_eq!(*slots.iter().max().unwrap(), 366); // 4. Intentional carve-out gap (278, 279) is absent. assert!(!slots.contains(&278), "slot 278 must be absent (carve-out gap)"); assert!(!slots.contains(&279), "slot 279 must be absent (carve-out gap)"); - // 5. Set equals {174..278} ∪ {280..361} — no holes (other than the + // 5. Set equals {174..278} ∪ {280..367} — no holes (other than the // carve-out gap 278..280), no overlaps. Layer D D1 extended the // upper end 286 → 290; D2 extended it 290 → 294; D3 extends it // 294 → 297; SP7 T1 extended it 297 → 313; SP7 activation-flag // fix extends it 313 → 321; SP8 (Fix 36) extends it 321 → 330; // SP9 (Fix 37) extends it 330 → 339; SP10 (Fix 38) extends it // 339 → 340; SP11 (Fix 39) extends it 340 → 360; SP11 B1b fix-up - // (2026-05-04) extends it 360 → 361 with POPART_COMPONENT_MAG_EMA_INDEX. - let expected: HashSet = (174..278).chain(280..361).collect(); + // (2026-05-04) extends it 360 → 361 with POPART_COMPONENT_MAG_EMA_INDEX; + // SP11 B1b smoke-recovery (2026-05-04) extends it 361 → 367 with + // `REWARD_COMPONENT_VAR_EMA_BASE..+6` for the z-score normalised + // mag-ratio canary. + let expected: HashSet = (174..278).chain(280..367).collect(); assert_eq!(slots, expected, - "slot set does not match expected {{174..278}} ∪ {{280..360}}"); + "slot set does not match expected {{174..278}} ∪ {{280..367}}"); } #[test] @@ -635,11 +646,11 @@ mod tests { // is cross-fold-persistent — the contracts must remain disjoint). assert!(PNL_TOTAL_INDEX > LOSS_RATE_SMOOTH_INDEX); // SP5_SLOT_END must reflect the post-SP8 end-of-block. - assert_eq!(SP5_SLOT_END, 361); + assert_eq!(SP5_SLOT_END, 367); // SP5_PRODUCER_COUNT is the wiener-buffer linear span (slot-range // width including the 2-slot carve-out gap), NOT the unique-slot // count. See SP5_PRODUCER_COUNT docstring for the rationale. - assert_eq!(SP5_PRODUCER_COUNT, 187); + assert_eq!(SP5_PRODUCER_COUNT, 193); assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE); } @@ -658,7 +669,7 @@ mod tests { // 4-slot block is internally contiguous. assert_eq!(GRAD_NORM_NORM_INDEX - HEALTH_SCORE_INDEX, 3); // SP5_SLOT_END must reflect the post-SP8 end-of-block. - assert_eq!(SP5_SLOT_END, 361); + assert_eq!(SP5_SLOT_END, 367); // SP5_PRODUCER_COUNT linear-span check matches the new end. assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE); } @@ -679,12 +690,12 @@ mod tests { assert_eq!(MAX_DD_EMA_INDEX - TRAINING_SHARPE_EMA_INDEX, 1); assert_eq!(LOW_DD_RATIO_INDEX - MAX_DD_EMA_INDEX, 1); // SP5_SLOT_END must reflect the post-SP8 end-of-block. - assert_eq!(SP5_SLOT_END, 361); + assert_eq!(SP5_SLOT_END, 367); // SP5_PRODUCER_COUNT linear-span check matches the new end. The // wiener buffer must cover the entire linear span — including the // 2-slot carve-out gap (278..280) and the Pearl 6 reserved-but- // unused 6-float block (slots 280..286 don't call apply_pearls). - assert_eq!(SP5_PRODUCER_COUNT, 187); + assert_eq!(SP5_PRODUCER_COUNT, 193); assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE); } @@ -704,7 +715,7 @@ mod tests { assert_eq!(lb_c51_active(b), LB_C51_ACTIVE_BASE + b); } // SP5_SLOT_END / SP5_PRODUCER_COUNT cover the new range. - assert_eq!(SP5_SLOT_END, 361); + assert_eq!(SP5_SLOT_END, 367); assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE); } @@ -724,7 +735,7 @@ mod tests { assert_eq!(lb_max_budget_c51(b), LB_MAX_BUDGET_C51_BASE + b); } // Layout: 1 + 4 + 4 = 9 contiguous slots (321..330). - assert_eq!(SP5_SLOT_END, 361); + assert_eq!(SP5_SLOT_END, 367); assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE); } @@ -757,7 +768,7 @@ mod tests { // Strictly above the SP8 MAX_BUDGET block. assert!(KELLY_WARMUP_FLOOR_INDEX > LB_MAX_BUDGET_C51_BASE + 3); // Block end at ISV[340) post-SP10 (Fix 38). - assert_eq!(SP5_SLOT_END, 361); + assert_eq!(SP5_SLOT_END, 367); // Linear span matches producer count (no gaps in SP9 block). assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE); } @@ -771,9 +782,9 @@ mod tests { // Strictly above the SP9 eval_dist block (last entry at 338). assert!(EVAL_THOMPSON_TEMP_INDEX > EVAL_DIST_F_INDEX); // SP5_SLOT_END reflects the post-SP10 end-of-block. - assert_eq!(SP5_SLOT_END, 361); + assert_eq!(SP5_SLOT_END, 367); // Linear span matches producer count (single new slot at top). assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE); - assert_eq!(SP5_PRODUCER_COUNT, 187); + assert_eq!(SP5_PRODUCER_COUNT, 193); } } diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index ca518095b..53013bd97 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -900,6 +900,11 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "ISV[POPART_COMPONENT_MAG_EMA_INDEX=360] — SP11 Fix 39 B1b fix-up popart-component-specific magnitude EMA. Produced by `popart_component_ema_kernel` (block tree-reduce over the per-bar `popart_component_per_sample` GPU buffer populated by `experience_env_step` at the `r_popart = capped_pnl;` voluntary-exit branch); smoothed by Pearls A+D via chained `apply_pearls_ad_kernel` (n_slots=1). Consumed by the SP11 mag-ratio canary as the popart axis input — replaces ISV[REWARD_POPART_EMA_INDEX=63]'s overloaded role as 'popart's signal' with a popart-component-specific value. Pre-fix-up slot 63 was overloaded (PopArt's normalisation input AND popart-component magnitude collapsed to the same value because reward composition was inline accumulation); B1b's structural decomposition surfaced the contamination — controller redistributed weight toward popart (w_pop ≈ 2.0) based on a contaminated ratio → 10× sharpe drop in B1b smoke. Slot 63 unchanged (PopArt's input = total reward magnitude, pre-SP11 invariant preserved). FoldReset sentinel 0; Pearl A bootstraps from the first observation. Spec §4 amendment 'Why slot 360'.", }, + RegistryEntry { + name: "sp11_reward_component_var_emas", + category: ResetCategory::FoldReset, + description: "ISV[REWARD_COMPONENT_VAR_EMA_BASE=361..367) — SP11 Fix 39 B1b smoke-recovery (2026-05-04, spec §4 amendment 'Why z-score' lines 564-619) per-reward-component variance EMAs (6 contiguous slots: popart at 361, then cf 362, trail 363, micro 364, opp_cost 365, bonus 366). Slot 361 is produced alongside the popart-component magnitude EMA by the extended `popart_component_ema_kernel` via single-pass two-tree-reduce Welford (mean → var); slots 362..367 are produced alongside the non-popart magnitudes by the extended `reward_component_ema_kernel` via per-thread Welford (component 0/popart variance is split off because slot 63 is the pre-SP11 total-reward magnitude EMA, not popart-component magnitude — see slot-360 fix-up rationale). All 6 variances are smoothed by Pearls A+D via chained `apply_pearls_ad_kernel` (n_slots=1 for popart variance at slot 361, n_slots=5 for the contiguous block at slots 362..367). Consumed by `reward_component_mag_ratio_compute_kernel` for z-score normalisation: `z[c] = mag[c] / max(sqrt(var[c]), EPS_DIV)` then `ratio[c] = z[c] / sum_z`. Replaces the linear `winner_weight = mag_ratio` formula that amplified popart's intrinsic O(100) magnitude over the other components' O(0.1-2) — w_pop saturated to MAX_WEIGHT, curiosity_b exploded, sharpe collapsed in B1b smoke `smoke-test-6wd2c` on commit `61b2fa962`. The single registry entry covers the whole 6-slot block — the dispatch arm zeroes all 6 with one loop, matching the `sp11_reward_component_mag_ratios` block-reset pattern. FoldReset sentinel 0; Pearl A bootstraps the variance from the first observation alongside its paired magnitude.", + }, // 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 diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index f590e36b6..198558ad4 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -7342,6 +7342,32 @@ impl DQNTrainer { fused.trainer().write_isv_signal_at(POPART_COMPONENT_MAG_EMA_INDEX, 0.0); } } + // SP11 Fix 39 B1b smoke-recovery (2026-05-04, spec §4 amendment + // 'Why z-score' lines 564-619): per-reward-component variance + // EMAs at ISV[REWARD_COMPONENT_VAR_EMA_BASE=361..367). Block- + // reset all 6 slots so the new fold's first popart_component_ema + // and reward_component_ema launches see Pearl A's sentinel and + // first-observation-replace the variance from the new fold's + // reward-component samples (paired with the magnitude EMA's + // bootstrap path). Companion Wiener-state for offsets + // (SP4_PRODUCER_COUNT*3 + (361..367 - SP5_SLOT_BASE)*3) is + // reset by the bulk memset of `wiener_state_buf` in the + // `sp5_wiener_state` registry entry's dispatch arm — both halves + // of the Pearl A sentinel contract reset together per + // `feedback_no_partial_refactor.md`. + "sp11_reward_component_var_emas" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp5_isv_slots::{ + REWARD_COMPONENT_VAR_EMA_BASE, REWARD_COMPONENT_VAR_EMA_COUNT, + }; + for i in 0..REWARD_COMPONENT_VAR_EMA_COUNT { + fused.trainer().write_isv_signal_at( + REWARD_COMPONENT_VAR_EMA_BASE + i, + 0.0, + ); + } + } + } // SP11 Task A2 (2026-05-04): novelty visit-count hash table // reset arm — closes the A0 deferral per the registry entry's // docstring. 1M f32 slots, zero-filled via mapped-pinned host diff --git a/crates/ml/tests/sp11_producer_unit_tests.rs b/crates/ml/tests/sp11_producer_unit_tests.rs index 00522956a..4a5c6804c 100644 --- a/crates/ml/tests/sp11_producer_unit_tests.rs +++ b/crates/ml/tests/sp11_producer_unit_tests.rs @@ -35,6 +35,7 @@ use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKe use ml::cuda_pipeline::mapped_pinned::MappedF32Buffer; use ml::cuda_pipeline::sp11_isv_slots::{ PNL_REWARD_MAGNITUDE_EMA_INDEX, REWARD_COMPONENT_MAG_RATIO_BASE, + REWARD_COMPONENT_VAR_EMA_BASE, SABOTEUR_ENGAGEMENT_RATE_INDEX, VAL_SHARPE_DELTA_EMA_INDEX, VAL_SHARPE_VAR_EMA_INDEX, }; @@ -46,7 +47,10 @@ use ml::cuda_pipeline::sp11_isv_slots::{ /// /// Bumped 360 → 361 by SP11 B1b fix-up (2026-05-04, spec §4 amendment /// "Why slot 360") which adds POPART_COMPONENT_MAG_EMA_INDEX=360. -const ISV_TOTAL_DIM: usize = 361; +/// Bumped 361 → 367 by SP11 B1b smoke-recovery (2026-05-04, spec §4 +/// amendment "Why z-score" lines 564-619) which adds 6 per-component +/// variance EMA slots at REWARD_COMPONENT_VAR_EMA_BASE=361..367. +const ISV_TOTAL_DIM: usize = 367; /// REWARD_POPART_EMA_INDEX from `gpu_dqn_trainer.rs` (the SP4 reward- /// component magnitude EMA base slot, mirrored here as a const so the @@ -71,6 +75,11 @@ const SP11_MAG_RATIO_CUBIN: &[u8] = /// SP11 A2 (2026-05-04): controller cubin for the 3 oracle tests below. const SP11_REWARD_SUBSYSTEM_CONTROLLER_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/reward_subsystem_controller_kernel.cubin")); +/// SP11 B1b smoke-recovery (2026-05-04): popart_component_ema cubin — +/// extended with Welford variance pass. Used by the variance-correctness +/// oracle test below. +const SP11_POPART_COMPONENT_EMA_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/popart_component_ema_kernel.cubin")); /// Build a CUDA stream against device 0. Mirrors `make_test_stream` in /// `sp4_producer_unit_tests.rs` and `distributional_q_tests.rs`. @@ -117,7 +126,7 @@ fn val_sharpe_delta_first_observation_writes_raw_delta() { .expect("alloc scratch (2 f32)"); let isv = unsafe { MappedF32Buffer::new(ISV_TOTAL_DIM) } - .expect("alloc isv (361 f32)"); + .expect("alloc isv (367 f32)"); // ISV[VAL_SHARPE_DELTA_EMA_INDEX] is constructor-zero — sentinel state // for Pearl A's first-observation replacement. No host writes needed. @@ -180,7 +189,7 @@ fn saboteur_engagement_counts_above_threshold_only() { .expect("alloc scratch (1 f32)"); let mut isv = unsafe { MappedF32Buffer::new(ISV_TOTAL_DIM) } - .expect("alloc isv (361 f32)"); + .expect("alloc isv (367 f32)"); isv.host_slice_mut()[PNL_REWARD_MAGNITUDE_EMA_INDEX] = 1.0; let delta_dev = delta_reward.dev_ptr; @@ -213,19 +222,60 @@ fn saboteur_engagement_counts_above_threshold_only() { ); } -/// SP11 A1.3 + B1b fix-up — per-component magnitude-ratio normaliser -/// preserves Σ=1 and mirrors popart-component magnitude into scratch[6]. +/// Helper to launch the post-B1b-smoke-recovery mag-ratio kernel with +/// the four slot-arg signature: (popart_mag_slot=360, cf_others_mag_base=64, +/// popart_var_slot=361, cf_others_var_base=362). Synchronises before +/// returning. +fn launch_mag_ratio_kernel( + stream: &Arc, + f: &CudaFunction, + isv: &MappedF32Buffer, + scratch: &MappedF32Buffer, +) { + let isv_dev = isv.dev_ptr; + let scratch_dev = scratch.dev_ptr; + let popart_specific_slot_i32: i32 = POPART_COMPONENT_MAG_EMA_INDEX as i32; + let cf_others_base_slot_i32: i32 = (REWARD_POPART_EMA_INDEX + 1) as i32; + let popart_var_slot_i32: i32 = REWARD_COMPONENT_VAR_EMA_BASE as i32; + let cf_others_var_base_slot_i32: i32 = (REWARD_COMPONENT_VAR_EMA_BASE + 1) as i32; + + unsafe { + stream + .launch_builder(f) + .arg(&isv_dev) + .arg(&scratch_dev) + .arg(&popart_specific_slot_i32) + .arg(&cf_others_base_slot_i32) + .arg(&popart_var_slot_i32) + .arg(&cf_others_var_base_slot_i32) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (6, 1, 1), + shared_mem_bytes: 0, + }) + .expect("launch reward_component_mag_ratio_compute_kernel"); + } + stream.synchronize().expect("sync after mag_ratio launch"); +} + +/// SP11 A1.3 + B1b fix-up + B1b smoke-recovery — per-component +/// magnitude-ratio normaliser preserves Σ=1 (with z-score normalisation) +/// and mirrors popart-component RAW magnitude into scratch[6]. /// -/// Post-B1b-fix-up the kernel takes TWO slot indices: -/// popart_specific_slot = POPART_COMPONENT_MAG_EMA_INDEX = 360 -/// cf_others_base_slot = REWARD_POPART_EMA_INDEX + 1 = 64 +/// Post-B1b-smoke-recovery the kernel takes FOUR slot indices: +/// popart_specific_slot = POPART_COMPONENT_MAG_EMA_INDEX = 360 +/// cf_others_base_slot = REWARD_POPART_EMA_INDEX + 1 = 64 +/// popart_var_slot = REWARD_COMPONENT_VAR_EMA_BASE = 361 +/// cf_others_var_base_slot = REWARD_COMPONENT_VAR_EMA_BASE + 1 = 362 /// /// Inputs: ISV[POPART_COMPONENT_MAG_EMA_INDEX] = 1.0, /// ISV[REWARD_POPART_EMA_INDEX+1..+6) = [2.0, 3.0, 4.0, 5.0, 6.0] -/// (cf=64, trail=65, micro=66, opp_cost=67, bonus=68). Sum = 21. -/// Expected ratios = [1/21, 2/21, 3/21, 4/21, 5/21, 6/21]. Mirror at -/// scratch[6] = 1.0 (popart-component-specific magnitude). Σ ratios = -/// 1.0 ± 1e-5. +/// (cf=64, trail=65, micro=66, opp_cost=67, bonus=68). +/// All variances set to 1.0 ⇒ σ=1 ⇒ z[c] = mag[c] ⇒ z-score reduces to +/// raw magnitude. Sum = 21. Expected ratios = [1/21, 2/21, 3/21, 4/21, +/// 5/21, 6/21]. Mirror at scratch[6] = 1.0 (popart-component RAW +/// magnitude — NOT z-score; consumed by curiosity bound + saboteur +/// engagement which need physical scale). #[test] #[ignore = "requires GPU"] fn reward_component_mag_ratio_normalises_and_mirrors_popart() { @@ -236,7 +286,7 @@ fn reward_component_mag_ratio_normalises_and_mirrors_popart() { .expect("alloc scratch (7 f32 — 6 ratios + 1 popart mirror)"); let mut isv = unsafe { MappedF32Buffer::new(ISV_TOTAL_DIM) } - .expect("alloc isv (361 f32)"); + .expect("alloc isv (367 f32)"); { let s = isv.host_slice_mut(); // B1b fix-up: popart axis sourced from slot 360, not 63. @@ -247,28 +297,18 @@ fn reward_component_mag_ratio_normalises_and_mirrors_popart() { s[REWARD_POPART_EMA_INDEX + 3] = 4.0; // micro s[REWARD_POPART_EMA_INDEX + 4] = 5.0; // opp_cost s[REWARD_POPART_EMA_INDEX + 5] = 6.0; // bonus + // B1b smoke-recovery: variance = 1.0 for every component ⇒ σ=1 ⇒ + // z-score reduces to raw magnitude (verifying the legacy + // ratio shape works under the new formula when σ is unit). + s[REWARD_COMPONENT_VAR_EMA_BASE] = 1.0; // popart var + s[REWARD_COMPONENT_VAR_EMA_BASE + 1] = 1.0; // cf var + s[REWARD_COMPONENT_VAR_EMA_BASE + 2] = 1.0; // trail var + s[REWARD_COMPONENT_VAR_EMA_BASE + 3] = 1.0; // micro var + s[REWARD_COMPONENT_VAR_EMA_BASE + 4] = 1.0; // opp_cost var + s[REWARD_COMPONENT_VAR_EMA_BASE + 5] = 1.0; // bonus var } - let isv_dev = isv.dev_ptr; - let scratch_dev = scratch.dev_ptr; - let popart_specific_slot_i32: i32 = POPART_COMPONENT_MAG_EMA_INDEX as i32; - let cf_others_base_slot_i32: i32 = (REWARD_POPART_EMA_INDEX + 1) as i32; - - unsafe { - stream - .launch_builder(&f) - .arg(&isv_dev) - .arg(&scratch_dev) - .arg(&popart_specific_slot_i32) - .arg(&cf_others_base_slot_i32) - .launch(LaunchConfig { - grid_dim: (1, 1, 1), - block_dim: (6, 1, 1), - shared_mem_bytes: 0, - }) - .expect("launch reward_component_mag_ratio_compute_kernel"); - } - stream.synchronize().expect("sync after mag_ratio launch"); + launch_mag_ratio_kernel(&stream, &f, &isv, &scratch); let out = scratch.read_all(); let sum_in: f32 = 1.0 + 2.0 + 3.0 + 4.0 + 5.0 + 6.0; @@ -291,6 +331,9 @@ fn reward_component_mag_ratio_normalises_and_mirrors_popart() { ); // Popart mirror at scratch[6] — sourced from slot 360 post-B1b-fix. + // This is RAW magnitude (not z-score), per spec §3.4.2 / §3.3.1 + // — slot 359 consumers (curiosity bound + saboteur engagement) + // need physical P&L scale. assert!( (out[6] - 1.0_f32).abs() < 1e-6, "expected popart mirror = 1.0 (= isv[POPART_COMPONENT_MAG_EMA_INDEX=360]), got {}", @@ -298,6 +341,201 @@ fn reward_component_mag_ratio_normalises_and_mirrors_popart() { ); } +/// SP11 B1b smoke-recovery — z-score canary produces uniform 1/6 +/// ratios when all components have equal z-scores (equal mag, equal +/// variance ⇒ equal σ ⇒ equal z = mag/σ). +/// +/// Inputs: +/// ISV[360]=10, ISV[64..68]=10 (equal magnitudes, popart split). +/// ISV[361]=4, ISV[362..366]=4 (equal variances ⇒ σ=2 each). +/// z[c] = 10/2 = 5 ∀c ⇒ Σz = 30 ⇒ ratio[c] = 5/30 = 1/6 ∀c. +#[test] +#[ignore = "requires GPU"] +fn mag_ratio_z_score_uniform_inputs_yields_uniform_ratios() { + let stream = make_test_stream(); + let f = load_kernel(&stream, SP11_MAG_RATIO_CUBIN, "reward_component_mag_ratio_compute_kernel"); + + let scratch = unsafe { MappedF32Buffer::new(7) } + .expect("alloc scratch (7 f32)"); + + let mut isv = unsafe { MappedF32Buffer::new(ISV_TOTAL_DIM) } + .expect("alloc isv (367 f32)"); + { + let s = isv.host_slice_mut(); + // Magnitudes: popart at slot 360, cf/trail/micro/opp_cost/bonus at 64..68 — all 10.0. + s[POPART_COMPONENT_MAG_EMA_INDEX] = 10.0; + for c in 1..6 { + s[REWARD_POPART_EMA_INDEX + c] = 10.0; + } + // Variances: popart at slot 361, others at 362..366 — all 4.0 (σ=2). + s[REWARD_COMPONENT_VAR_EMA_BASE] = 4.0; + for c in 1..6 { + s[REWARD_COMPONENT_VAR_EMA_BASE + c] = 4.0; + } + } + + launch_mag_ratio_kernel(&stream, &f, &isv, &scratch); + + let out = scratch.read_all(); + let expected = 1.0_f32 / 6.0; + + for c in 0..6_usize { + assert!( + (out[c] - expected).abs() < 1e-5, + "component {c}: expected ratio 1/6 = {expected}, got {}", + out[c] + ); + } + + // Σ ratios = 1.0 (sanity). + let sum_out: f32 = out[0..6].iter().sum(); + assert!( + (sum_out - 1.0_f32).abs() < 1e-5, + "expected Σratios = 1.0, got {sum_out}" + ); +} + +/// SP11 B1b smoke-recovery — THE ACTUAL FIX VERIFICATION. +/// Asymmetric input (popart 100× bigger linear magnitude AND 100× bigger +/// variance) produces a balanced ratio under the z-score formula +/// instead of the linear formula's pathological 95% popart share. +/// +/// Inputs: +/// ISV[360]=100, ISV[64..68]=1 (popart 100× bigger linear) +/// ISV[361]=400, ISV[362..366]=4 (popart variance 100× bigger ⇒ +/// σ_popart=20 vs σ_others=2) +/// Z-scores: +/// z[0] = 100 / 20 = 5 +/// z[c] = 1 / 2 = 0.5 for c=1..5 +/// Σz = 5 + 5×0.5 = 7.5 +/// Ratios: +/// ratio[0] = 5 / 7.5 ≈ 0.6667 +/// ratio[c] = 0.5 / 7.5 ≈ 0.0667 (c=1..5) +/// +/// The pre-z-score linear formula would give: +/// ratio[0] = 100 / (100+5) ≈ 0.952 ← saturates w_pop +/// ratio[c] = 1 / (100+5) ≈ 0.0095 (c=1..5) +/// — exactly the B1b smoke pathology. +#[test] +#[ignore = "requires GPU"] +fn mag_ratio_z_score_asymmetric_input_yields_balanced_ratios() { + let stream = make_test_stream(); + let f = load_kernel(&stream, SP11_MAG_RATIO_CUBIN, "reward_component_mag_ratio_compute_kernel"); + + let scratch = unsafe { MappedF32Buffer::new(7) } + .expect("alloc scratch (7 f32)"); + + let mut isv = unsafe { MappedF32Buffer::new(ISV_TOTAL_DIM) } + .expect("alloc isv (367 f32)"); + { + let s = isv.host_slice_mut(); + s[POPART_COMPONENT_MAG_EMA_INDEX] = 100.0; + for c in 1..6 { + s[REWARD_POPART_EMA_INDEX + c] = 1.0; + } + s[REWARD_COMPONENT_VAR_EMA_BASE] = 400.0; // σ=20 + for c in 1..6 { + s[REWARD_COMPONENT_VAR_EMA_BASE + c] = 4.0; // σ=2 + } + } + + launch_mag_ratio_kernel(&stream, &f, &isv, &scratch); + + let out = scratch.read_all(); + + // ratio[0] = 5/7.5 = 0.6667; ratio[1..6] = 0.5/7.5 = 0.0667 each. + let expected_popart = 5.0_f32 / 7.5; + let expected_others = 0.5_f32 / 7.5; + + assert!( + (out[0] - expected_popart).abs() < 1e-4, + "popart ratio = {} (expected ≈ {expected_popart}); the linear formula \ + would have given ≈ 0.952 — z-score keeps the canary balanced.", + out[0] + ); + + for c in 1..6_usize { + assert!( + (out[c] - expected_others).abs() < 1e-4, + "component {c} ratio = {} (expected ≈ {expected_others})", + out[c] + ); + } + + // Σ = 1.0 still. + let sum_out: f32 = out[0..6].iter().sum(); + assert!( + (sum_out - 1.0_f32).abs() < 1e-5, + "expected Σratios = 1.0, got {sum_out}" + ); +} + +/// SP11 B1b smoke-recovery — EPS_DIV floor honored on zero variance +/// (cold start before Pearl A bootstraps any variance EMA). +/// +/// Inputs: equal magnitudes, all variances = 0 ⇒ σ_floor = EPS_DIV = +/// 1e-6 ⇒ z[c] = mag[c] / 1e-6 = huge but proportional. Ratios still +/// normalize correctly. +#[test] +#[ignore = "requires GPU"] +fn mag_ratio_z_score_zero_variance_eps_floor_honored() { + let stream = make_test_stream(); + let f = load_kernel(&stream, SP11_MAG_RATIO_CUBIN, "reward_component_mag_ratio_compute_kernel"); + + let scratch = unsafe { MappedF32Buffer::new(7) } + .expect("alloc scratch (7 f32)"); + + let mut isv = unsafe { MappedF32Buffer::new(ISV_TOTAL_DIM) } + .expect("alloc isv (367 f32)"); + { + let s = isv.host_slice_mut(); + // Magnitudes [1, 2, 3, 4, 5, 6] → Σmag = 21; under uniform σ-floor + // ratios reduce to (mag[c] / sum_mag). + s[POPART_COMPONENT_MAG_EMA_INDEX] = 1.0; + for c in 1..6 { + s[REWARD_POPART_EMA_INDEX + c] = (c + 1) as f32; + } + // All variances zero — kernel must apply EPS_DIV floor on σ. + s[REWARD_COMPONENT_VAR_EMA_BASE] = 0.0; + for c in 1..6 { + s[REWARD_COMPONENT_VAR_EMA_BASE + c] = 0.0; + } + } + + launch_mag_ratio_kernel(&stream, &f, &isv, &scratch); + + let out = scratch.read_all(); + + // Σ ratios = 1.0 ± epsilon — the canary stays well-defined under + // cold-start zero variance. + let sum_out: f32 = out[0..6].iter().sum(); + assert!( + (sum_out - 1.0_f32).abs() < 1e-4, + "expected Σratios ≈ 1.0 under zero-variance EPS_DIV floor, got {sum_out}" + ); + + // Each ratio ∈ [0, 1]. + for c in 0..6_usize { + assert!( + out[c] >= 0.0 && out[c] <= 1.0 + 1e-5, + "component {c} ratio = {} out of [0, 1]", + out[c] + ); + } + + // Under uniform EPS_DIV floor σ all components get the same scale + // multiplier ⇒ ratio[c] reduces to mag[c] / Σmag. + let sum_mag: f32 = 1.0 + 2.0 + 3.0 + 4.0 + 5.0 + 6.0; + for c in 0..6_usize { + let expected = (c as f32 + 1.0) / sum_mag; + assert!( + (out[c] - expected).abs() < 1e-4, + "component {c}: expected ratio {expected} (uniform EPS_DIV σ), got {}", + out[c] + ); + } +} + // ── SP11 Task A2 — reward-subsystem controller GPU oracle tests ────────── /// Populate ISV with the 5 controller canary inputs the kernel reads. @@ -397,7 +635,7 @@ fn controller_z_score_at_zero_yields_midpoint_outputs() { ); let mut isv = unsafe { MappedF32Buffer::new(ISV_TOTAL_DIM) } - .expect("alloc isv (361 f32)"); + .expect("alloc isv (367 f32)"); populate_controller_isv( &mut isv, /*delta*/ 0.0, @@ -487,7 +725,7 @@ fn controller_weights_renormalize_after_floor() { ); let mut isv = unsafe { MappedF32Buffer::new(ISV_TOTAL_DIM) } - .expect("alloc isv (361 f32)"); + .expect("alloc isv (367 f32)"); populate_controller_isv( &mut isv, /*delta*/ 1.0, // any non-zero @@ -576,7 +814,7 @@ fn controller_saboteur_post_clamp_holds_min() { ); let mut isv = unsafe { MappedF32Buffer::new(ISV_TOTAL_DIM) } - .expect("alloc isv (361 f32)"); + .expect("alloc isv (367 f32)"); populate_controller_isv( &mut isv, /*delta*/ -100.0, // extreme regression @@ -611,3 +849,98 @@ fn controller_saboteur_post_clamp_holds_min() { s[6] ); } + +// ── SP11 B1b smoke-recovery — popart_component_ema Welford variance ──── + +/// Helper to launch the popart_component_ema kernel with the standard +/// (popart_buf, scratch, n_samples) signature. BLOCK_SIZE=256. +fn launch_popart_component_ema( + stream: &Arc, + f: &CudaFunction, + popart_buf: &MappedF32Buffer, + scratch: &MappedF32Buffer, + n_samples: i32, +) { + let popart_dev = popart_buf.dev_ptr; + let scratch_dev = scratch.dev_ptr; + unsafe { + stream + .launch_builder(f) + .arg(&popart_dev) + .arg(&scratch_dev) + .arg(&n_samples) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 256 * std::mem::size_of::() as u32, + }) + .expect("launch popart_component_ema_kernel"); + } + stream.synchronize().expect("sync after popart_component_ema launch"); +} + +/// SP11 B1b smoke-recovery — popart_component_ema_kernel writes both +/// mean(|x|) AND var(|x|) via single-pass two-tree-reduce Welford. +/// +/// Constant-input case: input = [1.0, 1.0, 1.0, 1.0] ⇒ |x| = 1 ∀i ⇒ +/// mean = 1.0, var = 0.0 (no deviation). +#[test] +#[ignore = "requires GPU"] +fn popart_component_ema_welford_constant_input_zero_variance() { + let stream = make_test_stream(); + let f = load_kernel(&stream, SP11_POPART_COMPONENT_EMA_CUBIN, "popart_component_ema_kernel"); + + let popart = unsafe { MappedF32Buffer::new(4) } + .expect("alloc popart buf (4 f32)"); + popart.write_from_slice(&[1.0_f32, 1.0, 1.0, 1.0]); + + let scratch = unsafe { MappedF32Buffer::new(2) } + .expect("alloc scratch (2 f32 — mean + var)"); + + launch_popart_component_ema(&stream, &f, &popart, &scratch, 4); + + let out = scratch.read_all(); + assert!( + (out[0] - 1.0_f32).abs() < 1e-6, + "expected mean = 1.0 on constant input, got {}", + out[0] + ); + assert!( + out[1].abs() < 1e-6, + "expected var = 0.0 on constant input, got {}", + out[1] + ); +} + +/// SP11 B1b smoke-recovery — Welford variance correctness on bimodal +/// input. +/// +/// Input = [0.0, 2.0, 0.0, 2.0] ⇒ |x| = [0, 2, 0, 2] ⇒ mean = 1.0; +/// deviations = [-1, 1, -1, 1] ⇒ var = (1 + 1 + 1 + 1) / 4 = 1.0. +#[test] +#[ignore = "requires GPU"] +fn popart_component_ema_welford_bimodal_input_unit_variance() { + let stream = make_test_stream(); + let f = load_kernel(&stream, SP11_POPART_COMPONENT_EMA_CUBIN, "popart_component_ema_kernel"); + + let popart = unsafe { MappedF32Buffer::new(4) } + .expect("alloc popart buf (4 f32)"); + popart.write_from_slice(&[0.0_f32, 2.0, 0.0, 2.0]); + + let scratch = unsafe { MappedF32Buffer::new(2) } + .expect("alloc scratch (2 f32 — mean + var)"); + + launch_popart_component_ema(&stream, &f, &popart, &scratch, 4); + + let out = scratch.read_all(); + assert!( + (out[0] - 1.0_f32).abs() < 1e-6, + "expected mean = 1.0 on bimodal input, got {}", + out[0] + ); + assert!( + (out[1] - 1.0_f32).abs() < 1e-5, + "expected var = 1.0 on bimodal input (Welford), got {}", + out[1] + ); +} diff --git a/docs/isv-slots.md b/docs/isv-slots.md index e58378eb7..38ef948ab 100644 --- a/docs/isv-slots.md +++ b/docs/isv-slots.md @@ -286,6 +286,7 @@ htod_copy/dtoh_sync_copy/alloc_zeros per | [358] | `SABOTEUR_ENGAGEMENT_RATE_INDEX` | scalar canary | Per-bar saboteur engagement rate EMA | | [359] | `PNL_REWARD_MAGNITUDE_EMA_INDEX` | scalar canary | `|reward_total|` EMA (curiosity-bound scale factor) | | [360] | `POPART_COMPONENT_MAG_EMA_INDEX` | scalar canary | popart-component-specific magnitude EMA (B1b fix-up; mag-ratio canary popart axis input) | +| [361..367) | `REWARD_COMPONENT_VAR_EMA_BASE` (+6) | per-component canary | Per-reward-component variance EMA (B1b smoke-recovery z-score normalisation; popart var at 361, then cf/trail/micro/opp_cost/bonus at 362..367) | Invariant-1 fraction constants (rate-limiter anchors per spec §3.4.2): `SP11_EPS_DIV=1e-6`, `SP11_WEIGHT_HARD_FLOOR=0.01`, @@ -310,6 +311,23 @@ unchanged (PopArt's input = total reward magnitude, pre-SP11 invariant preserved); slot 360 is the new dedicated popart-component magnitude produced by `popart_component_ema_kernel`. ISV total: 360 → 361. +**B1b smoke-recovery (2026-05-04, spec §4 amendment "Why z-score" lines +564-619):** ISV[361..367) = `REWARD_COMPONENT_VAR_EMA_BASE..+6` added +to drive z-score normalisation in `reward_component_mag_ratio_compute_kernel`. +Pre-z-score the linear `winner_weight = mag_ratio` formula amplified +popart's intrinsic O(100) magnitude over the other 5 components' +O(0.1-2). Result: w_pop saturated toward MAX_WEIGHT, curiosity_b +exploded, sharpe collapsed in B1b smoke `smoke-test-6wd2c` on commit +`61b2fa962`. Z-score makes ratios scale-invariant — +z[c] = mag[c] / max(sqrt(var[c]), EPS_DIV); ratio[c] = z[c] / Σz. +Variance EMAs are produced via Welford's online algorithm in the +extended `popart_component_ema_kernel` (slot 361) and +`reward_component_ema_kernel` (slots 362..367); the canary now takes +4 slot indices (popart mag/var split + cf-others mag/var bases). The +mirror at scratch[6] still emits popart RAW magnitude (slot 359 input +for curiosity bound + saboteur engagement which need physical scale). +ISV total: 361 → 367. + All 21 slots are FoldReset (sentinel 0; Pearl A bootstraps from first producer launch on each fold per `pearl_first_observation_bootstrap.md`). After A2 all 20 slots have producers; B1b wires the on-policy reward