Commit Graph

5197 Commits

Author SHA1 Message Date
jgrusewski
a4127935a8 feat(ml-backtesting): BacktestHarness orchestrator + Ring 1b parity (C8)
BacktestHarness::new(cfg, &dev, trunk) constructs a MultiHorizonLoader
(inference_only=true), captures the trunk's perception Graph A using
loader.peek_first() as the template, then allocates a LobSimCuda.
run() walks the chronological snapshot stream, calls apply_snapshot on
every event, and at decision-stride boundaries:
  trunk.update_input_buffers(raw)
  → trunk.perception_forward_captured() → (probs[N_HORIZONS], proj)
  → sim.broadcast_alpha(&probs)
  → sim.step_decision(ts, target_vol, ann_factor, max_lots)

Returns RunStats { events_processed, decisions_taken }. The harness
deliberately accepts an externally-built CfcTrunk (random-init in v1)
because a checkpoint format isn't pinned in ml-alpha yet; when one
lands, the binary CLI (C9) can switch from new_random to load_checkpoint.

trainer_parity.rs Ring 1b — two ignored tests:

  peek_first_byte_equal_across_modes
    Verifies the Mbp10RawInput produced by the loader path used by the
    backtest harness (inference_only=true) is BYTE-EQUAL to what the
    trainer's loader (inference_only=false) produces from the same
    source. Guards against any future refactor accidentally diverging
    the two paths (e.g. someone special-casing inference path to skip
    regime feature computation). All 50 f32 fields + scalars compared
    via .to_bits() equality.

  inference_iteration_matches_chronological_snapshots
    Verifies next_inference_input() yields monotone-ts snapshots with
    correct cur==prev semantics on the first read and prev_ts==prior_ts
    afterwards.

Both tests skip gracefully when FOXHUNT_TEST_DATA fixtures lack
populated sidecars (the placeholder-empty bins committed in the
test_data/ tree).

GPU-side inference parity (probs[N_HORIZONS] bit-equal across loader
modes) is deferred until ml-alpha pins a checkpoint format and we have
a small checkpoint fixture to gate it on FOXHUNT_TEST_CKPT.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-18 08:52:49 +02:00
jgrusewski
25f43a4268 feat(ml-backtesting): decision_policy + per-horizon ISV-Kelly (C7)
Two new kernels in cuda/decision_policy.cu:

  decision_policy_default — alpha[N_HORIZONS] × per-horizon IsvKellyState
  → market target. Per-horizon Kelly fraction × signal magnitude × ISV-cap
  (target_annual_vol / sqrt(realised_return_var × annualisation_factor)),
  aggregated via WeightedByRealizedSharpe (weights = max(0, recent_sharpe)
  / Σ; auto-shifts capital toward empirically winning horizons per spec §5
  + §6). No floor — sub-1-lot intents become no-op. Round-to-nearest on
  the final lot count to dodge an f32-truncation off-by-one where
  0.8(f32) * 0.75 * 5.0 ≈ 2.99999976 → trunc-int 2.

  isv_kelly_update_on_close — for every horizon flagged in
  closed_horizon_mask[b], updates pnl_ema_{win,loss} via Wiener-α (floor
  0.4 per pearl_wiener_alpha_floor_for_nonstationary), win_rate_ema,
  Welford-ish realised_return_var, and the recent_sharpe composite.
  First-observation bootstrap (per pearl_first_observation_bootstrap):
  sentinel n_trades_seen=0 → direct EMA replacement, no zero-bias warmup.

The full bytecode VM from spec §6 is NOT in this commit — the default
policy is hardcoded in the kernel as the path Strategy::default_for()
already produces. Bytecode plumbing in src/policy/mod.rs stays put for
v2 expansion (custom RegimeSwitch / Portfolio compositions).

IsvKellyState struct added to lob_state.cuh (24 bytes per horizon × 5
per backtest); host mirror IsvKellyStateHost from C3 cast-compatible
via bytemuck::Pod. LobSimCuda gains broadcast_alpha + step_decision +
read_isv_kelly + write_isv_kelly (warm-start). step_decision chains:
  decision_policy → merge_open_mask → submit_market_immediate
  → pnl_track_step → host close-detect → isv_kelly_update_on_close.

PRE-submit pos/pnl/mask snapshots feed the host close-detection;
captured via three small DtoH copies (cold path, 24 bytes × n_backtests).

decision_alpha_buy_close fixture: warm-start h4 with positive Kelly
state (n=50, recent_sharpe=1.0), broadcast alpha[4]=0.9 → buy 3 lots
@ ask top 5500.00. Snapshot moves to bid 5505.00, broadcast alpha[4]=0.1
→ sell 3 lots → close at +15 P&L. Verify ISV-Kelly h4: n_trades 50→51
exactly, others unchanged. PASS — all 6 Ring 1 fixtures green on RTX 3050.

Out-of-scope for this commit (defer to follow-ups, per plan trim notes):
  - stop_trigger / oco_one_cancels_other / submission_overflow fixtures
    (need resting-order LimitSlot[32] machinery deferred from C5)
  - Bytecode VM dispatch (RegimeSwitch / Portfolio compositions)

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-18 08:50:37 +02:00
jgrusewski
1b679b5e40 feat(ml-backtesting): pnl_track kernel + segment_complete TradeRecord emission
pnl_track_step runs after each matching pass, compares per-block
position-state-now against persisted OpenTradeState (24 B scratch)
and either:
  - records entry context (entry_ts_ns, entry_px_x100, entry_size,
    realised_at_open) on open transition (prev==0, now!=0); or
  - emits a 40-byte TradeRecord into the per-block trade-log buffer
    on close transition (prev!=0, now==0), reconstructing implied
    exit_px from the realized P&L delta and converting to USD ×100
    fixed-point ($50/index-point × 100 = ×5000 multiplier).

Multi-fill averaging (scale-in then partial close) deferred to v2 —
v1 covers the clean open→close case the spec calls out as primary.

LobSimCuda owns three new buffers: open_trade_state_d (n × 24),
trade_log_d (n × TRADE_LOG_CAP × 40), trade_log_head_d (n × u32).
submit_market now takes current_ts_ns and chains pnl_track_step
internally; step_pnl_track() exposed for caller-driven orchestration.

read_trade_records(backtest_idx) drains the per-block ring as
Vec<TradeRecord>; LSP-pinned 40-byte Pod struct from C2 lines up
1:1 with the kernel's hand-rolled byte writes.

pnl_accounting_buy_close fixture: buy 4 lots @ ask top (5500.00),
book moves +5 to bid top 5505.00, sell 4 to close. Expected
realized_pnl = (5505 − 5500.00) × 4 = $20 in price-units, which is
$20 × $50/contract × 100 = 100000 USD ×100 fixed-point. PASS within
$1 fixed-point tolerance.

All 5 Ring 1 fixtures green on RTX 3050.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-18 08:37:24 +02:00
jgrusewski
78f0618294 feat(ml-backtesting): order_match immediate market-fill kernel + Pos accounting
submit_market_immediate walks the ask/bid book for a buy/sell market
order, computes notional cost + filled lots across up to 10 levels,
and updates per-backtest Pos { position_lots, vwap_entry, realized_pnl }
with correct VWAP scale-in + counter-direction realized-P&L accounting.
Single-writer (thread 0) per block — no atomics per
feedback_no_atomicadd.md. Each backtest gets its own target slot
{ side, size } in the device-global targets array (side=2 = no-op).

Pos struct added to lob_state.cuh (24 bytes); host mirror PosFlat in
src/lob/mod.rs is repr(C) bytemuck::Pod for direct host↔device cast.

LobSimCuda gains submit_market(backtest_idx, side, size) +
read_pos(backtest_idx) -> PosFlat. Fixture harness extended with the
SubmitMarket event variant + ExpectedPos check (vwap_tol + realized_pnl_tol
for FP tolerance).

market_order_consumes_top fixture verifies a buy of 5 lots into
ask[3@5500.00, 10@5500.25] fills 3+2 → position_lots=5,
vwap_entry=5500.10 within 1e-3 tolerance. All 4 Ring 1 fixtures
(book_update × 3 + market_order × 1) green on RTX 3050.

Scope deliberately trimmed from plan C5: this commit lands market-fill
matching only. Queue decay on resting limits, in-flight latency
promotion, stop-trigger detection, and OCO leg-cancellation will land
in follow-up commits — each in its own kernel addition. This keeps the
incremental delta reviewable.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-18 08:33:16 +02:00
jgrusewski
33636d250b feat(ml-backtesting): build.rs + book_update CUDA kernel + sim skeleton
build.rs compiles every .cu under crates/ml-backtesting/cuda/ to
\$OUT_DIR/<name>.cubin via nvcc (no nvrtc per feedback_no_nvrtc.md);
pairs every env::var with rerun-if-env-changed per
pearl_build_rs_rerun_if_env_changed.md. Default arch sm_86 (RTX 3050
local); production sets FOXHUNT_CUDA_ARCH=sm_89 (L40S) or sm_90 (H100).

First kernel: book_update_apply_snapshot — block-per-backtest replace
of the 10-level book from a broadcast MBP-10 snapshot. Single-writer
per level (thread tid = level idx), no atomics per feedback_no_atomicadd.md.

lob_state.cuh defines the canonical per-block shared-mem layout that
all subsequent kernels share (Book struct in this commit; Orders, Pos,
IsvKellyState[5] added in C5-C7).

LobSimCuda host struct (in src/sim.rs) constructed from an &MlDevice;
owns per-backtest device book state + mapped-pinned input snapshots.
apply_snapshot launches the kernel and synchronises; read_books drains
device state for tests.

Three Ring 1 fixture tests (book_update_replace, _two_step,
_n_backtests=4) — N≤4 bit-exact GPU oracle assertions loaded from JSON.
All three pass on RTX 3050 locally. Verifies the build-script → cubin →
cudarc.load_cubin → kernel launch → DtoH read pipeline end-to-end.

cudarc added as direct path dep (../../vendor/cudarc) since it's not in
[workspace.dependencies] — same vendored fork ml-alpha uses.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-18 08:26:53 +02:00
jgrusewski
1ca034e22f feat(ml-backtesting): policy tree + bytecode flatten
Strategy (recursive Leaf | Ensemble | Portfolio | RegimeSwitch) +
StrategyConfig (horizon_idx + sizing_policy + sl_tp_rules +
max_concurrent_lots). Strategy::default_for(max_lots) returns the v1
default: per-horizon adaptive Ensemble with WeightedByRealizedSharpe
aggregator. No static horizon mask — capital auto-shifts via per-horizon
ISV-Kelly recent_sharpe weights (see spec §5 + §6).

Strategy::flatten() walks the tree → linear bytecode Program
(repr(C) Pod Instruction { op, arg0, arg1, arg2 }, 8 bytes/instruction)
for upload to device-global memory at slot blockIdx.x. RegimeSwitch
emits BranchIfRegime chains with arg2 patched to the post-child jump
offset.

Tests cover: 5-horizon default tree shape, flatten output for default
+ single-leaf + 2-leaf Portfolio (verifying ApplyConflict op emission),
Instruction layout size pin (8 bytes), LatencyConfig default 100ms
(IBKR+Scaleway baseline), Empirical wrap-around + empty case,
Decomposed summation.

IsvKellyStateHost host-mirror of cuda IsvKellyState (24 bytes, Pod)
defined in policy::sizing for warm-start file I/O in later commits.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-18 08:15:07 +02:00
jgrusewski
0b8dfa946f feat(ml-backtesting): host-side order types + SlotTag
OrderIntent (ergonomic host-side enum for policy YAML config + audit
reconstruction), OrderEvent (24-byte repr(C) Pod for the device-side
audit ring), TradeRecord (40-byte repr(C) Pod for the device-side
trade-log buffer), SlotTag (u32-packed {SlotKind, index, generation}
for cancel/modify against slot reuse), plus Side / LimitKind /
SlotKind / LimitParams.

