diff --git a/crates/ml-backtesting/cuda/decision_policy.cu b/crates/ml-backtesting/cuda/decision_policy.cu index f80368228..ea9d383bf 100644 --- a/crates/ml-backtesting/cuda/decision_policy.cu +++ b/crates/ml-backtesting/cuda/decision_policy.cu @@ -425,6 +425,16 @@ extern "C" __global__ void decision_policy_default( // this kernel updates the intra-trade fields each event so the composite // exit check has fresh inputs on the next pass. unsigned char* __restrict__ open_trade_state_per_b, // [n_backtests * 64] + // CRT.diag Group A: per-horizon signal persistence counters. + // Observe-only; zero behavior impact. Single-writer-per-block + // (caller already gated on threadIdx.x == 0). End-of-run dump only. + unsigned int* __restrict__ diag_flip_count, // [n_backtests * N_HORIZONS] + unsigned int* __restrict__ diag_current_run_length, // [n_backtests * N_HORIZONS] + unsigned long long* __restrict__ diag_sum_run_length, // [n_backtests * N_HORIZONS] + unsigned int* __restrict__ diag_run_length_hist, // [n_backtests * N_HORIZONS * 5] + 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] int n_backtests ) { int b = blockIdx.x; @@ -465,6 +475,44 @@ extern "C" __global__ void decision_policy_default( conviction_diff_var_ema_per_b, conviction_sample_var_ema_per_b); + // CRT.diag Group A: per-horizon direction-flip + run-length tracking. + // Observe-only. Single-writer-per-block (already gated on thread 0). + #pragma unroll + for (int h = 0; h < N_HORIZONS; ++h) { + const float p = alpha_probs[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_prev_dir_signed[b * N_HORIZONS + h]; + const unsigned int prev_run = diag_current_run_length[b * N_HORIZONS + h]; + if (prev_dir != 0 && new_dir != 0 && prev_dir != new_dir) { + // Direction flipped — record the prev run length. + diag_flip_count[b * N_HORIZONS + h] += 1u; + diag_sum_run_length[b * N_HORIZONS + h] += (unsigned long long)prev_run; + int bk; + if (prev_run < 10u) bk = 0; + else if (prev_run < 100u) bk = 1; + else if (prev_run < 1000u) bk = 2; + else if (prev_run < 10000u) bk = 3; + else bk = 4; + diag_run_length_hist[(b * N_HORIZONS + h) * 5 + bk] += 1u; + diag_current_run_length[b * N_HORIZONS + h] = 1u; // start new run + } else { + diag_current_run_length[b * N_HORIZONS + h] = prev_run + 1u; + } + diag_prev_dir_signed[b * N_HORIZONS + h] = new_dir; + } + + // CRT.diag Group B: smoothed-conviction histogram (10 buckets over [0, 1)). + { + float c = conv_ema; + if (c < 0.0f) c = 0.0f; + if (c >= 1.0f) c = 0.999999f; + int bk = (int)(c * 10.0f); + if (bk < 0) bk = 0; + if (bk > 9) bk = 9; + diag_conv_hist[b * 10 + bk] += 1u; + } + // CRT.1 C1.4: refresh intra-trade trajectory state (offsets 24, 44, 48, // 56) for the composite exit_signal. Skipped internally when flat. update_open_trade_trajectory(b, positions, books_per_b, alpha_probs, @@ -596,6 +644,14 @@ extern "C" __global__ void decision_policy_program( // 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] + // CRT.diag Group A: per-horizon signal persistence counters (observe-only). + unsigned int* __restrict__ diag_flip_count, // [n_backtests * N_HORIZONS] + unsigned int* __restrict__ diag_current_run_length, // [n_backtests * N_HORIZONS] + unsigned long long* __restrict__ diag_sum_run_length, // [n_backtests * N_HORIZONS] + unsigned int* __restrict__ diag_run_length_hist, // [n_backtests * N_HORIZONS * 5] + 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] int n_backtests ) { int b = blockIdx.x; @@ -643,6 +699,43 @@ extern "C" __global__ void decision_policy_program( conviction_diff_var_ema_per_b, conviction_sample_var_ema_per_b); + // CRT.diag Group A: per-horizon direction-flip + run-length tracking. + // Mirrors decision_policy_default. Observe-only; single-writer-per-block. + #pragma unroll + for (int h = 0; h < N_HORIZONS; ++h) { + const float p = alpha_probs[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_prev_dir_signed[b * N_HORIZONS + h]; + const unsigned int prev_run = diag_current_run_length[b * N_HORIZONS + h]; + if (prev_dir != 0 && new_dir != 0 && prev_dir != new_dir) { + diag_flip_count[b * N_HORIZONS + h] += 1u; + diag_sum_run_length[b * N_HORIZONS + h] += (unsigned long long)prev_run; + int bk; + if (prev_run < 10u) bk = 0; + else if (prev_run < 100u) bk = 1; + else if (prev_run < 1000u) bk = 2; + else if (prev_run < 10000u) bk = 3; + else bk = 4; + diag_run_length_hist[(b * N_HORIZONS + h) * 5 + bk] += 1u; + diag_current_run_length[b * N_HORIZONS + h] = 1u; + } else { + diag_current_run_length[b * N_HORIZONS + h] = prev_run + 1u; + } + diag_prev_dir_signed[b * N_HORIZONS + h] = new_dir; + } + + // CRT.diag Group B: smoothed-conviction histogram. + { + float c = conv_ema; + if (c < 0.0f) c = 0.0f; + if (c >= 1.0f) c = 0.999999f; + int bk = (int)(c * 10.0f); + if (bk < 0) bk = 0; + if (bk > 9) bk = 9; + diag_conv_hist[b * 10 + bk] += 1u; + } + // CRT.1 C1.4: refresh intra-trade trajectory state for composite exit. update_open_trade_trajectory(b, positions, books_per_b, alpha_probs, conv_ema, open_trade_state_per_b); diff --git a/crates/ml-backtesting/cuda/pnl_track.cu b/crates/ml-backtesting/cuda/pnl_track.cu index b36b753d1..acea14b52 100644 --- a/crates/ml-backtesting/cuda/pnl_track.cu +++ b/crates/ml-backtesting/cuda/pnl_track.cu @@ -68,7 +68,13 @@ extern "C" __global__ void pnl_track_step( // 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] + const float* __restrict__ conviction_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). + unsigned int* __restrict__ diag_outcome_n, // [n_backtests * 10] + float* __restrict__ diag_outcome_sum_pnl, // [n_backtests * 10] — segment_realized (price-units × lots) + unsigned int* __restrict__ diag_outcome_n_wins // [n_backtests * 10] ) { int b = blockIdx.x; if (b >= n_backtests || threadIdx.x != 0) return; @@ -120,6 +126,40 @@ extern "C" __global__ void pnl_track_step( // Realized P&L delta = pos.realized_pnl_now − pos.realized_pnl_at_open. // In price-units × lots. const float segment_realized = pos->realized_pnl - realized_at_open; + + // CRT.diag Group C: hold-time histogram. Observe-only. The buckets + // (< 1s, 1-10s, 10-60s, 1-10m, 10-60m, > 1h) capture log-scale + // distribution of how long trades stay open. + { + const unsigned long long hold_ns = current_ts_ns - entry_ts; + const double hold_s = (double)hold_ns / 1.0e9; + int bk; + if (hold_s < 1.0) bk = 0; + else if (hold_s < 10.0) bk = 1; + else if (hold_s < 60.0) bk = 2; + else if (hold_s < 600.0) bk = 3; + else if (hold_s < 3600.0) bk = 4; + else bk = 5; + diag_hold_hist[b * 6 + bk] += 1u; + } + + // CRT.diag Group D: outcome by entry-conviction (10 conviction + // buckets over [0, 1)). Reads conviction_at_entry from open_trade_state + // (offset 20, populated by the open branch above on a prior call). + { + const float conv_at_entry = *reinterpret_cast(st + 20); + float c = conv_at_entry; + if (c < 0.0f) c = 0.0f; + if (c >= 1.0f) c = 0.999999f; + int bk = (int)(c * 10.0f); + if (bk < 0) bk = 0; + if (bk > 9) bk = 9; + diag_outcome_n[b * 10 + bk] += 1u; + diag_outcome_sum_pnl[b * 10 + bk] += segment_realized; + if (segment_realized > 0.0f) { + diag_outcome_n_wins[b * 10 + bk] += 1u; + } + } // Convert to USD ×100 fixed-point: × $50/index-point × 100 = ×5000. const int realised_usd_fp = (int)(segment_realized * 5000.0f + (segment_realized >= 0.0f ? 0.5f : -0.5f)); // Exit price reconstructed from realized delta: realized = (exit − entry) × dir × |size|. diff --git a/crates/ml-backtesting/src/harness.rs b/crates/ml-backtesting/src/harness.rs index bb33ea0f9..30a803f47 100644 --- a/crates/ml-backtesting/src/harness.rs +++ b/crates/ml-backtesting/src/harness.rs @@ -354,6 +354,101 @@ impl BacktestHarness { Err(e) => eprintln!("stop_ctrl_counters read failed: {e}"), } + // CRT.diag: empirical measurement battery (Groups A-D). Drives the + // next round of CRT.1 threshold design — measure structural truth + // BEFORE another tuning pass. Single end-of-run dump, observe-only. + match self.sim.read_diagnostics() { + Ok(diag) => { + let n_b = self.cfg.n_parallel; + // Horizons match crates/ml-alpha/src/heads.rs::HORIZONS. + let horizons: [&str; 5] = ["30", "100", "300", "1000", "6000"]; + let n_horizons = horizons.len(); + + // Group A — per-horizon signal persistence. + for h in 0..n_horizons { + let flips: u64 = (0..n_b) + .map(|b| diag.flip_count[b * n_horizons + h] as u64) + .sum(); + let sum_run: u64 = (0..n_b) + .map(|b| diag.sum_run_length[b * n_horizons + h]) + .sum(); + let mean_run = if flips > 0 { + sum_run as f64 / flips as f64 + } else { + f64::NAN + }; + eprintln!( + "crt_diag h{}: flips={} mean_run_len={:.1} events", + horizons[h], flips, mean_run, + ); + let bucket_labels = ["1-9", "10-99", "100-999", "1k-9.9k", "10k+"]; + let mut hist_str = String::new(); + for bk in 0..5 { + let count: u64 = (0..n_b) + .map(|b| { + diag.run_length_hist[(b * n_horizons + h) * 5 + bk] as u64 + }) + .sum(); + hist_str.push_str(&format!(" {}:{}", bucket_labels[bk], count)); + } + eprintln!("crt_diag h{} run_length_hist:{}", horizons[h], hist_str); + } + + // Group B — smoothed-conviction histogram. + let mut conv_str = String::new(); + for bk in 0..10 { + let count: u64 = (0..n_b) + .map(|b| diag.conv_hist[b * 10 + bk] as u64) + .sum(); + let lo = bk as f32 / 10.0; + let hi = (bk + 1) as f32 / 10.0; + conv_str.push_str(&format!(" [{:.1}-{:.1}]:{}", lo, hi, count)); + } + eprintln!("crt_diag conv_ema_hist:{}", conv_str); + + // Group C — hold-time histogram. + let hold_labels = ["<1s", "1-10s", "10-60s", "1-10m", "10-60m", ">1h"]; + let mut hold_str = String::new(); + for bk in 0..6 { + let count: u64 = (0..n_b) + .map(|b| diag.hold_hist[b * 6 + bk] as u64) + .sum(); + hold_str.push_str(&format!(" {}:{}", hold_labels[bk], count)); + } + eprintln!("crt_diag hold_time_hist:{}", hold_str); + + // Group D — outcome by entry-conviction. + for bk in 0..10 { + let n_total: u64 = (0..n_b) + .map(|b| diag.outcome_n[b * 10 + bk] as u64) + .sum(); + let n_wins: u64 = (0..n_b) + .map(|b| diag.outcome_n_wins[b * 10 + bk] as u64) + .sum(); + let sum_pnl: f64 = (0..n_b) + .map(|b| diag.outcome_sum_pnl[b * 10 + bk] as f64) + .sum(); + let mean_pnl = if n_total > 0 { + sum_pnl / n_total as f64 + } else { + f64::NAN + }; + let win_rate = if n_total > 0 { + n_wins as f64 / n_total as f64 + } else { + f64::NAN + }; + let lo = bk as f32 / 10.0; + let hi = (bk + 1) as f32 / 10.0; + eprintln!( + "crt_diag outcome_by_entry_conv[{:.1}-{:.1}]: n={} win_rate={:.2}% mean_pnl_pu={:.3}", + lo, hi, n_total, win_rate * 100.0, mean_pnl, + ); + } + } + Err(e) => eprintln!("crt_diag read failed: {e}"), + } + Ok(RunStats { events_processed: self.event_count, decisions_taken: total_decisions, diff --git a/crates/ml-backtesting/src/sim/mod.rs b/crates/ml-backtesting/src/sim/mod.rs index 37f3effda..babf181c6 100644 --- a/crates/ml-backtesting/src/sim/mod.rs +++ b/crates/ml-backtesting/src/sim/mod.rs @@ -94,6 +94,38 @@ pub struct NanCountersV2 { pub last_bad_path: Vec, } +/// CRT.diag: end-of-run empirical measurement battery dump. +/// +/// All vectors are flat per-backtest layouts; the harness aggregates +/// across backtests for the log summary. See `LobSimCuda::read_diagnostics` +/// for the read-back path (single batched DtoV at smoke termination). +pub struct CrtDiagnostics { + /// Per-horizon flip count. Layout: `[backtest * N_HORIZONS + h]`. + pub flip_count: Vec, + /// Per-horizon cumulative run length over completed runs (u64 to avoid + /// overflow at multi-million-event smoke). Layout: `[backtest * N_HORIZONS + h]`. + pub sum_run_length: Vec, + /// Per-horizon 5-bucket log-scale run-length histogram. + /// Layout: `[(backtest * N_HORIZONS + h) * 5 + bucket]`. Buckets: + /// 0=1-9, 1=10-99, 2=100-999, 3=1k-9.9k, 4=10k+. + pub run_length_hist: Vec, + /// 10-bucket histogram of smoothed conviction at decision time. + /// Layout: `[backtest * 10 + bucket]`. Bucket k = [k/10, (k+1)/10). + pub conv_hist: Vec, + /// 6-bucket log-scale hold-time histogram at trade close. + /// Layout: `[backtest * 6 + bucket]`. Buckets: + /// 0=<1s, 1=1-10s, 2=10-60s, 3=1-10m, 4=10-60m, 5=>1h. + pub hold_hist: Vec, + /// Trade count per entry-conviction bucket. Layout: `[backtest * 10 + bucket]`. + pub outcome_n: Vec, + /// Cumulative realized PnL per entry-conviction bucket (price-units × lots, + /// pre USD-fp conversion). Layout: `[backtest * 10 + bucket]`. + pub outcome_sum_pnl: Vec, + /// Win count per entry-conviction bucket (segment_realized > 0). + /// Layout: `[backtest * 10 + bucket]`. + pub outcome_n_wins: Vec, +} + /// S2.1/S2.2: max_hold enforcement diagnostic counters. /// /// After S2.2, max_hold enforcement lives in `resting_orders_step` (event-rate). @@ -261,6 +293,26 @@ pub struct LobSimCuda { // block). Read by the kernel on every event; zero device-to-host traffic // on the per-event hot path. pub(crate) delta_floor_d: CudaSlice, // [n_backtests] — no-trade band in lots + + // CRT.diag: empirical measurement battery (observe-only, end-of-run dump). + // All counters allocated via alloc_zeros and incremented kernel-side; no + // host roundtrip per event. `read_diagnostics()` materialises them at + // smoke termination via a single batch of memcpy_dtov calls. + // + // Group A — per-horizon signal persistence: + pub(crate) diag_flip_count_d: CudaSlice, // [n_backtests * N_HORIZONS] + pub(crate) diag_current_run_length_d: CudaSlice, // [n_backtests * N_HORIZONS] + pub(crate) diag_sum_run_length_d: CudaSlice, // [n_backtests * N_HORIZONS] + pub(crate) diag_run_length_hist_d: CudaSlice, // [n_backtests * N_HORIZONS * 5] + pub(crate) diag_prev_dir_signed_d: CudaSlice, // [n_backtests * N_HORIZONS] + // Group B — conviction-EMA histogram: + pub(crate) diag_conv_hist_d: CudaSlice, // [n_backtests * 10] + // Group C — hold-time histogram at trade close: + pub(crate) diag_hold_hist_d: CudaSlice, // [n_backtests * 6] + // Group D — outcome by entry-conviction: + 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] } impl LobSimCuda { @@ -525,6 +577,40 @@ impl LobSimCuda { .alloc_zeros::(n_backtests) .context("alloc delta_floor_d")?; + // CRT.diag: empirical measurement battery (observe-only). Per-backtest + // memory footprint ≈ 4*(4+4+8+5*4)+10*4+6*4+10*(4+4+4) = 144+40+24+120 + // = 328 bytes. Negligible vs other state. All counters zero-init. + let diag_flip_count_d = stream + .alloc_zeros::(n_backtests * N_HORIZONS) + .context("alloc diag_flip_count_d")?; + let diag_current_run_length_d = stream + .alloc_zeros::(n_backtests * N_HORIZONS) + .context("alloc diag_current_run_length_d")?; + let diag_sum_run_length_d = stream + .alloc_zeros::(n_backtests * N_HORIZONS) + .context("alloc diag_sum_run_length_d")?; + let diag_run_length_hist_d = stream + .alloc_zeros::(n_backtests * N_HORIZONS * 5) + .context("alloc diag_run_length_hist_d")?; + let diag_prev_dir_signed_d = stream + .alloc_zeros::(n_backtests * N_HORIZONS) + .context("alloc diag_prev_dir_signed_d")?; + let diag_conv_hist_d = stream + .alloc_zeros::(n_backtests * 10) + .context("alloc diag_conv_hist_d")?; + let diag_hold_hist_d = stream + .alloc_zeros::(n_backtests * 6) + .context("alloc diag_hold_hist_d")?; + let diag_outcome_n_d = stream + .alloc_zeros::(n_backtests * 10) + .context("alloc diag_outcome_n_d")?; + let diag_outcome_sum_pnl_d = stream + .alloc_zeros::(n_backtests * 10) + .context("alloc diag_outcome_sum_pnl_d")?; + let diag_outcome_n_wins_d = stream + .alloc_zeros::(n_backtests * 10) + .context("alloc diag_outcome_n_wins_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 @@ -619,6 +705,16 @@ impl LobSimCuda { conviction_diff_var_ema_d, conviction_sample_var_ema_d, delta_floor_d, + diag_flip_count_d, + diag_current_run_length_d, + diag_sum_run_length_d, + diag_run_length_hist_d, + diag_prev_dir_signed_d, + diag_conv_hist_d, + diag_hold_hist_d, + diag_outcome_n_d, + diag_outcome_sum_pnl_d, + diag_outcome_n_wins_d, }) } @@ -847,6 +943,52 @@ impl LobSimCuda { }) } + /// CRT.diag: end-of-run materialisation of the empirical measurement + /// battery. Reads all 8 host-visible diagnostic slots in one batch (the + /// 9th — `diag_current_run_length_d` — and 10th — `diag_prev_dir_signed_d` + /// — are transient kernel-internal state, not part of the summary). + /// + /// Called ONCE at smoke termination from `BacktestHarness::run`. Each + /// `memcpy_dtoh` is a single device-to-host hop; per-event hot path + /// 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]; + 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()) + .context("diag_sum_run_length DtoH")?; + self.stream.memcpy_dtoh(&self.diag_run_length_hist_d, run_length_hist.as_mut_slice()) + .context("diag_run_length_hist DtoH")?; + self.stream.memcpy_dtoh(&self.diag_conv_hist_d, conv_hist.as_mut_slice()) + .context("diag_conv_hist DtoH")?; + self.stream.memcpy_dtoh(&self.diag_hold_hist_d, hold_hist.as_mut_slice()) + .context("diag_hold_hist DtoH")?; + self.stream.memcpy_dtoh(&self.diag_outcome_n_d, outcome_n.as_mut_slice()) + .context("diag_outcome_n DtoH")?; + self.stream.memcpy_dtoh(&self.diag_outcome_sum_pnl_d, outcome_sum_pnl.as_mut_slice()) + .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")?; + Ok(CrtDiagnostics { + flip_count, + sum_run_length, + run_length_hist, + conv_hist, + hold_hist, + outcome_n, + outcome_sum_pnl, + outcome_n_wins, + }) + } + /// Upload per-backtest price-range bounds. Call ONCE before the first /// apply_snapshot. min=0.0 / max=f32::INFINITY disables the check for /// that backtest (default after LobSimCuda::new). @@ -992,6 +1134,11 @@ impl LobSimCuda { // CRT.1 C1.4: open branch snapshots conviction_at_entry + // initializes conviction_ema mirror in open_trade_state. .arg(&self.conviction_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) + .arg(&mut self.diag_outcome_sum_pnl_d) + .arg(&mut self.diag_outcome_n_wins_d) .launch(cfg)?; } self.stream.synchronize()?; @@ -1238,6 +1385,13 @@ impl LobSimCuda { .arg(&mut self.conviction_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. + .arg(&mut self.diag_flip_count_d) + .arg(&mut self.diag_current_run_length_d) + .arg(&mut self.diag_sum_run_length_d) + .arg(&mut self.diag_run_length_hist_d) + .arg(&mut self.diag_prev_dir_signed_d) + .arg(&mut self.diag_conv_hist_d) .arg(&n) .launch(cfg)?; } @@ -1271,6 +1425,13 @@ impl LobSimCuda { .arg(&mut self.conviction_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. + .arg(&mut self.diag_flip_count_d) + .arg(&mut self.diag_current_run_length_d) + .arg(&mut self.diag_sum_run_length_d) + .arg(&mut self.diag_run_length_hist_d) + .arg(&mut self.diag_prev_dir_signed_d) + .arg(&mut self.diag_conv_hist_d) .arg(&n) .launch(cfg)?; } @@ -1328,6 +1489,11 @@ impl LobSimCuda { // CRT.1 C1.4: open branch snapshots conviction_at_entry + // initializes conviction_ema mirror in open_trade_state. .arg(&self.conviction_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) + .arg(&mut self.diag_outcome_sum_pnl_d) + .arg(&mut self.diag_outcome_n_wins_d) .launch(cfg)?; } self.stream.synchronize()?;