fix(sp11): B1b smoke-recovery — z-score normalization for mag-ratio canary
Linear magnitude ratios in reward_component_mag_ratio_compute_kernel amplified popart's intrinsic O(100) magnitude over the other 5 components' O(0.1-2) magnitudes, causing controller to saturate w_pop toward MAX_WEIGHT regardless of actual signal quality. Replaced with z-score: z[c] = mag[c] / max(sqrt(var[c]), EPS_DIV). 6 new ISV slots [361..367) for per-component variance EMAs computed via Welford's online algorithm in extended popart_component_ema_kernel and reward_component_ema_kernel. Atomic per feedback_no_partial_refactor: slot allocation + state-reset registry + 2 producer kernels + canary signature + launcher Pearls A+D + tests. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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,
|
||||
)?;
|
||||
|
||||
@@ -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(())
|
||||
|
||||
@@ -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 <cuda_runtime.h>
|
||||
@@ -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();
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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 */
|
||||
}
|
||||
|
||||
@@ -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();
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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<usize> = (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<usize> = (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);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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<CudaStream>,
|
||||
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<CudaStream>,
|
||||
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::<f32>() 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]
|
||||
);
|
||||
}
|
||||
|
||||
@@ -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
|
||||
|
||||
Reference in New Issue
Block a user