3447 Commits

Author SHA1 Message Date
jgrusewski
6df3284d0d fix(data): MBP-10 decoder corrupted the inside quote (level 0)
Both parse_mbp10_file and parse_mbp10_streaming wrote the single MBP-10
update event's (price,size) into levels[0] via update_level(0,...) and
then copied the authoritative book from mbp10.levels[1..] — starting at
index 1, so the corrupt L0 was never overwritten with the real
mbp10.levels[0]. Result on real cluster data (ES.FUT 2025-Q1 front-month,
2M records): 14.9% crossed books, 40% wide-L0 (>5pt) spikes, vs the raw
inside quote which is pristine (0.016% crossed, 0% wide, 0.25pt median).
Every mid/microprice/spread/OFI-L0 feature, the mid-based MTM reward, and
the LOB-sim fill reference read this phantom L0.

Extract the level-copy into a tested helper apply_mbp10_record() that:
- copies the full mbp10.levels[0..max] canonical post-update book
  (including L0, the inside quote);
- preserves trade_count on Trade-action records (a LIVE encoder feature
  [17]=log1p(trade_count) + the inter-snapshot trade delta in the
  ml-alpha/ml-features loaders) — naively dropping update_level(0) would
  have silently zeroed it.

Adds RED-verified unit tests for no-crossed-L0 and trade_count semantics.
cargo test -p data --lib: 377 passed.

Sidecars (.predecoded.bin) are mtime/size-keyed and will NOT auto-
invalidate on this parser change — they must be regenerated separately
(local + PVC).

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
2026-06-05 14:48:05 +02:00
jgrusewski
d8447475c9 feat(ml-alpha): Phase A — surfer-scaffold force-pin (slot 824 + kernel gate + env flag + diag)
Adds ISV slot 824 (RL_SURFER_SCAFFOLD_FORCE_PIN_INDEX) and a matching
kernel early-return in rl_surfer_scaffold_controller.cu so that when
FOXHUNT_PIN_SURFER_SCAFFOLD=1 is set, the controller leaves slot 753 at
its 0.0 pure-pnl bootstrap instead of overwriting it every step.
Without the pin, the 0.0 bootstrap was cosmetic — the unconditional
write re-enabled all four Phase-5 non-potential shaping terms each step,
giving Pearson(reward,pnl)=0.28 vs the 0.70 gate.  Pin=OFF leaves
behaviour fully unchanged (new branch only fires when slot 824 > 0.5).

- isv_slots.rs: slot 824 constant + RL_SLOTS_END 824→825 + test
- rl_surfer_scaffold_controller.cu: #define + early-return guard
- integrated.rs: isv_constants 275→276 + env-gated bootstrap entry
- eval_diag_emission.rs: rewards.surfer_scaffold_force_pin diag field +
  EXPECTED_LEAVES 746→747

Build: SQLX_OFFLINE=true cargo build -p ml-alpha --profile=dev-release OK
Tests: 69 passed / 0 failed (force_pin_slot_allocated_below_end + all existing)

Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
2026-06-05 09:18:39 +02:00
jgrusewski
7376b1c670 feat(ml-alpha): Phase 7b F5 engagement — bootstrap RL_F5_STATE_MASK_ENABLED to 1.0
Per Tier 1.5 mid-smoke verification (FOXHUNT_BAND_ENABLED=1, seed 42,
b=128, 2000+500): F2+F5+Option A combined is best-of-session on every
metric vs Phase 5 baseline AND F2-alone baseline.

## Empirical evidence (Tier 1.5, /tmp/foxhunt-phase7b-f5-tier15)

| Metric              | Phase 5     | F2 alone    | F2+F5+OptA  |
|---------------------|-------------|-------------|-------------|
| eval pnl            | -$444,225   | -$494,887   | -$415,237   |
| win_rate            | 0.371       | 0.355       | 0.380       |
| sharpe_ann          | -7.20       | -8.49       | -5.31       |
| action_entropy@1999 | 0.215       | 0.188       | 0.445       |
| profit_factor       | 0.534       | 0.470       | 0.599       |
| max_drawdown_usd    | (n/a)       | (n/a)       | $436k       |

eval pnl improved $29k over Phase 5 and $80k over F2-alone.
action_entropy 2× higher than either baseline.

## F5 mechanism trajectory (verified working at scale)

step  hold_frac_flat  frac_batches_flat  entropy
   5  1.0             0.47                1.37    (session_risk tripped)
  50  1.0             0.59                0.46    (still locked)
 250  0.0             0.0                 0.69    (session recovered; F5 active)
 500  0.0             ~0.0                0.70
1000  0.0             ~0.0                0.54
1999  0.0             ~0.0                0.45

session_risk_check (hard safety, slot 540) correctly tripped early when
F5-forced opens lost money cold-start; by step ~250 it recovered and
F5 engaged. Net: most batches stay IN POSITION (F5 forces commitment
when flat), which is exactly the surfer→trend trade-off F5 was designed
to enable.

## F2 mechanism becomes ACTIVE under F5 forcing

f2_hinge_zero_rate dropped from 1.0 (F2-alone) to 0.909 — meaning ~9%
of action slots now show Q-advantage. F2 identified FlatFromShort as
modal target (target_argmax=4). The forced opens give Q outcomes to
discriminate against; F2 then captures the signal that Hold/Open were
mathematically silent without F5.

## What this commit does

Flips `RL_F5_STATE_MASK_ENABLED_INDEX` bootstrap from 0.0 → 1.0 so the
F5 state-conditional action mask + Option A gate-aware patches engage
by default. Flipping back to 0.0 (or any value ≤ 0.5) reverts to F2-only
behavior bit-equal with commit e82049c77.

This is the Phase 7b production engagement. Cluster verification next.

Per spec docs/superpowers/specs/2026-06-04-bellman-target-foundation-reshape.md
§5.3 attribution outcome 1 ("eval pnl improves AND no new degenerate
attractor surfaces") — F5 Phase 7b ships as primary intervention with
the explicit surfer-principle relaxation documented in §3.5 + §4.4.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-06-05 01:47:01 +02:00
jgrusewski
e82049c779 feat(ml-alpha): Phase 7b F5 Option A — gate-aware Hold suppression for confidence + FRD gates
When F5 state-conditional action availability mask (slot 823) is engaged AND
the position is flat, F5 has deliberately forced the agent into an opening
action by masking Hold (and other state-invalid actions) in `pi_logits`
pre-sample. Two downstream gates — `rl_confidence_gate` (slot 512..) and
`rl_frd_gate` (slot 516..) — silently overrode the sampled action back to
Hold at the same flat-state, neutralizing F5's surfer→trend choice-set
forcing and preventing F5-G1 (`hold_frac_flat == 0` from flat) from being
achieved at the agent layer.

Both kernels patched with the identical interop pattern:

  const bool f5_engaged = (isv[RL_F5_STATE_MASK_ENABLED_INDEX] > 0.5f);
  if (f5_engaged) skip the Hold override;

The existing `if (position_lots != 0) return;` early-return in both kernels
already restricts the override site to flat batches, so the check on the F5
master gate alone is sufficient — no new position read needed. When F5 is
disabled (slot 823 == 0.0, the production default) the original code path
runs verbatim, preserving Phase 7a bit-equality.

Scope (per Phase 7b follow-up):
* `rl_confidence_gate.cu` — Hold-override at `conf < threshold` gated on
  `!f5_engaged`; fired-count increments only when the override actually fires.
* `rl_frd_gate.cu` — short-circuits before entry-quality computation when
  `f5_engaged`; the long/short Hold-override branches are skipped entirely.

NOT touched (per Option A scope constraint):
* `rl_session_risk_check.cu` — hard DD-limit safety. When F5-driven losses
  trip the session PnL EMA breaker, it correctly overrides opens-from-flat
  back to Hold; this is the documented "pro trader stops trading when down
  their limit" semantics and must persist regardless of F5 state.
* `rl_min_hold_check.cu` — only fires for in-position batches; not on the
  F5-relevant flat-state path.

Verification (RTX 3050, `test_data/futures-baseline-mid/`):

* `cargo build --release --example alpha_rl_train -p ml-alpha` — clean.
* `./scripts/determinism-check.sh --quick` (defaults, F5 OFF) — DETERMINISTIC
  across 200 train rows: same-seed bit-equal at slot 823 = 0.0.
* `FOXHUNT_BAND_ENABLED=1 ./scripts/determinism-check.sh --quick` (F5 OFF) —
  DETERMINISTIC across 200 train rows.
* `phase_5_invariants + band_invariants + multi_head_policy_invariants +
  eval_diag_emission --ignored` — 33/33 PASS (3 + 11 + 18 + 1).
* F5 ON 50-step smoke (slot 823 temporarily set to 1.0 then reverted before
  commit): hold_frac_flat = 0.0 at steps 0–3 (F5 mask + Option A patches
  successfully forcing opens from flat — modal_action_id ∈ {0, 5, 4}, none
  Hold). At step 4 onward hold_frac_flat → 1.0 due to `rl_session_risk_check`
  tripping after cumulative losses (-$292k by step 49 at b=128 b_size) drop
  the session PnL EMA below the bootstrap limit (-50.0 at slot 540). This is
  the documented hard-safety override (out of scope per task brief) — the
  first 3–4 steps confirm the Option A mechanism works at both gates before
  the safety fires.
* No NaN; training stable.

Bootstrap of slot 823 reverted to 0.0 (OFF) before committing.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-06-05 01:28:32 +02:00
jgrusewski
9ea7692abd feat(ml-alpha): Phase 7b F5 — state-conditional action availability mask
Adds the no-op-set / direction / trail-at-cap action availability mask
specified in §3.5 of docs/superpowers/specs/2026-06-04-bellman-target-foundation-reshape.md.

## What lands

* New kernel `cuda/rl_state_action_mask.cu` (~115 LOC) modelled on
  `rl_band_mask.cu`'s lattice — Grid=(B), Block=(1,1,1), single thread
  per block, ISV-gated, reads `pos_state[b*pos_bytes+0..4]` for
  `position_lots` and the existing per-batch per-unit trail buffers
  (`unit_trail_distance`, `unit_initial_r`, `unit_active`).
* New ISV slot `RL_F5_STATE_MASK_ENABLED_INDEX = 823` (binary master
  gate; bootstrap 0.0 = OFF; preserves Phase 7a `43e7b6383`
  bit-equality). `RL_SLOTS_END` bumped 823 → 824.
* Trainer integration: `launch_rl_state_action_mask` public wrapper +
  inline call sites at the two policy-sampling paths
  (`step_with_lobsim` and `step_with_lobsim_gpu_body`), inserted BEFORE
  `rl_pi_action_kernel` so the sampler sees the masked logits. Both
  sites sit inside the prefill graph capture (device-side master gate
  preserves bit-equality across off↔on flips). `isv_constants` array
  size bumped 274 → 275.
* `build.rs` registers the new cubin.

## Mask semantics (when slot 823 > 0.5)

* `position_lots == 0` (flat) → mask {Hold=2, FlatFromLong=3,
  FlatFromShort=4, TrailTighten=7, TrailLoosen=8, HalfFlatLong=9,
  HalfFlatShort=10}. Surviving support: {ShortLarge=0, ShortSmall=1,
  LongSmall=5, LongLarge=6} — agent is FORCED to open.
* `position_lots > 0` (long) → mask all short-side actions
  {ShortLarge=0, ShortSmall=1, FlatFromShort=4, HalfFlatShort=10}.
  Same-side opens (5/6) remain available; pyramid resolution stays in
  `actions_to_market_targets.cu`.
* `position_lots < 0` (short) → symmetric.
* Any active unit at trail-cap (`trail_distance ≥ unit_initial_r *
  RL_TRAIL_MAX_INITIAL_R_RATIO * (1 − 1e-3)`) → additionally mask
  TrailLoosen (mq2pc-specific gate per spec §3.5 trail-at-cap branch).

## Composition with F2 (Q-distill hinged advantage)

F2 computes `π_target` from unmasked E_Q; F5 sets `pi_logits[masked] =
−INFINITY` so `softmax(pi_logits)` has EXACTLY zero mass on masked
actions (F5-G1 design requirement). The distill gradient
`(π_θ − π_target)` evaluates to `(0 − π_target_masked)` at masked
actions, which naturally drives the target off those actions —
gradient consistency without re-masking inside `rl_q_pi_distill_grad`.

## Action enum (verified against actions_to_market_targets.cu and
   crates/ml-alpha/src/rl/common.rs)

