diff --git a/docs/plans/2026-02-28-gpu-pipeline-phase2b-design.md b/docs/plans/2026-02-28-gpu-pipeline-phase2b-design.md
new file mode 100644
index 000000000..d422454f4
--- /dev/null
+++ b/docs/plans/2026-02-28-gpu-pipeline-phase2b-design.md
@@ -0,0 +1,288 @@
+# 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.