feat(sp5): Task A5 — Pearl 5 per-branch IQN τ schedule GPU producer (Layer A)
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 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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
|
||||
|
||||
@@ -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,
|
||||
|
||||
46
crates/ml/src/cuda_pipeline/pearl_5_iqn_tau_kernel.cu
Normal file
46
crates/ml/src/cuda_pipeline/pearl_5_iqn_tau_kernel.cu
Normal file
@@ -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 <cuda_runtime.h>
|
||||
|
||||
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();
|
||||
}
|
||||
81
crates/ml/src/cuda_pipeline/q_skew_kurtosis_kernel.cu
Normal file
81
crates/ml/src/cuda_pipeline/q_skew_kurtosis_kernel.cu
Normal file
@@ -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 <cuda_runtime.h>
|
||||
|
||||
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();
|
||||
}
|
||||
@@ -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 }
|
||||
}
|
||||
|
||||
@@ -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");
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
@@ -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<CudaStream>) -> 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<CudaStream>) -> 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<CudaStream>) -> 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}"
|
||||
);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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.
|
||||
|
||||
Reference in New Issue
Block a user