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.

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)
+    }
+}