gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

fix(decision-policy): use signed-conviction EMA, not magnitude-only EMA

Browse Source 
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-22 00:25:45 +02:00
parent fe5d7a1ad7
commit 0384f76743
4 changed files with 407 additions and 153 deletions
Download Patch File Download Diff File Expand all files Collapse all files

275
crates/ml-backtesting/cuda/decision_policy.cu
View File

@@ -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<float*>(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<float*>(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<const float*>(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) {

30
crates/ml-backtesting/cuda/pnl_track.cu
View File

@@ -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<float*>(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<float*>(st + 20) = conv_now; // conviction_at_entry (magnitude)
*reinterpret_cast<float*>(st + 24) = conv_now; // conviction_ema (init = entry)
*reinterpret_cast<float*>(st + 44) = 0.0f; // pnl_adjusted_conviction_ema
*reinterpret_cast<float*>(st + 48) = 0.0f; // peak_unrealized_pnl

93
crates/ml-backtesting/src/sim/mod.rs
View File

@@ -286,15 +286,20 @@ pub struct LobSimCuda {
mh_kernel_calls_d: CudaSlice<u32>, // [n_backtests]
mh_force_flat_seen_by_seed_d: CudaSlice<u32>, // [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<f32>, // [n_backtests] — smoothed conviction
conviction_diff_var_ema_d: CudaSlice<f32>, // [n_backtests] — 2nd-order EMA of (new−prev)²
conviction_sample_var_ema_d: CudaSlice<f32>, // [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<f32>, // [n_backtests] — smoothed signed conviction ∈ [-1, +1]
conv_signed_diff_var_ema_d: CudaSlice<f32>, // [n_backtests] — 2nd-order EMA of (new−prev)²
conv_signed_sample_var_ema_d: CudaSlice<f32>, // [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::<u32>(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::<f32>(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::<f32>(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::<f32>(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<f32> {
/// 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<f32> {
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)

162
crates/ml-backtesting/tests/stop_controller.rs
View File

@@ -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::<Vec<_>>()
);
Ok(())
}