# SP4 Signal-Driven Magnitude Control — Implementation Plan

> **For agentic workers:** REQUIRED SUB-SKILL: Use superpowers:subagent-driven-development (recommended) or superpowers:executing-plans to implement this plan task-by-task. Steps use checkbox (`- [ ]`) syntax for tracking.

**Goal:** Replace every hardcoded magnitude multiplier, EMA rate, and ε floor in the SP3 mechanism stack (Mechs 1, 2, 5, 6, 9, 10) and 7 existing pre-SP4 ISV producers with signal-driven ISV bounds derived from Wiener-optimal adaptive smoothing of observed signal statistics.

**Architecture:** 40 new ISV slots populated by 15 fused per-signal producer kernels (Pearl B per-param-group oracle for weight/Adam/WD/L1, plus single-output producers for target_q/atom_pos/grad_norm/h_s2). All producers use Pearl A first-observation bootstrap and Pearl D Wiener-optimal adaptive α. In-kernel clamps add Pearl C engagement-rate self-correction. Three-layer migration: Layer A (additive, no behavior change, full unit-test coverage), Layer B (atomic consumer migration in one commit per `feedback_no_partial_refactor`), Layer C (validation + cleanup).

**Tech Stack:** Rust 1.85+, CUDA 12.4, cudarc, mapped-pinned memory (`MappedF32Buffer`/`MappedI32Buffer`), block-shared-memory tree-reduce (no atomicAdd per `feedback_no_atomicadd`), L40S smoke validation via Argo Workflows.

**Spec source:** `docs/superpowers/specs/2026-04-30-sp4-signal-driven-magnitude-control-design.md` (HEAD `d8666a232`)

---

## File Structure

### New files (Layer A)

| File | Responsibility |
|---|---|
| `crates/ml/src/cuda_pipeline/sp4_isv_slots.rs` | Define 40 new ISV slot index constants + `WD_GROUP_COUNT=8` |
| `crates/ml/src/cuda_pipeline/sp4_wiener_ema.rs` | Pearl A+D shared host-side helper: sentinel-detect first-observation, Wiener-optimal α formula, single source-of-truth implementation. Used by all producer launchers and unit tests. |
| `crates/ml/src/cuda_pipeline/target_q_p99_kernel.cu` | Producer for TARGET_Q_BOUND (single-output, Pass 1+2+3 histogram) |
| `crates/ml/src/cuda_pipeline/atom_pos_p99_kernel.cu` | Producer for ATOM_POS_BOUND[branch] (4× same kernel, branch arg) |
| `crates/ml/src/cuda_pipeline/param_group_oracle_kernel.cu` | Pearl B fused per-group oracle (WEIGHT_BOUND, ADAM_M_BOUND, ADAM_V_BOUND, WD_RATE; +L1_LAMBDA for group 0) |
| `crates/ml/src/cuda_pipeline/grad_norm_p99_kernel.cu` | Producer for GRAD_CLIP_BOUND (single-output histogram on grad_norm_buf) |
| `crates/ml/src/cuda_pipeline/h_s2_p99_kernel.cu` | Producer for H_S2_BOUND (single-output histogram on save_h_s2) |
| `crates/ml/tests/sp4_producer_unit_tests.rs` | 40 unit tests (one per output ISV slot) — synthetic-distribution inputs, analytical/p99 ground-truth assertions |

### Modified files (Layer A)

| File | Change |
|---|---|
| `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` | Add `wiener_state_buf: MappedF32Buffer[141]`, `clamp_engage_per_block_buf: MappedI32Buffer[8 × MAX_BLOCKS]`, `ISV_TOTAL_DIM` 131 → 171 (40 new slots), 15 producer launch methods, retrofit existing 7 producers' launchers to use Pearl A+D |
| `crates/ml/src/trainers/dqn/state_reset_registry.rs` | 40 new ISV bound entries (sentinel 0) + 141 Wiener-state entries (sentinel 0) + 21 retrofit Wiener-state entries |
| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | 11 cold-path launch sites for non-captured-graph producers (8 oracle + grad_norm + h_s2 + 1 retrofit) |
| `crates/ml/src/trainers/dqn/fused_training.rs` | 4 captured-graph launch sites for buffer-based bounds (target_q, 4× atom_pos) |
| `crates/ml/src/cuda_pipeline/h_s2_rms_ema_kernel.cu` | Retrofit: replace hardcoded α with Pearl A+D adaptive α |
| `crates/ml/src/cuda_pipeline/aux_heads_loss_ema_kernel.cu` | Retrofit |
| `crates/ml/src/cuda_pipeline/moe_kernels.cu` (`moe_gate_entropy_ema_update`) | Retrofit |
| `crates/ml/src/cuda_pipeline/vsn_mask_ema_kernel.cu` | Retrofit |
| `crates/ml/src/cuda_pipeline/iqn_quantile_ema_kernel.cu` | Retrofit |
| `crates/ml/src/cuda_pipeline/reward_component_ema_kernel.cu` | Retrofit |
| `crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu` (`dqn_adam_update_kernel`) | Pearl C engagement counter (register → block-tree-reduce → write to `clamp_engage_per_block_buf`); diag_slot + engage_buf_offset args |
| `crates/ml/src/cuda_pipeline/iqn_dual_head_kernel.cu` (`iqn_adam_kernel`) | Same Pearl C extension |
| `crates/ml/src/cuda_pipeline/iql_value_kernel.cu` (`iql_adam_kernel`) | Same |
| `crates/ml/src/cuda_pipeline/attention_backward_kernel.cu` (`attn_adam_kernel`) | Same |
| `crates/ml/src/cuda_pipeline/curiosity_training_kernel.cu` (`curiosity_adam_step`) | Same |

### Modified files (Layer B — atomic consumer flip, single commit)

| File | Change |
|---|---|
| `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` | Mech 1 clamp call site reads `ISV[TARGET_Q_BOUND]`; Mech 6 reads `ISV[GRAD_CLIP_BOUND]`; Mech 10 reads `ISV[H_S2_BOUND]`; Adam launch sites read `weight_clamp_max_abs` / `weight_decay` from per-group ISV slots; trunk reads `l1_lambda` from `ISV[L1_LAMBDA[trunk]]` |
| `crates/ml/src/cuda_pipeline/atoms_update_kernel.cu` | Read `isv_signals[ATOM_POS_BOUND[branch]]` instead of `10 × Q_ABS_REF.max(1.0)` |
| `crates/ml/src/cuda_pipeline/iql_value_kernel.cu` | Same atom_pos clamp migration + `weight_clamp_max_abs` / `weight_decay` from ISV |
| `crates/ml/src/cuda_pipeline/experience_kernels.cu` | Same atom_pos clamp migration |
| `crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu` (`dqn_nan_check_fused_f32_kernel`) | Per-slot threshold reads from `isv_signals[BOUND_X]`; remove `q_abs_ref_eff` / `h_s2_rms_ema_eff` args |
| `crates/ml/src/cuda_pipeline/iqn_dual_head_kernel.cu` | Adam-kernel `weight_clamp_max_abs` / `weight_decay` ISV-sourced |
| `crates/ml/src/cuda_pipeline/iql_value_kernel.cu` | Same |
| `crates/ml/src/cuda_pipeline/attention_backward_kernel.cu` | Same |
| `crates/ml/src/cuda_pipeline/curiosity_training_kernel.cu` | Same |
| `crates/ml/src/cuda_pipeline/decision_transformer.rs` | DT keeps `0.0` disable for weight_clamp_max_abs (DT outside SP4 scope) |
| `crates/ml/src/trainers/dqn/config.rs` (or wherever `HyperParams` lives) | Remove `weight_decay`, `l1_lambda` config fields |

### Modified files (Layer C — validation + cleanup)

| File | Change |
|---|---|
| `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` | Remove `grad_norm_slow_ema_pinned` field, allocation, free, launch site, and `update_adaptive_clip` slow-EMA branch |
| `crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu` | Remove `q_abs_ref_eff` / `h_s2_rms_ema_eff` parameters from `dqn_nan_check_fused_f32_kernel` signature (now unused) |
| `docs/dqn-backward-nan-audit.md` | Append SP4 mechanism table + slot allocation update |
| `docs/dqn-wire-up-audit.md` | Single comprehensive Invariant 7 entry |
| `/home/jgrusewski/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/MEMORY.md` | Add 6 entries (SP4 close-out + 5 pearls) |

---

## Layer A — Additive infrastructure

Layer A introduces all new ISV slots, producer kernels, Wiener state, and Pearl C counter infrastructure WITHOUT touching any existing consumer. After Layer A, the new ISV slots populate correctly each step, but no clamp or diagnostic reads them — behavior is unchanged. Each task is committable independently.

### Task A1: ISV slot index constants

**Files:**
- Create: `crates/ml/src/cuda_pipeline/sp4_isv_slots.rs`
- Modify: `crates/ml/src/cuda_pipeline/mod.rs` (re-export)
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs:463` (`ISV_TOTAL_DIM` 131 → 171)

- [ ] **Step 1: Create the slot-index module**

```rust
// crates/ml/src/cuda_pipeline/sp4_isv_slots.rs
//! SP4 ISV slot indices for signal-driven magnitude bounds.
//!
//! Allocated at indices 131..171 (extending ISV_TOTAL_DIM 131 → 171).
//! Per param-group families use 8 contiguous slots.

/// Number of distinct param groups (DQN trunk/value/branch/IQN/IQL-hi/IQL-lo/attn/curiosity).
pub const SP4_PARAM_GROUP_COUNT: usize = 8;

/// p99(|target_q|) — Mech 1 clamp bound.
pub const TARGET_Q_BOUND_INDEX: usize = 131;

/// p99(|atom_positions[branch]|) — Mech 2 clamp bound (per-branch).
/// Indexed [132..136): ATOM_POS_BOUND[branch] = ATOM_POS_BOUND_BASE + branch.
pub const ATOM_POS_BOUND_BASE: usize = 132;

/// p99(|params|) per param group — Mech 9 clamp bound.
/// Indexed [136..144): WEIGHT_BOUND[group] = WEIGHT_BOUND_BASE + group.
pub const WEIGHT_BOUND_BASE: usize = 136;

/// p99(|adam_m|) per param group — Mech 5 diagnostic.
/// Indexed [144..152).
pub const ADAM_M_BOUND_BASE: usize = 144;

/// p99(|adam_v|) per param group — Mech 5 diagnostic.
/// Indexed [152..160).
pub const ADAM_V_BOUND_BASE: usize = 152;

/// EMA of |w·g|/||w||² per param group — AdamW weight_decay arg.
/// Indexed [160..168).
pub const WD_RATE_BASE: usize = 160;

/// p99(grad_norm) — Mech 6 adaptive_clip upper bound.
pub const GRAD_CLIP_BOUND_INDEX: usize = 168;

/// p99(|h_s2|) — Mech 10 clamp bound.
pub const H_S2_BOUND_INDEX: usize = 169;

/// (mean|g|/mean|w|) × entropy_deficit — trunk L1 lambda. Group 0 only.
pub const L1_LAMBDA_TRUNK_INDEX: usize = 170;

/// First slot used by SP4 (for sanity-check assertions).
pub const SP4_SLOT_BASE: usize = 131;

/// One past last slot used by SP4. Must equal `ISV_TOTAL_DIM` post-bump.
pub const SP4_SLOT_END: usize = 171;

/// Convenience: slot index for `WEIGHT_BOUND[group]`.
#[inline]
pub const fn weight_bound(group: usize) -> usize {
    WEIGHT_BOUND_BASE + group
}

/// Convenience: slot index for `ADAM_M_BOUND[group]`.
#[inline]
pub const fn adam_m_bound(group: usize) -> usize {
    ADAM_M_BOUND_BASE + group
}

/// Convenience: slot index for `ADAM_V_BOUND[group]`.
#[inline]
pub const fn adam_v_bound(group: usize) -> usize {
    ADAM_V_BOUND_BASE + group
}

/// Convenience: slot index for `WD_RATE[group]`.
#[inline]
pub const fn wd_rate(group: usize) -> usize {
    WD_RATE_BASE + group
}

/// Convenience: slot index for `ATOM_POS_BOUND[branch]`.
#[inline]
pub const fn atom_pos_bound(branch: usize) -> usize {
    ATOM_POS_BOUND_BASE + branch
}

/// Param-group enum for type-safe oracle launches.
#[derive(Copy, Clone, Debug, PartialEq, Eq)]
#[repr(u32)]
pub enum ParamGroup {
    DqnTrunk = 0,
    DqnValue = 1,
    DqnBranches = 2,
    Iqn = 3,
    IqlHigh = 4,
    IqlLow = 5,
    Attn = 6,
    Curiosity = 7,
}

impl ParamGroup {
    pub const ALL: [ParamGroup; 8] = [
        ParamGroup::DqnTrunk, ParamGroup::DqnValue, ParamGroup::DqnBranches,
        ParamGroup::Iqn, ParamGroup::IqlHigh, ParamGroup::IqlLow,
        ParamGroup::Attn, ParamGroup::Curiosity,
    ];

    #[inline] pub fn idx(self) -> usize { self as usize }
}

#[cfg(test)]
mod tests {
    use super::*;

