5bd8897880084005097ac47b2ed4722f6dc7d4a0
5256 Commits
| Author | SHA1 | Message | Date | |
|---|---|---|---|---|
jgrusewski
|
5bd8897880 |
refactor(ml-alpha): move GRN heads from PerceptionTrainer to CfcTrunk (X9)
Adds the 4 missing GRN head fields (heads_w1, heads_b1, heads_w2, heads_b2) that X1's skeleton under-modeled. Trunk now owns all 10 GRN head tensors: input projection (HIDDEN→HEAD_MID), mid→mid layer, gate/main/skip outputs. Trainer's 10 head fields removed. Trainer init still draws heads from its ChaCha8Rng chain at the same call position, then memcpy_htod's them into the trunk's now-allocated slots. Forward, backward, and AdamW access sites redirect to self.trunk.heads_*. Gradient buffers + AdamW state stay at trainer level. Verification: - perception_forward_golden: PASS (max_diff = 0.000000) - ml-alpha lib tests: 34 pass Per spec §2.2 (X9). After this commit, the trunk owns ALL v2 inference weights — the source of truth for downstream checkpoint serialization. |
||
|
jgrusewski
|
ef29e13c74 |
refactor(ml-alpha): move CfC weights from PerceptionTrainer to CfcTrunk (X8)
Adds CfcConfig.cfc_n_in field (default HIDDEN_DIM) so the trunk's CfC
layer is sized for v2 usage (CfC input = LN_b output = HIDDEN_DIM) rather
than V1 usage (CfC input = snap features = FEATURE_DIM). The previous
CfcConfig.n_in field stays as "raw snap feature dim" for any V1 callers
still in the tree (fxt-backtest, trunk_forward.rs, graph_a_replay.rs);
their forward paths will be cleaned up in X10/X11 when the v2 forward
graph lives natively on the trunk.
Trainer's w_in_d / w_rec_d / b_d / tau_d fields are removed. Trainer
init still draws CfC weights from its ChaCha8Rng chain at the same
call position, then memcpy_htod's them into the trunk's now-v2-shaped
slots. Forward, backward, and AdamW access sites redirect to
self.trunk.{w_in,w_rec,b,tau}_d.
Verification:
- perception_forward_golden: PASS (max_diff = 0.000000)
- ml-alpha lib tests: 34 pass
Per spec §2.2 (X8).
|
||
|
jgrusewski
|
2bd774ad14 |
refactor(ml-alpha): move attention-pool weight to CfcTrunk (X7)
attn_q_d moved from PerceptionTrainer to self.trunk.attn_q_d. Verification: golden bit-exact; ml-alpha lib green. Per spec §2.2 (X7). |
||
|
jgrusewski
|
3357699431 |
refactor(ml-alpha): move LN_b weights from PerceptionTrainer to CfcTrunk (X6)
LN_b (formerly ln_gain_d / ln_bias_d on trainer) now lives at self.trunk.ln_b_gain_d / ln_b_bias_d. LN_b is the LayerNorm after Mamba2 stack 2. Verification: golden bit-exact; ml-alpha lib green. Per spec §2.2 (X6). |
||
|
jgrusewski
|
5a534e9972 |
refactor(ml-alpha): move Mamba2 stack 2 from PerceptionTrainer to CfcTrunk (X5)
Same pattern as X3: trainer constructs mamba2_l2 + Mamba2AdamW (against &mamba2_l2), then moves the block into trunk.mamba2_stack_2. All forward/backward/AdamW access sites redirect to self.trunk.mamba2_l2_mut(). Gradient buffers + AdamW state remain at trainer level. Verification: golden bit-exact; ml-alpha lib green. Per spec §2.2 (X5). |
||
|
jgrusewski
|
868021e818 |
refactor(ml-alpha): move LN_a weights from PerceptionTrainer to CfcTrunk (X4)
LN_a gain/bias tensors now live at self.trunk.ln_a_gain_d / ln_a_bias_d. Trainer-side initialization values still drive the upload, preserving PRNG-driven init values. Gradient buffers + AdamW state remain at trainer level. Verification: golden bit-exact (max_diff = 0.000000); ml-alpha lib green. Per spec 2026-05-19-ml-alpha-v2-trunk-grows-and-deployability-design.md §2.2 (X4). |
||
|
jgrusewski
|
2d849dd5e3 |
refactor(ml-alpha): move Mamba2 stack 1 from PerceptionTrainer to CfcTrunk (X3)
PerceptionTrainer no longer owns the Mamba2 stack-1 block; the trunk's mamba2_stack_1 Option is filled in after Mamba2Block::new (which still runs at the same point in PerceptionTrainer::new so the Mamba2 weight init order is unchanged). Mamba2AdamW was already constructed against &mamba2 before the move, so optimizer state is preserved. CfcTrunk gains mamba2_l1_mut() / mamba2_l1() / mamba2_l2_mut() / mamba2_l2() accessors that unwrap the Options. Forward, backward, and AdamW step sites in evaluate_batched / step_batched / evaluate redirect through the new accessors. Gradient buffers (mamba2_grads_buffers) and AdamW state (mamba2_adamw) remain at trainer level — training-only state stays with the trainer. Verification: - perception_forward_golden: PASS (max_diff = 0.000000) - ml-alpha lib tests: 34 pass Per spec 2026-05-19-ml-alpha-v2-trunk-grows-and-deployability-design.md §1.1, §2.2 (X3). EOF ) |
||
|
jgrusewski
|
fdef6efe98 |
refactor(ml-alpha): move VSN weights from PerceptionTrainer to CfcTrunk
PerceptionTrainer gains a `trunk: CfcTrunk` field constructed at the top of `new()` (after the determinism seed guard). VSN's weight tensors (`vsn_w_d`, `vsn_b_d`) now live on `self.trunk`; the trainer-owned copies are removed. Forward + backward + AdamW access sites redirect to `self.trunk.vsn_w_d` / `self.trunk.vsn_b_d`. PRNG-state preservation: the trainer's ChaCha8Rng chain still draws VSN values at the same call position as before, then memcpy_htod's them into the trunk's zero-initialised VSN slots from X1. This keeps every downstream weight (attn_q, ...) bit-identical to the pre-X2 layout — perception_forward_golden continues to pass with max_diff = 0.000000. Gradient buffers (`grad_vsn_w_d`, `grad_vsn_b_d`) and AdamW state (`opt_vsn_w`, `opt_vsn_b`) remain at trainer level — they're training-only and shouldn't move to the inference trunk. Verification: - perception_forward_golden: PASS (max_diff = 0.000000) - ml-alpha lib tests: 34 pass - alpha_train example builds clean Per spec 2026-05-19-ml-alpha-v2-trunk-grows-and-deployability-design.md §1.1, §2.2 (X2). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
e338000eec |
feat(ml-alpha): CfcTrunk v2 weight skeleton (no callers yet)
Adds zero-initialised v2 weight tensors (VSN, LN_a/b, attn_q, GRN heads) and Option<Mamba2Block> slots for stacks 1 and 2 to CfcTrunk. Extends CfcConfig with mamba2_state_dim (default 16, matches PerceptionTrainerConfig::default). Imports HEAD_MID_DIM for GRN head sizing. No forward path changes — fields allocated, not yet read. Existing new_random init for V1 fields (CfC + simple heads + projection) is preserved unchanged so the trunk's forward kernels still produce the same output. Skeleton for X2..X9 weight-group migrations. Verification: - ml-alpha lib tests: 34 pass - perception_forward_golden bit-equivalence: PASS (max_diff = 0.000000) Per spec 2026-05-19-ml-alpha-v2-trunk-grows-and-deployability-design.md §1.1, §2.2 (X1). |
||
|
jgrusewski
|
b47b2fabfb |
fix(ml-core): deterministic GPU weight init via scoped_init_seed
ml_core::cuda_autograd::init::generate_uniform (backing xavier_uniform,
kaiming_uniform, bias_uniform, near_zero_xavier) defaulted to seeding
from SystemTime::now() + thread_id, producing non-reproducible weights
across processes. Mamba2 stacks initialise via OwnedGpuLinear::xavier,
which routes through this helper — so PerceptionTrainer.evaluate output
diverged 5-30% across fresh-process runs with identical cfg.seed.
Fix: thread-local seedable RNG override. New API:
let _g = ml_core::cuda_autograd::init::scoped_init_seed(seed);
// ... all xavier/kaiming/bias/near_zero calls draw from
// StdRng::seed_from_u64(seed) chain while _g is alive ...
// _g dropped here -> restores default time-based seeding
PerceptionTrainer::new now installs the guard before any Mamba2Block
construction, so the trainer is reproducible from cfg.seed end-to-end.
CfC/VSN/heads already used explicit ChaCha8Rng::seed_from_u64 — only
Mamba2 was affected.
Production behavior unchanged when no guard is set. ml-core: 306 tests
pass, ml-alpha: 34 lib tests pass.
Regression test: crates/ml-alpha/tests/perception_forward_golden.rs
captures bit-exact PerceptionTrainer.evaluate output (loss + 160 probs
on a deterministic seed=42 fixture) into a 644-byte golden file.
Three consecutive runs now produce max_abs_diff=0; pre-fix runs varied
by 0.1-0.3 absolute on individual probs.
.gitignore: added exception for crates/ml-alpha/tests/fixtures/*.bin
so deterministic test fixtures land in repo.
Per pearl_scoped_init_seed_for_reproducibility in project memory.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
d45dde8458 |
plan(ml-alpha): trunk-grows refactor + deployability validation roadmap
20-commit atomic ladder (X0–X19) + Phase 1→2 gate + Phase 2 runtime
runbook. Implements spec
|
||
|
jgrusewski
|
da1dd92bf8 |
spec(ml-alpha): trunk-grows refactor + deployability validation (supersedes prior)
Supersedes the 2026-05-19 deployability spec (commit
|
||
|
jgrusewski
|
4938ac2ec5 |
plan(ml-alpha): v2 deployability validation — atomic-commit roadmap
Five-commit implementation plan + two-step runtime runbook for the
deployability spec committed at
|
||
|
jgrusewski
|
07d5de5048 |
spec(ml-alpha): v2 deployability — stress anchor + max-dd gate + diagnostics
User revisions to design spec from this session: 1. §2.1 — split anchor into realistic (200 ms, 1 tick) and stress (400 ms, 1.5 tick). Realistic remains the hard verdict gate; stress grades the Pass into Pass-robust vs Pass-nominal. 2. §2.2 — expand metrics from Sharpe-only to four: annualized daily Sharpe and max-drawdown are hard gates (median across windows > 1.0 and < 20% respectively); Sortino and profit factor are diagnostics. Per-window summary.json schema extended. 3. §2.6 — verdict emitter rewrites to two-anchor logic, tiered output: Pass-robust / Pass-nominal / Fail-inconclusive / Fail / Fail-degenerate. 4. §3.3 — added max-dd computation and zero-trade-window failure modes. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
0809390cd5 |
spec(ml-alpha): v2 deployability validation — falsifiable LOB-backtest verdict
Closes the wiring gap between the existing real-LOB backtest system (C1–C19 on this branch) and the v2 ml-alpha model. alpha_train.rs currently never calls save_checkpoint, so the LOB harness/sweep/aggregate machinery has never been pointed at a real trained model. Design defines a single-pass falsifiable deployability gate: produce a production checkpoint (cv-n-folds=1, cv-train-window=4 → train on 2024 quarters, val on 2025-Q1, hold out 2025-Q2..2026-Q1), pre-register one threshold on the W0 val window, then evaluate median Sharpe across 4 held-out walk-forward quarters at the realistic Scaleway→IBKR anchor (200ms RTT, 1-tick all-in cost). Pass iff median > 1.0; inconclusive in [0.8, 1.0] counts as fail; smoke gate halts the full sweep on any wiring failure. Approved through brainstorming. Awaiting user spec-review before plan handoff. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> |
||
|
jgrusewski
|
6b7920474d |
spec(ml-alpha): mark AUC-regret controller SUPERSEDED — empirically falsified
The motivating "50% saturation hit-rate" observation came from gm67g fold 0 (Option-2 config, commit |
||
|
jgrusewski
|
e004d6c217 |
Revert "arch(ml-alpha): restore count-delta redundancy at feature slot [26]"
This reverts commit
|
||
|
jgrusewski
|
1fc100ae74 |
spec(ml-alpha): AUC-regret controller — replace BCE-z-score signal with per-horizon-regret
The current λ controller boosts horizons by BCE-EMA z-score, which conflates
three distinct causes of high BCE — only one of which (under-trained
horizon) benefits from boosting. The other two (intrinsically harder
horizon, calibration drift) are unaffected by gradient-magnitude lifts.
Empirical motivation (gm67g fold-0, this branch's 3-fold run):
when λ_h6000 saturated at 2.0, h6000's next-epoch AUC went UP 2/4
times and DOWN 2/4 times. 50% hit rate ⇒ the BCE saturation signal
is misaligned with the optimization objective.
This spec replaces BCE-z-score with AUC-regret:
best_auc[h] = running max of per-horizon validation AUC
regret[h] = max(0, best_auc[h] - current_auc[h])
regret_max_ema = EMA of max_h regret[h]
λ[h] = clamp(1.0, 2.0, 1 + regret[h] / regret_max_ema)
Properties:
- Aligned with the objective (AUC, not BCE)
- Naturally bounded (regret ∈ [0, 1])
- "At personal best" → λ=1.0 (no wasted boost)
- Auto-saturation by design (max-regret horizon → ceiling)
- Cold-start clean (e0: best=current, regret=0, uniform λ)
- Zero hardcoded magic beyond bootstrap epsilons
Implementation surface ~250 LOC:
- Split horizon_lambda kernel: horizon_loss_ema (per-step) +
horizon_lambda (per-epoch, AUC-regret math)
- Trainer state: drop z_max_ema, add best_auc + regret_max_ema
- Per-epoch entry point: trainer.update_lambda_from_auc()
- Extend isv snapshot log line with best_auc_h* + regret_h*
Test plan:
- Local 9/9 perception_overfit
- New synthetic-AUC unit test (controller correctness)
- Cluster 5-epoch smoke + 3-fold A/B vs Option-2 baseline (
|
||
|
jgrusewski
|
004b662c80 |
obs(ml-alpha): per-N-step train_loss log line for liveness + intra-epoch monitor
Adds a single tracing::info!("step") emitted every 500 training steps:
if epoch_train_steps % 500 == 0 {
tracing::info!(epoch, step = epoch_train_steps, loss = loss, "step");
}
Consumed by /tmp/alpha_monitor.py (v2: per-epoch trajectories + ISV +
liveness) to render intra-epoch loss trajectory and detect stalls
faster than the once-per-epoch granularity allowed.
At 8000 steps/epoch and ~36s/epoch on L40S → ~16 step lines per epoch
per fold, well below the cluster log volume budget.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
008f65d894 |
arch(ml-alpha): restore count-delta redundancy at feature slot [26]
The C16 tick-rule swap (
|
||
|
jgrusewski
|
17eb825113 |
arch(ml-alpha): σ becomes sidecar — BCE drops Kendall damping (upgraded path)
Empirical record at this data + this architecture:
Phase 1+2+3 (no σ in loss): mean_auc 0.7749 ± 0.024
v2 (σ + axes B/C/D/E): mean_auc 0.7541 ± 0.005
σ-only (kept σ, dropped B/C/D/E): mean_auc 0.7506 ± 0.008
Perf-fix (σ-only math): mean_auc 0.7499 ± 0.010
ISV-σ (closed-form σ + adaptive λ): mean_auc ~0.75 (2/3 folds)
Every architecture with σ-in-loss lands at 0.75. Removing σ is the
only thing that hits 0.77. That is a framework mismatch, not a tuning
problem.
Kendall+Gal+Cipolla 2018 frames σ as TASK NOISE level — damp the noisy
task, trust the clean one. Our horizons don't have different label
noise; they have different intrinsic difficulty (longer horizon = more
price-walk uncertainty = lower achievable AUC). σ-Kendall sees "high
BCE on h6000" and interprets it as "h6000 is unreliable, back off" —
precisely the opposite of what we want. h6000 is the deployment
target; damping it is a self-inflicted wound. With mean_bce ~ 0.7,
σ ≈ √0.7 ≈ 0.84 ⇒ w_h ≈ 0.71, uniformly attenuating gradient by
~30% across every horizon. λ's z-score boost (max 2×) can rebalance
relative-per-horizon but cannot recover the absolute magnitude.
Per-horizon prioritization remains via the ISV-driven λ controller
(grad scaler in heads_grn_bwd, per pearl_adam_normalizes_loss_weights).
That controller IS appropriate for our problem: it boosts hard
horizons rather than damping them, and it operates on the gradient
into the trunk rather than on the loss aggregate (Adam-cancellation
safe).
Changes to bce_loss_multi_horizon.cu (six lines):
w_h = bw (was: 0.5 * bw * exp(-2 * log_sigma_h))
d_log_sigma_h[h] = 0.0 (was: 1 - 2 * w_h * mean_bce)
total_loss += bw * mean_bce (was: + w_h * mean_bce + log_sigma_h)
σ infrastructure preserved unchanged:
- horizon_ema_and_lambda still computes log_sigma_h closed-form from
loss_ema (Kendall equilibrium) for telemetry / future label-noise
estimation use cases
- log_sigma_h kernel arg still in BCE signature (zero churn at
callsite); ignored inside
All 9 perception_overfit tests pass — including
horizon_ema_and_lambda_track_after_training which validates the
per-horizon controller end-to-end through 64 K-loop iterations of
capture/replay.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
410ab6b0ea |
arch(ml-alpha): ISV-driven σ + adaptive Z_SCALE — both controllers anchor on loss_ema
Per pearl_controller_anchors_isv_driven, every controller anchor/target/
cap derives from a tracked signal, not hardcoded constants. The σ-only
revert kept Kendall σ as a free Adam-learned scalar — that violated ISV
discipline and fought Adam's m/√v normalization
(pearl_adam_normalizes_loss_weights).
Single source of truth for both per-horizon controllers:
log_sigma_h[h] ← max(log(0.5), 0.5 * log(loss_ema[h]))
Kendall equilibrium (∂L/∂log σ = 0 ⟹ σ_h² = mean_bce_h)
in closed form. Asymmetric floor at log(0.5) prevents
collapse. No Adam state, no gradient delay.
lambda[h] ← clamp(1.0, 2.0, 1.0 + Z_SCALE_ISV * z_h)
Z_SCALE_ISV = (LAMBDA_CEILING - LAMBDA_FLOOR) / z_max_ema
Adaptive scale auto-uses the full clamp envelope:
the historical-max-z horizon maps exactly to
LAMBDA_CEILING. Replaces hardcoded Z_SCALE=0.5 which
rarely engaged on real data (max observed λ ~1.04).
Both anchor on the same ISV (loss_ema). z_max_ema is a new single-scalar
EMA state tracking max |z| across horizons, with first-obs bootstrap.
Removes:
- opt_log_sigma AdamW optimizer (σ no longer learned)
- grad_log_sigma_h_d memset (BCE kernel writes; output ignored — kept
only to preserve BCE kernel signature)
Kernel signature change (horizon_ema_and_lambda):
+z_max_ema [1] (read+write EMA state)
+log_sigma_h [5] (closed-form output, overwrite)
Discipline:
- First-obs bootstrap (sentinel <= 0) per pearl_first_observation_bootstrap
- Permanent floor (max(real, floor)) per pearl_blend_formulas_must_have_permanent_floor
- Asymmetric clamp per pearl_audit_unboundedness_for_implicit_asymmetry
- Z-score normalisation per pearl_zscore_normalization_for_magnitude_asymmetric_signals
- No nvrtc, no atomicAdd, no host branches in graph capture
All 9 perception_overfit tests pass — including
horizon_ema_and_lambda_track_after_training which validates the kernel
end-to-end through 64 K-loop iterations of capture/replay.
Submit local smoke; cluster A/B vs σ-only baseline (0.7506/0.7519) and
vs Phase 1+2+3 (0.7749/0.7591) follows once the perf-only 3-fold A/B
confirms no regression at
|
||
|
jgrusewski
|
b23f8f2efa |
perf(ml-alpha): NVIDIA-grade rewrite of CfC K-loop hot kernels — 2.15× faster
Local L40S profile (perception_overfit smoke) GPU kernel time:
1589ms → 739ms (53.5% reduction, 2.15× speedup).
Wall-clock smoke: 9.6s → 4.94s (1.94× faster).
Per-kernel deltas (nsys --cuda-graph-trace=node):
reduce_axis0: 362ms → 10ms (36× faster)
Block layout: per-column (1 block / output) → 32-wide column tile
(block_dim = 32 × 8). Cross-thread reads were strided by n_tail
(~40K floats = 160KB stride) — one cache line per thread, 8× HBM
bandwidth wasted. New tile gives coalesced 128B transactions per
warp. Block tree-reduce kept (no atomicAdd, per feedback_no_atomicadd).
+1 shared-mem pad to eliminate 32-way bank conflict on the ty reduce.
multi_horizon_heads_grn_bwd_batched: 540ms → 113ms (4.8× faster)
1. Stage h_row[HIDDEN] and a1[5,HEAD_MID] in shared at block entry.
Eliminates ~28K redundant DRAM reads/block across Pass 3 + Pass 5.
2. Pass 5 reorder: k outer / i inner with d_z1[k,m] pinned in
register; writes to grad_w1_scratch are sequential per-thread.
3. Block size 64 → 128 threads. Pass 5/6 now partition over i
(output column): cross-thread writes become COALESCED 128B/warp
(was stride-128 = 512B). Passes 2/3/4 gate on (tid < HEAD_MID).
4. Pass 3 thread role: m_out → m_in/n. Same coalescing fix on
grad_w2 writes AND w2 reads in the d_eta_2 sum.
cfc_step_backward_batched: 351ms → 271ms (1.3× faster)
1. Stage x_b[n_in] and h_old_b[n_hid] in shared (was 128× redundant
DRAM reads per block; now 1× cooperative load).
2. Pass 1 thread role: i (output row) → k (output col). For each
i loop iteration, the warp writes grad_w_in[..., tid] /
grad_w_rec[..., tid] — COALESCED 128B/warp (was stride-128
non-coalesced).
multi_horizon_heads_grn_fwd_batched: 211ms → 204ms
Stage h_row[HIDDEN] in shared — Pass 1 and Pass 3 both consume.
cfc_step_batched (fwd): 95ms → 94ms
Stage x_b and h_old_b in shared.
Shared-mem budgets fit comfortably under the 48KB SM cap (~6KB / ~2KB
respectively). All 9 perception_overfit tests pass — gradient
correctness validated end-to-end (constant-signal overfit, K-loop
capture/replay, stride-4 path, evaluate-only paths).
Discipline:
- Block tree-reduce only, never atomicAdd
- No nvrtc; pre-compiled cubins via build.rs
- Mapped-pinned-only is unaffected (CPU↔GPU contract untouched)
- Single source of truth: replaced kernels in place, no v2 suffixes
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
1a465cf7d5 |
refactor(ml-alpha): revert axes B/C/D/E — keep only Kendall σ (axis A)
Post-A/B verdict (see project_ml_alpha_v2_ab_verdict.md): v2 with all
5 axes was marginally tied on h6000 (+0.0013 vs 0.7591 baseline mean,
fails the +0.01 win threshold) and slightly below on mean_auc
(−0.0208 vs 0.7749 baseline mean, within 1σ) at ~5× the wall-time
cost. Per `feedback_v7_gem_methodology` (measure before delete or
wire), the architecture has been measured — it doesn't earn its
compute cost. This commit reverts the axes that didn't lift:
- axis B (L2 anchor + Wiener-α controller) — DROPPED
- axis C (horizon-token attention pool) — DROPPED
- axis D (regime-MoE gate + experts) — DROPPED
- axis E (inverted cross-variate attn) — DROPPED
- axis A (Kendall σ-weighted BCE) — KEPT
Files deleted (kernels, host bindings, numgrad tests, trainer state):
- cuda/{horizon_token_attention_pool, inverted_attention_pool,
inv_pooled_merge, regime_moe_gate, anchor_l2,
horizon_mean_collapse}.cu
- src/{horizon_token_attention_pool, inverted_attention_pool,
inv_pooled_merge, regime_moe_gate, anchor_l2,
horizon_mean_collapse}.rs
- src/trainer/{multi_horizon_attention, anchor_controller}.rs
- tests/{horizon_token_attention_pool_numgrad,
inverted_attention_pool_numgrad,
regime_moe_gate_numgrad,
anchor_l2_numgrad}.rs
Files restored (from V1 commit
|
||
|
jgrusewski
|
4ae9a27f48 |
fix(ml-alpha): wire axis E into loss — real add_inv_broadcast kernel
CRITICAL ARCHITECTURAL FIX discovered while planning fused kernel:
The previous Stage 2 `add_inv_broadcast` helper in
`multi_horizon_attention.rs` was a STUB that returned Ok(()) without
doing anything. Practical consequences:
- Forward: `inv_pooled_d` (output of inverted_attention_pool.forward)
was never added into `ctx_h_d`. Downstream MoE + heads never saw
the inverted-attention signal. Axis E contributed ZERO to the
forward output and the loss.
- Backward: `inv_pool.backward` was being fed `grad_ctx_mean` as
its "upstream gradient", but that's the gradient at the CHAIN
TERMINUS — not the gradient w.r.t. inv_pooled_d (which is zero
by construction since inv_pooled wasn't in the loss). The bwd
was injecting incorrect noise into `grad_ln_out`.
Net: paying inverted_attention compute for no gain, plus polluting
ln_b's gradient. Two perf rewrite rounds earlier today showed no
wall-time movement precisely because the slow path was wired into
training while the optimized one was dead.
FIX:
(a) New kernel `cuda/inv_pooled_merge.cu`:
fwd: ctx_h[b, h, d] += inv_pooled[b, d] (broadcast over h)
bwd: grad_inv_pooled[b, d] = Σ_h grad_ctx_h[b, h, d]
Tiny — single block-per-batch, no syncthreads, coalesced reads.
(b) Host binding `src/inv_pooled_merge.rs` (InvPooledMerge).
(c) `MultiHorizonAttention` adds:
- `merge: InvPooledMerge` field.
- `grad_inv_pooled_d [B, H]` buffer for the real upstream of inv_pool.bwd.
(d) `MultiHorizonAttention.forward` now calls `merge.forward(...)`
between inv_pool.forward and moe.forward. Axis E is now actually
in the model's forward output.
(e) `MultiHorizonAttention.backward` now calls `merge.backward` after
moe.bwd writes `grad_ctx_h_d`, producing `grad_inv_pooled_d`.
`inv_pool.backward` consumes the REAL upstream gradient
(`grad_inv_pooled_d`) instead of the prior `grad_ctx_mean` fake.
(f) Old stub `add_inv_broadcast` deleted.
CORRECTNESS:
- perception_overfit 9/9 PASS after the wiring. Loss still shrinks
on the constant-signal test (0.67 → -0.99 over 250 steps), now
with axis E actually contributing.
- All numgrad kernel tests still pass (kernels themselves unchanged
in this commit; only the wiring).
PERF IMPACT (expected):
- +2 tiny launches per step (merge fwd + bwd). Negligible.
- The inverted-attention compute that was previously dead now
actually feeds the loss → same wall-time, but it's earning the
cost. This unblocks meaningful axis-E perf measurement on next
smoke.
NEXT: re-run cluster smoke to confirm wall-time + verify axis E
gradient flow is healthy. After that, consider full MHA forward
fusion as a follow-up.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
11b964359b |
perf(ml-alpha): MoE bwd compact scratch + scatter kernel + inv-attn loop interchange
Two perf optimizations bundled:
(1) MoE backward scratch compaction
Old: `grad_w_scratch_d [B, N_H, N_E, H, H]` = 1.3 MB per step (B=1)
New: `grad_w_scratch_d [B, N_H, H, H]` = 320 KB per step
4× memory reduction. Since each batch's top-1 router selects ONE
expert, only that expert's slot was ever non-zero in the prior
layout — the N_E axis was entirely wasteful.
The shape change required:
- Updated `regime_moe_gate_bwd` to write the compact layout.
- New `regime_moe_gate_scatter` kernel scatters per-(b, h)
rank-1 contributions into `grad_experts_W[e]` / `grad_experts_b[e]`
based on `top_e[b]`. Grid (N_E, H, ceil(H/32)) × block (32) —
one warp per (e, d_out, d_in_chunk). 65536 → 16384 grid cells
(4× fewer blocks dispatched).
- Dropped the previously-naive 65536-block `reduce_axis0` for
`grad_experts_w` from `perception.rs` (the scatter kernel
produces the final per-expert grad directly).
- `tests/regime_moe_gate_numgrad.rs` reads `grad_experts_w` from
the scatter output instead of host-side reducing the 5D scratch.
(2) inverted_attention bwd loop interchange
Phase 2's tight loop:
for k:
for j:
ds_myh_j = d_scores[my_h, j] // doesn't depend on k!
ds_j_myh = d_scores[j, my_h] // doesn't depend on k!
...
Hoisted d_scores reads out of the K-loop into J-outer with
per-thread `q_arr[K_MAX]` / `k_arr[K_MAX]` register accumulators.
Net: 32× fewer DRAM reads of d_scores per thread per bwd.
CORRECTNESS:
- regime_moe_gate numgrad PASSES (1/1, 11 numgrad checks).
- inverted_attention numgrad PASSES (1/1, 6 numgrad checks).
- perception_overfit 9/9 PASS — including loss-shrinks tests.
NEXT: re-run cluster smoke to measure the new wall-time vs the 17 s
baseline. Prior smoke at
|
||
|
jgrusewski
|
a263cd5446 |
perf(ml-alpha): inverted_attention_pool — cache mean_k, pre-compute pooled
The cluster smoke at
|
||
|
jgrusewski
|
9170d24fe3 |
refactor(ml-alpha): replace legacy attention_pool with MultiHorizonAttention [Stage 2]
Single source of truth for the attention path. Deletes the legacy
single-Q `attention_pool.cu` and all `attn_*` fields from
`PerceptionTrainer`; wires `MultiHorizonAttention` (the bundle
introduced in Stage 1) into `step_batched` + `evaluate_batched` as
THE attention summary that seeds CfC's `h_old` at k=0.
Deletions:
cuda/attention_pool.cu (244 lines)
perception.rs::attn_q_d/attn_context_d/
attn_weights_d/grad_attn_q_d/opt_attn_q/
attn_fwd_fn/attn_bwd_fn/_attn_module/
attn_grad_q_scratch_d (all struct fields)
perception.rs::ATTENTION_POOL_CUBIN (include_bytes constant)
Their corresponding init + struct-construction lines.
build.rs::KERNELS (drops "attention_pool")
New kernel + binding:
cuda/horizon_mean_collapse.cu (53 lines)
- `horizon_mean_collapse_fwd/_bwd`: collapses [B, N_H, H] → [B, H]
by averaging over the horizon axis. Single-pass, no reductions.
src/horizon_mean_collapse.rs (host binding)
MHA additions:
- `collapse` field + `ctx_mean_d` + `grad_ctx_mean_d` for the seed.
- `grad_ctx_h_d` scratch (split from grad_horizon_tokens_scratch to
avoid aliasing when MoE bwd writes d_ctx_h while pool bwd writes
d_horizon_tokens).
- `forward(ln_b_out)`: horizon-token pool → inverted pool → MoE
dispatch → mean-collapse → ctx_mean_d.
- `backward(ln_b_out, grad_ctx_mean, grad_ln_out)`: full reverse
chain.
- `apply_anchor()`: launches anchor_l2 on horizon_tokens, Q,
experts_w.
- `adamw_step()`: steps all 6 owned optimizer groups.
PerceptionTrainer integration:
- Section 2d (forward): `self.mha.forward(&self.ln_out_d)` replaces
the legacy attention_pool launch. CfC's h_old at k=0 now reads
`self.mha.ctx_mean_d.device_ptr` (was `self.attn_context_d`).
- Section 7c-pre (backward): `self.mha.backward(ln_out, grad_h_carry,
grad_h_enriched_seq)` replaces the legacy attn_bwd_fn launch.
- Four `reduce_axis0` launches collapse MHA's per-batch scratches
into shared gradient buffers: grad_horizon_tokens, grad_q,
grad_experts_w, grad_experts_b.
- `self.mha.apply_anchor()` adds L2 anchor grad contributions.
- Section 9 (AdamW): `self.mha.adamw_step()` replaces opt_attn_q.
- `evaluate_batched`: `self.mha.forward` replaces the legacy fwd
launch; h_old at k=0 reads `mha.ctx_mean_d`.
- `self.mha.zero_grads()` at step start (capture-safe memset_zeros).
BUG CAUGHT DURING WIRING (NVIDIA-grade discipline): first wiring
attempt mis-sized the reduce_axis0 launches for the MoE
`grad_w_scratch_d` ([B, N_H, N_E, H, H]). Initial `n_tail = N_H * N_E
* H * H = 327680` would have made reduce_axis0 read 5× past the end
of the buffer → CUDA_ERROR_ILLEGAL_ADDRESS. Fix: `n_tail = N_E * H *
H = 65536` with `n_batch = B * N_H`, treating the leading two axes
together as the reduction dimension. Caught by stacked_trainer test
on RTX 3050; would have caused silent corruption then a hard fault
on L40S/H100 later.
LOCAL VERIFICATION (RTX 3050 sm_86):
- ml-alpha builds clean (cuda feature).
- All 38+ tests PASS serially with --test-threads=1:
perception_overfit (8 tests incl. loss-shrinks)
trunk_forward (5)
stacked_loss_shrinks (multiple)
bce_grad_finite_diff (4)
snap_feature_assemble (9)
... (full suite green)
- Numgrad parity for the 4 new MHA kernels (horizon_token, inv_attn,
regime_moe_gate, anchor_l2) PASSES at 5e-2 rel.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
6a3f45d872 |
refactor(ml-alpha): consolidate BCE — Kendall σ is THE BCE, wire MHA into trainer
Single source of truth for the multi-horizon BCE per the new
feedback_single_source_of_truth_no_duplicates pearl. Eliminates the
`bce_loss_multi_horizon_sigma.cu` / `loss_sigma.rs` duplication
introduced earlier today and folds the Kendall σ-weighting into the
canonical kernel.
Deletions:
cuda/bce_loss_multi_horizon_sigma.cu (folded into legacy)
src/trainer/loss_sigma.rs (helper merged)
tests/bce_sigma_numgrad.rs (subsumed)
Rewrites:
cuda/bce_loss_multi_horizon.cu — replaced with σ-aware
implementation; kernel function name kept as
`bce_multi_horizon_forward_backward` so cubin symbol stays stable.
NVIDIA-grade warp-shuffle reduce (4 warps, 1 cross-warp barrier);
new args `log_sigma_h` (per-horizon Kendall σ logarithm) and
`d_log_sigma_h` (its gradient).
src/trainer/loss.rs — standalone helper updated to new 11-arg
kernel signature. Exposes optional `log_sigma_h` in `BceInput`
(None → zeros / passthrough Kendall init) and returns
`mean_bce_per_h` + `d_log_sigma_h` in `BceOutput`.
tests/bce_grad_finite_diff.rs — adds `log_sigma_h: None` to the
test inputs; all 4 numgrad tests PASS unchanged.
build.rs — drops `bce_loss_multi_horizon_sigma` entry from
KERNELS. The single canonical `bce_loss_multi_horizon` cubin
now contains the σ-aware kernel.
Wiring:
PerceptionTrainer gains a single `pub mha: MultiHorizonAttention`
field. Owns `log_sigma_h_d` + `grad_log_sigma_h_d` (and the rest
of the multi-horizon attention path, anchored on Stage 2 to fully
replace the legacy `attention_pool` path).
step_batched + evaluate_batched BCE callsites now thread
`mha.log_sigma_h_d` and `mha.grad_log_sigma_h_d` through the
11-arg kernel signature.
This commit keeps the legacy `attention_pool` callsite in place; the
Stage 2 commit will replace it with `mha.pool` + `mha.inv_pool` +
`mha.moe` and delete the `attn_*` fields entirely.
Verified locally: ml-alpha builds clean with the cuda feature,
bce_grad_finite_diff (4/4) PASS on RTX 3050.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
996e61b6ef |
rename(ml-alpha): MultiHorizonAttention (no version suffix in identifiers)
Removes the version suffix from the trainer-state bundle introduced
in
|
||
|
jgrusewski
|
5aad1eb846 |
feat(ml-alpha): PerceptionV2State bundle (axes A+B+C+D+E) [V9]
Bundles all v2 components into one cohesive trainer field that
PerceptionTrainer wires through in V10. Owns:
(C) horizon_tokens + Q for the horizon-token attention pool
+ per-batch grad scratches + 2 AdamWs
(E) inverted attention pool + saved attn + d_scores scratch
(D) MoE gate weights + experts (W, b) + per-batch sparse-by-expert
grad scratches + 3 AdamWs + top_e / gate_probs / aux_loss
(A) log_sigma_h per-horizon Kendall scalar + AdamW (0.25× LR)
(B) AnchorController (Wiener-α host-side) + anchor_l2 kernel +
init snapshots of horizon_tokens, Q, experts_w, experts_b
for the L2 anchor
Init scale: 1/√HIDDEN_DIM Xavier for horizon_tokens/Q/experts_W;
zeros for experts_b and log_sigma_h. Anchor controller bootstrap
uses ‖horizon_tokens_init‖₂ + ‖Q_init‖₂ + their total numel to
derive the signal-driven floor (no tuned constants).
zero_grads uses memset_zeros only — capture-safe.
Compile-clean. V10 wires this state into perception.rs::step_batched
forward / backward / AdamW.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
55aeddaebd |
feat(ml-alpha): anchor_l2 kernel + Wiener-α controller (v2 B) [V8]
L2 anchor regularization toward initialization (axis B). Anchors
horizon_tokens + Q + MoE experts toward their init values to prevent
the calibration drift observed in v1 (where val_loss climbed as α
opened past epoch 1 in 2 of 3 folds).
KERNEL (`anchor_l2_fwd_bwd`):
loss_out = λ · Σ_i (p[i] − p_init[i])²
grad_p[i] += 2λ · (p[i] − p_init[i])
- Warp-shuffle reduce; one block per parameter group; strided thread
loop over n. Cross-warp reduce uses one __syncthreads.
- Coalesced grad write via stride loop.
- λ passed as device-side [1]-buffer (host writes scalar before launch
— capture-safe).
CONTROLLER (`trainer::anchor_controller::AnchorController`):
- Signal-driven λ floor: λ_floor = ‖p_init‖ / (100 · √numel).
Cross-fold-persistent per pearl_kelly_cap_signal_driven_floors.
- Wiener-α smoother (α = diff_var / (diff_var + sample_var + ε))
on val_loss change; α floored at 0.4 per
pearl_wiener_alpha_floor_for_nonstationary.
- λ blends toward target = |ema_change|·scale with α; floored at
λ_floor per pearl_blend_formulas_must_have_permanent_floor.
- First-observation bootstrap (sentinel state replaced directly on
first step) per pearl_first_observation_bootstrap.
- 4/4 unit tests PASS: signal-floor init, bootstrap returns floor,
floor protection across 1000 steps, λ_max cap.
NUMGRAD VERIFICATION (RTX 3050 sm_86):
anchor_l2_numgrad PASSES with closed-form parity (machine precision)
and central-difference parity (4 random positions) within 5e-2 rel.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
11b96dac6a |
feat(ml-alpha): regime_moe_gate kernel + numgrad (v2 D) [V6]
Top-1 Mixture-of-Experts gate per Switch Transformer. Three kernels
in one .cu file:
- regime_moe_gate_fwd: select top-1 expert from gate_logits, apply
its [H, H] linear to fused_ctx → routed_ctx [B, N_H, H].
- regime_moe_gate_bwd: chain-rule grads through the SELECTED
expert's W and bias (sparse-by-expert scratches), plus
d_fused_ctx accumulator. Inactive experts get zero contribution.
- regime_moe_gate_aux: softmax of gate_logits + load-balancing
auxiliary loss (frac · prob_mean × N_EXPERTS).
ARCHITECTURE:
- N_EXPERTS = 4. Each expert is a [H, H] linear with bias.
- Total expert params: 4 · 128 · 128 + 4 · 128 = 66 KB. Cheap.
- STE on gate: gate logit grad is zero from the expert path (top-1
is non-differentiable); the load-balance aux loss provides the
differentiable signal that pushes routing toward balanced usage.
PERFORMANCE:
- Grid (B, N_H, 1), block (HIDDEN_DIM). One block per (b, h).
- Forward: each thread computes one output channel via a dot
product over HIDDEN_DIM input dims (#pragma unroll 8).
- Backward d_fused_ctx: each thread accumulates over d_out
sequentially (HIDDEN_DIM iterations) since the weight matrix
column is naturally aligned to the thread's d_in index.
- Backward d_W/d_b scratches are sparse-by-expert; downstream
reduce_axis0 collapses over (B, N_H).
- Top-1 chosen by thread 0 per block, broadcast via shared mem.
NUMGRAD VERIFICATION (RTX 3050 sm_86):
forward_matches_host_reference_and_backward_matches_numgrad
PASSES 11 checks (4 on d_W, 3 on d_b, 4 on d_fused_ctx) within
5e-2 rel / 5e-3 abs.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
d94696620d |
feat(ml-alpha): inverted_attention_pool kernel + numgrad (v2 E) [V3]
iTransformer-style cross-variate attention pass: each of HIDDEN_DIM
features becomes a "variate token" with its K-trajectory as embedding.
FORWARD:
X_inv[h, k] = ln_out[k, h] # transpose
scores[h, j] = inv_scale · Σ_k X_inv[h, k] · X_inv[j, k]
attn[h, j] = softmax_j(scores)
pooled[h] = Σ_j attn[h, j] · mean_k_X_inv[j] # mean-K commutes out
BACKWARD: three independent chains into ln_out:
- value path: (1/K) · attn[h', my_h] · d_pooled[h'] (per-k constant)
- query path: inv_scale · Σ_j d_scores[my_h, j] · X_inv[j, k]
- key path: inv_scale · Σ_h' d_scores[h', my_h] · X_inv[h', k]
Softmax bwd: d_scores[h, j] = attn · (d_attn - Σ_l attn · d_attn)
IMPLEMENTATION NOTES:
- First attempt cached attn [H, H] = 64 KB in shared mem → tripped
the 48 KB dynamic-shared limit on sm_86 (CUDA_ERROR_INVALID_VALUE).
- Fixed by moving d_scores to a DRAM scratch buffer; shared mem
holds only X_inv [H, K] (≤ 16 KB at K = 32). One block-wide barrier
between the d_scores write and the value/query/key accumulation.
- All per-batch slice writes; no atomicAdd, no cross-block race.
- Pooled computation uses the mean-K commute (Σ_k attn · X_inv =
attn · mean_k_X_inv), saving an entire H×K accumulation pass.
LOCAL VERIFICATION (RTX 3050 sm_86):
forward_then_backward_matches_central_difference PASSES 6 numgrad
checks on ln_out positions within 5e-2 rel / 5e-3 abs.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
e8f6c4721f |
feat(ml-alpha): horizon_token_attention_pool kernel + numgrad (v2 C) [V2]
Replaces the falsified per-horizon Q_h pool with a single shared
query Q over an extended key sequence [horizon_tokens; LN_b_out],
producing per-horizon outputs via TFT-style horizon-token mixing.
FORWARD:
scores[i] = Σ_d Q[d] · ext[i, d] (i ∈ [0, N_H + K))
attn = softmax_i(scores)
S[d] = Σ_k attn[N_H + k] · ln_out[k, d] (shared time agg)
ctx[h, d] = attn[h] · horizon_tokens[h, d] + S[d] (per-horizon out)
BACKWARD: full chain rule with softmax-centring; gradients to
horizon_tokens, Q, and ln_out via the saved attn weights.
NVIDIA-grade implementation per feedback_nvidia_grade_perf_for_kernels:
- Warp-shuffle reduce (block_reduce_sum / block_reduce_max helpers)
for all per-d dot products and softmax aggregates.
- Cross-warp reduce uses exactly one __syncthreads.
- Non-divergent shuffles: inactive lanes contribute 0 via ternary.
- Block-per-batch + horizon-loop inside block → grad_ln_out += is
race-free without atomicAdd.
- Smem layout computed at launch: [s_attn(N_H+K); s_warp(N_WARPS);
s_d_S(H) on bwd]. No over-allocation.
LOCAL VERIFICATION (RTX 3050 sm_86):
forward_then_backward_matches_central_difference PASSES 12 numgrad
checks (4 each on horizon_tokens / Q / ln_out) at 5e-2 rel / 5e-3
abs envelope. First-try pass.
NOTE: .gitignore adjusted with narrow allow-rules for crates/ml-alpha/{
cuda,src,tests}/horizon_token_* paths — the broad "*token*" rule
intended for auth tokens was hiding these source files. Explicit
allow keeps the security rule intact while exempting these specific
files.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
679ab3f5eb |
feat(ml-alpha): Kendall sigma-weighted BCE kernel + numgrad (v2 A) [V7]
New `bce_multi_horizon_sigma_forward_backward` kernel implementing the
Kendall homoscedastic uncertainty weighting per spec axis A:
raw_bce_h = Σ_{i in h, m_i=1} L_i
count_h = #{i in h : m_i = 1}
mean_bce_h = raw_bce_h / count_h
w_h = base_weight_h / (2 · exp(2 · log_sigma_h))
total_loss = Σ_h [ w_h · mean_bce_h + log_sigma_h ]
d L / d p_i = m_i · (w_h / count_h) · (p − y) / (p (1 − p))
d L / d log_sigma_h = 1 − 2 · w_h · mean_bce_h
NVIDIA-grade implementation per feedback_nvidia_grade_perf_for_kernels:
- Warp-shuffle reduction (`__shfl_xor_sync`) for both per-horizon
sums and the global valid count, replacing block tree-reduce.
- One `__syncthreads` for the cross-warp aggregate; no inner-loop
barriers.
- Non-divergent shuffles: inactive lanes contribute 0 via ternary,
never via `if (tid < N) shuffle`.
- Coalesced strided access in both forward and gradient passes.
- Pre-compiled cubin via build.rs; no nvrtc.
Independent of the legacy `bce_loss_multi_horizon` kernel — that one
stays untouched so eval/smoke paths are unaffected. The v2 trainer
wires this kernel in via commit V10.
Standalone helper `bce_sigma_loss_and_grad_gpu` in `trainer::loss_sigma`
for numgrad parity tests. Three numgrad tests all PASS on RTX 3050
(sm_86) within 5e-2 rel / 5e-3 abs:
- d_log_sigma_h ↔ central-difference (numgrad on log_sigma)
- grad_probs ↔ central-difference (8 random positions)
- total_loss ↔ closed-form reconstructed from mean_bce_per_h
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
41292303dc |
refactor(ml-alpha): remove per-horizon Q_h path (C21-C25 falsified) [V1]
3-fold A/B sweep 2026-05-18 at commit
|
||
|
jgrusewski
|
4425a77844 |
plan(ml-alpha): v2 multi-horizon implementation plan (V1-V13)
Concrete TDD-driven commit map for the v2 spec (2026-05-18-ml-alpha-v2-multi-horizon-design.md). Thirteen commits ordered by dependency: delete falsified path, build new kernels with numgrad parity (V2-V8), trainer state bundle (V9), wire into PerceptionTrainer (V10), local smoke (V11), cluster smoke (V12), 3-fold A/B (V13). Per-commit verification gates and explicit stop conditions. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
1c2ee1d7c3 |
spec(ml-alpha): v2 multi-horizon redesign (A+B+C+D+E integrated)
Integrated design spec for the post-A/B redesign: Kendall sigma BCE (A), L2 anchor on horizon tokens + shared Q (B), horizon-token K-prepend replacing per-horizon Q_h (C), regime-aware MoE gate (D), and inverted-axis attention pass (E). Bundled per pearl_no_deferrals_for_complementary_fixes — all five axes have orthogonal architectural scope and independent kill criteria. Spec drops the C21-C25 per-horizon Q_h path (falsified by sweep 2026-05-18: mean_auc -0.019 vs single-Q baseline) and migrates the existing init buffers into the new horizon-tokens prefix. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
286ea26e2a |
perf(ml-alpha): warp-shuffle reduce in per-horizon kernels
Cluster A/B sweep with C25 wiring showed 86 s/epoch vs 17 s/epoch baseline = ~5x regression. Root cause: per-horizon attention pool + residual head used block tree-reduce with 8 __syncthreads per K-step in a serialised K-loop, repeated H=5 times in both fwd and bwd = ~3200 barriers/step. Plus the prob_blend bwd reduce kernel ran with a single thread per block, fully serialising over K*B. Replacements: - per_horizon_attention_pool fwd/bwd: introduce block_reduce_sum / block_reduce_max helpers using intra-warp __shfl_xor_sync + cross-warp shuffle (1 syncthread per K instead of 8). Smem shrinks to [K + N_WARPS] / [2K + N_WARPS]. - per_horizon_residual_head fwd: same warp-shuffle reduce pattern. - per_horizon_prob_blend_reduce_alpha_residual: 1 thread → 1 warp per horizon, lane-strided reduction over K*B via shfl_xor_sync. Launch config updated to block_dim=(32,1,1). Tricky bug found while implementing: the cross-warp reduce in the residual head originally guarded `__shfl_xor_sync(0xffffffff, ...)` with `if (tid < PHR_N_WARPS)`, leaving 28 of 32 lanes in warp 0 outside the call. Mask 0xffffffff requires all 32 lanes to participate — divergence is UB and hung the full-pipeline smoke on Ampere/Ada. Fix: read s_warp via ternary into all 32 lanes, then shuffle inside `if (tid < 32)`. Matches the pattern used in block_reduce_sum. Verified locally on RTX 3050 (sm_86): per_horizon_attention_pool numgrad, per_horizon_residual_head numgrad, and per_horizon_full_pipeline_smoke (zero-init identity + non-zero end-to-end) all PASS. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
83546b5c37 |
fix(ml-alpha): drop synchronize() in per-horizon pool + head bindings
After
|
||
|
jgrusewski
|
adaf275af3 |
fix(ml-alpha): C25 per-horizon path graph-capture safety
Four code paths in PerHorizonTrainState violated CUDA Graph capture invariants and tripped CUDA_ERROR_STREAM_CAPTURE_UNSUPPORTED on smoke alpha-perception-qk2p9: 1. stream.synchronize() inside forward_with_blend (illegal in capture) 2. stream.synchronize() inside backward_through_blend (idem) 3. reduce_per_batch_scratches_to_shared used host-side vec allocations + memcpy_dtoh + CPU summation + memcpy_htod (forbidden during capture per pearl_no_host_branches_in_captured_graph) 4. zero_grads allocated host zero vectors + memcpy_htod each step (idem) Fix: - Remove both synchronizes; same-stream issue order is sufficient. - Replace host-side reduction with three reduce_axis0 GPU kernel launches (q_h, w_res, bias_res). PerHorizonTrainState now owns its own reduce_axis0 cubin handle. - Replace host-zero memcpy with stream.memset_zeros for all nine gradient buffers plus d_alpha_reduced. Verified locally: per_horizon_full_pipeline_smoke (2/2) and both numgrad parity tests (attention pool + residual head) pass on RTX 3050. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
f50b466f77 |
feat(ml-alpha): PerceptionTrainer wires C21+C22 fwd+bwd+AdamW (C25)
The per-horizon attention pool from C21+C22 is now fully integrated
into step_batched's hot loop. Forward and backward both flow; AdamW
updates all four new parameter groups (Q_h, w_res, bias_res, α) every
training step. At α=0 init the contribution is byte-identical to
baseline; training discovers whether α should grow.
New kernel: cuda/per_horizon_prob_blend.cu
per_horizon_prob_blend_fwd
Reads logit_per_k_d (already stored by GRN forward) +
sigmoid(logit_baseline + tanh(α[h]) * residual[b, h]) →
overwrites probs_per_k_d in place. At α=0: r_contrib=0, output
== sigmoid(logit_baseline) == probs_baseline → bit-identical.
per_horizon_prob_blend_reduce_alpha_residual
Reads probs_per_k (= p_final, post-blend) + grad_probs_per_k
(= ∂L/∂p_final from BCE) and computes:
d_logit[k,b,h] = grad_probs[k,b,h] * p_final * (1 - p_final)
d_residual[b,h] = tanh(α[h]) * Σ_k d_logit[k,b,h]
d_alpha[h] = sech²(α[h]) * Σ_{k,b} d_logit[k,b,h] * residual[b,h]
No separate prob_blend_bwd needed — chain-rule equivalence
∂L/∂logit_baseline = ∂L/∂r_contrib (both flow through the same
sigmoid derivative) means the existing GRN backward is UNCHANGED.
trainer/per_horizon_state.rs extensions:
forward_with_blend(ln_b_out, logit_per_k, probs_per_k)
Pool fwd → context_h; head fwd → residual; prob_blend fwd
in-place rewrites probs_per_k.
backward_through_blend(probs_per_k, grad_probs_per_k, ln_b_out,
grad_ln_b_out_target)
Reduce kernel → d_residual + d_alpha. Then:
head bwd → d_w_res_scratch, d_bias_res_scratch, d_context.
pool bwd → d_q_h_scratch, += grad_h_enriched_seq_d.
Per-batch scratches reduced to shared grads host-side
(n_batch ≤ 64 → sub-millisecond on host).
adamw_step()
Steps the four optimizers using the shared grad buffers.
zero_grads()
Called once per step before forward to clear scratch.
trainer/perception.rs step_batched integration:
── 4.5 (after GRN K-loop, before BCE): zero_grads + forward_with_blend
overwrites probs_per_k_d with p_final.
── 5 (existing BCE consumes probs_per_k_d as today; grad_probs is
now ∂L/∂p_final automatically).
── 5a (after BCE, before ISV-lambda + heads bwd): backward_through_blend.
Existing GRN bwd path is UNTOUCHED — the chain rule absorbs
the bias.
── 9 (after existing 17 AdamW group steps): per_horizon.adamw_step
updates Q_h, w_res, bias_res, α.
Verification:
- 34 ml-alpha lib tests still green.
- Per-horizon kernel numgrad parity (C21, C22) still green.
- Per-horizon end-to-end pipeline smoke (C23, including the
alpha=0 byte-identity invariant) still green.
- Full workspace builds clean.
Closes the kernels+wiring portion of #203 (per-horizon attention pool
kernels + wiring). What remains (#204): 30-epoch × 3-fold A/B vs
single-Q baseline. The branch is ready for that sweep when GPU time
is budgeted; the implementation is structurally adoption-safe
(α=0 → identity to baseline) so it can be merged before the A/B if
desired.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
5eb6567c4c |
feat(ml-alpha): PerceptionTrainer per-horizon trainer state (C24)
Allocates the device buffers + AdamW optimizers + kernel bindings for
the per-horizon attention-pool training path from C21 + C22, bundled
into a single PerceptionTrainer.per_horizon field.
crates/ml-alpha/src/trainer/per_horizon_state.rs (NEW):
PerHorizonTrainState — owns:
Learnable params (zero-init for α + bias, Xavier-scale for Q_h,
0.1× scale for w_res):
q_h_d [N_HORIZONS, HIDDEN_DIM] — attention queries
w_res_d [N_HORIZONS, HIDDEN_DIM] — residual head weights
bias_res_d [N_HORIZONS] — residual bias
alpha_d [N_HORIZONS] — learnable gate (init 0)
Forward intermediates (per-batch):
context_d [B, N_HORIZONS, HIDDEN_DIM]
attn_weights_d [B, N_HORIZONS, K]
residual_d [B, N_HORIZONS]
Backward grad scratch + reduced grads + AdamW state.
Kernel bindings: PerHorizonAttentionPool + PerHorizonResidualHead.
PerHorizonTrainState::new(dev, n_batch, k_seq, lr, seed)
Allocates all buffers, runs Xavier-style init, constructs four
AdamW optimizers (q_h, w_res, bias_res at param-LR; α at 0.25× LR
per spec §5 open Q2 default — slow gate ramp). Captures the seed
via wrapping_add(0xA110C00A) from cfg.seed for determinism.
PerHorizonTrainState::zero_grads()
Clears all grad-scratch buffers between training steps.
trainer/mod.rs:
pub mod per_horizon_state — module export.
trainer/perception.rs:
PerceptionTrainer gains one field:
pub per_horizon: PerHorizonTrainState
Initialised in new() with cfg.n_batch + cfg.seq_len + cfg.lr_cfc.
α-gate init=0 ⇒ tanh(0)=0 ⇒ contribution to per-batch logits is exactly
0 at step 0 (per C23 alpha_zero_init_is_identity_to_baseline byte-equality
test). Adopting this commit produces bit-identical training behaviour
to the previous commit until C25 wires the forward+backward calls into
step_batched; even then the α=0 init means a one-epoch smoke against
existing baseline should match within FP rounding noise.
All 34 ml-alpha lib tests + 4 per-horizon GPU tests (numgrad pair +
end-to-end pipeline pair) green.
Next:
C25 — forward + backward integration into step_batched. The new path
runs as ADDITIONAL kernel launches before/after the existing
captured graph (not inside it) so the graph stays unchanged
and the integration risk is contained.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
69c64f2266 |
test(ml-alpha): per-horizon pipeline end-to-end smoke + α-gate (C23)
Composes the C21 + C22 kernels with a host-side learnable α-gate to
prove the full per-horizon contribution path works end-to-end without
yet doing the captured-graph integration in PerceptionTrainer.
Pipeline:
LNb [B, K, HIDDEN_DIM]
→ per_horizon_attention_pool_fwd → context_h [B, N_HORIZONS, HIDDEN_DIM]
→ per_horizon_residual_head_fwd → residual [B, N_HORIZONS]
→ final[b, h] = baseline[b, h] + tanh(α[h]) * residual[b, h]
Two tests cover the critical invariants for adoption-safety:
alpha_zero_init_is_identity_to_baseline
With α = [0, 0, 0, 0, 0] and any random Q_h / w_res / bias_res,
final_logit MUST be bit-identical to baseline_logit (because
tanh(0) = 0). Verified via to_bits() byte equality. Proves that
initialising the new variant with α=0 makes it a strict superset
of the existing path — switching to AttentionPoolVariant::PerHorizon
cannot regress before any training has happened.
alpha_nonzero_changes_output_and_grads_flow_end_to_end
With α = [0.5, -0.3, 0.2, -0.1, 0.4]:
* final ≠ baseline (residual contributing) ✓
* all final logits finite ✓
* full backward chain (residual_head_bwd → attention_pool_bwd)
produces finite d_Q_h_scratch + finite d_LNb with at least
one non-zero entry in each → gradients flow back to both the
attention queries and the LN_b input ✓
This closes the kernel-side correctness story. The remaining
integration commits (C24+) are operational:
C24: extend CheckpointV1 → V2 (add q_h, w_res, bias_res, alpha
fields; V1 files load as Variant::SharedQuery)
C25: PerceptionTrainer wiring — allocate the device buffers, fold
attention + residual + gate into the captured graph, plumb
gradients into AdamW's param list
C26: 1-epoch smoke (assert no NaN, loss decreases vs baseline) —
needs real training data + multi-GPU time
C27: 30-epoch × 3-fold A/B (task #204) — decision gate per
docs/superpowers/specs/2026-05-18-per-horizon-attention-pool-design.md
§0 falsifiable claim
C24-C25 are 1-2 day work even when carefully scoped; C26-C27 need
real GPU-hours + result analysis. C21-C23 land the validatable kernel
correctness piece without committing to that time investment yet.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
8c335caef7 |
feat(ml-alpha): per-horizon residual head kernel + numgrad parity (C22)
Companion kernel to C21's per_horizon_attention_pool. Computes a
per-horizon scalar residual from each horizon's context vector:
residual[b, h] = Σ_d w_res[h, d] * context_h[b, h, d] + bias_res[h]
Designed to be added (behind a learnable α-gate) to the existing
multi_horizon_heads logit output — keeps the existing GRN head kernel
completely unchanged. The per-horizon attention pool's contribution
flows through this lightweight projection without weight-shape
changes elsewhere or checkpoint-V2-bumping.
Path A integration sketch (deferred to follow-up commit C23):
alpha_logit_per_horizon = existing_head(h_K)[h] # from current path
+ tanh(α[h]) * residual_kernel(context_h)[h]
where α[h] is a learnable 5-vector init'd to 0 (no effect at start).
Training discovers per-horizon whether the residual contributes.
This is a strict superset of the existing path — α=0 → bit-identical
to today.
Backward kernel produces:
d_w_res — per-block scratch [B, N_HORIZONS, HIDDEN_DIM]
for host reduce_axis0 → shared [N_HORIZONS, HIDDEN_DIM]
d_bias_res — per-block scratch [B, N_HORIZONS], same reduction
d_context_h — per-batch indexed; += chained with attention bwd
Single-writer discipline preserved (no atomicAdd per
feedback_no_atomicadd.md); horizon loop inside the per-batch block.
Numgrad parity test:
- B=3, N_HORIZONS=5, HIDDEN_DIM=128 fixture.
- Loss = Σ residual_out (so d_residual = 1).
- Probes 8 random w_res indices, all 5 bias_res entries, 8 random
context_h indices via central-difference at ±eps=1e-2.
- All within 5e-2 rel-tol or 5e-3 abs-floor.
- Passes on RTX 3050.
Same scope discipline as C21: kernel + binding + numgrad first;
trainer wiring + α-gate + smoke training + A/B sweep follow once
both kernels are individually validated (now done).
Closes the second kernel-correctness portion of #203. Remaining:
C23: trainer wiring (capture attn_pool fwd into the graph; sum
residual into existing head output with α-gate)
C24: CheckpointV1 → V2 bump (add q_h, w_res, bias_res, alpha fields)
C25: 1-epoch smoke (assert no NaN, loss decreases vs baseline)
C26: 30-epoch × 3-fold A/B (#204) — decision gate per spec §0
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
edc449eecf |
feat(ml-alpha): per-horizon attention pool kernel + numgrad parity (C21)
First implementation slice of the per-horizon attention pool design
(docs/superpowers/specs/2026-05-18-per-horizon-attention-pool-design.md).
Lands the kernel + Rust binding + numgrad verification; downstream
wiring into PerceptionTrainer's captured graph + CheckpointV2 bump +
A/B sweep are follow-up commits gated on this proving correctness.
Kernel (cuda/per_horizon_attention_pool.cu):
per_horizon_attention_pool_fwd
Q_h[N_HORIZONS, HIDDEN_DIM] × LNb[B, K, HIDDEN_DIM]
→ context_h[B, N_HORIZONS, HIDDEN_DIM]
attn_h_weights[B, N_HORIZONS, K]
Per-block math identical to the single-Q variant, looped over
N_HORIZONS sequentially within each batch's block. Grid stays
(B, 1, 1) so backward grad_ln_out writes are race-free
(per feedback_no_atomicadd.md — no cross-block contention).
Per-batch shared mem ~k_seq + BLOCK + HIDDEN_DIM floats.
per_horizon_attention_pool_bwd
Same chain-rule pattern as attention_pool_bwd but with the horizon
loop inside the block: each (b, h) slice updates grad_ln_out in
place (sequential horizon accumulation), grad_Q_h is written as
per-block scratch [B, N_HORIZONS, HIDDEN_DIM] for host reduce.
Single-writer discipline preserved.
Rust binding (src/per_horizon_attention_pool.rs):
PerHorizonAttentionPool::{new, forward, backward}. Self-contained;
doesn't yet touch PerceptionTrainer or CfcTrunk. Loads the cubin
via the standard env!("OUT_DIR") path. Dynamic shared-mem byte
count computed per launch from k_seq.
Numgrad parity test (tests/per_horizon_attention_pool_numgrad.rs):
- B=2, K=8, HIDDEN_DIM=128, N_HORIZONS=5 fixture.
- Loss = Σ context_h (so d_context = 1 everywhere — clean analytical).
- Backward kernel produces analytical grads; central-difference of
forward kernel at ±eps=1e-2 across 8 random Q_h indices + 8 random
LNb indices verifies analytical matches CD within 5e-2 rel-tol or
5e-3 abs-floor.
- Passes on RTX 3050.
build.rs picks up the new .cu file automatically via the existing
KERNELS list; cubin compiles cleanly at sm_86 + sm_89.
Same scope discipline as Phase 2D.2 (VSN numgrad) — kernel correctness
first, integration second. The follow-up commit set per the spec §3
appendix:
C22: extend multi_horizon_heads.cu signature to accept per-horizon
context input + bump head_w shape to [N_HORIZONS, 2*HIDDEN_DIM]
C23: wire PerHorizonAttentionPool into PerceptionTrainer + CfcTrunk
captured graph behind AttentionPoolVariant config flag
C24: CheckpointV1 → V2 bump with discriminant + optional q_h field
C25: smoke training (one epoch, no NaN, loss decreases)
C26: 30-epoch × 3-fold A/B sweep (#204) — decision gate per spec §0
Closes the kernel-correctness portion of #203.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
f5632649ca |
spec(ml-alpha): per-horizon attention pool design (C20)
Captures the brainstormed "alternative attention pool variants"
follow-on from the original real-LOB integration brainstorm (Axis 1,
deferred from the LOB workstream as a separate model-side spec).
Design:
Replace shared learned query Q[HIDDEN_DIM] with per-horizon queries
Q_h[N_HORIZONS, HIDDEN_DIM]. Per-horizon softmax + context vectors
feed multi-horizon heads directly (PATH A) — each horizon attends
to a different part of the K=6000 LN_b output sequence. CfC k=0
state is initialised by the MEAN of per-horizon contexts so the
K-loop recurrence + state amplification (per
pearl_state_amplifies_short_horizon_into_long_horizon) survives.
Heads consume per-horizon context concat CfC h_K (residual) with a
default 75/25 weight split.
Falsifiable claim (§0): A/B-tested win means h6000 mean_auc lifts by
≥ +0.01 absolute OR per-horizon distribution shifts toward short
horizons (h1000, h300) with no net h6000 loss. The 3-fold variance
band on the current architecture (mean_auc 0.7749 ± 0.024) means a
+0.01 lift is within noise — a meaningful effect needs ≥ +0.024 or
qualitative distributional shift.
Two new kernels (per_horizon_attention_pool_fwd + _bwd) + signature
extension on multi_horizon_heads_{fwd,bwd}. Variant-toggle config flag
(SharedQuery vs PerHorizonQuery) keeps the existing path fully
functional; new variant is opt-in. CheckpointV1 → V2 with explicit
discriminant + optional q_h field; V1 files load as SharedQuery, new
V2 training writes the discriminant.
Three validation rings:
1. Per-(b,h) numgrad parity at K=16
2. One-epoch smoke (no NaN, loss decreases)
3. 30-epoch × 3-fold A/B (#204) — decision gate per §0 falsifiable claim
Implementation explicitly deferred. The decision to invest depends on
(a) GPU time budget (~3-6 hrs on L40S × 5 GPUs for the A/B), (b)
whether per-horizon cost-frontier sweeps (#202 follow-ups) surface
viable horizons beyond h6000 that would benefit from per-horizon
specialisation, and (c) the 3-fold variance noise floor making the
expected effect size visible.
Next step when ready: invoke superpowers:writing-plans against this
spec for the ~6-8 commit implementation plan.
Closes the "good to have" question from the recent brainstorm with a
concrete decision framework rather than ad-hoc implementation.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
62b1fc0965 |
infra(argo): lob-backtest-sweep workflow + argo-lob-sweep.sh (C19)
Cluster fan-out for the `fxt-backtest sweep` single-machine path.
Reads the same grid YAML format as the binary; runs each cell on a
dedicated GPU pod in parallel; aggregates at the end on a CPU pod.
infra/k8s/argo/lob-backtest-sweep-template.yaml:
WorkflowTemplate `lob-backtest-sweep` with three job templates:
ensure-binary — cache-or-compile fxt-backtest by short-SHA into
/mnt/training-data/bin/<sha>/. Mirrors the
train-multi-seed-template.yaml ensure-binary
shape but for a single binary.
run-cell — single GPU pod (ci-training-l40s default per
feedback_default_to_l40s_pool.md). Receives
cell-name + every Run arg via inputs.parameters.
Writes artifacts to <sweep-root>/<sweep-tag>/<cell>/.
aggregate — CPU pod runs `fxt-backtest aggregate <sweep-dir>`
producing aggregate.parquet + pareto_frontier.json
at the sweep root.
DAG marker `# __SWEEP_CELLS__` replaced at submission time with N
WorkflowTask stanzas (one per cell), and the aggregate's
`dependencies: [ensure-binary]` is rewritten to include every
run-cell-* dep — so aggregate waits for ALL cells.
scripts/argo-lob-sweep.sh:
Companion submission script following the argo-train.sh pattern.
Parses the grid YAML via python3 + PyYAML (no `yq` dependency —
yq isn't used elsewhere in foxhunt scripts; python3+PyYAML is
universal in our CI images). Emits per-cell WorkflowTask stanzas
+ aggregate dependency list, awks them into the template, then
`kubectl apply` + `argo submit`. Supports --dry-run for offline
rendering and --watch for live log following.
Defaults match the spec / pearl set:
- sm_89 / ci-training-l40s default (override via --gpu-pool
ci-training-h100 for sm_90 + 80 GB)
- data root /mnt/training-data/futures-baseline/ES.FUT
- sweep results under /mnt/training-data/sweeps/lob-backtest/<tag>/
- sweep-tag defaults to <basename of grid>-<short-sha>
Verified locally:
- bash -n syntax-check passes
- --help renders
- --dry-run against the existing
config/ml/sweep_decision_stride_example.yaml renders a valid
workflow with 4 cells + correct aggregate dependency list
Live submission is operational work that needs cluster access to
verify; the rendered YAML follows the same conventions as the
existing argo-train.sh workflows that ship in this repo.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
8d3efa450e |
test(ml-backtesting): integrated Ring 2 fuzz over full pipeline (C18)
Existing lob_sim_fuzz only exercised book_update + market orders. This new test suite drives the FULL integrated pipeline under random adversarial conditions: apply_snapshot (random-walk book) → step_resting_orders (random signed trade-flow signal) → broadcast_alpha (random per-horizon probs every 4th event) → step_decision_with_latency (mixed latency=0 + latency=50ms cells) → submit_market_immediate (immediate path) → pnl_track_step + isv_kelly_update_on_close Warm-starts every backtest with random-but-plausible Kelly state (at least h4 set credibly profitable so decisions actually open positions). Random target_annual_vol + annualisation_factor + max_lots per decision to vary the Kelly cap. Per-50-event invariants: • book monotonicity (bid_px[k] ≤ bid_px[k-1], ask_px[k] ≥ ask_px[k-1]) • no-crossed-book (ask[0] > bid[0]) • Pos.realized_pnl + Pos.vwap_entry finite (no NaN/Inf leaks) • Pos.position_lots in plausible range (|≤100|) • All 5 per-horizon IsvKellyState fields finite Three test sizes: integrated_fuzz_n1_short — N=1, 200 events integrated_fuzz_n8_medium — N=8, 500 events integrated_fuzz_n64_long — N=64, 1000 events All three pass on RTX 3050. The N=64 × 1000 case exercises 64,000 event-snapshots × 16,000 decisions × ~12,800 trade attempts without any invariant violation or NaN propagation. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |