# Unified cuBLAS Training Graph

## Goal

Replace the fragmented graph capture system (graph_forward + graph_aux + graph_spectral + graph_mega + graph_adam) with a single cudarc-captured unified graph. Rewrite IQL, IQN, and Attention auxiliary trainers from per-sample CUDA kernels to batched cuBLAS GEMMs. Place temporal ops (Mamba2, ISV temporal route, regime dropout, predictive coding, risk budget) inside the unified graph. Target: ~60s epochs (from 685s), aux_frequency removed, zero dead code.

## Context

### Current State

The training pipeline uses up to 6 separate CUDA graphs:
- `graph_forward`: cuBLAS trunk forward + ISV + plan head + loss + grad + backward
- `graph_forward_ddqn`: Double DQN Pass 3
- `graph_adam`: grad_norm + Adam + unflatten
- `graph_aux`: HER + EMA + IQL + IQN + attention + CQL + recursive_confidence_backward (cudarc wrapper)
- `graph_spectral`: spectral norm (cudarc wrapper)
- `graph_mega`: all-in-one (disabled — raw cuda_sys capture corrupted graphs on Hopper)

Aux ops (IQL x2, IQN, Attention) use per-sample CUDA kernels that do manual matrix multiplication in thread blocks — 16,384 kernel launches per GEMM layer. This takes ~3.5 seconds per training step, dominating epoch time (682s of 685s per epoch).

Temporal ops (Mamba2, ISV temporal route, regime dropout, predictive coding, risk budget) cannot run between graph replays — they corrupt graph_aux replay on Hopper. Currently disabled.

`aux_frequency=4` workaround skips aux ops 3 out of 4 steps.

### Root Cause

1. **Per-sample kernels**: IQL/IQN/Attention launch one thread block per sample (B=16,384). Each block does a manual dot product. cuBLAS does the same work in ONE call.
2. **Graph fragmentation**: Temporal ops between separate graph replays corrupt cudarc's internal state on Hopper. A single unified graph eliminates the gap.
3. **Raw cuda_sys**: `capture_graph_mega` and `capture_graph_spectral` originally used raw `cuda_sys::cuStreamBeginCapture_v2` which bypasses cudarc's state tracking. Fixed for graph_aux/spectral but graph_mega still hangs even with cudarc — likely due to combining cuBLAS with per-sample kernels in one capture.

## Design

### Composable Child Graph Architecture

NVIDIA's recommended pattern for complex pipelines: independently captured child
sub-graphs composed into a single parent graph via `cudaGraphAddChildGraphNode`.
Each child is captured, validated, and testable in isolation. The parent graph
launches all children in one call.

```
parent_graph (single launch per training step)
  ├── child: spectral_subgraph
  │     Spectral norm (10 weight matrices)
  │
  ├── child: forward_subgraph
  │     cuBLAS trunk forward (Pass 1 + 2)
  │     ISV forward + feature gate + plan head + recursive confidence
  │     Mamba2 temporal scan + ISV temporal route + regime dropout
  │     Predictive coding + risk budget
  │     C51/MSE loss + gradient computation
  │     cuBLAS trunk backward
  │
  ├── child: ddqn_subgraph
  │     cuBLAS DDQN Pass 3 (Double DQN on next_states)
  │
  ├── child: aux_subgraph
  │     HER in-place relabel
  │     EMA target update
  │     cuBLAS IQL high-tau: forward + loss + backward + adam
  │     cuBLAS IQL low-tau: same
  │     cuBLAS IQN: embed + forward + loss + backward + adam + EMA
  │     cuBLAS Attention: Q/K/V/O projections + backward + adam
  │     CQL gradient + SAXPY
  │     Recursive confidence backward
  │
  └── child: adam_subgraph
        ISV signal update
        Mamba2 BPTT + SGD
        Grad norm (two-phase reduction)
        Adam update + unflatten
```

**Why child graphs instead of flat capture:**
1. Each child is captured with its own `begin_capture`/`end_capture` — if one fails, the others still work
2. `cudaGraphExecChildGraphNodeSetParams` lets you update individual children without recapturing the parent
3. graph_mega hung because flat capture of mixed cuBLAS + NVRTC kernels corrupts on Hopper — child graphs isolate each subsystem
4. Each child can be tested independently before composing
5. Future changes to one subsystem only recapture that child

**Replaces ALL existing graphs:**
- `graph_forward` → forward_subgraph (child)
- `graph_forward_ddqn` → ddqn_subgraph (child)
- `graph_adam` → adam_subgraph (child)
- `graph_aux` → aux_subgraph (child)
- `graph_spectral` → spectral_subgraph (child)
- `graph_mega` → DELETED (parent graph IS the mega graph)

Single `parent_graph.launch()` per step. Single ownership in FusedTrainingCtx.

### IQL cuBLAS Rewrite

