From 3bd339b5a3ce5dec8c646743b4652ea9d77b7331 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Wed, 25 Mar 2026 19:43:26 +0100 Subject: [PATCH] =?UTF-8?q?fix:=20backtest=20evaluator=20=E2=80=94=20corre?= =?UTF-8?q?ct=20annualization,=20window=20sizing,=20objective=20calibratio?= =?UTF-8?q?n?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Root cause: backtest windows of 300K bars produced ±billions% returns via multiplicative compounding, and sqrt(252) annualization was wrong for 1-minute bars. Fixes: - Window size capped to 10K bars (~25 trading days), evenly distributed across the full validation set (was clustered in first 6%) - Annualization: configurable bars_per_day field in GpuBacktestConfig (default 390.0 for 1-min), produces sqrt(98280) ≈ 313.5 - tanh normalization recalibrated: Sharpe/5, Sortino/8 (was /2, /3) - CVaR threshold scaled to per-bar: 0.003 with slope 1400 (was 0.05/200) - VaR/CVaR strided sampling covers full window (was first 4096 only) - financials.rs + ab_testing.rs: sqrt(252) → sqrt(98280) for consistency Co-Authored-By: Claude Opus 4.6 (1M context) --- crates/ml-ensemble/src/ab_testing.rs | 10 +- crates/ml/src/hyperopt/adapters/dqn.rs | 75 ++- crates/ml/src/trainers/dqn/financials.rs | 11 +- .../2026-03-25-boundary-trade-counting.md | 494 ++++++++++++++++++ ...26-03-25-boundary-trade-counting-design.md | 265 ++++++++++ 5 files changed, 830 insertions(+), 25 deletions(-) create mode 100644 docs/superpowers/plans/2026-03-25-boundary-trade-counting.md create mode 100644 docs/superpowers/specs/2026-03-25-boundary-trade-counting-design.md diff --git a/crates/ml-ensemble/src/ab_testing.rs b/crates/ml-ensemble/src/ab_testing.rs index 5be69e493..bd84529ee 100644 --- a/crates/ml-ensemble/src/ab_testing.rs +++ b/crates/ml-ensemble/src/ab_testing.rs @@ -113,8 +113,9 @@ impl GroupMetrics { if std_dev < 1e-10 { 0.0 } else { - // Annualize: assume 252 trading days, multiply by sqrt(252) - mean_return / std_dev * 15.874 // sqrt(252) ≈ 15.874 + // Annualize: 1-minute bars, 390 bars/day × 252 days/year + const BARS_PER_YEAR: f64 = 390.0 * 252.0; + mean_return / std_dev * BARS_PER_YEAR.sqrt() } } } @@ -823,9 +824,10 @@ mod tests { // Should be positive with these returns assert!(sharpe > 0.0, "Sharpe ratio should be positive"); - // Rough check: mean ~0.0085, std ~0.013, annualized Sharpe ~10 + // Rough check: mean ~0.008, std ~0.013, annualized Sharpe ~195 + // (1-minute bar annualization: sqrt(390 * 252) ≈ 313.5) assert!( - sharpe > 5.0 && sharpe < 15.0, + sharpe > 100.0 && sharpe < 300.0, "Sharpe ratio {} out of expected range", sharpe ); diff --git a/crates/ml/src/hyperopt/adapters/dqn.rs b/crates/ml/src/hyperopt/adapters/dqn.rs index 4a1ec620d..53419111e 100644 --- a/crates/ml/src/hyperopt/adapters/dqn.rs +++ b/crates/ml/src/hyperopt/adapters/dqn.rs @@ -3368,20 +3368,42 @@ impl HyperparameterOptimizable for DQNTrainer { // used: EVAL_CHUNK_SIZE bars per GPU forward pass → sequential // trade simulation on CPU → portfolio sync between chunks. const EVAL_CHUNK_SIZE: usize = 1024; - // Sliding walk-forward: window_size = total/3 (same statistical power), - // 50% overlap stride produces ~5 windows for robust aggregation. - const WINDOW_OVERLAP: f64 = 0.5; + // Cap window size to prevent extreme multiplicative compounding. + // 10K bars ≈ 25 trading days of 1-minute data — enough for stable + // Sharpe estimates without astronomical cumulative returns. + // (Previous: total/3 ≈ 300K bars → (1.00002)^300K = billions %) + const MAX_WINDOW_BARS: usize = 10_000; + // Target window count for walk-forward cross-validation. + // More windows = better statistical estimate but slower evaluation. + const TARGET_WINDOWS: usize = 10; let device = internal_trainer.device().clone(); let total_bars = val_close_prices.len(); - let window_size = total_bars / 3; // Same per-window size as before - #[allow(clippy::cast_possible_truncation, clippy::cast_sign_loss)] - let stride = ((window_size as f64) * (1.0 - WINDOW_OVERLAP)).max(1.0) as usize; - let window_count = if window_size == 0 || stride == 0 { - 0 + let window_size = (total_bars / 3).min(MAX_WINDOW_BARS); + + // Distribute windows EVENLY across the full validation set. + // Old approach: 50% overlap stride → clustered at the start, + // only 6% of data evaluated when window_count was capped. + // New approach: stride = (total - window) / (N-1) so windows + // span the entire time series for representative evaluation. + let (stride, window_count) = if window_size == 0 || total_bars < window_size { + (1, 0) } else { - (total_bars.saturating_sub(window_size)) / stride + 1 + let available = total_bars - window_size; + // How many windows can we fit with 50% overlap? + let overlap_stride = (window_size / 2).max(1); + let overlap_windows = available / overlap_stride + 1; + + if overlap_windows <= TARGET_WINDOWS { + // Few enough — use 50% overlap for maximum coverage + (overlap_stride, overlap_windows) + } else { + // Too many — redistribute evenly across the full dataset + let n = TARGET_WINDOWS; + let even_stride = if n <= 1 { available } else { available / (n - 1) }; + (even_stride.max(1), n) + } }; // Need at least 1 bar per window to produce meaningful metrics @@ -3707,17 +3729,34 @@ impl HyperparameterOptimizable for DQNTrainer { trade_penalty } else { // Component 1: Multi-objective composite score (60% weight) - // tanh normalization: maps each ratio to [-1,1] to prevent - // scale dominance (Calmar with tiny drawdowns was 1000x+). + // + // tanh normalization: maps each ratio to [-1,1] to prevent scale + // dominance. Divisors calibrated for 1-minute bar annualization + // (sqrt(98280) ≈ 313.5), which produces Sharpe ~2-10 for good models. + // + // Divisor calibration (tanh has useful gradient in [-2, 2]): + // Sharpe /5: range [-5,15] → tanh [-1.0, 3.0] → [-0.76, 0.995] + // Sortino /8: range [-5,25] → tanh [-0.6, 3.1] → [-0.54, 0.996] + // Calmar /5: range [0, 15] → tanh [0, 3.0] → [0.0, 0.995] + // Omega /2: range [0, 5] → tanh [0, 2.5] → [0.0, 0.987] let composite_score = - 0.40 * (backtest.sharpe_ratio / 2.0).tanh() + - 0.25 * (backtest.sortino_ratio / 3.0).tanh() + + 0.40 * (backtest.sharpe_ratio / 5.0).tanh() + + 0.25 * (backtest.sortino_ratio / 8.0).tanh() + 0.20 * (backtest.calmar_ratio / 5.0).tanh() + 0.15 * (backtest.omega_ratio / 2.0).tanh(); - // Tail risk penalty: smooth ramp instead of hard cliff at -5% CVaR. - // Gives PSO gradient signal near the threshold instead of a binary 0/10 jump. - let cvar_penalty = ((-backtest.cvar_95 - 0.05).max(0.0) * 200.0).min(10.0); + // Tail risk penalty: smooth ramp based on per-bar CVaR. + // + // CVaR is the mean of per-bar returns below the 5th percentile. + // For 1-minute bars, typical CVaR is -0.001 to -0.003 (-0.1% to -0.3%/bar). + // Threshold scaled from daily (0.05) to per-bar: 0.05 / sqrt(390) ≈ 0.003. + // + // Penalty ramp: + // CVaR = -0.002 → no penalty (normal) + // CVaR = -0.005 → penalty ≈ 2.8 (elevated) + // CVaR = -0.010 → penalty ≈ 9.8 (max, dangerous) + let cvar_threshold = 0.003; + let cvar_penalty = ((-backtest.cvar_95 - cvar_threshold).max(0.0) * 1400.0).min(10.0); // Component 2: HFT activity score (25% weight) let hft_activity = calculate_hft_activity_score_wave10(buy_pct, sell_pct, hold_pct); @@ -4697,8 +4736,8 @@ mod tests { total_trades: 200, sortino_ratio: 4.0, calmar_ratio: 2.4, - var_95: -0.02, - cvar_95: -0.03, + var_95: -0.002, // 5th percentile per-bar return (1-min data) + cvar_95: -0.002, // Mean of tail returns below VaR (1-min data) beta: 0.1, alpha: 0.05, information_ratio: 1.5, diff --git a/crates/ml/src/trainers/dqn/financials.rs b/crates/ml/src/trainers/dqn/financials.rs index 26c1026db..766c6a967 100644 --- a/crates/ml/src/trainers/dqn/financials.rs +++ b/crates/ml/src/trainers/dqn/financials.rs @@ -68,7 +68,12 @@ pub(crate) fn compute_epoch_financials( 0.0 }; - // Sharpe ratio from step_returns (per-bar, annualized x sqrt(252)) + // Sharpe ratio from step_returns (per-bar, annualized) + // Data is 1-minute bars: 390 bars/day × 252 days/year = 98280 bars/year. + // sqrt(98280) ≈ 313.5 — consistent with GPU backtest evaluator annualization. + const BARS_PER_YEAR: f64 = 390.0 * 252.0; + let annualization = BARS_PER_YEAR.sqrt(); + let returns = &trade_stats.step_returns; let n = returns.len(); let (sharpe, sortino) = if n > 1 { @@ -77,7 +82,7 @@ pub(crate) fn compute_epoch_financials( let std = variance.sqrt(); let s = if std > 1e-10 { - (mean / std) * (252.0_f64).sqrt() + (mean / std) * annualization } else { 0.0 }; @@ -89,7 +94,7 @@ pub(crate) fn compute_epoch_financials( / downside.len() as f64; let down_std = down_var.sqrt(); if down_std > 1e-10 { - (mean / down_std) * (252.0_f64).sqrt() + (mean / down_std) * annualization } else { 0.0 } diff --git a/docs/superpowers/plans/2026-03-25-boundary-trade-counting.md b/docs/superpowers/plans/2026-03-25-boundary-trade-counting.md new file mode 100644 index 000000000..0f83d8ca1 --- /dev/null +++ b/docs/superpowers/plans/2026-03-25-boundary-trade-counting.md @@ -0,0 +1,494 @@ +# Boundary-Aware Trade Counting + Dead Code Cleanup — Implementation Plan + +> **For agentic workers:** REQUIRED SUB-SKILL: Use superpowers:subagent-driven-development (recommended) or superpowers:executing-plans to implement this plan task-by-task. Steps use checkbox (`- [ ]`) syntax for tracking. + +**Goal:** Fix fragmented trade counting across GPU thread chunk boundaries, remove 3 dead parallel reduction arrays + dead per-thread drawdown variables, and extract BARS_PER_YEAR constant in coordinator_extended.rs. + +**Architecture:** The backtest metrics kernel uses 256 threads processing consecutive chunks. Trades spanning chunk boundaries were fragmented (returns lost, counts wrong). Replace with hybrid: parallel per-thread boundary metadata export + thread-0 O(256) stitching pass. Simultaneously remove dead code from prior exact-drawdown fix (s_max_dd, s_trades, s_wins arrays + local_peak/local_max_dd per-thread vars). Shared memory layout shrinks from 9 to 6 reduction arrays (25KB → 22KB). + +**Tech Stack:** CUDA kernel (C), Rust (cudarc 0.19.3), nvcc PTX compilation. + +--- + +## Task 1: Dead code cleanup — remove 3 dead reduction arrays and per-thread drawdown vars + +**Files:** +- Modify: `crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu` +- Modify: `crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs:1795-1798` + +This task removes the dead code FIRST so that Task 2 starts from a clean baseline. After this task, the kernel still has the OLD (broken) trade counting, but the dead drawdown/trades/wins reduction arrays are gone. + +- [ ] **Step 1: Update kernel header comment** + +In `backtest_metrics_kernel.cu`, replace lines 1-18 (the header comment) with: + +```cuda +// Per-window metrics reduction kernel. +// One block per window. Threads cooperate to reduce step_returns. +// +// Output per window [14 floats]: +// [0] sharpe_ratio (annualized via configurable annualization_factor) +// [1] total_pnl (multiplicatively compounded fractional return, e.g. 0.05 = +5%) +// [2] max_drawdown (exact sequential scan, worst peak-to-trough fraction) +// [3] sortino_ratio +// [4] win_rate (winning_trades / total_trades, exact via boundary stitching) +// [5] total_trades (exact position changes, boundary-aware) +// [6] var_95 (5th-percentile return — Value at Risk at 95% confidence) +// [7] cvar_95 (mean of returns below VaR — Expected Shortfall) +// [8] calmar_ratio (annualized mean return / max drawdown) +// [9] omega_ratio (sum of gains / sum of losses) +// [10] buy_count (exposure_idx 5-8 — Long25..Long100) +// [11] sell_count (exposure_idx 0-3 — Short100..Short25) +// [12] hold_count (action 4 — Flat) +// [13] unique_action_count (popcount of 9-bit mask — how many exposure levels 0..8 were used) +``` + +- [ ] **Step 2: Update shared memory layout comment and pointers** + +Replace lines 38-61 (shared memory layout comment + pointer declarations) with: + +```cuda + // Shared memory layout (6 reduction arrays + sort scratch): + // [0 .. stride) : s_sum + // [stride .. 2*stride) : s_sq_sum + // [2*stride..3*stride) : s_down_sq + // [3*stride..4*stride) : s_buys + // [4*stride..5*stride) : s_sells + // [5*stride..6*stride) : s_holds + // [6*stride .. 6*stride + 4096) : s_sorted (bitonic sort scratch + boundary data) + extern __shared__ float shmem[]; + float* s_sum = shmem; // [blockDim.x] + float* s_sq_sum = shmem + stride; // [blockDim.x] + float* s_down_sq = shmem + 2*stride; // [blockDim.x] (downside deviation) + // Action distribution: exposure 0-3=sell, 4=flat, 5-8=buy + float* s_buys = shmem + 3*stride; // [blockDim.x] + float* s_sells = shmem + 4*stride; // [blockDim.x] + float* s_holds = shmem + 5*stride; // [blockDim.x] + float* s_sorted = shmem + 6*stride; // [4096] for bitonic sort + boundary stitching +``` + +- [ ] **Step 3: Remove dead per-thread variables** + +Replace lines 63-71 (per-thread local accumulators) with: + +```cuda + // Pass 1: per-thread local accumulators + float local_sum = 0.0f, local_sq = 0.0f, local_down = 0.0f; + // Multiplicative compounding for equity product reduction (total_pnl) + float local_cum = 1.0f; + int local_buys = 0, local_sells = 0, local_holds = 0; + int local_action_mask = 0; // 9-bit bitmask: bit i set if exposure_idx i was seen +``` + +Note: `local_peak`, `local_max_dd`, `local_wins`, `local_trades`, `prev_action`, `trade_cum_return` are all removed. `local_cum` is kept for the equity product reduction. + +- [ ] **Step 4: Remove dead per-bar drawdown tracking and old trade counting from loop body** + +Replace lines 82-129 (the entire for loop body + post-loop trade close) with: + +```cuda + for (int i = chunk_start; i < chunk_end; i++) { + float r = step_returns[base + i]; + // Sharpe/Sortino use arithmetic mean of per-bar returns + local_sum += r; + local_sq += r * r; + if (r < 0.0f) local_down += r * r; + + // Equity product for total_pnl (drawdown computed exactly by thread 0 later) + local_cum *= (1.0f + r); + + // Action distribution: decode exposure from factored action. + // Factored action = exposure * (b1*b2) + order * b2 + urgency. + // exposure_idx: 0-3=sell, 4=flat, 5-8=buy (for 9-action space) + int act = actions_history[base + i]; + int exp_idx = act / (DQN_ORDER_ACTIONS * DQN_URGENCY_ACTIONS); + if (exp_idx < DQN_NUM_ACTIONS / 2) local_sells++; + else if (exp_idx == DQN_NUM_ACTIONS / 2) local_holds++; + else local_buys++; + + // Unique action tracking: set bit per exposure index (0..8) + if (exp_idx >= 0 && exp_idx < DQN_NUM_ACTIONS) + local_action_mask |= (1 << exp_idx); + } +``` + +Note: the entire drawdown tracking block (lines 90-95: `local_cum *= ..., local_peak = fmaxf(...), dd = ..., local_max_dd = ...`) is removed. Only `local_cum *= (1.0f + r)` is kept for the equity product. The entire old trade counting block (lines 97-106 + 120-129) is removed — will be replaced by boundary tracking in Task 2. + +- [ ] **Step 5: Update shared memory stores** + +Replace lines 131-141 (store to shmem) with: + +```cuda + // Store local values to shared memory for parallel reduction + s_sum[tid] = local_sum; + s_sq_sum[tid] = local_sq; + s_down_sq[tid] = local_down; + s_buys[tid] = (float)local_buys; + s_sells[tid] = (float)local_sells; + s_holds[tid] = (float)local_holds; + __syncthreads(); +``` + +- [ ] **Step 6: Update block reduction from 9 arrays to 6** + +Replace lines 143-157 (block reduction loop) with: + +```cuda + // Block-level parallel reduction for 6 arrays + for (int s = stride / 2; s > 0; s >>= 1) { + if (tid < s) { + s_sum[tid] += s_sum[tid + s]; + s_sq_sum[tid] += s_sq_sum[tid + s]; + s_down_sq[tid] += s_down_sq[tid + s]; + s_buys[tid] += s_buys[tid + s]; + s_sells[tid] += s_sells[tid + s]; + s_holds[tid] += s_holds[tid + s]; + } + __syncthreads(); + } +``` + +- [ ] **Step 7: Update thread-0 metrics output (remove s_trades/s_wins references)** + +In the thread-0 section (lines 176-214), replace lines 207-213 with: + +```cuda + // Trade count and win rate: placeholder zeros (replaced by boundary stitching in next commit) + metrics_out[out_base + 4] = 0.0f; // win_rate (TODO: boundary stitching) + metrics_out[out_base + 5] = 0.0f; // total_trades (TODO: boundary stitching) + metrics_out[out_base + 10] = s_buys[0]; // buy action count + metrics_out[out_base + 11] = s_sells[0]; // sell action count + metrics_out[out_base + 12] = s_holds[0]; // hold action count +``` + +- [ ] **Step 8: Update shmem_bytes in gpu_backtest_evaluator.rs** + +In `gpu_backtest_evaluator.rs:1795-1798`, replace: + +```rust + // Shared memory: 9 reduction arrays × 256 threads × 4 bytes = 9 KB + // + 4096 floats for bitonic sort scratch = 16 KB + // Total = 25 KB (well within the 48 KB L1/shmem limit) + let shmem_bytes = (256_u32 * 9 + 4096) * std::mem::size_of::() as u32; +``` + +with: + +```rust + // Shared memory: 6 reduction arrays × 256 threads × 4 bytes = 6 KB + // + 4096 floats for sort scratch / boundary data = 16 KB + // Total = 22 KB (well within the 48 KB L1/shmem limit) + let shmem_bytes = (256_u32 * 6 + 4096) * std::mem::size_of::() as u32; +``` + +- [ ] **Step 9: Build + test** + +Run: `SQLX_OFFLINE=true cargo check -p ml --lib` +Expected: compiles clean + +Run: `SQLX_OFFLINE=true cargo test -p ml --lib -- backtest` +Expected: all backtest tests pass (trade count/win rate will be 0 temporarily — that's expected, fixed in Task 2) + +- [ ] **Step 10: Commit** + +```bash +git add crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs +git commit -m "refactor: remove 3 dead reduction arrays + per-thread drawdown vars from metrics kernel + +s_max_dd, s_trades, s_wins parallel reductions were made redundant by +the exact sequential drawdown scan and upcoming boundary stitching. +Reduces shared memory from 25KB to 22KB and removes 3 dead ops/bar +from the per-thread loop. Trade count/win rate temporarily zeroed — +restored by boundary stitching in next commit." +``` + +--- + +## Task 2: Boundary-aware parallel trade counting with stitching + +**Files:** +- Modify: `crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu` + +This task adds the hybrid parallel trade counting: per-thread boundary metadata + thread-0 stitching. It replaces the temporary zeros from Task 1 with exact values. + +- [ ] **Step 1: Add boundary-tracking variables to per-thread accumulators** + +After the existing per-thread variables (from Task 1), add: + +```cuda + // Boundary-aware trade tracking: each thread exports metadata about + // trades at chunk boundaries. Thread 0 stitches these to produce + // exact trade counts and win rates across the full window. + int bnd_first_action = -1; // -1 signals empty chunk + int bnd_last_action = -1; + float bnd_prefix_return = 0.0f; + float bnd_suffix_return = 0.0f; + int bnd_complete_trades = 0; + int bnd_complete_wins = 0; + int bnd_num_changes = 0; + int bnd_cur_action = -1; + float bnd_cur_return = 0.0f; +``` + +- [ ] **Step 2: Add boundary tracking to the per-bar loop** + +Inside the `for (int i = chunk_start; i < chunk_end; i++)` loop, AFTER the action distribution block and BEFORE the closing brace, add: + +```cuda + // Boundary-aware trade tracking + if (i == chunk_start) { + bnd_first_action = act; + bnd_cur_action = act; + } + + if (act != bnd_cur_action) { + bnd_num_changes++; + if (bnd_num_changes == 1) { + // First change: everything accumulated so far is the prefix + bnd_prefix_return = bnd_cur_return; + } else { + // Interior change: segment since last change is a complete trade + bnd_complete_trades++; + if (bnd_cur_return > 0.0f) bnd_complete_wins++; + } + bnd_cur_return = 0.0f; + bnd_cur_action = act; + } + + bnd_cur_return += r; + bnd_last_action = act; +``` + +- [ ] **Step 3: Add post-loop suffix/prefix assignment** + +After the closing brace of the for loop, add: + +```cuda + // Post-loop: assign suffix and handle no-change edge case + bnd_suffix_return = bnd_cur_return; + if (bnd_num_changes == 0) { + // Entire chunk is one partial trade segment (prefix == suffix) + bnd_prefix_return = bnd_cur_return; + } +``` + +- [ ] **Step 4: Export boundary data to shared memory** + +After the equity product reduction (after the `// s_sorted[0] now holds the full-window equity` comment and its `__syncthreads()`), add: + +```cuda + // ── Export boundary trade data to shared memory ────────────────────── + // Reuse s_sorted[stride..8*stride] for 7 boundary arrays (1792 floats). + // This region is free between the equity product reduction (which only + // uses s_sorted[0]) and the bitonic sort (which starts later). + // Int values stored via __int_as_float for bit-exact round-trip. + { + float* s_bnd = &s_sorted[stride]; + s_bnd[0 * stride + tid] = __int_as_float(bnd_first_action); + s_bnd[1 * stride + tid] = __int_as_float(bnd_last_action); + s_bnd[2 * stride + tid] = bnd_prefix_return; + s_bnd[3 * stride + tid] = bnd_suffix_return; + s_bnd[4 * stride + tid] = __int_as_float(bnd_complete_trades); + s_bnd[5 * stride + tid] = __int_as_float(bnd_complete_wins); + s_bnd[6 * stride + tid] = __int_as_float(bnd_num_changes); + } + __syncthreads(); +``` + +- [ ] **Step 5: Add stitching pass in thread-0 section** + +In the thread-0 section (`if (tid == 0)`), AFTER the exact drawdown sequential scan and BEFORE the metrics output writes, add the stitching code: + +```cuda + // ── Boundary stitching: exact trade count + win rate ───────────── + // Sum interior complete trades from all threads, then stitch + // boundary trades by walking the per-thread metadata left-to-right. + float* s_bnd = &s_sorted[stride]; + int total_trades = 0; + int total_wins = 0; + + // Sum interior trades (parallel-counted, always correct) + for (int t = 0; t < stride; t++) { + total_trades += __float_as_int(s_bnd[4 * stride + t]); + total_wins += __float_as_int(s_bnd[5 * stride + t]); + } + + // Stitch boundary trades + float open_return = 0.0f; + int open_action = -1; + + for (int t = 0; t < stride; t++) { + int fa = __float_as_int(s_bnd[0 * stride + t]); + int la = __float_as_int(s_bnd[1 * stride + t]); + float pr = s_bnd[2 * stride + t]; + float sr = s_bnd[3 * stride + t]; + int nc = __float_as_int(s_bnd[6 * stride + t]); + + if (fa < 0) continue; // empty chunk + + // Step 1: connect this chunk's start to the open trade + if (open_action < 0) { + // First valid chunk — open the first trade (no closure yet) + open_action = fa; + open_return = 0.0f; + } else if (fa != open_action) { + // Action changed at chunk boundary — close the open trade + total_trades++; + if (open_return > 0.0f) total_wins++; + open_action = fa; + open_return = 0.0f; + } + + // Step 2: process this chunk's contribution + if (nc == 0) { + // No action changes — entire chunk extends the open trade + open_return += pr; + } else { + // Prefix extends and closes the open trade + open_return += pr; + total_trades++; + if (open_return > 0.0f) total_wins++; + // Interior trades already counted above + // Suffix becomes the new open trade + open_action = la; + open_return = sr; + } + } + + // Close the final open trade (window end) + if (open_action >= 0) { + total_trades++; + if (open_return > 0.0f) total_wins++; + } +``` + +- [ ] **Step 6: Update metrics output to use stitched values** + +Replace the temporary trade count/win rate output (from Task 1 Step 7) with: + +```cuda + float total_trades_f = (float)total_trades; + metrics_out[out_base + 4] = (total_trades_f > 0.0f) ? (float)total_wins / total_trades_f : 0.0f; + metrics_out[out_base + 5] = total_trades_f; + metrics_out[out_base + 10] = s_buys[0]; + metrics_out[out_base + 11] = s_sells[0]; + metrics_out[out_base + 12] = s_holds[0]; +``` + +- [ ] **Step 7: Build + test** + +Run: `SQLX_OFFLINE=true cargo check -p ml --lib` +Expected: compiles clean + +Run: `SQLX_OFFLINE=true cargo test -p ml --lib -- backtest` +Expected: all backtest tests pass + +Run: `SQLX_OFFLINE=true cargo test -p ml --lib -- objective` +Expected: all objective function tests pass + +Run: `SQLX_OFFLINE=true cargo test -p ml --lib` +Expected: 887 pass, 0 fail + +- [ ] **Step 8: Commit** + +```bash +git add crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu +git commit -m "feat: boundary-aware parallel trade counting with exact stitching + +Hybrid approach: each of 256 threads exports boundary metadata (first/last +action, prefix/suffix return, interior trade count). Thread 0 stitches +boundaries in O(256) to produce exact trade count and win rate. + +Fixes: trades spanning chunk boundaries were fragmented (returns lost, +counts incorrect). Now every trade's cumulative return is tracked exactly +regardless of which thread chunks it spans. + +7 boundary values stored in s_sorted[stride..8*stride] — zero extra +shared memory (reuses sort scratch between equity reduction and bitonic sort)." +``` + +--- + +## Task 3: coordinator_extended.rs — extract BARS_PER_YEAR constant + +**Files:** +- Modify: `crates/ml-ensemble/src/coordinator_extended.rs:338-340` + +- [ ] **Step 1: Replace inline formula with named constant** + +In `coordinator_extended.rs`, replace lines 338-340: + +```rust + // Annualize assuming 252 trading days, 6.5 hours per day, predictions every minute + let sharpe = if std_dev > 1e-10 { + let annualization_factor = (252.0 * 6.5 * 60.0_f64).sqrt(); +``` + +with: + +```rust + // Annualize: 1-minute bars, 390 bars/day × 252 days/year + const BARS_PER_YEAR: f64 = 390.0 * 252.0; + let sharpe = if std_dev > 1e-10 { + let annualization_factor = BARS_PER_YEAR.sqrt(); +``` + +- [ ] **Step 2: Build + test** + +Run: `SQLX_OFFLINE=true cargo check -p ml-ensemble --lib` +Expected: compiles clean + +Run: `SQLX_OFFLINE=true cargo test -p ml-ensemble --lib` +Expected: 118 pass, 0 fail + +- [ ] **Step 3: Commit** + +```bash +git add crates/ml-ensemble/src/coordinator_extended.rs +git commit -m "refactor: extract BARS_PER_YEAR constant in coordinator_extended.rs + +Replaces inline (252.0 * 6.5 * 60.0).sqrt() with named constant +matching the pattern used in financials.rs and ab_testing.rs." +``` + +--- + +## Task 4: Full integration test + +**Files:** +- No new files — validation only + +- [ ] **Step 1: Run full ml test suite** + +Run: `SQLX_OFFLINE=true cargo test -p ml --lib` +Expected: 887 pass, 0 fail + +- [ ] **Step 2: Run full ml-dqn test suite** + +Run: `SQLX_OFFLINE=true cargo test -p ml-dqn --lib` +Expected: 359 pass, 0 fail + +- [ ] **Step 3: Run full ml-ensemble test suite** + +Run: `SQLX_OFFLINE=true cargo test -p ml-ensemble --lib` +Expected: 118 pass, 0 fail + +- [ ] **Step 4: Workspace compilation check** + +Run: `SQLX_OFFLINE=true cargo check --workspace` +Expected: compiles clean, no errors + +- [ ] **Step 5: Verify PTX compilation test** + +Run: `SQLX_OFFLINE=true cargo test -p ml --lib -- test_metrics_ptx_compilation` +Expected: passes (verifies the modified kernel compiles to valid PTX) + +--- + +## Execution Order + +``` +Task 1 (dead code cleanup) → Task 2 (boundary stitching) → Task 3 (coordinator constant) → Task 4 (integration test) +``` + +Tasks 1 and 2 MUST be sequential (Task 2 builds on Task 1's cleaned-up kernel). Task 3 is independent but ordered last for clean git history. Task 4 validates everything. diff --git a/docs/superpowers/specs/2026-03-25-boundary-trade-counting-design.md b/docs/superpowers/specs/2026-03-25-boundary-trade-counting-design.md new file mode 100644 index 000000000..937db867c --- /dev/null +++ b/docs/superpowers/specs/2026-03-25-boundary-trade-counting-design.md @@ -0,0 +1,265 @@ +# Boundary-Aware Parallel Trade Counting — Design Spec + +## Problem + +The GPU backtest metrics kernel (`backtest_metrics_kernel.cu`) uses 256 threads to process consecutive chunks of a 10K-bar evaluation window. Trades that span chunk boundaries are fragmented: each thread tracks trades independently, so a single trade held across 3 chunks gets split into 3 partial segments. The cumulative return of the fragmented portions is lost, producing incorrect trade counts and corrupted win rates. + +Additionally, `coordinator_extended.rs` uses a correct but non-parameterized annualization formula (`(252.0 * 6.5 * 60.0).sqrt()`). This should use a named constant for consistency with the rest of the codebase. + +## Design: Hybrid Parallel + Boundary Stitching + +### Architecture + +**Phase 1 — Parallel per-thread pass (all 256 threads):** +Each thread processes its consecutive chunk of bars and computes: +- Standard accumulators (sum, sq_sum, downside, equity, drawdown, action distribution) — unchanged +- **NEW**: Boundary metadata instead of local trade counting + +**Phase 2 — Boundary export (all threads → shared memory):** +Each thread writes 7 boundary values to `s_sorted[stride..8*stride]` (reuses existing sort scratch — zero extra shared memory). + +**Phase 3 — Stitching (thread 0 only, O(256)):** +Thread 0 walks boundary metadata left-to-right, merging partial trades across chunk boundaries, accumulating exact per-trade returns, and counting trades/wins. + +### Per-Thread Boundary Metadata (7 values per thread) + +| Field | Type | Description | +|-------|------|-------------| +| `first_action` | int | Action ID at chunk's first bar (-1 if chunk is empty) | +| `last_action` | int | Action ID at chunk's last bar | +| `prefix_return` | float | Cumulative return from chunk start → first action change (or total chunk return if no changes) | +| `suffix_return` | float | Cumulative return from last action change → chunk end (or total chunk return if no changes) | +| `complete_trades` | int | Trades fully contained within chunk (between first and last change) | +| `complete_wins` | int | Winning interior trades (complete_return > 0) | +| `num_changes` | int | Number of action changes in chunk | + +**Edge case — no action changes:** When `num_changes == 0`, the entire chunk is one partial trade segment. `prefix_return == suffix_return == total chunk return`. The stitching pass merges it with adjacent chunks. + +**Edge case — empty window (`wlen == 0`):** All threads get empty chunks (`chunk_start >= chunk_end`), so `first_action == -1`. The stitching pass skips all threads, closes no trades, and produces `total_trades = 0`, `win_rate = 0.0`. + +### Per-Thread Loop Code + +The old trade-counting variables (`prev_action`, `trade_cum_return`, `local_trades`, `local_wins`) are **replaced** by the 7 boundary-tracking variables: + +```cuda +// NEW: boundary-aware trade tracking variables +int bnd_first_action = -1; // -1 signals empty chunk +int bnd_last_action = -1; +float bnd_prefix_return = 0.0f; +float bnd_suffix_return = 0.0f; +int bnd_complete_trades = 0; +int bnd_complete_wins = 0; +int bnd_num_changes = 0; +int bnd_cur_action = -1; +float bnd_cur_return = 0.0f; + +for (int i = chunk_start; i < chunk_end; i++) { + float r = step_returns[base + i]; + int act = actions_history[base + i]; + + // ... (existing Sharpe/drawdown/action-distribution accumulators unchanged) ... + + // Boundary-aware trade tracking + if (i == chunk_start) { + bnd_first_action = act; + bnd_cur_action = act; + } + + if (act != bnd_cur_action) { + bnd_num_changes++; + if (bnd_num_changes == 1) { + // First change: everything accumulated so far is the prefix + bnd_prefix_return = bnd_cur_return; + } else { + // Interior change: the segment since the last change is a complete trade + bnd_complete_trades++; + if (bnd_cur_return > 0.0f) bnd_complete_wins++; + } + bnd_cur_return = 0.0f; + bnd_cur_action = act; + } + + bnd_cur_return += r; + bnd_last_action = act; +} + +// Post-loop: assign suffix and handle no-change case +bnd_suffix_return = bnd_cur_return; +if (bnd_num_changes == 0) { + // Entire chunk is one partial trade segment + bnd_prefix_return = bnd_cur_return; +} +``` + +### Stitching Algorithm + +The stitching matches the existing kernel's trade semantics: a "trade" is a contiguous run of bars with the same action. The first bar of the window **opens** the first trade; it is **closed** only when the action changes or the window ends. This matches the old `if (act != prev_action && prev_action >= 0)` logic — the initial position is not counted as a closure until it changes. + +``` +total_trades = sum(complete_trades[t] for all threads) +total_wins = sum(complete_wins[t] for all threads) +open_return = 0.0 +open_action = -1 + +for each thread t (0..blockDim.x): + fa = first_action[t] + la = last_action[t] + pr = prefix_return[t] + sr = suffix_return[t] + nc = num_changes[t] + + if fa < 0: skip (empty chunk) + + // Step 1: connect this chunk's start to the open trade + if open_action < 0: + // First valid chunk — open the first trade (no closure yet) + open_action = fa + open_return = 0.0 + else if fa != open_action: + // Action changed at chunk boundary — close the open trade + total_trades++ + if open_return > 0.0: total_wins++ + open_action = fa + open_return = 0.0 + + // Step 2: process this chunk's contribution + if nc == 0: + // No action changes — entire chunk extends the open trade + open_return += pr // (pr == sr == total chunk return) + else: + // Prefix extends and closes the open trade (first intra-chunk change ends it) + open_return += pr + total_trades++ + if open_return > 0.0: total_wins++ + // Interior trades already counted by per-thread code + // Suffix becomes the new open trade + open_action = la + open_return = sr + +// Close the final open trade (window end) +if open_action >= 0: + total_trades++ + if open_return > 0.0: total_wins++ +``` + +**Key properties:** +- `open_return` is always reset to 0.0 when `open_action` changes (no return leakage across trades) +- The first chunk's prefix opens a trade but does not close it (matching the old `prev_action >= 0` guard) +- Interior complete trades are summed from per-thread data (parallel), boundary trades are stitched (sequential) + +### Dead Code Removal + +With this fix (stitched trade counting) and the prior fix (exact sequential drawdown scan), three of the nine parallel reduction arrays are now dead: + +| Array | Was | Dead because | +|-------|-----|-------------| +| `s_max_dd` | Parallel drawdown max-reduction | Overridden by thread-0 exact sequential scan | +| `s_trades` | Parallel trade count sum | Replaced by boundary stitching | +| `s_wins` | Parallel win count sum | Replaced by boundary stitching | + +**Remove from the kernel:** +- Per-thread variables: `local_peak`, `local_max_dd` (drawdown tracking per-thread is dead — `local_cum` is still needed for the equity product reduction) +- Shared memory arrays: `s_max_dd`, `s_trades`, `s_wins` +- Reduction: remove the 3 dead arrays from the block reduction loop (9 arrays → 6) +- Per-thread drawdown tracking: remove the `local_peak`/`local_max_dd`/dd computation from the per-bar loop (saves 1 `fmaxf` + 1 division + 1 comparison per bar) + +### Shared Memory Layout (after cleanup) + +Reduction shrinks from 9 arrays to 6. Total shared memory drops from 25KB to 22KB: + +``` +s_sum [0*stride .. 1*stride) // Sharpe mean +s_sq_sum [1*stride .. 2*stride) // Sharpe variance +s_down_sq [2*stride .. 3*stride) // Sortino downside +s_buys [3*stride .. 4*stride) // action distribution +s_sells [4*stride .. 5*stride) +s_holds [5*stride .. 6*stride) +s_sorted [6*stride .. 6*stride + 4096) // sort scratch + boundary data +``` + +The 7 × 256 = 1,792 boundary values are stored in `s_sorted[stride..8*stride]` (7 arrays of `stride` floats, starting at offset `stride`). This memory is written after the equity product reduction and read before the bitonic sort — no conflict. + +**Update `shmem_bytes` in `gpu_backtest_evaluator.rs`**: change from `(256 * 9 + 4096) * 4` to `(256 * 6 + 4096) * 4`. + +### Int-as-Float Storage + +Boundary ints (action IDs, counts) are stored in float shared memory via `__int_as_float()` / `__float_as_int()` — bit-exact round-trip, no precision loss. + +### Correctness Verification + +**Test case 1: Single trade spanning all chunks** +Window `[5,5,5,...,5]` (10K bars, all same action) +- Every thread: `num_changes=0`, `prefix_return = chunk_total` +- Stitching: + - t=0: `open_action=-1` → set `open_action=5`, `open_return=0`; nc=0 → `open_return += chunk_0_total` + - t=1..255: `fa=5 == open_action` → nc=0 → `open_return += chunk_t_total` + - Final close → **1 trade, return = sum of all bars** ✓ + +**Test case 2: Trades at chunk boundaries** +Window `[5,5,5, 3,3, 7,7,7, 2,2]`, 2 chunks of 5 bars +- Thread 0 `[5,5,5,3,3]`: first=5, last=3, prefix=r[0]+r[1]+r[2], suffix=r[3]+r[4], nc=1, interior=0 +- Thread 1 `[7,7,7,2,2]`: first=7, last=2, prefix=r[5]+r[6]+r[7], suffix=r[8]+r[9], nc=1, interior=0 +- Stitching step-by-step: + 1. t=0: `open=-1` → set open=5. nc=1 → `open_return += prefix(r[0..2])` → CLOSE trade 1 (action=5, return=r[0..2]). Open=3, return=suffix(r[3..4]) + 2. t=1: `fa=7 != open=3` → CLOSE trade 2 (action=3, return=r[3..4]). Open=7. nc=1 → `open_return += prefix(r[5..7])` → CLOSE trade 3 (action=7, return=r[5..7]). Open=2, return=suffix(r[8..9]) + 3. Final close → trade 4 (action=2, return=r[8..9]) +- Result: **4 trades, each with exact cumulative return** ✓ + +**Test case 3: Multiple changes in one chunk** +Actions `[5,3,7,2,8]` in one chunk +- nc=4, prefix=r[0] (action=5), suffix=r[4] (action=8), interior trades: action=3(r[1]), action=7(r[2]), action=2(r[3]) → 3 complete interior trades +- Stitching: open=5 → prefix(r[0]) → CLOSE trade (action=5). Open=8, return=suffix(r[4]). Final close trade (action=8). +- Total: 3 interior + 1 prefix close + 1 final close = **5 trades** ✓ + +**Test case 4: Alternating actions every bar (stress test)** +Actions `[0,1,2,3,4,5,6,7,8,0,1,...]` for 10K bars +- Each thread chunk of ~39 bars: nc=38, prefix=r[0], suffix=r[38], interior=37 complete trades +- Stitching: each boundary has `fa != la`, so boundary trades close normally +- Total: ~10K trades, each 1 bar long, each with exact single-bar return ✓ + +**Test case 5: Empty window (`wlen == 0`)** +- All threads get empty chunks (`first_action = -1`) +- Stitching: all skipped, `open_action` stays -1, no final close +- Result: **0 trades, win_rate = 0.0** ✓ + +### Kernel Header Comment Update + +Update the kernel file's header comment (lines 4-5) to reflect exact counting: +``` +// [4] win_rate (winning_trades / total_trades, exact via boundary stitching) +// [5] total_trades (exact position changes, boundary-aware) +``` + +### coordinator_extended.rs Fix + +Replace the inline formula `(252.0 * 6.5 * 60.0_f64).sqrt()` with a named constant: + +```rust +const BARS_PER_YEAR: f64 = 390.0 * 252.0; +let annualization_factor = BARS_PER_YEAR.sqrt(); +``` + +This matches the pattern already used in `financials.rs` and `ab_testing.rs`. + +## Files Modified + +| File | Change | +|------|--------| +| `crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu` | Replace trade counting with boundary-aware hybrid; remove 3 dead reduction arrays + per-thread drawdown vars; update header comment | +| `crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs` | Update `shmem_bytes` from `(256*9+4096)*4` to `(256*6+4096)*4` | +| `crates/ml-ensemble/src/coordinator_extended.rs` | Extract `BARS_PER_YEAR` constant | + +## Performance Impact + +- **Per-thread pass**: +4 registers for boundary tracking, −4 registers from removed drawdown+trade vars. Net: 0 registers. +- **Reduction**: 6 arrays instead of 9 → 33% fewer reduction operations. +- **Per-bar loop**: Removes 3 operations/bar (fmaxf peak, drawdown division, drawdown comparison) from the parallel pass. Still computed exactly by thread-0 sequential scan. +- **Shared memory**: 22KB instead of 25KB → 12% reduction, more headroom for L1 cache. +- **Boundary export**: 7 shared memory writes per thread + 1 `__syncthreads()`. Negligible. +- **Stitching**: O(256) loop in thread 0. Negligible vs existing O(10K) sequential scan. +- **Net effect**: Kernel is slightly FASTER due to dead code removal, despite added stitching complexity. + +## Risks + +- **Complexity**: The stitching algorithm has edge cases (empty chunks, single-bar chunks, all-same-action windows, empty windows). Each is covered by the 5 test cases above and will have dedicated unit tests. +- **Shared memory aliasing**: `s_sorted` is reused for boundary data between equity reduction and bitonic sort. Execution order enforced by `__syncthreads()` barriers — safe.