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

feat(sp4): Task A13.4 — retrofit iqn_quantile_ema with Pearls A+D

Browse Source 
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-01 02:18:55 +02:00
parent 0e9d69787d
commit 5f800fe5c8
4 changed files with 232 additions and 53 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -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::<f32>() 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(())
}

87
crates/ml/src/cuda_pipeline/iqn_quantile_ema_kernel.cu
View File

@@ -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 <cuda_runtime.h>
@@ -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();
}
}

120
crates/ml/tests/sp4_producer_unit_tests.rs
View File

@@ -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<CudaStream>) -> 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::<f32>() 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)",
);
}

2
docs/dqn-wire-up-audit.md
View File

@@ -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).