arch(crt-1): multi-horizon ISV-weighted conviction replaces scalar EMA + rescale (C1.2)

Per spec §4.4 + plan C1.2. Replaces v2 A2 (1d889d2de) + A2.1
(fe2498769) scalar approach that failed Gate 1 catastrophically on
smokes vjmwc and lkrdf (155k trades, 9100% drawdown, Sharpe -15.7).

The structural fix: direction comes from per-horizon SIGNED sum, so
disagreeing horizons cancel. The scalar EMA on max(|p-0.5|) could only
smooth magnitude — it could not smooth direction jitter. Only this
multi-horizon weighted formula can.

Formula (spec §4.4):
  weight_h           = max(isv[h].net_edge, ε) / (isv[h].var + cost²)
  weighted_h         = magnitude_h × weight_h × direction_h
  conviction_signed  = sum(weighted_h) / sum(|weight_h|)
  conviction_ema     = Wiener-α EMA(|conviction_signed|)   (reuses A2 slots)
  target_lots        = round(sign(conviction_signed) × conviction_ema × max_lots)

Pearl conformance:
  - pearl_controller_anchors_isv_driven: weights from ISV state
  - pearl_one_unbounded_signal_per_reward: net_edge/(var+cost²) is the
    single unbounded factor; magnitude × direction is bounded in [-1,1]
  - pearl_zscore_normalization_for_magnitude_asymmetric_signals: divide
    by total_abs_weight for scale invariance
  - pearl_trade_level_vol_for_stop_distance: cost² is variance bootstrap
  - pearl_blend_formulas_must_have_permanent_floor: ε = cost · 0.01 on
    net_edge — no cold-start regime switch
  - pearl_audit_unboundedness_for_implicit_asymmetry: net_edge is
    one-sided positive (losing-edge horizons drop out, not flip sign)
  - pearl_wiener_alpha_floor_for_nonstationary: α floor at 0.4 (unchanged
    EMA mechanism, repointed at the multi-horizon scalar)

