feat(sp14-c6): h_s2_aux_rms_ema producer — ISV[449] per-collector-step

Single-block 256-thread CUDA kernel computing RMS(h_s2_aux [B, SH2])
and EMA-blending the step observation into ISV[H_S2_AUX_RMS_EMA_INDEX=449]
directly. Pearl-A first-observation bootstrap embedded in kernel body
(sentinel 0.0 → replace); fixed α=0.05 EMA blend thereafter.

ISV slot 449 is outside the SP4/SP5 wiener buffer linear span so the
scratch+apply_pearls_ad_kernel path is not available — self-contained
Pearl-A logic mirrors the avg_win_hold_time_update_kernel precedent
(slot 451). No atomicAdd; shmem block-tree-reduce only. Launched after
aux_trunk_forward in the collector per-step hot path.

- h_s2_aux_rms_ema_kernel.cu — new CUDA kernel (81 lines)
- build.rs — cubin manifest entry
- gpu_dqn_trainer.rs — H_S2_AUX_RMS_EMA_CUBIN static
- gpu_aux_trunk.rs — HS2AuxRmsEmaOps struct + launch()
- gpu_experience_collector.rs — field + constructor + hot-path launch
- aux_trunk_oracle_tests.rs — h_s2_aux_rms_ema_pearl_a_bootstrap test
- dqn-wire-up-audit.md — Phase C.6 audit entry

cargo check -p ml --tests: clean (only pre-existing warnings)
Oracle test: 1 new test added (requires GPU to run)

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>

crates/ml/build.rs

@@ -786,6 +786,18 @@ fn main() {
         // (no atomicAdd) per `feedback_no_atomicadd.md`. Pearl-A
         // first-observation bootstrap; α=0.05 EMA thereafter. Plan: §C.4b.
         "avg_win_hold_time_update_kernel.cu",
+        // SP14 Layer C Phase C.6 (2026-05-08): h_s2_aux RMS EMA producer.
+        // Single-block 256-thread kernel computing RMS = sqrt(mean(x²))
+        // over `h_s2_aux [B, SH2]` (aux trunk final output, no activation)
+        // and EMA-blending into `ISV[H_S2_AUX_RMS_EMA_INDEX=449]`. Pearl-A
+        // first-observation bootstrap embedded in kernel (sentinel 0.0 →
+        // replace directly); fixed α=0.05 EMA blend thereafter. ISV slot 449
+        // is outside the SP4/SP5 wiener buffer linear span so the shared
+        // scratch+apply_pearls_ad_kernel path is unavailable — bootstrap
+        // logic is self-contained per `avg_win_hold_time_update_kernel`
+        // precedent. Block-tree-reduce (no atomicAdd) per
+        // `feedback_no_atomicadd.md`. Per-collector-step launch cadence.
+        "h_s2_aux_rms_ema_kernel.cu",
         // SP14 Layer B Task B.3 (2026-05-05): Earned Gradient Flow
         // producer kernel. Per-step computes the K=4↔K=2 mapped argmax
         // mismatch between the Q-head's 4-way direction action

crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs

@@ -52,7 +52,7 @@ use crate::MLError;
 
 use super::gpu_dqn_trainer::{
     AUX_HORIZON_UPDATE_CUBIN, AUX_TRUNK_BACKWARD_CUBIN, AUX_TRUNK_FORWARD_CUBIN,
-    AVG_WIN_HOLD_TIME_UPDATE_CUBIN,
+    AVG_WIN_HOLD_TIME_UPDATE_CUBIN, H_S2_AUX_RMS_EMA_CUBIN,
 };
 
 /// Hidden width of the aux trunk's first internal layer (Linear_1 → ELU
@@ -617,3 +617,89 @@ impl AvgWinHoldTimeUpdateOps {
         Ok(())
     }
 }
