diff --git a/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu index 041ceb7a2..25a055a18 100644 --- a/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu +++ b/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu @@ -2,19 +2,19 @@ // One block per window. Threads cooperate to reduce step_returns. // // Output per window [14 floats]: -// [0] sharpe_ratio (annualized, sqrt(252)) +// [0] sharpe_ratio (annualized via configurable annualization_factor) // [1] total_pnl (multiplicatively compounded fractional return, e.g. 0.05 = +5%) -// [2] max_drawdown (worst peak-to-trough as fraction, e.g. 0.10 = -10%) +// [2] max_drawdown (exact sequential scan, worst peak-to-trough fraction) // [3] sortino_ratio -// [4] win_rate (winning_trades / total_trades) -// [5] total_trades (position changes) +// [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 2 — Flat) +// [12] hold_count (exposure_idx 4 — Flat) // [13] unique_action_count (popcount of 9-bit mask — how many exposure levels 0..8 were used) extern "C" __global__ void compute_backtest_metrics( @@ -35,40 +35,30 @@ extern "C" __global__ void compute_backtest_metrics( int stride = blockDim.x; int base = w * max_len; - // Shared memory layout: - // [0 .. stride) : s_sum (9 reduction arrays of size stride) + // 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_max_dd - // [4*stride..5*stride) : s_wins - // [5*stride..6*stride) : s_trades - // [6*stride..7*stride) : s_buys (action distribution) - // [7*stride..8*stride) : s_sells - // [8*stride..9*stride) : s_holds - // [9*stride .. 9*stride + 4096) : s_sorted (bitonic sort scratch, up to 4096 returns) + // [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) - float* s_max_dd = shmem + 3*stride; // [blockDim.x] (max drawdown) - // wins and trades stored as float for reduction compatibility - float* s_wins = shmem + 4*stride; // [blockDim.x] - float* s_trades = shmem + 5*stride; // [blockDim.x] - // Action distribution: buy (3,4), sell (0,1), hold (2) - float* s_buys = shmem + 6*stride; // [blockDim.x] - float* s_sells = shmem + 7*stride; // [blockDim.x] - float* s_holds = shmem + 8*stride; // [blockDim.x] - float* s_sorted = shmem + 9*stride; // [4096] for bitonic sort + // 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 // Pass 1: per-thread local accumulators float local_sum = 0.0f, local_sq = 0.0f, local_down = 0.0f; - // Multiplicative compounding: equity starts at 1.0 (= 100%) - float local_cum = 1.0f, local_peak = 1.0f, local_max_dd = 0.0f; - int local_wins = 0, local_trades = 0; + // 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 - int prev_action = -1; - float trade_cum_return = 0.0f; // cumulative return during current trade // Process consecutive chunks (not strided) for correct drawdown tracking. // Strided access (bars 0, 256, 512...) skips intermediate peaks/troughs, @@ -81,33 +71,18 @@ extern "C" __global__ void compute_backtest_metrics( for (int i = chunk_start; i < chunk_end; i++) { float r = step_returns[base + i]; - // Sharpe/Sortino use arithmetic mean of per-bar returns (correct even - // for multiplicatively compounded portfolios) + // Sharpe/Sortino use arithmetic mean of per-bar returns local_sum += r; local_sq += r * r; if (r < 0.0f) local_down += r * r; - // Drawdown tracking with multiplicative compounding + // Equity product for total_pnl (drawdown computed exactly by thread 0 later) local_cum *= (1.0f + r); - local_peak = fmaxf(local_peak, local_cum); - // Fractional drawdown: (peak - current) / peak - float dd = (local_peak - local_cum) / fmaxf(local_peak, 1e-8f); - local_max_dd = fmaxf(local_max_dd, dd); - - // Trade counting and per-trade win/loss tracking - int act = actions_history[base + i]; - if (act != prev_action && prev_action >= 0) { - // Position changed — close the previous trade - local_trades++; - if (trade_cum_return > 0.0f) local_wins++; - trade_cum_return = 0.0f; // reset for next trade - } - trade_cum_return += r; // accumulate return during this trade - prev_action = act; // 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++; @@ -117,33 +92,22 @@ extern "C" __global__ void compute_backtest_metrics( if (exp_idx >= 0 && exp_idx < DQN_NUM_ACTIONS) local_action_mask |= (1 << exp_idx); } - // Close the final trade (if any bars were processed) - if (chunk_start < chunk_end && prev_action >= 0) { - local_trades++; - if (trade_cum_return > 0.0f) local_wins++; - } - // Store ALL local values to shared memory + // 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_max_dd[tid] = local_max_dd; - s_wins[tid] = (float)local_wins; - s_trades[tid] = (float)local_trades; s_buys[tid] = (float)local_buys; s_sells[tid] = (float)local_sells; s_holds[tid] = (float)local_holds; __syncthreads(); - // Block-level parallel reduction for ALL 9 arrays + // 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_max_dd[tid] = fmaxf(s_max_dd[tid], s_max_dd[tid + s]); // max reduction - s_wins[tid] += s_wins[tid + s]; - s_trades[tid] += s_trades[tid + s]; s_buys[tid] += s_buys[tid + s]; s_sells[tid] += s_sells[tid + s]; s_holds[tid] += s_holds[tid + s]; @@ -176,19 +140,35 @@ extern "C" __global__ void compute_backtest_metrics( float std = sqrtf(fmaxf(var, 1e-10f)); float down_std = sqrtf(fmaxf(s_down_sq[0] / n, 1e-10f)); + // ── Exact max drawdown via sequential scan ────────────────────── + // The per-thread parallel reduction misses cross-chunk drawdowns + // because each thread tracks equity from local_peak=1.0, unaware + // of the global peak. Sequential scan over step_returns gives the + // TRUE maximum peak-to-trough drawdown. + // Cost: O(wlen) ≈ 10K iterations ≈ microseconds on GPU. + float exact_equity = 1.0f; + float exact_peak = 1.0f; + float exact_max_dd = 0.0f; + for (int i = 0; i < wlen; i++) { + float r = step_returns[base + i]; + exact_equity *= (1.0f + r); + exact_peak = fmaxf(exact_peak, exact_equity); + float dd = (exact_peak - exact_equity) / fmaxf(exact_peak, 1e-8f); + exact_max_dd = fmaxf(exact_max_dd, dd); + } + int out_base = w * 14; metrics_out[out_base + 0] = (mean / std) * annualization_factor; // Sharpe // Total return: (final_equity / initial_equity) - 1.0 as fraction metrics_out[out_base + 1] = s_sorted[0] - 1.0f; - metrics_out[out_base + 2] = s_max_dd[0]; // max drawdown (fractional, reduced across threads) + metrics_out[out_base + 2] = exact_max_dd; // max drawdown (exact sequential scan) metrics_out[out_base + 3] = (mean / down_std) * annualization_factor; // Sortino - // Win rate: winning trades / total trades (not per-bar) - float total_trades_f = s_trades[0]; - metrics_out[out_base + 4] = (total_trades_f > 0.0f) ? s_wins[0] / total_trades_f : 0.0f; - metrics_out[out_base + 5] = total_trades_f; // trade count (reduced) - 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 + // Trade count and win rate: placeholder zeros (restored by boundary stitching in next commit) + metrics_out[out_base + 4] = 0.0f; // win_rate + metrics_out[out_base + 5] = 0.0f; // total_trades + metrics_out[out_base + 10] = s_buys[0]; + metrics_out[out_base + 11] = s_sells[0]; + metrics_out[out_base + 12] = s_holds[0]; } // ── Unique action count via bitmask OR-reduction ───────────────────────── @@ -222,9 +202,14 @@ extern "C" __global__ void compute_backtest_metrics( int sort_len = wlen < 4096 ? wlen : 4096; - // Load returns into sort scratch (stride-strided load) + // Load returns into sort scratch. + // For windows > 4096 bars, use strided sampling to cover the full window + // evenly instead of only the first 4096 bars (which biases VaR/CVaR + // toward early-window volatility characteristics). + int sample_step = (wlen + sort_len - 1) / sort_len; // ceil(wlen / sort_len) for (int i = tid; i < sort_len; i += stride) { - s_sorted[i] = step_returns[base + i]; + int src = i * sample_step; + s_sorted[i] = (src < wlen) ? step_returns[base + src] : 0.0f; } // Compute next power-of-two for bitonic sort diff --git a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs index e85cc8dea..68bda2efa 100644 --- a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs +++ b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs @@ -238,6 +238,10 @@ pub struct GpuBacktestConfig { /// Typical value: 8 (VPIN, Kyle's Lambda, OFI L1/L5, depth imbalance, bid/ask slope, /// trade imbalance). Default: 0 (no OFI). pub ofi_dim: usize, + /// Number of bars in one trading day (390 for 1-minute, 1 for daily). + /// Used to compute the correct Sharpe annualization: `sqrt(bars_per_day * 252)`. + /// Default: 390.0 (1-minute bar frequency). + pub bars_per_day: f32, } impl Default for GpuBacktestConfig { @@ -250,6 +254,7 @@ impl Default for GpuBacktestConfig { contract_multiplier: 50.0, max_leverage: 2.0, ofi_dim: 0, + bars_per_day: 390.0, // 1-minute bar frequency (6.5h × 60min) } } } @@ -317,6 +322,8 @@ pub struct GpuBacktestEvaluator { /// Aligned state dimension (feature_dim + portfolio_dim, 8-aligned for tensor cores). state_dim: usize, config: GpuBacktestConfig, + /// Precomputed annualization factor: sqrt(bars_per_day * 252). + annualization_factor: f32, /// Actions buffer for forward kernel output (reused across steps). forward_actions_buf: CudaSlice, @@ -629,6 +636,7 @@ impl GpuBacktestEvaluator { feature_dim, portfolio_dim: PORTFOLIO_AND_MTF_DIM, state_dim, + annualization_factor: (config.bars_per_day * 252.0).sqrt(), config, forward_actions_buf, cublas_forward: None, @@ -1784,10 +1792,10 @@ impl GpuBacktestEvaluator { /// /// Shared by all evaluation paths. fn launch_metrics_and_download(&self) -> Result, MLError> { - // 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; + // 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; let metrics_cfg = LaunchConfig { grid_dim: (self.n_windows as u32, 1, 1), block_dim: (256, 1, 1), @@ -1795,7 +1803,7 @@ impl GpuBacktestEvaluator { }; let n_win_i32 = self.n_windows as i32; let max_len_i32 = self.max_len as i32; - let annualization: f32 = 252.0_f32.sqrt(); + let annualization: f32 = self.annualization_factor; // Safety: argument order matches `compute_backtest_metrics` signature exactly. unsafe { @@ -1945,6 +1953,18 @@ mod tests { assert_eq!(c.tx_cost_bps, 0.1); assert!((c.spread_cost - 0.0001).abs() < 1e-8); assert_eq!(c.initial_capital, 100_000.0); + // 1-minute bar frequency: 390 bars/day (6.5h × 60min) + assert_eq!(c.bars_per_day, 390.0); + } + + #[test] + fn test_annualization_factor_1min_bars() { + let c = GpuBacktestConfig::default(); + // For 1-minute bars: sqrt(390 * 252) = sqrt(98280) ≈ 313.5 + let expected = (390.0_f32 * 252.0).sqrt(); + let actual = (c.bars_per_day * 252.0).sqrt(); + assert!((actual - expected).abs() < 0.01, "annualization: {actual} vs {expected}"); + assert!(actual > 300.0 && actual < 320.0, "annualization should be ~313: {actual}"); } #[test]