Current architecture (per-sample kernels):
- Forward: `iql_forward_loss_kernel` — grid=(B,1,1), block=(256,1,1). Each block manually computes 3-layer MLP + loss for one sample.
- Backward: `iql_backward_per_sample_kernel` — tiled, min(B,256) samples per tile. Manual Jacobian computation.
- Reduce: `iql_weight_grad_reduce_kernel` — sums per-sample gradients.

New architecture (batched cuBLAS):

**Forward (3 cuBLAS GEMMs + 2 activation kernels + 1 loss kernel):**
```
h1_pre = cublasLtMatmul(W1[128, state_dim], states[state_dim, B], ldb=state_dim_padded)
h1 = silu_kernel(h1_pre)                    // cubin element-wise
h2_pre = cublasLtMatmul(W2[128, 128], h1)
h2 = silu_kernel(h2_pre)
v = cublasLtMatmul(W3[1, 128], h2)
loss = expectile_loss_kernel(v, q_taken)     // cubin element-wise
```

Save h1_pre, h1, h2_pre, h2 for backward. Total: 3 GEMMs + 3 cubin kernels = 6 ops.

**Backward (6 cuBLAS GEMMs + activation backward kernels):**
```
dv = d_expectile_loss_kernel(v, q_taken)     // [1, B]
dW3 = cublasLtMatmul(dv[1, B], h2^T[B, 128])        // dW: [1, 128]
dh2 = cublasLtMatmul(W3^T[128, 1], dv[1, B])        // dX: [128, B]
dh2_pre = dh2 * d_silu(h2_pre)
dW2 = cublasLtMatmul(dh2_pre[128, B], h1^T[B, 128])
dh1 = cublasLtMatmul(W2^T[128, 128], dh2_pre)
dh1_pre = dh1 * d_silu(h1_pre)
dW1 = cublasLtMatmul(dh1_pre[128, B], states^T[B, state_dim])
```

Total: 6 GEMMs + 2 activation backward kernels = 8 ops.

**Adam + auxiliary kernels:**
```
grad_norm_phase1 + grad_norm_phase2 + adam_update = 3 kernels
loss_reduce = 1 kernel
```

**Total IQL per step: 18 ops** (vs current 7 fused but 16,384x slower each).

### IQN cuBLAS Rewrite

Similar pattern. Current: fused per-sample quantile forward+loss.

New forward:
```
tau_embed = cublasLtMatmul(W_embed[256, 64], cos_features[64, B*Q])  // Q=num_quantiles
combined = h_s2 * sigmoid(tau_embed)                                  // element-wise
// Per-branch outputs (4 branches, each 3 actions):
logits_b0 = cublasLtMatmul(W_b0[3, 256], combined)
logits_b1 = cublasLtMatmul(W_b1[3, 256], combined)
logits_b2 = cublasLtMatmul(W_b2[3, 256], combined)
logits_b3 = cublasLtMatmul(W_b3[3, 256], combined)
```

Backward: corresponding dW and dX GEMMs. Total: ~20 ops.

### Attention cuBLAS Rewrite

Current: manual Q/K/V projection per sample.

New:
```
Q = cublasLtMatmul(W_Q[256, 256], h_s2[256, B])
K = cublasLtMatmul(W_K[256, 256], h_s2[256, B])
V = cublasLtMatmul(W_V[256, 256], h_s2[256, B])
// Multi-head attention scores (per-head: 64-dim, 4 heads)
// Softmax + V aggregation via cubin kernel (cross-sample, not matmul)
O = cublasLtMatmul(W_O[256, 256], attn_out[256, B])
```

Backward: 4 dW GEMMs + 4 dX GEMMs. Total: ~12 ops.

### Shared cuBLAS Handle

All aux cuBLAS GEMMs use the existing `SharedCublasHandle::lt_handle`. GEMM descriptors pre-cached at init time (same pattern as trunk forward). Each IQL/IQN/Attention gets its own `CublasGemmSet` with cached descriptors for their specific shapes.

IQL GEMM shapes (cached at init):
- (128, B, state_dim, state_dim_padded) — h1 forward
- (128, B, 128, 128) — h2 forward
- (1, B, 128, 128) — v forward
- Plus transposed variants for backward dW and dX

### Activation Save Buffers

Each IQL/IQN/Attention allocates activation save buffers at init:
- IQL: `save_h1_pre[B, 128]`, `save_h1[B, 128]`, `save_h2_pre[B, 128]`, `save_h2[B, 128]`
- IQN: `save_embed[B*Q, 256]`, `save_combined[B*Q, 256]`
- Attention: `save_Q[B, 256]`, `save_K[B, 256]`, `save_V[B, 256]`

### Temporal Ops Placement

Inside the unified graph, between step 3 (ISV/plan) and step 5 (loss):

