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

arch(crt-a): Wiener-α conviction-EMA smoothing in decision_policy

Browse Source 
After A1 (commit 045850e8f) the controller fires every event. Without
smoothing, target lots would oscillate as alpha probabilities jitter
event-to-event, generating hyperactive spread-bleed from back-and-forth
trades.

Ships per spec §4.2 and §3.7: Wiener-α adaptive EMA on max-conviction-
across-horizons. Formula:
  α_raw    = diff_var / (diff_var + sample_var + ε)
  α_active = max(α_raw, 0.4)        per pearl_wiener_alpha_floor_for_nonstationary
  ema      = α_active × new + (1 − α_active) × prev
First-observation bootstrap: prev_ema == 0 → replace directly per
pearl_first_observation_bootstrap.

Three new device slots (per-backtest):
  - conviction_ema_d:             smoothed conviction (used for final-aggregate rescale)
  - conviction_diff_var_ema_d:    second-order EMA, drives adaptive α
  - conviction_sample_var_ema_d:  second-order EMA, drives adaptive α

Architectural choice: rescale at the final aggregate (final_size *= conv_ema /
raw_max_conv), not at per-horizon sig_mag. Per-horizon sizing keeps raw
|alpha-0.5|*2 so horizons with no signal (alpha=0.5) contribute zero —
the spec §4.2 "smoothed conviction" is a unified gain over the aggregate,
not a substitute for per-horizon signal strength. At bootstrap the scale
is exactly 1.0 (raw_max_conv == conv_ema) so the very first decision is
bit-identical to the pre-A2 kernel. Direction recovery (sign of
alpha-0.5) remains per-horizon — genuine reversals respond at event rate.

Threaded through both decision_policy_default and decision_policy_program
kernels' signatures and launches in sim/mod.rs.

Full multi-horizon conviction *aggregation* (spec §4.4) remains Phase B;
A2 ships ONLY the smoothing operator on the existing scalar conviction.

Pure on-device: no memcpy_htod / dtoh / dtov / synchronize in hot path.
Single test-only DtoH accessor (read_conviction_ema) for unit-test
inspection of the smoothed value.

Tests:
- conviction_ema_smooths_micro_oscillations — sentinel→bootstrap→bounded
  EMA across 10 alternating high/low all-bullish alpha drives; direction
  stable.
- conviction_ema_does_not_lag_reversals — sign flip in alpha → target
  side flips next event.
- All 22 stop_controller + 5 decision_floor_coldstart + 3 threshold_and_cost
  + parallel_sim + ring3_replay + trainer_parity + 6 fuzz tests pass.
- 2 lob_sim_fixtures tests (fix_decision_alpha_buy_close, fix_decision_program_h4_only)
  were already failing on baseline 045850e8f pre-A2; not a regression from
  this commit.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-20 19:01:09 +02:00