Size pinning tests (24/40 bytes) ensure any future field-add forces
the kernel-side struct to migrate in lockstep. Serde roundtrips cover
OCO + Cancel variants. Renamed `.gen()` accessor to `.generation()`
to dodge the Rust 2024 reserved-keyword clash.

OCO is two limit legs linked by oco_pair byte on-device (not recursive
in OrderIntent) per spec §3 — keeps device-side layout flat for the
matching kernel coming in C5.

Adds ml-alpha + bytemuck + serde_json + tracing + thiserror to deps.
cudarc + parquet + arrow added later when CUDA + sweep aggregator
land (C4 / C9). No CUDA dependency in this commit.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-18 08:11:50 +02:00
jgrusewski
e9c4acfb12 feat(ml-alpha): MultiHorizonLoader inference-only mode
Add inference_only flag to MultiHorizonLoaderConfig that skips per-file
forward-label precomputation (~half the file-load cost), plus
peek_first() and next_inference_input() chronological-streaming methods
for the ml-backtesting LOB harness.

- min_size relaxed to cfg.seq_len when inference_only=true (training
  still requires seq_len + max_horizon + 1 for label generation)
- New cursor fields (inference_file_idx, inference_snap_idx) walk every
  loaded snapshot in chronological order; reset() zeros both
- peek_first() seeds CfcTrunk::capture_graph_a with cur==prev semantics
  (prev_ts_ns==ts_ns, trade_signed_vol=0) — natural stream-start
- next_inference_input() errors if cfg.inference_only=false (guard
  against accidental mixing of training/inference paths)
- All trainer call-sites (alpha_train example + multi_horizon_loader
  tests) updated with inference_only: false (zero behaviour change)
- Inline test module exercises both modes; tests skip gracefully when
  fixture data isn't populated rather than panicking

See docs/superpowers/specs/2026-05-18-real-lob-integration-design.md
§1 (trainer parity) + §7 (orchestrator).

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-18 08:08:38 +02:00
jgrusewski
fab2d70c92 plan(ml-backtesting): real-LOB integration implementation plan
10 atomic commits, ~40 sub-tasks, TDD per superpowers:writing-plans:
  C1 ml-alpha loader inference_only mode
  C2 ml-backtesting order types + SlotTag
  C3 policy tree + bytecode flatten
  C4 build.rs + book_update kernel + sim skeleton (3 fixtures)
  C5 order_match kernel + latency-aware fills (4 fixtures)
  C6 pnl_track kernel (1 fixture)
  C7 decision_policy kernel + per-horizon ISV-Kelly (3 fixtures)
  C8 BacktestHarness orchestrator + Ring 1b trainer-parity test
  C9 artifacts + aggregate + fxt-backtest binary
  C10 Ring 2 invariant fuzz (N in {1, 16, 256})

Plan ends with green Ring 1 (11 fixtures) + Ring 1b + Ring 2 (3 fuzz)
and a working fxt-backtest run|aggregate CLI. Ring 3 (production-day
replay), per-horizon cost-frontier sweep, model-checkpoint sweeps,
IBKR live adapter, and multi-fill P&L explicitly deferred.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-18 02:18:18 +02:00
jgrusewski
bfbaea2661 spec(ml-backtesting): real-LOB integration design
Greenfield LOB simulator + execution policy + backtest harness inside
existing ml-backtesting crate. Turns the AUC predictor into a P&L
producer at IBKR+Scaleway latency profile (100ms baseline).

Key design decisions:
- Reuse trainer's MultiHorizonLoader, Mbp10RawInput, snap_feature_assemble
  cubin, CfcTrunk captured graph verbatim (zero train-vs-deploy skew).
- Per-horizon ISV-Kelly + adaptive WeightedByRealizedSharpe aggregator
  (no static horizon mask; policy self-shifts capital).
- Block-per-backtest CUDA (32 threads/warp/block), ~1.4 KB shared mem.
- Three captured graphs (perception, step-event, decision); host branches
  only between captures.
- Mapped-pinned for all CPU/GPU buffers (hard requirement per
  feedback_no_htod_htoh_only_mapped_pinned).
- Three validation rings: N=1 bit-exact fixtures, trainer-parity check,
  N>1 invariant fuzz; production-day replay as Ring 3.

Single binary fxt-backtest with run + aggregate subcommands. No new
crates; extends existing ml-backtesting alongside barrier_backtest.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-18 01:51:11 +02:00
jgrusewski
73925b15d3 fix(ml-alpha): eval-path cfc_step launch config — block-per-batch (Phase B fix-up)
Phase B commit 1 (cfc_step block-per-batch refactor) updated the
TRAINING-path cfg_cfc to grid=(B,1,1) but missed the same change in
evaluate_batched. The eval path silently kept the legacy grid=(1,1,1)
launch config, which against the refactored kernel meant only block 0
ran — batches 1..B-1 got GARBAGE h_new (whatever was in scratch
memory), which propagated through CfC + GRN to produce garbage probs,
which BCE-eval'd to chance-level AUC.

Caught by t6z89-vs-txftz cluster A/B at L40S:
  baseline (t6z89, pre-Phase-B): mean_auc=0.7428 / h6000=0.7211 @ E0
  broken  (txftz, Phase B):      mean_auc=0.4973 / h6000=0.5136 @ E0
  train_loss matched within noise (0.6232 vs 0.6258) — the smoking gun
  for "training works, eval is broken".

Why the perception_overfit smokes didn't catch it: those tests check
that loss converges on a synthetic up-ramp signal, exercising only the
training path. Eval is exercised by `evaluate_works_after_*` smokes
but those use n_batch=1, where grid=(1,1,1) and grid=(B,1,1)=(1,1,1)
are bit-identical — the bug only manifests at B>1.

Fix: eval cfg_cfc → grid=(b_sz, 1, 1), matching training. Plus an
explanatory comment so future eyes don't repeat the mistake.

Phase B perf gains are unchanged (training path was correct). Only
eval's wall time may grow slightly because of the now-correct
per-batch parallelism doing the work it should have done.

Follow-up: re-submit cluster A/B vs t6z89 to confirm AUC trajectory
recovers to ±0.005 of baseline.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-18 00:23:42 +02:00
jgrusewski
a478ba3d84 perf(ml-alpha): block-per-batch attention pool bwd refactor (Phase B commit 4)
attention_pool_bwd refactored from grid=(1,1,1) to grid=(B,1,1). The
existing per-batch grad_ln_out writes were already uniquely indexed;
only grad_Q needed scratch+reducer (1 scratch, 1 reducer launch).

Adds 1 per-batch grad scratch buffer + 1 reduce_axis0 launch:
  attn_grad_q_scratch_d  [B, HIDDEN_DIM]
~16 KB scratch at B=32 — trivially small.

attn_pool bwd runs 1×/step (not in K-loop) so the absolute wall-time
win here is tiny. With this commit every single-SM bwd kernel in the
trainer has been refactored to block-per-batch + scratch+reducer.

Phase B kernel work complete. Next: local + cluster A/B perf benchmark
to verify acceptance gates 6, 7, 8 from the spec.

All 9 perception_overfit smokes pass.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-18 00:01:31 +02:00
jgrusewski
9607f33518 perf(ml-alpha): block-per-row VSN bwd refactor (Phase B commit 3)
variable_selection_bwd refactored from grid=(1,1,1) to grid=(B*K,1,1).
VSN's n_rows = B*K positions (one row per (batch, K-position) pair);
block-per-row matches the existing fwd kernel's layout.

Adds 2 per-row grad scratch buffers + 2 reduce_axis0 launches:
  vsn_grad_w_scratch_d  [B*K, FEATURE_DIM, FEATURE_DIM]
  vsn_grad_b_scratch_d  [B*K, FEATURE_DIM]
~210 KB scratch at B=32, K=64.

VSN bwd runs 1×/step (not K×) so the absolute wall-time win here is
small versus commits 1+2. Done for pattern uniformity — every per-batch
or per-row bwd in the trainer now uses scratch+reducer.

All 9 perception_overfit smokes pass.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 23:57:32 +02:00
jgrusewski
5c2c3b65a8 perf(ml-alpha): block-per-batch GRN bwd refactor (Phase B commit 2)
multi_horizon_heads_grn_bwd_batched refactored from grid=(1,1,1) to
grid=(B,1,1). Removes the single-SM bottleneck on the second-most-called
K-loop kernel (64×/step like cfc_bwd).

Adds 10 per-batch grad scratch buffers (one per GRN param tensor) + 10
reduce_axis0 launches collapsing B → final grad after the K-loop:
  grn_grad_w1_scratch_d    [B, 5, HEAD_MID, HIDDEN]
  grn_grad_b1_scratch_d    [B, 5, HEAD_MID]
  grn_grad_w2_scratch_d    [B, 5, HEAD_MID, HEAD_MID]
  grn_grad_b2_scratch_d    [B, 5, HEAD_MID]
  grn_grad_w_gate_scratch_d  [B, 5, HEAD_MID]
  grn_grad_b_gate_scratch_d  [B, 5]
  grn_grad_w_main_scratch_d  [B, 5, HEAD_MID]
  grn_grad_b_main_scratch_d  [B, 5]
  grn_grad_w_skip_scratch_d  [B, 5, HIDDEN]
  grn_grad_b_skip_scratch_d  [B, 5]
Total: ~8 MB scratch at B=32.

All 9 perception_overfit smokes pass (including stacked_trainer_loss_
shrinks_at_batch_32 which exercises the cross-batch reducer path on
both cfc and GRN grads).

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 23:53:43 +02:00
jgrusewski
494a2e4827 perf(ml-alpha): block-per-batch cfc_step + reduce_axis0 reducer (Phase B commit 1)
Per docs/superpowers/specs/2026-05-17-kloop-parallelization-design.md.

cfc_step_batched (fwd + bwd) refactored from grid=(1,1,1) with internal
n_batch loop to grid=(B,1,1) — each block handles one batch. Removes
the single-SM bottleneck on the K-loop's most-called kernel (64×/step).

Param-grad accumulation moves to per-batch scratch:
  cfc_grad_w_in_scratch_d  [B, n_hid, n_in]
  cfc_grad_w_rec_scratch_d [B, n_hid, n_hid]
  cfc_grad_b_scratch_d     [B, n_hid]
  cfc_grad_tau_scratch_d   [B, n_hid]

Zeroed once per training step, K-loop's 64 bwd calls += into them, then
4 reduce_axis0 launches collapse B → final grad buffers (OVERWRITE)
before AdamW. New AdamW-after-reducer invariant: final grads are
meaningful only after the reducer has run in the current step.

New reduce_axis0 kernel: single parameterised reducer [B, N] → [N] via
block tree-reduce (no atomicAdd per feedback_no_atomicadd.md). Same
pattern as layer_norm_reduce_param_grads — CUDA-Graph-safe.

cfc_step_backward_batched shared-mem dropped from (B+1)*n_hid*4 to
2*n_hid*4 bytes per block (only one row of sd_pre needed per block bi).

Tests:
- New stacked_trainer_loss_shrinks_at_batch_32: FIRST test that
  actually exercises the cross-batch reduction code path; existing
  perception_overfit suite was all B=1. Initial 0.24 → final 0.00.
- Scratch-clears test removed (explanatory comment kept): structurally
  hard to assert directly due to begin_capture/end_capture not
  executing kernels; the B=32 convergence smoke implicitly validates
  scratch zeroing since divergence would otherwise be immediate.

All 9 perception_overfit smokes + 4 backward_finite_diff tests pass.

build.rs:
- KERNELS list adds "reduce_axis0"
- Cache-bust → v11

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 23:47:40 +02:00
jgrusewski
c508d101c1 plan(perf): K-loop block-per-batch parallelization implementation plan
Bite-sized 22-task plan implementing the design at
docs/superpowers/specs/2026-05-17-kloop-parallelization-design.md.

