Three review-driven fixes:
1. Slot 24/25 cite the actual production buffers d_value_logits_buf
(gpu_dqn_trainer.rs:3123) and d_adv_logits_buf (:3125) — the f32
atomicAdd buffers that launch_c51_grad writes — not the staging
buffers at :3127/:3129. Pointer expressions exist at the launch
site (:17707/:17708), making new accessors optional. Task 3 summary
and priority-list entries updated to match.
2. Explicit plan-supersession note inserted immediately above the
per-slot table: the audit's per-slot allocation supersedes the plan
Task 2 step 6 name table. The plan's allocation was placeholder;
the audit's is grounded in per-kernel inspection. Task 2 should use
the audit's names (d_value_logits_buf, d_adv_logits_buf, iqn_trunk_m,
iqn_d_h_s2_ptr, d_branch_logits_buf, etc.).
3. Summary suspicion-ranking row #3 reconciled with slot-28 dormancy:
iqn_backward_per_sample has no Rust caller (verified via
grep crates/ml/src/), so row #3 is re-pointed at the production
kernel iqn_quantile_huber_loss (iqn_dual_head_kernel.cu:1346-1413,
loaded at gpu_iqn_head.rs:2114). Same unsafe-write pattern (no
isfinite guard at line 1410's d_q_online[idx] = qw*d_huber/Q),
now attributed to the live path. apply_iqn_trunk_gradient at #1
stands — its reasoning (orchestrator consuming iqn_d_h_s2_ptr) is
unchanged by which specific kernel writes that buffer.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Quality-review fixes for the SP1 Phase A audit:
1. nan_flags_buf [i32;24] declaration: quad-cite (field 2659, alloc
11178, constructor 12563, read_nan_flags signature 14982). Task 2's
24->48 expansion must touch all four atomically per
feedback_no_partial_refactor; consumer site in
trainers/dqn/fused_training.rs and the two name-table sites in
trainers/dqn/trainer/training_loop.rs are still listed in Task 2's
plan body.
2. Per-slot Rust buffer-pointer expression added for slots 24-35.
Already-accessible (no new accessor): slots 26, 32, 33, 34, 35
(4 via self.ptrs, 1 via existing bn_d_concat_buf accessor). Need
new accessor on GpuDqnTrainer: slots 24 (d_value_logits), 25
(d_adv_logits), 30 (aux_dh_s2_nb_buf), and 29 (cql_d_value_logits,
only if un-deferred). Need new accessor on GpuIqnHead: slot 28
(d_branch_logits_buf — note: production IQN backward uses
iqn_quantile_huber_loss, NOT iqn_backward_per_sample which is
declared but never loaded). Optional: slot 27 (d_h_s2_buf_ptr) —
can be inlined inside apply_iqn_trunk_gradient instead. Need new
accessor on FusedDqnTraining: slot 31 (only if un-deferred).
Becomes input for Task 3.
3. Citation typo: c51_loss_kernel.cu line 274 reference removed —
line 274 is __syncthreads() in the projection-reduction warp loop;
the a_std=sqrtf reference belongs to c51_grad_kernel.cu:274.
Loss-kernel sqrtf sites are 779 and 804.
Verified by reading each cited line; SQLX_OFFLINE=true cargo check
--workspace passes (docs-only).
Read-only γ audit of every backward-path kernel writing to bw_d_h_s2 /
grad_buf / save_h_s2 accumulators. Per-kernel inventory against
unsafe-pattern checklist (sqrtf-neg, 1/0, logf-≤0, expf-large, EMA
variance, atomicAdd/saxpy NaN-propagation). Each section includes:
identified pattern + line citation, severity, proposed guard form,
ISV bound option, F0 risk, Phase B flag-slot allocation.
Cross-referenced session_2026-04-05_nan_investigation.md residual 8%
step-2 NaN in apply_iqn_trunk_gradient — never closed; ranked top
suspect for SP1.
Output drives SP1 Phase B (12 new flag slots in nan_flags_buf 24→48)
and Phase C surgical fix(es). Durable artifact for SP2 framework
codification + SP3 structural-fix scoping.
Implementation plan for SP1 (Sub-project 1 of 3) of the numerical
stability investigation. Follows the γ + β methodology from the spec
at docs/superpowers/specs/2026-04-29-numerical-stability-investigation-design.md.
8 tasks across 4 phases:
- Phase A (Task 1): γ read-only audit producing docs/dqn-backward-nan-audit.md
- Phase B (Tasks 2-5): always-landing β instrumentation expanding
nan_flags_buf 24→48 with 12 new backward-kernel NaN check slots
+ 12 reserved slots for future coverage
- Phase C (Task 6): surgical fix(es), content-driven by audit + smoke
topology, ISV-driven for any dynamic bound (mandatory)
- Phase D (Task 7): multi-fold L40S smoke validation against 7 pass
criteria (F0 ≥ 95% baseline, F1+F2 monotone improvement, zero
NaN-CLAMPED-TO-ZERO, all 48 NaN flag slots remain at zero)
- Closure (Task 8): audit doc closure, memory entry, SP2/SP3 handoff
Operating principles (mandatory per spec):
- No deferrals — anomalies discovered during investigation get fixed
within SP1, not punted to SP2/SP3
- Combined RELATED fixes ship as rich commits (per
feedback_no_partial_refactor)
- ISV-driven design for any dynamic bound
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
9 substantive issues addressed inline:
1. ISV-driven design elevated from 'if applicable' to MANDATORY for
all dynamic bounds in SP1 fixes. Numerical-stability ε bounds are
the only carve-out (Invariant 1). Hardcoded tuning constants for
dynamic ranges explicitly rejected.
2. F0 Sharpe regression criterion changed from absolute (≥55) to
ratio-based (≥95% of latest baseline; floor 53.08 currently).
Prevents iterative erosion across multiple fix commits.
3. 'F1 trending positive' replaced with concrete monotone-improvement
test: Best Sharpe at last epoch ≥ Best Sharpe at first epoch of
the same fold.
4. Pass criterion distinguishes NaN-CLAMPED-TO-ZERO (failure) from
'Genuine grad collapse' (legitimate observation, permitted) per
the existing infrastructure from commit d1808df14.
5. Multi-source NaN scenarios explicitly supported — γ + β may
identify multiple kernels; SP1 fixes ALL within the same cycle.
6. F0-safety paper-review gate added BEFORE smoke validation. Audit
doc carries 'F0 risk' (low/medium/high) per proposed fix; high-
risk fixes get math-on-paper inspection before consuming L40S.
7. Audit doc structure now requires 'ISV bound option' column —
forces ISV-first thinking at audit stage, not as afterthought.
8. Anti-patterns expanded: micro-clamping (per-op clamping that hides
upstream causes); combining unrelated fixes (anti-pattern of the
rich-commit principle); hardcoded constants for dynamic bounds.
9. 48-slot allocation explicitly justified (24 used + 12 new + 12
headroom) AND marked reviewable by SP2 if right-size differs.
All 5 design sections preserved structurally; revisions integrated
inline. Per brainstorming skill: spec self-review fixes applied
without re-review cycle.
R1: K's hardcoded shrink-and-perturb (m×0.1, v×0.01) at fold boundary
violated feedback_adaptive_not_tuned (untracked tunable knobs) AND
created a downstream pathology: tiny v_hat denominator → oversized
Adam updates 50+ steps post-reset → trunk param overshoot → save_h_s2
NaN at F1 ~step 1745 (smoke-test-bkdx5 diagnostic).
K was introduced (commit 4ef1d8ebb) BEFORE fold_warmup_factor existed
in the same commit's "K + adaptive warmup" pair. With warmup_factor
in place — ISV-driven, dampens lr+clip via lr_eff = lr_base ×
max(MIN_WARMUP_LR_FRAC, fold_warmup_factor) — K is redundant. Single
mechanism, ISV-driven, no hardcoded constants. Eliminating K leaves
m,v reset to 0 at fold boundary; warmup_factor handles cold-start.
P: expanded nan_flags_buf 16→24 with 5 GRN-stage checks
(linear_a_out, elu_out, linear_b_out, glu_sigmoid_out,
layernorm_var/out) for finer-grained source identification if R1
alone doesn't fix F1.
Predicted outcomes:
- If K's tiny-v_hat was the cause: F1 trains successfully (R1 alone)
- If different mechanism: new GRN-stage flags pinpoint which sub-
stage produces NaN, enabling layer-level fix
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Existing nan_flags_buf [16] covers 13 buffers with per-step NaN checks
inside the captured training graph. But read_nan_flags() was only
called when training guard's halt_nan fired — which checks pinned-
readback grad_norm. NaN-clamped-to-zero gradients reach the pinned
scalar as 0, not NaN, so halt_nan never fires for the explosion case.
The collapse path (halt_grad_collapse, grad < 0.01) was firing instead
without reading the flags. Add flag readback in that path when
gr.raw_grad_norm < 1e-6 (suspicious zero) — logs NaN-CLAMPED-TO-ZERO
with flagged buffer names. Genuine near-zero gradients get a separate
"no NaN flags set" log so we can distinguish the two cases.
This should pinpoint which kernel produces the F1 ep2 NaN first
(C51 KL projection? IQN aux? CQL? aux heads?). Diagnostic only —
keep after fix lands; correct gate for future regressions.
After 6+ intervention layers (A.1+A.2+A.3+F+H+K+warmup+N) F1 ep2 still
explodes with grad_norm=2.66T while F1 ep1 trains cleanly. None of our
defenses catch the explosion path. Need data to pinpoint WHICH signal
explodes first.
Per-step FOLD_EXPLOSION_DIAG in fold >= 1: when grad_norm > 1000 OR
jumps 5x from prior guard step, fire diagnostic warn with:
- loss decomposition: total / c51 / mse / iqn (pinned readback, no DtoH)
- Q range: q_min / q_mean / q_max via cold-path reduce
- atom positions: per_sample_support[0,d0] + [0,d2] (v_min, v_max, dz)
- prev grad norm + ratio for context
- trigger flag + post-trigger trajectory countdown
After trigger, continues emitting for 5 guard steps so the explosion
trajectory is captured (not just the first crossing).
Plus an unconditional FOLD_EXPLOSION_DIAG[F1_END_EP1] baseline log at
the end of fold 1 epoch 0 — healthy state immediately preceding the
explosion. Compare-and-contrast with per-step explosion frames pins
the runaway driver.
CQL and ensemble losses are not pinned-readback (transient device
buffers consumed inside the training graph) and are explicitly absent
from the decomposition. If none of {c51, iqn, mse} is the runaway
driver but total_loss still explodes, that implicates the unexposed
CQL/ens path — the absence is itself diagnostic information.
Implementation:
- Three diagnostic-only fields on DQNTrainer: current_fold,
last_logged_grad_norm, explosion_diag_steps_remaining. Reset
last_logged + remaining in reset_for_fold; current_fold
overwritten unconditionally at fold-loop entry.
- Three accessors on FusedTrainingCtx: explosion_diag_loss_components
(pinned), explosion_diag_atom_range (12-elem cold DtoH),
explosion_diag_q_range (reduce + 28B DtoH).
- Three accessors on GpuDqnTrainer: c51_loss_pinned_value,
mse_loss_pinned_value, stream_for_diag.
Diagnostic-only — remove once root cause identified. Not for prod.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Smoke smoke-test-xw4c6 showed F1 ep2 grad_norm=8.4B polluting the
adaptive_clip EMA -> next adaptive_clip threshold became huge ->
subsequent extreme grads passed unclipped -> NaN propagation -> grad
collapse -> fold 1 fails despite K+warmup fixing the first-step problem.
Add Winsorized clamp on raw_grad_norm before the adaptive_clip EMA
update: clamped_sample = min(raw, K * previous_adaptive_clip) where
K = 100 (numerical-stability bound, not a tuned constant — explicit
"single sample can grow EMA at most 100x in one step" semantics).
Companion observability log emits GRAD_CLIP_OUTLIER warning per clamp
event so we can see when this fires in subsequent smokes.
Fast/slow grad_norm EMAs (driving warmup factor) intentionally NOT
winsorized — they're a stability signal that SHOULD respond to
outliers, providing extra warmup damping when the system is unstable.
Predicted impact on Plan C smoke F1 ep2: 8.4B grad -> clamped to
~3000 for EMA -> next adaptive_clip caps at ~1000 -> subsequent F1
batches get bounded gradients passed to Adam -> no NaN propagation.
F1+F2 should now train through, exposing whether the underlying
Q-target inflation requires further structural work or whether
clipping alone suffices.
Composes with K+warmup (4ef1d8ebb), A.2, A.1/A.3, F+H — fourth layer
of the adaptive gradient-stability scaffolding (now five total:
A.2 target-sync drift dampening, K+warmup post-reset transient
dampening, N persistent clip-EMA outlier defense).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Closes the F1 ep1 catastrophic-overshoot gap exposed by smoke-test-s9h4h
(F0 succeeded with Best Sharpe = 36.03; F1 ep1 grad_norm = 355,009 — 5
orders of magnitude larger than F0 steady-state ~10 — leading to NaN
propagation and grad-clamp-to-zero early stop). Combined fix: (1) Adam
shrink-and-perturb at fold boundary (replace m=0/v=0 with m*=0.1, v*=0.01
to preserve direction while damping magnitude), and (2) a single
adaptive ISV signal driving BOTH lr_eff and clip_eff dampening over
the fold's first ~50 steps.
K — Adam shrink-and-perturb in `GpuDqnTrainer::reset_adam_state`:
- m *= 0.1, v *= 0.01 via existing `dqn_scale_f32_kernel` (loaded as
`scale_f32_ungraphed`); t_pinned still zeroed so bias correction
restarts. Architectural constants (preserve direction / lose magnitude
history) per `feedback_isv_for_adaptive_bounds.md` Invariant 1
carve-out — not tuned. Composes with existing param shrink-and-perturb
(`alpha=0.8`) in `FusedTrainingCtx::reset_for_fold`. Root cause for
F1 overshoot: m=0,v=0 → first Adam step ≈ lr × g / ε → 6 OoM
amplification.
New CUDA kernel `fold_warmup_factor_kernel.cu`:
- Single-block single-thread cold-path producer mirroring
`q_drift_rate_ema_kernel.cu` / `moe_lambda_eff_kernel.cu` shape.
- Reads two grad-norm EMAs (fast α=0.1, slow α=0.001) plus host-passed
step counter; writes ISV[FOLD_WARMUP_FACTOR_INDEX=130] = clamp(fast/slow, 0, 1).
- Bootstrap branches (steps_observed < 200, slow EMA < 1e-6) emit
factor=1.0 (no damping during cold-start). No atomicAdd; no DtoH.
New ISV slot Q_DRIFT_RATE_INDEX → FOLD_WARMUP_FACTOR_INDEX = 130:
- ISV_TOTAL_DIM 130 → 131; layout fingerprint shifts (checkpoint-
incompatible per `feedback_no_legacy_aliases.md`, expected for a
real architecture change).
- FoldReset entries: `isv_fold_warmup_factor` → 0.0 and companion
`isv_grad_norm_fast_ema` → 0.0 (lockstep reset per
`feedback_no_partial_refactor.md`); slow EMA persists across folds
as the cross-fold steady-state baseline.
- Two new mapped-pinned scalars on GpuDqnTrainer (grad_norm_fast_ema_pinned,
grad_norm_slow_ema_pinned) fed by `update_adaptive_clip` from the
same `gr.raw_grad_norm` observation source as the existing adaptive
clip EMA.
Two consumers, both monotone (only dampen, never excite):
- lr_eff = cosine_effective_lr × max(MIN_WARMUP_LR_FRAC=0.05, factor)
via `set_lr` per-step. New `cosine_effective_lr_base` field
on DQNTrainer composes the cosine schedule's per-epoch
baseline with the warmup factor's per-step damping (rather
than overriding the cosine schedule).
- clip_eff = clip_base × (MIN_CLIP_FRAC=0.1 + 0.9 × factor) via new
`set_active_clip` setter on FusedTrainingCtx + GpuDqnTrainer.
Composes with the EMA-derived `clip_base = grad_norm_ema × 2`
that `update_adaptive_clip` just wrote to the pinned slot.
Numerical-stability bounds 0.05 / 0.1 are Invariant 1
carve-outs.
Steady-state behaviour unchanged: factor=1 → lr_eff=lr_base,
clip_eff=clip_base. Fold-boundary behaviour: factor starts at 0 →
lr_eff = 0.05 × lr_base, clip_eff ≈ 0.1 × clip_base; rises to 1 over
~50 steps as the fast EMA catches up to the slow steady-state EMA.
Predicted impact on Plan C smoke F1: 355,009-magnitude transient grad
clipped to ~clip_base × 0.1 ≈ 1.0 (vs 10), Adam state shrunk instead
of zeroed → first step update bounded; grad recovers normally over
~50 steps. Companion to A.1 (prev_epoch_q_mean reset), A.2 (adaptive
Polyak-tau), A.3 (gradient_collapse_counter reset), F+H (kill-criterion
robustness) — completes the fold-boundary state-reset family.
Per `pearl_adaptive_moe_lambda.md` (kernel + ISV slot + bootstrap +
reset + observability template), `pearl_cold_path_no_exception_to_gpu_drives.md`
(GPU-stays-on-GPU even at cold-path cadence),
`pearl_blend_formulas_must_have_permanent_floor.md` (lr/clip floors
are permanent minimums), `feedback_adaptive_not_tuned.md` (lr+clip
ISV-driven), `feedback_isv_for_adaptive_bounds.md` (factor IS the
bound; consumers compose at runtime), `feedback_no_atomicadd.md`
(single-thread reduce), `feedback_cudarc_f64_f32_abi.md` (slot index
passed as i32), `feedback_no_partial_refactor.md` (kernel + slot +
reset + producer + 2 consumers all land together),
`feedback_no_quickfixes.md` (replaces brittle full-reset with
adaptive damping; not threshold relaxation).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Smoke smoke-test-vh9bj revealed: after F+H Q-drift kill fix, fold 0
trains successfully (5 epochs, Best Sharpe 36.03), but fold 1 fails
at epoch 1 with "Gradient collapse detected for 5 consecutive epochs"
despite epoch 1 grad_norm=296,417 (healthy). Root cause: the per-step
gradient_collapse_counter on DQN was already near patience (5) from
fold 0's late-epoch near-zero-grad steps, and DQNTrainer::reset_for_fold
didn't clear it. Plus, training_steps accumulating across folds made
the `past_warmup` gate always-true in fold 1+, removing the warmup
grace period for data distribution shifts.
Add DQN::reset_for_fold zeroing both. Wire from DQNTrainer::reset_for_fold
alongside the A.1 prev_epoch_q_mean reset. Pure additive — same
fold-boundary state-reset gap pattern as A.1, A.2, F.
Predicted impact: fold 1+ now starts with counter=0 and warmup window
restored; gradient collapse check has its full per-fold grace period.
Current criterion uses |q_mean| / |prev_q_mean| which explodes near
zero crossings (legitimate cold-start has |prev_q_mean| ≈ 0.01-0.1,
making any non-tiny current trigger the ratio threshold).
smoke-test-n9xzr fired with prev=−0.0795, curr=0.7647, ratio=9.62×
despite this being natural cold-start growth, not geometric runaway.
Replace with rolling 5-epoch window of q_means. Compute median +
MAD (Median Absolute Deviation, a 50%-breakdown estimator robust
to single outliers); kill condition becomes
|q_mean − median(window)| > 4.0 × max(MAD(window), 0.01)
AND the existing adaptive floor.
Constants are declared `const` near the kill block:
- Q_DRIFT_WINDOW_SIZE=5 (matches smoke fold length, gives MAD a
meaningful estimator without averaging across regime shifts)
- Q_DRIFT_WARMUP_SAMPLES=3 (skip until ≥ 3 priors — smaller window
degenerates to half-range MAD that trips on monotonic
trajectories)
- Q_DRIFT_DEVIATION_THRESHOLD=4.0 (4 MADs ≈ 2.7σ Gaussian-
equivalent; clear outlier without firing on every legitimate
dip; literal MAD count rather than σ because q_mean is non-
Gaussian during cold-start)
- Q_DRIFT_MIN_DEVIATION=0.01 (numerical-stability floor when
window is constant)
Window resets at fold boundary alongside prev_epoch_q_mean /
adaptive_tau (A.1 pattern — cross-fold q-stats are independent
training runs, mixing them would inflate MAD or shift the
median).
Predicted impact on smoke-test-n9xzr:
- Plan C smoke fold 0 ep2: window has only 2 priors, warmup gate
not yet satisfied → kill stays silent (was firing on ratio=9.62×)
- Genuine geometric runaway (q_mean → 62.5 from baseline 0.5 over
4 epochs): ep3 deviation = |62.5 − 2.5|/MAD=2.0 = 30 > 4 AND
|62.5| > floor=3×12=36 → kill fires correctly
Both conditions ANDed (floor + deviation), preserving the
production-safety semantics of the original criterion. The floor
check is unchanged; only the divergence detector is replaced.
Per feedback_no_quickfixes.md: this is a principled robust-
statistics replacement, not a threshold relaxation.
Per feedback_no_partial_refactor.md: window field, constants,
criterion site, and fold-boundary reset land in lockstep — the
kill criterion's contract is internally consistent across
trainer/mod.rs, constructor.rs, and training_loop.rs.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The Q-drift kill criterion's adaptive floor formula is
max(0.5, 3 × max(ISV[Q_ABS_REF=16], ISV[Q_DIR_ABS_REF=21]))
but smoke-test-n9xzr showed both ISVs stuck at bootstrap 0.05 even
when q_mean reached 0.76 — the producers (q_mag_bin_means_reduce
and q_dir_bin_means_reduce) only fired inside reduce_current_q_stats
at epoch boundary, while the per-step captured update_isv_signals
consumed the resulting scratch buffers each training step.
Concrete failure mode at α=0.05 EMA with raw ≈ 1.0:
- Epoch 0: scratch is zero-initialized, all per-step EMA pulls
toward zero, ISV[16,21] stay 0.0
- Epoch 0 boundary: reduce_current_q_stats fires, scratch becomes
q_abs_ref ≈ 1.0, single update_isv_signals writes ISV[16] = 0.05
- Epoch 1+: with stale scratch held constant between boundaries,
per-step EMA over hundreds of steps would saturate — but smoke
step counts are small enough (small batch=64, buffer=256) that
the slot stays near 0.05
- kill_floor degenerates to max(0.5, 3 × 0.05) = 0.5 forever, so
the criterion fires on legitimate cold-start growth
Fix: launch q_mag_bin_means_reduce + q_dir_bin_means_reduce in
both fused_training paths (captured adam_update_child at ~line 2228
AND ungraphed step-0 fallback at ~line 1465) immediately before
update_isv_signals. Per-step graph replays now read q_out_buf that
forward_child populated this same step, write fresh scratch, and
the captured update_isv_signals consumes it on the same stream.
Per pearl_cold_path_no_exception_to_gpu_drives.md: cold-path EMAs
fire alongside their consumers. Per feedback_no_partial_refactor.md:
graphed and ungraphed paths migrate together — same producer-
consumer contract.
After this fix, the floor will adapt: ISV[16,21] saturate to the
policy's actual q_abs_ref scale within ~60 steps at α=0.05, so by
epoch 2 the floor becomes 3 × ~0.5 = ~1.5, no longer firing on
legitimate cold-start growth.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
C.4 timing bonus (experience_kernels.cu) and D.4b regime penalty
multiplied by the trade-cumulative |reward| / |final_pnl| as an
unbounded multiplicand. Across trades this compounds — bonus values
inflate as Q values inflate, then bonus rewards inflate Q values
further (Bellman bootstraps off shaped reward).
Per pearl_one_unbounded_signal_per_reward + feedback_isv_for_adaptive_bounds:
replace |reward|/|final_pnl| with ISV[Q_DIR_ABS_REF_INDEX]-bounded
variant. The unbounded multiplicand becomes the direction-branch
Q-scale EMA (already-tracked, gradient-decoupled), not the
trade-cumulative shaping output. Cap is adaptive (matches Q magnitude
as it evolves) and breaks the multi-trade compounding loop.
Discovered during Plan C Phase 2 smoke diagnosis (researcher report
2026-04-29 a25f669e9df953174). a52d99613 baseline also hits
bonus=235 — pre-existing structural pathology, not Plan-C-specific.
Other shaping sites (D.4a persist, B.2 novelty, D.4c stable) already
have all factors bounded — no fix needed.
Surgical change set:
- experience_env_step gains final `int q_dir_abs_ref_idx` param.
- C.4 (line 2521): pnl_unit = |final_pnl|/max(ISV[21], eps);
pnl_capped = min(pnl_unit, 1) * ISV[21].
- D.4b (line 2581): same pattern, reward_unit/reward_capped.
- gpu_experience_collector.rs:3800 wires Q_DIR_ABS_REF_INDEX as i32.
No new ISV slot, no new producer kernel, no layout-fingerprint shift —
ISV[21] already produced by q_stats_kernel.cu since Plan 1.
Predicted impact:
- bonus EMA drops O(100-256) -> O(0.05-2.0)
- rc[5] -> Bellman-target optimism loop broken
- Plan C smoke: Q-drift kill at F0 ep2 likely no longer fires
- a52d99613 baseline: same effect; validates pre-existing fix
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Wires the q_drift_rate_ema kernel + ISV slot Q_DRIFT_RATE_INDEX=129
(landed in fc8dbb0a8) into the production training path:
- StateResetRegistry FoldReset entry isv_q_drift_rate -> 0.0
(companion to A.1's prev_epoch_q_mean reset; both ensure no
cross-fold leakage of drift state).
- reset_named_state dispatch arm in training_loop.rs.
- Per-epoch launch_q_drift_rate_ema call after the Q-drift kill
check, gated on the same `prev_epoch_q_mean.abs() > 1e-6`
cold-start guard. Both q_mean_curr and q_mean_prev are in scope
BEFORE the `prev_epoch_q_mean = q_mean` update, so the producer
sees the correct delta.
- tau_update_kernel.cu multiplies the cosine-scheduled,
health-coupled tau_eff by `1 / (1 + clip(ISV[129], 0, 4))` so
tau ranges [tau_base/5, tau_base] — monotone dampening under
drift; healthy runs (drift ≈ 0) unaffected.
- reset_for_fold comment in trainer/mod.rs notes the registry
entry handles the new ISV slot.
Predicted impact: Plan C fold 0 ep2 with q_mean(t)=0.82,
q_mean(t-1)=-0.018, ISV[16]+ISV[21]≈0.1 -> drift_rate ≈ 8.4 ->
clipped to 4 -> tau_eff = tau_base × 0.2 (5× slower target sync,
dampens Q-target optimism through the inflation spike).
a52d99613 fold 0 with q_mean staying ~0.21 -> drift_rate ≈ 0 ->
tau_eff = tau_base (unchanged).
Per pearl_adaptive_moe_lambda.md — canonical "EMA-tracked
diagnostic drives a controller" pattern. Per
pearl_cold_path_no_exception_to_gpu_drives.md — cold-path scalar
arithmetic stays on GPU. Per feedback_no_partial_refactor.md —
consumer (tau_update + state_reset + producer launch) all migrate
together in this commit.
Layout fingerprint already shifted by fc8dbb0a8 (slot + ISV_TOTAL_DIM
bump); no additional fingerprint shift in this commit.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Single-block single-thread ISV producer mirrors h_s2_rms_ema /
moe_lambda_eff pattern. Computes per-epoch
drift_rate = |q_mean(t) - q_mean(t-1)| /
max(|q_mean(t-1)|, ISV[Q_ABS_REF] + ISV[Q_DIR_ABS_REF], 1e-6)
clipped to [0, 4] and writes to ISV[Q_DRIFT_RATE_INDEX=129].
Includes the full ISV-contract shift required for the kernel to load:
- ISV_TOTAL_DIM 129 -> 130
- Q_DRIFT_RATE_INDEX = 129 (tail-appended after MOE_LAMBDA_EFF=128)
- Cold-start ISV[129] = 0.0 in constructor (no-op dampening factor)
- layout_fingerprint_seed entry Q_DRIFT_RATE=129 + ISV_TOTAL_DIM=130
- Cubin static, kernel field, kernel load, struct assignment
Per feedback_no_partial_refactor.md the ISV slot + dim + fingerprint
all migrate together (the wrapper references Q_DRIFT_RATE_INDEX so they
cannot be split). Tau consumer + state reset registry + per-epoch
producer launch land in the next commit.
Audit doc dqn-wire-up-audit.md updated with the kernel + ISV slot
description per Invariant 7.
No callers in this commit; layout fingerprint shifts so existing
checkpoints will fail-fast at load per feedback_no_legacy_aliases.md
(expected for a real architecture change).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Plan C Phase 2 T9. Verifies the production kernel's eval-mode argmax
exactly matches a Rust ground-truth E[Q] argmax on a 256-sample batch
with realistic per-direction non-degenerate C51 distributions, and
confirms Thompson explores Long+Short ≥ 40% in training mode.
The plan-prescribed approach (reuse Phase 0 Test 0.F's converged
checkpoint loader) was deferred — Test 0.F itself already exercises
the safetensors load + branching forward path. Replacement strategy
from the dispatch brief: synthetic batch via direct buffer write.
Setup:
- 256 samples, 21 atoms, per-(sample, direction, atom) C51 logits
drawn from a deterministic hash → uniform [-1, 1] (post-Xavier-init
scale of fresh C51 head outputs)
- Per-sample adaptive support [v_min ∈ [-1.0, -0.2], v_max ∈ [0.2, 1.0]]
matching the layout produced by `update_per_sample_support`
- Uniform q_values (mag/ord/urg fall through Boltzmann uniformly)
Rust ground-truth: softmax(b_logits[i, d]) · atom_vals[i, d] computed
with the numerically-stable subtract-max softmax matching the kernel's
softmax_c51_inline.
Assertions:
- Eval mode: per-sample dir_idx EXACTLY matches ground-truth argmax E[Q]
- Train mode: count(d ∈ {Long, Short}) / batch > 0.40
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Plan C Phase 2 T8. The plan-prescribed pre-T2 snapshot approach was
impossible (T2 had already landed); replacement strategy (ii) from the
dispatch brief — behavioral parity vs an analytical Boltzmann reference
computed in Rust — is used.
Setup forces direction = Long (d=2) deterministically via peaked C51
logits (Long peaked at v=+0.8 atom; other directions at v=-0.5), so
the kernel's Hold/Flat → mag_idx=0 short-circuit doesn't mask the
magnitude branch's Boltzmann sampling. q_values are crafted with each
branch (mag/ord/urg) peaked at a single bin with magnitude 1.0:
Mag Q = [0.0, 0.0, 1.0] peak at Full (mag=2)
Order Q = [1.0, 0.0, 0.0] peak at Market (ord=0)
Urgency Q = [0.0, 1.0, 0.0] peak at urg=1
With q_range=1.0 in all three branches, tau collapses to 1.0 and the
analytical Boltzmann probabilities are:
P(best) = 1/(1 + 2/e) ≈ 0.5767
P(other) = 1/e/(1 + 2/e) ≈ 0.2117
Tolerance: at batch=8192 the 1-σ Bernoulli noise is ~0.0055 for p≈0.58;
±5% absolute tolerance covers ~9σ. Algorithmic divergence (e.g. an
inadvertent strict-argmax substitution) would shift P(best) to 1.0 —
trivially detected by the ±5% tolerance.
Assertions:
- dir_idx == Long for every sample (eval argmax E[Q] over peaked C51)
- mag/ord/urg histograms each within ±5% of the Boltzmann reference
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Plan C Phase 2 T7. #[ignore]-gated GPU test that runs BOTH the
production experience_action_select kernel AND the standalone
direction_thompson_v2_test (added in commit 5de5e546a) on identical
Phase 0 Test 0.D failure-mode inputs and asserts the resulting
direction histograms are statistically indistinguishable.
Path (a) — bit-identical comparison via a shared pre-computed uniform
array — would require editing the standalone v2 kernel's signature to
accept a `float* uniforms` instead of generating its own LCG draws.
Out of scope for the 1-2-hour dispatch. Path (b) — KS-style histogram
comparison — is used.
Setup:
- Single per-direction logits tile (Phase 0 Test 0.D failure mode)
- Production: tile replicated across batch=100 000 samples; Philox
keyed on (i, timestep=0, ctr) per thread
- Standalone v2: same tile fed to single-tile launcher with n_seeds=100 000;
LCG keyed on (base_seed=0 + seed_idx) per thread
Assertion:
- KS distance over the {Short, Hold, Long, Flat} histograms ≤ 0.02
(n=100k empirical-CDF noise floor for matching distributions is
~4·sqrt(2/n) ≈ 0.018; algorithm divergences would shift mass by
O(10%) → KS ≈ O(0.1), easily detected)
- Sanity: both histograms have P(Long+Short) ≥ 0.20 (rules out
coincidental shared-mode collapse with KS ≈ 0)
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
reset_for_fold cleared q_value_history (the q_mean source) but missed
the kill criterion's prev_epoch_q_mean baseline AND the adaptive_tau
modulation state. So the kill ratio at fold-N epoch 0 was computed
against fold-(N-1) epoch 5's q_mean — a stale cross-fold baseline.
Discovered while diagnosing why Plan C Phase 2 Thompson smoke triggers
Q-drift kill at fold 0 epoch 2 (researcher report 2026-04-29).
Pairs with the existing q_value_history clear; matches the
project_fold_boundary_q_drift_resolved.md pattern (kill criterion's
companion resets were partially missed).
Independent of Plan C — this is a bug in the kill criterion's
fold-boundary contract regardless of which exploration mechanism is
active.
Plan C Phase 2 T6. Adds an #[ignore]-gated GPU test that exercises
the production experience_action_select kernel directly through a
minimum-viable inline cudarc fixture (ProdActionSelectFixture). The
Plan B audit fixture (GpuExperienceCollector::new_for_test) was never
built, so each test in this Phase 2 batch builds its own buffers,
launches the kernel, and reads back action picks for assertions.
Setup mirrors Phase 0 Test 0.D failure-mode:
- Flat/Hold C51 logits = peaked at v=0 atom (δ(v=0) shape)
- Long/Short C51 logits = log-Gaussian centred at v=-0.001, σ=0.05
- Linear adaptive support [v_min=-0.5, v_max=+0.5, delta_z=0.05]
- Uniform q_values (mag/ord/urg fall through Boltzmann uniformly)
Assertions:
- Eval mode (eps=0): 10 launches at distinct timesteps produce
IDENTICAL per-sample dir_idx (deterministic argmax E[Q]).
Every sample picks Hold or Flat (E[Q] argmax bias reproduces).
- Training mode (eps_start=0.5): >=30% of 10 000 samples pick
Long or Short (Thompson explores directional alternatives).
Shared helpers (PROD_B0..3, decode_dir/mag/ord/urg, fill_peaked_logits,
fill_linear_support, ProdActionSelectFixture) added at the top of the
Phase 2 section to be reused by Tests 2.B-2.D.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The prior `test_eval_action_select_boltzmann_bounded` asserted
Boltzmann theory (P(best)≈0.366 at tau=q_range) on the direction
branch. Plan C T2 (commit 52a2663a2) replaced the direction-branch
Boltzmann path with single-distribution C51 Thompson + argmax-E[Q]
eval; the T2 amendment (5de5e546a) added 4 args to the kernel ABI
(b_logits_dir, per_sample_support, atom_positions, n_atoms). The old
test was launching with an outdated arg count and asserting a
distribution shape that no longer applies.
Migration (per feedback_no_partial_refactor):
- Rename to `test_eval_action_select_eval_argmax_picks_best`
- Build per-direction C51 logits PEAKED at distinct E[Q] values:
Short=-0.5, Hold=-0.1, Long=+0.5 (best), Flat=+0.1
- Linear adaptive support [v_min=-1, v_max=+1, delta_z=0.1] uniform
across all (sample, direction) pairs
- Run kernel in eval mode (eps_start=eps_end=0): assert P(Long)>=0.99
i.e. the kernel deterministically picks argmax(E[Q]) per sample
This validates the T2 amended ABI runs end-to-end and exercises the
production kernel's argmax-E[Q] eval branch with a controlled
distinct-E[Q] setup. Companion to Phase 2 Tests 2.A-2.D in
distributional_q_tests.rs which add training-mode assertions and
production-vs-standalone parity checks.
Verification: `SQLX_OFFLINE=true cargo check -p ml --lib --tests` clean.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Documents the project-wide invariant: any new distributional Q-head
(atoms, quantiles, ensembles) must expose sampleable interface and
wire into Thompson (training) + argmax (eval) action selection.
Pre-commit hook (Invariant 7) will enforce.
Architecture reflects Plan C Phase 2 amendment (commit 5de5e546a):
single-distribution C51 Thompson is the production pattern; joint
C51+IQN forward-looking under the same contract.
Counterpart to existing val_active_frac. Reports the fraction of
TRAINING rollout actions that were Long or Short. Required for L3
verification gate of Plan D (Phase 3) which checks Thompson is
generating directional exploration during training.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Threads b_logits_dir / per_sample_support / atom_positions / n_atoms
through experience_action_select kernel launch in
gpu_backtest_evaluator.rs's chunked val pipeline (line ~1722, inside
submit_dqn_step_loop_cublas).
The evaluator sources Q-values via the QValueProvider trait
(delegates forward to the trainer's CUDA-Graphed cuBLAS), so this
change extends the trait surface rather than duplicating the
forward:
- New trait method compute_q_and_b_logits_to(states_ptr, batch,
q_out_ptr, b_logits_out_ptr) — DtoD-copies trainer's
on_b_logits_buf into caller's chunked buffer per sub-batch
iteration alongside the existing q_out copy.
- New trait accessors per_sample_support_ptr / atom_positions_ptr /
num_atoms / total_branch_atoms (stable trainer-owned buffers).
- FusedTrainingCtx implements all of the above; trainer gains
on_b_logits_buf_ptr / atom_positions_buf_ptr /
per_sample_support_ptr_get pub accessors.
- Evaluator gains chunked_b_logits_buf field (sized
[n_windows * CHUNK_SIZE, total_actions * num_atoms] + 32*3 tail
safety), allocated in ensure_action_select_ready.
Phase 6's last-step plan_params forward keeps using the original
compute_q_values_to (b_logits not consumed there).
Plan C Task 4 — evaluator companion to T3 collector wire-up. After
this commit, all production callers of experience_action_select use
the amended kernel ABI (T2 amendment 5de5e546a) end-to-end
(feedback_no_partial_refactor).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Threads b_logits_dir (= exp_b_logits, first N*b0*NA region for direction
branch), per_sample_support, atom_positions (NULL = linear), and n_atoms
through experience_action_select kernel launch in gpu_experience_collector.
After amendment commit 5de5e546a, the kernel uses single-distribution C51
Thompson with production's adaptive per-direction support — no IQN
inference required at rollout.
Plan C Task 3 (corrected file path: collector, not trainer per
gpu_dqn_trainer.rs/gpu_experience_collector.rs distinction).
Plan C as authored assumed:
- c51_probs_dir = post-softmax probabilities (didn't exist on collector;
only raw exp_b_logits is materialised)
- atom_values = single global linear support (production uses per-sample
per-direction adaptive [v_min, v_max, delta_z] + optional atom_positions)
- iqn_quantiles_dir = rollout-side IQN inference (doesn't exist; IQN is
training-only)
Amended kernel signature uses the production architecture:
- b_logits_dir [N, b0_size, n_atoms] — raw direction-branch logits
- per_sample_support [N, b0_size, 3] — adaptive per-direction support
- atom_positions [b0_size, n_atoms] — non-linear positions (NULL = linear)
- n_atoms
Direction-branch Thompson is now single-distribution over C51 (with adaptive
support), not joint C51+IQN. Single-distribution still provides the
principled posterior sample that fixes the UCB selector/target asymmetry —
the goal of Plan C is preserved, just sourced from the existing rollout-time
distribution instead of a non-existent IQN inference path.
New device-inline helpers: softmax_c51_inline, compute_atom_values_inline.
Plan + audit doc updated. Phase 0 standalone test kernel gained two new
entry points (direction_thompson_v2_test, argmax_eq_v2_test) matching the
amended production API; original Plan A entry points retained for tests
0.B-0.F. Tasks 3+4 (buffer wiring) unblocked — collector/evaluator already
have per_sample_support_buf and exp_b_logits in production.
Plan C T2 commit 52a2663a2 migrated out_conviction to read from
e_dir[] (joint E[C51 + IQN]) per spec "Conviction stays E[Q]-based"
but left the out_q_gaps writer reading raw q_b0[]. Both signals feed
env_step's Kelly-adjacent position sizing; mixing q_b0-scale gaps
with e_dir-scale conviction inside the same downstream consumer
violates feedback_no_partial_refactor.
Migrate out_q_gaps to read from e_dir[] (already hoisted to outer
scope by T2). Contrarian-mode sign flip preserved symmetrically.
Comment at ~line 1289 updated to reflect joint E[Q] source.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Replaces eps-greedy + Boltzmann on direction with:
- Training: argmax of (sample_C51 + sample_IQN) per direction (Thompson)
- Eval: argmax of (E_C51 + E_IQN) per direction (no exploration)
Conviction calculation now reads e_dir[] (joint E[Q] computed inline)
preserving the spec's requirement that conviction uses raw E[Q] not
samples (avoids Kelly cap jitter).
Magnitude/order/urgency branches: unchanged.
Per-function purpose comments + parameter-contract notes were present
in the Phase 0 reference (thompson_test_kernel.cu) but dropped during
the port to experience_kernels.cu. Restoring them addresses code
review feedback. No functional change.
The original PAUSED state was motivated by measurement-artefact hunt
exit. The bug-hunt cycle is complete (commits a86fba2b1, b8788511c).
A new pathology has surfaced: ff00af68a's UCB count bonus activation
causes selector/target asymmetry → F0 Q-drift kill at epoch 2 →
F1+F2 cascade. Verified by paired DIAG smokes (smoke-test-qlz7t fail,
smoke-test-wmsht pass).
Thompson sampling on C51+IQN distributions eliminates the asymmetry
by construction (sample from learned distribution; no augment-then-
argmax step). Net code-surface decrease — replaces eps-greedy +
Boltzmann + UCB with one principled mechanism.
Plan C Phase 2 execution begins on branch plan-c-phase-2-thompson.
Phase 2A of the HEALTH_DIAG GPU port. Lands the simplest member of the
kernel family (single-block single-thread ISV scalar copies) plus the
orchestrator scaffolding (cubin loader, mapped-pinned snapshot owner,
launcher) that subsequent phases will reuse.
Why simplest first: 5 kernels each with different masking semantics
landing in one commit risks silent numerical mismatches that downstream
parsers (aggregate-multi-seed-metrics.py, smoke summarisers) would
mis-parse. Per feedback_no_quickfixes.md, kernel-by-kernel with
parallel-shadow validation. Phase 4 deletes CPU paths in one commit
once all kernels are bit-identical.
What this kernel does: copies 22 ISV signal-bus slots into the matching
HealthDiagSnapshot fields — reward-component EMAs (slots 63-68), VSN
attention focus EMAs (87, 88), target-drift EMAs (92, 93), aux-head
loss EMAs (113, 114, 117), MoE expert utilisation (118-126), gate
entropy (126), and λ_eff (128). 22 of 147 snapshot words populated
(reward_split[6] + noisy_vsn[2] + noisy_drift[2] + aux[3] + aux_moe[10]).
Producer-only — no consumer reads health_diag.snapshot() yet; Phase 3
wires the call site, Phase 4 makes the snapshot the sole emit source.
Architecture rules upheld:
- No HtoD / DtoH (writes through mapped-pinned device pointer).
- No atomicAdd (single-thread kernel; family-wide rule for 2B-2E too).
- No DtoD memcpy on the diag path (kernel pointer-load + store).
- Kernel WORD-index table inline; static_assert pins WORD_TOTAL=147 in
lockstep with snapshot_size_is_stable test.
- ISV indices passed as kernel args (not #define'd) so the kernel
never embeds the numbering — single source of truth lives in
gpu_dqn_trainer.rs constants.
Changes:
- new crates/ml/src/cuda_pipeline/health_diag_kernel.cu (311 lines).
- new crates/ml/src/cuda_pipeline/gpu_health_diag.rs (190 lines)
holds GpuHealthDiag orchestrator (cubin module + isv_mirror_kernel
handle + MappedHealthDiagSnapshot).
- mod.rs: pub mod gpu_health_diag.
- build.rs: registers health_diag_kernel.cu in kernels_with_common.
- gpu_dqn_trainer.rs:
- imports super::gpu_health_diag::GpuHealthDiag;
- adds health_diag: Option<GpuHealthDiag> field next to aux_heads_fwd;
- constructor calls GpuHealthDiag::new() after the cubin loads;
- adds pub fn launch_health_diag_isv_mirror() (uses compile-time
ISV index constants, single source of truth);
- adds pub fn health_diag_snapshot() accessor.
- docs/health_diag_inventory.md: Phase 2A entry + Phase 4 deletion
list pinned per Invariant 7.
- docs/dqn-wire-up-audit.md: new row for health_diag_kernel.cu under
"CUDA Kernels (.cu files)" — classified Pending wire-up (Phase 3).
Verification:
- cargo check -p ml: 0 new warnings (12 lib warnings pre-existing,
none touch health_diag/gpu_health_diag).
- cargo test -p ml --no-run: clean.
- cargo test -p ml --lib cuda_pipeline::health_diag::tests: 3/3 pass
(snapshot_size_is_stable, default_is_zeroed, alignment_is_4_bytes
— the WORD_TOTAL=147 static_assert in the kernel matches the host
size_of test bit-for-bit).
Scope discipline (per prompt's "if kernel #3 takes too long, stop"):
This commit lands the orchestrator scaffolding + simplest kernel only.
Phases 2B-2E (eval-histogram, q-mag-reduce, per-sample-reduce,
finalise) intentionally deferred — masking semantics for q_mag_reduce
in particular need careful per-CPU-getter inspection before kernel-side
implementation. Producer-only by design; no production caller of
launch_health_diag_isv_mirror in this commit (matches AuxHeadsForwardOps
landing pattern from Plan 4 Task 6 Commit A).
Phase 0 (commit 333ea7184): exhaustive inventory of every numeric value
emitted across the 7 HEALTH_DIAG sites in training_loop.rs and metrics.rs.
Classified each as GPU-already / CPU-bound / Mixed / Host-state /
ISV-already. Aggregate cost: ~50-150 MB/epoch DtoH + ~70 sec/epoch CPU
compute (the original observation that motivated the port). Inventory
lives at docs/health_diag_inventory.md.
Phase 1 (commit 60fd7de96): foundation only — no kernel, no producer
launch, no CPU code deletion yet.
* HealthDiagSnapshot #[repr(C)] POD with 147 fields (all f32 or u32)
covering every numeric value in the HEALTH_DIAG line family. Field
order matches the existing log format strings (downstream parsers
like aggregate-multi-seed-metrics.py depend on this).
* Controller fire-bools promoted from u8 to u32 to avoid #[repr(C)]
padding mismatch when followed by f32 fields. ~24-byte cost.
* MappedHealthDiagSnapshot wrapper around cuMemHostAlloc(DEVICEMAP|
PORTABLE) + cuMemHostGetDevicePointer_v2, mirroring the
MappedF32Buffer pattern.
* Default impl via MaybeUninit::zeroed() (allocator-free, valid
bit-pattern).
* 3 unit tests pinning size (588 bytes), alignment (4 bytes),
zeroed-default — all passing on local CPU host.
Phases 2-5 deferred (kernel family + wiring + CPU emit rewrite + old
code deletion). The 5 kernels each read different device buffers with
different masking semantics — getting any one wrong would silently
produce numerically-different HEALTH_DIAG output that downstream
parsers would mis-parse. Per the prompt's 'don't wedge yourself'
guidance and feedback_no_quickfixes.md, the safer move was to land
the typed wrapper cleanly so subsequent commits can iterate kernel
by kernel against the actual buffer-by-buffer reduction.
nsys profile of multi_fold_convergence on L40S identified compute_expected_q
as the #3 GPU consumer at 12.9% (207ms / 1382 calls — ~150 µs per call,
compute-bound at typical batch=8192 × num_atoms=51). The kernel computed each
per-action softmax over atoms in 3 passes (max → sum_exp → normalize+expected
+var+entropy+util), re-reading v_row[z]+adv_a[z] 3 times per atom for each of
the 4 × ~3.25 = 13 actions per sample.
Replaced with online softmax (running-max + running-sum) so expected_q,
sum_z_sq, and entropy all accumulate in a single forward pass over atoms.
Atom utilisation still needs a 2nd pass because it depends on the final
S = sum exp(logit-M) being known to compute prob_z = exp(logit_z-M)/S for
the threshold test.
Inner loop reduction: 3 → 2 passes, ~33% fewer global loads of branch
advantage logits (which dominate because v_row[z] is reused across 13
(action, branch) pairs and may stay in L1, while adv_a[z] flips per
action and is pure global).
Online accumulation pattern (Page-Olshen):
M, S, TZ, TZ², TLM = -inf, 0, 0, 0, 0
for z in 0..num_atoms:
logit = v_row[z] + adv_a[z]
if logit > M:
scale = exp(M - logit) # ≤ 1, no overflow
S, TZ, TZ², TLM *= scale
M = logit
w = exp(logit - M)
S += w
TZ += w * z_val
TZ² += w * z_val²
TLM += w * (logit - M)
expected_q = TZ / S
sum_z_sq = TZ² / S
entropy = log(S) - TLM / S # analytical: -sum(p log p)
Correctness: bit-stable when atoms processed in fixed order
(z = 0..num_atoms-1). The analytical entropy form is mathematically
identical to -sum(p log p):
-sum(p log p) = log(S) - (1/S) * sum(exp(logit-M) * (logit-M))
= log(S) - TLM / S
The previous code's `if (prob > 1e-10f)` underflow guard is no longer needed:
exp(logit - M) underflows cleanly to 0 when logit << M, and the analytical
form does not multiply tiny probs by very negative logs. Bit-stable per
feedback_stop_on_anomaly.md — no fuzzy tolerance change.
Atom-stat block-sum reduction unchanged (warp shfl + shared-mem bank,
deterministic).
Expected wall-clock saving: ~33% of 207ms = ~70ms across the smoke run; on
production batch=8192 × 1382 calls per fold, roughly proportional. Lower-bound
estimate — at low batch (smoke) the kernel may be launch-bound rather than
load-bound. nsys re-profile after deployment will quantify.
ABI unchanged; no Rust caller changes required. Per Invariant 7 the audit
doc dqn-gpu-hot-path-audit.md is updated with Fix 19 entry.
Build: SQLX_OFFLINE=true cargo check -p ml --lib clean (12 warnings, baseline).
Tests: cargo test -p ml --lib --no-run clean.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Phase 1 of the HEALTH_DIAG GPU port (Phase 0 inventory landed in
333ea7184). Lands the typed snapshot struct and the mapped-pinned
allocation that subsequent phases write into and read from — no kernel
and no caller wiring yet, deliberately so the type can stabilise before
Phase 2 invests in CUDA shmem reductions against its byte layout.
`HealthDiagSnapshot` (#[repr(C)] POD, 147×4 = 588 bytes) holds every
numeric field emitted across the 7 HEALTH_DIAG log sites. All fields
are `f32` or `u32` so CPU and GPU agree byte-for-byte with no padding
to reason about. Controller fire booleans (`fire_lr`/`fire_tau`/etc.)
are promoted from u8 → u32 per the Phase 0 design review's open-
question #3 — a `[u8; 6]` followed by `f32` would risk implementation-
defined padding mismatch between Rust's #[repr(C)] and CUDA's struct
layout rules; the +24 bytes of cost is acceptable.
`MappedHealthDiagSnapshot` wraps `cuMemHostAlloc(DEVICEMAP|PORTABLE)`
of `sizeof(HealthDiagSnapshot)` plus `cuMemHostGetDevicePointer_v2` —
mirrors the existing `MappedF32Buffer` allocator pattern but typed.
The kernel chain in Phase 2 will write through `dev_ptr()`; the host
emit code in Phase 4 reads through `host_ref()`. No `memcpy_dtoh`,
no `Vec` allocator on the path — consistent with
`feedback_no_htod_htoh_only_mapped_pinned.md`.
Three unit tests pin the layout (`snapshot_size_is_stable`,
`alignment_is_4_bytes`, `default_is_zeroed`) so any future field
add/reorder requires an explicit test update — guards the kernel-side
offset table from silent drift between commits. The size assertion
records the field count breakdown line by line in a comment so the
next maintainer can audit the math without re-deriving it.
Touched: cuda_pipeline/health_diag.rs (+339 LOC new),
cuda_pipeline/mod.rs (+8 re-exports). cargo check clean at 12 warnings
(workspace baseline). Three new unit tests pass on local CPU host (no
GPU required). Audit doc note added per Invariant 7 — explicitly notes
producer kernel + caller migrations land in Phase 2 / 4 commits.
No fingerprint change — separate mapped-pinned alloc, not part of the
param-tensor layout the fingerprint guards.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
nsys profile of multi_fold_convergence on L40S identified backtest_state_gather as
the #1 GPU consumer at 37.2% (596 ms / 60,588 calls — kernel launch latency
dominated compute) and the per-step TLOB cuBLAS gemvx calls as #2 at 25%
(242K calls). Both share the same per-step amplification: chunk_len=512 separate
gather launches + 512 separate TLOB.forward calls (4 SGEMMs each) per chunk
before any Q-values can be computed.
This commit replaces the per-step gather + DtoD pattern with a single batched
launch, and reuses the same chunked buffer for a single chunk-wide TLOB forward.
Per-chunk launch reduction: from 2*chunk_len + chunk_len*7 to 1 + 7 for the
gather+TLOB phase (4608 -> 8 with chunk_len=512, a 576x reduction).
New kernel `backtest_state_gather_chunk` (experience_kernels.cu):
- Writes [chunk_len, N, padded_sd] directly into chunked_states_buf
- chunk_len * N threads, 1 thread per output row
- Mathematically identical to per-step gather: portfolio_buf and plan_isv_buf
are CONSTANT within a chunk (env_step + plan_state_isv update at chunk
boundary only). Each thread reads independent feature offsets, no atomics,
no reordering.
Chunked val TLOB (gpu_backtest_evaluator.rs + metrics.rs):
- Val TLOB instance now sized to DQN_BACKTEST_CHUNK_SIZE * n_windows via new
GpuBacktestEvaluator::val_tlob_batch_size() helper.
- submit_dqn_step_loop_cublas calls tlob.forward(chunked_states, batch) ONCE
per chunk on the chunk-wide buffer instead of chunk_len times on states_buf.
- Partial last chunks reuse the same buffers (forward(b) accepts any
b <= construction_batch).
Borrow restructure:
- Removed top-of-function `let ch_states = self.chunked_states_buf.as_ref()?`
binding (TLOB needs &mut). Replaced with per-chunk ch_states_base raw u64
device pointer extracted in tight scope, reused by Phase 1 (gather) and
Phase 2+3 (compute_q_values_to + last_step_states_ptr). The pointer is
stable across the chunk because the Option<CudaSlice<f32>> does not
reallocate.
Per-step gather kernel `backtest_state_gather` retained unchanged for
evaluate() / evaluate_ppo() / evaluate_supervised() paths that still need
a single-step writer (closure-based callers with no chunked buffer).
Audit doc dqn-gpu-hot-path-audit.md updated with Fix 18 entry per Invariant 7.
Build: SQLX_OFFLINE=true cargo check -p ml --lib clean (12 warnings, baseline).
Tests: cargo test -p ml --lib --no-run clean.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Catalogue every numeric field in the HEALTH_DIAG family of log lines
(7 emit sites across training_loop.rs + metrics.rs) and classify each
source as GPU-already / CPU-bound / Mixed / Host-state / ISV-already.
Identifies the dominant cost contributors driving the observed
~70 s/epoch HEALTH_DIAG overhead on L40S after the eval async-split
(commit f815f7239) — they are all CPU-side reductions over multi-million
element per-sample buffers fetched via memcpy_dtoh:
- update_q_mag_means_cached (~786 KB DtoH + B×total_actions loop)
- var_scale_epoch_mean (N-element DtoH + conditional avg)
- trail_fire_and_hold_per_mag (4× N-element DtoH + per-mag loop)
- per_magnitude_winrate_and_variance (4× DtoH + sumsq loop)
- reward_contrib_fractions (6× large-buffer DtoH + 6 passes)
- read_eval_action_distribution_* (4 sites; CPU iter on pinned)
- branch_noisy_sigma_mean (NoisyNet param DtoH per branch)
Aggregate per-emit DtoH ≈ 50-150 MB depending on alloc_episodes ×
alloc_timesteps; fully consistent with the 70 s/epoch observation.
All inputs are already device-resident, so the migration is "stop
reducing on the host" plus a small kernel family writing into a single
mapped-pinned HealthDiagSnapshot struct.
Phase 0 deliverable per the dispatching prompt; subsequent commits
(struct + mapped wrapper, kernel family, CPU emit rewrite, audit-doc
updates) follow the phased plan documented at the bottom of the file
and gate-review on this inventory before touching code.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per feedback_mbp10_mandatory.md, MBP-10 + trades are mandatory inputs.
multi_fold_convergence now reads FOXHUNT_MBP10_DATA / FOXHUNT_TRADES_DATA
env vars, forwards them to train_baseline_rl as --mbp10-data-dir /
--trades-data-dir, and hard-fails if either path is missing.
Argo sanitizer-test + nsys-test templates set the env vars to PVC paths
/data/test-data/{mbp10,trades} and FOXHUNT_TEST_DATA points at
/data/test-data/ohlcv (PVC layout has lost+found/trades/mbp10/ohlcv as
top-level subdirs; the symbol ES.FUT lives under ohlcv/).
Without this, the smoke would silently fall back to tick-rule proxy
classification and validate a degraded model variant (no real OFI),
defeating the purpose of L40S validation. Audit doc updated per
Invariant 7.
After Fix 1..16 migrated all 80+ production callers off
`super::htod_f32` and `super::clone_htod_f32`, the helper bodies in
`cuda_pipeline/mod.rs:129-145` had zero non-test consumers. Deleted
both function definitions per `feedback_no_legacy_aliases.md` (no
deprecated wrappers).
Per `feedback_no_partial_refactor.md` (when a shared contract is
deleted, every consumer migrates together — including tests), the
two surviving test-block callers in `gpu_tlob.rs::tests` (lines
1017 and 1132) are migrated to `mapped_pinned::upload_f32_via_pinned`
in the same commit. The other test-only callers in
`signal_adapter.rs::tests`, `gpu_action_selector.rs::tests`, and
`cuda_pipeline/mod.rs::tests` use bare `stream.memcpy_htod` /
`stream.memcpy_stod` against the cudarc handle directly (not the
deleted helpers) — no change needed.
A docstring was added at the deletion site recording when and why
the helpers were removed, pointing future readers at the canonical
replacements `mapped_pinned::clone_to_device_f32_via_pinned` and
`mapped_pinned::upload_f32_via_pinned`.
Final state of the HtoD migration sequence:
- production callers of `stream.memcpy_htod` / `memcpy_stod`: 0
- production callers of `htod_f32` / `clone_htod_f32`: 0
- helper definitions: removed from `mod.rs`
docs/dqn-gpu-hot-path-audit.md updated with Fix 17 entry.
cargo check -p ml --lib clean at 12 warnings.
cargo check -p ml --tests clean at 23 warnings (12 lib duplicates +
11 test-specific, baseline unchanged).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The two `crate::cuda_pipeline::clone_htod_f32` callsites in
`training_loop.rs::set_raw_market_data` (lines 1250 and 1254) are
rewritten to `mapped_pinned::clone_to_device_f32_via_pinned` per
`feedback_no_htod_htoh_only_mapped_pinned.md`.
These run once per fold immediately before the step loop begins; the
`_raw` suffix indicates the un-normalized arrays kept on-GPU for the
eval-time critic.
Error-type continuity: the helper already returns `Result<_, String>`,
which feeds `anyhow::anyhow!("…raw upload: {e}")` directly — the prior
.map_err wrapper is preserved verbatim. 1:1 path swap.
docs/dqn-gpu-hot-path-audit.md updated with Fix 16 entry.
cargo check -p ml --lib clean at 12 warnings.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The remaining `super::clone_htod_f32` call in `upload_ofi_features`
(`:4258`) is rewritten to `mapped_pinned::clone_to_device_f32_via_pinned`
per `feedback_no_htod_htoh_only_mapped_pinned.md`. Called from
training_loop's data-load phase before each fold's training run, so
the OFI tensor (4M bars × 32 dims) now stages through mapped-pinned
+ DtoD just like all other large data uploads.
Error-type shift: `mapped_pinned::clone_to_device_f32_via_pinned`
returns `Result<CudaSlice<f32>, String>` whereas `clone_htod_f32`
returned `Result<CudaSlice<f32>, MLError>`. Wrapped with
`.map_err(|e| MLError::ModelError(format!("upload_ofi_features: {e}")))`
to preserve the call-site label in error messages.
docs/dqn-gpu-hot-path-audit.md updated with Fix 15 entry.
cargo check -p ml --lib clean at 12 warnings.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Residual COLD ctor sites missed by the original audit are now migrated
off explicit HtoD per `feedback_no_htod_htoh_only_mapped_pinned.md`:
- `:9922` sel_clip_buf 1-element init (sigmoid head clip-norm seed)
- `:10075-10078` spectral norm init_u_s1/v_s1/u_s2/v_s2 (4 calls)
- `:10095-10096` `alloc_spec_pair!` macro body — both u/v init uploads
inside the macro now go through `upload_f32_via_pinned`; the macro's
`$lbl_u`/`$lbl_v` are reused in the error-message format so per-pair
failures stay diagnostically distinct
- `:11276` graph_params_host (60 floats: cross-branch graph message
passing weights)
- `:11330` denoise_params_host (1800 floats: 2-step diffusion Q-refinement)
- `:11476` qlstm_weights_host (528 floats: QLSTM Xavier init)
All sites use `mapped_pinned::upload_f32_via_pinned` (the canonical
mapped-pinned + DtoD staging helper). The helper returns
`Result<_, String>` whereas this constructor returns
`Result<_, MLError>`, so each site wraps the error via
`.map_err(|e| MLError::ModelError(format!("<site> upload via pinned: {e}")))`.
Site labels preserved so backtraces remain readable.
docs/dqn-gpu-hot-path-audit.md updated with Fix 14 entry.
cargo check -p ml --lib clean at 12 warnings.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Agent 1 — bulk COLD/WARM migration across init/loader files.
Sites: 39 production HtoD calls migrated to MappedF32Buffer / MappedI32Buffer /
new MappedU32Buffer. Cross-file infrastructure includes new staging helpers
(clone_to_device_{f32,i32}_via_pinned, upload_{f32,i32,u32}_via_pinned),
MappedU32Buffer type, and host_slice_mut accessor on MappedF32/U32Buffer.
Per feedback_no_htod_htoh_only_mapped_pinned.md — first of 3 parallel batches.
# Conflicts:
# docs/dqn-gpu-hot-path-audit.md
HOT sites:
- L2992 hindsight bar_indices: was `alloc + memcpy_htod(total i32s)` on
every collect_experiences_gpu when hindsight_fraction > 0. Added
persistent `bar_indices_pinned: MappedI32Buffer` (capacity =
alloc_episodes * alloc_timesteps * 2) allocated at constructor.
Per-call host writes go through host_ptr; relabel kernel reads via
dev_ptr u64.
- L3186 feature_mask: was alloc + memcpy_htod each epoch at t==0.
Promoted `feature_mask_buf` to `Option<MappedF32Buffer>`. Reallocates
only when mask size changes; CPU writes via host_ptr, state_gather
reads via dev_ptr.
- L3816 update_per_sample_support: per_sample_support_buf is read-only
by the kernel — promoted to MappedF32Buffer (kernel args now
dev_ptr u64). per-epoch tile fill becomes a direct host_ptr write.
WARM sites:
- L3097 episode_starts upload: episode_starts_buf is GPU-mutated by
domain_rand_episode_starts so must remain CudaSlice. Stage via
MappedI32Buffer + memcpy_dtod_async through new helper
`upload_host_to_cuda_i32_via_pinned`.
- L2009 upload_expert_actions: promoted expert_actions_gpu from
Option<CudaSlice<i32>> to Option<MappedI32Buffer>. Direct host_ptr
write replaces alloc + memcpy_htod.
COLD sites:
- L1076,3848 portfolio_states init/reset (GPU-mutated): use
upload_host_to_cuda_f32_via_pinned (mapped-pinned staging + DtoD).
- L1093 epoch_state init (GPU-mutated): replace clone_htod_f32 with
explicit alloc_zeros + upload_host_to_cuda_f32_via_pinned.
- L1341 saboteur_base init (GPU-mutated): use the same helper.
- L2784 trade_stats_buf pre-reduction zero: replaced htod_f32 of an
all-zeros vec with stream.memset_zeros — fully GPU-side, no PCIe.
Imports: added DevicePtrMut for memcpy_dtod_async pointer extraction.
Per `feedback_no_partial_refactor.md`, every consumer of the migrated
buffers is updated in this commit:
- per_sample_support_buf kernel args at L3426 (compute_expected_q) and
L3546 (quantile_q_select) now consume dev_ptr u64.
- feature_mask_buf accessor at L3195 reads via .dev_ptr.
- expert_actions_gpu field type change is contained (no external
consumers in this crate).
cargo check -p ml --lib: 11 warnings (unchanged from prior commit on
this branch; 2 below 13-warning baseline because two unrelated trivial
warnings disappeared as a side-effect of the refactor).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per `feedback_no_htod_htoh_only_mapped_pinned.md`, mapped pinned
(cuMemHostAlloc DEVICEMAP) is the only allowed CPU↔GPU path.
Migrates 3 WARM HtoD sites:
- `fused_training.rs:505` HER source_indices Vec<i32> upload
- `fused_training.rs:511` HER reward_ones Vec<f32> upload
- `training_loop.rs:470` curriculum episode_starts Vec<i32> upload
For `fused_training.rs` adds module-level `staging_upload_{i32,f32}`
helpers. For `training_loop.rs` uses an inline staging+DtoD block
because the destination is reached via the out-of-scope accessor
`collector.episode_starts_buf_mut() -> &mut CudaSlice<i32>`.
All destination buffer types remain `CudaSlice<T>` so no downstream
consumer (incl. `gpu_her::relabel_batch_with_strategy` and any
`episode_starts_buf` reader) needs to change.
Note: if Agent 2's `gpu_experience_collector` merge changes
`episode_starts_buf_mut`'s return type to a mapped pinned buffer,
the inline staging block in `training_loop.rs:470` can be simplified
to a direct `write_from_slice` (follow-up).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
