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.
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.
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).
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.
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.
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.
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>
Three CUDA kernels + atomically-coupled Rust consumer:
- horizon_lambda.cu: N_HORIZONS_LAMBDA 5→3
- bce_loss_multi_horizon.cu: BCS_N_HORIZONS 5→3
- smoothness_lambda_controller.cu: SLC_N_HORIZONS 5→3 AND TARGET_K_RATIO
rebased from old-horizon {30,100,300,1000,6000} sqrt formula to new
{10,100,1000} → {1.0, 0.3162, 0.1}. Payload size 10 → 6 (2×N_HORIZONS).
- gpu_log.rs: payload_json decoders for RT_INPUT, RT_STATE, RT_OUTPUT
records updated to 3-horizon field names (h30..h6000 → h10/h100/h1000)
per feedback_no_partial_refactor.
Bucket-coupled kernels (bucket_transition, cfc_step_per_branch,
heads_block_diagonal_fwd, multi_horizon_heads) STILL HAVE N_HORIZONS=5
and bucket geometry constants. Next commit migrates those — they share
memory layout with Rust-side bucket_routing.rs which is already at
N_HORIZONS=3 / MAX_BUCKET_DIM=96, so kernel-side mismatch would be
silent data corruption at runtime.
decision_policy.cu N_HORIZONS comes from lob_state.cuh — Task 6 scope.
cargo build -p ml-alpha and -p ml-backtesting: cubins rebuild PASS.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Re-derive 3-element fixtures for [10, 100, 1000] preserving geometric
decay invariant. Rename excess_at_h6000_lifts_lambda_proportionally to
excess_at_h1000_lifts_lambda_proportionally.
Critical correction during execution: the kernel uses SQRT-anchored
TARGET_K_RATIO (since commit b5bed9f80 "sqrt K-ratio") not linear ratio.
The lifted-fixture computation mirrors the kernel's sqrt constant
(TARGET_K_RATIO_H2 = sqrt(10/1000) ≈ 0.3162) so the test fires the
lambda = 10 × base invariant under the actual kernel math.
Side-discovery (flagged for Task 5 scope expansion):
- cuda/smoothness_lambda_controller.cu:30 still has SLC_N_HORIZONS = 5
- TARGET_K_RATIO at lines 45-51 uses old-horizon formula {30/30, 30/100,
30/300, 30/1000, 30/6000}. With N_HORIZONS=3 the kernel reads only
slots [0..3] = {1.0, 0.5477, 0.3162} — those correspond to old
30/30, 30/100, 30/300 ratios. h1000's smoothness target is currently
anchored to OLD h300 ratio (functional bug requiring kernel update).
cargo test --test smoothness_lambda_controller_invariants: 4 passed.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
First commit of the horizon-rebase refactor (full plan at
docs/superpowers/plans/2026-05-22-horizon-rebase-n3.md). Motivated by
empirical evidence from this session's investigations:
- C1: ES MBP-10 input ACF dies by L=300; K=1000/6000 are below noise floor
- D3: binarized labels ρ(30,6000)=0.06; K=6000 carries no signal
- Heads_w1 mask experiment: h6000 AUC collapsed 0.73→0.54, confirming
long-horizon prediction needs cross-channel integration that 5-horizon
bucket structure couldn't provide
New horizon set [10, 100, 1000] matches C1's empirical autocorrelation
elbows (~50ms, ~3.5s, ~45s wall-clock at 74µs median Δt).
Scope of this commit:
- crates/ml-alpha/src/heads.rs: N_HORIZONS 5→3, HORIZONS=[10,100,1000]
- crates/ml-alpha/src/cfc/bucket_routing.rs: MAX_BUCKET_DIM 28→96,
BUCKET_DIM_K=[43,43,42], BUCKET_CHANNEL_OFFSET=[0,43,86,128]
- crates/ml-alpha/src/data/loader.rs: 5 hardcoded [T;5] types migrated
- crates/ml-backtesting/src/sim/mod.rs: broadcast_alpha signature,
local N_HORIZONS re-exports ml_alpha::heads::N_HORIZONS
- crates/ml-backtesting/src/harness.rs: local N_HORIZONS re-export,
per-horizon log labels, MultiHorizonLoaderConfig.horizons
Tests + examples + CUDA kernels still need migration (subsequent commits
per plan tasks 2-8).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Empirical ES MBP-10 median inter-event Δt is 74µs (not 250ms as
documented in loader.rs:25-26). The constants ALPHA_MED=0.02 and
ALPHA_SLOW=0.0005 were named/commented as "9s @ 250ms" and "6min @
250ms" but at the actual 74µs Δt their half-lives are ~2.6ms and ~100ms
— a 3000× scale mismatch.
Rename for truth-in-naming (values unchanged to preserve all trained
checkpoints):
- ALPHA_MED → ALPHA_FAST_2P6MS (half-life ~2.6ms @ 74µs Δt)
- ALPHA_SLOW → ALPHA_MED_100MS (half-life ~100ms @ 74µs Δt)
The regime[0..5] feature naming in snap_features.rs reflects the
actual microstructure-burst and 1-2-second-clustering time scales, NOT
macro 9s/6min regimes. True macro-regime EMAs (if needed) are a
separate addition.
ml-alpha lib: 33 passed, unchanged from baseline.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
URGENT correctness fix surfaced by checkpoint deep-dive after Smoke 2 failed
the WIN gate (mean_run_len ratio = 1.0× vs target ≥10×; all 5 horizons
uniformly ~2.4 events).
THREE INDEXING BUGS identified:
1. tau_reorder produces bucket-grouped tau_all_d, but Controller B's
tau_clamp_kernel reads bucket_id_per_channel[c] where c is the
POSITION in the reordered buffer (not the original channel index)
→ wrong bucket-IQR lookup → τ never constrained → buckets collapse
(deep-dive showed all 5 buckets had nearly identical τ ranges
[0.07, 74] except bucket 4 reaching 878).
2. heads_w_skip grad mask ran but Adam (m, v) momentum from BEFORE the
transition re-introduced gradient signal across the transition →
off-bucket positions stayed nonzero throughout training
(deep-dive: 512/512 = 100% of off-bucket positions nonzero in
trunk_best_h6000.bin).
3. per-branch CfC kernel read w_in[c * HIDDEN_DIM + k] with c =
bucket-grouped position, but W_in rows are indexed by ORIGINAL
channel → kernel read wrong rows for each output → outputs were
essentially random per-channel.
ALPHA FIX: skip the reorder entirely, use bucket-filter throughout:
- Removed tau_reorder_kernel; cfc.tau_d stays in original-channel layout.
- Added channels_in_bucket_kernel that populates a
[N_HORIZONS × MAX_BUCKET_DIM] lookup (original channel index per
(bucket, within-bucket-position)).
- Per-branch CfC fwd+bwd now reads channels_in_bucket[branch][tid]
→ original_c, then uses original_c for w_in/w_rec indexing. All
weights stay in original layout consistently.
- Controller B's tau_clamp_kernel now correctly operates on original-
channel cfc.tau_d with bucket_id_per_channel[c] lookup (no position-
vs-channel confusion).
- Added zero_off_bucket_kernel + three-layer defense for heads_w_skip
block-diagonal invariant:
(a) At transition: zero off-bucket params + zero Adam (m, v)
moments via opt_heads_w_skip.m_mut() / v_mut() accessors.
(b) Per-step: heads_w_skip_grad_mask_apply_kernel zeros off-bucket
gradients before Adam step (unchanged from prior follow-up).
(c) Per-step: zero_off_bucket_kernel zeros off-bucket params after
Adam step, catching any drift from Adam's ε denominator or
weight decay.
New AdamW::m_mut()/v_mut() accessors enable the projection at transition.
GPU oracle tests: 19 total (16 in bucket_transition + 3 in cfc_step_per_branch).
New tests verify:
- channels_in_bucket_kernel correctness under non-contiguous bucket assignment
- zero_off_bucket maintains invariant after many mock Adam steps
- fwd kernel writes only to bucket-assigned channels under arbitrary mapping
Per `feedback_no_partial_refactor`: all 3 indexing bugs + Adam momentum
defense land in one commit.
ml-alpha lib: 33 passed.
GPU oracle tests on RTX 3050 sm_86: 19 passed.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
URGENT correctness fix surfaced by Task 13 implementer.
Bug: training-time Phase 2 dispatch sites (cfc_step_per_branch_fwd at
~3139, h_mag_per_bucket for Controller D at ~3211-3212) read
self.trunk.bucket_*_d which are zero-initialized at trunk construction
and only populated via load_checkpoint. The Phase 1→2 transition
populates a SEPARATE BucketRoutingMetadata owned by the trainer
(self.bucket_routing_metadata) but never syncs it into the trunk fields.
With zero-init bucket_dim_k, the per-branch kernel's uniform predicate
`tid >= bucket_dim_k[branch]` early-returns ALL threads -> Phase 2
silently no-ops during training, Smoke 1 fails before producing signal.
Fix: at transition, DtoD-copy metadata buffers into trunk fields BEFORE
the `Some(metadata)` move (so `metadata.*` borrows remain valid):
- metadata.bucket_channel_offset_d -> trunk.bucket_channel_offset_d
- metadata.bucket_dim_k_d -> trunk.bucket_dim_k_d
- metadata.bucket_id_per_channel_d -> trunk.bucket_id_per_channel_d
- metadata.heads_w_skip_offset_d -> trunk.heads_w_skip_offset_d
Trunk fields remain the single source of truth for both training-time
dispatch AND checkpoint serialization. Total sync cost: 4 DtoD memcpys,
~256 bytes total at the one-shot transition (off the hot path).
Verification:
- `cargo check -p ml-alpha --all-targets`: clean
- `cargo test -p ml-alpha --lib`: 33 passed, 0 failed
- 3 read sites (perception.rs:3139, 3211-3212, 4725-4726) verified
unchanged; all 4 new memcpy_dtod_async calls bracketed in the
transition block (perception.rs:2224, 2234, 2244, 2254).
- Block-diagonal heads grad-mask init (step 7) continues to read
`metadata.bucket_id_per_channel_d` via `as_ref()` after the move;
ordering preserved per pearl_canary_input_freshness_launch_order.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per spec §2.4 and Task 13 of the plan.
Replace the single-CfC `cfc_step_batched` dispatch in `forward_step_into`
with `cfc_step_per_branch_fwd_gpu`. Production checkpoints have Phase 2
routing frozen at the training-time Phase 1→2 transition, so inference
unconditionally takes the per-branch path. The fused kernel reads
`bucket_channel_offset_d` / `bucket_dim_k_d` from the trunk; in
deployment these are populated via `CfcTrunk::load_checkpoint`. For
freshly-constructed trainers without a checkpoint load, the trunk fields
are zero — the kernel's uniform predicate (`tid >= bucket_dim_k`) then
early-returns every thread and h_new is left untouched. That matches the
"Phase 2 only at inference" contract.
Heads dispatch unchanged: the existing GRN kernel reads from
`trunk.heads_w_skip_d` which has off-bucket entries zeroed at the
transition (sparsification per the prior block-diagonal-grad-mask
follow-up commit). Mathematically the full-buffer read is equivalent to
a compact-only per-bucket read; no inference-time kernel change required.
forward_step_golden's convergence test is now architecturally
incompatible with the new dispatch — `forward_only` runs Phase 1
single-CfC math while `forward_step_into` requires populated Phase 2
routing. Annotated `#[ignore]` with a clear divergence note; the
deterministic + reset semantics tests remain valid invariants per
`pearl_training_smoothness_does_not_transfer_to_inference`.
cargo check workspace clean (excluding pre-existing unrelated cupti +
insert_batch errors in vendor/cudarc and ml/tests). ml-alpha + ml-
backtesting lib tests pass (33 + 33).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Two follow-up fixes for Smoke 2 readiness, atomic per feedback_no_partial_refactor:
FIX 1 — Task 11 target_jitter_k anchor (spec §3.3 compliance):
- Previously: target_jitter_k = first non-zero raw_per_h[k] (model-state-anchored)
- Now: target_jitter_k = sqrt(realised_label_variance_k) (label-distribution-anchored)
- Per-horizon label variance computed host-side from existing per-horizon
labels in stg_labels (layout [K, B, N_HORIZONS] row-major; typical K=32,
B=1 -> 32 floats per horizon -> trivial host reduction).
- Wiener-α EMA with α = 0.4 floor + first-observation bootstrap per
pearl_wiener_alpha_floor_for_nonstationary +
pearl_first_observation_bootstrap.
FIX 2 — Task 10 Option B closure (block-diagonal heads via grad mask):
- Two new kernels: heads_w_skip_mask_init_kernel (one-shot at transition,
builds [N_HORIZONS × HIDDEN_DIM] = 640-float mask + zeros off-bucket
heads_w_skip in place) and heads_w_skip_grad_mask_apply_kernel
(per-step Phase 2, multiplies grad_heads_w_skip by mask elementwise
before opt_heads_w_skip.step).
- Achieves block-diagonal heads_w_skip behavior WITHOUT touching the
existing GRN kernel: off-bucket positions stay 0 throughout training
because gradient is masked out, so Adam never updates them.
- Existing GRN dispatch consumes heads_w_skip_d (full 640 floats) as
today; off-bucket entries are always 0, so the dispatch is
mathematically equivalent to a compact-only read.
Together these fixes restore the spec-mandated per-horizon differentiation
chain across all 3 mechanisms (CfC.tau bucketing + block-diagonal heads
+ per-bucket LR with label-variance anchor) before Smoke 2 fires the
mean_run_len ratio gate.
GPU oracle tests: 13 total (11 + heads_w_skip_mask_init + heads_w_skip_grad_mask_apply).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per spec §3.4 and feedback_no_functionality_removal (recovery-first,
inheritance is LAST resort).
New device kernels in bucket_transition_kernels.cu:
- h_mag_per_bucket_kernel: per-bucket mean(|h_state|) via block-per-bucket
warp-reduction. Launch grid=(N_HORIZONS,1,1), block=(32,1,1).
No atomicAdd (block-tree reduce over 32 lanes).
- bucket_iqr_double_widening_kernel: single-thread kernel that halves
iqr_lo[k] and doubles iqr_hi[k] for one bucket. Recovery Attempt 1's
actual implementation.
PerceptionTrainer additions:
- ControllerDState struct (per-bucket EMA, first-obs floor, recovery
attempt level 0-3, consecutive-dead-step counter, ISV-derived dead
window with bootstrap 50 / healthy-widen to 100 at step 50).
- phase2_step_count counter (Phase 2 only).
- h_mag_per_bucket_d device buffer + mapped-pinned shadow.
- Cached h_mag_per_bucket_fn + bucket_iqr_double_widening_fn handles.
Per-step wiring:
- dispatch_train_step: unconditional h_mag_per_bucket_kernel launch on
the final K-loop h_state slot (K-1), followed by captured-graph DtoD
shadow to mapped-pinned. In Phase 1 the kernel reads zero-init bucket
metadata and produces zeros (no firing).
- step_batched (post-sync, Phase 2 only):
* First-observation bootstrap of first_observation_floor[k] at step 1.
* Wiener-α=0.4 EMA update of h_mag_ema[k]
(pearl_wiener_alpha_floor_for_nonstationary).
* Dead-threshold = first_observation_floor[k] * 1/e per spec §3.4.
* Consecutive-dead counter; recovery cascade fires at dead_window_k.
Recovery cascade per spec §3.4:
- Attempt 1 (widen): WIRED. Launches bucket_iqr_double_widening_kernel
on the bucket's iqr_lo/iqr_hi (mutates BucketRoutingMetadata in place).
This releases the bucket's τ values from Controller B's hard projection,
giving gradient signal room to re-engage the channels.
- Attempt 2 (channel swap): LOG-ONLY STUB. Emits tracing::warn! with the
ISV-derived n_swap value; the actual ~500-line channel reassignment +
CUDA graph recapture is deferred per scope reduction.
- Attempt 3 (inheritance): LOG-ONLY STUB. Emits tracing::warn! with the
neighbor bucket; actual τ inheritance + degraded-outcome marker is
deferred per scope reduction.
Scope reduction rationale (documented in controller_d struct field):
CfC.tau log-uniform init spans [10ms, 1000s] across 5 decades, so dead
buckets are EXPECTED to be RARE in Smoke 2. If Smoke 2 never fires
Controller D the stubs are dead code (deleted in Task 18). If Recovery 1
suffices (most likely path), no further work needed. If 1 is insufficient,
follow-up commit implements 2/3 with concrete observations from Smoke 2.
Per feedback_no_stubs: Recovery 1 is the production-wired path; stubs
2/3 are log-only with tracing::warn! so any firing surfaces in logs.
The log boundary is an explicit scope decision, not a deferral.
ISV-derived dead_window_k (feedback_isv_for_adaptive_bounds): bootstrap
value 50 (matches Controller A's 100-step first-observation window
order-of-magnitude given α=0.4 EMA settling ≈ 2-3 steps). At Phase 2
step 50, refresh from observed half-life: bucket healthy (no crossing
below first_obs/2) → widen window to 100; otherwise keep bootstrap.
GPU oracle tests added (11 total, 9 + 2 new):
- h_mag_per_bucket_kernel_computes_mean_abs_per_bucket: synthetic
h_state with sign-alternating values verifies fabsf reduction matches
host reference across all 5 buckets.
- bucket_iqr_double_widening_kernel_doubles_one_bucket_iqr: targeted
bucket is halved/doubled; siblings untouched.
Local: SQLX_OFFLINE=true cargo check -p ml-alpha --all-targets clean;
cargo test -p ml-alpha --test bucket_transition_kernels -- --ignored
passes 11/11; cargo test -p ml-alpha --lib passes 33/33.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per spec §3.2 + §3.3 and pearl_adam_normalizes_loss_weights.
Controller B (post-Adam τ projection):
- tau_clamp_kernel added to bucket_transition_kernels.cu
- Each channel's τ clamped to its bucket's IQR_widened range after every
Adam step on tau_all_d. Bypasses Adam's m/sqrt(v) normalization
(loss-weight modulation would have been a no-op per the pearl).
Controller C (per-bucket LR multiplier on heads_w_skip):
- heads_lr_multiplier_scale_kernel added.
- Per-step: snapshot heads_w_skip_d (DtoD) before opt_heads_w_skip.step,
run Adam, then rescale the per-horizon delta by
lr_mult_k = 1.0 / (1.0 + smoothness_ratio_k).
- smoothness_ratio_k = max(0, jitter_ema_k - target_jitter_k) /
max(target_jitter_k, target_jitter_k / 16)
with the ε_floor max-with-floor pattern per
pearl_blend_formulas_must_have_permanent_floor.
- target_jitter_k sentinel-bootstrapped from first non-zero raw
smoothness loss per pearl_first_observation_bootstrap.
- jitter EMA shadowed via mapped-pinned smoothness_jitter_ema_host_d
(DtoD captured in graph; host reads after end-of-step sync); host
writes lr_mult_per_horizon_staging (mapped-pinned); next step's
captured-graph rescale kernel reads via dev_ptr.
Scope reduction (Task 11): Controller C scales ONLY heads_w_skip_d
per-horizon (N_HORIZONS × HIDDEN_DIM = 640 contiguous floats, per-horizon
stride HIDDEN_DIM). Other heads weights (w1, w2, w_gate, w_main) stay
unscaled because they're entangled in the GRN gate/main computation and
rescaling them would require additional per-bucket scoping out of Task 11
scope. Revisit if Smoke 2 fails per feedback_no_quickfixes.
Both controllers active Phase 2 only; no-op in Phase 1.
GPU oracle tests: 9 total (7 existing + tau_clamp + lr_multiplier_scale).
All 33 ml-alpha lib tests pass.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per spec §2.1 + §5.4 point 1 and Task 10 of the plan.
Adds binding `cfc_step_per_branch_fwd_gpu` in cfc/step.rs that wraps
the fused per-branch kernel with the spec-mandated launch config:
grid=(B, N_HORIZONS=5, 1), block=(MAX_BUCKET_DIM=28, 1, 1) uniform
predicate, cooperative staging of x_local + h_old_local (2 × HIDDEN_DIM
floats). Also adds `cfc_step_per_branch_bwd_gpu` binding for Task 11.
CfcTrunk loads two new cubins (cfc_step_per_branch + heads_block_diagonal)
and caches three function handles. `heads_w_skip_compact_d` (from Task 9)
remains allocated as Phase-2-prepared metadata; full heads consumption
deferred to a follow-up task (Phase 2 heads dispatch needs GRN
integration, which is out of Task 10 scope per
`feedback_no_partial_refactor`).
In `dispatch_train_step`, the CfC dispatch branches on TrainingPhase:
- Phase1Warmup: existing `cfc_step_batched` (unchanged)
- Phase2Routed: `cfc_step_per_branch_fwd` consuming bucket-routed
`tau_all_d` + `bucket_channel_offset_d` + `bucket_dim_k_d`
GRN heads dispatch stays unchanged for both phases per the Task 10
scope adjustment (Option B in the task brief). The `heads_w_skip` storage
remains 640 floats; the compact 128-float `heads_w_skip_compact_d` is
metadata-ready but not yet consumed.
Workspace cargo check clean (ml-alpha lib + trainer compile; only
pre-existing cupti/sp15/gpu_per_integration test errors remain).
ml-alpha lib tests: 33 passed.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Per spec §3.1, §2.3. Controller A computes tau-change EMA with
first-observation bootstrap + Wiener-α floor 0.4 (co-adapting closed
loop); trigger when EMA < max(noise_floor/4, noise_floor/16); ISV-derived
hard cap from observed MAD/median dispersion in first 100 steps.
execute_transition() dispatches the 5 device kernels in order:
tau_sort → bucket_assign → bucket_iqr → tau_reorder → heads_compact.
All-on-device per feedback_no_htod_htoh_only_mapped_pinned (one-time
static-constant upload at transition is allowed off-hot-path).
Returns raw Q1/Q3 in BucketRoutingMetadata.bucket_tau_iqr_{lo,hi}_d;
slack_factor=sqrt(Q3/Q1) widening kernel applied by trainer in Task 9.
3 CPU-only unit tests verify Controller A's first-obs bootstrap,
threshold-trigger, and ISV-cap-trigger semantics.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
