4399a56d7612eae786f97f4b57985367cea2e7d3
4064 Commits
| Author | SHA1 | Message | Date | |
|---|---|---|---|---|
jgrusewski
|
4399a56d76 |
test(smoke): strengthen weak assertions across DQN smoke suite
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> |
||
|
jgrusewski
|
ed4b30b493 |
fix(smoke): controller_activity — V7 intervention-based fire detection
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> |
||
|
jgrusewski
|
a5f23b28f0 |
fix(diag): move GPU summary download before HEALTH_DIAG emit
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
|
||
|
jgrusewski
|
2fb30f098e |
fix(monitoring): cascade 9→12 action bins for 4-direction action space
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.
|
||
|
jgrusewski
|
2ac956298b |
fix(data-loading): truncate OFI in lockstep with bars on max_bars path
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. |
||
|
jgrusewski
|
83d524f866 |
diag(policy-quality): remediate stub return values per no-stubs rule
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.
|
||
|
jgrusewski
|
aa16ca31f2 |
diag(policy-quality): finish Tasks 0.3 + 0.6 partial wirings
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
|
||
|
jgrusewski
|
e1ac2c7578 |
docs(policy-quality): Task 0.17 baseline metrics scaffold (PENDING GPU run)
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.
|
||
|
jgrusewski
|
f4ea0f6d13 |
test(policy-quality): Task 0.16 surrogate_noise_check smoke (mandatory gate)
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. |
||
|
jgrusewski
|
bf8415c5d6 |
diag(policy-quality): Task 0.8 reward-term contribution diag (Track 2)
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). |
||
|
jgrusewski
|
48d962ee24 |
diag(policy-quality): Task 0.5 trailing-stop per-mag fire/hold (H6 signal)
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). |
||
|
jgrusewski
|
93d8c5ae4a |
test(policy-quality): Task 0.12 — reward_component_audit smoke
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. |
||
|
jgrusewski
|
999bb2fa0c |
diag(policy-quality): NoisyNets σ per branch — partial Task 0.6, completes 0.10
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).
|
||
|
jgrusewski
|
bb42c99636 |
diag(policy-quality): Task 0.4 — per-branch grad norm via pinned readback
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.
|
||
|
jgrusewski
|
0310b1d1eb |
diag(policy-quality): Task 0.7 — eval-mode action distribution (H10 signal)
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.
|
||
|
jgrusewski
|
4e1a937cec |
feat+test(policy-quality): Task 0.15 — multi_fold_convergence smoke + --max-folds
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.
|
||
|
jgrusewski
|
93471d85a2 |
diag+test(policy-quality): Task 0.9 + 0.13 — controller fire-rate tracking
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.
|
||
|
jgrusewski
|
9004c9b0a2 |
test(policy-quality): Task 0.14 — exploration_coverage smoke
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
|
||
|
jgrusewski
|
ec035ca3a6 |
diag(policy-quality): wire per-branch action entropy (ent_mag, ent_dir)
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). |
||
|
jgrusewski
|
691a7669a3 |
test(policy-quality): add magnitude_distribution smoke + accessor
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.
|
||
|
jgrusewski
|
6ce0382cf7 |
diag(policy-quality): Task 0.3a — wire action_dist_* from monitor
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. |
||
|
jgrusewski
|
592aaaf505 |
diag(policy-quality): scaffold HEALTH_DIAG fields for 4-track V7 audit
All new fields wired with zero/false values. Later tasks replace zeros with real GPU-sourced measurements. Per plan Task 0.2 at docs/superpowers/plans/2026-04-21-policy-quality.md. 37 new field placeholders across 4 V7 tracks: mag [10 f32] — Track 1 magnitude diagnosis (H1-H5) trail [6 f32] — Track 1 trailing stop per bin (H6) noisy [6 f32] — Track 1 VSN/NoisyNets/target-drift (H7-H8) eval_dist [3 f32] — Track 1 eval-mode action dist (H10) reward_contrib [6 f32] — Track 2 reward component audit controller [6 bool + 1 f32] — Track 3 firing rates explore [3 f32] — Track 4 entropy + sigma No behavioral change — all values zero/false until later tasks wire real readbacks. cargo check clean. |
||
|
jgrusewski
|
944fecd913 |
docs(policy-quality): main-branch workflow — cut feat/wip branches
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
|
||
|
jgrusewski
|
1456094a5a |
plan(policy-quality): implementation plan for 4-track V7 audit
Produced by superpowers:writing-plans from the approved spec at
docs/superpowers/specs/2026-04-21-policy-quality-design.md.
Plan structure:
* Phase 0 (tasks 0.1-0.17): measurement substrate — TDD-detailed,
exact code + commands. Produces 6 new smoke tests + ~25 HEALTH_DIAG
fields across 4 tracks.
* Phase 1 (tasks 1.0-1.5): parallel investigation tracks — process
tasks (run instrumentation, record evidence, classify hypotheses).
No code to pre-write; output is 4 triage documents + a Phase-2
planning doc.
* Phase 2 (template only): convergence fix. Concrete tasks produced
by a Phase-2-specific plan AFTER Phase 1 triage — can't pre-write
the fixes without knowing which hypotheses confirm.
* Phase 3 (tasks 3.1-3.4): L40S validation gates, surrogate-noise,
branch decision (merge or iterate).
* Phase 4 (tasks 4.1-4.4): merge via --no-ff, tag, cleanup.
* Rollback procedure: git reset --hard policy-quality-baseline.
|
||
|
jgrusewski
|
ad677466c8 |
design(policy-quality): revision 4 — functional gaps
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)
|
||
|
jgrusewski
|
b5b70fdd8e |
design(policy-quality): revision 3 — internal consistency pass
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.
|
||
|
jgrusewski
|
5552517f61 |
design(policy-quality): caveat Criterion A for H9 (delete-magnitude-branch) case
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. |
||
|
jgrusewski
|
26be2b8abe |
design(policy-quality): revision 2 — full 4-track V7 + outside-the-box
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. |
||
|
jgrusewski
|
bee81b4967 |
design(policy-quality): self-review revisions
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
|
||
|
jgrusewski
|
3c24c7f49e |
design(policy-quality): V7 audit + magnitude collapse fix spec
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.
|
||
|
jgrusewski
|
4085831452 |
fix(dqn): asymmetric anti-LR — fast response to good, slow to bad
Symmetric 5-window mean (
|
||
|
jgrusewski
|
8afe9562e5 |
infra(argo): wire sccache into ensure-binary — PVC-local per-crate cache
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.
|
||
|
jgrusewski
|
f1359f3dcc |
fix(dqn): temporal smoothing for anti-intuitive LR controller
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.
|
||
|
jgrusewski
|
a1346dac1b |
fix(dqn/fold-boundary): reset adaptive v_range EMAs between folds
Fold 1 of the 50-epoch L40S smoke (train-92xbj) NaN'd at Epoch 12 even
with the config.v_range clamp from commit
|
||
|
jgrusewski
|
f988ca384c |
fix(train): emit per-fold norm_stats.json for evaluate_baseline
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.
|
||
|
jgrusewski
|
679881fa53 |
fix(dqn/c51): clamp adaptive v_range to config bounds — Fold 1 explosion
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
|
||
|
jgrusewski
|
932ac2bda8 |
fix(graph-capture): eliminate dtoh in PER diversity path + align evaluate CLI
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
|
||
|
jgrusewski
|
ed317d6f89 |
Merge wip/env-unification: Phase 3 env unification + determinism + gems
Brings the env-unification branch (26 commits) into main:
* Phase 3 env unification: one __device__ helper (unified_env_step_core)
called by both training and val kernels — structural train/val parity
* Determinism: all training-path atomicAdds removed, CUBLAS_WORKSPACE_CONFIG
set for smoke tests, 2-phase deterministic reductions
* V7 gems: G12 predictive_coding wired (6× variance reduction), G6/G10
measured sub-noise and deleted, regime-adaptive trailing stop restored
with symmetric train/val wiring
* Multi-trial TD-propagation smoke test for stability measurement
* Measurement-honest step_returns (pure P&L, no shaping drift)
Smoke: 5/5 trials pass, median_q_gap=2.00, Best Sharpe 27 locally (RTX).
L40S validation pending.
|
||
|
jgrusewski
|
cb2015ab2f |
gem(trail): regime-adaptive trailing stop — symmetric train/val wiring
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
|
||
|
jgrusewski
|
c823865008 |
determinism(atom_stats): 2-phase reduction — final training-path atomicAdd removed
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>
|
||
|
jgrusewski
|
d6d846f896 |
cleanup(gems): delete G6/G10 scaffolding — 263 LOC of measured-sub-noise dead code
Follow-up to commit |
||
|
jgrusewski
|
69211af40b |
phase3(env-unification): full unification — all 3 kernels now call unified_env_step_core
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>
|
||
|
jgrusewski
|
a3b6bc2f3b |
phase3(env-unification): extract unified_env_step_core to trade_physics.cuh + port val
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>
|
||
|
jgrusewski
|
bd022154f4 |
phase3(env-unification): ABI-align backtest signature with future unified kernel
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
|
||
|
jgrusewski
|
1961857c22 |
gem(G6+G10): measurement says redundant — unwire (V7 methodology applied)
Empirical data from commit |
||
|
jgrusewski
|
61ab27ff34 |
gem(G6+G10): graph-safe forward + HEALTH_DIAG readback (measurement-first)
Following the V7-gem methodology you flagged: instead of blindly wiring
backward gradients for two possibly-redundant features (G10 redundant
with spectral_norm, G6 redundant with NoisyNets/ensemble-KL), this commit
makes both forward-only and exposes their values via HEALTH_DIAG so we
can MEASURE whether they're producing meaningful signal before committing
to the full backward integration.
Three coordinated changes:
1) G10 temporal_consistency_penalty rewritten graph-safe
(experience_kernels.cu)
Was: multi-block kernel with cross-block atomicAdd into a single scalar,
and a Rust wrapper that called cuMemsetD32Async to zero it before each
launch. cuMemsetD32Async is NOT capturable in CUDA Graph, so wiring
into submit_aux_ops would break graph capture.
Now: per-sample loss buffer [B], one thread per sample, no atomic, no
memset. Caller reduces with the existing c51_loss_reduce_kernel
(single-thread sequential sum) for the scalar. Same pattern as G12
predictive_coding_loss + reduce. Fully graph-safe, fully deterministic.
2) Rust wrappers updated (gpu_dqn_trainer.rs)
- compute_branch_independence: drop the cuMemsetD32Async (G6 was
already graph-safe — single-block kernel writes scalar via
overwrite, not accumulation; the memset was unnecessary)
- compute_temporal_consistency: switch to two-step (per-sample +
reduce) to match the new kernel signature; drop cuMemsetD32Async
- New temporal_per_sample_buf field [B] alongside the existing
temporal_penalty_buf scalar
- Three new readback methods (sync DtoH, epoch-boundary only):
branch_indep_loss_value() — G6
temporal_loss_value() — G10
predictive_loss_value() — G12 (was unread previously)
3) Wired into the training loop
- submit_aux_ops calls compute_branch_independence + compute_temporal_consistency
right after compute_predictive_coding_loss (G12). Forward-only — no
backward gradient flows yet.
- FusedTrainingCtx::read_gem_losses() — pass-through accessor that
calls all three trainer-level readbacks at once.
- HEALTH_DIAG line gained a `gems [g6_branch_indep=X g10_temporal=Y
g12_predictive=Z]` suffix so each epoch's penalty magnitudes are
visible. Sync DtoH happens once per epoch (batch boundary), so the
overhead is negligible (~3 × ~1µs).
What this gives us:
- Empirical evidence of whether G6/G10 gem signals are nonzero
- A clean baseline for deciding whether to wire backward (V7
methodology: measure, then commit)
- G12 predictive loss is now also visible (was wired backward in
earlier commit but the loss scalar itself was never logged)
Smoke test:
- 6 trials passed
- Best Sharpe variance: [15.72, 31.94] (wider than G12-only [17.99,
19.56]) — likely cuBLAS algorithm reselection from the new kernel
launches changing graph timing; not a correctness issue
- Tests pass; HEALTH_DIAG now logs gem values per epoch
Next session can run a 5-trial multi-trial test, look at the HEALTH_DIAG
gems line, and decide:
- If g10_temporal ≈ 0 → G10 redundant with spectral_norm; delete
- If g10_temporal nonzero → wire backward
- Same logic for g6_branch_indep
Files touched:
crates/ml/src/cuda_pipeline/experience_kernels.cu (-30 / +48)
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs (+87 / -31)
crates/ml/src/trainers/dqn/fused_training.rs (+25)
crates/ml/src/trainers/dqn/trainer/training_loop.rs (+10 / +3)
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
e72885e8b6 |
gem(G12): finish predictive_coding — write backward + wire into training loop
Found by V7-gem audit: predictive_coding_loss kernel existed (forward only,
per-sample MSE on consecutive trunk activations), and compute_predictive_coding_loss
Rust wrapper existed, but no backward and no caller. Pure scaffolding.
Self-supervised temporal smoothness on the enriched trunk h_s2 IS in the
"better form" by the V7 methodology — it operates at the gradient level on
the network representation, not as a reward shaping term. Worth wiring up.
Three changes:
1) New CUDA kernel `predictive_coding_backward` (experience_kernels.cu)
Loss: L = sum_{i=0..N-2} (lambda/SH2) * sum_j (h[i,j] - h[i+1,j])^2
Grad: dL/dh[i,j] = (2*lambda/SH2) * sum-of-affected-loss-terms
Each h[i,j] appears in TWO loss terms (interior i): "current" of term i
and "next" of term i-1. Boundaries (i=0, i=N-1) have one. Closed-form
gradient written directly with no atomicAdd — one thread per (sample,
feature) cell, each writes to a unique slot, plain += accumulates into
bw_d_h_s2. Bit-deterministic.
2) Rust glue (gpu_dqn_trainer.rs)
- Load `predictive_coding_backward` kernel via existing exp_module_for_mag
- New field on DQNTrainer (predictive_coding_backward_kernel)
- Extend `compute_predictive_coding_loss` to also launch backward as
step 3 (after forward + reduce). Now the function name accurately
describes what it does — both compute and accumulate gradient.
3) Integration (fused_training.rs::submit_aux_ops)
Inserted the call right after `launch_recursive_confidence_backward`,
before regime_scale_td_errors. Both spots accumulate into bw_d_h_s2 via
plain +=, so ordering is irrelevant for correctness — what matters is
that this runs INSIDE the aux_child CUDA-graph capture window AND
BEFORE the trunk W_s2 → W_s1 backward GEMMs read bw_d_h_s2.
Why not also G6 (branch_independence) and G10 (temporal_consistency)?
Per V7-gem methodology — check for redundancy first:
- G10 wants Lipschitz on Q for similar states. Spectral normalization
(already wired on all 12 weight tensors) achieves *global* Lipschitz.
G10 adds *local-pair* Lipschitz on top. Possibly redundant — needs
a measurement before wiring blindly.
- G6 wants the 4 advantage branches to stay diverse. NoisyNets already
adds different parameter noise per layer → naturally diverse heads.
Ensemble KL gradient pushes ensemble heads apart (different mechanism
but similar intent). Possibly redundant for the 4-branch case.
G12 is unambiguously useful — trunk smoothness is a genuine gem with no
existing equivalent in the codebase. G6/G10 land separately if measurement
shows they add real value beyond spectral_norm + NoisyNets + ensemble KL.
Files touched:
crates/ml/src/cuda_pipeline/experience_kernels.cu (+45)
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs (+33 / -3)
crates/ml/src/trainers/dqn/fused_training.rs (+9)
Verified: cargo check passes. Smoke test running to verify training is
stable with G12 active (lambda_pred=0.1 — small enough that any regression
is from a real bug, not dominance over the C51/IQN gradient).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
863eebb82a |
determinism: CUBLAS_WORKSPACE_CONFIG for tests + remove 2 more atomicAdds
Three coordinated changes for catching real regressions locally on the
RTX 3050 Ti before deploying to L40S:
1) Smoke tests now run in CUDA deterministic mode
(crates/ml/src/trainers/dqn/smoke_tests/helpers.rs)
New init_deterministic_cuda() runs once before the first cuBLAS handle
is created in the test process. Sets:
CUBLAS_WORKSPACE_CONFIG=:4096:8 # standard PyTorch determinism knob
NVIDIA_TF32_OVERRIDE=0 # belt-and-braces against TF32
Wired into cuda_device() via OnceLock so every smoke test gets the
deterministic path. Trade-off: ~1.5-3× slower per run (cuBLAS picks
slower-but-deterministic algorithms instead of heuristic-fastest).
For tests this is a great trade — bit-stable results enable real A/B
regression detection across kernel changes.
Production code is NOT affected — it doesn't call this helper.
2) trade_stats_reduce: deterministic per-warp scratch
(crates/ml/src/cuda_pipeline/trade_stats_kernel.cu)
Single-block kernel was atomicAdd-ing 6 floats from per-warp lane 0s
into __shared__ scalars. Replaced with one __shared__ slot per warp
(max 32 warps) + sequential reduction by tid 0 in fixed warp-id order.
Bit-stable across runs, same arithmetic, slightly more shared mem
(~768 bytes additional, well under 48 KB limit).
Result is consumed by HEALTH_DIAG (Kelly stats accumulated across
episodes) — so determinism here matters for downstream training
decisions, not just logs.
3) causal_q_delta_reduce: plain += (single thread writer to unique slot)
(crates/ml/src/cuda_pipeline/dqn_utility_kernels.cu)
Single-block kernel where tid 0 is the only writer, sensitivity_out
buffer is zeroed before the loop, and feature_k is the loop variable
so each call writes to a unique slot. The atomicAdd was unnecessary
— plain += is sequential within one block + serialized across kernel
launches on the same stream. No race possible.
After these three commits the live atomicAdd inventory is:
experience_kernels.cu :: atom_stats[0..1] (cross-block, multi-pass needed)
experience_kernels.cu :: penalty_out (cross-block; result also looks
unused — possible dead path)
branch_indep_penalty_buf in gpu_dqn_trainer.rs is also written but never
read — same dead-code pattern as the just-deleted ensemble_diversity_kernel.
Verified: SQLX_OFFLINE=true cargo check -p ml --lib --tests passes.
Determinism verification (two consecutive smoke runs, byte diff) is the
next step locally before any L40S deploy.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
379cc446e2 |
determinism: deterministic per-warp reduction in monitoring_reduce kernel
The monitoring kernel computed reward/action statistics for HEALTH_DIAG and
log output. Each of 8 warps used per-warp reduction via shfl, then the
warp's lane-0 atomicAdd-ed its partial into 14 shared scalars/arrays:
atomicAdd(&s_sum, ...) // per-trade reward sum
atomicAdd(&s_sq_sum, ...) // sum of squares
atomicMinFloat(&s_min, ...) // CAS-based atomic min
atomicMaxFloat(&s_max, ...)
atomicAdd(&s_nonzero, ...)
atomicAdd(&s_exp[i], ...) // per-action histogram
atomicAdd(&s_ord[i], ...)
atomicAdd(&s_urg[i], ...)
Atomic-into-shared is order-of-arrival → for floats, a few-ULP variance
across runs in mean/std/sharpe/min/max LOG values, even when the underlying
inputs were bit-identical.
Diagnostic non-determinism is still bad: it makes A/B comparisons of kernel
changes unreliable, and bisecting a numeric regression by training-log diff
becomes impossible.
Fix: standard "warp-id-keyed scratch + sequential reduction by tid 0":
__shared__ float w_sum[32], w_sq_sum[32], ... // one slot per warp
__shared__ int w_exp[32][9], ... // arrays too
if ((tid & 31) == 0) w_sum[warp_id] = local_sum; // one writer per slot
__syncthreads();
if (tid == 0) {
for (w = 0; w < num_warps; w++) total += w_sum[w]; // fixed order
...
}
Results are now bit-identical across runs given identical inputs. Shared
memory cost: ~2.6 KB (32 warps × ~80 bytes), well under the 48 KB limit.
Also deleted the now-unused atomicMinFloat / atomicMaxFloat helper
__device__ functions (top of the file). They were the only callers.
The 4 remaining call-site atomicAdds in the codebase (3 in
experience_kernels for atom_stats/penalty_out, 1 in dqn_utility for
sensitivity_out, 1 in trade_stats) are similar diagnostic-only paths.
Each follows the same cross-block hierarchical-atomicAdd pattern used in
the just-deleted ensemble_diversity_kernel; same fix would apply but
they're individual cleanup follow-ups.
Net behavior change: zero (training trajectory unaffected — diagnostic
output values are now stable across same-input runs).
Files touched:
crates/ml/src/cuda_pipeline/monitoring_kernel.cu (-26 / -33 +62)
Verified: SQLX_OFFLINE=true cargo check -p ml --lib --tests passes.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
25ceb3b8d5 |
cleanup: delete dead ensemble_diversity_kernel + readback path
Found while auditing the remaining atomicAdds in the codebase:
ensemble_diversity_kernel computed pairwise KL divergence between K
ensemble heads and wrote the result to ensemble_diversity_loss_buf.
A readback function (`readback_diversity_loss`) was defined to consume
this scalar — but **nothing in the codebase ever called it**. The
kernel ran every training step, allocated a buffer per step, and the
result was discarded.
Removed (-195 LOC net):
- ensemble_diversity_kernel CUDA kernel (~100 LOC C++)
- kernel load + cubin pin in compile_ensemble_kernels
- ensemble_diversity_kernel field on FusedTrainingCtx
- ensemble_diversity_loss_buf field + alloc
- pending_diversity_loss_ptr + pending_diversity_normalizer fields
- readback_diversity_loss() function (the would-be consumer)
- launch site in run_ensemble_step (zero+launch+pending_ptr setup)
Kept (these ARE used):
- ensemble_aggregate_kernel (Q-value mean/variance for exploration bonus)
- ensemble_kl_gradient_kernel (computes diversity gradient → SAXPY into
grad_buf → adam — this is the actual training-path mechanism)
- apply_ensemble_diversity_backward (calls kl_gradient_kernel)
- ensemble_diversity_weight (scales the gradient)
Net effect on training: zero (the deleted code's output was unused).
Net effect on per-step cost: small but non-zero — saves one kernel
launch + memset per step + a CudaSlice<f32> alloc per training context.
On L40S/H100 this is microseconds; on RTX 3050 Ti slightly more.
Effect on determinism: zero. The atomicAdd in the deleted kernel was
in a code path whose output didn't feed training, so removing it
doesn't change the training trajectory. The remaining 9 atomicAdds
in the codebase break down as: 5 in monitoring_kernel (diagnostic
stats only), 3 in experience_kernels (atom_stats / penalty_out —
diagnostic-ish), 1 in dqn_utility (sensitivity_out feature attribution),
1 in trade_stats. None are in the gradient hot path.
Files touched:
crates/ml/src/cuda_pipeline/ensemble_kernels.cu (-100 lines)
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs (-13 lines)
crates/ml/src/trainers/dqn/fused_training.rs (-100 lines)
Verified: SQLX_OFFLINE=true cargo check -p ml --lib --tests passes.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|