```
[cuBLAS forward → h_s2]
  → ISV forward + feature gate + plan head
  → Mamba2 temporal scan (enriches h_s2 with K=8 history)
  → ISV temporal route (per-feature temporal weights)
  → Regime dropout (regime-conditioned)
  → Predictive coding loss (temporal smoothness)
  → Risk budget forward
[loss computation uses enriched h_s2]
```

No gap between graph replays. All captured in one continuous cudarc session.

### ISV Signal Update

Inside the unified graph, after Adam update (step 18). Reads from pinned device-mapped scalars (loss, grad_norm, Q-mean) written by earlier nodes in the same graph. The graph's stream ordering guarantees visibility.

### What Gets Deleted

**Structs/fields removed from FusedTrainingCtx:**
- `graph_forward: Option<SendSyncGraph>`
- `graph_forward_ddqn: Option<SendSyncGraph>`
- `graph_adam: Option<SendSyncGraph>`
- `graph_aux: Option<SendSyncGraph>`
- `graph_spectral: Option<SendSyncGraph>`
- `graph_mega: Option<SendSyncGraph>`
- `aux_frequency: usize`

**Replaced with:**
- `parent_graph: Option<CudaGraphExec>` — the composed parent
- `spectral_subgraph: Option<CudaGraph>` — child: spectral norm
- `forward_subgraph: Option<CudaGraph>` — child: trunk forward + temporal + loss + backward
- `ddqn_subgraph: Option<CudaGraph>` — child: Double DQN Pass 3
- `aux_subgraph: Option<CudaGraph>` — child: IQL + IQN + attention + CQL + HER + EMA
- `adam_subgraph: Option<CudaGraph>` — child: ISV + mamba2 bwd + grad_norm + adam

**Functions removed:**
- `capture_graph_aux()`
- `capture_graph_spectral()`
- `capture_graph_mega()`
- `submit_aux_ops()` (inlined into unified capture)
- `pre_replay_state_update()` (host-side state updates move before unified launch)

**From GpuDqnTrainer:**
- `capture_training_graphs()` — replaced by unified capture in FusedTrainingCtx
- `replay_forward()` / `replay_adam()` — replaced by single unified launch
- `graph_forward` / `graph_forward_ddqn` / `graph_adam` fields

**From gpu_iql_trainer.rs:**
- `forward_loss_kernel` + `backward_per_sample_kernel` + `weight_grad_reduce_kernel` — replaced by cuBLAS GEMMs + small cubin kernels
- `grads_per_sample` buffer (tiled backward scratch) — not needed with cuBLAS backward
- `grad_tile_size` — not needed

**From gpu_iqn_head.rs:**
- Same per-sample kernel removal pattern

**From gpu_attention.rs:**
- Same per-sample kernel removal pattern

**Removed from config:**
- `aux_frequency` field from `DQNHyperparameters`

**CUDA kernel files — removed kernels:**
- `iql_value_kernel.cu`: `iql_forward_loss_kernel`, `iql_backward_per_sample_kernel`, `iql_weight_grad_reduce_kernel` — replaced by cuBLAS + `iql_silu_kernel`, `iql_expectile_loss_kernel`, `iql_d_silu_kernel`
- Similar removals in IQN and attention kernel files

### New cubin Kernels (element-wise, small)

Added to existing .cu files:

**iql_value_kernel.cu:**
- `iql_silu_fwd(x, out, N)` — element-wise SiLU activation
- `iql_silu_bwd(x, d_out, d_x, N)` — SiLU backward
- `iql_expectile_loss(v, q_taken, tau, loss, d_v, B)` — expectile loss + gradient
- `iql_bias_add(x, bias, out, M, N)` — add bias after GEMM

**iqn_dual_head_kernel.cu:**
- `iqn_cos_embed_fwd(tau, W_embed, out, B, Q, embed_dim)` — cosine embedding forward
- `iqn_hadamard_sigmoid(h_s2, embed, out, N)` — element-wise h_s2 * sigmoid(embed)
- `iqn_quantile_huber_loss(...)` — quantile Huber loss + gradient

**attention kernel file:**
- `attn_softmax_fwd(scores, out, B, num_heads, seq_len)` — multi-head softmax
- `attn_softmax_bwd(...)` — softmax backward
- `attn_layer_norm(...)` — LayerNorm forward+backward

### File Structure

Modified files:
- `crates/ml/src/cuda_pipeline/gpu_iql_trainer.rs` — rewrite to cuBLAS
- `crates/ml/src/cuda_pipeline/gpu_iqn_head.rs` — rewrite to cuBLAS
- `crates/ml/src/cuda_pipeline/gpu_attention.rs` — rewrite to cuBLAS
- `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs` — remove graph_forward/adam/ddqn, expose submit ops
- `crates/ml/src/trainers/dqn/fused_training.rs` — unified graph capture + launch, delete all split-graph code
- `crates/ml/src/trainers/dqn/config.rs` — remove aux_frequency
- `crates/ml/src/cuda_pipeline/iql_value_kernel.cu` — remove per-sample kernels, add element-wise cubin kernels
- `crates/ml/src/cuda_pipeline/iqn_dual_head_kernel.cu` — same
- `crates/ml/src/cuda_pipeline/shared_cublas_handle.rs` — no changes (reuse existing handle)