parent 045850e8f3
commit 1d889d2de9
3 changed files with 350 additions and 25 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -158,6 +158,71 @@ __device__ static int stop_check_isv(
return 0;
}
// CRT-A2: Wiener-α adaptive EMA on max-conviction-across-horizons.
// Reads alpha_probs[h] (broadcast across backtests), computes the current
// max |alpha-0.5|*2, updates per-backtest second-order EMAs of sample
// variance + difference variance, derives adaptive α with floor 0.4 per
// pearl_wiener_alpha_floor_for_nonstationary, and blends.
// First-observation bootstrap (prev_ema == 0): replace directly per
// pearl_first_observation_bootstrap.
//
// Writes new EMA to conviction_ema[b], updates the two second-order
// variance EMAs, and returns the smoothed max-conviction. Direction is
// recovered separately at each per-horizon sizing call (raw alpha sign);
// only the magnitude is smoothed so genuine reversals respond at event
// rate while micro-jitter is suppressed.
//
// Single-thread-per-backtest invariant: caller must already have
// gated on threadIdx.x == 0.
__device__ static float update_conviction_ema(
int b,
const float* __restrict__ alpha_probs,
float* __restrict__ conviction_ema,
float* __restrict__ conviction_diff_var_ema,
float* __restrict__ conviction_sample_var_ema
) {
float max_conv = 0.0f;
#pragma unroll
for (int h = 0; h < N_HORIZONS; ++h) {
const float c = fabsf(alpha_probs[h] - 0.5f) * 2.0f;
if (c > max_conv) max_conv = c;
}
// Defensive clamp to [0, 1] — alpha_probs should be in [0, 1] by
// contract but bound it before feeding into a variance EMA.
if (max_conv > 1.0f) max_conv = 1.0f;
if (max_conv < 0.0f) max_conv = 0.0f;
// Update second-order EMAs (diff_var, sample_var) with fixed α=0.1.
// These track the variance of the signal and the variance of changes,
// and are themselves bootstrapped with the first observation.
const float prev_ema = conviction_ema[b];
const float diff = max_conv - prev_ema;
const float diff_var_prev = conviction_diff_var_ema[b];
const float new_diff_var = (diff_var_prev == 0.0f)
? (diff * diff) // first-observation bootstrap
: (0.9f * diff_var_prev + 0.1f * diff * diff);
conviction_diff_var_ema[b] = new_diff_var;
const float sample_var_prev = conviction_sample_var_ema[b];
const float new_sample_var = (sample_var_prev == 0.0f)
? (max_conv * max_conv) // first-observation bootstrap
: (0.9f * sample_var_prev + 0.1f * max_conv * max_conv);
conviction_sample_var_ema[b] = new_sample_var;
// Wiener α with floor at 0.4 (pearl_wiener_alpha_floor_for_nonstationary).
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.4f) ? alpha_raw : 0.4f;
// First-observation bootstrap: prev_ema==0 → replace directly. Otherwise
// blend with Wiener-α.
const float new_ema = (prev_ema == 0.0f)
? max_conv
: (alpha_active * max_conv + (1.0f - alpha_active) * prev_ema);
conviction_ema[b] = new_ema;
return new_ema;
}
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)]
@@ -187,6 +252,10 @@ extern "C" __global__ void decision_policy_default(
const Book* __restrict__ books_per_b, // [n_backtests * BOOK_LEVELS] (via lob_state.cuh)
// S2.1: mh_kernel_calls diagnostic counter (non-zero-position entries).
unsigned int* __restrict__ mh_kernel_calls_per_b, // [n_backtests]
// CRT-A2: Wiener-α conviction-EMA state (per-backtest).
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
int n_backtests
) {
int b = blockIdx.x;
@@ -202,23 +271,33 @@ extern "C" __global__ void decision_policy_default(
return;
}
// Compute raw max conviction once (used by threshold gate AND by the
// CRT-A2 magnitude-scale below).
float raw_max_conv = 0.0f;
#pragma unroll
for (int h = 0; h < N_HORIZONS; ++h) {
raw_max_conv = fmaxf(raw_max_conv, fabsf(alpha_probs[h] - 0.5f) * 2.0f);
}
// CRT-A2: Wiener-α adaptive EMA on max-conviction. Runs ahead of the
// threshold gate so the EMA tracks all events (gating off an unsmoothed
// tick would let the EMA grow stale on long gated stretches).
const float conv_ema = update_conviction_ema(
b, alpha_probs,
conviction_ema_per_b,
conviction_diff_var_ema_per_b,
conviction_sample_var_ema_per_b);
// P4: threshold gate. Skip entirely if max conviction below threshold.
// Pre-Kelly placement keeps the gate deterministic from alpha alone,
// so the threshold pre-registration step (compute p60-p95 absolute
// values from observed convictions) reflects exactly what gets gated
// in deployment.
{
float max_conv = 0.0f;
#pragma unroll
for (int h = 0; h < N_HORIZONS; ++h) {
max_conv = fmaxf(max_conv, fabsf(alpha_probs[h] - 0.5f) * 2.0f);
}
if (max_conv < threshold_per_b[b]) {
market_targets[b * 2 + 0] = 2;
market_targets[b * 2 + 1] = 0;
open_horizon_masks[b] = 0u;
return;
}
if (raw_max_conv < threshold_per_b[b]) {
market_targets[b * 2 + 0] = 2;
market_targets[b * 2 + 1] = 0;
open_horizon_masks[b] = 0u;
return;
}
// Per-backtest scalar reads (host-uploaded arrays).
@@ -295,6 +374,26 @@ extern "C" __global__ void decision_policy_default(
}
}
// CRT-A2: rescale the aggregate magnitude by (conv_ema / raw_max_conv).
// The smoothed conviction acts as a unified scalar gain over the raw
// per-horizon contributions — per-horizon sizing logic (Kelly, cap_units,
// sharpe weights, attribution mask) is unchanged. Spec §4.2.
//
// The scale factor is in [0, 1] when raw_max_conv > 0 (since the EMA
// is a convex blend of historical raw_max_conv values, bounded by their
// running max). At the first observation the EMA == raw_max_conv exactly
// (bootstrap), so the scale is 1.0 and behaviour is bit-identical to the
// pre-A2 kernel for the very first decision.
//
// Guard: when raw_max_conv ≈ 0 the threshold gate above (when configured
// to non-zero) would already have returned. Under threshold=0.0 a raw_max_conv
// of 0 means every alpha is exactly 0.5 → final_size is also 0 (any dir
// contributions cancel in expectation under uniform Kelly/cap); the scale
// doesn't change that. We still divide-guard with eps to be safe.
if (raw_max_conv > 1e-9f) {
final_size *= (conv_ema / raw_max_conv);
}
// Round-to-nearest (truncation alone bites here: 0.8(f32) * 0.75 * 5.0
// = 2.99999976 → (int)2, off by one. Floor truncation for sub-1
// suppression is preserved via the explicit |lots| < 1 check.
@@ -375,6 +474,10 @@ extern "C" __global__ void decision_policy_program(
const Book* __restrict__ books_per_b, // [n_backtests * BOOK_LEVELS]
// S2.1: mh_kernel_calls diagnostic counter (non-zero-position entries).
unsigned int* __restrict__ mh_kernel_calls_per_b, // [n_backtests]
// CRT-A2: Wiener-α conviction-EMA state (per-backtest).
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
int n_backtests
) {
int b = blockIdx.x;
@@ -395,19 +498,30 @@ extern "C" __global__ void decision_policy_program(
return;
}
// Compute raw max conviction once (used by threshold gate AND by the
// CRT-A2 final-aggregate rescale in OP_WRITE_ORDER).
float raw_max_conv = 0.0f;
#pragma unroll
for (int h = 0; h < N_HORIZONS; ++h) {
raw_max_conv = fmaxf(raw_max_conv, fabsf(alpha_probs[h] - 0.5f) * 2.0f);
}
// CRT-A2: Wiener-α adaptive EMA on max-conviction. See decision_policy_default
// for rationale. The smoothed magnitude rescales final_size at the
// OP_WRITE_ORDER terminal — per-horizon attribution + Kelly remain
// bit-identical with raw |alpha-0.5|*2.
const float conv_ema = update_conviction_ema(
b, alpha_probs,
conviction_ema_per_b,
conviction_diff_var_ema_per_b,
conviction_sample_var_ema_per_b);
// P4: threshold gate.
{
float max_conv = 0.0f;
#pragma unroll
for (int h = 0; h < N_HORIZONS; ++h) {
max_conv = fmaxf(max_conv, fabsf(alpha_probs[h] - 0.5f) * 2.0f);
}
if (max_conv < threshold_per_b[b]) {
market_targets[b * 2 + 0] = 2;
market_targets[b * 2 + 1] = 0;
open_horizon_masks[b] = 0u;
return;
}
if (raw_max_conv < threshold_per_b[b]) {
market_targets[b * 2 + 0] = 2;
market_targets[b * 2 + 1] = 0;
open_horizon_masks[b] = 0u;
return;
}
// Per-backtest scalar reads (host-uploaded arrays).
@@ -564,8 +678,13 @@ extern "C" __global__ void decision_policy_program(
case OP_WRITE_ORDER: {
if (sp <= 0) break;
sp -= 1;
const float final_size = stack_val[sp];
float final_size = stack_val[sp];
const unsigned int attribution = stack_mask[sp];
// CRT-A2: rescale the aggregate magnitude by (conv_ema / raw_max_conv).
// See decision_policy_default for rationale.
if (raw_max_conv > 1e-9f) {
final_size *= (conv_ema / raw_max_conv);
}
int lots = (int)truncf(final_size + (final_size >= 0.0f ? 0.5f : -0.5f));
if (lots == 0) {
market_targets[b * 2 + 0] = 2; // no-op

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

@@ -244,6 +244,16 @@ pub struct LobSimCuda {
// S2.1/S2.2: max_hold diagnostic counters (2 remain after S2.2 prune).
mh_kernel_calls_d: CudaSlice<u32>, // [n_backtests]
mh_force_flat_seen_by_seed_d: CudaSlice<u32>, // [n_backtests]
// CRT-A2: Wiener-α adaptive EMA state on max-conviction. Smooths
// event-rate alpha jitter so sizing magnitude doesn't oscillate
// bar-by-bar. All three slots start at zero (sentinel = "no
// observation yet") so the first decision-kernel call replaces them
// directly per pearl_first_observation_bootstrap. Updated in-place
// by both decision_policy_default and decision_policy_program.
conviction_ema_d: CudaSlice<f32>, // [n_backtests] — smoothed conviction
conviction_diff_var_ema_d: CudaSlice<f32>, // [n_backtests] — 2nd-order EMA of (new−prev)²
conviction_sample_var_ema_d: CudaSlice<f32>, // [n_backtests] — 2nd-order EMA of x²
}
impl LobSimCuda {
@@ -488,6 +498,19 @@ impl LobSimCuda {
.alloc_zeros::<u32>(n_backtests)
.context("alloc mh_force_flat_seen_by_seed_d")?;
// CRT-A2: Wiener-α conviction-EMA state. Zero-init means
// "first observation pending" — the decision kernel branches on
// prev_ema == 0 to bootstrap directly (no warmup zero-bias).
let conviction_ema_d = stream
.alloc_zeros::<f32>(n_backtests)
.context("alloc conviction_ema_d")?;
let conviction_diff_var_ema_d = stream
.alloc_zeros::<f32>(n_backtests)
.context("alloc conviction_diff_var_ema_d")?;
let conviction_sample_var_ema_d = stream
.alloc_zeros::<f32>(n_backtests)
.context("alloc conviction_sample_var_ema_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
@@ -578,6 +601,9 @@ impl LobSimCuda {
last_bad_path_d,
mh_kernel_calls_d,
mh_force_flat_seen_by_seed_d,
conviction_ema_d,
conviction_diff_var_ema_d,
conviction_sample_var_ema_d,
})
}
@@ -713,6 +739,20 @@ impl LobSimCuda {
Ok(buf[backtest_idx])
}
/// CRT-A2: read `conviction_ema_d[backtest_idx]`. Test-only accessor —
/// production paths inspect smoothed conviction on-device only. One DtoH.
pub fn read_conviction_ema(&self, backtest_idx: usize) -> Result<f32> {
anyhow::ensure!(
backtest_idx < self.n_backtests,
"backtest_idx {} >= n_backtests {}",
backtest_idx,
self.n_backtests
);
let mut buf = vec![0.0f32; self.n_backtests];
self.stream.memcpy_dtoh(&self.conviction_ema_d, buf.as_mut_slice())?;
Ok(buf[backtest_idx])
}
/// Read `snapshots_skipped_d`. Returns per-backtest count of snapshots skipped
/// due to corrupt top-of-book (NaN/Inf price or zero/negative size at level 0).
/// Used by tests to verify the skip logic fires; production code can sum for
@@ -1170,6 +1210,10 @@ impl LobSimCuda {
.arg(&self.books_d)
// S2.1: mh_kernel_calls diagnostic counter.
.arg(&mut self.mh_kernel_calls_d)
// CRT-A2: Wiener-α conviction-EMA state.
.arg(&mut self.conviction_ema_d)
.arg(&mut self.conviction_diff_var_ema_d)
.arg(&mut self.conviction_sample_var_ema_d)
.arg(&n)
.launch(cfg)?;
}
@@ -1195,6 +1239,10 @@ impl LobSimCuda {
.arg(&self.books_d)
// S2.1: mh_kernel_calls diagnostic counter.
.arg(&mut self.mh_kernel_calls_d)
// CRT-A2: Wiener-α conviction-EMA state.
.arg(&mut self.conviction_ema_d)
.arg(&mut self.conviction_diff_var_ema_d)
.arg(&mut self.conviction_sample_var_ema_d)
.arg(&n)
.launch(cfg)?;
}

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

@@ -1234,3 +1234,161 @@ fn max_hold_counters_initialize_to_zero() -> Result<()> {
assert_eq!(c.mh_force_flat_seen_by_seed[0], 0, "mh_force_flat_seen_by_seed must initialize to 0");
Ok(())
}
/// CRT-A2: Wiener-α adaptive EMA on max-conviction-across-horizons.
///
/// After A1 the controller fires every event. Without smoothing, sizing
/// magnitude would oscillate with the per-event alpha jitter (spread-bleed).
/// A2 adds a per-backtest Wiener-α EMA with floor 0.4 (non-stationary
/// control loop per pearl_wiener_alpha_floor_for_nonstationary) on the
/// max |alpha-0.5|*2 across horizons. Direction continues to come from
/// the raw alpha sign — only magnitude is smoothed.
///
/// This test drives alternating high/low all-bullish alpha probs and
/// asserts:
/// (a) the on-device EMA is in fact populated (sentinel 0 → nonzero
/// after first observation);
/// (b) the EMA value sits strictly between the high and low input
/// convictions (smoothing produces a bounded mix, never extrapolates);
/// (c) under high enough kelly-floor to clear the integer-rounding bar,
/// direction (side) stays stable across alternation — no sign flips
/// on co-directional inputs.
fn cfg_high_kelly_floor(n: usize) -> BatchedSimConfig {
// kelly_frac_floor=1.0 keeps |signed_size| ≥ ~sig_mag * max_lots = 0.1*5 = 0.5
// for the low-conv alpha (sig_mag ~ EMA), rounding to ±1 lots — so the
// *direction* of the target is stable across alternation regardless of
// sizing oscillation.
BatchedSimConfig::from_uniform(n, &UniformSimParams {
target_annual_vol_units: 50.0,
annualisation_factor: 825.0,
max_lots: 5,
latency_ns: 0,
kelly_frac_floor: 1.0,
sharpe_weight_floor: 0.10,
threshold: 0.0,
cost_per_lot_per_side: 0.0,
max_hold_ns: 0,
min_reasonable_px: 0.0,
max_reasonable_px: f32::INFINITY,
})
}
#[test]
#[ignore = "requires CUDA"]
fn conviction_ema_smooths_micro_oscillations() -> 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)?;
let (bp, bs, ap, az) = level_book(5500.0, 0.25);
sim.apply_snapshot(&bp, &bs, &ap, &az)?;
let alphas_high: [f32; N_HORIZONS] = [0.8; N_HORIZONS]; // strong bullish, conv=0.6
let alphas_low: [f32; N_HORIZONS] = [0.55; N_HORIZONS]; // weak bullish, conv=0.1
let conv_high = 0.6_f32;
let conv_low = 0.1_f32;
let cfg = cfg_high_kelly_floor(1);
// Pre-condition: EMA slot starts at sentinel 0 (no observation yet).
assert_eq!(sim.read_conviction_ema(0)?, 0.0,
"conviction_ema must initialize to 0 (first-observation sentinel)");
// First decision: alphas_high. Bootstraps the EMA directly per
// pearl_first_observation_bootstrap (prev_ema == 0 → replace).
sim.broadcast_alpha(&alphas_high)?;
sim.step_decision_with_latency(1_000_000_000u64, &cfg)?;
let ema_after_first = sim.read_conviction_ema(0)?;
assert!((ema_after_first - conv_high).abs() < 1e-5,
"first observation must replace directly (bootstrap); got EMA={ema_after_first}, expected {conv_high}");
// Drive 10 alternations. Capture each EMA + target sign.
let mut prev_target_signed: Option<i32> = None;
let mut sign_flips = 0;
let mut ema_min = ema_after_first;
let mut ema_max = ema_after_first;
for i in 0..10 {
let probs = if i % 2 == 0 { alphas_low } else { alphas_high };
sim.broadcast_alpha(&probs)?;
let ts = 1_000_000_000u64 * ((i as u64) + 2);
sim.step_decision_with_latency(ts, &cfg)?;
let (side, size) = sim.read_market_target(0)?;
// side ∈ {0=buy, 1=sell, 2=noop, 3=force-flat}. For all-bullish
// alphas the raw direction is +; smoothed magnitude shouldn't flip
// it under the high-kelly-floor config.
let target_signed: i32 = match side {
0 => size,
1 => -size,
_ => 0,
};
if let Some(p) = prev_target_signed {
if target_signed != 0 && p != 0 && (target_signed > 0) != (p > 0) {
sign_flips += 1;
}
}
prev_target_signed = Some(target_signed);
let ema = sim.read_conviction_ema(0)?;
if ema < ema_min { ema_min = ema; }
if ema > ema_max { ema_max = ema; }
}
// (b): EMA stays bounded between the two input convictions (post-bootstrap,
// the blend formula α·new + (1-α)·old is a strict convex combination of
// observed values, so the running EMA is bounded by the observed range).
assert!(ema_min > conv_low - 1e-4 && ema_max < conv_high + 1e-4,
"EMA must stay between conv_low={conv_low} and conv_high={conv_high}; got min={ema_min}, max={ema_max}");
// The EMA should also have moved away from the bootstrap value (proves
// the second-and-onward update path executed, not just bootstrap).
assert!(ema_min < conv_high - 1e-4,
"EMA must drop below initial bootstrap conv_high={conv_high} after seeing alphas_low; got min={ema_min}");
// (c): direction stable under co-directional alternation. The raw alpha
// sign is positive in both cases (>0.5); smoothed magnitude shouldn't
// produce a sign flip.
assert_eq!(sign_flips, 0,
"direction must stay stable under EMA smoothing of co-directional alphas; flips={sign_flips}");
Ok(())
}
/// CRT-A2 corollary: when a true reversal arrives (alpha sign flips from
/// >0.5 to <0.5), the target direction must respond at the next event —
/// the smoothing is on magnitude only, not on sign. This is the property
/// that distinguishes Wiener-α smoothing on max-conviction from smoothing
/// the alpha probabilities themselves (which would lag reversals).
#[test]
#[ignore = "requires CUDA"]
fn conviction_ema_does_not_lag_reversals() -> 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)?;
let (bp, bs, ap, az) = level_book(5500.0, 0.25);
sim.apply_snapshot(&bp, &bs, &ap, &az)?;
let cfg = cfg_high_kelly_floor(1);
// Bootstrap with strong bullish.
sim.broadcast_alpha(&[0.8; N_HORIZONS])?;
sim.step_decision_with_latency(1_000_000_000, &cfg)?;
let (side_buy, size_buy) = sim.read_market_target(0)?;
assert_eq!(side_buy, 0, "bootstrap bullish → buy; got side={side_buy} size={size_buy}");
assert!(size_buy >= 1, "bootstrap bullish → size>=1; got {size_buy}");
// Hard reversal: strong bearish. Direction must flip on next event.
sim.broadcast_alpha(&[0.2; N_HORIZONS])?;
sim.step_decision_with_latency(2_000_000_000, &cfg)?;
let (side_sell, size_sell) = sim.read_market_target(0)?;
// The position may now be long (from the prior buy fill), but the
// *target* side written by the decision kernel reflects raw alpha sign.
// Stop checks only fire on open positions with SL/trail breach; with
// identical book + zero ATR they won't fire here, so the alpha-driven
// target dictates.
assert_eq!(side_sell, 1,
"post-reversal target must flip to sell; got side={side_sell} size={size_sell}");
Ok(())
}