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>
Adds 4 new kernel symbols alongside the existing single-sample ones —
zero changes to current call sites, so the in-flight qf5mj baseline is
unaffected. The next commit wires these into PerceptionTrainer's K
loop and exposes --batch-size in the CLI.
cfc_step_batched processes [n_batch, n_in/n_hid] tensors
cfc_step_backward_batched same; shared mem holds sd_pre[B, n_hid]
+ sdecay[n_hid] (~16 KiB at B=32, well
under L40S 48 KiB block limit). Param
grads (grad_b/grad_w_in/grad_w_rec/
grad_tau) accumulated via += — thread i
is sole writer to its row across all
samples, so no atomicAdd and no per-
batch scratch buffer.
multi_horizon_heads_batched [n_batch, 5] sigmoid outputs from
[n_batch, 128] hidden inputs.
multi_horizon_heads_backward_batched
shared mem holds sd_z[B, 5]. grad_w
/ grad_b += across batch (thread tid
sole writer). grad_h carries the
optional per-sample grad_h_carry
(cfc recurrence chain).
Design notes:
- Threading: one block of n_hid threads. Each thread loops over
b ∈ 0..B internally. This avoids cross-block races on grad_*
buffers and keeps the existing "no atomicAdd" discipline. Cost:
less raw parallelism than grid-batching, but the bottleneck is
Mamba2 (already batch-parallel via its own kernel grid).
- Per-thread accumulators: grad_b / grad_tau land in registers,
flushed once at end. grad_w_in / grad_w_rec written += per-b
(thread sole writer to its row, safe).
- All B samples processed in stream order inside one kernel launch
— saves K * (B-1) launches per sequence vs serialising B
independent calls.
77 ml-alpha tests pass (kernels not yet exercised — wiring is the
next commit).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
For seq_len K and horizon h with h ≫ K, the K position-supervised
labels in a single sequence are near-identical (sequential positions'
forward windows overlap by ~(h-1)/h). Per-position BCE therefore
treats ~K highly-correlated labels as independent samples, inflating
gradient pressure on long horizons by a factor of K.
Concretely at K=96:
h=30 → ~3 effective samples per seq (forward windows overlap ~97%)
h=100 → ~1 (~99%)
h=6000 → ~1 (~99.98%)
Per-position supervision was paying 96× the natural signal density on
h=6000, pulling the model toward fitting noise at long horizons.
Fix: the fused BCE kernel now accepts an optional
`loss_weights[N_HORIZONS]` (nullptr → uniform = no-op). Each (k, h)
loss + grad contribution is multiplied by w_h; the normaliser is the
sum of weighted valid entries instead of the raw valid count.
`auto_horizon_weights(K, horizons)` computes `w_h = min(1.0, K/h)` so
short horizons stay at full weight and long horizons collapse to
their independent-sample density. Exposed via CLI:
--auto-horizon-weights # K/h auto-derived
--horizon-weights "1,1,0.5,0.1,0.02" # explicit floats
Default behaviour is uniform (1.0) — apples-to-apples with the
in-flight qf5mj baseline. Synthetic overfit still 0.6268 → 0.1144 in
250 steps (82% drop). 77 tests pass.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Comprehensive fix for the issues identified after the BPTT-unroll cluster
run plateaued at val_loss ~0.692 with oscillating AUCs:
ARCHITECTURE
- CfC h_old is now RECURRENT across positions. Previously reset to
zero every step → CfC degenerated to a per-cell tanh-FC layer.
New: h_old at step k IS h_new at step k-1. Heads still operate
on h_new_k, but now the CfC actually carries state. Reverse-order
backward through the K positions accumulates grad_h_old → grad_h_new
via the new optional `grad_h_carry` arg on multi_horizon_heads_backward.
- tau is TRAINED. cfc_step_backward now writes grad_tau (per-cell decay
constant derivative), trainer gets a 7th AdamW group at 0.1× cfc lr.
- 6 NEW regime features (EMA cascade computed loader-side per file)
fill slots out[20..26] of snap_features. Gives the model multi-minute
trend / volatility / liquidity context that is structurally unreachable
inside the K-snapshot BPTT window. Slots: mid-z (med/slow), trend
signal, log-vol slow, log-spread med, log-trade-rate med. All bounded
via log1p / signed-log so no tuned constants leak in.
PERFORMANCE (NVIDIA-style)
- GPU-fused multi-horizon BCE for the entire [K, N_HORIZONS] grid in
ONE launch (was K host roundtrips). Native NaN-label masking.
- K-loop is fully GPU-resident: pre-allocated per-K scratch
(h_new_per_k, probs_per_k, labels_per_k, grad_probs_per_k), zero
device allocs inside step(). Only TWO syncs per sequence (after
forward, after backward) vs previously 2K+1.
- Stream-ordered kernel launches with pointer-offset addressing into
per-K buffers — host doesn't wait between K iterations.
- cfc_step_backward / multi_horizon_heads_backward both use += grad
semantics; trainer pre-zeroes accumulators once per step().
- MAMBA2_ALPHA_MAX_K capped at 96 (was temporarily at 256). 96 covers
h=30/100/300 with room; regime features handle h=1000/h=6000.
TRAINING DISCIPLINE
- LR schedule: linear warmup (default 200 steps) + cosine decay to
lr * lr_min_factor (default 0.1). Applied per training step to both
CfC and Mamba2 AdamW groups via new set_lr_cfc/set_lr_mamba2.
- Best-checkpoint tracking by val_loss; recorded in summary
(best_epoch, best_val_loss, best_val_auc).
- Early stopping on val_loss plateau (default patience = 3).
- CRITICAL BUG FIX: validation now uses new `evaluate()` method
(forward-only) instead of `step()`. Previous CLI called step()
on val data, which ran the full backward + AdamW update on the
validation set. With per-step BPTT that's ~K× more pressure than
the old comment ("statistically negligible") assumed.
Synthetic overfit: 0.6442 → 0.1233 in 250 steps (81% drop, sharper
than the previous 70%). 77 ml-alpha tests pass.
Local 2Q smoke (seq_len=64, 600 train seqs/epoch, 4 epochs):
val_loss 0.7011 → 0.6990, best epoch=1, h300 AUC 0.565 in epoch 0.
Phase E.3 callers (ml/examples/alpha_baseline.rs,
alpha_dqn_h600_smoke.rs) use the LEGACY Mamba2 forward_train +
backward_from_h_enriched path — unaffected by these changes (their
kernels are pre-zeroed via alloc_zeros, so the += grad semantics
remain correct in single-call mode).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Previous BPTT-unroll run had val_loss trending (0.6941→0.6922 over 5
epochs) but AUCs oscillating around 0.50 — the architecture lacked
context for medium/long horizons (h300, h1000, h6000 ≫ seq_len=32).
Phase 1d.2 validated the SSM at seq_len=6000; this is a step toward
restoring useful sequence depth.
Bumps:
- MAMBA2_ALPHA_MAX_K constant: 32 → 256
- x_hist per-thread replay buffer: 2KB → 16KB (spills to
DRAM-backed per-thread local memory; L2-cached, acceptable
perf cost vs the 8x context gain)
- Mamba2BlockConfig::validate updates the cap
Backward compat: legacy Phase E.3 callers (alpha_baseline,
alpha_dqn_h600_smoke) only use K=12 / K=32 — unaffected by the
larger compile-time max.
Synthetic overfit still converges 0.7135 → 0.2079 in 250 steps.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The final-step-only trainer (one BCE prediction per 32-snapshot
window) trained flat at chance on real ES data despite working on
synthetic overfit: train_loss=0.6953, val_loss=0.6943 across 40k
gradient steps. Gradient density was the bottleneck — one supervised
position per sequence × ~8K seqs/epoch isn't enough signal for the
SSM to find the alpha.
This commit supervises the model at EVERY position in the sequence:
mamba2_alpha_scan_fwd_seq — emits h_enriched at every t step
([N, K, sh2] instead of [N, sh2])
mamba2_alpha_scan_bwd_seq — accepts d_h_enriched_seq, injects
gradient at each t before propagating
d_state through the gate chain.
d_w_c and d_h_s2 accumulate across t.
PerceptionTrainer.step() — loop k=0..K; cfc + heads + BCE at
each valid label; cfc/heads grads
accumulate via += in kernel writes.
One Mamba2 backward call consumes the
full grad_h_enriched_seq.
cfc_step_backward — grad_w_in/w_rec/b writes changed
to += (callers MUST pre-zero).
multi_horizon_heads_backward — grad_w/grad_b writes changed to +=.
alpha_train.rs — passes per-position label rows to
step(); AUC still scored from
last-position predictions.
Phase E.3 callers (alpha_baseline.rs, alpha_dqn_h600_smoke.rs) use
the LEGACY Mamba2 forward_train + backward path with `alloc_zeros`
grad buffers — unaffected.
Synthetic overfit still converges 0.6664 → 0.1976 in 250 steps.
Local 2-quarter ES.FUT smoke shows the val AUC at h300 climbing
0.513 → 0.566 over 3 epochs (was flat-at-chance before). First
gradient signal we've gotten through the new architecture.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
First L40S run (alpha-perception-s6hqv, commit 586d1e782) trained flat
at chance: train_loss=0.6953, val_loss=0.6943, AUCs all near 0.50
across 5 epochs and 40k gradient steps. Synthetic overfit on the same
trainer hit 0.59→0.19 in 250 steps, so wiring was sound.
The signal-killer was feature scale: raw size deltas were ±100, dt_ms
could exceed 1e4, and log-returns sat at ~1e-4. Mamba2's W_in
projection saturates on those extremes, gradient bleeds out.
Now in the kernel:
out[0] = (mid - prev_mid) / tick_size [tick-return]
out[12..17]= sgn(d) * log1p(|d|) [signed-log OFI]
out[18] = sgn(v) * log1p(|v|) [signed-log vol]
out[19] = log1p(max(dt_ms, 0)) [log dt]
No tuned constants — tick_size is a market quantity; log1p and
signed-log are monotone bounded transforms. Bit-equiv tests updated
to assert the new closed-form expressions (still GPU-only, no CPU
oracle).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The MultiHorizonLoader was calling load_or_predecode_mbp10 on every
next_sequence() call, deserializing millions of MBP-10 snapshots per
sequence. With 8000 sequences and 9 files this gave ~50+ hour
training time on what should be IO-trivial work.
Now keep one file cached (LoadedFile { snapshots, labels_full }),
yield ceil(n_max_sequences / n_files) sequences from it before
advancing. Per-horizon labels are computed once per file load and
sliced cheaply per anchor. With 8000/9: ~890 loads → 9 loads.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
alpha_train --seed is clap-typed as u64, and clap's default u64
parser only accepts decimal digits. Previous default "0x4242"
hit "invalid digit found in string" at startup.
Replace with decimal equivalent 16962 (= 0x4242) in both the
submission script default and the workflow template default.
Long-term: could add a custom clap value_parser that accepts
hex/dec/oct prefixes, but for now decimal-only matches the
foxhunt convention in other CLIs (alpha_baseline, etc.).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The training-data PVC actually has its MBP-10 .dbn.zst files at
/data/futures-baseline-mbp10/ES.FUT (9 files for ES futures), NOT at
/data/futures-baseline/mbp10. The latter path doesn't exist; the
prev run's bash check fired exit 1 with "MBP-10 data directory not
found".
The PVC root layout (from a probe pod):
/data/bin/ <- compiled binaries by SHA
/data/feature-cache/ <- predecoded sidecar cache
/data/futures-baseline/{ES,NQ,ZN,6E}.FUT/ <- legacy multi-asset
/data/futures-baseline-mbp10/ES.FUT/ <- ES MBP-10 (this is what we want)
/data/futures-baseline-trades/ES.FUT/ <- ES trades
/data/futures-baseline-1s/ES.FUT/ <- 1-second OHLCV
/data/trained-models/ <- model checkpoints
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The CI pipeline (.gitlab-ci.yml stage `build-foxhunt-training-runtime`)
builds and pushes the image as foxhunt-training-runtime:latest (with
the foxhunt- prefix). Other consumers in the repo agree:
- infra/k8s/training/image-prepuller.yaml
- infra/k8s/training/job-template.yaml
The alpha-perception template (and the older alpha-cv template) used
training-runtime:latest without the prefix — broken since the image
never existed at that path. Kubelet kept hitting ErrImagePull /
ImagePullBackOff with "not found".
This was the next blocker after the CPU oversizing fix. Stopping the
in-flight workflow and resubmitting on the corrected template.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Adds a tiny alpine pod (check-cache) that runs first on the platform
pool (no autoscaler delay, ~3 sec end-to-end) and probes the
training-data PVC for /data/bin/$SHA/alpha_train. Outputs:
- sha: short SHA used for binary cache keying
- cache: "hit" or "miss"
ensure-binary now has `when: cache == miss` — when the binary is
already cached for the current SHA, the entire ~4.8GB ci-builder
image pull + sccache compile cycle is skipped. Re-runs on the same
SHA now go straight from submission to training in ~30 seconds
instead of ~3 minutes.
train depends on check-cache + ensure-binary; sources the SHA from
check-cache's output (works whether ensure-binary ran or was
skipped — Argo treats `when:` skip as a satisfied dependency).
Submission script (scripts/argo-alpha-perception.sh) now pre-resolves
commit-sha=HEAD to an actual git SHA via `git rev-parse origin/<branch>`
before submission. This lets the alpine check-cache pod work without
installing git in the container.
Also removed the now-stale `ci-training-h100x2|ci-training-h100-sxm`
case branch from the SM-arch detection — those pools no longer exist
post-pool-cleanup commit a252119fd.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Previous cpu request was 8 (limit 16), but L40S-1-48G allocatable cpu
is 7800m (8 vCPU total minus kubelet overhead). Pod couldn't fit on
the very node we wanted it on — autoscaler provisioned the L40S
cleanly but the scheduler then rejected the pod with
"Insufficient cpu" forever.
Fix: requests cpu=6 / mem=16Gi, limits cpu=7 / mem=64Gi. Leaves
~1.8 vCPU and ~27Gi memory headroom for the system daemonsets
(cilium, csi-node, nvidia driver/device-plugin/dcgm-exporter,
gpu-feature-discovery, node-bootstrap, prometheus-node-exporter,
promtail) that the L40S node hosts.
The trainer itself is GPU-bound (kernels do the work), so 6 vCPU is
plenty for the orchestration host process.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Reasons:
- h100-sxm: Scaleway account quota is 0/0 for the SXM instance type
(cp_servers_type_H100_SXM_2_80G). The pool was perpetually trying
to maintain size=1 with a creation_error node, which JAMMED the
cluster autoscaler. That blocked L40S scale-up entirely during
today's alpha-perception cluster run.
- h100x2: more expensive than current workloads justify. Every
training path (alpha perception, future PPO) fits on either the
single-GPU H100 or L40S.
Removed:
- Two `scaleway_k8s_pool` resources from infra/modules/kapsule/main.tf
- Six variables (enable + type + max_size for each pool)
- Two outputs (pool_id for each)
- Corresponding inputs in infra/live/production/kapsule/terragrunt.hcl
The live cluster has the SXM pool stuck node manually deleted via
`scw k8s node delete` (this commit-session); the pool resource itself
will be destroyed on next `terragrunt apply`.
Post-cleanup pool inventory:
- platform (DEV1-L × 3)
- ci-training-h100 (H100-1-80G, max 1) <- regular single-GPU
- ci-training-l40s (L40S-1-48G, max 1) <- primary training target
- ci-compile-cpu (POP2-HC, max 4)
- ci-compile-cpu-hm (POP2-HM, max 1)
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The actual Kapsule autoscaler_config (infra/modules/kapsule/main.tf
lines 46-52) is scale_down_delay_after_add="10m" +
scale_down_unneeded_time="10m". Combined, a freshly-provisioned node
won't be eligible for scaledown for 10m + another 10m unneeded
before action, so effective grace window is ~20m.
Update the comment in the warmup-gpu template to cite the actual
config rather than the speculative "15m" value. Behavior is
unchanged — the warmup pod still exits immediately and relies on
the grace window to keep the node warm for train.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Removed the 30s sleep. The warmup pod's purpose is to make the L40S
pool autoscaler scale 0 → 1; once the pod lands on the new node
(Scheduled → Running → Succeeded), the node enters Scaleway Kapsule's
scaledown-grace window (~15 min). That single window covers the
entire range of ensure-binary durations (sccache-hit ~10s through
cold ~15 min), so train always lands on a hot node without holding
the warmup pod open.
Trimmed cpu request 100m → 50m and mem 64Mi → 32Mi: the pod runs
~one shell command then exits; tiny resource footprint = faster
scheduling and no kubelet noise.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Adds a parallel warmup-gpu task that runs concurrently with
ensure-binary. The warmup pod is a tiny CPU-only alpine container
scheduled on the gpu-pool's nodeSelector — its presence triggers
cluster-autoscaler scale-up of the L40S pool. After a 30s sleep, the
warmup pod exits; the node enters Scaleway's scaledown grace window
(~10 min), so the train pod lands on a hot node without waiting for
autoscaler provisioning.
No GPU resource request on the warmup pod — that would serialise
warmup and train on the same GPU. nodeSelector + nvidia.com/gpu
toleration are sufficient to force placement on the L40S pool.
Expected savings: ~3-5 min per cold cluster submission. First run
(this session, alpha-perception-4vl7c) showed compile took 113s
(sccache warm) with serial GPU provisioning following; subsequent
submissions should overlap the two stages.
DAG topology:
ensure-binary ──┐
├──> train
warmup-gpu ────┘ (only depends on ensure-binary; warmup is
fire-and-forget infrastructure)
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Per user direction "no gating, this is the new default": the stacked
Mamba2 -> CfC -> heads design is THE production architecture. There's
no competing-baseline comparison to run. Validation reduces to normal
training metrics (per-horizon val AUC, train loss curve, sanity floor
of >0.5 AUC).
Deletions:
- crates/ml-alpha/src/gate/cfc_vs_mamba2.rs (gate verdict logic)
- crates/ml-alpha/src/gate/mod.rs
- crates/ml-alpha/examples/alpha_gate.rs (gate runner binary)
Renames:
- crates/ml-alpha/src/gate/auc.rs -> crates/ml-alpha/src/eval/auc.rs
- lib.rs: pub mod gate -> pub mod eval (gate implied comparison;
eval doesn't)
Spec amendments:
- Drop the "Gate baseline strategy" amendment (committed earlier
this session)
- Reframe the stacked-architecture amendment as a "decision" not a
"gate"; production path is unambiguous
- Reframe Section 4 "Validation gate: CfC must meet Mamba2" -> just
"Validation: per-horizon val AUC" with the >0.5 sanity floor
Doc cleanups: stale "Mamba2 gate baseline" mentions in build.rs and
pinned_mem.rs replaced with neutral wording. The Argo template
comment about "downstream gate consumption" becomes "for monitoring".
Test status: all 26+ ml-alpha tests pass. AUC tests (6/6) still pass
under the eval:: namespace.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Defines the Mamba2-only baseline for the stacked-vs-baseline gate
verdict as an ablation of the SAME PerceptionTrainer (a --bypass-cfc
flag), not a separate model. Apples-to-apples; same data window,
same hyperparameters, same code path. The only difference is whether
the CfC step is in the loop.
Three ablation options evaluated:
1. --bypass-cfc flag (recommended): Mamba2 -> heads directly
2. --mamba2-state-dim 2 (crippled Mamba2, CfC stays)
3. Frozen CfC initialized to identity (no code branch needed)
Option 1 wins on clarity: it answers "is CfC additive on top of
Mamba2" unambiguously, with the same Mamba2 capacity and same
training regime in both arms.
Concrete next-session work documented (1-2 hours):
- PerceptionTrainerConfig.bypass_cfc: bool + step() branch
- alpha_train --bypass-cfc CLI flag
- alpha-perception-template.yaml workflow parameter + bash branch
- submit both runs, fetch summaries, alpha_gate, commit verdict
gate_verdict logic unchanged — the cfc/mamba2 naming in the report
becomes stacked/bypass at the binding layer; the verdict math is
generic.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Argo WorkflowTemplate at infra/k8s/argo/alpha-perception-template.yaml
runs the stacked Mamba2 -> CfC -> heads PerceptionTrainer on a single
L40S in fr-par-2. Two-stage DAG:
ensure-binary (ci-compile-cpu pool, sccache-backed cargo build of
alpha_train example, SHA-keyed binary cache under
/data/bin/$SHORT_SHA/)
train (ci-training-l40s pool, runs the cached binary against
/data/futures-baseline/mbp10 with predecoded sidecar
cache at /feature-cache/predecoded, writes
alpha_train_summary.json to
/feature-cache/alpha-perception-runs/$SHA/)
Defaults mirror the validated synthetic-overfit smoke config:
epochs=5, seq_len=32, mamba2_state_dim=16, lr_cfc=3e-3,
lr_mamba2=1e-3, n_train_seqs=8000, n_val_seqs=1000, seed=0x4242
Submission script scripts/argo-alpha-perception.sh wraps argo submit
with the standard L40S/H100 cuda-compute-cap mapping. --watch
follows logs.
Workflow nodeSelector pinned to fr-par-2 (consistent with the cluster
topology constraint). ttlStrategy 1h after completion;
activeDeadlineSeconds 4h cap (well above expected ~30-90 min wall).
This is the cluster entrypoint for the stacked perception design.
Once it lands a summary on MinIO, the gate runner (alpha_gate, Task
17) can consume it.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The merged Mamba2 -> CfC -> heads design from the 2026-05-16 spec
amendment supersedes the CfC-alone path. Removing the old CfC-only
PerceptionTrainer and renaming the stacked Mamba2CfcTrainer to
PerceptionTrainer (one trainer, clean naming).
Deletions:
- src/trainer/perception.rs (the OLD CfC-only trainer)
- tests/perception_overfit.rs (CfC-only smoke)
- tests/perception_debug_dump.rs (CfC-only trajectory print)
- tests/stacked_overfit.rs (replaced by perception_overfit pointing
at the renamed module)
Renames:
- src/trainer/stacked.rs -> src/trainer/perception.rs
- Mamba2CfcTrainer -> PerceptionTrainer
- Mamba2CfcTrainerConfig -> PerceptionTrainerConfig
- tests/stacked_overfit.rs content -> tests/perception_overfit.rs
CLI rewrite:
examples/alpha_train.rs now drives the stacked PerceptionTrainer.
Per-step inputs are sequences (Vec<Mbp10RawInput>) of length
seq_len; labels come from the LAST position of the window
(per-horizon). Flags: --seq-len, --mamba2-state-dim, --lr-cfc,
--lr-mamba2 (no more --n-hid since hidden_dim is fixed at 128 to
match Mamba2 and CfC by design).
Test status:
- 64 GPU tests pass on local sm_86 (31 lib unit + 33 integration)
- synthetic-overfit (rebranded perception_overfit): 250 steps,
initial=0.5951 -> final=0.1917 (68% drop, well above 40% gate)
- All bit-equiv / finite-diff / invariant tests still PASS
- One ignored test: the gate_artifact integration (waiting for
cluster-trained summary inputs)
The cluster gate (Task 18) now compares stacked-trained AUC vs a
Mamba2-baseline AUC (TBD: stacked vs a simpler "Mamba2 only" config
or an external reference baseline).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Realizes the 2026-05-16 spec amendment merging Mamba2 + CfC into one
stacked architecture (vs the original "compete via gate" framing).
Forward chain:
snap_features × seq_len
-> window pack [1, seq_len, FEATURE_DIM]
-> Mamba2Block.forward_train -> (logit, cache.h_enriched [1, hidden_dim])
-> cfc_step(x=h_enriched, h_old=0) -> h_new
-> heads -> probs [5]
-> BCE(probs, labels)
Backward chain:
BCE -> grad_probs
-> heads_backward -> grad_h_new + grad_W_heads, grad_b_heads
-> cfc_step_backward -> grad_W_in, grad_W_rec, grad_b + grad_x (=grad_h_enriched)
-> Mamba2.backward_from_h_enriched(&cache, &grad_h_enriched_tensor)
-> Mamba2BackwardGrads (full 9-tensor gradient set)
Optimizers (6 total):
- 5 CfC AdamWs (W_in, W_rec, b, heads_w, heads_b) — reused from
PerceptionTrainer's per-param-group pattern
- 1 Mamba2AdamW for all 9 Mamba2 parameter tensors (existing
implementation in mamba2_block.rs)
Synthetic-overfit on constant +1 direction (seq_len=16, state_dim=8,
lr_cfc=3e-3, lr_mamba2=1e-3, 250 steps):
initial_avg=0.5951 → final_avg=0.1917 (68% drop, well past 40% gate).
Monotone descent at all 5 progress checkpoints.
Architectural note (v1): CfC runs with h_old=0 each step (no inter-
step recurrence). With h_old=0, the CfC layer is effectively per-cell
tau-scaled tanh FC. Inter-step CfC state (h_old carrying between
calls) is a v2 extension once the cluster gate validates the v1
foundation.
The cluster gate (Task 18) now has the actual stacked production
trainer to deploy, not a CfC-alone-vs-Mamba2-alone bench.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Adds grad_x[k] = sum_i d_pre[i] * W_in[i,k] computed by thread 0 of
the cfc_step_backward kernel (after the existing __syncthreads in
the shared-mem sd_pre relay). Required by the stacked Mamba2 -> CfC
design: Mamba2.backward_from_h_enriched needs grad on h_enriched,
which is the CfC's "x" input in the stacked topology.
For the existing CfC-only PerceptionTrainer (x = snap_features, no
upstream learnable layer), grad_x is computed but discarded into a
preallocated buffer.
backward_finite_diff tests still pass (4/4) — the new arg is the
14th positional kernel arg; existing callers updated. perception_
overfit smoke still passes (loss 0.5669 -> 0.0665 in 200 steps).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
