# GPU Pipeline Phase 2b: Zero-Roundtrip CUDA Experience Collection

**Date:** 2026-02-28
**Branch:** `feature/cuda-pipeline-phase2b`
**Status:** Design approved
**Depends on:** Phase 2 (complete — GPU batch states + structural portfolio sim)

## Problem

Phase 2 achieved GPU batch state construction (eliminating 1.8M Vec allocations/epoch) and structural GPU portfolio simulation. However, the DQN trainer's inner loop still:

1. Runs portfolio simulation on CPU (duplicating GPU work)
2. Computes barrier tracking, curiosity, diversity, and reward combination on CPU per bar
3. Performs 15,600+ kernel launches per epoch with PCIe roundtrips each
4. Achieves only ~23% GPU utilization on L4

## Solution

A single monolithic CUDA kernel that runs the **entire experience collection pipeline** — portfolio simulation, Q-network inference (Dueling architecture), epsilon-greedy action selection, barrier tracking, diversity entropy, curiosity inference, and reward combination — with zero CPU roundtrips during execution.

128 parallel episodes × 500 timesteps = 64,000 experiences per kernel launch.

## Architecture

### Approach: Extended Monolithic Kernel + Parallel Episodes

**Why monolithic:** The kernel loop is inherently sequential per episode (portfolio state depends on previous bar). Adding more sequential logic (barrier, diversity, Q-network forward) per iteration costs negligible extra ALU ops. One kernel launch eliminates 15,600 launches worth of overhead.

**Why parallel episodes:** Each episode has independent portfolio state. N=128 threads process 128 episodes simultaneously. This is the standard GPU RL environment pattern (Isaac Gym, EnvPool).

### Data Flow (One Training Round)

```
CPU:  Generate 128 random episode_starts
CPU:  Upload episode_starts [128] to GPU                  (512B)
GPU:  Launch dqn_full_experience_kernel(N=128, L=500)     (ONE launch)
      |
      +-- Each of 128 threads independently:
      |     for t in 0..500:
      |       1. Read market features[bar, 51]
      |       2. Compute portfolio features [3] from local state
      |       3. Concatenate -> state [54]
      |       4. Dueling Q-network forward (online) -> 45 Q-values
      |       5. Epsilon-greedy action selection (LCG RNG)
      |       6. Portfolio simulation (trade execution)
      |       7. Barrier tracking (profit/stop/time)
      |       8. Diversity entropy (sliding window)
      |       9. Curiosity inference (forward model MSE)
      |      10. Mark-to-market -> next_state
      |      11. Dueling Q-network forward (target) -> target Q-values
      |      12. Compute TD error for priority replay
      |      13. Combine reward (PnL + barrier + risk + diversity + curiosity)
      |      14. Write experience to output buffers
      |      15. Handle episode reset if done
      |
CPU:  Download outputs (~16MB)                            (ONE transfer, <1ms)
CPU:  Add 64,000 experiences to replay buffer (with TD priorities)
CPU:  Sample batch, Candle gradient update (standard DQN)
CPU:  Every K rounds: Candle curiosity training batch
CPU:  After gradient update: sync weights to GPU           (1.2MB, <1ms)
```

### Performance Comparison

| Metric | Current (Phase 2) | Phase 2b |
|--------|-------------------|----------|
| Kernel launches/epoch | ~15,600 | ~31 |
| PCIe roundtrips/epoch | ~31,200 | ~62 |
| Experiences/launch | 128 | 64,000 |
| GPU compute utilization | ~23% | ~95%+ |
| CPU inner loop | Full per-bar processing | Replay buffer + gradient only |

## Kernel Components

### 1. Dueling Q-Network (Online + Target)

Production architecture: Dueling DQN with 54-dim input, 45-action output.

```
Shared pathway:  [54] -> [256] LeakyReLU(0.01) -> [256] LeakyReLU(0.01)
                         |                               |
         Value stream: [256]->[128] LReLU -> [1]   Advantage: [256]->[128] LReLU -> [45]
                         |                                      |
                     Q(s,a) = V(s) + [A(s,a) - mean(A)]
```

