From b1ef312a40e4d3ba615b4e1a36eaac6720c2ae65 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 1 May 2026 22:50:06 +0200 Subject: [PATCH] =?UTF-8?q?feat(sp5):=20Task=20A5=20=E2=80=94=20Pearl=205?= =?UTF-8?q?=20per-branch=20IQN=20=CF=84=20schedule=20GPU=20producer=20(Lay?= =?UTF-8?q?er=20A)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Two new CUDA kernels land as SP5 Layer A additive producers feeding ISV[250..270) with per-branch IQN quantile-τ schedules derived from Q-distribution skew. q_skew_kurtosis_update: single-block 4-thread, reads save_q_online[B×13], two-pass central moments → skew clamped [-3,+3] + ex_kurt clamped [-3,+30]; writes scratch[171..179). EPS_DIV=1e-12 (Invariant 1 anchor). No atomicAdd. pearl_5_iqn_tau_update: single-block 4-thread, shifts symmetric default τ {0.05,0.25,0.5,0.75,0.95} by skew×SKEW_SHIFT=0.05, clamps to [0.01,0.99]; writes scratch[179..199). SKEW_SHIFT and envelope are Invariant 1 anchors. launch_sp5_pearl_5_iqn_tau fires both kernels + 20 apply_pearls_ad calls (ALPHA_META=1e-3) → ISV[IQN_TAU_BASE=250..270). SP5_SCRATCH_TOTAL 171→199. StateResetRegistry +1 FoldReset entry (sp5_iqn_tau, ISV[250..270)). training_loop.rs wired after Pearl 4 with tracing::warn on error. Two GPU-only unit tests (10+11): zero-skew symmetric default + left-skew floor clamp. Module docstring updated A1-A5. No consumer migration — Layer A additive only per spec. cargo check -p ml --offline clean (11 pre-existing warnings, none new). cargo test --no-run clean. Co-Authored-By: Claude Sonnet 4.6 --- crates/ml/build.rs | 13 + .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 189 ++++++++++++- .../cuda_pipeline/pearl_5_iqn_tau_kernel.cu | 46 ++++ .../cuda_pipeline/q_skew_kurtosis_kernel.cu | 81 ++++++ .../src/trainers/dqn/state_reset_registry.rs | 9 + .../src/trainers/dqn/trainer/training_loop.rs | 6 + crates/ml/tests/sp5_producer_unit_tests.rs | 253 +++++++++++++++++- docs/dqn-wire-up-audit.md | 2 + 8 files changed, 596 insertions(+), 3 deletions(-) create mode 100644 crates/ml/src/cuda_pipeline/pearl_5_iqn_tau_kernel.cu create mode 100644 crates/ml/src/cuda_pipeline/q_skew_kurtosis_kernel.cu diff --git a/crates/ml/build.rs b/crates/ml/build.rs index f3ecb024c..f631f52b7 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -442,6 +442,19 @@ fn main() { // apply_pearls_ad_kernel then smooths via Pearls A+D (ALPHA_META=5e-4) // into ISV[226..250). Must run AFTER grad_cosine_sim_kernel. "pearl_4_adam_hparams_kernel.cu", + // SP5 Task A5 (2026-05-01): auxiliary per-branch Q-distribution skew + + // kurtosis reducer. Single-block 4-thread kernel reads save_q_online [B × 13]; + // computes per-branch (skew, excess kurtosis) via two-pass central moments; + // writes skew[4]→scratch[171..175), ex_kurt[4]→scratch[175..179). + // EPS_DIV=1e-12 Invariant 1 anchor. Structural envelopes [-3,+3]/[-3,+30]. + "q_skew_kurtosis_kernel.cu", + // SP5 Task A5 (2026-05-01): Pearl 5 per-branch IQN τ schedule controller. + // Single-block 4-thread kernel reads skew from scratch[skew_idx_base..+4); + // computes τ = default + skew × SKEW_SHIFT=0.05 (Invariant 1 anchor), + // clamped to [0.01, 0.99]; writes τ[4 branches × 5 quantiles]→scratch[179..199). + // apply_pearls_ad_kernel then smooths via Pearls A+D (ALPHA_META=1e-3) + // into ISV[IQN_TAU_BASE=250..270). Must run AFTER q_skew_kurtosis_kernel. + "pearl_5_iqn_tau_kernel.cu", // SP4 GPU-only Pearls A+D applicator (2026-05-01). Single-thread // device-side loop applies Pearls A+D to N consecutive ISV slots // (each with its own Wiener-state triple) on the producer's diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 03e8c0514..57fcd993b 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -336,6 +336,28 @@ static SP5_GRAD_COSINE_SIM_CUBIN: &[u8] = static SP5_PEARL_4_ADAM_HPARAMS_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/pearl_4_adam_hparams_kernel.cubin")); +/// SP5 Task A5 (2026-05-01): q_skew_kurtosis auxiliary cubin. +/// Single-block 4-thread kernel (one per branch). Reads save_q_online [B × 13] +/// row-major; computes per-branch (skew, excess kurtosis) via numerically-stable +/// two-pass central moments; writes 8 floats to scratch_buf[171..179) +/// (skew[4] at [171..175), ex_kurt[4] at [175..179)). +/// EPS_DIV=1e-12 Invariant 1 anchor (m2^1.5 / m2^2 denominator stability). +/// Envelope [-3,+3] / [-3,+30]. No atomicAdd (feedback_no_atomicadd). +static SP5_Q_SKEW_KURTOSIS_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/q_skew_kurtosis_kernel.cubin")); + +/// SP5 Task A5 (2026-05-01): Pearl 5 per-branch IQN τ schedule cubin. +/// Single-block 4-thread kernel (one per branch). Reads per-branch skew from +/// scratch_buf[skew_idx_base..skew_idx_base+4) (q_skew_kurtosis output). +/// Computes 5-quantile τ schedule per branch: symmetric default +/// {0.05, 0.25, 0.5, 0.75, 0.95} shifted by skew × SKEW_SHIFT=0.05 +/// (Invariant 1 anchor), clamped to structural envelope [0.01, 0.99]. +/// Writes 20 floats to scratch_buf[179..199) (branch b, quantile q at +/// 179 + b*5 + q). Chained apply_pearls_ad_kernel (20 launches, ALPHA_META=1e-3) +/// smooths into ISV[IQN_TAU_BASE=250..270). No atomicAdd. +static SP5_PEARL_5_IQN_TAU_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/pearl_5_iqn_tau_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 @@ -939,8 +961,30 @@ pub const SP5_WIENER_TOTAL_FLOATS: usize = /// 8 floats for pearl_4_beta1 output (SCRATCH_PEARL_4_BETA1=147, β1 × 8 groups) /// 8 floats for pearl_4_beta2 output (SCRATCH_PEARL_4_BETA2=155, β2 × 8 groups) /// 8 floats for pearl_4_eps output (SCRATCH_PEARL_4_EPS=163, ε × 8 groups) -/// Combined: 71 + 16 + 16 + 4 + 4 + 20 + 8 + 8 + 8 + 8 + 8 = 171 scratch slots [0..171). -pub const SP5_SCRATCH_TOTAL: usize = 171; +/// Task A5 adds: +/// 4 floats for q_skew_kurtosis skew[4] output (SCRATCH_PEARL_5_SKEW=171, skew × 4 branches) +/// 4 floats for q_skew_kurtosis ex_kurt[4] output (SCRATCH_PEARL_5_EX_KURT=175, ex_kurt × 4 branches) +/// 20 floats for pearl_5_iqn_tau τ output (SCRATCH_PEARL_5_TAU=179, [4 branches × 5 quantiles] at [179..199)) +/// Combined: 71 + 16 + 16 + 4 + 4 + 20 + 8 + 8 + 8 + 8 + 8 + 4 + 4 + 20 = 199 scratch slots [0..199). +pub const SP5_SCRATCH_TOTAL: usize = 199; + +/// SP5 Task A5: scratch index base for q_skew_kurtosis skew[4] output block. +/// Slots [171..175): per-branch Q-distribution standardized skewness. +/// Written by `q_skew_kurtosis_update`; consumed by `pearl_5_iqn_tau_update`. +pub const SCRATCH_PEARL_5_SKEW: usize = 171; + +/// SP5 Task A5: scratch index base for q_skew_kurtosis ex_kurt[4] output block. +/// Slots [175..179): per-branch Q-distribution excess kurtosis (clamped [-3, +30]). +/// Written by `q_skew_kurtosis_update`. Reserved for future heavy-tail-aware +/// τ schedule extension; not consumed by pearl_5_iqn_tau_update in this commit. +pub const SCRATCH_PEARL_5_EX_KURT: usize = 175; + +/// SP5 Task A5: scratch index base for pearl_5_iqn_tau τ schedule output block. +/// Slots [179..199): per-branch × per-quantile τ values (4 branches × 5 quantiles). +/// Layout: branch b, quantile q → scratch[179 + b*5 + q]. +/// Written by `pearl_5_iqn_tau_update`; consumed by apply_pearls_ad_kernel → +/// ISV[IQN_TAU_BASE=250..270). +pub const SCRATCH_PEARL_5_TAU: usize = 179; /// SP5 Task A2: scratch index base for pearl_3_sigma sigma[4] output block. /// Slots [103..107): per-branch NoisyNet sigma. @@ -4170,6 +4214,21 @@ pub struct GpuDqnTrainer { /// No atomicAdd, no CPU compute. /// Loaded from `pearl_4_adam_hparams_kernel.cubin`. pearl_4_adam_hparams_kernel: CudaFunction, + /// SP5 Task A5 (2026-05-01): q_skew_kurtosis auxiliary kernel. + /// Single-block 4-thread kernel (one per branch). Reads save_q_online [B × 13] + /// row-major; computes per-branch (skew, excess kurtosis) via two-pass central + /// moments; writes 8 floats to scratch_buf[171..179) (skew[4] + ex_kurt[4]). + /// Loaded from `q_skew_kurtosis_kernel.cubin`. + q_skew_kurtosis_kernel: CudaFunction, + /// SP5 Task A5 (2026-05-01): Pearl 5 per-branch IQN τ schedule controller kernel. + /// Single-block 4-thread kernel (one per branch). Reads per-branch skew from + /// scratch_buf[skew_idx_base..skew_idx_base+4); shifts symmetric default τ schedule + /// {0.05, 0.25, 0.5, 0.75, 0.95} by skew × SKEW_SHIFT=0.05 (Invariant 1 anchor); + /// clamps to [0.01, 0.99]; writes 20 floats to scratch_buf[179..199). + /// Chained apply_pearls_ad_kernel (20 single-slot launches, ALPHA_META=1e-3) smooths + /// into ISV[IQN_TAU_BASE=250..270). No atomicAdd, no consumer migration in this commit. + /// Loaded from `pearl_5_iqn_tau_kernel.cubin`. + pearl_5_iqn_tau_kernel: CudaFunction, /// SP5 Task A4 (2026-05-01): previous-step gradient buffer for cosine similarity. /// [total_params f32] mapped-pinned (feedback_no_htod_htoh_only_mapped_pinned). /// Zeroed at construction and at each fold boundary (Pearl A sentinel: cosine_sim=0 @@ -11012,6 +11071,109 @@ impl GpuDqnTrainer { Ok(()) } + /// SP5 Pearl 5: per-branch IQN τ schedule from per-branch Q-skew. + /// + /// Two-kernel chain on the producer's stream: + /// 1. `q_skew_kurtosis_update` — computes per-branch (skew, ex_kurt) from + /// save_q_online ([B × 13] row-major); writes 8 floats to + /// scratch[171..179) (skew[4] at [171..175), ex_kurt[4] at [175..179)). + /// 2. `pearl_5_iqn_tau_update` — derives 5-quantile τ schedule per branch + /// with skew-driven shift + structural-envelope clamp; writes 20 floats + /// to scratch[179..199) (branch b, quantile q at 179 + b*5 + q). + /// + /// 20 apply_pearls calls follow (4 branches × 5 quantiles), each chained on + /// the same stream with ALPHA_META=1e-3 (default SP4/SP5 rate). + /// + /// Pearl 5 reads save_q_online (forward-pass output); independent of Pearls + /// 1/2/3/4 ISV outputs. + /// + /// Layer A only — IQN τ-quantile-selection consumer migration deferred to + /// Layer B. + /// + /// Wiener offset formula: `base_wiener_offset + (isv_slot - SP5_SLOT_BASE) * 3` + /// where `base_wiener_offset = SP4_PRODUCER_COUNT * 3 = 213`. + pub(crate) fn launch_sp5_pearl_5_iqn_tau(&self) -> Result<(), MLError> { + use crate::cuda_pipeline::sp5_isv_slots::{IQN_TAU_BASE, SP5_SLOT_BASE}; + use crate::cuda_pipeline::sp4_wiener_ema::{launch_apply_pearls, ALPHA_META}; + + debug_assert!(self.isv_signals_dev_ptr != 0, + "launch_sp5_pearl_5_iqn_tau: isv_signals_dev_ptr must be allocated by constructor"); + + let scratch_dev = self.producer_step_scratch_buf.dev_ptr; + let isv_dev = self.isv_signals_dev_ptr; + let wiener_dev = self.wiener_state_buf.dev_ptr; + let action_counts_dev = self.action_counts_buf.dev_ptr; + let action_offsets_dev = self.action_offsets_buf.dev_ptr; + let q_out_ptr = self.q_out_buf.raw_ptr(); + let b_i32 = self.config.batch_size as i32; + + // Wiener base offset for SP5 slots: SP4_PRODUCER_COUNT * 3 slots consumed. + let base_wiener_offset: i32 = (SP4_PRODUCER_COUNT * 3) as i32; + + let skew_base_i32 = SCRATCH_PEARL_5_SKEW as i32; // 171 + let ex_kurt_base_i32 = SCRATCH_PEARL_5_EX_KURT as i32; // 175 + let tau_base_i32 = SCRATCH_PEARL_5_TAU as i32; // 179 + + let cfg_4t = LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (4, 1, 1), + shared_mem_bytes: 0, + }; + + // Step 1: q_skew_kurtosis_update — reads save_q_online [B × 13]; + // writes skew[4] → scratch[171..175), ex_kurt[4] → scratch[175..179). + unsafe { + self.stream + .launch_builder(&self.q_skew_kurtosis_kernel) + .arg(&q_out_ptr) + .arg(&b_i32) + .arg(&action_counts_dev) + .arg(&action_offsets_dev) + .arg(&scratch_dev) + .arg(&skew_base_i32) + .arg(&ex_kurt_base_i32) + .launch(cfg_4t) + .map_err(|e| MLError::ModelError(format!("q_skew_kurtosis_update: {e}")))?; + } + + // Step 2: pearl_5_iqn_tau_update — reads scratch[171..175) (skew); + // writes τ[4 branches × 5 quantiles] → scratch[179..199). + // Stream-ordered with Step 1 (same stream, no explicit sync needed). + unsafe { + self.stream + .launch_builder(&self.pearl_5_iqn_tau_kernel) + .arg(&scratch_dev) + .arg(&skew_base_i32) + .arg(&scratch_dev) + .arg(&tau_base_i32) + .launch(cfg_4t) + .map_err(|e| MLError::ModelError(format!("pearl_5_iqn_tau_update: {e}")))?; + } + + // Step 3: 20 apply_pearls calls (4 branches × 5 quantiles). + // τ[b*5 + q]: scratch[179 + b*5 + q] → ISV[IQN_TAU_BASE + b*5 + q]. + for b in 0..4_usize { + for q in 0..5_usize { + let isv_slot = (IQN_TAU_BASE + b * 5 + q) as i32; + let scratch_idx = (SCRATCH_PEARL_5_TAU + b * 5 + q) as i32; + let wiener_off = base_wiener_offset + (isv_slot - SP5_SLOT_BASE as i32) * 3; + unsafe { + launch_apply_pearls( + &self.stream, + &self.apply_pearls_ad_kernel, + scratch_dev, scratch_idx, + isv_dev, isv_slot, + wiener_dev, wiener_off, + 1, + ALPHA_META, + )?; + } + } + } + + 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]. /// @@ -13074,6 +13236,27 @@ impl GpuDqnTrainer { .map_err(|e| MLError::ModelError(format!("pearl_4_adam_hparams_update load: {e}")))? }; + // SP5 Task A5 (2026-05-01): load q_skew_kurtosis_kernel (auxiliary + // per-branch Q-distribution skew + kurtosis reducer). Single-block + // 4-thread kernel. Reads save_q_online [B × 13]; writes skew[4] → + // scratch[171..175) and ex_kurt[4] → scratch[175..179). + let q_skew_kurtosis_kernel = { + let module = stream.context().load_cubin(SP5_Q_SKEW_KURTOSIS_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("sp5 q_skew_kurtosis cubin load: {e}")))?; + module.load_function("q_skew_kurtosis_update") + .map_err(|e| MLError::ModelError(format!("q_skew_kurtosis_update load: {e}")))? + }; + + // SP5 Task A5 (2026-05-01): load pearl_5_iqn_tau_kernel. + // Single-block 4-thread kernel. Reads skew from scratch; writes + // τ schedule [4 branches × 5 quantiles] → scratch[179..199). + let pearl_5_iqn_tau_kernel = { + let module = stream.context().load_cubin(SP5_PEARL_5_IQN_TAU_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("sp5 pearl_5_iqn_tau cubin load: {e}")))?; + module.load_function("pearl_5_iqn_tau_update") + .map_err(|e| MLError::ModelError(format!("pearl_5_iqn_tau_update load: {e}")))? + }; + // SP5 Task A4 (2026-05-01): allocate grad_prev_buf_per_group [total_params f32]. // Zeroed at construction; fold-boundary reset zeroes it again (Pearl A sentinel: // zero grad_prev → cosine_sim=0 on first step → β1/β2 start at envelope midpoints). @@ -16221,6 +16404,8 @@ impl GpuDqnTrainer { pearl_2_budget_kernel, grad_cosine_sim_kernel, pearl_4_adam_hparams_kernel, + q_skew_kurtosis_kernel, + pearl_5_iqn_tau_kernel, grad_prev_buf_per_group, group_param_offsets_buf, atoms_clip_count_buf, diff --git a/crates/ml/src/cuda_pipeline/pearl_5_iqn_tau_kernel.cu b/crates/ml/src/cuda_pipeline/pearl_5_iqn_tau_kernel.cu new file mode 100644 index 000000000..620182ca8 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/pearl_5_iqn_tau_kernel.cu @@ -0,0 +1,46 @@ +// crates/ml/src/cuda_pipeline/pearl_5_iqn_tau_kernel.cu +// +// SP5 Task A5: Pearl 5 — per-branch IQN τ schedule from per-branch Q-skew. +// +// Reads q_skew_per_branch[b] from scratch_buf at skew_idx_base + b (written by +// q_skew_kurtosis_update). Computes per-branch 5-quantile τ schedule: +// - Symmetric default: {0.05, 0.25, 0.5, 0.75, 0.95} +// - Skewed: shift each τ by skew × SKEW_SHIFT toward the dense tail +// Clamps every τ to the structural envelope [0.01, 0.99]. +// +// 4-thread single-block (one per branch); each thread emits 5 contiguous +// τ values to scratch_out at tau_idx_base + b*5 + q. +// +// SKEW_SHIFT=0.05 is an Invariant 1 anchor (modest schedule shift; preserves +// quantile spread; max envelope reach is ±0.15 at clamped skew=±3). +// kurtosis is computed by q_skew_kurtosis_update and lives in scratch at +// ex_kurt_idx_base — reserved for future heavy-tail-aware schedule extension; +// this kernel reads only skew. The excess kurtosis slots are not consumed here. + +#include + +extern "C" __global__ void pearl_5_iqn_tau_update( + const float* __restrict__ scratch_buf, // scratch buffer for reading skew + int skew_idx_base, // skew[4] written by q_skew_kurtosis_update + float* __restrict__ scratch_out, // scratch output (same buffer, non-overlapping) + int tau_idx_base // 20 floats: branch b, quantile q at base + b*5 + q +) { + int b = threadIdx.x; + if (blockIdx.x != 0 || b >= 4) return; + + float skew = scratch_buf[skew_idx_base + b]; + // SKEW_SHIFT=0.05 is an Invariant 1 anchor (modest schedule shift). + const float SKEW_SHIFT = 0.05f; + + // Symmetric default offsets (Invariant 1 anchors — quantile interpolation grid). + float tau_defaults[5] = {0.05f, 0.25f, 0.5f, 0.75f, 0.95f}; + + for (int q = 0; q < 5; ++q) { + float t = tau_defaults[q] + skew * SKEW_SHIFT; + // Structural envelope [0.01, 0.99] (Invariant 1 anchor — IQN requires + // τ strictly in (0,1); clamp prevents degenerate boundary quantiles). + t = fmaxf(0.01f, fminf(0.99f, t)); + scratch_out[tau_idx_base + b * 5 + q] = t; + } + __threadfence_system(); +} diff --git a/crates/ml/src/cuda_pipeline/q_skew_kurtosis_kernel.cu b/crates/ml/src/cuda_pipeline/q_skew_kurtosis_kernel.cu new file mode 100644 index 000000000..510ab40b5 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/q_skew_kurtosis_kernel.cu @@ -0,0 +1,81 @@ +// crates/ml/src/cuda_pipeline/q_skew_kurtosis_kernel.cu +// +// SP5 Task A5 auxiliary: per-branch Q-distribution skew + kurtosis reduction. +// Reads save_q_online[B × 13] row-major; computes per-branch standardized +// skewness and excess kurtosis. Used by Pearl 5 to adapt the IQN τ schedule. +// +// Single-block, 4-thread (one per branch: Direction(4) / Magnitude(3) / +// Order(3) / Urgency(3)). Each thread reduces over its branch's +// (batch × n_actions) slice using a numerically-stable two-pass approach: +// Pass 1: mean +// Pass 2: m2 (variance), m3 (third central moment), m4 (fourth) +// skew = m3 / m2^1.5 (with EPS_DIV denom; clamped to [-3, +3]) +// ex_kurt = m4 / m2^2 - 3.0 (excess kurtosis; clamped to [-3, +30]) +// +// Writes (skew[4], ex_kurt[4]) — 8 floats — to scratch_out at provided bases. +// __threadfence_system() after writes. +// +// No atomicAdd (feedback_no_atomicadd) — one thread per branch, independent +// scratch writes. + +#include + +extern "C" __global__ void q_skew_kurtosis_update( + const float* __restrict__ save_q_online_ptr, // [B × 13] row-major + int batch_size, + const int* __restrict__ action_counts, // [4] = {4, 3, 3, 3} + const int* __restrict__ action_offsets, // [4] = {0, 4, 7, 10} + float* __restrict__ scratch_out, + int skew_idx_base, // skew[4] at [base..base+4) + int ex_kurt_idx_base // excess kurtosis[4] at [base..base+4) +) { + int b = threadIdx.x; + if (blockIdx.x != 0 || b >= 4) return; + + int n_actions = action_counts[b]; + int act_off = action_offsets[b]; + int total = batch_size * n_actions; + if (total == 0) return; + + // Pass 1: mean + float sum = 0.0f; + for (int i = 0; i < batch_size; ++i) { + for (int a = 0; a < n_actions; ++a) { + sum += save_q_online_ptr[i * 13 + act_off + a]; + } + } + float mean = sum / (float)total; + + // Pass 2: central moments m2, m3, m4 + float m2 = 0.0f, m3 = 0.0f, m4 = 0.0f; + for (int i = 0; i < batch_size; ++i) { + for (int a = 0; a < n_actions; ++a) { + float d = save_q_online_ptr[i * 13 + act_off + a] - mean; + float d2 = d * d; + m2 += d2; + m3 += d2 * d; + m4 += d2 * d2; + } + } + m2 /= (float)total; + m3 /= (float)total; + m4 /= (float)total; + + // EPS_DIV = 1e-12 is an Invariant 1 anchor — numerical stability for + // m2^1.5 / m2^2 denominators. Prevents 0/0 when Q-distribution is constant. + const float EPS_DIV = 1e-12f; + float m2_15 = m2 * sqrtf(m2) + EPS_DIV; + float m2_2 = m2 * m2 + EPS_DIV; + + float skew = m3 / m2_15; + float ex_kurt = m4 / m2_2 - 3.0f; + + // Structural envelopes (Invariant 1 anchors): clamp to keep downstream τ + // shifts within [0.01, 0.99] and prevent runaway under degenerate input. + skew = fmaxf(-3.0f, fminf(3.0f, skew)); + ex_kurt = fmaxf(-3.0f, fminf(30.0f, ex_kurt)); + + scratch_out[skew_idx_base + b] = skew; + scratch_out[ex_kurt_idx_base + b] = ex_kurt; + __threadfence_system(); +} diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index b2e42e689..7192c8b1c 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -651,6 +651,15 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "GpuDqnTrainer.grad_prev_buf_per_group [total_params f32s] — previous-step gradient direction for cosine similarity computation (Pearl 4). Zero at fold boundary so the new fold's first grad_cosine_sim_update sees cosine_sim=0 (low-stability bootstrap) and Pearl A fires first-observation replacement (Task A4).", }, + // SP5 Task A5: per-branch IQN τ schedule (Pearl 5). + // ISV[IQN_TAU_BASE=250..270) = 20 floats (4 branches × 5 quantiles). + // Pearl A sentinel 0 at fold boundary → first step fires first-observation + // replacement: zero skew → default symmetric schedule {0.05,0.25,0.5,0.75,0.95}. + RegistryEntry { + name: "sp5_iqn_tau", + category: ResetCategory::FoldReset, + description: "ISV[IQN_TAU_BASE=250..270) — per-branch IQN τ schedule (4 branches × 5 quantiles). SP5 Pearl A sentinel 0 at fold boundary so the new fold's first pearl_5_iqn_tau_update launch fires the first-observation replacement back to the symmetric default {0.05, 0.25, 0.5, 0.75, 0.95} at zero skew (Task A5).", + }, ]; Self { entries } } diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index f8b86f1cf..7d6a5de69 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -3610,6 +3610,12 @@ impl DQNTrainer { if let Err(e) = fused.trainer().launch_sp5_pearl_4_adam_hparams() { tracing::warn!(error = %e, "SP5 Pearl 4 producer launch failed"); } + // SP5 Task A5: pearl_5_iqn_tau → apply_pearls_ad ×20 + // Reads save_q_online (forward-pass output); independent of Pearls 1/2/3/4 ISV outputs. + // Writes ISV[250..270) = τ[4 branches × 5 quantiles] (ALPHA_META=1e-3). + if let Err(e) = fused.trainer().launch_sp5_pearl_5_iqn_tau() { + tracing::warn!(error = %e, "SP5 Pearl 5 producer launch failed"); + } } } diff --git a/crates/ml/tests/sp5_producer_unit_tests.rs b/crates/ml/tests/sp5_producer_unit_tests.rs index 5fe6c84ec..0db77f949 100644 --- a/crates/ml/tests/sp5_producer_unit_tests.rs +++ b/crates/ml/tests/sp5_producer_unit_tests.rs @@ -2,7 +2,7 @@ //! SP5 Layer A producer kernel unit tests. //! -//! Tests cover the SP5 Layer A CUDA kernels across Tasks A1–A4: +//! Tests cover the SP5 Layer A CUDA kernels across Tasks A1–A5: //! //! Task A1 (Pearl 1 atom-span + Q-stats source): //! 1. `sp5_q_branch_stats_uniform_input_matches_analytical_ground_truth` @@ -17,6 +17,9 @@ //! 7. `pearl_4_adam_hparams_high_stability_yields_long_memory` //! 8. `pearl_4_adam_hparams_low_stability_yields_short_memory` //! 9. `grad_cosine_sim_per_group_dot_norm_and_writeback` +//! Task A5 (Pearl 5 per-branch IQN τ schedule): +//! 10. `pearl_5_iqn_tau_symmetric_q_yields_uniform_default_schedule` +//! 11. `pearl_5_iqn_tau_left_skew_shifts_quantiles_negatively` //! //! All tests use analytically-known synthetic inputs with no CPU reference //! implementation, per `feedback_no_cpu_test_fallbacks.md`. All mapped-pinned @@ -51,6 +54,12 @@ const SP5_PEARL_4_ADAM_HPARAMS_CUBIN: &[u8] = const SP5_GRAD_COSINE_SIM_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/grad_cosine_sim_kernel.cubin")); +const SP5_Q_SKEW_KURTOSIS_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/q_skew_kurtosis_kernel.cubin")); + +const SP5_PEARL_5_IQN_TAU_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/pearl_5_iqn_tau_kernel.cubin")); + // ── Helpers ─────────────────────────────────────────────────────────────────── fn load_pearl_3_sigma_kernel(stream: &Arc) -> CudaFunction { @@ -1260,3 +1269,245 @@ fn grad_cosine_sim_per_group_dot_norm_and_writeback() { ); } } + +// ── Test 10 + 11 helpers (Task A5) ──────────────────────────────────────────── + +fn load_q_skew_kurtosis_kernel(stream: &Arc) -> CudaFunction { + let module = stream + .context() + .load_cubin(SP5_Q_SKEW_KURTOSIS_CUBIN.to_vec()) + .expect("load q_skew_kurtosis_kernel cubin"); + module + .load_function("q_skew_kurtosis_update") + .expect("load q_skew_kurtosis_update function") +} + +fn load_pearl_5_iqn_tau_kernel(stream: &Arc) -> CudaFunction { + let module = stream + .context() + .load_cubin(SP5_PEARL_5_IQN_TAU_CUBIN.to_vec()) + .expect("load pearl_5_iqn_tau_kernel cubin"); + module + .load_function("pearl_5_iqn_tau_update") + .expect("load pearl_5_iqn_tau_update function") +} + +// ── Test 10: zero skew → symmetric default τ schedule ───────────────────────── + +/// SP5 Task A5 unit test: constant Q-values (zero variance → zero skew) produce +/// the symmetric default τ schedule {0.05, 0.25, 0.5, 0.75, 0.95} on all branches. +/// +/// Synthetic input: B=4 batch, q_out_buf [4 × 13] all set to constant 3.0. +/// +/// Analytical ground truth for constant input: +/// mean = 3.0, all deviations = 0 → m2 = m3 = m4 = 0 +/// m2_15 = 0 + EPS_DIV = 1e-12 +/// m3 / m2_15 = 0 / 1e-12 = 0 → skew = 0 (before clamp, still 0) +/// τ[q] = tau_defaults[q] + 0 × SKEW_SHIFT = tau_defaults[q] +/// Expected τ per branch: {0.05, 0.25, 0.5, 0.75, 0.95} +/// +/// This test chains q_skew_kurtosis_update → pearl_5_iqn_tau_update exactly as +/// the production launcher does (scratch[171..175) is the skew bridge slot). +/// No CPU reference computation per `feedback_no_cpu_test_fallbacks.md`. +#[test] +#[ignore = "requires GPU"] +fn pearl_5_iqn_tau_symmetric_q_yields_uniform_default_schedule() { + const BATCH_SIZE: usize = 4; + const TOTAL_ACTIONS: usize = 13; + const SKEW_IDX_BASE: usize = 171; // SCRATCH_PEARL_5_SKEW + const EX_KURT_IDX_BASE: usize = 175; // SCRATCH_PEARL_5_EX_KURT + const TAU_IDX_BASE: usize = 179; // SCRATCH_PEARL_5_TAU + const SCRATCH_SIZE: usize = TAU_IDX_BASE + 20; // 199 + + let stream = make_test_stream(); + let skew_kernel = load_q_skew_kurtosis_kernel(&stream); + let tau_kernel = load_pearl_5_iqn_tau_kernel(&stream); + + // q_out_buf [B × 13]: all 3.0 — constant Q across all branches. + let q_buf = unsafe { MappedF32Buffer::new(BATCH_SIZE * TOTAL_ACTIONS) } + .expect("alloc q_buf"); + let q_data = vec![3.0_f32; BATCH_SIZE * TOTAL_ACTIONS]; + q_buf.write_from_slice(&q_data); + + // action_counts = {4, 3, 3, 3} + let action_counts = [4_i32, 3, 3, 3]; + let action_counts_buf = unsafe { MappedI32Buffer::new(4) } + .expect("alloc action_counts_buf"); + action_counts_buf.write_from_slice(&action_counts); + + // action_offsets = {0, 4, 7, 10} + let action_offsets = [0_i32, 4, 7, 10]; + let action_offsets_buf = unsafe { MappedI32Buffer::new(4) } + .expect("alloc action_offsets_buf"); + action_offsets_buf.write_from_slice(&action_offsets); + + let scratch_buf = unsafe { MappedF32Buffer::new(SCRATCH_SIZE) } + .expect("alloc scratch_buf"); + + let batch_size_i32 = BATCH_SIZE as i32; + let skew_idx_i32 = SKEW_IDX_BASE as i32; + let ex_kurt_idx_i32 = EX_KURT_IDX_BASE as i32; + let tau_idx_i32 = TAU_IDX_BASE as i32; + + let q_dev = q_buf.dev_ptr; + let counts_dev = action_counts_buf.dev_ptr; + let offsets_dev = action_offsets_buf.dev_ptr; + let scratch_dev = scratch_buf.dev_ptr; + + // Step 1: q_skew_kurtosis_update + unsafe { + stream + .launch_builder(&skew_kernel) + .arg(&q_dev) + .arg(&batch_size_i32) + .arg(&counts_dev) + .arg(&offsets_dev) + .arg(&scratch_dev) + .arg(&skew_idx_i32) + .arg(&ex_kurt_idx_i32) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (4, 1, 1), + shared_mem_bytes: 0, + }) + .expect("launch q_skew_kurtosis_update (zero-skew test)"); + } + stream.synchronize().expect("sync after q_skew_kurtosis_update (zero-skew test)"); + + // Step 2: pearl_5_iqn_tau_update + unsafe { + stream + .launch_builder(&tau_kernel) + .arg(&scratch_dev) + .arg(&skew_idx_i32) + .arg(&scratch_dev) + .arg(&tau_idx_i32) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (4, 1, 1), + shared_mem_bytes: 0, + }) + .expect("launch pearl_5_iqn_tau_update (zero-skew test)"); + } + stream.synchronize().expect("sync after pearl_5_iqn_tau_update (zero-skew test)"); + + let scratch = scratch_buf.read_all(); + + // Verify skew[4] ≈ 0 (constant input → zero skew) + for b in 0..4_usize { + let skew = scratch[SKEW_IDX_BASE + b]; + println!("branch={b} skew={skew:.8}"); + assert!( + skew.abs() < 1e-5, + "branch={b}: expected skew≈0 for constant Q, got {skew}" + ); + } + + // Verify τ[4×5] ≈ symmetric defaults {0.05, 0.25, 0.5, 0.75, 0.95} + let defaults = [0.05_f32, 0.25, 0.5, 0.75, 0.95]; + for b in 0..4_usize { + for q in 0..5_usize { + let tau = scratch[TAU_IDX_BASE + b * 5 + q]; + let expected = defaults[q]; + println!("branch={b} q={q} tau={tau:.6} expected={expected:.6}"); + assert!( + (tau - expected).abs() < 1e-5, + "branch={b} q={q}: tau={tau:.6} expected={expected:.6}" + ); + } + } +} + +// ── Test 11: negative skew shifts τ schedule down ───────────────────────────── + +/// SP5 Task A5 unit test: negative Q-distribution skew shifts the τ schedule +/// downward (toward lower quantiles), verifying the SKEW_SHIFT=0.05 mechanism. +/// +/// Synthetic input: inject skew = -1.0 directly into the scratch buffer +/// (SCRATCH_PEARL_5_SKEW slot) and launch only pearl_5_iqn_tau_update. +/// This isolates the τ-schedule kernel from the skew-computation kernel. +/// +/// Analytical ground truth for skew = -1.0: +/// τ[q] = tau_defaults[q] + (-1.0) × 0.05 = tau_defaults[q] - 0.05 +/// Expected τ: {0.0, 0.20, 0.45, 0.70, 0.90} +/// Clamp check: 0.05 - 0.05 = 0.00 → clamped to 0.01 (structural envelope [0.01, 0.99]) +/// So expected τ[0] = 0.01, τ[1..5] = {0.20, 0.45, 0.70, 0.90} +/// +/// Checks that τ values are strictly less than the symmetric defaults (shifted down) +/// and that the 0.01 floor clamp fires on the first quantile. +/// No CPU reference computation per `feedback_no_cpu_test_fallbacks.md`. +#[test] +#[ignore = "requires GPU"] +fn pearl_5_iqn_tau_left_skew_shifts_quantiles_negatively() { + const SKEW_IDX_BASE: usize = 171; // SCRATCH_PEARL_5_SKEW + const TAU_IDX_BASE: usize = 179; // SCRATCH_PEARL_5_TAU + const SCRATCH_SIZE: usize = TAU_IDX_BASE + 20; // 199 + const INJECTED_SKEW: f32 = -1.0; + + let stream = make_test_stream(); + let tau_kernel = load_pearl_5_iqn_tau_kernel(&stream); + + let scratch_buf = unsafe { MappedF32Buffer::new(SCRATCH_SIZE) } + .expect("alloc scratch_buf"); + + // Inject skew = -1.0 for all 4 branches directly into the scratch slot. + let mut scratch_init = vec![0.0_f32; SCRATCH_SIZE]; + for b in 0..4_usize { + scratch_init[SKEW_IDX_BASE + b] = INJECTED_SKEW; + } + scratch_buf.write_from_slice(&scratch_init); + + let scratch_dev = scratch_buf.dev_ptr; + let skew_idx_i32 = SKEW_IDX_BASE as i32; + let tau_idx_i32 = TAU_IDX_BASE as i32; + + unsafe { + stream + .launch_builder(&tau_kernel) + .arg(&scratch_dev) + .arg(&skew_idx_i32) + .arg(&scratch_dev) + .arg(&tau_idx_i32) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (4, 1, 1), + shared_mem_bytes: 0, + }) + .expect("launch pearl_5_iqn_tau_update (left-skew test)"); + } + stream.synchronize().expect("sync after pearl_5_iqn_tau_update (left-skew test)"); + + let scratch = scratch_buf.read_all(); + + // Analytical expected τ for skew = -1.0: + // tau_defaults[q] + (-1.0) × 0.05, clamped to [0.01, 0.99] + let expected = [0.01_f32, 0.20, 0.45, 0.70, 0.90]; + + for b in 0..4_usize { + for q in 0..5_usize { + let tau = scratch[TAU_IDX_BASE + b * 5 + q]; + let exp = expected[q]; + println!("branch={b} q={q} tau={tau:.6} expected={exp:.6}"); + assert!( + (tau - exp).abs() < 1e-5, + "branch={b} q={q}: tau={tau:.6} expected={exp:.6}" + ); + } + // τ[1..5] must be strictly below their symmetric defaults (skew shifted them down). + let defaults = [0.05_f32, 0.25, 0.5, 0.75, 0.95]; + for q in 1..5_usize { + let tau = scratch[TAU_IDX_BASE + b * 5 + q]; + assert!( + tau < defaults[q], + "branch={b} q={q}: tau={tau:.6} should be below default={:.6}", + defaults[q] + ); + } + // τ[0] must be at the floor (clamp fired: 0.05 - 0.05 = 0.00 → clamped to 0.01). + let tau0 = scratch[TAU_IDX_BASE + b * 5]; + assert!( + (tau0 - 0.01_f32).abs() < 1e-5, + "branch={b} q=0: expected clamp floor 0.01, got {tau0:.6}" + ); + } +} diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 740a423dc..7bdc9462e 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,8 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +SP5 Task A5 — Pearl 5 per-branch IQN τ schedule GPU producer (Layer A, 2026-05-01): two new CUDA kernels (`q_skew_kurtosis_kernel.cu`, `pearl_5_iqn_tau_kernel.cu`) land as Layer A additive producers that feed ISV slots [250..270) with per-branch IQN quantile-τ schedules. No IQN τ-selection consumer migration in this commit — Layer B wires the τ consumer when distributional action-selection is migrated. `q_skew_kurtosis_update`: single-block 4-thread kernel (one thread per branch), reads `save_q_online[B × 13]` row-major; two-pass central moments (mean then m2/m3/m4); computes `skew = m3 / (m2^1.5 + EPS_DIV)` and `ex_kurt = m4 / (m2^2 + EPS_DIV) - 3`; EPS_DIV=1e-12 is an Invariant 1 anchor (numerical stability for degenerate constant-Q distribution); clamps skew to [-3,+3] and ex_kurt to [-3,+30] (structural envelopes); writes skew[4] to scratch[171..175) and ex_kurt[4] to scratch[175..179); `__threadfence_system()` after writes. No atomicAdd — one thread per branch with independent scratch writes. `pearl_5_iqn_tau_update`: single-block 4-thread kernel, reads scratch_buf[skew_idx_base + b]; shifts symmetric default schedule {0.05,0.25,0.5,0.75,0.95} by `skew × SKEW_SHIFT=0.05` (Invariant 1 anchor — modest schedule shift, max ±0.15 at clamped skew=±3); clamps every τ to structural envelope [0.01,0.99] (Invariant 1 anchor — IQN requires τ strictly in (0,1)); writes 20 τ floats to scratch_out[tau_idx_base + b×5 + q] at scratch[179..199). `launch_sp5_pearl_5_iqn_tau` fires both kernels sequentially then 20 `launch_apply_pearls` calls (4 branches × 5 quantiles → ISV[IQN_TAU_BASE=250..270)). ALPHA_META=1e-3 (same as SP4 default; Pearl 5's τ signal is direct per-step output, not an EMA-of-EMA, so slower smoothing relative to Pearl 4's 5e-4 not needed). Wiener offset formula: `base_wiener_offset + (isv_slot - SP5_SLOT_BASE) × 3` where `base_wiener_offset = SP4_PRODUCER_COUNT × 3 = 213`; for ISV slot 250 (first IQN_TAU slot): wiener_off = 213 + (250-174)×3 = 441. Buffer growth: `producer_step_scratch_buf` 171→199 slots (SP5_SCRATCH_TOTAL); 3 new scratch constants `SCRATCH_PEARL_5_SKEW=171`, `SCRATCH_PEARL_5_EX_KURT=175`, `SCRATCH_PEARL_5_TAU=179`. wiener_state_buf stays at 543 (SP5_WIENER_TOTAL_FLOATS — sized at Task A1 for all 110 SP5 slots, no growth). StateResetRegistry gains 1 new FoldReset entry: `sp5_iqn_tau` (ISV[250..270)) — Pearl A sentinel 0 at fold boundary triggers first-observation replacement to zero-skew symmetric default {0.05,0.25,0.5,0.75,0.95}. Wire-up in training_loop.rs: `launch_sp5_pearl_5_iqn_tau()` called immediately after `launch_sp5_pearl_4_adam_hparams()` with `tracing::warn!(error = %e, ...)` on error. Two GPU-only `#[ignore = "requires GPU"]` unit tests: (10) constant Q input → zero skew → symmetric default τ (chains both kernels); (11) injected skew=-1.0 → τ shifted down by 0.05, τ[0] clamped to 0.01 floor (exercises SKEW_SHIFT mechanism and structural envelope clamp). No CPU compute, no HtoD/HtoH, no atomicAdd, no stubs, no consumer migration. Touched: `cuda_pipeline/q_skew_kurtosis_kernel.cu` (new), `cuda_pipeline/pearl_5_iqn_tau_kernel.cu` (new), `build.rs` (+2 cubin entries), `cuda_pipeline/gpu_dqn_trainer.rs` (SP5_SCRATCH_TOTAL 171→199 + 3 new scratch constants + 2 static cubins + 2 struct fields + cubin loaders + struct initializer + launch method), `trainers/dqn/state_reset_registry.rs` (+1 FoldReset entry), `trainers/dqn/trainer/training_loop.rs` (+5 LOC after Pearl 4 launch), `tests/sp5_producer_unit_tests.rs` (+2 GPU-only tests + 2 cubin constants + 2 loader helpers + module docstring). Hard rules: feedback_no_cpu_compute_strict, feedback_no_atomicadd, feedback_no_htod_htoh_only_mapped_pinned, feedback_no_stubs, feedback_no_partial_refactor (no consumer migration — Layer A additive only). + SP5 docs/comment-only fix-up — close two minor review findings before Layer A Task A5 (2026-05-01): two minor code-review nits accumulated across A1-A4 reviews; closed in a single docs/comment-only commit. (1) `tests/sp5_producer_unit_tests.rs` module docstring (A4 review) updated to enumerate all 9 SP5 Layer A tests (q_branch_stats + pearl_1 + pearl_3 × 2 + pearl_2 × 2 + pearl_4 × 2 + grad_cosine_sim) so a reader cold to the file can locate each producer's coverage without scanning the body. (2) `gpu_dqn_trainer.rs:2854` stale field comment `// [B, TOTAL_ACTIONS(11)]` (A1 review) updated to `// [B, TOTAL_ACTIONS(13)] — 4 direction + 3 magnitude + 3 order + 3 urgency` so the field declaration matches the SP5 producer docstrings (e.g. q_branch_stats_kernel comments at line 277/282) and the 4-branch action layout established in commit `2fb30f098`. Comment-only changes; no behavior change. cargo check + cargo test --no-run both clean (11 pre-existing warnings, none new). SP5 Task A4 fix-up — grad_cosine_sim_update unit test added (2026-05-01): code-quality review caught that `pearl_4_adam_hparams_kernel` had direct GPU tests but the auxiliary `grad_cosine_sim_kernel` had none. Test 7 and Test 8 launched the hparams kernel directly with pre-baked cosine inputs, never exercising the per-group reduction (dot + 2× L2 norm sums) or the writeback (grad_curr → grad_prev for next step's comparison). The writeback in particular is load-bearing for cross-step cosine evolution — if broken, every subsequent step's cosine_sim is computed against a stale or mis-aligned previous gradient. Fix adds Test 9 `grad_cosine_sim_per_group_dot_norm_and_writeback` with 8 analytically-known synthetic group cases: identical (cos=+1), orthogonal (0), antiparallel (-1), Pythagorean (cos=+1, |c|=5), cold-start curr (zero grad → cos=0, |c|=0), cold-start prev (Pearl A sentinel state → cos=0, |c|=1), and 2 repeats. NO CPU oracle — expected values are unit-vector cosines from the kernel formula. Writeback verified by reading grad_prev_buf after the kernel runs and asserting bit-identity with grad_curr_buf for all 32 elements. Touched: `crates/ml/tests/sp5_producer_unit_tests.rs` (+1 cubin const + 1 loader + 120 LOC test). cargo test --no-run clean.