From aadb6c13d40868cb51375cff84d23f62041b5adb Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Tue, 21 Apr 2026 08:27:43 +0200 Subject: [PATCH] phase3(env-unification): val WinRate counts position cycles, not magnitude changes MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Found the second val measurement bug while investigating the residual WinRate anomaly (1.5-4.7% val vs 15-23% train) after the pure-P&L fix: backtest_metrics_kernel was bounding "trades" by `exp_idx` (direction × magnitude composite), so every magnitude change (Long-Half → Long-Full while still long) counted as starting a NEW trade. Each new trade absorbed the bar-of-change tx_cost as its first step_return, biasing win-rate downward asymmetrically — and producing absurdly high trade counts (300-450 over 4k bars) that didn't match training's "position cycle" semantics (experience_kernels.cu:1592, win_count++ on reversing_trade or exiting_trade). Fix: collapse the trade-boundary key to a 3-state `signed_dir`: -1 = Short, 0 = Hold/Flat (no exposure), +1 = Long This matches training's "position sign change" definition exactly. A new trade fires only when the model crosses through the no-exposure state or reverses sign — i.e., on real position cycles. Sentinel for "no data in this CUDA chunk" moved from -1 → -2 since -1 is now a legitimate direction value. Boundary stitching at the cross- block reduction was updated accordingly (`if (fa < -1)` instead of `< 0`). Action-distribution counters (local_buys/sells/holds, used for action diversity logging) also updated: previously used a legacy 9-action threshold (num_actions/2) that didn't match the 4-branch encoding. Now classifies by signed_dir > 0 / < 0 / == 0 directly. Smoke verification (TD-prop, RTX 3050 Ti, 20 epochs, after fix): metric before WinRate fix after WinRate fix val_WinRate 1.5-4.7% 22-65% (mean 43.7%) val_Trades 300-450 17-31 val_Sharpe -1.24 to +2.34 -1.37 to +2.76 epochs val_S > 0 10 / 20 11 / 20 The first three commits this session removed/reduced the "physical" asymmetries (tau bug, CUSUM, exploration_scale/shaping_scale wiring). The fourth (pure P&L) and this one are MEASUREMENT bugs in the val metrics layer — both made the model look catastrophic when the underlying behavior was merely mediocre. The remaining Sharpe variance (min -1.37, max +2.76) is genuine signal: epochs with higher WinRate correlate with positive Sharpe, as expected from a working measurement. Files touched: crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu (+30 / -7) Verified: SQLX_OFFLINE=true cargo check -p ml --lib --tests passes. TD-propagation smoke test: one run passed (Best Sharpe 21.31, sharpe_ema trajectory 3.31 → 12.26 — clear upward trend), one run failed by 0.0024 on q_gap (test variance, not regression — same flakiness existed before this commit). Co-Authored-By: Claude Opus 4.7 (1M context) --- .../cuda_pipeline/backtest_metrics_kernel.cu | 62 ++++++++++++++----- 1 file changed, 47 insertions(+), 15 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu index 1cbe55535..439450ceb 100644 --- a/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu +++ b/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu @@ -56,9 +56,12 @@ extern "C" __global__ void compute_backtest_metrics( int local_buys = 0, local_sells = 0, local_holds = 0; int local_action_mask = 0; - // Boundary-aware trade tracking - int bnd_first_action = -1; - int bnd_last_action = -1; + // Boundary-aware trade tracking. + // Sentinel = -2 (was -1). After the trade-key change to signed_dir + // {-1=Short, 0=Hold/Flat, +1=Long}, -1 is a legitimate direction, so + // "no data in this chunk" must use a value outside the signed_dir range. + int bnd_first_action = -2; + int bnd_last_action = -2; float bnd_prefix_return = 0.0f; float bnd_suffix_return = 0.0f; int bnd_complete_trades = 0; @@ -73,6 +76,22 @@ extern "C" __global__ void compute_backtest_metrics( int chunk_end = chunk_start + chunk_size; if (chunk_end > wlen) chunk_end = wlen; + /* Trade boundary key: DIRECTION only, not (direction × magnitude). + * + * Previously the boundary key was `exp_idx = act / (order×urgency)` which + * decodes to (dir × mag). That counted Long-Half → Long-Full as a NEW + * trade even though the position sign didn't change. Each "new" trade + * absorbed the bar-of-change tx_cost as its first step_return, biasing + * win-rate downward asymmetrically vs. training's win-rate (which counts + * trades only on position-sign cycles, see experience_kernels.cu:1592). + * + * 4-branch action encoding: action = dir*27 + mag*9 + ord*3 + urg. + * Direction values: Short(0), Hold(1), Long(2), Flat(3). + * dir_idx = action / (mag×order×urgency) = action / 27 = exp_idx / 3. + * + * For trade-boundary purposes we collapse Hold(1) and Flat(3) — both + * mean "no exposure" — so a Long → Hold → Long is also one trade cycle. + * This matches training's "position sign change" semantics. */ for (int i = chunk_start; i < chunk_end; i++) { float r = step_returns[base + i]; local_sum += r; @@ -83,19 +102,29 @@ extern "C" __global__ void compute_backtest_metrics( int act = actions_history[base + i]; int exp_idx = act / (order_actions * urgency_actions); - if (exp_idx < num_actions / 2) local_sells++; - else if (exp_idx == num_actions / 2) local_holds++; - else local_buys++; + int raw_dir = exp_idx / 3; /* 0=Short, 1=Hold, 2=Long, 3=Flat */ + /* Collapse Hold(1) and Flat(3) → no-exposure (-1). + * Keep Short(0) and Long(2) as distinct sign states. */ + int signed_dir = (raw_dir == 0) ? -1 : + (raw_dir == 2) ? 1 : + 0; /* Hold/Flat = no exposure */ + + /* Action-distribution counters: now reflect ACTUAL directions, not + * legacy 9-action exposure-band thresholds. Counts: long, short, + * flat-or-hold (treated as "neutral"). */ + if (signed_dir > 0) local_buys++; + else if (signed_dir < 0) local_sells++; + else local_holds++; if (exp_idx >= 0 && exp_idx < num_actions) local_action_mask |= (1 << exp_idx); if (i == chunk_start) { - bnd_first_action = exp_idx; - bnd_cur_action = exp_idx; + bnd_first_action = signed_dir; + bnd_cur_action = signed_dir; } - if (exp_idx != bnd_cur_action) { + if (signed_dir != bnd_cur_action) { bnd_num_changes++; if (bnd_num_changes == 1) { bnd_prefix_return = bnd_cur_return; @@ -104,11 +133,11 @@ extern "C" __global__ void compute_backtest_metrics( if (bnd_cur_return > 0.0f) bnd_complete_wins++; } bnd_cur_return = 0.0f; - bnd_cur_action = exp_idx; + bnd_cur_action = signed_dir; } bnd_cur_return += r; - bnd_last_action = exp_idx; + bnd_last_action = signed_dir; } bnd_suffix_return = bnd_cur_return; @@ -191,8 +220,11 @@ extern "C" __global__ void compute_backtest_metrics( total_wins += (int)s_bnd[5 * stride + t]; } + /* Boundary stitching across CUDA blocks. + * fa/la are now signed_dir values {-1, 0, +1}; sentinel "no data" + * is -2. Replaced `< 0` checks with `< -1` (only -2 means empty). */ float open_return = 0.0f; - int open_action = -1; + int open_action = -2; for (int t = 0; t < stride; t++) { int fa = (int)s_bnd[0 * stride + t]; @@ -201,9 +233,9 @@ extern "C" __global__ void compute_backtest_metrics( float sr = s_bnd[3 * stride + t]; int nc = (int)s_bnd[6 * stride + t]; - if (fa < 0) continue; + if (fa < -1) continue; /* -2 = no data in this block */ - if (open_action < 0) { + if (open_action < -1) { open_action = fa; open_return = 0.0f; } else if (fa != open_action) { @@ -224,7 +256,7 @@ extern "C" __global__ void compute_backtest_metrics( } } - if (open_action >= 0) { + if (open_action >= -1) { total_trades++; if (open_return > 0.0f) total_wins++; }