feat(sp5): Task A3 — Pearl 2 per-branch loss budget
Per-branch C51/IQN/CQL/Ens loss budgets driven by per-branch flatness = var(Q[b]) / (σ[b]² + EPS_DIV). IQN dominates when flat, C51 yields proportionally. CQL stays gated by existing regime_stability allocator. Reads Q_VAR_PER_BRANCH (222..226 from Pearl 1's q_branch_stats) AND NOISY_SIGMA (210..214 from Pearl 3 — must run AFTER both). 20 ISV slots (BUDGET_C51[190..194), BUDGET_IQN[194..198), BUDGET_CQL[198..202), BUDGET_ENS[202..206), FLATNESS[206..210)). StateResetRegistry: 5 new FoldReset entries. producer_step_scratch_buf grew 111 → 131 (20 new outputs). wiener_state_buf already at 543 (A1 sized for entire SP5 block). Loss budget consumer migration deferred to Layer B. Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -421,6 +421,15 @@ fn main() {
|
||||
// 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",
|
||||
// SP5 Task A3 (2026-05-01): Pearl 2 per-branch loss budget controller.
|
||||
// Single-block 4-thread kernel reads Q_VAR_PER_BRANCH (Pearl 1's
|
||||
// q_branch_stats, ISV[222..226)) + NOISY_SIGMA (Pearl 3 output,
|
||||
// ISV[210..214)); computes per-branch flatness = clamp(var_q/(σ²+EPS_DIV),
|
||||
// 0, 1) and derives (budget_c51, budget_iqn, budget_cql=0, budget_ens,
|
||||
// flatness) per branch. Writes 20 floats to scratch[111..131).
|
||||
// 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",
|
||||
// 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
|
||||
|
||||
@@ -304,6 +304,20 @@ static SP5_PEARL_1_ATOM_CUBIN: &[u8] =
|
||||
static SP5_PEARL_3_SIGMA_CUBIN: &[u8] =
|
||||
include_bytes!(concat!(env!("OUT_DIR"), "/pearl_3_sigma_kernel.cubin"));
|
||||
|
||||
/// SP5 Task A3 (2026-05-01): Pearl 2 per-branch loss budget cubin.
|
||||
/// Single-block 4-thread kernel. 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 per-branch flatness = clamp(var_q / (σ²+EPS_DIV),
|
||||
/// 0, 1) and derives (budget_iqn, budget_c51, budget_cql=0, budget_ens,
|
||||
/// flatness) per branch. Writes 20 floats to
|
||||
/// producer_step_scratch_buf[111..131). Chained apply_pearls_ad_kernel
|
||||
/// (20 single-slot launches) smooths via Pearls A+D into ISV[190..210).
|
||||
/// Must run AFTER launch_sp5_pearl_1_atom (Q_VAR) AND
|
||||
/// launch_sp5_pearl_3_sigma (NOISY_SIGMA).
|
||||
/// Layer A only — loss-budget consumer migration deferred to Layer B.
|
||||
static SP5_PEARL_2_BUDGET_CUBIN: &[u8] =
|
||||
include_bytes!(concat!(env!("OUT_DIR"), "/pearl_2_budget_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
|
||||
@@ -887,15 +901,22 @@ 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: producer_step_scratch_buf total slot count after growth.
|
||||
/// SP5 Task A1+A2+A3: 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)
|
||||
/// 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;
|
||||
/// Task A3 adds:
|
||||
/// 20 floats for pearl_2_budget output (5 outputs × 4 branches at [111..131))
|
||||
/// budget_c51[4] → SCRATCH_PEARL_2_C51=111
|
||||
/// budget_iqn[4] → SCRATCH_PEARL_2_IQN=115
|
||||
/// 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;
|
||||
|
||||
/// SP5 Task A2: scratch index base for pearl_3_sigma sigma[4] output block.
|
||||
/// Slots [103..107): per-branch NoisyNet sigma.
|
||||
@@ -905,6 +926,27 @@ pub const SCRATCH_PEARL_3_SIGMA: usize = 103;
|
||||
/// Slots [107..111): per-branch sigma fraction (entropy-deficit controller state).
|
||||
pub const SCRATCH_PEARL_3_SF: usize = 107;
|
||||
|
||||
/// SP5 Task A3: scratch index base for pearl_2_budget budget_c51[4] output block.
|
||||
/// Slots [111..115): per-branch C51 loss budget share.
|
||||
pub const SCRATCH_PEARL_2_C51: usize = 111;
|
||||
|
||||
/// SP5 Task A3: scratch index base for pearl_2_budget budget_iqn[4] output block.
|
||||
/// Slots [115..119): per-branch IQN loss budget share.
|
||||
pub const SCRATCH_PEARL_2_IQN: usize = 115;
|
||||
|
||||
/// SP5 Task A3: scratch index base for pearl_2_budget budget_cql[4] output block.
|
||||
/// Slots [119..123): per-branch CQL loss budget share (zero in Pearl 2; gated by
|
||||
/// existing regime_stability allocator).
|
||||
pub const SCRATCH_PEARL_2_CQL: usize = 119;
|
||||
|
||||
/// SP5 Task A3: scratch index base for pearl_2_budget budget_ens[4] output block.
|
||||
/// Slots [123..127): per-branch ensemble loss budget share (1 - IQN - C51 - CQL, ≥0).
|
||||
pub const SCRATCH_PEARL_2_ENS: usize = 123;
|
||||
|
||||
/// SP5 Task A3: scratch index base for pearl_2_budget flatness[4] output block.
|
||||
/// Slots [127..131): per-branch flatness signal (var_q / (σ² + EPS_DIV), clamped [0,1]).
|
||||
pub const SCRATCH_PEARL_2_FLATNESS: usize = 127;
|
||||
|
||||
/// 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).
|
||||
@@ -4039,6 +4081,20 @@ pub struct GpuDqnTrainer {
|
||||
/// 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 A3 (2026-05-01): Pearl 2 per-branch loss budget controller kernel.
|
||||
/// Single-block 4-thread kernel (one thread per branch). Reads Q_VAR_PER_BRANCH
|
||||
/// (Pearl 1's q_branch_stats, ISV[222..226)) + NOISY_SIGMA (Pearl 3 output,
|
||||
/// ISV[210..214)); computes per-branch flatness = clamp(var_q / (σ²+EPS_DIV), 0, 1)
|
||||
/// and derives (budget_c51, budget_iqn, budget_cql=0, budget_ens, flatness) per
|
||||
/// branch. Writes 20 floats to `producer_step_scratch_buf[111..131)`.
|
||||
/// Chained `apply_pearls_ad_kernel` (20 single-slot launches) then smooths via
|
||||
/// Pearls A+D into ISV[190..210).
|
||||
/// Must run AFTER launch_sp5_pearl_1_atom (Q_VAR) AND
|
||||
/// launch_sp5_pearl_3_sigma (NOISY_SIGMA). Layer A only — loss-budget consumer
|
||||
/// migration deferred to Layer B.
|
||||
/// 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 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.
|
||||
@@ -10608,6 +10664,106 @@ impl GpuDqnTrainer {
|
||||
Ok(())
|
||||
}
|
||||
|
||||
/// SP5 Task A3 (2026-05-01): Pearl 2 — per-branch loss budget from per-branch
|
||||
/// flatness signal.
|
||||
///
|
||||
/// Chained AFTER `launch_sp5_pearl_1_atom` (reads Q_VAR_PER_BRANCH written by
|
||||
/// Pearl 1's q_branch_stats) AND AFTER `launch_sp5_pearl_3_sigma` (reads
|
||||
/// NOISY_SIGMA written by Pearl 3). Two-step sequence:
|
||||
/// 1. `pearl_2_budget_update` (1×4 kernel) → scratch[111..131).
|
||||
/// 2. `apply_pearls_ad_kernel` × 20 (one per output slot):
|
||||
/// budget_c51[4] → ISV[BUDGET_C51_BASE=190..194)
|
||||
/// budget_iqn[4] → ISV[BUDGET_IQN_BASE=194..198)
|
||||
/// budget_cql[4] → ISV[BUDGET_CQL_BASE=198..202)
|
||||
/// budget_ens[4] → ISV[BUDGET_ENS_BASE=202..206)
|
||||
/// flatness[4] → ISV[FLATNESS_BASE=206..210)
|
||||
///
|
||||
/// Wiener offsets: (SP4_PRODUCER_COUNT=71 + (isv_slot − SP5_SLOT_BASE=174)) × 3.
|
||||
/// Layer A only — loss-budget consumer migration deferred to Layer B.
|
||||
pub(crate) fn launch_sp5_pearl_2_budget(&self) -> Result<(), MLError> {
|
||||
use crate::cuda_pipeline::sp4_wiener_ema::launch_apply_pearls;
|
||||
use crate::cuda_pipeline::sp5_isv_slots::{
|
||||
SP5_SLOT_BASE,
|
||||
BUDGET_C51_BASE, BUDGET_IQN_BASE, BUDGET_CQL_BASE, BUDGET_ENS_BASE,
|
||||
FLATNESS_BASE, NOISY_SIGMA_BASE, Q_VAR_PER_BRANCH_BASE,
|
||||
};
|
||||
|
||||
debug_assert!(self.isv_signals_dev_ptr != 0,
|
||||
"launch_sp5_pearl_2_budget: 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;
|
||||
|
||||
// Step 1: pearl_2_budget_update → scratch[111..131).
|
||||
// Kernel signature: (isv_signals, q_var_isv_base, sigma_isv_base,
|
||||
// scratch_out, c51_idx_base, iqn_idx_base,
|
||||
// cql_idx_base, ens_idx_base, flatness_idx_base)
|
||||
let q_var_base_i32 = Q_VAR_PER_BRANCH_BASE as i32; // 222
|
||||
let sigma_base_i32 = NOISY_SIGMA_BASE as i32; // 210
|
||||
let c51_base_i32 = SCRATCH_PEARL_2_C51 as i32; // 111
|
||||
let iqn_base_i32 = SCRATCH_PEARL_2_IQN as i32; // 115
|
||||
let cql_base_i32 = SCRATCH_PEARL_2_CQL as i32; // 119
|
||||
let ens_base_i32 = SCRATCH_PEARL_2_ENS as i32; // 123
|
||||
let flatness_base_i32 = SCRATCH_PEARL_2_FLATNESS as i32; // 127
|
||||
|
||||
unsafe {
|
||||
self.stream
|
||||
.launch_builder(&self.pearl_2_budget_kernel)
|
||||
.arg(&isv_dev)
|
||||
.arg(&q_var_base_i32)
|
||||
.arg(&sigma_base_i32)
|
||||
.arg(&scratch_dev)
|
||||
.arg(&c51_base_i32)
|
||||
.arg(&iqn_base_i32)
|
||||
.arg(&cql_base_i32)
|
||||
.arg(&ens_base_i32)
|
||||
.arg(&flatness_base_i32)
|
||||
.launch(LaunchConfig {
|
||||
grid_dim: (1, 1, 1),
|
||||
block_dim: (4, 1, 1),
|
||||
shared_mem_bytes: 0,
|
||||
})
|
||||
.map_err(|e| MLError::ModelError(format!("pearl_2_budget_update: {e}")))?;
|
||||
}
|
||||
|
||||
// Step 2: apply_pearls_ad_kernel × 20 — one single-slot launch per ISV output.
|
||||
// Wiener offset: (SP4_PRODUCER_COUNT + (isv_slot - SP5_SLOT_BASE)) * 3.
|
||||
//
|
||||
// budget_c51[4] → scratch [111..115), ISV [190..194) (BUDGET_C51_BASE)
|
||||
// budget_iqn[4] → scratch [115..119), ISV [194..198) (BUDGET_IQN_BASE)
|
||||
// budget_cql[4] → scratch [119..123), ISV [198..202) (BUDGET_CQL_BASE)
|
||||
// budget_ens[4] → scratch [123..127), ISV [202..206) (BUDGET_ENS_BASE)
|
||||
// flatness[4] → scratch [127..131), ISV [206..210) (FLATNESS_BASE)
|
||||
let base_wiener_offset = SP4_PRODUCER_COUNT as i32 * 3; // 213
|
||||
|
||||
for (isv_base, scratch_base) in [
|
||||
(BUDGET_C51_BASE, SCRATCH_PEARL_2_C51),
|
||||
(BUDGET_IQN_BASE, SCRATCH_PEARL_2_IQN),
|
||||
(BUDGET_CQL_BASE, SCRATCH_PEARL_2_CQL),
|
||||
(BUDGET_ENS_BASE, SCRATCH_PEARL_2_ENS),
|
||||
(FLATNESS_BASE, SCRATCH_PEARL_2_FLATNESS),
|
||||
] {
|
||||
for b in 0..4_usize {
|
||||
let scratch_idx = (scratch_base + b) as i32;
|
||||
let isv_idx = (isv_base + b) 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,
|
||||
)?;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
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].
|
||||
///
|
||||
@@ -12617,6 +12773,17 @@ impl GpuDqnTrainer {
|
||||
.map_err(|e| MLError::ModelError(format!("pearl_3_sigma_update load: {e}")))?
|
||||
};
|
||||
|
||||
// SP5 Task A3 (2026-05-01): load pearl_2_budget loss budget controller kernel.
|
||||
// Single-block 4-thread kernel; reads Q_VAR_PER_BRANCH (ISV[222..226)) +
|
||||
// NOISY_SIGMA (ISV[210..214)) from ISV; writes 20 floats to scratch[111..131).
|
||||
// Must be invoked AFTER pearl_1_atom (Q_VAR) AND pearl_3_sigma (NOISY_SIGMA).
|
||||
let pearl_2_budget_kernel = {
|
||||
let module = stream.context().load_cubin(SP5_PEARL_2_BUDGET_CUBIN.to_vec())
|
||||
.map_err(|e| MLError::ModelError(format!("sp5 pearl_2_budget cubin load: {e}")))?;
|
||||
module.load_function("pearl_2_budget_update")
|
||||
.map_err(|e| MLError::ModelError(format!("pearl_2_budget_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.
|
||||
@@ -15738,6 +15905,7 @@ impl GpuDqnTrainer {
|
||||
q_branch_stats_kernel,
|
||||
pearl_1_atom_kernel,
|
||||
pearl_3_sigma_kernel,
|
||||
pearl_2_budget_kernel,
|
||||
atoms_clip_count_buf,
|
||||
action_counts_buf,
|
||||
action_offsets_buf,
|
||||
|
||||
69
crates/ml/src/cuda_pipeline/pearl_2_budget_kernel.cu
Normal file
69
crates/ml/src/cuda_pipeline/pearl_2_budget_kernel.cu
Normal file
@@ -0,0 +1,69 @@
|
||||
// crates/ml/src/cuda_pipeline/pearl_2_budget_kernel.cu
|
||||
//
|
||||
// SP5 Task A3: Pearl 2 — per-branch C51/IQN/CQL/Ens loss budgets driven by
|
||||
// per-branch flatness signal.
|
||||
//
|
||||
// Reads Q_VAR_PER_BRANCH[b] (ISV[222..226), Pearl 1's q_branch_stats output) and
|
||||
// NOISY_SIGMA[b] (ISV[210..214), Pearl 3 output). Computes:
|
||||
//
|
||||
// flatness[b] = clamp(var_q[b] / (σ[b]² + EPS_DIV), 0, 1)
|
||||
//
|
||||
// When flatness ≈ 1 (Q values are flat relative to noise σ), IQN dominates;
|
||||
// when flatness ≈ 0 (Q values are well-separated), C51 takes a proportional share.
|
||||
//
|
||||
// budget_iqn[b] = BASE_IQN + (1 - flatness[b]) × (1 - BASE_IQN)
|
||||
// budget_c51[b] = flatness[b] × (1 - BASE_IQN) × BASE_C51_SHARE
|
||||
// budget_cql[b] = 0 // CQL stays gated by existing regime_stability allocator
|
||||
// budget_ens[b] = max(0, 1 - iqn - c51 - cql)
|
||||
//
|
||||
// Writes per-branch (budget_c51, budget_iqn, budget_cql, budget_ens, flatness)
|
||||
// to scratch_out. apply_pearls_ad_kernel then smooths these via Pearls A+D.
|
||||
//
|
||||
// 4-thread single-block (one thread per branch). No atomicAdd.
|
||||
// __threadfence_system() after writes.
|
||||
//
|
||||
// BASE_IQN=0.11, BASE_C51_SHARE=0.84/(1-BASE_IQN), and EPS_DIV=1e-8 are
|
||||
// Invariant 1 structural anchors (preserve SP4-baseline budget under non-flat
|
||||
// regime; small numerical-stability epsilon).
|
||||
|
||||
#include <cuda_runtime.h>
|
||||
|
||||
extern "C" __global__ void pearl_2_budget_update(
|
||||
const float* __restrict__ isv_signals,
|
||||
int q_var_isv_base, // Q_VAR_PER_BRANCH_BASE = 222
|
||||
int sigma_isv_base, // NOISY_SIGMA_BASE = 210
|
||||
float* __restrict__ scratch_out,
|
||||
int c51_idx_base,
|
||||
int iqn_idx_base,
|
||||
int cql_idx_base,
|
||||
int ens_idx_base,
|
||||
int flatness_idx_base
|
||||
) {
|
||||
int b = threadIdx.x;
|
||||
if (blockIdx.x != 0 || b >= 4) return;
|
||||
|
||||
// BASE_IQN = 0.11 is an Invariant 1 anchor (SP4 baseline budget share).
|
||||
const float BASE_IQN = 0.11f;
|
||||
// BASE_C51_SHARE = 0.84 / (1 - BASE_IQN) is an Invariant 1 anchor
|
||||
// (proportional remainder preserving SP4 baseline C51 share under flat regime).
|
||||
const float BASE_C51_SHARE = 0.84f / (1.0f - BASE_IQN);
|
||||
// EPS_DIV = 1e-8 is an Invariant 1 anchor (numerical stability denominator floor).
|
||||
const float EPS_DIV = 1e-8f;
|
||||
|
||||
float var_q = isv_signals[q_var_isv_base + b];
|
||||
float sigma = isv_signals[sigma_isv_base + b];
|
||||
float sigma_sq = sigma * sigma + EPS_DIV;
|
||||
float flatness = fminf(1.0f, fmaxf(0.0f, var_q / sigma_sq));
|
||||
|
||||
float budget_iqn = BASE_IQN + (1.0f - flatness) * (1.0f - BASE_IQN);
|
||||
float budget_c51 = flatness * (1.0f - BASE_IQN) * BASE_C51_SHARE;
|
||||
float budget_cql = 0.0f;
|
||||
float budget_ens = fmaxf(0.0f, 1.0f - budget_iqn - budget_c51 - budget_cql);
|
||||
|
||||
scratch_out[c51_idx_base + b] = budget_c51;
|
||||
scratch_out[iqn_idx_base + b] = budget_iqn;
|
||||
scratch_out[cql_idx_base + b] = budget_cql;
|
||||
scratch_out[ens_idx_base + b] = budget_ens;
|
||||
scratch_out[flatness_idx_base + b] = flatness;
|
||||
__threadfence_system();
|
||||
}
|
||||
@@ -572,6 +572,31 @@ 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_budget_c51",
|
||||
category: ResetCategory::FoldReset,
|
||||
description: "ISV[BUDGET_C51_BASE=190..194) — per-branch C51 loss budget share (driven by flatness × (1 - BASE_IQN) × BASE_C51_SHARE). SP5 Pearl A sentinel 0 at fold boundary so the new fold's first pearl_2_budget_update launch fires the first-observation replacement (Task A3).",
|
||||
},
|
||||
RegistryEntry {
|
||||
name: "sp5_budget_iqn",
|
||||
category: ResetCategory::FoldReset,
|
||||
description: "ISV[BUDGET_IQN_BASE=194..198) — per-branch IQN loss budget share (BASE_IQN + (1-flatness)×(1-BASE_IQN), dominates when Q is flat relative to noise σ). SP5 Pearl A sentinel 0 at fold boundary (Task A3).",
|
||||
},
|
||||
RegistryEntry {
|
||||
name: "sp5_budget_cql",
|
||||
category: ResetCategory::FoldReset,
|
||||
description: "ISV[BUDGET_CQL_BASE=198..202) — per-branch CQL loss budget share (zero in Pearl 2; CQL gating remains under the existing regime_stability allocator). SP5 Pearl A sentinel 0 at fold boundary (Task A3).",
|
||||
},
|
||||
RegistryEntry {
|
||||
name: "sp5_budget_ens",
|
||||
category: ResetCategory::FoldReset,
|
||||
description: "ISV[BUDGET_ENS_BASE=202..206) — per-branch ensemble loss budget share (1 - IQN - C51 - CQL, clamped to ≥0). SP5 Pearl A sentinel 0 at fold boundary (Task A3).",
|
||||
},
|
||||
RegistryEntry {
|
||||
name: "sp5_flatness",
|
||||
category: ResetCategory::FoldReset,
|
||||
description: "ISV[FLATNESS_BASE=206..210) — per-branch flatness signal (var(Q[b]) / (σ[b]² + EPS_DIV), clamped to [0, 1]). Drives Pearl 2's per-branch IQN/C51 split. SP5 Pearl A sentinel 0 at fold boundary (Task A3).",
|
||||
},
|
||||
RegistryEntry {
|
||||
name: "sp5_noisy_sigma",
|
||||
category: ResetCategory::FoldReset,
|
||||
|
||||
@@ -3596,6 +3596,13 @@ impl DQNTrainer {
|
||||
if let Err(e) = fused.trainer().launch_sp5_pearl_3_sigma() {
|
||||
tracing::warn!("SP5 A2 launch_sp5_pearl_3_sigma failed: {e}");
|
||||
}
|
||||
// SP5 Task A3: pearl_2_budget → apply_pearls_ad ×20
|
||||
// Must run AFTER launch_sp5_pearl_1_atom (Q_VAR) AND
|
||||
// launch_sp5_pearl_3_sigma (NOISY_SIGMA).
|
||||
// Producer-only (Layer A) — no consumer reads ISV[190..210) in this commit.
|
||||
if let Err(e) = fused.trainer().launch_sp5_pearl_2_budget() {
|
||||
tracing::warn!(error = %e, "SP5 Pearl 2 producer launch failed");
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
@@ -30,6 +30,9 @@ const SP5_PEARL_1_ATOM_CUBIN: &[u8] =
|
||||
const SP5_PEARL_3_SIGMA_CUBIN: &[u8] =
|
||||
include_bytes!(concat!(env!("OUT_DIR"), "/pearl_3_sigma_kernel.cubin"));
|
||||
|
||||
const SP5_PEARL_2_BUDGET_CUBIN: &[u8] =
|
||||
include_bytes!(concat!(env!("OUT_DIR"), "/pearl_2_budget_kernel.cubin"));
|
||||
|
||||
// ── Helpers ───────────────────────────────────────────────────────────────────
|
||||
|
||||
fn load_pearl_3_sigma_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
|
||||
@@ -42,6 +45,16 @@ fn load_pearl_3_sigma_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
|
||||
.expect("load pearl_3_sigma_update function")
|
||||
}
|
||||
|
||||
fn load_pearl_2_budget_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
|
||||
let module = stream
|
||||
.context()
|
||||
.load_cubin(SP5_PEARL_2_BUDGET_CUBIN.to_vec())
|
||||
.expect("load pearl_2_budget_kernel cubin");
|
||||
module
|
||||
.load_function("pearl_2_budget_update")
|
||||
.expect("load pearl_2_budget_update function")
|
||||
}
|
||||
|
||||
fn make_test_stream() -> Arc<CudaStream> {
|
||||
let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?");
|
||||
ctx.default_stream()
|
||||
@@ -563,3 +576,271 @@ fn pearl_3_sigma_fraction_increases_when_entropy_below_target() {
|
||||
assert!(sf > 0.1_f32, "branch={b}: SF should have increased from 0.1, got {sf}");
|
||||
}
|
||||
}
|
||||
|
||||
// ── Tests 5 + 6: pearl_2_budget_update ────────────────────────────────────────
|
||||
|
||||
/// SP5 Task A3 unit test: `pearl_2_budget_update` on flat-Q regime.
|
||||
///
|
||||
/// Synthetic ISV: σ=1.0 (all branches), var_q=2.0 (all branches).
|
||||
/// flatness = clamp(2.0 / (1.0² + 1e-8), 0, 1) = clamp(~2.0, 0, 1) = 1.0
|
||||
///
|
||||
/// Analytical expected (flatness = 1.0):
|
||||
/// BASE_IQN = 0.11
|
||||
/// BASE_C51_SHARE = 0.84 / (1 - 0.11) = 0.84 / 0.89 ≈ 0.94382...
|
||||
/// budget_iqn = 0.11 + (1 - 1.0) * 0.89 = 0.11
|
||||
/// budget_c51 = 1.0 * 0.89 * 0.94382 ≈ 0.84
|
||||
/// budget_cql = 0
|
||||
/// budget_ens = max(0, 1 - 0.11 - 0.84 - 0) = 0.05
|
||||
///
|
||||
/// All 4 branches receive identical inputs, so identical outputs.
|
||||
/// No CPU reference oracle per `feedback_no_cpu_test_fallbacks.md`.
|
||||
#[test]
|
||||
#[ignore = "requires GPU"]
|
||||
fn pearl_2_budget_flat_regime_iqn_floor_c51_takes_remainder() {
|
||||
// Scratch layout mirrors SCRATCH_PEARL_2_* constants:
|
||||
// [111..115) = budget_c51[4]
|
||||
// [115..119) = budget_iqn[4]
|
||||
// [119..123) = budget_cql[4]
|
||||
// [123..127) = budget_ens[4]
|
||||
// [127..131) = flatness[4]
|
||||
const SCRATCH_C51: usize = 111;
|
||||
const SCRATCH_IQN: usize = 115;
|
||||
const SCRATCH_CQL: usize = 119;
|
||||
const SCRATCH_ENS: usize = 123;
|
||||
const SCRATCH_FLATNESS: usize = 127;
|
||||
const SCRATCH_SIZE: usize = 131;
|
||||
|
||||
// ISV: Q_VAR_PER_BRANCH_BASE=222, NOISY_SIGMA_BASE=210.
|
||||
// Need ISV size >= 222 + 4 = 226.
|
||||
const ISV_SIZE: usize = 230;
|
||||
// Q_VAR_PER_BRANCH_BASE = 222, NOISY_SIGMA_BASE = 210.
|
||||
const Q_VAR_BASE: usize = 222;
|
||||
const SIGMA_BASE: usize = 210;
|
||||
|
||||
let stream = make_test_stream();
|
||||
let kernel = load_pearl_2_budget_kernel(&stream);
|
||||
|
||||
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[Q_VAR_BASE + b] = 2.0_f32; // var_q = 2.0 (above σ² → clamp to 1.0 flatness)
|
||||
isv_data[SIGMA_BASE + b] = 1.0_f32; // σ = 1.0
|
||||
}
|
||||
isv_buf.write_from_slice(&isv_data);
|
||||
|
||||
let scratch_buf = unsafe { MappedF32Buffer::new(SCRATCH_SIZE) }
|
||||
.expect("alloc scratch_buf");
|
||||
|
||||
let isv_dev = isv_buf.dev_ptr;
|
||||
let q_var_base_i32 = Q_VAR_BASE as i32;
|
||||
let sigma_base_i32 = SIGMA_BASE as i32;
|
||||
let scratch_dev = scratch_buf.dev_ptr;
|
||||
let c51_base_i32 = SCRATCH_C51 as i32;
|
||||
let iqn_base_i32 = SCRATCH_IQN as i32;
|
||||
let cql_base_i32 = SCRATCH_CQL as i32;
|
||||
let ens_base_i32 = SCRATCH_ENS as i32;
|
||||
let flatness_base_i32 = SCRATCH_FLATNESS as i32;
|
||||
|
||||
unsafe {
|
||||
stream
|
||||
.launch_builder(&kernel)
|
||||
.arg(&isv_dev)
|
||||
.arg(&q_var_base_i32)
|
||||
.arg(&sigma_base_i32)
|
||||
.arg(&scratch_dev)
|
||||
.arg(&c51_base_i32)
|
||||
.arg(&iqn_base_i32)
|
||||
.arg(&cql_base_i32)
|
||||
.arg(&ens_base_i32)
|
||||
.arg(&flatness_base_i32)
|
||||
.launch(LaunchConfig {
|
||||
grid_dim: (1, 1, 1),
|
||||
block_dim: (4, 1, 1),
|
||||
shared_mem_bytes: 0,
|
||||
})
|
||||
.expect("launch pearl_2_budget_update (flat regime)");
|
||||
}
|
||||
stream.synchronize().expect("sync after pearl_2_budget_update (flat regime)");
|
||||
|
||||
let scratch = scratch_buf.read_all();
|
||||
|
||||
// Analytical values for flatness = 1.0 (saturated ceiling):
|
||||
// BASE_IQN = 0.11, BASE_C51_SHARE = 0.84 / 0.89
|
||||
const BASE_IQN: f32 = 0.11_f32;
|
||||
const BASE_C51_SHARE: f32 = 0.84_f32 / (1.0_f32 - BASE_IQN);
|
||||
let expected_flatness = 1.0_f32;
|
||||
let expected_iqn = BASE_IQN + (1.0_f32 - expected_flatness) * (1.0_f32 - BASE_IQN);
|
||||
let expected_c51 = expected_flatness * (1.0_f32 - BASE_IQN) * BASE_C51_SHARE;
|
||||
let expected_cql = 0.0_f32;
|
||||
let expected_ens = (1.0_f32 - expected_iqn - expected_c51 - expected_cql).max(0.0_f32);
|
||||
|
||||
for b in 0..4_usize {
|
||||
let c51 = scratch[SCRATCH_C51 + b];
|
||||
let iqn = scratch[SCRATCH_IQN + b];
|
||||
let cql = scratch[SCRATCH_CQL + b];
|
||||
let ens = scratch[SCRATCH_ENS + b];
|
||||
let flatness = scratch[SCRATCH_FLATNESS + b];
|
||||
|
||||
println!(
|
||||
"branch={b} flatness={flatness:.6} iqn={iqn:.6} c51={c51:.6} \
|
||||
cql={cql:.6} ens={ens:.6}"
|
||||
);
|
||||
|
||||
let rel_flatness = (flatness - expected_flatness).abs();
|
||||
assert!(
|
||||
rel_flatness < 1e-5,
|
||||
"branch={b}: flatness={flatness:.6} expected={expected_flatness:.6}"
|
||||
);
|
||||
let rel_iqn = (iqn - expected_iqn).abs() / expected_iqn.max(1e-9);
|
||||
assert!(
|
||||
rel_iqn < 0.01,
|
||||
"branch={b}: iqn={iqn:.6} expected={expected_iqn:.6} rel={rel_iqn:.5}"
|
||||
);
|
||||
let rel_c51 = (c51 - expected_c51).abs() / expected_c51.max(1e-9);
|
||||
assert!(
|
||||
rel_c51 < 0.01,
|
||||
"branch={b}: c51={c51:.6} expected={expected_c51:.6} rel={rel_c51:.5}"
|
||||
);
|
||||
assert!(
|
||||
cql.abs() < 1e-7,
|
||||
"branch={b}: cql={cql:.6} expected=0"
|
||||
);
|
||||
let rel_ens = (ens - expected_ens).abs() / expected_ens.max(1e-9);
|
||||
assert!(
|
||||
rel_ens < 0.01,
|
||||
"branch={b}: ens={ens:.6} expected={expected_ens:.6} rel={rel_ens:.5}"
|
||||
);
|
||||
}
|
||||
}
|
||||
|
||||
/// SP5 Task A3 unit test: `pearl_2_budget_update` on sharp-Q regime.
|
||||
///
|
||||
/// Synthetic ISV: σ=2.0 (all branches), var_q=0.04 (all branches).
|
||||
/// flatness = clamp(0.04 / (4.0 + 1e-8), 0, 1) = 0.04/4.0 = 0.01
|
||||
///
|
||||
/// Analytical expected (flatness = 0.01):
|
||||
/// BASE_IQN = 0.11
|
||||
/// BASE_C51_SHARE = 0.84 / 0.89 ≈ 0.94382
|
||||
/// budget_iqn = 0.11 + (1 - 0.01) * 0.89 = 0.11 + 0.8811 = 0.9911
|
||||
/// budget_c51 = 0.01 * 0.89 * 0.94382 ≈ 0.0084
|
||||
/// budget_cql = 0
|
||||
/// budget_ens = max(0, 1 - 0.9911 - 0.0084 - 0) ≈ 0.0005 (rounds to ~0)
|
||||
///
|
||||
/// IQN dominates in the sharp-Q regime, C51 share is minimal.
|
||||
/// No CPU reference oracle per `feedback_no_cpu_test_fallbacks.md`.
|
||||
#[test]
|
||||
#[ignore = "requires GPU"]
|
||||
fn pearl_2_budget_sharp_regime_iqn_dominates() {
|
||||
const SCRATCH_C51: usize = 111;
|
||||
const SCRATCH_IQN: usize = 115;
|
||||
const SCRATCH_CQL: usize = 119;
|
||||
const SCRATCH_ENS: usize = 123;
|
||||
const SCRATCH_FLATNESS: usize = 127;
|
||||
const SCRATCH_SIZE: usize = 131;
|
||||
const ISV_SIZE: usize = 230;
|
||||
const Q_VAR_BASE: usize = 222;
|
||||
const SIGMA_BASE: usize = 210;
|
||||
|
||||
let stream = make_test_stream();
|
||||
let kernel = load_pearl_2_budget_kernel(&stream);
|
||||
|
||||
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[Q_VAR_BASE + b] = 0.04_f32; // var_q small
|
||||
isv_data[SIGMA_BASE + b] = 2.0_f32; // σ large → flatness small
|
||||
}
|
||||
isv_buf.write_from_slice(&isv_data);
|
||||
|
||||
let scratch_buf = unsafe { MappedF32Buffer::new(SCRATCH_SIZE) }
|
||||
.expect("alloc scratch_buf");
|
||||
|
||||
let isv_dev = isv_buf.dev_ptr;
|
||||
let q_var_base_i32 = Q_VAR_BASE as i32;
|
||||
let sigma_base_i32 = SIGMA_BASE as i32;
|
||||
let scratch_dev = scratch_buf.dev_ptr;
|
||||
let c51_base_i32 = SCRATCH_C51 as i32;
|
||||
let iqn_base_i32 = SCRATCH_IQN as i32;
|
||||
let cql_base_i32 = SCRATCH_CQL as i32;
|
||||
let ens_base_i32 = SCRATCH_ENS as i32;
|
||||
let flatness_base_i32 = SCRATCH_FLATNESS as i32;
|
||||
|
||||
unsafe {
|
||||
stream
|
||||
.launch_builder(&kernel)
|
||||
.arg(&isv_dev)
|
||||
.arg(&q_var_base_i32)
|
||||
.arg(&sigma_base_i32)
|
||||
.arg(&scratch_dev)
|
||||
.arg(&c51_base_i32)
|
||||
.arg(&iqn_base_i32)
|
||||
.arg(&cql_base_i32)
|
||||
.arg(&ens_base_i32)
|
||||
.arg(&flatness_base_i32)
|
||||
.launch(LaunchConfig {
|
||||
grid_dim: (1, 1, 1),
|
||||
block_dim: (4, 1, 1),
|
||||
shared_mem_bytes: 0,
|
||||
})
|
||||
.expect("launch pearl_2_budget_update (sharp regime)");
|
||||
}
|
||||
stream.synchronize().expect("sync after pearl_2_budget_update (sharp regime)");
|
||||
|
||||
let scratch = scratch_buf.read_all();
|
||||
|
||||
// Analytical values for flatness = 0.04 / (4.0 + 1e-8) ≈ 0.01:
|
||||
let expected_sigma_sq = 2.0_f32 * 2.0_f32 + 1e-8_f32;
|
||||
let expected_flatness = (0.04_f32 / expected_sigma_sq).clamp(0.0, 1.0);
|
||||
const BASE_IQN: f32 = 0.11_f32;
|
||||
const BASE_C51_SHARE: f32 = 0.84_f32 / (1.0_f32 - BASE_IQN);
|
||||
let expected_iqn = BASE_IQN + (1.0_f32 - expected_flatness) * (1.0_f32 - BASE_IQN);
|
||||
let expected_c51 = expected_flatness * (1.0_f32 - BASE_IQN) * BASE_C51_SHARE;
|
||||
let expected_cql = 0.0_f32;
|
||||
let expected_ens = (1.0_f32 - expected_iqn - expected_c51 - expected_cql).max(0.0_f32);
|
||||
|
||||
for b in 0..4_usize {
|
||||
let c51 = scratch[SCRATCH_C51 + b];
|
||||
let iqn = scratch[SCRATCH_IQN + b];
|
||||
let cql = scratch[SCRATCH_CQL + b];
|
||||
let ens = scratch[SCRATCH_ENS + b];
|
||||
let flatness = scratch[SCRATCH_FLATNESS + b];
|
||||
|
||||
println!(
|
||||
"branch={b} flatness={flatness:.6} iqn={iqn:.6} c51={c51:.6} \
|
||||
cql={cql:.6} ens={ens:.6}"
|
||||
);
|
||||
|
||||
let rel_flatness = (flatness - expected_flatness).abs() / expected_flatness.max(1e-9);
|
||||
assert!(
|
||||
rel_flatness < 0.01,
|
||||
"branch={b}: flatness={flatness:.6} expected={expected_flatness:.6} rel={rel_flatness:.5}"
|
||||
);
|
||||
let rel_iqn = (iqn - expected_iqn).abs() / expected_iqn.max(1e-9);
|
||||
assert!(
|
||||
rel_iqn < 0.01,
|
||||
"branch={b}: iqn={iqn:.6} expected={expected_iqn:.6} rel={rel_iqn:.5}"
|
||||
);
|
||||
// IQN must dominate: iqn >> c51
|
||||
assert!(
|
||||
iqn > 0.9_f32,
|
||||
"branch={b}: iqn={iqn:.6} should exceed 0.9 in sharp regime"
|
||||
);
|
||||
let rel_c51 = (c51 - expected_c51).abs() / expected_c51.max(1e-9);
|
||||
assert!(
|
||||
rel_c51 < 0.01,
|
||||
"branch={b}: c51={c51:.6} expected={expected_c51:.6} rel={rel_c51:.5}"
|
||||
);
|
||||
assert!(
|
||||
cql.abs() < 1e-7,
|
||||
"branch={b}: cql={cql:.6} expected=0"
|
||||
);
|
||||
// ens is near-zero in sharp regime; just verify it is non-negative
|
||||
assert!(
|
||||
ens >= 0.0_f32,
|
||||
"branch={b}: ens={ens:.6} must be non-negative"
|
||||
);
|
||||
let _ = expected_ens; // expected is ~0.0005, just verify non-negative above
|
||||
}
|
||||
}
|
||||
|
||||
@@ -2,6 +2,8 @@
|
||||
|
||||
**Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7.
|
||||
|
||||
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).
|
||||
|
||||
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.
|
||||
|
||||
Reference in New Issue
Block a user