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>
This commit is contained in:
jgrusewski
@@ -4,122 +4,148 @@
|
||||
> on branch `ml-alpha-phase-f-g-flawed` (commits 99a125cdb..b3808a5ac).
|
||||
> Baseline: `ml-alpha-phase-a` reset to commit `9114374d2`
|
||||
> (Phase E.3b activation).
|
||||
>
|
||||
> **Discipline:** Smokes are expensive ($5–20 / 30–120 min L40S each)
|
||||
> and producing G8 numbers from a buggy trainer wastes both cluster
|
||||
> credit and gate-evaluation cycles. Every rule below is non-negotiable.
|
||||
> Every gate is falsifiable locally before any `argo submit`.
|
||||
|
||||
## Why this rebuild
|
||||
|
||||
The prior Phase F + G attempt landed 9 commits that shipped a trainer
|
||||
with multiple production-blocking defects. The defects were not bugs in
|
||||
the new code so much as **propagation of pre-existing violations** in
|
||||
`step_with_lobsim` plus **incomplete wiring of the ISV adaptive-control
|
||||
infrastructure**. The convergence-gate fixtures (`dqn_reward_signal`,
|
||||
`ppo_reward_signal`, then named `dqn_toy` / `ppo_toy`) hid every defect
|
||||
because their `MockLobEnv` has `done = true` every step, reward in
|
||||
`[-1, +1]`, and horizon = 1. The cluster-bound smoke was the only
|
||||
mechanism that would have exposed the problems — at $5–20 / 30–120 min
|
||||
each.
|
||||
The prior Phase F + G attempt landed 9 commits with multiple
|
||||
production-blocking defects. None of the defects were exotic bugs;
|
||||
all of them were direct violations of pre-existing project memory
|
||||
that I missed because the `MockLobEnv` "toy bandit" fixture has
|
||||
`done = true` every step, reward ∈ `[-1, +1]`, and horizon = 1 —
|
||||
which hides every failure mode the production cluster will see.
|
||||
|
||||
### Defects in the flawed branch
|
||||
### Defect table — each row maps to a memory rule the flawed branch violated
|
||||
|
||||
1. **ISV[400..406] uninitialised in production paths.** The trainer
|
||||
`alloc_zeros` the ISV array; only `rl_lr_controller` was launched.
|
||||
Production consumers read sentinel `0`:
|
||||
- `bellman_target_projection.cu:133` reads γ = 0 → Bellman target
|
||||
reduces to `r + 0·V(s_{t+1}) = r`. Q-head learns immediate-reward
|
||||
only, zero credit assignment.
|
||||
- `ppo_clipped_surrogate.cu:205,305` reads ε = 0 → `clip(r, 1, 1) = 1`
|
||||
always. PPO trust region collapses.
|
||||
- `ppo_clipped_surrogate.cu:210` reads entropy coef = 0 → no entropy
|
||||
bonus, policy collapses to deterministic immediately.
|
||||
| # | Defect | Memory rule violated |
|
||||
|---|--------|----------------------|
|
||||
| 1 | ISV[400..406] uninitialised in production paths (γ=0, ε=0, entropy=0 in kernels) | `feedback_wire_everything_up`, `feedback_isv_for_adaptive_bounds` |
|
||||
| 2 | `rl_reward_scale_controller` drifts to 1e3 on no-trade steps | `pearl_first_observation_bootstrap` |
|
||||
| 3 | 6 controllers exist as .cu files but never launched | `feedback_wire_everything_up`, `feedback_no_partial_refactor` |
|
||||
| 4 | Target net never soft-updated (τ has no consumer) | `feedback_no_stubs`, `feedback_wire_everything_up` |
|
||||
| 5 | `step_with_lobsim` has 5 DtoH + host argmax + host EMA + host advantage loops | `feedback_cpu_is_read_only`, `pearl_cold_path_no_exception_to_gpu_drives` |
|
||||
| 6 | PER buffer exists in `src/rl/replay.rs` but trainer never samples from it | `feedback_always_per` |
|
||||
| 7 | "Toy bandit" framing leaked into production (`alpha_rl_train.rs` shipped with `next_snapshots = snapshots` until caught mid-review) | `feedback_no_stubs`, `feedback_no_quickfixes`, `feedback_extending_existing_code_audits_for_existing_violations` |
|
||||
| 8 | No NaN / divergence abort in production CLI | `feedback_stop_on_anomaly`, `feedback_kill_runs_on_anomaly_quickly` |
|
||||
| 9 | Convergence-gate fixtures (`dqn_toy.rs`, `ppo_toy.rs`) used `MockLobEnv` with state-invariant rewards — intrinsically can't catch defects #1, #2, #5 | `pearl_tests_must_prove_not_lock_observations`, `feedback_no_cpu_test_fallbacks` |
|
||||
|
||||
2. **`rl_reward_scale_controller` drifts to `1e3` on no-trade steps.**
|
||||
The launcher fired every step regardless of `mean_abs_pnl_ema`
|
||||
readiness; with `ema = 0` and `prev = REWARD_SCALE_BOOTSTRAP = 1.0`,
|
||||
the kernel computed `target = 1/max(0, 1e-3) = 1000`, Wiener-blended
|
||||
toward it, and ramped the scale to `1e3` over the first ~10 hold
|
||||
steps. Toy fixture hid this because `done = true` every step.
|
||||
### Root-cause pattern
|
||||
|
||||
3. **6 controllers exist as .cu files but never launched.**
|
||||
`rl_gamma`, `rl_target_tau`, `rl_ppo_clip`, `rl_entropy_coef`,
|
||||
`rl_rollout_steps`, `rl_per_alpha`. The build system compiles their
|
||||
cubins; no Rust launcher fires them. ISV bootstrap (defect #1)
|
||||
would have masked this even if the controllers had been wired.
|
||||
Defects #1, #3, #5, #6 are all the same shape: **the pre-existing
|
||||
Phase E.3b code already violated current memory rules, and the
|
||||
flawed Phase F extended the pattern instead of flagging it.** The new
|
||||
memory pearl `feedback_extending_existing_code_audits_for_existing_violations`
|
||||
exists specifically to prevent this in future work.
|
||||
|
||||
4. **Target-net soft update missing.** `DqnHead` owns `target_w_d /
|
||||
target_b_d` but no kernel updates them. The "target network" is
|
||||
actually the initial random init for the entire run — no soft Polyak
|
||||
update, no hard sync. Double-DQN Bellman target uses an
|
||||
untrained-static target.
|
||||
## Strict memory rules this rebuild MUST honor
|
||||
|
||||
5. **`feedback_cpu_is_read_only` violation in `step_with_lobsim`.**
|
||||
Host code:
|
||||
- 5 DtoH copies per step (Q / V / π logits at h_t, V / Q at h_{t+1})
|
||||
- Host-side softmax + Thompson categorical sample + argmax loops
|
||||
over `[B][N_ACTIONS][Q_N_ATOMS]`
|
||||
- Host EMA updates for `mean_abs_pnl_ema`
|
||||
- Host advantage / return computation
|
||||
- HtoD copies for `actions`, `rewards`, `dones`
|
||||
The pre-existing Phase E.3b code already had this pattern; the
|
||||
flawed Phase F extended it instead of rebuilding it as GPU-resident
|
||||
per `pearl_cold_path_no_exception_to_gpu_drives`.
|
||||
Explicit catalog, every rule named below is gating. If the rebuild
|
||||
diverges from any of these the work does NOT ship.
|
||||
|
||||
6. **"Toy" framing in production code.** Test files `dqn_toy.rs` /
|
||||
`ppo_toy.rs`, fixture constructor `MockLobEnv::toy_bandit()`, and
|
||||
pervasive "toy bandit" / "for the toy bandit smoke" comments leaked
|
||||
into production source. Worse: `alpha_rl_train.rs` shipped with
|
||||
`next_snapshots = snapshots` — production code applying a
|
||||
toy-fixture simplification — until that one issue got caught
|
||||
mid-review.
|
||||
### GPU/CPU contract
|
||||
- `feedback_cpu_is_read_only` — no CPU compute, no CPU forwards, no GPU↔CPU roundtrips on hot path
|
||||
- `feedback_no_htod_htoh_only_mapped_pinned` — mapped-pinned only; tests not exempt
|
||||
- `feedback_no_atomicadd` — block tree-reduce / per-batch scratch + `reduce_axis0`, never atomicAdd
|
||||
- `feedback_no_cpu_test_fallbacks` — GPU oracle tests only; no CPU reference implementations
|
||||
- `feedback_no_nvrtc` — pre-compiled cubins via `build.rs`
|
||||
- `pearl_no_host_branches_in_captured_graph` — host branches inside captured graphs break recording
|
||||
- `pearl_cold_path_no_exception_to_gpu_drives` — "cold path" is not a license for CPU compute
|
||||
- `pearl_fused_per_group_statistics_oracle` — K groups × N stats = ONE fused kernel
|
||||
- `pearl_build_rs_rerun_if_env_changed` — every `std::env::var()` paired with `cargo:rerun-if-env-changed`
|
||||
|
||||
7. **No NaN / divergence abort in `alpha_rl_train.rs`.** A NaN-poisoned
|
||||
run would burn the full 14400-second `activeDeadlineSeconds` budget
|
||||
silently.
|
||||
### Code quality
|
||||
- `feedback_no_stubs` — wire it for real or delete
|
||||
- `feedback_no_todo_fixme` — no TODO/FIXME/XXX markers
|
||||
- `feedback_no_hiding` — never suppress with `_` prefix or `#[allow]`
|
||||
- `feedback_no_legacy_aliases` — no deprecated wrappers
|
||||
- `feedback_no_feature_flags` — no `enable_*` / `use_*` booleans
|
||||
- `feedback_no_quickfixes` — every issue gets the proper fix
|
||||
- `feedback_no_functionality_removal` — never delete features; fix what's broken
|
||||
- `feedback_fix_everything_aggressively` — zero tolerance, no deferrals
|
||||
- `feedback_extending_existing_code_audits_for_existing_violations` — when extending existing code, audit it FIRST against current rules
|
||||
|
||||
### Pattern lesson
|
||||
### Architecture discipline
|
||||
- `feedback_single_source_of_truth_no_duplicates` — no `v2_` / parallel impls
|
||||
- `feedback_no_partial_refactor` — contract change = every consumer migrates atomically
|
||||
- `feedback_wire_everything_up` — every kernel/module wired in same commit; no orphans
|
||||
- `feedback_isv_for_adaptive_bounds` — adaptive bounds live in ISV, not hardcoded constants
|
||||
- `feedback_adaptive_not_tuned` — signal-driven (EMA/windowed), not tuned constants
|
||||
- `feedback_magnitude_must_be_useful` — DQN magnitude branch never deleted
|
||||
|
||||
The convergence-gate fixtures verified the gradient mechanics
|
||||
(`reward_signal_drives_argmax_to_rewarded_action`) but their
|
||||
state-invariant fixture intrinsically *can't* test:
|
||||
- Sparse-reward credit assignment (defect #1's γ = 0 effect)
|
||||
- Reward-magnitude drift (defect #2)
|
||||
- Trust-region clipping (defect #1's ε = 0 effect)
|
||||
- Target-network staleness (defect #4)
|
||||
### Training discipline
|
||||
- `feedback_always_per` — PER always enabled; non-PER paths are dead code
|
||||
- `feedback_mbp10_mandatory` — `--mbp10-data-dir` mandatory; trades-data-dir paired for OFI completeness
|
||||
- `feedback_default_to_l40s_pool` — defaults to `ci-training-l40s`, not H100
|
||||
- `feedback_stop_on_anomaly` — terminate on training anomaly, diagnose, fix, re-run
|
||||
- `feedback_kill_runs_on_anomaly_quickly` — kill smoke at first useful signal
|
||||
|
||||
**Pre-cluster falsifiability requires fixtures that exercise the same
|
||||
failure modes the cluster will see.** A short bull-market loop with
|
||||
state-dependent rewards + per-step holds + multi-step trades is the
|
||||
minimum.
|
||||
### Workflow
|
||||
- `feedback_push_before_deploy` — always `git push` before `argo submit`
|
||||
- `feedback_argo_template_must_apply` — `kubectl apply -f <template>` before `argo submit --from=wftmpl/<name>`
|
||||
- `feedback_no_concurrent_agents_shared_tree` — worktrees or serial dispatch
|
||||
- `feedback_registry_entries_need_dispatch_arms` — registry entries require matching reset arms in same commit
|
||||
- `feedback_smoke_validation_before_structural_priors` — validate wiring with neutral config first
|
||||
|
||||
### Reward + controller pearls
|
||||
- `pearl_first_observation_bootstrap` — sentinel=0; first observation replaces directly
|
||||
- `pearl_wiener_alpha_floor_for_nonstationary` — α-floor 0.4 in control loops
|
||||
- `pearl_wiener_optimal_adaptive_alpha` — α = diff_var / (diff_var + sample_var + ε); no tuned α
|
||||
- `pearl_blend_formulas_must_have_permanent_floor` — `max(real, floor)`, not blend
|
||||
- `pearl_one_unbounded_signal_per_reward` — exactly ONE unbounded multiplicand per reward term
|
||||
- `pearl_symmetric_clamp_audit` — kernel-derived bounded scalars use bilateral `fmaxf(lo, fminf(x, hi))`
|
||||
- `pearl_audit_unboundedness_for_implicit_asymmetry` — symmetric bounds may erase implicit asymmetry; use asymmetric bounded clamp
|
||||
- `pearl_event_driven_reward_density_alignment` — event-driven objectives need event-density reward
|
||||
- `pearl_bounded_modifier_outputs_require_structural_activation` — bounded outputs feeding multiplicative modifiers MUST use structural activation (sigmoid/tanh), not runtime clamps
|
||||
- `pearl_controller_anchors_isv_driven` — every controller anchor/target/cap is ISV-driven
|
||||
- `pearl_thompson_for_distributional_action_selection` — Thompson at rollout, argmax for Bellman target only
|
||||
|
||||
### Backtest / OOS discipline
|
||||
- `pearl_single_window_oos_is_not_oos` — ≥3 walk-forward folds, report mean ± SD
|
||||
- `pearl_q_spread_kill_criterion_misaligned_for_distributional_q` — DON'T use Q_SPREAD as kill metric (spikes signal GOOD balance under C51+Thompson)
|
||||
- `pearl_tests_must_prove_not_lock_observations` — assert invariants, not observed values
|
||||
- `pearl_action_pruning_falsified` — KEEP 9 actions (pruning 9→4 regressed Sharpe 2.4×)
|
||||
- `pearl_c51_thompson_closed_phase_e3_gap` — C51 + Thompson is the proven combo
|
||||
|
||||
### Kernel hygiene pearls
|
||||
- `pearl_kernel_input_filter_must_match_validation_surface` — per-depth validation
|
||||
- `pearl_behavioral_sentinel_values_poison_downstream_consumers` — per-emission-site counters localize bugs
|
||||
- `pearl_cooperative_staging_eliminates_redundant_reads` — stage per-batch rows once
|
||||
- `pearl_coalesce_via_thread_role_swap` — output-row → output-column thread role for non-coalesced writes
|
||||
- `pearl_reduce_axis0_column_tile` — block-per-column 32-wide column tile
|
||||
|
||||
## Architectural decisions
|
||||
|
||||
### A1. GPU-pure action submission via the existing decision-policy path
|
||||
### A1 — Delete `LobEnv` trait; drive `LobSimCuda` via kernel-driven decision path
|
||||
|
||||
The pre-existing production decision pipeline (`step_decision_with_latency`
|
||||
in `crates/ml-backtesting/src/sim/mod.rs`) already submits actions
|
||||
The production decision pipeline (`step_decision_with_latency` in
|
||||
`crates/ml-backtesting/src/sim/mod.rs`) already submits actions
|
||||
GPU-pure: `alpha_probs_d` (device) feeds `decision_policy_kernel`
|
||||
which writes `market_targets_d` (device) which the fill kernel reads.
|
||||
No host roundtrip.
|
||||
which writes `market_targets_d` (device). No host roundtrip.
|
||||
|
||||
**Decision:** the integrated RL trainer drives the same kernel path.
|
||||
Action sampling happens IN THE DECISION KERNEL (or a new
|
||||
`rl_action_kernel` that lives parallel to it), reading Q / π logits +
|
||||
ISV[entropy, ε] from device + writing to `market_targets_d`. The
|
||||
`LobEnv` trait + `LobSimEnvAdapter` from the flawed branch are
|
||||
DELETED. The trainer interacts with `LobSimCuda` via the existing
|
||||
decision-policy entry point.
|
||||
**Decision:** the RL trainer drives the same kernel path. Action
|
||||
sampling lives in a new `rl_action_kernel` (Thompson over Q logits)
|
||||
parallel to the production decision kernel; both write to
|
||||
`market_targets_d`. `LobEnv` trait, `LobSimEnvAdapter`, and
|
||||
`MockLobEnv` are DELETED. Convergence-gate fixtures use real
|
||||
`LobSimCuda` with synthetic book data.
|
||||
|
||||
**Consequence:** the trait-abstraction layer that let `MockLobEnv` mock
|
||||
the env is removed. Convergence-gate tests must run against a real
|
||||
`LobSimCuda` with synthetic book data (still feasible — fixtures
|
||||
inject snapshots).
|
||||
Rules honored: `feedback_cpu_is_read_only`,
|
||||
`pearl_cold_path_no_exception_to_gpu_drives`,
|
||||
`feedback_single_source_of_truth_no_duplicates` (no duplicate
|
||||
action-submission paths), `feedback_no_legacy_aliases` (no trait
|
||||
adapter left as a shim).
|
||||
|
||||
### A2. All EMAs live in ISV
|
||||
### A2 — All EMAs live in ISV
|
||||
|
||||
Currently `mean_abs_pnl_ema` lives as a host `f32` field on
|
||||
`IntegratedTrainer`. **Decision:** every EMA the controllers consume
|
||||
lives in an ISV slot, updated by a dedicated per-step kernel that reads
|
||||
the source signal from device buffers (rewards_d, dones_d, etc.).
|
||||
Every controller-input EMA lives in an ISV slot, updated by a
|
||||
dedicated per-step kernel that reads the source signal from device
|
||||
buffers (rewards_d, dones_d, actions_d, etc.).
|
||||
|
||||
New ISV slots needed for the EMA inputs:
|
||||
New ISV slots:
|
||||
- `RL_MEAN_TRADE_DURATION_EMA_INDEX` (input to `rl_gamma_controller`)
|
||||
- `RL_Q_DIVERGENCE_EMA_INDEX` (input to `rl_target_tau_controller`)
|
||||
- `RL_KL_PI_EMA_INDEX` (input to `rl_ppo_clip_controller`)
|
||||
@@ -128,254 +154,498 @@ New ISV slots needed for the EMA inputs:
|
||||
- `RL_TD_KURTOSIS_EMA_INDEX` (input to `rl_per_alpha_controller`)
|
||||
- `RL_MEAN_ABS_PNL_EMA_INDEX` (input to `rl_reward_scale_controller`)
|
||||
|
||||
That's 7 new slots → `RL_SLOTS_END = 424`.
|
||||
7 new slots → `RL_SLOTS_END = 424`.
|
||||
|
||||
### A3. ISV bootstrap at trainer construction
|
||||
All EMA updates use `pearl_wiener_optimal_adaptive_alpha` (adaptive α
|
||||
from observed diff/sample variance ratio) with floor 0.4 per
|
||||
`pearl_wiener_alpha_floor_for_nonstationary`. Sentinel-zero bootstrap
|
||||
per `pearl_first_observation_bootstrap`.
|
||||
|
||||
Defect #1 and the Wiener-blend behavior of the controllers both
|
||||
require slots to be at known-good values before any kernel reads them.
|
||||
**Decision:** trainer `new()` does ONE memcpy_htod writing canonical
|
||||
defaults to ISV[400..406]:
|
||||
Rules honored: `feedback_isv_for_adaptive_bounds`,
|
||||
`pearl_controller_anchors_isv_driven`, `feedback_adaptive_not_tuned`.
|
||||
|
||||
| Slot | Default | Source |
|
||||
|------|---------|--------|
|
||||
| RL_GAMMA_INDEX (400) | 0.99 | DQN convention |
|
||||
| RL_TARGET_TAU_INDEX (401) | 0.005 | DQN convention |
|
||||
| RL_PPO_CLIP_INDEX (402) | 0.2 | PPO original paper |
|
||||
| RL_ENTROPY_COEF_INDEX (403) | 0.01 | PPO convention |
|
||||
| RL_N_ROLLOUT_STEPS_INDEX (404) | 2048 | PPO convention |
|
||||
| RL_PER_ALPHA_INDEX (405) | 0.6 | Schaul et al. 2016 |
|
||||
| RL_REWARD_SCALE_INDEX (406) | 1.0 | identity pass-through |
|
||||
### A3 — Bootstrap controllers via launch-once at init, NOT host memcpy_htod
|
||||
|
||||
Controllers then ADAPT from these defaults rather than bootstrapping
|
||||
from sentinel zero. The sentinel-zero bootstrap path in each kernel
|
||||
stays as a safety net but is never the production cold-start path.
|
||||
The prior plan said "memcpy_htod canonical defaults to ISV[400..406]".
|
||||
That violated `feedback_no_htod_htoh_only_mapped_pinned` (tests not
|
||||
exempt) and short-circuited `pearl_first_observation_bootstrap`.
|
||||
|
||||
### A4. Target-net soft update consumer for ISV[401]
|
||||
**Decision:** in `IntegratedTrainer::new`, after the ISV alloc_zeros,
|
||||
launch each controller ONCE with a zero (sentinel) input EMA. Each
|
||||
kernel sees `prev = 0` (sentinel), executes its first-observation
|
||||
bootstrap path, and writes its canonical value (γ=0.99, τ=0.005, ε=0.2,
|
||||
entropy_coef=0.01, n_rollout_steps=2048, per_α=0.6, reward_scale=1.0).
|
||||
|
||||
New `dqn_target_soft_update` kernel: per-step element-wise
|
||||
Subsequent in-loop launches see non-sentinel ISV → execute Wiener
|
||||
blend with real EMA input.
|
||||
|
||||
No host write to ISV. Bootstrap path is canonical, observable in HEALTH
|
||||
diagnostics, and follows the same `pearl_first_observation_bootstrap`
|
||||
pattern as every other adaptive controller in the codebase.
|
||||
|
||||
Rules honored: `pearl_first_observation_bootstrap`,
|
||||
`feedback_isv_for_adaptive_bounds`,
|
||||
`feedback_no_htod_htoh_only_mapped_pinned`,
|
||||
`feedback_wire_everything_up`.
|
||||
|
||||
### A4 — Target-net soft update consumer for ISV[401]
|
||||
|
||||
New `dqn_target_soft_update` kernel: element-wise per-weight
|
||||
`target_w[i] = (1-τ)·target_w[i] + τ·w[i]` reading τ from `ISV[401]`.
|
||||
Parallel across weight elements (no contention; no atomicAdd needed).
|
||||
Launched from `step_after_encoder_forward` after the Q-head Adam
|
||||
update.
|
||||
|
||||
### A5. GPU-resident action sampling
|
||||
Rules honored: `feedback_wire_everything_up`,
|
||||
`feedback_no_atomicadd`, `feedback_no_stubs`.
|
||||
|
||||
### A5 — GPU-resident action sampling kernels
|
||||
|
||||
New kernels:
|
||||
- `thompson_sample_categorical` — input: Q logits, ISV[entropy_coef],
|
||||
device PRNG state. Output: per-batch action index in `market_targets`
|
||||
format directly.
|
||||
- `argmax_expected_q` — input: Q logits. Output: per-batch
|
||||
next-action index.
|
||||
- `log_pi_at_action` — input: π logits + action_d. Output: per-batch
|
||||
log π(a_t) — feeds the PPO importance-ratio path.
|
||||
- `thompson_sample_categorical` — per-batch: softmax over atoms,
|
||||
sample one atom per action via device PRNG, pick argmax over
|
||||
sampled per-action returns. Writes action index to `market_targets_d`
|
||||
directly. Single-writer-per-block. Device PRNG state =
|
||||
`curandStatePhilox4_32_10_t` allocated at trainer init, advanced per
|
||||
step. Reads ISV[entropy_coef] if exploration-temperature scaling
|
||||
applies.
|
||||
- `argmax_expected_q` — per-batch: softmax + expected value + argmax.
|
||||
Writes next_actions for the Bellman target argmax. Single-writer.
|
||||
- `log_pi_at_action` — per-batch: softmax + log-prob lookup at action.
|
||||
Writes log_pi_old for PPO importance ratio. Single-writer.
|
||||
|
||||
Device PRNG via a per-batch `curandStatePhilox4_32_10_t` allocated at
|
||||
trainer init + advanced per step.
|
||||
Per `pearl_thompson_for_distributional_action_selection`: Thompson at
|
||||
rollout (rl_action_kernel uses `thompson_sample_categorical`), argmax
|
||||
only for Bellman target (uses `argmax_expected_q`).
|
||||
|
||||
### A6. GPU-resident advantage / return / TD computation
|
||||
Rules honored: `feedback_cpu_is_read_only`, `feedback_no_atomicadd`,
|
||||
`pearl_thompson_for_distributional_action_selection`,
|
||||
`pearl_no_host_branches_in_captured_graph`.
|
||||
|
||||
New `compute_advantage_return` kernel: per-batch `r + γ(1-done)·V_tp1`
|
||||
and `r_t - V_t`, reading γ from `ISV[400]`. Replaces the host loop in
|
||||
the flawed branch's `step_with_lobsim`.
|
||||
### A6 — GPU-resident advantage / return / TD computation
|
||||
|
||||
### A7. Fee model
|
||||
New `compute_advantage_return` kernel: per-batch
|
||||
`returns[b] = r_b + γ(1 − done_b) · v_tp1[b]`,
|
||||
`advantages[b] = returns[b] − v_t[b]`. Reads γ from `ISV[400]`,
|
||||
v_t / v_tp1 from device buffers. Element-wise; trivially parallel.
|
||||
|
||||
The flawed F.3a's approach (`order_match.cu` deduct fee + new
|
||||
`upload_cost_per_lot_per_side` API) was correct in principle, mirroring
|
||||
the existing `apply_fill_to_pos` pattern from `resting_orders.cu`. The
|
||||
rebuild keeps this change.
|
||||
Replaces the host advantage/return loop in the flawed branch's
|
||||
`step_with_lobsim`.
|
||||
|
||||
### A8. Loader pair API
|
||||
Rules honored: `feedback_cpu_is_read_only`,
|
||||
`pearl_controller_anchors_isv_driven`.
|
||||
|
||||
The flawed Phase G's `MultiHorizonLoader::next_sequence_pair` was
|
||||
correct (true `(s_t, s_{t+1})` consecutive sampling). The rebuild keeps
|
||||
this addition.
|
||||
### A7 — Fee model: port F.3a unchanged
|
||||
|
||||
## Falsifiability gates — pre-cluster
|
||||
`order_match.cu::submit_market_immediate` adds per-fill cost
|
||||
deduction mirroring `apply_fill_to_pos` from `resting_orders.cu:213-219`.
|
||||
`upload_cost_per_lot_per_side` API on `LobSimCuda`. Default zero,
|
||||
opt-in (CLI defaults to ES $1.25/side, $2.50 round-trip).
|
||||
|
||||
Each gate is a local CUDA test that exercises the failure mode it
|
||||
covers. None of these existed in the flawed branch.
|
||||
Single-writer-per-block plain `+=` for `total_fees_per_b[b]` (line 79
|
||||
of `submit_market_immediate` has `threadIdx.x != 0 → return`).
|
||||
|
||||
| Gate | Asserts | Catches defect |
|
||||
|------|---------|----------------|
|
||||
| G1 | ISV[400..406] held at canonical defaults after `IntegratedTrainer::new` | #1 |
|
||||
| G2 | `RL_REWARD_SCALE_INDEX` stays at 1.0 for first N hold-only steps | #2 |
|
||||
| G3 | Each controller's launcher visibly mutates its ISV slot after first non-sentinel input | #3 |
|
||||
| G4 | `target_w` Frobenius norm drifts from `w` only after the soft-update kernel fires | #4 |
|
||||
| G5 | `cudaMemcpy` count per `step_with_lobsim` is 0 host-initiated (other than ISV diagnostic readback) | #5 |
|
||||
| G6 | Bull-market forced-long 50-trade loop: mean reward per trade > 0 (already in flawed F.3b's `reward_calibration.rs`) | (regression-coverage for fee + scale) |
|
||||
| G7 | Run for N steps; abort + non-zero exit code if any per-head loss is non-finite | #7 |
|
||||
Rules honored: `feedback_no_atomicadd` (single-writer),
|
||||
`feedback_single_source_of_truth_no_duplicates` (mirrors existing
|
||||
`apply_fill_to_pos` pattern).
|
||||
|
||||
G5 is the load-bearing gate for the GPU-pure decision (A1). Tooling:
|
||||
hook `cudaMemcpyAsync` via a CUPTI counter or count Rust-side
|
||||
`memcpy_htod` / `memcpy_dtoh` calls during the step.
|
||||
### A8 — Loader pair API: port G fix unchanged
|
||||
|
||||
`MultiHorizonLoader::next_sequence_pair` returns
|
||||
`(LabeledSequence, LabeledSequence)` at adjacent anchors in the same
|
||||
file. Refactor `next_sequence`'s body into `build_sequence_at(file_idx,
|
||||
anchor)` helper shared by both API surfaces.
|
||||
|
||||
Rules honored: `feedback_no_partial_refactor` (helper shared by both
|
||||
call sites; no duplicate windowing logic).
|
||||
|
||||
### A9 — Wire PER replay into the training step
|
||||
|
||||
`feedback_always_per` says PER is always on; the flawed branch had
|
||||
the `ReplayBuffer` struct but never sampled from it in
|
||||
`step_with_lobsim`. The rebuild WIRES the buffer:
|
||||
|
||||
- Every step pushes the current transition to the replay buffer.
|
||||
- Every step samples a batch from the buffer (priority^α) for the
|
||||
off-policy Q-head update.
|
||||
- TD-error magnitudes from the Q backward feed
|
||||
`update_priorities`.
|
||||
- `RL_PER_ALPHA_INDEX` (already in ISV[405], adaptive via
|
||||
`rl_per_alpha_controller`) drives the priority exponent.
|
||||
|
||||
PER buffer storage: device-resident `h_t`/`next_h_t` per
|
||||
`pearl_cold_path_no_exception_to_gpu_drives`. Sum-tree starts
|
||||
naive O(N) per `replay.rs` doc (acceptable at capacity ≤ 4096); GPU
|
||||
sum-tree is a Phase R-future enhancement.
|
||||
|
||||
Rules honored: `feedback_always_per`, `feedback_no_stubs` (replay
|
||||
buffer becomes load-bearing instead of orphan), `feedback_wire_everything_up`.
|
||||
|
||||
## Falsifiability gates — local CUDA BEFORE cluster
|
||||
|
||||
Each gate has a local CUDA test (`#[ignore = "requires CUDA"]`,
|
||||
runnable on the dev RTX 3050 Ti per CLAUDE.md) that exercises the
|
||||
failure mode the convergence-gate fixtures couldn't catch. Per
|
||||
`feedback_no_cpu_test_fallbacks`: every test oracle is either
|
||||
analytical synthetic, property-based, or cross-kernel GPU oracle.
|
||||
NEVER a CPU reference implementation.
|
||||
|
||||
| Gate | Asserts (invariant, not observed value) | Test design (GPU oracle only) | Catches defect |
|
||||
|------|----------------------------------------|-------------------------------|----------------|
|
||||
| G1 | ISV[400..406] equal canonical values after `IntegratedTrainer::new` | Construct trainer, read ISV slice via HEALTH_DIAG mapped-pinned, assert each slot equals its bootstrap value | #1 |
|
||||
| G2 | `RL_REWARD_SCALE_INDEX` stays in `[1e-3, 1e3]` and remains at 1.0 across N hold-only env steps (no closed trades) | Drive trainer for 100 steps with synthetic book where market never crosses (no fills); assert ISV[406] never moves from 1.0 | #2 |
|
||||
| G3 | Each controller's launcher visibly mutates its ISV slot after first non-sentinel EMA input | Construct trainer, manually launch one controller with a non-zero scalar input via the EMA-update kernel, assert ISV slot moved off bootstrap value | #3 |
|
||||
| G4 | `target_w` Frobenius distance from `w` is zero before `dqn_target_soft_update` ever fires, and >0 after one fire | Construct trainer, snapshot `target_w_d`, fire 1 step (which fires soft-update once), snapshot again, assert any element changed | #4 |
|
||||
| G5 | Zero host-initiated `cudaMemcpyAsync` calls in `step_with_lobsim`'s hot path (other than HEALTH_DIAG readback) | Hook `cudaMemcpyAsync` via CUPTI counter OR count Rust-side `memcpy_htod`/`memcpy_dtoh` calls; assert count = expected ISV-readback count | #5 |
|
||||
| G6 | PER buffer's `len()` grows by exactly 1 per `step_with_lobsim` call; `sample_indices` returns batch of correct size | Drive trainer for 100 steps, assert `replay.len() == min(100, capacity)`; sample, assert batch size | #6 |
|
||||
| G7 | Bull-market 50-cycle forced-long produces mean per-trade reward > 0 on synthetic `LobSimCuda` with ES-realistic fees | Port the flawed branch's `reward_calibration.rs` to drive `LobSimCuda` directly (no `LobEnv` adapter) | (regression coverage for fee + scale wiring) |
|
||||
| G8 | NaN abort: trainer exits with non-zero code if any per-head loss is non-finite | Inject NaN into a synthetic reward stream, assert process exits 2 within 1 step | #8 |
|
||||
|
||||
G5 is the load-bearing gate for the GPU-pure decision (A1).
|
||||
Implementation: a thin Rust wrapper around `CudaStream::memcpy_*`
|
||||
that increments a counter; tests assert the counter after one
|
||||
`step_with_lobsim` call. CUPTI is a follow-up enhancement.
|
||||
|
||||
## Cluster smoke discipline (Phase H equivalent)
|
||||
|
||||
Per `pearl_single_window_oos_is_not_oos`: the G8 backtest gate
|
||||
(`profit_factor > 1.0`) MUST be evaluated across ≥3 walk-forward
|
||||
folds, reporting mean ± SD. A single-window pass is NOT G8 pass.
|
||||
|
||||
Per `feedback_kill_runs_on_anomaly_quickly`: the cluster smoke
|
||||
script tails logs and kills the run if any of the following fire:
|
||||
- Any per-head loss non-finite (G8)
|
||||
- `RL_REWARD_SCALE_INDEX` clipped at `1e-3` or `1e3` (instrumentation
|
||||
saturation indicator)
|
||||
- Per-step wall time > 10× the rolling median (kernel hang indicator)
|
||||
|
||||
DO NOT kill on Q-distribution spread (per
|
||||
`pearl_q_spread_kill_criterion_misaligned_for_distributional_q`:
|
||||
Q_SPREAD spikes signal GOOD balance under C51+Thompson).
|
||||
|
||||
Per `feedback_push_before_deploy`: dispatcher script verifies
|
||||
`git push` succeeded before `argo submit`. Hard-error if local HEAD
|
||||
differs from `origin/<branch>` after push.
|
||||
|
||||
Per `feedback_argo_template_must_apply`: dispatcher prints a
|
||||
top-of-output reminder `kubectl apply -n foxhunt -f
|
||||
infra/k8s/argo/alpha-rl-template.yaml` and refuses to submit if the
|
||||
template has changed since last apply (compare git SHA of the
|
||||
template file vs the cluster CRD's `app.kubernetes.io/version` label
|
||||
— or just hard-print the reminder).
|
||||
|
||||
Per `feedback_default_to_l40s_pool`: dispatcher defaults to
|
||||
`--gpu-pool ci-training-l40s` (NOT H100). H100 pool is opt-in for
|
||||
production scale-up, not smoke. Per `feedback_h100_gpu`: when H100
|
||||
IS chosen, max-utilization config; hard-error on under-utilization.
|
||||
|
||||
## Required new files
|
||||
|
||||
| File | Purpose |
|
||||
|------|---------|
|
||||
| `crates/ml-alpha/cuda/thompson_sample_categorical.cu` | A5 |
|
||||
| `crates/ml-alpha/cuda/argmax_expected_q.cu` | A5 |
|
||||
| `crates/ml-alpha/cuda/log_pi_at_action.cu` | A5 |
|
||||
| `crates/ml-alpha/cuda/compute_advantage_return.cu` | A6 |
|
||||
| `crates/ml-alpha/cuda/dqn_target_soft_update.cu` | A4 |
|
||||
| `crates/ml-alpha/cuda/ema_update_on_done.cu` | A2 — generic done-gated EMA (parameterised slot) |
|
||||
| `crates/ml-alpha/cuda/ema_update_per_step.cu` | A2 — generic per-step EMA (parameterised slot) |
|
||||
### Kernels
|
||||
- `crates/ml-alpha/cuda/rl_action_kernel.cu` — A5 Thompson sampler writing market_targets
|
||||
- `crates/ml-alpha/cuda/argmax_expected_q.cu` — A5
|
||||
- `crates/ml-alpha/cuda/log_pi_at_action.cu` — A5
|
||||
- `crates/ml-alpha/cuda/compute_advantage_return.cu` — A6
|
||||
- `crates/ml-alpha/cuda/dqn_target_soft_update.cu` — A4
|
||||
- `crates/ml-alpha/cuda/ema_update_on_done.cu` — A2 (generic done-gated EMA, parameterised by slot)
|
||||
- `crates/ml-alpha/cuda/ema_update_per_step.cu` — A2 (generic per-step EMA, parameterised by slot)
|
||||
|
||||
Reused / extended:
|
||||
- `rl_reward_scale_controller.cu` (port from flawed branch)
|
||||
### Ported from flawed branch
|
||||
- `crates/ml-alpha/cuda/rl_reward_scale_controller.cu`
|
||||
- `crates/ml-backtesting/cuda/order_match.cu` modifications (A7)
|
||||
|
||||
### Already in baseline, need Rust launchers added (R5)
|
||||
- `rl_gamma_controller.cu`, `rl_target_tau_controller.cu`,
|
||||
`rl_ppo_clip_controller.cu`, `rl_entropy_coef_controller.cu`,
|
||||
`rl_rollout_steps_controller.cu`, `rl_per_alpha_controller.cu`
|
||||
(already in baseline tree — port their Rust launchers from the
|
||||
flawed branch concept, but feed device-resident EMA inputs)
|
||||
- `crates/ml-backtesting/cuda/order_match.cu` (port the F.3a fee
|
||||
deduction)
|
||||
|
||||
Deleted:
|
||||
- `crates/ml-alpha/src/rl/reward.rs::LobEnv` trait
|
||||
- `crates/ml-backtesting/src/sim/mod.rs::LobSimEnvAdapter`
|
||||
### Rust modules
|
||||
- Extensions to `crates/ml-alpha/src/trainer/integrated.rs` (rebuild
|
||||
`step_with_lobsim` as GPU-pure, add launchers)
|
||||
- Extensions to `crates/ml-alpha/src/rl/isv_slots.rs` (7 new EMA slots)
|
||||
- New `crates/ml-alpha/src/rl/per_step.rs` (wires PER buffer into the step)
|
||||
- Extensions to `crates/ml-alpha/src/data/loader.rs` (A8 port)
|
||||
- Extensions to `crates/ml-backtesting/src/sim/mod.rs` (A7 port)
|
||||
|
||||
### Deleted
|
||||
- `crates/ml-alpha/src/rl/reward.rs::LobEnv` trait + `MockLobEnv`
|
||||
- (Flawed branch's `LobSimEnvAdapter` was on `ml-backtesting`; not in
|
||||
baseline since baseline reset removed it)
|
||||
|
||||
### Renamed (per the "no toys" feedback)
|
||||
- `crates/ml-alpha/tests/dqn_toy.rs` → `dqn_reward_signal.rs`
|
||||
(rewritten to drive synthetic-book `LobSimCuda` directly)
|
||||
- `crates/ml-alpha/tests/ppo_toy.rs` → `ppo_reward_signal.rs`
|
||||
(same)
|
||||
- All "toy bandit" / "bandit" comments in production source rewritten
|
||||
to "reward-signal fixture" or equivalent
|
||||
|
||||
## Phased execution
|
||||
|
||||
Each phase = one commit. Each phase has a green-build + green-gate
|
||||
requirement before the next phase starts.
|
||||
Each R-phase = one atomic commit. Each phase has a green-build +
|
||||
green local-CUDA-gate requirement before the next phase starts. No
|
||||
phase is allowed to defer work to a later phase (per
|
||||
`feedback_no_partial_refactor` and the lesson from the flawed
|
||||
branch).
|
||||
|
||||
### R1 — ISV bootstrap + slot extension (1 commit)
|
||||
### R1 — ISV slot extension + controller bootstrap launches (1 commit)
|
||||
- Add 7 new ISV slot constants (`RL_*_EMA_INDEX`) → `RL_SLOTS_END = 424`.
|
||||
- In `IntegratedTrainer::new`: alloc_zeros ISV, then launch each of
|
||||
the 7 controllers ONCE with sentinel-zero EMA input. Each kernel's
|
||||
first-observation-bootstrap path fires → ISV slot gets canonical
|
||||
value (A3).
|
||||
- **Gate G1** local CUDA test.
|
||||
|
||||
- Add new ISV slot constants for the 7 EMA inputs (A2) → `RL_SLOTS_END = 424`.
|
||||
- In `IntegratedTrainer::new`, one memcpy_htod writing canonical
|
||||
defaults to ISV[400..406] (A3).
|
||||
- Gate G1.
|
||||
### R2 — Fee model + loader pair API (1 commit, was 2)
|
||||
- Port `order_match.cu` fee deduction + `upload_cost_per_lot_per_side`
|
||||
host API (A7).
|
||||
- Port `MultiHorizonLoader::next_sequence_pair` + `build_sequence_at`
|
||||
refactor (A8).
|
||||
- Existing tests pass; no new tests in this phase (covered by G7 later).
|
||||
|
||||
### R2 — Fee model + loader pair API (2 commits)
|
||||
### R3 — GPU-resident EMA + advantage/return kernels (1 commit)
|
||||
- `ema_update_on_done.cu` + `ema_update_per_step.cu` generic kernels
|
||||
(A2).
|
||||
- `compute_advantage_return.cu` (A6).
|
||||
- Local CUDA tests for each kernel via GPU-oracle pattern:
|
||||
- EMA: feed constant input, assert ISV converges to input (property:
|
||||
EMA of constant = constant).
|
||||
- EMA bootstrap: feed sentinel-zero start + one observation, assert
|
||||
slot equals observation (per `pearl_first_observation_bootstrap`).
|
||||
- Advantage/return: feed all-zero rewards + done=0 + v_t=v_tp1=k,
|
||||
assert returns = γ·k and advantages = γ·k − k = (γ−1)·k (property).
|
||||
|
||||
- R2a: port `order_match.cu` fee deduction + `upload_cost_per_lot_per_side`
|
||||
host API (from F.3a). Gate: existing `lob_env_adapter_buy_close_cycle`
|
||||
passes with zero fees; new with-fees variant lands.
|
||||
- R2b: port `MultiHorizonLoader::next_sequence_pair` (from G fix).
|
||||
|
||||
### R3 — GPU-pure action sampling (3 commits)
|
||||
|
||||
- R3a: `thompson_sample_categorical` kernel + Rust launcher. Local CUDA
|
||||
test: sample 10000 actions from a known Q distribution, assert
|
||||
empirical frequency matches softmax probabilities within tolerance.
|
||||
- R3b: `argmax_expected_q` kernel. Local CUDA test: feed adversarial
|
||||
Q distributions, assert argmax matches host reference.
|
||||
- R3c: `log_pi_at_action` kernel. Local CUDA test: feed known π
|
||||
distribution, assert log-prob within 1e-5 of host softmax.
|
||||
|
||||
### R4 — GPU-resident EMA + advantage / return (2 commits)
|
||||
|
||||
- R4a: `ema_update_on_done.cu` + `ema_update_per_step.cu` generic
|
||||
kernels. Local CUDA test: feed deterministic reward stream, assert
|
||||
EMA tracks the analytic Wiener-α trajectory.
|
||||
- R4b: `compute_advantage_return.cu`. Local CUDA test: feed known
|
||||
rewards / dones / V values, assert advantage / return match host
|
||||
reference.
|
||||
### R4 — GPU-resident action sampling kernels (1 commit)
|
||||
- `rl_action_kernel.cu` (Thompson) + `argmax_expected_q.cu` +
|
||||
`log_pi_at_action.cu` + curand PRNG state allocation.
|
||||
- Local CUDA tests via GPU-oracle:
|
||||
- `thompson_sample_categorical`: sample 10000 actions from a Q
|
||||
distribution where atom probability mass is concentrated on a
|
||||
known atom; assert empirical action frequency matches the
|
||||
softmax-weighted expected value within tolerance. Cross-kernel
|
||||
oracle: compare against `argmax_expected_q` on the same input
|
||||
when the distribution is sharp enough that Thompson ≈ argmax.
|
||||
- `argmax_expected_q`: feed Q distribution where one action has
|
||||
strictly higher expected value than all others by a known margin;
|
||||
assert that action is chosen for every batch.
|
||||
- `log_pi_at_action`: feed π logits where one action has logit
|
||||
massively above others; assert log π(that action) ≈ 0
|
||||
(probability ≈ 1).
|
||||
|
||||
### R5 — Wire all 7 controllers + target-net soft update (1 commit)
|
||||
- Launcher fields + helpers for all 7 controllers (gamma, target_tau,
|
||||
ppo_clip, entropy_coef, rollout_steps, per_alpha, reward_scale).
|
||||
Each pulls input from the corresponding EMA slot (R3a outputs).
|
||||
- `dqn_target_soft_update.cu` + launcher (A4).
|
||||
- **Gate G3** local CUDA test (per-controller).
|
||||
- **Gate G4** local CUDA test (soft-update fires + moves target_w).
|
||||
|
||||
- Add launcher fields + helpers for all 7 controllers (gamma,
|
||||
target_tau, ppo_clip, entropy_coef, rollout_steps, per_alpha,
|
||||
reward_scale). Each launcher pulls its input from the corresponding
|
||||
EMA slot (A2 from R4a). Gate G3.
|
||||
- Add `dqn_target_soft_update` kernel + launcher (A4). Gate G4.
|
||||
### R6 — Replace `LobEnv` with kernel-driven decision path (2 commits)
|
||||
- R6a: delete `LobEnv` trait + `MockLobEnv` from
|
||||
`src/rl/reward.rs`. Direct `LobSimCuda` integration on
|
||||
`IntegratedTrainer` (no trait abstraction). Wire action submission
|
||||
via the existing decision-policy kernel pattern; `market_targets_d`
|
||||
receives the Thompson-sampled index from `rl_action_kernel`.
|
||||
- R6b: rewrite convergence-gate fixtures (`dqn_reward_signal.rs`,
|
||||
`ppo_reward_signal.rs`, `integrated_trainer_smoke.rs`) to drive a
|
||||
real `LobSimCuda` with synthetic book data instead of MockLobEnv.
|
||||
Reward signal: "synthetic monotonic uptick → forced long action →
|
||||
trade close pays > 0". Per `pearl_tests_must_prove_not_lock_observations`,
|
||||
asserts invariants (reward signal flows, sign correctness), not
|
||||
specific Q-value magnitudes.
|
||||
- **Gate G5** local CUDA test (zero host memcpy in step_with_lobsim
|
||||
hot path).
|
||||
- **Gate G7** local CUDA test (bull-market mean per-trade reward > 0).
|
||||
|
||||
### R6 — Replace `LobEnv` with kernel-driven decision path (3 commits)
|
||||
|
||||
- R6a: delete `LobEnv` trait + `LobSimEnvAdapter`. Add direct
|
||||
`LobSimCuda` integration on `IntegratedTrainer` (no trait).
|
||||
- R6b: wire action submission via the decision-policy kernel pattern
|
||||
(mirror `step_decision_with_latency`). Adapt `market_targets_d` to
|
||||
accept the Thompson-sampled action index directly.
|
||||
- R6c: rewrite the trainer's convergence-gate fixtures (formerly
|
||||
`dqn_toy.rs` / `ppo_toy.rs`) to drive a real `LobSimCuda` with
|
||||
synthetic book data instead of a `MockLobEnv`. Reward signal becomes
|
||||
"synthetic monotonic uptick → forced long → trade close pays > 0".
|
||||
- Gate G5 (no host memcpy in step_with_lobsim other than ISV
|
||||
diagnostics).
|
||||
|
||||
### R7 — Reward composition (1 commit)
|
||||
|
||||
- Per-step reward = `realized_pnl_delta` (already net of fees + slippage
|
||||
from R2a) scaled by `ISV[RL_REWARD_SCALE_INDEX]`. Optionally subtract
|
||||
exposure penalty (deferred unless smoke shows it's needed).
|
||||
- Gate G6.
|
||||
### R7 — PER wiring (1 commit)
|
||||
- Push every step's transition to the replay buffer.
|
||||
- Sample `batch_size` indices per step for the off-policy Q update.
|
||||
- Update priorities from per-sample TD-error magnitudes after Q
|
||||
backward.
|
||||
- **Gate G6** local CUDA test (buffer grows, sample returns correct
|
||||
size).
|
||||
|
||||
### R8 — CLI binary + Argo template + dispatcher (1 commit)
|
||||
|
||||
- `crates/ml-alpha/examples/alpha_rl_train.rs` driving the rebuilt
|
||||
trainer. Uses `next_sequence_pair` for true `(s_t, s_{t+1})` windows.
|
||||
- Mandatory CLI args: `--mbp10-data-dir`. Plan-time reconciliation
|
||||
with `feedback_mbp10_mandatory`: ml-alpha's `MultiHorizonLoader`
|
||||
does not currently consume `--trades-data-dir` (the OFI features
|
||||
in `Mbp10RawInput` are derived from MBP-10 snapshot deltas, not
|
||||
from a separate trades stream). If the project chooses to wire a
|
||||
trades stream into ml-alpha's feature path, that's a separate
|
||||
plan; this rebuild ships with mbp10-only and flags the gap.
|
||||
- NaN abort: per-step `is_finite()` check on `l_total`; exit code 2
|
||||
on divergence. (**Gate G8**)
|
||||
- `infra/k8s/argo/alpha-rl-template.yaml` + `scripts/argo-alpha-rl.sh`
|
||||
(port from flawed branch).
|
||||
- NaN abort: per-step `is_finite()` check on `l_total`, exit code 2 on
|
||||
divergence. Gate G7.
|
||||
- Per `feedback_default_to_l40s_pool`: defaults to
|
||||
`ci-training-l40s`, not H100.
|
||||
- Per `feedback_argo_template_must_apply`: top-of-output reminder
|
||||
+ refuse-to-submit if template not applied since last edit.
|
||||
- Per `feedback_push_before_deploy`: verify `git push` succeeded
|
||||
before `argo submit`; hard-error if local HEAD ≠ origin HEAD.
|
||||
|
||||
### R9 — Pre-cluster smoke + first cluster smoke (out-of-band)
|
||||
### R9 — Pre-cluster validation + first cluster smoke (out-of-band)
|
||||
- Run all R1..R8 local CUDA tests on a GPU host (the dev RTX 3050 Ti
|
||||
per CLAUDE.md or any L40S). Confirm green: G1, G2, G3, G4, G5, G6,
|
||||
G7, G8.
|
||||
- `git push`. `kubectl apply -n foxhunt -f
|
||||
infra/k8s/argo/alpha-rl-template.yaml`.
|
||||
- `./scripts/argo-alpha-rl.sh --n-steps 1000 --instrument-mode
|
||||
front-month` (5-min validation smoke). If green: scale to full 50k.
|
||||
- For the G8 backtest gate: per `pearl_single_window_oos_is_not_oos`,
|
||||
run ≥3 walk-forward lob-sweep folds, report mean ± SD. A single
|
||||
pass is NOT G8 pass.
|
||||
|
||||
- Run all R1..R8 local CUDA tests on a GPU host. Confirm green.
|
||||
- `git push`. `kubectl apply -f` template. `./scripts/argo-alpha-rl.sh
|
||||
--n-steps 1000` for a 5-min validation smoke. If green, full 50k.
|
||||
## Pre-cluster validation checklist (R9 prerequisite)
|
||||
|
||||
Per `feedback_smoke_validation_before_structural_priors`, validate the
|
||||
wiring with neutral config BEFORE shipping production scale. The
|
||||
following MUST be green on a local GPU host (RTX 3050 Ti is
|
||||
sufficient for everything except wall-time benchmarks):
|
||||
|
||||
```bash
|
||||
# Build clean
|
||||
SQLX_OFFLINE=true cargo build --release --workspace 2>&1 | tail -10
|
||||
SQLX_OFFLINE=true cargo build --release -p ml-alpha --example alpha_rl_train
|
||||
|
||||
# All local-CUDA tests green (R1..R8 gates)
|
||||
SQLX_OFFLINE=true cargo test -p ml-alpha --test isv_bootstrap -- --ignored --nocapture # G1
|
||||
SQLX_OFFLINE=true cargo test -p ml-alpha --test reward_scale_hold_steady -- --ignored --nocapture # G2
|
||||
SQLX_OFFLINE=true cargo test -p ml-alpha --test controllers_emit -- --ignored --nocapture # G3
|
||||
SQLX_OFFLINE=true cargo test -p ml-alpha --test target_soft_update -- --ignored --nocapture # G4
|
||||
SQLX_OFFLINE=true cargo test -p ml-alpha --test no_host_memcpy -- --ignored --nocapture # G5
|
||||
SQLX_OFFLINE=true cargo test -p ml-alpha --test per_wired -- --ignored --nocapture # G6
|
||||
SQLX_OFFLINE=true cargo test -p ml-backtesting --test reward_calibration -- --ignored --nocapture # G7
|
||||
SQLX_OFFLINE=true cargo test -p ml-alpha --test nan_abort -- --ignored --nocapture # G8
|
||||
|
||||
# Convergence-gate fixtures green
|
||||
SQLX_OFFLINE=true cargo test -p ml-alpha --test dqn_reward_signal -- --ignored --nocapture
|
||||
SQLX_OFFLINE=true cargo test -p ml-alpha --test ppo_reward_signal -- --ignored --nocapture
|
||||
SQLX_OFFLINE=true cargo test -p ml-alpha --test integrated_trainer_smoke -- --ignored --nocapture
|
||||
```
|
||||
|
||||
Only after ALL of these pass green: `git push` → `kubectl apply` →
|
||||
`argo submit --n-steps 1000`.
|
||||
|
||||
Only after the 1000-step smoke completes Succeeded with finite losses
|
||||
+ ISV slots in healthy ranges + > 0 trades closed: scale to 50k.
|
||||
|
||||
Only after the 50k run produces a checkpoint that achieves G6
|
||||
calibration + survives 3 walk-forward lob-sweep folds: claim Phase H
|
||||
G8 pass.
|
||||
|
||||
## What we keep from the flawed branch
|
||||
|
||||
These specific elements landed correctly and are ported directly:
|
||||
- `rl_reward_scale_controller.cu` kernel (Phase F.2)
|
||||
- `order_match.cu` fee deduction pattern (Phase F.3a)
|
||||
- `MultiHorizonLoader::next_sequence_pair` (Phase G fix)
|
||||
- The 6 unlaunched controller .cu files (Phases C, D, E) — they're
|
||||
fine kernels, just need launchers
|
||||
- F.4's true-h_{t+1} architectural insight (Bellman target on
|
||||
h_{t+1}, V bootstrap on h_{t+1}) — the structure becomes simpler in
|
||||
the GPU-pure rebuild because state lives on device
|
||||
Direct ports:
|
||||
- `rl_reward_scale_controller.cu` kernel
|
||||
- `order_match.cu` fee deduction + `upload_cost_per_lot_per_side` API
|
||||
- `MultiHorizonLoader::next_sequence_pair` + `build_sequence_at`
|
||||
refactor
|
||||
|
||||
Architectural insights that survive the rebuild:
|
||||
- True h_{t+1} Bellman target (F.4 close-out idea) — in the rebuild
|
||||
this falls out for free because action sampling is GPU-resident, so
|
||||
the next-state Q forward is just another kernel call with the next
|
||||
snapshot, no host roundtrip needed.
|
||||
- V-bootstrap from V(s_{t+1}) — same.
|
||||
- Toy rename (`dqn_toy.rs` → `dqn_reward_signal.rs`, etc.) — applied
|
||||
in R6b on top of the rewritten fixtures.
|
||||
|
||||
## What we discard
|
||||
|
||||
- `LobEnv` trait + `LobSimEnvAdapter` (replaced by kernel-driven path)
|
||||
- `LobEnv` trait + `LobSimEnvAdapter` (kernel-driven path replaces
|
||||
the trait abstraction entirely)
|
||||
- `MockLobEnv` "toy bandit" fixture (replaced by synthetic-book
|
||||
`LobSimCuda` fixture)
|
||||
- Host-side action sampling, EMA tracking, advantage / return loops
|
||||
- The `dqn_toy.rs` / `ppo_toy.rs` test file names (renamed to
|
||||
`dqn_reward_signal.rs` / `ppo_reward_signal.rs` in the rebuild)
|
||||
`LobSimCuda` fixture in R6b)
|
||||
- Host-side action sampling, EMA tracking, advantage/return loops
|
||||
- Host `mean_abs_pnl_ema: f32` field on `IntegratedTrainer` (moved
|
||||
to ISV slot, updated by `ema_update_on_done` kernel)
|
||||
- Single-window OOS as a G8 pass criterion (replaced by ≥3
|
||||
walk-forward folds per `pearl_single_window_oos_is_not_oos`)
|
||||
|
||||
## Pearls honored
|
||||
## Memory pearls honored
|
||||
|
||||
- `feedback_cpu_is_read_only` — A1, A5, A6 are direct compliance.
|
||||
- `feedback_isv_for_adaptive_bounds` — A2, A3, R5.
|
||||
- `feedback_no_partial_refactor` — each R-phase is atomic.
|
||||
- `feedback_no_stubs` — every controller has a launcher AND a
|
||||
consumer at landing time.
|
||||
- `feedback_no_atomicadd` — new kernels use single-writer-per-block
|
||||
or per-batch scratch + reduce_axis0.
|
||||
- `feedback_no_htod_htoh_only_mapped_pinned` — applies to the
|
||||
remaining ISV diagnostic readback.
|
||||
- `pearl_first_observation_bootstrap` — combined with A3, the host
|
||||
bootstrap takes precedence; kernel sentinel-zero path stays as a
|
||||
safety net.
|
||||
- `pearl_wiener_alpha_floor_for_nonstationary` — every EMA uses
|
||||
α-floor = 0.4.
|
||||
- `feedback_no_hiding` — `step_event`'s `_ts_ns` / `_trade_signed_vol`
|
||||
go away with the `LobEnv` trait deletion.
|
||||
- `feedback_no_feature_flags` — controllers are unconditional.
|
||||
- `feedback_registry_entries_need_dispatch_arms` — controllers go in
|
||||
the state-reset registry (if applicable to ml-alpha).
|
||||
- `pearl_tests_must_prove_not_lock_observations` — convergence-gate
|
||||
fixtures assert invariants (reward signal flows, sign correctness),
|
||||
not specific learned-checkpoint values.
|
||||
Catalog of pearls the rebuild explicitly applies (line numbers and
|
||||
specific implementation pointers will be filled in at R-phase commit
|
||||
time per `feedback_no_partial_refactor`):
|
||||
|
||||
- `feedback_cpu_is_read_only` → A1, A5, A6, G5
|
||||
- `feedback_isv_for_adaptive_bounds` → A2, A3, R5
|
||||
- `feedback_wire_everything_up` → A3, A4, A9, R1, R5, R7
|
||||
- `feedback_no_partial_refactor` → every R-phase atomic; controllers
|
||||
AND their consumers land together
|
||||
- `feedback_no_stubs` → A4 (target soft update), A9 (PER wired)
|
||||
- `feedback_no_atomicadd` → all new kernels single-writer-per-block
|
||||
or `reduce_axis0`
|
||||
- `feedback_no_cpu_test_fallbacks` → every gate's test uses GPU
|
||||
oracles only (analytical synthetic, property assertion, or
|
||||
cross-kernel)
|
||||
- `feedback_no_htod_htoh_only_mapped_pinned` → A3 (no host memcpy
|
||||
bootstrap), mapped-pinned for any HEALTH_DIAG readback
|
||||
- `feedback_always_per` → A9
|
||||
- `feedback_mbp10_mandatory` → R8 (CLI requires `--mbp10-data-dir`;
|
||||
reconciles trades-data-dir gap explicitly)
|
||||
- `feedback_default_to_l40s_pool` → R8 dispatcher
|
||||
- `feedback_push_before_deploy` → R8 dispatcher
|
||||
- `feedback_argo_template_must_apply` → R8 dispatcher
|
||||
- `feedback_kill_runs_on_anomaly_quickly` → cluster smoke discipline
|
||||
- `feedback_extending_existing_code_audits_for_existing_violations`
|
||||
→ the lesson driving this entire rebuild
|
||||
- `pearl_first_observation_bootstrap` → A3, R3 EMA bootstrap tests
|
||||
- `pearl_wiener_alpha_floor_for_nonstationary` → all EMA controllers
|
||||
- `pearl_wiener_optimal_adaptive_alpha` → A2 EMA design
|
||||
- `pearl_thompson_for_distributional_action_selection` → A5 split
|
||||
(Thompson at rollout, argmax for Bellman)
|
||||
- `pearl_action_pruning_falsified` → 9-action grid preserved
|
||||
- `pearl_c51_thompson_closed_phase_e3_gap` → C51 + Thompson combo
|
||||
preserved
|
||||
- `pearl_single_window_oos_is_not_oos` → cluster smoke discipline
|
||||
(≥3 walk-forward folds for G8)
|
||||
- `pearl_q_spread_kill_criterion_misaligned_for_distributional_q` →
|
||||
cluster smoke kill criteria explicitly exclude Q_SPREAD
|
||||
- `pearl_tests_must_prove_not_lock_observations` → every gate
|
||||
asserts invariants, not specific learned-checkpoint values
|
||||
- `pearl_cold_path_no_exception_to_gpu_drives` → A1, A2, A5, A6 are
|
||||
all GPU-resident even for "cold path" diagnostic computations
|
||||
- `pearl_no_host_branches_in_captured_graph` → A5 PRNG state lives
|
||||
on device; no host RNG advance inside any captured kernel chain
|
||||
- `pearl_one_unbounded_signal_per_reward` → reward = scale × pnl_delta
|
||||
has scale bounded `[1e-3, 1e3]` (R5 reward_scale controller) and
|
||||
pnl_delta as the single unbounded multiplicand
|
||||
- `pearl_symmetric_clamp_audit` → reward_scale clamp is bilateral
|
||||
`fmaxf(REWARD_SCALE_MIN, fminf(x, REWARD_SCALE_MAX))`
|
||||
- `pearl_event_driven_reward_density_alignment` → per-trade reward
|
||||
fires on `done=true` (event density = trade close density)
|
||||
|
||||
## Calendar estimate
|
||||
|
||||
- R1: 0.5 days
|
||||
- R2: 1 day
|
||||
- R3: 1.5 days
|
||||
- R4: 1 day
|
||||
- R5: 1 day
|
||||
- R6: 2 days
|
||||
- R7: 0.5 days
|
||||
- R8: 0.5 days
|
||||
- R9: 0.5 days + 1h cluster smoke
|
||||
Roughly 2× the flawed Phase F+G calendar because the rebuild does the
|
||||
work the flawed branch deferred or simplified.
|
||||
|
||||
Total ~8.5 dev days. Roughly 2× the flawed Phase F + G calendar
|
||||
because the rebuild does the work the flawed branch deferred or
|
||||
simplified.
|
||||
- R1: 0.5 days (controller bootstrap launches + ISV slot extension)
|
||||
- R2: 1 day (port F.3a + G fixes)
|
||||
- R3: 1 day (EMA + advantage/return kernels + GPU-oracle tests)
|
||||
- R4: 1.5 days (action sampling kernels + GPU-oracle tests)
|
||||
- R5: 1 day (controller wiring + soft update)
|
||||
- R6: 2 days (LobEnv removal + fixture rewrite)
|
||||
- R7: 0.5 days (PER wiring)
|
||||
- R8: 0.5 days (CLI + Argo)
|
||||
- R9: 0.5 days local validation + 5-min smoke + 1.5h × 3 folds = ~5h cluster
|
||||
|
||||
Total ~8.5 dev days + ~6h cluster smoke wall.
|
||||
|
||||
## Execution discipline
|
||||
|
||||
- Per `feedback_no_concurrent_agents_shared_tree`: R-phases execute
|
||||
serially in this branch (`ml-alpha-phase-a`). If parallel execution
|
||||
is needed, use git worktrees.
|
||||
- Per `feedback_fix_everything_aggressively`: no defer-to-later. Each
|
||||
R-phase ships its work complete.
|
||||
- Per `feedback_extending_existing_code_audits_for_existing_violations`:
|
||||
when each R-phase touches pre-existing code, the audit comes FIRST.
|
||||
If the existing code already violates a rule, the R-phase's scope
|
||||
expands to include the fix, not silently extends the violation.
|
||||
|
||||
End of plan.
|
||||
|
||||
Reference in New Issue
Block a user