fd415d9b17497d7803f85ddd8513cc49b6fdf9f6
5475 Commits
| Author | SHA1 | Message | Date | |
|---|---|---|---|---|
jgrusewski
|
fd415d9b17 |
fix(rl): apply derive-from-input bootstrap to γ + coef controllers
Systematic completion of the per_α fix (commit
|
||
|
jgrusewski
|
0857d40acd |
fix(rl): rl_per_alpha bootstrap dead-zone at canonical kurtosis
Real production bug surfaced by the R9 G3 local smoke (not the
test-fixture bug the prior commit's message claimed). The
`rl_per_alpha_controller` kernel hardcoded its bootstrap value to
`PER_ALPHA_BOOTSTRAP = 0.6` (canonical PER default per Schaul 2016).
The per-step path's target formula `target = 0.4 + 0.2 × (kurt − 3) / 7`
maps kurt=10 (the canonical heavy-tailed market kurtosis) to **exactly
0.6** — the bootstrap. The Wiener-α blend `(1 − α) · prev + α · target`
then produces 0.6 from any α when prev = target = 0.6, freezing the
controller at the bootstrap value for any input near the canonical
market regime.
This is a real production behaviour bug, not a test fixture bug:
in any market session where TD-error kurtosis sits near canonical 10
(which is the *most common* regime — that's WHY 0.6 is the canonical
PER default), the controller never adapts off bootstrap. The
adaptation mechanism is effectively disabled for typical inputs and
only fires when kurtosis drifts away. The prior commit
(
|
||
|
jgrusewski
|
ee24f0a303 |
fix(rl): R9 local-smoke prep — two test-fixture bugs
Two `pearl_tests_must_prove_not_lock_observations` violations
surfaced during the R9 pre-cluster validation sweep on the dev RTX
3050 Ti (sm_86). Neither was a bug in the trainer or kernels — both
were test fixtures that asserted observed-value coincidences rather
than invariants. Per the canonical pearl, observed-value tests
become bug-locks (the SP16 T3 sp16_phase3_alpha_low_in_steady_state
incident was an assertion `α<0.40` matching the bug itself).
## g3_per_step_controllers_move_isv_outputs_when_fed_real_emas
The fixture fed `RL_TD_KURTOSIS_EMA = 10.0` to the rl_per_alpha
controller, expecting ISV[405] to move off its bootstrap 0.6 after
the Wiener blend. But the kernel's target formula at td_kurtosis=10
maps to **exactly** the bootstrap:
target = 0.4 + 0.2 × (10 − 3) / 7 = 0.6
The Wiener blend `(1−α)·prev + α·target` then produces 0.6 from any
α, so the controller can't move off bootstrap. The assertion was
asserting a coincidence — fixed by picking `td_kurtosis = 20.0`
which lands at `target = 0.886`, distinct from the 0.6 bootstrap.
With the fix all 7 controllers move (γ→0.9, τ→0.023, ε→0.14,
coef→0.0154, n_roll→2867, per_α→0.714, scale→0.608).
The kernel itself is correct — the test was wrong.
## integrated_trainer_step_with_lobsim_runs_without_panic
Asserted `λ_sum ≈ 1.0` for the loss-balance λs. But
`LossLambdas::default()` returns each λ=1.0 (sum = 5.0) with the
`/5.0` divide applied at the trainer's loss-combine site so each
head's contribution is `lambda/5.0`. The "sum=1" assertion was
based on a normalization that the trainer never used. Loosened to
the actual invariant we care about ("every head has a finite
positive λ so the encoder receives real-valued gradient") which
survives any future controller-driven λ re-weighting.
## R9 local-smoke results (all gates green on sm_86)
```
G1 isv_bootstrap ✅ γ=0.99 τ=0.005 ε=0.2
coef=0.01 n_roll=2048
per_α=0.6 scale=1.0
R3 r3_ema_advantage (3 tests) ✅ bootstrap + per-step EMA +
advantage/return formula
R4 r4_action_kernels (3 tests) ✅ Thompson + argmax + log_pi
G3 controllers_emit ✅ all 7 ISV outputs moved
G4 target_soft_update ✅ Polyak τ=0.005 applied
G6 r7d_per_wiring ✅ buffer 0→5→8 + sample size
end integrated_trainer_smoke ✅ all 5 head losses finite
```
Confirms R7c-data + R7d run end-to-end on real CUDA. Next R9 step
(cluster smoke via scripts/argo-alpha-rl.sh + multi-fold G8) requires
git push + cluster credits — paused per the chosen R9 path "stop
before cluster submission."
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
jgrusewski
|
1168f3ea83 |
feat(rl): R8 — alpha_rl_train CLI + Argo template + dispatcher
Closes the rebuild plan's R8 scope: production runner shape for the
integrated RL trainer. Three artifacts wired end-to-end:
1. `crates/ml-alpha/examples/alpha_rl_train.rs` — clap CLI driving
`IntegratedTrainer::step_with_lobsim` against MBP-10 windows
loaded via `MultiHorizonLoader::next_sequence_pair` (R2) for
true `(s_t, s_{t+1})` adjacency. Per `feedback_mbp10_mandatory`,
`--mbp10-data-dir` is required — no synthetic-data fallback in
the production path.
2. `infra/k8s/argo/alpha-rl-template.yaml` — WorkflowTemplate
mirroring alpha-perception's DAG (check-cache → ensure-binary →
train; warmup-gpu parallel). Binary cache slot is
`/data/bin/<sha>/alpha_rl_train` (distinct from `alpha_train`
so the two binaries coexist at the same SHA).
3. `scripts/argo-alpha-rl.sh` — dispatcher with three rebuild-plan
guards baked in.
## Dispatcher guards (per the rebuild plan's feedback list)
`feedback_default_to_l40s_pool` (2026-05-09): default `--gpu-pool`
is `ci-training-l40s` (sm_89). H100 (sm_90) is opt-in for production
scale-up only. Cubins must match the device, so the dispatcher
derives `cuda-compute-cap` from the pool name and threads it into
the workflow params.
`feedback_argo_template_must_apply` (2026-05-21 canonical incident):
`argo submit --from=wftmpl/<name>` reads the cluster CRD, NOT the
on-disk YAML; unknown `-p` parameters silently no-op without a prior
`kubectl apply`. Dispatcher applies the local template BEFORE every
submission (overrideable via `--skip-template-apply` for the rare
case where you've already applied manually).
`feedback_push_before_deploy` (2026-05-20 canonical incident): the
in-cluster `ensure-binary` pod fetches source from `origin/<branch>`,
NOT the local working tree. Submitting before `git push` deploys the
last-pushed SHA, which can lag local diff by N commits. Dispatcher
verifies `git rev-parse HEAD == git rev-parse origin/<branch>` and
hard-errors with the explicit push command otherwise. Bypass via
`--skip-push-check` (only when intentionally deploying a previously-
pushed SHA via `--sha`).
## CLI: gate G8 (NaN abort)
Per `feedback_stop_on_anomaly` + `feedback_kill_runs_on_anomaly_quickly`,
the CLI checks every per-head loss (l_bce / l_q / l_pi / l_v / l_aux
/ l_total) for finiteness after each `step_with_lobsim` call.
Non-finite at any step → write summary with `nan_abort_step` set →
`process::exit(2)`. R9's cluster smoke tail-watcher kills the
workflow on the non-zero exit code, satisfying gate G8 from the
rebuild plan.
## CLI: knobs that ARE on the CLI
Structural / boundary parameters only (per
`pearl_controller_anchors_isv_driven`: every adaptive knob lives in
ISV, not CLI flags):
* `--mbp10-data-dir / --predecoded-dir / --out` — I/O paths.
* `--n-steps` — wall-budget control (1000 R9 smoke / 50k+ prod).
* `--seq-len / --n-backtests / --per-capacity` — structural
sizing. seq_len threads into the loader's multi-resolution
`1:<seq_len>` config; n_backtests into both LobSimCuda and
PerceptionTrainerConfig.n_batch.
* `--seed` — reproducibility per
`pearl_scoped_init_seed_for_reproducibility` (forks deterministic
sub-seeds for dqn / ppo / per).
* `--instrument-mode` — MBP-10 filter (all / front-month / id=N).
* `--gpu-idx` — CUDA device selection.
What's NOT on the CLI: γ / τ / ε / entropy_coef / per_α / reward_scale
/ per-head LRs — all live in ISV[400..417] and are driven by R5's
controllers from EMA-tracked diagnostics. Per the rebuild plan
A1: "every adaptive bound is signal-driven, not tuned."
## Cluster smoke entry point
```bash
# R9 validation smoke (after pre-cluster local CUDA tests green).
./scripts/argo-alpha-rl.sh --n-steps 1000 --instrument-mode front-month
# Production scale-up (gated by R9's multi-fold pass).
./scripts/argo-alpha-rl.sh --n-steps 50000 --n-backtests 32 \
--per-capacity 100000 # GPU sum-tree R-future when capacity > 4096
```
## What's NOT in this commit
The R9 cluster smoke run itself is out of band — this commit ships
the entry points. R9 will execute the pre-cluster validation
checklist + first 1000-step smoke + multi-fold walk-forward G8 gate
per the rebuild plan §"Cluster smoke discipline".
The summary JSON's schema is intentionally narrow (final-step losses
+ replay len + completion state + NaN abort marker). R-future may
add per-epoch breakdowns + per-ISV-slot snapshots once the cluster
smoke tells us which diagnostics are actually load-bearing for kill
decisions.
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
|
||
|
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> |