+
+/// SP14 Layer C Phase C.6 (2026-05-08): h_s2_aux RMS EMA producer.
+///
+/// Computes `RMS(h_s2_aux) = sqrt(mean(h_s2_aux²))` over the
+/// `[B, SH2]` aux trunk output and EMA-blends the result into
+/// `ISV[H_S2_AUX_RMS_EMA_INDEX=449]`.
+///
+/// ISV slot 449 is in the SP14 Layer C block (outside the SP4/SP5 wiener
+/// buffer linear span), so Pearl-A bootstrap logic is embedded in the
+/// kernel body rather than delegated to `apply_pearls_ad_kernel`. Fixed
+/// α=0.05 (per-step cadence blend rate; no Wiener-optimal adaptive α
+/// available without a wiener_state_buf triple for this slot).
+///
+/// Single-block 256-thread kernel; shmem block-tree-reduce (no atomicAdd
+/// per `feedback_no_atomicadd.md`). `CudaFunction` pre-loaded at
+/// construction per `pearl_no_host_branches_in_captured_graph.md`.
+///
+/// # Pearls applied
+/// - `feedback_no_atomicadd.md` — block-tree-reduce in shmem, one global
+///   write from thread 0.
+/// - `pearl_first_observation_bootstrap.md` — sentinel 0.0 → REPLACE.
+/// - `pearl_no_host_branches_in_captured_graph.md` — `CudaFunction`
+///   pre-loaded at construction; no per-launch cubin load.
+/// - `feedback_no_stubs.md` — full body, no placeholder.
+#[allow(missing_debug_implementations)]
+pub(crate) struct HS2AuxRmsEmaOps {
+    update_kernel: CudaFunction,
+}
+
+impl HS2AuxRmsEmaOps {
+    /// Block dim — must match the kernel's `BLK_DIM` of 256.
+    const BLK_DIM: u32 = 256;
+    /// Per-step EMA blend rate. Matches the aux trunk forward launch
+    /// cadence — one observation per collector step.
+    pub(crate) const ALPHA: f32 = 0.05;
+
+    pub(crate) fn new(stream: &Arc<CudaStream>) -> Result<Self, MLError> {
+        let context = stream.context();
+        let module = context
+            .load_cubin(H_S2_AUX_RMS_EMA_CUBIN.to_vec())
+            .map_err(|e| MLError::ModelError(format!("h_s2_aux_rms_ema cubin load: {e}")))?;
+        let update_kernel = module
+            .load_function("h_s2_aux_rms_ema_update")
+            .map_err(|e| MLError::ModelError(format!("h_s2_aux_rms_ema_update load: {e}")))?;
+        Ok(Self { update_kernel })
+    }
+
+    /// Launch the RMS EMA producer.
+    ///
+    /// Args:
+    ///   - `h_s2_aux_ptr`: aux trunk output `[B, SH2]` device pointer.
+    ///   - `b`: batch size (B).
+    ///   - `sh2`: aux trunk output width (SH2 = AUX_HIDDEN_DIM = 256).
+    ///   - `isv_ptr`: ISV[ISV_TOTAL_DIM] device pointer (mapped-pinned).
+    ///   - `isv_target_idx`: H_S2_AUX_RMS_EMA_INDEX (449).
+    ///   - `alpha`: EMA blend rate (use `Self::ALPHA = 0.05`).
+    pub(crate) fn launch(
+        &self,
+        stream: &Arc<CudaStream>,
+        h_s2_aux_ptr: u64,
+        b: i32,
+        sh2: i32,
+        isv_ptr: u64,
+        isv_target_idx: i32,
+        alpha: f32,
+    ) -> Result<(), MLError> {
+        let smem_bytes = Self::BLK_DIM * std::mem::size_of::<f32>() as u32;
+        unsafe {
+            stream
+                .launch_builder(&self.update_kernel)
+                .arg(&h_s2_aux_ptr)
+                .arg(&b)
+                .arg(&sh2)
+                .arg(&isv_ptr)
+                .arg(&isv_target_idx)
+                .arg(&alpha)
+                .launch(LaunchConfig {
+                    grid_dim: (1, 1, 1),
+                    block_dim: (Self::BLK_DIM, 1, 1),
+                    shared_mem_bytes: smem_bytes,
+                })
+                .map_err(|e| MLError::ModelError(format!("h_s2_aux_rms_ema_update: {e}")))?;
+        }
+        Ok(())
+    }
+}

crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs

