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

feat(sp5): Task A4 — Pearl 4 per-group Adam β1/β2/ε

Browse Source 
Per-param-group Adam hyperparameters from per-group gradient
direction-stability EMA + L2 norm. 8 SP4 param groups × 3 hparams = 24
ISV slots [226..250).

Two-kernel chain:
  grad_cosine_sim_update — per-group cosine_sim of curr vs prev grads
                          + L2 norm; writes back grad_curr → grad_prev
                          for next step's comparison.
  pearl_4_adam_hparams_update — β1/β2/ε from cosine + l2_norm.

Structural envelopes (Invariant 1 anchors per spec line 88):
  β1 ∈ [0.85, 0.95]; β2 ∈ [0.99, 0.9995]; ε ∈ [1e-10, 1e-6]

ALPHA_META migration (Option A — atomic shared-contract migration per
feedback_no_partial_refactor): apply_pearls_ad_kernel.cu and the
launch_apply_pearls Rust wrapper now take alpha_meta as a parameter.
All 45 existing call sites (SP4 + SP5 producers) pass the prior default
1.0e-3 explicitly via crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META.
Pearl 4's apply_pearls calls pass 5.0e-4 (half the default per spec
line 89) to limit β change rate.

Theoretical caveat: adaptive β breaks Adam's constant-β convergence
proof. Mitigations: structural envelopes + halved ALPHA_META + Pearls
A+D smoothing. Fall-back path defined: revert + ε-only adaptive
variant if Layer C destabilization observed.

New mapped-pinned buffer: grad_prev_buf_per_group [TOTAL_PARAMS] for
cosine-sim previous-step direction storage. Mapped-pinned i32 mirror
of grad_buf layout.

producer_step_scratch_buf grew 131 → 171 (40 new outputs: 8 cosine_sim
+ 8 l2_norm + 24 β1/β2/ε). wiener_state_buf already at 543 (A1 sized
for entire SP5 block).

StateResetRegistry: 4 new FoldReset entries: sp5_adam_beta1,
sp5_adam_beta2, sp5_adam_eps, sp5_grad_prev_buf. All sentinel 0 →
bootstrap to envelope midpoint via Pearls A+D + cosine_sim=0 fall-back
on first step of new fold.

Adam-launcher consumer migration deferred to Layer B.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-01 22:18:10 +02:00
parent 71a0275f54
commit 67dd414cb3
11 changed files with 744 additions and 9 deletions
Download Patch File Download Diff File Expand all files Collapse all files

12
crates/ml/build.rs
View File

