9ba08ff609ace717cbb24ec6c1c4392e8564b3d3
2194 Commits
| Author | SHA1 | Message | Date | |
|---|---|---|---|---|
jgrusewski
|
9ba08ff609 |
chore(ml): gate ZN.FUT data-loader tests + migrate test HtoD/DtoH to mapped-pinned
ZN.FUT tests in crates/ml/src/data_loader.rs were failing because no
valid ZN.FUT DBN data is available locally; gated with #[ignore].
ml-asset-selection's universe definition and backtesting's
zn_futures() slippage profile remain untouched — those are production
references to ZN as a candidate symbol, distinct from data availability.
Migrated 4 deprecated cudarc memcpy_stod/memcpy_dtov sites in the
test function test_eval_action_select_eval_argmax_picks_best in
crates/ml/src/cuda_pipeline/mod.rs to mapped-pinned per
feedback_no_htod_htoh_only_mapped_pinned:
- 3x memcpy_stod (f32 input uploads) → MappedF32Buffer::new +
write_from_slice + dev_ptr as raw u64 kernel arg; kernel reads
directly from mapped-pinned pages, no DtoD copy needed
- 1x memcpy_dtov (i32 output readback) → MappedI32Buffer::new +
dev_ptr as kernel arg + read_all() after stream sync
The cudarc deprecation suggested clone_htod/clone_dtoh as replacements
but those still perform HtoD/DtoH copies — violating the strict rule.
Mapped-pinned with direct dev_ptr kernel args is the correct pattern
(matches distributional_q_tests.rs).
Note: DqnGpuData/PpoGpuData upload paths also in mod.rs still use
clone_to_device_f32_via_pinned; migrating those requires changing
CudaSlice<f32> struct fields to MappedF32Buffer which is blocked until
gpu_dqn_trainer.rs consumers are also updated (separate scope).
Workspace cargo check warnings: 15 → 15 (test-only deprecated calls
not visible to cargo check; ZN gate adds 3 to ignored count).
cargo test -p ml --lib failures: 16 → 13 (3 ZN tests now ignored).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
d6eca73e52 |
feat(dqn): #212 — intent-side magnitude distribution HEALTH_DIAG metric
Adds intent_dist_q/h/f alongside eval_dist_q/h/f to separate policy-
learning quality from Kelly-enforcement reality. Per memory pearl
project_magnitude_eval_collapse_kelly_capped.md: EVAL_DIST Quarter
dominance is a downstream artefact of Kelly cap × warmup_floor ×
safety_multiplier math, NOT a policy-learning failure. The intent
metric exposes the policy's pre-Kelly-cap chosen mag bucket so
operators can distinguish "policy isn't learning Full" from "Kelly
cap is suppressing Full" — the latter being an operationally-correct
state during cold-start positions.
Pure observability: no Kelly-math tweaks, no new tuned constants, no
reward-bias mechanisms (all explicitly forbidden by the memory pearl).
The intent_mag bucket comes straight from the factored action's
mag_idx (0/1/2) which already maps 1:1 to Quarter/Half/Full buckets.
The pre-cap mag_idx is already captured by the action-select kernel
into intent_mag_buf and exposed via the trainer's pre-existing
last_eval_intent_magnitude_dist host field — this commit only wires
that signal into HEALTH_DIAG.
Wired through (one coordinated commit per feedback_no_partial_refactor):
- health_diag.rs: HealthDiagSnapshot gains intent_dist_q/h/f fields
adjacent to eval_dist_*; size assertion bumped 147 → 150 fields.
- health_diag_kernel.cu: 3 new WORD_INTENT_DIST_* slots at [77..80);
every downstream WORD_* shifted by +3 in lockstep; WORD_TOTAL +
static_assert bumped 147 → 150. (The slots are reserved identically
to WORD_EVAL_DIST_* — both currently inert; HEALTH_DIAG GPU port
Phases 2/3 will populate them in lockstep with their sibling slots.)
- training_loop.rs: new adjacent tracing::info!() emitting
"intent_dist [iq=... ih=... if=...]" from the existing host-side
last_eval_intent_magnitude_dist field, immediately after the big
HEALTH_DIAG line containing eval_dist [eq=... eh=... ef=...].
- docs/dqn-wire-up-audit.md: new top-of-file entry per Invariant 7.
Behavior: zero change. Only adds observability.
Verification:
- cargo check -p ml --offline: clean.
- cargo test -p ml --lib --offline -- health_diag: 3/3 pass.
- cargo test -p ml --test sp4_producer_unit_tests --release
--ignored: 16/16 GPU tests pass.
Closes #212.
|
||
|
jgrusewski
|
6d0ac7beb3 |
fix(sp4): GPU-port calibrate_homeostatic_targets per feedback_no_cpu_compute_strict
Per-step host-side EMA loop at gpu_dqn_trainer.rs:4671 over 6 mapped-
pinned homeostatic-target slots was a feedback_no_cpu_compute_strict
violation discovered during the sweep audit (commit
|
||
|
jgrusewski
|
a385c1d2be |
chore(sp4): delete compute_adaptive_tau orphan helper per feedback_wire_everything_up
`compute_adaptive_tau` (gpu_dqn_trainer.rs:7045) had zero production
callers — its q_div_ema field's doc-comment explicitly flagged it
"NOTE: only consumed by the orphan helper compute_adaptive_tau (zero..)".
Pre-SP3 attempt at adaptive-tau control, superseded by the SP3/SP4 ISV-
driven tau controller (`tau_kernel.cu` + ISV[TAU_INDEX]).
Removed:
- compute_adaptive_tau method (CPU EMA + ratio-scaled tau formula)
- q_divergence_readback method (only caller was compute_adaptive_tau)
- q_div_ema field + initialiser + fold-reset assignment
- q_divergence_pinned + q_divergence_dev_ptr struct fields, alloc, free
- q_divergence_dev_ptr memsets at C51 launch sites (2)
- q_divergence_dev_ptr arg in launch_c51_loss kernel call
- c51_loss_batched kernel parameter `q_divergence` (never written by
kernel body — comment "removed from hot path — zero atomicAdd"
confirmed buffer was inert; deleting parameter keeps kernel ABI clean
per feedback_no_partial_refactor)
- Stale doc-comments referencing q_divergence in c51 reduction kernel
- readback_pinned [12]=q_divergence layout doc (slot was never read)
cargo check clean. SP4 + state_reset_registry lib tests pass (11/11).
16/16 SP4 producer GPU tests pass on RTX 3050 Ti. No behavior change —
dead code only.
Refs: feedback_no_cpu_compute_strict sweep audit (commit
|
||
|
jgrusewski
|
1112abc2a4 |
fix(sp4): migrate winsorized adaptive grad-clip update to GPU per feedback_no_cpu_compute_strict
Layer C close-out C4 — the most architecturally substantive site of the feedback_no_cpu_compute_strict sweep. GpuDqnTrainer::update_adaptive_clip was running a 6-step host-side compute chain on GPU-produced inputs: winsor (1) + cold-start sentinel + EMA (3) + scalar reduction (4) + ISV upper-bound clamp (5) + mapped-pinned write (6). Per feedback_no_partial_refactor the chain must migrate coherently — splitting EMA-only into a kernel and leaving the surrounding scalar reductions on host would be a partial migration violating the rule. Migration: new update_adaptive_clip_kernel.cu (single-thread, single-block). Takes host-passed observed_grad_norm (already a mapped-pinned readback) + 6 fixed structural constants (legacy values preserved per feedback_no_quickfixes) + 4 mapped-pinned dev_ptrs + ISV[GRAD_CLIP_BOUND_INDEX]. Writes the full output chain (adaptive_clip_pinned, grad_norm_ema_pinned, outlier_diag_pinned). Storage migration: - grad_norm_ema migrated from host-resident f32 to mapped-pinned scalar (matches C1/C2/C3 pattern). - New outlier_diag_pinned mapped-pinned slot for the GRAD_CLIP_OUTLIER warn diagnostic. The kernel writes `delta = observed - clamped` and the host reads it post-launch to format the warn log without running scalar arithmetic on the host. All structural constants preserved: EMA_BETA=0.95, CLIP_MULTIPLIER=2.0, MIN_CLIP=1.0, GRAD_CLIP_OUTLIER_K=100, EPS_CLAMP_FLOOR (SP4). Same cold-start sentinel `prev_ema <= 0.0 ⇒ assign clamped directly`; same Mech 6 (SP3) + Layer B (SP4) bound design. Same outlier-warn log format reconstructed from the mapped-pinned diagnostic slot. Host-side early-return guard preserved (the pre-existing pattern from C1 redesigned). grad_norm_emas_step_count counter unchanged (scalar control-flow metadata, not compute, per the rule's explicit carve-out). Verification: SP4 lib tests + 16 SP4 GPU producer unit tests pass on RTX 3050 Ti. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
ca63158606 |
fix(sp4): migrate atom utilization EMA to GPU per feedback_no_cpu_compute_strict
Layer C close-out C3 — the host-side adaptive-α EMA over result.atom_utilization
(GPU-produced via q_readback_pinned mapped-pinned readback) at the bottom of
GpuDqnTrainer::reduce_current_q_stats violated feedback_no_cpu_compute_strict.
Migration: new update_utilization_ema_kernel.cu (single-thread, single-block,
mirrors C2's update_iqn_readiness_kernel shape). Takes the host-passed
atom_util scalar and updates two mapped-pinned slots in lockstep:
- utilization_ema_pinned (the EMA's new storage)
- homeostatic_obs_pinned[1] (the homeostatic mirror previously written by
update_homeostatic_observables host-side; kernel takes over)
Storage: utilization_ema migrated from host-resident f32 field to mapped-pinned
device-mapped scalar (matches C1/C2 pattern). Constructor init=1.0 preserves
the deleted host code's `if self.utilization_ema > 0.99 { assign obs }`
cold-start sentinel; same adaptive-α formula clamp(|err|/(|err|+0.1), 0.01, 0.30)
and EMA recurrence.
Consumer-chain coherence per feedback_no_partial_refactor:
- `update_homeostatic_observables` slot [1] write removed (kernel takes over).
- `utilization_ema()` accessor reads through pinned host_ptr.
- `adaptive_entropy` host derivation in `launch_c51_grad` reads via accessor.
- collector `set_utilization_ema` callsite uses accessor unchanged.
Discovered during the audit (deferred to separate tasks):
- `compute_adaptive_tau` q_div_ema is host-side EMA on a helper with ZERO
production callers — flagged per feedback_wire_everything_up.
- `calibrate_homeostatic_targets` runs a per-step host-side EMA loop over 6
mapped-pinned slots — NEW violation not in original 9-site list, flagged
for separate Layer C task.
- collector::set_utilization_ema was misclassified in original audit (it's a
setter, not EMA arithmetic).
Verification: SP4 lib tests + 16 SP4 GPU producer unit tests pass on RTX 3050 Ti.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
605a8f5268 |
fix(sp4): migrate IQN readiness gauge update to GPU per feedback_no_cpu_compute_strict
Layer C close-out C2 — the host-side EMA arithmetic block in GpuDqnTrainer::update_iqn_readiness violated feedback_no_cpu_compute_strict (any compute — EMA, reduction, mean — belongs on GPU regardless of frequency). The cold-start sentinel + adaptive-α EMA + improvement-gauge formula now run GPU-side via update_iqn_readiness_kernel (single-thread, single-block, mirrors update_grad_norm_emas_kernel shape). The kernel takes the host-passed iqn_loss scalar (from gpu_iqn.read_total_loss() mapped-pinned readback) and updates three coupled mapped-pinned scalars (iqn_loss_initial_pinned, iqn_loss_ema_pinned, iqn_readiness_pinned) in-place. __threadfence_system() guarantees PCIe-visibility to mapped-pinned host_ptr accessors and to other-stream kernel reads via dev_ptrs. Storage: iqn_loss_initial and iqn_loss_ema migrated from host-resident f32 fields to mapped-pinned device-mapped scalars (matches grad_norm_fast/slow_ema pattern from C1 redesigned). New accessors iqn_loss_ema_value() and iqn_loss_initial_value() read through the host_ptrs. iqn_readiness pinned slot remains the same — c51_loss_kernel CVaR α consumer dev_ptr unchanged. Cold-start sentinel (prev_initial < 1e-12 ⇒ assign loss directly + readiness=0) preserves the deleted host-side bootstrap branch exactly. Same adaptive-α formula clamp(|err|/(|err|+0.01), 0.01, 0.30) and improvement gauge clamped to [0,1]. state_reset_registry entries updated to reflect the new mapped-pinned storage and producer relocation; reset_iqn_readiness_state writes 0.0 through the pinned slots so the next kernel launch re-enters the bootstrap branch. Verification: SP4 lib tests + 16 SP4 GPU producer unit tests pass on RTX 3050 Ti. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
24accea774 |
fix(sp4): migrate fast/slow grad-norm EMA update to GPU per feedback_no_cpu_compute_strict
Layer C close-out C1 redesigned. Original plan claimed
grad_norm_slow_ema_pinned was orphan post-Mech-6 migration; verification
surfaced a SECOND live consumer (fold_warmup_factor_update, commit
|
||
|
jgrusewski
|
45da3eac69 |
refactor(sp4): #260 follow-ups — symbolic doc anchors, dedicated q_dir_grad table, p99 helper
3 IMPORTANT items from #260 code-quality review (commit |
||
|
jgrusewski
|
88ae74ca73 |
fix(sp4): #260 — eliminate SP1-era 1e6 × isv multipliers via ISV-driven producers
Two SP1-Phase-C cuBLAS backward sanitisers at gpu_dqn_trainer.rs:7615
(bw_d_h_s2 input) and :20679 (d_value_logits + d_adv_logits) used
hardcoded `1e6 × signal.max(1.0)` formulas — outside SP4 mechanisms
1/2/5/6/9/10 but flagged by feedback_isv_for_adaptive_bounds review as
the last remaining hardcoded multipliers in the cuBLAS-backward
sanitiser consumer path after Layer A/B closed out the rest of SP4.
Migrated to the proper SP4 pattern:
- 2 new SP4 ISV slots [171, 172) extending [131..171) → [131..173):
BW_D_H_S2_BOUND_INDEX=171 (single-buffer p99) and
Q_DIR_GRAD_BOUND_INDEX=172 (multi-sub-buffer p99 over
d_value_logits ∪ d_adv_logits, n_sub=2).
- 2 new producer kernels:
bw_d_h_s2_p99_kernel.cu (mirrors h_s2_p99 single-buffer pattern)
q_dir_grad_p99_kernel.cu (mirrors param_group_oracle multi-sub
convention via shared sp4_histogram_p99_multi<256> template;
reuses oracle_subbuf_table_buf + oracle_subbuf_counts_buf).
- Pearls A+D wired via existing apply_pearls_ad_kernel chained on
same stream (GPU-only, graph-capture-compatible).
- Consumers read ISV[slot].max(EPS_CLAMP_FLOOR) directly.
Buffer growth (single source of truth in gpu_dqn_trainer.rs):
SP4_PRODUCER_COUNT 69 → 71, SP4_WIENER_TOTAL_FLOATS 207 → 213.
ISV_TOTAL_DIM 171 → 173. LAYOUT_FINGERPRINT_SEED extended (existing
checkpoints will fail-fast at load — re-train required).
StateResetRegistry: 2 new FoldReset entries (sp4_bw_d_h_s2_bound,
sp4_q_dir_grad_bound) + 2 matching reset_named_state dispatch arms;
both halves of Pearl A's sentinel contract reset together per
feedback_no_partial_refactor.
Unit tests: 2 new GPU-gated tests in sp4_producer_unit_tests.rs
exercising (single-buffer Pearl-A→Pearl-D convergence, multi-sub-
buffer p99 vs analytical |N(0,1)|). All 16 GPU tests pass on local
RTX 3050 Ti (rel-err 0.00000 / 0.00526).
Behaviour change: bound is now Pearls-smoothed p99 of the actual
gradient distribution, tighter than the pre-existing 1e6 ceiling but
appropriate-scale for observed values. Smoke validation must verify
F0/F1/F2 still converge.
Refs: task #260, baseline commit
|
||
|
jgrusewski
|
1c150d1901 |
fix(sp4): Layer B fix-up — restore Adam coverage for tensors [33..163)
Layer B's 3-way DQN main Adam split (commit |
||
|
jgrusewski
|
58ffb3a48e |
feat(sp4): Layer B — atomic consumer migration to ISV-driven bounds
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) <noreply@anthropic.com> |
||
|
jgrusewski
|
1389d1c810 |
refactor(sp4): GPU-only Pearls A+D — eliminate all host-side compute paths
Layer A L40S smoke smoke-test-v9kjv revealed CUDA Graph capture failure
at aux_label_scale_ema mid-step host sync inside aux_heads_forward Step 2b
(CUDA_ERROR_STREAM_CAPTURE_UNSUPPORTED). Per feedback_no_cpu_compute_strict
and pearl_cold_path_no_exception_to_gpu_drives: CPU compute path is
strictly forbidden — any formula (EMA / reduction / Pearls A+D / adaptive α)
belongs on GPU regardless of frequency.
This commit migrates ALL 11 Pearls A+D applications + 1 inline application
to a single GPU apply_pearls_ad_kernel:
- 5 SP4 producers (target_q, atom_pos×4, param_group_oracle, grad_norm, h_s2)
- 6 A13 retrofits (h_s2_rms, aux_heads_loss, moe_expert_util,
vsn_mask, iqn_quantile, reward_component)
- 1 inline aux_label_scale block in aux_heads_forward Step 2b
Eliminates: stream.synchronize() + host_ptr read_volatile + host arithmetic
+ host_ptr write_volatile pattern from all 11 launchers + 1 inline.
apply_pearls_to_slot host helper deleted. pearls_ad_update kept as
test-only reference implementation (also retained for the post-Adam
pearl_c_post_adam_engagement_check host-side diagnostic which reduces
mapped-pinned i32 counters outside any captured graph). New GPU unit
test asserts kernel output matches reference within fp32 ULP for 5
representative cases plus a multi-slot batch sanity check.
The refactor is graph-capture-compatible by construction: the kernel
runs single-thread in the same stream as the producer kernel; no host
synchronisation needed between producer and applicator.
Build clean (cargo check --workspace), 6 host sp4_wiener_ema tests pass,
13 SP4 GPU tests + 1 new oracle test pass on RTX 3050 Ti. L40S smoke
re-validation deferred to A17 redo.
Refs: smoke-test-v9kjv graph-capture failure (terminated), commit
|
||
|
jgrusewski
|
4c231fa812 |
refactor(sp4): A13 code-quality pass — DRY helper, named constant, doc fix, param cleanup
Addresses 5 IMPORTANT items from A13 code-quality review: 1. apply_pearls_to_slot helper extracted into sp4_wiener_ema.rs. Collapses ~30 lines of read_volatile / pearls_ad_update / write_volatile per-slot block into a single unsafe fn. 12 launchers now consume the helper (5 SP4 producers A5-A9, 6 A13 retrofits A13.0-A13.5, plus the inline label-scale block in aux_heads_forward Step 2b — including 2 apply_pearls closures inside multi-slot launchers and the cross-boundary GpuExperienceCollector consumer). Pearl C rate_deficit site untouched (different mapped-pinned buffer + Rust ema array, doesn't share the helper's pointer-based contract). 2. REWARD_COMPONENT_COUNT named constant added next to REWARD_POPART_EMA_INDEX in gpu_dqn_trainer.rs. Replaces 3 hardcoded `6` literals across collector launcher (block_dim, Pearls A+D loop) and training_loop fold-reset range arithmetic. Mirrors MOE_NUM_EXPERTS / SL_NUM_FEATURE_GROUPS invariant-guard pattern with debug_assert_eq! in the collector launcher. Kernel literal `6` retained (allows nvcc full unroll); kernel comment now documents the host-side invariant. 3. SP4_PRODUCER_COUNT, SP4_WIENER_FLOATS_PER_SLOT, SP4_WIENER_TOTAL_FLOATS promoted from fn-local consts inside `pub fn new` to module-level `pub const`s in gpu_dqn_trainer.rs, re-exported via cuda_pipeline::mod. All 13 redeclarations in tests/sp4_producer_unit_tests.rs replaced with single `use ml::cuda_pipeline::SP4_PRODUCER_COUNT;` import. Future buffer growth requires single-file edit. 4. _ema_alpha_unused: f32 caller-compat shim removed from 6 retrofitted launchers (launch_h_s2_rms_ema, launch_aux_heads_loss_ema, launch_vsn_mask_ema, launch_moe_expert_util_ema, launch_iqn_quantile_ema, launch_reward_component_ema_inplace) and the FusedTrainingCtx proxy. All callers in training_loop.rs + fused_training.rs updated to drop the unused argument per feedback_no_legacy_aliases (no soft-deprecated wrappers; rename all call sites directly). Doc-comments updated from "α dropped per SP4 — argument preserved so callers compile unchanged" to "α derived adaptively from per-slot Pearls A+D Wiener state — see sp4_wiener_ema::pearls_ad_update". 5. Stale doc reference fixed at gpu_aux_heads.rs:391 — comment referenced non-existent launch_label_scale_ema_with_pearls function. Now correctly points at the inline Pearls A+D block in GpuDqnTrainer::aux_heads_forward Step 2b (which now consumes apply_pearls_to_slot). Pure refactor — no spec or behaviour change. Same kernel launches, same Pearls A+D semantics, same ISV slot writes. cargo check -p ml --offline clean (12 pre-existing warnings, no new); cargo test -p ml --lib --offline sp4_wiener_ema 7/7 passing (6 originals + apply_pearls_to_slot_pearl_a_bootstrap_path); cargo test -p ml --lib --offline state_reset_registry 3/3 passing. Refs: A13 review (commits aada419de..c5add566d). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
c5add566db |
fix(sp4): A13 spec gap — fold-reset retrofit ISV slots to Pearl A sentinel (0.0)
A13's per-producer commits (A13.0..A13.5) zeroed the Wiener triples at
fold boundary via the bulk `sp4_wiener_state` reset, but left the
companion ISV-slot dispatch arms writing pre-SP4 cold-start values. With
`prev_x_mean != 0 && state.x_lag == 0`, `pearls_ad_update` enters Pearl D's
formula at step 0 with `prev = stale_cold_start` — defeating Pearl A's
sentinel-detection contract documented at design-spec L262/284/286.
Updated 5 retrofit producer dispatch arms in `training_loop.rs` to write
0.0 at fold boundary (matching the Wiener-state half of the contract):
- isv_h_s2_rms_ema (was 1.0)
- isv_vsn_mask_g{0..5}_ema (was 1/SL_NUM_FEATURE_GROUPS = 1/6 per group)
- isv_aux_label_scale_ema (was 1.0; consumer 1e-6 floor guards
the divide-by-zero window between fold
boundary and first producer fire)
- isv_moe_expert_util_ema (was 1/MOE_EXPERT_UTIL_EMA_COUNT = 1/8 per
expert)
- isv_moe_gate_entropy_ema (was ln(MOE_EXPERT_UTIL_EMA_COUNT) = ln(8))
Already-correct retrofit arms verified writing 0.0 (no change):
- isv_iqn_q_p{05,25,75,95}_ema (A13.4) — already 0.0
- isv_aux_next_bar_mse_ema, isv_aux_regime_ce_ema (A13.1) — already 0.0
- isv_reward_component_emas (A13.5) — already 0.0
Registry descriptions in `state_reset_registry.rs` updated in the same
commit per `feedback_no_partial_refactor.md` (doc + dispatch arm are two
halves of the same contract — both must reset together).
Also fixed stale `141`-float and `2048`-int Wiener-buffer comments to
`207` / `2816` in:
- state_reset_registry.rs (block comment around the SP4 fold-reset
contract block)
- training_loop.rs (sp4_wiener_state bulk-reset comment 141 -> 207)
- gpu_dqn_trainer.rs (wiener_state_buf field doc 47 -> 69 ISV-bound
EMAs; scratch-layout comment 40..47/7 -> 40..69/29)
Verification:
- cargo check -p ml --offline: clean (11 pre-existing warnings)
- cargo test -p ml --lib --offline sp4_wiener_ema: 6/6 passing
- cargo test -p ml --lib --offline state_reset_registry: 3/3 passing
Spec: docs/superpowers/specs/2026-04-30-sp4-signal-driven-magnitude-control-design.md L262/284/286
Plan: docs/superpowers/plans/2026-04-30-sp4-signal-driven-magnitude-control.md L440
Refs: A13.0..A13.5 (aada419de..4f82b74a5)
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
4f82b74a51 |
feat(sp4): Task A13.5 — retrofit reward_component_ema with Pearls A+D + cross-boundary wiring + orphan deletion
Replaces the kernel's hardcoded `ema_alpha` with the shared `pearls_ad_update`
host-side helper, wires the host-side update across the trainer/collector
boundary (mirrors the A14/A15 Pearl C wiring path), and deletes the
trainer's orphan `launch_reward_component_ema` per
`feedback_wire_everything_up.md`.
Kernel + collector launcher:
- Kernel `reward_component_ema` signature: drops `(isv_out, ema_alpha,
isv_reward_base_slot)` for `(scratch_buf, scratch_first_index=63)`.
Single-block 6-thread (one per component); each writes mean|r_c| to
`scratch_buf[scratch_first_index + c]` with `__threadfence_system()`.
- 6 ISV slots wired with Pearls A+D: ISV[63..69) (Wiener offsets
189..207 — last slots in the post-A13 207-float wiener_state_buf).
- `GpuExperienceCollector::launch_reward_component_ema_inplace` now:
launches kernel → syncs stream → applies Pearls A+D in 6-iteration
loop → memsets reward_components_per_sample to zero (preserves
original behaviour). Degenerate-zero short-circuit per slot covers
the always-zero placeholder components (c=2 trail, c=5 bonus) plus
cold-start.
Cross-boundary wiring (mirrors A14/A15 precedent):
- 4 new fields on `GpuExperienceCollector`:
`reward_component_pearls_{wiener_host_ptr, scratch_dev_ptr,
scratch_host_ptr, isv_pinned_ptr}` — all NULL/0 until wired.
- New setter `set_reward_component_pearls_buffers(...)` on collector.
- New accessors on `GpuDqnTrainer`: `wiener_state_buf_host_ptr()`,
`producer_step_scratch_buf_dev_ptr()`,
`producer_step_scratch_buf_host_ptr()`, `isv_signals_pinned_ptr()`.
- New wire helper `FusedTrainingCtx::wire_reward_component_pearls_buffers`
pulls all 4 pointers from trainer, pushes into collector.
- Wired once in `training_loop.rs::init_gpu_experience_collector`
immediately after `set_curiosity_pearl_c_buffers`.
Orphan deletion (per `feedback_wire_everything_up.md`):
- Removed `GpuDqnTrainer::launch_reward_component_ema` (zero call sites
pre-deletion).
- Removed trainer-side `reward_component_ema_kernel: CudaFunction`
field (zero consumers post-launcher-deletion).
- Removed cubin loader + struct-init line.
- `REWARD_COMPONENT_EMA_CUBIN` static remains because the collector
still loads from it.
Tests:
- `sp4_reward_component_ema_writes_step_obs_via_pearl_a_then_converges_pearl_d`:
drives kernel with N=128 reward_components where r[i*6+c] = sign(i) ×
(c+1), asserts each slot ∈ ±1e-5 of (c+1), non-target slots remain 0,
Pearl A bootstrap + Pearl D convergence verified.
- The cross-boundary wiring path exercised in production by integration
smoke harness; this kernel-direct test isolates kernel signature +
Pearls A+D semantics.
- `cargo test -p ml --lib sp4_wiener_ema --offline` 6/6 passing.
Per `feedback_no_atomicadd.md`,
`feedback_no_htod_htoh_only_mapped_pinned.md`,
`feedback_no_partial_refactor.md`, `feedback_wire_everything_up.md`.
Build: `cargo check -p ml --lib --tests --offline` clean (11 pre-existing
warnings, no new warnings).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
5f800fe5c8 |
feat(sp4): Task A13.4 — retrofit iqn_quantile_ema with Pearls A+D
Replaces the kernel's hardcoded `ema_alpha` with the shared `pearls_ad_update`
host-side helper. 4 ISV slots retrofit (off-median IQN quantiles).
- Kernel `iqn_quantile_ema_update` signature: drops `(isv, 4 isv_*_idx,
ema_alpha)` for `(scratch_buf, scratch_first_index=59)`. Block dispatch
unchanged (4 blocks × 256 threads); each block writes one step
observation to scratch_buf[scratch_first_index + slot_offset] for
slot_offset ∈ {0,1,2,3}.
- 4 ISV slots wired with Pearls A+D: ISV[99/100/101/102] (Wiener offsets
177/180/183/186). Median tau_idx=2 intentionally skipped (already
surfaced via greedy-Q).
- Wrapper `GpuDqnTrainer::launch_iqn_quantile_ema(..., _ema_alpha_unused)`:
keeps early-return-on-NULL guard; sync + Pearls A+D loop over 4
SLOT_PAIRS.
Behavior: stationary signals converge to the same value at adaptive rate.
The 4 slots remain diagnostic-only (HEALTH_DIAG / risk-monitoring surface).
Tests: `sp4_iqn_quantile_ema_writes_step_obs_via_pearl_a_then_converges_pearl_d`
drives kernel with B=32 Q=5 TBA=12 controlled q surface where
Q[a, b*Q+tau_idx] = (tau_idx+1)*0.7. Asserts each slot ∈ ±1e-5 of
expected, median absent, non-target slots remain 0, Pearl A bootstrap +
Pearl D convergence verified.
Per `feedback_no_atomicadd.md`,
`feedback_no_htod_htoh_only_mapped_pinned.md`. Build: `cargo check -p ml
--lib --tests --offline` clean (11 pre-existing warnings, no new warnings).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
0e9d69787d |
feat(sp4): Task A13.3 — retrofit vsn_mask_ema with Pearls A+D
Replaces the kernel's hardcoded `ema_alpha` with the shared `pearls_ad_update`
host-side helper. 6 ISV slots retrofit (one per VSN feature group).
- Kernel `vsn_mask_ema_update` signature: drops `(isv, isv_first_index,
ema_alpha)` for `(scratch_buf, scratch_first_index=53)`. Per-group mean
writes to `scratch_buf[53..59)`. Single `__threadfence_system()` after
all 6 groups write.
- 6 ISV slots wired with Pearls A+D: ISV[105..111) (Wiener offsets
159..177). Wiener offset for group g: (53+g)*3.
- Wrapper `GpuDqnTrainer::launch_vsn_mask_ema(_ema_alpha_unused)`: sync
+ Pearls A+D loop over the 6 groups.
- `debug_assert_eq!(SL_NUM_FEATURE_GROUPS, 6)` guards against group-count
changes silently breaking the contiguous scratch layout.
Behavior: stationary signals converge to the same value at adaptive rate.
The 6 slots stay semantically identical (per-group mean of VSN softmax
mask); HEALTH_DIAG mirror + VSN focus monitor consume them unchanged.
Tests: `sp4_vsn_mask_ema_writes_step_obs_via_pearl_a_then_converges_pearl_d`
drives kernel with B=128 num_groups=6 mask `mask[b,g]=(g+1)/21` (rows sum
to 1, per-group mean = (g+1)/21). Asserts each slot ∈ ±1e-5, non-target
slots remain 0; verifies Pearl A bootstrap + Pearl D convergence.
Per `feedback_no_atomicadd.md`,
`feedback_no_htod_htoh_only_mapped_pinned.md`. Build: `cargo check -p ml
--lib --tests --offline` clean (11 pre-existing warnings, no new warnings).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
04fcbd27cf |
feat(sp4): Task A13.2 — retrofit moe_expert_util_ema with Pearls A+D
Replaces the kernel's hardcoded `alpha` with the shared `pearls_ad_update`
host-side helper. 9 ISV slots retrofit: 8 per-expert util + 1 gate entropy.
- Kernel `moe_expert_util_ema_update` signature: drops `(isv,
isv_util_base, isv_entropy_index, alpha)` for `(scratch_buf,
scratch_idx_util_base=44, scratch_idx_entropy=52)`.
- Single-block single-thread; computes per-expert col_mean in a single
forward pass (collapses prior dual-pass), writes each to scratch slots
[44..52) and Shannon entropy to slot 52.
- 9 ISV slots wired with Pearls A+D: ISV[118..126) (per-expert util,
Wiener offsets 132..156) + ISV[126] (gate entropy, Wiener offset 156).
- Launcher `gpu_moe_head.rs::launch_expert_util_ema`: drops alpha + isv
args; takes scratch_dev_ptr + 2 scratch_idx args.
- Wrapper `GpuDqnTrainer::launch_moe_expert_util_ema(_ema_alpha_unused)`:
sync + Pearls A+D loop over 9 slots via `apply_pearls(scratch_idx,
isv_idx)` closure.
- `debug_assert_eq!(MOE_NUM_EXPERTS, 8)` guards against K changes
silently breaking the contiguous scratch layout.
Behavior: stationary signals converge to the same value at adaptive rate.
The 9 slots stay semantically identical (per-expert col_mean, gate
entropy); HEALTH_DIAG mirror + the adaptive λ controller
`moe_lambda_eff_update` (which reads ISV[MOE_GATE_ENTROPY_EMA_INDEX])
continue to consume them unchanged.
Tests: `sp4_moe_expert_util_ema_writes_step_obs_via_pearl_a_then_converges_pearl_d`
drives kernel with B=128 K=8 perfect-uniform gate → analytical col_mean=1/8,
entropy=ln(8)≈2.0794. Asserts slot values ∈ ±1e-6/1e-5, non-target slots
remain 0; verifies Pearl A bootstrap + Pearl D convergence.
Per `feedback_no_atomicadd.md`,
`feedback_no_htod_htoh_only_mapped_pinned.md`. Build: `cargo check -p ml
--lib --tests --offline` clean (11 pre-existing warnings, no new warnings).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
74ed2f5008 |
feat(sp4): Task A13.1 — retrofit aux_heads_loss + aux_label_scale with Pearls A+D
Both kernels in `aux_heads_loss_ema_kernel.cu` retrofit in the same commit
per `feedback_no_partial_refactor.md` (single shared cubin, single producer
family).
- Kernel `aux_heads_loss_ema_update`: writes nb_loss/rg_loss scalars to
`producer_step_scratch_buf[41..43)` (slots 41=next-bar, 42=regime).
- Kernel `aux_label_scale_ema_update`: writes mean(|label|) to
`producer_step_scratch_buf[43]`.
- 3 ISV slots wired with Pearls A+D: ISV[113] (Wiener 123..126),
ISV[114] (Wiener 126..129), ISV[117] (Wiener 129..132).
- Launcher `AuxHeadsForwardOps::launch_loss_ema`: drops isv_dev_ptr +
isv_*_index + ema_alpha; takes scratch_dev_ptr + 2 scratch_idx args.
- Launcher `AuxHeadsForwardOps::launch_label_scale_ema`: same retrofit
pattern with a single scratch_idx arg.
- Wrapper `GpuDqnTrainer::launch_aux_heads_loss_ema(_ema_alpha_unused)`:
sync + Pearls A+D loop over both slots.
- `aux_heads_forward` mid-step launch (Step 2b — runs BEFORE
next_bar_loss_reduce + backward consume ISV[117]) gains inline sync +
Pearls A+D update so consumers see the up-to-date scale this step.
Behavior: stationary signals converge to the same value at adaptive rate
(Pearl D's α* derived from per-slot signal-vs-noise variance);
non-stationary signals respond Wiener-optimally faster. Slots stay
semantically identical (next-bar MSE, regime CE, label-scale mean_abs);
only the EMA blending logic changes.
Tests:
- `sp4_aux_heads_loss_ema_writes_step_obs_via_pearl_a_then_converges_pearl_d`:
nb_loss=0.5, rg_loss=1.2 stationary, both slots converge within 1%
after 1000 observations.
- `sp4_aux_label_scale_ema_writes_step_obs_via_pearl_a_then_converges_pearl_d`:
B=256 mixed-sign ±3.0 labels (mean_abs=3.0), step_obs ∈ ±1e-4 of
analytical mean, Pearl A bootstrap + Pearl D convergence verified.
Per `feedback_no_atomicadd.md`,
`feedback_no_htod_htoh_only_mapped_pinned.md`,
`feedback_no_partial_refactor.md`. Build: `cargo check -p ml --lib --tests
--offline` clean (11 pre-existing warnings, no new warnings).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
aada419de3 |
feat(sp4): Task A13.0 — retrofit h_s2_rms_ema with Pearls A+D + buffer growth
Replaces the kernel's hardcoded `ema_alpha` with the shared `pearls_ad_update`
host-side helper (Task A3). Buffer growth in same commit so subsequent A13.x
retrofits land on a stable scratch/Wiener layout.
- Kernel `h_s2_rms_ema_update` signature: `(h_s2, B, SH2, scratch_buf,
scratch_idx)`. Reduces RMS = sqrt(sum_sq/(B*SH2)) via the existing
256-thread shmem tree (no atomicAdd) and writes step_obs to
`producer_step_scratch_buf[40]` with `__threadfence_system()`.
- Launcher `launch_h_s2_rms_ema(_ema_alpha_unused: f32)` syncs the stream,
then applies Pearls A+D host-side via zero-copy mapped-pinned reads/
writes of `isv_signals_pinned[H_S2_RMS_EMA_INDEX=96]` and
`wiener_state_buf[120..123)` (= scratch slot 40 × 3). Degenerate-zero
short-circuit before mutating ISV/Wiener state.
- Buffer growth: `SP4_PRODUCER_COUNT 47 → 69` (40 SP4 + 29 Task A13
retrofit producers); `wiener_state_buf 141 → 207` floats;
`producer_step_scratch_buf` grows to 69 entries.
- Stable-layout doc-comments updated: `producer_step_scratch_buf` field
comment, `launch_sp4_target_q_p99` slot-table comment, 5 launcher
`wiener_offset + 2 < 141` safety comments (now `< 207`),
`reset_sp4_wiener_state` doc + body comment, and
`state_reset_registry.rs::sp4_wiener_state` description.
- Test cubin reference `SP4_PRODUCER_COUNT: usize = 47` in
`tests/sp4_producer_unit_tests.rs` updated to 69 across all 6
occurrences.
Behavior: stationary signals converge to the same RMS at adaptive rate
(Pearl D's α* derived from per-slot signal-vs-noise variance);
non-stationary signals respond Wiener-optimally faster. Cold-path producer
with no consumer-facing change beyond the EMA mechanism — slot 96 is read
by `mag_concat_qdir`'s adaptive-scale path and stays semantically identical
(RMS of `save_h_s2`).
Tests: new `sp4_h_s2_rms_ema_writes_step_rms_via_pearl_a_then_converges_pearl_d`
unit test (`#[ignore]`-gated for GPU) drives the production kernel kernel-
direct on a constant-5.0 stationary signal of B=4 SH2=64, asserts
step_rms ≈ analytical RMS=5.0 ± 1e-4, asserts non-target scratch slots
remain 0, then exercises Pearl A bootstrap (returns step_rms directly +
seeds x_lag) and Pearl D convergence (1000 stationary observations →
x_mean within 1% of 5.0).
Per `feedback_no_atomicadd.md`, `feedback_no_htod_htoh_only_mapped_pinned.md`,
`feedback_no_partial_refactor.md`. Build: `cargo check -p ml --lib --tests
--offline` clean (11 pre-existing warnings, no new warnings); `cargo test
-p ml --lib sp4_wiener_ema --offline` 6/6 passing.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
138d41b761 |
refactor(ml): drop 8 leaf sub-crates from trading + backtesting services
Two-part fix that finally removes ~35% of ml-* sub-crates from
service-binary build closures:
1. crates/ml/Cargo.toml: 8 leaf-level ml-* sub-crates become optional
behind feature flags (one feature per `pub mod` in lib.rs):
ml-backtesting ↔ feature `backtest-mod` (ml::backtesting)
ml-paper-trading ↔ feature `paper-trading-mod` (ml::paper_trading)
ml-stress-testing ↔ feature `stress-testing-mod` (ml::stress_testing)
ml-explainability ↔ feature `explainability-mod` (ml::explainability)
ml-universe ↔ feature `universe-mod` (ml::universe)
ml-regime-detection ↔ feature `regime-detection-mod`(ml::regime_detection)
ml-validation ↔ feature `validation-mod` (ml::validation)
ml-data-validation ↔ feature `data-validation-mod` (ml::data_validation)
Aggregated into the umbrella `full-stack` feature, which is in
`default = [...]` so existing `ml.workspace = true` callers see
identical behavior (testing/integration, crates/backtesting).
2. services/trading_service + services/backtesting_service: switched
from `ml.workspace = true` to explicit `path = "../../crates/ml"`
form with `default-features = false`. This is necessary because
cargo 1.89 silently ignores `default-features = false` when
combined with `workspace = true` (a known cargo limitation).
Path form bypasses the workspace dep and applies the opt-out.
Verified by `cargo tree -p X | grep ml-* | wc -l`:
trading-service 23 → 16 (-30%)
backtesting-service 23 → 16 (-30%)
Source-grep verified neither service references any of the 8 gated
modules (`use ml::backtesting`, `use ml::explainability`, etc. all
zero hits in src/). The only `ml::explainability` mention in
trading-service is a string in an error log — not a code path.
ml-training-service and trading-agent-service were already on path
form with default-features = false; they're unaffected here.
cargo check --workspace passes.
|
||
|
jgrusewski
|
ea31ebfe38 |
feat(sp4): Tasks A14+A15 — Pearl C engagement counter in 5 Adam kernels
Layer A scaffolding for post-clamp feedback-loop self-correction. A14: All 5 Adam kernels (dqn/iqn/iql/attn/curiosity) extended with register-then-tree-reduce engagement counter. Per-thread `local_engage` set on clamp engagement; block tree-reduce (no atomicAdd, mirrors dqn_grad_norm_kernel pattern); single block-leader writes per-block count to clamp_engage_per_block_buf[engage_buf_offset + blockIdx.x]. `engage_buf_offset == -1 (SP4_ENGAGE_OFFSET_DISABLED)` skips writeback. A15: Host-side rate-deficit check in pearl_c_post_adam_engagement_check. Sums per-block counts via mapped-pinned zero-copy reads, computes engagement_rate, rate_deficit = rate - 0.01. Applies Pearls A+D via pearl_c_rate_deficit_ema (host-only [f32;8]) + pearl_c_rate_deficit_state_buf (MappedF32Buffer[24] = 8 groups × 3 Wiener floats). Wired post-graph in FusedTrainingCtx::run_full_step for groups 0/3/4/5/6 and post-curiosity_train in training_loop for group 7. 3 design issues resolved per Layer A scope: 1. Curiosity sub-launches: SP4_ENGAGE_BUF_LEN extended 2048->2816 (= 11 x 256). Curiosity sub-launches use offsets 1792/2048/2304/2560 (W1/b1/W2/b2). Host-side check sums all 4. Allocation + reset-registry updated. 2. TLOB/Attn sharing attn_adam_kernel: TLOB passes SP4_ENGAGE_OFFSET_DISABLED=-1 (silently skips Pearl C). Attn writes to its slot (offset 6x256=1536). Layer B can refactor if needed. 3. DQN main Adam covers groups 0/1/2 in single launch: Layer A accounts only under group 0 (DqnTrunk). Groups 1/2 Pearl C tracking deferred to Layer B's per-group sub-launch decision. Wiring path: GpuDqnTrainer exposes nan_flags_buf_ptr() + clamp_engage_per_block_buf_dev_ptr(). FusedTrainingCtx adds wire_aux_trainer_pearl_c_buffers() for IQN/IQL hi+lo/Attn. GpuExperienceCollector adds set_curiosity_pearl_c_buffers() for the curiosity trainer. Wiring fires once in init_gpu_experience_collector. State reset registry (Task A12 follow-up): - sp4_clamp_engage_counters description updated to reflect 2816 buffer length and 4 distinct curiosity sub-launch offsets - New entries: sp4_pearl_c_rate_deficit_state (24 mapped-pinned floats, Pearls A+D Wiener state) and sp4_pearl_c_rate_deficit_ema (host-only [f32; 8] EMA surrogate). Both reset to zero at fold boundary so Pearl A's first-observation sentinel fires on the new fold's first engagement-rate-deficit observation. Dispatch arms wired in reset_named_state alongside existing sp4_* entries. Layer A: Pearl C is observability scaffolding. Mech 9's clamp still uses hardcoded 100xQ_ABS_REF. Engagement counters fire correctly, rate_deficit EMA tracks. Force-bump branch logs via tracing::debug only; Layer B's atomic flip activates the actual ISV mutation. cargo check --lib --tests clean. state_reset_registry tests pass. sp4_isv_slots::tests::pearl_c_engage_buf_layout pass. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
ce11d56eda |
feat(sp4): Task A12 — StateResetRegistry entries for SP4 + Wiener + Pearl C
Layer A additive: at fold boundary, all SP4 ISV bound slots, Wiener
state, and Pearl C engagement counters reset to 0 (Pearl A sentinel).
Pearl A's first-observation replacement requires both `prev_x_mean`
(the ISV bound slot) AND `state.x_lag` (Wiener triple offset 2) to be
exactly 0 when a producer first fires in a new fold. Without these
resets, the new fold's first producer launch would EMA against fold-N's
stale state — exactly the cross-fold anchor staleness Mech 8's reverted
slow_ema reset was trying to fix imperfectly.
Per `feedback_no_partial_refactor.md`, every half of the SP4 fold-reset
contract migrates in the same commit:
- 9 SP4 ISV bound family entries added to StateResetRegistry, dispatched
by `reset_named_state` to write 0.0 to every slot in [131..171):
- sp4_target_q_bound (slot 131)
- sp4_atom_pos_bounds (slots 132..136 = 4 branches)
- sp4_weight_bounds (slots 136..144 = 8 param groups)
- sp4_adam_m_bounds (slots 144..152)
- sp4_adam_v_bounds (slots 152..160)
- sp4_wd_rate_bounds (slots 160..168)
- sp4_grad_clip_bound (slot 168)
- sp4_h_s2_bound (slot 169)
- sp4_l1_lambda_trunk (slot 170)
- sp4_wiener_state entry — bulk write_bytes(0) on `wiener_state_buf`
(141 mapped-pinned f32 = 47 producers × {sample_var, diff_var, x_lag}).
New `GpuDqnTrainer::reset_sp4_wiener_state` helper mirrors the existing
`reset_fast_grad_norm_ema` pattern.
- sp4_clamp_engage_counters entry — bulk write_bytes(0) on
`clamp_engage_per_block_buf` (2048 mapped-pinned i32 = 8 groups ×
256 blocks). New `GpuDqnTrainer::reset_sp4_clamp_engage_counters`
helper. Each fold restarts Pearl C engagement-rate tracking from a
clean 0 baseline.
Total 11 registry entries (group-keyed, not per-slot) — follows the
existing convention of `isv_grad_balance_targets` (1 name → 4 slots)
and `isv_q_quantiles` (1 name → 8 slots). Family-level entries keep
the dispatch table bounded while still providing per-family auditability.
Audit doc `docs/dqn-wire-up-audit.md` extended with SP4 Task A12 entry.
cargo check --lib --tests clean. state_reset_registry + soft_reset
smoke tests pass.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
392fc7d698 |
feat(sp4): Tasks A10+A11 — wire all 5 SP4 producer launches in training_loop
Layer A: cold-path SP4 producer launches added next to launch_h_s2_rms_ema
and the surrounding ISV producers (mag_concat h_s2_rms_ema, fold_warmup,
iqn_quantile_ema, vsn_mask_ema, aux_heads_loss_ema, moe_expert_util_ema,
moe_lambda_eff_update). All five fire once per training step before the
HEALTH_DIAG line emit per the post-cascade-fix invariant (
|
||
|
jgrusewski
|
8f93c11525 |
fix(sp4): Task A7 fix-up #2 — Curiosity sub-buffer support, all 8 groups wired
Curiosity stores params/grads/Adam state as 4 non-contiguous sub-buffers (w1, b1, w2, b2 — each its own CudaSlice<f32>). The single-pointer kernel signature couldn't describe them. The first fix-up wired 4 of 5 aux groups; Curiosity (group 7) was deferred. This commit extends the param_group_oracle kernel to accept a sub-buffer table (mapped-pinned u64 ptr-arrays + i32 counts), with `n_sub=1` for groups 0-6 (existing behavior) and `n_sub=4` for Curiosity. Pass A/B/C (p99 histograms) iterate sub-buffers via a new `sp4_histogram_p99_multi<BLOCK_SIZE>` template that mirrors the original three-pass structure but loops sub-buffers within the max-reduce + binning passes; Pass 3 (cumulative-from-top) divides by `total_count`. Pass D (4-way reduce) iterates sub-buffers in the accumulator loop; Pass E (L1 trunk only) reads `grads_ptrs[0]` (group 0 has n_sub=1 by construction). `Sp4ParamGroupBufs` redefined from a flat quartet to `Vec<Sp4SubBuffer>`. Convenience constructors `Sp4ParamGroupBufs::single(...)` and `Sp4ParamGroupBufs::empty()` keep call sites ergonomic; `total_count()` for kernel arg. Switched Copy → Clone since the descriptor now owns a Vec. `SP4AuxBuffers` extended with `curiosity` field (5-tuple from 4-tuple). Added 16 public accessors to GpuCuriosityTrainer (grad + Adam state) and 8 to CuriosityWeightSet (weights + lengths) for the 4 sub-buffers × 4 signal types. `oracle_subbuf_table_buf: MappedU64Buffer` (4×4 = 16 u64s) and `oracle_subbuf_counts_buf: MappedI32Buffer` (4 i32s) allocated at construction. Launcher overwrites entries [0..n_sub) per group launch via volatile writes; zeros the unused tail (defence-in-depth so a kernel bug reading past `n_sub` lands on count=0 no-op). Inter-launch `stream.synchronize()` added so the next iteration's host writes don't race with the in-flight kernel's coalesced loads from the persistent table. Cold-path producer; per-launch sync cost is negligible vs the kernel work. `build_sp4_aux_buffers` signature changed: takes `Option<&CuriosityWeightSet>` and `Option<&GpuCuriosityTrainer>` since Curiosity state lives outside FusedTrainingCtx (owned by GpuExperienceCollector). Layer B's training-loop caller threads them in from `collector.curiosity_weight_set()` + the collector's `curiosity_trainer` field; both must be Some together (caller responsibility — they live on the same collector). Passing None for either yields an empty curiosity descriptor and the launcher silently skips group 7. `param_group_buffers` return type changed from `Option<(u64, u64, u64, u64, usize, i32, i32)>` to `Option<(Sp4ParamGroupBufs, i32, i32)>`. All groups now return Some(...) (Curiosity included); None reserved for forward-compat. GPU test extended: group 7 exercises 4 sub-buffers of distinct shapes [1024, 32, 1024, 32] (w1/b1/w2/b2-like sizes scaled to keep test runtime small while still exercising the multi-sub-buffer iteration), each with its own seed offset so distinct sub-buffers have distinct distributions — catches buffer-mixup bugs in the kernel's sub-buffer iteration. Reference computation builds union vectors and computes p99/WD_RATE over the union. Test launcher refactored to `&[TestSubBuffer]` slice matching the production kernel's table-packing layout. All 8 SP4 param-groups now produce real outputs in Layer A. The launcher's `count == 0` short-circuit retained for the optional aux- trainer fallback (init failures on gpu_iqn / gpu_attention / curiosity). `MappedU64Buffer` gained a manual Debug impl (warn missing_debug_impls) for parity with MappedU32Buffer. cargo check -p ml --lib --tests clean. Workspace clean. Layer B can now safely consume all 8 ISV[WEIGHT_BOUND/ADAM_M_BOUND/ADAM_V_BOUND/WD_RATE[group]] slots. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
64298b34c0 |
fix(sp4): Task A7 fix-up — wire 4 of 5 aux param-groups (IQN/IQL/Attn)
A7 (commit
|
||
|
jgrusewski
|
a72468666f |
feat(sp4): Task A9 — h_s2_p99 producer for ISV[H_S2_BOUND=169]
New separate producer kernel reading save_h_s2 [B × SH2] via sp4_histogram_p99<256>. Output → ISV[169] via host-side Pearls A+D applied at scratch slot 39. Mirrors Task A5's pattern exactly. NB: distinct from existing h_s2_rms_ema_update (slot 96) which is a different signal (RMS, not p99). Task A13 will retrofit that existing kernel to also use Pearls A+D, in a separate commit. GPU test verifies p99 ≈ 2.576 for 4096 |N(0,1)| samples within 5% tolerance (rel_err = 0.336% on local RTX 3050 Ti) and exercises Pearl A's sentinel branch (first observation replaces sentinel directly, x_lag seeds to step_p99, both variances stay zero). No consumer wired yet. Behaviour unchanged. cargo check --lib --tests clean. 6/6 SP4 producer GPU tests passing locally (A4 + A5 + A6 + A7 + A8 + A9). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
611d2663cb |
feat(sp4): Task A8 — grad_norm_p99 producer (degenerate single-element)
grad_norm_buf is single-element f32; p99 == the element. Kernel copies to scratch slot 38 with __threadfence_system; host applies Pearls A+D to update ISV[GRAD_CLIP_BOUND_INDEX=168]. Trivial pattern — establishes the launcher template for Mech 6's eventual ISV read in Layer B. GPU test verifies first-observation Pearl A replacement (input=42.0 → scratch[38]=42.0 bit-identical → new_x_mean=42.0, x_lag=42.0, both variances stay zero) and that all non-target scratch slots remain 0. No consumer wired yet. Behaviour unchanged. cargo check --lib --tests clean. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
4f13e2ca37 |
feat(sp4): Task A7 — Pearl B fused per-param-group statistics oracle
Single kernel per group reads (params, grads, adam_m, adam_v) once and fuses 5 passes (4 for non-trunk): Pass A: WEIGHT_BOUND[group] = p99(|params|) via sp4_histogram_p99 Pass B: ADAM_M_BOUND[group] = p99(|adam_m|) via sp4_histogram_p99 Pass C: ADAM_V_BOUND[group] = p99(|adam_v|) via sp4_histogram_p99 Pass D: WD_RATE[group] = |Σ w·g| / max(Σ w², ε_div); side: mean|g|, mean|w| Pass E (group 0 only): L1_LAMBDA[trunk] via gradient-direction entropy deficit × (mean|g|/mean|w|) magnitude scale Launcher loops over 8 groups (DQN trunk/value/branches, IQN, IQL hi/lo, attn, curiosity), one launch per group. Phase 2 applies Pearls A+D to each of the 4-5 outputs per group via host-side pearls_ad_update. Producer-scratch slots 5..38 reserved for this task's 33 outputs: [5..13)=WEIGHT, [13..21)=ADAM_M, [21..29)=ADAM_V, [29..37)=WD_RATE, 37=L1. Three of eight groups wired (DQN trunk/value/branches); five aux groups (IQN, IQL hi/lo, attn, curiosity) skip silently because their backing trainers live on FusedTrainingCtx, not GpuDqnTrainer. Wiring those requires either threading buffer pointers through the launcher signature or hoisting the launcher onto FusedTrainingCtx — both follow-up changes scoped beyond Task A7. Documented in launcher + param_group_buffers docstrings. Per-group GPU tests verify outputs within 5% rel_err (p99 quantization) or 2% rel_err (analytical formulas). Local RTX 3050 Ti: max WEIGHT rel_err=0.589%, max ADAM_M=0.589%, max ADAM_V=0.554%, max WD_RATE=0.001% across all 8 group shapes; group 0 L1 within tolerance. Pearl B 4× memory-bandwidth reduction vs naive 4 separate per-group producer kernels: each buffer read once for all 4-5 outputs. No consumer wired yet — Adam kernels still take hardcoded weight_decay + weight_clamp_max_abs config args. Behavior unchanged. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
b4c28811ab |
feat(sp4): Task A6 — atom_pos_p99 producer × 4 branches
Single kernel parameterized by branch slice. Launcher loops 0..4, launching once per branch with distinct scratch slot (1..5). Phase 2 applies Pearls A+D host-side per branch, writing to ISV[ATOM_POS_BOUND[ branch]] + wiener_state. Same end-to-end pattern as Task A5. GPU test verifies per-branch independence: 4 scales of |N(0,1)| samples (1×, 10×, 100×, 1000×) → 4 distinct p99 values within 5% tolerance. No consumer wired yet. Behavior unchanged. cargo check --lib --tests clean. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
cd037084b2 |
feat(sp4): Task A5 — target_q_p99 producer kernel + Pearls A/D wire-up
First end-to-end SP4 producer. Kernel reads denoise_target_q_buf, computes p99(|target_q|) via sp4_histogram_p99<256>, writes step_obs to producer_step_scratch_buf[0] with __threadfence_system. Launcher syncs, applies Pearls A+D via pearls_ad_update (zero-copy mapped-pinned reads of ISV[TARGET_Q_BOUND_INDEX=131] + wiener_state_buf[(131-base)*3]), writes new x_mean back to ISV + state back to wiener_state. Cold-path launch (no captured graph in this task; Task A10 may move to captured). Producer-step-scratch slot 0 reserved for TARGET_Q_BOUND (stable layout documented in launcher comment for Tasks A6-A11 to extend). GPU unit test verifies kernel writes step_p99 ≈ p99(|N(0,1)|) within 5% rel_err on 4096 Box-Muller samples, then exercises Pearl A's sentinel branch (sentinel ISV + zero Wiener state → first-observation replacement). Helper asserts non-target scratch slots stay zero. No consumer wired yet — Mech 1's clamp still uses 10 × Q_ABS_REF.max(1.0). Behavior unchanged. cargo check --lib --tests clean. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
7540df1ecf |
feat(sp4): Task A4 — linear-histogram p99 device function + GPU unit test
Header-only `__device__` function `sp4_histogram_p99<BLOCK_SIZE>(buf, count)`
returning p99(|buf|) via three-pass single-block algorithm:
Pass 1: block-wide max-reduce → step_max (0.0 short-circuit on degenerate)
Pass 2: linear-spaced [0, step_max] binning into 256 bins via per-warp
tiles (no atomicAdd; lane-collisions cost <0.012% precision)
Pass 3: cumulative-from-top → p99 = bin upper-edge
Returns 0.0 for degenerate step_max=0 (caller skips ISV update). Linear
spacing chosen over log because SP4 producers care about resolution at
the top of the distribution — top bin width = step_max/256 ≈ 0.39%, well
within the 1% quantile precision budget.
Wrapper kernel `sp4_histogram_p99_test_kernel` exposes the device fn for
Rust testing; mapped-pinned scalar output with __threadfence_system() so
host read_volatile sees the result (no memcpy_dtoh). Build.rs
registration mirrors thompson_test_kernel.cu pattern + adds
rerun-if-changed for the .cuh header.
Unit test: 4096 deterministic |N(0,1)| Box-Muller samples, sorted
ground-truth p99 ≈ 2.576 z-score one-tailed, asserts rel_err < 5%
against device output. GPU-gated (#[ignore]). Local L40S run:
true_p99=2.59758, computed_p99=2.59858, rel_err=0.039% — passes.
No producer wired yet — header is library code included only by the
test wrapper; SP4 Tasks A5-A9 add the magnitude-bound producer kernels
that #include "sp4_histogram_p99.cuh" directly. Behaviour unchanged.
cargo check --lib --tests clean.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
9ece1a4daa |
feat(sp4): Task A3 — Pearls A+D shared host-side helper
Single source-of-truth implementation of: - Pearl A (first-observation bootstrap): sentinel-detect at fold reset, replace x_mean directly with first observation when prev_x_mean=0 AND state.x_lag=0. Bypass Pearl D's Wiener math. - Pearl D (Wiener-optimal adaptive α): for t≥1, α* = diff_var / (diff_var + sample_var + ε_div); variances tracked at uniform meta-α. 6 unit tests: Pearl A sentinel replacement; Pearl D anchors at stationary signal; Pearl D tracks step-change; Pearl D does NOT subsume Pearl A at t=0 (mathematical correctness check from spec self-review); meta-constants are structural; Pearl A only fires when both x_mean and x_lag are zero (does not re-fire post-Pearl-D). ALPHA_META = 1e-3 (structural — single uniform meta-rate, no per-signal tuning). EPS_DIV = 1e-8 (Adam-ε numerical category). EPS_CLAMP_FLOOR = 1.0 (consumer cold-start floor). No consumers wired yet — helper is library code unused by the producer pipeline. Behavior unchanged. cargo check clean. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
a2c14cb063 |
feat(sp4): Task A2 — allocate 3 mapped-pinned buffers for Pearls B/C/D
Layer A additive: three mapped-pinned buffers added to GpuDqnTrainer. - wiener_state_buf (141 floats) — Pearl D state (47 producers × 3 floats: sample_var, diff_var, x_lag). Per `feedback_no_htod_htoh_only_mapped_pinned`. - clamp_engage_per_block_buf (2048 ints) — Pearl C engagement counters (8 param-groups × 256 max blocks per Adam launch). - producer_step_scratch_buf (47 floats) — per-producer per-step step_observation output. Host applies Pearls A+D to map step_obs to ISV bound slot via pearls_ad_update (Task A3, pending). All three zero-initialized at construction (Pearl A sentinels). Reset registry entries follow in Task A12. No consumers wired yet — buffers are reserved but unread. Behavior unchanged. cargo check clean. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
206ebd3558 |
fix(sp4): Task A1 review fix-ups — fingerprint seed + branch count + doc refresh
3 important + 1 optional findings from code quality review: 1. layout_fingerprint_seed() now lists all 40 SP4 slots and bumps `ISV_TOTAL_DIM=131` -> `ISV_TOTAL_DIM=171`. Without this, the fail-fast checkpoint-load guard would not detect the SP4 layout extension — a binary built against new code would falsely compare-equal to old checkpoints' fingerprints. 2. Added `SP4_BRANCH_COUNT=4` constant and `debug_assert!(branch < SP4_BRANCH_COUNT)` in `atom_pos_bound`. Test extended to verify atom_pos_bound max does not alias into WEIGHT_BOUND family. 3. Refreshed stale content in docs/isv-slots.md: ISV_TOTAL_DIM 96->171, fingerprint location [37..39)/[47..49) -> [115..117), table entries [94]/[95] -> [115]/[116]. 4. Added `debug_assert!(group < SP4_PARAM_GROUP_COUNT)` to weight_bound, adam_m_bound, adam_v_bound, wd_rate accessors (symmetric with the atom_pos_bound change). cargo check clean, cargo test slot_layout_is_contiguous_and_total_40 passes. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
c724fa99be |
feat(sp4): Task A1 — define 40 ISV slot indices for signal-driven bounds
Layer A additive: ISV_TOTAL_DIM 131 → 171. New module sp4_isv_slots.rs defines TARGET_Q_BOUND, ATOM_POS_BOUND[4], WEIGHT_BOUND[8], ADAM_M_BOUND[8], ADAM_V_BOUND[8], WD_RATE[8], GRAD_CLIP_BOUND, H_S2_BOUND, L1_LAMBDA_TRUNK. 40 slots total. ParamGroup enum for type-safe per-group launches. No consumers wired yet — slots are reserved but unused. Behavior unchanged. docs/isv-slots.md updated with the SP4 reservation table per Invariant 7. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
6e71576537 |
Revert "diag(dqn): SP3 Mech 10 — investigate F0=21.09 plateau via 1e30 multiplier"
This reverts commit
|
||
|
jgrusewski
|
074ff8c32a |
diag(dqn): SP3 Mech 10 — investigate F0=21.09 plateau via 1e30 multiplier
Mech 10 with 100× multiplier (commit
|
||
|
jgrusewski
|
55576d047b |
fix(dqn): SP3 Mech 10 — ISV-driven h_s2 activation clamp + slot 49 diag
Real root cause of the F1 step-1000 NaN, identified by smoke-test-2xrxk diagnostic. With Mech 9 clamping weights at 100×Q_ABS_REF=5000, F0/F2 trained clean (weights stay at natural scale ~0.5) but F1's regime shift drove weights toward the clamp ceiling. With weights pinned at 5000 and K=256 inner-dim, forward GEMM accumulators compound by ~80000× per layer; trunk depth ≥6 (encoder + VSN + GRN-blocks + bottleneck) hits f32_max ≈ 3.4e38 in the chain. Slot 12 (save_h_s2) firing was the smoking gun — h_s2 itself goes non-finite mid-forward. Slots 39+43+48 (IQN-side) stayed clean, disproving the close-out v2 hypothesis (IQN partial refactor) and confirming the surface that wasn't protected: ACTIVATIONS. Mech 1+2 clamp targets+atoms. Mech 9 clamps weights. Nothing clamped the trunk's intermediate hidden state. Mech 10 closes that gap by clamping h_s2 at 100 × H_S2_RMS_EMA.max(1.0) immediately after the trunk + temporal pipeline finalises it, before any downstream consumer reads it (branches, IQN, value head, replay save). Implementation: - Reuses existing launch_clamp_finite_f32 (dqn_utility_kernels.cu:462) which does isfinite(v) ? clamp(v) : 0.0f — NaN→0, Inf→0, finite clamped. No new kernel. - ISV-driven via H_S2_RMS_EMA_INDEX=96 (already consumed by mag_concat adaptive scale per P4.T2c.3c.6 — no new slot). - ε on multiplier per SP1 pearl: isv[96].max(1.0). Cold-start RMS near 0 floors the bound at 100×1=100, not at 0. - Inline NaN check (slot 12) preserved BEFORE the clamp per SP1 diagnostic-ordering pearl — sentinel sees the original Inf/NaN before sanitization replaces it with 0. - Insertion site: end of submit_forward_ops_main 1c temporal-pipeline block (after mamba2_step + apply_regime_dropout — the last writers to save_h_s2), immediately before launch_curiosity_inference and the loss path. Captured in the same forward child graph. New diagnostic slot 49 = h_s2_max_abs at threshold 1e3 × H_S2_RMS_EMA.max(1.0). Mirrors Mech 5 family pattern (clamp at 100×, diagnostic at 1000×) — sentinel that NEVER fires under normal Mech 10 operation. If it does fire, the activation clamp is disabled or H_S2_RMS_EMA itself drifted. Buffer bumps: nan_flags_buf 49→50, metadata 25→26, fused-kernel block count 25→26. Kernel signature gains second ISV-derived float param h_s2_rms_ema_eff (distinct from q_abs_ref_eff — h_s2 lives in activation space, not Q space; don't conflate the two ISV bases). Both name tables in training_loop.rs (halt_nan + halt_grad_collapse) extended to 50 entries. No new ISV slots, no new kernels, no graph-topology surprise (one extra launch in captured graph). cargo check clean. Audit docs docs/dqn-wire-up-audit.md (Mech 10 entry prepended) and docs/dqn-backward-nan-audit.md (Mech 10 row in mechanism table + slot 49 row in slot-allocation table + ISV bindings note + full Mech 10 section) updated per Invariant 7. Validation: deferred to one L40S smoke. Expected: F1 trains past step 1000; slots 36-43 + slot 49 stay quiet. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
0d7e27d612 |
fix(dqn): SP3 v2 follow-up — sync slot 48 name to halt_grad_collapse table
The SP3 close-out v2 commit
|
||
|
jgrusewski
|
1bcc703920 |
fix(dqn): SP3 close-out v2 — Mech 9 to all 5 Adam kernels + IQN weight diag
Real fix for the F1 step-3540 NaN. Commit
|
||
|
jgrusewski
|
8956c2fb77 |
fix(dqn): SP3 close-out — Mech 9 post-Adam weight clamp + Mech 8 revert
Coordinated close-out of SP3 Q-learning numerical-stability per
feedback_no_partial_refactor.
REVERT Mech 8 (slow_ema fold-boundary reset). smoke-test-rxhjh on
|
||
|
jgrusewski
|
b8a7ac6f70 |
fix(dqn): SP3 Mech 8 — slow_ema fold-boundary reset (real root cause)
User-identified root cause: grad_norm_slow_ema (α=0.001, half-life
~693 steps) persists across fold boundaries while every other
distribution-tracking signal (Adam m/v, target nets, atom positions)
gets reset. Mech 6's upper_bound formula (100 × slow_ema × isv) was
anchored to the WRONG scale during F1 ramp-up, driving the
non-monotonic multiplier-tuning dance across smokes:
- smoke-test-fxvkk (mult=100×): F0=44, F1 NaN @ 3720
- smoke-test-ftdjz (mult=100×, +Mech7): F0=38, F1 collapse @ 2040
- smoke-test-d25vq (mult=5×): F0=21, F1 collapse @ 2820
The multiplier was searching the wrong dimension — the anchor itself
was stale.
Three coordinated changes (per feedback_no_partial_refactor):
1. RESTORE Mech 6 multiplier 5× → 100×. The original 100× was correct
for steady-state; the F1 saturation was driven by anchor staleness,
not multiplier looseness. Tightening it harmed F0 (over-clip during
ramp-up when slow_ema lagged grad_norm_ema).
2. ADD reset_grad_norm_slow_ema method on GpuDqnTrainer. Zeroes the
mapped-pinned scalar. First step of new fold builds up slow_ema
fresh, with MIN_CLIP=1.0 floor active during the brief transient.
3. WIRE Mech 8 call in fused_training.rs::reset_for_fold, alongside
Mech 4's existing Adam resets. Mech 6's anchor now aligns with
the new fold's grad scale from step 1 — same philosophy as Mech 3
(target net hard-sync) and Mech 4 (Adam EMA reset).
Net SP3 design: Mech 6 stays at 100× multiplier (broad, principled
headroom), Mech 8 keeps the anchor honest. The pair is more robust
than either change alone:
- Without Mech 8: anchor is stale, multiplier tuning has no winning
setting (5× hurts F0 ramp, 100× lets F1 saturate at slot 36).
- With Mech 8: anchor is fresh per fold; 100× multiplier provides
legitimate per-step headroom over CURRENT fold's grad norm.
Mech 7 stays reverted (per-element clip was misdiagnosis — over-clipped
legitimate gradient outliers without addressing the saturation root
cause).
F0 risk: low — F0 starts with slow_ema=0 anyway (cold start), so Mech 8
is a no-op on F0. Only changes F1+F2 fold-boundary behavior.
F1+F2 expectation: Mech 6's upper_bound now scales with the CURRENT
fold's grad norm, providing legitimate ~10× headroom per step without
allowing Adam EMA saturation. Slot 36-42 should stay quiet.
|
||
|
jgrusewski
|
f67ede94fb |
fix(dqn): SP3 Mech 6 v2 — tighten clip multiplier 100x -> 5x; revert Mech 7
Smoke smoke-test-ftdjz (commit |
||
|
jgrusewski
|
d9a4d98a3d |
feat(dqn): SP3 Mech 7 — per-element gradient clip in Adam kernel
Smoke smoke-test-fxvkk (commit
|
||
|
jgrusewski
|
48c25d9997 |
feat(dqn): SP3 Mech 6 — anchored upper bound on adaptive grad clip
Adds an upper bound to update_adaptive_clip's new_clip formula to
prevent clip-drift pathology: consecutive elevated samples were
ratcheting the EMA-driven clip threshold upward without bound,
eventually rendering clipping a no-op against in-distribution drift.
Smoke smoke-test-5rqzs (commit
|
||
|
jgrusewski
|
5f7db7d8d8 |
feat(dqn): SP3 — name table entries for slots 36-47 diagnostic
Replaces rsv36-rsv47 placeholders with the SP3 Mech 5 diagnostic slot names. Both name table sites (halt_nan + halt_grad_collapse) updated identically per shared-contract migration principle (feedback_no_partial_refactor). Slot names mirror the audit doc table: - 36-39: Adam m max-abs (trunk/value/branch/IQN) - 40-43: Adam v max-abs (trunk/value/branch/IQN) - 44-45: Weight max-abs (trunk/heads) - 46: target_q post-clip - 47: atom_span_max When the fused kernel sets a slot bit, the readback log line names the buffer that exceeded its ISV-derived threshold — providing direct observability for SP3 mechanism effectiveness. |
||
|
jgrusewski
|
45e077188f |
feat(dqn): SP3 Mech 5 — fused kernel extended for slots 36-47 threshold checks
Extends dqn_nan_check_fused_f32_kernel to handle 24 slots (12 NaN-only + 12 ISV-threshold). Per-slot thresholds computed inline in the kernel from a single q_abs_ref_eff arg (no HtoD per step; no separate thresholds buffer needed). Slot index → threshold mapping: - 12-15 (slots 36-39): Adam m ≥ 100 × q_abs_ref_eff - 16-19 (slots 40-43): Adam v ≥ 1e6 × q_abs_ref_eff² - 20-21 (slots 44-45): Weight max ≥ 1e3 × q_abs_ref_eff - 22 (slot 46): target_q ≥ 95 × q_abs_ref_eff (= 9.5 × max_abs_target_q) - 23 (slot 47): atom span ≥ 190 × q_abs_ref_eff (= 9.5 × max_atom_abs × 2) nan_check_buf_ptrs/lens resized 12→24 entries (mapped-pinned host write at construction). populate_nan_check_meta extended with 4 new args for IQN Adam m/v ptr+len (Option<u64>/Option<usize> — null when IQN inactive). Grid: 12 → 24 blocks. Single launch covers all 24 backward-path diagnostic slots. Slot 31 (deferred) + slots 33-35 (inline elsewhere) + slots with null IQN entries no-op via the kernel's null-pointer guard. q_abs_ref_eff = max(isv[Q_ABS_REF=16], 1.0) — ε on multiplier per SP1 pearl; cold-start ISV (~0) gives q_abs_ref_eff = 1, so all thresholds floor at their ε-floor multiplier × 1. SP3 Mech 5 closes the diagnostic instrumentation loop — slots 36-47 fire when their threshold is exceeded, providing observability for SP3's other 4 mechanisms' effectiveness. Fail-safe: if SP3 fix doesn't fully resolve F1 NaN, slot 36-47 firing pattern guides the next iteration. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
ef429c25d8 |
feat(dqn): SP3 Mech 4 — comprehensive Adam EMA reset (audit + GpuAttention wire)
Audit of all Adam optimizer state in crates/ml/src/ against the existing reset_adam_state calls at the fold-boundary in fused_training.rs::reset_for_fold: - DQN main (line 971) — covers m_buf/v_buf, IQN-trunk, q_attn, sel, denoise, mamba2, ofi_embed, PopArt, Q-div EMA via GpuDqnTrainer::reset_adam_state. Also covers VSN params (in main params buffer slots) and aux next_bar/regime heads (slots 119-126). - IQN head (line 1011) - TLOB (line 1026) - IQL-high (line 1031) - IQL-low (line 1036) Missing wire identified and added: - GpuAttention (4-head feature attention on h_s2): owns a separate (attn_m, attn_v, attn_adam_step) tuple in gpu_attention.rs, called every training step via attn.adam_step at fused_training.rs:1962/1995, but was never zeroed at fold boundary. Same pathology as IQN/TLOB/IQL — fold N momentum oversizes fold N+1's first-epoch SDP + output projection updates, compounding through the trunk gradient. Implementation: - Add GpuAttention::reset_adam_state mirroring the gpu_iqn_head and gpu_tlob pattern (memset_zeros m/v + zero step counter + write 0 through pinned t_pinned for next adam_step launch). - Wire it under `if let Some(ref mut attn) = self.gpu_attention` in reset_for_fold, immediately after the IQL-low reset, with the same warn-on-failure / info-on-success log pattern. - Append SP3 Task B5 entry to docs/dqn-wire-up-audit.md documenting the audit + fix (Invariant 7 contract). Out-of-scope optimizers (documented in audit, not changed): - DecisionTransformer: scoped within DT pretrain block, dropped at end of pretrain (no fold leak possible). - GpuCuriosityTrainer: owned by GpuExperienceCollector, external to FusedTrainingCtx. Reset belongs at the collector layer if needed. - MetaQNetwork: host-side observability MLP for collapse prediction (does not feed LearningHealth or trunk gradients). - GpuMoeHead: forward-only, no Adam state of its own. Comprehensive coverage now per feedback_no_partial_refactor — every in-context Adam state buffer zeros at fold transition. SP3 Mech 4 closes the "Adam EMA accumulation drives weight pathology" pathway. |