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

feat(audit): per-regime val WR instrumentation (T/R/V buckets)

Browse Source 
Adds 6 output slots to compute_backtest_metrics_kernel — per-bucket
(win_rate, trade_count) for the {Trending, Ranging, Volatile} regime
split — so the val backtest surfaces whether the long-running ~46%
aggregate WR hides regime-conditional edge. Trades are bucketed at
trade-OPEN by feature[40] (ADX-norm) per the structural thresholds
(T:ADX>0.4, R:ADX<0.2, V:otherwise), mirroring
gpu_walk_forward.rs::classify_regime_from_features. Block tree-reduce
only per feedback_no_atomicadd. Observability-only emission via new
HEALTH_DIAG[N]: val_regime [wr_T=... n_T=... wr_R=... n_R=... wr_V=...
n_V=...] line; thresholds remain kernel constants per
feedback_isv_for_adaptive_bounds (no controller consumer yet).

Implementation atomic (kernel + launcher + WindowMetrics + HEALTH_DIAG
emit + 3 GPU oracle tests + audit doc):
- backtest_metrics_kernel.cu: per-thread per-regime trade counters,
  2-slot boundary buffer extension carrying open-bar regime through
  block stitch, output stride 13 → 19, shmem 5 → 11 reduction tiles
- gpu_backtest_evaluator.rs: WindowMetrics +6 fields, metrics_buf
  size 13 → 19, launcher passes features_buf + feature_dim, consume
  populates per-regime fields
- metrics.rs: val_regime HEALTH_DIAG line in consume_validation_loss
- regime_wr_oracle_tests.rs (NEW): 1 CPU sanity + 3 GPU oracle tests
  (bit-exact match ε=1e-5 vs CPU oracle on stratified 30/40/30 batch)

Validation: cargo check --workspace clean; 17/17 gpu_backtest_evaluator
unit tests pass; 1+3/4 regime WR tests pass on local RTX 3050 Ti
(1.89s); audit_sp18_consumers.sh --check exit 0.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-05-09 13:38:48 +02:00
parent 5ee5a1b655
commit e140392f86
5 changed files with 878 additions and 25 deletions
Download Patch File Download Diff File Expand all files Collapse all files

257
crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu
View File

