552d91bf45dc11eed4da1c7bcd07ed35ba644a57
5787 Commits
| Author | SHA1 | Message | Date | |
|---|---|---|---|---|
|
|
552d91bf45 |
feat(ml-alpha): Tier 3 Phase 4-B — Band turnover controller + backward chain wiring
Closes Phase 4-A's open loop by making the no-transaction band LEARN its width. Two coupled mechanisms wired in: 1. Adaptive turnover-target controller — single-thread CUDA kernel (rl_band_turnover_controller.cu) reads per-step `frac_not_masked` and adjusts the turnover target via asymmetric Schulman-bounded adapter: tighten (×0.97/step) above 0.60 threshold, loosen (×1.05/step) below 0.20. Asymmetric LOOSEN rate is faster per `pearl_dead_signal_resurrection_discipline`. First-observation bootstrap on slot 809 with dedicated boot-done sentinel slot 810. 2. Backward chain wiring — rl_band_head_backward.cu propagates `grad_band_per_b [B × 2]` through the asymmetric ±|tanh|·N_max activation and linear projection into per-batch weight/bias scratch plus `grad_h_t [B × HIDDEN_DIM]` (OVERWRITE). Trainer reduces via `reduce_axis0` and folds into the encoder grad combiner alongside π/V/FRD/outcome heads using the existing `grad_h_accumulate_scaled` pattern. Band-weight Adam steps share LR with π (RL_LR_PI_INDEX) — same plateau dynamics as policy. Plus a GPU reducer rl_band_frac_aggregate.cu (single-block tree-reduce over batch → writes to slot 812) and a host-side launch sequence in step_with_lobsim_gpu_body that runs frac_aggregate → controller → turnover_loss → backward outside graph capture, gated by RL_BAND_ENABLED_INDEX > 0.5 so Phase 3D bit-equality is preserved when the master gate is OFF. ISV slots added (809-812): turnover_ema, controller_boot_done, turnover_target_adaptive, frac_not_masked_observed. The turnover loss kernel now reads slot 811 (controller output) instead of the static slot 803. Bootstrap target 0.05 matches the Phase 4-A static default so first-step behavior is identical. Tests: 4 new band_invariants tests cover the controller bootstrap, the escalate-on-over-trading path, the loosen-on-saturation path, and the backward chain producing non-zero `grad_h_t` from a turnover/target mismatch. All 9 band tests + 18 multi_head_policy regression tests pass. Determinism: `determinism-check.sh --quick` exits 0 for the three relevant configurations (band off, band on, band + multi-head on). The Schulman controllers are deterministic by construction (no PRNG); backward kernels are sole-writer per (batch, slot) so reduce_axis0 delivers bit-equal weight grads across replays. Spec: docs/superpowers/specs/2026-06-03-no-transaction-band-architecture.md §3.1 Option (c), §3.3 (backward chain), §3.4 (loss weight controller), §5 (diag emission), §9.1 Mitigation 3 (resurrection discipline). Pearls applied: * pearl_bootstrap_must_respect_clamp_range — bootstrap 0.05 ∈ [0.01, 0.20] * pearl_dead_signal_resurrection_discipline — asymmetric tighten/loosen rates * pearl_welford_trade_count_is_step_not_trade — bootstrap-done sentinel slot * pearl_determinism_achieved — single-thread controller, deterministic reducers Single-seed local smoke (b=128, 2000+500 steps, FOXHUNT_BAND_ENABLED=1) runs end-to-end without NaN. The band saturates to `frac_masked = 1.0` because position state stays at 0 throughout — the sigmoid surrogate gradient is zero deep inside the band, so the band's upper boundary cannot be pulled toward 0 even though the controller drives the adaptive target to its MAX (0.20). Spec §9.1 Mitigation 2 (`RL_BAND_MAX_MASK_FRAC_INDEX` exploration-slot bypass in rl_band_mask.cu) was declared but not shipped in Phase 4-A; it is the structural fix for this chicken-and-egg trap. Phase 4-B mechanism itself is verified correct via the four new tests; the saturation is an architectural-grafting gap surfaced for Phase 4-A2 follow-up rather than a Phase 4-B implementation defect. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
|
e41a732081 |
feat(ml-alpha): Phase 4-A — No-transaction-band foundation (ISV slot 799 gate)
Foundation of the no-transaction-band architectural turnover regulator
described in
docs/superpowers/specs/2026-06-03-no-transaction-band-architecture.md.
Davis-Norman (1990) / Imaki-Imajo-Ito (2021, arXiv:2103.01775) — when the
current position lies inside a learned band [b_l, b_u], the architectural
default is "do nothing", complementing Phase 3D's reward-side fixes which
the literature (Goodhart-Skalse 2024) bounds the effectiveness of.
Scope: ISV slots 799-808 + BandHead forward + ±|tanh|·N_max_eff
activation + rl_band_mask action override + rl_band_turnover_loss
(Option b, fixed-target) + 5 GPU-oracle invariants + diag emission.
Master gate `RL_BAND_ENABLED_INDEX` (slot 799) bootstraps to 0.0 (OFF)
so the foundation preserves bit-equality with Phase 3D `bd811a774` until
operator opt-in via `FOXHUNT_BAND_ENABLED=1`.
Adaptive controller, encoder backward chain, and cluster verification
are Phase 4-B / 4-C, not in scope here.
ISV slots (799-808, bumps RL_SLOTS_END to 809):
799 RL_BAND_ENABLED_INDEX (master gate, bootstrap 0.0)
800 RL_BAND_LOWER_INIT_INDEX (b_l init, -0.5 in tanh space)
801 RL_BAND_UPPER_INIT_INDEX (b_u init, +0.5 in tanh space)
802 RL_BAND_LOSS_WEIGHT_INDEX (λ_turnover, 0.01)
803 RL_BAND_TURNOVER_TARGET_INDEX (target frac unmasked, 0.05)
804 RL_BAND_GRAD_SHARPNESS_INDEX (sigmoid surrogate, 4.0)
805 RL_BAND_FLAT_RECENTER_RATE_INDEX (reserved for 4-B controller)
806 RL_BAND_WIDTH_MIN_INDEX (collapse detection, 0.1)
807 RL_BAND_WIDTH_MAX_INDEX (saturation detection, 1.8)
808 RL_BAND_DIAG_LAUNCH_EVERY_INDEX (diag cadence, 1.0)
CUDA kernels (3 new):
rl_band_head_forward.cu — 2-stage forward: linear projection
(tree-reduce over HIDDEN_DIM, matches
ppo_policy_logits_fwd shape) + asymmetric
±|tanh|·N_max_eff activation enforcing
b_l ≤ 0 ≤ b_u (Davis-Norman invariant).
rl_band_mask.cu — overrides actions[b]→Hold when
position_lots[b] ∈ [b_l, b_u]. Master-
gated at slot 799; no atomicAdd; runs
OUTSIDE graph capture per spec §9.5.
rl_band_turnover_loss.cu — Option (b) turnover regularizer with
sigmoid surrogate. Per-batch loss +
per-batch grad on (b_l, b_u). Phase 4-A
wires kernel + test; full encoder grad
fold is Phase 4-B.
Rust glue:
crates/ml-alpha/src/rl/band_head.rs — BandHead struct, forward,
launch_mask, launch_turnover_loss.
Xavier-0.01 init under
scoped_init_seed(seed+0xBA_5EED).
trainer/integrated.rs — BandHead field + construction
(bias init = atanh(0.5)) +
ISV bootstrap row + FOXHUNT_BAND_
ENABLED override + forward call
at both step_with_lobsim and
step_with_lobsim_gpu_body sites
+ mask launch BEFORE confidence
gate + per-step band aggregate
diag (lower/upper/width means,
frac_in_band, frac_masked,
collapse_warning).
Tests (5 GPU-oracle invariants, all PASS):
band_activation_clamps_correctly — b_l ≤ 0 ≤ b_u, |·| ≤ N_max_eff
band_mask_forces_hold_when_in_band — pos 0 ∈ [-4,+4] → action becomes Hold
band_mask_passes_through_when_out_of_band — pos 5 ∉ [-1,+1] → action unchanged
band_turnover_loss_correct — loss + grad match analytical form
band_disabled_means_no_mask — slot 799 = 0.0 → mask no-op
Verification:
cargo build --release --example alpha_rl_train: exit 0
cargo test band_invariants --release -- --ignored: 5/5 PASS
cargo test multi_head_policy_invariants -- --ignored: 18/18 PASS (regression)
determinism-check.sh --quick (band OFF): exit 0
FOXHUNT_BAND_ENABLED=1 determinism-check.sh --quick: exit 0
FOXHUNT_USE_MULTI_HEAD_POLICY=1 determinism-check.sh --quick: exit 0
FOXHUNT_USE_MULTI_HEAD_POLICY=1 FOXHUNT_BAND_ENABLED=1 …: exit 0
Phase 4-A local mid-smoke (b=128, 2000+500 steps, band enabled):
exit 0, no NaN observed
frac_masked = 1.0 throughout (band wide at [-4,+4] init)
total_trades = 0 (vs Phase 3D 11,767) — primary kill criterion PASS
band width drifts 8.02 → 7.58 over 2000 steps (slight narrowing from
encoder shared-h_t shift; band-head's own backward chain is Phase 4-B)
G_no_band_collapse FAILS (frac_masked saturates at 1.0) — expected at
Phase 4-A because turnover-loss gradient is not yet folded into the
encoder (per spec §7.1 / §9.1 Mitigation 1, which requires Phase 4-B
adaptive controller + backward chain).
Pearls:
pearl_bootstrap_must_respect_clamp_range (every bootstrap ∈ clamp)
pearl_scoped_init_seed_for_reproducibility (band-head init guard)
pearl_determinism_achieved (no PRNG / no atomicAdd in kernels)
pearl_foxhunt_pi_trained_by_q_distillation_not_ppo (band is structural,
not reward-side — sidesteps Goodhart-Skalse attenuation)
pearl_fleet_fraction_not_aggregate (frac_in_band / frac_masked emitted)
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
|
bd811a7748 |
feat(ml-alpha): Phase 3D — three-intervention overtrading fix (A+B+C)
Combined atomic attack on foxhunt's structural overtrading pathology
per omnisearch (2026-06-03) RL HFT literature: foxhunt uniquely uses
Q-distill as the SOLE policy-training mechanism, making it susceptible
to the four-stage Q→π attenuation chain (Goodhart-Skalse 2024) that
blinds π to small persistent fees. Three concurrent fixes attack
different layers:
A. Hold-action logit bias (+log(4) ≈ 1.386 on action 2)
* crates/ml-alpha/src/rl/ppo.rs PolicyHead::new
* crates/ml-alpha/src/rl/multi_head_policy.rs all K heads + build_priors
Counter-balances the structural 4:1 open-vs-hold action prior (4
open variants 0,1,5,6 vs 1 Hold=2). Pre-bias P(open)=36% / P(hold)=9%;
post-bias P(hold)≈29% / P(any open)≈7%. Mid-smoke seed=42 confirms
Hold rises to 71/128 = 55.5% by step 1999.
B. Quadratic-in-trade-size impact-aware cost (Cao et al. 2026,
arXiv:2603.29086 §4)
* crates/ml-alpha/cuda/rl_fused_reward_pipeline.cu Phase 1.5
* 3 new ISV slots 794-796 (α=0.5, β=2.0, enabled=1.0)
cost = α·|Δlots| + β·(Δlots)². 1-lot=2.5; 4-lot flip=34; 8-lot=132
(superlinear). Applied BEFORE shaping so the surfer-scaffold weight
does not amplify or mute. Trail actions (7,8) are no-ops in the
position kernel → cost=0 for them as expected.
C. PPO surrogate gradient restoration with adaptive blend (Cao 2026 §4)
* crates/ml-alpha/cuda/rl_pi_grad_blend.cu (new — element-wise
scale-or-zero operator)
* crates/ml-alpha/src/trainer/integrated.rs Step 7 (π gradient blend)
* 2 new ISV slots 797-798 (weight=0.005, enabled=1.0)
Previously π was trained ONLY by Q-distillation (line 5850 header).
Now: pi_grad = w_ppo·grad_PPO + grad_Q_distill + grad_SAC_entropy.
Restores the direct fee-aware policy-gradient channel that Q-distill
alone cannot transmit. Blend kernel runs BETWEEN surrogate_backward
and rl_q_pi_distill_grad (which uses +=).
Diag emission (E):
* crates/ml-alpha/src/trainer/integrated.rs rewards.{quadratic_cost_alpha,
quadratic_cost_beta, quadratic_cost_enabled, ppo_surrogate_weight,
ppo_surrogate_enabled}
Tests (D):
* multi_head_policy_invariants: updated k1_reduces_to_single_head for
Phase 3D-A bias; new phase_3d_a_hold_bias_propagates_all_heads
invariant verifies Hold dominance in every head at h_t=0. 18/18 pass.
* phase_3d_blend_kernel_invariants (new): 3 GPU-oracle invariants on
rl_pi_grad_blend (disabled-zeros, enabled-scales-linearly, weight-0-
equivalent-to-disabled). 3/3 pass.
* reward_alignment_invariants: 2 new tests (phase_3d_diag_emission +
phase_3d_quadratic_cost_visible_in_rewards) + fix to the existing
surfer_scaffold test (relaxed bootstrap check for Phase 3B-Y pure-pnl
mode default 0.0; absorbs eval drain row). 3/3 pass.
* All 68 ml-alpha lib tests pass.
Verification:
* SQLX_OFFLINE=true cargo build --release --example alpha_rl_train -p
ml-alpha: exit 0
* SQLX_OFFLINE=true cargo check --workspace: exit 0
* ./scripts/determinism-check.sh --quick: DETERMINISTIC (200/200 rows
bit-equal across two same-seed runs)
* FOXHUNT_USE_MULTI_HEAD_POLICY=1 ./scripts/determinism-check.sh --quick:
DETERMINISTIC
* Local Tier 1.5 mid-smoke (seed=42, b=128, 2000 train + 500 eval):
exit 0, completed_clean=true, no NaN, no abort.
Primary kill criterion (total_trades final < 5,000): NOT MET.
Result: 11,767 trades vs 14,691 baseline = 20% reduction. Cao 2026
forecast 96% reduction for pure-PPO/SAC architectures was not
achieved — foxhunt's Q-distill dominance (q_pi_agree_ema = 0.948 in
this run) attenuates the PPO surrogate's fee signal even with the
blend operator. The behavioral signature IS present (Hold dominance
rises from baseline ~36% structural prior to 55.5% at step 1999;
action_entropy = 1.748 within healthy [1.2, 2.04] target).
Regression guard (eval pnl ≥ -$5M): MARGINAL PASS at -$4.96M (-$36k
inside threshold). The Tier 1.5 verdict flags KILL on Pearson +
wr_train + wr_eval + eval_pnl. Pre-cluster, the literature
recommendation is to A/B-ablate each intervention (slots 794-798
individually gated). Mid-smoke architecturally validates that the
three interventions PROPAGATE and do not crash; cluster b=1024 with
longer runs (20k steps) will surface whether the 20% reduction
compounds into a viable policy.
Architectural references:
* pearl_foxhunt_pi_trained_by_q_distillation_not_ppo
* pearl_reward_signal_anti_aligned_with_pnl
* pearl_bootstrap_must_respect_clamp_range
* feedback_no_atomicadd / feedback_no_htod_htoh_only_mapped_pinned
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
|
17e453a1d5 |
fix(ml-alpha): gate rl_popart_v_correct by slot 793 (Phase 3B-Y companion fix)
Phase 3B-Y gated PopArt reward normalization (slot 793 = 0) but left the V-correction kernel running unconditionally. Mid-cluster diagnostic (alpha-rl-k9lz9 step 77, ~80 min into Phase 3C): catastrophic pnl bleed (-$6.3M cum) driven by V-head unit mismatch. ## Mechanism * `rl_popart_normalize` running mean/sigma stats keep updating even when the rewards-write path is gated (so re-enabling produces sensible values mid-session). With rewards in raw pnl-units, sigma grew to ~2,666 in 77 cluster steps. * `rl_popart_v_correct` unconditionally applies v_pred = (sigma_old/sigma_new)·v_pred + (mean_old-mean_new)/sigma_new every step — scaling V outputs against a running statistic that no longer matches the V regression target. * Net effect: V predictions scaled toward zero while V regression target was raw pnl-units → broken V → broken Bellman Q-target → broken π via Q-distill → catastrophic policy. ## Fix Add a short-circuit at the top of rl_popart_v_correct: ```c if (isv[RL_POPART_NORMALIZE_ENABLED_INDEX] <= 0.5f) return; ``` When normalization is off (slot 793 = 0), V is being trained against raw returns and needs no correction. The popart calibration machinery is bypassed end-to-end. ## Verification * Build clean (1m46s incremental). * FOXHUNT_USE_MULTI_HEAD_POLICY=0/1 ./scripts/determinism-check.sh --quick: both exit 0 (200 rows match, rel-tol=1e-05). ## Discovery sequence * Phase 2A-E cluster ran for 2h53m with shaped_reward=+$201M / pnl=-$330M (Phase 5 shaping anti-aligned) * Phase 3B+Y disabled Phase 5 shaping + PopArt normalize + hindsight * Phase 3C cluster (b=1024) at step 77 showed pnl=-$6.3M already * Diagnosis: V-correction running against stale sigma created unit mismatch that broke V regression chain * This fix is the natural companion to Phase 3B-Y's gate Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
|
093feac3da |
feat(ml-alpha): Phase 3B+Y — reward-pnl alignment foundation
Three coordinated changes to remove anti-aligned components from the reward pipeline. Predicted Pearson trajectory based on Phase 3A audit was wrong (PopArt + hindsight only delivered +0.04 not +0.4), but the deep Phase 3B-Z audit revealed the "low Pearson" was largely a measurement artifact: raw_reward tracks ALL PnL events (including fees + partial closes via pos.realized_pnl_delta) while diag's realized_pnl_usd_cum tracks only full closures. The actual training reward IS aligned with total realized PnL evolution. ## Changes * Slot 753 RL_SURFER_SCAFFOLD_WEIGHT_INDEX bootstrap 1.0 → 0.0 (Phase 5 hold-bonus/entry-cost/short-hold-penalty/long-ride-bonus shaping disabled — was the dominant anti-aligned contaminant per Phase 3A audit and cluster 2A-E evidence of +$201M shaped vs -$330M actual). * New ISV slot 793 RL_POPART_NORMALIZE_ENABLED_INDEX bootstrap 0.0 (PopArt batch normalization disabled — was found to sign-flip rewards under adverse batch composition). * Hindsight injection (rl_hindsight_inject.cu wants_inject) extended with scaffold_w > 0.5 condition (when slot 753 = 0, hindsight off). Semantically pairs the two training scaffolds. ## Verification * 13/13 multi_head_policy_invariants PASS (no regression). * FOXHUNT_USE_MULTI_HEAD_POLICY=0/1 ./scripts/determinism-check.sh --quick: both exit 0. * Local single-seed mid-smoke (b=128, seed 42): - eval pnl: -$4.6M (vs Phase 0 baseline -$4.46M, within noise) - 14,691 trades training / 3,268 eval (healthy, not paralyzed) - action_entropy 1.82 at step 1999 (below uniform 2.40, learning) - 0 NaN, exit 0 ## Pearson interpretation (Phase 3B-Z audit) The Pearson 0.38 between raw_reward and pnl_step_close_usd_delta is a structural mismatch between two valid PnL views (all events vs close events only), NOT an anti-alignment. raw_reward at slot 753 = 0 is the delta of pos.realized_pnl which already aggregates the events we care about for total-USD optimization. Cluster verification (Phase 3C, next) will reveal whether the removal of Phase 5 + popart + hindsight produces a base policy that can learn at scale. The 2A-E cluster catastrophe (-$330M) was driven by Phase 5 shaping that's now disabled. Plan: docs/superpowers/plans/2026-06-03-reward-alignment-gae- combined.md Linked pearls: * pearl_reward_signal_anti_aligned_with_pnl * pearl_phase5_term_4_is_almost_potential * pearl_popart_blind_to_session_signals * pearl_pure_pnl_mode_starves_b16_controllers Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
|
c5036af030 |
feat(infra): plumb FOXHUNT_USE_MULTI_HEAD_POLICY through alpha-rl template
Adds use-multi-head-policy workflow parameter (default "0") wired into the container env block as FOXHUNT_USE_MULTI_HEAD_POLICY. Adds --use-multi-head-policy flag to argo-alpha-rl.sh that passes through to the workflow. Enables Phase 2A-E cluster verification with the multi-head policy mixture (commits |
||
|
|
c8c81ab7e4 |
feat(ml-alpha): Phase 2A-D B1.3 — adaptive gate LR controller
Replaces the fixed 5.0 bootstrap of slot 790 with a real adaptive
controller. ISV slot 790 is now driven by signal observation per
foxhunt-discipline ("no hardcoded constants that should be ISV-
driven"). The constant value remains the bootstrap; the controller
takes over once the signal observation begins.
## Mechanism
* 2 new ISV slots (793 total now):
- 791 RL_POLICY_GATE_ENTROPY_EMA_INDEX (Wiener-α=0.02 on slot 781)
- 792 RL_POLICY_GATE_CONTROLLER_BOOTSTRAP_DONE_INDEX (0→1 latch)
* New kernel rl_gate_lr_multiplier_controller.cu — single-thread
launch (1,1,1)/(1,1,1) matching rl_surfer_scaffold_controller
pattern. No atomicAdd.
* Control law (per pearl_bootstrap_must_respect_clamp_range +
pearl_dead_signal_resurrection_discipline):
if !boot_done: ema = entropy_now; boot_done = 1
else: ema = (1-α)·ema + α·entropy_now
over = 0.85·log(K) (~0.934 for K=3)
collapse = 0.20·log(K) (~0.220 for K=3)
if ema > over: mult *= 1.003 (escalate +0.3%/step)
elif ema < collapse: mult *= 0.95 (decay -5%/step, emergency)
mult = clamp(mult, 1.0, 50.0)
* Per-step launch wired flag-gated AFTER MultiHeadPolicy::
emit_diag_stats (writes slot 781) and BEFORE next-step host-side
lr_gate = lr_pi * isv[790] read. Captured inside the flag-on
CUDA graph; flag-off graph never includes the launch.
* Bootstrap: slot 790=5.0 (B1's empirically-validated starting
point), 791=0.0 (overwritten on first observation), 792=0.0.
* Diag emit: 3 new fields gate_lr_multiplier, gate_entropy_ema,
gate_controller_bootstrap_done — inside the existing flag-gated
multi_head_policy JSON block (flag-off schema unchanged).
## Test results
* 17/17 multi_head_policy_invariants PASS (13 pre-existing + 4 new):
- bootstrap_initializes_ema_to_first_observation
- escalates_when_entropy_above_threshold (5.0 → 6.747 in 100
steps; matches analytical 5·1.003^100 within ±0.1)
- decays_when_entropy_below_collapse (5.0 → 1.0 in 100 steps,
clamped to MIN_FLOOR)
- clamps_at_min_floor_and_max_ceiling (50.0 ceiling, 1.0 floor,
long-escalate 5.0 → 50.0 over 800 steps)
* Flag-off determinism: exit 0, byte-equal to pre-B1.3 modulo
elapsed_s (0/200 non-elapsed_s diffs).
* Flag-on determinism: exit 0 (verified 3 consecutive runs).
Note: one transient first-run FAIL surfaced during verification
(mamba2_l2_out step 0 Δ≈4773) that disappeared after clean
rebuild — same build-cache pattern documented in Phase 2A-C
reports. Three confirmed clean post-rebuild.
* Pre-commit: 0 atomicAdd, 0 raw memcpy_htod/dtoh, 0 TODO.
## Single-seed seed-42 trajectory (b=128, 2000 train steps)
step entropy_ema multiplier notes
0 1.0986 5.015 bootstrap
100 1.0972 6.75 escalating
500 1.0967 22.36
1000 1.0944 50.00 hit MAX_CEIL
1500 1.0820 50.00
1999 1.0830 50.00 stuck at ceiling
Final gate_probs_mean = [0.253, 0.383, 0.365] — close to fixed 5×
endpoint [0.252, 0.382, 0.366]. The controller saturated at the
50× ceiling because entropy_ema never crossed below the 0.934
threshold — but the gate end-state matched 5× regardless. This
is the controller correctly diagnosing "no value of multiplier in
[1, 50] is sufficient to specialize the gate at b=128" — the
binding constraint is scale (signal/noise ratio), not LR. Cluster
b=1024 / 20k steps provides ~80× more gradient signal and is the
right next test.
## Linked
* Phase 2A-A:
|
||
|
|
6f3639bfbf |
feat(ml-alpha): Phase 2A-D B1 — gate LR multiplier (ISV slot 790)
Addresses the gate-not-learning failure mode observed at 3-seed local mid-smoke verification of HEAD |
||
|
|
5f10bcde3a |
feat(ml-alpha): Phase 2A-C+ — device-aggregated gate diag
Adds the specialization-signal diagnostics that Phase 2A-C deferred.
4 new stats are emitted under policy_diagnostic.multi_head_policy.*
when FOXHUNT_USE_MULTI_HEAD_POLICY=1: gate_probs_mean[K],
gate_argmax_mass[K], gate_entropy_mean, per_head_entropy_mean[K].
These are the load-bearing signals for Phase 2A-D's verdict —
without them we'd see a pnl delta but couldn't distinguish "the
mixture genuinely specializes by regime" from "the mixture
accidentally acts as a single-head regularizer".
## ISV slots (25 new)
* 765-772 RL_POLICY_GATE_PROBS_MEAN_BASE (8 slots, MAX_K_HEADS=8 stride)
* 773-780 RL_POLICY_GATE_ARGMAX_MASS_BASE
* 781 RL_POLICY_GATE_ENTROPY_MEAN
* 782-789 RL_POLICY_PER_HEAD_ENTROPY_MEAN_BASE
* RL_SLOTS_END = 790
## Aggregator kernel (multi_head_policy_aggregate_diag.cu)
* Grid = (MAX_K_HEADS + 1, 1, 1) = 9 blocks. Block = (128, 1, 1).
* Per-head blocks (k_block 0..7): if k_block ≥ runtime K, thread 0
writes 0.0 to its two ISV destinations. Else parallel-sum over B
in shared memory, tree-reduce by halving stride, thread 0 writes
gate_probs_mean[k] + per_head_entropy_mean[k].
* Global block (k_block = MAX_K_HEADS): single block computes gate
entropy + argmax-mass in one pass. Per-thread argmax counter
array in registers scatters to s_am[MAX_K_HEADS][BLOCK_THREADS]
shared mem; tree-reduce; thread 0 writes 9 ISV scalars.
* No-atomicAdd: every ISV destination has a single writer thread.
Tree-reduce via shared memory + __syncthreads(). Deterministic
fixed-order pairwise sum.
## Trainer wiring
* MultiHeadPolicy::emit_diag_stats(isv_dev_ptr) launches the
aggregator on the struct stream. Called at all 3 train-path
forward sites in integrated.rs immediately after mhp.forward()
(gate_probs and pi_probs_k are forward outputs — must aggregate
before next step's forward overwrites them).
* build_diag_value emits the 4 fields inside the existing
multi_head_policy.* block under if self.use_multi_head_policy.
Flag-off schema unchanged (multi_head_policy key absent).
## Verification (all gates pass)
* 13/13 invariants: 11 pre-existing + 2 new
- aggregate_diag_gate_probs_mean_correct: uniform 1/K + asymmetric
ramp case, max abs err < 1e-5
- aggregate_diag_entropy_pins_top_and_bottom: top (uniform gate
→ entropy = log(K) = 1.0986; uniform pi → per-head = log(11) =
2.398); bottom (one-hot gate → 0; one-hot pi → 0). Tie-broken-
low argmax verified.
* FOXHUNT_USE_MULTI_HEAD_POLICY=0 determinism-check.sh --quick:
exit 0. Flag-off diag.jsonl preserved (multi_head_policy key
ABSENT, EXPECTED_LEAVES=712 unchanged).
* FOXHUNT_USE_MULTI_HEAD_POLICY=1 determinism-check.sh --quick:
exit 0. Aggregator is deterministic.
* 200-step flag-on smoke shows the diag working as designed:
- gate_probs_mean ≈ [0.328, 0.339, 0.333, 0,0,0,0,0] sum ≈ 1.0
- gate_argmax_mass ≈ [0, 1.0, 0, ...] (Head 1 Long-bias dominant
at init — matches init bias asymmetry)
- gate_entropy_mean ≈ 1.0985 (at max log(3) ≈ 1.0986 — gate has
NOT collapsed)
- per_head_entropy_mean ≈ [2.36, 2.28, 2.29] vs max 2.398 —
heads slightly less than uniform, expected at random init
* Pre-commit: 0 atomicAdd, 0 raw memcpy_htod/dtoh, 0 TODO.
## Surprises handled
None — all STOP-on-surprise predictions held:
1. Slot bootstrap zero-init contributes 0² to isv_state checksum,
so unconditional bootstrap preserves flag-off bit-equality
without needing a flag-gated block (unlike the 2A-C ISV
surprise).
2. EXPECTED_LEAVES (712) preserved — flag-off schema unchanged.
3. K consistency (Rust MAX_K_HEADS=8 + CUDA #define) maintained.
4. emit_diag_stats placed after mhp.forward() and before any
downstream consumer that would overwrite forward outputs.
5. CUDA graph capture: aggregator launch is deterministic (fixed
grid/block, fixed reductions). Replay single-path. No conflict.
## Linked
* Phase 2A-A:
|
||
|
|
3e36f4a0e6 |
feat(ml-alpha): Phase 2A-C — trainer integration behind FOXHUNT_USE_MULTI_HEAD_POLICY flag
Wires MultiHeadPolicy into the live trainer's π forward + backward + Adam path, gated by FOXHUNT_USE_MULTI_HEAD_POLICY (default off). Mirrors the FOXHUNT_USE_ROLLOUT precedent (commit |
||
|
|
e22da61cf8 |
feat(ml-alpha): Phase 2A-B — MultiHeadPolicy backward + aux KL prior
Backward pass through the K-head mixture + per-head KL prior
regularization. Components still inert — trainer integration is
Phase 2A-C.
* multi_head_policy_backward.cu: two kernels.
- backward_pi (grid=(B), block=(HIDDEN_DIM=128)): full chain rule
log_grad → mixture decomposition → per-head softmax-Jacobian
→ W_heads/b_heads/h_t grads. Per-batch grad scratch +
reduce_axis0 reduction; no atomicAdd. Sole-writer per (b,k,a,h).
- backward_gate (grid=(B), block=(K=8)): gating softmax-Jacobian +
W_gate/b_gate grad scratch + grad_regime_h (computed but NOT
routed upstream — regime_h is loader-precomputed).
* multi_head_policy_aux_prior.cu: per-head KL prior softmax-Jacobian
grad β·π·(log_ratio − KL) additively accumulated into
grad_pi_logits_k. Read-modify-write race-free (sequential same
stream after backward_pi).
* MultiHeadPolicy::backward(grad_pi_logits, h_t, regime_h, isv) —
launches backward_pi → aux_prior → backward_gate → 4× reduce_axis0
on the same stream. Priors are computed at new() as softmax of
the per-head init bias vectors, so KL=0 at initialization and
the aux term only fires once Adam moves W_heads off zero.
* 4 new GPU-oracle invariant tests (9/9 total):
- backward_gradcheck_w_heads (W_heads central difference)
- backward_gradcheck_w_gate (W_gate central difference)
- aux_prior_grad_sums_to_zero_per_head (softmax-Jacobian invariant)
- backward_is_deterministic_across_contexts
* 2 diagnostic tests (no asserts, print-only):
- backward_gradcheck_w_heads_eps_sweep
- backward_gradcheck_w_gate_eps_sweep
## Math investigation (gradcheck error chase)
Initial gradcheck at ε=1e-3 showed 1.7–1.9% relative error which
SHOULD have meant a chain-rule bug. Investigation:
1. Re-derived the analytical backward chain from first principles
(8 chain steps) and diff'd against actual kernel source line
by line. No discrepancies — math is implemented exactly as
derived. ε in log-divisor (1e-12) matches forward.
2. ε-sweep characterization (b=2, k=2):
| ε | max_rel_err W_heads | max_rel_err W_gate |
|------|---------------------|--------------------|
| 1e-2 | 1.6e-3 | 5.0e-3 |
| 5e-3 | 1.3e-3 | 1.8e-2 |
| 1e-3 | 1.9e-2 | 3.3e-2 |
| 5e-4 | 1.0e-2 | 4.3e-2 |
| 1e-4 | 1.0e-1 | 4.7e-1 |
Error INCREASES as ε shrinks — opposite of O(ε²) truncation.
Unambiguous fingerprint of fp32 catastrophic cancellation in
(L_plus − L_minus)/2ε with |L| ≈ 6.6, |grad| ≈ 1e-2:
noise ≈ ulp(L)/(2ε·|grad|) ≈ 2⁻²³·6.6/(2ε·0.01)
ε=1e-3 → 4% (matches 1.9% observed), ε=1e-2 → 0.4%, ε=1e-4 → 40%.
3. Max-error indices random (not init-bias positions, not saturated
softmax) — consistent with noise scatter, not systematic bug.
Verdict: TRUNCATION_CONFIRMED (specifically fp32 cancellation, not
chain-rule truncation). Math is correct.
Fix: ε bumped 1e-3 → 1e-2 (the cancellation/truncation sweet
spot for fp32 chained softmax). Tolerance restored to foxhunt-
standard 1e-2 / 1e-3 (was 5e-2 / 5e-4 — that was a workaround for
the noise-dominated ε, not justified). Measured error now
1.6e-3 / 5.0e-3 — comfortably under tolerance. ε-sweep diagnostic
tests added as a permanent reproducibility anchor so future
gradcheck regressions are easy to diagnose.
## Verification
* `cargo test multi_head_policy_invariants --release -- --ignored`:
11/11 PASS (9 invariants + 2 diagnostics).
* `./scripts/determinism-check.sh --quick`: exit 0
(pipeline unchanged — backward not yet wired).
* 0 atomicAdd, 0 raw memcpy_htod/dtoh, 0 TODO/FIXME.
## Linked
* Phase 2A-A:
|
||
|
|
0b3e401500 |
feat(ml-alpha): Phase 2A-A — MultiHeadPolicy foundation (inert)
K=3 policy mixture + regime-gated routing head, with gate reading
REGIME_DIM=6 features DIRECTLY (bypassing the VSN softmax bottleneck
per the regime-attenuation empirical finding). Components are
unconditional but not yet wired into the trainer — Phase 2A-B
adds backward + aux KL prior, 2A-C wires Q-distill grad routing
to the mixture.
* 4 new ISV slots 761-764 (K, gating entropy floor, head entropy
floor, aux prior β) + RL_SLOTS_END 761→765.
* multi_head_policy_forward.cu: per-batch K-head logits + mixture
combination. Grid=(B), Block=(N_ACTIONS=11). Persists pi_logits_k
+ pi_probs_k for backward.
* multi_head_policy_gate_forward.cu: per-batch K-thread softmax
over W·regime + b. Reads regime_h directly (parallel channel —
bypass VSN). Stores pre-softmax logits to gmem BEFORE the
in-place exp (caught during impl — plan pseudocode would have
corrupted gate_logits).
* MultiHeadPolicy struct with Option A asymmetry-break init:
Head 0 → ShortLarge bias (+0.5), Head 1 → LongSmall+LongLarge
bias (+0.5 each), Head 2 → Hold bias (+0.5). Indices verified
against rl/common.rs:56-70 N_ACTIONS=11 enum. htod via local
upload() helper using MappedF32Buffer + raw_memcpy_dtod_async
(mirrors rl/ppo.rs, rl/dueling_q.rs).
* 5 GPU-oracle invariant tests, all PASS:
- gate_probs_sum_to_one
- pi_probs_mixture_sums_to_one
- k1_reduces_to_single_head (bit-equal vs reference Linear)
- gate_responds_to_regime_change (TVD > 0.5, modal head flips)
- forward_is_deterministic_across_contexts (bit-equal across
two fresh CUDA contexts)
* Determinism preserved: ./scripts/determinism-check.sh --quick
exits 0 (pipeline unchanged, components inert).
* No memcpy_htod/memcpy_dtoh on regular slices, no atomicAdd, no
scoped-init-seed bypass.
Empirical motivation (3-seed mid-smoke at HEAD
|
||
|
|
779c03b9d7 |
feat(ml-alpha): Phase 1B-B — rollout collection loop + GAE behind FOXHUNT_USE_ROLLOUT flag
Spec: docs/superpowers/specs/2026-06-02-trainer-rollout-buffer-gae.md Plan: docs/superpowers/plans/2026-06-02-trainer-rollout-buffer-gae-implementation.md §Phase 1B-B Builds on commit |
||
|
|
969caf26c3 |
docs(ml-alpha): spec + plan for Phase 1B trainer rollout-buffer + GAE refactor
Companion docs for commit |
||
|
|
5cd2f87039 |
feat(ml-alpha): Phase 1B-A — RolloutBuffer + GAE kernel foundation
Spec: docs/superpowers/specs/2026-06-02-trainer-rollout-buffer-gae.md
Plan: docs/superpowers/plans/2026-06-02-trainer-rollout-buffer-gae-implementation.md
Foundation phase of the trainer rollout-buffer + GAE refactor (Phase 1B).
Empirically required after Phase 1A regression confirmed math-agent's
atomicity claim: dropping Phase 5 shaping WITHOUT GAE makes things worse
(eval_pnl regressed -$691k from -$4.46M to -$5.15M, popart sigma CV
0.96 -> 1.58, sign agreement 60% -> 46%). See pearl_reward_signal_
anti_aligned_with_pnl ADDENDUM 2026-06-02d for the mechanism analysis.
This commit establishes the components WITHOUT trainer integration —
subsequent phases (1B-B through 1B-E) wire them into the actual training
loop. Subdivides the multi-week refactor into atomically-verifiable
phases per feedback_investigation_first_falsification_methodology.
New components:
* 3 ISV slots (758-760):
RL_PPO_ROLLOUT_HORIZON_INDEX = 758 (T_rollout, bootstrap 256)
RL_PPO_N_EPOCHS_INDEX = 759 (K_ppo, bootstrap 4)
RL_PPO_N_MINIBATCHES_INDEX = 760 (minibatches/epoch, bootstrap 8)
RL_SLOTS_END 757 -> 761
* crates/ml-alpha/cuda/gae_backward_sweep.cu (58 LOC):
Single-thread-per-batch sequential backward sweep computing
A_t = δ_t + γλ·A_{t+1}·(1-done), returns_t = A_t + V_t.
Deterministic by construction (no parallel reductions, no atomicAdd,
no nvrtc). Reset on done. v_T_bootstrap parameter for trajectory-end
V estimate.
* crates/ml-alpha/src/trainer/rollout_buffer.rs (210 LOC):
RolloutBuffer struct with [B × T] device buffers for rewards, dones,
v_t, actions, log_pi_old, h_t; plus [B] v_T_bootstrap; plus output
advantages, returns. compute_gae(γ, λ) invokes the kernel using
cudarc raw-pointer pattern (`device_ptr(stream).0` resolved into
local before .arg() — idiomatic for codebase). Loaded module +
kernel handle stay private; struct fields exposed per spec.
* crates/ml-alpha/tests/rollout_buffer_invariants.rs (466 LOC):
5 GPU-oracle invariant tests (#[ignore = "requires CUDA"]):
1. gae_terminal_only_matches_close_event_pnl — single done at T-1
2. gae_dense_reward_geometric_decay — Σ(γλ)^k geometric series
3. gae_done_resets_credit — done reset propagation
4. gae_deterministic_across_runs — bit-equality across contexts
5. rollout_buffer_alloc_sizes_match_spec — alloc verification
All 5 PASS in 2.15s.
Validation gates (Phase 1B-A complete when all pass):
* cargo build --release --example alpha_rl_train -p ml-alpha: exit 0
(1m 00s release build, gae_backward_sweep.cubin compiled with -O3
--use_fast_math --ftz=true --fmad=true for sm_86)
* cargo test -p ml-alpha --test rollout_buffer_invariants --release: 5/5 PASS
* ./scripts/determinism-check.sh --quick: exit 0 — DETERMINISTIC: all
checksums.* leaves match across all 200 rows (rel-tol=1e-5, abs-tol=1e-7).
Pipeline output bit-equal to baseline since new struct/kernel are
loaded but unused in the training loop.
* Pre-commit hook on staged diff: PASS (0 memcpy_dtoh, 0 atomicAdd).
Next: Phase 1B-B (rollout collection loop behind FOXHUNT_USE_ROLLOUT env
flag) per plan §Phase 1B-B. Dispatch after this commit lands.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
|
cfaa420bd4 |
fix(ml-alpha): Phase 0 — per-step authoritative USD pnl in diag + fix mislabeled fields
The diag's `trading.*_pnl_*_usd` fields were systematically mislabeled, making
the eval verdict signal untrustworthy. Three quantities shared the `_usd`
suffix but were not in USD:
- `eval_summary.total_pnl_usd` (USD ground truth, $50/pt ES applied)
- `trading.realized_pnl_cum_usd` (cumulative SHAPED reward at close events;
pts × lots × Phase-5 shaping; NOT USD despite the suffix)
- `trading.pnl_cum_usd` (mathematically nonsense — reads rewards_d after
apply_reward_scale AND rl_popart_normalize whiten it in-place; dividing
by current_scale un-does only the first transform)
At baseline mid-smoke (b=128, 2000+500, seed=42): eval_summary.total_pnl_usd
= -$4,457,625 while diag.trading.realized_pnl_cum_usd = +$469k and
diag.trading.pnl_cum_usd = +$12.9k. 346x ratio, sign-opposite. Same
mechanism as `pearl_grwwh_eval_catastrophic_collapse` ($234M cluster
discrepancy). Every Pearson(reward, Δpnl_usd) computation against these
fields was shaped-vs-shaped tautology, not pnl alignment.
This commit:
DELETE `trading.pnl_cum_usd` (mathematically nonsense post-popart).
RENAME `trading.realized_pnl_cum_usd` -> `trading.shaped_reward_close_event_cum`.
Truthful name; the underlying values are post-Phase-5 shaped reward
summed at close events, useful for gradient-signal diagnostics but never
to be confused with USD pnl.
ADD `trading.realized_pnl_usd_delta` and `trading.realized_pnl_usd_cum` in
diag.jsonl / eval_diag.jsonl. Source: new per-batch float buffer
pnl_step_close_usd_d on LobSimCuda, zeroed via raw_memset_d8_zero before
each pnl_track_step launch, written by pnl_track.cu's close branch as
realised_pnl_usd_fp / 100.0 (same arithmetic as eval_summary's
TradeRecord aggregator, applying ES $50/pt). Single-writer per block —
no atomicAdd per feedback_no_atomicadd.
Direct readback via read_slice_d_into<f32> from sim.pnl_step_close_usd_d()
after step_with_lobsim_gpu returns (main stream already synchronized via
self.stream.synchronize() at pnl_track_step exit). NOT routed through
diag_staging double-buffer — a prior implementation tried this and
broke determinism (cross-stream race between main-stream
raw_memset_d8_zero + pnl_track_step write and diag-stream
raw_memcpy_dtod_async read). Direct same-stream read avoids the race
entirely and stays mega-graph compatible (the readback runs AFTER
per-step pipeline finishes, outside any captured CUDA graph).
Cumulative tracked Rust-side in alpha_rl_train.rs; reset at train->eval
boundary so realized_pnl_usd_cum tracks the eval window only.
scripts/tier1_5_verdict.py upgrade:
- signal_reward_alignment reads trading.realized_pnl_usd_cum
- signal_eval_pnl falls back to realized_pnl_usd_cum
- signal_consistency upgraded from WARN-at-5% to KILL-at-$50 with a new
consistency_kill_usd threshold key; cross-source disagreement is now
a hard kill not a warning
tests/eval_diag_emission.rs: bumped EXPECTED_LEAVES 711 -> 712 (-1 deleted,
-0 renamed, +2 added); added invariant that cum == running_sum(delta) and
regression check that legacy pnl_cum_usd / realized_pnl_cum_usd fields are
absent from the schema; allowed n_eval_steps + 1 row count for the
drain-row pattern.
Falsification gate (load-bearing, must hold every smoke from this commit
forward): `diag.last_eval_row.trading.realized_pnl_usd_cum` matches
`eval_summary.total_pnl_usd` within $50.
Validation (mid-smoke b=128, 2000 train + 500 eval, seed=42, RTX 3050):
- cargo build --release --example alpha_rl_train -p ml-alpha: exit 0
- cargo test -p ml-alpha --test eval_diag_emission: PASS (712 leaves)
- local-mid-smoke.sh: exit 0, drain row at step 2500
- cross-source consistency: eval_summary=-$4,457,625.00 vs
diag.realized_pnl_usd_cum=-$4,457,625.18 -> delta=$0.18 (exactly fp/100
f32 rounding noise; well within $50 tolerance)
- ./scripts/determinism-check.sh --quick: PASS — all checksums.* leaves
match across all 200 rows (rel-tol 1e-5, abs-tol 1e-7)
- Pre-commit hook on staged diff: PASS
Memory pearls created in this session document the diagnostic chain:
- pearl_diag_pnl_fields_are_shaped_reward_not_usd (the mislabeling root)
- pearl_reward_signal_anti_aligned_with_pnl ADDENDUM 2026-06-02b
(correction to the 2026-06-02 ADDENDUM — the "+0.93 Pearson at HEAD"
finding was a measurement artifact because both sides were in shaped
pts x lots units, not USD)
- pearl_advantage_kernel_is_done_gated_td_not_gae (related signal-density
gap — Phase 1.5 follow-up)
- pearl_phase5_term_4_is_almost_potential (Ng-Harada-Russell analysis of
Phase 5 shaping terms — Phase 1 follow-up)
Unlocks: TRUE Pearson(rewards.sum, Δpnl_usd) can now be measured per-step
at HEAD. The eval-collapse investigation (next phases: Ng-Harada-Russell
shaping cleanup + GAE upgrade) is now falsifiable with bit-deterministic
verdicts grounded in authoritative USD pnl.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
|
629ebd667c |
feat(ml-alpha): deterministic same-seed training + Tier 1.5 fast-dev-cycle
Two same-seed runs now produce bit-equal eval_summary.json, alpha_rl_train_summary.json,
and diag.jsonl (modulo wall-clock elapsed_s). The 5-phase falsification chain landed:
Phase 2 PER tree-rebuild: __threadfence is NOT a grid-wide barrier; multiple blocks
raced across sum-tree levels. Fix: Grid=(1) Block=(1024) + __syncthreads
in rl_per_tree_rebuild.cu.
Phase 2.3 cuBLAS GEMM_DFALT + TF32 default-math allowed split-K non-deterministic
accumulation at 3 sites. New crates/ml-alpha/src/cublas_determinism.rs
applies CUBLAS_PEDANTIC_MATH via FOXHUNT_DETERMINISTIC env toggle
(0=TF32 prod, 1=PEDANTIC dev default, 2=DEFAULT_MATH control).
Phase 2.6 Two bugs surfaced sequentially in the backward kernel chain:
(1) rl_iqn_tau_cos_features had a multi-block r/w race on prng_state[batch]
— all N_TAU=32 blocks read seed; only tau_idx==0 wrote back; no
inter-block barrier. Fix: split into READ-ONLY rl_iqn_tau_cos_features
+ new sibling rl_iqn_advance_prng_state launched on same stream
(kernel-launch ordering = grid-wide barrier).
(2) OutcomeHead::new called near_zero_xavier without scoped_init_seed,
falling back to time+thread-id RNG. Stayed dormant until first done
event activated non-sentinel labels and divergent weights flowed via
grad_h_t_outcome into encoder gradient. Fix: add seed param + install
scoped_init_seed(dqn_seed.wrapping_add(0x0CE0)) guard.
Validation (./scripts/determinism-check.sh --quick, RTX 3050, b=128, 200+50 steps):
- All 200 rows of checksums.* leaves match (rel-tol 1e-5, abs-tol 1e-7)
- eval_summary.json, alpha_rl_train_summary.json byte-equal between runs
- diag.jsonl byte-equal modulo elapsed_s
- Eval pnl identical run-A vs run-B at seed 42
Pre-fix baseline (Phase 2.5 measurement): same-seed eval pnl spread $450k
($187k vs -$261k). Post-fix: $0 spread.
Speed cost: ~1.5ms/step amortised; ~10-15% slower than TF32 production
(PEDANTIC tax — acceptable in dev, toggle to FOXHUNT_DETERMINISTIC=0 for prod).
Mapped-pinned discipline: all 11 NEW memcpy_dtoh sites in diagnostic dump methods
+ per-step checksum readback use a new pub(crate) helper
read_slice_d_into<T: Copy>(stream, src, dst) — MappedRecordBuffer + raw
memcpy_dtod_async + raw_stream_sync + volatile read. Generic over T (f32, f64,
i32, u32, u8). Satisfies feedback_no_htod_htoh_only_mapped_pinned + hook guard.
Bundled Tier 1.5 fast-dev-cycle infrastructure (spec
docs/superpowers/specs/2026-06-02-fast-dev-cycle.md):
- scripts/local-mid-smoke.sh b=128, 2000+500, ~10min on RTX 3050
- scripts/determinism-check.sh runs mid-smoke twice, diffs checksums
- scripts/tier1_5_verdict.py behavioral kill verdict
- AdamW checkpoint save/load (crates/ml-alpha/src/trainer/optim.rs)
- IntegratedTrainer checkpoint save/load (resume from checkpoint)
- 15 Phase 1 checksum leaves in build_diag_value
- Env-gated dump methods (FOXHUNT_DETERMINISM_DEBUG_PER/MAMBA2/RL/BACKWARD)
for future divergence-chasing — never run in production
Documentation:
- docs/superpowers/specs/2026-06-02-determinism-foundation.md
- docs/superpowers/specs/2026-06-02-fast-dev-cycle.md
- docs/superpowers/plans/2026-06-02-determinism-foundation-implementation.md
- docs/superpowers/notes/2026-06-02-determinism-phase{1,2,2.2,2.5,2.6}-*.md
- Adjacent specs/plans/notes from the analytical chain that surfaced determinism
as the load-bearing blocker (eval-summary, eval-boundary, regime-observer,
multi-head policy, regime-invariance, Phase 3 IQN-complement post-mortem)
Unlocks: every controller / architecture / reward-shaping A/B from this commit
onward attributes outcome differences to the change, not random-init kernel-race
drift cascading through training x eval LOB-sim trajectories. The eval-collapse
investigation (pearl_reward_signal_anti_aligned_with_pnl, multi-head spec,
regime-invariance spec) is now testable with trustworthy verdicts.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
|
63fc16f173 |
fix(rl): surfer-scaffold v5.2 — drop wr_excess clamp, full scaffold when novice
alpha-rl-mjsmr (
|
||
|
|
4a4953b01f |
fix(rl): surfer-scaffold v5.1 — decay re-engagement requires EXCESS alertedness
alpha-rl-zf6s5 (
|
||
|
|
38a4aa15b3 |
feat(rl): adaptive surfer-scaffold reward shaping (spec v5 A+B combined)
Diagnosis from alpha-rl-8fb55 ( |
||
|
|
fa347e4812 |
feat(rl): reward-policy alignment — pure-pnl mode default (spec 2026-06-01)
Empirical diagnosis (local analysis of two 20k+5k cluster runs on PVC, SHAs |
||
|
|
b45c2fb108 |
fix(rl): edge-decay PH — flip sign for downside detection
Canonical Page-Hinkley as cited in spec/pearl uses m = Σ(x − μ̄ − δ), which detects mean INCREASE not decrease. First cluster smoke alpha-rl-n5x87 exposed the sign error: at train step 1277 the detector showed 77% fleet alerted (ph_mean=237.8 vs λ=5) — but training-time improvement shouldn't trigger the decay detector. The current formula was firing on "recent x > long-run μ̄" (improvement) rather than "x < μ̄" (decay). Fix: m = Σ(μ̄_pre − x_t − δ). Now grows when x_t < μ̄ − δ (signal below mean by more than tolerance = degradation). M still tracks min(m) and ph_stat = m − M still triggers when cumulative deviation rises above recent minimum. Local b=16 smoke (100+50) post-fix: train[99] ph_mean=40, alert=20%, warmup=0.69 — detector firing on early-training noise. Eval ph_mean=0 (too few closes for warmup). The cluster scale will give the real test. Also renamed kernel param `ph_M_per_batch` → `ph_mmin_per_batch` to match Rust snake_case field naming. (Was a name mismatch between Rust struct and CUDA kernel param that snuck through the v1 build because both sides were edited consistently within their own languages but the cross-language contract wasn't.) Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
|
34806b6b62 |
diag(rl): edge-decay detector Phase 1 — Page-Hinkley on per-trade EV
Pure observability. No behavior change. Implements the missing
"short-run trust adjustment" layer between Kelly long-run sizing and
CMDP tail kill, per pearl_edge_decay_detection_is_a_missing_abstraction_layer.
The 64 commits since
|
||
|
|
a5e0d00794 |
diag(rl): CMDP fleet-fraction observability — 4 new ISV slots
The existing CMDP diag emits AGGREGATES (worst_consec, max_cool_remain,
mean session_pnl) that, at b=1024, hide per-batch behavior. The aggregate
"worst" stays at the trip limit even when only a single batch is
constrained — which misled today's investigation into thinking the fleet
was in a perma-trap when local b=16 smoke shows only 31% of batches in
cooldown during train, 0% during eval.
This commit adds 4 new ISV slots that count what FRACTION of the fleet is
constrained at each step, computed in the existing rl_cmdp_constraints_check
per-batch loop. Pure observability — no behavior change.
ISV slots (RL_SLOTS_END 743 → 747):
- RL_CMDP_FRAC_IN_COOLDOWN_INDEX = 743: cooldown_remaining > 0
- RL_CMDP_FRAC_CONSEC_NEAR_LIMIT_INDEX = 744: consec >= limit - 1
- RL_CMDP_FRAC_DD_TRIGGERED_INDEX = 745: dd_triggered flag set
- RL_CMDP_FRAC_SESSION_NEG_INDEX = 746: session_pnl < 0
Diag emit: risk_stack.cmdp.frac_{in_cooldown, consec_near_limit,
dd_triggered, session_neg}. EXPECTED_LEAVES 671 → 675.
Bootstrap array 230 → 234 (all 4 default 0.0; kernel overwrites every step).
Predetermined falsification criteria (cluster smoke b=1024):
- frac_in_cooldown > 0.5 sustained → cooldown trap is real, ship a
resurrection fix to rl_cmdp_constraints_check.
- frac_in_cooldown < 0.1 throughout → cooldown is fine, look elsewhere
(overfit / training scale / data signal).
- 0.1-0.5 → ambiguous, refine.
Local b=16 smoke (100+50): leaves=675, schema parity, kernel writes
fractions correctly. Train peaks at frac_in_cooldown=0.31 / frac_dd=0.31.
Eval stays at 0.0 throughout. Aggregate `cooldown_remaining_steps=495`
at train[99] coexists with frac_in_cooldown=0.31, confirming the aggregate
hides the fleet-level reality.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
|
f428be794b |
Revert "feat(rl): B-11-β Q-distill informativeness gate"
This reverts commit
|
||
|
|
b93971726d |
feat(rl): B-11-β Q-distill informativeness gate
First behavior change since B-7 (preceding B-8/B-9/B-10 were
observability-only). Attenuates the distill gradient when softmax(Q/τ)
is near-uniform — when target_entropy → ln(N_ACTIONS), the distill term
is pure max-entropy regularization with no informational signal, so it
fights PPO instead of carrying Q's preferences into π.
Verified cluster cause @ alpha-rl-88f5c (B-10 smoke,
|
||
|
|
c1dc84a345 |
fix(rl): B-10 wire G1 Q-distribution launch into step_with_lobsim_gpu
Initial B-10 commit placed `launch_q_distribution_stats` after the `dqn_head.forward(h_t)` at line 6615 — but that's inside `step_with_lobsim`, the legacy non-GPU path. The actual GPU path exercised by `alpha_rl_train` is `step_with_lobsim_gpu` (line 8402), which has its own h_t-based Q forward at line 8632. The G1 diag fields (q_dist_entropy_mean, q_value_range_mean, q_value_abs_max) returned 0 locally because the launch never ran in the GPU path. Added the same launch_q_distribution_stats call after line 8632. The non-GPU launch at 6615 is left in place — `step_with_lobsim` is still callable; per `feedback_no_partial_refactor.md` both paths are instrumented consistently. Validated locally (RTX 3050 Ti, 200+50 b=16 fold-1): q_dist_entropy_mean = 2.6482 (was 0) q_value_range_mean = 4.9257 (was 0) q_value_abs_max = 3.8715 (was 0) The entropy value 2.65 is consistent with a near-uniform softmax over Q_N_ATOMS=21 atoms with some learned concentration (ln(21)=3.04 = full uniform, 0 = one-hot delta). Range and abs_max in 4-5 unit scale match the locally-adapted atom support. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
|
8c7ce02da9 |
feat(rl): B-10 policy-quality cascade diagnostic
alpha-rl-8gtk2 (B-7+B-8+B-9 run at SHA
|
||
|
|
033906f213 |
docs(rl): fix stale C51 V_MAX/V_MIN "ratchet" claims + B-9 test bug
While alpha-rl-8gtk2 (the B-7+B-8+B-9 20k+5k cluster run) was training, B-9's saturation diag exposed V_MAX_eff dropping from 856.82 (step 482) to 464.61 (step 817) within the same run. The kernel docstrings and slot doc-comments uniformly claimed "ratchet (monotone-grow)" semantics — contradicting the observation by 392 units. Root cause: wwcsz followup 2026-05-24 (commit landed in rl_reward_clamp_controller.cu Step 5 only) REPLACED the original ratchet with a slow symmetric EWMA (α=0.001, half-life ~700 steps) on `win_bound`/`loss_bound`. Static ratchet was wasting atom resolution on rare tails (avg rewards in [-5,+5] with span at [-60,+20] → Δz=4, Q couldn't distinguish "slightly winning" from "slightly losing"). The EWMA refocuses atom resolution on the ACTIVE range. The controller's own header was correctly updated at the time. The documentation drift was in 3 other places — fixed here: - bellman_target_projection.cu header (lines 47-49): "ratchet (monotone-grow)" → accurate EWMA description with cross-references + pearl_c51_v_max_freeze_required_for_surfer warning (V_MAX in 100-200 → trend-follower; past 1000 → degraded). - rl_atom_support_update.cu line 4: "Companion to the C51 atom-span ratchet" → "Companion to the C51 atom-span EWMA". - isv_slots.rs slot allocation table line 43: "C51 atom-span ratchet slots" → "C51 atom-span EWMA slots (α=0.001)". - isv_slots.rs RL_C51_V_MAX_INDEX / RL_C51_V_MIN_INDEX doc-comments (lines 649-668): replaced with accurate EWMA description, observed 857 → 465 drop example, and asymmetric tracking note (V_MIN_eff EWMAs -loss_bound NOT -V_MAX_eff — they only coincide when ratio≈1). Latent test bug also surfaced + fixed: c51_atom_saturation_diagnostic assumed `V_MIN_eff = -V_MAX_eff` (line 136). With observed Kelly EMAs avg_loss=$436 vs avg_win=$872 → ratio ≈ 0.5 → V_MIN_eff ≈ -0.5 × V_MAX_eff, the test's bot-rate consistency check would false-fail if saturation ever became non-zero. Dormant so far (sat=0% in both smoke and full run). Relaxed the bot-rate assertion to use the v_bound_floor (-1.0) as the necessary lower bound — correct for any asymmetric ratio and catches gross drift without false-failing on legitimate adaptation. No behavior change; pure docstring drift + dormant test bug repair per feedback_trust_code_not_docs. Compile clean. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
|
87a8259c6e |
fix(build): ml-backtesting honors CUDA_COMPUTE_CAP for sm_90 on H100
alpha-rl-mbg2n on H100 failed at LobSimCuda::new with `CUDA_ERROR_NO_BINARY_FOR_GPU` loading book_update.cubin. Root cause: ml-backtesting/build.rs read only `FOXHUNT_CUDA_ARCH` (set by lob-backtest-sweep-template) but NOT `CUDA_COMPUTE_CAP` (set by alpha-rl-template from `nvidia-smi --query-gpu=compute_cap` inside the compile pod). On H100 it silently fell through to default sm_86; the sm_86 cubin contains no PTX → no JIT path on sm_90 device. The docstring claimed "Mirrors crates/ml-alpha/build.rs" — it didn't (ml-alpha reads CUDA_COMPUTE_CAP). Completing the mirror now: detect_arch returns sm_<NN> honoring (in order): 1. CUDA_COMPUTE_CAP numeric env (alpha-rl-template) 2. FOXHUNT_CUDA_ARCH sm_-prefixed env (lob-backtest-sweep-template) 3. nvidia-smi --query-gpu=compute_cap at build time 4. Default sm_86 (RTX 3050 Ti local dev) Both production argo templates now produce sm_90 cubins on H100 and sm_89 on L40S without further changes. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
|
29b5acad55 |
feat(rl): B-9 C51 Bellman-target saturation observability
Under B-7 (clamp default-disabled), per-transition rewards in bellman_target_projection / bellman_fused_select_project can exceed the adaptive atom support [V_MIN_eff, V_MAX_eff]; each t_z overshoot is silently clamped before being mapped onto the discrete support. B-9 publishes the per-step saturation rate + pre-clamp t_z extremes so we can decide if the atom span has become a bottleneck — without re-attempting the reverted Fix F atom-widening (pearl_c51_v_max_freeze _required_for_surfer). Changes: - 4 new ISV slots (726-729): top/bot saturation rate + max/min pre-proj. - bellman_target_projection.cu: both entry points (bellman_target_projection AND bellman_fused_select_project per feedback_no_partial_refactor) gain 4 new [B] f32 pointer params; thread-0 sequential reduction over Q_N_ATOMS=21 (odd count rules out symmetric tree-reduce — matches the kernel's existing softmax max/sum pattern at lines 157/171). - New cross-batch reducer cuda/rl_bellman_target_saturation_reduce.cu: single block, grid-stride gather + power-of-2 tree reduce, no atomicAdd per feedback_no_atomicadd. - dqn.rs: load reducer cubin, add saturation_reduce_fn handle, launch_saturation_reduce method, 4 scratch pointer params on both bellman methods. - integrated.rs: allocate 4 [B] f32 scratch buffers; pass through both fused_select_and_project_bellman call sites + launch reducer after each. 4 new bootstrap entries (213 → 217 fixed-size array). - build.rs: register new kernel. - 4 new diag leaves under risk_stack.atom_calibration.target_*. Comment distinguishes them from popart.max_abs_reward_ema (different signal: Bellman target = r + γ·atom_value, can exceed reward by γ·V_MAX_eff). - EXPECTED_LEAVES 653 → 657. - tests/c51_atom_saturation_diagnostic.rs: GPU-oracle test asserts 4 invariants over 249 rows — rates in [0,1], max≥min, rate>0 ⇒ overshoot exists, top+bot ≤ 1. Validation: - 200+50 b=16 fold-1 smoke clean. Locally V_MAX_eff adapts to ~19.3 (atom support is ISV-driven via rl_atom_support_update), so saturation is 0% in the smoke; both diag leaves emit + invariants hold. - popart_disaggregation_invariants: still passes (249 rows, identity). - eval_diag_emission: train = eval = 657 leaves. - Determinism preserved: kernel adds shared-mem reduction over per-thread t_z values that were already computed; target_dist output unchanged. Spec: docs/superpowers/specs/2026-06-01-b9-c51-atom-saturation-diagnostic.md Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
|
1739d9c173 |
feat(rl): B-8 popart σ_welford disaggregation + identity invariant
Under B-7 (clamp default-disabled), popart's Welford state now updates against unclamped magnitudes; σ_effective = max(σ_welford, env.max) can spike from either source but diag only emitted the combined value. This blocks attribution of any future σ shocks. Changes: - RL_POPART_SIGMA_WELFORD_INDEX (slot 725): Welford-only σ, BEFORE the F4 envelope floor at rl_popart_normalize.cu:156. Pure observability — no computation change. - rl_popart_normalize.cu: insert one ISV write between σ_welford computation (line 141) and envelope-floor application (line 156). Mirrors the existing #define-local-then-write pattern (POPART_SIGMA_INDEX). - build_diag_value: new leaf popart.sigma_welford in the canonical popart block. NOT duplicating max_abs_reward_ema (already emitted at risk_stack.regime.popart_envelope.max_abs_reward_ema per feedback_single_source_of_truth_no_duplicates). - EXPECTED_LEAVES 652 → 653. - Bootstrap array [(usize, f32); 212] → 213 with sentinel 0.0 (overwritten every step by popart kernel). - tests/popart_disaggregation_invariants.rs: GPU-oracle test asserts identity popart.sigma == max(σ_welford, env.max) across 249 rows (200 train + 50 eval), skipping step 0 bootstrap. - tests/eval_diag_emission.rs: migrate fold-idx 0/n_folds 2 → 1/3 (the n_folds=2 split picks the first 4 files which don't satisfy loader's 1033-snapshot minimum after test_data grew from 2 → 9 files). Validation: - popart_disaggregation_invariants passes locally (249 rows OK). - eval_diag_emission passes locally: train=653 eval=653 leaves. - B-9 spec at docs/superpowers/specs/2026-06-01-b9-c51-atom-saturation-diagnostic.md builds on slots 726-729 next (uncommitted, awaiting implementation). Spec: docs/superpowers/specs/2026-06-01-b8-popart-calibration-observability-and-floor.md Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
|
55d049ecf4 |
feat(rl): B-7 ISV-toggle reward clamp (default disabled)
apply_reward_scale.cu's post-scale [-3,+1] clamp was masking ~99.99% of realized tail magnitudes from the agent's reward signal: local b=16 smoke showed train pnl_cum_usd −$0.63 (clamp-truncated) vs realized_pnl_cum_usd −$8,574.63 (raw raw_rewards sum), a 13,600× compression. Per van Hasselt 2016, popart standardization + F4 envelope are designed to handle tail magnitudes; the clamp fights them. Changes: - RL_REWARD_CLAMP_ENABLED_INDEX (slot 724): default 0 (disabled). When 0 the kernel skips the asymmetric clamp; scaled rewards pass through to rewards[b] unchanged. Legacy behavior restored by setting to 1. - DiagInputs.realized_pnl_cum_usd: parallel counter computed from raw_rewards (pre-scale, pre-clamp shaped pnl). Compare against trading.pnl_cum_usd to surface clamp-truncation gaps. - Trainer accumulates realized_pnl_cum_usd in both train + eval loops per closed-trade done-step (same pattern as pnl_cum_usd). - EXPECTED_LEAVES 651 → 652 for the new diag leaf. Validation: - 200+100 b=16 fold-1 smoke clean; train leaves = eval leaves = 652; realized_pnl_cum_usd diverges from pnl_cum_usd as expected when clamp disabled (the reward-hacking gap is now observable). - compute-sanitizer pending (cluster). B-8 (popart σ_welford disaggregation) and B-9 (C51 Bellman-target saturation observability) specs at docs/superpowers/specs/2026-06-01-* build on this slot allocation (725, 726-729 next). Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
|
912f33c6fc |
test(rl): B-6 invariant regression test — asymmetric Wiener-α / Bayesian shrinkage
GPU-oracle test (per `feedback_no_cpu_test_fallbacks`) validating: 1. Cold-start EMA values match spec defaults (avg_w=1, avg_l=1, wr_ema=0.5 from B-3 cold_start bootstrap) 2. ISV slot 721/722/723 defaults exposed in diag (α_slow_min=0.001, n_full_threshold=30000, cv_gain=1.0) 3. Asymmetric direction holds at cold-start: avg_loss EMA grows ~50× faster than avg_win EMA per equivalent observation (because α_fast/α_slow_min = 0.05/0.001 = 50). Verified locally: avg_l=$425 vs avg_w=$4.36 at train_end (100× empirical ratio — matches expected math under volatile batch=16 data) 4. Boundary reset: at eval[1], avg_w=avg_l=1.0 and wr_ema=0.5 (cold_start resumed via reset_session_state) Test result locally: cold-start: avg_w=1 avg_l=1 wr_ema=0.5 ISV slots: alpha_slow_min=0.001 trust_full=30000 cv_gain=1 train_end: avg_w=4.36 avg_l=425.39 dones=131 eval[1]: avg_w=1.00 avg_l=1.00 wr_ema=0.500 dones=0 TEST PASS Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
|
bc9eaac89d |
fix(rl): B-6 — ISV-driven adaptive asymmetric Wiener-α (Bayesian shrinkage)
B-5 (asymmetric α with static α_slow=0.001) revealed the static parameter problem: provably bounds cascades (avg_win peak $2k vs B-4's $40k) BUT over-conservative in train (dckcc step 800: avg_l > avg_w → Kelly says don't trade → model can't discover edges; wr_ema crashed to 0.145). The fundamental tension: static α_slow can't satisfy BOTH - Train convergence: asymmetry must FADE so model learns from real data - Boundary safety: asymmetry must ENGAGE at every fold to prevent cascade B-6 RESOLVES this via Bayesian shrinkage: trust(n) = min(1, cum_dones / n_full_threshold) [Phase 1] stability = exp(-CV × cv_gain) [Phase 2] trust_eff = trust(n) × stability α_slow_eff = α_slow_min + (α_fast − α_slow_min) × trust_eff Phase 1 (data-quantity): trust grows with cum_dones; reset_session_state zeroes cum_dones → asymmetry RESUMES at every boundary. Math: at n=0 α_slow_eff = α_slow_min = 0.001 (full skepticism). At n=n_full = 30k trades: α_slow_eff = α_fast = 0.05 (full standard Wiener). Phase 2 (data-quality): Welford CV of reward magnitude gates trust. Stable signal (CV→0): stability=1, trust opens normally. Volatile signal (CV high): stability→0, asymmetry persists. cv_gain=0 disables Phase 2. Per-EMA asymmetry direction encodes Kelly safety semantics: avg_win: slow-up (skeptical of wins), fast-down avg_loss: fast-up (admit losses), slow-down (slow forget) wr_ema: slow-up (skeptical of high WR), fast-down ISV slots (all signal-derived from cum_dones + Welford): 721 RL_EMA_ALPHA_SLOW_MIN_INDEX = 0.001 722 RL_EMA_TRUST_FULL_THRESHOLD_INDEX = 30000 723 RL_EMA_CV_GAIN_INDEX = 1.0 Threshold calibration (n_full=30k): - Train: ~1000 cluster steps for trust to fully open → asymmetry active during cold-start (first 30 steps, cascade prevention) then fades. - Eval: 30 dones/step × 500 eval steps = 15k dones → trust climbs to 0.5 by eval end → partial protection throughout eval. Composes: - B-3 Kelly fractional-trust (Kelly SIZING gated by cum_dones) - B-6 EMA asymmetric-α (Kelly INPUTS biased conservative by cum_dones) Both fade as data accumulates; both reset at boundary. Spec: docs/superpowers/specs/2026-06-01-ema-asymmetric-trust-with-cv-gain.md Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
|
d725b77031 |
fix(rl): B-5 — asymmetric Wiener-α replaces B-4 caps (math gap fix)
B-4 multiplicative cap (1.5/step) and additive cap (0.05) failed math gap: 1. Multiplicative cap: 1.5^N grows unboundedly over many steps. zh96b confirmed avg_win peak still $40,911 (B-4) vs $44,821 (B-3) — only 9% reduction at peak despite 12× reduction at early steps. 2. Additive cap on wr_ema: 0.05 > natural Wiener step at α=0.05 (max 0.035 from p=0.3 to obs=1.0). Cap NEVER engages → no-op. Fundamental math: per-step rate-caps CANNOT bound EMA convergence to E[X]. For sustained observations, ema → X regardless of any per-step rate-cap (EMA's natural property). B-5 ASYMMETRIC WIENER-α — provably biased estimator for Kelly safety: avg_win: alpha = (step_avg > prev) ? α_slow : α_fast avg_loss: alpha = (step_avg > prev) ? α_fast : α_slow // mirror wr_ema: alpha = (step_wr > prev) ? α_slow : α_fast With α_slow=0.001, α_fast=0.05: EMA_N (sustained X in slow direction) = X × (1 - 0.999^N) N=100: ema = 9.5% × X N=200: 18% N=500: 39% N=1000: 63% For avg_win: $40k sustained → ema reaches only $3,800 by step 100 (vs B-4's $40k peak). Provably biased low → Kelly's b̂ underestimated → smaller f_safe by construction. For wr_ema: 100% wr sustained from prev=0.3 → reaches 0.87 in N=694 steps (vs prior cascade in 140 steps). Asymmetry direction chosen per-EMA for Kelly safety: - avg_win: slow up (skeptical of wins), fast down (correct quickly) - avg_loss: fast up (admit losses), slow down (slow to forget pessimism) - wr_ema: slow up (skeptical of high WR), fast down Single new ISV slot: RL_EMA_ALPHA_SLOW_INDEX = 0.001. Replaces B-4's RL_WR_EMA_MAX_DELTA_INDEX (same slot 721, repurposed). Removed: B-4 cap logic from both kernels (multiplicative + additive). Single uniform asymmetric-alpha rule. Cleaner than B-4 patchwork. Math validation prediction: avg_win peak should be ≤ $5k (vs B-4's $40k); wr_ema peak should be ≤ 0.50 (vs B-4's 0.88). Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
|
cbe4869375 |
fix(rl): B-4 — extend Winsorization rate-cap to Kelly input EMAs
Deep JSONL analysis of x56wn revealed the B-2/B-3 cold-start fixes still left Kelly inputs (avg_win/loss, wr_ema) vulnerable to Wiener-α cascade: - avg_win_ema spiked to $44,820 by step 219 (vs current $1.6k stable) - wr_ema oscillated 0.32 ↔ 0.60 across 500-step windows - Per-step wr stable 0.29-0.33 but EMA admitted 27-percentage-point swings Root cause: B-3 fixed the cold_start=1.0 anti-pattern, but Wiener-α=0.05 still admits 5% of arbitrarily large observations into the EMA. From cold_start=1, a single $45k tail-win lifts ema by 5% × $45k = $2250 in ONE step. Heavy-tailed magnitude → unbounded EMA variance. B-4 fix — apply Winsorization rate-caps with EMA-type-appropriate form: ISSUE A — avg_win/loss EMAs (heavy-tailed positive magnitudes) ============================================================== Math: Hoeffding bound requires bounded support; heavy-tailed sample mean is unboundedly biased. Winsorize observations at c × prev. Same multiplicative rate-cap as pos_max_ema (B-2). REUSES slot 718 — single source of truth for "magnitude EMA growth cap". ema_raw = (1-α) × prev + α × step_avg ema_new = min(ema_raw, prev × growth_cap) // growth_cap = isv[718] At growth_cap=1.225 (B-2 default): adapts from cold_start=1.0 to 10000× in ~50 steps. Single $45k observation now caps at $1.225 first step. ISSUE B — wr_ema (proportion [0,1]) ==================================== Multiplicative cap wrong for bounded proportion. Use ADDITIVE cap: |ema_new - prev| ≤ max_delta (default 0.05) Math (Hoeffding): at batch=1024, σ_binomial = √(p(1-p)/b) ≈ 0.014 for p=0.3. Cap 0.05 = 3.5σ → P(admit | IID) ≈ 1.2%. Steady-state EMA noise std ≈ 0.004 (α=0.05) so cap = 12σ_EMA — never noise-triggered. ISIV-driven: new slot 721 RL_WR_EMA_MAX_DELTA_INDEX = 0.05 (tunable). Also REMOVED first-observation bootstrap branch in win_rate_ema_update (was: if prev==SENTINEL → ema = step_wr direct). SENTINEL=0.5 is a conservative neutral; Wiener-α blend from it is safe. Generalized principle (deeper than B-2/B-3): every adaptive EMA that gates a magnitude-sensitive downstream controller needs a rate-limit on per-step change. Form depends on domain: - Unbounded ℝ⁺ (magnitudes): multiplicative cap (Winsorize at c × prev) - Bounded proportion [0,1]: additive cap (|Δ| ≤ ε) - Signed unbounded: additive cap scaled by EMA magnitude Validation: cargo check clean. Will validate at cluster against x56wn (same SHA except B-4 additions) for direct comparison. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
|
42b4898239 |
fix(rl): B-3 — Kelly fractional-trust schedule + avg_win/loss cold-start
alpha-rl-4xmxm eval analysis revealed the residual -$100M loss was OVERCONFIDENT
SIZING, not σ-explosion. At eval[1]: wr_ema=0.575, avg_win=$1180, avg_loss=$117
(b=10:1), kelly_fraction=1.0 (warmup gate). The controller was sizing at 53% of
capital based on n=1 sample — Hoeffding ε at n=1 is √(ln(40)/2) = 1.36, so the
empirical win-rate has 95% CI half-width of 100%+. Kelly is mathematically
unidentified from this sample.
ISSUE A — binary warmup gate at f=1.0
=====================================
rl_kelly_fraction_controller.cu:45 + rl_fused_controllers.cu:858:
```
if (cumulative_dones < min_trades) kelly = 1.0; // MAX SIZE during warmup
```
Intent was anti-trade-death (pearl_kelly_trade_stream_death). But forces
maximum Kelly at every fold boundary where parent fix resets cum_dones=0.
ISSUE B — avg_win/loss bootstrap to first observation
=====================================================
rl_avg_win_loss_ema_update.cu:61-65,71-75:
```
if (prev == 0.0f) ema_new = step_avg; // bootstrap = first observation
```
Same anti-pattern as pos_max_ema (B-2). At eval[1] the first trade pair's
INSTANCE ratio (b=10:1) became the EMA, making Kelly's b̂ wildly biased.
B-3 FIX — fractional-trust schedule with bounded floor:
trust(n) = max(f_floor, min(1, n / N_full))
f_safe = max(f_floor·safety, f_kelly · trust · safety)
Where:
N_full = 200 trades (Hoeffding-derived: ε=0.10 at 95% confidence)
f_floor = 0.05 (5% of full Kelly, prevents trade-death)
safety = 0.5 (existing fractional-Kelly multiplier)
At n=0: f_safe = 0.05 × 0.5 = 0.025 (2.5% of capital — 21× smaller than
the prior f=1.0 catastrophe)
At n=N_full: f_safe = f_kelly · safety (asymptotic optimality)
Plus B-2 extension: avg_win, avg_loss cold-start to 1.0 (was 0), in both
`with_controllers_bootstrapped` AND `reset_session_state`. Single uniform
Wiener-α update — no first-observation branch.
New ISV slot 720: RL_KELLY_BOOTSTRAP_FLOOR_INDEX = 0.05.
Mirror change to fused_controllers.cu Kelly block (twice-implementation rule).
Local validation (b=16 800+200 fold-1):
- eval[1]: avg_win=1.0, avg_loss=1.0, wr_ema=0.5 (NEUTRAL — NOT bootstrapped
to first-observation $1180/$117 ratio)
- eval[100]: avg_win=153, avg_loss=12 (controller smoothly adapted via Wiener-α)
- eval pnl: -$169k (similar to prior smokes; local scale doesn't trigger
the b=1024 catastrophe pattern)
Cluster prediction: Kelly's max sizing during eval should be ~2.5% during
first ~200 trades, then asymptote to safety × kelly. Expected eval pnl
improvement vs 4xmxm's -$100M: at least 5×, target 10×.
Spec: docs/superpowers/specs/2026-05-31-pos-max-ema-cold-start-redesign.md
(B-3 follow-up appended; full Hoeffding math in spec body)
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
|
16cf9f260c |
fix(rl): B-2 — pos_max_ema cold-start cascade eliminated, ISV-driven cap
The addendum's rate-cap (B-1) only protected the Wiener-α path; the first-
observation bootstrap branch let cold-start fat-tail events seed pos_max_ema
unbounded. At alpha-rl-4xmxm step 5, a single $947 scaled reward bootstrapped
pos_max_ema=879 directly, cascading through clamp_win → unclamped subsequent
rewards → env.max=11375 by step 37 (1500× the eventual steady-state σ).
B-2 fix per docs/superpowers/specs/2026-05-31-pos-max-ema-cold-start-redesign.md:
1. Bootstrap RL_POS/NEG_SCALED_REWARD_MAX_EMA_INDEX to MIN_WIN=1.0 (was 0)
in `with_controllers_bootstrapped`. Conservative neutral value → clamp_win
starts at MARGIN × 1.0 = 1.5 → rewards heavily clipped until adaptation.
2. Remove the `if (ema_prev == 0.0f) ema_new = pos_max;` branch from
`rl_reward_clamp_controller.cu`. Single uniform update rule (Wiener-α +
rate-cap) applies from cold-start onward. Mirror change for neg_max_ema.
3. Replace `#define POS_MAX_EMA_MAX_GROWTH_PER_STEP 1.5f` with ISV-driven
reads (per feedback_isv_for_adaptive_bounds). Three new slots:
717 RL_POS_MAX_EMA_COLD_START_INDEX = 1.0
718 RL_POS_MAX_EMA_GROWTH_CAP_BASE_INDEX = 1.225 (√1.5 for
twice-per-step inv)
719 RL_POS_MAX_EMA_GROWTH_CAP_CV_GAIN_INDEX = 0.0 (adaptive layer
disabled by default)
4. Adaptive growth_cap from Welford CV of reward magnitude (slot 615-617)
when cv_gain > 0: stable regime → tight cap, volatile regime → loose.
Disabled by default; opt-in via ISV tuning.
Local smoke validation (800+200 fold-1 b=16):
- step 1: pos_ema=1.0 (initialized, NOT bootstrapped from observation)
- step 5: pos_ema=1.0 (no observation yet, sparse-skip working)
- step 25: pos_ema=63.8 (vs 1314 without B-2 — 21× reduction)
- step 37: pos_ema=32.5 (vs 7074 without B-2 — 218× reduction)
- env.max @ step 37: 126 (vs 11375 without B-2 — 90× reduction)
Generalizes the pattern: NORMALIZATION EMAs that gate signal magnitudes
should bootstrap to CONSERVATIVE neutral values, NEVER to first observation.
Cross-references pearl_first_observation_bootstrap which needs revision.
Phase 4 audit (other controllers with same anti-pattern) deferred to
follow-up — see spec §4 Phase 4 for the grep target list.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
|
1aa92f57f0 |
fix(rl): eval-boundary addendum — reward_scale warmed_flag + pos_max_ema rate-cap
The parent eval-boundary fix (
|
||
|
|
72684ed3ef |
fix(rl): preserve normalization EMAs across eval boundary
reset_session_state was resetting four EMAs that calibrate signal scale rather than predict behavior: - RL_POS_SCALED_REWARD_MAX_EMA_INDEX (clamp scale anchor) - RL_NEG_SCALED_REWARD_MAX_EMA_INDEX (clamp scale anchor) - RL_REWARD_CLAMP_CLIP_RATE_EMA_INDEX (clamp feedback) - RL_POPART_MAX_ABS_REWARD_EMA_INDEX (F4 envelope) At the train→eval fold boundary this disabled the reward clamp (cap = pos_max_ema × ratio = 0) and let eval's first ~10 trade outliers blow up the popart envelope σ 1500× (57 → 4471) over ~14 steps. PPO surrogate (A_unnorm = σ × A_norm) ran with catastrophically mis-scaled gradients for the next ~1000 eval steps, accounting for the bulk of the -$185M eval loss at alpha-rl-6kghr fold 1. Refines pearl_adaptive_carryover_discipline: RESET PREDICTIVE EMAs (Kelly, dd, recency) but PRESERVE NORMALIZATION EMAs. Spec: docs/superpowers/specs/2026-05-31-eval-boundary-normalization-preservation-design.md Pearl: pearl_popart_reset_at_eval_boundary_shocks_normalization.md Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
|
22e6ddbcac |
fix(rl): eval_summary aggregates across all b_size accounts (E.4-E.5)
Replaces the broken single-account + scale-mismatch slicing block in
alpha_rl_train.rs eval phase with proper per-account aggregation.
Pre-fix (cluster v11 alpha-rl-8ll7j, b=1024):
head_before_eval = sim.read_total_trade_count() // aggregate=1,203,376
all_records = sim.read_trade_records(0) // backtest 0 only, ≤1024
// 1.2M > 1024 → wrap branch fires → eval_records = ALL 1024 records
// from account 0, mixing train+eval trades
eval_summary: n_trades=1024 pnl=$61,513 wr=0.217 — MEANINGLESS
Post-fix (this commit, b=16 local smoke):
head_before_per_b = sim.read_per_backtest_trade_counts()
all_per_b = sim.read_trade_records_all()
// for each backtest: slice ring by per-account head_before vs head_after,
// handle wrap correctly, aggregate eval-only records
eval_summary: n_trades=226 pnl=$-41,137 wr=0.376
n_eval_trades_seen=226 (= captured) n_dropped=0 n_pre_eval_wrapped=0
Local smoke validation (b=16, 1k train + 250 eval, fold 1 ES.FUT all):
- n_trades jumped 58 → 226 (4× — confirms aggregation across 16 accts)
- per-account pre-eval heads: min=16 max=59 sum=596 (correct scale)
- n_eval_trades_seen == n_trades (no wrap at this scale)
- eval_summary.json gained 4 new fields:
n_eval_trades_seen, n_eval_trades_dropped,
n_pre_eval_trades_wrapped, b_size
Existing keys (n_trades, total_pnl_usd, profit_factor, sharpe_ann,
max_drawdown_usd, win_rate) preserved for downstream G8-gate / Argo
aggregator consumers (per spec §3).
Cluster validation pending Phase A cluster (alpha-rl-jz48s) completion
to avoid alpha_rl_train.rs branch conflict. After Phase A merges,
this fix can submit alongside.
Spec: docs/superpowers/specs/2026-05-31-eval-summary-trade-aggregation-design.md
Plan: docs/superpowers/plans/2026-05-31-eval-summary-trade-aggregation.md
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
|
fa0858f0b1 |
feat(lobsim): per-backtest trade-count read + 4× TRADE_LOG_CAP (E.1-E.3)
Adds two new public methods to LobSimCuda + bumps TRADE_LOG_CAP to prevent eval-phase ring-buffer wrap at cluster scale. - read_per_backtest_trade_counts() -> Vec<u32>: cumulative trade counters per backtest (length n_backtests). Replaces the broken pattern in alpha_rl_train.rs where head_before_eval = aggregate across batch was compared against all_records = single-account ring. - read_trade_records_all() -> Vec<Vec<TradeRecord>>: all backtests' rings in one call. Mapped-pinned staging per feedback_no_htod_htoh_only_mapped_pinned: allocate MappedRecordBuffer<u8> for the payload + MappedRecordBuffer<u32> for heads, DtoD copy from device buffers into mapped-pinned dev_ptrs, sync, read host_ptrs. - TRADE_LOG_CAP 1024 → 4096: cluster v11 (alpha-rl-8ll7j) showed ~342 eval trades/account mean with peaks toward 1000. 4096 gives 4× headroom; memory cost b=1024 × cap × 40 B = 167 MB (was 41 MB), comfortable on L40S 48GB / H100 80GB. Both methods synchronize after DtoD so host reads see the data. E.4 (alpha_rl_train.rs aggregation block replacement) lands separately after Phase A cluster validates to avoid alpha_rl_train.rs conflict. See spec docs/superpowers/specs/2026-05-31-eval-summary-trade-aggregation-design.md and plan docs/superpowers/plans/2026-05-31-eval-summary-trade-aggregation.md. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
|
ad16e9d941 |
fix(trainer): address Phase A code review minors (8 fixes)
Per `feedback_always_fix_minor_review_findings`, all 8 minor review findings from Phase A (eval-diag emission) are applied: 1. Restore aliasing comment for trail_fired_step / conf_gate_step (both read RL_CONF_GATE_FIRED_COUNT_INDEX deliberately). 2. Replace act_hist[7/8/9/10] magic indices with `Action::*` enum variants in both `IntegratedTrainer::build_diag_value` and the train/eval loops in `alpha_rl_train.rs`. 3. Rewrite `recursion_limit = "256"` comment in `lib.rs` to describe macro recursion DEPTH (~30 nested object blocks), not leaf count. 4. Move `RL_CONF_GATE_FIRED_COUNT_INDEX`, `RL_PYRAMID_ADD_COUNT_INDEX`, `RL_FRD_GATE_FIRED_COUNT_INDEX`, `RL_HEAT_CAP_FIRED_COUNT_INDEX` to the top-of-file `use` block (each appears ≥2× as a fully- qualified path); 10 call sites switched to bare names. 5. Harmonise leaf-count documentation to 643 (642 scalars + 1 bool `pyramid.max_units_reached`) across `build_diag_value` docstring, `--eval-diag-jsonl` arg docs, eval-phase comment, and the test module-level docstring. 6. Drop `_host` suffix from every `DiagInputs` field (14 fields). The struct's docstring already states all slices are host-side; the suffix was redundant. `DiagStaging.*_host_ptr` fields are NOT renamed — those still refer to actual host pointers. 28 internal trainer references and 28 call-site references updated in lockstep. 7. Align eval_diag_emission test default data dir with foxhunt convention: `test_data/futures-baseline/ES.FUT` (resolved from `CARGO_MANIFEST_DIR`) instead of `/tmp/rl-smoke-lpi-diag/data`. Switch `--instrument-mode` from `front-month` to `all` to match the all-instrument predecoded files at that path. Env override `FOXHUNT_EVAL_DIAG_DATA` still wins. 8. Add eval-step monotone assertion: verify `eval[0].step == n_steps` and `eval[N-1].step == n_steps + n_eval_steps - 1`. Catches a regression where the eval loop would reuse the training step counter instead of continuing past it. Verification: `SQLX_OFFLINE=true cargo check -p ml-alpha` clean + `cargo test --release -p ml-alpha --test eval_diag_emission -- --ignored --nocapture` PASS with curated data (100 train + 50 eval lines, 643 leaves both phases, schema parity, step axis 100..=149). |
||
|
|
210794626a |
feat(rl): emit eval-phase per-step diag to eval_diag.jsonl
Cluster run alpha-rl-8ll7j ended with +$61,513 pnl and max_dd -$444,512
but the eval phase emitted ZERO per-step diag, leaving the drawdown
trajectory invisible. Phase A of the 2026-05-31 checkpoints+eval-diag
plan wires the eval loop into the same diag pipeline as train using
the builder extracted in the previous commit.
Changes:
* `--eval-diag-jsonl <PATH>` CLI flag (defaults to
`<out>/eval_diag.jsonl`).
* Eval loop now calls `diag_staging.sync_and_swap` +
`snapshot_async` after every `step_with_lobsim_gpu`, builds a
`DiagInputs` from the staging reads, and writes a JSONL line via
the same `IntegratedTrainer::build_diag_value` the train loop
uses. Step indices continue past the train phase
(`cli.n_steps + eval_step`) so post-hoc tooling can concatenate
train + eval JSONL into a monotone step axis.
* Eval-phase running counters (pnl_cum_usd, trades, gates, …) are
independent of train counters so the eval JSONL reflects the
eval window only — mirrors the trade-record checkpoint that
eval_summary.json uses.
* New integration test `eval_diag_emission` validates schema
parity: same 643 leaf paths in `diag.jsonl` and `eval_diag.jsonl`,
correct line counts (n_steps / n_eval_steps). Ignored by default
because it requires CUDA + the pre-built release binary.
Verification (locally on RTX 3050 Ti):
100 train + 50 eval @ b=16, n_folds=2 →
`diff <(head -1 diag.jsonl | jq 'paths(scalars)|sort')
<(head -1 eval_diag.jsonl | jq 'paths(scalars)|sort')`
returns empty (schema parity confirmed).
|
||
|
|
8a93a77adc |
refactor(trainer): extract json!{} into IntegratedTrainer::build_diag_value
The per-step diag JSONL `json!{...}` block had grown to 642 leaf paths
across ~30 nested objects, duplicating every ISV slot read into the
example binary. Phase A of the 2026-05-31 checkpoints+eval-diag plan
extracts it into a single builder on the trainer so the eval phase
can reuse it (next commit) without duplicating the schema.
Changes:
* `IntegratedTrainer::build_diag_value(step, elapsed_s, &DiagInputs)
-> Result<serde_json::Value>` — same 642-leaf schema as before,
bit-equivalent ISV reads (all from `self.isv_host_slice()`).
* `DiagInputs<'a>` struct bundles the host-side per-step state the
trainer doesn't own (DiagStaging reads + running counters +
windowed act histogram), so the call site stays a 1-liner.
* Train loop in `alpha_rl_train.rs` swaps the inline json! for the
builder call; the ~140-slot ISV-imports wall collapses to four
slots still read by the stderr ticker.
* `#![recursion_limit = "256"]` moves from the example into
`ml-alpha/src/lib.rs` since the builder now lives in the library.
Schema parity verified: `head -1 diag.jsonl | jq 'paths(scalars)|sort'`
yields the same 642 keys as before this refactor (no schema drift).
|
||
|
|
13d8ed76da | docs: checkpoints + eval-diag spec + plans (v1 superseded by v2) | ||
|
|
7b3309edcc |
fix(rl): eliminate atomicAdd in PPO+DQN loss reduction (F4.1)
Replaces atomicAdd accumulators in `ppo_clipped_surrogate_fwd` and
`dqn_distributional_q_bwd` with per-batch [B] outputs + a dedicated
single-block tree-reduce kernel. Per feedback_no_atomicadd.
Root cause: `ss_pi_loss_dev_ptr` was zero-init at trainer construction
but never reset between steps; the PPO atomicAdd accumulated across
every step since startup. Result: `loss.pi` = step_count × mean_per_step,
bit-exact step-count fingerprints:
- local 1k × ~9 ≈ 9080 ✓
- cluster 20k × ~1200 ≈ 24M ✓
The DQN distributional Q kernel had the same atomicAdd pattern. Its
trainer caller happened to memset between launches in dqn_replay_step
so the symptom was masked, but the anti-pattern was identical. Fixed
in the same commit per the user's "atomicAdd should not be used at all"
reminder.
Local smoke (RTX 3050, b=16, 1k steps, seed=16962):
step | l_pi (pre → post) | l_q (pre → post)
100 | 27.5 → 0.56 | 18.1 → 0.19
500 | 743.1 → 4.59 | 93.3 → 0.19
999 | 9080.5 → 65.1 | 186.1 → 0.19
l_q now bit-flat at per-step mean (~0.19) — no step-count fingerprint.
l_pi grows organically with policy excursion (PPO surrogate magnitude
when ratio→clamp_max under Q-distillation-driven policy updates),
which is the legitimate diagnostic the prior staleness was burying.
Changes:
- ppo_clipped_surrogate_fwd: scalar [1] outputs → per-batch [B]
- dqn_distributional_q_bwd: remove atomicAdd; per_batch[B] only
- ppo_loss_reduce_b.cu (NEW): two block tree-reduce kernels
• ppo_loss_reduce_b (dual: PPO loss + entropy loss)
• mean_reduce_b_f32 (single, reused by DQN head)
- PolicyHead + DqnHead: load reducer cubin, add reduce_loss_to_scalar
- Trainer: allocate ss_pi_loss_per_b_d + ss_pi_loss_entropy_per_b_d,
invoke reducer after each backward (PPO + DQN replay paths)
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
|
30982963ef |
feat(rl): IQN τ tail-recency consumer + G24/G25 invariants (F5)
When TAIL_EVENT_RECENCY < N_window (default 100), boost τ_min by factor (default 1.5) — agent uses more pessimistic action selection during tail-recent regimes. Defense-in-depth alongside Kelly resurrection (F2): F2 catches the sizing-layer absorbing state; F5 makes action selection more risk-averse right after a shock. Tests: - G24: TAIL_EVENT_RECENCY < 100 → τ_action ≥ τ_min × 1.5 - G25: TAIL_EVENT_RECENCY ≥ 100 → τ_action behavior unchanged 22 risk_stack_invariants now pass on RTX 3050 Ti. Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com> |
||
|
|
6dacde95ef |
feat(rl): popart per-account max-magnitude envelope (F4) — Problem 3 fix
Adds envelope-detector floor on popart_sigma so single-account tail events within a batch are captured instead of diluted by the batch-mean variance. Kernel changes (rl_popart_normalize.cu): - warp_reduce_max helper alongside existing warp_reduce_sum - Pass 1 extension: per-thread local_max_abs via fmaxf(fabsf(r)) - 3rd shared-mem bank for max_abs reduction (s_max_abs[]) - envelope-detector inside tid==0 block: fast-up, slow-down (α_d=0.01) - new_sigma = fmaxf(new_sigma, max_r_ema) before write - CRITICAL (Issue β): Pass 3 broadcast slots moved block_dim*2 → block_dim*3 Trainer launch update: smem_bytes = (block_x * 3 + 3) * sizeof(f32). Per Theorem 6 (spec v3): at Run #8 step 7377 with r_tail=-19.45, the envelope captures 19.45 instantly via fast-up path; sigma jumps from 2.327 → 19.45 in one step. The popart_v_correct kernel handles the σ discontinuity affinely so V regression doesn't destabilize. Decay phase: α_d=0.01 (69-step half-life); max_r_ema returns to typical batch-max baseline over ~200 steps after a single shock. Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com> |