# Boundary-Aware Parallel Trade Counting — Design Spec

## Problem

The GPU backtest metrics kernel (`backtest_metrics_kernel.cu`) uses 256 threads to process consecutive chunks of a 10K-bar evaluation window. Trades that span chunk boundaries are fragmented: each thread tracks trades independently, so a single trade held across 3 chunks gets split into 3 partial segments. The cumulative return of the fragmented portions is lost, producing incorrect trade counts and corrupted win rates.

Additionally, `coordinator_extended.rs` uses a correct but non-parameterized annualization formula (`(252.0 * 6.5 * 60.0).sqrt()`). This should use a named constant for consistency with the rest of the codebase.

## Design: Hybrid Parallel + Boundary Stitching

### Architecture

**Phase 1 — Parallel per-thread pass (all 256 threads):**
Each thread processes its consecutive chunk of bars and computes:
- Standard accumulators (sum, sq_sum, downside, equity, drawdown, action distribution) — unchanged
- **NEW**: Boundary metadata instead of local trade counting

**Phase 2 — Boundary export (all threads → shared memory):**
Each thread writes 7 boundary values to `s_sorted[stride..8*stride]` (reuses existing sort scratch — zero extra shared memory).

**Phase 3 — Stitching (thread 0 only, O(256)):**
Thread 0 walks boundary metadata left-to-right, merging partial trades across chunk boundaries, accumulating exact per-trade returns, and counting trades/wins.

### Per-Thread Boundary Metadata (7 values per thread)

| Field | Type | Description |
|-------|------|-------------|
| `first_action` | int | Action ID at chunk's first bar (-1 if chunk is empty) |
| `last_action` | int | Action ID at chunk's last bar |
| `prefix_return` | float | Cumulative return from chunk start → first action change (or total chunk return if no changes) |
| `suffix_return` | float | Cumulative return from last action change → chunk end (or total chunk return if no changes) |
| `complete_trades` | int | Trades fully contained within chunk (between first and last change) |
| `complete_wins` | int | Winning interior trades (complete_return > 0) |
| `num_changes` | int | Number of action changes in chunk |

**Edge case — no action changes:** When `num_changes == 0`, the entire chunk is one partial trade segment. `prefix_return == suffix_return == total chunk return`. The stitching pass merges it with adjacent chunks.

**Edge case — empty window (`wlen == 0`):** All threads get empty chunks (`chunk_start >= chunk_end`), so `first_action == -1`. The stitching pass skips all threads, closes no trades, and produces `total_trades = 0`, `win_rate = 0.0`.

### Per-Thread Loop Code

The old trade-counting variables (`prev_action`, `trade_cum_return`, `local_trades`, `local_wins`) are **replaced** by the 7 boundary-tracking variables:

```cuda
// NEW: boundary-aware trade tracking variables
int bnd_first_action = -1;     // -1 signals empty chunk
int bnd_last_action = -1;
float bnd_prefix_return = 0.0f;
float bnd_suffix_return = 0.0f;
int bnd_complete_trades = 0;
int bnd_complete_wins = 0;
int bnd_num_changes = 0;
int bnd_cur_action = -1;
float bnd_cur_return = 0.0f;

for (int i = chunk_start; i < chunk_end; i++) {
    float r = step_returns[base + i];
    int act = actions_history[base + i];

    // ... (existing Sharpe/drawdown/action-distribution accumulators unchanged) ...

    // Boundary-aware trade tracking
    if (i == chunk_start) {
        bnd_first_action = act;
        bnd_cur_action = act;
    }

    if (act != bnd_cur_action) {
        bnd_num_changes++;
        if (bnd_num_changes == 1) {
            // First change: everything accumulated so far is the prefix
            bnd_prefix_return = bnd_cur_return;
        } else {
            // Interior change: the segment since the last change is a complete trade
            bnd_complete_trades++;
            if (bnd_cur_return > 0.0f) bnd_complete_wins++;
        }
        bnd_cur_return = 0.0f;
        bnd_cur_action = act;
    }

    bnd_cur_return += r;
    bnd_last_action = act;
}

// Post-loop: assign suffix and handle no-change case
bnd_suffix_return = bnd_cur_return;
if (bnd_num_changes == 0) {
    // Entire chunk is one partial trade segment
    bnd_prefix_return = bnd_cur_return;
}
```

### Stitching Algorithm

The stitching matches the existing kernel's trade semantics: a "trade" is a contiguous run of bars with the same action. The first bar of the window **opens** the first trade; it is **closed** only when the action changes or the window ends. This matches the old `if (act != prev_action && prev_action >= 0)` logic — the initial position is not counted as a closure until it changes.