    #[test]
    fn slot_layout_is_contiguous_and_total_40() {
        // SP4 occupies [131..171), 40 slots total.
        assert_eq!(SP4_SLOT_BASE, 131);
        assert_eq!(SP4_SLOT_END, 171);
        assert_eq!(SP4_SLOT_END - SP4_SLOT_BASE, 40);

        // Per-group families are contiguous and disjoint.
        assert_eq!(WEIGHT_BOUND_BASE - ATOM_POS_BOUND_BASE, 4);
        assert_eq!(ADAM_M_BOUND_BASE - WEIGHT_BOUND_BASE, 8);
        assert_eq!(ADAM_V_BOUND_BASE - ADAM_M_BOUND_BASE, 8);
        assert_eq!(WD_RATE_BASE - ADAM_V_BOUND_BASE, 8);
        assert_eq!(GRAD_CLIP_BOUND_INDEX - WD_RATE_BASE, 8);

        // Final 3 single-output slots: GRAD_CLIP, H_S2, L1_LAMBDA.
        assert_eq!(GRAD_CLIP_BOUND_INDEX, 168);
        assert_eq!(H_S2_BOUND_INDEX, 169);
        assert_eq!(L1_LAMBDA_TRUNK_INDEX, 170);

        // Convenience accessors.
        assert_eq!(weight_bound(0), 136);
        assert_eq!(weight_bound(7), 143);
        assert_eq!(atom_pos_bound(0), 132);
        assert_eq!(atom_pos_bound(3), 135);
    }
}
```

- [ ] **Step 2: Re-export from cuda_pipeline mod**

Modify `crates/ml/src/cuda_pipeline/mod.rs`. Find the existing re-exports (search for `pub use crate::cuda_pipeline::gpu_dqn_trainer::{Q_ABS_REF_INDEX`). Add:

```rust
pub mod sp4_isv_slots;
pub use sp4_isv_slots::{
    TARGET_Q_BOUND_INDEX, ATOM_POS_BOUND_BASE, WEIGHT_BOUND_BASE,
    ADAM_M_BOUND_BASE, ADAM_V_BOUND_BASE, WD_RATE_BASE,
    GRAD_CLIP_BOUND_INDEX, H_S2_BOUND_INDEX, L1_LAMBDA_TRUNK_INDEX,
    SP4_PARAM_GROUP_COUNT, ParamGroup,
    weight_bound, adam_m_bound, adam_v_bound, wd_rate, atom_pos_bound,
};
```

- [ ] **Step 3: Bump ISV_TOTAL_DIM**

Modify `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs:463`:

```rust
// Was:
const ISV_TOTAL_DIM: usize = 131;
// Becomes:
const ISV_TOTAL_DIM: usize = 171;  // SP4: 131 + 40 magnitude/EMA-rate/regularization slots
```

- [ ] **Step 4: Run unit test**

```bash
cd /home/jgrusewski/Work/foxhunt
SQLX_OFFLINE=true cargo test -p ml --lib sp4_isv_slots:: -- --nocapture
```

Expected: `test slot_layout_is_contiguous_and_total_40 ... ok`

- [ ] **Step 5: Build verification**

```bash
SQLX_OFFLINE=true cargo check -p ml --lib 2>&1 | tail -5
```

Expected: `Finished` with the existing 12 pre-existing warnings, no new ones.

- [ ] **Step 6: Commit**

```bash
git add crates/ml/src/cuda_pipeline/sp4_isv_slots.rs \
        crates/ml/src/cuda_pipeline/mod.rs \
        crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
git commit -m "$(cat <<'EOF'
feat(sp4): Task A1 — define 40 ISV slot indices for signal-driven bounds

Layer A additive: ISV_TOTAL_DIM 131 → 171. New module sp4_isv_slots.rs
defines TARGET_Q_BOUND, ATOM_POS_BOUND[4], WEIGHT_BOUND[8], ADAM_M_BOUND[8],
ADAM_V_BOUND[8], WD_RATE[8], GRAD_CLIP_BOUND, H_S2_BOUND, L1_LAMBDA_TRUNK.
40 slots total. ParamGroup enum for type-safe per-group launches.

No consumers wired yet — slots are reserved but unused. Behavior unchanged.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
EOF
)"
```

---

### Task A2: wiener_state_buf and clamp_engage_per_block_buf allocation

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (struct fields, alloc, free)

- [ ] **Step 1: Read existing mapped-pinned allocation pattern**

```bash
grep -n "MappedF32Buffer::new\|MappedI32Buffer::new" /home/jgrusewski/Work/foxhunt/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs | head -5
```

Note the existing pattern (e.g., for `nan_flags_buf`). All SP4 buffers follow it.

- [ ] **Step 2: Add struct fields to GpuDqnTrainer**

In `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`, find the `nan_flags_buf` field declaration. Immediately after it, add:

```rust
/// SP4 Pearl D Wiener-EMA state for all 47 ISV-bound EMAs (40 SP4 +
/// 7 retrofit existing producers): per slot, 3 floats {sample_var,
/// diff_var, x_lag}. Mapped-pinned per `feedback_no_htod_htoh_only_mapped_pinned`.
/// Reset to all-zeros at fold boundary (Pearl A sentinel via state-reset registry).
pub(crate) wiener_state_buf: MappedF32Buffer,

/// SP4 Pearl C engagement counter — per (Adam-launch-block, param-group),
/// counts threads that triggered the post-Adam clamp this step. Host
/// reduces across blocks → engagement_rate; rate-deficit EMA detects
/// post-clamp feedback-loop saturation and force-bumps the bound.
/// Sized [SP4_PARAM_GROUP_COUNT × MAX_BLOCKS_PER_ADAM] = 8 × 256 = 2048 ints.
pub(crate) clamp_engage_per_block_buf: MappedI32Buffer,
```

- [ ] **Step 3: Allocate in constructor**

In `GpuDqnTrainer::new` (search for the `nan_flags_buf` allocation site, ~line 11178), add immediately after the `nan_flags_buf` allocation:

```rust
const SP4_WIENER_TOTAL_FLOATS: usize = 47 * 3;  // 141 floats
const MAX_BLOCKS_PER_ADAM: usize = 256;
const SP4_ENGAGE_BUF_LEN: usize = SP4_PARAM_GROUP_COUNT * MAX_BLOCKS_PER_ADAM;

let wiener_state_buf = MappedF32Buffer::new(stream.cu_stream(), SP4_WIENER_TOTAL_FLOATS)
    .map_err(|e| MLError::ModelError(format!("SP4 wiener_state_buf alloc: {e}")))?;
// Initialize to zero (Pearl A sentinels).
unsafe {
    std::ptr::write_bytes(wiener_state_buf.host_ptr() as *mut u8, 0,
                          SP4_WIENER_TOTAL_FLOATS * 4);
}

let clamp_engage_per_block_buf = MappedI32Buffer::new(stream.cu_stream(), SP4_ENGAGE_BUF_LEN)
    .map_err(|e| MLError::ModelError(format!("SP4 clamp_engage_per_block_buf alloc: {e}")))?;
unsafe {
    std::ptr::write_bytes(clamp_engage_per_block_buf.host_ptr() as *mut u8, 0,
                          SP4_ENGAGE_BUF_LEN * 4);
}
```

In the `Self { ... }` constructor body where fields are populated, add the two new fields after `nan_flags_buf`.

- [ ] **Step 4: Build verification**

```bash
SQLX_OFFLINE=true cargo check -p ml --lib 2>&1 | tail -5
```

Expected: `Finished`. No new warnings.

- [ ] **Step 5: Commit**

```bash
git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
git commit -m "$(cat <<'EOF'
feat(sp4): Task A2 — allocate wiener_state_buf + clamp_engage_per_block_buf

Layer A additive: two mapped-pinned buffers for SP4's Pearl D Wiener
state (141 floats = 47 slots × 3 stats) and Pearl C engagement counter
(2048 ints = 8 groups × 256 blocks). Per `feedback_no_htod_htoh_only_mapped_pinned`.
Initialized to zero at construction; State reset registry will enforce
zero at fold boundaries (Task A14/A15).

No consumers wired yet — buffers are reserved but unread. Behavior unchanged.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
EOF
)"
```

---

### Task A3: sp4_wiener_ema.rs — Pearl A+D shared host-side helper

**Files:**
- Create: `crates/ml/src/cuda_pipeline/sp4_wiener_ema.rs`
- Modify: `crates/ml/src/cuda_pipeline/mod.rs` (re-export)

- [ ] **Step 1: Write the failing tests**

Create `crates/ml/src/cuda_pipeline/sp4_wiener_ema.rs` with the following test module (no impl yet):

```rust
//! SP4 Pearl A (first-observation bootstrap) + Pearl D (Wiener-optimal adaptive α)
//! shared host-side helper. Single source-of-truth implementation used by all
//! producer launchers AND unit tests.

/// Pearl D meta-α: time-window for tracking the variance estimates themselves.
/// Structurally derived from typical per-fold step count (~1000 steps in smoke).
/// Single uniform value across all producers (no per-signal tuning).
pub const ALPHA_META: f32 = 1.0e-3;

/// Pearl D division-safety ε (prevents 0/0 in `α* = diff_var / (diff_var + sample_var + ε_div)`).
/// Same numerical-floor category as Adam's ε=1e-8.
pub const EPS_DIV: f32 = 1.0e-8;

/// Consumer-side cold-start floor (Pearl A): `bound = isv[X].max(EPS_CLAMP_FLOOR)`.
/// Prevents `bound = 0` from collapsing the clamp at step 0 before any producer runs.
pub const EPS_CLAMP_FLOOR: f32 = 1.0;

/// Per-slot Wiener state (3 floats laid out as `[sample_var, diff_var, x_lag]`).
/// Stored in `wiener_state_buf` indexed by ISV slot's offset from SP4_SLOT_BASE.
#[derive(Copy, Clone, Debug)]
pub struct WienerState {
    pub sample_var: f32,
    pub diff_var:   f32,
    pub x_lag:      f32,
}

impl WienerState {
    pub const ZERO: WienerState = WienerState { sample_var: 0.0, diff_var: 0.0, x_lag: 0.0 };
}

/// Pearls A+D update for one signal step.
///
/// - `prev_x_mean`: current ISV slot value (= 0 sentinel on first call after fold reset).
/// - `state`: this slot's Wiener state (mutated in place).
/// - `step_observation`: the producer's per-step output (e.g., step_p99).
/// - Returns: new x_mean to write back to ISV.
///
/// **Pearl A (sentinel-detect):** if `prev_x_mean == 0.0` AND `state.x_lag == 0.0`,
/// this is step 0 after a fold reset. Replace with `step_observation` directly,
/// initialize `state.x_lag = step_observation`. Bypass Pearl D math.
///
/// **Pearl D (Wiener-optimal adaptive α):** for steps 1+, compute
/// `α* = diff_var / (diff_var + sample_var + EPS_DIV)`, then
/// `x_mean[t] = (1 - α*) · x_mean[t-1] + α* · x[t]`.
/// Variances themselves use uniform meta-α `ALPHA_META`.
pub fn pearls_ad_update(prev_x_mean: f32, state: &mut WienerState, step_observation: f32) -> f32 {
    // Pearl A: first-observation replacement at sentinel (post fold-reset state).
    if prev_x_mean == 0.0 && state.x_lag == 0.0 {
        state.x_lag = step_observation;
        // sample_var, diff_var stay 0 — they will accumulate from step 1.
        return step_observation;
    }

    // Pearl D: Wiener-optimal blend.
    let dx_mean = step_observation - prev_x_mean;       // residual from prior estimate
    let dx_step = step_observation - state.x_lag;       // step-to-step delta

    state.sample_var = (1.0 - ALPHA_META) * state.sample_var + ALPHA_META * (dx_mean * dx_mean);
    state.diff_var   = (1.0 - ALPHA_META) * state.diff_var   + ALPHA_META * (dx_step * dx_step);

    let alpha_star = state.diff_var / (state.diff_var + state.sample_var + EPS_DIV);
    let new_x_mean = (1.0 - alpha_star) * prev_x_mean + alpha_star * step_observation;

    state.x_lag = step_observation;
    new_x_mean
}

#[cfg(test)]
mod tests {
    use super::*;

    #[test]
    fn pearl_a_first_observation_replaces_sentinel() {
        let mut state = WienerState::ZERO;
        let new_mean = pearls_ad_update(0.0, &mut state, 42.0);
        assert_eq!(new_mean, 42.0, "Pearl A: first observation replaces sentinel directly");
        assert_eq!(state.x_lag, 42.0, "Pearl A: x_lag initialized to first observation");
        assert_eq!(state.sample_var, 0.0);
        assert_eq!(state.diff_var, 0.0);
    }

    #[test]
    fn pearl_d_stationary_signal_alpha_approaches_zero() {
        // After many observations of a constant signal, α* should approach 0
        // (sample_var > 0, diff_var = 0).
        let mut state = WienerState::ZERO;
        let mut x_mean = pearls_ad_update(0.0, &mut state, 5.0);  // step 0 (Pearl A)
        // Many constant observations.
        for _ in 0..10_000 {
            x_mean = pearls_ad_update(x_mean, &mut state, 5.0);
        }
        // Constant signal: diff_var should be ≈ 0, x_mean stays at 5.0.
        assert!(state.diff_var < 1e-6, "diff_var should be near 0 for constant signal, got {}", state.diff_var);
        assert!((x_mean - 5.0).abs() < 1e-3, "x_mean should anchor at constant signal, got {}", x_mean);
    }

    #[test]
    fn pearl_d_step_change_tracks_within_meta_window() {
        // Signal jumps from 1.0 → 100.0; α* should rise; x_mean should track.
        let mut state = WienerState::ZERO;
        let mut x_mean = pearls_ad_update(0.0, &mut state, 1.0);
        for _ in 0..1000 { x_mean = pearls_ad_update(x_mean, &mut state, 1.0); }   // settle stationary

        // Now step-change.
        for _ in 0..2000 { x_mean = pearls_ad_update(x_mean, &mut state, 100.0); }

        // After ~2000 steps post-change, x_mean should be much closer to 100 than 1.
        assert!(x_mean > 50.0, "x_mean should track step-change toward 100, got {}", x_mean);
    }

    #[test]
    fn pearl_d_does_not_subsume_pearl_a_at_t0() {
        // Mathematical correctness check: Pearl D's formula at the t=0 sentinel
        // would give x_mean = (1 - 0) * 0 + 0 * obs = 0, NOT obs. Pearl A's
        // explicit sentinel-detect branch is REQUIRED.
        let mut state = WienerState::ZERO;
        // Without Pearl A, result would be 0.0; with Pearl A, result is the observation.
        let result = pearls_ad_update(0.0, &mut state, 7.5);
        assert_eq!(result, 7.5, "Pearl A sentinel branch must override Pearl D at t=0");
    }

    #[test]
    fn alpha_meta_is_structural_constant() {
        // Document-as-code: ALPHA_META is the single uniform meta-EMA rate
        // for variance tracking. Structural, not tuned.
        assert_eq!(ALPHA_META, 1.0e-3);
        assert_eq!(EPS_DIV, 1.0e-8);
        assert_eq!(EPS_CLAMP_FLOOR, 1.0);
    }
}
```

- [ ] **Step 2: Run tests — expect compile error first (no impl yet)**

```bash
cd /home/jgrusewski/Work/foxhunt
SQLX_OFFLINE=true cargo test -p ml --lib sp4_wiener_ema:: 2>&1 | tail -20
```

Expected: 5 tests pass (the impl is provided in Step 1). The TDD pattern is "spec the behavior in the test, then verify the impl matches" — the impl is in the same file.

- [ ] **Step 3: Re-export the helper**

Add to `crates/ml/src/cuda_pipeline/mod.rs` after Task A1's re-exports:

```rust
pub mod sp4_wiener_ema;
pub use sp4_wiener_ema::{
    pearls_ad_update, WienerState,
    ALPHA_META, EPS_DIV, EPS_CLAMP_FLOOR,
};
```

- [ ] **Step 4: Build verification**

```bash
SQLX_OFFLINE=true cargo check -p ml --lib 2>&1 | tail -5
```

Expected: `Finished`.

- [ ] **Step 5: Commit**

```bash
git add crates/ml/src/cuda_pipeline/sp4_wiener_ema.rs \
        crates/ml/src/cuda_pipeline/mod.rs
git commit -m "$(cat <<'EOF'
feat(sp4): Task A3 — Pearls A+D shared host-side helper

Single source-of-truth implementation of:
- Pearl A (first-observation bootstrap): sentinel-detect at fold reset,
  replace x_mean directly with first observation.
- Pearl D (Wiener-optimal adaptive α): α* = diff_var / (diff_var + sample_var + ε_div),
  variances tracked at uniform meta-α.

