From 0e17fd4f2dbeb9fbebcbc0f3fb702bfe0eade078 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Wed, 20 May 2026 22:34:35 +0200 Subject: [PATCH] =?UTF-8?q?diag(crt-2):=20per-horizon=20alpha-input=20EMA?= =?UTF-8?q?=20test=20=E2=80=94=20hypothesis=20investigation?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Driven by ffr59 (commit b44a97ff9) findings: all 5 horizons flip direction every 2.5 events; win rate flat 24% across conviction; mean PnL anti-correlated with conviction. Hypothesis: per-event alpha output is high-frequency noise on top of a slower signal. If true, smoothing input alpha BEFORE the conviction formula should reduce direction flips and recover signal. Adds Wiener-α adaptive EMA on raw alpha_probs[h] for each horizon, applied BEFORE the multi-horizon conviction formula. Floor at 0.1 (stronger than the 0.4 floor on the output-side conviction EMA — this tests whether INPUT smoothing has different impact than OUTPUT smoothing). Three new device slots: - alpha_ema_per_b_per_h (per-horizon EMA state) - alpha_diff_var_per_b_per_h (variance of changes) - alpha_sample_var_per_b_per_h (variance of value) The CRT.diag Group A direction-flip counter still reads RAW alpha_probs so we have a head-to-head comparison: raw flip rate vs smoothed flip rate. Group E adds the smoothed-direction counter + mean run length. End-of-run log adds one line per horizon: crt_diag h smoothed: flips=Y mean_run_len=Z events (vs raw F / M) If smoothed mean_run_len >> raw mean_run_len: hypothesis is RIGHT, the input had signal under the noise. Next step would be to make this an operational EMA in the controller. If smoothed and raw are similar: hypothesis is WRONG, per-event output is genuinely noisy. Next step would be to investigate model training (horizon collapse) OR the AUC=0.66 measurement definition. Either way, definitive result from one smoke run. Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml-backtesting/cuda/decision_policy.cu | 211 +++++++++++++++++- crates/ml-backtesting/src/harness.rs | 21 ++ crates/ml-backtesting/src/sim/mod.rs | 111 ++++++++- 3 files changed, 329 insertions(+), 14 deletions(-) diff --git a/crates/ml-backtesting/cuda/decision_policy.cu b/crates/ml-backtesting/cuda/decision_policy.cu index ea9d383bf..03e9e2a55 100644 --- a/crates/ml-backtesting/cuda/decision_policy.cu +++ b/crates/ml-backtesting/cuda/decision_policy.cu @@ -394,6 +394,127 @@ __device__ static float compute_multi_horizon_conviction( : 0.0f; } +// CRT.diag.2: per-horizon Wiener-α adaptive EMA on RAW alpha probabilities. +// Smooths alpha_probs[h] BEFORE compute_multi_horizon_conviction consumes +// it — testing whether per-event output noise is hiding a slower signal +// (ffr59 found all 5 horizons flip every 2.5 events, win rate flat ~24%, +// mean PnL anti-correlated with conviction; either smoothing recovers +// 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. +// Floor at 0.1 — stronger smoothing than the 0.4 conviction-level floor, +// testing whether INPUT-side smoothing differs from OUTPUT-side. +// +// Writes alpha_smoothed[h] for h in 0..N_HORIZONS; updates the three per- +// horizon EMA-state slots in place. Single-thread-per-backtest invariant: +// caller must already have gated on threadIdx.x == 0. +__device__ static void compute_alpha_ema_per_horizon( + int b, + const float* __restrict__ alpha_probs, + float* __restrict__ alpha_ema_per_b_per_h, + float* __restrict__ alpha_diff_var_per_b_per_h, + float* __restrict__ alpha_sample_var_per_b_per_h, + float (&alpha_smoothed)[N_HORIZONS] +) { + #pragma unroll + for (int h = 0; h < N_HORIZONS; ++h) { + const float p_raw = alpha_probs[h]; + const float prev_ema = alpha_ema_per_b_per_h[b * N_HORIZONS + h]; + const float diff = p_raw - prev_ema; + const float diff_var_prev = alpha_diff_var_per_b_per_h[b * N_HORIZONS + h]; + const float new_diff_var = (diff_var_prev == 0.0f) + ? (diff * diff) // first-observation bootstrap + : (0.9f * diff_var_prev + 0.1f * diff * diff); + alpha_diff_var_per_b_per_h[b * N_HORIZONS + h] = new_diff_var; + + // sample_var approximates the variance of p around its natural + // no-edge mean (0.5). Centering on 0.5 (instead of the running + // mean) gives a Wiener-α that reads "how far from the no-edge + // anchor the signal sits" rather than "how dispersed it is". + const float sample_var_prev = alpha_sample_var_per_b_per_h[b * N_HORIZONS + h]; + const float p_centered = p_raw - 0.5f; + const float new_sample_var = (sample_var_prev == 0.0f) + ? (p_centered * p_centered) // first-observation bootstrap + : (0.9f * sample_var_prev + 0.1f * p_centered * p_centered); + alpha_sample_var_per_b_per_h[b * N_HORIZONS + h] = new_sample_var; + + // Wiener α with FLOOR 0.1 (stronger smoothing than the 0.4 floor + // on the output-side conviction EMA — testing input-side smoothing + // differs from output-side). + 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.1f) ? alpha_raw : 0.1f; + + const float new_ema = (prev_ema == 0.0f) + ? p_raw // first-observation bootstrap + : (alpha_active * p_raw + (1.0f - alpha_active) * prev_ema); + alpha_ema_per_b_per_h[b * N_HORIZONS + h] = new_ema; + alpha_smoothed[h] = new_ema; + } +} + +// CRT.diag.2: variant of compute_multi_horizon_conviction that reads from +// a thread-local alpha_smoothed[h] array instead of the broadcast +// alpha_probs[h] pointer. The §4.4 weight math is identical — only the +// per-horizon p_h source differs. +__device__ static float compute_multi_horizon_conviction_smoothed( + const float (&alpha_smoothed)[N_HORIZONS], + const IsvKellyState* __restrict__ isv, + float cost +) { + const float eps_edge = cost * 0.01f; + + float weighted_sum_signed = 0.0f; + float total_abs_weight = 0.0f; + #pragma unroll + for (int h = 0; h < N_HORIZONS; ++h) { + const float p_h = alpha_smoothed[h]; + const float direction_h = (p_h > 0.5f) ? 1.0f : -1.0f; + const float magnitude_h = fabsf(p_h - 0.5f) * 2.0f; + const float net_edge_h = fmaxf(isv[h].pnl_ema_win - isv[h].pnl_ema_loss, eps_edge); + const float weight_h = net_edge_h / (isv[h].realised_return_var + cost * cost); + weighted_sum_signed += magnitude_h * weight_h * direction_h; + total_abs_weight += fabsf(weight_h); + } + return (total_abs_weight > 1e-9f) + ? (weighted_sum_signed / total_abs_weight) + : 0.0f; +} + +// CRT.diag.2 Group E: per-horizon SMOOTHED direction-flip + mean-run-length +// tracking. Mirrors Group A (which reads RAW alpha_probs[h]) but reads +// alpha_smoothed[h]. Single-writer-per-block; caller must already have +// gated on threadIdx.x == 0. Skips the 5-bucket histogram (Group A keeps +// the histogram; Group E only needs flip count + mean run length for the +// head-to-head comparison). +__device__ static void update_smoothed_flip_counters( + int b, + const float (&alpha_smoothed)[N_HORIZONS], + unsigned int* __restrict__ diag_smoothed_flip_count, + unsigned int* __restrict__ diag_smoothed_current_run_length, + unsigned long long* __restrict__ diag_smoothed_sum_run_length, + signed char* __restrict__ diag_smoothed_prev_dir +) { + #pragma unroll + for (int h = 0; h < N_HORIZONS; ++h) { + const float p = alpha_smoothed[h]; + const signed char new_dir = (p > 0.5f) ? (signed char)1 + : ((p < 0.5f) ? (signed char)-1 : (signed char)0); + const signed char prev_dir = diag_smoothed_prev_dir[b * N_HORIZONS + h]; + const unsigned int prev_run = diag_smoothed_current_run_length[b * N_HORIZONS + h]; + if (prev_dir != 0 && new_dir != 0 && prev_dir != new_dir) { + diag_smoothed_flip_count[b * N_HORIZONS + h] += 1u; + diag_smoothed_sum_run_length[b * N_HORIZONS + h] += (unsigned long long)prev_run; + diag_smoothed_current_run_length[b * N_HORIZONS + h] = 1u; + } else { + diag_smoothed_current_run_length[b * N_HORIZONS + h] = prev_run + 1u; + } + diag_smoothed_prev_dir[b * N_HORIZONS + h] = new_dir; + } +} + extern "C" __global__ void decision_policy_default( const float* __restrict__ alpha_probs, // [N_HORIZONS] broadcast const unsigned char* __restrict__ isv_kelly_base, // [n_backtests * N_HORIZONS * sizeof(IsvKellyState)] @@ -435,6 +556,18 @@ extern "C" __global__ void decision_policy_default( signed char* __restrict__ diag_prev_dir_signed, // [n_backtests * N_HORIZONS] // CRT.diag Group B: conviction-EMA histogram. unsigned int* __restrict__ diag_conv_hist, // [n_backtests * 10] + // CRT.diag.2: per-horizon alpha-input EMA state. Smooths raw + // alpha_probs[h] BEFORE the §4.4 conviction formula consumes it. + // Three slots mirror the conviction-level EMA pattern. + float* __restrict__ alpha_ema_per_b_per_h, // [n_backtests * N_HORIZONS] + float* __restrict__ alpha_diff_var_per_b_per_h, // [n_backtests * N_HORIZONS] + float* __restrict__ alpha_sample_var_per_b_per_h, // [n_backtests * N_HORIZONS] + // CRT.diag.2 Group E: smoothed-direction-flip counters (mirror Group A + // but reading alpha_smoothed[h] instead of alpha_probs[h]). No histogram. + unsigned int* __restrict__ diag_smoothed_flip_count, // [n_backtests * N_HORIZONS] + unsigned int* __restrict__ diag_smoothed_current_run_length, // [n_backtests * N_HORIZONS] + unsigned long long* __restrict__ diag_smoothed_sum_run_length, // [n_backtests * N_HORIZONS] + signed char* __restrict__ diag_smoothed_prev_dir, // [n_backtests * N_HORIZONS] int n_backtests ) { int b = blockIdx.x; @@ -450,18 +583,36 @@ extern "C" __global__ void decision_policy_default( return; } + // CRT.diag.2: smooth RAW alpha_probs[h] per-horizon BEFORE the §4.4 + // conviction formula consumes it. Wiener-α EMA with floor 0.1 (stronger + // than the 0.4 output-side floor). The alpha_smoothed[] array stays in + // registers — fed into both compute_multi_horizon_conviction_smoothed + // and the Group E flip counter. RAW alpha_probs[] is still passed + // through to Group A so the head-to-head comparison is apples-to-apples. + float alpha_smoothed[N_HORIZONS]; + compute_alpha_ema_per_horizon( + b, alpha_probs, + alpha_ema_per_b_per_h, + alpha_diff_var_per_b_per_h, + alpha_sample_var_per_b_per_h, + alpha_smoothed); + // CRT.1 C1.2: multi-horizon ISV-weighted conviction (spec §4.4). // Replaces v2 A2 (1d889d2de) + A2.1 (fe2498769) scalar approach that // failed Gate 1 catastrophically on smokes vjmwc + lkrdf (155k trades, // 9100% drawdown, Sharpe -15.7). Direction now comes from the // per-horizon SIGNED sum so disagreeing horizons cancel — the // structural fix the scalar EMA could not provide. + // + // CRT.diag.2: the §4.4 formula now reads alpha_smoothed[] (per-horizon + // EMA) instead of alpha_probs[]. Hypothesis test: does input smoothing + // recover signal hidden under per-event noise? const IsvKellyState* isv = reinterpret_cast( isv_kelly_base + (size_t)b * N_HORIZONS * sizeof(IsvKellyState) ); const float cost = cost_per_lot_per_side_per_b[b]; - const float conviction_signed = compute_multi_horizon_conviction( - alpha_probs, isv, cost); + 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; @@ -477,6 +628,8 @@ extern "C" __global__ void decision_policy_default( // CRT.diag Group A: per-horizon direction-flip + run-length tracking. // Observe-only. Single-writer-per-block (already gated on thread 0). + // Reads RAW alpha_probs[h] — head-to-head comparison with Group E + // (which reads alpha_smoothed[h]). Apples-to-apples vs ffr59 baseline. #pragma unroll for (int h = 0; h < N_HORIZONS; ++h) { const float p = alpha_probs[h]; @@ -502,6 +655,15 @@ extern "C" __global__ void decision_policy_default( diag_prev_dir_signed[b * N_HORIZONS + h] = new_dir; } + // CRT.diag.2 Group E: per-horizon SMOOTHED direction-flip + mean-run-length. + // Mirrors Group A reading alpha_smoothed[] instead of alpha_probs[]. + update_smoothed_flip_counters( + b, alpha_smoothed, + diag_smoothed_flip_count, + diag_smoothed_current_run_length, + diag_smoothed_sum_run_length, + diag_smoothed_prev_dir); + // CRT.diag Group B: smoothed-conviction histogram (10 buckets over [0, 1)). { float c = conv_ema; @@ -515,6 +677,9 @@ extern "C" __global__ void decision_policy_default( // CRT.1 C1.4: refresh intra-trade trajectory state (offsets 24, 44, 48, // 56) for the composite exit_signal. Skipped internally when flat. + // 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. update_open_trade_trajectory(b, positions, books_per_b, alpha_probs, conv_ema, open_trade_state_per_b); @@ -539,11 +704,13 @@ extern "C" __global__ void decision_policy_default( // Build per-horizon attribution mask: every horizon whose signed // contribution agrees with the multi-horizon direction gets credit. // This replaces the v2 sharpe-weight attribution because the v3 - // weights come from the §4.4 formula directly. + // weights come from the §4.4 formula directly. CRT.diag.2: reads + // alpha_smoothed[] so attribution matches the smoothed conviction the + // size formula used. unsigned int attribution_mask = 0u; #pragma unroll for (int h = 0; h < N_HORIZONS; ++h) { - const float dir_h = (alpha_probs[h] > 0.5f) ? 1.0f : -1.0f; + const float dir_h = (alpha_smoothed[h] > 0.5f) ? 1.0f : -1.0f; if (dir_h == conviction_dir) attribution_mask |= (1u << h); } @@ -652,6 +819,15 @@ extern "C" __global__ void decision_policy_program( signed char* __restrict__ diag_prev_dir_signed, // [n_backtests * N_HORIZONS] // CRT.diag Group B: conviction-EMA histogram. unsigned int* __restrict__ diag_conv_hist, // [n_backtests * 10] + // CRT.diag.2: per-horizon alpha-input EMA state. + float* __restrict__ alpha_ema_per_b_per_h, // [n_backtests * N_HORIZONS] + float* __restrict__ alpha_diff_var_per_b_per_h, // [n_backtests * N_HORIZONS] + float* __restrict__ alpha_sample_var_per_b_per_h, // [n_backtests * N_HORIZONS] + // CRT.diag.2 Group E: smoothed-direction-flip counters. + unsigned int* __restrict__ diag_smoothed_flip_count, // [n_backtests * N_HORIZONS] + unsigned int* __restrict__ diag_smoothed_current_run_length, // [n_backtests * N_HORIZONS] + unsigned long long* __restrict__ diag_smoothed_sum_run_length, // [n_backtests * N_HORIZONS] + signed char* __restrict__ diag_smoothed_prev_dir, // [n_backtests * N_HORIZONS] int n_backtests ) { int b = blockIdx.x; @@ -672,6 +848,16 @@ extern "C" __global__ void decision_policy_program( return; } + // CRT.diag.2: smooth RAW alpha_probs[h] per-horizon BEFORE the §4.4 + // conviction formula consumes it. Mirrors decision_policy_default. + float alpha_smoothed[N_HORIZONS]; + compute_alpha_ema_per_horizon( + b, alpha_probs, + alpha_ema_per_b_per_h, + alpha_diff_var_per_b_per_h, + alpha_sample_var_per_b_per_h, + alpha_smoothed); + // CRT.1 C1.2: multi-horizon ISV-weighted conviction (spec §4.4). // Mirrors decision_policy_default — see that kernel for the full // rationale. The VM's per-horizon emit + aggregate opcodes still @@ -679,6 +865,9 @@ extern "C" __global__ void decision_policy_program( // signed `final_size` and uses the §4.4 conviction-driven target // instead. The attribution mask from the stack IS still consumed // at OP_WRITE_ORDER so the VM continues to credit horizons on entry. + // + // CRT.diag.2: feeds the §4.4 formula alpha_smoothed[] instead of + // alpha_probs[]. const Instruction* prog = programs + (size_t)b * max_instructions; const IsvKellyState* isv = reinterpret_cast( isv_kelly_base + (size_t)b * N_HORIZONS * sizeof(IsvKellyState) @@ -686,8 +875,8 @@ extern "C" __global__ void decision_policy_program( const unsigned int my_regime = regimes[b]; const float cost = cost_per_lot_per_side_per_b[b]; - const float conviction_signed = compute_multi_horizon_conviction( - alpha_probs, isv, cost); + 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; @@ -701,6 +890,7 @@ extern "C" __global__ void decision_policy_program( // CRT.diag Group A: per-horizon direction-flip + run-length tracking. // Mirrors decision_policy_default. Observe-only; single-writer-per-block. + // Reads RAW alpha_probs[h] — head-to-head comparison with Group E. #pragma unroll for (int h = 0; h < N_HORIZONS; ++h) { const float p = alpha_probs[h]; @@ -725,6 +915,14 @@ extern "C" __global__ void decision_policy_program( diag_prev_dir_signed[b * N_HORIZONS + h] = new_dir; } + // CRT.diag.2 Group E: smoothed-direction-flip + mean-run-length. + update_smoothed_flip_counters( + b, alpha_smoothed, + diag_smoothed_flip_count, + diag_smoothed_current_run_length, + diag_smoothed_sum_run_length, + diag_smoothed_prev_dir); + // CRT.diag Group B: smoothed-conviction histogram. { float c = conv_ema; @@ -737,6 +935,7 @@ 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. update_open_trade_trajectory(b, positions, books_per_b, alpha_probs, conv_ema, open_trade_state_per_b); diff --git a/crates/ml-backtesting/src/harness.rs b/crates/ml-backtesting/src/harness.rs index 30a803f47..c31c72d36 100644 --- a/crates/ml-backtesting/src/harness.rs +++ b/crates/ml-backtesting/src/harness.rs @@ -392,6 +392,27 @@ impl BacktestHarness { hist_str.push_str(&format!(" {}:{}", bucket_labels[bk], count)); } eprintln!("crt_diag h{} run_length_hist:{}", horizons[h], hist_str); + + // CRT.diag.2 Group E: SMOOTHED-direction flip rate head-to-head + // vs raw. If smoothed mean_run_len ≫ raw mean_run_len, the + // hypothesis (per-event output noise masks a slower signal) + // is supported. If similar, hypothesis is falsified. + let smoothed_flips: u64 = (0..n_b) + .map(|b| diag.smoothed_flip_count[b * n_horizons + h] as u64) + .sum(); + let smoothed_sum_run: u64 = (0..n_b) + .map(|b| diag.smoothed_sum_run_length[b * n_horizons + h]) + .sum(); + let smoothed_mean_run = if smoothed_flips > 0 { + smoothed_sum_run as f64 / smoothed_flips as f64 + } else { + f64::NAN + }; + eprintln!( + "crt_diag h{} smoothed: flips={} mean_run_len={:.1} events (vs raw {} / {:.1})", + horizons[h], smoothed_flips, smoothed_mean_run, + flips, mean_run, + ); } // Group B — smoothed-conviction histogram. diff --git a/crates/ml-backtesting/src/sim/mod.rs b/crates/ml-backtesting/src/sim/mod.rs index babf181c6..4b0648ddb 100644 --- a/crates/ml-backtesting/src/sim/mod.rs +++ b/crates/ml-backtesting/src/sim/mod.rs @@ -124,6 +124,14 @@ pub struct CrtDiagnostics { /// Win count per entry-conviction bucket (segment_realized > 0). /// Layout: `[backtest * 10 + bucket]`. pub outcome_n_wins: Vec, + /// CRT.diag.2 Group E: per-horizon SMOOTHED-direction flip count. + /// Layout: `[backtest * N_HORIZONS + h]`. Compare head-to-head with + /// `flip_count` (which counts RAW alpha_probs direction flips). + pub smoothed_flip_count: Vec, + /// CRT.diag.2 Group E: per-horizon cumulative SMOOTHED run length over + /// completed runs (u64 to avoid overflow at multi-million-event smoke). + /// Layout: `[backtest * N_HORIZONS + h]`. + pub smoothed_sum_run_length: Vec, } /// S2.1/S2.2: max_hold enforcement diagnostic counters. @@ -313,6 +321,28 @@ pub struct LobSimCuda { pub(crate) diag_outcome_n_d: CudaSlice, // [n_backtests * 10] pub(crate) diag_outcome_sum_pnl_d: CudaSlice, // [n_backtests * 10] pub(crate) diag_outcome_n_wins_d: CudaSlice, // [n_backtests * 10] + + // CRT.diag.2: per-horizon Wiener-α EMA on raw alpha_probs. Smoothed + // alpha enters the multi-horizon §4.4 conviction formula instead of + // the raw broadcast values — hypothesis test for whether per-event + // output noise is hiding a slower signal. All zero-init: sentinel "no + // observation yet" so the first decision-kernel call bootstraps per + // pearl_first_observation_bootstrap. Floor on Wiener α = 0.1 (stronger + // smoothing than the 0.4 output-side floor; testing input vs output + // smoothing differ). Touched by both decision_policy_default and + // decision_policy_program every event; no host roundtrip. + pub(crate) alpha_ema_per_b_per_h_d: CudaSlice, // [n_backtests * N_HORIZONS] + pub(crate) alpha_diff_var_per_b_per_h_d: CudaSlice, // [n_backtests * N_HORIZONS] + pub(crate) alpha_sample_var_per_b_per_h_d: CudaSlice, // [n_backtests * N_HORIZONS] + + // CRT.diag.2 Group E: per-horizon SMOOTHED direction-flip + mean-run-length. + // Mirrors Group A but reading alpha_smoothed[h] instead of alpha_probs[h]. + // Skipping the 5-bucket histogram (Group A keeps it). Head-to-head + // comparison: raw flip rate vs smoothed flip rate, per horizon. + pub(crate) diag_smoothed_flip_count_d: CudaSlice, // [n_backtests * N_HORIZONS] + pub(crate) diag_smoothed_current_run_length_d: CudaSlice, // [n_backtests * N_HORIZONS] + pub(crate) diag_smoothed_sum_run_length_d: CudaSlice, // [n_backtests * N_HORIZONS] + pub(crate) diag_smoothed_prev_dir_d: CudaSlice, // [n_backtests * N_HORIZONS] } impl LobSimCuda { @@ -611,6 +641,35 @@ impl LobSimCuda { .alloc_zeros::(n_backtests * 10) .context("alloc diag_outcome_n_wins_d")?; + // CRT.diag.2: per-horizon alpha-input EMA state. Zero-init = "no + // observation yet" — first decision-kernel call bootstraps per + // pearl_first_observation_bootstrap. Same pattern as the + // conviction-level Wiener-α EMA above (zero-init = sentinel). + let alpha_ema_per_b_per_h_d = stream + .alloc_zeros::(n_backtests * N_HORIZONS) + .context("alloc alpha_ema_per_b_per_h_d")?; + let alpha_diff_var_per_b_per_h_d = stream + .alloc_zeros::(n_backtests * N_HORIZONS) + .context("alloc alpha_diff_var_per_b_per_h_d")?; + let alpha_sample_var_per_b_per_h_d = stream + .alloc_zeros::(n_backtests * N_HORIZONS) + .context("alloc alpha_sample_var_per_b_per_h_d")?; + + // CRT.diag.2 Group E: per-horizon SMOOTHED direction-flip counters. + // Same shape as Group A's flip/run-length slots (minus the histogram). + let diag_smoothed_flip_count_d = stream + .alloc_zeros::(n_backtests * N_HORIZONS) + .context("alloc diag_smoothed_flip_count_d")?; + let diag_smoothed_current_run_length_d = stream + .alloc_zeros::(n_backtests * N_HORIZONS) + .context("alloc diag_smoothed_current_run_length_d")?; + let diag_smoothed_sum_run_length_d = stream + .alloc_zeros::(n_backtests * N_HORIZONS) + .context("alloc diag_smoothed_sum_run_length_d")?; + let diag_smoothed_prev_dir_d = stream + .alloc_zeros::(n_backtests * N_HORIZONS) + .context("alloc diag_smoothed_prev_dir_d")?; + // CRT Phase A0.5 corrective: on-device conviction history buffer. // Capacity = 5M decisions × 4B/f32 = 20MB. Matches the prior // host-side `Vec::with_capacity(3_000_000)` plus headroom for @@ -715,6 +774,13 @@ impl LobSimCuda { diag_outcome_n_d, diag_outcome_sum_pnl_d, diag_outcome_n_wins_d, + alpha_ema_per_b_per_h_d, + alpha_diff_var_per_b_per_h_d, + alpha_sample_var_per_b_per_h_d, + diag_smoothed_flip_count_d, + diag_smoothed_current_run_length_d, + diag_smoothed_sum_run_length_d, + diag_smoothed_prev_dir_d, }) } @@ -953,14 +1019,17 @@ impl LobSimCuda { /// touches none of these. pub fn read_diagnostics(&self) -> Result { let n_b = self.n_backtests; - let mut flip_count = vec![0u32; n_b * N_HORIZONS]; - let mut sum_run_length = vec![0u64; n_b * N_HORIZONS]; - let mut run_length_hist = vec![0u32; n_b * N_HORIZONS * 5]; - let mut conv_hist = vec![0u32; n_b * 10]; - let mut hold_hist = vec![0u32; n_b * 6]; - let mut outcome_n = vec![0u32; n_b * 10]; - let mut outcome_sum_pnl = vec![0f32; n_b * 10]; - let mut outcome_n_wins = vec![0u32; n_b * 10]; + let mut flip_count = vec![0u32; n_b * N_HORIZONS]; + let mut sum_run_length = vec![0u64; n_b * N_HORIZONS]; + let mut run_length_hist = vec![0u32; n_b * N_HORIZONS * 5]; + let mut conv_hist = vec![0u32; n_b * 10]; + let mut hold_hist = vec![0u32; n_b * 6]; + let mut outcome_n = vec![0u32; n_b * 10]; + let mut outcome_sum_pnl = vec![0f32; n_b * 10]; + let mut outcome_n_wins = vec![0u32; n_b * 10]; + // CRT.diag.2 Group E: smoothed-direction flip/run-length. + let mut smoothed_flip_count = vec![0u32; n_b * N_HORIZONS]; + let mut smoothed_sum_run_length = vec![0u64; n_b * N_HORIZONS]; self.stream.memcpy_dtoh(&self.diag_flip_count_d, flip_count.as_mut_slice()) .context("diag_flip_count DtoH")?; self.stream.memcpy_dtoh(&self.diag_sum_run_length_d, sum_run_length.as_mut_slice()) @@ -977,6 +1046,12 @@ impl LobSimCuda { .context("diag_outcome_sum_pnl DtoH")?; self.stream.memcpy_dtoh(&self.diag_outcome_n_wins_d, outcome_n_wins.as_mut_slice()) .context("diag_outcome_n_wins DtoH")?; + self.stream.memcpy_dtoh(&self.diag_smoothed_flip_count_d, + smoothed_flip_count.as_mut_slice()) + .context("diag_smoothed_flip_count DtoH")?; + self.stream.memcpy_dtoh(&self.diag_smoothed_sum_run_length_d, + smoothed_sum_run_length.as_mut_slice()) + .context("diag_smoothed_sum_run_length DtoH")?; Ok(CrtDiagnostics { flip_count, sum_run_length, @@ -986,6 +1061,8 @@ impl LobSimCuda { outcome_n, outcome_sum_pnl, outcome_n_wins, + smoothed_flip_count, + smoothed_sum_run_length, }) } @@ -1392,6 +1469,15 @@ impl LobSimCuda { .arg(&mut self.diag_run_length_hist_d) .arg(&mut self.diag_prev_dir_signed_d) .arg(&mut self.diag_conv_hist_d) + // CRT.diag.2: per-horizon alpha-input EMA state. + .arg(&mut self.alpha_ema_per_b_per_h_d) + .arg(&mut self.alpha_diff_var_per_b_per_h_d) + .arg(&mut self.alpha_sample_var_per_b_per_h_d) + // CRT.diag.2 Group E: smoothed-direction flip counters. + .arg(&mut self.diag_smoothed_flip_count_d) + .arg(&mut self.diag_smoothed_current_run_length_d) + .arg(&mut self.diag_smoothed_sum_run_length_d) + .arg(&mut self.diag_smoothed_prev_dir_d) .arg(&n) .launch(cfg)?; } @@ -1432,6 +1518,15 @@ impl LobSimCuda { .arg(&mut self.diag_run_length_hist_d) .arg(&mut self.diag_prev_dir_signed_d) .arg(&mut self.diag_conv_hist_d) + // CRT.diag.2: per-horizon alpha-input EMA state. + .arg(&mut self.alpha_ema_per_b_per_h_d) + .arg(&mut self.alpha_diff_var_per_b_per_h_d) + .arg(&mut self.alpha_sample_var_per_b_per_h_d) + // CRT.diag.2 Group E: smoothed-direction flip counters. + .arg(&mut self.diag_smoothed_flip_count_d) + .arg(&mut self.diag_smoothed_current_run_length_d) + .arg(&mut self.diag_smoothed_sum_run_length_d) + .arg(&mut self.diag_smoothed_prev_dir_d) .arg(&n) .launch(cfg)?; }