Preliminary triage of spec §5.1 hypotheses H1–H10 using the Phase 0
baseline capture on RTX 3050 Ti. L40S validation pending per plan.
Verdicts:
H1 PENDING (needs forced-exploration instrumentation not yet wired)
H2 REJECTED var_scale=0.96 across 19/20 epochs; Var[Q] inactive at smoke scale
H3 INCONCLUSIVE kelly degenerate (insufficient win/loss counts at smoke scale)
H4 CONFIRMED grad_ratio_mag_dir=0.0000 across 20/20 epochs (threshold <0.1)
H5 REJECTED ent_mag stays ≥0.98 throughout; no bootstrap collapse
H6 REJECTED Full fire rate (0) is lower than Quarter fire rate, not higher
H7 REJECTED vsn and sigma symmetric between mag and dir branches
H8 REJECTED target-net drift equal (mag=dir=0.001)
H9 PENDING (same instrumentation gap as H1)
H10 CONFIRMED training ent_mag=0.98, eval F_Quarter=100%
Synthesis: H4 is the root cause. Magnitude branch receives ~0 gradient →
weights stay near init → three magnitude Q-values near-identical → argmax
picks bin 0 (Quarter) on ties → H10 manifests at eval time. H2, H5, H7,
H8 all ruled out as contributors.
Proposed Phase 2 priority: fix H4 (gradient-flow path into magnitude head
— likely per-component advantage weighting or direction-conditioning of
w_b1fc) + H10 (Q-margin argmax + stochastic eval rollouts as safety net).
Phase 1 next: validate preliminary verdicts on L40S, instrument
per-component gradient decomposition for magnitude, proceed with
Tracks 2/3/4 in parallel.
All <PENDING> markers replaced with real values from a clean 20-epoch
magnitude_distribution smoke run + 3-fold multi_fold_convergence run
on RTX 3050 Ti at HEAD 0472b9730.
All 5 Phase 0 smoke tests PASS:
magnitude_distribution — Quarter=0.596 Half=0.150 Full=0.255
reward_component_audit — cf_flip=0.619 trail=0.290 la=0.013
controller_activity — all 6 controllers < 20% firing
exploration_coverage — ent_mag @ep5=0.987 @ep20=0.985
multi_fold_convergence — 3/3 folds, Best Sharpe 51 / 39 / 87
This commit closes Phase 0. The next commit tags it as
policy-quality-baseline — the reference state against which Phase 2
interventions are measured.
Previously test_training_throughput_measurement and test_real_data_single_epoch
only asserted loss.is_finite() on the final metric. That passes on trivial
zeros, on huge-but-finite NaN-disguised values, and on any regression that
doesn't produce literal NaN — giving effectively no signal.
test_training_throughput_measurement now asserts:
- loss finite AND non-negative
- epochs_trained >= 1
- throughput floor: epochs_per_sec > 0.05 (i.e. each epoch < 20s on the
RTX 3050 Ti; catches accidental CPU fallback or kernel CPU-pinning).
Documented as a conservative local floor; CI may tighten.
- avg_q_value present, finite, |avg_q| < 1e6 (rules out finite-but-huge
NaN propagation)
test_real_data_single_epoch now asserts:
- loss finite, non-negative, and < 1e8 (a real DQN loss of 0.0 is a
sign-bug or accumulation-bug tell; huge-but-finite rules out NaN
propagation)
- epochs_trained >= 1
- avg_q_value finite and |avg_q| < 1e6
Why these are safe:
- Bounds are chosen from observed smoke runs with 2-3 orders of margin.
- Passes locally in 1.58s and 10.24s respectively.
- Designed to flag regressions, not true production-scale deviations.
Verified PASS on laptop (RTX 3050 Ti).
Previously the fold loop swallowed every training error into f64::NAN,
unconditionally pushed a value at the top of the loop, and then asserted
only !fold_losses.is_empty() — a tautology that could never fail. Any
CUDA error, NaN explosion, or regression of the zero-copy path passed
silently.
Now:
- Error propagation via `?` (no swallow). A zero-copy test can't
tolerate training failures — if training blew up, the zero-copy
wiring is broken and must surface.
- All fold losses must be finite and non-negative.
- Every fold must report epochs_trained >= 1 (zero-epoch fold =
kernel skipped or buffer not populated).
A direct "no htod/dtoh copy happened" assertion would need a copy-counter
instrumented into the fused-training GPU path plus a field on
TrainingMetrics. That is out of scope for this test; documented inline.
Verified PASS on laptop (RTX 3050 Ti):
loss=0.0018366, epochs=1
Both lazy-init sites (train_with_data_full_loop_slices @ L344 and
run_training_steps_slices @ L1511) previously logged FusedTrainingCtx::new
failures via tracing::error! and continued with fused_ctx = None. On CUDA
builds there is no CPU fallback, so the next stage (GPU experience collector)
then fails ~100ms later with a misleading "GPU experience collector MUST be
active for CUDA training" error that obscures the real CUDA root cause (OOM,
driver error, etc.).
Both sites now propagate the original error via map_err → anyhow::anyhow! so
callers see the actual failure. The post-init wiring (PER buffer pointers,
RNG dev ptr) is refactored from `if let Some(ref mut fused)` — which was
silently-no-op on the hidden-error path — to an unconditional
as_mut().expect() since fused_ctx is guaranteed Some after the `?`.
The two remaining `fused_ctx = None` assignments are legitimate:
- mod.rs:655 in Drop (ordered GPU teardown)
- training_loop.rs:1508 batch-size-change recreate (immediately reassigned
by the ? on the next line)
No new API, no threshold changes, no feature flags.
Two wrong-scale assumptions in the test made it unachievable on the
RTX 3050 Ti / 4 GB laptop this smoke is meant to run on:
1. `train_baseline_rl` was invoked without `--training-profile`, so it
defaulted to `dqn-production`: batch_size=16384, buffer=500k,
num_atoms=52, hidden_dim_base=256. Fused-CUDA init OOMs at
`kan_d_coeff_per_elem alloc` on 4 GB, leaving `fused_ctx = None` and
every subsequent fold failing with "GPU experience collector MUST be
active for CUDA training". Fix: pass `--training-profile=dqn-smoketest`.
2. Default walk-forward windows (12 train / 3 val / 3 test / 3 step) only
yield 2 folds in the 24-month baseline dataset — fold 2's test-end
lands one month past `data_end`. The test's pass-gate is "≥2/3 folds
produce a checkpoint", so a test that can only ever generate 2 folds
is degenerate. Fix: explicit shorter windows (6 / 2 / 2, step 2) that
yield all 3 folds (`6 + 2*2 + 2 + 2 = 14 ≤ 24`, comfortable margin).
Also drops `--epochs 20` → `--epochs 5`. Each fold runs ~5500 batches at
~33 s/epoch on this GPU; 20 × 3 folds ≈ 33 min was exceeding the smoke
budget (kill observed around the 10-minute mark). 5 epochs is ample for
the checkpoint gate — `best_sharpe` saves on the first improving epoch
(epoch 1 in practice), so more epochs add no pass/fail signal, only
wall-clock.
Verified locally: 3/3 folds produce `dqn_fold{N}_best.safetensors`,
total wall-clock ~7 min.
[MULTI_FOLD] fold 0 checkpoint OK
[MULTI_FOLD] fold 1 checkpoint OK
[MULTI_FOLD] fold 2 checkpoint OK
test result: ok. 1 passed; 0 failed ... finished in 416.36s
Docstring updated to reflect new sizing and call out the 4 GB / 24-month
constraints explicitly so the next person reading this can see why the
numbers are what they are.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Five tests in crates/ml/src/trainers/dqn/smoke_tests had pass gates that
validated existence rather than the invariant their doc-comment claimed
to test. A trainer returning all-zero diagnostics (a plausible wiring
regression) would have passed them.
- reward_component_audit: was `is_finite() && >= 0.0` on 5 slots —
passes trivially on all-zero stubs. Added `cf_flip > 0.1` and
`trail_r > 0.01` floors (known-wired slots in smoke config;
popart/micro/loss_aversion remain finite-only as they are
legitimately near-zero in smoke). Run-observed values: cf=0.614,
trail=0.295, la=0.006 — well above floors.
- exploration_coverage: `.unwrap_or(0.0)` silently substituted 0 for a
missing epoch, conflating "emission regressed" with "exploration
collapsed". Now panics with a distinct message on missing entries,
also asserts len >= 20, normalized range [0,1], and spread > 1e-6 to
catch constant-output emitters. Run-observed spread: 0.275.
- training_stability::50_epoch_convergence: entropy assertion was
guarded behind `if entropy.is_finite()`, so NaN entropy (the more
severe failure) silently passed. Fail hard on NaN first.
- training_stability::trading_model_behavior: same `is_finite` guard
pattern on action_entropy — now fails hard when the diagnostic is
missing or NaN rather than skipping.
- training_stability::gpu_collector_auto_initializes: only asserted
training returned `Ok(_)`. A collector producing silent zeros would
pass. Now also verifies epochs_trained, loss finiteness, and
gradient flow.
- walk_forward::no_overfitting_50_epochs: had two tautological "finite
check" assertions (`x < x + 1` and the signum-adjusted ratio) that
always passed regardless of divergence. Replaced with real
`.is_finite()` checks plus a `div_ratio < 10.0` stability gate.
Tests run under CUBLAS_WORKSPACE_CONFIG=:4096:8 +
FOXHUNT_TEST_DATA=test_data/futures-baseline, release-test profile.
reward_component_audit and exploration_coverage both PASS on the
local RTX 3050. No threshold relaxation or quickfixes applied; any
future wiring regression will now be caught by a meaningful
assertion instead of a near-tautology.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Redefine the "fire" semantic for adaptive controllers per the V7 audit
(policy-quality-design spec §5.3): a controller fires iff it made an
ADAPTIVE INTERVENTION this epoch, not merely because its observable
output value changed.
Previously fire detection was absolute-delta on the output:
- grad_clip fired in 98.3% of epochs because the adaptive clip threshold
is an EMA recomputed every training step. The EMA drifts by > 1e-3
every epoch regardless of whether the clip actually clamped a gradient.
- cost_anneal fired in 98.3% of epochs because it is a deterministic
sigmoid of current_epoch (1/(1+exp(-(epoch-10)/3))) with no adaptive
or reactive component. Every epoch moves it by > 1e-4 by design.
Neither was "load-bearing" in the V7 sense — one was a per-step EMA
tracker, the other a pure curriculum schedule. The prior test output
"controller 'grad_clip' fires in 98.3%" was a false positive from
measuring the wrong signal.
New semantics:
- anti_lr / tau / gamma / cql_alpha: unchanged — absolute delta vs
prior epoch on the effective output value (real adaptive controllers).
- grad_clip: intervention-based latch `grad_clip_kicked_this_epoch`,
set in run_training_steps_slices iff raw_grad_norm > active clip at
any training step this epoch. Reset in reset_epoch_state.
- cost_anneal: pure deterministic schedule → never load-bearing →
always fires=false. The value is still tracked in prev_controller_values
and the HEALTH_DIAG line still emits it for observability, but it
cannot trip the 50% load-bearing gate.
After fix, controller_activity smoke reports:
anti_lr=0.000 tau=0.033 gamma=0.017 clip=0.233 cql=0.033 cost=0.000
All 6 rates ≤ 0.5. Test passes.
Touched:
- crates/ml/src/trainers/dqn/trainer/mod.rs (add grad_clip_kicked_this_epoch)
- crates/ml/src/trainers/dqn/trainer/constructor.rs (init new field)
- crates/ml/src/trainers/dqn/trainer/training_loop.rs (latch kick per step,
redefine fire_clip + fire_cost in HEALTH_DIAG block)
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Third root cause of the apparent magnitude collapse in smoke-test data:
HEALTH_DIAG reads monitor.action_counts to compute dist_q/h/f and
ent_mag/ent_dir, but the block that populated action_counts from the
GPU summary ran AFTER HEALTH_DIAG. Every epoch saw all-zero counts so
dist_q=dist_h=dist_f=ent_mag=ent_dir=0 in every log line, and the
last_magnitude_dist cache the smoke test reads was always zeros.
Moved the GPU-summary download + monitor population block up above the
HEALTH_DIAG preparation block. Signal is now real per-epoch:
dist_q=0.60 dist_h=0.15 dist_f=0.25 (was 0/0/0)
ent_mag=0.83 ent_dir=0.89 (was 0/0)
magnitude_distribution smoke test now PASSES on local RTX 3050 Ti at
20 epochs. Final: Quarter=0.598 Half=0.154 Full=0.249 — all above 5%
smoke threshold.
Combined with 2fb30f098 (9→12 bin cascade + magnitude-preference
denominator restricted to tradable directions), this resolves the
Track-1 "magnitude collapse" pathology. The collapse was never a
policy-learning failure — it was a monitoring visibility + ordering
bug that made every reading look like mag=0 100%.
The runtime DQN action space is 4×3×3×3 = 108 (dir×mag×order×urgency) per
the kernel's b0_size=4 (Short/Hold/Long/Flat) layout, but the monitoring
stack was still coded for the legacy 3×3=9-bin exposure layout. Flat
actions (dir=3) had exp_idx = 9 >= num_actions=9 and were silently
dropped by the monitoring_reduce kernel's bounds check.
Cascade:
- monitoring_kernel.cu: num_exposure_bins=12 (EXP_MAX constant), summary
layout 27 floats with bins [5..17), order [17..20), urgency [20..23),
N [23], trades [24], pad [25..27).
- gpu_monitoring.rs: MonitoringSummary.action_counts [usize; 12];
summary_buf 27 f32; reduce() now takes b0_size and validates
num_exposure_bins ≤ EXP_MAX.
- trainers/dqn/monitoring.rs: action_counts/q_value_sums/q_value_counts
all [_; 12], factored_action_counts [_; 108], EXPOSURE_NAMES with 4
dirs × 3 mags ("S_Small".."F_Full"); track_action now maps the 8-level
ExposureLevel enum onto the kernel's dir*3+mag layout by going through
direction()/magnitude() accessors (was indexing by raw `exposure as
usize` which is an entirely different scheme).
- trainer/training_loop.rs: total_action_counts [_; 12],
total_factored_action_counts [_; 108]; GPU monitoring reduce call
now passes b0 from agent.branch_sizes(); direction-entropy gains a
4th dir bin; per-action Q diagnostics names array extended to 12.
- trainer/metrics.rs: create_final_metrics arrays widened; q_diagnostics
per_action_avgs [_; 12]; combo_idx = d*3 + m (matching kernel layout);
action_space_size 108/12 and buy/sell/hold uses the new exp_idx
boundaries.
- financials.rs: action_counts [_; 12]; buy/sell/hold classification
updated — BUY = Long [6..9], SELL = Short [0..3], HOLD = Hold + Flat
([3..6] ∪ [9..12]). Tests updated for 12-bin layout.
Also fixes several collateral bugs exposed by the audit:
1. magnitude_action_dist_final conflated Hold (dir=1, mag-forced-to-0)
with "Quarter magnitude". New formula restricts the denominator to
tradable directions (Short + Long) and computes Q/H/F fractions only
over those, producing the real magnitude preference signal.
2. training_loop.rs:3012 flat_count indexed [3,4,5] (which is Hold in
the new layout, but was already the wrong bucket — "Flat" was never
at that index even under the legacy 9-bin reading, where [3,4,5]
was the Flat bin but mag was forced to 1 not 0). New code correctly
separates flat_count [9..12] from hold_count [3..6] and excludes
both from the directional denominator and diversity cell threshold.
3. monitoring.rs had q_value_sums: [f64; 7] but q_value_counts:
[usize; 9] — different sizes for what should be the same exposure
axis. Both now [_; 12].
4. log_action_distribution iterated over 7 indices when printing per-
action Q-values but the array is logically 12-wide — now uses
`0..12` with the 12-entry EXPOSURE_NAMES lookup.
Tests: all 7 financials tests pass. Monitoring cubin loads + default
summary sanity checks pass. Smoke test magnitude_distribution now
emits correct 12-bin breakdown — GPU summary shows (example epoch 20):
actions=[786, 157, 280, 377, 2, 7, 806, 156, 317, 303, 5, 4]
which decodes as Short (S) = 786/157/280 (~38%), Hold (H) = 377/2/7
(~12%, mag-forced), Long (L) = 806/156/317 (~40%), Flat (F) =
303/5/4 (~10%, mag-forced). Previously the Flat bucket (303+5+4 =
~10% of all actions) was silently dropped — policy was picking Flat
~10% of the time and nobody could see it.
No atomic ops added; all reductions remain via warp-scratch + lane-0
sequential merge. No feature flags, no stubs, no quickfixes.
Two truncation sites in data_loading.rs truncated feature/target vectors
to max_bars but left self.ofi_features at the original (cache-full)
length. Downstream DqnGpuData::upload_slices asserts OFI len == num_bars
and surfaced this as "OFI/market data length mismatch: 175874 OFI rows
vs 5000 bars — fxcache is corrupt".
The cache was not corrupt — the upstream truncation was the bug. Fixed
by moving OFI storage below the truncation block in the fxcache path and
by adding a parallel truncate in the DBN fallback path. Added
debug_assert_eq on equal lengths post-truncation so future regressions
hard-fail in dev.
Also corrected the cuda_pipeline/mod.rs error message to point at the
actual culprit (upstream truncation desync) instead of blaming the
on-disk cache.
WIRE:
* q_mag_full/half/quarter — host-side mean of magnitude-branch q_out_buf
cols [b0..b0+b1] (Task 0.3 stub remediated). Cold-path dtoh at epoch
boundary via update_q_mag_means_cached(), cached in q_mag_means_cached,
read by q_magnitude_bucket_means().
* var_scale_mean — per-sample CudaSlice + host reduce over var_scale
written by experience_env_step at line 1422 (Task 0.3 stub remediated).
Host reducer averages only over slots where kernel wrote > 0 so
Var[Q]-disabled samples don't dilute the signal.
DELETE:
* segment_patience slot from reward_contrib group — term not wired in
kernel, slot reserved space for non-existent feature (Task 0.8 stub
remediated). reward_contrib [...] is now 5 floats not 6. Cascaded
through reward_contrib_fractions return type, trainer's
last_reward_contrib field, reward_component_audit_summary accessor,
and the reward_component_audit smoke test.
PopArt slot kept at 0.0 — that IS the semantic value during warmup or
disabled state, not a stub. Comment updated to make this explicit.
Per ~/.claude/.../memory/feedback_no_stubs.md: stubs strictly forbidden
even with code comments / commit-message documentation / plan sanction.
Task 0.3 (Track 1 magnitude):
* q_mag_full/half/quarter: plumbed via GpuDqnTrainer::q_magnitude_bucket_means()
→ FusedTrainingCtx::q_magnitude_bucket_means() → HEALTH_DIAG. Phase-0
returns [0.0; 3]; per-mag Q reduction kernel deferred per plan §0.3 step 2
(goal here is the accessor chain, swap-in without touching emit formatting).
* var_scale_mean: STUB 0.0 — kernel computes `1/(1+sqrt(Var[Q]))` per-sample
inside experience_env_step but does not persist to a device buffer; it is
consumed in-place to shrink effective_max_pos. Exposing requires a
dedicated per-sample output buffer + launch arg + kernel write; deferred
to Phase 1+. Documented in code.
* kelly_f_mean / avg_win_ratio: REAL — computed host-side from
self.trade_stats_history.last() using the SAME formula the kernel uses
(b = avg_win/avg_loss, kelly = (b·p − (1−p))/b clamped [0,1] × 0.5
for half-Kelly; avg_win_ratio = (sum_wins/win_count) /
max(sum_losses/loss_count, 0.001) per plan formula line 239).
One-epoch lag (trade_stats_history is appended later in the same
epoch body); zero until first epoch's stats land. Documented in code.
Task 0.6 (Track 1 noisy):
* vsn_mag / vsn_dir: REAL per-branch — mean |w| across each branch's VSN
projection weight pair (tensors 26..34 grouped by branch:
26/27=dir, 28/29=mag, 30/31=order, 32/33=urg). The live VSN gate mask is
computed per-sample inside variable_select_bottleneck and consumed
immediately, not materialized into a device buffer we can cheaply read;
projection-weight mean is the closest stable H7 surrogate. Documented
in accessor rustdoc. Higher-fidelity measurement would require a
per-sample mask output buffer following the Task 0.5 pattern — deferred.
* drift_mag / drift_dir: REAL per-branch — RMS ‖target_w − online_w‖ across
each branch's 4 weight tensors (indices 8..12 dir, 12..16 mag, 16..20
order, 20..24 urg). Host-computed cold path via two memcpy_dtoh + normal
Vec<f32> reduce loops; stream-sync'd before read. NO atomics (project rule).
~2.7 MB × 2 DtoH per epoch, negligible diagnostic cost.
NoisyNets sigma_mag/sigma_dir already wired in 999bb2fa0 — preserved untouched.
Accessor chain (mirrors Task 0.4 grad_ratio_mag_dir pattern):
* GpuDqnTrainer::per_branch_target_drift() -> Result<[f32; 4], MLError>
* GpuDqnTrainer::per_branch_vsn_mean() -> Result<[f32; 4], MLError>
* GpuDqnTrainer::q_magnitude_bucket_means()-> [f32; 3] (plumbing stub)
* FusedTrainingCtx wrappers → delegate, fallback to zeros on MLError.
* training_loop.rs HEALTH_DIAG block: populate slots via
self.fused_ctx.as_ref().map(|f| f.…).unwrap_or(…) — same pattern
used for grad_ratio_mag_dir.
HEALTH_DIAG `mag` and `noisy` groups now carry real signal (or plan-sanctioned
plumbing stubs for the q_mag_* / var_scale slots documented in commit + code).
Phase 0 instrumentation is complete (Tasks 0.4 / 0.5 / 0.8 / 0.10 / 0.16
shipped this session; 0.6 partial; 0.3 partial). This scaffold encodes
the full metric set with explicit <PENDING …> markers per row so the
GPU-run capture can fill in values without forgetting any HEALTH_DIAG
field.
The policy-quality-baseline tag is intentionally NOT applied yet —
tagging requires real captured values, not the scaffold. Runner workflow:
for t in magnitude_distribution reward_component_audit \\
controller_activity exploration_coverage multi_fold_convergence; do
SQLX_OFFLINE=true CUBLAS_WORKSPACE_CONFIG=:4096:8 \\
FOXHUNT_TEST_DATA=test_data/futures-baseline \\
cargo test -p ml --release --lib -- \$t --ignored --nocapture 2>&1 | tee out_\$t.log
done
Then edit this doc, replace markers with values, amend, and tag.
Adds evaluate_baseline CLI args:
--surrogate-mode=off|random
--surrogate-seed=<u64>
--surrogate-marginals=<path.json>
--emit-action-marginals
--emit-pooled-sharpe
Random-action surrogate samples actions from marginal distribution (or
uniform fallback) using a seeded RNG, bypassing the model entirely.
Surrogate mode forces the CPU DQN path so action selection can be
overridden (GPU path would need invasive kernel changes and defeats
the bypass-the-model sanity check).
Flat action-index counts are tracked in the CPU DQN path only; the
ACTION_MARGINALS: line emits those as a JSON distribution, or emits
an honest stub marker when only the GPU path was exercised. Pooled
Sharpe is computed from concatenated per-fold returns (CPU path) or
falls back to mean-of-fold-Sharpes with a warning.
Smoke test runs 30 surrogate seeds + 1 trained run via evaluate_baseline
subprocess, asserts trained pooled Sharpe exceeds the 95th percentile of
surrogate Sharpes. Test is #[ignore]'d and requires a trained checkpoint
at /workspace/output/dqn_fold0_best.safetensors (Phase 3 deliverable);
FOXHUNT_SURROGATE_CKPT env var overrides for local testing. Uses the
compiled target/release/examples/evaluate_baseline if present, otherwise
falls back to cargo run.
Will pass once Phase 3 produces a checkpoint.
Per-sample contribution arrays for additive terms (micro, loss_aversion,
segment_patience) + sample-selector flags (cf_flip), summed host-side
and divided by total reward magnitude. Trail rate reuses Task 0.5
buffers (peek without reset).
PopArt slot: REAL — computed from FusedTrainingCtx::read_popart_variance()
as |Δvar| / |var_pre| tracked across epochs. Returns 0.0 during warmup
(first 10 batches of PopArt stats) or when PopArt is disabled.
segment_patience slot: STUB (0.0) — the kernel has no patience term
wired. The sparse-reward formula comment (experience_kernels.cu:1049)
mentions trade_return * patience_mult but no multiplier lives in the
live code path. Slot preserved for HEALTH_DIAG schema stability; will
carry real signal if/when segment patience ships.
No atomicAdd — each thread writes its own [i*L+t] slot (or [cf_off]
for cf_flip). Buffers reset via memset_zeros after readback. Order
matters: reward_contrib_fractions peeks trail/traded buffers without
reset so trail_fire_and_hold_per_mag can read+reset them afterward.
HEALTH_DIAG reward_contrib group now carries real signal for 5 of 6
slots. Smoke test reward_component_audit_summary returns the cached
latest reading (was all-zero stub).
Per-sample i32/f32 output arrays sized [alloc_episodes*alloc_timesteps]
populated by experience_env_step at every (i,t). Magnitude bin computed
from pre_trade_position (the position that just got trailed-out), NOT
post-enforcement actual_mag_core (which is always 0 on trail fire).
Host-side deterministic reduction in trail_fire_and_hold_per_mag().
NO atomicAdd — each thread owns its own [i*L+t] slot. Buffers reset
via memset_zeros after each epoch readback.
Updates HEALTH_DIAG trail group with real trail_fire_rate[3] and
hold_at_exit_mean[3] (Quarter/Half/Full bins = 0/1/2).
Adds smoke test asserting the reward-contribution diagnostic produces
finite, non-negative values for all 6 reward terms (popart, cf_flip,
trail, micro, loss_aversion, segment_patience).
The smoke test validates the DIAGNOSTIC INFRASTRUCTURE — that the
per-term contribution accessor doesn't return NaN/Inf/negative values
that would break log parsing in Phase 1 audit. The actual triage
(KEEP/DELETE per term per spec §5.2) happens against a multi-fold
L40S training log in Phase 1, not in this smoke.
DQNTrainer.reward_component_audit_summary() returns [f32; 6] with
explicit measurement-class semantics per spec §4.1:
- additive: fraction of |total reward|
- transform (popart): |delta| / |pre|
- sample-selector (cf_flip, trail): firing rate
Currently returns zeros — Task 0.8 wires kernel-side instrumentation
to populate real values. Smoke remains green throughout (zeros pass
finite + non-negative gates), and ensures any future kernel-side
breakage that produces NaN/Inf is caught immediately.
Wires Track 1 noisy_mag/noisy_dir + Track 4 sigma_mean fields with the
mean |sigma| across each action branch's NoisyLinear fc + out layers.
Branch 0 = direction, Branch 1 = magnitude. H7 detection signal — if
magnitude branch has 2x larger σ than direction, NoisyNets noise is
dominating the magnitude head's effective signal.
Cross-crate API chain (no shortcuts):
* ml-dqn::noisy_layers::NoisyLinear::sigma_mean() -> mean |weight_σ| + |bias_σ|
* ml-dqn::branching::BranchingDuelingQNetwork::branch_noisy_sigma_mean(idx)
averages fc + out σ for the named branch
* ml-dqn::dqn::DQN::branch_noisy_sigma_mean(idx) — None-tolerant proxy
* ml::trainers::dqn::config::DQNAgentType::branch_noisy_sigma_mean(idx)
delegates through primary_head
* HEALTH_DIAG reads via self.agent.read().await at epoch boundary
Pinned-readback pattern not used here — NoisyLinear weights live in
candle-managed CudaSlices, not flat trainer params buffer. Per-call
dtoh of ~256 + 768 floats × 2 layers × 4 branches = ~8KB total per
epoch. Negligible.
Track 0.10 (exploration entropy + sigma_mean) is now COMPLETE — the
sigma_mean field that was previously stubbed is now real.
Task 0.6 remains partial: VSN mask (vsn_mag, vsn_dir) and target drift
(drift_mag, drift_dir) still stubbed — they require separate accessors
on different layer types (VSN module + target_params_buf reductions).
Wires grad_ratio_mag_dir HEALTH_DIAG field with the magnitude head's
gradient L2 norm divided by the direction head's. H4 detection signal
(magnitude-head gradient starvation).
Implementation:
* GpuDqnTrainer.grad_readback_pinned_ptr — pinned host buffer sized
TOTAL_PARAMS f32, allocated once at construction. Plain pinned
(not device-mapped) — written via memcpy_dtoh.
* per_branch_grad_norms() -> [f32; 4]: sync stream, dtoh full grad
buffer to pinned host, compute L2 norm per branch slice using
compute_param_sizes + padded_byte_offset (branch tensors at indices
8-11 / 12-15 / 16-19 / 20-23 for direction/magnitude/order/urgency).
* grad_ratio_mag_dir() -> f32: convenience accessor for HEALTH_DIAG.
* fused_training::FusedTrainingCtx::grad_ratio_mag_dir(&mut) — exposes
via the wrapper used by training_loop.
Performance: pinned dtoh of ~2.7 MB (667K params × 4 bytes) takes ~0.5ms
on PCIe 4. Stream-sync cost is per-epoch (not per-step), acceptable for
diagnostic readback. Pageable-memory dtoh would be ~3-5x slower.
Drop free of the pinned buffer added alongside other pinned slots.
Per plan Task 0.4. Complete — no stubs.
Wires the eval_dist HEALTH_DIAG group with per-magnitude action distribution
read from the validation backtest. H10 detection signal — training-mode
entropy may look uniform while eval-mode argmax collapses to Quarter.
Implementation:
* GpuBacktestEvaluator::read_eval_action_distribution_per_magnitude()
reads actions_history_buf to host, decodes magnitude bin per sample
(action layout: dir*27 + mag*9 + ord*3 + urg → mag = (a/9) % 3),
returns [Quarter, Half, Full] normalized over non-skipped entries.
* DQNTrainer.last_eval_magnitude_dist field, populated in metrics.rs
after each validation backtest. Non-fatal on readback error.
* HEALTH_DIAG eval_dist group [eq, eh, ef] now shows real values.
Per plan Task 0.7. Complete — no stubs.
Adds --max-folds CLI flag to train_baseline_rl (0 = no cap). When set,
truncates the generated walk-forward fold list to the first N folds.
Used by the new multi_fold_convergence smoke test to run a 3-fold × 20-epoch
local variant of the Phase 3 L40S 6-fold × 50-epoch validation gate
(~5 min on RTX 3050 vs ~1 hour on L40S).
Test asserts:
* train_baseline_rl subprocess exits 0 (no NaN/Inf)
* ≥ 2/3 folds produce dqn_fold{N}_best.safetensors checkpoint
(checkpoint is only written when Best Sharpe improves during the
fold, so presence = policy learned something on that window)
Per plan Task 0.15 at docs/superpowers/plans/2026-04-21-policy-quality.md.
Adds delta-based adaptive-controller fire detection. A controller 'fired'
in epoch N iff its observable output changed from epoch N-1 beyond a small
numerical threshold. Applied to all 6 controllers per spec §5.3:
- anti_lr (threshold 1e-10 on LR)
- tau (1e-6 on target-net tau)
- gamma (1e-4 on discount factor)
- grad_clip (1e-3 on adaptive clip threshold)
- cql_alpha (1e-5 on CQL pessimism weight)
- cost_anneal (1e-4 on tx-cost anneal factor)
Adds:
* DQNTrainer.prev_controller_values + .controller_fire_counts
* ControllerPrevValues (NaN sentinel on first epoch = no fire)
* ControllerFireCounts (running u32 per controller)
* pub fn controller_fire_rates_final() -> [f32; 6]
* HEALTH_DIAG 'controller' group shows per-epoch fire bools + max
running fire rate across all controllers (fire_frac)
Adds smoke test controller_activity.rs asserting no controller fires in
> 50% of epochs. Expected to FAIL on current main — anti-LR fires every
epoch per today's session observation. Documents the load-bearing issue.
Completes Tasks 0.9 + 0.13 per plan. No partial state — all 6 controllers
are delta-tracked, all fire rates exposed, the smoke test asserts against
all of them.
Asserts magnitude-branch action entropy stays meaningful during training:
- entropy @ epoch 5 (0-idx 4) ≥ 0.5
- entropy @ epoch 20 (0-idx 19) ≥ 0.3
Uses the ent_mag values populated in HEALTH_DIAG by commit ec035ca3a.
Adds DQNTrainer.explore_entropy_mag_history Vec + public accessor.
Expected to FAIL on current main (documents current collapse speed).
Per plan Task 0.14 at docs/superpowers/plans/2026-04-21-policy-quality.md.
Adds module-level shannon_entropy_normalized helper and computes per-branch
action entropy from the existing monitor.action_counts data:
* ent_mag = entropy over Quarter/Half/Full distribution
* ent_dir = entropy over direction-bin distribution (indices 0-2 / 3-5 / 6-8)
Normalized to [0, 1] so 0 = full collapse, 1 = uniform.
Does not complete Task 0.10: sigma_mean (NoisyNets σ overall) remains 0.0
until Task 0.6 wires the per-branch NoisyNets σ readback (separate commit).
Phase 0 Task 0.11 (and completes Task 0.3 action_dist wiring):
* DQNTrainer.last_magnitude_dist: [f32; 3] — [Quarter, Half, Full]
cached at end of each epoch from monitor.action_counts
* pub fn magnitude_action_dist_final() → [f32; 3] accessor
* new smoke test magnitude_distribution.rs — asserts F_Half/F_Full ≥ 5%
after 20-epoch train on fxcache smoke data
Expected to FAIL on current main (observed: F_Half ≈ 0.2%, F_Full ≈ 0.5%)
— documents the bug per spec §1.1. Will pass after Phase 2 fixes land.
Tracked as Task 0.11 in docs/superpowers/plans/2026-04-21-policy-quality.md.
Replaces 3 zero stubs in the HEALTH_DIAG mag group with real values
computed from monitor.action_counts. Quarter = indices 0,3,6 /
Half = 1,4,7 / Full = 2,5,8, all divided by total epoch actions.
Other 7 mag-group fields (q_full/half/quarter, var_scale, kelly_f,
avg_win_ratio, grad_ratio_mag_dir) still stubbed — require GPU readback
infrastructure (per-magnitude Q reduction, var_scale running mean,
per-branch grad norm) landed in follow-up commits.
Partial Task 0.3 per docs/superpowers/plans/2026-04-21-policy-quality.md.
User preference: everything lands on main, ~500 LOC is small enough
that feature-branch isolation isn't needed.
Spec changes:
* §2.3 outcome paths: drop "unmerged branch" vocabulary
* §3 architecture: main-branch timeline, no feat/, no wip/
* §4 Phase 0: commits on main, tag policy-quality-baseline at end
* §5 Phase 1: experiments are worktree edits reverted after, only
triage docs commit (to main)
* §6 Phase 2: commits land on main, author chooses granularity,
each leaves smokes green
* §7.3 outcome handling: iterate on main, baseline tag for rollback
* §7.4 closing: tag policy-quality-v1, update status
Plan changes:
* Task 0.1: verify-starting-state (not branch creation)
* Task 1.x: temp-edit-then-revert (not wip branches)
* Task 1.5: triage docs commit directly to main
* Phase 4: tag + close (not merge + cleanup)
* Rollback: one git reset --hard command
Addressed functional concerns from third review:
1. Surrogate-noise gate was statistically weak — "random Sharpe ≤ 0.5×
trained Sharpe on same val data" could be satisfied by sample
noise. Replaced with percentile-based test:
* Pool all 6 val folds (~600 trades) for statistical power.
* Run N=30 surrogate random-action rollouts (matched marginal
action distribution — so difference is state-action mapping,
not action frequency).
* Assert trained pooled Sharpe > 95th-percentile of surrogate
distribution. Explicit false-positive control.
2. Escalation path ("scope exceeds sub-project A v2") was hand-wavy.
Now concrete: failure after 3 iterations → paused, not closed →
triage spec opened with evidence summary + diagnosis (policy/
reward/state/architecture/data bottleneck) + explicit decision
(continue on feat/policy-quality or fork feat/policy-quality-v2).
Triage itself is a full brainstorming cycle.
3. HEALTH_DIAG §4.2 Track 1 fields only covered H1–H5 — H6–H10
needed their own detection signals. Added:
- trail_fire_rate_{quarter,half,full} (H6)
- hold_time_at_exit_{quarter,half,full} (H6)
- vsn_mask_{magnitude,direction} (H7)
- noisy_sigma_{mag_head,dir_head} (H7)
- target_drift_{mag_head,dir_head} (H8)
- action_dist_eval_{quarter,half,full} (H10)
Fixed inconsistencies between sections after the rev-2 feature-branch
refactor (sections had gotten out of sync):
1. §2 (success criteria) rewritten — was stale from rev 1. Now has
clean mandatory/soft split matching §7.2, corrects the trade-count
gate (per-fold not averaged), references the surrogate-noise gate.
New §2.3 enumerates the outcome paths.
2. §4 and §6 no longer say "commit on main" — both land on
feat/policy-quality per the feature-branch architecture. Phase 4
merges to main at §7.4.
3. §7.3 outcome handling rewritten — was written for single-branch
model ("Phase 2 commit is NOT reverted"). New wording matches
feature-branch semantics: failing mandatory = don't merge; failing
soft = merge + follow-up.
4. §6 smoke-tests rule now lists all six Phase 0 smokes, not just
Track 1/2.
5. §7.4 Phase-4 merge-strategy added — --no-ff merge commit, tag
policy-quality-v1, cleanup sequence documented.
6. §5.5 conflict-resolution added — tracks can reach contradictory
conclusions (e.g. T1 says fix, T2 says delete). Resolution rules:
mandatory-gate dominance, simplification wins, escalation.
7. §4.3 baseline metrics now committed to the feature branch (not
"investigation-only"), matching the crash-safety discipline.
8. Budget / risk register aligned — 5.5–6.5 hrs plan + 1 buffer,
hard-capped at 3 Phase-3 validation runs.
If H9 confirms and Phase 2 removes the magnitude branch, F_Half/F_Full
gates become vacuous. Substitute: direction bin distribution healthy
(F_Short + F_Long each >= 20% of non-Flat actions). Preserves the
'model actually trades' intent of Criterion A.
Addresses critique from second self-review:
1. SCOPE GAP fixed: spec now covers all 4 V7 categories the user asked
for. Added Track 3 (controllers) and Track 4 (exploration) with their
own smoke tests (controller_activity, exploration_coverage) and
HEALTH_DIAG fields. Previous version only had action-space + reward.
2. H1-H5 expanded to H1-H10 with 5 new codebase-specific hypotheses:
- H6: Regime-adaptive trailing stop closing Full-positions early
- H7: VSN / NoisyNets masking magnitude features
- H8: Target tau too slow for magnitude propagation
- H9: Data genuinely favors Quarter (no bug, delete magnitude branch)
- H10: Entropy regularization + argmax tie ordering bias
3. H3 REFRAMED from "Kelly bug" to design critique — Kelly as hard
multiplier vs soft gate. Fix sketch updated to offer replacement.
4. Feature branch + tag strategy replaces "single branch" containment.
feat/policy-quality isolates work from main until Phase 4 merge via
PR. git tag policy-quality-baseline gives named rollback anchor.
5. SURROGATE-NOISE CHECK added as MANDATORY validation gate — catches
look-ahead/leakage bugs that would otherwise produce false-positive
edges. Permanent smoke artifact.
6. Validation gate restructured: B + surrogate-noise MANDATORY;
A criteria soft; controller-activity soft. Partial success path
explicit per-gate.
7. Risk register expanded — added scenarios for B-mandatory failure
(feat branch stays unmerged), all-10-hypotheses-REJECTED (fall back
to H9 delete-branch), surrogate-noise failure (halt spec, open
diagnostic spec), multi-load-bearing-controllers (tech debt doc).
8. Budget corrected to 5.5-6.5 hrs L40S (was underestimated 3.5-4),
with per-track breakdown.
Spec now: 464 lines, 10 sections, covers 4 V7 tracks × ~10 hypotheses
+ 8 reward terms + 7 controllers + 5 exploration mechanisms = ~30
audit items total.
Critical fixes after honest review:
1. reward-audit measurement: split into 3 classes (additive / transform
/ sample-selector) — one-size-fits-all metric couldn't work for
PopArt normalization or CF flip mirroring
2. H1 detection: replace undefined "Sharpe/state" with concrete
forced-exploration protocol (epsilon=1.0 for 1 epoch at mid-train,
sample states × magnitudes, compare Sharpe-per-trade)
3. H4 scope: clarify this is the per-action-branch gradient (direction
vs magnitude head), NOT the per-loss-component (IQN/CQL/C51/Ens)
budget. Fix sketch updated to target the correct knob.
4. Validation gate: criterion B (multi-fold Sharpe + WinRate) is now
MANDATORY; criterion A gates are soft. Prevents "partial success"
hiding a policy-quality failure.
5. Trade-count gate: per-fold ≥100 on ≥5/6, not an average — 5 lean
folds pooled with 1 fat one should not pass.
Minor:
- Budget trimmed to 3.5–4 hrs L40S (was 6–10)
- L40S pool made explicit as target
- Phase 1 branches now pushed as wip/* for crash safety
- Hyperopt-params mismatch handling documented
Sub-project A of the production-readiness roadmap. Single-branch,
phase-gated: measurement substrate → parallel investigation tracks
(magnitude collapse + reward V7 audit) → single convergence commit →
6-fold × 50-epoch L40S validation.
Target criteria at close:
A (action dist): F_Half/F_Full each ≥15% on ≥4/6 folds, WinRate
55-65%, ≥100 trades/val window
B (multi-fold): Best Sharpe > 10 on ≥5/6 folds, WinRate > 55% on
≥5/6 folds
Greenfield checkpoint policy — any action-space or weight-shape change
is allowed. All confirmed bugs fix together in one Phase-2 commit.
V7 discipline gates inclusion (measurement confirms the bug), no
ranking across findings.
Out of scope for this spec: paper trading (sub-project C), canary
deployment (sub-project D), hyperopt re-runs, architectural changes
beyond action space.
Awaiting user review before invoking writing-plans.
Symmetric 5-window mean (f1359f3dc) filtered noise but also filtered
exploration. RL Sharpe is plentiful-bad and rare-good in early training,
so the mean always sees the plentiful side — controller locked at 0.3×
LR and the model couldn't escape its initial bad minimum
(train-v82b2: sharpe stuck at −8 to −21 throughout Fold 0, never
peaked positive like prior runs had).
Root fix: the two decisions have different evidence requirements.
* Boost LR: low cost if wrong (clamp caps runaway), high value if
right (kicks out of overfit). Accept weak evidence — ANY of the
last short_window epochs clearly positive fires the boost.
* Dampen LR: high cost if wrong (stuck model), low value if right
(stability we didn't need). Demand strong evidence — MEAN over
long_window must be clearly negative.
Both windows derive from anti_lr_warmup (one knob):
long_window = anti_lr_warmup (full window for sustained-bad)
short_window = anti_lr_warmup / 2 (half window for recent-good)
No new hyperparameter. Asymmetry is structural (max vs mean over
differently-sized windows), not tuned.
Verified: multi-trial smoke 5/5 pass, median_q_gap=2.77, mean_sharpe_ema=12.24.
Compared to symmetric smoothing (2.13 / 11.03) and raw-signal (1.13 / 2.94),
the asymmetric version is the best on all three multi-trial metrics.
Tie-breaking: "good wins" when both signals fire in the same step —
matches the controller's original intent (exploration over dampening)
and our diagnosis (model needs LR headroom to escape bad starting
states).
Every new commit triggered a full workspace recompile (~6min for a 2-file
change) in ensure-binary. The binary cache at /data/bin/$SHA is SHA-keyed
so it misses on every new commit. The cargo incremental cache at
/cargo-target/target is file-mtime keyed, which git checkout invalidates
whenever it retouches files.
sccache (already present in ci-builder image, verified via `sccache --version`
in Dockerfile) is content-hash keyed — identical source → identical hit
regardless of mtime. Sits in front of rustc via RUSTC_WRAPPER. Cache dir
on the same cargo-target-cuda PVC (SCCACHE_DIR=/cargo-target/sccache) so:
- Combines with cargo's target/ cache on one volume, no separate PVC
- Disk-local hits, no network round-trip per rustc invocation
- Survives across pods and commits — the git-checkout mtime flip that
breaks cargo's fingerprints doesn't affect content-hash caching
40G cache size cap — the cargo-target-cuda PVC is large enough that
sccache won't evict useful entries under normal use.
Expected behaviour: first build after this commit is still a full rebuild
(sccache cache empty). Subsequent builds for the same CUDA_COMPUTE_CAP
should hit compiled objects at near-zero cost for crates whose source
hasn't changed.
Applied to cluster via `kubectl apply -n foxhunt -f` — WorkflowTemplate
is live before the next ./scripts/argo-train.sh invocation.
Closes task #34.
The anti-LR adjuster (config.rs#24) reacts to Sharpe swings by multiplying
the learning rate — good Sharpe → ×3 to escape overfit minima, bad Sharpe
→ ×0.3 to stabilize. Previously it fed on raw `sharpe_history.last()`,
which is a single noisy epoch value. In 30-epoch L40S smoke (train-br8cb
Fold 0) per-epoch Sharpe oscillated between −20 and +30, so the controller
flipped multipliers every epoch and amplified its own input noise —
gradient norm spiked to 1.16M at Epoch 18 from a baseline of ~3000.
Fix: feed the controller a rolling mean of the last `anti_lr_warmup`
epochs (the same knob that already gates the controller on — no new
hyperparameter). This filters per-epoch oscillation at the frequency the
anti-LR logic wants to react on (multi-epoch trends), while still letting
genuine sustained improvement or degradation trigger adjustments. Keeps
the original 3.0 / 0.3 multipliers and [0.1, 5.0] clamp — the problem
was the signal, not the magnitudes.
Why temporal instead of tightening magnitudes:
* Shrinking 3.0/0.3 → 1.5/0.7 reduces the symptom but keeps the
structure — still amplifies noise, just less.
* Smoothing removes the noise before the controller sees it, so the
controller stays as aggressive as designed.
* No new magic numbers — reuses anti_lr_warmup (=5) for both
"don't-tune-yet" and "this-is-a-stable-horizon".
Verified: multi-trial smoke 5/5 finite, 4/5 q_pass, median_q_gap=1.13.
Slight reduction vs the fold-reset baseline (2.13) is expected — the
smoother trades responsiveness for stability. Real validation is the
L40S 20-epoch behaviour test queued after this commit.
Fold 1 of the 50-epoch L40S smoke (train-92xbj) NaN'd at Epoch 12 even
with the config.v_range clamp from commit 679881fa5. Root cause: the
adaptive C51 v_range state (eval_q_mean_ema, eval_q_std_ema,
eval_ema_initialized, and the pinned eval_v_range_pinned buffer) was
initialised at construction but never reset at fold boundaries.
Reset checklist before this commit (fused_training::reset_for_fold):
- graph children destroyed ✓
- shrink-and-perturb applied to weights ✓
- Adam momentum reset ✓
- replay buffer cleared (via DQNTrainer::reset_for_fold) ✓
- PopArt running stats reset ✓
- eval_v_range EMA state ✗ ← the gap
Fold N+1 therefore inherited Fold N's final tight atom support. When
Fold 1's new data distribution produced Q-values that didn't fit Fold
0's narrow range, TD errors saturated the atom bins; gradient norm
blew up geometrically (170K → 2.3e15 → inf) and the NaN guard bailed
the trainer out at Epoch 12.
Fix: new GpuDqnTrainer::reset_eval_v_range_state() zeros the EMAs,
clears eval_ema_initialized (so the next update_eval_v_range call
reinitialises from observed statistics), and restores the pinned
range buffer to [config.v_min, config.v_max] so any kernel that reads
the range before the first update sees the wide safe support. Invoked
right after reset_adam_state() in fused.reset_for_fold().
Verified: multi-trial smoke 5/5 pass, median_q_gap=2.13, mean_sharpe_ema=11.03
(best observed on this branch). No regression on single-fold training.
50-epoch L40S retest will confirm Fold 1 converges without divergence.
Closes task #33.
evaluate_baseline looks for <output_dir>/norm_stats_fold{N}.json per fold
(matching the convention written by train_baseline_supervised.rs:812).
train_baseline_rl never produced this file — the fxcache has a single
<hex_key>.norm_stats.json written by precompute_features, which RL
training uses implicitly for pre-normalized features but never copies to
the per-fold paths evaluate wants. Result (L40S train-7r9zf):
Error: NormStats not found at /workspace/output/norm_stats_fold0.json
- cannot evaluate without training-set statistics (computing from test
data would introduce lookahead bias). Run training first to generate
this file.
WARN: Evaluation failed, continuing
Evaluate fails gracefully but no report is produced.
Fix: new helper `fxcache::norm_stats_path_for_key(data_dir, override, cache_key)`
returns the canonical <hex_key>.norm_stats.json path that sits next to
the .fxcache file. train_baseline_rl resolves this once after the fxcache
load (falling back to None if the key is zero — i.e., DBN fallback path
with no precomputed stats) and copies it into the output dir as
norm_stats_fold{N}.json for every fold processed.
The RL-side fxcache norm stats are fold-independent (one z-score per
cache key, applied to all walk-forward windows), so the per-fold copies
are intentional duplicates of the same file — this matches evaluate's
per-fold lookup semantics without diverging from supervised convention.
Closes task #32.
Root cause of the L40S 10-epoch smoke Fold 1 loss explosion
(train-7r9zf, final_loss=1.6e16, grad_norm=62B): update_eval_v_range
recomputed [v_min, v_max] every epoch from eval_q_mean_ema ± gap_width
with NO upper bound, while the bounded theoretical range in
config.v_min/v_max (derived from reward_scale/(1-gamma)*1.2, clamp [20, 300])
was only honored at buffer initialization.
The runaway path (visible in the epoch-by-epoch trace):
Epoch 3: Q-range=[-10, 10] — pegged at config bounds ✓
Epoch 4: Q-range=[-10, 14.16] — Q escapes the configured window
Epoch 5: Q-range=[-10, 22.30], loss=43,660 (up from 30)
Epoch 10: loss=1.6e16, grad_norm=62B, Fold 1 converged=false
Positive feedback loop:
Q overestimate → eval_q_mean_ema drifts up → v_max follows unbounded
→ C51 atom support widens → TD targets grow → Q targets grow
→ network chases → gradient explodes → repeat.
Fold 0 stayed in the safe region only because its reward dynamics never
pushed Q far enough to escape ±10; Fold 1's larger window (3.7M bars vs
2.9M) contained the trigger event.
Fix: gpu_dqn_trainer::update_eval_v_range now clamps BOTH the half-width
to (config.v_max - config.v_min) / 2 AND the centre to the range
[config.v_min + half, config.v_max - half], guaranteeing the adaptive
window always fits inside the theoretical bounds. Single source of truth
— the config-level clamp is now actually authoritative.
Verified: multi-trial smoke 5/5 pass, median_q_gap=2.07 (beats 2.00
baseline from commit cb2015ab2), mean_sharpe_ema=7.82. No regression on
Fold 0's normal-range behaviour.
Closes task #31.
Two bugs caught by the L40S smoke (train-qhgj6) that couldn't surface on
local RTX-3050 single-fold runs:
1. PER dtoh inside CUDA Graph capture (Fold 1 crash)
Failure: CUDA_ERROR_STREAM_CAPTURE_INVALIDATED at per_prefix_scan on
Fold 1 re-capture. Chain: fused_training parent graph captures →
memcpy_dtoh + cuStreamSynchronize in gpu_replay_buffer::update_priorities_gpu
(health<0.8 diversity path) poisons the stream → subsequent per_sample
kernel on the same stream sees an invalidated capture context.
The prior comment claimed "runs once per epoch, DtoH cost acceptable"
— wrong, it runs every priority update when health<0.8 (common during
Fold handoff when health_cache is re-seeded low). Any dtoh inside
capture invalidates regardless of latency.
Proper fix (no shortcut):
* New kernel actions_sum_scale_reduce_u32 — single-block deterministic
tree reduction over sample_actions (u32) → writes (sum*1000)/n as i32
to a device-accessible slot. No atomics (consistent with the 1/N
determinism policy from commit c82386500).
* mean_action_scaled storage is pinned + device-mapped (cuMemAllocHost
+ cuMemHostGetDevicePointer — same pattern as rng_step_dev_ptr and
size_dev_ptr elsewhere in the file). Zero-copy between host and
device, graph-safe, no explicit free needed (process-exit cleanup,
matches existing pattern).
* pow_alpha_diverse_f32 now takes const int* mean_action_scaled_ptr
and does a plain global load — NOT __ldg. The read-only cache used
by __ldg is not guaranteed coherent with device-mapped host memory;
multi-trial smoke regression caught it (median q_gap collapsed
from 2.0 → 0.15 with __ldg, recovered to 2.8 with plain load).
Verified: multi-trial smoke 5/5 pass, median_q_gap=2.80 (beats 2.00
baseline), Best Sharpe peaks 19-30 per trial. No stream capture
invalidation.
2. evaluate step CLI drift in Argo template
evaluate_baseline's Args struct uses --models-dir and --output (single
file path). Template was passing --checkpoint-dir and --output-dir,
causing clap to reject the invocation. Fixed argument names + added
mkdir for the eval subdir + updated the comment to pin the source of
truth for future drift catches.
Both fixes are graph-capture-clean and match the "wire properly or delete"
discipline. No masking, no feature flags, no dead params.
Restores regime adaptation to the 0.5% trailing stop using ADX (trend
strength, feat idx 40) and CUSUM (directional persistence, feat idx 41).
Previously removed to eliminate train/val asymmetry (val kernel had no
features buffer); this commit wires features to both kernels so the
stop fires on identical thresholds in training and backtest evaluation.
Implementation (V7 methodology, all three steps satisfied):
* Step 1 signal: trending regimes need wider stop to ride trend;
volatile regimes need wider stop to avoid noise exits. Original
formulation from reward v5 (commit 263997ad3).
* Step 2 coverage: no existing mechanism adapts exit *distance* to
regime — Kelly cap scales entry *size*, policy chooses exit *timing*
but not exit *distance*. So this is a genuine gap.
* Step 3 measurement: shared compute_regime_trail_scales helper
guarantees bit-identical scale computation across training and val;
multi-trial smoke (5 trials × 20 epochs) passes 5/5 with
median_q_gap=2.00, Best Sharpe 27.05 in last trial. No regression
vs fixed-width (pre: median_q_gap=2.24; post: 2.00, within variance).
Shared helper (trade_physics.cuh::compute_regime_trail_scales):
trend_scale = min(2.5, 1 + max(ADX - 0.25, 0) * 2.0)
vol_scale = min(2.5, 1 + max(|CUS| - 0.50, 0) * 2.0)
Pass NULL features to fall back to 1.0/1.0 (fixed-width). Thresholds
match existing regime_conditional convention (adx>0.25=Trending).
Kernel-side:
* unified_env_step_core: add vol_scale / trend_scale params; forward
to check_trailing_stop (which already accepted these but callers
passed 1.0/1.0). Updated step-6 docstring.
* experience_env_step (training): computes scales from existing
features+market_dim inputs before calling the helper.
* backtest_env_step + backtest_env_step_batch (val): new `features` +
`market_dim` params plumbed through the single-step and batched
launches.
Rust-side (gpu_backtest_evaluator.rs):
* launch_env_step: pass self.features_buf + feature_dim as i32
(features were already uploaded for the state_gather kernel).
* Batched launcher: same.
Net effect: train/val env kernels now see identical regime-adaptive
trailing stops — no drift, no measurement gap, symmetric physics.
Verified:
cargo check/build clean (release); multi-trial smoke 5/5 pass.
Closes task #25.
compute_expected_q had 2 atomicAdds accumulating per-block entropy + utilization
into atom_stats[0..1]. These feed HEALTH_DIAG's atoms=X%ent/Y%util field AND
the adaptive_atom_positions kernel (which reshapes C51 atom positions based on
utilization) — so non-determinism here propagated into training dynamics.
Replaced with standard 2-phase deterministic reduction:
Phase 1 (inside compute_expected_q):
Per-block warp→block shared-mem reduce unchanged, but the final write is
atom_stats_block_sums[blockIdx.x * 2 + i] = value (unique slot per block,
no atomic). Kernel signature: `atom_stats` param renamed to
`atom_stats_block_sums` to reflect new semantics.
Phase 2 (new kernel atom_stats_finalize):
Sequentially sums block_sums[0..num_blocks*2] into atom_stats[0..1].
Launch grid=(1,1,1) block=(1,1,1) — bit-stable across runs.
Rust glue:
- New atom_stats_block_sums_buf [max_blocks * 2] allocated at trainer init.
max_blocks = ceil(batch_size / 256); smoke uses 1, production 64.
- New atom_stats_finalize_kernel loaded from experience_kernels cubin.
- populate_q_out() now launches phase 1 + phase 2 sequentially.
- Other launch sites (replay_forward_for_q_values, compute_denoise_target_q)
already passed NULL atom_stats and don't need changes — the kernel skips
accumulation on NULL.
This was the FINAL training-path atomicAdd call site. Post-commit, the entire
crates/ml/src/cuda_pipeline/ has ZERO atomicAdd call sites — full CUDA-level
determinism for the training path (cuBLAS algorithm selection remains the
main remaining nondeterminism source, mitigated via CUBLAS_WORKSPACE_CONFIG
for tests).
Smoke verification (post-commit, 5-trial multi-trial):
q_gaps: 2.24 / 0.98 / 2.75 / 5.58 / 1.08 (all > 0.02, 5/5 pass)
Best Sharpe: 19.41 / 38.71 / 25.33 / 35.94 / 24.22 (median 25.3)
mean_degradation = -10.66 (NEGATIVE means sharpe_ema IMPROVED over epochs)
All assertions pass.
Session atomicAdd tally:
69 string occurrences → 13 real call sites (most "hits" were comments).
Today removed: 4 (CQL barrier+IB) + 2 (monitoring_reduce) + 2 (trade_stats
+ dqn_utility) + deleted-kernel (ensemble_diversity's atomic + G6/G10's
atomics removed with scaffolding) + 2 (atom_stats, this commit) = 13 total.
Zero remaining.
Files touched:
crates/ml/src/cuda_pipeline/experience_kernels.cu (+38 / -9 net)
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs (+40 / -10 net)
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Completes Phase 3 of the env-unification effort: experience_env_step (training),
backtest_env_step (val single-step), and backtest_env_step_batch (val batched)
all call the same __device__ helper `unified_env_step_core` in trade_physics.cuh.
Drift between train/val becomes structurally impossible at compile time — any
change to the canonical step applies atomically to both.
Three structural fixes were required to make the training port work:
1) TWO max_position params in unified_env_step_core
Training scales `max_position` by Var[Q] (Kelly-from-atoms variance sizing)
→ `effective_max_pos`. The original training kernel used this for
compute_target_position_4branch but the UNSCALED `max_position` for
apply_kelly_cap and execute_trade. Initial port collapsed both into one
helper arg, tightening Kelly cap aggressively → tiny positions → Q-values
converge → q_gap=0.00 collapse. Split into `max_position_target` (scaled)
and `max_position_physics` (unscaled). Val passes the same value for both.
2) Remove double hold_time update
Helper now owns `update_hold_time(...)` inside the step. Training's
inline `hold_time = update_hold_time(...)` at line ~1574 was a double-
increment. Fix: capture `saved_hold_time` BEFORE helper for the
patience-multiplier in segment reward, then trust the helper's in-place
update of hold_time. Segment-hold-time semantics preserved.
3) Remove triple Kelly-stat update
Helper calls `record_kelly_trade_outcome(...)`. Training ALSO had Kelly
stat updates inside the reversing_trade block (line 1558) and the
exiting_trade block (line 1638). TRIPLE update → win_count/loss_count/
sum_wins/sum_losses inflated 3× → Kelly cap tightened proportionally →
positions starved → Q-gap collapse. Fix: only the helper now touches
Kelly win/loss/sum_wins/sum_losses. Training still updates its own
sum_returns / sum_sq_returns (continuous-Kelly stats, distinct
buffers not touched by the helper).
Smoke test result (RTX 3050 Ti, 20 epochs, single-trial):
Before port: Best Sharpe ~15-25 (varies, range 10-31)
After broken port: Best Sharpe 1.17 (q_gap=0.00 collapse)
After fix: Best Sharpe 39.25 (HIGHEST ever recorded)
q_gap 0.00 → 0.27 (growing, healthy separation)
sharpe_ema trajectory: -5.3 → 12.7 → 26.9 (rising)
Multi-trial (partial visible): q_gap rising 0.00 → 0.86 by epoch 6,
sharpe_ema 15-25 consistently positive. Both tests passed.
Training kernel now has the CORE STEP in a single helper call — the
kernel body downstream continues to handle training-only concerns
(counterfactuals, plan_params, reward shaping via shaping_scale,
saboteur via exploration_scale, replay buffer writes).
Net −99 LOC across three kernels replaced by shared helper calls.
Files touched:
crates/ml/src/cuda_pipeline/trade_physics.cuh (+15/-5)
crates/ml/src/cuda_pipeline/backtest_env_kernel.cu (+53/-108 net)
crates/ml/src/cuda_pipeline/experience_kernels.cu (+48/-102 net)
docs/superpowers/specs/2026-04-21-unified-train-val-env-design.md
(Phase 2 documented as helper-extraction; Phase 3 kernel-deletion
remains deferred — both kernel entry points kept for their distinct
threading models; the physics is what's now truly unified.)
Verified: cargo check + smoke test pass.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Core of the Phase 3 unification: instead of writing a brand-new
unified_env_kernel.cu (3-day rewrite per the design doc), extract the
canonical step logic into a single __device__ __forceinline__ helper
that BOTH kernels call. Drift between train/val becomes structurally
impossible — any change to the helper applies to both atomically.
unified_env_step_core (trade_physics.cuh) encapsulates:
1. Action decode (4-branch: dir, mag, order, urgency)
2. Hold action passthrough
3. Target position via compute_target_position_4branch
4. Margin cap (apply_margin_cap)
5. Kelly cap with health-coupled safety (apply_kelly_cap)
6. Trailing stop check (fixed 0.005/1.0/1.0 — regime-adaptive remains
a deferred gem, see trade_physics.cuh comment block)
7. execute_trade with sqrt-impact (spread_scale = -1.0)
8. record_kelly_trade_outcome (Kelly stats update)
9. entry_price update (new entry / reversal / flat)
10. hold_time tick via update_hold_time
11. step_return = (new_value − prev_equity) / prev_equity (PURE P&L)
12. Post-enforcement actual_dir + actual_mag (for actions_history)
Pass-by-pointer for all mutable state (position, cash, entry_price,
hold_time, max_equity, Kelly stats). Outputs: step_return, new_value,
prev_position_sign, actual_dir, actual_mag, trail_triggered.
Ported backtest_env_step (single-step variant) to call the helper —
replaced ~100 lines of inline step logic with a single call. Kept the
capital-floor pre-check / post-check / actions_history stitching as
per-kernel logic (output buffer formats differ between train and val,
so these stay per-kernel).
Files touched:
crates/ml/src/cuda_pipeline/trade_physics.cuh (+172)
crates/ml/src/cuda_pipeline/backtest_env_kernel.cu (-99 / +27 net)
Follow-up in a separate commit:
- Port backtest_env_step_batch (same refactor, batched variant)
- Port experience_env_step to call unified_env_step_core
(subset — training's body has many more layers: counterfactuals,
plan_params, reward shaping bundle — all of which STAY in the
caller; only the core step gets unified)
Verified: cargo check passes. Smoke test in flight to confirm
behavioral equivalence (should be bit-identical — same arithmetic in
same order, just refactored into a helper).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Stepping stone toward the full unified_env_kernel: add exploration_scale +
shaping_scale ptr params to backtest_env_step + backtest_env_step_batch
signatures. Both are NO-OP in backtest today (val has no saboteur, no
plan_params, no counterfactuals; step_returns are pure P&L since commit
1ffdf38dd) — but the signature now matches what the unified kernel will
take, so a future swap is mechanical.
Updated launch sites in gpu_backtest_evaluator.rs to pass NULL (0u64)
for both pointers — kernel ignores via (void) suppression.
Files touched:
crates/ml/src/cuda_pipeline/backtest_env_kernel.cu (+22)
crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs (+8)
Verified: cargo check passes.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
