From 5f800fe5c871ed66fca08122b2a066dee03b5d19 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 1 May 2026 02:18:55 +0200 Subject: [PATCH] =?UTF-8?q?feat(sp4):=20Task=20A13.4=20=E2=80=94=20retrofi?= =?UTF-8?q?t=20iqn=5Fquantile=5Fema=20with=20Pearls=20A+D?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Replaces the kernel's hardcoded `ema_alpha` with the shared `pearls_ad_update` host-side helper. 4 ISV slots retrofit (off-median IQN quantiles). - Kernel `iqn_quantile_ema_update` signature: drops `(isv, 4 isv_*_idx, ema_alpha)` for `(scratch_buf, scratch_first_index=59)`. Block dispatch unchanged (4 blocks × 256 threads); each block writes one step observation to scratch_buf[scratch_first_index + slot_offset] for slot_offset ∈ {0,1,2,3}. - 4 ISV slots wired with Pearls A+D: ISV[99/100/101/102] (Wiener offsets 177/180/183/186). Median tau_idx=2 intentionally skipped (already surfaced via greedy-Q). - Wrapper `GpuDqnTrainer::launch_iqn_quantile_ema(..., _ema_alpha_unused)`: keeps early-return-on-NULL guard; sync + Pearls A+D loop over 4 SLOT_PAIRS. Behavior: stationary signals converge to the same value at adaptive rate. The 4 slots remain diagnostic-only (HEALTH_DIAG / risk-monitoring surface). Tests: `sp4_iqn_quantile_ema_writes_step_obs_via_pearl_a_then_converges_pearl_d` drives kernel with B=32 Q=5 TBA=12 controlled q surface where Q[a, b*Q+tau_idx] = (tau_idx+1)*0.7. Asserts each slot ∈ ±1e-5 of expected, median absent, non-target slots remain 0, Pearl A bootstrap + Pearl D convergence verified. Per `feedback_no_atomicadd.md`, `feedback_no_htod_htoh_only_mapped_pinned.md`. Build: `cargo check -p ml --lib --tests --offline` clean (11 pre-existing warnings, no new warnings). Co-Authored-By: Claude Opus 4.7 (1M context) --- .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 76 +++++++++-- .../cuda_pipeline/iqn_quantile_ema_kernel.cu | 87 +++++++------ crates/ml/tests/sp4_producer_unit_tests.rs | 120 ++++++++++++++++++ docs/dqn-wire-up-audit.md | 2 + 4 files changed, 232 insertions(+), 53 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index b11304ed0..01a06e9fb 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -10854,8 +10854,19 @@ impl GpuDqnTrainer { save_q_online_ptr: u64, tba: usize, num_quantiles: usize, - ema_alpha: f32, + _ema_alpha_unused: f32, ) -> Result<(), MLError> { + // SP4 Layer A Task A13.4 retrofit (2026-05-01): kernel writes 4 + // step observations (mean |Q| at p05/p25/p75/p95) to + // `producer_step_scratch_buf[59..63)`. Host applies Pearls A+D + // host-side, mapping each scratch slot to its ISV slot (99/100/ + // 101/102) and updating Wiener-EMA state. + // + // **α dropped per SP4 — Pearls A+D adapt α from per-slot signal- + // vs-noise variance.** The `_ema_alpha_unused` argument is preserved + // so callers (training_loop.rs / fused_training.rs) compile unchanged. + use crate::cuda_pipeline::sp4_wiener_ema::{pearls_ad_update, WienerState}; + if save_q_online_ptr == 0 { return Ok(()); } @@ -10864,14 +10875,14 @@ impl GpuDqnTrainer { debug_assert!(num_quantiles == 5, "launch_iqn_quantile_ema: Plan 4 Task 3 (E.3) pins IQN_NUM_QUANTILES = 5; \ got num_quantiles={num_quantiles}"); + + const SCRATCH_FIRST: usize = 59; + let b_i = self.config.batch_size as i32; let q_i = num_quantiles as i32; let tba_i = tba as i32; - let isv_dev_ptr = self.isv_signals_dev_ptr; - let p05 = IQN_Q_P05_EMA_INDEX as i32; - let p25 = IQN_Q_P25_EMA_INDEX as i32; - let p75 = IQN_Q_P75_EMA_INDEX as i32; - let p95 = IQN_Q_P95_EMA_INDEX as i32; + let scratch_dev = self.producer_step_scratch_buf.dev_ptr; + let scratch_first_i = SCRATCH_FIRST as i32; const BLOCK_DIM: u32 = 256; let smem_bytes = BLOCK_DIM * std::mem::size_of::() as u32; unsafe { @@ -10880,12 +10891,8 @@ impl GpuDqnTrainer { .arg(&b_i) .arg(&q_i) .arg(&tba_i) - .arg(&isv_dev_ptr) - .arg(&p05) - .arg(&p25) - .arg(&p75) - .arg(&p95) - .arg(&ema_alpha) + .arg(&scratch_dev) + .arg(&scratch_first_i) .launch(LaunchConfig { grid_dim: (4, 1, 1), block_dim: (BLOCK_DIM, 1, 1), @@ -10893,6 +10900,51 @@ impl GpuDqnTrainer { }) .map_err(|e| MLError::ModelError(format!("iqn_quantile_ema_update: {e}")))?; } + self.stream.synchronize() + .map_err(|e| MLError::ModelError(format!("iqn_quantile_ema sync: {e}")))?; + + // ── Pearls A+D host-side update for all 4 quantile slots ── + // Order matches the kernel's `slot_offset` mapping: + // +0 → p05 (ISV[99]), +1 → p25 (ISV[100]), + // +2 → p75 (ISV[101]), +3 → p95 (ISV[102]). + const SLOT_PAIRS: [(usize, usize); 4] = [ + (0, IQN_Q_P05_EMA_INDEX), + (1, IQN_Q_P25_EMA_INDEX), + (2, IQN_Q_P75_EMA_INDEX), + (3, IQN_Q_P95_EMA_INDEX), + ]; + for &(off, isv_idx) in SLOT_PAIRS.iter() { + let scratch_idx = SCRATCH_FIRST + off; + let step_obs = unsafe { + std::ptr::read_volatile(self.producer_step_scratch_buf.host_ptr.add(scratch_idx)) + }; + // Skip degenerate-zero observations — IQN's per-quantile + // mean|Q| is non-zero in steady state. A 0 here means the + // kernel hasn't run (cold start before IQN forward populates + // `save_q_online`). + if step_obs == 0.0 { continue; } + + let wiener_offset = scratch_idx * 3; + let prev_x_mean = unsafe { + std::ptr::read_volatile(self.isv_signals_pinned.add(isv_idx)) + }; + let mut state = unsafe { + let p = self.wiener_state_buf.host_ptr; + WienerState { + sample_var: std::ptr::read_volatile(p.add(wiener_offset)), + diff_var: std::ptr::read_volatile(p.add(wiener_offset + 1)), + x_lag: std::ptr::read_volatile(p.add(wiener_offset + 2)), + } + }; + let new_x_mean = pearls_ad_update(prev_x_mean, &mut state, step_obs); + unsafe { + std::ptr::write_volatile(self.isv_signals_pinned.add(isv_idx), new_x_mean); + let wp = self.wiener_state_buf.host_ptr; + std::ptr::write_volatile(wp.add(wiener_offset), state.sample_var); + std::ptr::write_volatile(wp.add(wiener_offset + 1), state.diff_var); + std::ptr::write_volatile(wp.add(wiener_offset + 2), state.x_lag); + } + } Ok(()) } diff --git a/crates/ml/src/cuda_pipeline/iqn_quantile_ema_kernel.cu b/crates/ml/src/cuda_pipeline/iqn_quantile_ema_kernel.cu index 68710757a..6f8550cf9 100644 --- a/crates/ml/src/cuda_pipeline/iqn_quantile_ema_kernel.cu +++ b/crates/ml/src/cuda_pipeline/iqn_quantile_ema_kernel.cu @@ -1,35 +1,42 @@ /* - * iqn_quantile_ema_update — Plan 4 Task 3 (E.3) producer kernel. + * iqn_quantile_ema_update — Plan 4 Task 3 (E.3) producer kernel; SP4 Layer + * A Task A13.4 retrofit (2026-05-01). * * Reads the IQN online forward's per-quantile Q surface * (`save_q_online [TBA, B*Q]` col-major; populated by - * `GpuIqnHead::execute_training_pipeline` after each step) and EMAs the - * mean Q at each off-median fixed-τ position into ISV[99..103): + * `GpuIqnHead::execute_training_pipeline` after each step) and writes the + * mean |Q(s,a;τ)| step observation at each off-median fixed-τ position + * to `producer_step_scratch_buf[scratch_first..scratch_first+4)` for + * Pearls A+D consumption: * * FIXED_TAUS = [0.05, 0.25, 0.50, 0.75, 0.95] * - * blockIdx.x = 0 → τ_idx 0 (τ=0.05) → ISV[IQN_Q_P05_EMA_INDEX = 99] - * blockIdx.x = 1 → τ_idx 1 (τ=0.25) → ISV[IQN_Q_P25_EMA_INDEX = 100] - * blockIdx.x = 2 → τ_idx 3 (τ=0.75) → ISV[IQN_Q_P75_EMA_INDEX = 101] - * blockIdx.x = 3 → τ_idx 4 (τ=0.95) → ISV[IQN_Q_P95_EMA_INDEX = 102] + * blockIdx.x = 0 → τ_idx 0 (τ=0.05) → scratch[scratch_first+0] → ISV[99] + * blockIdx.x = 1 → τ_idx 1 (τ=0.25) → scratch[scratch_first+1] → ISV[100] + * blockIdx.x = 2 → τ_idx 3 (τ=0.75) → scratch[scratch_first+2] → ISV[101] + * blockIdx.x = 3 → τ_idx 4 (τ=0.95) → scratch[scratch_first+3] → ISV[102] * * The median (τ=0.50, idx=2) is intentionally skipped — it is already * surfaced via the existing greedy-Q diagnostic, no duplication needed. * - * Reduction layout (one block per ISV slot): + * Reduction layout (one block per scratch slot): * For target τ_idx t, walk every (sample b, action a) and read - * save_q_online[a + (b*Q + t) * TBA]. Sum across (B × TBA) entries, - * divide by N = B × TBA, EMA into ISV[target_slot]. Single block, + * save_q_online[a + (b*Q + t) * TBA]. Sum |Q| across (B × TBA) entries, + * divide by N = B × TBA, write to scratch slot. Single block, * 256-thread shmem-tree reduce — no atomicAdd * (per `feedback_no_atomicadd.md`). * - * Cold-path-cadence: launched once per training step alongside - * `h_s2_rms_ema_update`. Diagnostic only — no consumer kernel reads - * slots [99..103) in this commit. ISV is pinned device-mapped, so host - * HEALTH_DIAG sees the new values without an explicit DtoH. + * **α dropped per SP4 — Pearls A+D adapt α from per-slot signal-vs-noise + * variance.** The host launcher + * (`GpuDqnTrainer::launch_iqn_quantile_ema`) syncs the stream and applies + * Pearls A+D host-side via `pearls_ad_update`, mapping each scratch slot + * to its ISV bound slot (99/100/101/102) and updating Wiener-EMA state + * for each. * - * The EMA α parameter is passed by the host (matches the `ema_alpha` - * already plumbed into the other Plan 4 producers from `training_loop.rs`). + * Cold-path-cadence: launched once per training step alongside + * `h_s2_rms_ema_update`. Diagnostic only — no consumer kernel reads slots + * [99..103). ISV is pinned device-mapped, so host HEALTH_DIAG sees the + * Pearls-A+D-updated values without an explicit DtoH. */ #include @@ -39,30 +46,28 @@ extern "C" __global__ void iqn_quantile_ema_update( int batch_size, /* B */ int num_quantiles, /* Q (= 5 under fixed-τ) */ int tba, /* total branch actions */ - float* __restrict__ isv, /* pinned device-mapped ISV bus */ - int isv_q_p05_idx, /* IQN_Q_P05_EMA_INDEX = 99 */ - int isv_q_p25_idx, /* = 100 */ - int isv_q_p75_idx, /* = 101 */ - int isv_q_p95_idx, /* = 102 */ - float ema_alpha /* per-step EMA rate (0,1) */ + float* __restrict__ scratch_buf, /* producer_step_scratch_buf */ + int scratch_first_index /* slot 59 — first of 4 contiguous slots */ ) { - /* Map blockIdx.x → (τ_idx within FIXED_TAUS, target ISV slot). - * Skip the median (idx=2) — that's the existing greedy-Q diagnostic. */ + /* Map blockIdx.x → (τ_idx within FIXED_TAUS, target scratch slot). + * Skip the median (idx=2) — that's the existing greedy-Q diagnostic. + * Scratch slot order matches blockIdx.x: p05 → +0, p25 → +1, p75 → +2, + * p95 → +3. */ int tau_idx; - int target_slot; + int slot_offset; switch (blockIdx.x) { - case 0: tau_idx = 0; target_slot = isv_q_p05_idx; break; - case 1: tau_idx = 1; target_slot = isv_q_p25_idx; break; - case 2: tau_idx = 3; target_slot = isv_q_p75_idx; break; - case 3: tau_idx = 4; target_slot = isv_q_p95_idx; break; + case 0: tau_idx = 0; slot_offset = 0; break; + case 1: tau_idx = 1; slot_offset = 1; break; + case 2: tau_idx = 3; slot_offset = 2; break; + case 3: tau_idx = 4; slot_offset = 3; break; default: return; } /* Bounds check — caller must launch with grid.x = 4 and num_quantiles - * matching FIXED_TAUS.len() = 5. The defensive guard returns the EMA - * unchanged (no write) if the layout is incompatible, surfacing the - * misconfiguration to the host via a stale ISV value rather than a - * silent bad-index read. */ + * matching FIXED_TAUS.len() = 5. The defensive guard skips the + * scratch write if the layout is incompatible, surfacing the + * misconfiguration to the host via a stale (zero or prior-step) value + * rather than a silent bad-index read. */ if (tau_idx >= num_quantiles) return; extern __shared__ float smem[]; @@ -92,15 +97,15 @@ extern "C" __global__ void iqn_quantile_ema_update( __syncthreads(); } - /* Thread 0 writes the EMA. The host caller's trainer config requires - * batch_size > 0 and tba > 0, so N > 0 is invariant; no defensive - * `(N > 0) ? N : 1` here — masking the divide would silently EMA - * toward 0 if the invariant ever broke; the NaN that division-by-zero - * produces propagates to the ISV slot and downstream HEALTH_DIAG - * loudly visible. Same convention as h_s2_rms_ema_update. */ + /* Thread 0 writes the step observation. The host caller's trainer + * config requires batch_size > 0 and tba > 0, so N > 0 is invariant; + * no defensive `(N > 0) ? N : 1` here — masking the divide would + * silently emit 0 if the invariant ever broke; the NaN that + * division-by-zero produces propagates loudly through Pearls A+D to + * the ISV slot. */ if (tid == 0) { const float mean_abs_q = smem[0] / (float)N; - const float prev = isv[target_slot]; - isv[target_slot] = (1.0f - ema_alpha) * prev + ema_alpha * mean_abs_q; + scratch_buf[scratch_first_index + slot_offset] = mean_abs_q; + __threadfence_system(); } } diff --git a/crates/ml/tests/sp4_producer_unit_tests.rs b/crates/ml/tests/sp4_producer_unit_tests.rs index b38ec192e..651fe542c 100644 --- a/crates/ml/tests/sp4_producer_unit_tests.rs +++ b/crates/ml/tests/sp4_producer_unit_tests.rs @@ -1914,3 +1914,123 @@ fn sp4_vsn_mask_ema_writes_step_obs_via_pearl_a_then_converges_pearl_d() { "Pearl D convergence: got {x_mean}, expected {expected_group_0}", ); } + +// ── SP4 Task A13.4: iqn_quantile_ema retrofit ──────────────────────────────── + +const SP4_IQN_QUANTILE_EMA_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/iqn_quantile_ema_kernel.cubin")); + +fn load_iqn_quantile_ema_kernel(stream: &Arc) -> CudaFunction { + let module = stream + .context() + .load_cubin(SP4_IQN_QUANTILE_EMA_CUBIN.to_vec()) + .expect("load iqn_quantile_ema cubin"); + module + .load_function("iqn_quantile_ema_update") + .expect("load iqn_quantile_ema_update function") +} + +/// SP4 Task A13.4: `iqn_quantile_ema_update` writes 4 mean-|Q| step +/// observations to scratch[59..63) (one per non-median quantile). Pearls +/// A+D bootstrap+convergence verified host-side. +/// +/// Test surface: `save_q_online [TBA, B*Q]` col-major where +/// Q[a, (b*Q + tau_idx)] = (tau_idx + 1) × 0.7. Each tau's mean|Q| +/// matches its (tau_idx + 1) × 0.7 setpoint. Skipped median (tau_idx=2) +/// is consequently ignored by the kernel. +#[test] +#[ignore = "requires GPU"] +fn sp4_iqn_quantile_ema_writes_step_obs_via_pearl_a_then_converges_pearl_d() { + use ml::cuda_pipeline::sp4_wiener_ema::{pearls_ad_update, WienerState}; + + const B: usize = 32; + const Q: usize = 5; + const TBA: usize = 12; + const SCRATCH_FIRST: usize = 59; + const SP4_PRODUCER_COUNT: usize = 69; + const BLOCK_DIM: u32 = 256; + const SHARED_BYTES: u32 = BLOCK_DIM * std::mem::size_of::() as u32; + + // save_q_online[a + (b*Q + tau_idx)*TBA] = (tau_idx+1) * 0.7 + let mut q_data = vec![0.0_f32; TBA * B * Q]; + for b in 0..B { + for tau in 0..Q { + let val = (tau as f32 + 1.0) * 0.7; + for a in 0..TBA { + let col = b * Q + tau; + let flat = a + col * TBA; + q_data[flat] = val; + } + } + } + + let stream = make_test_stream(); + let kernel = load_iqn_quantile_ema_kernel(&stream); + + let q_buf = unsafe { MappedF32Buffer::new(q_data.len()) }.expect("alloc q buf"); + q_buf.write_from_slice(&q_data); + let scratch_buf = unsafe { MappedF32Buffer::new(SP4_PRODUCER_COUNT) } + .expect("alloc scratch"); + + let q_dev = q_buf.dev_ptr; + let scratch_dev = scratch_buf.dev_ptr; + let b_i32 = B as i32; + let q_i32 = Q as i32; + let tba_i32 = TBA as i32; + let scratch_first_i32 = SCRATCH_FIRST as i32; + + unsafe { + stream.launch_builder(&kernel) + .arg(&q_dev) + .arg(&b_i32) + .arg(&q_i32) + .arg(&tba_i32) + .arg(&scratch_dev) + .arg(&scratch_first_i32) + .launch(LaunchConfig { + grid_dim: (4, 1, 1), + block_dim: (BLOCK_DIM, 1, 1), + shared_mem_bytes: SHARED_BYTES, + }) + .expect("launch iqn_quantile_ema_update"); + } + stream.synchronize().expect("sync"); + + let host = scratch_buf.read_all(); + for (i, &v) in host.iter().enumerate() { + let in_target = i >= SCRATCH_FIRST && i < SCRATCH_FIRST + 4; + if !in_target { + assert_eq!(v, 0.0, "wrote outside target slots: scratch[{i}]={v}"); + } + } + + // Slot order: +0 → p05 (tau_idx=0, expected 0.7), + // +1 → p25 (tau_idx=1, expected 1.4), + // +2 → p75 (tau_idx=3, expected 2.8), + // +3 → p95 (tau_idx=4, expected 3.5). + let expected: [f32; 4] = [0.7, 1.4, 2.8, 3.5]; + for (off, &exp) in expected.iter().enumerate() { + let actual = host[SCRATCH_FIRST + off]; + assert!( + (actual - exp).abs() < 1e-5, + "slot offset {off}: got {actual}, expected {exp}", + ); + } + + // Pearl A bootstrap on slot +0 (p05). + let p05_obs = host[SCRATCH_FIRST]; + let mut state = WienerState::ZERO; + let new_x_mean = pearls_ad_update(0.0, &mut state, p05_obs); + assert_eq!(new_x_mean, p05_obs); + assert_eq!(state.x_lag, p05_obs); + + // Pearl D convergence on the same slot. + let mut x_mean = new_x_mean; + for _ in 0..1000 { + x_mean = pearls_ad_update(x_mean, &mut state, p05_obs); + } + assert!( + ((x_mean - 0.7) / 0.7).abs() < 0.01, + "Pearl D: stationary p05 converged to {x_mean} (expected 0.7)", + ); +} diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 6d84d8749..ef5fbbf7a 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2332,3 +2332,5 @@ SP4 Layer A Task A13.1 — aux_heads_loss_ema + aux_label_scale_ema retrofit (Pe SP4 Layer A Task A13.2 — moe_expert_util_ema retrofit (Pearls A+D) (2026-05-01): retrofit the existing per-step MoE expert-utilisation + gate-entropy producer in `crates/ml/src/cuda_pipeline/moe_kernels.cu::moe_expert_util_ema_update`. Kernel signature changed: drops `(isv, isv_util_base, isv_entropy_index, alpha)` for `(scratch_buf, scratch_idx_util_base=44, scratch_idx_entropy=52)`. Single-block single-thread; computes per-expert col_mean of `gate [B, K]` in a single forward pass, writing each to `scratch_buf[scratch_idx_util_base + k]` for k in 0..K, and accumulating Shannon entropy (`-Σ col_mean·ln(col_mean)` with 1e-9 floor) for slot `scratch_idx_entropy`. The dual-pass formulation in the prior implementation (one pass for util, one re-read for entropy) collapses into a single pass — same output, no algorithmic change beyond the scratch-write routing. **9 new ISV slots wired with Pearls A+D**: `ISV[MOE_EXPERT_UTIL_EMA_BASE+k=118+k]` for k in 0..8 (Wiener offsets `(44+k)*3` = 132..156), plus `ISV[MOE_GATE_ENTROPY_EMA_INDEX=126]` (Wiener offset 156..159). Launcher `gpu_moe_head.rs::launch_expert_util_ema` drops `(isv_dev_ptr, isv_util_base, isv_entropy_index, alpha)` for `(scratch_dev_ptr, scratch_idx_util_base, scratch_idx_entropy)`. Wrapper `GpuDqnTrainer::launch_moe_expert_util_ema(_ema_alpha_unused: f32)` syncs the stream after launch and applies Pearls A+D in a per-slot loop via `apply_pearls(scratch_idx, isv_idx)` closure: 8 expert util pairs + 1 entropy pair. Degenerate-zero short-circuit at every slot prevents Wiener-state advancement when the kernel hasn't run (cold start before forward populates `moe_gate_softmax_buf`). `debug_assert_eq!(MOE_NUM_EXPERTS, 8)` guards against future K changes silently breaking the contiguous scratch layout. **Unit test** `sp4_moe_expert_util_ema_writes_step_obs_via_pearl_a_then_converges_pearl_d` (`#[ignore]`-gated for GPU): drives kernel with B=128 K=8 perfect-uniform gate (all entries 1/8) → analytical col_mean=1/8 per expert, entropy=ln(8)≈2.0794. Asserts each util slot ∈ ±1e-6, entropy slot ∈ ±1e-5, all non-target scratch slots remain 0; exercises Pearl A bootstrap (entropy slot returns directly + seeds x_lag) + Pearl D convergence (1000 stationary observations within 1% of analytical). Behaviour: cold-path producer with no consumer-facing change beyond the EMA mechanism — the 9 slots stay semantically identical (per-expert col_mean, gate entropy) and continue to be read by HEALTH_DIAG mirror + the adaptive λ controller `moe_lambda_eff_update` (which reads ISV[MOE_GATE_ENTROPY_EMA_INDEX]). Per `feedback_no_atomicadd.md`, `feedback_no_htod_htoh_only_mapped_pinned.md`. `cargo check -p ml --lib --tests --offline` clean (11 pre-existing warnings, no new warnings). SP4 Layer A Task A13.3 — vsn_mask_ema retrofit (Pearls A+D) (2026-05-01): retrofit the existing per-step VSN feature-group mask producer in `crates/ml/src/cuda_pipeline/vsn_mask_ema_kernel.cu`. Kernel signature changed: drops `(isv, isv_first_index, ema_alpha)` for `(scratch_buf, scratch_first_index=53)`. Single-block, 256-thread shmem-tree reduce per group; `num_groups` separate reductions sharing the same 256-float scratchpad — total smem footprint identical to the prior implementation. Per-group mean is written to `scratch_buf[scratch_first_index + g]` for g in 0..num_groups; single `__threadfence_system()` after all 6 groups write (PCIe-visible to mapped-pinned host_ptr). **6 ISV slots wired with Pearls A+D**: `ISV[VSN_MASK_GROUP_0_EMA_INDEX..+6) = ISV[105..111)`, with Wiener offsets `(53+g)*3` = 159..177 for g in 0..6. Wrapper `GpuDqnTrainer::launch_vsn_mask_ema(_ema_alpha_unused: f32)` syncs the stream after launch and applies Pearls A+D in a `for` loop over the 6 groups. `debug_assert_eq!(num_groups, 6)` guards against `SL_NUM_FEATURE_GROUPS` changes silently breaking the contiguous scratch layout. Degenerate-zero short-circuit at every slot prevents Wiener-state advancement when the kernel hasn't run (cold start before VSN forward populates `vsn_mask_buf`). **Unit test** `sp4_vsn_mask_ema_writes_step_obs_via_pearl_a_then_converges_pearl_d` (`#[ignore]`-gated for GPU): drives kernel with B=128 num_groups=6 mask `mask[b,g] = (g+1)/21` (rows sum to 1, per-group mean = (g+1)/21). Asserts each scratch slot ∈ ±1e-5, non-target slots remain 0; verifies Pearl A bootstrap on group 0 + Pearl D convergence (1000 stationary observations within 1% of analytical). Behaviour: cold-path producer with no consumer-facing change beyond the EMA mechanism — slots [105..111) stay semantically identical (per-group mean of VSN softmax mask) and continue to be read by HEALTH_DIAG mirror + VSN focus monitor. Per `feedback_no_atomicadd.md`, `feedback_no_htod_htoh_only_mapped_pinned.md`. `cargo check -p ml --lib --tests --offline` clean (11 pre-existing warnings, no new warnings). + +SP4 Layer A Task A13.4 — iqn_quantile_ema retrofit (Pearls A+D) (2026-05-01): retrofit the existing per-step IQN off-median-quantile diagnostic producer in `crates/ml/src/cuda_pipeline/iqn_quantile_ema_kernel.cu`. Kernel signature changed: drops `(isv, isv_q_p05_idx, isv_q_p25_idx, isv_q_p75_idx, isv_q_p95_idx, ema_alpha)` for `(scratch_buf, scratch_first_index=59)`. Block-dispatch unchanged: 4 blocks × 256 threads, blockIdx.x ∈ {0,1,2,3} maps to (τ_idx ∈ {0,1,3,4}, slot_offset ∈ {0,1,2,3}); the median (τ_idx=2) is intentionally skipped. Each block reduces mean |Q(s,a;τ)| over (B × TBA) entries via shmem-tree (no atomicAdd) and writes the step observation to `scratch_buf[scratch_first_index + slot_offset]` with `__threadfence_system()`. **4 ISV slots wired with Pearls A+D**: ISV[IQN_Q_P05_EMA_INDEX=99] (Wiener offset 177), ISV[IQN_Q_P25_EMA_INDEX=100] (Wiener offset 180), ISV[IQN_Q_P75_EMA_INDEX=101] (Wiener offset 183), ISV[IQN_Q_P95_EMA_INDEX=102] (Wiener offset 186). Wrapper `GpuDqnTrainer::launch_iqn_quantile_ema(save_q_online_ptr, tba, num_quantiles, _ema_alpha_unused: f32)` retains its early-return `if save_q_online_ptr == 0` guard (no IQN active in this trainer config), and now syncs the stream after launch and applies Pearls A+D in a `for` loop over the 4 SLOT_PAIRS const array. Degenerate-zero short-circuit at every slot prevents Wiener-state advancement when the kernel hasn't run (cold start before IQN forward populates `save_q_online`). **Unit test** `sp4_iqn_quantile_ema_writes_step_obs_via_pearl_a_then_converges_pearl_d` (`#[ignore]`-gated for GPU): drives kernel with B=32 Q=5 TBA=12 controlled `save_q_online` where `Q[a, b*Q+tau_idx] = (tau_idx+1) × 0.7` (analytical mean per tau = (tau_idx+1) × 0.7). Asserts each slot ∈ ±1e-5 of expected (0.7/1.4/2.8/3.5 for p05/p25/p75/p95), median (tau_idx=2 expected 2.1) intentionally absent, non-target scratch slots remain 0; verifies Pearl A bootstrap on p05 + Pearl D convergence (1000 stationary observations within 1% of 0.7). Behaviour: cold-path producer with no consumer-facing change beyond the EMA mechanism — slots [99..103) remain diagnostic-only (HEALTH_DIAG / risk-monitoring surface). Per `feedback_no_atomicadd.md`, `feedback_no_htod_htoh_only_mapped_pinned.md`. `cargo check -p ml --lib --tests --offline` clean (11 pre-existing warnings, no new warnings).