From de1dbba91038d01347ded5c6ee2b334c3ac1a10f Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Wed, 11 Mar 2026 10:57:52 +0100 Subject: [PATCH] feat(cuda): add per-window backtest metrics reduction kernel One block per window. Parallel reduction for Sharpe, Sortino, total PnL, max drawdown, win rate, trade count. Single kernel launch reduces all windows simultaneously. --- .../cuda_pipeline/backtest_metrics_kernel.cu | 106 ++++++++++++++++++ 1 file changed, 106 insertions(+) create mode 100644 crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu diff --git a/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu new file mode 100644 index 000000000..e985fcc24 --- /dev/null +++ b/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu @@ -0,0 +1,106 @@ +// Per-window metrics reduction kernel. +// One block per window. Threads cooperate to reduce step_returns. +// +// Output per window [6 floats]: +// [0] sharpe_ratio (annualized, sqrt(252)) +// [1] total_pnl (cumulative return) +// [2] max_drawdown (worst peak-to-trough, positive number) +// [3] sortino_ratio +// [4] win_rate +// [5] total_trades (approximated from position changes) + +extern "C" __global__ void compute_backtest_metrics( + const float* __restrict__ step_returns, // [n_windows * max_len] + const float* __restrict__ portfolio_state, // [n_windows * 8] + const int* __restrict__ window_lens, // [n_windows] + const int* __restrict__ actions_history, // [n_windows * max_len] for trade counting + float* metrics_out, // [n_windows * 6] + int n_windows, + int max_len, + float annualization_factor // sqrt(252) for daily +) { + int w = blockIdx.x; + if (w >= n_windows) return; + + int wlen = window_lens[w]; + int tid = threadIdx.x; + int stride = blockDim.x; + int base = w * max_len; + + // Shared memory for parallel reduction — 6 arrays + 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] + + // Pass 1: per-thread local accumulators + float local_sum = 0.0f, local_sq = 0.0f, local_down = 0.0f; + float local_cum = 0.0f, local_peak = 0.0f, local_max_dd = 0.0f; + int local_wins = 0, local_trades = 0; + int prev_action = -1; + + for (int i = tid; i < wlen; i += stride) { + float r = step_returns[base + i]; + local_sum += r; + local_sq += r * r; + if (r < 0.0f) local_down += r * r; + + // Drawdown tracking (NOTE: strided — approximate per thread, + // then take max across threads for worst-case estimate) + local_cum += r; + local_peak = fmaxf(local_peak, local_cum); + float dd = local_peak - local_cum; + local_max_dd = fmaxf(local_max_dd, dd); + + // Win/loss counting + if (r > 0.0f) local_wins++; + + // Trade counting (position changes) + int act = actions_history[base + i]; + if (act != prev_action && i > 0) local_trades++; + prev_action = act; + } + + // Store ALL local values to shared memory + 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; + __syncthreads(); + + // Block-level parallel reduction for ALL 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]; + } + __syncthreads(); + } + + // Thread 0 computes final metrics from fully reduced values + if (tid == 0) { + float n = (float)wlen; + float mean = s_sum[0] / n; + float var = s_sq_sum[0] / n - mean * mean; + float std = sqrtf(fmaxf(var, 1e-10f)); + float down_std = sqrtf(fmaxf(s_down_sq[0] / n, 1e-10f)); + + int out_base = w * 6; + metrics_out[out_base + 0] = (mean / std) * annualization_factor; // Sharpe + metrics_out[out_base + 1] = s_sum[0]; // total cumulative return + metrics_out[out_base + 2] = s_max_dd[0]; // max drawdown (reduced across all threads) + metrics_out[out_base + 3] = (mean / down_std) * annualization_factor; // Sortino + metrics_out[out_base + 4] = (n > 0.0f) ? s_wins[0] / n : 0.0f; // win rate (reduced) + metrics_out[out_base + 5] = s_trades[0]; // trade count (reduced) + } +}