4-task atomic ladder (Q1-Q4) implementing spec 566e8bcb0. Each task
maps to one commit per feedback_no_partial_refactor:
Q1: Cold-start stopgap (bytecode max-confidence policy upload)
Q2: CBSW kernel aggregator + Tier 1 revert (atomic)
Q3: Memory pearl pearl_conviction_bootstrap_for_kelly_aggregation
Q4: Parallelism spec §3.4 cross-reference
Kernel ABI unchanged across all 4 commits — Q2 modifies only
decision_policy_default's body (decision_policy_program left as a
pluggable bytecode-VM experiment surface). Q2 atomically deletes
the Q1 stopgap field/CLI/YAML/harness branch in the same commit
that lands the kernel fix (feedback_no_legacy_aliases compliance:
grep returns zero hits for use_cold_start_stopgap post-Q2).
Validation gates per spec §9:
- Q1: cluster smoke n_trades > 100, total_pnl != 0 (diagnosis
confirmation). If n_trades = 0 still, STOP — downstream bug.
- Q2: pre-existing 11 CUDA tests still pass; 7 new cbsw_*
regression tests pass; cluster smoke n_trades > 100; ev/s
rate ratio Q2/Q1 ≥ 0.95.
Self-review confirms all 11 spec sections covered, no placeholders,
type/signature consistency across tasks.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Critical review surfaced 11 actionable issues in the v1 spec; 10 fixed
inline (#11 — legacy-comparison test — dropped per user decision since
empirical legacy behavior is already known: n_trades=0).
Performance:
1. expf → piecewise-linear ramp (~5× cheaper on GPU, no transcendental
in the hot path). Operationally equivalent: monotonic, saturating,
midpoint at K. Side-benefit: pure max-confidence at cold-start
(sq=0 instead of sigmoid's 11.9% leakage).
2. Single fused per-horizon pass replaces the two-loop sketch.
3. Tier 2 scope narrowed to decision_policy_default ONLY. The bytecode
VM (decision_policy_program) stays unchanged — runs only for custom
strategy experiments, never in production policy. Halves Tier 2's
surface area + test cost.
4. __device__ helpers cbsw_signal_quality + cbsw_weight introduced for
single source of truth.
Correctness bugs (silently present in v1 sketch, would have shipped):
5. strong_h and sq_min_active initializers added (were referenced
before init in the per-horizon loop).
6. Attribution mask at cold-start was setting all 5 bits because every
w[h] >= floor > 1e-9; trades would pollute ALL horizons'
recent_sharpe. Fix: binary split — at sq_min_active < 0.5 attribute
only to strong_h; at mature, attribute per weights > floor + epsilon.
7. sq_min was "min over h", which permanently locked the aggregator
into cold-start mode if any horizon never gets attributed (e.g.,
h6000-only-trading regime). Fix: sq_min_active = min over h with
n_trades_seen > 0; cold horizons don't gate maturity.
8. Opposing-horizons case now correctly fires on the strongest single
horizon at cold-start (piecewise-linear sq=0 → pure max-conf), with
explicit design-choice note explaining why this conservatism trade-
off favors firing.
Spec hygiene:
9. Rate validation committed to a measurement gate (Q2 ev/s ≥ 95% of
Q1 ev/s) instead of back-of-envelope estimate.
10. Kernel ABI explicitly stated as unchanged → feedback_no_partial_
refactor compliance is trivial at the kernel boundary.
12. Tier 1's use_cold_start_stopgap field is DROPPED in Q2 atomically
(not orphaned), and DoD checklist includes a grep-zero gate for
feedback_no_legacy_aliases compliance.
Risks updated: removed the sigmoid-narrowing risk (irrelevant now);
added register-pressure risk with the rate-gate mitigation; added
explicit "cold-start may lose on noisy trades" risk with empirical
escalation path (bump kelly_floor 0.20 → 0.40 if Q2 smoke shows large
negative PnL).
Math walk-throughs rewritten for piecewise-linear (different numbers
at cold-start): cold = pure max-conf, mature = pure weighted-sharpe,
clean separation at sq_min_active threshold.
Awaiting review of the revised spec before writing-plans dispatch.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
The post-trunk-grows threshold-tuning smoke (81decf40f) produced
n_trades=0 despite the model having a HEALTHY max-conviction
distribution (74.6% of decisions ≥ 0.30, 26.7% ≥ 0.70, full spread
across [0,1]). Diagnosis: the linear-weighted-mean aggregator in
decision_policy_default is structurally dilution-bound at cold-start
— single-horizon strong signals get washed out when uniform-floor
weights produce mean-over-horizons aggregation.
Solution: Conviction-Bootstrapped Sharpe Weighting (CBSW). Hybrid
max-confidence × weighted-sharpe with a per-horizon sigmoid transition
keyed on n_trades_seen vs MIN_TRADES_FOR_VAR_CAP. Cold-start: sig_mag
(decision-time conviction) drives weights AND max-confidence aggregator
fires single-horizon trades. Mature: recent_sharpe (historical) drives
weights AND linear-mean aggregator emphasizes strong-Sharpe horizons.
Permanent floor preserved per pearl_blend_formulas_must_have_permanent_floor.
Mirrors DQN bootstrap pattern (pearl_thompson_for_distributional_action_
selection): when historical estimates are uncertain, use available
signal as the bootstrap. Sig_mag is the ISV signal at decision time;
recent_sharpe is the ISV signal at trade-close time. Transition
self-terminates based on data accumulation, not time constants.
3-tier delivery (one spec, atomic commits):
Q1: Bytecode VM stopgap — upload max-confidence 7-instruction
program per backtest. Validates diagnosis; zero kernel work.
Q2: Kernel CBSW — replace weight + aggregator in both decision
kernels. 5 new regression tests covering cold/mature/transition.
Q3: New memory pearl pearl_conviction_bootstrap_for_kelly_aggregation
capturing the lesson.
Q4: Cross-reference from parallelism spec (deployability sweep
depends on CBSW being live to produce meaningful verdict).
The parallelism work (P1-P6) is fully working — confirmed by the
threshold-tuning smoke completing end-to-end (Succeeded status, 500k
decisions, artifacts written, aggregator parquet emitted). What's
blocked is the deployability VERDICT, because the dilution bug means
all variants would show n_trades=0 regardless of cost/latency/threshold.
CBSW unblocks the verdict.
Awaiting review before transitioning to writing-plans for Q1-Q4
implementation plan.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Critical review surfaced 9 issues; all fixed inline:
1. §1 — Reframed "140× amortization" as "amortization on forward; sim
runs parallel". True win is on the ~2ms forward shared across 140
cells, not on sim work which scales linearly with n_backtests.
2. §2 — Made the 1h target vs firm bound explicit (≤1h target, ≤2h
firm). Acknowledges Graph capture realistic speedup is 1.2-1.5×,
not 2×.
3. §5 — Dropped atomicAdd. Plain `+=` under the single-writer-per-block
convention (existing pattern across sim kernels). No race.
4. §4 — Documented max-over-horizons threshold rationale (vs per-
horizon or aggregate-conviction). Flagged per-horizon as a follow-up
tweak if dilution pathway matters in the verdict.
5. §7 P5 + §10 risk — Captured-vs-uncaptured tolerance is 1e-5 relative,
NOT strict bit-identity. CUDA Graph capture can reorder reductions
harmlessly by 1 ULP; strict bit-identity would be a false-positive.
6. §8 — Made explicit that parallel_sim_equivalence + independence
tests are BOTH required. Equivalence alone is necessary but not
sufficient (a shadow-backtest[0] bug still passes equivalence).
7. §3.3 — Specified output schema: `cell_W{n}/sim_<variant_name>/`
with summary.json carrying a resolved `sim_config` block (verdict
emitter reads that, not the directory name).
8. §7 P4 — Enumerated apply_fill_to_pos call sites + added grep-verify
step before commit, so no fill path silently loses cost.
9. §9 — Tightened rate-validation gates with hard targets (P2 ≤90s,
P5 ≤60s) instead of generous minute envelopes. Added §9.1 stride=8
fallback as explicit Plan-B if Graph capture under-delivers.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Brainstormed live with claude-opus-4-7. Targets ≤1h wall-clock per
4-quarter sweep on 1-3 L40S pods (vs ~750 GPU-hours naive). Three
multiplicative levers: per-backtest sim parameter matrix (~140×),
CUDA Graph capture (~2×), pod-level parallelism (~3×). Verified by
code-read during scoping:
- forward_only has no graph capture (X11 plan said so, X11 commit
didn't ship it)
- two host-roundtrip loops in sim.rs need GPU-ization for batching
to pay off (dispatch_latent_market_orders + close-detection)
- cost system is a literal-zero placeholder in pnl_track.cu:92
7-commit atomic ladder (P1-P7), each gated by tests + a measured rate
target. bf16 explicitly deferred to follow-up.
Awaiting review before transitioning to writing-plans for the
implementation plan.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Supersedes the 2026-05-19 deployability spec (commit 07d5de504). The
prior spec assumed CfcTrunk::save_checkpoint was the producer-side
wiring point — discovered at execution time that alpha_train trains via
PerceptionTrainer (full v2: VSN + Mamba2 ×2 + LN ×2 + attn-pool + CfC +
heads), not the simpler CfcTrunk. The existing LOB backtester loads
CheckpointV1 envelopes that only know about CfC weights, so there is no
producer for a checkpoint containing the full v2 model.
New scope: one bigger spec covering refactor + deployability end-to-end.
Phase 1 (X0–X19, code commits): grow CfcTrunk to own the full v2
inference graph; restructure PerceptionTrainer to wrap a trunk + add
training-only state (grads, AdamW). Discipline: bit-equivalence golden
fixture (X0) gates every refactor commit (X1–X11). CheckpointV2
envelope (X12) replaces V1. Verdict emitter (X17) reuses the tiered
classification (Pass-robust / Pass-nominal / Fail-inconclusive / Fail /
Fail-degenerate) from the superseded spec.
Phase 2 (Argo runtime): production training → smoke gate → threshold
pre-registration → 560-cell deployability sweep → verdict + memory
update.
Hard gate before Phase 2: post-refactor fold-0 smoke must reproduce
recorded 3-fold A/B numbers (best_mean_auc 0.7529, best_h6000 0.7639,
both within ±0.010 absolute) from project_ml_alpha_v2_ab_verdict
memory. Prior spec marked SUPERSEDED in its header, kept in history as
audit trail.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
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>
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>
The motivating "50% saturation hit-rate" observation came from gm67g
fold 0 (Option-2 config, commit 004b662c8 — itself a -0.003 mean_auc
regression vs ISV-σ at 410ab6b0e). Re-running the same diagnostic on
the actual production baseline (ISV-σ 3 folds: rxm5t/r57lx/x24d6,
logs retrieved from MinIO argo-logs bucket) on 215 horizon-epoch
observations:
saturated λ→AUC up: 8/13 = 61.5% median Δauc = +0.0025
non-saturated→AUC up: 96/202 = 47.5% median Δauc = -0.0012
difference: +14pp in favor of the controller working
The BCE-z-score controller is empirically correlated with the
optimization target. No evidence it's misaligned with AUC. Spec
premise falsified.
Spec retained in tree as historical record. The diagnostic
methodology itself (saturation→Δauc analysis on archived MinIO logs)
is the durable artifact and is documented in the supersede block.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
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 (004b662c8)
- Success: mean_auc lifts ≥ +0.005 AND median sat→next-AUC delta positive
Awaiting user review before plan handoff.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
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>
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>
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>
Revises all 12 issues from the self-review pass:
1. (HIGH) Soften bit-equivalence claim — single-sample helper and
batched kernel at B=1 are different CUDA kernels; FP order may
differ. Acceptance is relative_eq ≤ 1e-6, not bit-exact.
2. (HIGH) Explicit asymmetry: only cfc_step has a single-sample GPU
oracle. GRN/VSN/attn bwd rely on smoke + chain-rule preservation.
feedback_no_cpu_test_fallbacks.md forbids a CPU reference oracle.
3. (HIGH) Realistic targets — 3× floor, 5× stretch. Drops 15× claim
which was Amdahl-bounded under any reasonable assumption.
4. (MED) AdamW-after-reducer invariant stated explicitly: final grad
buffer is OVERWRITE by reducer, meaningful only after reducer ran.
5. (MED) New B=32 smoke test (stacked_trainer_loss_shrinks_at_batch_32)
actually exercises the cross-batch reduction path; existing
perception_overfit suite is all B=1.
6. (MED) Rollout commit 1 bundles reduce_axis0 + first consumer
(cfc_step refactor) to avoid feedback_wire_everything_up.md
orphan-kernel anti-pattern.
7. (LOW) Drop "merge to ml-alpha-phase-a" — user already chose direct
commits to that branch; clarify in Rollout.
8. (LOW) Add explicit scratch-sizing formula:
scratch_bytes ≈ B × Σ(param tensor sizes per kernel).
9. (LOW) Remove resolved open question (memset_zeros ordering).
10. (STRUCT) Post-refactor bottleneck analysis section — names the
next L40S floor (Mamba2 scan, launch latency, cuBLAS).
11. (STRUCT) cudaFuncSetAttribute note — refactored cfc_bwd's per-
block shared mem drops to 1 KiB; no attribute change needed.
12. (STRUCT) Explicit Rollback section — atomic-commit-per-kernel +
revert strategy; rollback baseline is the spec commit
(54aa69c10) on ml-alpha-phase-a.
Plus locks the target pool to L40S — speedup must be attributable to
the refactor, not to a hardware bump. H100 / BF16 / larger batch
become candidates for a follow-up spec once the L40S floor is known.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Documents the design for fixing the single-SM bottleneck in five
backward/K-loop kernels: cfc_step (fwd+bwd), GRN bwd, VSN bwd, and
attention_pool bwd. All currently use grid=(1,1,1) with an internal
n_batch loop — on L40S (142 SMs) with B=32 this puts <1% of the GPU
to work in the K-loop critical path.
Architecture: block-per-batch (grid=(B,1,1)) for the kernel body, plus
per-batch grad scratch buffers reduced via a single parameterised
reduce_axis0 kernel (block tree-reduce, no atomicAdd per
feedback_no_atomicadd.md). Same pattern as the existing LayerNorm bwd
reducer — CUDA-Graph-safe, debuggable, and consistent with foxhunt's
no-cooperative-groups discipline.
Target: ≥3× epoch wall speedup (stretch 8-15×). Makes 3-fold CV
tractable (10.5h → 2-3h) and unblocks decision-stride / state-dim
sweeps that compound the gain.
Acceptance gates: (a) all 8 perception_overfit smokes still converge,
(b) new B=1 bit-equivalence test asserts the refactored batched bwd
kernel at B=1 matches the existing single-sample helper byte-for-byte,
(c) cluster A/B vs t6z89 baseline shows AUC trajectory within ±0.005
and epoch wall ≥3× faster.
Atomic refactor per kernel — one commit per kernel covering the
kernel rewrite, scratch buffer alloc, reducer launch wiring, and
smoke. No "_legacy" parallel kernels per feedback_no_legacy_aliases.md
+ feedback_no_partial_refactor.md.
Open implementation-plan decisions: exact memset_zeros ordering inside
the captured graph, batch-vs-per-tensor reducer launches, optimal
block_dim for reduce_axis0 itself.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
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>
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>
Per-horizon P(up) at h ∈ {30, 100, 300, 1000, 6000} snapshots forward.
Single-block 5-thread kernel; each thread is its own 128-dim dot
product + sigmoid. No atomicAdd.
Tests (5/5 pass on sm_86) assert invariants only:
- sigmoid output ∈ [0, 1] for all heads
- zero weights + zero bias → 0.5 exactly
- bias = +20 → saturates near 1
- bias = -20 → saturates near 0
- per-head independence (mixed-bias configuration)
Addendum updated to explicitly state no-CPU-mirror discipline per
feedback_no_cpu_test_fallbacks.md.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Mid-execution architecture revision: Mamba2 stays as a sequence
encoder; CfC becomes the layer on top (replacing the Phase 1d.3 MLP
stacker). Gate becomes 'stacked AUC >= Mamba2-only stacker AUC at
every horizon' — proves the CfC layer is additive, rather than CfC
alone beating Mamba2 alone.
Plan 1 kernels (cfc_step, heads, projection, BCE, AdamW, Graph A)
are unchanged. Only CfcTrunk's forward path gains a Mamba2 prefix
that consumes the snapshot stream and emits a 128-dim h_mamba which
CfC reads. The Mamba2 kernel (mamba2_alpha_kernel.cubin) is already
in the build.
Option B (parallel + fused dual-stream) is documented as the Plan 2
fallback if the stacked gate fails.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Cargo.toml: drops gbdt; adds memmap2 + approx; keeps ml-core only
(cannot depend on ml: would cycle since ml depends on ml-alpha for
the Mamba2 gate baseline).
build.rs: compiles 7 cubins (mamba2_alpha + 6 new placeholders)
with -O3 --use_fast_math --ftz --fmad. Skips kernels whose source
isn't present yet so partial check-ins work. Every env::var paired
with rerun-if-env-changed per the canonical build pearl.
src/pinned_mem.rs: local copy of MappedF32Buffer (mirrors
ml::cuda_pipeline::mapped_pinned::MappedF32Buffer). Drives the only
permitted CPU<->GPU path per feedback_no_htod_htoh_only_mapped_pinned.
Eventually the move-to-ml-core refactor will deduplicate; out of
scope for the Phase A branch.
Addendum: updates the import path to ml_alpha::pinned_mem.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Plan 1 was written against an older cudarc device-centric API. cudarc 0.19
moved alloc/launch ownership to the stream (partly for CUDA Graph capture
hygiene). This addendum pins:
- MlDevice -> CudaContext -> CudaStream construction
- Cubin load + module + function caching
- MappedF32Buffer staging -> DtoD-async -> CudaSlice (canonical CPU->GPU)
- Slow-path readback via DtoD into a staging MappedF32Buffer
- launch_builder(&func).arg(...).launch(cfg) idiom
- IsvBus and MappedPinnedSnapshotSlot/FillSlot using the real
MappedF32Buffer API (host_slice_mut, read_all, dev_ptr field)
- CUDA Graph A capture via stream.begin_capture / end_capture / instantiate
Plan 1 kernel .cu source, CPU oracles, finite-diff thresholds, smoke
criteria, and gate logic are unchanged. Only Rust binding code uses
these patterns.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Cleanup of compiler warnings flagged by both local cargo check and the
cluster ensure-binary log. Per `feedback_no_hiding`, every site is
either deleted or wired up — no #[allow] suppressions.
Lib (5 sites):
- gpu_backtest_evaluator.rs:34 — drop unused DevicePtrMut.
- gpu_dqn_trainer.rs:49 — drop unused DevicePtrMut (8 device_ptr_mut
calls don't need the trait import in current cudarc). Line 19852:
drop unnecessary parens around `b * sh2`.
- training_loop.rs:20 — drop unused DevicePtrMut; unbrace single-
symbol use at 5766.
- state_reset_registry.rs:4 — delete the 10-symbol use-block of slot
constants. Names appear in description strings (documentation only),
symbols are never referenced.
Examples (3 sites):
- alpha_dqn_h600_smoke.rs:181, 186 — drop COL_RAW_CLOSE, FEAT_DIM,
FillCoeffs, FillModel imports.
- alpha_baseline.rs:79 — delete unused MappedI32::read. Batched path
uses read_all for N-element action readback; the single-element
method was leftover from the pre-batched legacy path.
Lib + examples now have zero removable warnings. The remaining
unsafe_block lints (each cudarc kernel launch needs unsafe) are
structural and not actionable under the project's -W unsafe-code
policy.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Closes the TrainingPersist loop for the regime defense per the
KELLY_F_SMOOTH precedent (pearl_kelly_cap_signal_driven_floors).
Three layers:
(1) Per-step: alpha_regime_vol_update_kernel tracks min(vol_obs > 1e-12)
in new slot 553 (REGIME_VOL_OBS_MIN_INDEX). Filtered against artifact-
zero observations from stationary snapshots where curr == prev.
(2) Per-cell: stacker_threshold_controller_update gains a fifth branch
that reads vol_ref (slot 550) at cell-end, computes
`target = floor_target_ratio × vol_ref`, slow-EMAs the floor anchor
(slot 552) toward the target with rate `floor_update_rate` (0.1 per
cell). Subfloor 1e-12 inside the kernel guards against the anchor
collapsing to zero.
(3) Per-invocation: alpha_baseline reads
`config/ml/alpha_baseline_state.json` at startup and seeds slot 552
from the `regime_vol_ref_floor` field. At end of main(), the learned
floor is written back via tmp+rename atomic write so concurrent
walk-forward invocations see a consistent file. Matches the
cross-fold-persistent shape of KELLY_F_SMOOTH.
Block extended to 15 slots (539..=553). Smoke + kernel unit test
pass -1/-1 for the new floor-controller indices (backward compat).
Walk-forward CV verdict (Q1 fxcache, 3 sequential folds):
iteration fold-A fold-B fold-C mean ± SD
pre-defense (no regime) +91.52 -21.44 +46.74 +38.94 ± 56.88
hardcoded 1e-9 floor -19.77 +65.04 +6.45 +17.24 ± 43.42
learned floor (0.5 × cell_min) +74.78 -12.72 +8.80 +23.62 ± 45.59
learned floor (0.1 × vol_ref) -19.53 -26.51 +15.46 -10.20 ± 22.49
Controller infrastructure is structurally correct (loop closes, floor
persists across invocations, kernel + disk + ISV all roundtrip). The
TUNING is data-dependent — single-quarter CV doesn't have enough
regime diversity to anchor the floor against. Multi-quarter fxcache
validation is the next step (built cluster-side on the 9-quarter
2024-Q1..2026-Q1 ES futures dataset, downloaded as a single artifact
for local CV).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Two coupled fixes to the vol-EMA regime detector exposed by walk-forward
CV after all features were promoted to always-on:
(1) Bootstrap window for vol_ref (slot 551 = REGIME_VOL_REF_SAMPLES)
- Replace the Pearl-A "first observation replaces directly" bootstrap
with a running mean over the first N=100 vol_ema observations, then
switch to β-tracking. For IID observations the running-mean estimator
has variance σ²/N — a 100-sample mean is 10× less noisy than the
single-shot replace.
(2) Permanent floor on vol_ref (slot 552 = REGIME_VOL_REF_FLOOR)
- The bootstrap alone exposed the asymmetric deadband-deadlock: if
vol_ref converged to a tiny value during a calm initial stretch,
vol_ref / vol_ema fired the moment any realistic vol resumed and
Kelly stayed trapped at regime_scale_floor=0.25 forever. Floor
lives in ISV slot (TrainingPersist) with a hardcoded 1e-12 sub-floor
inside the kernel as numerical-underflow guard.
- Host seeds slot 552 with 1e-9. Future controller kernel will refine
this from observed cell-level vol minima with cross-fold persistence.
Walk-forward CV on Q1 fxcache (3 folds, window=700K, train_frac=0.6):
cost fold-A fold-B fold-C mean ± SD
0.00 -19.77 +65.04 (100%) +6.45 +17.24 ± 43.42
Fold B turnaround is the headline: -47.64 (bootstrap-only) → +65.04
(bootstrap + floor) confirms the floor is the load-bearing fix.
Cross-fold std-dev compressed 24% at cost=0; mean dropped from +38.94
(pre-defense) to +17.24 (with defense). Classic mean/variance trade.
Block extended to 14 slots (539..=552). Kernel sig: vol_ref_floor
moved from f32 scalar to vol_ref_floor_index i32, so the anchor is
named/addressable in ISV.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Adds two coupled interventions on the regime fragility exposed by
walk-forward CV (mean Sharpe +27 ± 56 at half-tick across 3 folds —
std-dev ≈ mean means the strategy is regime-dependent).
(1) Vol-EMA regime detector (new ISV slots 549/550)
- New alpha_regime_vol_update.cu kernel: per inference step, reads B
parallel-env mid prices, computes cross-env mean squared log-return,
and maintains ISV[549]=REGIME_VOL_EMA (Wiener-α with 0.4 floor +
Pearl A bootstrap) and ISV[550]=REGIME_VOL_REF (slow tracker β=0.005,
≈200-step horizon).
- Block-tree-reduce (no atomicAdd), guards against zero/non-finite mids.
(2) Pre-emptive Kelly attenuation (modified stacker controller)
- stacker_threshold_controller.cu takes 3 new args: regime_vol_ema_idx,
regime_vol_ref_idx, regime_scale_floor.
- Multiplies its reactive Sharpe-error Kelly output by
regime_scale = clamp(vol_ref / vol_ema, 0.25, 1.0)
- Disabled when indices = -1 (backward-compatible smoke + kernel test).
(3) Cost-aware training (--train-cost-hi)
- alpha_compose_backtest --train-cost-hi: when > --train-cost, each
training epoch samples cost ~ U[lo, hi] so the Q-network learns
cost-conservative behaviour across the realistic ES range.
(4) Wiring
- alpha_compose_backtest --regime-scale enables both per-step regime
kernel firing during eval AND the regime hookup in the per-episode
controller call. Mapped-pinned mids buffers, all compute device-side.
- ExecutionEnv exposes current_mid() so the host gather reads the
active snapshot mid per env without leaking the private cursor field.
Smoke + test sites pass -1/-1 for regime indices (backward compat).
Doc: docs/isv-slots.md ledger for slots 549/550.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
T10 wires the C51 → Mamba2 backward chain in both binaries:
- New alpha_train_window_store_batched_kernel captures per-step windows
for end-of-epoch re-forward (Mamba2 cache required for backward).
- Mamba2Block::backward_from_h_enriched lets the C51 grad-input feed
directly into Mamba2 backward, skipping the unused W_out projection.
- alpha_c51_grad_input → backward_from_h_enriched → Mamba2AdamW.step
closes the loop in alpha_dqn_h600_smoke and alpha_compose_backtest.
Smoke (--temporal --c51): all 4 KCs PASS. R_mean -6.3 → +4.2 vs
Phase E.3 close R_mean -4.7 (no-temporal). EARLY_Q_MOVEMENT
calibration (mamba2_snapshot + mamba2_weight_distance) lifts the
diagnostic from 0.0023 (head-only) to 0.0590 (head + encoder),
giving an honest learning signal when the encoder absorbs gradient.
Backtest (--c51 --temporal --window-k 16 --isv-continual):
cost=0.0000 best τ=0.250 Sharpe_ann=+34.56 (was +10.41 head-only,
-22.54 frozen-Mamba2)
cost=0.0625 best τ=0.250 Sharpe_ann=+33.22
cost=0.1250 best τ=0.250 Sharpe_ann=+30.85 (Phase 1d.4 baseline: -4.0)
cost=0.2500 best τ=0.250 Sharpe_ann=+27.73
cost=0.5000 best τ=0.250 Sharpe_ann=+15.68
Caveat: in-sample results; OOS gate next.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
T15: training rewrite mirroring T14 eval. N_par parallel envs
lockstep H steps per epoch; ONE batched C51 update at B = N_par * H.
Expected ~15× speedup vs sequential.
- NEW kernel alpha_h_enriched_store_batched_kernel for batched
h_enriched slot writes
- Training section greenfielded: legacy sequential loop deleted
- CLI flag --n-train-par (default 50)
- Terminal next-state slot zeroed; done=1 at horizon masks Q_next
contribution in Bellman projection — no terminal Mamba2 forward
- docs/isv-slots.md updated per kernel-audit-doc hook
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The Phase E.4.A T14 backtest at B=1 with per-step stream.synchronize()
was running ~150μs/step × 9M steps = ~22 min — dominated by sync
overhead, not GPU compute. RTX 3050 Ti to L40S swap wouldn't help
(launch overhead is the bottleneck, not FLOPS).
Solution: batched parallel envs. N=cli.n_eval_episodes environments
run in LOCKSTEP per cell — ONE sync per step (instead of N syncs).
Expected ~30× speedup at N=500.
Changes:
1. ExecutionEnv snapshots → Arc<Vec<SnapshotRow>>
- new() wraps Vec into Arc internally (backward compat)
- new_arc() takes pre-existing Arc (for parallel envs)
- snapshots_arc() accessor for snapshot sharing
- 50MB × N memory duplication avoided
2. alpha_window_push_batched_kernel (NEW CUDA)
- Same chronological shift+insert semantics as single-env kernel
- Grid (state_dim_blocks, B, 1): one thread per (batch, feature)
- launcher: launch_alpha_window_push_batched
3. MappedI32 (per-binary) gains len param + read_all()
- smoke & backtest pass len=1 for existing single-int use
- backtest passes len=N for batched action readback
4. backtest binary eval loop GREENFIELDED
- Legacy sequential 'for ep in 0..N { for step in ... }' loop
body deleted entirely
- New: 'for step in 0..horizon' outer, lockstep over N envs
- Build N envs sharing snapshots_arc at cell start
- Per step: gather N states (CPU loop, <100μs for N=500) →
write to mapped-pinned [N, STATE_DIM] → push kernel B=N →
Mamba2 batched forward → C51 batched forward → Thompson
batched → ONE sync → read N actions → step N envs on CPU
- ISV-continual moved from per-episode to per-cell (single fire
with aggregate stats)
5. docs/isv-slots.md updated per kernel-audit hook
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The Phase E.4.A T8 wiring stored Mamba2's per-step cache.h_enriched
into h_enriched_buf_dev via a dtoh+htod sequence:
let h_host = stream.clone_dtoh(cache.h_enriched.cuda_data())?;
let mut buf_host = stream.clone_dtoh(&h_enriched_buf_dev)?; // <- whole buffer
for j in 0..hidden_dim { buf_host[slot_offset + j] = h_host[j]; }
stream.memcpy_htod(&buf_host, &mut h_enriched_buf_dev)?; // <- whole buffer
This violates feedback_cpu_is_read_only AND
feedback_no_htod_htoh_only_mapped_pinned. Worse, the buffer-wide
dtoh+htod every step is ~20K floats × 600 steps × 500 eps × 30 cells
= ~9M roundtrips totaling significant PCIe latency in the backtest.
Fix: new tiny CUDA kernel alpha_h_enriched_store_kernel in
alpha_window_push.cu (one thread per hidden-dim feature, writes
src[j] → buf[slot_offset + j]). Replaces the dtoh/htod sequence
in both smoke and backtest binaries.
Estimated speed-up at backtest scale: 3-6× on the temporal eval
path. Pure-GPU per-step inference restored — no synchronisation
points on the hot path.
docs/isv-slots.md updated per kernel-audit-doc hook.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Two documentation deliverables produced while T14 backtest runs:
1. Plan update (specs/2026-05-15-phase-e-4-a-temporal-foundation.md):
adds 'Execution Status' section reflecting actual T1-T14
progression. T5 deferred (real MBP-10 peek), T9 skipped (GRN
moved to E.4.B per integration notes), T10 partial (new C51
grad-input kernel landed but Mamba2 backward wiring deferred),
T14 in flight. Documents the 4 execution learnings:
research-first saved a week of duplicate kernel work; cheap
falsification experiments (Path 2, Path 3) avoided expensive
investments; C51 borrow was the largest single Sharpe-lift in
the session; GpuTensor/CudaSlice interop friction is the real
integration cost.
2. T10 patch sketch (specs/2026-05-15-t10-mamba2-backward-from-h-enriched.md):
ready-to-apply patch for ml-alpha::Mamba2Block adding a new
public method backward_from_h_enriched(cache, d_h_enriched).
Bypasses the W_out projection backward, accepts the
[B, hidden_dim] gradient from C51's grad-input kernel directly,
zero-initialises dw_out/db_out (AdamW step on zero grad is a
no-op with correct moment decay — effectively freezes W_out
params which is correct semantics since Phase E never uses
them). Includes the smoke binary wiring snippet that consumes
the new method via launch_alpha_c51_grad_input → Mamba2
backward → AdamW step. Application gated on T14 backtest
validation — if frozen Mamba2 already lifts Sharpe, T10
becomes optimisation rather than prerequisite.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Phase E.4.A Task 10 partial: adds the C51 gradient-w.r.t.-input
kernel needed to chain the C51 head's loss gradient back into the
Mamba2 temporal encoder.
Kernel signature: alpha_c51_grad_input_kernel reads probs[B, A, K],
m[B, K], actions[B], W[A*K, in_dim] and writes d_input[B, in_dim].
Math: dL/d_input[b,j] = Σ_k (p[b,a_taken,k] − m[b,k]) · W[a_taken*K+k, j]
(only the taken-action column of W contributes; restricted by the
sparse-over-actions C51 CE gradient structure).
Status: kernel + Rust launcher in place. Mamba2 backward NOT yet
wired in the smoke binary because ml_alpha::Mamba2Block::backward
takes d_logit [B, 1] (post-W_out scalar gradient), not
d_h_enriched [B, hidden_dim]. Two paths to complete:
1. Modify ml-alpha to expose backward_from_h_enriched(cache,
d_h_enriched) bypassing W_out.
2. Replicate the post-W_out backward sequence inline.
Both deferred pending backtest validation (T14): if frozen Mamba2
already lifts backtest Sharpe meaningfully, T10 becomes
optimisation rather than prerequisite. Smoke gate already passed
gate 1 (R_mean +3.9 vs C51-flat -1.1) with frozen weights.
docs/isv-slots.md updated.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Switch alpha_window_push from circular-buffer-with-head_idx layout
to shift+insert layout matching production mamba2_update_history.
Slot 0 = oldest, slot K-1 = newest after each push, matching
Mamba2Block's [B, K, in_dim] input contract directly (no reorder).
Cost: O(K-1) shifts per state_dim feature per push. For K=16,
state_dim=10: 10 threads × ~15 ops each = trivial.
Kernel signature: drops head_idx, adds K. Test updated to verify
chronological shift across 3 pushes into 4-slot buffer.
docs/isv-slots.md updated per kernel-audit-doc hook.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Phase E.4.A Task 6: tiny CUDA kernel that pushes a state[state_dim]
vector into slot `head_idx` of a circular window buffer
[K, state_dim]. Host tracks head_idx and zeros buffer on episode
reset. This is the GPU-side append primitive that the smoke
binary's per-step inference will call (T7) before Mamba2 over the
buffer (T8).
GPU smoke (alpha_window_push_circular_writes_to_indexed_slot):
writes state_a at slot 0, state_b at slot 2, verifies non-targeted
slots stay zero. PASS.
docs/isv-slots.md: documents kernel as slot-agnostic per
kernel-audit-doc hook requirement.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
C51 distributional Q-network with GPU Thompson selection borrowed
minimally from production (alpha_c51.cu: forward, project, grad,
expected_q, thompson_select kernels; ~260 lines). Uses Huber
negative-tail compression in projection per production
block_bellman_project_f. Action selection 100% GPU via mapped-pinned
i32 output + __threadfence_system + host volatile read (matches
gpu_training_guard MappedBuffer pattern).
Backtest result (2D sweep, 500 episodes per cell, 30 cells):
cost=0 C51 +10.41 vs linear-Q -15.72 (+26pt, BEATS Phase 1d.4
no-RL baseline +4.4 by 6pt)
cost=0.125 C51 -13.81 vs -29.17 (+15pt closes half-tick gap)
Win rate at cost=0 best τ: linear-Q 0.008 → C51 0.552.
Calibration hypothesis vindicated; documented in
memory/pearl_c51_thompson_closed_phase_e3_gap.md.
Also in this commit (Phase E.3 follow-up cleanup):
- --pruned-actions falsified (2.4× worse Sharpe). Documented in
memory/pearl_action_pruning_falsified.md.
- --real-spread falsified for ES futures (76% of bars at 1-tick floor).
- SnapshotRow bid_l/ask_l extended from [f32; 3] to [f32; 10].
L4-L10 synthesized in this commit; real MBP-10 peek lands in E.4.A T5.
- docs/isv-slots.md updated per kernel-audit-doc hook requirement.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Design doc (specs/): TFT-style architecture for Phase E execution
policy — sliding window → Mamba2 SSM → GRN trunk → MoE regime gate
→ C51 head → Thompson selector. Two core pillars added per user:
A) Full L1-L10 LOB depth input via hybrid MBP-10 peek
B) ISV-continual-learning: controllers fire at training AND
inference; Q-net weights frozen at inference but effective
policy adapts via ISV modulation
Plan doc (plans/): 14-task implementation plan for E.4.A foundation
(window buffer + L1-L10 depth + Mamba2 forward+backward + ISV-eval
controllers). Falsification gates: smoke R_mean improvement ≥ 50%,
backtest cost=0 Sharpe ≥ +8 (no regression vs C51-flat +10.41),
half-tick Sharpe ≥ -8 (closes 5pt+ of 10pt gap to Phase 1d.4
baseline -4.0).
TGGN (foxhunt Temporal Graph Gated Network) explicitly deferred to
Phase E.5+: existing CPU graph implementation + GPU adapter at
ml-supervised/src/tgnn/ — marginal benefit for single-instrument ES
futures vs the TFT-Mamba2 stack; revisit for multi-instrument
extension or production HFT inference layer.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Findings:
- Production Mamba2 (gpu_dqn_trainer) is coupled to SH2=256 trunk +
ofi_embed + ISV[8] temporal routing — not portable to Phase E.
- ml-alpha::mamba2_block::Mamba2Block is from-scratch, fully
configurable (in_dim/hidden_dim/state_dim/seq_len), GPU-pure with
forward_train/backward/AdamW. Used in Phase 1d.2 to lift AUC 0.50
to 0.66. ml-alpha is already a workspace dep of ml.
- GRN skipped for E.4.A — Mamba2 output goes straight to C51 head.
Reintroduce GRN in E.4.B if Sharpe gates don't pass.
- Controller-at-inference: kernel has no training-mode branches;
Wiener state preserved across episodes/cost cells for natural
live-deployment simulation.
Revises Tasks 8-10 of the plan: use Mamba2Block API instead of
writing custom kernels. Only new CUDA needed: alpha_c51_grad_input
(C51 gradient w.r.t. input features, for Mamba2 backward chain).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Phase E.2 Task 16. Engagement-rate self-correcting controller per
pearl_engagement_rate_self_correction. Single-block, single-thread
kernel; runs once per rollout-end boundary.
ISV slots driven:
543 STACKER_THRESHOLD_INDEX clamp [0, 0.5] P-controller on rate
545 TRADE_RATE_OBSERVED_EMA_INDEX Pearl A+D floored Wiener-α
546 STACKER_KELLY_ATTENUATION_INX clamp [0.1, 1.0] P-controller on Sharpe
Reads ISV[544] TRADE_RATE_TARGET_INDEX (TrainingPersist anchor, set once
at training start).
Control law:
observed = trade_count / max(decisions, 1)
ISV[545] ← Pearl_A+D_floored(observed, prev, x_lag)
[α* floor = 0.4 per pearl_wiener_alpha_floor_for_nonstationary;
controller co-adapts with policy → need responsive EMA]
err_rate = ISV[545] - ISV[544]
ISV[543] ← clamp(0, 0.5, ISV[543] + k_threshold · err_rate)
err_sharpe = rollout_sharpe - target_sharpe
ISV[546] ← clamp(0.1, 1.0, prev_atten + k_atten · err_sharpe)
where prev_atten = 1.0 if ISV[546] == 0.0 (sentinel-start)
else ISV[546]
Wiener-α is INLINE (not via canonical apply_pearls_ad_kernel chain)
because the EMA is part of the control loop, not a separate diagnostic
slot. Lower latency, fewer kernels per step.
Floor at 0.1 on Kelly attenuation per
pearl_blend_formulas_must_have_permanent_floor — can't be 0, would
zero out position sizing permanently.
Pub launcher `launch_stacker_threshold_controller` in alpha_kernels.rs
with full safety asserts. Slot indices passed as i32 args (decouple
slot numbering from kernel).
GPU smoke test `stacker_threshold_controller_smoke_matches_hand_
computation` verifies 2-iteration sequence:
Iter 1 (Pearl A): observed=0.30 → ISV[545]=0.30; ISV[543]: 0.05 → 0.052
Iter 2 (Pearl D): observed=0.05 → ISV[545]=0.175 (α* hit floor 0.5)
ISV[543]: 0.052 → 0.05275
Both within 1e-5 tolerance. Anchor slot 544 unchanged.
`cargo test -p ml --lib alpha_kernels`: 6 pass (compile witness + 5 GPU
smokes including this one) on RTX 3050 Ti in 2.18s.
Audit doc docs/isv-slots.md updated per Invariant 7.
Three stabilizers applied to the H=600 DQN smoke after initial run showed
unstable training (early_mvmt=2268× at lr=1e-6, NaN at lr=1e-4):
1. Reward normalization (--reward-scale, default 1000)
Rewards divided by scale BEFORE the Munchausen target. TD error
drops from ~1000 (raw reward magnitude at H=600) into O(1) target /
gradient / weight-update scale. Action selection + rollout-R
reporting use ORIGINAL rewards (so rvr math stays correct against
the Task 7c baseline).
2. Target network (--target-update-every, default 10 episodes)
Separate w_target_dev / b_target_dev buffers. Q_next(s') forward
uses target weights; SGD updates online only. Hard-update copies
online → target every K episodes. Breaks the V_soft(s') chase-its-
own-tail divergence of online-only Munchausen.
3. Gradient clipping (--grad-clip, default 1.0)
New `alpha_clip_inplace_kernel` in alpha_linear_q.cu (element-wise
clamp). Applied to dW and db after grad, before SGD. Safety net.
Diagnostic fix: weight_norm was direction-insensitive — orthogonal
rotations don't change ||W||_F, so early_mvmt read ≈0 even when training.
Switched to weight_distance_from_init = ||W_now − W_init||_F +
||b_now − b_init||_F (captures rotation). q_early = q_init + distance
so kernel's |q_early − q_init| / |q_init| ratio = distance / ||W_init||_F.
With lr bumped back up to 1e-4 (default for the stabilized config),
verified at horizon=100, n_episodes=200:
Q_SPREAD_EMA = 23.64 (≥ 0.05) PASS
ACTION_ENTROPY_EMA = 1.86 (≥ 1.0986) PASS
RETURN_VS_RANDOM_EMA = +0.586 (≥ 0.0) PASS
EARLY_Q_MOVEMENT_EMA = 0.0212 (≥ 0.01) PASS
Overall: PASS (H=6000 scale-up VIABLE)
early_mvmt grew monotonically (0.005 → 0.021) across the 200-episode
run — direction-sensitive diagnostic confirms genuine policy learning.
Audit doc docs/isv-slots.md updated per Invariant 7.
Next: H=600 / 1000-episode run on full data; if PASS holds, Task 13
(H=6000 scale-up) unlocks.
