diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 16a4d2caa..011456050 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -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 { diff --git a/crates/ml/src/cuda_pipeline/pearl_2_budget_kernel.cu b/crates/ml/src/cuda_pipeline/pearl_2_budget_kernel.cu index 8bd2aa894..d00a0ccfe 100644 --- a/crates/ml/src/cuda_pipeline/pearl_2_budget_kernel.cu +++ b/crates/ml/src/cuda_pipeline/pearl_2_budget_kernel.cu @@ -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 @@ -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(); } diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index a669fa78a..a3ca42e77 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -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.