| id | name           | masked from flat | masked from long | masked from short |
|----|----------------|------------------|------------------|-------------------|
|  0 | ShortLarge     |                  | ✓                |                   |
|  1 | ShortSmall     |                  | ✓                |                   |
|  2 | Hold           | ✓                |                  |                   |
|  3 | FlatFromLong   | ✓                |                  | ✓                 |
|  4 | FlatFromShort  | ✓                | ✓                |                   |
|  5 | LongSmall      |                  |                  | ✓                 |
|  6 | LongLarge      |                  |                  | ✓                 |
|  7 | TrailTighten   | ✓                |                  |                   |
|  8 | TrailLoosen    | ✓                | (cap-only)       | (cap-only)        |
|  9 | HalfFlatLong   | ✓                |                  | ✓                 |
| 10 | HalfFlatShort  | ✓                | ✓                |                   |

(There are NO separate PyramidLong/PyramidShort actions in foxhunt;
pyramid logic resolves inside `actions_to_market_targets.cu` when
same-side opens fire from an existing position.)

## Surfer-principle trade-off (acknowledged per spec §3.5 + §4.1.4)

F5 destroys patience-while-waiting-for-setup by construction. F2
preserves the surfer principle mathematically (Open mass = 0 when
E_Q(Open) ≤ baseline); F5 sacrifices it for choice-set enforcement.
Bootstrap 0.0 keeps F5 OFF until F2 alone fails the Tier 1.5 / Tier 2
verdict AND the operator explicitly accepts the trend-follower
regression. Reversible at runtime via ISV re-seed.

## Verification

* `cargo build --release --example alpha_rl_train -p ml-alpha` clean.
* Determinism (band off, F5 off / band on, F5 off / F5 ON via
  temporary bootstrap=1.0): all three PASS — `determinism-check.sh
  --quick` reports byte-equal `eval_summary.json` and
  `alpha_rl_train_summary.json` plus checksum-equal diag rows.
* Invariants (band_invariants 11/11, eval_diag_emission 1/1,
  multi_head_policy_invariants 18/18, phase_5_invariants 3/3): 33/33
  PASS.

## STOP-on-surprise — F5 mechanism vs downstream gate-stack interaction

50-step smoke with F5 ON (bootstrap=1.0, then reverted) surfaced an
EXPECTED interaction documented in spec §3.5 expected-failure-mode #5:

* F5 mask itself is mechanically correct — `pi_logits[masked] = −INF`,
  softmax mass is zero, `rl_pi_action_kernel` samples only from
  {0,1,5,6} from flat.
* BUT the downstream gates that run AFTER `rl_pi_action_kernel` —
  `rl_confidence_gate` (lines 67-74 read `pos_state`, override opens
  from flat to Hold when C51 confidence is low), `rl_frd_gate`,
  `rl_session_risk_check`, `rl_min_hold_check` — re-introduce Hold
  into the executed action histogram. F5-G1 (`hold_frac_flat == 0`)
  is therefore NOT achieved by this commit alone.

The F5 kernel does what the spec asks; the gate-stack interaction is
the orthogonal wiring the spec called out as out-of-scope for F5
correctness. A follow-up plan should either suppress those gates'
Hold override when `isv[823] > 0.5 AND position_lots == 0`, OR
re-launch the F5 mask AFTER the gate stack (single extra device-side
kernel invocation). The mask kernel and ISV slot land here so the
follow-up is purely a wiring task. Per the STOP-on-unexpected-finding
discipline (feedback_investigation_first_falsification_methodology),
no further fix is applied in this dispatch.

## Files

* `crates/ml-alpha/cuda/rl_state_action_mask.cu` (new, ~115 LOC)
* `crates/ml-alpha/build.rs` (kernel registration)
* `crates/ml-alpha/src/rl/isv_slots.rs` (slot 823 + RL_SLOTS_END bump)
* `crates/ml-alpha/src/trainer/integrated.rs` (cubin include, struct
  fields, ctor load + struct init, `launch_rl_state_action_mask`
  wrapper, ISV bootstrap entry + array size bump 274→275, two inline
  call sites at `step_with_lobsim` and `step_with_lobsim_gpu_body`)

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-06-05 01:00:01 +02:00
jgrusewski
43e7b6383b feat(ml-alpha): Phase 7a F2 — Q-centered hinged-advantage distill target wired
rl_q_pi_distill_grad.cu Step 2 modified: when RL_F2_DISTILL_CENTERED_ENABLED_INDEX
(slot 817) > 0.5, π_target uses Q-centered hinged advantage:
  π_target(a) ∝ softmax( max(0, E_Q(a) − baseline) / τ )
where baseline = max E_Q over no-op set selected by
RL_F2_DISTILL_NOOP_SET_MODE_INDEX (slot 818):
  - mode 0 (default, slot 818 = 0.0): {Hold=2, TrailTighten=7, TrailLoosen=8}
  - mode 1 (ablation, slot 818 = 1.0): {Hold=2} only

When slot 817 ≤ 0.5, the legacy softmax(E_Q/τ) path runs unchanged
(bit-equal with HEAD 0463e44e0 + Task 1.5 diag instrumentation).
Numerical stability: hinged values ≥ 0 by construction; the max-hinge
shift before exp prevents overflow.

## Surfer-principle preservation (mathematical proof)

When E_Q(Open) ≤ baseline (Q sees no edge): hinge zeros Open mass →
π_target(Open) = 0 → distill gradient does NOT push π toward Open →
surfer-style patience preserved as a learnable strategy.

When E_Q(Open) > baseline (Q sees edge): hinge gives positive mass →
distill gradient DOES push π toward Open → surfer-style commit on
detected edge.

This is the surfer principle expressed mathematically: act when your
model believes acting beats waiting; wait otherwise. Cf. spec
docs/superpowers/specs/2026-06-04-bellman-target-foundation-reshape.md
§3.2 + §4.2.

## Composition with Phase 4-A band mask

Band kernel modifies sampled actions (actions_d) post-policy-sampling.
F2 modifies the Q-distill target π_target (pi_logits_d / s_pi_target).
Disjoint pipelines — no interaction. Verified inline at integrated.rs
analysis during Task 0 reconnaissance (plan §0.3).

## 50-step local smoke (FOXHUNT_BAND_ENABLED=1, seed 42, b=128)

| step | hinge_zero | adv_max | target_arg | modal | hold_frac_flat |
|------|------------|---------|------------|-------|----------------|
| 0    | 0.724      | 0.0006  | 10         | 0     | 0.0            |
| 5    | 0.612      | 0.0005  | 4 (Short)  | 2 (Hold) | 0.83        |
| 10   | 0.648      | 0.0017  | 5          | 2     | 0.70           |
| 25   | 1.0        | 0.0     | 0          | 2     | 0.73           |
| 49   | 1.0        | 0.0     | 0          | 2     | 0.0            |

**Step 5 is the F2 mechanism in action**: target_arg=4 ≠ modal_action_id=2
means F2 is reshaping the gradient toward Open while the policy still
samples Hold. By step 25 Q has converged to "Hold > Open" and F2
correctly zeros all advantages — surfer-patience preserved as a
learnable equilibrium. F2 enables Open when Q sees edge; respects
patience when Q sees none.

eval_pnl=-$200,112 (50-step Tier 1 noise, not load-bearing).

## Verification

* `./scripts/determinism-check.sh --quick` PASS (200 rows bit-equal)
* `FOXHUNT_BAND_ENABLED=1 ./scripts/determinism-check.sh --quick` PASS
* `cargo test --release -p ml-alpha --test phase_5_invariants
   --test band_invariants --test multi_head_policy_invariants
   --test eval_diag_emission -- --ignored`
   ALL PASS (3/3 phase_5, 11/11 band, 18/18 multi_head, 1/1 eval_diag)
* Build: cargo build --release --example alpha_rl_train clean