Four atomic commits per the spec's Rollout section:
  Commit 1: reduce_axis0 kernel + cfc_step refactor (Tasks 1-10)
  Commit 2: GRN bwd refactor (Tasks 11-14)
  Commit 3: VSN bwd refactor (Tasks 15-17)
  Commit 4: attention_pool bwd refactor (Tasks 18-20)

Each commit covers kernel rewrite + per-batch scratch buffers +
reducer launches + memset_zeros wiring + smoke validation.

Tests added across commits:
  - cfc_bwd_b1_oracle.rs (Task 6): B=1 oracle vs single-sample helper
    within relative_eq 1e-5 (FP-tolerant, not bit-exact)
  - stacked_trainer_loss_shrinks_at_batch_32 (Task 7): first test
    that exercises the cross-batch reducer path
  - cfc_bwd_scratch_clears_between_steps (Task 8): scratch-zero
    regression guard

Acceptance gates 1-8 from the spec are mapped to Tasks 6, 7, 8, 9,
21, 22 (local + cluster A/B perf). Reference baseline for gate #8
is t6z89's per-epoch AUC trajectory.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 23:27:20 +02:00
jgrusewski
ea1d8703b7 spec(perf): revise K-loop parallelization design — full critical pass
Revises all 12 issues from the self-review pass:

1. (HIGH) Soften bit-equivalence claim — single-sample helper and
   batched kernel at B=1 are different CUDA kernels; FP order may
   differ. Acceptance is relative_eq ≤ 1e-6, not bit-exact.
2. (HIGH) Explicit asymmetry: only cfc_step has a single-sample GPU
   oracle. GRN/VSN/attn bwd rely on smoke + chain-rule preservation.
   feedback_no_cpu_test_fallbacks.md forbids a CPU reference oracle.
3. (HIGH) Realistic targets — 3× floor, 5× stretch. Drops 15× claim
   which was Amdahl-bounded under any reasonable assumption.
4. (MED) AdamW-after-reducer invariant stated explicitly: final grad
   buffer is OVERWRITE by reducer, meaningful only after reducer ran.
5. (MED) New B=32 smoke test (stacked_trainer_loss_shrinks_at_batch_32)
   actually exercises the cross-batch reduction path; existing
   perception_overfit suite is all B=1.
6. (MED) Rollout commit 1 bundles reduce_axis0 + first consumer
   (cfc_step refactor) to avoid feedback_wire_everything_up.md
   orphan-kernel anti-pattern.
7. (LOW) Drop "merge to ml-alpha-phase-a" — user already chose direct
   commits to that branch; clarify in Rollout.
8. (LOW) Add explicit scratch-sizing formula:
        scratch_bytes ≈ B × Σ(param tensor sizes per kernel).
9. (LOW) Remove resolved open question (memset_zeros ordering).
10. (STRUCT) Post-refactor bottleneck analysis section — names the
    next L40S floor (Mamba2 scan, launch latency, cuBLAS).
11. (STRUCT) cudaFuncSetAttribute note — refactored cfc_bwd's per-
    block shared mem drops to 1 KiB; no attribute change needed.
12. (STRUCT) Explicit Rollback section — atomic-commit-per-kernel +
    revert strategy; rollback baseline is the spec commit
    (54aa69c10) on ml-alpha-phase-a.

Plus locks the target pool to L40S — speedup must be attributable to
the refactor, not to a hardware bump. H100 / BF16 / larger batch
become candidates for a follow-up spec once the L40S floor is known.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 23:13:32 +02:00
jgrusewski
54aa69c108 spec(perf): K-loop block-per-batch parallelization design
Documents the design for fixing the single-SM bottleneck in five
backward/K-loop kernels: cfc_step (fwd+bwd), GRN bwd, VSN bwd, and
attention_pool bwd. All currently use grid=(1,1,1) with an internal
n_batch loop — on L40S (142 SMs) with B=32 this puts <1% of the GPU
to work in the K-loop critical path.

Architecture: block-per-batch (grid=(B,1,1)) for the kernel body, plus
per-batch grad scratch buffers reduced via a single parameterised
reduce_axis0 kernel (block tree-reduce, no atomicAdd per
feedback_no_atomicadd.md). Same pattern as the existing LayerNorm bwd
reducer — CUDA-Graph-safe, debuggable, and consistent with foxhunt's
no-cooperative-groups discipline.

Target: ≥3× epoch wall speedup (stretch 8-15×). Makes 3-fold CV
tractable (10.5h → 2-3h) and unblocks decision-stride / state-dim
sweeps that compound the gain.

Acceptance gates: (a) all 8 perception_overfit smokes still converge,
(b) new B=1 bit-equivalence test asserts the refactored batched bwd
kernel at B=1 matches the existing single-sample helper byte-for-byte,
(c) cluster A/B vs t6z89 baseline shows AUC trajectory within ±0.005
and epoch wall ≥3× faster.

Atomic refactor per kernel — one commit per kernel covering the
kernel rewrite, scratch buffer alloc, reducer launch wiring, and
smoke. No "_legacy" parallel kernels per feedback_no_legacy_aliases.md
+ feedback_no_partial_refactor.md.

Open implementation-plan decisions: exact memset_zeros ordering inside
the captured graph, batch-vs-per-tensor reducer launches, optimal
block_dim for reduce_axis0 itself.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 23:06:28 +02:00
jgrusewski
7a558b88b7 feat(ml-alpha): wire attention pool into PerceptionTrainer (Phase 3.2)
Replaces CfC's zero-initialised h_old at k=0 with the attention-pooled
context vector — a learned content-addressable summary over all K LN_b
output positions. The K-loop's recurrent semantics (h_old at k+1 = h_new
at k) are preserved; only the INITIAL state at k=0 changes from zero to
the pooled context.

Forward chain change:
  ... → m2 → LN_b → ln_out_d [B, K, HIDDEN_DIM]
       → attention_pool_fwd(Q, ln_out_d)
            → attn_context_d [B, HIDDEN_DIM]  (used as h_old@k=0)
            → attn_weights_d [B, K]            (saved for bwd)
       → K-loop CfC (h_old@k=0 = attn_context, not zero)

Backward chain change:
  K-loop bwd ends with grad_h_carry_d holding the gradient that would
  have flowed into the initial h_old = grad on attn_context.
  attention_pool_bwd consumes:
    Q, ln_out_d, attn_weights_d (forward state)
    grad_h_carry_d = grad_context
  Writes (BOTH `+=`):
    grad_attn_q_d  (accumulates Q gradient — pre-zeroed at step start)
    grad_h_enriched_seq_d (ADDS attn-path contribution onto LN_b output
                           gradient — chains with K-loop contribution)
  LN_b bwd then consumes the now-summed grad_h_enriched_seq_d.

Trainer state additions (8 fields):
  attn_q_d, attn_context_d, attn_weights_d, grad_attn_q_d,
  attn_fwd_fn, attn_bwd_fn, _attn_module, opt_attn_q

Q is tiny (HIDDEN_DIM=128 floats); initialised near zero so initial
attention ≈ uniform 1/K (context ≈ mean of LN_b output). Model learns
content addressing from a near-uniform starting point.

Eval path mirrors training: attn_pool_fwd runs after LN_b fwd,
attn_context_d feeds the eval K-loop at k=0.

Trainer now manages 22 AdamW: CfC×4 + GRN heads×10 + LN×2 + LN_a×2 +
VSN×2 + Attn Q×1 + Mamba2×2 grouped.

Synthetic overfit smoke: stride=1 0.29 → 0.0000 in 50 steps (faster
than pre-attn 0.30), stride=4 0.30 → 0.0000. All 8 perception_overfit
tests PASS. Demonstrates the full Phase 1+2+3 stack (VSN → m1 → LN_a →
m2 → LN_b → attn pool → CfC + GRN heads) is wired forward + backward
end-to-end with every gradient flowing through every learned param.

Phase 1+2+3 capacity scale-up complete. The cumulative architectural
lift over the 3-fix-stack baseline (496q7 mean_auc=0.716 / h6000=0.704)
will be measured by deploying this stack head-to-head against bsml6.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 22:38:05 +02:00
jgrusewski
f76437c0c7 feat(ml-alpha): attention pool forward+backward CUDA kernels (Phase 3.1)
Single-head attention pool over Mamba2 K-positions, designed to replace
the CfC's zero-initialised `h_old` at k=0 with a learned content-
addressable summary over all K LN_b output positions. Forward math:

  scores[k]   = Q · keys[b, k, :]              # [K]
  attn[k]     = softmax_k(scores)              # [K]
  context[h]  = sum_k attn[k] * values[b, k, h]   # [HIDDEN_DIM]

For our attention pool, keys == values == LN_b output [B, K, HIDDEN_DIM].
Single learned param: Q [HIDDEN_DIM]. Tiny (128 params).

Forward layout: grid = (B, 1, 1), block = HIDDEN_DIM=128 threads. Three
passes: (1) K dot-products with tree-reduce over HIDDEN_DIM, (2)
softmax over K with max-subtract+sum, (3) weighted sum into context.

Backward chain rule:
  d_attn[k]   = sum_h grad_context[h] * values[b, k, h]
  d_scores[k] = attn[k] * (d_attn[k] - sum_kp attn[kp] * d_attn[kp])
  d_Q[h]     += sum_{b, k} d_scores[k] * values[b, k, h]
  d_values[b, k, h] += grad_context[h] * attn[k] + d_scores[k] * Q[h]

Both `d_Q` and `d_values` use += semantics:
- d_Q: accumulates across batch (single block, internal n_batch loop).
- d_values: writes ADD onto whatever grad_ln_out already holds, so the
  trainer can chain it on top of the K-loop's contribution to the LN_b
  output gradient (no separate add-kernel needed).

Single-writer (no atomicAdd): one block per launch, thread h owns
column h of grad_ln_out for ALL (b, k). Internal n_batch loop matches
the GRN / VSN bwd pattern.

build.rs:
  - "attention_pool" added to KERNELS
  - Cache bust → v10

Wiring into PerceptionTrainer (Phase 3.2) is the follow-up commit:
add attn_q_d learned param + per-batch context + attn_weights buffers,
run attn_pool_fwd between LN_b fwd and the K-loop, use attn_context as
the K-loop's k=0 h_old (instead of zero_h_d), and chain attn_pool_bwd
after the K-loop reverse pass.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 22:32:00 +02:00
jgrusewski
8f5f22fe4d feat(ml-alpha): 2-stack Mamba2 with inter-stack LayerNorm (Phase 2B)
Doubles the trunk capacity. Forward chain:
  snap_features → VSN → m1 → LN_a → m2 → LN_b → CfC → GRN heads

m1 = Mamba2Block { in_dim=FEATURE_DIM=40, hidden_dim=128 }
m2 = Mamba2Block { in_dim=128, hidden_dim=128 }

Both stacks share the SAME state_dim (cfg.mamba2_state_dim) and the
SAME hidden_dim. m2 reads m1's output (post LN_a). LN_a is a separate
LayerNorm instance from LN_b (the existing trunk-to-CfC normaliser).

Backward chain reverses the forward:
  ... grad_ln_in_d → m2.bwd → m2_grads_buffers.d_x_from_in (= LN_a output grad)
  → LN_a.bwd → grad_ln_a_in_d (= m1 output grad)
  → m1.bwd → m1_grads_buffers.d_x_from_in (= VSN output grad)
  → VSN.bwd → ...

Both Mamba2 stacks emit `d_x_from_in` (Phase 2D refactor already
exposed it on m1; m2 uses the same code path). LN_a uses the existing
layer_norm_fwd / layer_norm_bwd / layer_norm_reduce_param_grads
kernels — no new CUDA work, just a second instance with its own
gain/bias/stats/grad scratches.

New trainer state: ~17 fields (mamba2_l2 + its scratch + LN_a + LN_a
grads + opt_ln_a_*). All initialised in the construction order that
respects the `stream` move-into-Self at the end of new().

