Commit Graph

4182 Commits

Author SHA1 Message Date
jgrusewski
2ecf36b94c refactor(dqn): Plan 1 Task 4E — DIR_* direction sub-index named constants (Invariant 8)
Define DIR_SHORT=0, DIR_HOLD=1, DIR_LONG=2, DIR_FLAT=3, NUM_DIRECTIONS=4 in
state_layout.cuh. Migrate all literal dir_idx/raw_dir comparisons to named
constants across experience_kernels.cu (15 sites), trade_physics.cuh (3 sites),
and backtest_metrics_kernel.cu (2 sites). Raw integer comparisons (== 0/1/2/3)
eliminated from all direction-branch consumers.

Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com>
2026-04-24 11:37:00 +02:00
jgrusewski
fc2d65c1a3 refactor(dqn-v2): Invariant 8 — named constants for branch indices
Add BRANCH_DIR/BRANCH_MAG/BRANCH_ORD/BRANCH_URG/NUM_BRANCHES to
state_layout.cuh. Migrate the 4 clear literal branch-index accesses
in branch_grad_balance_kernel.cu (branch_norms_dev[0..3] in
grad_balance_isv_update).

Other branch-keyed arrays (branch_starts, branch_lens, branch_norms,
branch_scales) use the runtime `branch = blockIdx.y` variable or loop
counter — no raw literal index, so no migration needed. Rust call sites
use struct fields (branch_0_size etc.) or loop vars — not literals.

No behavioural change; pure refactor.

Plan 1 Task 4D. Spec §3 Invariant 8.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-24 11:37:00 +02:00
jgrusewski
a5a67ad19a refactor(dqn-v2): Invariant 8 — named constants for plan_params[0..6)
Add PLAN_PARAM_* constants to state_layout.cuh and replace raw
plan_params/pp slot accesses in experience_kernels.cu and
backtest_plan_kernel.cu.

Sites migrated: pp[0..5] in both files (12 sites), plan_params_ptr
indexed accesses at slots 0, 1, 4 in experience_kernels.cu (4 sites),
plan_params[w*6+4] in backtest_plan_kernel.cu (1 site). Loop variable
pp[p] in noisy-plan path left as-is (p is a runtime loop counter, not
a literal slot index).

No behavioural change; pure refactor.

Plan 1 Task 4C. Spec §3 Invariant 8.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-24 11:37:00 +02:00
jgrusewski
f36b1bde7a refactor(dqn-v2): Invariant 8 — named constants for plan_isv[0..6)
Add PLAN_ISV_* constants to state_layout.cuh and replace raw
plan_isv[N] and pisv[N] accesses in experience_kernels.cu and
backtest_plan_kernel.cu. backtest_plan_kernel.cu gains an
#include "state_layout.cuh" so the constants are visible.

6 code sites migrated (plan_isv[0..5] in experience_kernels.cu),
plus 6 pisv[N] local-array accesses in backtest_plan_kernel.cu
(same semantic slots).

No behavioural change; pure refactor.

Plan 1 Task 4B. Spec §3 Invariant 8.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-24 11:37:00 +02:00
jgrusewski
7a2adeb374 refactor(dqn-v2): Invariant 8 — named constants for portfolio state ps[0..PS_STRIDE=38)
Replace raw ps[N] access across all CUDA kernels with PS_* named
constants defined in state_layout.cuh. 32 constants total covering
the full 38-slot portfolio state: position/cash/value, DSR stats,
Kelly stats, plan fields, and OFI scratch range.

Notable deviations from pre-written plan mapping: the plan's
PS_VALUE/PS_POSITION/PS_CASH values (0,1,2) had the wrong semantics.
Code wins (feedback_trust_code_not_docs): ps[0]=position, ps[1]=cash,
ps[2]=portfolio_value. All 148 raw ps[digit] accesses in
experience_kernels.cu and trade_stats_kernel.cu migrated. In-comment
range references (e.g. "ps[30..37]") left as documentation.

trade_stats_kernel.cu: migrated base+14 offset to PS_KELLY_WIN_COUNT.
docs/dqn-named-dims.md: PS_* table corrected to match actual code layout.

No behavioural change; pure refactor.

Plan 1 Task 4A. Spec §3 Invariant 8.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-24 11:37:00 +02:00
jgrusewski
46e99c5ccc feat(dqn-v2): A.1 wire StateResetRegistry into fold-boundary reset
Replaces the ad-hoc scattered fold-boundary reset calls with a single
registry-driven iteration. Adding new fold-reset state now requires
adding a registry entry AND a dispatch arm in reset_named_state in the
same commit (Invariant 2 Wire-It-Up).

Step 3.4 correction: plan_state entry removed from the registry —
plan_state_buf exists only in GpuBacktestEvaluator (val path), not in
the training-path fused ctx. No training-side fold reset is applicable.

New behaviour: isv_learning_health, isv_sharpe_ema, isv_q_means are
now properly reset to baseline at each fold boundary (previously unset,
which allowed signals from fold N to bias fold N+1 initialisation).

Tests: 3 registry unit tests pass; cargo check -p ml clean (8 pre-existing
warnings only, no new).

Authority: spec §4.A.1. Plan 1 Task 3.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-24 11:10:11 +02:00
jgrusewski
b688827d66 feat(dqn-v2): A.1 state reset registry
Typed classification of every piece of training state by reset
lifecycle: FoldReset, WindowReset, SoftReset(decay_bars),
TrainingPersist, SchemaContract.

Replaces the previous vibes-driven fold-boundary reset logic with an
explicit registry. Adding new state requires adding an entry in the
same commit (Invariant 2 Wire-It-Up).

Consumer wiring (reset_fold_state, reset_window_state helpers that
iterate the registry) follows in Task 3 of Plan 1.

Tests: 3 unit tests covering category lookup, fold-reset filtering,
unknown-name handling.

Authority: docs/superpowers/specs/2026-04-24-dqn-v2-unified-design.md §4.A.1

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-24 10:38:15 +02:00
jgrusewski
06989cfdf9 infra(dqn-v2): audit doc scaffolding + pre-commit enforcement
Plan 1 Task 1. Creates the five audit docs plus config/metric-bands.toml
that track Invariants 2, 7, 8 per the DQN v2 spec, and extends the
pre-commit hook with two checks:

  - component-adding commits must touch an audit doc (Invariant 7)
  - added code may not contain TODO/FIXME/XXX/HACK/TBD/unimplemented!/
    todo! markers (Invariant 9)

Tests: manually verified by staging a TODO-marked file; commit
rejected with the correct error message.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-24 10:25:27 +02:00
jgrusewski
3ab87d7672 plan(dqn-v2): Plan 3 self-review fix — remove unimplemented!() from example
Self-review caught that the CqlAlphaSeedCoupledController read_signals
example used unimplemented!("plumbed by caller"), which violates the very
Invariant 9 the plan enforces. Replaced with a concrete implementation
that reads SEED_FRACTION_EMA_INDEX from the ISV bus (a Plan 1 reserved
controller-signal slot).

No other plan had placeholder violations. The remaining TBD/TODO/FIXME
grep hits across Plans 1-5 are either the pre-commit hook's rejection
patterns (correctly quoted) or enforcement scans that DETECT the markers
in audit docs.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-24 10:07:10 +02:00
jgrusewski
74071af908 plan(dqn-v2): Plan 5 — validation substrate + final pass implementation plan
Fifth and final plan decomposing the DQN v2 unified spec
(docs/superpowers/specs/2026-04-24-dqn-v2-unified-design.md §2, §4.A.3,
§4.A.4, §4.A.4.1). No new policy mechanisms — orchestration, automation,
regression detection, and the final sign-off run.

Covers spec sections:
- §4.A.3 Multi-seed × multi-fold Argo harness (N=5, K=6 defaults via
  --multi-seed / --folds flags on argo-train.sh)
- §4.A.4 Regression detection hard-stop (2N consecutive error-band
  metrics → self-SIGTERM with TerminationReason)
- §4.A.4.1 nsys profile harness with --profile flag and
  compare-nsys-profiles.py for 20% per-epoch regression detection
- §2 Tier 1 + Tier 2 + Tier 3 final exit criteria via scripts/validation/
  suite (tiered tier1/tier2/tier3 check scripts + wrapper)
- §2 DQN v2 rollout close-out (spec marked LANDED, retrospective,
  plan-level validation docs archived)

Plan structure: 6 tasks — 4 infrastructure (harness, regression-detect,
nsys, validation-scripts) + final-run task + rollout-closeout task. The
final-run task is the sole exit-gate for the DQN v2 rollout: ALL THREE
tiers must PASS on the first final-validation run (per stop-on-anomaly:
retrying with different seeds hoping for a different outcome is anti-
pattern).

Preserves all 9 invariants. Every mechanism wired, no stubs, no TODO/FIXME.
Retrospective written from actual implementation experience after Plans
1-5 all land.

This is the terminal plan. Plans 1-5 together constitute the complete
DQN v2 rollout. After Plan 5's final-run exit passes, the spec is marked
LANDED and the rollout is declared complete.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-24 10:06:25 +02:00
jgrusewski
691d769bb3 plan(dqn-v2): Plan 4 — supervised→DQN concept adoption implementation plan
Fourth of five sequential plans decomposing the DQN v2 unified spec
(docs/superpowers/specs/2026-04-24-dqn-v2-unified-design.md §4.E).

Covers the six Part E IN decisions:
- §4.E.1 TFT Variable Selection Network across 6 feature groups
  (market/OFI/TLOB/MTF/portfolio/plan_isv) with vsn_feature_selection kernel
- §4.E.2 Gated Residual Network — CONDITIONAL branch (ADOPT vs CANONICALISE)
  based on docs/ml-supervised-to-dqn-concept-audit.md row
- §4.E.3 Multi-quantile IQN heads (5/25/50/75/95) replacing single CVaR output
- §4.E.4 Encoder-Decoder separation — explicit StateEncoder + per-branch
  ValueDecoder with D.3 horizon-decomposed V_short/V_long sub-heads
- §4.E.5 Attention-weight interpretability — 7 new ISV slots [65..72) for
  per-group attention focus EMAs + Mamba2 retention proxy
- §4.E.6 Multi-task auxiliary heads — next-bar return MSE + 5-bar regime CE,
  ISV-coupled aux-weight schedule (sharpe-reactive)

ISV_TOTAL_DIM seals at 72 with this plan's allocations. Part E audit doc
closes out all rows (zero TBD/evaluate remain per Invariant 9).

Plan structure: 8 tasks (6 feature tasks + audit close-out + validation).
TDD-disciplined steps. Task 2 documents the CONDITIONAL branch decision
pathway explicitly (ADOPT vs CANONICALISE Branch A/B structure).

Plan 4 exit gate: Tier 1 convergence RETAINED (no regression vs Plan 3
baseline) + aux heads produce measurable signal. Tier 2 + Tier 3 checked
by Plan 5.

