phase3(env-unification): val WinRate counts position cycles, not magnitude changes

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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-04-21 08:27:43 +02:00
parent 1ffdf38ddc
commit aadb6c13d4

View File

@@ -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++;
}