gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

feat(sp15-wave3b): host-side wire-up — eliminates 5 orphan launchers via GpuBacktestEvaluator constructor signature change

Browse Source 
Second half of the Wave 3 val-cost-streams refactor (3a kernel-side
foundation landed at e968f4ded). Atomically migrates the
GpuBacktestEvaluator::new contract; 3 call sites + 6 test call sites
+ 5 WindowMetrics fields + 11 new buffers + launch sequence wiring
all in this commit.

Constructor signature change: GpuBacktestEvaluator::new gains
window_lob_bars: &[Vec<LobBar>] parameter alongside existing
window_prices + window_features. Three production call sites migrated
atomically:
  - trainers/dqn/trainer/metrics.rs:651 (val_evaluator construction)
  - trainers/dqn/trainer/metrics.rs:1166 (extra_eval Dev/Test)
  - hyperopt/adapters/dqn.rs:1493 (full LobBar with real OFI)
  - hyperopt/adapters/ppo.rs:1364 (zero-OFI LobBar — PPO lacks per-bar
    OFI features; cost-net OFI-impact term degrades to 0; commission +
    half-spread × position still apply)

11 new mapped-pinned buffers on GpuBacktestEvaluator:
  Input (3): close_prices_buf, half_spread_buf, ofi_scalar_buf
  Derivation (3): position_history_buf, side_ind_buf, rt_ind_buf
  Output (5): cost_net_sharpe_buf, baseline_{buyhold,hold_only,
    momentum,reversion}_sharpe_buf

5 new WindowMetrics fields (host-annualised via × annualization_factor):
  - sharpe_cost_net (1.2.b cost-net sharpe)
  - baseline_{buyhold,hold_only,momentum,reversion}_sharpe (1.4.b)

Per-window eval flow now: existing fused metrics kernel → 1.1.b sharpe →
position_history_derivation (one launch over all windows) →
cost_net_sharpe (per-window) → 4 × baseline_* (per-window).

commission_per_rt: host-computed constant per D3 resolution, formula
config.tx_cost_bps × 0.0001 × config.initial_capital, passed by-value
at each baseline + cost_net launcher invocation.

Wave 3a kernel signature follow-up: cost_net_sharpe_kernel side_ind /
rt_ind switched from unsigned int* → float* so the cost-net kernel
chains directly with the f32 streams emitted by
position_history_derivation_kernel (no u32→f32 adapter buffer; bit-pun
mismatch fixed). cost_net oracle test migrated MappedU32Buffer →
MappedF32Buffer accordingly.

Atomic per feedback_no_partial_refactor: constructor sig change + 3
production + 6 test call site migrations + 5 orphan launchers
eliminated + WindowMetrics field additions + cost_net kernel sig fix
+ audit doc all in this commit.

Closes 1.2.b + 1.4.b + position_history_derivation orphan launchers
per feedback_wire_everything_up.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-06 23:39:45 +02:00
parent e968f4ded9
commit 4320820ae2
9 changed files with 562 additions and 27 deletions
Download Patch File Download Diff File Expand all files Collapse all files

34
crates/ml/src/cuda_pipeline/cost_net_sharpe_kernel.cu
View File