Preserves all 9 invariants. Zero stubs. Zero TODO/FIXME. Every Part E item
either lands fully or is explicit OUT (xLSTM, KAN); no third path.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-24 10:03:13 +02:00
jgrusewski
c70bbd5d15 plan(dqn-v2): Plan 3 — behavioural core implementation plan
Third of five sequential plans decomposing the DQN v2 unified spec
(docs/superpowers/specs/2026-04-24-dqn-v2-unified-design.md).

Covers spec sections:
- §4.C.2 Reward-component attribution (6 ISV slots [52..58), reward_split[..]
  in HEALTH_DIAG) — lands first as diagnostic substrate for the rest
- §4.B.1 ISV-driven continuous Flat opp-cost (scales by isv[Q_DIR_ABS_REF])
- §4.B.2 Novelty-weighted trade-attempt bonus (2 ISV slots [50, 51])
- §4.B.4 Adaptive plan-MLP activation threshold (1 ISV slot [49] +
  PlanMlpThresholdController via the AdaptiveController trait)
- §4.C.4 Temporal timing bonus (trade-exit, bars_early / hold_time)
- §4.D.4 Generalised temporal reward coupling (3 ISV slots [59..62):
  persistence credit, regime-shift penalty, conviction consistency)
- §4.C.3 State-distribution divergence signal (3 ISV slots [58, 63, 64])
- §4.B.3 Replay buffer seeded warm-start (4 scripted policies,
  ≥100K experiences, CQL α decoupled from seed phase)
- §4.C.5 CQL alpha schedule coupled to seed-phase decay via
  CqlAlphaSeedCoupledController

Plan structure: 9 implementation tasks + pre-plan verification + exit-gate
validation task (Task 10). Each task TDD-disciplined with bite-sized
steps. Task 1 (C.2 diagnostic audit) ordered first so subsequent
reward-shape tasks can verify contributions rather than log-archaeology.

Plan 3 exit gate: multi-seed (N=5) × multi-fold (K=6) run passes Tier 1
(convergence) + Tier 2 behavioural subset (val trades/bar ≥ 0.005,
val_active_frac > 0.2). Tier 3 profitability deferred to Plan 5.

Preserves all 9 invariants. Zero stubs. Zero TODO/FIXME. Every new module
wired to its consumer in the same task it lands.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-24 09:59:32 +02:00
jgrusewski
a5101eb2f8 plan(dqn-v2): Plan 2 — temporal core implementation plan
Second of five sequential plans decomposing the DQN v2 unified spec
(docs/superpowers/specs/2026-04-24-dqn-v2-unified-design.md).

Covers spec sections:
- §4.C.1 Quantile-based atom support (8 new ISV slots, q_quantile_reduce kernel)
- §4.D.1 Mamba2 backward pipeline completion (no atomicAdd — per-sample arrays + host reduce)
- §4.D.2 Per-branch gamma via AdaptiveController (4 new ISV slots)
- §4.D.5 Soft fold-boundary transitions (extends StateResetRegistry with SoftReset category)
- §4.D.7 Liquid Time-constant audit (trace fire rate, decide wire-or-delete)
- §4.D.3 + §4.D.6 + §4.D.8 coordinated state-layout migration (horizon-decomposed V +
  plan_isv[6] + TLOB integration with atomic ISV schema version bump 1→2)

Plan structure: 7 tasks + pre-plan verification, all TDD-disciplined with
bite-sized steps (test → fail → implement → pass → commit). Task 6 is
explicitly atomic to preserve Invariant 5 (state-layout consistency).

Plan 2 prerequisites: Plan 1 landed (StateResetRegistry, AdaptiveController
trait, ISV schema version, audit docs). Plan 2 exit: all 9 invariants
preserved; convergence-scaffolding run passes with 16 new ISV slots
populated; ISV schema version == 2.

Preserves all 9 invariants. No stubs. No TODO/FIXME. Every new module
wired to production path in the same task it lands in.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-24 09:53:11 +02:00
jgrusewski
e5f6afca8d plan(dqn-v2): Plan 1 — substrate & refactor implementation plan
First of 5 implementation plans for the DQN v2 unified integrated
policy system spec (d13b53586, 336ee40b9). Plan 1 covers:

  Task 1: Audit doc scaffolding + pre-commit hook (Invariant 7, 9)
  Task 2: StateResetRegistry definition (A.1)
  Task 3: Wire StateResetRegistry into fold-boundary reset (A.1)
  Task 4: Named-dimension refactor — ps[], plan_isv[], plan_params[],
          branch indices, dir/mag sub-indices (Invariant 8)
  Task 5: ISV schema version at ISV[0], fail-fast on checkpoint
          mismatch (A.2)
  Task 6: Orphan audit populated (A.5)
  Task 7: Hot-path purity audit populated; MIGRATE calls fixed (A.6)
  Task 8: AdaptiveController trait + harness (C.6)
  Tasks 9–17: Migrate each adaptive controller to the trait in spec-
            specified order (atoms → gamma → Kelly → cql_alpha → tau
            → epsilon → conviction_floor → plan_threshold → balancer)
  Task 18: Plan 1 validation run (smoke + 3-epoch L40S)

Plan 1 landing criteria: all 9 invariants preserved across every
commit, all smoke tests pass, audit docs fully populated. Plan 2
(temporal core) starts only after Plan 1 passes exit criteria.

Plans 2–5 cover Parts B, C (minus C.6 already in Plan 1), D, E, and
final validation. Planned as separate documents in docs/superpowers/
plans/2026-04-24-dqn-v2-plan-{2..5}-*.md.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-24 09:41:22 +02:00
jgrusewski
336ee40b9d spec(dqn): add Invariant 9 — no deferred work, no stubs, concrete E decisions
User-mandated addition:

Invariant 9 — No deferred work, no stubs, no TODO/FIXME
Every commit lands complete. No stubs (placeholder return values), no
TODO/FIXME/XXX/HACK/TBD markers, no half-finished implementations. If
it can't finish in this commit, it doesn't start. Stubs and deferrals
train the network on semantic emptiness — invisible in convergence
metrics but burn GPU time optimising against partly-fake signal.
Authority: feedback_no_stubs.md, feedback_no_todo_fixme.md,
feedback_no_quickfixes.md elevated to first-class spec enforcement.
Pre-commit hook greps for forbidden markers.

Part E decisions made concrete (per Invariant 9):
- E.1 TFT VSN: IN (full VSN extension)
- E.2 GRN: IN if A.5 audit finds absent (otherwise mark existing as canonical)
- E.3 Multi-quantile heads: IN (5/25/50/75/95 quantile decomposition)
- E.4 Encoder-decoder: IN (extends D.3 to full trunk/head separation)
- E.5 Interpretability: IN (attention-weight ISV diagnostics)
- E.6 Auxiliary heads: IN (next-return + regime-classification)
- xLSTM: OUT (redundant with Mamba2+TLOB; YAGNI)
- KAN: OUT (function approx not a bottleneck)
No "evaluate later" items. Every concept either lands or explicitly
doesn't, with rationale.

§8 decisions tightened:
- D.8 TLOB mode: decision criterion = per-step cost benchmark ≤ 5ms
- C.6 controller order: fixed in spec (atoms→gamma→Kelly→cql_alpha→
  tau→epsilon→conviction_floor→plan_threshold→balancer)
- A.3 resource plan: auto-switch L40S→H100 at 12 GPU-hour threshold
- A.4 metric band init: [mean − 3σ, mean + 3σ] from last 3 good runs

No discretionary "we'll see" remain. All decisions data/order/budget/
statistic-driven.

Counter updates: "seven invariants" → "nine" in §3 header and §5
cleanup contract.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-24 09:32:22 +02:00
jgrusewski
3a62e23056 spec(dqn): review fixes — clarify hot-path scope, add Invariant 8 named dims
Self-review pass on the DQN v2 unified design spec (d13b53586). Fixes
applied inline:

Ambiguity fixes:
- §3 Invariant 3: explicit definition of "training step loop" (per-step
  code in fused_training::step_fused and captured graph children; NOT
  per-epoch / per-fold / per-checkpoint code).
- §2 Tier 1: "stable epochs" defined as epochs [warmup_end..run_end)
  with warmup_end from B.3's seed-phase decay or explicit flag.
- §4 A.2: migration mechanism clarified as fail-fast initially; helpers
  added only when a specific migration need arises (no speculative
  scaffolding).
- §4 A.3: N=5, K=6 stated as DEFAULTS with ≥ MINIMUMS per Invariant 6;
  resource budget flagged (~8-10 GPU-hours per validation pass).
- §4 B.3: "≥ 100K experiences" (minimum), warm-restart clarified as
  runtime condition not feature flag.
- §4 C.3: adaptive KL threshold mechanism made explicit (second ISV
  slot for threshold EMA, third for amplification multiplier).
- §4 C.6: cleanup timing explicit — old scaffolding removed in the SAME
  commit that migrates its last consumer (no deferred pass).
- §4 D.1: DQN-path-specific scope clarified; validation via grad-norm
  smoke + integration test.
- §4 D.6: plan_isv index naming made consistent with existing layout;
  coordinated state-layout migration commit called out.

New Invariant 8 — Named dimensions, not indices:
Every dimension, slot, offset, or semantic position in a multi-field
buffer has a named constant. Raw numeric indices (`[0]`, `[6]`, `[23]`)
appear ONLY in the definition site. Named constants specified for ISV
slots (existing), portfolio state ps[0..30), plan_isv[0..7), plan_params
[0..6), state-vector offsets, branch indices (BRANCH_DIR/MAG/ORD/URG),
action sub-indices (DIR_SHORT/HOLD/LONG/FLAT, MAG_QUARTER/HALF/FULL).
Enforcement: audit pass during A.1/A.2, lint via grep for raw-index
access outside definition modules.

Landing order revision:
- Dependency graph explicit (A.2 before new ISV slots, D.1 before
  D.2-4-8, state-layout changes in ONE coordinated commit).
- Spec decomposition note added — writing-plans may produce multiple
  sequential plans rather than one monolithic plan.

New doc tracked by pre-commit: docs/dqn-named-dims.md (Invariant 8).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-24 09:28:45 +02:00
jgrusewski
d13b535862 spec(dqn): DQN v2 unified integrated policy system design
Design spec consolidating every hard lesson from the val-Flat-collapse
investigation, the task #92 gradient-pathology triad, the ISV v-range
unification work, the f64→f32 ABI cleanup, and many sessions of
fold-1 grad explosions and orphan-feature accumulation.

Seven non-negotiable invariants:
  1. ISV-driven adaptive bounds (no tuned constants)
  2. Wire-It-Up (no orphans)
  3. GPU-only hot path (zero memcpy DtoH; pinned-zero-copy only)
  4. Pinned-only host memory
  5. Diagnostic-first
  6. Convergence discipline (primary guardrail, multi-seed × multi-fold)
  7. Wire-up + diagnostic audit per-commit