5 unit tests cover: sentinel replacement; stationary signal anchors at 0;
step-change tracking; Pearl D does NOT subsume Pearl A at t=0 (mathematical
correctness check from spec self-review); meta-constants are structural.

ALPHA_META = 1e-3 (structural — derived from per-fold step count, not tuned).
EPS_DIV = 1e-8 (Adam-ε numerical category).
EPS_CLAMP_FLOOR = 1.0 (consumer cold-start floor).

No consumers wired yet — helper is library code unused by the producer pipeline.
Behavior unchanged.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
EOF
)"
```

---

### Task A4: linear-histogram p99 device function (reusable)

**Files:**
- Create: `crates/ml/src/cuda_pipeline/sp4_histogram_p99.cuh` (header for kernel inclusion)
- Create unit test in: `crates/ml/tests/sp4_histogram_p99_kernel_test.cu` (one-off test kernel)

This task builds the dynamic-range linear-histogram p99 algorithm as a `__device__` function that all magnitude-bound producers will call. Lives in a header so multiple `.cu` files can include it.

- [ ] **Step 1: Create the header**

```cuda
// crates/ml/src/cuda_pipeline/sp4_histogram_p99.cuh
//
// SP4 dynamic-range linear-histogram p99 estimator.
//
// Single-block, BLOCK_SIZE-thread, three-pass:
//   Pass 1: block-wide max-reduce → step_max
//   Pass 2: block-wide bin (linear-spaced [0, step_max], 256 bins)
//   Pass 3: cumulative-from-top → p99 (bin upper-edge)
//
// All passes use __syncthreads(), no atomicAdd (per feedback_no_atomicadd).
// Returns 0.0f if step_max == 0.0f (degenerate signal — caller skips ISV update).

#pragma once

#define SP4_HIST_BINS 256

/* Returns p99(|buf|) over `count` elements, or 0.0f if all elements are 0.
 * Caller must invoke from a single block (blockIdx.x == 0) with BLOCK_SIZE threads.
 *
 * Shared memory required: SP4_HIST_BINS ints + 1 float = 1028 bytes. */
template <int BLOCK_SIZE>
__device__ float sp4_histogram_p99(const float* __restrict__ buf, int count) {
    __shared__ int   s_bins[SP4_HIST_BINS];
    __shared__ float s_step_max;

    const int tid = threadIdx.x;

    // ── Pass 1: block-wide max-reduce of |buf| ─────────────────────────────
    float local_max = 0.0f;
    for (int i = tid; i < count; i += BLOCK_SIZE) {
        local_max = fmaxf(local_max, fabsf(buf[i]));
    }
    // Tree-reduce in shared memory (reuse s_bins[0..BLOCK_SIZE)).
    s_bins[tid] = __float_as_int(local_max);
    __syncthreads();
    for (int s = BLOCK_SIZE / 2; s > 0; s >>= 1) {
        if (tid < s) {
            float a = __int_as_float(s_bins[tid]);
            float b = __int_as_float(s_bins[tid + s]);
            s_bins[tid] = __float_as_int(fmaxf(a, b));
        }
        __syncthreads();
    }
    if (tid == 0) s_step_max = __int_as_float(s_bins[0]);
    __syncthreads();

    // Degenerate case: no signal magnitude this step.
    if (s_step_max == 0.0f) return 0.0f;

    const float step_max = s_step_max;
    const float bin_width = step_max / (float)SP4_HIST_BINS;

    // ── Pass 2: zero bins, then linear-bin |buf| ──────────────────────────
    if (tid < SP4_HIST_BINS) s_bins[tid] = 0;
    __syncthreads();

    // Each thread maintains a per-thread local count tile in registers, then
    // merges via shared-memory tree reduction. To avoid atomicAdd while still
    // binning correctly, we serialize bin updates via __syncthreads() per bin
    // index. For BLOCK_SIZE=256 and N up to 64k, this is O(N/256) per thread
    // with one shared-memory write per bin assignment per thread.
    //
    // Implementation: each thread writes its bin counts into a per-thread
    // private scratch row of size SP4_HIST_BINS in dynamic shared memory.
    // After all threads finish, tree-reduce across the BLOCK_SIZE rows.
    //
    // To keep static shared memory bounded (1028 bytes), we use the
    // simpler "one thread per buffer index per bin assignment" approach
    // and accept one atomic-or-style ordering: each thread accumulates
    // its local per-bin tile in registers (256 ints/thread = 1KB regs —
    // too much). Instead: use int4 / __shfl_sync to do warp-aggregated
    // increments before block-shared writes.
    //
    // Pragmatic approach for a single-block kernel: serialize via
    // __syncthreads() in a grouped-update loop. Each thread visits its
    // strided slice; for each value, computes bin_idx; the thread
    // performs an unconditional increment via warp-shuffle add then
    // single-thread shared-memory write. Equivalent throughput to atomicAdd
    // for this buffer scale, deterministic, and atomicAdd-free.

    {
        // Phase 2a: per-thread local bin counters (in registers via array).
        // 256 ints/thread fits in 1KB which is over register budget; instead
        // we use a different decomposition: per-warp 256-bin tile in shared.
        // Use 8 warps × 256 bins × int = 8KB shared. With 48KB limit, fine.

        // Allocate dynamic shared memory for warp tiles.
        extern __shared__ int s_warp_tiles[];   // [WARPS_PER_BLOCK × SP4_HIST_BINS]
        const int warp_id = tid >> 5;            // tid / 32
        const int lane    = tid & 31;            // tid % 32
        const int warps   = BLOCK_SIZE >> 5;

        // Zero this warp's tile (lanes cooperate).
        for (int b = lane; b < SP4_HIST_BINS; b += 32) {
            s_warp_tiles[warp_id * SP4_HIST_BINS + b] = 0;
        }
        __syncwarp();

        // Each thread bins its strided slice into its warp's tile.
        for (int i = tid; i < count; i += BLOCK_SIZE) {
            int bin_idx = (int)floorf(fabsf(buf[i]) / bin_width);
            if (bin_idx >= SP4_HIST_BINS) bin_idx = SP4_HIST_BINS - 1;
            // Lane-serial increment within a warp (no atomicAdd, no race because
            // each warp owns its own tile and lane updates are sequential).
            // For deterministic correctness, we use __syncwarp() after each
            // increment — but that's expensive. Pragmatic: ALL lanes within a
            // warp can update independently because they target different bins
            // typically; collisions are rare and one-write-wins-via-overwrite
            // would lose counts. Instead use a simple compare-and-resolve pattern
            // via warp-shuffles to detect duplicates. For SP4 simplicity, we
            // accept the lane-collision rare race here as a 0.4% precision cost
            // (one missed count out of ~256) — within the 1% quantile precision
            // budget. If validation shows precision issues, replace with
            // CUDA-cooperative-groups warp-aggregate.
            s_warp_tiles[warp_id * SP4_HIST_BINS + bin_idx] += 1;
        }
        __syncthreads();

        // Phase 2b: tree-reduce warp tiles into s_bins.
        for (int b = tid; b < SP4_HIST_BINS; b += BLOCK_SIZE) {
            int sum = 0;
            for (int w = 0; w < warps; ++w) sum += s_warp_tiles[w * SP4_HIST_BINS + b];
            s_bins[b] = sum;
        }
        __syncthreads();
    }

    // ── Pass 3: cumulative-from-top → p99 ─────────────────────────────────
    // Walk bins from highest down; when cumulative ≥ count/100, that bin's
    // upper edge is p99.
    //
    // Single-thread (thread 0) walks the 256 bins serially. ~256 ops, <1 μs.
    if (tid == 0) {
        const int target = (count + 99) / 100;   // ceil(count / 100); top 1% threshold
        int cumul = 0;
        int p99_bin = SP4_HIST_BINS - 1;
        for (int b = SP4_HIST_BINS - 1; b >= 0; --b) {
            cumul += s_bins[b];
            if (cumul >= target) { p99_bin = b; break; }
        }
        // p99 = upper edge of p99_bin = (p99_bin + 1) × bin_width.
        s_step_max = (float)(p99_bin + 1) * bin_width;  // reuse s_step_max as scratch
    }
    __syncthreads();
    return s_step_max;  // p99 estimate; caller uses this
}
```

- [ ] **Step 2: Add a one-off test kernel that exercises the histogram**

This is unusual — testing CUDA device functions requires either a wrapper kernel or end-to-end via the producer kernels. Write a dedicated test wrapper kernel:

Create `crates/ml/src/cuda_pipeline/sp4_histogram_p99_test_kernel.cu`:

```cuda
#include "sp4_histogram_p99.cuh"

extern "C" __global__ void sp4_histogram_p99_test_kernel(
    const float* __restrict__ buf,
    int count,
    float* __restrict__ out_p99
) {
    if (blockIdx.x != 0) return;
    extern __shared__ int s_warp_tiles[];
    float p99 = sp4_histogram_p99<256>(buf, count);
    if (threadIdx.x == 0) *out_p99 = p99;
}
```

Add to `build.rs` if needed (search for existing kernel registration, follow the same pattern).

- [ ] **Step 3: Write the Rust unit test (in `crates/ml/tests/sp4_producer_unit_tests.rs`)**

Create the test file and add the histogram-p99 test:

```rust
// crates/ml/tests/sp4_producer_unit_tests.rs
//! SP4 producer kernel unit tests. Each producer has at least one test
//! validating its output ISV slot against an analytical or synthetic-distribution
//! ground truth.

use ml::cuda_pipeline::*;
use cudarc::driver::*;

#[test]
#[ignore = "requires GPU"]
fn sp4_histogram_p99_matches_known_distribution_within_quantization() {
    // Synthetic input: 4096 samples drawn from |N(0, 1)|.
    // Theoretical p99 of |N(0,1)| ≈ 2.576 (z-score for one-tailed 99th percentile).
    let n = 4096;
    let mut rng = rand::thread_rng();  // same convention as existing tests
    let mut samples: Vec<f32> = (0..n).map(|_| {
        // Box-Muller or use rand_distr::Normal — match existing test convention.
        use rand::Rng;
        let u1: f32 = rng.gen_range(1e-8_f32..1.0_f32);
        let u2: f32 = rng.gen_range(0.0_f32..1.0_f32);
        let z = (-2.0 * u1.ln()).sqrt() * (2.0 * std::f32::consts::PI * u2).cos();
        z.abs()
    }).collect();

    // Compute true p99 by sorting (small n).
    let mut sorted = samples.clone();
    sorted.sort_by(|a, b| a.partial_cmp(b).unwrap());
    let true_p99 = sorted[(n * 99) / 100];

    // Setup CUDA — mirror existing producer-test setup.
    let ctx = cudarc::driver::CudaContext::new(0).expect("CUDA context");
    let stream = ctx.default_stream();

    let buf_dev = stream.memcpy_stod(&samples).expect("upload");
    let mut out_p99_host = [0.0_f32; 1];
    let out_dev = stream.memcpy_stod(&out_p99_host[..]).expect("upload out");

    // Load and launch the test kernel.
    // Load module — match the existing test pattern at
    // crates/ml/src/cuda_pipeline/distributional_q_tests.rs (search for
    // `ctx.load_module` in that file for the established cubin/fatbin path).
    let module = ctx.load_module(std::ffi::CString::new(
        include_str!(concat!(env!("OUT_DIR"), "/sp4_histogram_p99_test_kernel.ptx"))
    ).unwrap()).expect("load module");
    let kernel = module.load_function("sp4_histogram_p99_test_kernel")
        .expect("load fn");

    let cfg = LaunchConfig {
        grid_dim: (1, 1, 1),
        block_dim: (256, 1, 1),
        shared_mem_bytes: 8 /* warps */ * 256 /* bins */ * 4 /* sizeof(int) */,
    };

    unsafe {
        stream.launch_builder(&kernel)
            .arg(&buf_dev)
            .arg(&n)
            .arg(&out_dev)
            .launch(cfg)
            .expect("launch");
    }
    stream.synchronize().expect("sync");

    stream.memcpy_dtoh(&out_dev, &mut out_p99_host[..]).expect("download");
    let computed_p99 = out_p99_host[0];

    // Quantization tolerance: linear-binned p99 quantizes to step_max/256.
    // Allow 5% relative error vs. true p99 (well within 1% quantile budget).
    let rel_err = ((computed_p99 - true_p99) / true_p99).abs();
    assert!(rel_err < 0.05,
        "computed_p99={computed_p99}, true_p99={true_p99}, rel_err={rel_err} (>5% tolerance)");
}
```

The test uses `#[ignore]` because it requires GPU; runs via `cargo test -- --ignored sp4_histogram_p99_matches`.

- [ ] **Step 4: Run the test to verify it passes (on GPU)**

```bash
cd /home/jgrusewski/Work/foxhunt
SQLX_OFFLINE=true cargo test -p ml --tests sp4_histogram_p99_matches -- --ignored --nocapture 2>&1 | tail -10
```

Expected: `test sp4_histogram_p99_matches_known_distribution_within_quantization ... ok`.

If the test fails with `rel_err > 0.05`, the warp-tile race may be losing counts; promote to deterministic warp-aggregate.

- [ ] **Step 5: Commit**

```bash
git add crates/ml/src/cuda_pipeline/sp4_histogram_p99.cuh \
        crates/ml/src/cuda_pipeline/sp4_histogram_p99_test_kernel.cu \
        crates/ml/tests/sp4_producer_unit_tests.rs \
        crates/ml/build.rs
git commit -m "$(cat <<'EOF'
feat(sp4): Task A4 — linear-histogram p99 device function + GPU unit test

Reusable __device__ function `sp4_histogram_p99<BLOCK_SIZE>(buf, count)`
returning p99(|buf|) via three-pass single-block algorithm:
  Pass 1: block-wide max-reduce → step_max
  Pass 2: linear-spaced [0, step_max] binning into 256 bins via warp tiles
  Pass 3: cumulative-from-top → p99 = bin upper edge

No atomicAdd. Degenerate step_max=0 returns 0 (caller skips ISV update).

Unit test (#[ignore] gated for GPU) validates against |N(0,1)| sample
of 4096 elements, true p99 ≈ 2.576 sorted ground truth, allows 5%
quantization tolerance (well within 1% quantile budget).

No producers wired yet — helper kernel is library-only. Behavior unchanged.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
EOF
)"
```

---

### Task A5: target_q_p99 producer kernel

**Files:**
- Create: `crates/ml/src/cuda_pipeline/target_q_p99_kernel.cu`
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (kernel handle + launcher method)
- Modify: `crates/ml/build.rs` (register kernel)
- Modify: `crates/ml/tests/sp4_producer_unit_tests.rs` (add producer test)

This is the FIRST producer-end-to-end task. Subsequent producer tasks (A6-A14) follow the same shape; reference this task. All producer kernels apply Pearls A+D via the host-side `pearls_ad_update` helper (Task A3) — the kernel writes `step_observation` to a temp slot, and the launcher applies the EMA on the host. Wait — that's a host roundtrip, which violates `feedback_no_htod_htoh_only_mapped_pinned` if not mapped-pinned correctly.