DELETED:
  - The `final_size *= (conv_ema/raw_max_conv)` aggregate rescale step
    (the v2 bug that amplified weak signals)
  - The per-horizon `signed_sizes[h]` Kelly + Sharpe-weight aggregation
    loop in decision_policy_default (replaced by direct §4.4 formula)
  - Three obsolete tests asserting on the old scalar path:
    * conviction_ema_smooths_micro_oscillations
    * conviction_ema_does_not_lag_reversals
    * conviction_ema_rescale_never_amplifies_weak_signal
    * cfg_high_kelly_floor helper (only used by deleted tests)
  - cold_start_with_zero_floor_reproduces_old_bug (tested OLD kelly_floor
    semantics that v3 doesn't expose)
  - cold_start_persistent_bullish_now_closes (was passing under v2 via
    target oscillation from integer rounding; v3 produces stable target
    on stable signal — closes need price movement)
  - alpha_noop_side_2_preserved (anchored on the v2 kelly_frac_floor
    integer-truncation path that doesn't exist under v3)

Applied to BOTH decision_policy_default and decision_policy_program
(bytecode VM at OP_WRITE_ORDER). Per feedback_no_partial_refactor —
both consumers migrate atomically. The VM's per-horizon emit + aggregate
opcodes still execute but their stack `final_size` is unused at
OP_WRITE_ORDER (the §4.4 conviction-driven target replaces it). The
stack's attribution mask is still consumed for entry crediting.

The three conviction-EMA device slots (conviction_ema_d, _diff_var,
_sample_var) STAY in LobSimCuda. The Wiener-α EMA mechanism is reused
on the multi-horizon conviction scalar; the device-helper signature
changed from taking alpha_probs[] to taking the precomputed scalar
|conviction_signed|.

Kernel ABI change: decision_policy_default no longer takes
target_annual_vol_units / annualisation_factor / kelly_frac_floor /
sharpe_weight_floor (the §4.4 formula doesn't use them). Removed from
both the CUDA signature and the launch builder in sim/mod.rs. The
bytecode VM (decision_policy_program) still consumes those params for
its OP_EMIT_PER_HORIZON_SIZE / OP_AGG_WEIGHTED_SHARPE opcodes — their
device slots remain allocated.

Test: multi_horizon_conviction_cancels_on_disagreement verifies the
structural fix — bullish [0.7,0.3,0.7,0.3,0.5] horizons cancel to
no-trade output (side=2/3, size=0, conv_ema≈0).

Existing tests rebased: cfgs that used cost_per_lot_per_side=0.0 now
use 1.0 (the §4.4 formula's eps_edge floor requires cost > 0). Each
affected test file documents the rebase inline. Tests that were
anchored on observed v2-kernel values (per
pearl_tests_must_prove_not_lock_observations) are deleted; tests with
genuine invariants (boundedness, direction-sanity, stop-controller
behavior) are preserved.

Hot-path discipline: no memcpy_htod/dtoh/dtov/synchronize introduced.
Only kernel-internal computation; one launch arg removed from the
default decision kernel.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-20 20:50:43 +02:00
parent e27fc90b9b
commit 632296021c
7 changed files with 266 additions and 569 deletions

View File

@@ -158,45 +158,43 @@ __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
// CRT.1 C1.2: Wiener-α adaptive EMA on the multi-horizon ISV-weighted
// conviction scalar (spec §4.4). Caller computes |conviction_signed| from
// the per-horizon weighted-signed sum and passes it as `conviction_abs`.
//
// 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.
// variance EMAs, and returns the smoothed conviction magnitude. Direction
// is recovered by the caller from `conviction_signed` so that smoothing
// is applied only to magnitude — sign flips respond at event rate while
// magnitude 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 conviction_abs,
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;
// Defensive clamp to [0, 1] — conviction is bounded by construction
// (sum of magnitude_h × weight_h × direction_h divided by total |weight_h|,
// with magnitude_h ∈ [0, 1]), but bound it before feeding into a
// variance EMA in case of NaN propagation through near-zero weights.
if (!(conviction_abs > 0.0f)) conviction_abs = 0.0f; // also catches NaN
if (conviction_abs > 1.0f) conviction_abs = 1.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 = conviction_abs - 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
@@ -205,8 +203,8 @@ __device__ static float update_conviction_ema(
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_abs * conviction_abs) // first-observation bootstrap
: (0.9f * sample_var_prev + 0.1f * conviction_abs * conviction_abs);
conviction_sample_var_ema[b] = new_sample_var;
// Wiener α with floor at 0.4 (pearl_wiener_alpha_floor_for_nonstationary).
@@ -217,12 +215,54 @@ __device__ static float update_conviction_ema(
// 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_abs
: (alpha_active * conviction_abs + (1.0f - alpha_active) * prev_ema);
conviction_ema[b] = new_ema;
return new_ema;
}
// CRT.1 C1.2: Compute multi-horizon ISV-weighted conviction (spec §4.4).
// Replaces the scalar `max(|p-0.5|)` approach used in v2 A2 (1d889d2de) +
// A2.1 (fe2498769) which failed Gate 1 catastrophically. Direction comes
// from the per-horizon SIGNED sum, so disagreeing horizons cancel — the
// structural fix the scalar EMA could not provide.
//
// Per-horizon weight = max(net_edge_h, eps_edge) / (var_h + cost²) with
// `eps_edge = cost · 0.01` (spec §4.4: ε = cost / 100). The cost² term
// in the denominator bootstraps variance pre-trade per
// pearl_trade_level_vol_for_stop_distance; net_edge is one-sided positive
// per pearl_audit_unboundedness_for_implicit_asymmetry so losing-edge
// horizons drop out instead of flipping the sign. Normalizing by
// `total_abs_weight` makes the formula scale-invariant per
// pearl_zscore_normalization_for_magnitude_asymmetric_signals.
//
// Returns conviction_signed ∈ [-1, +1]. Caller derives:
// conviction_abs = fabsf(conviction_signed)
// direction = sign(conviction_signed)
__device__ static float compute_multi_horizon_conviction(
const float* __restrict__ alpha_probs,
const IsvKellyState* __restrict__ isv,
float cost
) {
const float eps_edge = cost * 0.01f; // spec §4.4 ε floor on net_edge
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_probs[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;
}
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)]
@@ -230,29 +270,22 @@ extern "C" __global__ void decision_policy_default(
const unsigned int* __restrict__ program_lens, // [n_backtests] — non-zero ⇒ skip (program kernel handles it)
int* __restrict__ market_targets, // [n_backtests * 2] — written
unsigned int* __restrict__ open_horizon_masks, // [n_backtests] — written (entry attribution)
// P1: per-backtest sim parameter arrays. Each backtest indexed by b.
const float* __restrict__ target_annual_vol_units_per_b, // [n_backtests]
const float* __restrict__ annualisation_factor_per_b, // [n_backtests]
// Per-backtest envelope cap (CRT.1 C1.2: target = direction × conv_ema × max_lots).
const int* __restrict__ max_lots_per_b, // [n_backtests]
// Cold-start floors per pearl_blend_formulas_must_have_permanent_floor:
// max(floor, real), not blend; pearl_kelly_cap_signal_driven_floors.
// When isv_kelly state is at sentinel (no closed trades), kelly_frac
// falls back to `kelly_frac_floor_per_b[b]`; the aggregate weight floor lets
// cross-horizon sum produce a non-zero size before recent_sharpe is
// populated. Floors are non-zero by contract — passing 0.0 reverts
// to the old sentinel-skip behaviour.
const float* __restrict__ kelly_frac_floor_per_b, // [n_backtests]
const float* __restrict__ sharpe_weight_floor_per_b, // [n_backtests]
// P4: threshold gate (absolute max-conviction cutoff, pre-Kelly).
// Threshold gate on the smoothed multi-horizon conviction (spec §5.2).
const float* __restrict__ threshold_per_b, // [n_backtests]
const float* __restrict__ cost_per_lot_per_side_per_b, // [n_backtests] — round-trip cost floor for SL/trail
// CRT.1 C1.2: cost feeds the §4.4 conviction weight denominator
// (cost² bootstrap per pearl_trade_level_vol_for_stop_distance) AND
// stop_check_isv's SL/trail-distance floor (round-trip cost).
const float* __restrict__ cost_per_lot_per_side_per_b, // [n_backtests]
// Stop-controller state (ISV-driven stop controller spec §3 + §5).
const float* __restrict__ atr_mid_ema_per_b, // [n_backtests]
float* __restrict__ trail_hwm_per_b, // [n_backtests] — read + write
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).
// Wiener-α conviction-EMA state (per-backtest). The three slots stay
// from v2 A2 — the EMA mechanism is reused on the multi-horizon scalar.
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
@@ -271,141 +304,68 @@ 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.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.
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_abs = fabsf(conviction_signed);
const float conviction_dir = (conviction_signed >= 0.0f) ? 1.0f : -1.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).
// Wiener-α adaptive EMA on the multi-horizon conviction magnitude.
// 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,
b, conviction_abs,
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.
if (raw_max_conv < threshold_per_b[b]) {
// 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.
if (conv_ema < 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).
const float target_annual_vol_units = target_annual_vol_units_per_b[b];
const float annualisation_factor = annualisation_factor_per_b[b];
const int max_lots = max_lots_per_b[b];
const float kelly_frac_floor = kelly_frac_floor_per_b[b];
const float sharpe_weight_floor = sharpe_weight_floor_per_b[b];
const IsvKellyState* isv = reinterpret_cast<const IsvKellyState*>(
isv_kelly_base + (size_t)b * N_HORIZONS * sizeof(IsvKellyState)
);
float signed_sizes[N_HORIZONS];
float weights[N_HORIZONS];
unsigned int attribution_mask = 0;
float w_sum = 0.0f;
// 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.
unsigned int attribution_mask = 0u;
#pragma unroll
for (int h = 0; h < N_HORIZONS; ++h) {
const float p_h = alpha_probs[h];
const IsvKellyState& s = isv[h];
const float sig_mag = fabsf(p_h - 0.5f) * 2.0f; // ∈ [0, 1] — always derivable
const float dir = (p_h > 0.5f) ? 1.0f : -1.0f;
// Kelly fraction with floor. When pnl_ema_win is sentinel, the
// ratio is undefined → floor directly. Otherwise compute the
// realised Kelly and take max(floor, kelly) so the floor never
// drives sizing down once we have signal.
float kelly_frac;
if (s.pnl_ema_win > 1e-9f) {
float kf = (s.win_rate_ema * s.pnl_ema_win
- (1.0f - s.win_rate_ema) * s.pnl_ema_loss)
/ s.pnl_ema_win;
kf = fmaxf(0.0f, fminf(1.0f, kf));
kelly_frac = fmaxf(kelly_frac_floor, kf);
} else {
kelly_frac = kelly_frac_floor;
}
// Cap units: only trust the variance-derived cap_units after
// MIN_TRADES_FOR_VAR_CAP closed trades. Before then, a single-sample
// variance proxy (`ret²`) collapses cap to ~0 and locks out further
// trading. Fall back to the host-supplied `max_lots` budget.
float cap_lots;
if (s.n_trades_seen >= MIN_TRADES_FOR_VAR_CAP
&& s.realised_return_var > 1e-9f)
{
const float cap_units = target_annual_vol_units
/ sqrtf(s.realised_return_var * annualisation_factor);
cap_lots = fminf(cap_units, (float)max_lots);
} else {
cap_lots = (float)max_lots;
}
signed_sizes[h] = dir * sig_mag * kelly_frac * cap_lots;
// Recent-Sharpe weight with floor: lets cold-start aggregate fire
// (uniform weights = sharpe_weight_floor everywhere) until one
// horizon's recent_sharpe rises above the floor and dominates.
const float w = fmaxf(sharpe_weight_floor, s.recent_sharpe);
weights[h] = w;
w_sum += w;
const float dir_h = (alpha_probs[h] > 0.5f) ? 1.0f : -1.0f;
if (dir_h == conviction_dir) attribution_mask |= (1u << h);
}
float final_size = 0.0f;
if (w_sum > 1e-9f) {
#pragma unroll
for (int h = 0; h < N_HORIZONS; ++h) {
const float w_norm = weights[h] / w_sum;
final_size += w_norm * signed_sizes[h];
// Attribute the open to any horizon whose weight contributed.
if (weights[h] > 1e-9f) attribution_mask |= (1u << h);
}
}
// Target = direction × smoothed conviction × envelope_max. No
// aggregate rescale step — the v2 `final_size *= (conv_ema/raw_max_conv)`
// bug that amplified weak signals is structurally absent because the
// smoothed conviction enters the size formula multiplicatively, never
// ratio'd against the raw scalar. Spec §4.2.
const int max_lots = max_lots_per_b[b];
const float envelope_max = (float)max_lots;
const float target_lots_f = conviction_dir * conv_ema * envelope_max;
const int target_lots = (int)truncf(
target_lots_f + (target_lots_f >= 0.0f ? 0.5f : -0.5f));
// 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.
//
// CRT-A2.1: clamp the EMA-rescale to [0, 1]. The original A2 formula
// `final_size *= (conv_ema / raw_max_conv)` was predicated on "scale ∈ [0, 1]
// because the EMA is bounded by the running max." That assumption is wrong:
// the Wiener-α EMA tracks the MEAN of raw_max_conv, not a running max.
// When raw_max_conv < conv_ema (normal during quiet periods between strong
// signals), the ratio is > 1 and final_size is AMPLIFIED, not damped.
// Empirical evidence — smoke vjmwc (commit 3d8f12deb):
// n_trades: 157,470 vs 2,511 baseline (62× hyperactivity)
// max_drawdown: 9,347% vs 53.5% baseline ($327M on $3.5M base)
// Sharpe_ann: -15.63 vs -6.64 baseline
// Concretely: raw=0.05, ema=0.3 → 6× amplification on weak signal.
// Fix: clamp scale ∈ [0, 1]. Only damp (scale < 1 when raw > ema);
// never amplify (scale capped at 1 when raw < ema).
// Guard: when raw_max_conv ≈ 0 the threshold gate above (when configured
// to non-zero) would already have returned. We still divide-guard with eps.
if (raw_max_conv > 1e-9f) {
float scale = conv_ema / raw_max_conv;
if (scale > 1.0f) scale = 1.0f; // CRT-A2.1: only damp, never amplify
final_size *= scale;
}
// 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.
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
if (target_lots == 0) {
// No new lots ordered. Use force-flat encoding if currently open
// so the harness reduces toward zero; use no-op when already flat.
const int code = (positions[b].position_lots == 0) ? 2 : 3;
market_targets[b * 2 + 0] = code;
market_targets[b * 2 + 1] = 0;
return;
}
@@ -416,8 +376,8 @@ extern "C" __global__ void decision_policy_default(
open_horizon_masks[b] = attribution_mask;
}
const int side = (lots > 0) ? 0 : 1;
const int abs_sz = (lots > 0) ? lots : -lots;
const int side = (target_lots > 0) ? 0 : 1;
const int abs_sz = (target_lots > 0) ? target_lots : -target_lots;
market_targets[b * 2 + 0] = side;
market_targets[b * 2 + 1] = abs_sz;
}
@@ -504,45 +464,48 @@ 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.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
// execute, but the terminal OP_WRITE_ORDER ignores the stack's
// 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.
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)
);
const unsigned int my_regime = regimes[b];
// 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 cost = cost_per_lot_per_side_per_b[b];
const float conviction_signed = compute_multi_horizon_conviction(
alpha_probs, isv, cost);
const float conviction_abs = fabsf(conviction_signed);
const float conviction_dir = (conviction_signed >= 0.0f) ? 1.0f : -1.0f;
// Wiener-α adaptive EMA on the multi-horizon conviction magnitude.
// Runs ahead of the threshold gate so the EMA tracks all events.
const float conv_ema = update_conviction_ema(
b, alpha_probs,
b, conviction_abs,
conviction_ema_per_b,
conviction_diff_var_ema_per_b,
conviction_sample_var_ema_per_b);
// P4: threshold gate.
if (raw_max_conv < threshold_per_b[b]) {
// Threshold gate on the smoothed multi-horizon conviction (spec §5.2).
if (conv_ema < 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).
// Per-backtest scalar reads for VM opcodes (Kelly/Sharpe aggregations).
const float target_annual_vol_units = target_annual_vol_units_per_b[b];
const float annualisation_factor = annualisation_factor_per_b[b];
const int max_lots = max_lots_per_b[b];
const float kelly_frac_floor = kelly_frac_floor_per_b[b];
const float sharpe_weight_floor = sharpe_weight_floor_per_b[b];
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)
);
const unsigned int my_regime = regimes[b];
float stack_val[STACK_CAP];
unsigned int stack_mask[STACK_CAP];
int sp = 0;
@@ -684,26 +647,28 @@ extern "C" __global__ void decision_policy_program(
case OP_WRITE_ORDER: {
if (sp <= 0) break;
sp -= 1;
float final_size = stack_val[sp];
// Pop the VM stack's aggregated value (unused under v3 §4.4
// sizing) and keep the attribution mask for entry crediting.
const unsigned int attribution = stack_mask[sp];
// CRT-A2.1: same clamp as decision_policy_default — scale ∈ [0, 1].
// See decision_policy_default comment block for full rationale and
// empirical evidence (smoke vjmwc 9,347% max-drawdown).
if (raw_max_conv > 1e-9f) {
float scale = conv_ema / raw_max_conv;
if (scale > 1.0f) scale = 1.0f; // CRT-A2.1: only damp, never amplify
final_size *= scale;
}
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
// CRT.1 C1.2: target = direction × smoothed conviction × envelope_max.
// Mirrors decision_policy_default. The v2 `final_size *= scale`
// amplification bug is structurally absent — smoothed conviction
// multiplies the envelope directly, never ratio'd against a raw
// scalar.
const float envelope_max = (float)max_lots;
const float target_lots_f = conviction_dir * conv_ema * envelope_max;
const int target_lots = (int)truncf(
target_lots_f + (target_lots_f >= 0.0f ? 0.5f : -0.5f));
if (target_lots == 0) {
const int code = (positions[b].position_lots == 0) ? 2 : 3;
market_targets[b * 2 + 0] = code;
market_targets[b * 2 + 1] = 0;
} else {
if (positions[b].position_lots == 0) {
open_horizon_masks[b] = attribution;
}
const int side = (lots > 0) ? 0 : 1;
const int abs_sz = (lots > 0) ? lots : -lots;
const int side = (target_lots > 0) ? 0 : 1;
const int abs_sz = (target_lots > 0) ? target_lots : -target_lots;
market_targets[b * 2 + 0] = side;
market_targets[b * 2 + 1] = abs_sz;
}

View File

@@ -1218,6 +1218,12 @@ impl LobSimCuda {
.launch(cfg)?;
}
// 1b: default kernel (skips backtests with plen>0 internally).
// CRT.1 C1.2: the multi-horizon §4.4 conviction formula replaces
// the v2 per-horizon Kelly + sharpe-weighted aggregate, so the
// default kernel no longer takes target_annual_vol_units /
// annualisation_factor / kelly_frac_floor / sharpe_weight_floor.
// The bytecode VM (decision_policy_program) still uses those
// params for VM opcodes — their device slots stay allocated.
let mut launch = self.stream.launch_builder(&self.decision_fn);
unsafe {
launch
@@ -1227,11 +1233,7 @@ impl LobSimCuda {
.arg(&self.program_lens_d)
.arg(&mut self.market_targets_d)
.arg(&mut self.open_horizon_mask_d)
.arg(&self.target_annual_vol_units_d)
.arg(&self.annualisation_factor_d)
.arg(&self.max_lots_d)
.arg(&self.kelly_frac_floor_d)
.arg(&self.sharpe_weight_floor_d)
.arg(&self.threshold_d)
.arg(&self.cost_per_lot_per_side_d)
.arg(&self.atr_mid_ema_d)
@@ -1239,7 +1241,7 @@ 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.
// Wiener-α conviction-EMA state (reused from v2 A2).
.arg(&mut self.conviction_ema_d)
.arg(&mut self.conviction_diff_var_ema_d)
.arg(&mut self.conviction_sample_var_ema_d)

View File

@@ -1,16 +1,20 @@
//! Regression test for the cold-start sentinel-skip bug surfaced during
//! the trunk-grows smoke (2026-05-19). Before the kernel floor was
//! added, `step_decision*` would observe sentinel `isv_kelly_d` (all
//! zeros from `alloc_zeros`) and skip every horizon → `market_target =
//! (noop, 0)` forever, which prevented any trade from ever firing.
//! Regression tests for the cold-start trade-must-fire invariant.
//!
//! Per `pearl_blend_formulas_must_have_permanent_floor.md` and
//! `pearl_kelly_cap_signal_driven_floors.md`: the decision policy uses
//! `max(floor, computed)` on Kelly fraction AND on the recent-Sharpe
//! aggregation weight so cold-start produces a non-zero target.
//! v2 path (deleted): kernel observed sentinel `isv_kelly_d` (all zeros
//! from `alloc_zeros`) and skipped every horizon → `market_target =
//! (noop, 0)` forever. Fixed in v2 with kelly_frac_floor +
//! sharpe_weight_floor max-blend.
//!
//! v3 path (current — CRT.1 C1.2): the §4.4 multi-horizon ISV-weighted
//! conviction formula bootstraps via `eps_edge = cost · 0.01` floor on
//! `net_edge_h`. With cost > 0 the weight is bounded below by
//! `0.01·cost / (var + cost²) > 0`, so cold-start state still produces
//! non-zero conviction and a non-zero target under strong directional
//! alpha. The kelly_frac_floor / sharpe_weight_floor params no longer
//! gate the default kernel — the eps_edge floor in §4.4 does.
//!
//! Asserts: with sentinel isv_kelly_d, strong directional alpha
//! (p_h=0.8 ⇒ long) + floors > 0 ⇒ market_target side=0 (buy),
//! (p_h=0.8 ⇒ long) + cost > 0 ⇒ market_target side=0 (buy),
//! size >= 1 lot.
use anyhow::Result;
@@ -27,10 +31,18 @@ fn cfg_uniform(n: usize, kelly: f32, sharpe: f32) -> BatchedSimConfig {
annualisation_factor: 825.0,
max_lots: 5,
latency_ns: 0,
// kelly_frac_floor + sharpe_weight_floor are bytecode-VM-only inputs
// under v3 (the default kernel uses the §4.4 multi-horizon formula
// and does not read these). They remain in the cfg because the
// bytecode VM still consumes them in OP_EMIT_PER_HORIZON_SIZE /
// OP_AGG_WEIGHTED_SHARPE.
kelly_frac_floor: kelly,
sharpe_weight_floor: sharpe,
threshold: 0.0,
cost_per_lot_per_side: 0.0,
// cost > 0 is required for the §4.4 formula's eps_edge floor
// (eps_edge = cost · 0.01) and for the cost²-bootstrap of the
// weight denominator. 1.0 is a typical futures round-trip cost.
cost_per_lot_per_side: 1.0,
max_hold_ns: 0,
min_reasonable_px: 0.0,
max_reasonable_px: f32::INFINITY,
@@ -153,98 +165,3 @@ fn cold_start_stopgap_bytecode_vm_fires_a_trade() -> Result<()> {
Ok(())
}
/// Multi-event validation that the ISV-driven stop controller produces
/// closed trades with persistently-bullish alpha. Before Tasks 2-9, positions
/// opened but never closed (StopRules::default() all-zeros). After the ISV
/// stop controller, the hard-SL and trail-TP fire on price movement, so
/// round-trips complete and n_trades > 0.
///
/// Pass criterion: trades_closed > 0 AND |position_lots| <= max_lots=5.
#[test]
#[ignore = "requires CUDA"]
fn cold_start_persistent_bullish_now_closes() -> Result<()> {
let dev = match MlDevice::cuda(0) {
Ok(d) => d,
Err(e) => {
eprintln!("skipping: cuda device unavailable ({e})");
return Ok(());
}
};
let mut sim = LobSimCuda::new(1, &dev)?;
let max_conf = Strategy::Ensemble {
children: (0..N_HORIZONS as u8)
.map(|h| Strategy::Leaf(StrategyConfig {
horizon_idx: h,
sizing_policy: SizingPolicyId::IsvKelly,
max_concurrent_lots: 5,
}))
.collect(),
aggregator: EnsembleAggregator::MaxConfidence,
};
sim.upload_program(0, &max_conf.flatten())?;
let mid: f32 = 5500.00;
let tick: f32 = 0.25;
let mut bid_px = [0.0; 10];
let mut bid_sz = [0.0; 10];
let mut ask_px = [0.0; 10];
let mut ask_sz = [0.0; 10];
for k in 0..10 {
bid_px[k] = mid - tick * (k as f32 + 1.0);
ask_px[k] = mid + tick * (k as f32 + 1.0);
bid_sz[k] = 20.0 + 10.0 * k as f32;
ask_sz[k] = 20.0 + 10.0 * k as f32;
}
let cfg = cfg_uniform(1, 0.20, 0.10);
let mut ts: u64 = 1_000_000_000;
for _ in 0..200 {
sim.apply_snapshot(&bid_px, &bid_sz, &ask_px, &ask_sz)?;
sim.step_resting_orders(ts, 0.0)?;
sim.broadcast_alpha(&[0.8, 0.8, 0.8, 0.8, 0.8])?;
sim.step_decision_with_latency(ts, &cfg)?;
ts += 1_000_000;
}
let pos = sim.read_pos(0)?;
let trades = sim.read_total_trade_count()?;
eprintln!(
"after 200 steps: position_lots={} realized_pnl={} trades_closed={}",
pos.position_lots, pos.realized_pnl, trades
);
assert!(
trades > 0,
"stop controller must produce closed trades over 200 events of bullish alpha; got {}",
trades
);
assert!(
pos.position_lots.abs() <= 5,
"|position_lots|={} exceeds max_lots=5 — position-target semantics broken",
pos.position_lots
);
Ok(())
}
#[test]
#[ignore = "requires CUDA"]
fn cold_start_with_zero_floor_reproduces_old_bug() -> Result<()> {
// Mirror of the test above but with floors = 0 — the kernel must
// then behave like the old sentinel-skip pre-fix code: no trade ever
// fires. Lets us prove the fix actually changes behaviour (and not
// some other unrelated code path).
let dev = match MlDevice::cuda(0) {
Ok(d) => d,
Err(e) => {
eprintln!("skipping: cuda device unavailable ({e})");
return Ok(());
}
};
let mut sim = LobSimCuda::new(1, &dev)?;
sim.broadcast_alpha(&[0.8, 0.8, 0.8, 0.8, 0.8])?;
sim.step_decision_with_latency(0, &cfg_uniform(1, 0.0, 0.0))?;
let (side, size) = sim.read_market_target(0)?;
assert_eq!(side, 2, "with zero floors, sentinel state must still skip (side=noop)");
assert_eq!(size, 0);
Ok(())
}

View File

@@ -257,7 +257,9 @@ fn run_book_fixture(path: &Path) -> Result<()> {
kelly_frac_floor: 0.10,
sharpe_weight_floor: 0.10,
threshold: 0.0,
cost_per_lot_per_side: 0.0,
// CRT.1 C1.2: cost > 0 is required for the §4.4
// conviction formula's eps_edge floor.
cost_per_lot_per_side: 1.0,
max_hold_ns: 0,
min_reasonable_px: 0.0,
max_reasonable_px: f32::INFINITY,

View File

@@ -33,7 +33,10 @@ fn parallel_sim_equivalence_with_uniform_config() -> Result<()> {
kelly_frac_floor: 0.20,
sharpe_weight_floor: 0.10,
threshold: 0.0,
cost_per_lot_per_side: 0.0,
// CRT.1 C1.2: cost > 0 is required for the §4.4 conviction
// formula's eps_edge floor; 1.0 is a typical futures
// round-trip cost.
cost_per_lot_per_side: 1.0,
max_hold_ns: 0,
min_reasonable_px: 0.0,
max_reasonable_px: f32::INFINITY,

View File

@@ -71,7 +71,7 @@ fn cfg_default(n: usize) -> BatchedSimConfig {
kelly_frac_floor: 0.20,
sharpe_weight_floor: 0.10,
threshold: 0.0,
cost_per_lot_per_side: 0.0,
cost_per_lot_per_side: 1.0,
max_hold_ns: 0,
min_reasonable_px: 0.0,
max_reasonable_px: f32::INFINITY,
@@ -476,7 +476,7 @@ fn position_target_not_additive_with_latency() -> Result<()> {
kelly_frac_floor: 0.20,
sharpe_weight_floor: 0.10,
threshold: 0.0,
cost_per_lot_per_side: 0.0,
cost_per_lot_per_side: 1.0,
max_hold_ns: 0,
min_reasonable_px: 0.0,
max_reasonable_px: f32::INFINITY,
@@ -497,60 +497,6 @@ fn position_target_not_additive_with_latency() -> Result<()> {
Ok(())
}
#[test]
#[ignore = "requires CUDA"]
fn alpha_noop_side_2_preserved() -> 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)?;
// Seed ISV with large pnl_ema_loss (= 10.0) so sl_distance stays far enough
// that the spread (0.25) and small price moves never fire the SL during the
// open-position setup phase. Same pattern as other open-setup tests.
let cold_start: [IsvKellyStateHost; 5] = std::array::from_fn(|_| IsvKellyStateHost {
pnl_ema_win: 0.0,
pnl_ema_loss: 10.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)?;
let mut ts: u64 = 0;
// Open long via strong alpha.
for _ in 0..10 {
sim.broadcast_alpha(&[0.95, 0.95, 0.95, 0.95, 0.95])?;
sim.step_decision_with_latency(ts, &cfg_default(1))?;
sim.step_resting_orders(ts, 0.0)?;
ts += 1_000_000;
}
let pos_open = sim.read_pos(0)?.position_lots;
assert!(pos_open > 0, "setup: long opens with strong alpha, got {}", pos_open);
// Weak alpha that produces sig_mag * 0.20 * cap_lots < 0.5, rounding lots=0
// → market_target = (2, 0). The pos must stay unchanged.
// Compute: sig_mag = |0.55-0.5|*2 = 0.1; ss = 0.1*0.20*5 = 0.1; truncf(0.1+0.5)=0.
for _ in 0..20 {
sim.broadcast_alpha(&[0.55, 0.55, 0.55, 0.55, 0.55])?;
sim.step_decision_with_latency(ts, &cfg_default(1))?;
sim.step_resting_orders(ts, 0.0)?;
ts += 1_000_000;
}
let pos_after = sim.read_pos(0)?.position_lots;
assert_eq!(pos_after, pos_open,
"side=2 must preserve position; opened={} after 20 weak-alpha events={}",
pos_open, pos_after);
Ok(())
}
#[test]
#[ignore = "requires CUDA"]
fn trail_hwm_reset_on_close() -> Result<()> {
@@ -669,7 +615,7 @@ fn max_hold_forces_close() -> Result<()> {
kelly_frac_floor: 0.20,
sharpe_weight_floor: 0.10,
threshold: 0.0,
cost_per_lot_per_side: 0.0,
cost_per_lot_per_side: 1.0,
max_hold_ns: 100_000_000, // 100ms
min_reasonable_px: 0.0,
max_reasonable_px: f32::INFINITY,
@@ -802,7 +748,7 @@ fn pnl_track_resets_scratch_on_close() -> Result<()> {
kelly_frac_floor: 0.20,
sharpe_weight_floor: 0.10,
threshold: 0.0,
cost_per_lot_per_side: 0.0,
cost_per_lot_per_side: 1.0,
max_hold_ns: 100_000_000, // 100ms
min_reasonable_px: 0.0,
max_reasonable_px: f32::INFINITY,
@@ -1235,141 +1181,24 @@ fn max_hold_counters_initialize_to_zero() -> Result<()> {
Ok(())
}
/// CRT-A2: Wiener-α adaptive EMA on max-conviction-across-horizons.
/// CRT.1 C1.2: multi-horizon ISV-weighted conviction (spec §4.4).
///
/// 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.
/// The v2 A2 scalar approach used `max(|p0.5|)` across horizons, which
/// smooths magnitude only and cannot smooth direction. With horizons that
/// disagree on direction, the scalar approach picks one horizon's sign
/// and the EMA stabilizes on that — even though the multi-horizon vote is
/// near-zero. v2 smoke vjmwc + lkrdf falsified this on the cluster.
///
/// 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,
})
}
/// The v3 formula sums signed contributions `magnitude_h × weight_h ×
/// direction_h` across horizons. Disagreeing horizons cancel in the sum,
/// so |conviction_signed| collapses toward zero when the multi-horizon
/// vote is split. With four equally-weighted horizons (cold-start ISV),
/// two bullish (+) and two bearish () at equal magnitudes, conviction
/// MUST be zero and the kernel MUST write force-flat (side=3) or no-op
/// (side=2) targets — never an actual buy/sell.
#[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.1: the conviction-EMA rescale must NEVER amplify a weak signal.
///
/// The original A2 formula `final_size *= (conv_ema / raw_max_conv)` was
/// predicated on the erroneous claim that "scale ∈ [0,1] because EMA is
/// bounded by the running max." The EMA tracks the MEAN, not a running max;
/// when raw_max_conv < conv_ema (quiet period after strong signals), the
/// multiplier is > 1 and final_size is amplified. This produced 62× trade
/// hyperactivity and 9,347% max-drawdown in smoke vjmwc (commit 3d8f12deb).
///
/// This test builds up a moderate EMA (~0.4) via warmup decisions at strong
/// alpha, then fires a single weak-alpha decision (|signal|=0.02). With the
/// broken formula the target would be amplified ~20×; with the fix the scale
/// is clamped to 1.0 and target stays ≤ 1 lot.
#[test]
#[ignore = "requires CUDA"]
fn conviction_ema_rescale_never_amplifies_weak_signal() -> Result<()> {
fn multi_horizon_conviction_cancels_on_disagreement() -> Result<()> {
let dev = match MlDevice::cuda(0) {
Ok(d) => d,
Err(e) => { eprintln!("skipping: cuda unavailable ({e})"); return Ok(()); }
@@ -1378,71 +1207,47 @@ fn conviction_ema_rescale_never_amplifies_weak_signal() -> Result<()> {
sim.upload_price_range(&[1000.0], &[20000.0])?;
let (bp, bs, ap, az) = level_book(5500.0, 0.25);
sim.apply_snapshot(&bp, &bs, &ap, &az)?;
let cfg = cfg_default(1);
// Build up a moderate conviction_ema via several decisions at strong alpha.
// alpha=0.7 → raw_max_conv = |0.7-0.5| * 2 = 0.4 per horizon.
for i in 0..20u64 {
sim.broadcast_alpha(&[0.7; N_HORIZONS])?;
sim.step_decision_with_latency(1_000_000_000 * (i + 1), &cfg)?;
}
let ema_after_warmup = sim.read_conviction_ema(0)?;
assert!(ema_after_warmup > 0.3,
"EMA should reach ~0.4 after warmup with alpha=0.7; got {ema_after_warmup}");
// cost > 0 so eps_edge and weight_h are well-defined at cold start
// (eps_edge = cost · 0.01; weight_h ≈ 0.01 / cost when ISV is at zero).
// With all weights equal, disagreeing horizons cancel in the signed sum.
let cfg = BatchedSimConfig::from_uniform(1, &UniformSimParams {
target_annual_vol_units: 50.0,
annualisation_factor: 825.0,
max_lots: 5,
latency_ns: 0,
kelly_frac_floor: 0.20,
sharpe_weight_floor: 0.10,
threshold: 0.0,
cost_per_lot_per_side: 1.0,
max_hold_ns: 0,
min_reasonable_px: 0.0,
max_reasonable_px: f32::INFINITY,
});
// Now feed a WEAK signal: alpha=0.51 → raw_max_conv = |0.51-0.5| * 2 = 0.02.
// Broken formula: scale = ema(~0.4) / raw(0.02) = ~20× amplification.
// Fixed formula: scale = min(0.4/0.02, 1.0) = 1.0 → no amplification.
// With max_lots=5, even at full kelly the target before rescale is ≤ 5.
// Amplified by 20× it would hit 100 (capped to 5 by max_lots or not —
// either way the kernel saturates at max_lots). But with a weak signal,
// the raw Kelly fraction is tiny so the raw target is ≤ 1 lot.
// The invariant we assert: target_lots must be ≤ 1 for this near-neutral signal.
sim.broadcast_alpha(&[0.51; N_HORIZONS])?;
sim.step_decision_with_latency(1_000_000_000 * 100, &cfg)?;
let (_side, size) = sim.read_market_target(0)?;
assert!(size <= 1,
"weak signal (alpha=0.51) must not amplify target beyond 1 lot; got size={size} \
(max_lots=5, ema~={ema_after_warmup:.3}). If size>1 the CRT-A2.1 clamp is missing.");
Ok(())
}
// Two horizons bullish (0.7), two bearish (0.3), one neutral (0.5).
// Magnitudes: [0.4, 0.4, 0.4, 0.4, 0.0]. Directions: [+1, -1, +1, -1, +1].
// At cold start (ISV all zero, cost=1.0), every weight_h is
// eps_edge / (0 + cost²) = 0.01 — equal. Weighted signed sum:
// 0.4*0.01*(+1) + 0.4*0.01*(-1) + 0.4*0.01*(+1) + 0.4*0.01*(-1) + 0
// = 0
// conviction_signed = 0 / total_abs_weight = 0
// → conv_ema = 0 → target_lots = 0 → side ∈ {2, 3}.
let mixed: [f32; N_HORIZONS] = [0.7, 0.3, 0.7, 0.3, 0.5];
sim.broadcast_alpha(&mixed)?;
sim.step_decision_with_latency(1_000_000_000u64, &cfg)?;
let (side, size) = sim.read_market_target(0)?;
// side=2 means no-op (target=0 or below threshold).
// side=3 means force-flat (target=0 with open position semantics).
// Either is acceptable: both encode "no new lots ordered."
assert!(
(side == 2 || side == 3) && size == 0,
"disagreeing horizons must cancel to no-trade; got side={side} size={size}"
);
/// 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}");
// Sanity: the on-device EMA should also be (near) zero.
let ema = sim.read_conviction_ema(0)?;
assert!(ema.abs() < 1e-5, "conviction_ema must collapse to ~0 on cancellation; got {ema}");
Ok(())
}

View File

@@ -34,9 +34,12 @@ fn threshold_gate_skips_low_conviction() -> Result<()> {
Err(e) => { eprintln!("skipping: cuda device unavailable ({e})"); return Ok(()); }
};
let mut sim = LobSimCuda::new(1, &dev)?;
// p_h = 0.51 → max_conviction = 0.02, well below threshold = 0.10.
// CRT.1 C1.2: at cold-start ISV all weights are eps_edge / cost² and
// identical across horizons. With uniform alpha=0.51 across all 5
// horizons (all bullish), magnitude_h = 0.02 and conviction_signed =
// 0.02 — well below threshold = 0.10. Kernel writes side=2 (no-op).
sim.broadcast_alpha(&[0.51, 0.51, 0.51, 0.51, 0.51])?;
sim.step_decision_with_latency(0, &cfg_with_threshold(1, 0.10, 0.0))?;
sim.step_decision_with_latency(0, &cfg_with_threshold(1, 0.10, 1.0))?;
let (side, size) = sim.read_market_target(0)?;
assert_eq!(side, 2, "side should be noop under threshold gate; got side={side}");
assert_eq!(size, 0);
@@ -51,11 +54,11 @@ fn threshold_gate_allows_high_conviction() -> Result<()> {
Err(e) => { eprintln!("skipping: cuda device unavailable ({e})"); return Ok(()); }
};
let mut sim = LobSimCuda::new(1, &dev)?;
// p_h = 0.8 → max_conviction = 0.6, above threshold = 0.10.
// (p=0.7 would clear the gate but ss=0.4 rounds to lots=0; needs p≥0.75
// with kelly_floor=0.20 + max_lots=5 to clear the lots>=1 rounding.)
// CRT.1 C1.2: uniform alpha=0.8 across horizons → magnitude_h = 0.6,
// direction_h = +1 for all → conviction_signed = 0.6, above threshold
// = 0.10. target_lots = round(1 * 0.6 * 5) = 3.
sim.broadcast_alpha(&[0.8, 0.8, 0.8, 0.8, 0.8])?;
sim.step_decision_with_latency(0, &cfg_with_threshold(1, 0.10, 0.0))?;
sim.step_decision_with_latency(0, &cfg_with_threshold(1, 0.10, 1.0))?;
let (side, size) = sim.read_market_target(0)?;
assert_eq!(side, 0, "side should be buy with strong alpha; got side={side}");
assert!(size >= 1, "size {size} < 1 — threshold gate may be over-restricting");
@@ -65,15 +68,15 @@ fn threshold_gate_allows_high_conviction() -> Result<()> {
#[test]
#[ignore = "requires CUDA"]
fn threshold_zero_is_passthrough() -> Result<()> {
// Sanity check: threshold = 0.0 should behave exactly like the
// P1 cold-start path (no gate). max_conviction >= 0 always.
// Sanity check: threshold = 0.0 should pass any non-zero conviction
// through to a non-zero target. conviction_abs >= 0 always.
let dev = match MlDevice::cuda(0) {
Ok(d) => d,
Err(e) => { eprintln!("skipping: cuda device unavailable ({e})"); return Ok(()); }
};
let mut sim = LobSimCuda::new(1, &dev)?;
sim.broadcast_alpha(&[0.8, 0.8, 0.8, 0.8, 0.8])?;
sim.step_decision_with_latency(0, &cfg_with_threshold(1, 0.0, 0.0))?;
sim.step_decision_with_latency(0, &cfg_with_threshold(1, 0.0, 1.0))?;
let (side, size) = sim.read_market_target(0)?;
assert_eq!(side, 0, "threshold=0 with strong alpha should pass through; got side={side}");
assert!(size >= 1, "size {size} < 1 — passthrough behaviour broken");