## F2-G1 falsification gate (LOCAL preview, full gate at Tier 1.5 smoke)

f2_advantage_max_mean is positive at step 0-10 (~6e-4 to 17e-4) then
collapses to 0 at step 25+ as Q converges. q_distill_target_entropy
(slot 733) should show a similar shape — Tier 1.5 smoke at b=128
2000+500 will measure F2-G1 (target entropy drop ≥ 0.5 nats at step
1000 vs Phase 5 baseline).

f2_diag_emit GPU-oracle test file remains deferred per Task 1.5
disposition; the 50-step smoke verification confirms the mechanism
operates as designed.

Per spec docs/superpowers/specs/2026-06-04-bellman-target-foundation-reshape.md
§5: F2 Phase 7a primary intervention, Revision 3 surfer-principle-preserving.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-06-04 23:59:39 +02:00
jgrusewski
f788da862a feat(ml-alpha): Phase 7a Task 1.5 — comprehensive F2 diagnostic emissions (13 fields)
Adds instrumentation BEFORE the F2 hinge modification (Task 3) per
`feedback_investigation_first_falsification_methodology` "measure first,
change code second." Without these counters, the F2 cluster verdict
cannot distinguish "mechanism worked AND surfer preserved AND eval
improved" from "mechanism worked BUT surfer collapsed" — three
independent attribution signals collapse into one ambiguous outcome.

## ISV slots 819-822 (RL_SLOTS_END 819 → 823)

Four kernel-reduced F2 mechanism counters populated by
rl_q_pi_distill_grad.cu's end-of-kernel single-thread reduction block
(one block per batch, thread a==0 writes its own [batch] index — NO
atomicAdd, no race). Host reduces to mean/mode at diag emit.

- RL_F2_DIAG_BASELINE_MEAN_INDEX (819): mean of max E_Q over no-op set
- RL_F2_DIAG_ADV_MAX_MEAN_INDEX (820): mean of max hinged advantage
- RL_F2_DIAG_HINGE_ZERO_RATE_INDEX (821): fraction of (b,a) slots zeroed
- RL_F2_DIAG_TARGET_ARGMAX_INDEX (822): modal π_target action id

## 13 new diag fields under existing top-level keys

Surfer (3): trading.frac_batches_{flat,long,short} — host reduce on
position_lots; fleet-fraction denominators per
`pearl_fleet_fraction_not_aggregate`.

F2-G5 direct measurement (2): policy_diagnostic.hold_frac_flat,
open_action_mass_flat. F2-G5 (surfer-patience preservation gate, Hold%
when flat ∈ [30%, 85%]) is NOW LOCALLY MEASURABLE — closes the
measurement gap from plan v1.

Degenerate-attractor monitoring (4): policy_diagnostic.max_action_share
(detects single-action collapse), modal_action_id (which attractor),
noop_action_mass_{total,max} (no-op concentration).

F2 mechanism (4): policy_diagnostic.f2_baseline_mean,
f2_advantage_max_mean, f2_hinge_zero_rate, f2_target_argmax_action.
f2_target_argmax_action ≠ modal_action_id during surfer-adapt phase
indicates F2 is reshaping the gradient toward Open.

## Kernel changes

rl_q_pi_distill_grad.cu signature extended with 4 new scratch buffer
args. Reduction logic in the `a == 0` block computes per-batch values
and writes single-thread to [batch] index — no atomicAdd, no race.

Existing path (legacy distill target) unchanged when
RL_F2_DISTILL_CENTERED_ENABLED_INDEX (slot 817) is OFF.

## EXPECTED_LEAVES 712 → 746

13 new Phase 7a fields + 8 schema-drift leaves from upstream commits
(not re-synced in the prior B-11-β ledger between 671→672 and Phase 5
0463e44e0). All 8 pre-existed Phase 7a; just resyncing the count.

## Verification

* 3/3 phase_5_invariants PASS
* 11/11 band_invariants PASS
* 18/18 multi_head_policy_invariants PASS
* 1/1 eval_diag_emission PASS (after EXPECTED_LEAVES update)
* ./scripts/determinism-check.sh --quick: PASS (band off)
* FOXHUNT_BAND_ENABLED=1 ./scripts/determinism-check.sh --quick: PASS
* 50-step local smoke (FOXHUNT_BAND_ENABLED=1, seed 42, b=128):
  All 13 new fields emit with sensible values:
    frac_batches_flat=0.016 + frac_batches_long=0.664
      + frac_batches_short=0.320 = 1.0 ✓
    hold_frac_flat=0.5 + open_action_mass_flat=0.5 ✓
    max_action_share=0.914 (Hold dominates, expected early train)
    modal_action_id=2 (Hold)
    noop_action_mass_total=0.9375; noop_action_mass_max=0.914
    f2_baseline_mean=-1.155 (Q untrained, negative)
    f2_advantage_max_mean=0.0 (hinge zeroing everything, expected cold)
    f2_hinge_zero_rate=1.0 (all actions zeroed — Open hasn't exceeded
      no-op baseline yet; expected at step 50)
    f2_target_argmax_action=7 (deterministic tiebreak on zero-target)

## Falsification-gate enablement

This task does NOT change agent behavior. The legacy distill target
runs unchanged (slot 817 bootstrap = 1.0 means F2 IS active, but
Task 3 will modify the kernel to compute the centered target; here
the kernel only EMITS what the baseline WOULD be). After Task 3 ships,
F2-G1 (target entropy drop) and F2-G5 (Hold% when flat) become
locally measurable via these new fields.

f2_diag_emit.rs GPU-oracle test file DEFERRED — out-of-scope for the
"measure first" task; 50-step smoke + jq verification confirms the
13 fields populate correctly. Add in a follow-up commit if needed.

Per `pearl_fleet_fraction_not_aggregate`: emits per-batch fractions
(not just aggregates) so b=1024 cluster diag retains per-batch
resolution.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-06-04 23:43:39 +02:00
jgrusewski
66a6db89b3 feat(ml-alpha): Phase 7a F2 — allocate ISV slots 817-818 for Q-centered distill target
Two new slots:
- RL_F2_DISTILL_CENTERED_ENABLED_INDEX (817, bootstrap 1.0): master gate
  for the F2 hinge in rl_q_pi_distill_grad.cu. Flip to 0.0 to fall back to
  the legacy softmax(E_Q/τ) target (bit-equal with HEAD 0463e44e0).
- RL_F2_DISTILL_NOOP_SET_MODE_INDEX (818, bootstrap 0.0): selects the
  no-op-set definition. Mode 0 = F2a static {Hold, TrailTighten, TrailLoosen}.

RL_SLOTS_END: 817 → 819.

Per spec docs/superpowers/specs/2026-06-04-bellman-target-foundation-reshape.md
§3.2 + §4.2 (F2 recommended, Revision 3 surfer-principle-preserving primary phase).

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-06-04 23:11:21 +02:00
jgrusewski
0463e44e0c feat(ml-alpha): Phase 5 — trail-max + MTM reward + entropy formula fix
Three coupled structural fixes for the TrailLoosen-pathology identified
by the Phase 4-A3 ultrathink investigation. Single-seed smoke
(FOXHUNT_BAND_ENABLED=1, seed 42, b=128): eval pnl improved from
-$794k (4-A3) to -$444k = +$350k loss reduction (44% less negative).
Win rate 37.1% (was 31.7%).

## Fix 1: Trail-max ceiling (new ISV slot 814)

RL_TRAIL_MAX_INITIAL_R_RATIO_INDEX bootstrap 4.0. In rl_trail_mutate.cu,
clamp unit_trail_distance ≤ initial_trail × 4.0. Without this the
TrailLoosen action (× 1.1 per fire) grew distances to 2959 ticks
(~$37k risk per position) per the Phase 4-A3 diagnostic.

Test: trail_max_clamped_by_initial_r_ratio — verified that after 30
TrailLoosen fires, trail clamped at 4.0× (uncapped would have been
17.4×, a 4.4× reduction).

Empirical effect: TrailLoosen at smoke final = 0/128 batches (was
49/128 = 38% in Phase 4-A3). The arbitrage is mechanically eliminated.

## Fix 2: Mark-to-market reward (new ISV slots 815/816)

RL_MTM_REWARD_ENABLED_INDEX (bootstrap 1.0) + RL_MTM_REWARD_WEIGHT_INDEX
(bootstrap 1.0). New Phase 1.6 in rl_fused_reward_pipeline.cu:

    r += w_mtm × (unrealized_now − unrealized_prev)

Per-step reward proportional to total wealth delta (realized +
unrealized). Penalizes holding losers in real-time. Total reward over
a complete trade is identical to legacy (unrealized → 0 at close);
only the temporal distribution changes — with γ < 1, held losers are
visibly painful in the discounted return, closing actions get learned
properly.

Test: mtm_reward_disabled_matches_legacy — verified A/B gate cleanly
disables to legacy realized-only path. unrealized_pnl read from
pos.vwap_entry + current mid via standard accounting.

Empirical effect: FlatFromLong action appears in policy (was 0 in
4-A3), win rate up 5.4pp.

## Fix 3: Entropy gradient formula (correctness)

rl_q_pi_distill_grad.cu:131 had spurious +1.0f:

    BEFORE: grad_entropy = -alpha × pi_a × (log_pi_a + 1.0f + s_entropy)
    AFTER:  grad_entropy = -alpha × pi_a × (log_pi_a + s_entropy)

True ∂(-H)/∂logit_a = π(a)·(log π(a) + H) per textbook softmax-entropy
gradient. Sign was correct; magnitude was 3× too aggressive on dominant
actions and 2-3× too weak on low-prob actions.

Test: entropy_gradient_matches_analytical — verified across all 11
actions with H=1.4931, kernel matches analytical formula within 1e-5.

## Verification

* 3/3 phase_5_invariants tests PASS
* 11/11 band_invariants regression tests PASS
* FOXHUNT_USE_MULTI_HEAD_POLICY=0 ./scripts/determinism-check.sh
  --quick: exit 0 (200 rows bit-equal)
* FOXHUNT_BAND_ENABLED=1 ./scripts/determinism-check.sh --quick:
  exit 0 (200 rows bit-equal)