set_lr_mamba2 now updates BOTH stacks' AdamW configs. Total AdamW
instances on the trainer: 21 (CfC×4 + GRN heads×10 + LN×2 + LN_a×2 +
VSN×2 + Mamba2×2 grouped × 9 params each).

All 8 perception_overfit smokes pass: synthetic constant-direction
signal converges 0.31 → 0.0000 by step 100 (matches single-stack
trajectory — proves both stacks are wired forward + backward and all
21 AdamW optimisers move weights).

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 22:28:08 +02:00
jgrusewski
73d68ab786 feat(ml-alpha): wire TFT VSN into PerceptionTrainer (Phase 2D.3)
VSN sits between snap_feature_assemble and Mamba2 input:
  snap_assemble → window_tensor_d [B, K, 40] (raw)
    → VSN fwd      → vsn_out_d [B, K, 40] (gated) + vsn_gates_d [B*K, 40]
    → Mamba2 fwd   → ...

Forward path: per-position softmax over FEATURE_DIM features, output[i] =
x[i] * gates[i]. Gates initialised near-uniform (W_vsn ~ N(0, 1/√FEATURE_DIM),
b_vsn = 0) so the model starts from "all features matter equally" and
learns regime-conditional gates.

Backward path: Mamba2 backward already computed d_x_from_in (gradient
w.r.t. its input) on its grads_buffers — previously labeled "unused but
allocated", now consumed by VSN bwd as grad_y. Zero refactor to
mamba2_block.rs.

VSN bwd writes:
  grad_W_vsn [40, 40]  → opt_vsn_w  (AdamW, default wd)
  grad_b_vsn [40]      → opt_vsn_b  (AdamW, wd=0)
  vsn_grad_x_d [B*K, 40] → discarded (snap_features are non-trainable
                            transforms of raw MBP-10 data)

