1fb2decb7928d0a92d79b7f659dcfca6f118391a
5200 Commits
| Author | SHA1 | Message | Date | |
|---|---|---|---|---|
jgrusewski
|
1fb2decb79 |
feat(ml-backtesting): resting orders + stops + OCO + audit ring (C11)
Picks up the deferred work from C5/C7 trim notes. Adds:
cuda/lob_state.cuh — LimitSlot (32 B) + StopSlot (32 B) + Orders
(limits[32] + stops[16] = 1536 B/backtest). u64-aligned with
explicit _pad[6] on both slot structs so the Rust mirrors
(LimitSlotFlat / StopSlotFlat) bytemuck::Pod-derive without
panics about implicit padding.
cuda/resting_orders.cu — three kernels:
resting_orders_step (runs per event after book_update):
1. In-flight → resting promotion at arrival_ts_ns. Queue position
initialised pessimistically (full level depth ahead) per spec §9.
2. Queue decay against same-side trade_signed_vol at level: ahead-
of-us first then us; partial-fill emits OrderEvent + pos update.
3. Marketability check: book moved through our price → cross at
the just-arrived book (slippage-aware fill walks levels).
IOC cancels remainder; FOK does too (partial-fill not allowed).
4. Stop trigger detection: best ask up for buy-stop, best bid
down for sell-stop. StopMarket walks book; StopLimit converts
to a new LimitSlot (overflow → submission_overflow++).
5. OCO mutual exclusion: oco_pair byte (0..31 = limit, 32..47 =
stop) — when one leg fills/triggers, the paired slot is freed.
seed_limit_slot / seed_stop_slot — host-side single-slot seeders
for fixture testing + as a v1 entry point until the decision
kernel learns to emit resting orders. Both set an overflow_flag
if the target slot is already in use (gen_counter bumped on
successful seed for SlotTag freshness).
src/sim.rs — wires three new methods:
seed_limit_order(b, slot, side, kind, active_state, oco_pair,
price, size, queue_position_init, arrival_ts_ns)
seed_stop_order(b, slot, side, kind, active_state, oco_pair,
trigger_price, limit_price, size, arrival_ts_ns)
step_resting_orders(current_ts_ns, trade_signed_vol)
→ launches resting_orders_step kernel + chains step_pnl_track
so any fill that closes a position emits a TradeRecord.
Plus read_limit_slot / read_stop_slot / read_audit_records for
fixture inspection.
Audit-ring buffers (deferred from C2 originally) now allocated:
audit_d (n × 256 × 24 B) + audit_head_d (n × u32). Both
resting_orders.cu and the decision kernel's WriteOrder path
populate it via the inlined pack_slot_tag + emit_audit helpers.
Four new GPU fixtures:
limit_rest_marketable_fill — resting buy at 5500 with book at
ask 5500 fills immediately on step_resting (3 lots, vwap=5500).
stop_trigger — buy 4 lots @ 5500, seed sell-stop at trigger=5495,
book moves so bid=5494.50: stop triggers, walks book, position
closes flat, TradeRecord emitted, stop slot active=0.
oco_one_cancels_other — pair {buy@5495, sell@5505} with oco_pair
cross-linked: book moves so bid=5505.50, sell leg fills (pos=-2),
buy leg auto-cancelled. Both slots end active=0.
submission_overflow — host loop fills all 32 limit slots, then 33rd
seed_limit_order call must return Err (slot 0 already in use).
All 10 GPU fixtures green on RTX 3050 (6 original + 4 new). Ring 2
fuzz at N ∈ {1, 16, 256} still green. 33 lib unit tests still green.
The decision kernel (C7) and submit_market_immediate (C5) paths
are unchanged — they continue to use the "immediate fill" route.
A follow-up commit can teach the decision kernel to emit resting
limits via submit_limit_alloc (already defined in resting_orders.cu).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
771faac723 |
test(ml-backtesting): Ring 2 invariant fuzz at N ∈ {1, 16, 256} (C10)
Property-based fuzz tests with random-walk MBP-10 sequences applied to
the LOB simulator at three backtest counts; assert per-snapshot
invariants that must hold regardless of input or block scheduling:
fuzz_n1_book_only (200 events, no orders)
Pure book-update kernel — verifies mid-drift random walks preserve
book monotonicity and never produce a crossed book.
fuzz_n16_with_orders (200 events, market orders every 8th event)
16 backtests in parallel, each submitting random buy/sell market
orders 1-3 lots at every 8th event. Asserts book invariants on
each backtest's state PLUS:
- position_lots stays within ±30 (plausible given fixture book depth)
- realized_pnl + vwap_entry finite (no NaN/Inf leaks)
fuzz_n256_with_orders (100 events, market orders)
Production-scale parallelism. Same invariants as N=16. Each block
has its own per-backtest Pos + OpenTradeState + TradeLog, exercising
the per-block isolation discipline established in C5-C7.
All 3 pass on RTX 3050. Spec §8 Ring 2 confidence gate hit.
Adds rand + rand_chacha to ml-backtesting dev-dependencies.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
63ed6d0217 |
feat(ml-backtesting): artifacts + sweep aggregate + fxt-backtest CLI (C9)
artifacts.rs:
- Summary struct (total_pnl_usd, sharpe_ann, sortino_ann,
max_drawdown_usd, calmar, n_trades, win_rate, avg_win/avg_loss,
profit_factor, total_fees_usd, exposure_pct,
kelly_cap_history_sample).
- compute_summary(records, pnl_curve_usd) — non-overlapping
annualisation × √825 per pearl_phase1d4_backtest_cost_edge_frontier
(K=6000 holding × 250 trading days ≈ 825 trades/year).
Sharpe + Sortino + max drawdown + Calmar.
- write_summary (JSON pretty-printed), write_trades_csv (with USD
conversion from fp ×100), write_pnl_curve_bin (bytemuck-cast f32
slice). 5 unit tests with tempdir.
aggregate.rs:
- aggregate_sweep_dir walks <root>/<cell>/summary.json, builds an
arrow RecordBatch (cell name + 9 stats columns), writes
SNAPPY-compressed aggregate.parquet.
- pareto_frontier: cells are kept unless another cell weakly
dominates on all three of (sharpe_ann maxed, max_drawdown_usd
minimised, total_fees_usd minimised) AND strictly improves on one.
Written to pareto_frontier.json (Vec<cell-name>).
- 2 unit tests (3-cell mutual-non-dominance; B-dominates-A).
harness.rs:
- run() now samples Pos.realized_pnl × $50/index-pt per event into
self.pnl_curves[b], so the per-cell P&L curve is ready for
write_artifacts() without an extra sim pass.
- write_artifacts(out_dir) — per-cell <out>/cell_NNNN/{summary.json,
trades.csv, pnl_curve.bin}.
bin/fxt-backtest:
- clap-derive CLI with two subcommands:
run --data <dir> [--predecoded-dir <dir>] [--policy-grid <yaml>]
[--n-parallel N] [--decision-stride S] [--latency-ns N]
[--target-annual-vol-units F] [--annualisation-factor F]
[--max-lots N] [--max-events N] [--seed N] --out <dir>
aggregate <sweep_dir>
- Constructs MlDevice::cuda(0) + CfcTrunk::new_random for the trunk
(v1 — ml-alpha has no checkpoint format yet; --seed gates init).
- Parses --policy-grid YAML if given but doesn't yet plumb to the
LobSimCuda decision kernel (the v1 kernel hardcodes the
Strategy::default_for path; bytecode VM is C7's deferred follow-up).
Parse step kept end-to-end so the YAML format is validated now.
- --latency-ns parsed but not consumed — reserved for follow-up
resting-order in-flight promotion (deferred from C5).
Adds parquet + arrow + arrow-array + arrow-schema + serde_yaml to
ml-backtesting deps; bin/fxt-backtest added to workspace members.
All 33 lib tests + 6 GPU fixture tests green. CLI --help renders both
subcommands correctly.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
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>
|
||
|
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>
|
||
|
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>
|
||
|
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>
|
||
|
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> |
||
|
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>
|
||
|
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>
|
||
|
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> |
||
|
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>
|
||
|
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> |
||
|
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> |
||
|
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> |
||
|
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> |
||
|
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> |
||
|
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> |
||
|
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>
|
||
|
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
(
|
||
|
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> |
||
|
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>
|
||
|
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>
|
||
|
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>
|
||
|
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>
|
||
|
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> |
||
|
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>
|
||
|
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> |
||
|
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>
|
||
|
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> |
||
|
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> |
||
|
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> |
||
|
jgrusewski
|
59e236b4e8 | feat(ml-alpha): 2-layer heads backward kernel with ISV lambda (Phase 1.4) | ||
|
jgrusewski
|
e0a497da9e | feat(ml-alpha): 2-layer GELU MLP heads — forward kernel (Phase 1.3) | ||
|
jgrusewski
|
167f065647 | feat(ml-alpha): LayerNorm backward + per-row param-grad reducer kernels | ||
|
jgrusewski
|
d8cd90c130 | feat(ml-alpha): LayerNorm forward kernel for trunk-pre-CfC normalisation | ||
|
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>
|
||
|
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>
|
||
|
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 ( |
||
|
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>
|
||
|
jgrusewski
|
0171c8c0ea |
fix(ml-alpha): asymmetric lambda clamp [1.0, 2.0] — boost-only ISV
3-fold A/B CV ( |
||
|
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 |
||
|
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 |
||
|
jgrusewski
|
eb51c0f9cd |
feat(ml-alpha): walk-forward CV via file-list-driven loader
mhzs7 reported val mean_auc=0.726 on
|
||
|
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>
|
||
|
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>
|
||
|
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> |
||
|
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>
|
||
|
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> |
||
|
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> |