Five Parts, landing as one coordinated rollout:
  A. Foundation — reset registry, ISV contract, multi-seed harness,
     regression detection, orphan audit, hot-path purity audit.
  B. Flat-trap escape — ISV-driven reward shaping, trade-attempt bonus,
     replay warm-start via scripted-policy portfolio, adaptive plan
     threshold.
  C. Refinement — quantile atom support, reward attribution, state-dist
     divergence signal, temporal timing bonus, CQL-seed coupling,
     adaptive controller unification refactor.
  D. Temporal architecture — Mamba2 backward completion, per-branch
     gamma, horizon-decomposed value function, generalised temporal
     reward, soft fold transitions, plan temporal dynamics, liquid
     audit, TLOB-DQN integration.
  E. Supervised→DQN concept audit — TFT VSN, GRN, multi-quantile,
     xLSTM, KAN, encoder-decoder, interpretability, auxiliary heads.

Tiered success criteria:
  Tier 1 — convergence discipline (must pass first).
  Tier 2 — behavioural parity (val trades escape Flat-trap).
  Tier 3 — profitability (val Sharpe > 1.0 on window).

ISV slots grow from 37 to 72. Five new audit docs tracked by
pre-commit hook.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-24 09:22:43 +02:00
jgrusewski
d64adc14f5 refactor(dqn): f64 → f32 for kernel-facing hyperparams
Eliminates the f64→f32 cudarc ABI trap (feedback_cudarc_f64_f32_abi.md,
task #82) at the type level: hyperparameters consumed by CUDA kernels
now live as f32 in Rust, cast once at the TOML/PSO ingest boundary
instead of at every kernel call site.

Structs changed:
  - DQNHyperparameters (crates/ml/src/trainers/dqn/config.rs) —
    ~85 scalar fields migrated from f64 → f32. Covers all
    kernel-facing scalars: reward weights (w_pnl/w_dd/w_idle,
    dd_threshold, cea_weight, micro_reward_*, price_confirm_weight,
    book_aggression_weight, hold_quality_weight), exploration
    (epsilon_* and the 4 branch mults, noisy_sigma_*, count_bonus,
    noise_sigma, q_gap_threshold), distributional RL (v_min, v_max,
    reward_scale, iqn_lambda, qr_kappa, spectral_*,
    gradient_collapse_multiplier), fill simulation (5 fill_*
    fields), risk/Kelly (kelly_fractional, kelly_max_fraction,
    max_leverage, max_position_absolute, minimum_profit_factor),
    ensemble/curiosity (curiosity_weight,
    curiosity_q_penalty_lambda, ensemble_*, beta_*, variance_cap),
    anti-LR (anti_lr_*, adversarial_dd_threshold,
    beta_penalty_strength), walk-forward (wf_*),
    experience (avg_spread, transaction_cost_multiplier,
    holding_cost_rate, churn_penalty_scale, contract_multiplier,
    margin_pct, tick_size, bars_per_day, cash_reserve_percent),
    misc kernel scalars (gamma, tau, huber_delta, q_clip_*,
    shrink_perturb_*, regime_replay_decay, per_alpha,
    per_beta_start, dt_target_return, etc.). Also
    `noisy_epsilon_floor: Option<f32>` and
    `count_bonus_coefficient: Option<f32>`.
  - `computed_v_min` / `computed_v_max` now return f32.
  - `compute_max_position` returns f32 (f64 internally for the
    notional division).

Fields preserved as f64 (precision-sensitive, NOT kernel-facing
scalars — per task spec and feedback_cudarc_f64_f32_abi.md):
  - `learning_rate` — tested at 1e-10 tolerance; f32 rounds to
    2e-12 for a 1e-5 LR.
  - `entropy_coefficient` — tested at 1e-9 tolerance.
  - `weight_decay` — tiny 1e-5..1e-3 range.
  - `adam_epsilon` — 1e-8 default; f32 preserves denorms here but
    paired with weight_decay/learning_rate for symmetry.
  - `gradient_clip_norm: Option<f64>` — grad norms are f64
    accumulators by project convention.
  - `min_loss_improvement_pct`, `q_value_floor` — early-stopping
    long-horizon stats.
  - `cql_alpha` — flows into DQNConfig (ml-dqn) still f64.
  - `min_learning_rate`, `lr_min` — paired with learning_rate.
  - Family intensity scalars (6 `*_intensity` fields) — f64 PSO
    search space; the intensity applies via `as f32` at each call
    site in `apply_family_scaling`.

No checkpoint format change: DQNHyperparameters has
`#[derive(Debug, Clone)]` only (not Serialize/Deserialize), so the
TOML-ingest path is the only serde boundary and already casts
explicitly via `hp.field = v as f32;` in
`DqnTrainingProfile::apply_to`. PSO hyperopt bounds stay f64 in
`SearchSpaceSection` and cast at the adapter boundary.

Call-site impact:
  - ~30 `as f32` casts removed from hot paths (fused_training
    FusedConfig builder, training_loop kernel launches,
    constructor DQNConfig builder, action.rs GPU action selector,
    trainer/mod.rs WF config). Kernels now receive the hp field
    directly via `&hp.x`.
  - ~75 `as f32` casts added at the `apply_to` / hyperopt-adapter
    ingest boundary — the single conversion point.
  - Cross-crate contracts (DQNConfig in ml-dqn, PortfolioTracker
    in ml-core, GAECalculator, DropoutScheduler, NoisySigmaScheduler,
    KellyOptimizerConfig, RewardConfig) retain their f64 signatures;
    ml calls cast at the boundary with `f64::from(hp.x)` so the
    contract is explicit and greppable.

Two test-side adjustments:
  - `test_kelly_fields_are_public` now asserts `f32` for
    kelly_fractional / kelly_max_fraction (these migrated).
  - `test_early_stopping_termination` uses `f64::from(...)` to
    preserve the f64 threshold computation against the now-f32
    `gradient_collapse_multiplier`.

Verified:
  - `SQLX_OFFLINE=true CARGO_INCREMENTAL=0 RUSTC_WRAPPER=sccache
    cargo check --workspace` — clean, zero new warnings.
  - `cargo check --workspace --tests` — clean.
  - `cargo test -p ml --lib training_profile::tests` — all 18
    migrated tests pass. The one pre-existing failure
    (`test_production_profile_applies_all_sections` n_steps
    mismatch, 1 vs 5) reproduces on stashed baseline, so it is
    unrelated to this change.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-24 01:55:14 +02:00
jgrusewski
b12483d91b feat(dqn): val plan_isv parity — hot-path integration
Final commit toward task #94 val plan_isv parity. Wires the plan
machinery inside the chunked val backtest loop so state positions
[86..92) carry real plan signals matching training instead of zero-
fill. This is the structural distribution-shift fix behind the
observed val trade count of 21 / 214,654 bars (0.0098%).

Phase 6 inside `submit_dqn_step_loop_cublas`, per chunk, after env
step advances portfolio:

  1. Forward on the LAST-STEP N rows of chunked_states — a targeted
     `compute_q_values_to(last_step_states, N, scratch_q)` rewrites
     `save_h_s2` to exactly [N, SH2] for the final step. Needed
     because the batched `compute_q_values_to` for N*chunk_len rows
     leaves save_h_s2 at an arbitrary last-sub-batch slice.
  2. `compute_plan_params(plan_params_buf, N)` runs the trade-plan
     MLP on that clean save_h_s2, producing [N, 6] plan_params.
  3. `backtest_plan_state_isv` — Flat↔Positioned activation /
     deactivation on plan_state[N, 7] and writes plan_isv_buf[N, 6]
     consumed by the next chunk's first `launch_gather` for state
     positions [86..92).

ch_q_values is reused as the Q-scratch for step (1) — its original
chunk Q-values were already consumed by Phase 4 action-select.

Epoch-boundary diagnostic `launch_plan_diag_and_log` after the step
loop: launches `backtest_plan_diag_reduce` (1-block, 256 threads),
DtoH-copies the 8-float summary, syncs, and emits a single
`tracing::info!(target: "val_plan_diag", ...)` line with the six
labelled plan_isv means plus active_frac / n_active.

Scalars passed with exact i32 types (current_step, max_len,
feat_dim, n_windows) per feedback_cudarc_f64_f32_abi. Kernels
pulled via `plan_state_isv_kernel.as_ref().ok_or_else(...)` — they
are loaded in `ensure_action_select_ready` on first eval.

Not captured inside a CUDA Graph: the backtest step loop uses direct
kernel launches (evaluator comment: graphs SIGSEGV on 500K+ launches).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-24 01:20:30 +02:00
jgrusewski
39a353495d feat(dqn): val plan_isv parity — evaluator buffer + kernel loading
Second commit toward task #94 val plan_isv parity. Adds structural
wiring in GpuBacktestEvaluator:

  - plan_params_buf [N, 6]  : live plan MLP output per step
  - plan_state_buf [N, 7]   : persistent stored plan across bars
  - plan_diag_buf  [8]      : epoch-end diagnostic reduction output
  - plan_state_isv_kernel   : lazy-loaded CudaFunction from
                              backtest_plan_kernel.cubin
  - plan_diag_reduce_kernel : diagnostic reduction kernel

Buffers allocated zero-initialised in the constructor; kernels loaded
alongside the action_select and scatter_intent kernels in
`ensure_action_select_ready` (same module-load pattern).

Remaining work (follow-up commit):
  - Wire QValueProvider::compute_plan_params into the chunked loop to
    populate plan_params_buf from the trainer's save_h_s2 after each
    forward pass.
  - Launch backtest_plan_state_isv after the env_batch_kernel each
    chunk to update plan_state_buf and plan_isv_buf for the next
    chunk's state gather.
  - Launch backtest_plan_diag_reduce + DtoH of plan_diag_buf at epoch
    boundary for HEALTH_DIAG emission.

The chunked batch layout (N*chunk_len samples per forward pass)
requires careful slicing to extract the last-step's plan_params for
the plan_state update — handled in the follow-up commit to avoid
rushing a subtle integration.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-24 01:13:17 +02:00
jgrusewski
29ac4180ab feat(dqn): val plan_isv parity — kernel + provider trait foundation
Val-Flat-collapse fix #5 infrastructure (task #94, 2026-04-24). The
val backtest's `plan_isv_buf` was zero-filled, causing state positions
[86..92) to be OOD relative to training and biasing val argmax toward
Flat/Hold on 99.99% of bars. Research-agent audit confirmed this is a
documented architectural gap, not a bug per se: val's portfolio state
has no plan slots (8 vs training's 30+), so there was no mechanism to
carry plan signals across bars in val.

This commit lays the foundation for closing that gap:

1. `backtest_plan_kernel.cu` (new) — two kernels:
   - `backtest_plan_state_isv`: per-window Flat→Positioned plan
     activation + plan deactivation on return to Flat + plan_isv[0..5]
     computation matching training's formula at
     experience_kernels.cu:596-619. Reads live plan_params from the
     plan MLP, writes plan_state (persistent across bars) and
     plan_isv_out. Extracts raw_close from features buffer for
     unrealized P&L ratios.
   - `backtest_plan_diag_reduce`: single-block reduction emitting 8
     diagnostic floats for plan activity logging (active fraction,
     per-slot mean magnitudes, raw active count).

2. `QValueProvider::compute_plan_params` trait method — allows the val
   evaluator to run the plan MLP on the trainer's most recent trunk
   hidden state (save_h_s2), filling plan_params[N, 6] for use in the
   state-update kernel.

3. `FusedTrainingQValueProvider::compute_plan_params` impl — runs
   `launch_trade_plan_forward` then DtoD-copies `plan_params_buf` to
   the caller's output in the same chunk-loop pattern as
   `compute_q_values_to`.

4. build.rs — register `backtest_plan_kernel.cu` for compilation.

Next commit will wire these into `GpuBacktestEvaluator`:
  - Add `plan_params_buf[N, 6]`, `plan_state_buf[N, 7]`,
    `plan_diag_buf[8]` device allocations.
  - Load the two kernels into CudaFunction slots.
  - Call sequence per step: `compute_q_values_to` →
    `compute_plan_params` → action_select → env_step →
    `backtest_plan_state_isv` (reads updated portfolio, writes
    plan_state + plan_isv_buf for next step).
  - Epoch boundary: `backtest_plan_diag_reduce` + DtoH + HEALTH_DIAG
    emission.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-24 01:10:00 +02:00
jgrusewski
33376525ac fix(dqn): Flat opportunity cost scales with ISV-driven conviction, not tuned constant
Val-Flat-collapse fix #3 revision (task #94, 2026-04-24). Prior
commit 543e3c11b used `-0.5 × holding_cost_rate × vol_proxy` — the
0.5 is a hardcoded tuned constant violating
`feedback_isv_for_adaptive_bounds.md` and
`feedback_adaptive_not_tuned.md`.

Replace with ISV-driven per-sample conviction, already computed by
the action-select kernel and threaded into env_step via
`conviction_ptr → conviction_core`:

    reward_flat = -shaping_scale * holding_cost_rate
                * conviction_core * vol_proxy_flat

`conviction_core ∈ [0, 1]` = direction-branch Q-range normalised by
`ISV[21]` (q_dir_abs_ref EMA). Self-adapting properties:

  - Cold start / ISV[21] uninitialised → fallback 1.0 → full penalty,
    encourages early exploration out of the flat equilibrium.
  - Low conviction (uncertain direction) → penalty scales toward 0
    → doing nothing is acceptable when there's no signal (matches
    real-world: flat cost is only real when there's opportunity).
  - High conviction (strong directional edge) → penalty scales up
    → Flat becomes expensive ONLY where the model itself says
    there's an edge. Forces the policy to take action exactly
    where it has belief, not blindly.

Continuity: conviction is continuous ∈ [0, 1], no step function. The
temporal Mamba2 layers in the trunk feed the Q-values that drive
conviction, so conviction inherits temporal history — the model can
learn "market has been signalling for N bars, time to try" without
an explicit time-since-last-trade feature.

Training-time exploration (Boltzmann sampling, epsilon floor 2%,
NoisyNets σ, count bonus) is unchanged. The Flat-cost shifts the
learned Q-values so that deterministic val argmax picks trade-
actions where the training-time exploration already found edge.

Hold semantics retained:
  - Hold while position≠0 (in-trade stance) → positioned-bar
    holding-cost branch (unchanged)
  - Hold while position=0 AND Flat (no-op outcomes) → this branch
    → conviction-scaled opportunity cost.
The distinction is structural by portfolio state, not by dir label.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-24 00:54:32 +02:00
jgrusewski
29e54a5bb9 fix(dqn): invert adaptive gamma direction on low atom utilization
Val-Flat-collapse fix #4 (task #94, 2026-04-24). Research-agent audit
of train-bv2n5 identified a positive-feedback loop driving atom
utilisation to 1%:

    util < 0.4 → γ -= 0.01 (toward 0.90 floor)
    γ × delta_z shrinks → TD targets concentrate on fewer bins
    → more collapse → util lower → γ lower ...

The original controller had the right structure but the wrong
direction on the low-util branch. Flipping the sign (`-=` → `+=`)
breaks the loop: when atoms collapse, raising γ widens the effective
TD-target support `γ × atom_support`, pushing targets across more
bins of the C51 grid and giving the network more distributional
signal to learn from.

Evidence from train-bv2n5 (pre-fix): atom_util stuck at 0.01 for 16
epochs; γ slowly drifted toward 0.90. No mechanism recovered.

Change localised to a single `else if util < 0.4` branch at
`training_loop.rs:660-675`. Same step size (0.005) as the healthy
branch avoids overshoot on the recovery path. Clamps retained
[0.90, 0.95] at this layer; the fused trainer re-clamps to
[0.90, 0.995] after regime adjustment downstream.

Other pending atom-util fixes from the audit (deferred pending
validation):
- Atom-entropy regularisation term in C51 loss (medium risk)
- Absolute v_half floor in update_eval_v_range (low risk)
- TD-target dithering before Bellman projection (low risk)
This single-line fix addresses the dominant mechanism — research
agent labelled it "lowest risk, highest leverage". If util remains
low after this, the others stack additively.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-24 00:26:39 +02:00
jgrusewski
543e3c11b9 fix(dqn): Flat opportunity cost breaks val Flat-equilibrium
Val-Flat-collapse fix #3 (task #94, 2026-04-24). Research agent
traced the dominant val failure mode: reward = 0.0 for Flat bars
produces structural Flat-beats-trading-actions in the argmax —
Flat accumulates Q ≈ 0 while CQL pulls trading-action Q's toward
pessimistic negative estimates. Result: val argmax picks Flat for
~all states (22 trades/epoch vs 30K+ in training via Boltzmann).

Prior comment defended the design: "No signal is correct — the
model should not be rewarded or penalized for correctly staying
flat when there's no edge." That reasoning was correct in isolation
but broken in context: the trading-action branch is penalized by
CQL's lower-bound estimator anyway, so Flat's neutral Q becomes the
de-facto upper bound on all action Q-values → Flat wins argmax
regardless of the policy's actual directional belief.

Add a per-bar opportunity cost for Flat bars, proportional to
realised volatility (ATR fraction):

  reward = -shaping_scale * holding_cost_rate * 0.5 * vol_proxy

This:
- Symmetric to the positioned-bar holding penalty (positioned
  pays `holding × |pos|`; Flat pays `holding × 0.5 × vol_proxy`).
- Vol-scaled: Flat during quiet markets approaches zero cost;
  Flat during volatile markets has a small negative signal
  representing missed-move expectation.
- Caps vol_proxy at 0.01 (1%) as numerical safety against
  single-bar vol spikes dominating reward.
- Recomputes the ATR locally (the segment-complete branch that
  normally does this is mutually exclusive with this Flat branch).
- Guarded on `features != NULL` and `bar_idx` range for smoke-test
  compatibility.

The `0.5` factor is a numerical-safety symmetric-to-holding-cost
bound per the carve-out in `feedback_isv_for_adaptive_bounds.md`,
not a tuned hyperparameter. The `holding_cost_rate` itself flows
from config (already adaptive via shaping_scale annealing).

Not part of the triad yet:
- plan_isv train/val parity (task #94 item #2): documented
  architectural gap requiring val-time plan-state generator; no
  trivial safe fix.
- Atom utilization (task #94 item #4): deep C51 representation
  work (atom-level entropy regularisation or target-variance
  engineering); deferred pending tie-break + opp-cost validation.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-24 00:08:55 +02:00
jgrusewski
97a6c55805 fix(dqn): adaptive tie-break epsilon in eval argmax via ISV half-width
Val-Flat-collapse fix #1 (task #94, 2026-04-24). Research agent audit
of train-bv2n5 identified that val picks Flat for ~all states while
training picks diverse directions via Boltzmann sampling. Evidence:
val produced 16-27 trades per epoch while training produced 30K-200K
with the same policy. val_Sharpe sat at 0.00 for 16 consecutive
epochs (raw magnitudes 1e-6 to 3e-5 — noise floor).

Root cause: C51 atom utilization collapsed to ~1% (all probability
mass on a handful of atoms). With v_range≈240 and 50 atoms, each
atom spans ~9.6 units; direction E[Q] values drifted by 1e-4 to
1e-5 from floating-point precision alone — below atom resolution
but above the existing `1e-6f` tie-break threshold. The strict
argmax then gave one bin (typically Flat via initialization /
accumulated-drift bias) a consistent-but-meaningless edge.

Fix: widen the tie-break epsilon adaptively using the per-branch
Q-support half-width stored in the ISV bus (V_HALF_{DIR,MAG,ORD,URG}
at indices 24, 26, 28, 30). `tie_eps = max(1e-6, 0.01 × isv_half)` —
Q-values within 1% of the learned support are numerically
indistinguishable from atom jitter and fire the existing Philox
uniform-among-tied sampler. NULL ISV falls back to 1e-6 (pre-fix
behaviour) for smoke-test callers that don't wire the bus.

Applied to all four branches in `experience_action_select`:
- direction: half_dir from isv[24]
- magnitude: half_mag from isv[26]
- order:     half_ord from isv[28]
- urgency:   half_urg from isv[30]

The 1% fraction is a numerical-safety bound per
`feedback_isv_for_adaptive_bounds.md` (safety carve-out: "Safety
hard limits can stay as constants; adaptive training-dynamics bounds
cannot"). The bound IS a safety floor; the adaptation flows through
isv_half which is itself an EMA-driven adaptive signal.

Complementary fixes still pending (task #94 remainder):
- opportunity_cost in training reward (structural Flat-Q-≈-0 fix)
- atom utilization / v-range improvements (deep C51 fix)
The plan_isv train/val gap is a known documented trade-off
(gpu_backtest_evaluator.rs:524-535) deferred pending infrastructure
to compute plan signals at val time.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-24 00:03:43 +02:00
jgrusewski
18da01bf23 fix(dqn): ISV-driven per-sample floor on IQL branch_scales
Root-cause bug #3 from the C-audit (task #92). The IQL advantage-weighting
kernel `iql_per_branch_advantage` computed:

    iql_scale   = a_branch[d] / max_a          // per-branch advantage / max
    branch_scales[b, d] = r * iql_scale + (1-r) * 0.25

When IQL readiness `r` → 1 and a sample's branch advantage `a_branch[d]`
is near zero, the scale collapsed to ~0. That zero multiplies the C51
per-sample gradient to zero, starving (b, d) of learning signal. For
direction branch this happens disproportionately on Hold (d==1 action)
and Flat (d==3 action) samples where the direction advantage is
structurally small. Combined with the ~22% Hold+Flat sample fraction,
direction accumulates gradient from only a subset of the batch — part of
the compound 150x grad_ratio_mag_dir observed in train-5wb4n.

Add an ISV-driven floor:
  - New ISV slot (36): IQL_BRANCH_SCALE_FLOOR_INDEX
  - ISV_TOTAL_DIM 36 → 37
  - Bootstrap at 0.1 (safety bound, 10% of uniform=0.25 — per-sample
    scale can never fall below this)
  - `iql_per_branch_advantage` kernel takes ISV pointer + slot index,
    reads floor, takes `max(scale, floor)` before writing branch_scales
  - `compute_branch_scales` in `gpu_iql_trainer.rs` threads ISV ptr +
    slot index through (mirrors existing `compute_per_sample_support`
    pattern)
  - `fused_training.rs` step 7 passes `self.trainer.isv_signals_dev_ptr()`
    + `IQL_BRANCH_SCALE_FLOOR_INDEX` at the call site

The floor at 0.1 is a numerical-safety bound per the carve-out in
`feedback_isv_for_adaptive_bounds.md` ("Safety hard limits can stay as
constants"). It's addressable via ISV for future adaptive tuning if
observed starvation patterns warrant, but is not actively EMA-updated
today.

Completes the three-part fix triad for task #92:
  1. d61aefe2b — ISV-driven grad balancer (safety net, per-branch norm)
  2. dc6a034a4 — C51 backward standardization (math bug, per-atom)
  3. this commit — IQL branch_scales floor (per-sample starvation)

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-23 23:23:01 +02:00
jgrusewski
dc6a034a4c fix(dqn): C51 backward matches forward advantage-standardization for d==1
Root-cause bug from C-audit of train-5wb4n magnitude-branch Q-saturation
(task #92). Forward kernel `c51_loss_batched` (c51_loss_kernel.cu
lines 736-752) applies per-atom advantage standardization for the
magnitude branch:

    centered[j] = (adv[a_d,j] - a_mean[j]) / (a_std[j] + 1e-6)   // d==1 only
    combined[j] = value[j] + centered[j]
    lp[j]       = log_softmax(combined)

But the backward kernel `c51_grad_kernel` treated the forward as if
`combined[j] = value[j] + (adv[a_d,j] - a_mean[j])` — no a_std divisor.
That is chain-rule-inconsistent: the computed `d_combined / d_adv[a,j]`
is the gradient of a *different* loss than the one forward-pass'd.
Gradient points in the wrong direction relative to the actual loss
surface.

Fix: pass the magnitude-branch online advantage logits pointer
(`on_adv_logits_b1`) to the backward kernel. For each (b, j) in the
d==1 dueling-grad block, recompute a_std from the 3 advantage values
and multiply `grad_val *= 1/(a_std + 1e-6)`. Counterfactual magnitude
gradient (Hold samples) inherits the same factor automatically —
single correction point.

Approximation: treats a_std as constant w.r.t. advantage values (skips
`da_std/d_adv` chain-rule terms). This is the standard simplification
used in batch-norm with `track_running_stats=False`; produces the
dominant scale correction without the full Jacobian complexity.

Call site updates the single launch in `launch_c51_grad`. Kernel
signature guards on NULL for backward compatibility with any smoke-
test launcher that may not wire the argument; d==1 block falls back
to pre-fix behaviour (identity) in that case.

Part of task #92 fix triad. Previous commit d61aefe2b addressed the
balancer (safety-net fix). Next: IQL branch_scales floor for
direction-branch Hold/Flat per-sample starvation.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-23 23:19:11 +02:00
jgrusewski
d61aefe2b8 fix(dqn): ISV-driven per-branch grad balancer — equalise via median target
The pre-spec balancer used `cap = 4 × median` to suppress outlier
branches. On L40S train-5wb4n (epoch 9 norms `[dir=130, mag=20000,
ord=18000, urg=16000]`), that made cap = 68000 — exceeding every
branch's norm. No cap ever fired; the kernel was a geometric no-op
for 11 consecutive epochs while direction starved at 1/150 the
magnitude gradient. The outlier-cap algorithm can only handle "one
branch above the pack", not "three co-elevated, one starved".

Replace with ISV-driven equalisation:
  - 4 new ISV slots (31..34): per-branch grad-norm target
  - 1 new ISV slot (35): scale clamp limit
  - ISV_TOTAL_DIM 31 → 36
  - New on-GPU producer kernel `grad_balance_isv_update` runs after
    `branch_grad_norm_reduce`, before `branch_grad_rescale`. Computes
    current-step median + max/min spread, applies adaptive-rate EMA
    to the ISV slots. Single-threaded (1 warp), graph-capture safe,
    no atomicAdd, no host syncs.
  - Rescale kernel now reads ISV target + limit, computes per-branch
    `scale[b] = clamp(target[b] / norm[b], 1/limit, limit)`. Every
    branch equalises toward the shared median target over time;
    starved branches boost, elevated branches suppress. Limit is the
    observed max/min spread EMA, hard-bounded [2.0, 1e4] as a
    numerical-safety cap (not tuning).

Adaptive-rate EMA: `alpha = clamp(err / (err + baseline), 0.01, 0.3)`.
Fast response when error is large relative to current value, slow
drift when stable. No fixed alpha constant.

Bootstrap at construction + fold reset: targets = 1.0, limit = 2.0.
Pre-first-EMA-update behaviour: scale clamped to [0.5, 2.0], effectively
the original identity on warm-up, until real observations flow in.

Complies with feedback_isv_for_adaptive_bounds.md (2026-04-23):
adaptive bounds always live in ISV, never hardcoded. No partial
refactor — all consumers of the balancer contract (launch, kernels,
rescale algorithm) migrate together.

Still pending from the task-#92 triad: (a) C51 backward kernel must
replicate the forward advantage-standardization for d==1 (chain-rule
violation produces wrong gradient; separate commit), (b) IQL
branch_scales floor for direction-branch Hold/Flat starvation
(separate commit). This commit is the safety net; the two root-cause
fixes come next.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-23 23:12:53 +02:00
jgrusewski
69c663d360 infra(argo): disable CARGO_INCREMENTAL in ensure-binary for sccache hit rate
The project's Cargo.toml and .cargo/config.toml both set
`incremental = true`, which makes rustc emit per-query save-analysis
artifacts that sccache does not cache. Result: even though
RUSTC_WRAPPER=sccache was set, the rustc calls wrote incremental
state that subsequent builds saw as stale and recompiled anyway.

Setting CARGO_INCREMENTAL=0 in the ensure-binary env forces rustc
to emit pure object output that sccache hashes uniformly on
(source + args). Same SHA → full cache hit; small source change →
only dirty crates recompile.

Does not change the service compile-and-deploy-template path,
which intentionally uses cargo-incremental against a persistent
/cargo-target-cpu PVC for repeated rebuilds of a narrow set of
service binaries.

Applies from the next workflow submission onward (train-5wb4n
already compiled under the old config and is running).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-23 22:38:12 +02:00
jgrusewski
15818dce01 fix(dqn): break cold-start q_std latch in update_eval_v_range
Root cause of Q-value saturation at +/-50 seen in train-6nbx5 after
ISV v-range unification (9deda5f65, 11df03785): cold-start path in
`update_eval_v_range` latched `q_std_ema = q_std.max(0.01)` on the
first epoch. But that first `q_std` is dominated by the +/-50
bootstrap atom-spread (scaffolding set in construct/reset), NOT by
real Q-distribution spread. Result: `3*std_ema` exceeds
`min_half_floor=10` and approaches `abs_half=50` immediately, atoms
stay wide next epoch, next `q_std` confirms that width, EMA never
escapes. Q saturated at +/-abs_half every run.

Fix 1 (gpu_dqn_trainer.rs:3106-3120): seed
`eval_q_std_ema = min_half_floor / 3.0` at cold start so initial
`half = 3 * std_ema = min_half_floor` exactly. Adaptive-rate EMA
(alpha clamped to [0.01, 0.3]) then relaxes upward only if genuine
Q-spread warrants it. Breaks the self-confirming initialization.

Fix 2 (training_loop.rs:479-494): remove leftover pre-clamp of
reward quantiles to `config.v_{min,max}`. That was from the earlier
quantile-clamp fix (d38a8cf99). Phase 2c (9deda5f65) moved the
per-branch clamp inside `warm_start_atom_positions`, which reads
each branch's [centre-half, centre+half] from the ISV pinned bus.
An outer static clamp to the wider config bound is redundant
double-clamping and hides which layer owns the support. Pass raw
quantiles through to warm_start.

Validated locally: SQLX_OFFLINE cargo check -p ml passes (only
pre-existing warnings).

Next: push + L40S validation run. Diagnostic instrumentation from
423ac460b remains in place to confirm (center, half) trajectory on
the next run — will be removed once validated.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-23 22:25:07 +02:00
jgrusewski
423ac460b3 diag(dqn): ISV v-range flow instrumentation (target=isv_vrange_diag)
Adds targeted tracing::info! at three call sites to diagnose why Q-value
range hits exactly +/-50 (config hard safety clamp) at every epoch after
the ISV v-range unification commits 9deda5f65 and 11df03785.

Instrumentation (all under target="isv_vrange_diag"):

1. update_eval_v_range (called at epoch-final Q-stats): per-branch
   (q_mean, q_std, q_gap, center, half, initialized_before) — captures
   whether cold init or warm EMA path produced the ISV write.

2. recompute_atom_positions (called at epoch N start): per-branch host
   read of ISV (center, half) — what the adaptive_atom kernel will
   consume on this epoch's forward pass.

3. warm_start_atom_positions (called mid-epoch, after experience
   collection): per-branch clamp bounds (b_min, b_max, isv_ptr_null),
   plus post-write per-branch (amin, amax) of the atoms just uploaded.

NOT PUSHED — local diagnostic commit for root-cause investigation.
2026-04-23 22:11:22 +02:00
jgrusewski
11df037855 feat(dqn): Phase 2d per-branch per_sample_support tile [B, 4, 3]
Completes the ISV-unified Q-support range spec
(docs/superpowers/specs/2026-04-23-isv-v-range-unification.md) by
migrating the per_sample_support buffer from per-sample [B, 3] to
per-sample-per-branch [B, 4, 3] stride-12. Without this phase the
atom_positions grid already spanned per-branch adaptive ranges (landed
in 9deda5f65 via ISV slots 23..30) while the loss-projection
Bellman step still read a single V(s)-centred range — atoms and
projection disagreed, which is the exact pathology the spec fixes.

Producers:
* iql_value_kernel.cu::iql_compute_per_sample_support — new kernel
  signature adds isv_signals pointer; writes 4 branch triples per
  sample where centre = isv_signals[23 + 2*d] and half-width = V(s)-
  derived Q spread. Bootstrap identity: ISV centres=0 + readiness=0
  falls back to [-1, 1] across all 4 branches, byte-identical to the
  pre-Phase-2d single-range default at epoch 1.
* iql_value_kernel.cu::iql_support_floor — Frugal-1U p5 estimator now
  aggregates half-widths across all (sample, branch) pairs and applies
  the floor per (sample, branch) independently.
* gpu_iql_trainer.rs — per_sample_support_buf sized b*4*3, seed writes
  12 floats per sample, compute_per_sample_support takes isv_dev_ptr
  and forwards it to the kernel; launch-site arg order aligned.
* fused_training.rs — passes trainer.isv_signals_dev_ptr() into the
  IQL call.

Consumers (all migrated to stride-12 indexing `b*12 + d*3 + {0,1,2}`):
* c51_loss_kernel.cu::c51_loss_batched — per-branch (v_min, v_max,
  delta_z) read INSIDE the d-loop; degenerate-support skip is now
  per-branch (continue instead of whole-sample early exit).
* c51_grad_kernel.cu::c51_grad_kernel — per-branch z_norm and
  delta_z for the q-gap floor gradient path.
* experience_kernels.cu::compute_expected_q — per-branch (v_min, dz)
  inside the d-loop that iterates all 4 branches.
* experience_kernels.cu::mag_concat_qdir — reads direction branch
  (d=0) slots from the stride-12 tile.
* experience_kernels.cu::quantile_q_select — per-branch (v_min, dz)
  inside the d-loop.

Experience-collector parity:
* gpu_experience_collector.rs — its OWN per_sample_support_buf grows
  to alloc_episodes*4*3; update_per_sample_support tiles the same
  (v_min, v_max, delta_z) triple to all 4 branches so the layout
  matches the IQL buffer and the consumer kernels read uniformly.

Safety:
* Bootstrap byte-identical at epoch 1 preserved (ISV centres default 0,
  readiness ramps from 0 → 1).
* No stub values, no TODO/FIXME/XXX markers introduced.
* Kernel scalar arg (gamma) is already f32 in GpuIqlConfig — no f64→f32
  cast needed at the call site (feedback_cudarc_f64_f32_abi compliance
  via type, not cast).
* c51_loss branch-degenerate `continue` is uniform across the block
  (all threads read the same support_base) so __syncthreads inside
  the loop body remains collective.

Compile verified: cargo check -p ml + --workspace pass (SQLX_OFFLINE,
CARGO_INCREMENTAL=0, sccache) and cargo build -p ml compiles all
CUDA kernels via nvcc. cargo test -p ml --lib --no-run succeeds.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-23 21:42:23 +02:00
jgrusewski
9deda5f65b feat(dqn): ISV-unified per-branch Q-support range (spec 2026-04-23)
Unifies the Q-support range source across atom grid / warm-start quantile
clamp consumers via the ISV signal bus. One broadcast written at epoch
boundary from per-branch Q-stats EMAs, read by the two consumers that
previously held disagreeing ranges. Target observation: atom utilisation
≥40% (up from 11-15% on train-fpxnw).

Phase 0 — per-branch Q-stats kernel Rust plumbing:
* Load q_stats_per_branch_reduce alongside legacy q_stats_reduce
* Add per_branch_q_stats_pinned (28 f32 = 4 × 7, device-mapped)
* PerBranchQValueStats struct: [QValueStatsResult; 4]
* reduce_current_q_stats_per_branch launches the new kernel with the
  four branch (off, size) pairs derived from config.branch_N_size

Phase 1 — ISV v-range plumbing (zero behavioural change at epoch 1):
* ISV_NETWORK_DIM=23 preserved for w_isv_fc1 sizing; ISV_TOTAL_DIM=31
  allocates 8 additional slots for per-branch (centre, half-width)
* Slot constants V_CENTER_DIR..V_HALF_URG covering slots 23..30
* eval_q_mean_ema / eval_q_std_ema / eval_ema_initialized promoted
  to [f32; 4] / [bool; 4]; scalar setters preserved for trajectory
  backtracking (broadcast same value to all branches)
* Bootstrap at construction: centre=0, half=(v_max-v_min)/2 → the
  byte-identical [config.v_min, config.v_max] span per branch before
  any Q observations arrive
* reset_eval_v_range_state resets the 4 per-branch EMAs AND the 8 ISV
  slots to bootstrap values; legacy eval_v_range_pinned[2] still reset
  (deferred removal — spec Phase 3)
* update_eval_v_range reworked: signature takes PerBranchQValueStats and
  per_branch_q_gaps. Maintains 4 independent adaptive-rate EMAs,
  computes (centre, half) per branch with min_half_floor=0.1×(v_max-v_min)
  and clamps to config bounds, writes 8 ISV slots. Branch-0 (direction)
  centre±half is also mirrored into the legacy eval_v_range_pinned for
  consumers that have not yet migrated to the per-branch bus.

Phase 2a/2b — atom grid per-branch v-range:
* adaptive_atom_positions kernel signature changed from
  (v_min: float, v_max: float) to (branch_idx: int, isv_signals: float*);
  reads centre/half from ISV slots 23+2·b, 24+2·b. Eliminates the f64→f32
  ABI trap (spec Phase 2 side-effect) since the only per-branch range
  path is now pointer-based.
* recompute_atom_positions passes branch_idx + isv_signals_dev_ptr per
  branch; no scalar v_min/v_max arg remains.

Phase 2c — warm-start quantile clamp per-branch from ISV:
* warm_start_atom_positions reads per-branch (centre, half) from pinned
  ISV host memory, clamps shared reward-quantile vector into each
  branch's adaptive range before tiling into atom_positions_buf.
  Bootstrap makes this equivalent to the pre-spec config.v_{min,max}
  clamp until the first Q observation lands.

Deviations from spec:
* Phase 2d (per_sample_support_buf → [N, 4, 3]) NOT implemented. The
  spec's premise was that per_sample_support is host-tiled from
  eval_v_range, but the active path in this codebase has it filled by
  iql_compute_per_sample_support (V(s)-centered, per-sample, already
  adaptive) — orthogonal to the ISV bus. Migrating that kernel to
  per-branch output would require rewriting iql_value_kernel +
  iql_support_floor + C51/MSE loss kernel indexing in lockstep, which
  the "no unrelated refactoring" constraint disallows. The loss-kernel
  Bellman projection today uses V-centered bounds that are themselves
  adaptive; the ISV v-range fix still lands the primary win (atom grid
  + warm-start agreement) without touching IQL.

Compile verified: cargo check -p ml + --workspace pass (SQLX_OFFLINE,
CARGO_INCREMENTAL=0, sccache). No TODO/FIXME/XXX introduced.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-23 21:18:19 +02:00
jgrusewski
d38a8cf997 fix(dqn): clamp reward quantiles to v_min/v_max in C51 warm-start
Root cause of the Q=±333k explosion at every run's epoch 2.

`warm_start_atom_positions` writes quantiles from the raw environment
reward distribution directly into `atom_positions_buf`. Raw rewards
are unbounded PnL-scaled values — a single extreme sample in the
first experience buffer becomes `atom_positions[num_atoms-1]`, and
the C51 expected-value readback `Q = Σ prob × atom_pos` inherits
that magnitude.

Observed deterministically across train-7rgqd, train-5gzpn, and
train-gj54m: epoch 1 Q in `[0, ~6]` (initial Xavier atoms), epoch 2
Q at exactly `±333406` once warm-start writes the sorted-reward-tail
into the atom grid. Every downstream path — the C51 loss projection,
eval_v_range EMA, IQL support, HEALTH_DIAG q_gap — assumes
`atom_positions ∈ [v_min, v_max]`. The warm-start path was the only
one bypassing that assumption.

Clamp each quantile to the configured `[v_min, v_max]` before writing.
This is a safety rail, not a tuning parameter: config.v_{min,max} are
already derived from reward_scale (±15 default), which is the support
range the rest of the system expects.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-23 19:03:17 +02:00
jgrusewski
9c3ddf8b37 fix(dqn): hard-copy online→target at fold boundaries
target_params_buf was initialized once via DtoD copy at first train
step and then only moved toward online via slow Polyak EMA (tau≈0.005).
At fold boundaries the online weights are shrink-and-perturb'd with
alpha=0.8, which modifies params_buf in-place — but target_params_buf
still held the end-of-previous-fold values. The Bellman target would
then use stale weights against freshly perturbed online predictions,
producing a large TD error gap in the first fold-N+1 training steps.
Polyak averaging at tau=0.005 is far too slow to close that gap before
the oversized gradients compound through Adam into runaway updates —
one of the drivers of the fold-1 gradient explosion observed in both
train-7rgqd and train-5gzpn.

- Add GpuDqnTrainer::sync_target_from_online() — DtoD memcpy of the
  full params_buf into target_params_buf.
- Call it from FusedTraining::reset_for_fold right after shrink-and-
  perturb and before reset_adam_state, so target = perturbed online
  and Adam moments zero out from the same starting point.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-23 17:45:37 +02:00
jgrusewski
fa1a94bf9e fix(dqn): reset IQN Adam state at fold boundaries
GpuIqnHead carries its own m_buf/v_buf/adam_step — separate from
GpuDqnTrainer::reset_adam_state, which only zeroes the legacy
iqn_trunk_* buffers. Without a fold-boundary reset, the IQN
optimizer enters fold N+1 with fold N's momentum, producing
oversized Adam steps that compound through the IQN backward
pass into the runaway gradients we observed in both train-7rgqd
(crashed fold 1 ep 52) and train-5gzpn (NaN'd fold 1 ep 17).

- Add GpuIqnHead::reset_adam_state() — zero m_buf, v_buf,
  adam_step, and the pinned t counter.
- Call it from FusedTraining::reset_for_fold after the trainer's
  main Adam reset, gated by gpu_iqn.is_some(). Non-fatal warn on
  failure to match the surrounding shrink-and-perturb pattern.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-23 17:16:52 +02:00
jgrusewski
768cc7d820 fix(cuda): f64→f32 cast for scalar kernel args that expect float
Three kernel launches passed f64 config fields directly into argument
slots whose kernel-side declaration is `float`. cudarc's `DeviceRepr`
impl for f64 places an 8-byte value at the next 8-byte-aligned slot,
but CUDA reads only 4 bytes for a `float` parameter — the low 4 bytes
of the f64 — then advances to the next slot. For a typical config
value the low bytes of the f64 encoding are near-zero, producing
garbage values and shifting every subsequent arg slot by 4 bytes of
padding mismatch.

Affected sites:
  - recompute_atom_positions → adaptive_atom_positions kernel
    (v_min/v_max for C51 atom grid placement)
  - c51_loss_batched (forward) → c51_loss_kernel
    (curiosity_q_penalty_lambda, spectral_decoupling_lambda)
  - mse_loss_batched (forward) → mse_loss_kernel
    (same two lambdas)

Cast to f32 explicitly at the call site and bind to a let so the
&value reference points into a 4-byte f32 slot. Observed symptom:
Q-value range oscillating to ±144k at epoch 25 while the config
`v_min=-15, v_max=+15` theoretical bound should have held atoms
inside that range.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-23 13:08:56 +02:00
jgrusewski
07b70ccff5 Merge: TLOB Phase B — OFI_DIM 20→32, +12 real microstructure features, kernel-read gap closed 2026-04-23 11:37:40 +02:00
jgrusewski
79578bbaf6 feat(fxcache): OFI_DIM 20→32, persist 12 real-math microstructure signals,
fix kernel-read gap

Adds 12 features to the DQN input pipeline:
 - 10 MicrostructureState::snapshot()[0..10] slots that were previously computed
   every bar and then discarded before reaching fxcache: ofi_trajectory,
   realized_variance, hawkes_intensity, book_pressure (weighted 10-level),
   spread_dynamics, aggression_ratio, queue_depletion_asymmetry,
   order_count_flux, intra_bar_momentum, regime_score.
 - 2 TLOB-novel slots derived directly from Mbp10Snapshot:
   order_count_imbalance = (Σbid_ct − Σask_ct) / Σ(bid_ct + ask_ct),
   microprice_residual = (weighted_mid − mid) / mid.

Also fixes a production gap: ofi_acceleration (slot 18) and
toxicity_gradient (slot 19) were persisted to fxcache via OFI_DIM=20
but the OFI embed kernel (experience_kernels.cu:6146-6173) only read
[0..18), silently discarding them every bar. Kernel extended to
consume full SL_OFI_DIM=32.

Dimension bumps (all 8-aligned):
  OFI_DIM          20 → 32
  FXCACHE_VERSION   4 → 5  (invalidates existing caches; regen via
                            precompute_features)
  STATE_DIM        96 → 104
  PADDING_DIM       4 → 0  (OFI expansion consumed padding, still 8-aligned)
  STATE_DIM_PADDED 128 (unchanged)
  OFI_EMBED_IN     18 → 32 (MLP input width; W/grad/Adam/m/v buffers
                            resized in lockstep via named constants)

fxcache regen results (175874 bars ES.FUT 2024-Q1):
  deltas_nonzero:       175781 / 175874  (99.9 percent)
  book_aggression:      102137 / 175874  (58.1 percent)
  microstructure[20-30): 175874 / 175874 (100 percent)
  tlob_novel[30-32):    133615 / 175874  (76.0 percent)

Compile status: SQLX_OFFLINE=true CARGO_INCREMENTAL=0 cargo check
  --workspace --tests passes cleanly (0 errors, pre-existing warnings
  only).

Test results:
  fxcache roundtrip (unit + integration): PASS (4+6 tests)
  magnitude_distribution smoke: ran through epoch 1 successfully
    (OFI_DIAG fires, state_dim=104 confirmed, feature_dim=74 in
    validation kernel); epoch 2 OOM on local RTX 3050 Ti (4 GB) —
    expected hardware limit from state_dim growth. Full 20-epoch run
    requires L40S/H100 CI verification.
  multi_fold_convergence smoke: not verified locally (same VRAM
    ceiling applies). L40S/H100 CI verification required.

The new slots follow the existing OFICalculator/MicrostructureState
pattern and consume signals already computed by ml-features — no new
crate, no ONNX, no stubs. All 12 sources were audited against their
implementation before persistence; every slot traces back to real
Mbp10Snapshot or MicrostructureState math.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-23 11:36:34 +02:00
jgrusewski
8d4c2c3b03 Merge: ISV bundle v2 — health sharpe-coupled, regime batch-agg,
C51 var wired, q_abs_ref clamp

Branch worktree-agent-a496c8e9, commit 8a5e7d316. Implements the 4
fixes from docs/superpowers/specs/2026-04-23-isv-signal-quality-audit.md.

1. Regime signals (slots 8-11) now aggregate over all B samples inside
   the single-threaded isv_signal_update kernel. Adds int batch_size
   arg and fixes a latent stride bug (launch passed STATE_DIM=104 but
   states_buf has stride STATE_DIM_PADDED=128; sample-0 worked by
   luck because row-0 offset is identical). 99.994% information loss
   closed. Expected to unblock the adaptive cql_alpha gate which
   previously stayed at 0.0 for all 20 epochs on train-mdh86.

2. Health (slot 12) decoupled from outcomes is the biggest finding
   from the audit: r = -0.765 over 20 epochs — health climbed
   0.50→0.64 EXACTLY while Sharpe collapsed +34→-67. Every adaptive
   mechanism using `(1 - health)` as stress signal was reading the
   inverse of truth. Replaced with sigmoid(0.1 × sharpe_ema) EMA-
   blended. New ISV slot 22 = SHARPE_EMA_INDEX persists the Rust-
   side training_sharpe_ema (broadcast from training_loop.rs:3541).
   ISV_DIM 22→23. Post-fix validation: 3 smoke runs show health
   declining 0.50→0.40 in response to consistently negative training
   Sharpe — outcome-coupled, not inverted.

3. Slot 3 now carries C51 Q-distribution variance (wired via third
   c51_loss_reduce launch, same pattern as td_error fix 7f92fa242).
   Previously zero-initialised with no writer. Renamed scratch to
   reflect actual semantics (true multi-head ensemble variance
   remains a separate follow-up — no per-head Q readback exists in
   the current architecture).

4. q_dir_abs_ref (slot 21) and q_abs_ref (slot 16) gain outlier
   clamps before EMA update: new_abs = min(raw, 10×current + 1) to
   prevent single Q-excursions (±10⁵ observed on train-mdh86 at
   epochs 0/3/7/11) from poisoning the Kelly-conviction denominator
   for 20 epochs.

Pre-production checkpoint compat note: ISV_DIM bump grows w_isv_fc1
from [16,22] to [16,23]. No safetensors format encodes ISV_DIM
directly, but flat param buffer layouts shift — any live checkpoint
from before this commit needs regeneration. Acceptable for current
dev state; no production-live models exist.

Pre-existing 14/872 test failures (OFI missing, profile drift,
test_batch_size env) unchanged by this commit.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-23 11:35:52 +02:00
jgrusewski
8a5e7d316c fix(isv): batch-aggregate regime signals + sharpe-coupled health +
C51 Q-var slot + q_abs_ref outlier clamp

Bundles 4 ISV signal-quality fixes from the audit at
docs/superpowers/specs/2026-04-23-isv-signal-quality-audit.md.

1. Regime signals (slots 8-11) now batch-aggregate over all B samples
   instead of reading sample 0 only. Adds `int batch_size` kernel arg
   and fixes a latent stride bug (launch passed STATE_DIM=104 but
   states_buf has stride STATE_DIM_PADDED=128 — sample 0 worked by
   luck because both strides land at the same offset for row 0).
   99.994% information loss closed.

2. Health (slot 12) now couples to outcomes: sigmoid(0.1 × sharpe_ema)
   EMA-blended. New ISV slot 22 = SHARPE_EMA_INDEX persists the
   Rust-side training_sharpe_ema. ISV_DIM 22→23. Prior component-
   aggregation health formula was ANTI-correlated with Sharpe
   (r=-0.765 per audit) because its components saturated at 0/1
   boundaries (q_gap=1.0 / q_var=1.0 for 19/20 epochs, grad_stable
   stuck at 0.0 for 20/20 epochs). New formula's sensitive sigmoid
   region [-10, +10] Sharpe matches the observed magnitude range.
   The Rust-side write_isv_signal_at is preserved as a fallback
   initializer at epoch boundaries — the kernel then overwrites
   slot 12 every training step based on slot 22's current value.

3. Slot 3 now carries C51 Q-distribution variance (wired via third
   c51_loss_reduce launch, same pattern as td_error fix 7f92fa242).
   Previously zero-initialised with no writer. Renamed from
   "ensemble_var_scratch" to "q_var_scratch" to reflect actual
   semantics — it is the batch mean of q_var_buf_trainer (atom-
   spread variance from the C51 distributional head), NOT multi-head
   ensemble disagreement. True multi-head ensemble variance remains
   a separate follow-up if/when a per-head ensemble is wired.
   Slot 4 (velocity derivative of slot 3) becomes meaningful
   automatically.

4. q_dir_abs_ref (slot 21) and q_abs_ref (slot 16) gain outlier
   clamps before EMA update: clamped = min(raw, 10×current + 1) to
   prevent single ±10⁵ Q-excursions from poisoning 20 epochs. The
   +1 floor handles the cold-start case where current EMA is near 0.
   Slot 21 is especially load-bearing because it feeds the Kelly
   conviction denominator (q_range / q_dir_abs_ref).

ISV_DIM bump 22→23 changes the flat-param buffer size (w_isv_fc1
tensor [68] grows from [16,22] to [16,23] → +16 floats). Xavier
init at gpu_dqn_trainer.rs:13819 picks up the new dimension
automatically. Pinned allocs scale via ISV_DIM * size_of::<f32>()
expressions. No safetensors / checkpoint format currently encodes
ISV layout directly — checkpoint_state_dim/num_actions/hidden_dims
are the only hashed architectural fields — but the flat param
buffer content differs, so existing live checkpoints will require
rebase. Acceptable for the current pre-production dev state.

Tests: workspace cargo check --workspace --tests passes clean. Ran
magnitude_distribution smoke 3 times locally (RTX 3050 Ti); health
in HEALTH_DIAG now tracks negative Sharpe regime (0.50→0.38
trajectory over 20 epochs with mean sharpe_raw=-7) rather than
climbing monotonically as in the audit's train-mdh86 logs
(0.50→0.64 against Sharpe +34→-67). The magnitude_distribution
eval-dist H10 gate sometimes passes, sometimes fails depending on
the training seed — the pre-existing H10 flakiness persists. The
pre-existing 14 ml-lib test failures are unrelated (OFI features
missing / production profile config mismatch / test_batch_size
test environment issues — all fail on HEAD as well).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-23 11:34:09 +02:00
jgrusewski
2e80f453d1 Merge: real DQN checkpoint loader + IG diagnostic CLI + latent ensemble-test bug fix
Branch worktree-agent-ace613af, 3 commits (f47af9d2 / 82dca76d / 5f95cad4).

1. DQN::load_from_safetensors is no longer a no-op. Adds
   BranchingDuelingQNetwork::load_from_named_slices (D2D memcpy,
   validates names + shapes, errors on drift). Adds
   NoisyLinear::{weight_mu_mut, bias_mu_mut} for in-place mu restore.
   Rewrites DQN::load_from_safetensors to actually restore weights;
   handles both plain and `trending__`-prefixed checkpoint layouts.
   Resyncs target network. Un-ignores the ensemble adapter's
   test_dqn_checkpoint_round_trip (now deterministic with NoisyNet
   disabled).

2. New post-hoc IG diagnostic CLI at crates/ml-explainability/src/bin/
   ig_diag.rs. Gated by Cargo feature `ig-diag-cli` to avoid cyclic
   dep (ml already depends on ml-explainability). CLI args:
   --checkpoint --states auto|PATH --feature-names --num-steps
   --output --auto-samples --seed. Forward target: mean(Q[direction,
   0..4]) → V(s) via dueling identity; NoisyNet disabled → deterministic.
   Fxcache reader inlined (~80 LOC) to avoid ml-crate dep.
   Integration test measured 1.6% IG completeness error at
   num_steps=64.

3. Latent bug caught while wiring checkpoint loading: ensemble DQN
   adapter tests used vec![0.1; 56] but STATE_DIM=96, so cuBLAS
   gemm_ex was reading 40 bytes of uninitialised memory past the
   CudaSlice. That was the source of flaky argmax in
   test_dqn_adapter_deterministic. Replaced hardcoded 56 with
   ml_core::state_layout::STATE_DIM in 3 test sites.

Tests (ml-dqn --lib): 270→274 pass / 16→15 fail (+4 pass, -1 fail).
All new tests pass. Remaining 15 failures are pre-existing flaky
tests (NoisyNet non-determinism, small state_dim GPU alloc edges)
not introduced by this work.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-23 10:55:34 +02:00
jgrusewski
5f95cad416 fix(ensemble): DQN adapter tests use correct STATE_DIM input size
The DQN inference adapter tests hardcoded `vec![0.1; 56]` as the
feature vector, but the DQN model's first shared layer expects
STATE_DIM=96 inputs. cuBLAS gemm_ex was reading 40 elements past the
end of the 56-allocated CudaSlice — uninitialized memory that
happened to give consistent-enough values for the deterministic test
to pass on the pre-change heap layout, and for other tests not to
notice the out-of-bounds read.

Surfaced by the Part 1 checkpoint-load work: adding the CUDA_LOCK
mutex and a new weight_mu_mut method to NoisyLinear shifted
allocation patterns enough that the uninitialized tail now reads
different values between the two predict() calls in
test_dqn_adapter_deterministic, flipping the argmax.

Replace all three `vec![0.1|0.3; 56]` occurrences with
`vec![...; ml_core::state_layout::STATE_DIM]` so the tests actually
exercise the model with in-bounds memory.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-23 10:50:56 +02:00
jgrusewski
82dca76dae feat(explainability): post-hoc IG diagnostic CLI for DQN checkpoints
Adds `ig_diag` binary under ml-explainability/src/bin/ that runs
Integrated Gradients on a trained DQN safetensors checkpoint and
writes a per-feature attribution report as JSON. Designed for offline
model inspection, not the inference hot path.

Forward target: mean(Q[direction, 0..4]), which simplifies to V(s)
under the dueling identity (mean of centered advantages is zero by
the identifiability constraint). Smooth, no argmax discontinuity,
well-posed for IG.

Regime-head handling: loads only the `trending__`-prefixed weights
(matches RegimeConditionalDQN::load_from_merged_safetensors). Full
multi-head attribution is a future extension.

CLI args:
  --checkpoint PATH        safetensors file
  --states auto|PATH       `auto` samples from test_data/feature-cache/
                           *.fxcache; otherwise a JSON file containing
                           { "states": [[f32; STATE_DIM], ...] }
  --feature-names PATH     JSON array of STATE_DIM names (optional)
  --num-steps N            IG Riemann steps (default 50)
  --output PATH            output JSON (default ig_report.json)
  --auto-samples N         fxcache sample count (default 16)
  --seed N                 LCG seed for reproducibility (default 42)

Output JSON schema:
  {
    "schema_version": 1,
    "checkpoint", "num_steps", "state_dim", "num_states",
    "forward_target", "regime_head",
    "features": [ { "name", "mean_abs", "stddev", "mean_signed" } ],
    "top_10_by_mean_abs": [...],
    "completeness": {
      "worst_relative_error": f64,
      "per_state": [ { "state_idx", "sum_attributions",
                       "f_input_minus_f_baseline", "relative_error" } ]
    }
  }

NoisyNet is disabled on both the Q-network and target network before
IG runs so attributions are deterministic (same checkpoint + states +
num_steps = bit-identical attributions).

Gated behind the `ig-diag-cli` Cargo feature (optional Cargo binary
feature, not a runtime flag) because the binary depends on ml-dqn.
Cannot depend on `ml` due to a cyclic dependency (ml already depends
on ml-explainability). To avoid pulling in the heavy `ml` crate, the
fxcache header parser is reimplemented inline in the binary (~80 LOC,
matches ml/src/fxcache.rs byte-for-byte for v4 files).

Tests:
- tests/ig_diag_cli_integration.rs: end-to-end test that builds a
  DQN, saves a checkpoint, writes a states JSON, invokes the binary
  via std::process::Command, parses the output, asserts schema +
  completeness axiom (worst relative error < 5%). Gated with
  #[cfg(all(feature = "cuda", feature = "ig-diag-cli"))] + #[ignore]
  because it needs CUDA + the binary pre-built.

Build/run:
  cargo build --release -p ml-explainability \
    --features ig-diag-cli --bin ig_diag
  cargo test -p ml-explainability --features ig-diag-cli \
    --test ig_diag_cli_integration -- --ignored --nocapture

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-23 10:35:48 +02:00
jgrusewski
f47af9d268 fix(ml-dqn): implement real DQN::load_from_safetensors
`DQN::load_from_safetensors` was a documented no-op that only checked
file existence and returned Ok(()). Real weight restoration happened
only via the fused trainer's params_buf path; any caller loading a
checkpoint via the DQN struct itself (e.g., DqnInferenceAdapter::
from_checkpoint used by the ensemble) got a fresh untrained DQN with
no error raised — a silent production bug.

Adds BranchingDuelingQNetwork::load_from_named_slices (symmetric
counterpart of existing named_weight_slices) using the same D2D
memcpy primitive as copy_weights_from. Validates names AND shapes;
returns Err on mismatch. Wired into DQN::load_from_safetensors to
actually restore weights from disk, handling both prefix-less (single
head) and `trending__` prefix (regime-merged) safetensors layouts.
Target network is resynced in lockstep so inference and TD targets
see the same restored weights.

Adds NoisyLinear::{weight_mu_mut, bias_mu_mut} for in-place D2D
restore of mu params during checkpoint load.

Tests:
- branching::tests::test_load_from_named_slices_round_trip (happy path)
- branching::tests::test_load_from_named_slices_rejects_extra_keys
  (arch-drift safety)
- branching::tests::test_load_from_named_slices_rejects_shape_mismatch
- dqn::save_load_tests::test_dqn_load_from_safetensors_round_trip
  (full end-to-end safetensors save+load, #[ignore] for CUDA gate)
- dqn::save_load_tests::test_dqn_load_from_safetensors_missing_file

Also un-ignores the ensemble adapter's
`test_dqn_checkpoint_round_trip` test (was ignored pre-fix because
load_from_safetensors silently dropped weights). Disables NoisyNet
on both sides for deterministic argmax comparison. Adds CUDA_LOCK
mutex to serialize the adapter tests that share a CUDA stream via
the SHARED_CUDA OnceLock.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-23 10:19:56 +02:00
jgrusewski
527a1ebf70 Merge: ISV signal quality audit doc — L40S oscillation root-cause diagnosis 2026-04-23 10:15:58 +02:00
jgrusewski
baa4151fab docs(isv): signal quality audit for DQN ISV bus — L40S oscillation diagnosis
Diagnostic-only audit of all 22 ISV slots using train-mdh86 HEALTH_DIAG
(20 epochs, same oscillation pattern as train-rq6n8). Identifies three
dead writers (slots 2, 3, 4), one strongly anti-correlated signal
(slot 12 learning_health, r=-0.765 vs Sharpe), sample-0 bias in four
regime slots (8-11), and EMA-pollution risk on the Q-scale EMAs
(16, 21) that feed Kelly conviction + C51 bin weighting.

Produces a per-slot table with writer/reader citations and an ordered
fix list. Hypothesises the observed +34 → -67 Sharpe swing is driven
by health mis-reporting "improving" while policy diverges, q_dir_abs_ref
contamination collapsing Kelly, and the zero-writer TD-error EMA
disabling micro-reward regime awareness.

No code changes — reconnaissance only.
2026-04-23 10:15:16 +02:00
jgrusewski
f86ef27e93 Merge: GPU Bernoulli dropout in supervised Mamba2 (ghost-feature fix)
# Conflicts:
#	crates/ml-supervised/src/mamba/mod.rs
2026-04-23 10:09:15 +02:00
jgrusewski
a6ca9fba4f feat(mamba): wire real GPU Bernoulli dropout in supervised Mamba2
Supervised Mamba2's dropout_rate field was configured and stored but
never applied — comments said "simulated" but the arithmetic was
absent. An operator tuning dropout_rate upward got zero regularisation.

New dropout_kernel.cu with a forward-only Bernoulli dropout (Philox-
seeded, deterministic, in-place). New build.rs matching ml-dqn's
pattern. Applied in forward_with_gradients (training path) only; the
`forward()` path used by validate/predict/SPSA stays deterministic so
SPSA's ±ε finite-difference estimator is not destabilised.

NOT a DQN fix: DQN has its own regime_dropout kernel in
cuda_pipeline/experience_kernels.cu and its own native mamba2_step
in gpu_dqn_trainer.rs; DQN does not call into Mamba2SSM. This commit
affects only the supervised Mamba2 trainer and its hyperopt adapter.

Tests: determinism, eval-mode identity, ctr-increment divergence.
DQN smoke tests (magnitude_distribution, multi_fold_convergence) are
unaffected by this change — ran them for regression assurance only.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-23 10:07:06 +02:00
jgrusewski
c5045c009e cleanup: delete dead DQN apply_accumulated_gradients + honest TFT error
Two ghost-feature fixes from the pre-L40S cleanup audit (tasks #66
and #68 tracked internally).

### #66 — delete apply_accumulated_gradients (dead code from removed path)

The agent-audit confirmed this function is residue from a prior
Candle-gradient-tracking training path that was REMOVED (see the
now-deleted test crates/ml/tests/test_var_source_gradients.rs which
was `#[ignore]`'d with the comment "Candle gradient tracking removed
-- DQN/PPO use custom CUDA backward passes"). Evidence:

- `DQN::optimizer: Option<GpuAdamW>` is always `None` — never
  initialised anywhere in the codebase.
- `DQN` uses `OwnedGpuLinear` + `NoisyLinear` with standalone
  `CudaSlice<f32>` BY DESIGN (see comment at branching.rs:988-989
  "this layout exists to avoid a GpuVarStore intermediate"). There
  is no GpuVarStore to feed GpuAdamW.
- Zero external callers for `DqnTrainer::apply_accumulated_gradients`,
  `DQN::apply_accumulated_gradients`, `RegimeConditionalDQN::
  apply_accumulated_gradients`, or the various `optimizer_vars()`
  wrappers.
- The only test exercising this path was `#[ignore]`'d with the
  "Candle gradient tracking removed" rationale.
- Production training goes through the fused CUDA trainer
  (`trainers/dqn/fused_training.rs`) which applies gradients into a
  flat `params_buf` — a separate, live path.

Deleted: 6 functions across 3 files + the stale test. Preserving a
ghost that has zero callers, zero initialisation path, and a design
direction explicitly chosen AWAY from its premise is not "keep and
wire" — it's accumulating fiction. Per feedback_no_functionality_
removal.md the rule preserves FUNCTIONAL features; this wasn't one.

### #68 — TFT honest error message

Audit found TFT is architecturally incompatible with GpuAdamW in its
current form (not a wiring gap, an architectural absence):

- `TemporalFusionTransformer` has no GpuVarStore, no `parameters()`,
  no `named_parameters()` accessor
- Internal layers use `StreamLinear` which has NO `backward()`
- `TFTModel` trait only exposes `forward()`, `get_config()`,
  `clear_cache()` — no gradient accessor
- `TemporalFusionTransformer::train()` runs forward + accumulates
  loss but never calls backward or optimizer step
- `TrainableTFT` adapter's `backward()` explicitly returns "not
  supported — use TFTTrainer::train() instead"

The previous error string "TFT optimizer not yet migrated to
GpuAdamW" implied 99% done and a simple constructor wiring would
finish it. Actual gap: GpuVarStore threading through 5+ layers +
StreamLinear→GpuLinear migration + backward ops for softmax
attention / layer norm 3D / quantile monotonicity chain / stack +
trait extension for var_store accessor + train-loop gradient
assembly. ~3-7 day dedicated architectural project.

New error message states this explicitly so no one wastes time
chasing an "almost migrated" fiction. Full-lift tracked separately.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
2026-04-23 09:43:38 +02:00