From 58ffb3a48e87c5223be5f1f74068aa77468578bb Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 1 May 2026 11:10:41 +0200 Subject: [PATCH] =?UTF-8?q?feat(sp4):=20Layer=20B=20=E2=80=94=20atomic=20c?= =?UTF-8?q?onsumer=20migration=20to=20ISV-driven=20bounds?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Single coordinated commit per `feedback_no_partial_refactor`. All SP3-era hardcoded magnitude multipliers (10×, 100×, 1e3×, 1e6×) and ε floors (.max(1.0)) replaced by per-slot ISV reads with consumer-side EPS_CLAMP_FLOOR=1.0 numerical safety. Mechanism mapping: - Mech 1 target_q clamp: 10 × Q_ABS_REF.max(1.0) → ISV[TARGET_Q_BOUND] - Mech 2 atom-position clamps (3 sites × 4 branches): 10 × Q_ABS_REF → ISV[ATOM_POS_BOUND[branch]] - Mech 5 fused diagnostic: per-slot ISV reads in `dqn_nan_check_fused_f32_kernel` (kernel takes `isv_signals*` instead of `q_abs_ref_eff` / `h_s2_rms_ema_eff` host args) - Mech 6 adaptive_clip upper_bound: 100 × slow_ema × Q_ABS_REF → ISV[GRAD_CLIP_BOUND] - Mech 9 post-Adam weight_clamp (5 Adam kernels): 100 × Q_ABS_REF → ISV[WEIGHT_BOUND[group]] - Mech 10 h_s2 clamp: 100 × H_S2_RMS_EMA → ISV[H_S2_BOUND] - AdamW weight_decay (5 kernels): config field → ISV[WD_RATE[group]] - L1 lambda (trunk only): 1e-3 → ISV[L1_LAMBDA_TRUNK_INDEX] DQN main Adam split into 3 per-group sub-launches (DqnTrunk / DqnValue / DqnBranches) per `feedback_no_quickfixes`. Overrides the plan's "max/min-of-3 single-launch shortcut" recommendation. Each sub-launch reads its own WEIGHT_BOUND[group], WD_RATE[group], and (trunk only) L1_LAMBDA_TRUNK_INDEX. Pearl C engagement-counter deferral from A14/A15 resolved in this same commit — per-group split means each sub-launch writes per-block counts at its own offset, and `pearl_c_post_adam_engagement_check` is invoked per group from fused_training.rs (DqnTrunk/DqnValue/DqnBranches separate calls). `weight_decay` field removed from: - DQNHyperparameters (crates/ml/src/trainers/dqn/config.rs) - GpuDqnTrainConfig (crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs) - GpuIqnConfig (crates/ml/src/cuda_pipeline/gpu_iqn_head.rs) - GpuIqlConfig (crates/ml/src/cuda_pipeline/gpu_iql_trainer.rs) - TrialOverrides + PSO search-space (crates/ml/src/training_profile.rs) - apply_family_scaling (`weight_decay *= li` line removed) Aux trainers outside SP4 8-group taxonomy (DT, ofi_embed, denoise, sel/recursive_conf) keep `weight_clamp_max_abs = 0.0` disable — mirrors the existing DT pattern. They have no individual ISV producer, so they don't read SP4 bounds. Files-touched (17): atoms_update_kernel.cu, iql_value_kernel.cu, experience_kernels.cu, dqn_utility_kernels.cu, gpu_dqn_trainer.rs, gpu_iqn_head.rs, gpu_iql_trainer.rs, gpu_attention.rs, gpu_tlob.rs, fused_training.rs, training_loop.rs, constructor.rs, config.rs, generalization.rs (smoke), training_profile.rs, train_baseline_rl.rs, dqn-wire-up-audit.md. Verification (local, RTX 3050 Ti): - `cargo check -p ml --offline`: clean. - `git grep -nE "10\.0_f32 \* q_abs_ref|10\.0f \* q_abs_ref|100\.0_f32 \* q_abs_ref|100\.0f \* q_abs_ref|1e6_f32 \* q_abs_ref|1e3_f32 \* q_abs_ref|100\.0_f32 \* h_s2|100\.0f \* h_s2_rms" crates/ml/src/`: ZERO matches. - `git grep -nE "weight_decay:\s*f64|l1_lambda:" crates/ml/src/trainers/dqn/`: ZERO matches. - `git grep -n "self.config.weight_decay" crates/ml/src/`: only TFT remains (separate trainer outside SP4 scope). - `git grep -n "q_abs_ref_eff|h_s2_rms_ema_eff" crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu`: ZERO matches. - 8 SP4 lib tests pass (sp4_wiener_ema, sp4_isv_slots, state_reset_registry). - 14 SP4 producer GPU tests pass on RTX 3050 Ti (no behavior change at producer level — consumer-side migration only). - `cargo test -p ml --lib --offline`: 928 passed, 14 failed (all 14 pre-existing on HEAD `1389d1c81`; no new failures). Validation deferred to Layer C smoke. Expected: F0/F1/F2 all complete 5 epochs; F1 trains past step 1000; F0 ≥ 37.5; F2 ≥ 55; slot 49 quiet. Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/examples/train_baseline_rl.rs | 3 +- .../src/cuda_pipeline/atoms_update_kernel.cu | 20 +- .../src/cuda_pipeline/dqn_utility_kernels.cu | 147 ++++--- .../src/cuda_pipeline/experience_kernels.cu | 20 +- crates/ml/src/cuda_pipeline/gpu_attention.rs | 13 +- .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 358 +++++++++++------- .../ml/src/cuda_pipeline/gpu_iql_trainer.rs | 22 +- crates/ml/src/cuda_pipeline/gpu_iqn_head.rs | 34 +- crates/ml/src/cuda_pipeline/gpu_tlob.rs | 17 +- .../ml/src/cuda_pipeline/iql_value_kernel.cu | 22 +- crates/ml/src/trainers/dqn/config.rs | 28 +- crates/ml/src/trainers/dqn/fused_training.rs | 164 +++++--- .../dqn/smoke_tests/generalization.rs | 6 +- .../src/trainers/dqn/trainer/constructor.rs | 5 +- .../src/trainers/dqn/trainer/training_loop.rs | 14 +- crates/ml/src/training_profile.rs | 14 +- docs/dqn-wire-up-audit.md | 104 +++++ 17 files changed, 671 insertions(+), 320 deletions(-) diff --git a/crates/ml/examples/train_baseline_rl.rs b/crates/ml/examples/train_baseline_rl.rs index 3e76aaf52..6d3da7ed6 100644 --- a/crates/ml/examples/train_baseline_rl.rs +++ b/crates/ml/examples/train_baseline_rl.rs @@ -444,7 +444,8 @@ fn build_dqn_hyperparams( huber_delta: hp_f32(hp, "huber_delta").unwrap_or(10.0), entropy_coefficient: hp_f64(hp, "entropy_coefficient").unwrap_or(0.01), curiosity_weight: hp_f32(hp, "curiosity_weight").unwrap_or(0.1), - weight_decay: hp_f64(hp, "weight_decay").unwrap_or(1e-4), + // SP4 Layer B: `weight_decay` removed from DQNHyperparameters; per-group + // rates are sourced from ISV[WD_RATE[group]] at runtime. kelly_fractional: hp_f32(hp, "kelly_fractional").unwrap_or(0.5), kelly_max_fraction: hp_f32(hp, "kelly_max_fraction").unwrap_or(0.25), mbp10_data_dir: args.mbp10_data_dir.as_ref().map(|p| p.to_string_lossy().into_owned()).unwrap_or_else(|| "test_data/futures-baseline-mbp10".to_string()), diff --git a/crates/ml/src/cuda_pipeline/atoms_update_kernel.cu b/crates/ml/src/cuda_pipeline/atoms_update_kernel.cu index 208d2d470..16afce46d 100644 --- a/crates/ml/src/cuda_pipeline/atoms_update_kernel.cu +++ b/crates/ml/src/cuda_pipeline/atoms_update_kernel.cu @@ -20,6 +20,11 @@ * with v_min = isv[23+2*b] - isv[24+2*b], v_max = isv[23+2*b] + isv[24+2*b]. */ +/* SP4 Layer B: per-branch atom-position bound from ISV[ATOM_POS_BOUND_BASE + branch]. + * Mirrors `crates/ml/src/cuda_pipeline/sp4_isv_slots.rs::ATOM_POS_BOUND_BASE = 132`. + */ +#define ATOM_POS_BOUND_BASE 132 + extern "C" __global__ void atoms_update( unsigned long long spacing_ptr_b0, /* device ptr to spacing_raw[dir] */ unsigned long long spacing_ptr_b1, /* device ptr to spacing_raw[mag] */ @@ -101,14 +106,13 @@ extern "C" __global__ void atoms_update( float frac = (tid == 0) ? 0.0f : s_cumsum[tid - 1]; float new_pos = v_min + range * frac; - /* SP3 Mech 2: bound atom positions to ±10 × ISV[Q_ABS_REF=16].max(1.0). - * Atoms span the Q range; capping at 10× allows full distribution coverage - * without extreme-tail growth. Same ISV bound as Mech 1 (target_q clip) - * — atoms and target_q share the magnitude scale. ε on the multiplier - * (`isv.max(1.0)`) per the SP1 ε-floor pearl; bound ≥ 10 even with - * cold-start ISV. */ - const float q_abs_ref_eff = fmaxf(isv_signals[16], 1.0f); - const float max_atom_abs = 10.0f * q_abs_ref_eff; + /* SP4 Layer B (Mech 2): bound atom positions to ±ISV[ATOM_POS_BOUND[branch]]. + * Per-branch p99(|atom_positions|) populated by the SP4 producer + * `launch_sp4_atom_pos_p99` (Layer A Task A6) per branch ∈ [0,4). + * ε floor on the bound itself (`max(.., 1.0)`) per the SP1 cold-start + * pearl — pre-first-observation ISV reads as 0 and the consumer must + * not collapse the bound to zero. */ + const float max_atom_abs = fmaxf(isv_signals[ATOM_POS_BOUND_BASE + branch], 1.0f); new_pos = fminf(fmaxf(new_pos, -max_atom_abs), max_atom_abs); out[tid] = new_pos; diff --git a/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu b/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu index c7c647a4f..1c3aa1779 100644 --- a/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu +++ b/crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu @@ -1792,62 +1792,85 @@ extern "C" __global__ void popart_normalize_robust( // // SP3 Mech 5 (Task B6) extension: grid extended from 12 → 24 blocks. Slots // 0-11 (= absolute flag indices 24-35) keep NaN-only check (threshold = +inf -// trivially false). Slots 12-23 (= absolute 36-47) add ISV-derived magnitude -// threshold checks against backward-path Adam state, weights, target_q, and -// atom-positions buffers. Threshold is computed inline per-slot from a single -// `q_abs_ref_eff` arg (no per-slot threshold buffer, no per-step HtoD copies). +// trivially false). Slots 12-25 (= absolute 36-49) add SP4-driven magnitude +// threshold checks against backward-path Adam state, weights, target_q, +// atom-positions, IQN online_params, and post-trunk h_s2 buffers. Each +// threshold reads ISV[SP4_*_BOUND[group/branch]] directly — no magnitude +// multiplier, no per-slot threshold buffer, no per-step HtoD copies. // // SP3 close-out v2 extension: grid extended 24 → 25 blocks. Relative slot 24 // (= absolute 48) covers IQN's separate online_params buffer — the F1 step- // 3540 NaN was traced to this buffer drifting through `iqn_adam_kernel` (its // own param + Adam state, never reached by the single-kernel Mech 9 in -// `dqn_adam_update_kernel`). Same `1e3 × q_abs_ref_eff` weight-slice -// threshold as slots 44-45 — IQN online_params share the trunk weight regime. +// `dqn_adam_update_kernel`). SP4 Layer B reads ISV[WEIGHT_BOUND[Iqn=3]] for +// this slot. // // SP3 Mech 10 extension: grid extended 25 → 26 blocks. Relative slot 25 // (= absolute 49) covers post-trunk h_s2 activation max-abs — the regression -// sentinel for Mech 10's ISV-driven activation clamp. Mech 10 clamps h_s2 at -// 100 × H_S2_RMS_EMA.max(1.0); slot 49 fires at 1e3 × the same base, so it -// stays quiet under normal Mech 10 operation and only triggers if the clamp -// is disabled or `H_S2_RMS_EMA` itself drifted. Mirrors the Mech 5 family -// pattern (clamp at 100×, diagnostic at 1000×). The threshold base is -// `h_s2_rms_ema_eff = max(ISV[H_S2_RMS_EMA_INDEX=96], 1.0)` — distinct from -// `q_abs_ref_eff` because h_s2 lives in pre-readout activation space, not Q -// space. Don't conflate the two ISV bases. +// sentinel for Mech 10's ISV-driven activation clamp. SP4 Layer B reads +// ISV[H_S2_BOUND_INDEX=169] for this slot (no magnitude multiplier; the +// p99 producer kernel `h_s2_p99_update` already tracks the upper bound that +// Mech 10 clamps to). Mirrors the Mech 5 family pattern. // -// Slot-index → threshold mapping (relative slot ≥ 12; absolute = slot + 24): -// 12-15 (abs 36-39): Adam m max-abs ≥ 100 × q_abs_ref_eff -// 16-19 (abs 40-43): Adam v max-abs ≥ 1e6 × q_abs_ref_eff² -// 20-21 (abs 44-45): Weight max-abs ≥ 1e3 × q_abs_ref_eff -// 22 (abs 46) : target_q post-clip ≥ 95 × q_abs_ref_eff -// (= 9.5 × max_abs_target_q, where max_abs_target_q = -// 10 × q_abs_ref_eff per `clamp_target_q_buf`) -// 23 (abs 47) : atom-span max ≥ 190 × q_abs_ref_eff -// (= 9.5 × max_atom_abs × 2, max_atom_abs = 10 × -// q_abs_ref_eff) -// 24 (abs 48) : IQN online_params max-abs ≥ 1e3 × q_abs_ref_eff -// (SP3 close-out v2 — slot-26 GEMM regression sentinel) -// 25 (abs 49) : h_s2 post-clamp max-abs ≥ 1e3 × h_s2_rms_ema_eff -// (SP3 Mech 10 — activation-clamp regression sentinel) +// SP4 Layer B (2026-05-01): per-slot threshold sourced directly from the +// SP4 ISV bound for that slot. No magnitude multipliers; the diagnostic +// fires when |buf[i]| ≥ ISV[bound_idx] (i.e. when a buffer overshoots its +// own clamp threshold). Mech 5's regression-sentinel headroom is now +// supplied by Pearl D's adaptive α tracking the underlying p99 directly. // -// q_abs_ref_eff = max(ISV[Q_ABS_REF=16], 1.0) (slots 12-24) -// h_s2_rms_ema_eff = max(ISV[H_S2_RMS_EMA_INDEX=96], 1.0)(slot 25 only) -// Both computed host-side at launch, passed by-value as float kernel args. -// ε on multiplier per SP1 pearl: cold-start ISV (~0) gives *_eff = 1, so -// all thresholds floor at their ε-floor multiplier × 1. No per-slot -// threshold buffer, no new ISV slots. +// Slot-index → ISV index mapping (relative slot ≥ 12; absolute = slot + 24): +// 12-15 (abs 36-39): Adam m max-abs ≥ ISV[ADAM_M_BOUND[group_idx]] +// group_idx = abs_slot - 36 (groups 0..3) +// 16-19 (abs 40-43): Adam v max-abs ≥ ISV[ADAM_V_BOUND[group_idx]] +// group_idx = abs_slot - 40 (groups 0..3) +// 20-21 (abs 44-45): Weight max-abs ≥ ISV[WEIGHT_BOUND[group_idx]] +// group_idx = abs_slot - 44 (groups 0..1 — trunk/value) +// 22 (abs 46) : target_q post-clip ≥ ISV[TARGET_Q_BOUND] +// 23 (abs 47) : atom-position span ≥ ISV[ATOM_POS_BOUND[branch=0]] +// (any branch is representative since the 4-branch +// diagnostic checks the full positions buffer +// union; branch 0 = direction) +// 24 (abs 48) : IQN online_params max-abs ≥ ISV[WEIGHT_BOUND[Iqn=3]] +// 25 (abs 49) : h_s2 post-clamp max-abs ≥ ISV[H_S2_BOUND] +// +// SP4 ISV slot constants (mirrored from `sp4_isv_slots.rs`): +// ATOM_POS_BOUND_BASE = 132 +// WEIGHT_BOUND_BASE = 136 +// ADAM_M_BOUND_BASE = 144 +// ADAM_V_BOUND_BASE = 152 +// TARGET_Q_BOUND_INDEX = 131 +// H_S2_BOUND_INDEX = 169 // // Invariants preserved from dqn_nan_check_f32: // - sticky-flag semantics (writes 1 only; never clears) // - block-local reduce (no atomicAdd) // - graph-capture safe (pure kernel launch, no host ops) // ============================================================================ + +#ifndef SP4_TARGET_Q_BOUND_INDEX +#define SP4_TARGET_Q_BOUND_INDEX 131 +#endif +#ifndef SP4_ATOM_POS_BOUND_BASE +#define SP4_ATOM_POS_BOUND_BASE 132 +#endif +#ifndef SP4_WEIGHT_BOUND_BASE +#define SP4_WEIGHT_BOUND_BASE 136 +#endif +#ifndef SP4_ADAM_M_BOUND_BASE +#define SP4_ADAM_M_BOUND_BASE 144 +#endif +#ifndef SP4_ADAM_V_BOUND_BASE +#define SP4_ADAM_V_BOUND_BASE 152 +#endif +#ifndef SP4_H_S2_BOUND_INDEX +#define SP4_H_S2_BOUND_INDEX 169 +#endif + extern "C" __global__ void dqn_nan_check_fused_f32_kernel( - const float* const* buf_ptrs, // [N_SLOTS] (26 entries: 12 NaN-only + 14 threshold) - const int* buf_lens, // [N_SLOTS] - float q_abs_ref_eff, // SP3 Mech 5: ISV-derived multiplier (slots 12-24) - float h_s2_rms_ema_eff, // SP3 Mech 10: ISV-derived multiplier (slot 25 only) - int base_flag_idx, // first slot index (24 for backward checks) + const float* const* buf_ptrs, // [N_SLOTS] (26 entries: 12 NaN-only + 14 threshold) + const int* buf_lens, // [N_SLOTS] + const float* __restrict__ isv_signals, // SP4 Layer B: per-slot threshold source (mapped-pinned) + int base_flag_idx, // first slot index (24 for backward checks) int* nan_flags_buf ) { const int slot = blockIdx.x; @@ -1855,39 +1878,39 @@ extern "C" __global__ void dqn_nan_check_fused_f32_kernel( const int n = buf_lens[slot]; if (buf == nullptr || n == 0) return; // deferred slot or unused entry - // SP3 Mech 5: per-slot threshold by relative slot index. Slots 0-11 + // SP4 Layer B: per-slot threshold by relative slot index. Slots 0-11 // (= absolute 24-35) are NaN-only checks (threshold = +inf so the // magnitude check is trivially false). Slots 12-25 (= absolute 36-49) - // check magnitude exceedance per the audit table above. Slots 12-24 - // derive thresholds from q_abs_ref_eff = max(isv[Q_ABS_REF=16], 1.0); - // slot 25 uses h_s2_rms_ema_eff = max(isv[H_S2_RMS_EMA_INDEX=96], 1.0) - // because h_s2 lives in activation space, not Q space. + // read directly from the SP4 ISV bound for that slot. The bound is + // floored at 1.0 to handle pre-first-observation (Pearl A sentinel) + // ISV reads as 0 — same cold-start convention as the Rust consumers. float thr = INFINITY; if (slot >= 12 && slot < 16) { - // slots 36-39: Adam m - thr = 100.0f * q_abs_ref_eff; + // slots 36-39: Adam m, groups 0..3 + const int group_idx = slot - 12; // abs_slot - 36 + thr = fmaxf(isv_signals[SP4_ADAM_M_BOUND_BASE + group_idx], 1.0f); } else if (slot >= 16 && slot < 20) { - // slots 40-43: Adam v - thr = 1.0e6f * q_abs_ref_eff * q_abs_ref_eff; + // slots 40-43: Adam v, groups 0..3 + const int group_idx = slot - 16; // abs_slot - 40 + thr = fmaxf(isv_signals[SP4_ADAM_V_BOUND_BASE + group_idx], 1.0f); } else if (slot >= 20 && slot < 22) { - // slots 44-45: weight slices (trunk / heads) - thr = 1.0e3f * q_abs_ref_eff; + // slots 44-45: weight slices, groups 0..1 (trunk / value) + const int group_idx = slot - 20; // abs_slot - 44 + thr = fmaxf(isv_signals[SP4_WEIGHT_BOUND_BASE + group_idx], 1.0f); } else if (slot == 22) { - // slot 46: target_q post-clip (max_abs_target_q = 10 × q_abs_ref_eff) - thr = 95.0f * q_abs_ref_eff; + // slot 46: target_q post-clip + thr = fmaxf(isv_signals[SP4_TARGET_Q_BOUND_INDEX], 1.0f); } else if (slot == 23) { - // slot 47: atom-positions span (max_atom_abs = 10 × q_abs_ref_eff) - thr = 190.0f * q_abs_ref_eff; + // slot 47: atom-positions (branch 0 = direction; positions buffer + // union diagnostic — any branch's bound is representative). + thr = fmaxf(isv_signals[SP4_ATOM_POS_BOUND_BASE], 1.0f); } else if (slot == 24) { - // slot 48 (SP3 close-out v2): IQN online_params — same weight-slice - // regime as slots 44-45. Captures the slot-26 GEMM blind spot. - thr = 1.0e3f * q_abs_ref_eff; + // slot 48 (SP3 close-out v2): IQN online_params — uses the IQN + // (group 3) weight bound. + thr = fmaxf(isv_signals[SP4_WEIGHT_BOUND_BASE + 3], 1.0f); } else if (slot == 25) { - // slot 49 (SP3 Mech 10): post-trunk h_s2 max-abs sentinel. Mech 10 - // clamps h_s2 at 100 × h_s2_rms_ema_eff; this fires at 1e3 ×, so - // it never trips under normal operation. If it does, the clamp was - // disabled or H_S2_RMS_EMA itself drifted. - thr = 1.0e3f * h_s2_rms_ema_eff; + // slot 49 (SP3 Mech 10): post-trunk h_s2 max-abs sentinel. + thr = fmaxf(isv_signals[SP4_H_S2_BOUND_INDEX], 1.0f); } int flag_set = 0; diff --git a/crates/ml/src/cuda_pipeline/experience_kernels.cu b/crates/ml/src/cuda_pipeline/experience_kernels.cu index 55637c657..a893c6dde 100644 --- a/crates/ml/src/cuda_pipeline/experience_kernels.cu +++ b/crates/ml/src/cuda_pipeline/experience_kernels.cu @@ -165,6 +165,13 @@ __device__ __forceinline__ int argmax_n(const float* arr, int n) { #define N_IQN_QUANTILES 5 #endif +/* SP4 Layer B: per-branch atom-position bound from ISV[ATOM_POS_BOUND_BASE + branch]. + * Mirrors `crates/ml/src/cuda_pipeline/sp4_isv_slots.rs::ATOM_POS_BOUND_BASE = 132`. + */ +#ifndef ATOM_POS_BOUND_BASE +#define ATOM_POS_BOUND_BASE 132 +#endif + /** * Inverse-CDF sample from a C51 atom distribution. * @@ -5974,12 +5981,13 @@ extern "C" __global__ void adaptive_atom_positions( float frac = (tid == 0) ? 0.0f : s_cumsum[tid - 1]; float new_pos = v_min + range * frac; - /* SP3 Mech 2: bound atom positions to ±10 × ISV[Q_ABS_REF=16].max(1.0). - * Mirrors atoms_update_kernel.cu — same ISV bound as Mech 1 (target_q - * clip). Kept in lockstep with the active atoms_update kernel so this - * legacy single-branch entry point stays bound-consistent if re-wired. */ - const float q_abs_ref_eff = fmaxf(isv_signals[16], 1.0f); - const float max_atom_abs = 10.0f * q_abs_ref_eff; + /* SP4 Layer B (Mech 2): bound atom positions to ±ISV[ATOM_POS_BOUND[branch_idx]]. + * Per-branch p99(|atom_positions|) populated by SP4 producer + * `launch_sp4_atom_pos_p99` (Layer A Task A6). Mirrors atoms_update_kernel.cu; + * kept in lockstep with the active atoms_update kernel so this legacy + * single-branch entry point stays bound-consistent if re-wired. + * ε floor on the bound itself (`max(.., 1.0)`) per the SP1 cold-start pearl. */ + const float max_atom_abs = fmaxf(isv_signals[ATOM_POS_BOUND_BASE + branch_idx], 1.0f); new_pos = fminf(fmaxf(new_pos, -max_atom_abs), max_atom_abs); positions_out[tid] = new_pos; diff --git a/crates/ml/src/cuda_pipeline/gpu_attention.rs b/crates/ml/src/cuda_pipeline/gpu_attention.rs index 7801cf8d6..3d107f5f0 100644 --- a/crates/ml/src/cuda_pipeline/gpu_attention.rs +++ b/crates/ml/src/cuda_pipeline/gpu_attention.rs @@ -1057,10 +1057,13 @@ impl GpuAttention { max_grad_norm: f32, override_stream: Option<&Arc>, override_workspace: Option<(u64, usize)>, - /* SP3 Mech 9 (close-out v2): ISV-driven |p| bound for `attn_adam_kernel` - * — caller computes `100 × ISV[Q_ABS_REF].max(1.0)` host-side. + /* SP4 Layer B (Mech 9): ISV-driven |p| bound for `attn_adam_kernel` + * — caller reads `ISV[WEIGHT_BOUND[Attn=6]]` with EPS_CLAMP_FLOOR ε. * 0=disabled (debug only). */ weight_clamp_max_abs: f32, + /* SP4 Layer B (Mech 9): ISV-driven AdamW weight_decay rate. + * Caller reads `ISV[WD_RATE[Attn=6]]` with `.max(0.0)`. */ + weight_decay_value: f32, ) -> Result<(), MLError> { let stream = override_stream.unwrap_or(&self.stream); // override_workspace is accepted for API symmetry; adam_step does not use @@ -1111,7 +1114,11 @@ impl GpuAttention { let beta1 = 0.9_f32; let beta2 = 0.999_f32; let eps = 1e-8_f32; - let wd = 1e-5_f32; + // SP4 Layer B (Mech 9): AdamW weight_decay sourced from + // ISV[WD_RATE[Attn=6]] by the caller (`FusedTrainingCtx`) and + // threaded through `adam_step`. `.max(0.0)` upstream enforces the + // AdamW non-negativity contract. + let wd: f32 = weight_decay_value; let t_ptr = self.t_dev_ptr; // SP4 Task A14: Pearl C engagement-counter args. Both diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 78aa42479..6117f91fc 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -1537,8 +1537,11 @@ pub struct GpuDqnTrainConfig { pub beta2: f32, /// Adam ε (numerical stability). pub epsilon: f32, - /// Decoupled weight decay (AdamW). - pub weight_decay: f32, + // SP4 Layer B: `weight_decay` field removed; per-group AdamW rates are + // sourced from ISV[WD_RATE[group]] inside `launch_adam_update`'s 3-way + // sub-launch loop (DqnTrunk / DqnValue / DqnBranches) and threaded + // through to IQN/IQL/Attn/Curiosity launchers via per-launch + // `weight_decay_value: f32` parameters. /// Maximum gradient L2 norm for clipping. pub max_grad_norm: f32, /// Spectral norm clipping target (σ_max). Constrains ||W||_σ ≤ σ_max. @@ -1659,7 +1662,7 @@ impl Default for GpuDqnTrainConfig { beta1: 0.9, beta2: 0.999, epsilon: 1e-8, - weight_decay: 1e-4, // Standard L2 regularization strength + // SP4 Layer B: weight_decay field removed; sourced from ISV at runtime. max_grad_norm: 10000.0, // Safety-only — per-component clip removed, raw norms ~4000 spectral_norm_sigma_max: 3.0, spectral_decoupling_lambda: 0.01, @@ -6146,9 +6149,13 @@ impl GpuDqnTrainer { let wd_mask_ptr = self.weight_decay_mask.raw_ptr(); let l1_end_aux: i32 = 0; let l1_lambda_aux: f32 = 0.0; - // SP3 Mech 9: ISV-driven post-Adam weight clamp (see `launch_adam_update`). - let q_abs_ref_eff = self.read_isv_signal_at(Q_ABS_REF_INDEX).max(1.0_f32); - let weight_clamp_max_abs: f32 = 100.0_f32 * q_abs_ref_eff; + // SP4 Layer B: aux trainer (OFI embed) is outside the SP4 8-group + // taxonomy — disable Mech 9 weight clamp here (mirrors DT pattern). + // Setting `weight_clamp_max_abs = 0.0` triggers the kernel's + // `if (bound > 0.0f)` short-circuit; Adam behaves as before for + // ofi_embed without unbounded ISV-driven clamps that would + // require an SP4 producer this aux head doesn't have. + let weight_clamp_max_abs: f32 = 0.0_f32; // SP4 Task A14: aux trainer (OFI embed) is outside the SP4 8-group // taxonomy — disable Pearl C engagement counter for this launch. let _aux_nan_flags_null: u64 = 0; @@ -7488,6 +7495,24 @@ impl GpuDqnTrainer { self.total_params } + /// SP4 Layer B: per-group param counts for the 3 DQN-internal groups + /// (Trunk / Value / Branches). Mirrors the byte-offset slicing in + /// `param_group_buffers` and `launch_adam_update`. Returned counts feed + /// `pearl_c_post_adam_engagement_check` per group, replacing Layer A's + /// "total_params under group 0 only" approximation. + pub fn dqn_group_param_counts(&self) -> (usize, usize, usize) { + let f32_sz = std::mem::size_of::() as u64; + let param_sizes = compute_param_sizes(&self.config); + let trunk = self.trunk_param_count; + let value_off = padded_byte_offset(¶m_sizes, 13); + let value_end = padded_byte_offset(¶m_sizes, 17); + let value = ((value_end - value_off) / f32_sz) as usize; + let branches_off = padded_byte_offset(¶m_sizes, 17); + let branches_end = padded_byte_offset(¶m_sizes, 33); + let branches = ((branches_end - branches_off) / f32_sz) as usize; + (trunk, value, branches) + } + /// Apply IQN auxiliary gradient to the shared trunk via SAXPY into `grad_buf`. /// /// After `graph_forward` replay (C51 forward → backward) and IQN backward @@ -17070,21 +17095,11 @@ impl GpuDqnTrainer { let nan_flags_ptr = self.nan_flags_buf.raw_ptr(); let buf_ptrs_dev = self.nan_check_buf_ptrs.dev_ptr; let buf_lens_dev = self.nan_check_buf_lens.dev_ptr; - // SP3 Mech 5: read ISV[Q_ABS_REF=16] for slot 36-48 threshold - // computation. Inline ε on multiplier (`max(.., 1.0)`) per the SP1 - // pearl — cold-start ISV (~0) gives q_abs_ref_eff = 1, so all - // thresholds floor at their ε-floor multiplier × 1. Pass as f32 to - // match the kernel's `float q_abs_ref_eff` arg (cudarc f64→float ABI - // bug pearl: passing f64 to a `float` param reads only low 4 bytes). - let q_abs_ref = self.read_isv_signal_at(Q_ABS_REF_INDEX); - let q_abs_ref_eff: f32 = q_abs_ref.max(1.0_f32); - // SP3 Mech 10: read ISV[H_S2_RMS_EMA_INDEX=96] for slot 49 threshold - // (= 1e3 × h_s2_rms_ema_eff). Distinct ISV base from `q_abs_ref_eff` - // because h_s2 lives in activation space, not Q space — don't conflate - // the two. Same ε floor convention (`max(.., 1.0)`) so cold-start RMS - // (~0) gives h_s2_rms_ema_eff = 1 (threshold floors at 1e3). - let h_s2_rms_ema = self.read_isv_signal_at(H_S2_RMS_EMA_INDEX); - let h_s2_rms_ema_eff: f32 = h_s2_rms_ema.max(1.0_f32); + // SP4 Layer B: per-slot threshold sourced directly from the SP4 ISV + // bound for that slot. Kernel reads `isv_signals[SP4_*_BOUND[...]]` + // inline — no host-side multiplier, no per-step HtoD copy. The + // mapped-pinned ISV buffer is graph-capture-safe via its dev_ptr. + let isv_signals_dev_ptr: u64 = self.isv_signals_dev_ptr; const BLOCKS: u32 = 26; const THREADS: u32 = 256; const BASE_FLAG_IDX: i32 = 24; @@ -17098,8 +17113,7 @@ impl GpuDqnTrainer { .launch_builder(&self.nan_check_fused_f32_kernel) .arg(&buf_ptrs_dev) .arg(&buf_lens_dev) - .arg(&q_abs_ref_eff) - .arg(&h_s2_rms_ema_eff) + .arg(&isv_signals_dev_ptr) .arg(&BASE_FLAG_IDX) .arg(&nan_flags_ptr) .launch(cfg) @@ -17870,18 +17884,19 @@ impl GpuDqnTrainer { .map_err(|e| MLError::ModelError(format!("compute_denoise_target_q expected_q: {e}")))?; } - // SP3 Mech 1: clip target_q to ±10 × ISV[Q_ABS_REF=16].max(1.0). - // Hard clamp via dqn_clamp_finite_f32_kernel (reused from SP1). Single - // point — all downstream losses (C51 / IQN / MSE) consume the same - // denoise_target_q_buf, so one clip covers every consumer. ISV-driven - // bound per feedback_isv_for_adaptive_bounds; ε on the multiplier - // (`isv.max(1.0)`) per the SP1 ε-floor pearl. - // - // F0 paper-review: F0-typical |target_q| bounded by reward × horizon - // ≤ ~10. With Q_ABS_REF EMA ≈ 1-5 at F0 convergence, max_abs = 10-50. - // F0 guard is no-op. F1 inflation (thousands+) gets clamped. - let q_abs_ref = self.read_isv_signal_at(Q_ABS_REF_INDEX); - let max_abs_target_q = 10.0_f32 * q_abs_ref.max(1.0_f32); + // SP4 Layer B (Mech 1): clip target_q to ±ISV[TARGET_Q_BOUND=131]. + // Hard clamp via dqn_clamp_finite_f32_kernel (reused from SP1). + // Single point — all downstream losses (C51 / IQN / MSE) consume the + // same denoise_target_q_buf, so one clip covers every consumer. + // SP4 producer `launch_sp4_target_q_p99` (Layer A Task A5) tracks + // p99(|target_q|) into the ISV bound directly; the multiplier-driven + // bound from SP3 (`10 × Q_ABS_REF.max(1.0)`) is replaced with this + // adaptive p99 bound per `feedback_no_quickfixes` and + // `feedback_isv_for_adaptive_bounds`. EPS_CLAMP_FLOOR=1.0 ε on the + // bound itself (SP1 ε-floor pearl) handles cold-start ISV (= 0). + use crate::cuda_pipeline::sp4_isv_slots::TARGET_Q_BOUND_INDEX; + use crate::cuda_pipeline::sp4_wiener_ema::EPS_CLAMP_FLOOR; + let max_abs_target_q = self.read_isv_signal_at(TARGET_Q_BOUND_INDEX).max(EPS_CLAMP_FLOOR); self.launch_clamp_finite_f32( self.denoise_target_q_buf.raw_ptr(), self.denoise_target_q_buf.len(), @@ -18554,9 +18569,9 @@ impl GpuDqnTrainer { let wd_mask_ptr = self.weight_decay_mask.raw_ptr(); let l1_end_aux: i32 = 0; let l1_lambda_aux: f32 = 0.0; - // SP3 Mech 9: ISV-driven post-Adam weight clamp (see `launch_adam_update`). - let q_abs_ref_eff = self.read_isv_signal_at(Q_ABS_REF_INDEX).max(1.0_f32); - let weight_clamp_max_abs: f32 = 100.0_f32 * q_abs_ref_eff; + // SP4 Layer B: aux trainer (denoise) is outside the SP4 8-group + // taxonomy — disable Mech 9 weight clamp here (mirrors DT pattern). + let weight_clamp_max_abs: f32 = 0.0_f32; // SP4 Task A14: aux trainer (denoise) is outside the SP4 8-group // taxonomy — Pearl C engagement counter disabled for this launch. let _aux_nan_flags_null: u64 = 0; @@ -18707,15 +18722,11 @@ impl GpuDqnTrainer { // consumer (loss path, IQN, value head, replay save, eval-action // select) reads the buffer. // - // Bound: `100 × ISV[H_S2_RMS_EMA_INDEX=96].max(1.0)`. ISV slot 96 - // already tracks the per-batch RMS EMA of h_s2 (per Plan 4 Task - // 2c.3c.5 — same ISV slot consumed by `mag_concat_qdir` adaptive - // scale). Reusing the existing slot keeps Mech 10 consistent with - // `feedback_isv_for_adaptive_bounds.md` (no new ISV slots, no - // hardcoded multipliers). The `max(1.0)` ε floor on the multiplier - // honours the SP1 cold-start convention — at fold boundary the EMA - // resets to 1.0 (neutral RMS) per `state_reset_registry.rs`, so the - // clamp floor is always at least 100 × 1 = 100, never zero. + // SP4 Layer B (Mech 10): bound = ISV[H_S2_BOUND=169]. Producer + // `launch_sp4_h_s2_p99` (Layer A Task A9) tracks p99(|save_h_s2|) + // through Pearls A+D into ISV[169] directly — no magnitude + // multiplier, no separate RMS scaling. The EPS_CLAMP_FLOOR ε on + // the bound itself preserves SP1 cold-start convention. // // Diagnostic ordering (SP1 pearl): the inline NaN check on slot 12 // fires BEFORE the clamp sanitises, so the sentinel sees the original @@ -18731,9 +18742,10 @@ impl GpuDqnTrainer { // The clamp adds a single kernel launch in the captured graph // (expected — captured graph topology grows by one node). { + use crate::cuda_pipeline::sp4_isv_slots::H_S2_BOUND_INDEX; + use crate::cuda_pipeline::sp4_wiener_ema::EPS_CLAMP_FLOOR; self.check_nan_f32(self.save_h_s2.raw_ptr(), self.save_h_s2.len(), 12)?; - let h_s2_rms_ema_eff = self.read_isv_signal_at(H_S2_RMS_EMA_INDEX).max(1.0_f32); - let max_abs_h_s2 = 100.0_f32 * h_s2_rms_ema_eff; + let max_abs_h_s2 = self.read_isv_signal_at(H_S2_BOUND_INDEX).max(EPS_CLAMP_FLOOR); self.launch_clamp_finite_f32( self.save_h_s2.raw_ptr(), self.save_h_s2.len(), @@ -21071,25 +21083,44 @@ impl GpuDqnTrainer { /// Argument order matches `dqn_adam_update_kernel` in `dqn_utility_kernels.cu`: /// params, grads, m, v, grad_norm_sq, lr, beta1, beta2, epsilon, /// weight_decay, max_grad_norm, t_ptr, total_params, - /// weight_decay_mask, l1_end, l1_lambda, weight_clamp_max_abs = 17 args. + /// weight_decay_mask, l1_end, l1_lambda, weight_clamp_max_abs, + /// nan_flags_buf, diag_slot, clamp_engage_per_block_buf, engage_buf_offset + /// = 21 args. /// /// Must be launched AFTER `launch_grad_norm` so that `grad_norm_buf` /// contains the completed sum of squares (no race condition). /// `t_ptr` is a device buffer (not scalar) so the CUDA Graph can be /// replayed with an updated step counter. /// - /// Single Adam launch covers ALL parameters `[0..total_params)`, - /// including the VSN tensors at indices `[95..119)`. VSN gradient noise + /// **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. + /// 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, + }; + 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; let beta2 = self.config.beta2; let epsilon = self.config.epsilon; - let weight_decay = self.config.weight_decay; let adaptive_clip_ptr = self.ptrs.adaptive_clip_buf; let norm_ptr = self.ptrs.grad_norm_buf; let t_ptr = self.ptrs.t_buf; @@ -21097,84 +21128,129 @@ impl GpuDqnTrainer { // G1+G8: weight_decay_mask + L1 proximal on w_s1. let wd_mask_base = self.weight_decay_mask.raw_ptr(); let param_sizes = compute_param_sizes(&self.config); - let l1_end: i32 = align4(param_sizes[0]) as i32; // end of w_s1 - let l1_lambda: f32 = 1e-3; // L1 penalty strength on w_s1 + let trunk_l1_end: i32 = align4(param_sizes[0]) as i32; // end of w_s1 - // SP3 Mech 9: ISV-driven post-Adam weight clamp. - // Bound = 100 × Q_ABS_REF.max(1.0). ε on the multiplier per the - // SP1 pearl (`isv.max(1.0)` — never floor the bound itself, that - // re-introduces the cold-start clamp pathology). One OOM below the - // slot 44-45 diagnostic threshold (1e3 × Q_ABS_REF) so the - // regression sentinel retains 10× firing headroom. - let q_abs_ref_eff = self.read_isv_signal_at(Q_ABS_REF_INDEX).max(1.0_f32); - let weight_clamp_max_abs: f32 = 100.0_f32 * q_abs_ref_eff; - - let count = self.total_params as i32; - // SP4 Task A14: cap blocks at MAX_BLOCKS_PER_ADAM (256). Per-block - // engagement-counter buffer is sized for that ceiling; running with - // more blocks would have the kernel write past the per-group slice. - // For DQN main-trainer total_params can exceed 256×256=65 536, so we - // assert the bound here. If it ever fires, the buffer needs sizing - // up before this Adam launch can grow. - let blocks = ((self.total_params + 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 \ - or chunk the launch.", - blocks, MAX_BLOCKS_PER_ADAM - ); - let cfg = LaunchConfig { - grid_dim: (blocks, 1, 1), - block_dim: (256, 1, 1), - shared_mem_bytes: 0, - }; - let params_ptr = self.ptrs.params_ptr; - let grad_ptr = self.ptrs.grad_buf; - let m_ptr = self.ptrs.m_buf; - let v_ptr = self.ptrs.v_buf; - - // SP4 Task A14: Pearl C engagement-counter args. The single main - // DQN Adam launch covers param-groups 0/1/2 (trunk/value/branches); - // we account engagement under group 0 (DqnTrunk) for Layer A — see - // `pearl_c_post_adam_engagement_check` doc-comment for the - // sub-launch deferral rationale. + // SP4 Layer B (Mech 9 + Pearl C): per-group sub-launches replace the + // 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. 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; - let engage_buf_offset: i32 = - (ParamGroup::DqnTrunk.idx() * MAX_BLOCKS_PER_ADAM) as i32; + + // 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 + // threading `SP4AuxBuffers` through this launcher (DQN-internal + // groups never need aux pointers). + let f32_sz = std::mem::size_of::() as u64; + let trunk_count = self.trunk_param_count; + let trunk_byte_off: u64 = 0; + let value_byte_off = padded_byte_offset(¶m_sizes, 13); + let value_byte_end = padded_byte_offset(¶m_sizes, 17); + let value_count = ((value_byte_end - value_byte_off) / f32_sz) as usize; + 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; + + 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), + ]; // Safety: argument order matches the extern "C" kernel signature exactly. // grad_norm_buf contains L2 norm from launch_grad_norm_finalize(). - unsafe { - self.stream - .launch_builder(&self.adam_update_kernel) - .arg(¶ms_ptr) - .arg(&grad_ptr) - .arg(&m_ptr) - .arg(&v_ptr) - .arg(&norm_ptr) - .arg(&lr_ptr) - .arg(&beta1) - .arg(&beta2) - .arg(&epsilon) - .arg(&weight_decay) - .arg(&adaptive_clip_ptr) // device buffer — EMA-based adaptive clip - .arg(&t_ptr) // device pointer — not baked scalar - .arg(&count) - .arg(&wd_mask_base) // G1: per-param weight decay mask - .arg(&l1_end) // G8: L1 boundary (end of w_s1) - .arg(&l1_lambda) // G8: L1 penalty strength - .arg(&weight_clamp_max_abs) // SP3 Mech 9: post-Adam |p_val| clamp - .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 A14: group offset (0 = DqnTrunk) - .launch(cfg) - .map_err(|e| { - MLError::ModelError(format!("dqn_adam_update_kernel launch: {e}")) - })?; + for (group, byte_off, count) 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); + + // 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 { + let l1_lambda = self.read_isv_signal_at(L1_LAMBDA_TRUNK_INDEX).max(0.0); + (trunk_l1_end, l1_lambda) + } else { + (0_i32, 0.0_f32) + }; + + // Sub-buffer pointers via byte-offset — same arithmetic as + // `param_group_buffers`. + let params_ptr = self.ptrs.params_ptr + byte_off; + let grad_ptr = self.ptrs.grad_buf + byte_off; + let m_ptr = self.m_buf.raw_ptr() + byte_off; + let v_ptr = self.v_buf.raw_ptr() + byte_off; + // wd_mask is a per-param byte-aligned buffer matching params layout. + let wd_mask_ptr = wd_mask_base + byte_off; + + let count_i32 = count as i32; + 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 \ + or chunk the launch.", + group, blocks, MAX_BLOCKS_PER_ADAM + ); + let cfg = LaunchConfig { + grid_dim: (blocks, 1, 1), + block_dim: (256, 1, 1), + 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) + .arg(¶ms_ptr) + .arg(&grad_ptr) + .arg(&m_ptr) + .arg(&v_ptr) + .arg(&norm_ptr) + .arg(&lr_ptr) + .arg(&beta1) + .arg(&beta2) + .arg(&epsilon) + .arg(&weight_decay_value) // SP4 Layer B: per-group ISV-driven + .arg(&adaptive_clip_ptr) // device buffer — EMA-based adaptive clip + .arg(&t_ptr) // device pointer — not baked scalar + .arg(&count_i32) + .arg(&wd_mask_ptr) // G1: per-param weight decay mask (sub-slice) + .arg(&l1_end_arg) // G8: L1 boundary (trunk only) + .arg(&l1_lambda_arg) // G8: SP4 Layer B: ISV-driven for trunk; 0 elsewhere + .arg(&weight_clamp_max_abs) // SP4 Layer B: per-group ISV-driven + .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) + .launch(cfg) + .map_err(|e| { + MLError::ModelError(format!("dqn_adam_update_kernel[{:?}] launch: {e}", group)) + })?; + } } Ok(()) @@ -21198,10 +21274,14 @@ impl GpuDqnTrainer { /// engagement rate. All other groups read the single contiguous /// `[group_idx × MAX_BLOCKS_PER_ADAM, +MAX_BLOCKS_PER_ADAM)` range. /// - /// **Layer A limitation #3** (groups 1/2 accounting): the main DQN - /// Adam covers param groups 0 (Trunk), 1 (Value), and 2 (Branches) - /// in a single fused launch — engagement is recorded under group 0 - /// only. Layer B will split per-group sub-launches. + /// **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 + /// `WEIGHT_BOUND[group]` consumer migration. pub(crate) fn pearl_c_post_adam_engagement_check( &mut self, group: ParamGroup, @@ -22176,9 +22256,19 @@ impl GpuDqnTrainer { // 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. - let abs_mult = self.read_isv_signal_at(Q_ABS_REF_INDEX).max(1.0_f32); - let prev_slow_ema = unsafe { *self.grad_norm_slow_ema_pinned }; - let upper_bound = (prev_slow_ema * 100.0_f32 * abs_mult).max(MIN_CLIP); + // + // 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 + // Pearl D–smoothed p99(grad_norm) into ISV[168] directly; the + // hardcoded `100 × prev_slow_ema × Q_ABS_REF.max(1.0)` bound from + // SP3 is replaced with this signal-driven upper-bound. The + // EPS_CLAMP_FLOOR ε on the bound itself preserves cold-start + // semantics (pre-first-observation ISV reads as 0). Note: + // `grad_norm_slow_ema_pinned` is no longer read by this consumer + // — its retirement is deferred to Layer C per the audit doc. + use crate::cuda_pipeline::sp4_isv_slots::GRAD_CLIP_BOUND_INDEX; + use crate::cuda_pipeline::sp4_wiener_ema::EPS_CLAMP_FLOOR; + let upper_bound = self.read_isv_signal_at(GRAD_CLIP_BOUND_INDEX).max(EPS_CLAMP_FLOOR); let new_clip = new_clip.min(upper_bound); // Write directly to pinned memory — GPU sees it on next kernel read @@ -22773,9 +22863,9 @@ impl GpuDqnTrainer { let wd_mask_ptr = self.weight_decay_mask.raw_ptr(); let l1_end_aux: i32 = 0; let l1_lambda_aux: f32 = 0.0; - // SP3 Mech 9: ISV-driven post-Adam weight clamp (see `launch_adam_update`). - let q_abs_ref_eff = self.read_isv_signal_at(Q_ABS_REF_INDEX).max(1.0_f32); - let weight_clamp_max_abs: f32 = 100.0_f32 * q_abs_ref_eff; + // SP4 Layer B: aux trainer (sel/recursive_conf) is outside the SP4 + // 8-group taxonomy — disable Mech 9 weight clamp here (mirrors DT pattern). + let weight_clamp_max_abs: f32 = 0.0_f32; // SP4 Task A14: aux trainer (sel/recursive_conf) is outside the SP4 // 8-group taxonomy — Pearl C engagement counter disabled. let _aux_nan_flags_null: u64 = 0; diff --git a/crates/ml/src/cuda_pipeline/gpu_iql_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_iql_trainer.rs index 26f816dc2..712a8be8a 100644 --- a/crates/ml/src/cuda_pipeline/gpu_iql_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_iql_trainer.rs @@ -96,8 +96,9 @@ pub struct GpuIqlConfig { pub beta2: f32, /// Adam epsilon (numerical stability). pub epsilon: f32, - /// Decoupled weight decay (AdamW). - pub weight_decay: f32, + // SP4 Layer B: `weight_decay` field removed; per-step AdamW rate is + // sourced from ISV[WD_RATE[group]] (group = IqlHigh=4 or IqlLow=5) + // and threaded into `train_value_step` as `weight_decay_value`. /// Maximum gradient L2 norm for clipping. pub max_grad_norm: f32, /// Number of C51 atoms. @@ -122,7 +123,7 @@ impl Default for GpuIqlConfig { beta1: 0.9, beta2: 0.999, epsilon: 1e-8, - weight_decay: 1e-5, + // SP4 Layer B: weight_decay field removed; sourced from ISV at runtime. max_grad_norm: 1.0, num_atoms: 51, total_actions: 108, @@ -580,10 +581,13 @@ impl GpuIqlTrainer { pub fn train_value_step( &mut self, states_f32: &CudaSlice, - /* SP3 Mech 9 (close-out v2): ISV-driven |p| bound for `iql_adam_kernel` - * — caller computes `100 × ISV[Q_ABS_REF].max(1.0)` host-side. - * 0=disabled (debug only). */ + /* SP4 Layer B (Mech 9): ISV-driven |p| bound for `iql_adam_kernel` + * — caller reads `ISV[WEIGHT_BOUND[group]]` (group = IqlHigh=4 or + * IqlLow=5) with EPS_CLAMP_FLOOR ε. 0=disabled (debug only). */ weight_clamp_max_abs: f32, + /* SP4 Layer B (Mech 9): ISV-driven AdamW weight_decay rate. + * Caller reads `ISV[WD_RATE[group]]` with `.max(0.0)` non-negativity. */ + weight_decay_value: f32, ) -> Result<(), MLError> { let b = self.config.batch_size; let h = self.config.value_hidden_dim; @@ -1035,7 +1039,11 @@ impl GpuIqlTrainer { let beta1 = self.config.beta1; let beta2 = self.config.beta2; let eps = self.config.epsilon; - let wd = self.config.weight_decay; + // SP4 Layer B (Mech 9): AdamW weight_decay sourced from + // ISV[WD_RATE[IqlHigh|IqlLow]] by the caller (`FusedTrainingCtx`) + // and threaded through `train_value_step`. `.max(0.0)` upstream + // enforces the AdamW non-negativity contract. + let wd: f32 = weight_decay_value; let mgn = self.config.max_grad_norm; let t_ptr = self.t_dev_ptr; diff --git a/crates/ml/src/cuda_pipeline/gpu_iqn_head.rs b/crates/ml/src/cuda_pipeline/gpu_iqn_head.rs index ea299eccf..e8f311ecd 100644 --- a/crates/ml/src/cuda_pipeline/gpu_iqn_head.rs +++ b/crates/ml/src/cuda_pipeline/gpu_iqn_head.rs @@ -100,8 +100,9 @@ pub struct GpuIqnConfig { pub beta2: f32, /// Adam epsilon (numerical stability). pub epsilon: f32, - /// Decoupled weight decay (AdamW). - pub weight_decay: f32, + // SP4 Layer B: `weight_decay` field removed; per-step AdamW rate is + // sourced from ISV[WD_RATE[Iqn=3]] and threaded into + // `execute_training_pipeline` as `weight_decay_value`. /// Maximum gradient L2 norm for clipping. pub max_grad_norm: f32, /// Discount factor for Bellman projection. @@ -128,7 +129,7 @@ impl Default for GpuIqnConfig { beta1: 0.9, beta2: 0.999, epsilon: 1e-8, - weight_decay: 1e-5, + // SP4 Layer B: weight_decay field removed; sourced from ISV at runtime. max_grad_norm: 1.0, gamma: 0.99, } @@ -816,11 +817,13 @@ impl GpuIqnHead { dqn_actions_buf: &CudaSlice, dqn_rewards_buf: &CudaSlice, dqn_dones_buf: &CudaSlice, - /* SP3 Mech 9 (close-out v2): ISV-derived |p| bound for `iqn_adam_kernel` - * — caller computes `100 × ISV[Q_ABS_REF].max(1.0)` and passes it - * through. 0=disabled (debug only); production callers always set - * a non-zero bound. */ + /* SP4 Layer B (Mech 9): ISV-derived |p| bound for `iqn_adam_kernel` + * — caller reads `ISV[WEIGHT_BOUND[Iqn=3]]` and passes it through. + * 0=disabled (debug only); production callers always set a non-zero + * bound. */ weight_clamp_max_abs: f32, + /* SP4 Layer B (Mech 9): ISV-derived AdamW weight_decay rate. */ + weight_decay_value: f32, ) -> Result { // 1-2. Decode actions + DtoD copy rewards/dones self.prepare_buffers(dqn_actions_buf, dqn_rewards_buf, dqn_dones_buf)?; @@ -833,6 +836,7 @@ impl GpuIqnHead { None, None, weight_clamp_max_abs, + weight_decay_value, ) } @@ -868,10 +872,13 @@ impl GpuIqnHead { target_dueling: &DuelingWeightSet, override_stream: Option<&Arc>, override_workspace: Option<(u64, usize)>, - /* SP3 Mech 9 (close-out v2): ISV-derived |p| bound for the trailing - * `iqn_adam_kernel` arg. Caller computes `100 × ISV[Q_ABS_REF].max(1.0)` - * host-side. 0=disabled (debug only). */ + /* SP4 Layer B (Mech 9): ISV-derived |p| bound for the trailing + * `iqn_adam_kernel` arg. Caller reads `ISV[WEIGHT_BOUND[Iqn=3]]` + * with EPS_CLAMP_FLOOR ε. 0=disabled (debug only). */ weight_clamp_max_abs: f32, + /* SP4 Layer B (Mech 9): ISV-derived AdamW weight_decay rate. + * Caller reads `ISV[WD_RATE[Iqn=3]]` with `.max(0.0)` non-negativity. */ + weight_decay_value: f32, ) -> Result { let effective_stream = override_stream.unwrap_or(&self.stream); let (lt_ws_ptr, lt_ws_size) = override_workspace @@ -1495,7 +1502,12 @@ impl GpuIqnHead { let beta1 = self.config.beta1; let beta2 = self.config.beta2; let eps = self.config.epsilon; - let wd = self.config.weight_decay; + // SP4 Layer B (Mech 9): AdamW weight_decay sourced from + // ISV[WD_RATE[Iqn=3]] by the caller (`FusedTrainingCtx`) and + // threaded through `execute_training_pipeline` alongside + // `weight_clamp_max_abs`. `.max(0.0)` upstream enforces the + // AdamW non-negativity contract. + let wd = weight_decay_value; let mgn = self.config.max_grad_norm; let t_ptr = self.t_dev_ptr; diff --git a/crates/ml/src/cuda_pipeline/gpu_tlob.rs b/crates/ml/src/cuda_pipeline/gpu_tlob.rs index 92f854e69..3ff6164c6 100644 --- a/crates/ml/src/cuda_pipeline/gpu_tlob.rs +++ b/crates/ml/src/cuda_pipeline/gpu_tlob.rs @@ -647,15 +647,18 @@ impl GpuTlob { /// Adam step: grad norm clip + Adam update on TLOB params. /// - /// SP3 Mech 9 (close-out v2): `weight_clamp_max_abs` is the ISV-driven - /// |p| bound for `attn_adam_kernel` (TLOB shares this kernel via the - /// attention cubin). Caller computes `100 × ISV[Q_ABS_REF].max(1.0)`. - /// 0=disabled (debug only). + /// SP4 Layer B (Mech 9): `weight_clamp_max_abs` is the ISV-driven |p| + /// bound for `attn_adam_kernel` (TLOB shares this kernel via the + /// attention cubin). Caller reads `ISV[WEIGHT_BOUND[Attn=6]]` (TLOB + /// co-lives with attention in the SP4 group taxonomy). 0=disabled + /// (debug only). `weight_decay_value` is sourced from + /// `ISV[WD_RATE[Attn=6]]` with `.max(0.0)`. pub(crate) fn adam_step( &mut self, lr: f32, max_grad_norm: f32, weight_clamp_max_abs: f32, + weight_decay_value: f32, ) -> Result<(), MLError> { let tp = TLOB_TOTAL_PARAMS; let tp_i32 = tp as i32; @@ -695,7 +698,11 @@ impl GpuTlob { let beta1 = 0.9_f32; let beta2 = 0.999_f32; let eps = 1e-8_f32; - let wd = 1e-5_f32; + // SP4 Layer B (Mech 9): AdamW weight_decay sourced from + // ISV[WD_RATE[Attn=6]] by the caller and threaded through + // `adam_step`. `.max(0.0)` upstream enforces the AdamW + // non-negativity contract. + let wd: f32 = weight_decay_value; let t_ptr = self.t_dev_ptr; // SP4 Task A14: Pearl C engagement-counter args. TLOB shares // the `attn_adam_kernel` cubin with `GpuAttention`. For Layer A diff --git a/crates/ml/src/cuda_pipeline/iql_value_kernel.cu b/crates/ml/src/cuda_pipeline/iql_value_kernel.cu index e3e7789b5..36977fc8e 100644 --- a/crates/ml/src/cuda_pipeline/iql_value_kernel.cu +++ b/crates/ml/src/cuda_pipeline/iql_value_kernel.cu @@ -29,6 +29,13 @@ #define VALUE_HIDDEN_DIM 128 #endif +/* SP4 Layer B: per-branch atom-position bound from ISV[ATOM_POS_BOUND_BASE + branch]. + * Mirrors `crates/ml/src/cuda_pipeline/sp4_isv_slots.rs::ATOM_POS_BOUND_BASE = 132`. + */ +#ifndef ATOM_POS_BOUND_BASE +#define ATOM_POS_BOUND_BASE 132 +#endif + /* Parameter sizes that don't depend on state_dim */ #define V_B1_SIZE (VALUE_HIDDEN_DIM) #define V_W2_SIZE (VALUE_HIDDEN_DIM * VALUE_HIDDEN_DIM) @@ -693,17 +700,14 @@ void iql_compute_per_sample_support( float v_max = r * iql_vmax + (1.0f - r) * ( 1.0f); float delta_z = r * iql_dz + (1.0f - r) * default_dz; - /* SP3 Mech 2: bound atom positions to ±10 × ISV[Q_ABS_REF=16].max(1.0). - * Per-sample support [v_min, v_max] spans the C51 atom grid for this - * (sample, branch); capping at 10× allows full distribution coverage - * without extreme-tail growth. Same ISV bound as Mech 1 (target_q clip) - * — atoms and target_q share the magnitude scale. ε on the multiplier - * (`isv.max(1.0)`) per the SP1 ε-floor pearl; bound ≥ 10 even with - * cold-start ISV. delta_z is recomputed from the clamped span so the + /* SP4 Layer B (Mech 2): bound per-sample support [v_min, v_max] to + * ±ISV[ATOM_POS_BOUND[d]]. Per-branch p99(|atom_positions|) populated + * by SP4 producer `launch_sp4_atom_pos_p99` (Layer A Task A6). + * ε floor on the bound itself (`max(.., 1.0)`) per the SP1 cold-start + * pearl. delta_z is recomputed from the clamped span so the * (v_min, v_max, delta_z) triple stays self-consistent. */ if (isv_signals != nullptr) { - const float q_abs_ref_eff = fmaxf(isv_signals[16], 1.0f); - const float max_atom_abs = 10.0f * q_abs_ref_eff; + const float max_atom_abs = fmaxf(isv_signals[ATOM_POS_BOUND_BASE + d], 1.0f); v_min = fminf(fmaxf(v_min, -max_atom_abs), max_atom_abs); v_max = fminf(fmaxf(v_max, -max_atom_abs), max_atom_abs); delta_z = (v_max - v_min) / (float)(num_atoms - 1); diff --git a/crates/ml/src/trainers/dqn/config.rs b/crates/ml/src/trainers/dqn/config.rs index 9ffc62872..fd1417d54 100644 --- a/crates/ml/src/trainers/dqn/config.rs +++ b/crates/ml/src/trainers/dqn/config.rs @@ -869,12 +869,13 @@ pub struct DQNHyperparameters { /// Steps for sigma annealing (default: 10000) pub noisy_sigma_anneal_steps: usize, - // WAVE 30: L2 Weight Decay for Overfitting Prevention - /// L2 weight decay for AdamW optimizer (default: 1e-4, range: [1e-5, 1e-3]). - /// Prevents overfitting by penalizing large weights via L2 regularization. - /// Recommended: 1e-4 for standard training, 1e-3 for aggressive regularization. - /// f64: tiny values (1e-5..1e-3) benefit from f64 precision during decay math. - pub weight_decay: f64, + // SP4 Layer B (2026-05-01): `weight_decay` field removed from + // hyperparams. Per-group AdamW weight_decay rates are now sourced + // from the SP4 ISV bus at `ISV[WD_RATE_BASE..+SP4_PARAM_GROUP_COUNT)` + // (slot indices 160..168) — see + // `crates/ml/src/cuda_pipeline/sp4_isv_slots.rs::wd_rate(group)`. + // Per-step rate adapts via the `param_group_oracle` Pearl B kernel + // (Layer A Task A7) and Pearls A+D smoothing (Task A12). /// Adam epsilon — BF16 requires ≥1e-4 (1e-8 rounds to 0 in f32 → div-by-zero). /// f64: tiny values (1e-8..1e-4) preserve denorm behaviour. @@ -1209,11 +1210,13 @@ impl DQNHyperparameters { pub fn apply_family_scaling(&mut self) { // ── Core families ── - // Learning family: scales learning_rate, tau, weight_decay + // Learning family: scales learning_rate, tau. + // (SP4 Layer B: `weight_decay` family-scaling removed — per-group + // weight_decay rates are now ISV-driven; PSO learning-intensity no + // longer perturbs the static config field that no longer exists.) let li = self.learning_intensity; self.learning_rate *= li; self.tau *= li as f32; - self.weight_decay *= li; // Exploration family: scales entropy_coefficient, noisy_sigma_init let exi = self.exploration_intensity; @@ -1519,8 +1522,8 @@ impl DQNHyperparameters { noisy_sigma_final: 0.4, // Default: 40% final noise noisy_sigma_anneal_steps: 10000, // Default: 10K steps for annealing - // WAVE 30: L2 Weight Decay - weight_decay: 1e-4, // Default: 0.0001 (standard regularization strength) + // SP4 Layer B: `weight_decay` field removed; per-group rates + // are sourced from ISV[WD_RATE[group]] at runtime. adam_epsilon: 1e-8, // f32 Adam (master weights are f32) // Conservative Q-Learning (CQL) @@ -1749,8 +1752,9 @@ pub(crate) fn dqn_default_config() -> DQNConfig { cvar_alpha: 0.05, minimum_profit_factor: 1.5, - weight_decay: 1e-4, - adam_epsilon: 1e-8, // f32 Adam (master weights are f32) + // SP4 Layer B: `weight_decay` field removed; per-group rates are + // sourced from ISV[WD_RATE[group]] at runtime. + adam_epsilon: 1e-8, // f32 Adam (master weights are f32) ..Default::default() } } diff --git a/crates/ml/src/trainers/dqn/fused_training.rs b/crates/ml/src/trainers/dqn/fused_training.rs index f124e708d..a0f1fff37 100644 --- a/crates/ml/src/trainers/dqn/fused_training.rs +++ b/crates/ml/src/trainers/dqn/fused_training.rs @@ -37,8 +37,7 @@ use crate::cuda_pipeline::mapped_pinned::{MappedF32Buffer, MappedI32Buffer}; use crate::cuda_pipeline::gpu_attention::{GpuAttention, GpuAttentionConfig}; use crate::cuda_pipeline::gpu_tlob::GpuTlob; use crate::cuda_pipeline::gpu_dqn_trainer::{ - GpuDqnTrainConfig, GpuDqnTrainer, Q_ABS_REF_INDEX, TAU_EFF_INDEX, - TLOB_REGIME_FOCUS_EMA_INDEX, + GpuDqnTrainConfig, GpuDqnTrainer, TAU_EFF_INDEX, }; use crate::cuda_pipeline::gpu_her::{GpuHer, GpuHerConfig, parse_her_strategy}; use crate::cuda_pipeline::gpu_iql_trainer::{GpuIqlConfig, GpuIqlTrainer}; @@ -458,7 +457,8 @@ impl FusedTrainingCtx { beta1: 0.9, beta2: 0.999, epsilon: hyperparams.adam_epsilon as f32, - weight_decay: hyperparams.weight_decay as f32, + // SP4 Layer B: `weight_decay` field removed from GpuDqnTrainConfig + // and DQNHyperparameters; per-group rates flow from ISV at runtime. max_grad_norm: resolved_grad_norm as f32, spectral_norm_sigma_max: hyperparams.spectral_norm_sigma_max, spectral_decoupling_lambda: hyperparams.spectral_decoupling_lambda, @@ -1615,15 +1615,26 @@ impl FusedTrainingCtx { self.gpu_tlob .backward(d_concat_ref, self.batch_size, concat_dim, tlob_concat_off) .map_err(|e| anyhow::anyhow!("TLOB backward: {e}"))?; - // SP3 Mech 9 (close-out v2): TLOB shares `attn_adam_kernel` - // (cubin-shared with GpuAttention). Same ISV-derived bound. - let q_abs_ref_eff_tlob = self.trainer.read_isv_signal_at(Q_ABS_REF_INDEX).max(1.0_f32); - let tlob_weight_clamp_max_abs: f32 = 100.0_f32 * q_abs_ref_eff_tlob; + // SP4 Layer B (Mech 9): TLOB shares `attn_adam_kernel` + // (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); self.gpu_tlob .adam_step( self.trainer.config().lr, self.trainer.config().max_grad_norm, tlob_weight_clamp_max_abs, + tlob_weight_decay_value, ) .map_err(|e| anyhow::anyhow!("TLOB Adam: {e}"))?; } @@ -1687,18 +1698,17 @@ impl FusedTrainingCtx { // ── Outside graph: kernel-free host ops only ───────────────── self.pending_vaccine_batch = None; - // SP4 Task A15: Pearl C post-Adam engagement-rate check (Layer A). + // SP4 Task A15: Pearl C post-Adam engagement-rate check. // Reads per-block engagement counts from the mapped-pinned buffer // (zero-copy host access), reduces → engagement_rate, applies - // Pearls A+D to the rate-deficit time series. Layer A logs only; - // Layer B will mutate ISV[WEIGHT_BOUND[group]] on sustained - // over-engagement. + // Pearls A+D to the rate-deficit time series. // - // The 5 main Adam launches captured in this step's graph-replay - // each contribute to a distinct param-group offset: - // group 0 (DqnTrunk) — covers DQN Trunk/Value/Branches - // (Layer A simplification: counted under - // group 0 only — Layer B will split). + // SP4 Layer B (2026-05-01): the main DQN Adam now splits into 3 + // sub-launches (Trunk/Value/Branches) — each writes per-block + // counts at its own group offset, so Pearl C is now per-group: + // group 0 (DqnTrunk) — DQN trunk Adam sub-launch. + // group 1 (DqnValue) — DQN value-head Adam sub-launch. + // group 2 (DqnBranches) — DQN branch-head Adam sub-launch. // group 3 (Iqn) — IQN dual-head trainer. // group 4 (IqlHigh) — IQL high-tau trainer. // group 5 (IqlLow) — IQL low-tau trainer. @@ -1706,13 +1716,25 @@ impl FusedTrainingCtx { // but disables Pearl C). // group 7 (Curiosity) — handled in `train_curiosity_gpu`, // not here (separate caller path). - let total_params_dqn = self.trainer.total_params(); + let (trunk_n, value_n, branches_n) = self.trainer.dqn_group_param_counts(); if let Err(e) = self.trainer.pearl_c_post_adam_engagement_check( crate::cuda_pipeline::ParamGroup::DqnTrunk, - total_params_dqn, + trunk_n, ) { tracing::warn!("SP4 Pearl C check (DQN trunk) failed (non-fatal): {e}"); } + if let Err(e) = self.trainer.pearl_c_post_adam_engagement_check( + crate::cuda_pipeline::ParamGroup::DqnValue, + value_n, + ) { + tracing::warn!("SP4 Pearl C check (DQN value) failed (non-fatal): {e}"); + } + if let Err(e) = self.trainer.pearl_c_post_adam_engagement_check( + crate::cuda_pipeline::ParamGroup::DqnBranches, + branches_n, + ) { + tracing::warn!("SP4 Pearl C check (DQN branches) failed (non-fatal): {e}"); + } if let Some(ref iqn) = self.gpu_iqn { let n = iqn.total_params(); if let Err(e) = self.trainer.pearl_c_post_adam_engagement_check( @@ -1900,18 +1922,33 @@ impl FusedTrainingCtx { // 3. Train V(s) — both tau=high and tau=low. // - // SP3 Mech 9 (close-out v2): ISV-driven post-Adam |p| clamp for - // `iql_adam_kernel`. Same `100 × Q_ABS_REF.max(1.0)` bound as - // the main DQN Adam — IQL's separate value-net params drift via - // its own Adam state and were never reached by the single-kernel - // Mech 9 in commit 8956c2fb7. ε on multiplier per SP1 pearl. - let q_abs_ref_eff = self.trainer.read_isv_signal_at(Q_ABS_REF_INDEX).max(1.0_f32); - let weight_clamp_max_abs: f32 = 100.0_f32 * q_abs_ref_eff; + // 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); let states_f32 = self.trainer.states_buf(); - self.gpu_iql.train_value_step(states_f32, weight_clamp_max_abs) + self.gpu_iql.train_value_step(states_f32, iql_high_clamp, iql_high_wd) .map_err(|e| anyhow::anyhow!("IQL high train: {e}"))?; let states_f32 = self.trainer.states_buf(); - self.gpu_iql_low.train_value_step(states_f32, weight_clamp_max_abs) + self.gpu_iql_low.train_value_step(states_f32, iql_low_clamp, iql_low_wd) .map_err(|e| anyhow::anyhow!("IQL low train: {e}"))?; // 4. Compute advantage weights (from high-tau V(s)). @@ -2003,19 +2040,26 @@ impl FusedTrainingCtx { let online_h_s2 = self.trainer.save_h_s2(); let next_states_buf = self.trainer.next_states_buf(); - // SP3 Mech 9 (close-out v2): ISV-driven post-Adam |p| clamp - // for `iqn_adam_kernel`. Same `100 × Q_ABS_REF.max(1.0)` - // bound as the main DQN Mech 9 launch (`launch_adam_update`) - // — IQN's separate online_params drives the slot-26 GEMM and - // was never reached by the single-kernel Mech 9 in commit - // 8956c2fb7. ε on multiplier per SP1 pearl. - let q_abs_ref_eff = self.trainer.read_isv_signal_at(Q_ABS_REF_INDEX).max(1.0_f32); - let weight_clamp_max_abs: f32 = 100.0_f32 * q_abs_ref_eff; + // 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); match iqn.execute_training_pipeline( online_h_s2, next_states_buf, &self.target_dueling, Some(&iqn_stream_ref), ws_override, weight_clamp_max_abs, + weight_decay_value, ) { Ok(_) => { iqn_ok = true; } Err(e) => { @@ -2084,11 +2128,21 @@ impl FusedTrainingCtx { // Adam on attn_stream. let lr = self.trainer.config().lr; let mgn = self.trainer.config().max_grad_norm; - // SP3 Mech 9 (close-out v2): ISV-driven post-Adam |p| clamp - // for `attn_adam_kernel`. Same pattern as IQN/IQL/main DQN. - let q_abs_ref_eff_attn = self.trainer.read_isv_signal_at(Q_ABS_REF_INDEX).max(1.0_f32); - let attn_weight_clamp_max_abs: f32 = 100.0_f32 * q_abs_ref_eff_attn; - attn.adam_step(lr, mgn, Some(&attn_stream_ref), ws_override, attn_weight_clamp_max_abs) + // 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); + 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}"))?; // Record completion on attn_stream. @@ -2121,11 +2175,18 @@ impl FusedTrainingCtx { .map_err(|e| anyhow::anyhow!("Attention backward: {e}"))?; let lr = self.trainer.config().lr; let mgn = self.trainer.config().max_grad_norm; - // SP3 Mech 9 (close-out v2): same ISV-derived bound as - // parallel arm. Sequential-fallback path. - let q_abs_ref_eff_attn = self.trainer.read_isv_signal_at(Q_ABS_REF_INDEX).max(1.0_f32); - let attn_weight_clamp_max_abs: f32 = 100.0_f32 * q_abs_ref_eff_attn; - attn.adam_step(lr, mgn, None, None, attn_weight_clamp_max_abs) + // 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); + attn.adam_step(lr, mgn, None, None, + attn_weight_clamp_max_abs, attn_weight_decay_value) .map_err(|e| anyhow::anyhow!("Attention Adam: {e}"))?; } @@ -2134,15 +2195,22 @@ impl FusedTrainingCtx { let online_h_s2 = self.trainer.save_h_s2(); let next_states_buf = self.trainer.next_states_buf(); - // SP3 Mech 9 (close-out v2): same ISV-derived bound as the + // SP4 Layer B (Mech 9): same ISV-driven bound as the // parallel arm above. Mirror the dqn_adam pattern. - let q_abs_ref_eff = self.trainer.read_isv_signal_at(Q_ABS_REF_INDEX).max(1.0_f32); - let weight_clamp_max_abs: f32 = 100.0_f32 * q_abs_ref_eff; + 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); match iqn.execute_training_pipeline( online_h_s2, next_states_buf, &self.target_dueling, None, None, weight_clamp_max_abs, + weight_decay_value, ) { Ok(_) => { iqn_ok = true; } Err(e) => { diff --git a/crates/ml/src/trainers/dqn/smoke_tests/generalization.rs b/crates/ml/src/trainers/dqn/smoke_tests/generalization.rs index 731246b9c..696a431c6 100644 --- a/crates/ml/src/trainers/dqn/smoke_tests/generalization.rs +++ b/crates/ml/src/trainers/dqn/smoke_tests/generalization.rs @@ -37,7 +37,11 @@ fn test_generalization_components_smoke() -> anyhow::Result<()> { params.epochs = 3; params.batch_size = 4096; // Larger batch to catch OOB that only trigger at scale params.buffer_size = 8192; - params.weight_decay = 1e-4; // G1: AdamW active + // SP4 Layer B: `weight_decay` removed from hyperparams. Per-group AdamW + // rates flow from ISV[WD_RATE[group]] at runtime; smoke tests inherit + // the bootstrap value (Pearl A sentinel = 0 → zero-decay until the + // first real observation, identical to the historical "AdamW active" + // contract once the EMA spins up). let mut trainer = smoke_trainer_with(params)?; let rt = tokio::runtime::Builder::new_current_thread() diff --git a/crates/ml/src/trainers/dqn/trainer/constructor.rs b/crates/ml/src/trainers/dqn/trainer/constructor.rs index 5dbf6a286..d972ced64 100644 --- a/crates/ml/src/trainers/dqn/trainer/constructor.rs +++ b/crates/ml/src/trainers/dqn/trainer/constructor.rs @@ -298,7 +298,10 @@ impl DQNTrainer { #[allow(clippy::cast_possible_truncation)] minimum_profit_factor: hyperparams.minimum_profit_factor, - weight_decay: hyperparams.weight_decay, + // SP4 Layer B: weight_decay no longer flows from hyperparams; + // ml-dqn `DQNConfig.weight_decay` retains its own default + // (currently unread by any consumer post-migration; cleanup + // deferred to a separate ml-dqn refactor pass). dropout_rate: f64::from(hyperparams.dropout_initial), entropy_coefficient: hyperparams.entropy_coefficient, noisy_epsilon_floor: hyperparams.noisy_epsilon_floor.unwrap_or(0.0), // C2: NoisyNet handles exploration diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 7ca234abc..dd2bfb007 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -1771,14 +1771,16 @@ impl DQNTrainer { } if count > 0 { - // SP3 Mech 9 (close-out v2): ISV-driven post-Adam |p| clamp for - // `curiosity_adam_step`. Same `100 × Q_ABS_REF.max(1.0)` bound as - // the main DQN / IQN / IQL / attention Adam launches. Curiosity + // SP4 Layer B (Mech 9): ISV-driven post-Adam |p| clamp for + // `curiosity_adam_step`. Curiosity params live in group 7 + // (`ParamGroup::Curiosity`); read `ISV[WEIGHT_BOUND[Curiosity]]` + // with EPS_CLAMP_FLOOR = 1.0 ε for cold-start. Curiosity // collector doesn't own ISV — read from fctx here and thread the - // bound through. ε on multiplier per SP1 pearl. + // bound through to the 4 sub-launches (W1/b1/W2/b2). let weight_clamp_max_abs: f32 = if let Some(ref fused) = self.fused_ctx { - let q_abs_ref_eff = fused.read_isv_signal_at(Q_ABS_REF_INDEX).max(1.0_f32); - 100.0_f32 * q_abs_ref_eff + fused.read_isv_signal_at(crate::cuda_pipeline::weight_bound( + crate::cuda_pipeline::ParamGroup::Curiosity.idx())) + .max(crate::cuda_pipeline::EPS_CLAMP_FLOOR) } else { // No fctx → no ISV → disable Mech 9. Reachable only on the // CPU-only build path (CUDA build asserts fctx Some at diff --git a/crates/ml/src/training_profile.rs b/crates/ml/src/training_profile.rs index cbf0aa1e8..df3ce5d75 100644 --- a/crates/ml/src/training_profile.rs +++ b/crates/ml/src/training_profile.rs @@ -59,7 +59,9 @@ pub struct TrainingSection { pub tau: Option, pub warmup_steps: Option, pub gradient_clip_norm: Option, - pub weight_decay: Option, + // SP4 Layer B: `weight_decay` removed from trial overrides. Per-group + // AdamW weight_decay rates are now sourced from ISV[WD_RATE[group]]; + // PSO no longer searches over a static config field. pub hidden_dim_base: Option, /// Reward scale factor for v_range computation. /// v_range = (reward_scale / (1 - gamma) * 1.2).clamp(20.0, 300.0) @@ -389,7 +391,8 @@ pub struct SearchSpaceSection { pub dueling_hidden_dim: Option<[f64; 2]>, pub n_steps: Option<[f64; 2]>, pub num_atoms: Option<[f64; 2]>, - pub weight_decay: Option<[f64; 2]>, + // SP4 Layer B: `weight_decay` removed from PSO search bounds. Per-group + // weight_decay rates are ISV-driven at runtime — no static knob. pub kelly_fractional: Option<[f64; 2]>, pub kelly_max_fraction: Option<[f64; 2]>, pub volatility_window: Option<[f64; 2]>, @@ -558,7 +561,7 @@ impl HyperoptProfile { "dueling_hidden_dim" => ss.dueling_hidden_dim, "n_steps" => ss.n_steps, "num_atoms" => ss.num_atoms, - "weight_decay" => ss.weight_decay, + // SP4 Layer B: `weight_decay` removed from PSO search-space. "kelly_fractional" => ss.kelly_fractional, "kelly_max_fraction" => ss.kelly_max_fraction, "volatility_window" => ss.volatility_window, @@ -746,9 +749,8 @@ impl DqnTrainingProfile { if let Some(v) = t.gradient_clip_norm { hp.gradient_clip_norm = Some(v); } - if let Some(v) = t.weight_decay { - hp.weight_decay = v; - } + // SP4 Layer B: `weight_decay` removed from trial overrides; per-group + // rates flow through ISV[WD_RATE[group]] at runtime. if let Some(v) = t.hidden_dim_base { hp.hidden_dim_base = Some(v); } diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index a10f89ea7..87f6f9e42 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2676,3 +2676,107 @@ Verification: `git grep -nE "_ema_alpha_unused|launch_reward_component_ema\b|lau asserting kernel output matches host-reference `pearls_ad_update` within fp32 ULP across 5 representative trajectories + 1 multi-slot batch). 13 existing tests + 1 new test pass on RTX 3050 Ti. + +### Layer B atomic consumer migration (2026-05-01) + +**Layer B atomic consumer migration — all SP3 hardcoded multipliers +replaced by ISV reads, DQN main Adam per-group split, Pearl C engagement +deferral resolved.** Single coordinated commit per +`feedback_no_partial_refactor`. The SP3-era multiplier-driven bounds +(`10 × Q_ABS_REF`, `100 × Q_ABS_REF`, `1e3 × Q_ABS_REF`, +`1e6 × Q_ABS_REF²`, `100 × H_S2_RMS_EMA`) and `.max(1.0)` ε floors are +replaced by per-slot ISV reads (`ISV[TARGET_Q_BOUND]`, +`ISV[ATOM_POS_BOUND[branch]]`, `ISV[WEIGHT_BOUND[group]]`, +`ISV[ADAM_M_BOUND[group]]`, `ISV[ADAM_V_BOUND[group]]`, +`ISV[GRAD_CLIP_BOUND]`, `ISV[H_S2_BOUND]`, +`ISV[L1_LAMBDA_TRUNK_INDEX]`, `ISV[WD_RATE[group]]`) with +consumer-side `EPS_CLAMP_FLOOR = 1.0` ε for cold-start. + +DQN main Adam launch split into 3 sub-launches (DqnTrunk / DqnValue / +DqnBranches) per `feedback_no_quickfixes` — overrides the plan's +"max/min-of-3 single-launch shortcut" recommendation. Each sub-launch +reads its own `WEIGHT_BOUND[group]`, `WD_RATE[group]`, and (trunk only) +`L1_LAMBDA_TRUNK_INDEX`. Per-group split also resolves the Pearl C +engagement-counter accounting deferral from A14/A15 (groups 1+2 were +silently rolled into group 0's offset; now each writes per-block +counts at its own `group_idx × MAX_BLOCKS_PER_ADAM` offset, and +`pearl_c_post_adam_engagement_check` is invoked per group from +`fused_training.rs`). + +Mech 5 fused diagnostic kernel (`dqn_nan_check_fused_f32_kernel`) +takes `isv_signals` device pointer instead of two host-side `*_eff` +floats — per-slot ISV reads inside the kernel eliminate the per-step +HtoD copy and remove `q_abs_ref_eff` / `h_s2_rms_ema_eff` kernel args +entirely. Mapped-pinned ISV pointer is graph-capture-safe. + +`weight_decay` field removed from: +- `DQNHyperparameters` (`crates/ml/src/trainers/dqn/config.rs`) +- `GpuDqnTrainConfig` (`crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`) +- `GpuIqnConfig` (`crates/ml/src/cuda_pipeline/gpu_iqn_head.rs`) +- `GpuIqlConfig` (`crates/ml/src/cuda_pipeline/gpu_iql_trainer.rs`) +- `TrialOverrides` / search-space (`crates/ml/src/training_profile.rs`) +- `DQNHyperparameters::apply_family_scaling` + (`weight_decay *= li` line — no static field to scale anymore) + +DT (`decision_transformer.rs`) and aux trainers (ofi_embed, denoise, +sel/recursive_conf in `gpu_dqn_trainer.rs`) keep `weight_clamp_max_abs += 0.0` disable — they're outside the SP4 8-group taxonomy and have no +ISV producer. Mirrors DT's existing pattern. + +Files-touched manifest: +- `crates/ml/src/cuda_pipeline/atoms_update_kernel.cu` (Mech 2 per-branch + ATOM_POS_BOUND read) +- `crates/ml/src/cuda_pipeline/iql_value_kernel.cu` (Mech 2 per-branch + ATOM_POS_BOUND read) +- `crates/ml/src/cuda_pipeline/experience_kernels.cu` (Mech 2 per-branch + ATOM_POS_BOUND read for the legacy single-branch entry point) +- `crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu` (Mech 5 fused + diagnostic per-slot ISV reads + drop `q_abs_ref_eff`/`h_s2_rms_ema_eff` + args; doc updates) +- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (Mech 1 + 6 + 9 + 10 + consumer migrations; `launch_adam_update` 3-way split; aux trainers + disable clamp; `weight_decay` field removed; nan_check_fused launcher + drops `q_abs_ref_eff`/`h_s2_rms_ema_eff`; `dqn_group_param_counts` + accessor for fused_training Pearl C invocation) +- `crates/ml/src/cuda_pipeline/gpu_iqn_head.rs` (`weight_decay` field + removed; `execute_training_pipeline` + `train_iqn_step_gpu` accept + `weight_decay_value: f32`) +- `crates/ml/src/cuda_pipeline/gpu_iql_trainer.rs` (`weight_decay` field + removed; `train_value_step` accepts `weight_decay_value: f32`) +- `crates/ml/src/cuda_pipeline/gpu_attention.rs` (`adam_step` accepts + `weight_decay_value: f32`) +- `crates/ml/src/cuda_pipeline/gpu_tlob.rs` (`adam_step` accepts + `weight_decay_value: f32`; co-lives in Attn group taxonomy) +- `crates/ml/src/trainers/dqn/fused_training.rs` (per-launch ISV reads + for IQN/IQL/Attn/TLOB clamp + weight_decay; `weight_decay` removed + from `GpuDqnTrainConfig` builder; per-group `pearl_c_post_adam_engagement_check` + invocation for DqnTrunk/DqnValue/DqnBranches) +- `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (Curiosity + weight_clamp reads `ISV[WEIGHT_BOUND[Curiosity]]`) +- `crates/ml/src/trainers/dqn/trainer/constructor.rs` (`weight_decay: + hyperparams.weight_decay` line removed) +- `crates/ml/src/trainers/dqn/config.rs` (field + default + family-scaling + removed) +- `crates/ml/src/trainers/dqn/smoke_tests/generalization.rs` + (`params.weight_decay = 1e-4` line removed) +- `crates/ml/src/training_profile.rs` (TrialOverrides + search-space + + apply_to weight_decay branch removed) +- `crates/ml/examples/train_baseline_rl.rs` (`weight_decay:` initialiser + line removed) + +Verification (local, RTX 3050 Ti, post-commit): +- `SQLX_OFFLINE=true cargo check -p ml --offline`: clean, 11 warnings + (all pre-existing). +- `git grep -nE "10\.0_f32 \* q_abs_ref|10\.0f \* q_abs_ref|100\.0_f32 + \* q_abs_ref|100\.0f \* q_abs_ref|1e6_f32 \* q_abs_ref|1e3_f32 \* + q_abs_ref|100\.0_f32 \* h_s2|100\.0f \* h_s2_rms" crates/ml/src/`: + ZERO matches. +- `git grep -nE "weight_decay:\s*f64|l1_lambda:" crates/ml/src/trainers/dqn/`: + ZERO matches. +- `git grep -n "self.hyperparams.weight_decay|self.config.weight_decay| + self.hyperparams.l1_lambda" crates/ml/src/`: only TFT remains + (separate trainer outside SP4 scope). +- `git grep -n "q_abs_ref_eff|h_s2_rms_ema_eff" crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu`: + ZERO matches. +- `cargo test -p ml --lib --offline`: 928 passed, 14 failed (all 14 + pre-existing on HEAD `1389d1c81`; no new failures introduced).