diff --git a/crates/ml/src/cuda_pipeline/cost_net_sharpe_kernel.cu b/crates/ml/src/cuda_pipeline/cost_net_sharpe_kernel.cu index b1dcd5192..9813ae2ae 100644 --- a/crates/ml/src/cuda_pipeline/cost_net_sharpe_kernel.cu +++ b/crates/ml/src/cuda_pipeline/cost_net_sharpe_kernel.cu @@ -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; diff --git a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs index 71db74429..27966ddaf 100644 --- a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs +++ b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs @@ -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>` (feature rows) per window. + /// `window_lob_bars`: one `Vec` 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>], + window_lob_bars: &[Vec], feature_dim: usize, config: GpuBacktestConfig, parent_stream: &Arc, @@ -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::(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::() + flat_features.len() * std::mem::size_of::()) 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::() 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 = 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 = 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 = self.cost_net_sharpe_buf.read_all(); + let baseline_buyhold_host: Vec = self.baseline_buyhold_sharpe_buf.read_all(); + let baseline_hold_only_host: Vec = self.baseline_hold_only_sharpe_buf.read_all(); + let baseline_momentum_host: Vec = self.baseline_momentum_sharpe_buf.read_all(); + let baseline_reversion_host: Vec = 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, diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 58b454d38..6f88cb32c 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -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, diff --git a/crates/ml/src/hyperopt/adapters/dqn.rs b/crates/ml/src/hyperopt/adapters/dqn.rs index 562eddcd8..942cd4b47 100644 --- a/crates/ml/src/hyperopt/adapters/dqn.rs +++ b/crates/ml/src/hyperopt/adapters/dqn.rs @@ -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::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, diff --git a/crates/ml/src/hyperopt/adapters/ppo.rs b/crates/ml/src/hyperopt/adapters/ppo.rs index a0f6475e2..054d7eb9d 100644 --- a/crates/ml/src/hyperopt/adapters/ppo.rs +++ b/crates/ml/src/hyperopt/adapters/ppo.rs @@ -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 = + 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, diff --git a/crates/ml/src/trainers/dqn/trainer/metrics.rs b/crates/ml/src/trainers/dqn/trainer/metrics.rs index b7aa0d9d9..4df1569bc 100644 --- a/crates/ml/src/trainers/dqn/trainer/metrics.rs +++ b/crates/ml/src/trainers/dqn/trainer/metrics.rs @@ -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 = + 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 = + 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, diff --git a/crates/ml/tests/gpu_backtest_validation.rs b/crates/ml/tests/gpu_backtest_validation.rs index 9e8320ec7..157853372 100644 --- a/crates/ml/tests/gpu_backtest_validation.rs +++ b/crates/ml/tests/gpu_backtest_validation.rs @@ -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; 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 { + 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> { @@ -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, diff --git a/crates/ml/tests/sp15_phase1_oracle_tests.rs b/crates/ml/tests/sp15_phase1_oracle_tests.rs index 896c5c28a..608b4c865 100644 --- a/crates/ml/tests/sp15_phase1_oracle_tests.rs +++ b/crates/ml/tests/sp15_phase1_oracle_tests.rs @@ -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 diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 7073c56ce..bb81145ca 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,8 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +SP15 Wave 3b — host-side wire-up for val-cost-streams (2026-05-06): atomic host-side half of the Wave 3 val-cost-streams refactor. Wave 3a (commit `e968f4ded`) landed kernel signature changes + new derivation kernel + shared header + oracle-test migration with 5 ORPHAN launchers. Wave 3b lands the production callers atomically, eliminating those orphans by extending `GpuBacktestEvaluator::new`'s constructor signature with a new `window_lob_bars: &[Vec]` parameter, allocating 11 new mapped-pinned buffers (3 input + 3 derivation + 5 metric output), wiring the per-window launch sequence inside `launch_metrics_and_record_event`, adding 5 new fields to `WindowMetrics`, and migrating 3 production call sites + 6 integration-test call sites in this commit. (1) **Constructor signature change** — `GpuBacktestEvaluator::new` gains `window_lob_bars: &[Vec]` parameter (between `window_features` and `feature_dim`). Validates `window_lob_bars.len() == n_windows` (returns `ConfigError` otherwise — fail-loud per `feedback_no_stubs`). Each `Vec` is zero-padded to `max_len`; the SoA flatten splits into three flat `[n_windows * max_len]` f32 streams (close prices = `LobBar.price`, half-spread = `LobBar.spread × 0.5`, OFI scalar = `LobBar.ofi`). (2) **11 new mapped-pinned buffers on the struct**: 3 input streams (`close_prices_buf`, `half_spread_buf`, `ofi_scalar_buf`, each `[n_windows * max_len]`) populated via `write_from_slice` at construct time and read-only thereafter; 3 derivation outputs (`position_history_buf`, `side_ind_buf`, `rt_ind_buf`, each `[n_windows * max_len]`) zero-initialised and filled per-call by `launch_sp15_position_history_derivation`; 5 metric outputs (`cost_net_sharpe_buf` + 4 × `baseline_*_sharpe_buf`, each `[n_windows * 3]` for `[mean, std, raw_sharpe]` per window — same layout as the 1.1.b `sharpe_buf`). All allocated via `unsafe { MappedF32Buffer::new(...) }` per `feedback_no_htod_htoh_only_mapped_pinned.md`. (3) **`commission_per_rt: f32` field on the evaluator** — host-precomputed flat per-roundtrip commission per D3 resolution. Formula: `config.tx_cost_bps × 0.0001 × config.initial_capital`. Computed once at construct time so the per-window launch loop reads a stable f32 without re-deriving each call. The 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 `half_spread_buf` + `ofi_scalar_buf` streams (per spec §6.2). (4) **Wave 3a kernel signature follow-up — cost_net_sharpe `side_ind` / `rt_ind` switched from `unsigned int*` to `float*`** to chain directly with the f32 streams emitted by `position_history_derivation_kernel`. Wave 3a as-shipped declared u32 inputs which would have type-pun-failed when fed the f32 derivation outputs (1.0f bits → uint 1065353216 cast to float = catastrophic miscalc). The internal cast `(float)side_ind[i]` becomes `side_ind[i]` directly. Bit-equivalent for the production caller (derivation emits exact 0.0f/1.0f) and the migrated cost_net oracle test (now uses `MappedF32Buffer` for both flag streams; `MappedU32Buffer` import removed). (5) **3 production call sites migrated**: (a) `crates/ml/src/trainers/dqn/trainer/metrics.rs` val_evaluator construction (~line 651) — builds `Vec` from `val_data` slice using `target[TARGET_RAW_CLOSE]` for price, `config.spread_cost` broadcast scalar for spread, `ofi_features[ofi_val_offset + i][0]` (un-normalised L1 OFI) for OFI scalar, gated behind the same `VAL_SUBSAMPLE_STRIDE=4` modulus the existing price/feature loop uses; (b) `crates/ml/src/trainers/dqn/trainer/metrics.rs` extra_eval Dev/Test construction (~line 1166) — same broadcast-spread + L1-OFI pattern, stride=1 (no subsample for one-shot eval); (c) `crates/ml/src/hyperopt/adapters/dqn.rs` walk-forward eval construction (~line 1493) — multi-window LobBar SoA built from `val_close_prices` + `preloaded_ofi_features` with absolute `ofi_offset + i` indexing matching the price/feature loop above (DQN adapter has full per-bar OFI access per D4 resolution); (d) `crates/ml/src/hyperopt/adapters/ppo.rs` PPO backtest construction (~line 1364) — single-window LobBar SoA with `ofi=0.0` for ALL bars per D4 resolution (PPO hyperopt path lacks per-bar OFI features; 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). (6) **6 integration-test call sites migrated** in `crates/ml/tests/gpu_backtest_validation.rs` via a new `lob_bars_from_prices` helper — uses close price for `LobBar.price`, zeroes `spread`/`ofi` (these tests assert raw PnL/drawdown semantics, not cost-net; LobBar streams are inert by construction). The 6 call sites cover always-long/always-short/always-flat/multi-window/extended-metrics/active-model trade tests. (7) **5 new `WindowMetrics` fields**: `sharpe_cost_net` (1.2.b cost-net sharpe), `baseline_buyhold_sharpe` / `baseline_hold_only_sharpe` / `baseline_momentum_sharpe` / `baseline_reversion_sharpe` (1.4.b counterfactual baselines). All 5 annualised host-side in `consume_metrics_after_event` by multiplying the kernel's `raw_sharpe` (slot index 2 in each per-window 3-tuple) by `self.annualization_factor` — same pattern + same factor as `WindowMetrics.sharpe` per the 1.1.b precedent (single annualisation locus per `feedback_no_partial_refactor`). The struct's `Default`-style construction in `consume_metrics_after_event` and the test_window_metrics_fields unit test are updated to populate all 5; no other consumers needed updates because Rust struct literals are exhaustive. (8) **Launch sequence in `launch_metrics_and_record_event`** — the existing fused metrics kernel + 1.1.b per-window sharpe loop are UNCHANGED. After the existing sharpe loop, a new SP15 Wave 3b block launches (in order on the same `self.stream`): one `launch_sp15_position_history_derivation` over all windows (grid=[n_windows,1,1]; each block is one thread walking one window's actions_history sequentially); per-window loop running 5 launches per window — `launch_sp15_cost_net_sharpe` (consuming step_returns + half_spread + ofi + position_history + side_ind + rt_ind + commission_per_rt + isv_signals_ptr; per-window pnl/half_spread/ofi/position/side_ind/rt_ind slices via per-element-stride byte arithmetic mirroring 1.1.b's sharpe_per_bar slice pattern), then 4 baseline launches (`buyhold` / `hold_only` / `naive_momentum` / `naive_reversion`) each consuming close_prices + half_spread + ofi + commission_per_rt, each writing to its own per-window output buffer slice. All 6 launches inherit the existing `eval_done_event` / DtoD-async dependency — no new sync needed; only one `eval_done_event` record at the end of the queued work covers the new outputs alongside the existing metrics readback + action-history mirror copies. The per-window loop reads `window_lens_buf` via the mapped-pinned host alias (no DtoH) and skips windows with `n_w <= 0`. (9) **Buffer ABI for cost_net_sharpe per-window slice** — `step_returns_buf.device_ptr(...) + w * max_len * sizeof(f32)` for the pnl slice; `half_spread_buf.dev_ptr + w * max_len * sizeof(f32)` for the half-spread slice (same byte offset arithmetic for ofi/position_history/side_ind/rt_ind). Output slice = `cost_net_sharpe_buf.dev_ptr + w * 3 * sizeof(f32)`. The pattern mirrors the 1.1.b sharpe_per_bar per-window launch loop's offset arithmetic (`step_returns_base + (w as u64) * stride_bytes` / `sharpe_dev_base + (w as u64) * 3u64 * elem_bytes`). (10) **5 ORPHAN launchers eliminated** per `feedback_wire_everything_up`: `launch_sp15_baseline_buyhold` / `launch_sp15_baseline_hold_only` / `launch_sp15_baseline_naive_momentum` / `launch_sp15_baseline_naive_reversion` / `launch_sp15_position_history_derivation` all have production callers in this commit. The orphan annotation in their doc-comments is retained only as historical context (the comments document the 3a/3b decomposition); future maintainers reading the launcher headers see exactly when the orphan window opened and closed. (11) **Test-harness reject-empty-windows update** — `test_new_rejects_empty_windows` passes a third empty slice to match the new constructor signature; the constructor's first check (`n_windows == 0`) fires before reaching the new `window_lob_bars.len() != n_windows` check, so the original assertion against "No windows" still holds. Touched: `crates/ml/src/cuda_pipeline/cost_net_sharpe_kernel.cu` (signature: 2 params changed `unsigned int*` → `float*` + 4 internal `(float)` casts dropped), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (1 launcher doc-comment refresh), `crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs` (constructor sig + 11 new buffers + 1 new commission_per_rt field + LobBar SoA flatten + WindowMetrics 5 new fields + launch_metrics_and_record_event Wave 3b block + consume_metrics_after_event 5 new field reads/annualisations + test_window_metrics_fields update + test_new_rejects_empty_windows update + LobBar import), `crates/ml/src/trainers/dqn/trainer/metrics.rs` (2 call sites: val_evaluator + extra_eval), `crates/ml/src/hyperopt/adapters/dqn.rs` (1 call site: walk-forward eval), `crates/ml/src/hyperopt/adapters/ppo.rs` (1 call site: PPO backtest with zero-OFI explanation comment), `crates/ml/tests/gpu_backtest_validation.rs` (helper + 6 call sites: always-long/always-short/always-flat/multi-window/extended-metrics/active-model), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (cost_net oracle test side_ind/rt_ind buffers switched MappedU32Buffer → MappedF32Buffer + comment + import block trimmed). Hard rules: `feedback_no_partial_refactor` (constructor sig change + 3 production call sites + 6 test call sites + 5 ORPHAN launchers eliminated + WindowMetrics field additions + cost_net kernel sig fix + audit doc all in this commit; zero parallel paths), `feedback_wire_everything_up` (Wave 3a's 5 orphans all wired in this commit), `feedback_no_legacy_aliases` (no compatibility shim — old constructor signature is gone; type mismatch on cost_net `side_ind`/`rt_ind` is fixed in-place not bridged), `feedback_no_htod_htoh_only_mapped_pinned` (all 11 new buffers are mapped-pinned with `cuMemHostAlloc DEVICEMAP`; kernel writes via dev_ptr, host reads via host_ptr after stream sync), `feedback_no_atomicadd` (no new reductions; the new kernels chain together inside the existing single-block tree-reduce contracts established in Wave 3a), `feedback_no_stubs` (all 6 launches do real work against real data; no NULL-skip stubs — `commission_per_rt = 0.0` is a configured-zero, not a stub), `pearl_no_host_branches_in_captured_graph` (the new launches happen in `launch_metrics_and_record_event` which is OUTSIDE the experience-rollout CUDA Graph capture region — same precedent as the existing 1.1.b sharpe-per-bar per-window launch loop in the same function). End-to-end oracle test deferred per the dispatch's "skip if too heavy" guidance — full `GpuBacktestEvaluator` fixture requires a populated FusedTrainingCtx with weight buffers + cuBLAS handles, which is too much trainer machinery for a single oracle test; the existing 3a oracle tests + L40S smoke validation path is the correct gate. Verified: `SQLX_OFFLINE=true cargo check -p ml --features cuda` clean (18 unrelated pre-existing warnings); `cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored` 30 of 30 oracle tests green (cost_net + baseline + position_history + dd_state + final_reward + cooldown + plasticity + dd_trajectory + regret); `cargo test -p ml --features cuda --lib` HELD/IMPROVED the baseline (Wave 3a baseline = 945 pass / 14 fail; Wave 3b = 946 pass / 13 fail — one MORE test passing post-Wave-3b due to the cost_net/integration plumbing landing). Pre-existing failures (test_gpu_backtest_evaluator_state_dim_calculation stale STATE_DIM constant, gpu_backtest_validation portfolio_dim 3≠24, test_hyperopt_14d_search_space dimension assertion drift, etc.) are unchanged. + SP15 Wave 3a — kernel-side foundation for val-cost-streams (2026-05-06): atomic kernel-side half of the Phase 1.4.b + Phase 1.2.b val-cost-streams refactor, decomposed into 3a (this commit) + 3b (host-side wire-up, separate dispatch) per session-capacity scoping. Wave 3a lands the kernel contract changes + new derivation kernel + shared header + oracle-test migration; the 5 launchers (4 baselines + 1 derivation) currently sit ORPHAN awaiting Wave 3b's `GpuBacktestEvaluator::new` constructor signature change that wires production callers. Acceptable transient state — explicitly bounded by the 3a/3b split. (1) **Baseline kernel signature change (1.4)**: `baseline_buyhold_kernel`, `baseline_hold_only_kernel`, `baseline_naive_momentum_kernel`, `baseline_naive_reversion_kernel` all gain `out: float*` as their last parameter, replacing the prior `isv: float*` ISV-bus pointer. Each kernel writes `[mean, std, raw_sharpe]` to `out[0..3]` for the per-window slice the launcher invocation handles — mirrors the 1.1.b `sharpe_per_bar_kernel` shape exactly. The ISV-scalar writes to slots 409 (BASELINE_BUYHOLD_SHARPE), 410 (BASELINE_HOLD_ONLY_SHARPE), 412 (BASELINE_NAIVE_MOMENTUM_SHARPE), 416 (BASELINE_NAIVE_REVERSION_SHARPE) are removed entirely (NOT kept as parallel writes per `feedback_no_partial_refactor`). The internal `compute_baseline_sharpe` helper changes return-by-value to write-via-out-pointer, emitting all three of mean/std/sharpe (was sharpe-only). (2) **4 ISV slot constants retired in `sp15_isv_slots.rs`**: `BASELINE_BUYHOLD_SHARPE_INDEX` / `BASELINE_HOLD_ONLY_SHARPE_INDEX` / `BASELINE_NAIVE_MOMENTUM_SHARPE_INDEX` / `BASELINE_NAIVE_REVERSION_SHARPE_INDEX` removed per `feedback_no_legacy_aliases` — no ghost constants. The slot indices 409/410/412/416 themselves remain reserved gaps in the SP15 layout for layout-fingerprint stability across the transitional release; re-allocation deferred to a future SP. The 4 trunk-shared baseline slots (411 random_dir_kelly, 413 aux_only, 414 mag_quarter_fixed, 415 trail_only) remain allocated — they need partial-policy-forward access from the main eval pass and stay on the ISV-scalar path until follow-up Task 1.4.b (out of Wave 3 scope). (3) **`grep -rn` audit verified zero non-baseline consumers** of the 4 retired slots: only `baseline_kernels.cu` (writes), `sp15_isv_slots.rs` (constants + layout regression test), `gpu_dqn_trainer.rs` (launcher doc-comments + layout-fingerprint string), and `tests/sp15_phase1_oracle_tests.rs` (asserts) referenced them — every consumer is migrated in this commit. The layout-fingerprint string in `gpu_dqn_trainer.rs::layout_fingerprint_seed` drops the 4 retired entries (BASELINE_BUYHOLD_SHARPE / BASELINE_HOLD_ONLY_SHARPE / BASELINE_NAIVE_MOMENTUM_SHARPE / BASELINE_NAIVE_REVERSION_SHARPE) but keeps the 4 trunk-shared entries (BASELINE_RANDOM_DIR_KELLY_SHARPE / BASELINE_AUX_ONLY_SHARPE / BASELINE_MAG_QUARTER_FIXED_SHARPE / BASELINE_TRAIL_ONLY_SHARPE) — the fingerprint changes with this contract change (intentional, layout-break-class) but is locked again post-commit. (4) **`state_reset_registry` change — no entries to remove**: a `grep -n` audit on `state_reset_registry.rs` + `training_loop.rs::reset_named_state` confirmed the 4 retired baseline slots NEVER had registry entries or dispatch arms in the first place (they were always single-fold-aggregate scalars defaulted at every fold start by the constructor-write that initialises the ISV bus). Task #4 from the dispatch is therefore a no-op — there are no fold-reset arms to remove. The `every_fold_and_soft_reset_entry_has_dispatch_arm` regression test continues to pass unchanged. (5) **4 launcher signatures updated in `gpu_dqn_trainer.rs`**: `launch_sp15_baseline_buyhold` / `launch_sp15_baseline_hold_only` / `launch_sp15_baseline_naive_momentum` / `launch_sp15_baseline_naive_reversion` each replace their `isv: CUdeviceptr` parameter with `out: CUdeviceptr` (per-window output buffer of `[mean, std, raw_sharpe]`). Production callers (Wave 3b) will allocate an `[n_windows × 3]` MappedF32Buffer per baseline and invoke each launcher per-window with the per-window slice (mirror of 1.1.b sharpe_per_bar per-window launch sequence). Doc-comments updated to reflect the contract change and explicitly mark the launchers as ORPHAN (transient 3a-only state). (6) **`action_decoding_helpers.cuh` (NEW)** — single-source-of-truth for the factored-action → direction → position mapping. Two `__device__ __forceinline__` helpers: `factored_action_to_dir_idx(action, b1, b2, b3)` mirrors `trade_physics.cuh::decode_direction_4b` exactly (returns DIR_SHORT=0 / DIR_HOLD=1 / DIR_LONG=2 / DIR_FLAT=3), and `factored_action_to_position(action, b1, b2, b3, prev_position)` lifts the dir-idx to a signed position scalar with explicit Hold-keeps-prev semantics. The header `#include "state_layout.cuh"` for the canonical DIR_* macros, so this and `trade_physics.cuh` share the source-of-truth for direction encoding. Existing consumers (`backtest_env_kernel.cu`, `experience_kernels.cu`) already use `decode_direction_4b` inline; the new derivation kernel is the only consumer of the new helper today. The header does NOT replace `trade_physics.cuh`'s primitives — it adds a higher-level helper on top so derivation kernels avoid pulling in the full Kelly-cap machinery. (7) **`position_history_derivation_kernel.cu` (NEW)** — post-loop derivation of per-bar `position_history` (-1/0/+1), `side_ind` (1.0 on position change), and `rt_ind` (1.0 on transition-to-flat from non-flat) from the recorded `actions_history_buf`. Single-block-per-window launch (`grid=[n_windows,1,1]`, `block=[1,1,1]`): the position state machine is sequential per window because `position(t)` depends on `position(t-1)`, so within-window parallelism is zero by construction; the across-windows axis carries all parallelism. No atomicAdd per `feedback_no_atomicadd` (the kernel is a pure scan, not a reduction). Sentinel handling: `actions_history_buf` is initialised to -1 by `experience_kernels.cu` (per the actions_history measurement-bug fix from a prior commit). Bars where `action == -1` are treated as "no position change" — keep prior position, emit zero side_ind / rt_ind. Downstream Wave 3b consumers gate on `window_lens[w]` to suppress the trailing-sentinel tail. (8) **`build.rs` cubin manifest** gains `position_history_derivation_kernel.cu`; the existing `baseline_kernels.cu` entry's comment block is updated to reflect the per-window output-buffer contract (was ISV-scalar pre-3a). Both compile clean on local RTX 3050 Ti (sm_86) via `CUDA_COMPUTE_CAP=86`. (9) **`launch_sp15_position_history_derivation` (NEW launcher)** in `gpu_dqn_trainer.rs`: free function (matches `launch_sp15_dd_state` / `launch_sp15_baseline_*` precedent) so unit/oracle tests can drive the kernel without the trainer struct. Caches no `CudaFunction` — loads the cubin per-call, same precedent as the baseline launchers (this is a once-per-eval call, not a hot-path producer). New static `SP15_POSITION_HISTORY_DERIVATION_CUBIN` follows the existing static-pub include-bytes pattern. Marked ORPHAN in doc-comment with explicit Wave 3b reference. (10) **5 oracle tests migrated / added in `tests/sp15_phase1_oracle_tests.rs`**: the 4 existing baseline tests (`baseline_buyhold_positive_on_drift`, `baseline_hold_only_emits_zero`, `baseline_momentum_reversion_symmetry`) now allocate a 3-f32 `MappedF32Buffer` per baseline (or two for the symmetry test) and assert against `out[2]` (raw_sharpe) instead of `isv[BASELINE_*_SHARPE_INDEX]`; the hold_only test additionally asserts `out[0] == 0` (hold_only's per-bar PnL stream is identically zero, not just zero-mean noise — exercises the new mean/std outputs that pure-sharpe coverage missed). One new test `position_history_derivation_walks_action_sequence`: 1 window × 8 bars, action sequence Short→Hold→Long→Hold→Flat→Long→Flat→Short (factored ints 0, 27, 54, 27, 81, 54, 81, 0 with b1=b2=b3=3); expected `position = [-1,-1,+1,+1,0,+1,0,-1]`, `side_ind = [1,0,1,0,1,1,1,1]`, `rt_ind = [0,0,0,0,1,0,1,0]`. All 5 tests pass on RTX 3050 Ti via `CUDA_COMPUTE_CAP=86 cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored` (8 oracle tests total, 0 failures). (11) **Transient state — 5 ORPHAN launchers awaiting Wave 3b**: `launch_sp15_baseline_buyhold` / `launch_sp15_baseline_hold_only` / `launch_sp15_baseline_naive_momentum` / `launch_sp15_baseline_naive_reversion` / `launch_sp15_position_history_derivation` are all reachable only from oracle tests in this commit. Wave 3b lands the production callers atomically: `GpuBacktestEvaluator::new` constructor signature gains the per-baseline output buffers; eval-time per-window launch sequence becomes `… → launch_sp15_baseline_*` (×4 per window) `→ launch_sp15_position_history_derivation` (once over all windows) `→ launch_sp15_cost_net_sharpe` (×n_windows, consuming the position_history outputs). The 3a/3b decomposition is documented here so reviewers can connect the kernel-side foundation to the host-side wire-up across the two commits. Touched: `crates/ml/src/cuda_pipeline/baseline_kernels.cu` (4 kernel signatures + helper template + extern-C entry-point bodies), `crates/ml/src/cuda_pipeline/sp15_isv_slots.rs` (4 const removals + 4 layout-test assertions removed + 4 retained-slot assertions added), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (4 launcher param-rename + doc-comment updates + layout-fingerprint string update + new SP15_POSITION_HISTORY_DERIVATION_CUBIN static + new launch_sp15_position_history_derivation), `crates/ml/src/cuda_pipeline/action_decoding_helpers.cuh` (NEW — 2 device helpers, 60 LOC), `crates/ml/src/cuda_pipeline/position_history_derivation_kernel.cu` (NEW — single extern-C kernel, ~70 LOC), `crates/ml/build.rs` (1 new cubin entry + baseline_kernels comment refresh), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (4 baseline-test bodies migrated to output-buffer assertion + 1 new derivation oracle test + import block trimmed). Hard rules: `feedback_no_partial_refactor` (4 ISV slots removed + 4 launcher signatures + 4 oracle tests + new derivation kernel + new shared header + audit doc all in this commit; the only "partial" element is the 5 ORPHAN launchers, explicitly bounded by the documented 3a/3b split), `feedback_no_legacy_aliases` (4 ISV slot constants gone — no compatibility-shim aliases), `feedback_no_atomicadd` (baseline kernels still tree-reduce; derivation kernel is a pure scan with one thread per window), `feedback_no_cpu_compute_strict` (the new derivation kernel runs entirely on the GPU; output buffers are mapped-pinned for the host-readback path used by Wave 3b's HEALTH_DIAG emission), `feedback_no_stubs` (the 5 launchers are real GPU launches against real data; the ORPHAN status is the absence of a production caller, not a stub-return), `pearl_no_host_branches_in_captured_graph` (the derivation kernel is launched OUTSIDE the experience-collection CUDA Graph capture region — Wave 3b will launch it once at end-of-fold from a per-fold scratch path mirroring 1.1.b's sharpe-per-bar end-of-fold launch), `pearl_first_observation_bootstrap` (every window starts with `prev_position = 0` per derivation kernel convention; matches `experience_env_step`'s portfolio-state init that every window starts flat). Verified: `SQLX_OFFLINE=true cargo check -p ml --features cuda` clean (18 unrelated pre-existing warnings); `cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored baseline` 3 tests green; `cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored position_history` 1 test green; `cargo test -p ml --test sp15_phase1_oracle_tests --features cuda -- --ignored unified_sharpe` 3 tests green; `cargo test -p ml --features cuda --lib` holds the 946 pass / 13 fail baseline (no regression). SP15 Wave 2 — fused post-SP11 reward-axis composer (Option β layered architecture, 2026-05-06): closes the deferred-consumer gap left by Phase 3.1 (`r_quality_discipline_split_kernel`, commit `5d36f3238`), Phase 3.3 (`dd_penalty_kernel`), and Phase 3.5.2 (`dd_asymmetric_reward_kernel`) — three SP15 reward-axis kernels that all landed only as standalone scalar producers awaiting deferred consumer wiring per `feedback_no_partial_refactor.md`. Wave 2 chooses Option β (layered, composable post-modifier) over Option α (replace SP11 entirely): SP11 B1b stays canonical "trader-quality" composer that writes `out_rewards`; the SP15 reward-axis composition becomes a fused PER-(i,t) post-modifier read-modify-writing the same buffer in place. This preserves SP11's z-score mag-ratio contract (canary tests untouched) while landing all three deferred SP15 consumers atomically with zero parallel paths. (1) **Architectural decision — fused per-(i,t) post-modifier, NOT three sequential scalar launches**: the three deferred kernels (split composer + dd_penalty + dd_asymmetric_reward) each produced one f32 from ISV reads + a couple of arithmetic ops. Launching three single-thread/single-block kernels per training step plus three ↔-buffer round-trips between them was wasted launch overhead and added three CPU-sync/scratch-pointer plumbing chains. The fused `compute_sp15_final_reward_kernel` does the four-stage composition (α-blend → DD asymmetric reward → quadratic DD penalty → SP12 v3 cap) as inline `__device__` helpers in `sp15_reward_axis_helpers.cuh`, parallel over `[N*2*L]` slots — 3 launches → 1 launch, ~50× more parallelism per step. (2) **Per-step launch sequence in `gpu_experience_collector.rs::collect_experiences_gpu`** (after step 5 dd_state launch from Phase 1.3.b): `launch_experience_env_step` → `launch_sp15_alpha_split_producer` (per-step scalar producer of α from grad-norm ratio, OWNS warm-count increment per Q4) → `launch_sp15_final_reward` (the new fused composer that reads `out_rewards` + `r_discipline_per_sample` + `slot_completed_normally_per_sample` + ISV and writes back). Both new SP15 launches gated on `isv_signals_dev_ptr != 0 && sp15_alpha_warm_count_dev_ptr != 0` — test scaffolds may run env_step without ISV wiring, in which case both buffers receive their default values and the fused composer would be a no-op anyway. (3) **`compute_sp15_final_reward_kernel.cu` (new)**: `extern "C" __global__ void compute_sp15_final_reward_kernel(float* out_rewards, const float* r_discipline_out, const int* slot_completed_normally, float* isv, int N, int L)`. Per-(i,t) parallel over `[N*2*L]`; thread `idx` reads `out_rewards[idx]`, applies `sp15_alpha_blend` → `sp15_dd_asymmetric_reward` → `sp15_dd_penalty` → `sp15_apply_sp12_cap` from `sp15_reward_axis_helpers.cuh`, writes back. NaN/Inf guard skips the write (preserves SP11's value). (4) **`sp15_reward_axis_helpers.cuh` (new)**: 4 `__device__` helpers — `sp15_alpha_blend(r_in, r_disc, isv)` reads ALPHA_SPLIT_INDEX=417, returns `α × r_in + (1 − α) × r_disc`; `sp15_dd_asymmetric_reward(r_in, isv, boost_diag_out)` reads DD_PCT_INDEX=406 + DD_ASYMMETRY_LAMBDA_INDEX=430, returns `r × (1 + λ × dd_pct)` for r>0 else r unchanged, optionally writes R_GAIN_DD_BOOST_INDEX=431; `sp15_dd_penalty(r_in, isv)` reads DD_CURRENT_INDEX=401 + DD_THRESHOLD_INDEX=421 + LAMBDA_DD_INDEX=420, returns `r − λ_dd × max(0, dd − thr)²`; `sp15_apply_sp12_cap(r_in)` returns `fmaxf(REWARD_NEG_CAP, fminf(REWARD_POS_CAP, r))` where the macros come from `state_layout.cuh`. The diagnostic write to slot 431 happens ONLY from a single representative thread (idx==0) to avoid concurrent writes. (5) **Q1 resolution — SP12 caps are `state_layout.cuh` macros, NOT lifted to ISV**: per the user resolution, `REWARD_NEG_CAP=-10.0f` / `REWARD_POS_CAP=+5.0f` are spec'd constants per the SP12 v3 design (asymmetric loss aversion); lifting them to ISV would be a separate refactor outside Wave 2 scope. The fused kernel `#include "state_layout.cuh"` and uses the macros directly via the `sp15_apply_sp12_cap` helper, mirroring the existing `compute_asymmetric_capped_pnl` pattern in `trade_physics.cuh`. (6) **Q2 resolution — uniform shaping over on-policy AND CF slots**: `out_rewards[cf_off]` contains `cf_reward_weighted = w_cf × r_cf` (SP11 controller weight already applied). The fused kernel's per-(i,t) parallelism iterates over the full `[N*2*L]` range; both on-policy and CF slots get the same DD-aware shaping. CF slots map back to the per-(i,t) `r_discipline_out` and `slot_completed_normally` via `idx % (N*L)` — the DD context is per-step, not per-action (`dd_pct` doesn't depend on which action you took, it depends on what state you're in). The model learns "in this DD context, this counterfactual action would have been better/worse" with consistent SP15 shaping. (7) **Q3 resolution — `slot_completed_normally_per_sample[N*L]` flag preserves early-return semantics**: SP11's reward composer takes early-return paths at `experience_kernels.cu:2142` (data-end → `out_rewards = 0.0`) and `:2203` (blown-account → `out_rewards = -10.0`). Those terminal-episode sentinel rewards MUST NOT be SP15-shaped (the bounded clamp would still pass them through unchanged at +0/-10, but applying α-blend to them with a regret-EMA r_discipline would corrupt the sentinel). The new flag buffer defaults to 0 at every kernel entry (alongside the other `*_per_sample` defaults) and is set to 1 ONLY at the end of the normal reward-composition path (right before `out_rewards[out_off] = reward`). The fused kernel `if (slot_completed_normally[idx % (N*L)] == 0) return;` skips early-return slots. The flag is shared between on-policy and CF: if on-policy at (i,t) was an early-return, the CF slot at (i,t) is also skipped. (8) **Q4 resolution — `alpha_split_producer_kernel` OWNS the warm-count increment**: pre-Wave-2 the deleted scalar composer owned the increment. Moving the composer into the fused per-(i,t) kernel would over-tick by N×2×L per step. Per Q4 the warm-count increment moves to the `alpha_split_producer_kernel` — the per-step scalar producer launched once per step from `gpu_experience_collector.rs`, preserving the original "exactly one increment per step" semantics. The producer kernel was renamed from `r_quality_discipline_split_kernel.cu` to `alpha_split_producer_kernel.cu` and dropped the deleted scalar composer; it now runs the warm-count increment unconditionally and gates the formula write on `warm >= N_WARM=100`. (9) **`experience_env_step` signature change — 2 new output params**: `r_discipline_out: float*` ([N*L] f32, NULL-tolerant) gets `isv_signals_ptr[REGRET_EMA_INDEX=423]` written at the end of normal reward composition; `slot_completed_normally_out: int*` ([N*L] i32, NULL-tolerant) gets 0 default at entry and 1 at the same finalisation point. Both writes happen RIGHT BEFORE `out_rewards[out_off] = reward` so the fused composer's reads see consistent values. The CudaSlice/i32 buffers are allocated on `GpuExperienceCollector` mirroring the existing `cf_flip_per_sample` / `slot_live_per_sample` pattern (per-step ephemeral; defaulted at every kernel entry; no fold-boundary registry reset needed). i32 (not u8) for `slot_completed_normally` because no `MappedU8Buffer` exists in `mapped_pinned.rs` and the existing `cf_flip_per_sample` / `trade_close_per_sample` int-flag pattern is the established choice. (10) **3 deleted kernel files**: `r_quality_discipline_split_kernel.cu` (composer + producer; replaced by renamed `alpha_split_producer_kernel.cu` keeping ONLY the producer), `dd_penalty_kernel.cu` (replaced by `sp15_dd_penalty` `__device__` helper), `dd_asymmetric_reward_kernel.cu` (replaced by `sp15_dd_asymmetric_reward` helper). 3 deleted launchers in `gpu_dqn_trainer.rs`: `launch_sp15_r_quality_discipline_split` (composer, scalar variant), `launch_sp15_dd_penalty`, `launch_sp15_dd_asymmetric_reward`. 3 deleted CUBIN statics: `SP15_R_QUALITY_DISCIPLINE_SPLIT_CUBIN`, `SP15_DD_PENALTY_CUBIN`, `SP15_DD_ASYMMETRIC_REWARD_CUBIN`. 2 new CUBIN statics: `SP15_ALPHA_SPLIT_PRODUCER_CUBIN`, `SP15_FINAL_REWARD_CUBIN`. 1 new launcher: `launch_sp15_final_reward`. The retained `launch_sp15_alpha_split_producer` now loads the new `SP15_ALPHA_SPLIT_PRODUCER_CUBIN`. (11) **State reset registry — no new entries**: `r_discipline_per_sample` and `slot_completed_normally_per_sample` are per-step ephemeral (the env_step kernel defaults them at every entry), so cross-fold leakage is impossible without an explicit registry reset. The existing Phase 3.1 ISV slot 417/418/419 + `sp15_alpha_warm_count` registry entries already cover the α path; the existing 420/421/422 entries cover the DD-penalty path; the existing 430/431/432 entries cover the DD-asymmetric path. None of those need new entries — only their CONSUMERS changed (from deleted scalar kernels to inline helpers). (12) **6 new oracle tests** in `sp15_phase1_oracle_tests.rs::mod gpu`: `final_reward_alpha_blend_at_cold_start` (Stage 1: α=0.5 sentinel produces -0.5 from r=1, r_disc=-2), `final_reward_dd_penalty_above_threshold` (Stage 3: dd=0.10, thr=0.05, λ=10 → r=1.0 - 0.025 = 0.975), `final_reward_dd_asymmetric_gain` (Stage 2: r=4 + dd_pct=0.5 + λ=0.5 → r=5.0 boosted; R_GAIN_DD_BOOST=1.25 written by idx==0 thread), `final_reward_dd_asymmetric_loss` (Stage 2: r=-3 unchanged through gain branch), `final_reward_sp12_cap_clamps_both_directions` (Stage 4: r=+20 → +5, r=-20 → -10 via macros), `final_reward_skips_early_return_slots` (Q3: slot_completed=0 + r=-10 sentinel UNCHANGED; slot_completed=1 + r=2.5 SP15-shaped). All 6 drive `compute_sp15_final_reward_kernel` directly with hand-crafted `out_rewards` + `r_discipline_out` + ISV bundles. The 9 OLD oracle tests for the deleted scalar kernels (`r_split_uses_sentinel_alpha_at_cold_start`, `r_quality_subtracts_explicit_cost`, `dd_penalty_quadratic_above_threshold`, `dd_penalty_zero_below_threshold`, 3 `dd_asymmetric_reward_*`) are deleted with the kernels — the new fused-kernel tests cover the same behavioral surface end-to-end. The 2 retained `regret_*` tests (regret_signal_kernel) and 4 `cooldown_*` / 3 `plasticity_*` / 2 `dd_trajectory_*` tests are unchanged (those kernels still have orphan-launcher status pending their own consumer wiring). (13) **Phase 3.2 cost-on-trade-close subtraction NOT replicated in Wave 2**: pre-Wave-2 the deleted composer subtracted `cost_t × trade_close_indicator` from `r_quality` BEFORE the α-blend. The Wave 2 layered architecture re-uses SP11's `out_rewards[out_off]` as `r_in` — and SP11 already charges the model at evaluation-time costs via `cost_anneal_ptr` + the existing inventory / churn / spread-cost shaping. Re-inserting a cost subtraction at the SP15 layer would double-count. The `r_quality_subtracts_explicit_cost` oracle test is deleted accordingly. Touched: `crates/ml/src/cuda_pipeline/r_quality_discipline_split_kernel.cu` (DELETED), `crates/ml/src/cuda_pipeline/dd_penalty_kernel.cu` (DELETED), `crates/ml/src/cuda_pipeline/dd_asymmetric_reward_kernel.cu` (DELETED), `crates/ml/src/cuda_pipeline/alpha_split_producer_kernel.cu` (NEW — producer only, with warm-count increment moved here per Q4), `crates/ml/src/cuda_pipeline/sp15_reward_axis_helpers.cuh` (NEW — 4 `__device__` helpers + slot index `#define`s), `crates/ml/src/cuda_pipeline/compute_sp15_final_reward_kernel.cu` (NEW — fused per-(i,t) composer), `crates/ml/src/cuda_pipeline/experience_kernels.cu` (signature: 2 new output params; defaults at entry; writes at end-of-normal-path), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (3 deleted CUBIN statics, 3 deleted launchers, 2 new CUBIN statics, 1 new launcher; producer launcher renamed/refactored), `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (2 new buffer fields, 2 new alloc lines, 2 new env_step args, 1 new setter `set_sp15_alpha_warm_count_ptr`, +30-line Step 5b block invoking both SP15 launches), `crates/ml/src/trainers/dqn/trainer/training_loop.rs` (1 new wire-up block setting the warm-count dev_ptr on the collector), `crates/ml/build.rs` (3 deleted cubin manifest entries replaced with REMOVED comment blocks; 2 new entries for `alpha_split_producer_kernel.cu` + `compute_sp15_final_reward_kernel.cu`), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (9 deleted scalar-kernel tests; 6 new fused-kernel oracle tests). Hard rules: `feedback_no_partial_refactor` (3 phases' deferred consumers + 3 deletions + 2 new files + 6 new tests + audit doc all in this commit; no parallel paths, no feature flags), `feedback_wire_everything_up` (closes 3 SP15 phase orphan launchers in one commit), `feedback_no_legacy_aliases` (3 scalar-kernel files + their CUBIN+launcher scaffolding all deleted in same commit as their replacement; no compatibility shim left), `feedback_no_atomicadd` (fused kernel is per-(i,t) parallel, pure scalar arithmetic; no reductions, no atomics), `pearl_audit_unboundedness_for_implicit_asymmetry` (gain-only DD multiplier + asymmetric NEG/POS caps preserve loss aversion through the layered pipeline), `pearl_symmetric_clamp_audit` (SP12 cap is bilateral via `fmaxf(NEG, fminf(POS, x))` even though the bounds are intentionally asymmetric per spec), `pearl_no_host_branches_in_captured_graph` (the new SP15 launches happen inside `collect_experiences_gpu::launch_timestep_loop` per-step, OUTSIDE the experience-fwd CUDA Graph capture region — same precedent as Phase 1.3.b's dd_state launch). Verified: NEEDS_VERIFICATION pending L40S CUDA build + oracle-test pass (RTX 3050 Ti smoke verification on next push); SQLX_OFFLINE cargo check + ml-lib non-CUDA suite expected to hold the 946 pass / 13 fail baseline (SP15 buffers are CUDA-feature-gated paths).