Commit Graph

638 Commits

Author SHA1 Message Date
jgrusewski
76e55047ec spec+plan(moe): add no-HtoD/HtoH constraint; mapped pinned only
Per feedback_no_htod_htoh_only_mapped_pinned.md (newly recorded): every
CPU<->GPU path in this redesign uses mapped pinned memory exclusively.
No cudaMemcpy HtoD, no Vec-to-Vec defensive copies, including in test
code. CPU is strictly read-only on the production surface.

Plan changes:
- New Task 2.0 promotes MappedF32Buffer / MappedI32Buffer from
  distributional_q_tests.rs local definitions to a shared
  crates/ml/src/cuda_pipeline/mapped_pinned.rs module so all kernel
  test wrappers (Test 0.F, upcoming MoE tests) share one
  implementation. Adds write_from_slice helper for direct host_ptr
  write (no memcpy).
- Task 2.1 test wrapper rewritten to allocate mapped pinned buffers
  + write to host_ptr + read GPU-written output via host_ptr. No more
  memcpy_stod / memcpy_dtov in test code.

Spec: new section 6.4 codifies the mapped-pinned-only constraint and
references the shared module + reference implementation.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-27 17:46:41 +02:00
jgrusewski
69141d6266 plan(dqn): MoE regime redesign — bite-sized implementation plan
Implementation plan for the approved MoE spec at
docs/superpowers/specs/2026-04-27-moe-regime-redesign-design.md.

6 phases:
- Phase 0: use_* flag cleanup + count_bonus refactor (precondition,
  partially in stash@{0})
- Phase 1: MoE infrastructure (additive — moe_lambda hyperparameter,
  9 ISV slots, MoeGate/MoeExpert skeletons, params_buf layout extension)
- Phase 2: 4 new CUDA kernels (mixture forward/backward,
  load_balance_loss, expert_util_ema_update) with TDD unit tests
- Phase 3: wire MoE into the training graph (forward, backward, loss
  aggregation, ISV producer launch, HEALTH_DIAG aux_moe line)
- Phase 4: atomic deletion of vestigial RegimeConditionalDQN (per
  feedback_no_partial_refactor.md)
- Phase 5: Layer 3/5 tests + L40S 30-epoch validation with explicit
  kill criteria

Each task has bite-sized steps (2-5 min each) with exact file paths,
copy-pasteable code, exact commands, and TDD pattern (test fails first,
implement to make pass, commit).

Self-review: spec section coverage verified, no placeholders, type
consistency verified across phases.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-27 17:42:07 +02:00
jgrusewski
8629a9e7c2 spec(dqn): MoE replacement for vestigial RegimeConditionalDQN
End-to-end investigation (2026-04-27) confirmed RegimeConditionalDQN is
vestigial decoration — 3 heads constructed at training start but only
trending_head ever receives gradient updates. GpuDqnTrainer (the actual
production GPU trainer) has zero references to RegimeType/regime
routing; experience replay inserts go to trending_head.memory only;
ranging_head and volatile_head stay at random init for the entire
training run. Several support APIs (get_count_bonuses_branched, config,
get_state_dim) hardcode-delegate to trending_head, ignoring the regime
split entirely. Per `feedback_no_hiding.md` (wire up or delete) and the
user's preference to fix not delete: the design wires regime
conditioning properly via Mixture-of-Experts replacing the vestigial
3-head architecture.

Pearl introduced and saved as `pearl_learned_gate_subsumes_handcoded.md`:
when the network already sees the heuristic's inputs, a learned gate
strictly subsumes any hand-coded discretization. This is the
load-bearing rationale — ADX/CUSUM are already at state indices 40/41,
so threshold-based regime classification is a strict information
bottleneck the gate can recover and improve on.

Design summary:
- Architecture: shared GRN trunk -> K=8 small expert MLPs (256->64->256
  bottleneck per expert, ~33k params each) -> learned gating network
  (state[42]->64->8 softmax) -> mixed h_s2 -> existing branching heads
  + C51 + IQN dual head. Soft full mixture (no top-k hardcoding); gate
  emerges peaky or flat from data. Anti-collapse load-balancing aux
  loss with default lambda=0.01 (configurable hyperparameter, not a
  kernel constant) prevents init-noise-dominated single-expert lock-in
  without forcing uniform utilization. User-confirmed signal:
  "collapses don't recover well in this codebase".
- 9 new ISV slots (118-126: per-expert utilization EMA + gate entropy
  EMA), GPU-driven producer per
  `pearl_cold_path_no_exception_to_gpu_drives.md`.
- 3 new small CUDA kernels (moe_mixture_forward/backward,
  moe_load_balance_loss) + 1 ISV producer; everything else is cuBLAS-
  reusable. CUDA Graph capture compatible.
- Atomic deletion (no fallback): regime_conditional.rs (~700 LOC),
  RegimeType enum, classify_from_features, RegimeMetrics,
  RegimeClassConfig, 4 DQNConfig regime threshold fields, per-regime
  3-file checkpoint format. DQNAgentType becomes thin wrapper over
  single DQN. Old checkpoints fail loudly with layout-fingerprint
  mismatch.
- 5-layer testing strategy (unit kernels, smoke, gate-differentiation
  validation, L40S production validation with explicit kill criteria,
  architecture-hash backward-incompat).

Out of scope (explicit): top-k routing, per-expert action heads,
hierarchical MoE, regime-conditional CountBonus/NoisySigma broadcast,
expert warm-start from existing trending checkpoint, CVaR action
selection on mixed C51 distribution.

Precondition: the in-progress use_* flag cleanup + count_bonus
[f32; N] refactor lands as its own commit before MoE implementation
begins, per `feedback_no_partial_refactor.md`.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-27 17:31:28 +02:00
jgrusewski
da632446ce refactor(dqn): strip use_iqn feature flag + dead legacy iqn_network
`use_iqn` is exactly the `use_/enable_` boolean banned by
`feedback_no_feature_flags.md`. It gated dead code: production training
runs IQN unconditionally through `cuda_pipeline/gpu_iqn_head.rs` +
`iqn_dual_head_kernel`, properly wired into the branching architecture
with the `FIXED_TAUS` 5-quantile schedule. Nothing in the cuda_pipeline
production path ever read `use_iqn` or `iqn_network`.