@@ -430,6 +430,18 @@ fn main() {
// apply_pearls_ad_kernel then smooths via Pearls A+D into ISV[190..210).
// Must run AFTER pearl_1_atom (Q_VAR) AND pearl_3_sigma (NOISY_SIGMA).
"pearl_2_budget_kernel.cu",
// SP5 Task A4 (2026-05-01): per-group gradient cosine similarity + L2 norm.
// Single-block 8-thread kernel (one per param group: DqnTrunk/Value/Branches/
// IQN/IqlHigh/IqlLow/Attn/Curiosity). Reads grad_buf + grad_prev_buf;
// writes cosine_sim[8]→scratch[131..139) and l2_norm[8]→scratch[139..147).
// Also writebacks grad_curr → grad_prev_buf for the next step's comparison.
"grad_cosine_sim_kernel.cu",
// SP5 Task A4 (2026-05-01): Pearl 4 per-param-group Adam β1, β2, ε.
// Single-block 8-thread kernel reads cosine_sim + l2_norm from scratch;
// writes β1[8]→scratch[147..155), β2[8]→scratch[155..163), ε[8]→scratch[163..171).
// 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",
// 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

18
crates/ml/src/cuda_pipeline/apply_pearls_kernel.cu
View File

@@ -20,8 +20,13 @@
// pattern broke graph capture at `aux_label_scale_ema` mid-step
// (CUDA_ERROR_STREAM_CAPTURE_UNSUPPORTED — smoke `smoke-test-v9kjv`).
//
// Constants (`ALPHA_META`, `EPS_DIV`) MUST match
// `cuda_pipeline::sp4_wiener_ema::{ALPHA_META, EPS_DIV}` bit-for-bit.
// `ALPHA_META` is now passed as a kernel parameter so Pearl 4 can use
// a halved rate (5e-4) while all SP4 and other SP5 producers continue
// to use the SP4 default (1e-3). `EPS_DIV` remains a compile-time
// constant — it is numerical-stability-only and must never vary.
//
// Constants (`EPS_DIV`) MUST match
// `cuda_pipeline::sp4_wiener_ema::EPS_DIV` bit-for-bit.
// Single-source-of-truth lives in `sp4_wiener_ema.rs` for documentation;
// this duplicate is required because CUDA C++ can't include the Rust
// const directly. The new GPU oracle test asserts agreement within
@@ -34,14 +39,15 @@ extern "C" __global__ void apply_pearls_ad_kernel(
int isv_idx_base,
float* __restrict__ wiener_state_buf, /* [wiener_offset_base + 3*n_slots) packed [sv, dv, xl] triples */
int wiener_offset_base,
int n_slots
int n_slots,
float alpha_meta /* meta-EMA rate for Wiener variance tracking;
SP4 default = 1e-3; Pearl 4 uses 5e-4 */
) {
// Single-thread device-side loop. Every slot has its own state; no
// cross-slot parallelism is meaningful. Guard avoids redundant work
// if the launcher accidentally requests >1 thread.
if (blockIdx.x != 0 || threadIdx.x != 0) return;
constexpr float ALPHA_META = 1.0e-3f;
constexpr float EPS_DIV = 1.0e-8f;
for (int i = 0; i < n_slots; ++i) {
@@ -85,9 +91,9 @@ extern "C" __global__ void apply_pearls_ad_kernel(
const float x_lag_diff = step_obs - x_lag;
const float new_sample_var =
(1.0f - ALPHA_META) * sample_var + ALPHA_META * sample_diff * sample_diff;
(1.0f - alpha_meta) * sample_var + alpha_meta * sample_diff * sample_diff;
const float new_diff_var =
(1.0f - ALPHA_META) * diff_var + ALPHA_META * x_lag_diff * x_lag_diff;
(1.0f - alpha_meta) * diff_var + alpha_meta * x_lag_diff * x_lag_diff;
// EPS_DIV protects 0/0 at t=1 when both variances start near 0.
const float alpha_opt =

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

@@ -318,6 +318,24 @@ static SP5_PEARL_3_SIGMA_CUBIN: &[u8] =
static SP5_PEARL_2_BUDGET_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/pearl_2_budget_kernel.cubin"));
/// SP5 Task A4 (2026-05-01): auxiliary gradient cosine similarity cubin.
/// Single-block 8-thread kernel (one thread per SP4 param group). Reads
/// grad_buf + grad_prev_buf_per_group + group_param_offsets_buf [9 i32];
/// computes per-group cos_sim and l2_norm; writes to scratch_buf[131..147).
/// Also writes grad_curr → grad_prev_buf_writeback for next step's comparison.
/// No atomicAdd (feedback_no_atomicadd).
static SP5_GRAD_COSINE_SIM_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/grad_cosine_sim_kernel.cubin"));
/// SP5 Task A4 (2026-05-01): Pearl 4 per-group Adam β1/β2/ε cubin.
/// Single-block 8-thread kernel (one thread per param group). Reads
/// scratch_buf[131..147) (cosine_sim + l2_norm); computes β1/β2/ε within
/// structural envelopes (Invariant 1 anchors); writes to scratch_buf[147..171).
/// Chained apply_pearls_ad_kernel (24 launches, ALPHA_META=5e-4) smooths into
/// ISV[ADAM_BETA1_BASE..ADAM_EPS_BASE+8). No atomicAdd.
static SP5_PEARL_4_ADAM_HPARAMS_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/pearl_4_adam_hparams_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
@@ -901,7 +919,7 @@ pub const SP5_WIENER_TOTAL_FLOATS: usize =
(SP4_PRODUCER_COUNT + crate::cuda_pipeline::sp5_isv_slots::SP5_PRODUCER_COUNT)
* SP4_WIENER_FLOATS_PER_SLOT;
/// SP5 Task A1+A2+A3: producer_step_scratch_buf total slot count after growth.
/// SP5 Task A1+A2+A3+A4: producer_step_scratch_buf total slot count after growth.
/// SP4 had 71 scratch slots [0..71); Task A1 adds:
/// 16 floats for q_branch_stats output (SCRATCH_Q_STATS=71, 4 outputs × 4 branches)
/// 16 floats for pearl_1_atom output (SCRATCH_ATOM_OUT=87, 4 outputs × 4 branches)
@@ -915,8 +933,14 @@ pub const SP5_WIENER_TOTAL_FLOATS: usize =
/// budget_cql[4] → SCRATCH_PEARL_2_CQL=119
/// budget_ens[4] → SCRATCH_PEARL_2_ENS=123
/// flatness[4] → SCRATCH_PEARL_2_FLATNESS=127
/// Combined: 71 + 16 + 16 + 4 + 4 + 20 = 131 scratch slots [0..131).
pub const SP5_SCRATCH_TOTAL: usize = 131;
/// Task A4 adds:
/// 8 floats for grad_cosine_sim output (SCRATCH_PEARL_4_COSINE=131, cos_sim × 8 groups)
/// 8 floats for grad_l2_norm output (SCRATCH_PEARL_4_L2=139, l2_norm × 8 groups)
/// 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;
/// SP5 Task A2: scratch index base for pearl_3_sigma sigma[4] output block.
/// Slots [103..107): per-branch NoisyNet sigma.
@@ -947,6 +971,34 @@ pub const SCRATCH_PEARL_2_ENS: usize = 123;
/// Slots [127..131): per-branch flatness signal (var_q / (σ² + EPS_DIV), clamped [0,1]).
pub const SCRATCH_PEARL_2_FLATNESS: usize = 127;
/// SP5 Task A4: scratch index base for grad_cosine_sim_kernel cosine_sim[8] output block.
/// Slots [131..139): per-param-group gradient direction cosine similarity vs previous step.
/// Written by `grad_cosine_sim_update`; consumed by `pearl_4_adam_hparams_update`.
pub const SCRATCH_PEARL_4_COSINE: usize = 131;
/// SP5 Task A4: scratch index base for grad_cosine_sim_kernel l2_norm[8] output block.
/// Slots [139..147): per-param-group L2 norm of current-step gradient.
/// Written by `grad_cosine_sim_update`; consumed by `pearl_4_adam_hparams_update`.
pub const SCRATCH_PEARL_4_L2: usize = 139;
/// SP5 Task A4: scratch index base for pearl_4_adam_hparams_kernel β1[8] output block.
/// Slots [147..155): per-param-group adaptive Adam β1 ∈ [0.85, 0.95].
/// Written by `pearl_4_adam_hparams_update`; consumed by apply_pearls_ad_kernel →
/// ISV[ADAM_BETA1_BASE..ADAM_BETA1_BASE+8).
pub const SCRATCH_PEARL_4_BETA1: usize = 147;
/// SP5 Task A4: scratch index base for pearl_4_adam_hparams_kernel β2[8] output block.
/// Slots [155..163): per-param-group adaptive Adam β2 ∈ [0.99, 0.9995].
/// Written by `pearl_4_adam_hparams_update`; consumed by apply_pearls_ad_kernel →
/// ISV[ADAM_BETA2_BASE..ADAM_BETA2_BASE+8).
pub const SCRATCH_PEARL_4_BETA2: usize = 155;
/// SP5 Task A4: scratch index base for pearl_4_adam_hparams_kernel ε[8] output block.
/// Slots [163..171): per-param-group adaptive Adam ε ∈ [1e-10, 1e-6].
/// Written by `pearl_4_adam_hparams_update`; consumed by apply_pearls_ad_kernel →
/// ISV[ADAM_EPS_BASE..ADAM_EPS_BASE+8).
pub const SCRATCH_PEARL_4_EPS: usize = 163;
/// SP5 Task A1: scratch index base for q_branch_stats 16-float output block.
/// Slots [71..87): mean, max_abs_dev, var, entropy × 4 branches.
/// SP4 slot 70 = q_dir_grad_p99; SP5 starts at 71 (no collision).
@@ -4095,6 +4147,40 @@ pub struct GpuDqnTrainer {
/// No atomicAdd (feedback_no_atomicadd), no CPU compute (feedback_no_cpu_compute_strict).
/// Loaded from `pearl_2_budget_kernel.cubin`.
pearl_2_budget_kernel: CudaFunction,
// ── SP5 Task A4: Pearl 4 per-group Adam β1/β2/ε kernels ─────────────
/// SP5 Task A4 (2026-05-01): auxiliary gradient cosine similarity kernel.
/// Single-block 8-thread kernel (one thread per SP4 param group).
/// Reads `grad_buf [total_params]` + `grad_prev_buf_per_group [total_params]`
/// + `group_param_offsets_buf [9]` (prefix sums in element units); computes
/// per-group cos_sim = dot(g,g_prev)/(|g|×|g_prev|+EPS_COS) and l2_norm = |g|;
/// writes to `producer_step_scratch_buf[131..147)`. Also writebacks g_curr →
/// grad_prev_buf_per_group for the next step.
/// No atomicAdd (feedback_no_atomicadd), no CPU compute (feedback_no_cpu_compute_strict).
/// Loaded from `grad_cosine_sim_kernel.cubin`.
grad_cosine_sim_kernel: CudaFunction,
/// SP5 Task A4 (2026-05-01): Pearl 4 per-group Adam β1/β2/ε controller kernel.
/// Single-block 8-thread kernel (one thread per param group). Reads per-group
/// cosine_sim + l2_norm from scratch_buf[131..147); computes
/// β1=0.85+0.10×stability ∈ [0.85,0.95],
/// β2=0.99+0.0095×stability ∈ [0.99,0.9995],
/// ε=clamp(grad_norm×1e-7, 1e-10, 1e-6); writes to scratch_buf[147..171).
/// Chained apply_pearls_ad_kernel (24 single-slot launches, ALPHA_META=5e-4)
/// smooths via Pearls A+D into ISV[ADAM_BETA1_BASE..ADAM_EPS_BASE+8).
/// Structural envelopes (Invariant 1 anchors per spec line 88).
/// No atomicAdd, no CPU compute.
/// Loaded from `pearl_4_adam_hparams_kernel.cubin`.
pearl_4_adam_hparams_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
/// on first step → β1/β2 start at envelope midpoints).
pub(crate) grad_prev_buf_per_group: MappedF32Buffer,
/// SP5 Task A4 (2026-05-01): param-group boundary prefix sums for cosine-sim kernel.
/// [9 i32] with group g in element range [offsets[g], offsets[g+1]).
/// Values are in f32-element units (not bytes): divide padded byte offsets by 4.
/// Host-written once at construction; static for the lifetime of the trainer.
pub(crate) group_param_offsets_buf: MappedI32Buffer,
/// SP5 Task A1 (2026-05-01): per-branch clip-count buffer for Pearl 1
/// headroom controller. [4 i32], one element per branch. Zero-initialized
/// at construction; populated by `atoms_update_kernel` in Layer B.
@@ -9491,6 +9577,7 @@ impl GpuDqnTrainer {
self.wiener_state_buf.dev_ptr,
(SCRATCH_IDX * 3) as i32,
1,
crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
)?;
}
Ok(())
@@ -9575,6 +9662,7 @@ impl GpuDqnTrainer {
self.wiener_state_buf.dev_ptr,
wiener_offset,
1,
crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
)?;
}
Ok(())
@@ -9734,6 +9822,7 @@ impl GpuDqnTrainer {
self.wiener_state_buf.dev_ptr,
((ATOM_POS_BOUND_BASE - SP4_SLOT_BASE) * 3) as i32,
SP4_BRANCH_COUNT as i32,
crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
)?;
}
Ok(())
@@ -9976,6 +10065,7 @@ impl GpuDqnTrainer {
self.wiener_state_buf.dev_ptr,
((isv_idx - SP4_SLOT_BASE) * 3) as i32,
1,
crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
)
}
};
@@ -10169,6 +10259,7 @@ impl GpuDqnTrainer {
self.wiener_state_buf.dev_ptr,
((isv_idx - SP4_SLOT_BASE) * 3) as i32,
1,
crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
)?;
}
Ok(())
@@ -10384,6 +10475,7 @@ impl GpuDqnTrainer {
self.wiener_state_buf.dev_ptr,
((isv_idx - SP4_SLOT_BASE) * 3) as i32,
1,
crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
)?;
}
Ok(())
@@ -10524,6 +10616,7 @@ impl GpuDqnTrainer {
isv_dev, isv_idx,
wiener_dev, wiener_off,
1,
crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
)?;
}
}
@@ -10545,6 +10638,7 @@ impl GpuDqnTrainer {
isv_dev, isv_entropy,
wiener_dev, wiener_entropy,
1,
crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
)?;
}
// q_var → ISV Q_VAR_PER_BRANCH_BASE + b
@@ -10559,6 +10653,7 @@ impl GpuDqnTrainer {
isv_dev, isv_var,
wiener_dev, wiener_var,
1,
crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
)?;
}
}
@@ -10643,6 +10738,7 @@ impl GpuDqnTrainer {
isv_dev, isv_sigma,
wiener_dev, wiener_sigma,
1,
crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
)?;
}
// SF → ISV SIGMA_FRACTION_BASE + b
@@ -10657,6 +10753,7 @@ impl GpuDqnTrainer {
isv_dev, isv_sf,
wiener_dev, wiener_sf,
1,
crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
)?;
}
}
@@ -10756,6 +10853,7 @@ impl GpuDqnTrainer {
isv_dev, isv_idx,
wiener_dev, wiener_off,
1,
crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
)?;
}
}
@@ -10764,6 +10862,156 @@ impl GpuDqnTrainer {
Ok(())
}
/// SP5 Task A4 (2026-05-01): launch the two-kernel Pearl 4 chain +
/// 24 `apply_pearls_ad_kernel` launches to smooth β1/β2/ε into ISV.
///
/// Execution sequence (all on `self.stream`):
/// 1. `grad_cosine_sim_update` (8 threads, 1 block): reads
/// `grad_buf [total_params]` + `grad_prev_buf_per_group [total_params]`
/// + `group_param_offsets_buf [9]`; writes per-group cosine_sim[8]
/// to scratch[131..139) and l2_norm[8] to scratch[139..147); also
/// writes grad_curr → grad_prev_buf_per_group for the next step's comparison.
/// 2. `pearl_4_adam_hparams_update` (8 threads, 1 block): reads
/// scratch[131..147); writes β1[8]→scratch[147..155),
/// β2[8]→scratch[155..163), ε[8]→scratch[163..171).
/// 3. 24 `apply_pearls_ad_kernel` launches (ALPHA_META=5e-4, half the
/// SP4 default): smooths β1[8]/β2[8]/ε[8] via Pearls A+D into
/// ISV[ADAM_BETA1_BASE..ADAM_BETA1_BASE+8),
/// ISV[ADAM_BETA2_BASE..ADAM_BETA2_BASE+8),
/// ISV[ADAM_EPS_BASE..ADAM_EPS_BASE+8).
///
/// Pearl 4 is independent of Pearls 1/2/3 outputs — depends only on
/// grad_buf state. Can run in any order relative to other SP5 launchers.
///
/// Structural envelopes (Invariant 1 anchors per spec line 88):
/// β1 ∈ [0.85, 0.95]; β2 ∈ [0.99, 0.9995]; ε ∈ [1e-10, 1e-6].
/// Fold-boundary reset zeroes grad_prev_buf_per_group (Pearl A sentinel):
/// first step cosine_sim=0 → β1=0.85, β2=0.99 (low-stability envelope).
/// Layer A only — Adam optimizer consumer migration deferred to Layer B.
pub(crate) fn launch_sp5_pearl_4_adam_hparams(&self) -> Result<(), MLError> {
use crate::cuda_pipeline::sp5_isv_slots::{
ADAM_BETA1_BASE, ADAM_BETA2_BASE, ADAM_EPS_BASE, SP5_SLOT_BASE,
};
use crate::cuda_pipeline::sp4_wiener_ema::launch_apply_pearls;
const ALPHA_META_PEARL_4: f32 = 5.0e-4_f32; // half SP4 default → slow β change rate
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;
// Wiener base offset for SP5 slots: SP4_PRODUCER_COUNT * 3 Wiener slots
// already consumed; SP5 slots start at offset (SP4_PRODUCER_COUNT * 3).
let base_wiener_offset: i32 = (SP4_PRODUCER_COUNT * 3) as i32;
let grad_buf_ptr = self.ptrs.grad_buf;
let grad_prev_ptr = self.grad_prev_buf_per_group.dev_ptr;
let group_offsets_ptr = self.group_param_offsets_buf.dev_ptr;
let cosine_idx_base = SCRATCH_PEARL_4_COSINE as i32; // 131
let l2_norm_idx_base = SCRATCH_PEARL_4_L2 as i32; // 139
let beta1_idx_base = SCRATCH_PEARL_4_BETA1 as i32; // 147
let beta2_idx_base = SCRATCH_PEARL_4_BETA2 as i32; // 155
let eps_idx_base = SCRATCH_PEARL_4_EPS as i32; // 163
// Step 1: grad_cosine_sim_update — 8 threads, 1 block.
// Reads grad_buf + grad_prev_buf_per_group + group_param_offsets;
// writes cosine_sim[8]→scratch[131..139) and l2_norm[8]→scratch[139..147).
// Also writebacks grad_curr → grad_prev_buf_per_group.
let cfg_8t = LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (8, 1, 1),
shared_mem_bytes: 0,
};
unsafe {
self.stream
.launch_builder(&self.grad_cosine_sim_kernel)
.arg(&grad_buf_ptr)
.arg(&grad_prev_ptr)
.arg(&group_offsets_ptr)
.arg(&scratch_dev)
.arg(&cosine_idx_base)
.arg(&l2_norm_idx_base)
.arg(&grad_prev_ptr) // writeback = same buffer
.launch(cfg_8t)
.map_err(|e| MLError::ModelError(format!("grad_cosine_sim_update launch: {e}")))?;
}
// Step 2: pearl_4_adam_hparams_update — 8 threads, 1 block.
// Reads scratch[131..147) (cosine_sim + l2_norm); writes
// β1[8]→scratch[147..155), β2[8]→scratch[155..163), ε[8]→scratch[163..171).
// Stream-ordered with Step 1 (same stream, no explicit sync needed).
let cosine_ptr = scratch_dev + (SCRATCH_PEARL_4_COSINE * std::mem::size_of::<f32>()) as u64;
let l2_ptr = scratch_dev + (SCRATCH_PEARL_4_L2 * std::mem::size_of::<f32>()) as u64;
unsafe {
self.stream
.launch_builder(&self.pearl_4_adam_hparams_kernel)
.arg(&cosine_ptr)
.arg(&l2_ptr)
.arg(&scratch_dev)
.arg(&beta1_idx_base)
.arg(&beta2_idx_base)
.arg(&eps_idx_base)
.launch(cfg_8t)
.map_err(|e| MLError::ModelError(format!("pearl_4_adam_hparams_update launch: {e}")))?;
}
// Step 3: 24 apply_pearls_ad_kernel launches (ALPHA_META=5e-4).
// β1[8]: scratch[147..155) → ISV[ADAM_BETA1_BASE..ADAM_BETA1_BASE+8)
for g in 0..8_usize {
let scratch_idx = (SCRATCH_PEARL_4_BETA1 + g) as i32;
let isv_idx = (ADAM_BETA1_BASE + g) as i32;
let wiener_off = base_wiener_offset + (isv_idx - SP5_SLOT_BASE as i32) * 3;
unsafe {
launch_apply_pearls(
&self.stream,
&self.apply_pearls_ad_kernel,
scratch_dev, scratch_idx,
isv_dev, isv_idx,
wiener_dev, wiener_off,
1,
ALPHA_META_PEARL_4,
)?;
}
}
// β2[8]: scratch[155..163) → ISV[ADAM_BETA2_BASE..ADAM_BETA2_BASE+8)
for g in 0..8_usize {
let scratch_idx = (SCRATCH_PEARL_4_BETA2 + g) as i32;
let isv_idx = (ADAM_BETA2_BASE + g) as i32;
let wiener_off = base_wiener_offset + (isv_idx - SP5_SLOT_BASE as i32) * 3;
unsafe {
launch_apply_pearls(
&self.stream,
&self.apply_pearls_ad_kernel,
scratch_dev, scratch_idx,
isv_dev, isv_idx,
wiener_dev, wiener_off,
1,
ALPHA_META_PEARL_4,
)?;
}
}
// ε[8]: scratch[163..171) → ISV[ADAM_EPS_BASE..ADAM_EPS_BASE+8)
for g in 0..8_usize {
let scratch_idx = (SCRATCH_PEARL_4_EPS + g) as i32;
let isv_idx = (ADAM_EPS_BASE + g) as i32;
let wiener_off = base_wiener_offset + (isv_idx - SP5_SLOT_BASE as i32) * 3;
unsafe {
launch_apply_pearls(
&self.stream,
&self.apply_pearls_ad_kernel,
scratch_dev, scratch_idx,
isv_dev, isv_idx,
wiener_dev, wiener_off,
1,
ALPHA_META_PEARL_4,
)?;
}
}
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].
///
@@ -10889,6 +11137,7 @@ impl GpuDqnTrainer {
self.wiener_state_buf.dev_ptr,
(SCRATCH_FIRST * 3) as i32,
num_groups as i32,
crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
)?;
}
Ok(())
@@ -10965,6 +11214,7 @@ impl GpuDqnTrainer {
self.wiener_state_buf.dev_ptr,
(SCRATCH_IDX_NEXT_BAR * 3) as i32,
2,
crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
)?;
}
Ok(())
@@ -11231,6 +11481,7 @@ impl GpuDqnTrainer {
self.wiener_state_buf.dev_ptr,
(SCRATCH_IDX_UTIL_BASE * 3) as i32,
(MOE_NUM_EXPERTS + 1) as i32,
crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
)?;
}
@@ -11555,6 +11806,7 @@ impl GpuDqnTrainer {
self.wiener_state_buf.dev_ptr,
(SCRATCH_IDX_LABEL_SCALE * 3) as i32,
1,
crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
)?;
}
}
@@ -11869,6 +12121,7 @@ impl GpuDqnTrainer {
self.wiener_state_buf.dev_ptr,
(SCRATCH_FIRST * 3) as i32,
4,
crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
)?;
}
Ok(())
@@ -12801,6 +13054,66 @@ impl GpuDqnTrainer {
.map_err(|e| MLError::ModelError(format!("SP5 action_offsets_buf alloc: {e}")))?;
action_offsets_buf.write_from_slice(&[0_i32, 4, 7, 10]);
// SP5 Task A4 (2026-05-01): load grad_cosine_sim_kernel (auxiliary
// gradient cosine similarity). Single-block 8-thread, one thread per
// SP4 param group. Reads grad_buf + grad_prev_buf_per_group.
let grad_cosine_sim_kernel = {
let module = stream.context().load_cubin(SP5_GRAD_COSINE_SIM_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!("sp5 grad_cosine_sim cubin load: {e}")))?;
module.load_function("grad_cosine_sim_update")
.map_err(|e| MLError::ModelError(format!("grad_cosine_sim_update load: {e}")))?
};
// SP5 Task A4 (2026-05-01): load pearl_4_adam_hparams_kernel.
// Single-block 8-thread, one thread per SP4 param group. Reads
// cosine_sim + l2_norm from scratch; writes β1/β2/ε to scratch.
let pearl_4_adam_hparams_kernel = {
let module = stream.context().load_cubin(SP5_PEARL_4_ADAM_HPARAMS_CUBIN.to_vec())
.map_err(|e| MLError::ModelError(format!("sp5 pearl_4_adam_hparams cubin load: {e}")))?;
module.load_function("pearl_4_adam_hparams_update")
.map_err(|e| MLError::ModelError(format!("pearl_4_adam_hparams_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).
// `total_params` was computed at constructor entry (line ~12337).
let grad_prev_buf_per_group = unsafe { MappedF32Buffer::new(total_params) }
.map_err(|e| MLError::ModelError(format!("SP5 grad_prev_buf_per_group alloc: {e}")))?;
// SP5 Task A4 (2026-05-01): allocate group_param_offsets_buf [9 i32].
// Prefix sums in f32-element units (not bytes): [offsets[0], ..., offsets[8]]
// where group g covers [offsets[g], offsets[g+1]).
// Groups match SP4_PARAM_GROUP_COUNT=8:
// 0=DqnTrunk [0..trunk_param_count)
// 1=DqnValue [trunk_param_count..trunk+value_count)
// 2=DqnBranches [trunk+value..trunk+value+branch_count)
// 3=Iqn, 4=IqlHigh, 5=IqlLow, 6=Attn, 7=Curiosity
// For aux groups 3-7: they are separate grad buffers (not slices of grad_buf)
// so their offsets are set to [total_params, total_params] — zero-width range
// signals the kernel to skip those groups (thread loop produces dot=0,
// norm=0, cosine_sim=0 → β1/β2 snap to stable envelope midpoints).
// Only DqnTrunk/Value/Branches (groups 0-2) cover the contiguous main grad_buf.
let param_sizes = compute_param_sizes(&config);
let trunk_end = (padded_byte_offset(&param_sizes, 13) / 4) as i32;
let value_end = (padded_byte_offset(&param_sizes, 17) / 4) as i32;
let branch_end = (padded_byte_offset(&param_sizes, 33) / 4) as i32;
// Aux groups 3-7 sit outside the contiguous grad_buf slab → zero-width sentinel.
let tp = total_params as i32;
let group_param_offsets_buf = unsafe { MappedI32Buffer::new(9) }
.map_err(|e| MLError::ModelError(format!("SP5 group_param_offsets_buf alloc: {e}")))?;
group_param_offsets_buf.write_from_slice(&[
0, // group 0 DqnTrunk start
trunk_end, // group 1 DqnValue start / DqnTrunk end
value_end, // group 2 DqnBranches start / DqnValue end
branch_end, // group 3 Iqn start / DqnBranches end
tp, // group 4 IqlHigh start — zero-width sentinel for aux groups
tp, // group 5 IqlLow start
tp, // group 6 Attn start
tp, // group 7 Curiosity start
tp, // sentinel: offsets[8] = end of group 7
]);
// 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`.
@@ -15906,6 +16219,10 @@ impl GpuDqnTrainer {
pearl_1_atom_kernel,
pearl_3_sigma_kernel,
pearl_2_budget_kernel,
grad_cosine_sim_kernel,
pearl_4_adam_hparams_kernel,
grad_prev_buf_per_group,
group_param_offsets_buf,
atoms_clip_count_buf,
action_counts_buf,
action_offsets_buf,

1
crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
View File

@@ -2399,6 +2399,7 @@ impl GpuExperienceCollector {
self.reward_component_pearls_wiener_dev_ptr,
(SCRATCH_FIRST * 3) as i32,
REWARD_COMPONENT_COUNT as i32,
crate::cuda_pipeline::sp4_wiener_ema::ALPHA_META,
)?;
}

62
crates/ml/src/cuda_pipeline/grad_cosine_sim_kernel.cu Normal file
View File

@@ -0,0 +1,62 @@
// crates/ml/src/cuda_pipeline/grad_cosine_sim_kernel.cu
//
// SP5 Task A4 auxiliary: per-group gradient cosine similarity.
// Computes cos_sim(grad_curr, grad_prev) for each of 8 SP4 param groups.
// Also computes per-group L2 norm of grad_curr.
// Writes grad_curr direction as grad_prev for the next step.
//
// Single-block, 8-thread (one per group). Each thread reduces over its
// group's parameter range to compute:
// cos_sim = (g_curr · g_prev) / (|g_curr| × |g_prev| + EPS_COS)
// l2_norm = |g_curr|
//
// No atomicAdd per feedback_no_atomicadd — one thread per group with
// independent writes. grad_prev_buf_writeback is written by the same
// thread that reads it (no race).
#include <cuda_runtime.h>
extern "C" __global__ void grad_cosine_sim_update(
const float* __restrict__ grad_buf, // [total_params] current step's gradient
const float* __restrict__ grad_prev_buf, // [total_params] previous step's gradient
const int* __restrict__ group_param_offsets, // [9] prefix sums: group g in [offsets[g], offsets[g+1])
float* __restrict__ scratch_out, // cosine_sim[8] at [cosine_idx_base..+8)
// l2_norm[8] at [l2_norm_idx_base..+8)
int cosine_idx_base,
int l2_norm_idx_base,
float* __restrict__ grad_prev_buf_writeback // [total_params] overwrite prev with curr for next step
) {
int g = threadIdx.x;
if (blockIdx.x != 0 || g >= 8) return;
int start = group_param_offsets[g];
int end = group_param_offsets[g + 1];
const float EPS_COS = 1e-12f;
float dot = 0.0f, norm_curr_sq = 0.0f, norm_prev_sq = 0.0f;
for (int i = start; i < end; ++i) {
float gc = grad_buf[i];
float gp = grad_prev_buf[i];
dot += gc * gp;
norm_curr_sq += gc * gc;
norm_prev_sq += gp * gp;
}
float norm_curr = sqrtf(norm_curr_sq);
float norm_prev = sqrtf(norm_prev_sq);
float cos_sim = dot / (norm_curr * norm_prev + EPS_COS);
scratch_out[cosine_idx_base + g] = cos_sim;
scratch_out[l2_norm_idx_base + g] = norm_curr;
// Writeback grad_curr to grad_prev for the next step's comparison.
// Same thread that read grad_prev writes back grad_curr — no race.
for (int i = start; i < end; ++i) {
grad_prev_buf_writeback[i] = grad_buf[i];
}
// Make all writes globally visible before pearl_4_adam_hparams_update
// reads scratch_out and before the next step reads grad_prev_buf.
__threadfence_system();
}

54
crates/ml/src/cuda_pipeline/pearl_4_adam_hparams_kernel.cu Normal file
View File

@@ -0,0 +1,54 @@
// crates/ml/src/cuda_pipeline/pearl_4_adam_hparams_kernel.cu
//
// SP5 Task A4: Pearl 4 — per-param-group Adam β1, β2, ε from gradient
// direction-stability + L2 norm.
// 8-thread single-block (one per param group: DqnTrunk/Value/Branches/IQN/
// IqlHigh/IqlLow/Attn/Curiosity).
//
// Producer signal: per-group cosine similarity of current vs previous-step
// gradient direction (written by grad_cosine_sim_kernel). High stability
// (cosine→1) → β2 → 0.9995 (long memory); low stability (cosine→0 or →-1)
// → β2 → 0.99 (short memory). β1 likewise scales with stability.
//
// Structural envelopes (Invariant 1 anchors per spec line 88):
// β1 ∈ [0.85, 0.95]; β2 ∈ [0.99, 0.9995]; ε ∈ [1e-10, 1e-6]
//
// THEORETICAL CAVEAT: adaptive β breaks Adam's constant-β convergence
// proof. Mitigations:
// - Structural β envelopes (Invariant 1 anchors)
// - ALPHA_META=5e-4 in apply_pearls_ad_kernel (half the SP4 default → slow change rate)
// - Pearl A sentinel bootstrap on fold boundary (cosine_sim=0 → envelope midpoints)
// If Layer C destabilization observed, fall-back path: revert to ε-only adaptive.
#include <cuda_runtime.h>
extern "C" __global__ void pearl_4_adam_hparams_update(
const float* __restrict__ grad_cosine_sim_per_group, // [8] from grad_cosine_sim_kernel
const float* __restrict__ grad_l2_norm_per_group, // [8] from grad_cosine_sim_kernel
float* __restrict__ scratch_out,
int beta1_idx_base,
int beta2_idx_base,
int eps_idx_base
) {
int g = threadIdx.x;
if (blockIdx.x != 0 || g >= 8) return;
float stability = fmaxf(0.0f, fminf(1.0f, grad_cosine_sim_per_group[g]));
float grad_norm = grad_l2_norm_per_group[g];
// Structural envelopes (Invariant 1 anchors per spec line 88).
// β1 scales linearly with stability: [0.85, 0.95]
float beta1 = 0.85f + 0.10f * stability;
// β2 scales linearly with stability: [0.99, 0.9995]
float beta2 = 0.99f + 0.0095f * stability;
// ε scales with grad norm to maintain numerical headroom; envelope clamp.
// grad_norm * 1e-7 maps typical grad norms [1e-3, 10.0] to [1e-10, 1e-6].
float eps = fmaxf(1e-10f, fminf(1e-6f, grad_norm * 1e-7f));
scratch_out[beta1_idx_base + g] = beta1;
scratch_out[beta2_idx_base + g] = beta2;
scratch_out[eps_idx_base + g] = eps;
// Make writes globally visible before apply_pearls_ad_kernel consumes them.
__threadfence_system();
}

6
crates/ml/src/cuda_pipeline/sp4_wiener_ema.rs
View File

@@ -146,6 +146,10 @@ pub fn pearls_ad_update(
/// typically 1..9. Single-thread device-side loop is correct because each
/// slot has its own state — no cross-slot parallelism is meaningful.
///
/// `alpha_meta`: meta-EMA rate for Wiener variance tracking. SP4 default
/// and all SP5 producers except Pearl 4 use [`ALPHA_META`] (1e-3). Pearl 4
/// uses 5e-4 (half the default per spec line 89 to limit β change rate).
///
/// # Safety
///
/// All three buffer pointers MUST be valid device pointers. Index ranges
@@ -168,6 +172,7 @@ pub(crate) unsafe fn launch_apply_pearls(
wiener_state_dev: u64,
wiener_offset_base: i32,
n_slots: i32,
alpha_meta: f32,
) -> Result<(), crate::MLError> {
use cudarc::driver::{LaunchConfig, PushKernelArg};
@@ -197,6 +202,7 @@ pub(crate) unsafe fn launch_apply_pearls(
.arg(&wiener_state_dev)
.arg(&wiener_offset_base)
.arg(&n_slots)
.arg(&alpha_meta)
.launch(cfg)
.map_err(|e| {
crate::MLError::ModelError(format!("apply_pearls_ad_kernel launch: {e}"))

24
crates/ml/src/trainers/dqn/state_reset_registry.rs
View File

@@ -627,6 +627,30 @@ impl StateResetRegistry {
category: ResetCategory::FoldReset,
description: "GpuDqnTrainer.wiener_state_buf SP5 block [330 floats = 110 SP5 producers × {sample_var, diff_var, x_lag}] at offsets [213..543). SP5 Pearl D Wiener-EMA state for Task A1's 24 producers (16 Pearl 1 + 8 shared signal). Bulk write_bytes(0) at fold boundary covers the full wiener_state_buf [0..543) (SP4 + SP5 together) so Pearl A's sentinel branch fires on the new fold's first producer launch. Both SP4 and SP5 blocks reset atomically — wiener_state_buf grew from 213 to 543 floats in Task A1.",
},
// SP5 Task A4: per-group Adam β1/β2/ε ISV slots (Pearl 4).
// ISV[226..250) = 24 floats (8 groups × 3 hyperparams).
// Pearl A sentinel 0 at fold boundary → first step fires first-observation
// replacement (cosine_sim=0 → low-stability envelope midpoints).
RegistryEntry {
name: "sp5_adam_beta1",
category: ResetCategory::FoldReset,
description: "ISV[ADAM_BETA1_BASE=226..234) — per-group Adam β1 (gradient direction stability). Pearl 4 producer: cosine_sim→[0,1] → β1 = 0.85 + 0.10×stability ∈ [0.85, 0.95]. SP5 Pearl A sentinel 0 at fold boundary (Task A4).",
},
RegistryEntry {
name: "sp5_adam_beta2",
category: ResetCategory::FoldReset,
description: "ISV[ADAM_BETA2_BASE=234..242) — per-group Adam β2 (gradient direction stability). Pearl 4 producer: cosine_sim→[0,1] → β2 = 0.99 + 0.0095×stability ∈ [0.99, 0.9995]. SP5 Pearl A sentinel 0 at fold boundary (Task A4).",
},
RegistryEntry {
name: "sp5_adam_eps",
category: ResetCategory::FoldReset,
description: "ISV[ADAM_EPS_BASE=242..250) — per-group Adam ε (grad L2 norm scaled). Pearl 4 producer: ε = clamp(grad_norm×1e-7, 1e-10, 1e-6). SP5 Pearl A sentinel 0 at fold boundary (Task A4).",
},
RegistryEntry {
name: "sp5_grad_prev_buf",
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).",
},
];
Self { entries }
}

7
crates/ml/src/trainers/dqn/trainer/training_loop.rs
View File

@@ -3603,6 +3603,13 @@ impl DQNTrainer {
if let Err(e) = fused.trainer().launch_sp5_pearl_2_budget() {
tracing::warn!(error = %e, "SP5 Pearl 2 producer launch failed");
}
// SP5 Task A4: pearl_4_adam_hparams → apply_pearls_ad ×24
// Reads grad_buf (populated in this step's backward pass) + grad_prev_buf.
// No ordering dependency on Pearls 1/2/3 — reads gradient state, not ISV.
// Writes ISV[226..250) = β1[8] + β2[8] + ε[8] (ALPHA_META_PEARL_4=5e-4).
if let Err(e) = fused.trainer().launch_sp5_pearl_4_adam_hparams() {
tracing::warn!(error = %e, "SP5 Pearl 4 producer launch failed");
}
}
}

244
crates/ml/tests/sp5_producer_unit_tests.rs
View File

@@ -33,6 +33,9 @@ const SP5_PEARL_3_SIGMA_CUBIN: &[u8] =
const SP5_PEARL_2_BUDGET_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/pearl_2_budget_kernel.cubin"));
const SP5_PEARL_4_ADAM_HPARAMS_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/pearl_4_adam_hparams_kernel.cubin"));
// ── Helpers ───────────────────────────────────────────────────────────────────
fn load_pearl_3_sigma_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
@@ -55,6 +58,16 @@ fn load_pearl_2_budget_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
.expect("load pearl_2_budget_update function")
}
fn load_pearl_4_adam_hparams_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
let module = stream
.context()
.load_cubin(SP5_PEARL_4_ADAM_HPARAMS_CUBIN.to_vec())
.expect("load pearl_4_adam_hparams_kernel cubin");
module
.load_function("pearl_4_adam_hparams_update")
.expect("load pearl_4_adam_hparams_update function")
}
fn make_test_stream() -> Arc<CudaStream> {
let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?");
ctx.default_stream()
@@ -861,3 +874,234 @@ fn pearl_2_budget_sharp_regime_iqn_dominates() {
);
}
}
// ── Tests 7 + 8: pearl_4_adam_hparams_update ──────────────────────────────────
/// SP5 Task A4 unit test: `pearl_4_adam_hparams_update` with high stability input.
///
/// Synthetic input: cosine_sim=[1.0; 8] (maximum stability), grad_norm=[1.0; 8].
///
/// Analytical expected:
/// stability = fmaxf(0, fminf(1, cosine_sim)) = 1.0
/// β1 = 0.85 + 0.10 × 1.0 = 0.95 (long-memory ceiling)
/// β2 = 0.99 + 0.0095 × 1.0 = 0.9995 (long-memory ceiling)
/// ε = clamp(1.0 × 1e-7, 1e-10, 1e-6) = 1e-7 → clamped to 1e-7 (within [1e-10, 1e-6])
///
/// Structural envelopes (Invariant 1 anchors): β1∈[0.85,0.95], β2∈[0.99,0.9995].
/// No CPU reference oracle per `feedback_no_cpu_test_fallbacks.md`.
#[test]
#[ignore = "requires GPU"]
fn pearl_4_adam_hparams_high_stability_yields_long_memory() {
// Scratch layout (Pearl 4 outputs):
// [0..8) = β1[8] (beta1_idx_base=0)
// [8..16) = β2[8] (beta2_idx_base=8)
// [16..24) = ε[8] (eps_idx_base=16)
const BETA1_BASE: usize = 0;
const BETA2_BASE: usize = 8;
const EPS_BASE: usize = 16;
const SCRATCH_SIZE: usize = 24;
const N_GROUPS: usize = 8;
let stream = make_test_stream();
let kernel = load_pearl_4_adam_hparams_kernel(&stream);
// cosine_sim = 1.0 for all 8 groups → maximum stability.
let cosine_buf = unsafe { MappedF32Buffer::new(N_GROUPS) }
.expect("alloc cosine_buf");
cosine_buf.write_from_slice(&[1.0_f32; N_GROUPS]);
// grad_norm = 1.0 for all 8 groups → ε = 1e-7 (within envelope).
let l2_norm_buf = unsafe { MappedF32Buffer::new(N_GROUPS) }
.expect("alloc l2_norm_buf");
l2_norm_buf.write_from_slice(&[1.0_f32; N_GROUPS]);
let scratch_buf = unsafe { MappedF32Buffer::new(SCRATCH_SIZE) }
.expect("alloc scratch_buf");
let cosine_dev = cosine_buf.dev_ptr;
let l2_dev = l2_norm_buf.dev_ptr;
let scratch_dev = scratch_buf.dev_ptr;
let beta1_base_i32 = BETA1_BASE as i32;
let beta2_base_i32 = BETA2_BASE as i32;
let eps_base_i32 = EPS_BASE as i32;
unsafe {
stream
.launch_builder(&kernel)
.arg(&cosine_dev)
.arg(&l2_dev)
.arg(&scratch_dev)
.arg(&beta1_base_i32)
.arg(&beta2_base_i32)
.arg(&eps_base_i32)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (8, 1, 1),
shared_mem_bytes: 0,
})
.expect("launch pearl_4_adam_hparams_update (high stability)");
}
stream.synchronize().expect("sync after pearl_4_adam_hparams_update (high stability)");
let scratch = scratch_buf.read_all();
// Analytical values for stability = 1.0:
let expected_beta1: f32 = 0.85 + 0.10 * 1.0; // = 0.95
let expected_beta2: f32 = 0.99 + 0.0095 * 1.0; // = 0.9995
// ε = clamp(1.0 × 1e-7, 1e-10, 1e-6) = 1e-7
let expected_eps: f32 = 1.0_f32 * 1e-7_f32; // within [1e-10, 1e-6]
for g in 0..N_GROUPS {
let beta1 = scratch[BETA1_BASE + g];
let beta2 = scratch[BETA2_BASE + g];
let eps = scratch[EPS_BASE + g];
println!("group={g} beta1={beta1:.6} beta2={beta6:.7} eps={eps:.2e}",
beta6 = beta2);
// β1 ceiling: 0.95 ± 1 ULP.
let rel_b1 = (beta1 - expected_beta1).abs() / expected_beta1.max(1e-9);
assert!(
rel_b1 < 1e-5,
"group={g}: beta1={beta1:.6} expected={expected_beta1:.6} rel={rel_b1:.2e}"
);
// β2 ceiling: 0.9995 ± 1 ULP.
let rel_b2 = (beta2 - expected_beta2).abs() / expected_beta2.max(1e-9);
assert!(
rel_b2 < 1e-5,
"group={g}: beta2={beta2:.7} expected={expected_beta2:.7} rel={rel_b2:.2e}"
);
// ε must be within structural envelope.
assert!(
eps >= 1e-10_f32 && eps <= 1e-6_f32,
"group={g}: eps={eps:.2e} outside structural envelope [1e-10, 1e-6]"
);
let rel_eps = (eps - expected_eps).abs() / expected_eps.max(1e-15);
assert!(
rel_eps < 1e-4,
"group={g}: eps={eps:.2e} expected={expected_eps:.2e} rel={rel_eps:.2e}"
);
// β1 must be at structural ceiling for maximum stability.
assert!(
(beta1 - 0.95_f32).abs() < 1e-5,
"group={g}: beta1={beta1:.6} should be at ceiling 0.95 for stability=1.0"
);
// β2 must be at structural ceiling for maximum stability.
assert!(
(beta2 - 0.9995_f32).abs() < 1e-5,
"group={g}: beta2={beta2:.7} should be at ceiling 0.9995 for stability=1.0"
);
}
}
/// SP5 Task A4 unit test: `pearl_4_adam_hparams_update` with low stability input.
///
/// Synthetic input: cosine_sim=[0.0; 8] (minimum stability), grad_norm=[0.01; 8].
///
/// Analytical expected:
/// stability = fmaxf(0, fminf(1, 0.0)) = 0.0
/// β1 = 0.85 + 0.10 × 0.0 = 0.85 (low-stability floor)
/// β2 = 0.99 + 0.0095 × 0.0 = 0.99 (low-stability floor)
/// ε = clamp(0.01 × 1e-7, 1e-10, 1e-6) = 1e-9 → clamped to 1e-10
///
/// No CPU reference oracle per `feedback_no_cpu_test_fallbacks.md`.
#[test]
#[ignore = "requires GPU"]
fn pearl_4_adam_hparams_low_stability_yields_short_memory() {
const BETA1_BASE: usize = 0;
const BETA2_BASE: usize = 8;
const EPS_BASE: usize = 16;
const SCRATCH_SIZE: usize = 24;
const N_GROUPS: usize = 8;
let stream = make_test_stream();
let kernel = load_pearl_4_adam_hparams_kernel(&stream);
// cosine_sim = 0.0 for all 8 groups → minimum stability.
let cosine_buf = unsafe { MappedF32Buffer::new(N_GROUPS) }
.expect("alloc cosine_buf");
cosine_buf.write_from_slice(&[0.0_f32; N_GROUPS]);
// grad_norm = 0.01 → ε = 0.01 × 1e-7 = 1e-9 → clamped to 1e-10.
let l2_norm_buf = unsafe { MappedF32Buffer::new(N_GROUPS) }
.expect("alloc l2_norm_buf");
l2_norm_buf.write_from_slice(&[0.01_f32; N_GROUPS]);
let scratch_buf = unsafe { MappedF32Buffer::new(SCRATCH_SIZE) }
.expect("alloc scratch_buf");
let cosine_dev = cosine_buf.dev_ptr;
let l2_dev = l2_norm_buf.dev_ptr;
let scratch_dev = scratch_buf.dev_ptr;
let beta1_base_i32 = BETA1_BASE as i32;
let beta2_base_i32 = BETA2_BASE as i32;
let eps_base_i32 = EPS_BASE as i32;
unsafe {
stream
.launch_builder(&kernel)
.arg(&cosine_dev)
.arg(&l2_dev)
.arg(&scratch_dev)
.arg(&beta1_base_i32)
.arg(&beta2_base_i32)
.arg(&eps_base_i32)
.launch(LaunchConfig {
grid_dim: (1, 1, 1),
block_dim: (8, 1, 1),
shared_mem_bytes: 0,
})
.expect("launch pearl_4_adam_hparams_update (low stability)");
}
stream.synchronize().expect("sync after pearl_4_adam_hparams_update (low stability)");
let scratch = scratch_buf.read_all();
// Analytical values for stability = 0.0:
let expected_beta1: f32 = 0.85; // low-stability floor
let expected_beta2: f32 = 0.99; // low-stability floor
// ε = clamp(0.01 × 1e-7, 1e-10, 1e-6) = clamp(1e-9, 1e-10, 1e-6) = 1e-9
// Actually 1e-9 is within [1e-10, 1e-6], so no clamping needed.
let expected_eps: f32 = (0.01_f32 * 1e-7_f32).clamp(1e-10, 1e-6); // = 1e-9
for g in 0..N_GROUPS {
let beta1 = scratch[BETA1_BASE + g];
let beta2 = scratch[BETA2_BASE + g];
let eps = scratch[EPS_BASE + g];
println!("group={g} beta1={beta1:.6} beta2={beta2:.7} eps={eps:.2e}");
// β1 floor: 0.85 ± 1 ULP.
let rel_b1 = (beta1 - expected_beta1).abs() / expected_beta1.max(1e-9);
assert!(
rel_b1 < 1e-5,
"group={g}: beta1={beta1:.6} expected={expected_beta1:.6} rel={rel_b1:.2e}"
);
// β2 floor: 0.99 ± 1 ULP.
let rel_b2 = (beta2 - expected_beta2).abs() / expected_beta2.max(1e-9);
assert!(
rel_b2 < 1e-5,
"group={g}: beta2={beta2:.7} expected={expected_beta2:.7} rel={rel_b2:.2e}"
);
// ε must be within structural envelope.
assert!(
eps >= 1e-10_f32 && eps <= 1e-6_f32,
"group={g}: eps={eps:.2e} outside structural envelope [1e-10, 1e-6]"
);
let rel_eps = (eps - expected_eps).abs() / expected_eps.max(1e-15);
assert!(
rel_eps < 1e-4,
"group={g}: eps={eps:.2e} expected={expected_eps:.2e} rel={rel_eps:.2e}"
);
// β1 must be at structural floor for zero stability.
assert!(
(beta1 - 0.85_f32).abs() < 1e-5,
"group={g}: beta1={beta1:.6} should be at floor 0.85 for stability=0.0"
);
// β2 must be at structural floor for zero stability.
assert!(
(beta2 - 0.99_f32).abs() < 1e-5,
"group={g}: beta2={beta2:.7} should be at floor 0.99 for stability=0.0"
);
}
}

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

