diff --git a/crates/ml/build.rs b/crates/ml/build.rs index c2baba98e..523be9ab1 100644 --- a/crates/ml/build.rs +++ b/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 diff --git a/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs b/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs index 1a7854ae5..96129b4ce 100644 --- a/crates/ml/src/cuda_pipeline/gpu_aux_trunk.rs +++ b/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) -> Result { + 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, + 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::() 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(()) + } +} diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index d00222e3f..8537ce09c 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/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 diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index 31b7e1a20..f75d526d8 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/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 diff --git a/crates/ml/src/cuda_pipeline/h_s2_aux_rms_ema_kernel.cu b/crates/ml/src/cuda_pipeline/h_s2_aux_rms_ema_kernel.cu new file mode 100644 index 000000000..2d85b1fe1 --- /dev/null +++ b/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 */ + } +} diff --git a/crates/ml/tests/aux_trunk_oracle_tests.rs b/crates/ml/tests/aux_trunk_oracle_tests.rs index db933beb3..326ec0d73 100644 --- a/crates/ml/tests/aux_trunk_oracle_tests.rs +++ b/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) -> 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, + 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::()) 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})" + ); + } } diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 6cdb6b224..be5d08d3c 100644 --- a/docs/dqn-wire-up-audit.md +++ b/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.