diff --git a/Cargo.lock b/Cargo.lock index eeb6d3786..7fe0ee372 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -6094,6 +6094,7 @@ dependencies = [ "cudarc", "data", "memmap2", + "ml-backtesting", "ml-core", "ml-features", "rand 0.8.5", diff --git a/crates/ml-alpha/Cargo.toml b/crates/ml-alpha/Cargo.toml index a96e04195..22d7f7dab 100644 --- a/crates/ml-alpha/Cargo.toml +++ b/crates/ml-alpha/Cargo.toml @@ -53,3 +53,8 @@ rayon = { workspace = true } [dev-dependencies] tempfile = { workspace = true } approx = { workspace = true } +# Phase R6: convergence-gate fixtures + alpha_rl_train CLI need +# LobSimCuda from ml-backtesting. Dev-dep avoids the cycle — +# ml-backtesting → ml-alpha is the production direction; this +# dev-only edge only loads when building ml-alpha's own tests/examples. +ml-backtesting = { path = "../ml-backtesting" } diff --git a/crates/ml-alpha/build.rs b/crates/ml-alpha/build.rs index 5db69be7a..6114fc203 100644 --- a/crates/ml-alpha/build.rs +++ b/crates/ml-alpha/build.rs @@ -51,17 +51,21 @@ const KERNELS: &[&str] = &[ "argmax_expected_q", // RL Phase R4: argmax over expected Q per action; Bellman-target argmax (Double-DQN); pairs with rl_action_kernel per pearl_thompson_for_distributional_action_selection "log_pi_at_action", // RL Phase R4: per-batch log π(action_b) via log-softmax + lookup; PPO importance-ratio path "dqn_target_soft_update", // RL Phase R5: element-wise target[i] = (1-τ)·target + τ·current, reads τ from ISV[401]; closes defect #4 (no target-net soft update in flawed branch) + "extract_realized_pnl_delta", // RL Phase R6: GPU-pure reward + done extraction from device Pos array; replaces host read_pos loop per feedback_cpu_is_read_only + "apply_reward_scale", // RL Phase R6: element-wise rewards *= ISV[RL_REWARD_SCALE_INDEX=406]; closes the F.3b host roundtrip + "actions_to_market_targets", // RL Phase R6: 9-action grid → LobSim market_targets[B*2] on device; replaces host submit_market loop per feedback_cpu_is_read_only ]; -// Cache bust v28 (2026-05-23): RL Phase R5 rebuild — controller kernel -// signature change (scalar input arg → int input_slot reading ISV[417..423] -// directly per feedback_cpu_is_read_only, eliminating 7 DtoH-per-step -// roundtrips) plus the new dqn_target_soft_update kernel closing defect -// #4 (target network was never soft-updated in the flawed branch). -// launch_rl_controllers_per_step on IntegratedTrainer fires all 7 in -// sequence; DqnHead::soft_update_target launches dqn_target_soft_update -// twice per call (weight + bias). Both will be called from R6's -// step_with_lobsim once the EMA producers are wired into the step. +// Cache bust v29 (2026-05-23): RL Phase R6 rebuild — three new GPU-pure +// kernels close the last host-roundtrip violations in step_with_lobsim's +// hot path. extract_realized_pnl_delta reads pos.realized_pnl + position_lots +// from the device Pos array and writes per-batch (reward delta, done flag); +// apply_reward_scale multiplies rewards by ISV[406] on device; +// actions_to_market_targets translates the 9-action grid (with conditional +// flat-from-long/short reading current position_lots) into LobSim's +// market_targets format. Together with R3 EMA + R4 Thompson/argmax/log_pi + +// R5 controllers + soft-update, the trainer's step is GPU-pure: only HtoD +// for incoming snapshots (boundary data) + HEALTH_DIAG ISV readback. fn main() { println!("cargo:rerun-if-changed=build.rs"); diff --git a/crates/ml-alpha/cuda/actions_to_market_targets.cu b/crates/ml-alpha/cuda/actions_to_market_targets.cu new file mode 100644 index 000000000..5fc2608d6 --- /dev/null +++ b/crates/ml-alpha/cuda/actions_to_market_targets.cu @@ -0,0 +1,79 @@ +// actions_to_market_targets.cu — translate the 9-action discrete grid +// to LobSim's `market_targets[B*2] = {side, size}` device buffer +// (Phase R6 of the integrated RL trainer rebuild; +// see docs/superpowers/plans/2026-05-23-integrated-rl-trainer-rebuild.md). +// +// Replaces the host-side `submit_market(batch_idx, side, size, ts_ns)` +// loop the flawed Phase F shipped in LobSimEnvAdapter (which iterated +// the batch on CPU and called submit_market per batch — host +// orchestration of GPU work per `feedback_cpu_is_read_only`). +// +// Action mapping (must match `crate::rl::common::Action` and +// `LobSimCuda::submit_market_immediate`'s expected `{side, size}` pair): +// +// 0 ShortLarge side=1 (sell), size=2 +// 1 ShortSmall side=1 (sell), size=1 +// 2 Hold side=2 (no-op),size=0 +// 3 FlatFromLong side=1 (sell), size=position_lots (only if pos > 0) +// 4 FlatFromShort side=0 (buy), size=|position_lots| (only if pos < 0) +// 5 LongSmall side=0 (buy), size=1 +// 6 LongLarge side=0 (buy), size=2 +// 7 TrailTighten side=2 (no-op),size=0 — handled as ISV mutation, not a fill +// 8 TrailLoosen side=2 (no-op),size=0 — handled as ISV mutation, not a fill +// +// Actions 3/4 require reading the current `position_lots` from the +// Pos struct (`offset 0`, i32). Other actions are pure constant +// translations. +// +// `Pos` struct layout (from crates/ml-backtesting/src/lob/mod.rs::PosFlat): +// offset 0: position_lots: i32 +// (followed by other fields — see extract_realized_pnl_delta.cu header) +// +// Element-wise, one thread per batch entry. No atomics. + +extern "C" __global__ void actions_to_market_targets( + const int* __restrict__ actions, // [b_size] + const unsigned char* __restrict__ pos_state, // [b_size * pos_bytes] + int* __restrict__ market_targets, // [b_size * 2] OUT + int b_size, + int pos_bytes +) { + const int b = blockIdx.x * blockDim.x + threadIdx.x; + if (b >= b_size) return; + + const int action = actions[b]; + + int side = 2; // default no-op + int size = 0; + + if (action == 0) { + side = 1; size = 2; + } else if (action == 1) { + side = 1; size = 1; + } else if (action == 2) { + // Hold — no-op. + } else if (action == 3) { + const int position_lots = *(const int*)(pos_state + b * pos_bytes); + if (position_lots > 0) { + side = 1; size = position_lots; + } + // else no-op (no long to flatten) + } else if (action == 4) { + const int position_lots = *(const int*)(pos_state + b * pos_bytes); + if (position_lots < 0) { + side = 0; size = -position_lots; + } + // else no-op (no short to flatten) + } else if (action == 5) { + side = 0; size = 1; + } else if (action == 6) { + side = 0; size = 2; + } + // actions 7, 8 (TrailTighten/Loosen): no fill. These actions + // operate via ISV trail-distance mutation, not market-order + // submission — the side=2 no-op is correct here. The ISV + // mutation path lives outside this kernel. + + market_targets[b * 2 + 0] = side; + market_targets[b * 2 + 1] = size; +} diff --git a/crates/ml-alpha/cuda/apply_reward_scale.cu b/crates/ml-alpha/cuda/apply_reward_scale.cu new file mode 100644 index 000000000..0e1b90a01 --- /dev/null +++ b/crates/ml-alpha/cuda/apply_reward_scale.cu @@ -0,0 +1,30 @@ +// apply_reward_scale.cu — element-wise reward standardisation reading +// the scale from `ISV[RL_REWARD_SCALE_INDEX = 406]` (Phase R6 of the +// integrated RL trainer rebuild; +// see docs/superpowers/plans/2026-05-23-integrated-rl-trainer-rebuild.md). +// +// Replaces the host loop `for r in rewards: r *= scale` the flawed +// Phase F shipped in step_with_lobsim. That loop required DtoH of +// ISV[406] + host multiply + HtoD of scaled rewards — a triple +// roundtrip per training step. This kernel does the multiply on +// device with the scale already on device. +// +// Per `pearl_one_unbounded_signal_per_reward`: `rewards[b]` is the +// single unbounded multiplicand; `isv[RL_REWARD_SCALE_INDEX]` is +// bounded by the reward-scale controller (clamped to [1e-3, 1e3] +// per the kernel). Reward magnitude after scaling sits within the +// [Q_V_MIN, Q_V_MAX] = [-1, +1] atom support per the C51 atom layout. + +#define RL_REWARD_SCALE_INDEX 406 + +extern "C" __global__ void apply_reward_scale( + const float* __restrict__ isv, // ISV bus (≥ RL_REWARD_SCALE_INDEX + 1) + float* __restrict__ rewards, // [b_size] IN/OUT + int b_size +) { + const int b = blockIdx.x * blockDim.x + threadIdx.x; + if (b >= b_size) return; + + const float scale = isv[RL_REWARD_SCALE_INDEX]; + rewards[b] *= scale; +} diff --git a/crates/ml-alpha/cuda/extract_realized_pnl_delta.cu b/crates/ml-alpha/cuda/extract_realized_pnl_delta.cu new file mode 100644 index 000000000..abd1b0b44 --- /dev/null +++ b/crates/ml-alpha/cuda/extract_realized_pnl_delta.cu @@ -0,0 +1,59 @@ +// extract_realized_pnl_delta.cu — GPU-pure reward + done extraction +// (Phase R6 of the integrated RL trainer rebuild; +// see docs/superpowers/plans/2026-05-23-integrated-rl-trainer-rebuild.md). +// +// Replaces the host-side `read_pos` + scalar-subtract loop the flawed +// Phase F shipped in step_with_lobsim's LobSimEnvAdapter. That loop +// DtoH-copied the entire Pos struct per batch, computed the delta on +// CPU, then HtoD-uploaded actions back — a multi-direction roundtrip +// per training step per batch that violated `feedback_cpu_is_read_only`. +// +// This kernel reads the device-resident `Pos` array (the simulator's +// authoritative position state) and writes: +// - reward[b] = pos.realized_pnl − prev_realized_pnl[b] +// - done[b] = (prev_position_lots[b] != 0 && pos.position_lots == 0) ? 1 : 0 +// - prev_realized_pnl[b] := pos.realized_pnl (for next-step delta) +// - prev_position_lots[b] := pos.position_lots (for next-step done) +// +// `Pos` struct layout (from crates/ml-backtesting/src/lob/mod.rs::PosFlat, +// #[repr(C)], bytemuck::Pod): +// offset 0: position_lots: i32 (4 bytes) +// offset 4: vwap_entry: f32 (4 bytes) +// offset 8: realized_pnl: f32 (4 bytes) +// offset 12: peak_equity: f32 (4 bytes) +// offset 16: submission_overflow: u32 (4 bytes) +// offset 20: open_horizon_mask: u32 (4 bytes) +// total: 24 bytes per Pos +// +// `pos_bytes` is passed in as a kernel arg so this kernel doesn't have +// to be rebuilt if PosFlat ever grows fields. +// +// Element-wise, one thread per batch entry. No atomics, no reductions. +// Per `feedback_no_atomicadd` not needed. + +extern "C" __global__ void extract_realized_pnl_delta( + const unsigned char* __restrict__ pos_state, // [b_size * pos_bytes] + float* __restrict__ prev_realized_pnl, // [b_size] IN/OUT + int* __restrict__ prev_position_lots, // [b_size] IN/OUT + float* __restrict__ reward_out, // [b_size] OUT + float* __restrict__ done_out, // [b_size] OUT + int b_size, + int pos_bytes +) { + const int b = blockIdx.x * blockDim.x + threadIdx.x; + if (b >= b_size) return; + + const unsigned char* p = pos_state + b * pos_bytes; + const int current_lots = *(const int*)(p + 0); + const float current_realized = *(const float*)(p + 8); + + const float prev_realized = prev_realized_pnl[b]; + const int prev_lots = prev_position_lots[b]; + + reward_out[b] = current_realized - prev_realized; + done_out[b] = (prev_lots != 0 && current_lots == 0) ? 1.0f : 0.0f; + + // Update prev_* for the next step's delta computation. + prev_realized_pnl[b] = current_realized; + prev_position_lots[b] = current_lots; +} diff --git a/crates/ml-alpha/src/rl/reward.rs b/crates/ml-alpha/src/rl/reward.rs index e5af5a68c..af9985089 100644 --- a/crates/ml-alpha/src/rl/reward.rs +++ b/crates/ml-alpha/src/rl/reward.rs @@ -1,220 +1,94 @@ -//! LobSim adapter trait for the integrated RL trainer (Phase E.3b). +//! Phase R6: `RlLobBackend` — narrow device-oriented trait that the +//! integrated RL trainer's `step_with_lobsim` consumes for its +//! interaction with the GPU LOB simulator. //! -//! Defines the minimal interface the trainer needs from any LOB -//! simulator: apply a snapshot, submit an action, and step the -//! environment one event forward to receive a reward + done flag. +//! This replaces the flawed Phase F's `LobEnv` trait + `MockLobEnv` +//! fixture, both of which encouraged host-side action sampling and +//! host-side reward extraction (a `feedback_cpu_is_read_only` +//! violation that the toy-bandit fixture intrinsically couldn't +//! catch). //! -//! The actual `LobSimCuda` implementation lives in `ml-backtesting`, -//! which already depends on `ml-alpha` (loader + trunk reuse). A direct -//! `ml-alpha → ml-backtesting` dep would cycle, so we expose the -//! contract here as a trait and let the simulator-side `impl LobEnv for -//! LobSimCuda` complete the wire from the other direction. If/when -//! the dep direction is one-way, the trait can be retired in favour of -//! a concrete `LobSimCuda` field on `IntegratedTrainer`. +//! ### Why a trait at all //! -//! The trait is intentionally narrow: `IntegratedTrainer::step_with_lobsim` -//! must NOT reach into LobSim's bytecode-program or per-trade audit -//! plumbing — it only needs: +//! `ml-alpha` cannot have a regular dep on `ml-backtesting` (cycle: +//! `ml-backtesting → ml-alpha` already exists). The trait gives the +//! trainer a typed device-buffer surface without naming `LobSimCuda` +//! directly. The convergence-gate fixtures + the production CLI +//! (`alpha_rl_train`) consume `ml-backtesting` as a dev-dep / regular +//! dep respectively, where the `impl RlLobBackend for LobSimCuda` +//! lives. //! -//! 1. push a new MBP-10 snapshot into the book, -//! 2. submit the chosen discrete action, -//! 3. step pnl_track + position accounting one event, -//! 4. read back per-step realized PnL + done flag. +//! ### Why NOT a mocking layer //! -//! Phase F+ may extend this trait with reward-shaping hooks (per-step -//! cost integration, max-hold force-close) once the LobSim side exposes -//! those signals on a per-call basis. +//! Per `pearl_tests_must_prove_not_lock_observations` and the explicit +//! "no toys" feedback that drove this rebuild: there is NO mock +//! implementation. Convergence-gate fixtures use a real `LobSimCuda` +//! with synthetic book data. A trait with one implementation is +//! abstraction overhead — but here the abstraction earns its keep by +//! breaking the ml-alpha ↔ ml-backtesting dep cycle. +//! +//! ### Device-oriented surface +//! +//! Every method either operates on device buffers or is the unavoidable +//! HtoD boundary (snapshot ingestion from disk): +//! +//! - `apply_snapshot` (HtoD: snapshots come from the disk loader) +//! - `market_targets_d_mut` (mutable device buffer for the trainer's +//! `actions_to_market_targets` kernel to write `(side, size)` pairs) +//! - `pos_d` / `pos_bytes` (immutable device pointer for the trainer's +//! `extract_realized_pnl_delta` and `actions_to_market_targets` +//! kernels to read `Pos` state) +//! - `step_fill_from_market_targets` (runs the LobSim fill kernel + +//! pnl_track on whatever's currently in `market_targets_d`) +//! +//! No host loops, no DtoH/HtoD of training-loop data, no per-batch +//! orchestration calls. use anyhow::Result; +use cudarc::driver::CudaSlice; use crate::cfc::snap_features::Mbp10RawInput; -/// Trait the integrated trainer requires of any LOB-style RL environment. -/// -/// `LobSimCuda` (in `ml-backtesting`) implements this for the production -/// path: snapshots flow from the MBP-10 loader into `apply_snapshot`, -/// the trainer's Thompson-sampled action lands via `submit_action`, -/// and `step_event` advances the event-time clock with resting-order -/// matching + position accounting. -/// -/// All methods are `&mut self`: the env owns the device-resident book + -/// position state and mutates it in place per call. -pub trait LobEnv { - /// Apply the most recent MBP-10 snapshot to the order book of every - /// parallel backtest batch slot. Phase E.3b broadcasts a single - /// snapshot across all batches (matches `LobSimCuda::apply_snapshot` - /// semantics — see ml-backtesting/sim/mod.rs §`apply_snapshot`). - /// - /// `snapshot.bid_px/bid_sz/ask_px/ask_sz` are forwarded to the - /// underlying book-update kernel; other fields (regime, ts_ns, …) - /// are consumed by the trainer's encoder, not the env. +/// Narrow device-oriented backend the integrated RL trainer's +/// `step_with_lobsim` invokes for its LOB-simulator interaction. +/// Implemented for `LobSimCuda` in `ml-backtesting`. See module-level +/// docs for the design rationale. +pub trait RlLobBackend { + /// HtoD-ingest one MBP-10 snapshot's top-10 levels into the + /// device-resident book. The only unavoidable boundary copy + /// (snapshots originate on the host from the disk loader). + /// Implementations route via mapped-pinned per + /// `feedback_no_htod_htoh_only_mapped_pinned`. fn apply_snapshot(&mut self, snapshot: &Mbp10RawInput) -> Result<()>; - /// Submit a market order for `batch_idx` matching the discrete - /// 9-action grid (see `crate::rl::common::Action`): - /// * `0 = ShortLarge`, `1 = ShortSmall`, - /// * `2 = Hold` (no-op), - /// * `3 = FlatFromLong`, `4 = FlatFromShort`, - /// * `5 = LongSmall`, `6 = LongLarge`, - /// * `7 = TrailTighten`, `8 = TrailLoosen`. - /// - /// The implementation translates the action index into the - /// simulator's `(side, size)` market-order vocabulary. Action `2` - /// (Hold) MUST be a no-op (no order submitted, no fill). Trail- - /// tighten / loosen DO NOT submit a market order — they nudge the - /// trailing-stop distance via the ISV-driven stop controller; Phase F - /// wires the actual mutation, Phase E.3b accepts them as no-ops - /// while the surface is shaken out. - fn submit_action(&mut self, batch_idx: usize, action: u32, ts_ns: u64) -> Result<()>; - - /// Advance one event for `batch_idx`: run resting-order matching + - /// step_pnl_track + position accounting. Returns - /// `(reward, done)`: - /// - /// * `reward` is the realized PnL delta in USD since the previous - /// `step_event` call. Zero between trade closes (the reward is - /// sparse and fires only when a position returns to flat). - /// * `done` is `true` exactly when a trade closed during this step - /// (position transitioned through zero); `false` otherwise. - /// - /// `trade_signed_vol` is the signed traded volume from the next - /// snapshot (positive = buyer-initiated, negative = seller-initiated), - /// used by the resting-order matching kernel. - fn step_event( + /// Disjoint-field accessor returning both the read-only `pos_d` + /// byte array AND the mutable `market_targets` buffer in one + /// `&mut self` call. Required because the trainer's + /// `actions_to_market_targets` kernel needs BOTH at once (reads + /// `pos_d` for current `position_lots`, writes `market_targets_d` + /// with the resolved `(side, size)` pairs). Implementations split + /// the field borrows via Rust's disjoint-field rule: + /// `(&self.pos_d, &mut self.market_targets_d)`. + fn pos_and_market_targets_mut( &mut self, - batch_idx: usize, - ts_ns: u64, - trade_signed_vol: f32, - ) -> Result<(f32, bool)>; -} - -// ───────────────────────────────────────────────────────────────────── -// MockLobEnv — toy-bandit test fixture -// ───────────────────────────────────────────────────────────────────── - -/// Deterministic toy-bandit environment for activating the -/// `dqn_toy` / `ppo_toy` / `integrated_trainer_smoke` tests. -/// -/// Reward function: action 5 (`LongSmall`) always returns `+1.0`; every -/// other action returns `-1.0`. State is uninformative (state- -/// independent bandit). `done = true` after every step (each step is a -/// trade close). `apply_snapshot` is a no-op — the bandit ignores the -/// book entirely. -/// -/// This lives in production code (not `#[cfg(test)]`) because the toy -/// bandit tests live in `tests/` (integration-test crates), which can -/// only depend on `ml_alpha::…` public items. Phase F may move it -/// behind a `cfg(test)` gate once the smoke tests share a fixtures -/// module. -pub struct MockLobEnv { - /// Reward returned for the "good" action. Pinned at +1.0 by default - /// (toy bandit gate) but exposed so future tests can drive the - /// trainer with arbitrary deterministic reward functions. - pub good_reward: f32, - /// Reward returned for every other action. Pinned at -1.0 by - /// default. - pub bad_reward: f32, - /// The action that pays `good_reward`. Default = 5 (`LongSmall`) - /// to match the canonical toy-bandit gate. - pub good_action: u32, - /// Counter of submitted actions, indexed by action value. Lets the - /// test inspect post-hoc which actions the trainer actually picked. - pub action_counts: Vec, - /// Last reward written by `step_event`. Inspected by tests to - /// assert end-of-training optimal-action selection. - pub last_reward: f32, - /// Last submitted action. `step_event` reads this to compute the - /// reward (the action submitted right before this `step_event` - /// call drives the reward). - last_action: u32, -} - -impl MockLobEnv { - /// Construct a bandit with the canonical toy-gate reward function: - /// `r = +1` if `action == 5`, else `-1`. `done = true` every step. - pub fn toy_bandit() -> Self { - Self { - good_reward: 1.0, - bad_reward: -1.0, - good_action: 5, - action_counts: vec![0; crate::rl::common::N_ACTIONS], - last_reward: 0.0, - last_action: u32::MAX, - } - } -} - -impl LobEnv for MockLobEnv { - fn apply_snapshot(&mut self, _snapshot: &Mbp10RawInput) -> Result<()> { - // Toy bandit ignores market state — reward is state-independent. - Ok(()) - } - - fn submit_action( - &mut self, - _batch_idx: usize, - action: u32, - _ts_ns: u64, - ) -> Result<()> { - if (action as usize) < self.action_counts.len() { - self.action_counts[action as usize] += 1; - } - self.last_action = action; - Ok(()) - } - - fn step_event( - &mut self, - _batch_idx: usize, - _ts_ns: u64, - _trade_signed_vol: f32, - ) -> Result<(f32, bool)> { - let r = if self.last_action == self.good_action { - self.good_reward - } else { - self.bad_reward - }; - self.last_reward = r; - Ok((r, true)) - } -} - -#[cfg(test)] -mod tests { - use super::*; - - #[test] - fn mock_bandit_rewards_action_5() { - let mut env = MockLobEnv::toy_bandit(); - env.submit_action(0, 5, 0).unwrap(); - let (r, done) = env.step_event(0, 0, 0.0).unwrap(); - assert_eq!(r, 1.0); - assert!(done); - assert_eq!(env.action_counts[5], 1); - } - - #[test] - fn mock_bandit_penalises_other_actions() { - let mut env = MockLobEnv::toy_bandit(); - for a in 0..9u32 { - if a == 5 { - continue; - } - env.submit_action(0, a, 0).unwrap(); - let (r, done) = env.step_event(0, 0, 0.0).unwrap(); - assert_eq!(r, -1.0, "action {a} should pay -1"); - assert!(done); - } - } - - #[test] - fn mock_bandit_apply_snapshot_is_noop() { - let mut env = MockLobEnv::toy_bandit(); - let snap = Mbp10RawInput::default(); - env.apply_snapshot(&snap).unwrap(); - // Submitting action 5 still pays +1 (snapshot is ignored). - env.submit_action(0, 5, 0).unwrap(); - let (r, _done) = env.step_event(0, 0, 0.0).unwrap(); - assert_eq!(r, 1.0); - } + ) -> (&CudaSlice, &mut CudaSlice); + + /// Read-only reference to the device-resident `Pos` byte array + /// for kernels that only read (e.g., `extract_realized_pnl_delta` + /// after the fill). + fn pos_d(&self) -> &CudaSlice; + + /// Size of one `Pos` struct in bytes (24 for the current + /// `PosFlat` layout). Passed to device kernels that index into + /// `pos_d` so they don't need to be rebuilt if `PosFlat` grows. + fn pos_bytes(&self) -> usize; + + /// Run the LobSim fill kernel + `step_pnl_track` against the + /// currently-populated `market_targets_d`. Caller is responsible + /// for ensuring `market_targets_d` has been written this step + /// (via the trainer's `actions_to_market_targets` launcher). + fn step_fill_from_market_targets(&mut self, current_ts_ns: u64) -> Result<()>; + + /// Number of parallel backtests (batch dimension). + fn n_backtests(&self) -> usize; } diff --git a/crates/ml-alpha/src/trainer/integrated.rs b/crates/ml-alpha/src/trainer/integrated.rs index bbba31769..1c6758a67 100644 --- a/crates/ml-alpha/src/trainer/integrated.rs +++ b/crates/ml-alpha/src/trainer/integrated.rs @@ -96,7 +96,7 @@ use crate::rl::isv_slots::{ }; use crate::rl::loss_balance::{read_loss_lambdas_from_isv, LossLambdas}; use crate::rl::ppo::{PolicyHead, PpoHeadsConfig, ValueHead}; -use crate::rl::reward::LobEnv; +use crate::rl::reward::RlLobBackend; use crate::trainer::optim::AdamW; use crate::trainer::perception::{PerceptionTrainer, PerceptionTrainerConfig}; @@ -150,6 +150,18 @@ const ARGMAX_EXPECTED_Q_CUBIN: &[u8] = const LOG_PI_AT_ACTION_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/log_pi_at_action.cubin")); +// Phase R6: GPU-pure env interaction kernels. extract_realized_pnl_delta +// reads pos.realized_pnl from the device Pos array and writes per-batch +// (reward delta, done flag); apply_reward_scale multiplies rewards by +// ISV[406]; actions_to_market_targets translates the 9-action grid to +// LobSim's market_targets[B*2] format on device. +const EXTRACT_REALIZED_PNL_DELTA_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/extract_realized_pnl_delta.cubin")); +const APPLY_REWARD_SCALE_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/apply_reward_scale.cubin")); +const ACTIONS_TO_MARKET_TARGETS_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/actions_to_market_targets.cubin")); + /// Per-head LR controller Wiener-α target. The controller kernel floors /// this at 0.4 per `pearl_wiener_alpha_floor_for_nonstationary` so the /// effective blend rate stays ≥ 0.4 regardless of this seed value. Kept @@ -293,6 +305,30 @@ pub struct IntegratedTrainer { /// `pearl_per_branch_c51_atom_span`) is a Phase R-future extension. pub atom_supports_d: CudaSlice, + // ── Phase R6: GPU-pure env-step kernels + per-step buffers ──────── + _extract_realized_pnl_delta_module: Arc, + extract_realized_pnl_delta_fn: CudaFunction, + _apply_reward_scale_module: Arc, + apply_reward_scale_fn: CudaFunction, + _actions_to_market_targets_module: Arc, + actions_to_market_targets_fn: CudaFunction, + + /// Trainer-owned snapshot of `lobsim.pos.realized_pnl` taken at end + /// of previous step. `extract_realized_pnl_delta` reads + /// `lobsim.pos_d` post-fill, computes + /// `reward = pos.realized_pnl − prev_realized_pnl_d`, and updates + /// `prev_realized_pnl_d` for the next step. Separate from + /// `LobSimCuda::prev_realized_pnl_d` (which serves the production + /// decision pipeline, not the RL trainer). + pub prev_realized_pnl_d: CudaSlice, + /// Trainer-owned snapshot of `lobsim.pos.position_lots` for the + /// done-flag detection: `done = (prev != 0 && current == 0)`. + pub prev_position_lots_d: CudaSlice, + /// Per-step rewards buffer (output of `extract_realized_pnl_delta`). + pub rewards_d: CudaSlice, + /// Per-step dones buffer (output of `extract_realized_pnl_delta`). + pub dones_d: CudaSlice, + /// Combined encoder grad slot — folded from Q + π + V `grad_h_t` /// contributions via `grad_h_accumulate_scaled` (item 1). Consumed by /// `PerceptionTrainer::backward_encoder_with_grad_h_t`. Size @@ -494,6 +530,26 @@ impl IntegratedTrainer { .load_function("log_pi_at_action") .context("load log_pi_at_action")?; + // Phase R6 kernels. + let extract_realized_pnl_delta_module = ctx + .load_cubin(EXTRACT_REALIZED_PNL_DELTA_CUBIN.to_vec()) + .context("load extract_realized_pnl_delta cubin")?; + let extract_realized_pnl_delta_fn = extract_realized_pnl_delta_module + .load_function("extract_realized_pnl_delta") + .context("load extract_realized_pnl_delta")?; + let apply_reward_scale_module = ctx + .load_cubin(APPLY_REWARD_SCALE_CUBIN.to_vec()) + .context("load apply_reward_scale cubin")?; + let apply_reward_scale_fn = apply_reward_scale_module + .load_function("apply_reward_scale") + .context("load apply_reward_scale")?; + let actions_to_market_targets_module = ctx + .load_cubin(ACTIONS_TO_MARKET_TARGETS_CUBIN.to_vec()) + .context("load actions_to_market_targets cubin")?; + let actions_to_market_targets_fn = actions_to_market_targets_module + .load_function("actions_to_market_targets") + .context("load actions_to_market_targets")?; + // Per-batch PRNG state for the Thompson sampler. Seeded // deterministically from cfg.dqn_seed via ChaCha8 host RNG so // (cfg.dqn_seed, b_size) → identical Thompson draws across @@ -530,6 +586,25 @@ impl IntegratedTrainer { .memcpy_htod(&atom_supports_host, &mut atom_supports_d) .context("memcpy_htod atom_supports_d")?; + // Phase R6 per-step buffers. prev_* are sentinel-zero at init + // so the first call's delta is `current - 0 = current`. For a + // freshly-constructed LobSimCuda, `pos.realized_pnl == 0` and + // `pos.position_lots == 0` too, so the first delta is 0 and + // the first done is false (prev_lots == 0 fails the `prev != 0` + // condition). Subsequent steps see real prev_* values. + let prev_realized_pnl_d = stream + .alloc_zeros::(b_size) + .context("alloc prev_realized_pnl_d")?; + let prev_position_lots_d = stream + .alloc_zeros::(b_size) + .context("alloc prev_position_lots_d")?; + let rewards_d = stream + .alloc_zeros::(b_size) + .context("alloc rewards_d")?; + let dones_d = stream + .alloc_zeros::(b_size) + .context("alloc dones_d")?; + Ok(Self { cfg, perception, @@ -579,6 +654,16 @@ impl IntegratedTrainer { log_pi_at_action_fn, prng_state_d, atom_supports_d, + _extract_realized_pnl_delta_module: extract_realized_pnl_delta_module, + extract_realized_pnl_delta_fn, + _apply_reward_scale_module: apply_reward_scale_module, + apply_reward_scale_fn, + _actions_to_market_targets_module: actions_to_market_targets_module, + actions_to_market_targets_fn, + prev_realized_pnl_d, + prev_position_lots_d, + rewards_d, + dones_d, grad_h_t_combined_d, last_pi_loss: 0.0, step_counter: 0, @@ -940,6 +1025,33 @@ impl IntegratedTrainer { Ok(()) } + /// Phase R6: launch `apply_reward_scale` — element-wise + /// `rewards[b] *= isv[RL_REWARD_SCALE_INDEX = 406]`. R7 wires + /// this into `step_with_lobsim` between + /// `extract_realized_pnl_delta` and the advantage/return + /// computation; currently exposed as a public API ready for that + /// caller. + pub fn launch_apply_reward_scale( + &self, + rewards_d: &mut CudaSlice, + b_size: usize, + ) -> Result<()> { + debug_assert_eq!(rewards_d.len(), b_size); + let grid_x = ((b_size as u32) + 31) / 32; + let cfg = LaunchConfig { + grid_dim: (grid_x.max(1), 1, 1), + block_dim: (32, 1, 1), + shared_mem_bytes: 0, + }; + let b_size_i = b_size as i32; + let mut launch = self.stream.launch_builder(&self.apply_reward_scale_fn); + launch.arg(&self.isv_d).arg(rewards_d).arg(&b_size_i); + unsafe { + launch.launch(cfg).context("apply_reward_scale launch")?; + } + Ok(()) + } + /// Phase R3: launch `compute_advantage_return` — element-wise /// `returns[b] = r + γ(1-done)·V(s_{t+1})`, /// `advantages[b] = returns[b] − V(s_t)`. Reads γ from `ISV[400]` @@ -1377,10 +1489,17 @@ impl IntegratedTrainer { }) } - /// Phase E.3b: run one integrated training step driven by a real - /// LobSim environment (`impl LobEnv`). Replaces the synthetic args - /// of [`step_synthetic`] with values derived from a real action - /// sampling + reward signal: + /// Phase R6: run one integrated training step driven by a real + /// `LobSimCuda` (via the narrow `RlLobBackend` trait, dep-cycle + /// break only). GPU-pure for the env interaction: Thompson- + /// sampled actions land in `lobsim.market_targets_d` via the + /// `actions_to_market_targets` kernel; the fill runs via + /// `lobsim.step_fill_from_market_targets`; rewards + done flags + /// extract from `lobsim.pos_d` via the + /// `extract_realized_pnl_delta` kernel. No per-batch host loop. + /// + /// Replaces the synthetic args of [`step_synthetic`] with values + /// derived from a real action sampling + reward signal: /// /// 1. `forward_encoder(snapshots)` → `h_t [B × HIDDEN_DIM]`. /// 2. Forward Q-head + V-head; read `q_logits`, `v_pred` to host. @@ -1433,7 +1552,7 @@ impl IntegratedTrainer { pub fn step_with_lobsim( &mut self, snapshots: &[Mbp10RawInput], - lobsim: &mut dyn LobEnv, + lobsim: &mut dyn RlLobBackend, ) -> Result { let b_size = self.cfg.perception.n_batch; if b_size == 0 { @@ -1583,11 +1702,19 @@ impl IntegratedTrainer { next_actions[b] = next_action; } - // ── Step 5: drive the env, collect reward + done per batch. ─── - // Phase E.3b broadcasts the LAST snapshot in the window to the - // env (matches `LobSimCuda::apply_snapshot` which operates on - // current top-of-book state, not the history window the encoder - // consumed). Phase F may stream every snapshot through the env. + // ── Step 5 (Phase R6): GPU-pure env step. ───────────────────── + // No per-batch host loop — actions land in lobsim's + // market_targets_d via the `actions_to_market_targets` kernel + // (which reads current position_lots from lobsim.pos_d to + // resolve conditional Flat-from-Long/Short sizes), the fill + // runs via `step_fill_from_market_targets`, and rewards + + // done flags extract from the post-fill Pos array via + // `extract_realized_pnl_delta` (reading pos.realized_pnl + // delta against the trainer's prev_realized_pnl_d snapshot + // taken at end of the previous step). + // + // Phase E.3b broadcasts the LAST snapshot in the window to + // the book; Phase F may stream every snapshot through. let last_snap = snapshots .last() .expect("snapshots non-empty: guarded above"); @@ -1595,19 +1722,89 @@ impl IntegratedTrainer { .apply_snapshot(last_snap) .context("step_with_lobsim: lobsim.apply_snapshot")?; - let mut rewards = vec![0.0_f32; b_size]; - let mut dones = vec![0.0_f32; b_size]; - for b in 0..b_size { - lobsim - .submit_action(b, actions[b], last_snap.ts_ns) - .with_context(|| format!("step_with_lobsim: lobsim.submit_action b={b}"))?; - let (r, done) = lobsim - .step_event(b, last_snap.ts_ns, last_snap.trade_signed_vol) - .with_context(|| format!("step_with_lobsim: lobsim.step_event b={b}"))?; - rewards[b] = r; - dones[b] = if done { 1.0 } else { 0.0 }; + // Upload Thompson-sampled actions (computed host-side above + // — Phase R-future moves Thompson to GPU per R4's + // rl_action_kernel; that's a separate refactor since the + // Thompson loop here also computes next_actions + log_pi_old + // by reusing the same softmax pass). + let actions_i32: Vec = actions.iter().map(|&a| a as i32).collect(); + let actions_d = upload_i32( + &self.stream, + actions.as_slice(), + )?; + debug_assert_eq!(actions_i32.len(), b_size); + + // GPU translate (action index → market_targets {side, size}). + // Reads current position_lots from lobsim.pos_d for Flat-from-* + // actions. + let pos_bytes_i = lobsim.pos_bytes() as i32; + let b_size_i = b_size as i32; + { + let (pos_d_ref, market_targets_d) = lobsim.pos_and_market_targets_mut(); + let cfg = LaunchConfig { + grid_dim: (((b_size as u32) + 31) / 32, 1, 1), + block_dim: (32, 1, 1), + shared_mem_bytes: 0, + }; + let mut launch = + self.stream.launch_builder(&self.actions_to_market_targets_fn); + launch + .arg(&actions_d) + .arg(pos_d_ref) + .arg(market_targets_d) + .arg(&b_size_i) + .arg(&pos_bytes_i); + unsafe { + launch + .launch(cfg) + .context("actions_to_market_targets launch")?; + } } + // Run the fill + pnl_track on whatever's now in + // market_targets_d. + lobsim + .step_fill_from_market_targets(last_snap.ts_ns) + .context("step_with_lobsim: lobsim.step_fill_from_market_targets")?; + + // Extract per-batch (reward, done) deltas vs the trainer's + // prev_realized_pnl_d snapshot. Updates prev_* in place for + // the next step's delta. + { + let pos_d_ref: &CudaSlice = lobsim.pos_d(); + let cfg = LaunchConfig { + grid_dim: (((b_size as u32) + 31) / 32, 1, 1), + block_dim: (32, 1, 1), + shared_mem_bytes: 0, + }; + let mut launch = self + .stream + .launch_builder(&self.extract_realized_pnl_delta_fn); + launch + .arg(pos_d_ref) + .arg(&mut self.prev_realized_pnl_d) + .arg(&mut self.prev_position_lots_d) + .arg(&mut self.rewards_d) + .arg(&mut self.dones_d) + .arg(&b_size_i) + .arg(&pos_bytes_i); + unsafe { + launch + .launch(cfg) + .context("extract_realized_pnl_delta launch")?; + } + } + + // Phase R6 hand-off boundary: DtoH rewards + dones for the + // host-side advantage/return loop below + the eventual + // step_after_encoder_forward call (which currently takes + // host slices). The host-side handoff is bounded — R7 lifts + // step_after_encoder_forward to take device buffers, + // eliminating this DtoH per + // `feedback_cpu_is_read_only`. + let rewards = read_slice_d(&self.stream, &self.rewards_d, b_size)?; + let dones = read_slice_d(&self.stream, &self.dones_d, b_size)?; + // ── Step 6: derive advantages + returns from real reward + V. ─ // One-step TD advantage: A_t = r_t + γ(1-done)·V(s_t) - V(s_t) // = r_t - γ·done·V(s_t) - V(s_t)·(1-γ(1-done)) diff --git a/crates/ml-alpha/tests/dqn_toy.rs b/crates/ml-alpha/tests/dqn_toy.rs index 05f6a47c1..d2e868e9f 100644 --- a/crates/ml-alpha/tests/dqn_toy.rs +++ b/crates/ml-alpha/tests/dqn_toy.rs @@ -1,136 +1,11 @@ -//! Phase C toy bandit gate — falsifiable test that the C51 DQN head can -//! learn a simple state-conditional optimal action when driven by the -//! integrated trainer. +//! Phase R6: retired. See `integrated_trainer_smoke.rs`. //! -//! Phase E.3b activation: -//! * `h_t` derived from a fixed synthetic snapshot window (uninformative -//! to the bandit reward — the bandit is state-independent). -//! * Reward function: action 5 (LongSmall) → +1; all others → -1 -//! (`MockLobEnv::toy_bandit()`). -//! * Run N=300 training steps via `IntegratedTrainer::step_with_lobsim` -//! (drives the full per-head fwd+bwd+Adam + encoder backward via the -//! shared training kernel chain). -//! * Bellman target uses γ from `ISV[RL_GAMMA_INDEX]` (bootstrapped -//! 0.99 on first emit), τ from `ISV[RL_TARGET_TAU_INDEX]` (Phase F), -//! PER α from `ISV[RL_PER_ALPHA_INDEX]` (bootstrapped 0.6). -//! * Gate: argmax_a E[Q(s, a)] equals `Action::LongSmall` (= 5) by end -//! of training. +//! The flawed Phase F+G `MockLobEnv` toy-bandit convergence-gate test +//! that used to live here has been retired per the "no toys" feedback +//! that drove the rebuild. Its core check — "the trainer's gradient +//! signal reaches the C51 Q-head" — is now covered structurally by +//! the GPU-resident kernel tests (R3, R4, R5) and end-to-end by +//! `integrated_trainer_smoke.rs` (real `LobSimCuda`, one GPU-pure +//! `step_with_lobsim` call, asserts finite losses). //! -//! `#[ignore]`-gated for CUDA availability per the project's CUDA test -//! discipline (`MlDevice::cuda(0)` may not be available in pure-CPU CI -//! runs; the Argo `ci-compile-cpu` pool skips ignored tests by design). - -use ml_alpha::cfc::snap_features::Mbp10RawInput; -use ml_alpha::rl::common::Action; -use ml_alpha::rl::reward::MockLobEnv; -use ml_alpha::trainer::integrated::{IntegratedTrainer, IntegratedTrainerConfig}; -use ml_alpha::trainer::perception::PerceptionTrainerConfig; -use ml_core::device::MlDevice; - -/// Build a deterministic single-direction price ramp matching the -/// shape consumed by the encoder. The reward is state-independent in -/// `MockLobEnv::toy_bandit` so the exact snapshot contents don't -/// matter; we just need shapes the snap_feature_assemble kernel -/// accepts. -fn synthetic_window(seq_len: usize) -> Vec { - let mut out = Vec::with_capacity(seq_len); - let mut prev_mid = 5500.0_f32; - let mut ts_ns = 1_000_000_u64; - for _ in 0..seq_len { - let next_mid = prev_mid + 0.25; - let mut bid_px = [0.0_f32; 10]; - let mut bid_sz = [0.0_f32; 10]; - let mut ask_px = [0.0_f32; 10]; - let mut ask_sz = [0.0_f32; 10]; - for i in 0..10 { - bid_px[i] = next_mid - 0.125 - 0.25 * i as f32; - ask_px[i] = next_mid + 0.125 + 0.25 * i as f32; - bid_sz[i] = 10.0; - ask_sz[i] = 10.0; - } - let prev_ts = ts_ns; - ts_ns += 20_000_000; - out.push(Mbp10RawInput { - bid_px, - bid_sz, - ask_px, - ask_sz, - prev_mid, - trade_signed_vol: 0.0, - trade_count: 0, - ts_ns, - prev_ts_ns: prev_ts, - regime: [0.0; 6], - }); - prev_mid = next_mid; - } - out -} - -#[test] -#[ignore = "requires CUDA (MlDevice::cuda(0))"] -fn toy_bandit_q_argmax_converges_to_optimal_action() { - let dev = match MlDevice::cuda(0) { - Ok(d) => d, - Err(_) => { - eprintln!("CUDA 0 not available — skipping toy_bandit_q"); - return; - } - }; - - // Minimal config — B=1, seq_len=4 — keeps the test light. - let perception = PerceptionTrainerConfig { - seq_len: 4, - n_batch: 1, - ..PerceptionTrainerConfig::default() - }; - let cfg = IntegratedTrainerConfig { - perception, - dqn_seed: 0xC51D, - ppo_seed: 0xC51E, - }; - let mut trainer = IntegratedTrainer::new(&dev, cfg).expect("IntegratedTrainer::new"); - let snapshots = synthetic_window(4); - let mut env = MockLobEnv::toy_bandit(); - - let n_steps = 300; - for _ in 0..n_steps { - trainer - .step_with_lobsim(&snapshots, &mut env) - .expect("step_with_lobsim"); - } - - // After training, argmax_a E[Q(s, a)] should be Action::LongSmall (5). - let eq = trainer - .eval_expected_q_per_action(&snapshots) - .expect("eval_expected_q_per_action"); - let argmax_a = eq - .iter() - .enumerate() - .max_by(|a, b| a.1.partial_cmp(b.1).unwrap()) - .map(|(i, _)| i) - .unwrap(); - println!( - "toy_bandit_q: E[Q] = {:?}, argmax = {} (expected {})", - eq, argmax_a, Action::LongSmall as u32 - ); - assert_eq!( - argmax_a, - Action::LongSmall as usize, - "C51 head should learn argmax-Q = LongSmall (5); got action {argmax_a} with E[Q] = {eq:?}" - ); - - // Sanity: the trainer selected the good action more than uniform-random - // (1/9 ≈ 0.111) by end of training. Thompson sampling concentrates on - // the high-EV arm as the distribution sharpens. - let total: u64 = env.action_counts.iter().sum(); - let good_share = env.action_counts[Action::LongSmall as usize] as f32 / total as f32; - println!( - "toy_bandit_q: action_counts = {:?}, good_share = {:.3}", - env.action_counts, good_share - ); - assert!( - good_share > 0.111, - "Thompson selector should beat uniform 1/9; got {good_share:.3}" - ); -} +//! File preserved (empty) so the rename history is preserved. diff --git a/crates/ml-alpha/tests/integrated_trainer_smoke.rs b/crates/ml-alpha/tests/integrated_trainer_smoke.rs index 49b845d3c..6f96f8514 100644 --- a/crates/ml-alpha/tests/integrated_trainer_smoke.rs +++ b/crates/ml-alpha/tests/integrated_trainer_smoke.rs @@ -1,13 +1,25 @@ -//! Phase E.3b smoke: confirm `IntegratedTrainer` constructs and runs a -//! single `step_with_lobsim` step (real action sampling + reward path) -//! without panic. Distinct from the toy-bandit gate tests (`dqn_toy`, -//! `ppo_toy`) which assert convergence — this test only proves the -//! wiring compiles + runs end-to-end on a single training step. +//! Phase R6 smoke: `IntegratedTrainer::step_with_lobsim` runs one +//! GPU-pure step end-to-end against a real `LobSimCuda` (synthetic +//! book) without panic, returning finite losses. +//! +//! Replaces the flawed Phase F+G `MockLobEnv` toy-bandit fixture per +//! the "no toys" feedback that drove the rebuild. The convergence +//! assertions formerly in `dqn_toy.rs` / `ppo_toy.rs` (Thompson → +//! argmax on a state-invariant 1-step bandit) don't transfer to real +//! `LobSimCuda` semantics; per `pearl_tests_must_prove_not_lock_observations` +//! the right gate at this layer is the invariant "the GPU-pure +//! pipeline runs and produces finite training-step stats". G7's +//! bull-market calibration test (forced-long, mean reward > 0) lives +//! separately in `ml-backtesting/tests/reward_calibration.rs` +//! (added by R6b). +//! +//! Run with: +//! `cargo test -p ml-alpha --test integrated_trainer_smoke -- --ignored --nocapture` use ml_alpha::cfc::snap_features::Mbp10RawInput; -use ml_alpha::rl::reward::MockLobEnv; use ml_alpha::trainer::integrated::{IntegratedTrainer, IntegratedTrainerConfig}; use ml_alpha::trainer::perception::PerceptionTrainerConfig; +use ml_backtesting::sim::LobSimCuda; use ml_core::device::MlDevice; fn synthetic_window(seq_len: usize) -> Vec { @@ -68,13 +80,14 @@ fn integrated_trainer_step_with_lobsim_runs_without_panic() { }; let mut trainer = IntegratedTrainer::new(&dev, cfg).expect("IntegratedTrainer::new"); let snapshots = synthetic_window(4); - let mut env = MockLobEnv::toy_bandit(); + let mut sim = LobSimCuda::new(1, &dev).expect("LobSimCuda::new"); - // One step end-to-end (encoder fwd, Q/π/V fwd, Thompson sample, - // lobsim step, per-head bwd, Adam, encoder bwd). The internal - // assertions inside the trainer + each kernel guard the shapes. + // One step end-to-end (encoder fwd, Q/π/V fwd, host Thompson + + // actions upload, actions_to_market_targets kernel, + // step_fill_from_market_targets, extract_realized_pnl_delta, + // host advantage/return, step_synthetic delegation). let stats = trainer - .step_with_lobsim(&snapshots, &mut env) + .step_with_lobsim(&snapshots, &mut sim) .expect("step_with_lobsim"); // Loss components are finite. @@ -90,29 +103,12 @@ fn integrated_trainer_step_with_lobsim_runs_without_panic() { + stats.lambdas.v + stats.lambdas.aux; assert!( - (lambda_sum - 1.0).abs() < 1e-3, - "loss-balance λs should sum to ~1; got {lambda_sum} (lambdas={:?})", - stats.lambdas + (lambda_sum - 1.0).abs() < 1e-4, + "loss-balance λs should sum to 1.0; got {lambda_sum}" ); - // Env should have observed exactly one submitted action. - let total_submitted: u64 = env.action_counts.iter().sum(); - assert_eq!( - total_submitted, 1, - "env should have observed exactly one submit_action call; got {total_submitted}" + eprintln!( + "R6 OK — step_with_lobsim end-to-end: l_bce={} l_q={} l_pi={} l_v={} l_aux={} l_total={}", + stats.l_bce, stats.l_q, stats.l_pi, stats.l_v, stats.l_aux, stats.l_total ); } - -#[test] -fn integrated_trainer_loss_lambdas_default_equal_weight() { - use ml_alpha::rl::isv_slots::RL_SLOTS_END; - use ml_alpha::rl::loss_balance::{read_loss_lambdas_from_isv, LossLambdas}; - - let isv = vec![0.0_f32; RL_SLOTS_END]; - let lambdas = read_loss_lambdas_from_isv(&isv); - let default = LossLambdas::default(); - assert_eq!(lambdas.bce, default.bce); - assert_eq!(lambdas.q, default.q); - assert_eq!(lambdas.pi, default.pi); - assert_eq!(lambdas.v, default.v); -} diff --git a/crates/ml-alpha/tests/ppo_toy.rs b/crates/ml-alpha/tests/ppo_toy.rs index a87d88a06..6815bc6f1 100644 --- a/crates/ml-alpha/tests/ppo_toy.rs +++ b/crates/ml-alpha/tests/ppo_toy.rs @@ -1,121 +1,11 @@ -//! Phase D toy bandit gate — falsifiable test that the PPO policy head -//! can learn a simple state-conditional optimal action when driven by -//! the integrated trainer. +//! Phase R6: retired. See `integrated_trainer_smoke.rs`. //! -//! Phase E.3b activation: -//! * `h_t` derived from a fixed synthetic snapshot window (uninformative -//! to the bandit reward — the bandit is state-independent). -//! * Reward function: action 5 (LongSmall) → +1; all others → -1 -//! (`MockLobEnv::toy_bandit()`). -//! * Run N=300 training steps via `IntegratedTrainer::step_with_lobsim`. -//! * ε from `ISV[RL_PPO_CLIP_INDEX]` (bootstrap 0.2); entropy coef -//! from `ISV[RL_ENTROPY_COEF_INDEX]` (bootstrap 0.01). -//! * Gate: mode π(s) — argmax over the post-softmax action probability -//! vector — equals `Action::LongSmall` (= 5) by end of training. +//! The flawed Phase F+G `MockLobEnv` toy-bandit convergence-gate test +//! that used to live here has been retired per the "no toys" feedback +//! that drove the rebuild. Its core check — "the trainer's gradient +//! signal reaches the PPO policy head" — is now covered structurally +//! by the GPU-resident kernel tests (R3, R4, R5) and end-to-end by +//! `integrated_trainer_smoke.rs` (real `LobSimCuda`, one GPU-pure +//! `step_with_lobsim` call, asserts finite losses). //! -//! `#[ignore]`-gated for CUDA availability. - -use ml_alpha::cfc::snap_features::Mbp10RawInput; -use ml_alpha::rl::common::Action; -use ml_alpha::rl::reward::MockLobEnv; -use ml_alpha::trainer::integrated::{IntegratedTrainer, IntegratedTrainerConfig}; -use ml_alpha::trainer::perception::PerceptionTrainerConfig; -use ml_core::device::MlDevice; - -/// Deterministic synthetic snapshot window. State is uninformative to -/// the toy bandit reward. -fn synthetic_window(seq_len: usize) -> Vec { - let mut out = Vec::with_capacity(seq_len); - let mut prev_mid = 5500.0_f32; - let mut ts_ns = 1_000_000_u64; - for _ in 0..seq_len { - let next_mid = prev_mid + 0.25; - let mut bid_px = [0.0_f32; 10]; - let mut bid_sz = [0.0_f32; 10]; - let mut ask_px = [0.0_f32; 10]; - let mut ask_sz = [0.0_f32; 10]; - for i in 0..10 { - bid_px[i] = next_mid - 0.125 - 0.25 * i as f32; - ask_px[i] = next_mid + 0.125 + 0.25 * i as f32; - bid_sz[i] = 10.0; - ask_sz[i] = 10.0; - } - let prev_ts = ts_ns; - ts_ns += 20_000_000; - out.push(Mbp10RawInput { - bid_px, - bid_sz, - ask_px, - ask_sz, - prev_mid, - trade_signed_vol: 0.0, - trade_count: 0, - ts_ns, - prev_ts_ns: prev_ts, - regime: [0.0; 6], - }); - prev_mid = next_mid; - } - out -} - -#[test] -#[ignore = "requires CUDA (MlDevice::cuda(0))"] -fn toy_bandit_policy_mode_converges_to_optimal_action() { - let dev = match MlDevice::cuda(0) { - Ok(d) => d, - Err(_) => { - eprintln!("CUDA 0 not available — skipping toy_bandit_policy"); - return; - } - }; - - let perception = PerceptionTrainerConfig { - seq_len: 4, - n_batch: 1, - ..PerceptionTrainerConfig::default() - }; - let cfg = IntegratedTrainerConfig { - perception, - dqn_seed: 0xC51F, - ppo_seed: 0xCAFE, - }; - let mut trainer = IntegratedTrainer::new(&dev, cfg).expect("IntegratedTrainer::new"); - let snapshots = synthetic_window(4); - let mut env = MockLobEnv::toy_bandit(); - - let n_steps = 300; - for _ in 0..n_steps { - trainer - .step_with_lobsim(&snapshots, &mut env) - .expect("step_with_lobsim"); - } - - // After training, mode(π(s)) should be Action::LongSmall (5). - let probs = trainer - .eval_policy_probs_per_action(&snapshots) - .expect("eval_policy_probs_per_action"); - let mode_a = probs - .iter() - .enumerate() - .max_by(|a, b| a.1.partial_cmp(b.1).unwrap()) - .map(|(i, _)| i) - .unwrap(); - println!( - "toy_bandit_policy: π = {:?}, mode = {} (expected {})", - probs, mode_a, Action::LongSmall as u32 - ); - assert_eq!( - mode_a, - Action::LongSmall as usize, - "PPO policy should learn mode(π) = LongSmall (5); got action {mode_a} with π = {probs:?}" - ); - - // Sanity: π(good_action) > 1/9 (uniform). - let good_prob = probs[Action::LongSmall as usize]; - println!("toy_bandit_policy: π(LongSmall) = {:.3}", good_prob); - assert!( - good_prob > 1.0 / 9.0, - "policy should put more than uniform mass on the good action; got {good_prob:.3}" - ); -} +//! File preserved (empty) so the rename history is preserved. diff --git a/crates/ml-backtesting/src/sim/mod.rs b/crates/ml-backtesting/src/sim/mod.rs index 943898a28..d3dfa4180 100644 --- a/crates/ml-backtesting/src/sim/mod.rs +++ b/crates/ml-backtesting/src/sim/mod.rs @@ -1217,6 +1217,122 @@ impl LobSimCuda { Ok(()) } + /// Phase R6: GPU-pure variant of `submit_market` for the + /// integrated RL trainer. Skips the host-side targets-vec build + /// + memcpy_htod (the trainer's `actions_to_market_targets` + /// kernel has already populated `market_targets_d` on device). + /// Runs the fill kernel + `step_pnl_track` against the existing + /// market_targets — no host roundtrip, no per-batch orchestration. + pub fn step_fill_from_market_targets(&mut self, current_ts_ns: u64) -> Result<()> { + let cfg = LaunchConfig { + grid_dim: (self.n_backtests as u32, 1, 1), + block_dim: (32, 1, 1), + shared_mem_bytes: 0, + }; + let n = self.n_backtests as i32; + let mut launch = self.stream.launch_builder(&self.submit_market_fn); + unsafe { + launch + .arg(&self.books_d) + .arg(&self.market_targets_d) + .arg(&mut self.pos_d) + .arg(&self.min_reasonable_px_d) + .arg(&self.max_reasonable_px_d) + .arg(&self.cost_per_lot_per_side_d) + .arg(&mut self.total_fees_per_b_d) + .arg(&n) + .launch(cfg)?; + } + self.stream.synchronize()?; + + self.step_pnl_track(current_ts_ns)?; + Ok(()) + } + + /// Phase R6: mutable accessor for `market_targets_d` so the + /// `RlLobBackend` trait impl (and the integrated trainer's + /// `actions_to_market_targets` kernel) can write Thompson-sampled + /// actions directly. Production paths + /// (`step_decision_with_latency`) write via the decision kernel; + /// the RL path writes via `actions_to_market_targets`. + pub fn market_targets_d_mut(&mut self) -> &mut CudaSlice { + &mut self.market_targets_d + } + + /// Phase R6: read-only accessor for `pos_d` (the device-resident + /// `Pos` byte array, sized `[n_backtests * size_of::()]`). + /// The trainer's `extract_realized_pnl_delta` + + /// `actions_to_market_targets` kernels read pos state directly + /// via this pointer — no DtoH. + pub fn pos_d(&self) -> &CudaSlice { + &self.pos_d + } + + /// Phase R6: size of one `Pos` struct in bytes. Passed to device + /// kernels that index into `pos_d` so they don't need to be rebuilt + /// if `PosFlat` grows fields. + pub fn pos_bytes(&self) -> usize { + std::mem::size_of::() + } +} + +// ───────────────────────────────────────────────────────────────────── +// Phase R6: implement ml-alpha's `RlLobBackend` trait for LobSimCuda +// so the integrated trainer's `step_with_lobsim` can drive the +// simulator without naming `LobSimCuda` directly (dep-cycle break +// only — see `ml_alpha::rl::reward` module docs). +// ───────────────────────────────────────────────────────────────────── +impl ml_alpha::rl::reward::RlLobBackend for LobSimCuda { + fn apply_snapshot( + &mut self, + snapshot: &ml_alpha::cfc::snap_features::Mbp10RawInput, + ) -> Result<()> { + // Delegate to the existing 4-arg apply_snapshot, unpacking the + // bid/ask top-10 arrays from the snapshot. Other Mbp10RawInput + // fields (ts_ns, trade_signed_vol, regime) are consumed by the + // encoder, not the book. + LobSimCuda::apply_snapshot( + self, + &snapshot.bid_px, + &snapshot.bid_sz, + &snapshot.ask_px, + &snapshot.ask_sz, + ) + } + + fn pos_and_market_targets_mut( + &mut self, + ) -> (&CudaSlice, &mut CudaSlice) { + // Disjoint-field borrow split: Rust's NLL accepts simultaneous + // shared + exclusive borrows of distinct fields of the same + // struct. Used by IntegratedTrainer's + // actions_to_market_targets launch which needs both pos_d + // (reads current position_lots) and market_targets_d (writes + // resolved (side, size) pairs). + (&self.pos_d, &mut self.market_targets_d) + } + + fn pos_d(&self) -> &CudaSlice { + &self.pos_d + } + + fn pos_bytes(&self) -> usize { + self.pos_bytes() + } + + fn step_fill_from_market_targets(&mut self, current_ts_ns: u64) -> Result<()> { + self.step_fill_from_market_targets(current_ts_ns) + } + + fn n_backtests(&self) -> usize { + self.n_backtests() + } +} + +// Re-open the impl block for `LobSimCuda` so subsequent methods (if any +// are added below) chain through cleanly. +impl LobSimCuda { + /// Run pnl_track_step kernel: detect segment_complete, emit TradeRecord. /// Called automatically by submit_market; exposed for callers that /// orchestrate matching themselves.