Each thread performs matrix-vector multiplies (not matrix-matrix):
- Layer 1: 54x256 = 13,824 FMA
- Layer 2: 256x256 = 65,536 FMA
- Value/Advantage streams: 2x(256x128 + 128x1/45) = ~66K FMA
- Total per forward pass: ~145K FMA per thread
- Two passes (online + target): ~290K FMA per thread per timestep

With 128 threads x 500 timesteps: ~18.6B FMA per kernel launch. Trivial for L4.

**Weight layout:** Candle convention `[out_features, in_features]`. Kernel accesses as:
```c
output[j] = bias[j];
for (int i = 0; i < in_dim; i++)
    output[j] += input[i] * weight[j * in_dim + i];
```

### 2. Curiosity Forward Model (Inference Only)

Architecture: 35 -> 64 LeakyReLU -> 32

Input encoding: 32 state features + 3 action one-hot (Long/Flat/Short mapping from 45 actions).

```c
int exposure = action_idx / 9;  // 0-4
int cat = (exposure <= 1) ? 0 : (exposure == 2) ? 1 : 2;  // Short/Flat/Long
```

Curiosity bonus = clamped MSE between predicted and actual next state embedding.

**Training:** Runs periodically on Candle (every K rounds). Download batch of transitions, run backward + AdamW step, re-upload weights to GPU.

### 3. Portfolio Simulation

Existing logic from Phase 2 kernel: trade execution with two-phase reversals, cash reserves, transaction costs. Per-thread portfolio state [8]: cash, position, entry_price, initial_capital, spread, last_price, reserve_pct, cum_costs.

### 4. Barrier Tracking

Per-thread barrier state [5]: entry_price_raw, upper_barrier, lower_barrier, expiry_step, barrier_result.

On position open: set barriers based on config (profit_target_bps, stop_loss_bps, max_holding_bars).
Per bar: check price against upper/lower barriers, check step against expiry.
Output: barrier_label (-1 stop, 0 time/none, +1 profit) for reward scaling.

### 5. Diversity Entropy

Per-thread sliding window [100 action indices] with circular buffer.
Shannon entropy: H = -sum(p_i * log2(p_i)) over 3-action histogram (Long/Flat/Short).
Penalty: -0.1 if H < 1.0 (low diversity).

### 6. Reward Combination

```c
float reward = raw_pnl_return;                          // (next_norm - curr_norm) / curr_norm
reward *= (1.0f + barrier_scale);                       // +0.5 profit, -0.5 stop, 0.0 time/none
if (abs_pos_norm > 0.8f)
    reward -= (abs_pos_norm - 0.8f) * 5.0f * 0.1f;     // risk penalty
reward += diversity_bonus;                               // -0.1 if low entropy
reward += fminf(curiosity_mse, curiosity_max_reward);   // novelty bonus
```

### 7. TD Error for Priority Replay

```c
float td_target = reward + gamma * max_target_q * (1 - done);
float td_error = fabsf(td_target - online_q_selected);
td_error_out[idx] = td_error;
```

### 8. Epsilon-Greedy RNG

LCG per thread, seeded from host:
```c
__device__ float gpu_random(unsigned int* state) {
    *state = *state * 1664525u + 1013904223u;
    return (float)(*state & 0x7FFFFF) / (float)0x7FFFFF;
}
```

## Memory Budget

### GPU Buffers

| Category | Size |
|----------|------|
| Market features [total_bars, 51] | ~400MB (2M bars) |
| Targets [total_bars, 4] | ~32MB |
| Online Q-network weights | 592KB |
| Target Q-network weights | 592KB |
| Curiosity weights | 17KB |
| Per-episode state (128 episodes: portfolio + barrier + diversity) | 57KB |
| Output buffers (128 x 500 x (54+1+1+1+1+1)) | ~16MB |
| RNG states [128] | 512B |
| Barrier config [4] | 16B |
| **Total** | **~450MB** |

Fits comfortably in L4 24GB or RTX 3050 Ti 4GB (with smaller dataset).