@@ -2099,6 +2099,17 @@ pub(crate) static AUX_HORIZON_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!(
 /// `gpu_aux_trunk::AvgWinHoldTimeUpdateOps`. Per-epoch boundary launch.
 pub(crate) static AVG_WIN_HOLD_TIME_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/avg_win_hold_time_update_kernel.cubin"));
 
+/// SP14 Layer C Phase C.6 (2026-05-08): h_s2_aux RMS EMA producer cubin.
+/// Single-block 256-thread kernel computing `sqrt(mean(h_s2_aux²))` over
+/// `h_s2_aux [B, SH2]` (aux trunk final output) and EMA-blending into
+/// `ISV[H_S2_AUX_RMS_EMA_INDEX=449]`. Pearl-A first-observation bootstrap
+/// embedded in kernel (sentinel 0.0 → replace); fixed α=0.05 EMA blend
+/// thereafter. Slot 449 is outside the SP4/SP5 wiener buffer linear span
+/// so the scratch+apply_pearls_ad_kernel path is not available. Loaded by
+/// `gpu_aux_trunk::HS2AuxRmsEmaOps`. Per-collector-step launch (same
+/// cadence as `aux_trunk_forward`).
+pub(crate) static H_S2_AUX_RMS_EMA_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/h_s2_aux_rms_ema_kernel.cubin"));
+
 /// Plan C Phase 2 follow-up A.2 (2026-04-29): q-drift rate ISV producer.
 /// Single-thread single-block cold-path kernel mirroring
 /// `moe_lambda_eff_kernel.cu` / `kelly_cap_update_kernel.cu`. Reads

crates/ml/src/cuda_pipeline/gpu_experience_collector.rs

@@ -1054,6 +1054,13 @@ pub struct GpuExperienceCollector {
     #[allow(dead_code)]
     exp_aux_trunk_forward_ops: super::gpu_aux_trunk::AuxTrunkForwardOps,
 
+    /// SP14 Layer C Phase C.6 (2026-05-08): h_s2_aux RMS EMA producer.
+    /// Launched immediately after `exp_aux_trunk_forward_ops` so that
+    /// `h_s2_aux [B, SH2]` is fully populated before the RMS reduction.
+    /// Writes `ISV[H_S2_AUX_RMS_EMA_INDEX=449]` via embedded Pearl-A
+    /// bootstrap (sentinel 0.0 → replace) + fixed α=0.05 EMA blend.
+    exp_h_s2_aux_rms_ema_ops: super::gpu_aux_trunk::HS2AuxRmsEmaOps,
+
     /// SP14 β-migration step 3 (2026-05-07): per-rollout-step label
     /// producer kernel handle. Sibling of the trajectory-wide
     /// `aux_sign_label_kernel` (line 3776 above) — same per-thread O(1)
@@ -2123,6 +2130,12 @@ impl GpuExperienceCollector {
         let exp_aux_trunk_forward_ops =
             super::gpu_aux_trunk::AuxTrunkForwardOps::new(&stream)?;
 
+        // SP14 Layer C Phase C.6 (2026-05-08): pre-load the h_s2_aux RMS EMA
+        // producer. Launched after `exp_aux_trunk_forward_ops` in the
+        // per-step hot path to update ISV[H_S2_AUX_RMS_EMA_INDEX=449].
+        let exp_h_s2_aux_rms_ema_ops =
+            super::gpu_aux_trunk::HS2AuxRmsEmaOps::new(&stream)?;
+
         // SP14 β-migration step 3 (2026-05-07): per-rollout-step label
         // producer kernel. Loaded on the collector's stream so its launch
         // chains directly after `forward_online_f32` (which populates
@@ -2589,6 +2602,7 @@ impl GpuExperienceCollector {
             aux_trunk_w3_ptr: 0,
             aux_trunk_b3_ptr: 0,
             exp_aux_trunk_forward_ops,
+            exp_h_s2_aux_rms_ema_ops,
             exp_aux_sign_label_per_step_kernel,
             // SP15 Wave 5 follow-up (2026-05-07): pre-loaded SP15 kernel handles.
             exp_sp15_dd_state_kernel,
@@ -4771,6 +4785,29 @@ impl GpuExperienceCollector {
                     sh2,
                 )?;
 
+                // Step B-1b (SP14-C.6, 2026-05-08): h_s2_aux RMS EMA producer.
+                // Reads `exp_h_s2_aux [n_episodes, SH2]` (written above by
+                // aux trunk forward) and EMA-blends RMS into
+                // `ISV[H_S2_AUX_RMS_EMA_INDEX=449]`. Pearl-A first-observation
+                // bootstrap (sentinel 0.0 → replace) + fixed α=0.05.
+                // Launched on the same stream as aux_trunk_forward so
+                // producer→consumer ordering is stream-implicit (no explicit
+                // sync required). `isv_signals_dev_ptr` is NULL at cold
+                // construction; the guard inside the kernel (sentinel check)
+                // is on-device so no host-side guard is needed.
+                if self.isv_signals_dev_ptr != 0 {
+                    use crate::cuda_pipeline::sp14_isv_slots::H_S2_AUX_RMS_EMA_INDEX;
+                    self.exp_h_s2_aux_rms_ema_ops.launch(
+                        &self.stream,
+                        self.exp_h_s2_aux.raw_ptr(),
+                        n as i32,
+                        sh2 as i32,
+                        self.isv_signals_dev_ptr,
+                        H_S2_AUX_RMS_EMA_INDEX as i32,
+                        super::gpu_aux_trunk::HS2AuxRmsEmaOps::ALPHA,
+                    )?;
+                }
+
                 // Step B-2: aux next-bar forward. SP14-C.5b: input redirected
                 // from `exp_h_s2_f32` (Q trunk output) to `exp_h_s2_aux`
                 // (aux trunk output, written above). Writes hidden + logits

crates/ml/src/cuda_pipeline/h_s2_aux_rms_ema_kernel.cu

@@ -0,0 +1,90 @@
+/* h_s2_aux_rms_ema — GPU-driven RMS(h_s2_aux) step observation into
+ * `ISV[H_S2_AUX_RMS_EMA_INDEX=449]` for Pearls A+D consumption.
+ *
+ * SP14 Layer C Phase C.6 (2026-05-08): mirrors `h_s2_rms_ema_kernel.cu`
+ * (slot 96, Q's GRN trunk output) for the separate auxiliary trunk output.
+ * Input: `h_s2_aux [B, SH2]` (aux trunk final output, no activation).
+ * Output: writes directly to `ISV[isv_target_idx]` (slot 449).
+ *
+ * GPU-drives-CPU-reads (spec §4.C.6) and
+ * pearl_cold_path_no_exception_to_gpu_drives.md compliant.
+ *
+ * Single-block kernel (256 threads, shmem-reduction). No atomicAdd
+ * (per `feedback_no_atomicadd.md`); the in-block tree reduction is
+ * standard `s >>= 1` with `__syncthreads()`. Only thread 0 writes the
+ * step observation.
+ *
+ * Pearl-A first-observation bootstrap: sentinel = 0.0; first valid
+ * observation REPLACES directly (no blend). Subsequent observations
+ * use fixed α=0.05 EMA blend (per `pearl_first_observation_bootstrap.md`).
+ * ISV slot 449 is outside the SP4/SP5 wiener_state_buf linear span so
+ * Pearl-D Wiener-optimal adaptive α is not available — fixed α=0.05 is
+ * the correct per-step cadence blend rate consistent with the aux trunk
+ * producer launch frequency.
+ *
+ * Args:
+ *   h_s2_aux      — [B, SH2] row-major f32 device ptr.
+ *   B             — batch size.
+ *   SH2           — aux trunk output width (AUX_HIDDEN_DIM = 256).
+ *   isv           — [ISV_TOTAL_DIM] f32 device ptr (mapped-pinned ISV bus).
+ *   isv_target_idx — H_S2_AUX_RMS_EMA_INDEX (449).
+ *   alpha         — EMA blend rate (0.05 for per-step producer cadence).
+ *
+ * Cost: one launch per collector step (same cadence as aux_trunk_forward).
+ * 256 threads × ceil(B*SH2/256) strided loads + log2(256)=8 reduce steps
+ * + 1 global write.
+ */
+
+extern "C" __global__ void h_s2_aux_rms_ema_update(
+    const float* __restrict__ h_s2_aux,   /* [B, SH2] row-major */
+    int          B,
+    int          SH2,
+    float* __restrict__ isv,              /* [ISV_TOTAL_DIM] mapped-pinned */
+    int          isv_target_idx,          /* H_S2_AUX_RMS_EMA_INDEX (449) */
+    float        alpha                    /* EMA blend rate (0.05) */
+) {
+    /* Single-block contract — no atomicAdd; only block 0 runs. */
+    if (blockIdx.x != 0) return;
+
+    extern __shared__ float smem[];
+    const int tid   = (int)threadIdx.x;
+    const int block = (int)blockDim.x;
+    const int N     = B * SH2;
+
+    /* Per-thread sum of squares over a strided slice. */
+    float local_sumsq = 0.0f;
+    for (int i = tid; i < N; i += block) {
+        const float v = h_s2_aux[i];
+        local_sumsq += v * v;
+    }
+    smem[tid] = local_sumsq;
+    __syncthreads();
+
+    /* In-block tree reduction (no atomicAdd). */
+    for (int s = block / 2; s > 0; s >>= 1) {
+        if (tid < s) smem[tid] += smem[tid + s];
+        __syncthreads();
+    }
+
+    /* Thread 0 writes the EMA-blended step observation.
+     * B > 0 and SH2 > 0 are constructor-guaranteed invariants; N = B*SH2 > 0
+     * always holds. No defensive ternary — a zero N would produce NaN which
+     * propagates loudly through downstream consumers.
+     *
+     * Pearl-A first-observation bootstrap: sentinel 0.0 → REPLACE directly.
+     * Steady-state: EMA blend with caller-supplied α. */
+    if (tid == 0) {
+        const float rms     = sqrtf(smem[0] / (float)N);
+        const float current = isv[isv_target_idx];
+
+        float next;
+        if (current <= 0.0f) {
+            /* Pearl-A: sentinel or cold-start zero — replace outright. */
+            next = rms;
+        } else {
+            next = (1.0f - alpha) * current + alpha * rms;
+        }
+        isv[isv_target_idx] = next;
+        __threadfence_system();   /* PCIe-visible write for mapped pinned host_ptr */
+    }
+}

crates/ml/tests/aux_trunk_oracle_tests.rs

@@ -1165,4 +1165,131 @@ mod gpu {
         );
         eprintln!("No-winning-trades guard: H={h_after} (sentinel preserved)");
     }
+
+    // ════════════════════════════════════════════════════════════════════
+    // SP14 Layer C Phase C.6 oracle tests (2026-05-08).
+    //
+    // Added 1 test exercising the h_s2_aux RMS EMA producer:
+    //
+    //   - h_s2_aux_rms_ema_pearl_a_bootstrap — sentinel REPLACED on first
+    //     observation; EMA blend applied on second.
+    //
+    // Plan: docs/superpowers/plans/2026-05-07-sp14-layer-c-separate-aux-trunk.md §C.6
+    // ════════════════════════════════════════════════════════════════════
+
+    const H_S2_AUX_RMS_EMA_CUBIN: &[u8] =
+        include_bytes!(concat!(env!("OUT_DIR"), "/h_s2_aux_rms_ema_kernel.cubin"));
+
+    /// H_S2_AUX_RMS_EMA_INDEX — must match `sp14_isv_slots.rs`.
+    const H_S2_AUX_RMS_EMA_IDX: usize = 449;
+
+    fn load_h_s2_aux_rms_ema(stream: &Arc<CudaStream>) -> CudaFunction {
+        let module = stream
+            .context()
+            .load_cubin(H_S2_AUX_RMS_EMA_CUBIN.to_vec())
+            .expect("load h_s2_aux_rms_ema cubin");
+        module
+            .load_function("h_s2_aux_rms_ema_update")
+            .expect("load h_s2_aux_rms_ema_update function")
+    }
+
+    /// Launch `h_s2_aux_rms_ema_update` once and synchronize.
+    fn run_h_s2_aux_rms_ema(
+        stream: &Arc<CudaStream>,
+        kernel: &CudaFunction,
+        h_s2_aux_buf: &MappedF32Buffer,
+        b: i32,
+        sh2: i32,
+        isv_buf: &MappedF32Buffer,
+        alpha: f32,
+    ) {
+        let h_ptr = h_s2_aux_buf.dev_ptr;
+        let isv_ptr = isv_buf.dev_ptr;
+        let isv_idx = H_S2_AUX_RMS_EMA_IDX as i32;
+        let smem = (256 * std::mem::size_of::<f32>()) as u32;
+        unsafe {
+            stream
+                .launch_builder(kernel)
+                .arg(&h_ptr)
+                .arg(&b)
+                .arg(&sh2)
+                .arg(&isv_ptr)
+                .arg(&isv_idx)
+                .arg(&alpha)
+                .launch(LaunchConfig {
+                    grid_dim: (1, 1, 1),
+                    block_dim: (256, 1, 1),
+                    shared_mem_bytes: smem,
+                })
+                .expect("h_s2_aux_rms_ema_update launch");
+        }
+        stream.synchronize().expect("sync h_s2_aux_rms_ema");
+    }
+
+    /// Step C.6 oracle: Pearl-A first-observation bootstrap.
+    ///
+    /// Scenario A — sentinel → replace:
+    ///   Set ISV[449] = 0.0 (sentinel). Fill `h_s2_aux [B, SH2]` with
+    ///   constant value 1.5. Expected RMS = sqrt(1.5² × N / N) = 1.5.
+    ///   After one launch the kernel MUST replace (not blend): ISV[449] ≈ 1.5.
+    ///
+    /// Scenario B — steady-state EMA blend:
+    ///   Fill `h_s2_aux` with 3.0 (expected RMS = 3.0). Relaunch with
+    ///   α=0.05 from current=1.5.
+    ///   Expected: (1-0.05)*1.5 + 0.05*3.0 = 1.5*0.95 + 0.15 = 1.575.
+    ///   ISV[449] must be strictly between 1.5 and 3.0 and close to 1.575.
+    #[test]
+    #[ignore = "requires GPU"]
+    fn h_s2_aux_rms_ema_pearl_a_bootstrap() {
+        const B: i32 = 4;
+        const SH2: i32 = 256;
+        const ALPHA: f32 = 0.05;
+
+        let stream = make_test_stream();
+        let kernel = load_h_s2_aux_rms_ema(&stream);
+
+        // Allocate ISV buffer; zero-initialise to set sentinel.
+        let isv_buf = unsafe { MappedF32Buffer::new(ISV_TEST_SIZE) }
+            .expect("ISV buf alloc");
+        let host = unsafe {
+            std::slice::from_raw_parts_mut(isv_buf.host_ptr, ISV_TEST_SIZE)
+        };
+        // sentinel = 0.0 (zero-init already)
+
+        // Allocate h_s2_aux and fill with constant 1.5.
+        let n = (B * SH2) as usize;
+        let h_buf = unsafe { MappedF32Buffer::new(n) }.expect("h_s2_aux buf alloc");
+        let h_host = unsafe { std::slice::from_raw_parts_mut(h_buf.host_ptr, n) };
+        for v in h_host.iter_mut() { *v = 1.5; }
+
+        // ── Scenario A: Pearl-A bootstrap ─────────────────────────────
+        run_h_s2_aux_rms_ema(&stream, &kernel, &h_buf, B, SH2, &isv_buf, ALPHA);
+
+        let slot_a = host[H_S2_AUX_RMS_EMA_IDX];
+        assert!(
+            (slot_a - 1.5).abs() < 1e-5,
+            "Pearl-A bootstrap: expected ISV[449] = 1.5 (replace sentinel), got {slot_a}",
+        );
+        eprintln!("Pearl-A bootstrap (h=1.5): ISV[449] = {slot_a:.6} (expected 1.5)");
+
+        // ── Scenario B: steady-state EMA blend ────────────────────────
+        for v in h_host.iter_mut() { *v = 3.0; }
+
+        run_h_s2_aux_rms_ema(&stream, &kernel, &h_buf, B, SH2, &isv_buf, ALPHA);
+
+        let slot_b = host[H_S2_AUX_RMS_EMA_IDX];
+        let expected = (1.0 - ALPHA) * 1.5 + ALPHA * 3.0;  // 1.575
+        assert!(
+            (slot_b - expected).abs() < 1e-4,
+            "EMA blend: expected ISV[449] ≈ {expected:.6} after α={ALPHA} step from 1.5→3.0, \
+             got {slot_b:.6}",
+        );
+        assert!(
+            slot_b > 1.5 && slot_b < 3.0,
+            "EMA blend: ISV[449]={slot_b} must be strictly between 1.5 and 3.0",
+        );
+        eprintln!(
+            "EMA blend (h=3.0, α={ALPHA}): ISV[449] = {slot_b:.6} (expected {expected:.6})"
+        );
+    }
 }

docs/dqn-wire-up-audit.md

@@ -7618,7 +7618,56 @@ BEFORE the aux trunk forward runs.
 
 C.5b only wires the trunk forward+backward+Adam mechanics. Phases C.6+
 (separate commits) layer on:
-- C.6: `h_s2_aux_rms_ema` producer kernel feeding ISV[449].
+- C.6: `h_s2_aux_rms_ema` producer kernel feeding ISV[449]. **DONE — see below.**
 - C.7-C.8: ISV-driven aux trunk Adam β1/β2/ε.
 - C.9: synthetic-data smoke + audit close-out + memory pearls.
 - C.10: L40S 30-epoch validation.
+
+---
+
+## 2026-05-08 — SP14 Layer C Phase C.6: h_s2_aux_rms_ema producer kernel
+
+**Commit:** (this commit) **Branch:** sp15-phase1-honest-numbers
+
+### What was wired
+
+| File | Change |
+|------|--------|
+| `h_s2_aux_rms_ema_kernel.cu` | New single-block 256-thread CUDA kernel. Computes `RMS = sqrt(mean(h_s2_aux²))` over `[B, SH2]`, writes directly to `ISV[isv_target_idx]` with Pearl-A bootstrap (sentinel 0.0 → replace) + fixed α=0.05 EMA blend. `__threadfence_system()` after write for PCIe visibility. No atomicAdd — shmem tree-reduce. |
+| `build.rs` | Added `"h_s2_aux_rms_ema_kernel.cu"` to cubin manifest after `avg_win_hold_time_update_kernel.cu`. |
+| `gpu_dqn_trainer.rs` | Added `pub(crate) static H_S2_AUX_RMS_EMA_CUBIN` after `AVG_WIN_HOLD_TIME_UPDATE_CUBIN`. |
+| `gpu_aux_trunk.rs` | Added `HS2AuxRmsEmaOps` struct (mirrors `AvgWinHoldTimeUpdateOps`). Owns `CudaFunction` pre-loaded at construction. `launch()` takes `h_s2_aux_ptr, b, sh2, isv_ptr, isv_target_idx, alpha`. Added `H_S2_AUX_RMS_EMA_CUBIN` to the import list. |
+| `gpu_experience_collector.rs` | Added `exp_h_s2_aux_rms_ema_ops: HS2AuxRmsEmaOps` field. Initialized in constructor. Launched after `exp_aux_trunk_forward_ops.launch()` (Step B-1b) when `isv_signals_dev_ptr != 0`. |
+| `aux_trunk_oracle_tests.rs` | Added `h_s2_aux_rms_ema_pearl_a_bootstrap` test: Scenario A verifies sentinel→replace (ISV[449]≈1.5 from constant h=1.5), Scenario B verifies EMA blend (ISV[449]≈1.575 from h=3.0 with α=0.05). |
+
+### Design decision: no scratch+Pearls-A+D path
+
+ISV slot 449 is in the SP14 Layer C block (outside the SP4/SP5 wiener_state_buf
+linear span). The wiener buffer is indexed by `(SP4_PRODUCER_COUNT + SP5_PRODUCER_COUNT) * 3 = 840`
+floats covering ISV slots up to `174 + 209 - 1 = 382`. Slot 449 is 66 positions
+beyond the SP5 end, making scratch_idx=N → wiener_offset=N*3 out-of-bounds for any
+scratch slot outside the SP4/SP5 range. The correct pattern — established by
+`avg_win_hold_time_update_kernel.cu` (slot 451) — embeds Pearl-A bootstrap and
+EMA blend logic directly in the kernel body and writes to ISV in-place. Fixed
+α=0.05 replaces the Wiener-optimal adaptive α (which requires wiener state storage).
+
+### Verification
+
+- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --tests` — clean.
+- `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml --test aux_trunk_oracle_tests --release -- --ignored --nocapture` — 13/13 pass (12 prior + 1 new `h_s2_aux_rms_ema_pearl_a_bootstrap`).
+
+### Wire status — POST-C.6
+
+| Component | Status |
+|-----------|--------|
+| `h_s2_aux_rms_ema_kernel.cu` | LIVE — compiled to cubin via build.rs |
+| `HS2AuxRmsEmaOps` (gpu_aux_trunk.rs) | LIVE — struct with pre-loaded CudaFunction |
+| `H_S2_AUX_RMS_EMA_CUBIN` (gpu_dqn_trainer.rs) | LIVE — cubin static |
+| `exp_h_s2_aux_rms_ema_ops` (collector) | LIVE — launched post aux_trunk_forward when ISV ptr set |
+| `ISV[H_S2_AUX_RMS_EMA_INDEX=449]` | LIVE — written per collector step |
+
+### Outstanding for C.7+
+
+- C.7-C.8: ISV-driven aux trunk Adam β1/β2/ε/LR/grad-clip.
+- C.9: synthetic-data smoke + audit close-out + memory pearls.
+- C.10: L40S 30-epoch validation.