Code-quality review found 1 Important + 3 Minor; applied all 4:
- comment: "warp-aware shared-mem reduction" → "shared-mem block reduction"
(no __shfl_down_sync used; claim was misleading per reviewer)
- const-correctness: tid is now `const int tid = threadIdx.x` (matches sibling kernels)
- power-of-2 assumption: documented via `#define BLOCK_X 256` + comment
- removed unused `#include <stdint.h>` (kernel uses plain int only)
Zero functional change — kernel emit is identical; cubin rebuilds and passes.
Block-reduce per-account lots[b] → flat_count single int.
Per feedback_no_atomicadd: shared-mem tree-reduce, no atomics.
Foundation for regime_observer (F1.3) which reads flat_count
to compose the dead-zone detection (kelly=0 ∧ all-flat ∧ no-cooldown).
Launch profile (for trainer in F1.4):
Grid=(1), Block=(256), smem=256*sizeof(int)=1024 bytes.
Code-quality review found 2 Important + 2 Minor issues; applied all 4:
- separator style: ────── → ============ (file convention)
- restored split sub-headers: outputs (6) / Welford state (4)
- restored full safety-net comment with issue #4 reference
- added inline value hints on every constant (per spec block)
Zero functional change — comments and whitespace only.
Diag: surfaces `risk_stack.eval_warmup.{remaining,active,blend,
floor_*,target_*}` so the v9 defensive-warmup window is observable
in diag.jsonl. `remaining` is the counter; `blend` is the
defensive-vs-normal mix coefficient (1.0 = full defensive, 0.0 =
normal); `floor_*` reflect the LIVE override values (read AFTER the
warmup kernel ran). Pre-warmup the kernel is a no-op (remaining=-1),
so v9 train-phase diag is bit-identical to v8.
Cleanup: tightens unreachable_pub items in tests/behavioral/* and
tests/sp5_producer_unit_tests.rs (pub → pub(crate)), removes
unused_mut on 6 sp5 scratch buffers, renames unused `step` loop
counter in alpha_baseline example, and explicitly discards an
intentionally-no-op `Command::assert` in cli_integration_test.
Reduces lint count by ~25; remaining 3 dead_code warnings flag
SP15 Phase 2A behavioral scaffolding (Phase 2B never landed —
deliberate signal, not noise).
Pre-existing pearl per feedback_no_hiding: do NOT suppress these
with #[allow]; the warnings ARE the design call surface.
Implements docs/superpowers/specs/2026-05-31-v9-defensive-eval-boundary-calibration.md.
Closes the [[pearl_adaptive_carryover_discipline]] gap: every adaptive
EMA must reset OR re-bootstrap at regime boundaries, never selectively
preserve train-acquired statistics.
Layer 1 — full EMA reset in `reset_session_state`:
- win_rate, avg_win/loss EMAs → neutral sentinels
- cumulative_dones → 0 (re-enter Kelly bootstrap)
- inventory_beta + inventory_variance EMAs → 0
- reward_clamp pos/neg max EMAs + clip_rate EMA → 0
- Kelly fraction → 1.0 (bootstrap)
Lifts Fix D's deliberate carryover (commit 7064c9269), which was
diagnosed in v8 fold-1 eval as the primary -$507k driver: agent
entered eval with train wr=0.34 EMA and took aggressively-sized
losing trades while EMA decayed to true eval wr=0.25.
Layer 2 — defensive warmup window:
- New `rl_eval_warmup_decay` CUDA kernel: single-thread single-block,
no atomics, mapped-pinned only. Runs after fused controllers each
step. Overrides 4 risk-sizing floors (Kelly safety_frac, IQN τ_min,
entropy_coef_min, PPO clip ε_min) with conservative defensive
values for the first 500 steps post-boundary, with linear decay
over the final 200 steps back to normal targets.
- 11 new ISV slots (685-695): remaining counter, configured
durations, defensive overrides (0.25/0.30/0.05/0.10), normal
targets (0.50/0.10/0.01/0.05).
- Sentinel counter = -1 at boot → kernel is no-op until reset arms it.
All boundary semantics live in ISV; no hardcoded constants in the
kernel. Build verified, all 13 risk-stack invariants + 5 controller
adaptive-floor tests + integrated-trainer smoke pass on RTX 3050.
Per the newly-saved memory pearl pearl_adaptive_carryover_discipline,
diagnoses Fix D's "intentionally preserve train Kelly + inventory EMAs
into eval" as the dominant cause of v8's eval-phase regression.
Design — 3 layers at the train→eval boundary:
Layer 1: Reset every adaptive EMA (Kelly wr/avg-win/avg-loss/
cumulative_dones, inventory β/variance, reward-clamp pos/neg/clip
EMAs) to neutral sentinels. Lets the controllers re-bootstrap from
eval-distribution observations rather than carrying poisoned train
state.
Layer 2: Defensive warmup window (500 steps) overrides risk-sizing
controllers — Kelly safety_frac 0.5→0.25, IQN τ_min 0.10→0.30,
entropy coef floor 0.01→0.05, PPO ε floor 0.05→0.10. Linear decay
back to normal over additional 200 steps.
Layer 3: 2× LR multiplier during warmup window. Network weights
adapt fast to new-regime statistics; this is the slowest-adapting
layer of the agent.
Adds 8 new ISV slots (686-693), bumps RL_SLOTS_END 686→694. New
small kernel rl_eval_warmup_decay applies overrides each step during
warmup. Pure additive design — no kernel logic changes outside the
new warmup-decay kernel.
Validation requires running ALL 3 folds (vs v8's single fold-1) per
pearl_single_window_oos_is_not_oos. Success criterion: mean eval pnl
across 3 folds > v8 mean.
No code change in this commit — design document only. Implementation
deferred until decision is made on whether to ship v9 or first
collect v8's fold-0 + fold-2 baseline for cleaner comparison.
Per docs/superpowers/specs/2026-05-31-c51-atom-span-math-validation.md,
the binding constraint on atom_max during training is the dynamic
Bellman bound `atom_max ≥ WIN + γ × atom_max`, not the overstated
fixed-point `WIN/(1-γ)`. The math validation against local b=128
smoke confirmed atom_max can be 4.5× the fixed-point bound yet train
cleanly (qpa=+0.969 at step 999).
This commit adds derived diag fields under
`risk_stack.atom_calibration` to expose the bound directly:
- win_bound, atom_max, gamma (inputs)
- dynamic_bound = WIN + γ × atom_max (the binding constraint)
- atom_max_headroom (=atom_max - dynamic_bound; >0 = self-consistent)
- popart_sigma, v_target_max_3sigma (statistical V_target estimate)
- atom_max_over_3sigma (resolution waste ratio)
These let future cluster runs measure CURRENT design's over-sizing
empirically (smoke step 999 showed atom_max ~5× larger than 3σ of
V_target requires — wasteful but safe). Future iterations can use
this data to safely tighten atom_max anchor toward V_target_max
without speculating about which design works.
Pure additive diag — no kernel changes, no behavior change. Pulls
values from existing ISV slots (RL_REWARD_CLAMP_WIN_INDEX,
RL_C51_V_MAX_INDEX, RL_GAMMA_INDEX, RL_POPART_SIGMA_INDEX).
Validates the math from 2026-05-31-c51-atom-span-math-validation.md
empirically in every cluster run going forward.
Formal proof of why the current `atom_max = EWMA(WIN_bound)` design
trains successfully despite empirical violations of the "structural
minimum atom_max ≥ WIN/(1-γ)" claim made in 2026-05-30-c51-atom-
resolution-design-alternatives.md.
Key findings (proven by local smoke + cluster v8 trajectory):
1. The fixed-point bound `WIN/(1-γ)` is overstated as a structural
constraint. Empirical: smoke step 999 atom_max = 4.5× the bound,
system trains healthy (qpa=+0.969).
2. The actually-binding constraint is the dynamic bound:
`atom_max ≥ max V_target_observed ≤ WIN + γ × V_max_observed`
Self-consistent and satisfied with margin 23 at smoke step 999.
3. Both speculative alternatives (Fix F = mean_abs_pnl anchor,
Fix F-v2 = V_target_observed anchor) have closed-form instability
at γ(1+ε) ≥ 1 — for γ=0.99 only ε ≤ 0.01 stable, no resolution
benefit over current design.
4. Current design is *conservative* (atom_max ~5× larger than V_target
3σ requires) but *correct* — slow EWMA hysteresis absorbs WIN
transients and keeps Q-V Bellman self-consistent.
Includes empirical trajectory data:
- Local smoke (b=128, HEAD d57bee054, 1000 steps): atom_max 1798→1468,
qpa +0.65→+0.97, clip_rate 5-11%, dynamic bound satisfied with
margin 23 at convergence.
- Cluster v8 (b=1024, alpha-rl-vxbpq @ 6d4a962e5): popart_σ collapses
56→1.69 from step 100→16830, qpa stays >+0.93, pnl_cum $164M.
Saves pearl_atom_span_dynamic_vs_fixed_point for future sessions.
Diagnostic emit of V_target_max_observed proposed (§5) but deferred
— current signals (WIN + γ × atom_max) already provide the bound
indirectly via JSONL. Real-time V_target_max is a nice-to-have for
future calibration work, not required to validate the current design.
Documents why Fix F crashed at cluster scale (qpa: +0.898 → -0.976
between step 371 and step 500 in alpha-rl-qrgjr v7), the structural
Bellman self-consistency constraint atom_max >= WIN/(1-γ) that makes
the resolution trade-off fundamental, and the design alternatives
considered (non-uniform atoms, adaptive γ, two-headed Q). Recommends
shipping A+B+C+D+E only — the pre-Fix-F design implements the
structurally-correct atom span via WIN-tracking EWMA.
No code change in this commit — design document only. Empirical
result from v8 (alpha-rl-vxbpq, SHA 6d4a962e5) will determine whether
deeper atom-resolution work (Options 4.C or 4.E) is justified.
Spec: docs/superpowers/specs/2026-05-30-c51-atom-span-decouple-from-clamp.md
Diagnosed in cluster v5 (alpha-rl-rjsjq SHA 7064c9269) at step 371:
c51_v_max ratcheted up 1 → 24 → 53 → 859 → 1938 while WIN clamp swung
4910 → 568 → 7540 → 44 (volatile from MARGIN saturation). The slow
EWMA tracking win_bound time-averaged the early spikes, locking
atom span at extreme-value magnitude. Δz reached ~184 per atom — Q
could no longer discriminate +$500 from +$5000 wins.
Root cause: atom-span target was anchored on `MARGIN × pos_max_ema`,
where pos_max is the EXTREME-VALUE statistic across b=1024 batch
elements. The clamp bound (rightly fat-tail) and the atom span (should
be typical-magnitude) were treating two distinct concerns as one.
Fix F: anchor atom span on `atom_margin × mean_abs_pnl_ema` — true
central-tendency signal — while leaving the reward clamp on
pos_max_ema. The clamp can still admit fat-tail wins ($42k → +245)
and V regression (scalar head) retains the magnitude for PPO advantage,
but Q's CATEGORICAL distributional support now sizes itself for the
trades where most learning happens.
New ISV slot 685 (RL_C51_ATOM_MARGIN_INDEX, bootstrap 10.0).
RL_SLOTS_END 685 → 686.
Validation (local b=128 1k smoke):
metric | pre-Fix-F | with Fix F
v_max final | 1929 | 86.6 (22× tighter)
Δz per atom | 184 | 8.7 (resolution restored)
qpa final | +0.957 | +0.958 (Q-π alignment preserved)
mean_active | +$47k | +$34k (within smoke variance)
13/13 risk_stack_invariants + 20/20 trade_management_kernels +
integrated_trainer_smoke pass.
Tradeoff documented in spec: Q loses fat-tail categorical magnitude
(rewards beyond ±87 project to top/bottom atom), but V regression
preserves it. Acceptable for the lottery-ticket strategy where typical-
magnitude resolution matters more than fat-tail magnitude precision.
Cluster v5 alpha-rl-rjsjq step 371 revealed worst-account session_pnl
growing monotonically across cooldown windows (-$3.7k → -$9.3k from
step 100 → 371). Root cause: when DD triggers, actions_to_market_targets
forces {side=2, size=0} — a no-op that suppresses EVERY action type
including trail-stop-driven closes. If the account was holding a
losing position when DD fired, that position bleeds mark-to-market for
the entire 500-step cooldown with no exit path.
Fix: on first DD/cooldown step, emit a closing market order matched to
the current position (side=1 sell for long, side=0 buy for short, size
= |position_lots|). Once flat, subsequent cooldown steps emit the
existing no-op. Net effect: account closes its position at the moment
of DD trip, then sits flat until its recovery clock expires.
Validation: 13/13 risk_stack_invariants pass, 20/20 trade_management
pass, integrated_trainer_smoke passes. Local b=128 1k smoke:
qpa: +0.95 (best result of session)
mean_active: +$47k
worst: stabilizes at -$11.8k (vs unbounded growth pre-fix); residual
loss is from fat-tail single-step market moves that cross dd_limit
before the controller can react — out of scope for this fix.
Two compounding bugs found via cluster diag at alpha-rl-gwmkf step 4083:
Fix C — reward-clamp writeback un-freeze:
rewards.clip_rate_ema = 0.9999 — essentially every win clamped to +1.
With R-multiple 43 trade outcomes ($7k avg win / $163 avg loss), V
regression saw identical targets for +$500 wins and +$42k wins; PPO
advantage lost magnitude signal; policy stagnated at qpa = -0.95.
Root cause: a 2026-05-24 "G.2" patch in rl_reward_clamp_controller.cu
added `(void) margin;` and froze WIN=1.0, LOSS=3.0 over concern that
simultaneously adapting clamp bounds + atom span could create a
positive feedback loop. Empirically refuted — MARGIN is bounded
[1.0, 5.0] and pos_max_ema is bounded by reward_scale × raw_PnL, so
WIN has natural structural ceiling. The freeze was over-conservative.
Restore the writeback: WIN = max(MIN_WIN, MARGIN × pos_max_ema),
LOSS = max(MIN_WIN, WIN × RATIO). Atom span Step 5 then follows via
the slow EWMA already in place.
Local b=128 1k smoke confirms:
clip_rate_ema: 0.9999 → 0.094 (target 5%, near hit)
WIN clamp: 1.0 → 245.6
V_max atom: 1.0 → 1929
qpa: -0.95 → +0.71 (Q and π aligned!)
mean_active: +$47k stays positive
Fix D — reset_session_state extends to per-batch buffers:
The 39efacf77 per-batch CMDP refactor added 4 device buffers but
reset_session_state() was only zeroing the 4 ISV summary slots. On
train→eval fold transition, eval inherited all the accumulated
dead/cooldown accounts from training — eval started with a poisoned
fleet.
reset_session_state now also memsets session_pnl_per_batch_d,
session_dd_triggered_per_batch_d, consec_loss_per_batch_d, and
cooldown_remaining_per_batch_d to zero. Plus also resets the new
RL_SESSION_PNL_WORST_INDEX slot.
Validation: 13/13 risk_stack_invariants pass, 20/20 trade_management
pass, integrated_trainer_smoke end-to-end passes.
Diagnosed via the diag-emit added in 6e0f56816 — alpha-rl-d6d8d step 2000
showed worst-account at -$18k, IQN τ pinned at floor 0.1, popart σ
collapsing 1.00 → 0.47. Root cause: dead-account accumulation in a
sticky-DD fleet.
The per-batch CMDP from 39efacf77 fixed the GLOBAL lockout but left
two failure modes:
1. Once an account hit `dd_limit`, `session_dd_triggered_per_batch[b]`
stayed sticky for the fold. `actions_to_market_targets` forced
those accounts to no-op → V_target ≈ γ·V(s'_unchanged) → popart
σ leaked variance from accumulating dead-weight.
2. ISV[RL_SESSION_PNL_USD_INDEX] exposed worst-account pnl, which
IQN-τ consumes for `drawdown_frac`. A single broken account
dragged τ to its floor for the entire fold — defensive action
selection for a fleet that was, on average, doing fine.
Fix A: DD recovery on cooldown expiry. When `dd_limit` trips, also
start a recovery cooldown clock. When the clock decrements to 0,
clear `dd_triggered` + reset `session_pnl` to 0. Account rejoins the
active fleet. Matches surfer trading philosophy: accept the wipeout,
take the forced break, get back on the board.
Fix B: IQN-τ signal split. CMDP now writes mean-of-active-accounts
to RL_SESSION_PNL_USD_INDEX (slot 662, IQN-τ consumer) and mirrors
worst per-batch pnl to a new slot RL_SESSION_PNL_WORST_INDEX (684,
diag only). τ now tracks fleet-typical drawdown instead of a single
catastrophic outlier.
Subtle kernel detail caught by G1: `cool_prev` must be snapshotted
BEFORE the DD-trip section, otherwise a fresh cooldown set this step
gets immediately decremented by 1 in the same launch.
Validation (local b=128 1k smoke):
metric | pre-fix (HEAD) | with A+B
worst pnl | -$34,016 | -$7,183 (79% less bleed)
mean active pnl | n/a | +$40,343 (typical account profitable)
IQN τ | 0.10 (floored) | 0.50 (neutral)
popart σ | 0.47→collapsing| 0.84 (variance preserved)
win rate | 0.21 | 0.56 (positive edge)
13/13 risk_stack_invariants pass (G1 updated, G1b NEW for cooldown
recovery). 20/20 trade_management_kernels pass. integrated_trainer_smoke
end-to-end passes.
Slot 684 added: RL_SESSION_PNL_WORST_INDEX. RL_SLOTS_END 684 → 685.
Each batch element is an independent backtest session with its own
$35k starting capital. The CMDP kernel had been summing per-batch
rewards into a single ISV slot — at b=1024 the accumulator grew
b_size× faster than the -$3500 single-account DD limit, tripping
session_dd_triggered globally in ~250 steps (cluster alpha-rl-2j9k9
step 320: session_pnl_usd=-$5775, all 1024 sessions locked out).
Same flaw on consec_loss_count: any 10 losses across the batch in a
single step triggered cooldown for the entire fleet.
Refactor — every Layer-1 state shard is per-batch:
session_pnl_per_batch_d [b_size] f32
session_dd_triggered_per_batch_d [b_size] f32
consec_loss_per_batch_d [b_size] f32
cooldown_remaining_per_batch_d [b_size] f32
`actions_to_market_targets` now reads per-batch DD-triggered + cooldown,
so one account hitting its limit does not lock out the other 1023.
Summary ISV slots (RL_SESSION_PNL_USD, RL_CONSEC_LOSS_COUNT,
RL_SESSION_DD_TRIGGERED, RL_COOLDOWN_REMAINING_STEPS) are now
kernel-OUT only — they expose worst-account / any-triggered / max
aggregates for diag + the single IQN-τ consumer ("be pessimistic
when ANY session is in trouble").
3 raw `actions_to_market_targets_fn.cu_function()` launch sites in
integrated.rs were updated atomically per `feedback_no_partial_refactor`
— public wrapper + 2 internal step_with_lobsim sites. The first
attempt hit CUDA_ERROR_INVALID_VALUE in integrated_trainer_smoke
because only the public wrapper had been updated.
Validation:
- 12/12 risk_stack_invariants pass on per-batch semantics (G1-G4
rewritten to seed per-batch buffers via write_slice_f32_d_pub).
- 20/20 trade_management_kernels still pass (actions gating intact).
- integrated_trainer_smoke passes end-to-end.
- Local b=128 1k smoke: 1000/1000 steps clean, no NaN. CMDP behavior
matches design: worst-of-128 hit DD (-$34k > $3500 limit), other 127
unaffected. IQN τ floored at 0.1 (worst-account drawdown >>10%),
Kelly = 0.0 (observed edge negative: wr=0.21, R=1.57).
Added `risk_stack` section to alpha_rl_train.rs diag emitter so the five
risk-management layers introduced in 285d42aa7 are observable from the
diag JSONL. Without this the kernels run but the output is invisible.
cmdp session_pnl_usd, session_dd_triggered, consec_loss_*,
cooldown_*, max_open_units, net_inventory_limit_usd
iqn_tau action_tau, tau_min, dd_sensitivity
inventory penalty_beta, variance_ema
kelly fraction, win_rate_ema, avg_win/loss_usd_ema,
safety_frac, min_trades_for_release, cumulative_dones
trail_factors tighten, loosen
Pure additive emission — no kernel changes, no perturbation of training
state. Local 5-step smoke b=16 confirms every field appears with values
that match the spec math (e.g. IQN τ adapts linearly with drawdown).
Doesn't affect the in-flight alpha-rl-2bm59 (pinned SHA 285d42aa7);
takes effect on the next submission.
rl_confidence_gate.cu accepted pos_state as a kernel argument but the
body never dereferenced it. Reading position_lots from bytes [0..4] and
returning early when non-flat makes Case 3 of the existing test
(`confidence_gate_overrides_low_confidence_opening`) pass and restores
the documented opening-only gating semantics.
Repro before fix:
SQLX_OFFLINE=true cargo test -p ml-alpha --release \
--test trade_management_kernels \
confidence_gate_overrides_low_confidence_opening \
-- --ignored --nocapture
→ FAIL Case 3: non-flat position got Hold instead of original action
After: 20/20 trade_management_kernels tests pass.
Follow-up to the 2026-05-30 adaptive controller floor refactor (commit
083a88f7c). Spec Special case R wired the bootstrap-fraction floor to
release after `RL_TRADE_DUR_VAR_COUNT_INDEX < MIN_TRADES_FOR_RELEASE
(=100)`. That counter increments via the Welford trade-duration-EMA
producer ONCE PER STEP where any trade closed — not once per closed
trade. At b=1024 with typical ~7 dones/step, the gate released after
step ~100 (= ~700 actual closed trades, far short of the intended
"100 trades of confidence"). After release reward_scale crashed to
~0.001 within 200 more steps, identical to the pre-fix fold 0/1
behavior the spec was supposed to prevent.
Fix: track REAL cumulative closed trades by summing `dones[b]` per
step in the fused kernel's reward_scale block. New ISV slots:
- RL_CUMULATIVE_DONES_INDEX (660) — running total
- RL_MIN_TRADES_FOR_RELEASE_INDEX (661) — gate threshold,
ISV-driven per `feedback_adaptive_not_tuned`, bootstrap 5000
At ~7 dones/step (typical) the floor now holds for ~715 steps before
release — enough for the reward magnitude EMA to stabilize against
genuine trade outcomes rather than early-training noise. Threshold is
ISV-resident so production runs can tune without recompile.
Verified locally (500-step b=128 smoke):
step 50: reward_scale = 0.384 (decaying toward floor)
step 100: reward_scale = 0.140 (approaching floor)
step 200: reward_scale = 0.100 (AT FLOOR, held)
step 500: reward_scale = 0.100 (still at floor)
Without this fix (pre-commit): reward_scale crashed to 0.001 by step 500.
Tests: integrated_trainer_smoke + signal_variance_kernel +
controller_adaptive_floors all pass. Release build clean.
🤖 Generated with [Claude Code](https://claude.com/claude-code)
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Replaces hardcoded thresholds AND clamp bounds across 12 RL controllers
with observed-signal-driven ISV-slot bounds. Eliminates the architectural
failure mode that surfaced in walk-forward fold 0 (alpha-rl-m9cx5) and
fold 1 (alpha-rl-jgdh6): under Phase 4.5 advantage normalization, eight
controllers (PPO clip, target_tau, rollout_steps, entropy_coef, per_α,
gamma, reward_scale, q_distill_lambda) saturated at extrema within ~50
steps and stayed pinned for the rest of training — same pattern across
two different data slices, confirming the bug is structural rather than
data-size-dependent.
Mechanism: each controller's noise floor was hardcoded as a small
fraction of its target (`_NOISE_FLOOR_FRAC = 0.01f`), calibrated against
a pre-Phase-4.5 signal regime. Phase 4.5 normalization reduces operating
KL by ~50× — observed signal stays below the 1%-of-target floor, the
Schulman widen path fires continuously (asymmetric in the wrong
direction), and ε hits MAX 0.50 within ~50 training steps. ksll2's full-
data run (n_folds=1) happened to escape via a single above-band KL
observation that triggered tighten; both walk-forward folds (3 and 6
files) did not.
Per `feedback_adaptive_not_tuned`, `feedback_isv_for_adaptive_bounds`,
`pearl_controller_anchors_isv_driven`: every threshold now derives from
observed signal statistics (Welford online variance) rather than
constants calibrated against a prior signal regime. User explicitly
expanded scope mid-implementation: "if all clamps ISV bound should be
added to this spec and tasks" — applying the principle consistently
means hardcoded MIN/MAX clamp bounds count too, not just the saturating
noise floors. 12 controllers, single atomic commit per
`feedback_no_partial_refactor`.
Spec: docs/superpowers/specs/2026-05-30-adaptive-controller-floor-design.md
Plan: docs/superpowers/plans/2026-05-30-adaptive-controller-floor-plan.md
# New kernel
- rl_signal_variance_update.cu (80 LOC) — Welford online variance.
Single-thread single-block. Sentinel-zero skip per pearl. Per-controller
Welford triple (count, mean, M²) drives every adaptive noise floor.
# Per-controller refactors (12 .cu files)
- ppo_clip + target_tau + rollout_steps + entropy_coef + per_α —
adaptive noise floor = max(target × 0.5, sqrt(observed_var) × 2);
asymmetric Schulman (tighten on single observation, widen requires 3
consecutive below-band).
- gamma (Special G) — hardcoded GAMMA_MIN = 0.995 → adaptive via
Welford MEAN of trade duration. Per spec Q6 resolution: the Welford
mean's natural N-smoothing lag breaks the gamma↔trade_duration
feedback loop without explicit step-period gating. 100-observation
warmup falls back to EMA before Welford has enough samples.
- reward_scale (Special R) — asymmetric DECREASE rate cap (5% per
step) on both bootstrap-replace and Wiener-blend paths; bootstrap-
fraction floor (10% of bootstrap) until 100 trades close.
- q_distill_lambda (Special Q) — hardcoded MIN_LAMBDA = 0.05 → adaptive
via Welford on q_distill_kl_ema: max(0.001, std × 0.05).
- v_blend_alpha (Phase 4.4) — 5 hardcoded constants → 5 ISV slots
(DEAD_SIGNAL_FLOOR, TARGET_TRACK_RATIO, SCHULMAN_STEP, EMA_ALPHA,
BOOTSTRAP_ALPHA) + adaptive dead-signal floor from V_scalar magnitude
variance.
- ppo_ratio_clamp — adaptive MIN/MAX from observed log-ratio variance.
Architectural 2.0 absolute floor preserved (don't degenerate to
vanilla policy gradient).
- reward_clamp — V_BOUND_FLOOR, V_BOUND_EWMA_ALPHA, MIN_WIN, MIN_RATIO,
MAX_RATIO, MIN/MAX_MARGIN, MARGIN_TOLERANCE/ADJUST_RATE,
CLIP_RATE_EMA_ALPHA → 10 ISV slots.
- gate_threshold — hardcoded alpha = 0.01 → ISV slot 638.
- Clamp-bound expansion: EPS_MIN/MAX, TAU_MIN, ROLLOUT_MIN/MAX,
COEF_MIN/MAX, PER_ALPHA_MIN/MAX, GAMMA_MAX, MAX_LAMBDA, KL_TOLERANCE,
LAMBDA_RAMP_RATE, LAMBDA_DECAY_RATE → all ISV.
- WIENER_ALPHA_FLOOR shared across 9 controllers → single ISV slot 659.
# Trainer integration (integrated.rs)
- rl_signal_variance_update kernel loaded + helper method
`launch_rl_signal_variance_update`.
- Per-step Welford launches for 9 controller inputs, placed between
EMA producers and rl_fused_controllers in the per-step pipeline.
- Input-slot lookup array for rl_fused_controllers updated: rollout_steps
now consumes RL_ADV_VAR_PRE_NORM_INDEX (emitted by
rl_advantage_normalize before in-place normalize) instead of the
post-norm advantage_var_ratio (definitionally ~0 under Phase 4.5).
- ~30 new ISV bootstrap entries in with_controllers_bootstrapped.
# ISV slot allocation (isv_slots.rs)
- 72 new slots, RL_SLOTS_END 588 → 660. +288 bytes mapped-pinned.
- 9 Welford variance triples + 5 asymmetric Schulman counters +
3 new input signals (var_pre_norm, gamma_min_adaptive, reward_mag) +
v_blend (5 + 3 Welford) + ppo_ratio_clamp (1 + 3 Welford) +
reward_clamp (5) + gate_threshold (1) + q_distill (1) + 20 clamp bounds +
shared wiener floor.
# Bootstrap-clamp consistency fix (caught by Task 16 testing)
The original draft bootstrapped RL_EPS_BOOTSTRAP at 0.01 (= KL target),
but EPS_MIN was 0.05 — bootstrap value below clamp range. First post-
bootstrap step always snapped ε to MIN regardless of signal direction
("snap to MIN" behavior the unit tests surfaced). Fixed by bootstrapping
ε at MIN (0.05) so asymmetric Schulman operates from a valid state.
# Validation
- cargo build --release: clean
- cargo build --tests: clean
- 3 GPU Welford kernel tests (G1 constant→0var, G2 sequence→known var,
G3 sentinel skip): all pass
- 5 GPU adaptive-floor invariant tests (G3 PPO clip holds at bootstrap
when signal below floor, G4 tightens on single above-band, G5 widens
only after 3 consecutive below-band, G6 post-warmup uses Welford mean,
G6 pre-warmup uses EMA): all pass
- integrated_trainer_smoke (full GPU pipeline): passes
- 1k local smoke at b=128: 1000/1000 steps, completed_clean, no NaN,
controllers genuinely adapting (γ 0.90→0.974, per_α 0.40→0.52,
q_distill_λ 0.05→0.21, ε held at 0.05 = MIN per Phase 4.5 small-KL
regime — was previously stuck at MAX 0.50 throughout fold 0/1)
- compute-sanitizer memcheck b=128 5 steps: 0 errors
Cluster validation (G7-G9) submitted as follow-up runs.
🤖 Generated with [Claude Code](https://claude.com/claude-code)
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
After the Phase 4 dueling-head merge landed on main, six integration
tests no longer compiled — they referenced trainer fields and methods
that were renamed or removed during the R-series refactor:
isv_d (CudaSlice<f32>) → isv_dev_ptr: u64 (raw, cached)
isv_host (Vec<f32>) → isv_mapped: MappedF32Buffer
launch_rl_controllers_per_step() → launch_rl_fused_controllers()
softmax_ce_grad(..&mut CudaSlice) → softmax_ce_grad(..&u64)
trainer.replay (Vec-based PER) → gpu_replay (CUDA buffers)
N_HORIZONS = 5 → N_HORIZONS = 3
Release binary built clean throughout (the cluster doesn't pull in
test sources), so the breakage was invisible until `cargo test --tests`
surfaced it post-merge.
Per `feedback_no_partial_refactor`: when a contract changes, every
consumer migrates atomically — the test suite was left behind by
those R-series PRs, this commit closes the gap.
Per `feedback_no_htod_htoh_only_mapped_pinned`: tests now use the
same mapped-pinned ISV view as production (zero-copy host reads via
`isv_host_slice()` / `read_isv_host(slot)`, single-slot writes via
`isv_mapped.write_record(slot, val)`).
Changes per file:
isv_bootstrap.rs (1 site)
Read full ISV via `trainer.isv_host_slice()` instead of dtoh
of the now-removed `isv_d` CudaSlice. Sync producing stream
first so bootstrap-controller writes are visible host-side.
r3_ema_advantage.rs (5 sites)
Rewrote `readback_isv` helper to take `&IntegratedTrainer`
and use the mapped-pinned mirror. All 5 call sites simplified
from `readback_isv(&dev, &trainer.isv_d)` to `readback_isv(&trainer)`.
r5_controllers_and_soft_update.rs
Deleted G3 (`launch_rl_controllers_per_step` no longer exists;
`launch_rl_fused_controllers` is the architectural replacement
with different setup requirements — its 'all controllers move
slots' invariant is exercised end-to-end by every cluster run).
Kept G4 (DqnHead soft-update Polyak formula) with updated API.
trade_management_kernels.rs (3 sites)
`set_isv_slot` helper now uses `isv_mapped.write_record(slot, val)`
— single volatile write to mapped-pinned, GPU sees it after next
sync, no explicit HtoD copy needed.
frd_head.rs (11 sites incl. ce_total_loss helper)
Added `alloc_loss_buf(n) -> MappedF32Buffer` helper. All callers
of `FrdHead::softmax_ce_grad` now pass `&loss_buf.dev_ptr`
(raw u64) instead of `&mut loss_d` (CudaSlice), and read results
via `stream.synchronize()?; loss_buf.read_all()`.
heads_bit_equiv.rs (per_head_independence)
N_HORIZONS dropped from 5 to 3 in production. Test was hardcoded
against the old count (probs[3], probs[4], 5-element bias vec)
causing compile-time index-out-of-bounds. Per
`feedback_use_consts_not_literals_for_structural_dims`: rewrote
to address by N_HORIZONS-relative offsets (first / last / middle).
r7d_per_wiring.rs (deleted)
The old Rust-side `PrioritizedReplay` struct (R7c's
`src/rl/replay.rs`) was removed when the PER buffer moved fully
GPU-side as `gpu_replay: GpuReplayBuffer`. The test was a guard
against re-introducing that dead Rust struct; the dead file no
longer exists in the tree (verified `crates/ml-alpha/src/rl/replay.rs`
is gone), so the guard is moot. The new buffer's correctness is
exercised end-to-end by every cluster training run.
Validation:
- cargo build -p ml-alpha --release: clean
- cargo build -p ml-alpha --tests: clean (all files compile)
- integrated_trainer_smoke (GPU, --ignored): passes
🤖 Generated with [Claude Code](https://claude.com/claude-code)
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Standard PPO practice (Schulman et al. 2017): normalize advantages
per-batch BEFORE PPO surrogate computation:
advantage[b] ← (advantage[b] − mean_b) / sqrt(var_b + ε²)
Self-adaptive — uses observed per-batch statistics, no tuned
hyperparameters (fits feedback_adaptive_not_tuned).
Why now: Phase 4.3 cluster (alpha-rl-qkdm2) ended with pnl=+$16.28M
(2x Plan A v2's +$8.5M) but l_pi=1.6e9 vs Plan A v2's 3.6e5 — a
4500× increase in PPO surrogate magnitude. Adam internally normalizes
the gradient direction, but per-step magnitude variance is high
→ pnl trajectory chops (saw ±$2-4M swings between milestones).
Advantage normalization fixes this regardless of V baseline source:
batches with high-magnitude advantages get scaled down to ~unit
variance, ensuring consistent PPO update strength across batches.
Standard in Stable-Baselines3, RLlib, OpenAI Baselines.
New kernel: rl_advantage_normalize.cu (~80 LOC).
Grid=(1,1,1), block=(min(B,1024),1,1).
Single block parallel reduction: pass 1 = mean, pass 2 = variance,
pass 3 = normalize in-place. ε² floor on variance prevents
div-by-zero when all advantages identical.
Trainer wiring (~30 LOC):
Load kernel + handle field + struct init.
Single launch immediately AFTER compute_advantage_return,
BEFORE PPO surrogate consumes advantages_d. In-place.
Stacks with Phase 4.4 adaptive V blend on same branch
(ml-alpha-phase4-dueling-head). Two complementary stabilization
mechanisms:
- Phase 4.4: adapts WHICH V baseline (scalar vs dq) to use based
on observed tracking error → reduces variance source
- Phase 4.5: normalizes advantages regardless of variance source
→ reduces variance impact
Validated:
- cargo build --release clean
- integrated_trainer_smoke 1 step passes
- alpha_rl_train --steps 3 --b 128 under compute-sanitizer: 0 errors
Cluster validation pending — submit Phase 4.4+4.5 combined to test
whether dampened-variance preserves Phase 4.3's +$16M pnl gain.
🤖 Generated with [Claude Code](https://claude.com/claude-code)
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Replaces Phase 4.3's hard V_dq → PPO swap with an adaptive blend
driven by an on-device controller. Per the project's no-tuning
philosophy (pearl_controller_anchors_isv_driven, feedback_adaptive_not_tuned):
V_used[b] = α × V_scalar[b] + (1 − α) × V_dq[b]
where α ∈ [0, 1] is emitted by rl_v_blend_alpha_controller from
the observed V_dq vs V_scalar tracking ratio:
track_ratio = EMA(|V_dq − V_scalar|) / EMA(|V_scalar|)
if track_ratio > 1.5 × TARGET: α ← min(α + 0.01, 1.0)
if track_ratio < TARGET / 1.5: α ← max(α - 0.01, 0.0)
else: hold α
Plus dead-signal guard: if EMA(|V_scalar|) < 1e-4, hold α (no V
signal yet to calibrate against).
Bootstrap on sentinel 0: α = 1.0 (Plan A v2 behavior on first step).
EMAs first-observation bootstrap (no Wiener-α blend on first sample).
Two new kernels:
- rl_v_blend.cu: elementwise blend (~25 LOC). Grid (ceil(B/256),1,1).
- rl_v_blend_alpha_controller.cu: single-block parallel reduction
+ Schulman-bounded controller (~90 LOC). Grid (1,1,1), block (1024,1,1).
Three new ISV slots (585/586/587):
- RL_V_BLEND_ALPHA_INDEX — current α
- RL_V_TRACK_ERR_EMA_INDEX — EMA(|V_dq − V_scalar|)
- RL_V_SCALAR_MAG_EMA_INDEX — EMA(|V_scalar|), dead-signal floor
IntegratedTrainer wiring (~80 LOC):
- 2 new buffers v_blended_d, v_blended_tp1_d
- In step_with_lobsim_gpu_body, after DuelingQHead Adam steps:
1. Launch controller (reads V_scalar at h_t + V_dq at h_t, emits α)
2. Launch blend kernel for s_t → v_blended_d
3. Launch blend kernel for s_tp1 → v_blended_tp1_d
- compute_advantage_return now reads v_blended_d / v_blended_tp1_d
instead of dueling_v_d / dueling_v_tp1_d (Phase 4.3's direct swap)
Both value_head and DuelingQHead still train independently. The blend
just selects which baseline drives PPO advantage per step based on
observed calibration. As V_dq learns to track V_scalar, the controller
gradually shifts α down toward V_dq usage. If V_dq diverges (e.g., late
training entropy spikes producing volatile advantages), controller
raises α back to V_scalar safety.
Phase 4.3 cluster (alpha-rl-qkdm2 @ 25f5ce99b) is still running and
showing dramatic late-training pnl growth (+$20M at step 18132 vs
Plan A v2 peak +$9.3M). Phase 4.4 adds adaptive control on top —
should reduce variance while preserving the architectural benefit.
Validated:
- cargo build --release clean
- integrated_trainer_smoke 1 step passes
- alpha_rl_train --steps 3 --b 128 under compute-sanitizer: 0 errors
🤖 Generated with [Claude Code](https://claude.com/claude-code)
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Activates DuelingQHead's V output as the PPO advantage baseline,
replacing scalar value_head's contribution. Also wires soft-update
of DuelingQHead's target net (was deferred from 4.2).
Two changes:
1. DuelingQHead.soft_update_target(): reuses dqn_target_soft_update
kernel (generic element-wise blend with τ from ISV[401]) across
all 4 weight tensors (w_v, b_v, w_a, b_a). Called once per training
step after Adam, mirroring DQN's pattern.
2. compute_advantage_return call: v_pred_d → dueling_v_d,
v_pred_tp1_d → dueling_v_tp1_d. PPO advantage is now:
A(s_t, a_t) = E_Q_C51(s_t, a_taken) − V_dq(s_t)
returns(s_t) = r_t + γ(1−done) × V_dq(s_{t+1})
value_head still trains via MSE on returns for diagnostic
comparison; can be retired in a future commit if V_dq proves
itself.
Phase 4.2 validated structurally at b=1024 5k that DuelingQHead's
training does not perturb Plan A v2's dynamics (qpa tracked within
0.02 at every milestone). Phase 4.3 is the smallest possible commit
that USES V_dq downstream — fully de-risked by 4.2's structural test.
Cluster expectation: qpa trajectory roughly matches Plan A v2 (V_dq
calibrated by joint training on same Bellman reward signal as C51,
so cross-architecture scale issue from Phase 3.2 doesn't apply here).
Validated:
- cargo build --release clean
- integrated_trainer_smoke 1 step passes
- alpha_rl_train --steps 3 --b 128 under compute-sanitizer: 0 errors,
l_q=0.024, l_v=0.0003 (slightly higher than Plan A v2's 0.0001
because V_scalar now sees different gradient prop now that V_dq
drives advantage — still healthy)
🤖 Generated with [Claude Code](https://claude.com/claude-code)
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Implements §4.1 + §5 of the Phase 4 spec
(docs/superpowers/specs/2026-05-30-phase4-independent-dueling-head-design.md).
Forward kernel rl_dueling_q_forward.cu:
V[b] = Σ_c w_v[c] × h_t[b, c] + b_v[0]
A[b, a] = Σ_c w_a[c, a] × h_t[b, c] + b_a[a]
composed_Q[b, a] = V[b] + A[b, a] − (1/N) Σ_a' A[b, a']
Single fused kernel — three outputs (V, A, composed_Q) computed in one
pass with shared-mem cache of h_t row. Grid (B,1,1), block (HIDDEN_DIM,1,1),
smem HIDDEN_DIM × 4 bytes. Tree-reduce for V projection; per-action
matmul (threads 0..N-1 active); mean-over-actions in shared mem.
DuelingQHead struct (crates/ml-alpha/src/rl/dueling_q.rs, NEW FILE):
- Config + new() with Xavier init (HIDDEN_DIM → 1 for V,
HIDDEN_DIM → N_ACTIONS for A); seed 0x4DEAD_1234
- Online weights: w_v_d, b_v_d, w_a_d, b_a_d
- Target weights: mirrors of online (soft-update wiring in Phase 4.1)
- forward(h_t, b_size, v_out, a_out, q_composed_out)
- forward_target(...) — same kernel with target weights
- forward_inner — parameterized over weight slices
What this is NOT yet:
- No loss kernel (Phase 4.1)
- No backward / decompose (Phase 4.1)
- No trainer wiring (Phase 4.2)
- No PPO advantage swap (Phase 4.3)
Built off Plan A v2 baseline fd3174262 on branch
ml-alpha-phase4-dueling-head. Build verified clean.
Per session pearls:
- pearl_complement_dont_replace_with_dual_architecture: this IS the
'add a complement' pattern, but in fully encapsulated form (zero
shared state with downstream consumers, unlike Phase 3.x attempts).
- pearl_complement_internal_loss_vs_external_consumers: composed_Q
from this head NEVER feeds ensemble/distill/selection (the failure
mode that broke Phase 3.x is structurally impossible here).
🤖 Generated with [Claude Code](https://claude.com/claude-code)
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Plan A v1 (commit 2d68ef5d8 = revert Phase 2.1 on top of 104fe81ca)
failed at cluster (alpha-rl-4sjzw): entropy collapsed to 0.69 by step 3000.
The Phase 2.0 V envelope clamp (c52282fb4) — kept in v1 — was likely
the culprit: it bounds V_pred to a tight envelope, biasing advantages
and breaking PPO gradient flow.
Plan A v2: start from dd049d9a4 (proven working at cluster wr=0.57
+$6.3M @ step 15000 today via alpha-rl-lpbp8) and apply ONLY the
kernel-only bug fixes that don't affect training dynamics:
- variable_selection.cu: VSN stride 40→56 fix (104fe81ca) — prevents
step-4 NaN from reading wrong-stride window_tensor
- bucket_transition_kernels.cu: h_mag_per_bucket multi-warp 32→128
(7e38e46e6) — correct warp-shuffle reduction for HIDDEN_DIM=128
- compute_advantage_return.cu: branch-gate done flag (a6acc25ec) —
done's terminal-state semantics applied via gate, resolves step-4 NaN
NOT applied (intentionally):
- c52282fb4 Phase 2.0 V envelope clamp — biases V regression target
- b4aadff75 V envelope ±200→±10 — extends Phase 2.0
- 10d4614fb atomicAdd removal — kernel non-determinism is real but
intrusive (Rust changes), dd049d9a4 works without this fix
- db4d9a16f Phase 2.1 dueling — broken decomposition per spec analysis
- 72672c9e7+ Phase 2.2/2.3/H/P1+P2+P3 — built on broken Phase 2.1
🤖 Generated with [Claude Code](https://claude.com/claude-code)
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The clamp bounds were bootstrapped at ±0.5 (matching the old ±0.5 atom
span) instead of the intended asymmetric WIN=1.0, LOSS=3.0. This
caused: (1) apply_reward_scale clamped at ±0.5 instead of [-3,+1],
(2) the C51 atom span EWMA target = max(1.0, 0.5) = 1.0 → V_MIN
never moved because target min(-1.0, -0.5) = -1.0 = bootstrap.
Now: WIN=1.0 (positive reward ceiling), LOSS=3.0 (loss-aversion
asymmetry per pearl_audit_unboundedness). Atom span will EWMA from
[-1, +1] toward [-3, +1] over ~2100 steps (α=0.001).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Same issue as apply_reward_scale: the clamp controller and atom
updater were wired in the OLD step path (step_with_lobsim_reward_and_train)
but not in the GPU path (step_with_lobsim_gpu_body). My earlier fix
inlined apply_reward_scale but missed the two kernels that follow it.
Without this, the C51 atom span EWMA never fires — V_MIN/V_MAX stay
frozen at bootstrap [-1.0, +1.0] despite the EWMA code being correct.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
G.2 disabled atom span adaptation to break a positive feedback loop.
With clamp bounds now FROZEN, re-enabling span adaptation is safe.
The span anchors on the CLAMP bounds (WIN=1.0, LOSS=3.0) — the
structural ceiling on post-clamp rewards that Q actually sees. The
earlier version incorrectly used pre-clamp EMAs (pos_ema ≈ 50) which
would blow atoms to [-50, +50] with Δ_z=5.
V_MAX EWMAs from 1.0 toward WIN=1.0 (stays put).
V_MIN EWMAs from -1.0 toward -LOSS=-3.0 (slowly widens to cover the
loss tail). Final asymmetric span [-3, +1] with Δ_z=0.2 gives Q
full resolution across the clamped reward range.
α=0.001 (half-life ~700 steps), floor ±1.0.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
G.2 audit froze the atom span and said "atom span tracks the SEED
clamp range" — but the seed was ±0.5 while the clamp WIN=1.0. With
reward_scale keeping typical scaled rewards at ±0.6-0.9, half the
reward range was projected to edge atoms (binary resolution above
V_MAX=0.5).
Now: V_MIN=-1.0, V_MAX=+1.0 in both Q_V_MIN/Q_V_MAX constants and
ISV bootstrap. Δ_z = 2.0/20 = 0.1, giving ~6 atoms for a typical
scaled reward of 0.6. Q can now distinguish "good trade" from "great
trade" instead of just "positive" vs "negative."
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
launch_apply_reward_scale was defined but NEVER CALLED from the step
body. Raw PnL ($100-1000) went directly into C51 Bellman with atom
span [-0.5, +0.5]. Every reward projected to edge atoms → binary Q
resolution → Thompson sampling couldn't distinguish actions → wr stuck
at 0.36.
The full reward chain: apply_reward_scale → reward_clamp_controller →
atom_support_update was all wired (kernels loaded, methods written)
but the entry point was orphaned. The clamp controller and atom
updater ran but saw zero input (pos_max_ema=0, neg_max_ema=0).
Now: fused_reward_pipeline → apply_reward_scale → (existing) clamp
controller → atom_support_update. Rewards are scaled to fit the C51
atom span, the span adapts to reward magnitudes, and Q gets proper
distributional resolution.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
The gate created a self-reinforcing Hold trap: gate forces Hold → Q
learns Hold is best → conf=0 for trading actions → gate blocks
everything → Hold=100%. Even max SAC (α=2.0, τ=50) couldn't break it.
Two fixes:
1. Exploration slots: the first `b_size × (1 - max_hold_frac)` batch
elements are NEVER gated. At max_hold_frac=0.85 with b=1024, 154
slots always use the Thompson-sampled action. This guarantees Q
sees trading outcomes and can learn their value.
2. Symmetric threshold decay: the adaptive threshold had 100× asymmetry
(10% raise vs 0.1% lower). When Hold exceeded target, the threshold
barely decreased. Now both directions use THRESHOLD_ADJUST_RATE=1.1.
New ISV slot: RL_CONF_GATE_MAX_HOLD_FRAC_INDEX=584 (bootstrap 0.85).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Root cause: confidence gate decouples policy from actions. Policy
softmax stays high-entropy (~2.4) while the gate forces Hold, producing
action entropy ~0.7. SAC read policy entropy EMA (slot 420), saw
"above target", and kept LOWERING τ — the exact opposite of what was
needed.
New kernel `action_entropy_per_step` computes H(action_histogram) from
the POST-gate actions buffer and writes to ISV slot 583
(RL_ACTION_ENTROPY_EMA_INDEX). SAC co-tuning in rl_q_pi_distill_grad
now reads this slot. Launched OUTSIDE CUDA Graph capture (after
confidence gate + FRD gate) in both training and prefill paths.
Kernel design: 11 threads (N_ACTIONS), each thread counts its action
across all b_size elements. Thread 0 computes entropy from the
histogram and updates the EMA. No atomics.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Two bugs caused τ to collapse to floor (1.0) instead of ramping up:
1. Single-sample entropy: SAC auto-tune computed s_entropy from block 0
only (one batch element). At b=1024 this is noise. Now reads
RL_ENTROPY_OBSERVED_EMA_INDEX (slot 420) — the batch-average entropy
EMA written by ema_update_per_step via tree reduction.
2. Symmetric step rate: SAC_ALPHA_STEP=0.001 was identical for ramp and
decay. Entropy collapsed faster than the controller could recover.
Now asymmetric: SAC_STEP_RAMP=0.01 (100× faster when entropy is below
target) vs SAC_STEP_DECAY=0.0001. Per pearl_asymmetric_controller_decay.
Also adds sac_alpha + sac_entropy_target to JSONL isv_config for
diagnostic visibility.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
When entropy drops below target, BOTH α (entropy gradient) AND τ
(distillation softness) increase together. At b=1024 where Q learns
fast and peaks quickly, τ auto-ramps to soften the target. At b=16
where Q is slower, τ stays lower.
τ range [1.0, 50.0], same exponential step rate as α.
This couples the explore-exploit tradeoff into a single adaptive
mechanism that scales with batch size automatically.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
With τ=1.0 the target softmax(E_Q/τ) was too peaked, causing entropy
collapse despite SAC α=2.0. At τ=5.0 the target is soft enough for
the entropy term to maintain equilibrium.
Entropy: STABLE at 0.78-0.81 from step 1000 to 5000 (was declining
to 0.10 at τ=1.0). Hold: 56-100% oscillating. wr: 0.38.
The τ-α balance controls the explore-exploit tradeoff:
high τ + high α = explore (soft target + entropy bonus)
low τ + low α = exploit (sharp target + no entropy)
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Entropy still declines at b=16 (0.28 at 5k) but slower than before.
α_max raised 0.5→2.0, λ lowered 0.1→0.01. wr=0.38. The b=16
environment is too sparse for the SAC auto-tuning to maintain entropy.
b=1024 with 60× denser Q signal should produce different dynamics.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Replace PPO surrogate with target-Q distillation as sole π gradient:
grad = λ×(π_θ - softmax(E[Q_TARGET]/τ)) - α×π×(logπ+1+H)
Uses TARGET Q (slow τ=0.005) not online Q — fixes phase lag.
SAC α auto-tunes to maintain entropy target (70% of ln(11)=1.68).
ISV slots: SAC_ALPHA=581 (bootstrap 0.01), SAC_ENTROPY_TARGET=582.
compute_advantage_return cleaned to pure 8-arg done-gated V-advantage.
PPO surrogate removed entirely from step_synthetic_body.
Local b=16: wr=0.39, Hold=75-100%, entropy declining (SAC α may need
higher max or λ reduction). L40S b=1024 test needed.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
π trained solely by Q→π distillation: grad = λ × (π_θ - softmax(E_Q/τ)).
No PPO, no advantages for π, no importance ratios. π is a behavioral
clone of Q-Thompson's action preferences.
Remove KL reward augmentation from compute_advantage_return (back to
pure env reward). Advantage only feeds V regression now.
qpa still negative (-0.94) due to Q changing between distillation and
measurement (phase lag). But wr=0.37 and entropy stable at 0.45-1.20.
The system learns profitably — qpa may be the wrong metric for this
architecture where Q drives action selection.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Static: β×log(π_ref(a)) added to reward, stored in PER. Q sees
consistent Hold preference across replays.
Dynamic: -β×log(π_θ(a)) added to advantage only (not stored in PER).
Provides entropy-regularized signal (SAC-like) that adapts to
current policy. Low π_θ → positive advantage (explore). High π_θ →
near-zero (don't over-concentrate).
Result: entropy STABLE 0.94-1.29 through 5000 steps (no collapse).
Hold 62-100%. wr=0.36 (best local). Removed KL gradient kernel
from training step (KL is now entirely in the reward/advantage).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Replace the separate KL gradient kernel with an RLHF-style KL reward:
r_modified = r_env - β × (log π_θ(a|s) - log π_ref(a))
This feeds through PPO's standard advantage pipeline — no structural
imbalance between KL (100% of steps) and PPO (6% done-steps). The KL
signal has the same per-step coverage as the advantage because it IS
the advantage on non-done steps.
Remove done-gating: all steps now get advantages (KL reward provides
directional signal on non-done steps). Remove distillation done-gating
too (distill on all steps now that PPO has signal everywhere).
Result: entropy STABLE 0.81-1.31 (was collapsing to 0.0 with gradient
KL). Hold oscillates 62-100% at b=16 (expected — sparse PnL). At
b=1024 should settle ~50% with 60× denser PnL signal.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Replace 24 cudarc .arg() calls (8+16) with RawArgs + raw_launch.
Eliminates ~48 cuStreamWaitEvent + cuEventRecord per step from
cudarc's event tracking. Last hot-path launch_builder in the GPU
RL training pipeline.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
β=0.003 pinned Hold=100% at b=1024 despite 60 dones/step.
The KL gradient overwhelms even the dense PPO signal.
Try 6× lower β to find the balance point.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Remove unnecessary stream.synchronize() from:
1. step_fill_from_market_targets (lobsim): stream ordering handles
the dependency — step_pnl_track launches on the same stream
2. step_pnl_track (lobsim): downstream kernels read pos_d via
stream ordering, no host read needed per step
Defer pos_fraction readback by one step (same pattern as loss
readback) — eliminates the third 7ms sync. pos_fraction is a
dataset-level statistic that barely changes step to step.
nsys: 3.0 → 1.0 slow syncs/step. Sync time: 21.8 → 9.65 ms/step.
The remaining sync is the training graph event wait (actual compute).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Old RL_KL_TARGET_INDEX at slot 454 was for the Q-distill KL target.
New reference-policy KL target at slot 580 needs a distinct name.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