```
total_trades = sum(complete_trades[t] for all threads)
total_wins = sum(complete_wins[t] for all threads)
open_return = 0.0
open_action = -1

for each thread t (0..blockDim.x):
    fa = first_action[t]
    la = last_action[t]
    pr = prefix_return[t]
    sr = suffix_return[t]
    nc = num_changes[t]

    if fa < 0: skip (empty chunk)

    // Step 1: connect this chunk's start to the open trade
    if open_action < 0:
        // First valid chunk — open the first trade (no closure yet)
        open_action = fa
        open_return = 0.0
    else if fa != open_action:
        // Action changed at chunk boundary — close the open trade
        total_trades++
        if open_return > 0.0: total_wins++
        open_action = fa
        open_return = 0.0

    // Step 2: process this chunk's contribution
    if nc == 0:
        // No action changes — entire chunk extends the open trade
        open_return += pr   // (pr == sr == total chunk return)
    else:
        // Prefix extends and closes the open trade (first intra-chunk change ends it)
        open_return += pr
        total_trades++
        if open_return > 0.0: total_wins++
        // Interior trades already counted by per-thread code
        // Suffix becomes the new open trade
        open_action = la
        open_return = sr

// Close the final open trade (window end)
if open_action >= 0:
    total_trades++
    if open_return > 0.0: total_wins++
```

**Key properties:**
- `open_return` is always reset to 0.0 when `open_action` changes (no return leakage across trades)
- The first chunk's prefix opens a trade but does not close it (matching the old `prev_action >= 0` guard)
- Interior complete trades are summed from per-thread data (parallel), boundary trades are stitched (sequential)

### Dead Code Removal

With this fix (stitched trade counting) and the prior fix (exact sequential drawdown scan), three of the nine parallel reduction arrays are now dead:

| Array | Was | Dead because |
|-------|-----|-------------|
| `s_max_dd` | Parallel drawdown max-reduction | Overridden by thread-0 exact sequential scan |
| `s_trades` | Parallel trade count sum | Replaced by boundary stitching |
| `s_wins` | Parallel win count sum | Replaced by boundary stitching |

**Remove from the kernel:**
- Per-thread variables: `local_peak`, `local_max_dd` (drawdown tracking per-thread is dead — `local_cum` is still needed for the equity product reduction)
- Shared memory arrays: `s_max_dd`, `s_trades`, `s_wins`
- Reduction: remove the 3 dead arrays from the block reduction loop (9 arrays → 6)
- Per-thread drawdown tracking: remove the `local_peak`/`local_max_dd`/dd computation from the per-bar loop (saves 1 `fmaxf` + 1 division + 1 comparison per bar)

### Shared Memory Layout (after cleanup)

Reduction shrinks from 9 arrays to 6. Total shared memory drops from 25KB to 22KB:

```
s_sum      [0*stride .. 1*stride)     // Sharpe mean
s_sq_sum   [1*stride .. 2*stride)     // Sharpe variance
s_down_sq  [2*stride .. 3*stride)     // Sortino downside
s_buys     [3*stride .. 4*stride)     // action distribution
s_sells    [4*stride .. 5*stride)
s_holds    [5*stride .. 6*stride)
s_sorted   [6*stride .. 6*stride + 4096)  // sort scratch + boundary data
```

The 7 × 256 = 1,792 boundary values are stored in `s_sorted[stride..8*stride]` (7 arrays of `stride` floats, starting at offset `stride`). This memory is written after the equity product reduction and read before the bitonic sort — no conflict.

**Update `shmem_bytes` in `gpu_backtest_evaluator.rs`**: change from `(256 * 9 + 4096) * 4` to `(256 * 6 + 4096) * 4`.

### Int-as-Float Storage

Boundary ints (action IDs, counts) are stored in float shared memory via `__int_as_float()` / `__float_as_int()` — bit-exact round-trip, no precision loss.

### Correctness Verification

**Test case 1: Single trade spanning all chunks**
Window `[5,5,5,...,5]` (10K bars, all same action)
- Every thread: `num_changes=0`, `prefix_return = chunk_total`
- Stitching:
  - t=0: `open_action=-1` → set `open_action=5`, `open_return=0`; nc=0 → `open_return += chunk_0_total`
  - t=1..255: `fa=5 == open_action` → nc=0 → `open_return += chunk_t_total`
  - Final close → **1 trade, return = sum of all bars** ✓