### Constraints

- All cuBLAS GEMMs use `CUBLAS_COMPUTE_32F_FAST_TF32` (same as trunk)
- All cuBLAS descriptors cached at init time via `cublasLtMatmulAlgoGetHeuristic`
- Deterministic: zero atomicAdd in all new cubin kernels (same policy as existing code)
- IQL reads states_buf with `ldb=state_dim_padded` (correct stride, zero-copy)
- graph_unified captured via cudarc `stream.begin_capture()` / `stream.end_capture()`
- `disable_event_tracking()` before capture, `enable_event_tracking()` after
- No `RawCudaGraph`, no raw `cuda_sys` graph APIs anywhere

### Testing

- Smoke test on RTX 3050 (sm_80): `FOXHUNT_TEST_DATA=test_data/futures-baseline cargo test -p ml --lib -- smoke_tests --ignored --nocapture`
- H100 validation: 3-epoch baseline run, compare val_Sharpe/WinRate/Trades against pre-rewrite metrics
- Epoch timing: target <80s (from 685s)
- Determinism: two consecutive runs must produce identical val_Sharpe_raw

### Performance Target (Phase 1: Single-stream)

| Metric | Before | Phase 1 |
|--------|--------|---------|
| Epoch time | 685s | <80s |
| Per-step aux ops | ~3.5s | ~5ms |
| Graph launches/step | 4 | 1 |
| Temporal ops | disabled | active |
| aux_frequency | 4 (workaround) | 1 (every step) |
| Total kernel launches/step | ~71 (per-sample) | ~60 (cuBLAS batched) |

---

## Phase 2: Multi-stream Graph Parallelism (designed now, implemented after Phase 1 validates)

### Problem

Single-stream execution wastes H100's 132 SMs. When a 1-thread ISV signal
update runs, 131 SMs sit idle. The branch head GEMMs (direction, magnitude,
order, urgency) are independent but execute sequentially. IQL high/low are
independent but serialized.

### Architecture

The child graph architecture enables this naturally. Each child sub-graph
manages its own internal parallelism via fork-join events captured as graph
edges. The parent graph replays the full multi-stream topology in one launch.

```
forward_child (internal multi-stream):
    main_stream: states → h_s1 → h_s2 → [FORK EVENT]
                                              │
    branch_stream[0]: ───── h_v → v_logits ───┤
    branch_stream[1]: ───── h_b0 → b0_logits ─┤
    branch_stream[2]: ───── h_b1 → b1_logits ─┤
    branch_stream[3]: ───── h_b2 → b2_logits ─┤
                                              │
    main_stream: ──────── [JOIN EVENT] → loss → grad → backward

aux_child (internal multi-stream):
    stream_a: ── IQL high-tau (forward + backward + adam) ──┐
    stream_b: ── IQL low-tau (forward + backward + adam)  ──┤
                                                            ├── [JOIN] → CQL → conf_bwd
    stream_c: ── IQN (forward + backward + adam)          ──┤
    stream_d: ── Attention (forward + backward + adam)    ──┘
```

### CUDA Graph Multi-stream Capture

Use `CU_STREAM_CAPTURE_MODE_RELAXED` instead of `GLOBAL`. Fork-join events
within the capture are recorded as graph dependency edges:

```rust
// Inside forward_child capture:
stream.begin_capture(CU_STREAM_CAPTURE_MODE_RELAXED)?;

// Trunk (main stream)
submit_trunk_forward()?;  // states → h_s1 → h_s2

// Fork: record event on main stream, wait on branch streams
let fork_event = stream.record_event()?;
for branch_stream in &branch_streams {
    branch_stream.wait_event(&fork_event)?;
}

// Branch GEMMs (parallel on 4 streams)
for (i, branch_stream) in branch_streams.iter().enumerate() {
    submit_branch_forward(branch_stream, i)?;  // h_s2 → h_bi → logits_bi
    let done = branch_stream.record_event()?;
    stream.wait_event(&done)?;  // Join back to main
}

// Loss + backward (main stream, after all branches joined)
submit_loss_and_backward()?;

stream.end_capture()?;  // Graph captures the fork-join topology
```

The captured graph replays all 4 branches in parallel automatically — the
graph scheduler sees the dependency edges and dispatches to multiple SMs.

### Implementation Details

#### 2a. Parallel Branch Forward (forward_child)

**cuBLAS workspace isolation:** Each branch stream needs its own cuBLAS
workspace to avoid contention. `batched_forward.rs` already allocates:
- `branch_workspace_ptrs: [u64; 4]` — 4 × 32MB = 128MB total
- `branch_streams: [Arc<CudaStream>; 4]` — forked from main stream
- `trunk_done_event` + `branch_done_events[4]` — fork/join synchronization