@@ -9,7 +9,20 @@
// 14 → 13 floats per window; old slot 0 (sharpe) is removed and every
// remaining metric shifts down by one.
//
// Output per window [13 f32 values written to metrics_out]:
// 2026-05-09 — per-regime WR instrumentation (audit follow-up). Output
// stride grows 13 → 19 floats per window: 6 new slots [13..19) carrying
// per-regime win-rate + trade-count for the {Trending, Ranging, Volatile}
// bucketing introduced by `gpu_walk_forward.rs::classify_regime_from_features`.
// Regime is read from feature[40] (ADX-norm, post-`RollingPercentileRank`)
// per state_layout.cuh's market block; bucket thresholds are empirically
// grounded structural defaults (NOT tuned constants) — Trending: ADX>0.4,
// Ranging: ADX<0.2, Volatile otherwise. Each trade is attributed to the
// regime at its OPENING bar; boundary-stitched trades that span chunks
// inherit the regime from the chunk where they first appeared (block
// boundary buffers carry the open-bar regime alongside the existing
// prefix/suffix pair).
//
// Output per window [19 f32 values written to metrics_out]:
// [0] total_pnl (multiplicative compounding)
// [1] max_drawdown (exact sequential scan)
// [2] sortino_ratio
@@ -23,13 +36,59 @@
// [10] sell_count
// [11] hold_count
// [12] unique_action_count
// [13] win_rate_trending (regime T: ADX > 0.4)
// [14] trade_count_trending
// [15] win_rate_ranging (regime R: ADX < 0.2)
// [16] trade_count_ranging
// [17] win_rate_volatile (regime V: otherwise)
// [18] trade_count_volatile
// Per-regime ADX thresholds. Empirically-grounded structural defaults
// (post-`RollingPercentileRank` normalisation, where ADX-norm ∈ [0,1]):
// - ADX > 0.4 ⇒ Trending (raw ADX > 25 in the upper 60th percentile)
// - ADX < 0.2 ⇒ Ranging (raw ADX < 10 in the lower 20th percentile)
// - else ⇒ Volatile (mid-ADX, often choppy)
// Per `feedback_isv_for_adaptive_bounds` these are conceptually adaptive
// (could derive from observed distribution). They live as kernel-local
// constants only because per-regime WR is observability-only (does not
// feed any controller); if a controller ever consumes a per-regime WR
// signal, lift the thresholds to ISV slots.
#define REGIME_TRENDING_THRESHOLD 0.4f
#define REGIME_RANGING_THRESHOLD 0.2f
// Feature index for ADX-norm in the 42-dim market feature block. Mirrors
// `gpu_walk_forward.rs::classify_regime_from_features` (uses index 40)
// and `state_layout.cuh`'s SL_MARKET block — slot 40 is the post-rolling-
// percentile-ranked ADX, ∈ [0,1].
#define REGIME_ADX_FEATURE_IDX 40
// Regime bucket labels. -1 sentinel = "no data" (chunk-empty / boundary).
#define REGIME_BUCKET_T 0
#define REGIME_BUCKET_R 1
#define REGIME_BUCKET_V 2
#define REGIME_BUCKET_NONE -1
// Classify a bar's regime from its feature vector at `bar_idx`. Returns
// REGIME_BUCKET_T / R / V (never NONE — that sentinel is reserved for
// "no observation" in boundary buffers).
__device__ __forceinline__ int classify_bar_regime(
const float* __restrict__ features,
int bar_idx,
int feature_dim
) {
float adx = features[bar_idx * feature_dim + REGIME_ADX_FEATURE_IDX];
if (adx > REGIME_TRENDING_THRESHOLD) return REGIME_BUCKET_T;
if (adx < REGIME_RANGING_THRESHOLD) return REGIME_BUCKET_R;
return REGIME_BUCKET_V;
}
extern "C" __global__ void compute_backtest_metrics(
const float* __restrict__ step_returns, // [n_windows * max_len]
const float* __restrict__ portfolio_state, // [n_windows * 8]
const int* __restrict__ window_lens, // [n_windows]
const int* __restrict__ actions_history, // [n_windows * max_len]
float* metrics_out, // [n_windows * 13]
const float* __restrict__ features, // [n_windows * max_len * feature_dim]
int feature_dim,
float* metrics_out, // [n_windows * 19]
int n_windows,
int max_len,
float annualization_factor, // f32 scalar (sortino × factor; calmar × factor²)
@@ -44,20 +103,34 @@ extern "C" __global__ void compute_backtest_metrics(
int tid = threadIdx.x;
int stride = blockDim.x;
int base = w * max_len;
// Per-window features base offset. Each row at `bar_idx` lives at
// `features[feat_base + bar_idx * feature_dim + k]`.
int feat_base = w * max_len * feature_dim;
// Shared memory layout (5 reduction arrays + sort scratch).
// SP15 Phase 1.1.b removed s_sq_sum (was sharpe-only) — variance/std
// are now computed by the dedicated `sharpe_per_bar_kernel`. The
// remaining metrics use s_sum (mean for sortino + calmar daily_mean),
// s_down_sq (sortino's down_std), and the action-distribution
// counters (buys/sells/holds).
// Shared memory layout (5 reduction arrays + 6 per-regime arrays +
// sort scratch). SP15 Phase 1.1.b removed s_sq_sum (was sharpe-only)
// — variance/std are now computed by the dedicated
// `sharpe_per_bar_kernel`. The remaining metrics use s_sum (mean for
// sortino + calmar daily_mean), s_down_sq (sortino's down_std), and
// the action-distribution counters (buys/sells/holds).
//
// 2026-05-09 — per-regime WR adds 6 reduction arrays for internal
// (chunk-local) per-regime trade/win counters. Boundary-stitched
// trades are attributed by thread 0 using the open-regime carried
// through `s_bnd[7..9*stride]`.
extern __shared__ float shmem[];
float* s_sum = shmem;
float* s_down_sq = shmem + stride;
float* s_buys = shmem + 2*stride;
float* s_sells = shmem + 3*stride;
float* s_holds = shmem + 4*stride;
float* s_sorted = shmem + 5*stride; // [4096] for sort + boundary data
float* s_trades_T = shmem + 5*stride;
float* s_wins_T = shmem + 6*stride;
float* s_trades_R = shmem + 7*stride;
float* s_wins_R = shmem + 8*stride;
float* s_trades_V = shmem + 9*stride;
float* s_wins_V = shmem + 10*stride;
float* s_sorted = shmem + 11*stride; // [4096] for sort + boundary data
// Per-thread local accumulators — native F32
float local_sum = 0.0f;
@@ -66,6 +139,14 @@ extern "C" __global__ void compute_backtest_metrics(
int local_buys = 0, local_sells = 0, local_holds = 0;
int local_action_mask = 0;
// Per-regime trade counters — populated for INTERNAL complete trades
// (i.e., trades whose open AND close both fall within this thread's
// chunk). Boundary-stitched trades are handled by thread 0 below using
// the open-bar regime carried in `bnd_open_regime` shared-memory slot.
int local_trades_T = 0, local_wins_T = 0;
int local_trades_R = 0, local_wins_R = 0;
int local_trades_V = 0, local_wins_V = 0;
// 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
@@ -79,6 +160,17 @@ extern "C" __global__ void compute_backtest_metrics(
int bnd_num_changes = 0;
int bnd_cur_action = -1;
float bnd_cur_return = 0.0f;
// Regime at the bar where THIS chunk's first trade opens (i.e., the
// chunk's first valid action). Carried across the boundary stitch so
// thread 0's trade-counting loop attributes stitched trades to the
// correct regime bucket. REGIME_BUCKET_NONE if the chunk has no
// valid action (entire chunk is `act < 0` sentinel).
int bnd_open_regime = REGIME_BUCKET_NONE;
// Regime at the open bar of the trade currently being accumulated
// INSIDE this thread's chunk. Tracked locally so internal-trade
// closes can attribute themselves to the correct bucket; reset on
// each transition.
int bnd_cur_open_regime = REGIME_BUCKET_NONE;
// Consecutive chunks for correct drawdown tracking
int chunk_size = (wlen + stride - 1) / stride;
@@ -139,9 +231,17 @@ extern "C" __global__ void compute_backtest_metrics(
if (exp_idx >= 0 && exp_idx < num_actions)
local_action_mask |= (1 << exp_idx);
// Classify regime at the current bar (for trade-open attribution).
int cur_regime = classify_bar_regime(features, feat_base / feature_dim + i, feature_dim);
if (i == chunk_start) {
bnd_first_action = signed_dir;
bnd_cur_action = signed_dir;
// Open-bar regime for this chunk's prefix trade. Persists
// across the boundary stitch so thread 0 can attribute the
// stitched trade to the correct bucket.
bnd_open_regime = cur_regime;
bnd_cur_open_regime = cur_regime;
}
if (signed_dir != bnd_cur_action) {
@@ -149,11 +249,27 @@ extern "C" __global__ void compute_backtest_metrics(
if (bnd_num_changes == 1) {
bnd_prefix_return = bnd_cur_return;
} else {
// Internal complete trade close. Attribute to the regime
// captured at the trade's OPEN bar (held in
// `bnd_cur_open_regime` since the last transition).
bnd_complete_trades++;
if (bnd_cur_return > 0.0f) bnd_complete_wins++;
bool is_win = (bnd_cur_return > 0.0f);
if (is_win) bnd_complete_wins++;
if (bnd_cur_open_regime == REGIME_BUCKET_T) {
local_trades_T++;
if (is_win) local_wins_T++;
} else if (bnd_cur_open_regime == REGIME_BUCKET_R) {
local_trades_R++;
if (is_win) local_wins_R++;
} else if (bnd_cur_open_regime == REGIME_BUCKET_V) {
local_trades_V++;
if (is_win) local_wins_V++;
}
}
bnd_cur_return = 0.0f;
bnd_cur_action = signed_dir;
// New trade just opened at bar `i` — capture its regime.
bnd_cur_open_regime = cur_regime;
}
bnd_cur_return += r;
@@ -164,6 +280,10 @@ extern "C" __global__ void compute_backtest_metrics(
if (bnd_num_changes == 0) {
bnd_prefix_return = bnd_cur_return;
}
// Regime of the trade currently "open" at chunk-end. Equals
// `bnd_open_regime` if no transition occurred (single ongoing trade),
// else `bnd_cur_open_regime` (regime of the last transition's open).
int bnd_last_open_regime = bnd_cur_open_regime;
// Store to shared memory for parallel reduction
s_sum[tid] = local_sum;
@@ -171,6 +291,12 @@ extern "C" __global__ void compute_backtest_metrics(
s_buys[tid] = (float)local_buys;
s_sells[tid] = (float)local_sells;
s_holds[tid] = (float)local_holds;
s_trades_T[tid] = (float)local_trades_T;
s_wins_T[tid] = (float)local_wins_T;
s_trades_R[tid] = (float)local_trades_R;
s_wins_R[tid] = (float)local_wins_R;
s_trades_V[tid] = (float)local_trades_V;
s_wins_V[tid] = (float)local_wins_V;
__syncthreads();
// Block-level parallel reduction — native f32 += f32
@@ -181,6 +307,12 @@ extern "C" __global__ void compute_backtest_metrics(
s_buys[tid] += s_buys[tid + s];
s_sells[tid] += s_sells[tid + s];
s_holds[tid] += s_holds[tid + s];
s_trades_T[tid] += s_trades_T[tid + s];
s_wins_T[tid] += s_wins_T[tid + s];
s_trades_R[tid] += s_trades_R[tid + s];
s_wins_R[tid] += s_wins_R[tid + s];
s_trades_V[tid] += s_trades_V[tid + s];
s_wins_V[tid] += s_wins_V[tid + s];
}
__syncthreads();
}
@@ -195,7 +327,12 @@ extern "C" __global__ void compute_backtest_metrics(
__syncthreads();
}
// Export boundary trade data to shared memory.
// Export boundary trade data to shared memory. Slots 7-8 added for
// per-regime stitching: `bnd_open_regime` carries the regime at this
// chunk's first action (used when a stitched trade ENDS at this
// chunk's first action), `bnd_last_open_regime` carries the regime
// at this chunk's last trade-open bar (used when a stitched trade
// CONTINUES from this chunk's suffix into a subsequent chunk).
{
float* s_bnd = &s_sorted[stride];
s_bnd[0 * stride + tid] = (float)bnd_first_action;
@@ -205,6 +342,8 @@ extern "C" __global__ void compute_backtest_metrics(
s_bnd[4 * stride + tid] = (float)bnd_complete_trades;
s_bnd[5 * stride + tid] = (float)bnd_complete_wins;
s_bnd[6 * stride + tid] = (float)bnd_num_changes;
s_bnd[7 * stride + tid] = (float)bnd_open_regime;
s_bnd[8 * stride + tid] = (float)bnd_last_open_regime;
}
__syncthreads();
@@ -226,10 +365,26 @@ extern "C" __global__ void compute_backtest_metrics(
exact_max_dd = fmaxf(exact_max_dd, dd);
}
// Boundary stitching: exact trade count + win rate
// Boundary stitching: exact trade count + win rate.
//
// 2026-05-09 — also accumulate per-regime trade/win counts.
// INTERNAL trade closes (within a chunk's body) were attributed
// per-thread; here we sum those across threads. STITCHED trade
// closes (boundary-spanning) are attributed using the open-bar
// regime carried via `bnd_open_regime` / `bnd_last_open_regime`
// in `s_bnd[7..]` and `s_bnd[8..]`.
float* s_bnd = &s_sorted[stride];
int total_trades = 0;
int total_wins = 0;
// Per-regime accumulators. INTERNAL counts come from the
// pre-reduced `s_trades_T/R/V` + `s_wins_T/R/V` arrays
// (slots 0 after the parallel reduction above).
int trades_T = (int)s_trades_T[0];
int wins_T = (int)s_wins_T[0];
int trades_R = (int)s_trades_R[0];
int wins_R = (int)s_wins_R[0];
int trades_V = (int)s_trades_V[0];
int wins_V = (int)s_wins_V[0];
for (int t = 0; t < stride; t++) {
total_trades += (int)s_bnd[4 * stride + t];
@@ -241,6 +396,10 @@ extern "C" __global__ void compute_backtest_metrics(
* is -2. Replaced `< 0` checks with `< -1` (only -2 means empty). */
float open_return = 0.0f;
int open_action = -2;
// Open-bar regime of the running stitched trade. Sentinel = NONE
// (carries through if the stitch spans regime-less blocks; in
// practice every block with `fa != -2` has a valid regime).
int open_regime = REGIME_BUCKET_NONE;
for (int t = 0; t < stride; t++) {
int fa = (int)s_bnd[0 * stride + t];
@@ -248,36 +407,84 @@ extern "C" __global__ void compute_backtest_metrics(
float pr = s_bnd[2 * stride + t];
float sr = s_bnd[3 * stride + t];
int nc = (int)s_bnd[6 * stride + t];
int blk_open_regime = (int)s_bnd[7 * stride + t];
int blk_last_open_regime = (int)s_bnd[8 * stride + t];
if (fa < -1) continue; /* -2 = no data in this block */
if (open_action < -1) {
// First active block — capture initial open regime.
open_action = fa;
open_return = 0.0f;
open_regime = blk_open_regime;
} else if (fa != open_action) {
// Stitched trade closes here. Attribute to the regime
// captured at its open (carried in `open_regime`).
total_trades++;
if (open_return > 0.0f) total_wins++;
bool is_win = (open_return > 0.0f);
if (is_win) total_wins++;
if (open_regime == REGIME_BUCKET_T) {
trades_T++;
if (is_win) wins_T++;
} else if (open_regime == REGIME_BUCKET_R) {
trades_R++;
if (is_win) wins_R++;
} else if (open_regime == REGIME_BUCKET_V) {
trades_V++;
if (is_win) wins_V++;
}
// New stitched trade opens at this block's first action;
// its regime is the block's open-bar regime.
open_action = fa;
open_return = 0.0f;
open_regime = blk_open_regime;
}
if (nc == 0) {
open_return += pr;
} else {
// Block has internal transitions. Prefix portion closes
// the running stitched trade; the suffix opens a NEW
// stitched trade at the block's last transition bar
// (regime = `blk_last_open_regime`).
open_return += pr;
total_trades++;
if (open_return > 0.0f) total_wins++;
bool is_win = (open_return > 0.0f);
if (is_win) total_wins++;
if (open_regime == REGIME_BUCKET_T) {
trades_T++;
if (is_win) wins_T++;
} else if (open_regime == REGIME_BUCKET_R) {
trades_R++;
if (is_win) wins_R++;
} else if (open_regime == REGIME_BUCKET_V) {
trades_V++;
if (is_win) wins_V++;
}
open_action = la;
open_return = sr;
open_regime = blk_last_open_regime;
}
}
if (open_action >= -1) {
// Final trailing trade — attribute to its tracked open regime.
total_trades++;
if (open_return > 0.0f) total_wins++;
bool is_win = (open_return > 0.0f);
if (is_win) total_wins++;
if (open_regime == REGIME_BUCKET_T) {
trades_T++;
if (is_win) wins_T++;
} else if (open_regime == REGIME_BUCKET_R) {
trades_R++;
if (is_win) wins_R++;
} else if (open_regime == REGIME_BUCKET_V) {
trades_V++;
if (is_win) wins_V++;
}
}
int out_base = w * 13;
int out_base = w * 19;
metrics_out[out_base + 0] = s_sorted[0] - 1.0f;
metrics_out[out_base + 1] = exact_max_dd;
metrics_out[out_base + 2] = (mean / down_std) * annualization_factor;
@@ -290,6 +497,20 @@ extern "C" __global__ void compute_backtest_metrics(
metrics_out[out_base + 9] = s_buys[0];
metrics_out[out_base + 10] = s_sells[0];
metrics_out[out_base + 11] = s_holds[0];
// Per-regime WR + trade-count outputs (slots 13-18).
// WR is `wins / trades` per bucket; emits 0.0 when the bucket
// has zero trades (matches the aggregate WR's empty-trades
// convention at slot 3 above).
float trades_T_f = (float)trades_T;
float trades_R_f = (float)trades_R;
float trades_V_f = (float)trades_V;
metrics_out[out_base + 13] = (trades_T > 0) ? (float)wins_T / trades_T_f : 0.0f;
metrics_out[out_base + 14] = trades_T_f;
metrics_out[out_base + 15] = (trades_R > 0) ? (float)wins_R / trades_R_f : 0.0f;
metrics_out[out_base + 16] = trades_R_f;
metrics_out[out_base + 17] = (trades_V > 0) ? (float)wins_V / trades_V_f : 0.0f;
metrics_out[out_base + 18] = trades_V_f;
}
// Unique action count via bitmask OR-reduction
@@ -302,7 +523,7 @@ extern "C" __global__ void compute_backtest_metrics(
__syncthreads();
}
if (tid == 0) {
int out_base = w * 13;
int out_base = w * 19;
metrics_out[out_base + 12] = (float)__popc(s_action_mask[0]);
}
}
@@ -345,7 +566,7 @@ extern "C" __global__ void compute_backtest_metrics(
// Thread 0: extended metrics from sorted array
if (tid == 0) {
int out_base = w * 13;
int out_base = w * 19;
int var_idx = (int)(0.05f * (float)sort_len);
if (var_idx < 0) var_idx = 0;

68
crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs
View File

@@ -184,6 +184,21 @@ pub struct WindowMetrics {
pub baseline_hold_only_sharpe: f32,
pub baseline_momentum_sharpe: f32,
pub baseline_reversion_sharpe: f32,
// 2026-05-09 — per-regime WR instrumentation (audit follow-up).
// Trades are bucketed at trade-OPEN by the bar's ADX-norm value
// (feature[40], post-`RollingPercentileRank`) per
// `gpu_walk_forward.rs::classify_regime_from_features`:
// - Trending (T): ADX > 0.4
// - Ranging (R): ADX < 0.2
// - Volatile (V): otherwise
// `win_rate_*` is `wins / trades` per bucket (0.0 when zero trades,
// matching the aggregate `win_rate` empty-trade convention).
pub win_rate_trending: f32,
pub trade_count_trending: f32,
pub win_rate_ranging: f32,
pub trade_count_ranging: f32,
pub win_rate_volatile: f32,
pub trade_count_volatile: f32,
}
/// Configuration for the GPU backtest evaluator.
@@ -942,8 +957,12 @@ impl GpuBacktestEvaluator {
// zero-copy path per `feedback_no_htod_htoh_only_mapped_pinned.md`.
// SP15 Phase 1.1.b: stride dropped 14 → 13 (sharpe split out into
// the dedicated `sharpe_per_bar_kernel` consumed via `sharpe_buf`).
// 2026-05-09 — stride grew 13 → 19 with per-regime WR slots
// [13..19) (audit follow-up). Layout mirrored in
// `consume_metrics_after_event` and the kernel's `metrics_out`
// contract docstring.
let metrics_buf = unsafe {
MappedF32Buffer::new(n_windows * 13)
MappedF32Buffer::new(n_windows * 19)
.map_err(|e| MLError::ModelError(format!("metrics mapped-pinned alloc: {e}")))?
};
// SP15 Phase 1.1.b: per-window mean/std/sharpe output for the
@@ -2825,13 +2844,18 @@ impl GpuBacktestEvaluator {
///
/// Shared by all evaluation paths via the thin `evaluate_*` wrappers.
pub fn launch_metrics_and_record_event(&mut self) -> Result<(), MLError> {
// Shared memory: 5 reduction arrays × 256 threads × 4 bytes = 5 KB
// Shared memory: 11 reduction arrays × 256 threads × 4 bytes = 11 KB
// + 4096 floats for sort scratch / boundary data = 16 KB
// Total = 21 KB (well within the 48 KB L1/shmem limit).
// Total = 27 KB (well within the 48 KB L1/shmem limit).
// SP15 Phase 1.1.b shrunk this from 6 → 5 reduction arrays after the
// s_sq_sum reduction (sharpe-only) was removed; sharpe is now computed
// by the dedicated `sharpe_per_bar_kernel` launched per-window below.
let shmem_bytes = (256_u32 * 5 + 4096) * std::mem::size_of::<f32>() as u32;
// 2026-05-09 — per-regime WR adds 6 reduction arrays
// (s_trades_T/R/V + s_wins_T/R/V), bumping reduction-array count
// 5 → 11. The boundary buffer slot count also grows 7 → 9 to
// carry the open-bar regime + last-open-bar regime per chunk;
// those still live within the 4096-float `s_sorted` reservation.
let shmem_bytes = (256_u32 * 11 + 4096) * std::mem::size_of::<f32>() as u32;
let metrics_cfg = LaunchConfig {
grid_dim: (self.n_windows as u32, 1, 1),
block_dim: (256, 1, 1),
@@ -2845,6 +2869,11 @@ impl GpuBacktestEvaluator {
let num_actions_i32 = self.b0_size;
let order_actions_i32 = self.b1_size;
let urgency_actions_i32 = self.b2_size;
// 2026-05-09 — per-regime WR reads `features[bar_idx*feature_dim+40]`
// (ADX-norm) per bar. The features buffer was uploaded to GPU via
// mapped-pinned at construct time; pass its device pointer + the
// feature_dim scalar so the kernel can index per-window per-bar.
let feature_dim_i32 = self.feature_dim as i32;
// Pass mapped-pinned device pointer as kernel arg (f32* slot). The
// kernel writes through `dev_ptr`; the host reads `host_ptr` after
// sync. Same pattern as the IQN total_loss readback in `gpu_iqn_head.rs`.
@@ -2856,6 +2885,8 @@ impl GpuBacktestEvaluator {
.arg(&self.portfolio_buf.dev_ptr)
.arg(&self.window_lens_buf.dev_ptr)
.arg(&self.actions_history_buf)
.arg(&self.features_buf.dev_ptr)
.arg(&feature_dim_i32)
.arg(&metrics_dev_ptr)
.arg(&n_win_i32)
.arg(&max_len_i32)
@@ -3171,8 +3202,9 @@ impl GpuBacktestEvaluator {
// Parse flat metrics into per-window structs.
// Layout matches compute_backtest_metrics kernel output AFTER the
// SP15 Phase 1.1.b sharpe split (13 floats per window — sharpe at
// old slot 0 was removed, every metric below it shifted down by 1):
// SP15 Phase 1.1.b sharpe split + the 2026-05-09 per-regime WR
// expansion (19 floats per window — 13 base metrics + 6
// per-regime WR/trade-count slots):
// [0] total_pnl
// [1] max_drawdown
// [2] sortino
@@ -3186,13 +3218,19 @@ impl GpuBacktestEvaluator {
// [10] sell_count
// [11] hold_count
// [12] unique_actions
// [13] win_rate_trending
// [14] trade_count_trending
// [15] win_rate_ranging
// [16] trade_count_ranging
// [17] win_rate_volatile
// [18] trade_count_volatile
// Sharpe is sourced from `sharpe_host` (sharpe_per_bar_kernel) and
// annualised host-side: `sharpe = raw_sharpe * annualization_factor`,
// matching the in-kernel multiplication done in the pre-split fused
// kernel (was `(mean / std_val) * annualization_factor`).
let results: Vec<WindowMetrics> = (0..self.n_windows)
.map(|w| {
let base = w * 13;
let base = w * 19;
let sharpe_base = w * 3;
let raw_sharpe = sharpe_host[sharpe_base + 2];
let sharpe_annualised = raw_sharpe * annualization;
@@ -3225,6 +3263,12 @@ impl GpuBacktestEvaluator {
baseline_hold_only_sharpe: baseline_hold_only_raw * annualization,
baseline_momentum_sharpe: baseline_momentum_raw * annualization,
baseline_reversion_sharpe: baseline_reversion_raw * annualization,
win_rate_trending: metrics_host[base + 13],
trade_count_trending: metrics_host[base + 14],
win_rate_ranging: metrics_host[base + 15],
trade_count_ranging: metrics_host[base + 16],
win_rate_volatile: metrics_host[base + 17],
trade_count_volatile: metrics_host[base + 18],
}
})
.collect();
@@ -3287,12 +3331,22 @@ mod tests {
baseline_hold_only_sharpe: 0.0,
baseline_momentum_sharpe: 0.3,
baseline_reversion_sharpe: -0.1,
win_rate_trending: 0.6,
trade_count_trending: 14.0,
win_rate_ranging: 0.45,
trade_count_ranging: 18.0,
win_rate_volatile: 0.55,
trade_count_volatile: 10.0,
};
assert!(m.sharpe > 0.0);
assert!(m.win_rate > 0.5);
assert!(m.total_trades > 0.0);
assert!(m.calmar > 0.0);
assert!(m.omega_ratio > 1.0);
// Per-regime aggregates should sum to total trades.
let regime_total = m.trade_count_trending + m.trade_count_ranging + m.trade_count_volatile;
assert!((regime_total - m.total_trades).abs() < 1e-3,
"regime totals must sum to total_trades: {} vs {}", regime_total, m.total_trades);
}
#[test]

20
crates/ml/src/trainers/dqn/trainer/metrics.rs
View File

@@ -908,6 +908,26 @@ impl DQNTrainer {
window_bars,
);
// 2026-05-09 — per-regime val WR (audit follow-up). Trades
// bucketed at trade-OPEN by ADX-norm (feature[40]) per the
// structural thresholds in `backtest_metrics_kernel.cu`:
// T (Trending): ADX > 0.4 — directional regime
// R (Ranging): ADX < 0.2 — mean-reverting / chop
// V (Volatile): otherwise — transitional
// Surfaces whether the aggregate WR (~46% baseline) hides a
// regime-conditional edge. Top-level aggregator keys
// `val_regime_wr_T`, `val_regime_n_T`, etc. (single-`_` join).
tracing::info!(
"HEALTH_DIAG[{}]: val_regime [wr_T={:.4} n_T={:.0} wr_R={:.4} n_R={:.0} wr_V={:.4} n_V={:.0}]",
epoch,
m.win_rate_trending,
m.trade_count_trending,
m.win_rate_ranging,
m.trade_count_ranging,
m.win_rate_volatile,
m.trade_count_volatile,
);
// Plan C Task 5 — training-rollout counterpart to val_active_frac.
// SP8 (Fix 36, 2026-05-03): source migrated from host-side
// computation to ISV[TRAIN_ACTIVE_FRAC_INDEX] (Pearls A+D smoothed,

483
crates/ml/tests/regime_wr_oracle_tests.rs Normal file
View File

@@ -0,0 +1,483 @@
//! Per-regime WR GPU oracle tests (audit follow-up 2026-05-09).
//!
//! Validates that `compute_backtest_metrics` correctly bucketises
//! per-trade win rate by ADX-norm regime, matching a CPU oracle to
//! ε=1e-5. Trades are attributed at trade-OPEN by the bar's ADX-norm
//! value (feature[40]) per the structural thresholds:
//! - Trending (T): ADX > 0.4
//! - Ranging (R): ADX < 0.2
//! - Volatile (V): otherwise
//!
//! Run on a GPU host:
//!
//! SQLX_OFFLINE=true CUDA_COMPUTE_CAP=86 \
//! cargo test -p ml --test regime_wr_oracle_tests \
//! --features cuda -- --ignored --nocapture
//!
//! Per `feedback_no_htod_htoh_only_mapped_pinned`, all buffers are
//! `MappedF32Buffer` / `MappedI32Buffer` and reads happen through the
//! mapped host alias after stream sync.
#![cfg(feature = "cuda")]
use std::sync::Arc;
use cudarc::driver::{CudaContext, CudaFunction, CudaStream, LaunchConfig, PushKernelArg};
use ml::cuda_pipeline::mapped_pinned::{MappedF32Buffer, MappedI32Buffer};
/// Cubin handle for the per-window metrics kernel.
const METRICS_CUBIN: &[u8] =
include_bytes!(concat!(env!("OUT_DIR"), "/backtest_metrics_kernel.cubin"));
const KERNEL_NAME: &str = "compute_backtest_metrics";
/// 4-branch action encoding constants (must match production).
const NUM_ACTIONS: i32 = 108; // 4 dir × 3 mag × 3 ord × 3 urg
const ORDER_ACTIONS: i32 = 3;
const URGENCY_ACTIONS: i32 = 3;
/// Direction divisor: action / (mag × ord × urg) = action / 27 = exp_idx,
/// then dir_idx = exp_idx / 3. So dir-only multiplier = 27.
const DIR_DIV: i32 = 27;
/// Direction encoding (must match `state_layout.cuh`). Tests exercise
/// only Short/Long transitions — Hold and Flat are kernel-side
/// "no-exposure" sentinels that never become trade-open bars under the
/// boundary key (`signed_dir != bnd_cur_action` collapses both into
/// the 0 bucket which doesn't open a trade).
const DIR_SHORT: i32 = 0;
const DIR_LONG: i32 = 2;
/// Per-regime ADX threshold mirrors of the kernel constants.
const REGIME_TRENDING_THRESHOLD: f32 = 0.4;
const REGIME_RANGING_THRESHOLD: f32 = 0.2;
const REGIME_ADX_FEATURE_IDX: usize = 40;
/// Output stride per window (must match kernel `metrics_out` layout).
const METRICS_STRIDE: usize = 19;
fn make_test_stream() -> Arc<CudaStream> {
let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?");
ctx.default_stream()
}
fn load_metrics_kernel(stream: &Arc<CudaStream>) -> CudaFunction {
let module = stream
.context()
.load_cubin(METRICS_CUBIN.to_vec())
.expect("load backtest_metrics_kernel cubin");
module
.load_function(KERNEL_NAME)
.expect("load compute_backtest_metrics function")
}
/// Encode a packed action from a direction (short=0/long=2/flat=3).
/// Magnitude/order/urgency = 0 (canonical "Quarter Market Aggressive"
/// — irrelevant for trade-boundary purposes since the kernel collapses
/// magnitude into the direction key).
fn encode_dir(dir: i32) -> i32 {
dir * DIR_DIV
}
/// Classify a bar's regime from its ADX value (CPU oracle, mirrors the
/// kernel's `classify_bar_regime`).
fn classify_regime(adx: f32) -> usize {
if adx > REGIME_TRENDING_THRESHOLD {
0 // T
} else if adx < REGIME_RANGING_THRESHOLD {
1 // R
} else {
2 // V
}
}
/// Generic per-window oracle launcher. Builds a single-window batch
/// from `(actions, step_returns, adx_per_bar)`, launches the metrics
/// kernel, and returns the 19-float metrics output for that window.
///
/// The `adx_per_bar` slice produces a synthetic feature row per bar:
/// `feature[40] = adx`, all other slots zero. Feature dim = 42.
fn launch_metrics_for_window(
actions: &[i32],
step_returns: &[f32],
adx_per_bar: &[f32],
) -> Vec<f32> {
assert_eq!(actions.len(), step_returns.len());
assert_eq!(actions.len(), adx_per_bar.len());
let wlen = actions.len();
let max_len = wlen;
let n_windows = 1;
let feature_dim = 42usize;
// CUDA context must be initialized before MappedF32Buffer alloc
// (the cuMemHostAlloc DEVICEMAP path needs an active context).
// Stream creation binds the context to this thread; reuse the
// resulting stream for the kernel launch below.
let stream = make_test_stream();
let kernel = load_metrics_kernel(&stream);
// ── Buffers ──────────────────────────────────────────────────────
let step_returns_buf = unsafe { MappedF32Buffer::new(wlen) }
.expect("alloc step_returns");
step_returns_buf.write_from_slice(step_returns);
// 8 floats per window for portfolio_state (unused by the metrics
// kernel beyond `[w * 8]` indexing — kernel reads but doesn't act
// on portfolio fields).
let portfolio_buf = unsafe { MappedF32Buffer::new(8) }
.expect("alloc portfolio");
portfolio_buf.write_from_slice(&[0.0_f32; 8]);
let window_lens_buf = unsafe { MappedI32Buffer::new(1) }
.expect("alloc window_lens");
window_lens_buf.write_from_slice(&[wlen as i32]);
let actions_buf = unsafe { MappedI32Buffer::new(wlen) }
.expect("alloc actions_history");
actions_buf.write_from_slice(actions);
// Features [n_windows × max_len × feature_dim] — slot 40 = ADX.
let mut flat_features = vec![0.0_f32; n_windows * max_len * feature_dim];
for (i, &adx) in adx_per_bar.iter().enumerate() {
flat_features[i * feature_dim + REGIME_ADX_FEATURE_IDX] = adx;
}
let features_buf = unsafe { MappedF32Buffer::new(flat_features.len()) }
.expect("alloc features");
features_buf.write_from_slice(&flat_features);
let metrics_buf = unsafe { MappedF32Buffer::new(n_windows * METRICS_STRIDE) }
.expect("alloc metrics");
metrics_buf.write_from_slice(&vec![f32::NAN; n_windows * METRICS_STRIDE]);
// ── Launch ────────────────────────────────────────────────────────
// Shared memory: 11 reduction arrays × 256 threads × 4 bytes
// + 4096 floats for sort scratch / boundary data
let shmem_bytes: u32 = (256 * 11 + 4096) * std::mem::size_of::<f32>() as u32;
let cfg = LaunchConfig {
grid_dim: (n_windows as u32, 1, 1),
block_dim: (256, 1, 1),
shared_mem_bytes: shmem_bytes,
};
let n_win_i32 = n_windows as i32;
let max_len_i32 = max_len as i32;
let annualisation: f32 = 1.0; // not exercised by these tests
let num_actions = NUM_ACTIONS;
let order_actions = ORDER_ACTIONS;
let urgency_actions = URGENCY_ACTIONS;
let feature_dim_i32 = feature_dim as i32;
unsafe {
stream
.launch_builder(&kernel)
.arg(&step_returns_buf.dev_ptr)
.arg(&portfolio_buf.dev_ptr)
.arg(&window_lens_buf.dev_ptr)
.arg(&actions_buf.dev_ptr)
.arg(&features_buf.dev_ptr)
.arg(&feature_dim_i32)
.arg(&metrics_buf.dev_ptr)
.arg(&n_win_i32)
.arg(&max_len_i32)
.arg(&annualisation)
.arg(&num_actions)
.arg(&order_actions)
.arg(&urgency_actions)
.launch(cfg)
.expect("launch compute_backtest_metrics");
}
stream.synchronize().expect("sync after metrics kernel");
metrics_buf.read_all()
}
/// CPU oracle for per-regime WR. Walks the `(actions, step_returns,
/// adx)` series, replicates the kernel's trade-boundary key (collapses
/// Hold/Flat into "no exposure"), and computes per-regime
/// `(wins_T/R/V, trades_T/R/V)` attributing each trade to the regime
/// at its OPEN bar.
fn cpu_oracle_per_regime(
actions: &[i32],
step_returns: &[f32],
adx: &[f32],
) -> [u32; 6] {
let mut wins = [0u32; 3];
let mut trades = [0u32; 3];
let mut cur_signed_dir: i32 = 0;
let mut cur_open_regime: usize = 0;
let mut cur_return: f32 = 0.0;
let mut active = false;
for (i, &act) in actions.iter().enumerate() {
if act < 0 { continue; }
let exp_idx = act / (ORDER_ACTIONS * URGENCY_ACTIONS);
let raw_dir = exp_idx / 3;
let signed_dir = match raw_dir {
d if d == DIR_SHORT => -1,
d if d == DIR_LONG => 1,
_ => 0, // Hold(1) or Flat(3)
};
if !active {
cur_signed_dir = signed_dir;
cur_open_regime = classify_regime(adx[i]);
cur_return = step_returns[i];
active = true;
continue;
}
if signed_dir != cur_signed_dir {
// Close previous trade, attribute to its open regime.
trades[cur_open_regime] += 1;
if cur_return > 0.0 {
wins[cur_open_regime] += 1;
}
// Open new trade at this bar; its regime = bar's regime.
cur_signed_dir = signed_dir;
cur_open_regime = classify_regime(adx[i]);
cur_return = step_returns[i];
} else {
cur_return += step_returns[i];
}
}
// Final trailing trade.
if active {
trades[cur_open_regime] += 1;
if cur_return > 0.0 {
wins[cur_open_regime] += 1;
}
}
[
wins[0], trades[0],
wins[1], trades[1],
wins[2], trades[2],
]
}
/// CPU oracle test — sanity check for the oracle itself. Mixes 3
/// trades each in T/R/V buckets with known outcomes and verifies the
/// classifier + accounting.
#[test]
fn regime_wr_cpu_oracle_sanity() {
// 9 bars, 3 distinct trades per regime (transitions between
// long/short directions force a new trade each time).
// Bar layout (action, step_return, adx):
// bar 0: long +0.10 adx=0.5 (T)
// bar 1: short -0.05 adx=0.5 (T) ← closes trade 0 (T win=1)
// bar 2: long +0.20 adx=0.5 (T) ← closes trade 1 (T win=0, ret -0.05)
// bar 3: short +0.10 adx=0.1 (R) ← closes trade 2 (T win=1, ret +0.20)
// bar 4: long -0.10 adx=0.1 (R) ← closes trade 3 (R win=1, ret +0.10)
// bar 5: short -0.20 adx=0.1 (R) ← closes trade 4 (R win=0, ret -0.10)
// bar 6: long +0.30 adx=0.3 (V) ← closes trade 5 (R win=0, ret -0.20)
// bar 7: short +0.10 adx=0.3 (V) ← closes trade 6 (V win=1, ret +0.30)
// bar 8: long +0.05 adx=0.3 (V) ← closes trade 7 (V win=1, ret +0.10)
// final trailing trade 8: V at bar 8, ret +0.05 → win=1
let actions: Vec<i32> = vec![
encode_dir(DIR_LONG), // 0 — opens trade 0 (T)
encode_dir(DIR_SHORT), // 1 — closes 0 (+0.10), opens 1 (T)
encode_dir(DIR_LONG), // 2 — closes 1 (-0.05), opens 2 (T)
encode_dir(DIR_SHORT), // 3 — closes 2 (+0.20), opens 3 (R)
encode_dir(DIR_LONG), // 4 — closes 3 (+0.10), opens 4 (R)
encode_dir(DIR_SHORT), // 5 — closes 4 (-0.10), opens 5 (R)
encode_dir(DIR_LONG), // 6 — closes 5 (-0.20), opens 6 (V)
encode_dir(DIR_SHORT), // 7 — closes 6 (+0.30), opens 7 (V)
encode_dir(DIR_LONG), // 8 — closes 7 (+0.10), opens 8 (V)
];
let step_ret: Vec<f32> = vec![
0.10, -0.05, 0.20, 0.10, -0.10, -0.20, 0.30, 0.10, 0.05,
];
let adx: Vec<f32> = vec![0.5, 0.5, 0.5, 0.1, 0.1, 0.1, 0.3, 0.3, 0.3];
let result = cpu_oracle_per_regime(&actions, &step_ret, &adx);
// T: trades 0,1,2 → outcomes +0.10, -0.05, +0.20 → 2 wins / 3 trades
assert_eq!(result[1], 3, "T trades count");
assert_eq!(result[0], 2, "T wins count");
// R: trades 3,4,5 → outcomes +0.10, -0.10, -0.20 → 1 win / 3 trades
assert_eq!(result[3], 3, "R trades count");
assert_eq!(result[2], 1, "R wins count");
// V: trades 6,7 (closed) + trade 8 (trailing, +0.05) → outcomes
// +0.30, +0.10, +0.05 → 3 wins / 3 trades
assert_eq!(result[5], 3, "V trades count");
assert_eq!(result[4], 3, "V wins count");
}
/// GPU oracle test — bit-exact match of the kernel's per-regime WR
/// against the CPU oracle on the same 9-bar synthetic batch from the
/// CPU sanity test.
#[test]
#[ignore = "requires GPU"]
fn regime_wr_gpu_oracle_matches_cpu() {
let actions: Vec<i32> = vec![
encode_dir(DIR_LONG),
encode_dir(DIR_SHORT),
encode_dir(DIR_LONG),
encode_dir(DIR_SHORT),
encode_dir(DIR_LONG),
encode_dir(DIR_SHORT),
encode_dir(DIR_LONG),
encode_dir(DIR_SHORT),
encode_dir(DIR_LONG),
];
let step_ret: Vec<f32> = vec![
0.10, -0.05, 0.20, 0.10, -0.10, -0.20, 0.30, 0.10, 0.05,
];
let adx: Vec<f32> = vec![0.5, 0.5, 0.5, 0.1, 0.1, 0.1, 0.3, 0.3, 0.3];
let cpu = cpu_oracle_per_regime(&actions, &step_ret, &adx);
let metrics = launch_metrics_for_window(&actions, &step_ret, &adx);
let wr_t = metrics[13];
let n_t = metrics[14];
let wr_r = metrics[15];
let n_r = metrics[16];
let wr_v = metrics[17];
let n_v = metrics[18];
let exp_wr_t = cpu[0] as f32 / cpu[1].max(1) as f32;
let exp_wr_r = cpu[2] as f32 / cpu[3].max(1) as f32;
let exp_wr_v = cpu[4] as f32 / cpu[5].max(1) as f32;
let eps = 1e-5_f32;
assert!((wr_t - exp_wr_t).abs() < eps,
"T WR: kernel={wr_t:.6} cpu={exp_wr_t:.6}");
assert!((wr_r - exp_wr_r).abs() < eps,
"R WR: kernel={wr_r:.6} cpu={exp_wr_r:.6}");
assert!((wr_v - exp_wr_v).abs() < eps,
"V WR: kernel={wr_v:.6} cpu={exp_wr_v:.6}");
assert!((n_t - cpu[1] as f32).abs() < eps,
"T trades: kernel={n_t} cpu={}", cpu[1]);
assert!((n_r - cpu[3] as f32).abs() < eps,
"R trades: kernel={n_r} cpu={}", cpu[3]);
assert!((n_v - cpu[5] as f32).abs() < eps,
"V trades: kernel={n_v} cpu={}", cpu[5]);
// Per-regime trade counts must sum to the aggregate trade count.
let total_trades = metrics[4];
let regime_total = n_t + n_r + n_v;
assert!((regime_total - total_trades).abs() < eps,
"regime total {regime_total} != aggregate trade count {total_trades}");
}
/// GPU oracle test — known-distribution stratified batch.
///
/// 30% trending / 40% ranging / 30% volatile by bar count, with
/// per-regime trade outcomes that fix exactly:
/// T: 2/3 win rate (2 wins, 1 loss)
/// R: 1/4 win rate (1 win, 3 losses)
/// V: 1/3 win rate (1 win, 2 losses)
///
/// Verifies the kernel honours the regime-conditional skew. Uses a
/// 30-bar window — large enough to exercise the bitonic-sort path
/// without being so large that boundary stitching dominates the
/// regime-attribution test.
#[test]
#[ignore = "requires GPU"]
fn regime_wr_known_distribution_30_40_30() {
// 30 bars total. Trade transitions every 3 bars (10 trades).
// Bars 0-8: T regime (3 trades). Bars 9-20: R regime (4 trades).
// Bars 21-29: V regime (3 trades).
let mut actions: Vec<i32> = Vec::with_capacity(30);
let mut step_ret: Vec<f32> = Vec::with_capacity(30);
let mut adx: Vec<f32> = Vec::with_capacity(30);
// Helper: push a trade chunk with `dir` direction and per-bar
// returns. The trade closes when the next chunk has a different
// direction.
let push_trade = |actions: &mut Vec<i32>,
step_ret: &mut Vec<f32>,
adx: &mut Vec<f32>,
dir: i32,
returns: &[f32],
adx_value: f32| {
for &r in returns {
actions.push(encode_dir(dir));
step_ret.push(r);
adx.push(adx_value);
}
};
// T regime, ADX=0.7 (>0.4 → Trending). 3 trades.
// Trade T0: long, returns sum to +0.15 (win)
// Trade T1: short, returns sum to +0.10 (win)
// Trade T2: long, returns sum to -0.05 (loss)
push_trade(&mut actions, &mut step_ret, &mut adx, DIR_LONG, &[0.05, 0.05, 0.05], 0.7);
push_trade(&mut actions, &mut step_ret, &mut adx, DIR_SHORT, &[0.05, 0.04, 0.01], 0.7);
push_trade(&mut actions, &mut step_ret, &mut adx, DIR_LONG, &[-0.02, -0.02, -0.01], 0.7);
// R regime, ADX=0.1 (<0.2 → Ranging). 4 trades.
// Trade R0: short, returns sum to +0.05 (win)
// Trade R1: long, returns sum to -0.05 (loss)
// Trade R2: short, returns sum to -0.10 (loss)
// Trade R3: long, returns sum to -0.05 (loss)
push_trade(&mut actions, &mut step_ret, &mut adx, DIR_SHORT, &[0.02, 0.02, 0.01], 0.1);
push_trade(&mut actions, &mut step_ret, &mut adx, DIR_LONG, &[-0.02, -0.02, -0.01], 0.1);
push_trade(&mut actions, &mut step_ret, &mut adx, DIR_SHORT, &[-0.04, -0.03, -0.03], 0.1);
push_trade(&mut actions, &mut step_ret, &mut adx, DIR_LONG, &[-0.02, -0.02, -0.01], 0.1);
// V regime, ADX=0.3 (between → Volatile). 3 trades.
// Trade V0: short, +0.15 (win)
// Trade V1: long, -0.05 (loss)
// Trade V2: short, -0.10 (loss)
push_trade(&mut actions, &mut step_ret, &mut adx, DIR_SHORT, &[0.05, 0.05, 0.05], 0.3);
push_trade(&mut actions, &mut step_ret, &mut adx, DIR_LONG, &[-0.02, -0.02, -0.01], 0.3);
push_trade(&mut actions, &mut step_ret, &mut adx, DIR_SHORT, &[-0.04, -0.03, -0.03], 0.3);
assert_eq!(actions.len(), 30, "30-bar window");
// Sanity-check the CPU oracle matches the design.
let cpu = cpu_oracle_per_regime(&actions, &step_ret, &adx);
assert_eq!(cpu[1], 3, "T trades = 3 (CPU oracle)");
assert_eq!(cpu[0], 2, "T wins = 2 (CPU oracle)");
assert_eq!(cpu[3], 4, "R trades = 4 (CPU oracle)");
assert_eq!(cpu[2], 1, "R wins = 1 (CPU oracle)");
assert_eq!(cpu[5], 3, "V trades = 3 (CPU oracle)");
assert_eq!(cpu[4], 1, "V wins = 1 (CPU oracle)");
let metrics = launch_metrics_for_window(&actions, &step_ret, &adx);
let eps = 1e-5_f32;
let wr_t = metrics[13];
let n_t = metrics[14];
let wr_r = metrics[15];
let n_r = metrics[16];
let wr_v = metrics[17];
let n_v = metrics[18];
let exp_wr_t = 2.0 / 3.0;
let exp_wr_r = 1.0 / 4.0;
let exp_wr_v = 1.0 / 3.0;
assert!((wr_t - exp_wr_t).abs() < eps, "T WR: kernel={wr_t} expected={exp_wr_t}");
assert!((wr_r - exp_wr_r).abs() < eps, "R WR: kernel={wr_r} expected={exp_wr_r}");
assert!((wr_v - exp_wr_v).abs() < eps, "V WR: kernel={wr_v} expected={exp_wr_v}");
assert!((n_t - 3.0).abs() < eps, "T trades: kernel={n_t} expected=3");
assert!((n_r - 4.0).abs() < eps, "R trades: kernel={n_r} expected=4");
assert!((n_v - 3.0).abs() < eps, "V trades: kernel={n_v} expected=3");
}
/// GPU oracle test — empty-trade resilience. When all bars have the
/// same direction (single trade), only one regime bucket is non-zero
/// and the others must report exactly 0 trades / 0 WR (matches the
/// aggregate WR's empty-trades convention).
#[test]
#[ignore = "requires GPU"]
fn regime_wr_single_regime_only_one_bucket_populated() {
// 5-bar long-only trade in the trending regime.
let actions: Vec<i32> = (0..5).map(|_| encode_dir(DIR_LONG)).collect();
let step_ret: Vec<f32> = vec![0.01, 0.01, 0.01, 0.01, 0.01];
let adx: Vec<f32> = vec![0.6; 5]; // ADX>0.4 → all T
let metrics = launch_metrics_for_window(&actions, &step_ret, &adx);
let eps = 1e-5_f32;
// 1 trade, T bucket only.
assert!((metrics[14] - 1.0).abs() < eps, "T trades = 1");
assert!((metrics[13] - 1.0).abs() < eps, "T WR = 1.0 (winner)");
assert!(metrics[16].abs() < eps, "R trades = 0");
assert!(metrics[15].abs() < eps, "R WR = 0 (no trades)");
assert!(metrics[18].abs() < eps, "V trades = 0");
assert!(metrics[17].abs() < eps, "V WR = 0 (no trades)");
}

75
docs/dqn-wire-up-audit.md
View File

@@ -2,6 +2,81 @@
**Status:** Populated during Plan 1 Task 6 (A.5 orphan audit). Updated on every commit per Invariant 7.
## 2026-05-09 — Per-regime val WR instrumentation (audit follow-up)
Adds 6 output slots [13..19) to `compute_backtest_metrics_kernel` —
per-bucket `(win_rate, trade_count)` triples for the {Trending,
Ranging, Volatile} regime split — so the val backtest surfaces
whether the long-running ~46% aggregate win rate hides a
regime-conditional edge. Trades are bucketed at trade-OPEN by the
bar's ADX-norm value (`feature[40]`, post-`RollingPercentileRank`)
per the structural thresholds: T:ADX>0.4, R:ADX<0.2, V:otherwise
(mirrors `gpu_walk_forward.rs::classify_regime_from_features` —
identical feature index, slightly tighter threshold split per the
2026-05-09 audit findings call-out for Trending=0.4 / Ranging=0.2).
Empirically-grounded structural defaults per
`feedback_isv_for_adaptive_bounds`; observability-only emission
(does not feed any controller), so the thresholds remain kernel
constants rather than ISV slots.
Atomic commit per `feedback_no_partial_refactor` /
`feedback_wire_everything_up`:
- `crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu` —
per-regime trade tracking: 6 new reduction arrays
(`s_trades_T/R/V`, `s_wins_T/R/V`), 2 new boundary-buffer slots
(`bnd_open_regime`, `bnd_last_open_regime`) carrying open-bar
regime through the block-stitch loop. Output stride 13 → 19.
Block tree-reduce only, no atomicAdd per `feedback_no_atomicadd`.
- `crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs` —
`WindowMetrics` struct grows 6 fields
(`win_rate_trending`/`trade_count_trending` × 3 regimes);
`metrics_buf` size 13 → 19; launcher passes
`features_buf.dev_ptr` + `feature_dim_i32` after the actions
arg; `consume_metrics_after_event` populates the new fields;
shmem 21 KB → 27 KB (5 → 11 reduction tiles).
- `crates/ml/src/trainers/dqn/trainer/metrics.rs` — new
`HEALTH_DIAG[N]: val_regime [wr_T={:.4} n_T={:.0} wr_R={:.4}
n_R={:.0} wr_V={:.4} n_V={:.0}]` line emitted alongside the
existing `val [...]` block in `consume_validation_loss`.
Aggregator parser keys: `val_regime_wr_T`, `val_regime_n_T`, etc.
- `crates/ml/tests/regime_wr_oracle_tests.rs` (NEW) — 1 CPU
sanity test + 3 GPU oracle tests:
* `regime_wr_cpu_oracle_sanity` — verifies the CPU oracle
classifier + accounting math
* `regime_wr_gpu_oracle_matches_cpu` — bit-exact match (ε=1e-5)
of kernel vs. CPU oracle on a 9-bar synthetic batch
* `regime_wr_known_distribution_30_40_30` — 30%/40%/30%
stratified batch with known per-regime WR (2/3, 1/4, 1/3)
* `regime_wr_single_regime_only_one_bucket_populated` —
empty-trade resilience for unused buckets
Wire-up audit: kernel definition + cubin load (already in build.rs
manifest, semantic stride change only) + launcher arg list updated +
3 ops struct field reads (`features_buf.dev_ptr`, `feature_dim`,
new metrics layout) + WindowMetrics struct grew + HEALTH_DIAG emit
in metrics.rs + GPU oracle test. No orphans.
Lookahead-audit closure (`docs/lookahead-bias-audit-2026-04-28.md`):
this is the **observability** half of the per-regime conditional
edge investigation flagged at the bottom of §6 (CVaR / Q-variance
feedback). The instrumentation does NOT modify training or
validation behaviour — it only surfaces a per-regime breakdown of
the existing aggregate metric. The actual policy fix (if a regime
turns out to dominate the edge) is downstream of this telemetry.
Pearls applied: `feedback_no_atomicadd` (block tree-reduce per-regime
counters), `feedback_no_htod_htoh_only_mapped_pinned` (test buffers
are `MappedF32Buffer` / `MappedI32Buffer`), `feedback_isv_for_adaptive_bounds`
(thresholds documented as conceptually-adaptive structural defaults,
liftable to ISV when a controller consumes the per-regime signal),
`feedback_wire_everything_up` (kernel + launcher + diag emit + tests
in same commit).
Validation: `cargo check --workspace` clean (1m 04s); 17/17
gpu_backtest_evaluator unit tests pass; 1/1 CPU sanity + 3/3 GPU
oracle tests pass on local RTX 3050 Ti (1.89s wallclock);
`bash scripts/audit_sp18_consumers.sh --check` exit 0.
## 2026-05-09 — SP18 v2 Phase 3 Tasks 3.1-3.5: adaptive HOLD_REWARD_POS/NEG_CAP producer
D-leg adaptive cap producer kernel + launcher + per-epoch wiring + GPU