c7ccf0c301e30f8fa0ead9dd4801ad83d7f3364a
5471 Commits
| Author | SHA1 | Message | Date | |
|---|---|---|---|---|
jgrusewski
|
c7ccf0c301 |
feat(rl): R7d — PER wired + off-policy DQN with stop-grad on encoder
Closes plan A9 (rebuild plan's "PER wiring" R7 scope second half;
R7c-data shipped the first half last commit). The `ReplayBuffer` in
`src/rl/replay.rs` has sat as dead code since Phase C — this commit
makes it load-bearing per `feedback_always_per` ("PER always enabled;
non-PER paths are dead code").
## Architecture: off-policy Q + on-policy PPO + V + stop-grad encoder
Shared-encoder pattern with the canonical off-policy + shared-encoder
discipline: the Q head trains from PER-sampled past transitions,
PPO + V train on current-step on-policy data, and the encoder receives
gradient signal ONLY from PPO + V (and BCE/aux via the perception
trainer's separate `step_batched` path). Standard pattern in SAC,
R2D2, IMPALA.
Stop-grad is implemented by computing Q's `grad_h_t` (via
`backward_to_w_b_h(sampled_h_t, ...)`) but NOT accumulating it into
`grad_h_t_combined_d` — the encoder backward only sees π + V
contributions. Per `feedback_no_hiding` the discarded buffer is
allocated and written (the kernel API requires the writeback target);
the discard is a deliberate design call documented at the
accumulation site.
## Wiring summary
### Kernel: `dqn_distributional_q_bwd`
* New `loss_per_batch [B]` output. Atom 0 of each block writes the
per-sample CE loss (non-atomic — single writer per batch).
`loss_out [1]` continues to atomicAdd the scalar sum for the
diagnostic total. Build.rs cache bust v30.
### `DqnHead::backward_logits` (Rust wrapper)
* New `loss_per_batch: &mut CudaSlice<f32>` arg. Migrated atomically
in the same commit per `feedback_no_partial_refactor` — only
caller is the integrated trainer.
### `IntegratedTrainerConfig`
* New `per_capacity: usize` (default 4096, matches `replay.rs` doc
ceiling for naive O(N) sampling).
* New `per_seed: u64` (default 0x9E37_79B9_7F4A_7C15).
* `Default` impl added so test fixtures forward-compat via
`..IntegratedTrainerConfig::default()`. All 5 existing test
fixtures migrated.
### `IntegratedTrainer`
* New fields: `replay: ReplayBuffer`, `sampled_h_t_d`,
`sampled_h_tp1_d`, `sampled_actions_d`, `sampled_rewards_d`,
`sampled_dones_d`, `sampled_next_actions_d`, `td_per_sample_d`.
* New methods: `push_to_replay(b_size)` — DtoH per-batch metadata
(action/reward/done/log_pi_old) + alloc per-transition
`CudaSlice<f32>(HIDDEN_DIM)` ×2 + DtoD per-batch slice copies +
push to `ReplayBuffer`. `sample_and_gather(b_size)` — read
per_α from ISV[405], call `replay.sample_indices`, gather sampled
transitions' h_t/h_tp1 device payloads via per-batch DtoD into
`sampled_h_t_d` / `sampled_h_tp1_d`, HtoD upload action/reward/done.
### `step_with_lobsim` orchestration
After `compute_advantage_return` and BEFORE `step_synthetic`:
1. DtoH full ISV slice to refresh `isv_host` (so PER reads ISV[405]
for per_α).
2. `push_to_replay(b_size)` — push current step's transitions.
3. `sample_and_gather(b_size)` — return `per_indices` for the
priority update.
4. `step_synthetic(snapshots)` — runs π + V on current-step h_t,
Q on SAMPLED h_t (off-policy).
5. DtoH `td_per_sample_d` → host; `replay.update_priorities(
per_indices, td_per_sample_host)`.
6. Target-net soft update (unchanged from R5).
### `step_synthetic` redirects (Q path → sampled, π/V stay on-policy)
* Q forward: `forward(&self.sampled_h_t_d)` (was `h_t_borrow`).
* New: forward online Q on `&self.sampled_h_tp1_d` → local scratch +
`argmax_expected_q` → `self.sampled_next_actions_d`. The
Double-DQN argmax MUST be recomputed each step (online net weights
drift faster than transitions recycle through replay; storing
argmax at push time would feed stale-action data into the
projection).
* `forward_target(&self.sampled_h_tp1_d)` (was `&self.h_tp1_d`).
* `select_action_atoms(..., &self.sampled_next_actions_d, ...)`
(was `&self.next_actions_d`).
* `project_bellman_target(..., &self.sampled_rewards_d,
&self.sampled_dones_d, ...)` (was `rewards_d` / `dones_d`).
* `backward_logits(..., &self.sampled_actions_d, ...,
&mut self.td_per_sample_d, ...)` (added per-sample loss output).
* `backward_to_w_b_h(&self.sampled_h_t_d, ...)` (was `h_t_borrow`).
* Q grad_h_t accumulation REMOVED from Step 10 (stop-grad).
## Test: r7d_per_wiring.rs (gate G6)
Three invariants per `pearl_tests_must_prove_not_lock_observations`:
1. `replay.len()` grows by exactly `b_size` per `step_with_lobsim`
call (push semantics).
2. `sample_indices(b_size, α)` returns vec of length `b_size` on a
non-empty buffer.
3. Buffer caps at `per_capacity` (ring-with-random-replacement).
Drives 15 steps with `per_capacity=8`, asserts growth 0→5→8 across
the cap boundary.
## Acceptable host traffic this commit adds
* Per-step DtoH of 4 × b_size scalars (action/reward/done/log_pi_old)
for PER push metadata.
* Per-step DtoH of b_size floats (td_per_sample_d) for
update_priorities.
* Per-step HtoD of 3 × b_size scalars (sampled action/reward/done)
for sampled metadata gather.
* Per-step DtoD of 2 × b_size × HIDDEN_DIM floats (per-batch h_t /
h_tp1 slices) for PER push + gather.
PER bookkeeping is a control-plane operation by design (host-side
priority/index management); the device-side training hot path
(encoder, Q/π/V forward/backward, Adam) stays GPU-pure. GPU sum-tree
+ device-resident transitions are a Phase R-future optimization
flagged in `replay.rs`'s doc.
## What's NOT in this commit
* Q `loss_per_batch [B]` is now wired through `backward_logits` but
the DtoH happens inside step_with_lobsim (not inside
step_synthetic). Earlier R7d sketches considered a separate
`dqn_offpolicy_step` method; the in-step_synthetic redirect
approach landed because it touches fewer lines + reuses the
existing scratch buffer allocations + matches the trainer's
established λ-weighted multi-head pattern. A future refactor
could split for clarity.
Local sm_86 smoke gates: `cargo test -p ml-alpha --test
r7d_per_wiring -- --ignored --nocapture` (G6) +
`integrated_trainer_smoke` (end-to-end). Cluster smoke deferred to
R9.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
acde2e8932 |
fix(rl): R7c-data — true h_{t+1}/V(s_{t+1}) closes Bellman approximation
Closes the long-standing "h_t as proxy for s_{t+1}'s encoder
representation" approximation introduced in Phase E.2's Bellman target
build (canonical comment at the call site: "A future enhancement
(Phase E.3 LobSim integration) will pass next_h_t separately"). The
approximation also leaked into compute_advantage_return's V(s_{t+1})
input — R7b's `v_tp1_d_ref = &v_pred_d` alias — and into R4's
argmax_expected_q kernel call, which had been computing the
Double-DQN argmax on online Q at h_t since R4 first wired it.
Three downstream consumers now read TRUE h_{t+1}:
* `value_head.forward(&self.h_tp1_d) → v_pred_tp1_d`, fed to
`compute_advantage_return` as the canonical TD target V(s_{t+1}).
Was an alias of v_pred_d (V(s_t)) — bootstrap was wrong by one
time index.
* `dqn_head.forward(&self.h_tp1_d) → q_logits_tp1_d`, fed to
`argmax_expected_q` for `next_actions_d` (Double-DQN online-Q
argmax on h_{t+1}, not h_t).
* `dqn_head.forward_target(&self.h_tp1_d)` inside step_synthetic's
Bellman target build. Replaces the h_t-as-proxy comment with the
R7c data-correctness lift inline-doc.
Wiring mechanics:
* New trainer field `h_tp1_d: CudaSlice<f32>` (`[B × HIDDEN_DIM]`).
Zero-initialised; populated each step by step_with_lobsim.
* `step_with_lobsim` signature gains a `next_snapshots:
&[Mbp10RawInput]` parameter (caller — R8's CLI binary — uses
`MultiHorizonLoader::next_sequence_pair` from R2 to load adjacent
`(s_t, s_{t+1})` windows from real MBP-10 data).
* Encoder is now called TWICE in step_with_lobsim:
1. `forward_encoder(next_snapshots)` first → DtoD copy
`perception.h_t_d → self.h_tp1_d` immediately (before any
consumer reads the slot).
2. `forward_encoder(snapshots)` second → leaves perception's
internal forward state (`h_new_per_k_d`, CfC `h_state_d`)
primed for step_synthetic's encoder backward (which still
redundantly re-runs `forward_encoder(snapshots)` per the
pre-existing pattern — separate compute-redundancy fuse for
Phase R-future).
* Three encoder forwards total per step (down from R7b's 2: one
in step_with_lobsim, one in step_synthetic — R7c adds the
second-snapshot forward in step_with_lobsim). The CfC encoder is
deterministic given its input window (per the canonical
step_with_lobsim header comment), so calling forward_encoder
twice on different inputs in a row yields independent h_t and
h_{t+1} via the trainer's own DtoD copy.
Test impact:
* `integrated_trainer_smoke.rs` passes a `next_snapshots` second
window (synthesised as a +1-tick shift of the snapshot window
— production callers will use R2's `next_sequence_pair` on real
MBP-10 data). Smoke continues to assert finite losses across the
five heads.
Scope split note: this commit handles ONLY the data-correctness
half of plan A9 (rebuild plan's "PER buffer + true V(s_{t+1})
Bellman target" R7 scope). The off-policy DQN-via-PER half lands
in the next commit (R7d): Replay-buffer push, sample,
stop-grad-on-encoder Q redirect, and per-sample |TD| → update_priorities.
Split is per `pearl_no_deferrals_for_complementary_fixes`'s
sequencing-with-architectural-justification carve-out: R7d's PER
push requires the correct h_{t+1} this commit provides, so it MUST
sequence after — and the data-correctness fix is independently
useful (the on-policy DQN path now produces correct Bellman
targets even without the off-policy replay).
Local sm_86 smoke: `cargo test -p ml-alpha --test
integrated_trainer_smoke -- --ignored --nocapture` is the gate.
Cluster smoke deferred to R9.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
c34650a241 |
refactor(rl): R7b — host Thompson/argmax/log_pi → GPU; V stays on device
Closes the last host orchestration in `step_with_lobsim`'s policy path
per `feedback_cpu_is_read_only` + `pearl_no_host_branches_in_captured_graph`.
The flawed Phase F+G shipped three host loops (Thompson sampling over
C51 atoms, argmax over expected Q, log-softmax for log π_old) that each
required a DtoH staging copy of the relevant Q/V/π device buffer per
step. R7b deletes them all in favour of three R4 kernels:
* `rl_action_kernel` — Thompson sampler with per-batch xorshift32
PRNG state (device-resident, no salt-based
host RNG seeding per step).
* `argmax_expected_q` — deterministic Bellman-target argmax over
expected Q per action; distinct selector
from Thompson per
`pearl_thompson_for_distributional_action_selection`.
* `log_pi_at_action` — per-batch log π(action_b) via log-softmax
+ lookup; feeds the PPO importance ratio.
Side effects of the lift:
* The host `read_slice_d` DtoH calls for `q_logits_host`,
`v_pred_host`, `pi_logits_host` are deleted (each was 4 × b_size ×
Q_N_ATOMS or b_size × N_ACTIONS floats per step). Three full
device→host→device roundtrips per step → zero.
* The host γ readback (ISV[400] mirror DtoH + bootstrap fallback
that the flawed branch kept "just in case") is also gone:
`compute_advantage_return` already reads γ on-device from
ISV[400] (R3). The fallback was a smell that R7a planned to
eliminate; R7b actually removes it.
* V(s_t) stays in `v_pred_d` throughout the hot path. The host
upload via the (now-deleted) `upload_f32` helper that round-tripped
V back to device for `compute_advantage_return` is gone. V(s_{t+1})
still reuses `v_pred_d` (h_t bootstrap) — true V(s_{t+1}) via
`forward_encoder(next_snapshots)` is explicitly scoped to R7c.
* Step-counter increment retained (drives the Thompson-vs-argmax
diagnostic ordering in the trainer's stats record), but the
ChaCha8Rng-from-step_counter that seeded the host loop is gone.
Per `feedback_no_hiding` the three retired helper functions
(`upload_f32`, `upload_i32`, `argmax_f32`) are DELETED rather than
`#[allow(dead_code)]`'d. The matching unused imports
(`MappedI32Buffer`, `DevicePtrMut`) are removed in the same commit.
The `actions_d` allocation that used to live in step_with_lobsim's
local scope (used only to bridge the host Thompson loop → the
GPU `actions_to_market_targets` kernel) is gone too — that kernel now
reads directly from `self.actions_d` written by `rl_action_kernel`,
closing the last action-path host upload.
What's still HtoD in the hot path (intentional, boundary data):
* `apply_snapshot(last_snap)` — the trainer-fed MBP-10 input
crosses the I/O boundary; mapped-pinned per
`feedback_no_htod_htoh_only_mapped_pinned`.
* HEALTH_DIAG ISV readback inside `step_synthetic` — explicit
diagnostic surface, not hot-path orchestration.
Falsifiability gate G4 (R4 kernel oracle tests + R5 controller +
soft-update tests) already passing on this branch. The R6/R7a smoke
(`integrated_trainer_step_with_lobsim_runs_without_panic`) re-builds
clean and the only host code remaining in `step_with_lobsim` is the
bounded `LaunchConfig`/`b_size_i` setup before each kernel launch.
R7c (next): PER buffer wiring (push transitions, sample, update
priorities) per `feedback_always_per`, and `next_snapshots` +
`forward_encoder(next_snapshots)` for true V(s_{t+1}) Bellman target
(replaces the h_t bootstrap reuse above).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
aba8ec61b2 |
feat(rl): R7a — lift remaining host work in step_with_lobsim to GPU
Honors R6's commit-message promise to close the host-work boundary
the partial GPU-purity left behind. After R7a, step_with_lobsim's
post-fill pipeline is GPU-resident through the entire training step
except for 3 small host-slice uploads (actions, next_actions,
log_pi_old) that R7b lifts via R4's Thompson/argmax/log_pi kernels.
CHANGES:
1. New cuda/abs_copy.cu — element-wise dst[b] = fabsf(src[b]).
Feeds the |reward| signal into ema_update_on_done for the
MEAN_ABS_PNL_EMA slot without mutating signed rewards_d.
2. New trainer-owned per-step device buffers (allocated once in
new(), reused every step — no per-call churn):
reward_abs_d, actions_d, next_actions_d, log_pi_old_d,
advantages_d, returns_d
3. step_synthetic signature change: drops the 7 synthetic_* host
slice args (synthetic_actions, synthetic_rewards, synthetic_dones,
synthetic_next_actions, synthetic_advantages, synthetic_returns,
synthetic_log_pi_old). The body now reads from trainer-owned
device buffers (self.actions_d, self.rewards_d, etc.) via the
disjoint-field borrow rule. The 7 upload_i32/upload_f32 calls
are deleted — caller (step_with_lobsim) populates the buffers
via GPU kernels before invoking step_synthetic.
4. step_with_lobsim Step 6 rewritten end-to-end:
- Upload host Thompson outputs (actions, next_actions, log_pi_old)
to trainer buffers — R7b removes these via R4's GPU kernels.
- GPU abs_copy(rewards_d) → reward_abs_d.
- GPU ema_update_on_done(MEAN_ABS_PNL_EMA, reward_abs_d, dones_d).
- GPU launch_rl_controllers_per_step (R5) — all 7 controllers
adapt to this step's EMA inputs.
- GPU apply_reward_scale(rewards_d) in-place — reads the freshly
updated ISV[406] from the controller.
- GPU compute_advantage_return(rewards_d, dones_d, v_t, v_tp1)
→ returns_d, advantages_d.
- step_synthetic(snapshots) consumes all trainer device buffers,
runs the training kernel chain, returns stats.
- dqn_head.soft_update_target(isv_d) — R5 target-net Polyak
update with τ from ISV[401].
R7a PARTIAL — work R7b lifts:
- Host Thompson sampler still produces actions/next_actions/log_pi
(R7b replaces with R4 rl_action_kernel + argmax_expected_q +
log_pi_at_action — eliminating the 3 remaining HtoD uploads).
- v_pred_host → v_pred_d HtoD round-trip (the host already has
v_pred_host from the action-sampling Thompson read; R7b keeps V
on device throughout).
- v_tp1_d uses v_pred_host (V(s_t) approximated as V(s_{t+1}) until
R7b wires next_snapshots + forward_encoder(next_snapshots)).
The 4 final DtoH copies the R6 commit message warned about are
GONE. Hot path: snapshot upload (boundary HtoD), ISV diagnostic
readback (HEALTH_DIAG), and 3 host-Thompson uploads (R7b removes).
cargo check + cargo build --tests on ml-alpha green. Tests still
compile against the new step_synthetic signature (no test calls it
directly — only step_with_lobsim does).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
7e7c8b5d90 |
feat(rl): R6 — replace LobEnv with kernel-driven GPU-pure env step
Replaces the flawed Phase F+G LobEnv trait + LobSimEnvAdapter +
MockLobEnv fixture with a GPU-pure env-interaction path.
DELETED:
- LobEnv trait (host-method-per-batch surface)
- MockLobEnv fixture (the toy-bandit pattern that hid the production defects)
ADDED:
- RlLobBackend trait (src/rl/reward.rs): narrow device-oriented
surface (apply_snapshot, pos_and_market_targets_mut, pos_d,
pos_bytes, step_fill_from_market_targets, n_backtests). Single
purpose: break the ml-alpha ↔ ml-backtesting dep cycle. No
mocking layer.
- 3 new GPU-pure kernels (cuda/):
* extract_realized_pnl_delta.cu — reads pos.realized_pnl +
position_lots from device Pos array, writes per-batch
(reward delta, done flag), updates trainer's prev_* buffers.
* apply_reward_scale.cu — element-wise rewards *= ISV[406].
* actions_to_market_targets.cu — 9-action grid → market_targets
{side, size} on device, reading current position_lots for
conditional Flat-from-Long/Short.
- LobSimCuda impls RlLobBackend (ml-backtesting/src/sim/mod.rs)
plus a new step_fill_from_market_targets entry that runs
submit_market_immediate + step_pnl_track without the host-side
targets-vec build. pos_and_market_targets_mut returns disjoint
field borrows (&pos_d, &mut market_targets_d).
- IntegratedTrainer: 3 cubin includes, 3 module/function fields,
launch_apply_reward_scale launcher, prev_realized_pnl_d /
prev_position_lots_d / rewards_d / dones_d buffers,
step_with_lobsim signature switched to RlLobBackend, body's
Step 5 rewritten as kernel-driven (no per-batch host loop, no
individual submit_action calls).
R6 PARTIAL — work R7 lifts:
- Thompson + log_pi stay host (R7 uses R4 kernels)
- mean_abs_pnl EMA stays host (R7 uses R3 ema_update_on_done)
- Advantage/return stays host (R7 uses R3 compute_advantage_return)
- 4 final DtoH copies for step_synthetic's host-slice signature
(R7 lifts step_after_encoder_forward to device-buffer args)
The "DtoH the device rewards/dones at end" comment in
step_with_lobsim documents the boundary. After R7, hot path has
zero host loops other than the now-unused Thompson host loop
(which R7 retires in favour of rl_action_kernel from R4).
Tests:
- integrated_trainer_smoke.rs: rewritten — real LobSimCuda with
synthetic book, one step_with_lobsim call, asserts finite losses
+ λs sum to 1. Smoke gate, not a convergence test.
- dqn_toy.rs, ppo_toy.rs: retired (empty stubs documenting
rationale). Files preserved so rename history survives. The
MockLobEnv toy-bandit pattern intrinsically couldn't catch the
production defects #1, #2, #5 that motivated this rebuild.
ml-backtesting added as ml-alpha dev-dep (cycle-safe: production
direction is ml-backtesting → ml-alpha; dev edge only loads when
building ml-alpha's own tests/examples).
Build cache-bust v29. cargo check + cargo build for all R-phase
tests green (pre-existing heads_bit_equiv index-OOB persists).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
0a32a3bb89 |
feat(rl): R5 — wire 7 controllers per-step + target-net soft update
Closes defects #3 (controllers never launched) and #4 (target net never soft-updated) from the flawed Phase F+G arc. CONTROLLER SIGNATURE CHANGE (all 7 .cu files): Scalar input arg → int input_slot. Each controller now reads its EMA input from ISV[input_slot] directly inside the kernel, eliminating the 7 DtoH-per-step host roundtrips a scalar-arg signature would have required — per feedback_cpu_is_read_only (hot-path must be GPU-pure). Bootstrap path unchanged: kernel reads its output slot, sees sentinel zero, writes *_BOOTSTRAP, and returns BEFORE the input_slot read. So R1's launch_isv_controller_3arg(controller_fn, alpha=0.4, input_slot) works both at bootstrap (input read deferred via early return) and at per-step (input slot has real EMA observation from R3 producers). PER-STEP CONTROLLER LAUNCHER: New IntegratedTrainer::launch_rl_controllers_per_step() fires all 7 controllers in sequence, each with its dedicated EMA input slot: ISV[400] γ ← ISV[417] MEAN_TRADE_DURATION_EMA ISV[401] τ ← ISV[418] Q_DIVERGENCE_EMA ISV[402] ε ← ISV[419] KL_PI_EMA ISV[403] entropy_coef ← ISV[420] ENTROPY_OBSERVED_EMA ISV[404] n_rollout_steps← ISV[421] ADVANTAGE_VAR_RATIO_EMA ISV[405] per_α ← ISV[422] TD_KURTOSIS_EMA ISV[406] reward_scale ← ISV[423] MEAN_ABS_PNL_EMA R1's with_controllers_bootstrapped also updated to pass the input slot indices (the bootstrap path still ignores them via early return). TARGET-NET SOFT UPDATE (defect #4): New cuda/dqn_target_soft_update.cu — element-wise target[i] = (1-τ)·target[i] + τ·current[i] reading τ from ISV[401]. Trivially parallel, no atomicAdd. DqnHead gains target_soft_update_fn + _target_soft_update_module fields + soft_update_target(&isv_d) method that fires the kernel twice (weights + biases). R6 calls this from step_with_lobsim after the Q-head Adam update. GATE TESTS (tests/r5_controllers_and_soft_update.rs): G3: g3_per_step_controllers_move_isv_outputs_when_fed_real_emas - Verifies R1 bootstrap pre-conditions (all 7 output slots at documented bootstrap values; all 7 EMA-input slots at sentinel 0). - Populates each EMA-input slot with a distinct non-zero value via R3's ema_update_per_step bootstrap path (different values per slot so a wrong-slot wiring bug would produce out-of-range outputs). - Verifies the EMA producers wrote what we expected (sanity). - Fires launch_rl_controllers_per_step. - Asserts each output slot moved off its bootstrap value (catches "controller doesn't fire" / "reads wrong slot" / "dead kernel"). G4: g4_dqn_target_soft_update_implements_polyak_formula - Force-overwrite w_d with all-ones (breaks the w==target init symmetry so soft_update has something to blend). - Snapshot w_target (Xavier init values). - Fire dqn_head.soft_update_target with R1-bootstrapped τ=0.005. - For sample indices: assert target_after[i] equals (1-τ)·target_before[i] + τ·1.0 within 1e-6 (exact algebraic identity, not a CPU reference — kernel IS the kernel). - Negative invariant: at least one element changed. Per feedback_no_cpu_test_fallbacks: G3 oracle is the invariant "output != bootstrap after non-trivial input"; G4 oracle is the algebraic identity (1-τ)·a + τ·b applied to the SAME numbers the kernel saw — not a parallel CPU implementation. Build cache-bust v28. cargo check + cargo build --tests on ml-alpha green for all R-phase tests. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
27feb94a49 |
feat(rl): R4 — GPU-resident action sampling kernels
Closes defect #5 prerequisites for the GPU-pure step_with_lobsim that lands in R6. Replaces the host Thompson + host argmax + host log-softmax loops the flawed Phase F shipped in step_with_lobsim (which violated feedback_cpu_is_read_only with 5 DtoH copies + host per-action loops + HtoD action upload per training step). Three new kernels: 1. rl_action_kernel.cu — Thompson sampler over the C51 atom distribution. One block per batch, N_ACTIONS=9 threads. Each thread softmaxes its action's Q_N_ATOMS=21 atoms, samples one atom via CDF walk, writes sampled return to shared mem. Thread 0 argmaxes over per-action sampled returns + writes actions[b] + advances per-batch PRNG state. PRNG: per-batch xorshift32 state in prng_state_d (allocated + host-seeded from cfg.dqn_seed via ChaCha8 at trainer init per pearl_scoped_init_seed_for_reproducibility, with .max(1) guard since xorshift32 freezes at 0). Each per-action thread XORs its action index (golden-ratio mixed) into a thread-local copy of the per-batch state — no inter-thread race, reproducible by (cfg.dqn_seed, b_size, step_count). No cuRAND dep. 2. argmax_expected_q.cu — Bellman-target argmax over expected Q per action. Same layout as rl_action_kernel but deterministic (no PRNG). Per pearl_thompson_for_distributional_action_selection: Thompson for rollout (rl_action_kernel), argmax for Bellman target (this kernel) — distinct kernels, distinct ISV consumers. 3. log_pi_at_action.cu — per-batch log π(actions[b] | s_b) via log-softmax + lookup. One thread per batch entry (N_ACTIONS=9 is small enough for a per-thread sequential loop). Feeds the PPO importance ratio in R6. IntegratedTrainer gains: - 3 cubin includes (rl_action_kernel, argmax_expected_q, log_pi_at_action) - 3 module/function field pairs - 2 new device buffers populated at init: prng_state_d: CudaSlice<u32> of length n_batch atom_supports_d: CudaSlice<f32> of length Q_N_ATOMS=21, values [Q_V_MIN, Q_V_MIN + step, …, Q_V_MAX] = linspace(-1, +1, 21) - 3 launcher methods: launch_rl_action_kernel(q_logits_d, actions_d, b_size) launch_argmax_expected_q(q_logits_d, next_actions_d, b_size) launch_log_pi_at_action(pi_logits_d, actions_d, log_pi_out_d, b_size) GPU-oracle tests in tests/r4_action_kernels.rs (per feedback_no_cpu_test_fallbacks every oracle is analytical, not a CPU reference): R4.1: Thompson under sharp distribution (action 5 has logit=20 on atom 20 / support +1.0; others have logit=20 on atom 0 / support −1.0) collapses to argmax — per-action dominant-atom probability ≈ 1 − 4e-8, so 100/100 trials should pick action 5. Assert ≥99/100 (tolerates one fp-rounding edge near u ≈ 1.0 in CDF walk). R4.2: argmax_expected_q picks the rewarded action under the same sharp distribution. Negative invariant: swap dominant atom to action 2 → next_action follows. R4.3: log_pi_at_action with π logits dominant at action 3 (logit=20, others=0) → log π(3) ≈ 0 within 1e-4. Negative invariant: log π(other action) ≈ −20 within 1e-3. Build cache-bust v27. cargo check + cargo build --tests on ml-alpha green (heads_bit_equiv pre-existing failure persists). Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
6d433784f4 |
feat(rl): R3 — GPU-resident EMA + advantage/return kernels
Closes defect #5 from the flawed Phase F+G arc (feedback_cpu_is_read_only violation in step_with_lobsim's host advantage + EMA loops) by landing the GPU primitives those loops will become in R6. Three new kernels, each with a GPU-oracle gate test: 1. ema_update_on_done.cu — done-gated EMA producer. - Slot-parameterised (one kernel, 3 callers in R5 covering mean_abs_pnl_ema, q_divergence_ema, td_kurtosis_ema). - Shared-mem tree reduce, no atomicAdd (feedback_no_atomicadd). - Per pearl_first_observation_bootstrap: sentinel-zero ISV → first observation replaces directly. Defers bootstrap if mean_obs == 0 to avoid writing a degenerate sentinel that would be re-bootstrapped next call. - Per pearl_wiener_alpha_floor_for_nonstationary: Wiener-α blend on subsequent calls; caller pre-floors α at 0.4. 2. ema_update_per_step.cu — per-step EMA producer (no done-gate). - Slot-parameterised (kl_pi_ema, entropy_observed_ema, advantage_var_ratio_ema, mean_trade_duration_ema in R5). - Same shared-mem tree reduce + bootstrap discipline as ema_update_on_done. 3. compute_advantage_return.cu — element-wise returns[b] = r + γ(1-done)·V(s_{t+1}); advantages[b] = returns − V(s_t). - Reads γ from ISV[400] (R1 bootstrap = 0.99). - Trivially parallel, one thread per batch entry; no atomics. Rust launchers added to IntegratedTrainer: - launch_ema_update_on_done(slot, alpha, obs_d, dones_d, b_size) - launch_ema_update_per_step(slot, alpha, obs_d, b_size) - launch_compute_advantage_return(rewards_d, dones_d, v_t_d, v_tp1_d, returns_d, advantages_d, b_size) 3 cubin includes, 3 module/function fields, loaders in new() between the rl_reward_scale_controller load and the with_controllers_bootstrapped call so the new fields are populated by struct construction. GPU-oracle tests in tests/r3_ema_advantage.rs (per feedback_no_cpu_test_fallbacks every oracle is either the kernel's documented bootstrap behaviour or an analytical property of the formula, not a CPU reference): R3.1: ema_update_on_done bootstrap path — sentinel-zero ISV + one observation k → ISV[slot] == k exactly. Negative invariant: hold-only step (dones all zero) preserves the EMA. R3.2: ema_update_per_step convergence — feed obs=5.0 for 50 steps with α=0.4 → ISV[slot] → 5.0 within 1e-4 (EMA of constant = constant). R3.3: compute_advantage_return formula — r=0, done=0, v_t=v_tp1=k, γ=0.99 → returns=γk=4.95, advantages=(γ−1)k=−0.05. Negative invariant: done=1 + r=0 zeros the future-value bootstrap (returns=0, advantages=−k). Build cache-bust v26. cargo check + cargo build --test r3_ema_advantage on ml-alpha green. Pre-existing heads_bit_equiv.rs index-out-of-bounds failure persists (unrelated; pre-Phase E). Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
da2ad438ca |
feat(rl): R2 — fee model + loader pair API
Ports two scope-complete fixes from the ml-alpha-phase-f-g-flawed reference branch: A7 fee model: - order_match.cu::submit_market_immediate gains 2 new kernel args (cost_per_lot_per_side, total_fees_per_b) mirroring the per-fill fee deduction in resting_orders.cu::apply_fill_to_pos:213-219. Fee deducts from pos.realized_pnl on EVERY fill (open, scale-in, counter); total_fees_per_b accumulates for telemetry. Single-writer-per-block plain += is safe — line 79 has `threadIdx.x != 0 → return` (feedback_no_atomicadd). - LobSimCuda::submit_market launch updated to pass the 2 new args. - LobSimCuda::upload_cost_per_lot_per_side host API lets callers configure ES-realistic fees (≈$1.25/contract/side). Default alloc_zeros = $0; production decision-policy path is unchanged (uploads its own cost via step_decision_with_latency). A8 loader pair API: - MultiHorizonLoader::next_sequence_pair returns (LabeledSequence, LabeledSequence) at adjacent anchors in the same source file. anchor_t sampled from [min_anchor, max_anchor−1) so anchor+1 also fits the upper-bound. Counts as ONE yielded sequence against n_max_sequences. - next_sequence_random and next_sequence_pair share a new private helper build_sequence_at(lf, anchor) -> LabeledSequence that contains the multi-resolution windowing logic. Single source of truth for the build (feedback_single_source_of_truth_no_duplicates). - next_sequence's caller-facing contract (random anchor, one sequence per call) is unchanged — alpha_train.rs supervised pipeline keeps working as-is. No new local tests this phase per the rebuild plan (R2): the fee deduction with default cost=0 is a no-op for existing callers, and G7 in R6 covers the with-fees path via a rebuilt reward_calibration test driving LobSimCuda directly (no LobEnv adapter). cargo check -p ml-alpha -p ml-backtesting + cargo build --tests on ml-backtesting both green; baseline test suite unaffected. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
0840fcfe64 |
feat(rl): R1 — ISV slot extension + 7-controller bootstrap (G1 gate)
Closes defect #1 from the flawed Phase F+G arc: ISV[400..406] were left at alloc_zeros sentinel 0 in production, causing bellman_target_projection (γ=0), ppo_clipped_surrogate (ε=0, entropy=0), and the C51 backward to train against degenerate targets that the MockLobEnv toy fixture (done=true every step, horizon=1) intrinsically could not detect. Three changes: 1. Port crates/ml-alpha/cuda/rl_reward_scale_controller.cu from the ml-alpha-phase-f-g-flawed reference branch (93 lines, unchanged). Add to build.rs KERNELS list; bump cache-bust to v25. 2. Extend src/rl/isv_slots.rs: add 7 new EMA-input slot constants (RL_MEAN_TRADE_DURATION_EMA_INDEX..RL_MEAN_ABS_PNL_EMA_INDEX), RL_SLOTS_END goes 417 -> 424. These are reserved for the EMA producer kernels Phase R3 lands; in R1 they stay at sentinel 0 (asserted by the G1 test). 3. Wire all 7 RL adaptive controllers (γ / τ / ε / entropy_coef / n_rollout_steps / per_α / reward_scale) into IntegratedTrainer: - 7 cubin includes + 7 module/function fields - All 7 loaded in new() via the existing load_cubin pattern - New fn launch_isv_controller_3arg() centralises the shared (isv*, alpha, scalar_input) launch signature - New fn with_controllers_bootstrapped() consumes self and fires each controller once against the freshly-zeroed isv_d; each kernel's first-observation-bootstrap path (per pearl_first_observation_bootstrap) sees sentinel zero in its slot and writes its canonical *_BOOTSTRAP value: ISV[400] γ = 0.99 ISV[401] τ = 0.005 ISV[402] ε = 0.2 ISV[403] entropy_coef = 0.01 ISV[404] n_rollout_steps= 2048 ISV[405] per_α = 0.6 ISV[406] reward_scale = 1.0 - new() ends with `.with_controllers_bootstrapped()?` so every trainer construction site picks this up automatically. This replaces the flawed Phase F approach of host memcpy_htod-ing canonical constants into ISV, which violated feedback_no_htod_htoh_only_mapped_pinned (tests not exempt) AND short-circuited the canonical pearl_first_observation_bootstrap pattern every other adaptive controller in the codebase uses. The launch_isv_controller_3arg helper is reused by Phase R5's per-step controller launches with real EMA inputs sourced from ISV[417..424] — at that point the Wiener-α blend kicks in and the slots adapt away from the R1 bootstrap defaults. Gate G1 (crates/ml-alpha/tests/isv_bootstrap.rs): - Construct IntegratedTrainer - memcpy_dtoh full ISV slice to host - Assert ISV[400..406] equal each kernel's #define *_BOOTSTRAP - Assert ISV[417..424] still at sentinel 0 (R3 wires producers) Per feedback_no_cpu_test_fallbacks: the oracle is the kernel's own *_BOOTSTRAP constant, not a CPU computation. Per pearl_tests_must_prove_not_lock_observations: the test asserts an invariant (bootstrap path wrote the canonical value defined by the kernel), not a tuned magic number. Build clean: cargo check + cargo build --test isv_bootstrap on ml-alpha both green. CUDA-required, #[ignore]'d for non-GPU CI. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com> |
||
|
jgrusewski
|
0efdee4b0c |
docs(plans): rebuild plan v2 — strict memory + GPU-oracle gates
Audited the v1 plan against the full project memory catalog. Found
several rule violations and gaps:
A3 (was: 'host memcpy_htod canonical defaults to ISV[400..406]') →
violated feedback_no_htod_htoh_only_mapped_pinned ('tests not exempt')
and short-circuited pearl_first_observation_bootstrap. Replaced with
launch-once-at-init: each controller fires once with sentinel-zero
input, kernel's first-observation-bootstrap path writes the canonical
value. No host write to ISV. Canonical pattern across the codebase.
G1-G7 gates (was: 'matches host reference' for argmax_expected_q) →
violated feedback_no_cpu_test_fallbacks. Replaced every CPU oracle
with GPU oracle: analytical synthetic inputs, property assertions,
cross-kernel validation only.
Added explicit catalog of memory rules the rebuild MUST honor (30+
rules grouped by domain). Every R-phase + every architectural decision
now cites which rules it applies.
Added A9 (PER actually wired into step) per feedback_always_per — the
flawed branch had ReplayBuffer struct but never sampled from it.
Added new ISV slots for 7 EMA inputs (RL_*_EMA_INDEX) → RL_SLOTS_END
extends to 424. The controllers' inputs live on device, not host.
Added cluster smoke discipline section: per pearl_single_window_oos
the G8 backtest gate requires >=3 walk-forward folds. Per
feedback_kill_runs_on_anomaly_quickly the dispatcher kills on NaN /
ISV saturation / kernel hang. Per pearl_q_spread misaligned, the
dispatcher MUST NOT kill on Q_SPREAD. Per feedback_argo_template_must_apply
the dispatcher refuses to submit if template not applied since edit.
Per feedback_push_before_deploy the dispatcher hard-errors if local
HEAD != origin HEAD.
Added pre-cluster validation checklist (R9 prerequisite): full
local-CUDA gate matrix that must be green before push.
Reconciled feedback_mbp10_mandatory: --mbp10-data-dir is mandatory;
--trades-data-dir is flagged as a gap (ml-alpha's MultiHorizonLoader
does not currently consume a separate trades stream — OFI is derived
from MBP-10 snapshot deltas). Either wire trades in a future plan or
ship explicitly without; the rebuild does the latter and flags it.
Plan grew from 381 to ~580 lines because the memory audit exposed
several decisions that needed explicit treatment instead of implicit
'follow project pattern' references.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
e4c3cc60d2 |
docs(plans): integrated RL trainer GPU-pure rebuild
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> |
||
|
jgrusewski
|
9114374d25 |
feat(rl): LobEnv trait + step_with_lobsim + toy bandit activation (E.3b)
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 |
||
|
jgrusewski
|
4b5ef093de |
refactor(rl): per-K hidden-grad buffer + split fused K-loop (E.3a)
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
|
||
|
jgrusewski
|
7356e3c7bb |
fix(rl): close 3 of 4 deferred items from Phase E.2 (greenfield)
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
|
||
|
jgrusewski
|
2665669b58 |
feat(rl): real Q/π/V forward+backward + per-head Adam (Phase E.2)
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 |
||
|
jgrusewski
|
729f110e00 |
feat(rl): IntegratedTrainer skeleton + loss-balance λ (Phase E.1)
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 (
|
||
|
jgrusewski
|
9732a667cc |
feat(rl): PPO π/V heads + clipped surrogate + 2 ISV controllers (Phase D)
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
|
||
|
jgrusewski
|
56efd96cb2 |
feat(rl): C51 distributional Q-head + PER replay + 2 ISV controllers (Phase C)
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> |
||
|
jgrusewski
|
9ec43fdb9d |
feat(rl): expose forward_encoder() for RL head consumption (Phase B)
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). |
||
|
jgrusewski
|
6a46ded7d3 |
feat(rl): module skeleton for integrated RL trainer (Phase A)
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.
|
||
|
jgrusewski
|
5930d9586e |
perf(loader): cache computed labels arrays to disk (~50x first-run / ~∞ subsequent)
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. |
||
|
jgrusewski
|
57de1a8b4e |
docs(plans): integrated RL trainer (DQN + PPO + BCE + aux) plan
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. |
||
|
jgrusewski
|
dab4794fb0 |
chore(sweep): point backtest at jvv7d clean front-month checkpoint
Updates sweep_smoke_perhoriz_cfc.yaml to use the FIRST clean (front- month-filtered) checkpoint (commit |
||
|
jgrusewski
|
f68e0a1d0d |
revert(loader): multi-resolution default '1:32' (single-scale) after htpp6 falsification
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. |
||
|
jgrusewski
|
30db01ccc8 |
perf(loader): parallel per-file load via rayon par_iter (~4-8x speedup)
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. |
||
|
jgrusewski
|
955613d02d |
perf(ml-alpha): online Welford in generate_outcome_labels_ab (~400x speedup)
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. |
||
|
jgrusewski
|
458d678e9f |
docs(plans): multi-resolution input architecture migration plan
Greenfield Phase 1 plan addressing temporal receptive field mismatch (auc_h1000=0.576 plateau): replace seq_len=32 raw-tick input with 3-scale aggregation [(1,10), (30,10), (100,12)] covering 1510 ticks. 10 tasks total, executed via subagent-driven development. Falsifiable gate: auc_h1000 >= 0.65 on clean front-month data. |
||
|
jgrusewski
|
a06abf60df |
test(ml-alpha): synthetic multi-resolution pipeline test
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. |
||
|
jgrusewski
|
b1bfae2367 |
feat(argo): replace --seq-len with --multi-resolution in alpha-perception
Default '1:10,30:10,100:12' (32 positions, 1510-tick context). Greenfield migration — no legacy seq-len fallback. Operators override per-run via './scripts/argo-alpha-perception.sh --multi-resolution 1:32' for the parity-check config. |
||
|
jgrusewski
|
7e438aeba0 |
feat(alpha_train): --multi-resolution CLI flag
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. |
||
|
jgrusewski
|
8d72c14a8b |
fix(ml): migrate test uploads from clone_htod to mapped-pinned
Per feedback_no_htod_htoh_only_mapped_pinned: mapped-pinned only for
CPU↔GPU; tests not exempt. Previous commit
|
||
|
jgrusewski
|
ab64c0412b |
chore(ml): migrate deprecated cudarc memcpy_* calls
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. |
||
|
jgrusewski
|
7abfd3b0b6 |
chore(ml-alpha-tests): migrate test fixtures to MultiResolutionConfig
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). |
||
|
jgrusewski
|
99982cc8fb |
chore(ml-backtesting): migrate to MultiResolutionConfig
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. |
||
|
jgrusewski
|
dcbc51f432 |
fix(loader): anchor sampler enforces lookback lower bound
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. |
||
|
jgrusewski
|
3b1ed72eaa |
feat(ml-alpha): replace seq_len with MultiResolutionConfig in loader
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. |
||
|
jgrusewski
|
cc40780b80 |
chore(ml): file-level allow(unsafe_code) on 12 CUDA-launch files
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. |
||
|
jgrusewski
|
326fa526ce |
feat(ml-alpha): aggregate_window helper
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. |
||
|
jgrusewski
|
bd60aa6afe |
feat(ml-alpha): MultiResolutionConfig for multi-scale input aggregation
Adds a config struct + CLI grammar for stacking multiple time-scale aggregations into the input window. default_three_scale() returns (1,10) + (30,10) + (100,12) = 32 positions covering 1510 ticks of context. Greenfield design — no legacy single-scale default. |
||
|
jgrusewski
|
20aa345a7c |
fix(loader): soft-validate FrontMonth symbol when no streaming SymbolMapping
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. |
||
|
jgrusewski
|
78a9e08358 |
feat(loader): InstrumentFilter::FrontMonth for cross-quarter ES.FUT data
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.
|
||
|
jgrusewski
|
11c658d3fb |
feat(argo): plumb --instrument-id flag through alpha-perception script + 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. |
||
|
jgrusewski
|
783297e002 | feat(loader): instrument_id filter + outdated test fix + sp18 fingerprint | ||
|
jgrusewski
|
1764cc394b |
revert(aux): F2 mid_price_f32 NaN-on-one-sided was a wrong hypothesis
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> |
||
|
jgrusewski
|
85d3ca5c72 |
fix(aux): three bugs causing systematic anti-correlation (F1+F2)
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> |
||
|
jgrusewski
|
f065cfbaa2 | diag(aux): log pos_fraction once at epoch=0, step=0 (diagnose pw bit-identity) | ||
|
jgrusewski
|
24289f40a5 |
tune(aux-bce): lower POS_WEIGHT_MAX 50→10 (gentler class-weight clamp)
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> |
||
|
jgrusewski
|
0f5d5c7b4a |
feat(aux-supervision): wire BCE + conditional-Huber actual calls (CB5)
CB3+CB4 shipped the kernels with a holding-pattern (grad-zero) call site.
CB5 wires the real calls + per-loss ISV signals.
perception.rs:
- step_batched body: pos_weight computed host-side from
self.last_pos_fraction (n_neg/n_pos clamped to [1.0, 50.0] per E3),
uploaded mapped-pinned to stg_aux_pos_weight_{long,short}
- 4 kernel calls per step (slab mode, K*B*N_AUX_HORIZONS):
* aux_bce_loss_gpu × 2 (prof_long, prof_short) — class-weighted
* aux_huber_masked_loss_gpu × 2 (size_long, size_short) — NaN-mask
from CB1's y_size=NaN at y_prof=0 gives conditional-Huber for free
- Per-loss EMAs replace single aux_huber_ema:
* aux_prof_bce_ema_per_h (BCE EMA per horizon)
* aux_size_huber_ema_per_h (Huber EMA per horizon)
* aux_dir_acc_ema_per_h (unchanged)
- Stop-grad lift condition: aux_prof_bce_ema < 0.4 AND aux_dir_acc > 0.85
for ALL horizons (uses BCE not Huber per E3 — size Huber is
observability-only since the regression scale varies more than the
binary classification quality)
- aux_lift_huber_threshold renamed to aux_lift_prof_bce_threshold
alpha_train.rs:
- AlphaTrainSummary: final_aux_huber_ema_per_h split into
final_aux_prof_bce_ema_per_h + final_aux_size_huber_ema_per_h
- Per-epoch tracing: aux_prof_bce_h{100,300,1000} + aux_size_huber_h{...}
+ aux_dir_acc_h{...} + stop_grad_aux_to_encoder
perception_overfit.rs: synthetic test asserts BCE finite + below ln(2)
chance baseline, size Huber finite, dir_acc >= 0.5.
Synthetic test on RTX 3050:
aux_prof_bce_ema = [0.0142, 0.0142, 0.0142] (30× below threshold)
aux_size_huber_ema = [0.1213, 0.1213, 0.1213] (small)
aux_dir_acc_ema = [1.0, 1.0, 1.0] (perfect)
stop_grad lifted = false (lift fired)
Known limitation (follow-up CB6): both kernels return single joint scalar
over the [K × B × N_AUX_HORIZONS] slab — all per-horizon EMA entries
carry the broadcast joint mean. Per-h gating would require kernel split.
cargo check --workspace --all-targets clean.
cargo test -p ml-alpha --lib: 43 passed.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
|
||
|
jgrusewski
|
842e90aeb0 |
feat(aux-heads): rewrite for A+B paired (BCE + conditional-Huber) (CB3+CB4)
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> |