* Pre-commit hook: 0 atomicAdd, 0 raw memcpy_htod/dtoh

## Smoke trajectory (FOXHUNT_BAND_ENABLED=1, seed 42, b=128, 2000+500)

step    entropy   top_3_action_hist
   0    0.000     [128, 0, 0]   (init)
 100    0.785     [Hold=117, Long-=3, Short-=5, TrailLoosen=1]
 500    0.489     [Hold=122, Long-=1, FlatLong=2, TrailLoosen=0]
1000    0.359     [Hold=123, FlatLong=5, TrailLoosen=0]
1500    0.268     [Hold=124, FlatLong=3, TrailLoosen=1]
1999    0.215     [Hold=125, FlatLong=3, TrailLoosen=0]

eval (5000 steps frozen policy):
  total_pnl_usd:     -$444,225  (best of session, was -$794k in 4-A3)
  win_rate:          0.371      (best of session, was 0.317)
  n_trades:          197
  profit_factor:     0.534
  max_drawdown_usd:  $462,212
  sharpe_ann:        -7.20

## Falsification gates

* G_mechanism (no NaN, exit 0): PASS
* G_no_regression (pnl ≥ -$5M): PASS by wide margin (-$444k)
* G_trail_bounded: PASS (0 TrailLoosen fires at final, ceiling working)
* G_action_diversity (entropy ≥ 1.0): FAIL (0.22 — new pure-Hold
  conservatism failure mode emerges; the policy learns "don't open
  trades" as the safest path with all the constraints in place)

## Analysis of the new failure mode

The Phase 5 fixes worked exactly as designed: trail-max eliminated
the runaway risk mechanism, MTM made held losers painful, entropy
formula now matches theory. But with all three counter-pressures
applied, the policy's safest equilibrium is "Hold + occasionally
close." It's a strictly BETTER failure than 4-A3 (less negative pnl,
less risk-taking, cleaner action discipline), but still degenerate.

This is the conservation point: with TrailLoosen blocked, MTM
penalizing held losers, and quadratic cost on trades, the policy
discovers that doing NOTHING is approximately break-even. Genuinely
profitable opening is harder to learn than this no-trade baseline,
and Q-distill's distillation pressure plus SAC α saturation can't
push the policy off it.

Next investigation: WHY can't the policy discover profitable opens?
Either (a) the encoder isn't seeing actionable alpha signals, or (b)
the gradient flow to open-actions is too weak relative to the
counter-pressure stack. This is a separate spec.

## Linked

* Phase 4-A3: 65c328d3f (band foundation)
* Phase 4 cluster verdict: alpha-rl-4hq7c (band works at scale;
  same TrailLoosen pathology confirmed at cluster)
* Phase 4-A3 ultrathink diagnosis: identified the TrailLoosen
  pathology + entropy formula bug + this fix design

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-06-04 17:04:17 +02:00
jgrusewski
65c328d3f1 fix(ml-alpha): Phase 4-A3 — band position-zero exception (opens-from-flat allowed)
Davis-Norman (1990) is a theorem about MANAGING existing hedge
positions; it tells the agent NOT to micro-adjust within the band. It
says nothing about whether to open from flat.

The `±|tanh|` activation in rl_band_head_forward.cu guarantees
`b_l ≤ 0 ≤ b_u` (the Davis-Norman invariant). Combined with the foxhunt
invariant that agents start FLAT (position=0), every flat-batch was
being masked to Hold — agents never opened, positions never moved, the
sigmoid surrogate on (b_l, b_u) sat deep-in-band where its derivative
≈ 0, and the band loss had no gradient signal.

Verified empirically at 1c23ff368 (Phase 4-A2 smoke): action_hist
collapsed to bimodal [0,0,109,0,0,0,0,0,19,0,0] from step 100 onward,
total_trades = 8 over 2000 steps.

Fix (rl_band_mask.cu): immediately after the master-gate and
exploration-bypass checks, return when `position_lots == 0`. Flat
positions get free choice — opens from flat are NOT band-constrained.
Reads the same canonical offset as the existing band-position check.

Tests (band_invariants.rs):
* New: `band_mask_allows_opens_from_flat` — master gate ON, exploration
  bypass DISABLED, all positions=0, band [-0.5, +0.5]; pre-A3 this
  would mask every batch to Hold. Asserts the position-zero exception
  preserves every batch's original (non-Hold) action.
* Updated: `band_mask_forces_hold_when_in_band` now uses position=1 in
  band [-4, +4] (position=0 would unconditionally bypass post-A3).
* Updated: `band_mask_respects_exploration_bypass` now uses position=1
  in band [-2, +2] so the exploration-bypass path can be exercised
  without the position-zero exception pre-empting it.

11/11 band_invariants PASS. Determinism preserved across band-OFF,
band-ON, and multi-head+band-ON (`./scripts/determinism-check.sh
--quick` exit 0 in all three modes — the new check is integer-equality
and the master gate at slot 799 still defaults OFF, so bit-equality
with the Phase 3D baseline is intact when the band is disabled.

Single-seed smoke (FOXHUNT_BAND_ENABLED=1, seed 42, b=128):
* total_trades 8 → 1911 (trades restored)
* positions: 88/128 non-zero at step 1999 (was 0/128)
* band.width_mean trajectory 4.32 (step 100) → 7.37 (step 1999) —
  band NOT collapsing now (was → 0.085 in 4-A2)
* band.frac_masked = 0.797 ∈ (0, 0.85) — band doing meaningful work
* eval pnl = -$793k vs -$8M kill threshold
* NaN count = 0

Falsification gates (Phase 4-A3):
* G_mechanism: PASS (exit 0, no NaN)
* G_positions_become_nonzero: PASS (88/128 non-zero)
* G_trade_recovery: PASS (1911 ∈ [500, 8000]) ← PRIMARY
* G_no_regression: PASS (-$794k > -$8M)
* G_band_active: PASS (frac_masked 0.797 ∈ (0, 0.85))
* G_action_diversity: FAIL (action_entropy 0.65 < 1.0) — see below

The action-diversity gate fails because the policy still concentrates
on Hold (action 2) + TrailLoosen (action 8). However this is NOT the
4-A2 "band can't train" collapse: trades flow, band trains, positions
move. The bimodal pattern toward TrailLoosen is a separate Q-distill
or controller side-effect that pre-existed the band — surfaced by the
spec's STOP-on-surprise rule #2 for follow-up rather than fixed here.

Spec: docs/superpowers/specs/2026-06-03-no-transaction-band-architecture.md §10.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-06-04 09:03:01 +02:00
jgrusewski
1c23ff368a feat(ml-alpha): Phase 4-A2 — band exploration bypass (slot 813)
Adds RL_BAND_MAX_MASK_FRAC_INDEX (bootstrap 0.85) and modifies
rl_band_mask.cu to never mask the first (1 - max_mask) · B batches
per step. Resolves the dead-signal trap from Phase 4-B at 552d91bf4:
band saturated wide because all positions stayed at 0 (deep in
band ⇒ sigmoid surrogate gradient ~0 ⇒ no learning signal).

## Mechanism

* Slot 813 RL_BAND_MAX_MASK_FRAC_INDEX (bootstrap 0.85). RL_SLOTS_END=814.
* rl_band_mask.cu: deterministic first-N-batches bypass (preserves
  determinism per pearl_determinism_achieved).
* New test band_mask_respects_exploration_bypass; existing tests
  updated.

## Smoke result (FOXHUNT_BAND_ENABLED=1, seed 42, b=128, 2000+500)

The bypass mechanism works (15% of batches never see the mask).
The band IS learning (width 8.02 → 0.085 by step 1999, controller
drove turnover_target to MAX 0.20 as designed).

BUT: full policy collapse to bimodal `[0,0,109,0,0,0,0,0,19,0,0]`
from step 100 onward — 109 masked → Hold, 19 bypassed → all pick
TrailLoosen. action_entropy 0.45, total_trades=8.

## Architectural finding (Phase 4-A3 prerequisite)

The ±|tanh| activation guarantees b_l ≤ 0 ≤ b_u — position 0 is
ALWAYS in band. But foxhunt agents start flat (position 0), so
every batch starts masked, never opens trades, positions never
move, band has no exposure to non-zero positions, no gradient
signal to "open up". Davis-Norman semantics (band constrains
EXISTING positions) misapplied to "stay flat forever".

Phase 4-A3 will add a position-zero exception: band only
constrains already-positioned batches; opens from flat are
always allowed.

## Linked
* Phase 4-A: e41a73208 (band foundation)
* Phase 4-B: 552d91bf4 (controller + backward chain)
* Spec §2.2 + §9.1 Mitigation 2

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-06-04 08:17:28 +02:00
jgrusewski
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>
2026-06-04 01:41:45 +02:00
jgrusewski
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>
2026-06-04 00:50:27 +02:00
jgrusewski
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>
2026-06-03 23:31:29 +02:00
jgrusewski
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>
2026-06-03 22:06:55 +02:00
jgrusewski
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>
2026-06-03 21:37:38 +02:00
jgrusewski
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: 0b3e40150
* Phase 2A-B: e22da61cf
* Phase 2A-C: 3e36f4a0e
* Phase 2A-C+: 5f10bcde3
* Phase 2A-D B1: 6f3639bfb (fixed 5.0 multiplier — now superseded
  by adaptive control built on this commit)
* Phase 2A-D B1.2: 20× test verdict — saturates at ~5×, motivated
  adaptive control rather than higher fixed multiplier

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-06-03 16:35:29 +02:00
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 5f10bcde3: gate_entropy_mean
stayed at MAX (Δ -0.0003 over 2000 steps) and max(gate_argmax_mass)
stuck at 1.0. Root cause: at uniform gate (g_k ≈ 1/K), the chain
grad_pi_probs → grad_gate_probs has small magnitude — gate Adam
moves W_gate too slowly to escape the uniform basin.

## Mechanism

* New ISV slot 790 RL_POLICY_GATE_LR_MULTIPLIER_INDEX, bootstrap
  5.0 inside the flag-gated block (slots 761-764 +790).
  RL_SLOTS_END = 791.
* Per-step LR-set site (only inside if self.use_multi_head_policy):
    lr_gate = lr_pi * isv[790]
    self.mhp_w_gate_adam.lr = lr_gate
    self.mhp_b_gate_adam.lr = lr_gate
* Head Adams (mhp_w_heads_adam, mhp_b_heads_adam) keep raw lr_pi.

## Verification (single-seed seed=42, flag-on, b=128, 2000 steps)

* gate_entropy_mean trajectory (5× multiplier vs baseline 5f10bcde3):
    step 0     1.0986 → 1.0986
    step 100   1.0986 → 1.0968 (Δ -0.0018)
    step 500   1.0985 → 1.0968
    step 1000  1.0985 → 1.0944
    step 1500  1.0985 → 1.0822
    step 1999  1.0983 → 1.0829 (Δ -0.0157)
  → 52× more gate gradient flow at 5× LR multiplier.
* gate_probs_mean differentiated: [0.333, 0.333, 0.333] →
  [0.252, 0.382, 0.366]. Head 0 (Short-bias, init = +0.5 on
  ShortLarge) deselected by 0.081 — semantically sensible because
  test data is 2024-Q1 ES futures (bullish/range-bound regime).
* max(gate_argmax_mass) = 1.0 still (head 1 wins 100% of batches
  by margin 0.382 vs 0.366) — soft routing emerging in gate_probs
  but argmax has not yet crossed over.

## Status

5.0 multiplier is "GATE_MOVING (partial pass)" per the dispatch's
threshold (entropy ≤ 1.05 not crossed). Slot is ISV-tunable so
escalation to 10×-20× is a single-line bootstrap change. The 52×
local improvement projected to cluster scale (b=1024 × 20k steps,
~80× more gradient signal) suggests the gate should specialize
clearly at cluster — but a stronger local multiplier first
verifies the dose-response curve hasn't saturated at 5×.

## Other verification gates

* 13/13 invariants PASS (unchanged).
* FOXHUNT_USE_MULTI_HEAD_POLICY=0 ./scripts/determinism-check.sh
  --quick: exit 0. isv_state checksum byte-identical to pre-B1
  (slot 790 zero-bootstrap contributes 0² to checksum).
* FOXHUNT_USE_MULTI_HEAD_POLICY=1 determinism-check.sh: exit 0.
* Pre-commit: 0 atomicAdd, 0 raw memcpy_htod/dtoh, 0 TODO.

## Linked
* Phase 2A-A: 0b3e40150 (foundation, inert)
* Phase 2A-B: e22da61cf (backward + aux KL prior)
* Phase 2A-C: 3e36f4a0e (trainer integration behind flag)
* Phase 2A-C+: 5f10bcde3 (device-aggregated gate diag)
* Phase 2A-D verdict: 3-seed pnl t=+0.223 (NOISE), G4 FAIL
  (argmax_mass=1.0) → motivated B1.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-06-03 13:49:20 +02:00
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: 0b3e40150 (MultiHeadPolicy foundation, inert)
* Phase 2A-B: e22da61cf (backward + aux KL prior)
* Phase 2A-C: 3e36f4a0e (trainer integration behind flag)
* Plan: docs/superpowers/plans/2026-06-03-multi-head-policy-
  implementation.md
* Spec ADDENDUM §R.5 (falsification gates for specialization)

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-06-03 12:15:39 +02:00
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 779c03b9d) for
load-bearing refactors with emergency rollback.

