From 1e656948b3054911d6bf336f9d75101659cfee31 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Wed, 20 May 2026 21:15:40 +0200 Subject: [PATCH] arch(crt-1): composite exit_signal safety circuit-breaker (C1.4) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Per spec §4.3 + plan C1.4. Safety backup for the primary exit path (target → 0 from collapsed multi-horizon conviction, C1.2). Catches the edge case where conviction stays bounded above zero but the trade is deep in unrealized loss, or short/long horizons disagree for many events while threshold-gate logic still permits sizing. Three terms, max-of-three composite: degradation_term = (1 − conv_ema_now / conviction_at_entry) / 2.0 disagreement_term = disagree_consecutive_events / 32.0 pnl_decay_term = max(−pnl_adjusted_conv_ema, 0) / 1.0 exit_signal = max(deg, dis, pnl_decay) Fires force-flat (market_targets = (3, 0)) when exit_signal >= 1.0. State-update flow: - pnl_track open branch: write conviction_at_entry (offset 20), conviction_ema (offset 24, initialised to entry value), pnl_adjusted_conviction_ema (offset 44, = 0), peak_unrealized_pnl (offset 48, = 0). Takes new conviction_ema_per_b read-only arg to snapshot at open. - decision kernel intra-event (when pos open): update offsets 24/44/48/56 (disagreement counter) via update_open_trade_trajectory. - composite_exit_check device helper: reads offsets 20/24/44/56, writes (3, 0) on fire. Sibling of stop_check_isv with same return semantics; called after trajectory update so all four reads see the freshly-written values. Fields not read in CRT.1 (offsets 28-43 per-horizon EMA, 52 degradation counter, 60 horizon_idx_dominant) are NOT written by the open branch — alloc_zeros covers init. Per feedback_no_hiding: don't populate unused fields. Also fixes a latent C1.1 bug: pnl_track close branch read horizon_idx from offset 20 (now conviction_at_entry, f32); moved the read to offset 60 per the documented layout. Threshold-bootstrap defaults (will be ISV-derived in CRT.2 alongside the rest of the controller anchors per pearl_controller_anchors_isv_driven): degradation_factor = 2.0 (50% conviction decay vs entry → ≤0.5 term) disagreement_factor = 32.0 (32 events of short/long disagreement) pnl_decay_factor = 1.0 (pnl_adj_conv structurally ∈ [-1, +1]) Under these defaults disagreement_term is the only term that can fire on its own — degradation_term caps at 0.5 (conv_ema/conviction_at_entry ≥ 0) and pnl_decay_term caps at 1.0 in the limit. The three terms still add evidence in superposition; CRT.2 will retune factors against ISV percentiles. Unit-test coverage validates the disagreement path (composite_exit_signal_fires_on_disagreement). Per pearl_no_host_branches_in_captured_graph — all composite logic lives inside the decision kernel; no host roundtrip in the per-event hot path. Per feedback_no_htod_htoh_only_mapped_pinned — no new memcpy or stream.synchronize() in step_pnl_track or dispatch_latent_market_orders. Co-Authored-By: Claude Opus 4.7 (1M context) --- crates/ml-backtesting/cuda/decision_policy.cu | 161 ++++++++++++++++++ crates/ml-backtesting/cuda/pnl_track.cu | 26 ++- crates/ml-backtesting/src/sim/mod.rs | 10 ++ .../ml-backtesting/tests/stop_controller.rs | 90 ++++++++++ 4 files changed, 283 insertions(+), 4 deletions(-) diff --git a/crates/ml-backtesting/cuda/decision_policy.cu b/crates/ml-backtesting/cuda/decision_policy.cu index 44223eb7e..f80368228 100644 --- a/crates/ml-backtesting/cuda/decision_policy.cu +++ b/crates/ml-backtesting/cuda/decision_policy.cu @@ -176,6 +176,137 @@ __device__ static int stop_check_isv( // // Single-thread-per-backtest invariant: caller must already have // gated on threadIdx.x == 0. +// CRT.1 C1.4: composite exit_signal safety circuit-breaker per spec §4.3. +// Backup exit path for cases the primary `target → 0 from collapsed +// conviction` mechanism (C1.2) misses — most importantly when conviction +// stays bounded above zero but the trade is deep in unrealized loss. +// +// Three risk dimensions, max-of-three composite: +// degradation_term = (1 − conv_ema_now / conviction_at_entry) / 2.0 +// (50% conviction decay relative to entry fires) +// disagreement_term = disagree_consecutive_events / 32.0 +// (32 events of short/long horizon disagreement) +// pnl_decay_term = max(−pnl_adjusted_conv_ema, 0) / 1.0 +// (pnl-adjusted conviction flips and grows < −1) +// +// Reads the trajectory fields populated by the decision kernel above. +// Writes force-flat (3, 0) on fire. Returns 1 if fired, 0 otherwise. +// +// Thresholds (2.0 / 32.0 / 1.0) are spec-bootstrap defaults; they migrate +// to ISV-driven values in CRT.2 alongside the rest of the controller anchors +// (pearl_controller_anchors_isv_driven). Single-thread-per-backtest invariant. +__device__ static int composite_exit_check( + int b, + const Pos* __restrict__ positions, + const unsigned char* __restrict__ open_trade_state, + int* __restrict__ market_targets +) { + if (positions[b].position_lots == 0) return 0; + + const unsigned char* st_ro = open_trade_state + (size_t)b * 64; + const float conviction_at_entry = *reinterpret_cast(st_ro + 20); + const float conv_ema_now = *reinterpret_cast(st_ro + 24); + const float pnl_adjusted_conv_ema = *reinterpret_cast(st_ro + 44); + const unsigned int disagree_n = *reinterpret_cast(st_ro + 56); + + const float degradation_factor = 2.0f; + const float disagreement_factor = 32.0f; + const float pnl_decay_factor = 1.0f; + + const float degradation_term = (conviction_at_entry > 1e-6f) + ? (1.0f - conv_ema_now / conviction_at_entry) / degradation_factor + : 0.0f; + const float disagreement_term = (float)disagree_n / disagreement_factor; + const float pnl_decay_term = fmaxf(-pnl_adjusted_conv_ema, 0.0f) / pnl_decay_factor; + + float exit_signal = degradation_term; + if (disagreement_term > exit_signal) exit_signal = disagreement_term; + if (pnl_decay_term > exit_signal) exit_signal = pnl_decay_term; + + if (exit_signal >= 1.0f) { + market_targets[b * 2 + 0] = 3; + market_targets[b * 2 + 1] = 0; + return 1; + } + return 0; +} + +// CRT.1 C1.4: update intra-trade trajectory fields in open_trade_state. +// 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 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) +// +// Skipped when no position is open — the composite check is a guard +// against losing-trade edge cases and trajectory state is meaningless +// pre-entry. Single-thread-per-backtest invariant: caller must already +// have gated on threadIdx.x == 0. +__device__ static void update_open_trade_trajectory( + int b, + const Pos* __restrict__ positions, + const Book* __restrict__ books, + const float* __restrict__ alpha_probs, + float conv_ema_now, + unsigned char* __restrict__ open_trade_state +) { + const Pos& pos = positions[b]; + if (pos.position_lots == 0) return; + + unsigned char* st = open_trade_state + (size_t)b * 64; + + // Mirror the current smoothed conviction into open_trade_state so + // composite_exit_check reads it without needing the conviction_ema_per_b + // pointer plumbed in. + *reinterpret_cast(st + 24) = conv_ema_now; + + // Unrealized P&L = (close_px − vwap_entry) × dir × |lots|. + // Long: close at best bid; short: close at best ask. Matches the + // close_px used by stop_check_isv's SL/trail computation. + const Book& bk = books[b]; + const bool is_long = pos.position_lots > 0; + const float close_px = is_long ? bk.bid_px[0] : bk.ask_px[0]; + const float lots_abs = is_long ? (float)pos.position_lots + : -(float)pos.position_lots; + 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 + // 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). + const float pnl_sign = (unrealized >= 0.0f) ? 1.0f : -1.0f; + const float prev_pnl_adj = *reinterpret_cast(st + 44); + const float new_pnl_adj_inst = pnl_sign * conv_ema_now; + const float new_pnl_adj = (prev_pnl_adj == 0.0f) + ? new_pnl_adj_inst + : (0.9f * prev_pnl_adj + 0.1f * new_pnl_adj_inst); + *reinterpret_cast(st + 44) = new_pnl_adj; + + // peak_unrealized_pnl tracker. + const float prev_peak = *reinterpret_cast(st + 48); + if (unrealized > prev_peak) { + *reinterpret_cast(st + 48) = unrealized; + } + + // disagreement_consecutive_events — short-horizon (h=0) vs long-horizon + // (h=N_HORIZONS-1) directional sign disagreement. Reset to 0 when they + // align, increment when they oppose. Used by composite_exit_check's + // disagreement_term. + const float p_short = alpha_probs[0]; + const float p_long = alpha_probs[N_HORIZONS - 1]; + const bool dir_short_up = (p_short > 0.5f); + const bool dir_long_up = (p_long > 0.5f); + const unsigned int prev_disagree = *reinterpret_cast(st + 56); + if (dir_short_up != dir_long_up) { + *reinterpret_cast(st + 56) = prev_disagree + 1u; + } else { + *reinterpret_cast(st + 56) = 0u; + } +} + __device__ static float update_conviction_ema( int b, float conviction_abs, @@ -289,6 +420,11 @@ extern "C" __global__ void decision_policy_default( 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 + // 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 + // exit check has fresh inputs on the next pass. + unsigned char* __restrict__ open_trade_state_per_b, // [n_backtests * 64] int n_backtests ) { int b = blockIdx.x; @@ -329,6 +465,19 @@ extern "C" __global__ void decision_policy_default( conviction_diff_var_ema_per_b, conviction_sample_var_ema_per_b); + // 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, + conv_ema, 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 + // the trade is deep in unrealized loss — the primary `target → 0 from + // collapsed conviction` path (C1.2) cannot reach. + if (composite_exit_check(b, positions, open_trade_state_per_b, market_targets)) { + 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. @@ -444,6 +593,9 @@ extern "C" __global__ void decision_policy_program( 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 + // 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] int n_backtests ) { int b = blockIdx.x; @@ -491,6 +643,15 @@ extern "C" __global__ void decision_policy_program( conviction_diff_var_ema_per_b, conviction_sample_var_ema_per_b); + // 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); + + // 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]) { market_targets[b * 2 + 0] = 2; diff --git a/crates/ml-backtesting/cuda/pnl_track.cu b/crates/ml-backtesting/cuda/pnl_track.cu index 28ea93a8e..b36b753d1 100644 --- a/crates/ml-backtesting/cuda/pnl_track.cu +++ b/crates/ml-backtesting/cuda/pnl_track.cu @@ -53,7 +53,7 @@ extern "C" __global__ void pnl_track_step( unsigned char* pos_base, // [n_backtests * 24 bytes] - unsigned char* open_trade_state, // [n_backtests * 24 bytes] + unsigned char* open_trade_state, // [n_backtests * 64 bytes] unsigned char* trade_log_base, // [n_backtests * cap * 40 bytes] unsigned int* trade_log_head, // [n_backtests] unsigned long long current_ts_ns, @@ -63,7 +63,12 @@ extern "C" __global__ void pnl_track_step( // S1.20: NaN instrumentation counters (per-backtest). 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 + 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] ) { int b = blockIdx.x; if (b >= n_backtests || threadIdx.x != 0) return; @@ -86,7 +91,18 @@ extern "C" __global__ void pnl_track_step( *reinterpret_cast(st + 8) = (int)(vwap_at_open * 100.0f + 0.5f); *reinterpret_cast(st + 12) = now_size; *reinterpret_cast(st + 16) = pos->realized_pnl; - st[20] = 0; // horizon_idx; set by decision kernel in C7 + // CRT.1 C1.4: snapshot conviction_at_entry + seed the intra-trade + // EMA mirrors that the composite exit_signal reads (spec §4.3). + // Only the four offsets actually consumed by stop_check_isv's + // composite check are written here — per feedback_no_hiding, + // don't populate fields nothing reads. Other slots (per-horizon + // EMA at 28–43, degradation counter at 52, horizon_idx at 60) + // remain zero from the close-branch reset / alloc_zeros init. + const float conv_now = conviction_ema_per_b[b]; + *reinterpret_cast(st + 20) = conv_now; // conviction_at_entry + *reinterpret_cast(st + 24) = conv_now; // conviction_ema (init = entry) + *reinterpret_cast(st + 44) = 0.0f; // pnl_adjusted_conviction_ema + *reinterpret_cast(st + 48) = 0.0f; // peak_unrealized_pnl } else if (prev_size != 0 && now_size == 0) { // Close — emit TradeRecord. const unsigned int idx = trade_log_head[b] % (unsigned int)trade_log_cap; @@ -97,7 +113,9 @@ extern "C" __global__ void pnl_track_step( const int entry_px_x100 = *reinterpret_cast(st + 8); const int entry_size = *reinterpret_cast(st + 12); const float realized_at_open = *reinterpret_cast(st + 16); - const unsigned char horizon_idx = st[20]; + // CRT.1 C1.4: horizon_idx_dominant_at_entry moved from offset 20 + // (now conviction_at_entry, f32) to offset 60 in the 64-byte layout. + const unsigned char horizon_idx = st[60]; // Realized P&L delta = pos.realized_pnl_now − pos.realized_pnl_at_open. // In price-units × lots. diff --git a/crates/ml-backtesting/src/sim/mod.rs b/crates/ml-backtesting/src/sim/mod.rs index 5ee76a198..37f3effda 100644 --- a/crates/ml-backtesting/src/sim/mod.rs +++ b/crates/ml-backtesting/src/sim/mod.rs @@ -989,6 +989,9 @@ impl LobSimCuda { .arg(&mut self.zero_vwap_at_open_d) .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) .launch(cfg)?; } self.stream.synchronize()?; @@ -1233,6 +1236,8 @@ impl LobSimCuda { .arg(&mut self.conviction_ema_d) .arg(&mut self.conviction_diff_var_ema_d) .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) .arg(&n) .launch(cfg)?; } @@ -1264,6 +1269,8 @@ impl LobSimCuda { .arg(&mut self.conviction_ema_d) .arg(&mut self.conviction_diff_var_ema_d) .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) .arg(&n) .launch(cfg)?; } @@ -1318,6 +1325,9 @@ impl LobSimCuda { .arg(&mut self.zero_vwap_at_open_d) .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) .launch(cfg)?; } self.stream.synchronize()?; diff --git a/crates/ml-backtesting/tests/stop_controller.rs b/crates/ml-backtesting/tests/stop_controller.rs index e6cdac87c..7fed7cd25 100644 --- a/crates/ml-backtesting/tests/stop_controller.rs +++ b/crates/ml-backtesting/tests/stop_controller.rs @@ -1331,3 +1331,93 @@ fn open_trade_state_64_byte_layout() { assert_eq!(ml_backtesting::lob::OPEN_TRADE_STATE_BYTES, 64, "spec §7: open_trade_state must be 64 bytes for CRT.1 trajectory tracking"); } + +/// CRT.1 C1.4: composite exit_signal — disagreement_term fires force-flat. +/// +/// With the spec-default `disagreement_factor = 32.0`, the composite +/// disagreement_term reaches 1.0 after 32 consecutive events of short +/// (h=0) vs long (h=N-1) horizon directional sign disagreement. Once +/// `exit_signal >= 1.0`, the composite check writes the (3, 0) force-flat +/// encoding. This is the only one of the three composite terms reachable +/// with bounded conv_ema ∈ [0, 1] under spec-default factors — both +/// degradation_term and pnl_decay_term cap at 0.5 / 1.0 by construction, +/// so the disagreement counter is the canonical edge-case trigger. +/// +/// Test geometry: +/// 1. Cold-start ISV + strong-bullish alpha → open a long position. The +/// pnl_track open-branch snapshots conviction_at_entry into +/// open_trade_state[20], seeds open_trade_state[24] = entry. +/// 2. Switch to a SPLIT alpha (h0 bullish, h4 bearish, middle bullish so +/// a position-holding majority keeps conviction_signed positive and +/// SL doesn't widen). Every event with h0/h4 sign disagreement +/// increments the disagreement counter at open_trade_state[56]. +/// 3. After 32 such events, composite_exit_check sees +/// disagreement_term = 32/32 = 1.0 → writes force-flat (side=3). +/// +/// The SL/trail check runs first; it must not fire over the test horizon, +/// so pnl_ema_loss is set large (1000.0) and the book stays at the entry +/// price level (no SL pressure). +#[test] +#[ignore = "requires CUDA"] +fn composite_exit_signal_fires_on_disagreement() -> 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])?; + + // Cold-start ISV with large pnl_ema_loss so SL/trail can't fire during + // the 32-event disagreement window. + let cold_start: [IsvKellyStateHost; N_HORIZONS] = std::array::from_fn(|_| IsvKellyStateHost { + pnl_ema_win: 0.0, + pnl_ema_loss: 1000.0, + win_rate_ema: 0.0, + n_trades_seen: 0, + realised_return_var: 0.0, + recent_sharpe: 0.0, + }); + sim.write_isv_kelly(0, &cold_start)?; + + let (bp, bs, ap, az) = level_book(5500.0, 0.25); + sim.apply_snapshot(&bp, &bs, &ap, &az)?; + let (bp2, bs2, ap2, az2) = level_book(5500.1, 0.25); + sim.apply_snapshot(&bp2, &bs2, &ap2, &az2)?; + + // Open long with strong unanimous bullish alpha. + let bullish: [f32; N_HORIZONS] = [0.95; N_HORIZONS]; + sim.broadcast_alpha(&bullish)?; + let mut ts: u64 = 1_000_000; + sim.step_decision_with_latency(ts, &cfg_default(1))?; + for _ in 0..3 { + sim.step_resting_orders(ts, 0.0)?; + sim.step_pnl_track(ts)?; + ts += 1_000_000; + } + let pos_open = sim.read_pos(0)?; + assert!(pos_open.position_lots > 0, + "setup: long position must open before disagreement test; got {}", + pos_open.position_lots); + + // Switch to split alpha: h=0 bullish (0.6), h=N-1 bearish (0.4), middle + // horizons bullish so the signed conviction stays positive and the + // threshold gate doesn't shut down sizing each event. The trajectory + // update bumps disagreement_consecutive_events because the h=0 vs + // h=N-1 directional sign disagrees. + let mut split: [f32; N_HORIZONS] = [0.6; N_HORIZONS]; + split[N_HORIZONS - 1] = 0.4; + + // 32 events at spec-default factor produces exactly 1.0 — to be safe + // against any single-event slack, run 33 events. + for _ in 0..33 { + sim.broadcast_alpha(&split)?; + sim.step_decision_with_latency(ts, &cfg_default(1))?; + ts += 1_000_000; + } + + let (side, size) = sim.read_market_target(0)?; + assert_eq!(side, 3, + "composite exit_signal must fire force-flat after 32+ disagreement events; got side={side}"); + assert_eq!(size, 0, "force-flat encoding must have abs_sz=0; got size={size}"); + Ok(()) +}