Param grads are explicitly memset_zeros before each VSN bwd call (the
kernel uses += semantics like the GRN bwd, but VSN runs ONCE per step
not K times, so zeroing makes the += a clean overwrite — matches
Adam's `step` expectations).

Eval path mirrors training (VSN fwd applied between snap_assemble and
Mamba2 fwd) so eval sees the gated features layers were trained on.

Trainer now manages 19 AdamW: CfC×4 + GRN heads×10 + LN×2 + VSN×2 +
Mamba2 grouped.

Synthetic overfit smoke: stride=1 initial=0.30 → final=0.00 in 50 steps
(faster than pre-VSN's 0.32, consistent with VSN's near-uniform init
giving a small head start). All 8 perception_overfit tests pass.

Note: the smoke proves VSN's W/b actually move via the Mamba2 input
gradient — if `d_x_from_in` were zero, VSN params wouldn't update and
the chain still converges but VSN remains identity. The fact that
initial loss DIFFERS (0.30 vs 0.32) shows VSN is in the forward path
end-to-end.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 22:21:47 +02:00
jgrusewski
c363a7e94c feat(ml-alpha): TFT VSN forward+backward CUDA kernels (Phase 2D.1+2D.2)
Per-position softmax-normalised feature gating for the trunk entry.
Per (b, k) sample:
  gate_logit[i] = sum_j W_vsn[i, j] * x[j] + b_vsn[i]
  gates         = softmax(gate_logit)         # [FEATURE_DIM]
  y[i]          = x[i] * gates[i]

Backward chain rule (cleanly factored from the softmax Jacobian):
  d_gates[i] = grad_y[i] * x[i]
  d_logit[i] = gates[i] * (d_gates[i] - sum_j gates[j] * d_gates[j])
  grad_W[i,j] += d_logit[i] * x[j]
  grad_b[i]   += d_logit[i]
  grad_x[j]   = grad_y[j] * gates[j] + sum_i d_logit[i] * W[i,j]

Single-writer (no atomicAdd): thread tid owns row tid of grad_W and
column tid of d_x_via_W. ONE block per launch (loops n_rows internally),
same pattern as 2-layer / GRN bwd kernels.

Softmax uses standard max-subtract + sum trick for numerical
stability. Block dim = 64 (one warp + 24 idle threads at
FEATURE_DIM=40).

Wiring blocked on: Mamba2 backward needs to emit `d_input` (currently
dropped at line 1413 of mamba2_block.rs via `_d_input`). Next commit
exposes that so VSN bwd has the right grad_y signal — and the same
refactor unblocks Phase 2B (2-stack Mamba2 needs the inter-stack LN
to backprop through the 2nd stack's d_input).

build.rs:
  - "variable_selection" added to KERNELS
  - Cache bust → v9

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 22:16:58 +02:00
jgrusewski
005ded9722 feat(ml-alpha): TGN Δt Fourier features in snap_feature_assemble (Phase 2C)
Bumps FEATURE_DIM 32→40. Slots [32..40] now carry 8 TGN-style Fourier
features encoding the elapsed time Δt = ts_ns - prev_ts_ns:
  (cos(ω_k · Δt_ns), sin(ω_k · Δt_ns))_{k=0..3}
at log-spaced periods [60s, 6s, 600ms, 60ms].

This gives Mamba2's input vector explicit Δt encoding that's
discriminative across temporal scales — particularly important once
decision-stride > 1 (Phase 2A) lands and the gap between K-positions
becomes irregular. Without these features the model has no way to
distinguish "1ms gap" from "1s gap" between consecutive K-positions.

Slots [0..32] unchanged (bit-equivalent for the first 32 features).
Reserved-zero slots [26..32] kept for future macro context. Slots
[20..26] still hold the loader-precomputed EMA regime cascade.

Frequencies stored in __constant__ memory (SNAP_DT_OMEGAS[4]) — small
table, broadcast read pattern, no register pressure. Frequency
selection rationale (one per log-decade):
  60s    — minute-scale macro session context
  6s     — 10s-scale liquidity windows
  600ms  — sub-second microstructure
  60ms   — tick-cluster spacing

Δt clamped to >= 0 so the rare out-of-order timestamp doesn't produce
nonsense angles. Each (cos, sin) pair satisfies cos²+sin² = 1
(verified by new test `dt_fourier_features_are_bounded`).

New tests in snap_feature_bit_equiv.rs:
  - dt_fourier_features_are_bounded: |slot| <= 1 + cos²+sin² == 1
  - dt_fourier_discriminates_scales: Δt=1ms vs Δt=1s produce L2-distinct
    Fourier vectors (>0.1)
  - reserved_slots_are_zero updated to check [20..32] (regime + reserved)
    instead of [20..FEATURE_DIM]

All 8 perception_overfit smokes still pass (synthetic stride=1 and
stride=4 both converge 0.32 → 0.0000) — proves the wider FEATURE_DIM=40
input doesn't break the Mamba2+LN+GRN chain.

build.rs cache-bust → v8.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 22:09:08 +02:00
jgrusewski
70d5fc29cf feat(ml-alpha): decision-stride loader + CLI + Mamba2 dt_s scaling (Phase 2A)
Decision-stride S lets a length-K sequence span ((K-1)*S + 1) raw
snapshots instead of K consecutive ones — expands the effective
time-window covered by each sequence at the same K-positions compute
cost. With K=64 and S=4, the window covers 256 ticks (~5s on ES MBP-10
at 20ms-tick) instead of 64 ticks (~1.3s).

Loader (crates/ml-alpha/src/data/loader.rs):
- `MultiHorizonLoaderConfig.decision_stride: usize` (default 1, must
  pre-existing call sites add the new field).
- `next_sequence` reads snapshot at `anchor + k * stride`; labels at the
  same indices (labels stay in absolute-snapshot horizons regardless of
  stride, e.g. h=6000 always means "predict 6000 raw snapshots forward").
- `prev` snapshot for microstructure features (prev_mid, prev_ts_ns)
  now points to the prior K-position (`anchor + (k-1)*stride`), NOT the
  consecutive-snapshot prior, so `Δt = ts_ns - prev_ts_ns` carries the
  actual elapsed time between K-positions (consumed by Mamba2's dt_s and
  the planned Phase 2C TGN Fourier features).
- New `#[ignore]` real-data test: `loader_stride_4_yields_correct_spacing`
  asserts Δt monotonicity at stride=4.

Mamba2 dt_s (crates/ml-alpha/src/trainer/perception.rs):
- `PerceptionTrainerConfig.decision_stride: usize` plumbs the stride
  through. dispatch_train_step + evaluate_batched now use
  `dt_s = decision_stride as f32` so Mamba2's selective scan
  `exp(-dt * sigmoid(a))` reflects the real elapsed time. With stride=1
  the behaviour is identical to before.

CLI (crates/ml-alpha/examples/alpha_train.rs):
- `--decision-stride <S>` flag (default 1) wired into both train and val
  loaders + PerceptionTrainerConfig.

Argo workflow:
- `decision-stride` parameter on the template (default "1") +
  `--decision-stride` script flag + propagation into the train pod's
  alpha_train invocation.

Synthetic smoke (tests/perception_overfit.rs):
- `stacked_trainer_loss_shrinks_with_stride_4` proves the trainer-level
  dt_s=4.0 keeps the Mamba2+LN+CfC+GRN chain numerically stable.
  Converges 0.32 → 0.0000 (matches stride=1 smoke trajectory — dt_s
  scaling didn't break the SSM dynamics).

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 22:04:38 +02:00
jgrusewski
b6dfd89903 feat(ml-alpha): wire TFT GRN heads into PerceptionTrainer (Phase 1.7)
Replaces the single-linear heads (5 × 128) with per-horizon TFT GRN:
  trunk h[B, 128]
    → W1 (HIDDEN→HEAD_MID) + GELU      = eta_2 [B, 5, HEAD_MID]
    → W2 (HEAD_MID→HEAD_MID)            = eta_1 [B, 5, HEAD_MID]
    → (W_gate, W_main) split            = (gate_lin, main) [B, 5]
    → W_skip (HIDDEN→1)                 = skip [B, 5]
    → skip + sigma(gate_lin) * main     = logit [B, 5]
    → sigma(logit)                      = p [B, 5]

10 trainable param tensors per horizon (5 weight + 5 bias), 10 AdamW
state objects (biases get wd=0), 6 per-K-position intermediate buffers
sized [K, B, 5, HEAD_MID] (z1, a1, z2) and [K, B, 5] (gate_logit, main,
logit). All saved during fwd, consumed by GRN bwd for the chain rule.

K-loop forward replaces `multi_horizon_heads_batched` with the new
`multi_horizon_heads_grn_fwd_batched` kernel. K-loop backward replaces
`multi_horizon_heads_backward_batched` with `multi_horizon_heads_grn_bwd_batched`.
ISV `lambda_d` is consumed by the GRN bwd kernel to scale ONLY the
trunk gradient (per pearl_adam_normalizes_loss_weights.md — the
effective lever is the trunk gradient, not loss weight).

Eval path also uses the GRN fwd kernel so eval sees the same
distribution downstream layers trained on (same pattern as Phase 1.6
LN wiring).

Xavier-uniform init: scale = sqrt(1 / fan_in) per tensor:
  W1, W_skip:        scale = sqrt(1/HIDDEN)        ≈ 0.088
  W2, W_gate, W_main: scale = sqrt(1/HEAD_MID)     ≈ 0.125

Parameter count per horizon: W1 (8192) + b1 (64) + W2 (4096) + b2 (64)
  + W_gate (64) + b_gate (1) + W_main (64) + b_main (1) + W_skip (128)
  + b_skip (1) = 12,675 params × 5 horizons = 63,375 head params total
  (vs single-linear's 5*(128+1)=645). Trunk grad dimensionality is
  unchanged.

Synthetic overfit smoke: loss 0.32 → 0.0000 by step 50 (faster than
single-linear's 0.37 → 0.0006 in 250 steps); all 7 perception_overfit
tests PASS. Confirms the full Mamba2 + LN + CfC + GRN backward chain
is wired correctly and all 17 AdamW optimisers move weights.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 21:54:52 +02:00
jgrusewski
5e23005dea feat(ml-alpha): TFT GRN forward+backward kernels for multi-horizon heads (Phase 1.7a)
Per-horizon GRN structure (Lim et al. 2021 §3.3 adapted to scalar output):
  eta_2[k, m] = GELU(W1[k, m, :] @ h + b1[k, m])             # [HIDDEN] → [HEAD_MID]
  eta_1[k, m] = W2[k, m, :] @ eta_2[k, :] + b2[k, m]          # [HEAD_MID] → [HEAD_MID]
  gate_lin[k] = W_gate[k, :] @ eta_1[k, :] + b_gate[k]        # → scalar
  main[k]     = W_main[k, :] @ eta_1[k, :] + b_main[k]        # → scalar
  skip[k]     = W_skip[k, :] @ h + b_skip[k]                  # [HIDDEN] → scalar
  logit[k]    = skip[k] + sigmoid(gate_lin[k]) * main[k]
  p[k]        = sigmoid(logit[k])

Gated residual lets each per-horizon head learn "linear vs deeper-transform"
gating, matching the regime-conditional alpha pattern from
pearl_snapshot_alpha_is_regime_conditional (~20% of book states carry the
edge; spread-Q4 hits 75% acc, middle quintiles below chance).

Backward chain rule covers all 10 parameter tensors + the trunk gradient
(skip-path direct + main-path through W2→GELU→W1, lambda-scaled).

Single-writer discipline (no atomicAdd per feedback_no_atomicadd.md):
- Thread m owns row m of grad_w1 (col i in 0..HIDDEN), row m of grad_w2
  (col m_in in 0..HEAD_MID), and column m of d_eta_2.
- Threads 0..4 own per-horizon scalar grads (skip/gate/main biases).
- Trunk grad_h tiles i over 2 strides of HEAD_MID for HIDDEN=128 coverage.

Shared mem: ~6.5KB (s_a1 + s_z2 + s_d_eta1 + s_d_eta2 + s_d_z1 + scalars),
well within 48KB limit.

Existing 2-layer MLP kernels (Tasks 1.3/1.4) stay in the cubin as
ablation baseline; the wired path becomes GRN once perception.rs lands.

build.rs cache-bust → v7.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 21:47:09 +02:00
jgrusewski
010445b5df docs(plans): pivot Phase 1.7 to TFT GRN heads; add Phase 2C (TGN Δt Fourier) + 2D (TFT VSN)
User directive 2026-05-17: borrow TFT GRN over the planned 2-layer MLP heads.
GRN structure: 2-layer GELU MLP body (eta_2 → eta_1) + GLU gate + main +
skip-projection from trunk → final sigmoid. Gives per-horizon "linear vs
deeper-transform" gating, matches the regime-conditional alpha pattern
(pearl_snapshot_alpha_is_regime_conditional). 5x parameter count vs the
2-layer MLP but the gated residual is exactly what TFT empirically wins on.

Phase 2C (TGN Δt Fourier features): 8 sin/cos features of Δt at log-spaced
periods [60s, 6s, 600ms, 60ms] appended to snap_features. Critical with
decision-stride>1 where Δt varies across positions. Bumps FEATURE_DIM 32→40.

Phase 2D (TFT VSN): per-feature softmax-normalised gates at the trunk entry,
replacing raw concat of snap_features. Learns to down-weight noisy
features per regime (canonical: trade-flow in low-volume, OFI in
spread-Q4). 2 new param tensors, 1 new cuda kernel (fwd+bwd).

Existing 2-layer MLP kernels from Tasks 1.3/1.4 stay in the cubin as
ablation baseline; wired path becomes GRN.

Phase 1.7 plan now spells out the full GRN forward + backward chain rule
(skip + sigmoid(gate) * main → outer sigmoid), kernel signatures,
parameter Xavier init, AdamW × 10 setup, and an extended numerical-grad
check covering all 10 new param tensors.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 21:41:02 +02:00
jgrusewski
c24cf7423b feat(ml-alpha): wire LayerNorm into PerceptionTrainer (Phase 1.6)
LN sits between Mamba2 encoder output and the K-loop CfC consumer:
  Mamba2.h_enriched_seq [B,K,H] → LN fwd → ln_out_d [B,K,H]
  → transpose → h_enriched_seq_t_d [K,B,H] → CfC + heads

Forward path: LN consumes raw Mamba2 output, writes ln_out_d + per-row
stats (mean, inv_std). Transpose now reads from ln_out_d.

Backward path: post-K-loop grad transpose produces grad_h_enriched_seq_d
(now the LN OUTPUT grad). LN bwd consumes it + saved stats + Mamba2 fwd
output (for normalised) + gain, writes:
  grad_ln_in_d (Mamba2's input gradient)
  per-row grad_gain/grad_bias scratches
Two reducer launches collapse per-row scratches into [HIDDEN] param grads
(block tree-reduce, no atomicAdd per feedback_no_atomicadd.md).

AdamW state for ln_gain + ln_bias added (wd=0, lr=lr_cfc). Mamba2 bwd
now consumes grad_ln_in_d, NOT grad_h_enriched_seq_d.

Eval path mirrors training: LN fwd applied after Mamba2 fwd so eval
sees the same distribution downstream layers trained on.

Synthetic overfit smoke passes: BCE 0.37 → 0.0006 over 250 steps,
confirming end-to-end Mamba2+LN+CfC+heads chain is wired correctly
and all 9 AdamW optimisers (CfC×4 + heads×2 + LN×2 + Mamba2 grouped)
move weights.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 21:38:15 +02:00
jgrusewski
59e236b4e8 feat(ml-alpha): 2-layer heads backward kernel with ISV lambda (Phase 1.4) 2026-05-17 21:24:49 +02:00
jgrusewski
e0a497da9e feat(ml-alpha): 2-layer GELU MLP heads — forward kernel (Phase 1.3) 2026-05-17 21:23:46 +02:00
jgrusewski
167f065647 feat(ml-alpha): LayerNorm backward + per-row param-grad reducer kernels 2026-05-17 21:22:44 +02:00
jgrusewski
d8cd90c130 feat(ml-alpha): LayerNorm forward kernel for trunk-pre-CfC normalisation 2026-05-17 21:22:06 +02:00
jgrusewski
b5530b551b docs(plans): model capacity scale-up — 3-phase implementation plan
Three sequenced architectural capacity additions to push h6000 AUC
from current ~0.72-0.74 cross-fold plateau toward ≥0.78 deployment
target. Each phase independently deployable + measurable.

Phase 1 — Per-horizon specialisation (~1.5hr code):
  - LayerNorm between Mamba2 trunk and CfC K-loop (with backward
    + per-row param-grad reduction kernel; no atomicAdd).
  - 2-layer GELU MLP heads [hidden=128 → mid=64 → 1] per horizon;
    ISV lambda integrates into the trunk-grad component of head
    backward. Tasks 1.1-1.8 fully detailed (kernel source +
    integration + numerical grad check + smoke + deploy).

Phase 2A — Decision-stride sampling (~1.5hr code):
  - User-prioritised lever. Yield every S-th snapshot per training
    sequence; K=64 with stride=4 covers 256 ticks of context for
    the same compute as 64 ticks at stride=1. Mamba2's dt_s scalar
    becomes stride-aware. Tasks 2A.1-2A.5 fully detailed (loader
    refactor + test + CLI plumbing + synthetic smoke + deploy).

Phase 2B — 2-stack Mamba2 (~2hr code, sketch):
  - Two Mamba2Block instances; forward chain
    snap_feat → mamba2_l1 → LN → mamba2_l2 → LN → CfC.
  - Acceptance criteria + key implementation notes documented;
    bite-sized tasks elaborated when Phase 1+2A results land.

Phase 3 — Attention pool over Mamba2 K-positions (~4-6hr code,
sketch):
  - Replace CfC's zero-init initial state with an attention-
    pooled context vector over all K Mamba2 outputs. Design
    decision documented (Option A: attention sets initial state,
    preserving CfC's recurrent path).

Cross-phase deployment loop documented: fold-1 smoke → 3-fold
CV → per-fold comparison + isv-snapshot trajectory archive.

Honors `superpowers:writing-plans` skill: exact file paths,
complete code in every step, exact commands with expected output,
TDD-style steps, frequent commits, no placeholders in Phase 1
tasks. Self-review pass complete.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 21:18:49 +02:00
jgrusewski
00da163078 feat(ml-alpha): h6000-aligned ISV — uniform BCE + z-score lambda + K=64
Three correlated fixes addressing the architectural inconsistency
surfaced by the 3-fold ISV CV: we built a horizon-aware gradient
controller (ISV) but suppressed its target horizon (h6000) to 0.36%
of the loss via auto-horizon-weights, then used a ratio formula
that never approached its own clamp ceiling. ISV's lambda was
operating on rounding error.

(1) Uniform BCE weights as auto-default
    trainer/perception.rs: `auto_horizon_weights` now returns
    [1.0; 5] regardless of seq_len. Prior schedule `min(1, K/h)`
    gave h6000 weight 0.0053 at K=32 — combined with lambda ~1.04,
    h6000's effective loss contribution was ~0.37%, indistinguishable
    from zero. With uniform weights, each horizon contributes 20% and
    ISV's lambda actually has something to scale.

(2) Z-score lambda derivation
    cuda/horizon_lambda.cu: replace `ratio = ema_h / mean(ema)` with
    `z_h = (ema_h - mean) / std(ema); lambda = clamp(1.0 + 0.5*z, 1.0, 2.0)`.
    Per `pearl_zscore_normalization_for_magnitude_asymmetric_signals.md`
    z-score makes lambda spread scale-invariant of the absolute EMA
    level. The ratio formula gave lambdas ≤ 1.04 in our data because
    per-horizon BCE clusters tightly (range ~0.04) while mean is
    ~0.65. Z-score fills the [1.0, 2.0] envelope: 1σ → 1.5, 2σ →
    ceiling. Boost-only asymmetric clamp preserved.

    Test verification on the existing smoke (after 5 steps):
      ema    = [0.526, 0.522, 0.608, 0.641, 0.553]
      lambda = [1.00, 1.00, 1.40, 1.76, 1.00]
    Previously with ratio formula, max lambda on the same data
    would have been ~1.05. h1000 (1.5σ above mean BCE here) now
    gets a 76% trunk-gradient boost vs uniform.

(3) Default --seq-len 32 → 64
    examples/alpha_train.rs: K=32 gives the model 0.5% of the
    h6000 prediction window as in-window context. K=64 doubles
    that, giving Mamba2's SSM state more material to build
    long-horizon predictions. Within the kernel's MAMBA2_KERNEL_SEQ_MAX
    cap of 96. Per-epoch wall scales ~K (more K-loop launches in
    the captured graph, ~2× wall at K=64 vs K=32 for the K-loop
    portion of dispatch).

Cache-bust v5 in build.rs to force nvcc recompile against the new
horizon_lambda.cu formula on the cluster's /cargo-target PVC. Old
cubins compute a numerically different lambda; running them against
the new Rust loop would silently apply the wrong gradient scaler.

Validation: 7 perception_overfit tests + 26 lib + 23 integration
ml-alpha tests pass. Synthetic overfit still converges to 0.0006.
horizon_ema_and_lambda_track_after_training observes the new
lambda spread (1.0-1.76) and asserts the asymmetric clamp envelope.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 20:58:44 +02:00
jgrusewski
a45fd85986 feat(ml-alpha): raise Mamba2 state cap 16→32 + add auc_h6000 early-stop
Three correlated changes for the next CV round:

1. Mamba2 state_dim cap: 16 → 32
   cuda/mamba2_alpha_kernel.cu: MAMBA2_ALPHA_MAX_STATE_D 16 → 32.
   Per-thread state register `float x[32]` (128 B/thread) and
   per-thread x_hist replay cache `float x_hist[K*32]` (up to
   12 KiB/thread of local memory at K=96). L40S/H100 register file
   (256 KiB/SM) absorbs this without occupancy collapse for our
   block dims (32-128 threads). Update Rust-side
   MAMBA2_KERNEL_STATE_MAX + validation message + test name. Kernel
   header doc updated.

2. New early-stop option: auc_h6000
   examples/alpha_train.rs: add the long-horizon AUC as a third
   early-stop metric. The ISV CV (3a196382f, 5d42ab0e9, 0171c8c0e)
   showed mean_auc-best-epoch and h6000-best-epoch can differ by
   1-2 epochs and the h6000 gap can be 5-6pt within a single run
   (fblb2 fold-1: saved E10 h6000=0.681, but E11 h6000=0.739 — we
   threw away the deployment-better checkpoint). For multi-minute
   trading deployment we want the h6000-best checkpoint directly.

3. Summary JSON: best_auc_h6000_epoch / best_auc_h6000 /
   best_auc_h6000_per_horizon
   So the analysis tooling can see the h6000-best checkpoint
   independently of mean_auc / val_loss bests.

Test rename: test_mamba2_config_rejects_state_over_16 →
test_mamba2_config_rejects_state_over_32 (tests now reject state_dim=33).

build.rs cache-bust v4 forces cluster nodes to recompile the kernel
against the new MAMBA2_ALPHA_MAX_STATE_D — old cubins from previous
SHAs were sized for state_d=16 and would silently truncate state_d=32
state arrays.

Validation: 26 lib + 23 integration ml-alpha tests pass. Mamba2 block
tests use state_dim=8 or 16 (well below the new cap), exercise both
the forward + backward + AdamW paths.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 20:25:47 +02:00
jgrusewski
d220755309 obs(alpha_train): per-epoch ISV controller snapshot (EMA + lambda)
Fold-2 of the asymmetric-ISV 3-fold CV regressed -2.0pt mean_auc
vs no-ISV (0.696 vs 0.716) while folds 0 and 1 gained +1.3pt and
+2.5pt respectively. To understand why the controller hurts that
specific regime, log the per-horizon BCE EMA and lambda multiplier
each epoch via the trainer's `loss_ema_snapshot()` /
`lambda_snapshot()` accessors (mapped-pinned reads; per-epoch
budget, not hot-path).

Single fold-2 re-run on `--cv-fold 2 --cv-n-folds 3
--cv-train-window 4` will surface:
  - which horizon's BCE the controller flagged as hardest each epoch
  - whether lambda saturated at the [1.0, 2.0] ceiling for one horizon
  - whether the lambda trajectory has more variance vs the
    stable-fold runs (suggests the regime drifts during training)

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 20:12:39 +02:00
jgrusewski
0171c8c0ea fix(ml-alpha): asymmetric lambda clamp [1.0, 2.0] — boost-only ISV
3-fold A/B CV (5d42ab0e9 vs eb51c0f9c, same data splits) showed
ISV winning the aggregate (+0.9pt mean_auc, +1.8pt h6000 across
folds) but FOLD-2 regressed -1.3pt on h6000 while folds 0 and 1
both gained (+2.0pt and +4.6pt respectively).

Per-fold per-horizon breakdown showed exactly the failure mode:
  Fold 0 (val 2025-Q1): h6000 no-ISV 0.714 → ISV 0.734 (+2.0pt)
  Fold 1 (val 2025-Q2): h6000 no-ISV 0.682 → ISV 0.728 (+4.6pt)
  Fold 2 (val 2025-Q3): h6000 no-ISV 0.698 → ISV 0.685 (-1.3pt)

In folds 0 and 1, h6000 was the hardest horizon — ISV correctly
boosted it (lambda > 1). In fold 2 the regime made h6000 relatively
easy at no-ISV (0.698 vs the worst horizon at ~0.70). ISV's
SYMMETRIC clamp [0.5, 2.0] then computed ratio = ema_h6000 /
mean_ema < 1 and DEMOTED h6000's trunk-gradient pull below uniform,
starving further learning on the horizon we actually deploy.

Per `pearl_audit_unboundedness_for_implicit_asymmetry.md`: when a
control signal serves an asymmetric goal (here: we never want to
de-prioritize h6000, only ever boost it OR leave it alone), encode
that asymmetry in the clamp. LAMBDA_FLOOR 0.5 → 1.0 makes the
controller boost-only: under-trained horizons get more pull, but
no horizon is ever demoted below its uniform contribution.

Expected effect with asymmetric clamp:
  - Fold 0 and 1: lambda for the hardest horizon stays at 2.0
    (ceiling-clamped), trunk-gradient lift unchanged. The gains
    +2.0pt and +4.6pt should hold.
  - Fold 2: h6000's lambda was being demoted to ~0.74; now floored
    at 1.0 — the -1.3pt h6000 regression should disappear. h6000
    trains at uniform weight, recovering toward 0.698.
  - Cross-fold mean projected: ~0.724 (+1.3pt vs no-ISV).

Test update: `horizon_ema_and_lambda_track_after_training` now
asserts lambda ∈ [1.0, 2.0] (the asymmetric envelope) and
lambda mean ∈ [1.0, 2.0] (boost-only guarantee). Observed values
on the test data: lambda = [1.13, 1.00, 1.08, 1.08, 1.00] — the
two easier horizons correctly floor at 1.00 instead of demoting.

Validation: 7 perception_overfit tests pass, synthetic overfit
trajectory unchanged (0.36 → 0.0006), 26 ml-alpha lib + 23
integration tests green.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 19:16:37 +02:00
jgrusewski
5d42ab0e98 feat(ml-alpha): ISV horizon weighting Phase 3 — wire lambda into trunk grad
Connects the per-horizon lambda computed by horizon_ema_and_lambda
(landed in 37c3a8f4d) to the actual trunk gradient flow. This is the
model-behavior change. mhzs7 (3a196382f) hit mean_auc=0.726 but with
the long-horizon h6000 stuck at 0.694 — exactly the horizon the
auto-horizon-weights formula `min(1, K/h)` over-weights down (h6000
weight = 0.0053). ISV detects "this horizon is hard, give it more
trunk-gradient pull" and lambda[h] scales the per-horizon `d_z`
contribution into the trunk in heads_bwd.

Kernel change (cuda/multi_horizon_heads.cu):
  multi_horizon_heads_backward_batched(): new arg
    `const float* __restrict__ lambda` (5 elements, between
    grad_h_carry and n_batch in arg order).
  - lambda is broadcast into __shared__ float s_lambda[5] once per
    block so every thread reads the 5 floats without repeated
    global loads.
  - Sentinel: zero buffer (init state, EMA hasn't run yet) is read
    as 1.0 so the kernel reduces to pre-ISV behavior. After step 1
    every entry is clamped to [0.5, 2.0] by horizon_lambda.cu and
    the > 0 check is always true.
  - lambda[k] multiplies the `acc += s_lambda[k] * sd_z * w[k]`
    term that produces grad_h (trunk gradient).
  - grad_w and grad_b updates are UNCHANGED. The horizon heads keep
    learning their own weight/bias normally; lambda only biases how
    much each horizon's error signal leaks into the shared trunk.
    Per `pearl_adam_normalizes_loss_weights.md`: scaling the
    effective gradient bypasses Adam's m/sqrt(v) normalization that
    would otherwise cancel a loss-weight lift.

Trainer change (trainer/perception.rs):
  - heads_bwd_batched launch now passes `&self.lambda_d` between
    grad_h_carry_d and n_batch_i. lambda_d is refreshed each step
    by the horizon_ema_and_lambda kernel that ran just after BCE.
  - Whole chain lives inside the captured CUDA Graph.

Validation:
  - All 7 perception_overfit tests pass.
  - Synthetic overfit converges marginally FASTER than pre-Phase-3
    (final loss 0.0007 → 0.0005). Plausible explanation: in a
    constant-signal smoke, the per-horizon BCE values are similar
    enough that lambda mostly stays near 1.0, but the EMA
    bootstrap on step 1 still produces non-uniform initial lambdas
    that nudge the trunk toward whichever horizon converges
    slowest. Either way, no regression.
  - horizon_ema_and_lambda_track_after_training test still passes
    with the lambda values now flowing through heads_bwd:
      ema    = [0.50, 0.87, 0.64, 0.49, 0.60]
      lambda = [0.80, 1.41, 1.03, 0.79, 0.97]
      → h100 hardest (BCE 0.87) gets 1.41× trunk grad; h300 easiest
        (BCE 0.49) gets 0.79× — exactly the intended ISV semantics.
  - 26 lib + 23 integration ml-alpha tests pass.

Honors:
  - feedback_isv_for_adaptive_bounds.md (no hardcoded constants;
    lambda derives from runtime BCE signal).
  - pearl_adam_normalizes_loss_weights.md (scaling trunk gradient,
    not BCE coefficient, to bypass Adam normalization).
  - pearl_audit_unboundedness_for_implicit_asymmetry.md (lambda is
    bounded by horizon_lambda.cu's clamp [0.5, 2.0]).
  - pearl_no_deferrals_for_complementary_fixes.md — Phase 3 wires
    up the infrastructure from Phase 1+2 in the same logical
    delivery rather than waiting a CV cycle.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 16:52:36 +02:00
jgrusewski
37c3a8f4d7 feat(ml-alpha): ISV-driven per-horizon EMA + lambda (Phase 1+2)
Foundation for replacing the static `--auto-horizon-weights` formula
(`min(1, K/h)`) with a signal-driven per-horizon gradient scaler. Per
`feedback_isv_for_adaptive_bounds.md`: adaptive bounds live in ISV,
not hardcoded constants. Per `pearl_adam_normalizes_loss_weights.md`:
Adam normalizes per-loss weight lifts (SP13 saw 13× aux_w produce
only 0.6%/epoch divergence), so the effective lever is scaling the
GRADIENT into the shared trunk, not the BCE coefficient. This commit
sets up the EMA + lambda infrastructure; Phase 3 (wiring lambda into
heads_bwd to actually scale the trunk gradient) is gated on the
3-fold CV results from eb51c0f9c.

Phase 1 — BCE kernel emits per-horizon UNWEIGHTED mean BCE:
  cuda/bce_loss_multi_horizon.cu:
    - New output buffer `loss_per_horizon[N_HORIZONS=5]`.
    - Per-horizon shared-mem accumulators (sloss_h, svalid_h) with
      block tree-reduce — no atomicAdd, per `feedback_no_atomicadd.md`.
    - Hardcoded N_HORIZONS_BCE=5; total shared-mem usage ~13 KiB
      (comfortable under any SM smem limit).
    - The aggregate `loss_out` is still the externally-weighted mean
      callers use for reporting; the new buffer is the UNWEIGHTED
      signal an EMA layer needs.

Phase 2 — EMA + lambda kernel:
  cuda/horizon_lambda.cu (new):
    - Single-thread kernel (5 horizons, fixed-size loop — trivial).
    - First-observation bootstrap via sentinel = 0 per
      `pearl_first_observation_bootstrap.md`; replaces directly when
      `loss_ema_h <= 0` (safer than `== 0` under --use_fast_math).
    - Fixed α = 0.1 EMA for now; Wiener-optimal α follow-up flagged
      (`pearl_wiener_optimal_adaptive_alpha.md`).
    - lambda_h = clamp(loss_ema_h / mean(loss_ema), 0.5, 2.0).
      - Ratio gives natural "under-trained → boost" signal.
      - Clamp prevents winner-take-all per
        `pearl_controller_amplifies_dominant_magnitude_trap.md`
        and bounded-modifier safety per
        `pearl_audit_unboundedness_for_implicit_asymmetry.md`.

Trainer wiring (trainer/perception.rs):
  - 3 new fields: `loss_per_horizon_d`, `loss_ema_d`, `lambda_d`
    (all 5-element f32 CudaSlices; pre-allocated, zero-initialised).
  - Cached `horizon_lambda_fn` + module handle per the
    `BiasKernels`-style pattern (no per-call cuModuleLoadData).
  - BCE callsite (train + eval paths) updated to pass
    `loss_per_horizon_d`.
  - `dispatch_train_step` launches `horizon_ema_and_lambda` right
    after BCE, BEFORE the K-loop backward. Inside the captured
    graph; per-step launch overhead is ~1 µs.
  - `loss_ema_snapshot()` + `lambda_snapshot()` test-only accessors
    (mapped-pinned readback, not for hot path) for diagnostics.

Smoke test — `horizon_ema_and_lambda_track_after_training`:
  - Verifies pre-step EMA + lambda are zero (sentinel).
  - After 5 training steps:
    loss_ema = [0.59, 0.66, 0.45, 0.62, 0.50] — finite + positive.
    lambda   = [1.05, 1.16, 0.80, 1.10, 0.89] — mean ≈ 1.0, all
    inside the [0.5, 2.0] clamp envelope.
  - Lower per-horizon BCE → lower lambda (de-emphasize); higher
    BCE → higher lambda (boost). Exactly the ISV semantics we want.

Validation: 6 perception_overfit tests pass, synthetic overfit still
shrinks (0.33 → 0.0006), 26 ml-alpha lib + 23 integration tests
green. lambda_d is computed every step but NOT YET CONSUMED by
heads_bwd; training behavior is bit-identical to 3a196382f. Phase 3
(consume lambda_d in heads_bwd_batched to scale the per-horizon
gradient into the trunk) follows once CV confirms the foundation is
stable.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 16:48:37 +02:00
jgrusewski
eb51c0f9cd feat(ml-alpha): walk-forward CV via file-list-driven loader
mhzs7 reported val mean_auc=0.726 on 3a196382f — but the trainer
constructed both train and val MultiHorizonLoader with the SAME
`mbp10_root: cli.mbp10_data_dir`. The two loaders only differed by
seed. So val sequences were held-out-by-anchor from the same files
train sampled from; not temporally OOS. Per
`pearl_single_window_oos_is_not_oos.md` a single-window result that
doesn't enforce time-ordered separation can collapse across true
walk-forward folds.

Refactor: drop `mbp10_root` from `MultiHorizonLoaderConfig` (which
forced caller to share the dir between train and val). New API takes
an explicit `files: Vec<PathBuf>` — the loader preserves the order
given and does no internal shuffle, so callers control temporal
ordering. Added `discover_mbp10_files_sorted(root)` helper that
enumerates a dir and sorts by filename (chronological under the
`ES.FUT_<YEAR>-Q<n>.dbn.zst` convention).

alpha_train.rs splits the discovered files by 3 new CLI flags:
  --cv-fold <k>            (default 0)
  --cv-n-folds <N>         (default 1 — single fold)
  --cv-train-window <W>    (default 0 — auto)

Single-fold default (cv_n_folds=1): train on all files except the
last, val on the last file. This replaces the old "same files for
both" bug; even runs that don't think about CV now get a temporal
split by default.

Sliding-window CV (cv_n_folds > 1): fold k trains on files
[k..k+W] and validates on file [k+W]. With 9 quarterly files
(2024-Q1..2026-Q1) and `--cv-n-folds 3`, the natural layout is:

  fold 0: train 2024-Q1..2024-Q4 (W=4) → val 2025-Q1
  fold 1: train 2024-Q2..2025-Q1       → val 2025-Q2
  fold 2: train 2024-Q3..2025-Q2       → val 2025-Q3
  blind holdout: 2025-Q4, 2026-Q1

Threaded the flags through scripts/argo-alpha-perception.sh and
infra/k8s/argo/alpha-perception-template.yaml so each fold submits
as an independent workflow.

Updated tests/multi_horizon_loader.rs to the new API:
  loader_yields_seq_with_valid_labels — exercises discover + load.
  loader_errors_on_empty_files       — replaces missing-root test.
  discover_errors_on_missing_root    — pinpoints the discover step.

Honors:
  - feedback_no_partial_refactor.md — every consumer migrated atomically.
  - feedback_no_legacy_aliases.md — no `mbp10_root` shim left behind.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 16:36:48 +02:00
jgrusewski
3a196382f0 fix(ml-alpha): captured graph poisoned eval via SyncOnDrop events
z2w9w cluster run hit CUDA_ERROR_INVALID_VALUE at "eval snap_batched
fwd" the first time validation ran after a captured training step.
Training itself succeeded (epoch 0 train_loss=0.69 over 250 captured
graph replays); only the subsequent direct eval kernel launch failed.

Root cause (vendor/cudarc/src/driver/safe/core.rs:920):
  CudaSlice::device_ptr_mut() returns a `SyncOnDrop::Record` guard.
  On drop, that guard UNCONDITIONALLY calls `event.record(stream)` on
  the slice's `.read` event (the check at line 953 only gates the
  cuStreamWaitEvent on .write — the unconditional event.record at the
  end runs no matter what). Inside a stream-capture region, those
  event.record(stream) calls turn the CudaEvents into "captured
  events" per the CUDA Driver API. Captured events can ONLY be waited
  on by streams in the same capture sequence; any later
  cuStreamWaitEvent from outside fails with CUDA_ERROR_INVALID_VALUE.

  The trainer's eval path then called `device_ptr_mut()` again to
  stage the eval DtoDs — which inserted exactly that
  cuStreamWaitEvent on the now-captured `.write` event of every
  trainer CudaSlice. First kernel launch after the dtods failed.

Why this hit z2w9w now: a) all our work is on a SINGLE stream
(`self.stream`), so the event-based multi-stream sync that cudarc
inserts is pure overhead, b) the capture region is exactly where
those overhead events become poisonous.

Fix: disable cudarc's read/write event tracking BEFORE the trainer
allocates ANY device memory. With tracking off at alloc time,
CudaSlice::new returns `read: None, write: None` (core.rs:1283).
SyncOnDrop::record_event with `event: None` produces a `Record(None)`
that does nothing on drop. launch_builder skips its event waits/
records too. The capture region runs clean; the eval direct launches
have no stale captured events to wait on.

Validation:
  - 6 perception_overfit tests pass, including 3 NEW regression tests
    that pin the exact failure modes:
    * evaluate_alone_succeeds — eval with no prior training
    * evaluate_works_after_warmup_only — eval after 1 uncaptured step
    * evaluate_works_after_capture_no_replay — eval right after capture
      (the minimal repro that pinpointed `device_ptr_mut` inside capture)
    * evaluate_works_after_captured_training_step — full warmup +
      capture + replay + eval
  - 26 ml-alpha lib + 23 ml-alpha integration + 306 ml-core lib all pass.

Also drops the now-redundant disable/enable_event_tracking dance
around the capture region — events are globally disabled for the
trainer's lifetime so no per-capture flipping needed.

Honors: feedback_no_quickfixes.md (root-cause traced through
cudarc's safe wrappers to the SyncOnDrop record contract, not a
symptom-suppress sleep/retry hack).

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 16:01:02 +02:00
jgrusewski
87ca1f5f55 perf(ml-alpha): CUDA Graph capture of training step (#162 finale)
Captures the full Mamba2 fwd → K-loop fwd → BCE → K-loop bwd →
Mamba2 bwd → AdamW × 7 → loss DtoD chain into a single CUDA Graph.
First step_batched call runs uncaptured (cuBLAS warmup); second
call captures; third+ replays. Replaces ~155 individual kernel
launches per step with one graph launch.

Four root-cause issues had to be fixed in concert to make the
capture region capture-compatible:

1. cuBLAS lazy workspace allocation
   crates/ml-alpha/src/mamba2_block.rs — pre-allocate an 8 MiB
   workspace via cublasSetWorkspace_v2 at Mamba2Block::new.
   cuBLAS would otherwise allocate on first gemm call with each
   new shape, breaking capture.

2. cuBLAS heuristic plan-cache lookup
   crates/ml-core/src/cuda_autograd/linear.rs — switch
   gemm_ex_f32 from CUBLAS_GEMM_DEFAULT_TENSOR_OP to
   CUBLAS_GEMM_DFALT. The heuristic algo path triggers
   plan-cache allocs; DFALT is deterministic with negligible
   perf delta for our shapes (Mamba2: 128×32, 128×state_dim).

3. Per-call cuModuleLoadData in bias kernels
   crates/ml-core/src/cuda_autograd/linear.rs — add a
   `BiasKernels` struct (add_bias_2d_kernel + reduce_sum_axis0
   handles) cached at construction. `BiasKernels::shared(stream)`
   uses a per-context OnceLock cache so the cubin loads exactly
   once per CUDA context for the process lifetime. Every
   `GpuLinear` and `OwnedGpuLinear` constructor now stores its
   `BiasKernels`. Removed the per-call `get_bias_kernels`
   helper entirely (no legacy aliases — greenfield).

4. Per-call CudaSlice::clone() in Mamba2 fwd + bwd
   crates/ml-alpha/src/mamba2_block.rs — three sites cloned
   the input slice to build a fresh `GpuTensor` view. Each
   `CudaSlice::clone()` does cuMemAlloc + dtod copy
   (vendor/cudarc/src/driver/safe/core.rs:1437); cuMemAlloc
   is forbidden during capture.

   Refactored `forward_with_slices_into` and
   `backward_with_slices_into` to take raw `x_data: &CudaSlice<f32>,
   batch: usize` instead of `&GpuTensor` / `&LinearActivations`.
   All three Mamba2 fwd call sites + three bwd call sites now
   pass the underlying CudaSlice directly.

Trainer changes (trainer/perception.rs):
  - Three-state machine in step_batched: warmup → capture → replay.
  - Labels staging fill moved BEFORE the captured region (host writes
    only; replays read whatever the host last wrote).
  - Removed the post-snap_batched sync that was inside dispatch (it
    would trip STREAM_CAPTURE_ISOLATION; kernels are stream-ordered
    so the sync was unnecessary).
  - Dispatch extracted into `dispatch_train_step` method; the
    captured region brackets exactly this method.
  - Final sync + mapped-pinned loss read happens once per step in
    step_batched, outside the captured region.

Validation:
  - Synthetic overfit smoke: 0.397 → 0.0007 (matches pre-refactor
    trajectory; capture+replay produces equivalent loss).
  - 26 ml-alpha lib tests + 23 integration tests pass.
  - 306 ml-core lib tests pass.

Also fixed (orthogonal but required for green workspace):
  crates/ml-core/src/action_space.rs — tests assumed 7-exposure
  layout (63 actions). Production `ExposureLevel` enum has 8
  variants (Hold inserted at idx 3, Flat moved to idx 7 → 72
  total actions). Updated tests + `get_valid_action_mask` to
  match the 8-level layout.

Honors:
  - feedback_no_legacy_aliases.md (no `add_bias_2d_with_fn` /
    legacy `get_bias_kernels` fallback; one canonical API).
  - feedback_no_partial_refactor.md (signature change propagated
    to every caller atomically; no half-migrated state).
  - feedback_wire_everything_up.md (BiasKernels wired into all
    GpuLinear/OwnedGpuLinear constructors in same commit).

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 15:21:00 +02:00
jgrusewski
ab94ce2a49 perf(ml-alpha): device-resident AdamW step counter (capture prep)
Stage 1+2 of #162 (CUDA Graph capture of training step). The AdamW
kernels previously took the step counter as a host scalar arg, which
gets baked into kernel args at CUDA Graph capture time — replays would
freeze the counter and produce wrong bias-correction values.

Both AdamW variants now read the step from a device pointer, advanced
by a tiny 1-thread `increment_counter` kernel that goes inside the
captured region. Each replay correctly increments and observes the
new step value.

Kernel changes:
  adamw_step.cu:
    - adamw_step:                int step → const int* step_ptr
    - adamw_increment_counter:   new, +=1 on step_ptr[0]
  mamba2_alpha_kernel.cu:
    - mamba2_alpha_adamw_step_devscale: int t → const int* step_ptr
    - mamba2_alpha_increment_step_counter: new

Rust changes:
  trainer/optim.rs (AdamW):
    - host `step: i32` → device `step_count_d: CudaSlice<i32>`
    - step(): launch increment kernel BEFORE adamw kernel; both read
      device counter via pointer arg.
    - step_count(): test-only accessor, mapped-pinned readback (sync).

  mamba2_block.rs (Mamba2AdamW):
    - kept host `step_count: i32` for legacy paths (`step`,
      `step_from_buffers`) which aren't capture-compatible anyway
      (host grad-norm dtoh, host scalar grad_scale).
    - added device `step_count_d: CudaSlice<i32>` for the production
      gpu_clip path; advances via `kernel_increment_step` kernel
      inside the captured region.
    - adamw_apply_devscale: `t: i32` → `step_d: &CudaSlice<i32>`.

Validation:
  - 4 adamw_invariants tests pass (step_count_increments specifically
    exercises the device counter).
  - 10 mamba2_block lib tests pass (training_loop_decreases_loss
    exercises legacy host-counter path).
  - Synthetic overfit smoke: initial=0.25 → final=0.0006 (matches
    pre-refactor trajectory bit-for-bit-equivalent).

Stage 3+4 (capture brackets + first-call-capture / subsequent-replay
state machine in step_batched) follows in the next commit.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 14:39:38 +02:00
jgrusewski
b6fb720acd perf(ml-alpha): eliminate all dtoh from training hot path
GPU-resident grad-norm + clip-scale; mapped-pinned loss readback.
Replaces 9× memcpy_dtoh per Mamba2 AdamW step (grad-norm host roundtrip)
+ 1× per-step download() (loss). Saves ~10 stream-sync barriers/step.

New kernels (cuda/grad_norm.cu):
  - grad_norm_sq_phase1: per-block tree-reduce of x[i]^2 (no atomicAdd)
  - grad_norm_sq_phase2: cross-tensor accumulator (sequential stream-ordered)
  - grad_clip_scale: writes min(1, max_norm/sqrt(norm_sq)) to device ptr
  - mamba2_alpha_adamw_step_devscale: reads grad_scale from device pointer
    instead of host scalar, allowing AdamW kernels to launch async without
    waiting for a CPU-side norm computation.

Trainer changes (perception.rs):
  - loss_d kept device-side; mapped-pinned MappedF32Buffer shadow.
  - Single stream.synchronize() at end of step (was 2: post-bwd + download).
  - DtoD copy loss_d → loss_host_d queued, then sync flushes both kernels
    + copy in one barrier. Loss read via host_ptr (no dtoh).

Mamba2 AdamW (mamba2_block.rs):
  - step_from_buffers_gpu_clip(): all grad-norm tensors processed via
    phase1+phase2 chain, scale computed on-device, AdamW launches with
    devscale variant. Zero host roundtrips.
  - Pre-allocated block_partials_d, grad_norm_sq_d, grad_scale_d.

Optimizer (optim.rs): removed redundant stream.synchronize() per AdamW step.
Each per-tensor AdamW kernel is stream-ordered; sync only needed before
host reads, which the trainer handles centrally.

Synthetic overfit smoke: initial=0.30 → final=0.0006 (matches pre-refactor
trajectory). Full ml-alpha test suite passes (45 tests across lib +
integration).

Honors:
  - feedback_no_htod_htoh_only_mapped_pinned.md (MappedF32Buffer only)
  - feedback_no_atomicadd.md (block tree-reduce only)
  - feedback_no_legacy_aliases.md (step_from_buffers replaced, not aliased)

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 13:57:02 +02:00
jgrusewski
eb0e4b6328 perf(ml-alpha): full zero-alloc training step (#3 foundation)
Eliminates ALL per-step allocations from the training hot path —
foundation for CUDA Graph capture (next commit). Before this commit,
each step_batched call allocated:

  Mamba2 forward:    input_2d view, x, a_proj, b_proj, h_s2, h_enriched_seq
  Mamba2 backward:   d_a_per_channel/d_b_per_channel/d_w_c/d_h_s2 (#2 covered)
                     d_a_proj_flat, d_b_proj_flat, dw_c
                     LinearGrads.{dw,db,dx} × 3 projections (cuBLAS internal)
                     d_x_from_a + d_x_from_b + d_x (elementwise add)
                     dw_out, db_out (zero-init shells)
  Trainer wrapper:   window_tensor, h_enriched_seq_t, grad_h_enriched_seq_t,
                     grad_h_enriched_seq

~20-25 cudaMalloc / GpuTensor::zeros calls per step × 2000 steps/epoch =
40-50K allocations per epoch.

This commit adds zero-alloc `_into` variants throughout the chain:

  ml-core/cuda_autograd/linear.rs:
    OwnedGpuLinear::forward_with_slices_into
    OwnedGpuLinear::backward_with_slices_into
    reduce_sum_axis0_into

  ml-core/cuda_autograd/elementwise.rs + gpu_tensor.rs:
    ElementwiseKernels::binary_into
    GpuTensor::add_into

  ml-alpha/mamba2_block.rs:
    Mamba2BlockForwardScratch (pre-allocated forward cache)
    Mamba2BackwardGradsBuffers (pre-allocated backward outputs)
    Mamba2Block::forward_train_seq_into (zero-alloc forward)
    Mamba2Block::backward_from_h_enriched_seq_full_into (zero-alloc backward)
    Mamba2AdamW::step_from_buffers (reads grads_buffers directly)

  ml-alpha/trainer/perception.rs:
    PerceptionTrainer pre-allocates: window_tensor_d, h_enriched_seq_t_d,
      grad_h_enriched_seq_t_d, grad_h_enriched_seq_d, mamba2_fwd_scratch,
      mamba2_grads_buffers
    step_batched + evaluate_batched fully wired through _into variants

Original `forward_with_slices` / `backward_with_slices` / `binary` / `add` /
`backward_from_h_enriched_seq` paths preserved unchanged — Phase E.3
ml/examples callers unaffected.

The captured-graph commit (next) only needs to wrap this zero-alloc
training step in cuGraph capture/replay; no further refactoring of
buffer management.

77 ml-alpha tests pass. Synthetic overfit converges identically
(0.29 → 0.0007 in 250 steps) — gradients are bit-identical to the
allocating path.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 13:34:56 +02:00
jgrusewski
c70c5cdf21 perf(ml-alpha): fused batched snap_feature_assemble kernel (#4)
Previously the per-step snap_feature path did B*K = 768 single-snapshot
kernel launches (at B=8, K=96) + 768 DtoD copies into the window
tensor. New `snap_feature_assemble_batched` processes all B*K
snapshots in a SINGLE launch and writes outputs directly into the
window tensor's storage.

Per-step CPU work: pack 12 mapped-pinned staging buffers (~150 KB
total host writes), then 10 DtoD copies of the staging → device
buffers. Per-step GPU work: 1 batched kernel launch with B*K
threads (each writes 32 floats to its output row).

Mapped-pinned staging buffers cover the full B*K capacity at trainer
init — no per-step allocation. New `MappedI32Buffer` and
`MappedI64Buffer` types parallel `MappedF32Buffer` to stage
`trade_count` (i32) and `ts_ns` / `prev_ts_ns` (i64) without
violating the no-htod rule (`feedback_no_htod_htoh_only_mapped_pinned.md`).

Dead per-snapshot scratch + helpers (`bid_px_d`, `snap_feat_d`,
`stg_bid_px`, `snap_fn`, `upload_into`, etc.) removed per
`feedback_no_legacy_aliases.md` — the only callers were the
per-snapshot path, gone.

Expected per-step savings: ~5-10 ms launch + DtoD overhead at
B=8, K=96. Over 2000 steps/epoch = 10-20 sec/epoch.

77 ml-alpha tests pass. Synthetic overfit unchanged.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 13:12:32 +02:00
jgrusewski
ebae67cb6b perf(ml-alpha): pre-allocate Mamba2 backward scratch (#2)
The four biggest per-call scratch buffers in
Mamba2Block::backward_from_h_enriched_seq were allocated fresh on
every training step:

  d_a_per_channel  [N, sh2, K, state_d]  ~6 MB at B=8, K=96
  d_b_per_channel  [N, sh2, K, state_d]  ~6 MB
  d_w_c_per_sample [N, sh2, state_d]     ~64 KB
  d_h_s2           [N, sh2]              ~4 KB

For 2000 optimizer steps/epoch × 15 epochs = 30 000 alloc_zeros
calls per training run, all on the hot path.

New `Mamba2BackwardScratch` struct holds these as device-resident
buffers, constructed once per (n_batch, seq_len, hidden_dim,
state_dim) at trainer init. New `backward_from_h_enriched_seq_into`
method takes the scratch by &mut and reuses the buffers each call.

The smaller per-call buffers (d_a_proj_flat, d_b_proj_flat, dw_c)
still allocate per call — they feed into ownership-transferring
LinearGrads outputs, where pre-allocation would require refactoring
ml-core's cuBLAS wrappers without proportional gain.

The original `backward_from_h_enriched_seq` is preserved (Phase E.3
callers still use it). Trainer switches to the `_into` variant.

Expected per-step savings: ~10-20 ms on L40S (4 cudaMalloc latencies
per call × 2.5-5 μs each + cache pressure reduction). Over 2000
steps/epoch that's 20-40 sec/epoch.

77 ml-alpha tests pass. Synthetic overfit unchanged (0.27 → 0.0006).

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 13:00:42 +02:00
jgrusewski
3f089210f3 perf(ml-alpha): preload all MBP-10 files into RAM, reset loaders per epoch
cnjfl wall-time analysis: 80% of training time was disk IO. The
MultiHorizonLoader was constructed fresh per epoch in the CLI loop,
forcing 9 file deserializations from bincode (~50s/file × 9 = 7-8
min) per epoch. For a 6-epoch run that was ~45 min of pure IO out
of ~50 min total.

Now loaders preload ALL files into RAM at construction (one-time
~7 min startup) and expose a `reset(seed: u64)` method that
re-seeds anchor sampling per epoch — no disk IO between epochs.

Memory cost: ~13-15 GB for the 9-quarter ES.FUT dataset (45M
snapshots × ~280 bytes). Well under the training pod's 64 GB
limit; current 16 GB request remains sufficient since the resident
set fits.

Expected wall-time impact at K=96, B=8, 16K seqs:
  6 epochs:  ~50 min → ~13 min  (~3.7×)
  15 epochs: ~120 min → ~23 min (~5×)

The internal LoadedFile-cache-with-cycling logic is gone — the
loader now holds Vec<LoadedFile> with all files resident. Per-call
`next_sequence` picks a uniformly random file + uniformly random
anchor inside it (instead of cycling files with a per-file budget).
Distribution is equivalent: each file contributes ~n_max_sequences /
n_files samples per epoch in expectation.

77 ml-alpha tests pass.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 11:57:07 +02:00
jgrusewski
26d91a816c feat(alpha_train): configurable early-stop metric (default mean_auc)
cnjfl evidence: val_loss and mean_auc disagree.
  val_loss best at e3 (0.5592)
  mean_auc best at e4 (0.7670 — new h300 + h6000 peaks)

For downstream trading, ranking quality (AUC) matters more than
probability calibration (BCE loss). New default is mean_auc-based
early stopping, but val_loss/none remain selectable.

AUC is noisier than loss epoch-to-epoch (1-2pt bounces are common
even when long-horizon AUCs are still drifting up under
auto-horizon-weights), so patience defaults bump from 3 → 5.

CLI:  --early-stop-metric {val_loss|mean_auc|none}   default mean_auc
      --early-stop-patience N                         default 5

Argo template parameters added with matching defaults.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-17 11:48:21 +02:00