**Test case 2: Trades at chunk boundaries**
Window `[5,5,5, 3,3, 7,7,7, 2,2]`, 2 chunks of 5 bars
- Thread 0 `[5,5,5,3,3]`: first=5, last=3, prefix=r[0]+r[1]+r[2], suffix=r[3]+r[4], nc=1, interior=0
- Thread 1 `[7,7,7,2,2]`: first=7, last=2, prefix=r[5]+r[6]+r[7], suffix=r[8]+r[9], nc=1, interior=0
- Stitching step-by-step:
  1. t=0: `open=-1` → set open=5. nc=1 → `open_return += prefix(r[0..2])` → CLOSE trade 1 (action=5, return=r[0..2]). Open=3, return=suffix(r[3..4])
  2. t=1: `fa=7 != open=3` → CLOSE trade 2 (action=3, return=r[3..4]). Open=7. nc=1 → `open_return += prefix(r[5..7])` → CLOSE trade 3 (action=7, return=r[5..7]). Open=2, return=suffix(r[8..9])
  3. Final close → trade 4 (action=2, return=r[8..9])
- Result: **4 trades, each with exact cumulative return** ✓

**Test case 3: Multiple changes in one chunk**
Actions `[5,3,7,2,8]` in one chunk
- nc=4, prefix=r[0] (action=5), suffix=r[4] (action=8), interior trades: action=3(r[1]), action=7(r[2]), action=2(r[3]) → 3 complete interior trades
- Stitching: open=5 → prefix(r[0]) → CLOSE trade (action=5). Open=8, return=suffix(r[4]). Final close trade (action=8).
- Total: 3 interior + 1 prefix close + 1 final close = **5 trades** ✓

**Test case 4: Alternating actions every bar (stress test)**
Actions `[0,1,2,3,4,5,6,7,8,0,1,...]` for 10K bars
- Each thread chunk of ~39 bars: nc=38, prefix=r[0], suffix=r[38], interior=37 complete trades
- Stitching: each boundary has `fa != la`, so boundary trades close normally
- Total: ~10K trades, each 1 bar long, each with exact single-bar return ✓

**Test case 5: Empty window (`wlen == 0`)**
- All threads get empty chunks (`first_action = -1`)
- Stitching: all skipped, `open_action` stays -1, no final close
- Result: **0 trades, win_rate = 0.0** ✓

### Kernel Header Comment Update

Update the kernel file's header comment (lines 4-5) to reflect exact counting:
```
//   [4] win_rate (winning_trades / total_trades, exact via boundary stitching)
//   [5] total_trades (exact position changes, boundary-aware)
```

### coordinator_extended.rs Fix

Replace the inline formula `(252.0 * 6.5 * 60.0_f64).sqrt()` with a named constant:

```rust
const BARS_PER_YEAR: f64 = 390.0 * 252.0;
let annualization_factor = BARS_PER_YEAR.sqrt();
```

This matches the pattern already used in `financials.rs` and `ab_testing.rs`.

## Files Modified

| File | Change |
|------|--------|
| `crates/ml/src/cuda_pipeline/backtest_metrics_kernel.cu` | Replace trade counting with boundary-aware hybrid; remove 3 dead reduction arrays + per-thread drawdown vars; update header comment |
| `crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs` | Update `shmem_bytes` from `(256*9+4096)*4` to `(256*6+4096)*4` |
| `crates/ml-ensemble/src/coordinator_extended.rs` | Extract `BARS_PER_YEAR` constant |

## Performance Impact

- **Per-thread pass**: +4 registers for boundary tracking, −4 registers from removed drawdown+trade vars. Net: 0 registers.
- **Reduction**: 6 arrays instead of 9 → 33% fewer reduction operations.
- **Per-bar loop**: Removes 3 operations/bar (fmaxf peak, drawdown division, drawdown comparison) from the parallel pass. Still computed exactly by thread-0 sequential scan.
- **Shared memory**: 22KB instead of 25KB → 12% reduction, more headroom for L1 cache.
- **Boundary export**: 7 shared memory writes per thread + 1 `__syncthreads()`. Negligible.
- **Stitching**: O(256) loop in thread 0. Negligible vs existing O(10K) sequential scan.
- **Net effect**: Kernel is slightly FASTER due to dead code removal, despite added stitching complexity.

## Risks

- **Complexity**: The stitching algorithm has edge cases (empty chunks, single-bar chunks, all-same-action windows, empty windows). Each is covered by the 5 test cases above and will have dedicated unit tests.
- **Shared memory aliasing**: `s_sorted` is reused for boundary data between equity reduction and bitonic sort. Execution order enforced by `__syncthreads()` barriers — safe.
