gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity
5,053 Commits 13 Branches 0 Tags
fa30c2dd66c14a94cd166a810cb9eba6ef8ccb73
Commit Graph

2574 Commits

Author SHA1 Message Date
jgrusewski
fa30c2dd66 feat(alpha): alpha_dqn_h600_smoke — runnable Task 12 DQN smoke 
Phase E.1 Task 12. Linear Q-network (W [9×10] + b [9], no hidden layer)
trained with ε-greedy + Munchausen target on the Phase E ExecutionEnv.
End-to-end runnable: load env, train, periodically launch
alpha_kill_criteria + apply_pearls_ad chain at episode boundaries, emit
PASS/FAIL verdict against the 4 kill criteria thresholds.

Pipeline per training step (all on GPU):
  1. forward Q_current on s_batch via alpha_linear_q_forward
  2. forward Q_next on s'_batch via alpha_linear_q_forward
  3. alpha_munchausen_target → targets[batch]
  4. alpha_linear_q_grad → dW, db (sparse over taken actions)
  5. alpha_linear_q_sgd_step on W and b (separate launches)
  6. every K episodes: kill_criteria + apply_pearls_ad chain → ISV[539..542]

Pipeline visibility bumps so examples can reach launchers:
  - cuda_pipeline::alpha_kernels module → pub
  - All launch_alpha_* fns → pub
  - launch_apply_pearls → pub
  - ALPHA_LINEAR_Q_CUBIN → pub

These are appropriate pub exports (Phase E.1 public API surface).

Initial micro-smoke (horizon=100, n_episodes=50, lr=1e-6):
  Q_SPREAD_EMA         = 3.12   (≥ 0.05)    PASS
  ACTION_ENTROPY_EMA   = 2.12   (≥ 1.0986)  PASS
  RETURN_VS_RANDOM_EMA = +1.03  (≥ 0.0)     PASS
  EARLY_Q_MOVEMENT_EMA = 2268   (≥ 0.01)    PASS [unphysical scale]
  Overall: PASS (uncalibrated)

Known stability issues — flagged in the binary's CLI docstring:
  - lr=1e-4 diverges to NaN (Q grows, Munchausen target explodes)
  - lr=1e-6 stays finite but Q grows 2000× over 50 episodes
  - Follow-ups: gradient clipping, target network, reward normalisation

Bug fixed during development: `stream.memcpy_htod(&host, &mut buf.clone())`
was uploading to a TEMPORARY clone (dropped immediately) — `kc_scalar_dev`
and `kc_action_counts_dev` never got their host data → entropy=0, early_mvmt=0,
rvr stuck at the alloc-zeros default. Fixed by removing `.clone()` and using
direct `&mut` refs.

Reads:
  config/ml/alpha_fill_coeffs.json   (Task 5c)
  ISV slots 547/548                  (Task 7c baseline)

Writes:
  config/ml/alpha_dqn_h600_smoke.json (verdict + per-checkpoint KC trajectory)

Reproduction:
  cargo run -p ml --release --example alpha_dqn_h600_smoke -- \
    --mbp10-dir /home/jgrusewski/Work/foxhunt/test_data/futures-baseline-mbp10/ES.FUT \
    --horizon 600 --n-episodes 1000

Audit doc docs/isv-slots.md updated per Invariant 7.
 2026-05-15 15:39:30 +02:00
jgrusewski
36ab50814e feat(alpha): alpha_linear_q kernels + launchers for Task 12 DQN smoke 
Phase E.1 Task 12a. Three new CUDA kernels for the H=600 DQN smoke
(Task 12 proper) that lands in a follow-up commit:

  alpha_linear_q_forward_kernel    Q = X · W^T + b
  alpha_linear_q_grad_kernel       dW, db sparse MSE-TD over taken actions
  alpha_linear_q_sgd_step_kernel   element-wise params -= lr · grad

Architecture: single linear layer, no hidden layer. The Phase E state
vector has meaningful direct features (alpha_logit, spread_bps, position,
ofi_sum_5, …) so linear Q can capture real relations like Q[Buy] ∝
alpha_logit. If linear can't pass the kill-criteria gate, no architecture
upgrade will save it — and the smoke proceeds with NoisyNet escalation
per the plan.

Sparse gradient: only the taken action contributes (standard DQN TD
loss). No atomicAdd needed — one thread per (i, j) loops over the batch
and adds only when actions[b] == i.

GPU contract:
  - No host branches inside any kernel (graph-capture compatible)
  - No atomicAdd (per feedback_no_atomicadd)
  - All compute on GPU (forward, grad, weight update)
  - Tiny launch overhead — fits per-step (batch=64 forward = 576 threads,
    1 block; grad = 99 threads, 1 block)

Three pub(crate) Rust launchers in alpha_kernels.rs match the
launch_apply_pearls pattern. Cubin embedded via include_bytes!.

Smoke test `linear_q_forward_grad_sgd_round_trip_matches_hand_math`
exercises all three kernels end-to-end on a small (batch=2, state_dim=2,
n_actions=3) case with full hand-math:

  Forward:  Q = [[2.1, 3.2, 0.3], [4.1, 5.2, 0.3]] ✓
  Grad:     dW = [[-1.8, -2.7], [0.8, 1.0], [0, 0]]
            db = [-0.9, 0.2, 0] ✓
  SGD:      W' = [[1.18, 0.27], [-0.08, 0.90], [0, 0]]
            b' = [0.19, 0.18, 0.30] ✓

All within 1e-4 tolerance. `cargo test -p ml --lib alpha_kernels`:
5/5 pass on RTX 3050 Ti in 2.04s (compile witness + 4 GPU smokes).
Audit doc docs/isv-slots.md updated per Invariant 7.
 2026-05-15 15:25:23 +02:00
jgrusewski
ebbd28437a test(alpha): chained pipeline smoke — Task 12 wiring validation 
End-to-end test for the kernel COMPOSITION that the H=600 DQN smoke
(Task 12 proper) will use at each rollout boundary:

  t+0  alpha_kill_criteria_compute_kernel  → scratch[0..4]
  t+1  apply_pearls_ad_kernel(n_slots=4)   → ISV[539..543]

The two prior alpha_kernels smokes (munchausen + kill_criteria) validated
kernels in isolation. This smoke validates the COMPOSITION on the same
stream — failures here are different (stream-ordering, Pearls index base,
Wiener offset base, scratch visibility) and would silently break Task 12.

Two iterations with stationary synthetic inputs:

  Iter 1 (Pearl A bootstrap)
    prev_x_mean=0 AND x_lag=0 → ISV[539..542] populated with raw scratch
    observations = [0.2041, 0.8980, 1.0670, 0.1] within 0.01 tolerance.
    Anchor slots 547/548 remain at Task 7c values (-5185, 4953).

  Iter 2 (Pearl D stationary)
    dx_mean = dx_step = 0 → α* = 0 → ISV unchanged from iter 1 within
    1e-4. Stationary signal stays at the bootstrap value.

Test would catch:
  - Producer's scratch write not visible to applicator (stream-ordering)
  - Wrong Pearls isv_idx_base / wiener_offset_base
  - Pearl A sentinel detection broken (formula yields 0 at t=0)
  - Wiener state corruption (iter 2 drifts from iter 1)

Reuses launch_apply_pearls from sp4_wiener_ema.rs (pub(crate)). Helper
fn run_chained_iter factors the producer→applicator sequence so the two
iterations are byte-identical apart from the Wiener state's evolution.

`cargo test -p ml --lib alpha_kernels`: 4 pass (compile witness + 2 prior
GPU smokes + this chained smoke) on RTX 3050 Ti in 1.89s total. Audit doc
docs/isv-slots.md updated per Invariant 7.

Remaining Task 12 work (full H=600 DQN trainer integration with this
pipeline at rollout boundaries) is queued for a dedicated session.
 2026-05-15 15:14:56 +02:00
jgrusewski
697bb586d0 test(alpha): kill-criteria GPU smoke matching hand-computed observations 
Adds the second of the two Phase E.1 kernel smoke tests in
alpha_kernels.rs (companion to the munchausen_target smoke from
91d1a52b9). End-to-end exercises the alpha_kill_criteria_compute_kernel
launcher with synthetic inputs covering all 4 outputs:

  q_values = [[1, 2, 3], [5, 5, 5]]   →  q_spread        ≈ 0.2041
  action_counts = [10, 30, 60]        →  action_entropy  ≈ 0.8980
  rollout_R=100, isv[547]=-5185,
                 isv[548]=4953       →  return_vs_random ≈ 1.0670
  q_init=50, q_early=55              →  early_movement   = 0.1000

ISV buffer is sized to 552 floats with slots 547/548 populated using the
committed Task 7c baseline values — this exercises the production
slot-indexing path through the kernel's
`isv[random_baseline_mean_slot]` / `isv[random_baseline_std_slot]` reads,
not just isolated kernel arithmetic.

Does NOT chain apply_pearls_ad_kernel afterward — the smoothing path is
canonical SP4 applicator territory already covered elsewhere. This test
isolates the kill-criteria producer arithmetic.

Tolerance 0.01 on all four observations; passes on RTX 3050 Ti in <2s
including kernel JIT.

`cargo test -p ml --lib alpha_kernels`: 3 pass (compile witness + both
GPU smokes). Audit doc docs/isv-slots.md updated per Invariant 7.
 2026-05-15 14:52:06 +02:00
jgrusewski
91d1a52b9c refactor(alpha): rename phase_e_* → alpha_* — system-scoped naming 
The kill-criteria producer, Munchausen target kernel, Rust launchers,
fit/baseline binaries, and their output JSON artifacts are *durable
infrastructure* of the alpha trading system (live across Phase E/F/G/...),
not milestone-scoped to Phase E specifically. Aligns with the earlier
`phase_e_isv_slots.rs` → `alpha_isv_slots.rs` rename rationale.

What was renamed:

  Code files:
    crates/ml/src/cuda_pipeline/phase_e_kill_criteria.cu       → alpha_kill_criteria.cu
    crates/ml/src/cuda_pipeline/phase_e_munchausen_target.cu   → alpha_munchausen_target.cu
    crates/ml/src/cuda_pipeline/phase_e_kernels.rs             → alpha_kernels.rs
    crates/ml/examples/phase_e_fit_fill_model.rs               → alpha_fit_fill_model.rs
    crates/ml/examples/phase_e_random_baseline.rs              → alpha_random_baseline.rs

  Artifacts:
    config/ml/phase_e_fill_coeffs.json                         → alpha_fill_coeffs.json
    config/ml/phase_e_random_baseline.json                     → alpha_random_baseline.json

  Kernel function names:
    phase_e_kill_criteria_compute_kernel  → alpha_kill_criteria_compute_kernel
    phase_e_munchausen_target_kernel      → alpha_munchausen_target_kernel

  Rust launcher names:
    launch_phase_e_kill_criteria          → launch_alpha_kill_criteria
    launch_phase_e_munchausen_target      → launch_alpha_munchausen_target

  Static cubin names:
    PHASE_E_MUNCHAUSEN_TARGET_CUBIN       → ALPHA_MUNCHAUSEN_TARGET_CUBIN

Historical milestone tags in doc-comments ("Phase E.1 Task N (2026-05-15)")
are RETAINED — they record WHEN the work landed and what plan it
implemented, which doesn't change with the system-scoped rename.

Plus: ADDS the alpha_munchausen_target GPU smoke test in alpha_kernels.rs.
End-to-end validates the launcher + kernel against hand-computed expected
values: batch=2 with one terminal sample; expected targets [29.8, 1.1];
got match within 0.05 tolerance on RTX 3050 Ti. PROVES the Task 9/10
kernels actually run on GPU.

All affected references updated in:
  - build.rs (kernel compile list)
  - mod.rs (module registration)
  - state_reset_registry.rs (4 RegistryEntry descriptions for slots 539-542)
  - alpha_isv_slots.rs (slot table comment)
  - docs/isv-slots.md (audit-doc cross-references)

Verified:
  cargo test -p ml --lib alpha_kernels: 2/2 pass (including GPU smoke)
  cargo test -p ml --lib state_reset_registry: 10/10 pass
  cargo build -p ml --release --example alpha_fit_fill_model --example alpha_random_baseline: clean
 2026-05-15 14:30:40 +02:00
jgrusewski
feb2e8cc34 feat(alpha): phase_e_kernels — Rust launchers for Task 9 + Task 10 cubins 
Phase E.1 Task 11. Two pub(crate) launchers expose the kill-criteria
producer (Task 9) and Munchausen target augmentation (Task 10) for use
by future trainer integration:

  launch_phase_e_kill_criteria(stream, kernel, q_values_dev, action_counts_dev,
                               scalar_inputs_dev, isv_dev, batch, n_actions,
                               scratch_out_dev)
    → kicks the kill_criteria producer; caller chains apply_pearls_ad_kernel
      (n_slots=4, isv_idx_base=ALPHA_ISV_BLOCK_LO=539) to smooth into slots
      539..542 via Pearl A bootstrap + Pearl D Wiener-α.

  launch_phase_e_munchausen_target(stream, kernel, q_next_dev, q_current_dev,
                                   actions_dev, rewards_dev, dones_dev,
                                   gamma, alpha_m, tau, log_clip_min,
                                   target_out_dev, batch, n_actions)
    → one-thread-per-sample target augmentation; α_m/τ/log_clip_min are
      scalar args so a downstream ISV-driven controller can tune them.