@@ -2,6 +2,8 @@
**Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7.
SP5 Task A4 — Pearl 4 per-group Adam β1/β2/ε GPU producers (Layer A, 2026-05-01): two new CUDA kernels (`grad_cosine_sim_kernel.cu`, `pearl_4_adam_hparams_kernel.cu`) and one `MappedF32Buffer` (`grad_prev_buf_per_group`) land as Layer A additive producers that feed ISV slots [226..250) with per-param-group Adam hyperparameters. `grad_cosine_sim_update`: single-block 8-thread kernel (one thread per SP4 param group: DqnTrunk/Value/Branches/IQN/IqlHigh/IqlLow/Attn/Curiosity), reads `grad_buf[total_params]` + `grad_prev_buf[total_params]` + `group_param_offsets[9]` (prefix sums in element units); computes `dot = Σ gc·gp`, `norm_curr_sq = Σ gc²`, `norm_prev_sq = Σ gp²` per group; writes `cos_sim[8]→scratch[131..139)` and `l2_norm[8]→scratch[139..147)`; writebacks `grad_curr → grad_prev_buf` for the next step's comparison (same thread — no race). `pearl_4_adam_hparams_update`: single-block 8-thread kernel, reads `cos_sim[8]` + `l2_norm[8]` from scratch; clips stability to `[0, 1]`; computes `β1 = 0.85 + 0.10×stability ∈ [0.85, 0.95]`, `β2 = 0.99 + 0.0095×stability ∈ [0.99, 0.9995]`, `ε = clamp(grad_norm × 1e-7, 1e-10, 1e-6)`; writes `β1[8]→scratch[147..155)`, `β2[8]→scratch[155..163)`, `ε[8]→scratch[163..171)`. Structural envelopes (Invariant 1 anchors per spec): β1∈[0.85, 0.95], β2∈[0.99, 0.9995], ε∈[1e-10, 1e-6]. `launch_sp5_pearl_4_adam_hparams` fires 24 `launch_apply_pearls` calls: β1[8] → ISV[226..234), β2[8] → ISV[234..242), ε[8] → ISV[242..250). Uses `ALPHA_META_PEARL_4=5e-4` (half SP4 default 1e-3) per theoretical-caveat mitigation — slow β change rate reduces instability risk from adaptive β breaking Adam's constant-β convergence proof. Pearl A sentinel 0 at fold boundary fires first-observation replacement (cosine_sim=0 → low-stability envelope endpoints). Aux groups 3-7 (IQN/IqlHigh/IqlLow/Attn/Curiosity) have non-contiguous grad allocations; zero-width sentinel offsets `[tp, tp, tp, tp, tp]` in `group_param_offsets_buf` prevent the kernel inner loop from executing for those groups, yielding cosine_sim=0 → β1/β2 start at low-stability floors. `apply_pearls_kernel.cu` signature gained `float alpha_meta` parameter (migration of all 18 call sites atomic per `feedback_no_partial_refactor.md`: 1 in `gpu_experience_collector.rs`, 17 in `gpu_dqn_trainer.rs`). Buffer growth: `producer_step_scratch_buf` 131→171 slots (SP5_SCRATCH_TOTAL); 5 new scratch constants `SCRATCH_PEARL_4_{COSINE=131, L2=139, BETA1=147, BETA2=155, EPS=163}`. StateResetRegistry gains 4 new FoldReset entries: `sp5_adam_beta1` (ISV[226..234)), `sp5_adam_beta2` (ISV[234..242)), `sp5_adam_eps` (ISV[242..250)), `sp5_grad_prev_buf` (grad_prev_buf_per_group — zero at fold boundary so cosine_sim=0 fires Pearl A bootstrap). Wire-up in training_loop.rs: `launch_sp5_pearl_4_adam_hparams()` called immediately after `launch_sp5_pearl_2_budget()` with `tracing::warn!` on error. Two GPU-only `#[ignore = "requires GPU"]` unit tests: (7) high-stability (cosine_sim=1.0 → β1=0.95, β2=0.9995, ε within envelope); (8) low-stability (cosine_sim=0.0 → β1=0.85, β2=0.99, ε within envelope). No CPU compute, no HtoD/HtoH, no atomicAdd, no stubs, no consumer migration. Touched: `cuda_pipeline/grad_cosine_sim_kernel.cu` (new), `cuda_pipeline/pearl_4_adam_hparams_kernel.cu` (new), `cuda_pipeline/apply_pearls_kernel.cu` (+alpha_meta param), `cuda_pipeline/sp4_wiener_ema.rs` (+alpha_meta arg), `cuda_pipeline/gpu_experience_collector.rs` (+1 call site ALPHA_META arg), `build.rs` (+2 cubin entries), `cuda_pipeline/gpu_dqn_trainer.rs` (SP5_SCRATCH_TOTAL 131→171 + 5 new constants + 2 static cubins + 4 struct fields + constructor allocs + struct initializer + 17 call-site ALPHA_META args + launch method), `trainers/dqn/state_reset_registry.rs` (+4 FoldReset entries), `trainers/dqn/trainer/training_loop.rs` (+5 LOC after Pearl 2 launch), `tests/sp5_producer_unit_tests.rs` (+2 GPU-only tests + cubin constant + loader helper). Hard rules: feedback_no_cpu_compute_strict, feedback_no_atomicadd, feedback_no_htod_htoh_only_mapped_pinned, feedback_no_stubs, feedback_no_partial_refactor (apply_pearls 18 call sites all migrated).
SP5 Task A3 fix-up — Pearl 2 budget tests now assert sum-to-1 normalization invariant (2026-05-01): code-quality review caught that the two A3 unit tests verified each output (c51, iqn, cql, ens) against its analytical expected value within 1% relative tolerance, but did NOT assert the structural invariant `c51 + iqn + cql + ens ≈ 1.0` that the kernel maintains by construction (`ens = max(0, 1 - iqn - c51 - cql)`). A coefficient typo such as `BASE_IQN = 0.111` instead of `0.11` would produce individually-plausible per-component values that all still pass the relative checks while quietly summing to 0.97 or 1.04. Fix adds `assert!((c51+iqn+cql+ens - 1.0).abs() < 1e-4)` inside both per-branch loops. Touched: `crates/ml/tests/sp5_producer_unit_tests.rs` (+12 LOC). cargo test --no-run clean.
SP5 Task A3 — Pearl 2 per-branch loss budget GPU producer (Layer A, 2026-05-01): one new CUDA kernel (`pearl_2_budget_kernel.cu`) lands as a Layer A additive producer that feeds ISV slots [190..210) with per-branch loss budget and flatness signals. No loss-budget consumer migration in this commit — Layer B wires consumers. `pearl_2_budget_update`: single-block 4-thread kernel (one thread per branch), reads Q_VAR_PER_BRANCH[b] (Pearl 1's q_branch_stats output, ISV[222..226)) + NOISY_SIGMA[b] (Pearl 3 output, ISV[210..214)); computes `flatness[b] = clamp(var_q[b] / (σ[b]² + EPS_DIV), 0, 1)` where EPS_DIV=1e-8 is an Invariant 1 anchor (numerical stability); derives `budget_iqn = BASE_IQN + (1 - flatness) * (1 - BASE_IQN)` (IQN dominates when Q is flat relative to noise), `budget_c51 = flatness * (1 - BASE_IQN) * BASE_C51_SHARE` (C51 takes proportional remainder when Q is well-separated), `budget_cql = 0` (gated by existing regime_stability allocator), `budget_ens = max(0, 1 - iqn - c51 - cql)`; writes 20 floats to scratch[111..131) (SCRATCH_PEARL_2_C51=111, SCRATCH_PEARL_2_IQN=115, SCRATCH_PEARL_2_CQL=119, SCRATCH_PEARL_2_ENS=123, SCRATCH_PEARL_2_FLATNESS=127). BASE_IQN=0.11 and BASE_C51_SHARE=0.84/(1-BASE_IQN) are Invariant 1 anchors (preserve SP4-baseline budget under non-flat regime). `launch_sp5_pearl_2_budget` fires 20 `launch_apply_pearls` calls: 5 output groups × 4 branches each, mapping to ISV[BUDGET_C51_BASE=190..194), ISV[BUDGET_IQN_BASE=194..198), ISV[BUDGET_CQL_BASE=198..202), ISV[BUDGET_ENS_BASE=202..206), ISV[FLATNESS_BASE=206..210). Wiener offset formula: `base_wiener_offset + (isv_slot - SP5_SLOT_BASE) * 3` where `base_wiener_offset = SP4_PRODUCER_COUNT * 3 = 213`. Dependencies: Pearl 2 must chain AFTER Pearl 1 (reads Q_VAR_PER_BRANCH) AND AFTER Pearl 3 (reads NOISY_SIGMA). Buffer growth: `producer_step_scratch_buf` 111→131 slots (SP5_SCRATCH_TOTAL updated 111→131; 5 new module-level constants SCRATCH_PEARL_2_{C51,IQN,CQL,ENS,FLATNESS}). wiener_state_buf already at 543 (A1 sized for entire SP5 block — no further growth). StateResetRegistry gains 5 new FoldReset entries: sp5_budget_c51 (ISV[190..194)), sp5_budget_iqn (ISV[194..198)), sp5_budget_cql (ISV[198..202)), sp5_budget_ens (ISV[202..206)), sp5_flatness (ISV[206..210)) — Pearl A sentinel 0 at fold boundary triggers first-observation replacement on new fold's first launch. Wire-up in training_loop.rs: `launch_sp5_pearl_2_budget()` called immediately after `launch_sp5_pearl_3_sigma()` with `tracing::warn!(error = %e, ...)` on error. Two GPU-only `#[ignore = "requires GPU"]` unit tests: (5) flat-Q regime (σ=1.0, var_q=2.0 → flatness=1.0 → iqn=0.11, c51≈0.84); (6) sharp-Q regime (σ=2.0, var_q=0.04 → flatness≈0.01 → iqn>0.9, c51≈0.0084). No CPU compute, no HtoD/HtoH, no atomicAdd, no stubs, no consumer migration. Touched: `cuda_pipeline/pearl_2_budget_kernel.cu` (new), `build.rs` (+1 cubin entry), `cuda_pipeline/gpu_dqn_trainer.rs` (static cubin + SP5_SCRATCH_TOTAL 111→131 + 5 new scratch constants + docstring update + struct field + cubin loader + struct initializer + launch method), `trainers/dqn/state_reset_registry.rs` (+5 FoldReset entries), `trainers/dqn/trainer/training_loop.rs` (+5 LOC after Pearl 3 launch), `tests/sp5_producer_unit_tests.rs` (+2 GPU-only tests + cubin loader helper). Hard rules: feedback_no_cpu_compute_strict, feedback_no_atomicadd, feedback_no_stubs, feedback_no_partial_refactor (all consumers of existing budget allocator untouched).