sp7(pearl_2): contract change — drop CQL/C51/ENS args (atomic)
Kernel signature reduced from 9 to 6 args; pearl_2_budget_update now writes only budget_iqn[4] and flatness[4]. Launcher migrated atomically: arg list shrinks, apply_pearls smoothing loop shrinks from 5 slot-blocks to 2. SCRATCH_PEARL_2_C51, SCRATCH_PEARL_2_CQL, SCRATCH_PEARL_2_ENS deleted as orphan constants per feedback_wire_everything_up. The corresponding producer scratch slots (111..115, 119..127) become reserved-for-future inside the buffer. CQL/C51/ENS budget ownership now lives in the SP7 loss-balance controller (loaded in T5; wired in T7). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -1070,10 +1070,9 @@ pub const SP5_WIENER_TOTAL_FLOATS: usize =
|
||||
/// 4 floats for pearl_3_sf output (SCRATCH_PEARL_3_SF=107, sigma_fraction × 4 branches)
|
||||
/// 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
|
||||
/// [111..115) reserved-for-future (T6 retired SCRATCH_PEARL_2_C51)
|
||||
/// budget_iqn[4] → SCRATCH_PEARL_2_IQN=115
|
||||
/// budget_cql[4] → SCRATCH_PEARL_2_CQL=119
|
||||
/// budget_ens[4] → SCRATCH_PEARL_2_ENS=123
|
||||
/// [119..127) reserved-for-future (T6 retired SCRATCH_PEARL_2_CQL/_ENS)
|
||||
/// flatness[4] → SCRATCH_PEARL_2_FLATNESS=127
|
||||
/// Task A4 adds:
|
||||
/// 8 floats for grad_cosine_sim output (SCRATCH_PEARL_4_COSINE=131, cos_sim × 8 groups)
|
||||
@@ -1213,26 +1212,14 @@ 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 A3 + SP7 (2026-05-03): scratch index base for pearl_2_budget
|
||||
/// budget_iqn[4] output block. C51/CQL/ENS scratch slots were retired in
|
||||
/// SP7 — the loss-balance controller owns those budgets directly.
|
||||
/// Producer scratch slots 111..115, 119..127 are now reserved-for-future.
|
||||
const SCRATCH_PEARL_2_IQN: usize = 115; // 115..119
|
||||
/// SP5 Task A3: scratch index base for pearl_2_budget flatness[4] output
|
||||
/// block.
|
||||
const SCRATCH_PEARL_2_FLATNESS: usize = 127; // 127..131
|
||||
|
||||
/// SP5 Task A4: scratch index base for grad_cosine_sim_kernel cosine_sim[8] output block.
|
||||
/// Slots [131..139): per-param-group gradient direction cosine similarity vs previous step.
|
||||
@@ -4406,17 +4393,19 @@ 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.
|
||||
/// SP5 Task A3 + SP7 (2026-05-03): Pearl 2 IQN budget + flatness producer.
|
||||
/// 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).
|
||||
/// and derives (budget_iqn, flatness) per branch. Writes 8 floats to
|
||||
/// `producer_step_scratch_buf[115..119, 127..131)` (IQN + flatness; slots
|
||||
/// 111..115 and 119..127 are reserved-for-future after T6 retirement).
|
||||
/// Chained `apply_pearls_ad_kernel` (8 single-slot launches) then smooths via
|
||||
/// Pearls A+D into ISV[BUDGET_IQN_BASE, FLATNESS_BASE].
|
||||
/// SP7: C51/CQL/ENS budget ownership transferred to loss_balance_controller_kernel.
|
||||
/// Kernel signature: 6 args (isv, q_var_base, sigma_base, scratch, iqn_idx, flatness_idx).
|
||||
/// 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.
|
||||
/// launch_sp5_pearl_3_sigma (NOISY_SIGMA).
|
||||
/// 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,
|
||||
@@ -11229,28 +11218,26 @@ impl GpuDqnTrainer {
|
||||
Ok(())
|
||||
}
|
||||
|
||||
/// SP5 Task A3 (2026-05-01): Pearl 2 — per-branch loss budget from per-branch
|
||||
/// flatness signal.
|
||||
/// SP5 Task A3 + SP7 (2026-05-03): Pearl 2 — IQN budget + flatness producer.
|
||||
///
|
||||
/// 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)
|
||||
/// 1. `pearl_2_budget_update` (1×4 kernel, 6-arg signature) →
|
||||
/// scratch[115..119) (IQN) + scratch[127..131) (flatness).
|
||||
/// 2. `apply_pearls_ad_kernel` × 8 (one per output slot):
|
||||
/// 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.
|
||||
/// SP7: C51/CQL/ENS budget smoothing removed — owned by
|
||||
/// `launch_loss_balance_controller`. Wiener offsets unchanged:
|
||||
/// (SP4_PRODUCER_COUNT=71 + (isv_slot − SP5_SLOT_BASE=174)) × 3.
|
||||
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,
|
||||
BUDGET_IQN_BASE, FLATNESS_BASE,
|
||||
NOISY_SIGMA_BASE, Q_VAR_PER_BRANCH_BASE,
|
||||
};
|
||||
|
||||
debug_assert!(self.isv_signals_dev_ptr != 0,
|
||||
@@ -11260,17 +11247,14 @@ impl GpuDqnTrainer {
|
||||
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)
|
||||
// SP7 (2026-05-03): kernel signature reduced to (isv, q_var, sigma,
|
||||
// scratch, iqn_idx, flatness_idx). The c51/cql/ens scratch slots
|
||||
// and apply_pearls smoothing for those slots are owned by the SP7
|
||||
// loss-balance controller; Pearl 2 only writes IQN + flatness now.
|
||||
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
|
||||
let iqn_base_i32 = SCRATCH_PEARL_2_IQN as i32;
|
||||
let flatness_base_i32 = SCRATCH_PEARL_2_FLATNESS as i32;
|
||||
|
||||
unsafe {
|
||||
self.stream
|
||||
@@ -11279,10 +11263,7 @@ impl GpuDqnTrainer {
|
||||
.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),
|
||||
@@ -11292,21 +11273,13 @@ impl GpuDqnTrainer {
|
||||
.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
|
||||
// SP7: smoothing chain reduced from 5 slot-blocks to 2 (IQN +
|
||||
// flatness). The SP7 controller owns the BUDGET_{C51,CQL,ENS}_BASE
|
||||
// smoothing chain.
|
||||
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 {
|
||||
|
||||
@@ -1,30 +1,23 @@
|
||||
// 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.
|
||||
// SP5 Task A3 + SP7 (2026-05-03): Pearl 2 IQN budget + flatness producer.
|
||||
//
|
||||
// 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.
|
||||
// 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). Writes:
|
||||
//
|
||||
// flatness[b] = clamp(var_q[b] / (σ[b]² + EPS_DIV), 0, 1)
|
||||
// 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.
|
||||
// SP7 takeover: BUDGET_C51_BASE, BUDGET_CQL_BASE, BUDGET_ENS_BASE are
|
||||
// now driven by `loss_balance_controller_kernel.cu` (outcome-driven
|
||||
// ratio controller using flatness as the target modulator). Pearl 2's
|
||||
// kernel signature is reduced to drop the c51/cql/ens scratch-index
|
||||
// arguments — the launcher in gpu_dqn_trainer.rs is updated atomically.
|
||||
//
|
||||
// 4-thread single-block (one thread per branch). No atomicAdd.
|
||||
// __threadfence_system() after writes.
|
||||
// 4-thread single-block (one thread per branch). No atomicAdd
|
||||
// (feedback_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).
|
||||
// BASE_IQN=0.11 and EPS_DIV=1e-8 are Invariant 1 structural anchors.
|
||||
|
||||
#include <cuda_runtime.h>
|
||||
|
||||
@@ -33,37 +26,24 @@ extern "C" __global__ void pearl_2_budget_update(
|
||||
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;
|
||||
const float BASE_IQN = 0.11f;
|
||||
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 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();
|
||||
}
|
||||
|
||||
@@ -3904,6 +3904,12 @@ extended. No producer kernel yet — that arrives in the next commit.
|
||||
T5 polish landed (commit ⟨pending⟩) — `cql_dev` renamed to `cql_sx_dev`
|
||||
for clarity (the launcher uses the cql_sx post-budget delta) + `// 213`
|
||||
annotation on `base_wiener_offset` for sibling-launcher consistency.
|
||||
- T6: Pearl 2 contract change (drop CQL/C51/ENS args).
|
||||
- T6 (commit ⟨pending⟩): Pearl 2 contract change. Kernel signature
|
||||
reduced from 9-arg to 6-arg (drops c51/cql/ens scratch idx args).
|
||||
Launcher migrated atomically — kernel-arg list shrinks from 9 to 6,
|
||||
apply_pearls smoothing loop shrinks from 5 slot-blocks to 2 (IQN +
|
||||
flatness). SCRATCH_PEARL_2_C51/_CQL/_ENS deleted as orphans per
|
||||
feedback_wire_everything_up. CQL/C51/ENS budget ownership now lives
|
||||
entirely in the SP7 controller.
|
||||
- T7: launch site + sentinel-aware bootstrap with bootstrap constants matching the kernel's cold-start basis (defined in T7) + stale doc.
|
||||
- T9–T10: smoke + 50-epoch verification.
|
||||
|
||||
Reference in New Issue
Block a user