Per Class A audit P1 wiring batch — only 1 of 4 items was wireable as
spec'd. The other 3 are deferred per the spec's "DEFER not BLOCK" rule
because their target ISV slots either don't exist or are unit-incompatible
with the consumer site.
Items:
1. trade_physics.cuh:548 (DD floor_dd 0.25f) — DEFERRED. Slot 421
(DD_THRESHOLD_INDEX) holds the DD trigger threshold (~0.05), NOT the
penalty saturation floor (~0.25). Different parameters; substituting
would break the ramp formula `(dd-thr)/(floor-thr)`. Needs a new slot
for the saturation floor.
2. gpu_experience_collector.rs:343-344 (dd_threshold + w_dd config) —
DEFERRED. No W_DD slot exists in any sp*_isv_slots.rs. The legacy
compute_drawdown_penalty path (experience_kernels.cu:3769) uses
config scalars; the SP15 reward axis (compute_sp15_final_reward_kernel)
already reads slot 421 directly. Wiring needs a new W_DD slot.
3. plan_threshold_update_kernel.cu:44 (plan threshold floor) — DEFERRED.
Unit mismatch: ema and plan_threshold are probabilities ∈[0,1];
Q_DIR_ABS_REF (slot 21) is a Q-value magnitude EMA (5–50). Scaling
a probability by a Q-magnitude is dimensionally meaningless.
4. experience_kernels.cu:2471 (var_floor formula) — DONE. Drop the
hardcoded 0.25f lower bound. Was `fminf(fmaxf(q_gap*0.5f, 0.25f), 1.0f)`;
now `fminf(q_gap*0.5f, 1.0f)`. NULL-q_gaps fallback uses ε=1e-6f for
numerical safety (vs prior 0.25f). var_scale itself remains (0,1] from
`1/(1+sqrt(var_q))` so the multiplicative chain stays bounded.
Cumulative WR-plateau fix series:
- Class C bug 1 + P0-B (8f218cab2)
- P0-C (316db416b)
- P0-A (394de7d43)
- P1 wiring (this commit, partial — Items 1/2/3 deferred to producer batch)
Per feedback_isv_for_adaptive_bounds.md.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per Class A audit ranking, the highest-suspected-impact fix for the
months-long WR-stuck-at-46-48% plateau across 11 superprojects.
Hardcoded REWARD_POS_CAP=+5.0f / REWARD_NEG_CAP=-10.0f
(state_layout.cuh:266-267) was structurally clipping the upper tail of
realized alpha. Controller signals (sharpe EMA, var_q, q_gap) all
derive from this CAPPED buffer, so the controller cannot select for
trades it cannot see. Selectivity gradient evaporates — small wins
clip to +5 alongside large wins also clipping to +5.
The state_layout comment lines 261-264 explicitly deferred Phase 2
(ISV-driven) "IF Phase 1 validation reveals adaptive need". Phase 1
has been running 11 SPs without budging WR — adaptive need revealed.
Architecture:
- 2 new ISV slots [452..454): REWARD_POS_CAP_ADAPTIVE,
REWARD_NEG_CAP_ADAPTIVE.
- Producer kernel reward_cap_update_kernel.cu: block-tree-reduce
Welford `mean + Z_99 × sigma` p99 estimator over winning realized
returns + conservative `max(p99, max_win)` takeover, × 1.5 safety
factor → POS cap. NEG cap = -2 × POS cap (preserves Kahneman 2:1
asymmetry per pearl_audit_unboundedness_for_implicit_asymmetry —
asymmetry stays, but moved from hardcoded scalar to producer-time
multiplier; single source of truth, no consumer applies the 2× ratio
itself).
- Pearl-A first-observation bootstrap from sentinel (5.0 / -10.0,
matching pre-P0-A hardcoded values for bit-identical cold-start).
Welford EMA α=0.01 thereafter (slow blend — reward distribution is
the foundation of training and shouldn't move fast).
- Bounds: POS in [1, 50], NEG in [-100, -2] (Category-1 dimensional
safety per feedback_isv_for_adaptive_bounds, NOT tuning).
- 3 consumer sites migrated atomically per
feedback_no_partial_refactor: experience_kernels.cu:3112-3114
(segment_complete cap), compute_sp15_final_reward_kernel.cu:163
(Stage 4 helper invocation), sp15_reward_axis_helpers.cuh:211
(sp15_apply_sp12_cap device fn signature change to take isv ptr).
- Cold-start fallback: when ISV slot at sentinel OR outside
[REWARD_POS_CAP_MIN_BOUND=1, REWARD_POS_CAP_MAX_BOUND=50],
consumers fall back to original macros (still defined in
state_layout.cuh).
- 2 new device-ptr accessors on the experience collector
(step_ret_per_sample_dev_ptr, trade_close_per_sample_dev_ptr) —
reuses existing per-sample buffers; no new buffer allocated.
- Per-epoch boundary launch (cold path) in training_loop.rs alongside
launch_aux_horizon_chain.
- Reset registry entries + dispatch arms in reset_named_state per the
C.10 lesson (missing dispatch causes runtime crash).
- Layout fingerprint seed updated: ISV_TOTAL_DIM 452→454 +
AUX_PRED_HORIZON_BARS=450 + AVG_WIN_HOLD_TIME_BARS=451 (previously
missing from seed) + REWARD_POS_CAP_ADAPTIVE=452 +
REWARD_NEG_CAP_ADAPTIVE=453.
Per feedback_isv_for_adaptive_bounds: every adaptive bound in ISV.
Verification:
- cargo check -p ml --tests --all-targets: clean (19 pre-existing
warnings, 0 new).
- sp14_isv_slots tests: 8/8 pass (4 layout + 4 fits-within).
- sp14_oracle_tests with --features cuda --ignored: 8/8 pass (4
existing q_disagreement/dir_concat + 4 new P0-A tests covering
Pearl-A bootstrap, no-winning-trades preservation, bounds clamping
to [1,50], Welford α=0.01 EMA blend).
- sp15_phase1_oracle_tests with --features cuda --ignored: 36/36 pass
(no regression from sp15_apply_sp12_cap signature change).
Cumulative WR-plateau fix series:
- Class C bug 1 (8f218cab2): replay buffer intent→realized.
- Class A P0-B (8f218cab2): Kelly warmup floor wiring.
- Class A P0-C (316db416b): MIN_HOLD_TARGET adaptive.
- Class A P0-A (this commit): REWARD_POS/NEG_CAP adaptive — restores
upper-tail alpha to the controller's view.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per Class A audit: MIN_HOLD_TARGET=30.0f hardcoded was creating a
deterministic gradient pushing trades toward 30-bar holds regardless of
edge expiry. User's trading frequency is between HFT-MFT and varies by
regime; a 30-bar fixed target kills MFT-frequency alpha when the
optimal hold for current data is shorter (or longer).
The producer slot ISV[AVG_WIN_HOLD_TIME_BARS_INDEX=451] already exists
from SP14 Layer C Phase C.4b (commit 3b71d2183) — Pearl-A-bootstrapped
Welford EMA of observed winning trade hold times. Wiring fix only.
Cold-start fallback: when slot still at sentinel (no winning trades
observed yet), use MIN_HOLD_TARGET=30.0f as safety floor. Once a
winning trade closes and the EMA bootstraps, the adaptive value
takes over.
Validity window: isv_hold_target > 0.0f && < 240.0f; outside window
falls back to min_hold_target param (= MIN_HOLD_TARGET=30).
Added #define AVG_WIN_HOLD_TIME_BARS_INDEX 451 to state_layout.cuh
(C-side mirror of sp14_isv_slots.rs:97).
Per feedback_isv_for_adaptive_bounds: every adaptive bound in ISV.
Fixes the third Class A P0 hardcoded constant.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Two independent bugs surfaced by Class C (frame-shift) + Class A (hardcoded
bounds) audits, both implicated in the months-long WR-stuck-at-46-48%
plateau across 11 superprojects.
## Bug 1: Replay buffer Sutton's deadly triad
experience_kernels.cu:2275 was writing the original INTENT action to
out_actions, but the reward in the replay buffer was computed from the
REALIZED position (post-enforcement: Kelly cap, capital floor, trail-stop,
broker cap can clamp Long→Flat). Replay buffer stored (s, intent, r_realized,
s'). Q(s, Long) was therefore trained against r(s, Flat) whenever env
clamped the intent.
This explains the train_active_frac=0.40 vs val_active_frac=0.05 gap:
train measures intent (40% Long/Short), eval measures realized (5%
Long/Short). The 8× gap is env physics draining intent.
Fix: after unified_env_step_core resolves actual_dir_core/actual_mag_core,
overwrite out_actions[out_off] with the realized action (same encoding as
backtest_env_kernel.cu:323-330, which has been doing it correctly all
along). Order/urgency preserved from intent.
## Bug 2: Kelly cap update kernel ignored existing ISV warmup floor
kelly_cap_update_kernel.cu:53 hardcoded the kelly_f floor at 0.0f. Cold
path (per-epoch boundary). Per project_magnitude_eval_collapse_kelly_capped,
this collapses kelly_cap to 0 → max position pinned to Quarter for cold
start. The val-mag pathology.
The warmup floor producer (ISV[KELLY_WARMUP_FLOOR_INDEX=330], SP9 Fix 37)
was already populated and consumed by the per-step path at
trade_physics.cuh:377-384, but this cold-path kernel never read it.
Partial wiring.
Fix: replace fmaxf(kelly_f, 0.0f) with fmaxf(kelly_f, isv[330]). One-line
change.
## Predicted effect
- train_active_frac and val_active_frac should converge (Bug 1 inflated
train by counting overridden intents)
- Magnitude distribution should escape Quarter-only (Bug 2 was pinning it)
- WR ceiling at 46-48% may finally move (Bug 1 broke Bellman consistency;
Bug 2 prevented edge realization)
Falsification: 5-epoch L40S smoke. If unmoved by ep5, the plateau is
deeper still (Class A P0-A REWARD_POS_CAP/NEG_CAP next).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Phase C.1's atomic α deletion preserved Q_DISAGREEMENT_VARIANCE_EMA_INDEX=389
(HEALTH_DIAG diagnostic) and its state_reset_registry entry, but the
corresponding dispatch arm in reset_named_state was missing — even though
the inline comment at training_loop.rs:8024 falsely claimed it existed.
train-rqd8r (commit 10e647c14) failed both folds at fold-reset boundary:
Model error: StateResetRegistry reset dispatch:
unknown name 'sp14_q_disagreement_variance_ema'
Same class as SP5 Layer A bug-fix #281. Dispatch arm restored using
SENTINEL_VARIANCE (0.0) per the registry entry's documented sentinel.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
C.8 (ISV-driven aux trunk Adam β1/β2/ε/LR/grad-clip) was already complete in C.5a
commit c90de9859 — all 5 ISV reads and fold-boundary StateResetRegistry defaults were
wired atomically with the Adam launcher. No new code required; noted in audit doc.
C.9 adds `aux_trunk_learns_synthetic_uptrend` to aux_trunk_oracle_tests.rs:
- B=16, ENC=32, H1=32, H2=16, SH2=32, H_HEAD=32, K=2, 100 steps
- Backward kernel invocations corrected to match actual signatures:
- aux_trunk_bwd_dh_pre(d_logits, w3, w2, h_aux1, h_aux2, dh_pre2, dh_pre1, B, H1, H2, SH2)
shmem = H2 floats (sh_dh2_pre cache), NOT (H1+H2+SH2)
- aux_trunk_bwd_dW_reduce called 3×: dW3/dW2/dW1 each with (A, B_grad, dW_out, B, Krows, Jcols)
- aux_trunk_bwd_db_reduce called 3×: db3/db2/db1 each with (B_grad, db_out, B, Jcols)
- Head params trained via host-side SGD (test orchestration only; reads mapped-pinned partials)
- Trunk params trained via dqn_adam_update_kernel (GPU Adam)
- Pass gate: CE loss < 0.1 AND dir_acc > 0.95 after 100 steps
Near-random baseline (ln(2)≈0.693) = broken gradient chain, L40S dispatch blocked
Memory pearl pearl_separate_aux_trunk_when_shared_starves.md added and indexed.
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Single-block 256-thread CUDA kernel computing RMS(h_s2_aux [B, SH2])
and EMA-blending the step observation into ISV[H_S2_AUX_RMS_EMA_INDEX=449]
directly. Pearl-A first-observation bootstrap embedded in kernel body
(sentinel 0.0 → replace); fixed α=0.05 EMA blend thereafter.
ISV slot 449 is outside the SP4/SP5 wiener buffer linear span so the
scratch+apply_pearls_ad_kernel path is not available — self-contained
Pearl-A logic mirrors the avg_win_hold_time_update_kernel precedent
(slot 451). No atomicAdd; shmem block-tree-reduce only. Launched after
aux_trunk_forward in the collector per-step hot path.
- h_s2_aux_rms_ema_kernel.cu — new CUDA kernel (81 lines)
- build.rs — cubin manifest entry
- gpu_dqn_trainer.rs — H_S2_AUX_RMS_EMA_CUBIN static
- gpu_aux_trunk.rs — HS2AuxRmsEmaOps struct + launch()
- gpu_experience_collector.rs — field + constructor + hot-path launch
- aux_trunk_oracle_tests.rs — h_s2_aux_rms_ema_pearl_a_bootstrap test
- dqn-wire-up-audit.md — Phase C.6 audit entry
cargo check -p ml --tests: clean (only pre-existing warnings)
Oracle test: 1 new test added (requires GPU to run)
Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
Atomic flip of the aux heads' input from Q's GRN trunk output `save_h_s2`
to the SEPARATE aux trunk's output `h_s2_aux`. The aux trunk now trains
its own w1/w2/w3/b1/b2/b3 from CE loss (next-bar + regime); Q's encoder
is structurally protected by `aux_trunk_backward`'s missing `dx_in`
output param (encoder boundary stop-grad enforced at the kernel-set
level).
Reverts the C.0 stop-grad band-aid commits (`872bd7392`, `411a30473`):
the zero-fills in `aux_next_bar_backward` + `aux_regime_backward` Step 3
are replaced with the genuine SAXPY-back-to-input gradient
(`dh_s2_aux[b,j] = sum_k sh_dh_pre[k] * w1[k,j]`). The leak that
motivated stop-grad is now blocked structurally rather than by data
zero-fill — aux gradient flows through the aux trunk's own params, never
into Q's encoder.
Wired in this commit (atomic, ~330 LOC):
- 4 kernel signatures renamed `h_s2 → h_s2_aux` / `dh_s2_out → dh_s2_aux_out`
(`aux_next_bar_forward`, `aux_regime_forward`, `aux_next_bar_backward`,
`aux_regime_backward`); Rust wrappers in `gpu_aux_heads.rs` follow
- Trainer fwd: insert `aux_trunk_forward_ops.launch(...)` in
`aux_heads_forward` Step 0, populating `h_s2_aux` from `save_h_s1`
(encoder layer-1 output, dim=shared_h1=256). Both head fwds redirect
input pointer from `save_h_s2` to `h_s2_aux`
- Trainer bwd: SAXPY both `aux_dh_s2_*_buf` into `dh_s2_aux_accum`
(pre-zeroed each step via graph-safe `cuMemsetD32Async`); then
`aux_trunk_backward_ops.launch(...)` propagates through w3/w2/w1 +
b3/b2/b1; then `launch_aux_trunk_adam_update` applies global L2-norm
clip + per-tensor Adam updates over 6 grad tensors
- Collector fwd: insert `exp_aux_trunk_forward_ops.launch(...)` after
`forward_online_f32`, reading `exp_h_s1_f32` and writing `exp_h_s2_aux`;
redirect `exp_aux_heads_fwd.forward_next_bar` input from
`exp_h_s2_f32` to `exp_h_s2_aux`
- Pre-capture host-write of ISV-driven LR + grad-clip + step counter
into mapped-pinned buffers in `launch_cublas_backward_to` (BEFORE
`aux_heads_backward`); same `&mut self` pattern as `step_ofi_embed_adam`
Verification:
- `cargo check -p ml --tests` clean (1m02s, only pre-existing warnings)
- `aux_trunk_oracle_tests` + `sp14_oracle_tests` 12/12 pass:
- aux_trunk gradient check: max_rel_err=1.33e-2 (tol=2e-2) — matches C.4 baseline
- aux_trunk_backward_does_not_write_dx: kernel source clean of dx_in/dx_in_out
- aux_sign_label_lookahead_mask: 60/100 masked, 40/100 valid
- 9 other oracle tests pass bit-identically
Plan: docs/superpowers/plans/2026-05-07-sp14-layer-c-separate-aux-trunk.md §C.5b
Audit: docs/dqn-wire-up-audit.md "SP14 Layer C Phase C.5b" section
Phase C.5a-fixup — completes C.5a's additive infrastructure so C.5b can be
a true atomic contract flip with zero scaffolding work mixed in:
- Allocate dh_aux1_pre_scratch [B_max, AUX_TRUNK_H1=256] +
dh_aux2_pre_scratch [B_max, AUX_TRUNK_H2=128] (missed in C.5a — both
required by aux_trunk_backward.launch per gpu_aux_trunk.rs:266-267).
- Collector struct gains 6× u64 aux_trunk_{w1,b1,w2,b2,w3,b3}_ptr fields,
exp_aux_trunk_forward_ops: AuxTrunkForwardOps field (constructed in
ctor on collector's stream), and set_trainer_aux_trunk_param_ptrs
setter — mirrors the existing set_trainer_params_ptr zero-copy pattern.
- Trainer gains aux_trunk_param_ptrs() -> (u64×6) accessor returning
raw_ptr() for all 6 aux trunk parameter tensors.
- training_loop wires the new setter at both existing
set_trainer_params_ptr call sites (initial fused-ctx init + fold-boundary
re-init).
NO contract change: wire sites still call save_h_s2. The setter is called
and the 6 aux trunk param ptrs are populated, but no collector-side launch
reads them yet — C.5b atomically inserts aux_trunk_forward.launch(...)
post-forward_online_f32 and switches the aux head input pointer.
Graph-capture audit: aux_heads_backward IS INSIDE the captured `forward`
child graph (call chain: submit_forward_ops_main → launch_cublas_backward
→ launch_cublas_backward_to → aux_heads_backward; capture begins at
fused_training.rs:2964 / capture_child_graph). The existing function body
is fully device-side (zero host writes) — capture-safe by construction.
C.5b's new aux_trunk Adam launch is also fully device-side and will sit
inside the same captured region. The host writes for aux_trunk_t_pinned
must use the existing GPU-side increment_step_counters kernel chain
(submit_counters_ops, line 22799) — NOT host-side aux_trunk_adam_step
+= 1 inside capture. ISV-driven LR/clip writes happen pre-capture
(cold-path); the captured graph reads via aux_trunk_lr_dev_ptr /
aux_trunk_grad_clip_dev_ptr. This avoids the &self → &mut self ripple on
aux_heads_backward (gap 4 in C.5b implementer's blocker report). Full
wiring strategy + alternative (pre-capture host-write) documented in
docs/dqn-wire-up-audit.md C.5a-fixup section.
Verification:
- cargo check -p ml --tests --all-targets: clean (no new warnings).
- cargo test -p ml --test aux_trunk_oracle_tests --test sp14_oracle_tests
--release -- --ignored --nocapture: 12/12 pass (8 aux_trunk + 4 sp14;
bit-identical to C.5a baseline — pure scaffolding, no regression).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Phase C.5a — additive infrastructure for the aux trunk wire-up.
Allocates saved-fwd buffers (h_s2_aux, h_aux1, h_aux2),
gradient buffers (6× aux_trunk_*_grad), accumulator
(dh_s2_aux_accum), and dedicated Adam launcher
(launch_aux_trunk_adam_update) reading β1/β2/ε/LR/grad-clip from
ISV[444..449).
No contract change. No call sites for the new launcher yet.
C.5b atomically wires these in.
Phase C.5 was split (authorized 2026-05-08) after the original
implementer flagged ~600-800 LOC scope across 4 files with
correctness windows. C.5a is purely additive; C.5b is the genuine
~300 LOC atomic migration.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Aux's original label was (p_{t+1} > p_t) — pure HFT-scale microstructure
noise that's unlearnable at our HFT-MFT trading frequency. Migrated to
(p_{t+H} > p_t) where H is read from ISV[AUX_PRED_HORIZON_BARS_INDEX=450].
Adaptive producer drives H from observed avg winning hold time:
- Pearl-A first-observation bootstrap: replace sentinel H=60 directly
on first valid observation
- Steady-state Wiener-α EMA blend, slow (α=0.01) for stable horizon
(no target-variance EMA available, fallback per
pearl_wiener_optimal_adaptive_alpha)
- "No winning trades yet" guard keeps sentinel until first valid observation
Lookahead truncation: labels at t where t+H >= total_bars are masked
(sentinel -1, loss-reduce skips). The existing aux_next_bar_loss_reduce
in aux_heads_kernel.cu already supports the -1 mask convention via the
B_valid count — no new valid_mask parameter needed.
Step 5b finding: Case B — existing per-sample buffers
(hold_at_exit_per_sample, trade_profitable_per_sample) populated by
unified_env_step_core, but no aggregate ISV slot. Added new aggregator
slot AVG_WIN_HOLD_TIME_BARS_INDEX=451 + new producer kernel
avg_win_hold_time_update_kernel.cu (block-tree-reduce, no atomicAdd).
ISV_TOTAL_DIM bumped 450 → 452.
ATOMIC migration per feedback_no_partial_refactor: both label kernels
(aux_sign_label_kernel.cu trajectory + aux_sign_label_per_step_kernel.cu
per-rollout-step) migrated together to the new
(targets, bar_indices, isv, isv_h_idx, out_labels, total, total_bars)
signature. The lookahead host-passed scalar argument is removed; H is
read from ISV inside the kernel (broadcast value, single read per
thread, on-device clamp [1, 240]).
Producer chain (per-epoch boundary): new
GpuDqnTrainer::launch_aux_horizon_chain orchestrates
avg_win_hold_time_update → aux_horizon_update sequentially alongside
launch_kelly_cap_update at the existing epoch-boundary slot in
training_loop.rs.
Trunk math (C.2/C.3/C.4) unchanged — separate aux trunk is label-
agnostic. Validation in C.10 will use H=60 cold-start; the adaptive
producer drives H from real winning-trade observations.
Tests (8 oracle, 5 new + 3 preserved):
- aux_trunk_forward_matches_numpy_reference (C.3) ✓
- aux_trunk_backward_gradient_check (C.4) ✓
- aux_trunk_backward_does_not_write_dx (C.4) ✓
- aux_sign_label_h_bar_horizon (NEW) ✓
- aux_sign_label_lookahead_mask (NEW) ✓
- aux_horizon_pearl_a_bootstrap (NEW) ✓
- aux_horizon_converges_to_steady_target (NEW) ✓
- aux_horizon_holds_sentinel_with_no_winning_trades (NEW) ✓
8/8 pass on RTX 3050 Ti.
Phase C.4b of SP14 Layer C separate-aux-trunk refactor.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Backward propagates dh_s2_aux through w3/w2/w1 with block-tree-reduce
(no atomicAdd per feedback_no_atomicadd). Critical: kernel set does NOT
write dx_in — encoder gradient remains Q-shaped only. Stop-grad
invariant verified via parameter-list structural enforcement (kernels
literally cannot reference an `dx_in_out` pointer they don't accept) +
kernel source inspection that strips comments and asserts no `dx_in`
write pattern.
Three kernels in aux_trunk_backward_kernel.cu:
- aux_trunk_bwd_dh_pre: per-sample, computes dh_aux2_pre [B, H2] +
dh_aux1_pre [B, H1] using ELU' from POST-activation form
(`(y > 0) ? 1 : (1 + y)` mirrors aux_elu_bwd_from_post in
aux_heads_kernel.cu).
- aux_trunk_bwd_dW_reduce: generic outer-product reduce
`dW[k, j] = sum_b A[b, k] * B[b, j]`. One block per output
element, shmem-tree reduce over batch. Used 3× (dW3, dW2, dW1).
- aux_trunk_bwd_db_reduce: generic batch-reduce `db[j] = sum_b
B[b, j]`. One block per output element. Used 3× (db3, db2, db1).
Memory-efficient: no per-sample partials (avoids B×163,072 floats for
production topology). Per-element reduction means O(P) blocks each
doing O(B) work in shmem.
Rust wrapper AuxTrunkBackwardOps in gpu_aux_trunk.rs orchestrates seven
launches in fixed sequence (capture-friendly, no host branches per
pearl_no_host_branches_in_captured_graph). All three CudaFunction
handles pre-loaded once at construction. Field added to GpuDqnTrainer
alongside aux_trunk_forward_ops; constructor mirrors C.3 pattern.
Tests (all pass on RTX 3050 Ti, sub-ULP forward, 1.33e-2 max rel-err
backward gradient at smallest sampled gradient):
- aux_trunk_forward_matches_numpy_reference (C.3 — preserved).
- aux_trunk_backward_gradient_check (NEW): central-difference
numerical gradient at 16 sampled dW3 indices vs analytic from
backward kernel. Loss = 0.5 * ||h_s2_aux||^2 so dh_s2_aux =
h_s2_aux. EPS=1e-3, B=4, ENC=H1=H2=AUX=32 (33 forwards in ~2s).
REL_TOL = 2e-2 (f32 finite-difference noise floor for
small-gradient tail; production topology is dimension-independent
given runtime args).
- aux_trunk_backward_does_not_write_dx (NEW): reads kernel source,
strips C-style comments (so design-discussion text mentioning
`dx_in` doesn't false-positive), asserts no `dx_in` / `dx_in_out`
symbol survives in code. Complements the structural enforcement
(kernel signatures don't accept `dx_in_out` pointer).
Phase C.4 of SP14 Layer C separate-aux-trunk refactor. Module is
additive — wire-up into collector backward chain + Adam updates lands
in Phase C.5 (atomic).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
3-layer MLP forward (Linear→ELU→Linear→ELU→Linear). Pre-loaded
CudaFunction for graph-capture safety per pearl_no_host_branches_in_captured_graph.
Oracle test verifies bit-for-bit match against numpy reference within
1e-4 tol. Saves h_aux1 and h_aux2 to global memory for backward.
Phase C.3 of SP14 Layer C separate-aux-trunk refactor (plan:
docs/superpowers/plans/2026-05-07-sp14-layer-c-separate-aux-trunk.md).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
3-layer MLP: encoder_out_dim → 256 → 128 → AUX_HIDDEN_DIM. Kaiming-He
weights (Box-Muller from LCG-uniform), zero biases, separate Adam m/v
buffers (12 state tensors). Allocated in trainer constructor — collector
borrows via raw_ptr at wire-up time per existing OFI-embed / q-attn
ownership pattern (no parallel param mirror needed). Not yet wired to
forward/backward — pure allocation per Phase C.2 design.
Topology dimensions resolved against actual codebase:
- encoder_out_dim = config.shared_h1 (= SH1 = 256 in production)
- AUX_HIDDEN_DIM = config.shared_h2 (= SH2 = 256, matches existing aux
head's input dim per aux_heads_kernel.cu:118 `h_s2 [B, SH2]`)
Total params: 65,536 + 256 + 32,768 + 128 + 32,768 + 256 = 131,712.
Audit doc updated per Invariant 7. Phase C.2 of SP14 Layer C
separate-aux-trunk refactor (plan:
docs/superpowers/plans/2026-05-07-sp14-layer-c-separate-aux-trunk.md).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Bug audit Pattern 2 — 6 SP15 launchers in gpu_dqn_trainer.rs were doing
load_cubin + load_function PER CALL inside the launcher body, called
from gpu_experience_collector.rs's per-rollout-step body (~thousands
of times per epoch x 4096 envs x 1000 timesteps).
Same architectural bug class as commits 5d63762ab (bn_tanh_concat_dd)
and 1396b62ec (sp15_baseline + cost_net). load_cubin/load_function
are host-side driver API calls — not capturable inside graph capture,
and prone to CUDA_ERROR_ILLEGAL_ADDRESS in subprocess child contexts
(documented failure mode in 1396b62ec).
Migration:
- launch_sp15_dd_state, launch_sp15_dd_state_reduce,
launch_sp15_dd_trajectory_decreasing, launch_sp15_alpha_split_producer,
launch_sp15_final_reward, launch_sp15_plasticity_injection —
signatures changed to accept &CudaFunction parameter.
- Collector struct gains 6 new pre-loaded CudaFunction fields,
populated once in collector::new() (mirrors SP14 EGF kernel loading
added in commit 09202aa99).
- All 6 hot-path call sites updated to pass &self.exp_<kernel>_kernel
reference (5 in collect_experiences_gpu, 1 in plasticity-injection
trigger path).
- Oracle tests in sp15_phase1_oracle_tests.rs pass through a small
load_sp15_kernel test helper that inlines the per-call load (graph
capture isn't a concern in oracle scaffolds; production callers use
struct-cached handles via the collector).
- docs/dqn-gpu-hot-path-audit.md gains Fix 29 documenting the
pattern, the 6 migrated sites, deferred evaluator launchers, and
out-of-scope orphan launchers.
Per feedback_no_partial_refactor: all 6 launchers + their callers
migrate atomically. Per pearl_no_host_branches_in_captured_graph:
zero load_cubin in collector per-rollout-step body after this commit.
Deferred (separate scope): launch_sp15_sharpe_per_bar and
launch_sp15_position_history_derivation are evaluator-path callers
(GpuBacktestEvaluator) — fixing them requires adding fields to a
different struct, deferred to a separate atomic commit. Orphan
launchers (regret_signal, cooldown) have no production callers; left
untouched.
Cubin static visibility: all 6 affected SP15 cubins were already
pub static (collector imports work without visibility promotion).
Verification: cargo check clean (only pre-existing warnings); 7/7
sp14_oracle_tests pass; 36/36 sp15_phase1_oracle_tests pass.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Bug audit finding #2 (post train-d2b2s diagnostic — same Pattern 1
class as SP14 B.11 commit 200f05fce and the 4-producer batch in
5608b866b). SP5 Pearl 2 budget producer and SP7 loss-balance
controller currently launch from process_epoch_boundary (fires once
per epoch), but the loss-balance budget output (ISV[BUDGET_CQL_BASE..]
/ ISV[BUDGET_C51_BASE..]) is consumed EVERY training step via
apply_c51_budget_scale (fused_training.rs:1941) and the dispatch
kernel that resolves the cached value into lb_budget_effective_buf
(fused_training.rs:3631).
The dispatch kernel is per-step, but the underlying flatness signal
ISV[FLATNESS_BASE..] is per-epoch. SP7's controller therefore reads
(steps_per_epoch − 1)-step-stale flatness — the same failure mode
that broke SP14's ALPHA_GRAD_SMOOTHED. The entire loss-balance
budget system has been operating on stale-flatness state for the
duration of training.
Migration (atomic, preserves Pearl 2 → SP7 dependency):
- Pearl 2 budget launch moved to submit_aux_ops (just-after the
producer-cadence batch's MoE chain, just-before the IQL gather
block).
- SP7 loss-balance controller follows immediately (reads Pearl 2
output via ISV).
- Same captured-into-aux_child graph-replay semantics as SP14 B.11.
- training_loop.rs lines 4312, 4336 deleted; replaced with redirect
comment.
Per feedback_no_partial_refactor: this is the 7th cadence-fix in
this branch since v8ztm. Other Pattern 1 candidates (SP5 Pearl 1
atom, Pearl 3 σ — Pearl 2 inputs, AND SP8 Fix 36 launch_max_budget_compute
— SP7 controller input) deferred per scope-tightening rule. They
sit one-epoch-stale at fold start; Pearl A bootstrap + the
controller's internal cold-start branch keep behavior functional.
Tracked in the audit doc top-of-file entry; independent migrations,
land separately.
Verification: cargo check -p ml --lib clean (dev profile);
sp14_oracle_tests 7/7 pass on RTX 3050 Ti.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Mirrors commit 872bd7392 (aux next_bar stop-gradient) for the
5-class regime CE head. Same architectural conflict: regime backward
propagated dh_s2 SAXPY back to shared trunk h_s2, conflicting with
Q-loss for trunk representation control.
Today's next_bar fix landed with regime documented as a deferred
follow-up. Audit confirmed regime head has the IDENTICAL kernel
structure (same Linear→ELU→Linear→softmax→CE topology, same dh_s2
SAXPY at lines 732-742). Fix is mechanical — zero-fill the dh_s2
write block.
Combined with 872bd7392, this completes the trunk-isolation pair:
both auxiliary heads now train their own params from CE loss without
pulling on shared h_s2.
Verification: cargo check clean; sp14_oracle_tests 7/7 pass.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Root cause from train-v8ztm 9-epoch HEALTH_DIAG aux next_bar_mse trajectory:
- Ep 0: 0.352 (learnable signal — below random baseline ln(2)≈0.693)
- Ep 9: 0.717 (above random baseline — aux is now WORSE than random)
- aux_dir_acc_long stuck at 0.19 (anti-correlated with truth)
Aux head's backward gradient was flowing back to shared trunk activation
h_s2 via dh_s2_out write at aux_heads_kernel.cu:599-613. Q-loss
gradient on h_s2 dominates (larger magnitude, structurally different
objective: cumulative discounted reward vs next-bar direction). h_s2
evolves to support Q's task; aux's CE loss climbs as h_s2 features
become anti-aligned with direction prediction.
Fix: stop-gradient. Aux reads h_s2 via forward, trains its own w1/b1/w2/b2
from CE loss, but does NOT propagate to h_s2. Q-loss is the sole shaping
force on h_s2. Aux must adapt to whatever h_s2 happens to be — if the
representation has direction signal, aux's params will extract it; if
not, aux can't learn (separate-trunk Option 2 deferred for that case).
This was the SEVENTH fix in today's chain (after 6 SP14 EGF cadence/
gate/saturation fixes). The EGF was a scaffold over a broken aux head;
fixing aux first is the architectural prerequisite for EGF to route
useful signal.
Verification: cargo check clean; sp14_oracle_tests 7/7 pass.
Validation: aux next_bar_mse should now DECREASE during training in
the next L40S smoke (vs the rising-from-0.35-to-0.72 pattern in v8ztm).
Deferred follow-up: aux_regime_backward has the same architecture
(propagates dh_s2 to trunk). Same fix is a candidate once next_bar
result validates the approach.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Atomic step 5 of β migration. SP14 EGF producer chain now fires
ONLY from the experience collector (commit c691bd381).
Training-time launches removed; the collector-native chain is the
single production caller.
Trainer's pub(crate) launcher methods retained (no Rust dead-code
warnings on pub(crate) — atomic-rollback potential preserved).
Oracle tests in sp14_oracle_tests.rs exercise the SAME kernels
directly via load_cubin / load_function, not through these launchers
— deletion of the launchers would not affect oracle coverage. They
stay for the possibility that a future curriculum stage reverses
the rollout-only signal-quality assumption.
gradient_hack_detect (per-epoch circuit breaker) unchanged —
lockout-counter decrement is one-per-epoch by design.
Verification: cargo check clean; sp14_oracle_tests 2/2 non-GPU pass
(7 GPU tests ignored on RTX 3050 Ti host); L40S smoke validation
pending.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Step 4 of β migration: 3 SP14/SP13-EGF producers fire per-rollout-step
in collector using collector-owned kernel handles + collector stream.
Reads rollout-time q_values (post-expected-Q, pre-IQR/ensemble/noise)
+ exp_aux_nb_softmax. Writes to shared ISV.
Producer order preserved (matches trainer submit_aux_ops chain):
1. SP13 dir-acc reduce → 2 fixed-α EMAs → aux_pred to ISV[375]
2. SP14 q_disagreement_update (reads aux softmax + q_values)
3. SP14 alpha_grad_compute (pure ISV state machine)
Same kernel gate (commit 9d0c124ce) preserves EMAs across
no-contribution rollout steps.
q_logits semantic note: collector's q_values buffer is the
post-expected-Q output, BEFORE IQR/ensemble/noise SAXPY bonuses (those
run after this block). The kernel's argmax-over-K=4 finds Q's intended
direction; this matches the trainer's q_out_buf semantic exactly. If a
future audit shows noise-induced argmax flips matter, the launch site
is one indirection from the noise-free expected_q_kernel output.
Gated on isv_signals_dev_ptr != 0 && trainer_params_ptr != 0. No
seed_phase_active_cache gate — EGF Gate 1 needs to observe both
seed-phase scripted-policy and post-seed Q-policy actions across the
curriculum.
Compile clean; sp14_oracle_tests 2/2 non-GPU pass (7 GPU tests
ignored on RTX 3050 Ti host).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Step 3 of β migration: collector now runs aux_next_bar_forward on
rollout state every step. Label producer (new thin variant
aux_sign_label_per_step_kernel) derives sign(price[t+1] - price[t])
per env using bar = episode_starts[ep] + t. Aux predictions feed
the EGF kernel chain (step 4), NOT the Q-head's input (rollout
Q-head still sees raw h_s2, dir_qaux_concat_ptr remains 0u64).
Placement: AFTER captured forward graph, BEFORE expected_q kernel.
Same-stream serial ordering reads exp_h_s2_f32 populated by
forward_online_f32 inside the captured graph. Cold-start gated on
trainer_params_ptr != 0 to skip the test-scaffold path where the
trainer hasn't wired its params yet.
Files added: aux_sign_label_per_step_kernel.cu (66 lines).
Files modified: build.rs (+8 lines, register cubin),
gpu_experience_collector.rs (+106 lines: struct field, cubin static,
load in new(), per-step launch block).
Compile clean; sp14_oracle_tests 2/2 non-GPU pass.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Step 1 of β migration — Option B (collector-native): collector
loads its own CudaFunction handles for the 3 SP14 producer kernels
plus their sub-kernels (4 total: aux_dir_acc_reduce,
aux_pred_to_isv_tanh, q_disagreement_update, alpha_grad_compute).
Mirrors SP13 hold_rate pattern at gpu_experience_collector.rs:1820.
Cleaner than cross-component launcher calls (avoids trainer-stream /
collector-stream race; no signature surgery on the existing trainer
launchers).
Cubin static decls flipped to pub(crate) so the collector can
re-load on its own stream. Compile clean; sp14_oracle_tests pass
(2/2 non-GPU; GPU-gated 7 ignored on RTX 3050 Ti host).
No launches yet — additive infrastructure commit.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Root cause from train-6fcml 5-epoch trajectory (commit 5608b866b after
producer cadence migration): HEALTH_DIAG[0] (post-experience-collection)
showed q_dis_s=0.0595 q_dis_l=0.1329 var_q=0.00091 — meaningful rollout
signal. HEALTH_DIAG[1+] (post-training, per-step launches) all showed
q_dis_s=0.0000 q_dis_l=0.0000 var_q=0.00000 — signal decayed to zero
inside ONE epoch.
The kernel's ISV write block ran unconditionally even when total_cnt
(non-masked-row count after Hold/Flat masking) was 0. Empty-batch
launches blended `batch_mean = 0/1 = 0` into the EMA, decaying the
rollout signal to 0 over ~178 training steps × 0.7^n. Per-step training
launches read replay batches whose Q-direction picks are dominated by
Hold/Flat (the natural distribution); so total_cnt = 0 was the common
case, not a corner case.
Fix (atomic, single kernel):
- Wrap the ISV write block in `if (total_cnt > 0.0f) { ... }`. When the
training batch has no non-masked rows, the kernel is a no-op for that
step — EMAs stay at the prior step's values. Stream-ordered launches
still run; only the ISV write is skipped.
- Remove redundant `&& (total_cnt > 0.0f)` clause from the `is_first`
bootstrap check (now guaranteed by the outer gate).
Per pearl_first_observation_bootstrap semantics: "no observation"
preserves prior; only "first observation" replaces sentinel. Decay-on-
empty was inconsistent with both rules.
Other EGF-chain kernels audited:
- alpha_grad_compute_kernel.cu — operates on persistent ISV state,
no batch concept; var_aux/var_alpha Welford updates use `diff` of
persistent EMAs, not batch means. No empty-batch path. SAFE.
- aux_dir_acc_reduce_kernel.cu — emits out_6[0..3] with sentinel
fallback (0.5) when denom==0; downstream apply_fixed_alpha_ema then
blends 0.5 toward EMA. The sentinel is the random-baseline (target
threshold lies above it), so empty-batch pulls EMA toward harmless
baseline rather than zero. Different semantics from q_disagreement
(which has 0 — far below baseline 0.5). SAFE.
- gradient_hack_detect_kernel.cu — single-thread state machine on
persistent ISV, no batch. SAFE.
Verification:
- 6 existing sp14_oracle_tests pass.
- New q_disagreement_empty_batch_preserves_ema test asserts bit-exact
preservation of pre-seeded EMAs (0.0595, 0.1329, 0.0009 — the
train-6fcml HEALTH_DIAG[0] values) across an all-Hold batch. Catches
the regression the existing all-hold test missed (its bound
`[0.0, 0.5]` accepted both decay-to-blend and preserve-prior; the new
test is strict bit-equality).
- L40S smoke validation pending — train-6fcml symptoms (alpha_smoothed
stuck at 0.0002, gate1 closed forever) expected to resolve.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Continuation of SP14 B.11 cadence fix (commit 200f05fce). The same per-
epoch staleness bug affected 4 additional producers whose docstrings
explicitly claim "per-step" cadence but whose launches lived in
`process_epoch_boundary` (fires once per epoch at line ~780):
* launch_h_s2_rms_ema → ISV[H_S2_RMS_EMA_INDEX=96]
* launch_fold_warmup_factor → ISV[FOLD_WARMUP_FACTOR_INDEX=130]
* launch_moe_expert_util_ema → ISV[MOE_EXPERT_UTIL_EMA_BASE..+8) +
ISV[MOE_GATE_ENTROPY_EMA_INDEX=126]
* launch_moe_lambda_eff_update → ISV[MOE_LAMBDA_EFF_INDEX=128]
Each producer's per-step consumer reads the ISV slot every training
step (thousands per epoch). With the producers stuck per-epoch, consumers
saw (steps_per_epoch − 1)-step-stale values — exactly the failure mode
SP14 B.11 documented for ALPHA_GRAD_SMOOTHED. Verified consumers:
* h_s2_rms_ema → mag_concat_kernel (forward graph) +
dqn_clamp_finite_f32_kernel (cuBLAS backward IQN-trunk sanitiser
bound = 1e6 × ISV[96])
* fold_warmup_factor → training_loop.rs:2662 per-step host read
deriving lr_eff and clip_eff
* moe_expert_util/entropy → consumed by launch_moe_lambda_eff_update
immediately below (same step)
* moe_lambda_eff → moe_load_balance_loss kernel reads ISV[128] per
step at gpu_dqn_trainer.rs:16864
Atomic migration per `feedback_no_partial_refactor`. Producers leaving
process_epoch_boundary are replaced with one-line redirect comments
pointing to the new hook. Ordering dependencies preserved
(launch_moe_lambda_eff_update still runs AFTER launch_moe_expert_util_ema
per the kernel docstring at gpu_dqn_trainer.rs:16962). All 4 launchers
use pre-loaded CudaFunction fields (gpu_dqn_trainer.rs:12646, 30983,
16901, 16969) — graph-capture safe per
`pearl_no_host_branches_in_captured_graph`.
Cold-start ordering: h_s2_rms_ema reads `save_h_s2` populated by THIS
step's online forward (forward_child runs BEFORE aux_child per
`capture_training_graph`). MoE producer reads `moe_gate_softmax_buf`
populated by `launch_moe_forward` (in `submit_forward_ops_main`,
forward_child). Both inputs are valid by the time submit_aux_ops runs.
State-reset registry coverage already in place
(state_reset_registry.rs:304/427/380/387/398/442 — isv_h_s2_rms_ema,
isv_fold_warmup_factor, isv_moe_expert_util_ema, isv_moe_gate_entropy_ema,
isv_moe_lambda_eff, isv_grad_norm_fast_ema). HEALTH_DIAG read points
unchanged — per-epoch reads of per-step-updated slots get the latest
value (strict improvement over per-epoch reads of per-epoch-stale slots).
Producers DELIBERATELY KEPT per-epoch (audit trail):
Collector EMAs (input data only updates per-epoch — running per-step
computes the same EMA repeatedly off the same data):
* launch_trade_attempt_rate_ema_inplace
* launch_plan_threshold_update_inplace
* launch_seed_step_counter_update_inplace
* launch_cql_alpha_seed_update_inplace
HEALTH_DIAG-only consumers (no per-step kernel reads these slots):
* launch_iqn_quantile_ema (ISV[99..103))
* launch_vsn_mask_ema (ISV[105..111))
* launch_aux_heads_loss_ema (ISV[113], ISV[114])
DEFERRED follow-up (NOT changed in this commit per the user-specified
strict heuristic "docstring per-step + per-step consumer ⇒ migrate";
docstrings of the SP4/SP5 producers below don't explicitly claim
per-step cadence — they're silent — so they fall outside the heuristic
even though their consumers ARE per-step):
* launch_sp4_target_q_p99 + atom_pos_p99 + grad_norm_p99 + h_s2_p99
* launch_sp4_param_group_oracles_all_groups
(consumed by `read_group_adam_bounds` per-step in submit_aux_ops)
* launch_sp5_pearl_1_atom + pearl_3_sigma + pearl_2_budget +
pearl_4_adam_hparams + pearl_5_iqn_tau + pearl_8_trail +
pearl_1_ext_num_atoms
* launch_max_budget_compute + launch_loss_balance_controller (SP7/SP8)
This staleness will be addressed in a follow-up audit pass; for now the
SP4/SP5 epoch-block producers stay where they are.
Verification:
* SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --tests
--all-targets clean (only pre-existing unused-var warnings).
* sp14_oracle_tests 6/6 pass (alpha_grad_adaptive_beta,
alpha_grad_schmitt_hysteresis, dir_concat_qaux_correct,
gradient_hack_circuit_breaker_fires,
q_disagreement_all_hold_no_contribution, q_disagreement_k4_k2_mapping).
* HEALTH_DIAG validation pending L40S re-dispatch.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Root cause from train-v8ztm 10-ep validation (commit 1396b62ec): HEALTH_DIAG
showed alpha_smoothed=0.0002 (vs ~0.5 expected steady-state), gate1=closed,
var_aux:var_q ratio 290:1 — symptoms of an EGF producer chain firing < 1%
as often as the consumer.
The original B.11 wire-up (commit 857722e77) placed `launch_sp14_q_disagreement_update`,
`launch_sp14_alpha_grad_compute`, and the prerequisite `launch_sp13_aux_dir_metrics`
in `process_epoch_boundary` — which runs ONCE per epoch (single call site at
training_loop.rs:780, called from the per-epoch loop, not from the per-step
loop in `run_training_steps_slices`). The captured backward consumer
`launch_sp14_scale_wire_col` (inside launch_cublas_backward_to, replays every
training step via parent graph) reads ISV[ALPHA_GRAD_SMOOTHED=393] per step,
but the producer was firing only at epoch boundary — every step inside the
epoch observed (steps_per_epoch − 1)-step-stale alpha values, with the EMA
chain barely accumulating past sentinel between rare per-epoch updates. The
plan §2550 explicitly specifies per-step cadence; the existing wire violated
the plan.
Fix (atomic, graph-capture-safe):
- MOVED launch_sp13_aux_dir_metrics, launch_sp14_q_disagreement_update,
launch_sp14_alpha_grad_compute from process_epoch_boundary into
fused_training.rs:submit_aux_ops, immediately after populate_q_out.
submit_aux_ops captures into the aux_child sub-graph, so each parent-graph
replay re-fires the full producer chain — restoring per-step cadence.
- launch_sp13_aux_dir_metrics had to migrate alongside the SP14 launches:
alpha_grad_compute_kernel consumes its outputs (ISV[373/374]); leaving
sp13 per-epoch while moving SP14 per-step would re-introduce the same
staleness bug for aux_dir_acc reads (atomic dependency migration per
feedback_no_partial_refactor).
- Per-epoch launch_sp14_gradient_hack_detect circuit breaker stays in
process_epoch_boundary — its lockout decrement IS one-per-epoch by design.
- Forward consumers (6 launch_sp14_dir_concat_qaux sites) and backward
consumers (2 launch_sp14_scale_wire_col sites) unchanged — they read the
same ISV[393], but now see live per-step values instead of per-epoch
staleness.
Verification:
- cargo check -p ml --tests --all-targets clean (no errors, no new warnings).
- All 6 SP14 oracle GPU tests pass (alpha_grad_adaptive_beta,
alpha_grad_schmitt_hysteresis, dir_concat_qaux_correct,
gradient_hack_circuit_breaker_fires, q_disagreement_all_hold_no_contribution,
q_disagreement_k4_k2_mapping).
- HEALTH_DIAG validation pending L40S re-dispatch — expect alpha_smoothed
to track real EGF-driven values (~0.5 in steady state).
Invariants:
- pearl_no_host_branches_in_captured_graph (kernels are pure GPU state
machines using launch_builder + pre-loaded CudaFunction; no per-call
load_cubin)
- feedback_no_partial_refactor (sp13 + 2 SP14 launches migrated atomically)
- feedback_wire_everything_up (all 3 producers now production hot-path,
re-fire on every parent-graph replay)
- feedback_isv_for_adaptive_bounds (no warmup_gate parameter — variance-
driven k_aux/k_q in alpha_grad_compute_kernel handles cold-start
adaptively, per c0fc28e45)
Refs: train-v8ztm trajectory analysis 2026-05-07T15:59:49 HEALTH_DIAG[10]
showed dir_entropy=0.6545 kill-fast breach with model converging to 64%
Hold + 84% Quarter magnitude — exactly the pathology B.11 was designed
to prevent by routing aux's directional signal into Q.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Root cause: 5 SP15 evaluation launchers (cost_net_sharpe + 4 baseline_*)
were doing `load_cubin` + `load_function` PER-CALL inside
`GpuBacktestEvaluator`'s eval hot loop. Pattern is fragile across CUDA
context lifetimes — works in single-pass train-best context (smoke
train-9bcwm verified), fails in hyperopt-trial child stream context
(workflow train-xggfc trial 1 failed at "load sp15_baseline_kernels
cubin: ILLEGAL_ADDRESS"; after the host-side load corrupted the trial's
context, trials 2-20 all cascade-failed at "Fork CUDA stream for trial").
Fix (atomic, matches 5d63762ab precedent for bn_tanh_concat_dd):
- Add 5 `CudaFunction` fields on `GpuBacktestEvaluator`.
- Pre-load both `SP15_BASELINE_KERNELS_CUBIN` (4 functions) and
`SP15_COST_NET_SHARPE_CUBIN` (1 function) once in
`GpuBacktestEvaluator::new()`, alongside the existing `env_module` /
`metrics_module` loads.
- Change all 5 launcher signatures in `gpu_dqn_trainer.rs` to take
`&CudaFunction` instead of doing per-call cubin load.
- Update all 5 call sites in `gpu_backtest_evaluator.rs:2877..2922` to
pass `&self.sp15_*_kernel`.
- Update 4 oracle test call sites in
`crates/ml/tests/sp15_phase1_oracle_tests.rs` (cost_net + 4 baselines)
to pre-load and pass the kernel handle directly, matching 5d63762ab's
pattern for `DQN_UTILITY_CUBIN`.
- Update Invariant 7 audit doc (`docs/dqn-wire-up-audit.md`).
Verification: `SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml
--tests --all-targets` clean (warnings only, no errors).
Refs: train-xggfc failure 2026-05-07T13:11:43, 5d63762ab precedent,
pearl_no_host_branches_in_captured_graph (this is the eval analogue),
feedback_no_partial_refactor, feedback_wire_everything_up.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Root cause: SP15 Wave 4.1b (eb9515e41) migrated the trainer's 3 production
bottleneck-concat call sites (online forward, target forward, DDQN argmax)
from the pre-loaded `bn_tanh_concat_kernel: CudaFunction` field to a
`launch_sp15_bn_concat_dd` free function that resolved the symbol via
`load_cubin` + `load_function` ON EVERY CALL. That pattern is incompatible
with CUDA Graph capture: `cuModuleLoadData` and `cuModuleGetFunction` are
HOST-side driver API calls that allocate memory and mutate driver state,
and they are NOT capturable inside a `cuStreamBeginCapture` region. Calling
the launcher from `submit_forward_ops_main` (graph-captured forward child)
caused a SEGV at exit 139 — the loader raced with the capture-mode driver
state, the host-side corruption surfaced as a segfault before
`CAPTURE_PHASE_FORWARD_DONE` could print.
Diagnostic evidence (L40S smoke `train-vg5f7` on commit `bfc3ffa9d`):
ALL 16 step-0 ungraphed checkpoints printed clean. Capture begins:
CAPTURE_PHASE_BEGIN
CAPTURE_PHASE_PER_SAMPLE_DONE
CAPTURE_PHASE_COUNTERS_DONE
CAPTURE_PHASE_SPECTRAL_DONE
Then: SEGV. The next checkpoint that didn't print was
CAPTURE_PHASE_FORWARD_DONE -> SEGV is inside `submit_forward_ops_main`'s
`forward` child capture. The same function ran cleanly ungraphed in
step 0 because the loader-host-branch was harmless without active
capture.
The experience collector's equivalent caller (gpu_experience_collector.rs:4127)
was already doing this correctly — pre-loaded `bn_tanh_concat_fn` field
populated at construction. Wave 4.1b's mistake was asymmetric: it kept the
collector's pre-load pattern but introduced an on-demand loader for the
trainer's call sites.
Fix (atomic, restores graph-safe contract):
- Add `bn_tanh_concat_dd_kernel: CudaFunction` field on `GpuDqnTrainer`
(back-fills what Wave 4.1b removed, with updated docstring naming the
new dd_pct-aware kernel).
- Pre-load the symbol in `compile_training_kernels` from the same `module`
as `bn_tanh_bw` / `bn_bias_grad` (forward-child captured replay group).
Tuple grows 43 -> 44 CudaFunctions; struct ctor wires the field.
- Change `launch_sp15_bn_concat_dd` signature to take `&CudaFunction` as
parameter; remove the per-call `load_cubin` + `load_function`. All 3
trainer call sites pass `&self.bn_tanh_concat_dd_kernel`.
- Promote `DQN_UTILITY_CUBIN` from `pub(crate)` to `pub` so the oracle
tests in `crates/ml/tests/sp15_phase1_oracle_tests.rs` can pre-load
the kernel handle (the launcher no longer hides this for them).
- Update the 2 test call sites (kernel-level oracle + Wave 4.1c behavioral
KL test) to load the kernel handle once up-front and pass it through.
Verification:
- `cargo check -p ml --tests` clean (release + dev profile).
- `cargo test -p ml --test sp15_phase1_oracle_tests -- --ignored`:
36/36 pass, including the Wave 4.1c behavioral KL test that exercises
two end-to-end forward passes through the modified launcher.
Refs: SP15 Wave 4.1b (eb9515e41), pearl_no_host_branches_in_captured_graph,
feedback_no_partial_refactor, feedback_wire_everything_up.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Smoke train-xq9hg got past all 13 original checkpoints (BEGIN through
NAN_CHECKS_DONE) cleanly. SEGV is downstream in Phase 6/7/8.
Adds 16 more checkpoints: TLOB_BWD_ADAM, MAMBA2_BWD, MAMBA2_ADAM,
OFI_EMBED_BWD, OFI_EMBED_ADAM, PRUNING, BRANCH_GRAD_BALANCE, GRAD_NORM,
ADAM_OPS, Q_MAG_BIN, Q_DIR_BIN, ISV_UPDATE, MAINTENANCE, IQL_MODULATE,
PER_PRIORITY, CAPTURE.
Diagnostic-only — no behavior change.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
L40S smoke train-jfbzr (commit 23e9a1f78) segfaults at exit 139 in the
first run_full_step invocation, after "GPU training guard initialized
(epoch loop)" and before any HEALTH_DIAG output. Static analysis
debugger couldn't reproduce locally (test_data/futures-baseline/ lacks
.fxcache).
This commit adds eprintln! checkpoints at each phase boundary in the
ungraphed step-0 path of FusedTrainer::run_full_step (PER sample,
PopArt, counters, spectral norm, TLOB forward, forward_main, DDQN,
aux_ops, post_aux, NaN checks).
stderr is line-flushed (vs stdout buffered through tracing JSON
formatter), so the last printed checkpoint identifies the SEGV site.
Each fires once per fold's first step (graph capture absorbs subsequent
steps), so log overhead is minimal.
Diagnostic-only commit — no behavior change. Will be reverted after the
next L40S smoke localizes the SEGV and the root-cause fix lands.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Pre-existing latent bug surfaced by compute-sanitizer after the SP15
NULL-pointer fix at 2e37af29d unblocked the cascade. Threads 60-63
were writing past denoise_target_q_buf's end every batch.
Sanitizer evidence (pre-fix on 2e37af29d):
Invalid __global__ write of size 4 bytes
at compute_expected_q+0x2480
by thread (60..63, 0, 0) in block (0, 0, 0)
Access at <addr> is out of bounds (45/97/149/201 bytes after
nearest allocation of size B*12*4 bytes)
Host backtrace: GpuDqnTrainer::compute_denoise_target_q
→ submit_post_aux_ops → run_full_step
Root cause — semantic mismatch between writer stride and buffer size:
- compute_expected_q writes q_values[i*total_actions + a] with
total_actions = b0+b1+b2+b3 = 4+3+3+3 = 13 (4-direction factored)
- denoise_target_q_buf was allocated b*12 (legacy 3+3+3+3 layout)
- threads 60..63 wrote slot 12 of samples (60..63 % batch_size)
past the buffer end every step
The downstream q_denoise_backward + denoise_loss_grad kernels read
Q_target[b * D + i] with hardcoded D=12 because the diffusion denoiser
MLP itself only has 12 output slots (W2[12,24] + b2[12]) — its 12 are
the q_coord_buf's post-cross-branch-attention narrowed output, NOT a
clean subset of the 13 raw Q-actions.
The exact same fix pattern was already applied to the sibling
q_var_buf_trainer allocation in the prior SP4 audit (see comment block
at gpu_dqn_trainer.rs:21620-21630 referencing the identical OOB at
threads 60..63); denoise_target_q_buf 8 lines below was missed because
it was guarded by the SP15 NULL-pointer ILLEGAL_ADDRESS that fault-
stopped the cascade before this OOB could fire — 2e37af29d removed the
upstream ILLEGAL_ADDRESS, surfacing the latent OOB.
Fix architecture (Option A — pad buffer to total_actions, pass stride
to consumer; same pattern as q_var_buf_trainer):
1. Widen denoise_target_q_buf from b*12 to b*total_actions (= b*13)
to match compute_expected_q's writer stride.
2. Add refined_stride / target_stride / input_stride parameters to
q_denoise_backward kernel; the kernel still computes D=12 per-
sample (denoiser MLP fixed width) but addresses each input buffer
at its own per-sample stride.
3. Add refined_stride / target_stride parameters to denoise_loss_grad
kernel (same pattern; used by launch_q_denoise_backward_cublas).
4. Update both Rust launchers (launch_q_denoise_backward,
launch_q_denoise_backward_cublas) to pass refined_stride=12 (q_coord
and q_input are post-attention narrowed),
target_stride=total_actions=13.
5. denoise_q_input_buf STAYS at b*12 — it's a snapshot of q_coord_buf
(also b*12) via snapshot_pre_denoise_q's DtoD copy; never written
by compute_expected_q.
6. Flat-scan consumers (SP4 target_q_p99_update producer + SP3 slot
46 threshold-check + dqn_clamp_finite_f32) UNCHANGED — they
consume the buffer flat via .len(); widening from b*12 to b*13 is
monotone (one more valid Q-value per sample in the histogram).
Atomic per feedback_no_partial_refactor: 2 kernel signatures + buffer
allocation + 2 launch sites + struct doc comments + SP3 slot 46 doc +
audit doc — all in one commit. Every consumer of the writer's stride
migrates simultaneously.
Verification:
- SQLX_OFFLINE=true cargo check -p ml --features cuda: clean
- compute-sanitizer (RTX 3050 Ti): 0 Invalid __global__ errors at
compute_expected_q post-fix (down from 4 per training step in
baseline — re-verified on 2e37af29d to confirm the diff)
- SQLX_OFFLINE=true cargo test -p ml --features cuda --lib: holds
parent baseline 947 pass / 12 fail (same 12 pre-existing failures)
Files touched:
crates/ml/src/cuda_pipeline/experience_kernels.cu
— q_denoise_backward + denoise_loss_grad kernel signatures
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
— alloc widen + 2 launcher updates + 4 doc comment updates
docs/dqn-wire-up-audit.md
— audit entry
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Root cause: Wave 4.1b (s1_input_dim 102→103, eb9515e41) introduced
`bn_tanh_concat_dd_kernel` which reads `isv[DD_PCT_INDEX=406]` from the
collector's `isv_signals_dev_ptr`. Wave 4.1b's atomic refactor docstring
claimed "guarded by the captured-graph wiring in training_loop.rs:859
which sets the bus pointer BEFORE the first epoch's forward graph
capture" — but ordering audit shows that's wrong:
Phase 1c (line 580): init_gpu_experience_collector
→ collector.isv_signals_dev_ptr = 0 (NULL)
Phase 2 (line 730): collect_gpu_experiences_slices
→ bn_tanh_concat_dd_kernel reads isv[406]
= NULL+1624 = 0x658 → ILLEGAL_ADDRESS
Phase 4 (line 859+): set_isv_signals_ptr (TOO LATE)
The error surfaces as "load sp15_dd_state cubin: ILLEGAL_ADDRESS" via
sticky-cascade (the OOB happened in the prior kernel; the next CUDA
call — dd_state cubin load — surfaces the cascade).
Reproduced locally via compute-sanitizer on
test_no_hang_single_epoch (RTX 3050 Ti). Sanitizer pinpointed
"Invalid __global__ read of size 4 bytes at bn_tanh_concat_dd_kernel
+0x4d0 ... Access at 0x658 is out of bounds" — exactly slot 406 * 4.
Fix: move the ISV / SP15 control pointer wiring (set_isv_signals_ptr,
set_sp15_alpha_warm_count_ptr, set_sp15_plasticity_target, PER
buffer's set_isv_signals_ptr + set_sp15_dd_trajectory_per_env_ptr +
set_sp15_per_env_dims) into init_gpu_experience_collector BEFORE the
collector is moved into self.gpu_experience_collector. This guarantees
the FIRST collect_experiences_gpu call (epoch 0, Phase 2) sees
non-NULL pointers — required because cudarc's graph capture bakes the
arg values into the captured graph at capture time, so subsequent
re-wiring has no effect on the captured replay path.
The per-epoch re-wiring at lines 855-957 of the outer epoch loop is
now redundant (idempotent setters re-applying the same stable
pointers) but kept in place per feedback_no_partial_refactor for
explicit per-epoch refresh semantics — these underlying pointers are
stable for the trainer's lifetime so re-setting is a no-op, and
removing the per-epoch call would scatter the wire-up logic across a
non-obvious pre-init / per-epoch split.
Verification:
* 30/30 SP15 phase 1 oracle tests pass under compute-sanitizer
memcheck with 0 errors (unchanged baseline)
* test_no_hang_single_epoch under sanitizer: SP15 dd_state OOB is
GONE; the first OOB now surfaces at compute_expected_q+0x2480
(denoise_target_q_buf stride 12 vs total_actions=13 — a SEPARATE
pre-existing bug previously masked by the SP15 OOB; reported but
out of scope for this commit per single-atomic-fix discipline).
* compute-sanitizer "Access at 0x658" matches the
bn_tanh_concat_dd_kernel isv[406] read path; production training
on L40S/H100 should now reach the next-layer issue.
Refs: SP15 Wave 4.1b (eb9515e41), feedback_no_partial_refactor,
feedback_wire_everything_up, pearl_no_host_branches_in_captured_graph.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Wave 3b (4320820ae) added a 6th window_lob_bars parameter to
GpuBacktestEvaluator::new and migrated 3 production lib call sites + 6
test call sites, but the lib-only `cargo check -p ml --features cuda`
validation didn't catch the 3 example-binary call sites in
evaluate_baseline.rs (lines 1344, 1641, 1802). Argo's ensure-binary
step compiles all 4 example binaries (hyperopt_baseline_rl,
train_baseline_rl, evaluate_baseline, precompute_features), and
evaluate_baseline failed with E0061 (wrong number of args).
This commit migrates all 3 sites to construct zero-OFI LobBar SoA
inline (eval-path lacks per-bar OFI features — same degradation
pattern as the PPO hyperopt adapter Wave 3b D4 resolution; OFI-impact
term degrades to 0 while commission + half-spread × position still
apply).
Validated: `cargo check -p ml --features cuda --all-targets` clean
(was --features cuda only before — now expanded to catch
example/test/bin compile-breaks).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Phase 1.3.b-followup (5b394f103) landed the main per-env redesign but
deferred dd_trajectory + per_insert_pa migration as Followup-B because
the contract shape is per-(env, t) buffer [N*L], not per-env tile.
This commit:
- Reshapes dd_trajectory_decreasing_kernel to per-env grid
[n_envs, 1, 1] x [1, 1, 1]; writes dd_trajectory_per_env[env]
- Migrates per_insert_pa to per-transition lookup via
env_id = (j % (n_envs * lookback)) / lookback; reads
dd_trajectory_per_env[env_id]
- Allocates two collector-owned per-env tiles
(sp15_dd_trajectory_per_env + sp15_dd_trajectory_prev_dd_per_env);
deletes the trainer-owned [1] sp15_dd_trajectory_prev_dd scratch +
its setter wiring (replaced by unconditional collector ownership,
mirrors sp15_dd_state_per_env pattern)
- Extends dd_state_reduce_kernel to mean-aggregate per-env
trajectory tile -> ISV[DD_TRAJECTORY_DECREASING_INDEX=439] for
HEALTH_DIAG diagnostic preservation
- Wires per-env tile dev_ptr + (n_envs, lookback) dims into
GpuReplayBuffer via new setters (set_sp15_dd_trajectory_per_env_ptr,
set_sp15_per_env_dims) called from training_loop
- Migrates 4 dd_trajectory + 1 PER oracle tests to per-env contract
- Adds NEW behavioral test
per_sampler_weights_per_transition_not_uniform_across_batch:
n_envs=2 batch with env-0 trajectory=1, env-1 trajectory=0; asserts
priorities[env-0 slots]=3.0 and priorities[env-1 slots]=1.0 in the
SAME insert batch (Wave 4.3 would produce uniform 3.0 OR uniform
1.0 across all 8 priorities — the test directly fails the old
implementation)
- Layout fingerprint marker DD_TRAJECTORY_PER_ENV=sp15_phase_1_3_b_followup_B
(greenfield checkpoints OK per spec Q1)
Fixes the Wave 4.3 PER limitation: ISV[DD_TRAJECTORY_DECREASING_INDEX=
439] was read ONCE per insert call and applied uniformly to ALL
n_envs * lookback * 2 transitions in the batch, so the recovery
oversample was statistically biased (whichever env wrote ISV[439]
most recently determined the boost for ALL inserted transitions).
Per-transition lookup fixes this — each transition's weight reflects
its own env's recovery context.
Atomic per feedback_no_partial_refactor: kernel reshape + 2 collector-
owned per-env tiles + trainer struct cleanup + reduction kernel
extension + per_insert_pa kernel signature change + 3 new GPU PER
replay buffer fields/setters/accessors + per_insert_pa launch site
update + collector launch sequence update + state-reset registry
rename (1->2 entries) + 2 dispatch arms + training_loop wiring
delete/replace + 4 dd_trajectory test migrations + 1 PER test
migration + 1 NEW behavioral test + 2 dd_state_reduce callsite
null updates + audit doc all in this commit.
Closes the per-env DD redesign chain end-to-end. SP15 reward shaping
+ recovery-curriculum PER oversample now correctly apply per-env
context everywhere they're consumed; no more "whichever env wrote
last wins" paths.
Verified: all dd_state, dd_trajectory, per_sampler, final_reward,
plasticity oracle tests green; ml lib HOLDS the Phase 1.3.b-followup
baseline (947 pass / 12 fail) on RTX 3050 Ti — no new regressions.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Closes the Phase 1.3.b deferred per-env redesign per
feedback_no_partial_refactor. Path A (env-0-canonical, commit 132609724)
made every downstream consumer read env-0's DD context; production envs
each have their own DD trajectory, so when env-3 was at 30% DD and
env-0 was at ATH, the model learning from env-3's transitions saw
dd_pct=0 and silently skipped the recovery shaping. Path B threads
each env's actual DD context through the reward shaping atomically:
(1) dd_state_kernel reshape — grid [n_envs, 1, 1], one thread per env,
writes 6 scalars per env to a new per-env tile dd_state_per_env
[n_envs * 6]. No more ISV scalar writes.
(2) NEW dd_state_reduce_kernel — single-block tree-reduce (no
atomicAdd; BLOCK=256 with strided initial pass for n_envs up to
32768 on H100). Mean-aggregates per-env tile → 6 scalar ISV slots
[401..407) for HEALTH_DIAG diagnostic. Max-aggregates DD_PERSISTENCE
→ new slot DD_PERSISTENCE_MAX_INDEX=443 for plasticity injection
trigger (one shared advantage-head ⇒ global firing ⇒ max-aggregate
is the only correct rule). ISV_TOTAL_DIM 443→444; SP15_SLOT_END
443→444; SP15_SLOT_COUNT 46→47.
(3) compute_sp15_final_reward_kernel migration — per-(i,t) per-env DD
lookup via env_id = (idx % (N*L)) / L. Helpers sp15_dd_asymmetric_reward
and sp15_dd_penalty migrated to take dd_pct + dd_current as scalar
parameters (the kernel reads from the per-env tile, threads scalars
in). On-policy + CF threads at the same (i,t) read the SAME tile
entry (one DD trajectory per env, shared across slot kinds).
(4) plasticity_injection_kernel migration — persistence read switched
from ISV[404] (mean) to ISV[443] (max). One set of advantage
weights ⇒ ANY env exceeding the threshold should arm the gate.
(5) Per-env tile owned by GpuExperienceCollector (not the trainer) —
the collector knows alloc_episodes (= n_envs); the trainer's
batch_size is a different quantity. Reset to zero via the
sp15_dd_state_per_env registry-arm dispatch.
(6) Per-step launch order: dd_state → dd_state_reduce →
alpha_split_producer → final_reward, all on the same stream
(CUDA serialises producer→consumer without explicit event sync).
(7) HEALTH_DIAG semantic shift (documented breaking change): slots
401-406 now report cross-env mean, not env-0 value. For n_envs=1
smoke configs the mean equals env-0's value (bit-stable migration).
(8) Layout fingerprint break: added markers DD_PERSISTENCE_MAX=443;
ISV_TOTAL_DIM=444; DD_STATE_PER_ENV=sp15_phase_1_3_b_followup.
Pre-followup checkpoints will not load (greenfield OK per spec Q1).
(9) 6 oracle tests migrated + 1 NEW behavioral test
`dd_state_per_env_diverge_independently` — two-env config (env-0
in recovery, env-1 deepening) verifies independent trajectories
+ mean-aggregate + max-aggregate semantics.
Phase 1.3.b-followup-B (separate split): dd_trajectory_decreasing_kernel
+ per_insert_pa migration is NOT in scope. The current Wave 4.3 reads
ISV[439] at insert-batch time (epoch end) — applied uniformly to ALL
inserted transitions (a known pre-existing limitation). Proper fix
requires per-(env, t) trajectory buffer [N*L] + env_id-aware lookup
in per_insert_pa via env_id = (j % (N*L)) / L. That's a different
contract change; splitting preserves no-partial-refactor within each
migration.
Atomic per feedback_no_partial_refactor: all 5 consumers of single-env-
canonical DD slots (final_reward kernel + plasticity kernel +
HEALTH_DIAG diagnostic + 6 oracle tests + new behavioral test) migrate
to per-env tile lookup in this commit; the new DD_PERSISTENCE_MAX
ISV slot lands with its sole consumer (plasticity).
Verified: SQLX_OFFLINE=true cargo check -p ml --features cuda clean;
cargo check -p ml --features cuda --tests clean;
CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests
--features cuda -- --ignored: 17/17 oracle tests green (2 dd_state
incl. new per-env behavioral + 5 plasticity + 3 dd_trajectory + 6
final_reward + 1 per_sampler); cargo test -p ml --features cuda --lib:
947 pass / 12 fail HOLDS the 483cef454 baseline (no new regressions).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Wave 4.2 (ef08611d3) landed cuRAND Kaiming-He weight reset as orphan-
with-tests-only. Phase 3.5.4.c creates the production caller and the
action-selection consumer atomically per spec lines 4346-4448.
Production caller in gpu_experience_collector.rs step 4-pre (BEFORE
experience_action_select):
- fan_in = cfg.adv_h = 128 (CORRECTED from Wave 4.2's audit-doc
error claiming 6528 = adv_h × num_atoms; that's total weight
count of one branch's projection, not per-unit input dim. Per
feedback_trust_code_not_docs the audit-doc error is also fixed
inline in this commit.)
- n_weights = branch_0_size × num_atoms × adv_h = 4 × 51 × 128 =
26112 (default config). Resets last 10% = 2611 directional
advantage-head weights via cuRAND curand_normal × sqrt(2/128) ≈
0.125 stddev (was wrongly documented as 0.0175).
- Branch: w_b0out (directional) only — plasticity is about
escaping stuck-in-flat regimes; targeted at directional choice.
- Seed: mix_seed(Phase-3.5.4.c-unique base) ^ t per-step
deterministic source; same (FOXHUNT_SEED, t) always yields the
same Kaiming-He samples per kernel thread.
- Trainer exposes target via new pub fn sp15_w_b0out_target()
returning (dev_ptr, n_weights, fan_in); collector consumes via
new set_sp15_plasticity_target(); training_loop wires the two.
OR-gate consumer in experience_action_select:
- New kernel arg float plasticity_m_warm threaded into the
cooldown-mask code path.
- Wave 1.B's cooldown_active = (cooldown_remaining > 0) extended
to cooldown_active = (cooldown_remaining > 0) ||
(plasticity_warm_remaining > 0).
- Flat-on-fire-bar detection — option (a), warm ≥ m_warm − 1.5f
per the trigger-then-decrement convention. The kernel sees
warm = m_warm − 1 on the fire bar; subsequent warm-up bars
observe warm ≤ m_warm − 2. New if (plasticity_fire_bar) branch
BEFORE the existing else if (cooldown_active) branch — the
fire-bar gets DIR_FLAT, subsequent warm bars get DIR_HOLD via
the OR-gate cooldown.
Two-step recovery semantics:
- Bar T (fire): plasticity launches → ISV[436]=1, ISV[438]=
m_warm then decrements to m_warm−1. action_select detects
fire-bar → dir_idx = DIR_FLAT.
- Bars T+1..T+M_warm−1 (warm): action_select OR-gate forces
dir_idx = DIR_HOLD.
- Bar T+M_warm: warm transitions to 0, OR-gate inactive, normal
Thompson/argmax resumes.
When the production caller is unwired (test scaffold path):
plasticity_m_warm passes 0.0f, the kernel's fire-bar predicate is
dead (warm == 0 too), and the OR-gate degenerates to cooldown-only
— bit-identical to the pre-3.5.4.c behaviour. Existing 3 SP15
3.5.b/3.5.3.b action_select oracle tests pass unchanged at
m_warm = 0.0f.
New behavioral oracle test plasticity_fires_force_flat_then_cooldown_
holds verifies the full sequence on RTX 3050 Ti with M_warm = 5
(shrunk from production's 200 for test runtime): bar 0 → DIR_FLAT,
bars 1-3 → DIR_HOLD, bar 4 (warm boundary) → DIR_LONG. ISV side-
effects (fired flips 0→1 then debounces, warm decrements with
underflow guard) verified bar-by-bar.
Eval/backtest path passes m_warm = 0.0f because plasticity is a
training-only mechanism (during deterministic eval the weights must
remain frozen at their checkpoint values).
Atomic per feedback_no_partial_refactor: kernel + caller + consumer +
trainer plumbing + 4 launcher-call-site updates (1 production
collector + 1 eval-path + 2 test scaffolds) + new behavioral test +
audit-doc fan_in inline correction + new audit-doc entry all in this
commit. No fallback. SP15 Phase 3.5 recovery-dynamics chain is now
end-to-end production-wired (3.5.2 + 3.5.3 + 3.5.4 + 3.5.4.c +
3.5.5 + 3.5.5.b all firing).
Verified: cargo check -p ml --features cuda clean; cargo check -p ml
--features cuda --tests clean; CUDA_COMPUTE_CAP=86 cargo test -p ml
--test sp15_phase1_oracle_tests --features cuda -- --ignored
--nocapture 34 of 34 SP15 oracle tests green (5 plasticity + 3
action_select + 3 cooldown + 23 others); cargo test -p ml --features
cuda --lib HOLDS the Wave 4.3 baseline (946 pass / 13 fail) on RTX
3050 Ti — no new regressions.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Phase 3.5.5 (69b8fdb61) landed dd_trajectory_kernel writing
ISV[DD_TRAJECTORY_DECREASING_INDEX=439] per-step but DEFERRED the PER
sampler integration AND the production launch. Wave 4.3 wires both
atomically per feedback_wire_everything_up.
Investigation finding — Case B (existing TD-error-driven priority
sampler): the GPU PER architecture in
crates/ml-dqn/src/gpu_replay_buffer.rs already weights replay by
per-transition priority (raw priority + priorities_pa = priority^alpha
feeding the prefix-sum sampler). Recovery boost slots in cleanly as a
multiplier on priorities[idx] at insert time — no architectural
refactor, no new sampling-path machinery.
Recovery-oversample formula:
effective = base_priority × (1.0 + ISV[440] × ISV[439])
priorities[idx] = effective
priorities_pa[idx] = effective^alpha
When a transition is in a recovery (DD shrinking from non-trivial DD,
ISV[439]=1.0), it lands in the buffer at 1 + ω(2.0) × δ(1.0) = 3.0×
the baseline max_priority, then sampled 3.0^0.6 ≈ 1.93× more often
than baseline (alpha=0.6 PER convention) until per_update_pa
overwrites priority based on TD-error and normal PER takes over.
Recovery transitions get amplified gradient signal — the agent learns
recovery dynamics over typical "average-DD" Bellman noise.
Atomic landings (per feedback_no_partial_refactor):
* crates/ml-dqn/src/per_kernels.cu — per_insert_pa kernel signature
extended (+priorities, +isv pointers); kernel writes both columns
of the (priority, priority^alpha) row pair; nullable ISV ptr falls
back to boost_factor=1.0
* crates/ml-dqn/src/gpu_replay_buffer.rs — new isv_signals_dev_ptr
field + setter + accessor + priorities_pa_slice accessor; redundant
pre-3.5.5.b scatter_insert_f32 broadcast of max_priority REMOVED
(now subsumed by per_insert_pa's priorities[idx]=effective store)
* crates/ml/src/cuda_pipeline/gpu_experience_collector.rs — new
sp15_dd_trajectory_prev_dd_dev_ptr field + setter; per-step
launch_sp15_dd_trajectory_decreasing invocation gated on both
ISV ptr + prev_dd ptr being non-zero, placed immediately after
launch_sp15_dd_state in the env-step loop
* crates/ml/src/trainers/dqn/trainer/training_loop.rs — two new
wiring blocks for the collector's prev_dd ptr and the replay
buffer's ISV ptr, mirroring the existing
set_isv_signals_ptr / set_sp15_alpha_warm_count_ptr plumbing
* crates/ml/tests/sp15_phase1_oracle_tests.rs — new oracle test
per_sampler_weights_recovery_transitions_higher verifying both
columns of the priority-buffer write equal exactly the
boosted/baseline values (3.0 vs 1.0; 3.0^0.6 vs 1.0^0.6) and the
boosted/baseline ratio = 3.0 within 1e-5 — the recovery oversample
factor by construction
Phase 3.5.5 deferred consumer eliminated per
feedback_wire_everything_up. This closes the SP15 Phase 3.5
recovery-dynamics chain (3.5.2 + 3.5.3 + 3.5.4 + 3.5.4.b + 3.5.5 +
3.5.5.b all wired). DD_TRAJECTORY_FLOOR (slot 441) and ISV-driven
RECOVERY_OVERSAMPLE_WEIGHT (slot 440) producers remain documented
follow-ups per feedback_isv_for_adaptive_bounds — kernel reads from
slots rather than literals so consumer migration is a no-op when the
producers land.
Verified on RTX 3050 Ti (CUDA 12.9, sm_86):
* cargo check -p ml-dqn / -p ml --features cuda clean
* 3 of 3 dd_trajectory oracle tests still green
* 1 of 1 new per_sampler oracle test green (3.0 vs 1.0 priority
bias, ratio = 3.000... within 1e-5)
* 8 of 8 (1 ignored) gpu_residency replay-buffer + adamw smoke
tests green
* 4 of 4 PER smoke tests green
* Full ml lib suite IMPROVES Wave 4.2 baseline: 947 pass / 12 fail
(was 946 pass / 13 fail; the previously-failing
test_dqn_checkpoint_round_trip now passes — likely the redundant
pre-3.5.5.b scatter_insert_f32 was racing with the immediate
per_insert_pa overwrite under particular timing conditions; the
merged single-kernel write closes that race)
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Phase 3.5.4 (e0e0abfb2) landed the trigger + warm-up tracker but
deferred the actual weight reset (kernel accepted advantage_head_weights
+ n_weights but no-op'd via (void) cast). Wave 4.2 lands the real
reset.
When DD_PERSISTENCE exceeds threshold AND not yet fired this fold,
the kernel now resets the last 10% of advantage-head weights to
Kaiming-He init: Normal(0, sqrt(2/fan_in)) sampled via cuRAND
curand_normal() with per-thread state initialized from a host-passed
seed (Option A: per-call curand_init(seed, tid, 0, &state) for full
determinism — bit-identical samples for identical (seed, tid) pairs).
Single kernel, two phases:
1. Every thread independently re-evaluates fire_now from the same
ISV reads (DD_PERSISTENCE / threshold / fired_flag); the trigger
condition is a pure function of these reads so all threads
converge without cross-block synchronisation. Block 0 / thread 0
also runs the trigger + warm-bars decrement (single-thread ISV
write path).
2. If fire_now: each tid < reset_count writes Kaiming-He sample to
advantage_head_weights[reset_start + tid]. Per-thread independent
write, no atomic, no reduction (feedback_no_atomicadd clean).
New launcher params: fan_in (i32), seed (u64). Grid:
((n_weights/10 + 255) / 256).max(1) x [256, 1, 1]. The .max(1) floor
ensures block 0 always exists even when n_weights/10 == 0.
cuRAND device functions (curand_init, curand_normal) are inlined into
the cubin by nvcc from <curand_kernel.h> in the standard CUDA toolkit
include path — no host-side cuRAND linker dependency required, no
build.rs link change needed.
Oracle test plasticity_injection_kernel_resets_last_10pct_kaiming_he
verifies: (a) ISV[fired] flipped 0->1, (b) ISV[warm] = m_warm - 1,
(c) first 90% bit-identical to 1.0, (d) last 10% all moved off 1.0,
(e) sample mean |mean| < 0.05, (f) sample std within +-20% of
sqrt(2/fan_in), (g) determinism re-check produces bit-identical
samples for identical seed.
3 existing trigger tests migrated to the new launcher signature; the
debounced + no-fire variants additionally assert weight-stability
(early-out path skips the reset region).
Atomic per feedback_no_partial_refactor: kernel + launcher + 4 tests
+ audit doc + build.rs comment + 2 docstrings land together.
Eliminates Phase 3.5.4 deferred consumer per feedback_wire_everything_up
(the (void) casts are gone).
Action-selection consumer wiring (Phase 3.5.4.c) remains the separate
follow-up per the established Phase 3.5.X pattern.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Closes out Wave 4.1 (Phase 1.5.b consumer migration). Wave 4.1a (a8da1cb9c)
landed bn_tanh_concat_dd_kernel that fuses dd_pct into the trunk input;
Wave 4.1b (eb9515e41) wired s1_input_dim 102→103 through GRN reshape + 4
forward + 3 backward call sites. Wave 4.1c proves the wiring actually
changes the policy: a synthetic GPU forward composing launch_sp15_bn_concat_dd
with cublasSgemm_v2 against random Xavier-init weights W[proj_h=4, 103]
yields measurably different action distributions when ISV[DD_PCT]=0.0 vs
0.10 — observed mean KL=1.158e-4 (max=3.005e-4) vs threshold 1e-6
(~100× headroom).
Why a synthetic projection vs the real GRN trunk: the seeded forward_trunk_for_test
helper from Wave 4.1a noted (lines 2304-2319) that exposing the trainer's
trunk forward for tests would require either (a) a public surface change on
DQNTrainer exposing internal cuBLAS handles + GRN scratch + weights (the
trainer's fused_ctx is pub(crate) and only initialised inside the training
loop at training_loop.rs:547 — DQNTrainer::new returns with fused_ctx: None),
or (b) duplicating the trunk's cuBLAS setup in a test (≥200 lines of buffer
plumbing). Both options are architecturally heavier than the test's purpose
justifies. Per the spec dispatch ("the test's purpose is 'non-zero KL proves
the wire is connected' not 'verifies trained behavior'"), the synthetic
single-layer projection is the right scope: it exercises the new column-102
weights on the dd_pct value — exactly the path the real GRN's Linear_a first
GEMM takes for w_a_h_s1[:, 102] (the dd_pct column added by Wave 4.1b's
reshape).
Test contract:
- Two passes through launch_sp15_bn_concat_dd + cublasSgemm_v2 differ ONLY in
ISV[DD_PCT_INDEX=406] (0.0 at-ATH vs 0.10 in-DD).
- Inputs (bn_hidden, states) deterministic; weights deterministic via LCG
seed=42 with Xavier-uniform bound = sqrt(6 / (103+4)) ≈ 0.237.
- KL > 1e-6 (set 100× below the observed magnitude so a real wiring break
fires this test, not silently passing).
What this test does NOT verify: the full GRN composition (ELU/GLU/LN/residual)
propagating dd_pct through h_s2 + the branch advantage heads. That end-to-end
behavior is exercised by the L40S smoke + production training runs.
Phase 1.5.b orphan launcher chain fully eliminated per
feedback_wire_everything_up: kernel landed (4.1a) → consumer migration
(4.1b) → behavioral verification (4.1c) — three atomic commits, the
3a/3b/3c split-pattern matching Wave 3's a/b decomposition. The Wave 4.1a
transient orphan window opened in a8da1cb9c → closed in eb9515e41 →
behavioral coverage added here.
Wave 4.1a's seeded helpers consumed: kl_divergence (used) and
minimal_trainer_for_tests (retained but unused — the seeded comment
correctly identified that exposing the trunk forward via the trainer
surface is non-trivial, so the helper waits for a future cargo-cult test
that needs trainer construction without GPU forward, e.g. weight-shape
introspection).
Touched: crates/ml/tests/sp15_phase1_oracle_tests.rs (+1 module
sp15_wave_4_1c_behavioral with 1 ignored test, 2 helper fns, 1 assertion
block — purely additive, no kernel or production-code changes), docs/dqn-wire-up-audit.md
(Wave 4.1c entry at top of audit doc).
Verified:
- SQLX_OFFLINE=true cargo check -p ml --features cuda --tests clean (18
pre-existing unrelated warnings, no new warnings).
- CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests
--features cuda -- --ignored bn_concat dd_pct --nocapture: 2 of 2 oracle
tests green (Wave 4.1a bn_tanh_concat_dd_kernel_writes_dd_pct_column +
Wave 4.1c dd_pct_trunk_input_shifts_policy_distribution).
- cargo test -p ml --features cuda --lib: 947 passed / 12 failed —
exactly matches Wave 4.1b baseline (test addition is in the --test
integration target, not lib target).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The standalone dd_pct_concat_kernel from Phase 1.5 was bottleneck-
incompatible — it operated on raw [B, 128] state, but production trunk
consumes [B, s1_input_dim] = [B, 102] post-bottleneck. Wave 4.1a fixes
this at the kernel level; Wave 4.1b lands the consumer migration
(s1_input_dim 102→103, GRN w_s1 reshape, 3 forward + 3 backward sites).
Spec correction (per feedback_trust_code_not_docs): the spec's
'state_dim 48→49' is stale terminology pre-STATE_DIM 48→112→128
evolution. Production s1_input_dim is bottleneck_dim + (STATE_DIM −
market_dim) = 16 + (128 − 42) = 102. Wave 4.1b will bump this to 103.
NEW bn_tanh_concat_dd_kernel in dqn_utility_kernels.cu:
- Fuses dd_pct append into the same launch as bn_tanh + portfolio
concat (output shape [B, bn_dim + portfolio_dim + 1])
- Reads isv[DD_PCT_INDEX=406] (set by Wave 1.3.b dd_state_kernel
per-step), broadcasts the scalar across batch as the appended
last column
DELETED standalone dd_pct_concat_kernel.cu + launch_sp15_dd_pct_concat
+ cubin manifest entry per feedback_no_legacy_aliases (zero production
callers — only test consumer; bottleneck-on path is canonical).
Test helpers added (used by Wave 4.1c behavioral KL test).
Phase 1.5 oracle test migrated to bn_tanh_concat_dd_kernel contract:
test name bn_tanh_concat_dd_kernel_writes_dd_pct_column passes on
RTX 3050 Ti.
Layout fingerprint already covers Phase 1.5 via the existing
TRUNK_INPUT_DD_PCT=sp15_phase_1_5; marker — pre-SP15 checkpoints
already break.
fxcache schema_hash auto-bumps from file content hashes (per task
P5T5 Phase F mechanism); no manual schema bump needed.
Atomic per feedback_no_partial_refactor for the kernel-signature
contract change. Consumer wiring (s1_input_dim propagation, GRN
reshape, forward/backward call sites) deferred to Wave 4.1b's atomic
commit per the established 3a/3b split precedent — kernel + launcher
land first.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Wave 3a half of the val-cost-streams refactor (3b host-side wire-up
follows). Atomically migrates the kernel-side contracts; 5 launchers
remain orphan transiently awaiting 3b production callers.
Baseline kernels (1.4):
- 4 baseline_*_kernel signatures gain 'out: float*' parameter writing
per-window [mean, std, raw_sharpe] (matches 1.1.b sharpe_per_bar shape)
- ISV writes to slots 409, 410, 412, 416 removed entirely
(per-window output is correct for WindowMetrics consumption;
ISV-scalar writes were spec scaffolding for a single-fold-aggregate
version that 1.4.b's per-window contract supersedes)
- 4 ISV slot constants removed from sp15_isv_slots.rs
- state_reset_registry: NO entries to remove (verified via grep —
the 4 slots never had registry entries / dispatch arms in the first
place; they were single-fold-aggregate scalars defaulted at every
fold start by the constructor-write that initialises the ISV bus).
Task 4 from the dispatch is a no-op; the
every_fold_and_soft_reset_entry_has_dispatch_arm regression test
continues to pass unchanged.
- 4 oracle tests migrated to output-buffer assertion
- layout_fingerprint_seed string updated (4 retired entries removed,
4 trunk-shared entries retained; layout-break-class change)
New action_decoding_helpers.cuh:
- Extracts factored_action_to_dir_idx + factored_action_to_position
__device__ helpers (the latter is a higher-level position state-
machine helper not previously available)
- Mirrors trade_physics.cuh::decode_direction_4b semantics exactly so
on-policy and counterfactual paths agree on factored-action meaning
- Single source of truth for action→direction→position mapping;
consumers #include the header
New position_history_derivation_kernel.cu (post-loop derivation for
cost_net_sharpe consumer in Wave 3b):
- Reads actions_history_buf, reconstructs per-bar position_history
(-1/0/+1), side_ind (1.0 on position change), rt_ind (1.0 on
transition-to-flat from non-flat) via sequential walk (single
block per window, no atomicAdd per feedback_no_atomicadd)
- New launcher launch_sp15_position_history_derivation in
gpu_dqn_trainer.rs
- New cubin manifest entry in build.rs
- 1 oracle test covering 8-bar Short→Hold→Long→Hold→Flat→Long→Flat→
Short sequence; expected position/side_ind/rt_ind triples match
hand-computed values
Atomic per feedback_no_partial_refactor for the ISV-contract change
(every consumer of slots 409/410/412/416 migrated in this commit; their
consumers were the 4 oracle tests, all migrated). The orphan launcher
transient state for the 5 baselines + derivation kernel is explicitly
the 3a/3b split point — production callers land in 3b's
GpuBacktestEvaluator::new constructor signature change.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Closes the deferred-consumer gap left by Phase 3.1
(r_quality_discipline_split_kernel, commit 5d36f3238), Phase 3.3
(dd_penalty_kernel), and Phase 3.5.2 (dd_asymmetric_reward_kernel) —
three SP15 reward-axis kernels that landed only as standalone scalar
producers awaiting deferred consumer wiring per
feedback_no_partial_refactor.md.
Wave 2 chooses Option β (layered, composable post-modifier) over
Option α (replace SP11 entirely): SP11 B1b stays canonical
"trader-quality" composer that writes out_rewards; the SP15 reward-axis
composition becomes a fused PER-(i,t) post-modifier read-modify-writing
the same buffer in place. This preserves SP11's z-score mag-ratio
contract (canary tests untouched) while landing all three deferred SP15
consumers atomically with zero parallel paths.
Architecture (Q1-Q4 user resolutions):
Q1: SP12 caps stay as state_layout.cuh macros (REWARD_NEG_CAP=-10,
REWARD_POS_CAP=+5), NOT lifted to ISV — spec'd constants per the
SP12 v3 design, not adaptive bounds.
Q2: Both on-policy + CF slots get the same DD-aware shaping. CF reward
= w_cf × r_cf (SP11 controller weight already applied) is composed
via the same helpers as on-policy. DD context is per-step, not
per-action.
Q3: New slot_completed_normally[N*L] flag preserves SP11's early-return
semantics (data-end at experience_kernels.cu:2142 → reward=0.0;
blown-account at :2203 → reward=-10.0). Fused kernel skips slots
where flag==0.
Q4: alpha_split_producer_kernel OWNS the warm-count increment (per-step
scalar producer; the fused kernel has 2*N*L threads and would
over-tick by that factor if it owned the increment).
Per-step launch sequence in gpu_experience_collector.rs (after Phase
1.3.b's dd_state launch): experience_env_step → alpha_split_producer
(reads grad-norm slots 418/419, writes ALPHA_SPLIT slot 417, increments
warm-count) → compute_sp15_final_reward (fused per-(i,t) over [N*2*L]
slots, applies α-blend → DD asymmetric → DD penalty → SP12 cap, writes
back to out_rewards in place).
experience_env_step signature change — 2 new output params:
r_discipline_out: float* [N*L] — per-step REGRET_EMA mirror, written
at end of normal reward composition.
slot_completed_normally_out: int* [N*L] — 0 default at entry, set to
1 only on the path that reaches out_rewards[out_off] = reward.
3 deleted kernel files:
- r_quality_discipline_split_kernel.cu (composer + producer; replaced
by renamed alpha_split_producer_kernel.cu keeping ONLY the producer
with the moved warm-count increment).
- dd_penalty_kernel.cu (replaced by sp15_dd_penalty __device__ helper).
- dd_asymmetric_reward_kernel.cu (replaced by sp15_dd_asymmetric_reward
helper).
3 new files:
- alpha_split_producer_kernel.cu (per-step scalar producer of α from
grad-norm ratio, with warm-count increment moved here per Q4).
- sp15_reward_axis_helpers.cuh (4 __device__ inline helpers:
sp15_alpha_blend, sp15_dd_asymmetric_reward, sp15_dd_penalty,
sp15_apply_sp12_cap).
- compute_sp15_final_reward_kernel.cu (fused per-(i,t) parallel over
[N*2*L] slots — α-blend + DD-asymmetric + DD-penalty + SP12 cap,
skips early-return sentinel slots).
3 deleted launchers + 3 deleted CUBIN statics in gpu_dqn_trainer.rs:
- launch_sp15_r_quality_discipline_split (composer scalar variant).
- launch_sp15_dd_penalty.
- launch_sp15_dd_asymmetric_reward.
- SP15_R_QUALITY_DISCIPLINE_SPLIT_CUBIN.
- SP15_DD_PENALTY_CUBIN.
- SP15_DD_ASYMMETRIC_REWARD_CUBIN.
2 new launchers + 2 new CUBIN statics:
- launch_sp15_final_reward + SP15_FINAL_REWARD_CUBIN.
- SP15_ALPHA_SPLIT_PRODUCER_CUBIN (the retained launch_sp15_alpha_split_producer
now loads this).
GpuExperienceCollector: 2 new CudaSlice fields
(r_discipline_per_sample, slot_completed_normally_per_sample) +
sp15_alpha_warm_count_dev_ptr field + set_sp15_alpha_warm_count_ptr
setter + Step 5b launch block.
State reset registry — no new entries: r_discipline +
slot_completed_normally are per-step ephemeral (defaulted at every
kernel entry); cross-fold leakage impossible. Existing Phase 3.1/3.3/
3.5.2 ISV slot entries cover the rest.
6 new oracle tests in sp15_phase1_oracle_tests.rs::mod gpu drive
compute_sp15_final_reward_kernel directly with hand-crafted buffers:
- final_reward_alpha_blend_at_cold_start (Stage 1)
- final_reward_dd_penalty_above_threshold (Stage 3)
- final_reward_dd_asymmetric_gain (Stage 2 + R_GAIN_DD_BOOST diag)
- final_reward_dd_asymmetric_loss (Stage 2 asymmetric guard)
- final_reward_sp12_cap_clamps_both_directions (Stage 4)
- final_reward_skips_early_return_slots (Q3 sentinel preservation)
9 deleted scalar-kernel oracle tests:
- r_split_uses_sentinel_alpha_at_cold_start
- r_quality_subtracts_explicit_cost
- dd_penalty_quadratic_above_threshold + dd_penalty_zero_below_threshold
- dd_asymmetric_reward_gain_amplified_by_dd_pct +
dd_asymmetric_reward_loss_unchanged + dd_asymmetric_reward_no_op_at_ath
(the new fused-kernel tests cover the same behavioral surface
end-to-end through the in-place RMW path).
Verified: SQLX_OFFLINE=true cargo check -p ml --features cuda clean (18
unrelated warnings); all 29 SP15 phase1 oracle tests pass on RTX 3050 Ti
(includes the 6 new fused-kernel tests + 17 retained tests + 6 old);
SP11 mag-ratio canary tests still untouched (no canary-test renames or
deletions); ml lib suite holds 946 pass / 13 fail = baseline.
Hard rules: feedback_no_partial_refactor (3 phases' deferred consumers +
3 deletions + 3 new files + 6 new tests + audit doc all in this commit;
no parallel paths, no feature flags), feedback_wire_everything_up
(closes 3 SP15 phase orphan launchers atomically), feedback_no_legacy_aliases
(deletions land in same commit as replacement; no compatibility shim),
feedback_no_atomicadd (fused kernel is per-(i,t) parallel, pure scalar
arithmetic), pearl_audit_unboundedness_for_implicit_asymmetry (gain-only
DD multiplier + asymmetric NEG/POS caps preserve loss aversion),
pearl_symmetric_clamp_audit (SP12 cap is bilateral via fmaxf/fminf even
though bounds are intentionally asymmetric per spec),
pearl_no_host_branches_in_captured_graph (new launches happen inside
collect_experiences_gpu::launch_timestep_loop per-step, OUTSIDE the
experience-fwd CUDA Graph capture region — same precedent as Phase
1.3.b's dd_state launch).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Phase 3.5 (hold_floor_kernel) + Phase 3.5.3 (cooldown_kernel) landed
the producer + state machinery; both deferred the action-selection
consumer wiring. This task wires both atomically.
Architectural decision: hold_floor is now an INLINE __device__
computation inside experience_action_select reading ISV slots
426/427/428/429 directly. The standalone hold_floor_kernel.cu +
launch_sp15_hold_floor + HOLD_FLOOR_CUBIN are deleted — launching a
kernel to write one f32 just to read it back was unnecessary. ISV
slots + state_reset_registry entries remain; only the launch path is
removed per feedback_wire_everything_up + feedback_no_legacy_aliases.
Entropy source: per-step Shannon entropy of softmax(e_dir) computed
inline from the 4 e_dir floats already in registers (Pass 1 of the
Thompson direction selector). High entropy = uncertain policy → Hold
gets the floor lift; low entropy = confident policy → floor ≈ 0.
q_eff_dir scratch preserves e_dir for downstream consumers
(out_conviction, out_q_gaps, out_magnitude_conviction) — adding
hold_floor there would corrupt the Kelly-cap warmup floor with a
meta-confidence mask.
cooldown mask: when ISV[COOLDOWN_BARS_REMAINING=435] > 0,
action_select hard short-circuits to dir_idx = DIR_HOLD before
Pass 2 — sidesteps the temperature-blend numerics where a
finite-sentinel-on-non-Hold approach would let pure-Thompson (τ=1)
samples dominate the masked direction. Cooldown supersedes
hold_floor — when forcing Hold the floor is moot.
3 new oracle tests:
- action_select_applies_hold_floor_inline (no cooldown)
- action_select_forces_hold_during_cooldown
- action_select_no_force_hold_when_cooldown_zero
Atomic per feedback_no_partial_refactor: action_select changes +
hold_floor_kernel deletion + cubin manifest update + 3 oracle tests +
audit doc all in this commit. No parallel paths, no feature flags.
Eliminates Phase 3.5 + Phase 3.5.3 deferred consumers. The
cooldown_kernel itself remains (it maintains the consecutive_losses
streak + decrements COOLDOWN_BARS_REMAINING per bar); only its
consumer is now wired.
Verified: cargo check -p ml --features cuda clean; ml lib suite
holds 946 pass / 13 fail = baseline; all 6 oracle tests pass on
RTX 3050 Ti (3 pre-existing cooldown + 3 new action_select).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Phase 1.6 (CLI flags + dev_features/holdout_features stash) and Phase
1.7 (set_test_data_from_slices observer + test_features stash) landed
the data-flow scaffolding; both deferred the actual eval consumer.
This task wires both atomically as parallel evaluator instances:
- dev_evaluator: Option<GpuBacktestEvaluator> -- lazy-init after final
fold, fires once against Q8 dev_features (when dev_quarters > 0)
- test_evaluator: Option<GpuBacktestEvaluator> -- lazy-init per fold,
fires inside the fold loop against the WF test slice (when
fold.test_end > fold.test_start)
Architectural choice: parallel evaluator instances (NOT window-swap on
val_evaluator). Window-swap would require invalidating the CUDA graph
between val and dev/test runs -- fragile, and a direct violation of
pearl_no_host_branches_in_captured_graph. Parallel instances mirror
val_evaluator's lazy-init pattern (TLOB sync, ISV signal pointer,
training_mode = false toggle). Implementation lives behind a single
shared helper `launch_extra_eval` keyed on an `ExtraEvalKind` enum so
Dev / Test share TLOB / ISV / config setup verbatim.
HEALTH_DIAG additions:
HEALTH_DIAG[N]: dev_eval dev_sharpe_net=... dev_calmar=... dev_max_dd=... dev_trades=...
HEALTH_DIAG[N]: test_slice fold=K test_sharpe_net=... test_calmar=... test_max_dd=... test_trades=...
The *_sharpe_net key uses the fused-metrics-kernel cost-aware Sharpe
(post-Phase-1.1.b split -- already includes tx_cost_bps + spread_cost
via the env-step PnL feed); when Phase 1.2.b cost-net sharpe lands, the
key name is preserved so the aggregator-script contract holds.
Atomic per feedback_no_partial_refactor: both eval calls + both
evaluator fields + both HEALTH_DIAG lines + audit doc all in this
commit. No parallel paths, no feature flags. Dev_eval runs synchronously
via evaluate_dqn_graphed (one-shot, no async pipelining benefit since
it doesn't fire per-epoch); val path stays async.
Sealed Q9 holdout remains untouched -- Phase 4.3 will load Q9 via a
separate eval-only entry point (NOT train_walk_forward). The Phase 1.6
debug_assert sealed-slice guard catches accidental future refactors.
Verified: cargo check -p ml --features cuda clean; ml lib suite holds
946 pass / 13 fail baseline; the existing Phase 1.7 oracle test
set_test_data_from_slices_fires_observer_and_stashes still passes.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Path A of the blocked 1.3.b investigation: fixes two architectural
issues atomically and wires the launcher.
(1) Bug fix: dd_state_kernel.cu was recomputing new_equity =
PS_PREV_EQUITY + pnl_step and writing it back, but experience_env_step
already maintains PS_PREV_EQUITY (experience_kernels.cu:3473-3475) —
wiring as-is would silently double-accumulate equity every step.
Kernel now READS PS_PREV_EQUITY / PS_PEAK_EQUITY only; does not
modify them. pnl_step parameter dropped from both kernel and
launcher signatures.
(2) Per-env shape decision: kernel is single-thread/single-block;
production has N envs but DD ISV slots [401..407) are scalars.
Picks 'env 0 as canonical observable' — kernel reads
pos_state[0 * PS_STRIDE + ...]. Per-env redesign (per-env tiles +
reduction kernel) deferred to Phase 1.3.b-followup if L40S smoke
shows single-env DD aggregation is insufficient.
(3) Wire-up: launch added at gpu_experience_collector.rs step 5b in
launch_timestep_loop, immediately after env_step writes PS_PREV_EQUITY,
outside the exp-fwd graph capture region (which ends at line ~3829,
well before env_step). Atomic per feedback_no_partial_refactor:
kernel signature change + oracle test update + launcher call site
update all in this commit.
Eliminates the Phase 1.3 orphan launcher per feedback_wire_everything_up.
Downstream Phase 3.3 / 3.5.2 / 3.5.4 / 3.5.5 readers will receive live
DD values when their consumer wiring lands.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
jgrusewski