**cublasLt handle sharing:** `cublasLtMatmul` is thread-safe when using
different streams + different workspaces. The single `lt_handle` is shared
across all 5 streams. Each branch calls `cublasLtMatmul(..., branch_ws_ptr,
branch_ws_size, branch_stream)` — no internal state conflict because cublasLt
is stateless per-call (unlike legacy cublas which has handle-bound stream).

**Buffer isolation:** Each branch writes to its own output buffer:
- Branch 0 (value): `on_v_logits_buf[B, NA]`
- Branch 1-3 (advantage): `on_b_logits_buf` at offsets `[B*b0*NA, B*b1*NA, ...]`
- No write conflicts — each branch has exclusive output region.

**Read sharing:** All branches READ from `save_h_s2[B, SH2]` (trunk output).
Read-read sharing is safe — no synchronization needed beyond the fork event.

**Backward parallelism:** Within each branch's backward:
- `dW_bi = d_logits_bi @ h_bi^T` (weight gradient) — writes to `grad_buf[offset_bi]`
- `dh_bi = W_bi^T @ d_logits_bi` (activation gradient) — writes to `d_h_bi_buf`
- These are independent for different branches. The trunk backward aggregates
  `d_h_s2 = sum(W_bi^T @ d_h_bi)` — this requires a join before the reduction.

#### 2b. Parallel Aux Trainers (aux_child)

**Workspace requirement:** IQL, IQN, and Attention each need their own cuBLAS
workspace. Current: all share `SharedCublasHandle::lt_workspace_ptr` (32MB).
Phase 2 adds per-trainer workspace buffers:

```rust
// In FusedTrainingCtx constructor:
let iql_workspace = stream.alloc_zeros::<u8>(32 * 1024 * 1024)?;   // 32MB
let iql_low_workspace = stream.alloc_zeros::<u8>(32 * 1024 * 1024)?;
let iqn_workspace = stream.alloc_zeros::<u8>(32 * 1024 * 1024)?;
let attn_workspace = stream.alloc_zeros::<u8>(32 * 1024 * 1024)?;
// Total: 128MB additional VRAM
```

Each trainer receives its own `(workspace_ptr, workspace_size)` at construction.

**Stream allocation:** 4 aux streams forked from main:
```rust
let aux_streams: [Arc<CudaStream>; 4] = [
    stream.fork()?,  // IQL high
    stream.fork()?,  // IQL low
    stream.fork()?,  // IQN
    stream.fork()?,  // Attention
];
```

**Dependency graph within aux_child:**
```
[fork from main after backward]
    │
    ├── stream_0: HER relabel → EMA target update → IQL high (fwd+bwd+adam)
    ├── stream_1: IQL low (fwd+bwd+adam)
    ├── stream_2: IQN (fwd+bwd+adam+EMA)
    └── stream_3: Attention (fwd+bwd+adam)
    │
[join all 4 → main stream]
    │
    └── CQL gradient (reads from all trainers) → recursive_confidence_bwd
```

HER + EMA run on stream_0 before IQL high because IQL reads from the updated
target weights. IQL low can start immediately (reads from online weights only).
CQL gradient reads from all trainers' outputs — must wait for all 4 to complete.

**Buffer isolation:** Each trainer has its own:
- `params_buf`, `m_buf`, `v_buf`, `grad_buf` — exclusive
- `states_buf` (read-only, shared) — safe
- `q_out_buf` (read by IQL gather) — written by forward_child, read by aux_child.
  Dependency enforced by child graph ordering (forward_child before aux_child).

#### 2c. Double-buffered Experience Collection

**Problem:** Experience collection (4096 episodes × 1000 timesteps = ~1.2s)
runs sequentially before training (~56s on single stream). The GPU idles
during collection because training graphs can't run simultaneously.

**Architecture:** Two sets of experience output buffers (ping-pong):

```rust
struct DoubleBufferedExperience {
    // Buffer A
    states_a: CudaSlice<f32>,    // [episodes * timesteps * state_dim]
    actions_a: CudaSlice<i32>,
    rewards_a: CudaSlice<f32>,
    dones_a: CudaSlice<f32>,

    // Buffer B (same sizes)
    states_b: CudaSlice<f32>,
    actions_b: CudaSlice<i32>,
    rewards_b: CudaSlice<f32>,
    dones_b: CudaSlice<f32>,

    // Which buffer is "training" (being read by graph) vs "collecting" (being written)
    active: usize,  // 0 = A training, B collecting. 1 = swap.
}
```

**Stream separation:**
- `training_stream`: runs the parent graph (reads from active buffer)
- `collection_stream`: runs experience collection (writes to inactive buffer)
- Synchronization via events:
  ```
  Epoch N:
    collection_stream: collect into buffer B
    training_stream: train from buffer A (parent_graph.launch())
    [event: collection B done]
    [event: training A done]
    swap active buffer
  Epoch N+1:
    collection_stream: collect into buffer A
    training_stream: train from buffer B
    ...
  ```

