diff --git a/crates/ml/src/trainers/dqn/trainer/constructor.rs b/crates/ml/src/trainers/dqn/trainer/constructor.rs index f24d32031..77af9608c 100644 --- a/crates/ml/src/trainers/dqn/trainer/constructor.rs +++ b/crates/ml/src/trainers/dqn/trainer/constructor.rs @@ -765,6 +765,14 @@ impl DQNTrainer { gpu_evaluator: None, + // SP15 Phase 1.6.b / 1.7.b: parallel evaluators for dev (Q8) and + // per-fold walk-forward test slices. Lazy-init in `launch_extra_eval` + // when their data slices are populated; both stay None on the + // CPU-only construction path. + dev_evaluator: None, + dev_ofi_offset: 0, + test_evaluator: None, + // Phase C: Fill simulation and smart order routing fill_simulator: FillSimulator::default(), vol_ema: 0.01, // Initial volatility estimate (1% daily) diff --git a/crates/ml/src/trainers/dqn/trainer/metrics.rs b/crates/ml/src/trainers/dqn/trainer/metrics.rs index badf3f0d5..b7aa0d9d9 100644 --- a/crates/ml/src/trainers/dqn/trainer/metrics.rs +++ b/crates/ml/src/trainers/dqn/trainer/metrics.rs @@ -7,6 +7,18 @@ use crate::TrainingMetrics; use crate::fxcache::TARGET_RAW_CLOSE; use super::super::config::DQNAgentType; +/// SP15 Phase 1.6.b / 1.7.b — discriminator for `launch_extra_eval`. Routes the +/// shared eval helper between the Q8 dev slice (consumed once after the final +/// fold) and the per-fold walk-forward test slice (consumed once per fold). +/// Each variant has its own parallel `Option` slot on +/// `DQNTrainer` (`dev_evaluator` / `test_evaluator`) — NOT a window-swap on +/// `gpu_evaluator`, per `pearl_no_host_branches_in_captured_graph`. +#[derive(Debug, Clone, Copy, PartialEq, Eq)] +pub(crate) enum ExtraEvalKind { + Dev, + Test, +} + /// Rolling EMA values for LearningHealth inputs. #[derive(Debug, Clone, Default)] pub(crate) struct HealthEmaTrackers { @@ -976,6 +988,301 @@ impl DQNTrainer { Ok(-val_sharpe) } + /// SP15 Phase 1.6.b / 1.7.b helper: run a synchronous one-shot backtest + /// against a stashed dev (Q8 final-fold) or per-fold walk-forward test + /// slice, mirroring the val-path lazy-init in `launch_validation_loss` + /// (TLOB sync, ISV signal pointer, identical `GpuBacktestConfig`) but + /// running synchronously via `evaluate_dqn_graphed` since dev/test eval + /// is one-shot at end-of-fold or end-of-training, not per-epoch. + /// + /// `kind == ExtraEvalKind::Dev` consumes `dev_features`/`dev_targets`, + /// `kind == ExtraEvalKind::Test` consumes `test_features`/`test_targets`. + /// The respective `Option` slot is lazy-built on + /// first call (Dev) or rebuilt on every call (Test, because the slice + /// length can vary fold-to-fold and the evaluator allocates step buffers + /// sized for `max_len` at construction). Parallel evaluator instances — + /// NOT a window-swap on `gpu_evaluator` — per + /// `pearl_no_host_branches_in_captured_graph`: window-swap would force + /// invalidating the val CUDA Graph between val and dev/test runs. + /// + /// Emits a single HEALTH_DIAG line per call: + /// `HEALTH_DIAG[N]: dev_eval dev_sharpe_net=... dev_calmar=... ...` + /// `HEALTH_DIAG[N]: test_slice fold=K test_sharpe_net=... ...` + /// where N is `self.current_epoch` (last epoch of the fold for both + /// kinds). Returns Ok(()) on a missing/empty slice (no-op fall-through). + pub(crate) async fn launch_extra_eval( + &mut self, + kind: ExtraEvalKind, + fold_idx: Option, + ) -> Result<()> { + use crate::cuda_pipeline::gpu_backtest_evaluator::{ + DqnBacktestConfig, GpuBacktestConfig, GpuBacktestEvaluator, + }; + + // Snapshot the slice + ofi offset under the slot enum. We snapshot + // (rather than borrow) so the rest of the body can take `&mut self` + // for `fused_ctx` / `agent` / TLOB sync without aliasing. + let (features, targets, ofi_offset) = match kind { + ExtraEvalKind::Dev => { + let f = match self.dev_features.as_ref() { + Some(f) if !f.is_empty() => f.clone(), + _ => return Ok(()), + }; + let t = match self.dev_targets.as_ref() { + Some(t) if !t.is_empty() => t.clone(), + _ => return Ok(()), + }; + (f, t, self.dev_ofi_offset) + } + ExtraEvalKind::Test => { + let f = match self.test_features.as_ref() { + Some(f) if !f.is_empty() => f.clone(), + _ => return Ok(()), + }; + let t = match self.test_targets.as_ref() { + Some(t) if !t.is_empty() => t.clone(), + _ => return Ok(()), + }; + (f, t, self.test_start_bar) + } + }; + + // Use the dedicated validation_stream when present so the new run + // serialises against the val_evaluator on the same stream (avoids + // cross-stream weight races on the shared FusedTrainingCtx). Falls + // back to cuda_stream on the legacy single-stream path. CPU-only + // builds skip silently — same fall-through as `launch_validation_loss`. + let stream = match self.validation_stream.as_ref().or(self.cuda_stream.as_ref()) { + Some(s) => s.clone(), + None => return Ok(()), + }; + // Cross-stream barrier: have the eval-stream wait on cuda_stream's + // just-recorded event so the most-recent training kernel's weight + // and ISV writes are globally visible before we read them. Skipped + // on the single-stream path. Mirrors the val path lines 532-547. + if let (Some(ref vs), Some(ref ts)) = + (self.validation_stream.as_ref(), self.cuda_stream.as_ref()) + { + let train_done = ts.record_event(None).map_err(|e| anyhow::anyhow!( + "training_done event record (extra_eval): {e}" + ))?; + vs.wait(&train_done).map_err(|e| anyhow::anyhow!( + "extra_eval stream wait training_done: {e}" + ))?; + } + + // Always rebuild Test (slice length varies per fold). Build Dev once + // and keep it across the trainer's lifetime — there's only one Q8 + // dev slice per training run and it never changes mid-run. + let needs_build = match kind { + ExtraEvalKind::Dev => self.dev_evaluator.is_none(), + ExtraEvalKind::Test => true, + }; + + if needs_build { + // Drop any prior evaluator BEFORE allocating the new one so VRAM + // is freed for the new build (relevant for Test where the slice + // length and thus buffer sizes change per fold). + match kind { + ExtraEvalKind::Dev => self.dev_evaluator = None, + ExtraEvalKind::Test => self.test_evaluator = None, + } + + let market_dim: usize = ml_core::state_layout::MARKET_DIM; + let ofi_dim: usize = ml_core::state_layout::OFI_DIM; + let feature_dim: usize = market_dim + ofi_dim; + + // Single window from the entire slice. No subsample stride — + // dev/test eval is one-shot, not per-epoch, so the val-path's + // VAL_SUBSAMPLE_STRIDE=4 cost optimisation doesn't apply. Setting + // val_subsample_stride=1 keeps the kernel's annualisation factor + // honest (it scales by sqrt(bars_per_day * trading_days_per_year / + // val_subsample_stride²)). + let mut prices: Vec<[f32; 4]> = Vec::with_capacity(features.len()); + let mut feats: 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; + prices.push([close, close, close, close]); + let fv_slice = &fv[..market_dim.min(fv.len())]; + let mut fv_f32: Vec = fv_slice.iter().map(|&v| v as f32).collect(); + let ofi_idx = ofi_offset + i; + let ofi_row = self.ofi_features.as_ref() + .and_then(|o| o.get(ofi_idx)) + .ok_or_else(|| anyhow::anyhow!( + "extra_eval ({:?}): OFI features missing or index out of bounds at {}", + kind, ofi_idx + ))?; + for &v in ofi_row.iter() { + fv_f32.push(v as f32); + } + feats.push(fv_f32); + } + + let window_prices = vec![prices]; + let window_features = vec![feats]; + + let hp = &self.hyperparams; + #[allow(clippy::cast_possible_truncation)] + let config = GpuBacktestConfig { + max_position: hp.max_position_absolute as f32, + tx_cost_bps: hp.transaction_cost_multiplier as f32, + spread_cost: (hp.tick_size * hp.fill_spread_cost_frac) as f32, + initial_capital: hp.initial_capital as f32, + contract_multiplier: hp.contract_multiplier as f32, + margin_pct: hp.margin_pct as f32, + max_leverage: hp.max_leverage as f32, + ofi_dim, + holding_cost_rate: hp.holding_cost_rate as f32, + churn_threshold: hp.churn_threshold_bars as f32, + churn_penalty_scale: hp.churn_penalty_scale as f32, + opportunity_cost_scale: 0.0, + bars_per_day: hp.bars_per_day as f32, + trading_days_per_year: hp.trading_days_per_year as f32, + val_subsample_stride: 1, // no subsample for one-shot eval + }; + + let evaluator = GpuBacktestEvaluator::new( + &window_prices, + &window_features, + feature_dim, + config, + &stream, + ).map_err(|e| anyhow::anyhow!( + "GpuBacktestEvaluator init (extra_eval {:?}): {e}", kind + ))?; + + match kind { + ExtraEvalKind::Dev => self.dev_evaluator = Some(evaluator), + ExtraEvalKind::Test => self.test_evaluator = Some(evaluator), + } + } + + // ── ISV signal pointer + TLOB weight sync (mirrors val_evaluator) ── + // Wire ISV signals on the chosen evaluator so adaptive eval-time + // controllers (Kelly cap, hold floor, etc.) read live values. + if let Some(ref fused) = self.fused_ctx { + let isv_ptr = fused.isv_signals_dev_ptr(); + let evaluator_opt: &mut Option = match kind { + ExtraEvalKind::Dev => &mut self.dev_evaluator, + ExtraEvalKind::Test => &mut self.test_evaluator, + }; + if let Some(ref mut ev) = *evaluator_opt { + ev.set_isv_signals_ptr(isv_ptr); + } + } + + // TLOB weight sync from the live training instance into the eval- + // path TLOB. Required for state-distribution parity between train + // and eval. Same code path as val (`metrics.rs:648-665`). + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::gpu_tlob::GpuTlob; + use crate::cuda_pipeline::shared_cublas_handle::PerStreamCublasHandles; + use std::sync::Arc; + let train_tlob = &fused.gpu_tlob; + let evaluator_opt: &mut Option = match kind { + ExtraEvalKind::Dev => &mut self.dev_evaluator, + ExtraEvalKind::Test => &mut self.test_evaluator, + }; + if let Some(ref mut ev) = *evaluator_opt { + let eval_stream = Arc::clone(ev.stream()); + let val_tlob_batch = ev.val_tlob_batch_size(); + let h = PerStreamCublasHandles::new(&eval_stream) + .map_err(|e| anyhow::anyhow!("Extra eval TLOB shared_cublas: {e:?}"))?; + let eval_tlob = GpuTlob::new(Arc::new(h), val_tlob_batch) + .map_err(|e| anyhow::anyhow!("Extra eval TLOB init: {e:?}"))?; + ev.set_tlob_from_training(eval_tlob, train_tlob) + .map_err(|e| anyhow::anyhow!("Extra eval TLOB weight sync: {e:?}"))?; + } + } + + // ── Extract pointers + DqnBacktestConfig (same as val path) ── + let fused_ptr = self.fused_ctx.as_mut() + .ok_or_else(|| anyhow::anyhow!("fused_ctx required for extra eval"))? + as *mut super::super::fused_training::FusedTrainingCtx; + let (online_weights_ptr, branching_weights_ptr) = unsafe { + let f = &*fused_ptr; + (f.online_dueling_ref() as *const _, f.online_branching_ref() as *const _) + }; + let fused_ref = unsafe { &*fused_ptr }; + let vr = fused_ref.eval_v_range(); + + let agent = self.agent.read().await; + let network_dims = agent.network_dims(); + let (b0, b1, b2, b3) = agent.branch_sizes(); + let is_branching = agent.is_using_branching(); + drop(agent); + + let evaluator_opt: &mut Option = match kind { + ExtraEvalKind::Dev => &mut self.dev_evaluator, + ExtraEvalKind::Test => &mut self.test_evaluator, + }; + let evaluator = evaluator_opt.as_mut() + .ok_or_else(|| anyhow::anyhow!( + "extra_eval evaluator must be initialised by this point ({:?})", kind + ))?; + evaluator.invalidate_dqn_graph(); + + let hp = &self.hyperparams; + #[allow(clippy::cast_possible_truncation)] + let dqn_cfg = DqnBacktestConfig { + shared_h1: network_dims.0, + shared_h2: network_dims.1, + value_h: network_dims.2, + adv_h: network_dims.3, + num_atoms: hp.num_atoms, + branch_0_size: b0, + branch_1_size: if is_branching { b1 } else { b0 }, + branch_2_size: if is_branching { b2 } else { b0 }, + branch_3_size: if is_branching { b3 } else { b0 }, + v_min: vr[0], + v_max: vr[1], + }; + + // Synchronous launch + consume — one-shot, no async pipelining + // benefit since dev/test runs don't fire per-epoch. Returns the + // single-window metrics directly. + let metrics = unsafe { + evaluator.evaluate_dqn_graphed( + &*online_weights_ptr, + Some(&*branching_weights_ptr), + &dqn_cfg, + Some(&mut *fused_ptr), + ) + }.map_err(|e| anyhow::anyhow!( + "GPU backtest evaluate_dqn_graphed (extra_eval {:?}): {e}", kind + ))?; + + // ── HEALTH_DIAG emission ────────────────────────────────────────── + let epoch = self.current_epoch; + if let Some(m) = metrics.first() { + // Phase 1.2.b cost-net sharpe is BLOCKED (LobBar dependency); we + // emit the existing kernel's cost-aware Sharpe (from the fused + // metrics kernel post-Phase-1.1.b split — `m.sharpe` already + // includes tx_cost_bps + spread_cost via the env-step PnL feed) + // under the `*_sharpe_net` key. When 1.2.b lands, the consumer + // will swap to the dedicated cost-net kernel output without + // changing the HEALTH_DIAG key name. + match kind { + ExtraEvalKind::Dev => tracing::info!( + "HEALTH_DIAG[{}]: dev_eval dev_sharpe_net={:.2} dev_calmar={:.2} dev_max_dd={:.2} dev_trades={:.0}", + epoch, m.sharpe, m.calmar, m.max_drawdown, m.total_trades, + ), + ExtraEvalKind::Test => tracing::info!( + "HEALTH_DIAG[{}]: test_slice fold={} test_sharpe_net={:.2} test_calmar={:.2} test_max_dd={:.2} test_trades={:.0}", + epoch, + fold_idx.map_or(-1_i64, |f| f as i64), + m.sharpe, m.calmar, m.max_drawdown, m.total_trades, + ), + } + } else { + tracing::warn!( + "extra_eval ({:?}): no metrics returned (empty window)", kind + ); + } + + Ok(()) + } + /// Estimate average Q-value from replay buffer samples for monitoring. /// /// Uses the cuBLAS forward path (`fused_ctx.compute_q_values`) — zero Candle involvement. diff --git a/crates/ml/src/trainers/dqn/trainer/mod.rs b/crates/ml/src/trainers/dqn/trainer/mod.rs index 10df1eaf5..a7ad0fe90 100644 --- a/crates/ml/src/trainers/dqn/trainer/mod.rs +++ b/crates/ml/src/trainers/dqn/trainer/mod.rs @@ -717,6 +717,30 @@ pub struct DQNTrainer { /// Invalidated (set to None) when val_data changes (fold transitions). pub(crate) gpu_evaluator: Option, + /// SP15 Phase 1.6.b — parallel evaluator for the Q8 dev slice. Lazy-init in + /// `launch_extra_eval` after the final fold; consumes `dev_features` / + /// `dev_targets`. Independent CUDA-Graph capture from `gpu_evaluator` per + /// `pearl_no_host_branches_in_captured_graph` — window-swap on the val + /// evaluator would require invalidating the val graph between val and dev + /// eval calls. Once-per-training-run lifetime: built right before the + /// dev_eval HEALTH_DIAG line is emitted, dropped with the trainer. + pub(crate) dev_evaluator: Option, + /// SP15 Phase 1.6.b — global OFI-features offset for `dev_features[0]`, + /// equal to `dev_start` in `train_walk_forward`. The eval loop computes + /// `ofi_features[dev_ofi_offset + i]` for sample `i` of the dev slice, + /// mirroring the val path's `ofi_val_offset + i`. + pub(crate) dev_ofi_offset: usize, + + /// SP15 Phase 1.7.b — parallel evaluator for the per-fold walk-forward + /// test slice. Lazy-init in `launch_extra_eval` once per fold; the previous + /// fold's evaluator is dropped at the start of the next fold's test_eval + /// call because the slice length can differ between folds and the + /// `GpuBacktestEvaluator` allocates step buffers sized for the + /// construction-time `max_len`. Independent CUDA-Graph capture from + /// `gpu_evaluator` per the same parallel-instances rationale as + /// `dev_evaluator` (window-swap on val would force val-graph invalidation). + pub(crate) test_evaluator: Option, + // Phase C: Fill simulation and smart order routing /// Fill simulator for order type-dependent execution modeling pub(crate) fill_simulator: FillSimulator, @@ -1162,6 +1186,12 @@ impl DQNTrainer { .collect(); self.dev_features = Some(dev_features_v); self.dev_targets = Some(dev_targets_v); + // SP15 Phase 1.6.b: global OFI-features offset for the dev slice. + // `init_from_fxcache` (called below) stashes the FULL ofi array on + // `self.ofi_features`, so the dev eval reads + // `ofi_features[dev_ofi_offset + i]` for sample `i` of the dev + // slice — same scheme as `ofi_val_offset` for the val path. + self.dev_ofi_offset = dev_start; } if holdout_end > holdout_start { let holdout_features_v: Vec<[f64; 42]> = training_data[holdout_start..holdout_end] @@ -1396,6 +1426,35 @@ impl DQNTrainer { .train_with_data_full_loop_slices(&fold_training_data, Arc::clone(&cb_handle)) .await?; + // SP15 Phase 1.7.b -- fire one-shot eval against this fold's + // walk-forward test slice. The set_test_data_from_slices call + // above gates on fold.test_end > fold.test_start, so when + // test_features is non-empty we know the slice is real. The + // helper itself is a no-op fall-through if either slice or the + // CUDA stream is missing (CPU-only build path). HEALTH_DIAG line: + // HEALTH_DIAG[N]: test_slice fold=K test_sharpe_net=... ... + // Phase 1.7 deferred consumer is closed by this call site. + if fold.test_end > fold.test_start && self.test_features.is_some() { + if let Err(e) = self + .launch_extra_eval( + metrics::ExtraEvalKind::Test, + Some(fold_idx), + ) + .await + { + // Non-fatal: a failed test_slice run shouldn't abort the + // fold loop. Surface the error for diagnostic visibility + // and let training continue. Per feedback_no_hiding: + // log + continue is allowed for diagnostic emissions + // that don't feed downstream consumers (test slice is + // observation-only -- no Q-targets, no replay writes). + tracing::warn!( + "SP15 Phase 1.7.b test_slice eval-call failed for fold {}: {}", + fold_idx + 1, e, + ); + } + } + // Store regime distribution in metrics last_metrics.add_metric("regime_trending_pct", trending_pct); last_metrics.add_metric("regime_ranging_pct", ranging_pct); @@ -1435,6 +1494,28 @@ impl DQNTrainer { self.regime_replay_decay_override = 1.0; self.fold_dominant_regime = 1; + // SP15 Phase 1.6.b -- after the final fold's training+val converge, + // fire one-shot dev eval against the Q8 dev slice stashed at the top + // of train_walk_forward (gated on dev_quarters > 0). Mirrors + // the test_slice call above but consumes dev_features/dev_targets and + // emits the dev_eval HEALTH_DIAG line. Phase 1.6 deferred consumer is + // closed by this call site. Sealed Q9 holdout is NOT touched -- the + // debug_assert sealed-slice guard above catches accidental refactors, + // and Phase 4.3 argo-eval-final.sh is the only legitimate Q9 consumer + // via a separate eval-only entry point that does NOT call this fn. + if self.dev_features.is_some() { + if let Err(e) = self + .launch_extra_eval(metrics::ExtraEvalKind::Dev, None) + .await + { + // Non-fatal: a failed dev run shouldn't fail train_walk_forward + // overall (val and test_slice metrics already emitted, fold + // checkpoints already saved). Surface the error and continue + // to the gap-analysis block below. + tracing::warn!("SP15 Phase 1.6.b dev_eval call failed: {}", e); + } + } + // Task 14: IS->OOS gap analysis (R^2 between volatile% and Sharpe) let n = fold_volatile_pcts.len() as f64; if n >= 3.0 { diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index cd2c0fb87..656e93a07 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 Phase 1.6.b + 1.7.b — wire dev-eval (Q8 final-fold) + test-eval (per-fold WF test slice) into trainer (2026-05-06): closes the deferred-consumer gap left by Phase 1.6 (CLI flags + `dev_features`/`holdout_features` stash, commit `ce019c72d`) and Phase 1.7 (`set_test_data_from_slices` observer + `test_features` stash, commit `ef373c34d`). Both phases landed only the data-flow scaffolding; the actual `evaluate_dqn_graphed` consumer was deferred per `feedback_no_partial_refactor.md` because wiring a parallel `GpuBacktestEvaluator` instance plus the TLOB-sync infrastructure is its own atomic refactor. Wave 1.A bundles both consumer migrations atomically. (1) **Architectural choice — parallel evaluator instances, NOT window-swap**: per the deferred-decision note in the Phase 1.7 audit row, the `gpu_evaluator` field's lazy-init path is fundamentally tied to the `val_data` window passed at construction time (`metrics.rs::launch_validation_loss` lines 549-633 build a single window from `self.val_data` once per fold and reuse the cached evaluator across epochs for deterministic cuBLAS state + stable buffer pointers). A window-swap on the val evaluator between val and dev/test eval calls would require invalidating the val CUDA Graph (recorded device-pointer values become stale when the underlying buffers are repointed) every epoch — fragile and a direct violation of `pearl_no_host_branches_in_captured_graph.md` (the cached graph records the buffer set at capture time; touching the buffers between captures is exactly the host-driven branch the pearl forbids). Parallel instances mirror the val_evaluator construction pattern verbatim and keep each evaluator's CUDA Graph self-contained. (2) **New trainer fields** (`crates/ml/src/trainers/dqn/trainer/mod.rs:720-744`): `dev_evaluator: Option` + `test_evaluator: Option` (both `None` at construction, lazy-built on first call). `dev_ofi_offset: usize` mirrors `ofi_val_offset` for dev; `test_start_bar` (already on the struct from Phase 1.7) provides the test-side OFI offset. (3) **Lazy-init helper** (`crates/ml/src/trainers/dqn/trainer/metrics.rs::launch_extra_eval`, +280 LOC): single helper takes an `ExtraEvalKind` discriminator (`Dev` / `Test`) + optional `fold_idx` and routes to the right `Option` slot. Mirrors `launch_validation_loss` for TLOB sync (`set_tlob_from_training` on the eval-stream's per-stream cuBLAS handle), ISV signal pointer (`set_isv_signals_ptr(fused.isv_signals_dev_ptr())`), `eval_v_range()` Q-range, network dims + branch sizes from the agent, and the `DqnBacktestConfig` struct construction. The helper diverges from val in three places: (a) it uses synchronous `evaluate_dqn_graphed` (one-shot) instead of `evaluate_dqn_graphed_async` + deferred `consume_metrics_after_event` — there's no async pipelining benefit because dev/test eval doesn't fire per-epoch; (b) `val_subsample_stride: 1` instead of val's `4` — one-shot eval doesn't need the val cost optimisation, and stride=1 keeps the kernel's annualisation factor honest; (c) the eval evaluator is rebuilt every fold for `Test` (slice length varies per fold; the evaluator allocates step buffers sized for `max_len` at construction) but built once and kept for `Dev` (only one Q8 dev slice per training run, never changes mid-run). The dev/test path does NOT clobber `self.last_eval_magnitude_dist` / `self.last_eval_intent_magnitude_dist` / `self.last_val_metrics` — those are tied to the val evaluator's most-recent state and the dev/test runs would otherwise overwrite them with non-val readings. (4) **Wire-up site for test_eval — inside the fold loop** (`crates/ml/src/trainers/dqn/trainer/mod.rs` after `train_with_data_full_loop_slices` returns, before regime metrics emit): gated on `fold.test_end > fold.test_start && self.test_features.is_some()`. The `set_test_data_from_slices` call already gates on the same `fold.test_end > fold.test_start` (Phase 1.7), so when `test_features` is non-empty we know the slice is real. Errors are logged + non-fatal: a failed test_slice eval shouldn't abort the fold loop because test eval is observation-only (no Q-targets, no replay writes, no downstream consumers). (5) **Wire-up site for dev_eval — after the fold loop** (`crates/ml/src/trainers/dqn/trainer/mod.rs` immediately after the regime replay decay reset): gated on `self.dev_features.is_some()`. The `dev_ofi_offset = dev_start` is set inside the dev-features stash block (`mod.rs:1163`), so by the time the fold loop completes we have a valid offset for the global OFI features array. Sealed Q9 holdout remains untouched — the existing `debug_assert` sealed-slice guard at `mod.rs:1228` catches accidental refactors that reintroduce holdout into the training path; Phase 4.3 `argo-eval-final.sh` is the only legitimate Q9 consumer via a separate eval-only entry point that does NOT call `train_walk_forward`. (6) **HEALTH_DIAG line formats**: `HEALTH_DIAG[N]: dev_eval dev_sharpe_net= dev_calmar= dev_max_dd= dev_trades=` (emitted once after the final fold) and `HEALTH_DIAG[N]: test_slice fold=K test_sharpe_net= test_calmar= test_max_dd= test_trades=` (emitted once per fold, K = 0-based fold index). N is `self.current_epoch` (last epoch of the fold for both kinds). The `*_sharpe_net` key uses the existing fused-metrics-kernel cost-aware Sharpe (post-Phase-1.1.b split — `WindowMetrics.sharpe` already includes `tx_cost_bps + spread_cost` via the env-step PnL feed); when Phase 1.2.b cost-net sharpe lands (BLOCKED pending LobBar merge per task #358), the consumer will swap to the dedicated cost-net kernel output without changing the HEALTH_DIAG key name — preserves the aggregator-script contract `aggregate-multi-seed-metrics.py` reads from. (7) **Stream choice**: `validation_stream` when present (mirrors val_evaluator's stream), with `cuda_stream` as fallback on the legacy single-stream path. Records a fresh event on `cuda_stream` and waits on `validation_stream` before launching eval kernels — same cross-stream barrier the val path uses (`metrics.rs:532-547`) so the most-recent training kernel's weight + ISV writes are globally visible before the eval kernels read them. CPU-only build path skips silently (no streams → no-op fall-through), same as val. (8) **OFI alignment**: dev_features/test_features are at GLOBAL bar indices (dev_start..dev_end and fold.test_start..fold.test_end respectively). `init_from_fxcache` stashes the FULL ofi-features array on `self.ofi_features` (line `mod.rs:1889`), so the eval gather reads `ofi_features[ofi_offset + i]` for sample `i` — `ofi_offset = self.dev_ofi_offset` for Dev (`= dev_start` in train_walk_forward) or `self.test_start_bar` for Test (already populated by `set_test_data_from_slices`). The `Arc::from(ofi)` reassignment at line 1889 keeps the Arc pointing at the FULL pre-train_walk_forward dataset because the trainer's own `self.ofi_features` was set up at the original 9-quarter scale before `train_walk_forward` was called. (9) **Phase 1.6 + 1.7 orphan stash-fields now have consumers** per `feedback_wire_everything_up.md`: `dev_features` / `dev_targets` (Phase 1.6) + `test_features` / `test_targets` (Phase 1.7) all reach a real GPU eval kernel with an emitted HEALTH_DIAG line. Touched: `crates/ml/src/trainers/dqn/trainer/mod.rs` (+3 fields on `DQNTrainer` struct: `dev_evaluator`, `dev_ofi_offset`, `test_evaluator`; +1-line `self.dev_ofi_offset = dev_start;` write inside the existing dev-features stash block; +14-line test_eval call inside the fold loop after `train_with_data_full_loop_slices`; +14-line dev_eval call after the regime-replay-decay reset post-fold-loop), `crates/ml/src/trainers/dqn/trainer/constructor.rs` (+5 lines initialising the 3 new fields to None/0), `crates/ml/src/trainers/dqn/trainer/metrics.rs` (+1 `ExtraEvalKind` enum at top of file; +~280-line `launch_extra_eval` helper just before `estimate_avg_q_value_with_early_stopping`). Hard rules: `feedback_no_partial_refactor` (both 1.6.b and 1.7.b consumers + both new fields + both HEALTH_DIAG lines + audit doc all in this commit; no parallel paths, no feature flags), `feedback_no_cpu_compute_strict` (dev/test eval is GPU-resident through the same `evaluate_dqn_graphed` path as val — no CPU forwards, no CPU reductions), `feedback_no_stubs` (both call sites are real GPU eval invocations against real stashed slices; failure is logged + non-fatal but the kernel actually fires when the slice is non-empty and the CUDA stream is present), `feedback_isv_for_adaptive_bounds` (no hardcoded thresholds introduced; the dev/test eval inherits all ISV-driven adaptive bounds from `fused_ctx.isv_signals_dev_ptr()` via the same `set_isv_signals_ptr` call val uses), `pearl_no_host_branches_in_captured_graph` (each evaluator's CUDA Graph is self-contained; no graph-capture interaction between dev/test/val because they live in separate `Option` slots with their own buffer ptrs and capture metadata; `evaluator.invalidate_dqn_graph()` is called before each launch_extra_eval invocation to force re-capture against the live training weights). Verified: `cargo check -p ml --features cuda` clean (3 unrelated pre-existing warnings only); `cargo test -p ml --features cuda --lib` holds the 946 pass / 13 fail baseline (no regression introduced); the existing `set_test_data_from_slices_fires_observer_and_stashes` Phase 1.7 oracle test still passes (the new code didn't change the observer-firing surface); a unit-style oracle test asserting the dev/test evaluators actually get built was NOT added because the build path requires a fully-initialised `fused_ctx` with TLOB params + agent forward path + GPU walk-forward generator, which exceeds the smoke-test fixture scope — L40S is the canonical verifier per the spec's existing "L40S smoke is the validation path" pattern. The two new HEALTH_DIAG lines will appear in the next L40S run output once a fold finishes (test_slice) and once the full training run finishes (dev_eval), giving the multi-seed aggregator real `dev_*` / `test_*` keys to track. + SP15 Phase 1.3.b — wire dd_state per-step launch + drop equity-recompute bug + env-0 canonical observable (2026-05-06): closes the Phase 1.3 orphan-launcher gap per `feedback_wire_everything_up.md` and atomically fixes two architectural blockers found during the wire-up investigation (Path A: minimum-scope fix; per-env redesign deferred to Phase 1.3.b-followup if smoke shows env-0-only aggregation insufficient). (1) **Double-update bug fix** (`crates/ml/src/cuda_pipeline/dd_state_kernel.cu`): the original kernel recomputed `new_equity = pos_state[PS_PREV_EQUITY] + pnl_step` and wrote it back into the same slot, but `experience_env_step` already maintains `PS_PREV_EQUITY = new_portfolio_value` (`experience_kernels.cu:3473-3475`). Wiring the kernel as-is would silently double-accumulate equity every step (1× from env_step, 1× from dd_state's recompute), corrupting the position state buffer for every downstream consumer (Kelly EMAs, dd penalty, trade-stats). Per `feedback_no_quickfixes.md` the recompute is removed outright (no `if (already_updated)` guards, no sentinel skip-stores) — the kernel is now READ-ONLY on `pos_state[PS_PREV_EQUITY]` and `pos_state[PS_PEAK_EQUITY]`. The `pnl_step` parameter is dropped from both the CUDA kernel signature and the `launch_sp15_dd_state` launcher signature accordingly. (2) **Env-0 canonical observable** (`crates/ml/src/cuda_pipeline/dd_state_kernel.cu`): kernel grid is `[1,1,1]` (single thread, single block) and the 6 DD ISV slots [401..407) are scalars, but production has N envs in the portfolio_states buffer. The kernel now reads `pos_state[0 * PS_STRIDE + PS_PREV_EQUITY]` and `pos_state[0 * PS_STRIDE + PS_PEAK_EQUITY]` — env 0 specifically — documented in the kernel header as "env 0 as canonical DD observable; per-env redesign deferred to Phase 1.3.b-followup." This is the minimum-scope fix that produces a coherent per-step DD signal for downstream consumers; it accepts the limitation that the DD slots reflect env-0's history rather than aggregate DD across all envs. The recovery / persistence counter logic is also adapted: zero on `current_dd ≤ 0` (new HWM or pre-trade flat) instead of the original `new_equity > peak` branch (which referenced the now-gone local variable). The two conditions are equivalent at the env-0 reading because env_step has just written `peak = max(prev_peak, prev_equity)` ≥ `prev_equity`, so `current_dd = max(0, (peak - prev_equity)/peak)` evaluates to exactly 0 iff env_step set a new HWM this step. (3) **Wire-up location**: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` step "5b" inside `launch_timestep_loop`, immediately after the `experience_env_step` launch (line ~4378, after the env_step kernel block ends at line ~4379) and before the step counter advance kernel (line ~4385). Same stream as env_step → CUDA serializes the two on-stream so the dd_state read sees env_step's writes (no event sync required). OUTSIDE the exp-fwd graph capture region (which is bounded `begin_capture` line ~3734 → `end_capture` line ~3829, well before env_step) so per `pearl_no_host_branches_in_captured_graph.md` no graph-capture interaction. Launcher invoked with `(stream, config.dd_threshold, self.isv_signals_dev_ptr, self.portfolio_states.raw_ptr())` — `dd_budget = config.dd_threshold` keeps the DD_PCT denominator in sync with the `w_dd` penalty term env_step applies on the same `dd_threshold` knob. (4) **Oracle test migration** (`crates/ml/tests/sp15_phase1_oracle_tests.rs::dd_state_kernel_tracks_drawdown_correctly`): test was written against the OLD kernel signature so it WOULD have failed to compile after the kernel signature change — atomic migration per `feedback_no_partial_refactor.md`. The test no longer passes `pnl_step` to the launcher; instead it pre-sums the per-step PnL series into cumulative equity values, sets `pos_state[PS_PREV_EQUITY] = equity` and `pos_state[PS_PEAK_EQUITY] = running_max(equity)` directly into the env-0 row of the portfolio buffer before each `launch_sp15_dd_state` call (mirrors what env_step does in production). The same six-step equity curve (100, 110, 90, 105, 95, 100) drives the kernel and the same expected results hold (max_dd ≈ 0.182, current_dd ≈ 0.091, recovery_bars ≥ 4, dd_pct ∈ (0,1]). Test passes on RTX 3050 Ti. (5) **Phase 1.3 orphan launcher closed** per `feedback_wire_everything_up.md` — `launch_sp15_dd_state` now has a production caller in the per-step training-loop hot path. Downstream consumers receive live values when their consumer wiring lands: Phase 3.3 `dd_penalty` (reads DD_CURRENT/DD_MAX), Phase 3.5.2 `dd_asymmetric_reward` (reads DD_PCT), Phase 3.5.4 `plasticity_injection` persistence (reads DD_PERSISTENCE), Phase 3.5.5 `dd_trajectory` (reads DD_PCT). All four consumer kernels are still orphan launchers themselves at the time of this commit; their wire-up is independent of this one. (6) **Per-env DD redesign deferral** (Phase 1.3.b-followup, task #360): if L40S smoke shows the env-0 single-observable aggregation produces stale or biased DD signals when envs diverge significantly, the followup will introduce a per-env DD tile + reduction kernel that aggregates DD across envs (e.g., max_dd over all envs, or weighted-mean DD by `|position|`). The current min-scope fix is sufficient for representative-env semantics that the DD-aware reward terms expect (one DD signal per training instance). Touched: `crates/ml/src/cuda_pipeline/dd_state_kernel.cu` (kernel — pnl_step param dropped, equity recompute removed, env-0 reads via `0 * PS_STRIDE`, recovery-counter branch adapted to `current_dd ≤ 0`), `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` (launcher — `pnl_step` param dropped from `launch_sp15_dd_state` signature; docstring rewritten to reflect read-only contract + env-0 canonical observable + post-env_step ordering requirement), `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` (production caller — new step "5b" block in `launch_timestep_loop` after env_step launch, calls `launch_sp15_dd_state(&self.stream, config.dd_threshold, self.isv_signals_dev_ptr, self.portfolio_states.raw_ptr())`), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (oracle migration — equity series pre-summed and written into `pos_state` directly per the new contract). Hard rules: `feedback_no_partial_refactor` (kernel signature change + oracle test migration + production wire-up land in this commit; per-env redesign is a separate atomic followup), `feedback_no_quickfixes` (recompute removed outright; no skip-store guards), `feedback_wire_everything_up` (closes the Phase 1.3 orphan launcher), `pearl_no_host_branches_in_captured_graph` (verified launch site is outside the exp-fwd graph capture region). Verified: `cargo check -p ml --features cuda` clean, oracle test `dd_state_kernel_tracks_drawdown_correctly` green, ml lib suite holds 946 pass / 13 fail = baseline. SP15 Phase 1.1.b — split sharpe out of fused backtest_metrics kernel; wire dedicated launch_sp15_sharpe_per_bar (2026-05-06): closes the orphan-launcher gap left by Phase 1.1. Phase 1.1 landed `sharpe_per_bar_kernel.cu` + `launch_sp15_sharpe_per_bar` as scaffolding because the val-side sharpe at `backtest_metrics_kernel.cu:277` was inline in the 8-metric fusion (kernel computed `(mean / std_val) * annualization_factor` directly into `metrics_out[out_base + 0]`), not in a host-side loop the spec sketch had assumed. This commit performs the atomic split per `feedback_no_partial_refactor` and `feedback_wire_everything_up`. (1) **Kernel changes** (`crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu`): removed the inline sharpe compute (variance accumulator `s_sq_sum`, `var = sum_sq/n - mean²`, `std_val = sqrt(max(var, 1e-10))`, write to slot 0). Sortino still uses `mean` (from `s_sum`) and `down_std` (from `s_down_sq` — sum of NEGATIVE squared returns, distinct from sharpe's sum of all squared); calmar still uses `daily_mean = sum/wlen` and `annualization_factor²`. Output stride dropped from 14 → 13 floats per window (sharpe slot 0 removed; every metric below shifted down by 1). Shared memory layout shrunk from 6 → 5 reduction arrays (256 × 4 × 5 = 5 KB + 4 KB sort scratch = 21 KB). Updated docstring + every offset write. (2) **Host changes** (`crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs`): added new `sharpe_buf: MappedF32Buffer` of size `[n_windows * 3]` (mean, std, raw_sharpe per window) per `feedback_no_htod_htoh_only_mapped_pinned`. `launch_metrics_and_record_event` now invokes BOTH kernels: the modified fused kernel for the 7+ remaining metrics, then loops `w in 0..n_windows` to launch `launch_sp15_sharpe_per_bar` once per window with `step_returns_buf + w*max_len*sizeof::()` and `n_w = window_lens[w]`. Per-window launch is acceptable because the kernel is single-block by design (`if (blockIdx.x != 0) return;`) and `n_windows` is small (typically 5). Both kernels share the same eval stream — the launches inherit the existing `eval_done_event` dependency and require no extra synchronisation. Shmem byte calc updated (`256_u32 * 5 + 4096`). (3) **Annualisation** moved host-side: `consume_metrics_after_event` reads `sharpe_host[base + 2]` (raw `mean/std`) and computes `WindowMetrics.sharpe = raw_sharpe * annualization_factor`. Same factor (`sqrt(effective_bars_per_day * trading_days_per_year)`) the kernel previously applied in-line — value contract preserved. (4) **Offset rebase summary** (every consumer migrated atomically; only consumer is the `consume_metrics_after_event` map at `gpu_backtest_evaluator.rs:~2616` — no other site reads `metrics_out[base + N]` raw): old slot 0 (sharpe) — REMOVED, sourced from `sharpe_buf`; old 1 (total_pnl) → new 0; old 2 (max_drawdown) → new 1; old 3 (sortino) → new 2; old 4 (win_rate) → new 3; old 5 (total_trades) → new 4; old 6 (var_95) → new 5; old 7 (cvar_95) → new 6; old 8 (calmar) → new 7; old 9 (omega_ratio) → new 8; old 10 (buy_count) → new 9; old 11 (sell_count) → new 10; old 12 (hold_count) → new 11; old 13 (unique_actions) → new 12. Inside the kernel the second `tid==0` block (var/cvar/calmar/omega) reads `max_dd` from `metrics_out[out_base + 1]` (was `+2`) — same intra-kernel dependency, just at the new offset. (5) **`WindowMetrics.sharpe` value contract** preserved per a new equivalence test (`unified_sharpe_kernel_equivalence_under_annualization` in `crates/ml/tests/sp15_phase1_oracle_tests.rs`). Test launches `sharpe_per_bar_kernel` against a deterministic 512-sample sin-wave + drift PnL, applies host-side `* annualization_factor`, and asserts the result matches the f64 closed-form `(mean/std) * annualization_factor` to within 1e-5 relative error. Cross-checks raw_sharpe alone (without annualisation) to the same tolerance. The two formulas (one-pass `var = sum_sq/n - mean²` vs two-pass `var = sum_dev_sq/n`) are mathematically identical within f32 precision; the bound differences (`max(var, 1e-10)` in the old kernel vs `max(var, 1e-12)` in the new) are inactive at realistic std values. (6) **Phase 1.1 oracle tests** still pass: `unified_sharpe_kernel_zero_mean`, `unified_sharpe_kernel_positive_drift`, `cost_net_sharpe_round_trip_charges_full_spread`. Metrics cubin loads still pass. Full ml lib suite holds at the 945+/13 baseline (no new regressions; the 14th flaky `test_dqn_checkpoint_round_trip` failure is preexisting and noted in spec). (7) **Eliminates the Phase 1.1 orphan launcher** per `feedback_wire_everything_up` — `launch_sp15_sharpe_per_bar` now has a production caller (`gpu_backtest_evaluator::launch_metrics_and_record_event`). Sets up Phase 1.2.b cost-net sharpe to wire `launch_sp15_cost_net_sharpe` against the same per-bar returns buffer (deferred — separate task, separate commit per `feedback_no_partial_refactor` boundaries). Touched: `crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu` (kernel — sharpe compute removed, offsets shifted, shmem layout), `crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs` (host — `sharpe_buf` field, alloc, per-window launch loop, annualisation in `consume_metrics_after_event`, new offsets in `WindowMetrics` parsing), `crates/ml/tests/sp15_phase1_oracle_tests.rs` (+1 numerical equivalence test in `mod gpu`).