From 611d2663cbd39a3b00a0c455ab742125e3b815b0 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Thu, 30 Apr 2026 23:45:21 +0200 Subject: [PATCH] =?UTF-8?q?feat(sp4):=20Task=20A8=20=E2=80=94=20grad=5Fnor?= =?UTF-8?q?m=5Fp99=20producer=20(degenerate=20single-element)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit grad_norm_buf is single-element f32; p99 == the element. Kernel copies to scratch slot 38 with __threadfence_system; host applies Pearls A+D to update ISV[GRAD_CLIP_BOUND_INDEX=168]. Trivial pattern — establishes the launcher template for Mech 6's eventual ISV read in Layer B. GPU test verifies first-observation Pearl A replacement (input=42.0 → scratch[38]=42.0 bit-identical → new_x_mean=42.0, x_lag=42.0, both variances stay zero) and that all non-target scratch slots remain 0. No consumer wired yet. Behaviour unchanged. cargo check --lib --tests clean. Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/build.rs | 11 ++ .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 155 ++++++++++++++++++ .../src/cuda_pipeline/grad_norm_p99_kernel.cu | 30 ++++ crates/ml/tests/sp4_producer_unit_tests.rs | 137 ++++++++++++++++ docs/dqn-wire-up-audit.md | 2 + 5 files changed, 335 insertions(+) create mode 100644 crates/ml/src/cuda_pipeline/grad_norm_p99_kernel.cu diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 8f522e00b..31a716579 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -285,6 +285,17 @@ fn main() { // tuples, then applies Pearls A+D host-side per output via // `pearls_ad_update`. "param_group_oracle_kernel.cu", + // SP4 Layer A Task A8 (2026-04-30): grad_norm_p99 producer + // (degenerate single-element). `grad_norm_buf` is a single-element + // f32 device buffer; p99 over 1 element IS the element itself, so + // the kernel reduces to a single-thread copy of `grad_norm_buf[0]` + // into `producer_step_scratch_buf[scratch_idx=38]` with + // `__threadfence_system()`. The host launcher + // (`GpuDqnTrainer::launch_sp4_grad_norm_p99`) syncs, applies Pearls + // A+D via `pearls_ad_update`, writes the new x_mean to + // ISV[GRAD_CLIP_BOUND_INDEX=168] + Wiener state. Establishes the + // launcher template that Mech 6's eventual Layer-B ISV read consumes. + "grad_norm_p99_kernel.cu", ]; // ALL kernels get common header (BF16 types + wrappers) diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 4cb71151b..e29b4d456 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -189,6 +189,18 @@ static SP4_ATOM_POS_P99_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), " static SP4_PARAM_GROUP_ORACLE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/param_group_oracle_kernel.cubin")); +/// SP4 Layer A Task A8 (2026-04-30): grad_norm_p99 producer cubin. +/// `grad_norm_buf` is a single-element f32 device buffer; p99 over 1 element +/// is the element itself, so the kernel reduces to a single-thread copy from +/// `grad_norm_buf[0]` into `producer_step_scratch_buf[scratch_idx]` with +/// `__threadfence_system()`. The launcher +/// (`GpuDqnTrainer::launch_sp4_grad_norm_p99`) syncs, applies Pearls A+D +/// (`pearls_ad_update`), and writes the new x_mean to +/// ISV[GRAD_CLIP_BOUND_INDEX=168] + Wiener state. No consumer wired yet — +/// Mech 6's adaptive_clip upper-bound consumer migrates in Layer B. Cold-path +/// launch in this task; Task A10 may move to the captured graph. +static SP4_GRAD_NORM_P99_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/grad_norm_p99_kernel.cubin")); + /// Plan 4 Task 3 (E.3): IQN multi-quantile diagnostic EMAs into ISV[99..103). /// 4-block (256 threads/block, shmem-reduce, no atomicAdd) kernel launched /// alongside `h_s2_rms_ema_update` from `training_loop.rs`. Reads the IQN @@ -3424,6 +3436,17 @@ pub struct GpuDqnTrainer { /// `SP4_PARAM_GROUP_ORACLE_CUBIN`. param_group_oracle_update: CudaFunction, + /// SP4 Layer A Task A8: grad_norm_p99 producer (degenerate single-element). + /// `grad_norm_buf` is a single-element f32 device buffer (populated by + /// `launch_grad_norm_finalize`). The kernel reduces to a single-thread + /// copy of `grad_norm_buf[0]` → `producer_step_scratch_buf[scratch_idx=38]` + /// with `__threadfence_system()`. The launcher + /// (`launch_sp4_grad_norm_p99`) syncs, applies Pearls A+D via + /// `pearls_ad_update` and writes back to ISV[GRAD_CLIP_BOUND_INDEX=168] + + /// Wiener state at `wiener_state_buf[(168-SP4_SLOT_BASE)*3 = 111..114)`. + /// Loaded from `SP4_GRAD_NORM_P99_CUBIN`. + grad_norm_p99_update: CudaFunction, + // ── Plan C Phase 2 follow-up A.2: q-drift rate ISV producer ─────── /// Single-thread single-block cold-path kernel. Reads ISV[Q_ABS_REF=16] + /// ISV[Q_DIR_ABS_REF=21] plus host-passed `q_mean_curr` / `q_mean_prev` @@ -9158,6 +9181,122 @@ impl GpuDqnTrainer { } } + /// SP4 Layer A Task A8: cold-path producer for ISV[GRAD_CLIP_BOUND_INDEX=168]. + /// + /// `grad_norm_buf` is a single-element f32 device buffer populated by + /// `launch_grad_norm_finalize` with the current step's L2 grad norm. + /// p99 over a single element is the element itself — the kernel + /// (`grad_norm_p99_update`) reduces to a single-thread copy of + /// `grad_norm_buf[0]` → `producer_step_scratch_buf[SCRATCH_IDX=38]` with + /// `__threadfence_system()`. No histogram pass is needed; the launcher + /// keeps the same shape as `launch_sp4_target_q_p99` for consistency + /// across SP4 producers, and Mech 6's eventual Layer-B ISV read in the + /// adaptive_clip path consumes ISV[168] directly. + /// + /// Synchronises the stream then applies Pearls A+D host-side via + /// `pearls_ad_update`, using zero-copy mapped-pinned reads of ISV[168] + /// + the per-slot Wiener state at `wiener_state_buf[(168 - + /// SP4_SLOT_BASE) * 3 = 111..114)`. Writes the new `x_mean` back to + /// ISV[168] and the updated Wiener state back to `wiener_state_buf` via + /// the same mapped-pinned host pointers — no HtoD/DtoH per + /// `feedback_no_htod_htoh_only_mapped_pinned`. + /// + /// **Cold-path launch** — NOT in the captured graph for now. Task A10 may + /// relocate to captured-graph cadence; this launcher mirrors A5's shape + /// so the migration is structural-only at that point. + /// + /// **No consumer wired in this task** — Mech 6's `adaptive_clip` upper + /// bound still uses the existing fast/slow grad-norm EMA path; SP4 + /// consumer migration follows once all producers (A5-A9) land. + pub fn launch_sp4_grad_norm_p99(&self) -> Result<(), MLError> { + use crate::cuda_pipeline::sp4_isv_slots::{GRAD_CLIP_BOUND_INDEX, SP4_SLOT_BASE}; + use crate::cuda_pipeline::sp4_wiener_ema::{pearls_ad_update, WienerState}; + + // Producer-step-scratch slot 38 is reserved for GRAD_CLIP_BOUND per + // the stable layout documented on `launch_sp4_target_q_p99`: + // 0 = TARGET_Q_BOUND (Task A5) + // 1..5 = ATOM_POS_BOUND[0..4] (Task A6) + // 5..29 = WEIGHT/ADAM_M/ADAM_V × 8 (Task A7) + // 29..37 = WD_RATE × 8 (Task A7) + // 37 = L1_LAMBDA_TRUNK (Task A7) + // 38 = GRAD_CLIP_BOUND (this task) + // 39 = H_S2_BOUND (Task A9) + // 40..47 = retrofit existing 7 producers + const SCRATCH_IDX: usize = 38; + + // Single-thread degenerate kernel — minimal launch config, no shared + // memory. Mirrors the `q_drift_rate_ema_kernel.cu` / + // `moe_lambda_eff_kernel.cu` cold-path single-thread footprint. + let cfg = LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }; + + let grad_norm_ptr = self.ptrs.grad_norm_buf; + let scratch_dev = self.producer_step_scratch_buf.dev_ptr; + let scratch_idx_arg: i32 = SCRATCH_IDX as i32; + + // Safety: kernel signature `(const float*, float*, int)` matches + // the three args below; both device pointers are valid for the + // lifetime of `self`. + unsafe { + self.stream + .launch_builder(&self.grad_norm_p99_update) + .arg(&grad_norm_ptr) + .arg(&scratch_dev) + .arg(&scratch_idx_arg) + .launch(cfg) + .map_err(|e| MLError::ModelError(format!("grad_norm_p99_update launch: {e}")))?; + } + // Cold-path producer — sync before reading the mapped-pinned scratch + // slot from the host. Captured-graph migration (Task A10) replaces + // this with the same fence the captured stream provides on graph + // launch boundaries. + self.stream.synchronize() + .map_err(|e| MLError::ModelError(format!("grad_norm_p99_update sync: {e}")))?; + + // ── Pearls A+D host-side update (zero-copy mapped-pinned reads) ── + let step_obs = unsafe { + std::ptr::read_volatile(self.producer_step_scratch_buf.host_ptr.add(SCRATCH_IDX)) + }; + // Degenerate-signal short-circuit: kernel writes 0.0 if grad_norm + // hasn't been populated yet (e.g., before first backward). Avoid + // advancing Wiener state with a zero observation, which would + // suppress the next genuine first-observation. Same semantics as + // Tasks A5/A6/A7. + if step_obs == 0.0 { return Ok(()); } + + let isv_idx = GRAD_CLIP_BOUND_INDEX; + let wiener_offset = (isv_idx - SP4_SLOT_BASE) * 3; + + // Safety: `isv_signals_pinned` is a mapped-pinned `*mut f32` of + // length `ISV_TOTAL_DIM`; `isv_idx < ISV_TOTAL_DIM` (171 > 168). + // `wiener_state_buf.host_ptr` is a mapped-pinned `*mut f32` of + // length `SP4_PRODUCER_COUNT * 3 = 141`; `wiener_offset + 2 < 141`. + let prev_x_mean = unsafe { std::ptr::read_volatile(self.isv_signals_pinned.add(isv_idx)) }; + + let mut state = unsafe { + let p = self.wiener_state_buf.host_ptr; + WienerState { + sample_var: std::ptr::read_volatile(p.add(wiener_offset)), + diff_var: std::ptr::read_volatile(p.add(wiener_offset + 1)), + x_lag: std::ptr::read_volatile(p.add(wiener_offset + 2)), + } + }; + + let new_x_mean = pearls_ad_update(prev_x_mean, &mut state, step_obs); + + unsafe { + std::ptr::write_volatile(self.isv_signals_pinned.add(isv_idx), new_x_mean); + let wp = self.wiener_state_buf.host_ptr; + std::ptr::write_volatile(wp.add(wiener_offset), state.sample_var); + std::ptr::write_volatile(wp.add(wiener_offset + 1), state.diff_var); + std::ptr::write_volatile(wp.add(wiener_offset + 2), state.x_lag); + } + Ok(()) + } + /// Plan C Phase 2 follow-up A.2 (2026-04-29): launch `q_drift_rate_ema_update` /// — single-thread single-block ISV producer for ISV[Q_DRIFT_RATE_INDEX=129]. /// @@ -10809,6 +10948,21 @@ impl GpuDqnTrainer { .map_err(|e| MLError::ModelError(format!("param_group_oracle_update load: {e}")))? }; + // SP4 Layer A Task A8: load grad_norm_p99 producer kernel (cold-path). + // Single-thread single-block kernel that copies `grad_norm_buf[0]` + // (single-element scalar populated by `launch_grad_norm_finalize`) to + // `producer_step_scratch_buf[38]` with `__threadfence_system()`. + // p99 over a single element IS the element itself, hence the + // degenerate-copy form (no histogram). Pearls A+D applied host-side + // via `pearls_ad_update` (Task A3). Producer-only — Mech 6's + // adaptive_clip upper bound consumer migrates in Layer B. + let grad_norm_p99_update = { + let module = stream.context().load_cubin(SP4_GRAD_NORM_P99_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("sp4 grad_norm_p99 cubin load: {e}")))?; + module.load_function("grad_norm_p99_update") + .map_err(|e| MLError::ModelError(format!("grad_norm_p99_update load: {e}")))? + }; + // Plan C Phase 2 follow-up A.2: load q_drift_rate_ema kernel // (cold-path, per-epoch). Single-thread single-block ISV producer // for ISV[Q_DRIFT_RATE_INDEX=129]; consumer is `tau_update_kernel`. @@ -13770,6 +13924,7 @@ impl GpuDqnTrainer { target_q_p99_update, atom_pos_p99_update, param_group_oracle_update, + grad_norm_p99_update, q_drift_rate_ema_kernel, fold_warmup_factor_kernel, grad_norm_fast_ema_pinned, diff --git a/crates/ml/src/cuda_pipeline/grad_norm_p99_kernel.cu b/crates/ml/src/cuda_pipeline/grad_norm_p99_kernel.cu new file mode 100644 index 000000000..df664070f --- /dev/null +++ b/crates/ml/src/cuda_pipeline/grad_norm_p99_kernel.cu @@ -0,0 +1,30 @@ +// crates/ml/src/cuda_pipeline/grad_norm_p99_kernel.cu +// +// SP4 Layer A Task A8 — grad_norm_p99 producer (degenerate single-element). +// +// `grad_norm_buf` is a single-element f32 device buffer containing the current +// step's L2 grad norm (populated by `launch_grad_norm_finalize`). p99 over a +// single element IS the element itself, so the "histogram" reduces to a copy: +// the single-thread kernel reads `grad_norm_buf[0]` and writes it to +// `producer_step_scratch_buf[scratch_idx]` with `__threadfence_system()` so +// the host pinned-mapped read sees the write. +// +// The host launcher (`GpuDqnTrainer::launch_sp4_grad_norm_p99`) syncs, applies +// Pearls A+D via `pearls_ad_update`, and writes the new x_mean to +// ISV[GRAD_CLIP_BOUND_INDEX=168] + Wiener state at slot 38 of +// `wiener_state_buf` (offset `(168 - SP4_SLOT_BASE) * 3 = 111`). +// +// Cold-path launch — NOT in the captured graph for now. Task A10 may relocate. +// Task A8 establishes the launcher template that Mech 6's eventual ISV read +// in Layer B will consume. Mirrors Task A5's launcher shape end-to-end with +// the kernel itself reduced to a copy. + +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]; + __threadfence_system(); // make host-mapped write visible +} diff --git a/crates/ml/tests/sp4_producer_unit_tests.rs b/crates/ml/tests/sp4_producer_unit_tests.rs index 38160762f..3ba7dbb45 100644 --- a/crates/ml/tests/sp4_producer_unit_tests.rs +++ b/crates/ml/tests/sp4_producer_unit_tests.rs @@ -998,3 +998,140 @@ fn sp4_param_group_oracle_per_group_writes_distinct_isv_slots() { assert!(wd_rate(g) >= 160 && wd_rate(g) < 168); } } + +// ── SP4 Task A8: grad_norm_p99 producer (degenerate single-element) ─────────── + +/// Test-only cubin for the SP4 grad_norm_p99 producer kernel. The same cubin +/// is loaded by `GpuDqnTrainer::new` in production via `SP4_GRAD_NORM_P99_CUBIN`; +/// this kernel-direct test bypasses the full trainer harness and exercises +/// the kernel's degenerate single-thread copy contract. +const SP4_GRAD_NORM_P99_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/grad_norm_p99_kernel.cubin")); + +/// Resolve the `grad_norm_p99_update` function handle from its cubin. +fn load_sp4_grad_norm_p99_kernel(stream: &Arc) -> CudaFunction { + let module = stream + .context() + .load_cubin(SP4_GRAD_NORM_P99_CUBIN.to_vec()) + .expect("load grad_norm_p99_kernel cubin"); + module + .load_function("grad_norm_p99_update") + .expect("load grad_norm_p99_update function") +} + +/// SP4 Task A8 unit test: the production `grad_norm_p99_update` kernel +/// performs a single-thread copy of `grad_norm_buf[0]` to +/// `producer_step_scratch_buf[scratch_idx=38]` and the host-side +/// `pearls_ad_update` helper (Task A3) replaces the ISV sentinel directly +/// with that step observation on the first call after a fold-boundary +/// reset (Pearl A bootstrap). +/// +/// `grad_norm_buf` is a single-element f32 device buffer in production; +/// p99 over 1 element IS the element itself, so the kernel is degenerate +/// (no histogram). This test sets that single element to a known scalar +/// (42.0) and verifies the kernel writes exactly that value to slot 38 +/// while leaving every other slot zero. +/// +/// Production launcher contract validated: +/// - Kernel writes the input scalar to `scratch[SCRATCH_IDX_GRAD_CLIP=38]`. +/// - All other scratch slots remain zero (helper assertion guards the +/// `2fb30f098`-style write-cascade pattern). +/// - `GRAD_CLIP_BOUND_INDEX == 168` and `(168 - SP4_SLOT_BASE) * 3 == 111` +/// (Wiener state offset for this slot). +#[test] +#[ignore = "requires GPU"] +fn sp4_grad_norm_p99_pearl_a_first_observation_replaces_scalar() { + use ml::cuda_pipeline::sp4_isv_slots::{GRAD_CLIP_BOUND_INDEX, SP4_SLOT_BASE}; + use ml::cuda_pipeline::sp4_wiener_ema::{pearls_ad_update, WienerState}; + + // Production scratch shape (47 = SP4 producers count). Slot 38 is the + // production-allocated GRAD_CLIP_BOUND slot (per the documented stable + // layout in `launch_sp4_target_q_p99` / `launch_sp4_grad_norm_p99`). + const SP4_PRODUCER_COUNT: usize = 47; + const SCRATCH_IDX_GRAD_CLIP: usize = 38; + + // Single-element grad_norm scalar — known synthetic value chosen to be + // distinguishable from the Pearl-A sentinel (0.0) and from a stale + // initialisation (1.0) so a "did the kernel actually copy?" question + // is answered unambiguously. + let grad_norm_scalar: f32 = 42.0; + + let stream = make_test_stream(); + let kernel = load_sp4_grad_norm_p99_kernel(&stream); + + // Safety: CUDA context active on this thread (resolved via stream). + let in_buf = unsafe { MappedF32Buffer::new(1) } + .expect("alloc MappedF32Buffer for SP4 grad_norm_p99 input"); + in_buf.write_from_slice(&[grad_norm_scalar]); + + let scratch_buf = unsafe { MappedF32Buffer::new(SP4_PRODUCER_COUNT) } + .expect("alloc MappedF32Buffer for SP4 producer scratch"); + + let scratch_idx_arg: i32 = SCRATCH_IDX_GRAD_CLIP as i32; + let in_dev_ptr = in_buf.dev_ptr; + let scratch_dev_ptr = scratch_buf.dev_ptr; + + // Single-thread degenerate kernel — minimal launch config, no shared + // memory. Mirrors the production `launch_sp4_grad_norm_p99` config. + unsafe { + stream + .launch_builder(&kernel) + .arg(&in_dev_ptr) + .arg(&scratch_dev_ptr) + .arg(&scratch_idx_arg) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }) + .expect("launch grad_norm_p99_update"); + } + stream + .synchronize() + .expect("synchronize after grad_norm_p99_update launch"); + + // Verify ALL non-target slots remained zero — the kernel must only + // touch the slot at `scratch_idx=38`. + let host = scratch_buf.read_all(); + for (i, &v) in host.iter().enumerate() { + if i != SCRATCH_IDX_GRAD_CLIP { + assert_eq!( + v, 0.0, + "grad_norm_p99 producer wrote to scratch[{i}] outside target slot {SCRATCH_IDX_GRAD_CLIP}: {v}", + ); + } + } + let step_obs = host[SCRATCH_IDX_GRAD_CLIP]; + println!( + "SP4 grad_norm_p99 — input={grad_norm_scalar:.5}, step_obs={step_obs:.5}", + ); + // Degenerate copy — bit-identical equality (no histogram quantization). + assert_eq!( + step_obs, grad_norm_scalar, + "kernel must copy grad_norm_buf[0] verbatim to scratch[{SCRATCH_IDX_GRAD_CLIP}]; got {step_obs} vs {grad_norm_scalar}", + ); + + // ── Pearl A bootstrap (host-side, mirrors `launch_sp4_grad_norm_p99`) ── + // Sentinel ISV slot + zero Wiener state — replicates the cold-start + // condition immediately after a fold-boundary reset. + let prev_x_mean: f32 = 0.0; + let mut state = WienerState::ZERO; + let new_x_mean = pearls_ad_update(prev_x_mean, &mut state, step_obs); + + assert_eq!( + new_x_mean, step_obs, + "Pearl A: first observation must replace sentinel directly; got {new_x_mean} vs step_obs={step_obs}", + ); + assert_eq!( + state.x_lag, step_obs, + "Pearl A: x_lag must seed to first observation", + ); + assert_eq!(state.sample_var, 0.0, "Pearl A: sample_var stays zero"); + assert_eq!(state.diff_var, 0.0, "Pearl A: diff_var stays zero"); + + // Slot-layout invariants — guard against off-by-one drift in + // GRAD_CLIP_BOUND_INDEX or SP4_SLOT_BASE. + assert_eq!(GRAD_CLIP_BOUND_INDEX, 168); + assert_eq!(SP4_SLOT_BASE, 131); + assert_eq!((GRAD_CLIP_BOUND_INDEX - SP4_SLOT_BASE) * 3, 111); +} diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index db6778126..994ff72ec 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2303,6 +2303,8 @@ SP4 Layer A Task A2 — mapped-pinned buffers for Pearls B/C/D (2026-04-30): all SP4 Layer A Task A4 — linear-histogram p99 device function + GPU unit test (2026-04-30): created `crates/ml/src/cuda_pipeline/sp4_histogram_p99.cuh` providing the header-only `__device__` template `sp4_histogram_p99(buf, count) -> float` that returns p99(|buf|) for a single-block, BLOCK_SIZE-thread launch. Three-pass algorithm: (1) block-wide max-reduce of |buf| → `step_max` (degenerate-zero short-circuits to 0.0 so callers skip the ISV update); (2) linear-spaced [0, step_max] binning into 256 bins via per-warp tiles in dynamic shared memory (no atomicAdd per `feedback_no_atomicadd` — lane-collisions within a warp cost <0.012% expected count loss, well within the 1% quantile precision budget; cross-warp collisions are eliminated by the per-warp tiling); (3) cumulative-from-top → p99 = bin upper-edge `(p99_bin + 1) × bin_width`. Linear spacing chosen over log because SP4 producers care about resolution at the *top* of the |buf| distribution; top bin width ≈ step_max/256 ≈ 0.39%, comfortably within the 1% quantile precision budget. Caller contract documented in the header: `grid_dim=(1,1,1)`, `block_dim=(BLOCK_SIZE,1,1)`, dynamic shared memory ≥ `(BLOCK_SIZE/32) × 256 × sizeof(int)` (8192 bytes for BLOCK_SIZE=256). Wrapper kernel `sp4_histogram_p99_test_kernel.cu` exposes the device fn for Rust testing — single-block kernel that calls `sp4_histogram_p99<256>` and writes the scalar result to a mapped-pinned `f32` with `__threadfence_system()` so the host reads via `read_volatile` (no `memcpy_dtoh` per `feedback_gpu_cpu_roundtrip` and `feedback_no_htod_htoh_only_mapped_pinned`). Build.rs registration follows the `thompson_test_kernel.cu` pattern (Plan A Task 1): added to `kernels_with_common`, plus a `cargo:rerun-if-changed` directive on the `.cuh` header so cubins rebuild when the device fn evolves. Unit test `sp4_histogram_p99_matches_known_distribution_within_quantization` (in `crates/ml/tests/sp4_producer_unit_tests.rs`) drives the wrapper with 4096 deterministic |N(0,1)| Box-Muller samples (`StdRng::seed_from_u64(0xDEAD_BEEF)`), computes ground-truth p99 by sorting (analytical reference ≈ 2.576 z-score one-tailed), and asserts `rel_err < 5%` against the device output. `#[ignore]`-gated for GPU per the existing `distributional_q_tests.rs` convention. Local L40S run: true_p99=2.59758, computed_p99=2.59858, rel_err=0.039% — passes. The 5% tolerance leaves ample headroom over the worst-case sum of (linear-bin quantization 0.4%) + (per-warp lane-collision <0.012%) + (finite-sample sorted-p99 jitter ≈0.5% at n=4096) ≈ <2% true budget. No producer wired yet — header is library code included only by the test wrapper; SP4 Tasks A5-A9 add the magnitude-bound producer kernels that `#include "sp4_histogram_p99.cuh"` directly. Zero new ISV slots; zero new HtoD/DtoD/HtoH copies; zero behaviour change. `cargo check -p ml --lib --tests` clean (12 pre-existing warnings, no new warnings); GPU test 1/1 passing locally. +SP4 Layer A Task A8 — `grad_norm_p99` producer (degenerate single-element) (2026-04-30): trivial cold-path producer establishing the Mech 6 launcher template that Layer B's adaptive_clip upper-bound consumer will eventually read. Created `crates/ml/src/cuda_pipeline/grad_norm_p99_kernel.cu` — single-thread single-block kernel that copies `grad_norm_buf[0]` (a single-element f32 device buffer populated by `launch_grad_norm_finalize`) to `producer_step_scratch_buf[scratch_idx=38]` with `__threadfence_system()`. p99 over 1 element IS the element itself, so no histogram pass is needed; the kernel is a degenerate copy. Registered in `crates/ml/build.rs` after `param_group_oracle_kernel.cu`. Cubin embedded as `SP4_GRAD_NORM_P99_CUBIN`; `grad_norm_p99_update: CudaFunction` field added next to `param_group_oracle_update`; loaded in the constructor immediately after `param_group_oracle_update`. New launcher `GpuDqnTrainer::launch_sp4_grad_norm_p99(&self) -> Result<(), MLError>` mirrors Task A5's shape end-to-end with the kernel itself reduced to a copy: `grid_dim=(1,1,1)`, `block_dim=(1,1,1)`, `shared_mem_bytes=0` (mirroring `q_drift_rate_ema_kernel` / `moe_lambda_eff_kernel` cold-path footprint); kernel arg list `(grad_norm_buf_ptr, scratch_dev_ptr, scratch_idx=38)`; `stream.synchronize()`; `read_volatile(producer_step_scratch_buf.host_ptr.add(38))`; degenerate-zero short-circuit before mutating ISV/Wiener state; `pearls_ad_update` (Task A3) host-side using zero-copy mapped-pinned reads of ISV[GRAD_CLIP_BOUND_INDEX=168] + Wiener triple at `wiener_state_buf.host_ptr.add((168-SP4_SLOT_BASE)*3 = 111)`; writes new `x_mean` to ISV[168] and updated `[sample_var, diff_var, x_lag]` to `wiener_state_buf` — all via mapped-pinned host pointers (no HtoD/DtoH per `feedback_no_htod_htoh_only_mapped_pinned`). **Cold-path launch** — NOT in the captured graph for now; Task A10 may relocate to captured-graph cadence. **No consumer wired yet** — Mech 6's `adaptive_clip` upper bound still uses the existing fast/slow grad-norm EMA path; SP4 consumer migration follows once all producers (A5-A9) land. Behaviour unchanged from before this commit. **Unit test** `sp4_grad_norm_p99_pearl_a_first_observation_replaces_scalar` (in `crates/ml/tests/sp4_producer_unit_tests.rs`) drives the production kernel kernel-direct with a known single-element scalar (42.0) chosen to be distinguishable from both the Pearl-A sentinel (0.0) and a stale init (1.0); asserts `scratch[38]` is bit-identical to the input (no histogram quantization for a degenerate copy), then exercises Pearl A's sentinel branch (`prev=0.0`, `state=ZERO`) and asserts `new_x_mean == step_obs`, `state.x_lag == step_obs`, `state.sample_var == 0.0`, `state.diff_var == 0.0`. Helper assertion guards every non-target slot ∈ {0..38} ∪ {39..47} stays zero (guards against the `2fb30f098`-style write-cascade). Layout invariants asserted directly: `GRAD_CLIP_BOUND_INDEX=168`, `SP4_SLOT_BASE=131`, `(168-SP4_SLOT_BASE)*3=111`. `#[ignore]`-gated for GPU. Zero new ISV slots (GRAD_CLIP_BOUND_INDEX=168 already allocated by Task A1); zero new HtoD/DtoD/HtoH copies (mapped-pinned only); one new kernel + one new launcher method; producer-only (no consumer). `cargo check -p ml --lib --tests` clean (12 pre-existing warnings, no new warnings). + SP4 Layer A Task A7 — Pearl B fused per-param-group statistics oracle (2026-04-30): single producer kernel per param-group, four-to-five fused passes per launch — `(params, grads, adam_m, adam_v)` read once per group, four scalar outputs (five for trunk) per launch. Created `crates/ml/src/cuda_pipeline/param_group_oracle_kernel.cu` — single-block, 256-thread kernel that runs Pass A (`WEIGHT_BOUND[g] = p99(|params|)` via `sp4_histogram_p99<256>`), Pass B (`ADAM_M_BOUND[g] = p99(|adam_m|)`), Pass C (`ADAM_V_BOUND[g] = p99(|adam_v|)`), Pass D (`WD_RATE[g] = |Σ w·g| / max(Σ w², EPS_DIV)` via 4 register-accumulator block-tree-reduces sharing one 256-float reuse of the histogram dynamic shmem region; side products `mean|g|`, `mean|w|` stashed in static `__shared__` for Pass E), and Pass E (group 0 trunk only — `L1_LAMBDA_TRUNK = (mean|g| / max(mean|w|, EPS_DIV)) × (log K − H) / log K` where H is the Shannon entropy of per-feature gradient L2 norms across the `[h_dim × k_in]` trunk weight gradient matrix; gated by `l1_lambda_scratch_idx >= 0`). Single helper `block_reduce_sum_256` (no `atomicAdd` per `feedback_no_atomicadd`); per-feature L2 norm uses cooperative thread-strided accumulation with sequential `block_reduce_sum_256` per feature to avoid atomicAdd while keeping cost linear in `k_in × h_dim`. `__threadfence_system()` after each scalar writeback so the host-mapped scratch slot is visible across the kernel/host boundary. `EPS_DIV = 1e-8f` matches `sp4_wiener_ema.rs::EPS_DIV` (Task A3). `K_IN_MAX = 1024` cap on Pass-E per-feature shared scratch (production trunk K_in is much smaller — 40 base + OFI/TLOB widening <128). Registered in `crates/ml/build.rs` after `atom_pos_p99_kernel.cu`. Cubin embedded as `SP4_PARAM_GROUP_ORACLE_CUBIN` in `gpu_dqn_trainer.rs`; `param_group_oracle_update: CudaFunction` field added next to `atom_pos_p99_update`; loaded in the constructor immediately after `atom_pos_p99_update`. **Pearl B 4× memory-bandwidth reduction** vs four naive single-stat producer kernels: each of `params`, `grads`, `adam_m`, `adam_v` is read once across all 4-5 outputs per group rather than four times. New launcher `GpuDqnTrainer::launch_sp4_param_group_oracles_all_groups(&self) -> Result<(), MLError>` loops `g_idx ∈ 0..SP4_PARAM_GROUP_COUNT=8`, calling `param_group_buffers(g)` to resolve `(params_ptr, grads_ptr, m_ptr, v_ptr, count, k_in, h_dim)` per group, launching the kernel once per group with distinct producer-step-scratch slot indices into the documented stable layout (slot `5 + g` for WEIGHT, `13 + g` for ADAM_M, `21 + g` for ADAM_V, `29 + g` for WD_RATE, slot 37 for L1_LAMBDA_TRUNK with `l1_idx = -1` for non-trunk groups). After a single end-of-loop `stream.synchronize()`, applies Pearls A+D host-side per output via `pearls_ad_update` (Task A3) — 4 outputs per group plus L1 for group 0, totalling 33 host-side updates when all groups launch (degenerate-zero short-circuit per output, mirroring Tasks A5/A6). All ISV reads/writes through `isv_signals_pinned` (mapped-pinned `*mut f32`, slots `weight_bound(g)`, `adam_m_bound(g)`, `adam_v_bound(g)`, `wd_rate(g)`, `L1_LAMBDA_TRUNK_INDEX`); all Wiener-state reads/writes through `wiener_state_buf.host_ptr` at `(isv_idx − SP4_SLOT_BASE) × 3`. Zero HtoD/DtoH per `feedback_no_htod_htoh_only_mapped_pinned`. **Three of eight groups wired** in this commit: `ParamGroup::DqnTrunk` (tensors [0..13), with `k_in = param_sizes[0] / shared_h1`, `h_dim = shared_h1` for Pass E), `ParamGroup::DqnValue` (tensors [13..17), Pass E gated off), `ParamGroup::DqnBranches` (tensors [17..33), Pass E gated off). All three slices derived from `compute_param_sizes` + `padded_byte_offset` + `f32` size, mirroring the existing `trunk_adam_m_ptr` / `value_adam_m_ptr` / `branch_adam_m_ptr` accessor offsets. **Five aux groups deferred**: `ParamGroup::Iqn`, `IqlHigh`, `IqlLow`, `Attn`, `Curiosity` — `param_group_buffers` returns `None`; the launcher silently skips. These trainers live on `FusedTrainingCtx`, not `GpuDqnTrainer`; wiring them requires either threading the buffer pointers through the launcher signature or hoisting the launcher onto `FusedTrainingCtx`. Both are follow-up changes scoped beyond Task A7 and documented in the launcher docstring + `param_group_buffers` docstring. IQN already has `online_params_ptr`/`online_params_len`/`adam_m_ptr`/`adam_v_ptr` accessors (added by SP3 close-out v2 for slot 48); it would just need an `online_grad_ptr`/`online_grad_len` pair. IQL/Attn/Curiosity will need full accessor sets. **Cold-path launch** — NOT in the captured graph for now; Task A10 may relocate select producers to the captured-graph cadence. **No consumer wired yet** — Mech 9's clamp still uses hardcoded `weight_clamp_max_abs = 100 × q_abs_ref_eff` config args; AdamW weight_decay and L1 unchanged. SP4 consumer migration follows once all producers (A5-A9) land. **Unit test** `sp4_param_group_oracle_per_group_writes_distinct_isv_slots` (in `crates/ml/tests/sp4_producer_unit_tests.rs`) drives the production kernel kernel-direct across all 8 group shapes (parameterised over `g_idx ∈ 0..SP4_PARAM_GROUP_COUNT`); each iteration generates 4 × 4096 deterministic |N(0, σ²)| Box-Muller samples for `(params, grads_mag, adam_m, adam_v)` with per-group sigmas spanning `0.001×..0.1×` and per-group seeds `0xA7_5A_F1_57 ^ (g_idx << 32) ^ buf_id` (4 distinct seeds per group → 32 buffers total). Grads are signed (alternating sign by index parity) so `Σ w·g` doesn't trivially collapse to `Σ w·|g|`. For group 0 the `(k_in, h_dim) = (64, 64)` shape exercises Pass E; other groups pass `(0, 0)` and the kernel skips Pass E via `l1_idx = -1`. Asserts: WEIGHT/ADAM_M/ADAM_V p99 within 5% rel_err of sorted-sample p99 (linear-bin quantization 0.4% + lane collision <0.012% + finite-sample jitter ≈0.5% ≈ <2% true budget); WD_RATE within 2% rel_err of analytical f64 reference `|Σ w·g| / max(Σ w², EPS_DIV)`; L1 (group 0) within 5% rel_err of analytical reference `(mean|g| / max(mean|w|, EPS_DIV)) × deficit` with deficit computed from per-feature L2-norm Shannon entropy. All non-target scratch slots stay zero (guards against the `2fb30f098`-class write-cascade). Layout invariants asserted directly: `SP4_PARAM_GROUP_COUNT=8`, `SP4_SLOT_BASE=131`, `WEIGHT_BOUND_BASE=136`, `ADAM_M_BOUND_BASE=144`, `ADAM_V_BOUND_BASE=152`, `WD_RATE_BASE=160`, `L1_LAMBDA_TRUNK_INDEX=170`, plus `weight_bound(g)`/`adam_m_bound(g)`/`adam_v_bound(g)`/`wd_rate(g)` const-fn drift guards. `#[ignore]`-gated for GPU. Local RTX 3050 Ti run, all 8 groups: maximum WEIGHT p99 rel_err = 0.589% (group 5); maximum ADAM_M p99 rel_err = 0.589% (group 6); maximum ADAM_V p99 rel_err = 0.554% (group 7); maximum WD_RATE rel_err = 0.001% (analytical formula reduces in f32 with no histogram quantization, hence sub-1%); group 0 L1 = 9.27e-5 vs reference 9.27e-5 (deficit=0.00092 reflects near-uniform per-feature half-normal distribution → small L1 by design). All assertions pass with substantial margin under tolerance. Zero new ISV slots (`SP4_PARAM_GROUP_COUNT=8` × 4 + 1 = 33 slots all allocated by Task A1 in [136..168) ∪ [170]); zero new HtoD/DtoD/HtoH copies (mapped-pinned only); one new kernel + one new launcher method + one new private accessor (`param_group_buffers`); producer-only (no consumer). `cargo check -p ml --lib --tests` clean (12 pre-existing warnings, no new warnings); 4/4 SP4 producer GPU tests passing locally (Task A4 + A5 + A6 + A7). SP4 Layer A Task A6 — `atom_pos_p99` producer × 4 branches (2026-04-30): single parameterised producer kernel + 4-branch launcher loop, mirroring Task A5's end-to-end pattern. Created `crates/ml/src/cuda_pipeline/atom_pos_p99_kernel.cu` — single-block, 256-thread kernel that `#include "sp4_histogram_p99.cuh"` directly (Task A4 header), reads ONE branch's slice of `atom_positions_buf` (canonical layout `[4 × num_atoms]` flat with branch `b` at offset `b × num_atoms` per `atoms_update_kernel.cu`'s `positions_out + (long long)branch * num_atoms` write contract), computes `p99(|atom_positions[branch]|)` via `sp4_histogram_p99<256>`, and writes the per-step observation to `producer_step_scratch_buf[scratch_idx]` with `__threadfence_system()`. Registered in `crates/ml/build.rs` after `target_q_p99_kernel.cu`. Cubin embedded as `SP4_ATOM_POS_P99_CUBIN`; `atom_pos_p99_update: CudaFunction` field added next to `target_q_p99_update`; loaded in the constructor immediately after `target_q_p99_update`. New launcher `GpuDqnTrainer::launch_sp4_atom_pos_p99_all_branches(&self) -> Result<(), MLError>` loops `branch ∈ 0..SP4_BRANCH_COUNT=4`, launching the same kernel four times with distinct `(branch_dev_ptr, scratch_idx)` pairs — `branch_dev_ptr = atom_positions_buf.raw_ptr() + branch × num_atoms × sizeof(f32)`, `scratch_idx ∈ {1, 2, 3, 4}` (SCRATCH_BASE=1 per Task A5's documented stable layout). After a single end-of-loop `stream.synchronize()`, applies Pearls A+D host-side per branch via `pearls_ad_update` (Task A3), reading `prev_x_mean` from `isv_signals_pinned.add(atom_pos_bound(branch) ∈ {132, 133, 134, 135})` and the per-slot Wiener triple from `wiener_state_buf.host_ptr.add((isv_idx − SP4_SLOT_BASE) × 3)`, writing the new `x_mean` back to ISV[ATOM_POS_BOUND_BASE + branch] and the updated `[sample_var, diff_var, x_lag]` back to `wiener_state_buf` — all via mapped-pinned host pointers (no HtoD/DtoH per `feedback_no_htod_htoh_only_mapped_pinned`). `debug_assert_eq!(total_len, SP4_BRANCH_COUNT × num_atoms)` guards against future `atom_positions_buf` layout drift; `debug_assert_eq!(isv_idx, ATOM_POS_BOUND_BASE + branch)` guards against `atom_pos_bound` const-fn drift. Same degenerate-zero short-circuit (`if step_obs == 0.0 { continue; }`) per branch as Task A5 — skips the ISV/Wiener update before the first `recompute_atom_positions` populates the branch slice. **Cold-path launch** — NOT in the captured graph for now; Task A10 may relocate atom_pos to captured-graph cadence. **No consumer wired yet** — Mech 2's atom-position clamp in `atoms_update_kernel.cu` still uses `±10 × ISV[Q_ABS_REF=16].max(1.0)`; SP4 consumer migration follows once all producers (A5-A9) land. Behaviour unchanged from before this commit. **Unit test** `sp4_atom_pos_p99_per_branch_writes_distinct_isv_slots` (in `crates/ml/tests/sp4_producer_unit_tests.rs`) drives the production kernel kernel-direct with 4 × 4096 deterministic |N(0,1)| Box-Muller samples (`StdRng::seed_from_u64(0xA6_5A_F1_57 ^ (branch << 32))`), each branch scaled by `10ⁿ` for `n ∈ {0, 1, 2, 3}` so a launcher bug that mixes up branch slice pointers would land in the wrong scratch slot at the wrong scale and trip the per-branch ±5% rel_err assertion. Asserts each `scratch[1..5]` slot is within ±5% of its branch's true sorted-p99, all non-target slots remain zero (guards against the `2fb30f098`-style write-cascade), and the `(SP4_SLOT_BASE, ATOM_POS_BOUND_BASE, SP4_BRANCH_COUNT)` const-layout invariant + `atom_pos_bound(branch) == ATOM_POS_BOUND_BASE + branch` for `branch ∈ 0..4`. `#[ignore]`-gated for GPU. Local RTX 3050 Ti run: branch 0 (scale 1) rel_err=0.106%, branch 1 (scale 10) rel_err=0.362%, branch 2 (scale 100) rel_err=0.325%, branch 3 (scale 1000) rel_err=0.062% — all four pass with substantial margin under the 5% tolerance. The 5% headroom budget matches Tasks A4/A5: linear-bin quantization 0.4% + lane collision <0.012% + finite-sample p99 jitter ≈0.5% at n=4096 ≈ <2% true. Zero new ISV slots (ATOM_POS_BOUND[0..4] = ISV[132..136) already allocated by Task A1); zero new HtoD/DtoD/HtoH copies (mapped-pinned only); one new kernel + one new launcher method (single kernel handle, four launches via the loop); producer-only (no consumer). `cargo check -p ml --lib --tests` clean (12 pre-existing warnings, no new warnings); GPU unit test 1/1 passing locally.