**VRAM cost:** Doubles the experience output buffer: ~1.5GB × 2 = 3GB additional.
On H100 (80GB) with ~47GB used, this fits within headroom.

**Replay buffer interaction:** The replay buffer is written by collection and
read by training. With double buffering, the PER segment tree updates happen
on the collection stream AFTER experience is written. Training reads from
the PREVIOUS epoch's segment tree state. One-epoch staleness in priorities
is acceptable — PER priorities are approximate by nature.

**Graph capture implication:** The parent graph reads from fixed buffer
addresses (captured at step 0). With double buffering, the parent graph must
be re-instantiated each epoch to point at the swapped buffer — or use
`cudaGraphExecKernelNodeSetParams` to update the buffer pointers. The child
graph architecture helps: only `forward_child` reads from the experience
buffer (states/rewards/dones). Re-capture just that one child, re-compose
parent. Other children unchanged.

### Expected Performance

| Metric | Phase 1 | Phase 2a+b | Phase 2c |
|--------|---------|------------|----------|
| Forward pass | ~3ms | ~1.2ms | ~1.2ms |
| Aux ops | ~5ms | ~2ms | ~2ms |
| Experience collection | 1.2s (blocking) | 1.2s (blocking) | 0s (overlapped) |
| SM utilization | ~40% | ~85% | ~90% |
| **Epoch time** | **<80s** | **<40s** | **<38s** |

### What Already Exists

`batched_forward.rs` infrastructure ready for Phase 2a:
- `branch_streams: [Arc<CudaStream>; 4]` — 4 dedicated branch streams
- `branch_workspace_ptrs: [u64; 4]` — per-branch cuBLAS workspace (no contention)
- `trunk_done_event: CudaEvent` — fork event after trunk
- `branch_done_events: [CudaEvent; 4]` — join events per branch
- `sgemm_f32_branch()` — already routes GEMMs to per-branch stream+workspace

These are currently unused during graph capture (single-stream mode). Phase 2a
activates them inside the `forward_child` capture with RELAXED mode.

### Determinism Guarantee

Multi-stream execution is non-deterministic in general (kernel scheduling
order varies). But CUDA Graph replay is deterministic — the graph captures
the exact dependency topology and replays it identically every time. The
fork-join events become fixed graph edges. As long as:
1. No atomicAdd across streams (we have zero atomicAdd policy)
2. No floating-point reduction across streams (each stream reduces independently)
3. All cross-stream joins happen via captured events (not ad-hoc synchronization)

...the multi-stream graph replay is bit-identical across runs.

### Implementation Order

1. Phase 1: single-stream child graphs (this spec) — validate correctness
2. Phase 2a: multi-stream forward_child — parallelize 4 branch heads
3. Phase 2b: multi-stream aux_child — parallelize IQL/IQN/attention (128MB additional workspace)
4. Phase 2c: double-buffered experience collection (3GB additional VRAM)

Each phase is additive — Phase 1 child graphs don't change, they just gain
internal parallelism. The parent graph composition is unchanged.

---

## Phase 3: Cost-Driven Hold Timing (replaces min_hold_bars)

### Problem

`min_hold_bars=50` is a hardcoded gate that blocks exits before 50 bars.
`enforce_hold()` in `trade_physics.cuh` overrides the model's action. This
prevents the temporal attention from learning hold timing — the model never
sees the consequences of early exits because they're blocked.

### Design: Continuous Cost Signal

Remove `min_hold_bars` and `enforce_hold()` entirely. Replace with three
cost components in the reward function:

```
reward = base_reward
       - round_trip_cost * |trade_occurred|        # per-trade cost
       - holding_cost * |position| * dt            # continuous inventory penalty
       + conviction * plan_scaling                  # existing plan head
```

#### Component 1: Per-trade cost (Gem — Implementation Shortfall)

Each trade entry/exit incurs a fixed cost equal to the round-trip transaction
cost (commission + spread). This is already implemented in `experience_env_step`
via `compute_tx_cost()`. No change needed — just ensure it's always active.

The key: the cost is VISIBLE in the reward. The model learns "each trade
costs X bps" and naturally reduces frequency to maximize net reward.

#### Component 2: Opportunity cost (Pearl — Self-supervised inaction penalty)

When the model is flat, penalize inaction proportional to the MODEL'S OWN
predicted edge (Q-value spread between best long and best short):

```cuda
// In experience_env_step, when position ≈ 0:
// q_spread comes from the Q-values computed during action selection
float q_spread = fabsf(max_q_long - max_q_short);
float opp_cost_scale = 0.001f;  // calibrate to reward scale
float opportunity_cost = q_spread * opp_cost_scale * (fabsf(position) < 0.001f ? 1.0f : 0.0f);
reward -= opportunity_cost;
```

