From 1790a31b66d31e03567016e4edd4ae5e51dee3ce Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sun, 10 May 2026 20:13:13 +0200 Subject: [PATCH] =?UTF-8?q?feat(sp21):=20T3.1+T3.2=20ISV=20defrost=20?= =?UTF-8?q?=E2=80=94=20wr=5Fema=20+=20hold=5Fpct=5Fema=20(atomic)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit SP21 Tier-3 foundation: defrost two production-frozen ISV slots that pinned at 0 across every observed training epoch (d7bj7, xmd6b). Both bugs in the same aggregator kernel, same structural shape: per-step binary majority-vote indicators feeding fractional EMAs. T3.1 — WR_EMA defrost (sp20_aggregate_inputs_kernel.cu): - was: is_win_out = (2 * wins_count >= closed_count) ? 1 : 0 binary "majority won this step" indicator - now: win_fraction_out = (float)wins_count / (float)closed_count fractional in [0, 1] - With actual val WR≈0.46, the binary signal was 0 most of the time, pinning WR_EMA at 0 across 30+ epochs in xmd6b. EMA now converges to population win rate. T3.2 — HOLD_PCT_EMA defrost (same kernel, same pattern): - was: action_is_hold_out = (hold_count * 2 > n_envs) ? 1 : 0 strict-majority indicator - now: hold_fraction_out = (float)hold_count / (float)n_envs - Cascade: HOLD_COST_SCALE controller compared hold_pct_ema=0 to tgt±0.05, always saw < lower, ramped × 0.95 → clamped at floor 0.01 (the hold_cost_scale=0.0100 observation in d7bj7 logs). T3.5 expected to cascade-fix in next training run. - HOLD_REWARD_EMA gate preserves strict-majority semantic via hold_fraction > 0.5f test in the consumer kernel. Atomic across struct fields, kernel logic, Rust mirror, byte serialization, doc tables, and 4 test files (per feedback_no_partial_refactor): - sp20_aggregate_inputs_kernel.cu (struct + 2 computations) - sp20_emas_compute_kernel.cu (struct + reader + gate) - sp20_aggregate_inputs.rs (doc table) - sp20_emas_compute.rs (Rust struct + serialize + 3 tests) - sp20_aggregate_inputs_test.rs (reader sig + 4 test assertions) - sp20_emas_compute_test.rs (5 struct literal updates) - sp20_phase1_4_wireup_test.rs (HOLD_PCT_EMA expected 0.625) Verification: - cargo check -p ml --tests: passes (warnings only) - cargo test -p ml --lib sp20_emas: 6/6 unit tests pass Pearl candidate: binary-majority aggregator over a fractional underlying signal cannot serve as input to a fractional-target EMA. The previous fix (commit 64bbbe418) addressed the per-bar vs segment predicate at the producer site, but didn't notice the aggregator's binarization step still collapsed the fraction to {0, 1}. Two bugs in series, both now resolved. Plan reference: docs/plans/2026-05-10-sp21-train-eval-coherence-isv-defrost.md Tiers 3.3-3.6 remaining; T3.5 expected to cascade-fix. Co-Authored-By: Claude Opus 4.7 (1M context) --- .../cuda_pipeline/sp20_aggregate_inputs.rs | 4 +- .../sp20_aggregate_inputs_kernel.cu | 56 ++++++++--- .../ml/src/cuda_pipeline/sp20_emas_compute.rs | 61 +++++++----- .../cuda_pipeline/sp20_emas_compute_kernel.cu | 55 +++++++---- crates/ml/tests/sp20_aggregate_inputs_test.rs | 95 +++++++++++-------- crates/ml/tests/sp20_emas_compute_test.rs | 20 ++-- crates/ml/tests/sp20_phase1_4_wireup_test.rs | 22 +++-- docs/dqn-wire-up-audit.md | 64 +++++++++++++ 8 files changed, 266 insertions(+), 111 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs b/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs index c4f585094..70841ba0e 100644 --- a/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs +++ b/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs @@ -14,9 +14,9 @@ //! | Field | Rule | //! |------------------------|---------------------------------------------------------------------------------------| //! | `is_close` | `1` if any env's `trade_close[env] != 0`, else `0` | -//! | `is_win` | `1` if `is_close == 1` AND fraction of closed envs with `step_ret > 0` ≥ 0.5, else `0` | +//! | `win_fraction` | `wins_count / closed_count ∈ [0, 1]` when `is_close == 1`, else `0.0` (SP21 T3.1: was binary majority `is_win`) | //! | `trade_duration` | `round(mean(hold_at_exit) over closed envs)` if `is_close == 1`, else `0` | -//! | `action_is_hold` | `1` if `count(decoded_dir == HOLD) > n_envs/2`, else `0` | +//! | `hold_fraction` | `hold_count / n_envs ∈ [0, 1]` (SP21 T3.2: was binary `action_is_hold` strict-majority vote) | //! | `alpha` | `mean(alpha_per_env) over closed envs` if `is_close == 1`, else `0` — **SP20 Phase 2 Task 2.2** (R_event - hold_baseline). | //! | `per_bar_hold_reward` | `0.0` — **Phase 3.2 forward reference** (-aux_conf * cost_scale) | //! | `aux_logits_p50` | `aux_logits_p50_dev[0]` (sp20_stats output) | diff --git a/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu b/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu index a779364f3..af388a0da 100644 --- a/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu +++ b/crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu @@ -104,9 +104,22 @@ * `#[repr(C)]`) ensures both kernels see the same byte layout. */ struct SP20EmaInputs { int is_close; - int is_win; + /* SP21 T3.1 (2026-05-10): fractional win rate ∈ [0, 1] = wins_count / + * closed_count when is_close == 1, else 0.0f. Replaces the prior + * binary `int is_win` (majority-vote) field which structurally pinned + * WR_EMA at 0 whenever true WR < 0.5 (the binary "majority won this + * step" indicator was 0 most of the time when actual win rate was + * < 50%). The EMA blends over this fractional value so WR_EMA + * converges to the true population win rate. */ + float win_fraction; int trade_duration; - int action_is_hold; + /* SP21 T3.2 (2026-05-10): fractional hold rate ∈ [0, 1] = hold_count / + * n_envs. Replaces the prior binary `int action_is_hold` (strict + * majority vote) which structurally pinned HOLD_PCT_EMA at 0 whenever + * fewer than half the envs were in Hold (the canonical case during + * directional trading). HOLD_REWARD_EMA's gate preserves the strict- + * majority semantic via `hold_fraction > 0.5f` test in the consumer. */ + float hold_fraction; float alpha; float per_bar_hold_reward; float aux_logits_p50; @@ -283,12 +296,22 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel( const int is_close_out = (closed_count > 0) ? 1 : 0; - /* is_win: fraction of closed envs that won >= 0.5. */ - int is_win_out = 0; + /* SP21 T3.1 (2026-05-10): fractional win rate, NOT binary majority. + * `win_fraction = wins_count / closed_count ∈ [0, 1]`. The prior + * binary majority-vote `(2 * wins_count >= closed_count)` predicate + * was structurally incapable of representing WR < 0.5 — it emitted + * 0 whenever fewer than half the closed envs won, regardless of + * the actual ratio, pinning WR_EMA at 0 in production whenever + * true WR < 0.5 (the canonical case). The downstream EMA blends + * over this float directly so WR_EMA converges to the population + * win rate. + * + * Emit 0.0f when `is_close_out == 0` for cleanliness — the consumer + * gates the blend on `is_close == 1` so the value is moot in that + * branch. */ + float win_fraction_out = 0.0f; if (is_close_out == 1) { - /* (wins_count / closed_count) >= 0.5 - * ⇔ 2 * wins_count >= closed_count (avoids fp). */ - is_win_out = (2 * wins_count >= closed_count) ? 1 : 0; + win_fraction_out = (float)wins_count / (float)closed_count; } /* trade_duration: round(mean of hold_at_exit over closed envs). */ @@ -311,9 +334,18 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel( alpha_out = alpha_sum_f / (float)closed_count; } - /* action_is_hold: majority-vote over envs. > n/2 threshold - * matches the spec (strict majority, ties ⇒ 0). */ - const int action_is_hold_out = (hold_count * 2 > n_envs) ? 1 : 0; + /* SP21 T3.2 (2026-05-10): fractional hold rate, NOT binary majority. + * `hold_fraction = hold_count / n_envs ∈ [0, 1]`. The prior strict- + * majority `(hold_count * 2 > n_envs)` predicate was structurally + * incapable of representing hold rates < 0.5 — it emitted 0 whenever + * fewer than half the envs were in Hold, pinning HOLD_PCT_EMA at 0 + * during directional trading and cascade-pinning HOLD_COST_SCALE at + * its floor (the controller saw `hold_pct_ema < lower` every step + * and ramped × 0.95 to clamp). Downstream HOLD_REWARD_EMA gate + * preserves the strict-majority semantic via `hold_fraction > 0.5f`. */ + const float hold_fraction_out = (n_envs > 0) + ? (float)hold_count / (float)n_envs + : 0.0f; /* Read the upstream stats / dir-acc outputs. Stream-ordering * invariant guarantees these are visible to this kernel's @@ -330,9 +362,9 @@ extern "C" __global__ void sp20_aggregate_inputs_kernel( * reference until Component 2 lands the per-bar Hold opp-cost * dual emission. */ out_inputs->is_close = is_close_out; - out_inputs->is_win = is_win_out; + out_inputs->win_fraction = win_fraction_out; out_inputs->trade_duration = trade_duration_out; - out_inputs->action_is_hold = action_is_hold_out; + out_inputs->hold_fraction = hold_fraction_out; /* SP21 T3.2 */ out_inputs->alpha = alpha_out; /* Task 2.2 — was 0.0f */ out_inputs->per_bar_hold_reward = 0.0f; /* Phase 3.2 fwd ref */ out_inputs->aux_logits_p50 = p50_in; diff --git a/crates/ml/src/cuda_pipeline/sp20_emas_compute.rs b/crates/ml/src/cuda_pipeline/sp20_emas_compute.rs index 719d00aa9..269c49cc8 100644 --- a/crates/ml/src/cuda_pipeline/sp20_emas_compute.rs +++ b/crates/ml/src/cuda_pipeline/sp20_emas_compute.rs @@ -19,10 +19,10 @@ //! | 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, 1} | +//! | `WR_EMA` | ISV[512] | `is_close == 1` | `win_fraction` ∈ [0, 1] (SP21 T3.1) | //! | `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` | +//! | `HOLD_PCT_EMA` | ISV[515] | per-step | `hold_fraction ∈ [0, 1]` (SP21 T3.2) | +//! | `HOLD_REWARD_EMA`| ISV[516] | `hold_fraction > 0.5` | `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} | @@ -147,21 +147,30 @@ pub const OBS_COUNT_AUX_DIR_ACC: usize = 7; /// 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 if the closed trade was a win, 0 otherwise. -/// Only meaningful when `is_close == 1`. Stored as i32 to match -/// the device-struct field width. +/// - `win_fraction`: fractional win rate ∈ [0, 1] = wins_count / +/// closed_count, only meaningful when `is_close == 1`. SP21 T3.1 +/// (2026-05-10): replaced the prior binary `is_win` (i32 majority +/// vote) — the aggregator now emits the actual population win rate +/// as a float so WR_EMA converges to the true value. The prior +/// binary indicator structurally pinned WR_EMA at 0 whenever true +/// WR < 0.5. /// - `trade_duration`: bars in the closed trade. Only meaningful /// when `is_close == 1`. Stored as i32 (rounded mean across envs /// in the production aggregation path). -/// - `action_is_hold`: 1 if the (majority) action this step was -/// Hold, 0 otherwise. Gates the HOLD_REWARD_EMA; also drives the -/// HOLD_PCT_EMA observation. +/// - `hold_fraction`: fractional hold rate ∈ [0, 1] = hold_count / +/// n_envs. SP21 T3.2 (2026-05-10): replaced the prior binary +/// `action_is_hold` (i32 strict-majority vote). HOLD_PCT_EMA +/// observes this directly so the EMA converges to the actual hold +/// rate, not the long-run probability of strict-majority hold. +/// HOLD_REWARD_EMA gate preserves the strict-majority semantic via +/// `hold_fraction > 0.5f` test in the consumer kernel. /// - `alpha`: `R_event - hold_baseline` from the SP20 spec /// Component 1 reward kernel. Phase 2 wires the real producer; /// Phase 1.4 wire-up populates this with sentinel 0.0 (forward /// reference, see audit doc + `sp20_aggregate_inputs_kernel.cu`). /// - `per_bar_hold_reward`: per-bar Hold reward = `-aux_conf * -/// cost_scale`. Only meaningful when `action_is_hold == 1`. +/// cost_scale`. Only meaningful when `hold_fraction > 0.5` (the +/// strict-majority gate preserved from pre-SP21-T3.2 semantics). /// Phase 3.2 wires the real producer; Phase 1.4 populates this /// with sentinel 0.0 (forward reference). /// - `aux_logits_p50`: per-step `aux_logits_p50` from Phase 1.1's @@ -176,9 +185,9 @@ pub const OBS_COUNT_AUX_DIR_ACC: usize = 7; #[derive(Copy, Clone, Debug, Default)] pub struct EmaInputs { pub is_close: i32, - pub is_win: i32, + pub win_fraction: f32, pub trade_duration: i32, - pub action_is_hold: i32, + pub hold_fraction: f32, pub alpha: f32, pub per_bar_hold_reward: f32, pub aux_logits_p50: f32, @@ -215,15 +224,14 @@ pub fn pack_inputs_into_f32_view(out: &mut [f32], inputs: &EmaInputs) { "pack_inputs_into_f32_view: out len ({}) < SP20_EMA_INPUTS_F32_LEN ({})", out.len(), SP20_EMA_INPUTS_F32_LEN, ); - // Match the device-side `SP20EmaInputs` layout: 4 i32 + 5 f32. - // Bit-pattern transmute: the device kernel reads the f32-aliased - // buffer as i32 fields by struct offset (the .cu file's - // `int is_close` is a 4-byte read off the same offset our - // `f32::from_bits(...)` write lands at). + // Match the device-side `SP20EmaInputs` layout: 2 i32 + 7 f32 (SP21 + // T3.1 + T3.2 converted is_win → win_fraction and action_is_hold → + // hold_fraction). Bit-pattern transmute for the i32 fields; native + // f32 stores for the rest. out[0] = f32::from_bits(inputs.is_close as u32); - out[1] = f32::from_bits(inputs.is_win as u32); + out[1] = inputs.win_fraction; out[2] = f32::from_bits(inputs.trade_duration as u32); - out[3] = f32::from_bits(inputs.action_is_hold as u32); + out[3] = inputs.hold_fraction; out[4] = inputs.alpha; out[5] = inputs.per_bar_hold_reward; out[6] = inputs.aux_logits_p50; @@ -354,7 +362,7 @@ mod tests { #[test] fn ema_inputs_default_is_safe_sentinel() { // A `Default::default()` `EmaInputs` has all gates closed - // (is_close = 0, action_is_hold = 0) and all observations at + // (is_close = 0, hold_fraction = 0.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::default(); @@ -363,7 +371,7 @@ mod tests { // 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); + assert_eq!(inp.hold_fraction, 0.0); } #[test] @@ -381,11 +389,12 @@ mod tests { fn pack_inputs_round_trips() { // Sanity: pack into a 9-f32 buffer and read back the bit // patterns; verify i32 fields land in the right slots. + // SP21 T3.1: win_fraction is now a native f32 in slot [1]. let inp = EmaInputs { is_close: 1, - is_win: 1, + win_fraction: 0.46, trade_duration: 7, - action_is_hold: 0, + hold_fraction: 0.0, alpha: 0.5, per_bar_hold_reward: -0.05, aux_logits_p50: 0.2, @@ -396,9 +405,11 @@ mod tests { pack_inputs_into_f32_view(&mut buf, &inp); // Verify i32 slots come back via to_bits round trip. assert_eq!(buf[0].to_bits() as i32, 1); - assert_eq!(buf[1].to_bits() as i32, 1); + // win_fraction is a native f32 — direct comparison. + assert!((buf[1] - 0.46).abs() < 1e-6); assert_eq!(buf[2].to_bits() as i32, 7); - assert_eq!(buf[3].to_bits() as i32, 0); + // hold_fraction is a native f32 — direct comparison. + assert!(buf[3].abs() < 1e-6); // Float slots are direct. assert_eq!(buf[4], 0.5); assert_eq!(buf[5], -0.05); diff --git a/crates/ml/src/cuda_pipeline/sp20_emas_compute_kernel.cu b/crates/ml/src/cuda_pipeline/sp20_emas_compute_kernel.cu index 4ef597d7a..9a3ade0ff 100644 --- a/crates/ml/src/cuda_pipeline/sp20_emas_compute_kernel.cu +++ b/crates/ml/src/cuda_pipeline/sp20_emas_compute_kernel.cu @@ -134,9 +134,15 @@ * each 4 bytes. */ struct SP20EmaInputs { int is_close; /* 0 or 1 — gates per-trade EMAs */ - int is_win; /* 0 or 1 — only meaningful when is_close == 1 */ + /* SP21 T3.1 (2026-05-10): fractional win rate ∈ [0, 1] = wins_count / + * closed_count when is_close == 1, else 0.0f. Replaces the prior + * binary `int is_win` (majority-vote) field. */ + float win_fraction; int trade_duration; /* bars — only meaningful when is_close == 1 */ - int action_is_hold; /* 0 or 1 — gates HOLD_REWARD_EMA */ + /* SP21 T3.2 (2026-05-10): fractional hold rate ∈ [0, 1] = hold_count / + * n_envs. HOLD_PCT_EMA observes this directly. HOLD_REWARD_EMA gate + * preserves strict-majority semantic via `hold_fraction > 0.5f`. */ + float hold_fraction; float alpha; /* R_event - hold_baseline at trade close */ float per_bar_hold_reward; /* -aux_conf * cost_scale on Hold bars */ float aux_logits_p50; /* upstream sp20_stats output [0] */ @@ -207,22 +213,29 @@ extern "C" __global__ void sp20_emas_compute_kernel( * coalescing is moot, but the explicit local copies make the * dependency chain readable. */ const int is_close = ema_inputs->is_close; - const int is_win_i = ema_inputs->is_win; + /* SP21 T3.1 (2026-05-10): read fractional win rate as float directly. + * The aggregator emits `wins_count / closed_count ∈ [0, 1]` so the + * EMA converges to the population win rate. Was previously `int + * is_win` (binary majority vote) which structurally pinned WR_EMA + * at 0 whenever true WR < 0.5. */ + const float win_fraction = ema_inputs->win_fraction; const int trade_duration = ema_inputs->trade_duration; - const int action_is_hold = ema_inputs->action_is_hold; + /* SP21 T3.2 (2026-05-10): fractional hold rate. HOLD_PCT_EMA observes + * this directly; HOLD_REWARD_EMA gate uses `hold_fraction > 0.5f` to + * preserve the prior strict-majority semantic. */ + const float hold_fraction = ema_inputs->hold_fraction; const float alpha = ema_inputs->alpha; const float r_per_bar = ema_inputs->per_bar_hold_reward; const float aux_p50_in = ema_inputs->aux_logits_p50; const float aux_std_in = ema_inputs->aux_logits_std; const float aux_dir_acc_in = ema_inputs->aux_dir_acc; - const float is_win = (float)is_win_i; const float trade_duration_f = (float)trade_duration; /* ── 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, + /* WR_EMA — blends over win_fraction ∈ [0, 1] (SP21 T3.1). */ + isv[isv_idx_wr] = ema_step(isv[isv_idx_wr], win_fraction, &obs_count[OBS_IDX_WR]); /* ALPHA_EMA */ @@ -237,13 +250,16 @@ extern "C" __global__ void sp20_emas_compute_kernel( /* ── 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]); - } + /* HOLD_PCT_EMA — observation = hold_fraction ∈ [0, 1] (SP21 T3.2). + * Was previously the `(action_is_hold == 1) ? 1.0 : 0.0` binary + * indicator which structurally pinned HOLD_PCT_EMA at 0 whenever + * the policy held positions in less than half the envs. The + * downstream HOLD_COST_SCALE controller compares hold_pct_ema to + * `tgt ± 0.05` to decide ramp direction; with the binary signal it + * always saw `< lower` and clamped at floor 0.01. */ + isv[isv_idx_hold_pct] = ema_step(isv[isv_idx_hold_pct], + hold_fraction, + &obs_count[OBS_IDX_HOLD_PCT]); /* aux_conf_p50_ema — internal scratch */ internal[INT_IDX_AUX_P50] = @@ -260,8 +276,15 @@ extern "C" __global__ void sp20_emas_compute_kernel( 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) { + /* ── Hold-reward EMA: gated on hold_fraction > 0.5 (SP21 T3.2). ── + * Preserves the prior strict-majority semantic of `action_is_hold==1` + * (`hold_count*2 > n_envs` ⇔ `hold_fraction > 0.5`). The gate stays + * binary because `r_per_bar` is a per-step scalar (single value, not + * per-env); reading it only on majority-Hold steps approximates the + * "policy is in Hold this bar" semantic. T3.3 (per_bar_hold_reward + * producer is currently 0.0, Phase 3.2 forward reference) will + * unfreeze HOLD_REWARD_EMA fully when the producer wires real values. */ + if (hold_fraction > 0.5f) { isv[isv_idx_hold_reward] = ema_step(isv[isv_idx_hold_reward], r_per_bar, &obs_count[OBS_IDX_HOLD_REWARD]); diff --git a/crates/ml/tests/sp20_aggregate_inputs_test.rs b/crates/ml/tests/sp20_aggregate_inputs_test.rs index d27254573..cb50ad24c 100644 --- a/crates/ml/tests/sp20_aggregate_inputs_test.rs +++ b/crates/ml/tests/sp20_aggregate_inputs_test.rs @@ -69,23 +69,26 @@ mod gpu { } /// Read the post-launch `SP20EmaInputs` struct out of the f32-aliased - /// buffer. Returns `(is_close, is_win, trade_duration, action_is_hold, - /// alpha, per_bar_hold_reward, p50, std, dir_acc)`. - fn read_ema_inputs(buf: &MappedF32Buffer) -> (i32, i32, i32, i32, f32, f32, f32, f32, f32) { + /// buffer. Returns `(is_close, win_fraction, trade_duration, + /// hold_fraction, alpha, per_bar_hold_reward, p50, std, dir_acc)`. + /// SP21 T3.1+T3.2 (2026-05-10): slots 1 and 3 are native f32 + /// (`win_fraction`, `hold_fraction`) — were binary `is_win` / + /// `action_is_hold` previously. + fn read_ema_inputs(buf: &MappedF32Buffer) -> (i32, f32, i32, f32, f32, f32, f32, f32, f32) { let v = buf.read_all(); assert_eq!(v.len(), SP20_EMA_INPUTS_F32_LEN); - // The kernel writes int fields in the first 4 slots and float - // fields in the last 5. We cast back from f32::to_bits. + // 2 i32 fields + 7 f32 fields. i32 fields come back via to_bits; + // f32 fields are direct reads. let is_close = v[0].to_bits() as i32; - let is_win = v[1].to_bits() as i32; + let win_fraction = v[1]; let trade_duration = v[2].to_bits() as i32; - let action_is_hold = v[3].to_bits() as i32; + let hold_fraction = v[3]; let alpha = v[4]; let per_bar_hr = v[5]; let p50 = v[6]; let std = v[7]; let dir_acc = v[8]; - (is_close, is_win, trade_duration, action_is_hold, + (is_close, win_fraction, trade_duration, hold_fraction, alpha, per_bar_hr, p50, std, dir_acc) } @@ -102,7 +105,7 @@ mod gpu { aux_logits_p50: f32, aux_logits_std: f32, aux_dir_acc: f32, - ) -> (i32, i32, i32, i32, f32, f32, f32, f32, f32) { + ) -> (i32, f32, i32, f32, f32, f32, f32, f32, f32) { let stream = make_test_stream(); let kernel = load_aggregation_kernel(&stream); let n = trade_close.len(); @@ -169,7 +172,7 @@ mod gpu { aux_logits_p50: f32, aux_logits_std: f32, aux_dir_acc: f32, - ) -> (i32, i32, i32, i32, f32, f32, f32, f32, f32) { + ) -> (i32, f32, i32, f32, f32, f32, f32, f32, f32) { let stream = make_test_stream(); let kernel = load_aggregation_kernel(&stream); let n = trade_close.len(); @@ -243,7 +246,7 @@ mod gpu { aux_logits_p50: f32, aux_logits_std: f32, aux_dir_acc: f32, - ) -> (i32, i32, i32, i32, f32, f32, f32, f32, f32) { + ) -> (i32, f32, i32, f32, f32, f32, f32, f32, f32) { let stream = make_test_stream(); let kernel = load_aggregation_kernel(&stream); let n = trade_close.len(); @@ -314,9 +317,9 @@ mod gpu { run_kernel(&trade_close, &step_ret, &hold_at_exit, &actions, 0.3, 0.1, 0.55); assert_eq!(ic, 0, "is_close must be 0 (no trades closed)"); - assert_eq!(iw, 0, "is_win must be 0 when is_close=0"); + assert_eq!(iw, 0.0, "win_fraction must be 0.0 when is_close=0"); assert_eq!(td, 0, "trade_duration must be 0 when is_close=0"); - assert_eq!(ah, 0, "action_is_hold must be 0 (all envs picked Long)"); + assert_eq!(ah, 0.0, "hold_fraction must be 0.0 (all envs picked Long)"); assert_eq!(alpha, 0.0, "alpha is Phase 2 forward ref placeholder = 0.0"); assert_eq!(hr, 0.0, "per_bar_hold_reward is Phase 3.2 forward ref = 0.0"); assert!((p50 - 0.3).abs() < 1e-6, "aux_logits_p50 forwarded"); @@ -325,9 +328,11 @@ mod gpu { } /// Half envs win, half lose; all close. Expected: - /// is_close = 1, is_win = 1 (4 wins / 8 closed = 0.5, threshold). + /// is_close = 1, win_fraction = 0.5 (4 wins / 8 closed). /// trade_duration = round(mean(hold_at_exit)) = round(7.5) = 8 /// (round-to-nearest-even at 7.5 ⇒ 8). + /// SP21 T3.1: was `is_win = 1` (binary majority threshold); now + /// `win_fraction = 0.5` exact. #[test] #[ignore = "requires GPU"] fn half_wins_meets_threshold_and_duration_rounds() { @@ -340,7 +345,8 @@ mod gpu { run_kernel(&trade_close, &step_ret, &hold_at_exit, &actions, 0.0, 0.0, 0.0); assert_eq!(ic, 1); - assert_eq!(iw, 1, "fraction wins (4/8 = 0.5) meets >= 0.5 threshold"); + assert!((iw - 0.5).abs() < 1e-6, + "win_fraction = 4/8 = 0.5 exactly; got {iw}"); // round-to-nearest-even at 7.5 = 8 (banker's rounding); CUDA // __float2int_rn uses round-to-nearest-even. Either 7 or 8 are // structurally acceptable; the kernel uses CUDA's RN so we @@ -349,14 +355,15 @@ mod gpu { td == 7 || td == 8, "trade_duration must round 7.5 to 7 (down) or 8 (RN-even); got {td}", ); - assert_eq!(ah, 0); + assert_eq!(ah, 0.0, + "hold_fraction must be 0.0 (all envs picked Long, no Hold)"); } - /// Strict majority Hold over n_envs. With 5 of 8 picking Hold, - /// 5 > 8/2 = 4 ⇒ action_is_hold = 1. + /// 5 of 8 envs picking Hold → hold_fraction = 5/8 = 0.625 (SP21 T3.2: + /// was binary `action_is_hold = 1` under strict majority). #[test] #[ignore = "requires GPU"] - fn majority_hold_sets_action_is_hold() { + fn five_of_eight_hold_yields_fraction_0_625() { let trade_close = vec![0; 8]; let step_ret = vec![0.0; 8]; let hold_at_exit = vec![0.0; 8]; @@ -367,14 +374,19 @@ mod gpu { let (_, _, _, ah, _, _, _, _, _) = run_kernel(&trade_close, &step_ret, &hold_at_exit, &actions, 0.0, 0.0, 0.0); - assert_eq!(ah, 1, "5 of 8 = strict majority Hold => action_is_hold=1"); + assert!( + (ah - 0.625).abs() < 1e-6, + "5 of 8 Hold => hold_fraction = 5/8 = 0.625 exact; got {ah}", + ); } - /// Tie at exactly half Hold. 4 of 8 = 4 = n_envs/2; the kernel uses - /// `count * 2 > n_envs` (strict majority), so a tie ⇒ 0. + /// Half of envs picking Hold → hold_fraction = 4/8 = 0.5 exact (SP21 + /// T3.2: was `action_is_hold = 0` under strict majority — ties did + /// not trigger). The HOLD_REWARD_EMA gate `hold_fraction > 0.5f` + /// preserves the strict-majority semantic at the consumer site. #[test] #[ignore = "requires GPU"] - fn tie_hold_count_does_not_set_action_is_hold() { + fn half_hold_yields_fraction_0_5_exact() { let trade_close = vec![0; 8]; let step_ret = vec![0.0; 8]; let hold_at_exit = vec![0.0; 8]; @@ -385,11 +397,14 @@ mod gpu { let (_, _, _, ah, _, _, _, _, _) = run_kernel(&trade_close, &step_ret, &hold_at_exit, &actions, 0.0, 0.0, 0.0); - assert_eq!(ah, 0, "tie (4 of 8 Hold) => action_is_hold=0 (strict majority)"); + assert!( + (ah - 0.5).abs() < 1e-6, + "4 of 8 Hold => hold_fraction = 4/8 = 0.5 exact; got {ah}", + ); } - /// Single env closed with a loss: is_close=1, is_win=0 - /// (0/1 < 0.5), trade_duration = round(3.0) = 3. + /// Single env closed with a loss: is_close=1, win_fraction=0.0 + /// (0/1), trade_duration = round(3.0) = 3. #[test] #[ignore = "requires GPU"] fn single_env_loss_close() { @@ -401,7 +416,7 @@ mod gpu { run_kernel(&trade_close, &step_ret, &hold_at_exit, &actions, 0.0, 0.0, 0.0); assert_eq!(ic, 1); - assert_eq!(iw, 0, "loss (0 wins out of 1 closed) => is_win=0"); + assert_eq!(iw, 0.0, "loss (0 wins out of 1 closed) => win_fraction=0.0"); assert_eq!(td, 3, "trade_duration = round(3.0) = 3"); } @@ -557,14 +572,15 @@ mod gpu { ); assert_eq!(ic, 1, "is_close = 1 (4 envs closed)"); - assert_eq!( - iw, 1, - "is_win MUST be 1 — segment-level wins (is_win_per_env = \ - [1,1,1,0]) put 3 of 4 closed envs in the win bucket; \ - 2*3 >= 4 ⇒ majority threshold met. The legacy `step_ret \ - > 0` predicate would have given is_win = 0 (all step_rets \ - negative due to tx-cost domination at close bars), pinning \ - WR_EMA at 0 in production.", + assert!( + (iw - 0.75).abs() < 1e-6, + "win_fraction MUST be 3/4 = 0.75 — segment-level wins \ + (is_win_per_env = [1,1,1,0]) put 3 of 4 closed envs in the \ + win bucket. SP21 T3.1: aggregator now emits the actual \ + fraction, not a binary majority indicator. The legacy \ + `step_ret > 0` predicate would have given win_fraction = 0 \ + (all step_rets negative due to tx-cost domination at close \ + bars), pinning WR_EMA at 0 in production. got iw={iw}", ); } @@ -590,12 +606,13 @@ mod gpu { 0.0, 0.0, 0.0); assert_eq!(ic, 1); - assert_eq!( - iw, 1, + assert!( + (iw - 1.0).abs() < 1e-6, "NULL is_win_per_env falls back to `sr > 0` legacy \ predicate — all 4 step_rets positive ⇒ wins_count = 4 ⇒ \ - is_win = 1 (the pre-fix oracle-test contract preserved \ - via NULL-tolerance).", + win_fraction = 4/4 = 1.0 (the pre-fix oracle-test contract \ + preserved via NULL-tolerance). SP21 T3.1: aggregator emits \ + fraction not binary majority. got iw={iw}", ); } } diff --git a/crates/ml/tests/sp20_emas_compute_test.rs b/crates/ml/tests/sp20_emas_compute_test.rs index 79b5cd180..4fdeb94da 100644 --- a/crates/ml/tests/sp20_emas_compute_test.rs +++ b/crates/ml/tests/sp20_emas_compute_test.rs @@ -146,9 +146,9 @@ mod gpu { fn first_observation_replaces_sentinel() { let inp = EmaInputs { is_close: 1, - is_win: 1, + win_fraction: 1.0, trade_duration: 7, - action_is_hold: 0, + hold_fraction: 0.0, alpha: 0.5, per_bar_hold_reward: 0.0, aux_logits_p50: 0.0, @@ -216,9 +216,9 @@ mod gpu { let inputs: Vec = (0..100) .map(|i| EmaInputs { is_close: 1, - is_win: if i % 2 == 0 { 1 } else { 0 }, + win_fraction: if i % 2 == 0 { 1.0 } else { 0.0 }, trade_duration: 1, - action_is_hold: 0, + hold_fraction: 0.0, alpha: 0.0, per_bar_hold_reward: 0.0, aux_logits_p50: 0.0, @@ -253,9 +253,9 @@ mod gpu { for _ in 0..5 { inputs.push(EmaInputs { is_close: 0, - is_win: 0, + win_fraction: 0.0, trade_duration: 0, - action_is_hold: 0, // gated OFF + hold_fraction: 0.0, // gated OFF alpha: 0.0, per_bar_hold_reward: -0.05, aux_logits_p50: 0.0, @@ -266,9 +266,9 @@ mod gpu { for _ in 0..5 { inputs.push(EmaInputs { is_close: 0, - is_win: 0, + win_fraction: 0.0, trade_duration: 0, - action_is_hold: 1, // gate ON + hold_fraction: 1.0, // > 0.5 gate ON (SP21 T3.2) alpha: 0.0, per_bar_hold_reward: -0.05, aux_logits_p50: 0.0, @@ -317,9 +317,9 @@ mod gpu { let inputs: Vec = (0..10) .map(|_| EmaInputs { is_close: 0, - is_win: 0, + win_fraction: 0.0, trade_duration: 0, - action_is_hold: 0, + hold_fraction: 0.0, alpha: 0.0, per_bar_hold_reward: 0.0, aux_logits_p50: 0.2, diff --git a/crates/ml/tests/sp20_phase1_4_wireup_test.rs b/crates/ml/tests/sp20_phase1_4_wireup_test.rs index 7f381c0e2..9554b16ef 100644 --- a/crates/ml/tests/sp20_phase1_4_wireup_test.rs +++ b/crates/ml/tests/sp20_phase1_4_wireup_test.rs @@ -254,15 +254,23 @@ mod gpu { assert_eq!(isv[ALPHA_EMA_INDEX], 0.0, "ALPHA_EMA stays at 0.0 with NULL alpha_per_env producer (Task 2.2 NULL contract)"); - // WR_EMA: 1 win out of 1 close ⇒ is_win = 1, first-obs replace - // ⇒ WR_EMA = 1.0. + // WR_EMA: 1 win out of 1 close ⇒ win_fraction = 1.0, first-obs + // replace ⇒ WR_EMA = 1.0. SP21 T3.1: aggregator emits exact + // fraction (was binary majority). assert_eq!(isv[WR_EMA_INDEX], 1.0, - "WR_EMA = 1.0 after one winning close (1 win / 1 closed >= 0.5)"); + "WR_EMA = 1.0 after one winning close (1 win / 1 closed = 1.0)"); - // HOLD_PCT_EMA: 5 of 8 envs picked Hold ⇒ action_is_hold = 1 - // ⇒ HOLD_PCT_EMA = 1.0 on first observation. - assert_eq!(isv[HOLD_PCT_EMA_INDEX], 1.0, - "HOLD_PCT_EMA = 1.0 after one majority-Hold observation"); + // HOLD_PCT_EMA: 5 of 8 envs picked Hold ⇒ hold_fraction = 5/8 = + // 0.625, first-obs replace ⇒ HOLD_PCT_EMA = 0.625. SP21 T3.2: + // aggregator emits exact fraction (was binary majority indicator + // → 1.0). The fractional value lets the downstream HOLD_COST_SCALE + // controller make graded decisions instead of the cascade- + // pinning-at-floor that the binary signal produced. + let hold_pct = isv[HOLD_PCT_EMA_INDEX]; + assert!( + (hold_pct - 0.625).abs() < 1e-6, + "HOLD_PCT_EMA = 0.625 (5/8 of envs Hold); got {hold_pct}", + ); // HOLD_REWARD_EMA: per_bar_hold_reward = 0.0 (Phase 3.2 fwd // ref); fired on Hold ⇒ EMA = 0.0. diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index 208a1044d..6503cd69f 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -2,6 +2,70 @@ **Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7. +## 2026-05-10 — SP21 Tier-3 foundation: T3.1+T3.2 ISV defrost (wr_ema, hold_pct_ema) + +**Branch:** `sp20-aux-h-fixed` (off `6df44e4c6`). +**Commits:** 1 atomic (this commit). + +Defrosts two production-frozen ISV slots that pinned-at-zero across all +observed training epochs in d7bj7/xmd6b logs. Both bugs were in the same +aggregator kernel (`sp20_aggregate_inputs_kernel.cu`), same structural +shape: per-step binary majority-vote indicators feeding fractional EMAs. + +**T3.1: WR_EMA — `is_win` → `win_fraction` (struct field rename + type change).** +The aggregator computed `is_win_out = (2 * wins_count >= closed_count) ? 1 : 0` +— a binary "majority of closed envs won" indicator. With actual val WR +≈ 0.46 (xmd6b 30 epochs), the binary signal was 0 most of the time, and +WR_EMA → 0 by Wiener-α floor blend. The aggregator now emits +`win_fraction = wins_count / closed_count ∈ [0, 1]` and the EMA blends +over the float directly. WR_EMA will converge to the population win rate. + +**T3.2: HOLD_PCT_EMA — `action_is_hold` → `hold_fraction` (same fix shape).** +The aggregator computed `action_is_hold_out = (hold_count * 2 > n_envs) ? 1 : 0` +— strict-majority indicator. With actual hold pct in [0.10, 0.56], +majority was rarely met, HOLD_PCT_EMA → 0. **Cascade**: the +`HOLD_COST_SCALE` controller compares `hold_pct_ema` to `tgt ± 0.05` and +ramps; with hold_pct_ema=0 it always saw `< lower` and ramped × 0.95 +every step → clamped at floor 0.01 (the `hold_cost_scale = 0.0100` +observation in d7bj7 sp20_isv logs). Fixing T3.2 should cascade-fix +T3.5 — verify in next training run. HOLD_REWARD_EMA gate preserves +strict-majority semantic via `hold_fraction > 0.5f` test in the +consumer kernel. + +**Affected files (atomic per `feedback_no_partial_refactor`):** + +- `crates/ml/src/cuda_pipeline/sp20_aggregate_inputs_kernel.cu` — + struct fields, output computations. +- `crates/ml/src/cuda_pipeline/sp20_emas_compute_kernel.cu` — + struct fields, EMA reader, HOLD_REWARD_EMA gate. +- `crates/ml/src/cuda_pipeline/sp20_aggregate_inputs.rs` — doc table. +- `crates/ml/src/cuda_pipeline/sp20_emas_compute.rs` — Rust struct mirror, + byte serialization, doc table, default-sentinel test, round-trip test. +- `crates/ml/tests/sp20_aggregate_inputs_test.rs` — reader signature, + 4 fractional-comparison test assertions, 2 renamed test fns. +- `crates/ml/tests/sp20_emas_compute_test.rs` — 5 struct-literal field + renames (is_win → win_fraction, action_is_hold → hold_fraction). +- `crates/ml/tests/sp20_phase1_4_wireup_test.rs` — HOLD_PCT_EMA expected + value updated 1.0 → 0.625 (5/8 envs Hold). + +**Verification:** + +- `cargo check -p ml --tests` — passes (warnings only). +- `cargo test -p ml --lib sp20_emas` — 6/6 unit tests pass, including + `pack_inputs_round_trips` with `win_fraction: 0.46` + `hold_fraction: 0.0`. +- GPU oracle tests gated `#[ignore = "requires GPU"]` — compile cleanly. + +**Pearls applied / candidates:** + +- `pearl_first_observation_bootstrap` — unchanged, both EMAs still bootstrap + from sentinel. +- `pearl_wiener_alpha_floor_for_nonstationary` — unchanged, α=0.4 floor. +- New pearl candidate: **binary-majority aggregator over a fractional + underlying signal cannot serve as input to a fractional-target EMA**. + When the EMA is supposed to converge to a population fraction, the + observation must itself be the fraction or a 0/1 sample, not a + thresholded reduction. + ## 2026-05-10 — REVERT: aux_horizon_update_kernel sp20-aux-h-fixed experiment closed Reverts the forced-H=30 diagnostic that tested the bootstrap-failure hypothesis. **Hypothesis verdict: dead.** Across two epochs of `train-d7bj7`: