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

arch(crt-1): composite exit_signal safety circuit-breaker (C1.4)

Browse Source 
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-20 21:15:40 +02:00
parent 84793ea46e
commit 1e656948b3
4 changed files with 283 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -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<const float*>(st_ro + 20);
const float conv_ema_now = *reinterpret_cast<const float*>(st_ro + 24);
const float pnl_adjusted_conv_ema = *reinterpret_cast<const float*>(st_ro + 44);
const unsigned int disagree_n = *reinterpret_cast<const unsigned int*>(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<float*>(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<const float*>(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<float*>(st + 44) = new_pnl_adj;
// peak_unrealized_pnl tracker.
const float prev_peak = *reinterpret_cast<const float*>(st + 48);
if (unrealized > prev_peak) {
*reinterpret_cast<float*>(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<const unsigned int*>(st + 56);
if (dir_short_up != dir_long_up) {
*reinterpret_cast<unsigned int*>(st + 56) = prev_disagree + 1u;
} else {
*reinterpret_cast<unsigned int*>(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;

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

@@ -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<int*>(st + 8) = (int)(vwap_at_open * 100.0f + 0.5f);
*reinterpret_cast<int*>(st + 12) = now_size;
*reinterpret_cast<float*>(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<float*>(st + 20) = conv_now; // conviction_at_entry
*reinterpret_cast<float*>(st + 24) = conv_now; // conviction_ema (init = entry)
*reinterpret_cast<float*>(st + 44) = 0.0f; // pnl_adjusted_conviction_ema
*reinterpret_cast<float*>(st + 48) = 0.0f; // peak_unrealized_pnl
} 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<int*>(st + 8);
const int entry_size = *reinterpret_cast<int*>(st + 12);
const float realized_at_open = *reinterpret_cast<float*>(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.

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

@@ -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()?;

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

@@ -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(())
}