gitadmin/foxhunt
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

refactor: remove 3 dead reduction arrays + per-thread drawdown vars from metrics kernel

Browse Source 
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.
This commit is contained in:
jgrusewski
2026-03-25 19:25:10 +01:00
parent fb5d6a57a2
commit b6038ca709
2 changed files with 79 additions and 74 deletions
Download Patch File Download Diff File Expand all files Collapse all files

123
crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu
View File

@@ -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

30
crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs
View File

@@ -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<i32>,
@@ -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<Vec<WindowMetrics>, 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::<f32>() 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::<f32>() 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]