From 0384f76743f0fd8a56e3428975112dae6efbd7d9 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 22 May 2026 00:25:45 +0200 Subject: [PATCH] fix(decision-policy): use signed-conviction EMA, not magnitude-only EMA MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Anti-calibration evidence (smoke ckpts across architecture variants showed higher reported conviction → MORE negative PnL, -15.6 → -145.2) traced to a magnitude/direction mismatch in the sizing formula: conv_ema = EMA(|conviction_signed|) # magnitude smoothing target_lots = sign(conviction_signed) # INSTANTANEOUS sign × conv_ema × max_lots # × SMOOTHED magnitude When per-horizon directions disagree event-to-event, the magnitude EMA stays high (|x| EMA can't cancel) but the sign flips on noise. Result: big trades in random direction whenever the 5 horizons disagree. Replace with a single signed-conviction EMA: conv_signed_ema = EMA(conviction_signed) target_lots = sign(conv_signed_ema) × |conv_signed_ema| × max_lots Now BOTH sign AND magnitude come from the same smoothed signal. When directions disagree, signed EMA → 0 collapses size to zero naturally. Atomic migration across decision_policy.cu (2 kernels), pnl_track.cu (open-branch reads the SIGNED buffer and takes fabsf for the magnitude- semantics open_trade_state offsets that composite_exit_check forms ratios from), and src/sim/mod.rs (renamed device buffers + accessor + 4 launch sites). No legacy aliases per feedback_no_legacy_aliases; no parallel buffers per feedback_single_source_of_truth_no_duplicates; α floor 0.4 preserved per pearl_wiener_alpha_floor_for_nonstationary; first- observation bootstrap preserved per pearl_first_observation_bootstrap. ml-backtesting lib: 33 passed. GPU oracle: signed_conviction_ema_collapses_on_sign_flips passes — observed steady-state shows |signed_ema| ≈ 0.205 (post-fix) producing |target_lots| = 2 every tail event, vs the pre-fix bug's predicted |target_lots| = 7 every tail event. Existing multi_horizon_conviction_cancels_on_disagreement still passes (zero- cancellation regime is even cleaner under signed EMA). ml-alpha lib: 33 passed (no regression). Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml-backtesting/cuda/decision_policy.cu | 275 +++++++++++------- crates/ml-backtesting/cuda/pnl_track.cu | 30 +- crates/ml-backtesting/src/sim/mod.rs | 93 +++--- .../ml-backtesting/tests/stop_controller.rs | 162 ++++++++++- 4 files changed, 407 insertions(+), 153 deletions(-) diff --git a/crates/ml-backtesting/cuda/decision_policy.cu b/crates/ml-backtesting/cuda/decision_policy.cu index 03e9e2a55..1e3594843 100644 --- a/crates/ml-backtesting/cuda/decision_policy.cu +++ b/crates/ml-backtesting/cuda/decision_policy.cu @@ -158,21 +158,36 @@ __device__ static int stop_check_isv( return 0; } -// CRT.1 C1.2: Wiener-α adaptive EMA on the multi-horizon ISV-weighted -// conviction scalar (spec §4.4). Caller computes |conviction_signed| from -// the per-horizon weighted-signed sum and passes it as `conviction_abs`. +// CRT.1 C1.2 (anti-calibration fix 2026-05-22): Wiener-α adaptive EMA on +// the SIGNED multi-horizon ISV-weighted conviction scalar (spec §4.4). The +// prior magnitude-only EMA had a structural anti-calibration bug: +// +// conv_ema_mag = EMA(|conviction_signed|) // magnitude smoothing +// target_lots = sign(conviction_signed) // INSTANTANEOUS sign +// × conv_ema_mag × max_lots // × SMOOTHED magnitude +// +// When per-horizon directions disagreed event-to-event, the magnitude EMA +// stayed high (|x| EMA cannot cancel) but the sign flipped on noise → +// big trades in random direction. Empirically: higher reported conviction +// correlated with MORE negative PnL (-15.6 → -145.2 in smoke ckpts). +// +// The signed-EMA fix collapses both sign and magnitude from the same +// smoothed signal: directions that disagree event-to-event cancel naturally +// in EMA(signed_conviction) → 0, which collapses |signed_ema| → 0 → size 0. // // Updates per-backtest second-order EMAs of sample variance + difference -// variance, derives adaptive α with floor 0.4 per -// pearl_wiener_alpha_floor_for_nonstationary, and blends. First-observation -// bootstrap (prev_ema == 0): replace directly per -// pearl_first_observation_bootstrap. +// variance (now both operating on the signed conviction), derives adaptive +// α with floor 0.4 per pearl_wiener_alpha_floor_for_nonstationary, and +// blends. First-observation bootstrap (prev_ema == 0): replace directly +// per pearl_first_observation_bootstrap. // -// Writes new EMA to conviction_ema[b], updates the two second-order -// variance EMAs, and returns the smoothed conviction magnitude. Direction -// is recovered by the caller from `conviction_signed` so that smoothing -// is applied only to magnitude — sign flips respond at event rate while -// magnitude jitter is suppressed. +// Writes new SIGNED EMA to conv_signed_ema[b]. Caller derives both +// direction and magnitude from `fabsf(returned_signed_ema)` and the sign +// of the returned signed EMA — guaranteeing they come from the SAME +// smoothed signal. The sample-variance EMA tracks E[signed_conviction²] +// = Var[signed_conviction] (since the signal is bounded in [-1, +1] and +// the variance-centering Welford pass would be a no-op against zero in +// the steady disagreement-cancellation case the EMA targets). // // Single-thread-per-backtest invariant: caller must already have // gated on threadIdx.x == 0. @@ -235,7 +250,9 @@ __device__ static int composite_exit_check( // Called by the decision kernel each event AFTER the multi-horizon // conviction is computed and EMA-smoothed. Writes the four offsets the // composite_exit_check reads on the next stop_check pass: -// offset 24: conviction_ema (mirror of conviction_ema_per_b[b]) +// offset 24: conviction_ema_mag (magnitude of conv_signed_ema_per_b[b]; +// composite_exit_check uses the ratio conv_ema_now / +// conviction_at_entry which only makes sense as magnitudes) // offset 44: pnl_adjusted_conviction_ema (signed by sign of unrealized) // offset 48: peak_unrealized_pnl (max tracker) // offset 56: disagreement_consecutive_events (short vs long horizon) @@ -249,7 +266,7 @@ __device__ static void update_open_trade_trajectory( const Pos* __restrict__ positions, const Book* __restrict__ books, const float* __restrict__ alpha_probs, - float conv_ema_now, + float conv_ema_mag, unsigned char* __restrict__ open_trade_state ) { const Pos& pos = positions[b]; @@ -257,10 +274,12 @@ __device__ static void update_open_trade_trajectory( unsigned char* st = open_trade_state + (size_t)b * 64; - // Mirror the current smoothed conviction into open_trade_state so - // composite_exit_check reads it without needing the conviction_ema_per_b - // pointer plumbed in. - *reinterpret_cast(st + 24) = conv_ema_now; + // Mirror the current smoothed conviction MAGNITUDE into open_trade_state. + // composite_exit_check forms `1 − conv_ema_now / conviction_at_entry` + // which is a degradation ratio — only magnitude (not signed) makes + // sense there, because the signed EMA can flip sign on a deep-loss + // reversal which would invert the ratio meaning. + *reinterpret_cast(st + 24) = conv_ema_mag; // Unrealized P&L = (close_px − vwap_entry) × dir × |lots|. // Long: close at best bid; short: close at best ask. Matches the @@ -273,13 +292,13 @@ __device__ static void update_open_trade_trajectory( const float dir_long = is_long ? 1.0f : -1.0f; const float unrealized = (close_px - pos.vwap_entry) * dir_long * lots_abs; - // pnl-adjusted conviction: sign of unrealized × conv_ema_now. EMA on + // pnl-adjusted conviction: sign of unrealized × conv_ema_mag. EMA on // top with the same fixed-α=0.1 used by the conviction second-order - // EMAs in update_conviction_ema. First-observation bootstrap per - // pearl_first_observation_bootstrap (prev == 0 ⇒ replace directly). + // EMAs in update_signed_conviction_ema. First-observation bootstrap + // per pearl_first_observation_bootstrap (prev == 0 ⇒ replace directly). const float pnl_sign = (unrealized >= 0.0f) ? 1.0f : -1.0f; const float prev_pnl_adj = *reinterpret_cast(st + 44); - const float new_pnl_adj_inst = pnl_sign * conv_ema_now; + const float new_pnl_adj_inst = pnl_sign * conv_ema_mag; const float new_pnl_adj = (prev_pnl_adj == 0.0f) ? new_pnl_adj_inst : (0.9f * prev_pnl_adj + 0.1f * new_pnl_adj_inst); @@ -307,48 +326,69 @@ __device__ static void update_open_trade_trajectory( } } -__device__ static float update_conviction_ema( +__device__ static float update_signed_conviction_ema( int b, - float conviction_abs, - float* __restrict__ conviction_ema, - float* __restrict__ conviction_diff_var_ema, - float* __restrict__ conviction_sample_var_ema + float conviction_signed, + float* __restrict__ conv_signed_ema, + float* __restrict__ conv_signed_diff_var_ema, + float* __restrict__ conv_signed_sample_var_ema ) { - // Defensive clamp to [0, 1] — conviction is bounded by construction - // (sum of magnitude_h × weight_h × direction_h divided by total |weight_h|, - // with magnitude_h ∈ [0, 1]), but bound it before feeding into a - // variance EMA in case of NaN propagation through near-zero weights. - if (!(conviction_abs > 0.0f)) conviction_abs = 0.0f; // also catches NaN - if (conviction_abs > 1.0f) conviction_abs = 1.0f; + // Defensive clamp to [-1, +1] — conviction is bounded by construction + // (weighted_sum / total_abs_weight with per-horizon magnitudes ∈ [0, 1] + // and per-horizon directions ∈ {−1, +1}), but bound it before feeding + // into a variance EMA in case of NaN propagation through near-zero + // total_abs_weight. NaN guard is reflexive — isnan(x) is true when + // x compares unequal with itself; the bilateral clamp below catches it. + if (!(conviction_signed == conviction_signed)) conviction_signed = 0.0f; // NaN + if (conviction_signed > 1.0f) conviction_signed = 1.0f; + if (conviction_signed < -1.0f) conviction_signed = -1.0f; // Update second-order EMAs (diff_var, sample_var) with fixed α=0.1. - // These track the variance of the signal and the variance of changes, - // and are themselves bootstrapped with the first observation. - const float prev_ema = conviction_ema[b]; - const float diff = conviction_abs - prev_ema; - const float diff_var_prev = conviction_diff_var_ema[b]; + // These track the variance of the signed conviction and the variance + // of its changes, and are themselves bootstrapped with the first + // observation. sample_var uses x² which is non-negative regardless of + // the sign of x, so the zero-sentinel bootstrap pattern still works + // (sample_var_prev == 0 ⇔ no observation seen yet, and a first signed + // observation produces a strictly positive x² that escapes the + // sentinel — assuming the multi-horizon formula returns non-zero on + // the first event, which is the normal case; the degenerate all-zero + // case stays bootstrapped which is correct). + const float prev_ema = conv_signed_ema[b]; + const float diff = conviction_signed - prev_ema; + const float diff_var_prev = conv_signed_diff_var_ema[b]; const float new_diff_var = (diff_var_prev == 0.0f) ? (diff * diff) // first-observation bootstrap : (0.9f * diff_var_prev + 0.1f * diff * diff); - conviction_diff_var_ema[b] = new_diff_var; + conv_signed_diff_var_ema[b] = new_diff_var; - const float sample_var_prev = conviction_sample_var_ema[b]; + const float sample_var_prev = conv_signed_sample_var_ema[b]; const float new_sample_var = (sample_var_prev == 0.0f) - ? (conviction_abs * conviction_abs) // first-observation bootstrap - : (0.9f * sample_var_prev + 0.1f * conviction_abs * conviction_abs); - conviction_sample_var_ema[b] = new_sample_var; + ? (conviction_signed * conviction_signed) // first-observation bootstrap + : (0.9f * sample_var_prev + 0.1f * conviction_signed * conviction_signed); + conv_signed_sample_var_ema[b] = new_sample_var; // Wiener α with floor at 0.4 (pearl_wiener_alpha_floor_for_nonstationary). + // Floor is critical here: the signed signal can drift around zero in + // the disagreement-cancellation regime where diff_var → 0 (consecutive + // events near the same near-zero value), which would push α_raw to + // ~0 and lock the EMA at its prev value. Floor 0.4 ensures the EMA + // continues to respond to genuine signal recovery. const float eps_v = 1e-6f; const float alpha_raw = new_diff_var / (new_diff_var + new_sample_var + eps_v); const float alpha_active = (alpha_raw > 0.4f) ? alpha_raw : 0.4f; - // First-observation bootstrap: prev_ema==0 → replace directly. Otherwise - // blend with Wiener-α. + // First-observation bootstrap: prev_ema==0 → replace directly per + // pearl_first_observation_bootstrap. Otherwise blend with Wiener-α. + // Note: a SIGNED 0.0 prev_ema is the sentinel (zero-init by host alloc), + // not a real observation. A real observation that happens to evaluate + // to exactly 0.0 (total_abs_weight==0 or perfect cancellation) is + // indistinguishable from the sentinel — but a perfect-cancellation + // event already produces the desired EMA=0 outcome via the bootstrap + // replace branch, so the ambiguity is harmless. const float new_ema = (prev_ema == 0.0f) - ? conviction_abs - : (alpha_active * conviction_abs + (1.0f - alpha_active) * prev_ema); - conviction_ema[b] = new_ema; + ? conviction_signed + : (alpha_active * conviction_signed + (1.0f - alpha_active) * prev_ema); + conv_signed_ema[b] = new_ema; return new_ema; } @@ -367,9 +407,11 @@ __device__ static float update_conviction_ema( // `total_abs_weight` makes the formula scale-invariant per // pearl_zscore_normalization_for_magnitude_asymmetric_signals. // -// Returns conviction_signed ∈ [-1, +1]. Caller derives: -// conviction_abs = fabsf(conviction_signed) -// direction = sign(conviction_signed) +// Returns conviction_signed ∈ [-1, +1]. The caller passes this signed +// scalar directly into update_signed_conviction_ema; sign and magnitude +// both come from the resulting smoothed signed EMA so the magnitude/sign +// mismatch documented in update_signed_conviction_ema's header cannot +// recur. __device__ static float compute_multi_horizon_conviction( const float* __restrict__ alpha_probs, const IsvKellyState* __restrict__ isv, @@ -402,8 +444,9 @@ __device__ static float compute_multi_horizon_conviction( // a signal, or the per-event output is genuinely noisy and the hypothesis // is wrong). // -// Pattern mirrors update_conviction_ema (which works on the multi-horizon -// scalar). Per pearl_wiener_optimal_adaptive_alpha + pearl_first_observation_bootstrap. +// Pattern mirrors update_signed_conviction_ema (which works on the +// multi-horizon scalar). Per pearl_wiener_optimal_adaptive_alpha + +// pearl_first_observation_bootstrap. // Floor at 0.1 — stronger smoothing than the 0.4 conviction-level floor, // testing whether INPUT-side smoothing differs from OUTPUT-side. // @@ -522,7 +565,7 @@ extern "C" __global__ void decision_policy_default( const unsigned int* __restrict__ program_lens, // [n_backtests] — non-zero ⇒ skip (program kernel handles it) int* __restrict__ market_targets, // [n_backtests * 2] — written unsigned int* __restrict__ open_horizon_masks, // [n_backtests] — written (entry attribution) - // Per-backtest envelope cap (CRT.1 C1.2: target = direction × conv_ema × max_lots). + // Per-backtest envelope cap (post-fix 2026-05-22: target = conv_signed_ema × max_lots). const int* __restrict__ max_lots_per_b, // [n_backtests] // Threshold gate on the smoothed multi-horizon conviction (spec §5.2). const float* __restrict__ threshold_per_b, // [n_backtests] @@ -536,11 +579,16 @@ extern "C" __global__ void decision_policy_default( const Book* __restrict__ books_per_b, // [n_backtests * BOOK_LEVELS] (via lob_state.cuh) // S2.1: mh_kernel_calls diagnostic counter (non-zero-position entries). unsigned int* __restrict__ mh_kernel_calls_per_b, // [n_backtests] - // Wiener-α conviction-EMA state (per-backtest). The three slots stay - // from v2 A2 — the EMA mechanism is reused on the multi-horizon scalar. - float* __restrict__ conviction_ema_per_b, // [n_backtests] — read + write - float* __restrict__ conviction_diff_var_ema_per_b, // [n_backtests] — read + write - float* __restrict__ conviction_sample_var_ema_per_b, // [n_backtests] — read + write + // Anti-calibration fix (2026-05-22): Wiener-α conviction-EMA state on + // the SIGNED multi-horizon conviction. The buffer name carries `_signed_` + // so call sites can't accidentally re-introduce the magnitude-only + // misuse that produced the higher-conviction-→-more-negative-PnL + // anti-calibration. Same three-slot layout as before (sample / diff / + // value EMAs); the kernel-internal update routine now operates on the + // signed signal end-to-end. + float* __restrict__ conv_signed_ema_per_b, // [n_backtests] — read + write + float* __restrict__ conv_signed_diff_var_ema_per_b, // [n_backtests] — read + write + float* __restrict__ conv_signed_sample_var_ema_per_b, // [n_backtests] — read + write // CRT.1 C1.4: trajectory state + composite exit_signal circuit-breaker. // pnl_track populated conviction_at_entry on the open transition; // this kernel updates the intra-trade fields each event so the composite @@ -613,18 +661,23 @@ extern "C" __global__ void decision_policy_default( const float cost = cost_per_lot_per_side_per_b[b]; const float conviction_signed = compute_multi_horizon_conviction_smoothed( alpha_smoothed, isv, cost); - const float conviction_abs = fabsf(conviction_signed); - const float conviction_dir = (conviction_signed >= 0.0f) ? 1.0f : -1.0f; - // Wiener-α adaptive EMA on the multi-horizon conviction magnitude. + // Anti-calibration fix (2026-05-22): Wiener-α adaptive EMA on the + // SIGNED conviction so disagreement-driven sign flips cancel in the + // EMA → |EMA| → 0 → no trade. Magnitude AND direction both come from + // the same smoothed signal — eliminates the previous magnitude/sign + // mismatch where size stayed large while sign flipped on noise. + // // Runs ahead of the threshold gate so the EMA tracks all events // (gating off an unsmoothed tick would let the EMA grow stale on long // gated stretches). - const float conv_ema = update_conviction_ema( - b, conviction_abs, - conviction_ema_per_b, - conviction_diff_var_ema_per_b, - conviction_sample_var_ema_per_b); + const float conv_signed_ema = update_signed_conviction_ema( + b, conviction_signed, + conv_signed_ema_per_b, + conv_signed_diff_var_ema_per_b, + conv_signed_sample_var_ema_per_b); + const float conv_ema_mag = fabsf(conv_signed_ema); + const float conviction_dir = (conv_signed_ema >= 0.0f) ? 1.0f : -1.0f; // CRT.diag Group A: per-horizon direction-flip + run-length tracking. // Observe-only. Single-writer-per-block (already gated on thread 0). @@ -664,9 +717,12 @@ extern "C" __global__ void decision_policy_default( diag_smoothed_sum_run_length, diag_smoothed_prev_dir); - // CRT.diag Group B: smoothed-conviction histogram (10 buckets over [0, 1)). + // CRT.diag Group B: smoothed-conviction MAGNITUDE histogram (10 buckets + // over [0, 1)). Magnitude (not signed) because the histogram tracks + // "strength of opinion" not directional bias — symmetric +0.7 vs −0.7 + // both belong in the same bucket. { - float c = conv_ema; + float c = conv_ema_mag; if (c < 0.0f) c = 0.0f; if (c >= 1.0f) c = 0.999999f; int bk = (int)(c * 10.0f); @@ -680,8 +736,10 @@ extern "C" __global__ void decision_policy_default( // NOTE: keeps RAW alpha_probs for the short-vs-long disagreement // counter (offset 56) — composite exit's disagreement signal is meant // to react fast (32-event window), opposite of the smoothing rationale. + // Passes the MAGNITUDE of the smoothed signed EMA — composite_exit_check + // forms degradation ratios that only make sense between two magnitudes. update_open_trade_trajectory(b, positions, books_per_b, alpha_probs, - conv_ema, open_trade_state_per_b); + conv_ema_mag, open_trade_state_per_b); // CRT.1 C1.4: composite exit_signal safety circuit-breaker (spec §4.3). // Catches the edge case where conviction stays bounded above zero but @@ -691,10 +749,12 @@ extern "C" __global__ void decision_policy_default( return; } - // Threshold gate on the smoothed multi-horizon conviction. Per spec §5.2: - // gating on the smoothed signal rather than the raw signal makes the - // pre-registered percentile thresholds robust to event-rate jitter. - if (conv_ema < threshold_per_b[b]) { + // Threshold gate on the smoothed multi-horizon conviction MAGNITUDE. + // Per spec §5.2: gating on the smoothed signal rather than the raw + // signal makes the pre-registered percentile thresholds robust to + // event-rate jitter. Magnitude (not signed) because the pre-registered + // thresholds are |conviction| percentiles — sign-agnostic by design. + if (conv_ema_mag < threshold_per_b[b]) { market_targets[b * 2 + 0] = 2; market_targets[b * 2 + 1] = 0; open_horizon_masks[b] = 0u; @@ -714,14 +774,15 @@ extern "C" __global__ void decision_policy_default( if (dir_h == conviction_dir) attribution_mask |= (1u << h); } - // Target = direction × smoothed conviction × envelope_max. No - // aggregate rescale step — the v2 `final_size *= (conv_ema/raw_max_conv)` - // bug that amplified weak signals is structurally absent because the - // smoothed conviction enters the size formula multiplicatively, never - // ratio'd against the raw scalar. Spec §4.2. + // Anti-calibration fix (2026-05-22): target = SIGNED smoothed conviction + // × envelope_max. Both sign AND magnitude come from `conv_signed_ema` + // — the SAME smoothed scalar. When per-horizon directions disagree + // event-to-event, conv_signed_ema → 0 collapses size to zero naturally, + // unlike the prior `sign(instantaneous) × EMA(|instantaneous|)` which + // multiplied a smoothed magnitude by an event-rate sign flip. const int max_lots = max_lots_per_b[b]; const float envelope_max = (float)max_lots; - const float target_lots_f = conviction_dir * conv_ema * envelope_max; + const float target_lots_f = conv_signed_ema * envelope_max; const int target_lots = (int)truncf( target_lots_f + (target_lots_f >= 0.0f ? 0.5f : -0.5f)); @@ -804,10 +865,11 @@ extern "C" __global__ void decision_policy_program( const Book* __restrict__ books_per_b, // [n_backtests * BOOK_LEVELS] // S2.1: mh_kernel_calls diagnostic counter (non-zero-position entries). unsigned int* __restrict__ mh_kernel_calls_per_b, // [n_backtests] - // CRT-A2: Wiener-α conviction-EMA state (per-backtest). - float* __restrict__ conviction_ema_per_b, // [n_backtests] — read + write - float* __restrict__ conviction_diff_var_ema_per_b, // [n_backtests] — read + write - float* __restrict__ conviction_sample_var_ema_per_b, // [n_backtests] — read + write + // Anti-calibration fix (2026-05-22): SIGNED conviction-EMA state. + // Mirrors decision_policy_default — see that kernel for the contract. + float* __restrict__ conv_signed_ema_per_b, // [n_backtests] — read + write + float* __restrict__ conv_signed_diff_var_ema_per_b, // [n_backtests] — read + write + float* __restrict__ conv_signed_sample_var_ema_per_b, // [n_backtests] — read + write // CRT.1 C1.4: trajectory state + composite exit_signal circuit-breaker. // Mirrors decision_policy_default — see that kernel for the contract. unsigned char* __restrict__ open_trade_state_per_b, // [n_backtests * 64] @@ -877,16 +939,19 @@ extern "C" __global__ void decision_policy_program( const float cost = cost_per_lot_per_side_per_b[b]; const float conviction_signed = compute_multi_horizon_conviction_smoothed( alpha_smoothed, isv, cost); - const float conviction_abs = fabsf(conviction_signed); - const float conviction_dir = (conviction_signed >= 0.0f) ? 1.0f : -1.0f; - // Wiener-α adaptive EMA on the multi-horizon conviction magnitude. - // Runs ahead of the threshold gate so the EMA tracks all events. - const float conv_ema = update_conviction_ema( - b, conviction_abs, - conviction_ema_per_b, - conviction_diff_var_ema_per_b, - conviction_sample_var_ema_per_b); + // Anti-calibration fix (2026-05-22): SIGNED conviction EMA. Sign and + // magnitude both come from the same smoothed scalar — directions that + // disagree event-to-event cancel in the EMA → size collapses to zero + // naturally. See update_signed_conviction_ema header for the bug this + // resolves. + const float conv_signed_ema = update_signed_conviction_ema( + b, conviction_signed, + conv_signed_ema_per_b, + conv_signed_diff_var_ema_per_b, + conv_signed_sample_var_ema_per_b); + const float conv_ema_mag = fabsf(conv_signed_ema); + const float conviction_dir = (conv_signed_ema >= 0.0f) ? 1.0f : -1.0f; // CRT.diag Group A: per-horizon direction-flip + run-length tracking. // Mirrors decision_policy_default. Observe-only; single-writer-per-block. @@ -923,9 +988,9 @@ extern "C" __global__ void decision_policy_program( diag_smoothed_sum_run_length, diag_smoothed_prev_dir); - // CRT.diag Group B: smoothed-conviction histogram. + // CRT.diag Group B: smoothed-conviction MAGNITUDE histogram. { - float c = conv_ema; + float c = conv_ema_mag; if (c < 0.0f) c = 0.0f; if (c >= 1.0f) c = 0.999999f; int bk = (int)(c * 10.0f); @@ -936,16 +1001,19 @@ extern "C" __global__ void decision_policy_program( // CRT.1 C1.4: refresh intra-trade trajectory state for composite exit. // RAW alpha_probs kept here — disagreement signal is meant to be fast. + // Passes the magnitude of the signed EMA — composite_exit_check forms + // degradation ratios that only make sense between magnitudes. update_open_trade_trajectory(b, positions, books_per_b, alpha_probs, - conv_ema, open_trade_state_per_b); + conv_ema_mag, open_trade_state_per_b); // CRT.1 C1.4: composite exit_signal safety circuit-breaker (spec §4.3). if (composite_exit_check(b, positions, open_trade_state_per_b, market_targets)) { return; } - // Threshold gate on the smoothed multi-horizon conviction (spec §5.2). - if (conv_ema < threshold_per_b[b]) { + // Threshold gate on the smoothed multi-horizon conviction MAGNITUDE + // (spec §5.2). Pre-registered percentile thresholds are sign-agnostic. + if (conv_ema_mag < threshold_per_b[b]) { market_targets[b * 2 + 0] = 2; market_targets[b * 2 + 1] = 0; open_horizon_masks[b] = 0u; @@ -1103,13 +1171,12 @@ extern "C" __global__ void decision_policy_program( // Pop the VM stack's aggregated value (unused under v3 §4.4 // sizing) and keep the attribution mask for entry crediting. const unsigned int attribution = stack_mask[sp]; - // CRT.1 C1.2: target = direction × smoothed conviction × envelope_max. - // Mirrors decision_policy_default. The v2 `final_size *= scale` - // amplification bug is structurally absent — smoothed conviction - // multiplies the envelope directly, never ratio'd against a raw - // scalar. + // Anti-calibration fix (2026-05-22): target = SIGNED smoothed + // conviction × envelope_max. Sign AND magnitude both come from + // conv_signed_ema — eliminates the prior magnitude/sign mismatch + // where size stayed large while direction flipped on noise. const float envelope_max = (float)max_lots; - const float target_lots_f = conviction_dir * conv_ema * envelope_max; + const float target_lots_f = conv_signed_ema * envelope_max; const int target_lots = (int)truncf( target_lots_f + (target_lots_f >= 0.0f ? 0.5f : -0.5f)); if (target_lots == 0) { diff --git a/crates/ml-backtesting/cuda/pnl_track.cu b/crates/ml-backtesting/cuda/pnl_track.cu index acea14b52..952ede0c5 100644 --- a/crates/ml-backtesting/cuda/pnl_track.cu +++ b/crates/ml-backtesting/cuda/pnl_track.cu @@ -39,8 +39,12 @@ // 8 4 entry_px_x100 (i32) // 12 4 entry_size (i32) // 16 4 realized_at_open (f32) -// 20 4 conviction_at_entry (f32) ← NEW in CRT.1 -// 24 4 conviction_ema (f32) ← NEW in CRT.1 +// 20 4 conviction_at_entry (f32, magnitude — pnl_track +// takes fabsf of the SIGNED +// EMA so degradation ratios +// in composite_exit_check +// compare magnitudes) +// 24 4 conviction_ema (f32, magnitude — same) // 28 16 conviction_per_horizon_ema [4 × f32] ← NEW in CRT.1 // 44 4 pnl_adjusted_conviction_ema (f32) ← NEW in CRT.1 // 48 4 peak_unrealized_pnl (f32) ← NEW in CRT.1 @@ -64,11 +68,14 @@ extern "C" __global__ void pnl_track_step( unsigned int* __restrict__ zero_vwap_at_open, // open branch: vwap_entry == 0 unsigned int* __restrict__ saturated_vwap_at_open, // open branch: |vwap_entry| > 21M or NaN/Inf unsigned int* __restrict__ defensive_exit_clamp, // close branch: defensive clamp fired - // CRT.1 C1.4: current smoothed multi-horizon conviction (from - // decision_policy kernels). Read on open transition to snapshot - // conviction_at_entry + initialize conviction_ema mirror in - // open_trade_state. - const float* __restrict__ conviction_ema_per_b, // [n_backtests] + // Post-fix 2026-05-22: current SIGNED smoothed multi-horizon conviction + // (from decision_policy kernels). Read on open transition to snapshot + // conviction_at_entry + initialize the in-trade conviction_ema mirror + // in open_trade_state. The kernel takes fabsf() before writing because + // composite_exit_check forms a degradation ratio (current/entry) that + // only makes sense as a ratio of magnitudes — a signed flip mid-trade + // would otherwise invert the ratio meaning. + const float* __restrict__ conv_signed_ema_per_b, // [n_backtests] // CRT.diag Group C: hold-time histogram at trade close (6 log buckets). unsigned int* __restrict__ diag_hold_hist, // [n_backtests * 6] // CRT.diag Group D: outcome by entry-conviction (10 conviction buckets). @@ -104,8 +111,13 @@ extern "C" __global__ void pnl_track_step( // don't populate fields nothing reads. Other slots (per-horizon // EMA at 28–43, degradation counter at 52, horizon_idx at 60) // remain zero from the close-branch reset / alloc_zeros init. - const float conv_now = conviction_ema_per_b[b]; - *reinterpret_cast(st + 20) = conv_now; // conviction_at_entry + // + // Post-fix 2026-05-22: the live buffer holds a SIGNED EMA (anti- + // calibration fix in decision_policy.cu); composite_exit_check + // reads offsets 20 and 24 as magnitudes (current/entry ratio is + // only meaningful magnitude-to-magnitude), so we take fabsf here. + const float conv_now = fabsf(conv_signed_ema_per_b[b]); + *reinterpret_cast(st + 20) = conv_now; // conviction_at_entry (magnitude) *reinterpret_cast(st + 24) = conv_now; // conviction_ema (init = entry) *reinterpret_cast(st + 44) = 0.0f; // pnl_adjusted_conviction_ema *reinterpret_cast(st + 48) = 0.0f; // peak_unrealized_pnl diff --git a/crates/ml-backtesting/src/sim/mod.rs b/crates/ml-backtesting/src/sim/mod.rs index 79aaca369..2e7204e19 100644 --- a/crates/ml-backtesting/src/sim/mod.rs +++ b/crates/ml-backtesting/src/sim/mod.rs @@ -286,15 +286,20 @@ pub struct LobSimCuda { mh_kernel_calls_d: CudaSlice, // [n_backtests] mh_force_flat_seen_by_seed_d: CudaSlice, // [n_backtests] - // CRT-A2: Wiener-α adaptive EMA state on max-conviction. Smooths - // event-rate alpha jitter so sizing magnitude doesn't oscillate - // bar-by-bar. All three slots start at zero (sentinel = "no - // observation yet") so the first decision-kernel call replaces them - // directly per pearl_first_observation_bootstrap. Updated in-place - // by both decision_policy_default and decision_policy_program. - conviction_ema_d: CudaSlice, // [n_backtests] — smoothed conviction - conviction_diff_var_ema_d: CudaSlice, // [n_backtests] — 2nd-order EMA of (new−prev)² - conviction_sample_var_ema_d: CudaSlice, // [n_backtests] — 2nd-order EMA of x² + // Anti-calibration fix (2026-05-22): Wiener-α adaptive EMA on the + // SIGNED multi-horizon conviction scalar. The prior magnitude-only + // EMA produced higher-conviction-→-more-negative-PnL anti-calibration + // because the sizing path multiplied `sign(instantaneous) × EMA(|x|)` + // — magnitude and direction came from different signals. Now both + // come from the same smoothed signed EMA: when per-horizon directions + // disagree event-to-event, EMA(signed_x) → 0 → |EMA| → 0 → size 0. + // All three slots start at zero (sentinel = "no observation yet") so + // the first decision-kernel call replaces them directly per + // pearl_first_observation_bootstrap. Updated in-place by both + // decision_policy_default and decision_policy_program. + conv_signed_ema_d: CudaSlice, // [n_backtests] — smoothed signed conviction ∈ [-1, +1] + conv_signed_diff_var_ema_d: CudaSlice, // [n_backtests] — 2nd-order EMA of (new−prev)² + conv_signed_sample_var_ema_d: CudaSlice, // [n_backtests] — 2nd-order EMA of x² (x is signed; x² is non-negative) // CRT.1 C1.3: no-trade band. seed_inflight_limits_batched skips seeding // when |target_signed − effective| < delta_floor[b]. Uploaded once per @@ -588,18 +593,21 @@ impl LobSimCuda { .alloc_zeros::(n_backtests) .context("alloc mh_force_flat_seen_by_seed_d")?; - // CRT-A2: Wiener-α conviction-EMA state. Zero-init means - // "first observation pending" — the decision kernel branches on - // prev_ema == 0 to bootstrap directly (no warmup zero-bias). - let conviction_ema_d = stream + // Anti-calibration fix (2026-05-22): SIGNED conviction-EMA state. + // Zero-init means "first observation pending" — the decision kernel + // branches on prev_ema == 0 to bootstrap directly (no warmup zero- + // bias). The fact that 0.0 is also a legitimate near-disagreement + // outcome is harmless: a first observation that lands at exactly + // zero produces the same bootstrap → zero outcome regardless. + let conv_signed_ema_d = stream .alloc_zeros::(n_backtests) - .context("alloc conviction_ema_d")?; - let conviction_diff_var_ema_d = stream + .context("alloc conv_signed_ema_d")?; + let conv_signed_diff_var_ema_d = stream .alloc_zeros::(n_backtests) - .context("alloc conviction_diff_var_ema_d")?; - let conviction_sample_var_ema_d = stream + .context("alloc conv_signed_diff_var_ema_d")?; + let conv_signed_sample_var_ema_d = stream .alloc_zeros::(n_backtests) - .context("alloc conviction_sample_var_ema_d")?; + .context("alloc conv_signed_sample_var_ema_d")?; // CRT.1 C1.3: no-trade band. Uploaded from BatchedSimConfig in the // config-upload block of step_decision_with_latency. Default zero @@ -761,9 +769,9 @@ impl LobSimCuda { last_bad_path_d, mh_kernel_calls_d, mh_force_flat_seen_by_seed_d, - conviction_ema_d, - conviction_diff_var_ema_d, - conviction_sample_var_ema_d, + conv_signed_ema_d, + conv_signed_diff_var_ema_d, + conv_signed_sample_var_ema_d, delta_floor_d, diag_flip_count_d, diag_current_run_length_d, @@ -917,9 +925,16 @@ impl LobSimCuda { Ok(buf[backtest_idx]) } - /// CRT-A2: read `conviction_ema_d[backtest_idx]`. Test-only accessor — - /// production paths inspect smoothed conviction on-device only. One DtoH. - pub fn read_conviction_ema(&self, backtest_idx: usize) -> Result { + /// Read the SIGNED conviction EMA for one backtest. Returns the raw + /// signed value in [-1, +1]; callers asserting on magnitude take + /// `.abs()`. Test-only accessor — production paths inspect smoothed + /// conviction on-device only. One DtoH per call. + /// + /// Post-fix 2026-05-22: returns the signed EMA, not the magnitude-only + /// EMA from the prior buggy implementation. Sign of the returned value + /// is the trade direction the kernel will commit to on the next event; + /// `|returned|` is the sizing magnitude (post-threshold-gate). + pub fn read_conv_signed_ema(&self, backtest_idx: usize) -> Result { anyhow::ensure!( backtest_idx < self.n_backtests, "backtest_idx {} >= n_backtests {}", @@ -927,7 +942,7 @@ impl LobSimCuda { self.n_backtests ); let mut buf = vec![0.0f32; self.n_backtests]; - self.stream.memcpy_dtoh(&self.conviction_ema_d, buf.as_mut_slice())?; + self.stream.memcpy_dtoh(&self.conv_signed_ema_d, buf.as_mut_slice())?; Ok(buf[backtest_idx]) } @@ -1210,8 +1225,10 @@ impl LobSimCuda { .arg(&mut self.saturated_vwap_at_open_d) .arg(&mut self.defensive_exit_clamp_d) // CRT.1 C1.4: open branch snapshots conviction_at_entry + - // initializes conviction_ema mirror in open_trade_state. - .arg(&self.conviction_ema_d) + // initializes the in-trade conviction_ema magnitude mirror + // in open_trade_state. Post-fix 2026-05-22: kernel reads + // the SIGNED EMA and takes fabsf before writing. + .arg(&self.conv_signed_ema_d) // CRT.diag Groups C + D: hold-time + outcome-by-entry-conviction. .arg(&mut self.diag_hold_hist_d) .arg(&mut self.diag_outcome_n_d) @@ -1492,10 +1509,10 @@ impl LobSimCuda { .arg(&self.books_d) // S2.1: mh_kernel_calls diagnostic counter. .arg(&mut self.mh_kernel_calls_d) - // CRT-A2: Wiener-α conviction-EMA state. - .arg(&mut self.conviction_ema_d) - .arg(&mut self.conviction_diff_var_ema_d) - .arg(&mut self.conviction_sample_var_ema_d) + // Anti-calibration fix 2026-05-22: SIGNED conviction-EMA state. + .arg(&mut self.conv_signed_ema_d) + .arg(&mut self.conv_signed_diff_var_ema_d) + .arg(&mut self.conv_signed_sample_var_ema_d) // CRT.1 C1.4: trajectory state for composite exit_signal. .arg(&mut self.open_trade_state_d) // CRT.diag Groups A + B: signal persistence + conviction histogram. @@ -1541,10 +1558,10 @@ impl LobSimCuda { .arg(&self.books_d) // S2.1: mh_kernel_calls diagnostic counter. .arg(&mut self.mh_kernel_calls_d) - // Wiener-α conviction-EMA state (reused from v2 A2). - .arg(&mut self.conviction_ema_d) - .arg(&mut self.conviction_diff_var_ema_d) - .arg(&mut self.conviction_sample_var_ema_d) + // Anti-calibration fix 2026-05-22: SIGNED conviction-EMA state. + .arg(&mut self.conv_signed_ema_d) + .arg(&mut self.conv_signed_diff_var_ema_d) + .arg(&mut self.conv_signed_sample_var_ema_d) // CRT.1 C1.4: trajectory state for composite exit_signal. .arg(&mut self.open_trade_state_d) // CRT.diag Groups A + B: signal persistence + conviction histogram. @@ -1618,8 +1635,10 @@ impl LobSimCuda { .arg(&mut self.saturated_vwap_at_open_d) .arg(&mut self.defensive_exit_clamp_d) // CRT.1 C1.4: open branch snapshots conviction_at_entry + - // initializes conviction_ema mirror in open_trade_state. - .arg(&self.conviction_ema_d) + // initializes the in-trade conviction_ema magnitude mirror + // in open_trade_state. Post-fix 2026-05-22: kernel reads + // the SIGNED EMA and takes fabsf before writing. + .arg(&self.conv_signed_ema_d) // CRT.diag Groups C + D: hold-time + outcome-by-entry-conviction. .arg(&mut self.diag_hold_hist_d) .arg(&mut self.diag_outcome_n_d) diff --git a/crates/ml-backtesting/tests/stop_controller.rs b/crates/ml-backtesting/tests/stop_controller.rs index 7fed7cd25..140558e88 100644 --- a/crates/ml-backtesting/tests/stop_controller.rs +++ b/crates/ml-backtesting/tests/stop_controller.rs @@ -1251,9 +1251,165 @@ fn multi_horizon_conviction_cancels_on_disagreement() -> Result<()> { "disagreeing horizons must cancel to no-trade; got side={side} size={size}" ); - // Sanity: the on-device EMA should also be (near) zero. - let ema = sim.read_conviction_ema(0)?; - assert!(ema.abs() < 1e-5, "conviction_ema must collapse to ~0 on cancellation; got {ema}"); + // Sanity: the on-device signed EMA should also be (near) zero — + // disagreement cancels in BOTH magnitude (|EMA|→0) and signed value + // (signed→0) under the post-2026-05-22 anti-calibration fix, so the + // strict |ema| < ε check still holds (and is in fact more correct, + // since the test previously asserted on the magnitude-EMA buffer). + let ema = sim.read_conv_signed_ema(0)?; + assert!(ema.abs() < 1e-5, "conv_signed_ema must collapse to ~0 on cancellation; got {ema}"); + Ok(()) +} + +/// Anti-calibration fix (2026-05-22): event-to-event sign flips on the +/// multi-horizon conviction signal must produce a SIGNED EMA whose +/// magnitude is structurally smaller than the per-event |conviction|, +/// AND `target_lots` derived from the signed EMA must be strictly +/// smaller (in magnitude) than the pre-fix size that the bug would have +/// produced. +/// +/// Bug history: the prior magnitude-only EMA combined `sign(instantaneous) +/// × EMA(|instantaneous|)`. When per-horizon directions disagreed +/// event-to-event, the magnitude EMA stayed high (|x| EMA cannot cancel) +/// while the sign flipped on noise → big trades in random direction. +/// Empirically (smoke ckpts across architecture variants): higher reported +/// conviction → MORE negative PnL (-15.6 → -145.2). The fix replaces the +/// magnitude-only EMA with a SIGNED EMA so directions that disagree +/// event-to-event cancel in the EMA — sign AND magnitude both derive from +/// the same smoothed scalar. +/// +/// Test setup: alternate alpha_probs across calls — all horizons bullish +/// (0.85) then all bearish (0.15) — producing a `conviction_signed` that +/// flips between ~+0.7 and ~−0.7 each event. With the Wiener-α floor of +/// 0.4, perfect ±M sign-flip alternation steady-state oscillates with +/// bounded amplitude `M · α / (2 - α)` ≈ 0.175 around zero (NOT exactly +/// to zero — the Wiener α can rise above 0.4 under diff-variance growth, +/// pushing amplitude somewhat higher). What matters structurally: +/// +/// 1. |signed_ema_steady_state| ≪ |conviction_signed_instantaneous| +/// (the EMA cancels what the magnitude EMA cannot) +/// 2. |target_lots_post_fix| < |target_lots_pre_fix| +/// = |sign(instantaneous) × EMA(|x|) × max_lots| +/// = 1 × ~0.7 × 10 = ~7 +/// +/// Both invariants pin the new behavior and would fail under the prior +/// magnitude-only EMA. +#[test] +#[ignore = "requires CUDA"] +fn signed_conviction_ema_collapses_on_sign_flips() -> Result<()> { + let dev = match MlDevice::cuda(0) { + Ok(d) => d, + Err(e) => { eprintln!("skipping: cuda unavailable ({e})"); return Ok(()); } + }; + let mut sim = LobSimCuda::new(1, &dev)?; + sim.upload_price_range(&[1000.0], &[20000.0])?; + let (bp, bs, ap, az) = level_book(5500.0, 0.25); + sim.apply_snapshot(&bp, &bs, &ap, &az)?; + + // threshold = 0.0 so sizing isn't gated off; max_lots = 10 so the + // bug's predicted |target_lots| under steady magnitude EMA = 7 is + // well separated from the post-fix expected magnitude (< 4 at the + // oscillation crest under α floor 0.4). + let max_lots: u16 = 10; + let cfg = BatchedSimConfig::from_uniform(1, &UniformSimParams { + target_annual_vol_units: 50.0, + annualisation_factor: 825.0, + max_lots, + latency_ns: 0, + kelly_frac_floor: 0.20, + sharpe_weight_floor: 0.10, + threshold: 0.0, + cost_per_lot_per_side: 1.0, + max_hold_ns: 0, + min_reasonable_px: 0.0, + max_reasonable_px: f32::INFINITY, + delta_floor: 0.0, + }); + + // All horizons strongly bullish: magnitude ≈ 2·|0.85−0.5| = 0.7, + // direction = +1 across the board → conviction_signed ≈ +0.7. + let bullish: [f32; N_HORIZONS] = [0.85; N_HORIZONS]; + // All horizons strongly bearish: magnitude ≈ 0.7, direction = −1 → + // conviction_signed ≈ −0.7. Per-event sign flip with constant + // magnitude — the exact scenario the bug-trace investigation + // identified (magnitude EMA stays at 0.7, sign flips event-to-event). + let bearish: [f32; N_HORIZONS] = [0.15; N_HORIZONS]; + + // Pre-fix predicted |target_lots| in the TAIL (event >= 1) under the + // magnitude-only EMA: + // round(sign(instantaneous) × EMA(|x|) × max_lots) + // = round(±1 × ~0.7 × 10) = ±7 every tail event. + // Post-fix, the EMA tracks the SIGNED signal. The bootstrap event + // (E0) lands at ±7 in BOTH implementations (first-observation + // replace). What pins the fix is the TAIL behavior: after Wiener + // stabilization, |signed_ema| converges to its steady-state + // oscillation amplitude (~0.467 at α≈0.8, perfectly symmetric ±0.7 + // alternation) so |target_lots| caps at ~5, strictly below the + // bug's tail value of 7. We assert |target| <= 5 across the tail — + // strict enough that the bug (constant 7) trips immediately. + const PRE_FIX_TAIL_ABS: i32 = 7; + const POST_FIX_TAIL_CAP: i32 = 5; // < 7 strictly; bug would always be 7 + + // Drive 16 alternating events. Track per-event signed EMA and + // |target_lots|. Bootstrap event (k=0) is exempt from the magnitude + // comparison (both implementations land at ±7 there). Tail = k >= 4 + // to ensure Wiener α has stabilized past the bootstrap transient. + let mut events: Vec<(f32, i32)> = Vec::with_capacity(16); + for k in 0_u32..16 { + let probs = if k.is_multiple_of(2) { &bullish } else { &bearish }; + sim.broadcast_alpha(probs)?; + sim.step_decision_with_latency(((k as u64) + 1) * 1_000_000_000u64, &cfg)?; + let ema = sim.read_conv_signed_ema(0)?; + let (_side, size) = sim.read_market_target(0)?; + events.push((ema, size)); + } + + // Print per-event diagnostics — helps localize regressions if a + // future implementation drift trips the cap. + for (k, (ema, size)) in events.iter().enumerate() { + eprintln!("event {k:2}: signed_ema = {ema:+.4} |target_lots| = {size}"); + } + + let tail: &[(f32, i32)] = &events[4..]; + + // Invariant 1: tail |target_lots| must be strictly smaller than the + // pre-fix value of 7. This is the load-bearing check — under the bug, + // every tail event would produce |target| = 7. + for (i, (ema, size)) in tail.iter().enumerate() { + assert!( + size.abs() <= POST_FIX_TAIL_CAP, + "tail event {i}: |target_lots| = {size} must be <= {POST_FIX_TAIL_CAP} \ + (pre-fix bug produces {PRE_FIX_TAIL_ABS} every tail event); \ + signed_ema = {ema}" + ); + } + + // Invariant 2: tail signed EMA magnitudes are structurally smaller + // than the per-event |conviction_signed| ≈ 0.7. Allowing a generous + // 0.6 ceiling — the theoretical steady-state amplitude is ~0.467 at + // α≈0.8 under perfect ±M alternation. + for (i, (ema, _)) in tail.iter().enumerate() { + assert!( + ema.abs() < 0.6, + "tail event {i}: |signed EMA| = {ema} must be < 0.6 \ + (per-event |conviction_signed| ≈ 0.7); the EMA is not \ + cancelling sign flips" + ); + } + + // Invariant 3: under perfect ±M alternation the signed EMA itself + // must FLIP SIGN across the tail samples (the EMA tracks the + // alternating mean). The pre-fix magnitude EMA was locked in [0, 1] + // — sign flipping is uniquely the signed-EMA's behavior. + let any_pos = tail.iter().any(|(e, _)| *e > 0.0); + let any_neg = tail.iter().any(|(e, _)| *e < 0.0); + assert!( + any_pos && any_neg, + "tail signed EMA must oscillate across signs under ±M alternation; \ + tail emas = {:?}", + tail.iter().map(|(e, _)| *e).collect::>() + ); + Ok(()) }