### Per-Thread Scratch (Registers/Local Memory)

| Buffer | Per Thread | 128 Threads |
|--------|-----------|-------------|
| shared_h1 [256] | 1KB | 128KB |
| shared_h2 [256] | 1KB | 128KB |
| value_h/adv_h [128+128] | 1KB | 128KB |
| q_values [45] | 180B | 22KB |
| target_q [45] | 180B | 22KB |
| curiosity_h [64] | 256B | 32KB |
| state [54] | 216B | 27KB |
| next_state [54] | 216B | 27KB |
| **Total** | **~4.1KB** | **~514KB** |

## Weight Synchronization

### Online Network Sync

After each gradient update round:
1. `vs.get("shared_layers.0.weight")` -> `Tensor` -> `.to_vec1::<f32>()`
2. Repeat for all 12 weight/bias tensors
3. `stream.memcpy_htod(&weights, &mut gpu_buffer)`
4. Total: 592KB upload per sync

### Target Network Sync

Every `target_update_freq` gradient steps:
1. Same extraction process from target VarMap
2. 592KB upload
3. Infrequent (typically every 1000-10000 steps)

### Curiosity Weight Sync

Every K training rounds:
1. Download batch of transitions from replay buffer
2. Run Candle backward + AdamW step on curiosity model
3. Re-upload 17KB weights to GPU
4. K chosen to balance curiosity signal freshness vs overhead

## Training Loop Changes

### Current (Phase 2)

```
for epoch in 0..num_epochs:
    for batch_idx in 0..num_batches:     // ~15,600 batches
        for bar in batch:                 // 128 bars per batch
            CPU: portfolio sim
            CPU: barrier tracking
            CPU: Q-network forward (Candle GPU)
            CPU: curiosity
            CPU: diversity
            CPU: reward combination
            CPU: store experience
        GPU: build_batch_states           // Phase 2 optimization
    for update in 0..gradient_steps:
        sample replay buffer
        Candle: forward + backward + step
```

### New (Phase 2b)

```
for epoch in 0..num_epochs:
    for round in 0..num_rounds:           // ~31 rounds
        CPU: generate 128 episode_starts
        GPU: ONE kernel launch -> 64,000 experiences
        CPU: download experiences (~16MB)
        CPU: add to replay buffer (with TD priorities)
    for update in 0..gradient_steps:
        sample replay buffer
        Candle: forward + backward + step
        every target_freq: sync target weights -> GPU
    sync online weights -> GPU
    every K rounds: Candle curiosity training + sync
```

## Concerns and Mitigations

### Training Dynamics (Sequential -> Random Episode Sampling)

**Risk:** Low. Random episode sampling reduces temporal correlation, which is desirable for DQN (replay buffer exists to break correlations). With ~31 rounds per epoch, data coverage is equivalent.

### Barrier Timing Precision

**Risk:** Solvable. GPU uses bar indices (i32) not nanosecond timestamps. `max_holding_bars` replaces `max_holding_period_ns`. No precision loss.

### Curiosity Signal Freshness

**Risk:** Moderate. Curiosity model trains periodically (every K rounds) instead of per-bar. More data per training step (batch of 1024 vs single transition) may compensate. Monitor curiosity bonus distribution during training.

### Done-Episode Idling

**Risk:** Minor. When a barrier terminates an episode at step 200/500, that thread idles for 300 steps. ~30% waste on average. Acceptable vs complexity of dynamic episode recycling.

### Epsilon Decay

**Risk:** None. Epsilon is a scalar parameter passed to kernel. CPU decrements between rounds. All threads use same epsilon per launch.

## Future Phase 3 Optimization

If >95% GPU utilization is insufficient:
- **Dynamic episode recycling:** When a thread's episode ends, start a new episode from a new random bar instead of idling. Eliminates 30% waste.
- **Multi-stream pipelining:** Launch experience collection on stream A while gradient updates run on stream B. Overlaps compute phases.
- **Larger N:** Scale from 128 to 512+ parallel episodes if GPU memory allows.
