diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 3c9a1b6c2..ffe264541 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -412,6 +412,15 @@ fn main() { // Writes 16 floats to producer_step_scratch_buf[87..103). // apply_pearls_ad_kernel then smooths via Pearls A+D into ISV[174..190). "pearl_1_atom_kernel.cu", + // SP5 Task A2 (2026-05-01): Pearl 3 per-branch NoisyNet sigma controller. + // Single-block 4-thread kernel reads ATOM_V_HALF (Pearl 1 output, + // ISV[178..182)) + BRANCH_ENTROPY (q_branch_stats, ISV[218..222)) + + // current SIGMA_FRACTION (ISV[214..218)); computes entropy-deficit + // controller (target_entropy = log(n_actions[b]) * 0.7, SF_LR=0.005) + // and new_sigma = new_sf * max(v_half, 1.0). Writes (sigma[4], SF[4]) + // to producer_step_scratch_buf[103..111). + // apply_pearls_ad_kernel then smooths via Pearls A+D into ISV[210..218). + "pearl_3_sigma_kernel.cu", // SP4 GPU-only Pearls A+D applicator (2026-05-01). Single-thread // device-side loop applies Pearls A+D to N consecutive ISV slots // (each with its own Wiener-state triple) on the producer's diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 4e3ba3ce9..1b10539dc 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -292,6 +292,18 @@ static SP5_Q_BRANCH_STATS_CUBIN: &[u8] = static SP5_PEARL_1_ATOM_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/pearl_1_atom_kernel.cubin")); +/// SP5 Task A2 (2026-05-01): Pearl 3 per-branch NoisyNet sigma controller cubin. +/// Single-block 4-thread kernel. Reads ATOM_V_HALF[b] (Pearl 1, ISV[178..182)) +/// + BRANCH_ENTROPY[b] (q_branch_stats, ISV[218..222)) + SIGMA_FRACTION[b] +/// (ISV[214..218), Pearl A sentinel 0 → bootstrap 0.1); applies entropy-deficit +/// controller (target=log(n_actions[b])×0.7, LR=0.005, envelope [0.001,0.5]); +/// writes (sigma[4], SF[4]) to producer_step_scratch_buf[103..111). +/// Chained apply_pearls_ad_kernel (8 single-slot launches) smooths via +/// Pearls A+D into ISV[210..218). +/// Layer A only — NoisyLinear consumer migration deferred to Layer B. +static SP5_PEARL_3_SIGMA_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/pearl_3_sigma_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 @@ -875,12 +887,23 @@ 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: producer_step_scratch_buf total slot count after growth. +/// SP5 Task A1+A2: 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) -/// Combined: 71 + 16 + 16 = 103 scratch slots [0..103). -pub const SP5_SCRATCH_TOTAL: usize = 103; +/// Task A2 adds: +/// 4 floats for pearl_3_sigma output (SCRATCH_PEARL_3_SIGMA=103, sigma × 4 branches) +/// 4 floats for pearl_3_sf output (SCRATCH_PEARL_3_SF=107, sigma_fraction × 4 branches) +/// Combined: 71 + 16 + 16 + 4 + 4 = 111 scratch slots [0..111). +pub const SP5_SCRATCH_TOTAL: usize = 111; + +/// SP5 Task A2: scratch index base for pearl_3_sigma sigma[4] output block. +/// Slots [103..107): per-branch NoisyNet sigma. +pub const SCRATCH_PEARL_3_SIGMA: usize = 103; + +/// SP5 Task A2: scratch index base for pearl_3_sigma sigma_fraction[4] output block. +/// Slots [107..111): per-branch sigma fraction (entropy-deficit controller state). +pub const SCRATCH_PEARL_3_SF: usize = 107; /// 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. @@ -4004,6 +4027,18 @@ pub struct GpuDqnTrainer { /// smooths via Pearls A+D into ISV[174..190) + ISV[218..226). /// Loaded from `pearl_1_atom_kernel.cubin`. pearl_1_atom_kernel: CudaFunction, + /// SP5 Task A2 (2026-05-01): Pearl 3 per-branch NoisyNet sigma controller kernel. + /// Single-block 4-thread kernel (one thread per branch). Reads ATOM_V_HALF + /// (Pearl 1, ISV[178..182)) + BRANCH_ENTROPY (ISV[218..222)) + SIGMA_FRACTION + /// (ISV[214..218), Pearl A sentinel 0 → bootstrap 0.1); entropy-deficit + /// controller (target = log(n_actions[b]) × 0.7, LR=0.005, envelope [0.001,0.5]); + /// writes (sigma[4], SF[4]) to `producer_step_scratch_buf[103..111)`. + /// Chained `apply_pearls_ad_kernel` (8 single-slot launches) then + /// smooths via Pearls A+D into ISV[210..218). + /// Layer A only — NoisyLinear consumer migration deferred to Layer B. + /// No atomicAdd (feedback_no_atomicadd), no CPU compute (feedback_no_cpu_compute_strict). + /// Loaded from `pearl_3_sigma_kernel.cubin`. + pearl_3_sigma_kernel: CudaFunction, /// 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. @@ -10475,6 +10510,104 @@ impl GpuDqnTrainer { Ok(()) } + /// SP5 Task A2 (2026-05-01): Pearl 3 — per-branch NoisyNet sigma from + /// per-branch v_half + entropy. + /// + /// Chained AFTER `launch_sp5_pearl_1_atom` (reads ATOM_V_HALF written by + /// Pearl 1). Two-step sequence: + /// 1. `pearl_3_sigma_update` (1×4 kernel) → scratch[103..111). + /// 2. `apply_pearls_ad_kernel` × 8 (one per output slot): + /// sigma[4] → ISV[NOISY_SIGMA_BASE=210..214) + /// SF[4] → ISV[SIGMA_FRACTION_BASE=214..218) + /// + /// Wiener offsets: (SP4_PRODUCER_COUNT=71 + (isv_slot − SP5_SLOT_BASE=174)) × 3. + /// Layer A only — NoisyLinear consumer migration deferred to Layer B. + pub(crate) fn launch_sp5_pearl_3_sigma(&self) -> Result<(), MLError> { + use crate::cuda_pipeline::sp4_wiener_ema::launch_apply_pearls; + use crate::cuda_pipeline::sp5_isv_slots::{ + SP5_SLOT_BASE, + ATOM_V_HALF_BASE, BRANCH_ENTROPY_BASE, + NOISY_SIGMA_BASE, SIGMA_FRACTION_BASE, + }; + + debug_assert!(self.isv_signals_dev_ptr != 0, + "launch_sp5_pearl_3_sigma: isv_signals_dev_ptr must be allocated by constructor"); + + let isv_dev = self.isv_signals_dev_ptr; + let scratch_dev = self.producer_step_scratch_buf.dev_ptr; + let wiener_dev = self.wiener_state_buf.dev_ptr; + let action_counts_dev = self.action_counts_buf.dev_ptr; + + // Step 1: pearl_3_sigma_update → scratch[103..107) sigma + scratch[107..111) SF. + // Kernel signature: (isv_signals, v_half_isv_base, branch_entropy_isv_base, + // sigma_fraction_isv_base, action_counts, + // scratch_out, sigma_idx_base, sigma_fraction_idx_base) + let v_half_base_i32 = ATOM_V_HALF_BASE as i32; // 178 + let branch_entropy_base_i32 = BRANCH_ENTROPY_BASE as i32; // 218 + let sigma_fraction_base_i32 = SIGMA_FRACTION_BASE as i32; // 214 + let sigma_idx_i32 = SCRATCH_PEARL_3_SIGMA as i32; // 103 + let sf_idx_i32 = SCRATCH_PEARL_3_SF as i32; // 107 + + unsafe { + self.stream + .launch_builder(&self.pearl_3_sigma_kernel) + .arg(&isv_dev) + .arg(&v_half_base_i32) + .arg(&branch_entropy_base_i32) + .arg(&sigma_fraction_base_i32) + .arg(&action_counts_dev) + .arg(&scratch_dev) + .arg(&sigma_idx_i32) + .arg(&sf_idx_i32) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (4, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!("pearl_3_sigma_update: {e}")))?; + } + + // Step 2: apply_pearls_ad_kernel × 8 — one single-slot launch per ISV output. + // Wiener offset: (SP4_PRODUCER_COUNT + (isv_slot - SP5_SLOT_BASE)) * 3. + // + // sigma[4] → scratch [103..107), ISV [210..214) (NOISY_SIGMA_BASE) + // SF[4] → scratch [107..111), ISV [214..218) (SIGMA_FRACTION_BASE) + let base_wiener_offset = SP4_PRODUCER_COUNT as i32 * 3; // 213 + + for b in 0..4_usize { + // sigma → ISV NOISY_SIGMA_BASE + b + let scratch_sigma = (SCRATCH_PEARL_3_SIGMA + b) as i32; + let isv_sigma = (NOISY_SIGMA_BASE + b) as i32; + let wiener_sigma = base_wiener_offset + (isv_sigma - SP5_SLOT_BASE as i32) * 3; + unsafe { + launch_apply_pearls( + &self.stream, + &self.apply_pearls_ad_kernel, + scratch_dev, scratch_sigma, + isv_dev, isv_sigma, + wiener_dev, wiener_sigma, + 1, + )?; + } + // SF → ISV SIGMA_FRACTION_BASE + b + let scratch_sf = (SCRATCH_PEARL_3_SF + b) as i32; + let isv_sf = (SIGMA_FRACTION_BASE + b) as i32; + let wiener_sf = base_wiener_offset + (isv_sf - SP5_SLOT_BASE as i32) * 3; + unsafe { + launch_apply_pearls( + &self.stream, + &self.apply_pearls_ad_kernel, + scratch_dev, scratch_sf, + isv_dev, isv_sf, + wiener_dev, wiener_sf, + 1, + )?; + } + } + + 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]. /// @@ -12474,6 +12607,16 @@ impl GpuDqnTrainer { .map_err(|e| MLError::ModelError(format!("pearl_1_atom_update load: {e}")))? }; + // SP5 Task A2 (2026-05-01): load pearl_3_sigma NoisyNet sigma controller kernel. + // Single-block 4-thread kernel; reads ATOM_V_HALF + BRANCH_ENTROPY + SIGMA_FRACTION + // from ISV; writes (sigma[4], SF[4]) to scratch[103..111). + let pearl_3_sigma_kernel = { + let module = stream.context().load_cubin(SP5_PEARL_3_SIGMA_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("sp5 pearl_3_sigma cubin load: {e}")))?; + module.load_function("pearl_3_sigma_update") + .map_err(|e| MLError::ModelError(format!("pearl_3_sigma_update load: {e}")))? + }; + // SP5 Task A1 (2026-05-01): allocate atoms_clip_count_buf [4 i32]. // Zero-initialized; populated by atoms_update_kernel in Layer B. // Until Layer B wires, clip_rate=0 → headroom holds at 6.0 bootstrap. @@ -15594,6 +15737,7 @@ impl GpuDqnTrainer { apply_pearls_ad_kernel, q_branch_stats_kernel, pearl_1_atom_kernel, + pearl_3_sigma_kernel, atoms_clip_count_buf, action_counts_buf, action_offsets_buf, diff --git a/crates/ml/src/cuda_pipeline/pearl_3_sigma_kernel.cu b/crates/ml/src/cuda_pipeline/pearl_3_sigma_kernel.cu new file mode 100644 index 000000000..4eeb8aecf --- /dev/null +++ b/crates/ml/src/cuda_pipeline/pearl_3_sigma_kernel.cu @@ -0,0 +1,61 @@ +// crates/ml/src/cuda_pipeline/pearl_3_sigma_kernel.cu +// +// SP5 Task A2: Pearl 3 — per-branch NoisyNet sigma scales with per-branch Q magnitude. +// Reads ATOM_V_HALF[b] (Pearl 1 from Task A1) + BRANCH_ENTROPY[b] (q_branch_stats from +// Task A1) + current SIGMA_FRACTION[b] (Pearl A sentinel 0 -> bootstrap to 0.1). +// Writes per-branch (NOISY_SIGMA, SIGMA_FRACTION) to scratch_out. +// apply_pearls_ad_kernel then smooths these via Pearls A+D into ISV. +// +// Controller: SIGMA_FRACTION[b] adapts via entropy-deficit: +// target_entropy = log(action_counts[b]) * 0.7 (Invariant 1 anchor: 70% of max) +// SF[b] += SF_LR * (target_entropy - actual_entropy[b]) +// New sigma = SF * max(v_half, EPS_CLAMP_FLOOR). +// +// 4-thread single-block (one thread per branch). +// No atomicAdd (feedback_no_atomicadd). __threadfence_system() after writes. +// +// SF_LR=0.005 (slow) and target_entropy_factor=0.7 are Invariant 1 structural anchors. +// Envelope [0.001, 0.5] prevents sigma runaway in either direction. + +#include + +extern "C" __global__ void pearl_3_sigma_update( + const float* __restrict__ isv_signals, + int v_half_isv_base, // ATOM_V_HALF_BASE = 178 + int branch_entropy_isv_base, // BRANCH_ENTROPY_BASE = 218 + int sigma_fraction_isv_base, // SIGMA_FRACTION_BASE = 214 + const int* __restrict__ action_counts, // [4] = {4, 3, 3, 3} + float* __restrict__ scratch_out, + int sigma_idx_base, // scratch_out base for sigma[4] + int sigma_fraction_idx_base // scratch_out base for SF[4] +) { + int b = threadIdx.x; + if (blockIdx.x != 0 || b >= 4) return; + + // EPS_CLAMP_FLOOR = 1.0 is an Invariant 1 anchor (structural lower bound on + // v_half when Pearl 1 has not yet produced a real observation). + const float EPS_CLAMP_FLOOR = 1.0f; + + float v_half = fmaxf(isv_signals[v_half_isv_base + b], EPS_CLAMP_FLOOR); + float entropy = isv_signals[branch_entropy_isv_base + b]; + float sigma_fraction = isv_signals[sigma_fraction_isv_base + b]; + // Pearl A sentinel: ISV slot initialised to 0 at fold boundary. + // First observation replaces directly with the bootstrap value. + if (sigma_fraction <= 0.0f) sigma_fraction = 0.1f; + + int n_actions = action_counts[b]; + // target_entropy_factor=0.7 is an Invariant 1 structural anchor (70% of max entropy). + float target_entropy = logf((float)n_actions) * 0.7f; + + // SF_LR=0.005 is an Invariant 1 structural anchor (slow controller). + const float SF_LR = 0.005f; + float new_sf = sigma_fraction + SF_LR * (target_entropy - entropy); + // Structural envelope [0.001, 0.5] prevents sigma runaway in either direction. + new_sf = fmaxf(0.001f, fminf(0.5f, new_sf)); + + float new_sigma = new_sf * v_half; + + scratch_out[sigma_idx_base + b] = new_sigma; + scratch_out[sigma_fraction_idx_base + b] = new_sf; + __threadfence_system(); +} diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index a675be450..ee6052044 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -572,6 +572,16 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "ISV[ATOM_CLIP_RATE_BASE=186..190) — per-branch atom clip rate (atoms_clip_count / (batch_size*13)). SP5 Pearl A sentinel 0 at fold boundary (Task A1).", }, + RegistryEntry { + name: "sp5_noisy_sigma", + category: ResetCategory::FoldReset, + description: "ISV[NOISY_SIGMA_BASE=210..214) — per-branch NoisyNet sigma (= sigma_fraction x v_half from Pearl 1). SP5 Pearl A sentinel 0 at fold boundary so the new fold's first pearl_3_sigma_update launch fires the first-observation replacement (Task A2).", + }, + RegistryEntry { + name: "sp5_sigma_fraction", + category: ResetCategory::FoldReset, + description: "ISV[SIGMA_FRACTION_BASE=214..218) — per-branch sigma fraction (entropy-deficit controller, init 0.1 via Pearl A bootstrap, envelope [0.001, 0.5]). SP5 Pearl A sentinel 0 at fold boundary so the controller re-bootstraps to 0.1 at the new fold's first step (Task A2).", + }, RegistryEntry { name: "sp5_branch_entropy", category: ResetCategory::FoldReset, diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 566cc33c9..279f66b55 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -3589,6 +3589,13 @@ impl DQNTrainer { if let Err(e) = fused.trainer().launch_sp5_pearl_1_atom() { tracing::warn!("SP5 A1 launch_sp5_pearl_1_atom failed: {e}"); } + // SP5 Task A2: pearl_3_sigma → apply_pearls_ad ×8 + // Reads ATOM_V_HALF (Pearl 1 output) + BRANCH_ENTROPY (q_branch_stats). + // Must run AFTER launch_sp5_pearl_1_atom (Pearl 3 reads Pearl 1's v_half). + // Producer-only (Layer A) — no consumer reads ISV[210..218) in this commit. + if let Err(e) = fused.trainer().launch_sp5_pearl_3_sigma() { + tracing::warn!("SP5 A2 launch_sp5_pearl_3_sigma failed: {e}"); + } } } diff --git a/crates/ml/tests/sp5_producer_unit_tests.rs b/crates/ml/tests/sp5_producer_unit_tests.rs index 6d2a64630..c96eac5ec 100644 --- a/crates/ml/tests/sp5_producer_unit_tests.rs +++ b/crates/ml/tests/sp5_producer_unit_tests.rs @@ -27,8 +27,21 @@ const SP5_Q_BRANCH_STATS_CUBIN: &[u8] = const SP5_PEARL_1_ATOM_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/pearl_1_atom_kernel.cubin")); +const SP5_PEARL_3_SIGMA_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/pearl_3_sigma_kernel.cubin")); + // ── Helpers ─────────────────────────────────────────────────────────────────── +fn load_pearl_3_sigma_kernel(stream: &Arc) -> CudaFunction { + let module = stream + .context() + .load_cubin(SP5_PEARL_3_SIGMA_CUBIN.to_vec()) + .expect("load pearl_3_sigma_kernel cubin"); + module + .load_function("pearl_3_sigma_update") + .expect("load pearl_3_sigma_update function") +} + fn make_test_stream() -> Arc { let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?"); ctx.default_stream() @@ -327,3 +340,226 @@ fn sp5_pearl_1_atom_headroom_bootstrap_and_controller_step() { ); } } + +// ── Test 3: pearl_3_sigma_update bootstrap at target entropy ────────────────── + +/// SP5 Task A2 unit test: `pearl_3_sigma_update` bootstraps SF=0.1 when ISV +/// sigma_fraction=0 (Pearl A sentinel), computes new_sigma = SF * v_half. +/// +/// Synthetic input: +/// - ISV[ATOM_V_HALF_BASE+b] = 2.0 (all branches) +/// - ISV[BRANCH_ENTROPY_BASE+b] = log(n_actions[b]) * 0.7 (target entropy exactly) +/// - ISV[SIGMA_FRACTION_BASE+b] = 0.0 (Pearl A sentinel → bootstrap to 0.1) +/// - action_counts = {4, 3, 3, 3} +/// +/// Analytical ground truth (target == actual → no SF update step, only bootstrap): +/// - sigma_fraction bootstraps to 0.1 (sentinel 0 detection) +/// - new_sf = 0.1 + 0.005 * (target - actual) = 0.1 + 0.005 * 0 = 0.1 (exact) +/// - new_sigma = 0.1 * max(2.0, 1.0) = 0.1 * 2.0 = 0.2 for each branch +/// +/// This verifies Pearl A sentinel detection, bootstrap value, and SF * v_half formula. +/// No CPU reference oracle per `feedback_no_cpu_test_fallbacks.md`. +#[test] +#[ignore = "requires GPU"] +fn pearl_3_sigma_bootstrap_at_target_entropy() { + const SCRATCH_SIGMA: usize = 103; + const SCRATCH_SF: usize = 107; + + // Need enough ISV slots: ATOM_V_HALF_BASE=178, SIGMA_FRACTION_BASE=214, + // BRANCH_ENTROPY_BASE=218. Max = 218+4 = 222. + const ISV_SIZE: usize = 230; + + let stream = make_test_stream(); + let kernel = load_pearl_3_sigma_kernel(&stream); + + 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); + + // Build ISV: v_half=2.0 at [178..182), SF=0.0 at [214..218) (sentinel), + // entropy=log(n)*0.7 at [218..222). + let isv_buf = unsafe { MappedF32Buffer::new(ISV_SIZE) } + .expect("alloc isv_buf"); + let mut isv_data = vec![0.0_f32; ISV_SIZE]; + for b in 0..4_usize { + isv_data[178 + b] = 2.0_f32; // ATOM_V_HALF_BASE = 178 + // SIGMA_FRACTION_BASE = 214: leave at 0.0 (sentinel) + // BRANCH_ENTROPY_BASE = 218: set to target_entropy = log(n) * 0.7 + let n = action_counts[b] as f32; + isv_data[218 + b] = n.ln() * 0.7_f32; + } + isv_buf.write_from_slice(&isv_data); + + let scratch_buf = unsafe { MappedF32Buffer::new(SCRATCH_SF + 4) } + .expect("alloc scratch_buf"); + + let isv_dev = isv_buf.dev_ptr; + let v_half_base_i32 = 178_i32; // ATOM_V_HALF_BASE + let branch_entropy_base_i32 = 218_i32; // BRANCH_ENTROPY_BASE + let sigma_fraction_base_i32 = 214_i32; // SIGMA_FRACTION_BASE + let counts_dev = action_counts_buf.dev_ptr; + let scratch_dev = scratch_buf.dev_ptr; + let sigma_idx_i32 = SCRATCH_SIGMA as i32; + let sf_idx_i32 = SCRATCH_SF as i32; + + unsafe { + stream + .launch_builder(&kernel) + .arg(&isv_dev) + .arg(&v_half_base_i32) + .arg(&branch_entropy_base_i32) + .arg(&sigma_fraction_base_i32) + .arg(&counts_dev) + .arg(&scratch_dev) + .arg(&sigma_idx_i32) + .arg(&sf_idx_i32) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (4, 1, 1), + shared_mem_bytes: 0, + }) + .expect("launch pearl_3_sigma_update (bootstrap test)"); + } + stream.synchronize().expect("sync after pearl_3_sigma_update (bootstrap test)"); + + let scratch = scratch_buf.read_all(); + + for b in 0..4_usize { + let sigma = scratch[SCRATCH_SIGMA + b]; + let sf = scratch[SCRATCH_SF + b]; + + println!("branch={b} sigma={sigma:.6} sf={sf:.6}"); + + // SF bootstraps to 0.1 (sentinel 0 → 0.1), no update because target=actual. + let rel_err_sf = (sf - 0.1_f32).abs() / 0.1_f32; + assert!( + rel_err_sf < 0.01, + "branch={b}: sf={sf:.6} expected=0.1 rel_err={rel_err_sf:.5}" + ); + + // sigma = sf * v_half = 0.1 * 2.0 = 0.2 + let expected_sigma = 0.1_f32 * 2.0_f32; + let rel_err_sigma = (sigma - expected_sigma).abs() / expected_sigma; + assert!( + rel_err_sigma < 0.01, + "branch={b}: sigma={sigma:.6} expected={expected_sigma:.6} rel_err={rel_err_sigma:.5}" + ); + } +} + +// ── Test 4: pearl_3_sigma_update fraction increases when entropy below target ── + +/// SP5 Task A2 unit test: when actual entropy is BELOW target, the entropy-deficit +/// controller increases SF. +/// +/// Synthetic input for branch 0 (Direction, n_actions=4): +/// - ISV[ATOM_V_HALF_BASE+0] = 2.0 +/// - ISV[BRANCH_ENTROPY_BASE+0] = 0.0 (extreme case: zero entropy) +/// - ISV[SIGMA_FRACTION_BASE+0] = 0.1 (already bootstrapped) +/// - target_entropy = log(4) * 0.7 = 1.3862944 * 0.7 = 0.97040609 +/// +/// Analytical ground truth for branch 0: +/// new_sf = 0.1 + 0.005 * (0.97040609 - 0.0) = 0.1 + 0.00485203 = 0.10485203 +/// new_sigma = 0.10485203 * 2.0 = 0.20970406 +/// +/// Same inputs for branches 1/2/3 (n=3, entropy=0, sf=0.1): +/// target_entropy = log(3) * 0.7 = 1.0986123 * 0.7 = 0.76902860 +/// new_sf = 0.1 + 0.005 * 0.76902860 = 0.10384514 +/// new_sigma = 0.20769028 +/// +/// Checks SF strictly increases when entropy is below target. 1% rel-err tolerance. +#[test] +#[ignore = "requires GPU"] +fn pearl_3_sigma_fraction_increases_when_entropy_below_target() { + const SCRATCH_SIGMA: usize = 103; + const SCRATCH_SF: usize = 107; + const ISV_SIZE: usize = 230; + + let stream = make_test_stream(); + let kernel = load_pearl_3_sigma_kernel(&stream); + + 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); + + // ISV: v_half=2.0, SF=0.1 (already bootstrapped), entropy=0 (below target). + let isv_buf = unsafe { MappedF32Buffer::new(ISV_SIZE) } + .expect("alloc isv_buf"); + let mut isv_data = vec![0.0_f32; ISV_SIZE]; + for b in 0..4_usize { + isv_data[178 + b] = 2.0_f32; // ATOM_V_HALF_BASE = 178 + isv_data[214 + b] = 0.1_f32; // SIGMA_FRACTION_BASE = 214 (already bootstrapped) + // BRANCH_ENTROPY_BASE = 218: leave at 0.0 (zero entropy) + } + isv_buf.write_from_slice(&isv_data); + + let scratch_buf = unsafe { MappedF32Buffer::new(SCRATCH_SF + 4) } + .expect("alloc scratch_buf"); + + let isv_dev = isv_buf.dev_ptr; + let v_half_base_i32 = 178_i32; + let branch_entropy_base_i32 = 218_i32; + let sigma_fraction_base_i32 = 214_i32; + let counts_dev = action_counts_buf.dev_ptr; + let scratch_dev = scratch_buf.dev_ptr; + let sigma_idx_i32 = SCRATCH_SIGMA as i32; + let sf_idx_i32 = SCRATCH_SF as i32; + + unsafe { + stream + .launch_builder(&kernel) + .arg(&isv_dev) + .arg(&v_half_base_i32) + .arg(&branch_entropy_base_i32) + .arg(&sigma_fraction_base_i32) + .arg(&counts_dev) + .arg(&scratch_dev) + .arg(&sigma_idx_i32) + .arg(&sf_idx_i32) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (4, 1, 1), + shared_mem_bytes: 0, + }) + .expect("launch pearl_3_sigma_update (entropy-below-target test)"); + } + stream.synchronize().expect("sync after pearl_3_sigma_update (entropy-below-target test)"); + + let scratch = scratch_buf.read_all(); + + // Branch 0: n=4, target=log(4)*0.7 + { + let b = 0_usize; + let sigma = scratch[SCRATCH_SIGMA + b]; + let sf = scratch[SCRATCH_SF + b]; + let target_entropy = (4.0_f32).ln() * 0.7_f32; + let expected_sf = 0.1_f32 + 0.005_f32 * target_entropy; + let expected_sigma = expected_sf * 2.0_f32; + println!("branch={b} n=4 sf={sf:.7} expected_sf={expected_sf:.7} \ + sigma={sigma:.7} expected_sigma={expected_sigma:.7}"); + let rel_sf = (sf - expected_sf).abs() / expected_sf; + assert!(rel_sf < 0.01, "branch={b}: sf={sf:.7} expected={expected_sf:.7} rel={rel_sf:.5}"); + let rel_sig = (sigma - expected_sigma).abs() / expected_sigma; + assert!(rel_sig < 0.01, "branch={b}: sigma={sigma:.7} expected={expected_sigma:.7} rel={rel_sig:.5}"); + // SF strictly increased from 0.1 + assert!(sf > 0.1_f32, "branch={b}: SF should have increased from 0.1, got {sf}"); + } + + // Branches 1-3: n=3, target=log(3)*0.7 + for b in 1..4_usize { + let sigma = scratch[SCRATCH_SIGMA + b]; + let sf = scratch[SCRATCH_SF + b]; + let target_entropy = (3.0_f32).ln() * 0.7_f32; + let expected_sf = 0.1_f32 + 0.005_f32 * target_entropy; + let expected_sigma = expected_sf * 2.0_f32; + println!("branch={b} n=3 sf={sf:.7} expected_sf={expected_sf:.7} \ + sigma={sigma:.7} expected_sigma={expected_sigma:.7}"); + let rel_sf = (sf - expected_sf).abs() / expected_sf; + assert!(rel_sf < 0.01, "branch={b}: sf={sf:.7} expected={expected_sf:.7} rel={rel_sf:.5}"); + let rel_sig = (sigma - expected_sigma).abs() / expected_sigma; + assert!(rel_sig < 0.01, "branch={b}: sigma={sigma:.7} expected={expected_sigma:.7} rel={rel_sig:.5}"); + assert!(sf > 0.1_f32, "branch={b}: SF should have increased from 0.1, got {sf}"); + } +} diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 625b8f18a..7f0e2d8f9 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,8 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +SP5 Task A2 — Pearl 3 per-branch NoisyNet sigma GPU producer (Layer A, 2026-05-01): one new CUDA kernel (`pearl_3_sigma_kernel.cu`) lands as a Layer A additive producer that feeds ISV slots [210..218) with per-branch NoisyNet sigma and sigma-fraction controller state. No NoisyLinear consumer migration in this commit — Layer B wires the noisy_linear_kernel.cu consumer. `pearl_3_sigma_update`: single-block 4-thread kernel (one thread per branch), reads ATOM_V_HALF[b] (Pearl 1, ISV[178..182)) + BRANCH_ENTROPY[b] (q_branch_stats, ISV[218..222)) + current SIGMA_FRACTION[b] (ISV[214..218), Pearl A sentinel 0 → bootstrap to 0.1); applies entropy-deficit controller `new_sf = clamp(sf + SF_LR * (target_entropy - entropy), 0.001, 0.5)` where `target_entropy = log(n_actions[b]) * 0.7` (SF_LR=0.005 and target_entropy_factor=0.7 are Invariant 1 structural anchors); computes `new_sigma = new_sf * max(v_half, EPS_CLAMP_FLOOR=1.0)` (EPS_CLAMP_FLOOR is an Invariant 1 anchor); writes (sigma[4], SF[4]) to scratch[103..111) (SCRATCH_PEARL_3_SIGMA=103, SCRATCH_PEARL_3_SF=107). `launch_sp5_pearl_3_sigma` fires 8 `launch_apply_pearls` calls: sigma[4] → ISV[210..214) (NOISY_SIGMA_BASE), SF[4] → ISV[214..218) (SIGMA_FRACTION_BASE). Wiener offset formula: `base_wiener_offset + (isv_slot - SP5_SLOT_BASE) * 3` where `base_wiener_offset = SP4_PRODUCER_COUNT * 3 = 213`. Pearl 3 must chain AFTER Pearl 1 (reads ATOM_V_HALF which Pearl 1 populates via apply_pearls_ad). Buffer growth: `producer_step_scratch_buf` 103→111 (SP5_SCRATCH_TOTAL updated; new constants SCRATCH_PEARL_3_SIGMA=103, SCRATCH_PEARL_3_SF=107). wiener_state_buf already at 543 (A1 sized for entire SP5 block — no further growth). StateResetRegistry gains 2 new FoldReset entries: sp5_noisy_sigma (ISV[210..214)) and sp5_sigma_fraction (ISV[214..218)) — Pearl A sentinel 0 at fold boundary triggers bootstrap to 0.1 on new fold's first launch. Wire-up in training_loop.rs: `launch_sp5_pearl_3_sigma()` called immediately after `launch_sp5_pearl_1_atom()` with `warn!` on error. Two GPU-only `#[ignore = "requires GPU"]` unit tests: (3) bootstrap + target-entropy equality case (sentinel ISV SF=0→bootstrap 0.1, target==actual → no update, sigma=0.1*2.0=0.2); (4) entropy-below-target case (SF=0.1, entropy=0 → SF increases by SF_LR*target_entropy, sigma increases proportionally). No CPU compute, no HtoD/HtoH, no atomicAdd, no stubs. Touched: `cuda_pipeline/pearl_3_sigma_kernel.cu` (new), `build.rs` (+1 cubin entry), `cuda_pipeline/gpu_dqn_trainer.rs` (static cubin + SP5_SCRATCH_TOTAL 103→111 + 2 new constants + struct field + cubin loader + struct initializer + launch method), `trainers/dqn/state_reset_registry.rs` (+2 FoldReset entries), `trainers/dqn/trainer/training_loop.rs` (+4 LOC after Pearl 1 launch), `tests/sp5_producer_unit_tests.rs` (+2 GPU-only tests + cubin loader helper). cargo check -p ml --lib clean. No consumer migration in this commit. + SP5 Task A1 fix-up — pearl_1 clip_rate denominator must be per-branch (2026-05-01): code-quality review caught the headroom controller's clip_rate computation using `batch_size × 13` (sum of per-branch action counts) instead of `batch_size × action_counts[b]` (per-branch). The TARGET_CLIP_RATE=0.01 anchor is calibrated against the per-branch event rate; using a denominator 3.25–4.3× too large made the controller systematically under-responsive — at actual 1% per-branch clipping it would read ~0.23–0.31% and contract headroom toward the 2.0 floor instead of holding the target. Currently masked because `atoms_clip_count` is zero in Layer A (Layer B's atoms_update_kernel migration populates it); without this fix Layer B would inherit a silently-wrong formula. Fix plumbs `action_counts[4]` ({4, 3, 3, 3}) through pearl_1_atom_kernel signature and uses `batch_size × action_counts[b]` as the per-branch denominator, matching q_branch_stats_kernel's existing parameter pattern. Test #2 launcher updated to allocate action_counts_buf and pass its dev_ptr through the new parameter slot. Touched: `cuda_pipeline/pearl_1_atom_kernel.cu` (+1 param, denominator math), `cuda_pipeline/gpu_dqn_trainer.rs` (+1 launch arg in `launch_sp5_pearl_1_atom`), `tests/sp5_producer_unit_tests.rs` (+5 LOC for action_counts_buf + extra arg). cargo check + cargo build --release + cargo test --no-run all clean. SP5 Task A1 — Pearl 1 atom-span + Q-stats GPU producers (Layer A, 2026-05-01): two new CUDA kernels (`q_branch_stats_kernel.cu`, `pearl_1_atom_kernel.cu`) land as Layer A additive producers that feed ISV slots [174..226) with per-branch C51 atom-span signals. No consumer migration in this commit — Layer B (atoms_update_kernel wiring) is a separate task. `q_branch_stats_update`: single-block 4-thread kernel (one thread per branch: Direction/Magnitude/Order/Urgency), reads `q_out_buf[B × 13]` row-major (confirmed layout), computes 3 serial passes over batch × branch-actions to produce (mean, max_abs_dev, var, entropy) per branch, writes 16 floats to `producer_step_scratch_buf[71..87)` (SCRATCH_Q_STATS=71; slot 70 = SP4 q_dir_grad_p99 was the last SP4 slot — collision-free). No atomicAdd per `feedback_no_atomicadd.md` — one thread per branch with independent scratch writes. `pearl_1_atom_update`: single-block 4-thread kernel, reads q_branch_stats scratch [71..87), reads current HEADROOM from ISV (Pearl A sentinel: ISV=0.0 → bootstrap to 6.0), reads `atoms_clip_count_buf[4]` (zero-init until Layer B wires atoms_update_kernel — clip_rate=0, headroom holds at bootstrap), runs headroom controller `new_headroom = clamp(headroom + 0.01*(clip_rate - 0.01), 2.0, 20.0)` (TARGET_CLIP_RATE=0.01 and HEADROOM_LR=0.01 are Invariant-1 structural anchors, not tuned constants), writes `(v_center, v_half, headroom, clip_rate)` to scratch [87..103) (SCRATCH_ATOM_OUT=87). `launch_sp5_pearl_1_atom` then fires 24 `launch_apply_pearls` calls: 4 groups × 4 contiguous Pearl 1 outputs (ISV[174..178)=v_center, [178..182)=v_half, [182..186)=headroom, [186..190)=clip_rate) looped, then 8 single-slot launches for per-branch entropy ([218..222)) and var ([222..226)) which are non-contiguous in scratch. Wiener offset: `(SP4_PRODUCER_COUNT=71 + (isv_slot − SP5_SLOT_BASE=174)) × 3` — correct SP5 block within wiener_state_buf. Buffer growth per `feedback_no_partial_refactor.md`: `wiener_state_buf` 213→543 floats ((71+110)×3), `producer_step_scratch_buf` 71→103 slots (SP4_PRODUCER_COUNT → SP5_SCRATCH_TOTAL), both in the same commit that adds the kernels that use them. New module-level constants: `SP5_WIENER_TOTAL_FLOATS=543`, `SP5_SCRATCH_TOTAL=103`, `SCRATCH_Q_STATS=71`, `SCRATCH_ATOM_OUT=87`. New struct fields: `q_branch_stats_kernel`, `pearl_1_atom_kernel` (CudaFunction), `atoms_clip_count_buf` (MappedI32Buffer[4], zero-init), `action_counts_buf` ([4,3,3,3]), `action_offsets_buf` ([0,4,7,10]). StateResetRegistry gains 7 new FoldReset entries: sp5_atom_v_center/v_half/headroom/clip_rate (ISV[174..190)), sp5_branch_entropy (ISV[218..222)), sp5_q_var_per_branch (ISV[222..226)), sp5_wiener_state (wiener_state_buf[213..543) — SP5 block, resets companion Wiener state with Pearl A sentinel per `pearl_first_observation_bootstrap.md`). Wire-up in training_loop.rs: `launch_sp5_pearl_1_atom()` called after `launch_sp4_param_group_oracles_all_groups` with `warn!` on error (Layer A — non-fatal until Layer B consumers land). Two GPU-only `#[ignore = "requires GPU"]` unit tests in `tests/sp5_producer_unit_tests.rs`: (1) uniform-input analytical ground truth for q_branch_stats (mean=2.0, max_abs_dev=0.0, entropy=log(n_actions)); (2) headroom bootstrap + controller step for pearl_1_atom (sentinel ISV=0→bootstrap 6.0, clip_rate=0, one step adjustment). No CPU compute per `feedback_no_cpu_compute_strict.md`, no HtoD/HtoH per `feedback_no_htod_htoh_only_mapped_pinned.md`, no atomicAdd per `feedback_no_atomicadd.md`, no stubs per `feedback_no_stubs.md`. Touched: `cuda_pipeline/q_branch_stats_kernel.cu` (new), `cuda_pipeline/pearl_1_atom_kernel.cu` (new), `build.rs` (+2 cubin entries), `cuda_pipeline/gpu_dqn_trainer.rs` (static cubins + SP5 constants + struct fields + buffer growth 213→543 + 71→103 + cubin loaders + buffer allocs + struct initializer + launch method + docstring), `trainers/dqn/state_reset_registry.rs` (+7 FoldReset entries + sp4_wiener_state description 213→543), `trainers/dqn/trainer/training_loop.rs` (+3 LOC wire-up after SP4 param_group_oracles), `tests/sp5_producer_unit_tests.rs` (new, 2 GPU-only tests). cargo check -p ml --lib clean (11 pre-existing warnings, none new). Validation deferred to one L40S smoke; expected sp5_atom_headroom stays [2.0,20.0], sp5_atom_v_half tracks C51 atom spread, no NaN flags.