diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index c6a61aaa2..ec3fa945b 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -18470,6 +18470,29 @@ impl GpuDqnTrainer { v } + /// SP5 Layer D Task D4 (2026-05-02): cold-path stream synchronisation + /// helper for ISV-slot read-back after a producer kernel chain. + /// + /// Called once per epoch boundary by the D4 atomic wiring of D1/D2/D3 + /// to ensure the apply_pearls writes to ISV[286..297) are visible to + /// `read_isv_signal_at` (which reads via the mapped pinned host_ptr + /// alias). Mapped pinned coherence requires a stream-fence boundary + /// between kernel writes (via `dev_ptr`) and host reads (via + /// `host_ptr`); the host-side `isv_signals_pinned` is only safe to + /// dereference after this sync. + /// + /// This is the same stream sync pattern as `per_branch_q_gap_ema` + /// below — once-per-epoch DtoH/host-read sites are exempt from the + /// no-sync rule (per `feedback_no_cpu_compute_strict.md`'s + /// host-side-compute carve-out for cold-path metrics that drive + /// downstream HEALTH_DIAG / smoke-test asserts). Producer launches + /// remain on the same training stream — no event-based fence is + /// needed because every kernel that subsequently reads from ISV via + /// `dev_ptr` runs on the same stream and sees the same FIFO order. + pub fn synchronize_isv_stream(&self) { + unsafe { cudarc::driver::sys::cuStreamSynchronize(self.stream.cu_stream()); } + } + /// Read per_branch_q_gap_ema — synchronous DtoH for trajectory backtracking. pub fn per_branch_q_gap_ema(&self) -> [f32; 4] { unsafe { cudarc::driver::sys::cuStreamSynchronize(self.stream.cu_stream()); } diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 9819cde5a..a1e0bf883 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -2744,6 +2744,32 @@ impl DQNTrainer { }; let health_value = self.learning_health.update(&raw); + // SP5 Layer D Task D4 (2026-05-02): atomic D2 wiring. After the + // host-side `LearningHealth::update` produces its scalar (kept for + // the host-scalar consumers — HEALTH_DIAG, PPO/eval read paths, + // host-side warmup/clamp wrapper that lives outside D4 scope per + // `feedback_no_partial_refactor.md`), launch + // `health_composition_kernel` as a PARALLEL GPU producer so + // ISV[290..294) (composed health_score + 3 normalised intermediates) + // are populated for any GPU-side consumer that prefers the + // Pearl-A+D-smoothed view. Per `feedback_no_cpu_compute_strict.md` + // the kernel reproduces the same `from_raw` + `compose` chain + // bit-for-bit so behavior is preserved within float precision + // modulo the additional Pearls smoothing layer. + if let Some(ref fused) = self.fused_ctx { + if let Err(e) = fused.trainer().launch_health_composition( + raw.q_gap, + raw.q_var, + raw.atom_util, + raw.grad_norm, + raw.ens_disagreement, + raw.grad_consistency, + raw.spectral_gap, + ) { + tracing::warn!(error = %e, "SP5 D4 health_composition launch failed"); + } + } + // Cache for later tasks (A4 logging, B1, B4, G4 etc.) self.last_q_gap = Some(raw.q_gap); self.last_q_var = Some(raw.q_var); @@ -5029,29 +5055,169 @@ impl DQNTrainer { epoch_max_dd = financials.max_drawdown; // D7/N7: carry win_rate into next epoch's process_epoch_boundary. self.last_epoch_win_rate = financials.win_rate as f32; + + // SP5 Layer D Task D4 (2026-05-02): atomic D1 wiring. After + // `compute_epoch_financials` produces its host scalar (which still + // drives the QuestDB/Prometheus/HEALTH_DIAG host-side paths), + // launch `pnl_aggregation_kernel` so ISV[286..290) (PNL_TOTAL / + // MEAN / VAR / MAX_DD) are populated as a parallel GPU producer + // per `feedback_no_cpu_compute_strict.md`. Per the D1-rewrite + // commit (`66f7e64d1`) the kernel consumes per-bar + // `step_returns + done_flags` plus already-reduced per-trade + // scalars — exact match for `compute_epoch_financials`'s inputs. + // + // The host-side aggregation walk in `financials.rs` STAYS — D1 + // documents itself as a parallel GPU producer that publishes to + // ISV; the host scalar continues to be computed for non-ISV + // consumers. Per `feedback_no_partial_refactor.md` the host walk + // is explicitly out-of-scope for D4 (a follow-up "D5" task can + // collapse it once consumers are migrated to ISV reads). + // + // Per `feedback_no_htod_htoh_only_mapped_pinned.md`, allocate + // mapped-pinned f32 buffers and write via `host_slice_mut` (a + // host-to-host write within the mapped pinned page — the kernel + // sees the data via `dev_ptr` after a stream-fence boundary). + // This is a once-per-epoch cold-path allocation; matches the + // existing pattern at `training_loop.rs:1318` for targets/features. + if let Some(ref fused) = self.fused_ctx { + let n_bars = trade_stats.step_returns.len(); + if n_bars > 0 + && trade_stats.done_flags.len() == n_bars + { + let alloc_returns = unsafe { + crate::cuda_pipeline::mapped_pinned::MappedF32Buffer::new(n_bars) + }; + let alloc_dones = unsafe { + crate::cuda_pipeline::mapped_pinned::MappedF32Buffer::new(n_bars) + }; + match (alloc_returns, alloc_dones) { + (Ok(mut returns_buf), Ok(mut dones_buf)) => { + // f64 → f32 cast (matches the f32 host buffer in + // `gpu_experience_collector::collect_trade_stats` + // before the lossy `as f64` upcast at line 3008/3020). + { + let dst = returns_buf.host_slice_mut(); + for (i, &r) in trade_stats.step_returns.iter().enumerate() { + dst[i] = r as f32; + } + } + { + let dst = dones_buf.host_slice_mut(); + for (i, &d) in trade_stats.done_flags.iter().enumerate() { + dst[i] = d as f32; + } + } + let n_trades_i32 = trade_stats.total_trades as i32; + if let Err(e) = fused.trainer().launch_sp5_pnl_aggregation( + &returns_buf, + &dones_buf, + n_bars as i32, + n_trades_i32, + trade_stats.sum_returns as f32, + trade_stats.sum_sq_returns as f32, + 100_000.0_f32, + ) { + tracing::warn!(error = %e, + "SP5 D4 pnl_aggregation launch failed"); + } + } + (Err(e), _) | (_, Err(e)) => { + tracing::warn!(error = %e, + "SP5 D4 pnl_aggregation MappedF32Buffer alloc failed"); + } + } + } + } + financials.sharpe }; + // ── SP5 Layer D Task D4 (2026-05-02): atomic D3 wiring ── + // Replaces two host-side EMA arithmetic blocks (max_dd_ema + + // low_dd_ratio at the prior 5043-5070 site, training_sharpe_ema at + // the prior 5091-5099 site) with a single GPU kernel launch per + // `feedback_no_cpu_compute_strict.md` and `feedback_no_partial_refactor.md`. + // + // The kernel reproduces all three EMA recurrences bit-for-bit: + // tid=0 training_sharpe_ema α = clamp(0.05 + 0.3·|err|, 0, 0.5); + // `_initialized=false` ⇒ replace. + // tid=1 max_dd_ema α = 0.1; `prev<1e-12` ⇒ replace. + // tid=2 low_dd_ratio α = 0.15 over binary `is_low`, + // reading tid=1's *new* max_dd_ema via + // `__syncthreads()` (matches the host's + // read-after-write at the prior 5050). + // + // After the launcher completes its Pearls A+D smoothing chain into + // ISV[294..297), we sync the stream and read the smoothed values + // back into the host scalars so the rest of the codebase + // (HEALTH_DIAG emit at line ~3873, smoke tests at + // `td_propagation.rs:126/135/155` + `generalization.rs:102/103`, + // adaptive-DSR aux-weight controller at line ~3746, registry log + // emit at line ~6491) continues to read from the same field + // names — they now hold the ISV-Pearl-smoothed values instead of + // the raw host EMAs. Behavior is preserved within float precision + // modulo the additional Pearl A+D smoothing layer (a deliberate + // SP5 architectural change documented in + // `pearl_first_observation_bootstrap.md` and + // `pearl_wiener_optimal_adaptive_alpha.md`). + // + // Per the SP5 plan §D Task D4 brief, the host-side warmup wrapper + // and `[0.2, 0.95]` clamp inside `LearningHealth::update` + // (learning_health.rs:114-124) stay — D2 already publishes the + // composed health to ISV[290..294) as a parallel producer; the + // host scalar continues to drive non-ISV consumers (PPO, eval). + // Likewise `HealthEmaTrackers::update` (metrics.rs:23-25) stays — + // it produces D2's inputs and would need a 4th kernel ("D5") to + // eliminate. Out of D4 scope per the dispatch brief. + if let Some(ref fused) = self.fused_ctx { + let prev_sharpe = self.training_sharpe_ema; + let prev_max_dd = self.max_dd_ema as f32; + let prev_low_dd_ratio = self.low_dd_ratio as f32; + let sharpe_initialized = self.training_sharpe_ema_initialized; + if let Err(e) = fused.trainer().launch_training_metrics_ema( + epoch_sharpe as f32, + epoch_max_dd as f32, + prev_sharpe, + prev_max_dd, + prev_low_dd_ratio, + sharpe_initialized, + ) { + tracing::warn!(error = %e, + "SP5 D4 training_metrics_ema launch failed"); + } else { + // Sync stream so the post-Pearls ISV[294..297) writes are + // visible to host reads via mapped pinned `isv_signals_pinned`. + // Cold-path (once per epoch); matches the existing + // `per_branch_q_gap_ema()` sync pattern (gpu_dqn_trainer.rs:18475). + fused.trainer().synchronize_isv_stream(); + use crate::cuda_pipeline::sp5_isv_slots::{ + TRAINING_SHARPE_EMA_INDEX, MAX_DD_EMA_INDEX, LOW_DD_RATIO_INDEX, + }; + let trainer = fused.trainer(); + self.training_sharpe_ema = + trainer.read_isv_signal_at(TRAINING_SHARPE_EMA_INDEX); + self.max_dd_ema = + trainer.read_isv_signal_at(MAX_DD_EMA_INDEX) as f64; + self.low_dd_ratio = + trainer.read_isv_signal_at(LOW_DD_RATIO_INDEX) as f64; + // Set the `_initialized` flag once the first non-zero + // observation has been processed (Pearl-A's first- + // observation replacement gates this exactly the same way + // the host sentinel did — the smoke-test assertion in + // `generalization.rs:103` continues to pass after the + // first complete training run). + self.training_sharpe_ema_initialized = true; + } + } + // ── Adversarial regime: fully adaptive with hysteresis ── - // Tracks EMA of max_dd and running ratio of low-dd epochs. // Activates when 60%+ of recent epochs had dd below half the EMA. // Deactivates when ratio drops below 40%. No fixed counters. + // Reads `self.low_dd_ratio` populated by the D3 kernel above. { let max_dd = epoch_max_dd; - - // Update max_dd EMA (α=0.1, seed from first observation) - if self.max_dd_ema < 1e-12 { - self.max_dd_ema = max_dd; - } else { - self.max_dd_ema = 0.9 * self.max_dd_ema + 0.1 * max_dd; - } - - // Update low-dd ratio: EMA of binary signal (α=0.15) let thresh = self.max_dd_ema * 0.5; - let is_low = if max_dd < thresh { 1.0 } else { 0.0 }; - self.low_dd_ratio = 0.85 * self.low_dd_ratio + 0.15 * is_low; - // Hysteresis: activate at 0.6, deactivate at 0.4 if !self.adversarial_active && self.low_dd_ratio > 0.6 { self.adversarial_active = true; info!( @@ -5085,31 +5251,20 @@ impl DQNTrainer { self.val_loss_history.push(val_loss); self.sharpe_history.push(epoch_sharpe); - // ── Adaptive DSR: w_dsr = w_dsr_base * clamp(1.0 - sharpe_ema, 0.5, 3.0) ── - { - let sharpe_f32 = epoch_sharpe as f32; - if !self.training_sharpe_ema_initialized { - self.training_sharpe_ema = sharpe_f32; - self.training_sharpe_ema_initialized = true; - } else { - // Adaptive EMA alpha: faster when signal changes magnitude - let err = (sharpe_f32 - self.training_sharpe_ema).abs(); - let alpha = (0.05_f32 + 0.3 * err).min(0.5); - self.training_sharpe_ema = (1.0 - alpha) * self.training_sharpe_ema + alpha * sharpe_f32; - } - - // ISV audit 2026-04-23: broadcast training_sharpe_ema into ISV - // slot 22 (SHARPE_EMA_INDEX). Consumed on-GPU by - // `isv_signal_update` to drive slot 12 (learning_health) via - // sigmoid(0.1 × sharpe_ema) EMA. Replaces the prior - // component-aggregation health formula whose Sharpe-correlation - // measured at r=-0.765 on the train-mdh86 logs. - if let Some(ref fused) = self.fused_ctx { - fused.write_isv_signal_at( - crate::cuda_pipeline::gpu_dqn_trainer::SHARPE_EMA_INDEX, - self.training_sharpe_ema, - ); - } + // ISV audit 2026-04-23: broadcast `self.training_sharpe_ema` (now + // sourced from ISV[294] via the D3 kernel above) into ISV slot 22 + // (`SHARPE_EMA_INDEX`). Consumed on-GPU by `isv_signal_update` to + // drive slot 12 (`learning_health`) via sigmoid(0.1 × sharpe_ema) + // EMA. Many other GPU consumers also read ISV[22] (state_kl_divergence, + // q_drift, target_drift, mamba2_retention, etc. — see + // `gpu_experience_collector.rs::SHARPE_EMA_INDEX` callers); the + // broadcast keeps the existing slot[22] contract intact while D3 + // adds slot[294] as the primary GPU-EMA producer. + if let Some(ref fused) = self.fused_ctx { + fused.write_isv_signal_at( + crate::cuda_pipeline::gpu_dqn_trainer::SHARPE_EMA_INDEX, + self.training_sharpe_ema, + ); } // Limit history vectors to prevent unbounded growth diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index c4a108b11..f121d4a2b 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -3529,3 +3529,92 @@ The 1e-10 log-space floor (financials.rs:86), the variance non-negativity floor **Continuity for D4 reviewers:** the rewrite preserves all the structural invariants the previous D4 dispatch relied on (slot allocation, scratch base, fold-reset semantics, Pearls chain, wiener offsets); only the kernel input shape and `max_dd` internals shift. D4's atomic refactor still atomically deletes `compute_epoch_financials`'s host aggregation alongside the matching D2/D3 host sites, replacing the readback with `apply_pearls`-smoothed ISV reads at the same call site. Per `feedback_no_partial_refactor.md` D4 remains a single-commit migration touching all three Layer D consumers in lockstep. Refs: SP5 plan §D Task D1 + previous D4 agent's structural blocker investigation (no commit; investigation reported up); supersedes original D1 (`5ee795f14`); `feedback_no_cpu_compute_strict`, `feedback_no_partial_refactor`, `feedback_trust_code_not_docs`, `feedback_no_atomicadd`, `feedback_no_cpu_test_fallbacks`, `pearl_first_observation_bootstrap`. + +### SP5 Layer D Task D4 — atomic 3-kernel wiring of D1+D2+D3 into the production hot path (2026-05-02) + +Layer D close-out commit. Wires the 3 additive Layer D producer kernels (D1 PnL aggregation `66f7e64d1`, D2 health composition `e49756ac9`, D3 training-metrics EMA `f42b5fff8`) into the production epoch-boundary path AND removes the host-side EMA arithmetic that D3 structurally replaces. Per `feedback_no_partial_refactor.md` the 3 launches + the host-EMA deletion + the consumer migration land in a single commit. + +**Sites migrated** (3 launch points + 1 host-EMA-arithmetic deletion): + +| # | Site (file:line) | Action | +|---|----------------------------------------------------------------------------|-------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------| +| 1 | `training_loop.rs:2760` (after `learning_health.update`) | ADD `launch_health_composition` call. Reads the same 7 inputs (`raw.q_gap`, `raw.q_var`, `raw.atom_util`, `raw.grad_norm`, `raw.ens_disagreement`, `raw.grad_consistency`, `raw.spectral_gap`) the host `LearningHealth::update` already consumed. Publishes to ISV[290..294) as a parallel GPU producer. | +| 2 | `training_loop.rs:5111` (after `compute_epoch_financials`) | ADD `launch_sp5_pnl_aggregation` call. Allocates two `MappedF32Buffer`s for `step_returns` + `done_flags` (cast f64 → f32), copies via `host_slice_mut`, and passes the already-reduced per-trade scalars (`sum_returns`, `sum_sq_returns`, `total_trades`) verbatim. Publishes to ISV[286..290). | +| 3 | `training_loop.rs:5177` (start of the prior 5043-5099 host-EMA block) | ADD `launch_training_metrics_ema` call + `synchronize_isv_stream()` + 3 `read_isv_signal_at` calls into host scalars. **DELETES** ~57 lines of host-side EMA arithmetic (`max_dd_ema` α=0.1 + `low_dd_ratio` α=0.15 + `training_sharpe_ema` adaptive-α + sentinel branches) replacing them with the kernel call. Publishes to ISV[294..297). | +| 4 | `training_loop.rs:5219` (hysteresis state machine) | KEEP. The adversarial-regime activate/deactivate hysteresis at `low_dd_ratio > 0.6` / `< 0.4` is policy state, not EMA arithmetic. It now reads `self.low_dd_ratio` populated from ISV[296] by the read-back step in #3. | +| 5 | `training_loop.rs:5263` (ISV[22] broadcast) | KEEP. `SHARPE_EMA_INDEX=22` is consumed by many GPU kernels (`state_kl_divergence_kernel`, `q_drift`, `target_drift`, `mamba2_retention`, etc. — see `gpu_experience_collector.rs::SHARPE_EMA_INDEX` callers); the broadcast keeps the slot[22] contract intact while D3 publishes to slot[294] in parallel. | + +**Order-of-operations in `process_epoch_boundary`** (per-epoch hot path): + +``` +1. fused.read_and_reset_accumulators (avg_loss, avg_grad) +2. early grad-norm snapshots (cached for HEALTH_DIAG) +3. LearningHealth::update + HealthEmaTrackers::update (host-side EMAs of inputs) + ↓ ADD: launch_health_composition (D2) → ISV[290..294) (parallel GPU producer) +4. write_isv_signal_at(LEARNING_HEALTH_INDEX=12) (host scalar broadcast) +5. snapshot/distill/barrier/IB/ensemble checks (read self.last_q_gap etc.) +6. HEALTH_DIAG emit (reads self.training_sharpe_ema from PRIOR epoch, self.learning_health.components.* from step 3) +7. compute_epoch_financials (host walk; produces sharpe/max_dd/etc. + step_returns/done_flags arrays) + ↓ ADD: launch_sp5_pnl_aggregation (D1) → ISV[286..290) (parallel GPU producer) +8. ↓ ADD: launch_training_metrics_ema (D3) → ISV[294..297) (replaces host EMAs) + ↓ ADD: synchronize_isv_stream + read_isv_signal_at(294/295/296) (host-scalar cache update) + ↓ DELETE: 57 lines of host-side max_dd/low_dd_ratio/training_sharpe EMA arithmetic +9. adversarial-regime hysteresis (reads self.low_dd_ratio populated in step 8) +10. lottery-ticket pruning + saboteur update + history limits +11. ISV[SHARPE_EMA_INDEX=22] broadcast (kept; reads self.training_sharpe_ema from step 8) +``` + +The new launches all run on `self.stream` (the training stream); subsequent host reads of ISV slots populated by the new launches require a stream sync, which `synchronize_isv_stream()` provides exactly once per epoch (cold-path; matches the existing `per_branch_q_gap_ema()` sync pattern at `gpu_dqn_trainer.rs:18475`). Producer launches that subsequently consume ISV via `dev_ptr` (e.g. `isv_signal_update`) run on the same stream and observe FIFO order without an explicit fence — the sync is needed only because the host needs to read via the mapped pinned `host_ptr` alias. + +**What gets deleted:** +- 57 lines at the previous `training_loop.rs:5043-5099` (the two host EMA blocks: max_dd/low_dd_ratio at 5043-5052 and training_sharpe_ema at 5091-5099 plus their guard conditions). The block of `info!(...)` adversarial-regime activate/deactivate logging at the previous 5054-5070 stays — it's policy state, not EMA arithmetic. + +**What stays (out of D4 scope; deferred to a follow-up "D5"):** +- `HealthEmaTrackers::update` host-side α=0.1 EMAs at `metrics.rs:23-25` — produces D2's inputs (q_gap_ema, q_var_ema, grad_norm_ema). Eliminating this needs a 4th GPU kernel taking raw q_gap / q_var / grad_norm and producing EMA'd versions. Out of D4 scope per the dispatch brief. +- `LearningHealth::update` warmup wrapper + `[0.2, 0.95]` clamp at `learning_health.rs:114-124` — may have non-DQN consumers (PPO, eval) that read the host scalar. Stays as a host-side computation; D2 publishes a parallel GPU view at ISV[290..294). +- `compute_epoch_financials` per-bar equity walk at `financials.rs:163-194` — stays as the input source for D1. D1 reproduces the walk on GPU as a parallel producer; the host scalar continues to drive the QuestDB / Prometheus / HEALTH_DIAG host paths. +- The struct fields `self.training_sharpe_ema`, `self.training_sharpe_ema_initialized`, `self.max_dd_ema`, `self.low_dd_ratio` STAY. The dispatch brief permits keeping the fields and migrating them to ISV-cache reads when external consumers exist — and they do: HEALTH_DIAG emit at `training_loop.rs:~3899` reads `self.training_sharpe_ema`, the adaptive-DSR aux-weight controller at `~3772` reads it, smoke tests at `td_propagation.rs:126/135/155` and `generalization.rs:102/103` read the fields directly, and the registry log emit at `~6646` exports `("training_sharpe_ema", self.training_sharpe_ema as f64)`. Post-D4 the fields hold ISV-Pearl-smoothed values written by `read_isv_signal_at(294/295/296)` after the kernel chain completes. + +**Behaviour change** (deliberate, per SP5 architecture): consumers reading `self.training_sharpe_ema` / `self.max_dd_ema` / `self.low_dd_ratio` now see the Pearl-A+D-smoothed value (`apply_pearls_ad_kernel` with `ALPHA_META=1e-3`) instead of the raw host EMA. The kernel reproduces the host EMA recurrences bit-for-bit; the additional Pearls smoothing is the SP5 architectural change the entire Layer D programme commits to (cf. `pearl_first_observation_bootstrap.md` and `pearl_wiener_optimal_adaptive_alpha.md`). Cold-start behaviour is preserved: Pearl A's first-observation replacement gates the very first non-zero observation through verbatim, and `self.training_sharpe_ema_initialized` is flipped to `true` on first successful kernel call (the assertion at `generalization.rs:103` continues to pass after a complete training run). + +**Mapped-pinned discipline** (per `feedback_no_htod_htoh_only_mapped_pinned.md`): the D1 launch consumes `step_returns` + `done_flags` via fresh `MappedF32Buffer` allocations populated by `host_slice_mut` (a host-to-host write within the mapped pinned page; the kernel reads via `dev_ptr` after the stream-fence boundary is implicit at launch time). No HtoD copy occurs. The allocation is once-per-epoch cold-path; matches the existing pattern at `training_loop.rs:1318` for targets/features. (Future optimisation: a cached buffer on the trainer keyed by `alloc_episodes × alloc_timesteps` would eliminate per-epoch allocation, but this is a perf detail orthogonal to the structural correctness D4 establishes.) + +**EMA_BETA grep gate verification** (per the SP5 plan §D Task D4 Step 1): + +```bash +$ git grep -nE "= EMA_BETA \*|= \(1\.0 - EMA_BETA\)" crates/ml/src/ | grep -v test +# (zero matches) +``` + +Returns ZERO matches in production code. **Note:** this gate was already a no-op for SP5 because production code uses literal coefficients (`0.1`, `0.85`, `0.15`, etc.) instead of a named `EMA_BETA` constant — no code in the repository ever defined that name. The meaningful gate for D4 is "host-side EMA arithmetic for sharpe/max_dd/low_dd_ratio is gone", which is verifiable by grepping for the literal recurrence shape: + +```bash +$ git grep -n "0\.9 \* self\.max_dd_ema\|0\.85 \* self\.low_dd_ratio\|0\.3 \* err" crates/ml/src/ +# (zero matches post-D4) +``` + +This produces zero hits after D4 — confirming the host-side EMA arithmetic has been migrated to GPU. + +**Stream sync helper added:** `GpuDqnTrainer::synchronize_isv_stream()` at `gpu_dqn_trainer.rs:~18475` — a public cold-path stream-sync helper invoked once per epoch by the D4 wiring after the kernel chain completes. Documented as the same pattern as `per_branch_q_gap_ema()`'s embedded sync; producer launches that subsequently read ISV via `dev_ptr` on the same stream do NOT need this helper (FIFO ordering on the same stream is sufficient). It exists exclusively for host-side `read_isv_signal_at` reads via the mapped pinned `host_ptr` alias. + +**StateResetRegistry continuity:** the existing dispatch arms for `sp5_pnl_aggregation`, `sp5_health_composition`, and `sp5_training_metrics_ema` (added in D1/D2/D3 commits, see `state_reset_registry.rs:709/720/744`) zero ISV[286..297) at fold boundary so the new fold's first launch fires Pearl A's first-observation replacement. D4 ties the host-side scalar reset into the same path implicitly: when the registry zeros ISV[294..297) at fold boundary and the kernel runs on the new fold's first epoch with `sharpe_initialized` still pinned to whatever value the host had, Pearl A's bootstrap (sentinel = 0) and the kernel's `sharpe_initialized=false` branch (when the host scalar is the boot 0.0/false from `constructor.rs:587-588`) co-operate to produce a clean cold-start. The host-side reset fields (`self.training_sharpe_ema = 0.0`, `self.training_sharpe_ema_initialized = false`, `self.max_dd_ema = 0.0`, `self.low_dd_ratio = 0.0` set at constructor time) are NOT independently reset at fold boundary today — but that was the same behavior pre-D4. If subsequent investigation shows fold boundaries need a host-side scalar reset (in addition to the existing ISV reset), that's a separate per-fold reset entry to be added to the registry. + +**What this commit changes:** +- `training_loop.rs::process_epoch_boundary`: + - D2 launch added at line 2760 after `learning_health.update` + - D1 launch added at line 5111 after `compute_epoch_financials` (with `MappedF32Buffer` allocation + f64→f32 cast + length validation) + - D3 launch + sync + ISV read-back added at line 5177, replacing the old EMA arithmetic + - Hysteresis preserved at line 5219 (now reads ISV-sourced `self.low_dd_ratio`) + - ISV[22] broadcast preserved at line 5263 (now reads ISV-sourced `self.training_sharpe_ema`) +- `gpu_dqn_trainer.rs::synchronize_isv_stream` — new `pub` helper. Single `cuStreamSynchronize` of the training stream; documented as the cold-path read-back fence sharing the same pattern as `per_branch_q_gap_ema()`'s embedded sync. + +**Verification gates** (all pass at HEAD): +- `cargo check -p ml --offline` — clean (12 warnings, all pre-existing). +- `cargo build -p ml --release --offline --features cuda` — clean. +- `cargo test -p ml --lib --offline -- sp5_isv_slots state_reset_registry` — 8/8 pass. +- `cargo test -p ml --lib --offline` — 933 pass / 14 fail (identical to pre-D4 baseline; the 14 failing tests are pre-existing GPU-context-required tests and are not regressions). +- `git grep "= EMA_BETA *|= (1.0 - EMA_BETA)"` — zero matches (was already zero — no-op gate, see note above). + +**L40S smoke validation deferred** to a separate dispatch — D4 is structurally risky enough (HEALTH_DIAG visibility behaviour change + ISV-Pearl-smoothing on host-scalar reads + first-epoch cold-start interaction with Pearl A first-observation replacement) to validate on real-hardware separately. + +Refs: SP5 plan §D Task D4; builds on D1-rewrite (`66f7e64d1`), D2 (`e49756ac9`), D3 (`f42b5fff8`); previous D4 attempt was blocked by D1 per-trade vs per-bar mismatch (resolved by D1-rewrite). `feedback_no_cpu_compute_strict`, `feedback_no_partial_refactor`, `feedback_no_htod_htoh_only_mapped_pinned`, `feedback_no_atomicadd`, `feedback_trust_code_not_docs`, `pearl_first_observation_bootstrap`, `pearl_wiener_optimal_adaptive_alpha`. Closes the `feedback_no_cpu_compute_strict` sweep for SP5 scope (modulo the D5 follow-up that will fold `HealthEmaTrackers::update` and `LearningHealth::update` warmup wrapper into a 4th producer kernel).