110d3b4125
Cleanup of compiler warnings flagged by both local cargo check and the cluster ensure-binary log. Per `feedback_no_hiding`, every site is either deleted or wired up — no #[allow] suppressions. Lib (5 sites): - gpu_backtest_evaluator.rs:34 — drop unused DevicePtrMut. - gpu_dqn_trainer.rs:49 — drop unused DevicePtrMut (8 device_ptr_mut calls don't need the trait import in current cudarc). Line 19852: drop unnecessary parens around `b * sh2`. - training_loop.rs:20 — drop unused DevicePtrMut; unbrace single- symbol use at 5766. - state_reset_registry.rs:4 — delete the 10-symbol use-block of slot constants. Names appear in description strings (documentation only), symbols are never referenced. Examples (3 sites): - alpha_dqn_h600_smoke.rs:181, 186 — drop COL_RAW_CLOSE, FEAT_DIM, FillCoeffs, FillModel imports. - alpha_baseline.rs:79 — delete unused MappedI32::read. Batched path uses read_all for N-element action readback; the single-element method was leftover from the pre-batched legacy path. Lib + examples now have zero removable warnings. The remaining unsafe_block lints (each cudarc kernel launch needs unsafe) are structural and not actionable under the project's -W unsafe-code policy. Co-Authored-By: Claude Opus 4.7 <noreply@anthropic.com>
4137 lines
209 KiB
Rust
4137 lines
209 KiB
Rust
#![allow(unsafe_code)]
|
||
|
||
//! Vectorized GPU backtest evaluator — single-stream sequential execution.
|
||
//!
|
||
//! Runs walk-forward evaluation entirely on GPU:
|
||
//! 1. Upload test window data once (prices + features)
|
||
//! 2. Step loop: GPU gather kernel → forward pass → env kernel (sequential on one stream)
|
||
//! 3. Metrics reduction kernel → single readback
|
||
//!
|
||
//! Four forward-pass paths:
|
||
//! - **`evaluate()`**: Generic closure-based path (Candle tensors). Works for any model.
|
||
//! - **`evaluate_dqn()`**: Pure-CUDA forward kernel path. Eliminates Candle dispatch
|
||
//! overhead by running `q_forward_dueling_warp_shmem` directly via NVRTC. Enables
|
||
//! future CUDA Graph capture.
|
||
//! - **`evaluate_ppo()`**: Pure-CUDA PPO actor forward + softmax + exposure collapse.
|
||
//! - **`evaluate_supervised()`**: Candle forward + CUDA signal→action conversion.
|
||
//!
|
||
//! The only GPU→CPU transfer is the final metrics readback (n_windows × 10 floats).
|
||
//! All per-step data (states, actions, rewards, action history) stays GPU-resident.
|
||
//! Done-flags are checked per-thread inside the env kernel (no mid-loop downloads).
|
||
//!
|
||
//! ## Single-stream design
|
||
//!
|
||
//! All kernels run sequentially on a single CUDA stream:
|
||
//! gather_states → forward pass → env_step → repeat
|
||
//!
|
||
//! The previous dual-stream design with CudaEvent synchronization added overhead
|
||
//! that negated any parallelism benefit on H100 (forward pass already occupies
|
||
//! 100+ SMs, leaving no room for env_step overlap). Removing the second stream
|
||
//! and event sync yields a 15-20% wall-time reduction.
|
||
|
||
use std::mem::ManuallyDrop;
|
||
use std::sync::Arc;
|
||
use cudarc::driver::{CudaEvent, CudaFunction, CudaGraph, CudaSlice, CudaStream, DevicePtr, LaunchConfig, PushKernelArg};
|
||
use tracing::info;
|
||
|
||
use ml_core::nvtx::NvtxRange;
|
||
use crate::MLError;
|
||
use super::gpu_weights::{BranchingWeightSet, DuelingWeightSet, PpoActorWeightSet};
|
||
use super::lob_bar::LobBar;
|
||
use super::mapped_pinned::{MappedF32Buffer, MappedI32Buffer};
|
||
|
||
// ── CUDA Graph wrapper ───────────────────────────────────────────────────────
|
||
//
|
||
// `CudaGraph` is `!Send + !Sync` because it contains raw CUDA pointers.
|
||
// Our `GpuBacktestEvaluator` is used single-threaded (created, used, and
|
||
// dropped on the same thread that owns the CUDA context), but all other
|
||
// fields are `Send + Sync` and existing callers rely on that.
|
||
//
|
||
// We wrap `CudaGraph` in a newtype with manual `Send + Sync` to preserve
|
||
// the struct's auto-derived `Send + Sync` status.
|
||
//
|
||
// # Safety
|
||
// The graph is only launched on the same stream/context that created it.
|
||
// The evaluator is not shared across threads in practice.
|
||
|
||
struct SendSyncGraph(CudaGraph);
|
||
|
||
// Safety: CudaGraph is bound to a specific CUDA context. The evaluator
|
||
// that owns it is always used from the thread that created the context.
|
||
// No concurrent access occurs.
|
||
unsafe impl Send for SendSyncGraph {}
|
||
// Safety: same reasoning — single-owner, no concurrent launch calls.
|
||
unsafe impl Sync for SendSyncGraph {}
|
||
|
||
// ── Precompiled kernel cubins ──────────────────────────────────────────────────
|
||
|
||
static ENV_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/backtest_env_kernel.cubin"));
|
||
static METRICS_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/backtest_metrics_kernel.cubin"));
|
||
static PPO_FORWARD_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/backtest_forward_ppo_kernel.cubin"));
|
||
static SUPERVISED_SIGNAL_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/backtest_forward_supervised_kernel.cubin"));
|
||
static BACKTEST_DQN_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/experience_kernels.cubin"));
|
||
/// Val plan_isv parity kernels (task #94, 2026-04-24).
|
||
static BACKTEST_PLAN_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/backtest_plan_kernel.cubin"));
|
||
|
||
/// Number of timesteps to batch per cuBLAS forward call in the chunked step loop.
|
||
///
|
||
/// With n_windows=5 and CHUNK_SIZE=64, each cuBLAS SGEMM operates on [320, state_dim]
|
||
/// instead of [5, state_dim] -- 64x better SM utilisation. Reduces total kernel
|
||
/// launches from ~576K to ~74K (7.8x fewer) for a 32K-step backtest.
|
||
///
|
||
/// Portfolio features (3 out of 48 dims) are stale within each chunk. For a
|
||
/// 512-step chunk at 1-minute bars, that is ~8.5 hours of stale portfolio data.
|
||
/// The Q-network primarily responds to the 42 market features; portfolio
|
||
/// features (position_size, unrealized_pnl, drawdown) have <2% feature
|
||
/// importance in causal sensitivity analysis. 8× fewer cuBLAS forward passes
|
||
/// (301 vs 2406 for 154K bars) outweighs the marginal staleness.
|
||
const DQN_BACKTEST_CHUNK_SIZE: usize = 512;
|
||
|
||
/// Per-window Kelly-stats buffer width (f32 slots). Stored in a buffer
|
||
/// separate from `portfolio_buf` so the `backtest_state_gather` kernel's
|
||
/// stride-8 indexing into portfolio state stays correct while the env
|
||
/// kernel still has the stats it needs for the shared `apply_kelly_cap`
|
||
/// and `record_kelly_trade_outcome` helpers in `trade_physics.cuh`.
|
||
/// Slots: win_count, loss_count, sum_wins, sum_losses. Must equal
|
||
/// `KELLY_STATS_SIZE` in `backtest_env_kernel.cu`.
|
||
const BACKTEST_KELLY_STATS_SIZE: usize = 4;
|
||
|
||
/// SP21 T2.2 Phase 1 (2026-05-10) — per-trade tape capacity per window.
|
||
/// 200_000 trades is a conservative upper bound: typical eval windows are
|
||
/// 200k bars (xmd6b: window_bars=214654) with at most ~1 close per bar
|
||
/// (observed trade_count=139251 at trades_per_bar=0.65). Setting capacity
|
||
/// equal to the bar count guarantees no trade is dropped under any
|
||
/// trading frequency. Memory cost per window: 4 × 4 bytes × 200k = 3.2MB
|
||
/// (5 × 4 with the per-window count u32). For typical n_windows=5
|
||
/// walk-forward setups: 16MB total — acceptable for eval (fold-boundary
|
||
/// alloc, not per-step). The kernel guards against overflow via
|
||
/// `if (slot < max_trades_per_window)` so larger-than-expected windows
|
||
/// gracefully truncate the tape instead of crashing.
|
||
pub const MAX_TRADES_PER_WINDOW: usize = 200_000;
|
||
|
||
/// SP21 T2.2 Phase 1 (2026-05-10) — host-side per-trade record decoded
|
||
/// from the GPU per-trade tape buffers. The kernel writes the SoA tape;
|
||
/// `read_per_trade_tape()` decodes packed `dir_mag` and zips the slices
|
||
/// into `Vec<EvalTrade>` for the enrichment phase. Field order matches
|
||
/// what the enrichment functions in `trainer/enrichment.rs` consume.
|
||
#[derive(Debug, Clone, Copy)]
|
||
pub struct EvalTrade {
|
||
/// SP21 T2.2 Phase 6.5 (2026-05-11) — eval-window membership.
|
||
/// Used by `compute_hindsight_labels` to recover which window's
|
||
/// state buffer to look up when synthesizing a counterfactual
|
||
/// replay tuple from `bar_index`. Set host-side by
|
||
/// `read_per_trade_tape` from the read loop's `w` variable —
|
||
/// the GPU per-trade tape buffers don't carry this (each window
|
||
/// writes to its own SoA slot, and the flatten loses the index).
|
||
pub window_index: u32,
|
||
/// Global bar index (0-based step within the window) when the trade
|
||
/// CLOSED. Used by E6 (winner indices) for priority-replay tagging.
|
||
pub bar_index: u32,
|
||
/// Realized per-trade return = `prev_position × (close − entry_price)
|
||
/// / pre_trade_equity`. Dimensionless (fraction of equity). Used by
|
||
/// E1 (Q-correction), E5 (agreement-Sharpe), E6/E7/E8.
|
||
pub pnl: f32,
|
||
/// SP21 T2.2 Phase 1.5 (2026-05-11) — predicted Q-value at trade
|
||
/// open. Captured by the env kernel from `q_values_per_window[w *
|
||
/// num_actions + action_taken]` at the open / reversal that started
|
||
/// this trade. Consumed by E1 (`compute_q_correction`: predicted_q
|
||
/// − pnl bias). 0.0 for trades closed under callers that don't
|
||
/// wire `q_values_per_window` (single-step `evaluate()` for
|
||
/// non-DQN paths).
|
||
pub predicted_q: f32,
|
||
/// Trade duration in bars (capture of `hold_time` BEFORE close). Used
|
||
/// by E3 (gamma from holding-bars) and E4 (urgency-error rate).
|
||
pub holding_bars: u32,
|
||
/// Direction at close: 0=Short, 1=Hold, 2=Long, 3=Flat. Decoded from
|
||
/// the kernel-packed `dir_mag` (hi-16 bits).
|
||
pub direction: u8,
|
||
/// Magnitude at close: 0=Quarter, 1=Half, 2=Full. Decoded from
|
||
/// `dir_mag` (lo-16 bits).
|
||
pub magnitude: u8,
|
||
}
|
||
|
||
// ── DQN backtest forward config ───────────────────────────────────────────────
|
||
|
||
/// Configuration for cuBLAS-based DQN forward pass in backtest evaluation.
|
||
///
|
||
/// Bundles network architecture dimensions, C51 distributional parameters,
|
||
/// and branching configuration needed by `evaluate_dqn()`.
|
||
#[derive(Debug, Clone)]
|
||
pub struct DqnBacktestConfig {
|
||
/// Shared trunk hidden layer 1 dimension.
|
||
pub shared_h1: usize,
|
||
/// Shared trunk hidden layer 2 dimension.
|
||
pub shared_h2: usize,
|
||
/// Value head hidden dimension.
|
||
pub value_h: usize,
|
||
/// Advantage head hidden dimension.
|
||
pub adv_h: usize,
|
||
/// C51 distributional atom count (typically 51).
|
||
pub num_atoms: usize,
|
||
/// Direction branch action count (branch 0, default 4 — Short/Hold/Long/Flat).
|
||
pub branch_0_size: usize,
|
||
/// Magnitude branch action count (branch 1, default 3 — 25%/50%/100%).
|
||
pub branch_1_size: usize,
|
||
/// Order-type branch action count (branch 2, default 3).
|
||
pub branch_2_size: usize,
|
||
/// Urgency branch action count (branch 3, default 3 — Patient/Normal/Aggressive).
|
||
pub branch_3_size: usize,
|
||
/// C51 minimum support value.
|
||
pub v_min: f32,
|
||
/// C51 maximum support value.
|
||
pub v_max: f32,
|
||
}
|
||
|
||
impl DqnBacktestConfig {
|
||
/// Create from the legacy `network_dims` tuple with standard C51 defaults.
|
||
///
|
||
/// Uses branch sizes [4, 3, 3, 3], num_atoms=51, v_range from reward_scale=10/gamma=0.95.
|
||
pub fn from_network_dims(network_dims: (usize, usize, usize, usize)) -> Self {
|
||
// v_min/v_max computed from reward_scale=10, gamma=0.95:
|
||
// Q* ~ 10 / (1-0.95) * 1.2 = 240, clamped to [20, 300] → ±240
|
||
let v_range = (10.0_f32 / (1.0 - 0.95) * 1.2).clamp(20.0, 300.0);
|
||
Self {
|
||
shared_h1: network_dims.0,
|
||
shared_h2: network_dims.1,
|
||
value_h: network_dims.2,
|
||
adv_h: network_dims.3,
|
||
num_atoms: 51,
|
||
branch_0_size: 4,
|
||
branch_1_size: 3,
|
||
branch_2_size: 3,
|
||
branch_3_size: 3,
|
||
v_min: -v_range,
|
||
v_max: v_range,
|
||
}
|
||
}
|
||
}
|
||
|
||
// ── Public types ──────────────────────────────────────────────────────────────
|
||
|
||
/// Per-window evaluation result returned after a full backtest run.
|
||
#[derive(Debug, Clone)]
|
||
pub struct WindowMetrics {
|
||
pub sharpe: f32,
|
||
pub total_pnl: f32,
|
||
pub max_drawdown: f32,
|
||
pub sortino: f32,
|
||
pub win_rate: f32,
|
||
pub total_trades: f32,
|
||
// Phase 3 extended metrics
|
||
pub var_95: f32,
|
||
pub cvar_95: f32,
|
||
pub calmar: f32,
|
||
pub omega_ratio: f32,
|
||
// Action distribution (GPU-counted, indices 10-12)
|
||
pub buy_count: f32,
|
||
pub sell_count: f32,
|
||
pub hold_count: f32,
|
||
// Unique action IDs observed (popcount of 9-bit mask, index 13)
|
||
pub unique_actions: f32,
|
||
// SP15 Wave 3b (2026-05-06) — cost-net sharpe (per spec §6.2) and
|
||
// four constant-policy counterfactual baselines (§6.4). All 5 fields
|
||
// are annualised host-side in `consume_metrics_after_event` by
|
||
// multiplying the kernel's `raw_sharpe` by `self.annualization_factor`
|
||
// (matches `sharpe`'s annualisation pattern from Phase 1.1.b).
|
||
pub sharpe_cost_net: f32,
|
||
pub baseline_buyhold_sharpe: f32,
|
||
pub baseline_hold_only_sharpe: f32,
|
||
pub baseline_momentum_sharpe: f32,
|
||
pub baseline_reversion_sharpe: f32,
|
||
// 2026-05-09 — per-regime WR instrumentation (audit follow-up).
|
||
// Trades are bucketed at trade-OPEN by the bar's ADX-norm value
|
||
// (feature[40], post-`RollingPercentileRank`) per
|
||
// `gpu_walk_forward.rs::classify_regime_from_features`:
|
||
// - Trending (T): ADX > 0.4
|
||
// - Ranging (R): ADX < 0.2
|
||
// - Volatile (V): otherwise
|
||
// `win_rate_*` is `wins / trades` per bucket (0.0 when zero trades,
|
||
// matching the aggregate `win_rate` empty-trade convention).
|
||
pub win_rate_trending: f32,
|
||
pub trade_count_trending: f32,
|
||
pub win_rate_ranging: f32,
|
||
pub trade_count_ranging: f32,
|
||
pub win_rate_volatile: f32,
|
||
pub trade_count_volatile: f32,
|
||
}
|
||
|
||
/// Configuration for the GPU backtest evaluator.
|
||
#[derive(Debug, Clone)]
|
||
pub struct GpuBacktestConfig {
|
||
pub max_position: f32,
|
||
pub tx_cost_bps: f32,
|
||
pub spread_cost: f32,
|
||
pub initial_capital: f32,
|
||
/// Contract multiplier for futures (e.g., 50 for ES, 20 for NQ).
|
||
/// Used with `max_leverage` to compute leverage-safe position limits.
|
||
pub contract_multiplier: f32,
|
||
/// Maximum leverage ratio relative to capital (e.g., 2.0 = 2× leverage).
|
||
/// When > 0, caps effective position at `capital * max_leverage / (price * multiplier)`.
|
||
pub max_leverage: f32,
|
||
/// Number of OFI features embedded after market features in the feature vector.
|
||
/// When > 0, the gather kernel reorders output to `[market, portfolio, OFI, pad]`
|
||
/// (model training order) instead of `[market, OFI, portfolio, pad]` (storage order).
|
||
/// Typical value: 8 (VPIN, Kyle's Lambda, OFI L1/L5, depth imbalance, bid/ask slope,
|
||
/// trade imbalance). Default: 0 (no OFI).
|
||
pub ofi_dim: usize,
|
||
/// Number of bars in one trading day (390 for 1-minute, 1 for daily).
|
||
/// Used to compute the correct Sharpe annualization: `sqrt(bars_per_day * 252)`.
|
||
/// Default: 390.0 (1-minute bar frequency).
|
||
pub bars_per_day: f32,
|
||
/// Per-step subsample stride applied to the upstream return series
|
||
/// (validation backtest subsamples every Nth bar at
|
||
/// `crates/ml/src/trainers/dqn/trainer/metrics.rs:517-518`). The
|
||
/// constructor divides `bars_per_day` by this stride before computing
|
||
/// `annualization_factor` so reported Sharpe/Sortino/Calmar reflect the
|
||
/// effective sampling cadence of the per-step return series, not the
|
||
/// raw bar cadence. Default: 1 (no subsampling). MUST be >= 1.
|
||
/// Per audit `docs/lookahead-bias-audit-2026-04-28.md` §5: prior code
|
||
/// used the unscaled `bars_per_day` for an internally subsampled series
|
||
/// and reported sharpe inflated by `sqrt(stride)` (= 2 for stride=4).
|
||
pub val_subsample_stride: u32,
|
||
/// Holding cost rate per bar per unit position. Continuous inventory penalty.
|
||
/// Default: 0.00000052 (~0.52 bps / 100 bars)
|
||
pub holding_cost_rate: f32,
|
||
/// Bars threshold for churn penalty. Trades within this window get graduated penalty.
|
||
pub churn_threshold: f32,
|
||
/// Scale factor for churn penalty.
|
||
pub churn_penalty_scale: f32,
|
||
/// Opportunity cost scale for flat position penalty.
|
||
/// Penalty = |q_spread| * opp_cost_scale when position ≈ 0.
|
||
/// Set to 0.0 during backtest evaluation (no Q-gap available post-hoc).
|
||
pub opportunity_cost_scale: f32,
|
||
/// Initial margin as fraction of notional value.
|
||
/// CME initial margin ~6% of notional for equity index futures.
|
||
/// Default: 0.06.
|
||
pub margin_pct: f32,
|
||
/// Trading days per year for annualization. CME equity/futures: 252. Crypto: 365.
|
||
pub trading_days_per_year: f32,
|
||
}
|
||
|
||
impl Default for GpuBacktestConfig {
|
||
fn default() -> Self {
|
||
Self {
|
||
max_position: 1.0,
|
||
tx_cost_bps: 0.1,
|
||
spread_cost: 0.0001,
|
||
initial_capital: 100_000.0,
|
||
contract_multiplier: 50.0,
|
||
max_leverage: 2.0,
|
||
ofi_dim: 0,
|
||
bars_per_day: 390.0, // 1-minute bar frequency (6.5h × 60min)
|
||
holding_cost_rate: 0.00000052,
|
||
churn_threshold: 10.0,
|
||
churn_penalty_scale: 0.001,
|
||
opportunity_cost_scale: 0.0, // disabled during backtest: no Q-gap available
|
||
trading_days_per_year: 252.0,
|
||
margin_pct: 0.06,
|
||
val_subsample_stride: 1,
|
||
}
|
||
}
|
||
}
|
||
|
||
impl GpuBacktestConfig {
|
||
/// Compute leverage-safe max position given a reference price.
|
||
///
|
||
/// Returns `min(max_position, capital * max_leverage / (price * multiplier))`.
|
||
/// If `max_leverage <= 0` or `price <= 0`, returns `max_position` uncapped.
|
||
pub fn effective_max_position(&self, reference_price: f32) -> f32 {
|
||
if self.max_leverage <= 0.0 || reference_price <= 0.0 || self.contract_multiplier <= 0.0 {
|
||
return self.max_position;
|
||
}
|
||
let notional_per_contract = reference_price * self.contract_multiplier;
|
||
let leverage_cap = self.initial_capital * self.max_leverage / notional_per_contract;
|
||
self.max_position.min(leverage_cap)
|
||
}
|
||
}
|
||
|
||
// ── Main struct ───────────────────────────────────────────────────────────────
|
||
|
||
/// GPU backtest evaluator — runs walk-forward evaluation without CPU roundtrips.
|
||
///
|
||
/// Upload window data once via `new()`, then call `evaluate()` with a model
|
||
/// forward function. The entire step loop runs on GPU; only the final metrics
|
||
/// are downloaded (n_windows × 14 floats).
|
||
///
|
||
/// All kernels execute sequentially on a single CUDA stream:
|
||
/// ```text
|
||
/// gather_states → forward pass → env_step → repeat
|
||
/// ```
|
||
#[allow(missing_debug_implementations)]
|
||
pub struct GpuBacktestEvaluator {
|
||
stream: Arc<CudaStream>,
|
||
|
||
env_kernel: CudaFunction,
|
||
env_batch_kernel: CudaFunction,
|
||
metrics_kernel: CudaFunction,
|
||
gather_kernel: CudaFunction,
|
||
/// Chunked gather: writes `[chunk_len, N, padded_sd]` rows into
|
||
/// `chunked_states_buf` in a single launch. Replaces the per-step
|
||
/// `gather_kernel` + DtoD-into-chunk loop in `submit_dqn_step_loop_cublas`.
|
||
/// Cuts per-chunk launches from `2 * chunk_len` (gather + DtoD) to **1**.
|
||
/// Loaded from the same `experience_kernels.cubin` as `gather_kernel`;
|
||
/// the per-step `gather_kernel` is retained for `evaluate()` /
|
||
/// `evaluate_ppo()` / `evaluate_supervised()` where the closure-based
|
||
/// caller still needs a single-step output.
|
||
gather_chunk_kernel: CudaFunction,
|
||
|
||
/// SP15 Wave 5 follow-up (2026-05-07): pre-loaded handle for the
|
||
/// `cost_net_sharpe_kernel` (cost-net Sharpe per-window). Originally
|
||
/// `launch_sp15_cost_net_sharpe` resolved the symbol via `load_cubin`
|
||
/// + `load_function` ON EVERY CALL inside the per-window eval hot loop
|
||
/// (`gpu_backtest_evaluator.rs:2877`), which is fragile across CUDA
|
||
/// context lifetimes — works in single-pass train-best context (smoke
|
||
/// `train-9bcwm` verified), fails in hyperopt-trial child stream
|
||
/// context (workflow `train-xggfc` trial 1 failed at `load
|
||
/// sp15_baseline_kernels cubin: ILLEGAL_ADDRESS`; after the host-side
|
||
/// load corrupted trial 1's context, trials 2-20 cascade-failed at
|
||
/// `Fork CUDA stream for trial`). Pre-loading once in `new()` mirrors
|
||
/// the `5d63762ab` precedent fix for `bn_tanh_concat_dd_kernel`.
|
||
sp15_cost_net_sharpe_kernel: CudaFunction,
|
||
/// SP15 Wave 5 follow-up (2026-05-07): pre-loaded handle for
|
||
/// `baseline_buyhold_kernel` (counterfactual buyhold baseline). See
|
||
/// `sp15_cost_net_sharpe_kernel` docstring for root-cause / precedent.
|
||
sp15_baseline_buyhold_kernel: CudaFunction,
|
||
/// SP15 Wave 5 follow-up (2026-05-07): pre-loaded handle for
|
||
/// `baseline_hold_only_kernel` (counterfactual hold-only baseline).
|
||
/// See `sp15_cost_net_sharpe_kernel` docstring for root-cause /
|
||
/// precedent.
|
||
sp15_baseline_hold_only_kernel: CudaFunction,
|
||
/// SP15 Wave 5 follow-up (2026-05-07): pre-loaded handle for
|
||
/// `baseline_naive_momentum_kernel` (counterfactual naive-momentum
|
||
/// baseline). See `sp15_cost_net_sharpe_kernel` docstring for
|
||
/// root-cause / precedent.
|
||
sp15_baseline_naive_momentum_kernel: CudaFunction,
|
||
/// SP15 Wave 5 follow-up (2026-05-07): pre-loaded handle for
|
||
/// `baseline_naive_reversion_kernel` (counterfactual naive-reversion
|
||
/// baseline). See `sp15_cost_net_sharpe_kernel` docstring for
|
||
/// root-cause / precedent.
|
||
sp15_baseline_naive_reversion_kernel: CudaFunction,
|
||
|
||
// Uploaded data (read-only; persists across the step loop).
|
||
// MappedF32/I32Buffer (cuMemHostAlloc DEVICEMAP) — host writes via
|
||
// `host_ptr.write_from_slice(...)`, kernels read via `dev_ptr` after a
|
||
// stream-sync barrier. Zero-copy CPU↔GPU path per
|
||
// `feedback_no_htod_htoh_only_mapped_pinned.md`. The DtoD-via-pinned
|
||
// helpers are forbidden by the same-day pre-commit guard
|
||
// (`scripts/pre-commit-hook.sh::check_no_dtod_via_pinned`, commit
|
||
// `5275932f4`) — direct mapped-pinned storage replaces them.
|
||
prices_buf: MappedF32Buffer, // [n_windows * max_len * 4] raw OHLC
|
||
features_buf: MappedF32Buffer, // [n_windows * max_len * feat_dim] z-normed features
|
||
window_lens_buf: MappedI32Buffer, // [n_windows]
|
||
|
||
// Mutable state (written by kernels each step). MappedF32Buffer's
|
||
// `dev_ptr` aliases the same pinned pages as `host_ptr`, so kernels
|
||
// can read AND write through `dev_ptr` (see mapped_pinned.rs module
|
||
// docstring + `plan_diag_buf` precedent in this file). The host can
|
||
// refresh `portfolio_buf` between epochs via `write_from_slice` —
|
||
// see `reset_evaluation_state` below.
|
||
portfolio_buf: MappedF32Buffer, // [n_windows * 8]
|
||
/// Kelly win/loss stats — separate from portfolio_buf so the gather
|
||
/// kernel's stride-8 indexing stays correct. [n_windows * 4].
|
||
kelly_stats_buf: CudaSlice<f32>,
|
||
/// Plan/ISV features per window [n_windows * 7 f32]: plan_progress, pnl_vs_target, pnl_vs_stop, conviction, conviction_drift, regime_shift, remaining_fraction (D.6)
|
||
plan_isv_buf: CudaSlice<f32>,
|
||
/// Val plan_isv parity (task #94, 2026-04-24). Live plan_params from
|
||
/// the plan MLP, filled each step by `QValueProvider::compute_plan_params`
|
||
/// and consumed by `backtest_plan_state_isv`. Layout [N, 6]:
|
||
/// {target_bars, profit_target, stop_loss, scale_aggression,
|
||
/// conviction, asymmetry}. Zero-initialised so the first step
|
||
/// before any plan MLP call sees safe defaults.
|
||
plan_params_buf: CudaSlice<f32>,
|
||
/// Val plan_isv parity: persistent plan state across bars within a
|
||
/// trade. Layout [N, 7]: {target_bars, profit_target, stop_loss,
|
||
/// scale_aggression, conviction_at_entry, asymmetry,
|
||
/// entry_regime_stability}. Activated on Flat→Positioned transitions
|
||
/// (via the live plan_params), zeroed on return to Flat. Plays the
|
||
/// role training's ps[23..29] plays in the experience buffer.
|
||
plan_state_buf: CudaSlice<f32>,
|
||
/// Val plan_isv parity diagnostic scratch [8 f32]: {active_fraction,
|
||
/// mean_|isv[0..2]|, mean_isv[3], mean_|isv[4..5]|, raw_active_count}.
|
||
/// Written by `backtest_plan_diag_reduce` at epoch boundary, then read by
|
||
/// the host via mapped-pinned `host_ptr` for HEALTH_DIAG emission.
|
||
/// Mapped pinned (`cuMemHostAlloc DEVICEMAP`) so the kernel writes through
|
||
/// the device alias and the host reads via volatile load — zero DtoH copy
|
||
/// per `feedback_no_htod_htoh_only_mapped_pinned.md`. Caller synchronises
|
||
/// the eval stream before reading.
|
||
plan_diag_buf: MappedF32Buffer,
|
||
step_rewards_buf: CudaSlice<f32>, // [n_windows]
|
||
step_returns_buf: CudaSlice<f32>, // [n_windows * max_len]
|
||
done_buf: CudaSlice<i32>, // [n_windows]
|
||
actions_buf: CudaSlice<i32>, // [n_windows]
|
||
actions_history_buf: CudaSlice<i32>, // [n_windows * max_len]
|
||
|
||
/// SP21 T2.2 Phase 1 (2026-05-10) — per-trade tape device buffers.
|
||
/// Written by `backtest_env_step` / `backtest_env_step_batch` at
|
||
/// trade-close events (single thread per window writes its own slot
|
||
/// — race-free counter increment without atomicAdd per
|
||
/// `feedback_no_atomicadd`). Read host-side via
|
||
/// `read_per_trade_tape()` after each eval window completes; the
|
||
/// resulting `Vec<EvalTrade>` is the authoritative per-trade signal
|
||
/// for the enrichment phase (E1-E8) — replaces the prior
|
||
/// fake-trade synthesizer (`extract_eval_trades_from_metrics`)
|
||
/// per `feedback_no_stubs`. Reset to zero at `reset_per_trade_tape()`
|
||
/// at the start of each eval window.
|
||
///
|
||
/// Layout: `[n_windows × MAX_TRADES_PER_WINDOW]` for the SoA buffers,
|
||
/// `[n_windows]` for the count. SoA chosen over AoS for GPU coalescing
|
||
/// (each kernel thread writes one slot per buffer at the same index).
|
||
/// `dir_mag` packs direction (hi-16) and magnitude (lo-16) so the
|
||
/// per-trade emission stays at 4 buffers + count.
|
||
per_trade_pnl_buf: CudaSlice<f32>, // [n_windows * MAX_TRADES_PER_WINDOW]
|
||
per_trade_holding_bars_buf: CudaSlice<u32>, // [n_windows * MAX_TRADES_PER_WINDOW]
|
||
per_trade_bar_index_buf: CudaSlice<u32>, // [n_windows * MAX_TRADES_PER_WINDOW]
|
||
per_trade_dir_mag_buf: CudaSlice<u32>, // [n_windows * MAX_TRADES_PER_WINDOW]
|
||
per_trade_count_buf: CudaSlice<u32>, // [n_windows]
|
||
/// SP21 T2.2 Phase 1.5 (2026-05-11) — predicted Q-value at trade open.
|
||
/// 6th SoA buffer in the per-trade tape, parallel layout to
|
||
/// `per_trade_pnl_buf`. Written by both backtest kernels at trade
|
||
/// close (the `pre_entry_q` snapshot read from `entry_q_state_buf`).
|
||
/// Consumed host-side by `read_per_trade_tape()` and feeds E1
|
||
/// (`compute_q_correction`: predicted_q − pnl bias).
|
||
per_trade_predicted_q_buf: CudaSlice<f32>, // [n_windows * MAX_TRADES_PER_WINDOW]
|
||
/// SP21 T2.2 Phase 1.5 (2026-05-11) — per-window persistent storage
|
||
/// of the predicted Q-value at the OPEN of the currently-active
|
||
/// trade. Written by both backtest kernels on fresh open or
|
||
/// reversal (kernel reads `q_values_per_window[w * num_actions +
|
||
/// action_taken]`); read at the next close to populate the
|
||
/// per-trade tape's `predicted_q` field. Single float per window
|
||
/// — `entry_q` is per-trade, not per-bar, so one live slot is
|
||
/// enough at any time (a closed trade's entry_q is already on
|
||
/// the tape; the slot is overwritten on the next open).
|
||
/// Zeroed in `reset_evaluation_state()` so cross-fold trades
|
||
/// don't bleed stale entry_q values.
|
||
entry_q_state_buf: CudaSlice<f32>, // [n_windows]
|
||
/// Intent magnitude history — pure-policy magnitude argmax from each
|
||
/// eval step, recorded BEFORE the Hold/Flat dir_idx forcing in the
|
||
/// action-select kernel and BEFORE Kelly/margin caps modify exposure
|
||
/// inside env_step. Parallel diagnostic layout to `actions_history_buf`:
|
||
/// `[n_windows * max_len]`, indexed `w * max_len + step`. Populated
|
||
/// only when `evaluate_dqn_graphed` runs — stays zero otherwise.
|
||
/// `None` until `ensure_action_select_ready` allocates it.
|
||
intent_mag_buf: Option<CudaSlice<i32>>,
|
||
|
||
/// Persistent per-step history of the policy's RAW factored action picks
|
||
/// (PRE-env_step). `[n_windows, max_len]` window-major, parallel layout to
|
||
/// `actions_history_buf`. Filled by scattering `chunked_actions_buf` into
|
||
/// the corresponding step slots after each chunk's `experience_action_select`
|
||
/// (reuses the existing `scatter_intent_chunk` kernel — same scatter
|
||
/// semantics, different source/destination buffers). Diagnostic-only:
|
||
/// pairs with `actions_history_buf` (POST-env_step `actual_dir`) so the
|
||
/// reader can compare the kernel's intent vs the env's realisation
|
||
/// across the full val window — localises whether eval-mode collapse is
|
||
/// upstream (kernel produces collapsed picks) or downstream (env_step
|
||
/// drains active picks to Flat). Reset to zero by `reset_evaluation_state`.
|
||
picked_action_history_buf: Option<CudaSlice<i32>>,
|
||
|
||
// Gather kernel output buffer (overwritten every step)
|
||
states_buf: CudaSlice<f32>, // [n_windows * state_dim] (8-aligned, f32 for cublasSgemm)
|
||
|
||
// Output buffer (written by metrics kernel) — mapped pinned host memory
|
||
// visible to the GPU through `dev_ptr` (`cuMemHostAlloc DEVICEMAP`). The
|
||
// kernel writes through the device alias; the host reads via volatile
|
||
// load on `host_ptr` after stream synchronisation. No DtoH copy — the
|
||
// only allowed CPU↔CPU path per
|
||
// `feedback_no_htod_htoh_only_mapped_pinned.md`. Layout `[n_windows * 13]`
|
||
// matches the kernel's flat per-window stride after SP15 Phase 1.1.b
|
||
// (sharpe split out into the dedicated `sharpe_per_bar_kernel`; see
|
||
// `sharpe_buf` below).
|
||
metrics_buf: MappedF32Buffer,
|
||
|
||
/// SP15 Phase 1.1.b (2026-05-06): per-window output of the dedicated
|
||
/// `sharpe_per_bar_kernel` (`launch_sp15_sharpe_per_bar`). Layout
|
||
/// `[n_windows * 3]`: `(mean, std, raw_sharpe)` per window. Host-side
|
||
/// annualisation in `consume_metrics_after_event` multiplies
|
||
/// `raw_sharpe` by `annualization_factor` to produce
|
||
/// `WindowMetrics.sharpe`, preserving the value contract from the
|
||
/// pre-split fused kernel (which applied the same factor in-kernel).
|
||
/// Mapped-pinned per `feedback_no_htod_htoh_only_mapped_pinned.md`.
|
||
sharpe_buf: MappedF32Buffer,
|
||
|
||
// ── SP15 Wave 3b (2026-05-06) — val-cost-streams buffers ─────────────────
|
||
//
|
||
// Three input LobBar SoA streams populated at construction time from
|
||
// `window_lob_bars`, three derivation outputs filled by
|
||
// `launch_sp15_position_history_derivation` after the eval rollout, and
|
||
// five per-window `[mean, std, raw_sharpe]` outputs from the cost-net
|
||
// sharpe + 4 constant-policy baselines. All mapped-pinned per
|
||
// `feedback_no_htod_htoh_only_mapped_pinned.md`. Layout matches the
|
||
// 1.1.b sharpe_per_bar pattern (3 f32 slots per window for the metric
|
||
// outputs; flat `[n_windows * max_len]` for the stream inputs +
|
||
// derivation outputs).
|
||
|
||
/// LobBar.price stream `[n_windows * max_len]` — feeds the 4 baselines.
|
||
/// Populated at construct time via `write_from_slice` from the
|
||
/// `window_lob_bars` parameter. Read-only thereafter.
|
||
close_prices_buf: MappedF32Buffer,
|
||
|
||
/// LobBar.spread / 2 stream `[n_windows * max_len]` — feeds cost-net
|
||
/// sharpe + 4 baselines (per spec §6.2 the cost-net kernel multiplies
|
||
/// `half_spread × |position| × side_ind` so the +entry+exit pair
|
||
/// accumulates one full spread per round-trip). Populated at construct
|
||
/// time. Read-only thereafter.
|
||
half_spread_buf: MappedF32Buffer,
|
||
|
||
/// LobBar.ofi stream `[n_windows * max_len]` — feeds cost-net sharpe.
|
||
/// PPO hyperopt path lacks per-bar OFI features so populates this with
|
||
/// 0.0 for all bars (cost-net OFI-impact term degrades to 0 on PPO;
|
||
/// commission + half-spread × position still apply); DQN adapter and
|
||
/// val_evaluator construct with the real per-bar OFI value.
|
||
ofi_scalar_buf: MappedF32Buffer,
|
||
|
||
/// SP15 Wave 3b — derivation output `[n_windows * max_len]` populated
|
||
/// post-rollout by `launch_sp15_position_history_derivation`. Per-bar
|
||
/// derived position ∈ {-1.0, 0.0, +1.0} reconstructed from the recorded
|
||
/// factored-action history. Consumed by cost-net sharpe.
|
||
position_history_buf: MappedF32Buffer,
|
||
|
||
/// SP15 Wave 3b — derivation output `[n_windows * max_len]`. Per-bar
|
||
/// 1.0 when position changed from prev bar (entry OR exit), else 0.0.
|
||
/// Consumed by cost-net sharpe (`half_spread × |pos| × side_ind` and
|
||
/// `ofi_lambda × |pos| × |ofi| × side_ind` per spec §6.2).
|
||
side_ind_buf: MappedF32Buffer,
|
||
|
||
/// SP15 Wave 3b — derivation output `[n_windows * max_len]`. Per-bar
|
||
/// 1.0 when transitioning to flat from a non-flat position
|
||
/// (round-trip close), else 0.0. Consumed by cost-net sharpe
|
||
/// (`commission_per_rt × rt_ind` per spec §6.2).
|
||
rt_ind_buf: MappedF32Buffer,
|
||
|
||
/// SP15 Wave 3b — per-window `[mean_pnl_net, std, raw_sharpe_cost_net]`
|
||
/// emitted by `launch_sp15_cost_net_sharpe`. Annualised host-side in
|
||
/// `consume_metrics_after_event` to produce `WindowMetrics.sharpe_cost_net`.
|
||
/// Layout `[n_windows * 3]`. Mapped-pinned.
|
||
cost_net_sharpe_buf: MappedF32Buffer,
|
||
|
||
/// SP15 Wave 3b — per-window `[mean, std, raw_sharpe]` from the buyhold
|
||
/// counterfactual baseline (constant +1 position). Annualised host-side.
|
||
/// Layout `[n_windows * 3]`.
|
||
baseline_buyhold_sharpe_buf: MappedF32Buffer,
|
||
|
||
/// SP15 Wave 3b — per-window `[mean, std, raw_sharpe]` from the
|
||
/// hold-only baseline (action=Hold every bar). Annualised host-side.
|
||
/// Layout `[n_windows * 3]`.
|
||
baseline_hold_only_sharpe_buf: MappedF32Buffer,
|
||
|
||
/// SP15 Wave 3b — per-window `[mean, std, raw_sharpe]` from the
|
||
/// naive-momentum baseline (`sign(prices[i] - prices[i-1])`).
|
||
/// Annualised host-side. Layout `[n_windows * 3]`.
|
||
baseline_momentum_sharpe_buf: MappedF32Buffer,
|
||
|
||
/// SP15 Wave 3b — per-window `[mean, std, raw_sharpe]` from the
|
||
/// naive-reversion baseline (mirror of momentum). Annualised host-side.
|
||
/// Layout `[n_windows * 3]`.
|
||
baseline_reversion_sharpe_buf: MappedF32Buffer,
|
||
|
||
/// SP15 Wave 3b — host-precomputed flat per-roundtrip commission
|
||
/// (dollars per round-trip). Per D3 resolution: derived from the
|
||
/// hyperopt config as `tx_cost_bps × 0.0001 × initial_capital`; the
|
||
/// kernel charges this once per round-trip close (rt_ind=1). Stored
|
||
/// on the evaluator (not in `GpuBacktestConfig`) so the per-window
|
||
/// launch loop reads a stable f32 without re-computing each call.
|
||
commission_per_rt: f32,
|
||
|
||
// ── Async eval pipeline (perf-fix 2026-04-28) ─────────────────────────────
|
||
//
|
||
// Splits eval-stream readback from host-side wait. `launch_metrics_and_record_event`
|
||
// submits the metrics kernel + DtoD-async copies of the per-step action history
|
||
// buffers into the mapped-pinned mirrors below, then records `eval_done_event`.
|
||
// `consume_metrics_after_event` calls `event.synchronize()` (the ONLY host wait
|
||
// per epoch boundary) and reads through `host_ptr`. The split lets the host
|
||
// CPU thread launch the next epoch's training kernels while eval drains —
|
||
// restoring the pipelining the dedicated `validation_stream` is supposed to
|
||
// enable. Original async-diag commits (d9cb14f1b + 673b04a8d) wired the GPU-
|
||
// side cross-stream event correctly but left the host blocked inside
|
||
// `synchronize()` after the metrics launch, so pipelining never engaged.
|
||
|
||
/// Recorded on `self.stream` after the metrics kernel + action-history DtoD
|
||
/// copies are queued. Lazy-allocated on first launch (the evaluator is
|
||
/// constructed before the CUDA context is available for `new_event` here in
|
||
/// some early-init paths, so we defer creation). Reused across epochs —
|
||
/// `record` re-records into the same event handle.
|
||
eval_done_event: Option<CudaEvent>,
|
||
|
||
/// Mapped-pinned mirror of `actions_history_buf` ([n_windows * max_len] i32).
|
||
/// Filled via `memcpy_dtod_async` on `self.stream` from `actions_history_buf`
|
||
/// inside `launch_metrics_and_record_event`. Read by the per-direction and
|
||
/// per-magnitude distribution helpers after `consume_metrics_after_event`
|
||
/// has synchronised the event. Replaces the synchronous `memcpy_dtoh` into
|
||
/// a freshly-allocated `Vec<i32>` per call — that path violates
|
||
/// `feedback_no_htod_htoh_only_mapped_pinned.md` and forced a host wait
|
||
/// inside the eval flow.
|
||
actions_history_pinned: MappedI32Buffer,
|
||
|
||
/// Mapped-pinned mirror of `intent_mag_buf` ([n_windows * max_len] i32).
|
||
/// Lazy-allocated alongside `intent_mag_buf` in `ensure_action_select_ready`.
|
||
/// Same DtoD-async-on-eval-stream pattern as `actions_history_pinned`.
|
||
intent_mag_pinned: Option<MappedI32Buffer>,
|
||
|
||
/// Mapped-pinned mirror of `picked_action_history_buf` ([n_windows * max_len] i32).
|
||
/// Lazy-allocated alongside `picked_action_history_buf` in
|
||
/// `ensure_action_select_ready`. Same pattern.
|
||
picked_action_history_pinned: Option<MappedI32Buffer>,
|
||
|
||
// Dimensions and config
|
||
n_windows: usize,
|
||
max_len: usize,
|
||
feature_dim: usize,
|
||
/// Portfolio feature dimension used during construction (always 3).
|
||
portfolio_dim: usize,
|
||
/// Aligned state dimension (feature_dim + portfolio_dim, 8-aligned for tensor cores).
|
||
state_dim: usize,
|
||
config: GpuBacktestConfig,
|
||
/// Precomputed annualization factor: sqrt(bars_per_day * 252).
|
||
annualization_factor: f32,
|
||
/// ISV signals device pointer for adaptive hold. Initialised to point at
|
||
/// `default_isv_buf` (zero-filled). Production training overrides via
|
||
/// `set_isv_signals_ptr` to share the master training-ISV bus; closure-
|
||
/// path callers (eval-baseline) inherit the safe zero-default. Pre-Phase
|
||
/// 8.7 this was `0` (null) — every kernel that read `isv[slot]` OOBed in
|
||
/// the closure path (CUDA_ERROR_ILLEGAL_ADDRESS at
|
||
/// `cost_net_sharpe_kernel+0x90` reading slot 407 → byte 1628 = 0x65c).
|
||
isv_signals_ptr: u64,
|
||
|
||
/// Default zero-filled ISV buffer owned by this evaluator. Backs
|
||
/// `isv_signals_ptr` when no external ISV is wired (closure-path eval).
|
||
/// Sized at `ISV_TOTAL_DIM=536` f32 to cover all current ISV slot
|
||
/// indices. Zero-init means `ofi_lambda=0 → c_ofi=0` in cost-net sharpe
|
||
/// (degraded but valid — matches the "no MBP-10 / no OFI in eval"
|
||
/// semantics anyway).
|
||
default_isv_buf: CudaSlice<f32>,
|
||
|
||
/// Actions buffer for forward kernel output (reused across steps).
|
||
forward_actions_buf: CudaSlice<i32>,
|
||
|
||
// ── cuBLAS DQN forward pass state (lazy-initialised) ────────────────
|
||
|
||
/// `experience_action_select` kernel for greedy argmax.
|
||
action_select_kernel: Option<CudaFunction>,
|
||
|
||
/// Branch sizes for factored action decoding (from DqnBacktestConfig).
|
||
b0_size: i32,
|
||
b1_size: i32,
|
||
b2_size: i32,
|
||
b3_size: i32,
|
||
|
||
/// Q-value gap threshold for trade conviction filter (0.0 = disabled).
|
||
q_gap_threshold: f32,
|
||
|
||
/// Chunked Q-gap output buffer (chunked path: [n_windows * CHUNK_SIZE]).
|
||
chunked_q_gaps_buf: Option<CudaSlice<f32>>,
|
||
/// Per-sample direction-branch conviction ∈ [0, 1], chunked layout
|
||
/// `[chunk_len * n_windows]` matching `chunked_actions_buf`. Written by
|
||
/// `experience_action_select` each chunk; consumed by `backtest_env_step_batch`
|
||
/// as the adaptive Kelly warmup-floor signal (via `unified_env_step_core`).
|
||
chunked_conviction_buf: Option<CudaSlice<f32>>,
|
||
/// Per-sample magnitude-branch conviction ∈ [0, 1], chunked layout matching
|
||
/// `chunked_conviction_buf`. Parallel signal threaded through the same
|
||
/// kernels for the Kelly cap's `max(dir_conv, mag_conv)` composer — lets
|
||
/// the cap relax for confident magnitude picks even when direction
|
||
/// conviction is low (typical at smoke horizon when q_dir values cluster
|
||
/// while the magnitude head sharply prefers Full).
|
||
chunked_magnitude_conviction_buf: Option<CudaSlice<f32>>,
|
||
|
||
/// Cached CUDA graph for `evaluate_dqn_graphed()` step loop.
|
||
///
|
||
/// Captured on first call, replayed on subsequent calls with the same weights
|
||
/// and evaluator geometry. The graph encodes the full unrolled step loop
|
||
/// (gather + forward + DtoD + env_step for each step 0..max_len).
|
||
/// Because all buffer pointers are owned by `self` and remain stable,
|
||
/// the graph is valid as long as the evaluator exists and the same
|
||
/// weight buffers are used.
|
||
dqn_graph: Option<SendSyncGraph>,
|
||
|
||
// ── Chunked forward pass state (lazy-initialised) ─────────────────────
|
||
|
||
/// Q-values buffer for chunked C51 expected-Q (n_windows * CHUNK_SIZE rows).
|
||
chunked_q_values: Option<CudaSlice<f32>>,
|
||
|
||
/// Branch-major raw C51 logits buffer for the chunked forward pass:
|
||
/// `[n_windows * CHUNK_SIZE, total_actions * num_atoms]`. Populated each
|
||
/// chunk by `QValueProvider::compute_q_and_b_logits_to`, which DtoD-copies
|
||
/// the trainer's `on_b_logits_buf`. The direction-branch slice
|
||
/// (`[N, B0, num_atoms]` at the start of each row) is consumed by the
|
||
/// Thompson direction sampler in `experience_action_select`
|
||
/// (Plan C T2 amendment, T4 evaluator wire-up 2026-04-29).
|
||
chunked_b_logits_buf: Option<CudaSlice<f32>>,
|
||
|
||
/// SP17 Commit C: chunked V-head logits buffer for action_select. Layout:
|
||
/// `[n_windows * CHUNK_SIZE, num_atoms]` sample-major. Populated each chunk
|
||
/// by `QValueProvider::compute_q_and_b_logits_to`, which DtoD-copies the
|
||
/// trainer's `on_v_logits_buf`. Consumed by Thompson direction selection
|
||
/// in `experience_action_select` so action-selection sees the SAME
|
||
/// dueling-decomposed distribution `softmax(V + A_centered)` that the
|
||
/// backward path uses. No NULL fallback — kernel hard-requires V.
|
||
chunked_v_logits_buf: Option<CudaSlice<f32>>,
|
||
|
||
/// Batched states buffer for chunked gather: [n_windows * CHUNK_SIZE, state_dim].
|
||
chunked_states_buf: Option<CudaSlice<f32>>,
|
||
|
||
/// SP21 T2.2 Phase 6.5 (2026-05-11) — full-window state retention
|
||
/// for hindsight synthetic injection. Layout
|
||
/// `[max_len × n_windows × state_dim_padded]` flattened with the
|
||
/// SAME `[step, window, dim]` ordering as `chunked_states_buf`
|
||
/// uses within each chunk — each chunk's `chunked_states_buf` is
|
||
/// DtoD-copied into this buffer at offset `chunk_start × n_windows
|
||
/// × state_dim_padded` immediately after every `launch_gather_chunk`.
|
||
///
|
||
/// **Mapped-pinned** (`cuMemHostAlloc DEVICEMAP`) per
|
||
/// `feedback_no_htod_htoh_only_mapped_pinned`: the DtoD copy
|
||
/// writes through the device-aliased pointer; host-side reads
|
||
/// via `host_ptr` are zero-copy (no `memcpy_dtoh`). Consumed by
|
||
/// `training_loop::inject_hindsight_experiences` which synthesizes
|
||
/// replay tuples from `EnrichmentResult.hindsight` and pushes them
|
||
/// into PER via `GpuReplayBuffer::insert_batch` — the host read
|
||
/// at `(window_idx, bar_idx)` happens AFTER an eval stream sync,
|
||
/// so mapped-pinned coherence guarantees the kernel writes are
|
||
/// visible.
|
||
///
|
||
/// Memory cost at production cfg (max_len=200_000, n_windows=5,
|
||
/// state_dim_padded≈128): 5 × 200_000 × 128 × 4 ≈ 512 MB of
|
||
/// pinned (non-swappable) host RAM. Substantial but feasible on
|
||
/// the L40S host (192 GB+). Single-allocation cost per evaluator
|
||
/// (not per-step). Lives for the trainer's lifetime; not part of
|
||
/// the CUDA Graph capture (the DtoD copies run on the eval stream
|
||
/// outside the train graph).
|
||
retained_states_buf: Option<MappedF32Buffer>,
|
||
|
||
/// Batched forward actions buffer for chunked action select: [n_windows * CHUNK_SIZE].
|
||
chunked_actions_buf: Option<CudaSlice<i32>>,
|
||
|
||
/// Chunked intent-magnitude buffer mirroring `chunked_actions_buf`:
|
||
/// `[chunk_len, n_windows]` step-major layout (thread `i = s*n + w`).
|
||
/// Filled by `experience_action_select` on each chunk, then scattered
|
||
/// into the window-major `intent_mag_buf` by `scatter_intent_kernel`.
|
||
chunked_intent_mag_buf: Option<CudaSlice<i32>>,
|
||
|
||
/// `scatter_intent_chunk` kernel handle — copies `chunked_intent_mag_buf`
|
||
/// into `intent_mag_buf` at the current chunk's step range.
|
||
scatter_intent_kernel: Option<CudaFunction>,
|
||
|
||
/// Val plan_isv parity kernels (task #94, 2026-04-24).
|
||
/// Lazy-loaded with the other action-select machinery.
|
||
plan_state_isv_kernel: Option<CudaFunction>,
|
||
plan_diag_reduce_kernel: Option<CudaFunction>,
|
||
|
||
/// D.8 TLOB val-path: OFI[32]→TLOB[16] attention applied after each gather step.
|
||
///
|
||
/// Weights are synchronised from the training TLOB before `evaluate_dqn_graphed`
|
||
/// via `set_tlob_from_training`. `None` when TLOB is disabled or weight sync
|
||
/// has not yet been called.
|
||
tlob: Option<super::gpu_tlob::GpuTlob>,
|
||
}
|
||
|
||
impl Drop for GpuBacktestEvaluator {
|
||
fn drop(&mut self) {
|
||
// Synchronize forked stream before any CudaSlice fields drop.
|
||
// Without this, pending GPU work (graph replay, cuBLAS, kernel launches)
|
||
// may reference buffers that cudarc's CudaSlice::drop frees immediately.
|
||
#[allow(unsafe_code)]
|
||
unsafe {
|
||
cudarc::driver::sys::cuStreamSynchronize(self.stream.cu_stream());
|
||
}
|
||
// Destroy CUDA graph before buffer drops — graph references
|
||
// internal device pointers that become invalid after CudaSlice drops.
|
||
self.dqn_graph = None;
|
||
}
|
||
}
|
||
|
||
impl GpuBacktestEvaluator {
|
||
// ── Constructor ───────────────────────────────────────────────────────────
|
||
|
||
/// Create evaluator and upload window data to GPU.
|
||
///
|
||
/// `window_prices`: one `Vec<[f32; 4]>` (OHLC rows) per window.
|
||
/// `window_features`: one `Vec<Vec<f32>>` (feature rows) per window.
|
||
/// `window_lob_bars`: one `Vec<LobBar>` per window — feeds the SP15
|
||
/// Wave 3b cost-net sharpe + 4 counterfactual baselines via three
|
||
/// mapped-pinned SoA streams (close prices, half-spread, OFI scalar).
|
||
/// Per spec §4.4: bars must be aligned 1:1 with `window_prices` and
|
||
/// `window_features`. Shorter `window_lob_bars[w]` is zero-padded to
|
||
/// `max_len`. PPO hyperopt path passes `ofi=0.0` for all bars (per
|
||
/// D4 resolution: PPO lacks per-bar OFI features; cost-net OFI
|
||
/// impact term degrades to 0 — commission + half-spread × position
|
||
/// still apply). DQN adapter + val_evaluator pass real OFI.
|
||
/// Shorter windows are zero-padded to `max_len`.
|
||
pub fn new(
|
||
window_prices: &[Vec<[f32; 4]>],
|
||
window_features: &[Vec<Vec<f32>>],
|
||
window_lob_bars: &[Vec<LobBar>],
|
||
feature_dim: usize,
|
||
config: GpuBacktestConfig,
|
||
parent_stream: &Arc<CudaStream>,
|
||
) -> Result<Self, MLError> {
|
||
let n_windows = window_prices.len();
|
||
if n_windows == 0 {
|
||
return Err(MLError::ConfigError("No windows provided".to_owned()));
|
||
}
|
||
if window_lob_bars.len() != n_windows {
|
||
return Err(MLError::ConfigError(format!(
|
||
"window_lob_bars len {} != n_windows {n_windows}",
|
||
window_lob_bars.len()
|
||
)));
|
||
}
|
||
|
||
let max_len = window_prices.iter().map(|w| w.len()).max().unwrap_or(0);
|
||
if max_len == 0 {
|
||
return Err(MLError::ConfigError("All windows are empty".to_owned()));
|
||
}
|
||
|
||
// ── Apply leverage cap to max_position ────────────────────────────
|
||
let mut config = config;
|
||
if config.max_leverage > 0.0 {
|
||
// Use median close price across ALL windows for robust leverage cap.
|
||
// Single-window median can be skewed if first window is atypical.
|
||
let mut all_closes: Vec<f32> = window_prices
|
||
.iter()
|
||
.flat_map(|w| {
|
||
let mid = w.len() / 2;
|
||
w.get(mid).map(|ohlc| ohlc[3])
|
||
})
|
||
.filter(|p| *p > 0.0)
|
||
.collect();
|
||
all_closes.sort_unstable_by(|a, b| a.partial_cmp(b).unwrap_or(std::cmp::Ordering::Equal));
|
||
let ref_price = if all_closes.is_empty() {
|
||
0.0
|
||
} else {
|
||
all_closes[all_closes.len() / 2]
|
||
};
|
||
let effective = config.effective_max_position(ref_price);
|
||
if effective < config.max_position {
|
||
tracing::info!(
|
||
"Leverage cap: max_position {:.2} → {:.4} (price={:.0}, multiplier={}, leverage={}×, capital={:.0})",
|
||
config.max_position, effective, ref_price, config.contract_multiplier,
|
||
config.max_leverage, config.initial_capital
|
||
);
|
||
config.max_position = effective;
|
||
}
|
||
}
|
||
|
||
// ── Flatten prices (zero-pad shorter windows) ─────────────────────
|
||
let window_lens: Vec<i32> = window_prices.iter().map(|w| w.len() as i32).collect();
|
||
|
||
let mut flat_prices = vec![0.0_f32; n_windows * max_len * 4];
|
||
for (w, prices) in window_prices.iter().enumerate() {
|
||
for (t, ohlc) in prices.iter().enumerate() {
|
||
let base = (w * max_len + t) * 4;
|
||
flat_prices[base..base + 4].copy_from_slice(ohlc);
|
||
}
|
||
}
|
||
|
||
// ── Flatten features (zero-pad shorter windows / truncate wider) ──
|
||
let mut flat_features = vec![0.0_f32; n_windows * max_len * feature_dim];
|
||
for (w, feats) in window_features.iter().enumerate() {
|
||
for (t, fv) in feats.iter().enumerate() {
|
||
let base = (w * max_len + t) * feature_dim;
|
||
let copy_len = fv.len().min(feature_dim);
|
||
flat_features[base..base + copy_len].copy_from_slice(&fv[..copy_len]);
|
||
}
|
||
}
|
||
|
||
// ── CUDA stream ──────────────────────────────────────────────────
|
||
// Fork a dedicated stream for CUDA Graph capture.
|
||
// The default stream does NOT support begin_capture (returns
|
||
// CUDA_ERROR_STREAM_CAPTURE_UNSUPPORTED). Forking creates a non-default
|
||
// stream that supports capture and all async operations.
|
||
let stream = parent_stream.fork().map_err(|e| {
|
||
MLError::ModelError(format!("main stream fork: {e}"))
|
||
})?;
|
||
let context = stream.context();
|
||
|
||
// ── Compile / cache kernels ───────────────────────────────────────
|
||
let env_module = context
|
||
.load_cubin(ENV_CUBIN.to_vec())
|
||
.map_err(|e| MLError::ModelError(format!("env module load: {e}")))?;
|
||
let env_kernel = env_module
|
||
.load_function("backtest_env_step")
|
||
.map_err(|e| MLError::ModelError(format!("backtest_env_step load: {e}")))?;
|
||
let env_batch_kernel = env_module
|
||
.load_function("backtest_env_step_batch")
|
||
.map_err(|e| MLError::ModelError(format!("backtest_env_step_batch load: {e}")))?;
|
||
|
||
let metrics_module = context
|
||
.load_cubin(METRICS_CUBIN.to_vec())
|
||
.map_err(|e| MLError::ModelError(format!("metrics module load: {e}")))?;
|
||
let metrics_kernel = metrics_module
|
||
.load_function("compute_backtest_metrics")
|
||
.map_err(|e| MLError::ModelError(format!("compute_backtest_metrics load: {e}")))?;
|
||
|
||
let gather_module = context
|
||
.load_cubin(BACKTEST_DQN_CUBIN.to_vec())
|
||
.map_err(|e| MLError::ModelError(format!("gather module load: {e}")))?;
|
||
let gather_kernel = gather_module
|
||
.load_function("backtest_state_gather")
|
||
.map_err(|e| MLError::ModelError(format!("backtest_state_gather load: {e}")))?;
|
||
let gather_chunk_kernel = gather_module
|
||
.load_function("backtest_state_gather_chunk")
|
||
.map_err(|e| MLError::ModelError(format!("backtest_state_gather_chunk load: {e}")))?;
|
||
|
||
// SP15 Wave 5 follow-up (2026-05-07): pre-load the SP15 baseline +
|
||
// cost-net cubins ONCE here, in the evaluator's CUDA context, so
|
||
// the per-window eval hot loop (`for w in 0..self.n_windows` at
|
||
// `~line 2860`) calls `launch_sp15_*` against pre-resolved
|
||
// `&CudaFunction` handles instead of resolving the symbols on
|
||
// every iteration. The on-demand pattern caused
|
||
// `CUDA_ERROR_ILLEGAL_ADDRESS` in hyperopt-trial child stream
|
||
// contexts (workflow `train-xggfc`, 2026-05-07): trial 1 failed
|
||
// mid-eval with `load sp15_baseline_kernels cubin: ILLEGAL_ADDRESS`,
|
||
// and the corrupted context cascaded into trials 2-20 failing at
|
||
// `Fork CUDA stream for trial`. Mirrors `5d63762ab` precedent for
|
||
// `bn_tanh_concat_dd_kernel`.
|
||
let sp15_cost_net_module = context
|
||
.load_cubin(super::gpu_dqn_trainer::SP15_COST_NET_SHARPE_CUBIN.to_vec())
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"sp15_cost_net_sharpe module load: {e}"
|
||
)))?;
|
||
let sp15_cost_net_sharpe_kernel = sp15_cost_net_module
|
||
.load_function("cost_net_sharpe_kernel")
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"cost_net_sharpe_kernel load: {e}"
|
||
)))?;
|
||
let sp15_baseline_module = context
|
||
.load_cubin(super::gpu_dqn_trainer::SP15_BASELINE_KERNELS_CUBIN.to_vec())
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"sp15_baseline_kernels module load: {e}"
|
||
)))?;
|
||
let sp15_baseline_buyhold_kernel = sp15_baseline_module
|
||
.load_function("baseline_buyhold_kernel")
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"baseline_buyhold_kernel load: {e}"
|
||
)))?;
|
||
let sp15_baseline_hold_only_kernel = sp15_baseline_module
|
||
.load_function("baseline_hold_only_kernel")
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"baseline_hold_only_kernel load: {e}"
|
||
)))?;
|
||
let sp15_baseline_naive_momentum_kernel = sp15_baseline_module
|
||
.load_function("baseline_naive_momentum_kernel")
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"baseline_naive_momentum_kernel load: {e}"
|
||
)))?;
|
||
let sp15_baseline_naive_reversion_kernel = sp15_baseline_module
|
||
.load_function("baseline_naive_reversion_kernel")
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"baseline_naive_reversion_kernel load: {e}"
|
||
)))?;
|
||
|
||
// ── Upload read-only data via mapped-pinned (no HtoD, no DtoD) ────
|
||
// Direct MappedF32/I32Buffer storage: host writes payload via
|
||
// `host_ptr.write_from_slice`, kernels later read via `dev_ptr`.
|
||
// No memcpy_dtod_async, no stream.synchronize() at construction —
|
||
// mapped-pinned coherence guarantees the kernel sees the host
|
||
// writes after the next stream-sync barrier.
|
||
let prices_buf = unsafe { MappedF32Buffer::new(flat_prices.len()) }
|
||
.map_err(|e| MLError::ModelError(format!("prices mapped-pinned alloc: {e}")))?;
|
||
prices_buf.write_from_slice(&flat_prices);
|
||
let features_buf = unsafe { MappedF32Buffer::new(flat_features.len()) }
|
||
.map_err(|e| MLError::ModelError(format!("features mapped-pinned alloc: {e}")))?;
|
||
features_buf.write_from_slice(&flat_features);
|
||
let window_lens_buf = unsafe { MappedI32Buffer::new(window_lens.len()) }
|
||
.map_err(|e| MLError::ModelError(format!("window_lens mapped-pinned alloc: {e}")))?;
|
||
window_lens_buf.write_from_slice(&window_lens);
|
||
|
||
// ── Allocate mutable state buffers ────────────────────────────────
|
||
let portfolio_init = Self::init_portfolio_state(n_windows, config.initial_capital);
|
||
let portfolio_buf = unsafe { MappedF32Buffer::new(portfolio_init.len()) }
|
||
.map_err(|e| MLError::ModelError(format!("portfolio mapped-pinned alloc: {e}")))?;
|
||
portfolio_buf.write_from_slice(&portfolio_init);
|
||
|
||
// Kelly stats buffer — zero-initialized (no trades yet). The Kelly
|
||
// priors + warmup floor in trade_physics.cuh handle the cold start
|
||
// gracefully: during the first ~10 completed trades the cap is
|
||
// dominated by the 50% warmup floor, then transitions to pure
|
||
// data-driven Kelly.
|
||
let kelly_stats_buf = stream
|
||
.alloc_zeros::<f32>(n_windows * BACKTEST_KELLY_STATS_SIZE)
|
||
.map_err(|e| MLError::ModelError(format!("kelly_stats alloc: {e}")))?;
|
||
|
||
// SP21 T2.2 Phase 1 (2026-05-10) — per-trade tape device buffers.
|
||
// Sized at MAX_TRADES_PER_WINDOW per window (3.2MB per window for
|
||
// the 4 SoA buffers + count). Zero-init: per_trade_count[w]=0 so
|
||
// the first close in each window writes at slot 0. `reset_per_trade_tape`
|
||
// re-zeros at the start of each eval window so trade indices don't
|
||
// accumulate across folds.
|
||
let per_trade_pnl_buf = stream
|
||
.alloc_zeros::<f32>(n_windows * MAX_TRADES_PER_WINDOW)
|
||
.map_err(|e| MLError::ModelError(format!("per_trade_pnl alloc: {e}")))?;
|
||
let per_trade_holding_bars_buf = stream
|
||
.alloc_zeros::<u32>(n_windows * MAX_TRADES_PER_WINDOW)
|
||
.map_err(|e| MLError::ModelError(format!("per_trade_holding_bars alloc: {e}")))?;
|
||
let per_trade_bar_index_buf = stream
|
||
.alloc_zeros::<u32>(n_windows * MAX_TRADES_PER_WINDOW)
|
||
.map_err(|e| MLError::ModelError(format!("per_trade_bar_index alloc: {e}")))?;
|
||
let per_trade_dir_mag_buf = stream
|
||
.alloc_zeros::<u32>(n_windows * MAX_TRADES_PER_WINDOW)
|
||
.map_err(|e| MLError::ModelError(format!("per_trade_dir_mag alloc: {e}")))?;
|
||
// SP21 T2.2 Phase 1.5 (2026-05-11) — predicted-Q tape (6th SoA
|
||
// buffer) + per-window persistent entry_q state. Memory cost
|
||
// negligible vs the existing 4 buffers; same lifecycle.
|
||
let per_trade_predicted_q_buf = stream
|
||
.alloc_zeros::<f32>(n_windows * MAX_TRADES_PER_WINDOW)
|
||
.map_err(|e| MLError::ModelError(format!("per_trade_predicted_q alloc: {e}")))?;
|
||
let entry_q_state_buf = stream
|
||
.alloc_zeros::<f32>(n_windows)
|
||
.map_err(|e| MLError::ModelError(format!("entry_q_state alloc: {e}")))?;
|
||
let per_trade_count_buf = stream
|
||
.alloc_zeros::<u32>(n_windows)
|
||
.map_err(|e| MLError::ModelError(format!("per_trade_count alloc: {e}")))?;
|
||
|
||
// plan_isv_buf: [N, SL_PORTFOLIO_PLAN_DIM=7] — live plan_isv per window.
|
||
// Populated each step by backtest_plan_state_isv (val plan_isv parity,
|
||
// task #94) and consumed by the next step's backtest_state_gather.
|
||
// Zero-initialized so the first step (no active plan) emits zeros for
|
||
// all seven plan_isv slots (remaining_fraction included).
|
||
let plan_isv_buf = stream
|
||
.alloc_zeros::<f32>(n_windows * 7)
|
||
.map_err(|e| MLError::ModelError(format!("plan_isv alloc: {e}")))?;
|
||
|
||
// Val plan_isv parity (task #94, 2026-04-24). Three persistent
|
||
// buffers carrying the plan state machinery that replicates
|
||
// training's ps[23..29] + plan_params logic in the val backtest:
|
||
// plan_params_buf[N, 6] — live plan MLP output each step
|
||
// plan_state_buf[N, 7] — stored plan at entry, stable across
|
||
// bars within a trade
|
||
// plan_diag_buf[8] — epoch-end diagnostic reduction
|
||
let plan_params_buf = stream
|
||
.alloc_zeros::<f32>(n_windows * 6)
|
||
.map_err(|e| MLError::ModelError(format!("plan_params alloc: {e}")))?;
|
||
let plan_state_buf = stream
|
||
.alloc_zeros::<f32>(n_windows * 7)
|
||
.map_err(|e| MLError::ModelError(format!("plan_state alloc: {e}")))?;
|
||
// Mapped-pinned: kernel writes via `dev_ptr`, host reads via `host_ptr`
|
||
// after stream sync — zero DtoH copy. Caller must hold the CUDA context.
|
||
let plan_diag_buf = unsafe {
|
||
MappedF32Buffer::new(8)
|
||
.map_err(|e| MLError::ModelError(format!("plan_diag mapped-pinned alloc: {e}")))?
|
||
};
|
||
|
||
let step_rewards_buf = stream
|
||
.alloc_zeros::<f32>(n_windows)
|
||
.map_err(|e| MLError::ModelError(format!("step_rewards alloc: {e}")))?;
|
||
let step_returns_buf = stream
|
||
.alloc_zeros::<f32>(n_windows * max_len)
|
||
.map_err(|e| MLError::ModelError(format!("step_returns alloc: {e}")))?;
|
||
let done_buf = stream
|
||
.alloc_zeros::<i32>(n_windows)
|
||
.map_err(|e| MLError::ModelError(format!("done alloc: {e}")))?;
|
||
let actions_buf = stream
|
||
.alloc_zeros::<i32>(n_windows)
|
||
.map_err(|e| MLError::ModelError(format!("actions alloc: {e}")))?;
|
||
let actions_history_buf = stream
|
||
.alloc_zeros::<i32>(n_windows * max_len)
|
||
.map_err(|e| MLError::ModelError(format!("actions_history alloc: {e}")))?;
|
||
// Mapped-pinned: metrics_kernel writes through `dev_ptr` and the host
|
||
// reads via `host_ptr` after stream sync. Replaces the prior
|
||
// `clone_dtoh` readback (which violates the no-DtoH rule) with a
|
||
// zero-copy path per `feedback_no_htod_htoh_only_mapped_pinned.md`.
|
||
// SP15 Phase 1.1.b: stride dropped 14 → 13 (sharpe split out into
|
||
// the dedicated `sharpe_per_bar_kernel` consumed via `sharpe_buf`).
|
||
// 2026-05-09 — stride grew 13 → 19 with per-regime WR slots
|
||
// [13..19) (audit follow-up). Layout mirrored in
|
||
// `consume_metrics_after_event` and the kernel's `metrics_out`
|
||
// contract docstring.
|
||
let metrics_buf = unsafe {
|
||
MappedF32Buffer::new(n_windows * 19)
|
||
.map_err(|e| MLError::ModelError(format!("metrics mapped-pinned alloc: {e}")))?
|
||
};
|
||
// SP15 Phase 1.1.b: per-window mean/std/sharpe output for the
|
||
// dedicated `sharpe_per_bar_kernel`. Three f32 slots per window
|
||
// (mean, std, raw_sharpe) populated by N kernel launches (one per
|
||
// window) in `launch_metrics_and_record_event`. Host applies
|
||
// annualisation in `consume_metrics_after_event`.
|
||
let sharpe_buf = unsafe {
|
||
MappedF32Buffer::new(n_windows * 3)
|
||
.map_err(|e| MLError::ModelError(format!("sharpe mapped-pinned alloc: {e}")))?
|
||
};
|
||
// Mapped-pinned mirror for actions_history_buf — populated via
|
||
// `memcpy_dtod_async` on the eval stream before the eval_done_event is
|
||
// recorded. Replaces the per-call synchronous `memcpy_dtoh` into a
|
||
// freshly-allocated host `Vec<i32>`. See the field docstring above for
|
||
// the perf-fix rationale.
|
||
let actions_history_pinned = unsafe {
|
||
MappedI32Buffer::new(n_windows * max_len)
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"actions_history mapped-pinned alloc: {e}"
|
||
)))?
|
||
};
|
||
|
||
// State layout defined by state_layout.cuh — single source of truth.
|
||
// Canonical 96-dim: 42 market + 20 OFI + 16 MTF + 8 portfolio + 6 plan/ISV + 4 padding.
|
||
let state_dim = ml_core::state_layout::STATE_DIM;
|
||
let state_dim_padded = ml_core::state_layout::STATE_DIM_PADDED;
|
||
let states_buf = stream
|
||
.alloc_zeros::<f32>(n_windows * state_dim_padded)
|
||
.map_err(|e| MLError::ModelError(format!("states_buf alloc: {e}")))?;
|
||
|
||
// Pre-allocate forward kernel output buffer (used by evaluate_dqn path)
|
||
let forward_actions_buf = stream
|
||
.alloc_zeros::<i32>(n_windows)
|
||
.map_err(|e| MLError::ModelError(format!("forward_actions alloc: {e}")))?;
|
||
|
||
// Phase 8.7 (2026-05-12) — default zero-filled ISV bus for the
|
||
// closure-path callers (eval-baseline). Pre-fix, `isv_signals_ptr=0`
|
||
// (null) and every kernel reading any `isv[slot]` OOBed at
|
||
// `slot×4` bytes. compute-sanitizer caught it first at
|
||
// `cost_net_sharpe_kernel+0x90` reading slot 407 → 0x65c. Sized at
|
||
// ISV_TOTAL_DIM (536 f32s) to cover all current ISV slot indices.
|
||
// Production training overrides via `set_isv_signals_ptr` to share
|
||
// the master training ISV bus.
|
||
let default_isv_buf = stream
|
||
.alloc_zeros::<f32>(super::gpu_dqn_trainer::ISV_TOTAL_DIM)
|
||
.map_err(|e| MLError::ModelError(format!("default_isv_buf alloc: {e}")))?;
|
||
let default_isv_ptr = {
|
||
let (ptr, _g) = default_isv_buf.device_ptr(&stream);
|
||
ptr
|
||
};
|
||
|
||
// ── SP15 Wave 3b — val-cost-streams buffer alloc + LobBar SoA upload ──
|
||
//
|
||
// Three input streams sourced from `window_lob_bars` (zero-padded to
|
||
// `max_len`), three derivation outputs filled post-rollout by the
|
||
// position-history derivation kernel, and five per-window
|
||
// `[mean, std, raw_sharpe]` outputs (cost-net + 4 baselines).
|
||
// All mapped-pinned per `feedback_no_htod_htoh_only_mapped_pinned.md`.
|
||
|
||
let mut flat_close_prices = vec![0.0_f32; n_windows * max_len];
|
||
let mut flat_half_spread = vec![0.0_f32; n_windows * max_len];
|
||
let mut flat_ofi_scalar = vec![0.0_f32; n_windows * max_len];
|
||
for (w, bars) in window_lob_bars.iter().enumerate() {
|
||
for (t, b) in bars.iter().enumerate().take(max_len) {
|
||
let idx = w * max_len + t;
|
||
flat_close_prices[idx] = b.price;
|
||
flat_half_spread[idx] = b.spread * 0.5;
|
||
flat_ofi_scalar[idx] = b.ofi;
|
||
}
|
||
}
|
||
|
||
let close_prices_buf = unsafe { MappedF32Buffer::new(flat_close_prices.len()) }
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"close_prices mapped-pinned alloc: {e}"
|
||
)))?;
|
||
close_prices_buf.write_from_slice(&flat_close_prices);
|
||
let half_spread_buf = unsafe { MappedF32Buffer::new(flat_half_spread.len()) }
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"half_spread mapped-pinned alloc: {e}"
|
||
)))?;
|
||
half_spread_buf.write_from_slice(&flat_half_spread);
|
||
let ofi_scalar_buf = unsafe { MappedF32Buffer::new(flat_ofi_scalar.len()) }
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"ofi_scalar mapped-pinned alloc: {e}"
|
||
)))?;
|
||
ofi_scalar_buf.write_from_slice(&flat_ofi_scalar);
|
||
|
||
// Derivation outputs — zero-init; populated each
|
||
// `launch_metrics_and_record_event` by
|
||
// `launch_sp15_position_history_derivation`.
|
||
let position_history_buf = unsafe { MappedF32Buffer::new(n_windows * max_len) }
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"position_history mapped-pinned alloc: {e}"
|
||
)))?;
|
||
let side_ind_buf = unsafe { MappedF32Buffer::new(n_windows * max_len) }
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"side_ind mapped-pinned alloc: {e}"
|
||
)))?;
|
||
let rt_ind_buf = unsafe { MappedF32Buffer::new(n_windows * max_len) }
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"rt_ind mapped-pinned alloc: {e}"
|
||
)))?;
|
||
|
||
// Per-window `[mean, std, raw_sharpe]` outputs from the cost-net
|
||
// kernel + 4 counterfactual baselines. Same layout as `sharpe_buf`.
|
||
let cost_net_sharpe_buf = unsafe { MappedF32Buffer::new(n_windows * 3) }
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"cost_net_sharpe mapped-pinned alloc: {e}"
|
||
)))?;
|
||
let baseline_buyhold_sharpe_buf = unsafe { MappedF32Buffer::new(n_windows * 3) }
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"baseline_buyhold_sharpe mapped-pinned alloc: {e}"
|
||
)))?;
|
||
let baseline_hold_only_sharpe_buf = unsafe { MappedF32Buffer::new(n_windows * 3) }
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"baseline_hold_only_sharpe mapped-pinned alloc: {e}"
|
||
)))?;
|
||
let baseline_momentum_sharpe_buf = unsafe { MappedF32Buffer::new(n_windows * 3) }
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"baseline_momentum_sharpe mapped-pinned alloc: {e}"
|
||
)))?;
|
||
let baseline_reversion_sharpe_buf = unsafe { MappedF32Buffer::new(n_windows * 3) }
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"baseline_reversion_sharpe mapped-pinned alloc: {e}"
|
||
)))?;
|
||
|
||
// Host-precomputed flat per-roundtrip commission (per D3 resolution).
|
||
// Formula: `tx_cost_bps × 0.0001 × initial_capital`. The cost-net
|
||
// kernel charges this once per round-trip close (rt_ind=1); the
|
||
// half-spread + OFI-impact components scale with per-bar |position|
|
||
// and are charged via the separate spread/OFI streams. Stored on
|
||
// the evaluator so the per-window launch loop reads a stable f32
|
||
// without re-deriving each call. Mirrors the per-side commission
|
||
// semantics described in `cost_net_sharpe_kernel.cu` spec §6.2.
|
||
let commission_per_rt =
|
||
config.tx_cost_bps * 0.0001_f32 * config.initial_capital;
|
||
|
||
let upload_mb =
|
||
(flat_prices.len() * std::mem::size_of::<f32>()
|
||
+ flat_features.len() * std::mem::size_of::<f32>()) as f64
|
||
/ 1_048_576.0;
|
||
info!(
|
||
"GpuBacktestEvaluator: {} windows × {} max_len × {} features ({:.1} MB uploaded)",
|
||
n_windows, max_len, feature_dim, upload_mb,
|
||
);
|
||
|
||
Ok(Self {
|
||
stream,
|
||
env_kernel,
|
||
env_batch_kernel,
|
||
metrics_kernel,
|
||
gather_kernel,
|
||
gather_chunk_kernel,
|
||
// SP15 Wave 5 follow-up (2026-05-07): pre-loaded handles for
|
||
// SP15 baseline + cost-net launchers, see field-level
|
||
// docstrings in the struct definition.
|
||
sp15_cost_net_sharpe_kernel,
|
||
sp15_baseline_buyhold_kernel,
|
||
sp15_baseline_hold_only_kernel,
|
||
sp15_baseline_naive_momentum_kernel,
|
||
sp15_baseline_naive_reversion_kernel,
|
||
prices_buf,
|
||
features_buf,
|
||
window_lens_buf,
|
||
portfolio_buf,
|
||
kelly_stats_buf,
|
||
plan_isv_buf,
|
||
plan_params_buf,
|
||
plan_state_buf,
|
||
plan_diag_buf,
|
||
step_rewards_buf,
|
||
step_returns_buf,
|
||
done_buf,
|
||
actions_buf,
|
||
actions_history_buf,
|
||
// SP21 T2.2 Phase 1 (2026-05-10) — per-trade tape buffers
|
||
// (see field-decl docs for layout / lifecycle).
|
||
per_trade_pnl_buf,
|
||
per_trade_holding_bars_buf,
|
||
per_trade_bar_index_buf,
|
||
per_trade_dir_mag_buf,
|
||
per_trade_count_buf,
|
||
per_trade_predicted_q_buf,
|
||
entry_q_state_buf,
|
||
intent_mag_buf: None,
|
||
picked_action_history_buf: None,
|
||
states_buf,
|
||
metrics_buf,
|
||
sharpe_buf,
|
||
// SP15 Wave 3b — input streams + derivation outputs + 5 metric outputs.
|
||
close_prices_buf,
|
||
half_spread_buf,
|
||
ofi_scalar_buf,
|
||
position_history_buf,
|
||
side_ind_buf,
|
||
rt_ind_buf,
|
||
cost_net_sharpe_buf,
|
||
baseline_buyhold_sharpe_buf,
|
||
baseline_hold_only_sharpe_buf,
|
||
baseline_momentum_sharpe_buf,
|
||
baseline_reversion_sharpe_buf,
|
||
commission_per_rt,
|
||
// Lazy-init on first launch so the CUDA context is guaranteed
|
||
// active. Reused across epochs (re-record via `event.record(stream)`).
|
||
eval_done_event: None,
|
||
actions_history_pinned,
|
||
// Lazy-allocated alongside the underlying device buffers in
|
||
// `ensure_action_select_ready`. Stay `None` for non-DQN paths
|
||
// (evaluate / evaluate_ppo / evaluate_supervised) which never
|
||
// populate the underlying buffers.
|
||
intent_mag_pinned: None,
|
||
picked_action_history_pinned: None,
|
||
n_windows,
|
||
max_len,
|
||
feature_dim,
|
||
portfolio_dim: ml_core::state_layout::PORTFOLIO_BASE_DIM + ml_core::state_layout::MTF_DIM,
|
||
state_dim,
|
||
// Effective bars-per-day = configured bars-per-day / val_subsample_stride.
|
||
// The metrics kernel computes mean/std over the (already-subsampled)
|
||
// return series; the annualization scalar must therefore reflect the
|
||
// *effective* sampling cadence, not the underlying bar cadence.
|
||
// Per audit `docs/lookahead-bias-audit-2026-04-28.md` §5, using the
|
||
// unscaled bars_per_day inflates reported Sharpe by sqrt(stride).
|
||
annualization_factor: {
|
||
let stride = config.val_subsample_stride.max(1) as f32;
|
||
let effective_bars_per_day = config.bars_per_day / stride;
|
||
(effective_bars_per_day * config.trading_days_per_year).sqrt()
|
||
},
|
||
// Phase 8.7 (2026-05-12) — point at the zero-filled default ISV
|
||
// buffer. Production training overrides via `set_isv_signals_ptr`.
|
||
isv_signals_ptr: default_isv_ptr,
|
||
default_isv_buf,
|
||
config,
|
||
forward_actions_buf,
|
||
action_select_kernel: None,
|
||
b0_size: 0,
|
||
b1_size: 0,
|
||
b2_size: 0,
|
||
b3_size: 0,
|
||
q_gap_threshold: 0.0,
|
||
chunked_q_gaps_buf: None,
|
||
chunked_conviction_buf: None,
|
||
chunked_magnitude_conviction_buf: None,
|
||
dqn_graph: None,
|
||
chunked_q_values: None,
|
||
chunked_b_logits_buf: None,
|
||
chunked_v_logits_buf: None,
|
||
chunked_states_buf: None,
|
||
// SP21 T2.2 Phase 6.5: lazy-init alongside chunked_states_buf
|
||
// in `ensure_action_select_ready` once we know `state_dim`.
|
||
retained_states_buf: None,
|
||
chunked_actions_buf: None,
|
||
chunked_intent_mag_buf: None,
|
||
scatter_intent_kernel: None,
|
||
plan_state_isv_kernel: None,
|
||
plan_diag_reduce_kernel: None,
|
||
tlob: None,
|
||
})
|
||
}
|
||
|
||
/// Set ISV signals pointer for adaptive hold in validation backtest.
|
||
/// Pass the frozen ISV signals from the last training step.
|
||
pub fn set_isv_signals_ptr(&mut self, ptr: u64) {
|
||
self.isv_signals_ptr = ptr;
|
||
}
|
||
|
||
/// Install a val-path TLOB instance and immediately sync its weights from the training TLOB.
|
||
///
|
||
/// `eval_tlob` must already be constructed (with a `PerStreamCublasHandles` bound to
|
||
/// `self.stream`). `train_tlob` supplies the source weight parameters — copied via
|
||
/// DtoD async on `eval_tlob`'s stream.
|
||
///
|
||
/// Calling this before `evaluate_dqn_graphed` ensures every val forward pass runs the
|
||
/// same TLOB weights as the most recent training step, preserving train/val state-dist
|
||
/// parity for the TLOB slot (SL_TLOB_START) in the assembled state vector.
|
||
pub(crate) fn set_tlob_from_training(
|
||
&mut self,
|
||
mut eval_tlob: super::gpu_tlob::GpuTlob,
|
||
train_tlob: &super::gpu_tlob::GpuTlob,
|
||
) -> Result<(), MLError> {
|
||
eval_tlob.copy_params_from(train_tlob)?;
|
||
self.tlob = Some(eval_tlob);
|
||
Ok(())
|
||
}
|
||
|
||
// ── Private helpers ───────────────────────────────────────────────────────
|
||
|
||
/// Build the initial portfolio state vector.
|
||
///
|
||
/// Layout per window (8 floats):
|
||
/// `[value, position, cash, unrealised_pnl, max_equity, 0, 0, 0]`
|
||
fn init_portfolio_state(n_windows: usize, initial_capital: f32) -> Vec<f32> {
|
||
let mut state = vec![0.0_f32; n_windows * 8];
|
||
for w in 0..n_windows {
|
||
let base = w * 8;
|
||
state[base] = initial_capital; // value
|
||
// state[base + 1] = 0.0 // position (zero-initialised)
|
||
state[base + 2] = initial_capital; // cash
|
||
// state[base + 3] = 0.0 // unrealised_pnl
|
||
state[base + 4] = initial_capital; // max_equity (high-water mark)
|
||
}
|
||
state
|
||
}
|
||
|
||
/// Get the CUDA stream used by this evaluator.
|
||
///
|
||
/// Callers can use this to extract weight buffers on the same stream,
|
||
/// avoiding unnecessary synchronization between weight upload and
|
||
/// kernel launches.
|
||
pub fn stream(&self) -> &Arc<CudaStream> {
|
||
&self.stream
|
||
}
|
||
|
||
/// Number of evaluation windows (validation fold count).
|
||
///
|
||
/// Used by callers that need to size per-window buffers or TLOB instances
|
||
/// to match the evaluator's batch geometry.
|
||
pub(crate) fn n_windows(&self) -> usize {
|
||
self.n_windows
|
||
}
|
||
|
||
/// Set the four factored-action branch sizes from a `DqnBacktestConfig`.
|
||
///
|
||
/// The closure-based `evaluate()` path delegates the model forward pass
|
||
/// to a user-supplied closure but still uses the internal `env_step`
|
||
/// kernel to advance the simulator. `env_step` decodes flat action
|
||
/// indices into `(dir, mag, order, urgency)` via `decode_*_4b()` helpers
|
||
/// in `trade_physics.cuh`, which compute `action / (b1 * b2 * b3)`. If
|
||
/// any `b*_size` is zero the helpers divide by zero, the kernel reads
|
||
/// out-of-bounds, and CUDA raises `CUDA_ERROR_ILLEGAL_ADDRESS` at the
|
||
/// next synchronise — exactly the v7 smoke failure mode (2026-05-12).
|
||
///
|
||
/// The production `evaluate_dqn_graphed` path sets these via
|
||
/// `ensure_action_select_ready` (which also allocates intent buffers
|
||
/// the closure path doesn't need). This setter is the closure-path
|
||
/// equivalent — set b-sizes only, no buffer allocs.
|
||
pub fn set_branch_sizes(&mut self, dqn_cfg: &DqnBacktestConfig) {
|
||
self.b0_size = dqn_cfg.branch_0_size as i32;
|
||
self.b1_size = dqn_cfg.branch_1_size as i32;
|
||
self.b2_size = dqn_cfg.branch_2_size as i32;
|
||
self.b3_size = dqn_cfg.branch_3_size as i32;
|
||
}
|
||
|
||
/// Maximum batch size that the val TLOB instance must support.
|
||
///
|
||
/// `submit_dqn_step_loop_cublas` runs TLOB once per chunk on
|
||
/// `[chunk_len * n_windows, state_dim_padded]` rows of `chunked_states_buf`
|
||
/// instead of once per step on `[n_windows, ..]` rows of `states_buf`.
|
||
/// The val TLOB constructor (`metrics.rs::reset_evaluator_for_validation`)
|
||
/// must allocate buffers sized for this maximum so the per-chunk forward
|
||
/// can proceed without reallocation. Partial last chunks (chunk_len <
|
||
/// DQN_BACKTEST_CHUNK_SIZE) reuse the same buffers — `forward()` accepts
|
||
/// any `batch_size <= construction_batch_size`.
|
||
pub(crate) fn val_tlob_batch_size(&self) -> usize {
|
||
DQN_BACKTEST_CHUNK_SIZE * self.n_windows
|
||
}
|
||
|
||
/// Sharpe annualization factor used by `compute_backtest_metrics` kernel.
|
||
///
|
||
/// Returns the EXACT factor applied to `(mean/std)` inside the kernel
|
||
/// (`sqrt(bars_per_day × trading_days_per_year)` — see line 667 init).
|
||
/// Callers that need to un-annualize the kernel's reported Sharpe MUST
|
||
/// use this getter rather than recomputing from their own copy of
|
||
/// `bars_per_day`, otherwise the two numbers can drift if the evaluator's
|
||
/// config was overridden between construction and read (e.g. via
|
||
/// hyperopt). Source-of-truth alignment per `feedback_no_partial_refactor.md`.
|
||
pub fn annualization_factor(&self) -> f32 {
|
||
self.annualization_factor
|
||
}
|
||
|
||
/// C.3 Plan 3 Task 7: device pointer + sample count for the val-state
|
||
/// batch the most recent `evaluate_dqn_graphed` call left in
|
||
/// `chunked_states_buf`. Layout: `[chunk_len × n_windows, state_dim_padded]`
|
||
/// where `chunk_len = DQN_BACKTEST_CHUNK_SIZE` (rows are time-major
|
||
/// per-step blocks of `n_windows` rows; rows beyond the last actual chunk
|
||
/// length carry stale data from earlier chunks but the moments are not
|
||
/// sensitive to that within the OFI block).
|
||
///
|
||
/// Returns `(0, 0)` when the chunked buffer has not yet been allocated
|
||
/// (no val pass has run yet) — the caller's `state_kl` launcher no-ops in
|
||
/// that case.
|
||
pub(crate) fn val_state_sample(&self) -> (u64, usize) {
|
||
match self.chunked_states_buf.as_ref() {
|
||
Some(buf) => {
|
||
let (ptr, _guard) = buf.device_ptr(&self.stream);
|
||
(ptr, self.n_windows * DQN_BACKTEST_CHUNK_SIZE)
|
||
}
|
||
None => (0, 0),
|
||
}
|
||
}
|
||
|
||
// ── Public API ────────────────────────────────────────────────────────────
|
||
|
||
/// Build the state tensor for a given step: `[n_windows, feat_dim + portfolio_dim]`.
|
||
///
|
||
/// Launches the `gather_states` CUDA kernel which reads directly from the
|
||
/// pre-uploaded features buffer and the live portfolio state buffer on GPU,
|
||
/// then performs a zero-copy DtoD transfer from the kernel output `states_buf`
|
||
/// into a freshly-allocated Candle tensor. No data ever touches the CPU.
|
||
///
|
||
/// The kernel writes into `states_buf` (pre-allocated, `n_windows × state_dim`).
|
||
/// A `cuMemcpyDtoDAsync` then copies those bytes directly into the Candle tensor's
|
||
/// CUDA storage — eliminating the GPU→CPU→GPU roundtrip entirely.
|
||
///
|
||
/// # Panics
|
||
/// `portfolio_dim` must equal `self.portfolio_dim` (24: 8 portfolio + 16 multi-timeframe).
|
||
/// This matches the training experience_state_gather layout.
|
||
pub fn gather_states(
|
||
&self,
|
||
step: usize,
|
||
portfolio_dim: usize,
|
||
_device: &(),
|
||
) -> Result<CudaSlice<f32>, MLError> {
|
||
if portfolio_dim != self.portfolio_dim {
|
||
return Err(MLError::ConfigError(format!(
|
||
"gather_states: portfolio_dim={portfolio_dim} != expected {}",
|
||
self.portfolio_dim
|
||
)));
|
||
}
|
||
|
||
let state_dim = self.state_dim;
|
||
|
||
// Launch the gather kernel — one thread per window
|
||
// Shared memory: 256 threads × 8 portfolio f32 × 4 bytes = 8 KB
|
||
let grid = ((self.n_windows + 255) / 256) as u32;
|
||
let launch_cfg = LaunchConfig {
|
||
grid_dim: (grid.max(1), 1, 1),
|
||
block_dim: (256, 1, 1),
|
||
shared_mem_bytes: 256 * 8 * std::mem::size_of::<f32>() as u32,
|
||
};
|
||
let n_windows_i32 = self.n_windows as i32;
|
||
let max_len_i32 = self.max_len as i32;
|
||
let feat_dim_i32 = self.feature_dim as i32;
|
||
let step_i32 = step as i32;
|
||
let padded_sd_i32 = ml_core::state_layout::STATE_DIM_PADDED as i32;
|
||
|
||
// Safety: argument order matches `backtest_state_gather` signature
|
||
// exactly: features, portfolio, plan_isv, states_out, n_windows,
|
||
// max_len, feat_dim, current_step, padded_sd, aux_dir_prob_per_env
|
||
// (SP22 H6 Phase 2), aux_outcome_probs_per_env (Phase C-2).
|
||
//
|
||
// Phase 1 A3 (2026-05-12) + Phase C-2 (2026-05-14): eval has no
|
||
// aux producer infrastructure yet (neither K=2 head nor K=3 head
|
||
// runs in eval), so both pointers stay NULL → kernel uses 0.0f
|
||
// sentinel for slot 121 via the K=2 fallback path (since
|
||
// aux_outcome_probs_per_env is NULL, branch falls through to
|
||
// K=2 assemble_state which writes 0.0 to slot 121 when
|
||
// aux_dir_prob_per_env is also NULL). A2 follow-up would wire
|
||
// both buffers via the eval-side aux trunk forward.
|
||
let aux_dir_prob_null: u64 = 0;
|
||
let aux_outcome_probs_null: u64 = 0;
|
||
unsafe {
|
||
self.stream
|
||
.launch_builder(&self.gather_kernel)
|
||
.arg(&self.features_buf.dev_ptr)
|
||
.arg(&self.portfolio_buf.dev_ptr)
|
||
.arg(&self.plan_isv_buf)
|
||
.arg(&self.states_buf)
|
||
.arg(&n_windows_i32)
|
||
.arg(&max_len_i32)
|
||
.arg(&feat_dim_i32)
|
||
.arg(&step_i32)
|
||
.arg(&padded_sd_i32)
|
||
.arg(&aux_dir_prob_null) // SP22 H6 A3: NULL (no K=2 in eval)
|
||
.arg(&aux_outcome_probs_null) // SP22 H6 vNext Phase C-2: NULL (no K=3 in eval)
|
||
.launch(launch_cfg)
|
||
.map_err(|e| MLError::ModelError(format!("backtest_state_gather launch step {step}: {e}")))?;
|
||
}
|
||
|
||
// Return a copy of the states buffer as CudaSlice<f32> (padded stride).
|
||
let state_dim_padded = (state_dim + 127) & !127;
|
||
let n_elems = self.n_windows * state_dim_padded;
|
||
let dst = self.stream.alloc_zeros::<f32>(n_elems)
|
||
.map_err(|e| MLError::ModelError(format!("alloc states step {step}: {e}")))?;
|
||
let src_view = self.states_buf.slice(..n_elems);
|
||
{
|
||
let (src_ptr, _src_sync) = src_view.device_ptr(&self.stream);
|
||
let (dst_ptr, _dst_sync) = dst.device_ptr(&self.stream);
|
||
let num_bytes = n_elems * std::mem::size_of::<f32>();
|
||
unsafe {
|
||
cudarc::driver::result::memcpy_dtod_async(
|
||
dst_ptr, src_ptr, num_bytes, self.stream.cu_stream(),
|
||
).map_err(|e| MLError::ModelError(format!("states DtoD copy step {step}: {e}")))?;
|
||
}
|
||
}
|
||
Ok(dst)
|
||
}
|
||
|
||
/// Run the full backtest evaluation loop (closure-based, single-stream).
|
||
///
|
||
/// For each step up to `max_len`:
|
||
/// 1. Launch `gather_states` GPU kernel → state tensor `[n_windows, state_dim]`
|
||
/// 2. Call `forward_fn` to get Q-values `[n_windows, n_actions]`
|
||
/// 3. Greedy argmax → action indices
|
||
/// 4. Launch `backtest_env_step` kernel
|
||
///
|
||
/// After the loop, launches `compute_backtest_metrics`.
|
||
/// Returns one `WindowMetrics` per window.
|
||
pub fn evaluate<F>(
|
||
&mut self,
|
||
forward_fn: &F,
|
||
portfolio_dim: usize,
|
||
) -> Result<Vec<WindowMetrics>, MLError>
|
||
where
|
||
F: Fn(&CudaSlice<f32>, usize, usize) -> Result<CudaSlice<i32>, MLError>,
|
||
{
|
||
let _nvtx = NvtxRange::new("backtest_evaluate");
|
||
|
||
// Phase 8.5 (2026-05-12) — defensive guard: `env_step` decodes flat
|
||
// action indices into (dir, mag, order, urgency) via
|
||
// `decode_*_4b()` helpers that divide by `b1 * b2 * b3`. If any
|
||
// b-size is zero, the kernel divides by zero and raises
|
||
// CUDA_ERROR_ILLEGAL_ADDRESS asynchronously at the next event-sync.
|
||
// The production `evaluate_dqn_graphed` path sets these via
|
||
// `ensure_action_select_ready`; closure-path callers must invoke
|
||
// `set_branch_sizes` before `evaluate()`. Surface the requirement
|
||
// up-front rather than crashing the CUDA context downstream.
|
||
if self.b0_size == 0 || self.b1_size == 0 || self.b2_size == 0 || self.b3_size == 0 {
|
||
return Err(MLError::ConfigError(format!(
|
||
"evaluate: branch sizes unset (b0={}, b1={}, b2={}, b3={}); \
|
||
call `set_branch_sizes(&DqnBacktestConfig)` before `evaluate()` \
|
||
to wire factored-action decoding for env_step.",
|
||
self.b0_size, self.b1_size, self.b2_size, self.b3_size
|
||
)));
|
||
}
|
||
|
||
let state_dim = self.state_dim;
|
||
|
||
for step in 0..self.max_len {
|
||
// 1. Gather state CudaSlice via GPU kernel [n_windows * state_dim]
|
||
let states = self.gather_states(step, portfolio_dim, &())?;
|
||
|
||
// 2. Model forward pass + greedy argmax (on-device, no roundtrip)
|
||
// Closure returns action indices as CudaSlice<i32> [n_windows]
|
||
let action_indices = forward_fn(&states, self.n_windows, state_dim)?;
|
||
|
||
// DtoD copy: action indices → actions_buf (same stream, implicit ordering)
|
||
{
|
||
let src_view = action_indices.slice(..self.n_windows);
|
||
let (src_ptr, _src_sync) = src_view.device_ptr(&self.stream);
|
||
let (dst_ptr, _dst_sync) = self.actions_buf.device_ptr(&self.stream);
|
||
let num_bytes = self.n_windows * std::mem::size_of::<i32>();
|
||
unsafe {
|
||
cudarc::driver::result::memcpy_dtod_async(
|
||
dst_ptr,
|
||
src_ptr,
|
||
num_bytes,
|
||
self.stream.cu_stream(),
|
||
)
|
||
.map_err(|e| MLError::ModelError(format!("actions DtoD step {step}: {e}")))?;
|
||
}
|
||
}
|
||
|
||
// 3. Launch env step kernel on the same stream — sequential ordering
|
||
// guarantees actions_buf is ready. The kernel also scatter-writes
|
||
// actions into actions_history_buf (GPU-resident, no CPU roundtrip).
|
||
self.launch_env_step(step)?;
|
||
|
||
// No mid-loop GPU→CPU transfers. The env kernel checks done_flags per-thread
|
||
// and early-returns (no-op) for finished windows, so running all max_len steps
|
||
// is safe and avoids the latency of periodic done-flag downloads.
|
||
}
|
||
|
||
// Synchronous launch+consume — `evaluate()` is the closure-based path
|
||
// used by callers that don't pipeline (no caller currently does — the
|
||
// pipelined DQN path goes through `evaluate_dqn_graphed_async`).
|
||
self.launch_metrics_and_record_event()?;
|
||
self.consume_metrics_after_event()
|
||
}
|
||
|
||
/// Pure-GPU DQN evaluation via cuBLAS SGEMM — no candle, no closures.
|
||
///
|
||
/// Uses cuBLAS matrix multiplications for the Q-network forward pass,
|
||
/// replacing the old warp-cooperative shared-memory-tiling kernel that
|
||
/// crashed with `CUDA_ERROR_INVALID_VALUE` for `hidden_dim >= 768`
|
||
/// (shared memory exceeds 49 KB limit).
|
||
///
|
||
/// # Arguments
|
||
/// * `online_weights` - Dueling Q-network weight buffers (12 CudaSlice pointers).
|
||
/// For branching DQN, `w_a1`/`b_a1`/`w_a2`/`b_a2` map to branch 0 (exposure).
|
||
/// * `branching_weights` - Optional branch 1 (order) and branch 2 (urgency) weights.
|
||
/// When `None`, branch 0 weights are replicated for all 3 branches.
|
||
/// * `dqn_cfg` - Network dimensions and C51 distributional parameters.
|
||
// evaluate_dqn() removed — all evaluation routes through evaluate_dqn_graphed()
|
||
// with QValueProvider for deterministic results.
|
||
|
||
// ── CUDA Graph–accelerated DQN evaluation ──────────────────────────────────
|
||
|
||
/// Pure-GPU DQN evaluation with CUDA Graph capture and replay.
|
||
///
|
||
/// Identical to [`evaluate_dqn`] but wraps the step loop in a CUDA Graph.
|
||
/// On the first call, the entire step loop (gather + cuBLAS forward +
|
||
/// expected_q + action_select + env_step for each step 0..max_len) is
|
||
/// captured into a graph and then launched. On subsequent calls with the
|
||
/// same evaluator, the cached graph is replayed.
|
||
///
|
||
/// The metrics kernel runs OUTSIDE the graph (one-shot, not worth graphing).
|
||
///
|
||
/// # Fallback
|
||
///
|
||
/// If graph capture or instantiation fails (e.g., unsupported driver version,
|
||
/// or a captured operation that is not graph-safe), falls back to the
|
||
/// non-graphed `evaluate_dqn` path with a warning.
|
||
///
|
||
/// # Graph invalidation
|
||
///
|
||
/// The cached graph records the device pointer values of all buffers and
|
||
/// weight pointers at capture time. It is valid as long as:
|
||
/// - The same `GpuBacktestEvaluator` instance is used (buffer pointers stable)
|
||
/// - The same `online_weights` object is passed (weight pointers stable)
|
||
///
|
||
/// If you change weights (different `DuelingWeightSet`), call
|
||
/// [`invalidate_dqn_graph`] first or just use `evaluate_dqn` instead.
|
||
pub fn evaluate_dqn_graphed(
|
||
&mut self,
|
||
online_weights: &DuelingWeightSet,
|
||
branching_weights: Option<&BranchingWeightSet>,
|
||
dqn_cfg: &DqnBacktestConfig,
|
||
q_provider: Option<&mut dyn super::q_value_provider::QValueProvider>,
|
||
) -> Result<Vec<WindowMetrics>, MLError> {
|
||
// Synchronous wrapper: launch async + immediate consume. Used by
|
||
// hyperopt (single-shot) and the smoke-test path that expects metrics
|
||
// back from this call without a separate consume step. The pipelined
|
||
// training path uses `evaluate_dqn_graphed_async` + a deferred
|
||
// `consume_metrics_after_event` from the next epoch's start (see
|
||
// `compute_validation_loss_launch` / `_consume` in trainer/metrics.rs).
|
||
self.evaluate_dqn_graphed_async(
|
||
online_weights, branching_weights, dqn_cfg, q_provider,
|
||
)?;
|
||
self.consume_metrics_after_event()
|
||
}
|
||
|
||
/// Async (launch-only) variant of [`evaluate_dqn_graphed`]. Submits the
|
||
/// full eval rollout + metrics + diagnostic kernels to `self.stream` and
|
||
/// records `eval_done_event`, then returns **without** any host-side wait.
|
||
///
|
||
/// The companion [`consume_metrics_after_event`] performs the single host
|
||
/// `cuEventSynchronize` and parses the mapped-pinned `metrics_buf`. Splitting
|
||
/// the two unblocks the host CPU thread to submit the next training epoch's
|
||
/// kernels while eval drains on the dedicated `validation_stream` — the
|
||
/// pipelining the original async-diag commit (d9cb14f1b) wired the GPU side
|
||
/// for but the host-side `synchronize()` inside the now-removed
|
||
/// `launch_metrics_and_download` defeated.
|
||
///
|
||
/// Per-epoch wall-clock saving on L40S: ~25-30s (the original audit's
|
||
/// 30-40s estimate, give or take, depending on training-side launch
|
||
/// density).
|
||
///
|
||
/// # Arguments
|
||
/// Same as [`evaluate_dqn_graphed`]. The `_online_weights` / `_branching_weights`
|
||
/// parameters are placeholders for caller-side ABI symmetry — actual
|
||
/// weight wiring goes through the `q_provider`'s `FusedTrainingCtx` which
|
||
/// owns the live training weights.
|
||
pub fn evaluate_dqn_graphed_async(
|
||
&mut self,
|
||
_online_weights: &DuelingWeightSet,
|
||
_branching_weights: Option<&BranchingWeightSet>,
|
||
dqn_cfg: &DqnBacktestConfig,
|
||
q_provider: Option<&mut dyn super::q_value_provider::QValueProvider>,
|
||
) -> Result<(), MLError> {
|
||
let _nvtx = NvtxRange::new("backtest_evaluate_dqn_graphed_async");
|
||
|
||
// Reset mutable state: portfolio, done flags, step returns, actions history.
|
||
// Without this, the second call sees done_flags=1 from the previous evaluation
|
||
// and skips the entire simulation — producing constant metrics.
|
||
self.reset_evaluation_state()?;
|
||
|
||
// Drain stale CUDA errors without CPU sync.
|
||
// The caller ensures training completion via cuStreamWaitEvent (async validation).
|
||
let _ = self.stream.context().check_err();
|
||
|
||
// Lazy-initialise action selection kernels + chunked buffers
|
||
self.ensure_action_select_ready(dqn_cfg)?;
|
||
|
||
// Direct kernel launches (no CUDA Graph for the backtest step loop).
|
||
//
|
||
// The step loop captures 500K+ kernel launches (10K steps × 10 windows
|
||
// × ~5 ops/step). cuGraphLaunch SIGSEGVs on graphs this large due to
|
||
// CUDA driver-internal memory limits. Direct launches are fine because:
|
||
// - Eval runs ONCE per hyperopt trial (not per training step)
|
||
// - Training is 99%+ of hyperopt wall time; eval is negligible
|
||
// - cuBLAS forward passes dominate eval throughput, not launch overhead
|
||
//
|
||
// The training pipeline (graph_forward + graph_adam in GpuDqnTrainer)
|
||
// still uses CUDA Graphs — those capture ~20 operations per replay,
|
||
// well within driver limits.
|
||
self.submit_dqn_step_loop_cublas(dqn_cfg, q_provider)?;
|
||
|
||
// Val plan_isv parity diagnostic (task #94): reduce plan_state +
|
||
// plan_isv across windows → 8 floats, log for epoch-boundary
|
||
// visibility into plan activity. Cold path, runs once per eval.
|
||
// The host-side `read_all` inside this method reads the PREVIOUS
|
||
// epoch's diagnostic via the documented one-epoch lag — the value
|
||
// currently in mapped-pinned memory was already synchronised by the
|
||
// previous epoch's `consume_metrics_after_event`, so no additional
|
||
// host wait is required here.
|
||
self.launch_plan_diag_and_log()?;
|
||
|
||
self.launch_metrics_and_record_event()
|
||
}
|
||
|
||
/// Val plan_isv parity diagnostic emission (task #94).
|
||
///
|
||
/// Launches `backtest_plan_diag_reduce` on the final plan_state + plan_isv
|
||
/// buffers and logs a single `val_plan_diag` line with labelled stats.
|
||
/// Called once per `evaluate_dqn_graphed_async` invocation at the epoch
|
||
/// boundary. The host-side `read_all` here reads the **previous epoch's**
|
||
/// diagnostic via the documented one-epoch lag (the value currently in
|
||
/// mapped-pinned memory was already synchronised by the previous epoch's
|
||
/// `consume_metrics_after_event`); the just-launched kernel populates the
|
||
/// buffer for the **next** epoch's read.
|
||
fn launch_plan_diag_and_log(&mut self) -> Result<(), MLError> {
|
||
let kernel = self.plan_diag_reduce_kernel.as_ref().ok_or_else(|| {
|
||
MLError::ModelError("plan_diag_reduce_kernel not loaded".to_owned())
|
||
})?;
|
||
let n_i32 = self.n_windows as i32;
|
||
// Mapped-pinned device pointer for the diag scratch — kernel writes
|
||
// through `dev_ptr`, host reads via `host_ptr`. No DtoH per
|
||
// `feedback_no_htod_htoh_only_mapped_pinned.md`. The metrics readback
|
||
// (now `consume_metrics_after_event`) is the single host-side wait
|
||
// per epoch boundary.
|
||
let plan_diag_dev_ptr = self.plan_diag_buf.dev_ptr;
|
||
unsafe {
|
||
self.stream
|
||
.launch_builder(kernel)
|
||
.arg(&self.plan_state_buf)
|
||
.arg(&self.plan_isv_buf)
|
||
.arg(&plan_diag_dev_ptr)
|
||
.arg(&n_i32)
|
||
.launch(LaunchConfig {
|
||
grid_dim: (1, 1, 1),
|
||
block_dim: (256, 1, 1),
|
||
shared_mem_bytes: 0,
|
||
})
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"backtest_plan_diag_reduce launch: {e}"
|
||
)))?;
|
||
}
|
||
// Read mapped-pinned host buffer. The values currently in memory were
|
||
// produced by the PREVIOUS epoch's plan_diag launch and made globally
|
||
// visible by the previous epoch's `consume_metrics_after_event` (which
|
||
// calls `eval_done_event.synchronize()`); the kernel we just launched
|
||
// above will populate the buffer for the NEXT epoch's read. The first
|
||
// epoch returns zeros (safe default from `MappedF32Buffer::new`
|
||
// zero-init); subsequent epochs see the previous epoch's diagnostic —
|
||
// same one-epoch lag pattern as `pending_val_loss`. Acceptable for a
|
||
// diagnostic-only readout.
|
||
let diag = self.plan_diag_buf.read_all();
|
||
info!(
|
||
target: "val_plan_diag",
|
||
"val_plan_isv_diag: active_frac={:.3} n_active={:.0} \
|
||
mean|isv0|(progress)={:.4} mean|isv1|(pnl_vs_tgt)={:.4} mean|isv2|(pnl_vs_stop)={:.4} \
|
||
mean_isv3(conv_entry)={:.4} mean|isv4|(conv_drift)={:.4} mean|isv5|(regime_shift)={:.4}",
|
||
diag[0], diag[7], diag[1], diag[2], diag[3], diag[4], diag[5], diag[6],
|
||
);
|
||
Ok(())
|
||
}
|
||
|
||
|
||
/// Discard the cached CUDA graph.
|
||
///
|
||
/// Call this before `evaluate_dqn_graphed()` if you have changed the
|
||
/// `DuelingWeightSet` (different weight buffer pointers). The next call
|
||
/// to `evaluate_dqn_graphed` will re-capture a fresh graph.
|
||
pub fn invalidate_dqn_graph(&mut self) {
|
||
self.dqn_graph = None;
|
||
}
|
||
|
||
/// Task 0.7 — read back the actions_history buffer and compute the
|
||
/// per-magnitude action distribution observed during the most recent
|
||
/// evaluation. Action encoding is `dir*27 + mag*9 + ord*3 + urg`, so
|
||
/// magnitude = `(action / 9) % 3` (0=Quarter, 1=Half, 2=Full).
|
||
///
|
||
/// Returns `[quarter_frac, half_frac, full_frac]` summing to 1.0 over
|
||
/// non-negative action entries (negative entries indicate uninitialized
|
||
/// or skipped steps and are excluded).
|
||
///
|
||
/// **Sync contract**: caller MUST have invoked
|
||
/// [`consume_metrics_after_event`] (or the synchronous `evaluate_*`
|
||
/// wrappers, which call it internally) since the last
|
||
/// [`launch_metrics_and_record_event`]. The reader does no host wait —
|
||
/// it just `read_volatile`s the mapped-pinned mirror that
|
||
/// `launch_metrics_and_record_event` queued a DtoD-async copy into.
|
||
pub fn read_eval_action_distribution_per_magnitude(&self) -> Result<[f32; 3], MLError> {
|
||
let host = self.actions_history_pinned.read_all();
|
||
let mut counts = [0_u64; 3];
|
||
let mut total = 0_u64;
|
||
for &a in &host {
|
||
if a < 0 {
|
||
continue;
|
||
}
|
||
let mag = ((a / 9) % 3) as usize;
|
||
if mag < 3 {
|
||
counts[mag] += 1;
|
||
total += 1;
|
||
}
|
||
}
|
||
if total == 0 {
|
||
return Ok([0.0; 3]);
|
||
}
|
||
let t = total as f32;
|
||
Ok([
|
||
counts[0] as f32 / t,
|
||
counts[1] as f32 / t,
|
||
counts[2] as f32 / t,
|
||
])
|
||
}
|
||
|
||
/// SP9 (Fix 37, 2026-05-03): GPU-resident accessor for the
|
||
/// `intent_mag_buf` raw device pointer + element count. Replaces the
|
||
/// host-side reduction `read_eval_intent_magnitude_distribution()` that
|
||
/// was here pre-SP9 — that path performed a DtoH `read_all()` over the
|
||
/// full `intent_mag_pinned` mirror plus a host-side loop, violating
|
||
/// `feedback_no_cpu_compute_strict.md` and
|
||
/// `feedback_no_htod_htoh_only_mapped_pinned.md`.
|
||
///
|
||
/// Returns `(0, 0)` when `intent_mag_buf` has not been allocated yet
|
||
/// (lazy-init in `ensure_action_select_ready` — pre-DQN-eval). The
|
||
/// downstream SP9 GPU producer (`launch_eval_intent_dist_compute` on
|
||
/// `GpuDqnTrainer`) handles `n_total == 0` by writing zero fractions
|
||
/// to scratch — Pearl A's sentinel-bootstrap fires on the first
|
||
/// non-empty call.
|
||
///
|
||
/// **Sync contract**: caller must have synchronised on
|
||
/// `eval_done_event` (typically via `consume_metrics_after_event`)
|
||
/// before passing the device pointer through to a GPU launch — same
|
||
/// constraint as the deleted host helper.
|
||
pub fn intent_mag_dev_ptr_and_n(&self) -> (u64, i32) {
|
||
match self.intent_mag_buf.as_ref() {
|
||
Some(b) => {
|
||
let n_total = self.n_windows * self.max_len;
|
||
let (ptr, _guard) = b.device_ptr(&self.stream);
|
||
(ptr, n_total as i32)
|
||
}
|
||
None => (0_u64, 0_i32),
|
||
}
|
||
}
|
||
|
||
/// Task 2.X "make Full useful" diagnostic — read back the per-direction
|
||
/// action distribution from the most recent evaluation. Action encoding
|
||
/// is `dir*27 + mag*9 + ord*3 + urg`, so direction = `action / 27`
|
||
/// (0=Short, 1=Hold, 2=Long, 3=Flat).
|
||
///
|
||
/// Returns `[short, hold, long, flat]` summing to 1.0 over non-negative
|
||
/// action entries. When eval is strict-argmax and direction Q-values are
|
||
/// near-uniform, this distribution helps diagnose whether the eval
|
||
/// collapse is direction-side (e.g., `hold + flat > 0.95`) or whether
|
||
/// direction is healthy but magnitude itself is collapsing. Paired with
|
||
/// `read_eval_action_distribution_per_magnitude` to localise the
|
||
/// magnitude-only Task 2.X fix's leverage surface.
|
||
///
|
||
/// **Sync contract**: same as
|
||
/// [`read_eval_action_distribution_per_magnitude`].
|
||
pub fn read_eval_action_distribution_per_direction(&self) -> Result<[f32; 4], MLError> {
|
||
let host = self.actions_history_pinned.read_all();
|
||
let mut counts = [0_u64; 4];
|
||
let mut total = 0_u64;
|
||
for &a in &host {
|
||
if a < 0 {
|
||
continue;
|
||
}
|
||
let dir = (a / 27) as usize;
|
||
if dir < 4 {
|
||
counts[dir] += 1;
|
||
total += 1;
|
||
}
|
||
}
|
||
if total == 0 {
|
||
return Ok([0.0; 4]);
|
||
}
|
||
let t = total as f32;
|
||
Ok([
|
||
counts[0] as f32 / t,
|
||
counts[1] as f32 / t,
|
||
counts[2] as f32 / t,
|
||
counts[3] as f32 / t,
|
||
])
|
||
}
|
||
|
||
/// Read the per-direction histogram of the policy's RAW Boltzmann picks
|
||
/// across the FULL val window. Reads `picked_action_history_pinned` —
|
||
/// the mapped-pinned mirror of `picked_action_history_buf`, scattered
|
||
/// from `chunked_actions_buf` after each chunk's `experience_action_select`,
|
||
/// so the buffer accumulates the kernel's raw action_idx writes for every
|
||
/// bar before env_step touches anything.
|
||
///
|
||
/// `actions_history_buf` records the POST-physics `actual_dir`, so a
|
||
/// divergence between this reader and `read_eval_action_distribution_per_direction`
|
||
/// localises whether eval collapse is upstream (kernel produces collapsed
|
||
/// picks → both flat) or downstream (kernel diverse, env_step / Kelly cap
|
||
/// drains active picks to Flat → only post-physics flat).
|
||
///
|
||
/// Returns `[short, hold, long, flat]` summing to 1.0 over the full window.
|
||
/// `Ok([0.0; 4])` if `picked_action_history_pinned` has not been allocated yet
|
||
/// (first call before `ensure_action_select_ready` has run).
|
||
///
|
||
/// **Sync contract**: same as
|
||
/// [`read_eval_action_distribution_per_magnitude`].
|
||
pub fn read_chunked_actions_direction_distribution(&self) -> Result<[f32; 4], MLError> {
|
||
let buf = match self.picked_action_history_pinned.as_ref() {
|
||
Some(b) => b,
|
||
None => return Ok([0.0; 4]),
|
||
};
|
||
let host = buf.read_all();
|
||
let mut counts = [0_u64; 4];
|
||
let mut total = 0_u64;
|
||
for &a in &host {
|
||
// action_idx is always non-negative; the only invalid entry is
|
||
// a kernel-side bug. Slots beyond the window length still get
|
||
// written by experience_action_select (it runs for the full
|
||
// chunk_len × n_windows batch regardless of done_flags), so
|
||
// every slot in [0, max_len) carries a real Boltzmann pick.
|
||
// action_idx = 0 (Short Quarter Market Normal) is a legitimate
|
||
// pick; the Short bucket includes it.
|
||
if a < 0 {
|
||
continue;
|
||
}
|
||
let dir = (a / 27) as usize;
|
||
if dir < 4 {
|
||
counts[dir] += 1;
|
||
total += 1;
|
||
}
|
||
}
|
||
if total == 0 {
|
||
return Ok([0.0; 4]);
|
||
}
|
||
let t = total as f32;
|
||
Ok([
|
||
counts[0] as f32 / t,
|
||
counts[1] as f32 / t,
|
||
counts[2] as f32 / t,
|
||
counts[3] as f32 / t,
|
||
])
|
||
}
|
||
|
||
/// Reset mutable evaluation state: portfolio, done flags, step returns, actions history.
|
||
/// Must be called before each `evaluate_dqn_graphed` when reusing the evaluator
|
||
/// across epochs — otherwise done_flags=1 from the previous evaluation causes
|
||
/// the env_step to skip immediately, producing constant metrics.
|
||
fn reset_evaluation_state(&mut self) -> Result<(), MLError> {
|
||
// Re-init portfolio to initial capital via in-place host_ptr write.
|
||
// `portfolio_buf` is a `MappedF32Buffer` (cuMemHostAlloc DEVICEMAP):
|
||
// host writes through `host_ptr`, the env_step kernel reads/writes
|
||
// through `dev_ptr` aliasing the same pinned pages — no DtoD copy,
|
||
// no extra allocation. The kernel sees these writes after the next
|
||
// stream-sync barrier, which the caller (`evaluate_dqn_graphed`)
|
||
// performs before each chunk's first env_step launch.
|
||
let portfolio_init = Self::init_portfolio_state(self.n_windows, self.config.initial_capital);
|
||
self.portfolio_buf.write_from_slice(&portfolio_init);
|
||
|
||
// Zero done flags, step returns, step rewards
|
||
self.stream.memset_zeros(&mut self.done_buf)
|
||
.map_err(|e| MLError::ModelError(format!("reset done_buf: {e}")))?;
|
||
self.stream.memset_zeros(&mut self.step_returns_buf)
|
||
.map_err(|e| MLError::ModelError(format!("reset step_returns: {e}")))?;
|
||
self.stream.memset_zeros(&mut self.step_rewards_buf)
|
||
.map_err(|e| MLError::ModelError(format!("reset step_rewards: {e}")))?;
|
||
|
||
// Initialise actions_history_buf to -1 sentinel (NOT zero) so the
|
||
// reader (`read_eval_action_distribution_per_direction`) can distinguish
|
||
// unwritten slots from real action_idx=0 (Short Quarter Market Normal).
|
||
//
|
||
// Why this matters: after `done_flags[w]=1` (capital floor breach),
|
||
// backtest_env_step early-returns without writing actions_history for
|
||
// the remaining slots in [done_step, max_len). Zero-initialised
|
||
// unwritten slots decoded as Short via `dir = 0/27 = 0` and inflated
|
||
// val_dir_dist's Short bucket — a measurement artifact masking the
|
||
// real direction distribution. The reader already filters `if a < 0
|
||
// { continue; }`, so writing -1 (i32) = 0xFFFFFFFF (u32) lets it skip
|
||
// unwritten slots correctly.
|
||
let ah_ptr = self.actions_history_buf.raw_ptr();
|
||
let ah_len = self.actions_history_buf.len();
|
||
unsafe {
|
||
let rc = cudarc::driver::sys::cuMemsetD32Async(
|
||
ah_ptr,
|
||
0xFFFFFFFFu32,
|
||
ah_len,
|
||
self.stream.cu_stream(),
|
||
);
|
||
if rc != cudarc::driver::sys::cudaError_enum::CUDA_SUCCESS {
|
||
return Err(MLError::ModelError(format!(
|
||
"reset actions_history (cuMemsetD32Async -1): {rc:?}"
|
||
)));
|
||
}
|
||
}
|
||
|
||
// Zero the parallel intent-magnitude history (diagnostic) so a short
|
||
// eval rollout can't read stale values from a prior longer rollout.
|
||
if let Some(ref mut intent) = self.intent_mag_buf {
|
||
self.stream.memset_zeros(intent)
|
||
.map_err(|e| MLError::ModelError(format!("reset intent_mag_buf: {e}")))?;
|
||
}
|
||
// Zero the picked-action history for the same reason.
|
||
if let Some(ref mut picked) = self.picked_action_history_buf {
|
||
self.stream.memset_zeros(picked)
|
||
.map_err(|e| MLError::ModelError(format!("reset picked_action_history_buf: {e}")))?;
|
||
}
|
||
|
||
// Reset Kelly stats — each evaluation starts cold. Kelly priors +
|
||
// warmup floor in trade_physics.cuh handle the cold-start period.
|
||
self.stream.memset_zeros(&mut self.kelly_stats_buf)
|
||
.map_err(|e| MLError::ModelError(format!("reset kelly_stats: {e}")))?;
|
||
|
||
// SP21 T2.2 Phase 1 (2026-05-10) — reset per-trade tape counters
|
||
// so the first close in each window writes at slot 0. The SoA
|
||
// value buffers (pnl, holding_bars, bar_index, dir_mag,
|
||
// predicted_q) don't need zeroing — they're only read up to
|
||
// `count[w]` host-side, and stale values past `count[w]` are
|
||
// never visited by the kernel writer (which guards on
|
||
// `slot < max_trades_per_window`) or the host reader (which
|
||
// slices `[0..count[w]]`).
|
||
self.stream.memset_zeros(&mut self.per_trade_count_buf)
|
||
.map_err(|e| MLError::ModelError(format!("reset per_trade_count: {e}")))?;
|
||
|
||
// SP21 T2.2 Phase 1.5 (2026-05-11) — reset entry_q persistent
|
||
// state so trades from a prior fold don't bleed stale entry_q
|
||
// values into the new fold's tape. The first open of the new
|
||
// fold writes a real entry_q before any close reads it.
|
||
self.stream.memset_zeros(&mut self.entry_q_state_buf)
|
||
.map_err(|e| MLError::ModelError(format!("reset entry_q_state: {e}")))?;
|
||
|
||
Ok(())
|
||
}
|
||
|
||
/// SP21 T2.2 Phase 1 (2026-05-10) — read the per-trade tape from the
|
||
/// device into a host-side `Vec<EvalTrade>`. Called after the eval
|
||
/// window completes (post `launch_metrics_and_record_event` +
|
||
/// `consume_metrics_after_event`); flattens the per-window SoA
|
||
/// buffers into one chronological-by-window concatenated `Vec`.
|
||
///
|
||
/// The kernel emits trades in close-event order WITHIN each window
|
||
/// (single-threaded per-window writes preserve order). The flatten
|
||
/// concatenates by window — call sites that need per-window tapes
|
||
/// can reconstruct via `count[]` if needed.
|
||
///
|
||
/// Synchronizes the eval stream before reading — assumes caller has
|
||
/// already issued the eval kernels and just needs to wait for them
|
||
/// to complete. Returns total trade count + the flat `Vec`.
|
||
/// SP21 T2.2 Phase 6.5 (2026-05-11) — read the val state vector at
|
||
/// `(window_idx, bar_idx)` from the retained-states buffer.
|
||
/// Returns the unpadded state row (length = `state_layout::STATE_DIM`)
|
||
/// as a host-side `Vec<f32>`. Padded columns
|
||
/// `[STATE_DIM..STATE_DIM_PADDED)` are dropped — they're zero
|
||
/// (cuBLAS K-tile alignment, see `backtest_state_gather`).
|
||
///
|
||
/// **Zero-copy via mapped-pinned host_ptr** per
|
||
/// `feedback_no_htod_htoh_only_mapped_pinned`. The retained
|
||
/// buffer's `host_ptr` aliases the same pages the kernel writes
|
||
/// to via `dev_ptr`; reads go through `read_volatile` to defeat
|
||
/// compiler reorder. Caller MUST sync the eval stream BEFORE
|
||
/// invoking — mapped-pinned coherence requires the kernel writes
|
||
/// to have completed first (`cuStreamSynchronize` or an event
|
||
/// wait). The matching wait in production is the
|
||
/// `consume_metrics_after_event` call that gates the
|
||
/// post-enrichment training step.
|
||
///
|
||
/// Used by `training_loop::inject_hindsight_experiences` to build
|
||
/// the `state` field of each synthetic replay tuple.
|
||
/// Bounds-checks `bar_idx < max_len` and `window_idx < n_windows`.
|
||
/// Cold-start (buffer not allocated): returns `Ok(None)`.
|
||
pub fn read_retained_state(
|
||
&self,
|
||
window_idx: usize,
|
||
bar_idx: usize,
|
||
) -> Result<Option<Vec<f32>>, MLError> {
|
||
let retained = match self.retained_states_buf.as_ref() {
|
||
Some(r) => r,
|
||
None => return Ok(None),
|
||
};
|
||
if window_idx >= self.n_windows {
|
||
return Err(MLError::ModelError(format!(
|
||
"read_retained_state: window_idx={window_idx} out of range [0, {})",
|
||
self.n_windows
|
||
)));
|
||
}
|
||
if bar_idx >= self.max_len {
|
||
return Err(MLError::ModelError(format!(
|
||
"read_retained_state: bar_idx={bar_idx} out of range [0, {})",
|
||
self.max_len
|
||
)));
|
||
}
|
||
let state_dim = self.state_dim;
|
||
let state_dim_padded = (state_dim + 127) & !127;
|
||
let row_offset = bar_idx * self.n_windows * state_dim_padded
|
||
+ window_idx * state_dim_padded;
|
||
// Zero-copy host read via mapped-pinned host_ptr. Volatile
|
||
// reads defeat compiler reorder; the eval-stream sync the
|
||
// caller performs upstream is the actual coherence barrier.
|
||
let mut host = Vec::with_capacity(state_dim);
|
||
// Safety: `host_ptr` is valid for `retained.len` f32s by
|
||
// construction (`MappedF32Buffer::new`); offset+state_dim is
|
||
// bounds-checked above (row_offset+state_dim_padded ≤
|
||
// retained.len since the buffer is sized for max_len
|
||
// rows of state_dim_padded each).
|
||
unsafe {
|
||
for i in 0..state_dim {
|
||
host.push(std::ptr::read_volatile(
|
||
retained.host_ptr.add(row_offset + i),
|
||
));
|
||
}
|
||
}
|
||
Ok(Some(host))
|
||
}
|
||
|
||
pub fn read_per_trade_tape(&self) -> Result<Vec<EvalTrade>, MLError> {
|
||
// 1. Read per-window counts (cheap — n_windows u32 values).
|
||
let mut counts = vec![0u32; self.n_windows];
|
||
self.stream
|
||
.memcpy_dtoh(&self.per_trade_count_buf, &mut counts)
|
||
.map_err(|e| MLError::ModelError(format!("read per_trade_count: {e}")))?;
|
||
|
||
// 2. Compute total + early-return if no trades. Saves a 16MB+
|
||
// DtoH copy when the policy was inactive (rare but possible at
|
||
// cold-start before warmup).
|
||
let total: usize = counts.iter().map(|&c| c as usize).sum();
|
||
if total == 0 {
|
||
return Ok(Vec::new());
|
||
}
|
||
|
||
// 3. Read the four SoA buffers in full (n_windows * MAX_TRADES_PER_WINDOW
|
||
// each). Slicing window-side would need an extra index-list copy
|
||
// per window; reading whole buffers is simpler and stays cheap
|
||
// (16MB at PCIe ≈ 1ms — well below per-epoch overhead). The host
|
||
// loop below filters down to `count[w]` per window.
|
||
let buf_len = self.n_windows * MAX_TRADES_PER_WINDOW;
|
||
let mut pnl_host = vec![0.0_f32; buf_len];
|
||
let mut holding_host = vec![0u32; buf_len];
|
||
let mut bar_idx_host = vec![0u32; buf_len];
|
||
let mut dir_mag_host = vec![0u32; buf_len];
|
||
let mut predicted_q_host = vec![0.0_f32; buf_len];
|
||
|
||
self.stream.memcpy_dtoh(&self.per_trade_pnl_buf, &mut pnl_host)
|
||
.map_err(|e| MLError::ModelError(format!("read per_trade_pnl: {e}")))?;
|
||
self.stream.memcpy_dtoh(&self.per_trade_holding_bars_buf, &mut holding_host)
|
||
.map_err(|e| MLError::ModelError(format!("read per_trade_holding_bars: {e}")))?;
|
||
self.stream.memcpy_dtoh(&self.per_trade_bar_index_buf, &mut bar_idx_host)
|
||
.map_err(|e| MLError::ModelError(format!("read per_trade_bar_index: {e}")))?;
|
||
self.stream.memcpy_dtoh(&self.per_trade_dir_mag_buf, &mut dir_mag_host)
|
||
.map_err(|e| MLError::ModelError(format!("read per_trade_dir_mag: {e}")))?;
|
||
self.stream.memcpy_dtoh(&self.per_trade_predicted_q_buf, &mut predicted_q_host)
|
||
.map_err(|e| MLError::ModelError(format!("read per_trade_predicted_q: {e}")))?;
|
||
|
||
// 4. Flatten window-major: for each window, take the first
|
||
// count[w] trades and append.
|
||
let mut tape = Vec::with_capacity(total);
|
||
for w in 0..self.n_windows {
|
||
let count = counts[w] as usize;
|
||
// Defensive cap: if a future kernel change skips the
|
||
// overflow guard, count COULD exceed MAX_TRADES_PER_WINDOW.
|
||
// The min() avoids out-of-bounds slicing.
|
||
let count = count.min(MAX_TRADES_PER_WINDOW);
|
||
let base = w * MAX_TRADES_PER_WINDOW;
|
||
for i in 0..count {
|
||
let idx = base + i;
|
||
let dm = dir_mag_host[idx];
|
||
let direction = ((dm >> 16) & 0xFFFF) as u8;
|
||
let magnitude = (dm & 0xFFFF) as u8;
|
||
tape.push(EvalTrade {
|
||
window_index: w as u32,
|
||
bar_index: bar_idx_host[idx],
|
||
pnl: pnl_host[idx],
|
||
predicted_q: predicted_q_host[idx],
|
||
holding_bars: holding_host[idx],
|
||
direction,
|
||
magnitude,
|
||
});
|
||
}
|
||
}
|
||
Ok(tape)
|
||
}
|
||
|
||
/// Submit the cuBLAS-based DQN step loop to the stream (chunked).
|
||
///
|
||
/// Processes `DQN_BACKTEST_CHUNK_SIZE` steps per cuBLAS forward call to
|
||
/// amortise the ~3 us kernel-launch overhead. Within each chunk:
|
||
///
|
||
/// 1. Gather states for all steps in the chunk (per-step kernel + DtoD copy
|
||
/// into the chunked states buffer)
|
||
/// 2. Single cuBLAS forward on [n_windows * chunk_len, state_dim]
|
||
/// 3. Single `compute_expected_q` on the chunked logits
|
||
/// 4. Single `experience_action_select` on the chunked Q-values
|
||
/// 5. Per-step: DtoD copy chunk actions -> actions_buf + env_step (sequential)
|
||
///
|
||
/// Portfolio features (3/48 dims) are stale within a chunk because env_step
|
||
/// updates portfolio_buf but gather already ran for all steps in the chunk.
|
||
/// This is acceptable: market features (42 dims) dominate the Q-network and
|
||
/// a 64-bar (~1 hour) staleness window has negligible impact on action quality.
|
||
fn submit_dqn_step_loop_cublas(
|
||
&mut self,
|
||
dqn_cfg: &DqnBacktestConfig,
|
||
mut q_provider: Option<&mut dyn super::q_value_provider::QValueProvider>,
|
||
) -> Result<(), MLError> {
|
||
let as_kernel = self.action_select_kernel.as_ref().ok_or_else(|| {
|
||
MLError::ModelError("action_select_kernel unexpectedly None".to_owned())
|
||
})?;
|
||
|
||
let ch_q_values = self.chunked_q_values.as_ref().ok_or_else(|| {
|
||
MLError::ModelError("chunked_q_values unexpectedly None".to_owned())
|
||
})?;
|
||
// Note: `chunked_states_buf` is NOT borrowed up-front. Each chunk
|
||
// iteration re-derives a fresh borrow after the optional TLOB call
|
||
// (which needs `&mut self.chunked_states_buf`). The previous shared
|
||
// borrow held across the loop blocked the TLOB mutable borrow
|
||
// required by the chunk-batched TLOB refactor.
|
||
let ch_actions = self.chunked_actions_buf.as_ref().ok_or_else(|| {
|
||
MLError::ModelError("chunked_actions_buf unexpectedly None".to_owned())
|
||
})?;
|
||
let ch_q_gaps = self.chunked_q_gaps_buf.as_ref().ok_or_else(|| {
|
||
MLError::ModelError("chunked_q_gaps_buf unexpectedly None".to_owned())
|
||
})?;
|
||
let ch_intent = self.chunked_intent_mag_buf.as_ref().ok_or_else(|| {
|
||
MLError::ModelError("chunked_intent_mag_buf unexpectedly None".to_owned())
|
||
})?;
|
||
let intent_history = self.intent_mag_buf.as_ref().ok_or_else(|| {
|
||
MLError::ModelError("intent_mag_buf unexpectedly None".to_owned())
|
||
})?;
|
||
let picked_action_history = self.picked_action_history_buf.as_ref().ok_or_else(|| {
|
||
MLError::ModelError("picked_action_history_buf unexpectedly None".to_owned())
|
||
})?;
|
||
let scatter_intent = self.scatter_intent_kernel.as_ref().ok_or_else(|| {
|
||
MLError::ModelError("scatter_intent_kernel unexpectedly None".to_owned())
|
||
})?;
|
||
let ch_conviction = self.chunked_conviction_buf.as_ref().ok_or_else(|| {
|
||
MLError::ModelError("chunked_conviction_buf unexpectedly None".to_owned())
|
||
})?;
|
||
let ch_magnitude_conviction = self.chunked_magnitude_conviction_buf.as_ref().ok_or_else(|| {
|
||
MLError::ModelError("chunked_magnitude_conviction_buf unexpectedly None".to_owned())
|
||
})?;
|
||
let ch_b_logits = self.chunked_b_logits_buf.as_ref().ok_or_else(|| {
|
||
MLError::ModelError("chunked_b_logits_buf unexpectedly None".to_owned())
|
||
})?;
|
||
// SP17 Commit C: chunked V-head logits row for Thompson direction sampling.
|
||
let ch_v_logits = self.chunked_v_logits_buf.as_ref().ok_or_else(|| {
|
||
MLError::ModelError("chunked_v_logits_buf unexpectedly None".to_owned())
|
||
})?;
|
||
|
||
let n = self.n_windows;
|
||
let state_dim = self.state_dim;
|
||
let b0 = dqn_cfg.branch_0_size as i32;
|
||
let b1 = dqn_cfg.branch_1_size as i32;
|
||
let b2 = dqn_cfg.branch_2_size as i32;
|
||
let b3 = dqn_cfg.branch_3_size as i32;
|
||
// SP21 T2.2 Phase 1.5 (2026-05-11) — total_actions for
|
||
// `q_values_per_window` indexing in the batch env kernel.
|
||
// Mirrors `dqn_cfg`'s product (line ~2714 in
|
||
// `ensure_action_select_ready` uses the same product to size
|
||
// `chunked_q_values` at `cn * total_actions`).
|
||
let num_actions_i32 = b0 * b1 * b2 * b3;
|
||
let state_dim_padded = (state_dim + 127) & !127;
|
||
let state_row_bytes = n * state_dim_padded * std::mem::size_of::<f32>();
|
||
|
||
let total_chunks = (self.max_len + DQN_BACKTEST_CHUNK_SIZE - 1) / DQN_BACKTEST_CHUNK_SIZE;
|
||
for chunk_start in (0..self.max_len).step_by(DQN_BACKTEST_CHUNK_SIZE) {
|
||
let chunk_end = (chunk_start + DQN_BACKTEST_CHUNK_SIZE).min(self.max_len);
|
||
let chunk_len = chunk_end - chunk_start;
|
||
let batch = n * chunk_len; // total rows for this chunk's forward pass
|
||
let chunk_idx = chunk_start / DQN_BACKTEST_CHUNK_SIZE;
|
||
|
||
// ── Phase 1: Gather states for the entire chunk in a single launch ────
|
||
//
|
||
// `backtest_state_gather_chunk` writes [chunk_len, n_windows, padded_sd]
|
||
// directly into chunked_states_buf with one kernel launch + zero DtoD
|
||
// copies (replaces chunk_len gather launches + chunk_len DtoD copies).
|
||
//
|
||
// Within a chunk portfolio_buf and plan_isv_buf are CONSTANT — env_step
|
||
// updates them only at chunk boundary (Phase 5 + Phase 6). So the per-
|
||
// step gather sees identical inputs and the batched gather is bit-
|
||
// identical to the prior per-step + DtoD pattern (modulo float
|
||
// reorder, of which there is none — every thread writes its own row
|
||
// from independent feature reads).
|
||
//
|
||
// After gather, TLOB runs ONCE on `chunked_states_buf` with batch
|
||
// = chunk_len * n_windows (replacing chunk_len separate forward
|
||
// calls each at batch = n_windows). The val TLOB instance is sized
|
||
// to `val_tlob_batch_size()` = DQN_BACKTEST_CHUNK_SIZE * n_windows
|
||
// by `metrics.rs::reset_evaluator_for_validation`; partial last-
|
||
// chunks (chunk_len < DQN_BACKTEST_CHUNK_SIZE) reuse the same buffers.
|
||
//
|
||
// Borrow scope: ch_states_base is computed in a tight scope so the
|
||
// immutable borrow drops before we take &mut self.chunked_states_buf
|
||
// for TLOB below. The raw u64 device pointer survives the borrow
|
||
// (it's just a number copied into a stack local).
|
||
let ch_states_base = {
|
||
let ch_states = self.chunked_states_buf.as_ref().ok_or_else(|| {
|
||
MLError::ModelError("chunked_states_buf unexpectedly None".to_owned())
|
||
})?;
|
||
let (ptr, _guard) = ch_states.device_ptr(&self.stream);
|
||
ptr
|
||
};
|
||
|
||
self.launch_gather_chunk(ch_states_base, chunk_start, chunk_len)?;
|
||
|
||
// SP21 T2.2 Phase 6.5 (2026-05-11) — retain the just-gathered
|
||
// chunk's states for hindsight injection. The chunked layout
|
||
// is `[chunk_len, n_windows, padded_sd]` (matches retained
|
||
// buffer's layout); DtoD-copy as a single contiguous block
|
||
// starting at `chunk_start × n_windows × padded_sd`. No
|
||
// transpose needed — both buffers share `[step, window,
|
||
// dim]` ordering.
|
||
//
|
||
// The retained buffer is mapped-pinned (`MappedF32Buffer`)
|
||
// per `feedback_no_htod_htoh_only_mapped_pinned`; the DtoD
|
||
// writes through the device-aliased `dev_ptr` (same memory
|
||
// the host sees via `host_ptr` after stream sync).
|
||
//
|
||
// Skip the copy if the retained buffer isn't allocated
|
||
// (test scaffolds bypassing the standard init path —
|
||
// hindsight injection just degrades to a no-op in that case).
|
||
if let Some(ref retained) = self.retained_states_buf {
|
||
let bytes = chunk_len * n * state_dim_padded
|
||
* std::mem::size_of::<f32>();
|
||
let dst_offset_bytes = (chunk_start * n * state_dim_padded
|
||
* std::mem::size_of::<f32>()) as u64;
|
||
let dst_ptr = retained.dev_ptr + dst_offset_bytes;
|
||
unsafe {
|
||
cudarc::driver::result::memcpy_dtod_async(
|
||
dst_ptr, ch_states_base, bytes, self.stream.cu_stream(),
|
||
).map_err(|e| MLError::ModelError(format!(
|
||
"retain states chunk_start={chunk_start}: {e}"
|
||
)))?;
|
||
}
|
||
}
|
||
|
||
// D.8 TLOB val-path: single forward over the entire chunk at
|
||
// batch = chunk_len * n_windows. TLOB reads OFI[42..74) and writes
|
||
// TLOB[74..90) for each row independently — safe to batch.
|
||
//
|
||
// The kernel writes only the TLOB slot [SL_TLOB_START..+SL_TLOB_DIM);
|
||
// other slots remain as the chunked gather wrote them.
|
||
if self.tlob.is_some() {
|
||
let chunked_states = self.chunked_states_buf.as_mut().ok_or_else(|| {
|
||
MLError::ModelError("chunked_states_buf unexpectedly None".to_owned())
|
||
})?;
|
||
let tlob = self.tlob.as_mut().expect("checked is_some above");
|
||
tlob.forward(chunked_states, batch)
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"val TLOB forward chunk_start={chunk_start} batch={batch}: {e}"
|
||
)))?;
|
||
}
|
||
|
||
// Sync after Phase 1 gather — catch any gather errors
|
||
if chunk_idx == 0 {
|
||
self.stream.synchronize().map_err(|e| MLError::ModelError(format!(
|
||
"backtest chunk {chunk_idx}: sync after gather phase: {e}"
|
||
)))?;
|
||
}
|
||
|
||
// ── Phase 2+3: Q-values + raw branch C51 logits via QValueProvider ──
|
||
// ch_states_base is the device pointer to chunked_states_buf base —
|
||
// computed in Phase 1 and stable across the chunk (the Option does
|
||
// not reallocate). Reuse it instead of taking another borrow.
|
||
//
|
||
// Plan C T4 (2026-04-29): the val pipeline now also requires the
|
||
// trainer's `on_b_logits_buf` (branch-major raw C51 logits) for
|
||
// Thompson direction sampling in Phase 4 below. The
|
||
// compute_q_and_b_logits_to method DtoD-copies q_out_buf AND
|
||
// on_b_logits_buf per sub-batch iteration, keeping the caller's
|
||
// chunked_b_logits_buf consistent with chunked_q_values across
|
||
// the entire chunk batch (handles the case where batch >
|
||
// trainer.batch_size and the inner loop runs multiple sub-iters).
|
||
let states_ptr = ch_states_base;
|
||
let q_out_ptr = raw_f32_ptr(ch_q_values, &self.stream);
|
||
let b_logits_out_ptr = raw_f32_ptr(ch_b_logits, &self.stream);
|
||
// SP17 Commit C: V-head logits output pointer (chunked).
|
||
let v_logits_out_ptr = raw_f32_ptr(ch_v_logits, &self.stream);
|
||
match q_provider {
|
||
Some(ref mut qp) => {
|
||
qp.compute_q_and_b_logits_to(
|
||
states_ptr, batch, q_out_ptr, b_logits_out_ptr, v_logits_out_ptr,
|
||
)
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"QValueProvider chunk_start={chunk_start}: {e}"
|
||
)))?;
|
||
}
|
||
None => {
|
||
return Err(MLError::ModelError(
|
||
"evaluate_dqn_graphed requires a QValueProvider — pass Some(trainer) or use FusedTrainingCtx".to_owned()
|
||
));
|
||
}
|
||
}
|
||
|
||
let batch_i32 = batch as i32;
|
||
|
||
// ── Phase 4: Greedy action selection on chunked Q-values ─────────
|
||
let blocks_256 = ((batch + 255) / 256) as u32;
|
||
let launch_cfg = LaunchConfig {
|
||
grid_dim: (blocks_256.max(1), 1, 1),
|
||
block_dim: (256, 1, 1),
|
||
shared_mem_bytes: 0,
|
||
};
|
||
let null_portfolio: u64 = 0;
|
||
let eval_min_hold: i32 = 0; // hold enforcement removed — cost-driven
|
||
let eval_max_pos: f32 = 0.0;
|
||
// Plan C T4 (2026-04-29): direction-branch Thompson buffers.
|
||
// Mirror of the gpu_experience_collector wire-up (T3, 380d0ac0d).
|
||
// Eval mode (eps_start==eps_end==0) takes the argmax-E[Q] branch
|
||
// inside the kernel (no Philox draws), but the buffers must still
|
||
// be supplied with real production data — `feedback_no_stubs`.
|
||
// ch_b_logits is BRANCH-MAJOR `[batch, total_actions * num_atoms]`;
|
||
// the kernel reads only the first `batch * b0 * num_atoms` region
|
||
// (Branch 0 = direction). per_sample_support is the trainer's
|
||
// mapped-pinned `[max_batch, 4, 3]` stride-12 buffer (sized by
|
||
// trainer.batch_size ≥ chunked batch, safe to share). Atom
|
||
// positions live on the trainer (model-global `[B0, num_atoms]`).
|
||
let qp_for_buffers = q_provider.as_ref().ok_or_else(|| MLError::ModelError(
|
||
"QValueProvider required for action_select Thompson buffers".to_owned()
|
||
))?;
|
||
let per_sample_support_ptr = qp_for_buffers.per_sample_support_ptr();
|
||
let atom_positions_ptr = qp_for_buffers.atom_positions_ptr();
|
||
let n_atoms_i32 = qp_for_buffers.num_atoms() as i32;
|
||
unsafe {
|
||
self.stream
|
||
.launch_builder(as_kernel)
|
||
.arg(ch_q_values)
|
||
.arg(ch_actions)
|
||
.arg(ch_q_gaps)
|
||
.arg(&0.0_f32) // eps_start (greedy in backtest)
|
||
.arg(&0.0_f32) // eps_end
|
||
.arg(&0_i32) // current_epoch
|
||
.arg(&1_i32) // total_epochs
|
||
.arg(&batch_i32)
|
||
.arg(&b0)
|
||
.arg(&b1)
|
||
.arg(&b2)
|
||
.arg(&self.b3_size)
|
||
.arg(&self.q_gap_threshold)
|
||
.arg(&null_portfolio)
|
||
.arg(&eval_min_hold)
|
||
.arg(&eval_max_pos)
|
||
.arg(&1.0_f32) // eps_exp_mult (no-op at eps=0)
|
||
.arg(&1.0_f32) // eps_ord_mult (no-op at eps=0)
|
||
.arg(&1.0_f32) // eps_urg_mult (no-op at eps=0)
|
||
.arg(&(chunk_start as i32)) // timestep for stateless Philox RNG
|
||
.arg(&0u64) // per_sample_epsilon: NULL → use cosine schedule
|
||
.arg(&self.isv_signals_ptr) // ISV signals for adaptive hold
|
||
.arg(&0i32) // F6: contrarian_active — always off during backtest
|
||
.arg(ch_intent) // out_intent_mag: pure-policy magnitude intent diagnostic
|
||
.arg(ch_conviction) // out_conviction: direction-branch Kelly warmup signal
|
||
.arg(ch_magnitude_conviction) // out_magnitude_conviction: magnitude-branch Kelly warmup signal
|
||
// Plan C T2 amendment buffers (direction-branch Thompson):
|
||
.arg(ch_b_logits) // b_logits_dir: [batch, b0*n_atoms] used (rest is mag/ord/urg, untouched)
|
||
// SP17 Commit C: V-head logits row for softmax(V + A_centered) sampling.
|
||
.arg(ch_v_logits) // v_logits_dir: [batch, n_atoms]
|
||
.arg(&per_sample_support_ptr) // per_sample_support: [N, 4, 3] stride-12 mapped pinned (trainer-owned)
|
||
.arg(&atom_positions_ptr) // atom_positions: [b0, n_atoms] adaptive C51 positions (trainer-owned)
|
||
.arg(&n_atoms_i32) // n_atoms: C51 atom count (i32 to match `int` ABI)
|
||
// SP15 Phase 3.5.4.c (2026-05-07): plasticity_m_warm —
|
||
// pass 0.0f during eval/backtest. Plasticity injection
|
||
// is a TRAINING-only mechanism (it resets advantage-head
|
||
// weights to break out of stuck-in-flat regimes during
|
||
// experience collection); during deterministic eval the
|
||
// weights must remain frozen at their checkpoint values.
|
||
// m_warm = 0.0f makes the action_select fire-bar
|
||
// predicate dead, and PLASTICITY_WARM_BARS_REMAINING
|
||
// stays at its 0 sentinel through eval (no producer
|
||
// running) so the OR-gate degenerates to cooldown-only.
|
||
.arg(&0.0f32)
|
||
.launch(launch_cfg)
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"chunked action_select chunk_start={chunk_start}: {e}"
|
||
)))?;
|
||
}
|
||
|
||
// ── Phase 4b: Scatter chunked intent into window-major history
|
||
//
|
||
// experience_action_select writes ch_intent[s * n + w] for batch
|
||
// index i = s * n + w. scatter_intent_chunk copies into the
|
||
// window-major intent_mag_buf at [w * max_len + (start + s)].
|
||
// Diagnostic-only — writes are additive with the Phase 5 env
|
||
// kernel (which never touches intent_mag_buf).
|
||
let start_step_i32 = chunk_start as i32;
|
||
let chunk_len_i32 = chunk_len as i32;
|
||
let scatter_threads = n * chunk_len;
|
||
let scatter_blocks = ((scatter_threads as u32) + 255) / 256;
|
||
unsafe {
|
||
self.stream
|
||
.launch_builder(scatter_intent)
|
||
.arg(ch_intent)
|
||
.arg(intent_history)
|
||
.arg(&(n as i32))
|
||
.arg(&(self.max_len as i32))
|
||
.arg(&start_step_i32)
|
||
.arg(&chunk_len_i32)
|
||
.launch(LaunchConfig {
|
||
grid_dim: (scatter_blocks.max(1), 1, 1),
|
||
block_dim: (256, 1, 1),
|
||
shared_mem_bytes: 0,
|
||
})
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"scatter_intent_chunk chunk_start={chunk_start}: {e}"
|
||
)))?;
|
||
}
|
||
|
||
// ── Phase 4c: Parallel scatter of raw factored action picks ──
|
||
//
|
||
// Same scatter kernel, different src/dst — copies the kernel's
|
||
// RAW action_idx writes (chunked_actions_buf, before env_step
|
||
// touches anything) into picked_action_history_buf at the
|
||
// current chunk's step range. Pairs with actions_history_buf
|
||
// (POST-env_step actual_dir) so a downstream reader can compare
|
||
// the kernel's intent with the env's realisation across the
|
||
// full val window. Diagnostic-only.
|
||
unsafe {
|
||
self.stream
|
||
.launch_builder(scatter_intent)
|
||
.arg(ch_actions)
|
||
.arg(picked_action_history)
|
||
.arg(&(n as i32))
|
||
.arg(&(self.max_len as i32))
|
||
.arg(&start_step_i32)
|
||
.arg(&chunk_len_i32)
|
||
.launch(LaunchConfig {
|
||
grid_dim: (scatter_blocks.max(1), 1, 1),
|
||
block_dim: (256, 1, 1),
|
||
shared_mem_bytes: 0,
|
||
})
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"scatter picked_actions chunk_start={chunk_start}: {e}"
|
||
)))?;
|
||
}
|
||
|
||
// ── Phase 5: Batched env_step (single launch for entire chunk) ──
|
||
//
|
||
// All chunk_len steps processed in one kernel launch. The batched
|
||
// kernel loops over steps internally, reading actions directly from
|
||
// the chunked_actions buffer. Eliminates chunk_len individual launches
|
||
// (was 512 launches → 1 launch per chunk).
|
||
let env_blocks = ((n as u32) + 255) / 256;
|
||
let feature_dim_i32 = self.feature_dim as i32;
|
||
unsafe {
|
||
self.stream
|
||
.launch_builder(&self.env_batch_kernel)
|
||
.arg(&self.prices_buf.dev_ptr)
|
||
.arg(&self.window_lens_buf.dev_ptr)
|
||
.arg(&self.features_buf.dev_ptr) // regime-adaptive trail: ADX/CUSUM source
|
||
.arg(&feature_dim_i32)
|
||
.arg(ch_actions)
|
||
.arg(&self.portfolio_buf.dev_ptr)
|
||
.arg(&self.step_rewards_buf)
|
||
.arg(&self.step_returns_buf)
|
||
.arg(&self.done_buf)
|
||
.arg(&self.actions_history_buf)
|
||
.arg(&(n as i32))
|
||
.arg(&(self.max_len as i32))
|
||
.arg(&self.config.max_position)
|
||
.arg(&self.config.tx_cost_bps)
|
||
.arg(&self.config.spread_cost)
|
||
.arg(&start_step_i32)
|
||
.arg(&chunk_len_i32)
|
||
.arg(&b0)
|
||
.arg(&b1)
|
||
.arg(&b2)
|
||
.arg(&b3)
|
||
.arg(&self.config.holding_cost_rate)
|
||
.arg(&self.config.churn_threshold)
|
||
.arg(&self.config.churn_penalty_scale)
|
||
.arg(&self.config.contract_multiplier)
|
||
.arg(&self.config.margin_pct)
|
||
.arg(&self.config.opportunity_cost_scale)
|
||
// Kelly cap shared with training: stats buffer + ISV[12]=health
|
||
.arg(&self.kelly_stats_buf)
|
||
.arg(&self.isv_signals_ptr)
|
||
// Phase 3 unified-env scalars (NULL = no-op in val mode).
|
||
.arg(&0u64) // exploration_scale_ptr (NULL)
|
||
.arg(&0u64) // shaping_scale_ptr (NULL)
|
||
.arg(ch_conviction) // conviction_ptr: direction-branch Kelly warmup signal
|
||
.arg(ch_magnitude_conviction) // magnitude_conviction_ptr: magnitude-branch Kelly warmup signal
|
||
// SP21 T2.2 Phase 1 (2026-05-10): per-trade tape outputs.
|
||
// Real device buffers — same tape semantics as the
|
||
// single-step variant; the batched step loop writes
|
||
// per-step closes into the same per-window SoA slots
|
||
// (counter persists across the batched chunk via the
|
||
// shared `per_trade_count_buf`).
|
||
.arg(&self.per_trade_pnl_buf)
|
||
.arg(&self.per_trade_holding_bars_buf)
|
||
.arg(&self.per_trade_bar_index_buf)
|
||
.arg(&self.per_trade_dir_mag_buf)
|
||
.arg(&self.per_trade_count_buf)
|
||
.arg(&(MAX_TRADES_PER_WINDOW as i32))
|
||
// SP21 T2.2 Phase 1.5 (2026-05-11) — entry_q tracking.
|
||
// Production val path: real entry_q_state +
|
||
// chunked Q-values (`ch_q_values` written by
|
||
// `compute_q_and_b_logits_to` above; layout
|
||
// `[chunk_len, n_windows, total_actions]` matching
|
||
// the batch kernel's index formula `s * n * num_act
|
||
// + w * num_act + a`). Per-trade `predicted_q`
|
||
// populated for every close emitted in this chunk.
|
||
.arg(&self.entry_q_state_buf)
|
||
.arg(ch_q_values)
|
||
.arg(&num_actions_i32)
|
||
.arg(&self.per_trade_predicted_q_buf)
|
||
.launch(LaunchConfig {
|
||
grid_dim: (env_blocks.max(1), 1, 1),
|
||
block_dim: (256, 1, 1),
|
||
shared_mem_bytes: 0,
|
||
})
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"backtest_env_step_batch chunk {chunk_idx}: {e}"
|
||
)))?;
|
||
}
|
||
|
||
// ── Phase 6: Val plan_isv parity (task #94) ──────────────────────
|
||
//
|
||
// After Phase 5 has advanced portfolio_buf to the end-of-chunk
|
||
// state (last step = chunk_start + chunk_len - 1), compute:
|
||
// (a) plan_params[N, 6] for the LAST step's N windows via a
|
||
// targeted forward on ch_states[(chunk_len-1)*N..chunk_len*N]
|
||
// — this rewrites save_h_s2 = [N, SH2] for just those N
|
||
// rows (compute_q_values_to's earlier batched loop left
|
||
// save_h_s2 holding an arbitrary last-sub-batch slice, not
|
||
// aligned to the final-step N windows).
|
||
// (b) plan_state_isv: activation / deactivation / plan_isv
|
||
// compute so that next chunk's first launch_gather reads
|
||
// real plan_isv[7] at state positions [98..105).
|
||
//
|
||
// Using ch_q_values as scratch for the Q-output is safe — those
|
||
// Q-values were consumed by Phase 4 action-select and the Phase 5
|
||
// env step has already committed.
|
||
let plan_kernel = self.plan_state_isv_kernel.as_ref().ok_or_else(|| {
|
||
MLError::ModelError("plan_state_isv_kernel not loaded".to_owned())
|
||
})?;
|
||
let last_step = chunk_start + chunk_len - 1;
|
||
let last_step_row = chunk_len - 1;
|
||
let last_step_states_ptr = ch_states_base
|
||
+ (last_step_row * state_row_bytes) as u64;
|
||
let scratch_q_ptr = raw_f32_ptr(ch_q_values, &self.stream);
|
||
match q_provider {
|
||
Some(ref mut qp) => {
|
||
// (a.1) Forward on last-step N rows. Writes save_h_s2=[N,SH2].
|
||
qp.compute_q_values_to(last_step_states_ptr, n, scratch_q_ptr)
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"plan_params last-step forward chunk {chunk_idx}: {e}"
|
||
)))?;
|
||
// (a.2) Plan MLP on save_h_s2. Writes plan_params_buf[N,6].
|
||
let plan_params_ptr = {
|
||
let (p, g) = self.plan_params_buf.device_ptr(&self.stream);
|
||
let _ = ManuallyDrop::new(g);
|
||
p
|
||
};
|
||
qp.compute_plan_params(plan_params_ptr, n)
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"compute_plan_params chunk {chunk_idx}: {e}"
|
||
)))?;
|
||
}
|
||
None => {
|
||
return Err(MLError::ModelError(
|
||
"evaluate_dqn_graphed requires a QValueProvider for plan_params".to_owned()
|
||
));
|
||
}
|
||
}
|
||
|
||
// (b) Launch plan_state_isv kernel — updates plan_state in-place
|
||
// and writes plan_isv_buf[N, 7] consumed by the next chunk's
|
||
// first launch_gather() call.
|
||
//
|
||
// Fix 30 Stale-B (2026-05-02): the kernel previously took
|
||
// `features` + `feat_dim` to extract raw_close as
|
||
// `features[bar*feat_dim + 0]`, but post-Bug-1 that slot is
|
||
// z-normed log-return (NOT raw_close). Now passes the
|
||
// upload-once `prices` buffer (layout [n*max_len*4] OHLC,
|
||
// close at col 3) — the same source backtest_env_kernel.cu
|
||
// already uses for portfolio mark-to-market. `feat_dim` arg
|
||
// dropped because the kernel no longer needs feature stride.
|
||
let plan_blocks = ((n as u32) + 255) / 256;
|
||
let n_i32 = n as i32;
|
||
let last_step_i32 = last_step as i32;
|
||
let max_len_i32 = self.max_len as i32;
|
||
unsafe {
|
||
self.stream
|
||
.launch_builder(plan_kernel)
|
||
.arg(&self.plan_params_buf)
|
||
.arg(&self.plan_state_buf)
|
||
.arg(&self.portfolio_buf.dev_ptr)
|
||
.arg(&self.isv_signals_ptr)
|
||
.arg(&self.prices_buf.dev_ptr)
|
||
.arg(&last_step_i32)
|
||
.arg(&max_len_i32)
|
||
.arg(&self.plan_isv_buf)
|
||
.arg(&n_i32)
|
||
.launch(LaunchConfig {
|
||
grid_dim: (plan_blocks.max(1), 1, 1),
|
||
block_dim: (256, 1, 1),
|
||
shared_mem_bytes: 0,
|
||
})
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"backtest_plan_state_isv chunk {chunk_idx}: {e}"
|
||
)))?;
|
||
}
|
||
|
||
// Periodic sync: catch deferred CUDA errors without stalling
|
||
// every chunk. Errors in chunk N surface within 10 chunks.
|
||
if chunk_idx % 10 == 0 || chunk_idx + 1 == total_chunks {
|
||
self.stream.synchronize().map_err(|e| MLError::ModelError(format!(
|
||
"backtest chunk {chunk_idx}/{total_chunks} (steps {chunk_start}..{chunk_end}) sync: {e}"
|
||
)))?;
|
||
}
|
||
}
|
||
|
||
Ok(())
|
||
}
|
||
|
||
// ── PPO evaluation (pure CUDA) ───────────────────────────────────────────
|
||
|
||
/// Pure-GPU PPO evaluation — actor forward, softmax, exposure collapse, argmax.
|
||
///
|
||
/// Takes PPO actor weights directly, bypassing candle entirely for the
|
||
/// forward pass. The only GPU→CPU transfer is the final metrics readback.
|
||
///
|
||
/// Call `reset_mutable_state()` before re-running on the same evaluator
|
||
/// instance if portfolio state should be reinitialised.
|
||
pub fn evaluate_ppo(
|
||
&mut self,
|
||
actor_weights: &PpoActorWeightSet,
|
||
) -> Result<Vec<WindowMetrics>, MLError> {
|
||
let context = self.stream.context();
|
||
|
||
// Load precompiled PPO forward cubin
|
||
let ppo_module = context
|
||
.load_cubin(PPO_FORWARD_CUBIN.to_vec())
|
||
.map_err(|e| MLError::ModelError(format!("PPO forward module load: {e}")))?;
|
||
let ppo_kernel = ppo_module
|
||
.load_function("backtest_forward_ppo_kernel")
|
||
.map_err(|e| MLError::ModelError(format!("backtest_forward_ppo_kernel load: {e}")))?;
|
||
|
||
for step in 0..self.max_len {
|
||
// 1. Gather states via GPU kernel — writes into self.states_buf
|
||
self.launch_gather(step)?;
|
||
|
||
// 2. Launch PPO forward kernel: states → actions (pure CUDA, no candle)
|
||
// grid=(ceil(N/32),1,1), block=(32,1,1) — 32 independent windows/block
|
||
let n_i32 = self.n_windows as i32;
|
||
let ppo_cfg = LaunchConfig {
|
||
grid_dim: (((self.n_windows as u32) + 31) / 32, 1, 1),
|
||
block_dim: (32, 1, 1),
|
||
shared_mem_bytes: 0,
|
||
};
|
||
|
||
// Safety: argument order matches backtest_forward_ppo_kernel signature:
|
||
// states, pw1, pb1, pw2, pb2, pw3, pb3, out_actions, N, state_dim, num_actions
|
||
let state_dim_ppo_i32 = self.state_dim as i32;
|
||
let num_actions_ppo_i32 = 45_i32; // PPO_NUM_ACTIONS
|
||
unsafe {
|
||
self.stream
|
||
.launch_builder(&ppo_kernel)
|
||
.arg(&self.states_buf)
|
||
.arg(&actor_weights.pw1)
|
||
.arg(&actor_weights.pb1)
|
||
.arg(&actor_weights.pw2)
|
||
.arg(&actor_weights.pb2)
|
||
.arg(&actor_weights.pw3)
|
||
.arg(&actor_weights.pb3)
|
||
.arg(&self.actions_buf)
|
||
.arg(&n_i32)
|
||
.arg(&state_dim_ppo_i32)
|
||
.arg(&num_actions_ppo_i32)
|
||
.launch(ppo_cfg)
|
||
.map_err(|e| {
|
||
MLError::ModelError(format!(
|
||
"backtest_forward_ppo_kernel launch step {step}: {e}"
|
||
))
|
||
})?;
|
||
}
|
||
|
||
// 3. Launch env step kernel
|
||
self.launch_env_step(step)?;
|
||
}
|
||
|
||
// Synchronous launch+consume — `evaluate_ppo` is a single-shot path.
|
||
self.launch_metrics_and_record_event()?;
|
||
self.consume_metrics_after_event()
|
||
}
|
||
|
||
// ── Supervised evaluation (hybrid: candle forward + CUDA signal→action) ──
|
||
|
||
/// GPU-accelerated supervised model evaluation.
|
||
///
|
||
/// The model forward pass uses candle (unavoidable for complex architectures
|
||
/// like TFT, Mamba, etc.), but signal→action conversion is a pure CUDA kernel
|
||
/// replacing the candle `signal_to_action_scores` + `argmax` pipeline.
|
||
///
|
||
/// # Arguments
|
||
/// * `forward_fn` — model forward: `&CudaSlice<f32>[N * state_dim]` → `CudaSlice<f32>[N]` predictions
|
||
/// * `high_threshold_bps` — strong signal threshold (actions 0 and 4)
|
||
/// * `low_threshold_bps` — mild signal threshold (actions 1 and 3)
|
||
/// * `portfolio_dim` — must be 24 (8 portfolio + 16 multi-timeframe)
|
||
pub fn evaluate_supervised<F>(
|
||
&mut self,
|
||
forward_fn: &F,
|
||
high_threshold_bps: f32,
|
||
low_threshold_bps: f32,
|
||
portfolio_dim: usize,
|
||
) -> Result<Vec<WindowMetrics>, MLError>
|
||
where
|
||
F: Fn(&CudaSlice<f32>, usize, usize) -> Result<CudaSlice<f32>, MLError>,
|
||
{
|
||
if portfolio_dim != self.portfolio_dim {
|
||
return Err(MLError::ConfigError(format!(
|
||
"evaluate_supervised: portfolio_dim={portfolio_dim} != expected {}",
|
||
self.portfolio_dim
|
||
)));
|
||
}
|
||
|
||
let context = self.stream.context();
|
||
|
||
// Load precompiled supervised signal cubin
|
||
let sig_module = context
|
||
.load_cubin(SUPERVISED_SIGNAL_CUBIN.to_vec())
|
||
.map_err(|e| MLError::ModelError(format!("supervised signal module load: {e}")))?;
|
||
let sig_kernel = sig_module
|
||
.load_function("signal_to_action_kernel")
|
||
.map_err(|e| MLError::ModelError(format!("signal_to_action_kernel load: {e}")))?;
|
||
|
||
let state_dim = self.state_dim;
|
||
|
||
for step in 0..self.max_len {
|
||
// 1. Gather state CudaSlice via GPU kernel
|
||
let states = self.gather_states(step, portfolio_dim, &())?;
|
||
|
||
// 2. Model forward pass (pure cudarc, on-device)
|
||
let pred_slice = forward_fn(&states, self.n_windows, state_dim)?;
|
||
|
||
// 3. Launch signal_to_action_kernel: predictions → actions (pure CUDA)
|
||
let n_i32 = self.n_windows as i32;
|
||
let grid = ((self.n_windows + 255) / 256) as u32;
|
||
let sig_cfg = LaunchConfig {
|
||
grid_dim: (grid.max(1), 1, 1),
|
||
block_dim: (256, 1, 1),
|
||
shared_mem_bytes: 0,
|
||
};
|
||
|
||
// Safety: argument order matches signal_to_action_kernel:
|
||
// predictions, out_actions, high_threshold_bps, low_threshold_bps, N
|
||
unsafe {
|
||
self.stream
|
||
.launch_builder(&sig_kernel)
|
||
.arg(&pred_slice)
|
||
.arg(&self.actions_buf)
|
||
.arg(&high_threshold_bps)
|
||
.arg(&low_threshold_bps)
|
||
.arg(&n_i32)
|
||
.launch(sig_cfg)
|
||
.map_err(|e| {
|
||
MLError::ModelError(format!(
|
||
"signal_to_action_kernel launch step {step}: {e}"
|
||
))
|
||
})?;
|
||
}
|
||
|
||
// 4. Launch env step kernel
|
||
self.launch_env_step(step)?;
|
||
}
|
||
|
||
// Synchronous launch+consume — `evaluate_supervised` is a single-shot path.
|
||
self.launch_metrics_and_record_event()?;
|
||
self.consume_metrics_after_event()
|
||
}
|
||
|
||
// ── DQN forward helpers ──────────────────────────────────────────────
|
||
|
||
/// Lazily initialise action-selection kernels, RNG states, and chunked
|
||
/// Q-value / state / action buffers.
|
||
///
|
||
/// Called once on first `evaluate_dqn_graphed()`. The forward pass itself
|
||
/// is handled by the `QValueProvider` (trainer's CUDA-Graphed cuBLAS);
|
||
/// this method only sets up the action-selection and env-step infrastructure.
|
||
fn ensure_action_select_ready(&mut self, dqn_cfg: &DqnBacktestConfig) -> Result<(), MLError> {
|
||
if self.action_select_kernel.is_some() {
|
||
return Ok(());
|
||
}
|
||
|
||
let n = self.n_windows;
|
||
let total_actions = dqn_cfg.branch_0_size + dqn_cfg.branch_1_size + dqn_cfg.branch_2_size + dqn_cfg.branch_3_size;
|
||
|
||
// Compile experience kernels for action_select
|
||
let context = self.stream.context();
|
||
let dqn_module = context
|
||
.load_cubin(BACKTEST_DQN_CUBIN.to_vec())
|
||
.map_err(|e| MLError::ModelError(format!("backtest DQN module load: {e}")))?;
|
||
let as_kernel = dqn_module
|
||
.load_function("experience_action_select")
|
||
.map_err(|e| MLError::ModelError(format!("experience_action_select load: {e}")))?;
|
||
let scatter_intent = dqn_module
|
||
.load_function("scatter_intent_chunk")
|
||
.map_err(|e| MLError::ModelError(format!("scatter_intent_chunk load: {e}")))?;
|
||
|
||
// Val plan_isv parity (task #94, 2026-04-24): load the two kernels
|
||
// from the new backtest_plan_kernel cubin. See docstring in
|
||
// backtest_plan_kernel.cu for the activation / isv-compute /
|
||
// diagnostic-reduce pipeline.
|
||
let plan_module = context
|
||
.load_cubin(BACKTEST_PLAN_CUBIN.to_vec())
|
||
.map_err(|e| MLError::ModelError(format!("backtest_plan module load: {e}")))?;
|
||
let plan_state_isv_kernel_fn = plan_module
|
||
.load_function("backtest_plan_state_isv")
|
||
.map_err(|e| MLError::ModelError(format!("backtest_plan_state_isv load: {e}")))?;
|
||
let plan_diag_reduce_kernel_fn = plan_module
|
||
.load_function("backtest_plan_diag_reduce")
|
||
.map_err(|e| MLError::ModelError(format!("backtest_plan_diag_reduce load: {e}")))?;
|
||
self.plan_state_isv_kernel = Some(plan_state_isv_kernel_fn);
|
||
self.plan_diag_reduce_kernel = Some(plan_diag_reduce_kernel_fn);
|
||
|
||
// ── Chunked buffers (batch_size = n_windows * CHUNK_SIZE) ────────────
|
||
let chunk = DQN_BACKTEST_CHUNK_SIZE;
|
||
let cn = n * chunk; // chunked batch dimension
|
||
|
||
let ch_q_values = self.stream.alloc_zeros::<f32>(cn * total_actions)
|
||
.map_err(|e| MLError::ModelError(format!("alloc chunked q_values: {e}")))?;
|
||
|
||
// Plan C T4 (2026-04-29): chunked branch-major raw C51 logits buffer.
|
||
// Sized [cn, total_actions * num_atoms] — DtoD-populated each chunk
|
||
// by `compute_q_and_b_logits_to`. Direction-branch slice consumed by
|
||
// `experience_action_select` Thompson direction sampling. Padded by
|
||
// 32*3 floats to mirror the trainer's allocation pattern (cuBLAS
|
||
// tail safety; see gpu_dqn_trainer.rs:9759).
|
||
let total_branch_atoms_eval = total_actions * dqn_cfg.num_atoms;
|
||
let ch_b_logits = self.stream.alloc_zeros::<f32>(cn * total_branch_atoms_eval + 32 * 3)
|
||
.map_err(|e| MLError::ModelError(format!("alloc chunked b_logits: {e}")))?;
|
||
|
||
// SP17 Commit C (2026-05-08): chunked V-head logits buffer, sized
|
||
// `[cn, num_atoms]` sample-major. DtoD-populated each chunk by
|
||
// `compute_q_and_b_logits_to` and consumed by the Thompson direction
|
||
// selector in `experience_action_select` (closes the
|
||
// `softmax(V + A_centered)` contract gap). Same `+ 32*3` cuBLAS tail
|
||
// safety pad as ch_b_logits.
|
||
let ch_v_logits = self.stream.alloc_zeros::<f32>(cn * dqn_cfg.num_atoms + 32 * 3)
|
||
.map_err(|e| MLError::ModelError(format!("alloc chunked v_logits: {e}")))?;
|
||
|
||
let state_dim_padded = (self.state_dim + 127) & !127;
|
||
let ch_states_buf = self.stream.alloc_zeros::<f32>(cn * state_dim_padded)
|
||
.map_err(|e| MLError::ModelError(format!("alloc chunked states_buf: {e}")))?;
|
||
// SP21 T2.2 Phase 6.5 (2026-05-11): full-window retained-states
|
||
// buffer (mapped-pinned per `feedback_no_htod_htoh_only_mapped_
|
||
// pinned`). Size = max_len × n_windows × state_dim_padded.
|
||
// Layout matches the chunked layout `[step, window, dim]` so
|
||
// each chunk's `ch_states_buf` DtoD-copies into this buffer's
|
||
// `dev_ptr` at offset `chunk_start × n_windows × state_dim_
|
||
// padded` with no transpose. Host-side hindsight injection
|
||
// reads via `host_ptr` (zero-copy) after an eval stream sync.
|
||
// Lazy-init alongside other action-select scratch; freed when
|
||
// the evaluator is dropped (via `MappedF32Buffer`'s `Drop`).
|
||
let retained_states_buf = unsafe {
|
||
MappedF32Buffer::new(self.max_len * n * state_dim_padded)
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"alloc retained_states_buf (mapped-pinned): {e}"
|
||
)))?
|
||
};
|
||
let ch_actions_buf = self.stream.alloc_zeros::<i32>(cn)
|
||
.map_err(|e| MLError::ModelError(format!("alloc chunked actions_buf: {e}")))?;
|
||
let ch_intent_mag_buf = self.stream.alloc_zeros::<i32>(cn)
|
||
.map_err(|e| MLError::ModelError(format!("alloc chunked intent_mag_buf: {e}")))?;
|
||
|
||
// Intent-magnitude history: parallel layout to actions_history_buf.
|
||
let intent_mag_buf = self.stream.alloc_zeros::<i32>(n * self.max_len)
|
||
.map_err(|e| MLError::ModelError(format!("alloc intent_mag_buf: {e}")))?;
|
||
|
||
// Picked-action history: parallel layout to actions_history_buf,
|
||
// captures the kernel's RAW factored action_idx PRE-env_step. Filled
|
||
// by scattering chunked_actions_buf into per-step slots (reuses the
|
||
// scatter_intent_chunk kernel — same scatter semantics, different
|
||
// src/dst). Diagnostic-only.
|
||
let picked_action_history_buf = self.stream.alloc_zeros::<i32>(n * self.max_len)
|
||
.map_err(|e| MLError::ModelError(format!("alloc picked_action_history_buf: {e}")))?;
|
||
|
||
// Mapped-pinned mirrors for the two diagnostic history buffers above.
|
||
// Filled via `memcpy_dtod_async` on the eval stream before
|
||
// `eval_done_event` is recorded; read by the per-direction / intent /
|
||
// picked distribution helpers after the event syncs. See
|
||
// `actions_history_pinned` in the struct field docstring.
|
||
let intent_mag_pinned = unsafe {
|
||
MappedI32Buffer::new(n * self.max_len)
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"intent_mag mapped-pinned alloc: {e}"
|
||
)))?
|
||
};
|
||
let picked_action_history_pinned = unsafe {
|
||
MappedI32Buffer::new(n * self.max_len)
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"picked_action_history mapped-pinned alloc: {e}"
|
||
)))?
|
||
};
|
||
|
||
let ch_q_gaps = self.stream.alloc_zeros::<f32>(cn)
|
||
.map_err(|e| MLError::ModelError(format!("alloc chunked q_gaps: {e}")))?;
|
||
|
||
/* Per-sample conviction for adaptive Kelly warmup floor. Layout matches
|
||
* chunked_actions/chunked_q_gaps — [chunk_len * n_windows]. */
|
||
let ch_conviction = self.stream.alloc_zeros::<f32>(cn)
|
||
.map_err(|e| MLError::ModelError(format!("alloc chunked conviction: {e}")))?;
|
||
/* Parallel magnitude-branch conviction — same shape, threaded into
|
||
* action_select + env_step_batch alongside the direction-branch
|
||
* conviction so the Kelly cap's max(dir, mag) composer can relax
|
||
* for confident magnitude picks. */
|
||
let ch_magnitude_conviction = self.stream.alloc_zeros::<f32>(cn)
|
||
.map_err(|e| MLError::ModelError(format!("alloc chunked magnitude conviction: {e}")))?;
|
||
|
||
info!(
|
||
"GpuBacktestEvaluator: action-select initialised — {} windows, \
|
||
branches=[{},{},{},{}], chunk_size={chunk}, chunked_batch={cn}",
|
||
n, dqn_cfg.branch_0_size, dqn_cfg.branch_1_size, dqn_cfg.branch_2_size, dqn_cfg.branch_3_size,
|
||
);
|
||
|
||
self.b0_size = dqn_cfg.branch_0_size as i32;
|
||
self.b1_size = dqn_cfg.branch_1_size as i32;
|
||
self.b2_size = dqn_cfg.branch_2_size as i32;
|
||
self.b3_size = dqn_cfg.branch_3_size as i32;
|
||
self.action_select_kernel = Some(as_kernel);
|
||
|
||
self.chunked_q_values = Some(ch_q_values);
|
||
self.chunked_b_logits_buf = Some(ch_b_logits);
|
||
self.chunked_v_logits_buf = Some(ch_v_logits);
|
||
self.chunked_states_buf = Some(ch_states_buf);
|
||
self.retained_states_buf = Some(retained_states_buf);
|
||
self.chunked_actions_buf = Some(ch_actions_buf);
|
||
self.chunked_intent_mag_buf = Some(ch_intent_mag_buf);
|
||
self.chunked_q_gaps_buf = Some(ch_q_gaps);
|
||
self.intent_mag_buf = Some(intent_mag_buf);
|
||
self.picked_action_history_buf = Some(picked_action_history_buf);
|
||
self.intent_mag_pinned = Some(intent_mag_pinned);
|
||
self.picked_action_history_pinned = Some(picked_action_history_pinned);
|
||
self.scatter_intent_kernel = Some(scatter_intent);
|
||
self.chunked_conviction_buf = Some(ch_conviction);
|
||
self.chunked_magnitude_conviction_buf = Some(ch_magnitude_conviction);
|
||
|
||
Ok(())
|
||
}
|
||
|
||
/// Launch the `gather_states` GPU kernel for a given step.
|
||
///
|
||
/// Writes `[n_windows, state_dim]` floats into `self.states_buf`.
|
||
/// Does NOT produce a Candle Tensor — the pure-CUDA `evaluate_dqn` and
|
||
/// `evaluate_ppo` paths read directly from `states_buf`.
|
||
fn launch_gather(&self, step: usize) -> Result<(), MLError> {
|
||
// Shared memory: 256 threads × 8 portfolio f32 × 4 bytes = 8 KB
|
||
let grid = ((self.n_windows + 255) / 256) as u32;
|
||
let launch_cfg = LaunchConfig {
|
||
grid_dim: (grid.max(1), 1, 1),
|
||
block_dim: (256, 1, 1),
|
||
shared_mem_bytes: 256 * 8 * std::mem::size_of::<f32>() as u32,
|
||
};
|
||
let n_windows_i32 = self.n_windows as i32;
|
||
let max_len_i32 = self.max_len as i32;
|
||
let feat_dim_i32 = self.feature_dim as i32;
|
||
let step_i32 = step as i32;
|
||
let padded_sd_i32 = ml_core::state_layout::STATE_DIM_PADDED as i32;
|
||
|
||
// Safety: argument order matches `backtest_state_gather` signature exactly:
|
||
// features, portfolio, plan_isv, states_out, n_windows, max_len, feat_dim,
|
||
// current_step, padded_sd, aux_dir_prob_per_env (SP22 H6).
|
||
// SP22 H6 Phase 1 A3: NULL → kernel uses 0.5 sentinel.
|
||
let aux_dir_prob_null: u64 = 0;
|
||
unsafe {
|
||
self.stream
|
||
.launch_builder(&self.gather_kernel)
|
||
.arg(&self.features_buf.dev_ptr)
|
||
.arg(&self.portfolio_buf.dev_ptr)
|
||
.arg(&self.plan_isv_buf)
|
||
.arg(&self.states_buf)
|
||
.arg(&n_windows_i32)
|
||
.arg(&max_len_i32)
|
||
.arg(&feat_dim_i32)
|
||
.arg(&step_i32)
|
||
.arg(&padded_sd_i32)
|
||
.arg(&aux_dir_prob_null) // SP22 H6 A3: NULL → 0.5 sentinel
|
||
.launch(launch_cfg)
|
||
.map_err(|e| {
|
||
MLError::ModelError(format!("backtest_state_gather launch step {step}: {e}"))
|
||
})?;
|
||
}
|
||
Ok(())
|
||
}
|
||
|
||
/// Launch `backtest_state_gather_chunk` — batches an entire chunk of
|
||
/// `chunk_len` steps × `n_windows` rows into a single kernel launch,
|
||
/// writing directly into `chunked_states_out` in `[chunk_len, N, padded_sd]`
|
||
/// layout. Replaces the prior per-step `launch_gather` + DtoD copy loop.
|
||
///
|
||
/// Within a chunk `portfolio_buf` and `plan_isv_buf` are constant — the
|
||
/// env-batch kernel updates them only at chunk boundary. So the gather
|
||
/// for all `chunk_len` steps reads identical per-window portfolio /
|
||
/// plan_isv inputs and writes a unique state row per (window, step_offset)
|
||
/// pair. See the docstring on the kernel itself for layout details.
|
||
///
|
||
/// `chunked_states_dev_ptr` must be the device pointer to the start of
|
||
/// the chunk's portion of `chunked_states_buf` (i.e. base; the kernel
|
||
/// writes rows `0..chunk_len * n_windows`). Caller is responsible for
|
||
/// `chunk_len <= DQN_BACKTEST_CHUNK_SIZE` (the buffer was sized for
|
||
/// `cn = DQN_BACKTEST_CHUNK_SIZE * n_windows` rows).
|
||
fn launch_gather_chunk(
|
||
&self,
|
||
chunked_states_dev_ptr: u64,
|
||
start_step: usize,
|
||
chunk_len: usize,
|
||
) -> Result<(), MLError> {
|
||
let total_threads = (self.n_windows * chunk_len) as u32;
|
||
// Block size 256 — same as single-step gather, well within L40S
|
||
// 1024-thread per-block limit. Grid sized to cover all
|
||
// (window, step_offset) pairs.
|
||
let blocks = (total_threads + 255) / 256;
|
||
let launch_cfg = LaunchConfig {
|
||
grid_dim: (blocks.max(1), 1, 1),
|
||
block_dim: (256, 1, 1),
|
||
shared_mem_bytes: 0,
|
||
};
|
||
let n_windows_i32 = self.n_windows as i32;
|
||
let max_len_i32 = self.max_len as i32;
|
||
let feat_dim_i32 = self.feature_dim as i32;
|
||
let start_step_i32 = start_step as i32;
|
||
let chunk_len_i32 = chunk_len as i32;
|
||
let padded_sd_i32 = ml_core::state_layout::STATE_DIM_PADDED as i32;
|
||
|
||
// Safety: argument order matches `backtest_state_gather_chunk` signature exactly:
|
||
// features, portfolio, plan_isv, chunked_states_out (raw u64 ptr),
|
||
// n_windows, max_len, feat_dim, start_step, chunk_len, padded_sd,
|
||
// aux_dir_prob_per_env (SP22 H6).
|
||
// SP22 H6 Phase 1 A3: NULL → kernel uses 0.5 sentinel.
|
||
let aux_dir_prob_null: u64 = 0;
|
||
unsafe {
|
||
self.stream
|
||
.launch_builder(&self.gather_chunk_kernel)
|
||
.arg(&self.features_buf.dev_ptr)
|
||
.arg(&self.portfolio_buf.dev_ptr)
|
||
.arg(&self.plan_isv_buf)
|
||
.arg(&chunked_states_dev_ptr)
|
||
.arg(&n_windows_i32)
|
||
.arg(&max_len_i32)
|
||
.arg(&feat_dim_i32)
|
||
.arg(&start_step_i32)
|
||
.arg(&chunk_len_i32)
|
||
.arg(&padded_sd_i32)
|
||
.arg(&aux_dir_prob_null) // SP22 H6 A3: NULL → 0.5 sentinel
|
||
.launch(launch_cfg)
|
||
.map_err(|e| MLError::ModelError(format!(
|
||
"backtest_state_gather_chunk launch start_step={start_step} chunk_len={chunk_len}: {e}"
|
||
)))?;
|
||
}
|
||
Ok(())
|
||
}
|
||
|
||
/// Launch `backtest_env_step` kernel for a given step on `self.stream`.
|
||
fn launch_env_step(&self, step: usize) -> Result<(), MLError> {
|
||
// Shared memory: 256 threads × PORTFOLIO_STATE_SIZE(8) f32 × 4 bytes = 8 KB
|
||
let grid = ((self.n_windows + 255) / 256) as u32;
|
||
let env_cfg = LaunchConfig {
|
||
grid_dim: (grid.max(1), 1, 1),
|
||
block_dim: (256, 1, 1),
|
||
shared_mem_bytes: 256 * 8 * std::mem::size_of::<f32>() as u32,
|
||
};
|
||
let n_win_i32 = self.n_windows as i32;
|
||
let max_len_i32 = self.max_len as i32;
|
||
let step_i32 = step as i32;
|
||
// Branch sizes for factored action decoding (set from DqnBacktestConfig in ensure_action_select_ready).
|
||
let b0_i32 = self.b0_size;
|
||
let b1_i32 = self.b1_size;
|
||
let b2_i32 = self.b2_size;
|
||
let feature_dim_i32 = self.feature_dim as i32;
|
||
|
||
unsafe {
|
||
self.stream
|
||
.launch_builder(&self.env_kernel)
|
||
.arg(&self.prices_buf.dev_ptr)
|
||
.arg(&self.window_lens_buf.dev_ptr)
|
||
.arg(&self.features_buf.dev_ptr) // regime-adaptive trail: ADX/CUSUM source
|
||
.arg(&feature_dim_i32)
|
||
.arg(&self.actions_buf)
|
||
.arg(&self.portfolio_buf.dev_ptr)
|
||
.arg(&self.step_rewards_buf)
|
||
.arg(&self.step_returns_buf)
|
||
.arg(&self.done_buf)
|
||
.arg(&self.actions_history_buf)
|
||
.arg(&n_win_i32)
|
||
.arg(&max_len_i32)
|
||
.arg(&self.config.max_position)
|
||
.arg(&self.config.tx_cost_bps)
|
||
.arg(&self.config.spread_cost)
|
||
.arg(&step_i32)
|
||
.arg(&b0_i32)
|
||
.arg(&b1_i32)
|
||
.arg(&b2_i32)
|
||
.arg(&self.b3_size)
|
||
.arg(&self.config.holding_cost_rate)
|
||
.arg(&self.config.churn_threshold)
|
||
.arg(&self.config.churn_penalty_scale)
|
||
.arg(&self.config.contract_multiplier)
|
||
.arg(&self.config.margin_pct)
|
||
.arg(&self.config.opportunity_cost_scale)
|
||
// Kelly cap shared with training: stats buffer + ISV[12]=health
|
||
.arg(&self.kelly_stats_buf)
|
||
.arg(&self.isv_signals_ptr)
|
||
// Phase 3 unified-env scalars (NULL = no-op in val mode).
|
||
// ABI alignment with future unified_env_kernel; saboteur
|
||
// and shaping are not present in val so the values would
|
||
// be 0.0 anyway.
|
||
.arg(&0u64) // exploration_scale_ptr (NULL)
|
||
.arg(&0u64) // shaping_scale_ptr (NULL)
|
||
.arg(&0u64) // conviction_ptr (NULL) — fallback 1.0 in-kernel
|
||
.arg(&0u64) // magnitude_conviction_ptr (NULL) — fallback 0.0 in-kernel; max(dir, mag) → dir alone
|
||
// SP21 T2.2 Phase 1 (2026-05-10): per-trade tape outputs.
|
||
// Real device buffers wired — kernel emits per-trade
|
||
// records at every trade close into the SoA buffers, with
|
||
// race-free per-window counter increment. Buffers are
|
||
// reset to zero at the start of each eval window via
|
||
// `reset_per_trade_tape()`. Read host-side via
|
||
// `read_per_trade_tape()` after the eval window completes.
|
||
.arg(&self.per_trade_pnl_buf)
|
||
.arg(&self.per_trade_holding_bars_buf)
|
||
.arg(&self.per_trade_bar_index_buf)
|
||
.arg(&self.per_trade_dir_mag_buf)
|
||
.arg(&self.per_trade_count_buf)
|
||
.arg(&(MAX_TRADES_PER_WINDOW as i32))
|
||
// SP21 T2.2 Phase 1.5 (2026-05-11) — entry_q tracking.
|
||
// The single-step `evaluate()` path (used by hyperopt /
|
||
// smoke tests / non-DQN callers) receives only action
|
||
// indices from forward_fn — Q-values aren't exposed at
|
||
// the closure boundary. Pass real entry_q_state buffer
|
||
// (kernel still snapshots the state at potential close)
|
||
// and NULL q_values_per_window (kernel skips the
|
||
// open-time write). Per-trade `predicted_q` for trades
|
||
// closed by this path stays at buffer-init 0.0.
|
||
// Production val pipeline goes through
|
||
// `evaluate_dqn_graphed_async` → `submit_dqn_step_loop_cublas`,
|
||
// which DOES wire real q_values; see the batch launcher
|
||
// arg block in that method.
|
||
.arg(&self.entry_q_state_buf)
|
||
.arg(&0u64) // q_values_per_window = NULL
|
||
.arg(&0i32) // num_actions = 0 (NULL guard short-circuits)
|
||
.arg(&self.per_trade_predicted_q_buf)
|
||
.launch(env_cfg)
|
||
.map_err(|e| {
|
||
MLError::ModelError(format!("backtest_env_step launch step {step}: {e}"))
|
||
})?;
|
||
}
|
||
Ok(())
|
||
}
|
||
|
||
/// Launch the metrics reduction kernel and queue async DtoD copies of the
|
||
/// per-step action history buffers into their mapped-pinned mirrors. Records
|
||
/// `eval_done_event` on `self.stream` after the queued work, then returns
|
||
/// **without** any host-side wait.
|
||
///
|
||
/// This is the launch half of the perf-fix split (2026-04-28). The companion
|
||
/// `consume_metrics_after_event` calls `event.synchronize()` and reads through
|
||
/// the mapped-pinned `host_ptr`. Splitting the two lets the host CPU thread
|
||
/// submit the next epoch's training kernels on `cuda_stream` while eval drains
|
||
/// asynchronously on `validation_stream`.
|
||
///
|
||
/// Shared by all evaluation paths via the thin `evaluate_*` wrappers.
|
||
pub fn launch_metrics_and_record_event(&mut self) -> Result<(), MLError> {
|
||
// Shared memory: 11 reduction arrays × 256 threads × 4 bytes = 11 KB
|
||
// + 4096 floats for sort scratch / boundary data = 16 KB
|
||
// Total = 27 KB (well within the 48 KB L1/shmem limit).
|
||
// SP15 Phase 1.1.b shrunk this from 6 → 5 reduction arrays after the
|
||
// s_sq_sum reduction (sharpe-only) was removed; sharpe is now computed
|
||
// by the dedicated `sharpe_per_bar_kernel` launched per-window below.
|
||
// 2026-05-09 — per-regime WR adds 6 reduction arrays
|
||
// (s_trades_T/R/V + s_wins_T/R/V), bumping reduction-array count
|
||
// 5 → 11. The boundary buffer slot count also grows 7 → 9 to
|
||
// carry the open-bar regime + last-open-bar regime per chunk;
|
||
// those still live within the 4096-float `s_sorted` reservation.
|
||
let shmem_bytes = (256_u32 * 11 + 4096) * std::mem::size_of::<f32>() as u32;
|
||
let metrics_cfg = LaunchConfig {
|
||
grid_dim: (self.n_windows as u32, 1, 1),
|
||
block_dim: (256, 1, 1),
|
||
shared_mem_bytes: shmem_bytes,
|
||
};
|
||
let n_win_i32 = self.n_windows as i32;
|
||
let max_len_i32 = self.max_len as i32;
|
||
let annualization: f32 = self.annualization_factor;
|
||
|
||
// Safety: argument order matches `compute_backtest_metrics` signature exactly.
|
||
let num_actions_i32 = self.b0_size;
|
||
let order_actions_i32 = self.b1_size;
|
||
let urgency_actions_i32 = self.b2_size;
|
||
// 2026-05-09 — per-regime WR reads `features[bar_idx*feature_dim+40]`
|
||
// (ADX-norm) per bar. The features buffer was uploaded to GPU via
|
||
// mapped-pinned at construct time; pass its device pointer + the
|
||
// feature_dim scalar so the kernel can index per-window per-bar.
|
||
let feature_dim_i32 = self.feature_dim as i32;
|
||
// Pass mapped-pinned device pointer as kernel arg (f32* slot). The
|
||
// kernel writes through `dev_ptr`; the host reads `host_ptr` after
|
||
// sync. Same pattern as the IQN total_loss readback in `gpu_iqn_head.rs`.
|
||
let metrics_dev_ptr = self.metrics_buf.dev_ptr;
|
||
unsafe {
|
||
self.stream
|
||
.launch_builder(&self.metrics_kernel)
|
||
.arg(&self.step_returns_buf)
|
||
.arg(&self.portfolio_buf.dev_ptr)
|
||
.arg(&self.window_lens_buf.dev_ptr)
|
||
.arg(&self.actions_history_buf)
|
||
.arg(&self.features_buf.dev_ptr)
|
||
.arg(&feature_dim_i32)
|
||
.arg(&metrics_dev_ptr)
|
||
.arg(&n_win_i32)
|
||
.arg(&max_len_i32)
|
||
.arg(&annualization)
|
||
.arg(&num_actions_i32)
|
||
.arg(&order_actions_i32)
|
||
.arg(&urgency_actions_i32)
|
||
.launch(metrics_cfg)
|
||
.map_err(|e| MLError::ModelError(format!("compute_backtest_metrics launch: {e}")))?;
|
||
}
|
||
|
||
// SP15 Phase 1.1.b — dedicated per-bar Sharpe kernel.
|
||
// The fused metrics kernel above no longer computes Sharpe; the spec
|
||
// calls for the unified `sharpe_per_bar_kernel` to be the sole Sharpe
|
||
// producer for both train and val paths (see launcher docstring at
|
||
// gpu_dqn_trainer.rs::launch_sp15_sharpe_per_bar). The kernel is
|
||
// single-block by design (BLOCK=256, `if (blockIdx.x != 0) return;`)
|
||
// so per-window invocation requires N launches against
|
||
// `step_returns_buf + w*max_len*sizeof::<f32>()`. n_windows is small
|
||
// (typically 5) — the launches sit on the same stream as the metrics
|
||
// kernel and inherit its event-record dependency. No host wait
|
||
// between them; no extra synchronisation. Annualisation is applied
|
||
// host-side in `consume_metrics_after_event`.
|
||
{
|
||
let (step_returns_base, ptr_guard) =
|
||
self.step_returns_buf.device_ptr(&self.stream);
|
||
// Pin the device-ptr lifetime for the duration of the launches
|
||
// below — same precedent as `raw_f32_ptr` further down. The
|
||
// launches are queued synchronously to the stream, and the
|
||
// slice itself is owned by `self`, so the ptr remains valid
|
||
// past the borrow.
|
||
let _no_drop = ManuallyDrop::new(ptr_guard);
|
||
let elem_bytes = std::mem::size_of::<f32>() as u64;
|
||
let stride_bytes = (self.max_len as u64) * elem_bytes;
|
||
let sharpe_dev_base = self.sharpe_buf.dev_ptr;
|
||
// Read window lengths from the mapped-pinned host alias — the
|
||
// buffer was populated at construct time via `write_from_slice`
|
||
// and is otherwise read-only, so no sync barrier is required.
|
||
// `read_all()` performs `read_volatile` per entry to defeat any
|
||
// stale CPU cache line from before the construction store.
|
||
let window_lens_host: Vec<i32> = self.window_lens_buf.read_all();
|
||
for w in 0..self.n_windows {
|
||
let pnl_dev = step_returns_base + (w as u64) * stride_bytes;
|
||
let out_dev = sharpe_dev_base + (w as u64) * 3u64 * elem_bytes;
|
||
let n_w = window_lens_host[w];
|
||
super::gpu_dqn_trainer::launch_sp15_sharpe_per_bar(
|
||
&self.stream,
|
||
pnl_dev,
|
||
n_w,
|
||
out_dev,
|
||
)?;
|
||
}
|
||
}
|
||
|
||
// ── SP15 Wave 3b — val-cost-streams launch sequence ─────────────────
|
||
//
|
||
// 1) `launch_sp15_position_history_derivation` — single launch over
|
||
// all windows (grid = [n_windows, 1, 1]); each block walks one
|
||
// window's `actions_history_buf` sequentially, writing per-bar
|
||
// `position_history` / `side_ind` / `rt_ind` f32 streams.
|
||
//
|
||
// 2) Per-window `launch_sp15_cost_net_sharpe` — pnl from the
|
||
// per-window slice of `step_returns_buf`; half_spread / ofi /
|
||
// position / side_ind / rt_ind from the SoA streams produced
|
||
// above, sliced per window. Output `[mean, std, raw_sharpe]`
|
||
// written to `cost_net_sharpe_buf[w * 3..(w+1) * 3]`.
|
||
//
|
||
// 3) Per-window × 4 baseline launches (`buyhold` / `hold_only` /
|
||
// `naive_momentum` / `naive_reversion`). Each takes per-window
|
||
// slices of `close_prices_buf`, `half_spread_buf`,
|
||
// `ofi_scalar_buf` + the host-precomputed `commission_per_rt`,
|
||
// writing `[mean, std, raw_sharpe]` into its own per-window
|
||
// output buffer slot.
|
||
//
|
||
// All launches sit on `self.stream` and inherit the existing
|
||
// `eval_done_event` / DtoD-async dependency — no extra
|
||
// synchronisation needed. Annualisation is applied host-side in
|
||
// `consume_metrics_after_event`.
|
||
{
|
||
let n_win_i32_local = self.n_windows as i32;
|
||
let max_len_i32_local = self.max_len as i32;
|
||
|
||
// 1) Position-history derivation — one grid launch over all windows.
|
||
let (actions_hist_ptr, actions_hist_guard) =
|
||
self.actions_history_buf.device_ptr(&self.stream);
|
||
let _no_drop_actions = ManuallyDrop::new(actions_hist_guard);
|
||
super::gpu_dqn_trainer::launch_sp15_position_history_derivation(
|
||
&self.stream,
|
||
actions_hist_ptr,
|
||
n_win_i32_local,
|
||
max_len_i32_local,
|
||
self.b1_size,
|
||
self.b2_size,
|
||
self.b3_size,
|
||
self.position_history_buf.dev_ptr,
|
||
self.side_ind_buf.dev_ptr,
|
||
self.rt_ind_buf.dev_ptr,
|
||
)?;
|
||
|
||
// 2) + 3) Per-window cost-net + 4 baseline launches.
|
||
let (step_returns_base2, step_returns_guard2) =
|
||
self.step_returns_buf.device_ptr(&self.stream);
|
||
let _no_drop_returns = ManuallyDrop::new(step_returns_guard2);
|
||
let elem_bytes = std::mem::size_of::<f32>() as u64;
|
||
let stream_stride_bytes = (self.max_len as u64) * elem_bytes;
|
||
let isv_ptr_local = self.isv_signals_ptr;
|
||
let commission = self.commission_per_rt;
|
||
let cost_net_base = self.cost_net_sharpe_buf.dev_ptr;
|
||
let baseline_buyhold_base = self.baseline_buyhold_sharpe_buf.dev_ptr;
|
||
let baseline_hold_only_base = self.baseline_hold_only_sharpe_buf.dev_ptr;
|
||
let baseline_momentum_base = self.baseline_momentum_sharpe_buf.dev_ptr;
|
||
let baseline_reversion_base = self.baseline_reversion_sharpe_buf.dev_ptr;
|
||
let close_prices_base = self.close_prices_buf.dev_ptr;
|
||
let half_spread_base = self.half_spread_buf.dev_ptr;
|
||
let ofi_scalar_base = self.ofi_scalar_buf.dev_ptr;
|
||
let position_history_base = self.position_history_buf.dev_ptr;
|
||
let side_ind_base = self.side_ind_buf.dev_ptr;
|
||
let rt_ind_base = self.rt_ind_buf.dev_ptr;
|
||
let window_lens_host_v: Vec<i32> = self.window_lens_buf.read_all();
|
||
|
||
for w in 0..self.n_windows {
|
||
let n_w = window_lens_host_v[w];
|
||
if n_w <= 0 {
|
||
continue;
|
||
}
|
||
let off_stream = (w as u64) * stream_stride_bytes;
|
||
let off_three = (w as u64) * 3u64 * elem_bytes;
|
||
|
||
let pnl_dev = step_returns_base2 + off_stream;
|
||
let half_spread_dev = half_spread_base + off_stream;
|
||
let ofi_dev = ofi_scalar_base + off_stream;
|
||
let close_prices_dev = close_prices_base + off_stream;
|
||
let position_dev = position_history_base + off_stream;
|
||
let side_ind_dev = side_ind_base + off_stream;
|
||
let rt_ind_dev = rt_ind_base + off_stream;
|
||
|
||
// Cost-net sharpe (1.2.b).
|
||
super::gpu_dqn_trainer::launch_sp15_cost_net_sharpe(
|
||
&self.stream,
|
||
pnl_dev,
|
||
half_spread_dev,
|
||
ofi_dev,
|
||
side_ind_dev,
|
||
rt_ind_dev,
|
||
position_dev,
|
||
n_w,
|
||
commission,
|
||
isv_ptr_local,
|
||
cost_net_base + off_three,
|
||
&self.sp15_cost_net_sharpe_kernel,
|
||
)?;
|
||
|
||
// Buyhold baseline (1.4.b).
|
||
super::gpu_dqn_trainer::launch_sp15_baseline_buyhold(
|
||
&self.stream,
|
||
close_prices_dev,
|
||
half_spread_dev,
|
||
ofi_dev,
|
||
n_w,
|
||
commission,
|
||
baseline_buyhold_base + off_three,
|
||
&self.sp15_baseline_buyhold_kernel,
|
||
)?;
|
||
// Hold-only baseline.
|
||
super::gpu_dqn_trainer::launch_sp15_baseline_hold_only(
|
||
&self.stream,
|
||
close_prices_dev,
|
||
half_spread_dev,
|
||
ofi_dev,
|
||
n_w,
|
||
commission,
|
||
baseline_hold_only_base + off_three,
|
||
&self.sp15_baseline_hold_only_kernel,
|
||
)?;
|
||
// Naive-momentum baseline.
|
||
super::gpu_dqn_trainer::launch_sp15_baseline_naive_momentum(
|
||
&self.stream,
|
||
close_prices_dev,
|
||
half_spread_dev,
|
||
ofi_dev,
|
||
n_w,
|
||
commission,
|
||
baseline_momentum_base + off_three,
|
||
&self.sp15_baseline_naive_momentum_kernel,
|
||
)?;
|
||
// Naive-reversion baseline.
|
||
super::gpu_dqn_trainer::launch_sp15_baseline_naive_reversion(
|
||
&self.stream,
|
||
close_prices_dev,
|
||
half_spread_dev,
|
||
ofi_dev,
|
||
n_w,
|
||
commission,
|
||
baseline_reversion_base + off_three,
|
||
&self.sp15_baseline_naive_reversion_kernel,
|
||
)?;
|
||
}
|
||
}
|
||
|
||
// Queue DtoD-async copies from the device action-history buffers into
|
||
// their mapped-pinned mirrors. The kernel above (and the env-step kernel
|
||
// chain that produced these buffers) ordered before this point on the
|
||
// same stream, so the copies see the post-rollout values. The host reads
|
||
// through `host_ptr` after `eval_done_event.synchronize()` — a single
|
||
// barrier covers the metrics readback AND all four distribution readers.
|
||
//
|
||
// Mirrors are unconditional for `actions_history_buf` (always allocated)
|
||
// and conditional for `intent_mag_buf` / `picked_action_history_buf`
|
||
// (allocated only on DQN paths via `ensure_action_select_ready`). Skip
|
||
// the conditional copies when the source buffers are not yet wired —
|
||
// those readers return `Ok([0.0; *])` from the unwritten zero-init.
|
||
unsafe {
|
||
let n_bytes_full = self.n_windows * self.max_len * std::mem::size_of::<i32>();
|
||
// actions_history_buf → actions_history_pinned
|
||
{
|
||
let (src_ptr, _src_guard) =
|
||
self.actions_history_buf.device_ptr(&self.stream);
|
||
let dst_ptr = self.actions_history_pinned.dev_ptr;
|
||
cudarc::driver::result::memcpy_dtod_async(
|
||
dst_ptr, src_ptr, n_bytes_full, self.stream.cu_stream(),
|
||
).map_err(|e| MLError::ModelError(format!(
|
||
"actions_history dtod-to-pinned: {e}"
|
||
)))?;
|
||
}
|
||
// intent_mag_buf → intent_mag_pinned (DQN path only)
|
||
if let (Some(src_buf), Some(dst_buf)) =
|
||
(self.intent_mag_buf.as_ref(), self.intent_mag_pinned.as_ref())
|
||
{
|
||
let (src_ptr, _src_guard) = src_buf.device_ptr(&self.stream);
|
||
cudarc::driver::result::memcpy_dtod_async(
|
||
dst_buf.dev_ptr, src_ptr, n_bytes_full, self.stream.cu_stream(),
|
||
).map_err(|e| MLError::ModelError(format!(
|
||
"intent_mag dtod-to-pinned: {e}"
|
||
)))?;
|
||
}
|
||
// picked_action_history_buf → picked_action_history_pinned (DQN path).
|
||
// The picked buffer has the same shape `[n_windows * max_len]` (see
|
||
// its allocation in `ensure_action_select_ready`).
|
||
if let (Some(src_buf), Some(dst_buf)) = (
|
||
self.picked_action_history_buf.as_ref(),
|
||
self.picked_action_history_pinned.as_ref(),
|
||
) {
|
||
let (src_ptr, _src_guard) = src_buf.device_ptr(&self.stream);
|
||
cudarc::driver::result::memcpy_dtod_async(
|
||
dst_buf.dev_ptr, src_ptr, n_bytes_full, self.stream.cu_stream(),
|
||
).map_err(|e| MLError::ModelError(format!(
|
||
"picked_action_history dtod-to-pinned: {e}"
|
||
)))?;
|
||
}
|
||
}
|
||
|
||
// Lazy-create the event on first launch (so the CUDA context is
|
||
// guaranteed active). Re-record into the same event handle on every
|
||
// subsequent launch — `record` overwrites the prior recording per the
|
||
// CUDA driver semantics, no allocator churn.
|
||
if self.eval_done_event.is_none() {
|
||
let ev = self.stream.context().new_event(None).map_err(|e| {
|
||
MLError::ModelError(format!("eval_done_event create: {e}"))
|
||
})?;
|
||
self.eval_done_event = Some(ev);
|
||
}
|
||
let ev = self.eval_done_event.as_ref().expect("just created above");
|
||
ev.record(&self.stream).map_err(|e| {
|
||
MLError::ModelError(format!("eval_done_event record: {e}"))
|
||
})?;
|
||
|
||
Ok(())
|
||
}
|
||
|
||
/// Wait on `eval_done_event` (host-side `cuEventSynchronize`) and parse the
|
||
/// per-window metrics from the mapped-pinned `metrics_buf`. The event's GPU
|
||
/// dependency chain ensures the metrics kernel + all action-history mirror
|
||
/// copies queued by `launch_metrics_and_record_event` are globally visible
|
||
/// before the host load.
|
||
///
|
||
/// This is the consume half of the perf-fix split (2026-04-28). Pipelined
|
||
/// callers (`compute_validation_loss_consume`) invoke this at the start of
|
||
/// the next epoch — the host CPU thread already submitted the previous
|
||
/// epoch's training to `cuda_stream`, so by the time control reaches here
|
||
/// the eval stream has typically drained and the event is already complete.
|
||
/// Synchronous callers (the `evaluate_*` wrappers used by hyperopt / smoke
|
||
/// tests) invoke launch + consume back-to-back — no wall-clock change for
|
||
/// those paths since they had nothing to overlap with anyway.
|
||
pub fn consume_metrics_after_event(&self) -> Result<Vec<WindowMetrics>, MLError> {
|
||
let ev = self.eval_done_event.as_ref().ok_or_else(|| {
|
||
MLError::ModelError(
|
||
"consume_metrics_after_event called before launch_metrics_and_record_event"
|
||
.to_owned(),
|
||
)
|
||
})?;
|
||
ev.synchronize().map_err(|e| {
|
||
MLError::ModelError(format!("eval_done_event synchronize: {e}"))
|
||
})?;
|
||
let metrics_host: Vec<f32> = self.metrics_buf.read_all();
|
||
// SP15 Phase 1.1.b — per-window mean/std/raw_sharpe from the
|
||
// dedicated `sharpe_per_bar_kernel`. Three f32 slots per window:
|
||
// [0] mean, [1] std, [2] raw_sharpe (= mean / max(std, 1e-12)).
|
||
// Annualisation is applied below host-side, preserving the value
|
||
// contract from the pre-split fused kernel.
|
||
let sharpe_host: Vec<f32> = self.sharpe_buf.read_all();
|
||
// SP15 Wave 3b — per-window outputs from the cost-net sharpe kernel
|
||
// and the four counterfactual baselines. Same `[mean, std, raw_sharpe]`
|
||
// layout as `sharpe_buf`. Annualised host-side below using the same
|
||
// factor as `sharpe` — every Sharpe-style output goes through one
|
||
// multiplication by `annualization_factor` per
|
||
// `feedback_no_partial_refactor.md` (single annualisation locus).
|
||
let cost_net_host: Vec<f32> = self.cost_net_sharpe_buf.read_all();
|
||
let baseline_buyhold_host: Vec<f32> = self.baseline_buyhold_sharpe_buf.read_all();
|
||
let baseline_hold_only_host: Vec<f32> = self.baseline_hold_only_sharpe_buf.read_all();
|
||
let baseline_momentum_host: Vec<f32> = self.baseline_momentum_sharpe_buf.read_all();
|
||
let baseline_reversion_host: Vec<f32> = self.baseline_reversion_sharpe_buf.read_all();
|
||
let annualization = self.annualization_factor;
|
||
|
||
// Parse flat metrics into per-window structs.
|
||
// Layout matches compute_backtest_metrics kernel output AFTER the
|
||
// SP15 Phase 1.1.b sharpe split + the 2026-05-09 per-regime WR
|
||
// expansion (19 floats per window — 13 base metrics + 6
|
||
// per-regime WR/trade-count slots):
|
||
// [0] total_pnl
|
||
// [1] max_drawdown
|
||
// [2] sortino
|
||
// [3] win_rate
|
||
// [4] total_trades
|
||
// [5] var_95
|
||
// [6] cvar_95
|
||
// [7] calmar
|
||
// [8] omega_ratio
|
||
// [9] buy_count
|
||
// [10] sell_count
|
||
// [11] hold_count
|
||
// [12] unique_actions
|
||
// [13] win_rate_trending
|
||
// [14] trade_count_trending
|
||
// [15] win_rate_ranging
|
||
// [16] trade_count_ranging
|
||
// [17] win_rate_volatile
|
||
// [18] trade_count_volatile
|
||
// Sharpe is sourced from `sharpe_host` (sharpe_per_bar_kernel) and
|
||
// annualised host-side: `sharpe = raw_sharpe * annualization_factor`,
|
||
// matching the in-kernel multiplication done in the pre-split fused
|
||
// kernel (was `(mean / std_val) * annualization_factor`).
|
||
let results: Vec<WindowMetrics> = (0..self.n_windows)
|
||
.map(|w| {
|
||
let base = w * 19;
|
||
let sharpe_base = w * 3;
|
||
let raw_sharpe = sharpe_host[sharpe_base + 2];
|
||
let sharpe_annualised = raw_sharpe * annualization;
|
||
// SP15 Wave 3b — same `[mean, std, raw_sharpe]` layout
|
||
// and same annualisation pattern as `sharpe_buf`. Index
|
||
// 2 = raw_sharpe; multiply by `annualization_factor`
|
||
// for consistency with `WindowMetrics.sharpe`.
|
||
let cost_net_raw = cost_net_host[sharpe_base + 2];
|
||
let baseline_buyhold_raw = baseline_buyhold_host[sharpe_base + 2];
|
||
let baseline_hold_only_raw = baseline_hold_only_host[sharpe_base + 2];
|
||
let baseline_momentum_raw = baseline_momentum_host[sharpe_base + 2];
|
||
let baseline_reversion_raw = baseline_reversion_host[sharpe_base + 2];
|
||
WindowMetrics {
|
||
sharpe: sharpe_annualised,
|
||
total_pnl: metrics_host[base],
|
||
max_drawdown: metrics_host[base + 1].min(1.0), // Cap at 100% (liquidation)
|
||
sortino: metrics_host[base + 2],
|
||
win_rate: metrics_host[base + 3],
|
||
total_trades: metrics_host[base + 4],
|
||
var_95: metrics_host[base + 5],
|
||
cvar_95: metrics_host[base + 6],
|
||
calmar: metrics_host[base + 7],
|
||
omega_ratio: metrics_host[base + 8],
|
||
buy_count: metrics_host[base + 9],
|
||
sell_count: metrics_host[base + 10],
|
||
hold_count: metrics_host[base + 11],
|
||
unique_actions: metrics_host[base + 12],
|
||
sharpe_cost_net: cost_net_raw * annualization,
|
||
baseline_buyhold_sharpe: baseline_buyhold_raw * annualization,
|
||
baseline_hold_only_sharpe: baseline_hold_only_raw * annualization,
|
||
baseline_momentum_sharpe: baseline_momentum_raw * annualization,
|
||
baseline_reversion_sharpe: baseline_reversion_raw * annualization,
|
||
win_rate_trending: metrics_host[base + 13],
|
||
trade_count_trending: metrics_host[base + 14],
|
||
win_rate_ranging: metrics_host[base + 15],
|
||
trade_count_ranging: metrics_host[base + 16],
|
||
win_rate_volatile: metrics_host[base + 17],
|
||
trade_count_volatile: metrics_host[base + 18],
|
||
}
|
||
})
|
||
.collect();
|
||
|
||
Ok(results)
|
||
}
|
||
}
|
||
|
||
// ── Free helpers ──────────────────────────────────────────────────────────────
|
||
|
||
/// Extract raw F32 device pointer from a `CudaSlice<f32>`.
|
||
fn raw_f32_ptr(slice: &CudaSlice<f32>, stream: &Arc<CudaStream>) -> u64 {
|
||
let (ptr, guard) = slice.device_ptr(stream);
|
||
let _no_drop = ManuallyDrop::new(guard);
|
||
ptr
|
||
}
|
||
|
||
// ── Tests ─────────────────────────────────────────────────────────────────────
|
||
|
||
#[cfg(test)]
|
||
mod tests {
|
||
use super::*;
|
||
|
||
#[test]
|
||
fn test_init_portfolio_state() {
|
||
let state = GpuBacktestEvaluator::init_portfolio_state(3, 100_000.0);
|
||
assert_eq!(state.len(), 24); // 3 windows × 8 floats
|
||
// Window 0
|
||
assert_eq!(state[0], 100_000.0, "window 0 value");
|
||
assert_eq!(state[1], 0.0, "window 0 position");
|
||
assert_eq!(state[2], 100_000.0, "window 0 cash");
|
||
assert_eq!(state[3], 0.0, "window 0 unrealised_pnl");
|
||
assert_eq!(state[4], 100_000.0, "window 0 max_equity");
|
||
// Window 1 starts at index 8
|
||
assert_eq!(state[8], 100_000.0, "window 1 value");
|
||
assert_eq!(state[10], 100_000.0, "window 1 cash");
|
||
// Window 2 starts at index 16
|
||
assert_eq!(state[16], 100_000.0, "window 2 value");
|
||
}
|
||
|
||
#[test]
|
||
fn test_window_metrics_fields() {
|
||
let m = WindowMetrics {
|
||
sharpe: 1.5,
|
||
total_pnl: 0.05,
|
||
max_drawdown: 0.02,
|
||
sortino: 2.0,
|
||
win_rate: 0.55,
|
||
total_trades: 42.0,
|
||
var_95: 0.03,
|
||
cvar_95: 0.04,
|
||
calmar: 0.75,
|
||
omega_ratio: 1.2,
|
||
buy_count: 150.0,
|
||
sell_count: 120.0,
|
||
hold_count: 230.0,
|
||
unique_actions: 4.0,
|
||
sharpe_cost_net: 1.2,
|
||
baseline_buyhold_sharpe: 0.4,
|
||
baseline_hold_only_sharpe: 0.0,
|
||
baseline_momentum_sharpe: 0.3,
|
||
baseline_reversion_sharpe: -0.1,
|
||
win_rate_trending: 0.6,
|
||
trade_count_trending: 14.0,
|
||
win_rate_ranging: 0.45,
|
||
trade_count_ranging: 18.0,
|
||
win_rate_volatile: 0.55,
|
||
trade_count_volatile: 10.0,
|
||
};
|
||
assert!(m.sharpe > 0.0);
|
||
assert!(m.win_rate > 0.5);
|
||
assert!(m.total_trades > 0.0);
|
||
assert!(m.calmar > 0.0);
|
||
assert!(m.omega_ratio > 1.0);
|
||
// Per-regime aggregates should sum to total trades.
|
||
let regime_total = m.trade_count_trending + m.trade_count_ranging + m.trade_count_volatile;
|
||
assert!((regime_total - m.total_trades).abs() < 1e-3,
|
||
"regime totals must sum to total_trades: {} vs {}", regime_total, m.total_trades);
|
||
}
|
||
|
||
#[test]
|
||
fn test_gpu_backtest_config_default() {
|
||
let c = GpuBacktestConfig::default();
|
||
assert_eq!(c.max_position, 1.0);
|
||
assert_eq!(c.tx_cost_bps, 0.1);
|
||
assert!((c.spread_cost - 0.0001).abs() < 1e-8);
|
||
assert_eq!(c.initial_capital, 100_000.0);
|
||
// 1-minute bar frequency: 390 bars/day (6.5h × 60min)
|
||
assert_eq!(c.bars_per_day, 390.0);
|
||
}
|
||
|
||
#[test]
|
||
fn test_annualization_factor_1min_bars() {
|
||
let c = GpuBacktestConfig::default();
|
||
// For 1-minute bars: sqrt(390 * 252) = sqrt(98280) ≈ 313.5
|
||
let expected = (390.0_f32 * 252.0).sqrt();
|
||
let actual = (c.bars_per_day * 252.0).sqrt();
|
||
assert!((actual - expected).abs() < 0.01, "annualization: {actual} vs {expected}");
|
||
assert!(actual > 300.0 && actual < 320.0, "annualization should be ~313: {actual}");
|
||
}
|
||
|
||
/// Annualization compensates for val subsample stride — `effective_bars_per_day
|
||
/// = bars_per_day / stride`, so `annualization_factor =
|
||
/// sqrt(effective_bars_per_day * trading_days_per_year)`.
|
||
/// Per audit `docs/lookahead-bias-audit-2026-04-28.md` §5: prior code used
|
||
/// `bars_per_day` unscaled, inflating reported Sharpe by `sqrt(stride)`
|
||
/// (factor 2 for stride=4). Matches the constructor expression at
|
||
/// `gpu_backtest_evaluator.rs::new()`.
|
||
#[test]
|
||
fn test_annualization_factor_compensates_for_subsample_stride() {
|
||
// Replicates the constructor expression — single source of truth for
|
||
// the stride-aware annualization formula.
|
||
fn factor(c: &GpuBacktestConfig) -> f32 {
|
||
let stride = c.val_subsample_stride.max(1) as f32;
|
||
let effective_bars_per_day = c.bars_per_day / stride;
|
||
(effective_bars_per_day * c.trading_days_per_year).sqrt()
|
||
}
|
||
|
||
// stride=1 (default) — identical to legacy behavior.
|
||
let c1 = GpuBacktestConfig {
|
||
val_subsample_stride: 1,
|
||
..Default::default()
|
||
};
|
||
let f1 = factor(&c1);
|
||
let expected1 = (390.0_f32 * 252.0).sqrt();
|
||
assert!((f1 - expected1).abs() < 0.01, "stride=1: {f1} vs {expected1}");
|
||
|
||
// stride=4 — Sharpe annualization MUST shrink by factor 2.0.
|
||
let c4 = GpuBacktestConfig {
|
||
val_subsample_stride: 4,
|
||
..Default::default()
|
||
};
|
||
let f4 = factor(&c4);
|
||
let expected4 = ((390.0_f32 / 4.0) * 252.0).sqrt();
|
||
assert!((f4 - expected4).abs() < 0.01, "stride=4: {f4} vs {expected4}");
|
||
assert!(
|
||
(f1 / f4 - 2.0).abs() < 1e-3,
|
||
"stride=4 must halve annualization: f1={f1}, f4={f4}, ratio={}",
|
||
f1 / f4
|
||
);
|
||
|
||
// stride=0 is sanitized to 1 (defensive — no divide-by-zero).
|
||
let c0 = GpuBacktestConfig {
|
||
val_subsample_stride: 0,
|
||
..Default::default()
|
||
};
|
||
let f0 = factor(&c0);
|
||
assert!((f0 - expected1).abs() < 0.01, "stride=0 should sanitize to 1: {f0} vs {expected1}");
|
||
}
|
||
|
||
#[test]
|
||
fn test_new_rejects_empty_windows() {
|
||
let context = cudarc::driver::CudaContext::new(0);
|
||
if context.is_err() { return; } // skip on non-CUDA
|
||
let stream = context.ok().map(|c| c.default_stream());
|
||
if let Some(ref s) = stream {
|
||
let result = GpuBacktestEvaluator::new(
|
||
&[],
|
||
&[],
|
||
&[],
|
||
42,
|
||
GpuBacktestConfig::default(),
|
||
s,
|
||
);
|
||
assert!(result.is_err());
|
||
let msg = format!("{:?}", result.err());
|
||
assert!(msg.contains("No windows"), "expected 'No windows', got: {msg}");
|
||
}
|
||
}
|
||
|
||
/// Verify PTX sources compile without errors (requires CUDA — nvcc path).
|
||
#[test]
|
||
|
||
fn test_env_cubin_loads() {
|
||
let context = cudarc::driver::CudaContext::new(0).expect("CUDA context required");
|
||
let module = context.load_cubin(ENV_CUBIN.to_vec()).expect("load env cubin");
|
||
module.load_function("backtest_env_step").expect("load backtest_env_step");
|
||
}
|
||
|
||
#[test]
|
||
|
||
fn test_metrics_cubin_loads() {
|
||
let context = cudarc::driver::CudaContext::new(0).expect("CUDA context required");
|
||
let module = context.load_cubin(METRICS_CUBIN.to_vec()).expect("load metrics cubin");
|
||
module.load_function("compute_backtest_metrics").expect("load compute_backtest_metrics");
|
||
}
|
||
|
||
#[test]
|
||
|
||
fn test_gather_cubin_loads() {
|
||
let context = cudarc::driver::CudaContext::new(0).expect("CUDA context required");
|
||
let module = context.load_cubin(BACKTEST_DQN_CUBIN.to_vec()).expect("load backtest DQN cubin");
|
||
module.load_function("backtest_state_gather").expect("load backtest_state_gather");
|
||
}
|
||
|
||
/// Verify that gather_states() returns a ConfigError when portfolio_dim != 24.
|
||
#[test]
|
||
fn test_gather_states_portfolio_dim_mismatch() {
|
||
// We don't have a CUDA device in unit tests, so test the validation path
|
||
// indirectly by checking the expected portfolio_dim: 8 portfolio base + 16 MTF = 24.
|
||
let expected = ml_core::state_layout::PORTFOLIO_BASE_DIM + ml_core::state_layout::MTF_DIM;
|
||
assert_eq!(expected, 24, "portfolio_dim should be 8 portfolio + 16 MTF = 24");
|
||
}
|
||
|
||
/// Verify that the cuBLAS DQN helper kernels (compute_expected_q,
|
||
/// experience_action_select) compile from experience_kernels.cu.
|
||
#[test]
|
||
|
||
fn test_backtest_dqn_cubin_loads() {
|
||
let context = cudarc::driver::CudaContext::new(0).expect("CUDA context required");
|
||
let module = context.load_cubin(BACKTEST_DQN_CUBIN.to_vec()).expect("load backtest DQN cubin");
|
||
module.load_function("compute_expected_q").expect("load compute_expected_q");
|
||
}
|
||
|
||
/// Verify PPO forward cubin loads.
|
||
#[test]
|
||
|
||
fn test_ppo_forward_cubin_loads() {
|
||
let context = cudarc::driver::CudaContext::new(0).expect("CUDA context required");
|
||
let module = context.load_cubin(PPO_FORWARD_CUBIN.to_vec()).expect("load PPO forward cubin");
|
||
module.load_function("backtest_forward_ppo_kernel").expect("load backtest_forward_ppo_kernel");
|
||
}
|
||
|
||
/// Verify supervised signal cubin loads.
|
||
#[test]
|
||
|
||
fn test_supervised_signal_cubin_loads() {
|
||
let context = cudarc::driver::CudaContext::new(0).expect("CUDA context required");
|
||
let module = context.load_cubin(SUPERVISED_SIGNAL_CUBIN.to_vec()).expect("load supervised signal cubin");
|
||
module.load_function("signal_to_action_kernel").expect("load signal_to_action_kernel");
|
||
}
|
||
|
||
#[test]
|
||
fn test_gpu_backtest_evaluator_state_dim_calculation() {
|
||
// Canonical layout from ml_core::state_layout — 128-dim feature
|
||
// vector (42 market + 32 OFI + 14 TLOB + 16 MTF + 8 portfolio +
|
||
// 6 plan/ISV + 10 padding = 128). Already padded to 128 for cuBLAS
|
||
// K-tile alignment. The historical 96-dim layout (pre-2026-04
|
||
// refactor) was deprecated when the OFI / TLOB / MTF blocks were
|
||
// expanded; the canonical constant in `ml_core::state_layout`
|
||
// is the source of truth — code wins over docs (cf.
|
||
// `feedback_trust_code_not_docs`). Migrated under SP15 Wave 4.1b
|
||
// alongside the s1_input_dim 102→103 propagation that followed
|
||
// the same `feedback_trust_code_not_docs` correction in the
|
||
// wave's audit doc.
|
||
assert_eq!(ml_core::state_layout::STATE_DIM, 128);
|
||
assert_eq!(ml_core::state_layout::STATE_DIM_PADDED, 128);
|
||
}
|
||
|
||
/// Verify that CUDA Graph infrastructure compiles and struct fields are present.
|
||
///
|
||
/// This test validates:
|
||
/// - `SendSyncGraph` wrapper type is constructible (type-level check)
|
||
/// - `dqn_graph` field exists on `GpuBacktestEvaluator`
|
||
/// - `invalidate_dqn_graph` method is callable
|
||
/// - The struct remains `Send + Sync` despite wrapping `CudaGraph`
|
||
#[test]
|
||
fn test_cuda_graph_capture_smoke() {
|
||
// Verify GpuBacktestConfig still works (not affected by graph additions)
|
||
let c = GpuBacktestConfig::default();
|
||
assert_eq!(c.initial_capital, 100_000.0);
|
||
|
||
// Verify that the CudaGraph import resolves (type-level check).
|
||
// CudaGraph is !Send + !Sync, but SendSyncGraph wraps it safely.
|
||
fn assert_send<T: Send>() {}
|
||
fn assert_sync<T: Sync>() {}
|
||
assert_send::<Option<SendSyncGraph>>();
|
||
assert_sync::<Option<SendSyncGraph>>();
|
||
|
||
// The evaluator cannot be constructed without a CUDA device, but we
|
||
// can verify the method signatures compile by checking that
|
||
// GpuBacktestEvaluator has the expected methods (via trait object coercion).
|
||
// The evaluate_dqn_graphed and invalidate_dqn_graph methods are tested
|
||
// at runtime on GPU hardware via integration tests.
|
||
}
|
||
|
||
#[test]
|
||
fn test_dqn_backtest_config_from_network_dims() {
|
||
let cfg = DqnBacktestConfig::from_network_dims((256, 256, 128, 128));
|
||
assert_eq!(cfg.shared_h1, 256);
|
||
assert_eq!(cfg.shared_h2, 256);
|
||
assert_eq!(cfg.value_h, 128);
|
||
assert_eq!(cfg.adv_h, 128);
|
||
assert_eq!(cfg.num_atoms, 51);
|
||
assert_eq!(cfg.branch_0_size, 4);
|
||
assert_eq!(cfg.branch_1_size, 3);
|
||
assert_eq!(cfg.branch_2_size, 3);
|
||
assert_eq!(cfg.branch_3_size, 3);
|
||
// v_range = (10 / (1-0.95) * 1.2).clamp(20, 300) = 240
|
||
let expected_v_range = (10.0_f32 / (1.0 - 0.95) * 1.2).clamp(20.0, 300.0);
|
||
assert!((cfg.v_min - (-expected_v_range)).abs() < f32::EPSILON);
|
||
assert!((cfg.v_max - expected_v_range).abs() < f32::EPSILON);
|
||
}
|
||
|
||
/// Verify the chunked step loop parameters produce the expected kernel
|
||
/// launch reduction relative to the per-step loop.
|
||
#[test]
|
||
fn test_chunked_step_loop_launch_reduction() {
|
||
let chunk = DQN_BACKTEST_CHUNK_SIZE; // 512
|
||
let max_len = 32_000_usize;
|
||
let n_windows = 5_usize;
|
||
|
||
// Number of chunks (last chunk may be smaller)
|
||
let n_chunks = (max_len + chunk - 1) / chunk;
|
||
assert_eq!(n_chunks, 63); // 32000 / 512 = 62.5 → ceil = 63
|
||
|
||
// Per-step loop: 16 kernel launches per step (10 SGEMM + 6 bias)
|
||
// + 2 kernels (expected_q + action_select) + 1 DtoD + 1 env_step
|
||
// = 20 operations per step
|
||
let per_step_ops = max_len * 20;
|
||
assert_eq!(per_step_ops, 640_000);
|
||
|
||
// Chunked loop:
|
||
// Per chunk: 16 (cuBLAS forward) + 1 (expected_q) + 1 (action_select) = 18
|
||
// Per step: 1 (gather) + 1 (DtoD gather copy) + 1 (DtoD action copy) + 1 (env_step) = 4
|
||
let chunked_per_chunk_ops = n_chunks * 18; // 63 * 18 = 1134
|
||
let chunked_per_step_ops = max_len * 4; // 32000 * 4 = 128000
|
||
let chunked_total = chunked_per_chunk_ops + chunked_per_step_ops;
|
||
assert_eq!(chunked_total, 1_134 + 128_000);
|
||
|
||
// Verify the reduction factor (larger chunks → higher reduction)
|
||
let reduction = per_step_ops as f64 / chunked_total as f64;
|
||
assert!(reduction > 4.0, "expected >4x reduction, got {reduction:.1}x");
|
||
|
||
// Chunked batch size should be n_windows * chunk_size
|
||
let chunked_batch = n_windows * chunk;
|
||
assert_eq!(chunked_batch, 2_560); // 5 * 512
|
||
}
|
||
}
|