fix(sp4): Layer B fix-up — restore Adam coverage for tensors [33..163)
Layer B's 3-way DQN main Adam split (commit58ffb3a48) covered only tensors [0..33) (DqnTrunk + DqnValue + DqnBranches per param_group_buffers mapping). Tensors [33..163) — bottleneck, VSN bottleneck, GLU, KAN, regime gate, spacing, multi-horizon value heads, risk, ISV encoder, VSN per-group, aux heads, MoE — were silently no longer Adam-updated. Their grads still computed, m/v state still allocated, but no parameter step applied. Detected by code-quality reviewer; would have surfaced at L40S smoke as ~70% network freeze attributable to Sharpe noise. Fixes: 1. 4th sub-launch added to launch_adam_update: trunk-extras tail covering tensors [33..163), tagged ParamGroup::DqnTrunk for shared bounds, with SP4_ENGAGE_OFFSET_DISABLED for Pearl C engagement (rolls into trunk's accounting since they share WEIGHT_BOUND/ WD_RATE). 2. MAX_BLOCKS_PER_ADAM grown 256 → 4096 to accommodate trunk-extras's block count (production cfg ~2400 blocks). SP4_ENGAGE_BUF_LEN auto-updates: 11 × 4096 = 45056 ints. All 4 Curiosity-sub-launch offset constants auto-derive from MAX_BLOCKS_PER_ADAM. 3. Coverage debug_assert_eq added: trunk + value + branches + trunk_extras must equal total_params. Prevents future regressions. 4. read_group_adam_bounds(group) -> (clamp, wd) helper introduced on GpuDqnTrainer; 6 verbose 2-line ISV-read patterns in fused_training.rs collapsed to one-liners (TLOB + IQL high + IQL low + IQN parallel + Attn parallel + Attn sequential + IQN sequential = 7 call sites). The 4-way Adam loop in launch_adam_update also uses the helper. 5. 9 stale doc-comments referring to "100 × Q_ABS_REF.max(1.0)" updated to current "ISV[WEIGHT_BOUND[group]].max(EPS_CLAMP_FLOOR)" contract (curiosity_training_kernel.cu, gpu_curiosity_trainer.rs ×2, iqn_dual_head_kernel.cu ×3, iql_value_kernel.cu, attention_backward_kernel.cu, dqn_utility_kernels.cu ×2, gpu_dqn_trainer.rs historical SP3 reference). 6. launch_adam_update + pearl_c_post_adam_engagement_check + audit doc updated to reflect actual 4-way coverage and trunk-extras's engagement folding into DqnTrunk. Per feedback_no_partial_refactor: trunk-extras consumers (= no dedicated SP4 producer yet) and bound (= reused trunk's WEIGHT_BOUND/WD_RATE) ship together with the launch fix in this commit. A future task may introduce dedicated SP4 producers for sub-trunk regions (VSN, MoE, etc.) once warranted by signal analysis. Refs:58ffb3a48(Layer B atomic consumer migration). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -77,10 +77,10 @@ extern "C" __global__ void attn_adam_kernel(
|
||||
float weight_decay, float max_grad_norm,
|
||||
const int* __restrict__ adam_t_buf,
|
||||
int total_params,
|
||||
/* SP3 Mech 9 (close-out v2): ISV-driven |p| bound (100 × Q_ABS_REF.max(1.0))
|
||||
/* SP4 Layer B: ISV-driven |p| bound (`ISV[WEIGHT_BOUND[Attn]].max(EPS_CLAMP_FLOOR)`)
|
||||
* — same pattern wired into all five Adam kernels. Used by both
|
||||
* GpuAttention and GpuTlob (TLOB shares this kernel via the same cubin).
|
||||
* 0=disabled (debug only). */
|
||||
* GpuAttention and GpuTlob (TLOB shares this kernel via the same cubin
|
||||
* and co-lives in the Attn group taxonomy). 0=disabled (debug only). */
|
||||
float weight_clamp_max_abs,
|
||||
/* SP4 Task A14 (2026-04-30): Pearl C engagement counter via warp+
|
||||
* block tree-reduce (no atomicAdd). GpuAttention + GpuTlob both pass
|
||||
|
||||
@@ -321,7 +321,7 @@ extern "C" __global__ void curiosity_adam_step(
|
||||
float beta2,
|
||||
float eps,
|
||||
int step, /* 1-based step counter */
|
||||
/* SP3 Mech 9 (close-out v2): ISV-driven |p| bound (100 × Q_ABS_REF.max(1.0))
|
||||
/* SP4 Layer B: ISV-driven |p| bound (`ISV[WEIGHT_BOUND[group]].max(EPS_CLAMP_FLOOR)`)
|
||||
* — same pattern wired into all five Adam kernels. 0=disabled (debug only). */
|
||||
float weight_clamp_max_abs,
|
||||
/* SP4 Task A14 (2026-04-30): Pearl C engagement counter via warp+
|
||||
|
||||
@@ -142,7 +142,7 @@ extern "C" __global__ void dqn_adam_update_kernel(
|
||||
const float* __restrict__ weight_decay_mask, /* [TOTAL_PARAMS] 1.0=decay, 0.0=no decay */
|
||||
int l1_end, /* param index where L1 stops (end of w_s1) */
|
||||
float l1_lambda, /* L1 penalty (1e-3). 0=disabled */
|
||||
float weight_clamp_max_abs, /* SP3 Mech 9: ISV-driven |p_val| bound (100 * Q_ABS_REF.max(1.0)). 0=disabled. */
|
||||
float weight_clamp_max_abs, /* SP4 Layer B: ISV-driven |p_val| bound (`ISV[WEIGHT_BOUND[group]].max(EPS_CLAMP_FLOOR)`). 0=disabled. */
|
||||
/* SP4 Task A14: Pearl C engagement counter. */
|
||||
int* __restrict__ nan_flags_buf, /* sticky overflow flag store; nan_flags_buf[diag_slot]=1 on first thread that overshoots */
|
||||
int diag_slot, /* index into nan_flags_buf (one per Adam kernel) */
|
||||
@@ -196,13 +196,15 @@ extern "C" __global__ void dqn_adam_update_kernel(
|
||||
p_val = sign_val * fmaxf(fabsf(p_val) - lr * l1_lambda, 0.0f);
|
||||
}
|
||||
|
||||
/* SP3 Mech 9 + SP4 Task A14: post-Adam weight clamp + Pearl C
|
||||
* engagement counter. ISV-driven bound (100 × Q_ABS_REF.max(1.0))
|
||||
* computed host-side and passed in. Targets cuBLAS sgemm f32
|
||||
* accumulator overflow at slots 26+32 — bounds weights 1 OOM below
|
||||
* slot 44-45 diagnostic threshold (1e3 × Q_ABS_REF) so the
|
||||
* diagnostic retains regression-sentinel headroom. Disabled when
|
||||
* bound==0 (debug fast-path; production never sets 0).
|
||||
/* SP4 Layer B + SP4 Task A14: post-Adam weight clamp + Pearl C
|
||||
* engagement counter. ISV-driven bound
|
||||
* (`ISV[WEIGHT_BOUND[group]].max(EPS_CLAMP_FLOOR)`) computed host-side
|
||||
* via `read_group_adam_bounds` and passed in. Targets cuBLAS
|
||||
* sgemm f32 accumulator overflow at slots 26+32 — Wiener-EMA over
|
||||
* p99(|params|) keeps the bound just above legitimate weight
|
||||
* magnitudes; the diagnostic retains regression-sentinel
|
||||
* headroom. Disabled when bound==0 (debug fast-path; production
|
||||
* never sets 0).
|
||||
*
|
||||
* `local_engage = 1` when |p_val| would overshoot — captured into
|
||||
* the per-block tree-reduce below. The sticky `nan_flags_buf` write
|
||||
|
||||
@@ -505,10 +505,11 @@ pub struct GpuCuriosityTrainer {
|
||||
|
||||
/// Launch Adam optimizer step for one parameter group.
|
||||
///
|
||||
/// `weight_clamp_max_abs` is SP3 Mech 9 (close-out v2): ISV-driven |p|
|
||||
/// bound for `curiosity_adam_step`. Caller threads `100 × Q_ABS_REF.max(1.0)`
|
||||
/// from training_loop down through `train_on_collector_buffers`. 0=disabled
|
||||
/// (debug only).
|
||||
/// `weight_clamp_max_abs` is SP4 Layer B: ISV-driven |p| bound for
|
||||
/// `curiosity_adam_step`. Caller threads
|
||||
/// `ISV[WEIGHT_BOUND[ParamGroup::Curiosity]].max(EPS_CLAMP_FLOOR)`
|
||||
/// from training_loop down through `train_on_collector_buffers`.
|
||||
/// 0=disabled (debug only).
|
||||
#[allow(clippy::too_many_arguments)]
|
||||
fn launch_adam_step(
|
||||
stream: &CudaStream,
|
||||
@@ -753,9 +754,10 @@ impl GpuCuriosityTrainer {
|
||||
states: &CudaSlice<f32>,
|
||||
actions: &CudaSlice<i32>,
|
||||
n_samples: usize,
|
||||
/* SP3 Mech 9 (close-out v2): ISV-driven |p| bound for the four
|
||||
/* SP4 Layer B: ISV-driven |p| bound for the four
|
||||
* `curiosity_adam_step` launches (W1, b1, W2, b2). Caller computes
|
||||
* `100 × ISV[Q_ABS_REF].max(1.0)` host-side. 0=disabled (debug only). */
|
||||
* `ISV[WEIGHT_BOUND[Curiosity]].max(EPS_CLAMP_FLOOR)` host-side.
|
||||
* 0=disabled (debug only). */
|
||||
weight_clamp_max_abs: f32,
|
||||
) -> Result<(), MLError> {
|
||||
if n_samples < 2 {
|
||||
|
||||
@@ -15228,6 +15228,36 @@ impl GpuDqnTrainer {
|
||||
unsafe { *self.isv_signals_pinned.add(index) }
|
||||
}
|
||||
|
||||
/// SP4 Layer B helper: read both Adam bounds
|
||||
/// (`weight_clamp_max_abs`, `weight_decay_value`) for a given param
|
||||
/// group from ISV slots in one call. Replaces the verbose 2-line
|
||||
/// `read_isv_signal_at(weight_bound(...)).max(EPS_CLAMP_FLOOR);
|
||||
/// read_isv_signal_at(wd_rate(...)).max(0.0)` pattern repeated at
|
||||
/// every per-group Adam launcher (DQN main, IQN, IQL high+low,
|
||||
/// Attention, TLOB).
|
||||
///
|
||||
/// `weight_clamp_max_abs` enforces the SP1 cold-start ε floor
|
||||
/// (`EPS_CLAMP_FLOOR = 1.0`) on the post-Adam |p| bound. The
|
||||
/// returned value is always ≥ EPS_CLAMP_FLOOR; cold-start ISV reads
|
||||
/// (which can be 0.0 before the producer has run) are saturated to
|
||||
/// the floor so the kernel never silently disables clamping.
|
||||
///
|
||||
/// `weight_decay_value` enforces the AdamW non-negativity contract
|
||||
/// (`max(0.0)`) — the kernel applies `wd * p` decoupled from Adam
|
||||
/// state, so a negative wd would create a positive feedback loop.
|
||||
/// Returns `0.0` for cold-start.
|
||||
#[inline]
|
||||
pub fn read_group_adam_bounds(
|
||||
&self,
|
||||
group: crate::cuda_pipeline::sp4_isv_slots::ParamGroup,
|
||||
) -> (f32, f32) {
|
||||
use crate::cuda_pipeline::sp4_isv_slots::{weight_bound, wd_rate};
|
||||
use crate::cuda_pipeline::sp4_wiener_ema::EPS_CLAMP_FLOOR;
|
||||
let clamp = self.read_isv_signal_at(weight_bound(group.idx())).max(EPS_CLAMP_FLOOR);
|
||||
let wd = self.read_isv_signal_at(wd_rate(group.idx())).max(0.0);
|
||||
(clamp, wd)
|
||||
}
|
||||
|
||||
/// Read the ISV layout fingerprint from pinned memory and compare to the
|
||||
/// compile-time `LAYOUT_FINGERPRINT_CURRENT`.
|
||||
///
|
||||
@@ -21092,30 +21122,50 @@ impl GpuDqnTrainer {
|
||||
/// `t_ptr` is a device buffer (not scalar) so the CUDA Graph can be
|
||||
/// replayed with an updated step counter.
|
||||
///
|
||||
/// **SP4 Layer B (2026-05-01)**: split into 3 sub-launches per group
|
||||
/// (DqnTrunk / DqnValue / DqnBranches). Each sub-launch reads its own
|
||||
/// `ISV[WEIGHT_BOUND[group]]` clamp + `ISV[WD_RATE[group]]` AdamW rate.
|
||||
/// L1 lambda (`ISV[L1_LAMBDA_TRUNK_INDEX]`) applies to the trunk only;
|
||||
/// value/branches pass `l1_end=0 + l1_lambda=0`. Per-group split
|
||||
/// resolves the Layer A Pearl C engagement-counter accounting deferral
|
||||
/// — each sub-launch writes per-block counts at its own group offset.
|
||||
/// **SP4 Layer B (2026-05-01)**: split into 4 sub-launches over the
|
||||
/// 3-group SP4 taxonomy + a "trunk-extras" tail. The first 3
|
||||
/// (DqnTrunk / DqnValue / DqnBranches) cover tensors `[0..33)` with
|
||||
/// per-group ISV-driven `WEIGHT_BOUND[group]` clamp and `WD_RATE[group]`
|
||||
/// AdamW rate. L1 lambda (`ISV[L1_LAMBDA_TRUNK_INDEX]`) applies to
|
||||
/// the trunk only; value/branches pass `l1_end=0 + l1_lambda=0`.
|
||||
/// Per-group split also resolves the Layer A Pearl C engagement-
|
||||
/// counter accounting deferral — each sub-launch writes per-block
|
||||
/// counts at its own group offset.
|
||||
///
|
||||
/// **Layer B fix-up (2026-05-01)**: the 4th sub-launch covers the
|
||||
/// tail tensors `[33..163)` (bottleneck, VSN bottleneck, GLU gates,
|
||||
/// KAN, regime gate, spacing, multi-horizon value heads, risk, ISV
|
||||
/// encoder, VSN per-group MLP, aux heads, MoE) — collectively
|
||||
/// "trunk-extras". The original 3-way split silently no-op'd Adam
|
||||
/// for these ~70% of weight tensors (grads computed, m/v allocated,
|
||||
/// no parameter step). The fix-up co-locates trunk-extras under
|
||||
/// `ParamGroup::DqnTrunk` for the WEIGHT_BOUND / WD_RATE contract
|
||||
/// (these tensors share the trunk's gradient regime: pre-bottleneck,
|
||||
/// attention, encoder paths) and uses `SP4_ENGAGE_OFFSET_DISABLED`
|
||||
/// for Pearl C engagement (rolls into trunk's accounting since they
|
||||
/// share bounds). A future task may introduce dedicated SP4
|
||||
/// producers for sub-trunk regions (VSN, MoE, etc.) once warranted
|
||||
/// by signal analysis. Per `feedback_no_partial_refactor`: trunk-
|
||||
/// extras' single (pragmatic) consumer ships with the launch fix in
|
||||
/// this commit.
|
||||
///
|
||||
/// All parameters `[0..total_params)` are covered across the 4
|
||||
/// sub-launches:
|
||||
/// - trunk = `params[0..trunk_param_count]` (tensors [0..13))
|
||||
/// - value = byte-slice for tensors [13..17)
|
||||
/// - branches = byte-slice for tensors [17..33)
|
||||
/// - trunk-tail = byte-slice for tensors [33..163)
|
||||
/// A coverage debug_assert at the bottom of the function verifies
|
||||
/// `trunk + value + branches + trunk_extras == total_params` to
|
||||
/// prevent regressions of this kind.
|
||||
///
|
||||
/// Per `feedback_no_quickfixes`: NOT a max/min-of-3 single-launch
|
||||
/// shortcut. Per `feedback_no_partial_refactor`: every consumer of
|
||||
/// the SP4 contract migrates in this same commit.
|
||||
///
|
||||
/// All parameters `[0..total_params)` are covered across the 3
|
||||
/// sub-launches (trunk = `params[0..trunk_param_count]`, value =
|
||||
/// indices [13..17), branches = indices [17..33)). VSN tensors at
|
||||
/// indices `[95..119)` fall inside the trunk slice; gradient noise
|
||||
/// from aux paths is attenuated upstream by the in-place
|
||||
/// `VSN_DW_DAMP` scaler in `aux_bottleneck_vsn_backward_dispatch`
|
||||
/// (1B-iv-rc) before it reaches `grad_buf[95..119)` — no Adam-state
|
||||
/// isolation is required.
|
||||
fn launch_adam_update(&self) -> Result<(), MLError> {
|
||||
use crate::cuda_pipeline::sp4_isv_slots::{
|
||||
weight_bound, wd_rate, ParamGroup as SP4Group, L1_LAMBDA_TRUNK_INDEX,
|
||||
ParamGroup as SP4Group, L1_LAMBDA_TRUNK_INDEX, SP4_ENGAGE_OFFSET_DISABLED,
|
||||
};
|
||||
use crate::cuda_pipeline::sp4_wiener_ema::EPS_CLAMP_FLOOR;
|
||||
|
||||
let lr_ptr = self.lr_dev_ptr; // pinned device-mapped — graph-safe, value read at replay
|
||||
let beta1 = self.config.beta1;
|
||||
@@ -21134,24 +21184,23 @@ impl GpuDqnTrainer {
|
||||
// single fused launch. Each group reads its own ISV-driven weight
|
||||
// clamp (`ISV[WEIGHT_BOUND[group]]`) and AdamW weight_decay
|
||||
// (`ISV[WD_RATE[group]]`); L1 lambda is read from
|
||||
// `ISV[L1_LAMBDA_TRUNK_INDEX]` for the trunk only (groups 1+2 pass 0).
|
||||
// The split also resolves Layer A's Pearl C engagement-counter
|
||||
// limitation (engagement was previously accounted under group 0
|
||||
// only — A14/A15 deferral). Each sub-launch now writes per-block
|
||||
// counts at distinct offsets `group_idx × MAX_BLOCKS_PER_ADAM`,
|
||||
// and `pearl_c_post_adam_engagement_check` reads each group's
|
||||
// dedicated range. Per `feedback_no_quickfixes`: NOT a
|
||||
// max/min-of-3 single-launch shortcut. Per `feedback_no_partial_refactor`:
|
||||
// every consumer of the WEIGHT_BOUND/WD_RATE/L1_LAMBDA_TRUNK
|
||||
// contract migrates in this same commit.
|
||||
// `ISV[L1_LAMBDA_TRUNK_INDEX]` for the trunk only (groups 1+2+
|
||||
// trunk-extras pass 0). The split also resolves Layer A's Pearl C
|
||||
// engagement-counter limitation (engagement was previously
|
||||
// accounted under group 0 only — A14/A15 deferral). Each
|
||||
// SP4-canonical sub-launch writes per-block counts at distinct
|
||||
// offsets `group_idx × MAX_BLOCKS_PER_ADAM`; trunk-extras passes
|
||||
// `SP4_ENGAGE_OFFSET_DISABLED` because it shares trunk's
|
||||
// WEIGHT_BOUND / WD_RATE bounds (Pearl C engagement is a per-
|
||||
// bound diagnostic, not a per-launch one).
|
||||
let nan_flags_ptr = self.nan_flags_buf.raw_ptr();
|
||||
let diag_slot: i32 = SP4_DQN_ADAM_DIAG_SLOT;
|
||||
let engage_buf_ptr: u64 = self.clamp_engage_per_block_buf.dev_ptr;
|
||||
|
||||
// Per-group descriptor: (ParamGroup, byte_offset, count). Mirrors
|
||||
// `param_group_buffers` for the DQN-internal groups exactly (trunk
|
||||
// = full main params, value = bytes [13..17], branches = bytes
|
||||
// [17..33]). We compute byte-offsets here directly to avoid
|
||||
// Per-group descriptor: (ParamGroup, byte_offset, count, engage_offset).
|
||||
// Mirrors `param_group_buffers` for the DQN-internal groups exactly
|
||||
// for groups 0/1/2; the trunk-extras tail extends through end of
|
||||
// params buffer. We compute byte-offsets here directly to avoid
|
||||
// threading `SP4AuxBuffers` through this launcher (DQN-internal
|
||||
// groups never need aux pointers).
|
||||
let f32_sz = std::mem::size_of::<f32>() as u64;
|
||||
@@ -21163,33 +21212,63 @@ impl GpuDqnTrainer {
|
||||
let branches_byte_off = padded_byte_offset(¶m_sizes, 17);
|
||||
let branches_byte_end = padded_byte_offset(¶m_sizes, 33);
|
||||
let branches_count = ((branches_byte_end - branches_byte_off) / f32_sz) as usize;
|
||||
// Trunk-extras tail: tensors [33..163). Spans from end of branches
|
||||
// to end of params buffer.
|
||||
let trunk_extras_byte_off = padded_byte_offset(¶m_sizes, 33);
|
||||
let trunk_extras_count = self.total_params
|
||||
.saturating_sub((trunk_extras_byte_off / f32_sz) as usize);
|
||||
|
||||
let groups: [(SP4Group, u64, usize); 3] = [
|
||||
(SP4Group::DqnTrunk, trunk_byte_off, trunk_count),
|
||||
(SP4Group::DqnValue, value_byte_off, value_count),
|
||||
(SP4Group::DqnBranches, branches_byte_off, branches_count),
|
||||
// (group_tag, byte_off, count, engage_offset). The 4th entry is
|
||||
// tagged as DqnTrunk (pragma: shares trunk's WEIGHT_BOUND/WD_RATE
|
||||
// contract) but uses SP4_ENGAGE_OFFSET_DISABLED so its block
|
||||
// engagement does not overwrite trunk's slot in the Pearl C buf.
|
||||
let trunk_engage_off: i32 = (SP4Group::DqnTrunk.idx() * MAX_BLOCKS_PER_ADAM) as i32;
|
||||
let value_engage_off: i32 = (SP4Group::DqnValue.idx() * MAX_BLOCKS_PER_ADAM) as i32;
|
||||
let branches_engage_off: i32 = (SP4Group::DqnBranches.idx() * MAX_BLOCKS_PER_ADAM) as i32;
|
||||
let groups: [(SP4Group, u64, usize, i32); 4] = [
|
||||
(SP4Group::DqnTrunk, trunk_byte_off, trunk_count, trunk_engage_off),
|
||||
(SP4Group::DqnValue, value_byte_off, value_count, value_engage_off),
|
||||
(SP4Group::DqnBranches, branches_byte_off, branches_count, branches_engage_off),
|
||||
// Trunk-extras tail — shares DqnTrunk bounds; Pearl C disabled
|
||||
// for this sub-launch (engagement folds into trunk's
|
||||
// accounting since the same WEIGHT_BOUND drives both clamps).
|
||||
(SP4Group::DqnTrunk, trunk_extras_byte_off, trunk_extras_count, SP4_ENGAGE_OFFSET_DISABLED),
|
||||
];
|
||||
|
||||
// Coverage invariant: every of the 163 weight tensors must be
|
||||
// covered by exactly one of the 4 sub-launches above. Layer B
|
||||
// fix-up: prevents regression of the silent ~70% Adam freeze
|
||||
// that the original 3-way split introduced.
|
||||
debug_assert_eq!(
|
||||
trunk_count + value_count + branches_count + trunk_extras_count,
|
||||
self.total_params,
|
||||
"SP4 Layer B fix-up: 4-way Adam coverage must span all 163 weight tensors \
|
||||
(got trunk={} + value={} + branches={} + trunk_extras={} = {}, expected total_params={})",
|
||||
trunk_count, value_count, branches_count, trunk_extras_count,
|
||||
trunk_count + value_count + branches_count + trunk_extras_count,
|
||||
self.total_params
|
||||
);
|
||||
|
||||
// Safety: argument order matches the extern "C" kernel signature exactly.
|
||||
// grad_norm_buf contains L2 norm from launch_grad_norm_finalize().
|
||||
for (group, byte_off, count) in groups {
|
||||
for (group, byte_off, count, engage_buf_offset) in groups {
|
||||
if count == 0 {
|
||||
continue;
|
||||
}
|
||||
|
||||
// Per-group ISV-driven post-Adam weight clamp + AdamW
|
||||
// weight_decay rate. EPS_CLAMP_FLOOR=1.0 ε on the bound itself
|
||||
// (SP1 cold-start pearl); `.max(0.0)` on weight_decay enforces
|
||||
// the non-negativity contract on the AdamW rate.
|
||||
let weight_clamp_max_abs: f32 =
|
||||
self.read_isv_signal_at(weight_bound(group.idx())).max(EPS_CLAMP_FLOOR);
|
||||
let weight_decay_value: f32 =
|
||||
self.read_isv_signal_at(wd_rate(group.idx())).max(0.0);
|
||||
// weight_decay rate via the SP4 Layer B helper. Trunk-extras
|
||||
// sub-launch reuses DqnTrunk's bounds (its `group` tag).
|
||||
let (weight_clamp_max_abs, weight_decay_value) =
|
||||
self.read_group_adam_bounds(group);
|
||||
|
||||
// L1 proximal step is configured for the trunk only (G8 boundary
|
||||
// at end of `w_s1`). Value/branches pass `l1_end=0 + l1_lambda=0`
|
||||
// which deactivates the in-kernel L1 path.
|
||||
let (l1_end_arg, l1_lambda_arg): (i32, f32) = if group == SP4Group::DqnTrunk {
|
||||
// at end of `w_s1`). Value/branches/trunk-extras pass
|
||||
// `l1_end=0 + l1_lambda=0` which deactivates the in-kernel L1
|
||||
// path. Trunk-extras tensors [33..163) are well past `w_s1`'s
|
||||
// boundary and don't get L1 — only the canonical trunk
|
||||
// sub-launch (byte_off == 0) sets L1.
|
||||
let (l1_end_arg, l1_lambda_arg): (i32, f32) = if byte_off == 0 {
|
||||
let l1_lambda = self.read_isv_signal_at(L1_LAMBDA_TRUNK_INDEX).max(0.0);
|
||||
(trunk_l1_end, l1_lambda)
|
||||
} else {
|
||||
@@ -21209,10 +21288,10 @@ impl GpuDqnTrainer {
|
||||
let blocks = ((count + 255) / 256) as u32;
|
||||
debug_assert!(
|
||||
(blocks as usize) <= MAX_BLOCKS_PER_ADAM,
|
||||
"dqn_adam[{:?}] blocks={} exceeds MAX_BLOCKS_PER_ADAM={} — Pearl C \
|
||||
clamp_engage_per_block_buf cannot accommodate; resize buffer \
|
||||
"dqn_adam[{:?} byte_off={}] blocks={} exceeds MAX_BLOCKS_PER_ADAM={} — Pearl C \
|
||||
clamp_engage_per_block_buf cannot accommodate; grow MAX_BLOCKS_PER_ADAM \
|
||||
or chunk the launch.",
|
||||
group, blocks, MAX_BLOCKS_PER_ADAM
|
||||
group, byte_off, blocks, MAX_BLOCKS_PER_ADAM
|
||||
);
|
||||
let cfg = LaunchConfig {
|
||||
grid_dim: (blocks, 1, 1),
|
||||
@@ -21220,8 +21299,6 @@ impl GpuDqnTrainer {
|
||||
shared_mem_bytes: 0,
|
||||
};
|
||||
|
||||
let engage_buf_offset: i32 = (group.idx() * MAX_BLOCKS_PER_ADAM) as i32;
|
||||
|
||||
unsafe {
|
||||
self.stream
|
||||
.launch_builder(&self.adam_update_kernel)
|
||||
@@ -21245,10 +21322,13 @@ impl GpuDqnTrainer {
|
||||
.arg(&nan_flags_ptr) // SP4 A14: sticky engage flag store
|
||||
.arg(&diag_slot) // SP4 A14: diag slot (50 for dqn_adam)
|
||||
.arg(&engage_buf_ptr) // SP4 A14: per-block engagement buf
|
||||
.arg(&engage_buf_offset) // SP4 Layer B: per-group offset (resolves A14/A15 deferral)
|
||||
.arg(&engage_buf_offset) // SP4 Layer B: per-group offset; SP4_ENGAGE_OFFSET_DISABLED for trunk-extras tail
|
||||
.launch(cfg)
|
||||
.map_err(|e| {
|
||||
MLError::ModelError(format!("dqn_adam_update_kernel[{:?}] launch: {e}", group))
|
||||
MLError::ModelError(format!(
|
||||
"dqn_adam_update_kernel[{:?} byte_off={}] launch: {e}",
|
||||
group, byte_off
|
||||
))
|
||||
})?;
|
||||
}
|
||||
}
|
||||
@@ -21275,13 +21355,25 @@ impl GpuDqnTrainer {
|
||||
/// `[group_idx × MAX_BLOCKS_PER_ADAM, +MAX_BLOCKS_PER_ADAM)` range.
|
||||
///
|
||||
/// **Layer B (2026-05-01) — split sub-launches landed.** The main
|
||||
/// DQN Adam now runs as 3 sub-launches (Trunk/Value/Branches), each
|
||||
/// writing per-block engagement counts at its own group offset.
|
||||
/// `pearl_c_post_adam_engagement_check(group, count)` reads the
|
||||
/// dedicated range for any of the 8 SP4 groups directly. The Layer A
|
||||
/// "groups 1/2 accounted under group 0" limitation is resolved by
|
||||
/// the same Layer B atomic commit that introduces the
|
||||
/// DQN Adam now runs as 3 SP4-canonical sub-launches (Trunk / Value
|
||||
/// / Branches), each writing per-block engagement counts at its own
|
||||
/// group offset. `pearl_c_post_adam_engagement_check(group, count)`
|
||||
/// reads the dedicated range for any of the 8 SP4 groups directly.
|
||||
/// The Layer A "groups 1/2 accounted under group 0" limitation is
|
||||
/// resolved by the same Layer B atomic commit that introduces the
|
||||
/// `WEIGHT_BOUND[group]` consumer migration.
|
||||
///
|
||||
/// **Layer B fix-up (2026-05-01)** — the 4th sub-launch (trunk-
|
||||
/// extras tail covering tensors `[33..163)`) is tagged DqnTrunk for
|
||||
/// the WEIGHT_BOUND/WD_RATE bounds contract but uses
|
||||
/// `SP4_ENGAGE_OFFSET_DISABLED` for Pearl C. Engagement for those
|
||||
/// tensors **folds into the canonical DqnTrunk count** (same bound
|
||||
/// drives both clamps; Pearl C is a per-bound diagnostic). Callers
|
||||
/// passing `ParamGroup::DqnTrunk` here only see the canonical
|
||||
/// trunk's engagement counts in the buffer; trunk-extras' clamp
|
||||
/// engagements do not affect the rate-deficit EMA. This keeps
|
||||
/// engagement accounting faithful to the WEIGHT_BOUND[DqnTrunk]
|
||||
/// contract Pearl C is regulating.
|
||||
pub(crate) fn pearl_c_post_adam_engagement_check(
|
||||
&mut self,
|
||||
group: ParamGroup,
|
||||
@@ -22253,9 +22345,11 @@ impl GpuDqnTrainer {
|
||||
// - smoke-test-rxhjh (Mech 6 100× + Mech 8 anchor reset): F1 NaN
|
||||
// @ 2300 (worse than Mech 6 alone) — Mech 8 reverted; the
|
||||
// persistent slow_ema across folds was unintentional protection.
|
||||
// SP3 closes the cross-fold Adam pathology with Mech 9 (post-Adam
|
||||
// weight clamp at 100 × Q_ABS_REF.max(1.0)) rather than further
|
||||
// tuning the clip-anchor side of the chain.
|
||||
// SP3 closed the cross-fold Adam pathology with Mech 9 (post-Adam
|
||||
// weight clamp via host-side multiplier × Q_ABS_REF) rather than
|
||||
// further tuning the clip-anchor side of the chain. Layer B
|
||||
// generalised the bound to per-group `ISV[WEIGHT_BOUND[group]]`
|
||||
// (Wiener-EMA over p99(|params|), no host multiplier).
|
||||
//
|
||||
// SP4 Layer B (Mech 6): adaptive clip upper bound = ISV[GRAD_CLIP_BOUND=168].
|
||||
// Producer `launch_sp4_grad_norm_p99` (Layer A Task A8) writes the
|
||||
|
||||
@@ -208,7 +208,7 @@ void iql_adam_kernel(
|
||||
float max_grad_norm,
|
||||
const int* __restrict__ t_buf, /* [1] Adam step on device */
|
||||
int total_params,
|
||||
/* SP3 Mech 9 (close-out v2): ISV-driven |p| bound (100 × Q_ABS_REF.max(1.0))
|
||||
/* SP4 Layer B: ISV-driven |p| bound (`ISV[WEIGHT_BOUND[IqlHigh|IqlLow]].max(EPS_CLAMP_FLOOR)`)
|
||||
* — same pattern wired into all five Adam kernels. 0=disabled (debug only). */
|
||||
float weight_clamp_max_abs,
|
||||
/* SP4 Task A14 (2026-04-30): Pearl C engagement counter. Same warp+
|
||||
|
||||
@@ -844,13 +844,14 @@ void iqn_grad_norm_phase2(
|
||||
}
|
||||
|
||||
/* ── Adam Update Kernel ─────────────────────────────────────────────── */
|
||||
/* SP3 Mech 9 (close-out v2): trailing `weight_clamp_max_abs` arg matches the
|
||||
/* SP4 Layer B: trailing `weight_clamp_max_abs` arg matches the
|
||||
* pattern wired into all five Adam kernels (dqn/iqn/iql/attn/curiosity).
|
||||
* Bound = 100 × ISV[Q_ABS_REF].max(1.0) computed host-side. Targets the
|
||||
* IQN-internal weight buffer that the slot-26 cuBLAS GEMM reads —
|
||||
* `apply_iqn_trunk_gradient` saw NaN at slot 26 with finite inputs because
|
||||
* `online_params` here drifted via this Adam, never reached by the
|
||||
* single-kernel Mech 9 in `dqn_adam_update_kernel`. Disabled when bound==0.
|
||||
* Bound = `ISV[WEIGHT_BOUND[ParamGroup::Iqn]].max(EPS_CLAMP_FLOOR)` computed
|
||||
* host-side via `read_group_adam_bounds`. Targets the IQN-internal weight
|
||||
* buffer that the slot-26 cuBLAS GEMM reads — `apply_iqn_trunk_gradient`
|
||||
* saw NaN at slot 26 with finite inputs because `online_params` here
|
||||
* drifted via this Adam, never reached by the single-kernel Mech 9 in
|
||||
* `dqn_adam_update_kernel`. Disabled when bound==0.
|
||||
*
|
||||
* SP4 Task A14 (2026-04-30): Pearl C engagement counter via warp+block
|
||||
* tree-reduce mirroring `dqn_grad_norm_kernel` (no atomicAdd). Single
|
||||
@@ -876,7 +877,7 @@ void iqn_adam_kernel(
|
||||
int b1_size, /* runtime: unused, for consistent interface */
|
||||
int b2_size, /* runtime: unused, for consistent interface */
|
||||
int b3_size, /* runtime: unused, for consistent interface */
|
||||
float weight_clamp_max_abs, /* SP3 Mech 9: ISV-driven |p| bound (100 × Q_ABS_REF.max(1.0)). 0=disabled. */
|
||||
float weight_clamp_max_abs, /* SP4 Layer B: ISV-driven |p| bound (`ISV[WEIGHT_BOUND[Iqn]].max(EPS_CLAMP_FLOOR)`). 0=disabled. */
|
||||
/* SP4 Task A14: Pearl C engagement counter. */
|
||||
int* __restrict__ nan_flags_buf, /* sticky overflow flag */
|
||||
int diag_slot, /* nan_flags_buf index for this Adam kernel */
|
||||
@@ -918,11 +919,12 @@ void iqn_adam_kernel(
|
||||
float p = params[tid];
|
||||
p = p - lr * (m_hat / (sqrtf(v_hat) + eps) + weight_decay * p);
|
||||
|
||||
/* SP3 Mech 9 + SP4 Task A14: post-Adam |p| clamp +
|
||||
/* SP4 Layer B + SP4 Task A14: post-Adam |p| clamp +
|
||||
* engagement-counter increment. Bound is host-supplied
|
||||
* (100 × Q_ABS_REF.max(1.0)). One OOM below slot 48
|
||||
* diagnostic threshold (1e3 × Q_ABS_REF) so the regression
|
||||
* sentinel keeps 10× firing headroom. */
|
||||
* (`ISV[WEIGHT_BOUND[Iqn]].max(EPS_CLAMP_FLOOR)`). The
|
||||
* ISV producer (Wiener-EMA over p99(|params|)) keeps the
|
||||
* bound naturally above legitimate weight magnitudes;
|
||||
* `EPS_CLAMP_FLOOR=1.0` is a cold-start floor only. */
|
||||
if (weight_clamp_max_abs > 0.0f) {
|
||||
if (fabsf(p) > weight_clamp_max_abs) {
|
||||
if (nan_flags_buf != nullptr && diag_slot >= 0) {
|
||||
|
||||
@@ -52,22 +52,30 @@ pub const SP4_SLOT_END: usize = 171;
|
||||
|
||||
/// Maximum block grid size per Adam-kernel launch — the engagement
|
||||
/// counter buffer reserves `MAX_BLOCKS_PER_ADAM` slots per param group,
|
||||
/// indexed by `blockIdx.x`. 256 covers up to ≈65k params per group at
|
||||
/// blockDim=256 threads, comfortably above all current production
|
||||
/// param-group sizes (DQN trunk≈30k, DQN value≈4k, branches≈80k —
|
||||
/// branch group is the only one exceeding 65k, but the launcher caps
|
||||
/// `blocks` at `MAX_BLOCKS_PER_ADAM` defensively in launch code).
|
||||
pub const MAX_BLOCKS_PER_ADAM: usize = 256;
|
||||
/// indexed by `blockIdx.x`. 4096 covers up to ≈1M params per group at
|
||||
/// blockDim=256 threads.
|
||||
///
|
||||
/// **Layer B fix-up (2026-05-01)**: bumped 256 → 4096 to accommodate
|
||||
/// the trunk-extras tail sub-launch (tensors `[33..163)` co-located
|
||||
/// under DqnTrunk). Production cfg `shared_h2=256, num_atoms=51,
|
||||
/// adv_h=128, MOE_NUM_EXPERTS=8, MOE_EXPERT_BOTTLENECK=64` puts
|
||||
/// trunk-extras at ≈600k floats → ≈2400 blocks at blockDim=256, so
|
||||
/// 4096 leaves a comfortable safety margin (1.7×) for future config
|
||||
/// growth without wasted memory (each engage slot is 4 bytes × 11
|
||||
/// groups × 4096 = 180 KiB total — negligible vs. params buffer).
|
||||
/// Other groups (DqnValue ≈ 4k params, DqnBranches ≈ 80k, IQN/IQL/
|
||||
/// Attn/Curiosity bounded similarly) easily fit within 4096 blocks.
|
||||
pub const MAX_BLOCKS_PER_ADAM: usize = 4096;
|
||||
|
||||
/// Total length of the Pearl C engagement-counter mapped-pinned buffer
|
||||
/// (`clamp_engage_per_block_buf`). The 8 main param groups occupy
|
||||
/// offsets [0..2048) (8 × 256). Curiosity has 4 sub-launches (W1/b1/
|
||||
/// W2/b2) sharing the same kernel; to avoid block-index collisions
|
||||
/// between sub-launches, each gets its own offset slot. Sub-launch 0
|
||||
/// reuses Curiosity's main group offset (group_idx=7 → 1792). Sub-
|
||||
/// launches 1/2/3 occupy extra slots at offsets 2048/2304/2560 — i.e.
|
||||
/// 3 extra logical "groups" beyond the canonical 8. Total length:
|
||||
/// (8 + 3) × 256 = 11 × 256 = 2816 i32s. Host-side
|
||||
/// offsets `[0..8 × MAX_BLOCKS_PER_ADAM)`. Curiosity has 4 sub-launches
|
||||
/// (W1/b1/W2/b2) sharing the same kernel; to avoid block-index
|
||||
/// collisions between sub-launches, each gets its own offset slot.
|
||||
/// Sub-launch 0 reuses Curiosity's main group offset (group_idx=7).
|
||||
/// Sub-launches 1/2/3 occupy extra slots — 3 extra logical "groups"
|
||||
/// beyond the canonical 8. Total length: (8 + 3) × MAX_BLOCKS_PER_ADAM
|
||||
/// = 11 × 4096 = 45056 i32s. Host-side
|
||||
/// `pearl_c_post_adam_engagement_check` for `ParamGroup::Curiosity`
|
||||
/// sums all 4 sub-launch ranges.
|
||||
pub const SP4_ENGAGE_EXTRA_CURIOSITY_SUBLAUNCHES: usize = 3;
|
||||
@@ -212,18 +220,22 @@ mod tests {
|
||||
#[test]
|
||||
fn pearl_c_engage_buf_layout() {
|
||||
// 8 main groups + 3 extra curiosity sub-launch slots.
|
||||
assert_eq!(MAX_BLOCKS_PER_ADAM, 256);
|
||||
// Layer B fix-up (2026-05-01): MAX_BLOCKS_PER_ADAM grown 256 → 4096
|
||||
// to fit trunk-extras tail (≈2400 blocks at production cfg).
|
||||
assert_eq!(MAX_BLOCKS_PER_ADAM, 4096);
|
||||
assert_eq!(SP4_PARAM_GROUP_COUNT, 8);
|
||||
assert_eq!(SP4_ENGAGE_EXTRA_CURIOSITY_SUBLAUNCHES, 3);
|
||||
assert_eq!(SP4_ENGAGE_BUF_LEN, 11 * 256);
|
||||
assert_eq!(SP4_ENGAGE_BUF_LEN, 2816);
|
||||
assert_eq!(SP4_ENGAGE_BUF_LEN, 11 * MAX_BLOCKS_PER_ADAM);
|
||||
assert_eq!(SP4_ENGAGE_BUF_LEN, 11 * 4096);
|
||||
assert_eq!(SP4_ENGAGE_BUF_LEN, 45056);
|
||||
|
||||
// Curiosity sub-launch offsets are distinct + contiguous starting
|
||||
// from the canonical Curiosity group offset.
|
||||
assert_eq!(SP4_ENGAGE_OFFSET_CURIOSITY_W1, 7 * 256);
|
||||
assert_eq!(SP4_ENGAGE_OFFSET_CURIOSITY_B1, 8 * 256);
|
||||
assert_eq!(SP4_ENGAGE_OFFSET_CURIOSITY_W2, 9 * 256);
|
||||
assert_eq!(SP4_ENGAGE_OFFSET_CURIOSITY_B2, 10 * 256);
|
||||
// from the canonical Curiosity group offset. They're derived from
|
||||
// `MAX_BLOCKS_PER_ADAM` so they auto-track its growth.
|
||||
assert_eq!(SP4_ENGAGE_OFFSET_CURIOSITY_W1, 7 * MAX_BLOCKS_PER_ADAM as i32);
|
||||
assert_eq!(SP4_ENGAGE_OFFSET_CURIOSITY_B1, 8 * MAX_BLOCKS_PER_ADAM as i32);
|
||||
assert_eq!(SP4_ENGAGE_OFFSET_CURIOSITY_W2, 9 * MAX_BLOCKS_PER_ADAM as i32);
|
||||
assert_eq!(SP4_ENGAGE_OFFSET_CURIOSITY_B2, 10 * MAX_BLOCKS_PER_ADAM as i32);
|
||||
|
||||
// Highest sub-launch end must fit within the buffer.
|
||||
let highest_end = (SP4_ENGAGE_OFFSET_CURIOSITY_B2 as usize) + MAX_BLOCKS_PER_ADAM;
|
||||
|
||||
@@ -1619,16 +1619,11 @@ impl FusedTrainingCtx {
|
||||
// (cubin-shared with GpuAttention). TLOB params co-live with
|
||||
// attention in the Attn (group 6) param-group taxonomy, so
|
||||
// the same ISV-driven WEIGHT_BOUND[Attn] / WD_RATE[Attn]
|
||||
// contract applies. ε floor (EPS_CLAMP_FLOOR = 1.0) handles
|
||||
// cold-start; `.max(0.0)` enforces wd non-negativity.
|
||||
let tlob_weight_clamp_max_abs: f32 = self.trainer
|
||||
.read_isv_signal_at(crate::cuda_pipeline::weight_bound(
|
||||
crate::cuda_pipeline::ParamGroup::Attn.idx()))
|
||||
.max(crate::cuda_pipeline::EPS_CLAMP_FLOOR);
|
||||
let tlob_weight_decay_value: f32 = self.trainer
|
||||
.read_isv_signal_at(crate::cuda_pipeline::wd_rate(
|
||||
crate::cuda_pipeline::ParamGroup::Attn.idx()))
|
||||
.max(0.0);
|
||||
// contract applies. `read_group_adam_bounds` enforces the
|
||||
// EPS_CLAMP_FLOOR=1.0 ε for cold-start and AdamW
|
||||
// non-negativity.
|
||||
let (tlob_weight_clamp_max_abs, tlob_weight_decay_value) =
|
||||
self.trainer.read_group_adam_bounds(crate::cuda_pipeline::ParamGroup::Attn);
|
||||
self.gpu_tlob
|
||||
.adam_step(
|
||||
self.trainer.config().lr,
|
||||
@@ -1924,26 +1919,14 @@ impl FusedTrainingCtx {
|
||||
//
|
||||
// SP4 Layer B (Mech 9): ISV-driven post-Adam |p| clamp for
|
||||
// `iql_adam_kernel`. Per-group bounds: IqlHigh (group 4) and
|
||||
// IqlLow (group 5) each read `ISV[WEIGHT_BOUND[group]]`. ε
|
||||
// floor (EPS_CLAMP_FLOOR = 1.0) handles cold-start ISV. AdamW
|
||||
// weight_decay sourced from `ISV[WD_RATE[group]]` with
|
||||
// `.max(0.0)` non-negativity contract.
|
||||
let iql_high_clamp: f32 = self.trainer
|
||||
.read_isv_signal_at(crate::cuda_pipeline::weight_bound(
|
||||
crate::cuda_pipeline::ParamGroup::IqlHigh.idx()))
|
||||
.max(crate::cuda_pipeline::EPS_CLAMP_FLOOR);
|
||||
let iql_high_wd: f32 = self.trainer
|
||||
.read_isv_signal_at(crate::cuda_pipeline::wd_rate(
|
||||
crate::cuda_pipeline::ParamGroup::IqlHigh.idx()))
|
||||
.max(0.0);
|
||||
let iql_low_clamp: f32 = self.trainer
|
||||
.read_isv_signal_at(crate::cuda_pipeline::weight_bound(
|
||||
crate::cuda_pipeline::ParamGroup::IqlLow.idx()))
|
||||
.max(crate::cuda_pipeline::EPS_CLAMP_FLOOR);
|
||||
let iql_low_wd: f32 = self.trainer
|
||||
.read_isv_signal_at(crate::cuda_pipeline::wd_rate(
|
||||
crate::cuda_pipeline::ParamGroup::IqlLow.idx()))
|
||||
.max(0.0);
|
||||
// IqlLow (group 5) each read `ISV[WEIGHT_BOUND[group]]` and
|
||||
// `ISV[WD_RATE[group]]` via `read_group_adam_bounds`, which
|
||||
// applies the EPS_CLAMP_FLOOR=1.0 ε for cold-start ISV and
|
||||
// the AdamW non-negativity contract.
|
||||
let (iql_high_clamp, iql_high_wd) =
|
||||
self.trainer.read_group_adam_bounds(crate::cuda_pipeline::ParamGroup::IqlHigh);
|
||||
let (iql_low_clamp, iql_low_wd) =
|
||||
self.trainer.read_group_adam_bounds(crate::cuda_pipeline::ParamGroup::IqlLow);
|
||||
let states_f32 = self.trainer.states_buf();
|
||||
self.gpu_iql.train_value_step(states_f32, iql_high_clamp, iql_high_wd)
|
||||
.map_err(|e| anyhow::anyhow!("IQL high train: {e}"))?;
|
||||
@@ -2042,18 +2025,12 @@ impl FusedTrainingCtx {
|
||||
|
||||
// SP4 Layer B (Mech 9): ISV-driven post-Adam |p| clamp for
|
||||
// `iqn_adam_kernel`. IQN params live in group 3
|
||||
// (`ParamGroup::Iqn`); read `ISV[WEIGHT_BOUND[Iqn]]` with
|
||||
// EPS_CLAMP_FLOOR = 1.0 ε for cold-start. AdamW
|
||||
// weight_decay sourced from `ISV[WD_RATE[Iqn]]` with
|
||||
// `.max(0.0)` non-negativity contract.
|
||||
let weight_clamp_max_abs: f32 = self.trainer
|
||||
.read_isv_signal_at(crate::cuda_pipeline::weight_bound(
|
||||
crate::cuda_pipeline::ParamGroup::Iqn.idx()))
|
||||
.max(crate::cuda_pipeline::EPS_CLAMP_FLOOR);
|
||||
let weight_decay_value: f32 = self.trainer
|
||||
.read_isv_signal_at(crate::cuda_pipeline::wd_rate(
|
||||
crate::cuda_pipeline::ParamGroup::Iqn.idx()))
|
||||
.max(0.0);
|
||||
// (`ParamGroup::Iqn`); `read_group_adam_bounds` returns
|
||||
// `ISV[WEIGHT_BOUND[Iqn]]` (with EPS_CLAMP_FLOOR=1.0 ε
|
||||
// for cold-start) and `ISV[WD_RATE[Iqn]]` (with AdamW
|
||||
// non-negativity).
|
||||
let (weight_clamp_max_abs, weight_decay_value) =
|
||||
self.trainer.read_group_adam_bounds(crate::cuda_pipeline::ParamGroup::Iqn);
|
||||
|
||||
match iqn.execute_training_pipeline(
|
||||
online_h_s2, next_states_buf, &self.target_dueling,
|
||||
@@ -2130,17 +2107,12 @@ impl FusedTrainingCtx {
|
||||
let mgn = self.trainer.config().max_grad_norm;
|
||||
// SP4 Layer B (Mech 9): ISV-driven post-Adam |p| clamp for
|
||||
// `attn_adam_kernel`. Attention params live in group 6
|
||||
// (`ParamGroup::Attn`); read `ISV[WEIGHT_BOUND[Attn]]` with
|
||||
// EPS_CLAMP_FLOOR = 1.0 ε for cold-start. AdamW weight_decay
|
||||
// sourced from `ISV[WD_RATE[Attn]]` with `.max(0.0)`.
|
||||
let attn_weight_clamp_max_abs: f32 = self.trainer
|
||||
.read_isv_signal_at(crate::cuda_pipeline::weight_bound(
|
||||
crate::cuda_pipeline::ParamGroup::Attn.idx()))
|
||||
.max(crate::cuda_pipeline::EPS_CLAMP_FLOOR);
|
||||
let attn_weight_decay_value: f32 = self.trainer
|
||||
.read_isv_signal_at(crate::cuda_pipeline::wd_rate(
|
||||
crate::cuda_pipeline::ParamGroup::Attn.idx()))
|
||||
.max(0.0);
|
||||
// (`ParamGroup::Attn`); `read_group_adam_bounds` returns
|
||||
// `ISV[WEIGHT_BOUND[Attn]]` (with EPS_CLAMP_FLOOR=1.0 ε
|
||||
// for cold-start) and `ISV[WD_RATE[Attn]]` (with AdamW
|
||||
// non-negativity).
|
||||
let (attn_weight_clamp_max_abs, attn_weight_decay_value) =
|
||||
self.trainer.read_group_adam_bounds(crate::cuda_pipeline::ParamGroup::Attn);
|
||||
attn.adam_step(lr, mgn, Some(&attn_stream_ref), ws_override,
|
||||
attn_weight_clamp_max_abs, attn_weight_decay_value)
|
||||
.map_err(|e| anyhow::anyhow!("Attention Adam: {e}"))?;
|
||||
@@ -2177,14 +2149,8 @@ impl FusedTrainingCtx {
|
||||
let mgn = self.trainer.config().max_grad_norm;
|
||||
// SP4 Layer B (Mech 9): same ISV-driven bound as parallel arm.
|
||||
// Sequential-fallback path.
|
||||
let attn_weight_clamp_max_abs: f32 = self.trainer
|
||||
.read_isv_signal_at(crate::cuda_pipeline::weight_bound(
|
||||
crate::cuda_pipeline::ParamGroup::Attn.idx()))
|
||||
.max(crate::cuda_pipeline::EPS_CLAMP_FLOOR);
|
||||
let attn_weight_decay_value: f32 = self.trainer
|
||||
.read_isv_signal_at(crate::cuda_pipeline::wd_rate(
|
||||
crate::cuda_pipeline::ParamGroup::Attn.idx()))
|
||||
.max(0.0);
|
||||
let (attn_weight_clamp_max_abs, attn_weight_decay_value) =
|
||||
self.trainer.read_group_adam_bounds(crate::cuda_pipeline::ParamGroup::Attn);
|
||||
attn.adam_step(lr, mgn, None, None,
|
||||
attn_weight_clamp_max_abs, attn_weight_decay_value)
|
||||
.map_err(|e| anyhow::anyhow!("Attention Adam: {e}"))?;
|
||||
@@ -2197,14 +2163,8 @@ impl FusedTrainingCtx {
|
||||
|
||||
// SP4 Layer B (Mech 9): same ISV-driven bound as the
|
||||
// parallel arm above. Mirror the dqn_adam pattern.
|
||||
let weight_clamp_max_abs: f32 = self.trainer
|
||||
.read_isv_signal_at(crate::cuda_pipeline::weight_bound(
|
||||
crate::cuda_pipeline::ParamGroup::Iqn.idx()))
|
||||
.max(crate::cuda_pipeline::EPS_CLAMP_FLOOR);
|
||||
let weight_decay_value: f32 = self.trainer
|
||||
.read_isv_signal_at(crate::cuda_pipeline::wd_rate(
|
||||
crate::cuda_pipeline::ParamGroup::Iqn.idx()))
|
||||
.max(0.0);
|
||||
let (weight_clamp_max_abs, weight_decay_value) =
|
||||
self.trainer.read_group_adam_bounds(crate::cuda_pipeline::ParamGroup::Iqn);
|
||||
|
||||
match iqn.execute_training_pipeline(
|
||||
online_h_s2, next_states_buf, &self.target_dueling,
|
||||
|
||||
@@ -2780,3 +2780,57 @@ Verification (local, RTX 3050 Ti, post-commit):
|
||||
ZERO matches.
|
||||
- `cargo test -p ml --lib --offline`: 928 passed, 14 failed (all 14
|
||||
pre-existing on HEAD `1389d1c81`; no new failures introduced).
|
||||
|
||||
**Layer B fix-up (2026-05-01)**: code-quality reviewer surfaced a silent
|
||||
correctness regression in the original 3-way `launch_adam_update` split —
|
||||
tensors `[33..163)` (bottleneck, VSN bottleneck, GLU gates, KAN, regime
|
||||
gate, spacing, multi-horizon value heads, risk, ISV encoder, VSN per-group,
|
||||
aux heads, MoE — ~70% of the DQN's 163 weight tensors) were silently no
|
||||
longer Adam-updated (grads computed, m/v allocated, no parameter step).
|
||||
The 3-group SP4 taxonomy (DqnTrunk / DqnValue / DqnBranches via
|
||||
`param_group_buffers`) only maps tensors `[0..33)`. Resolution:
|
||||
- 4th sub-launch added — trunk-extras tail covering `[33..163)`,
|
||||
tagged `ParamGroup::DqnTrunk` for shared `WEIGHT_BOUND/WD_RATE`
|
||||
bounds, with `SP4_ENGAGE_OFFSET_DISABLED` so Pearl C engagement
|
||||
rolls into trunk's accounting (same bound drives both clamps).
|
||||
- `MAX_BLOCKS_PER_ADAM` grown 256 → 4096 to fit trunk-extras (~2400
|
||||
blocks at production cfg `shared_h2=256`, `MOE_NUM_EXPERTS=8`).
|
||||
`SP4_ENGAGE_BUF_LEN` auto-grows: 11 × 4096 = 45056 ints (≈180 KiB
|
||||
total; engage offsets are derived from `MAX_BLOCKS_PER_ADAM` so
|
||||
they auto-track).
|
||||
- `read_group_adam_bounds(group) -> (clamp, wd)` helper introduced
|
||||
on `GpuDqnTrainer`; 6 verbose 2-line ISV-read patterns in
|
||||
`fused_training.rs` collapsed to one-liners (TLOB + IQL high +
|
||||
IQL low + IQN parallel + Attn parallel + Attn sequential + IQN
|
||||
sequential = 7 call sites). The 4-way Adam loop in
|
||||
`launch_adam_update` also calls the helper.
|
||||
- Coverage `debug_assert_eq!` added: `trunk + value + branches +
|
||||
trunk_extras == total_params`. Prevents future regressions of
|
||||
this kind.
|
||||
- 9 stale doc-comments referring to "100 × Q_ABS_REF.max(1.0)"
|
||||
replaced by current "ISV[WEIGHT_BOUND[group]].max(EPS_CLAMP_FLOOR)"
|
||||
contract:
|
||||
- `curiosity_training_kernel.cu:324`
|
||||
- `gpu_curiosity_trainer.rs:509,758`
|
||||
- `iqn_dual_head_kernel.cu:847,879,923`
|
||||
- `iql_value_kernel.cu:211`
|
||||
- `attention_backward_kernel.cu:80`
|
||||
- `dqn_utility_kernels.cu:145,200`
|
||||
- `gpu_dqn_trainer.rs:22349` (historical SP3 reference)
|
||||
|
||||
Per `feedback_no_partial_refactor`: trunk-extras consumers (= no
|
||||
dedicated SP4 producer yet) and bound (= reused trunk's
|
||||
`WEIGHT_BOUND/WD_RATE`) ship together with the launch fix in this
|
||||
commit. A future task may split sub-trunk regions (VSN, MoE, etc.)
|
||||
into dedicated SP4 producers once warranted by signal analysis.
|
||||
|
||||
Verification (post-fix-up):
|
||||
- `cargo check -p ml --offline`: clean, same 11 pre-existing warnings.
|
||||
- `cargo test -p ml --lib --offline -- sp4_wiener_ema sp4_isv_slots
|
||||
state_reset_registry`: 11 passed (incl. updated
|
||||
`pearl_c_engage_buf_layout` for MAX_BLOCKS_PER_ADAM=4096).
|
||||
- `cargo test -p ml --lib --offline`: 928 passed, 14 failed (no new
|
||||
failures vs. Layer B initial run).
|
||||
- `git grep -nE "100 × Q_ABS_REF\\.max\\(1\\.0\\)|100\\.0f \\*
|
||||
q_abs_ref|100 \\* Q_ABS_REF" crates/ml/src/cuda_pipeline/
|
||||
crates/ml/src/trainers/`: ZERO matches.
|
||||
|
||||
Reference in New Issue
Block a user