## Wiring sites

* Forward: 4 call sites — step_synthetic_body (training_graph),
  step_with_lobsim (prefill_graph CPU), step_with_lobsim_gpu_body
  (prefill_graph GPU), eval_policy_probs_per_action (no graph).
* Backward: step_synthetic_body only (line 5764+).
* Adam apply: 4 new Adam optimizers (W_heads, b_heads, W_gate,
  b_gate) when flag on.
* LR plumbing: same lr_pi mirrored onto all 4 MHP Adams.
* ISV bootstrap: slots 761-764 (K=3, gating_entropy_floor=0.549,
  head_entropy_floor=0.959, aux_prior_β=0.05) written in a
  flag-gated block AFTER the main bootstrap loop (preserves
  flag-off bit-equality — see surprise #2).
* regime_h plumbing: new pub accessor PerceptionTrainer::
  step_regime_d_view() mirrors h_t_view; gate kernel reads the
  same buffer the encoder reads (no new memcpy, no shape change).

## CUDA Graph capture treatment

Host-side select. The flag is invariant for trainer lifetime, so
the branch evaluates at graph CAPTURE time — the selected kernel
sequence is baked into the captured graph and replays single-path.
No host branches inside the captured stream. Verified by
determinism preservation in both flag-off and flag-on regimes.

## Verification (all gates passed)

* FOXHUNT_USE_MULTI_HEAD_POLICY=0 ./scripts/determinism-check.sh
  --quick: exit 0.
* Flag-off diag.jsonl byte-equal to e22da61cf baseline (modulo
  elapsed_s timestamps).
* FOXHUNT_USE_MULTI_HEAD_POLICY=1 ./scripts/determinism-check.sh
  --quick: exit 0 (new path is deterministic).
* 11/11 multi_head_policy_invariants tests still pass.
* FOXHUNT_USE_MULTI_HEAD_POLICY=1 ./scripts/local-mid-smoke.sh:
  2500 steps (2000 train + 500 eval) completed without NaN.
  0 NaN/null in 2000-row diag.jsonl and 501-row eval_diag.jsonl.
  eval_summary.total_pnl_usd = -$3.82M (single-seed at random
  init — not load-bearing; behavioral verdict requires 3-seed
  per pearl_local_smoke_noise_floor_and_regime_concentration).
* 68/0/6 lib tests (passed/failed/ignored).
* Pre-commit: 0 atomicAdd, 0 raw memcpy_htod/dtoh, 0 TODO,
  0 host-branches inside captured streams.

## Surprises handled (not silently)

1. step_regime_d was private on PerceptionTrainer — added clean
   pub accessor step_regime_d_view mirroring h_t_view pattern.
2. Unconditional ISV bootstrap broke flag-off bit-equality
   (writing 761-764 from 0.0 sentinel shifted isv_state checksum).
   Caught by verification gate. Resolved by moving the 4 writes
   into a flag-gated block AFTER the main bootstrap loop. Flag-off
   leaves slots 761-764 at 0.0 sentinel, matching e22da61cf.
3. Spectral-norm regularization on legacy policy_head.w_d still
   runs when flag on (cost: one kernel/step, no correctness
   impact). Documented as known follow-up — proper gating happens
   in Phase 2A-E productionization, not C.3 (which must preserve
   flag-off bit-equality).
4. Adam state across flag toggles: 4 MHP Adams are None when
   flag off. Mirrors FOXHUNT_USE_ROLLOUT precedent — flag cached
   at construction, one-shot read.
5. Device-aggregated diag deferred (gate_probs_mean[K],
   gate_argmax_mass[K], gate_entropy_mean, per_head_entropy_mean
   [K]). These require either a new reduction kernel + ISV slots
   or host-readback (forbidden by feedback_no_htod_htoh_only_
   mapped_pinned). Emitted ISV-resident summary leaves only
   (multi_head_policy.{active,k,gating_entropy_floor,head_
   entropy_floor,aux_prior_beta}). Phase 2A-D scope.

## Linked
* Phase 2A-A: 0b3e40150 (MultiHeadPolicy foundation, inert)
* Phase 2A-B: e22da61cf (backward + aux KL prior)
* Plan: docs/superpowers/plans/2026-06-03-multi-head-policy-
  implementation.md
* Spec ADDENDUM: docs/superpowers/specs/2026-06-02-multi-head-
  policy-with-r-multiple.md
* Q-distill pearl: pearl_foxhunt_pi_trained_by_q_distillation_
  not_ppo.md

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-06-03 11:07:27 +02:00
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: 0b3e40150
* Plan: docs/superpowers/plans/2026-06-03-multi-head-policy-implementation.md
* Spec ADDENDUM: docs/superpowers/specs/2026-06-02-multi-head-policy-with-r-multiple.md

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-06-03 09:51:55 +02:00
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 779c03b9d, b=128):
  mean -$6.41M ± $1.39M σ; 62× spread in popart_envelope quartile
  (Q1 -$1.99M / Q4 -$32k); policy TVD asymmetry 0.17 on popart
  axis vs 0.02 on vol axis. Regime-direct gating targets the
  attenuation chokepoint; mixture provides regime-conditional
  capacity that single-head softmax cannot express.

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-06-03 01:41:57 +02:00
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 5cd2f8703 (1B-A foundation). Wires the RolloutBuffer
into the per-step training loop via a new RolloutCollection helper,
gated behind FOXHUNT_USE_ROLLOUT=1 env flag. Mode 0 (default unset) is
bit-equal to HEAD; Mode 1 (flag=1) snapshots per-step data into the
rollout buffer and invokes GAE every T_rollout steps.

This phase is a STRUCTURAL validation, not an empirical pnl test.
Phase 1B-C will replace the per-step training in Mode 1 with multi-
epoch PPO over the rollout buffer.

New components:

* crates/ml-alpha/cuda/rollout_pack.cu (42 LOC):
    Single-thread-per-batch deterministic f32 -> u8 packer for dones.
    Used because RolloutBuffer.dones_d is u8 (per spec for memory) but
    the trainer's dones_d is f32 (per existing pipeline).

* crates/ml-alpha/src/trainer/rollout_collection.rs (375 LOC):
    RolloutCollection helper with three methods:
      - new(ctx): loads rollout_pack cubin
      - snapshot_step(trainer, buf, t, b_size, t_max): DtoD scatter via
        cuMemcpy2DAsync_v2 for rewards/v_t/actions/log_pi/h_t at offset
        b*t_max+t in the [B × T] buffer; rollout_pack kernel for dones
      - copy_bootstrap_v(trainer, buf, b_size): DtoD V(s_T) from
        trainer.v_pred_tp1_d into buf.v_t_bootstrap_d at rollout end
    All transfers use mapped-pinned / DtoD only — no raw memcpy_dtoh
    on regular slices (`feedback_no_htod_htoh_only_mapped_pinned`).

