6f3639bfbf22acaaefefacdf9a0f0edf96b5126b
5780 Commits
| Author | SHA1 | Message | Date | |
|---|---|---|---|---|
jgrusewski
|
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 |
||
|
jgrusewski
|
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:
|
||
|
jgrusewski
|
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 |
||
|
jgrusewski
|
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:
|
||
|
jgrusewski
|
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
|
||
|
jgrusewski
|
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 |
||
|
jgrusewski
|
969caf26c3 |
docs(ml-alpha): spec + plan for Phase 1B trainer rollout-buffer + GAE refactor
Companion docs for commit |
||
|
jgrusewski
|
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>
|
||
|
jgrusewski
|
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>
|
||
|
jgrusewski
|
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>
|
||
|
jgrusewski
|
63fc16f173 |
fix(rl): surfer-scaffold v5.2 — drop wr_excess clamp, full scaffold when novice
alpha-rl-mjsmr (
|
||
|
jgrusewski
|
4a4953b01f |
fix(rl): surfer-scaffold v5.1 — decay re-engagement requires EXCESS alertedness
alpha-rl-zf6s5 (
|
||
|
jgrusewski
|
38a4aa15b3 |
feat(rl): adaptive surfer-scaffold reward shaping (spec v5 A+B combined)
Diagnosis from alpha-rl-8fb55 ( |
||
|
jgrusewski
|
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 |
||
|
jgrusewski
|
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> |
||
|
jgrusewski
|
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
|
||
|
jgrusewski
|
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>
|
||
|
jgrusewski
|
f428be794b |
Revert "feat(rl): B-11-β Q-distill informativeness gate"
This reverts commit
|
||
|
jgrusewski
|
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,
|
||
|
jgrusewski
|
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> |
||
|
jgrusewski
|
8c7ce02da9 |
feat(rl): B-10 policy-quality cascade diagnostic
alpha-rl-8gtk2 (B-7+B-8+B-9 run at SHA
|
||
|
jgrusewski
|
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> |
||
|
jgrusewski
|
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> |
||
|
jgrusewski
|
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> |
||
|
jgrusewski
|
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> |
||
|
jgrusewski
|
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> |
||
|
jgrusewski
|
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> |
||
|
jgrusewski
|
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> |
||
|
jgrusewski
|
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> |
||
|
jgrusewski
|
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> |
||
|
jgrusewski
|
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>
|
||
|
jgrusewski
|
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>
|
||
|
jgrusewski
|
1aa92f57f0 |
fix(rl): eval-boundary addendum — reward_scale warmed_flag + pos_max_ema rate-cap
The parent eval-boundary fix (
|
||
|
jgrusewski
|
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> |
||
|
jgrusewski
|
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>
|
||
|
jgrusewski
|
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> |
||
|
jgrusewski
|
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). |
||
|
jgrusewski
|
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).
|
||
|
jgrusewski
|
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).
|
||
|
jgrusewski
|
13d8ed76da | docs: checkpoints + eval-diag spec + plans (v1 superseded by v2) | ||
|
jgrusewski
|
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>
|
||
|
jgrusewski
|
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> |
||
|
jgrusewski
|
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> |
||
|
jgrusewski
|
cfe40f2f3f |
test(rl): G10-G12 back-compat + G15-G21 Kelly resurrection (F2.3+F2.4)
G10/G11/G12: explicit dead-zone disable (DEAD_ZONE_FLAG=0, TIMEOUT_FLAG=0) before Kelly kernel launch — verifies analytic-Kelly path stays unmodified. New invariants: - G15: DEAD_ZONE_FLAG = 1 on composite kelly=0 ∧ all-flat ∧ no-cooldown - G16: DEAD_ZONE_FLAG = 0 when any condition violated - G17: Kelly resurrection sets kelly_f = ε_recovery_live when flag set - G18: Kelly retains analytic value when flag NOT set - G19: ε_recovery_live ramps linearly from ε_min at T=0 to ε_max at T≥N - G20: TIMEOUT_FLAG fires when DURATION > MAX_DURATION - G21: Kelly resurrection NOT triggered when TIMEOUT_FLAG = 1 20 risk_stack_invariants now pass on RTX 3050 Ti. Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com> |
||
|
jgrusewski
|
ad05ffeb2a |
feat(rl): Kelly resurrection override (F2) — Problem 1 fix
When DEAD_ZONE_FLAG fires (kelly_f=0 ∧ all-flat ∧ no-cooldown), override Kelly with ε_recovery_live (regime_observer-computed value ramping from ε_min=0.05 to ε_max=0.50 as TAIL_EVENT_RECENCY grows). TIMEOUT_FLAG gates resurrection off after MAX_DURATION=1000 steps (safety net). Per Theorem 1 (spec v3): exit probability from absorbing state per step is > 1 - 10^-23 given π(any Open) ≥ 0.05 over 1024 accounts. The Kelly absorbing state no longer exists in the state graph (except the operator-intended TIMEOUT terminal). Block appended to end of rl_kelly_fraction_controller.cu (analytic Kelly clamp first, then conditional override) and mirrored in Layer-4 branch of rl_fused_controllers.cu. |
||
|
jgrusewski
|
bf214d1a09 |
refactor(rl): v9 slot rename — RL_EVAL_WARMUP_* → RL_REGIME_TRANSITION_* (F3)
Theorem 2 (v9 behavior preserved): pure constant identifier rename across 3 files; no physical migration. Slots 685/686/687 keep their addresses; v9 kernel, trainer reset_session_state, and diag emit all reference the same memory. - crates/ml-alpha/src/trainer/integrated.rs: 4 references renamed - crates/ml-alpha/cuda/rl_eval_warmup_decay.cu: 3 `#define`s renamed - crates/ml-alpha/examples/alpha_rl_train.rs: ~12 references renamed (import + diag emit) Unblocks the 4 pre-existing errors that F1.1 introduced. Branch now compiles cleanly with regime_observer foundation (F1) + v9 cleanup (F3). |
||
|
jgrusewski
|
b09feaf9e4 |
feat(rl): regime_observer diag emit (F1.8)
Adds 'regime' block to risk_stack diag JSON, surfacing 5 nested groups: - dead_zone: flag, duration, timeout_flag, max_duration - tail: recency, session_pnl_variance_ema, sigma_threshold, welford_count - kelly_eps_recovery: factor, live, min, max, n_recovery - popart_envelope: max_abs_reward_ema, decay_alpha - iqn_tau_boost: factor, n_window All values read directly from ISV slots (drift-free per spec issue #9). F3 will add the 'transition' group when v9 slot renames land in the diag. Co-Authored-By: Claude Sonnet 4.6 <noreply@anthropic.com> |
||
|
jgrusewski
|
0a9ff73b1f |
feat(rl): reset_session_state regime boundary policy (F1.7)
Extends reset_session_state with 8 new regime_observer slot resets per spec section "Regime observer reset policy at fold/eval boundaries": - Transient state (FLAG/DURATION/TIMEOUT) → 0 - RECOVERY_FACTOR → 1.0 - TAIL_EVENT_RECENCY → 1e6 sentinel - KELLY_EPS_RECOVERY_LIVE → 0.50 (= ε_max) - PREV_WORST_PNL → 0 (CRITICAL — prevents spurious 3σ tail at boundary) - POPART_MAX_ABS_REWARD_EMA → 0 (F4 envelope reset) Welford state (M2/MEAN/COUNT) intentionally NOT reset — variance of session_pnl_change is regime-invariant per pearl_adaptive_carryover_discipline. Issue γ fix: without PREV_WORST_PNL reset, the first regime_observer call after a fold boundary would see Δworst = 0 - (-\$15k_train) = +\$15k, fire a spurious 3σ tail event, and peg ε_recovery at ε_min for 100 steps. |
||
|
jgrusewski
|
67f862191f |
feat(rl): regime_observer bootstrap entries (F1.6) — array 179→199
Initial values for the 20 new regime_observer slots per spec v3: - Transient state: DEAD_ZONE_FLAG/DURATION/TIMEOUT = 0, RECOVERY_FACTOR = 1 - Sentinels: TAIL_EVENT_RECENCY = 1e6, PREV_WORST_PNL = 0, MAX_ABS_REWARD_EMA = 0 - Welford state: M2/MEAN/COUNT = 0 (bootstrap via first-observation, gated count>=10) - Configs: ε_recovery_min=0.05, ε_recovery_max=0.50, N_recovery=100 - Configs: TAIL_SIGMA_THRESHOLD=3.0, MAX_DURATION=1000, popart α_d=0.01 - IQN τ tail boost: factor=1.5, window=100 |
||
|
jgrusewski
|
9e97c2acaf |
feat(rl): wire regime_observer into step_with_lobsim_reward_and_train (F1.5)
Inserts flat_count helper + regime_observer kernel launches BEFORE the risk-stack controllers (CMDP/IQN/Inventory/Kelly). Uses prev_position_lots_d (current-step snapshot, just-updated by LobSim per existing pipeline) as the input to flat_count. One-step lag on kelly_f/cooldown/worst_pnl reads is benign per Theorem 1: absorbing state persists across step boundaries, detection at T+1 still fires before harm escalates. No consumer changes yet — F2/F3/F4/F5 land independently. |