# Financial Metrics Pipeline Fix — Implementation Plan

> **For agentic workers:** REQUIRED SUB-SKILL: Use superpowers:subagent-driven-development (recommended) or superpowers:executing-plans to implement this plan task-by-task. Steps use checkbox (`- [ ]`) syntax for tracking.

**Goal:** Make Sharpe/WinRate/MaxDD/Trades report real per-trade performance from the GPU experience collector's portfolio simulation, not epoch-averaged mean rewards.

**Architecture:** Add a GPU reduction kernel that reads per-episode Kelly stats from `portfolio_states[N, 14:19]` (win_count, loss_count, sum_wins, sum_losses, sum_returns, sum_sq_returns) and reduces them into a single epoch-level `TradeStats` struct. Download this 48-byte summary once per epoch. Feed real trade-level data to `compute_epoch_financials()`.

**Tech Stack:** CUDA kernel (inline in gpu_experience_collector.rs), Rust structs, existing monitoring download path.

---

## Root Cause

```
CURRENT (broken):
  experience_env_step → out_rewards[N,L] → monitoring_reduce → mean(N*L rewards) → ONE float per epoch
  pnl_history gets 1 entry per epoch → Sharpe=0.00, Trades=1

FIXED:
  experience_env_step → portfolio_states[N, 14:19] → trade_stats_reduce → TradeStats per epoch
  TradeStats has: total_trades, wins, losses, sum_pnl, sum_wins, sum_losses, per-step returns
  compute_epoch_financials uses REAL trade data → meaningful Sharpe/WinRate/Trades
```

The per-trade Kelly stats already exist in `portfolio_states` — the experience kernel populates `ps[14]` (win_count), `ps[15]` (loss_count), `ps[16]` (sum_wins), `ps[17]` (sum_losses), `ps[18]` (sum_returns), `ps[19]` (sum_sq_returns) at every trade completion event. These are never read.

## File Structure

| File | Action | Responsibility |
|------|--------|---------------|
| `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs` | Modify | Add `trade_stats_reduce` kernel, `TradeStats` struct, `collect_trade_stats()` method |
| `crates/ml/src/trainers/dqn/financials.rs` | Modify | Rewrite `compute_epoch_financials()` to accept `TradeStats` instead of `VecDeque<f64>` |
| `crates/ml/src/trainers/dqn/trainer/training_loop.rs` | Modify | Call `collect_trade_stats()` after experience collection, feed to financials |
| `crates/ml/src/trainers/dqn/trainer/mod.rs` | Modify | Add `trade_stats_history: Vec<TradeStats>` field |
| `crates/ml/src/trainers/dqn/trainer/metrics.rs` | Modify | Use TradeStats in final metrics |
| `crates/ml/src/trainers/dqn/smoke_tests/training_stability.rs` | Modify | Assert Trades > 1, WinRate > 0 |

---