* crates/ml-alpha/examples/alpha_rl_train.rs (+137 LOC at lines 59-66,
  410-466, 591-666):
    - FOXHUNT_USE_ROLLOUT env-flag parse at startup
    - Conditional RolloutBuffer + RolloutCollection allocation
    - Per-step snapshot_step() after step_with_lobsim_gpu
    - Every T_rollout steps: copy_bootstrap_v + compute_gae + stderr log
      `[rollout] step=N GAE complete (T=…, γ=…, λ=0.95): adv_mean=…
       adv_abs_mean=… returns_mean=…`

Adaptations from the dispatch (documented inline):

1. T_rollout fallback to 32 when slot 758 is unbootstrapped. Slot 758
   (RL_PPO_ROLLOUT_HORIZON_INDEX) was allocated in 1B-A but no kernel
   bootstraps it — `read_isv_host(758)` returns 0.0. Binary uses
   `.clamp(32, 1024)`, flooring to T=32 for the smoke. Phase 1B-C will
   add the trainer bootstrap entry (or a controller kernel) to make
   T_rollout=256 the production default.

2. Mode 1 still runs per-step training as a side effect; the rollout
   buffer is populated as a parallel snapshot. The "stop training in
   Mode 1" separation would require a ~1500 LOC refactor of
   step_with_lobsim_gpu (single-step/training-coupled per
   pearl_foxhunt_trainer_is_genuinely_single_step). The dispatch
   explicitly authorized this "skeleton with correct semantics"
   adaptation — the load-bearing 1B-B gate is "collection + GAE
   pipeline works structurally", which is verified. Phase 1B-C will
   replace the per-step training with multi-epoch PPO over the rollout
   buffer.

Validation (all gates PASS):
* cargo build --release --example alpha_rl_train -p ml-alpha: exit 0 (~60s)
* cargo test rollout_buffer_invariants --release: 5/5 PASS (no regression)
* Mode 0 (flag unset, default): ./scripts/determinism-check.sh --quick exit 0
  - DETERMINISTIC: all checksums.* leaves match across all 200 rows
  - pipeline output bit-equal to HEAD 5cd2f8703 baseline
* Mode 1 (FOXHUNT_USE_ROLLOUT=1): smoke completes in 2m 53s
  - 500 train + 100 eval steps, exit 0, completed_clean=true
  - 16 GAE windows logged (T=32 each)
  - adv_abs_mean range [4.4e-3, 1.92e-1] — always > 1e-3 sanity gate
  - eval_summary written: total_pnl_usd=-$101,487.70, n_trades=301, wr=0.346
* Cross-source consistency BOTH modes within $50 (Mode 0: $0.00 diff;
  Mode 1: $0.01 diff). Phase 0 contract preserved.
* Pre-commit hook PASS (0 memcpy_dtoh raw violations, 0 atomicAdd).

Next: Phase 1B-C will (a) bootstrap slot 758 to production T_rollout=256,
(b) replace Mode 1's per-step training with multi-epoch PPO over the
rollout buffer using the existing PPO surrogate + V regression kernels
operating on minibatches of [B × T] flattened to [B*T].

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-06-02 22:42:39 +02:00
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>
2026-06-02 22:09:47 +02:00
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>
2026-06-02 19:55:31 +02:00
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>
2026-06-02 17:56:00 +02:00
jgrusewski
63fc16f173 fix(rl): surfer-scaffold v5.2 — drop wr_excess clamp, full scaffold when novice
alpha-rl-mjsmr (4a4953b01 step 99-249, 2026-06-02): v5.1 with wr_ema=0.248
(below break_even 0.30) produced scaffold_weight = 0.500, not the expected
1.0. Math: `max(0, wr_ema - break_even) = 0` → sigmoid(0) = 0.5 → competence
= 0.5 → w_competence = 0.5. So a novice agent below break-even got
HALF-scaffold instead of full scaffold (entropy dropped to 1.27 at step 249
because of the early directional signal, then ramped back up to 1.96 once
w_decay took over — unstable).

Fix: drop the `max(0, ...)` clamp. Use signed `wr_distance = wr_ema -
break_even` so:

  wr=0.20 → sigmoid(30·-0.10) ≈ 0.05 → w_competence ≈ 0.95 (full scaffold ✓)
  wr=0.30 → sigmoid(0)        = 0.50  → w_competence = 0.50 (half at BE)
  wr=0.40 → sigmoid(30·+0.10) ≈ 0.95 → w_competence ≈ 0.05 (pure pnl ✓)

This is the proper symmetric fade around break_even; the asymmetry was a
v5 original-design error. Combined with v5.1 `w_decay = max(0, 2·frac-1)`,
the scaffold should now:
  - hold near 1.0 while wr < break_even (novice)
  - smoothly fade through 0.5 as wr crosses break_even
  - drop toward 0 as wr settles above break_even AND PH decay subsides

Local invariant test:
  reward_alignment_surfer_scaffold_invariants OK: 251 rows validated
  train[0] bootstrap = 0.966 (unchanged)

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-06-02 00:34:03 +02:00
jgrusewski
4a4953b01f fix(rl): surfer-scaffold v5.1 — decay re-engagement requires EXCESS alertedness
alpha-rl-zf6s5 (38a4aa15b b=1024 fold-1 step 5719 verdict, 2026-06-02):
the v5 scaffold_weight stayed pinned at 1.0 even when wr crossed
break-even (0.306 > 0.30). Root cause: `w_decay = min(1, 2 × frac_alerted)`
saturated at 1.0 because train-time Page-Hinkley naturally alerts on
75-95% of batches (per pearl_edge_decay_detector_phase1_validated_train_also_decays
discovered 2026-06-01); `max(w_competence, w_decay)` then kept the
scaffold engaged forever, defeating the fade mechanism.

FIX: `w_decay = max(0, 2 × frac_alerted − 1)` so re-engagement measures
EXCESS alertedness above 50% baseline, not absolute level:

  frac_alerted   v5 w_decay    v5.1 w_decay
  -----------   -----------   ------------
  0.50          1.0           0.00      ← v5 broken here
  0.78          1.0           0.56
  0.90          1.0           0.80
  1.00          1.0           1.00

At zf6s5 step 5719 (frac_alerted=0.78), v5.1 gives w_decay=0.56 not 1.0;
w_competence=0.455 wins via max(), so w drops to ~0.56 letting the fade
begin properly.

Health signals from zf6s5 (informative even though terminated):
  step 5719: wr=0.306, hold=14.4, entropy=1.82, realized=+$152M
  trajectory: wr crossing 0.30 ✓, hold growing 12→14 ✓, entropy ↓ ✓
  the agent IS becoming competent; the scaffold just couldn't fade.

Local invariant test:
  reward_alignment_surfer_scaffold_invariants OK: 251 rows validated
  train[0] bootstrap = 0.966 (unchanged — boot path identical)

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-06-02 00:22:23 +02:00
jgrusewski
38a4aa15b3 feat(rl): adaptive surfer-scaffold reward shaping (spec v5 A+B combined)
Diagnosis from alpha-rl-8fb55 (fa347e481 b=1024 fold-1, 4500 steps):
the v4 binary `pure_pnl_mode=1` correctly aligned the gradient (Pearson
0.92) but left the agent unable to LEARN. action_entropy stayed at
1.9-2.1 (near uniform log(11)=2.40) — the policy never converged on a
strategy, just oscillated between random trading and all-flat. pnl
bled −$3.9M, wr stuck at 0.30 (random-baseline). The Phase 5 hold-bonus
was NOT pure contamination — it was the INDUCTIVE BIAS SCAFFOLD that
taught a random-init network what to optimize. Pure pnl alone is too
sparse to bootstrap learning in foxhunt's reward-density regime.

REDESIGN — replace binary pure_pnl_mode with continuous adaptive
surfer-scaffold weight w ∈ [0,1] multiplying Phase 5 shaping:
  w=1 → identical to dd049d9a4 surfer-baseline shaping (full scaffold)
  w=0 → identical to v4 pure pnl reward (no shaping)
  0<w<1 → smooth lerp

Controller `rl_surfer_scaffold_controller.cu` writes w each step from:
  - RL_WIN_RATE_EMA_INDEX (677): agent's competence signal
  - RL_CUMULATIVE_DONES_INDEX (660): trade-count confidence gate
  - RL_EDGE_PH_FRAC_ALERTED_INDEX (751): edge-decay re-engagement

Math:
  confidence  = sigmoid((n_trades - warmup) / (warmup*0.3))
  competence  = confidence × sigmoid(k_sharp × max(0, wr_ema - break_even))
  w_competence = 1 - competence
  w_decay     = min(1, 2 × frac_alerted)
  w = clamp(max(w_competence, w_decay), 0, 1)

At random init (n_trades≈0, wr≈0): w → 1 (full scaffold, agent learns
to hold via Phase 5 like dd049d9a4 surfer baseline). After warmup with
wr crossing break-even: w → 0 (pure pnl reward, no Phase 5
contamination for eval-relevant strategies). Edge-decay PH alerts
force w back up if the policy degrades — Option A "rethink strategy"
mechanism without an explicit reset.

Slot reuse: 753 was `RL_REWARD_PURE_PNL_MODE_INDEX` (v4 binary, never
deployed beyond fa347e481 cluster smoke alpha-rl-kwppb/8fb55, both
verdict=fail). v5 renames to `RL_SURFER_SCAFFOLD_WEIGHT_INDEX` with
semantic flip (v5 w=1 ⇔ v4 mode=0; bootstrap 1.0 invariant). 3 new
config slots (754/755/756) for break_even_wr / k_sharpness /
warmup_trades. RL_SLOTS_END 754 → 757.

Phase 5 kernel modifications — multiplicative lerp form preserves
w=1 ⇔ legacy shaping:
  additive a:        r += w × a
  multiplicative m:  r *= (1 + w × (m - 1))

Diag: `rewards.surfer_scaffold_weight` replaces `rewards.pure_pnl_mode`
(1 leaf in, 1 leaf out → EXPECTED_LEAVES still 679).

Invariant test:
  - I1+I2: scaffold_weight ∈ [0, 1] across train + eval (250 rows)
  - I3: scaffold_weight ≥ 0.9 at step 0 (novice agent ⇒ full scaffold)

