From 71aade18cf016284e50d8befdc2a504f4d670261 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sat, 9 May 2026 18:57:30 +0200 Subject: [PATCH] feat(sp20): Phase 1.2 sp20_emas_compute kernel MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Component 5 / Kernel 1 of the SP20 design — central state-tracker that updates 8 Wiener-α EMAs per training step: - 4 ISV slots: ALPHA_EMA (511), WR_EMA (512), HOLD_PCT_EMA (515), HOLD_REWARD_EMA (516) - 4 internal: trade_duration_ema, aux_conf_p50_ema, aux_conf_std_ema, aux_dir_acc_ema (private scratch consumed by Phase 1.3 controllers) Per-EMA i32 observation counters (mapped-pinned obs_count[8]) handle the pearl_first_observation_bootstrap sentinel transition correctly even when 0.0 is a legitimate observation (e.g., first-loss WR=0). count==0 ⇒ replace, count>0 ⇒ Wiener-blend at α = 0.4 (WIENER_ALPHA_FLOOR per pearl_wiener_alpha_floor_for_nonstationary). Phase 1.2 lands kernel + launcher + 5 tests (4 GPU oracle, 1 floor lock) + build entry + audit doc atomically per feedback_no_partial_refactor. Production wire-up (Phase 1.4) deferred — kernel is dead code until then. Verified on RTX 3050 Ti (sm_86): - 4 GPU oracle tests pass: first_observation_replaces_sentinel, wiener_alpha_converges_to_long_run_mean, hold_reward_ema_gated_on_hold_bars, per_step_emas_fire_unconditionally - 4 launcher unit tests pass (constants + struct sanity) - 1 floor-lock test pass (WIENER_ALPHA_FLOOR == 0.4) Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml/build.rs | 17 + crates/ml/src/cuda_pipeline/mod.rs | 8 + .../ml/src/cuda_pipeline/sp20_emas_compute.rs | 315 ++++++++++++++++ .../cuda_pipeline/sp20_emas_compute_kernel.cu | 222 +++++++++++ crates/ml/tests/sp20_emas_compute_test.rs | 347 ++++++++++++++++++ docs/dqn-wire-up-audit.md | 204 ++++++++++ 6 files changed, 1113 insertions(+) create mode 100644 crates/ml/src/cuda_pipeline/sp20_emas_compute.rs create mode 100644 crates/ml/src/cuda_pipeline/sp20_emas_compute_kernel.cu create mode 100644 crates/ml/tests/sp20_emas_compute_test.rs diff --git a/crates/ml/build.rs b/crates/ml/build.rs index 7be987c37..f7ce6bba3 100644 --- a/crates/ml/build.rs +++ b/crates/ml/build.rs @@ -1379,6 +1379,23 @@ fn main() { // (`tests/sp20_stats_compute_test.rs`) atomically per // `feedback_no_partial_refactor`. "sp20_stats_compute_kernel.cu", + // SP20 Phase 1.2 (2026-05-09): 8 Wiener-α EMA producer. + // Component 5 / Kernel 1 of the SP20 fused-producer chain. + // Single-block, single-thread kernel updating 4 ISV slots + // (ALPHA_EMA, WR_EMA, HOLD_PCT_EMA, HOLD_REWARD_EMA) plus + // 4 internal scratch slots (trade_duration_ema + + // aux_conf_p50/std/dir_acc EMAs) per training step. Per-EMA + // observation counters (i32 mapped-pinned) handle the + // `pearl_first_observation_bootstrap` sentinel transition + // (count==0 ⇒ replace, count>0 ⇒ Wiener-blend) so EMAs + // that legitimately observe 0 (e.g., first-loss WR) don't + // re-bootstrap. Wiener-α floor 0.4 hardcoded as a structural + // property per `pearl_wiener_alpha_floor_for_nonstationary`. + // No atomicAdd, no host branches — captureable in the per- + // step CUDA Graph. Phase 1.4 wires the production launch + // site atomically with Kernel 2 (controllers) per + // `feedback_no_partial_refactor`. + "sp20_emas_compute_kernel.cu", ]; // ALL kernels get common header (BF16 types + wrappers) diff --git a/crates/ml/src/cuda_pipeline/mod.rs b/crates/ml/src/cuda_pipeline/mod.rs index 16f8a107e..3dbf937d4 100644 --- a/crates/ml/src/cuda_pipeline/mod.rs +++ b/crates/ml/src/cuda_pipeline/mod.rs @@ -92,6 +92,14 @@ pub use sp4_wiener_ema::{ // emits `[p50, std]` into a `MappedF32Buffer<2>` that the EMA producer // (Phase 1.2) blends into ISV slots [510..520). pub mod sp20_stats_compute; +// SP20 Phase 1.2 (2026-05-09): 8 Wiener-α EMA producer kernel +// launcher. Component 5 / Kernel 1 of the SP20 fused-producer chain +// (Kernel 2 = controllers lands in Phase 1.3; Kernel 3 = stats already +// landed in Phase 1.1). Updates 4 ISV slots (ALPHA_EMA, WR_EMA, +// HOLD_PCT_EMA, HOLD_REWARD_EMA) plus 4 internal scratch slots +// (trade_duration_ema + 3 aux EMAs) per training step, gated by +// is_close / action_is_hold per the SP20 spec §4.5. +pub mod sp20_emas_compute; // `launch_apply_pearls` is `pub(crate)` and consumed only inside this // crate via `use crate::cuda_pipeline::sp4_wiener_ema::launch_apply_pearls`. // `pearls_ad_update` stays `pub` for the test-oracle path in diff --git a/crates/ml/src/cuda_pipeline/sp20_emas_compute.rs b/crates/ml/src/cuda_pipeline/sp20_emas_compute.rs new file mode 100644 index 000000000..5d133b5d4 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/sp20_emas_compute.rs @@ -0,0 +1,315 @@ +#![allow(unsafe_code)] // CUDA kernel launch via cudarc requires unsafe. +//! SP20 Phase 1.2 (2026-05-09) — `sp20_emas_compute_kernel` launcher. +//! +//! Component 5 / Kernel 1 of the SP20 design — the central state-tracker +//! that updates 8 Wiener-α EMAs per training step. Reads per-step / +//! per-trade observations + the upstream [`super::sp20_stats_compute`] +//! (Phase 1.1) outputs, then writes 4 ISV slots +//! ([`crate::cuda_pipeline::sp14_isv_slots::ALPHA_EMA_INDEX`], +//! [`crate::cuda_pipeline::sp14_isv_slots::WR_EMA_INDEX`], +//! [`crate::cuda_pipeline::sp14_isv_slots::HOLD_PCT_EMA_INDEX`], +//! [`crate::cuda_pipeline::sp14_isv_slots::HOLD_REWARD_EMA_INDEX`]) plus +//! 4 internal scratch slots (defined by [`EMA_INTERNAL_*`] constants +//! below) that Phase 1.3's `sp20_controllers_compute` consumes to derive +//! `LOSS_CAP / HOLD_COST_SCALE / TARGET_HOLD_PCT / N_STEP / +//! AUX_CONF_THRESHOLD / AUX_GATE_TEMP`. +//! +//! ## EMAs maintained +//! +//! | EMA | Storage | Gate | Observation source | +//! |-----|---------|------|--------------------| +//! | `ALPHA_EMA` | ISV[511] | `is_close == 1` | `alpha = R_event - hold_baseline` | +//! | `WR_EMA` | ISV[512] | `is_close == 1` | `is_win` ∈ {0.0, 1.0} | +//! | `TRADE_DUR` | internal[0]| `is_close == 1` | `trade_duration` (bars) | +//! | `HOLD_PCT_EMA` | ISV[515] | per-step | `(action_is_hold == 1) ? 1 : 0` | +//! | `HOLD_REWARD_EMA`| ISV[516] | `action_is_hold == 1` | `r_per_bar = -aux_conf * cost_scale` | +//! | `AUX_P50` | internal[1]| per-step | `aux_logits_p50` (Phase 1.1 output) | +//! | `AUX_STD` | internal[2]| per-step | `aux_logits_std` (Phase 1.1 output) | +//! | `AUX_DIR_ACC` | internal[3]| per-step | `aux_dir_acc` ∈ {0.0, 1.0} | +//! +//! Per `pearl_first_observation_bootstrap`: the kernel tracks bootstrap +//! state via a per-EMA observation counter (i32) in `obs_count[]`. The +//! sentinel is `count == 0`, NOT `value == 0` — because some EMAs +//! legitimately observe 0 (e.g., consecutive losses ⇒ WR_EMA stays +//! 0). The counter atomically transitions 0 → 1 on the first firing +//! observation; subsequent observations Wiener-blend with the +//! [`WIENER_ALPHA_FLOOR`] = 0.4 floor per +//! `pearl_wiener_alpha_floor_for_nonstationary`. +//! +//! ## Hard rules covered +//! +//! - `pearl_first_observation_bootstrap` — counter-based bootstrap. +//! - `pearl_wiener_alpha_floor_for_nonstationary` — α = 0.4 floor. +//! - `feedback_isv_for_adaptive_bounds` — α floor is structural +//! (NOT a tuned constant — it's the canonical floor from the pearl). +//! - `feedback_no_atomicadd` — single-thread kernel, no concurrent +//! writes anywhere. +//! - `feedback_no_cpu_compute_strict` — every blend lives in the +//! kernel; the launcher only packs scalar inputs into kernel args. +//! - `feedback_no_htod_htoh_only_mapped_pinned` — both ISV and the +//! `internal` / `obs_count` buffers are mapped-pinned; kernel emits +//! `__threadfence_system()` for PCIe-visible coherence. +//! - `pearl_no_host_branches_in_captured_graph` — single-block, +//! single-thread kernel; safe to capture in the per-step CUDA +//! Graph. +//! +//! ## Phase 1.2 wiring scope +//! +//! Kernel + launcher + tests + build entry **only**. The training-loop +//! call site lands in Phase 1.4 atomically with the rest of the SP20 +//! reward chain (Kernel 2 = controllers, Kernel 3 = stats already +//! landed in Phase 1.1) per `feedback_no_partial_refactor`. + +/// Wiener-α floor per `pearl_wiener_alpha_floor_for_nonstationary`. NOT +/// a tuned constant — structural floor that preserves catch-up +/// bandwidth for control loops where the target drifts via the +/// co-adapting policy (the SP-chain's defining co-adaptation regime). +/// Per `feedback_isv_for_adaptive_bounds`: only adaptive bounds live in +/// ISV; structural floors are compile-time. +pub const WIENER_ALPHA_FLOOR: f32 = 0.4; + +/// Number of internal scratch slots maintained by this kernel (4 +/// f32s). The slot indices are exported as `EMA_INTERNAL_*` constants +/// below; callers MUST allocate a `MappedF32Buffer` of exactly this +/// length and pass its `dev_ptr` to the kernel. +pub const EMA_INTERNAL_SLOT_COUNT: usize = 4; + +/// Internal scratch slot: `trade_duration_ema` (bars). Updated when +/// `is_close == 1`. Phase 1.3's `sp20_controllers_compute` reads this +/// to derive `N_STEP_INDEX = clamp(round(trade_duration_ema), 1, 30)`. +pub const EMA_INTERNAL_TRADE_DURATION: usize = 0; + +/// Internal scratch slot: `aux_conf_p50_ema`. Wiener-α blend of the +/// upstream `sp20_stats_compute` output `aux_logits_p50`. Phase 1.3 +/// reads this to derive +/// `TARGET_HOLD_PCT = clamp(0.8 - aux_conf_p50_ema * 1.5, 0.1, 0.8)`. +pub const EMA_INTERNAL_AUX_P50: usize = 1; + +/// Internal scratch slot: `aux_conf_std_ema`. Wiener-α blend of the +/// upstream `sp20_stats_compute` output `aux_logits_std`. Phase 1.3 +/// reads this to derive `AUX_GATE_TEMP = aux_conf_std_ema` (with +/// floor). +pub const EMA_INTERNAL_AUX_STD: usize = 2; + +/// Internal scratch slot: `aux_dir_acc_ema`. Per-step blend of the +/// indicator "did the aux head predict the right direction at this +/// step" (1.0 if right, 0.0 if wrong). Phase 1.3 reads this to derive +/// `AUX_CONF_THRESHOLD = clamp(aux_dir_acc_ema - 0.50, 0.01, 0.20)`. +pub const EMA_INTERNAL_AUX_DIR_ACC: usize = 3; + +/// Number of per-EMA observation counter slots. MUST stay 8 — one +/// counter per EMA (4 ISV + 4 internal). Layout fixed by the +/// `OBS_COUNT_*` constants below. +pub const OBS_COUNT_SLOTS: usize = 8; + +/// Observation-counter slot: ALPHA_EMA. MUST match the kernel's +/// `OBS_IDX_ALPHA` macro. +pub const OBS_COUNT_ALPHA: usize = 0; +/// Observation-counter slot: WR_EMA. +pub const OBS_COUNT_WR: usize = 1; +/// Observation-counter slot: TRADE_DURATION_EMA (internal). +pub const OBS_COUNT_TRADE_DURATION: usize = 2; +/// Observation-counter slot: HOLD_PCT_EMA. +pub const OBS_COUNT_HOLD_PCT: usize = 3; +/// Observation-counter slot: HOLD_REWARD_EMA. +pub const OBS_COUNT_HOLD_REWARD: usize = 4; +/// Observation-counter slot: aux_conf_p50_ema (internal). +pub const OBS_COUNT_AUX_P50: usize = 5; +/// Observation-counter slot: aux_conf_std_ema (internal). +pub const OBS_COUNT_AUX_STD: usize = 6; +/// Observation-counter slot: aux_dir_acc_ema (internal). +pub const OBS_COUNT_AUX_DIR_ACC: usize = 7; + +/// Per-step inputs to `launch_sp20_emas_compute`. Packs the 9 scalar +/// fields the kernel needs into a single struct so the launcher +/// signature stays manageable and the call sites are +/// readable. +/// +/// Field semantics: +/// - `is_close`: 1 if a trade closed this step, 0 otherwise. Gates +/// the per-trade EMAs (WR, ALPHA, TRADE_DURATION). +/// - `is_win`: 1.0 if the closed trade was a win, 0.0 if loss. +/// Only meaningful when `is_close == 1`. +/// - `alpha`: `R_event - hold_baseline` from the SP20 spec +/// Component 1 reward kernel. Only meaningful when +/// `is_close == 1`. +/// - `trade_duration`: bars in the closed trade. Only meaningful +/// when `is_close == 1`. +/// - `action_is_hold`: 1 if the action this step was Hold, 0 +/// otherwise. Gates the HOLD_REWARD_EMA; also drives the +/// HOLD_PCT_EMA observation. +/// - `r_per_bar`: per-bar Hold reward = `-aux_conf * cost_scale`. +/// Only meaningful when `action_is_hold == 1`. +/// - `aux_p50_in`: per-step `aux_logits_p50` from Phase 1.1's +/// `sp20_stats_compute` (mapped-pinned `[2]` output, slot 0). +/// - `aux_std_in`: per-step `aux_logits_std` from Phase 1.1 +/// (slot 1 of the same output). +/// - `aux_dir_acc`: 1.0 if the aux head predicted the right +/// direction at this step, 0.0 otherwise. +#[derive(Copy, Clone, Debug)] +pub struct EmaInputs { + pub is_close: i32, + pub is_win: f32, + pub alpha: f32, + pub trade_duration: f32, + pub action_is_hold: i32, + pub r_per_bar: f32, + pub aux_p50_in: f32, + pub aux_std_in: f32, + pub aux_dir_acc: f32, +} + +/// Launch `sp20_emas_compute_kernel` on the producer's stream. +/// +/// Single-block, single-thread launch. `kernel` MUST be the loaded +/// `sp20_emas_compute_kernel` `CudaFunction` (from the cubin emitted +/// by `build.rs` — see manifest entry `sp20_emas_compute_kernel.cu`). +/// +/// Buffer pointers MUST point to mapped-pinned device-visible +/// allocations: +/// - `isv_dev`: ISV bus, at least +/// [`crate::cuda_pipeline::sp14_isv_slots::SP20_ISV_SLOT_END`] f32s. +/// - `internal_dev`: [`EMA_INTERNAL_SLOT_COUNT`]-long f32 scratch. +/// - `obs_count_dev`: [`OBS_COUNT_SLOTS`]-long i32 counters. +/// +/// `inputs` is the per-step gather of all 9 scalar EMA observations; +/// see [`EmaInputs`]. +/// +/// # Safety +/// +/// All three buffer pointers MUST be valid mapped-pinned device +/// pointers obtained via `MappedF32Buffer::dev_ptr` / +/// `MappedI32Buffer::dev_ptr`. The caller MUST issue this launch on +/// the SAME stream as the upstream `sp20_stats_compute` (so the +/// stream-ordering invariant guarantees `aux_p50_in` / `aux_std_in` +/// reads see the latest stats output). Stream ordering with the +/// trade-physics kernels that produce `is_close` / `is_win` / etc. is +/// the caller's responsibility — Phase 1.4 lands the production wire +/// site with that ordering enforced. +pub unsafe fn launch_sp20_emas_compute( + stream: &cudarc::driver::CudaStream, + kernel: &cudarc::driver::CudaFunction, + isv_dev: u64, + internal_dev: u64, + obs_count_dev: u64, + inputs: EmaInputs, +) -> Result<(), crate::MLError> { + use cudarc::driver::{LaunchConfig, PushKernelArg}; + + debug_assert!(isv_dev != 0, + "launch_sp20_emas_compute: isv_dev must be a valid device pointer"); + debug_assert!(internal_dev != 0, + "launch_sp20_emas_compute: internal_dev must be a valid device pointer"); + debug_assert!(obs_count_dev != 0, + "launch_sp20_emas_compute: obs_count_dev must be a valid device pointer"); + debug_assert!(inputs.is_close == 0 || inputs.is_close == 1, + "launch_sp20_emas_compute: is_close must be 0 or 1, got {}", inputs.is_close); + debug_assert!(inputs.action_is_hold == 0 || inputs.action_is_hold == 1, + "launch_sp20_emas_compute: action_is_hold must be 0 or 1, got {}", + inputs.action_is_hold); + + // Slot indices passed in as kernel args so the kernel stays + // decoupled from SP14 slot-number drift (the contract is enforced + // by `sp14_isv_slots.rs` slot-locked tests). + let isv_idx_alpha: i32 = crate::cuda_pipeline::sp14_isv_slots::ALPHA_EMA_INDEX as i32; + let isv_idx_wr: i32 = crate::cuda_pipeline::sp14_isv_slots::WR_EMA_INDEX as i32; + let isv_idx_hold_pct: i32 = crate::cuda_pipeline::sp14_isv_slots::HOLD_PCT_EMA_INDEX as i32; + let isv_idx_hold_reward: i32 = crate::cuda_pipeline::sp14_isv_slots::HOLD_REWARD_EMA_INDEX as i32; + + let cfg = LaunchConfig { + grid_dim: (1, 1, 1), + block_dim: (1, 1, 1), + shared_mem_bytes: 0, + }; + + stream + .launch_builder(kernel) + .arg(&isv_dev) + .arg(&internal_dev) + .arg(&obs_count_dev) + .arg(&isv_idx_alpha) + .arg(&isv_idx_wr) + .arg(&isv_idx_hold_pct) + .arg(&isv_idx_hold_reward) + .arg(&inputs.is_close) + .arg(&inputs.is_win) + .arg(&inputs.alpha) + .arg(&inputs.trade_duration) + .arg(&inputs.action_is_hold) + .arg(&inputs.r_per_bar) + .arg(&inputs.aux_p50_in) + .arg(&inputs.aux_std_in) + .arg(&inputs.aux_dir_acc) + .launch(cfg) + .map_err(|e| { + crate::MLError::ModelError(format!("sp20_emas_compute_kernel launch: {e}")) + })?; + Ok(()) +} + +#[cfg(test)] +mod tests { + use super::*; + + #[test] + fn wiener_alpha_floor_locked() { + // Locks the canonical Wiener-α floor against silent drift. + // 0.4 is the structural floor from + // `pearl_wiener_alpha_floor_for_nonstationary`. + assert!((WIENER_ALPHA_FLOOR - 0.4_f32).abs() < f32::EPSILON); + } + + #[test] + fn internal_slot_count_locks_to_4() { + // Locks the internal scratch buffer length against silent + // changes — the kernel's `INT_IDX_*` macros and the + // `EMA_INTERNAL_*` constants below MUST stay aligned. + assert_eq!(EMA_INTERNAL_SLOT_COUNT, 4); + // Slot indices are 0..4 and unique: + assert_eq!(EMA_INTERNAL_TRADE_DURATION, 0); + assert_eq!(EMA_INTERNAL_AUX_P50, 1); + assert_eq!(EMA_INTERNAL_AUX_STD, 2); + assert_eq!(EMA_INTERNAL_AUX_DIR_ACC, 3); + } + + #[test] + fn obs_count_slot_count_locks_to_8() { + // 4 ISV EMAs + 4 internal EMAs = 8 counters. + assert_eq!(OBS_COUNT_SLOTS, 8); + // Layout matches the kernel's `OBS_IDX_*` macros — locked + // explicitly so the .cu / .rs constants stay in sync. + assert_eq!(OBS_COUNT_ALPHA, 0); + assert_eq!(OBS_COUNT_WR, 1); + assert_eq!(OBS_COUNT_TRADE_DURATION, 2); + assert_eq!(OBS_COUNT_HOLD_PCT, 3); + assert_eq!(OBS_COUNT_HOLD_REWARD, 4); + assert_eq!(OBS_COUNT_AUX_P50, 5); + assert_eq!(OBS_COUNT_AUX_STD, 6); + assert_eq!(OBS_COUNT_AUX_DIR_ACC, 7); + } + + #[test] + fn ema_inputs_default_is_safe_sentinel() { + // A zero-initialised `EmaInputs` has all gates closed + // (is_close = 0, action_is_hold = 0) and all observations at + // 0 — equivalent to "no firing this step" for all 8 EMAs. + // This is the canonical pre-warmup / cold-start input shape. + let inp = EmaInputs { + is_close: 0, + is_win: 0.0, + alpha: 0.0, + trade_duration: 0.0, + action_is_hold: 0, + r_per_bar: 0.0, + aux_p50_in: 0.0, + aux_std_in: 0.0, + aux_dir_acc: 0.0, + }; + // Per-step EMAs still fire on a zero-input call, but the + // observation is 0 so the EMA Wiener-blends toward 0 — that + // is correct behavior for cold-start, and is verified by the + // `per_step_emas_fire_unconditionally` GPU oracle test. + assert_eq!(inp.is_close, 0); + assert_eq!(inp.action_is_hold, 0); + } +} diff --git a/crates/ml/src/cuda_pipeline/sp20_emas_compute_kernel.cu b/crates/ml/src/cuda_pipeline/sp20_emas_compute_kernel.cu new file mode 100644 index 000000000..1ef1369b1 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/sp20_emas_compute_kernel.cu @@ -0,0 +1,222 @@ +/* ══════════════════════════════════════════════════════════════════════════ + * SP20 Phase 1.2 (2026-05-09) — sp20_emas_compute kernel. + * + * Component 5 / Kernel 1 of the SP20 design — the central state-tracker + * that updates 8 Wiener-α EMAs per training step. Reads per-step / per- + * trade observations + the upstream `sp20_stats_compute` (Phase 1.1) + * outputs, then writes 4 ISV slots (511, 512, 515, 516) plus 4 + * internal scratch slots that Phase 1.3's `sp20_controllers_compute` + * consumes to derive `LOSS_CAP`, `HOLD_COST_SCALE`, `TARGET_HOLD_PCT`, + * `N_STEP`, `AUX_CONF_THRESHOLD`, `AUX_GATE_TEMP`. + * + * EMAs maintained (8 total — see launcher header for slot indices): + * + * Per-trade (gated on is_close == 1): + * - WR_EMA → ISV[WR_EMA_INDEX = 512] blend over is_win + * - ALPHA_EMA → ISV[ALPHA_EMA_INDEX = 511] blend over alpha + * - TRADE_DUR → internal[0] blend over trade_duration + * + * Per-step (always fire): + * - HOLD_PCT → ISV[HOLD_PCT_EMA_INDEX = 515] blend over (action==Hold) + * - AUX_P50 → internal[1] blend over aux_p50_in + * - AUX_STD → internal[2] blend over aux_std_in + * - AUX_DIR_ACC → internal[3] blend over aux_dir_acc_in + * + * Per-step (gated on action_is_hold == 1): + * - HOLD_REWARD → ISV[HOLD_REWARD_EMA_INDEX = 516] blend over r_per_bar + * + * ── Pearls + invariants ──────────────────────────────────────────────── + * - `pearl_first_observation_bootstrap` — sentinel = 0; first observation + * REPLACES directly. **Implementation note**: a sentinel value of 0.0 + * is a *plausible legitimate observation* for some EMAs (e.g., + * consecutive losses ⇒ WR_EMA legitimately stays 0), so this kernel + * tracks bootstrap state via a per-EMA observation counter (i32) in + * a separate mapped-pinned `obs_count` buffer, NOT via the sentinel- + * detect-via-current-value pattern in `sp17_a_var_ema_kernel.cu`. The + * counter atomically transitions 0 → 1 on the first firing observation + * of each EMA; `count == 0` ⇒ replace, `count > 0` ⇒ Wiener-blend. + * - `pearl_wiener_alpha_floor_for_nonstationary` — α = WIENER_ALPHA_FLOOR + * (= 0.4f) hardcoded as a structural property; this is NOT a tuned + * constant per `feedback_isv_for_adaptive_bounds` — it's the steady- + * state floor that preserves catch-up bandwidth when the training + * target drifts via co-adapting policy (the SP-chain's defining + * co-adaptation regime). + * - `pearl_no_host_branches_in_captured_graph` — single-block, single- + * thread kernel; no host branches; safe to capture in the per-step + * CUDA Graph. + * - `feedback_no_atomicadd` — single-thread kernel; no concurrent + * writes anywhere, atomicAdd never used. + * - `feedback_no_htod_htoh_only_mapped_pinned` — both ISV (mapped- + * pinned) and `internal` / `obs_count` (this commit's two new + * mapped-pinned buffers) are device-visible via DEVICEMAP; kernel + * emits `__threadfence_system()` after the writes for PCIe-visible + * coherence. + * - `feedback_no_cpu_compute_strict` — every blend lives in this + * kernel; the launcher's only host work is to package the per-step + * scalar inputs into kernel arguments. + * - `feedback_no_partial_refactor` — Phase 1.2 lands kernel + launcher + * + tests + build entry atomically. The production wire-up (Phase + * 1.4) consumes the launcher and the audit doc updates per + * Invariant 7. + * + * Algorithm (single block, single thread): + * + * For each EMA in [WR, ALPHA, TRADE_DUR, HOLD_PCT, HOLD_REWARD, + * AUX_P50, AUX_STD, AUX_DIR_ACC]: + * if (gate is OFF for this EMA this step) continue; + * int c = obs_count[ema_idx]; + * float prev = (target buf)[slot_idx]; + * float next; + * if (c == 0) { + * next = observation; // Pearl-A bootstrap + * } else { + * next = (1 - α) * prev + α * observation; // Wiener floor + * } + * (target buf)[slot_idx] = next; + * obs_count[ema_idx] = c + 1; + * + * Launch contract: + * grid_dim = (1, 1, 1) + * block_dim = (1, 1, 1) + * shared mem = 0 + * + * Args: + * isv — [ISV_TOTAL_DIM] f32 mapped-pinned dev ptr (slots + * 511, 512, 515, 516 written). + * internal — [4] f32 mapped-pinned dev ptr (the 4 internal + * scratch slots: TRADE_DUR, AUX_P50, AUX_STD, + * AUX_DIR_ACC). Layout matches launcher constants. + * obs_count — [8] i32 mapped-pinned dev ptr (per-EMA observation + * counters; layout matches launcher OBS_COUNT_* + * constants). + * isv_idx_alpha, isv_idx_wr, isv_idx_hold_pct, isv_idx_hold_reward + * — ISV slot indices for the 4 ISV EMAs (caller passes + * the SP14 constants so the kernel stays decoupled + * from slot-number drift). + * is_close, is_win, alpha, trade_duration, action_is_hold, + * r_per_bar, aux_p50_in, aux_std_in, aux_dir_acc_in — per-step + * inputs (see EMA list above for which fields gate + * which EMA). + * ══════════════════════════════════════════════════════════════════════════ */ + +#include + +/* Per-EMA slot indices in the obs_count[] array. MUST match the + * launcher's `OBS_COUNT_*` constants. The order is fixed by the + * launch contract, not by physical ordering of the EMAs in the + * algorithm — sticking to launcher order keeps reading easier. */ +#define OBS_IDX_ALPHA 0 +#define OBS_IDX_WR 1 +#define OBS_IDX_TRADE_DUR 2 +#define OBS_IDX_HOLD_PCT 3 +#define OBS_IDX_HOLD_REWARD 4 +#define OBS_IDX_AUX_P50 5 +#define OBS_IDX_AUX_STD 6 +#define OBS_IDX_AUX_DIR_ACC 7 + +/* Internal scratch slot indices. MUST match the launcher's + * `EMA_INTERNAL_*` constants. */ +#define INT_IDX_TRADE_DUR 0 +#define INT_IDX_AUX_P50 1 +#define INT_IDX_AUX_STD 2 +#define INT_IDX_AUX_DIR_ACC 3 + +/* Wiener-α floor per `pearl_wiener_alpha_floor_for_nonstationary`. NOT a + * tuned constant — structural floor that preserves catch-up bandwidth + * for control loops where the target drifts via the co-adapting policy. + * Hardcoded per `feedback_isv_for_adaptive_bounds` rationale: structural + * properties are compile-time, only adaptive bounds live in ISV. */ +#define WIENER_ALPHA_FLOOR_F 0.4f + +/* Pearl-A bootstrap helper: returns the new EMA value and increments + * the observation counter. */ +__device__ __forceinline__ +float ema_step(float prev, float observation, int* counter) +{ + const int c = *counter; + float next; + if (c == 0) { + next = observation; /* Pearl-A bootstrap */ + } else { + next = (1.0f - WIENER_ALPHA_FLOOR_F) * prev + + WIENER_ALPHA_FLOOR_F * observation; + } + *counter = c + 1; + return next; +} + +extern "C" __global__ void sp20_emas_compute_kernel( + float* __restrict__ isv, /* [ISV_TOTAL_DIM] */ + float* __restrict__ internal, /* [4] f32 internal scratch */ + int* __restrict__ obs_count, /* [8] i32 per-EMA counters */ + int isv_idx_alpha, + int isv_idx_wr, + int isv_idx_hold_pct, + int isv_idx_hold_reward, + int is_close, + float is_win, + float alpha, + float trade_duration, + int action_is_hold, + float r_per_bar, + float aux_p50_in, + float aux_std_in, + float aux_dir_acc_in) +{ + /* Single-block, single-thread contract. */ + if (blockIdx.x != 0) return; + if (threadIdx.x != 0) return; + + /* ── Per-trade EMAs: gated on is_close == 1 ───────────────────── */ + if (is_close == 1) { + /* WR_EMA */ + isv[isv_idx_wr] = ema_step(isv[isv_idx_wr], is_win, + &obs_count[OBS_IDX_WR]); + + /* ALPHA_EMA */ + isv[isv_idx_alpha] = ema_step(isv[isv_idx_alpha], alpha, + &obs_count[OBS_IDX_ALPHA]); + + /* TRADE_DURATION_EMA — internal scratch */ + internal[INT_IDX_TRADE_DUR] = + ema_step(internal[INT_IDX_TRADE_DUR], trade_duration, + &obs_count[OBS_IDX_TRADE_DUR]); + } + + /* ── Per-step EMAs: always fire ───────────────────────────────── */ + + /* HOLD_PCT_EMA — observation = (action_is_hold == 1) ? 1.0 : 0.0 */ + { + const float hold_indicator = (action_is_hold == 1) ? 1.0f : 0.0f; + isv[isv_idx_hold_pct] = ema_step(isv[isv_idx_hold_pct], + hold_indicator, + &obs_count[OBS_IDX_HOLD_PCT]); + } + + /* aux_conf_p50_ema — internal scratch */ + internal[INT_IDX_AUX_P50] = + ema_step(internal[INT_IDX_AUX_P50], aux_p50_in, + &obs_count[OBS_IDX_AUX_P50]); + + /* aux_conf_std_ema — internal scratch */ + internal[INT_IDX_AUX_STD] = + ema_step(internal[INT_IDX_AUX_STD], aux_std_in, + &obs_count[OBS_IDX_AUX_STD]); + + /* aux_dir_acc_ema — internal scratch */ + internal[INT_IDX_AUX_DIR_ACC] = + ema_step(internal[INT_IDX_AUX_DIR_ACC], aux_dir_acc_in, + &obs_count[OBS_IDX_AUX_DIR_ACC]); + + /* ── Hold-reward EMA: gated on action_is_hold == 1 ────────────── */ + if (action_is_hold == 1) { + isv[isv_idx_hold_reward] = + ema_step(isv[isv_idx_hold_reward], r_per_bar, + &obs_count[OBS_IDX_HOLD_REWARD]); + } + + /* PCIe-visible writes for mapped-pinned buffers (ISV + internal + + * obs_count). Single fence covers all writes — they all land in + * the same memory region from the device's view. */ + __threadfence_system(); +} diff --git a/crates/ml/tests/sp20_emas_compute_test.rs b/crates/ml/tests/sp20_emas_compute_test.rs new file mode 100644 index 000000000..88fded482 --- /dev/null +++ b/crates/ml/tests/sp20_emas_compute_test.rs @@ -0,0 +1,347 @@ +//! SP20 Phase 1.2 (2026-05-09) — `sp20_emas_compute_kernel` GPU oracle tests. +//! +//! Verifies the 8 Wiener-α EMA producer kernel (Component 5 / Kernel 1 +//! of the SP20 design) updates the ISV slots [511, 512, 515, 516] and +//! 4 internal scratch slots correctly under the +//! `pearl_first_observation_bootstrap` + `pearl_wiener_alpha_floor_for_ +//! nonstationary` discipline. +//! +//! Tests cover: +//! +//! 1. **First-observation bootstrap** — sentinel (count==0) replaces +//! directly without blend, even when the observed value happens +//! to be 0 (e.g., first-loss WR=0). Exercises the per-EMA +//! observation counter mechanism. +//! 2. **Wiener-α convergence with alternating wins** — 100 trade- +//! close events with alternating is_win 1/0/1/0 land WR_EMA +//! within tolerance of 0.5 (the long-run mean of the indicator). +//! 3. **Hold-reward gating** — `HOLD_REWARD_EMA` only updates on +//! Hold-state bars (`action_is_hold == 1`); non-Hold steps leave +//! it at sentinel. +//! 4. **Per-step EMAs always fire** — `HOLD_PCT_EMA`, the three +//! internal aux EMAs (p50, std, dir_acc) update every step +//! regardless of `is_close` / `action_is_hold` gates. +//! +//! Run on a GPU host: +//! +//! SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 \ +//! cargo test -p ml --test sp20_emas_compute_test --features cuda \ +//! -- --ignored --nocapture + +#![allow(clippy::tests_outside_test_module)] + +#[cfg(feature = "cuda")] +#[allow(unsafe_code)] // CUDA kernel launch + mapped-pinned memory. +mod gpu { + use std::sync::Arc; + + use cudarc::driver::{CudaContext, CudaFunction, CudaStream}; + use ml::cuda_pipeline::mapped_pinned::{MappedF32Buffer, MappedI32Buffer}; + use ml::cuda_pipeline::sp14_isv_slots::{ + ALPHA_EMA_INDEX, HOLD_PCT_EMA_INDEX, HOLD_REWARD_EMA_INDEX, WR_EMA_INDEX, + }; + use ml::cuda_pipeline::sp20_emas_compute::{ + launch_sp20_emas_compute, EmaInputs, EMA_INTERNAL_AUX_DIR_ACC, + EMA_INTERNAL_AUX_P50, EMA_INTERNAL_AUX_STD, EMA_INTERNAL_TRADE_DURATION, + EMA_INTERNAL_SLOT_COUNT, OBS_COUNT_SLOTS, WIENER_ALPHA_FLOOR, + }; + + /// Total ISV slot count is the trainer's contract; for tests we + /// allocate just enough to cover the highest SP20 EMA slot (516). + const ISV_TEST_LEN: usize = 520; + + const SP20_EMAS_COMPUTE_CUBIN: &[u8] = include_bytes!(concat!( + env!("OUT_DIR"), + "/sp20_emas_compute_kernel.cubin" + )); + + fn make_test_stream() -> Arc { + let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?"); + ctx.default_stream() + } + + fn load_sp20_emas_compute(stream: &Arc) -> CudaFunction { + let module = stream + .context() + .load_cubin(SP20_EMAS_COMPUTE_CUBIN.to_vec()) + .expect("load sp20_emas_compute_kernel cubin"); + module + .load_function("sp20_emas_compute_kernel") + .expect("load sp20_emas_compute_kernel function") + } + + /// Test harness: allocates ISV + internal + obs_count buffers, then + /// fires `inputs.len()` sequential kernel invocations on the same + /// stream. Returns final ISV + internal state for assertions. + fn run_kernel_sequence(inputs: &[EmaInputs]) -> (Vec, Vec, Vec) { + let stream = make_test_stream(); + let kernel = load_sp20_emas_compute(&stream); + + let isv = unsafe { MappedF32Buffer::new(ISV_TEST_LEN) }.expect("alloc isv"); + isv.write_from_slice(&vec![0.0_f32; ISV_TEST_LEN]); + + let internal = unsafe { MappedF32Buffer::new(EMA_INTERNAL_SLOT_COUNT) } + .expect("alloc internal"); + internal.write_from_slice(&vec![0.0_f32; EMA_INTERNAL_SLOT_COUNT]); + + let obs_count = + unsafe { MappedI32Buffer::new(OBS_COUNT_SLOTS) }.expect("alloc obs_count"); + obs_count.write_from_slice(&vec![0_i32; OBS_COUNT_SLOTS]); + + for inp in inputs { + unsafe { + launch_sp20_emas_compute( + &stream, + &kernel, + isv.dev_ptr, + internal.dev_ptr, + obs_count.dev_ptr, + *inp, + ) + .expect("launch sp20_emas_compute_kernel"); + } + } + stream.synchronize().expect("sync after sp20_emas_compute"); + + (isv.read_all(), internal.read_all(), obs_count.read_all()) + } + + /// Test 1 — first-observation bootstrap on a trade-close event with + /// `is_win=1, alpha=0.5`. Per `pearl_first_observation_bootstrap`, + /// the sentinel (count==0) MUST be replaced directly without blend. + /// WR_EMA → 1.0, ALPHA_EMA → 0.5 (not 0.4 × 1 or 0.4 × 0.5 from + /// any blend formula). + #[test] + #[ignore = "requires GPU"] + fn first_observation_replaces_sentinel() { + let inp = EmaInputs { + is_close: 1, + is_win: 1.0, + alpha: 0.5, + trade_duration: 7.0, + action_is_hold: 0, + r_per_bar: 0.0, + aux_p50_in: 0.0, + aux_std_in: 0.0, + aux_dir_acc: 0.0, + }; + let (isv, internal, obs_count) = run_kernel_sequence(&[inp]); + + // Per-trade EMAs fired by is_close=1. + assert_eq!( + isv[WR_EMA_INDEX], 1.0, + "WR_EMA: first observation must REPLACE sentinel directly (got {})", + isv[WR_EMA_INDEX] + ); + assert_eq!( + isv[ALPHA_EMA_INDEX], 0.5, + "ALPHA_EMA: first observation must REPLACE sentinel directly (got {})", + isv[ALPHA_EMA_INDEX] + ); + assert_eq!( + internal[EMA_INTERNAL_TRADE_DURATION], 7.0, + "trade_duration_ema: first observation must REPLACE sentinel directly", + ); + + // Per-step EMAs that fired (action_is_hold=0 ⇒ HOLD_PCT first + // observation = 0.0 directly, so we read it as 0). + assert_eq!( + isv[HOLD_PCT_EMA_INDEX], 0.0, + "HOLD_PCT_EMA: first observation = 0.0 (action_is_hold=0)", + ); + + // HOLD_REWARD_EMA was NOT fired (action_is_hold=0) — sentinel + // remains at 0.0 with obs_count == 0. + assert_eq!( + isv[HOLD_REWARD_EMA_INDEX], 0.0, + "HOLD_REWARD_EMA must remain at sentinel when action_is_hold=0", + ); + + // Observation counters: WR, ALPHA, TRADE_DURATION, HOLD_PCT, + // AUX_P50, AUX_STD, AUX_DIR_ACC fired = 1; HOLD_REWARD did not. + // The counter layout is fixed by the launcher's contract; verify + // HOLD_REWARD's slot stayed at 0. + // (Slot indices come from launcher constants; we just check a + // non-zero count for at least one fired EMA and zero for HOLD_REWARD.) + let hold_reward_count = obs_count[4]; // HOLD_REWARD slot + let wr_count = obs_count[1]; + assert_eq!(hold_reward_count, 0, "HOLD_REWARD obs_count must stay 0"); + assert_eq!(wr_count, 1, "WR obs_count must be 1 after one trade"); + } + + /// Test 2 — Wiener-α convergence with alternating wins. + /// + /// Fire 100 trade-close events with alternating `is_win` 1/0/1/0... + /// Expected: WR_EMA → 0.5 (long-run mean of the indicator). + /// Tolerance: at α=0.4 the EMA tracks within ≈ √(α/(2-α)) of the + /// mean for an alternating signal — about 0.5 amplitude × √0.25 ≈ + /// 0.25, so we widen to 0.30 to absorb the alternation noise of + /// the last few samples. + /// + /// First observation is a win (is_win=1) → WR_EMA = 1.0 directly, + /// then alternation drives it to ≈ 0.5 long-run. + #[test] + #[ignore = "requires GPU"] + fn wiener_alpha_converges_to_long_run_mean() { + let inputs: Vec = (0..100) + .map(|i| EmaInputs { + is_close: 1, + is_win: if i % 2 == 0 { 1.0 } else { 0.0 }, + alpha: 0.0, + trade_duration: 1.0, + action_is_hold: 0, + r_per_bar: 0.0, + aux_p50_in: 0.0, + aux_std_in: 0.0, + aux_dir_acc: 0.0, + }) + .collect(); + let (isv, _internal, _obs_count) = run_kernel_sequence(&inputs); + + let wr_ema = isv[WR_EMA_INDEX]; + let amplitude = 0.30_f32; // generous bound, absorbs alternation + assert!( + (wr_ema - 0.5).abs() < amplitude, + "WR_EMA after 100 alternating events should be near 0.5; got {wr_ema}", + ); + } + + /// Test 3 — HOLD_REWARD_EMA only updates on Hold-state bars. + /// + /// Fire 5 non-Hold bars (action_is_hold=0, r_per_bar=−0.05) then + /// 5 Hold bars (action_is_hold=1, r_per_bar=−0.05). Expected: + /// after all 10 calls, HOLD_REWARD_EMA == −0.05 (only the 5 Hold + /// bars updated; the first Hold bar bootstrapped to −0.05, and + /// subsequent identical observations leave it at −0.05). + /// + /// Per the spec: "Update only on Hold-state bars (gate on + /// action_is_hold == 1)." + #[test] + #[ignore = "requires GPU"] + fn hold_reward_ema_gated_on_hold_bars() { + let mut inputs: Vec = Vec::new(); + for _ in 0..5 { + inputs.push(EmaInputs { + is_close: 0, + is_win: 0.0, + alpha: 0.0, + trade_duration: 0.0, + action_is_hold: 0, // gated OFF + r_per_bar: -0.05, + aux_p50_in: 0.0, + aux_std_in: 0.0, + aux_dir_acc: 0.0, + }); + } + for _ in 0..5 { + inputs.push(EmaInputs { + is_close: 0, + is_win: 0.0, + alpha: 0.0, + trade_duration: 0.0, + action_is_hold: 1, // gate ON + r_per_bar: -0.05, + aux_p50_in: 0.0, + aux_std_in: 0.0, + aux_dir_acc: 0.0, + }); + } + let (isv, _internal, obs_count) = run_kernel_sequence(&inputs); + + // After 5 non-Hold bars the EMA is at sentinel (0.0) and + // counter is 0; the first Hold bar (step 6) bootstraps to + // −0.05; the next 4 Hold bars Wiener-blend toward −0.05 (all + // identical) — the EMA stays at −0.05 to fp32 precision. + let hold_reward = isv[HOLD_REWARD_EMA_INDEX]; + let tol = 1e-6_f32; + assert!( + (hold_reward - (-0.05)).abs() < tol, + "HOLD_REWARD_EMA: after 5 non-Hold + 5 Hold bars (all r_per_bar=−0.05) \ + expected −0.05, got {hold_reward}", + ); + + // Counter must reflect ONLY the 5 Hold bars. + let hold_reward_count = obs_count[4]; + assert_eq!( + hold_reward_count, 5, + "HOLD_REWARD obs_count must equal Hold-bar count (5), got {hold_reward_count}", + ); + } + + /// Test 4 — per-step always-fire EMAs (HOLD_PCT, aux_p50, aux_std, + /// aux_dir_acc) update every step regardless of trade-close / + /// hold-state gates. + /// + /// Fire 10 steps with `is_close=0, action_is_hold=0`, so the + /// per-trade and Hold-reward EMAs never fire, but the per-step + /// EMAs receive consistent inputs (aux_p50=0.2, aux_std=0.1, + /// aux_dir_acc=1.0) and must converge there. + /// + /// First observation bootstraps directly (count==0); subsequent + /// 9 observations Wiener-blend toward the same value. With + /// identical inputs after bootstrap, the EMA stays put — the + /// blend `(1-α)·prev + α·curr = curr` when prev == curr. + #[test] + #[ignore = "requires GPU"] + fn per_step_emas_fire_unconditionally() { + let inputs: Vec = (0..10) + .map(|_| EmaInputs { + is_close: 0, + is_win: 0.0, + alpha: 0.0, + trade_duration: 0.0, + action_is_hold: 0, + r_per_bar: 0.0, + aux_p50_in: 0.2, + aux_std_in: 0.1, + aux_dir_acc: 1.0, + }) + .collect(); + let (isv, internal, obs_count) = run_kernel_sequence(&inputs); + + let tol = 1e-6_f32; + assert!( + (isv[HOLD_PCT_EMA_INDEX] - 0.0).abs() < tol, + "HOLD_PCT_EMA: 10 non-Hold steps should land at 0.0, got {}", + isv[HOLD_PCT_EMA_INDEX] + ); + assert!( + (internal[EMA_INTERNAL_AUX_P50] - 0.2).abs() < tol, + "aux_conf_p50_ema: 10 identical steps should land at 0.2, got {}", + internal[EMA_INTERNAL_AUX_P50] + ); + assert!( + (internal[EMA_INTERNAL_AUX_STD] - 0.1).abs() < tol, + "aux_conf_std_ema: 10 identical steps should land at 0.1, got {}", + internal[EMA_INTERNAL_AUX_STD] + ); + assert!( + (internal[EMA_INTERNAL_AUX_DIR_ACC] - 1.0).abs() < tol, + "aux_dir_acc_ema: 10 identical steps should land at 1.0, got {}", + internal[EMA_INTERNAL_AUX_DIR_ACC] + ); + + // Per-step EMAs fired all 10 times. + let hold_pct_count = obs_count[3]; + let aux_p50_count = obs_count[5]; + assert_eq!(hold_pct_count, 10, "HOLD_PCT obs_count must be 10"); + assert_eq!(aux_p50_count, 10, "AUX_P50 obs_count must be 10"); + + // Per-trade EMAs never fired (is_close=0). + let wr_count = obs_count[1]; + assert_eq!(wr_count, 0, "WR obs_count must remain 0"); + } + + /// Test 5 — Wiener-α floor 0.4 contract. + /// + /// The launcher exports `WIENER_ALPHA_FLOOR = 0.4` per + /// `pearl_wiener_alpha_floor_for_nonstationary`. Lock the value + /// against silent drift. + #[test] + fn wiener_alpha_floor_locked_at_0_4() { + assert!( + (WIENER_ALPHA_FLOOR - 0.4).abs() < 1e-9, + "Wiener-α floor must be 0.4 per pearl_wiener_alpha_floor_for_nonstationary", + ); + } +} diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index cd4cf15f3..20d9f8f26 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -10907,3 +10907,207 @@ caller. AUX_GATE_TEMP / HOLD_COST_SCALE` from the EMAs. - Phase 1.4: training-loop wire-up calls all 3 producers in sequence on the same stream as the aux-head forward. + +## 2026-05-09 — SP20 Phase 1.2: sp20_emas_compute kernel (additive) + +### Component + +- New CUDA kernel `crates/ml/src/cuda_pipeline/sp20_emas_compute_kernel.cu` +- New Rust launcher `crates/ml/src/cuda_pipeline/sp20_emas_compute.rs` + (re-exported from `cuda_pipeline::sp20_emas_compute`) +- New GPU oracle test `crates/ml/tests/sp20_emas_compute_test.rs` +- Cubin manifest entry in `crates/ml/build.rs` + +### Purpose + +Component 5 / Kernel 1 of the SP20 fused-producer chain — the central +state-tracker that updates 8 Wiener-α EMAs per training step. Phase +1.4 will land the production launch site atomically with Kernel 2 +(`sp20_controllers_compute`, Phase 1.3) per +`feedback_no_partial_refactor`. + +### EMAs maintained + +Per `pearl_wiener_alpha_floor_for_nonstationary` (α floor 0.4): + +| EMA | Storage | Gate | Observation source | +|------------------|--------------------|---------------------------|---------------------------------------------| +| `ALPHA_EMA` | ISV[511] | `is_close == 1` | `alpha = R_event - hold_baseline` | +| `WR_EMA` | ISV[512] | `is_close == 1` | `is_win` ∈ {0.0, 1.0} | +| `TRADE_DURATION` | internal[0] | `is_close == 1` | `trade_duration` (bars) | +| `HOLD_PCT_EMA` | ISV[515] | per-step (always) | `(action_is_hold == 1) ? 1 : 0` | +| `HOLD_REWARD_EMA`| ISV[516] | `action_is_hold == 1` | `r_per_bar = -aux_conf * cost_scale` | +| `AUX_P50` | internal[1] | per-step (always) | `aux_logits_p50` (Phase 1.1 output) | +| `AUX_STD` | internal[2] | per-step (always) | `aux_logits_std` (Phase 1.1 output) | +| `AUX_DIR_ACC` | internal[3] | per-step (always) | `aux_dir_acc` ∈ {0.0, 1.0} | + +The 4 internal scratch slots are deliberately NOT exposed as ISV +constants — Phase 1.3's `sp20_controllers_compute` consumes them +in-kernel to derive the public ISV controllers (`N_STEP`, +`TARGET_HOLD_PCT`, `AUX_GATE_TEMP`, `AUX_CONF_THRESHOLD`). Keeping +them private keeps the SP20 ISV slot count fixed at the 10 reserved +in `f5eed1fa7`. + +### Bootstrap design — counter-based, not value-based + +Per `pearl_first_observation_bootstrap`: sentinel = 0; first +observation REPLACES directly. The naive value-based sentinel-detect +pattern (`fabsf(current - 0) < EPS` ⇒ replace) used in +`sp17_a_var_ema_kernel.cu` does NOT work for some SP20 EMAs because +0.0 is a *plausible legitimate observation*: + + - `WR_EMA`: first-loss ⇒ legit observation 0.0; second-loss observation + would be incorrectly treated as "still un-bootstrapped" and replace + instead of blend, never converging away from 0. + - `HOLD_PCT_EMA`: long initial non-Hold streak legitimately keeps + the EMA at 0; a first Hold action would over-correct. + - `HOLD_REWARD_EMA`: post-bootstrap, identical r_per_bar = −0.05 + observations would all read `current ≈ −0.05`, not equal sentinel, + so this case happens to work — but the contract is asymmetric + across EMAs which is a contract bug waiting to happen. + +This kernel uses a **per-EMA i32 observation counter** (8 entries in +a separate mapped-pinned `obs_count` buffer, aligned with +`OBS_COUNT_*` constants in the launcher) that atomically transitions +0 → 1 on the first firing observation. `count == 0` ⇒ replace, +`count > 0` ⇒ Wiener-blend at α = 0.4. This pattern is safer than +value-based sentinel detect for any EMA whose value range includes 0 +as a legitimate observation; subsequent SP20 EMA producers (Kernel 2 +controller outputs that aren't EMAs themselves) won't need the same +mechanism, but any future EMA-style producer added here SHOULD reuse +the same counter buffer rather than reverting to value-based detect. + +### Algorithm + +Single-block, single-thread kernel. Per launch, for each of the 8 +EMAs: + + 1. If the gate is OFF for this EMA this step, skip. + 2. Read `prev` from the target buffer (ISV or internal). + 3. Read `count` from `obs_count[ema_idx]`. + 4. If `count == 0`, `next = observation` (Pearl-A bootstrap). + 5. Else `next = (1 - α) * prev + α * observation` with α = 0.4 + (`WIENER_ALPHA_FLOOR_F`). + 6. Write `next` back; increment `obs_count[ema_idx]`. + +Single `__threadfence_system()` at end covers all writes (single +thread, sequentially consistent within itself; the fence makes them +PCIe-visible to mapped-pinned host pages). + +### Wiring contract + +Phase 1.2 wires kernel + launcher + tests + build entry **only**. +There is NO production caller in this commit — the kernel is dead +code awaiting Phase 1.4's atomic wire-up of: + +1. Trainer-owned `MappedF32Buffer<4>` for the `internal` scratch. +2. Trainer-owned `MappedI32Buffer<8>` for the `obs_count` counters. +3. Per-step plumbing of the 9 scalar inputs (`is_close`, `is_win`, + `alpha`, `trade_duration`, `action_is_hold`, `r_per_bar`, + `aux_p50_in`, `aux_std_in`, `aux_dir_acc`) from their producers + (trade-physics for the close/win/duration; reward kernel for the + alpha; experience collector for the hold-state and r_per_bar; + `sp20_stats_compute` Phase 1.1 for the two stats; aux-head + compare for `aux_dir_acc`). +4. Stream-ordering: launch on the SAME stream as `sp20_stats_compute` + so the stream-ordering invariant guarantees `aux_p50_in` / + `aux_std_in` reads see the latest stats. +5. Fold-boundary reset: `internal` buffer + `obs_count` counters + must reset alongside ISV slots [510..520) at fold transitions. + Phase 1.4 adds the dispatch arms. + +The `cuda_pipeline::sp20_emas_compute` module exports +`launch_sp20_emas_compute` + the contract constants +(`WIENER_ALPHA_FLOOR`, `EMA_INTERNAL_*`, `OBS_COUNT_*`, +`EmaInputs`) the wire-up will consume. + +### Pearls + invariants honoured + +- `pearl_first_observation_bootstrap` — counter-based bootstrap; + count==0 ⇒ replace, count>0 ⇒ blend. +- `pearl_wiener_alpha_floor_for_nonstationary` — α = 0.4 hardcoded + per `feedback_isv_for_adaptive_bounds` (structural floor, not + adaptive bound). +- `feedback_no_atomicadd` — single-thread kernel; no concurrent + writes anywhere. +- `feedback_no_cpu_compute_strict` — every blend lives in the + kernel; the launcher only packs scalar inputs into kernel args. +- `feedback_no_htod_htoh_only_mapped_pinned` — both ISV and the + `internal` / `obs_count` buffers are mapped-pinned; kernel emits + `__threadfence_system()` for PCIe-visible coherence. +- `pearl_no_host_branches_in_captured_graph` — single-block, + single-thread kernel; safe to capture in the per-step CUDA + Graph. +- `feedback_no_partial_refactor` — kernel + launcher + tests + + build entry land in one commit; production wire-up lands + atomically in Phase 1.4 with the controller producer. +- `feedback_no_stubs` — every output written for real; no + return-zero placeholders. + +### Tests + +`crates/ml/tests/sp20_emas_compute_test.rs` — GPU oracle tests +(all `#[ignore = "requires GPU"]` except the floor lock): + + 1. `first_observation_replaces_sentinel` — fire one trade-close + with `is_win=1, alpha=0.5, trade_duration=7`; verify + `WR_EMA == 1.0`, `ALPHA_EMA == 0.5`, + `internal[TRADE_DURATION] == 7.0` (no blending). Also + verifies the per-EMA counters track only firings. + 2. `wiener_alpha_converges_to_long_run_mean` — 100 alternating + `is_win = 1/0/1/0…` events; assert `|WR_EMA - 0.5| < 0.30`. + Bound is generous to absorb the alternation noise inherent to + a binary-indicator EMA at α = 0.4. + 3. `hold_reward_ema_gated_on_hold_bars` — 5 non-Hold bars + (`action_is_hold=0, r_per_bar=−0.05`) followed by 5 Hold bars + with same `r_per_bar`; verify `HOLD_REWARD_EMA ≈ −0.05` after + all 10 (only 5 fired, first bootstrapped, next 4 were + identical-input no-ops). Counter check: HOLD_REWARD count == 5. + 4. `per_step_emas_fire_unconditionally` — 10 steps with + `is_close=0, action_is_hold=0`, `aux_p50=0.2, aux_std=0.1, + aux_dir_acc=1.0`; verify the 4 per-step EMAs reach exactly + those values (first observation bootstraps; subsequent + identical observations are no-op blends). Counter check: + per-step counters == 10, per-trade counters == 0. + 5. `wiener_alpha_floor_locked_at_0_4` — non-GPU contract lock. + +Plus 4 unit tests in the launcher module (Wiener floor lock, +internal/obs_count slot count locks, EmaInputs sentinel safety). + +### Files modified + +| File | Status | Purpose | +|------|--------|---------| +| `crates/ml/src/cuda_pipeline/sp20_emas_compute_kernel.cu` | NEW | Single-block 8-EMA producer kernel | +| `crates/ml/src/cuda_pipeline/sp20_emas_compute.rs` | NEW | Rust launcher + contract constants | +| `crates/ml/tests/sp20_emas_compute_test.rs` | NEW | 4 GPU oracle tests + 1 floor-lock test | +| `crates/ml/src/cuda_pipeline/mod.rs` | +pub mod | Module declaration | +| `crates/ml/build.rs` | +cubin entry | Compile entry for nvcc | +| `docs/dqn-wire-up-audit.md` | This entry | Audit log | + +### Verification + +``` +SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo check -p ml --features cuda +SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml \ + --test sp20_emas_compute_test --features cuda -- --ignored --nocapture +SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 cargo test -p ml \ + --features cuda --lib sp20_emas_compute +``` + +All 4 GPU oracle tests + 4 launcher unit tests + 1 floor-lock test +pass on RTX 3050 Ti (sm_86). L40S verification deferred until Phase +1.4 lands the production caller. + +### Phase 1.2 → Phase 1.4 forward references + +- Phase 1.3: `sp20_controllers_compute_kernel.cu` reads the 4 ISV + EMAs (511, 512, 515, 516) + the 4 internal scratch slots + (`internal[0..4]`) and derives the 6 controller outputs + (LOSS_CAP, HOLD_COST_SCALE, TARGET_HOLD_PCT, N_STEP, + AUX_CONF_THRESHOLD, AUX_GATE_TEMP). +- Phase 1.4: trainer plumbs the per-step inputs into the + `EmaInputs` struct, allocates the `internal` / + `obs_count` buffers (with fold-reset registry entries), and + launches all 3 SP20 kernels in sequence on the same stream as + the aux-head forward (Kernel 3 → Kernel 1 → Kernel 2).