feat(rl): R6 — replace LobEnv with kernel-driven GPU-pure env step
Replaces the flawed Phase F+G LobEnv trait + LobSimEnvAdapter +
MockLobEnv fixture with a GPU-pure env-interaction path.
DELETED:
- LobEnv trait (host-method-per-batch surface)
- MockLobEnv fixture (the toy-bandit pattern that hid the production defects)
ADDED:
- RlLobBackend trait (src/rl/reward.rs): narrow device-oriented
surface (apply_snapshot, pos_and_market_targets_mut, pos_d,
pos_bytes, step_fill_from_market_targets, n_backtests). Single
purpose: break the ml-alpha ↔ ml-backtesting dep cycle. No
mocking layer.
- 3 new GPU-pure kernels (cuda/):
* extract_realized_pnl_delta.cu — reads pos.realized_pnl +
position_lots from device Pos array, writes per-batch
(reward delta, done flag), updates trainer's prev_* buffers.
* apply_reward_scale.cu — element-wise rewards *= ISV[406].
* actions_to_market_targets.cu — 9-action grid → market_targets
{side, size} on device, reading current position_lots for
conditional Flat-from-Long/Short.
- LobSimCuda impls RlLobBackend (ml-backtesting/src/sim/mod.rs)
plus a new step_fill_from_market_targets entry that runs
submit_market_immediate + step_pnl_track without the host-side
targets-vec build. pos_and_market_targets_mut returns disjoint
field borrows (&pos_d, &mut market_targets_d).
- IntegratedTrainer: 3 cubin includes, 3 module/function fields,
launch_apply_reward_scale launcher, prev_realized_pnl_d /
prev_position_lots_d / rewards_d / dones_d buffers,
step_with_lobsim signature switched to RlLobBackend, body's
Step 5 rewritten as kernel-driven (no per-batch host loop, no
individual submit_action calls).
R6 PARTIAL — work R7 lifts:
- Thompson + log_pi stay host (R7 uses R4 kernels)
- mean_abs_pnl EMA stays host (R7 uses R3 ema_update_on_done)
- Advantage/return stays host (R7 uses R3 compute_advantage_return)
- 4 final DtoH copies for step_synthetic's host-slice signature
(R7 lifts step_after_encoder_forward to device-buffer args)
The "DtoH the device rewards/dones at end" comment in
step_with_lobsim documents the boundary. After R7, hot path has
zero host loops other than the now-unused Thompson host loop
(which R7 retires in favour of rl_action_kernel from R4).
Tests:
- integrated_trainer_smoke.rs: rewritten — real LobSimCuda with
synthetic book, one step_with_lobsim call, asserts finite losses
+ λs sum to 1. Smoke gate, not a convergence test.
- dqn_toy.rs, ppo_toy.rs: retired (empty stubs documenting
rationale). Files preserved so rename history survives. The
MockLobEnv toy-bandit pattern intrinsically couldn't catch the
production defects #1, #2, #5 that motivated this rebuild.
ml-backtesting added as ml-alpha dev-dep (cycle-safe: production
direction is ml-backtesting → ml-alpha; dev edge only loads when
building ml-alpha's own tests/examples).
Build cache-bust v29. cargo check + cargo build for all R-phase
tests green (pre-existing heads_bit_equiv index-OOB persists).
Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
1
Cargo.lock
generated
1
Cargo.lock
generated
@@ -6094,6 +6094,7 @@ dependencies = [
|
||||
"cudarc",
|
||||
"data",
|
||||
"memmap2",
|
||||
"ml-backtesting",
|
||||
"ml-core",
|
||||
"ml-features",
|
||||
"rand 0.8.5",
|
||||
|
||||
@@ -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" }
|
||||
|
||||
@@ -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");
|
||||
|
||||
79
crates/ml-alpha/cuda/actions_to_market_targets.cu
Normal file
79
crates/ml-alpha/cuda/actions_to_market_targets.cu
Normal file
@@ -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;
|
||||
}
|
||||
30
crates/ml-alpha/cuda/apply_reward_scale.cu
Normal file
30
crates/ml-alpha/cuda/apply_reward_scale.cu
Normal file
@@ -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;
|
||||
}
|
||||
59
crates/ml-alpha/cuda/extract_realized_pnl_delta.cu
Normal file
59
crates/ml-alpha/cuda/extract_realized_pnl_delta.cu
Normal file
@@ -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;
|
||||
}
|
||||
@@ -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<u64>,
|
||||
/// 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<u8>, &mut CudaSlice<i32>);
|
||||
|
||||
/// 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<u8>;
|
||||
|
||||
/// 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;
|
||||
}
|
||||
|
||||
@@ -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<f32>,
|
||||
|
||||
// ── Phase R6: GPU-pure env-step kernels + per-step buffers ────────
|
||||
_extract_realized_pnl_delta_module: Arc<CudaModule>,
|
||||
extract_realized_pnl_delta_fn: CudaFunction,
|
||||
_apply_reward_scale_module: Arc<CudaModule>,
|
||||
apply_reward_scale_fn: CudaFunction,
|
||||
_actions_to_market_targets_module: Arc<CudaModule>,
|
||||
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<f32>,
|
||||
/// Trainer-owned snapshot of `lobsim.pos.position_lots` for the
|
||||
/// done-flag detection: `done = (prev != 0 && current == 0)`.
|
||||
pub prev_position_lots_d: CudaSlice<i32>,
|
||||
/// Per-step rewards buffer (output of `extract_realized_pnl_delta`).
|
||||
pub rewards_d: CudaSlice<f32>,
|
||||
/// Per-step dones buffer (output of `extract_realized_pnl_delta`).
|
||||
pub dones_d: CudaSlice<f32>,
|
||||
|
||||
/// 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::<f32>(b_size)
|
||||
.context("alloc prev_realized_pnl_d")?;
|
||||
let prev_position_lots_d = stream
|
||||
.alloc_zeros::<i32>(b_size)
|
||||
.context("alloc prev_position_lots_d")?;
|
||||
let rewards_d = stream
|
||||
.alloc_zeros::<f32>(b_size)
|
||||
.context("alloc rewards_d")?;
|
||||
let dones_d = stream
|
||||
.alloc_zeros::<f32>(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<f32>,
|
||||
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<IntegratedStepStats> {
|
||||
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<i32> = 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<u8> = 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))
|
||||
|
||||
@@ -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<Mbp10RawInput> {
|
||||
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.
|
||||
|
||||
@@ -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<Mbp10RawInput> {
|
||||
@@ -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);
|
||||
}
|
||||
|
||||
@@ -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<Mbp10RawInput> {
|
||||
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.
|
||||
|
||||
@@ -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<i32> {
|
||||
&mut self.market_targets_d
|
||||
}
|
||||
|
||||
/// Phase R6: read-only accessor for `pos_d` (the device-resident
|
||||
/// `Pos` byte array, sized `[n_backtests * size_of::<PosFlat>()]`).
|
||||
/// 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<u8> {
|
||||
&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::<crate::lob::PosFlat>()
|
||||
}
|
||||
}
|
||||
|
||||
// ─────────────────────────────────────────────────────────────────────
|
||||
// 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<u8>, &mut CudaSlice<i32>) {
|
||||
// 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<u8> {
|
||||
&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.
|
||||
|
||||
Reference in New Issue
Block a user