fix: backtest evaluator — correct annualization, window sizing, objective calibration
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) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -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
|
||||
);
|
||||
|
||||
@@ -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,
|
||||
|
||||
@@ -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
|
||||
}
|
||||
|
||||
494
docs/superpowers/plans/2026-03-25-boundary-trade-counting.md
Normal file
494
docs/superpowers/plans/2026-03-25-boundary-trade-counting.md
Normal file
@@ -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::<f32>() 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::<f32>() 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.
|
||||
@@ -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.
|
||||
Reference in New Issue
Block a user