Merge: sp7 complete chain — controller + observability + ISV bus fix + raw CQL norm
This commit is contained in:
jgrusewski
@@ -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",
|
||||
|
||||
97
crates/ml/src/cuda_pipeline/cql_raw_norm_kernel.cu
Normal file
97
crates/ml/src/cuda_pipeline/cql_raw_norm_kernel.cu
Normal file
@@ -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 <cuda_runtime.h>
|
||||
|
||||
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]);
|
||||
}
|
||||
}
|
||||
@@ -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"));
|
||||
@@ -825,7 +834,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).
|
||||
@@ -1100,8 +1109,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 +1183,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).
|
||||
@@ -1688,7 +1708,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;\
|
||||
@@ -3205,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
|
||||
@@ -5852,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::<f32>() 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.
|
||||
@@ -11318,6 +11395,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,
|
||||
@@ -11332,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::<f32>() 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.
|
||||
@@ -11347,18 +11434,23 @@ 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
|
||||
.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)
|
||||
@@ -11368,6 +11460,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 +11470,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 +11480,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 +11492,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;
|
||||
@@ -14038,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,
|
||||
@@ -17362,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,
|
||||
|
||||
@@ -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
|
||||
@@ -20,15 +29,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.
|
||||
@@ -39,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,
|
||||
@@ -54,6 +87,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 +98,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 +113,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 +147,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();
|
||||
}
|
||||
|
||||
@@ -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<usize> = (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<usize> = (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,44 @@ 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);
|
||||
}
|
||||
|
||||
#[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,
|
||||
);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -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,
|
||||
};
|
||||
@@ -2380,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
|
||||
@@ -2412,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()
|
||||
@@ -2419,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()
|
||||
@@ -2428,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),
|
||||
@@ -3370,10 +3413,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 +3437,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 +3457,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 +3497,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
|
||||
|
||||
@@ -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.
|
||||
|
||||
@@ -4072,6 +4072,121 @@ 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,
|
||||
);
|
||||
}
|
||||
|
||||
// SP7 observability (grad_decomp_pinned): surface the exact contents of
|
||||
// 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
|
||||
// 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_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_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[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,
|
||||
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_raw:[m={:.4} d={:.4} t={:.4}] c51:[m={:.4} d={:.4} t={:.4}]]",
|
||||
epoch,
|
||||
iqn_m, iqn_d, iqn_t,
|
||||
cql_raw_m, cql_raw_d, cql_raw_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.
|
||||
@@ -6570,6 +6685,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
|
||||
|
||||
@@ -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<CudaStream>) -> 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<CudaStream>) -> 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::<f32>() 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]
|
||||
);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -3935,9 +3935,190 @@ 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.
|
||||
- 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.
|
||||
- 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).
|
||||
- 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`).
|
||||
- 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 `"<name>"` 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).
|
||||
|
||||
Reference in New Issue
Block a user