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

fix(cuda): GPU-counted unique_actions via bitmask OR-reduction in metrics kernel

Browse Source 
unique_actions was hardcoded to 5 in the GPU backtest path, making the
diversity penalty (0.8 max weight) always return 0.0 — another dead
objective component. CPU backtest path correctly computed it via HashSet.

Added 5-bit bitmask OR-reduction: each thread sets bit(action_id) during
the existing per-step loop, then a single OR-reduction across the block
produces the union. __popc() gives unique count in one PTX instruction.

Memory efficiency: reuses s_sorted shared memory (sequential staging —
OR-reduction completes before bitonic sort overwrites it). No extra
shared memory arrays needed. Output expanded from 13→14 floats/window.

Combined with the previous commit, this restores gradient signal for
100% of the multi-objective function: composite (60%), HFT activity
(25%), stability (15%), and diversity penalty (soft signal).

Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>
This commit is contained in:
jgrusewski
2026-03-13 03:11:28 +01:00
parent e910f6dfa8
commit 3d04c24eaf
3 changed files with 50 additions and 12 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -15,13 +15,14 @@
// [10] buy_count (actions 3,4 — Long50, Long100)
// [11] sell_count (actions 0,1 — Short100, Short50)
// [12] hold_count (action 2 — Flat)
// [13] unique_action_count (popcount of 5-bit mask — how many of 0..4 were used)
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 * 13]
float* metrics_out, // [n_windows * 14]
int n_windows,
int max_len,
float annualization_factor // sqrt(252) for daily
@@ -64,6 +65,7 @@ extern "C" __global__ void compute_backtest_metrics(
float local_cum = 0.0f, local_peak = 0.0f, local_max_dd = 0.0f;
int local_wins = 0, local_trades = 0;
int local_buys = 0, local_sells = 0, local_holds = 0;
int local_action_mask = 0; // 5-bit bitmask: bit i set if action i was seen
int prev_action = -1;
for (int i = tid; i < wlen; i += stride) {
@@ -91,6 +93,9 @@ extern "C" __global__ void compute_backtest_metrics(
if (act <= 1) local_sells++;
else if (act == 2) local_holds++;
else local_buys++;
// Unique action tracking: set bit for each distinct action ID (0..4)
if (act >= 0 && act < 5) local_action_mask |= (1 << act);
}
// Store ALL local values to shared memory
@@ -129,7 +134,7 @@ 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));
int out_base = w * 13;
int out_base = w * 14;
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)
@@ -141,6 +146,28 @@ extern "C" __global__ void compute_backtest_metrics(
metrics_out[out_base + 12] = s_holds[0]; // hold action count
}
// ── Unique action count via bitmask OR-reduction ─────────────────────────
// Reuses s_sorted memory (safe: OR-reduction completes before bitonic sort).
// Each thread accumulated a 5-bit mask of observed action IDs. OR-reduce
// across the block, then popcount gives unique_action_count (1-5).
{
int* s_action_mask = (int*)s_sorted; // alias — sequential use, no conflict
s_action_mask[tid] = local_action_mask;
__syncthreads();
for (int s = stride / 2; s > 0; s >>= 1) {
if (tid < s) {
s_action_mask[tid] |= s_action_mask[tid + s];
}
__syncthreads();
}
if (tid == 0) {
int out_base = w * 14;
metrics_out[out_base + 13] = (float)__popc(s_action_mask[0]);
}
}
// ── Extended metrics: VaR, CVaR, Calmar, Omega via bitonic sort ──────────
//
// All threads cooperate to load and sort up to 4096 step_returns for this
@@ -185,7 +212,7 @@ extern "C" __global__ void compute_backtest_metrics(
// Thread 0 computes extended metrics from the sorted array
if (tid == 0) {
int out_base = w * 13;
int out_base = w * 14;
// VaR at 95% confidence = 5th-percentile of sorted returns
int var_idx = (int)(0.05f * (float)sort_len);

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

@@ -131,6 +131,8 @@ pub struct WindowMetrics {
pub buy_count: f32,
pub sell_count: f32,
pub hold_count: f32,
// Unique action IDs observed (popcount of 5-bit mask, index 13)
pub unique_actions: f32,
}
/// Configuration for the GPU backtest evaluator.
@@ -232,7 +234,7 @@ pub struct GpuBacktestEvaluator {
states_buf: CudaSlice<f32>, // [n_windows * state_dim] (8-aligned)
// Output buffer (written by metrics kernel)
metrics_buf: CudaSlice<f32>, // [n_windows * 13]
metrics_buf: CudaSlice<f32>, // [n_windows * 14]
// Dimensions and config
n_windows: usize,
@@ -432,7 +434,7 @@ impl GpuBacktestEvaluator {
.alloc_zeros::<i32>(n_windows * max_len)
.map_err(|e| MLError::ModelError(format!("actions_history alloc: {e}")))?;
let metrics_buf = stream
.alloc_zeros::<f32>(n_windows * 13)
.alloc_zeros::<f32>(n_windows * 14)
.map_err(|e| MLError::ModelError(format!("metrics alloc: {e}")))?;
// Portfolio dimension is always 3: (normalised value, position, spread_cost).
@@ -1457,20 +1459,20 @@ impl GpuBacktestEvaluator {
.map_err(|e| MLError::ModelError(format!("compute_backtest_metrics launch: {e}")))?;
}
// Single download: n_windows × 13 floats (the ONLY GPU→CPU transfer)
let mut metrics_host = vec![0.0_f32; self.n_windows * 13];
// Single download: n_windows × 14 floats (the ONLY GPU→CPU transfer)
let mut metrics_host = vec![0.0_f32; self.n_windows * 14];
self.stream
.memcpy_dtoh(&self.metrics_buf, &mut metrics_host)
.map_err(|e| MLError::ModelError(format!("metrics download: {e}")))?;
// Parse flat metrics into per-window structs.
// Layout matches compute_backtest_metrics kernel output (13 floats per window):
// Layout matches compute_backtest_metrics kernel output (14 floats per window):
// [0] sharpe, [1] total_pnl, [2] max_drawdown, [3] sortino, [4] win_rate,
// [5] total_trades, [6] var_95, [7] cvar_95, [8] calmar, [9] omega_ratio,
// [10] buy_count, [11] sell_count, [12] hold_count
// [10] buy_count, [11] sell_count, [12] hold_count, [13] unique_actions
let results: Vec<WindowMetrics> = (0..self.n_windows)
.map(|w| {
let base = w * 13;
let base = w * 14;
WindowMetrics {
sharpe: metrics_host[base],
total_pnl: metrics_host[base + 1],
@@ -1485,6 +1487,7 @@ impl GpuBacktestEvaluator {
buy_count: metrics_host[base + 10],
sell_count: metrics_host[base + 11],
hold_count: metrics_host[base + 12],
unique_actions: metrics_host[base + 13],
}
})
.collect();
@@ -1532,6 +1535,7 @@ mod tests {
buy_count: 150.0,
sell_count: 120.0,
hold_count: 230.0,
unique_actions: 4.0,
};
assert!(m.sharpe > 0.0);
assert!(m.win_rate > 0.5);

11
crates/ml/src/hyperopt/adapters/dqn.rs
View File

@@ -1865,12 +1865,19 @@ impl DQNTrainer {
let mean_cvar_95 = metrics.iter().map(|m| m.cvar_95 as f64).sum::<f64>() / n;
let mean_omega = metrics.iter().map(|m| m.omega_ratio as f64).sum::<f64>() / n;
// Action distribution from GPU kernel (indices 10-12) — summed across all windows
// Action distribution from GPU kernel (indices 10-13) — summed across all windows
let total_buys = metrics.iter().map(|m| m.buy_count as f64).sum::<f64>();
let total_sells = metrics.iter().map(|m| m.sell_count as f64).sum::<f64>();
let total_holds = metrics.iter().map(|m| m.hold_count as f64).sum::<f64>();
let action_denom = (total_buys + total_sells + total_holds).max(1.0);
// Unique actions: max across windows (union of observed action IDs)
let unique_actions = metrics
.iter()
.map(|m| m.unique_actions as usize)
.max()
.unwrap_or(1);
Ok(Some(BacktestMetrics {
sharpe_ratio: mean_sharpe,
total_return_pct: mean_pnl * 100.0,
@@ -1889,7 +1896,7 @@ impl DQNTrainer {
alpha: 0.0,
information_ratio: 0.0,
omega_ratio: mean_omega,
unique_actions: 5,
unique_actions,
buy_action_pct: total_buys / action_denom,
sell_action_pct: total_sells / action_denom,
hold_action_pct: total_holds / action_denom,