This creates a virtuous cycle:
- **Weak temporal attention** → noisy Q-values → small q_spread → low
  opportunity cost → model stays flat (correct — no signal)
- **Strong temporal attention** → accurate Q-values → large q_spread →
  high opportunity cost → model trades on signal (correct — has edge)
- **Model improves temporal attention** → better predictions → higher
  self-imposed penalty for missed trades → more trading → more reward

No inventory penalty. No hindsight bias (uses predicted edge, not actual
price change). The model teaches ITSELF when to trade. The temporal
attention has intrinsic motivation to improve — better predictions
increase the cost of doing nothing.

Calibration: `opp_cost_scale` should be set so that the opportunity cost
at typical Q-spread (~0.1-1.0) is comparable to the per-trade transaction
cost (~0.52 bps). This way the model trades when its predicted edge exceeds
the round-trip cost.

#### Component 3: Churn penalty (Novel — TFJ-DRL dual signal)

Add a supervised penalty for position flips within a short window:

```cuda
// Track bars_since_last_trade in portfolio state
float bars_since = ps[10];  // hold_time
float churn_threshold = 10.0f;  // minimum bars between trades
if (trade_occurred && bars_since < churn_threshold) {
    float churn_penalty = 0.001f * (churn_threshold - bars_since) / churn_threshold;
    reward -= churn_penalty;
}
```

This is NOT a hard gate (the model CAN exit early), but early exits get a
graduated penalty. The penalty is zero after `churn_threshold` bars. The
temporal attention learns that waiting past the threshold is free but
exiting before costs extra.

### What Gets Deleted

- `min_hold_bars` field from `DQNHyperparameters`
- `min_hold_bars` field from `GpuBacktestConfig`
- `enforce_hold()` function from `trade_physics.cuh`
- `enforce_hold` call in `experience_env_step`
- `enforce_hold` call in backtest evaluator
- `isv_signals` adaptive hold extension (ISV now modulates gamma/gating, not hold)
- All `min_hold_bars` references across config files

### What Gets Added

In `experience_env_step` (`experience_kernels.cu`):
- `inventory_penalty = holding_cost_rate * |position|` subtracted from reward
- `churn_penalty` for rapid position flips

In `DQNHyperparameters` (`config.rs`):
- `opportunity_cost_scale: f64` (default: 0.001, calibrated to Q-spread vs tx_cost)
- `churn_threshold_bars: usize` (default: 10)
- `churn_penalty_scale: f64` (default: 0.001)

In `GpuBacktestConfig`:
- Same three parameters propagated to validation backtest

### Temporal Attention Role

With cost-driven hold timing, the temporal attention (Mamba2 + ISV) learns:
- **When to enter**: conviction high enough to overcome round-trip cost
- **How long to hold**: until holding cost exceeds expected remaining profit
- **When to exit**: cost of staying > cost of round-trip to re-enter later

ISV signals shift from "extend hold time" to "modulate risk tolerance":
- High grad_norm → reduce position size (not extend hold)
- Low regime_stability → widen stops (not force hold)
- Negative reward_ema → cut losses (not block exit)

The plan head's `target_bars` becomes a PREDICTION, not an enforcement.
The model predicts how long it will hold, and the cost signals teach it
whether that prediction was optimal.

### Implementation Order

Phase 3 runs AFTER Phase 1 validates (child graphs working, <80s epochs).
It's a reward engineering change, not a graph/GPU change:

1. Add `holding_cost_rate` + `churn_penalty` to experience_env_step
2. Remove `enforce_hold()` and `min_hold_bars`
3. Add config parameters
4. Propagate to backtest evaluator
5. Validate: train 50 epochs, compare trade frequency and val_Sharpe vs Phase 1

---

## Phase 1 Results (2026-04-18) — Measured on H100 PCIe

### What was delivered

- **cuGraphClone eliminated**: `std::mem::forget` takes ownership of cudarc's CudaGraph handles directly. cuGraphClone corrupted cublasLtMatmul HMMA tile scheduling on Hopper.
- **5 child graphs**: spectral + forward + ddqn + aux + adam — launched individually (parent composition deferred).
- **All features inside graph**: temporal ops, ISV, attention (12 cuBLAS GEMMs), curiosity inference (2 cuBLAS GEMMs), ensemble (2 fwd + KL gradient + trunk backward), IQL×2, IQN — all captured in child graphs.
- **Zero atomicAdd on gradient paths**: curiosity training, KAN backward, DT kernels, recursive_confidence all rewritten to per-sample/per-block output + deterministic reduce.
- **Zero cuStreamSynchronize in training loop**: IQN loss, Q-stats, DQN loss/grad_norm all use pinned device-mapped memory. Q-stats moved to epoch boundary only.
- **Zero cuMemcpyHtoDAsync in graph capture**: IQN/IQL/Attention adam_step use pinned device-mapped.
- **bf16 naming purged**: all `bf16_*()` alias calls replaced with f32 intrinsics across 4 crates.