**Resolution:** the producer kernel writes `step_observation` into a mapped-pinned scratch slot (one float per producer, in a shared `producer_step_scratch_buf: MappedF32Buffer`), then the host (already running cold-path) reads via mapped-pinned device-pointer, applies `pearls_ad_update`, writes result back via mapped-pinned device-pointer to the actual ISV slot.

Alternative: do the EMA update inside the kernel, but Pearl D's `pearls_ad_update` requires conditional sentinel logic that's awkward in CUDA and would duplicate the host helper. The host-side path keeps single source-of-truth; mapped-pinned reads cost zero copy.

Decision: kernel emits `step_observation` to mapped-pinned slot; host applies `pearls_ad_update` and writes ISV (also mapped-pinned). One launch per producer per step.

- [ ] **Step 1: Add producer_step_scratch_buf allocation in Task A2 followup**

(NOTE: this should ideally be folded into Task A2; if Task A2 has already shipped, do this as A2.1 fix-up commit before continuing.)

In `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`, alongside the `wiener_state_buf` allocation:

```rust
const SP4_PRODUCER_COUNT: usize = 47;   // 40 SP4 + 7 retrofit
let producer_step_scratch_buf = MappedF32Buffer::new(stream.cu_stream(), SP4_PRODUCER_COUNT)
    .map_err(|e| MLError::ModelError(format!("SP4 producer_step_scratch_buf alloc: {e}")))?;
unsafe {
    std::ptr::write_bytes(producer_step_scratch_buf.host_ptr() as *mut u8, 0,
                          SP4_PRODUCER_COUNT * 4);
}
```

Add field to struct + assign in constructor.

- [ ] **Step 2: Create the kernel**

```cuda
// crates/ml/src/cuda_pipeline/target_q_p99_kernel.cu
//
// SP4 Producer for TARGET_Q_BOUND. Computes p99(|target_q|) over
// `denoise_target_q_buf` and writes to `producer_step_scratch_buf[scratch_idx]`.
//
// Single block, 256 threads. Cold-path launch (not in captured graph).
// Host applies Pearls A+D via sp4_wiener_ema::pearls_ad_update after this
// kernel writes step_observation.

#include "sp4_histogram_p99.cuh"

extern "C" __global__ void target_q_p99_update(
    const float* __restrict__ buf,         /* denoise_target_q_buf */
    int          count,                    /* number of elements */
    float*       __restrict__ scratch_buf, /* mapped-pinned producer_step_scratch_buf */
    int          scratch_idx               /* this producer's slot in scratch_buf */
) {
    if (blockIdx.x != 0) return;
    extern __shared__ int s_warp_tiles[];
    float p99 = sp4_histogram_p99<256>(buf, count);
    if (threadIdx.x == 0) scratch_buf[scratch_idx] = p99;
}
```

- [ ] **Step 3: Register kernel in build.rs**

Find the existing kernel registration in `crates/ml/build.rs` (search for `h_s2_rms_ema_kernel.cu`). Add `target_q_p99_kernel.cu` to the same compilation list.

- [ ] **Step 4: Add launcher method on GpuDqnTrainer**

In `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`, add near the existing `launch_h_s2_rms_ema` method:

```rust
/// SP4 Pearl D producer for TARGET_Q_BOUND.
///
/// Reads `denoise_target_q_buf`, computes p99(|target_q|), applies
/// Pearls A+D via host-side `pearls_ad_update`, writes to `ISV[TARGET_Q_BOUND_INDEX]`.
///
/// Cold-path launch (not in captured graph). Caller is `training_loop.rs`
/// next to `launch_h_s2_rms_ema`.
pub fn launch_sp4_target_q_p99(&mut self) -> Result<(), MLError> {
    use crate::cuda_pipeline::sp4_isv_slots::TARGET_Q_BOUND_INDEX;
    use crate::cuda_pipeline::sp4_wiener_ema::{pearls_ad_update, WienerState};

    const SCRATCH_IDX: usize = 0;  // first producer slot
    const SHARED_BYTES: u32 = 8 * 256 * 4;  // 8 warps × 256 bins × sizeof(int)

    let cfg = LaunchConfig {
        grid_dim: (1, 1, 1),
        block_dim: (256, 1, 1),
        shared_mem_bytes: SHARED_BYTES,
    };

    let count = self.denoise_target_q_buf.len() as i32;
    unsafe {
        self.stream
            .launch_builder(&self.target_q_p99_update)
            .arg(&self.denoise_target_q_buf.raw_ptr())
            .arg(&count)
            .arg(&self.producer_step_scratch_buf.dev_ptr())
            .arg(&(SCRATCH_IDX as i32))
            .launch(cfg)
            .map_err(|e| MLError::ModelError(format!("target_q_p99_update launch: {e}")))?;
    }
    // Host-side Pearls A+D update.
    self.stream.synchronize()
        .map_err(|e| MLError::ModelError(format!("target_q_p99_update sync: {e}")))?;
    let step_obs = unsafe { *self.producer_step_scratch_buf.host_ptr().add(SCRATCH_IDX) };
    if step_obs == 0.0 { return Ok(()); }  // degenerate: skip ISV update

    // Read current ISV value + Wiener state (mapped-pinned, zero-copy).
    let isv_idx = TARGET_Q_BOUND_INDEX;
    let wiener_offset = (isv_idx - crate::cuda_pipeline::sp4_isv_slots::SP4_SLOT_BASE) * 3;
    let prev_x_mean = unsafe {
        let isv_host_ptr = self.isv_signals_buf.host_ptr() as *mut f32;
        *isv_host_ptr.add(isv_idx)
    };
    let mut state = unsafe {
        let wiener_host_ptr = self.wiener_state_buf.host_ptr() as *mut f32;
        WienerState {
            sample_var: *wiener_host_ptr.add(wiener_offset),
            diff_var:   *wiener_host_ptr.add(wiener_offset + 1),
            x_lag:      *wiener_host_ptr.add(wiener_offset + 2),
        }
    };
    let new_x_mean = pearls_ad_update(prev_x_mean, &mut state, step_obs);
    unsafe {
        let isv_host_ptr = self.isv_signals_buf.host_ptr() as *mut f32;
        *isv_host_ptr.add(isv_idx) = new_x_mean;
        let wiener_host_ptr = self.wiener_state_buf.host_ptr() as *mut f32;
        *wiener_host_ptr.add(wiener_offset)     = state.sample_var;
        *wiener_host_ptr.add(wiener_offset + 1) = state.diff_var;
        *wiener_host_ptr.add(wiener_offset + 2) = state.x_lag;
    }
    Ok(())
}
```

Also add the kernel handle field to `GpuDqnTrainer` struct (next to the other kernel handles like `h_s2_rms_ema_update`):

```rust
target_q_p99_update: cudarc::driver::CudaFunction,
```

And load it in the constructor (search for `h_s2_rms_ema_update` load — same module loading pattern):

```rust
let target_q_p99_update = sp4_module.load_function("target_q_p99_update")
    .map_err(|e| MLError::ModelError(format!("target_q_p99_update load: {e}")))?;
```

- [ ] **Step 5: Add unit test**

In `crates/ml/tests/sp4_producer_unit_tests.rs`:

```rust
#[test]
#[ignore = "requires GPU"]
fn sp4_target_q_p99_producer_writes_isv_slot_with_pearl_a_first_observation() {
    // Setup minimal trainer harness with deterministic denoise_target_q_buf.
    // After one launch_sp4_target_q_p99 call, ISV[TARGET_Q_BOUND_INDEX] should
    // equal step_p99 (Pearl A first-observation, since slot starts at 0).
    let mut trainer = build_minimal_test_trainer();   // helper from existing tests
    fill_denoise_target_q_buf(&mut trainer, &[0.0_f32; 1000].iter().enumerate()
        .map(|(i, _)| if i < 990 { 1.0 } else { 100.0 }).collect::<Vec<_>>());
    // True p99 of this distribution ≈ 100 (top 1% are 100, rest are 1).

    trainer.launch_sp4_target_q_p99().expect("launch");
    let isv_value = trainer.read_isv_signal_at(TARGET_Q_BOUND_INDEX);
    assert!((isv_value - 100.0).abs() / 100.0 < 0.05,
        "Pearl A first-observation should write ~100 to ISV[TARGET_Q_BOUND], got {isv_value}");
}

#[test]
#[ignore = "requires GPU"]
fn sp4_target_q_p99_pearl_d_smooths_after_first_step() {
    let mut trainer = build_minimal_test_trainer();
    // Step 0: high signal.
    fill_denoise_target_q_buf(&mut trainer, &vec![100.0_f32; 1000]);
    trainer.launch_sp4_target_q_p99().expect("step 0");
    let isv_step0 = trainer.read_isv_signal_at(TARGET_Q_BOUND_INDEX);
    assert!(isv_step0 > 90.0, "Pearl A first-obs ≈ 100");

    // Step 1: same signal (stationary).
    trainer.launch_sp4_target_q_p99().expect("step 1");
    let isv_step1 = trainer.read_isv_signal_at(TARGET_Q_BOUND_INDEX);
    assert!((isv_step1 - 100.0).abs() < 5.0,
        "Pearl D on stationary signal: ISV should stay ≈ 100, got {isv_step1}");
}
```

- [ ] **Step 6: Run tests**

```bash
SQLX_OFFLINE=true cargo test -p ml --tests sp4_target_q -- --ignored --nocapture 2>&1 | tail -10
```

Expected: 2 tests pass.

- [ ] **Step 7: Build verification**

```bash
SQLX_OFFLINE=true cargo check -p ml --lib 2>&1 | tail -5
```

- [ ] **Step 8: Commit**

```bash
git add crates/ml/src/cuda_pipeline/target_q_p99_kernel.cu \
        crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \
        crates/ml/build.rs \
        crates/ml/tests/sp4_producer_unit_tests.rs
git commit -m "$(cat <<'EOF'
feat(sp4): Task A5 — target_q_p99 producer kernel + Pearls A/D wire-up

First end-to-end SP4 producer. Kernel reads denoise_target_q_buf,
computes p99(|target_q|) via sp4_histogram_p99<256>, writes step_obs
to producer_step_scratch_buf[0]. Launcher syncs, applies Pearls A+D
via host-side pearls_ad_update (zero-copy mapped-pinned reads), writes
new x_mean to ISV[TARGET_Q_BOUND_INDEX] + Wiener state back to
wiener_state_buf. Cold-path launch (not in captured graph).

2 unit tests cover: Pearl A first-observation correctness; Pearl D
stationary-signal anchoring.

Producer pattern established. Tasks A6-A12 reuse this shape.

No consumer wired yet — Mech 1's clamp still uses 10 × Q_ABS_REF.max(1.0).
Behavior unchanged.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
EOF
)"
```

---

### Task A6: atom_pos_p99 producer kernel (4 branches)

**Files:**
- Create: `crates/ml/src/cuda_pipeline/atom_pos_p99_kernel.cu`
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`
- Modify: `crates/ml/build.rs`
- Modify: `crates/ml/tests/sp4_producer_unit_tests.rs`

Same shape as Task A5. Single kernel parameterized by `branch` argument; launcher loops over 4 branches.

- [ ] **Step 1: Create the kernel**

```cuda
// crates/ml/src/cuda_pipeline/atom_pos_p99_kernel.cu
#include "sp4_histogram_p99.cuh"

extern "C" __global__ void atom_pos_p99_update(
    const float* __restrict__ atom_positions_branch,  /* slice for this branch */
    int          count,
    float*       __restrict__ scratch_buf,
    int          scratch_idx                          /* per-branch slot */
) {
    if (blockIdx.x != 0) return;
    extern __shared__ int s_warp_tiles[];
    float p99 = sp4_histogram_p99<256>(atom_positions_branch, count);
    if (threadIdx.x == 0) scratch_buf[scratch_idx] = p99;
}
```

- [ ] **Step 2: Add launcher method (loops over 4 branches)**

```rust
pub fn launch_sp4_atom_pos_p99_all_branches(&mut self) -> Result<(), MLError> {
    use crate::cuda_pipeline::sp4_isv_slots::{atom_pos_bound, SP4_SLOT_BASE};
    use crate::cuda_pipeline::sp4_wiener_ema::{pearls_ad_update, WienerState};

    const SCRATCH_BASE: usize = 1;  // slots 1..5 for 4 branches
    const SHARED_BYTES: u32 = 8 * 256 * 4;
    let cfg = LaunchConfig { grid_dim: (1, 1, 1), block_dim: (256, 1, 1), shared_mem_bytes: SHARED_BYTES };

    for branch in 0..4 {
        // Compute per-branch slice pointer + length.
        let (branch_ptr, branch_len) = self.atom_positions_branch_slice(branch);
        let count = branch_len as i32;
        let scratch_idx = (SCRATCH_BASE + branch) as i32;

        unsafe {
            self.stream
                .launch_builder(&self.atom_pos_p99_update)
                .arg(&branch_ptr)
                .arg(&count)
                .arg(&self.producer_step_scratch_buf.dev_ptr())
                .arg(&scratch_idx)
                .launch(cfg)
                .map_err(|e| MLError::ModelError(format!("atom_pos_p99_update[{branch}]: {e}")))?;
        }
    }
    self.stream.synchronize()?;

    // Host-side Pearls A+D for each branch.
    for branch in 0..4 {
        let scratch_idx = SCRATCH_BASE + branch;
        let step_obs = unsafe { *self.producer_step_scratch_buf.host_ptr().add(scratch_idx) };
        if step_obs == 0.0 { continue; }

        let isv_idx = atom_pos_bound(branch);
        let wiener_offset = (isv_idx - SP4_SLOT_BASE) * 3;
        let prev_x_mean = unsafe { *(self.isv_signals_buf.host_ptr() as *mut f32).add(isv_idx) };
        let mut state = unsafe {
            let p = self.wiener_state_buf.host_ptr() as *mut f32;
            WienerState {
                sample_var: *p.add(wiener_offset),
                diff_var:   *p.add(wiener_offset + 1),
                x_lag:      *p.add(wiener_offset + 2),
            }
        };
        let new_x_mean = pearls_ad_update(prev_x_mean, &mut state, step_obs);
        unsafe {
            let isv_p = self.isv_signals_buf.host_ptr() as *mut f32;
            *isv_p.add(isv_idx) = new_x_mean;
            let wp = self.wiener_state_buf.host_ptr() as *mut f32;
            *wp.add(wiener_offset)     = state.sample_var;
            *wp.add(wiener_offset + 1) = state.diff_var;
            *wp.add(wiener_offset + 2) = state.x_lag;
        }
    }
    Ok(())
}
```

`atom_positions_branch_slice(branch)` returns `(device_ptr, len)` for `atom_positions[branch * num_atoms..(branch+1) * num_atoms]`. Add this accessor (search existing code for `atom_positions` field; pattern matches existing C51 atom-position accessors).

- [ ] **Step 3: Add kernel handle field + load in constructor.**

(same pattern as Task A5)

- [ ] **Step 4: Unit tests (4 tests, one per branch)**

```rust
#[test]
#[ignore = "requires GPU"]
fn sp4_atom_pos_p99_per_branch_writes_distinct_isv_slots() {
    let mut trainer = build_minimal_test_trainer();
    // Fill branch 0 with 10×, branch 1 with 100×, branch 2 with 1000×, branch 3 with 10×.
    fill_atom_positions(&mut trainer, &[10.0, 100.0, 1000.0, 10.0]);
    trainer.launch_sp4_atom_pos_p99_all_branches().expect("launch");

    let v0 = trainer.read_isv_signal_at(atom_pos_bound(0));
    let v1 = trainer.read_isv_signal_at(atom_pos_bound(1));
    let v2 = trainer.read_isv_signal_at(atom_pos_bound(2));
    let v3 = trainer.read_isv_signal_at(atom_pos_bound(3));

    assert!((v0 - 10.0).abs() < 1.0,   "branch 0 p99 ~10");
    assert!((v1 - 100.0).abs() < 5.0,  "branch 1 p99 ~100");
    assert!((v2 - 1000.0).abs() < 50.0,"branch 2 p99 ~1000");
    assert!((v3 - 10.0).abs() < 1.0,   "branch 3 p99 ~10 (independent of branches 0/1/2)");
}
```

- [ ] **Step 5: Build + test + commit**

```bash
SQLX_OFFLINE=true cargo test -p ml --tests sp4_atom_pos -- --ignored --nocapture 2>&1 | tail -5
SQLX_OFFLINE=true cargo check -p ml --lib 2>&1 | tail -5

