diag(crt-2): per-horizon alpha-input EMA test — hypothesis investigation

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<X> 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) <noreply@anthropic.com>

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<const IsvKellyState*>(
         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<const IsvKellyState*>(
         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);
 

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.

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<u32>,
+    /// 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<u32>,
+    /// 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<u64>,
 }
 
 /// S2.1/S2.2: max_hold enforcement diagnostic counters.
@@ -313,6 +321,28 @@ pub struct LobSimCuda {
     pub(crate) diag_outcome_n_d:            CudaSlice<u32>,  // [n_backtests * 10]
     pub(crate) diag_outcome_sum_pnl_d:      CudaSlice<f32>,  // [n_backtests * 10]
     pub(crate) diag_outcome_n_wins_d:       CudaSlice<u32>,  // [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<f32>,  // [n_backtests * N_HORIZONS]
+    pub(crate) alpha_diff_var_per_b_per_h_d:   CudaSlice<f32>,  // [n_backtests * N_HORIZONS]
+    pub(crate) alpha_sample_var_per_b_per_h_d: CudaSlice<f32>,  // [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<u32>,  // [n_backtests * N_HORIZONS]
+    pub(crate) diag_smoothed_current_run_length_d: CudaSlice<u32>,  // [n_backtests * N_HORIZONS]
+    pub(crate) diag_smoothed_sum_run_length_d:     CudaSlice<u64>,  // [n_backtests * N_HORIZONS]
+    pub(crate) diag_smoothed_prev_dir_d:           CudaSlice<i8>,   // [n_backtests * N_HORIZONS]
 }
 
 impl LobSimCuda {
@@ -611,6 +641,35 @@ impl LobSimCuda {
             .alloc_zeros::<u32>(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::<f32>(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::<f32>(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::<f32>(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::<u32>(n_backtests * N_HORIZONS)
+            .context("alloc diag_smoothed_flip_count_d")?;
+        let diag_smoothed_current_run_length_d = stream
+            .alloc_zeros::<u32>(n_backtests * N_HORIZONS)
+            .context("alloc diag_smoothed_current_run_length_d")?;
+        let diag_smoothed_sum_run_length_d = stream
+            .alloc_zeros::<u64>(n_backtests * N_HORIZONS)
+            .context("alloc diag_smoothed_sum_run_length_d")?;
+        let diag_smoothed_prev_dir_d = stream
+            .alloc_zeros::<i8>(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<CrtDiagnostics> {
         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)?;
         }