Replaces the flawed Phase F + G arc preserved on branch
`ml-alpha-phase-f-g-flawed` (commits 99a125cdb..b3808a5ac). The prior
attempt shipped a trainer with multiple production-blocking defects:
1. ISV[400..406] uninitialised → kernels read γ=0, ε=0, entropy=0
2. rl_reward_scale_controller drifted to 1e3 on no-trade steps
3. 6 controllers exist as .cu but never launched
4. Target net never soft-updated (τ has no consumer)
5. step_with_lobsim violated feedback_cpu_is_read_only with host
Thompson sampling + EMA tracking + advantage/return loops
6. "toy" framing leaked into production (alpha_rl_train.rs shipped
with next_snapshots=snapshots — the F.4 next-state code path was
a no-op until that one issue got caught mid-review)
7. No NaN abort in production CLI
The convergence-gate fixtures (dqn_toy/ppo_toy → renamed
dqn_reward_signal/ppo_reward_signal on the flawed branch) hid every
defect because MockLobEnv is state-invariant with horizon=1.
The rebuild is GPU-pure: kernel-driven action sampling, kernel-driven
EMA tracking, kernel-driven advantage/return, ISV bootstrap at trainer
construction, all 7 controllers wired with device-resident EMA inputs,
target-net soft update consumer, NaN abort, no LobEnv trait (drives
LobSimCuda via the existing decision-policy kernel pattern).
Sequenced as R1..R9 with falsifiability gates G1..G7 that exercise
the specific failure modes the convergence-gate fixtures couldn't
catch. Calendar ~8.5 dev days.
Also adds memory pearl
feedback_extending_existing_code_audits_for_existing_violations
capturing the lesson: extending pre-existing CPU/orphan-controller
violations is how this happened.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Closes Phase E of the integrated RL trainer plan
(docs/superpowers/plans/2026-05-22-integrated-rl-trainer.md). The
integrated trainer can now be exercised end-to-end on a real (or mock)
LobSim environment.
What this commit lands:
- LobEnv trait in src/rl/reward.rs — narrow contract over apply_snapshot
+ submit_action + step_event. Chosen over a direct
ml-backtesting → ml-alpha dep because ml-backtesting already depends
on ml-alpha (loader + trunk reuse); a reverse direct dep would cycle.
The simulator-side `impl LobEnv for LobSimCuda` completes the wire
from ml-backtesting in a follow-up (the trait surface is intentionally
small and lobsim-agnostic so other env implementations can plug in).
- MockLobEnv in the same module — deterministic toy bandit
(action 5 → +1, else → -1, done = true every step). Powers the
dqn_toy / ppo_toy / integrated_trainer_smoke gate tests.
- IntegratedTrainer::step_with_lobsim — one training step driven by a
real LobEnv. Forwards encoder, forwards Q + V + π, reads logits to
host, Thompson-samples action per batch
(pearl_thompson_for_distributional_action_selection), drives the env,
derives real reward / advantages / returns / log_pi_old, and delegates
to step_synthetic for the per-head backward + Adam + encoder backward
(single source of truth per feedback_single_source_of_truth_no_duplicates).
- IntegratedTrainer::eval_expected_q_per_action +
IntegratedTrainer::eval_policy_probs_per_action — public eval helpers
the gate tests use to inspect the trained Q-distribution / policy
without re-implementing device readback in test crates.
- dqn_toy / ppo_toy / integrated_trainer_smoke — bodies filled with the
real training loop driven by MockLobEnv::toy_bandit. The convergence
gates assert argmax_a E[Q(s, a)] == LongSmall and mode π(s) == LongSmall
after 300 step_with_lobsim calls. Tests are #[ignore]-gated for CUDA
availability per the project's test discipline; activate via
`cargo test -p ml-alpha --test dqn_toy -- --ignored` on a CUDA host.
- rl_rollout_steps_controller.cu — ISV[404] producer. Rollout length
adapts to var(advantage)/|mean A|: noisy advantages → grow rollout
(more samples per PPO update), stable → shrink (fresher data).
Bootstrap 2048 (PPO default), bounds [256, 8192], Wiener-α blend with
floor 0.4 per pearl_wiener_alpha_floor_for_nonstationary.
- rl_per_alpha_controller.cu — ISV[405] producer. PER priority exponent
α adapts to TD-error kurtosis EMA: heavy tails → raise α to concentrate
on the informative tails, light tails → keep α near the canonical 0.6.
Bootstrap 0.6 (Schaul 2016), bounds [0.3, 1.0].
Both controllers honour pearl_first_observation_bootstrap (sentinel
0.0 → first-emit writes bootstrap, subsequent emits Wiener-α blend) and
pearl_controller_anchors_isv_driven (no hardcoded constants — every
adaptive hyperparameter sourced from ISV).
build.rs registers both kernels and bumps cache-bust v23 → v24.
Phase F follows with the reward-shaping ISV controller (RL_REWARD_SCALE_INDEX=406)
+ per-trade PnL extraction calibration. Phase G adds the Argo workflow
+ dispatcher. Phase H runs the actual training + backtest smoke and tests
the G8 gate (profit_factor > 1.0).
Verification:
- cargo check --workspace --lib clean (no warnings)
- cargo test -p ml-alpha --lib: 66 passed (was 63; +3 mock_bandit_*
tests in rl::reward), 0 failed, 6 ignored
- cargo build -p ml-alpha --test dqn_toy --test ppo_toy
--test integrated_trainer_smoke clean
- integrated_trainer_loss_lambdas_default_equal_weight (non-ignored)
still passes
Refactors the encoder backward to make its INPUT contract explicit so
the integrated RL trainer (and any future caller) can seed it without
going through the supervised GRN/aux backward path.
Refactor:
- New field PerceptionTrainer::grad_h_new_per_k_d: [K × B × HIDDEN_DIM]
buffer. The Loop-1 GRN head backward kernel writes pure per-K head
contributions into this buffer (called with grad_h_carry = nullptr,
the kernel's existing null-check path); the Loop-2 CfC step backward
reads slot k after folding in the recurrent grad_h_carry_d via
aux_vec_add_inplace. Replaces the pre-E.3a single-slot
grad_h_new_d ping-pong (field removed).
- dispatch_train_step's fused (GRN bwd → CfC bwd) reverse K-loop split
into two passes:
Loop 1 (head backwards) → grad_h_new_per_k_d (slot k, no carry)
Loop 2 (encoder backward) reads grad_h_new_per_k_d, folds in
grad_h_carry_d via the existing aux_vec_add_inplace kernel, then
runs cfc_step_bwd; CfC carry semantics byte-identical to before.
- New private helper dispatch_encoder_backward containing the entire
encoder backward kernel-launch chain: Loop-2 K-loop CfC bwd → 3D
transpose → attn-pool bwd + reducer → LN_b bwd + 2 reducers →
Mamba2 L2 bwd → LN_a bwd + 2 reducers → Mamba2 L1 bwd → VSN bwd +
2 reducers → reduce_axis0 for CfC param grads (4 launches) →
encoder Adam steps (CfC ×4 + Controller B + LN ×2 + VSN + attn_q +
Mamba2 L1/L2 grouped). Both supervised and RL paths call this same
helper — single source of truth per
feedback_single_source_of_truth_no_duplicates.
- New public method backward_encoder_with_grad_h_t seeds the per-K
buffer from a caller-provided grad_h_t at slot K-1 (the only slot
the RL heads consumed h_t from) and dispatches the shared encoder
backward helper. All other slots stay zero — those positions had
no downstream loss signal.
- Reducer closure at the old fused-loop site split per Option (a) in
the SDD: reduce_encoder lives in the helper (CfC ×4), reduce_heads
stays in dispatch_train_step (GRN ×10).
Integration:
- IntegratedTrainer::step_synthetic now calls
backward_encoder_with_grad_h_t with the accumulated
grad_h_t_combined_d (Q + π + V contributions with loss-balance λ
scaling). The encoder now learns from the full per-head gradient —
closes the last deferred item from Phase E.2-DEFER (item 1: encoder
backward integration).
Byte-identical supervised behavior preserved: all 63 ml-alpha lib
tests pass after the refactor (baseline 63, post-refactor 63). The
K-loop split + per-K buffer seeding contract is mathematically
equivalent to the prior fused-loop + single-slot ping-pong:
Old: GRN bwd writes grad_h[k] = head_contrib[k] + carry[k+1] (folded
in via kernel's grad_h_carry arg); cfc bwd reads grad_h.
New: Loop 1 writes grad_h_new_per_k[k] = head_contrib[k] (nullptr
carry); Loop 2 at slot k does grad_h_new_per_k[k] += carry[k+1]
via aux_vec_add_inplace, then cfc bwd reads the slot. Same final
value before cfc_step_bwd consumes it.
Capture-graph safety preserved: the new aux_vec_add launches per K
iteration are kernel-only (no host branches, no host mallocs, no
event tracking) per pearl_no_host_branches_in_captured_graph.
Companion to E.2 (commit 2665669b5) and E.2-DEFER (commit 7356e3c7b).
Phase E.3b follows with LobSim integration + toy bandit activation +
2 more controller kernels (rl_rollout_steps + rl_per_alpha).
Closes 3 of 4 deferred items from commit 2665669b5; item 1 (encoder
backward) is deferred to a follow-up commit (see Notes below).
Item 2 — Bellman target projection kernel:
New `bellman_target_projection.cu` replaces the host-side stand-in
`build_synthetic_bellman_target`. Reads γ from
ISV[RL_GAMMA_INDEX=400] per pearl_controller_anchors_isv_driven.
Standard C51 categorical projection with linear interpolation onto
the discrete support [V_MIN, V_MAX]; no atomicAdd (per-source-atom
serial shared writes bracketed by __syncthreads, Q_N_ATOMS=21
inner-loop syncs).
Also adds `dqn_select_action_atoms` in the same TU — selects per-batch
action-row atoms from the full target-net output. Both kernels exposed
via DqnHead::project_bellman_target / DqnHead::select_action_atoms +
DqnHead::forward_target (new — target-network forward using
w_target_d / b_target_d).
IntegratedTrainer::step_synthetic now consumes the kernel path:
forward_target → select_action_atoms → project_bellman_target →
backward_logits. The `build_synthetic_bellman_target` host function is
DELETED.
Item 3 — Per-head LR ISV controller:
New `rl_lr_controller.cu` emits per-head learning rates to ISV slots
[412..417]: RL_LR_BCE_INDEX, RL_LR_Q_INDEX, RL_LR_PI_INDEX,
RL_LR_V_INDEX, RL_LR_AUX_INDEX (RL_SLOTS_END bumped 412 → 417).
Bootstrap 1e-3 per pearl_first_observation_bootstrap; Wiener-α blend
with floor 0.4 per pearl_wiener_alpha_floor_for_nonstationary.
IntegratedTrainer launches the controller at the top of step_synthetic
AFTER the encoder forward; it then DtoH-mirrors ISV and mutates each
per-head AdamW.lr field before the Adam steps. The PHASE_E2_DEFAULT_LR
constant is DELETED. The `lr` parameter on step_synthetic is gone
(greenfield — no caller override).
The controller currently emits the constant LR_BOOTSTRAP target; the
signal-modulated variant (gradient-norm EMA-driven LR adaptation) is
reserved for Phase E.3+ — kernel header documents the upgrade path
and the 5 diagnostic-signal args are already plumbed.
Item 4 — PPO V-loss canonicalization:
Deletes `returns_/v_pred/loss_v` from ppo_clipped_surrogate_fwd
(kernel + ppo.rs::surrogate_forward signature). V gradient now flows
ONLY through the dedicated `v_head_fwd_bwd` kernels per
feedback_single_source_of_truth_no_duplicates. PPO surrogate now
handles policy loss + entropy bonus only. The IntegratedTrainer's
surrogate_forward call site no longer passes returns/v_pred/loss_v.
ISV slot table (rl/isv_slots.rs):
408..411 loss-balance λ (BCE/Q/π/V) ← existing
412 RL_LR_BCE_INDEX ← new
413 RL_LR_Q_INDEX ← new
414 RL_LR_PI_INDEX ← new
415 RL_LR_V_INDEX ← new
416 RL_LR_AUX_INDEX ← new
417 RL_SLOTS_END ← was 412
Notes — item 1 (PerceptionTrainer::backward_encoder_with_grad_h_t)
deferred:
The IntegratedTrainer.step_synthetic path now COMBINES the Q/π/V
grad_h_t contributions into a dedicated `grad_h_t_combined_d` slot
via the existing `grad_h_accumulate_scaled` kernel (loss-balance λ
weighted), which is the prerequisite plumbing for the encoder
backward entry point. The PerceptionTrainer-side method that consumes
this slot — CfC bwd + Mamba2 bwd + encoder Adam step on caller-
provided grad_h_t — is a multi-thousand-line refactor of
perception.rs::dispatch_train_step (~7,400 lines) and is being
sequenced as a standalone follow-up commit to keep the kernel
signature changes here (items 2/3/4) reviewable in isolation.
All three landed items are greenfield replacements per project policy
(feedback_no_legacy_aliases, feedback_no_partial_refactor): no
deprecated fields, no const-or-ISV fallback shims, no parameter
backward-compat. ISV is the single source of truth for all adaptive
parameters per pearl_controller_anchors_isv_driven.
Validation:
cargo check -p ml-alpha --lib ✓
cargo check -p ml-alpha --tests --examples ✓
cargo check --workspace --lib ✓
cargo test -p ml-alpha --lib (63 passed, 6 ignored) ✓
cargo build -p ml-alpha --features cuda (new cubins built) ✓
Companion to E.2 (commit 2665669b5). Encoder backward + Phase E.3
LobSim integration follow in subsequent commits.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Replaces Phase E.1's placeholder host-side loss scalars with actual GPU
kernel calls. Each of Q, π, V heads now runs:
- Forward kernel on h_t → head outputs (q_logits / pi_logits / v_pred)
- Backward kernel chain → grad_w / grad_b / grad_h_t (per-batch scratch
reduced by reduce_axis0)
- Adam step using the existing project-wide adamw_step cubin via the
AdamW wrapper (each head owns 2 instances — one for w, one for b)
New CUDA kernels:
- v_head_fwd_bwd.cu: scalar V head forward + MSE backward. Per-batch
loss scratch + per-batch grad_w / grad_b scratch — caller reduces
via reduce_axis0. Per-(batch, c) grad_h_t writes are sole-writer so
no atomicAdd needed.
- grad_h_accumulate.cu: element-wise grad_h_encoder += λ × grad_h_head
combiner. Loaded and ready; consumed by Phase E.3 once the encoder
backward entry point lands.
Extends existing kernels (additive, doesn't break prior callers):
- dqn_distributional_q.cu: adds `dqn_grad_w_b_h_t` that maps the C51
bwd kernel's grad_logits output into per-batch grad_w / grad_b
scratch + OVERWRITE grad_h_t. Mirrors aux_heads_bwd's thread-role
pattern (one thread per HIDDEN_DIM channel, sole writer per slot).
- ppo_clipped_surrogate.cu: adds `ppo_policy_logits_fwd` (standalone
linear-projection forward producing pi_logits for the surrogate
kernel to consume) AND `ppo_grad_w_b_h_t` (same backward pattern as
the DQN extension, with output dim N_ACTIONS=9).
All new kernels honour feedback_no_atomicadd: per-batch grad scratch
reduced by the existing reduce_axis0 path, never atomics. The Phase C/D
loss-scalar atomicAdds (dqn / ppo bwd accumulators) stay as their
loud-flagged deferral.
Per-head Adam state:
- DqnHead: 2 AdamW instances (w_d, b_d). Re-uses the existing
adamw_step cubin from `trainer::optim::AdamW`.
- PolicyHead: same.
- ValueHead: same.
- Each AdamW owns independent m / v buffers and a device-resident
step counter (per pearl_no_host_branches_in_captured_graph: counter
advancement is a captured kernel, not host scalar).
PerceptionTrainer: adds `h_t_view()` accessor — borrows the h_t_d
field that `forward_encoder` populates. IntegratedTrainer uses this
to dispatch Q/π/V head kernels on the same encoder representation
without re-borrowing self.perception mutably between the encoder
forward and the head dispatches (disjoint field borrows).
ENCODER BACKWARD: Phase E.2 DEFERS the encoder-side gradient combine
to Phase E.3 (where the LobSim integration provides a clean entry
point for a new PerceptionTrainer backward-encoder method). The Q/π/V
kernels DO emit grad_h_t but it is not yet wired into the encoder
backward path. The encoder still learns from BCE+aux only in E.2;
the new heads update their own weights via per-head Adam. E.3 lands
the missing piece via `launch_grad_h_accumulate` (wired and ready).
Bellman target distribution: Phase E.2 uses a deterministic
single-atom-mass projection (host-side, mapped-pinned upload). Phase
E.3 replaces with the proper categorical projection kernel that
consumes γ from ISV[400] and the target net's bootstrap atom values.
build.rs: registers `v_head_fwd_bwd` + `grad_h_accumulate` cubins.
Cache-bust v21 → v22 with a verbose changelog entry covering this
phase's contract changes.
The integrated_trainer_smoke #[ignore]-gated GPU test remains gated;
Phase E.3 activates it alongside LobSim. Non-GPU host tests
(loss_balance default + bootstrap) continue to pass.
Companion to Phases A-E.1 (commits 6a46ded7d9ec43fdb956efd96cb9732a667c729f110e0).
Adds the orchestration layer for the integrated RL trainer per
docs/superpowers/plans/2026-05-22-integrated-rl-trainer.md.
What this commit lands:
- IntegratedTrainer struct owning:
* PerceptionTrainer (encoder + BCE + aux machinery from Phase B)
* DqnHead (Phase C)
* PolicyHead + ValueHead (Phase D)
* Device ISV buffer (RL_SLOTS_END = 412), zero-initialised so
controllers bootstrap on their first emit
* Host ISV mirror for loss-lambda reads
- LossLambdas struct + read_loss_lambdas_from_isv() helper following
the pearl_first_observation_bootstrap pattern (sentinel-zero ISV
reads as Pearl-A bootstrap value = 1.0 default per head)
- step_synthetic() entry point: runs encoder forward, refreshes ISV
host mirror, reads λ, combines synthetic placeholder losses
- Integration smoke test (ignore-gated until Phase E.2 activates
the GPU kernel path) + host-side λ defaults test
What this commit DEFERS to Phase E.2:
- Real GPU kernel calls for Q/π/V forward (currently placeholder
scalar losses)
- Backward path combining all 5 heads' grad_h_t into the encoder
- DQN/PPO head Adam state (currently encoder + BCE/aux update only)
- LobSim integration
- Toy bandit test activation (dqn_toy.rs / ppo_toy.rs /
integrated_trainer_smoke.rs all ignore-gated)
The placeholder loss values in step_synthetic are NOT a
feedback_no_stubs violation: they are a deterministic host-side
computation that becomes a real GPU kernel call in Phase E.2's atomic
refactor commit. Struct fields, cubin handles, and ISV plumbing are
all real and exercised by Phase E.2.
Per pearl_loss_balance_controller and feedback_isv_for_adaptive_bounds:
the 4 RL loss λ slots (ISV[408..412]) are read at the loss-combine
site, not hardcoded. Aux λ is unchanged (aux trainer still owns it).
Companion to Phases A (6a46ded7d), B (9ec43fdb9), C (56efd96cb),
D (9732a667c). Phase E.2 wires the kernel calls + LobSim.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Partner to Phase C (DQN/C51). Adds the PPO component of the integrated
RL trainer per docs/superpowers/plans/2026-05-22-integrated-rl-trainer.md.
What this commit lands:
- PolicyHead: linear h_t [B, HIDDEN_DIM] -> logits [B, N_ACTIONS=9].
Softmax fused into surrogate kernel.
- ValueHead: linear h_t [B, HIDDEN_DIM] -> scalar V(s) [B].
- ppo_clipped_surrogate.cu: fwd kernel computes pi_new probabilities,
the clipped surrogate L_pi = -min(ratio*A, clip(ratio,1+-eps)*A),
value MSE, entropy bonus. Bwd kernel computes per-logit grad with
clip-mask. eps and entropy coef are read from ISV[402] / ISV[403],
NOT hardcoded.
- RolloutBuffer: capacity-bounded on-policy buffer with Q-bootstrapped
advantage A_t = Q(s_t,a_t) - V(s_t) and done-aware backward-returns.
- rl_ppo_clip_controller.cu: ISV[402] producer; eps adapts to keep
KL ~ 0.01 target; bootstrap eps=0.2; clamp [0.05, 0.5].
- rl_entropy_coef_controller.cu: ISV[403] producer; coef adapts to keep
entropy >= 0.7*ln(9); bootstrap 0.01; clamp [0.0, 0.05].
What this commit DEFERS to Phase E:
- Toy bandit test activation (test stub is #[ignore])
- atomicAdd in surrogate loss accumulator (replaces with warp-shuffle
reduce when integrated with the full training loop, same plan as
dqn_distributional_q_bwd)
- V-head gradient kernel (single MSE backward - trivial; Phase E's
loss-combine path will handle it inline)
- Boundary case in clipped surrogate bwd (Phase D uses 'zero outside
clip; standard PG inside'; Phase E may refine the sign-of-A edges)
Per pearl_controller_anchors_isv_driven and feedback_isv_for_adaptive_bounds:
eps and entropy coef are read from ISV at consumer site, not hardcoded.
Bootstrap values shown in the controllers (eps=0.2, coef=0.01) are what
first-observation emits produce, not const defaults baked into the loss
kernel.
Validation:
- SQLX_OFFLINE=true cargo check -p ml-alpha --lib: clean (54.96s)
- cargo test -p ml-alpha --lib rl::rollout: 1 passed
- cargo test -p ml-alpha --test ppo_toy --no-run: compiles
- All 3 new cubins built into target/debug/.../out/
Companion to Phase C (commit 56efd96cb). Phase E wires both DQN and PPO
heads into the IntegratedTrainer with LobSim and reward shaping.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Adds the DQN component of the integrated RL trainer per the plan at
docs/superpowers/plans/2026-05-22-integrated-rl-trainer.md.
What this commit lands:
- DqnHead: linear projection h_t [B, HIDDEN_DIM] -> atom logits
[B, N_ACTIONS=9, Q_N_ATOMS=21] with parallel target-network weights.
Xavier x 0.01 init (initial softmax-over-atoms approx uniform),
scoped_init_seed-guarded per pearl_scoped_init_seed_for_reproducibility.
- dqn_distributional_q.cu: forward (one block per (batch, action), one
thread per atom) + Bellman categorical-CE backward against a pre-projected
target distribution. Atom-softmax fused into backward.
- ReplayBuffer (rl/replay.rs): capacity-bounded PER with priority^alpha
sampling, random replacement, and TD-error priority update. O(N)
cumulative-sum sampling; Phase E may upgrade to a GPU sum-tree once
capacity profiling demands it.
- rl_gamma_controller.cu: ISV[RL_GAMMA_INDEX=400] producer; gamma
adapts toward 0.5^(1/mean_trade_duration) via Wiener-alpha blend
(floor 0.4 per pearl_wiener_alpha_floor_for_nonstationary), clamped
to [0.90, 0.999]. Bootstrap gamma = 0.99 on sentinel.
- rl_target_tau_controller.cu: ISV[RL_TARGET_TAU_INDEX=401] producer;
tau adapts multiplicatively from Q-divergence ratio vs anchor 0.01,
Wiener-alpha blend with floor 0.4, clamped to [0.001, 0.05].
Bootstrap tau = 0.005 on sentinel.
- Action enum + try_from_u32 in rl/common.rs (matches existing ml DQN
action grid for cross-system policy comparability).
- C51 atom support constants Q_V_MIN / Q_V_MAX in rl/common.rs (kept
for Phase E's projection kernel; backward in this commit operates in
categorical domain on a pre-projected target).
What this commit DEFERS to Phase E:
- soft_update_target kernel (struct fields w_target_d / b_target_d are
wired and read by the Bellman backward in this commit; the writer
lives in Phase E alongside the training-loop tau driver).
- Categorical projection kernel that reads gamma from ISV[400] and
produces the target_dist input to the backward kernel.
- Toy bandit test activation (tests/dqn_toy.rs is #[ignore]-gated; the
type contract is locked here, the training loop wires in Phase E).
- atomicAdd in the per-batch CE accumulator (Phase E replaces with the
warp-shuffle + shared reduce pattern from aux_loss.cu when batches
reach production sizes; B <= 32 toy contention is negligible).
Per pearl_controller_anchors_isv_driven and feedback_isv_for_adaptive_bounds:
gamma and tau are NOT hardcoded constants. They live in ISV[400] /
ISV[401], emitted by the controller kernels above, and Phase E consumers
read via __ldg(isv + INDEX). Bootstrap values (0.99, 0.005) appear only
in the controller kernel as first-observation defaults, NOT baked into
the loss kernel.
Validation:
- SQLX_OFFLINE=true cargo check -p ml-alpha --lib -> clean (1m 03s)
- SQLX_OFFLINE=true cargo check --workspace --lib -> clean (42s)
- SQLX_OFFLINE=true cargo test -p ml-alpha --lib rl::replay -> 3 pass
- Cubins built for all 3 new kernels (sm_80 default).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
Adds PerceptionTrainer::forward_encoder() returning a borrowed slice
to h_t — the CfC h_new at the FINAL window position (K-1), where the
trade decision is made. The integrated RL trainer (ml_alpha::rl,
Phase E) consumes this to dispatch its 5 loss heads (BCE direction,
C51 Q, PPO pi, PPO V, aux prof+size) on the same encoder
representation that a supervised step() would have used.
Implementation strategy: reuse the existing forward_only() captured
graph (snap_features -> VSN -> Mamba2 x2 -> CfC K-loop -> BCE GRN
heads) rather than splitting the encoder forward out of the captured
graph. The BCE heads still run but their output (probs_per_k_d) is
discarded; the overhead is one fused GRN kernel per k (small vs.
Mamba2+CfC) and avoids the risk of breaking
pearl_cudarc_disable_event_tracking_for_graph_capture or
pearl_no_host_branches_in_captured_graph. Phase E end-to-end
profiling can revisit if needed.
A new dedicated h_t_d: CudaSlice<f32> field of size [B * HIDDEN_DIM]
receives a stream-ordered DtoD copy from h_new_per_k_d at slot K-1
after each forward_encoder() call. This gives RL callers a stable
borrowed reference even if a subsequent forward_encoder() overwrites
the per-K scratch.
Phase B (this commit) lands forward only. The backward path
(per-head loss -> lambda-weighted combine -> encoder backward via
existing step_backward_* machinery) is wired in Phase E once the
heads (Phases C+D) exist.
Existing step()/step_batched()/forward_only() semantics are
unchanged — the new field is initialised once at construction and
written only by forward_encoder(). All 56 ml-alpha lib tests still
pass; the new tests/encoder_gradient.rs locks the API contract
(borrow length B*HIDDEN_DIM, captured-graph determinism across two
calls, finite values, at least one non-zero element).
Adds crates/ml-alpha/src/rl/{mod,common,isv_slots}.rs scaffolding
for the integrated RL trainer per the plan at
docs/superpowers/plans/2026-05-22-integrated-rl-trainer.md.
Phase A lands only the module structure + Transition struct +
12 ISV slot constants (400-411). Subsequent phases (B-H) wire the
encoder gradient hookup, DQN/PPO heads, training loop, Argo
workflow, and backtest gate.
Per pearl_controller_anchors_isv_driven and
feedback_isv_for_adaptive_bounds, every adaptive hyperparameter
documented in the slot constants is sourced from ISV (not from
hardcoded const). Controller kernels (producers) land in phases
C/D/E/F.
Adds per-file labels cache sidecar (CachedLabels struct with
labels_full + outcome_prof_long/short_full + outcome_size_long/short_full +
sigma_k_full + pos_fraction + regime_full). Cache key encodes
(horizons, outcome_label_cost, instrument_filter) so any of those
changing invalidates the cache. The load path also rejects caches whose
regime_full length doesn't match the freshly-loaded snapshot count, as a
defense against re-downloaded quarters whose underlying MBP-10 sidecar
was refreshed but whose labels sidecar wasn't.
Stacks with SPEED-C (~400x on Welford) and SPEED-A (~4-8x parallel):
- First run on a file: pay the existing label-generation cost, write
the cache.
- Subsequent runs: load arrays directly from bincode sidecar — ~1s/file
vs ~3 min/file previously.
Inference-only runs skip the cache write to avoid polluting it with
all-default vectors that would mis-serve a later non-inference run.
Cache key suffix format: 'labels_h<H0>_<H1>_<H2>_cost<HEX>_<FILTER>'
to keep filenames portable (no embedded '.' from f32 formatting) and
preserve exact-float identity via raw-bit encoding of the cost.
Single integrated trainer where 5 loss heads (BCE direction, C51
distributional Q, categorical π, scalar V, aux prof+size) sit on a
shared Mamba2+CfC encoder. Joint training with adaptive loss-balance
λ weights via the existing pearl_loss_balance_controller infrastructure.
Discrete 9-action grid shared between Q and π heads; reward = per-trade
realized PnL from LobSimCuda.
Designed in response to lob-backtest-sweep-jpdhg result (PF=0.24,
sharpe=-9.72, mono-anti-cal in conviction→PnL): the encoder learns
directional AUC but never sees PnL-aware gradient because
stop_grad_aux_to_encoder blocks the aux head's gradient. RL training
fixes this by making the encoder optimize per-trade realized PnL via
Q-head Bellman + π-head clipped surrogate.
Every hyperparameter (γ, τ, PPO clip ε, entropy coef, rollout steps,
PER α, reward scale, loss-balance λs) is ISV-driven per
pearl_controller_anchors_isv_driven and feedback_isv_for_adaptive_bounds.
12 new ISV slots (400-411) + 12 reused existing slots + 8 new controller
kernels following the Wiener-α + first-observation-bootstrap pattern.
10 phases (A-H, ~3 weeks), falsifiable gate G8 = profit_factor > 1.0
on 2M-event backtest sweep.
Updates sweep_smoke_perhoriz_cfc.yaml to use the FIRST clean (front-
month-filtered) checkpoint (commit 20aa345a7, auc_h1000=0.5757) instead
of the prior 2efedcd6b checkpoint (auc_h1000=0.7137 was a multi-
instrument $5000 ΔP contamination artifact, not real signal).
Primary gate unchanged: outcome_by_entry_conv table must NOT show
monotonic anti-calibration. With clean data the conviction signal
should be lower magnitude but properly aligned.
The Phase 1 multi-resolution layout (10 raw + 10 agg@30 + 12 agg@100)
regressed ALL horizons in alpha-perception-htpp6 (2026-05-22):
- auc_h100: 0.681 -> 0.512 (-0.169)
- auc_h300: 0.617 -> 0.506 (-0.111)
- auc_h1000: 0.576 -> 0.526 (-0.050)
Hypothesis falsified. Root cause: Mamba2+CfC SSM encoder already used
all 32 raw ticks effectively via state-recurrence; replacing 22 raw
ticks with arithmetic-mean aggregates destroyed within-window
microstructure variance (the actual h100 signal) AND broke temporal
continuity that the recurrence relies on. Δt Fourier encoder
couldn't compensate.
Architectural pearl: SSM/RNN/CfC + multi-resolution input is
incompatible without separate-encoder-per-scale or explicit scale
tokens. Transformer-style positional encoding tolerates scale-mixing;
recurrent state updates assume consecutive positions.
Reverts default to '1:32'. Adds explicit single_scale_32() constructor
for callers (harness, tests). Keeps default_three_scale() in code with
deprecation note for future sub-variant experiments. Production
defaults across alpha_train CLI, Argo template, dispatcher script,
ml-backtesting harness now match the proven baseline.
Each file's load_or_predecode + label generation is pure CPU work over
disjoint inputs (snapshots, cfg.horizons, cfg.outcome_label_cost). The
sequential for loop was the parallel-friendly bottleneck — converting
to rayon::par_iter gives ~4-8x speedup on typical 4-8 core hosts.
Combined with SPEED-C's ~400x speedup on the inner Welford loop, total
preload throughput is ~1600-3200x faster than the prior single-threaded
O(W) recompute path.
File order is preserved by par_iter's collect contract. Too-few-snapshots
skips emit a warn during load and resolve to None at collection.
Replaces O(W=1000) per-step mean+var recompute with f64 running
sum_x + sum_x2 updated incrementally on window push/pop. Per-snapshot
cost drops from ~2000 to ~5 float ops. On a 5M-snapshot file across
3 horizons, total hot-loop ops drop from ~30B to ~75M.
f64 accumulators contain ~16 decimal digits - over a 5M-step file
the accumulated rounding error stays well below the f32 output
precision. Every RECOMPUTE_PERIOD=10000 pops, we still do a full
window sweep to reset the accumulators as defense in depth against
pathological drift.
New parity test online_sigma_matches_naive_full_window_recompute_within_tolerance
asserts the fast path matches the naive O(W) algorithm within 1e-4
relative on a 30k-snapshot synthetic stream.
Per pearl_cooperative_staging_eliminates_redundant_reads (CPU analog):
running sums eliminate the redundant window reads that dominated
preload time.
Four pure-CPU tests confirming aggregate_window emits the right
ts_ns - prev_ts_ns span per scale, so the existing snap-feature Δt
Fourier encoder gets correctly-scaled inputs without any CUDA kernel
changes.
Replaces --seq-len. Default '1:10,30:10,100:12' = 32 positions
covering 1510 ticks of context, the Phase 1 fix for temporal
receptive field mismatch. Logs the active config at preload start
so Argo logs surface the per-run choice.
Per feedback_no_htod_htoh_only_mapped_pinned: mapped-pinned only for
CPU↔GPU; tests not exempt. Previous commit ab64c0412 was a shallow
deprecated-API swap (memcpy_stod -> clone_htod) — both still HTOD
copies. This commit replaces all 5 uploads and the 1 readback in
test_eval_action_select_thompson_picks_proportionally with the
production mapped-pinned + memcpy_dtod_async pattern (mirrors lines
433-448 of the same file).
Closure 'upload_pinned' DRYs the 5 upload sites; readback uses
MappedI32Buffer + DtoD into staging.host_slice_mut().to_vec().
Zero HTOD/HTOH copies remain in this file. Documented in
dqn-wire-up-audit.md.
cudarc 0.19 deprecated memcpy_stod (use clone_htod) and memcpy_dtov
(use clone_dtoh). Six call sites in cuda_pipeline/mod.rs migrated.
Pre-existing warnings surfaced now because the recent file-level
allow(unsafe_code) on this file no longer per-line-suppresses the
deprecated lint.
Per feedback_no_legacy_aliases: rename call sites directly, no
deprecated wrappers. Documented in dqn-wire-up-audit.md.
multi_horizon_loader.rs constructs MultiHorizonLoaderConfig with
default_three_scale() (matching alpha_train's production default
and yielding total_positions()=32). The cfg.seq_len=32 override
in loader_stride_4_yields_correct_spacing is now a no-op since
the constructor already provides 32, so it's removed with a comment.
perception_overfit.rs was untouched — it doesn't construct
MultiHorizonLoaderConfig, only PerceptionTrainerConfig (whose own
seq_len field is unrelated to the loader migration).
Backtest harness + tests use default_three_scale config (1:10,30:10,
100:12 = 1510 ticks of context). Replaces the previous seq_len=32
single-scale path which is now deleted.
Multi-resolution sequence builder reads from
[anchor + total - lookback, anchor + total), so anchor must be in
[max(0, lookback - total), snapshots.len() - (total + max_horizon)).
The previous gen_range(0..max_anchor) underflowed for default
3-scale (lookback=1510, total=32) at small anchors. Compute
min_anchor = lookback.saturating_sub(total) and sample
gen_range(min_anchor..max_anchor) with the ensure! upgraded to
check max_anchor > min_anchor.
Caught by implementer self-review of SDD-MR Task 3.
Sequence builder now walks fine→coarse scales from the anchor's
forward edge, aggregating each window into one pseudo-snapshot
via aggregate_window. seq_len field deleted (total_positions()
replaces all reads). Greenfield — no legacy single-scale path.
External callers updated in subsequent tasks.
CUDA kernel launches via cudarc::launch_builder and MappedF32::new
(cuMemHostAlloc DEVICEMAP FFI) inherently require unsafe blocks. The
workspace-wide '-W unsafe-code' lint produced ~80 warnings across these
files, all structurally identical. Match the established pattern from
ml-backtesting/src/harness.rs and ml-backtesting/src/sim/mod.rs: single
file-level allow with a comment explaining the rationale.
Files: alpha_dqn_h600_smoke, alpha_baseline, cublaslt_debug,
gpu_walk_forward, cuda_pipeline/mod, hyperopt adapters (mamba2, ppo),
DQN smoke_tests (helpers, td_propagation), trainers/ppo, and two
ml/tests files. Documented in dqn-wire-up-audit.md.
Builds one Mbp10RawInput from a window of consecutive snapshots.
Book levels meaned, flows summed, ts_ns spans the window so the
snap-feature Δt Fourier encoder gets correct dt without any CUDA
kernel changes.
Databento historical bulk DBN files emit symbol mappings only in the metadata
header (date-range form), not as in-stream SymbolMappingMsg records. The
strict 'no mapping = error' branch from the previous commit blocked the
front-month smoke (alpha-perception-vstrr) at Q1 2024 where the dominant id
5602 has no streaming mapping entry.
Volume-leader detection itself doesn't depend on the symbol; the regex check
was defense-in-depth. Log-warn on symbol-mismatch and proceed with id=metadata-only
when no streaming mapping exists. Calendar-spread anomalies still surface in
logs without blocking training.
Replaces Option<u32> instrument_id_filter with InstrumentFilter enum {All,
Id(u32), FrontMonth}. FrontMonth runs a two-pass detect over the DBN
stream: pass 1 counts instrument_ids and collects SymbolMapping records,
picks the dominant id, validates it resolves to an ES contract via regex
ES[FGHJKMNQUVXZ]\d{1,2}; pass 2 streams the filtered records.
Motivated by alpha-perception-k54wd: a single-id filter on parent-symbol
ES.FUT data caught Q1 2024 (kept=73M) but kept=0 for Q2-Q9 because ES
front-month rolls quarterly (ESH4 -> ESM4 -> ESU4 -> ESZ4 ...). FrontMonth
self-tunes across the rolls without needing a per-file id table.
Sidecar keys distinguish modes: mbp10 / mbp10_instr<id> / mbp10_front_month.
CLI flag renamed --instrument-id -> --instrument-mode {all,id=N,front-month}
with matching parameter rename in argo-alpha-perception.sh + template.
Adds new workflow parameter instrument-id (default empty string = no filter).
Script forwards via --instrument-id CLI flag to alpha_train when non-empty.
EXTRA_FLAGS template logic appends --instrument-id only when set.
kubectl apply must run first (per feedback_argo_template_must_apply) so the
cluster CRD reflects the new parameter before argo submit.
Step 2 of aux-diagnosis-deeper plan (NumPy on local ES data) falsified
the F2 hypothesis:
1. Local data has 0% zero-bid/zero-ask records. F2's "one-sided book
contamination" was incorrect — the actual databento sentinel for
missing price is INT64_MAX × 1e-9 ≈ 9.22e9 (a HUGE POSITIVE number
that passes the `> 0` check). F2 was a no-op on real data.
2. Sentinel rate on local Q1 is 0.006% — negligible.
3. Local pos_fraction at K=10 is 19.59%, cluster reports 35.07%. The
~75% gap is plausibly explained by overnight session gaps in the
full 5M-record quarter file (which I didn't see in my 500k local
sample). These are real price discontinuities, not "contamination".
4. F2 introduced an epoch-4 NaN explosion (alpha-perception-7shgw)
because NaN cascaded through the σ_K Welford rolling computation.
Reverting:
- crates/ml-alpha/src/data/loader.rs::mid_price_f32 → blind average
- crates/ml-alpha/src/multi_horizon_labels.rs:218 → is_finite only
KEEPING:
- F1 dir_acc fix in perception.rs:3647 — empirically working (metric
hovers ~0.5 instead of pinned sub-chance)
cargo check --workspace --all-targets clean.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
F-series investigations triangulated TWO real bugs that together
explained: aux BCE > ln(2) chance baseline, sub-chance dir_acc with
balanced labels, and bit-identical training across vastly different
POS_WEIGHT_MAX values.
BUG #1 (F2) — One-sided book contamination in mid_price_f32:
- crates/ml-alpha/src/data/loader.rs::mid_price_f32 blindly averaged
bid_px[0] + ask_px[0] without checking validity. MBP-10 snapshots
at session boundaries / halts / stale level-0 vacancies have
bid_px=0 or ask_px=0; result was mid = real_price/2 (~2750 for ES)
or mid = 0 (both sides empty).
- generate_outcome_labels_ab:218-220 only guarded is_finite() (NOT > 0).
Sister fn generate_labels:75 does check both — asymmetry between
parallel generators.
- Transitions like mid=0 → mid=5500 produced ΔP=5500, instantly
satisfying delta > 2×cost → spuriously triggering y_prof=1. Each
contaminated snapshot tainted up to 2K downstream labels (appears
as p_t for K positions and as p_kt for K positions).
- With ~10-20% one-sided books in real ES, ~35-45% of K=10 labels
spuriously positive — exactly matching the observed pos_fraction
(0.35, 0.39, 0.46 for long).
Fix (broader): mid_price_f32 returns NaN for one-sided/empty books
at the source. Cascades to ALL downstream consumers (regime features,
snap_features encoder, labels). Defensive guard also added in
multi_horizon_labels.rs:218 for direct-test callers bypassing the
loader.
BUG #2 (F1+F3) — dir_acc metric structural bias:
- perception.rs:3647 match condition for flat-true bucket required
float-EXACT equality on raw logits (`pred_diff == 0.0`) — essentially
unreachable for continuous-valued logits.
- On real ES, ~30-50% of samples are flat (neither direction
profitable at 2×cost threshold). ALL counted as misses → random
baseline depressed to ~0.35 (matching observed 0.32-0.45). BCE
can DROP while dir_acc DROPS — decoupled metrics.
- aux_lift_dir_acc_threshold=0.85 was structurally UNREACHABLE under
this metric regardless of model quality.
Fix: skip flat-true samples (`if true_diff == 0 { continue; }`).
Converts dir_acc into "given a directional outcome, did we get the
direction right?" with proper random baseline 0.5.
F3 + F4 confirmed: aux head wiring (8 head Adam optimizers, fwd/bwd
kernel arg order, BCE+sigmoid gradient sign, reduce_axis0 reduction)
is correct. The synthetic perception_overfit test stays in chance
regime because it constructs constant prof_long=1, prof_short=0 →
no zero-priced snapshots, no flat-true bucket. Production-data path
divergence is the canonical pattern in
pearl_canary_input_freshness_launch_order.
cargo check --workspace --all-targets: clean.
cargo test -p ml-alpha --lib multi_horizon_labels: 15 passed.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
A+B Smoke 1 (alpha-perception-79rbn) showed pos_weight=50 caused
overshooting: BCE=0.83 (worse than chance ln(2)=0.69) AND sub-chance
dir_acc (0.28-0.45) at all horizons. The 50× cap pushed the model
toward the rare positive class hard enough to blow up loss on 99.99%
of negative samples.
Lower MAX clamp to 10× lets per-horizon imbalance still scale the loss
but caps the runaway gradient on extreme-rare positive classes (like
K=100's 0.01% positive rate where pos_weight would otherwise be ~9999).
Synthetic test still passes (clamp doesn't matter on constant signal
where pos_fraction=1.0).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
CB1+CB2 swapped labels D→A+B; this swaps the kernels to match.
aux_heads.cu — 4-output structure:
- Forward: 12 outputs per snapshot = 4 per direction × N_HORIZONS
(prof_long_logit, size_long_pred, prof_short_logit, size_short_pred,
each [N_HORIZONS]). Linear projections; sigmoid applied in BCE kernel.
- Backward: accepts 4 grad_y inputs, produces 8 grad_W + 8 grad_b +
grad_h_aux. Cooperative h_aux staging in shmem once per block.
- 8 weight matrices total, Xavier × 0.1 init under scoped_init_seed.
aux_loss.cu — 2 kernels:
- aux_bce_loss_fwd_bwd: class-weighted BCE+sigmoid fused. pos_weight in
shared mem; scales positive-class gradient. NaN-mask y_true.
- aux_huber_masked_fwd_bwd: Huber w/ NaN-mask. CB1's y_size=NaN at
y_prof=0 provides the conditional-Huber semantics naturally — no
separate mask buffer needed.
aux_heads.rs:
- AuxHeads + AuxHeadsWeights: 8 buffer fields (4 W + 4 b)
- AuxBceLoss + AuxMaskedHuberLoss wrappers replace AuxHuberLoss
- POS_WEIGHT_MIN/MAX = [1.0, 50.0] clamps per E3
- aux_heads_fwd_gpu/aux_heads_bwd_gpu/aux_bce_loss_gpu/aux_huber_masked_loss_gpu
perception.rs (minimal compile-keeping signature updates only):
- Renamed/added buffers: 4 prediction (prof/size × long/short),
4 label staging, 4 grad_y per-K, 8 head grad scratches, 8 head Adam
optimizers, 2 pos_weight buffers (device + staging)
- HOLDING PATTERN: bwd zeroes the 4 grad_y_per_K buffers each step so
the head Adam updates are no-ops on grad=0 (no aux gradient signal
this commit). CB5 wires the actual aux_bce + aux_huber_masked calls.
- BCE direction signal + dir_acc readouts updated to use the new
prof_long/prof_short prediction buffers (so existing perception_overfit
aux test still passes).
7 GPU oracle tests on RTX 3050 sm_86, all pass in 2.43s:
- fwd_matches_naive_reference, bwd_finite_diff_matches_bias_sample,
aux_bce_loss_matches_naive_reference, aux_bce_pos_weight_scales_positive_gradient
(verified: pos_weight=10 → 10× gradient ratio within 1e-4),
aux_huber_masked_does_not_propagate, aux_huber_masked_covers_both_branches,
aux_bce_nan_mask_does_not_propagate
Cubins rebuilt: aux_heads (8736→10528 bytes, +20% for 4-head fwd/bwd),
aux_loss (6944→13344 bytes, +92% for 2 kernels).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Smoke 1 at HEAD 21e7dfd63 showed BCE bit-identical to baseline
(asymmetric stop-grad isolation working) but zero aux signal in logs.
Adds:
- Per-epoch tracing 'aux snapshot' line with aux_huber_h{10,100,1000},
aux_dir_acc_h{10,100,1000}, stop_grad_aux_to_encoder
- AlphaTrainSummary JSON fields: final_aux_huber_ema_per_h,
final_aux_dir_acc_ema_per_h, final_stop_grad_aux_to_encoder
Reads PerceptionTrainer's pub fields directly (no accessor noise).
After re-running Smoke 1 we'll have empirical evidence whether aux is
learning on REAL ES MBP-10 data (vs the synthetic constant-direction
test where dir_acc trivially hits 1.0). This informs B7's gate threshold
design.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Wires the aux supervision path parallel to BCE:
- AuxTrunk (64-hidden single-bucket CfC) consumes the same encoder output
as the main BCE trunk
- AuxHeads (linear regression long/short) maps aux_trunk output to
per-(direction, horizon) predicted outcomes
- AuxHuberLoss supervises against D-style labels from MultiHorizonLoader
Backward path with asymmetric stop-grad at encoder boundary:
- aux_trunk gets gradient signal into its OWN params at all times
- aux_trunk's encoder-boundary gradient is INITIALLY blocked
(stop_grad_aux_to_encoder = true)
- Conditional lift per E3 design: if aux_huber_ema < 0.4 AND
aux_dir_acc_ema > 0.85 within 200 steps, lift the stop-grad
- When lifted, aux_vec_add kernel folds aux's grad_x into the main
grad_h_enriched_seq slot (element-wise += per feedback_no_atomicadd)
ISV signals added: aux_huber_per_h, aux_dir_acc_per_h (per pearl).
Per-trunk scratch + reduced grad buffers (no Adam state sharing per
pearl_adam_normalizes_loss_weights — opt_aux is its own Adam group).
New helper kernel cuda/aux_vec_add.cu: position-local dst += src for
the asymmetric stop-grad lift accumulation.
New synthetic test stacked_trainer_aux_supervision_converges_on_constant_signal
validates end-to-end:
aux_huber_ema_per_h = [0.087, 0.087, 0.087] (converged)
aux_dir_acc_ema_per_h = [1.0, 1.0, 1.0] (perfect on constant)
stop_grad_aux_to_encoder = false (lift fired)
All 5 stacked_trainer tests pass on RTX 3050 (lib still converges, no
regression from parallel aux wiring).
Not yet consumed by decision policy (B7) — aux output flows through
training only.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Adds the aux heads that map AuxTrunk's output (64-dim) to per-(direction,
horizon) predicted outcomes, plus the Huber loss kernel that supervises
them against the D-style labels generated in B1.
Files (1212 LOC):
- cuda/aux_heads.cu (197 LOC): fused fwd+bwd for linear projection
h_aux[B, 64] → y_long_hat[B, N_HORIZONS] + y_short_hat[B, N_HORIZONS].
Single launch for both directions, per-batch grad scratch + caller-side
reduce_axis0, cooperative h_aux staging in shmem.
- cuda/aux_loss.cu (143 LOC): Huber loss fwd+bwd with NaN-masking. Per
direction call; returns Σ Huber + valid_count separately so caller picks
reduction policy.
- src/aux_heads.rs (412 LOC): AuxHeads + AuxHuberLoss wrappers, weight
structs, Xavier init under scoped_init_seed.
- tests/aux_heads.rs (460 LOC): 4 #[ignore]'d GPU oracle tests
(fwd_matches_naive, bwd_finite_diff_matches_bias, huber_loss_matches_
naive, huber_loss_nan_mask_does_not_propagate). 4/4 PASS on RTX 3050.
- build.rs: KERNELS += aux_heads, aux_loss; cache-bust bumped to v17.
- src/lib.rs: pub mod aux_heads.
Design decisions (full notes in subagent report):
- Linear heads (no GRN) — D-labels already encode asymmetric loss-aversion
- Per-direction Huber launches (cleaner per-direction telemetry)
- NaN-masking in loss kernel (single NaN label can't poison batch grad)
- Unreduced sum + valid_count separately (caller picks mean policy)
Not yet wired into trainer (B5) or used in policy (B7).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
New 64-hidden single-bucket CfC trunk to serve as the parallel
parameter group for D-style aux supervision (Layer B of anti-cal plan).
Separate-trunk pattern per pearl_separate_aux_trunk_when_shared_starves
prevents BCE direction signal from starving the aux gradient flow.
Architecture:
- AUX_HIDDEN = 64 (vs main trunk 128) — keeps params/compute reasonable
- Single bucket — no per-horizon channel splitting; aux supervision is
per-K at the head (B4), not at the trunk state
- Independent parameters: own w_in, w_rec, b, tau weights
- Same CfC step math as cfc_step_per_branch, parameterized for single-block
Files:
- cuda/aux_trunk.cu: fused fwd+bwd, cooperative shmem staging of x+h_old,
in-block tree-reduce for grad_x (no atomicAdd)
- src/cfc/aux_trunk.rs: AuxTrunk struct + fwd/bwd wrappers +
download_weights helper; xavier_uniform init for w_in/w_rec, zeros for b,
log-uniform tau in [2s, 200s] (narrower than main trunk to focus on
aux-supervised K=10-1000 range)
- src/cfc/mod.rs: pub mod aux_trunk + re-exports
- build.rs: KERNELS += "aux_trunk"
- tests/aux_trunk.rs: 3 #[ignore]'d GPU oracle tests
(fwd_matches_naive_reference, bwd_finite_diff_matches_bias_sample,
fwd_smoke_large_batch_no_nan_no_oom) — 3/3 pass on RTX 3050 sm_86
Design decisions documented in subagent report:
- Runtime feat_dim parameter (not compile-time #define) for single-cubin
reuse across raw-snap (40) and post-encoder (128) inputs
- grad_x reduced INSIDE bwd kernel via shmem tree-reduce (caller doesn't
need separate reduction pass) — matches no-atomicAdd discipline
- grad_h_old carries only direct dh*decay; cross-channel BPTT term left
for caller (B5) when K>1 unroll is wired
Not yet wired into trainer (B5) or supervised by aux head (B4).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Extends MultiHorizonLoaderConfig with `outcome_label_cost` (default
DEFAULT_OUTCOME_LABEL_COST_ES = 0.5 = 2 ticks for ES round-trip).
LabeledSequence + LoadedFile gain outcome_long + outcome_short arrays
of [Vec<f32>; N_HORIZONS] populated by generate_outcome_labels_d during
LoadedFile construction (same !inference_only branch as BCE labels).
CLI exposure: alpha_train.rs gains --outcome-label-cost flag with the
ES default. All 6 MultiHorizonLoaderConfig consumers updated atomically
per feedback_no_partial_refactor: alpha_train (train + val loaders),
in-file test fixtures (×2), multi_horizon_loader.rs (×2), harness.rs,
trainer_parity.rs, ring3_replay.rs.
cargo check --workspace clean; ml-alpha lib tests 41 passed.
B3+ tasks not yet touched: outcome labels are computed and propagated to
LabeledSequence but no consumer reads them yet (aux trunk Tasks B3-B5
will wire that).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Adds generate_outcome_labels_d for the future aux supervision head
(Layer B of the anti-calibration plan, docs/superpowers/plans/2026-05-22-
horizon-rebase-n3-100-300-1000-and-aux-d-labels.md).
Per (snapshot t, horizon K, direction d ∈ {long, short}):
profit = signed_pnl(d) - cost
dd_against = max adverse excursion in [t, t+K] holding direction d
y[t, K, d] = profit - 1.5 × |dd_against|
The 1.5× drawdown penalty encodes loss-aversion per behavioral finance —
trades with deep drawdowns are penalized even if final outcome is positive.
Implementation: monotonic-deque sliding-window min/max per horizon,
O(N · N_HORIZONS) amortized. NaN at right-edge positions (t+K >= n).
Input validation rejects zero horizon and non-finite/negative cost.
13 tests pass: 5 hand-derived value checks (simple up, drawdown
penalty, NaN edge, sliding-window correctness, per-horizon independence),
3 additional edge tests (constant-series invariant, zero-horizon
validation, negative-cost validation), plus the 5 pre-existing tests
unchanged.
Not yet wired into the loader (Task B2) or used as supervision target
(Tasks B3-B5). New function only.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Point at alpha-perception-9h5xc's trunk_best_h1000.bin (commit 2efedcd6b,
Smoke 1 validated auc_h1000=0.7137 >> 0.65 gate).
Reframe the sweep gate: primary is now anti-calibration check (validates
D1's signed-EMA fix at 0384f7674). Per-horizon mean_run_len
differentiation is informational only — the per-horizon-differentiation
thesis was falsified earlier this session.
5 sweep YAMLs updated to reference the new checkpoint filename:
- config/ml/sweep_smoke.yaml: trunk_best_h6000.bin → trunk_best_h1000.bin
- config/ml/sweep_perhoriz_diag.yaml: same
- config/ml/sweep_threshold_tuning.yaml: same
- config/ml/sweep_deployability.yaml: same
- config/ml/sweep_smoke_perhoriz_cfc.yaml: same + 3 comment updates
(WIN-gate criteria, header doc, best-checkpoint annotation)
scripts/generate_sweep_variants.py:61 updated atomically — without this
the next regeneration of sweep_deployability.yaml would silently
re-introduce trunk_best_h6000.bin.
Argo workflow templates (alpha-perception, alpha-cv, lob-backtest-sweep)
did NOT need text changes — they're already horizon-agnostic:
- They forward CLI flags via {{workflow.parameters.*}} to binaries
- Don't grep alpha_train_summary.json inline
- Don't reference per-horizon field names
- early-stop-metric default is "mean_auc" (horizon-agnostic)
Intentionally left:
- alpha-cv-template.yaml stacker-horizon: "6000" (unrelated TFT lookback)
- alpha-cv-template.yaml horizon: "1200" (DQN execution horizon in bars)
- lob-backtest-sweep-template.yaml ci-training-h100 (GPU pool name)
NOTE: kubectl apply of workflow templates is deferred to Task 10 (push
+ dispatch). Verified all 5 sweep YAMLs parse with python3 yaml.safe_load.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