Local validation:
  reward_alignment_surfer_scaffold_invariants OK: 251 rows validated;
    train[0] bootstrap = 0.966 (matches math: 1 - sigmoid(-3.33)×0.5 ≈ 0.98)
  eval_diag_emission OK: 679 leaves train + eval (schema parity)

Cluster smoke (next): fold-1 b=1024 20k+5k. Falsification:
  PASS: action_entropy drops below 1.7 by step 5000 (policy concentrating)
        AND wr_ema crosses 0.35 by step 10000
        AND surfer_scaffold_weight fades to < 0.5 by step 15000
  FAIL: any criterion below → revisit signal weights or anneal schedule

Spec: docs/superpowers/specs/2026-06-01-reward-policy-alignment-investigation.md §6
Pearls:
  - pearl_reward_signal_anti_aligned_with_pnl (the v4 diagnosis)
  - pearl_pure_pnl_mode_starves_b16_controllers (the v4 cluster verdict)

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-06-01 23:52:14 +02:00
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 22e6ddbca alpha-rl-6kghr and 87a8259c6 alpha-rl-8gtk2):

  Pearson(rewards.sum, Δrealized_pnl_cum_usd) = 0.16 train / 0.31 eval
  sign-agreement (rewards.sum vs Δpnl direction)   = 23-32% across both
  (random baseline 50%)

The shaped reward fed to PPO/Q/V losses is systematically anti-aligned
with profitable trade direction. wr_max=0.418 in train of 87a8259c6
proves the architecture can find profitable patterns; the misaligned
gradient un-learns them. THE single load-bearing bug behind 64 prior
negative-eval commits — controllers were rebalancing around a wrong
optimization target, not a wrong solver.

Mechanism: Phase 5 step 4 of rl_fused_reward_pipeline (per-step hold
bonus = hold_bonus × √hold_time, with hold_bonus=2.0) adds +20/step at
hold_time=100, dominating realized pnl (~$1k scale) by ~2600× per held
trade. V-regression learns inflated baseline → close-event advantages
wrong-signed → PPO un-trains profitable closes.

Implementation (ONE atomic commit, gated on a single ISV slot):
  - RL_REWARD_PURE_PNL_MODE_INDEX (slot 753), RL_SLOTS_END 753→754
  - Bootstrap = 1.0 (NEW path default; mode=0 preserves legacy shaping
    for ONE regression cluster smoke, deleted in 24-48h follow-up per
    feedback_no_feature_flags + feedback_single_source_of_truth)
  - rl_fused_reward_pipeline.cu: Phase 5 steps 1-4 AND Phase 5b
    inventory penalty gated by pure_pnl_mode > 0.5
  - rewards.pure_pnl_mode emitted in diag.jsonl (679 leaves)
  - tests/reward_alignment_invariants.rs: 250-row GPU-oracle invariant
    test (mode==1.0 propagation + abs_max < $25k inflation fence)

Fees: LobSimCuda apply_fill_to_pos deducts cost_per_lot_per_side per
fill into pos.realized_pnl (resting_orders.cu:213-219). realized_pnl_delta
is already net-of-fee; mode=1 cleanly delegates execution costs to
LobSim, inventory limits to Layers 1/2/4 (CMDP + IQN τ + Kelly sizing).

Local validation:
  reward_alignment_pure_pnl_mode_invariants OK: 250 rows validated
  eval_diag_emission OK: 679 leaves train + eval (schema parity)

Falsification (cluster smoke, fold-1 walk-forward 20k+5k b=1024):
  PASS: Pearson(rewards.sum, Δrealized_pnl) ≥ 0.7 train @ steps 1500-2000
        AND ≥ 0.5 eval; AND wr ≥ 0.30 in train post step 1000
  FAIL: any below → controller transient extended to step 3000-4000,
        re-check; if still fails, abandon reward path within 1 retry

Risk A acknowledged (HIGH): 12 adaptive controllers will rebalance
under new sparse-reward distribution. Dominant time constant is
RL_REWARD_CLAMP_V_BOUND_EWMA_ALPHA = 0.001 (τ=1000 steps; ~3τ to 95%
convergence). §2 robustness clause extends verdict window if clamp
bounds haven't stabilized.

Reviewer (sp-critical-reviewer 2 passes): FIX-AND-PROCEED — all v3
BLOCKERs structurally resolved; v4 + 5 LOW edits approved.

Refs:
  spec: docs/superpowers/specs/2026-06-01-reward-policy-alignment-investigation.md
  pearl: pearl_reward_signal_anti_aligned_with_pnl

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-06-01 22:38:47 +02:00
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>
2026-06-01 21:36:46 +02:00
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 dd049d9a4 added tail-event circuit breakers (CMDP DD,
Kelly trap, IQN-τ); none addressed slow-bleed via overtrading on degraded
edge. alpha-rl-glv6n showed eval pnl -$24.85M came from 47K trades at
wr=0.256 with 95 closes/step — 2000× higher close rate than train —
invisible to all existing risk machinery because no single trade tripped
a tail-event threshold.

Page-Hinkley change-point detector on σ_welford-normalized per-trade PnL:

  x_t           = rewards[b] / σ_welford   (slot 725, scale-invariant)
  μ̄_t           = pre-update Welford running mean
  m_t           = m_{t-1} + (x_t − μ̄_pre − δ)    [only after warmup]
  M_t           = min(M_{t-1}, m_t)               [only after warmup]
  ph_stat_t     = m_t − M_t

ISV slots (6 new, 747-752): δ=0.1, λ=5.0, warmup_min=10 (σ-relative
unit-less); fleet aggregates ph_mean (over active batches),
frac_ph_alerted (n_alerted/n_active), frac_ph_warmup (n_warmup/b_size).

Per-batch device buffers (5 new): ph_mu/count/m/mmin/stat. Read+written
only by rl_cmdp_constraints_check kernel.

Reset discipline: all 5 PH buffers added to reset_session_state memset
block alongside existing CMDP per-batch state. PH is PREDICTIVE (predicts
future edge degradation from recent trajectory), so it follows
"RESET PREDICTIVE, PRESERVE NORMALIZATION" per
pearl_popart_reset_at_eval_boundary_shocks_normalization. σ_welford
itself is normalization — it lives in popart EMA which IS preserved
across boundaries.

Kernel ordering note: rl_cmdp_constraints_check fires BEFORE
apply_reward_scale and rl_popart_normalize, so rewards[b] at kernel
entry is RAW USD and the σ slot read is the prior step's σ_welford
(one-step lag, acceptable for diagnostic). Slot 725 chosen over slot 555
(σ_effective) because slot 555 spikes 1000× at the eval shock window —
blinding the detector exactly when bleed is largest.

EXPECTED_LEAVES bump 675 → 678 (+3 diag leaves
edge_ph_{mean,frac_alerted,frac_warmup}).

Local b=16 smoke (100+50): wiring validated. Train phase by step 99
shows ph_mean=3.5 (close to λ=5), frac_alerted=0.25 — detector
firing as designed at active batches. Eval[0] frac_warmup snaps back to
1.0 — confirms reset_session_state extension correctly zeroed all 5
PH buffers at fold/eval boundary (BLOCKER #1 contract validated).

Predetermined falsification (cluster smoke b=1024, per spec §G5):
- frac_ph_warmup drops <0.4 by train step 1000 AND eval step 200
  (≥35% cooldown-suppressed tail acknowledged)
- frac_ph_warmup snaps to ≥0.95 at first eval row (reset wiring test)
- ph_mean rises from <1 to >3 over first 200 eval trades
- frac_ph_alerted > 0.3 by eval step 300

3-strike abandonment: if 3 V2 iterations fail for 3 different parameter
reasons, audit signal source (per feedback_going_in_circles_pattern).

Spec: docs/superpowers/specs/2026-06-01-edge-decay-detector-phase1-diagnostic.md
(v3 after 3 sp-critical-reviewer passes; verdict GO after final review).

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-06-01 21:04:47 +02:00
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>
2026-06-01 16:53:06 +02:00
jgrusewski
f428be794b Revert "feat(rl): B-11-β Q-distill informativeness gate"
This reverts commit b93971726d.
2026-06-01 14:57:41 +02:00
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, c1dc84a34):
target_entropy = 2.3976 / ln(11) = 2.398 max → 99.98% of max with Q_range
13.79 and τ=20.18. Distill at λ=0.224 was applying ~uniform-target KL
regularization across 20k steps → eval pnl -$218M at full run
(alpha-rl-8gtk2).

Gate (per-block, smooth):
  λ_eff = λ_base × max(0, 1 − h_target / ln(N_ACTIONS))^p
where λ_base is the existing KL-target Schulman controller's output
(slot 486, untouched) and p defaults to 2.0 (slot 744).

Bitwise disabled-mode: when RL_Q_DISTILL_INFO_GATE_ENABLED_INDEX = 0.0
the ternary forces gate_factor = 1.0 verbatim and s_lambda_eff = lambda
exactly. Same ISV-toggle pattern as B-7's reward-clamp control (slot 724).

Source delta:
- 3 new ISV slots: RL_Q_DISTILL_INFO_GATE_ENABLED_INDEX = 743,
  RL_Q_DISTILL_INFO_GATE_SENSITIVITY_INDEX = 744,
  RL_Q_DISTILL_LAMBDA_EFFECTIVE_INDEX = 745; RL_SLOTS_END 743 → 746.
- Bootstrap array 230 → 233 (defaults: ENABLED=1.0, SENSITIVITY=2.0,
  LAMBDA_EFFECTIVE=0.0 sentinel; kernel block-0 overwrites every step).
- rl_q_pi_distill_grad.cu: shared s_lambda_eff, per-block target-entropy
  reduction in the existing thread-0 s_pi_target block, gate compute,
  __syncthreads broadcast; grad_distill uses s_lambda_eff. Block-0 emits
  slot 745 alongside existing B-10 G2 emits.
- Diag: policy_diagnostic.q_distill_lambda_effective (672 leaves total).
- New invariant smoke: tests/q_distill_info_gate_invariants.rs (4 gates).