### Measured performance

| Metric | Before (pre-cuBLAS) | After Phase 1 |
|--------|---------------------|---------------|
| Epoch time | 685s | ~568s (GPU-bound) |
| Per-step CPU time | 3537ms (sync-bound) | 0.1ms (fully async) |
| Per-step GPU time | ~3.5s (ungraphed) | ~4.0s (graph replay, all features) |
| Graph replay launch | N/A | 0ms (async) |
| Q-stats overhead | 4s/step (15 kernel re-run + sync) | 0 (epoch boundary only) |
| Write lock overhead | ~4s/step (async yield) | 0 (acquired once/epoch) |
| Attention | disabled (wrong cubin) | enabled (12 GEMMs, LayerNorm arg fix) |
| Curiosity inference | per-sample matmul | 2 cuBLAS GEMMs |
| Training steps/epoch | 178 | 178 |

### Bottleneck analysis

The CPU pipeline is now fully async — zero overhead. The epoch time is 100% GPU compute:

```
178 steps × 4.0s per graph replay  = 712s
+ step 0 ungraphed:                    3.5s
+ experience collection:               2s
+ epoch boundary (validation + sync):  ~1s
= ~718s total (GPU compute bound)
```

**The 4.0s per graph replay is the floor on single-stream.** The graph replay time
matches the ungraphed time — graph eliminates CPU launch overhead but the GPU compute
is unchanged. cuBLAS replaced manual matmul but with all new features enabled (attention,
curiosity cuBLAS, ensemble, temporal), the total kernel count is higher than before.
Each replay executes ~177 kernel nodes:
- spectral_child: ~12 kernels (spectral norm on 10 weight tensors)
- forward_child: ~60 kernels (cuBLAS trunk fwd + temporal + ISV + loss + backward)
- ddqn_child: ~15 kernels (Double DQN pass 3)
- aux_child: ~80 kernels (attention + IQL×2 + IQN + ensemble + CQL + HER + EMA)
- adam_child: ~10 kernels (mamba2 bwd + pruning + grad_norm + Adam + ISV update)

### What Phase 2 multi-stream should target

Based on measurements, the parallelism opportunities are:

1. **forward_child branch parallelism** (Section 2a in spec):
   - Trunk (shared): h_s1 → h_s2 is SERIAL (cuBLAS SGEMM on [B, state_dim] → [B, SH2])
   - Branches (independent): 4 branches (direction, magnitude, order, urgency) can fork after h_s2
   - Each branch: 2 cuBLAS GEMMs + bias+ReLU + 2 cuBLAS backward GEMMs
   - Estimated branch time: ~25% of forward_child
   - **Expected speedup**: ~1.3x on forward_child (trunk is serial, branches parallelize)

2. **aux_child trainer parallelism** (Section 2b in spec):
   - IQL high, IQL low, IQN, Attention are fully independent
   - Each does: ~16-20 cuBLAS GEMMs (forward + backward + adam)
   - Currently serial: ~80 kernels × sequential launch
   - **Expected speedup**: ~3-4x on aux_child (4 trainers run in parallel)

3. **Kernel profiling needed** (nsys):
   - Which cuBLAS GEMMs dominate? Trunk [B, 256] × [256, 256] or branch [B, 128] × [128, 52]?
   - Are small kernels (ISV, bias_add, ReLU) creating SM idle gaps?
   - Is the spectral_child + adam_child contribution significant or negligible?

### Revised targets

| Phase | Target epoch time | Speedup | Method |
|-------|-------------------|---------|--------|
| Phase 1 (current) | 718s | 1.0x baseline | single-stream graph replay |
| Phase 2a: branch parallelism | ~550s | 1.3x | 4 branch streams in forward_child |
| Phase 2b: aux parallelism | ~250s | 2.9x | 4 aux streams (IQL×2 + IQN + Attention) |
| Phase 2a+2b combined | ~200s | 3.6x | all independent ops parallel |
| Phase 2c: double-buffer experience | ~198s | marginal | overlap collection with training |
| nsys-guided kernel optimization | ~120s | target | fuse small kernels, reduce node count |

**Priority order**: 2b (aux parallelism) > 2a (branch parallelism) > nsys profiling > 2c.
aux_child contains ~80 of 177 kernels and all 4 trainers are fully independent —
this is the highest-impact parallelism opportunity.

**nsys profiling** should run BEFORE implementing 2a/2b to validate assumptions:
- Are the aux kernels actually SM-bound (parallelism helps) or memory-bound (it won't)?
- What fraction of the 4s/step is trunk (serial) vs branches+aux (parallelizable)?
- Are there unexpected bottlenecks (cuBLAS workspace allocation, small kernel gaps)?
