feat(sp20): Phase 1.2 sp20_emas_compute kernel
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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)
|
||||
|
||||
@@ -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
|
||||
|
||||
315
crates/ml/src/cuda_pipeline/sp20_emas_compute.rs
Normal file
315
crates/ml/src/cuda_pipeline/sp20_emas_compute.rs
Normal file
@@ -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);
|
||||
}
|
||||
}
|
||||
222
crates/ml/src/cuda_pipeline/sp20_emas_compute_kernel.cu
Normal file
222
crates/ml/src/cuda_pipeline/sp20_emas_compute_kernel.cu
Normal file
@@ -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 <cuda_runtime.h>
|
||||
|
||||
/* 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();
|
||||
}
|
||||
347
crates/ml/tests/sp20_emas_compute_test.rs
Normal file
347
crates/ml/tests/sp20_emas_compute_test.rs
Normal file
@@ -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<CudaStream> {
|
||||
let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?");
|
||||
ctx.default_stream()
|
||||
}
|
||||
|
||||
fn load_sp20_emas_compute(stream: &Arc<CudaStream>) -> 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<f32>, Vec<f32>, Vec<i32>) {
|
||||
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<EmaInputs> = (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<EmaInputs> = 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<EmaInputs> = (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",
|
||||
);
|
||||
}
|
||||
}
|
||||
@@ -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).
|
||||
|
||||
Reference in New Issue
Block a user