feat(sp5): Task A2 — Pearl 3 per-branch NoisyNet sigma
Per-branch sigma scales with per-branch Q magnitude (v_half from Pearl 1's ATOM_V_HALF, populated in A1). SIGMA_FRACTION adapts via entropy-deficit controller targeting 70% of max action entropy (target_entropy = log(n_actions[b]) * 0.7). 8 ISV slots (NOISY_SIGMA[210..214), SIGMA_FRACTION[214..218)). Reuses BRANCH_ENTROPY (218..222) from Task A1's q_branch_stats_kernel. producer_step_scratch_buf grew 103 -> 111 (8 new outputs). wiener_state_buf already at 543 (A1 sized for entire SP5 block). NoisyLinear consumer migration deferred to Layer B. Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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
|
||||
|
||||
@@ -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,
|
||||
|
||||
61
crates/ml/src/cuda_pipeline/pearl_3_sigma_kernel.cu
Normal file
61
crates/ml/src/cuda_pipeline/pearl_3_sigma_kernel.cu
Normal file
@@ -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 <cuda_runtime.h>
|
||||
|
||||
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();
|
||||
}
|
||||
@@ -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,
|
||||
|
||||
@@ -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}");
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
@@ -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<CudaStream>) -> 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<CudaStream> {
|
||||
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}");
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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.
|
||||
|
||||
Reference in New Issue
Block a user