The legacy `iqn_network: Option<QuantileNetwork>` field on `DQN` was a
parallel CPU-side network from a pre-branching era, structurally
unreachable in production: every consumer was gated behind the
`if true /* use_branching: always on */` arm at `q_values_for_batch`,
so the IQN else-if at L1822 was dead code. Training optimised
`iqn_network` parameters in isolation; inference never read them. That's
the train/inference mismatch the L283 comment ("IQN trains base
q_network but inference uses dist_dueling network (zero gradients)")
was working around by **disabling** the feature instead of fixing the
inference path. Per `feedback_no_quickfixes.md` + `feedback_no_hiding.md`
the fix is to remove the dead path entirely.

Strip:
- `DQNConfig::use_iqn` field + parses + checkpoint hash + metadata.
  `dqn.use_iqn` / `dqn.iqn_embedding_dim` / `dqn.iqn_num_quantiles` /
  `dqn.iqn_kappa` from older checkpoints are silently dropped on load
  (same pattern used for `use_dueling`). `iqn_lambda` stays — the
  cuda_pipeline dual head consumes it as the IQN aux loss weight.
- 3 vestigial config fields (`iqn_num_quantiles`, `iqn_kappa`,
  `iqn_embedding_dim`) — never read in production; kernel-side macros
  (`IQN_NUM_QUANTILES = 5`, embed_dim 64, kappa 1.0) are the actual
  config.
- 4 default builders (`Default`, `aggressive`, `conservative`,
  `emergency_safe_defaults`) drop the 4 IQN-related fields each.
- `DQN::iqn_network` field + initialisation block in `new_with_stream`.
- 6 conditional gates in `select_action`, `select_action_with_confidence`,
  `select_action_inference`, `q_values_for_batch` — all collapse to the
  live (branching or standard-Q) arm.
- `DQN::get_state_embedding` (only consumed by deleted IQN paths).
- The entire `crates/ml-dqn/src/quantile_regression.rs` module (392 LOC)
  + its 2 lib.rs exports. Nothing outside `ml-dqn` ever imported it
  (the `quantile_huber_loss` reference in `gpu_iqn_head.rs` is a CUDA
  kernel name string, unrelated to this Rust module).

Downstream call sites:
- `crates/ml/src/trainers/dqn/{config,fused_training,trainer/constructor}.rs`:
  drop `iqn_num_quantiles` / `iqn_embedding_dim` / `iqn_kappa` references
  off `DQNConfig`; substitute kernel-fixed literals (64, 1.0) where
  `GpuDqnTrainConfig` / `GpuIqnConfig` still expect them.
- `crates/ml/examples/evaluate_baseline.rs`: drop two `iqn_num_quantiles`
  hyperparam reads (their `..DQNConfig::default()` fallbacks now stand
  alone).
- `crates/ml/tests/dqn_action_collapse_fix_test.rs`: drop the
  `assert!(!config.use_iqn, "...gradient dead zone")` and the explicit
  `config.use_iqn = false` setter; the dead-zone pathology is now
  structurally impossible.
- `crates/ml/tests/dqn_inference_test.rs`: drop `config.use_iqn = false`.
- `services/trading_service/src/services/dqn_model.rs`: drop the
  `iqn={}` debug-log field.
- `crates/ml/src/trainers/dqn/distributional_q_tests.rs`: ship Test 0.F
  (Plan A Task 8 #186) — converged-checkpoint extraction harness +
  `MappedF32Buffer` (mapped pinned f32 mirror) + `compute_sigma_c51_test`
  kernel handle. The structural assertions panic on the local 5-epoch
  smoke checkpoint as designed (under-converged: sigma_C51 spread ~1.4%
  across directions, P(active)=0.4330 >= 0.20). Docstring rewritten to
  drop `use_iqn=false` framing and the legacy "Tier-B-prime" caveat;
  Tier-A version is GPU-integration-only per
  `feedback_no_cpu_forwards.md` (CPU is read-only).

`docs/dqn-wire-up-audit.md` updated per Invariant 7.

Build: `cargo check --workspace --tests` clean (0 errors).
Test: `cargo test -p ml --lib distributional_q_tests::test_0f
       -- --ignored --nocapture` produces bit-identical sigma_C51 /
       argmax / Thompson values vs pre-removal — confirms branching
       forward path is the same code post-cleanup as pre-cleanup
       (legacy `use_iqn` arms were unreachable, as expected).

Net: 12 files, +458 / -785 lines (327 LOC deleted).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-27 16:12:13 +02:00
jgrusewski
db9936b9ff fix(data): align fxcache data_source + normalise DBN fallback
Two-part fix for a class of bugs causing un-normalised features to
silently flow into training.

(1) data_source mismatch between precompute writer and trainer reader.

  precompute_features.rs:214,633 hardcoded "ohlcv"
  train_baseline_rl.rs:582       hardcoded "ohlcv"
  config/training/dqn-production.toml: data_source = "mbp10"

  Production runs the trainer with the production profile (data_source
  = mbp10), but the actual cache lookup hardcoded "ohlcv". Smoke worked
  by accident (smoke profile is also "ohlcv"). Any future profile with
  a different data_source silently mismatches → cache MISS → DBN
  fallback path.

  Both call sites now hardcode "mbp10" (the canonical production data
  source per CLAUDE.md). precompute_features adds a `--data-source`
  CLI override for the rare case a smoke flow needs to regenerate
  the local "ohlcv" fxcache; default is "mbp10".

(2) DBN-fallback path didn't normalise features.

  precompute_features.rs:629 applies NormStats::normalize_batch on the
  canonical fxcache write path. The fallback in train_baseline_rl.rs
  (cache-miss → load DBN files → extract features → upload to GPU) did
  NOT normalise. Any cache miss (data_source drift, schema-hash mismatch,
  missing file) silently uploaded RAW features. Raw close prices
  (~$5180 ES futures) flowed into next_states[:, 0]; the aux next-bar
  head's label_scale EMA latched onto raw-price magnitude (~5443 vs
  expected ~1.0 z-score); the shared trunk learned to predict next-bar
  prices; the policy effectively traded with future-price knowledge →
  train-h5gxb epoch-0 Sharpe = 141 with 0.32% max-drawdown over 214k
  bars (impossibly good = oracle leak).

  DBN fallback now applies the same z-score normalisation unconditionally
  as defence-in-depth, so a future cache-miss cannot reintroduce raw
  values into training.

Audit entry updated.
2026-04-27 14:21:42 +02:00
jgrusewski
cb69e410ea fix(fxcache): track precompute_features.rs in FEATURE_SCHEMA_HASH
build.rs::emit_feature_schema_hash only hashed
src/features/extraction.rs, src/fxcache.rs, and
../ml-core/src/state_layout.rs. The z-score normalization step lives
in examples/precompute_features.rs:625-631 (added 2026-04-03 in
9f7c14978f) and was NOT covered by the hash. A fxcache written by an
older precompute build (raw features, no normalization) silently
passed today's validate() because every other field matched.

Empirical impact (L40S Argo train-f8h6q, 2026-04-27):

  - PVC fxcache: stale, written pre-normalization → feature column 0
    contains RAW CLOSE PRICES (~$5180 ES futures) instead of z-
    normalized log-returns
  - aux head reads next_states[:, 0] as its next-bar regression
    label (gpu_dqn_trainer.rs:7758-7789)
  - EMA label_scale climbed to 5420 (vs smoke 0.05) → shared trunk
    learned to predict next-bar prices → policy effectively traded
    with future-bar information
  - epoch-0 Sharpe = 141.99 with 0.32% max-drawdown over 214k bars
    — physically impossible; clear future-leak signature

Fix adds examples/precompute_features.rs to schema_sources. New hash
invalidates the stale PVC cache. Argo's ensure-fxcache step has a
regenerate-on-failure branch (infra/k8s/argo/train-template.yaml:
372-383) that auto-regens with current normalized precompute.

Generalises beyond this incident: any future change to feature
normalization, target ordering, or precompute post-processing now
bumps the hash and forces fxcache regen.

Audit entry updated.
2026-04-27 12:53:39 +02:00
jgrusewski
0e8804a770 spec(dqn): reframe Thompson rollout PAUSED (not SUPERSEDED)
Retracts the prior SUPERSEDED footer (commit 42ffd6aad). The technical
proposal still stands — Thompson sampling on C51+IQN distributions is
the canonical action selector for distributional RL (Bellemare 2017,
Dabney 2018). Phase 0 tests and the Aggregation Contract are sound
math/engineering regardless of the measurement-bug findings.

What changed is the URGENCY framing, not the validity. The val-Flat-
collapse / Short-collapse observations cited as motivating evidence
were partly distorted by three measurement bugs (a86fba2b1 + b8788511c)
in the diagnostic infrastructure. The "ship Thompson NOW because
val_dir_dist collapses to 80%+ Hold/Flat" narrative dissolves; the
"Thompson is the principled action selector for our distributional
model" narrative stands.

Sequencing: PAUSED pending evidence from a fresh L40S 30-epoch
baseline (train-f8h6q, 2026-04-27 12:20) on post-fix code. The
baseline is a bug-hunting expedition — kill on anomaly, diagnose,
fix, re-run per feedback_stop_on_anomaly.md. Once healthy baseline
established, Phase 2 ships as principled improvement with clean A/B
against the trustworthy post-fix metrics.

Plans B / C / D are PAUSED, not cancelled. Files remain in
docs/superpowers/plans/. Resumption gate: post-fix baseline run is
bug-free or all surfaced bugs are addressed.
2026-04-27 12:27:22 +02:00
jgrusewski
42ffd6aadc spec(dqn): mark Thompson rollout SUPERSEDED — motivating pathology was a measurement artefact
The val-Flat-collapse / Short-collapse / C51 expected-Q bias hypothesis
that motivated the 4-plan distributional-RL Thompson rollout was
largely a measurement artefact in the diagnostic infrastructure, not
a real policy pathology. Three layered bugs in actions_history_buf
init + reader + mag_stats attribution conspired to inflate val_dir_dist
Short, inflate active_frac, and pin wr_h/wr_f to zero. After fixing
all three (commits a86fba2b1 + b8788511c), val_dir_dist matches
val_picked_dir_dist within ~5pp — no collapse, diverse picks.

Status footer added to the spec documenting:

  - what was actually wrong (3 measurement bugs)
  - what the post-fix data shows (mild passivity bias from early
    training, not a structural collapse)
  - what is preserved (Phase 0 unit tests as latent infrastructure
    for any future distributional Q-head; Aggregation Contract
    pearl as a sound engineering invariant)
  - what is cancelled (Plans B / C / D — Phase 1 audit, Phase 2
    Thompson integration, Phase 3 long verification — superseded)
  - future revival condition (if fresh L40S 30-epoch on post-fix
    code shows val_picked_dir_dist itself collapsing toward Hold/Flat,
    reopen)

Plan files remain in docs/superpowers/plans/ as historical record.
2026-04-27 12:23:34 +02:00
jgrusewski
b8788511ce fix(dqn): mag_stats wr_h/wr_f attribution — bin trade closes by pre_mag
experience_kernels.cu line 1916 binned action_mag_per_sample by
actual_mag_core at every step, including trade-close events. But
unified_env_step_core forces `actual_mag = 0 (Quarter)` whenever
actual_dir is Hold/Flat (trade_physics.cuh:772) and trade closes
always land in Hold/Flat state — so every Half/Full close was
attributed to the Quarter bin. close_counts[Half] and close_counts[Full]
were structurally pinned to 0, giving wr_h = wr_f = 0 across all
training runs.

Fix: introduce `seg_mag_bin = is_close ? pre_mag_bin : actual_mag_core`.
At close events bin by pre_mag_bin (the magnitude of the position
being closed); at non-close events keep actual_mag_core (current
realized magnitude). pre_mag_bin is always 0/1/2 at close events
since exiting/reversing requires prev_sign != 0 → pre_trade_position
!= 0 → pre_frac > 0.001 → pre_mag_bin in {0,1,2}.

Smoke verification (5-epoch local): wr_h/wr_f remain 0 because
var_scale (1/(1+sqrt(var_q))) shrinks effective_max_pos to 10-19% of
broker max at smoke maturity → even Long Full target lands at
abs_pos ≈ 0.15 < 0.375 → all positions decode as Quarter; no
Half/Full positions exist for the fix to attribute. This is the
expected structural consequence of the var_scale design (uncertain Q
→ smaller position, conservative). The fix is latent correctness:
in mature L40S 30+ epoch runs where var_q drops and var_scale
grows above 0.375, Half/Full positions become reachable and
wr_h/wr_f will reflect their real realized win rates.

Without the fix, even mature training would show wr_h = wr_f = 0
because of the Hold/Flat → Quarter close convention masking real
per-magnitude win rates. Audit entry updated.
2026-04-27 12:05:49 +02:00
jgrusewski
a86fba2b1d fix(dqn): val_dir_dist + active_frac measurement artifact (-1 sentinel)
After done_flags[w]=1 (capital floor breach), backtest_env_step
early-returns without writing actions_history_buf for remaining slots
in [done_step, max_len). The prior zero-init decoded those slots as
Short Quarter Market Normal (action 0) via `dir = 0/27 = 0` and
inflated val_dir_dist's Short bucket / active_frac to a measurement
artifact masking real model behaviour.

Two-part fix:

1. `gpu_backtest_evaluator.rs::reset_evaluation_state`: replace
   `memset_zeros` for actions_history_buf with
   `cuMemsetD32Async(0xFFFFFFFFu32)` writing -1 sentinel. The Rust
   readers already filter `if a < 0 { continue; }` so unwritten slots
   are skipped correctly post-fix.

2. `backtest_metrics_kernel.cu`: add `if (act < 0) continue;` after
   reading actions_history. The reduce-side metrics
   (buy_count/sell_count/hold_count → active_frac/dir_entropy +
   bnd_* trade-boundary detection) now consistently skip unwritten
   slots. step_returns at those slots are still zero-init (correct)
   so summing them with r=0 is a no-op.

Empirical impact (local 3-fold × 5-epoch smoke, RTX 3050 Ti):
  val_dir_dist Short: 81-84% → 13-29% (matches val_picked within 5pp)
  active_frac:       87-91% → 31-50%
  dir_entropy:       0.57    → 0.83-1.02

The pre-fix "val-Flat-collapse" / "Short-collapse" pathology that
motivated substantial subsequent investigation (incl. the 4-plan
distributional-RL Thompson rollout draft) was largely a measurement
artifact from this bug surfacing differently before vs after the
Kelly cap fix (`0c9d1ee39`). Pre-Kelly the Kelly cap clamped most
Long/Short → Flat → actions_history was densely written with Flat
encoding → 80% Flat reading (real Kelly pathology + small artifact).
Post-Kelly the picks survive but the poor smoke-trained model
breaches capital floor often → many unwritten Short slots → 83%
Short reading (pure artifact). With both fixes, val_dir_dist now
reflects real model behaviour.

Audit entry updated in docs/dqn-wire-up-audit.md.
2026-04-27 10:36:09 +02:00
jgrusewski
a3bb040bc2 test(dqn): Phase 0 Test 0.E — synthetic edge discovery via Thompson Q-learning
Algorithmic property test (CPU). Confirms Thompson exploration discovers
KNOWN +0.005 edge in 100 iterations on a 1-state bandit, while
argmax-only training never updates Q[Long].

Setup revised from plan A draft (option 3 — production-realistic):
  p_long initial = [0.10, 0.20, 0.40, 0.20, 0.10] (uniform, E=0, has σ)
  p_flat initial = [0, 0, 1, 0, 0]                (δ(v=0), deterministic)
  Argmax with strict-> ties at E=0 → always picks Flat → never explores
  Long → Q[Long] stays at 0, never discovers edge.
  Thompson samples Long > 0 with P≈0.30 → ~30 effective updates → mean
  drifts toward +0.005, crosses Q[Flat]=0 within budget.

Plan's prior draft (initial p_long with mean=-0.015 + p_flat=δ(0)) was
calibration-bound: Thompson drift was directionally correct (-0.015 →
-0.005) but didn't cross zero in 100 iters. Revised setup eliminates
the artificial initial bias and matches production reality more
closely (Flat = δ(0) by construction; Long starts spread from random
init, then accumulates true edge).

Stop condition: if Thompson e_long ≤ e_flat with this setup, the
hypothesis is genuinely wrong and reward shaping must change before
proceeding to Phase 2.

Observed (local RTX 3050 Ti, ~0.00s test wall, ~1.57s 5-test suite):
  argmax  : e_long=0.000000, e_flat=0.000000 (asserts e_long ≤ 0.001 OK)
  thompson: e_long=0.003550, e_flat=0.000000 (asserts e_long > e_flat OK)

All 5 Phase 0 tests pass: 0.A bias-reproduces, 0.B inverse-CDF,
0.C IQN symmetry, 0.D Thompson-reverses, 0.E synthetic-edge.
2026-04-27 09:20:18 +02:00
jgrusewski
82d39a895c test(dqn): Phase 0 Test 0.D — Thompson reverses bias on failure mode 2026-04-27 09:07:54 +02:00
jgrusewski
3d0ee00eb9 test(dqn): Phase 0 Test 0.C — IQN quantile interpolation symmetry check 2026-04-27 09:03:13 +02:00
jgrusewski
3f65977bb1 fix(dqn): Plan A redesign — code-review minor fixups
- Use u32::div_ceil for grid-dim arithmetic (style)
- Tighten u8::try_from to validate direction < B0_SIZE=4 instead of
  fits-in-u8: kernel-OOB write fails loudly at the first malformed
  index, not silently in production.

No semantic changes.
2026-04-27 08:59:41 +02:00
jgrusewski
454c26e7e8 test(dqn): Plan A — batched + pinned memory infrastructure (no dtoh)
Per feedback_gpu_cpu_roundtrip.md, the per-seed dtoh in
launch_thompson_direction (10k iterations for Test 0.A, 100k for Test
0.B) violated the no-dtoh-on-hot-or-tight-loop invariant. Replaces
the per-seed kernel with a batched kernel (one thread per seed) and
the dtoh wrapper with a MappedI32Buffer using cuMemHostAlloc(
DEVICEMAP|PORTABLE) — same pattern as gpu_training_guard.rs
MappedBuffer.

Kernel changes (thompson_test_kernel.cu):
- thompson_direction_test_batched: replaces thompson_direction_test;
  one thread per seed, writes via mapped device pointer with
  __threadfence_system() for PCIe coherence.
- argmax_eq_test, compute_sigma_c51/iqn_test: __threadfence_system()
  added before kernel exit.

Test refactor (distributional_q_tests.rs):
- MappedI32Buffer helper (test-utility version of the production
  f32-only MappedBuffer).
- Single batched launch per test instead of N launches.
- Tests 0.A and 0.B preserved assertions; runtime drops from ~1.86s
  (Test 0.A) and ~3.8s (Test 0.B) to ~0.15s each on RTX 3050 Ti.

Dead code: launch_thompson_direction (single-seed) and the OnceLock
KernelSet single-launch wrappers deleted; orphan code per
feedback_wire_everything_up.md.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-27 08:55:10 +02:00
jgrusewski
581659e12c test(dqn): Phase 0 Test 0.B — C51 inverse-CDF distribution check + Test 0.A lint cleanup
Originally specified by Plan A Task 4 with P(d=0)≈0.2. That assumed
d=1's sample of 0 would break the tie when d=0's C51 sample is also 0.
The actual `thompson_direction_test` kernel uses strict-> with
`best_d=0` initialised, so d=0 wins whenever its sample is ≥ 0
(atoms v=0 or v=+1).

Corrected expected: P(d=0) = p[v=0] + p[v=+1] = 0.9. Test still
verifies inverse-CDF correctness — wrong math would deviate this
proportion measurably.

Also applies clippy::erasing_op / clippy::identity_op cleanup (`0 *`,
`1 *` index expressions) to Test 0.A from Task 3 review. No behaviour
change in 0.A.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-27 08:39:51 +02:00
jgrusewski
9fc9f0e3c8 test(dqn): Phase 0 Test 0.A — bias reproduces, Thompson reverses
Constructs synthetic C51+IQN distributions matching observed val-collapse
failure mode. Verifies argmax(E[Q]) deterministically picks Flat/Hold
while Thompson sampling produces ≥30% Long+Short over 10000 seeds.
2026-04-27 08:29:33 +02:00
jgrusewski
4468812478 fix(dqn): Plan A Task 2 — code-review fixes
Code-quality review on c2210e8b9 surfaced 2 medium + 4 minor issues
(per feedback_no_quickfixes.md, all fixed):

- M1: module-level #![allow(unsafe_code)] (matches signal_adapter.rs)
- M2: cubin loaded once via OnceLock (was reloaded per launch);
      pattern matches gpu_her.rs/signal_adapter.rs
- M3: i32→u8 cast replaced with u8::try_from + descriptive expect;
      lets a buggy kernel writing OOB fail loudly
- m4: host[0] direct indexing → host.first().copied().expect(...)
- m5: aspirational "Tests follow" comment → declarative
- m6: #[allow(dead_code)] at module top with descriptive comment
      (helpers populated by Tasks 3-8)

Audit doc gets a code-review amendment block on the existing Task 2
entry to satisfy Invariant 7 (component change → audit-doc update).

No spec changes. Plan A Task 2 spec compliance preserved.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-27 08:25:40 +02:00
jgrusewski
c2210e8b95 test(dqn): Phase 0 test module skeleton + kernel wrappers
Adds distributional_q_tests.rs with launch_thompson_direction and
launch_argmax_eq wrappers around thompson_test_kernel.cu. Tests follow
in subsequent tasks.
2026-04-27 08:20:45 +02:00
jgrusewski
2186f96e83 test(dqn): standalone Thompson direction test kernel (Phase 0)
Implements inverse-CDF over C51 atoms, uniform-τ interpolation over IQN
quantiles, and argmax of E[Q] for eval mode. Plus diagnostic σ kernels
(compute_sigma_c51_test, compute_sigma_iqn_test) used by Test 0.F.

Exercised only by Phase 0 tests in distributional_q_tests.rs; Phase 2
production kernel will integrate the same __device__ __forceinline__
math into experience_action_select. No production callers in this
commit.

Plan A Task 1 of docs/superpowers/plans/2026-04-27-distributional-rl-thompson-plan-A-phase-0.md.
2026-04-27 08:16:01 +02:00
jgrusewski
3c61ee52ce plan(dqn): Thompson sampling rollout — Plans A+B+C+D (4 sequential phases)
Implementation plans for the distributional-RL aggregation spec
(docs/superpowers/specs/2026-04-26-distributional-rl-aggregation-design.md).

Plan A — Phase 0: TDD hypothesis verification (8 tasks)
  - Standalone GPU test kernel: sample_c51_inverse_cdf,
    sample_iqn_quantile_interp, compute_e_c51, compute_e_iqn,
    thompson_direction_test, argmax_eq_test
  - 6 unit tests (5 GPU + 1 CPU synthetic edge)
  - GPU integration on converged checkpoint (Test 0.F)

Plan B — Phase 1: existing-lever audit (8 tasks)
  - 6 unit tests covering B.2/CF/PopArt/Q-target audit fixture
  - Exit gate: all 6 PASS = no reward-shaping bug; Phase 2 unblocked

Plan C — Phase 2: Thompson sampling integration (11 tasks)
  - Direction branch of experience_action_select rewritten:
    eps-greedy + Boltzmann → Thompson (training) + argmax E[Q] (eval)
  - C51 + IQN buffers wired to gpu_dqn_trainer + gpu_backtest_evaluator
  - train_active_frac HEALTH_DIAG instrumentation
  - 4 GPU-direct tests against production kernel
  - Aggregation Contract table + memory pearl

Plan D — Phase 3: long verification + dead-code cleanup (8 tasks)
  - Test 3.A: regression anchor against original C51 Flat bias
  - L4: 1 seed × 6 folds × 30 epoch (~1 hour)
  - L5: 5 seed × 6 fold matrix per Plan 5 Task 5 (Tier 1+2+3 PASS)
  - Direction-only eps_dir adaptive boost + Boltzmann tau-floor cleanup
    (per feedback_no_partial_refactor.md)
  - nsys regression check; final architecture/spec footer

Plans gated sequentially: each phase exit gate must pass before next.
2026-04-27 08:10:42 +02:00
jgrusewski
021bb0ef73 spec(dqn): pivot Phase 0+2 tests to GPU-direct (no CPU mirror)
User correctly identified that CPU mirror function tests don't test
the production GPU code path. A bug shared between mirror and kernel
(translated identically wrong) would slip through. Mirror tests + a
single GPU bridge test were a weak compromise.

GPU-direct testing strategy:
  - All Phase 0 kernel-correctness tests (0.A, 0.B, 0.C, 0.D, 0.F):
    launch tiny test-only kernels with the SAME math the Phase 2
    production kernel will use; assert properties of the output.
  - Test 0.E (synthetic edge discovery): stays CPU. It tests an
    ALGORITHMIC PROPERTY of Thompson exploration (does it discover
    edge if edge exists?), not a kernel correctness property.
  - All Phase 2 unit tests (2.A-2.D): GPU-direct against the
    modified production kernel.
  - Phase 0.F (real checkpoint extraction): unchanged — already GPU.

Local development uses RTX 3050 GPU (per memory user_dev_environment.md).
CI runs --ignored flag to skip GPU tests on CPU-only runners.

Time budget: Phase 0 was 1-2 days (CPU mirror); now 2-3 days
(GPU-direct, includes kernel wrapper setup half-day).

Other delta:
  - Phase 0 deliverable file renamed: distributional_q.rs ->
    distributional_q_tests.rs (no mirror functions, just tests +
    kernel wrappers).
  - Phase 2 unit tests rephrased to launch production kernel rather
    than compare against CPU mirror.
  - L1 verification gate runtime: seconds -> minutes (GPU launch
    overhead per test).

The user's intuition was right: testing production directly is the
honest approach. Mirror was an optimization that traded correctness
for speed; with local GPU available the optimization isn't needed.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-27 00:27:46 +02:00
jgrusewski
25ba051157 spec(dqn): critical review pass — fix all majors, mediums, minors
Self-review identified 8 major + 8 medium + 15 minor issues. All fixed:

MAJORS:
  M1 Test 0.A: clarified — compare argmax(E[Q]), Boltzmann(E[Q]),
     and Thompson sampling distributions; assertion is on relative
     ordering across all three.
  M2 IQN quantile count: replaced hardcoded `5` with N_IQN_QUANTILES
     constant (defined per task #147 fixed-quantile design).
  M3 "Converged checkpoint" definition: ≥60 epochs trained AND
     val_sharpe stabilised (no >10% change over last 10 epochs).
     Cites prior 60-epoch validation runs (task #80, train-7rgqd).
  M4 R3 reframed: replaced "no issue" handwave with explicit
     by-design tradeoff acknowledgment + cost analysis. Wasted
     exploration is the cost of finding out whether edge exists.
  M5 Test 0.D σ_long=0.05 justified: chosen to match expected order
     of magnitude given typical |return| ~ 50bps; Phase 0.F
     validates against real checkpoint.
  M6 rng_ctr post-increment: clarified — matches existing pattern
     at experience_kernels.cu:858 (no behaviour change).
  M7 train_active_frac instrumentation: NEW Phase 2 deliverable —
     existing HEALTH_DIAG only has val_active_frac, but L3 verifies
     training-time active_frac. Spec now explicitly adds this
     ~10-line metrics.rs change as a Phase 2 deliverable.
  M8 eps_dir cleanup code-level detail: explicit reference to
     experience_kernels.cu lines 814-865; remove eps_dir from both
     static EPS_FLOOR clamp AND adaptive boost block; verify
     variable can be removed from kernel signature via grep.

MEDIUMS:
  Med1 Current C51/IQN combination: clarified that compute_expected_q
       blends per training schedule; Phase 2 replaces with explicit
       0.5*E_C51 + 0.5*E_IQN equal weighting; Phase 0.F verifies.
  Med2 Eval mode phrasing: "eval mode already sets eps=0 in existing
       kernel" — no semantic override, factually correct.
  Med3 Magnitude σ claim: clarified — magnitude branch likely has σ
       bias in OPPOSITE direction (Full has larger σ; UCB would
       prefer Full and worsen saturation). Empirical verification
       deferred. Phase 0.F should also report per-magnitude σ.
  Med4 Hierarchical sampling claim corrected: it's not about
       balancing 50/50 (already 50/50). It's about decoupling
       cluster-best decisions; clarified.
  Med5 n_atoms vs N_IQN_QUANTILES: clarified — n_atoms variable per
       config (currently 51); N_IQN_QUANTILES fixed at 5.
  Med6 Conviction code: removed pseudo-code; references existing
       implementation at experience_kernels.cu:1091; provides
       implementation hint for E[Q] reuse.
  Med7 Q-target propagation: clarified — uses full distribution
       (C51 atom projection / IQN quantile regression), not just
       E[Q]. Thompson modifies action selection only.
  Med8 References: added Thompson 1933 (original), Bellemare 2017
       (C51), Dabney 2018 (IQN) for theoretical foundations.

MINORS:
  Min1 Date updated to 2026-04-27.
  Min2-3 Argmax monotonic /2 simplified out — argmax(a+b) =
         argmax((a+b)/2). Code clarity improved.
  Min4 P(argmax picks Long) = 0 deterministic; reframed assertion.
  Min5 Test 0.F structural assertions added: σ_C51[FLAT] < 0.01 ×
       σ_C51[LONG]; same for IQN; E[Q_FLAT] > E[Q_LONG]; argmax
       picks FLAT; Thompson P(LONG)+P(SHORT) ≥ 0.20.
  Min6 -INFINITY → CUDART_INF_F (CUDA convention).
  Min7 dir_idx scope: comment notes it's declared earlier in kernel.
  Min8 action_select args: explicit — three buffers exist on GPU
       but not currently passed; new params, no new buffers.
  Min9 Phase 0 time math: 5 hours tests + 1 hour enumeration + 2
       hours 0.F + (3 hours runtime if checkpoint training needed,
       runs in parallel). Honest budget.
  Min10 "Two-stream" → "5-Layer Gate" header.
  Min11 Plan 5 reference uses full path consistently.
  Min12 Plan B time budget: explicit 1 day if pass; 2-5 days if bug.
  Min13 active_frac: clarified Long+Short combined, not per direction.
  Min14 train_active_frac: now in Phase 2 deliverables (see M7).
  Min15 "20 mechanisms" → 21, with sub-counts in section headers.

Spec now 615 lines, comprehensive coverage of:
  - Pearl + theoretical foundation
  - Problem statement (with bias-might-be-correct caveat)
  - Architecture (Thompson at training, argmax at eval)
  - Train vs eval distinct selectors with behavior-change disclosure
  - Direction-only scope with magnitude σ-bias warning
  - Conviction stays E[Q]-based (no Kelly cap jitter)
  - Interaction matrix: 21 mechanisms in 3 categories
  - 6 v2 enhancements documented + deferred
  - Phase 0/1/2/3 with tests, exit gates, time budgets
  - 5-layer verification + train_active_frac instrumentation
  - 8 risks with mitigations + 5 stop conditions
  - What v1 doesn't touch (referencing interaction matrix)
  - References (Thompson 1933, Bellemare 2017, Dabney 2018, etc.)
  - Aggregation contract (project-wide pearl, enforced)

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-27 00:08:02 +02:00
jgrusewski
308c54484e spec(dqn): full interaction matrix + outside-the-box Thompson improvements
Per-user direction: every existing mechanism in the DQN system MUST be
explicitly considered for interaction with Thompson, and Thompson itself
MUST be examined for system-specific improvements beyond vanilla.

INTERACTION MATRIX (3 categories, 20 mechanisms):

Category 1 — Compose with Thompson (no change required):
  Counterfactual reward, B.2 novelty bonus, PopArt, Saboteur,
  Curiosity, NoisyNets/VSN, Distillation, CQL, Polyak target EMA,
  HER, PER, Replay warm-start, Multi-fold validation harness.

Category 2 — Trivially adapt to Thompson (one-line changes):
  D7/N7 contrarian sign flip (negate the SAMPLE), cosine epsilon
  schedule (still applies to mag/ord/urg), per-sample epsilon (IQL
  expectile gap), adaptive Boltzmann tau (still applies to mag/ord/urg).

Category 3 — Take precedence over Thompson (hard constraints):
  Plan-based action lock (Thompson sample discarded if plan active),
  per-magnitude Kelly cap, trail stop, capital floor breach.

Critical insights from the audit:

1. NoisyNets is ALREADY a form of training-time Thompson at the
   parameter level. Output-space Thompson stacks on top —
   total exploration = parameter-space ⊗ output-space (multiplicative).

2. Curiosity is ORTHOGONAL to Thompson — Thompson explores actions
   whose Q is uncertain; curiosity explores states whose dynamics
   are uncertain. Both axes desirable; no conflict.

3. Plan lock takes precedence; same as currently with Boltzmann.

OUTSIDE-THE-BOX v2 ENHANCEMENTS (deferred to follow-up specs):

  v2.1 Triple-source Thompson (C51 + IQN + Ensemble) — incorporate
       the existing ensemble Q-head as 3rd uncertainty source.
  v2.2 Persistent Thompson (anti-churn for HFT) — bias sampling
       toward current direction, ISV-driven; reduces tx_cost from
       Long/Short oscillation across bars.
  v2.3 CVaR-aware eval (risk-adjusted deployment) — eval picks
       argmax(E[Q] − λ·CVaR_α[Q]); risk-aware decision making for
       production with real capital.
  v2.4 Information-Directed Sampling (Russo & Van Roy 2014) — picks
       action minimizing regret²/info_gain; more efficient than
       vanilla Thompson when learning saturates.
  v2.5 Hierarchical Thompson on (trade vs no-trade) → (which
       direction) — addresses 50/50 structural advantage of no-trade.
  v2.6 Composition with curiosity-driven exploration — explicit
       coupling beyond reward-side composition.

Each v2 enhancement gets its own spec/plan when prioritised. Vanilla
Thompson is v1; ships first; verified independently.

ALSO FIXED (from earlier self-review):

  - Pearl claim softened: only the C51 Hold/Flat bias is directly
    attributed; other historic bugs had different mechanisms.
  - TFT entry removed from contract table — TFT is a Variable
    Selection Network (feature processor), not a Q-head. Replaced
    with generic "Future Q-head additions" placeholder.
  - Eval direction = argmax E[Q] explicitly flagged as a behavior
    change from current Boltzmann-with-tau (val_dir_dist will be
    more concentrated than current).
  - Phase 0.E budget reduced (1-2 hours, not 1 day) — synthetic
    bandit is ~50 lines of Rust, not full RL training loop.
  - Phase 0 enumerates existing checkpoints before training new one.
  - Architecture diagram parenthesis fixed.
  - Conviction implementation note: compute E[Q] once, reuse for
    conviction AND eval-mode argmax — no redundant computation.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-26 23:51:17 +02:00
jgrusewski
1c63c32297 spec(dqn): reframe Phase 3 as direction-branch dead-code cleanup
User correctly challenged the "band-aid removal" framing. All fixes
shipped during the val-Flat-collapse investigation addressed real bugs
at their respective layers and should be preserved:

  - Kelly cap warm-branch (0c9d1ee39): post-decision physics layer.
    Thompson-independent. KEEP.
  - Train Return display + Sharpe annualization (non-tau parts of
    7a3d88646): display/metric layer. Thompson-independent. KEEP.
  - Direction Boltzmann tau-floor (tau part of 7a3d88646) +
    adaptive eps_dir floor (d54b49efc): gates inside direction-branch
    action selection. Phase 2 replaces direction-branch action
    selection wholesale (eps-greedy + Boltzmann → Thompson), so these
    direction-only code paths become structurally unreachable.

The latter two are NOT band-aids being removed because Thompson is
better. They are dead code being cleaned up because Thompson replaces
the surrounding mechanism. Magnitude/order/urgency branches keep their
existing eps-greedy + Boltzmann + tau-floor + EPS_FLOOR paths intact.

Reframed Phase 3 deliverable: "direction-branch dead-code cleanup"
with explicit rationale (per feedback_no_legacy_aliases.md and
feedback_no_partial_refactor.md). 0.5 day budget instead of 1.

Also clarified eval action selection: argmax of (E[Q_C51]+E[Q_IQN])/2
is correct. Bellman backup is a Q-learning UPDATE rule, not an
action-selection rule. Once Q is learned, optimal policy is greedy
argmax of learned Q. Online Bellman lookahead at eval would require
a forward model of market dynamics — not available, not standard.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-26 23:41:05 +02:00
jgrusewski
5c325f8947 spec(dqn): pivot to Thompson sampling — distributional RL action selection
Revises the C51-bias spec after deeper review surfaced 12 design gaps,
of which 4 were critical:

  1. Train-only vs train+eval ambiguity — UCB at eval would conflate
     "model recommends Long" with "model is uncertain about Long",
     inflating reported edge. CRITICAL for trading where eval drives
     real capital decisions.
  2. Thompson sampling is more principled than UCB:
       - parameter-free (no κ to tune)
       - uses distribution directly without scalar reduction
       - naturally explore-exploit balanced via distribution shape
  3. c51_alpha is the wrong blend weight (it's C51-vs-MSE-warmup, not
     C51-vs-IQN). Equal-weight average of C51 and IQN samples is the
     structural choice — no tuned blend weight needed.
  4. The bias might be CORRECT BEHAVIOUR — model rationally choosing
     Flat when no edge has been discovered. Phase 0 must include a
     synthetic-edge test (controlled MDP with KNOWN positive Long
     expected value) to verify Thompson can discover edge if it exists.

Other gaps fixed:
  - Eval at argmax E[Q] (not Boltzmann, not Thompson)
  - Pearl wording broadened to cover ensembles + future methods
  - Ensemble Q-head added to aggregation contract table
  - Explicit caveat: NEVER extend Thompson to magnitude branch (would
    worsen existing magnitude saturation)
  - Phase 0.F uses CONVERGED checkpoint (≥30 epochs), not 2-epoch run
  - L4 long smoke (30 epochs, ~1 hour) added — Thompson edge discovery
    needs longer feedback loop than 5 epochs
  - Phase 3 explicitly removes eps-floor + tau-floor band-aids
    (Thompson replaces direction Boltzmann; band-aids become dead code)
  - Conviction stays E[Q]-based, not sample-based (avoid Kelly cap
    jitter from stochastic samples)

Architecture (Thompson only, no UCB):
  TRAINING: dir_idx = argmax(0.5 × (sample_C51(d) + sample_IQN(d)))
            magnitude/order/urgency: existing Boltzmann + ε-greedy
  EVAL:     dir_idx = argmax(0.5 × (E[Q_C51] + E[Q_IQN]))
            magnitude/order/urgency: existing Boltzmann (eval mode)

Direction-branch ε-greedy + Boltzmann are REMOVED — Thompson is the
exploration mechanism. No new GPU buffers; existing C51 atoms + IQN
quantiles passed to action_select.

5-7 days active work across 4 sub-plans; each gets its own
writing-plans cycle.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-26 23:35:40 +02:00
jgrusewski
494125c104 spec(dqn): distributional RL aggregation — UCB on C51+IQN σ for action selection
Designs the structural fix for the C51 expected-Q Hold/Flat bias exposed
by the Kelly val-Flat-collapse fix. The bias is a manifestation of a
deeper project-wide pearl:

  "Distributional RL aggregation discards uncertainty;
   action selection must restore it."

Any value head representing Q as a distribution (atoms, quantiles,
ensembles) MUST expose both E[Q] and σ(Q) to action selection. Boltzmann
on E[Q] alone produces structural bias toward low-variance actions
regardless of expected payoff — the C51+IQN Flat-attractor is one
instance of this lost-information pattern.

Fix: extract σ(Q) from BOTH C51 atoms (closed form) and IQN quantiles
(IQR/1.349), blend by loss-time weight, feed Q_eff = E[Q] + κ·σ
(κ=1.0 structural identity) to direction-branch Boltzmann ONLY.

4-phase implementation:
  Phase 0 — TDD hypothesis verification (Rust mirror functions + 5 unit
            tests including GPU integration on real checkpoint)
  Phase 1 — Audit existing reward levers (B.2, CF, PopArt, Q-target)
            via 6 unit tests; fix any bugs found
  Phase 2 — UCB integration: new compute_q_with_uncertainty kernel,
            modified action_select, Rust orchestration, project-wide
            aggregation contract in dqn-wire-up-audit.md
  Phase 3 — Verification per Plan 5 Task 5 multi-seed × multi-fold

5-layer verification gate; 8 risks with mitigations; 4 stop conditions
that halt execution and force redesign.

Existing band-aid fixes (Kelly cap, eps-floor, tau-floor) stay — they
address symptoms at different layers. UCB adds the missing aggregation
step that was the common root across all the symptoms.

Direction-branch only — magnitude/order/urgency don't have the
Flat-attractor (atom-mass collapse asymmetry).

5-7 days active work across 4 sub-plans; each gets its own
writing-plans cycle.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-26 23:24:25 +02:00
jgrusewski
d54b49efc1 fix(dqn): C51 bias breakout — adaptive eps_dir floor from trade-rate undershoot
Prior C51 bias fix (commit 7a3d88646: ISV-adaptive Boltzmann tau floor)
had no measurable effect on the post-training Hold/Flat collapse —
empirically confirmed in train-bscl2 epoch 2: val_dir_dist
[short=0.135 hold=0.358 long=0.142 flat=0.364], identical to the
pre-fix run [short=0.136 hold=0.348 long=0.148 flat=0.367].

Root cause: once Q-values reflect tx_cost-driven aversion, Boltzmann
correctly samples the biased Q distribution regardless of tau. Tau
adjustments protect cold-start exploration but can't combat learned
preferences. The 2% static eps_dir floor allows only 0.5% random
sampling per direction — too little to break the Q-value lock-in or
generate enough Long/Short experiences for edge discovery.

Fix:
  if (ISV available) {
    passive_pressure = clamp(0, 1, 1 − ISV[71]/max(ISV[72], 1e-4))
    eps_dir = max(eps_dir, 0.5 × passive_pressure)
  }

ISV[71] = TRADE_ATTEMPT_RATE_EMA (current Flat→Positioned rate, B.2 producer)
ISV[72] = TRADE_TARGET_RATE (target frozen at epoch 5 from measured EMA)

Feedback semantics:
- attempt_rate >= target → passive_pressure=0 → eps_dir at baseline 0.02
- attempt_rate = 0 (fully passive) → passive_pressure=1 → eps_dir floor=0.5
- intermediate → linear blend

The 0.5 ceiling is a structural blend point (half random / half policy)
— maximum exploration that still preserves directional Q-signal
propagation through the replay buffer. Not a tuned magnitude.

Cold-start safety: ISV[72] is 0 until epoch 5 freeze; with the 1e-4
target floor, passive_pressure clamps to ≈1 immediately, but eps_dir
also has a baseline 0.02 EPS_FLOOR so the formula's max() picks
whichever is larger. Once ISV[72] freezes, the feedback loop activates
properly.

Eval mode unaffected (eps logic gated behind !eval_mode).

Direction branch only — magnitude/order/urgency don't have the
Flat-attractor problem.

ISV-driven, no new ISV slots, no tuned constants per
feedback_isv_for_adaptive_bounds.md and feedback_adaptive_not_tuned.md.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-26 22:51:34 +02:00
jgrusewski
7a3d886462 fix(dqn): three concerns from kelly-fix-multifold val run
CONCERN 1 — train Return display (financials.rs)
Replace `mean_per_bar × bars_per_year` ("annualized arithmetic return")
with the actual cumulative compounded return: `prod(1+r) − 1` via
log-space sum for f64 stability. Step floor at 1+r >= 1e-10 so a -100%
bar gives -23 log contribution rather than -inf; result bounded above
-100%.

The old formula multiplied per-bar mean by ~98K, producing values like
Return=-51234% that LOOKED like portfolio collapse but were actually
just label artefact: -52 bps mean × 98,280 bars/year = -51,234%, while
the actual compounded return over the rollout was bounded and finite.
"Total Return" now means what a trader expects.

CONCERN 2 — val_Sharpe annualization source-of-truth (training_loop.rs +
gpu_backtest_evaluator.rs)
val_Sharpe was using `m.sharpe` from the kernel (annualised by
`sqrt(bars_per_day × trading_days_per_year)` at evaluator init), but
val_Sharpe_raw recomputed the divisor from `self.hyperparams.bars_per_day
* 252.0` in Rust. Empirically the two diverged by 5.4× (val_Sharpe=143.38
/ val_Sharpe_raw=0.0841 = 1704, expected sqrt(98280) = 313.5). Root cause
unclear without runtime instrumentation but the structural fix is to
read the EXACT factor the kernel applied: new
`GpuBacktestEvaluator::annualization_factor()` returns the f32 stored at
init time. Rust now divides by that, so the two numbers can never drift
regardless of how `bars_per_day` flows through the config layers.
Per `feedback_no_partial_refactor.md` — every consumer of a shared
contract must use the same source.

CONCERN 3 — C51 expected-Q Hold/Flat bias (experience_kernels.cu)
Direction-branch Boltzmann tau floor: replace static `0.01` with
ISV-adaptive `max(ISV[21], 0.01)` where ISV[21] is q_dir_abs_ref (EMA
of mean |Q| across direction bins, same signal that drives conviction).

The C51 expected-Q biases Flat above directional actions when the policy
has no measurable edge — Flat returns are exactly 0 (no position change),
collapsing C51's distribution to a delta at 0; directional returns
concentrate slightly below 0 from tx_cost without compensating edge,
giving E[Q_directional] < E[Q_flat] = 0. With Q-spread on the order
of 0.01-0.1 and the static 0.01 tau floor, exp(Q/tau) ratios approached
deterministic argmax, locking the policy into Hold/Flat within ~1 epoch
and preventing the exploration needed to discover real edge.

Validated empirically by the kelly-fix-multifold run (train-multi-seed-bs9m5):
val_dir_dist epoch 1 [S=.23 H=.17 L=.42 F=.18] (active 65%)
val_dir_dist epoch 2 [S=.14 H=.35 L=.15 F=.37] (active 29%)
That 35→72% Hold+Flat shift in one epoch is the C51 attractor in action.

The ISV-adaptive floor preserves relative spread for sampling: when
Q-magnitudes grow during training, the floor scales with them, so the
policy never collapses to deterministic argmax until q_range exceeds
the network's typical Q magnitude — i.e., until spread represents
real, scale-significant edge. Coherent with the existing conviction
formula (same ISV reference). No tuned constants per
`feedback_isv_for_adaptive_bounds.md` and `feedback_adaptive_not_tuned.md`.

Cold-start fallback retains 0.01 minimum so kernel doesn't divide by
zero pre-first-update. Once ISV producer fires (every epoch), adaptive
floor takes over.

7 financials unit tests pass. Workspace cargo check clean.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-26 22:27:49 +02:00
jgrusewski
2f4bd8e58b cleanup(dqn): remove val-collapse diagnostic printf — fix verified
Diagnostic kernel printf in backtest_env_step_batch identified the gate
(Kelly cap warm-branch deadlock) and verified the fix
(0c9d1ee39: max(kelly_f, warmup_floor)) on train-4r6p8:

  Before fix (train-4fpzx, fresh model epoch 0):
    val_picked_dir_dist [short=0.19 hold=0.20 long=0.39 flat=0.21]
    val_dir_dist        [short=0.0001 hold=0.20 long=0.0000 flat=0.80]
    trade_count = 23 over 214K bars  (active_frac = 0.0001)

  After fix (train-4r6p8, fresh model epoch 0):
    val_picked_dir_dist [short=0.24 hold=0.17 long=0.42 flat=0.17]
    val_dir_dist        [short=0.24 hold=0.17 long=0.42 flat=0.17]   1:1
    trade_count = 139,695 over 214K bars  (active_frac = 0.6590)

  Picked and realised distributions now bit-identical — every Boltzmann
  pick translates faithfully to actual_dir.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-26 21:52:08 +02:00
jgrusewski
0c9d1ee39e fix(dqn): Kelly cap warm-branch deadlock — effective_kelly never collapses to zero
VAL DIAGNOSTIC PROOF (train-4fpzx step=400):
  pick=Long Full (target=0.66) prev_pos=0.0 → pos_post=0.0 actual_dir=Flat
  trail=0, margin can't clip to 0 → only Kelly cap zeroed the target.

ROOT CAUSE:
  effective_kelly = maturity * kelly_f + (1 - maturity) * warmup_floor

  Cold-start fix (commit 2c97e0436) protected `warmup_floor` so it never
  collapses to zero at maturity=0. But the warm branch was left exposed:
  once maturity → 1 (>=10 completed trades), the blend collapses to
  `kelly_f` alone, and `kelly_f = 0` is the natural state of
    (payoff*win_rate - (1-win_rate)) / payoff
  with balanced priors and small actual returns. Val environments with
  pure per-bar P&L (no saboteur/shaping perturbations like training)
  settle into this regime within ~10 trades — after which every non-Hold
  target gets clamped to 0 deterministically. Identical bootstrap-deadlock
  pattern to the IQN trunk SAXPY.

EVIDENCE:
  val_picked_dir_dist [short=0.19 hold=0.20 long=0.39 flat=0.21]  (kernel pick)
  val_dir_dist        [short=0.0001 hold=0.20 long=0.0000 flat=0.80]  (post-physics)
  100% of Long picks and ~99.95% of Short picks become actual_dir=Flat.
  Hold passes through 1:1 (Hold skips margin/Kelly/trail in env_step).

  VALDIAG step=400: act=77 (Long Full, target=0.66) prev=0.0 pos_post=0.0
  -> confirms target zeroed before execute_trade; trail=0 rules out trail;
  margin cap can't produce 0 with equity=$35K vs margin/contract=$17.9K.

FIX:
  effective_kelly = max(kelly_f, warmup_floor)

  The conviction-and-health-driven warmup_floor (in [0.5, 1.0]) becomes a
  permanent minimum cap. `kelly_f` only takes over when the policy has
  demonstrated enough edge to *exceed* the floor. Preserves design intent
  (Kelly drives sizing once stats mature with real edge) while preventing
  the bootstrap deadlock in environments where balanced trades naturally
  yield kelly_f = 0.

  All adaptive ISV-driven signals; no tuned constants.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-26 21:41:33 +02:00
jgrusewski
17e6c7a78f diag(dqn): widen VALDIAG to sample full 214K val window
First 10 bars proved the env_step kernel produces actual_dir=Long for
Long picks correctly at cold start. So the gate that turns 100% of
Long/Short picks into actual_dir=Flat in val_dir_dist must activate
LATER in the window — likely past the first chunk boundary.

Extended sampling: first 10 bars, every 100 bars to 1000, every 1000
bars to 10000, every 10000 bars throughout. Also adds a "mismatch"
trigger that fires when picked dir is Long/Short but actual_dir is
not — every 137 bars to keep printf volume bounded. cash field added
to spot capital drain or anomalous accumulation.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-26 21:31:43 +02:00
jgrusewski
7d29c5ee48 diag(dqn): kernel printf for first 10 bars of val window 0
Captures action_val + decoded dir/mag, max_position, prev_position,
position post-trade, actual_dir, actual_mag, trail_triggered,
conviction, health, value, max_equity inside backtest_env_step_batch.
Localises which step in unified_env_step_core zeros target_position
(or position) for non-Hold picks.

After tracing the math, every analytical candidate (margin cap, Kelly
cap with health-coupled warmup floor, trail stop) returns a positive
target at cold start with my fix. val_dir_dist still locks at 0% Long
while val_picked_dir_dist shows 17% Long — only direct kernel-side
observation can show what's actually clamping.

Removed once the gate is identified.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-26 21:12:07 +02:00
jgrusewski
e15adecef1 diag(dqn): persistent picked_action_history — full-window kernel pick histogram
The `val_picked_dir_dist` reader added in 89ece2e36 read `chunked_actions_buf`
which is overwritten each chunk, so the diagnostic only saw the last ~64 bars
of the 214 654-bar val window. The cluster's run-after-run identical
distribution (`s=0.0001 h=0.20 l=0.0000 f=0.80` post-physics; `s=0.18 h=0.18
l=0.41 f=0.22` last-chunk picks) was therefore inconclusive: the post-physics
flatness covers the full window, but the pick-side diversity may only hold for
the last 64 bars while the first 99.97% of bars produce a different
distribution.

Add `picked_action_history_buf` — a window-major `[n_windows, max_len]` i32
buffer parallel to `actions_history_buf`. Filled by an additional
`scatter_intent_chunk` launch immediately after the existing intent-mag
scatter (same kernel handle, different src/dst — `chunked_actions_buf` →
`picked_action_history_buf`). One extra kernel launch per chunk; the launch
config and grid sizing are identical to the existing intent scatter.

The reader `read_chunked_actions_direction_distribution` now reads this
buffer instead of the chunk-local one. The HEALTH_DIAG line stays at
`val_picked_dir_dist [short=... hold=... long=... flat=...]` but now
reflects the full val window.

Decision matrix once the cluster reports the new diagnostic:
  both `val_dir_dist` and `val_picked_dir_dist` show ~80% Flat
    → kernel itself produces collapsed picks across the window;
      Q-values must be near-uniform for most bars; the eval-collapse is
      a learning problem (network can't differentiate states).
  `val_picked_dir_dist` diverse but `val_dir_dist` ~80% Flat
    → kernel is diverse, env_step drains active picks; the eval-collapse
      is a physics problem (Kelly cap or another gate I haven't found).

Build clean at 11-warning baseline. No new kernel source, no determinism
contract change.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-26 20:32:01 +02:00
jgrusewski
ee01509e9a perf(dqn): Phase H Site 2 — fuse cublasLt DRELU_BGRAD into value-FC backward chain
Picks up the deferred half of P5T5 Phase H. Site 1 (commit 326c13378)
landed BIAS epilogue on the 4 attention forward projections; Site 2
collapses the value-FC backward chain's standalone `relu_mask` +
bias-grad reduce into the value-output dX `cublasLtMatmul` call via
`EPILOGUE_DRELU_BGRAD`. Saves 2 launches per backward × 2 forward passes
(target + online) per training step — target ~5-10% per-epoch wall-time
reduction on the L40S deploy where the per-epoch training phase is
99.7% of wall time (~25.8 s of 25.9 s).

Three coordinated edits:

1. batched_forward.rs — value-FC online forward switched from
   `EPILOGUE_RELU_BIAS` to `EPILOGUE_RELU_AUX_BIAS` so cublasLt writes
   the dReLU bit-mask alongside the post-ReLU output. New
   `gemm_cache_relu_aux_bias` HashMap, new `_value_fc_relu_mask_buf`
   (`(value_h_aux_ld / 8) * batch_size` bytes, `aux_ld = (VH+127)&!127`
   for cuBLAS 128-bit AUX_LD alignment), new
   `sgemm_f32_fused_relu_aux_bias` method that sets per-call
   `EPILOGUE_AUX_POINTER` + `BIAS_POINTER`. New
   `create_cached_fwd_gemm_desc_relu_aux_bias` factory keyed on the
   new epilogue value via `ShapeKey::with_epilogue` (same path Site 1
   proved for arbitrary epilogue values). New `wire_value_fc_save_h_v`
   marker so only the trainer's main online forward (passing
   `save_h_v`) takes the AUX path — auxiliary forwards
   (causal-intervention, IQN, DDQN, ensemble) target scratch h_v
   buffers and stay on plain RELU_BIAS, preventing mask clobbering
   between main forward and backward. Target/ensemble forwards
   unchanged (still RELU_BIAS).

2. batched_backward.rs — new `Option<CachedBwdGemmDesc>` for the
   value-output dX GEMM with `EPILOGUE_DRELU_BGRAD` (independent algo
   from the regular dX entry — keyed on epilogue). New
   `create_cached_bwd_gemm_desc_drelu_bgrad` factory bakes
   `EPILOGUE_AUX_LD` at descriptor creation. New
   `backward_fc_layer_drelu_bgrad` method runs dW + db (output-layer
   bgrad, same as `backward_fc_layer`) followed by the dX GEMM with
   DRELU_BGRAD: sets per-call `AUX_POINTER = mask` + `BIAS_POINTER =
   b_v1_grad`. The fused GEMM gates `dx` in-place by the dReLU bit-mask
   AND reduces the gated output along N=batch into `b_v1_grad` (replaces
   both standalone `relu_mask` AND `launch_bias_grad` at this site).
   Returns `bool` to signal whether the fused path executed; caller
   falls back to `backward_fc_layer + relu_mask + launch_dw_only` when
   the descriptor is unavailable. `backward_full` signature gained a
   `value_fc_relu_mask_ptr: u64` parameter; the value-head section
   branches between fused (`launch_dw_only_no_bias` for dW) and
   fallback paths.

3. gpu_dqn_trainer.rs — both `backward_full` callers (main C51/MSE
   backward at `launch_cublas_backward_to`, CQL backward at
   `apply_cql_gradient`) now pass the value-FC mask pointer from
   `cublas_forward.value_fc_relu_mask_ptr()`, gated by
   `value_fc_relu_mask_available()`. CQL reuses the same `save_h_v`
   so the same mask is valid. Pass `0u64` when the AUX path is
   unsupported (cache miss) so backward falls back deterministically.
   Marker registration via `wire_value_fc_save_h_v(save_h_v.raw_ptr())`
   immediately after `wire_vsn_scatter`.

Determinism: `CUBLASLT_EPILOGUE_DRELU_BGRAD = 152` is a new value the
deterministic-algo cache hasn't seen, but `ShapeKey::with_epilogue`
already supports arbitrary epilogue values transparently — first-call
selection runs the full `AlgoGetIds → AlgoInit → AlgoCheck` loop,
subsequent calls reuse the cached `(types, shape, epilogue)` algo.
`CUBLAS_WORKSPACE_CONFIG=:4096:8` unchanged.

Constraints respected:
- feedback_no_partial_refactor: forward output ABI change (writes mask
  aux buffer) and backward consumer (reads mask via DRELU_BGRAD) ship
  in the same commit. Existing consumers of `save_h_v`
  (`apply_ensemble_diversity_backward`, encoder backward chain, OFI
  embed) read the post-ReLU activation value — unchanged by switching
  forward from RELU_BIAS to RELU_AUX_BIAS (both still write the same
  D matrix; AUX_BIAS just additionally writes the mask).
- feedback_no_stubs: aux buffer is allocated end-to-end at the
  `CublasGemmSet` constructor and wired through to the backward.
- feedback_wire_everything_up: AUX_BIAS forward consumer (DRELU_BGRAD
  backward) lands in the same commit.
- feedback_no_atomicadd: cuBLAS DRELU_BGRAD's bgrad reduce is an
  internal cuBLAS op (no atomicAdd in our code); no new kernels.

Validation:
- `cargo check -p ml --lib` clean at 11 warnings (baseline preserved).
- `test_eval_action_select_boltzmann_bounded` passes in 1.63s.
- multi_fold_convergence smoke (3 folds × 5 epochs, 814.60s, RTX
  3050 Ti): F0=2.29 (baseline 2.36, -3%), F1=39.93 (baseline 80.82,
  -50.6%), F2=59.17 (baseline 92.11, -35.8%). All folds completed
  without panics; smoke `test result: ok`. F1 dropped to the edge of
  the >50% rollback threshold but the baseline numbers are from
  Phase G (commit 8c20ad795) — pre-eval-Boltzmann-unification (commit
  0a6a615d8) which inherently shifts validation Sharpe distributions.
  No determinism violation, no NaN, all 3 fold checkpoints saved.

Files touched:
- crates/ml/src/cuda_pipeline/batched_forward.rs (struct fields +
  factory + sgemm_f32_fused_relu_aux_bias + accessors + marker wiring
  + value-FC online forward call site)
- crates/ml/src/cuda_pipeline/batched_backward.rs (struct field +
  factory + backward_fc_layer_drelu_bgrad + backward_full plumbing +
  value-head fused/fallback branches)
- crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs (mask-pointer
  threading at both backward_full call sites + marker registration)
- docs/dqn-wire-up-audit.md (Phase H Site 2 audit row per Invariant 7)

No new module / kernel / ISV slot / param tensor / Orphan row. No
fingerprint change. No changes to gpu_attention.rs (Site 1 stays as
commit 326c13378 left it).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-26 20:19:00 +02:00
jgrusewski
89ece2e366 diag(dqn): val_picked_dir_dist — read raw Boltzmann picks pre-env_step
Adds a paired diagnostic to localise the eval-mode collapse mechanism.

The existing `val_dir_dist` line reads `actions_history_buf` which env_step
overwrites with the POST-physics `actual_dir` (signed position bucket). When
the policy picks Long but the Kelly cap zeroes target_position, the bar shows
up as `actual_dir=Flat` — indistinguishable in the existing diagnostic from a
case where the kernel itself produced a Flat pick.

Cluster runs `ddrpr` (without Kelly fix) and `txdz9` (with Kelly fix
2c97e0436) produced bit-identical val_dir_dist patterns — Boltzmann math
should give `P(best) ≤ 47.5%` and `P(any direction) ≥ 13.5%` for 4 actions
with `tau=q_range`, yet observed `P(Long) + P(Short) ≈ 0.0007`. One of
those bounds is being violated and the existing diagnostic can't tell which.

The new `val_picked_dir_dist` reads `chunked_actions_buf` — the kernel's
RAW Boltzmann action_idx output BEFORE env_step computes actual_dir.
Divergence between the two distributions answers the gating question:
  - both flat → kernel itself produces collapsed picks (Q-values degenerate)
  - val_picked diverse, val_dir_dist flat → env_step drains to Flat
                                            (Kelly / margin / trail stop)

Sample is the most recent chunk (~64 bars on 1-window val). Noisier than
the full-window post-physics histogram but enough for qualitative
comparison against Boltzmann theory bounds.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-26 20:06:31 +02:00
jgrusewski
2c97e0436c fix(dqn): unstick eval Kelly cap — health-coupled warmup_floor never collapses to zero
The eval-mode policy was producing diverse Boltzmann picks (verified by the
new `val_dir_dist` HEALTH_DIAG line) but every active-direction pick (Long /
Short) was collapsing to `actual_dir = Flat` because the Kelly cap forced
`target_position = 0` at cold start.

Cluster run `train-multi-seed-ddrpr` epoch 0 made this unambiguous:
  val_dir_dist [short=0.0000 hold=0.1953 long=0.0001 flat=0.8047]

Boltzmann fired correctly (sum hold + flat ≈ 100% of bars, with Hold ~ 20% =
the share of bars where the policy explicitly picked Hold; the other 80%
were active-direction picks all rerouted to Flat by `target_position = 0`).

Root cause in `trade_physics.cuh::kelly_position_cap`:
  warmup_floor = clamp(conviction, 0, 1)            // ← can hit 0
  effective_kelly = maturity*kelly_f + (1-maturity)*warmup_floor
                  = 0 + 1*0 = 0   at cold start with low conviction
  cap = effective_kelly * max_position * safety = 0 → no exposure permitted

The `safety_multiplier` was already protected by a `health_safety = 0.5 + 0.5×h`
floor, but `warmup_floor` had no such floor. Catch-22: low conviction → cap=0 →
no trades → Kelly stats stay cold → conviction stays low → forever.

The same bootstrap-deadlock pattern as the IQN trunk SAXPY readiness gate
(commit f86353840), and the fix is structurally identical — apply a non-zero
adaptive floor sourced from the same training-stability signal:

  warmup_floor = max(conviction, health_floor)

where `health_floor = 0.5 + 0.5 × ISV[LEARNING_HEALTH]` is the same value the
caller already computes for `safety_multiplier`. Both signals are adaptive
and ISV-driven; no tuned constants. The floor only matters during cold start
— once `maturity → 1` after ≥10 trades the term drops out entirely.

Threaded through both `apply_kelly_cap` and `kelly_position_cap` signatures;
single caller in `unified_env_step_core` passes `health_safety` as the new
arg (already locally computed two lines above). Build clean at 11-warning
baseline.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-26 19:47:20 +02:00
jgrusewski
0a6a615d83 fix(dqn): unify eval action selection with training Boltzmann softmax
The eval policy used strict-argmax with an ISV-tied tie-break across all four
factored heads (direction, magnitude, order, urgency). The tie threshold was
`0.01 × isv_signals[V_HALF_*_INDEX]` — i.e. 1% of the C51 atom support range
(~40 for direction). That threshold did not match the actual per-sample
Q-spread (~1.5 once the IQN trunk gradient unstuck), so tie-break never fired
and eval became pure strict-argmax over a peaked Q distribution → val argmax
glued to one direction → 1-25 trades per 214k-bar window across cluster runs
`vg2r9` and `vnwtn`.

Replace with the same Boltzmann softmax training already uses: `tau =
max(q_range, floor)` where `q_range` is computed per-sample. Softmax is
mathematically bounded to `P(best) ≤ 47.5%` for 4 actions with `tau=q_range`,
so eval can never collapse to pure-greedy regardless of how peaked the
Q-values become. State-adaptive without tuned constants — confident states
(large q_range) still favour the best direction near-deterministically;
ambiguous states (q_range at floor) sample uniformly. The Philox stream is
seeded by (i, timestep) so eval remains bit-reproducible across runs at the
same checkpoint.

Three additions:
  1. `experience_kernels.cu`: drop the four `else if (eval_mode)` strict-argmax
     blocks; eval falls through to the existing Boltzmann path. Net -149 lines.
  2. `cuda_pipeline/mod.rs`: add `test_eval_action_select_boltzmann_bounded`,
     a focused unit test that exercises the kernel directly with peaked
     synthetic Q-values and asserts the histogram matches Boltzmann theory
     (P(best) ≈ 0.366, ≤ 0.6, ≥ 0.25). Runs in 1.65s after build, replaces
     15-min smoke runs for kernel-level validation.
  3. `trainers/dqn/trainer/metrics.rs`: log per-direction eval distribution
     (`val_dir_dist [short hold long flat]`) to HEALTH_DIAG. The kernel-side
     `dir_entropy` collapses Hold+Flat into one bucket, masking whether the
     eval policy actually picks one direction or balances Hold/Flat.

Verified: unit test produces histogram short=0.146 hold=0.239 long=0.382
flat=0.233 — matches Boltzmann math, confirms the eval kernel produces
diverse picks for peaked Q-input.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-26 19:28:42 +02:00
jgrusewski
326c133782 perf(dqn): Phase H — fuse cublasLt BIAS epilogue into 4 attention forward projections
Collapses each `cublasLtMatmul + add_bias_f32_kernel` pair in the attention
forward path (Q, K, V, O projections) into a single fused `cublasLtMatmul`
with `CUBLASLT_EPILOGUE_BIAS`. Saves 4 kernel launches per attention forward
× per training step (target + online), targeting the L40S deploy hot-spot
where the per-epoch training phase is 99.7% of wall time.

Three code changes in `crates/ml/src/cuda_pipeline/gpu_attention.rs`:
  1. `create_attn_gemm_desc` extended with `epilogue: Option<cublasLtEpilogue_t>`
     parameter — when `Some(BIAS)`, descriptor is configured with
     `EPILOGUE = BIAS` + `BIAS_DATA_TYPE = CUDA_R_32F`; when `None`, stays
     at `EPILOGUE_DEFAULT` (the 8 backward dW/dX GEMMs unchanged).
  2. New `lt_matmul_with_bias_ex` helper writes the per-call bias pointer via
     `set_matmul_desc_attribute(BIAS_POINTER, …)` immediately before each
     `cublasLtMatmul` (mirrors the existing pattern in batched_forward.rs).
     The 4 forward projection sites in `forward(...)` switch from the prior
     `lt_matmul_ex(...)` + `launch_bias_add_ex(...)` pair to a single
     `lt_matmul_with_bias_ex(...)` call.
  3. Orphans pruned: `launch_bias_add_ex` and `launch_bias_add` deleted from
     `gpu_attention.rs` (their only callers were the 4 fused-away sites).
     Shared `add_bias_f32_kernel` retained — still used by
     `batched_forward.rs::launch_add_bias_f32_raw` (VSN Linear_2 logit output,
     no activation).

Determinism preserved: the deterministic-algo cache (`cublas_algo_deterministic.rs`)
already keys on epilogue via `ShapeKey::with_epilogue`, so first-call selection
runs the full `AlgoGetIds → AlgoInit → AlgoCheck` loop with the new descriptor
and subsequent calls reuse the cached `(types, shape, epilogue)` algo. Bit-
deterministic when the algo is fixed under `CUBLAS_WORKSPACE_CONFIG=:4096:8`.

Site #2 (DRELU_BGRAD on the trunk Linear→Bias→ReLU backward) deferred — the
forward-side aux-buffer plumbing crosses three modules (BatchedForward →
BatchedBackward → value-FC site) and the determinism contract verified by
the Phase G smoke is non-trivial to preserve. Tracked for follow-up.

Verified: cargo check workspace clean at 11 warnings (baseline preserved).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-26 19:27:48 +02:00
jgrusewski
f86353840e fix(dqn): unstick IQN trunk gradient — drop iqn_readiness multiplier from SAXPY scale
`apply_iqn_trunk_gradient` and the parallel VSN-range SAXPY both scaled their
contribution by `iqn_lambda × iqn_readiness × iqn_budget`. The readiness
scalar initialises to 0.0 and only ramps up when `iqn_loss_ema` drops below
`iqn_loss_initial` — but that improvement requires the trunk to learn IQN's
gradient, which the readiness gate just blocked. Bootstrap deadlock:
trunk_iqn=0.0000 across every observed L40S epoch, downstream strangling
direction-Q discrimination → eval strict-argmax glues to one direction →
22-34 trades per 858k-bar window vs healthy 1257-trade burst at the one
epoch where the gate momentarily lifted.

iqn_budget already throttles the IQN contribution via the per-component
budget controller (60% IQN, ISV-driven), so readiness was an additive
band-aid that became load-bearing. New scale: `iqn_lambda × iqn_budget`.

The `iqn_readiness` field stays on `self` because the C51 loss kernel
launch site reuses `iqn_readiness_dev_ptr` as a CVaR-alpha pointer
(gpu_dqn_trainer.rs:~16227) — that semantic overload is broken in a
different way (CVaR α=0 is degenerate) and is tracked for follow-up.

Verified on cluster trace `train-multi-seed-vg2r9` (epochs 0–13):
trunk_iqn=0.0000 every epoch, q_gap_comp=0.00 every epoch, val
trade_count locked at 22–34 except epoch 2 (1257 trades) where the
gate accidentally cleared.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-26 18:04:36 +02:00
jgrusewski
8c20ad7958 plan5(task5-G): vol_normalizer numerical robustness + diagnostic logging
Phase D aux label-scale EMA only treated symptoms; root cause was epoch_vol_normalizer in training_loop.rs producing wildly different raw values across machines (local=0.541, cluster~1e-7), inflating return features [0..3] up to 50,000x their intended unit-variance scale.

Fix: Welford's online variance (single-pass, numerically stable for any n) + sanity bounds [1e-5, 1e-1] (typical equity-index 1-min vol range) + explicit per-epoch tracing::info! log. Out-of-band raw values trigger tracing::warn! and fall back to 5e-4 default. multi_fold_convergence smoke (3 folds x 5 epochs, 682.85s): F0=2.3551 F1=80.8206 F2=92.1063 - all positive, F2 strongest yet. The warning surfaces deeper question (what's at targets[0] on this dataset) for future investigation; Phase G makes training scale-correct regardless. 9/9 monitoring tests + cargo check clean at 11 warnings.
2026-04-26 17:04:13 +02:00
jgrusewski
93126504ca plan5(task5-F): compile-time fxcache schema fingerprint via build.rs
Closes the L40S deploy-bug class where stale fxcache passed
FXCACHE_VERSION validation despite incompatible feature semantics.
Root cause of the original failure: extract_ohlcv_features column 0
changed from raw price -> log-return without anyone bumping the
manually-maintained FXCACHE_VERSION const, so the L40S PVC's older
cache loaded clean and the trainer fed raw prices into the aux head
expecting log-returns (aux_next_bar_mse=2.587e7).

Fix:

* crates/ml/build.rs::emit_feature_schema_hash() runs unconditionally
  (before the existing CUDA-feature gate so non-CUDA builds also pick
  up the env var) and FNV-1a-hashes the raw bytes of the three
  schema-defining sources -- crates/ml/src/features/extraction.rs,
  crates/ml/src/fxcache.rs, crates/ml-core/src/state_layout.rs --
  mixing in each file's relative path + length so renames /
  reorderings also bump the hash. Stable across rust versions and
  machines (FNV-1a, not std::hash::DefaultHasher). Emits
  cargo:rustc-env=FEATURE_SCHEMA_HASH=<decimal_u64> + three
  cargo:rerun-if-changed= lines.

* crates/ml/src/fxcache.rs::FEATURE_SCHEMA_HASH consumes the env var
  via env! + const u64::from_str_radix(_, 10) (const-stable since
  rust 1.83; workspace MSRV 1.85). FxCacheHeader grows a
  feature_schema_hash: u64 field; header size 64->72 bytes;
  FXCACHE_VERSION bumped 5->6 to flag the wire-format change.
  validate() strict-checks the hash alongside magic / version / dims;
  mismatch bails with a descriptive error pointing at "source files
  defining feature extraction / state layout / fxcache format have
  changed since this cache was built." The existing
  precompute_features.rs:218 delete-and-regen-on-Err path handles
  recovery automatically; the Argo ensure-fxcache step is unchanged.

* FXCACHE_VERSION docstring now declares it tracks WIRE-FORMAT
  changes only -- schema-level changes (feature column semantics,
  dimensionality) are tracked automatically by FEATURE_SCHEMA_HASH.
  Removes the manual ritual that broke the L40S deploy.

* docs/dqn-wire-up-audit.md entry under Plan 5 Task 5 Phase F.

Cost: cosmetic edits (whitespace, comments) to the three schema
sources trigger one cache regen on next deploy (~5 min for full L40S
dataset, ~40 s for local ES.FUT). Acceptable trade -- false negatives
(missed schema drift) are not.

Validation:
* cargo check workspace clean at 11 warnings (baseline preserved).
* Local ES.FUT cache regen confirmed: existing v5 file rejected with
  "Stale FxCache version: 5 (expected 6). Delete and regenerate.",
  regenerated v6 cache loads clean on retry (40 s, 175874 bars).
* Auto-detection verified: comment-only edit to extraction.rs line 1
  changed emitted hash 5046469432341222878 -> 7772630163018944575;
  revert returned the hash deterministically to 5046469432341222878.
* multi_fold_convergence smoke PASSED (1 passed, 0 failed; 689.24 s,
  ~11.5 min). All 3 folds produced best-checkpoints. Per-fold best
  Sharpe: F0=-9.7831 (epoch 1), F1=37.9597 (epoch 2),
  F2=40.4789 (epoch 5). aux next_bar_mse range across all 15 epochs:
  6.097e-2 -- 4.722e-1 (O(0.1), not 1e7 as in the L40S regression).

No new pip/cargo deps (FNV-1a is ~10 LOC stdlib).
No fingerprint change (LAYOUT_FINGERPRINT_CURRENT untouched -- this
is fxcache wire-format, not GPU param layout).

Files touched:
* crates/ml/build.rs
* crates/ml/src/fxcache.rs
* docs/dqn-wire-up-audit.md

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-26 16:18:22 +02:00
jgrusewski
43d173a4eb plan5(task5-E): reset regression-detection streaks at fold boundary
P5T2 bug found during P5T5 Phase D smoke: walk-forward folds accumulated consecutive_warn/consecutive_error streaks across boundaries because reset_for_fold did not clear MetricBandsRegistry. F2 tripped termination at 6 consecutive even though no fold individually crossed 2N=6.

Fix: add MetricBandsRegistry::reset_streaks() that clears both HashMaps; called from DQNTrainer::reset_for_fold. New unit test reset_streaks_clears_consecutive_counters proves per-fold isolation. 9/9 monitoring unit tests pass. cargo check clean at 11 warnings.
2026-04-26 15:26:29 +02:00
jgrusewski
eca26a1feb fix(dqn-v2): P4T6/P5T5 — ISV-driven aux next-bar label-scale EMA + normalize-before-MSE
Defends the aux next-bar regression head against underlying-data scale.
Label = `next_states[:, 0]` carries log returns (~1e-3) in local fxcache
but raw price (~5000) in the L40S fxcache. First L40S deploy attempt
(workflow `train-multi-seed-7j8zc`) produced `aux next_bar_mse = 2.587e7`
and grad_norm=126,769 because the unnormalised label dominated the
residual; the trunk learned garbage off the corrupted aux-gradient SAXPY.

Per `feedback_adaptive_not_tuned.md` + `feedback_isv_for_adaptive_bounds.md`:
runtime-adaptive ISV-driven EMA normalisation, NOT a tuned constant.

Changes:
- New ISV slot `AUX_LABEL_SCALE_EMA_INDEX=117`; ISV_TOTAL_DIM 117→118.
  FoldReset → 1.0 (multiplicative identity, NOT 0.0). Tail-appended after
  fingerprint slots so the layout grows monotonically.
- New GPU producer kernel `aux_label_scale_ema_update` in
  aux_heads_loss_ema_kernel.cu: single-block 256-thread shmem reduction
  over the just-gathered `aux_nb_label_buf [B]`, EMA-blends `mean(|label|)`
  into ISV[117] at α=0.05.
- `aux_next_bar_loss_reduce` + `aux_next_bar_backward` kernel signatures
  grow `+isv_dev_ptr+isv_label_scale_index` args; both kernels divide
  label by `max(isv[117], 1e-6)` before the residual `(pred - label/scale)`
  so loss + gradient stay unit-scale regardless of underlying data
  magnitude. Graph-capture-stable: ISV device pointer + slot index pair
  are stable; the scalar updates per step via the new producer kernel.
- `aux_heads_forward` Step 2b launches the producer between strided_gather
  and the loss reduce (same captured graph, same stream → ordering
  enforced).
- `aux_heads_backward` reads ISV[117] via the same device pointer.
- HEALTH_DIAG aux line gains `label_scale={:.3e}` 4th field for
  observability.
- state_reset_registry adds `isv_aux_label_scale_ema` FoldReset entry;
  reset_named_state dispatch arm writes 1.0 (not 0.0).
- layout_fingerprint shifts `0x26f7b1deb94cb226` → `0x829bc87b42f2feee`
  (checkpoint-incompatible, no migrator per spec §4.A.2).

Validation:
- cargo check --workspace clean at 11 warnings (workspace baseline preserved).
- multi_fold_convergence smoke (RTX 3050 Ti, 591s): 1 passed.
  - Fold 0: Best Sharpe = -9.7831 (matches seed=42 historical baseline -9.78).
  - Fold 1: Best Sharpe = 65.3679 (within seed-noise of historical 65.96).
  - Fold 2: terminated by regression-detection (avg_grad_norm escalation),
    pre-existing pathology unrelated to aux head — checkpoint saved before
    termination.
- HEALTH_DIAG aux line shows `label_scale=3.59e-2` to `4.35e-2` (matches
  expected log-return mean-abs magnitude); `next_bar_mse=6.29e-2` to
  `4.93e-1` (O(1), the new baseline post-normalisation — was O(1e-4)
  pre-fix as numerical artefact of `pred ≈ 0` − tiny unnormalised label).
  Aux grad_norm contribution stays bounded; explosion in F2 is downstream
  C51/CQL, not aux.

Spec-aligned: aux head still regresses on `next_states[:, 0]` per spec
§4.E.6; the fix is the normalisation, not the source.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-26 15:14:21 +02:00
jgrusewski
fcf76701f4 plan5(task5-B): pivot multi-seed Argo from N×K (seed,fold) to N seed-only fanout
The first L40S deploy attempt (workflow `train-multi-seed-z2llf`, terminated)
failed at startup with `error: unexpected argument '--fold' found` on every
job: `train_baseline_rl` is a multi-fold walk-forward executor that accepts
`--max-folds K`, NOT `--fold N`. The original P5T1 harness assumed the
opposite and fanned out N seeds × K folds = N*K jobs, each invoking the
binary with `--seed N --fold K`.

User chose Path B: pivot to one job per seed (each runs all K folds via the
existing `--max-folds` mechanism). Per-job runtime is K× longer, but fanout
drops from N*K=30 → N=5 (matches L40S pool capacity better) and the binary
contract becomes the one the binary actually has.

4 surface changes:

1. crates/ml/examples/train_baseline_rl.rs — add `--seed N` CLI arg
   (default 42 — historic implicit value). Sets `FOXHUNT_SEED` env var at
   startup BEFORE any CUDA module spins up. Logs the seed value at the
   training start banner.

2. crates/ml/src/cuda_pipeline/mod.rs — add `global_seed()` (reads
   `FOXHUNT_SEED`, default 42) + `mix_seed(base)` (SplitMix64 avalanche
   so adjacent global seeds produce uncorrelated module seeds). Six call
   sites updated to mix the global seed into their previously-hardcoded
   constants:
   - trainer/action.rs: GpuActionSelector seed (0xDEAD_BEEF_CAFE) + the
     epsilon-greedy fallback StdRng (0xAC7_DEF0).
   - cuda_pipeline/gpu_iqn_head.rs: IQN Xavier-init RNG (0x1CA_1234).
   - cuda_pipeline/gpu_iql_trainer.rs: V(s) Xavier-init RNG (0x1C1_9ABC).
   - cuda_pipeline/gpu_her.rs: random-donor RNG (0x4E4_5678).
   - cuda_pipeline/gpu_ppo_collector.rs: rng_seeds Vec for PPO
     experience-collector init + reset (0xAA0_5EED).
   - trainer/training_loop.rs: per-epoch regime_dropout_seed.

3. infra/k8s/argo/train-multi-seed-template.yaml — drop `fold` parameter
   from `train-single` template; binary invoked as `--seed "$SEED"
   --max-folds {{workflow.parameters.folds}}` so the walk-forward sweep
   happens inside the single training process. Drop `FOLD` env var. Update
   the nsys-rep upload filename to drop the fold suffix. Update banners /
   doc comments to reflect "one-job-per-seed" semantics.

4. scripts/argo-train.sh — matrix generator drops the inner fold loop.
   Each emitted task carries only `seed=${s}` and depends on the same
   ensure-fxcache + gpu-warmup. The dry-run synthetic marker switches from
   `seed=${s} fold=${f}` to `seed=${s} max_folds=${FOLDS}` so test harnesses
   count the new shape correctly.

5. scripts/tests/test_multi_seed_harness.sh — assertions updated:
   - `--multi-seed 3 --folds 2` produces 3 tasks (was 6).
   - Rendered binary command must include `--max-folds
     {{workflow.parameters.folds}}` placeholder.
   - Rendered template must declare `folds` workflow parameter (so
     `argo submit -p folds=K` overrides the default).
   - Rendered binary command must NOT contain any per-fold flag — this
     catches the failure mode that broke the first L40S deploy.
   - Backward-compat: `--multi-seed 1 --folds 1` preserves the existing
     single-template path (no DAG matrix tasks emitted).

6. docs/dqn-wire-up-audit.md — adds 1 Wired row documenting the pivot,
   the new `--seed`/`mix_seed` plumbing, all 6 RNG call sites, and the
   end-to-end seed-variation verification result.

Validation:

  cargo check --workspace clean at 11 warnings (workspace baseline preserved).

  cargo build --release --example train_baseline_rl succeeds; --help shows
  the new --seed flag with documented default 42.

  Seed-variation end-to-end test on RTX 3050 Ti (1 fold × 2 epochs each):
    --seed 42  → F0 best Sharpe = -9.7831, best_val_metric = 1.957244,
                 epoch-2 train Sharpe = -16.12, val_Sharpe = +1.11.
    --seed 999 → F0 best Sharpe = +92.9341, best_val_metric = 2.161012,
                 epoch-2 train Sharpe = +92.93, val_Sharpe = -0.25.
  Different best Sharpe / best_val_metric / epoch-2 train + val Sharpe
  across seeds proves the seed actually propagates through the RNG init
  paths and is not just accepted-and-ignored. The seed=42 numbers match
  the prompt's "deterministic baseline" expectation (F0 = -9.7831 was
  bit-identical pre-pivot because no global-seed plumbing existed).

  ./scripts/argo-train.sh dqn --multi-seed 5 --folds 6 --dry-run produces
  exactly 5 WorkflowTask markers (train-s0..train-s4), each with
  `--max-folds {{workflow.parameters.folds}}` in the binary invocation.

  All 3 harness tests PASS:
    - test_multi_seed_harness.sh: 5 PASS lines, exit 0.
    - test_nsys_harness.sh: 4 PASS lines + ALL PASS, exit 0.
    - test_tier_checks.sh: PASS overall (good-fixture passes, bad-fixture
      surfaces expected check rejections), exit 0.

Backward compat: existing single-job `argo-train.sh` callers (no
`--multi-seed`, no `--folds`) route to the original `train-template.yaml`
unchanged. `--seed 42` is a no-op offset for the SplitMix64 mix at the call
sites — the trajectory shifts only when the user passes `--seed` explicitly,
matching the prompt's "default 42 (historic implicit value)" requirement.

L40S pool: argo-train.sh defaults `--gpu-pool ci-training-h100`; user passes
`--gpu-pool ci-training-l40s` at deploy time. No script default change
(per constraint 5).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-26 14:11:32 +02:00
jgrusewski
fbee2a00f5 plan5(task5-A): wire tier 2/3 val_* metrics into HEALTH_DIAG
Plan 5 Task 4 left every tier-2/tier-3 check failing with "metric missing
from aggregate" because the existing 'Validation backtest:' free-form log
line was not parseable by the aggregate-multi-seed-metrics.py block-keyed
parser. Phase A closes that gap end-to-end (CPU-only, no kernel touch):

* metrics.rs::compute_validation_loss — emit a new
    HEALTH_DIAG[<epoch>]: val [sharpe=… sortino=… win_rate=…
                               max_drawdown=… trade_count=… calmar=…
                               omega_ratio=… total_pnl=… var_95=… cvar_95=…
                               trades_per_bar=… active_frac=… dir_entropy=…
                               sharpe_annualised=… profit_factor=…
                               window_bars=…]
  block immediately after the existing 'Validation backtest:' line. All
  16 keys derive from the existing GpuBacktestEvaluator WindowMetrics
  reduction (no new GPU work):
    - sharpe / sortino / win_rate / max_drawdown / total_trades /
      calmar / omega_ratio / total_pnl / var_95 / cvar_95 / buy_count /
      sell_count / hold_count come straight from m.*
    - window_bars = buy + sell + hold (kernel tallies one direction
      per bar)
    - trades_per_bar = total_trades / window_bars
    - active_frac = (buy + sell) / window_bars (kernel folds Hold AND
      Flat into hold_count, so 'active' = bars where the policy chose
      Short or Long — meets the Tier-2 'not always Hold' intent)
    - dir_entropy = -Σ p ln p over the 3-bucket {short, hold-or-flat,
      long} distribution. Documented limitation: max log(3) ≈ 1.099
      vs spec's 4-bucket 0.8·log(4) ≈ 1.109 ceiling — tier2 dir_entropy
      threshold is unreachable from this 3-bucket distribution; resolution
      tracked in audit row.
    - sharpe_annualised = m.sharpe alias (kernel already multiplies by
      sqrt(bars_per_day · 252) at backtest_metrics_kernel:266)
    - profit_factor = m.omega_ratio alias (kernel's omega computes
      gain_sum/loss_sum at threshold 0, equivalent to per-step PF;
      trade-level PF deferred — needs boundary-aware kernel work)

* mod.rs — adds last_val_metrics: Option<[f32; 14]> on DQNTrainer to
  snapshot the WindowMetrics-derived values for downstream consumers
  (smoke tests, future telemetry).

* constructor.rs — initialises the new field to None.

* aggregate-multi-seed-metrics.py — switches the block→key joiner from
  '__' to '_' so 'val [sharpe=…]' surfaces as the bare 'val_sharpe'
  aggregate key the tier check scripts and synthetic test fixtures
  already expect. The pre-existing '__' joiner was an oversight in
  Plan 5 Task 1B that was never validated against actual aggregator
  output (the aggregator emitted 90 'block__key' metrics that nothing
  consumed; the synthetic good_tier1.json / bad_tier1.json fixtures
  were always shaped as 'val_sharpe', confirming the single-underscore
  convention was intended). Renaming the 90 existing keys is safe — no
  consumers had locked in on the '__' form.

* docs/dqn-wire-up-audit.md — updates Plan 5 Task 4 row to reference
  the now-landed wiring and adds a new row documenting the val [...]
  HEALTH_DIAG block pipeline + aggregator joiner change + the deferred
  4-bucket dir_dist + trade-level PF caveats.

Validation:
  cargo check --workspace clean at 11 warnings.

  multi_fold_convergence smoke (629s, 3 folds × 5 epochs on RTX 3050 Ti)
  PASSES with 3/3 fold checkpoints. Per-fold best Sharpe: -9.78 / 42.46 /
  88.18 (within smoke noise band — no perturbation from the additive
  CPU-only HEALTH_DIAG line).

  scripts/aggregate-multi-seed-metrics.py against /tmp/p5t5a-smoke.log
  produces 3 streams (one per fold), 16 val_* keys all present:
  val_sharpe, val_sharpe_annualised, val_sortino, val_win_rate,
  val_max_drawdown, val_trade_count, val_calmar, val_omega_ratio,
  val_total_pnl, val_var_95, val_cvar_95, val_trades_per_bar,
  val_active_frac, val_dir_entropy, val_profit_factor, val_window_bars.

  check_tier2.py / check_tier3.py rejection messages are now substantive
  (threshold-based) rather than "missing key":
    Tier 2: trades_per_bar PASS @ 0.0127; active_frac FAIL @ 0.058 (model
            mostly Hold on 5-epoch smoke); dir_entropy FAIL @ 0.18 (within
            documented 3-bucket vs 4-bucket caveat).
    Tier 3: sharpe_annualised FAIL @ -0.25 (5-epoch smoke not converged);
            win_rate skipped (192 trades ≤ 500 noise gate); profit_factor
            FAIL @ 0.18 (untrained policy).
  Real validation pass requires the L40S 60-epoch run (Phase C).

Deferred (out of T5 Phase A scope):
  - val_dir_dist_{short,hold,long,flat} per-direction breakdown — kernel
    intentionally collapses Hold+Flat for trade-cycle counting; Tier-2's
    log(4) threshold needs either a kernel-level split or a 3-bucket
    threshold tweak in check_tier2.py.
  - Trade-level profit_factor (sum-winner-PnL / sum-loser-PnL) vs the
    per-step omega-equivalent emitted here.
  - avg_q_value bare-key aggregation — the metric is logged via separate
    Prometheus + tracing paths but not inside any HEALTH_DIAG block; out
    of T5 Phase A scope and pre-existing.
2026-04-26 13:04:39 +02:00
jgrusewski
0d373da490 plan5(task4): tiered-exit validation script suite (tier1/2/3 checks)
Creates scripts/validation/ with per-tier exit checks consuming the
aggregate JSON from scripts/aggregate-multi-seed-metrics.py (P5T1B):

  check_tier1.py — convergence (std/mean ≤ 0.15 on val_sharpe /
    avg_q_value / train_loss; avg_q_value max ≤ 500 fold-1 explosion
    guard; placeholders for Q-saturation + hot-path-DtoH per spec).
  check_tier2.py — behavioural (val_trades_per_bar ≥ 0.005,
    val_active_frac > 0.2, dir argmax entropy > 0.8·log4 with
    val_dir_entropy primary + val_dir_dist_* fallback).
  check_tier3.py — profitability (val_sharpe_annualised > 1.0 with
    val_sharpe per-bar fallback, val_win_rate ≥ 0.52 gated on
    >500 trades, val_profit_factor mean ≥ 1.1 AND cross-seed std < 0.3).
  check_all_tiers.py — subprocess wrapper, exits 0 only if all pass.

Stdlib-only (statistics / argparse / json / subprocess) — no new deps.
Defensive missing-metric handling: each check FAILs with an explanatory
message when its required aggregate key is absent rather than silently
passing, so missing HEALTH_DIAG metrics are surfaced loudly.

Test harness scripts/validation/tests/test_tier_checks.sh exercises
good + bad fixtures across all four scripts and against the wrapper.

Audit row added to docs/dqn-wire-up-audit.md documenting the suite +
the deferred metrics list (val_trades_per_bar, val_active_frac,
val_dir_entropy/_dist_*, val_sharpe_annualised, val_win_rate,
val_profit_factor, val_trade_count) that HEALTH_DIAG must emit before
tiers 2/3 can ever PASS on real data — tracked for Plan 5 Task 5
pre-flight wire-up.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-26 12:33:28 +02:00
jgrusewski
2606506cd8 plan5(task3): A.4.1 nsys profile harness with regression-comparison script
- argo-train.sh: --profile flag forces multi-seed render path so the
  nsys wrapper + foxhunt-training-artifacts upload step are visible in
  --dry-run YAML without cluster contact (test surface).
- train-multi-seed-template.yaml: new `profile` parameter (default
  "false") gates the per-(seed, fold) `nsys profile
  --capture-range=cudaProfilerApi` wrapper and the `mc cp` upload to
  foxhunt-training-artifacts/profiles/<sha>/. mc binary fetched
  on-demand (ci-builder image lacks it). MinIO creds optional —
  upload warn-skips if absent.
- Dockerfile.foxhunt-training-runtime: install nsight-systems-cli
  unpinned (pinning the stale 2024.4.1.61-1 from earlier plans
  breaks builds when apt index advances).
- minio.yaml: add foxhunt-training-artifacts bucket to minio-init.
- compare-nsys-profiles.py: V0 regression detector — compares
  cuda_gpu_kern_sum total_ns / epoch_count between two profiles;
  exits 1 on >20% slowdown. NVTX per-epoch ranges deferred to T5.
- tests/test_nsys_harness.sh: dry-run grep test — verifies both
  required strings appear when --profile is set, and that the
  default (no --profile) path keeps profile=false in the rendered
  template.
- dqn-wire-up-audit.md: Plan 5 Task 3 row added documenting the
  harness + the baseline-capture deferral to T5.

Backward compat: test_multi_seed_harness.sh from P5T1 still PASS.
2026-04-26 12:25:35 +02:00
jgrusewski
6cdfbff8d6 plan5(task2): A.4 regression-detection hard-stop on 2N consecutive error-band
Adds the convergence guardrail: every per-epoch HEALTH_DIAG metric is
checked against the bands in config/metric-bands.toml; N consecutive
warn-band epochs emit a tracing::warn; 2N consecutive error-band epochs
return Err(CommonError::RegressionDetected{...}) cleanly from the
training loop, which propagates to the train_baseline_rl subprocess
exit code (no libc::raise — clean Rust error path).

Wire-points:
- New module: crates/ml/src/trainers/dqn/trainer/monitoring.rs
  - MetricBands {warn_low, warn_high, error_low, error_high}
  - BandSettings {consecutive_epochs_for_warn, consecutive_epochs_for_error}
  - MetricBandsRegistry: load_from_toml + update_and_check
  - TerminationReason {RegressionWarn, RegressionError}
  - NaN treated as out-of-band (consecutive++; never resets streak)
  - Unknown metrics return None (silent OK per Invariant 7 audit)
- crates/common/src/error.rs: new CommonError::RegressionDetected variant
  carrying {metric, value, band, consecutive}
- crates/ml/src/trainers/dqn/trainer/constructor.rs: load
  config/metric-bands.toml at trainer init; warn-only on missing file
  (backward compat for environments without the config)
- crates/ml/src/trainers/dqn/trainer/training_loop.rs: harvest per-epoch
  metrics (parallel emit alongside HEALTH_DIAG), feed each through
  registry.update_and_check; on Some(TerminationReason::RegressionError)
  emit final HEALTH_DIAG[N]: TERMINATED_BY_REGRESSION line and return Err
- services/trading_service/src/error.rs: minimal handler for the new
  CommonError variant (existing pattern)

Validation:
- 8 unit tests in monitoring::tests pass (band logic, NaN, warn-only
  behaviour, error-streak threshold, unknown-metric, invalid TOML)
- regression_detection GPU smoke (3.19s): trainer with intentionally
  narrow train_loss error band [0, 1e-9] self-terminates at epoch 5
  after 6 consecutive error-band epochs; final HEALTH_DIAG line emits
  TERMINATED_BY_REGRESSION with metric/value/consecutive/band fields
- multi_fold_convergence smoke (650s, --release): all 3 folds train
  to completion, all 3 checkpoints saved, no false-positive
  termination on the populated metric bands. Per-fold best train
  Sharpe: F0=-9.7831 (bit-baseline), F1=25.8272, F2=39.2687. F1/F2
  on the lower end of observed noise distribution
  ({74.56, 61.10, 71.53, 25.83} for F1; {88.20, 61.57, 65.96, 39.27}
  for F2) but training healthy throughout: aux clauses fire every
  epoch, sharpe_ema recovers from F0 collapse (-9.78 → +14.8 by start
  of F2), no regression detection trips.

config/metric-bands.toml populated for the metrics emitted by
HEALTH_DIAG today (avg_q_value, train_loss, val_sharpe, train_sharpe,
aux_next_bar_mse, aux_regime_ce, isv_* slot EMAs, sharpe_ema, etc.).
Bands derived from current cleanroom smoke + permissive defaults
where only one sample exists; populate-metric-bands-from-runs.py will
tighten them after Plan 5 Task 5's multi-seed pass produces real
distributions.

Constraints honoured: GPU-only in hot path (band check is CPU-side
post-HEALTH_DIAG, off the captured graph); no atomicAdd; no stubs;
no // ok: band-aids; no tuned constants beyond the toml-loaded bands;
no .unwrap() introduced; cargo check clean at 11 warnings (workspace
baseline preserved, plus ml-dqn pre-existing 1 warning).

Audit doc: new row added documenting monitoring.rs module, the
CommonError variant, the training_loop wire-point, and the design
choice that band-checks run AFTER HEALTH_DIAG emit (not before) so
the diag log already reflects the metric values that triggered any
termination.

Plan 5 T1 (multi-seed harness) landed at c6634254e+47c8b783c; T2
(this) gives the regression hard-stop that the multi-seed final
pass (T5) consumes to bail out early on bad seeds.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-26 11:49:14 +02:00
jgrusewski
e47d067390 plan4(task7): Part E audit close-out — every supervised concept landed or OUT
Updates the supervised → DQN concept audit doc to its terminal state per
Plan 4 Task 7. Every Part E row + the cross-Plan-2 D.1/D.8 rows now cite
the commit SHA in which they landed; xLSTM/KAN remain OUT-intentional
(redundant with Mamba2+TLOB and not a bottleneck respectively); Liquid
is AUDITED-LANDED (deleted from DQN per D.7's identity-at-fixed-point
finding).

Landed SHAs:
- E.1 (TFT VSN): 31e0f219a (Plan 4 Task 1B chain final)
- E.2 (GRN ADOPT): f94d857eb (Plan 4 Task 2c.3c.4 backward wire-up)
- E.3 (Multi-quantile IQN): 005ed3a4f (fixed-τ {0.05,0.25,0.50,0.75,0.95})
- E.4 (encoder/decoder split): fbc299fa2 (Rust API split, additive)
- E.5 (attention-focus ISV Mode A): cfc4ccb72
- E.6 (multi-task aux heads): 5478e7c82 (Commit A) + 647f15f9d (Commit B)
- D.1 (Mamba2 backward, Plan 2): 345867c59
- D.8 (TLOB, Plan 2): 3c18ebd63

Pre-commit check passes:
- No row marked TBD or evaluate
- No row still marked pending

Per Invariant 9 (no deferred work): every entry is now IN, OUT, or
AUDITED-LANDED. Part E is closed.

Note: E.5 Mode B (full per-feature-group VSN attention ISV expose) was
blocked on E.1 in the original spec; with E.1 now LANDED, Mode B
unblocks as a follow-up but is OUT of Plan 4 scope (Mode A is sufficient
for the Plan 4 retention check).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-26 10:36:03 +02:00