From 5e16b67ca68fd55b659ca3c4f318ba191cca614d Mon Sep 17 00:00:00 2001
From: jgrusewski <jgrusewski@localhost>
Date: Mon, 4 May 2026 10:29:47 +0200
Subject: [PATCH] =?UTF-8?q?fix(sp11):=20B1b=20follow-up=20=E2=80=94=20add?=
 =?UTF-8?q?=20slot=20360=20for=20popart-component=20mag=20EMA?=
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

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.
---
 crates/ml/build.rs                            |   9 +
 .../src/cuda_pipeline/experience_kernels.cu   |  45 +++-
 .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs   | 247 ++++++++++++++++--
 .../cuda_pipeline/gpu_experience_collector.rs |  43 +++
 .../popart_component_ema_kernel.cu            |  78 ++++++
 ...ward_component_mag_ratio_compute_kernel.cu |  95 ++++---
 crates/ml/src/cuda_pipeline/sp11_isv_slots.rs |  26 +-
 crates/ml/src/cuda_pipeline/sp5_isv_slots.rs  |  73 +++---
 .../src/trainers/dqn/state_reset_registry.rs  |   5 +
 .../src/trainers/dqn/trainer/training_loop.rs |  81 +++++-
 crates/ml/tests/sp11_producer_unit_tests.rs   |  65 +++--
 docs/isv-slots.md                             |  18 +-
 12 files changed, 660 insertions(+), 125 deletions(-)
 create mode 100644 crates/ml/src/cuda_pipeline/popart_component_ema_kernel.cu

diff --git a/crates/ml/build.rs b/crates/ml/build.rs
index ba78c6819..1677f1b3b 100644
--- a/crates/ml/build.rs
+++ b/crates/ml/build.rs
@@ -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
diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu
index 92d18443e..0502a6d53 100644
--- a/crates/ml/src/cuda_pipeline/experience_kernels.cu
+++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu
@@ -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
diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
index f89efef02..ce5a8b80f 100644
--- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
@@ -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
diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
index 5e9b8f4fd..cb0240396 100644
--- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
+++ b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
@@ -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}"
diff --git a/crates/ml/src/cuda_pipeline/popart_component_ema_kernel.cu b/crates/ml/src/cuda_pipeline/popart_component_ema_kernel.cu
new file mode 100644
index 000000000..9517377c1
--- /dev/null
+++ b/crates/ml/src/cuda_pipeline/popart_component_ema_kernel.cu
@@ -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();
+    }
+}
diff --git a/crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu b/crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu
index c14168a9e..d488791dc 100644
--- a/crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu
+++ b/crates/ml/src/cuda_pipeline/reward_component_mag_ratio_compute_kernel.cu
@@ -1,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();
diff --git a/crates/ml/src/cuda_pipeline/sp11_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp11_isv_slots.rs
index 5e203f947..367dad5e2 100644
--- a/crates/ml/src/cuda_pipeline/sp11_isv_slots.rs
+++ b/crates/ml/src/cuda_pipeline/sp11_isv_slots.rs
@@ -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
diff --git a/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs
index 0c499449b..319b0797b 100644
--- a/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs
+++ b/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs
@@ -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);
     }
 }
diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs
index 037c49c5f..05f967423 100644
--- a/crates/ml/src/trainers/dqn/state_reset_registry.rs
+++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs
@@ -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
diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
index 52a82eb51..f0fe81b0c 100644
--- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs
+++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs
@@ -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
diff --git a/crates/ml/tests/sp11_producer_unit_tests.rs b/crates/ml/tests/sp11_producer_unit_tests.rs
index 22a28598a..00522956a 100644
--- a/crates/ml/tests/sp11_producer_unit_tests.rs
+++ b/crates/ml/tests/sp11_producer_unit_tests.rs
@@ -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
diff --git a/docs/isv-slots.md b/docs/isv-slots.md
index 638d0943a..e58378eb7 100644
--- a/docs/isv-slots.md
+++ b/docs/isv-slots.md
@@ -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).