Plan deviation: Task 11 plan-spec said "replace hardcoded n_step=32,
gamma=0.999 literals" but grep across gpu_dqn_trainer.rs found ZERO such
literals — the trainer already reads gamma via read_isv_signal_at(
GAMMA_DIR_EFF_INDEX) and epsilon via read_isv_signal_at(AUX_TRUNK_EPS_
INDEX). The actual E.1 deliverable was Rust launchers for the new
cubins, which this commit lands.

Both launchers follow the launch_apply_pearls pattern in
sp4_wiener_ema.rs — pre-loaded CudaFunction as parameter, u64 device
pointers, debug_assert! guards.

Audit doc docs/isv-slots.md updated per Invariant 7.
Tested via `cargo test -p ml --lib phase_e_kernels` — 1 compile-witness
passes. Real GPU integration test in Task 12.
 2026-05-15 14:17:29 +02:00
jgrusewski
b1ba41d403 feat(alpha): Munchausen DQN target term kernel (Vieillard et al. 2020) 
Phase E.1 Task 10. Standalone target-augmentation kernel:

  m            = α_m · max(τ · log π(a|s), log_clip_min)
  V_soft(s')   = max(Q_next) + τ · log Σ exp((Q_next − max) / τ)
  target       = r + m + γ · V_soft(s')   (terminal: r + m)

π(a|s) ∝ exp(Q_online(s, a) / τ) — softmax policy from the online net.
Munchausen bonus is implicit KL regularisation between successive policies;
soft-V replaces the hard max bootstrap with a τ-weighted softmax average.

Both softmaxes are computed via log-sum-exp with the max-trick. This is
essential at τ ≈ 0.03 where raw exp(Q/τ) would overflow f32 for any
Q-spread > 25 nats. The kernel is one-thread-per-batch-sample, no
atomicAdd, no host branches.

α_m, τ, log_clip_min are kernel args (not hard-coded), so a Phase E.2+
controller can ISV-drive them. Typical Vieillard values: α_m=0.9, τ=0.03,
log_clip_min=-1.0.

Does NOT touch any ISV slot — pure target augmentation.

Cubin: target/release/build/ml-*/out/phase_e_munchausen_target.cubin (12.8 KB).

Launcher integration is Task 11 (consumes target_out where the C51/MSE
loss kernels currently consume `r + γ · max_a' Q_target`). Audit doc
docs/isv-slots.md updated per Invariant 7.
 2026-05-15 14:11:00 +02:00
jgrusewski
d493b729bf feat(alpha): phase_e_kill_criteria producer kernel (slots 539-542) 
Phase E.0 Task 9. Single-block, single-thread CUDA producer that writes 4
raw scalar observations into a contiguous scratch_out[0..4] block:

  scratch_out[0] → ISV[539] Q_SPREAD_EMA          (kill: ≥ 0.05)
  scratch_out[1] → ISV[540] ACTION_ENTROPY_EMA    (kill: ≥ 0.5·ln(9) ≈ 1.10)
  scratch_out[2] → ISV[541] RETURN_VS_RANDOM_EMA  (kill: ≥ 0, i.e. ≥ random)
  scratch_out[3] → ISV[542] EARLY_Q_MOVEMENT_EMA  (kill: ≥ 0.01, learned)

Downstream `apply_pearls_ad_kernel` (n_slots=4) chained on the same stream
applies Pearl A first-observation bootstrap + Pearl D Wiener-α smoothing —
composes with the canonical val_sharpe_delta_compute_kernel pattern rather
than reimplementing Wiener math (per feedback_no_cpu_compute_strict — one
Wiener implementation in the codebase, the GPU one).

Inputs:
  q_values[batch, n_actions]   most-recent Q forward output (device)
  action_counts[n_actions]     empirical action histogram (device)
  scalar_inputs[3]             [rollout_R_mean, q_init_norm, q_early_norm]
                               via mapped-pinned (host writes between rollouts)
  isv                          reads slots 547 + 548 (random baseline
                               mean/std — populated once by Task 7c,
                               TrainingPersist)

No atomicAdd (per feedback_no_atomicadd), no host branches
(per pearl_no_host_branches_in_captured_graph), no CPU compute
(per feedback_cpu_is_read_only). __threadfence_system() before exit for
the chained Pearls applicator's visibility guarantee.

Cubin produced: target/release/build/ml-*/out/phase_e_kill_criteria.cubin
(16.8 KB).

Launcher integration is Task 11 (separate commit so this task can stand
alone for cubin validation). Audit doc docs/isv-slots.md updated per
Invariant 7.
 2026-05-15 14:07:41 +02:00
jgrusewski
9cfc6d8502 feat(alpha): phase_e_random_baseline example + reset_at extension 
Phase E.0 Task 7b. Random-uniform policy reward baseline binary, plus a
small `ExecutionEnv::reset_at(seed, start_cursor)` extension so episodes
can sample random starting points across a long snapshot replay.

The binary loads MBP-10 snapshots, constructs SnapshotRow values (with
L2/L3 synthesized at ±0.25-tick offsets per the L1-only parser
limitation), loads the fitted FillModel from JSON, then runs N random
episodes from random start cursors. Reports mean / std / quintile
percentiles + kill threshold (mean + 2σ) for E.1 to exceed.

Smoke run (500 episodes, horizon 600, 100K snapshots):
  mean = -5600 (dominated by terminal force-close variance + market-order
                 over-reliance because fit converged to β_spread = -40
                 → limit fill probability ~0 at typical spreads)
  std  = 5383
  p95  = -895
  kill threshold (mean + 2σ) = +5167

The deeply negative baseline is correct *for this env* even though it
doesn't reflect realistic random-policy P&L. The DQN will face the same
env (same fill model, same cost structure), so the comparison stays
fair. Fitter regularisation (to prevent β_spread runaway) is a Phase E.1
follow-up.

Run:
  cargo run -p ml --release --example phase_e_random_baseline -- \
    --mbp10-dir /home/jgrusewski/Work/foxhunt/test_data/futures-baseline-mbp10/ES.FUT \
    --fill-coeffs config/ml/phase_e_fill_coeffs.json \
    --horizon 600 \
    --n-episodes 10000 \
    --out-path config/ml/phase_e_random_baseline.json

env.reset_at also called by reset() (1-line refactor); no behavior change.
 2026-05-15 13:36:43 +02:00
jgrusewski
12151ccf6a feat(alpha): fitted FillModel coefficients from 500K ES.FUT snapshots 
Phase E.0 Task 5c. Ran phase_e_fit_fill_model on the ES.FUT 2024-Q1 MBP-10
+ trade tape (5.2M trades, 3.9M MBP-10 events, 500K snapshots accumulated
at snapshot_interval=50 over a ~24-minute window). Total runtime ~80s.

Empirical fill rates within 60s window:
  - bid_l1: 4.97%   (matches L1 maker-side activity in trending market)
  - ask_l1: 71.34%  (high — most 60s windows see an aggressive buy)

Fitted L1 cloglog coefficients (all 5 features):
  BID L1: β_0=-0.213  β_spread=-2.064  β_imbal=-0.099  β_ofi=-0.006  β_logτ=-0.286
  ASK L1: β_0=+0.016  β_spread=-40.336 β_imbal=+0.041  β_ofi=+0.652  β_logτ=-0.055

Sanity (sign checks all pass):
  - β_spread < 0 both sides   (wider spread → fewer fills) ✓
  - bid β_imbal < 0           (more bid stack → harder to get hit by sell) ✓
  - ask β_ofi > 0             (buying pressure correlates with ask fills) ✓
  - β_logτ < 0 both sides     (quieter markets → slower execution) ✓

L1-only limitation: as documented in the binary header, the parser only
populates levels[0]; L2/L3 in the JSON are L1 with β_0 -= ln(L+1) attenuation.

Default --out-path bumped to config/ml/phase_e_fill_coeffs.json so future
re-runs land in the same committed location.
 2026-05-15 13:26:21 +02:00
jgrusewski
3bdf74018d feat(alpha): phase_e_fit_fill_model example — cloglog FillModel calibration 
Phase E.0 Task 5b. Calibrates FillModel coefficients from historical MBP-10
+ trade tape. Streams snapshots concurrently with time-sorted trades; per
snapshot, determines binary fill outcome ("would a posted L1 limit have
been hit within next --window-seconds?"), accumulates (FillFeatures, y),
calls fit_poisson (cloglog Bernoulli, see d08ab461d). Writes 6 fitted
coefficient sets to JSON.

L1-only limitation: DbnParser::parse_mbp10_streaming ignores
Mbp10Msg.levels[1..10] (only stores levels[0] via update_level(0, ...)),
so this binary fits L1 distributions only and replicates them across
L2/L3 with β_0 -= ln(L+1) attenuation. The parser bug is documented
inline; fixing it is out of Phase E.0 scope.

Scale-bug workaround: parser stores mbp10.price (1e9 fixed-point) directly
into BidAskPair.bid_px/ask_px, but BidAskPair::price_to_f64 divides by 1e12
(different convention). Net: helper returns prices 1000× too small. Binary
uses raw_price_to_f32 (i64 * 1e-9) directly — confirmed in smoke run
(bid_l1=0 with helper, bid_l1=4500-range with workaround).

Smoke run (2K snapshot cap):
  - 5.2M trades loaded, front-month filtered
  - 19.7M MBP-10 events in file → 2K accumulated via interval=10
  - bid_l1 empirical = 0.7%, ask_l1 empirical = 19.4% (uptrend bias in
    early-2024 file region; data, not bug)
  - bid β_spread = -4.05, ask β_spread = -0.39 (signs sane: wider
    spread reduces fill probability)
  - Full run pending (sequential mode per user)

Run:
  cargo run -p ml --release --example phase_e_fit_fill_model -- \
    --mbp10-dir /home/jgrusewski/Work/foxhunt/test_data/futures-baseline-mbp10/ES.FUT \
    --trades-dir /home/jgrusewski/Work/foxhunt/test_data/futures-baseline-trades/ES.FUT \
    --window-seconds 60 \
    --snapshot-interval 50 \
    --out-path phase_e_fill_coeffs.json
 2026-05-15 13:23:54 +02:00
jgrusewski
a1e3336b1f feat(alpha): ExecutionEnv — 10-dim state, 9-action gating, terminal-only reward 
Phase E Task 6. Core environment for the execution-policy DQN.

Episode lifecycle: position-open signal → up to `horizon_snapshots` →
forced close. Within the episode the policy chooses 9 actions; illegal
actions degrade to Wait inside step() (soft masking) — the network's
Q-output is never masked, preserving clean C51 and Munchausen targets.

Reward is **terminal-only**: realized PnL minus fees, accumulated over
all fills and force-closed at terminal mid if still open. No dense
per-snapshot shaping per pearl_event_driven_reward_density_alignment
(the canonical SP11→SP12 lesson — dense shaping on event-driven
objectives creates exposure-positive bias).

Bugs in the plan's sketch, caught + fixed during implementation:

  1. Synthetic test's bid/ask were CONSTANT while mid drifted — would
     have made the "uptrend" test lose money. Fixed: drift bid/ask
     together with mid.
  2. L1/L2 closing path used Side::None for fill lookup → 0% fill rate
     for closing limits. Fixed: closing-long → Side::Sell (post at ask),
     closing-short → Side::Buy (post at bid).
  3. `cursor`-based hash for fill PRNG was non-deterministic across
     replay seeds. Replaced with self-contained SplitMix64 RNG state
     (no `rand` dep added) — fill randomness is now bit-identical across
     train/eval replays at the same seed.
  4. `max_step_per_episode` and `mid_at_decision` fields stored but
     never read — dead per feedback_no_stubs. Removed.
  5. Test harness used n==horizon, which made the cursor-exhaustion
     guard (`cursor+1 >= len`) fire before the planned FlatMarket call.
     Fixed by setting n > horizon so the meaningful `ep.step >= horizon`
     path terminates the episode.
  6. `.unwrap()` in test bodies → `.expect("…")` per pre-commit policy.
  7. Volume-weighted entry_price across multi-fill scaling (the plan
     only handled the open-from-flat case).

5 unit tests:
  - state_has_expected_dim_and_is_finite (STATE_DIM=10, no NaN)
  - market_buy_then_market_sell_on_uptrend_profits (PnL math)
  - illegal_action_degrades_to_wait (soft mask, no fee, no fill counted)
  - terminal_force_close_pays_out_when_position_open (force-close at mid)
  - rng_is_deterministic_across_resets (same seed → identical fills)

`cargo test -p ml --lib env`: 16 passed (4 action_space + 7 fill_model
+ 5 execution_env); full module compiles clean with no new warnings.
 2026-05-15 12:42:00 +02:00
jgrusewski
d08ab461db feat(alpha): fit_poisson via cloglog likelihood + Serde derives on FillCoeffs 
Phase E Task 5 (code portion; the 5.2M-trade fit run lands separately).

Diagnosis: the plan's draft used pure Poisson NLL with Bernoulli y∈{0,1},
which converges to λ = empirical rate ȳ. But the runtime fill_prob() uses
`p = 1 − exp(−λ)`, so a trained λ=0.4 → predicted p=0.33 → systematic
~30pp under-fill bias on every passive backtest order. Training and
inference must agree on what λ means.

Fix: switch the fitter to the **cloglog (complementary log-log) binary
likelihood**. Per-sample gradient:

  ∂L/∂β_k = (p − y) · (μ/p) · x_k
  where μ = exp(β·x),  p = 1 − exp(−μ)

The μ/p factor is the link derivative; p.max(1e-7) handles the μ→0 limit
in f32 (the analytical limit μ/p → 1 is achieved automatically because
both numerator and denominator vanish proportionally).

Recovery test verifies the fix: 1000 deterministic 40% fills, all-zero
features → fitter recovers β_0 ≈ ln(0.5108) ≈ -0.672 (the value at which
1 − exp(−exp(β_0)) = 0.4), within tolerance 0.05. Slope coefficients
stay near zero (features uninformative).

Also adds Serde derives on FillCoeffs / FillFeatures / FillModel for JSON
serialization (downstream when the calibration example lands).

Deferred: the calibration example (Steps 3-7 of the plan task) is held
back until paired with the actual 5M-trade fit run — the plan's example
has a placeholder loop that violates feedback_no_stubs, and the loader
lift from precompute_features.rs deserves a dedicated commit.

4 new tests (7 total in env::fill_model now):
  - fit_recovers_baseline_when_features_uninformative
  - fit_rejects_empty_and_mismatched_inputs
  - coeffs_round_trip_through_json

`cargo test -p ml --lib env::fill_model`: 7 passed.
 2026-05-15 12:32:56 +02:00
jgrusewski
a5de50503f feat(alpha): Poisson regression fill model scaffold (coeffs fitted in Task 5) 
Phase E Task 4. Medium-tier fill simulator per the design memo.

For each (level ∈ {L1,L2,L3}, side ∈ {Bid,Ask}):
  λ(features) = exp(β · [1, spread_bps, L1_imb, OFI_5, log(τ+1)])

Per-snapshot Bernoulli fill probability for a posted limit order:
  p = 1 − exp(−λ)

Market orders fill immediately at the opposite-side L1 quote (no slippage
modeled at medium tier). Closing actions (Side::None) return λ=0 — they
are handled separately in the env.

Scaffolding only. Task 5 fits the 30 coefficients (6 distributions × 5
features) from the 5.2M-trade historical tape. The skeleton constructor
uses β=0 → λ=1 → p≈0.632, a stable sanity default for early smokes.

Numeric guard: `linear.exp().min(50.0)` caps λ to prevent f32 overflow
under outlier features before fitting lands. Fitted models should stay
well below this cap in practice.

4 unit tests:
  - skeleton_has_uniform_fill_prob (β=0 → p≈0.632 within 0.01)
  - fill_prob_bounded_under_outlier_features
  - lambda_cap_prevents_f32_overflow (β=100 outlier path)
  - side_none_returns_zero_rate

`cargo test -p ml --lib env::fill_model`: 4 passed.
 2026-05-15 12:27:57 +02:00
jgrusewski
94f8f65571 feat(alpha): 9-action discrete execution action space + legality gating 
Phase E Task 3. Introduces crates/ml/src/env/ for the execution-policy
environment. This commit lands the action space only; fill_model (Task 4)
and execution_env (Task 6) append to mod.rs in their respective commits
per feedback_wire_everything_up (no orphan declarations).

- 9-action discrete space: Wait + {Buy,Sell,Flat} × {Market,L1,L2 / L1}
- Soft action masking: illegal actions degrade to Wait in env::step
  (not by masking Q-output) per design memo — preserves clean C51
  categorical targets and Munchausen term (Task 10)
- `is_legal(position)` with position ∈ {-1, 0, +1} (sign only;
  magnitude is decoupled into the Kelly layer in Task 11)
- repr(u8) discriminants round-trip with from_u8 for replay-buffer
  storage by the existing GPU DQN trainer

4 unit tests:
  - n_actions_matches_enum_cardinality
  - legality_gating_by_position (covers flat/long/short × all 9 actions)
  - decode_closing_flag_is_only_flat
  - repr_u8_round_trip (replay-buffer contract)

`cargo test -p ml --lib env::action_space`: 4 passed.
 2026-05-15 12:25:23 +02:00
jgrusewski
8bf8bdf874 feat(alpha): state_reset_registry entries + dispatch arms for slots 539..548 
Per feedback_registry_entries_need_dispatch_arms — both must land in the
same commit; the test `every_fold_and_soft_reset_entry_has_dispatch_arm`
walks training_loop.rs::reset_named_state source and asserts coverage.

- 10 RegistryEntry rows appended after SP22 block (full producer/consumer
  rationale per existing dense-description pattern)
- 7 dispatch arms in reset_named_state (the 3 TrainingPersist anchors
  — slots 544/547/548 — are not dispatched per the existing convention)
- All sentinels 0.0 per pearl_first_observation_bootstrap

All 10 registry tests pass. Audit doc docs/isv-slots.md updated per
Invariant-7.
 2026-05-15 12:19:30 +02:00
jgrusewski
aa5908aa53 feat(alpha): reserve ISV slot block 539..550 for alpha trading system 
12 contiguous slots reserved for durable alpha-system infrastructure:
- 539-542: diagnostics (Q-spread, action entropy, return-vs-random,
  early-Q-movement EMAs) — read by Phase E kill criteria
- 543-546: stacker-threshold controller (threshold, target, observed-rate,
  Kelly attenuation) — engagement-rate self-correction
- 547-548: random-uniform baseline (mean, std) — anchor for kill criterion 541
- 549-550: reserved spare (absorb growth without bumping ISV_TOTAL_DIM)

Named `alpha_isv_slots` rather than `phase_e_isv_slots` because these slots
are intended to outlive any individual Phase E/F/G milestone — the SP4..SP22
naming was iteration-scoped, but this block is system-scoped infrastructure.

ISV_TOTAL_DIM bumped to 551. All indices validated in 4 unit tests
(bounds, membership, uniqueness, total-dim coverage). Audit doc
docs/isv-slots.md updated per Invariant-7.
 2026-05-15 10:53:41 +02:00
jgrusewski
db874b1841 feat(foxhuntq): Phase 1c snapshot-resolution alpha + leakage fix + variable-dim fxcache 
Three things landing atomically because they're load-bearing for each other:

1. **Trend-scanning leakage fix** — trend_scanning.rs was emitting OLS slope+t-stat
   over a *forward* window [t, t+L]. With the Phase 1a label = sign(price[t+60]
   − price[t]), the forward feature window overlaps the label window, contaminating
   it. Purged walk-forward only sterilizes forward-looking *labels* that cross
   the train/val split, not forward-looking *features* that peek inside the same
   horizon the label measures. The leak inflated MLP accuracy from 0.49
   (legacy 74-dim baseline) to 0.75 — vanished to 0.50 after switching to a
   trailing window. Bounded the perfect-fit t-stat sentinel from ±1e6 → ±20
   (p<1e-30 is already meaningless); eliminated the 16k corruption-cap drops.

2. **Variable-dim alpha column** — fxcache schema now carries the alpha-feature
   width via metadata (`alpha_feature_dim`), not a compile-time constant. Same
   on-disk format hosts the 134-dim bar-level stack OR the 81-dim snapshot stack.
   Reader + auto-detect honor the metadata-declared dim; downstream MLP auto-sizes
   `in_dim`. Single schema, no forks.

3. **Snapshot pipeline (Phase 1c falsification)** — `snapshot_pipeline.rs`: 81-dim
   per-MBP10-snapshot extractor reusing 10 snapshot-native alpha blocks + 6 new
   snapshot-specific features (time-since-trade, time-since-snap, event-rate,
   spread-bps, L1-imbalance, microprice-mid drift). `precompute_features` gets
   `--row-unit snapshot` flag; emits one fxcache row per LOB update (1.97M rows
   from MBP-10 data vs 206K for bar mode).

**Smoke verdict on real data** (ES.FUT, 1.97M snapshots, 384K val):
- Bar-level honest alpha: accuracy=0.5005, AUC=0.5043 (no signal)
- **Snapshot-level alpha**: accuracy=0.5241, AUC=0.6849 (real signal, 384K val)
- GBM corroboration: accuracy=0.5401 (non-linear partitioning sees more)
- Horizon decay: alpha peaks at K=20-50 snapshots (~5-25ms), gone by K=500
- Regime-conditional: spread-Q4 quintile hits 0.752 accuracy on 76k samples

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
 2026-05-15 01:01:15 +02:00
jgrusewski
19bb3bc3c8 fix(sp22-vnext): aux_outcome CF half mirror on-policy (NOT all -1) 
Smoke train-k95mj epoch 0 + epoch 1 printed
`trade_outcome_ce=0.000e0` — K=3 CE EMA stuck at Pearl A sentinel
across all training steps.

Root cause chain:

  1. insert_batch is called with bs = total = 4_096_000 (cf-mult-
     expanded), replay cap = 300_000 (L40S GpuProfile). Tail-clip:
     off = 3_796_000, slice(off..) takes the last 300K elements.
  2. Prior fix (256a5fa5a) filled CF half [base_total, total) with
     -1 mask via cuMemsetD32Async. Last 300K elements at
     [3_796_000, 4_096_000) are entirely in this CF half → 100%
     mask.
  3. Replay ring fills with 300K mask labels. PER gather samples
     batches → every batch has label == -1 for all samples →
     aux_trade_outcome_loss_reduce returns
     loss = 0 / fmaxf(valid=0, 1) = 0 every step.
  4. aux_outcome_ce_ema_update bootstrap requires `loss > 0.0f`;
     never fires → ISV[538] pinned at 0.0 sentinel.

Fix: CF half mirrors the on-policy half (matches K=2 sibling
aux_sign_labels which writes the same bar-resolved label to both
halves). Replace cuMemsetD32Async(-1) + single memcpy with TWO
memcpy_dtod calls (on-policy half + CF half), both sourced from
the same exp_aux_to_label_per_sample[base_total].

Semantic justification: the K=3 outcome label depends on the bar's
trade-close event, not the action. The CF action's hypothetical
trade outcome at the same bar approximates to the same label. The
B4b-1 kernel writes -1 to non-close slots (~97%) and 0/1/2 to
close slots (~3%); duplicating into CF preserves this sparsity →
ring tail retains ~7-9K real labels per 300K-element fill →
Pearl A bootstraps → K=3 head trains.

Lib suite 1016/0 green. Bug was runtime-only (small-batch lib
tests don't trigger the off > 0 tail-clip path). Audit doc
updated.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-14 16:19:15 +02:00
jgrusewski
878cc9ba72 feat(sp22-vnext): F-3c follow-up — K=3 CE EMA in stdout HEALTH_DIAG aux line 
The F-3c commit (cdd3dc6ed) pushed `aux_trade_outcome_ce_ema` into
the regression-detection metrics vec but never wired a `tracing::
info!` console emit. Smoke `train-q5k5k` ran clean past the
rollout + post-rollout phases — but `argo logs train-q5k5k | grep
aux_trade_outcome` returned zero hits, so the operator had no way
to read the K=3 head's CE EMA trajectory.

Fix: extend the existing K=2 aux HEALTH_DIAG print

  HEALTH_DIAG[N]: aux [next_bar_mse=… regime_ce=… w=…]

to include `trade_outcome_ce=…`:

  HEALTH_DIAG[N]: aux [next_bar_mse=… regime_ce=…
                       trade_outcome_ce=… w=…]

Reads ISV[538] via the same `trainer.read_isv_signal_at(...)`
accessor the regression-detection vec uses — single source of
truth, no duplicated mirroring.

Expected operator-visible trajectory across smoke epochs:
  Epoch 0 (cold-start, Pearl A sentinel):     0.000
  Epoch 1+ (first non-zero bootstrap):       ~1.099 (= ln(3))
  Healthy learning:                           0.5 - 0.7
  Pinned at ln(3) for many epochs:            head can't learn

cargo check -p ml clean. Audit doc updated.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-14 15:47:01 +02:00
jgrusewski
256a5fa5ab fix(sp22-vnext): K=3 aux_outcome_labels CF-half buffer underrun 
Phase B4b-2 introduced `exp_aux_to_label_per_sample` sized
`[alloc_episodes × alloc_timesteps]` (no cf-mult expansion — the
B4b-1 per-step kernel only writes the on-policy half). The batch
finalisation cloned into the emitted `aux_outcome_labels` with size
`total = base_total × 2` (cf-mult expanded), so
`dtod_clone_i32(src, total, ...)` invoked `src.slice(..total)` on a
half-sized source → `CudaSlice::try_slice` returned `None` → the
internal `unwrap()` panicked at cudarc safe/core.rs:1648.

Repro: workflow train-xzv56 panicked after rollout completed
(timestep=999) on fold 0; rollout itself ran clean, the OOM from
the prior commit is gone.

Fix: replace the single `dtod_clone_i32` with a 3-step build:

  1. alloc_zeros::<i32>(total)
  2. cuMemsetD32Async(ptr, 0xFFFFFFFFu32, total, stream) — fills
     all `total` slots with i32 mask sentinel -1 (byte pattern
     0xFFFFFFFF reinterprets as i32(-1))
  3. memcpy_dtod the first `base_total` real labels from
     exp_aux_to_label_per_sample into the on-policy half

CF half remains at -1. The K=3 sparse-CE loss masks `label == -1`
out of the mean and B_valid count, so CF samples contribute zero
gradient — exactly the semantic we want (CF actions have no
observed trade-close outcome to predict against).

Lib suite: 1016/0 green. Audit doc updated.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-14 15:26:46 +02:00
jgrusewski
0acf77e656 config(dqn): strip H100-tuned VRAM overrides from dqn-production.toml 
dqn-production.toml hard-coded three H100-tuned values that shadow
GpuProfile auto-detection: batch_size=16384, buffer_size=500K,
gpu_n_episodes=4096. After the L40S-default flip lands (prior commit
8b8bb1af7), workflow `train-ft8ph` deterministically OOMed at fold 0/1/2
with `build_next_states_f32` 4 GiB alloc — because the H100-sized
hyperparams.batch_size + buffer_size + gpu_n_episodes ate ~38 GB of the
L40S's 46 GB usable VRAM before the rollout step.

DqnTrainingProfile.apply_to() runs AFTER train_baseline_rl.rs populates
hyperparams from GpuProfile, so the production TOML always wins. All
three fields are `Option<...>` in the TOML schema — removing the lines
turns apply_to into a no-op for them, and the GpuProfile-detected values
flow through:

  field            | L40S  | H100   | (was forced)
  batch_size       | 4096  | 8192   | 16384
  buffer_size      | 300K  | 500K   | 500K
  gpu_n_episodes   | 2048  | 4096   | 4096

Two pinned assertions in training_profile.rs::tests checked the old
contract `hp.batch_size == 16384`. Rewritten to assert
`hp.batch_size == baseline_batch_size` — locks the new contract that
VRAM-tuned values stay GpuProfile-sourced.

Lib suite 1016/0 green. Audit doc updated.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-14 15:05:06 +02:00
jgrusewski
bbd52c3aa7 feat(sp22-vnext): FoldReset registry entries for B5b/C collector buffers 
Two collector-side device buffers used by the K=3 trade-outcome head
were missing FoldReset registry coverage:

  - prev_aux_outcome_probs [alloc_episodes × 3]
    TRUE stale-read risk. Producer writes end-of-step; consumer
    (experience_state_gather) reads start-of-next-step into
    state[121..124). Without FoldReset the new fold's step-0 state
    gather would inject the previous fold's last-step softmax probs
    into the first batch's state slots.

  - exp_aux_to_input_buf [alloc_episodes × 262]
    Cleanliness-only. Concat kernel overwrites all 262 columns every
    step before the K=3 forward reads them, so no steady-state stale-
    read risk. Registered for parity with the rest of the K=3
    pipeline + to satisfy feedback_registry_entries_need_dispatch_
    arms (the pin test asserts every registry entry has a matching
    dispatch arm in reset_named_state).

Both fields promoted to pub(crate) on GpuExperienceCollector so
reset_named_state can reach them. Matching dispatch arms added with
the standard memset_zeros pattern (is_win_per_env / hold_baseline_
buffer style).

Tests: All 10 state_reset_registry tests pass, including the critical
every_fold_and_soft_reset_entry_has_dispatch_arm pin test that walks
the dispatch body and validates parity with registry entries. Full
lib suite 1015/1 (the failing test is the pre-existing
test_dqn_checkpoint_round_trip NoisyLinear flake — pred1/pred2 sign
mismatch surfacing ~30-50% of full-suite runs, documented in
project_sp22_h6_vnext_resume memory as unrelated to this work).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-14 13:12:34 +02:00
jgrusewski
4b40710b7c feat(sp22-vnext): Phase B5b-2 collector trade plan forward — resolves K=3 asymmetry 
Phase B5b's K=3 input concat passed plan_params=NULL because the
collector had no trade plan launch. Trainer-side K=3 forward trained
on real plan_params while the collector queried at plan_params=0 —
documented train/inference asymmetry on the plan-conditioning surface.

Phase B5b-2 mirrors the (now-corrected) trainer
`launch_trade_plan_forward` chain on the collector inside the rollout
step:

  1. SGEMM:    hidden[N, AH] = h_s2_q[N, SH2] @ W_fc[AH, SH2]^T
  2. bias+relu in-place on hidden
  3. SGEMM:    pre_out[N, 6]  = hidden[N, AH] @ W_out[6, AH]^T
  4. trade_plan_activate → exp_plan_params[N, 6]

Weight resolution uses the same `aux_w_ptrs` array the K=3 forward
already consumes (`f32_weight_ptrs_from_base`); indices 91-94 match
the corrected trainer-side reads. The `trade_plan_activate` kernel is
loaded from `EXPERIENCE_KERNELS_CUBIN` (same cubin the rest of
`exp_module_extra` uses; the trainer loads it from there too).

The K=3 concat now takes `exp_plan_params.raw_ptr()` instead of NULL
— both sides see f(h_s2; W_plan_*_init), symmetry restored. Plan
tensors at [91..94] still have no backward (no Adam updates), so the
plan-head weights stay at Xavier cold-start forever. This commit
delivers the symmetry the K=3 head requires, not a learned plan
signal — adding a real plan-head backward is a follow-up project.

New struct fields on GpuExperienceCollector:
  - exp_trade_plan_hidden_buf  [alloc_episodes × adv_h]
  - exp_trade_plan_pre_out_buf [alloc_episodes × 6]
  - exp_plan_params            [alloc_episodes × 6]
  - exp_trade_plan_activate_kernel: CudaFunction

Lib test suite: 1016/0 green maintained. Audit doc updated.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-14 12:41:59 +02:00
jgrusewski
fb9b62a1f9 fix(plan-head): trainer trade plan forward reads from wrong param indices 
`launch_trade_plan_forward` read `w_fc/b_fc/w_out/b_out` from
`padded_byte_offset(&param_sizes, [82..86))` — i.e., the ISV-conditioning
+ recursive-confidence tensors (`b_isv_gate`, `w_isv_gamma`,
`b_isv_gamma`, `w_conf_fc`). The actual plan tensors live at
`[91..95)` per `compute_param_sizes` and the matching Xavier
init block. No backward exists for the plan head, so the plan
tensors at `[91..94]` sat at cold-start Xavier values forever
while `plan_params` was being driven by whichever ISV/conf
weights happened to occupy the wrong offsets.

Every downstream consumer (`backtest_plan_kernel`, `plan_isv` slots
in `experience_kernels`, `compute_plan_params` in `q_value_provider`,
regime gating, and the SP22 H6 vNext B5b concat path) was therefore
conditioning on noise correlated with ISV optimisation, not on a
learned plan. Discovered while preparing Phase B5b-2 (collector
trade plan launch). Two-line index fix + diagnostic comment +
audit doc entry.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-14 12:35:48 +02:00
jgrusewski
cdd3dc6edb feat(sp22-vnext): Phase F-3c — HEALTH_DIAG snap + console line for K=3 CE EMA 
Completes the F-3 observability story. Smoke runs now print
aux_trade_outcome_ce_ema every epoch in the standard HEALTH_DIAG output
— no ISV inspector required.

Changes:

health_diag.rs:
- New field aux_trade_outcome_ce: f32 appended at END of
  HealthDiagSnapshot per "Field order is stable; adding fields
  appends to the end" doc rule. Existing fields' byte offsets
  preserved → no kernel-side word offset re-validation.
- snapshot_size_is_stable pin: 149 × 4 → 150 × 4 = 600 bytes.

health_diag_kernel.cu:
- New WORD_AUX_TRADE_OUTCOME_CE = 149 (appended at end).
- WORD_TOTAL = 150 (was 149); static_assert bumped.
- New kernel arg aux_trade_outcome_ce_idx appended after
  moe_lambda_eff_idx.
- New mirror write after the existing MoE mirrors. Stream-implicit
  ordering: aux_outcome_ce_ema_update (F-3b) fires before
  health_diag_isv_mirror, so the read picks up the just-updated EMA.

gpu_health_diag.rs:
- launch_isv_mirror gets new aux_trade_outcome_ce_idx: i32 arg.

gpu_dqn_trainer.rs:
- launch_health_diag_isv_mirror passes
  AUX_TRADE_OUTCOME_CE_EMA_INDEX as the new arg.

training_loop.rs:
- Per-epoch metrics push appends ("aux_trade_outcome_ce_ema",
  ISV[538]) to the standard out vec. Console / CSV automatically
  includes the new column.

End-to-end F-3 chain now closed:
  K=3 fwd → aux_to_loss_scalar_buf → aux_outcome_ce_ema_update
  → ISV[538] → health_diag_isv_mirror → snap.aux_trade_outcome_ce
  → training_loop metrics → console.

Operator sees CE every epoch:
- Cold-start: 0.000
- After bootstrap: ~1.098 (= ln(3))
- After training: ideally 0.5-0.7 (head learning)

Phase F end-to-end ready. The vNext stack has:
- Full GPU kernel chain (A2-A5 + D + plan-conditioning)
- Full Rust wireup (B0-B4, B4b-1/2, C-1/C-2, B5b)
- Real labels reaching trainer
- K=3 → policy via state slots + atom-shift
- End-to-end CE observability

Verification:
- cargo check -p ml clean.
- cargo test -p ml --lib → 1016/0 green (incl. bumped
  snapshot_size_is_stable byte-size pin test).

Audit: docs/dqn-wire-up-audit.md Phase F-3c section.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-14 12:22:43 +02:00
jgrusewski
02479c885d feat(sp22-vnext): Phase F-3b — K=3 CE EMA launcher wireup + reset registry 
Completes the Phase F-3 observability chain. Kernel + ISV slot landed
in F-3a; this commit wires the launcher into the trainer's per-step
training graph and registers the FoldReset entry.

Changes:
- gpu_dqn_trainer.rs:
  - New pub(crate) static AUX_OUTCOME_CE_EMA_CUBIN embed
  - New aux_outcome_ce_ema_kernel: CudaFunction field on GpuDqnTrainer
  - Constructor loads kernel handle + struct-init
  - New launch_aux_outcome_ce_ema() method (single-thread launch,
    ISV slot 538, α=0.05)
- training_loop.rs:
  - Launch call appended after launch_aux_heads_loss_ema()
  - New dispatch arm "aux_trade_outcome_ce_ema" in reset_named_state
- state_reset_registry.rs:
  - New RegistryEntry for "aux_trade_outcome_ce_ema" (FoldReset
    sentinel 0.0)

End-to-end observability now live: K=3 head's batch-mean sparse-CE
flows into ISV[538] every step via Pearl A-bootstrapped EMA.
Smoke runs can read this slot to verify learning:
- Cold-start: 0.0
- After first step with B_valid > 0: bootstrap to first observation
  (~1.098 = ln(3) for uniform K=3 prediction)
- Healthy learning: monotonic decrease toward 0.5-0.7 over epochs
- Falsification: pinned at ~1.098 for many epochs

Phase F prep complete. The trade-outcome aux head's:
- Forward chain (A2-A5 + B0-B4)
- Real labels reach trainer (B4b-1/2)
- K=3 → policy state (C-1/C-2)
- K=3 → Q-target atom-shift (D)
- Plan-conditioning (B5b)
- Smoke observability (F-3 + F-3b)
are all wired. Phase F deployment (argo-train.sh smoke) is next.

Verification:
- cargo check -p ml clean.
- cargo test -p ml --lib → 1016/0 green (incl. every_fold_and_soft_
  reset_entry_has_dispatch_arm pin test).

Audit: docs/dqn-wire-up-audit.md Phase F-3b section.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-14 11:34:32 +02:00
jgrusewski
de64935b78 feat(sp22-vnext): Phase F-3 — K=3 CE EMA producer + ISV slot 
Adds observability infrastructure for Phase F smoke validation: new
ISV slot AUX_TRADE_OUTCOME_CE_EMA_INDEX = 538 tracks the K=3 head's
batch-mean sparse cross-entropy EMA. Producer kernel registered +
cubin-built; launcher wireup follows in Phase F-3b commit.

Changes:
- sp22_isv_slots.rs: new pub const AUX_TRADE_OUTCOME_CE_EMA_INDEX = 538
- gpu_dqn_trainer.rs ISV_TOTAL_DIM: 538 → 539
- NEW kernel aux_outcome_ce_ema_kernel.cu: single-thread single-block
  direct-to-ISV EMA writer with Pearl A first-observation bootstrap.
  Fixed α=0.05 (slow-moving observability). NULL-tolerant + NaN-guarded.
- build.rs: kernel registered; cubin compiles (3.2 KB)

Why dedicated kernel (not extension of aux_heads_loss_ema_update):
The K=2/K=5 EMA writes through Pearls A+D's 2-stage producer scratch.
Extending it would require allocating a new scratch slot, threading
Pearls A+D mapping, and touching apply_pearls_ad_kernel. The K=3
head's CE is OBSERVABILITY ONLY in this commit — no Pearls A+D
adaptive α needed. Minimum-scope direct-to-ISV path. Re-routable
through Pearls A+D if K=3 CE later becomes a controller anchor.

Smoke validation signal interpretation:
- Cold-start: ISV[538] = 0.0 (FoldReset sentinel)
- First step with B_valid > 0: Pearl A bootstrap → ISV[538] = first CE
- Typical uniform-K=3 cold-start CE: ln(3) ≈ 1.098
- Learning signal: CE drops from ~1.098 toward 0.5-0.7 over epochs
- Falsification signal: CE pinned at ln(3) for many epochs → head
  can't learn (label noise / under-capacity / outcomes unconditional)

Phase F-3b follow-up:
- Trainer launcher call after aux_heads_loss_ema_update
- HEALTH_DIAG snap layout extension + console column
- Reset registry entry (FoldReset sentinel 0)

Verification:
- cargo check -p ml clean.
- cargo test -p ml --lib → 1016/0 green.

Audit: docs/dqn-wire-up-audit.md Phase F-3 section.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-14 11:28:48 +02:00
jgrusewski
93aa4cd6a2 feat(sp22-vnext): Phase D — 12-weight W atom-shift (4 actions × 3 outcomes) 
THE K=3 HEAD NOW DIRECTLY MODULATES Q-TARGETS. Phase D extends the
Phase 3 atom-shift mechanism from per-action W[4] reading single state
slot 121 to per-(action, outcome) W[4, 3] = 12 weights reading 3 state
slots [121..124).

Mathematical change: shift[a, b] now sums Σ_k W[a*K+k] × state[121+k]
= W[a, 0]*p_Profit + W[a, 1]*p_Stop + W[a, 2]*p_Timeout.

5 atom-shift kernel sites updated (all coordinated):
1. experience_kernels.cu::compute_expected_q (replay path)
2. experience_kernels.cu::mag_concat_qdir (rollout path)
3. experience_kernels.cu::quantile_q_select
4. c51_loss_kernel.cu loss numerator (next_state CVaR side)
5. c51_loss_kernel.cu Bellman target (online + target combined)

Each site replaces `W[a] × state_121` with `Σ_k W[a*K+k] × state[121+k]`
unrolled 3 times. State hoist points lift 1 scalar → 3-element array.

5 supporting kernel updates:
- aux_w_prior_init_kernel.cu: writes 12 K=3 structural priors instead
  of 4. Spec prior matrix:
      Short × {Profit=+0.5, Stop=-0.5, Timeout=0}
      Hold  × {Profit= 0,    Stop=+0.5, Timeout=0}
      Long  × {Profit=+0.5, Stop=-0.5, Timeout=0}
      Flat  × {Profit= 0,    Stop=+0.5, Timeout=0}
  Block dim bumped 4 → 12.
- c51_aux_dw_kernel.cu: grid bumped (4,1,1) → (b0_size×K=12,1,1).
  blockIdx.x decoded as (a, k); reads state slot 121+k, writes
  dw_aux[a*K + k]. New kernel arg aux_outcome_k=3.
- adam_w_aux_kernel.cu: W_AUX_DIM 4 → 12. Block dim 12 threads.

Trainer-side buffer resizes:
  w_aux_to_q_dir   [4] → [12]
  adam_m_w_aux     [4] → [12]
  adam_v_w_aux     [4] → [12]
  dw_aux_buf       [4] → [12]

Rust launcher updates:
- c51_aux_dw_kernel launch: grid (4,1,1) → (12,1,1) + new aux_kto arg
- adam_w_aux_kernel launch: block (4,1,1) → (12,1,1)
- aux_w_prior_init launch: block (4,1,1) → (12,1,1)

Cold-start gracefulness preserved: state[121..124] = 0.0 at step 0
(no K=3 prediction yet from C-1 producer). Σ_k W[a*K+k] × 0 = 0 →
zero atom-shift across all actions. After step 1+ when C-1 producer
fires, real softmax probs activate the prior W's structural bias and
Adam refines from there.

End-to-end K=3 → Q-target chain now active:
  K=3 fwd (B3/B4) → softmax → C-1 producer → prev_aux_outcome_probs
  → C-2 state gather → state[121..124) → Phase D atom-shift
  → Q-target z_n + Σ_k W[a*K+k] × prob[k]
  → Bellman target + argmax + action_select all see aux's outcome
    prediction.

The K=3 head now influences policy via TWO paths: state input (Phase
C-2) AND Q-target modulation (Phase D).

Verification:
- cargo check -p ml clean.
- cargo test -p ml --lib → 1016/0 green.

Remaining vNext work:
- Phase E: dW backward gradient validation tests
- Phase F: Validation smoke at structural prior — decisive spec test
- B5b-2 (deferred): collector trade plan launch

Audit: docs/dqn-wire-up-audit.md Phase D section.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-14 11:20:56 +02:00
jgrusewski
d2331f2f62 feat(sp22-vnext): Phase B5b — full plan-conditioning integration 
The K=3 head's input is now [h_s2_aux (256) || plan_params (6)] = 262-
dim, matching the spec's intended architecture. Implements the input
concat in both the trainer's replay-batch path and the collector's
rollout-step path, with appropriate handling of the backward-side
stride mismatch.

NEW kernel strided_row_saxpy_kernel.cu: row-truncating SAXPY that
accumulates first n_cols_copy columns per row of src [B, src_cols]
into dst [B, dst_cols] scaled by alpha. Handles stride mismatch
(src_cols != dst_cols). Needed because backward emits dh_s2_aux_to_buf
[B, 262] but dh_s2_aux_accum [B, 256] only consumes first 256 cols.

PLAN_PARAM_DIM = 6 constant in gpu_aux_heads.rs.

Ops struct updates:
- AuxTradeOutcomeForwardOps gains concat_kernel + launch_concat()
- AuxTradeOutcomeBackwardOps gains strided_saxpy_kernel +
  launch_strided_row_saxpy()
- Both load new cubins in new()

Weight tensor resize:
- sizes[163] = H × (SH2 + PLAN_PARAM_DIM) = 128 × 262 = 33,536 floats
- fan_dims[163] = (H, SH2 + PLAN_PARAM_DIM)

Trainer changes:
- New aux_to_input_buf [B × 262] field
- aux_dh_s2_to_buf resized to [B × 262]
- aux_partial_to_w1 resized to [B × H × 262]
- max_aux_tensor_len bumped for param_grad_final scratch

Trainer forward (aux_heads_forward):
- Concat h_s2_aux + plan_params_buf → aux_to_input_buf
- forward() with SH2_TOTAL=262

Trainer backward (aux_heads_backward):
- backward() with SH2_TOTAL=262; reads aux_to_input_buf
- saxpy_f32_kernel SAXPY for dh_s2_aux REPLACED by
  launch_strided_row_saxpy: copies only first SH2=256 cols per row;
  trailing 6 cols (plan_params gradient) are dropped — STOP-GRAD on
  trade plan head from aux loss.

Collector forward (rollout):
- New exp_aux_to_input_buf [N × 262] field
- Concat with plan_params_ptr = 0 (NULL) → zero-fill trailing 6 cols
- forward() with SH2_TOTAL=262

Train/inference asymmetry (documented):
- Trainer: real plan_params from trade plan head output
- Collector: zeros (no trade plan launch in collector)
- The head is trained on real plan-conditional outcomes but queried
  at rollout time with plan_params=0. Phase B5b-2 follow-up would
  add a trade plan launch to the collector to resolve. Deferred —
  current state is functional, head still receives plan signal
  during training.

Stop-grad on plan_params: backward writes full [B, 262] gradient, but
strided SAXPY only copies first 256 cols. Trade plan head weights
NOT trained by K=3 aux loss in this commit.

Verification:
- cargo check -p ml clean.
- cargo test -p ml --lib → 1016/0 green.

Phase D next: 12-weight W atom-shift (4 actions × 3 outcomes).

Audit: docs/dqn-wire-up-audit.md Phase B5b section.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-14 10:53:23 +02:00
jgrusewski
53462a28d9 feat(sp22-vnext): Phase C-2 — state gather flip K=2 single-slot → K=3 3-slot 
THE K=3 HEAD NOW REACHES THE POLICY. Phase C-2 is the consumer-side
flip — slot 121's semantic changes from K=2's recentered p_up to K=3's
p_Profit, and slots [122, 123] gain new meaning as (p_Stop, p_Timeout).

Architectural constraint: aux_dir_prob_per_env (K=2 buffer) is ALSO
consumed by experience_env_step's β reward at lines 2335 + 3724-3725.
Cannot repurpose that pointer to point at the K=3 [N, 3] buffer —
env_step β consumer would read wrong-stride memory. This commit adds
SEPARATE arg aux_outcome_probs_per_env [N, 3] to state_gather kernels
with NULL fallback.

Branching semantic at state_gather read site:
- aux_outcome_probs_per_env != NULL → K=3 active path: writes slots
  [121..124) via new assemble_state_outcome_k3 helper
- aux_outcome_probs_per_env == NULL → K=2 fallback: writes slot 121
  via existing assemble_state from the recentered scalar

Changes:
- state_layout.cuh: NEW __device__ helper assemble_state_outcome_k3
  — mirrors assemble_state except padding slots [121..124) get
  p_Profit/p_Stop/p_Timeout (raw softmax probs [0, 1]), slots
  [124..128) zero for 8-alignment.
- experience_kernels.cu: training-side experience_state_gather +
  eval-side backtest_state_gather both get new trailing arg
  aux_outcome_probs_per_env (NULL-tolerant). Read site branches:
  K=3 reads 3 floats / env → assemble_state_outcome_k3; K=2 fallback
  preserves legacy assemble_state call.
- gpu_experience_collector.rs: training launcher passes
  self.prev_aux_outcome_probs.raw_ptr() → K=3 active in training.
- gpu_backtest_evaluator.rs: eval launcher passes NULL → K=2
  fallback in eval (eval has no aux producer infra yet).

K=2 head still alive:
- prev_aux_dir_prob still populated by aux_softmax_to_per_env_kernel
- experience_env_step still reads it for β reward (independent
  consumer untouched)
- EGF chain still reads exp_aux_nb_softmax_buf
- Only K=2's slot 121 contribution to policy state is suppressed

End-to-end K=3 chain now active:
  Label producer (A2) → per-(env, t) ring (B4b-1) → replay buffer
  scatter (B4b-2) → PER direct gather → trainer aux_to_label_buf
  → loss_reduce (B4) sparse CE on real labels → backward (B4)
  per-sample partials → Adam SAXPY (B1+B4) updates W1/b1/W2/b2 at
  [163..167)

  PARALLEL:
  Collector rollout K=3 forward (B3) → softmax tile → C-1 producer
  → prev_aux_outcome_probs [N, 3] → C-2 state gather → state[121..124]
  → policy reads in next step

The K=3 head closes the loop: learns from real labels via replay,
AND predictions reach policy via state assembly. Trainable +
observable.

Verification:
- cargo check -p ml clean.
- cargo test -p ml --lib → 1016/0 green.

Remaining vNext work:
- Phase D: 12-weight W atom-shift (4 actions × 3 outcomes)
- Phase E: dW backward + Adam for W[4, 3]
- Phase F: Validation smoke at structural prior
- B5b (deferred): plan_params input concat

Audit: docs/dqn-wire-up-audit.md Phase C-2 section.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-14 08:50:53 +02:00
jgrusewski
3286dc7dee feat(sp22-vnext): Phase C-1 — K=3 softmax → per-env 3-slot cache (producer) 
First half of Phase C. Lands the producer side of the K=3 trade-
outcome aux head's state bridge: new kernel populates a per-env
3-slot cache from the K=3 softmax tile every rollout step. The
consumer side (state gather reading from this cache → state slots
[121..124)) lands in Phase C-2.

Mirrors the K=2 head's existing aux_softmax_to_per_env_kernel exactly
at K=3:
- K=2: prev_aux_dir_prob[env] = 2*softmax[env, 1] - 1 (recentered)
- K=3: prev_aux_outcome_probs[env, k] = softmax[env, k] for k in [0, 3)

Changes:
- state_layout.rs: 3 new constants AUX_OUTCOME_PROFIT_INDEX = 121,
  AUX_OUTCOME_STOP_INDEX = 122, AUX_OUTCOME_TIMEOUT_INDEX = 123.
  PROFIT_INDEX aliases AUX_DIR_PROB_INDEX (same value, different
  semantic). Phase C-2 flips slot 121's meaning from K=2's recentered
  p_up to K=3's p_Profit.
- aux_outcome_softmax_to_per_env_kernel.cu: new kernel + cubin.
- gpu_dqn_trainer.rs: new SP22_AUX_OUTCOME_SOFTMAX_TO_PER_ENV_CUBIN
  embed.
- gpu_experience_collector.rs: 2 new struct fields (cache buffer +
  kernel handle); cubin load + alloc in constructor; struct-init;
  per-step launch in rollout loop after K=3 forward.
- build.rs: kernel registered.

Encoding shift K=2 → K=3: K=2 used recentered [-1, +1] to match
"no signal = 0" baseline of every other slot. K=3 keeps raw softmax
probabilities [0, 1]. Cold-start sentinel 0.0 for all 3 slots =
"no prediction yet" (mask). The 3-slot natural distribution is more
informative than a scalar.

Dead-code status: producer populates cache every step but
experience_state_gather doesn't read from it yet — state slot 121
still receives K=2's prev_aux_dir_prob write. Phase C-2 swaps the
state gather's source from K=2 cache to K=3 cache (3-slot write).

Why split C into C-1 + C-2: experience_state_gather is a hot-path
kernel with many consumers. Updating it touches training collector,
eval-side backtest evaluator, Rust launcher arg list. C-2 lands that
as an atomic state-semantic flip; C-1 lands the GPU-side scaffolding
independently so the producer chain can be validated first.

Verification:
- cargo check -p ml clean.
- cargo test -p ml --lib → 1016/0 green.

Audit: docs/dqn-wire-up-audit.md Phase C-1 section.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-14 02:26:39 +02:00
jgrusewski
68f0481a9e feat(sp22-vnext): Phase B5a — input concat kernel scaffolding 
Lands the plan-conditioning concat kernel for the K=3 trade-outcome
forward as reusable scaffolding. Phase B5b (integration) is deferred
with rationale: Phase C (state slots) is more critical for testing the
K=3 head's effect on policy behavior, and can land independently of
the plan-conditioning refinement.

NEW kernel aux_to_input_concat_kernel.cu:
- Writes [B, SH2+P] from h_s2_aux [B, 256] || plan_params [B, 6]
- Pure GPU map; one thread per output element, no atomicAdd
- NULL-tolerant on plan_params (zeros trailing P cols when source
  unavailable, e.g., collector cold-start where the trade plan head
  doesn't run)
- Registered in build.rs; cubin compiles (5.7 KB). Dead code at this
  commit — no Rust launcher yet.

Why B5 is split + B5b deferred:

Full Phase B5 (integration) requires three coordinated changes:
1. Forward path: bump aux_to_fwd.forward() to SH2=262 + 262-dim input
2. Backward stride mismatch: backward emits dh_s2_aux_to_buf [B, 262],
   but dh_s2_aux_accum (input to aux trunk backward) is [B, 256]. A
   direct SAXPY mismatches row strides (262 vs 256) and corrupts the
   trunk's upstream gradient. Needs a strided-SAXPY kernel.
3. Collector-path plan_params unavailability: trade plan head only runs
   trainer-side. Workarounds: zero-fill, add trade plan to collector,
   or skip K=3 forward in collector. All have trade-offs.

Phase B5b would need (1) strided-SAXPY kernel and (2) collector
plan_params decision. Real work but NOT on the critical path for
testing the K=3 head's effect on WR.

Why Phase C should land first:

The K=3 head currently trains on real labels (post-B4b) but doesn't
influence policy behavior. Phase C wires the head's softmax into state
slots [121..124) = (p_Profit, p_Stop, p_Timeout), replacing the K=2
single-slot 121 = 2*p_up - 1. WITH Phase C the policy reads aux's
outcome predictions as state features → behavior changes → testable.

Without Phase C, validation runs would show "K=3 head trains and
converges" but predictions don't reach the policy → WR signal isn't
a function of K=3 at all. We'd be testing nothing.

Recommendation: skip the full B5 for now, do Phase C next, then Phase
D (atom-shift). Phase B5b (plan-conditioning) is a refinement we add
IF Phase C/D's no-plan-params version shows promise but plateaus
below the WR ≥ 0.55 target.

Audit: docs/dqn-wire-up-audit.md Phase B5a section.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-14 02:19:27 +02:00
jgrusewski
da5e564ccf feat(sp22-vnext): Phase B4b-2 — replay-buffer scatter + trainer setter 
Second half of Phase B4b. Wires the trade-outcome label column through
the replay buffer (struct field, allocation, scatter on insert, gather
on sample, direct-to-trainer pointer + setter) and connects the
trainer's aux_to_label_buf to receive per-batch sampled labels.
Completes the end-to-end producer → ring → trainer i32 path the K=3
sparse CE consumer reads.

Replay buffer changes (crates/ml-dqn/src/gpu_replay_buffer.rs):
- New struct fields: aux_outcome_labels (capacity-sized ring),
  sample_aux_outcome_labels (mbs-sized fallback gather), trainer_aux_
  outcome_labels_ptr (direct-path destination)
- New aux_outcome_labels_ptr field on GpuBatchPtrs
- insert_batch signature: new aux_outcome arg between aux_conf_in and
  bs. Scatters via existing K-generic scatter_insert_i32 (same kernel
  the K=2 aux_sign_labels uses).
- sample_proportional direct gather when trainer_aux_outcome_labels_ptr
  != 0; fallback gather otherwise. Mirrors aux_sign_labels direct/
  fallback semantic exactly.
- New setter set_trainer_aux_outcome_labels_ptr mirrors
  set_trainer_aux_conf_ptr.

Collector emission:
- New aux_outcome_labels field on GpuExperienceBatch
- Populated at end of collect_experiences_gpu via dtod_clone_i32 from
  Phase B4b-1's per-(env, t) producer scratch.

Trainer wireup:
- aux_to_label_buf_ptr() accessor on GpuDqnTrainer (mirrors
  aux_nb_label_buf_ptr)
- trainer_aux_to_label_buf_ptr() delegating accessor on
  FusedTrainingCtx
- New set_trainer_aux_outcome_labels_ptr call in training_loop at the
  same site where set_trainer_buffers + set_trainer_aux_conf_ptr fire.
  Two call sites updated (lines ~835 + ~2857).
- insert_batch call in training_loop passes &gpu_batch.aux_outcome
  _labels as new arg.

Test fixtures updated: 4 smoke test files + 3 unit-test fixtures in
gpu_replay_buffer.rs alloc zero-init aux_outcome i32 arg.

End-to-end chain complete:
  trade_outcome_label_kernel (A2)
  → collector per-step launch (B3)
  → collector emission (B4b-1)
  → replay-buffer insert + scatter (B4b-2)
  → PER sample + direct gather (B4b-2)
  → trainer aux_heads_forward.loss_reduce (B4) reads sparse {-1,0,1,2}
  → trainer aux_heads_backward (B4) computes per-sample partials
  → Adam SAXPY (B1+B4) updates W1, b1, W2, b2 at [163..167)

The "degraded predict-Profit-everywhere" cold-start from Phase B4 is
resolved. K=3 head trains on real sparse trade-outcome labels.

Verification:
- cargo check -p ml clean (21 warnings, none new).
- cargo test -p ml --lib → 1016/0 on clean runs; pre-existing
  NoisyLinear flake still surfaces ~30-50% of runs (unrelated to
  vNext work — see ebc1b1502 / 20e1aea27 commit notes).

Remaining vNext work:
- B5: input concat 256 → 262 with plan_params
- Phase C: 3-slot state assembly (state[121..124])
- Phase D: 12-weight W atom-shift (4 actions × 3 outcomes)
- Phase E: dW backward + Adam for W[4, 3]
- Phase F: validation smoke at β=0.5 structural prior

Audit: docs/dqn-wire-up-audit.md Phase B4b-2 section.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-14 02:13:05 +02:00
jgrusewski
491bf7d3e6 feat(sp22-vnext): Phase B4b-1 — per-(env, t) label kernel output + collector buffer 
First half of Phase B4b (replay-buffer label scatter chain). Amends
the Phase A2 trade_outcome_label_kernel to emit a per-(env, t) output
column alongside the existing per-env tile, and adds the collector-side
buffer + per-step launch arg.

Kernel amendment (trade_outcome_label_kernel.cu):
- Added NULL-tolerant `out_labels_per_sample` arg after existing
  `out_labels`. When non-NULL, writes
  `out_labels_per_sample[env*L + t] = label` at the same offset as
  the `trade_close_per_sample[env*L + t]` read. NULL = no-op
  (preserves Phase A2/A3 contract for callers passing old signature).
- Pattern mirrors the K=2 head's `aux_sign_labels` per-(i, t) ring
  column that threads through the replay buffer.

Collector field + alloc + launch:
- New struct field `exp_aux_to_label_per_sample: CudaSlice<i32>`
  sized `[alloc_episodes × alloc_timesteps]`. Sentinel -1 (mask)
  populated by alloc_zeros + per-step kernel writes — survives
  until a trade-close event overwrites the env's slot at that t.
- Updated Phase B3 launcher in collect_experiences_gpu to pass
  `self.exp_aux_to_label_per_sample.raw_ptr()` as the new arg.

Why split B4b into B4b-1 + B4b-2: the full replay-buffer wireup
mirrors the K=2 head's aux_sign_labels pattern across ~8 distinct
code sites (replay-buffer struct field, sample destination buffer,
direct-to-trainer pointer, setter method, scatter on insert, gather
direct, gather fallback, GpuBatchPtrs field). Splitting lets us
validate the per-(i, t) producer in isolation before touching the
consumer pipeline.

B4b-1 (this commit) = producer chain complete. Per-(env, t) column
populated correctly every rollout step. Consumer wiring (replay-
buffer scatter + trainer setter) is B4b-2's scope.

Verification:
- cargo check -p ml clean (21 warnings, none new).
- cargo test -p ml --lib → 1016/0 on clean runs; pre-existing
  test_dqn_checkpoint_round_trip NoisyLinear flake still surfaces
  ~50-70% of full-suite runs (flake predates Phase B4b, unrelated
  to trade-outcome head — disable_noise() zeros ε but leaves some
  other randomness source intact).

Audit: docs/dqn-wire-up-audit.md Phase B4b-1 section.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-14 01:57:23 +02:00
jgrusewski
20e1aea27c feat(sp22-vnext): Phase B4 — trainer-side replay-batch chain wireup 
Wires the trade-outcome head's forward + loss reduce + backward +
per-sample partial reduce + SAXPY into the trainer's aux_heads_forward
and aux_heads_backward methods. Adam SAXPY for the 4 new weight tensors
at [163..167) extends uniformly via the existing aux_param_specs array
iteration.

Changes to aux_heads_forward:
- Steps 7 + 8 appended after K=2 next-bar head's loss reduce
- K=3 forward reads weights at [163..167) (Phase B1) → writes to
  aux_to_* save-for-backward tiles (Phase B2)
- K=3 loss_reduce writes aux_to_loss_scalar_buf + aux_to_valid_count_buf

Changes to aux_heads_backward:
- K=3 backward appended after K=5 regime backward, emits per-sample
  partials (dW1, db1, dW2, db2) + per-sample dh_s2_aux
- aux_param_specs array extended from 8 → 12 entries. Per-tensor
  reduce + SAXPY loop iterates uniformly, scaling each by aux_weight
- K=3 dh_s2_aux SAXPY appended after K=5's, all three heads'
  gradients flow into aux trunk's dh_s2_aux_accum (encoder stop-grad
  enforced structurally by aux_trunk_backward's missing dx_in output)

Label semantic (cold-start): aux_to_label_buf is alloc_zeros (all 0
= Profit) until Phase B4b lands replay-buffer label scatter. Model
trains on "predict Profit everywhere" — degraded but well-defined
(no NaN). Mirrors K=2 head's known-degraded state between B1.1a
(forward landed) and B1.1b (label producer wired).

Adam SAXPY: existing global SAXPY iterates 0..NUM_WEIGHT_TENSORS
(now 167) — 4 new weight slots get gradient SAXPYs followed by
Adam m/v updates uniformly. Architectural payoff of Phase B1's
NUM_WEIGHT_TENSORS bump.

Test flake mitigation: added bind_to_thread() to
ensemble::adapters::dqn::tests::shared_device() mirroring the
cuda_stream() test-helper pattern from the fix sweep at ebc1b1502.
test_dqn_checkpoint_round_trip had intermittent CUDA-context-thread-
state flakes under parallel test runs; the bind is idempotent and
forces context current on every test thread accessing the shared
device. Still occasionally non-deterministic at the prediction-
direction level (the test's disable_noise() zeros NoisyLinear
epsilon but may leave other randomness sources untouched; runs
alternate pred1=-1 pred2=1 ↔ pred1=1 pred2=-1). The underlying
NoisyLinear randomness has been flaky since pre-vNext; not a B4
regression.

Verification:
- cargo check -p ml clean (21 warnings, none new).
- cargo test -p ml --lib → 1016 passing / 0 failing.

Phase B4b next: replay-buffer label scatter populating trainer's
aux_to_label_buf from rollout's exp_aux_to_label_buf per (i, t).
Phase B5 (spec's actual "Phase B"): input concat 256→262 with
plan_params.

Audit: docs/dqn-wire-up-audit.md Phase B4 section.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-14 01:43:19 +02:00
jgrusewski
b28b349ac3 feat(sp22-vnext): Phase B3 — collector-side rollout buffers + forward chain wireup 
Adds collector-side trade-outcome head: 5 struct fields + allocations
+ per-step forward + per-step label producer launches in the rollout
loop. Mirrors the K=2 next-bar head's collector wireup at K=3.

Collector struct additions:
- exp_aux_to_fwd: AuxTradeOutcomeForwardOps (3 kernel handles)
- exp_aux_to_hidden_buf   [alloc_episodes × H=128] saved post-ELU
- exp_aux_to_logits_buf   [alloc_episodes × K=3]   saved logits
- exp_aux_to_softmax_buf  [alloc_episodes × K=3]   softmax tile
- exp_aux_to_label_buf    [alloc_episodes] i32     sparse {-1, 0, 1, 2}

Per-step launches in collect_experiences_gpu rollout loop:

1. aux_trade_outcome_forward — launched immediately after the K=2
   sibling's forward_next_bar, parallel on the same stream. Reads
   exp_h_s2_aux + weights at flat-buffer indices [163..167) (Phase
   B1 additions). Writes hidden/logits/softmax tiles. No consumer
   yet — Phase C wires state assembly; Phase B4 wires trainer
   scatter.

2. trade_outcome_label_kernel — launched immediately after
   experience_env_step on the same stream, reading the save-for-
   backward buffers (pnl_vs_target_at_close_per_env, pnl_vs_stop_at_
   close_per_env) that env_step just wrote at segment_complete.
   Stream-implicit producer→consumer ordering. Emits per-env
   {-1, 0, 1, 2} labels — sparse, ~95-99% bars produce -1 (mask).

Dead-code discipline per feedback_wire_everything_up: every kernel arg
+ producer site is real wiring (not NULL placeholder) — only the
absence of consumers reading the produced tiles is "dead". The smoke
run produces softmax tiles + labels every step bit-identical to
pre-vNext baseline (no consumer = no effect on training behavior).

Phase B4 next: trainer-side replay-batch chain (forward + loss_reduce
+ backward + Adam SAXPY for the 4 new weight tensors).

Audit: docs/dqn-wire-up-audit.md Phase B3 section.

Verification:
- cargo check -p ml clean (21 warnings, none new on aux_to_*).
- cargo test -p ml --lib → 1016 passing / 0 failing (unchanged from
  post-fix-sweep baseline at ebc1b1502).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-14 01:16:45 +02:00
jgrusewski
ebc1b15023 fix(tests): repair 14 pre-existing test failures across ml crate 
Lib test suite was at 14 failures from accumulated layout/contract drift.
Fixed each by tracing root cause; lib suite now 1016 passing / 0 failing.

Failures fixed (test → root cause → fix):

1. sp14_isv_slots::sp20_isv_slots_reserved_510_to_520 — ISV_TOTAL_DIM
   pin drifted from SP20-era 520 to current 538 via SP21/SP22 H6 bus
   growth. Slot positions 510-519 still pinned. Fix: relax total-dim
   check to >= 520; keep slot-position asserts tight.

2-3. gradient_budget::test_spectral_norm_all_heads_no_panic +
   test_spectral_norm_constrains_operator_norm — test config used
   cfg.state_dim=16 but trainer uses STATE_DIM=128 when bottleneck
   off; also DuelingWeightBacking slices [2]/[3] used pre-GRN w_s2
   shape, but post-GRN it's w_b_h_s1 [2*SH1, SH1] = 2048. Fix: add
   s1_input_dim_for_test helper; correct slice sizes to GRN layout.

4-9. dqn::trainer::tests::test_feature_vector_to_state +
   test_single_sample_batch + test_batched_action_selection +
   test_batched_vs_sequential_action_selection_consistency +
   test_batch_size_mismatch_larger/smaller_than_configured —
   feature_vector_to_state wrapper falsely advertised "no OFI" by
   signature but body required strict OFI. Contract violation —
   broke 6 unit tests + hyperopt's public convert_to_state APIs.
   Fix: wrapper now genuinely produces 45-dim no-OFI state; OFI-
   strict callers use _with_ofi with Some(idx).

10. state_kl_monitor::observe_tracks_fire_rate_on_change — last_amp
    initialized to 1.0 coincidentally equaled first observation value,
    silently suppressing first fire. Test expected first observation
    always fires (cold-start semantic). Fix: Option<f32> sentinel —
    None means "no prior baseline" → first observe ALWAYS fires.

11. cuda_pipeline::tests::test_eval_action_select_thompson_picks_
    proportionally — test launched experience_action_select kernel
    with 1 missing arg (v_logits_dir from SP17 Commit C). Kernel
    read garbage pointer → SIGSEGV. Also threshold 0.70 calibrated
    for pre-SP17 raw-A distribution; post-SP17 sampler uses mean-zero
    softmax(V + A_centered), Long still dominates ~4× but P(Long)
    drops to ~0.66. Fix: add zero-filled v_logits_buf at correct
    arg position; recalibrate threshold 0.70 → 0.60.

12. cuda_pipeline::tests::test_ppo_gpu_data_upload — flaked in parallel
    test run with CUDA_ERROR_INVALID_CONTEXT from MappedF32Buffer's
    cudaHostAlloc. cuda_stream() test helper used OnceLock to share
    a CudaStream but didn't bind_to_thread on subsequent calls. CUDA
    contexts are per-thread state. Fix: helper now calls bind_to_
    thread() on every invocation (idempotent — mirrors trainer/
    constructor.rs:111's pattern).

13. ppo::tests::test_reward_computation — test created hold action
    with ExposureLevel::Flat but is_hold() matches only Hold (Flat
    = close position = transaction with cost). So hold and buy got
    same transaction-cost reduction; reward_hold > reward_buy assert
    failed. Fix: use ExposureLevel::Hold for no-cost hold semantic.

14. training_profile::tests::test_production_profile_applies_all_
    sections — pinned n_steps==5 and tau==0.005 (pre-TD(0)); production
    flipped to n_steps=1, tau=0.01 per dqn-production.toml ("dense
    micro-rewards cancel over n>1 bars; faster target tracking for
    TD(0)"). Fix: update pins to current production values.

All 14 fixes target root causes — no #[ignore] masking. Verified:
- Lib-only run: cargo test -p ml --lib → 1016/0 (passed/failed)
- Stash-test pre-changes: 1001/15 (confirms 14 fixed atomically)

Remaining flake under `cargo test -p ml --tests`:
- ensemble::adapters::dqn::tests::test_dqn_checkpoint_round_trip —
  CUDA-context-race-adjacent flake under parallel `--tests` mode;
  passes in isolation and in `--lib` mode (single binary). Predates
  this commit; deferred as separate triage.

Audit: docs/dqn-wire-up-audit.md "Fix sweep" section.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-14 01:08:16 +02:00
jgrusewski
b98dc2730d feat(sp22-vnext): Phase B2 — trade-outcome trainer saved-tensor + partial buffers 
Adds 11 buffer fields + 2 orchestrator ops handles (fwd + bwd) to the
trainer struct, mirroring the existing aux_nb_* / aux_partial_nb_*
pattern at K=3 instead of K=2.

Trainer struct additions:
- aux_to_fwd: AuxTradeOutcomeForwardOps   (Phase B0 scaffold)
- aux_to_bwd: AuxTradeOutcomeBackwardOps
- aux_to_hidden_buf       [B, H=128]      saved post-ELU
- aux_to_logits_buf       [B, K=3]        saved logits
- aux_to_softmax_buf      [B, K=3]        saved softmax (3 future consumers)
- aux_to_label_buf        [B] i32         sparse {-1, 0, 1, 2}
- aux_to_loss_scalar_buf  [1]              mean CE
- aux_to_valid_count_buf  [1]              B_valid for backward
- aux_dh_s2_to_buf        [B, SH2]        SAXPYs into dh_s2_aux_accum
- aux_partial_to_w1       [B, H, SH2]     per-sample dW1
- aux_partial_to_b1       [B, H]          per-sample db1
- aux_partial_to_w2       [B, K=3, H]     per-sample dW2
- aux_partial_to_b2       [B, K=3]        per-sample db2

Memory: aux_partial_to_w1 = 256 MB at B=2048 — identical to K=2 head's
partial size (same SH2, same H). Total new aux-to footprint ≈ 260 MB.

The existing aux_param_grad_final_buf scratch is sized to the largest
tensor across all aux heads; trade-outcome head's largest is W1 [H, SH2]
= 32,768 floats — identical to K=2/K=5 W1s. No resize needed.

Cold-start label semantics: alloc_zeros yields label 0 (Profit) for
every sample. Until the producer wires in (B3), the trainer's CE loss
treats every sample as "should have predicted Profit" — degraded but
well-defined (no NaN). Mirrors the K=2 head's known-degraded state
between B1.1a and B1.1b.

No FoldReset registration: these buffers are overwritten every batch
— no stale-state-leak risk across folds (matches the existing aux_nb_*
pattern).

Phase B3 next: collector-side rollout buffers + forward chain wireup
into collect_experiences_gpu (per-env softmax → per-(i, t) fan-out
scatter for trainer's aux_to_softmax_buf population).

Audit: docs/dqn-wire-up-audit.md Phase B2 section.
Cargo check clean (21 warnings, none new on aux_to_* fields).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-14 00:37:56 +02:00
jgrusewski
205e46c171 feat(sp22-vnext): Phase B1 — trade-outcome head weight tensors + Xavier init 
Adds the 4 weight tensors (W1, b1, W2, b2) for the SP22 H6 vNext
trade-outcome aux head into the trainer's flat params_buf at indices
[163..167). Adam machinery (m/v moment buffers, SAXPY iteration over
0..NUM_WEIGHT_TENSORS) picks up the new tensors uniformly — no
per-tensor wiring needed.

Changes:
- NUM_WEIGHT_TENSORS bumped 163 → 167. Most of the 54 references are
  &[u64; NUM_WEIGHT_TENSORS] array-size generics that resize
  uniformly with the constant.
- compute_param_sizes() adds 4 new size entries:
    [163] aux_to_w1 [H=128, SH2=256]  = 32,768 floats
    [164] aux_to_b1 [H=128]            = 128 floats
    [165] aux_to_w2 [K=3, H=128]       = 384 floats
    [166] aux_to_b2 [K=3]               = 3 floats
  Total: 33,283 floats = ~133 KB params, ~266 KB Adam state.
- compute_param_sizes() debug_assert updated 163 → 167.
- Xavier fan_dims added: (H, SH2) for W1, (0, 0) for biases (zero-init),
  (K=3, H) for W2. Cold-start: logits ≈ 0 → softmax ≈ uniform 1/3 → no
  Profit/Stop/Timeout preference per pearl_first_observation_bootstrap.

SH2 stays at 256 in this commit (mirrors K=2 head exactly). The spec's
Phase B input concat (256 → 262 with plan_params 6-dim) will re-shape
slot [163] to 128 × 262 = 33,536 floats later — small touch-up vs the
full B1 commit.

Verification:
- cargo check -p ml clean (21 warnings, none new).
- 14 pre-existing test failures stashed-verified unrelated (OFI features
  missing, ISV slot count drift — independent of NUM_WEIGHT_TENSORS).

Phase B2 next: saved-tensor + per-sample partial buffers (hidden_post,
logits, softmax, valid_count, dW*_partial, dh_s2_aux_out).

Audit: docs/dqn-wire-up-audit.md Phase B1 section.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-14 00:33:38 +02:00
jgrusewski
108a426c38 feat(sp22-vnext): Phase B0 — AuxTradeOutcome ops struct scaffolding 
First Rust-side commit of Phase B for the SP22 H6 vNext trade-outcome
aux head. Pure orchestrator scaffolding mirroring AuxHeadsForwardOps /
AuxHeadsBackwardOps. Zero production callers — additive per the
gpu_grn / aux_trunk commit-by-commit ordering convention (scaffold
→ trainer fields → collector wireup).

Adds:
- gpu_dqn_trainer.rs: 4 cubin embeds (TRADE_OUTCOME_LABEL_CUBIN,
  AUX_TRADE_OUTCOME_FORWARD_CUBIN, AUX_TRADE_OUTCOME_LOSS_REDUCE_CUBIN,
  AUX_TRADE_OUTCOME_BACKWARD_CUBIN). All #[allow(dead_code)] until B1+.
- gpu_aux_heads.rs: AUX_OUTCOME_K = 3 constant (parallel to AUX_NEXT_BAR_K).
- gpu_aux_heads.rs: AuxTradeOutcomeForwardOps struct holding 3 kernel
  handles (forward + loss_reduce + label producer). Launch methods:
  forward(), loss_reduce(), compute_label(). Per-env label kernel uses
  grid ceil(n_envs/256) — pure per-env map, no reduction.
- gpu_aux_heads.rs: AuxTradeOutcomeBackwardOps struct holding 1 kernel
  handle (backward). Launch method: backward(). Reuses existing
  K-generic aux_param_grad_reduce from AuxHeadsBackwardOps — no
  separate reducer needed.

Contract shapes mirror AuxHeadsForwardOps/AuxHeadsBackwardOps so
subsequent wireup commits plug in with minimal contract drift. Phase B
proper (input concat 256→262 with plan_params) becomes a small change
touching only the buffer fill + W1 shape once B1-B4 land the wireup.

Phase B1 next: trainer struct fields for W1/b1/W2/b2 + Adam state +
Xavier init + reset registry entries.

Audit: docs/dqn-wire-up-audit.md Phase B0 section.
Cargo check clean (21 warnings, none new).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-14 00:06:03 +02:00
jgrusewski
9f4a25e623 feat(sp22-vnext): Phase A5 — aux_trade_outcome backward kernel (Phase A complete) 
K=3 backward kernel that closes the forward → loss → backward chain for
the trade-outcome aux head. Mirrors `aux_next_bar_backward` (K=2 sibling)
line-for-line because gradient flow is K-independent: `d_logits = (softmax
− one_hot)/B_valid` propagated through `Linear → ELU → Linear` chain via
standard softmax-CE derivative.

Per-sample partials (caller reduces via existing K-generic `aux_param_
grad_reduce` kernel):
  dW1_partial [B, H=128, SH2=256], db1_partial [B, H]
  dW2_partial [B, K=3, H],         db2_partial [B, K=3]
  dh_s2_aux_out [B, SH2]

Mask handling: labels[b] == -1 zeros the K-vector → all downstream
partials zero (chain rule's multiplicative zero). All-skip batch produces
valid_count=0 → d_logits=0 for every row → zero gradients across the
board, no NaN.

Sparse-label gradient amplification: B_valid is typically ~1-5% of
nominal batch (trade-close events are rare), so inv_B = 1/B_valid is
much larger than the K=2 sibling's inv_B = 1/(~B). Per-trade-close
gradients have proportionally higher magnitude — correct credit
assignment (rare signal speaks louder) but Phase E's Adam may need
class-weighted CE or per-group LR tuning.

ELU backward via post-activation identity: f'(x) = (h_post > 0) ? 1 :
1 + h_post — recovers derivative without re-evaluating x_pre.

SP14 Phase C.5b separation preserved: reads h_s2_aux (aux trunk output),
writes dh_s2_aux_out SAXPYing into dh_s2_aux_accum. Q's encoder
structurally protected (aux_trunk_backward has no dx_in output).

Phase A5 (this commit) is dead code — no Rust launcher. Phase B will
land the full launcher chain (gpu_aux_heads.rs parallel ops struct,
collector struct fields for W1/W2/b1/b2/Adam-state/saved-tensors/dW-
partials, wireup in collect_experiences_gpu).

Cubin: aux_trade_outcome_backward_kernel.cubin (24.8 KB).

═══ Phase A complete ═══

A1: ISV slots (none needed — reuses padding 121-123)
A2: trade_outcome_label_kernel.cu (label producer)            26ce7ba69
A3: aux_trade_outcome_forward + save-for-backward wireup       07728f9ef
A4: aux_trade_outcome_loss_reduce_kernel.cu (sparse CE)        3ddcfb886
A5: aux_trade_outcome_backward_kernel.cu (this commit)

All 5 GPU kernels compile, cubins built, build.rs registered. Contract
chain composes end-to-end on the GPU side. Phase B (Rust launcher
chain + 262-dim input concat) is next.

Audit: docs/dqn-wire-up-audit.md Phase A5 section + Phase A summary.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-14 00:01:39 +02:00
jgrusewski
3ddcfb8868 feat(sp22-vnext): Phase A4 — aux_trade_outcome loss reduce kernel 
K=3 sparse cross-entropy reduce over the trade-outcome softmax tile
produced by `aux_trade_outcome_forward` (Phase A3). Mirrors the K=2
sibling `aux_next_bar_loss_reduce` structurally: single-block shmem-tree
reduce, two parallel partial strips (loss_numer + valid_count) reduced
lockstep, fmaxf(p_tgt, 1e-30) numerical floor, fmaxf(valid, 1.0) all-
skip-batch guard, valid_count_out[1] save-for-backward.

Kept as SEPARATE kernel from the K=2 sibling:
- Diagnostic isolation (distinct HEALTH_DIAG slot, distinct cubin in
  profiles for clean per-loss-source attribution)
- Sparse-label semantic clarity (~95-99% mask=-1 vs ~50-100% valid for
  the K=2 next-bar head)
- Future per-class weighting headroom (Profit/Stop/Timeout 3:1-10:1
  imbalance will likely need class-weighted CE — surgical mod here
  without touching the K=2 head's contract)

Phase A4 (this commit) is dead code — no Rust launcher yet. Phase A5
lands backward; Phase B wires the full forward→loss→backward chain.

Discipline: feedback_no_atomicadd (single-block tree-reduce), feedback_
cpu_is_read_only (pure GPU), pearl_first_observation_bootstrap (sentinel
0 valid_count produces zero gradients gracefully on cold start).

Audit: docs/dqn-wire-up-audit.md Phase A4 section.
Cubin: aux_trade_outcome_loss_reduce_kernel.cubin (9.9 KB) compiles clean.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-13 23:41:59 +02:00
jgrusewski
07728f9efc feat(sp22-vnext): Phase A3 — aux_trade_outcome forward kernel + save-for-backward wireup 
Phase A3 of the SP22 H6 vNext trade-outcome aux head (per
docs/plans/2026-05-14-sp22-h6-vNext-trade-outcome-aux.md). Two atomic pieces:

1. NEW kernel `aux_trade_outcome_forward_kernel.cu` — K=3 softmax aux head
   forward (Linear → ELU → Linear → stable softmax). Mirrors
   `aux_next_bar_forward` (K=2) but emits {Profit, Stop, Timeout} probs.
   Saved tensors {hidden_out, logits_out, softmax_out} ready for A4 (loss
   reduce) and A5 (backward). Dead code at this commit — no Rust launcher
   yet. Registered in build.rs::kernels_with_common, cubin verified.

2. Save-for-backward buffers `pnl_vs_target_at_close_per_env` +
   `pnl_vs_stop_at_close_per_env` ([alloc_episodes] f32 device-resident).
   Producer: `experience_env_step::segment_complete` writes the trade's
   realized P&L ratios vs profit_target / stop_loss at close (inline-
   computed from `pre_trade_position × (raw_close − entry_price) /
   (ps[PS_PLAN_PROFIT_TARGET] × prev_equity)`, symmetric-clamped to
   [-2, +2] per pearl_symmetric_clamp_audit — same formula as the
   sibling experience_state_gather's plan_isv[PLAN_ISV_PNL_VS_TARGET/_STOP]
   slots). Consumer (eventual A4/A5 wireup): trade_outcome_label_kernel
   classifies each close into {Profit, Stop, Timeout} via the ≥1.0
   threshold-hit predicate.

Wireup discipline per feedback_registry_entries_need_dispatch_arms:
- New struct fields on GpuExperienceCollector
- stream.alloc_zeros at construct site
- Kernel-launch .arg() threading at experience_env_step launch
- StateResetRegistry entries (FoldReset sentinel 0.0)
- training_loop::reset_named_state dispatch arms
- All 10 registry pin tests pass including
  every_fold_and_soft_reset_entry_has_dispatch_arm

Audit doc updated: docs/dqn-wire-up-audit.md Phase A3 section.

Next phases (per spec): A4 = aux_trade_outcome_loss_reduce (sparse CE,
mask=-1), A5 = aux_trade_outcome_backward, Phase B = 262-dim input
(h_s2_aux || plan_params), Phase C = 3-slot state assembly, Phase D =
12-weight W atom-shift, Phase E = dW + Adam, Phase F = validation smoke.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-13 23:20:29 +02:00
jgrusewski
26ce7ba690 feat(sp22-vNext): Phase A2 — trade-outcome label producer kernel 
First foundation kernel for the H6 vNext trade-outcome aux head.

Per-env classification at trade-close events into K=3 outcomes:
- 0 = Profit (pnl_vs_target >= 1.0)
- 1 = Stop (pnl_vs_stop >= 1.0)
- 2 = Timeout (neither threshold hit)

Sparse labels — most bars get -1 mask. Priority: Profit > Stop > Timeout.

Pure per-env map; no atomicAdd, no reduction. Launch: grid(ceil(N/256)),
block(256).

Registered in build.rs. Cubin compiles clean. Currently dead code —
launcher wireup comes in Phase A3+ commits.

See docs/plans/2026-05-14-sp22-h6-vNext-trade-outcome-aux.md for the
full vNext architecture.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-13 21:55:34 +02:00
jgrusewski
d0c037a3d2 feat(sp22): H6 Phase 3 FALSIFIED + vNext spec (trade-outcome aux head) 
Decisive smoke train-xrkb7 @ ebc7144434: WR=0.4346 with full Phase 3
mechanism + enlarged aux head. Statistically identical to all baselines
(0.4338-0.4346). Hypothesis falsified.

Root cause investigation:
- Aux head's 70% accuracy was mostly majority-class prediction on
  imbalanced data (fold 0: 88% down labels)
- Fold 1 with balanced labels: aux accuracy dropped to 52% (near-random)
- Aux head has minimal per-bar discriminative power

Architectural insight: aux head predicts next-bar direction but the
system makes MULTI-BAR TRADE decisions with target_bars / profit_target
/ stop_loss / conviction plan parameters. The decision horizons mismatch.

This commit:
1. Revert mechanism to dormant (W=0, beta=0) - hypothesis falsified
2. Preserve infrastructure (kernels, NaN guards, enlarged aux head)
3. Write vNext spec: trade-outcome aux head with plan_params concat
   input + K=3 outputs {Profit, Stop, Timeout} + 12-weight W matrix

vNext reuses ~80% of current infrastructure. ~2-3 days focused work.
See docs/plans/2026-05-14-sp22-h6-vNext-trade-outcome-aux.md.

Cargo check clean.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-13 21:50:32 +02:00
jgrusewski
ebc7144434 feat(sp22): H6 Phase 3 RE-ACTIVATED with corrected aux head 
Smoke train-8zwtf @ 465abc7e9 validated that the AUX_HIDDEN_DIM 32→128
uplift fixed the aux head:
- aux_dir_acc: 0.28 (anti-predictive) → 0.70 (correctly predicting
  majority-down direction at H=60 bars)
- pred_tanh: +0.66 (UP-biased, wrong) → -0.52 (DOWN-biased, correct)
- val WR: 0.4345 (baseline preserved, no destabilization)

With aux head now informative, re-activate Phase 3 priors:
- aux_w_prior_init: W = [-0.5, 0, +0.5, 0]
- state_reset_registry dispatch: scale_beta = 0.5

When state_121 < 0 (aux predicts down, typical):
- atom_shift[Short] = -0.5 × neg = POSITIVE → encourages Short ✓
- atom_shift[Long]  = +0.5 × neg = NEGATIVE → discourages Long ✓

Next smoke is decisive H6 Phase 3 test: does mechanism move WR above
0.4345 baseline with the now-informative aux head?

Cargo check clean.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-13 21:03:02 +02:00
jgrusewski
465abc7e9b fix(sp14): dial back aux head/trunk enlargement after L40S OOM 
Initial 8x head + 2x trunk-H2 (commit 787ee7b86) OOM'd all 3 folds on
L40S with alloc next_states f32[1048576000] failure. Root cause: aux
backward stores [B, H, SH2] partial grad buffers per sample. At
H=256, SH2=256, B=16384: 4GB per head x 2 heads + 4GB trunk = +12GB
on top of existing DQN buffers, exceeding L40S 48GB budget.

Dial back:
- AUX_HIDDEN_DIM: 256 -> 128 (still 4x prior 32-unit bottleneck)
- AUX_TRUNK_H2: 256 -> 128 (back to original mild bottleneck)

At H=128: partial buffer = 2GB per head = 4GB total. Fits comfortably
alongside next_states ~4GB and other DQN buffers.

Cargo check clean. Smoke re-dispatch next.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-13 20:23:28 +02:00
jgrusewski
787ee7b86c feat(sp14/sp22): enlarge aux head + trunk capacity (8x head, 2x trunk-H2) 
After Path C confirmed aux head's 28% accuracy at H=60 is the H6 Phase 3
bottleneck (not the mechanism itself), enlarge aux capacity:
- AUX_HIDDEN_DIM: 32 -> 256 (8x, matches input dim, removes bottleneck)
- AUX_TRUNK_H2: 128 -> 256 (2x, uniform trunk width)

Architecture changes:
- Aux head: 256 -> Linear -> 256 -> ELU -> Linear -> 2 (was 256->32->2)
- Aux trunk: 256 -> 256 -> 256 -> 256 (was 256 -> 256 -> 128 -> 256)
- +160K params total (mostly aux_nb_w1 + aux_rg_w1: [256, 256] each)

Side effects:
- Checkpoint fingerprint change (intentional)
- Thread utilisation improves: AUX_BLOCK=256 threads x H=256 = 1:1
  (vs 8:1 at H=32 — most threads idle previously)

Phase 3 mechanism stays DORMANT (W=0, beta=0) for this validation
smoke. Verdict criteria: aux_dir_acc improves from 0.28 toward 0.50+
with the larger capacity. If yes, re-activate Phase 3 priors. If no,
the bottleneck is signal/horizon, not capacity.

Cargo build clean (full nvcc rebuild).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-13 19:59:16 +02:00
jgrusewski
2c0911981a feat(sp22): H6 Phase 3 DORMANT - mechanism falsified, infrastructure preserved 
Path C investigation revealed the aux head at epoch 1 is severely
anti-predictive at H=60 bars:
- Aux predicts UP 83% of the time
- Labels are 17% UP, 83% DOWN
- Accuracy = 28% (vs 50% random)

This means H6 Phase 3 hypothesis cannot help WR — atom-shift on an
anti-predictive signal produces no discriminative bias. Confirmed by
full-mechanism smoke at b4e26a3b4 producing WR=0.4337 (statistically
identical to dormant baseline 0.4338).

Path D: revert mechanism to dormant state:
- aux_w_prior_init_kernel: W = [0, 0, 0, 0]
- state_reset_registry dispatch: scale_beta = 0.0

Infrastructure preserved (kernels, NaN guards, Step 8/11, Phase C1
setter, scratch buffers, dispatch arms). Re-activation requires fixing
aux head's directional prediction quality first.

Next-step options documented in audit doc:
1. Shorter horizon (H=1 or H=8)
2. Regression target instead of binary classification
3. Direct market features into aux head
4. Multi-epoch aux-only training
5. Bigger aux trunk
6. Use aux as confidence gate only (not direction signal)

Cargo check clean.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
 2026-05-13 19:50:26 +02:00