From cb6a89714bc24b45bacf74587ac89940fd12b35f Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sun, 3 May 2026 11:34:21 +0200 Subject: [PATCH 1/5] =?UTF-8?q?sp7(controller):=20activation=20flag=20+=20?= =?UTF-8?q?Welford-=CE=B1=20to=20escape=20bootstrap=20trap?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Fixes the SP7 controller dormancy discovered in smoke-test-8556k: per-branch budgets stuck at exactly bootstrap constants because the kernel's cold_start_basis numerically equaled the consumer's bootstrap fallback, AND the Wiener-α was clamped at ALPHA_FLOOR=1e-4 from step 1. Architectural change: - 8 new ISV slots LB_{CQL,C51}_ACTIVE_BASE per (head × branch). Activation flag is monotonic per fold, FoldReset on boundary. - Kernel only sets active=1 when grads are populated AND on subsequent active-state computation; cold-start branch leaves active=0 (fresh branch) or holds prior budget steady (transient grad gate). - Consumer dispatches on activation: bootstrap when active<0.5, controller verbatim when active>=0.5. No more spurious bootstrap when controller writes legitimate small values. - Welford-α hybrid (max of 1/max(1,epoch_idx_in_fold) and Wiener-α) gives full update on first active step, falls off as 1/N until Wiener takes over with meaningful variance estimates. EPOCH_IDX_INDEX is the existing per-fold-reset counter (no new tuned constants). State reset registry: 2 new sp7_lb_*_active FoldReset entries + matching dispatch arms in reset_named_state. Contract test (every_fold_and_soft_reset_entry_has_dispatch_arm) gates compile. GPU unit test sp7_loss_balance_controller_activation_flag_transitions exercises 3 transitions (cold start → both flags 0; active → both flags 1 with controller-computed budget != bootstrap; transient grad- gate → flags hold at 1, prior budget held verbatim). Passes on local RTX 3050 Ti. Audit doc: Fix 31 sub-bullet describing the activation-flag fix. Memory pearl out-of-tree (controller will dispatch separately). Touched: crates/ml/src/cuda_pipeline/sp5_isv_slots.rs crates/ml/src/cuda_pipeline/loss_balance_controller_kernel.cu crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs crates/ml/src/trainers/dqn/fused_training.rs crates/ml/src/trainers/dqn/state_reset_registry.rs crates/ml/src/trainers/dqn/trainer/training_loop.rs crates/ml/tests/sp5_producer_unit_tests.rs docs/dqn-wire-up-audit.md Cargo check workspace clean. Cargo test ml --lib clean (incl. contract test + 6 sp5_isv_slots tests). 16 pre-existing failures unchanged. Co-Authored-By: Claude Opus 4.7 (1M context) --- .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 33 ++- .../loss_balance_controller_kernel.cu | 136 ++++++--- crates/ml/src/cuda_pipeline/sp5_isv_slots.rs | 99 +++++-- crates/ml/src/trainers/dqn/fused_training.rs | 83 ++++-- .../src/trainers/dqn/state_reset_registry.rs | 19 ++ .../src/trainers/dqn/trainer/training_loop.rs | 19 ++ crates/ml/tests/sp5_producer_unit_tests.rs | 276 ++++++++++++++++++ docs/dqn-wire-up-audit.md | 60 ++++ 8 files changed, 632 insertions(+), 93 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 7e4830635..ed4f7bffc 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -1100,8 +1100,11 @@ pub const SP5_WIENER_TOTAL_FLOATS: usize = /// 4 floats for LB sample_var_cql[4] (SCRATCH_LB_SAMPLE_VAR_CQL=230, at [230..234)) /// 4 floats for LB diff_var_c51[4] (SCRATCH_LB_DIFF_VAR_C51=234, at [234..238)) /// 4 floats for LB sample_var_c51[4] (SCRATCH_LB_SAMPLE_VAR_C51=238, at [238..242)) -/// Combined: 71 + 16 + 16 + 4 + 4 + 20 + 8 + 8 + 8 + 8 + 8 + 4 + 4 + 20 + 4 + 4 + 4 + 4 + 3 + 24 = 242 scratch slots [0..242). -pub const SP5_SCRATCH_TOTAL: usize = 242; +/// SP7 activation-flag fix (2026-05-03) adds: +/// 4 floats for LB cql_active[4] (SCRATCH_LB_ACTIVE_CQL=242, at [242..246)) +/// 4 floats for LB c51_active[4] (SCRATCH_LB_ACTIVE_C51=246, at [246..250)) +/// Combined: 71 + 16 + 16 + 4 + 4 + 20 + 8 + 8 + 8 + 8 + 8 + 4 + 4 + 20 + 4 + 4 + 4 + 4 + 3 + 24 + 8 = 250 scratch slots [0..250). +pub const SP5_SCRATCH_TOTAL: usize = 250; /// SP5 Layer D Task D1 (rewrite, 2026-05-02): scratch index base for /// `pnl_aggregation_update`. @@ -1171,6 +1174,14 @@ pub const SCRATCH_LB_SAMPLE_VAR_CQL: usize = SCRATCH_LB_BUDGET_CQL + 12; // 230. pub const SCRATCH_LB_DIFF_VAR_C51: usize = SCRATCH_LB_BUDGET_CQL + 16; // 234..238 /// SP7: sample_var_c51[4] → ISV[LB_SAMPLE_VAR_C51_BASE..+4). pub const SCRATCH_LB_SAMPLE_VAR_C51: usize = SCRATCH_LB_BUDGET_CQL + 20; // 238..242 +/// SP7 activation-flag fix (2026-05-03): cql_active[4] → ISV[LB_CQL_ACTIVE_BASE..+4). +/// Kernel writes 1.0 in the active path and when prior was_active=1, 0.0 in the +/// genuine cold-start branch. apply_pearls_ad_kernel's `step_obs == 0.0` short- +/// circuit means writing 0.0 is a no-op (slot stays at its current value); +/// the activation flag is therefore monotonic per fold. +pub const SCRATCH_LB_ACTIVE_CQL: usize = SCRATCH_LB_BUDGET_CQL + 24; // 242..246 +/// SP7 activation-flag fix: c51_active[4] → ISV[LB_C51_ACTIVE_BASE..+4). +pub const SCRATCH_LB_ACTIVE_C51: usize = SCRATCH_LB_BUDGET_CQL + 28; // 246..250 /// SP5 Task A7: scratch index base for pearl_8_trail_update trail_dist[4] output block. /// Slots [199..203): per-direction trail-stop distance (Short=0, Hold=1, Long=2, Flat=3). @@ -11318,6 +11329,7 @@ impl GpuDqnTrainer { BUDGET_CQL_BASE, BUDGET_C51_BASE, FLATNESS_BASE, LB_DIFF_VAR_CQL_BASE, LB_SAMPLE_VAR_CQL_BASE, LB_DIFF_VAR_C51_BASE, LB_SAMPLE_VAR_C51_BASE, + LB_CQL_ACTIVE_BASE, LB_C51_ACTIVE_BASE, }; debug_assert!(self.isv_signals_dev_ptr != 0, @@ -11347,12 +11359,17 @@ impl GpuDqnTrainer { let sample_var_cql_isv_base_i32 = LB_SAMPLE_VAR_CQL_BASE as i32; let diff_var_c51_isv_base_i32 = LB_DIFF_VAR_C51_BASE as i32; let sample_var_c51_isv_base_i32 = LB_SAMPLE_VAR_C51_BASE as i32; + let active_cql_isv_base_i32 = LB_CQL_ACTIVE_BASE as i32; + let active_c51_isv_base_i32 = LB_C51_ACTIVE_BASE as i32; + let epoch_idx_isv_index_i32 = EPOCH_IDX_INDEX as i32; let sb_budget_cql_i32 = SCRATCH_LB_BUDGET_CQL as i32; let sb_budget_c51_i32 = SCRATCH_LB_BUDGET_C51 as i32; let sb_diff_var_cql_i32 = SCRATCH_LB_DIFF_VAR_CQL as i32; let sb_sample_var_cql_i32 = SCRATCH_LB_SAMPLE_VAR_CQL as i32; let sb_diff_var_c51_i32 = SCRATCH_LB_DIFF_VAR_C51 as i32; let sb_sample_var_c51_i32 = SCRATCH_LB_SAMPLE_VAR_C51 as i32; + let sb_active_cql_i32 = SCRATCH_LB_ACTIVE_CQL as i32; + let sb_active_c51_i32 = SCRATCH_LB_ACTIVE_C51 as i32; unsafe { self.stream @@ -11368,6 +11385,9 @@ impl GpuDqnTrainer { .arg(&sample_var_cql_isv_base_i32) .arg(&diff_var_c51_isv_base_i32) .arg(&sample_var_c51_isv_base_i32) + .arg(&active_cql_isv_base_i32) + .arg(&active_c51_isv_base_i32) + .arg(&epoch_idx_isv_index_i32) .arg(&scratch_dev) .arg(&sb_budget_cql_i32) .arg(&sb_budget_c51_i32) @@ -11375,6 +11395,8 @@ impl GpuDqnTrainer { .arg(&sb_sample_var_cql_i32) .arg(&sb_diff_var_c51_i32) .arg(&sb_sample_var_c51_i32) + .arg(&sb_active_cql_i32) + .arg(&sb_active_c51_i32) .launch(LaunchConfig { grid_dim: (1, 1, 1), block_dim: (8, 1, 1), @@ -11383,8 +11405,9 @@ impl GpuDqnTrainer { .map_err(|e| MLError::ModelError(format!("loss_balance_controller_update: {e}")))?; } - // Step 2: apply_pearls_ad_kernel × 24 — one per ISV output slot. - // Wiener offset: (SP4_PRODUCER_COUNT + (isv_slot - SP5_SLOT_BASE)) * 3. + // Step 2: apply_pearls_ad_kernel × 32 — one per ISV output slot + // (8 slot-blocks × 4 branches). Wiener offset: + // (SP4_PRODUCER_COUNT + (isv_slot - SP5_SLOT_BASE)) * 3. let base_wiener_offset = SP4_PRODUCER_COUNT as i32 * 3; // 213 for (isv_base, scratch_base) in [ @@ -11394,6 +11417,8 @@ impl GpuDqnTrainer { (LB_SAMPLE_VAR_CQL_BASE, SCRATCH_LB_SAMPLE_VAR_CQL), (LB_DIFF_VAR_C51_BASE, SCRATCH_LB_DIFF_VAR_C51), (LB_SAMPLE_VAR_C51_BASE, SCRATCH_LB_SAMPLE_VAR_C51), + (LB_CQL_ACTIVE_BASE, SCRATCH_LB_ACTIVE_CQL), + (LB_C51_ACTIVE_BASE, SCRATCH_LB_ACTIVE_C51), ] { for b in 0..4_usize { let scratch_idx = (scratch_base + b) as i32; diff --git a/crates/ml/src/cuda_pipeline/loss_balance_controller_kernel.cu b/crates/ml/src/cuda_pipeline/loss_balance_controller_kernel.cu index ffd65b8ee..c49ac7a15 100644 --- a/crates/ml/src/cuda_pipeline/loss_balance_controller_kernel.cu +++ b/crates/ml/src/cuda_pipeline/loss_balance_controller_kernel.cu @@ -20,15 +20,38 @@ // // diff = candidate - old_budget // sample_sq = h_norm[slice]^2 -// alpha_eff = old_diff_var / (old_diff_var + old_sample_var + EPS_DIV) +// wiener_α = old_diff_var / (old_diff_var + old_sample_var + EPS_DIV) +// welford_α = 1 / max(1, epoch_idx_in_fold) (full update on first +// active step, falls off +// as 1/N within the fold) +// alpha_eff = max(welford_α, wiener_α) (hybrid: Welford bounds α +// from below early; Wiener +// takes over once accumulated +// variance is meaningful) // alpha_eff = clamp(alpha_eff, ALPHA_FLOOR, ALPHA_CEIL) // new_budget = clamp(old_budget + alpha_eff * diff, EPS_DIV, MAX_BUDGET) // -// Cold-start (per spec section "Cold start"): +// Cold-start (per "Activation flag" spec, 2026-05-03): // if h_norm[slice] < EPS_DIV || iqn_norm[slice] < EPS_DIV: -// new_budget = cold_start_basis (matches consumer bootstrap) -// else if old_budget < EPS_DIV: -// new_budget = clamp(cold_start_basis * correction, EPS_DIV, MAX_BUDGET) +// if prior was_active >= 0.5: +// # Transient grad gate-off; hold prior budget steady. +// new_budget = isv[budget_isv_base + branch] +// scratch_active = 1.0 (still active) +// else: +// # Genuine cold start; let consumer's bootstrap handle this step. +// # Don't touch scratch_budget/diff_var/sample_var (apply_pearls_ad +// # short-circuits on step_obs == 0.0). +// scratch_active = 0.0 (no-op via short-circuit) +// return +// +// SP7 activation-flag fix (2026-05-03): the controller writes 1.0 into +// `scratch_active` when the active path computes a real budget update OR +// when a transient cold-start hits a previously-active branch (prior +// budget held verbatim, controller stays "engaged"). It writes 0.0 in +// the genuine cold-start branch (was_active = 0). apply_pearls_ad_kernel +// short-circuits step_obs == 0.0, so writing 0 leaves the activation +// slot at its current value; activation is therefore monotonic per fold. +// FoldReset zeroes the slots so the new fold re-bootstraps cleanly. // // Single block, 8 threads (2 heads × 4 branches). No atomicAdd (feedback_no_atomicadd). // __threadfence_system() after writes. @@ -54,6 +77,9 @@ extern "C" __global__ void loss_balance_controller_update( int sample_var_cql_isv_base, // LB_SAMPLE_VAR_CQL_BASE = 301 int diff_var_c51_isv_base, // LB_DIFF_VAR_C51_BASE = 305 int sample_var_c51_isv_base, // LB_SAMPLE_VAR_C51_BASE = 309 + int active_cql_isv_base, // LB_CQL_ACTIVE_BASE = 313 + int active_c51_isv_base, // LB_C51_ACTIVE_BASE = 317 + int epoch_idx_isv_index, // EPOCH_IDX_INDEX = 39 (per-fold counter) // Producer scratch buffer (kernel writes here; apply_pearls_ad_kernel // smooths these into ISV downstream). float* __restrict__ scratch_out, @@ -62,7 +88,9 @@ extern "C" __global__ void loss_balance_controller_update( int scratch_diff_var_cql_base, int scratch_sample_var_cql_base, int scratch_diff_var_c51_base, - int scratch_sample_var_c51_base + int scratch_sample_var_c51_base, + int scratch_active_cql_base, + int scratch_active_c51_base ) { // ── Invariant 1 anchors ────────────────────────────────────────── // Modulated per-branch by flatness (spec section "Math"). @@ -75,8 +103,13 @@ extern "C" __global__ void loss_balance_controller_update( const float MAX_BUDGET = 1.0f; // Cold-start basis — must match consumer-side bootstrap in // fused_training.rs (CQL_BOOTSTRAP_BUDGET, C51_BOOTSTRAP_BUDGET). + // Used only for the seed-from-cold-start branch when was_active = 0 + // transitions to "first active step" (norms populated, old_budget = 0). const float COLD_START_FLOOR_CQL = 0.02f; const float COLD_START_FLOOR_C51 = 0.05f; + // Activation discriminator. Binary threshold (not a tuned knob); slot + // values are written as 0.0 or 1.0 and smoothed monotonically. + const float ACTIVE_THRESHOLD = 0.5f; // ── Thread layout: head × branch ───────────────────────────────── int tid = threadIdx.x; @@ -104,62 +137,93 @@ extern "C" __global__ void loss_balance_controller_update( ? ANCHOR_CQL_RATIO * (1.0f - flat_b) : ANCHOR_C51_RATIO * flat_b; - // ── Read prior budget + per-branch Wiener state ────────────────── + // ── Read prior budget + per-branch Wiener state + activation flag ─ int budget_isv_base = (head == 0) ? budget_cql_isv_base : budget_c51_isv_base; int diff_var_isv_base = (head == 0) ? diff_var_cql_isv_base : diff_var_c51_isv_base; int sample_var_isv_base = (head == 0) ? sample_var_cql_isv_base : sample_var_c51_isv_base; + int active_isv_base = (head == 0) ? active_cql_isv_base : active_c51_isv_base; int scratch_budget_base = (head == 0) ? scratch_budget_cql_base : scratch_budget_c51_base; int scratch_diff_base = (head == 0) ? scratch_diff_var_cql_base : scratch_diff_var_c51_base; int scratch_sample_base = (head == 0) ? scratch_sample_var_cql_base : scratch_sample_var_c51_base; + int scratch_active_base = (head == 0) ? scratch_active_cql_base : scratch_active_c51_base; float cold_start_basis = (head == 0) ? COLD_START_FLOOR_CQL : COLD_START_FLOOR_C51; float old_budget = isv_signals[budget_isv_base + branch]; float old_diff_var = isv_signals[diff_var_isv_base + branch]; float old_sample_var = isv_signals[sample_var_isv_base + branch]; + float was_active = isv_signals[active_isv_base + branch]; // ── Branch on cold-start regime ────────────────────────────────── + if (h_n < EPS_DIV || iqn_n < EPS_DIV) { + // Warmup hasn't fired or the loss head is gated off. + if (was_active >= ACTIVE_THRESHOLD) { + // Transient grad gate-off on a previously-active branch. + // Hold prior budget verbatim and re-assert active=1; don't + // drift the budget on what is purely a measurement gap. + scratch_out[scratch_budget_base + branch] = old_budget; + scratch_out[scratch_active_base + branch] = 1.0f; + // Do NOT update diff_var/sample_var: there's no new sample. + // (Writing 0 to those slots is a no-op via apply_pearls_ad's + // step_obs == 0.0 short-circuit; we must still write _something_ + // per the launcher loop, but 0 keeps Wiener state untouched.) + scratch_out[scratch_diff_base + branch] = 0.0f; + scratch_out[scratch_sample_base + branch] = 0.0f; + } else { + // Genuine cold start. Don't write a budget (apply_pearls_ad + // short-circuits on 0.0 → leaves slot untouched → consumer's + // bootstrap handles it). active stays 0. + scratch_out[scratch_budget_base + branch] = 0.0f; + scratch_out[scratch_active_base + branch] = 0.0f; + scratch_out[scratch_diff_base + branch] = 0.0f; + scratch_out[scratch_sample_base + branch] = 0.0f; + } + __threadfence_system(); + return; + } + + // ── Active path: norms populated, controller engages ───────────── + float actual_ratio = h_n / iqn_n; + float correction = target_ratio / fmaxf(actual_ratio, EPS_DIV); + float new_budget; float diff; float sample_sq = h_n * h_n; - if (h_n < EPS_DIV || iqn_n < EPS_DIV) { - // Warmup hasn't fired or the loss head is gated off. Seed budget - // at the consumer-side bootstrap value (effectively a no-op vs - // the existing behavior on the very first step). - new_budget = cold_start_basis; - diff = 0.0f; + if (old_budget < EPS_DIV) { + // Sentinel-0 read (cold start or fold boundary). The norm we + // observed was produced under cql/c51_budget = cold_start_basis + // (consumer-side bootstrap). Seed the controller at the value + // that achieves target_ratio next step. + float seed = cold_start_basis * correction; + new_budget = fminf(MAX_BUDGET, fmaxf(EPS_DIV, seed)); + diff = new_budget - cold_start_basis; } else { - float actual_ratio = h_n / iqn_n; - float correction = target_ratio / fmaxf(actual_ratio, EPS_DIV); + float candidate = old_budget * correction; + diff = candidate - old_budget; - if (old_budget < EPS_DIV) { - // Sentinel-0 read (cold start or fold boundary). The norm - // we observed was produced under cql/c51_budget = cold_start_basis - // (consumer-side bootstrap). Seed the controller at the value - // that achieves target_ratio next step. - float seed = cold_start_basis * correction; - new_budget = fminf(MAX_BUDGET, fmaxf(EPS_DIV, seed)); - diff = new_budget - cold_start_basis; - } else { - float candidate = old_budget * correction; - diff = candidate - old_budget; + // Welford-α: bounds α from below as 1/N within the fold, giving + // a full update on the first active step (epoch_idx = 0 → 1.0) + // and falling off as the fold progresses. Wiener-α takes over + // once accumulated variance is meaningful. + float epoch_idx = isv_signals[epoch_idx_isv_index]; + float welford_alpha = 1.0f / fmaxf(1.0f, epoch_idx); - // Wiener α from prior state (no in-step state mutation — - // we update the EMAs via scratch_out + apply_pearls_ad). - float wiener_num = old_diff_var; - float wiener_den = old_diff_var + old_sample_var + EPS_DIV; - float alpha_eff = wiener_num / wiener_den; - alpha_eff = fminf(ALPHA_CEIL, fmaxf(ALPHA_FLOOR, alpha_eff)); + float wiener_num = old_diff_var; + float wiener_den = old_diff_var + old_sample_var + EPS_DIV; + float wiener_alpha = wiener_num / wiener_den; - new_budget = old_budget + alpha_eff * diff; - new_budget = fminf(MAX_BUDGET, fmaxf(EPS_DIV, new_budget)); - } + float alpha_eff = fmaxf(welford_alpha, wiener_alpha); + alpha_eff = fminf(ALPHA_CEIL, fmaxf(ALPHA_FLOOR, alpha_eff)); + + new_budget = old_budget + alpha_eff * diff; + new_budget = fminf(MAX_BUDGET, fmaxf(EPS_DIV, new_budget)); } - // ── Write the three outputs to scratch (apply_pearls_ad smooths) ── + // ── Write the four outputs to scratch (apply_pearls_ad smooths) ─── scratch_out[scratch_budget_base + branch] = new_budget; scratch_out[scratch_diff_base + branch] = diff * diff; // raw observation scratch_out[scratch_sample_base + branch] = sample_sq; // raw observation + scratch_out[scratch_active_base + branch] = 1.0f; // controller active __threadfence_system(); } diff --git a/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs index 6e1693723..178b4c511 100644 --- a/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs +++ b/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs @@ -13,8 +13,10 @@ //! 286..290 Layer D D1: PnL aggregation outputs (4 slots, fold-reset) //! 290..294 Layer D D2: Health composition outputs (4 slots, fold-reset) //! 294..297 Layer D D3: Training metrics EMA outputs (3 slots, fold-reset) +//! 297..313 SP7 T1 loss-budget Wiener stats (16 slots, fold-reset) +//! 313..321 SP7 activation-flag fix per-(head, branch) (8 slots, fold-reset) //! -//! Total: 121 new SP5 ISV slots (52 + 24 + 20 + 4 + 4 + 6 + 4 + 4 + 3). +//! Total: 145 new SP5 ISV slots (52 + 24 + 20 + 4 + 4 + 6 + 4 + 4 + 3 + 16 + 8). pub const SP5_SLOT_BASE: usize = 174; @@ -180,7 +182,22 @@ pub const LB_SAMPLE_VAR_CQL_BASE: usize = 301; // [4] (cql_norm)² EMA pub const LB_DIFF_VAR_C51_BASE: usize = 305; // [4] (candidate_c51 − old_c51)² EMA pub const LB_SAMPLE_VAR_C51_BASE: usize = 309; // [4] (c51_norm)² EMA -pub const SP5_SLOT_END: usize = 313; +// ── SP7 activation-flag fix (2026-05-03): per-(head, branch) activation slots ─ +// +// 8 new ISV slots — 4 per managed head (CQL / C51), one per branch. +// The controller writes 1.0 to scratch_active when its active path +// computes a real budget update; 0.0 (a no-op via apply_pearls_ad's +// `step_obs == 0.0` short-circuit) when the genuine cold-start branch +// fires. The consumer dispatches on this flag instead of the prior +// numeric-equality bootstrap (CQL_BOOTSTRAP_BUDGET=0.02 was identical +// to the kernel's COLD_START_FLOOR_CQL — making the two states +// indistinguishable). Activation is monotonic per fold; fold-reset +// zeroes the slots so the new fold re-activates from scratch as +// gradients populate. +pub const LB_CQL_ACTIVE_BASE: usize = 313; // [4] activation flag for CQL controller +pub const LB_C51_ACTIVE_BASE: usize = 317; // [4] activation flag for C51 controller + +pub const SP5_SLOT_END: usize = 321; /// Wiener-buffer producer-count constant. Sizes `wiener_state_buf` via /// `(SP4_PRODUCER_COUNT + SP5_PRODUCER_COUNT) * SP4_WIENER_FLOATS_PER_SLOT`. @@ -221,13 +238,20 @@ pub const SP5_SLOT_END: usize = 313; /// at ISV[297..313). Unique-slot count grows 121 → 137; the linear span /// (and `SP5_PRODUCER_COUNT`) grows 123 → 139. Pearl 6 carve-out at /// [280..286) remains unchanged. -pub const SP5_PRODUCER_COUNT: usize = 139; -// linear span = SP5_SLOT_END - SP5_SLOT_BASE = 313 - 174 = 139 wiener triples -// unique-slot count = 137 (52 per-branch + 24 Adam + 20 IQN τ + 4 trail +/// +/// SP7 activation-flag fix (2026-05-03): allocates 8 new SP7 ISV slots +/// (`LB_{CQL,C51}_ACTIVE_BASE`, 4 each) at ISV[313..321). Unique-slot +/// count grows 137 → 145; the linear span (and `SP5_PRODUCER_COUNT`) +/// grows 139 → 147. Activation flags are monotonic per fold (FoldReset +/// at boundary), so the new fold re-bootstraps cleanly through Pearl A. +pub const SP5_PRODUCER_COUNT: usize = 147; +// linear span = SP5_SLOT_END - SP5_SLOT_BASE = 321 - 174 = 147 wiener triples +// unique-slot count = 145 (52 per-branch + 24 Adam + 20 IQN τ + 4 trail // + 4 num_atoms + 6 Kelly + 4 Layer D D1 PnL aggregation // + 4 Layer D D2 health composition // + 3 Layer D D3 training metrics EMA -// + 16 SP7 T1 loss-budget Wiener stats) +// + 16 SP7 T1 loss-budget Wiener stats +// + 8 SP7 activation-flag fix per-(head,branch) flags) // ── Convenience accessors ──────────────────────────────────────────── #[inline] pub const fn atom_v_center(b: usize) -> usize { ATOM_V_CENTER_BASE + b } @@ -253,6 +277,8 @@ pub const SP5_PRODUCER_COUNT: usize = 139; #[inline] pub const fn lb_sample_var_cql(b: usize) -> usize { LB_SAMPLE_VAR_CQL_BASE + b } #[inline] pub const fn lb_diff_var_c51(b: usize) -> usize { LB_DIFF_VAR_C51_BASE + b } #[inline] pub const fn lb_sample_var_c51(b: usize) -> usize { LB_SAMPLE_VAR_C51_BASE + b } +#[inline] pub const fn lb_cql_active(b: usize) -> usize { LB_CQL_ACTIVE_BASE + b } +#[inline] pub const fn lb_c51_active(b: usize) -> usize { LB_C51_ACTIVE_BASE + b } /// Layout fingerprint contribution. Appended to the existing /// LAYOUT_FINGERPRINT_SEED in gpu_dqn_trainer.rs so existing checkpoints @@ -270,7 +296,8 @@ pub const SP5_LAYOUT_FINGERPRINT_FRAGMENT: &str = TRAINING_SHARPE_EMA=294;MAX_DD_EMA=295;LOW_DD_RATIO=296;\ LB_DIFF_VAR_CQL_BASE=297;LB_SAMPLE_VAR_CQL_BASE=301;\ LB_DIFF_VAR_C51_BASE=305;LB_SAMPLE_VAR_C51_BASE=309;\ - ISV_TOTAL_DIM=313"; + LB_CQL_ACTIVE_BASE=313;LB_C51_ACTIVE_BASE=317;\ + ISV_TOTAL_DIM=321"; #[cfg(test)] mod tests { @@ -354,27 +381,33 @@ mod tests { slots.insert(lb_diff_var_c51(b)); slots.insert(lb_sample_var_c51(b)); } + // SP7 activation-flag fix: per-(head, branch) flags (b in 0..4) + for b in 0..4 { + slots.insert(lb_cql_active(b)); + slots.insert(lb_c51_active(b)); + } - // 1. Exactly 137 unique slots. - assert_eq!(slots.len(), 137, "expected 137 unique slots, got {}", slots.len()); + // 1. Exactly 145 unique slots. + assert_eq!(slots.len(), 145, "expected 145 unique slots, got {}", slots.len()); // 2. Min slot is SP5_SLOT_BASE = 174. assert_eq!(*slots.iter().min().unwrap(), 174); - // 3. Max slot is SP5_SLOT_END - 1 = 312. - assert_eq!(*slots.iter().max().unwrap(), 312); + // 3. Max slot is SP5_SLOT_END - 1 = 320. + assert_eq!(*slots.iter().max().unwrap(), 320); // 4. Intentional carve-out gap (278, 279) is absent. assert!(!slots.contains(&278), "slot 278 must be absent (carve-out gap)"); assert!(!slots.contains(&279), "slot 279 must be absent (carve-out gap)"); - // 5. Set equals {174..278} ∪ {280..313} — no holes (other than the + // 5. Set equals {174..278} ∪ {280..321} — no holes (other than the // carve-out gap 278..280), no overlaps. Layer D D1 extended the // upper end 286 → 290; D2 extended it 290 → 294; D3 extends it - // 294 → 297; SP7 extends it 297 → 313. - let expected: HashSet = (174..278).chain(280..313).collect(); + // 294 → 297; SP7 T1 extended it 297 → 313; SP7 activation-flag + // fix extends it 313 → 321. + let expected: HashSet = (174..278).chain(280..321).collect(); assert_eq!(slots, expected, - "slot set does not match expected {{174..278}} ∪ {{280..313}}"); + "slot set does not match expected {{174..278}} ∪ {{280..321}}"); } #[test] @@ -400,12 +433,12 @@ mod tests { // Strictly above the Kelly block (Layer D layout is per-fold; Kelly // is cross-fold-persistent — the contracts must remain disjoint). assert!(PNL_TOTAL_INDEX > LOSS_RATE_SMOOTH_INDEX); - // SP5_SLOT_END must reflect the post-D3 end-of-block. - assert_eq!(SP5_SLOT_END, 313); + // SP5_SLOT_END must reflect the post-SP7-activation-flag end-of-block. + assert_eq!(SP5_SLOT_END, 321); // SP5_PRODUCER_COUNT is the wiener-buffer linear span (slot-range // width including the 2-slot carve-out gap), NOT the unique-slot // count. See SP5_PRODUCER_COUNT docstring for the rationale. - assert_eq!(SP5_PRODUCER_COUNT, 139); + assert_eq!(SP5_PRODUCER_COUNT, 147); assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE); } @@ -423,8 +456,8 @@ mod tests { assert!(HEALTH_SCORE_INDEX > PNL_MAX_DD_INDEX); // 4-slot block is internally contiguous. assert_eq!(GRAD_NORM_NORM_INDEX - HEALTH_SCORE_INDEX, 3); - // SP5_SLOT_END must reflect the post-D3 end-of-block. - assert_eq!(SP5_SLOT_END, 313); + // SP5_SLOT_END must reflect the post-SP7-activation-flag end-of-block. + assert_eq!(SP5_SLOT_END, 321); // SP5_PRODUCER_COUNT linear-span check matches the new end. assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE); } @@ -444,13 +477,33 @@ mod tests { // 3-slot block is internally contiguous (no gaps). assert_eq!(MAX_DD_EMA_INDEX - TRAINING_SHARPE_EMA_INDEX, 1); assert_eq!(LOW_DD_RATIO_INDEX - MAX_DD_EMA_INDEX, 1); - // SP5_SLOT_END must reflect the new end-of-block (3-slot grow). - assert_eq!(SP5_SLOT_END, 313); + // SP5_SLOT_END must reflect the post-SP7-activation-flag end-of-block. + assert_eq!(SP5_SLOT_END, 321); // SP5_PRODUCER_COUNT linear-span check matches the new end. The // wiener buffer must cover the entire linear span — including the // 2-slot carve-out gap (278..280) and the Pearl 6 reserved-but- // unused 6-float block (slots 280..286 don't call apply_pearls). - assert_eq!(SP5_PRODUCER_COUNT, 139); + assert_eq!(SP5_PRODUCER_COUNT, 147); + assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE); + } + + #[test] + fn sp7_activation_flag_slots_contiguous_and_above_lb_block() { + // SP7 activation-flag fix (2026-05-03): 8 new ISV slots — + // 4 per managed head — at ISV[313..321) immediately after the + // SP7 T1 LB Wiener-stat block. + assert_eq!(LB_CQL_ACTIVE_BASE, 313); + assert_eq!(LB_C51_ACTIVE_BASE, 317); + // Strictly above the SP7 T1 LB Wiener-stat block (LB_*_VAR_C51_BASE + // ends at 313). + assert!(LB_CQL_ACTIVE_BASE > LB_SAMPLE_VAR_C51_BASE); + // Both 4-slot blocks internally contiguous. + for b in 0..4 { + assert_eq!(lb_cql_active(b), LB_CQL_ACTIVE_BASE + b); + assert_eq!(lb_c51_active(b), LB_C51_ACTIVE_BASE + b); + } + // SP5_SLOT_END / SP5_PRODUCER_COUNT cover the new range. + assert_eq!(SP5_SLOT_END, 321); assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE); } } diff --git a/crates/ml/src/trainers/dqn/fused_training.rs b/crates/ml/src/trainers/dqn/fused_training.rs index bce6d42c9..2aef0d4a9 100644 --- a/crates/ml/src/trainers/dqn/fused_training.rs +++ b/crates/ml/src/trainers/dqn/fused_training.rs @@ -41,6 +41,7 @@ use crate::cuda_pipeline::gpu_dqn_trainer::{ }; use crate::cuda_pipeline::sp5_isv_slots::{ BUDGET_C51_BASE, BUDGET_IQN_BASE, BUDGET_CQL_BASE, BUDGET_ENS_BASE, + LB_CQL_ACTIVE_BASE, LB_C51_ACTIVE_BASE, ADAM_BETA1_BASE, ADAM_BETA2_BASE, ADAM_EPS_BASE, IQN_TAU_BASE, }; @@ -3370,10 +3371,23 @@ impl FusedTrainingCtx { /// budgets) feed the full-buf trunk/value scaling call — preserving SP5 Layer B /// behaviour for shared parameters (D3 decision in SP6 spec). /// - /// Cold-start bootstrap: ISV slots read 0 (sentinel) before first observation. - /// IQN keeps a structural BASE_IQN=0.11 floor (reference budget, Invariant 1 - /// carve-out). C51/CQL/ENS use sentinel-aware bootstrap: 0.05/0.02/0.02 - /// on cold start, controller value verbatim otherwise (SP7 Task 7). + /// Cold-start bootstrap: per-(head, branch) activation flag at + /// `LB_{CQL,C51}_ACTIVE_BASE` discriminates "controller has fired" from + /// "controller hasn't fired yet". IQN keeps a structural BASE_IQN=0.11 + /// floor (reference budget, Invariant 1 carve-out). CQL/C51 dispatch on + /// the activation flag: bootstrap (0.02/0.05) when active < 0.5, + /// controller verbatim (with EPS_DIV minimum) when active >= 0.5. ENS + /// keeps the prior numeric-equality bootstrap because no controller + /// drives it (SP5 Pearl 2 stopped writing BUDGET_ENS_BASE in SP7 T6). + /// + /// SP7 activation-flag fix (2026-05-03): replaces the prior + /// numeric-equality bootstrap (`raw < 1e-8 → bootstrap`) for CQL/C51, + /// which couldn't distinguish "controller said bootstrap" from + /// "controller hasn't fired yet" because `CQL_BOOTSTRAP_BUDGET=0.02` + /// numerically equaled the kernel's `COLD_START_FLOOR_CQL=0.02`. The + /// activation flag is a binary discriminator written by the controller + /// itself: 1.0 when the active-path computes a real budget update, + /// 0.0 (no-op via apply_pearls_ad short-circuit) on genuine cold-start. pub(crate) fn compute_adaptive_budgets( &mut self, ) -> ([f32; 4], [f32; 4], [f32; 4], [f32; 4], f32, f32, f32, f32) { @@ -3381,22 +3395,18 @@ impl FusedTrainingCtx { let mut iqn = [0.0_f32; 4]; let mut cql = [0.0_f32; 4]; let mut ens = [0.0_f32; 4]; - // SP7 Task 7 (2026-05-03): sentinel-aware bootstrap. The SP7 - // loss-balance controller writes adaptive per-branch budgets - // to BUDGET_CQL_BASE and BUDGET_C51_BASE; the controller may - // legitimately drive budgets near zero when the structural - // target says so (CQL when Q is flat, C51 when Q is sharp). - // A hard floor here would clamp the controller's intent. - // Instead, on a sentinel-0 read (cold start / fold boundary) - // we use the bootstrap value; otherwise we take the - // controller's value verbatim. - // - // Bootstrap values must match the cold-start basis in - // loss_balance_controller_kernel.cu (CQL_BOOTSTRAP=0.02, - // C51_BOOTSTRAP=0.05). IQN keeps the structural BASE_IQN - // floor (reference budget — can't be 0 without breaking - // the ratio compute in the controller). + // SP7 activation-flag fix (2026-05-03): per-(head, branch) activation + // dispatch. ACTIVE_THRESHOLD=0.5 is a binary discriminator (the slot + // takes 0.0 or 1.0); not a tuning knob. Bootstrap values must match + // the cold-start basis in loss_balance_controller_kernel.cu + // (COLD_START_FLOOR_CQL=0.02, COLD_START_FLOOR_C51=0.05). IQN keeps + // the structural BASE_IQN floor (reference budget — can't be 0 + // without breaking the ratio compute in the controller). ENS still + // uses the prior numeric-equality bootstrap because no controller + // drives BUDGET_ENS_BASE; the indistinguishable-state pathology + // doesn't apply there. const SP7_EPS_DIV: f32 = 1e-8; + const ACTIVE_THRESHOLD: f32 = 0.5; const CQL_BOOTSTRAP_BUDGET: f32 = 0.02; const C51_BOOTSTRAP_BUDGET: f32 = 0.05; const ENS_BOOTSTRAP_BUDGET: f32 = 0.02; @@ -3405,11 +3415,19 @@ impl FusedTrainingCtx { if raw < SP7_EPS_DIV { bootstrap } else { raw } }; for b in 0..4_usize { - c51[b] = budget_or_bootstrap(self.read_isv_signal_at(BUDGET_C51_BASE + b), - C51_BOOTSTRAP_BUDGET); + let cql_active = self.read_isv_signal_at(LB_CQL_ACTIVE_BASE + b); + let c51_active = self.read_isv_signal_at(LB_C51_ACTIVE_BASE + b); + cql[b] = if cql_active >= ACTIVE_THRESHOLD { + self.read_isv_signal_at(BUDGET_CQL_BASE + b).max(SP7_EPS_DIV) + } else { + CQL_BOOTSTRAP_BUDGET + }; + c51[b] = if c51_active >= ACTIVE_THRESHOLD { + self.read_isv_signal_at(BUDGET_C51_BASE + b).max(SP7_EPS_DIV) + } else { + C51_BOOTSTRAP_BUDGET + }; iqn[b] = self.read_isv_signal_at(BUDGET_IQN_BASE + b).max(BASE_IQN); - cql[b] = budget_or_bootstrap(self.read_isv_signal_at(BUDGET_CQL_BASE + b), - CQL_BOOTSTRAP_BUDGET); ens[b] = budget_or_bootstrap(self.read_isv_signal_at(BUDGET_ENS_BASE + b), ENS_BOOTSTRAP_BUDGET); } @@ -3437,15 +3455,20 @@ impl FusedTrainingCtx { /// SP7 loss-balance controller's ratio compute and cannot be driven to 0. pub(crate) fn last_iqn_budget_eff(&self) -> f32 { self.trainer.last_iqn_budget_eff } /// SP7 (2026-05-03): last per-step CQL budget (mean of 4 per-branch - /// values driven by `loss_balance_controller_kernel`). The previous - /// docstring claimed a "0.10×(1−regime)×health" formula that was - /// never implemented; per `feedback_trust_code_not_docs`, code wins - /// over docs. + /// values driven by `loss_balance_controller_kernel`). Each per-branch + /// value dispatches on the activation flag at `LB_CQL_ACTIVE_BASE+b`: + /// `CQL_BOOTSTRAP_BUDGET=0.02` while controller has yet to fire on this + /// branch (active < 0.5), controller verbatim (clamped at SP7_EPS_DIV) + /// once activated (active >= 0.5). FoldReset zeroes the activation flag + /// so each fold re-bootstraps cleanly through Pearl A. pub(crate) fn last_cql_budget_eff(&self) -> f32 { self.trainer.last_cql_budget_eff } /// SP7 (2026-05-03): last per-step C51 budget (mean of 4 per-branch - /// values driven by `loss_balance_controller_kernel`). The previous - /// docstring claimed a "1−iqn−cql−ens" formula that was never - /// implemented; per `feedback_trust_code_not_docs`, code wins over docs. + /// values driven by `loss_balance_controller_kernel`). Each per-branch + /// value dispatches on the activation flag at `LB_C51_ACTIVE_BASE+b`: + /// `C51_BOOTSTRAP_BUDGET=0.05` while controller has yet to fire on this + /// branch (active < 0.5), controller verbatim (clamped at SP7_EPS_DIV) + /// once activated (active >= 0.5). FoldReset zeroes the activation flag + /// so each fold re-bootstraps cleanly through Pearl A. pub(crate) fn last_c51_budget_eff(&self) -> f32 { self.trainer.last_c51_budget_eff } /// SP7 (2026-05-03): last per-step ensemble budget (mean of 4 per-branch /// values). ENS has no active controller driver in SP7 (Pearl 2 no longer diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 22038e5d8..3c461a857 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -641,6 +641,25 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "ISV[LB_SAMPLE_VAR_C51_BASE=309..313) — per-branch EMA of (c51_norm)² used as Wiener α denominator for the C51 budget controller. SP7 Pearl A sentinel 0 at fold boundary.", }, + // SP7 activation-flag fix (2026-05-03): per-(head, branch) + // activation flags. Monotonic per fold (kernel writes 1.0 once, + // never 0; apply_pearls_ad short-circuits 0 writes). FoldReset + // to 0 forces the new fold to re-activate from scratch as + // gradients populate, so the consumer's bootstrap covers the + // fold's first few steps and the controller's first real + // budget update lands without the indistinguishable-state + // pathology that existed pre-fix (kernel cold_start_basis + // numerically equaled consumer's bootstrap). + RegistryEntry { + name: "sp7_lb_cql_active", + category: ResetCategory::FoldReset, + description: "ISV[LB_CQL_ACTIVE_BASE=313..317) — SP7 controller activation flag for CQL — fold-boundary reset to 0 forces fresh activation as grads populate in the new fold.", + }, + RegistryEntry { + name: "sp7_lb_c51_active", + category: ResetCategory::FoldReset, + description: "ISV[LB_C51_ACTIVE_BASE=317..321) — SP7 controller activation flag for C51 — fold-boundary reset to 0 forces fresh activation as grads populate in the new fold.", + }, // SP5 Task A1: Wiener-state companion reset. The wiener_state_buf // now covers SP4 (71 producers × 3 = 213 floats) + SP5 (110 × 3 = 330 // floats) = 543 floats total. The SP5 triples start at offset 213. diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 484970826..748d7f2dc 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -6570,6 +6570,25 @@ impl DQNTrainer { } } } + // SP7 activation-flag fix (2026-05-03): per-(head, branch) + // activation flags. Zero at fold boundary so the new fold + // re-activates from scratch as grads populate. + "sp7_lb_cql_active" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp5_isv_slots::LB_CQL_ACTIVE_BASE; + for b in 0..4 { + fused.trainer().write_isv_signal_at(LB_CQL_ACTIVE_BASE + b, 0.0); + } + } + } + "sp7_lb_c51_active" => { + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp5_isv_slots::LB_C51_ACTIVE_BASE; + for b in 0..4 { + fused.trainer().write_isv_signal_at(LB_C51_ACTIVE_BASE + b, 0.0); + } + } + } "sp5_pnl_aggregation" => { // Layer D Task D1: zero ISV[286..290) — total/mean/var/max_dd — // so the new fold's first `launch_sp5_pnl_aggregation` triggers diff --git a/crates/ml/tests/sp5_producer_unit_tests.rs b/crates/ml/tests/sp5_producer_unit_tests.rs index bf69a590e..b6c85e8a3 100644 --- a/crates/ml/tests/sp5_producer_unit_tests.rs +++ b/crates/ml/tests/sp5_producer_unit_tests.rs @@ -99,6 +99,9 @@ const SP5_HEALTH_COMPOSITION_CUBIN: &[u8] = const SP5_TRAINING_METRICS_EMA_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/training_metrics_ema_kernel.cubin")); +const SP7_LOSS_BALANCE_CONTROLLER_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/loss_balance_controller_kernel.cubin")); + // ── Helpers ─────────────────────────────────────────────────────────────────── fn load_pearl_3_sigma_kernel(stream: &Arc) -> CudaFunction { @@ -2590,3 +2593,276 @@ fn training_metrics_ema_kernel_correctness() { "low_dd_ratio [cold]: expected {exp_low_dd_ratio:.6}, got {:.6}", out[2]); } } + +// ── SP7 activation-flag fix unit test ───────────────────────────────────────── +// +// Verifies the activation-flag transitions in `loss_balance_controller_update`: +// 1. Cold start (zero ISV, zero grad_decomp_result): kernel takes the genuine +// cold-start branch (h_n < EPS_DIV), writes 0.0 to scratch_active. The +// apply_pearls_ad_kernel short-circuit on step_obs == 0.0 is observable +// via the scratch buffer directly (we don't run apply_pearls in this test). +// 2. Active (non-zero IQN/CQL/C51 norms): kernel takes the active path, +// writes 1.0 to scratch_active and a real budget to scratch_budget. +// 3. Resulting budget reflects controller compute (NOT bootstrap constant). +// +// Per `feedback_no_cpu_test_fallbacks.md`: assertions are on analytically-known +// kernel behaviour (active flag binary, budget structurally `cold_start_basis · +// correction` from the seed-from-cold-start branch since old_budget=0). + +fn load_loss_balance_controller_kernel(stream: &Arc) -> CudaFunction { + let module = stream + .context() + .load_cubin(SP7_LOSS_BALANCE_CONTROLLER_CUBIN.to_vec()) + .expect("load loss_balance_controller_kernel cubin"); + module + .load_function("loss_balance_controller_update") + .expect("load loss_balance_controller_update function") +} + +#[test] +#[ignore = "requires GPU"] +fn sp7_loss_balance_controller_activation_flag_transitions() { + use ml::cuda_pipeline::sp5_isv_slots::{ + BUDGET_CQL_BASE, BUDGET_C51_BASE, FLATNESS_BASE, + LB_DIFF_VAR_CQL_BASE, LB_SAMPLE_VAR_CQL_BASE, + LB_DIFF_VAR_C51_BASE, LB_SAMPLE_VAR_C51_BASE, + LB_CQL_ACTIVE_BASE, LB_C51_ACTIVE_BASE, + }; + + // ISV must cover up to SP5_SLOT_END (321) AND EPOCH_IDX_INDEX (39). + const ISV_SIZE: usize = 321; + const EPOCH_IDX_INDEX: usize = 39; + + // grad_decomp_result_pinned: 27 floats (9 components × 3 [mag, dir, trunk]). + // Component element offsets per kernel header: iqn=0, cql_sx=6, c51=9. + // We allocate the full 27-float layout to mirror production. + const GRAD_DECOMP_TOTAL: usize = 27; + const GRAD_OFFSET_IQN: usize = 0; + const GRAD_OFFSET_CQL_SX: usize = 6; + const GRAD_OFFSET_C51: usize = 9; + + // Scratch layout: 8 output slot-blocks × 4 branches = 32 floats. + // We use small-but-distinct base offsets for the test. + const SCRATCH_BUDGET_CQL: usize = 0; // [0..4) + const SCRATCH_BUDGET_C51: usize = 4; // [4..8) + const SCRATCH_DIFF_CQL: usize = 8; // [8..12) + const SCRATCH_SAMPLE_CQL: usize = 12; // [12..16) + const SCRATCH_DIFF_C51: usize = 16; // [16..20) + const SCRATCH_SAMPLE_C51: usize = 20; // [20..24) + const SCRATCH_ACTIVE_CQL: usize = 24; // [24..28) + const SCRATCH_ACTIVE_C51: usize = 28; // [28..32) + const SCRATCH_SIZE: usize = 32; + + let stream = make_test_stream(); + let kernel = load_loss_balance_controller_kernel(&stream); + + let isv_buf = unsafe { MappedF32Buffer::new(ISV_SIZE) }.expect("alloc isv_buf"); + let grad_buf = unsafe { MappedF32Buffer::new(GRAD_DECOMP_TOTAL) }.expect("alloc grad_buf"); + let scratch_buf = unsafe { MappedF32Buffer::new(SCRATCH_SIZE) }.expect("alloc scratch_buf"); + + let isv_dev = isv_buf.dev_ptr; + let scratch_dev = scratch_buf.dev_ptr; + let f32_size = std::mem::size_of::() as u64; + let iqn_dev = grad_buf.dev_ptr + (GRAD_OFFSET_IQN as u64) * f32_size; + let cql_sx_dev = grad_buf.dev_ptr + (GRAD_OFFSET_CQL_SX as u64) * f32_size; + let c51_dev = grad_buf.dev_ptr + (GRAD_OFFSET_C51 as u64) * f32_size; + + let flatness_isv_base_i32 = FLATNESS_BASE as i32; + let budget_cql_isv_base_i32 = BUDGET_CQL_BASE as i32; + let budget_c51_isv_base_i32 = BUDGET_C51_BASE as i32; + let diff_var_cql_isv_base_i32 = LB_DIFF_VAR_CQL_BASE as i32; + let sample_var_cql_isv_base_i32 = LB_SAMPLE_VAR_CQL_BASE as i32; + let diff_var_c51_isv_base_i32 = LB_DIFF_VAR_C51_BASE as i32; + let sample_var_c51_isv_base_i32 = LB_SAMPLE_VAR_C51_BASE as i32; + let active_cql_isv_base_i32 = LB_CQL_ACTIVE_BASE as i32; + let active_c51_isv_base_i32 = LB_C51_ACTIVE_BASE as i32; + let epoch_idx_isv_index_i32 = EPOCH_IDX_INDEX as i32; + let sb_budget_cql_i32 = SCRATCH_BUDGET_CQL as i32; + let sb_budget_c51_i32 = SCRATCH_BUDGET_C51 as i32; + let sb_diff_var_cql_i32 = SCRATCH_DIFF_CQL as i32; + let sb_sample_var_cql_i32 = SCRATCH_SAMPLE_CQL as i32; + let sb_diff_var_c51_i32 = SCRATCH_DIFF_C51 as i32; + let sb_sample_var_c51_i32 = SCRATCH_SAMPLE_C51 as i32; + let sb_active_cql_i32 = SCRATCH_ACTIVE_CQL as i32; + let sb_active_c51_i32 = SCRATCH_ACTIVE_C51 as i32; + + let cfg = LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (8, 1, 1), + shared_mem_bytes: 0, + }; + + // Helper: launch the controller kernel once with current ISV/grad state. + let launch_once = |isv_data: &[f32], grad_data: &[f32], scratch_seed: f32| { + isv_buf.write_from_slice(isv_data); + grad_buf.write_from_slice(grad_data); + // Seed scratch with a marker so we can detect "kernel didn't write". + let scratch_init = vec![scratch_seed; SCRATCH_SIZE]; + scratch_buf.write_from_slice(&scratch_init); + unsafe { + stream + .launch_builder(&kernel) + .arg(&iqn_dev) + .arg(&cql_sx_dev) + .arg(&c51_dev) + .arg(&isv_dev) + .arg(&flatness_isv_base_i32) + .arg(&budget_cql_isv_base_i32) + .arg(&budget_c51_isv_base_i32) + .arg(&diff_var_cql_isv_base_i32) + .arg(&sample_var_cql_isv_base_i32) + .arg(&diff_var_c51_isv_base_i32) + .arg(&sample_var_c51_isv_base_i32) + .arg(&active_cql_isv_base_i32) + .arg(&active_c51_isv_base_i32) + .arg(&epoch_idx_isv_index_i32) + .arg(&scratch_dev) + .arg(&sb_budget_cql_i32) + .arg(&sb_budget_c51_i32) + .arg(&sb_diff_var_cql_i32) + .arg(&sb_sample_var_cql_i32) + .arg(&sb_diff_var_c51_i32) + .arg(&sb_sample_var_c51_i32) + .arg(&sb_active_cql_i32) + .arg(&sb_active_c51_i32) + .launch(cfg) + .expect("launch loss_balance_controller_update"); + } + stream.synchronize().expect("sync after loss_balance_controller_update"); + scratch_buf.read_all() + }; + + // ── Step 1: cold start (all-zero ISV, all-zero grads). ─────────────────── + // Kernel hits cold-start branch: h_n=0 < EPS_DIV. was_active=0 (ISV is zero). + // Expect: scratch_active = 0.0 for all 4 branches × 2 heads. + // (apply_pearls_ad short-circuit on 0.0 means the consumer sees the + // bootstrap value, but we test the scratch directly here.) + let isv_cold = vec![0.0_f32; ISV_SIZE]; + let grad_cold = vec![0.0_f32; GRAD_DECOMP_TOTAL]; + + let scratch_after_cold = launch_once(&isv_cold, &grad_cold, /* seed = */ -1.0); + + for b in 0..4_usize { + let cql_active = scratch_after_cold[SCRATCH_ACTIVE_CQL + b]; + let c51_active = scratch_after_cold[SCRATCH_ACTIVE_C51 + b]; + assert!( + cql_active.abs() < 1e-7, + "cold start branch={b}: cql_active={cql_active:.6} should be 0.0 (genuine cold start)" + ); + assert!( + c51_active.abs() < 1e-7, + "cold start branch={b}: c51_active={c51_active:.6} should be 0.0 (genuine cold start)" + ); + // Budget slots also written 0.0 (no-op via apply_pearls_ad short-circuit). + assert!( + scratch_after_cold[SCRATCH_BUDGET_CQL + b].abs() < 1e-7, + "cold start branch={b}: budget_cql={:.6} should be 0.0 (no-op marker)", + scratch_after_cold[SCRATCH_BUDGET_CQL + b] + ); + assert!( + scratch_after_cold[SCRATCH_BUDGET_C51 + b].abs() < 1e-7, + "cold start branch={b}: budget_c51={:.6} should be 0.0 (no-op marker)", + scratch_after_cold[SCRATCH_BUDGET_C51 + b] + ); + } + + // ── Step 2: active (synthetic non-zero IQN/CQL/C51 norms). ─────────────── + // grad_decomp pinned layout per component: [mag, dir, trunk]. We populate + // all 3 slices so every branch has a non-zero h_n / iqn_n. + // IQN slice norms: [1.0, 1.0, 1.0] + // CQL_SX slice norms: [0.5, 0.5, 0.5] + // C51 slice norms: [0.3, 0.3, 0.3] + // ISV: epoch_idx = 0 (first step) → welford_alpha = 1/max(1,0) = 1.0 → full + // first-step update. Flatness defaults to 0 (clamped) → target_ratio for + // CQL = 2.0, target_ratio for C51 = 0.0. With actual_ratio_CQL = 0.5/1.0 = + // 0.5 → correction_CQL = 2.0/0.5 = 4.0. Old budget = 0 → seed-from-cold-start + // branch: new_budget = COLD_START_FLOOR_CQL · correction = 0.02 · 4.0 = 0.08. + // For C51: target=0 → correction=0 → new_budget = COLD_START_FLOOR_C51 · 0 = 0 + // (clamped at EPS_DIV = 1e-8). + let mut grad_active = vec![0.0_f32; GRAD_DECOMP_TOTAL]; + for slice in 0..3_usize { + grad_active[GRAD_OFFSET_IQN + slice] = 1.0; + grad_active[GRAD_OFFSET_CQL_SX + slice] = 0.5; + grad_active[GRAD_OFFSET_C51 + slice] = 0.3; + } + + let scratch_after_active = launch_once(&isv_cold, &grad_active, /* seed = */ -1.0); + + const EPS_DIV: f32 = 1e-8; + const COLD_START_FLOOR_CQL: f32 = 0.02; + + for b in 0..4_usize { + let cql_active = scratch_after_active[SCRATCH_ACTIVE_CQL + b]; + let c51_active = scratch_after_active[SCRATCH_ACTIVE_C51 + b]; + assert!( + cql_active >= 0.5, + "active branch={b}: cql_active={cql_active:.6} should be 1.0 after grads populated" + ); + assert!( + c51_active >= 0.5, + "active branch={b}: c51_active={c51_active:.6} should be 1.0 after grads populated" + ); + + // Budget slots reflect controller's seed-from-cold-start compute, not + // the bootstrap value. CQL: target_ratio=2.0, actual=0.5, + // correction=4.0, seed = 0.02·4.0 = 0.08. C51: target_ratio=0.0 + // (flatness=0), so seed = 0.05·0 = 0 → clamped at EPS_DIV. + let budget_cql = scratch_after_active[SCRATCH_BUDGET_CQL + b]; + let budget_c51 = scratch_after_active[SCRATCH_BUDGET_C51 + b]; + + let expected_cql = COLD_START_FLOOR_CQL * 4.0; + assert!( + (budget_cql - expected_cql).abs() / expected_cql < 0.01, + "active branch={b}: budget_cql={budget_cql:.6} expected {expected_cql:.6} (cold_start_basis · correction; not the {COLD_START_FLOOR_CQL:.6} bootstrap)" + ); + // C51 target_ratio = ANCHOR_C51_RATIO · flatness = 1.0 · 0 = 0; the + // seed becomes COLD_START_FLOOR_C51 · 0 = 0 → clamped at EPS_DIV. + // The load-bearing assertion: budget != bootstrap (0.05). + assert!( + budget_c51 < 0.04, + "active branch={b}: budget_c51={budget_c51:.6} should be ~EPS_DIV (target_ratio=0), not bootstrap 0.05" + ); + assert!( + budget_c51 >= EPS_DIV * 0.99, + "active branch={b}: budget_c51={budget_c51:.6} should be at least EPS_DIV={EPS_DIV:.1e}" + ); + } + + // ── Step 3: re-launch with prior was_active flags simulated by writing 1.0 + // into LB_*_ACTIVE_BASE. Even with grads gated off (back to zero), + // transient cold-start should hold prior budget verbatim and re-assert + // active=1, NOT collapse back to consumer bootstrap. + let mut isv_active = vec![0.0_f32; ISV_SIZE]; + for b in 0..4_usize { + isv_active[LB_CQL_ACTIVE_BASE + b] = 1.0; // was_active = 1 + isv_active[LB_C51_ACTIVE_BASE + b] = 1.0; + isv_active[BUDGET_CQL_BASE + b] = 0.07; // prior budget != bootstrap + isv_active[BUDGET_C51_BASE + b] = 0.04; // prior budget != bootstrap + } + + let scratch_after_transient = launch_once(&isv_active, &grad_cold, /* seed = */ -1.0); + for b in 0..4_usize { + // Active flag held at 1. + assert!( + scratch_after_transient[SCRATCH_ACTIVE_CQL + b] >= 0.5, + "transient grad-gate branch={b}: cql_active={:.6} should stay 1.0 (was_active=1)", + scratch_after_transient[SCRATCH_ACTIVE_CQL + b] + ); + assert!( + scratch_after_transient[SCRATCH_ACTIVE_C51 + b] >= 0.5, + "transient grad-gate branch={b}: c51_active={:.6} should stay 1.0 (was_active=1)", + scratch_after_transient[SCRATCH_ACTIVE_C51 + b] + ); + // Budget held verbatim at prior value (no drift on a measurement gap). + assert!( + (scratch_after_transient[SCRATCH_BUDGET_CQL + b] - 0.07).abs() < 1e-5, + "transient grad-gate branch={b}: budget_cql={:.6} should hold prior 0.07", + scratch_after_transient[SCRATCH_BUDGET_CQL + b] + ); + assert!( + (scratch_after_transient[SCRATCH_BUDGET_C51 + b] - 0.04).abs() < 1e-5, + "transient grad-gate branch={b}: budget_c51={:.6} should hold prior 0.04", + scratch_after_transient[SCRATCH_BUDGET_C51 + b] + ); + } +} diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 1d2c82d7a..87eb6fce2 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -3935,6 +3935,66 @@ extended. No producer kernel yet — that arrives in the next commit. test bug exposed by SP7 smoke; eval_dist is structurally Kelly-cold-start bound and shouldn't gate Q-learning checks per project_magnitude_eval_collapse_kelly_capped. +- Activation-flag fix (commit ⟨pending⟩, 2026-05-03): SP7 controller + was dormant at production scale — per-branch budgets sat at exactly + the bootstrap constants (CQL=0.02, C51=0.05, IQN=0.11) for every + branch, every epoch in `smoke-test-8556k`. Two architectural defects + combined: (1) the kernel's `cold_start_basis` (COLD_START_FLOOR_CQL= + 0.02 in `loss_balance_controller_kernel.cu`) numerically equaled the + consumer's bootstrap fallback (CQL_BOOTSTRAP_BUDGET=0.02 in + `compute_adaptive_budgets`) — the consumer couldn't distinguish + "controller said bootstrap" from "controller hasn't fired yet"; (2) + the Wiener α was clamped at ALPHA_FLOOR=1e-4 from step 1 because + sample_var = h_n² (gradient-norm scale) dominated diff_var (budget- + delta scale) — α was glacial. + + Fix in 8 atomic touches: + * 8 new ISV slots `LB_{CQL,C51}_ACTIVE_BASE` (313..317, 317..321) + per (head × branch). `SP5_SLOT_END` 313→321; `SP5_PRODUCER_COUNT` + 139→147; layout fingerprint and slot-contiguity tests updated + in lockstep. Activation flag is monotonic per fold (kernel writes + 1.0 once; apply_pearls_ad's `step_obs == 0.0` short-circuit means + writing 0.0 is a no-op). FoldReset zeroes the slots so each fold + re-bootstraps cleanly through Pearl A. + * 2 new SCRATCH offsets `SCRATCH_LB_ACTIVE_{CQL,C51}` (242..246, + 246..250). `SP5_SCRATCH_TOTAL` 242→250. + * Kernel signature 18→24 args: 5 new (active CQL/C51 ISV bases, + epoch_idx ISV index, scratch active CQL/C51 bases) plus 1 new + scratch_active write per active path. New cold-start branch logic: + was_active>=0.5 (transient grad-gate) holds prior budget verbatim + and re-asserts active=1; was_active<0.5 (genuine cold start) + writes 0.0 to all 4 scratch slots — apply_pearls_ad short-circuit + leaves ISV untouched. + * Welford-α hybrid added to active path: `welford_α = 1 / max(1, + epoch_idx)` (full update on first active step, falls off as 1/N + within the fold), `α_eff = clamp(max(welford_α, wiener_α), + ALPHA_FLOOR, ALPHA_CEIL)`. Wiener takes over once accumulated + variance is meaningful. EPOCH_IDX_INDEX (per-fold-reset epoch + counter, ISV[39]) is the per-fold step-count proxy. + * Launcher: 5 new i32 conversions, 5 new `.arg()` calls; the + apply_pearls_ad smoothing loop extended from 6 slot-blocks (× 4 + branches = 24 launches) to 8 (× 4 = 32 launches) covering the + 2 new active slot ranges. + * Consumer `compute_adaptive_budgets` (`fused_training.rs`): + CQL/C51 dispatch on `LB_{CQL,C51}_ACTIVE_BASE+b >= 0.5` instead + of the prior `raw < 1e-8` numeric-equality bootstrap. ENS keeps + the prior bootstrap (no controller drives BUDGET_ENS_BASE). + Stale docstrings on `last_cql_budget_eff` / + `last_c51_budget_eff` corrected to describe the activation-flag + semantics per `feedback_trust_code_not_docs.md`. + * 2 new `RegistryEntry` blocks (`sp7_lb_cql_active`, + `sp7_lb_c51_active`, both FoldReset) + matching dispatch arms in + `reset_named_state`. Contract test + (`every_fold_and_soft_reset_entry_has_dispatch_arm`) gates + compile. + * GPU unit test + `sp7_loss_balance_controller_activation_flag_transitions` + (added to `sp5_producer_unit_tests.rs`) exercises 3 transitions: + genuine cold start → both flags 0; first active step → both flags + 1 with controller-computed budget != bootstrap; transient grad- + gate (was_active=1, grads back to 0) → flags hold at 1, prior + budget held verbatim. No CPU reference oracle per + `feedback_no_cpu_test_fallbacks.md`. Passes on local RTX 3050 Ti. - T8 (out-of-tree): memory pearl `pearl_loss_balance_controller.md` + MEMORY.md index entry. Captures the two-layer (signal-modulated target × outcome-driven α) pattern for future controller designs. From bbf15ad690999cef94701581ed3c067d08a837af Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sun, 3 May 2026 12:33:56 +0200 Subject: [PATCH 2/5] diag(sp7): HEALTH_DIAG emit Q_VAR_PER_BRANCH signal MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The SP7 controller's flatness gate reads per-branch Q-variance from ISV[Q_VAR_PER_BRANCH_BASE=222..226) but that signal was not visible in HEALTH_DIAG. The existing "var_q" in the main HEALTH_DIAG line is realized step-return variance per magnitude bin from gpu_experience_collector — a trade-outcome metric, semantically distinct from per-branch Q-output variance. Added one emit line immediately after cql_budget_per_branch: HEALTH_DIAG[E]: q_var_per_branch [dir=X.XXXX mag=X.XXXX ord=X.XXXX urg=X.XXXX] Reads ISV[222..226) via read_isv_signal_at — the same slots written by q_branch_stats_kernel.cu (scratch slot 2 per branch) and routed into ISV by apply_pearls_ad_kernel in launch_sp5_pearl_1_atom. No new ISV slots, no kernel change, no StateResetRegistry entry. Class 2 signal (mag_concat_scale / q_rms): Option A infeasible — q_rms is a per-sample register variable in mag_concat_qdir with no existing ISV slot; h_s2_rms_ema at ISV[96] is the only available proxy. Option B (new ISV slot) blocked pending explicit controller OK. See audit doc for full stop-and-report rationale. Files touched: crates/ml/src/trainers/dqn/trainer/training_loop.rs (+28 LOC) docs/dqn-wire-up-audit.md (+13 LOC) [ISV slot decision: Option A reused existing Q_VAR_PER_BRANCH_BASE=222..226] Cargo check workspace clean. State-reset contract test passes. Co-Authored-By: Claude Sonnet 4.6 --- .../src/trainers/dqn/trainer/training_loop.rs | 28 +++++++++++++++++++ docs/dqn-wire-up-audit.md | 13 +++++++++ 2 files changed, 41 insertions(+) diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 748d7f2dc..e90b67b54 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -4072,6 +4072,34 @@ impl DQNTrainer { ); } + // SP7 observability: per-branch Q-variance HEALTH_DIAG line. + // Reads ISV[Q_VAR_PER_BRANCH_BASE..+4) = ISV[222..226), written by + // `q_branch_stats_kernel.cu` (scratch slot 2 per branch) and routed into + // ISV via `apply_pearls_ad_kernel` in `launch_sp5_pearl_1_atom`. + // This is the ACTUAL signal the SP7 controller's flatness gate reads + // from `FLATNESS_BASE = ISV[206..210)` (derived from Q-variance here). + // Semantically distinct from `mag_stats [var_q/h/f]` in the main + // HEALTH_DIAG line, which is the realized step-return variance per + // magnitude bin from `gpu_experience_collector.per_magnitude_winrate_and_variance()`. + { + use crate::cuda_pipeline::sp5_isv_slots::Q_VAR_PER_BRANCH_BASE; + let (qv_dir, qv_mag, qv_ord, qv_urg) = if let Some(ref fused) = self.fused_ctx { + let trainer = fused.trainer(); + ( + trainer.read_isv_signal_at(Q_VAR_PER_BRANCH_BASE), + trainer.read_isv_signal_at(Q_VAR_PER_BRANCH_BASE + 1), + trainer.read_isv_signal_at(Q_VAR_PER_BRANCH_BASE + 2), + trainer.read_isv_signal_at(Q_VAR_PER_BRANCH_BASE + 3), + ) + } else { + (0.0_f32, 0.0_f32, 0.0_f32, 0.0_f32) + }; + tracing::info!( + "HEALTH_DIAG[{}]: q_var_per_branch [dir={:.4} mag={:.4} ord={:.4} urg={:.4}]", + epoch, qv_dir, qv_mag, qv_ord, qv_urg, + ); + } + // C.2 Plan 3 Task 1 (spec §4.C.2): reward_split HEALTH_DIAG line. // Reads 6 ISV EMA slots updated by the GPU reward_component_ema kernel // launched just above. CPU-side code only reads; GPU wrote the values. diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 87eb6fce2..7481ee9e7 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -3999,5 +3999,18 @@ extended. No producer kernel yet — that arrives in the next commit. MEMORY.md index entry. Captures the two-layer (signal-modulated target × outcome-driven α) pattern for future controller designs. - T9–T10: smoke + 50-epoch verification. +- SP7 observability (2026-05-03, wt/sp7-observability): additive HEALTH_DIAG + emit for `q_var_per_branch [dir mag ord urg]` — the actual Q-variance signal + the SP7 controller reads from ISV[222..226). No new ISV slots, no kernel + change, no StateResetRegistry entry: purely reads the existing + `Q_VAR_PER_BRANCH_BASE` slots (written by `q_branch_stats_kernel.cu` via + `apply_pearls_ad_kernel`). Placed after the `cql_budget_per_branch` emit in + the per-epoch HEALTH_DIAG block (`training_loop.rs`). Semantically + distinguished from `mag_stats [var_q/h/f]` which is realized step-return + variance per magnitude bin, not Q-output variance. Class 2 signal + (`mag_concat_scale` / q_rms) is infeasible via Option A: `q_rms` in + `mag_concat_qdir` is a per-sample register variable with no existing ISV + slot; `h_s2_rms_ema` (ISV[96]) is the only available proxy. Option B blocked + pending controller OK (no new ISV slots without explicit approval). Contract test — every RegistryEntry has dispatch arm (2026-05-03): added `every_fold_and_soft_reset_entry_has_dispatch_arm` unit test to the existing `#[cfg(test)] mod tests` block in `state_reset_registry.rs`. Source-introspection design: `include_str!` both `state_reset_registry.rs` and `trainer/training_loop.rs`; brace-balance walk extracts the `match name {` body; for each FoldReset/SoftReset entry (97 total: 95 FoldReset + 2 SoftReset) asserts the literal `""` appears in that body. No production-code change. No new dependency (manual brace walk instead of regex). Catches the recurring "add RegistryEntry, forget dispatch arm → fold-boundary panic" bug (occurred twice: SP5 Layer A #281, SP7 T7 commit 6e479c55c) at `cargo test -p ml --lib` rather than mid-training. Test currently passes (all 97 dispatch arms present post SP7 T7 fix). Touched: `trainers/dqn/state_reset_registry.rs` (+59 LOC in cfg-test block). From 237b3dbfb0535ca21ccfc186038f067f48611732 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sun, 3 May 2026 12:45:59 +0200 Subject: [PATCH 3/5] =?UTF-8?q?fix(sp7):=20bump=20ISV=5FTOTAL=5FDIM=20294?= =?UTF-8?q?=E2=86=92321=20to=20match=20SP5=5FSLOT=5FEND?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit SP7 T1 added 24 ISV slots (LB_DIFF_VAR_CQL_BASE=297 ... LB_C51_ACTIVE through 321) without bumping ISV_TOTAL_DIM, leaving SP7's wiener-state and activation slots out-of-bounds of the allocated pinned buffer (294 * 4 = 1176 bytes; SP7 slots write to bytes 1188..1284). GPU direct-pointer writes silently corrupted memory in the next page; CPU read_isv_signal_at returned garbage in release builds. This explains the SP7 smoke's confusing zero-budget-everywhere observation: the activation flag was reading OOB memory, the consumer saw inconsistent values, and the controller's actual ISV state was never visible. Bumped ISV_TOTAL_DIM to 321 (max valid index 320, covering SP5_SLOT_END-1). Made pub(crate) so the new contract test can reference it from sp5_isv_slots.rs. Updated layout_fingerprint_seed()'s slot entries to include the previously missing D3 and SP7 slots (TRAINING_SHARPE_EMA=294 through LB_C51_ACTIVE_BASE=317) and updated ISV_TOTAL_DIM= literal to 321 in lockstep per feedback_no_partial_refactor. Added contract test all_sp5_slots_fit_within_isv_total_dim to permanently gate this class of bug at cargo test time. The test would have caught SP7 T1 instantly; future slot allocations cannot regress this. Files: gpu_dqn_trainer.rs (ISV_TOTAL_DIM + comment + fingerprint), sp5_isv_slots.rs (test), docs/dqn-wire-up-audit.md (Fix 31 sub-bullet). Cargo check workspace clean. Cargo test ml --lib 935 passed (934 baseline + 1 new contract test); 16 pre-existing GPU-hardware failures unchanged. Co-Authored-By: Claude Sonnet 4.6 --- .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 8 +++++-- crates/ml/src/cuda_pipeline/sp5_isv_slots.rs | 11 ++++++++++ docs/dqn-wire-up-audit.md | 21 +++++++++++++++++++ 3 files changed, 38 insertions(+), 2 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index ed4f7bffc..b61844e31 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -825,7 +825,7 @@ const ISV_NETWORK_DIM: usize = 23; /// (shifted 112→116 in Plan 4 Task 6 Commit A). /// Written by the constructor; checked at checkpoint load. Fail-fast only — no migration /// path exists. See spec §4.A.2 and `LAYOUT_FINGERPRINT_CURRENT` for structural-hash rationale. -const ISV_TOTAL_DIM: usize = 294; // SP5 + Layer D D1+D2: 173 + 121 (118 SP5 slots @ 174..278/280..294, with 2-slot gap before cross-fold-persistent Kelly block; Layer D D1 PnL outputs at [286..290); Layer D D2 health composition outputs at [290..294)) +pub(crate) const ISV_TOTAL_DIM: usize = 321; // SP5 + Layer D D1+D2+D3 + SP7: 173 + 148 (145 SP5 slots @ 174..278/280..321, with 2-slot gap before cross-fold-persistent Kelly block; Layer D D1 PnL outputs at [286..290); Layer D D2 health composition outputs at [290..294); Layer D D3 training metrics EMA at [294..297); SP7 T1 loss-balance Wiener stats at [297..313); SP7 activation-flag fix per-(head,branch) flags at [313..321)) /// Legacy alias preserved for call sites that haven't been audited for the /// network-vs-total split. New code should pick `ISV_NETWORK_DIM` (for weight /// tensor sizing) or `ISV_TOTAL_DIM` (for the broadcast bus buffer). @@ -1699,7 +1699,11 @@ const fn layout_fingerprint_seed() -> &'static [u8] { KELLY_SAMPLE_COUNT=283;WIN_RATE_SMOOTH=284;LOSS_RATE_SMOOTH=285;\ PNL_TOTAL=286;PNL_MEAN=287;PNL_VAR=288;PNL_MAX_DD=289;\ HEALTH_SCORE=290;Q_GAP_NORM=291;Q_VAR_NORM=292;GRAD_NORM_NORM=293;\ - ISV_TOTAL_DIM=294;\ + TRAINING_SHARPE_EMA=294;MAX_DD_EMA=295;LOW_DD_RATIO=296;\ + LB_DIFF_VAR_CQL_BASE=297;LB_SAMPLE_VAR_CQL_BASE=301;\ + LB_DIFF_VAR_C51_BASE=305;LB_SAMPLE_VAR_C51_BASE=309;\ + LB_CQL_ACTIVE_BASE=313;LB_C51_ACTIVE_BASE=317;\ + ISV_TOTAL_DIM=321;\ PARAM_W_A_H_S1=0;PARAM_B_A_H_S1=1;PARAM_W_B_H_S1=2;PARAM_B_B_H_S1=3;\ PARAM_W_RESIDUAL_H_S1=4;PARAM_GAMMA_H_S1=5;PARAM_BETA_H_S1=6;\ PARAM_W_A_H_S2=7;PARAM_B_A_H_S2=8;PARAM_W_B_H_S2=9;PARAM_B_B_H_S2=10;\ diff --git a/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs index 178b4c511..829565a01 100644 --- a/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs +++ b/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs @@ -506,4 +506,15 @@ mod tests { assert_eq!(SP5_SLOT_END, 321); assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE); } + + #[test] + fn all_sp5_slots_fit_within_isv_total_dim() { + use crate::cuda_pipeline::gpu_dqn_trainer::ISV_TOTAL_DIM; + assert!( + SP5_SLOT_END <= ISV_TOTAL_DIM, + "SP5_SLOT_END={} exceeds ISV_TOTAL_DIM={} — bus too small for SP5/SP7 slots; \ + bump ISV_TOTAL_DIM in gpu_dqn_trainer.rs (and update layout_fingerprint_seed()).", + SP5_SLOT_END, ISV_TOTAL_DIM, + ); + } } diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 7481ee9e7..735e3c71b 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -3995,6 +3995,27 @@ extended. No producer kernel yet — that arrives in the next commit. gate (was_active=1, grads back to 0) → flags hold at 1, prior budget held verbatim. No CPU reference oracle per `feedback_no_cpu_test_fallbacks.md`. Passes on local RTX 3050 Ti. +- ISV_TOTAL_DIM OOB fix (commit on wt/sp7-observability, 2026-05-03): + **Discovery**: `ISV_TOTAL_DIM=294` but `SP5_SLOT_END=321`; the pinned + ISV buffer was allocated at `294×4=1176 bytes`. SP7 T1 added 24 slots + (ISV[297..321)) and the activation-flag fix added 8 more (ISV[313..321)) + without bumping `ISV_TOTAL_DIM`. All SP7 ISV writes/reads from indices + 297..320 were out-of-bounds: GPU direct-pointer writes silently + corrupted memory in the next page-aligned region; CPU + `write_isv_signal_at` silently no-oped for `index >= ISV_TOTAL_DIM`; + CPU `read_isv_signal_at` returned garbage in release builds. + **Root cause**: T1 allocated slots in `sp5_isv_slots.rs` but did not + bump `ISV_TOTAL_DIM` in `gpu_dqn_trainer.rs`; the constant lived in a + different file with no compile-time linkage to `SP5_SLOT_END`. + **Fix**: `ISV_TOTAL_DIM` bumped `294→321` and made `pub(crate)`; + `layout_fingerprint_seed()` extended with the missing D3 and SP7 slot + entries (TRAINING_SHARPE_EMA=294 through LB_C51_ACTIVE_BASE=317) and + `ISV_TOTAL_DIM=321` in lockstep per `feedback_no_partial_refactor`. + **Contract test**: `all_sp5_slots_fit_within_isv_total_dim` added to + `sp5_isv_slots.rs` test module; asserts `SP5_SLOT_END <= ISV_TOTAL_DIM` + at `cargo test -p ml --lib`. This test would have caught the SP7 T1 + miss instantly; future slot allocations cannot silently regress the bus + size without breaking CI. - T8 (out-of-tree): memory pearl `pearl_loss_balance_controller.md` + MEMORY.md index entry. Captures the two-layer (signal-modulated target × outcome-driven α) pattern for future controller designs. From 6cf6f26ab90f9255034a7bf14885f1b9223c32a1 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sun, 3 May 2026 13:30:57 +0200 Subject: [PATCH 4/5] diag(sp7): emit grad_decomp_pinned + lb_active_per_branch HEALTH_DIAG MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The SP7 smoke at 237b3dbfb showed all 8 (head, branch) combinations stuck at bootstrap across 5 epochs, despite Q-variance and gradient norms being non-zero. The activation flag mechanism is monotonic per fold (once active, stays active), so 8/8 inactive across all observed epochs implies the kernel never hits the active path — but we can't confirm whether that's because (a) grad_decomp_result_pinned reads zero at SP7's epoch-boundary call site, or (b) something else. Added two HEALTH_DIAG lines: - grad_decomp_pinned — reads bytes [0..12,24..36,36..48] of the pinned buffer the SP7 kernel reads; surfaces iqn/cql_sx/c51 mag/ dir/trunk norms exactly as the kernel sees them. - lb_active_per_branch — reads LB_{CQL,C51}_ACTIVE_BASE+0..4 from ISV; surfaces the activation flag state per (head, branch). Together these disambiguate "grad_decomp not populated at SP7 read time" from "values populated but kernel cold-start branch fires for other reasons". Purely additive — no behavior change. Audit doc Fix 31 sub-bullet. Co-Authored-By: Claude Opus 4.7 (1M context) --- .../src/trainers/dqn/trainer/training_loop.rs | 81 +++++++++++++++++++ docs/dqn-wire-up-audit.md | 33 ++++++++ 2 files changed, 114 insertions(+) diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index e90b67b54..d222b720b 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -4100,6 +4100,87 @@ impl DQNTrainer { ); } + // SP7 observability (grad_decomp_pinned): surface the exact contents of + // grad_decomp_result_pinned[0..12, 24..36, 36..48] (iqn/cql_sx/c51 × + // [mag, dir, trunk]) that the SP7 loss-balance controller read in the + // last step of this epoch. The values are already cached in the + // DQNTrainer host arrays by `refresh_grad_component_norms` (called ~300 + // lines above in the epoch-boundary block); reading them here costs + // zero extra synchronisation. + // + // Component index mapping (matches `grad_decomp_result_pinned` layout): + // IQN = component 0 → pinned offsets [0, 1, 2] + // CQL_SX = component 2 → pinned offsets [6, 7, 8] + // C51 = component 3 → pinned offsets [9, 10, 11] + // + // This disambiguates "grad_decomp not populated at SP7 read time" + // (all zeros here) from "values non-zero but cold-start branch fires + // for another reason" (non-zeros here, lb_active stays 0.0). + { + let (iqn_m, iqn_d, iqn_t, + cql_sx_m, cql_sx_d, cql_sx_t, + c51_m, c51_d, c51_t) = if let Some(ref fused) = self.fused_ctx { + let trainer = fused.trainer(); + let mag = trainer.grad_component_norms_mag(); + let dir = trainer.grad_component_norms_dir(); + let trunk = trainer.grad_component_norms_trunk(); + (mag[0], dir[0], trunk[0], // IQN + mag[2], dir[2], trunk[2], // CQL_SX + mag[3], dir[3], trunk[3]) // C51 + } else { + (0.0_f32, 0.0_f32, 0.0_f32, + 0.0_f32, 0.0_f32, 0.0_f32, + 0.0_f32, 0.0_f32, 0.0_f32) + }; + tracing::info!( + "HEALTH_DIAG[{}]: grad_decomp_pinned [iqn:[m={:.4} d={:.4} t={:.4}] cql_sx:[m={:.4} d={:.4} t={:.4}] c51:[m={:.4} d={:.4} t={:.4}]]", + epoch, + iqn_m, iqn_d, iqn_t, + cql_sx_m, cql_sx_d, cql_sx_t, + c51_m, c51_d, c51_t, + ); + } + + // SP7 observability (lb_active_per_branch): surface the per-(head, + // branch) activation-flag state written by the SP7 controller and + // smoothed into ISV[LB_{CQL,C51}_ACTIVE_BASE..+4) by + // `apply_pearls_ad_kernel`. Values are 0.0 (cold-start, controller + // has not yet fired on this (head, branch)) or approaching 1.0 + // (active path executed at least once this fold). + // + // Paired with `grad_decomp_pinned` above: if both are all-zero the + // pinned buffer was empty at SP7's read site; if grad_decomp is + // non-zero but lb_active stays 0.0 the cold-start branch fired for + // a different reason (e.g. Welford-α cold-start or EPS_DIV guard). + { + use crate::cuda_pipeline::sp5_isv_slots::{ + LB_CQL_ACTIVE_BASE, LB_C51_ACTIVE_BASE, + }; + let (cql_d, cql_m, cql_o, cql_u, + c51_d, c51_m, c51_o, c51_u) = if let Some(ref fused) = self.fused_ctx { + let trainer = fused.trainer(); + ( + trainer.read_isv_signal_at(LB_CQL_ACTIVE_BASE), + trainer.read_isv_signal_at(LB_CQL_ACTIVE_BASE + 1), + trainer.read_isv_signal_at(LB_CQL_ACTIVE_BASE + 2), + trainer.read_isv_signal_at(LB_CQL_ACTIVE_BASE + 3), + trainer.read_isv_signal_at(LB_C51_ACTIVE_BASE), + trainer.read_isv_signal_at(LB_C51_ACTIVE_BASE + 1), + trainer.read_isv_signal_at(LB_C51_ACTIVE_BASE + 2), + trainer.read_isv_signal_at(LB_C51_ACTIVE_BASE + 3), + ) + } else { + (0.0_f32, 0.0_f32, 0.0_f32, 0.0_f32, + 0.0_f32, 0.0_f32, 0.0_f32, 0.0_f32) + }; + tracing::info!( + "HEALTH_DIAG[{}]: lb_active_per_branch [cql:[d={:.4} m={:.4} o={:.4} u={:.4}] c51:[d={:.4} m={:.4} o={:.4} u={:.4}]]", + epoch, + cql_d, cql_m, cql_o, cql_u, + c51_d, c51_m, c51_o, c51_u, + ); + } + // C.2 Plan 3 Task 1 (spec §4.C.2): reward_split HEALTH_DIAG line. // Reads 6 ISV EMA slots updated by the GPU reward_component_ema kernel // launched just above. CPU-side code only reads; GPU wrote the values. diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 735e3c71b..0bccaf898 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -4033,5 +4033,38 @@ extended. No producer kernel yet — that arrives in the next commit. `mag_concat_qdir` is a per-sample register variable with no existing ISV slot; `h_s2_rms_ema` (ISV[96]) is the only available proxy. Option B blocked pending controller OK (no new ISV slots without explicit approval). +- SP7 observability: grad_decomp_pinned + lb_active_per_branch (2026-05-03, + wt/sp7-observability): two additive HEALTH_DIAG lines to disambiguate the + L40S smoke (`smoke-test-kl4lw`, HEAD `237b3dbfb`) showing all 8 (head, + branch) SP7 activation flags stuck at bootstrap across 5 epochs of fold 0 + despite non-zero Q-variance (`q_var_per_branch [0.0024, 0.0021, 0.0013, + 0.0021]`) and non-zero `grad_split_bwd [cql=6.48 c51=17.87]`. The ambiguity: + unknown whether `grad_decomp_result_pinned[0..12, 24..36, 36..48]` was + zero at the SP7 kernel's epoch-boundary call site (producer timing issue) + or whether the values were present but cold-start branch fired for another + reason. + + `grad_decomp_pinned` reads components IQN=0, CQL_SX=2, C51=3 from + `grad_component_norms_{mag,dir,trunk}` — the same arrays already populated + by `refresh_grad_component_norms` ~300 lines above the HEALTH_DIAG emit + block (no extra stream sync needed; buffer already up-to-date at read site). + Surfaces pinned offsets [0..12] (iqn), [24..36] (cql_sx), [36..48] (c51) + exactly as the SP7 kernel read them on the last step of the epoch. + + `lb_active_per_branch` reads ISV[LB_CQL_ACTIVE_BASE+0..4] = ISV[313..317) + and ISV[LB_C51_ACTIVE_BASE+0..4] = ISV[317..321) via the existing + `read_isv_signal_at` pattern. Monotonic per fold: 0.0 = cold-start branch + always fired, approaching 1.0 = active path executed ≥ once this fold. + + Interpretation matrix: (a) grad_decomp_pinned all-zero + lb_active all-zero + → pinned buffer empty at SP7 read time (producer timing/ordering defect); + (b) grad_decomp_pinned non-zero + lb_active all-zero → cold-start branch + firing for another reason (Welford-α cold-start or EPS_DIV guard); + (c) grad_decomp_pinned non-zero + lb_active approaching 1.0 → controller + active, budget stuck for a consumer-side reason. + + No new ISV slots. No kernel change. No StateResetRegistry entries. + Purely additive. Both lines placed adjacent to `q_var_per_branch` in the + per-epoch HEALTH_DIAG block (`training_loop.rs`). Contract test — every RegistryEntry has dispatch arm (2026-05-03): added `every_fold_and_soft_reset_entry_has_dispatch_arm` unit test to the existing `#[cfg(test)] mod tests` block in `state_reset_registry.rs`. Source-introspection design: `include_str!` both `state_reset_registry.rs` and `trainer/training_loop.rs`; brace-balance walk extracts the `match name {` body; for each FoldReset/SoftReset entry (97 total: 95 FoldReset + 2 SoftReset) asserts the literal `""` appears in that body. No production-code change. No new dependency (manual brace walk instead of regex). Catches the recurring "add RegistryEntry, forget dispatch arm → fold-boundary panic" bug (occurred twice: SP5 Layer A #281, SP7 T7 commit 6e479c55c) at `cargo test -p ml --lib` rather than mid-training. Test currently passes (all 97 dispatch arms present post SP7 T7 fix). Touched: `trainers/dqn/state_reset_registry.rs` (+59 LOC in cfg-test block). From 6c3a91c8781525bb4ba4f8e311abd75a5df6685d Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sun, 3 May 2026 14:55:40 +0200 Subject: [PATCH 5/5] fix(sp7): wire raw CQL norm kernel reading cql_grad_scratch directly MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The SP7 controller's CQL reference signal was reading cql_sx (post-SAXPY, budget-scaled) which created a self-perpetuating deadlock at small budget values: cql_sx_norm = budget × raw_grad → small budget → small cql_sx → controller can't update → budget stays small. The earlier offset 6 → 3 attempt failed because grad_decomp_launch_cql was never effectively populating slot 3 — the snapshot pattern measures ‖grad_buf − snapshot‖, but apply_cql_gradient writes to cql_grad_scratch (separate buffer), not grad_buf, so the snapshot delta is always 0. This commit adds a real producer kernel cql_raw_norm_compute that reads cql_grad_scratch directly and computes ‖raw_cql‖ over mag/dir/trunk slices. Wired to fire AFTER apply_cql_gradient and BEFORE apply_cql_saxpy, populating grad_decomp_result_pinned[3..6] with the raw norm independent of cql_budget. SP7 launcher updated to read from offset 3 (now: real raw CQL norm, not the never-populated cql snapshot delta). HEALTH_DIAG label renamed cql_sx → cql_raw to reflect the new contract; component index in the cached grad_component_norms_* arrays switched 2 → 1. The historical grad_decomp_launch_cql() call is removed — keeping it would overwrite the slot with 0 after cql_raw_norm_compute fires. The paired grad_decomp_snapshot_cql snapshot is left in place to scope the diff to Path A; buffer cleanup (grad_snapshot_cql allocation + grad_decomp_launch_cql definition) belongs in a follow-up commit per feedback_no_partial_refactor. Files: cql_raw_norm_kernel.cu (new, 97 LOC), build.rs, gpu_dqn_trainer.rs (struct field + cubin static + load + launcher + SP7 read offset 6→3), loss_balance_controller_kernel.cu (docstring + arg comment), fused_training.rs (4 launch_cql_raw_norm call sites, 1 dead grad_decomp_launch_cql call removed), training_loop.rs (HEALTH_DIAG label + index), audit doc Fix 31 SP7 Path A entry. Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/build.rs | 6 ++ .../src/cuda_pipeline/cql_raw_norm_kernel.cu | 97 +++++++++++++++++++ .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 90 ++++++++++++++++- .../loss_balance_controller_kernel.cu | 18 +++- crates/ml/src/trainers/dqn/fused_training.rs | 60 ++++++++++-- .../src/trainers/dqn/trainer/training_loop.rs | 22 +++-- docs/dqn-wire-up-audit.md | 54 +++++++++++ 7 files changed, 322 insertions(+), 25 deletions(-) create mode 100644 crates/ml/src/cuda_pipeline/cql_raw_norm_kernel.cu diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 10bd72314..1ed09c9b1 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -101,6 +101,12 @@ fn main() { "mamba2_temporal_kernel.cu", "graph_utility_kernels.cu", "grad_decomp_kernel.cu", + // SP7 Path A (2026-05-03): raw CQL norm reading `cql_grad_scratch` + // directly so the SP7 controller has a budget-independent CQL + // reference signal. Replaces the never-populated `cql` slot at + // element offset 3 in `grad_decomp_result_pinned`. See + // `cql_raw_norm_kernel.cu` and audit doc Fix 31 SP7 Path A. + "cql_raw_norm_kernel.cu", "branch_grad_balance_kernel.cu", "backtest_plan_kernel.cu", "tau_update_kernel.cu", diff --git a/crates/ml/src/cuda_pipeline/cql_raw_norm_kernel.cu b/crates/ml/src/cuda_pipeline/cql_raw_norm_kernel.cu new file mode 100644 index 000000000..a7e82c545 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/cql_raw_norm_kernel.cu @@ -0,0 +1,97 @@ +/** + * cql_raw_norm_kernel — SP7 Path A (2026-05-03) raw CQL gradient norm. + * + * Computes L2 norms of `cql_grad_scratch` over the same trunk/dir/mag + * slices `grad_decomp_kernel` uses, and writes three f32 results to a + * 3-float view inside the shared `grad_decomp_result_pinned` buffer: + * result_out[0] = ‖cql_scratch‖ on the magnitude branch slice + * result_out[1] = ‖cql_scratch‖ on the direction branch slice + * result_out[2] = ‖cql_scratch‖ on the trunk slice + * + * Why this exists. The SP7 loss-balance controller needs a CQL reference + * signal that scales with the *raw* CQL gradient magnitude — independent + * of the per-branch `cql_budget` the consumer applies via + * `apply_cql_saxpy`. Reading the post-SAXPY delta (`cql_sx` slot) + * creates a self-perpetuating deadlock: cql_sx_norm = budget × raw_grad, + * so when budget bootstraps small (≈0.02) the controller sees ≈0, + * cold-starts every step, and never updates the budget. + * + * The earlier offset-3 (`cql`) attempt failed because + * `grad_decomp_launch_cql` measures `‖grad_buf − snapshot‖`, but + * `apply_cql_gradient` writes to `cql_grad_scratch` (a separate buffer); + * the snapshot delta on `grad_buf` is therefore always 0. + * + * This kernel reads `cql_grad_scratch` *directly* (no snapshot pair), + * so the norm reflects the raw gradient magnitude before budget scaling + * — exactly what the SP7 controller needs as a reference signal. + * + * Layout: + * cql_grad_scratch — f32 device pointer, length total_params (matches + * `grad_buf`'s element layout). Same trunk_start/ + * dir_start/mag_start indices used by the + * grad_decomp pipeline apply unmodified. + * trunk_start/len, dir_start/len, mag_start/len + * — element offsets + lengths into `cql_grad_scratch` + * for each slice (identical to grad_decomp_kernel + * contract). + * result_out — [3] output: (mag_norm, dir_norm, trunk_norm). The + * caller computes the destination pointer as + * `grad_decomp_result_dev_ptr + 3 * sizeof(f32)`, + * replacing the never-populated `cql` slot at + * element offset 3 (was always zero — see audit + * doc Fix 31 SP7 Path A). + * + * All reductions are block-internal via shared-memory tree reduction — + * NO atomicAdd (`feedback_no_atomicadd`). + * + * Launch config: one block, 256 threads. + */ +#include + +extern "C" __global__ void cql_raw_norm_compute( + const float* __restrict__ cql_grad_scratch, + int trunk_start, int trunk_len, + int dir_start, int dir_len, + int mag_start, int mag_len, + float* __restrict__ result_out +) { + __shared__ float sum_mag[256]; + __shared__ float sum_dir[256]; + __shared__ float sum_trunk[256]; + int tid = threadIdx.x; + sum_mag[tid] = 0.0f; + sum_dir[tid] = 0.0f; + sum_trunk[tid] = 0.0f; + + // Trunk slice — read cql_grad_scratch[trunk_start..+trunk_len) directly. + for (int i = tid; i < trunk_len; i += blockDim.x) { + float v = cql_grad_scratch[trunk_start + i]; + sum_trunk[tid] += v * v; + } + // Direction slice — cql_grad_scratch[dir_start..+dir_len). + for (int i = tid; i < dir_len; i += blockDim.x) { + float v = cql_grad_scratch[dir_start + i]; + sum_dir[tid] += v * v; + } + // Magnitude slice — cql_grad_scratch[mag_start..+mag_len). + for (int i = tid; i < mag_len; i += blockDim.x) { + float v = cql_grad_scratch[mag_start + i]; + sum_mag[tid] += v * v; + } + __syncthreads(); + + // Block-level tree reduction (256 → 1). + for (int s = blockDim.x / 2; s > 0; s >>= 1) { + if (tid < s) { + sum_mag[tid] += sum_mag[tid + s]; + sum_dir[tid] += sum_dir[tid + s]; + sum_trunk[tid] += sum_trunk[tid + s]; + } + __syncthreads(); + } + if (tid == 0) { + result_out[0] = sqrtf(sum_mag[0]); + result_out[1] = sqrtf(sum_dir[0]); + result_out[2] = sqrtf(sum_trunk[0]); + } +} diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index b61844e31..2b4a7ffa3 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -91,6 +91,15 @@ static CQL_GRAD_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/cql_gra static MAMBA2_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/mamba2_temporal_kernel.cubin")); pub(crate) static GRAPH_UTILITY_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/graph_utility_kernels.cubin")); static GRAD_DECOMP_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/grad_decomp_kernel.cubin")); +/// SP7 Path A (2026-05-03): raw CQL gradient norm reading `cql_grad_scratch` +/// directly. Populates `grad_decomp_result_pinned[3..6]` (the previously +/// never-populated `cql` slot at element offset 3) with `‖raw_cql‖` over the +/// same trunk/dir/mag slices `grad_decomp_kernel` uses. The SP7 +/// loss-balance controller reads this slot as a budget-independent CQL +/// reference; the prior `cql_sx` reference at offset 6 was budget-scaled +/// (cql_sx_norm = budget × raw_grad), creating a self-perpetuating +/// deadlock at small bootstrap budgets. +static CQL_RAW_NORM_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/cql_raw_norm_kernel.cubin")); static BRANCH_GRAD_BALANCE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/branch_grad_balance_kernel.cubin")); /// Plan 1 Task 13: GPU-driven Polyak-EMA tau coefficient. static TAU_UPDATE_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/tau_update_kernel.cubin")); @@ -3220,6 +3229,13 @@ pub struct GpuDqnTrainer { grad_decomp_snapshot_len: usize, /// Reduction kernel (one-block 256-thread launch, no atomics). grad_decomp_kernel: CudaFunction, + /// SP7 Path A (2026-05-03) raw CQL norm kernel. Reads + /// `cql_grad_scratch` directly (no snapshot pair) and writes 3 floats + /// to `grad_decomp_result_pinned[3..6]` — the `cql` slot at element + /// offset 3 that was never populated by the (defined-but-unused) + /// `grad_decomp_launch_cql` snapshot pattern. Single block, 256 + /// threads, shared-memory tree reduction (no atomicAdd). + cql_raw_norm_kernel: CudaFunction, /// Adaptive per-branch gradient-norm balancer — caps any branch whose /// L2 norm exceeds `num_branches × median_branch_norm` (see /// `branch_grad_balance_kernel.cu`). `num_branches = 4` is the @@ -5867,6 +5883,52 @@ impl GpuDqnTrainer { self.launch_grad_decomp(&self.grad_snapshot_ens, 24, "ens") } + /// SP7 Path A (2026-05-03) — launch `cql_raw_norm_compute` reading + /// `cql_grad_scratch` directly. Writes `‖raw_cql‖` over (mag, dir, + /// trunk) slices to `grad_decomp_result_pinned[3..6]` — the `cql` + /// slot at element offset 3. + /// + /// Why a dedicated launcher rather than reusing `launch_grad_decomp` + /// with the `grad_decomp_launch_cql` snapshot: that pattern produces + /// `‖grad_buf − snapshot_cql‖`, which is structurally always 0 + /// because `apply_cql_gradient` writes to `cql_grad_scratch` (a + /// separate buffer), never to `grad_buf`. The previously-defined + /// `grad_decomp_launch_cql` is unwired (no call site); this launcher + /// supplants it with a real producer that the SP7 controller can + /// read at element offset 3 as a budget-independent reference signal. + /// + /// Must run AFTER `apply_cql_gradient` populates `cql_grad_scratch` + /// and BEFORE `apply_cql_saxpy` consumes it (the saxpy doesn't + /// modify `cql_grad_scratch`, but ordering keeps the contract + /// "raw norm reflects what gets SAXPYed this step" obvious). + pub(crate) fn launch_cql_raw_norm(&self) -> Result<(), MLError> { + let scratch_ptr = self.cql_grad_scratch.raw_ptr(); + let f32_size = std::mem::size_of::() as u64; + // Offset 3 (cql slot) — was never populated by the unwired + // `grad_decomp_launch_cql` snapshot pattern. Now produces the + // raw-CQL-norm SP7 needs. + let result_ptr = self.grad_decomp_result_dev_ptr + 3 * f32_size; + let cfg = LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (256, 1, 1), + shared_mem_bytes: 0, + }; + unsafe { + self.stream.launch_builder(&self.cql_raw_norm_kernel) + .arg(&scratch_ptr) + .arg(&self.grad_decomp_trunk_start) + .arg(&self.grad_decomp_trunk_len) + .arg(&self.grad_decomp_dir_start) + .arg(&self.grad_decomp_dir_len) + .arg(&self.grad_decomp_mag_start) + .arg(&self.grad_decomp_mag_len) + .arg(&result_ptr) + .launch(cfg) + .map_err(|e| MLError::ModelError(format!("cql_raw_norm_compute: {e}")))?; + } + Ok(()) + } + /// Task 2.0 — populate the host-side `grad_component_norms_mag/_dir/ /// _trunk` caches from the pinned result slot. Called at epoch /// boundary, before HEALTH_DIAG emission. @@ -11348,11 +11410,20 @@ impl GpuDqnTrainer { // grad_decomp pinned layout: 27 floats total, 9 components × 3 floats // each ([mag, dir, trunk]). Component element offsets per // launch_grad_decomp documentation: iqn=0, cql=3, cql_sx=6, c51=9. - // We use cql_sx (post-budget delta) for ratio parity with what landed - // in grad_buf. + // + // SP7 Path A (2026-05-03): the CQL reference now reads slot 3 (`cql`) + // populated by `launch_cql_raw_norm` — a real producer that reads + // `cql_grad_scratch` directly. Previously SP7 read slot 6 (`cql_sx`, + // post-budget delta = budget × raw_grad), which created a self- + // perpetuating deadlock at small bootstrap budgets: the controller + // saw ≈0 norm → cold-started every step → never updated the budget. + // Slot 3 is budget-independent — exactly the reference SP7 needs. + // The historical "cql at offset 3" was zero-valued (`grad_decomp_launch_cql` + // measured a delta on `grad_buf`, but `apply_cql_gradient` writes + // to `cql_grad_scratch`); the unwired path is now fully replaced. let f32_size = std::mem::size_of::() as u64; let iqn_dev = self.grad_decomp_result_dev_ptr + 0 * f32_size; - let cql_sx_dev = self.grad_decomp_result_dev_ptr + 6 * f32_size; + let cql_dev = self.grad_decomp_result_dev_ptr + 3 * f32_size; let c51_dev = self.grad_decomp_result_dev_ptr + 9 * f32_size; // Step 1: producer kernel. @@ -11379,7 +11450,7 @@ impl GpuDqnTrainer { self.stream .launch_builder(&self.loss_balance_controller_kernel) .arg(&iqn_dev) - .arg(&cql_sx_dev) + .arg(&cql_dev) .arg(&c51_dev) .arg(&isv_dev) .arg(&flatness_isv_base_i32) @@ -14067,6 +14138,16 @@ impl GpuDqnTrainer { module.load_function("grad_component_delta_norm") .map_err(|e| MLError::ModelError(format!("grad_component_delta_norm load: {e}")))? }; + // SP7 Path A (2026-05-03) — raw CQL norm kernel. Same shape as + // grad_decomp_kernel but reads `cql_grad_scratch` directly (no + // snapshot pair). Loaded into its own CUmodule for the same + // CUfunction-isolation reason. + let cql_raw_norm_kernel = { + let module = stream.context().load_cubin(CQL_RAW_NORM_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("cql_raw_norm cubin load: {e}")))?; + module.load_function("cql_raw_norm_compute") + .map_err(|e| MLError::ModelError(format!("cql_raw_norm_compute load: {e}")))? + }; info!( "GpuDqnTrainer: grad_decomp kernel loaded — snapshot_len={} trunk=[{}..+{}] dir=[{}..+{}] mag=[{}..+{}]", grad_decomp_snapshot_len, @@ -17391,6 +17472,7 @@ impl GpuDqnTrainer { grad_decomp_mag_len, grad_decomp_snapshot_len, grad_decomp_kernel, + cql_raw_norm_kernel, branch_grad_balance_reduce, branch_grad_balance_isv_update, branch_grad_balance_rescale, diff --git a/crates/ml/src/cuda_pipeline/loss_balance_controller_kernel.cu b/crates/ml/src/cuda_pipeline/loss_balance_controller_kernel.cu index c49ac7a15..fda09bc1c 100644 --- a/crates/ml/src/cuda_pipeline/loss_balance_controller_kernel.cu +++ b/crates/ml/src/cuda_pipeline/loss_balance_controller_kernel.cu @@ -3,7 +3,16 @@ // SP7 (2026-05-03): per-branch loss-balance controller for CQL and C51 budgets. // // Reads per-loss decomp pinned slots [mag_norm, dir_norm, trunk_norm] for -// IQN (reference), CQL_SX (managed, post-budget), C51 (managed, pre-budget). +// IQN (reference), CQL (managed, RAW pre-budget — Path A 2026-05-03), +// C51 (managed, pre-budget). +// +// Path A note: prior to 2026-05-03 the kernel read CQL_SX (post-SAXPY, +// budget-scaled = budget × raw_grad) at element offset 6, which created +// a self-perpetuating deadlock at small bootstrap budgets — the controller +// saw ≈0 norm → cold-started → never updated the budget. The new +// `cql_raw_norm_compute` kernel populates element offset 3 (the `cql` slot) +// directly from `cql_grad_scratch`, giving SP7 a budget-independent CQL +// reference. Launcher updated atomically to read offset 3. // Reads per-branch flatness from ISV[FLATNESS_BASE..+4]. Reads prior // budgets from ISV[BUDGET_{CQL,C51}_BASE..+4] and prior Wiener state from // ISV[LB_*_VAR_{CQL,C51}_BASE..+4]. Writes new budgets + new Wiener @@ -62,9 +71,10 @@ extern "C" __global__ void loss_balance_controller_update( // Pointers into the shared `grad_decomp_result_pinned` buffer at the // 3-float component offsets (the launcher computes these by adding // (offset_elems * sizeof(float)) to grad_decomp_result_dev_ptr): - // iqn_decomp = base + 0 (slot iqn, layout [mag, dir, trunk]) - // cql_decomp = base + 24 (slot cql_sx, layout [mag, dir, trunk]) - // c51_decomp = base + 36 (slot c51, layout [mag, dir, trunk]) + // iqn_decomp = base + 0 (slot iqn, layout [mag, dir, trunk]) + // cql_decomp = base + 12 (slot cql, Path A 2026-05-03 — raw norm, + // layout [mag, dir, trunk]) + // c51_decomp = base + 36 (slot c51, layout [mag, dir, trunk]) const float* __restrict__ iqn_decomp, const float* __restrict__ cql_decomp, const float* __restrict__ c51_decomp, diff --git a/crates/ml/src/trainers/dqn/fused_training.rs b/crates/ml/src/trainers/dqn/fused_training.rs index 2aef0d4a9..c230b2122 100644 --- a/crates/ml/src/trainers/dqn/fused_training.rs +++ b/crates/ml/src/trainers/dqn/fused_training.rs @@ -2381,16 +2381,40 @@ impl FusedTrainingCtx { let temporal_amp = 1.0_f32 + (meta_q_pred - 0.5).max(0.0) * 2.0; let f5_barrier_weight = 0.20_f32 * temporal_amp; - // Task 2.0 — snapshot BEFORE the CQL gradient path. The snapshot runs - // whether or not `has_cql()` fires so the post-CQL reduction below - // can run unconditionally (keeps the pinned result slot's CQL field - // populated every step — 0.0 on steps where CQL was skipped). + // SP7 Path A (2026-05-03): the historical `grad_decomp_launch_cql` + // (paired with this snapshot) measured `‖grad_buf − snapshot_cql‖` + // over the CQL gradient path — structurally always 0 because + // `apply_cql_gradient` writes into `cql_grad_scratch` (a separate + // buffer), never touching `grad_buf` until `apply_cql_saxpy` fires. + // The launch is replaced below by `launch_cql_raw_norm`, which reads + // `cql_grad_scratch` directly and writes a real raw-CQL norm to + // `grad_decomp_result_pinned[3..6]`. The SP7 loss-balance controller + // reads slot 3 as a budget-independent CQL reference signal (was + // deadlocking on the budget-scaled `cql_sx` slot at offset 6). + // + // The snapshot itself is now functionally dead (nothing reads its + // output) but is left in place to preserve the `grad_snapshot_cql` + // buffer + `grad_decomp_snapshot_cql()` call site semantics — + // removing the buffer touches the constructor, struct definition, + // and 9 sibling `grad_snapshot_*` references; that cleanup belongs + // in a follow-up commit per `feedback_no_partial_refactor` (this + // commit is focused on Path A wire-up). self.trainer.grad_decomp_snapshot_cql() .map_err(|e| anyhow::anyhow!("Task 2.0 grad_decomp_snapshot_cql: {e}"))?; if self.trainer.has_cql() { match self.trainer.apply_cql_gradient(f5_barrier_weight) { Ok(true) => { + // SP7 Path A (2026-05-03) — populate `cql` slot at offset 3 + // with the raw CQL gradient norm read directly from + // `cql_grad_scratch`. Must run AFTER `apply_cql_gradient` + // populated the scratch and BEFORE `apply_cql_saxpy` + // consumes it. Replaces the unwired + // `grad_decomp_launch_cql` snapshot path (which always + // reported 0 because the snapshot was on `grad_buf` and + // CQL writes to `cql_grad_scratch`). + self.trainer.launch_cql_raw_norm() + .map_err(|e| anyhow::anyhow!("SP7 Path A cql_raw_norm: {e}"))?; // Task 2.0 extension — isolate `apply_cql_saxpy`. The // preceding `apply_cql_gradient` only writes into // `cql_grad_scratch`, not `grad_buf`, so the cql_sx @@ -2413,6 +2437,13 @@ impl FusedTrainingCtx { } Ok(false) => { // Keep cql_sx slot populated on skip path — delta = 0. + // SP7 Path A: also keep `cql` slot at offset 3 populated. + // `cql_grad_scratch` was zeroed at the top of + // `apply_cql_gradient` (`cuMemsetD32Async`) before the + // early-return, so the raw-norm kernel reads all-zeros + // and writes 0.0 to the slot — honest semantic. + self.trainer.launch_cql_raw_norm() + .map_err(|e| anyhow::anyhow!("SP7 Path A cql_raw_norm (skip): {e}"))?; self.trainer.grad_decomp_snapshot_cql_sx() .map_err(|e| anyhow::anyhow!("Task 2.0 grad_decomp_snapshot_cql_sx (skip): {e}"))?; self.trainer.grad_decomp_launch_cql_sx() @@ -2420,7 +2451,12 @@ impl FusedTrainingCtx { } Err(e) => { tracing::warn!("CQL gradient failed (non-fatal): {e}"); - // Keep cql_sx slot populated on error path — delta = 0. + // Keep both slots populated on error path. The scratch + // may be in an indeterminate state after the partial + // backward; the raw-norm kernel still produces a + // numerically-finite L2 norm, which is what SP7 needs. + self.trainer.launch_cql_raw_norm() + .map_err(|e| anyhow::anyhow!("SP7 Path A cql_raw_norm (err): {e}"))?; self.trainer.grad_decomp_snapshot_cql_sx() .map_err(|e| anyhow::anyhow!("Task 2.0 grad_decomp_snapshot_cql_sx (err): {e}"))?; self.trainer.grad_decomp_launch_cql_sx() @@ -2429,16 +2465,22 @@ impl FusedTrainingCtx { } } else { // Keep cql_sx slot populated on CQL-disabled path — delta = 0. + // SP7 Path A: scratch is left at its prior contents (no + // `apply_cql_gradient` call to zero it), so the raw norm reads + // whatever survived from the last step where CQL fired. In the + // CQL-permanently-disabled regime the buffer is still zero from + // construction, so the kernel writes 0.0. In a transient-disable + // regime the slot stays at the most-recent active value, which + // is the appropriate "carry forward" semantic for a controller + // reading a gradient-magnitude reference. + self.trainer.launch_cql_raw_norm() + .map_err(|e| anyhow::anyhow!("SP7 Path A cql_raw_norm (no-cql): {e}"))?; self.trainer.grad_decomp_snapshot_cql_sx() .map_err(|e| anyhow::anyhow!("Task 2.0 grad_decomp_snapshot_cql_sx (no-cql): {e}"))?; self.trainer.grad_decomp_launch_cql_sx() .map_err(|e| anyhow::anyhow!("Task 2.0 grad_decomp cql_sx (no-cql): {e}"))?; } - // Task 2.0 — CQL reduction. Delta = grad_buf − snapshot over branch 0+1. - self.trainer.grad_decomp_launch_cql() - .map_err(|e| anyhow::anyhow!("Task 2.0 grad_decomp CQL: {e}"))?; - // D1/N1: Temporal self-distillation — per-step kernel pull toward the // best historical snapshot. Fully GPU-native: alpha is computed // in-kernel from ISV[LEARNING_HEALTH_INDEX] (pinned device-mapped), diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index d222b720b..7858bb35c 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -4101,7 +4101,7 @@ impl DQNTrainer { } // SP7 observability (grad_decomp_pinned): surface the exact contents of - // grad_decomp_result_pinned[0..12, 24..36, 36..48] (iqn/cql_sx/c51 × + // grad_decomp_result_pinned[0..3, 12..15, 36..48] (iqn/cql_raw/c51 × // [mag, dir, trunk]) that the SP7 loss-balance controller read in the // last step of this epoch. The values are already cached in the // DQNTrainer host arrays by `refresh_grad_component_norms` (called ~300 @@ -4109,23 +4109,29 @@ impl DQNTrainer { // zero extra synchronisation. // // Component index mapping (matches `grad_decomp_result_pinned` layout): - // IQN = component 0 → pinned offsets [0, 1, 2] - // CQL_SX = component 2 → pinned offsets [6, 7, 8] - // C51 = component 3 → pinned offsets [9, 10, 11] + // IQN = component 0 → pinned offsets [0, 1, 2] + // CQL_RAW = component 1 → pinned offsets [3, 4, 5] (Path A 2026-05-03) + // C51 = component 3 → pinned offsets [9, 10, 11] + // + // SP7 Path A (2026-05-03): label changed `cql_sx` → `cql_raw` to + // reflect the new producer. `cql_sx` (component 2, post-SAXPY + // budget-scaled delta on `grad_buf`) is no longer the SP7 reference; + // `cql_raw` (component 1, raw norm of `cql_grad_scratch` written by + // `cql_raw_norm_compute`) is — see audit doc Fix 31 SP7 Path A. // // This disambiguates "grad_decomp not populated at SP7 read time" // (all zeros here) from "values non-zero but cold-start branch fires // for another reason" (non-zeros here, lb_active stays 0.0). { let (iqn_m, iqn_d, iqn_t, - cql_sx_m, cql_sx_d, cql_sx_t, + cql_raw_m, cql_raw_d, cql_raw_t, c51_m, c51_d, c51_t) = if let Some(ref fused) = self.fused_ctx { let trainer = fused.trainer(); let mag = trainer.grad_component_norms_mag(); let dir = trainer.grad_component_norms_dir(); let trunk = trainer.grad_component_norms_trunk(); (mag[0], dir[0], trunk[0], // IQN - mag[2], dir[2], trunk[2], // CQL_SX + mag[1], dir[1], trunk[1], // CQL_RAW (Path A — was [2] CQL_SX) mag[3], dir[3], trunk[3]) // C51 } else { (0.0_f32, 0.0_f32, 0.0_f32, @@ -4133,10 +4139,10 @@ impl DQNTrainer { 0.0_f32, 0.0_f32, 0.0_f32) }; tracing::info!( - "HEALTH_DIAG[{}]: grad_decomp_pinned [iqn:[m={:.4} d={:.4} t={:.4}] cql_sx:[m={:.4} d={:.4} t={:.4}] c51:[m={:.4} d={:.4} t={:.4}]]", + "HEALTH_DIAG[{}]: grad_decomp_pinned [iqn:[m={:.4} d={:.4} t={:.4}] cql_raw:[m={:.4} d={:.4} t={:.4}] c51:[m={:.4} d={:.4} t={:.4}]]", epoch, iqn_m, iqn_d, iqn_t, - cql_sx_m, cql_sx_d, cql_sx_t, + cql_raw_m, cql_raw_d, cql_raw_t, c51_m, c51_d, c51_t, ); } diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 0bccaf898..3d64bbb45 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -4066,5 +4066,59 @@ extended. No producer kernel yet — that arrives in the next commit. No new ISV slots. No kernel change. No StateResetRegistry entries. Purely additive. Both lines placed adjacent to `q_var_per_branch` in the per-epoch HEALTH_DIAG block (`training_loop.rs`). +- SP7 Path A: raw CQL norm kernel (2026-05-03, wt/sp7-cql-raw-kernel): + fixes a self-perpetuating deadlock in the SP7 controller's CQL reference + signal. Pre-fix: SP7 read `cql_decomp` from `grad_decomp_result_pinned` + at element offset 6 (the `cql_sx` slot, populated by + `grad_decomp_launch_cql_sx`). `cql_sx_norm` measures + `‖grad_buf_after_apply_cql_saxpy − grad_buf_before‖` = + `‖cql_budget × raw_cql_grad‖` = `cql_budget × ‖raw_cql_grad‖`. When + bootstrap `cql_budget ≈ 0.02` (`COLD_START_FLOOR_CQL` = + `CQL_BOOTSTRAP_BUDGET`), `cql_sx_norm ≈ 0`, the controller's + `h_n < EPS_DIV` cold-start guard fires every step, and the budget never + updates from its bootstrap value. Empirically observed on + `smoke-test-kl4lw` (HEAD `237b3dbfb`): all 4 CQL active flags stuck at + bootstrap across 5 epochs of fold 0 despite non-zero + `grad_split_bwd cql=6.48`. + + An earlier offset-3 attempt (read the `cql` slot instead of `cql_sx`) + was abandoned because: (a) `grad_decomp_launch_cql` measures + `‖grad_buf − snapshot_cql‖`, but `apply_cql_gradient` writes only into + `cql_grad_scratch` (a separate buffer) — the delta on `grad_buf` is + always 0; (b) the existing snapshot pattern can't be repointed at + `cql_grad_scratch` because the kernel takes a single `current` pointer + + `snapshot` pointer, designed around the `grad_buf` layout. + + Path A fix: new GPU kernel `cql_raw_norm_compute` reads + `cql_grad_scratch` directly and writes a 3-float `[mag, dir, trunk]` + L2-norm tuple to `grad_decomp_result_pinned[3..6]` — replacing the + always-zero `cql` slot with a real raw-CQL norm. Same 256-thread + single-block shared-memory tree-reduction shape as `grad_decomp_kernel` + (no atomicAdd per `feedback_no_atomicadd`), reusing the trainer's + `grad_decomp_trunk_start/_dir_start/_mag_start` slice indices for + layout parity. Wired AFTER `apply_cql_gradient` populates the scratch + and BEFORE `apply_cql_saxpy` consumes it in `fused_training.rs:2391-2470`, + unconditionally (active/skip/error/no-CQL paths) so the slot stays + populated every step (was the prior contract via the now-removed + `grad_decomp_launch_cql` call). SP7 launcher + (`launch_loss_balance_controller`) updated to read element offset 3 + (`cql_dev`); `loss_balance_controller_kernel.cu` docstring + + `cql_decomp` arg comment updated for the new contract. HEALTH_DIAG + `grad_decomp_pinned` line label renamed `cql_sx` → `cql_raw` and + switched from component index 2 → 1 in the cached + `grad_component_norms_*` arrays. The historical + `grad_decomp_launch_cql()` call is removed (would otherwise overwrite + the slot with 0 after `cql_raw_norm_compute` fires); the paired + `grad_decomp_snapshot_cql` snapshot is left in place to keep the + diff scoped to Path A — buffer cleanup belongs in a follow-up commit. + + Files: `cuda_pipeline/cql_raw_norm_kernel.cu` (new, 91 LOC), + `build.rs` (+5 LOC), `cuda_pipeline/gpu_dqn_trainer.rs` (struct field + + cubin static + load + launcher; +60 LOC), `cuda_pipeline/loss_balance_controller_kernel.cu` + (docstring + arg comment, +10 LOC), `trainers/dqn/fused_training.rs` + (3 new `launch_cql_raw_norm` call sites in active/skip/error/no-CQL + branches + comment block, +30 LOC, −3 LOC removing the dead + `grad_decomp_launch_cql` call), `trainers/dqn/trainer/training_loop.rs` + (HEALTH_DIAG label + index, +6 LOC, −5 LOC), this audit entry. Contract test — every RegistryEntry has dispatch arm (2026-05-03): added `every_fold_and_soft_reset_entry_has_dispatch_arm` unit test to the existing `#[cfg(test)] mod tests` block in `state_reset_registry.rs`. Source-introspection design: `include_str!` both `state_reset_registry.rs` and `trainer/training_loop.rs`; brace-balance walk extracts the `match name {` body; for each FoldReset/SoftReset entry (97 total: 95 FoldReset + 2 SoftReset) asserts the literal `""` appears in that body. No production-code change. No new dependency (manual brace walk instead of regex). Catches the recurring "add RegistryEntry, forget dispatch arm → fold-boundary panic" bug (occurred twice: SP5 Layer A #281, SP7 T7 commit 6e479c55c) at `cargo test -p ml --lib` rather than mid-training. Test currently passes (all 97 dispatch arms present post SP7 T7 fix). Touched: `trainers/dqn/state_reset_registry.rs` (+59 LOC in cfg-test block).