Local 200+50 b=16 fold-1 smoke validates all 4 invariants across 249 rows:
peak h_target/ln(N) = 1.000 (β cause confirmed at smoke scale too); gate
factor range [0, 5.99e-6]; 132 saturated rows (h_frac ≥ 0.999). l_q_b
healthy throughout.

NOT byte-identical to pre-B-11-β runs by design: gated grad_distill
changes π → action distribution → trade outcomes. V4 (spec §4) calls
this out explicitly. Cluster comparison is on aggregate signals (eval
pnl, l_pi trajectory), not exact reproducibility.

Falsification criteria (spec §2.G5) gate the next iteration:
- eval pnl > -$50M  → β was dominant, B-11-β is the fix
- -$50M to -$100M  → β contributed, write B-11-γ (advantage standardization)
- < -$100M          → β not dominant alone; B-11-α + B-11-γ combined

Risk D explicit threshold (spec §8): if λ_base (slot 486) exceeds 2× its
pre-B-11-β baseline EMA sustained 5k steps, the KL-target controller is
in compensatory-ramp territory and B-11-β.1 gates the controller's input
on the gated KL.

Spec: docs/superpowers/specs/2026-06-01-b11-beta-q-distill-informativeness-gate.md

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-06-01 13:00:06 +02:00
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>
2026-06-01 12:25:47 +02:00
jgrusewski
8c7ce02da9 feat(rl): B-10 policy-quality cascade diagnostic
alpha-rl-8gtk2 (B-7+B-8+B-9 run at SHA 87a8259c6) exposed a degenerate
equilibrium at step 6478 (~32% through 20k train): l_pi diverged 15×
from best (stale 478 steps), l_q + l_v bit-flat (stale 524/615 steps),
π near-uniform (entropy 2.36 / ln(11)=2.40 max, Hold-dominant 32%),
Kelly fraction floored at 0.025 with consistently-negative EV.

Three lit-confirmed candidate root causes (perplexity 2026-06-01):
  α Q-collapse via adaptive C51 EWMA atom support — Bellemare projection
    collapses target distribution to one-hot at center atom under
    EWMA-narrowed [V_MIN_eff, V_MAX_eff], yielding near-zero CE without
    Q having any information content.
  β Q→π distill amplification — softmax(Q/τ) is uniform for ANY τ when
    Q is uniform across actions, so high-τ distill = pure max-entropy
    regularization regardless of τ value.
  γ Heavy-tailed advantage normalization — per-batch standardization
    (Schulman 2017 canonical, computed in `rl_advantage_normalize.cu`,
    publishes var to slot 612) produces 5-10σ outliers under B-7's
    unclamped reward distribution → PPO surrogate magnitude explodes.

B-10 ships PURE OBSERVABILITY across 4 groups (G1 Q-dist informativeness,
G2 Q-distill effectiveness, G3 PPO advantage normalization, G4 PPO
surrogate decomposition). No controller, no behavior change, no atom-span
modification. The falsification criteria predetermine the B-11 path the
next cluster run motivates.

Changes:
- 13 new ISV slots (730-742) in isv_slots.rs; `RL_SLOTS_END = 743`.
- 13 bootstrap entries; fixed-size array `[(usize, f32); 217]` → 230 at
  integrated.rs:3394.
- New kernel `rl_q_distribution_stats.cu` — two entry points
  (`rl_q_distribution_per_batch` + `rl_q_distribution_reduce`) computing
  online Q distribution entropy + E_Q range + |E_Q| max via single-block
  tree-reduce (B-9 pattern). Reads online `q_logits_d` + `atom_supports_d`.
- New kernel `rl_ppo_diagnostic_stats_reduce.cu` — cross-batch reducer
  over 4 per-batch scratches (`ppo_a_norm_pb`, `ppo_ratio_dev_pb`,
  `ppo_ratio_clipped_pb`, `ppo_surrogate_pb`) written by surrogate_fwd.
  Reads slot 612 (var_pre_norm) for σ_used = sqrt(var); reconstructs
  |A_unnorm| stats from |A_norm| × σ_used.
- Modified `ppo_clipped_surrogate.cu` — 4 new pointer params for the
  scratches; writes happen in the same `act == 0` branch that already
  owns per-batch reductions, alongside existing loss_pi_per_b. Loss
  reduce path (`ppo_loss_reduce_b.cu`) and ratio path
  (`ppo_log_ratio_abs_max_b.cu`) untouched.
- Modified `rl_q_pi_distill_grad.cu` — adds 2 inline emits (slot 733
  target entropy, 734 target-π entropy diff) in the existing batch-0
  diagnostic block; controller writes (slot 486 λ, slot 487 τ) untouched.
- 5 new `[B] f32` scratch buffers + 2 new function handles in
  `PolicyHead` + 2 new function handles in `DqnHead`. New launch methods
  `launch_q_distribution_stats` (DqnHead) and
  `launch_ppo_diagnostic_stats_reduce` (PolicyHead) following B-9 pattern.
- Trainer wires: G1 launch after `dqn_head.forward(h_t)` at line 6615;
  G3+G4 reducer call right after `surrogate_forward` at line 5211
  (single call site verified — `surrogate_forward` is invoked once per
  step_with_lobsim_gpu pass).
- 14 new diag leaves under `policy_diagnostic` block: 13 from new slots
  + 1 from existing slot 407 (`rl_q_pi_agree_b` was writing every step
  but never reached diag; per B-10 spec §6 Open Decision 4 audit).
- `EXPECTED_LEAVES` 657 → 671 in `eval_diag_emission.rs`.

Local validation (RTX 3050 Ti, 200+50 b=16 fold-1 smoke):
  * Build clean, release binary 1m 43s.
  * Schema parity test path: train = eval = 671 leaves ✓
  * 11/14 new diag fields populating correctly:
    - q_distill_target_entropy = 2.395 (near-uniform ⇒ G2 working)
    - q_pi_agree_ema evolving 0.41 → -0.80 → 0.36 across training
    - PPO scratches show non-zero ppo_a_norm_mean=0.65 at step 50
      (when advantages have signal) and 0 elsewhere (correct: A=0 at
      most steps with replay still warming).
  * 3/14 fields (G1 q_dist_*) return 0 locally — possible sm_86-specific
    runtime issue with the new `rl_q_distribution_stats` kernel; will
    debug on the next cluster run (H100 sm_90) which is the canonical
    diagnostic environment for B-10's intended cluster validation.

Spec: docs/superpowers/specs/2026-06-01-b10-policy-quality-cascade-diagnostic.md

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-06-01 12:20:42 +02:00
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>
2026-06-01 11:31:31 +02:00
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>
2026-06-01 11:00:53 +02:00
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>
2026-06-01 10:31:18 +02:00
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>
2026-06-01 10:15:36 +02:00
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>
2026-06-01 10:05:59 +02:00
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>
2026-06-01 01:19:52 +02:00
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>
2026-06-01 01:13:27 +02:00
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>
2026-06-01 00:57:58 +02:00
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>
2026-06-01 00:32:06 +02:00
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>
2026-05-31 23:35:20 +02:00
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>
2026-05-31 23:09:59 +02:00
jgrusewski
1aa92f57f0 fix(rl): eval-boundary addendum — reward_scale warmed_flag + pos_max_ema rate-cap
The parent eval-boundary fix (72684ed3e) preserved env.max/clamp EMAs but
left the σ explosion intact. Diagnosis from alpha-rl-jnct8 (10k+2500 eval,
fold-1, -$106M eval pnl, σ → 4451 by eval step 5) revealed two pre-existing
mechanisms compounding the eval shock:

ISSUE A — reward_scale snaps to 0.10 at boundary
================================================
rl_fused_controllers' reward_scale block had a bootstrap-fraction-floor
gate: `(cumulative_dones < min_trades) → boot_floor = 0.10`. The intent
was cold-start protection (prevent scale crash before any closed-trade
ground truth). But cumulative_dones (slot 660) is reset by
reset_session_state for Kelly's PREDICTIVE-warmup purpose. At fold
boundary the gate fires again, clamping the preserved scale (0.0046 in
jnct8) UP to 0.10 — a 22× upward jump. Eval rewards are then 22× larger,
overwhelming env.max preservation; σ explodes regardless.

FIX A — decouple via monotonic warmed_flag (slot 716, never reset).
Set ONCE when cumulative_dones first crosses min_trades; persists across
all subsequent fold boundaries. boot_floor reads the flag instead of
re-evaluating trade_count. Math: at cold-start flag=0 → boot_floor=0.10
(original protection preserved). Post-warmup flag=1 → boot_floor=scale_min
(~1e-4) → preserved scale survives boundaries.

ISSUE B — pos_max_ema growth unbounded (train-phase fat-tail spike)
====================================================================
Pre-existing fat-tail behavior: at alpha-rl-jnct8 step 3895 a single
account had pre_clamp scaled reward 724. The clamp's Wiener-α EMA
(α=0.4 floor) admitted 40% of the observation, jumping pos_max_ema
112 → 834 in 5 steps. clamp_win = MARGIN × pos_max_ema followed
magnitude up rather than bounding it; env.max captured the unclamped
reward (121 → 1306). Recovery via slow-decay over 600 steps, but
during the spike PPO gradients were mis-scaled.

FIX B — asymmetric per-step growth cap on pos_max_ema (1.5× max).
Same Schulman-bounded-step pattern as reward_scale's 2% per-step
movement clamp. Decreases unbounded (allows fast recovery from spike).
Bootstrap path unchanged (ema_prev=0 → ema_new=pos_max, no cap).
Math: 5-step max growth = 1.5^5 ≈ 7.6× vs prior uncapped 7.4× in
practice — similar steady-state, bounded transient.

Local validation (b=16, 800+200 fold-1):
- σ preserved across boundary (51.9 → 51.4 at eval[1])
- σ stays bounded in 23-57 range across full eval phase (no 60× explosion)
- scale=0.10 stable across boundary
- warmed_flag stays 0 at b=16 (cumulative_dones never crosses min_trades
  at this scale — flip behavior tested at cluster b=1024)

Spec: docs/superpowers/specs/2026-05-31-eval-boundary-addendum-reward-scale-and-train-spike.md

Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
2026-05-31 22:25:24 +02:00
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>
2026-05-31 21:10:59 +02:00