From 3d04c24eaf1f2b2eaf60bc329ee0fd38cf27ef1d Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Fri, 13 Mar 2026 03:11:28 +0100 Subject: [PATCH] fix(cuda): GPU-counted unique_actions via bitmask OR-reduction in metrics kernel MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit 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 --- .../cuda_pipeline/backtest_metrics_kernel.cu | 33 +++++++++++++++++-- .../cuda_pipeline/gpu_backtest_evaluator.rs | 18 ++++++---- crates/ml/src/hyperopt/adapters/dqn.rs | 11 +++++-- 3 files changed, 50 insertions(+), 12 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu b/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu index 075b68906..193b70cc4 100644 --- a/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu +++ b/crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu @@ -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); diff --git a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs index c476c58e1..54113baab 100644 --- a/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs +++ b/crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs @@ -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, // [n_windows * state_dim] (8-aligned) // Output buffer (written by metrics kernel) - metrics_buf: CudaSlice, // [n_windows * 13] + metrics_buf: CudaSlice, // [n_windows * 14] // Dimensions and config n_windows: usize, @@ -432,7 +434,7 @@ impl GpuBacktestEvaluator { .alloc_zeros::(n_windows * max_len) .map_err(|e| MLError::ModelError(format!("actions_history alloc: {e}")))?; let metrics_buf = stream - .alloc_zeros::(n_windows * 13) + .alloc_zeros::(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 = (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); diff --git a/crates/ml/src/hyperopt/adapters/dqn.rs b/crates/ml/src/hyperopt/adapters/dqn.rs index 208f89a30..394236542 100644 --- a/crates/ml/src/hyperopt/adapters/dqn.rs +++ b/crates/ml/src/hyperopt/adapters/dqn.rs @@ -1865,12 +1865,19 @@ impl DQNTrainer { let mean_cvar_95 = metrics.iter().map(|m| m.cvar_95 as f64).sum::() / n; let mean_omega = metrics.iter().map(|m| m.omega_ratio as f64).sum::() / 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::(); let total_sells = metrics.iter().map(|m| m.sell_count as f64).sum::(); let total_holds = metrics.iter().map(|m| m.hold_count as f64).sum::(); 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,