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>
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>
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>
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>
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>
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>
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>
Revises all 12 issues from the self-review pass:
1. (HIGH) Soften bit-equivalence claim — single-sample helper and
batched kernel at B=1 are different CUDA kernels; FP order may
differ. Acceptance is relative_eq ≤ 1e-6, not bit-exact.
2. (HIGH) Explicit asymmetry: only cfc_step has a single-sample GPU
oracle. GRN/VSN/attn bwd rely on smoke + chain-rule preservation.
feedback_no_cpu_test_fallbacks.md forbids a CPU reference oracle.
3. (HIGH) Realistic targets — 3× floor, 5× stretch. Drops 15× claim
which was Amdahl-bounded under any reasonable assumption.
4. (MED) AdamW-after-reducer invariant stated explicitly: final grad
buffer is OVERWRITE by reducer, meaningful only after reducer ran.
5. (MED) New B=32 smoke test (stacked_trainer_loss_shrinks_at_batch_32)
actually exercises the cross-batch reduction path; existing
perception_overfit suite is all B=1.
6. (MED) Rollout commit 1 bundles reduce_axis0 + first consumer
(cfc_step refactor) to avoid feedback_wire_everything_up.md
orphan-kernel anti-pattern.
7. (LOW) Drop "merge to ml-alpha-phase-a" — user already chose direct
commits to that branch; clarify in Rollout.
8. (LOW) Add explicit scratch-sizing formula:
scratch_bytes ≈ B × Σ(param tensor sizes per kernel).
9. (LOW) Remove resolved open question (memset_zeros ordering).
10. (STRUCT) Post-refactor bottleneck analysis section — names the
next L40S floor (Mamba2 scan, launch latency, cuBLAS).
11. (STRUCT) cudaFuncSetAttribute note — refactored cfc_bwd's per-
block shared mem drops to 1 KiB; no attribute change needed.
12. (STRUCT) Explicit Rollback section — atomic-commit-per-kernel +
revert strategy; rollback baseline is the spec commit
(54aa69c10) on ml-alpha-phase-a.
Plus locks the target pool to L40S — speedup must be attributable to
the refactor, not to a hardware bump. H100 / BF16 / larger batch
become candidates for a follow-up spec once the L40S floor is known.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
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>
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>
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>
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>
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>
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>
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>
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>
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>
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>
Three correlated changes for the next CV round:
1. Mamba2 state_dim cap: 16 → 32
cuda/mamba2_alpha_kernel.cu: MAMBA2_ALPHA_MAX_STATE_D 16 → 32.
Per-thread state register `float x[32]` (128 B/thread) and
per-thread x_hist replay cache `float x_hist[K*32]` (up to
12 KiB/thread of local memory at K=96). L40S/H100 register file
(256 KiB/SM) absorbs this without occupancy collapse for our
block dims (32-128 threads). Update Rust-side
MAMBA2_KERNEL_STATE_MAX + validation message + test name. Kernel
header doc updated.
2. New early-stop option: auc_h6000
examples/alpha_train.rs: add the long-horizon AUC as a third
early-stop metric. The ISV CV (3a196382f, 5d42ab0e9, 0171c8c0e)
showed mean_auc-best-epoch and h6000-best-epoch can differ by
1-2 epochs and the h6000 gap can be 5-6pt within a single run
(fblb2 fold-1: saved E10 h6000=0.681, but E11 h6000=0.739 — we
threw away the deployment-better checkpoint). For multi-minute
trading deployment we want the h6000-best checkpoint directly.
3. Summary JSON: best_auc_h6000_epoch / best_auc_h6000 /
best_auc_h6000_per_horizon
So the analysis tooling can see the h6000-best checkpoint
independently of mean_auc / val_loss bests.
Test rename: test_mamba2_config_rejects_state_over_16 →
test_mamba2_config_rejects_state_over_32 (tests now reject state_dim=33).
build.rs cache-bust v4 forces cluster nodes to recompile the kernel
against the new MAMBA2_ALPHA_MAX_STATE_D — old cubins from previous
SHAs were sized for state_d=16 and would silently truncate state_d=32
state arrays.
Validation: 26 lib + 23 integration ml-alpha tests pass. Mamba2 block
tests use state_dim=8 or 16 (well below the new cap), exercise both
the forward + backward + AdamW paths.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
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>
3-fold A/B CV (5d42ab0e9 vs eb51c0f9c, same data splits) showed
ISV winning the aggregate (+0.9pt mean_auc, +1.8pt h6000 across
folds) but FOLD-2 regressed -1.3pt on h6000 while folds 0 and 1
both gained (+2.0pt and +4.6pt respectively).
Per-fold per-horizon breakdown showed exactly the failure mode:
Fold 0 (val 2025-Q1): h6000 no-ISV 0.714 → ISV 0.734 (+2.0pt)
Fold 1 (val 2025-Q2): h6000 no-ISV 0.682 → ISV 0.728 (+4.6pt)
Fold 2 (val 2025-Q3): h6000 no-ISV 0.698 → ISV 0.685 (-1.3pt)
In folds 0 and 1, h6000 was the hardest horizon — ISV correctly
boosted it (lambda > 1). In fold 2 the regime made h6000 relatively
easy at no-ISV (0.698 vs the worst horizon at ~0.70). ISV's
SYMMETRIC clamp [0.5, 2.0] then computed ratio = ema_h6000 /
mean_ema < 1 and DEMOTED h6000's trunk-gradient pull below uniform,
starving further learning on the horizon we actually deploy.
Per `pearl_audit_unboundedness_for_implicit_asymmetry.md`: when a
control signal serves an asymmetric goal (here: we never want to
de-prioritize h6000, only ever boost it OR leave it alone), encode
that asymmetry in the clamp. LAMBDA_FLOOR 0.5 → 1.0 makes the
controller boost-only: under-trained horizons get more pull, but
no horizon is ever demoted below its uniform contribution.
Expected effect with asymmetric clamp:
- Fold 0 and 1: lambda for the hardest horizon stays at 2.0
(ceiling-clamped), trunk-gradient lift unchanged. The gains
+2.0pt and +4.6pt should hold.
- Fold 2: h6000's lambda was being demoted to ~0.74; now floored
at 1.0 — the -1.3pt h6000 regression should disappear. h6000
trains at uniform weight, recovering toward 0.698.
- Cross-fold mean projected: ~0.724 (+1.3pt vs no-ISV).
Test update: `horizon_ema_and_lambda_track_after_training` now
asserts lambda ∈ [1.0, 2.0] (the asymmetric envelope) and
lambda mean ∈ [1.0, 2.0] (boost-only guarantee). Observed values
on the test data: lambda = [1.13, 1.00, 1.08, 1.08, 1.00] — the
two easier horizons correctly floor at 1.00 instead of demoting.
Validation: 7 perception_overfit tests pass, synthetic overfit
trajectory unchanged (0.36 → 0.0006), 26 ml-alpha lib + 23
integration tests green.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Connects the per-horizon lambda computed by horizon_ema_and_lambda
(landed in 37c3a8f4d) to the actual trunk gradient flow. This is the
model-behavior change. mhzs7 (3a196382f) hit mean_auc=0.726 but with
the long-horizon h6000 stuck at 0.694 — exactly the horizon the
auto-horizon-weights formula `min(1, K/h)` over-weights down (h6000
weight = 0.0053). ISV detects "this horizon is hard, give it more
trunk-gradient pull" and lambda[h] scales the per-horizon `d_z`
contribution into the trunk in heads_bwd.
Kernel change (cuda/multi_horizon_heads.cu):
multi_horizon_heads_backward_batched(): new arg
`const float* __restrict__ lambda` (5 elements, between
grad_h_carry and n_batch in arg order).
- lambda is broadcast into __shared__ float s_lambda[5] once per
block so every thread reads the 5 floats without repeated
global loads.
- Sentinel: zero buffer (init state, EMA hasn't run yet) is read
as 1.0 so the kernel reduces to pre-ISV behavior. After step 1
every entry is clamped to [0.5, 2.0] by horizon_lambda.cu and
the > 0 check is always true.
- lambda[k] multiplies the `acc += s_lambda[k] * sd_z * w[k]`
term that produces grad_h (trunk gradient).
- grad_w and grad_b updates are UNCHANGED. The horizon heads keep
learning their own weight/bias normally; lambda only biases how
much each horizon's error signal leaks into the shared trunk.
Per `pearl_adam_normalizes_loss_weights.md`: scaling the
effective gradient bypasses Adam's m/sqrt(v) normalization that
would otherwise cancel a loss-weight lift.
Trainer change (trainer/perception.rs):
- heads_bwd_batched launch now passes `&self.lambda_d` between
grad_h_carry_d and n_batch_i. lambda_d is refreshed each step
by the horizon_ema_and_lambda kernel that ran just after BCE.
- Whole chain lives inside the captured CUDA Graph.
Validation:
- All 7 perception_overfit tests pass.
- Synthetic overfit converges marginally FASTER than pre-Phase-3
(final loss 0.0007 → 0.0005). Plausible explanation: in a
constant-signal smoke, the per-horizon BCE values are similar
enough that lambda mostly stays near 1.0, but the EMA
bootstrap on step 1 still produces non-uniform initial lambdas
that nudge the trunk toward whichever horizon converges
slowest. Either way, no regression.
- horizon_ema_and_lambda_track_after_training test still passes
with the lambda values now flowing through heads_bwd:
ema = [0.50, 0.87, 0.64, 0.49, 0.60]
lambda = [0.80, 1.41, 1.03, 0.79, 0.97]
→ h100 hardest (BCE 0.87) gets 1.41× trunk grad; h300 easiest
(BCE 0.49) gets 0.79× — exactly the intended ISV semantics.
- 26 lib + 23 integration ml-alpha tests pass.
Honors:
- feedback_isv_for_adaptive_bounds.md (no hardcoded constants;
lambda derives from runtime BCE signal).
- pearl_adam_normalizes_loss_weights.md (scaling trunk gradient,
not BCE coefficient, to bypass Adam normalization).
- pearl_audit_unboundedness_for_implicit_asymmetry.md (lambda is
bounded by horizon_lambda.cu's clamp [0.5, 2.0]).
- pearl_no_deferrals_for_complementary_fixes.md — Phase 3 wires
up the infrastructure from Phase 1+2 in the same logical
delivery rather than waiting a CV cycle.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Foundation for replacing the static `--auto-horizon-weights` formula
(`min(1, K/h)`) with a signal-driven per-horizon gradient scaler. Per
`feedback_isv_for_adaptive_bounds.md`: adaptive bounds live in ISV,
not hardcoded constants. Per `pearl_adam_normalizes_loss_weights.md`:
Adam normalizes per-loss weight lifts (SP13 saw 13× aux_w produce
only 0.6%/epoch divergence), so the effective lever is scaling the
GRADIENT into the shared trunk, not the BCE coefficient. This commit
sets up the EMA + lambda infrastructure; Phase 3 (wiring lambda into
heads_bwd to actually scale the trunk gradient) is gated on the
3-fold CV results from eb51c0f9c.
Phase 1 — BCE kernel emits per-horizon UNWEIGHTED mean BCE:
cuda/bce_loss_multi_horizon.cu:
- New output buffer `loss_per_horizon[N_HORIZONS=5]`.
- Per-horizon shared-mem accumulators (sloss_h, svalid_h) with
block tree-reduce — no atomicAdd, per `feedback_no_atomicadd.md`.
- Hardcoded N_HORIZONS_BCE=5; total shared-mem usage ~13 KiB
(comfortable under any SM smem limit).
- The aggregate `loss_out` is still the externally-weighted mean
callers use for reporting; the new buffer is the UNWEIGHTED
signal an EMA layer needs.
Phase 2 — EMA + lambda kernel:
cuda/horizon_lambda.cu (new):
- Single-thread kernel (5 horizons, fixed-size loop — trivial).
- First-observation bootstrap via sentinel = 0 per
`pearl_first_observation_bootstrap.md`; replaces directly when
`loss_ema_h <= 0` (safer than `== 0` under --use_fast_math).
- Fixed α = 0.1 EMA for now; Wiener-optimal α follow-up flagged
(`pearl_wiener_optimal_adaptive_alpha.md`).
- lambda_h = clamp(loss_ema_h / mean(loss_ema), 0.5, 2.0).
- Ratio gives natural "under-trained → boost" signal.
- Clamp prevents winner-take-all per
`pearl_controller_amplifies_dominant_magnitude_trap.md`
and bounded-modifier safety per
`pearl_audit_unboundedness_for_implicit_asymmetry.md`.
Trainer wiring (trainer/perception.rs):
- 3 new fields: `loss_per_horizon_d`, `loss_ema_d`, `lambda_d`
(all 5-element f32 CudaSlices; pre-allocated, zero-initialised).
- Cached `horizon_lambda_fn` + module handle per the
`BiasKernels`-style pattern (no per-call cuModuleLoadData).
- BCE callsite (train + eval paths) updated to pass
`loss_per_horizon_d`.
- `dispatch_train_step` launches `horizon_ema_and_lambda` right
after BCE, BEFORE the K-loop backward. Inside the captured
graph; per-step launch overhead is ~1 µs.
- `loss_ema_snapshot()` + `lambda_snapshot()` test-only accessors
(mapped-pinned readback, not for hot path) for diagnostics.
Smoke test — `horizon_ema_and_lambda_track_after_training`:
- Verifies pre-step EMA + lambda are zero (sentinel).
- After 5 training steps:
loss_ema = [0.59, 0.66, 0.45, 0.62, 0.50] — finite + positive.
lambda = [1.05, 1.16, 0.80, 1.10, 0.89] — mean ≈ 1.0, all
inside the [0.5, 2.0] clamp envelope.
- Lower per-horizon BCE → lower lambda (de-emphasize); higher
BCE → higher lambda (boost). Exactly the ISV semantics we want.
Validation: 6 perception_overfit tests pass, synthetic overfit still
shrinks (0.33 → 0.0006), 26 ml-alpha lib + 23 integration tests
green. lambda_d is computed every step but NOT YET CONSUMED by
heads_bwd; training behavior is bit-identical to 3a196382f. Phase 3
(consume lambda_d in heads_bwd_batched to scale the per-horizon
gradient into the trunk) follows once CV confirms the foundation is
stable.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
mhzs7 reported val mean_auc=0.726 on 3a196382f — but the trainer
constructed both train and val MultiHorizonLoader with the SAME
`mbp10_root: cli.mbp10_data_dir`. The two loaders only differed by
seed. So val sequences were held-out-by-anchor from the same files
train sampled from; not temporally OOS. Per
`pearl_single_window_oos_is_not_oos.md` a single-window result that
doesn't enforce time-ordered separation can collapse across true
walk-forward folds.
Refactor: drop `mbp10_root` from `MultiHorizonLoaderConfig` (which
forced caller to share the dir between train and val). New API takes
an explicit `files: Vec<PathBuf>` — the loader preserves the order
given and does no internal shuffle, so callers control temporal
ordering. Added `discover_mbp10_files_sorted(root)` helper that
enumerates a dir and sorts by filename (chronological under the
`ES.FUT_<YEAR>-Q<n>.dbn.zst` convention).
alpha_train.rs splits the discovered files by 3 new CLI flags:
--cv-fold <k> (default 0)
--cv-n-folds <N> (default 1 — single fold)
--cv-train-window <W> (default 0 — auto)
Single-fold default (cv_n_folds=1): train on all files except the
last, val on the last file. This replaces the old "same files for
both" bug; even runs that don't think about CV now get a temporal
split by default.
Sliding-window CV (cv_n_folds > 1): fold k trains on files
[k..k+W] and validates on file [k+W]. With 9 quarterly files
(2024-Q1..2026-Q1) and `--cv-n-folds 3`, the natural layout is:
fold 0: train 2024-Q1..2024-Q4 (W=4) → val 2025-Q1
fold 1: train 2024-Q2..2025-Q1 → val 2025-Q2
fold 2: train 2024-Q3..2025-Q2 → val 2025-Q3
blind holdout: 2025-Q4, 2026-Q1
Threaded the flags through scripts/argo-alpha-perception.sh and
infra/k8s/argo/alpha-perception-template.yaml so each fold submits
as an independent workflow.
Updated tests/multi_horizon_loader.rs to the new API:
loader_yields_seq_with_valid_labels — exercises discover + load.
loader_errors_on_empty_files — replaces missing-root test.
discover_errors_on_missing_root — pinpoints the discover step.
Honors:
- feedback_no_partial_refactor.md — every consumer migrated atomically.
- feedback_no_legacy_aliases.md — no `mbp10_root` shim left behind.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
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>
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>
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>
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>
The four biggest per-call scratch buffers in
Mamba2Block::backward_from_h_enriched_seq were allocated fresh on
every training step:
d_a_per_channel [N, sh2, K, state_d] ~6 MB at B=8, K=96
d_b_per_channel [N, sh2, K, state_d] ~6 MB
d_w_c_per_sample [N, sh2, state_d] ~64 KB
d_h_s2 [N, sh2] ~4 KB
For 2000 optimizer steps/epoch × 15 epochs = 30 000 alloc_zeros
calls per training run, all on the hot path.
New `Mamba2BackwardScratch` struct holds these as device-resident
buffers, constructed once per (n_batch, seq_len, hidden_dim,
state_dim) at trainer init. New `backward_from_h_enriched_seq_into`
method takes the scratch by &mut and reuses the buffers each call.
The smaller per-call buffers (d_a_proj_flat, d_b_proj_flat, dw_c)
still allocate per call — they feed into ownership-transferring
LinearGrads outputs, where pre-allocation would require refactoring
ml-core's cuBLAS wrappers without proportional gain.
The original `backward_from_h_enriched_seq` is preserved (Phase E.3
callers still use it). Trainer switches to the `_into` variant.
Expected per-step savings: ~10-20 ms on L40S (4 cudaMalloc latencies
per call × 2.5-5 μs each + cache pressure reduction). Over 2000
steps/epoch that's 20-40 sec/epoch.
77 ml-alpha tests pass. Synthetic overfit unchanged (0.27 → 0.0006).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
cnjfl wall-time analysis: 80% of training time was disk IO. The
MultiHorizonLoader was constructed fresh per epoch in the CLI loop,
forcing 9 file deserializations from bincode (~50s/file × 9 = 7-8
min) per epoch. For a 6-epoch run that was ~45 min of pure IO out
of ~50 min total.
Now loaders preload ALL files into RAM at construction (one-time
~7 min startup) and expose a `reset(seed: u64)` method that
re-seeds anchor sampling per epoch — no disk IO between epochs.
Memory cost: ~13-15 GB for the 9-quarter ES.FUT dataset (45M
snapshots × ~280 bytes). Well under the training pod's 64 GB
limit; current 16 GB request remains sufficient since the resident
set fits.
Expected wall-time impact at K=96, B=8, 16K seqs:
6 epochs: ~50 min → ~13 min (~3.7×)
15 epochs: ~120 min → ~23 min (~5×)
The internal LoadedFile-cache-with-cycling logic is gone — the
loader now holds Vec<LoadedFile> with all files resident. Per-call
`next_sequence` picks a uniformly random file + uniformly random
anchor inside it (instead of cycling files with a per-file budget).
Distribution is equivalent: each file contributes ~n_max_sequences /
n_files samples per epoch in expectation.
77 ml-alpha tests pass.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
cnjfl evidence: val_loss and mean_auc disagree.
val_loss best at e3 (0.5592)
mean_auc best at e4 (0.7670 — new h300 + h6000 peaks)
For downstream trading, ranking quality (AUC) matters more than
probability calibration (BCE loss). New default is mean_auc-based
early stopping, but val_loss/none remain selectable.
AUC is noisier than loss epoch-to-epoch (1-2pt bounces are common
even when long-horizon AUCs are still drifting up under
auto-horizon-weights), so patience defaults bump from 3 → 5.
CLI: --early-stop-metric {val_loss|mean_auc|none} default mean_auc
--early-stop-patience N default 5
Argo template parameters added with matching defaults.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
cnjfl run showed val_loss and mean-AUC peak at different epochs:
epoch 3: val_loss=0.5592 (best) mean_auc=0.7608
epoch 4: val_loss=0.5609 (worse) mean_auc=0.7670 (best — new h300 + h6000 peaks)
val_loss tracks probability calibration; AUC tracks ranking quality.
For downstream trading the ranking profile matters more — so we now
publish both bests in alpha_train_summary.json and log a "new best
mean_auc" line whenever a new mean-AUC peak lands.
Early stopping still gates on val_loss (the two policies stay
decoupled — mean-AUC is reported-only).
New summary fields:
best_mean_auc_epoch
best_mean_auc
best_mean_auc_per_horizon
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
This kernel was referenced by PerceptionTrainer.step_batched /
evaluate_batched (commit c3ee5e165) but its CUDA source had never
actually been committed — the .cu edit lived only in my working
copy. Cluster runs at 8851e98bd and earlier failed during trainer
init with `CUDA_ERROR_NOT_FOUND` because the cubin lacked the
symbol the Rust code tried to load.
Discovered via `strings` on the cluster binary: 12 occurrences of
every other kernel name (cubin export string + Rust load string),
but only 1 occurrence of `transpose_3d_swap_01` (the Rust load
string alone).
Local builds passed because they were built against the working
copy which DID have the kernel. Hard cluster failure exposed the
gap.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Cluster ensure-binary at c3ee5e165 produced a binary whose cubin
lacked the batched kernel symbols (CUDA_ERROR_NOT_FOUND on
cfc_step_batched at trainer init). Local cubins have all symbols —
the cluster's persistent /cargo-target PVC appears to have cached
pre-batched build artifacts that cargo's incremental compilation
deemed up-to-date despite the .cu source changes.
Forcing a fresh build.rs run via a sentinel comment bumps the
build script's fingerprint, which invalidates all build.rs outputs
(cubins) and re-invokes nvcc for every .cu file in KERNELS.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Adds two new workflow parameters with backward-compatible defaults
(batch-size=1, auto-horizon-weights=false) so existing submissions
behave identically. Both flags are forwarded to alpha_train CLI.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Plumbs the batched cfc + heads kernels added in 829ddfa62 through
PerceptionTrainer, evaluator, and the alpha_train CLI:
PerceptionTrainerConfig.n_batch — batch size, default 1
PerceptionTrainer::step_batched — process B sequences per
optimizer step using
cfc_step_batched / heads_batched
PerceptionTrainer::evaluate_batched — forward-only batched eval
PerceptionTrainer::step / evaluate — thin B=1 wrappers preserving
existing single-sequence
test/inference APIs (assert
cfg.n_batch == 1)
alpha_train CLI: --batch-size N — accumulates B sequences per
optimizer step in train loop;
val loop also batches and uses
evaluate_batched
Per-K scratch buffers all grow to [K, B, dim] layout (K-major, slot-k
contiguous). Mamba2's [B, K, H] output is transposed once after
forward via the new transpose_3d_swap_01 kernel, and grad_h_enriched_seq_t
is transposed back to [B, K, H] before Mamba2 backward. Two transposes
per training step; negligible (1.5MB at B=32).
Dead unbatched kernel handles removed from the trainer (step_fn,
step_bwd_fn, heads_fn, heads_bwd_fn, grad_x_d) — all training and
inference now go through the batched variants for B ≥ 1. The
single-sample kernels remain in CUDA for the standalone test helpers
in cfc/step.rs and heads.rs.
Local 2Q smoke (seq_len=32, B=4, --auto-horizon-weights, 800 train
seqs × 2 epochs):
epoch 0: val_loss=0.7138 AUC h30/h100/h300/h1000/h6000 = .55/.55/.57/.61/.51
epoch 1: val_loss=0.6558 AUC h30/h100/h300/h1000/h6000 = .72/.68/.75/.68/.65
vs the in-flight qf5mj baseline (B=1, K=96, no horizon weighting) which
had val_loss=0.6933 best and AUCs oscillating at ~0.50 — this batched
run hits AUC 0.75 (h300) and 0.72 (h30) in just 2 epochs of 200
optimizer updates. Batching + horizon-weighting unblocks the model.
77 ml-alpha tests pass.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