## Task 1: TradeStats Struct + GPU Reduction Kernel (20 min)

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs`

The kernel reads `portfolio_states[N, PORTFOLIO_STRIDE]` and sums the per-episode Kelly fields into a single output vector.

- [ ] **Step 1:** Add `TradeStats` struct to `gpu_experience_collector.rs`:
```rust
/// Per-epoch trading statistics reduced from GPU portfolio states.
/// Downloaded once per epoch (48 bytes — 6 f64s).
#[derive(Debug, Clone, Default)]
pub struct TradeStats {
    /// Total completed trades (entries + exits) across all episodes
    pub total_trades: usize,
    /// Winning trades (positive trade return at exit)
    pub winning_trades: usize,
    /// Losing trades (negative trade return at exit)
    pub losing_trades: usize,
    /// Sum of P&L from winning trades (for profit factor)
    pub sum_wins: f64,
    /// Sum of |P&L| from losing trades (for profit factor)
    pub sum_losses: f64,
    /// Sum of all trade returns (for mean return)
    pub sum_returns: f64,
    /// Sum of squared trade returns (for Sharpe std computation)
    pub sum_sq_returns: f64,
    /// Per-step returns from all episodes (for drawdown/equity curve)
    /// These are the actual out_rewards[N*L] values, not averaged.
    pub step_returns: Vec<f64>,
}
```

- [ ] **Step 2:** Add inline CUDA kernel `trade_stats_reduce`:
```cuda
// Grid: 1 block, Block: 256 threads
// Parallel reduction over N episodes, summing portfolio_states[i, 14:19]
extern "C" __global__ void trade_stats_reduce(
    const float* __restrict__ portfolio_states,  // [N, PORTFOLIO_STRIDE]
    float* __restrict__ output,                   // [6] win_count, loss_count, sum_wins, sum_losses, sum_returns, sum_sq
    int N,
    int stride
) {
    __shared__ float shmem[6 * 32]; // 6 fields × 32 warps max
    int tid = threadIdx.x;

    // Each thread accumulates over a range of episodes
    float local[6] = {0};
    for (int i = tid; i < N; i += blockDim.x) {
        const float* ps = portfolio_states + i * stride;
        local[0] += ps[14]; // win_count
        local[1] += ps[15]; // loss_count
        local[2] += ps[16]; // sum_wins
        local[3] += ps[17]; // sum_losses
        local[4] += ps[18]; // sum_returns
        local[5] += ps[19]; // sum_sq_returns
    }

    // Warp-level reduction
    for (int offset = 16; offset > 0; offset >>= 1) {
        for (int f = 0; f < 6; f++)
            local[f] += __shfl_down_sync(0xffffffff, local[f], offset);
    }

    // Write warp results to shared memory
    int warp_id = tid / 32;
    int lane = tid % 32;
    if (lane == 0) {
        for (int f = 0; f < 6; f++)
            shmem[warp_id * 6 + f] = local[f];
    }
    __syncthreads();

    // Final reduction by first warp
    if (warp_id == 0) {
        int num_warps = (blockDim.x + 31) / 32;
        float sum[6] = {0};
        if (lane < num_warps) {
            for (int f = 0; f < 6; f++)
                sum[f] = shmem[lane * 6 + f];
        }
        for (int offset = 16; offset > 0; offset >>= 1) {
            for (int f = 0; f < 6; f++)
                sum[f] += __shfl_down_sync(0xffffffff, sum[f], offset);
        }
        if (lane == 0) {
            for (int f = 0; f < 6; f++)
                output[f] = sum[f];
        }
    }
}
```

- [ ] **Step 3:** Add `trade_stats_kernel: CudaFunction` field to `GpuExperienceCollector`. Compile the kernel in the constructor alongside the existing kernels. Add a `trade_stats_buf: CudaSlice<f32>` output buffer (6 floats).

- [ ] **Step 4:** Add `pub fn collect_trade_stats(&mut self) -> Result<TradeStats, MLError>` method:
  - Launch the `trade_stats_reduce` kernel with `portfolio_states`, N=alloc_episodes, stride=PORTFOLIO_STRIDE
  - Download the 6-float output buffer to host (24 bytes — single DtoH transfer)
  - Also download `rewards_out[N*L]` for per-step returns (this IS a larger transfer but only once per epoch, ~12KB for 3200 experiences)
  - Construct and return `TradeStats`

- [ ] **Step 5:** Run: `SQLX_OFFLINE=true cargo check -p ml` — verify clean build.

- [ ] **Step 6:** Commit: "feat: GPU trade stats reduction kernel — extracts per-trade metrics from portfolio states"

---

## Task 2: Rewrite compute_epoch_financials() (10 min)

**Files:**
- Modify: `crates/ml/src/trainers/dqn/financials.rs`

Replace the `pnl_history: &VecDeque<f64>` input with `TradeStats`. The financial metrics should use REAL trade data:

- [ ] **Step 1:** Change signature:
```rust
pub(crate) fn compute_epoch_financials(
    trade_stats: &TradeStats,
    action_counts: &[usize; 9],
    initial_capital: f64,
) -> EpochFinancials
```

- [ ] **Step 2:** Rewrite the body:
  - `total_trades = trade_stats.total_trades` (real trade count from GPU)
  - `win_rate = trade_stats.winning_trades as f64 / total_trades.max(1) as f64`
  - `profit_factor = trade_stats.sum_wins / trade_stats.sum_losses.max(1e-10)`
  - `total_return = trade_stats.sum_returns` (sum of all trade returns)
  - `avg_return = total_return / total_trades.max(1) as f64`
  - Sharpe: compute from `step_returns` (per-bar returns, annualized with √252)
  - Max drawdown: compute equity curve from `step_returns` starting at `initial_capital`
  - Sortino: from negative step_returns only
  - Action distribution: unchanged (from `action_counts`)

- [ ] **Step 3:** Update the existing tests to use `TradeStats` input:
  - `test_empty_history` → `TradeStats::default()` should produce zeros
  - `test_all_winning` → `TradeStats { winning_trades: 5, total_trades: 5, sum_wins: 100.0, ... }`
  - `test_mixed_pnl` → set win/loss counts and sums from trade data

- [ ] **Step 4:** Run: `SQLX_OFFLINE=true cargo test -p ml --lib -- financials` — all tests pass.

- [ ] **Step 5:** Commit: "refactor: compute_epoch_financials uses real TradeStats from GPU"

---

## Task 3: Wire TradeStats into Training Loop (10 min)

**Files:**
- Modify: `crates/ml/src/trainers/dqn/trainer/mod.rs`
- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs`
- Modify: `crates/ml/src/trainers/dqn/trainer/metrics.rs`

