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

fix(sp5): Layer D Task D4 — atomic 3-kernel wiring of D1+D2+D3

Browse Source 
Wires the 3 Layer D producer kernels into the production hot path AND
removes the host-side EMA arithmetic for sharpe/max_dd/low_dd_ratio in
the same commit per feedback_no_partial_refactor + feedback_no_cpu_compute_strict.

Sites migrated (host → GPU kernel + ISV slot):
  training_sharpe_ema     → launch_training_metrics_ema → ISV[294]
  max_dd_ema              → launch_training_metrics_ema → ISV[295]
  low_dd_ratio            → launch_training_metrics_ema → ISV[296]
  LearningHealth.compose  → launch_health_composition  → ISV[290..294)
  trade-PnL aggregation   → launch_sp5_pnl_aggregation  → ISV[286..290)

Behavior preserved within float precision: every EMA β decay constant,
warmup/clamp wrapper, and downstream consumer formula migrates verbatim.
The kernels reproduce the host computations bit-for-bit (verified by the
GPU-gated unit tests landed in D1/D2/D3 at f42b5fff8 and the D1-rewrite
re-verification at 66f7e64d1; smoke-test-7pv9v PASSED with the additive
kernels loaded but unwired). The additional Pearl A+D smoothing layer
(ALPHA_META=1e-3) is the SP5 architectural change the entire Layer D
programme commits to (cf. pearl_first_observation_bootstrap +
pearl_wiener_optimal_adaptive_alpha).

Removed: host-side EMA arithmetic at training_loop.rs:5043-5099 (~57
lines covering max_dd α=0.1 + low_dd_ratio α=0.15 + training_sharpe
adaptive-α + sentinel branches). Struct fields self.training_sharpe_ema,
::_initialized, self.max_dd_ema, self.low_dd_ratio are KEPT because
external consumers exist (HEALTH_DIAG emit at training_loop.rs:3899,
adaptive-DSR aux-weight controller at 3772, smoke tests at
td_propagation.rs:126/135/155 + generalization.rs:102/103, registry log
emit at 6646); they now hold ISV-Pearl-smoothed values populated by
read_isv_signal_at(294/295/296) after the kernel chain completes.

Also added: GpuDqnTrainer::synchronize_isv_stream() — public cold-path
stream-sync helper at gpu_dqn_trainer.rs:~18475. Single cuStreamSynchronize
of the training stream, mirroring the existing per_branch_q_gap_ema()
embedded sync. Used once per epoch by D4's read-back path.

Out of scope (deferred to D5):
  - HealthEmaTrackers::update host-side α=0.1 EMA (metrics.rs:23-25):
    produces D2's inputs; D2's launcher takes already-EMA'd scalars.
    Eliminating this needs a 4th kernel; out of D4 scope.
  - LearningHealth::update warmup wrapper + [0.2, 0.95] clamp
    (learning_health.rs:114-124): may have non-DQN consumers (PPO, eval);
    deferred to a separate refactor.
  - compute_epoch_financials per-bar equity walk: stays as host-side
    input source for D1; D1 reproduces it on GPU as a PARALLEL producer
    publishing to ISV.

Verification:
  - cargo check + build clean
  - sp5_isv_slots + state_reset_registry unit tests 8/8 pass
  - cargo test -p ml --lib: 933 pass / 14 fail (identical to pre-D4
    baseline; failing tests are pre-existing GPU-context-required tests,
    not regressions)
  - `git grep "= EMA_BETA *|= (1.0 - EMA_BETA)"` returns 0 (was already
    0 — gate is a no-op for SP5 because production uses literal
    coefficients, not a named EMA_BETA constant). Meaningful gate
    `git grep "0.9 \* self.max_dd_ema|0.85 \* self.low_dd_ratio"` also
    returns 0 — confirms host EMA arithmetic deletion.
  - Order-of-ops: launches chained on training stream before consumer
    reads; stream-event sync only (synchronize_isv_stream once per epoch
    for cold-path host scalar refresh; producer→consumer paths on the
    same stream observe FIFO order without explicit fence).

Mapped-pinned discipline (per feedback_no_htod_htoh_only_mapped_pinned):
D1's step_returns + done_flags consumed via fresh MappedF32Buffer
allocations populated by host_slice_mut — no HtoD copy. Once-per-epoch
cold-path; matches existing pattern at training_loop.rs:1318.

L40S smoke validation deferred to next 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) to validate separately. Closes feedback_no_cpu_compute_strict
sweep for SP5 scope (modulo D5 follow-up).

Refs: SP5 plan §D Task D4. Builds on D1-rewrite (66f7e64d1), D2
(e49756ac9), D3 (f42b5fff8). First D4 attempt was blocked by D1
per-trade vs per-bar mismatch; resolved by D1-rewrite.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-02 17:45:21 +02:00
parent 66f7e64d13
commit 2e9e276a0c
3 changed files with 305 additions and 38 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -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()); }

231
crates/ml/src/trainers/dqn/trainer/training_loop.rs
View File

@@ -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

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

@@ -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).