gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

fix(sp11): B1b follow-up — add slot 360 for popart-component mag EMA

Browse Source 
Per spec §4 amendment at 52c0b7521 on main: B1b smoke surfaced that
SP11 mag-ratio canary was reading slot 63 (REWARD_POPART_EMA_INDEX)
which is overloaded — pre-SP11 PopArt's normalization input (total
reward mag EMA) was the same value as popart-component magnitude
because composition was inline accumulation. B1b decomposition exposed
the overload; controller emitted w_pop ≈ 2.0 based on contaminated
ratio → 10× sharpe drop in smoke.

Resolution:
- Allocate ISV slot 360 = POPART_COMPONENT_MAG_EMA_INDEX
- Add popart_component_per_sample mapped-pinned buffer + write site
  in experience_env_step at the r_popart assignment
- New popart_component_ema_kernel.cu writes slot 360 (single-block
  block-tree-reduce per feedback_no_atomicadd)
- mag-ratio canary kernel signature changes from single
  popart_ema_base_slot to (popart_specific_slot, cf_others_base_slot)
  pair so it reads non-contiguous slot 360 + slots 64..68
- Reset registry: sp11_popart_component_mag_ema entry + dispatch arm
- Slot 63 (PopArt's input) UNCHANGED — pre-SP11 invariant preserved

ISV total: 360 → 361. SP5_SLOT_END = 361. SP5_PRODUCER_COUNT = 187.

cargo check + build clean; SP11 GPU oracle tests pass (6/6 including
updated mag_ratio test with 2 slot-index args); sp5_isv_slots layout
tests pass (10/10 with 185 unique slots / 187 linear span); state
reset registry tests pass (4/4 with new sp11_popart_component_mag_ema
entry + dispatch arm). Local multi_fold_convergence smoke gated on
data volume (175k bars on local fxcache vs 10-month walk-forward
requirement); validation deferred to L40S Argo run on PVC data per
the spec's pass criterion.
This commit is contained in:
jgrusewski
2026-05-04 10:29:47 +02:00
parent 034ba16801
commit 5e16b67ca6
12 changed files with 660 additions and 125 deletions
Download Patch File Download Diff File Expand all files Collapse all files

9
crates/ml/build.rs
View File

@@ -564,6 +564,15 @@ fn main() {
"val_sharpe_delta_compute_kernel.cu",
"saboteur_engagement_compute_kernel.cu",
"reward_component_mag_ratio_compute_kernel.cu",
// SP11 Fix 39 B1b fix-up (2026-05-04): popart-component-specific
// magnitude EMA producer. Reads the per-bar `r_popart` buffer
// populated by `experience_env_step` and block-tree-reduces
// mean(|x|) into ISV[POPART_COMPONENT_MAG_EMA_INDEX=360] via
// chained Pearls A+D. Resolves the slot-63 PopArt-input vs
// popart-component-magnitude overload that B1b smoke surfaced
// (contamination → controller over-weighed popart → 10× sharpe
// drop). Spec §4 amendment "Why slot 360".
"popart_component_ema_kernel.cu",
// SP11 Fix 39 (2026-05-04, Task A2): reward-subsystem controller +
// SimHash novelty buffer. Three new kernels in this layer:
// 1. reward_subsystem_controller — main producer (5 canaries → 10

45
crates/ml/src/cuda_pipeline/experience_kernels.cu
View File

@@ -1807,7 +1807,30 @@ extern "C" __global__ void experience_env_step(
* host-side — kernel-to-kernel dataflow only, so the no-CudaSlice
* rule from feedback_no_htod_htoh_only_mapped_pinned does not apply
* (that rule covers CPU↔GPU paths). */
float* __restrict__ saboteur_delta_reward_per_sample
float* __restrict__ saboteur_delta_reward_per_sample,
/* SP11 Fix 39 B1b fix-up (2026-05-04): per-bar popart-component-
* specific reward magnitude for the reward-as-controlled-subsystem
* chain. Layout [N*L] f32 mapped-pinned, row-major (i, t). Producer:
* this kernel writes the actual `r_popart` value (vol-normalized
* capped P&L on voluntary-exit bars; 0 on every other bar) — NOT
* the composed total reward. The kernel `popart_component_ema_kernel`
* consumes this buffer and computes mean(|x|) into ISV[POPART_
* COMPONENT_MAG_EMA_INDEX=360] via Pearls A+D.
*
* Why this is separate from `reward_components_per_sample[+0]`:
* pre-B1b that slot was overloaded (post-composition total reward
* AND popart-component magnitude collapsed to the same value because
* composition was inline accumulation). B1b's structural decomposition
* surfaces the contamination — slot 63 = REWARD_POPART_EMA_INDEX
* stays as PopArt's input (= total reward magnitude); slot 360 =
* POPART_COMPONENT_MAG_EMA_INDEX is the new popart-component-specific
* magnitude. Spec §4 amendment "Why slot 360".
*
* NULL = caller hasn't wired the buffer yet (early init); the kernel
* guards every write with the NULL check. Pure GPU buffer (the
* trainer holds it as a `MappedF32Buffer`); reads happen only inside
* `popart_component_ema_kernel` (kernel-to-kernel dataflow). */
float* __restrict__ popart_component_per_sample
) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i >= N) return;
@@ -1894,6 +1917,16 @@ extern "C" __global__ void experience_env_step(
if (saboteur_delta_reward_per_sample != NULL) {
saboteur_delta_reward_per_sample[out_off] = 0.0f;
}
/* SP11 Fix 39 B1b fix-up (2026-05-04): default per-bar popart-component
* magnitude to 0. Real value is overwritten only on voluntary-exit
* bars (where `r_popart = capped_pnl` — see line ~2680). All other
* bars (mid-trade, trail-triggered exits, hold/flat, early returns)
* keep the 0 default — popart only fires on voluntary segment_complete
* exits, so the kernel intentionally writes 0 elsewhere. NULL = buffer
* not wired yet (early init); the kernel guards every write. */
if (popart_component_per_sample != NULL) {
popart_component_per_sample[out_off] = 0.0f;
}
{
long long cf_off_default = out_off + (long long)N * L;
cf_flip_per_sample[cf_off_default] = 0;
@@ -2678,6 +2711,16 @@ extern "C" __global__ void experience_env_step(
r_trail = capped_pnl;
} else {
r_popart = capped_pnl;
/* SP11 Fix 39 B1b fix-up (2026-05-04): mirror r_popart to the
* dedicated per-bar buffer feeding ISV[POPART_COMPONENT_MAG_EMA_
* INDEX=360]. Voluntary-exit branch only (the trail-triggered
* branch above leaves popart-component at the entry-point 0
* default). The mag-ratio canary now reads slot 360 for the
* popart axis instead of slot 63's overloaded total-reward
* value. Spec §4 amendment "Why slot 360". */
if (popart_component_per_sample != NULL) {
popart_component_per_sample[out_off] = r_popart;
}
}
reward = capped_pnl;
/* C.2 attribution: rc[2] now carries the trail signal when the

247
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
View File

@@ -565,6 +565,18 @@ static SP11_MAG_RATIO_COMPUTE_CUBIN: &[u8] =
static SP11_REWARD_SUBSYSTEM_CONTROLLER_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/reward_subsystem_controller_kernel.cubin"));
/// SP11 Fix 39 B1b fix-up (2026-05-04): popart-component-specific
/// magnitude EMA producer. Single block, BLOCK_SIZE=256; reads the
/// per-bar `r_popart` mapped-pinned buffer populated by
/// `experience_env_step` and block-tree-reduces mean(|x|) into
/// `scratch[SCRATCH_SP11_POPART_COMPONENT_EMA]`. Downstream
/// `apply_pearls_ad_kernel` (n_slots=1) smooths into
/// ISV[POPART_COMPONENT_MAG_EMA_INDEX=360]. Resolves the slot-63
/// PopArt-input vs popart-component-magnitude overload that B1b smoke
/// surfaced. Spec §4 amendment "Why slot 360".
static SP11_POPART_COMPONENT_EMA_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/popart_component_ema_kernel.cubin"));
/// SP11 Fix 39 (2026-05-04, Task A2): one-shot GPU init for the SimHash
/// projection matrix. Philox-driven (deterministic from a host-passed seed
/// derived from the SP11 salt) writes ±1 values into a 42×16 = 672-element
@@ -971,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 = 360; // SP5 + Layer D D1+D2+D3 + SP7 + SP8 + SP9 + SP10 + SP11: 173 + 187 (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..360) — 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])
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])
/// 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).
@@ -1289,8 +1301,11 @@ pub const SP5_WIENER_TOTAL_FLOATS: usize =
/// body for explicit per-output traceability with the SP11 spec
/// table; functionally equivalent to a single n_slots=4 launch.
/// See `launch_sp11_reward_subsystem_controller` for the dispatch.
/// Combined: 259 + 6 + 1 + 10 + 10 = 286 scratch slots [0..286).
pub const SP5_SCRATCH_TOTAL: usize = 286;
/// 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;
/// SP5 Layer D Task D1 (rewrite, 2026-05-02): scratch index base for
/// `pnl_aggregation_update`.
@@ -1432,6 +1447,13 @@ pub const SCRATCH_SP11_MAG_RATIO_BASE: usize = 269; // [269..276) — f
/// matched-stride order.
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.
pub const SCRATCH_SP11_POPART_COMPONENT_EMA: usize = 286; // [286..287) → ISV[POPART_COMPONENT_MAG_EMA_INDEX=360]
/// 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 →
@@ -1966,7 +1988,8 @@ const fn layout_fingerprint_seed() -> &'static [u8] {
VAL_SHARPE_DELTA_EMA=350;VAL_SHARPE_VAR_EMA=351;\
REWARD_COMPONENT_MAG_RATIO_BASE=352;\
SABOTEUR_ENGAGEMENT_RATE=358;PNL_REWARD_MAGNITUDE_EMA=359;\
ISV_TOTAL_DIM=360;\
POPART_COMPONENT_MAG_EMA=360;\
ISV_TOTAL_DIM=361;\
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;\
@@ -4772,10 +4795,14 @@ pub struct GpuDqnTrainer {
/// apply_pearls_ad_kernel launch → ISV[SABOTEUR_ENGAGEMENT_RATE_INDEX=358].
/// Loaded from `saboteur_engagement_compute_kernel.cubin`.
sp11_saboteur_engagement_compute_kernel: CudaFunction,
/// SP11 Fix 39 (2026-05-04, A1.3): per-component reward-magnitude-ratio
/// canary producer. Single-block, 6-thread kernel reads 6 contiguous
/// EMAs at `ISV[REWARD_POPART_EMA_INDEX..+6)`, normalises to ratios,
/// and mirrors the popart magnitude into scratch[6] as a side-output.
/// SP11 Fix 39 (2026-05-04, A1.3 + B1b fix-up 2026-05-04):
/// per-component reward-magnitude-ratio canary producer. Single-block,
/// 6-thread kernel reads NON-CONTIGUOUS sources after the B1b fix-up:
/// axis 0 (popart) = ISV[POPART_COMPONENT_MAG_EMA_INDEX = 360]
/// axis 1..5 = ISV[REWARD_POPART_EMA_INDEX + 1 .. + 6)
/// (cf=64, trail=65, micro=66, opp_cost=67, bonus=68)
/// Normalises to ratios and mirrors the popart-component magnitude
/// into scratch[6] as the PnL signal scale (= ISV[359] post Pearls A+D).
/// Writes 7 floats to `scratch[SCRATCH_SP11_MAG_RATIO_BASE..+7)`.
/// Followed by 2 apply_pearls_ad_kernel launches:
/// (n_slots=6) → ISV[REWARD_COMPONENT_MAG_RATIO_BASE..+6)
@@ -4813,6 +4840,41 @@ pub struct GpuDqnTrainer {
/// (alloc_episodes × alloc_timesteps). Cached on the trainer because
/// the engagement kernel takes it as an i32 arg.
pub(crate) sp11_saboteur_delta_reward_len: usize,
/// SP11 Fix 39 B1b fix-up (2026-05-04): per-bar popart-component-
/// specific reward magnitude buffer. Mapped-pinned f32, sized
/// `MAX_BATCH_SIZE` for trainer-side use; the experience collector
/// owns the main per-bar buffer (alloc_episodes × alloc_timesteps).
/// The trainer field exists for the rare case where SP11 needs to
/// stage popart-component values inside the per-step path; the
/// canonical write site is `experience_env_step` operating on the
/// collector's buffer (kernel-to-kernel dataflow). Spec §4 amendment
/// "Why slot 360".
///
/// Per `feedback_no_htod_htoh_only_mapped_pinned.md`: mapped-pinned
/// (`MappedF32Buffer`), not `CudaSlice<f32>`. Initial value zero
/// (constructor); Pearl A's sentinel-bootstrap consumes the zero on
/// the first observation.
pub(crate) popart_component_per_sample: super::mapped_pinned::MappedF32Buffer,
/// SP11 Fix 39 B1b fix-up (2026-05-04): popart-component-specific
/// magnitude EMA producer kernel. Reads
/// `popart_component_per_sample_dev_ptr` (the collector's per-bar
/// buffer) and writes mean(|x|) to
/// `scratch[SCRATCH_SP11_POPART_COMPONENT_EMA]`. Followed by 1
/// chained `apply_pearls_ad_kernel` launch (n_slots=1) →
/// ISV[POPART_COMPONENT_MAG_EMA_INDEX=360]. Loaded from
/// `popart_component_ema_kernel.cubin`. Resolves the slot-63 PopArt-
/// input vs popart-component-magnitude overload that B1b smoke
/// surfaced. Spec §4 amendment "Why slot 360".
sp11_popart_component_ema_kernel: CudaFunction,
/// SP11 Fix 39 B1b fix-up (2026-05-04): device pointer + bar count
/// for the experience collector's `popart_component_per_sample`
/// buffer. Wired in the same commit by
/// `set_sp11_popart_component_buf` mirroring the saboteur Δreward
/// pattern at `set_sp11_saboteur_delta_reward_buf`. Initial 0 / 0
/// (kernel guarded — `launch_sp11_popart_component_ema` no-ops when
/// unwired).
pub(crate) sp11_popart_component_dev_ptr: u64,
pub(crate) sp11_popart_component_len: usize,
/// SP11 Fix 39 (2026-05-04, Task A2): reward-subsystem controller
/// kernel. Single-block, 10-thread producer reading 5 canary ISV slots
/// [350..360) and writing 10 floats to scratch[SCRATCH_SP11_CONTROLLER_BASE..+10).
@@ -12525,6 +12587,23 @@ impl GpuDqnTrainer {
self.sp11_saboteur_delta_reward_len = n_bars;
}
/// SP11 Fix 39 B1b fix-up (2026-05-04): wire the experience collector's
/// `popart_component_per_sample` device pointer into the trainer so
/// `launch_sp11_popart_component_ema` reads the right buffer. Mirror
/// of `set_sp11_saboteur_delta_reward_buf` — the collector owns the
/// per-bar buffer (alloc_episodes × alloc_timesteps); the trainer
/// caches the pointer + length for the per-step launcher path.
/// Per `feedback_wire_everything_up`: this is wired in the same commit
/// that adds the buffer + kernel.
pub(crate) fn set_sp11_popart_component_buf(
&mut self,
dev_ptr: u64,
n_bars: usize,
) {
self.sp11_popart_component_dev_ptr = dev_ptr;
self.sp11_popart_component_len = n_bars;
}
/// SP11 Fix 39 (2026-05-04, A1.1): launch the val-sharpe Δ + variance
/// canary producer.
///
@@ -12698,13 +12777,16 @@ impl GpuDqnTrainer {
Ok(())
}
/// SP11 Fix 39 (2026-05-04, A1.3): launch the per-component reward-
/// magnitude-ratio canary producer.
/// SP11 Fix 39 (2026-05-04, A1.3 + B1b fix-up 2026-05-04): launch the
/// per-component reward-magnitude-ratio canary producer.
///
/// Reads 6 contiguous EMAs at `ISV[REWARD_POPART_EMA_INDEX..+6)`,
/// normalises to ratios in `scratch[SCRATCH_SP11_MAG_RATIO_BASE..+6)`,
/// and mirrors the popart magnitude into `scratch[base+6]`. Followed
/// by two chained `apply_pearls_ad_kernel` launches:
/// B1b fix-up: reads NON-CONTIGUOUS sources after the slot-360 split.
/// axis 0 (popart) ← ISV[POPART_COMPONENT_MAG_EMA_INDEX = 360]
/// axis 1..5 ← ISV[REWARD_POPART_EMA_INDEX + 1 .. + 6)
/// (cf=64, trail=65, micro=66, opp_cost=67, bonus=68)
/// Normalises to ratios in `scratch[SCRATCH_SP11_MAG_RATIO_BASE..+6)`,
/// and mirrors the popart-component magnitude into `scratch[base+6]`.
/// Followed by two chained `apply_pearls_ad_kernel` launches:
/// 1. n_slots=6 → ISV[REWARD_COMPONENT_MAG_RATIO_BASE..+6)
/// 2. n_slots=1 → ISV[PNL_REWARD_MAGNITUDE_EMA_INDEX=359]
/// (The ISV slots are non-contiguous: 352..358 plus 359, so two
@@ -12715,7 +12797,7 @@ impl GpuDqnTrainer {
use crate::cuda_pipeline::sp5_isv_slots::SP5_SLOT_BASE;
use crate::cuda_pipeline::sp11_isv_slots::{
REWARD_COMPONENT_MAG_RATIO_BASE, REWARD_COMPONENT_MAG_RATIO_COUNT,
PNL_REWARD_MAGNITUDE_EMA_INDEX,
PNL_REWARD_MAGNITUDE_EMA_INDEX, POPART_COMPONENT_MAG_EMA_INDEX,
};
debug_assert!(self.isv_signals_dev_ptr != 0,
@@ -12729,14 +12811,19 @@ impl GpuDqnTrainer {
* std::mem::size_of::<f32>() as u64;
let scratch_out_dev: u64 = scratch_dev + scratch_offset_bytes;
let popart_ema_base_slot_i32: i32 = REWARD_POPART_EMA_INDEX as i32;
// B1b fix-up: split popart axis (slot 360) from cf/trail/micro/
// opp_cost/bonus axes (slots 64..68). Pre-fix the kernel read 6
// 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;
unsafe {
self.stream
.launch_builder(&self.sp11_mag_ratio_compute_kernel)
.arg(&isv_dev)
.arg(&scratch_out_dev)
.arg(&popart_ema_base_slot_i32)
.arg(&popart_specific_slot_i32)
.arg(&cf_others_base_slot_i32)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (6, 1, 1),
@@ -12785,6 +12872,93 @@ impl GpuDqnTrainer {
Ok(())
}
/// SP11 Fix 39 B1b fix-up (2026-05-04): launch the popart-component-
/// specific magnitude EMA producer + chained Pearls A+D singleton
/// smoothing.
///
/// Reads the experience collector's per-bar `popart_component_per_sample`
/// buffer (wired in via `set_sp11_popart_component_buf`), computes
/// mean(|x|) into `scratch[SCRATCH_SP11_POPART_COMPONENT_EMA]` via a
/// single-block tree-reduce (BLOCK_SIZE=256), then chains a Pearls
/// A+D singleton launch into ISV[POPART_COMPONENT_MAG_EMA_INDEX=360].
///
/// Runs BEFORE `launch_sp11_mag_ratio_compute` in the per-step path
/// so the mag-ratio canary's first read of slot 360 sees a populated
/// value (or sentinel 0 on the very first epoch — Pearl A bootstraps
/// from the first observation).
///
/// No-op when the popart-component buffer hasn't been wired yet
/// (early init before the collector exists). Mirrors the saboteur
/// engagement launcher's NULL-guard pattern.
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;
debug_assert!(self.isv_signals_dev_ptr != 0,
"launch_sp11_popart_component_ema: isv_signals_dev_ptr must be allocated");
// Precondition guard — the popart-component buffer is wired by
// `set_sp11_popart_component_buf` after collector construction.
// Until then, no-op (NOT a stub return value per
// `feedback_no_stubs.md` — a precondition guard that defers the
// launch). Once wired, every subsequent step launches.
if self.sp11_popart_component_dev_ptr == 0
|| self.sp11_popart_component_len == 0
{
return Ok(());
}
let isv_dev = self.isv_signals_dev_ptr;
let scratch_dev = self.producer_step_scratch_buf.dev_ptr;
let wiener_dev = self.wiener_state_buf.dev_ptr;
let popart_dev = self.sp11_popart_component_dev_ptr;
let scratch_offset_bytes = (SCRATCH_SP11_POPART_COMPONENT_EMA as u64)
* std::mem::size_of::<f32>() as u64;
let scratch_out_dev: u64 = scratch_dev + scratch_offset_bytes;
let scratch_idx_i32: i32 = SCRATCH_SP11_POPART_COMPONENT_EMA as i32;
let n_bars_i32: i32 = i32::try_from(self.sp11_popart_component_len)
.unwrap_or(i32::MAX);
// 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).
unsafe {
self.stream
.launch_builder(&self.sp11_popart_component_ema_kernel)
.arg(&popart_dev)
.arg(&scratch_out_dev)
.arg(&n_bars_i32)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (256, 1, 1),
shared_mem_bytes: 256 * std::mem::size_of::<f32>() as u32,
})
.map_err(|e| MLError::ModelError(format!("sp11 popart_component_ema: {e}")))?;
}
// Step 2: chained apply_pearls_ad_kernel singleton →
// ISV[POPART_COMPONENT_MAG_EMA_INDEX=360].
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;
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,
1,
ALPHA_META,
)?;
}
Ok(())
}
/// SP11 Fix 39 (2026-05-04, Task A2): launch the reward-subsystem
/// controller producer + chained Pearls A+D output smoothing.
///
@@ -16215,6 +16389,35 @@ impl GpuDqnTrainer {
module.load_function("reward_component_mag_ratio_compute_kernel")
.map_err(|e| MLError::ModelError(format!("sp11 mag_ratio_compute load: {e}")))?
};
// SP11 Fix 39 B1b fix-up (2026-05-04): popart-component-specific
// magnitude EMA producer kernel. Loaded alongside the other A1
// canary producers — Pearls A+D output target is ISV[POPART_
// COMPONENT_MAG_EMA_INDEX=360]. Resolves the slot-63 PopArt-input
// vs popart-component-magnitude overload that B1b smoke surfaced.
// Spec §4 amendment "Why slot 360".
let sp11_popart_component_ema_kernel = {
let module = stream.context()
.load_cubin(SP11_POPART_COMPONENT_EMA_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!("sp11 popart_component_ema cubin load: {e}")))?;
module.load_function("popart_component_ema_kernel")
.map_err(|e| MLError::ModelError(format!("sp11 popart_component_ema load: {e}")))?
};
// SP11 Fix 39 B1b fix-up (2026-05-04): allocate trainer-side
// popart-component-per-sample mapped-pinned placeholder buffer.
// The canonical per-bar buffer is owned by the experience
// collector (alloc_episodes × alloc_timesteps); the trainer
// launcher reads via raw device pointer wired in by
// `set_sp11_popart_component_buf` after collector construction.
// The trainer-side field exists to satisfy the buffer-lifetime
// contract on the trainer struct — it's a 1-element placeholder
// that's never directly read. Initial value zero (constructor);
// mapped-pinned per `feedback_no_htod_htoh_only_mapped_pinned`
// (the rule applies even to placeholder allocations).
let popart_component_per_sample = unsafe {
super::mapped_pinned::MappedF32Buffer::new(1)
}.map_err(|e| MLError::ModelError(
format!("SP11 popart_component_per_sample alloc: {e}")
))?;
// SP11 Fix 39 (2026-05-04, A1.1): allocate the val-sharpe history
// mapped-pinned buffer. 2-element [prev_epoch, curr_epoch] — the
@@ -19615,6 +19818,16 @@ impl GpuDqnTrainer {
// is constructed (see training_loop.rs init path).
sp11_saboteur_delta_reward_dev_ptr: 0,
sp11_saboteur_delta_reward_len: 0,
// SP11 Fix 39 B1b fix-up (2026-05-04): popart-component-
// specific magnitude EMA infrastructure. Kernel + trainer-side
// mapped-pinned buffer, plus dev_ptr/len cached for the
// canonical experience-collector buffer (wired in via
// `set_sp11_popart_component_buf` after collector construction
// — same pattern as the saboteur Δreward path above).
popart_component_per_sample,
sp11_popart_component_ema_kernel,
sp11_popart_component_dev_ptr: 0,
sp11_popart_component_len: 0,
// SP11 Fix 39 (Task A2): controller + SimHash novelty kernels
// and their backing buffers. Projection matrix was populated
// on-device above by `launch_novelty_simhash_proj_init`; hash

43
crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
View File

@@ -620,6 +620,14 @@ pub struct GpuExperienceCollector {
/// kernel reads via raw device pointer. See spec §3.3.1 + the param
/// docstring on `experience_env_step::saboteur_delta_reward_per_sample`.
saboteur_delta_reward_per_sample: CudaSlice<f32>, // [alloc_episodes * alloc_timesteps]
/// SP11 Fix 39 B1b fix-up (2026-05-04): per-bar popart-component-
/// specific reward magnitude. [alloc_episodes * alloc_timesteps] —
/// written by `experience_env_step` at the `r_popart` voluntary-exit
/// branch. Resolves the slot-63 PopArt-input vs popart-component-
/// magnitude overload that B1b smoke surfaced. Spec §4 amendment
/// "Why slot 360". Pure GPU buffer (kernel-to-kernel only); the SP11
/// popart-component EMA kernel reads via raw device pointer.
popart_component_per_sample: CudaSlice<f32>, // [alloc_episodes * alloc_timesteps]
/// C.2 Reward-component attribution (Plan 3 Task 1, spec §4.C.2).
/// [alloc_episodes * alloc_timesteps * 6] row-major.
/// Components: [0]=popart(total), [1]=cf, [2]=trail(0), [3]=micro, [4]=opp_cost, [5]=bonus.
@@ -1208,6 +1216,15 @@ impl GpuExperienceCollector {
let saboteur_delta_reward_per_sample = stream
.alloc_zeros::<f32>(total_output)
.map_err(|e| MLError::ModelError(format!("alloc saboteur_delta_reward_per_sample: {e}")))?;
// SP11 Fix 39 B1b fix-up (2026-05-04): per-bar popart-component-
// specific reward magnitude buffer. Same shape — one f32 per (i,t).
// Written by `experience_env_step` at the `r_popart = capped_pnl;`
// voluntary-exit branch (line ~2680); 0 elsewhere (default at
// entry-point). Read by the SP11 popart-component EMA kernel via
// raw device pointer. Spec §4 amendment "Why slot 360".
let popart_component_per_sample = stream
.alloc_zeros::<f32>(total_output)
.map_err(|e| MLError::ModelError(format!("alloc popart_component_per_sample: {e}")))?;
// C.2 Reward-component attribution (Plan 3 Task 1, spec §4.C.2).
// [total_output * 6] row-major: [popart, cf, trail, micro, opp_cost, bonus] per (i,t).
// Producer: experience_env_step (writes rc[0], rc[1], rc[3], rc[4], rc[5]; rc[2]=0 placeholder).
@@ -1707,6 +1724,7 @@ impl GpuExperienceCollector {
micro_reward_per_sample,
total_reward_per_sample,
saboteur_delta_reward_per_sample,
popart_component_per_sample,
reward_components_per_sample,
flat_to_pos_per_sample,
readiness_per_sample,
@@ -2136,6 +2154,22 @@ impl GpuExperienceCollector {
(dev_ptr, n_bars)
}
/// SP11 Fix 39 B1b fix-up (2026-05-04): expose the popart-component-
/// specific magnitude per-bar buffer to the trainer. Mirror of
/// `sp11_saboteur_delta_reward_dev_ptr` — same wire-up pattern. The
/// trainer caches this via
/// `GpuDqnTrainer::set_sp11_popart_component_buf` after the collector
/// is constructed (training_loop.rs init path) so
/// `launch_sp11_popart_component_ema` has the consumer-side pointer
/// ready every step. Per `feedback_wire_everything_up`: this getter +
/// the trainer setter + the env_step kernel write site land together
/// in this commit.
pub fn sp11_popart_component_dev_ptr(&self) -> (u64, usize) {
let dev_ptr = self.popart_component_per_sample.device_ptr(&self.stream).0;
let n_bars = self.alloc_episodes * self.alloc_timesteps;
(dev_ptr, n_bars)
}
/// Upload pre-computed expert demonstration actions to GPU.
///
/// `expert_actions[bar_index]` = expert exposure action index (-1 = no opinion).
@@ -4029,6 +4063,15 @@ impl GpuExperienceCollector {
// raw device pointer (see `set_sp11_saboteur_delta_reward_buf`
// wire-up in training_loop.rs init).
.arg(&mut self.saboteur_delta_reward_per_sample)
// SP11 Fix 39 B1b fix-up (2026-05-04): per-bar popart-
// component-specific reward magnitude output. Kernel
// writes `r_popart` (vol-normalized capped P&L on
// voluntary-exit bars; 0 elsewhere). Consumed by the
// SP11 popart-component EMA kernel via raw device
// pointer (see `set_sp11_popart_component_buf` wire-up
// in training_loop.rs init). Spec §4 amendment "Why
// slot 360".
.arg(&mut self.popart_component_per_sample)
.launch(launch_cfg)
.map_err(|e| MLError::ModelError(format!(
"experience_env_step t={t}: {e}"

78
crates/ml/src/cuda_pipeline/popart_component_ema_kernel.cu Normal file
View File

@@ -0,0 +1,78 @@
// crates/ml/src/cuda_pipeline/popart_component_ema_kernel.cu
//
// SP11 B1b fix-up (2026-05-04): popart-component-specific magnitude EMA.
//
// 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`.
//
// Why a separate kernel from `reward_component_ema_kernel`:
// The pre-SP11 SP4 reward_component_ema kernel writes
// `|reward_components_per_sample[+0]|` EMA to ISV[REWARD_POPART_EMA_INDEX=63].
// That slot is overloaded — it is BOTH PopArt's normalisation-target
// input (= total composed reward magnitude EMA) AND an attempted stand-
// in for popart-component-specific magnitude. Pre-SP11 the two values
// collapsed because reward composition was inline accumulation. B1b's
// structural decomposition (per-component locals + Σ w_i × r_i)
// surfaces the contamination — the SP11 mag-ratio canary reading slot
// 63 saw the inflated TOTAL reward magnitude as "popart's signal" and
// redistributed the controller's weight toward popart (w_pop ≈ 2.0)
// based on a contaminated ratio → 10× sharpe drop in B1b smoke. Spec
// §4 amendment "Why slot 360" documents the resolution: a separate ISV
// slot for popart-component-specific magnitude. Slot 63 is unchanged
// (PopArt's input = total reward magnitude, pre-SP11 invariant
// 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
// `feedback_no_cpu_compute_strict.md`.
#include <cuda_runtime.h>
#define BLOCK_SIZE 256
extern "C" __global__ void popart_component_ema_kernel(
/* Per-bar popart-component-specific reward magnitude. Producer
* `experience_env_step` writes the absolute popart-component value
* (matching the `r_popart = capped_pnl;` voluntary-exit branch,
* `0.0` everywhere else). Layout [n_samples] f32, mapped-pinned. */
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]. */
float* __restrict__ scratch_out,
const int n_samples)
{
/* Single-block reducer. */
if (blockIdx.x != 0) return;
extern __shared__ float sdata[];
/* 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;
for (int i = threadIdx.x; i < n_samples; i += blockDim.x) {
acc += fabsf(popart_component_per_sample[i]);
}
sdata[threadIdx.x] = acc;
__syncthreads();
/* Standard log2(BLOCK_SIZE) tree reduction — no atomicAdd. */
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) {
/* 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);
__threadfence_system();
}
}

95
crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu
View File

@@ -1,37 +1,47 @@
// crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu
//
// SP11 (Fix 39, 2026-05-04, A1.3): per-component reward-magnitude-ratio
// canary producer.
// 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 =
// `POPART_COMPONENT_MAG_EMA_INDEX` produced by `popart_component_ema_kernel`.
// This kernel reads non-contiguous sources after the fix:
//
// Reads 6 existing per-component reward-magnitude EMAs from ISV at the
// contiguous block `[popart_ema_base_slot..+6)` (= REWARD_POPART_EMA_INDEX
// = 63, populated by the SP4 reward-component EMA kernel). Components in
// fixed order (per `gpu_dqn_trainer.rs::REWARD_COMPONENT_COUNT`):
// 0 popart 1 cf 2 trail 3 micro 4 opp_cost 5 bonus
// axis 0 (popart) = isv[popart_specific_slot] ← slot 360 (B1b fix-up)
// axis 1 (cf) = isv[cf_others_base_slot + 0] ← slot 64
// axis 2 (trail) = isv[cf_others_base_slot + 1] ← slot 65
// axis 3 (micro) = isv[cf_others_base_slot + 2] ← slot 66
// 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.
//
// Writes 7 floats to `scratch_out`:
// scratch_out[0..6) = ratios (each / sum), Σ=1 modulo EPS_DIV floor
// scratch_out[6] = isv[popart_ema_base_slot] — popart magnitude
// mirrored to PNL_REWARD_MAGNITUDE_EMA_INDEX = 359
// via the chained Pearls-A+D launch (n_slots=7).
// 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).
//
// The mirror at scratch[6] avoids allocating a separate copy kernel:
// the popart-component EMA IS the PnL signal scale used to bound
// curiosity (spec §3.4.2) and to threshold saboteur engagement
// (spec §3.3.1). Promoting it to its own ISV slot keeps the controller's
// signal-relative bounds first-class.
// 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.
//
// Source EMAs are already populated by SP4's `reward_component_ema_kernel`
// — this kernel only reads, normalises, and forwards (no compute beyond
// sum + divide). Single block, 6 threads (one per component); thread 0
// also computes the sum and writes the popart mirror.
// 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`.
// The kernel emits all 7 scratch values in one shot.
//
// Pearls A+D applied chain-style by `apply_pearls_ad_kernel` (n_slots=7)
// targeting ISV[REWARD_COMPONENT_MAG_RATIO_BASE..+6) ∪ {PNL_REWARD_
// MAGNITUDE_EMA_INDEX}. The ISV slots are NOT contiguous (six at 352..358
// then a gap to 359), so the launcher chains two `launch_apply_pearls`
// calls — a 6-slot block then the singleton — rather than one n_slots=7
// call. The kernel signature is unchanged.
// Single block, 6 threads (one per component); thread 0 also computes
// the sum and writes the popart mirror.
#include <cuda_runtime.h>
@@ -39,23 +49,32 @@ extern "C" __global__ void reward_component_mag_ratio_compute_kernel(
/* ISV signal bus (read-only). */
const float* __restrict__ isv,
/* Producer scratch buffer. Caller indexes the base into this array
* via the `REWARD_COMPONENT_MAG_RATIO_BASE` constant. The kernel
* writes the 7 values relative to the base provided by the launcher
* (the launcher passes a sub-slice via pointer arithmetic at the
* `arg(&scratch_dev_with_offset)` site). */
* via the `SCRATCH_SP11_MAG_RATIO_BASE` constant via pointer
* arithmetic in the launcher. */
float* __restrict__ scratch_out,
/* ISV slot index of the first per-component magnitude EMA.
* = REWARD_POPART_EMA_INDEX = 63. */
int popart_ema_base_slot
/* ISV slot index of popart-component-specific magnitude EMA
* = POPART_COMPONENT_MAG_EMA_INDEX = 360 (B1b fix-up, 2026-05-04). */
int popart_specific_slot,
/* 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
) {
if (blockIdx.x != 0) return;
if (threadIdx.x >= 6) return;
/* Block-load all 6 component magnitudes into shared mem. */
/* 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;
values[threadIdx.x] = isv[popart_ema_base_slot + (int)threadIdx.x];
if (threadIdx.x == 0) {
values[0] = isv[popart_specific_slot];
} else {
/* Components 1..5 → ISV[cf_others_base_slot + (c - 1)] */
values[threadIdx.x] = isv[cf_others_base_slot + (int)threadIdx.x - 1];
}
__syncthreads();
if (threadIdx.x == 0) {
@@ -70,9 +89,11 @@ extern "C" __global__ void reward_component_mag_ratio_compute_kernel(
* 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 magnitude → ISV[359] via chained
* Pearls A+D singleton launch. Read directly from values[0]
* which thread 0 just loaded above. */
/* 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];
}
__syncthreads();

26
crates/ml/src/cuda_pipeline/sp11_isv_slots.rs
View File

@@ -1,12 +1,13 @@
//! SP11 — Reward-subsystem controller ISV slot constants.
//!
//! Slots [340..360) populated by the SP11 controller chain:
//! Slots [340..361) populated by the SP11 controller chain:
//! - 6 component weights [340..346)
//! - 4 controller scalar outputs [346..350)
//! - 2 val-sharpe canaries [350..352)
//! - 6 mag-ratio canaries [352..358)
//! - 1 saboteur engagement (358)
//! - 1 PnL magnitude EMA (359)
//! - 1 PnL magnitude EMA (359) — total reward magnitude (PopArt's input)
//! - 1 PopArt-component-specific magnitude EMA (360) — B1b fix-up
//!
//! Spec: docs/superpowers/specs/2026-05-04-sp11-reward-as-controlled-subsystem.md
//! Pearl: pearl_reward_as_controlled_subsystem.md
@@ -44,6 +45,25 @@ pub const REWARD_COMPONENT_MAG_RATIO_COUNT: usize = 6;
pub const SABOTEUR_ENGAGEMENT_RATE_INDEX: usize = 358;
pub const PNL_REWARD_MAGNITUDE_EMA_INDEX: usize = 359;
// Canary — popart-component-specific magnitude EMA (B1b fix-up, 2026-05-04).
//
// Pre-SP11 the SP4 reward_component_ema kernel wrote `|reward_components_per_sample[+0]|`
// EMA into ISV[REWARD_POPART_EMA_INDEX=63]. That slot was overloaded:
// it served BOTH as PopArt's normalisation-target input AND as a stand-in
// for popart-component-specific magnitude. Pre-SP11 reward composition was
// inline accumulation (popart was the dominant component, so the two values
// collapsed). B1b's structural decomposition (per-component locals + Σ w_i × r_i)
// surfaces the contamination: the SP11 mag-ratio canary reading slot 63
// sees the inflated TOTAL-reward magnitude as "popart's signal" and over-
// weighs popart (w_pop ≈ 2.0) → 10× sharpe drop in B1b smoke.
//
// Resolution: a separate ISV slot for popart-component-specific magnitude.
// Producer is `popart_component_ema_kernel.cu` chained through
// `apply_pearls_ad_kernel` (n_slots=1). Slot 63 stays as PopArt's input
// (= total reward magnitude), preserving the pre-SP11 invariant. Spec
// §4 amendment "Why slot 360".
pub const POPART_COMPONENT_MAG_EMA_INDEX: usize = 360;
// Invariant-1 anchors (rate-limiters, NOT regime thresholds).
//
// Per `feedback_isv_for_adaptive_bounds`: every adaptive bound lives on
@@ -62,7 +82,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 = 360;
pub const SP11_SLOT_END: usize = 361;
pub const SP11_PRODUCER_COUNT: usize = SP11_SLOT_END - SP11_SLOT_BASE;
/// Salt added to `config.seed` when seeding the novelty SimHash projection

73
crates/ml/src/cuda_pipeline/sp5_isv_slots.rs
View File

@@ -18,9 +18,10 @@
//! 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..360 SP11 (Fix 39) reward-subsystem controller (20 slots, fold-reset) — see sp11_isv_slots
//! 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]
//!
//! Total: 184 new SP5 ISV slots (52 + 24 + 20 + 4 + 4 + 6 + 4 + 4 + 3 + 16 + 8 + 9 + 9 + 1 + 20).
//! Total: 185 new SP5 ISV slots (52 + 24 + 20 + 4 + 4 + 6 + 4 + 4 + 3 + 16 + 8 + 9 + 9 + 1 + 21).
// Re-export SP11 slot constants so existing consumers can import every
// slot constant from a single location (`crate::cuda_pipeline::sp5_isv_slots::*`).
@@ -302,7 +303,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 = 360;
pub const SP5_SLOT_END: usize = 361;
/// Wiener-buffer producer-count constant. Sizes `wiener_state_buf` via
/// `(SP4_PRODUCER_COUNT + SP5_PRODUCER_COUNT) * SP4_WIENER_FLOATS_PER_SLOT`.
@@ -381,19 +382,23 @@ pub const SP5_SLOT_END: usize = 360;
/// two outputs); FoldReset (sentinel 0; Pearl A bootstrap on first
/// observation).
///
/// SP11 (2026-05-04, Fix 39): allocates 20 new SP11 ISV slots at
/// ISV[340..360) — 6 reward component weights + 4 controller scalar
/// outputs + 10 canaries (val-sharpe Δ/var, 6 mag-ratios, saboteur
/// engagement, PnL EMA). Unique-slot count grows 164 → 184; the linear
/// span (and `SP5_PRODUCER_COUNT`) grows 166 → 186. Drives the reward
/// subsystem controller per `pearl_reward_as_controlled_subsystem.md` +
/// `pearl_controller_anchors_isv_driven.md`. All 20 slots are FoldReset
/// SP11 (2026-05-04, Fix 39): allocates 21 new SP11 ISV slots at
/// ISV[340..361) — 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 = 186;
// linear span = SP5_SLOT_END - SP5_SLOT_BASE = 360 - 174 = 186 wiener triples
// unique-slot count = 184 (52 per-branch + 24 Adam + 20 IQN τ + 4 trail
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
// + 4 num_atoms + 6 Kelly + 4 Layer D D1 PnL aggregation
// + 4 Layer D D2 health composition
// + 3 Layer D D3 training metrics EMA
@@ -403,7 +408,7 @@ pub const SP5_PRODUCER_COUNT: usize = 186;
// + 8 SP8 LB_MAX_BUDGET per-(head,branch)
// + 9 SP9 Kelly cold-start warmup floor + targets + eval_dist
// + 1 SP10 eval Thompson temperature
// + 20 SP11 reward-subsystem controller)
// + 21 SP11 reward-subsystem controller [includes B1b fix-up slot 360])
// ── Convenience accessors ────────────────────────────────────────────
#[inline] pub const fn atom_v_center(b: usize) -> usize { ATOM_V_CENTER_BASE + b }
@@ -466,7 +471,8 @@ pub const SP5_LAYOUT_FINGERPRINT_FRAGMENT: &str =
VAL_SHARPE_DELTA_EMA=350;VAL_SHARPE_VAR_EMA=351;\
REWARD_COMPONENT_MAG_RATIO_BASE=352;\
SABOTEUR_ENGAGEMENT_RATE=358;PNL_REWARD_MAGNITUDE_EMA=359;\
ISV_TOTAL_DIM=360";
POPART_COMPONENT_MAG_EMA=360;\
ISV_TOTAL_DIM=361";
#[cfg(test)]
mod tests {
@@ -574,32 +580,33 @@ mod tests {
// SP10 (Fix 38): eval Thompson selector temperature
slots.insert(EVAL_THOMPSON_TEMP_INDEX);
// SP11 (Fix 39): include the 20 new reward-subsystem controller slots.
// SP11 (Fix 39): include the 21 new reward-subsystem controller slots.
for s in SP11_SLOT_BASE..SP11_SLOT_END {
slots.insert(s);
}
// 1. Exactly 184 unique slots (164 pre-SP11 + 20 SP11).
assert_eq!(slots.len(), 184, "expected 184 unique slots, got {}", slots.len());
// 1. Exactly 185 unique slots (164 pre-SP11 + 21 SP11).
assert_eq!(slots.len(), 185, "expected 185 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 = 359.
assert_eq!(*slots.iter().max().unwrap(), 359);
// 3. Max slot is SP5_SLOT_END - 1 = 360.
assert_eq!(*slots.iter().max().unwrap(), 360);
// 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..360} — no holes (other than the
// 5. Set equals {174..278} ∪ {280..361} — 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.
let expected: HashSet<usize> = (174..278).chain(280..360).collect();
// 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();
assert_eq!(slots, expected,
"slot set does not match expected {{174..278}} ∪ {{280..360}}");
}
@@ -628,11 +635,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, 360);
assert_eq!(SP5_SLOT_END, 361);
// 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, 186);
assert_eq!(SP5_PRODUCER_COUNT, 187);
assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE);
}
@@ -651,7 +658,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, 360);
assert_eq!(SP5_SLOT_END, 361);
// SP5_PRODUCER_COUNT linear-span check matches the new end.
assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE);
}
@@ -672,12 +679,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, 360);
assert_eq!(SP5_SLOT_END, 361);
// 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, 186);
assert_eq!(SP5_PRODUCER_COUNT, 187);
assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE);
}
@@ -697,7 +704,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, 360);
assert_eq!(SP5_SLOT_END, 361);
assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE);
}
@@ -717,7 +724,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, 360);
assert_eq!(SP5_SLOT_END, 361);
assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE);
}
@@ -750,7 +757,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, 360);
assert_eq!(SP5_SLOT_END, 361);
// Linear span matches producer count (no gaps in SP9 block).
assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE);
}
@@ -764,9 +771,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, 360);
assert_eq!(SP5_SLOT_END, 361);
// 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, 186);
assert_eq!(SP5_PRODUCER_COUNT, 187);
}
}

5
crates/ml/src/trainers/dqn/state_reset_registry.rs
View File

@@ -895,6 +895,11 @@ impl StateResetRegistry {
category: ResetCategory::FoldReset,
description: "ISV[PNL_REWARD_MAGNITUDE_EMA_INDEX=359] — SP11 Fix 39 PnL-reward magnitude EMA (running `|reward_total|` smoothed with Wiener-α). Produced by the unified-reward kernel chain (existing producer; SP11 wires the EMA output to this slot in A1 alongside the val-sharpe and engagement canaries). Consumed by the SP11 controller as the curiosity-bound scale factor (`curiosity_bound = SP11_CURIOSITY_BOUND_FRACTION=0.3 × this slot`) per `feedback_isv_for_adaptive_bounds.md` §3.4.2. FoldReset sentinel 0; Pearl A bootstraps. Spec §3.3.1, §3.4.2.",
},
RegistryEntry {
name: "sp11_popart_component_mag_ema",
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'.",
},
// 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

81
crates/ml/src/trainers/dqn/trainer/training_loop.rs
View File

@@ -819,6 +819,24 @@ impl DQNTrainer {
}
}
// SP11 Fix 39 B1b fix-up (2026-05-04): wire the per-bar popart-
// component-specific magnitude buffer from the experience
// collector into the trainer. Mirror of the saboteur Δreward
// wire-up above. The collector's `experience_env_step` writes
// `r_popart` to this buffer at the voluntary-exit branch; the
// trainer's `launch_sp11_popart_component_ema` reads via raw
// device pointer and reduces to mean(|x|) → Pearls A+D →
// ISV[POPART_COMPONENT_MAG_EMA_INDEX=360]. Resolves the slot-
// 63 PopArt-input vs popart-component-magnitude overload that
// B1b smoke surfaced. Spec §4 amendment "Why slot 360".
if let Some(ref collector) = self.gpu_experience_collector {
let (popart_dev_ptr, n_bars) = collector.sp11_popart_component_dev_ptr();
if let Some(ref mut fused_mut) = self.fused_ctx {
fused_mut.trainer_mut()
.set_sp11_popart_component_buf(popart_dev_ptr, n_bars);
}
}
// IQR exploration bonus from IQN quantile spread
if let Some(ref fused) = self.fused_ctx {
let iqr_ptr = fused.iqn_iqr_ptr().unwrap_or(0);
@@ -3405,18 +3423,44 @@ impl DQNTrainer {
"SP9 kelly_warmup_floor producer chain launch failed");
}
// SP11 Fix 39 (2026-05-04, A1.3): per-component
// reward-magnitude-ratio canary producer.
// Reads ISV[REWARD_POPART_EMA_INDEX..+6) (SP4
// reward-component EMAs populated by the
// experience collector's reward_component_ema
// call earlier in the epoch) and writes
// 6 normalised ratios + 1 popart mirror.
// SP11 Fix 39 B1b fix-up (2026-05-04): popart-
// component-specific magnitude EMA producer.
// Reads the per-bar `popart_component_per_sample`
// buffer (wired in via
// `set_sp11_popart_component_buf` above) and
// block-tree-reduces mean(|x|) into ISV[POPART_
// COMPONENT_MAG_EMA_INDEX=360] via Pearls A+D.
// MUST run BEFORE `launch_sp11_mag_ratio_compute`
// below so the canary's first read of slot 360
// sees a populated value (or sentinel 0 on the
// very first step — Pearl A bootstrap fires
// first-observation replacement on the second).
// Resolves the slot-63 PopArt-input vs popart-
// component-magnitude overload that B1b smoke
// surfaced. Spec §4 amendment "Why slot 360".
if let Err(e) = fused.trainer()
.launch_sp11_popart_component_ema()
{
tracing::warn!(error = %e,
"SP11 popart_component_ema launch failed");
}
// SP11 Fix 39 (2026-05-04, A1.3 + B1b fix-up):
// per-component reward-magnitude-ratio canary
// producer. Post-fix-up reads NON-CONTIGUOUS
// sources:
// axis 0 (popart) ← ISV[POPART_COMPONENT_MAG_EMA_INDEX=360]
// axis 1..5 ← ISV[REWARD_POPART_EMA_INDEX+1..+6)
// (cf=64, trail=65, micro=66,
// opp_cost=67, bonus=68)
// Pearls A+D smooth into
// ISV[REWARD_COMPONENT_MAG_RATIO_BASE..+6) and
// ISV[PNL_REWARD_MAGNITUDE_EMA_INDEX=359].
// Per-epoch cadence — matches the SP4 reward-
// component EMA producer firing pattern.
// ISV[PNL_REWARD_MAGNITUDE_EMA_INDEX=359]
// (PnL signal scale = popart-component
// magnitude post-fix-up, not contaminated
// total-reward magnitude). Per-epoch cadence —
// matches the SP4 reward-component EMA producer
// firing pattern.
if let Err(e) = fused.trainer()
.launch_sp11_mag_ratio_compute()
{
@@ -7277,6 +7321,23 @@ impl DQNTrainer {
fused.trainer().write_isv_signal_at(PNL_REWARD_MAGNITUDE_EMA_INDEX, 0.0);
}
}
// SP11 Fix 39 B1b fix-up (2026-05-04): popart-component-specific
// magnitude EMA. Pearl A sentinel-bootstrap fires on the first
// popart_component_ema_kernel launch of the new fold; the per-
// bar buffer accumulates fresh popart magnitudes from the new
// fold's voluntary-exit bars, mean(|x|) feeds Pearls A+D, slot
// 360 transitions from sentinel 0 → first-observation value.
// Companion Wiener-state slot for the offset
// (SP4_PRODUCER_COUNT*3 + (360-SP5_SLOT_BASE)*3) is reset by
// the existing bulk memset of `wiener_state_buf` covered by
// the `sp5_wiener_state` registry entry's dispatch arm. Spec
// §4 amendment "Why slot 360".
"sp11_popart_component_mag_ema" => {
if let Some(ref fused) = self.fused_ctx {
use crate::cuda_pipeline::sp5_isv_slots::POPART_COMPONENT_MAG_EMA_INDEX;
fused.trainer().write_isv_signal_at(POPART_COMPONENT_MAG_EMA_INDEX, 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

65
crates/ml/tests/sp11_producer_unit_tests.rs
View File

@@ -43,13 +43,22 @@ use ml::cuda_pipeline::sp11_isv_slots::{
/// Tests allocate a full-size mapped-pinned ISV buffer and write the
/// few slots they need; the kernels read by ISV slot index, so a
/// shorter buffer would be unsafe even with synthetic-only writes.
const ISV_TOTAL_DIM: usize = 360;
///
/// 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;
/// REWARD_POPART_EMA_INDEX from `gpu_dqn_trainer.rs` (the SP4 reward-
/// component magnitude EMA base slot, mirrored here as a const so the
/// test imports stay scoped to the SP11 producer surface).
const REWARD_POPART_EMA_INDEX: usize = 63;
/// POPART_COMPONENT_MAG_EMA_INDEX = 360 (SP11 B1b fix-up). The mag-ratio
/// kernel post-fix-up reads the popart axis from this slot instead of
/// slot 63's overloaded total-reward value. Mirrored here as a const so
/// the unit test stays scoped to the SP11 producer surface.
const POPART_COMPONENT_MAG_EMA_INDEX: usize = 360;
/// Cubin paths for each producer kernel under test. Each cubin is built
/// by `crates/ml/build.rs` from the `.cu` source and dropped into
/// `OUT_DIR` per the standard kernel-registration pattern.
@@ -108,7 +117,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 (360 f32)");
.expect("alloc isv (361 f32)");
// ISV[VAL_SHARPE_DELTA_EMA_INDEX] is constructor-zero — sentinel state
// for Pearl A's first-observation replacement. No host writes needed.
@@ -171,7 +180,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 (360 f32)");
.expect("alloc isv (361 f32)");
isv.host_slice_mut()[PNL_REWARD_MAGNITUDE_EMA_INDEX] = 1.0;
let delta_dev = delta_reward.dev_ptr;
@@ -204,13 +213,19 @@ fn saboteur_engagement_counts_above_threshold_only() {
);
}
/// SP11 A1.3 — per-component magnitude-ratio normaliser preserves Σ=1
/// and mirrors popart magnitude into scratch[6].
/// SP11 A1.3 + B1b fix-up — per-component magnitude-ratio normaliser
/// preserves Σ=1 and mirrors popart-component magnitude into scratch[6].
///
/// Inputs: ISV[REWARD_POPART_EMA_INDEX..+6) = [1.0, 2.0, 3.0, 4.0, 5.0,
/// 6.0]. Sum = 21. Expected ratios = [1/21, 2/21, 3/21, 4/21, 5/21,
/// 6/21]. Mirror at scratch[6] = 1.0 (popart magnitude unchanged).
/// Σ ratios = 1.0 ± 1e-5.
/// 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
///
/// 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.
#[test]
#[ignore = "requires GPU"]
fn reward_component_mag_ratio_normalises_and_mirrors_popart() {
@@ -221,27 +236,31 @@ 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 (360 f32)");
.expect("alloc isv (361 f32)");
{
let s = isv.host_slice_mut();
s[REWARD_POPART_EMA_INDEX + 0] = 1.0;
s[REWARD_POPART_EMA_INDEX + 1] = 2.0;
s[REWARD_POPART_EMA_INDEX + 2] = 3.0;
s[REWARD_POPART_EMA_INDEX + 3] = 4.0;
s[REWARD_POPART_EMA_INDEX + 4] = 5.0;
s[REWARD_POPART_EMA_INDEX + 5] = 6.0;
// B1b fix-up: popart axis sourced from slot 360, not 63.
s[POPART_COMPONENT_MAG_EMA_INDEX] = 1.0;
// cf, trail, micro, opp_cost, bonus at slots 64..68 (cf_others_base_slot = 64).
s[REWARD_POPART_EMA_INDEX + 1] = 2.0; // cf
s[REWARD_POPART_EMA_INDEX + 2] = 3.0; // trail
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
}
let isv_dev = isv.dev_ptr;
let scratch_dev = scratch.dev_ptr;
let popart_base_i32: i32 = REWARD_POPART_EMA_INDEX as i32;
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_base_i32)
.arg(&popart_specific_slot_i32)
.arg(&cf_others_base_slot_i32)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (6, 1, 1),
@@ -271,10 +290,10 @@ fn reward_component_mag_ratio_normalises_and_mirrors_popart() {
"expected Σratios = 1.0, got {sum_out}"
);
// Popart mirror at scratch[6].
// Popart mirror at scratch[6] — sourced from slot 360 post-B1b-fix.
assert!(
(out[6] - 1.0_f32).abs() < 1e-6,
"expected popart mirror = 1.0 (= isv[REWARD_POPART_EMA_INDEX]), got {}",
"expected popart mirror = 1.0 (= isv[POPART_COMPONENT_MAG_EMA_INDEX=360]), got {}",
out[6]
);
}
@@ -378,7 +397,7 @@ fn controller_z_score_at_zero_yields_midpoint_outputs() {
);
let mut isv = unsafe { MappedF32Buffer::new(ISV_TOTAL_DIM) }
.expect("alloc isv (360 f32)");
.expect("alloc isv (361 f32)");
populate_controller_isv(
&mut isv,
/*delta*/ 0.0,
@@ -468,7 +487,7 @@ fn controller_weights_renormalize_after_floor() {
);
let mut isv = unsafe { MappedF32Buffer::new(ISV_TOTAL_DIM) }
.expect("alloc isv (360 f32)");
.expect("alloc isv (361 f32)");
populate_controller_isv(
&mut isv,
/*delta*/ 1.0, // any non-zero
@@ -557,7 +576,7 @@ fn controller_saboteur_post_clamp_holds_min() {
);
let mut isv = unsafe { MappedF32Buffer::new(ISV_TOTAL_DIM) }
.expect("alloc isv (360 f32)");
.expect("alloc isv (361 f32)");
populate_controller_isv(
&mut isv,
/*delta*/ -100.0, // extreme regression

18
docs/isv-slots.md
View File

@@ -285,6 +285,7 @@ htod_copy/dtoh_sync_copy/alloc_zeros per
| [352..358) | `REWARD_COMPONENT_MAG_RATIO_BASE` (+6) | per-component canary | Per-component reward-magnitude ratio (winner/diversifier blend input) |
| [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) |
Invariant-1 fraction constants (rate-limiter anchors per spec §3.4.2):
`SP11_EPS_DIV=1e-6`, `SP11_WEIGHT_HARD_FLOOR=0.01`,
@@ -294,7 +295,22 @@ Invariant-1 fraction constants (rate-limiter anchors per spec §3.4.2):
`SP11_PROJ_SEED_SALT=0x_5511_0001` for the A2 novelty-SimHash projection-init
launcher's seed derivation.
All 20 slots are FoldReset (sentinel 0; Pearl A bootstraps from first
**B1b fix-up (2026-05-04, spec §4 amendment):** ISV[360] = `POPART_COMPONENT_MAG_EMA_INDEX`
added to resolve the slot-63 PopArt-input vs popart-component-magnitude
overload that B1b smoke surfaced. Pre-fix-up the SP11 mag-ratio canary
read 6 contiguous slots starting at slot 63 — but 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 contaminated
ratio → 10× sharpe drop. Post-fix-up the mag-ratio kernel takes two slot
indices: `popart_specific_slot=360` for the popart axis,
`cf_others_base_slot=64` for cf/trail/micro/opp_cost/bonus. Slot 63 is
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.
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
composition consumer for [340..346) (the 6 component weights).