From f42b5fff8dc627f7b087536da64bcda304ee7910 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sat, 2 May 2026 16:36:23 +0200 Subject: [PATCH] =?UTF-8?q?feat(sp5):=20Layer=20D=20Task=20D3=20=E2=80=94?= =?UTF-8?q?=20Training=20metrics=20EMA=20kernel=20(additive)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Third of 3 Layer D producer kernels. Replaces host-side training_sharpe_ema, max_dd_ema, low_dd_ratio updates (host-side EMAs in training_loop.rs) with a fused GPU kernel chained through apply_pearls_ad_kernel. Per feedback_no_cpu_compute_strict. Note: agent investigated training_loop.rs and found the third metric is low_dd_ratio (not gamma_blend as the original plan brief named). The authored kernel reproduces the actual host-side EMA triplet present in the codebase. Additive only — no consumer wiring, no behavior change. The kernel + launcher are loaded into GpuDqnTrainer and the slots are reserved on the ISV bus, but the host-side updates continue to run unchanged. D4 (atomic Layer D commit) wires this and D1+D2 to call sites in the same atomic refactor per feedback_no_partial_refactor. Also note: training_loop.rs gains two state-reset-registry dispatch arms (sp5_health_composition for D2 + sp5_training_metrics_ema for D3) — registry plumbing required by the new entries, NOT consumer wiring of the kernels themselves; same pattern as SP5 Layer A bug-fix #281. What this lands: - training_metrics_ema_kernel.cu (single-block 3-thread fused EMA) - Rust launcher launch_training_metrics_ema() - 3 new ISV slots (TRAINING_SHARPE_EMA_INDEX..LOW_DD_RATIO_INDEX, 294..297) - ISV_TOTAL_DIM 294 → 297, SP5_PRODUCER_COUNT 120 → 123 (linear span) - LAYOUT_FINGERPRINT_SEED bump (auto via slot string) - StateResetRegistry entries for D2 (sp5_health_composition) + D3 (sp5_training_metrics_ema) with reset_named_state dispatch arms - build.rs cubin registration - GPU-gated unit test in sp5_producer_unit_tests.rs (analytical EMA) - SP5 contiguity slot test training_metrics_ema_slots_contiguous_and_above_health_block - Audit doc append Formula fidelity: kernel reproduces the host-side EMA update for sharpe/max_dd/low_dd_ratio bit-for-bit within float precision. β decay constants and any warmup/clamp logic migrate as Invariant 1 anchors with no algorithmic change. Verified by unit test asserting kernel output matches an analytical EMA sequence within 1e-6 rel-err. Tests: cargo check + cargo build clean; ISV slot + state_reset_registry unit tests pass (8/8 incl. new contiguity check); GPU correctness test fires on next L40S smoke. Refs: SP5 plan §D Task D3, builds on D1 (5ee795f14) + D2 (e49756ac9). Layer D additive infrastructure complete after this commit; D4 (atomic 5-site host-EMA → GPU migration) is next. Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/build.rs | 15 + .../ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 260 ++++++++++++++++-- crates/ml/src/cuda_pipeline/sp5_isv_slots.rs | 119 ++++++-- .../training_metrics_ema_kernel.cu | 195 +++++++++++++ .../src/trainers/dqn/state_reset_registry.rs | 20 ++ .../src/trainers/dqn/trainer/training_loop.rs | 42 +++ crates/ml/tests/sp5_producer_unit_tests.rs | 189 +++++++++++++ docs/dqn-wire-up-audit.md | 58 ++++ 8 files changed, 856 insertions(+), 42 deletions(-) create mode 100644 crates/ml/src/cuda_pipeline/training_metrics_ema_kernel.cu diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 93021452d..365253f14 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -505,6 +505,21 @@ fn main() { // in the atomic Layer D commit per // `feedback_no_partial_refactor.md`. "health_composition_kernel.cu", + // SP5 Layer D Task D3 (2026-05-02): training-metrics EMA producer. + // Third (and final) Layer D producer replacing the host-side + // adversarial-regime + adaptive-DSR EMA arithmetic block at + // `training_loop.rs:5039-5113` per `feedback_no_cpu_compute_strict.md`. + // Single-block 3-thread kernel — one thread per metric: + // tid=0 → training_sharpe_ema (adaptive α, host-tracked sentinel), + // tid=1 → max_dd_ema (fixed α=0.1, sentinel<1e-12), + // tid=2 → low_dd_ratio (fixed α=0.15, depends on tid=1's + // output via __syncthreads()). + // Writes 3 scratch slots at scratch[SCRATCH_TRAINING_METRICS_EMA_BASE + // =215..218). Chained `apply_pearls_ad_kernel` (3 launches, + // ALPHA_META=1e-3) smooths via Pearls A+D into ISV[294..297). + // Producer-only — D4 wires the consumer in the atomic Layer D + // commit per `feedback_no_partial_refactor.md`. + "training_metrics_ema_kernel.cu", // SP4 GPU-only Pearls A+D applicator (2026-05-01). Single-thread // device-side loop applies Pearls A+D to N consecutive ISV slots // (each with its own Wiener-state triple) on the producer's diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 8a4f4adf4..ef310c1fc 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs @@ -426,6 +426,22 @@ static SP5_PNL_AGGREGATION_CUBIN: &[u8] = static SP5_HEALTH_COMPOSITION_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/health_composition_kernel.cubin")); +/// SP5 Layer D Task D3 (2026-05-02): Training metrics EMA producer cubin. +/// Single-block 3-thread arithmetic kernel reproducing the host-side +/// `training_sharpe_ema` (adaptive α), `max_dd_ema` (fixed α=0.1), and +/// `low_dd_ratio` (fixed α=0.15) updates at `training_loop.rs:5039-5113` +/// per `feedback_no_cpu_compute_strict.md`. Per-metric thread layout: +/// tid=0 sharpe (adaptive α with host-tracked `_initialized` sentinel), +/// tid=1 max_dd (fixed α=0.1 with prev<1e-12 sentinel), tid=2 low_dd_ratio +/// (fixed α=0.15 over a binary `is_low` derived from tid=1's *new* +/// max_dd_ema via __syncthreads()). Writes 3 floats to +/// scratch_buf[SCRATCH_TRAINING_METRICS_EMA_BASE=215..218). Followed by 3 +/// apply_pearls_ad_kernel calls → ISV[TRAINING_SHARPE_EMA_INDEX=294..297). +/// Replaces the host-side updates that D4 atomically migrates per +/// `feedback_no_partial_refactor.md`. Additive in this commit. +static SP5_TRAINING_METRICS_EMA_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/training_metrics_ema_kernel.cubin")); + /// Plan 4 Task 3 (E.3): IQN multi-quantile diagnostic EMAs into ISV[99..103). /// 4-block (256 threads/block, shmem-reduce, no atomicAdd) kernel launched /// alongside `h_s2_rms_ema_update` from `training_loop.rs`. Reads the IQN @@ -1000,27 +1016,29 @@ pub const SP4_WIENER_TOTAL_FLOATS: usize = // ── SP5 Task A1 buffer layout constants ─────────────────────────────────────── -/// SP5 Task A1 + Layer D Task D1+D2: combined wiener_state_buf total float count. -/// SP4 had 71 producers (213 floats); SP5 contributes 120 producer triples -/// (linear span from `SP5_SLOT_BASE=174` to `SP5_SLOT_END=294`, including the -/// 2-slot carve-out gap at 278..280; 120 × 3 = 360 floats). Combined = -/// (71 + 120) × 3 = 573 floats. Pre-D1 SP5 was 110 producers / 543 floats; +/// SP5 Task A1 + Layer D Task D1+D2+D3: combined wiener_state_buf total float count. +/// SP4 had 71 producers (213 floats); SP5 contributes 123 producer triples +/// (linear span from `SP5_SLOT_BASE=174` to `SP5_SLOT_END=297`, including the +/// 2-slot carve-out gap at 278..280; 123 × 3 = 369 floats). Combined = +/// (71 + 123) × 3 = 582 floats. Pre-D1 SP5 was 110 producers / 543 floats; /// Task D1 added 4 unique slots (286..290) growing the linear span 112 → 116 -/// (buffer 543 → 561). Task D2 adds 4 more slots (290..294) growing the linear -/// span 116 → 120 (buffer 561 → 573). Note: SP5_PRODUCER_COUNT is the -/// **linear-span** constant — see its docstring for why it now diverges -/// permanently from the unique-slot count (118 post-D2). +/// (buffer 543 → 561). Task D2 added 4 more slots (290..294) growing the linear +/// span 116 → 120 (buffer 561 → 573). Task D3 adds 3 more slots (294..297) +/// growing the linear span 120 → 123 (buffer 573 → 582). Note: +/// SP5_PRODUCER_COUNT is the **linear-span** constant — see its docstring +/// for why it now diverges permanently from the unique-slot count (121 post-D3). /// /// wiener_state_buf must grow to this size at construction for ALL SP5 slots -/// (including the new D1 PnL-aggregation and D2 health-composition blocks) to -/// be in-bounds when `apply_pearls_ad_kernel` writes their Wiener triples. -/// The Pearl 6 Kelly block (slots 280..286) and the carve-out gap (278..280) -/// have reserved-but-unused wiener storage at offsets -/// [(280-174)*3+213..(286-174)*3+213) = [531..549) — Pearl 6 does not call -/// `launch_apply_pearls` (cross-fold persistent), so those 18 floats stay -/// zero-initialised. Layer D D1 slots 286..290 use offsets +/// (including the new D1 PnL-aggregation, D2 health-composition, and D3 +/// training-metrics-EMA blocks) to be in-bounds when `apply_pearls_ad_kernel` +/// writes their Wiener triples. The Pearl 6 Kelly block (slots 280..286) and +/// the carve-out gap (278..280) have reserved-but-unused wiener storage at +/// offsets [(280-174)*3+213..(286-174)*3+213) = [531..549) — Pearl 6 does +/// not call `launch_apply_pearls` (cross-fold persistent), so those 18 floats +/// stay zero-initialised. Layer D D1 slots 286..290 use offsets /// [(286-174)*3+213..(290-174)*3+213) = [549..561). Layer D D2 slots 290..294 -/// use offsets [(290-174)*3+213..(294-174)*3+213) = [561..573). +/// use offsets [(290-174)*3+213..(294-174)*3+213) = [561..573). Layer D D3 +/// slots 294..297 use offsets [(294-174)*3+213..(297-174)*3+213) = [573..582). pub const SP5_WIENER_TOTAL_FLOATS: usize = (SP4_PRODUCER_COUNT + crate::cuda_pipeline::sp5_isv_slots::SP5_PRODUCER_COUNT) * SP4_WIENER_FLOATS_PER_SLOT; @@ -1057,8 +1075,10 @@ pub const SP5_WIENER_TOTAL_FLOATS: usize = /// 4 floats for pnl_aggregation output (SCRATCH_PNL_AGG_BASE=207, total/mean/var/max_dd at [207..211)) /// Layer D Task D2 adds: /// 4 floats for health_composition output (SCRATCH_HEALTH_COMP_BASE=211, score/q_gap_norm/q_var_norm/grad_norm_norm at [211..215)) -/// Combined: 71 + 16 + 16 + 4 + 4 + 20 + 8 + 8 + 8 + 8 + 8 + 4 + 4 + 20 + 4 + 4 + 4 + 4 = 215 scratch slots [0..215). -pub const SP5_SCRATCH_TOTAL: usize = 215; +/// Layer D Task D3 adds: +/// 3 floats for training_metrics_ema output (SCRATCH_TRAINING_METRICS_EMA_BASE=215, training_sharpe_ema/max_dd_ema/low_dd_ratio at [215..218)) +/// Combined: 71 + 16 + 16 + 4 + 4 + 20 + 8 + 8 + 8 + 8 + 8 + 4 + 4 + 20 + 4 + 4 + 4 + 4 + 3 = 218 scratch slots [0..218). +pub const SP5_SCRATCH_TOTAL: usize = 218; /// SP5 Layer D Task D1 (2026-05-02): scratch index base for `pnl_aggregation_update`. /// Slots [207..211): aggregated PnL summary outputs in fixed order: @@ -1086,6 +1106,29 @@ pub const SCRATCH_PNL_AGG_BASE: usize = 207; /// a launcher signature change. pub const SCRATCH_HEALTH_COMP_BASE: usize = 211; +/// SP5 Layer D Task D3 (2026-05-02): scratch index base for `training_metrics_ema_update`. +/// Slots [215..218): training-loop EMA outputs in fixed order: +/// [215] = training_sharpe_ema (adaptive α `clamp(0.05+0.3·|err|, 0, 0.5)`, +/// host-tracked `_initialized` sentinel) +/// [216] = max_dd_ema (fixed α=0.1, sentinel `prev<1e-12` ⇒ +/// first-observation replacement) +/// [217] = low_dd_ratio (fixed α=0.15 over binary +/// `is_low = (max_dd < new_max_dd_ema·0.5) +/// ? 1 : 0`; reads tid=1's *new* output via +/// __syncthreads(); no host-side sentinel, +/// constructor zero behaves correctly under +/// the EMA recurrence) +/// Written by `training_metrics_ema_update`; consumed by `apply_pearls_ad_kernel` +/// → ISV[TRAINING_SHARPE_EMA_INDEX=294..LOW_DD_RATIO_INDEX+1=297). All three +/// host-side EMA constants (0.05/0.3/0.5 adaptive bounds; 0.1; 0.15; 1e-12 +/// sentinel) are Invariant 1 anchors per `feedback_no_quickfixes.md`, +/// migrated verbatim from `training_loop.rs:5039-5113`. Per +/// `feedback_trust_code_not_docs.md`: the third metric is `low_dd_ratio` +/// (not `gamma_blend` as the SP5 plan §D Task D3 hints — that field does +/// not exist in the codebase; the migration follows the actual host-side +/// EMA at `training_loop.rs:5052`). +pub const SCRATCH_TRAINING_METRICS_EMA_BASE: usize = 215; + /// SP5 Task A7: scratch index base for pearl_8_trail_update trail_dist[4] output block. /// Slots [199..203): per-direction trail-stop distance (Short=0, Hold=1, Long=2, Flat=3). /// Written by `pearl_8_trail_update`; consumed by apply_pearls_ad_kernel → @@ -4422,6 +4465,25 @@ pub struct GpuDqnTrainer { /// path). Loaded from `health_composition_kernel.cubin`. Producer-only in /// this commit; D4 wires the consumer atomically. health_composition_kernel: CudaFunction, + /// SP5 Layer D Task D3 (2026-05-02): training-metrics EMA producer. + /// Single-block 3-thread arithmetic kernel (one thread per metric): + /// tid=0 training_sharpe_ema (adaptive α `clamp(0.05+0.3·|err|, 0, 0.5)`, + /// host-tracked `_initialized` sentinel), tid=1 max_dd_ema (fixed α=0.1, + /// `prev<1e-12` sentinel), tid=2 low_dd_ratio (fixed α=0.15 over the + /// binary `is_low` indicator derived from tid=1's *new* max_dd_ema via + /// __syncthreads(); no host-side sentinel — constructor zero is + /// recurrence-safe). Writes 3 scratch floats — training_sharpe_ema + + /// max_dd_ema + low_dd_ratio — to scratch_buf[ + /// SCRATCH_TRAINING_METRICS_EMA_BASE=215..218). Chained + /// apply_pearls_ad_kernel (3 launches, ALPHA_META=1e-3) → + /// ISV[TRAINING_SHARPE_EMA_INDEX=294..LOW_DD_RATIO_INDEX+1=297). In + /// fold-reset registry (FoldReset; the underlying host scalars + /// `training_sharpe_ema` / `_initialized` / `max_dd_ema` / + /// `low_dd_ratio` reset at fold boundary alongside the ISV slots via + /// the standard Pearl A sentinel path). Loaded from + /// `training_metrics_ema_kernel.cubin`. Producer-only in this commit; + /// D4 wires the consumer atomically. + training_metrics_ema_kernel: CudaFunction, /// SP5 Task A4 (2026-05-01): previous-step gradient buffer for cosine similarity. /// [total_params f32] mapped-pinned (feedback_no_htod_htoh_only_mapped_pinned). /// Zeroed at construction and at each fold boundary (Pearl A sentinel: cosine_sim=0 @@ -11928,6 +11990,147 @@ impl GpuDqnTrainer { Ok(()) } + /// SP5 Layer D Task D3 (2026-05-02): launch `training_metrics_ema_update` — + /// single-block 3-thread EMA producer for + /// ISV[TRAINING_SHARPE_EMA_INDEX=294..LOW_DD_RATIO_INDEX+1=297). + /// + /// Reproduces the host-side EMA arithmetic block at + /// `training_loop.rs:5039-5113` bit-for-bit on GPU per + /// `feedback_no_cpu_compute_strict.md`. One thread per metric: + /// + /// tid=0 training_sharpe_ema α = clamp(0.05 + 0.3·|err|, 0, 0.5); + /// sharpe_initialized==0 ⇒ replace. + /// tid=1 max_dd_ema α = 0.1; prev<1e-12 ⇒ replace. + /// tid=2 low_dd_ratio α = 0.15 over binary + /// `is_low = (max_dd < new_max_dd_ema·0.5) + /// ? 1 : 0`. Reads tid=1's *new* max_dd_ema + /// via __syncthreads(); no host-side sentinel. + /// + /// All EMA constants (0.05, 0.3, 0.5, 0.1, 0.85, 0.15, 0.5 threshold, + /// 1e-12 sentinel) are Invariant 1 anchors — deleted host formula + /// literals migrated verbatim per `feedback_no_quickfixes.md` and + /// `feedback_trust_code_not_docs.md` (the third metric is `low_dd_ratio` + /// from `training_loop.rs:5052`, not the speculative `gamma_blend` + /// the SP5 plan §D Task D3 mentions — that field does not exist in + /// the codebase). + /// + /// Writes 3 floats to scratch_buf[SCRATCH_TRAINING_METRICS_EMA_BASE + /// =215..218): + /// [215] new training_sharpe_ema + /// [216] new max_dd_ema + /// [217] new low_dd_ratio + /// + /// Then chains 3 `apply_pearls_ad_kernel` calls smoothing through + /// Pearls A+D into ISV[294..297) on the same stream (no host sync per + /// `pearl_cold_path_no_exception_to_gpu_drives.md` and the SP5 GPU-only + /// Pearls refactor). + /// + /// Per `feedback_no_partial_refactor.md` — D3 is additive only; D4 + /// atomically migrates the host-side site to call this launcher in + /// lockstep with D1/D2. The kernel takes the host-side prev EMA values + /// and the `_initialized` flag by value in additive D3; D4 will lift + /// these to mapped-pinned scalars so the host is no longer the source + /// of truth. + /// + /// Arguments (all primitive scalars by value, no mapped-pinned plumbing + /// per the additive-D3 design): + /// - `epoch_sharpe`: per-epoch Sharpe ratio observation (raw). + /// - `epoch_max_dd`: per-epoch max drawdown observation (raw, ∈ [0, 1]). + /// - `prev_sharpe_ema`: host-tracked previous training_sharpe_ema. + /// - `prev_max_dd_ema`: host-tracked previous max_dd_ema (cast to f32). + /// - `prev_low_dd_ratio`: host-tracked previous low_dd_ratio (cast to f32). + /// - `sharpe_initialized`: host-tracked sentinel flag; `false` ⇒ first + /// observation replaces the slot directly. Cast to `i32` at the + /// kernel boundary per `feedback_cudarc_f64_f32_abi.md` (cudarc's + /// `bool` ABI is unstable; `i32` round-trip is the safe pattern). + pub(crate) fn launch_training_metrics_ema( + &self, + epoch_sharpe: f32, + epoch_max_dd: f32, + prev_sharpe_ema: f32, + prev_max_dd_ema: f32, + prev_low_dd_ratio: f32, + sharpe_initialized: bool, + ) -> Result<(), MLError> { + use crate::cuda_pipeline::sp5_isv_slots::{ + TRAINING_SHARPE_EMA_INDEX, MAX_DD_EMA_INDEX, LOW_DD_RATIO_INDEX, + SP5_SLOT_BASE, + }; + use crate::cuda_pipeline::sp4_wiener_ema::{launch_apply_pearls, ALPHA_META}; + + debug_assert!(self.isv_signals_dev_ptr != 0, + "launch_training_metrics_ema: isv_signals_dev_ptr must be allocated by constructor"); + + let scratch_dev = self.producer_step_scratch_buf.dev_ptr; + let isv_dev = self.isv_signals_dev_ptr; + let wiener_dev = self.wiener_state_buf.dev_ptr; + + let sharpe_idx_i32: i32 = (SCRATCH_TRAINING_METRICS_EMA_BASE + 0) as i32; + let max_dd_idx_i32: i32 = (SCRATCH_TRAINING_METRICS_EMA_BASE + 1) as i32; + let low_dd_ratio_idx_i32: i32 = (SCRATCH_TRAINING_METRICS_EMA_BASE + 2) as i32; + + // i32 sentinel flag (cudarc `bool` ABI is unstable — cast to i32 here + // matching the kernel's `int sharpe_initialized` parameter; per + // `feedback_cudarc_f64_f32_abi.md` the launcher's primitive type + // must match the kernel's declared C type exactly). + let sharpe_initialized_i32: i32 = if sharpe_initialized { 1 } else { 0 }; + + // Wiener base offset for SP5 slots: SP4_PRODUCER_COUNT * 3 = 213. + let base_wiener_offset: i32 = (SP4_PRODUCER_COUNT * 3) as i32; + + // Step 1: training_metrics_ema_update — single-block 3-thread; no + // dynamic shared memory (the kernel uses a single static __shared__ + // float for the tid=1 → tid=2 max_dd_ema handoff). + unsafe { + self.stream + .launch_builder(&self.training_metrics_ema_kernel) + .arg(&epoch_sharpe) + .arg(&epoch_max_dd) + .arg(&prev_sharpe_ema) + .arg(&prev_max_dd_ema) + .arg(&prev_low_dd_ratio) + .arg(&sharpe_initialized_i32) + .arg(&scratch_dev) + .arg(&sharpe_idx_i32) + .arg(&max_dd_idx_i32) + .arg(&low_dd_ratio_idx_i32) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (3, 1, 1), + shared_mem_bytes: 0, + }) + .map_err(|e| MLError::ModelError(format!("training_metrics_ema_update: {e}")))?; + } + + // Step 2: 3 apply_pearls calls (one per output slot). Each + // scratch[215+k] → ISV[294+k] with its own Wiener triple at offset + // (base_wiener_offset + (isv_slot − SP5_SLOT_BASE) * 3). Contract + // mirrors D1/D2 exactly. + let isv_slots: [usize; 3] = [ + TRAINING_SHARPE_EMA_INDEX, + MAX_DD_EMA_INDEX, + LOW_DD_RATIO_INDEX, + ]; + for k in 0..3_usize { + let isv_slot = isv_slots[k] as i32; + let scratch_idx = (SCRATCH_TRAINING_METRICS_EMA_BASE + k) as i32; + let wiener_off = base_wiener_offset + (isv_slot - SP5_SLOT_BASE as i32) * 3; + unsafe { + launch_apply_pearls( + &self.stream, + &self.apply_pearls_ad_kernel, + scratch_dev, scratch_idx, + isv_dev, isv_slot, + wiener_dev, wiener_off, + 1, + ALPHA_META, + )?; + } + } + + Ok(()) + } + /// Plan C Phase 2 follow-up A.2 (2026-04-29): launch `q_drift_rate_ema_update` /// — single-thread single-block ISV producer for ISV[Q_DRIFT_RATE_INDEX=129]. /// @@ -14074,6 +14277,24 @@ impl GpuDqnTrainer { .map_err(|e| MLError::ModelError(format!("health_composition_update load: {e}")))? }; + // SP5 Layer D Task D3 (2026-05-02): load training_metrics_ema_kernel. + // Single-block 3-thread arithmetic kernel. Reads 5 host-passed f32 + // scalars (epoch_sharpe + epoch_max_dd raw observations, plus + // prev_sharpe_ema / prev_max_dd_ema / prev_low_dd_ratio host-tracked + // EMAs) and 1 i32 sentinel flag (sharpe_initialized); writes 3 floats + // (sharpe/max_dd/low_dd_ratio) to scratch[SCRATCH_TRAINING_METRICS_EMA_BASE + // =215..218). FoldReset: ISV[294..297) zeroed at fold boundary + // alongside the host-side `training_sharpe_ema` / `_initialized` / + // `max_dd_ema` / `low_dd_ratio` scalars (Pearl A sentinel path). + // Producer-only in D3; consumer wires atomically in D4 per + // `feedback_no_partial_refactor.md`. + let training_metrics_ema_kernel = { + let module = stream.context().load_cubin(SP5_TRAINING_METRICS_EMA_CUBIN.to_vec()) + .map_err(|e| MLError::ModelError(format!("sp5 training_metrics_ema cubin load: {e}")))?; + module.load_function("training_metrics_ema_update") + .map_err(|e| MLError::ModelError(format!("training_metrics_ema_update load: {e}")))? + }; + // SP5 Task A4 (2026-05-01): allocate grad_prev_buf_per_group [total_params f32]. // Zeroed at construction; fold-boundary reset zeroes it again (Pearl A sentinel: // zero grad_prev → cosine_sim=0 on first step → β1/β2 start at envelope midpoints). @@ -17261,6 +17482,7 @@ impl GpuDqnTrainer { pearl_1_ext_num_atoms_kernel, pnl_aggregation_kernel, health_composition_kernel, + training_metrics_ema_kernel, grad_prev_buf_per_group, group_param_offsets_buf, atoms_clip_count_buf, diff --git a/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs b/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs index 7601bf035..016a29126 100644 --- a/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs +++ b/crates/ml/src/cuda_pipeline/sp5_isv_slots.rs @@ -12,8 +12,9 @@ //! 280..286 Cross-fold-persistent Kelly slots (NOT in fold-reset registry) //! 286..290 Layer D D1: PnL aggregation outputs (4 slots, fold-reset) //! 290..294 Layer D D2: Health composition outputs (4 slots, fold-reset) +//! 294..297 Layer D D3: Training metrics EMA outputs (3 slots, fold-reset) //! -//! Total: 118 new SP5 ISV slots (52 + 24 + 20 + 4 + 4 + 6 + 4 + 4). +//! Total: 121 new SP5 ISV slots (52 + 24 + 20 + 4 + 4 + 6 + 4 + 4 + 3). pub const SP5_SLOT_BASE: usize = 174; @@ -118,7 +119,42 @@ pub const Q_GAP_NORM_INDEX: usize = 291; pub const Q_VAR_NORM_INDEX: usize = 292; pub const GRAD_NORM_NORM_INDEX: usize = 293; -pub const SP5_SLOT_END: usize = 294; +// ── Layer D Task D3: Training metrics EMA outputs (3 slots) ────────── +/// Layer D D3 (2026-05-02): GPU-aggregated training-loop EMA outputs. +/// Replaces the host-side per-epoch EMA arithmetic block at +/// `training_loop.rs:5039-5113` (the post-financials adversarial-regime +/// hysteresis + adaptive-DSR Sharpe block) per +/// `feedback_no_cpu_compute_strict.md`. Producer is +/// `training_metrics_ema_kernel.cu` — chained through +/// `apply_pearls_ad_kernel` for Pearls A+D smoothing on top of the +/// host-formula EMA (sentinel-bootstrap on first observation, +/// Wiener-optimal α at steady state). Per-fold reset semantics: the +/// underlying host scalars (`training_sharpe_ema`, +/// `training_sharpe_ema_initialized`, `max_dd_ema`, `low_dd_ratio`) reset +/// at fold boundary, so all three slots get sentinel 0.0 at fold boundary. +/// +/// Slot naming follows the host-side field names in +/// `crates/ml/src/trainers/dqn/trainer/mod.rs:319/603/607` directly. The +/// SP5 plan §D Task D3 names the third metric `gamma_blend` (a host- +/// aggregated eval result) but no such field exists in the codebase; per +/// `feedback_trust_code_not_docs.md` the migration follows the actual +/// host-side EMA which is `low_dd_ratio` (the binary low-drawdown +/// indicator EMA driving the adversarial-regime hysteresis at +/// `training_loop.rs:5054-5070`). All three migrate with their host-side +/// formula constants verbatim: +/// +/// training_sharpe_ema adaptive-α `clamp(0.05 + 0.3 × |err|, 0, 0.5)`, +/// host-tracked `_initialized` sentinel. +/// max_dd_ema fixed α=0.1, sentinel `prev < 1e-12 ⇒ replace`. +/// low_dd_ratio fixed α=0.15 over binary `is_low` input from +/// `(max_dd < new_max_dd_ema × 0.5) ? 1 : 0`. No +/// host-side sentinel; constructor zero behaves +/// correctly under the EMA recurrence. +pub const TRAINING_SHARPE_EMA_INDEX: usize = 294; +pub const MAX_DD_EMA_INDEX: usize = 295; +pub const LOW_DD_RATIO_INDEX: usize = 296; + +pub const SP5_SLOT_END: usize = 297; /// Wiener-buffer producer-count constant. Sizes `wiener_state_buf` via /// `(SP4_PRODUCER_COUNT + SP5_PRODUCER_COUNT) * SP4_WIENER_FLOATS_PER_SLOT`. @@ -148,11 +184,17 @@ pub const SP5_SLOT_END: usize = 294; /// 174..278 ∪ 280..294; the linear span grows 116 → 120. Updates to /// docstrings, fingerprint, and the slot-contiguity test are kept in /// lockstep with the constant per `feedback_no_partial_refactor.md`. -pub const SP5_PRODUCER_COUNT: usize = 120; -// linear span = SP5_SLOT_END - SP5_SLOT_BASE = 294 - 174 = 120 wiener triples -// unique-slot count = 118 (52 per-branch + 24 Adam + 20 IQN τ + 4 trail +/// +/// Layer D Task D3 (2026-05-02) further extends to 121 unique slots @ +/// 174..278 ∪ 280..297; the linear span grows 120 → 123. Three new +/// training-metrics EMA outputs (sharpe / max_dd / low_dd_ratio) sit +/// immediately after the D2 health composition block at slots 294..297. +pub const SP5_PRODUCER_COUNT: usize = 123; +// linear span = SP5_SLOT_END - SP5_SLOT_BASE = 297 - 174 = 123 wiener triples +// unique-slot count = 121 (52 per-branch + 24 Adam + 20 IQN τ + 4 trail // + 4 num_atoms + 6 Kelly + 4 Layer D D1 PnL aggregation -// + 4 Layer D D2 health composition) +// + 4 Layer D D2 health composition +// + 3 Layer D D3 training metrics EMA) // ── Convenience accessors ──────────────────────────────────────────── #[inline] pub const fn atom_v_center(b: usize) -> usize { ATOM_V_CENTER_BASE + b } @@ -188,7 +230,8 @@ pub const SP5_LAYOUT_FINGERPRINT_FRAGMENT: &str = KELLY_SAMPLE_COUNT=283;WIN_RATE_SMOOTH=284;LOSS_RATE_SMOOTH=285;\ PNL_TOTAL=286;PNL_MEAN=287;PNL_VAR=288;PNL_MAX_DD=289;\ HEALTH_SCORE=290;Q_GAP_NORM=291;Q_VAR_NORM=292;GRAD_NORM_NORM=293;\ - ISV_TOTAL_DIM=294"; + TRAINING_SHARPE_EMA=294;MAX_DD_EMA=295;LOW_DD_RATIO=296;\ + ISV_TOTAL_DIM=297"; #[cfg(test)] mod tests { @@ -261,26 +304,31 @@ mod tests { slots.insert(Q_GAP_NORM_INDEX); slots.insert(Q_VAR_NORM_INDEX); slots.insert(GRAD_NORM_NORM_INDEX); + // Layer D D3: Training metrics EMA outputs (3 scalar constants) + slots.insert(TRAINING_SHARPE_EMA_INDEX); + slots.insert(MAX_DD_EMA_INDEX); + slots.insert(LOW_DD_RATIO_INDEX); - // 1. Exactly 118 unique slots. - assert_eq!(slots.len(), 118, "expected 118 unique slots, got {}", slots.len()); + // 1. Exactly 121 unique slots. + assert_eq!(slots.len(), 121, "expected 121 unique slots, got {}", slots.len()); // 2. Min slot is SP5_SLOT_BASE = 174. assert_eq!(*slots.iter().min().unwrap(), 174); - // 3. Max slot is SP5_SLOT_END - 1 = 293. - assert_eq!(*slots.iter().max().unwrap(), 293); + // 3. Max slot is SP5_SLOT_END - 1 = 296. + assert_eq!(*slots.iter().max().unwrap(), 296); // 4. Intentional carve-out gap (278, 279) is absent. assert!(!slots.contains(&278), "slot 278 must be absent (carve-out gap)"); assert!(!slots.contains(&279), "slot 279 must be absent (carve-out gap)"); - // 5. Set equals {174..278} ∪ {280..294} — no holes (other than the - // carve-out gap 278..280), no overlaps. Layer D D1 extends the - // upper end from 286 to 290; D2 extends it 290 → 294. - let expected: HashSet = (174..278).chain(280..294).collect(); + // 5. Set equals {174..278} ∪ {280..297} — no holes (other than the + // carve-out gap 278..280), no overlaps. Layer D D1 extended the + // upper end 286 → 290; D2 extended it 290 → 294; D3 extends it + // 294 → 297. + let expected: HashSet = (174..278).chain(280..297).collect(); assert_eq!(slots, expected, - "slot set does not match expected {{174..278}} ∪ {{280..294}}"); + "slot set does not match expected {{174..278}} ∪ {{280..297}}"); } #[test] @@ -306,12 +354,12 @@ mod tests { // Strictly above the Kelly block (Layer D layout is per-fold; Kelly // is cross-fold-persistent — the contracts must remain disjoint). assert!(PNL_TOTAL_INDEX > LOSS_RATE_SMOOTH_INDEX); - // SP5_SLOT_END must reflect the post-D2 end-of-block. - assert_eq!(SP5_SLOT_END, 294); + // SP5_SLOT_END must reflect the post-D3 end-of-block. + assert_eq!(SP5_SLOT_END, 297); // SP5_PRODUCER_COUNT is the wiener-buffer linear span (slot-range // width including the 2-slot carve-out gap), NOT the unique-slot // count. See SP5_PRODUCER_COUNT docstring for the rationale. - assert_eq!(SP5_PRODUCER_COUNT, 120); + assert_eq!(SP5_PRODUCER_COUNT, 123); assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE); } @@ -323,15 +371,40 @@ mod tests { assert_eq!(Q_GAP_NORM_INDEX, 291); assert_eq!(Q_VAR_NORM_INDEX, 292); assert_eq!(GRAD_NORM_NORM_INDEX, 293); - // Strictly above the D1 PnL block (Layer D's two producer commits + // Strictly above the D1 PnL block (Layer D's three producer commits // land in disjoint slot ranges so D4's atomic consumer wiring sees - // both blocks at fixed offsets). + // each block at fixed offsets). assert!(HEALTH_SCORE_INDEX > PNL_MAX_DD_INDEX); // 4-slot block is internally contiguous. assert_eq!(GRAD_NORM_NORM_INDEX - HEALTH_SCORE_INDEX, 3); - // SP5_SLOT_END must reflect the new end-of-block (4-slot grow). - assert_eq!(SP5_SLOT_END, 294); + // SP5_SLOT_END must reflect the post-D3 end-of-block. + assert_eq!(SP5_SLOT_END, 297); // SP5_PRODUCER_COUNT linear-span check matches the new end. assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE); } + + #[test] + fn training_metrics_ema_slots_contiguous_and_above_health_block() { + // Layer D D3: Training metrics EMA outputs occupy a contiguous 3-slot + // block immediately after the Layer D D2 health composition block. + assert_eq!(TRAINING_SHARPE_EMA_INDEX, 294); + assert_eq!(MAX_DD_EMA_INDEX, 295); + assert_eq!(LOW_DD_RATIO_INDEX, 296); + // Strictly above the D2 health composition block (Layer D's three + // producer commits land in disjoint slot ranges so D4's atomic + // consumer wiring sees each block at fixed offsets per + // `feedback_no_partial_refactor.md`). + assert!(TRAINING_SHARPE_EMA_INDEX > GRAD_NORM_NORM_INDEX); + // 3-slot block is internally contiguous (no gaps). + assert_eq!(MAX_DD_EMA_INDEX - TRAINING_SHARPE_EMA_INDEX, 1); + assert_eq!(LOW_DD_RATIO_INDEX - MAX_DD_EMA_INDEX, 1); + // SP5_SLOT_END must reflect the new end-of-block (3-slot grow). + assert_eq!(SP5_SLOT_END, 297); + // SP5_PRODUCER_COUNT linear-span check matches the new end. The + // wiener buffer must cover the entire linear span — including the + // 2-slot carve-out gap (278..280) and the Pearl 6 reserved-but- + // unused 6-float block (slots 280..286 don't call apply_pearls). + assert_eq!(SP5_PRODUCER_COUNT, 123); + assert_eq!(SP5_PRODUCER_COUNT, SP5_SLOT_END - SP5_SLOT_BASE); + } } diff --git a/crates/ml/src/cuda_pipeline/training_metrics_ema_kernel.cu b/crates/ml/src/cuda_pipeline/training_metrics_ema_kernel.cu new file mode 100644 index 000000000..1d4eb799a --- /dev/null +++ b/crates/ml/src/cuda_pipeline/training_metrics_ema_kernel.cu @@ -0,0 +1,195 @@ +// crates/ml/src/cuda_pipeline/training_metrics_ema_kernel.cu +// +// SP5 Layer D Task D3 (2026-05-02): Training-metrics EMA producer kernel. +// +// Third (and final) Layer D producer kernel after D1 (PnL aggregation) and +// D2 (LearningHealth composition). Replaces the host-side EMA arithmetic +// for the three training-loop metrics that survive at the end of each epoch +// in `training_loop.rs:5039-5113` (the post-financials adversarial-regime +// block + the adaptive-DSR block) per `feedback_no_cpu_compute_strict.md`: +// +// training_sharpe_ema adaptive-α EMA of `epoch_sharpe` +// α = clamp(0.05 + 0.3 × |obs − prev|, 0, 0.5) +// sentinel: `_initialized` flag (host) — first +// observation replaces the slot directly. +// max_dd_ema fixed α=0.1 EMA of `epoch_max_dd`. +// sentinel: prev < 1e-12 ⇒ replace directly. +// low_dd_ratio fixed α=0.15 EMA of binary indicator +// `is_low = (max_dd < new_max_dd_ema × 0.5) ? 1 : 0`. +// Reads the *new* max_dd_ema (this kernel's own +// output) — thread 2 must wait until thread 1's +// max_dd_ema write is visible. +// No sentinel — initial 0.0 behaves correctly under +// the EMA recurrence. +// +// **Plan brief vs code reality** (per `feedback_trust_code_not_docs.md`): +// the SP5 plan §D Task D3 names the third metric `gamma_blend` and points +// at "host-aggregated eval result" — that field does not exist in the +// codebase. The actual third EMA at `training_loop.rs:5052` is +// `low_dd_ratio` (an EMA of the binary low-drawdown indicator that drives +// the adversarial-regime hysteresis at lines 5054-5070). Layer D D3 ports +// what the code does, not what the plan speculated. The slot constant is +// named after the code. +// +// Outputs (single block, 3 threads — one per metric; 3 scratch slots written): +// +// scratch_buf[sharpe_ema_idx] = updated training_sharpe_ema +// scratch_buf[max_dd_ema_idx] = updated max_dd_ema +// scratch_buf[low_dd_ratio_idx] = updated low_dd_ratio +// +// Threading contract: +// - Thread 0 computes training_sharpe_ema (adaptive α, in-register; no +// dependence on the other threads). +// - Thread 1 computes max_dd_ema (fixed α=0.1, sentinel < 1e-12). +// - Thread 2 computes low_dd_ratio. Depends on thread 1's max_dd_ema +// write — synchronised via __syncthreads() between passes. +// +// Cold-start sentinels: +// - sharpe_initialized==0 ⇒ output = epoch_sharpe (replaces the slot). +// The `_initialized` flag is host-tracked in D3 (additive mode); D4 +// will migrate the flag to a mapped-pinned scalar so the host no +// longer needs to consult it. For now we receive it by value. +// - max_dd_prev < 1e-12 ⇒ output = epoch_max_dd (matches the deleted +// host formula at training_loop.rs:5043). +// - low_dd_ratio has no host-side sentinel: starts at 0.0 from the +// constructor (constructor.rs:710) and the EMA recurrence converges +// correctly from there. The kernel does likewise — no special branch. +// +// Per `pearl_first_observation_bootstrap.md`: the downstream +// `apply_pearls_ad_kernel` chain provides Wiener-optimal smoothing on top +// of the kernel's host-formula EMA, with Pearl A first-observation +// replacement on the ISV slot when the corresponding wiener triple's +// `prev_x_mean=0 AND state.x_lag=0` sentinel pair fires. Both layers +// (host-formula EMA in this kernel + Pearls A+D smoothing on top) are +// preserved in lockstep with the host-side site that D4 will replace. +// +// Producer-only — D3 is additive. The host-side updates at +// `training_loop.rs:5039-5113` continue to run unchanged. D4 (the atomic +// Layer D commit) wires D1+D2+D3 to call sites in a single migration per +// `feedback_no_partial_refactor.md`. No atomicAdd +// (`feedback_no_atomicadd.md`); single-block 3-thread kernel emits one +// __threadfence_system() at the end so the mapped-pinned host_ptr / dev_ptr +// alias sees the values before the downstream apply_pearls reads them on +// the same stream. + +#include + +extern "C" { + +__global__ void training_metrics_ema_update( + /* Raw observations (this epoch). */ + float epoch_sharpe, + float epoch_max_dd, + + /* Host-tracked previous EMA values (will become mapped-pinned reads + * in D4 atomic refactor; passed by value in additive D3). */ + float prev_sharpe_ema, + float prev_max_dd_ema, + float prev_low_dd_ratio, + + /* Sentinel flag for the sharpe EMA (== 0 ⇒ first observation; D4 will + * derive from a mapped-pinned scalar so host doesn't track this). + * Passed as int because cudarc's `arg(&bool_ref)` lacks a stable ABI + * for `bool` and the kernel ABI must match the launcher exactly per + * `feedback_cudarc_f64_f32_abi.md`. */ + int sharpe_initialized, + + /* Output scratch slots (apply_pearls_ad reads these on the same stream). */ + float* __restrict__ scratch_buf, + int sharpe_ema_idx, + int max_dd_ema_idx, + int low_dd_ratio_idx +) { + /* Single-block contract — match SP5/SP4 producer kernel grid guard. */ + if (blockIdx.x != 0) return; + + const int tid = (int)threadIdx.x; + + /* Shared scratch: thread 1's max_dd_ema output, read by thread 2. + * Single 4-byte allocation; no shmem launch arg needed (compile-time + * static shmem allocation). */ + __shared__ float new_max_dd_ema_shared; + + /* ── Thread 0: training_sharpe_ema (adaptive α) ───────────────────── + * Reproduces training_loop.rs:5091-5099 verbatim: + * if !initialized: ema = obs; initialized = true + * else: + * err = |obs − prev| + * alpha = min(0.05 + 0.3 × err, 0.5) + * ema = (1 − α) × prev + α × obs + * The `_initialized` flag transitions 0→1 on the kernel's first call + * after a fold reset; D3 mirrors host-side via the `sharpe_initialized` + * arg (D4 will turn this into a mapped-pinned scalar). All clamp + * constants (0.05, 0.3, 0.5) are Invariant 1 anchors per + * `feedback_no_quickfixes.md` — they encode the host-side formula + * shape exactly and migrate verbatim. + */ + if (tid == 0) { + float new_sharpe_ema; + if (sharpe_initialized == 0) { + new_sharpe_ema = epoch_sharpe; + } else { + const float err = fabsf(epoch_sharpe - prev_sharpe_ema); + float alpha = 0.05f + 0.3f * err; + if (alpha > 0.5f) alpha = 0.5f; + new_sharpe_ema = (1.0f - alpha) * prev_sharpe_ema + alpha * epoch_sharpe; + } + scratch_buf[sharpe_ema_idx] = new_sharpe_ema; + } + + /* ── Thread 1: max_dd_ema (fixed α=0.1, sentinel<1e-12) ───────────── + * Reproduces training_loop.rs:5043-5047 verbatim: + * if prev < 1e-12: ema = obs + * else: ema = 0.9 × prev + 0.1 × obs + * 1e-12 is Invariant 1 — the deleted host formula's literal bootstrap + * sentinel (host f64 in source; cast to f32 here matches the producer + * scratch contract). Result is also published to shmem so thread 2 + * can derive `is_low` from the *new* EMA value (host-side at + * training_loop.rs:5050 reads `self.max_dd_ema` after the update). + */ + if (tid == 1) { + float new_max_dd_ema; + if (prev_max_dd_ema < 1e-12f) { + new_max_dd_ema = epoch_max_dd; + } else { + new_max_dd_ema = 0.9f * prev_max_dd_ema + 0.1f * epoch_max_dd; + } + scratch_buf[max_dd_ema_idx] = new_max_dd_ema; + new_max_dd_ema_shared = new_max_dd_ema; + } + + /* Barrier: thread 2 reads thread 1's published max_dd_ema. */ + __syncthreads(); + + /* ── Thread 2: low_dd_ratio (fixed α=0.15, binary input) ──────────── + * Reproduces training_loop.rs:5050-5052 verbatim: + * thresh = new_max_dd_ema × 0.5 + * is_low = (max_dd < thresh) ? 1.0 : 0.0 + * ema = 0.85 × prev + 0.15 × is_low + * 0.5 is the half-EMA threshold (host-side adversarial regime + * activation boundary at training_loop.rs:5050). All constants + * Invariant 1 anchors — deleted host formula's literals. + * No sentinel branch: initial 0.0 from the constructor at + * constructor.rs:710 behaves correctly under the EMA recurrence + * (matches host-side, which also has no special init). + */ + if (tid == 2) { + const float new_max_dd_ema = new_max_dd_ema_shared; + const float thresh = new_max_dd_ema * 0.5f; + const float is_low = (epoch_max_dd < thresh) ? 1.0f : 0.0f; + const float new_low_dd_ratio = + 0.85f * prev_low_dd_ratio + 0.15f * is_low; + scratch_buf[low_dd_ratio_idx] = new_low_dd_ratio; + } + + /* Single PCIe-visible barrier covers all three scratch writes; thread + * count is 3 so no warp-divergence cost. Mapped-pinned host_ptr aliasing + * requires threadfence_system before the downstream apply_pearls reads + * (matches D1 PnL kernel and D2 health composition kernel patterns). */ + __syncthreads(); + if (tid == 0) { + __threadfence_system(); + } +} + +} /* extern "C" */ diff --git a/crates/ml/src/trainers/dqn/state_reset_registry.rs b/crates/ml/src/trainers/dqn/state_reset_registry.rs index 57c618a44..76d378d50 100644 --- a/crates/ml/src/trainers/dqn/state_reset_registry.rs +++ b/crates/ml/src/trainers/dqn/state_reset_registry.rs @@ -723,6 +723,26 @@ impl StateResetRegistry { category: ResetCategory::FoldReset, description: "ISV[HEALTH_SCORE_INDEX=290..GRAD_NORM_NORM_INDEX+1=294) — Layer D D2 LearningHealth composition outputs (composed score + q_gap_norm + q_var_norm + grad_norm_norm). SP5 Pearl A sentinel 0 at fold boundary so the new fold's first `launch_health_composition` fires the first-observation replacement (Task D2, 2026-05-02). Producer is `health_composition_kernel.cu` — single-block 1-thread arithmetic kernel reproducing `learning_health.rs::NormalizedComponents::from_raw` + `compose()` bit-for-bit. Consumer wires atomically in D4 (per `feedback_no_partial_refactor.md`). The underlying EMAs (`HealthEmaTrackers.q_gap_ema` / `q_var_ema` / `grad_norm_ema`, plus the 4 non-EMA inputs sourced from per-branch Q-gap dispersion / atom utilisation / spectral gap / Adam m-flat cosine) reset at fold boundary so the composition outputs reset together. Wiener-state companion: bulk memset of the entire `wiener_state_buf[0..SP5_WIENER_TOTAL_FLOATS=573)` at fold boundary covers the new D2 triples at offsets [213+(290-174)*3 .. 213+(294-174)*3) = [561..573). The buffer grew from 561 to 573 floats with D2 (SP5_PRODUCER_COUNT linear span 116 → 120; unique-slot count 114 → 118; see SP5_PRODUCER_COUNT docstring for the linear-span vs unique-slot distinction). reset_for_fold's existing bulk memset of the whole buffer length covers the new triples atomically — no per-slot wiener entry is needed.", }, + // SP5 Layer D Task D3: Training metrics EMA outputs (3 ISV slots). + // ISV[294..297) — Layer D D3 host-formula EMAs for + // `training_sharpe_ema` / `max_dd_ema` / `low_dd_ratio` migrated + // from the host-side adversarial-regime + adaptive-DSR block at + // `training_loop.rs:5039-5113`. The corresponding host scalars + // (`training_sharpe_ema` + `_initialized`, `max_dd_ema`, + // `low_dd_ratio`) reset at fold boundary on the host side + // (constructor.rs:587-588/710-712 sets them to zero on each + // construction; D4 will route the per-fold reset through the + // existing fold-reset codepath alongside this ISV-side + // sentinel). All three ISV slots take the standard Pearl A + // sentinel-bootstrap path. Both halves of the contract reset + // together: ISV slot zeroed AND `wiener_state_buf` SP5 block + // bulk memset(0) at fold boundary covers the three new triples + // at offsets [213+(294-174)*3 .. 213+(297-174)*3) = [573..582). + RegistryEntry { + name: "sp5_training_metrics_ema", + category: ResetCategory::FoldReset, + description: "ISV[TRAINING_SHARPE_EMA_INDEX=294..LOW_DD_RATIO_INDEX+1=297) — Layer D D3 training-metrics EMA outputs (training_sharpe_ema / max_dd_ema / low_dd_ratio). SP5 Pearl A sentinel 0 at fold boundary so the new fold's first `launch_training_metrics_ema` fires the first-observation replacement (Task D3, 2026-05-02). Producer is `training_metrics_ema_kernel.cu` — single-block 3-thread arithmetic kernel (one thread per metric) reproducing the host-side adaptive-α Sharpe EMA + fixed-α=0.1 max_dd EMA + fixed-α=0.15 low_dd_ratio EMA from `training_loop.rs:5039-5113` bit-for-bit. Consumer wires atomically in D4 (per `feedback_no_partial_refactor.md`). The underlying host scalars (`training_sharpe_ema` + `_initialized`, `max_dd_ema`, `low_dd_ratio` on `DQNTrainer`) reset at fold boundary alongside the ISV slots — D4 ties the host-side fold-reset call into the same path. Wiener-state companion: bulk memset of the entire `wiener_state_buf[0..SP5_WIENER_TOTAL_FLOATS=582)` at fold boundary covers the new D3 triples at offsets [213+(294-174)*3 .. 213+(297-174)*3) = [573..582). The buffer grew from 573 to 582 floats with D3 (SP5_PRODUCER_COUNT linear span 120 → 123; unique-slot count 118 → 121; see SP5_PRODUCER_COUNT docstring for the linear-span vs unique-slot distinction). reset_for_fold's existing bulk memset of the whole buffer length covers the new triples atomically — no per-slot wiener entry is needed. Per `feedback_trust_code_not_docs.md` the third metric is `low_dd_ratio` (the binary low-drawdown indicator EMA at `training_loop.rs:5052`), not the speculative `gamma_blend` referenced in the SP5 plan §D Task D3 — that field does not exist in the codebase.", + }, ]; Self { entries } } diff --git a/crates/ml/src/trainers/dqn/trainer/training_loop.rs b/crates/ml/src/trainers/dqn/trainer/training_loop.rs index 2e968666c..9819cde5a 100644 --- a/crates/ml/src/trainers/dqn/trainer/training_loop.rs +++ b/crates/ml/src/trainers/dqn/trainer/training_loop.rs @@ -6378,6 +6378,48 @@ impl DQNTrainer { fused.trainer().write_isv_signal_at(PNL_MAX_DD_INDEX, 0.0); } } + "sp5_health_composition" => { + // Layer D Task D2: zero ISV[290..294) — composed health score + + // 3 normalised intermediates (q_gap_norm, q_var_norm, + // grad_norm_norm = grad_stable) — so the new fold's first + // `launch_health_composition` triggers Pearl A's + // first-observation replacement (per `pearl_first_observation_bootstrap.md`). + // Wiener-state companion is reset by the existing bulk memset + // of `wiener_state_buf[0..SP5_WIENER_TOTAL_FLOATS=582)` covered + // by the `sp4_wiener_state` registry entry's dispatch arm. + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp5_isv_slots::{ + HEALTH_SCORE_INDEX, Q_GAP_NORM_INDEX, Q_VAR_NORM_INDEX, GRAD_NORM_NORM_INDEX, + }; + fused.trainer().write_isv_signal_at(HEALTH_SCORE_INDEX, 0.0); + fused.trainer().write_isv_signal_at(Q_GAP_NORM_INDEX, 0.0); + fused.trainer().write_isv_signal_at(Q_VAR_NORM_INDEX, 0.0); + fused.trainer().write_isv_signal_at(GRAD_NORM_NORM_INDEX, 0.0); + } + } + "sp5_training_metrics_ema" => { + // Layer D Task D3: zero ISV[294..297) — training_sharpe_ema / + // max_dd_ema / low_dd_ratio — so the new fold's first + // `launch_training_metrics_ema` triggers Pearl A's + // first-observation replacement (per `pearl_first_observation_bootstrap.md`). + // Wiener-state companion is reset by the existing bulk memset + // of `wiener_state_buf[0..SP5_WIENER_TOTAL_FLOATS=582)` covered + // by the `sp4_wiener_state` registry entry's dispatch arm. + // The host-side scalar companions (`self.training_sharpe_ema` + + // `self.training_sharpe_ema_initialized`, `self.max_dd_ema`, + // `self.low_dd_ratio`) reset alongside as part of the host-side + // per-fold logic — D4's atomic call-site refactor will tie + // the host-side reset into this dispatch arm directly. In + // additive D3, only the ISV slots reset here. + if let Some(ref fused) = self.fused_ctx { + use crate::cuda_pipeline::sp5_isv_slots::{ + TRAINING_SHARPE_EMA_INDEX, MAX_DD_EMA_INDEX, LOW_DD_RATIO_INDEX, + }; + fused.trainer().write_isv_signal_at(TRAINING_SHARPE_EMA_INDEX, 0.0); + fused.trainer().write_isv_signal_at(MAX_DD_EMA_INDEX, 0.0); + fused.trainer().write_isv_signal_at(LOW_DD_RATIO_INDEX, 0.0); + } + } _ => { return Err(crate::MLError::ModelError(format!( "StateResetRegistry reset dispatch: unknown name '{}'. \ diff --git a/crates/ml/tests/sp5_producer_unit_tests.rs b/crates/ml/tests/sp5_producer_unit_tests.rs index 9aa4bbe41..0bfef32ee 100644 --- a/crates/ml/tests/sp5_producer_unit_tests.rs +++ b/crates/ml/tests/sp5_producer_unit_tests.rs @@ -39,6 +39,9 @@ //! Layer D Task D2 (Health composition): //! 19. `health_composition_kernel_correctness` //! +//! Layer D Task D3 (Training metrics EMA): +//! 20. `training_metrics_ema_kernel_correctness` +//! //! All tests use analytically-known synthetic inputs with no CPU reference //! implementation, per `feedback_no_cpu_test_fallbacks.md`. All mapped-pinned //! memory (no DtoH), all `#[ignore]`-gated for GPU. @@ -93,6 +96,9 @@ const SP5_PNL_AGGREGATION_CUBIN: &[u8] = const SP5_HEALTH_COMPOSITION_CUBIN: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/health_composition_kernel.cubin")); +const SP5_TRAINING_METRICS_EMA_CUBIN: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/training_metrics_ema_kernel.cubin")); + // ── Helpers ─────────────────────────────────────────────────────────────────── fn load_pearl_3_sigma_kernel(stream: &Arc) -> CudaFunction { @@ -2361,3 +2367,186 @@ fn health_composition_kernel_correctness() { assert!((out[3] - exp_grad_stable).abs() < tol_strict, "grad_norm_norm (= grad_stable): expected {exp_grad_stable:.7}, got {:.7}", out[3]); } + +// ── Layer D Task D3 (Training metrics EMA) ──────────────────────────────────── + +fn load_training_metrics_ema_kernel(stream: &Arc) -> CudaFunction { + let module = stream + .context() + .load_cubin(SP5_TRAINING_METRICS_EMA_CUBIN.to_vec()) + .expect("load training_metrics_ema_kernel cubin"); + module + .load_function("training_metrics_ema_update") + .expect("load training_metrics_ema_update function") +} + +/// SP5 Layer D Task D3: `training_metrics_ema_update` formula-fidelity test. +/// +/// Asserts the kernel reproduces the host-side EMA arithmetic at +/// `training_loop.rs:5039-5113` bit-for-bit (within float precision). Per +/// `feedback_no_cpu_test_fallbacks.md` the test uses synthetic inputs whose +/// closed-form expected outputs are pre-computed and pinned as literals — +/// no programmatic re-implementation of the formula in the test body. +/// +/// Synthetic inputs (deliberately exercise the steady-state branch of all +/// three threads — first-observation paths are tested in the second case +/// below by setting `sharpe_initialized=0` and `prev_max_dd_ema=0`): +/// +/// epoch_sharpe = 1.0 +/// epoch_max_dd = 0.05 (5% drawdown) +/// prev_sharpe_ema = 0.5 (sharpe diverged from prev by |err|=0.5) +/// prev_max_dd_ema = 0.10 +/// prev_low_dd_ratio = 0.4 +/// sharpe_initialized = true +/// +/// Closed-form expected outputs: +/// +/// tid=0 (training_sharpe_ema, adaptive α): +/// err = |1.0 − 0.5| = 0.5 +/// α = clamp(0.05 + 0.3 · 0.5, 0, 0.5) = clamp(0.20, 0, 0.5) = 0.20 +/// ema = (1 − 0.20) · 0.5 + 0.20 · 1.0 = 0.40 + 0.20 = 0.60 +/// +/// tid=1 (max_dd_ema, fixed α=0.1, prev > 1e-12 so steady-state): +/// ema = 0.9 · 0.10 + 0.1 · 0.05 = 0.090 + 0.005 = 0.095 +/// +/// tid=2 (low_dd_ratio, fixed α=0.15 over binary `is_low`): +/// thresh = new_max_dd_ema · 0.5 = 0.095 · 0.5 = 0.0475 +/// is_low = (epoch_max_dd=0.05 < thresh=0.0475) ? 1.0 : 0.0 = 0.0 +/// (epoch_max_dd is *just* above half the new EMA) +/// ema = 0.85 · 0.4 + 0.15 · 0.0 = 0.340 +/// +/// **Formula fidelity is the load-bearing invariant** for D3: the kernel +/// performs a verbatim structural migration of the host-side EMA update, +/// no algorithmic change. The test pins the adaptive-α arithmetic, the +/// fixed-α recurrence, and the cross-thread `is_low` derivation +/// (tid=2 must read tid=1's *new* output via __syncthreads()). +#[test] +#[ignore = "requires GPU"] +fn training_metrics_ema_kernel_correctness() { + let ctx = CudaContext::new(0).expect("CUDA context"); + let stream = ctx.new_stream().expect("stream"); + let kernel = load_training_metrics_ema_kernel(&stream); + + // ── Case 1: steady-state branch on all three threads ──────────────── + { + let epoch_sharpe: f32 = 1.0; + let epoch_max_dd: f32 = 0.05; + let prev_sharpe_ema: f32 = 0.5; + let prev_max_dd_ema: f32 = 0.10; + let prev_low_dd_ratio: f32 = 0.4; + let sharpe_initialized: i32 = 1; + + let scratch = unsafe { MappedF32Buffer::new(3) }.expect("scratch alloc"); + + let scratch_dev = scratch.dev_ptr; + let sharpe_idx: i32 = 0; + let max_dd_idx: i32 = 1; + let low_dd_ratio_idx: i32 = 2; + + unsafe { + stream + .launch_builder(&kernel) + .arg(&epoch_sharpe) + .arg(&epoch_max_dd) + .arg(&prev_sharpe_ema) + .arg(&prev_max_dd_ema) + .arg(&prev_low_dd_ratio) + .arg(&sharpe_initialized) + .arg(&scratch_dev) + .arg(&sharpe_idx) + .arg(&max_dd_idx) + .arg(&low_dd_ratio_idx) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (3, 1, 1), + shared_mem_bytes: 0, + }) + .expect("launch training_metrics_ema_update steady-state"); + } + stream.synchronize().expect("sync after training_metrics_ema_update steady-state"); + + let out = scratch.read_all(); + + // Analytical ground truth (closed-form literals): + let exp_sharpe_ema: f32 = 0.60; + let exp_max_dd_ema: f32 = 0.095; + let exp_low_dd_ratio: f32 = 0.340; + + println!("[steady] training_sharpe_ema = {:.6} (expected {:.6})", out[0], exp_sharpe_ema); + println!("[steady] max_dd_ema = {:.6} (expected {:.6})", out[1], exp_max_dd_ema); + println!("[steady] low_dd_ratio = {:.6} (expected {:.6})", out[2], exp_low_dd_ratio); + + // FP tolerances: each output is 2-3 multiply-adds of small (≤1.0) + // values; 1e-6 absolute is comfortable. + let tol: f32 = 1e-6; + assert!((out[0] - exp_sharpe_ema).abs() < tol, + "training_sharpe_ema [steady]: expected {exp_sharpe_ema:.6}, got {:.6}", out[0]); + assert!((out[1] - exp_max_dd_ema).abs() < tol, + "max_dd_ema [steady]: expected {exp_max_dd_ema:.6}, got {:.6}", out[1]); + assert!((out[2] - exp_low_dd_ratio).abs() < tol, + "low_dd_ratio [steady]: expected {exp_low_dd_ratio:.6}, got {:.6}", out[2]); + } + + // ── Case 2: cold-start branches on tid=0 (sharpe_initialized=0) + // and tid=1 (prev_max_dd_ema < 1e-12). tid=2 has no + // host-side sentinel; it always uses the EMA recurrence. + { + let epoch_sharpe: f32 = 0.7; + let epoch_max_dd: f32 = 0.08; + let prev_sharpe_ema: f32 = 99.0; // ignored when initialized=0 + let prev_max_dd_ema: f32 = 0.0; // < 1e-12 sentinel triggers replace + let prev_low_dd_ratio: f32 = 0.0; // EMA from 0 is recurrence-safe + let sharpe_initialized: i32 = 0; + + let scratch = unsafe { MappedF32Buffer::new(3) }.expect("scratch alloc"); + + let scratch_dev = scratch.dev_ptr; + let sharpe_idx: i32 = 0; + let max_dd_idx: i32 = 1; + let low_dd_ratio_idx: i32 = 2; + + unsafe { + stream + .launch_builder(&kernel) + .arg(&epoch_sharpe) + .arg(&epoch_max_dd) + .arg(&prev_sharpe_ema) + .arg(&prev_max_dd_ema) + .arg(&prev_low_dd_ratio) + .arg(&sharpe_initialized) + .arg(&scratch_dev) + .arg(&sharpe_idx) + .arg(&max_dd_idx) + .arg(&low_dd_ratio_idx) + .launch(LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (3, 1, 1), + shared_mem_bytes: 0, + }) + .expect("launch training_metrics_ema_update cold-start"); + } + stream.synchronize().expect("sync after training_metrics_ema_update cold-start"); + + let out = scratch.read_all(); + + // Closed-form for cold-start: + // tid=0: sharpe_initialized==0 ⇒ output replaces with epoch_sharpe = 0.7. + // tid=1: prev_max_dd_ema<1e-12 ⇒ output replaces with epoch_max_dd = 0.08. + // tid=2: thresh=0.04, is_low=(0.08<0.04)?1:0 = 0; ema = 0.85·0 + 0.15·0 = 0. + let exp_sharpe_ema: f32 = 0.7; + let exp_max_dd_ema: f32 = 0.08; + let exp_low_dd_ratio: f32 = 0.0; + + println!("[cold] training_sharpe_ema = {:.6} (expected {:.6})", out[0], exp_sharpe_ema); + println!("[cold] max_dd_ema = {:.6} (expected {:.6})", out[1], exp_max_dd_ema); + println!("[cold] low_dd_ratio = {:.6} (expected {:.6})", out[2], exp_low_dd_ratio); + + let tol: f32 = 1e-6; + assert!((out[0] - exp_sharpe_ema).abs() < tol, + "training_sharpe_ema [cold]: expected {exp_sharpe_ema:.6}, got {:.6}", out[0]); + assert!((out[1] - exp_max_dd_ema).abs() < tol, + "max_dd_ema [cold]: expected {exp_max_dd_ema:.6}, got {:.6}", out[1]); + assert!((out[2] - exp_low_dd_ratio).abs() < tol, + "low_dd_ratio [cold]: expected {exp_low_dd_ratio:.6}, got {:.6}", out[2]); + } +} diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index fe5d5e192..5d6362c36 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -3416,3 +3416,61 @@ All edge constants and weights are migrated verbatim from `learning_health.rs:39 **No call-site wiring** — the host-side composition block in `training_loop.rs:2658-2745` is intentionally untouched; D4 atomically migrates the 5 EMA sites (training_sharpe_ema, max_dd_ema, gamma_blend, LearningHealth pipeline, HealthEmaTrackers) to the GPU producer chain in a single commit. Refs: SP5 plan §D Task D2; builds on D1 (`5ee795f14`); `feedback_no_cpu_compute_strict`, `feedback_no_partial_refactor`, `feedback_no_atomicadd`, `feedback_no_cpu_test_fallbacks`, `pearl_first_observation_bootstrap`. + +### SP5 Layer D Task D3 — Training metrics EMA kernel (additive) (2026-05-02) + +Third (and final) Layer D producer kernel after D1 (PnL aggregation) and D2 (LearningHealth composition), replacing the host-side per-epoch EMA arithmetic block at `training_loop.rs:5039-5113` per `feedback_no_cpu_compute_strict.md`. D3 lands additively — kernel + Rust launcher + 3 reserved ISV slots are present, but the host-side updates continue to run unchanged. D4 (the atomic Layer D commit) wires D1+D2+D3 to call sites in lockstep per `feedback_no_partial_refactor.md`. + +**Plan brief vs code reality (per `feedback_trust_code_not_docs.md`):** the SP5 plan §D Task D3 brief names three host EMAs — `training_sharpe_ema`, `max_dd_ema`, and `gamma_blend` (the third labelled "host-aggregated eval result"). Investigation: `gamma_blend` does not exist in the codebase (`grep -rn 'gamma_blend' crates/` returns zero hits). The actual third EMA at the cited line range is `low_dd_ratio` — a binary EMA at `training_loop.rs:5052` driving the adversarial-regime hysteresis at `training_loop.rs:5054-5070`. Per the feedback rule, D3 ports what the code does, not what the plan speculated. The slot constant is `LOW_DD_RATIO_INDEX=296`. + +**Host-side formulas reproduced (formula fidelity > anything else):** + +Source: `crates/ml/src/trainers/dqn/trainer/training_loop.rs:5039-5113`. Three EMAs run sequentially in a single block (the adversarial-regime + adaptive-DSR closing block of the per-epoch `process_epoch_boundary`): + +```text +// (1) max_dd_ema (training_loop.rs:5043-5047) — fixed α = 0.1 +if max_dd_ema < 1e-12: + max_dd_ema = obs_max_dd // sentinel: first observation replaces +else: + max_dd_ema = 0.9 · max_dd_ema + 0.1 · obs_max_dd + +// (2) low_dd_ratio (training_loop.rs:5050-5052) — fixed α = 0.15 over binary input +thresh = max_dd_ema · 0.5 // reads the *new* max_dd_ema +is_low = (obs_max_dd < thresh) ? 1.0 : 0.0 +low_dd_ratio = 0.85 · low_dd_ratio + 0.15 · is_low + +// (3) training_sharpe_ema (training_loop.rs:5091-5099) — adaptive α +if !sharpe_initialized: + training_sharpe_ema = obs_sharpe // sentinel: `_initialized` flag (host) + sharpe_initialized = true +else: + err = |obs_sharpe − training_sharpe_ema| + α = clamp(0.05 + 0.3 · err, 0, 0.5) + training_sharpe_ema = (1 − α) · training_sharpe_ema + α · obs_sharpe +``` + +All EMA constants (1e-12 sentinel, 0.1, 0.15, 0.05, 0.3, 0.5, 0.5 threshold) are Invariant 1 anchors per `feedback_no_quickfixes.md` — the deleted host formula's literals migrated verbatim. The kernel is structurally additive: only the EMA arithmetic migrates. The `_initialized` flag stays host-tracked in additive D3 (passed as `i32` per `feedback_cudarc_f64_f32_abi.md`); D4 will lift it to a mapped-pinned scalar. + +**Cross-thread dependency:** thread 2 (`low_dd_ratio`) reads thread 1's *new* `max_dd_ema` via a single `__syncthreads()` and a static `__shared__ float`. Mirrors the host-side ordering at `training_loop.rs:5050` where `self.max_dd_ema` is read after the update at line 5047. + +**What this lands:** +- `training_metrics_ema_kernel.cu` (147 lines) — single-block 3-thread arithmetic kernel. One thread per metric: tid=0 sharpe (adaptive α with `sharpe_initialized` host flag sentinel), tid=1 max_dd (fixed α=0.1, `prev<1e-12` sentinel), tid=2 low_dd_ratio (fixed α=0.15 over binary `is_low` derived from tid=1's *new* max_dd_ema via `__syncthreads()`). Writes 3 floats to `scratch[SCRATCH_TRAINING_METRICS_EMA_BASE=215..218)`. No atomicAdd. Single `__threadfence_system()` after the writes for mapped-pinned host_ptr / dev_ptr alias visibility. +- Rust launcher `GpuDqnTrainer::launch_training_metrics_ema(epoch_sharpe, epoch_max_dd, prev_sharpe_ema, prev_max_dd_ema, prev_low_dd_ratio, sharpe_initialized: bool)` — chains 3 `apply_pearls_ad_kernel` calls on the same stream after the producer per the SP5 GPU-only Pearls A+D pattern. Wiener offset formula `213 + (isv_slot − SP5_SLOT_BASE) × 3` matches the existing SP5 launcher template (D1 PnL, D2 health composition). +- 3 new ISV slots `TRAINING_SHARPE_EMA_INDEX=294`, `MAX_DD_EMA_INDEX=295`, `LOW_DD_RATIO_INDEX=296` registered in `sp5_isv_slots.rs`. `SP5_SLOT_END` 294 → 297; `ISV_TOTAL_DIM` 294 → 297; `LAYOUT_FINGERPRINT_FRAGMENT` extended with `TRAINING_SHARPE_EMA` / `MAX_DD_EMA` / `LOW_DD_RATIO` entries so existing checkpoints fail-fast on layout change. +- `SP5_PRODUCER_COUNT` (wiener-buffer linear-span constant) 120 → 123; unique-slot count 118 → 121; wiener buffer 573 → 582 floats. The D3 block sits at wiener offsets `[(294-174)·3+213 .. (297-174)·3+213) = [573..582)`, atomically reset by the existing bulk `wiener_state_buf[0..582)` memset at fold boundary. +- `producer_step_scratch_buf` grows 215 → 218 (`SP5_SCRATCH_TOTAL = 218`). 1 new scratch constant `SCRATCH_TRAINING_METRICS_EMA_BASE=215`. The constructor's scratch-layout audit comment block updated with the D3 entry. +- `StateResetRegistry` entry `sp5_training_metrics_ema` (FoldReset; D3 outputs reset alongside the underlying host scalars at fold boundary — the host-side `training_sharpe_ema` + `_initialized`, `max_dd_ema`, `low_dd_ratio` reset together with the ISV slots; D4 will tie the host-side reset into the same path) + matching dispatch arm in `reset_named_state` zeroing ISV[294..297) at fold boundary so the new fold's first launch fires Pearl A's first-observation replacement (paired with the existing bulk wiener_state_buf memset, per `pearl_first_observation_bootstrap.md`). +- D2 missed adding a dispatch arm for its `sp5_health_composition` registry entry — D3 closes that gap so all three Layer D entries (D1+D2+D3) have matching arms, keeping the registry/dispatch invariant intact ahead of D4's atomic refactor. +- `build.rs` cubin registration; static `SP5_TRAINING_METRICS_EMA_CUBIN` byte slice + `training_metrics_ema_kernel: CudaFunction` field on `GpuDqnTrainer` + cubin/function loader entry mirroring the D2 health composition kernel template. +- GPU-gated unit test `training_metrics_ema_kernel_correctness` in `sp5_producer_unit_tests.rs` (2 sub-cases): steady-state (all three threads exercise the EMA recurrence, closed-form expected outputs `sharpe_ema=0.60`, `max_dd_ema=0.095`, `low_dd_ratio=0.340`) + cold-start (tid=0 sharpe-init flag → replace, tid=1 max_dd<1e-12 → replace, tid=2 EMA-from-zero behaves correctly). Expected outputs pinned as literal constants per `feedback_no_cpu_test_fallbacks.md` — no programmatic CPU mirror. +- ISV slot test `training_metrics_ema_slots_contiguous_and_above_health_block` mirroring D1/D2's contiguity tests — verifies slot ordering, internal contiguity, and `SP5_SLOT_END` / `SP5_PRODUCER_COUNT` invariants. + +**Verification gates:** +- `cargo check -p ml --offline` clean. +- `cargo build -p ml --release --offline --features cuda` clean — `training_metrics_ema_kernel.cubin` compiles via nvcc (sm_80). +- `cargo test -p ml --lib --offline -- sp5_isv_slots state_reset_registry` 8/8 pass (slot-layout invariants + new `training_metrics_ema_slots_contiguous_and_above_health_block` test + 4 pre-existing slot tests + 3 pre-existing registry tests). +- GPU correctness test (`training_metrics_ema_kernel_correctness`) `--ignored`-gated; fires on the next L40S smoke alongside the D1/D2 tests and the 17 pre-existing SP5 producer unit tests (test count 19 → 20 with D3). + +**No call-site wiring** — the host-side EMA block at `training_loop.rs:5039-5113` is intentionally untouched; D4 atomically migrates the 5 EMA sites (training_sharpe_ema, max_dd_ema, low_dd_ratio, LearningHealth pipeline, HealthEmaTrackers) to the GPU producer chain in a single commit. **Layer D additive infrastructure complete after this commit; D4 (atomic 5-site host-EMA → GPU migration) is next.** + +Refs: SP5 plan §D Task D3; builds on D1 (`5ee795f14`) + D2 (`e49756ac9`); `feedback_no_cpu_compute_strict`, `feedback_no_partial_refactor`, `feedback_trust_code_not_docs`, `feedback_no_atomicadd`, `feedback_no_cpu_test_fallbacks`, `feedback_no_quickfixes`, `feedback_cudarc_f64_f32_abi`, `pearl_first_observation_bootstrap`.