fdef6efe98f2b92ee91788ffc3dd32fc8187fbe8
2985 Commits
| Author | SHA1 | Message | Date | |
|---|---|---|---|---|
jgrusewski
|
fdef6efe98 |
refactor(ml-alpha): move VSN weights from PerceptionTrainer to CfcTrunk
PerceptionTrainer gains a `trunk: CfcTrunk` field constructed at the top of `new()` (after the determinism seed guard). VSN's weight tensors (`vsn_w_d`, `vsn_b_d`) now live on `self.trunk`; the trainer-owned copies are removed. Forward + backward + AdamW access sites redirect to `self.trunk.vsn_w_d` / `self.trunk.vsn_b_d`. PRNG-state preservation: the trainer's ChaCha8Rng chain still draws VSN values at the same call position as before, then memcpy_htod's them into the trunk's zero-initialised VSN slots from X1. This keeps every downstream weight (attn_q, ...) bit-identical to the pre-X2 layout — perception_forward_golden continues to pass with max_diff = 0.000000. Gradient buffers (`grad_vsn_w_d`, `grad_vsn_b_d`) and AdamW state (`opt_vsn_w`, `opt_vsn_b`) remain at trainer level — they're training-only and shouldn't move to the inference trunk. Verification: - perception_forward_golden: PASS (max_diff = 0.000000) - ml-alpha lib tests: 34 pass - alpha_train example builds clean Per spec 2026-05-19-ml-alpha-v2-trunk-grows-and-deployability-design.md §1.1, §2.2 (X2). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
e338000eec |
feat(ml-alpha): CfcTrunk v2 weight skeleton (no callers yet)
Adds zero-initialised v2 weight tensors (VSN, LN_a/b, attn_q, GRN heads) and Option<Mamba2Block> slots for stacks 1 and 2 to CfcTrunk. Extends CfcConfig with mamba2_state_dim (default 16, matches PerceptionTrainerConfig::default). Imports HEAD_MID_DIM for GRN head sizing. No forward path changes — fields allocated, not yet read. Existing new_random init for V1 fields (CfC + simple heads + projection) is preserved unchanged so the trunk's forward kernels still produce the same output. Skeleton for X2..X9 weight-group migrations. Verification: - ml-alpha lib tests: 34 pass - perception_forward_golden bit-equivalence: PASS (max_diff = 0.000000) Per spec 2026-05-19-ml-alpha-v2-trunk-grows-and-deployability-design.md §1.1, §2.2 (X1). |
||
|
jgrusewski
|
b47b2fabfb |
fix(ml-core): deterministic GPU weight init via scoped_init_seed
ml_core::cuda_autograd::init::generate_uniform (backing xavier_uniform,
kaiming_uniform, bias_uniform, near_zero_xavier) defaulted to seeding
from SystemTime::now() + thread_id, producing non-reproducible weights
across processes. Mamba2 stacks initialise via OwnedGpuLinear::xavier,
which routes through this helper — so PerceptionTrainer.evaluate output
diverged 5-30% across fresh-process runs with identical cfg.seed.
Fix: thread-local seedable RNG override. New API:
let _g = ml_core::cuda_autograd::init::scoped_init_seed(seed);
// ... all xavier/kaiming/bias/near_zero calls draw from
// StdRng::seed_from_u64(seed) chain while _g is alive ...
// _g dropped here -> restores default time-based seeding
PerceptionTrainer::new now installs the guard before any Mamba2Block
construction, so the trainer is reproducible from cfg.seed end-to-end.
CfC/VSN/heads already used explicit ChaCha8Rng::seed_from_u64 — only
Mamba2 was affected.
Production behavior unchanged when no guard is set. ml-core: 306 tests
pass, ml-alpha: 34 lib tests pass.
Regression test: crates/ml-alpha/tests/perception_forward_golden.rs
captures bit-exact PerceptionTrainer.evaluate output (loss + 160 probs
on a deterministic seed=42 fixture) into a 644-byte golden file.
Three consecutive runs now produce max_abs_diff=0; pre-fix runs varied
by 0.1-0.3 absolute on individual probs.
.gitignore: added exception for crates/ml-alpha/tests/fixtures/*.bin
so deterministic test fixtures land in repo.
Per pearl_scoped_init_seed_for_reproducibility in project memory.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
e004d6c217 |
Revert "arch(ml-alpha): restore count-delta redundancy at feature slot [26]"
This reverts commit
|
||
|
jgrusewski
|
004b662c80 |
obs(ml-alpha): per-N-step train_loss log line for liveness + intra-epoch monitor
Adds a single tracing::info!("step") emitted every 500 training steps:
if epoch_train_steps % 500 == 0 {
tracing::info!(epoch, step = epoch_train_steps, loss = loss, "step");
}
Consumed by /tmp/alpha_monitor.py (v2: per-epoch trajectories + ISV +
liveness) to render intra-epoch loss trajectory and detect stalls
faster than the once-per-epoch granularity allowed.
At 8000 steps/epoch and ~36s/epoch on L40S → ~16 step lines per epoch
per fold, well below the cluster log volume budget.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
008f65d894 |
arch(ml-alpha): restore count-delta redundancy at feature slot [26]
The C16 tick-rule swap (
|
||
|
jgrusewski
|
17eb825113 |
arch(ml-alpha): σ becomes sidecar — BCE drops Kendall damping (upgraded path)
Empirical record at this data + this architecture:
Phase 1+2+3 (no σ in loss): mean_auc 0.7749 ± 0.024
v2 (σ + axes B/C/D/E): mean_auc 0.7541 ± 0.005
σ-only (kept σ, dropped B/C/D/E): mean_auc 0.7506 ± 0.008
Perf-fix (σ-only math): mean_auc 0.7499 ± 0.010
ISV-σ (closed-form σ + adaptive λ): mean_auc ~0.75 (2/3 folds)
Every architecture with σ-in-loss lands at 0.75. Removing σ is the
only thing that hits 0.77. That is a framework mismatch, not a tuning
problem.
Kendall+Gal+Cipolla 2018 frames σ as TASK NOISE level — damp the noisy
task, trust the clean one. Our horizons don't have different label
noise; they have different intrinsic difficulty (longer horizon = more
price-walk uncertainty = lower achievable AUC). σ-Kendall sees "high
BCE on h6000" and interprets it as "h6000 is unreliable, back off" —
precisely the opposite of what we want. h6000 is the deployment
target; damping it is a self-inflicted wound. With mean_bce ~ 0.7,
σ ≈ √0.7 ≈ 0.84 ⇒ w_h ≈ 0.71, uniformly attenuating gradient by
~30% across every horizon. λ's z-score boost (max 2×) can rebalance
relative-per-horizon but cannot recover the absolute magnitude.
Per-horizon prioritization remains via the ISV-driven λ controller
(grad scaler in heads_grn_bwd, per pearl_adam_normalizes_loss_weights).
That controller IS appropriate for our problem: it boosts hard
horizons rather than damping them, and it operates on the gradient
into the trunk rather than on the loss aggregate (Adam-cancellation
safe).
Changes to bce_loss_multi_horizon.cu (six lines):
w_h = bw (was: 0.5 * bw * exp(-2 * log_sigma_h))
d_log_sigma_h[h] = 0.0 (was: 1 - 2 * w_h * mean_bce)
total_loss += bw * mean_bce (was: + w_h * mean_bce + log_sigma_h)
σ infrastructure preserved unchanged:
- horizon_ema_and_lambda still computes log_sigma_h closed-form from
loss_ema (Kendall equilibrium) for telemetry / future label-noise
estimation use cases
- log_sigma_h kernel arg still in BCE signature (zero churn at
callsite); ignored inside
All 9 perception_overfit tests pass — including
horizon_ema_and_lambda_track_after_training which validates the
per-horizon controller end-to-end through 64 K-loop iterations of
capture/replay.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
410ab6b0ea |
arch(ml-alpha): ISV-driven σ + adaptive Z_SCALE — both controllers anchor on loss_ema
Per pearl_controller_anchors_isv_driven, every controller anchor/target/
cap derives from a tracked signal, not hardcoded constants. The σ-only
revert kept Kendall σ as a free Adam-learned scalar — that violated ISV
discipline and fought Adam's m/√v normalization
(pearl_adam_normalizes_loss_weights).
Single source of truth for both per-horizon controllers:
log_sigma_h[h] ← max(log(0.5), 0.5 * log(loss_ema[h]))
Kendall equilibrium (∂L/∂log σ = 0 ⟹ σ_h² = mean_bce_h)
in closed form. Asymmetric floor at log(0.5) prevents
collapse. No Adam state, no gradient delay.
lambda[h] ← clamp(1.0, 2.0, 1.0 + Z_SCALE_ISV * z_h)
Z_SCALE_ISV = (LAMBDA_CEILING - LAMBDA_FLOOR) / z_max_ema
Adaptive scale auto-uses the full clamp envelope:
the historical-max-z horizon maps exactly to
LAMBDA_CEILING. Replaces hardcoded Z_SCALE=0.5 which
rarely engaged on real data (max observed λ ~1.04).
Both anchor on the same ISV (loss_ema). z_max_ema is a new single-scalar
EMA state tracking max |z| across horizons, with first-obs bootstrap.
Removes:
- opt_log_sigma AdamW optimizer (σ no longer learned)
- grad_log_sigma_h_d memset (BCE kernel writes; output ignored — kept
only to preserve BCE kernel signature)
Kernel signature change (horizon_ema_and_lambda):
+z_max_ema [1] (read+write EMA state)
+log_sigma_h [5] (closed-form output, overwrite)
Discipline:
- First-obs bootstrap (sentinel <= 0) per pearl_first_observation_bootstrap
- Permanent floor (max(real, floor)) per pearl_blend_formulas_must_have_permanent_floor
- Asymmetric clamp per pearl_audit_unboundedness_for_implicit_asymmetry
- Z-score normalisation per pearl_zscore_normalization_for_magnitude_asymmetric_signals
- No nvrtc, no atomicAdd, no host branches in graph capture
All 9 perception_overfit tests pass — including
horizon_ema_and_lambda_track_after_training which validates the kernel
end-to-end through 64 K-loop iterations of capture/replay.
Submit local smoke; cluster A/B vs σ-only baseline (0.7506/0.7519) and
vs Phase 1+2+3 (0.7749/0.7591) follows once the perf-only 3-fold A/B
confirms no regression at
|
||
|
jgrusewski
|
b23f8f2efa |
perf(ml-alpha): NVIDIA-grade rewrite of CfC K-loop hot kernels — 2.15× faster
Local L40S profile (perception_overfit smoke) GPU kernel time:
1589ms → 739ms (53.5% reduction, 2.15× speedup).
Wall-clock smoke: 9.6s → 4.94s (1.94× faster).
Per-kernel deltas (nsys --cuda-graph-trace=node):
reduce_axis0: 362ms → 10ms (36× faster)
Block layout: per-column (1 block / output) → 32-wide column tile
(block_dim = 32 × 8). Cross-thread reads were strided by n_tail
(~40K floats = 160KB stride) — one cache line per thread, 8× HBM
bandwidth wasted. New tile gives coalesced 128B transactions per
warp. Block tree-reduce kept (no atomicAdd, per feedback_no_atomicadd).
+1 shared-mem pad to eliminate 32-way bank conflict on the ty reduce.
multi_horizon_heads_grn_bwd_batched: 540ms → 113ms (4.8× faster)
1. Stage h_row[HIDDEN] and a1[5,HEAD_MID] in shared at block entry.
Eliminates ~28K redundant DRAM reads/block across Pass 3 + Pass 5.
2. Pass 5 reorder: k outer / i inner with d_z1[k,m] pinned in
register; writes to grad_w1_scratch are sequential per-thread.
3. Block size 64 → 128 threads. Pass 5/6 now partition over i
(output column): cross-thread writes become COALESCED 128B/warp
(was stride-128 = 512B). Passes 2/3/4 gate on (tid < HEAD_MID).
4. Pass 3 thread role: m_out → m_in/n. Same coalescing fix on
grad_w2 writes AND w2 reads in the d_eta_2 sum.
cfc_step_backward_batched: 351ms → 271ms (1.3× faster)
1. Stage x_b[n_in] and h_old_b[n_hid] in shared (was 128× redundant
DRAM reads per block; now 1× cooperative load).
2. Pass 1 thread role: i (output row) → k (output col). For each
i loop iteration, the warp writes grad_w_in[..., tid] /
grad_w_rec[..., tid] — COALESCED 128B/warp (was stride-128
non-coalesced).
multi_horizon_heads_grn_fwd_batched: 211ms → 204ms
Stage h_row[HIDDEN] in shared — Pass 1 and Pass 3 both consume.
cfc_step_batched (fwd): 95ms → 94ms
Stage x_b and h_old_b in shared.
Shared-mem budgets fit comfortably under the 48KB SM cap (~6KB / ~2KB
respectively). All 9 perception_overfit tests pass — gradient
correctness validated end-to-end (constant-signal overfit, K-loop
capture/replay, stride-4 path, evaluate-only paths).
Discipline:
- Block tree-reduce only, never atomicAdd
- No nvrtc; pre-compiled cubins via build.rs
- Mapped-pinned-only is unaffected (CPU↔GPU contract untouched)
- Single source of truth: replaced kernels in place, no v2 suffixes
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
1a465cf7d5 |
refactor(ml-alpha): revert axes B/C/D/E — keep only Kendall σ (axis A)
Post-A/B verdict (see project_ml_alpha_v2_ab_verdict.md): v2 with all
5 axes was marginally tied on h6000 (+0.0013 vs 0.7591 baseline mean,
fails the +0.01 win threshold) and slightly below on mean_auc
(−0.0208 vs 0.7749 baseline mean, within 1σ) at ~5× the wall-time
cost. Per `feedback_v7_gem_methodology` (measure before delete or
wire), the architecture has been measured — it doesn't earn its
compute cost. This commit reverts the axes that didn't lift:
- axis B (L2 anchor + Wiener-α controller) — DROPPED
- axis C (horizon-token attention pool) — DROPPED
- axis D (regime-MoE gate + experts) — DROPPED
- axis E (inverted cross-variate attn) — DROPPED
- axis A (Kendall σ-weighted BCE) — KEPT
Files deleted (kernels, host bindings, numgrad tests, trainer state):
- cuda/{horizon_token_attention_pool, inverted_attention_pool,
inv_pooled_merge, regime_moe_gate, anchor_l2,
horizon_mean_collapse}.cu
- src/{horizon_token_attention_pool, inverted_attention_pool,
inv_pooled_merge, regime_moe_gate, anchor_l2,
horizon_mean_collapse}.rs
- src/trainer/{multi_horizon_attention, anchor_controller}.rs
- tests/{horizon_token_attention_pool_numgrad,
inverted_attention_pool_numgrad,
regime_moe_gate_numgrad,
anchor_l2_numgrad}.rs
Files restored (from V1 commit
|
||
|
jgrusewski
|
4ae9a27f48 |
fix(ml-alpha): wire axis E into loss — real add_inv_broadcast kernel
CRITICAL ARCHITECTURAL FIX discovered while planning fused kernel:
The previous Stage 2 `add_inv_broadcast` helper in
`multi_horizon_attention.rs` was a STUB that returned Ok(()) without
doing anything. Practical consequences:
- Forward: `inv_pooled_d` (output of inverted_attention_pool.forward)
was never added into `ctx_h_d`. Downstream MoE + heads never saw
the inverted-attention signal. Axis E contributed ZERO to the
forward output and the loss.
- Backward: `inv_pool.backward` was being fed `grad_ctx_mean` as
its "upstream gradient", but that's the gradient at the CHAIN
TERMINUS — not the gradient w.r.t. inv_pooled_d (which is zero
by construction since inv_pooled wasn't in the loss). The bwd
was injecting incorrect noise into `grad_ln_out`.
Net: paying inverted_attention compute for no gain, plus polluting
ln_b's gradient. Two perf rewrite rounds earlier today showed no
wall-time movement precisely because the slow path was wired into
training while the optimized one was dead.
FIX:
(a) New kernel `cuda/inv_pooled_merge.cu`:
fwd: ctx_h[b, h, d] += inv_pooled[b, d] (broadcast over h)
bwd: grad_inv_pooled[b, d] = Σ_h grad_ctx_h[b, h, d]
Tiny — single block-per-batch, no syncthreads, coalesced reads.
(b) Host binding `src/inv_pooled_merge.rs` (InvPooledMerge).
(c) `MultiHorizonAttention` adds:
- `merge: InvPooledMerge` field.
- `grad_inv_pooled_d [B, H]` buffer for the real upstream of inv_pool.bwd.
(d) `MultiHorizonAttention.forward` now calls `merge.forward(...)`
between inv_pool.forward and moe.forward. Axis E is now actually
in the model's forward output.
(e) `MultiHorizonAttention.backward` now calls `merge.backward` after
moe.bwd writes `grad_ctx_h_d`, producing `grad_inv_pooled_d`.
`inv_pool.backward` consumes the REAL upstream gradient
(`grad_inv_pooled_d`) instead of the prior `grad_ctx_mean` fake.
(f) Old stub `add_inv_broadcast` deleted.
CORRECTNESS:
- perception_overfit 9/9 PASS after the wiring. Loss still shrinks
on the constant-signal test (0.67 → -0.99 over 250 steps), now
with axis E actually contributing.
- All numgrad kernel tests still pass (kernels themselves unchanged
in this commit; only the wiring).
PERF IMPACT (expected):
- +2 tiny launches per step (merge fwd + bwd). Negligible.
- The inverted-attention compute that was previously dead now
actually feeds the loss → same wall-time, but it's earning the
cost. This unblocks meaningful axis-E perf measurement on next
smoke.
NEXT: re-run cluster smoke to confirm wall-time + verify axis E
gradient flow is healthy. After that, consider full MHA forward
fusion as a follow-up.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
11b964359b |
perf(ml-alpha): MoE bwd compact scratch + scatter kernel + inv-attn loop interchange
Two perf optimizations bundled:
(1) MoE backward scratch compaction
Old: `grad_w_scratch_d [B, N_H, N_E, H, H]` = 1.3 MB per step (B=1)
New: `grad_w_scratch_d [B, N_H, H, H]` = 320 KB per step
4× memory reduction. Since each batch's top-1 router selects ONE
expert, only that expert's slot was ever non-zero in the prior
layout — the N_E axis was entirely wasteful.
The shape change required:
- Updated `regime_moe_gate_bwd` to write the compact layout.
- New `regime_moe_gate_scatter` kernel scatters per-(b, h)
rank-1 contributions into `grad_experts_W[e]` / `grad_experts_b[e]`
based on `top_e[b]`. Grid (N_E, H, ceil(H/32)) × block (32) —
one warp per (e, d_out, d_in_chunk). 65536 → 16384 grid cells
(4× fewer blocks dispatched).
- Dropped the previously-naive 65536-block `reduce_axis0` for
`grad_experts_w` from `perception.rs` (the scatter kernel
produces the final per-expert grad directly).
- `tests/regime_moe_gate_numgrad.rs` reads `grad_experts_w` from
the scatter output instead of host-side reducing the 5D scratch.
(2) inverted_attention bwd loop interchange
Phase 2's tight loop:
for k:
for j:
ds_myh_j = d_scores[my_h, j] // doesn't depend on k!
ds_j_myh = d_scores[j, my_h] // doesn't depend on k!
...
Hoisted d_scores reads out of the K-loop into J-outer with
per-thread `q_arr[K_MAX]` / `k_arr[K_MAX]` register accumulators.
Net: 32× fewer DRAM reads of d_scores per thread per bwd.
CORRECTNESS:
- regime_moe_gate numgrad PASSES (1/1, 11 numgrad checks).
- inverted_attention numgrad PASSES (1/1, 6 numgrad checks).
- perception_overfit 9/9 PASS — including loss-shrinks tests.
NEXT: re-run cluster smoke to measure the new wall-time vs the 17 s
baseline. Prior smoke at
|
||
|
jgrusewski
|
a263cd5446 |
perf(ml-alpha): inverted_attention_pool — cache mean_k, pre-compute pooled
The cluster smoke at
|
||
|
jgrusewski
|
9170d24fe3 |
refactor(ml-alpha): replace legacy attention_pool with MultiHorizonAttention [Stage 2]
Single source of truth for the attention path. Deletes the legacy
single-Q `attention_pool.cu` and all `attn_*` fields from
`PerceptionTrainer`; wires `MultiHorizonAttention` (the bundle
introduced in Stage 1) into `step_batched` + `evaluate_batched` as
THE attention summary that seeds CfC's `h_old` at k=0.
Deletions:
cuda/attention_pool.cu (244 lines)
perception.rs::attn_q_d/attn_context_d/
attn_weights_d/grad_attn_q_d/opt_attn_q/
attn_fwd_fn/attn_bwd_fn/_attn_module/
attn_grad_q_scratch_d (all struct fields)
perception.rs::ATTENTION_POOL_CUBIN (include_bytes constant)
Their corresponding init + struct-construction lines.
build.rs::KERNELS (drops "attention_pool")
New kernel + binding:
cuda/horizon_mean_collapse.cu (53 lines)
- `horizon_mean_collapse_fwd/_bwd`: collapses [B, N_H, H] → [B, H]
by averaging over the horizon axis. Single-pass, no reductions.
src/horizon_mean_collapse.rs (host binding)
MHA additions:
- `collapse` field + `ctx_mean_d` + `grad_ctx_mean_d` for the seed.
- `grad_ctx_h_d` scratch (split from grad_horizon_tokens_scratch to
avoid aliasing when MoE bwd writes d_ctx_h while pool bwd writes
d_horizon_tokens).
- `forward(ln_b_out)`: horizon-token pool → inverted pool → MoE
dispatch → mean-collapse → ctx_mean_d.
- `backward(ln_b_out, grad_ctx_mean, grad_ln_out)`: full reverse
chain.
- `apply_anchor()`: launches anchor_l2 on horizon_tokens, Q,
experts_w.
- `adamw_step()`: steps all 6 owned optimizer groups.
PerceptionTrainer integration:
- Section 2d (forward): `self.mha.forward(&self.ln_out_d)` replaces
the legacy attention_pool launch. CfC's h_old at k=0 now reads
`self.mha.ctx_mean_d.device_ptr` (was `self.attn_context_d`).
- Section 7c-pre (backward): `self.mha.backward(ln_out, grad_h_carry,
grad_h_enriched_seq)` replaces the legacy attn_bwd_fn launch.
- Four `reduce_axis0` launches collapse MHA's per-batch scratches
into shared gradient buffers: grad_horizon_tokens, grad_q,
grad_experts_w, grad_experts_b.
- `self.mha.apply_anchor()` adds L2 anchor grad contributions.
- Section 9 (AdamW): `self.mha.adamw_step()` replaces opt_attn_q.
- `evaluate_batched`: `self.mha.forward` replaces the legacy fwd
launch; h_old at k=0 reads `mha.ctx_mean_d`.
- `self.mha.zero_grads()` at step start (capture-safe memset_zeros).
BUG CAUGHT DURING WIRING (NVIDIA-grade discipline): first wiring
attempt mis-sized the reduce_axis0 launches for the MoE
`grad_w_scratch_d` ([B, N_H, N_E, H, H]). Initial `n_tail = N_H * N_E
* H * H = 327680` would have made reduce_axis0 read 5× past the end
of the buffer → CUDA_ERROR_ILLEGAL_ADDRESS. Fix: `n_tail = N_E * H *
H = 65536` with `n_batch = B * N_H`, treating the leading two axes
together as the reduction dimension. Caught by stacked_trainer test
on RTX 3050; would have caused silent corruption then a hard fault
on L40S/H100 later.
LOCAL VERIFICATION (RTX 3050 sm_86):
- ml-alpha builds clean (cuda feature).
- All 38+ tests PASS serially with --test-threads=1:
perception_overfit (8 tests incl. loss-shrinks)
trunk_forward (5)
stacked_loss_shrinks (multiple)
bce_grad_finite_diff (4)
snap_feature_assemble (9)
... (full suite green)
- Numgrad parity for the 4 new MHA kernels (horizon_token, inv_attn,
regime_moe_gate, anchor_l2) PASSES at 5e-2 rel.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
6a3f45d872 |
refactor(ml-alpha): consolidate BCE — Kendall σ is THE BCE, wire MHA into trainer
Single source of truth for the multi-horizon BCE per the new
feedback_single_source_of_truth_no_duplicates pearl. Eliminates the
`bce_loss_multi_horizon_sigma.cu` / `loss_sigma.rs` duplication
introduced earlier today and folds the Kendall σ-weighting into the
canonical kernel.
Deletions:
cuda/bce_loss_multi_horizon_sigma.cu (folded into legacy)
src/trainer/loss_sigma.rs (helper merged)
tests/bce_sigma_numgrad.rs (subsumed)
Rewrites:
cuda/bce_loss_multi_horizon.cu — replaced with σ-aware
implementation; kernel function name kept as
`bce_multi_horizon_forward_backward` so cubin symbol stays stable.
NVIDIA-grade warp-shuffle reduce (4 warps, 1 cross-warp barrier);
new args `log_sigma_h` (per-horizon Kendall σ logarithm) and
`d_log_sigma_h` (its gradient).
src/trainer/loss.rs — standalone helper updated to new 11-arg
kernel signature. Exposes optional `log_sigma_h` in `BceInput`
(None → zeros / passthrough Kendall init) and returns
`mean_bce_per_h` + `d_log_sigma_h` in `BceOutput`.
tests/bce_grad_finite_diff.rs — adds `log_sigma_h: None` to the
test inputs; all 4 numgrad tests PASS unchanged.
build.rs — drops `bce_loss_multi_horizon_sigma` entry from
KERNELS. The single canonical `bce_loss_multi_horizon` cubin
now contains the σ-aware kernel.
Wiring:
PerceptionTrainer gains a single `pub mha: MultiHorizonAttention`
field. Owns `log_sigma_h_d` + `grad_log_sigma_h_d` (and the rest
of the multi-horizon attention path, anchored on Stage 2 to fully
replace the legacy `attention_pool` path).
step_batched + evaluate_batched BCE callsites now thread
`mha.log_sigma_h_d` and `mha.grad_log_sigma_h_d` through the
11-arg kernel signature.
This commit keeps the legacy `attention_pool` callsite in place; the
Stage 2 commit will replace it with `mha.pool` + `mha.inv_pool` +
`mha.moe` and delete the `attn_*` fields entirely.
Verified locally: ml-alpha builds clean with the cuda feature,
bce_grad_finite_diff (4/4) PASS on RTX 3050.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
996e61b6ef |
rename(ml-alpha): MultiHorizonAttention (no version suffix in identifiers)
Removes the version suffix from the trainer-state bundle introduced
in
|
||
|
jgrusewski
|
5aad1eb846 |
feat(ml-alpha): PerceptionV2State bundle (axes A+B+C+D+E) [V9]
Bundles all v2 components into one cohesive trainer field that
PerceptionTrainer wires through in V10. Owns:
(C) horizon_tokens + Q for the horizon-token attention pool
+ per-batch grad scratches + 2 AdamWs
(E) inverted attention pool + saved attn + d_scores scratch
(D) MoE gate weights + experts (W, b) + per-batch sparse-by-expert
grad scratches + 3 AdamWs + top_e / gate_probs / aux_loss
(A) log_sigma_h per-horizon Kendall scalar + AdamW (0.25× LR)
(B) AnchorController (Wiener-α host-side) + anchor_l2 kernel +
init snapshots of horizon_tokens, Q, experts_w, experts_b
for the L2 anchor
Init scale: 1/√HIDDEN_DIM Xavier for horizon_tokens/Q/experts_W;
zeros for experts_b and log_sigma_h. Anchor controller bootstrap
uses ‖horizon_tokens_init‖₂ + ‖Q_init‖₂ + their total numel to
derive the signal-driven floor (no tuned constants).
zero_grads uses memset_zeros only — capture-safe.
Compile-clean. V10 wires this state into perception.rs::step_batched
forward / backward / AdamW.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
55aeddaebd |
feat(ml-alpha): anchor_l2 kernel + Wiener-α controller (v2 B) [V8]
L2 anchor regularization toward initialization (axis B). Anchors
horizon_tokens + Q + MoE experts toward their init values to prevent
the calibration drift observed in v1 (where val_loss climbed as α
opened past epoch 1 in 2 of 3 folds).
KERNEL (`anchor_l2_fwd_bwd`):
loss_out = λ · Σ_i (p[i] − p_init[i])²
grad_p[i] += 2λ · (p[i] − p_init[i])
- Warp-shuffle reduce; one block per parameter group; strided thread
loop over n. Cross-warp reduce uses one __syncthreads.
- Coalesced grad write via stride loop.
- λ passed as device-side [1]-buffer (host writes scalar before launch
— capture-safe).
CONTROLLER (`trainer::anchor_controller::AnchorController`):
- Signal-driven λ floor: λ_floor = ‖p_init‖ / (100 · √numel).
Cross-fold-persistent per pearl_kelly_cap_signal_driven_floors.
- Wiener-α smoother (α = diff_var / (diff_var + sample_var + ε))
on val_loss change; α floored at 0.4 per
pearl_wiener_alpha_floor_for_nonstationary.
- λ blends toward target = |ema_change|·scale with α; floored at
λ_floor per pearl_blend_formulas_must_have_permanent_floor.
- First-observation bootstrap (sentinel state replaced directly on
first step) per pearl_first_observation_bootstrap.
- 4/4 unit tests PASS: signal-floor init, bootstrap returns floor,
floor protection across 1000 steps, λ_max cap.
NUMGRAD VERIFICATION (RTX 3050 sm_86):
anchor_l2_numgrad PASSES with closed-form parity (machine precision)
and central-difference parity (4 random positions) within 5e-2 rel.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
11b96dac6a |
feat(ml-alpha): regime_moe_gate kernel + numgrad (v2 D) [V6]
Top-1 Mixture-of-Experts gate per Switch Transformer. Three kernels
in one .cu file:
- regime_moe_gate_fwd: select top-1 expert from gate_logits, apply
its [H, H] linear to fused_ctx → routed_ctx [B, N_H, H].
- regime_moe_gate_bwd: chain-rule grads through the SELECTED
expert's W and bias (sparse-by-expert scratches), plus
d_fused_ctx accumulator. Inactive experts get zero contribution.
- regime_moe_gate_aux: softmax of gate_logits + load-balancing
auxiliary loss (frac · prob_mean × N_EXPERTS).
ARCHITECTURE:
- N_EXPERTS = 4. Each expert is a [H, H] linear with bias.
- Total expert params: 4 · 128 · 128 + 4 · 128 = 66 KB. Cheap.
- STE on gate: gate logit grad is zero from the expert path (top-1
is non-differentiable); the load-balance aux loss provides the
differentiable signal that pushes routing toward balanced usage.
PERFORMANCE:
- Grid (B, N_H, 1), block (HIDDEN_DIM). One block per (b, h).
- Forward: each thread computes one output channel via a dot
product over HIDDEN_DIM input dims (#pragma unroll 8).
- Backward d_fused_ctx: each thread accumulates over d_out
sequentially (HIDDEN_DIM iterations) since the weight matrix
column is naturally aligned to the thread's d_in index.
- Backward d_W/d_b scratches are sparse-by-expert; downstream
reduce_axis0 collapses over (B, N_H).
- Top-1 chosen by thread 0 per block, broadcast via shared mem.
NUMGRAD VERIFICATION (RTX 3050 sm_86):
forward_matches_host_reference_and_backward_matches_numgrad
PASSES 11 checks (4 on d_W, 3 on d_b, 4 on d_fused_ctx) within
5e-2 rel / 5e-3 abs.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
d94696620d |
feat(ml-alpha): inverted_attention_pool kernel + numgrad (v2 E) [V3]
iTransformer-style cross-variate attention pass: each of HIDDEN_DIM
features becomes a "variate token" with its K-trajectory as embedding.
FORWARD:
X_inv[h, k] = ln_out[k, h] # transpose
scores[h, j] = inv_scale · Σ_k X_inv[h, k] · X_inv[j, k]
attn[h, j] = softmax_j(scores)
pooled[h] = Σ_j attn[h, j] · mean_k_X_inv[j] # mean-K commutes out
BACKWARD: three independent chains into ln_out:
- value path: (1/K) · attn[h', my_h] · d_pooled[h'] (per-k constant)
- query path: inv_scale · Σ_j d_scores[my_h, j] · X_inv[j, k]
- key path: inv_scale · Σ_h' d_scores[h', my_h] · X_inv[h', k]
Softmax bwd: d_scores[h, j] = attn · (d_attn - Σ_l attn · d_attn)
IMPLEMENTATION NOTES:
- First attempt cached attn [H, H] = 64 KB in shared mem → tripped
the 48 KB dynamic-shared limit on sm_86 (CUDA_ERROR_INVALID_VALUE).
- Fixed by moving d_scores to a DRAM scratch buffer; shared mem
holds only X_inv [H, K] (≤ 16 KB at K = 32). One block-wide barrier
between the d_scores write and the value/query/key accumulation.
- All per-batch slice writes; no atomicAdd, no cross-block race.
- Pooled computation uses the mean-K commute (Σ_k attn · X_inv =
attn · mean_k_X_inv), saving an entire H×K accumulation pass.
LOCAL VERIFICATION (RTX 3050 sm_86):
forward_then_backward_matches_central_difference PASSES 6 numgrad
checks on ln_out positions within 5e-2 rel / 5e-3 abs.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
e8f6c4721f |
feat(ml-alpha): horizon_token_attention_pool kernel + numgrad (v2 C) [V2]
Replaces the falsified per-horizon Q_h pool with a single shared
query Q over an extended key sequence [horizon_tokens; LN_b_out],
producing per-horizon outputs via TFT-style horizon-token mixing.
FORWARD:
scores[i] = Σ_d Q[d] · ext[i, d] (i ∈ [0, N_H + K))
attn = softmax_i(scores)
S[d] = Σ_k attn[N_H + k] · ln_out[k, d] (shared time agg)
ctx[h, d] = attn[h] · horizon_tokens[h, d] + S[d] (per-horizon out)
BACKWARD: full chain rule with softmax-centring; gradients to
horizon_tokens, Q, and ln_out via the saved attn weights.
NVIDIA-grade implementation per feedback_nvidia_grade_perf_for_kernels:
- Warp-shuffle reduce (block_reduce_sum / block_reduce_max helpers)
for all per-d dot products and softmax aggregates.
- Cross-warp reduce uses exactly one __syncthreads.
- Non-divergent shuffles: inactive lanes contribute 0 via ternary.
- Block-per-batch + horizon-loop inside block → grad_ln_out += is
race-free without atomicAdd.
- Smem layout computed at launch: [s_attn(N_H+K); s_warp(N_WARPS);
s_d_S(H) on bwd]. No over-allocation.
LOCAL VERIFICATION (RTX 3050 sm_86):
forward_then_backward_matches_central_difference PASSES 12 numgrad
checks (4 each on horizon_tokens / Q / ln_out) at 5e-2 rel / 5e-3
abs envelope. First-try pass.
NOTE: .gitignore adjusted with narrow allow-rules for crates/ml-alpha/{
cuda,src,tests}/horizon_token_* paths — the broad "*token*" rule
intended for auth tokens was hiding these source files. Explicit
allow keeps the security rule intact while exempting these specific
files.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
679ab3f5eb |
feat(ml-alpha): Kendall sigma-weighted BCE kernel + numgrad (v2 A) [V7]
New `bce_multi_horizon_sigma_forward_backward` kernel implementing the
Kendall homoscedastic uncertainty weighting per spec axis A:
raw_bce_h = Σ_{i in h, m_i=1} L_i
count_h = #{i in h : m_i = 1}
mean_bce_h = raw_bce_h / count_h
w_h = base_weight_h / (2 · exp(2 · log_sigma_h))
total_loss = Σ_h [ w_h · mean_bce_h + log_sigma_h ]
d L / d p_i = m_i · (w_h / count_h) · (p − y) / (p (1 − p))
d L / d log_sigma_h = 1 − 2 · w_h · mean_bce_h
NVIDIA-grade implementation per feedback_nvidia_grade_perf_for_kernels:
- Warp-shuffle reduction (`__shfl_xor_sync`) for both per-horizon
sums and the global valid count, replacing block tree-reduce.
- One `__syncthreads` for the cross-warp aggregate; no inner-loop
barriers.
- Non-divergent shuffles: inactive lanes contribute 0 via ternary,
never via `if (tid < N) shuffle`.
- Coalesced strided access in both forward and gradient passes.
- Pre-compiled cubin via build.rs; no nvrtc.
Independent of the legacy `bce_loss_multi_horizon` kernel — that one
stays untouched so eval/smoke paths are unaffected. The v2 trainer
wires this kernel in via commit V10.
Standalone helper `bce_sigma_loss_and_grad_gpu` in `trainer::loss_sigma`
for numgrad parity tests. Three numgrad tests all PASS on RTX 3050
(sm_86) within 5e-2 rel / 5e-3 abs:
- d_log_sigma_h ↔ central-difference (numgrad on log_sigma)
- grad_probs ↔ central-difference (8 random positions)
- total_loss ↔ closed-form reconstructed from mean_bce_per_h
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
41292303dc |
refactor(ml-alpha): remove per-horizon Q_h path (C21-C25 falsified) [V1]
3-fold A/B sweep 2026-05-18 at commit
|
||
|
jgrusewski
|
286ea26e2a |
perf(ml-alpha): warp-shuffle reduce in per-horizon kernels
Cluster A/B sweep with C25 wiring showed 86 s/epoch vs 17 s/epoch baseline = ~5x regression. Root cause: per-horizon attention pool + residual head used block tree-reduce with 8 __syncthreads per K-step in a serialised K-loop, repeated H=5 times in both fwd and bwd = ~3200 barriers/step. Plus the prob_blend bwd reduce kernel ran with a single thread per block, fully serialising over K*B. Replacements: - per_horizon_attention_pool fwd/bwd: introduce block_reduce_sum / block_reduce_max helpers using intra-warp __shfl_xor_sync + cross-warp shuffle (1 syncthread per K instead of 8). Smem shrinks to [K + N_WARPS] / [2K + N_WARPS]. - per_horizon_residual_head fwd: same warp-shuffle reduce pattern. - per_horizon_prob_blend_reduce_alpha_residual: 1 thread → 1 warp per horizon, lane-strided reduction over K*B via shfl_xor_sync. Launch config updated to block_dim=(32,1,1). Tricky bug found while implementing: the cross-warp reduce in the residual head originally guarded `__shfl_xor_sync(0xffffffff, ...)` with `if (tid < PHR_N_WARPS)`, leaving 28 of 32 lanes in warp 0 outside the call. Mask 0xffffffff requires all 32 lanes to participate — divergence is UB and hung the full-pipeline smoke on Ampere/Ada. Fix: read s_warp via ternary into all 32 lanes, then shuffle inside `if (tid < 32)`. Matches the pattern used in block_reduce_sum. Verified locally on RTX 3050 (sm_86): per_horizon_attention_pool numgrad, per_horizon_residual_head numgrad, and per_horizon_full_pipeline_smoke (zero-init identity + non-zero end-to-end) all PASS. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
83546b5c37 |
fix(ml-alpha): drop synchronize() in per-horizon pool + head bindings
After
|
||
|
jgrusewski
|
adaf275af3 |
fix(ml-alpha): C25 per-horizon path graph-capture safety
Four code paths in PerHorizonTrainState violated CUDA Graph capture invariants and tripped CUDA_ERROR_STREAM_CAPTURE_UNSUPPORTED on smoke alpha-perception-qk2p9: 1. stream.synchronize() inside forward_with_blend (illegal in capture) 2. stream.synchronize() inside backward_through_blend (idem) 3. reduce_per_batch_scratches_to_shared used host-side vec allocations + memcpy_dtoh + CPU summation + memcpy_htod (forbidden during capture per pearl_no_host_branches_in_captured_graph) 4. zero_grads allocated host zero vectors + memcpy_htod each step (idem) Fix: - Remove both synchronizes; same-stream issue order is sufficient. - Replace host-side reduction with three reduce_axis0 GPU kernel launches (q_h, w_res, bias_res). PerHorizonTrainState now owns its own reduce_axis0 cubin handle. - Replace host-zero memcpy with stream.memset_zeros for all nine gradient buffers plus d_alpha_reduced. Verified locally: per_horizon_full_pipeline_smoke (2/2) and both numgrad parity tests (attention pool + residual head) pass on RTX 3050. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
f50b466f77 |
feat(ml-alpha): PerceptionTrainer wires C21+C22 fwd+bwd+AdamW (C25)
The per-horizon attention pool from C21+C22 is now fully integrated
into step_batched's hot loop. Forward and backward both flow; AdamW
updates all four new parameter groups (Q_h, w_res, bias_res, α) every
training step. At α=0 init the contribution is byte-identical to
baseline; training discovers whether α should grow.
New kernel: cuda/per_horizon_prob_blend.cu
per_horizon_prob_blend_fwd
Reads logit_per_k_d (already stored by GRN forward) +
sigmoid(logit_baseline + tanh(α[h]) * residual[b, h]) →
overwrites probs_per_k_d in place. At α=0: r_contrib=0, output
== sigmoid(logit_baseline) == probs_baseline → bit-identical.
per_horizon_prob_blend_reduce_alpha_residual
Reads probs_per_k (= p_final, post-blend) + grad_probs_per_k
(= ∂L/∂p_final from BCE) and computes:
d_logit[k,b,h] = grad_probs[k,b,h] * p_final * (1 - p_final)
d_residual[b,h] = tanh(α[h]) * Σ_k d_logit[k,b,h]
d_alpha[h] = sech²(α[h]) * Σ_{k,b} d_logit[k,b,h] * residual[b,h]
No separate prob_blend_bwd needed — chain-rule equivalence
∂L/∂logit_baseline = ∂L/∂r_contrib (both flow through the same
sigmoid derivative) means the existing GRN backward is UNCHANGED.
trainer/per_horizon_state.rs extensions:
forward_with_blend(ln_b_out, logit_per_k, probs_per_k)
Pool fwd → context_h; head fwd → residual; prob_blend fwd
in-place rewrites probs_per_k.
backward_through_blend(probs_per_k, grad_probs_per_k, ln_b_out,
grad_ln_b_out_target)
Reduce kernel → d_residual + d_alpha. Then:
head bwd → d_w_res_scratch, d_bias_res_scratch, d_context.
pool bwd → d_q_h_scratch, += grad_h_enriched_seq_d.
Per-batch scratches reduced to shared grads host-side
(n_batch ≤ 64 → sub-millisecond on host).
adamw_step()
Steps the four optimizers using the shared grad buffers.
zero_grads()
Called once per step before forward to clear scratch.
trainer/perception.rs step_batched integration:
── 4.5 (after GRN K-loop, before BCE): zero_grads + forward_with_blend
overwrites probs_per_k_d with p_final.
── 5 (existing BCE consumes probs_per_k_d as today; grad_probs is
now ∂L/∂p_final automatically).
── 5a (after BCE, before ISV-lambda + heads bwd): backward_through_blend.
Existing GRN bwd path is UNTOUCHED — the chain rule absorbs
the bias.
── 9 (after existing 17 AdamW group steps): per_horizon.adamw_step
updates Q_h, w_res, bias_res, α.
Verification:
- 34 ml-alpha lib tests still green.
- Per-horizon kernel numgrad parity (C21, C22) still green.
- Per-horizon end-to-end pipeline smoke (C23, including the
alpha=0 byte-identity invariant) still green.
- Full workspace builds clean.
Closes the kernels+wiring portion of #203 (per-horizon attention pool
kernels + wiring). What remains (#204): 30-epoch × 3-fold A/B vs
single-Q baseline. The branch is ready for that sweep when GPU time
is budgeted; the implementation is structurally adoption-safe
(α=0 → identity to baseline) so it can be merged before the A/B if
desired.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
5eb6567c4c |
feat(ml-alpha): PerceptionTrainer per-horizon trainer state (C24)
Allocates the device buffers + AdamW optimizers + kernel bindings for
the per-horizon attention-pool training path from C21 + C22, bundled
into a single PerceptionTrainer.per_horizon field.
crates/ml-alpha/src/trainer/per_horizon_state.rs (NEW):
PerHorizonTrainState — owns:
Learnable params (zero-init for α + bias, Xavier-scale for Q_h,
0.1× scale for w_res):
q_h_d [N_HORIZONS, HIDDEN_DIM] — attention queries
w_res_d [N_HORIZONS, HIDDEN_DIM] — residual head weights
bias_res_d [N_HORIZONS] — residual bias
alpha_d [N_HORIZONS] — learnable gate (init 0)
Forward intermediates (per-batch):
context_d [B, N_HORIZONS, HIDDEN_DIM]
attn_weights_d [B, N_HORIZONS, K]
residual_d [B, N_HORIZONS]
Backward grad scratch + reduced grads + AdamW state.
Kernel bindings: PerHorizonAttentionPool + PerHorizonResidualHead.
PerHorizonTrainState::new(dev, n_batch, k_seq, lr, seed)
Allocates all buffers, runs Xavier-style init, constructs four
AdamW optimizers (q_h, w_res, bias_res at param-LR; α at 0.25× LR
per spec §5 open Q2 default — slow gate ramp). Captures the seed
via wrapping_add(0xA110C00A) from cfg.seed for determinism.
PerHorizonTrainState::zero_grads()
Clears all grad-scratch buffers between training steps.
trainer/mod.rs:
pub mod per_horizon_state — module export.
trainer/perception.rs:
PerceptionTrainer gains one field:
pub per_horizon: PerHorizonTrainState
Initialised in new() with cfg.n_batch + cfg.seq_len + cfg.lr_cfc.
α-gate init=0 ⇒ tanh(0)=0 ⇒ contribution to per-batch logits is exactly
0 at step 0 (per C23 alpha_zero_init_is_identity_to_baseline byte-equality
test). Adopting this commit produces bit-identical training behaviour
to the previous commit until C25 wires the forward+backward calls into
step_batched; even then the α=0 init means a one-epoch smoke against
existing baseline should match within FP rounding noise.
All 34 ml-alpha lib tests + 4 per-horizon GPU tests (numgrad pair +
end-to-end pipeline pair) green.
Next:
C25 — forward + backward integration into step_batched. The new path
runs as ADDITIONAL kernel launches before/after the existing
captured graph (not inside it) so the graph stays unchanged
and the integration risk is contained.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
69c64f2266 |
test(ml-alpha): per-horizon pipeline end-to-end smoke + α-gate (C23)
Composes the C21 + C22 kernels with a host-side learnable α-gate to
prove the full per-horizon contribution path works end-to-end without
yet doing the captured-graph integration in PerceptionTrainer.
Pipeline:
LNb [B, K, HIDDEN_DIM]
→ per_horizon_attention_pool_fwd → context_h [B, N_HORIZONS, HIDDEN_DIM]
→ per_horizon_residual_head_fwd → residual [B, N_HORIZONS]
→ final[b, h] = baseline[b, h] + tanh(α[h]) * residual[b, h]
Two tests cover the critical invariants for adoption-safety:
alpha_zero_init_is_identity_to_baseline
With α = [0, 0, 0, 0, 0] and any random Q_h / w_res / bias_res,
final_logit MUST be bit-identical to baseline_logit (because
tanh(0) = 0). Verified via to_bits() byte equality. Proves that
initialising the new variant with α=0 makes it a strict superset
of the existing path — switching to AttentionPoolVariant::PerHorizon
cannot regress before any training has happened.
alpha_nonzero_changes_output_and_grads_flow_end_to_end
With α = [0.5, -0.3, 0.2, -0.1, 0.4]:
* final ≠ baseline (residual contributing) ✓
* all final logits finite ✓
* full backward chain (residual_head_bwd → attention_pool_bwd)
produces finite d_Q_h_scratch + finite d_LNb with at least
one non-zero entry in each → gradients flow back to both the
attention queries and the LN_b input ✓
This closes the kernel-side correctness story. The remaining
integration commits (C24+) are operational:
C24: extend CheckpointV1 → V2 (add q_h, w_res, bias_res, alpha
fields; V1 files load as Variant::SharedQuery)
C25: PerceptionTrainer wiring — allocate the device buffers, fold
attention + residual + gate into the captured graph, plumb
gradients into AdamW's param list
C26: 1-epoch smoke (assert no NaN, loss decreases vs baseline) —
needs real training data + multi-GPU time
C27: 30-epoch × 3-fold A/B (task #204) — decision gate per
docs/superpowers/specs/2026-05-18-per-horizon-attention-pool-design.md
§0 falsifiable claim
C24-C25 are 1-2 day work even when carefully scoped; C26-C27 need
real GPU-hours + result analysis. C21-C23 land the validatable kernel
correctness piece without committing to that time investment yet.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
8c335caef7 |
feat(ml-alpha): per-horizon residual head kernel + numgrad parity (C22)
Companion kernel to C21's per_horizon_attention_pool. Computes a
per-horizon scalar residual from each horizon's context vector:
residual[b, h] = Σ_d w_res[h, d] * context_h[b, h, d] + bias_res[h]
Designed to be added (behind a learnable α-gate) to the existing
multi_horizon_heads logit output — keeps the existing GRN head kernel
completely unchanged. The per-horizon attention pool's contribution
flows through this lightweight projection without weight-shape
changes elsewhere or checkpoint-V2-bumping.
Path A integration sketch (deferred to follow-up commit C23):
alpha_logit_per_horizon = existing_head(h_K)[h] # from current path
+ tanh(α[h]) * residual_kernel(context_h)[h]
where α[h] is a learnable 5-vector init'd to 0 (no effect at start).
Training discovers per-horizon whether the residual contributes.
This is a strict superset of the existing path — α=0 → bit-identical
to today.
Backward kernel produces:
d_w_res — per-block scratch [B, N_HORIZONS, HIDDEN_DIM]
for host reduce_axis0 → shared [N_HORIZONS, HIDDEN_DIM]
d_bias_res — per-block scratch [B, N_HORIZONS], same reduction
d_context_h — per-batch indexed; += chained with attention bwd
Single-writer discipline preserved (no atomicAdd per
feedback_no_atomicadd.md); horizon loop inside the per-batch block.
Numgrad parity test:
- B=3, N_HORIZONS=5, HIDDEN_DIM=128 fixture.
- Loss = Σ residual_out (so d_residual = 1).
- Probes 8 random w_res indices, all 5 bias_res entries, 8 random
context_h indices via central-difference at ±eps=1e-2.
- All within 5e-2 rel-tol or 5e-3 abs-floor.
- Passes on RTX 3050.
Same scope discipline as C21: kernel + binding + numgrad first;
trainer wiring + α-gate + smoke training + A/B sweep follow once
both kernels are individually validated (now done).
Closes the second kernel-correctness portion of #203. Remaining:
C23: trainer wiring (capture attn_pool fwd into the graph; sum
residual into existing head output with α-gate)
C24: CheckpointV1 → V2 bump (add q_h, w_res, bias_res, alpha fields)
C25: 1-epoch smoke (assert no NaN, loss decreases vs baseline)
C26: 30-epoch × 3-fold A/B (#204) — decision gate per spec §0
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
edc449eecf |
feat(ml-alpha): per-horizon attention pool kernel + numgrad parity (C21)
First implementation slice of the per-horizon attention pool design
(docs/superpowers/specs/2026-05-18-per-horizon-attention-pool-design.md).
Lands the kernel + Rust binding + numgrad verification; downstream
wiring into PerceptionTrainer's captured graph + CheckpointV2 bump +
A/B sweep are follow-up commits gated on this proving correctness.
Kernel (cuda/per_horizon_attention_pool.cu):
per_horizon_attention_pool_fwd
Q_h[N_HORIZONS, HIDDEN_DIM] × LNb[B, K, HIDDEN_DIM]
→ context_h[B, N_HORIZONS, HIDDEN_DIM]
attn_h_weights[B, N_HORIZONS, K]
Per-block math identical to the single-Q variant, looped over
N_HORIZONS sequentially within each batch's block. Grid stays
(B, 1, 1) so backward grad_ln_out writes are race-free
(per feedback_no_atomicadd.md — no cross-block contention).
Per-batch shared mem ~k_seq + BLOCK + HIDDEN_DIM floats.
per_horizon_attention_pool_bwd
Same chain-rule pattern as attention_pool_bwd but with the horizon
loop inside the block: each (b, h) slice updates grad_ln_out in
place (sequential horizon accumulation), grad_Q_h is written as
per-block scratch [B, N_HORIZONS, HIDDEN_DIM] for host reduce.
Single-writer discipline preserved.
Rust binding (src/per_horizon_attention_pool.rs):
PerHorizonAttentionPool::{new, forward, backward}. Self-contained;
doesn't yet touch PerceptionTrainer or CfcTrunk. Loads the cubin
via the standard env!("OUT_DIR") path. Dynamic shared-mem byte
count computed per launch from k_seq.
Numgrad parity test (tests/per_horizon_attention_pool_numgrad.rs):
- B=2, K=8, HIDDEN_DIM=128, N_HORIZONS=5 fixture.
- Loss = Σ context_h (so d_context = 1 everywhere — clean analytical).
- Backward kernel produces analytical grads; central-difference of
forward kernel at ±eps=1e-2 across 8 random Q_h indices + 8 random
LNb indices verifies analytical matches CD within 5e-2 rel-tol or
5e-3 abs-floor.
- Passes on RTX 3050.
build.rs picks up the new .cu file automatically via the existing
KERNELS list; cubin compiles cleanly at sm_86 + sm_89.
Same scope discipline as Phase 2D.2 (VSN numgrad) — kernel correctness
first, integration second. The follow-up commit set per the spec §3
appendix:
C22: extend multi_horizon_heads.cu signature to accept per-horizon
context input + bump head_w shape to [N_HORIZONS, 2*HIDDEN_DIM]
C23: wire PerHorizonAttentionPool into PerceptionTrainer + CfcTrunk
captured graph behind AttentionPoolVariant config flag
C24: CheckpointV1 → V2 bump with discriminant + optional q_h field
C25: smoke training (one epoch, no NaN, loss decreases)
C26: 30-epoch × 3-fold A/B sweep (#204) — decision gate per spec §0
Closes the kernel-correctness portion of #203.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
8d3efa450e |
test(ml-backtesting): integrated Ring 2 fuzz over full pipeline (C18)
Existing lob_sim_fuzz only exercised book_update + market orders. This new test suite drives the FULL integrated pipeline under random adversarial conditions: apply_snapshot (random-walk book) → step_resting_orders (random signed trade-flow signal) → broadcast_alpha (random per-horizon probs every 4th event) → step_decision_with_latency (mixed latency=0 + latency=50ms cells) → submit_market_immediate (immediate path) → pnl_track_step + isv_kelly_update_on_close Warm-starts every backtest with random-but-plausible Kelly state (at least h4 set credibly profitable so decisions actually open positions). Random target_annual_vol + annualisation_factor + max_lots per decision to vary the Kelly cap. Per-50-event invariants: • book monotonicity (bid_px[k] ≤ bid_px[k-1], ask_px[k] ≥ ask_px[k-1]) • no-crossed-book (ask[0] > bid[0]) • Pos.realized_pnl + Pos.vwap_entry finite (no NaN/Inf leaks) • Pos.position_lots in plausible range (|≤100|) • All 5 per-horizon IsvKellyState fields finite Three test sizes: integrated_fuzz_n1_short — N=1, 200 events integrated_fuzz_n8_medium — N=8, 500 events integrated_fuzz_n64_long — N=64, 1000 events All three pass on RTX 3050. The N=64 × 1000 case exercises 64,000 event-snapshots × 16,000 decisions × ~12,800 trade attempts without any invariant violation or NaN propagation. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
dd8d731dcf |
test(ml-backtesting): Ring 3 production-day replay validation (C17)
Closes the spec §8 Ring 3 deferral. Two integration tests against
FOXHUNT_TEST_DATA real MBP-10 day files:
buy_and_hold_full_day
Walk the first .dbn.zst from open to close; submit a market buy
of 1 lot at the first event, run the full step-loop
(apply_snapshot + step_resting_orders) over every subsequent
event, then close with a market sell at the last event. Assert
realized P&L matches (close_mid - open_mid) within ±2 ticks per
spec slippage budget. Proves the book-walking + position
accounting + step-loop orchestration are wire-level correct
against real data, not just hand-crafted JSON fixtures.
walk_book_one_full_day
Same fixture, no trades — just iterates every event through
apply_snapshot + step_resting_orders and asserts the book
invariants (bid/ask monotonicity, no-crossed-book) hold at every
10_000-event checkpoint across the full day. Stress test of
the kernel against real market microstructure (gaps, locked
markets, regime shifts).
Both tests skip gracefully when FOXHUNT_TEST_DATA is absent OR the
predecoded sidecars are empty (the current 40-byte placeholder
fixtures in test_data/futures-baseline/ES.FUT/*.predecoded.bin will
trigger the skip path; populating those with a real Databento decode
makes both tests fire end-to-end). Ring 3 is operational, not
CI-mandatory.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
19986c8d96 |
feat(ml-alpha): tick-rule signed trade-flow inference at L1 (C16)
Replaces the placeholder `trade_signed_vol = trade_count_delta` (always
non-negative, sign-neutral) with a proper Databento-standard tick-rule
inference applied to L1 size + price deltas across consecutive MBP-10
snapshots:
• ask_px[0] unchanged AND ask_sz[0] decreased → aggressive buys
consumed ask depth; add (prev.ask_sz − cur.ask_sz).
• bid_px[0] unchanged AND bid_sz[0] decreased → aggressive sells hit
the bid; subtract (prev.bid_sz − cur.bid_sz).
• ask_px[0] moved up → previous best ask cleared; add prev.ask_sz.
• bid_px[0] moved down → previous best bid hit; subtract prev.bid_sz.
• Pure cancellations (size shrank but price moved AWAY from us) =
ambiguous; ignore.
Convention matches `Mbp10RawInput::trade_signed_vol`: positive =
buyer-initiated, negative = seller-initiated. This is a LOWER-BOUND
estimator — won't catch trades that cleared multiple levels (those
manifest only via deeper-level deltas) or trades against hidden /
off-book liquidity. Acceptable for v1 queue-decay signal; production
deployments can layer the trades-stream loader for ground-truth flow.
Wired into BacktestHarness::run() → sim.step_resting_orders(ts, vol)
so the queue-decay branch of resting_orders.cu finally fires with
non-zero input. Previously the harness passed 0.0 unconditionally,
which meant resting limits could only fill via the price-cross
marketability branch — same-price queue-decay was inert.
Six new unit tests cover each branch of the inference (pure cancel,
ask-shrank, bid-shrank, ask-px-up, bid-px-down, mixed both-sides).
34 ml-alpha lib tests + 33 ml-backtesting lib + 12 GPU fixtures + 3
fuzz still green.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
3fa215ad2e |
feat(ml-alpha): CfcTrunk save/load checkpoint + --checkpoint CLI (C14)
CfcTrunk::save_checkpoint(path) reads each device weight tensor back
via memcpy_dtoh and bincode-serialises into a CheckpointV1 envelope:
{ version, n_in, n_hid,
w_in, w_rec, b, tau,
heads_w, heads_b,
proj_w, proj_b, proj_g, proj_n }
Total ~22k-25k f32 = ~90 KB per trunk. Tiny.
CfcTrunk::load_checkpoint(dev, cfg, path) deserialises + validates
(version == 1, n_in/n_hid match the supplied CfcConfig — a model
trained for one arch can't silently load against another). Constructs
a fresh trunk via new_random (for kernel bindings + scratch buffers)
then overwrites every weight tensor via memcpy_htod. The random init
values are thrown away — marginally wasteful, but keeps the
construction code paths unified.
Roundtrip test (--ignored, CUDA-required): save trunk_A → load → read
back every device tensor and assert bit-equality between trunk_A and
the loaded trunk_B. Passed locally. Dim-mismatch rejection test runs
without CUDA (verifies bincode envelope serialise/deserialise).
bin/fxt-backtest --checkpoint <path>: when set, overrides --seed and
loads from disk. When absent, warns loudly that the trunk is
random-initialised and backtest results are noise. This makes the
binary genuinely useful as a deployment tool — point it at a trained
checkpoint and run real backtests.
Adds bincode workspace dep to ml-alpha (was already in workspace
dependencies, just not in ml-alpha's [dependencies] block). serde
features bumped to ["derive"] (was using workspace default which
omits derive macros).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
ed19985c95 |
feat(ml-backtesting): bytecode VM dispatcher for Strategy compositions (C13)
The hardcoded WeightedByRealizedSharpe path from C7 is now joined by
a stack-based bytecode interpreter that consumes Strategy::flatten()
output, unlocking RegimeSwitch / Portfolio / non-default Ensemble /
single-Leaf compositions specified via policy-grid YAML.
cuda/decision_policy.cu — new kernel `decision_policy_program`:
Stack-based VM with parallel (value, attribution_mask) stacks.
Opcodes mirror src/policy/mod.rs::OpCode exactly:
NoOp / PushScalar / EvalRegime / BranchIfRegime
EmitPerHorizonSize (computes sized intent from alpha[h] +
IsvKellyState[h] using same Kelly + ISV-cap formula as the
hardcoded default)
AggMean / AggWeightedSharpe / AggMaxConfidence (pop n values,
push aggregated, OR attribution masks)
ApplyConflict (v1 no-op, reserved for Portfolio)
WriteOrder (terminal — converts top-of-stack to market_target +
open_horizon_masks attribution if currently flat)
AggWeightedSharpe recovers the source horizon from a single-bit
attribution mask to look up recent_sharpe; multi-bit masks (nested
aggregators that collapsed horizons) fall back to uniform weight.
decision_policy_default extended with a program_lens param: skips any
backtest whose plen > 0 (the program kernel handled it). The two
kernels run sequentially in step_decision_with_latency with mutual
exclusivity on each backtest slot.
LobSimCuda gains:
upload_program(b, &Program) — uploads a Strategy::flatten() output
to backtest b's slot in program_table_d, updates program_lens_d.
set_regime(b, regime_id) — writes regimes_d for OP_EVAL_REGIME /
OP_BRANCH_IF_REGIME consumption.
BacktestHarnessConfig.strategies (Vec<Strategy>) — empty means every
cell uses the hardcoded default; non-empty len must equal n_parallel
and each strategy is flattened + uploaded at construction.
bin/fxt-backtest --policy-grid <yaml> path now actually plumbs through
to the kernel (was parsed-but-ignored in C9). Empty grid keeps the
default behaviour.
New fixture decision_program_h4_only: uploads Strategy::Leaf(h4_only)
flattened to (EmitPerHorizonSize, WriteOrder) — 2 instructions, 16
bytes — and verifies the bytecode kernel produces equivalent end-state
to the hardcoded default (3 lot buy at vwap=5500). Proves the VM
dispatcher works end-to-end.
12/12 GPU fixtures green. 33 lib unit tests + 3 Ring 2 fuzz tests
still green.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
344d7a67d7 |
feat(ml-backtesting): wire --latency-ns end-to-end + fixture (C12)
The in-flight machinery from C11 is now reachable from the harness +
CLI. New step_decision_with_latency on LobSimCuda:
latency_ns == 0 → existing immediate-match path (submit_market_immediate
kernel fills against current book).
latency_ns > 0 → host reads market_targets, converts each non-noop
into seed_limit_order(active=2, price=very-aggressive,
arrival_ts_ns = current + latency_ns). The
resting_orders kernel promotes + fills at arrival
time, so the order sees whatever book exists then —
not the book at decision time.
Aggressive-price heuristic: buy at 1e9 / sell at 0 — the marketability
check (price ≥ best ask for buy, ≤ best bid for sell) unconditionally
crosses at arrival; the kernel then walks book levels for the actual
fill price. This models "market order with latency" correctly because
slippage emerges from the book-walking at arrival, not from a limit
price boundary.
BacktestHarness gains a latency_ns field; harness::run() always calls
step_decision_with_latency now. Also calls sim.step_resting_orders
per event (with trade_signed_vol=0.0 — the Mbp10RawInput trade-flow
hookup is deferred to a trades-feed integration commit) so in-flight
orders get a chance to promote on every snapshot.
bin/fxt-backtest no longer ignores --latency-ns; the flag value
propagates through BacktestHarnessConfig.latency_ns into the kernel
path. Default stays at 100_000_000 (IBKR + Scaleway baseline per
spec §4). Setting --latency-ns 0 selects the legacy immediate path.
Adds latency_in_flight_miss GPU fixture: limit at 5510 submitted
active=2 at T=1s with arrival_ts=T+100ms. step_resting at T+50ms
sees no promotion (still in-flight). Book moves +5 against buyer
during the 100ms window. step_resting at T+150ms promotes the
limit and fills it at the WORSE post-move book (vwap=5505 vs the
original 5500 it would have hit without latency). Slot ends active=0.
11/11 GPU fixtures green on RTX 3050. 33 lib tests still green.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
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
|
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> |