diff --git a/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu index daa80daba..f0d7da74d 100644 --- a/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu +++ b/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu @@ -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; diff --git a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs index 29ae78d5a..2c39acd20 100644 --- a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs +++ b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs @@ -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::() 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::() 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 = (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] diff --git a/crates/ml/src/trainers/dqn/trainer/metrics.rs b/crates/ml/src/trainers/dqn/trainer/metrics.rs index 4df1569bc..723a7257b 100644 --- a/crates/ml/src/trainers/dqn/trainer/metrics.rs +++ b/crates/ml/src/trainers/dqn/trainer/metrics.rs @@ -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, diff --git a/crates/ml/tests/regime_wr_oracle_tests.rs b/crates/ml/tests/regime_wr_oracle_tests.rs new file mode 100644 index 000000000..8c68dc9a9 --- /dev/null +++ b/crates/ml/tests/regime_wr_oracle_tests.rs @@ -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 { + let ctx = CudaContext::new(0).expect("CUDA context — is a GPU available?"); + ctx.default_stream() +} + +fn load_metrics_kernel(stream: &Arc) -> 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 { + 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::() 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 = 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 = vec![ + 0.10, -0.05, 0.20, 0.10, -0.10, -0.20, 0.30, 0.10, 0.05, + ]; + let adx: Vec = 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 = 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 = vec![ + 0.10, -0.05, 0.20, 0.10, -0.10, -0.20, 0.30, 0.10, 0.05, + ]; + let adx: Vec = 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 = Vec::with_capacity(30); + let mut step_ret: Vec = Vec::with_capacity(30); + let mut adx: Vec = 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, + step_ret: &mut Vec, + adx: &mut Vec, + 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 = (0..5).map(|_| encode_dir(DIR_LONG)).collect(); + let step_ret: Vec = vec![0.01, 0.01, 0.01, 0.01, 0.01]; + let adx: Vec = 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)"); +} diff --git a/docs/dqn-wire-up-audit.md b/docs/dqn-wire-up-audit.md index f7f9e1ea0..c53feb9f7 100644 --- a/docs/dqn-wire-up-audit.md +++ b/docs/dqn-wire-up-audit.md @@ -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