@@ -26,14 +26,24 @@
// Single-block launch, BLOCK=256. No atomicAdd (per feedback_no_atomicadd):
// reduction uses shared-memory tree-reduce only. `__threadfence_system()`
// flushes mapped-pinned writes before the host reads via `read_volatile`.
//
// SP15 Wave 3b (2026-05-06): `side_ind` / `rt_ind` are `float*` (0.0/1.0)
// rather than `unsigned int*` to match the type emitted by the
// `sp15_position_history_derivation_kernel`. The Wave 3a kernel signature
// declared u32 inputs; Wave 3b's wire-up pipeline produces f32 indicators
// (one Welford-friendly type, no f32→u32 adapter kernel needed). Internal
// math reads them as floats directly — was `(float)side_ind[i]`, now
// `side_ind[i]`. Bit-equivalent for the production caller (derivation
// emits exact 0.0f/1.0f) and the cost-net oracle test (Wave 3b updates
// it to MappedF32Buffer in the same commit).
extern "C" __global__ void cost_net_sharpe_kernel(
const float* __restrict__ pnl,
const float* __restrict__ half_spread,
const float* __restrict__ ofi,
const unsigned int* __restrict__ side_ind,
const unsigned int* __restrict__ rt_ind,
const float* __restrict__ position,
const float* __restrict__ pnl,
const float* __restrict__ half_spread,
const float* __restrict__ ofi,
const float* __restrict__ side_ind,
const float* __restrict__ rt_ind,
const float* __restrict__ position,
int n,
float commission_per_rt,
float* __restrict__ isv,
@@ -56,10 +66,10 @@ extern "C" __global__ void cost_net_sharpe_kernel(
float local_cost = 0.0f;
for (int i = tid; i < n; i += BLOCK) {
local_pnl += pnl[i];
float c_comm = commission_per_rt * (float)rt_ind[i];
float c_comm = commission_per_rt * rt_ind[i];
float pos_abs = fabsf(position[i]);
float c_spread = half_spread[i] * pos_abs * (float)side_ind[i];
float c_ofi = ofi_lambda * pos_abs * fabsf(ofi[i]) * (float)side_ind[i];
float c_spread = half_spread[i] * pos_abs * side_ind[i];
float c_ofi = ofi_lambda * pos_abs * fabsf(ofi[i]) * side_ind[i];
local_cost += c_comm + c_spread + c_ofi;
}
reduce_a[tid] = local_pnl;
@@ -86,10 +96,10 @@ extern "C" __global__ void cost_net_sharpe_kernel(
// Pass 2: variance of (pnl_t - cost_t).
float local_sq = 0.0f;
for (int i = tid; i < n; i += BLOCK) {
float c_comm = commission_per_rt * (float)rt_ind[i];
float c_comm = commission_per_rt * rt_ind[i];
float pos_abs = fabsf(position[i]);
float c_spread = half_spread[i] * pos_abs * (float)side_ind[i];
float c_ofi = ofi_lambda * pos_abs * fabsf(ofi[i]) * (float)side_ind[i];
float c_spread = half_spread[i] * pos_abs * side_ind[i];
float c_ofi = ofi_lambda * pos_abs * fabsf(ofi[i]) * side_ind[i];
float net = pnl[i] - (c_comm + c_spread + c_ofi);
float d = net - mean_net;
local_sq += d * d;

380
crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs
View File

@@ -37,6 +37,7 @@ 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 ───────────────────────────────────────────────────────
@@ -173,6 +174,16 @@ pub struct WindowMetrics {
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,
}
/// Configuration for the GPU backtest evaluator.
@@ -393,6 +404,89 @@ pub struct GpuBacktestEvaluator {
/// 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`
@@ -554,10 +648,20 @@ impl GpuBacktestEvaluator {
///
/// `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>,
@@ -566,6 +670,12 @@ impl GpuBacktestEvaluator {
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 {
@@ -780,6 +890,92 @@ impl GpuBacktestEvaluator {
.alloc_zeros::<i32>(n_windows)
.map_err(|e| MLError::ModelError(format!("forward_actions alloc: {e}")))?;
// ── 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
@@ -815,6 +1011,19 @@ impl GpuBacktestEvaluator {
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,
@@ -2571,6 +2780,146 @@ impl GpuBacktestEvaluator {
}
}
// ── 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,
)?;
// 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,
)?;
// 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,
)?;
// 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,
)?;
// 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,
)?;
}
}
// 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
@@ -2672,6 +3021,17 @@ impl GpuBacktestEvaluator {
// 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.
@@ -2701,6 +3061,15 @@ impl GpuBacktestEvaluator {
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],
@@ -2716,6 +3085,11 @@ impl GpuBacktestEvaluator {
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,
}
})
.collect();
@@ -2773,6 +3147,11 @@ mod tests {
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,
};
assert!(m.sharpe > 0.0);
assert!(m.win_rate > 0.5);
@@ -2858,6 +3237,7 @@ mod tests {
let stream = context.ok().map(|c| c.default_stream());
if let Some(ref s) = stream {
let result = GpuBacktestEvaluator::new(
&[],
&[],
&[],
42,

13
crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
View File

@@ -775,11 +775,14 @@ pub static SP15_COST_NET_SHARPE_CUBIN: &[u8] =
/// `+ half_spread[t] × |position[t]| × side_ind[t]`
/// `+ ofi_lambda × |position[t]| × |ofi[t]| × side_ind[t]`
///
/// `pnl`, `half_spread`, `ofi`, `position`, `isv`, `out` are device f32
/// pointers. `side_ind` and `rt_ind` are device u32 pointers. `out` MUST
/// point to ≥3 f32 slots: `[mean_pnl_net, std, sharpe_net]`. `isv` MUST
/// be the ISV bus (≥443 f32 slots) — slot 407 is read for `ofi_lambda`,
/// slot 408 is written with `mean(cost_t)`.
/// All input streams are device f32 pointers (`pnl`, `half_spread`, `ofi`,
/// `side_ind`, `rt_ind`, `position`, `isv`, `out`). SP15 Wave 3b changed
/// `side_ind` / `rt_ind` from `unsigned int*` to `float*` (0.0/1.0
/// indicators) so the kernel chains directly with the f32 streams produced
/// by `sp15_position_history_derivation_kernel` — no u32 adapter buffer.
/// `out` MUST point to ≥3 f32 slots: `[mean_pnl_net, std, sharpe_net]`.
/// `isv` MUST be the ISV bus (≥443 f32 slots) — slot 407 is read for
/// `ofi_lambda`, slot 408 is written with `mean(cost_t)`.
#[allow(clippy::too_many_arguments)]
pub fn launch_sp15_cost_net_sharpe(
stream: &Arc<CudaStream>,

40
crates/ml/src/hyperopt/adapters/dqn.rs
View File

@@ -1454,9 +1454,49 @@ impl DQNTrainer {
let evaluator_stream = device.cuda_stream()
.map_err(|e| MLError::ConfigError(format!("GpuBacktestEvaluator needs CUDA stream: {e}")))?;
// SP15 Wave 3b — LobBar SoA built per-window from the same val
// close-price + ofi data already used to populate `window_prices`
// / `window_features`. DQN adapter has full OFI access (per
// D4 resolution); each LobBar carries the real per-bar L1 OFI
// (`ofi_row[0]`, un-normalised) so the cost-net OFI-impact term
// is fully populated. Spread is the hyperopt-configured
// `config.spread_cost` broadcast as a scalar (no per-bar L2
// spread on this path). Window/bar indexing mirrors the price
// /feature build loop above: window `w` covers the absolute
// val-data range `[w*stride .. min(w*stride+window_size, total_bars))`,
// and the OFI lookup is `ofi_offset + absolute_bar_idx`.
let scalar_spread_dqn_adapter = config.spread_cost;
let mut window_lob_bars: Vec<Vec<crate::cuda_pipeline::lob_bar::LobBar>> =
Vec::with_capacity(window_count);
for win_idx in 0..window_count {
let start = win_idx * stride;
let end = (start + window_size).min(total_bars);
let mut bars = Vec::with_capacity(end - start);
for i in start..end {
let close = *val_close_prices.get(i).ok_or_else(|| {
MLError::ConfigError(format!(
"lob_bar val_close_prices index {i} out of bounds (len={})",
total_bars
))
})? as f32;
let ofi_idx = ofi_offset + i;
let ofi_scalar = ofi_ref
.and_then(|o| o.get(ofi_idx))
.and_then(|row| row.first())
.copied()
.unwrap_or(0.0) as f32;
bars.push(crate::cuda_pipeline::lob_bar::LobBar::new(
close,
scalar_spread_dqn_adapter,
ofi_scalar,
));
}
window_lob_bars.push(bars);
}
let mut evaluator = GpuBacktestEvaluator::new(
&window_prices,
&window_features,
&window_lob_bars,
feature_dim,
config,
evaluator_stream,

20
crates/ml/src/hyperopt/adapters/ppo.rs
View File

@@ -1343,9 +1343,29 @@ impl PPOTrainer {
let stream = device.cuda_stream()
.map_err(|_| MLError::ConfigError("CUDA required for PPO backtest".to_owned()))?;
// SP15 Wave 3b — PPO hyperopt path lacks per-bar OFI features
// (per D4 resolution); construct LobBar SoA with `ofi=0.0` for
// ALL bars in PPO eval. Cost-net OFI-impact term degrades to 0
// here (commission + half-spread × position still apply). DQN
// adapter and val_evaluator get full LobBar with real OFI.
// Spread is broadcast from `config.spread_cost`. Single window
// matches `window_prices`/`window_features` shape above.
let scalar_spread_ppo = config.spread_cost;
let mut lob_bars_ppo: Vec<crate::cuda_pipeline::lob_bar::LobBar> =
Vec::with_capacity(val_data.len());
for (_fv, close_price) in val_data {
lob_bars_ppo.push(crate::cuda_pipeline::lob_bar::LobBar::new(
*close_price as f32,
scalar_spread_ppo,
0.0,
));
}
let window_lob_bars = vec![lob_bars_ppo];
let mut evaluator = GpuBacktestEvaluator::new(
&window_prices,
&window_features,
&window_lob_bars,
feature_dim,
config,
stream,

50
crates/ml/src/trainers/dqn/trainer/metrics.rs
View File

@@ -634,9 +634,38 @@ impl DQNTrainer {
val_subsample_stride: VAL_SUBSAMPLE_STRIDE,
};
// SP15 Wave 3b — LobBar SoA built from the same val_data slice
// already used to populate `prices` / `features`. Price = the
// raw close from `target[TARGET_RAW_CLOSE]`; spread is the
// hyperopt-configured `spread_cost` broadcast as a scalar (no
// per-bar L2 spread available in the val path); OFI is the
// un-normalised per-bar L1 OFI (`ofi_row[0]`) from the same
// mbp10 features the gather kernel already consumes — keeps
// the cost-net OFI-impact term meaningful on the val path.
// Mirrors the subsample-stride gating so each LobBar lines up
// with the corresponding (price, feature) pair.
let scalar_spread_val = config.spread_cost;
let mut lob_bars: Vec<crate::cuda_pipeline::lob_bar::LobBar> =
Vec::with_capacity(self.val_data.len() / VAL_SUBSAMPLE_STRIDE as usize + 1);
for (i, (_fv, target)) in self.val_data.iter().enumerate() {
if (i as u32) % VAL_SUBSAMPLE_STRIDE != 0 { continue; }
let close = target[TARGET_RAW_CLOSE] as f32;
let ofi_idx = self.ofi_val_offset + i;
let ofi_scalar = self.ofi_features.as_ref()
.and_then(|o| o.get(ofi_idx))
.and_then(|row| row.first())
.copied()
.unwrap_or(0.0) as f32;
lob_bars.push(crate::cuda_pipeline::lob_bar::LobBar::new(
close, scalar_spread_val, ofi_scalar,
));
}
let window_lob_bars = vec![lob_bars];
let evaluator = GpuBacktestEvaluator::new(
&window_prices,
&window_features,
&window_lob_bars,
feature_dim,
config,
stream,
@@ -1141,9 +1170,30 @@ impl DQNTrainer {
val_subsample_stride: 1, // no subsample for one-shot eval
};
// SP15 Wave 3b — LobBar SoA built from the same dev/test slice.
// Same broadcast-spread + L1-OFI source as the val path
// (one-shot eval, stride=1 so no subsample gate).
let scalar_spread_extra = config.spread_cost;
let mut lob_bars_extra: Vec<crate::cuda_pipeline::lob_bar::LobBar> =
Vec::with_capacity(features.len());
for (i, (_fv, target)) in features.iter().zip(targets.iter()).enumerate() {
let close = target[TARGET_RAW_CLOSE] as f32;
let ofi_idx = ofi_offset + i;
let ofi_scalar = self.ofi_features.as_ref()
.and_then(|o| o.get(ofi_idx))
.and_then(|row| row.first())
.copied()
.unwrap_or(0.0) as f32;
lob_bars_extra.push(crate::cuda_pipeline::lob_bar::LobBar::new(
close, scalar_spread_extra, ofi_scalar,
));
}
let window_lob_bars = vec![lob_bars_extra];
let evaluator = GpuBacktestEvaluator::new(
&window_prices,
&window_features,
&window_lob_bars,
feature_dim,
config,
&stream,

27
crates/ml/tests/gpu_backtest_validation.rs
View File

@@ -120,9 +120,23 @@ mod gpu_tests {
use ml::cuda_pipeline::gpu_backtest_evaluator::{
GpuBacktestConfig, GpuBacktestEvaluator,
};
use ml::cuda_pipeline::lob_bar::LobBar;
use ml::MLError;
use tracing::warn;
/// SP15 Wave 3b — build a deterministic `[Vec<LobBar>; 1]` shaped to
/// match the single-window OHLC `prices` slice passed to
/// `GpuBacktestEvaluator::new`. Uses the close price for `LobBar.price`
/// and zeroes `spread`/`ofi` (these tests don't exercise cost-net
/// sharpe — they assert raw PnL/drawdown semantics — so the LobBar
/// streams are inert by construction).
fn lob_bars_from_prices(prices: &[[f32; 4]]) -> Vec<LobBar> {
prices
.iter()
.map(|ohlc| LobBar::new(ohlc[3], 0.0, 0.0))
.collect()
}
/// Skip test gracefully if no CUDA device is available.
/// Returns the stream (owned Arc) on success.
fn try_cuda_stream() -> Option<Arc<CudaStream>> {
@@ -180,9 +194,11 @@ mod gpu_tests {
..Default::default()
};
let bars = lob_bars_from_prices(&prices);
let mut evaluator = GpuBacktestEvaluator::new(
&[prices],
&[features],
&[bars],
FEATURE_DIM,
config,
&stream,
@@ -240,9 +256,11 @@ mod gpu_tests {
..Default::default()
};
let bars = lob_bars_from_prices(&prices);
let mut evaluator = GpuBacktestEvaluator::new(
&[prices],
&[features],
&[bars],
FEATURE_DIM,
config,
&stream,
@@ -287,9 +305,11 @@ mod gpu_tests {
let config = GpuBacktestConfig::default();
let bars = lob_bars_from_prices(&prices);
let mut evaluator = GpuBacktestEvaluator::new(
&[prices],
&[features],
&[bars],
FEATURE_DIM,
config,
&stream,
@@ -339,9 +359,12 @@ mod gpu_tests {
..Default::default()
};
let bars_up = lob_bars_from_prices(&prices_up);
let bars_down = lob_bars_from_prices(&prices_down);
let mut evaluator = GpuBacktestEvaluator::new(
&[prices_up, prices_down],
&[features1, features2],
&[bars_up, bars_down],
FEATURE_DIM,
config,
&stream,
@@ -386,9 +409,11 @@ mod gpu_tests {
let features = generate_features(N_BARS, FEATURE_DIM);
let config = GpuBacktestConfig::default();
let bars = lob_bars_from_prices(&prices);
let mut evaluator = GpuBacktestEvaluator::new(
&[prices],
&[features],
&[bars],
FEATURE_DIM,
config,
&stream,
@@ -435,9 +460,11 @@ mod gpu_tests {
let features = generate_features(N_BARS, FEATURE_DIM);
let config = GpuBacktestConfig::default();
let bars = lob_bars_from_prices(&prices);
let mut evaluator = GpuBacktestEvaluator::new(
&[prices],
&[features],
&[bars],
FEATURE_DIM,
config,
&stream,

23
crates/ml/tests/sp15_phase1_oracle_tests.rs
View File

@@ -16,7 +16,7 @@ mod gpu {
launch_sp15_cost_net_sharpe, launch_sp15_dd_state,
launch_sp15_position_history_derivation, launch_sp15_sharpe_per_bar,
};
use ml::cuda_pipeline::mapped_pinned::{MappedF32Buffer, MappedI32Buffer, MappedU32Buffer};
use ml::cuda_pipeline::mapped_pinned::{MappedF32Buffer, MappedI32Buffer};
use ml::cuda_pipeline::sp15_isv_slots::{
ALPHA_SPLIT_INDEX, COST_PER_BAR_AVG_INDEX, DD_CURRENT_INDEX, DD_MAX_INDEX, DD_PCT_INDEX,
DD_RECOVERY_BARS_INDEX, OFI_IMPACT_LAMBDA_INDEX, SP15_SLOT_END,
@@ -141,11 +141,14 @@ mod gpu {
let half_spread = vec![0.25f32; n];
let ofi = vec![0.0f32; n];
let position = vec![1.0f32; n];
let mut side_ind = vec![0u32; n];
side_ind[0] = 1;
side_ind[5] = 1;
let mut rt_ind = vec![0u32; n];
rt_ind[5] = 1;
// SP15 Wave 3b: side_ind / rt_ind are now f32 (0.0/1.0) so the
// cost_net_sharpe kernel signature matches the f32 streams emitted
// by the position_history_derivation_kernel; no u32→f32 adapter.
let mut side_ind = vec![0.0f32; n];
side_ind[0] = 1.0;
side_ind[5] = 1.0;
let mut rt_ind = vec![0.0f32; n];
rt_ind[5] = 1.0;
// Safety: CUDA context active via `make_test_stream` above.
let pnl_buf = unsafe { MappedF32Buffer::new(n) }
@@ -160,11 +163,11 @@ mod gpu {
let position_buf = unsafe { MappedF32Buffer::new(n) }
.expect("alloc MappedF32Buffer for position");
position_buf.write_from_slice(&position);
let side_ind_buf = unsafe { MappedU32Buffer::new(n) }
.expect("alloc MappedU32Buffer for side_ind");
let side_ind_buf = unsafe { MappedF32Buffer::new(n) }
.expect("alloc MappedF32Buffer for side_ind");
side_ind_buf.write_from_slice(&side_ind);
let rt_ind_buf = unsafe { MappedU32Buffer::new(n) }
.expect("alloc MappedU32Buffer for rt_ind");
let rt_ind_buf = unsafe { MappedF32Buffer::new(n) }
.expect("alloc MappedF32Buffer for rt_ind");
rt_ind_buf.write_from_slice(&rt_ind);
// ISV bus seeded with OFI_IMPACT_LAMBDA=0 (mirrors the spec's

2
docs/dqn-wire-up-audit.md
View File

File diff suppressed because one or more lines are too long