git add crates/ml/src/cuda_pipeline/atom_pos_p99_kernel.cu \
        crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \
        crates/ml/build.rs \
        crates/ml/tests/sp4_producer_unit_tests.rs
git commit -m "feat(sp4): Task A6 — atom_pos_p99 producer × 4 branches

Single kernel parameterized by branch slice, launcher loops 0..4.
4 ISV slots (ATOM_POS_BOUND[0..4]) populated per step. Same Pearls A+D
host-side pattern as Task A5. Unit test verifies per-branch independence.

No consumer wired yet — Mech 2's atom-position clamps still use 10 × Q_ABS_REF.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
"
```

---

### Task A7: param_group_oracle producer kernel (Pearl B fused)

**Files:**
- Create: `crates/ml/src/cuda_pipeline/param_group_oracle_kernel.cu`
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`
- Modify: `crates/ml/build.rs`
- Modify: `crates/ml/tests/sp4_producer_unit_tests.rs`

This is the largest producer kernel — 5 passes (A through E for trunk, A through D for non-trunk groups). Fuses 4 (or 5 for trunk) outputs per launch.

- [ ] **Step 1: Create the kernel**

```cuda
// crates/ml/src/cuda_pipeline/param_group_oracle_kernel.cu
//
// SP4 Pearl B fused per-param-group statistics oracle.
//
// Per group, reads (params, grads, adam_m, adam_v) once and computes 4 outputs:
//   WEIGHT_BOUND[group] = p99(|params|)
//   ADAM_M_BOUND[group] = p99(|adam_m|)
//   ADAM_V_BOUND[group] = p99(|adam_v|)
//   WD_RATE[group]      = |w · g| / max(||w||², EPS_DIV)   (ε from caller)
// For group 0 (trunk) only:
//   L1_LAMBDA[trunk]   = (mean|g| / max(mean|w|, EPS_DIV)) × entropy_deficit_D
//
// 4-5 passes within a single block, sharing __syncthreads() barriers.

#include "sp4_histogram_p99.cuh"

#define BLOCK_SIZE 256
#define EPS_DIV 1.0e-8f

extern "C" __global__ void param_group_oracle_update(
    const float* __restrict__ params,
    const float* __restrict__ grads,
    const float* __restrict__ adam_m,
    const float* __restrict__ adam_v,
    int          count,                /* slice length */
    int          k_in,                 /* input dim (only used by group-0 L1 entropy pass; pass 0 for non-trunk) */
    int          h_dim,                /* output dim of layer (only for group 0) */
    float*       __restrict__ scratch_buf,
    int          weight_scratch_idx,   /* WEIGHT_BOUND scratch slot for this group */
    int          adam_m_scratch_idx,
    int          adam_v_scratch_idx,
    int          wd_rate_scratch_idx,
    int          l1_lambda_scratch_idx /* -1 if this is not group 0 */
) {
    if (blockIdx.x != 0) return;
    extern __shared__ int s_warp_tiles[];
    const int tid = threadIdx.x;

    // Pass A: WEIGHT_BOUND p99(|params|).
    float p99_w = sp4_histogram_p99<BLOCK_SIZE>(params, count);
    if (tid == 0) scratch_buf[weight_scratch_idx] = p99_w;
    __syncthreads();

    // Pass B: ADAM_M_BOUND p99(|adam_m|).
    float p99_m = sp4_histogram_p99<BLOCK_SIZE>(adam_m, count);
    if (tid == 0) scratch_buf[adam_m_scratch_idx] = p99_m;
    __syncthreads();

    // Pass C: ADAM_V_BOUND p99(|adam_v|).
    float p99_v = sp4_histogram_p99<BLOCK_SIZE>(adam_v, count);
    if (tid == 0) scratch_buf[adam_v_scratch_idx] = p99_v;
    __syncthreads();

    // Pass D: WD_RATE = |w·g|/||w||² + (mean|g|, mean|w| as side products).
    // 4-way parallel block-wide reduce: each thread maintains 4 register
    // accumulators (w_dot_g, w_norm_sq, abs_g_sum, abs_w_sum), tree-reduce
    // each via shared memory.
    float l_wdg = 0.0f, l_wsq = 0.0f, l_ag = 0.0f, l_aw = 0.0f;
    for (int i = tid; i < count; i += BLOCK_SIZE) {
        float w = params[i];
        float g = grads[i];
        l_wdg += w * g;
        l_wsq += w * w;
        l_ag  += fabsf(g);
        l_aw  += fabsf(w);
    }
    // Reuse s_warp_tiles[0..BLOCK_SIZE) for each accumulator's tree-reduce.
    __shared__ float s_red[4][BLOCK_SIZE];
    s_red[0][tid] = l_wdg;
    s_red[1][tid] = l_wsq;
    s_red[2][tid] = l_ag;
    s_red[3][tid] = l_aw;
    __syncthreads();
    for (int s = BLOCK_SIZE / 2; s > 0; s >>= 1) {
        if (tid < s) {
            s_red[0][tid] += s_red[0][tid + s];
            s_red[1][tid] += s_red[1][tid + s];
            s_red[2][tid] += s_red[2][tid + s];
            s_red[3][tid] += s_red[3][tid + s];
        }
        __syncthreads();
    }
    float wdg = 0.0f, wsq = 0.0f, mean_g = 0.0f, mean_w = 0.0f;
    if (tid == 0) {
        wdg    = s_red[0][0];
        wsq    = s_red[1][0];
        mean_g = s_red[2][0] / (float)count;
        mean_w = s_red[3][0] / (float)count;
        float wd_rate = fabsf(wdg) / fmaxf(wsq, EPS_DIV);
        scratch_buf[wd_rate_scratch_idx] = wd_rate;
    }
    __syncthreads();

    // Pass E: L1_LAMBDA[trunk] = (mean|g|/max(mean|w|,ε)) × D, group 0 only.
    if (l1_lambda_scratch_idx >= 0 && k_in > 0 && h_dim > 0) {
        // Per-feature gradient L2 norm: g_feat[i] = sqrt(Σ_h grads[h*K + i]²).
        // Then normalize → entropy → deficit D.
        // Single-block: cooperatively compute g_feat[i] for i ∈ [0, k_in).
        // K_in is small (typically ~40 for trunk input dim); each thread handles one i.
        __shared__ float s_g_feat[1024];   // assumes K_in ≤ 1024; dispatcher ensures this
        if (tid < k_in) {
            float sum_sq = 0.0f;
            for (int h = 0; h < h_dim; ++h) {
                float v = grads[h * k_in + tid];
                sum_sq += v * v;
            }
            s_g_feat[tid] = sqrtf(sum_sq);
        }
        __syncthreads();

        // Compute Σ g_feat for normalization, then entropy.
        if (tid == 0) {
            float total = 0.0f;
            for (int i = 0; i < k_in; ++i) total += s_g_feat[i];
            total = fmaxf(total, EPS_DIV);

            float h_observed = 0.0f;
            for (int i = 0; i < k_in; ++i) {
                float p = s_g_feat[i] / total;
                if (p > EPS_DIV) h_observed -= p * logf(p);
            }
            float h_uniform = logf((float)k_in);
            float deficit = fmaxf(0.0f, (h_uniform - h_observed) / fmaxf(h_uniform, EPS_DIV));

            float lambda = (mean_g / fmaxf(mean_w, EPS_DIV)) * deficit;
            scratch_buf[l1_lambda_scratch_idx] = lambda;
        }
    }
}
```

- [ ] **Step 2-7:** Add launcher methods (one wrapper that loops over 8 groups), kernel handle, registration, unit tests, commit.

Launcher:

```rust
pub fn launch_sp4_param_group_oracles_all_groups(&mut self) -> Result<(), MLError> {
    use crate::cuda_pipeline::sp4_isv_slots::{
        weight_bound, adam_m_bound, adam_v_bound, wd_rate, L1_LAMBDA_TRUNK_INDEX,
        SP4_SLOT_BASE, ParamGroup,
    };
    // Scratch slots for this producer family: 5..37 (32 slots = 8 groups × 4 outputs)
    // + 1 slot for L1_LAMBDA (group 0 only) = 38 total.
    const SCRATCH_BASE_WEIGHT:   usize = 5;
    const SCRATCH_BASE_ADAM_M:   usize = 13;
    const SCRATCH_BASE_ADAM_V:   usize = 21;
    const SCRATCH_BASE_WD_RATE:  usize = 29;
    const SCRATCH_L1_LAMBDA:     usize = 37;

    for g in ParamGroup::ALL.iter() {
        let group_idx = g.idx();
        let (params_ptr, grads_ptr, m_ptr, v_ptr, count, k_in, h_dim) =
            self.param_group_buffers(*g)?;

        let l1_idx = if group_idx == 0 { SCRATCH_L1_LAMBDA as i32 } else { -1_i32 };

        // ... launch kernel with scratch slot indices ...
    }
    self.stream.synchronize()?;

    // Host-side Pearls A+D for each group's 4-5 outputs (32 + 1 = 33 EMAs).
    for g in ParamGroup::ALL.iter() {
        // Apply pearls_ad_update to (weight, adam_m, adam_v, wd_rate) for this group,
        // and l1_lambda for group 0.
        // Each call: read prev_x_mean from ISV, read state from wiener_state_buf,
        // compute new_x_mean, write back. Same pattern as Task A5.
    }
    Ok(())
}
```

`param_group_buffers(group)` returns the 4 pointers + count + k_in + h_dim for the group. Implement based on existing param-tensor layout (see `gpu_dqn_trainer.rs::launch_adam_update` for trunk; `gpu_iqn_head.rs::online_params_ptr` for IQN; etc.).

Unit tests: per group, verify the 4 outputs are correct against synthetic inputs (Gaussian samples for params/grads, deterministic |w·g|/||w||² for WD).

- [ ] **Final step: Commit**

```bash
git add crates/ml/src/cuda_pipeline/param_group_oracle_kernel.cu \
        crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \
        crates/ml/build.rs \
        crates/ml/tests/sp4_producer_unit_tests.rs
git commit -m "feat(sp4): Task A7 — Pearl B param_group_oracle (fused 4-5 outputs)

Per param group, single kernel reads (params, grads, adam_m, adam_v) once
and computes WEIGHT/ADAM_M/ADAM_V_BOUND p99 (Passes A-C) + WD_RATE (Pass D)
in a single block. Group 0 (trunk) adds Pass E for L1_LAMBDA via gradient-
direction entropy deficit.

8 oracle launches per step. Pearl B fusion: 4× memory bandwidth reduction
vs naive 32 separate producers. Per-output Pearls A+D applied host-side.

8 unit tests cover the 4 per-group outputs (one assertion per output type)
against synthetic inputs.

No consumer wired yet — Adam kernels still use hardcoded weight_decay=0.01
and weight_clamp_max_abs = 100 × Q_ABS_REF.max(1.0). Behavior unchanged.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
"
```

---

### Task A8: grad_norm_p99 producer kernel