- [ ] **Step 1:** Add `trade_stats_history: Vec<TradeStats>` to `DQNTrainer` struct in `mod.rs`. Initialize to `Vec::new()` in constructor.

- [ ] **Step 2:** In `training_loop.rs`, after GPU experience collection (where `summary.mean_reward` is currently pushed to `pnl_history`), add:
```rust
// Collect real per-trade stats from GPU portfolio states
if let Some(ref mut collector) = self.gpu_experience_collector {
    match collector.collect_trade_stats() {
        Ok(stats) => {
            info!(
                "Epoch {}/{}: GPU trades={} wins={} losses={} PF={:.2}",
                epoch + 1, self.hyperparams.epochs,
                stats.total_trades, stats.winning_trades, stats.losing_trades,
                if stats.sum_losses > 1e-10 { stats.sum_wins / stats.sum_losses } else { 0.0 }
            );
            self.trade_stats_history.push(stats);
        }
        Err(e) => warn!("Failed to collect trade stats: {e}"),
    }
}
```

- [ ] **Step 3:** Replace the `compute_epoch_financials(&self.pnl_history, ...)` call with:
```rust
let trade_stats = self.trade_stats_history.last()
    .cloned()
    .unwrap_or_default();
let financials = compute_epoch_financials(
    &trade_stats,
    &monitor.action_counts,
    self.hyperparams.initial_capital,
);
```

- [ ] **Step 4:** In `metrics.rs`, update `create_final_metrics()` to aggregate trade stats across all epochs:
  - Total trades = sum of all epoch trade counts
  - Overall win rate = total wins / total trades
  - Overall Sharpe = from concatenated step_returns across all epochs

- [ ] **Step 5:** Remove the `pnl_history: VecDeque<f64>` field and all references. It's replaced by `trade_stats_history`.

- [ ] **Step 6:** Run: `SQLX_OFFLINE=true cargo check -p ml` — verify clean build.

- [ ] **Step 7:** Commit: "feat: wire GPU trade stats into training loop financials"

---

## Task 4: Update Smoke Tests to Assert Real Trading Metrics (5 min)

**Files:**
- Modify: `crates/ml/src/trainers/dqn/smoke_tests/training_stability.rs`

- [ ] **Step 1:** In `test_trading_model_behavior`, add:
```rust
// With 3200 experiences per epoch and position changes, we should see trades
let total_trades = metrics.additional_metrics.get("total_trades").copied().unwrap_or(0.0);
assert!(
    total_trades > 0.0,
    "BEHAVIORAL: Model must execute trades during experience collection. Got 0 trades."
);
```

- [ ] **Step 2:** In `test_50_epoch_convergence`, add:
```rust
// After 50 epochs × 3200 experiences, model must have executed real trades
let total_trades = metrics.additional_metrics.get("total_trades").copied().unwrap_or(0.0);
assert!(
    total_trades > 10.0,
    "CONVERGENCE: Model must execute >10 trades over 50 epochs. Got {total_trades}"
);
```

- [ ] **Step 3:** Run all smoke tests:
```bash
FOXHUNT_TEST_DATA=test_data/futures-baseline SQLX_OFFLINE=true cargo test -p ml --lib -- smoke_tests --ignored --nocapture
```

- [ ] **Step 4:** Commit: "test: assert real trade count in smoke tests"

---

## Key Design Decisions

### Why Reduce on GPU Instead of Downloading portfolio_states?

`portfolio_states` is `[N_episodes × PORTFOLIO_STRIDE=20]` = `N × 80 bytes`. With N=32, that's 2560 bytes — small enough to download. But the reduction kernel pattern is:
1. Future-proof: when N=1024 on H100, downloading 80KB per epoch wastes PCIe bandwidth
2. Consistent: follows the existing `monitoring_reduce` pattern
3. Zero-copy: the portfolio_states buffer is already on GPU, no staging needed

### Why Download step_returns (out_rewards)?

For Sharpe/MaxDD/Sortino, we need the per-step return distribution, not just the aggregates. With 3200 experiences, that's 12.8KB per epoch — acceptable for a once-per-epoch download. The alternative (computing Sharpe entirely on GPU) would require a second reduction kernel for mean/variance/drawdown, which is overengineering for a diagnostic metric.

### What Counts as a "Trade"?

A trade is an entry→exit cycle detected by the experience kernel's trade lifecycle logic. When `exiting_trade` fires (position goes from non-zero to zero), `win_count` or `loss_count` increments. Each episode can have 0-N trades depending on market conditions and model behavior. `total_trades = win_count + loss_count` summed across all episodes.