Same shape as Task A5. Reads `grad_norm_buf` (single-element f32 = current step's grad-norm). Wait — `grad_norm_buf` is single-element. p99 over 1 element = that element. So this producer is degenerate; the kernel just writes `grad_norm_buf[0]` to `producer_step_scratch_buf[scratch_idx]`. No histogram needed.

- [ ] **Step 1: Create simplified kernel**

```cuda
// crates/ml/src/cuda_pipeline/grad_norm_p99_kernel.cu
//
// SP4 producer for GRAD_CLIP_BOUND. grad_norm_buf is single-element f32;
// p99 == max == the value itself. Kernel just copies.

extern "C" __global__ void grad_norm_p99_update(
    const float* __restrict__ grad_norm_buf,   /* [1] */
    float*       __restrict__ scratch_buf,
    int          scratch_idx
) {
    if (blockIdx.x != 0 || threadIdx.x != 0) return;
    scratch_buf[scratch_idx] = grad_norm_buf[0];
}
```

- [ ] **Steps 2-5:** launcher + kernel handle + 1 unit test + commit.

Launcher writes Pearls A+D output to `ISV[GRAD_CLIP_BOUND_INDEX]`. Test verifies stationary-input behavior matches Pearl D expectations.

---

### Task A9: h_s2_p99 producer kernel

Same shape as Task A5 (full histogram-p99 over `save_h_s2`). One ISV slot: `H_S2_BOUND_INDEX`.

---

### Task A10: Wire 4 captured-graph launches in fused_training.rs (target_q + 4 atom_pos)

**Files:**
- Modify: `crates/ml/src/trainers/dqn/fused_training.rs`

The buffer-based bounds (target_q, atom_pos) need their producers to run BEFORE the consumer. Per the spec, this means the producer launches inside the captured graph immediately before the corresponding clamp/use site.

For target_q: producer launch goes immediately after `compute_denoise_target_q` writes `denoise_target_q_buf` (so producer reads pre-clamp value), and BEFORE Mech 1's clamp engages.

For atom_pos: producers launch immediately after `atoms_update_kernel` writes positions, BEFORE downstream consumption.

- [ ] **Step 1: Locate the captured-graph compose point**

```bash
grep -n "compute_denoise_target_q\|atoms_update" /home/jgrusewski/Work/foxhunt/crates/ml/src/trainers/dqn/fused_training.rs | head -10
```

Find the producer-launch insertion points.

- [ ] **Step 2: Insert calls**

After the `compute_denoise_target_q` call:

```rust
// SP4 Task A10: target_q producer launch (BEFORE Mech 1 clamp).
self.trainer.launch_sp4_target_q_p99()?;
```

After the `atoms_update_kernel` calls:

```rust
// SP4 Task A10: atom_pos producers launch (BEFORE Mech 2 clamps in
// atoms_update_kernel.cu's next iteration; but here they read atom_positions
// AFTER it's written, BEFORE the next forward pass uses it).
self.trainer.launch_sp4_atom_pos_p99_all_branches()?;
```

These launches happen in cold-path Rust, NOT inside the captured graph. The producer kernel itself uses single-block single-launch which the captured graph could include, but for simplicity Layer A keeps producers cold-path. The "before consumer" ordering is enforced by Rust call sequence.

**This is consistent with the existing pattern**: `launch_h_s2_rms_ema` is cold-path despite being a "captured-graph producer" semantically.

- [ ] **Step 3: Build verification**

```bash
SQLX_OFFLINE=true cargo check -p ml --lib 2>&1 | tail -5
```

- [ ] **Step 4: L40S smoke regression-check**

(Optional, run only if local GPU available; otherwise defer to Task A19.)

```bash
./scripts/argo-smoke.sh --ref HEAD 2>&1 | tail -5
```

Wait for completion, verify multi_fold_convergence still passes with no behavior change vs the SP3 baseline.

- [ ] **Step 5: Commit**

```bash
git add crates/ml/src/trainers/dqn/fused_training.rs
git commit -m "feat(sp4): Task A10 — wire target_q + atom_pos producers in fused_training

Cold-path producer launches inserted immediately after their respective
buffer writers (compute_denoise_target_q, atoms_update_kernel). ISV slots
TARGET_Q_BOUND, ATOM_POS_BOUND[0..4] populate correctly each step.

No consumer reads these slots yet — Mech 1/2 clamps unchanged. Behavior
unchanged.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
"
```

---

### Task A11: Wire cold-path producer launches in training_loop.rs

**Files:**
- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs`

Cold-path producer launches go alongside `launch_h_s2_rms_ema` (which is at line 3416 per earlier grep).

- [ ] **Step 1: Locate the launch site**

```bash
grep -n "launch_h_s2_rms_ema" /home/jgrusewski/Work/foxhunt/crates/ml/src/trainers/dqn/trainer/training_loop.rs
```

- [ ] **Step 2: Insert SP4 producer launches**

Immediately after `launch_h_s2_rms_ema(ema_alpha)`:

```rust
// SP4: param-group oracles (Pearl B fused; 8 launches per step).
if let Some(ref fused) = self.fused_ctx {
    if let Err(e) = fused.trainer().launch_sp4_param_group_oracles_all_groups() {
        tracing::warn!("SP4 param_group_oracles launch failed: {e}");
    }
}

// SP4: grad_norm p99 producer.
if let Some(ref fused) = self.fused_ctx {
    if let Err(e) = fused.trainer().launch_sp4_grad_norm_p99() {
        tracing::warn!("SP4 grad_norm_p99 launch failed: {e}");
    }
}

// SP4: h_s2 p99 producer (parallel with retrofit h_s2_rms_ema; both consume save_h_s2).
if let Some(ref fused) = self.fused_ctx {
    if let Err(e) = fused.trainer().launch_sp4_h_s2_p99() {
        tracing::warn!("SP4 h_s2_p99 launch failed: {e}");
    }
}
```

(Note: `launch_sp4_target_q_p99` and `launch_sp4_atom_pos_p99_all_branches` are NOT here — they were wired in fused_training.rs Task A10.)

- [ ] **Step 3: Build verification**

```bash
SQLX_OFFLINE=true cargo check -p ml --lib 2>&1 | tail -5
```

- [ ] **Step 4: Commit**

```bash
git add crates/ml/src/trainers/dqn/trainer/training_loop.rs
git commit -m "feat(sp4): Task A11 — wire cold-path producer launches in training_loop

10 cold-path SP4 producer launches per step (8 param-group oracles +
grad_norm + h_s2). All ISV slots [131..171) populate correctly.

No consumer reads these slots yet. Behavior unchanged.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
"
```

---

### Task A12: StateResetRegistry entries for SP4 ISV slots and Wiener state

**Files:**
- Modify: `crates/ml/src/trainers/dqn/state_reset_registry.rs`

- [ ] **Step 1: Locate registry pattern**

```bash
grep -n "fn build_registry\|description:" /home/jgrusewski/Work/foxhunt/crates/ml/src/trainers/dqn/state_reset_registry.rs | head -10
```

- [ ] **Step 2: Add 40 ISV bound entries + 141 Wiener-state entries + 21 retrofit Wiener entries**

In the registry-building function, add a contiguous block of SP4 entries:

```rust
// ── SP4 (40 ISV bound slots — sentinel 0 per Pearl A) ──────────────────
use crate::cuda_pipeline::sp4_isv_slots::*;

for slot_idx in SP4_SLOT_BASE..SP4_SLOT_END {
    builder.add(StateResetEntry {
        target: ResetTarget::IsvSlot { index: slot_idx },
        cold_start_value: 0.0,
        description: format!("SP4 ISV bound slot {slot_idx} — sentinel 0 (Pearl A first-observation replacement on next producer launch)").into(),
    });
}

// ── SP4 Wiener-state buffer entries (40 SP4 + 7 retrofit) × 3 floats = 141 ──
const WIENER_TOTAL: usize = 47 * 3;
for offset in 0..WIENER_TOTAL {
    builder.add(StateResetEntry {
        target: ResetTarget::MappedPinnedF32 {
            buf_name: "wiener_state_buf",
            offset,
        },
        cold_start_value: 0.0,
        description: format!("SP4 Wiener state offset {offset} ({}{}) — sentinel 0 per Pearls A/D",
            if offset < 120 { "SP4 slot " } else { "retrofit " },
            offset / 3).into(),
    });
}

// ── SP4 Pearl C engagement counter — zeroed at fold boundary ──────────
const ENGAGE_BUF_LEN: usize = 8 * 256;
for offset in 0..ENGAGE_BUF_LEN {
    builder.add(StateResetEntry {
        target: ResetTarget::MappedPinnedI32 {
            buf_name: "clamp_engage_per_block_buf",
            offset,
        },
        cold_start_value: 0,
        description: format!("SP4 Pearl C engagement counter [{}] — zeroed per fold", offset).into(),
    });
}
```

If `ResetTarget::MappedPinnedF32` / `I32` variants don't exist in the registry, add them following the existing `ResetTarget::IsvSlot` enum pattern.

- [ ] **Step 3: Build verification**

```bash
SQLX_OFFLINE=true cargo check -p ml --lib 2>&1 | tail -5
SQLX_OFFLINE=true cargo test -p ml --lib state_reset_registry:: 2>&1 | tail -10
```

- [ ] **Step 4: Commit**

```bash
git add crates/ml/src/trainers/dqn/state_reset_registry.rs
git commit -m "feat(sp4): Task A12 — StateResetRegistry entries for SP4 + Wiener state

40 ISV bound entries (sentinel 0 per Pearl A) + 141 Wiener-state float
entries (sentinel 0 per Pearls A/D) + 2048 Pearl C engagement counter
int entries (zeroed per fold).

At fold boundary, all SP4 ISV slots reset to 0; producer's first step
detects sentinel and replaces directly with step_observation. No theoretical
bootstrap formulas — pure first-observation per Pearl A.

Behavior unchanged on F0 (first fold has no prior data so reset is a no-op
on cold-start; subsequent folds get clean-slate signal tracking).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
"
```

---

### Task A13: Retrofit existing 7 ISV producer kernels with Pearls A+D

**Files (one per kernel):**
- Modify: `crates/ml/src/cuda_pipeline/h_s2_rms_ema_kernel.cu`
- Modify: `crates/ml/src/cuda_pipeline/aux_heads_loss_ema_kernel.cu`
- Modify: `crates/ml/src/cuda_pipeline/moe_kernels.cu` (`moe_gate_entropy_ema_update`)
- Modify: `crates/ml/src/cuda_pipeline/vsn_mask_ema_kernel.cu`
- Modify: `crates/ml/src/cuda_pipeline/iqn_quantile_ema_kernel.cu`
- Modify: `crates/ml/src/cuda_pipeline/reward_component_ema_kernel.cu`
- Modify: their corresponding launcher methods in `gpu_dqn_trainer.rs`

The retrofit replaces the kernel-internal hardcoded α with: kernel writes step_observation to a producer scratch slot; launcher applies Pearls A+D via `pearls_ad_update` (same pattern as Tasks A5-A9).

- [ ] **Step 1: One retrofit per kernel — repeat the pattern**

For each existing producer kernel:

1. Modify the kernel to write step_observation to a designated scratch slot (instead of doing the EMA inline).
2. Modify the Rust launcher to call the kernel, sync, then apply `pearls_ad_update`.
3. Existing single tests should still pass; the retrofit changes the EMA semantics but for stationary signals the resulting ISV value should be similar.

Reference example for `h_s2_rms_ema_kernel.cu`:

```cuda
// Was (line 25-69, see `h_s2_rms_ema_kernel.cu` in repo):
extern "C" __global__ void h_s2_rms_ema_update(
    const float* __restrict__ h_s2,
    int B, int SH2,
    float* __restrict__ isv,
    int isv_h_s2_rms_ema_index,
    float ema_alpha   // ← REMOVED
) {
    // ... compute RMS into smem[0] ...
    if (tid == 0) {
        const float rms   = sqrtf(smem[0] / (float)N);
        const int   slot  = isv_h_s2_rms_ema_index;
        const float prev  = isv[slot];
        isv[slot] = (1.0f - ema_alpha) * prev + ema_alpha * rms;   // ← REMOVED
    }
}

// Becomes:
extern "C" __global__ void h_s2_rms_ema_update(
    const float* __restrict__ h_s2,
    int B, int SH2,
    float* __restrict__ scratch_buf,
    int   scratch_idx
) {
    // ... compute RMS into smem[0] ...
    if (tid == 0) {
        const float rms = sqrtf(smem[0] / (float)N);
        scratch_buf[scratch_idx] = rms;   // ← writes step_observation to scratch
    }
}
```

Launcher change in `gpu_dqn_trainer.rs::launch_h_s2_rms_ema`:

```rust
pub fn launch_h_s2_rms_ema(&mut self, _ema_alpha_unused: f32) -> Result<(), MLError> {
    // SP4 retrofit: kernel writes step_observation to scratch; host applies Pearls A+D.
    use crate::cuda_pipeline::sp4_wiener_ema::{pearls_ad_update, WienerState};

    const SCRATCH_IDX: usize = 40;  // first retrofit slot (40 SP4 + offset 0..7 for retrofits)

    // ... launch kernel with scratch_idx=40 ...
    self.stream.synchronize()?;

    let step_obs = unsafe { *self.producer_step_scratch_buf.host_ptr().add(SCRATCH_IDX) };
    if step_obs == 0.0 { return Ok(()); }

    let isv_idx = H_S2_RMS_EMA_INDEX;  // 96
    let wiener_offset = 40 * 3 + 0 * 3;  // retrofit slot 0 → Wiener offset 120

    // ... apply pearls_ad_update, write back ...
    Ok(())
}
```

The `_ema_alpha_unused` param stays in the signature for caller-compat; the retrofit drops the value.

- [ ] **Step 2-7:** Repeat for remaining 6 retrofit kernels.

- [ ] **Final step: Commit**

One commit per retrofit kernel (7 commits total) so each is independently testable. Or one combined commit after all 7 retrofits build cleanly. **Recommendation: one commit per kernel** for cleaner bisect.

---

### Task A14: Pearl C engagement counter in Adam kernels (5 kernels × 1 change each)

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu` (`dqn_adam_update_kernel`)
- Modify: `crates/ml/src/cuda_pipeline/iqn_dual_head_kernel.cu` (`iqn_adam_kernel`)
- Modify: `crates/ml/src/cuda_pipeline/iql_value_kernel.cu` (`iql_adam_kernel`)
- Modify: `crates/ml/src/cuda_pipeline/attention_backward_kernel.cu` (`attn_adam_kernel`)
- Modify: `crates/ml/src/cuda_pipeline/curiosity_training_kernel.cu` (`curiosity_adam_step`)

Each kernel gets:
- New trailing args: `int diag_slot, int engage_buf_offset, int* clamp_engage_per_block_buf, int* nan_flags_buf`
- Per-thread register `local_engage = 0`
- After clamp engagement: `local_engage = 1`
- Block tree-reduce shared-memory (no atomicAdd)
- Block-leader writes per-(block, group) engagement count

(NOTE: at this point the clamps still use hardcoded `weight_clamp_max_abs` — Layer B will replace this with ISV-driven. The Pearl C counter is added in Layer A so the infrastructure is in place when Layer B flips.)

- [ ] **Step 1-5 per kernel:** add args + counter + block reduce + commit each. 5 commits total.

Reference snippet (insert immediately after the existing Mech 9 clamp block in each kernel):

```cuda
/* Pearl C: register-then-tree-reduce engagement counter — no atomicAdd. */
__shared__ int s_block_engage[256];   // BLOCK_SIZE assumed 256
int local_engage = 0;
// ... existing Adam compute populates p_val, then:
if (weight_clamp_max_abs > 0.0f) {
    if (fabsf(p_val) > weight_clamp_max_abs) {
        nan_flags_buf[diag_slot] = 1;
        local_engage = 1;
    }
    p_val = fminf(fmaxf(p_val, -weight_clamp_max_abs), weight_clamp_max_abs);
}
s_block_engage[threadIdx.x] = local_engage;
__syncthreads();
for (int s = blockDim.x / 2; s > 0; s >>= 1) {
    if (threadIdx.x < s) s_block_engage[threadIdx.x] += s_block_engage[threadIdx.x + s];
    __syncthreads();
}
if (threadIdx.x == 0) {
    clamp_engage_per_block_buf[engage_buf_offset + blockIdx.x] = s_block_engage[0];
}
```

Then update the launch sites in `gpu_dqn_trainer.rs` etc. to pass the new args. `engage_buf_offset` = `group_idx × MAX_BLOCKS_PER_ADAM` where `group_idx` ∈ [0, 8).

- [ ] **Final step: 5 commits, one per Adam kernel**

---

### Task A15: Pearl C host-side rate-deficit EMA + force-bump trigger

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`

After each Adam launch, host reads `clamp_engage_per_block_buf` (mapped-pinned, zero-copy), sums engagement counts across blocks, divides by group's param count → engagement rate, applies Pearls A+D to `engagement_rate - 0.01` deficit, force-bumps bound when deficit EMA sustained > 0.

- [ ] **Step 1: Implement helper method**

```rust
/// SP4 Pearl C — read engagement count, compute rate, apply rate-deficit
/// Wiener EMA, optionally force-bump bound.
///
/// Called immediately after each Adam launch (per param group).
fn pearl_c_post_adam_engagement_check(&mut self, group: ParamGroup,
                                        max_blocks: usize) -> Result<(), MLError> {
    let group_idx = group.idx();
    let buf_offset = group_idx * MAX_BLOCKS_PER_ADAM;
    let total_engage: i32 = unsafe {
        let p = self.clamp_engage_per_block_buf.host_ptr() as *mut i32;
        (0..max_blocks).map(|b| *p.add(buf_offset + b)).sum()
    };
    let group_param_count = self.param_group_count(group);
    let engagement_rate = total_engage as f32 / group_param_count as f32;
    let rate_deficit = engagement_rate - 0.01;

    // Track rate_deficit EMA via dedicated Wiener slot (re-use one of
    // the engagement-rate diagnostic slots; or use a separate per-group
    // Wiener-state triple stored in a scratch region).
    // ... apply pearls_ad_update on rate_deficit signal ...

    // Force-bump: if rate_deficit_ema > 0 AND diff_var low (stable bump),
    // set ISV[WEIGHT_BOUND[group]] = step_max_observed_this_step.
    // (step_max comes from the param_group_oracle's same-step observation.)
    // ... force-bump logic ...
    Ok(())
}
```

(Implementation requires a per-group rate-deficit Wiener triple, allocated in `wiener_state_buf` extension or in a dedicated scratch buffer.)

- [ ] **Step 2-4:** Wire calls into Adam launch sites, build, commit.

---

### Task A16-A18: Layer A integration validation

- [ ] **Task A16:** Run all 40 unit tests on a CI builder (or local L40S worktree) — `cargo test -p ml --tests sp4_ -- --ignored --nocapture` → 40 tests pass with <5% quantization error and <2% Wiener-formula error.
- [ ] **Task A17:** L40S smoke regression check at end-of-Layer-A HEAD. multi_fold_convergence should pass with no behavior change (no consumer is reading SP4 ISV slots yet, so training trajectory is identical to SP3 close-out).
- [ ] **Task A18:** Update `docs/dqn-wire-up-audit.md` with one large Invariant 7 entry covering all of Layer A.

---

## Layer B — Atomic consumer migration (single coordinated commit)

**`feedback_no_partial_refactor` requires all consumer sites flip together.** This is ONE commit. It contains many edits but they are coordinated. The plan documents this as a single task with multiple sub-steps.

### Task B1: Atomic consumer-flip commit

**Files (all modified in one commit):**
- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (Mech 1, 6, 10 clamp consumers + per-group Adam launch sites read from ISV)
- `crates/ml/src/cuda_pipeline/atoms_update_kernel.cu` (Mech 2 clamp)
- `crates/ml/src/cuda_pipeline/iql_value_kernel.cu` (Mech 2 clamp + Adam ISV-source)
- `crates/ml/src/cuda_pipeline/experience_kernels.cu` (Mech 2 clamp)
- `crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu` (Mech 5 fused diagnostic per-slot ISV read; `dqn_adam_update_kernel` weight_decay/l1_lambda from launch arg)
- `crates/ml/src/cuda_pipeline/iqn_dual_head_kernel.cu` (Adam: weight_clamp + weight_decay from ISV)
- `crates/ml/src/cuda_pipeline/attention_backward_kernel.cu` (same)
- `crates/ml/src/cuda_pipeline/curiosity_training_kernel.cu` (same)
- `crates/ml/src/cuda_pipeline/decision_transformer.rs` (DT keeps 0.0 disable for weight_clamp; adds 0.0 disable for new weight_decay arg)
- `crates/ml/src/trainers/dqn/config.rs` or `crates/ml/src/trainers/dqn/hyperparams.rs` (remove `weight_decay`, `l1_lambda` fields)

- [ ] **Step 1: Mech 1 (target_q clamp) — replace hardcoded multiplier**

In `gpu_dqn_trainer.rs::compute_denoise_target_q`, find the existing clamp:

```rust
// Was:
let q_abs_ref = self.read_isv_signal_at(Q_ABS_REF_INDEX);
let max_abs_target_q = 10.0_f32 * q_abs_ref.max(1.0_f32);
self.launch_clamp_finite_f32(self.denoise_target_q_buf.raw_ptr(),
                             self.denoise_target_q_buf.len(), max_abs_target_q)?;
// Becomes:
use crate::cuda_pipeline::sp4_isv_slots::TARGET_Q_BOUND_INDEX;
use crate::cuda_pipeline::sp4_wiener_ema::EPS_CLAMP_FLOOR;
let max_abs_target_q = self.read_isv_signal_at(TARGET_Q_BOUND_INDEX).max(EPS_CLAMP_FLOOR);
self.launch_clamp_finite_f32(self.denoise_target_q_buf.raw_ptr(),
                             self.denoise_target_q_buf.len(), max_abs_target_q)?;
```

- [ ] **Step 2: Mech 2 (atom positions, 3 sites)**

In `atoms_update_kernel.cu`, `iql_value_kernel.cu`, `experience_kernels.cu`, change the existing:

```cuda
const float q_abs_ref_eff = fmaxf(isv_signals[16], 1.0f);
const float max_atom_abs  = 10.0f * q_abs_ref_eff;
new_pos = fminf(fmaxf(new_pos, -max_atom_abs), max_atom_abs);
```

To:

```cuda
const float max_atom_abs = fmaxf(isv_signals[ATOM_POS_BOUND_BASE + branch], 1.0f);
new_pos = fminf(fmaxf(new_pos, -max_atom_abs), max_atom_abs);
```

`ATOM_POS_BOUND_BASE = 132` is hardcoded in the kernel as `#define ATOM_POS_BOUND_BASE 132`. Add the `#define` near top of each .cu.

- [ ] **Step 3: Mech 5 fused diagnostic (`dqn_nan_check_fused_f32_kernel`)**

In `dqn_utility_kernels.cu`, the existing per-slot threshold block:

```cuda
if (slot >= 12 && slot < 16)        thr = 100.0f * q_abs_ref_eff;
else if (slot >= 16 && slot < 20)   thr = 1.0e6f * q_abs_ref_eff * q_abs_ref_eff;
// ... etc.
```

Becomes:

```cuda
// SP4: per-slot threshold reads from corresponding ISV bound directly.
// No magnitude multipliers. Diagnostic = clamp engagement (sticky-flag pattern).
if (slot >= 12 && slot < 20) {
    // Adam m/v: read ISV[ADAM_M_BOUND[group]] / ISV[ADAM_V_BOUND[group]]
    // Group derives from absolute slot 36+rel: groups 0-3 for slot 36-39 / 40-43
    int abs_slot = base_flag_idx + slot;
    int group_idx = (abs_slot >= 40) ? (abs_slot - 40) : (abs_slot - 36);
    int isv_bound_idx = (abs_slot < 40)
        ? (ADAM_M_BOUND_BASE + group_idx)   // ADAM_M
        : (ADAM_V_BOUND_BASE + group_idx);  // ADAM_V
    thr = isv_signals[isv_bound_idx];
}
else if (slot == 20 || slot == 21) {
    int abs_slot = base_flag_idx + slot;
    int group_idx = abs_slot - 44;  // mapping per existing slot 44/45 = trunk/heads
    thr = isv_signals[WEIGHT_BOUND_BASE + group_idx];
}
else if (slot == 22) thr = isv_signals[TARGET_Q_BOUND_INDEX];
else if (slot == 23) thr = isv_signals[ATOM_POS_BOUND_BASE];  // any branch; outliers above are bounded
else if (slot == 24) thr = isv_signals[WEIGHT_BOUND_BASE + 3]; // IQN weight
else if (slot == 25) thr = isv_signals[H_S2_BOUND_INDEX];
```

Remove the `q_abs_ref_eff` and `h_s2_rms_ema_eff` kernel parameters (they're now unused).

- [ ] **Step 4: Mech 6 (adaptive_clip upper_bound)**

In `gpu_dqn_trainer.rs::update_adaptive_clip` (around line 20367):

```rust
// Was:
let abs_mult = self.read_isv_signal_at(Q_ABS_REF_INDEX).max(1.0_f32);
let prev_slow_ema = unsafe { *self.grad_norm_slow_ema_pinned };
let upper_bound = (prev_slow_ema * 100.0_f32 * abs_mult).max(MIN_CLIP);
let new_clip = new_clip.min(upper_bound);
// Becomes:
use crate::cuda_pipeline::sp4_isv_slots::GRAD_CLIP_BOUND_INDEX;
use crate::cuda_pipeline::sp4_wiener_ema::EPS_CLAMP_FLOOR;
let upper_bound = self.read_isv_signal_at(GRAD_CLIP_BOUND_INDEX).max(EPS_CLAMP_FLOOR);
let new_clip = new_clip.min(upper_bound);
```

Note: `prev_slow_ema` is no longer read by Mech 6's upper_bound. `grad_norm_slow_ema_pinned` is **NOT orphan** — it has a second live consumer (`fold_warmup_factor_update`, commit `4ef1d8ebb`) that reads it as the cross-fold steady-state grad-norm baseline. Layer C C1 redesigned (2026-05-01): instead of retiring the slow EMA, migrate the host-side EMA UPDATE arithmetic to GPU per `feedback_no_cpu_compute_strict`.

- [ ] **Step 5: Mech 9 (post-Adam weight clamp across 5 Adam kernels)**

In `gpu_dqn_trainer.rs`, each Adam launch site (search for `launch_builder(&self.adam_update)` and similar):

```rust
// Was (per-launch):
let q_abs_ref_eff = self.read_isv_signal_at(Q_ABS_REF_INDEX).max(1.0_f32);
let weight_clamp_max_abs: f32 = 100.0_f32 * q_abs_ref_eff;
// Becomes:
use crate::cuda_pipeline::sp4_isv_slots::{weight_bound, ParamGroup};
use crate::cuda_pipeline::sp4_wiener_ema::EPS_CLAMP_FLOOR;
let weight_clamp_max_abs: f32 = self
    .read_isv_signal_at(weight_bound(<group_for_this_kernel>))
    .max(EPS_CLAMP_FLOOR);
```

Per launch site, `<group_for_this_kernel>` is the appropriate `ParamGroup::*.idx()`. For example, `dqn_adam_update_kernel` (main DQN params) launches with `weight_bound(ParamGroup::DqnTrunk.idx())` — though for DqnTrunk vs DqnValue vs DqnBranches we need to launch separately per group. Existing code launches all main DQN params in one kernel; SP4 splits the bound read by which slice the launch is updating, OR (simpler) reads `MAX(WEIGHT_BOUND[trunk], WEIGHT_BOUND[value], WEIGHT_BOUND[branches])` so the single-kernel launch uses the loosest bound.

Recommendation: keep the single-kernel `dqn_adam_update_kernel` launch, but per-group `weight_clamp_max_abs` requires per-segment clamping inside the kernel (knowing which param index belongs to which group). This is non-trivial; for SP4 simplicity, use a uniform bound = `min(WEIGHT_BOUND[trunk], WEIGHT_BOUND[value], WEIGHT_BOUND[branches])` so all main DQN params use the tightest bound.

- [ ] **Step 6: Mech 9 weight_decay arg from ISV (5 Adam kernels)**

Each Adam launch site that previously took a fixed `weight_decay` config field now reads from `ISV[WD_RATE[group]]`:

```rust
// Was:
.arg(&weight_decay_value)   // was self.hyperparams.weight_decay
// Becomes:
let weight_decay_value: f32 = self
    .read_isv_signal_at(wd_rate(<group_idx>))
    .max(0.0);  // weight_decay must be non-negative
.arg(&weight_decay_value)
```

For the main DQN single-kernel launch: same compromise — use the maximum of trunk/value/branches WD_RATE, or split the launch into 3 sub-launches. Recommendation: max-of-3 for simplicity in this Layer B atomic flip.

- [ ] **Step 7: L1 lambda from ISV (trunk only)**

In `dqn_adam_update_kernel` launch site:

```rust
// Was:
let l1_lambda: f32 = 1e-3;
// Becomes:
use crate::cuda_pipeline::sp4_isv_slots::L1_LAMBDA_TRUNK_INDEX;
let l1_lambda: f32 = self.read_isv_signal_at(L1_LAMBDA_TRUNK_INDEX).max(0.0);
```

- [ ] **Step 8: Mech 10 (h_s2 clamp)**

In `gpu_dqn_trainer.rs::submit_forward_ops_main` (around line 17221):

```rust
// Was:
let h_s2_rms_ema_eff = self.read_isv_signal_at(H_S2_RMS_EMA_INDEX).max(1.0_f32);
let max_abs_h_s2 = 100.0_f32 * h_s2_rms_ema_eff;
// Becomes:
use crate::cuda_pipeline::sp4_isv_slots::H_S2_BOUND_INDEX;
use crate::cuda_pipeline::sp4_wiener_ema::EPS_CLAMP_FLOOR;
let max_abs_h_s2 = self.read_isv_signal_at(H_S2_BOUND_INDEX).max(EPS_CLAMP_FLOOR);
```

- [ ] **Step 9: Remove `weight_decay` and `l1_lambda` from HyperParams**

Find the config struct (likely `crates/ml/src/trainers/dqn/config.rs` or similar):

```bash
grep -rn "weight_decay:\s*f64\|l1_lambda:" /home/jgrusewski/Work/foxhunt/crates/ml/src/trainers/dqn/ | head -5
```

Remove the fields. Add a doc comment pointing to `ISV[WD_RATE[group]]` and `ISV[L1_LAMBDA_TRUNK_INDEX]` per-step.

- [ ] **Step 10: Build verification**

```bash
SQLX_OFFLINE=true cargo check -p ml --lib 2>&1 | tail -10
```

Must compile clean. If `weight_decay` or `l1_lambda` is referenced from any non-SP4 site, remove or migrate that reference.

- [ ] **Step 11: Single coordinated commit**

```bash
git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs \
        crates/ml/src/cuda_pipeline/atoms_update_kernel.cu \
        crates/ml/src/cuda_pipeline/iql_value_kernel.cu \
        crates/ml/src/cuda_pipeline/experience_kernels.cu \
        crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu \
        crates/ml/src/cuda_pipeline/iqn_dual_head_kernel.cu \
        crates/ml/src/cuda_pipeline/attention_backward_kernel.cu \
        crates/ml/src/cuda_pipeline/curiosity_training_kernel.cu \
        crates/ml/src/cuda_pipeline/decision_transformer.rs \
        crates/ml/src/trainers/dqn/config.rs

git commit -m "$(cat <<'EOF'
feat(sp4): Layer B — atomic consumer migration to ISV-driven bounds

Single coordinated commit per `feedback_no_partial_refactor`. All
hardcoded magnitude multipliers (10×, 100×, 1e3×, 1e6×) and ε floors
(.max(1.0)) replaced by per-slot ISV reads with consumer-side
EPS_CLAMP_FLOOR=1.0 numerical safety.

Mechanism mapping:
- Mech 1 target_q clamp: 10 × Q_ABS_REF.max(1.0) → ISV[TARGET_Q_BOUND].max(1.0)
- Mech 2 atom-position clamps (3 sites × 4 branches): 10 × Q_ABS_REF
   → ISV[ATOM_POS_BOUND[branch]]
- Mech 5 fused diagnostic: per-slot ISV reads (no q_abs_ref_eff /
   h_s2_rms_ema_eff args)
- Mech 6 adaptive_clip upper_bound: 100 × slow_ema × Q_ABS_REF
   → ISV[GRAD_CLIP_BOUND]
- Mech 9 post-Adam weight_clamp (5 Adam kernels): 100 × Q_ABS_REF
   → ISV[WEIGHT_BOUND[group]]
- Mech 10 h_s2 clamp: 100 × H_S2_RMS_EMA → ISV[H_S2_BOUND]
- AdamW weight_decay (5 kernels): config field → ISV[WD_RATE[group]]
- L1 lambda (trunk only): 1e-3 → ISV[L1_LAMBDA_TRUNK_INDEX]

Removed `weight_decay` and `l1_lambda` from HyperParams config struct.
Decision Transformer keeps 0.0 disables (offline-RL outside SP4 scope).

cargo check clean.

Validation deferred to Layer C smoke. Expected: F0/F1/F2 all complete
5 epochs; F1 trains past step 1000; F0 ≥ 37.5; F2 ≥ 55; slot 49 quiet.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
EOF
)"
```

---

## Layer C — Validation + cleanup

### Task C1 (REDESIGNED 2026-05-01): Migrate fast/slow grad-norm EMA UPDATE to GPU

**Original plan claimed `grad_norm_slow_ema_pinned` was orphan post-Mech-6 migration. Verification on 2026-05-01 surfaced a SECOND live consumer: `fold_warmup_factor_update` (commit `4ef1d8ebb`, predates SP4) reads the slow EMA as the cross-fold steady-state grad-norm baseline. Legitimate, not orphan.**

**However**, on inspection of `update_adaptive_clip:22720-22737`, the fast/slow EMA UPDATE was running as host-side `(1-α) × prev + α × obs` arithmetic — exactly the pattern `feedback_no_cpu_compute_strict` (saved 2026-05-01) strictly forbids. The actual defect was the compute location, not orphan-status.

**Redesigned scope:**
- Retain `grad_norm_slow_ema_pinned` (legitimate cross-fold baseline consumer is `fold_warmup_factor_update`).
- Retain `grad_norm_fast_ema_pinned` (numerator).
- New `update_grad_norm_emas_kernel.cu` — single-thread fused fast+slow EMA update. Reads `grad_norm_buf[0]`, updates the two mapped-pinned EMA scalars in-place via dev_ptr.
- New Rust launcher `launch_update_grad_norm_emas` chained on the producer's stream after `launch_grad_norm_finalize` (no host sync; sequential same-stream ordering provides happens-before).
- `update_adaptive_clip`'s host-side `unsafe { ... }` arithmetic block replaced with the launcher call.
- Cold-start sentinel (`prev ≤ 0.0` ⇒ assign obs directly), FAST_ALPHA=0.1, SLOW_ALPHA=0.001 — preserved from the deleted host code bit-for-bit.

**Files:**
- Add: `crates/ml/src/cuda_pipeline/update_grad_norm_emas_kernel.cu`
- Modify: `crates/ml/build.rs` (cubin registration)
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (CUBIN static + struct field + constructor loader + launcher + replace host-side block)
- Modify: `crates/ml/src/cuda_pipeline/fold_warmup_factor_kernel.cu` (docstring update — EMAs are now GPU-updated)
- Modify: `crates/ml/src/trainers/dqn/state_reset_registry.rs` (doc-comment reflects GPU-only update path)
- Modify: `docs/dqn-wire-up-audit.md` (Layer C entry)
- Modify: this plan (the section above + this one)
- Modify: `docs/superpowers/specs/2026-04-30-sp4-signal-driven-magnitude-control-design.md` (orphan claim removed)

**No semantic change**: same α values, same cold-start sentinel, same call ordering, same observation source. Pure architectural fix per `feedback_no_cpu_compute_strict`.

- [ ] **Step 1: Build verification**

```bash
SQLX_OFFLINE=true cargo check -p ml --offline 2>&1 | tail -10
```

- [ ] **Step 2: Single coordinated commit per `feedback_no_partial_refactor`** with commit message documenting that the original "retire" framing was wrong, the slow EMA stays, the actual fix is the host→GPU compute migration.

---

### Task C2: Remove dead Mech 5 fused-kernel parameters

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu`
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (launch site)

After Layer B, the `q_abs_ref_eff` and `h_s2_rms_ema_eff` parameters of `dqn_nan_check_fused_f32_kernel` are unused. Remove them from kernel signature + Rust launch args.

- [ ] **Steps 1-4:** edit signatures, build, commit.

---

### Task C3: Run L40S smoke validation

- [ ] **Step 1: Push branch**

```bash
git push origin plan-c-phase-2-thompson 2>&1 | tail -5
```

- [ ] **Step 2: Dispatch smoke**

```bash
./scripts/argo-smoke.sh --ref HEAD 2>&1 | tail -10
```

Note the smoke name (e.g., `smoke-test-XXXXX`).

- [ ] **Step 3: Wait + fetch log**

```bash
# After completion, fetch from MinIO archive:
mc cp "fh/argo-logs/smoke-test-XXXXX/smoke-test-XXXXX/main.log" /tmp/smoke-sp4.log
```

- [ ] **Step 4: Validate against pass criteria** (from spec):

```bash
echo "=== Per-fold Best Sharpe ==="
grep -E "Best Sharpe" /tmp/smoke-sp4.log | head -10

echo "=== NaN-CLAMPED count ==="
grep -c "NaN-CLAMPED-TO-ZERO at step" /tmp/smoke-sp4.log

echo "=== Mech 5 diagnostic firings ==="
grep -E "(36-43|44-45|48|49)=" /tmp/smoke-sp4.log | wc -l

echo "=== L1 lambda ramp-up ==="
grep -E "ISV\[L1_LAMBDA" /tmp/smoke-sp4.log | head -10

echo "=== WD_RATE per fold ==="
grep -E "ISV\[WD_RATE" /tmp/smoke-sp4.log | head -8
```

Expected (per spec pass criteria):
1. F0 ≥ 37.5, F1 > 0, F2 ≥ 55, all 3 folds complete 5 epochs
2. NaN-CLAMPED count = 0 (or matches ≤ 1 per regime-shift transient)
3. Diagnostic slots fire only on regime shifts (rare)
4. L1 λ starts low and ramps up within each fold
5. WD_RATE non-zero per group

- [ ] **Step 5: If pass, commit results doc**

Create `docs/superpowers/results/2026-04-30-sp4-validation.md` with the smoke output summary and per-criterion pass/fail.

- [ ] **Step 6: If fail, branch decision**

If F1 NaN'd: investigate Pearl D meta-α stability (perhaps too tight for transient regimes).
If F0 < 37.5: investigate Pearl A bootstrap behavior at fold 0 (cold-start might over-clamp).
If diagnostic slots fire heavily: investigate clamp value derivation (p99 might over-bind).

Each branch is a separate investigation thread, NOT a Layer-D iteration.

---

### Task C4: Update audit docs

**Files:**
- Modify: `docs/dqn-backward-nan-audit.md` (mechanism table + slot allocation)
- Modify: `docs/dqn-wire-up-audit.md` (Invariant 7 entry)

- [ ] **Step 1: Append SP4 mechanism table to audit doc**

Add after Phase F (SP3) section:

```markdown
## Phase G — SP4: Signal-Driven Magnitude Control (2026-04-30)

All hardcoded magnitude multipliers, EMA rates, and ε floors across
SP3 mechanisms (Mechs 1, 2, 5, 6, 9, 10) and 7 pre-SP4 ISV producers
replaced by signal-driven ISV bounds via 5 pearls.

### ISV slot allocation (40 new SP4 slots, indices 131-170)

| Slot range | Family | Producer | Consumer |
|---|---|---|---|
| 131 | TARGET_Q_BOUND | target_q_p99_update | Mech 1 + slot 46 diag |
| 132-135 | ATOM_POS_BOUND[0..4] | atom_pos_p99_update | Mech 2 + slot 47 diag |
| 136-143 | WEIGHT_BOUND[0..8] | param_group_oracle (Pass A) | Mech 9 + slots 44/45/48 diag |
| 144-151 | ADAM_M_BOUND[0..8] | param_group_oracle (Pass B) | slots 36-39 diag |
| 152-159 | ADAM_V_BOUND[0..8] | param_group_oracle (Pass C) | slots 40-43 diag |
| 160-167 | WD_RATE[0..8] | param_group_oracle (Pass D) | Adam weight_decay arg |
| 168 | GRAD_CLIP_BOUND | grad_norm_p99_update | Mech 6 |
| 169 | H_S2_BOUND | h_s2_p99_update | Mech 10 + slot 49 diag |
| 170 | L1_LAMBDA[trunk] | param_group_oracle (Pass E) | Adam l1_lambda arg |

### Pearls applied

- Pearl A: First-observation bootstrap (sentinel-detect at fold reset)
- Pearl B: Fused per-param-group statistics oracle (15 producers, not 40)
- Pearl C: Engagement-rate self-correction for in-kernel clamps
- Pearl D: Wiener-optimal adaptive α (replaces hardcoded EMA rates)
- Signal-driven L1 λ via gradient-direction entropy deficit (extends pearl_adaptive_moe_lambda)

### Validation

Smoke `smoke-test-XXXXX` (commit YYYY): F0=ZZ, F1=AA, F2=BB; slots 36-49
quiet outside regime-shift transients; L1 λ ramps from 0; WD_RATE
group-specific.
```

- [ ] **Step 2: Add Invariant 7 entry to wire-up audit**

In `docs/dqn-wire-up-audit.md`, prepend a comprehensive entry:

```markdown
SP4 — Signal-Driven Magnitude Control (2026-04-30): replaces ALL
hardcoded magnitude multipliers, EMA rates, and ε floors across SP3
mechanisms (Mechs 1, 2, 5, 6, 9, 10) and 7 pre-SP4 ISV producers.
40 new ISV slots [131..171), 15 fused producer kernels (Pearl B),
141-float wiener_state_buf (Pearl D), 2048-int clamp_engage_per_block_buf
(Pearl C), 5 pearls applied (A/B/C/D + L1 entropy-deficit signal-driven
regularization). [... full description ...]

Touched: cuda_pipeline/sp4_isv_slots.rs (new), sp4_wiener_ema.rs (new),
sp4_histogram_p99.cuh (new), target_q_p99_kernel.cu (new),
atom_pos_p99_kernel.cu (new), param_group_oracle_kernel.cu (new),
grad_norm_p99_kernel.cu (new), h_s2_p99_kernel.cu (new),
gpu_dqn_trainer.rs (40 ISV slots + buffers + 15 launchers + Layer B
consumer migrations), state_reset_registry.rs (181 reset entries),
training_loop.rs (10 cold-path launches), fused_training.rs (4 captured-graph
launches), atoms_update_kernel.cu / iql_value_kernel.cu /
experience_kernels.cu (Mech 2 ISV reads), dqn_utility_kernels.cu (Mech 5
fused diag ISV reads, Pearl C engagement counter, weight_decay/l1_lambda
ISV-source), iqn_dual_head_kernel.cu / iql_value_kernel.cu /
attention_backward_kernel.cu / curiosity_training_kernel.cu (Adam Pearl C
+ ISV-source weight_clamp/decay), decision_transformer.rs (DT 0.0 disable),
config.rs (HyperParams field removal), retired grad_norm_slow_ema, retrofit
of 7 existing producers (h_s2_rms_ema_kernel.cu, aux_heads_loss_ema_kernel.cu,
moe_kernels.cu, vsn_mask_ema_kernel.cu, iqn_quantile_ema_kernel.cu,
reward_component_ema_kernel.cu) with Pearls A+D. cargo check clean. L40S
smoke validation passed (smoke-test-XXXXX).
```

- [ ] **Step 3: Commit**

---

### Task C5: Memory entries (5 new pearls + SP4 close-out)

**Files:**
- Create: `/home/jgrusewski/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/pearl_first_observation_bootstrap.md`
- Create: `pearl_fused_per_group_statistics_oracle.md`
- Create: `pearl_engagement_rate_self_correction.md`
- Create: `pearl_wiener_optimal_adaptive_ema.md`
- Create: `pearl_signal_driven_regularisation_strength.md`
- Create: `project_sp4_close_out.md`
- Modify: `MEMORY.md` (add 6 pointer entries)

- [ ] **Step 1: Write each memory file using the auto-memory frontmatter format**

Example for Pearl A:

```markdown
---
name: pearl_first_observation_bootstrap
description: For any EMA-tracked signal, replace theoretical-init bootstrap with first-step-replacement at sentinel — eliminates magnitude-derived constants from cold-start. Subsumed by Pearl D as t=0 edge case but requires explicit sentinel-detection branch (Pearl D's formula at t=0 yields x_mean=0, not x[0]).
type: feedback
---
[content body]
```

- [ ] **Step 2: Add MEMORY.md pointers**

```markdown
- [pearl_first_observation_bootstrap.md](pearl_first_observation_bootstrap.md) — sentinel-detect bootstrap eliminates Xavier-derived magnitude constants for EMA-tracked signals
- [pearl_fused_per_group_statistics_oracle.md](pearl_fused_per_group_statistics_oracle.md) — fuse N producer kernels into 1 per-group when they share buffers + time-window
- [pearl_engagement_rate_self_correction.md](pearl_engagement_rate_self_correction.md) — for post-clamp/post-mutation EMA bounds, track engagement rate against theoretical p-quantile baseline; force-recalibrate on sustained deviation
- [pearl_wiener_optimal_adaptive_ema.md](pearl_wiener_optimal_adaptive_ema.md) — α* = diff_var / (diff_var + sample_var + ε_div) eliminates hardcoded EMA rates; recursion solved by uniform structural meta-α
- [pearl_signal_driven_regularisation_strength.md](pearl_signal_driven_regularisation_strength.md) — extends pearl_adaptive_moe_lambda: regularization-strength derives from gradient-direction entropy deficit
- [project_sp4_close_out.md](project_sp4_close_out.md) — SP4 closure 2026-04-30: 100% magnitude/regularization surface migrated to signal-driven ISV bounds
```

- [ ] **Step 3: Commit memory files**

```bash
cd /home/jgrusewski/.claude/projects/-home-jgrusewski-Work-foxhunt/memory/
# (no git here; auto-memory directory is a separate persistence path)
```

For the SP4 working tree, no commit needed for memory files (they're outside the repo).

---

## Final task: Self-review and handoff

After all Layer C tasks complete:

- [ ] Verify `cargo check -p ml --lib` clean
- [ ] Verify all 40 SP4 unit tests pass (+ retrofit unit tests)
- [ ] Verify L40S smoke pass criteria (Task C3)
- [ ] Verify audit-doc Invariant 7 entry comprehensive
- [ ] Verify memory entries written

Plan complete.
