#239 verify: stage the mode-2 PCIe share on the copy stream inside the captured window

描述

## What

Mode 2 (`--pcie-mode kernel`) staged a layer's PCIe-blob share with `fetch_blobs` + `rebase_ptrs` running on the **decode stream**, inside the captured verify window. The window therefore spent a stage in a `waitB` spin until that ~1.3 MB staging copy had landed - 31-39 ms per window on the profiled stage line before this change.

This moves the staging copy to the copy stream (`copy_`) so it runs beside the layers' own work, **as captured graph nodes**, with the rendezvous carried by device flags in the verifier's arena instead of host events.

## The three capture traps (all hit, all handled)

The window is a captured CUDA graph, so copy-stream work has to become graph nodes:

1. **Kernels on a stream that has not joined the capture are not captured** - they launch during capture itself, on garbage, which surfaced as an illegal access on first decode. `copy_` joins through an event record/wait (`joint_ev_`), which becomes a graph edge and replays with the graph.
2. **`EndCapture` refuses a graph whose secondary stream has no path back** to the capturing stream's end (`capturing stream has unjoined work`), so `end_copy_share` closes the loop with a second edge (`joint_bwd_`).
3. **A wait node inside a replayed graph cannot consult a host event's runtime state.** The cross-stream rendezvous therefore runs on device flags in the verifier's arena (`pipe_flags_`, 8 words: plan(2) | staged(2) | pcie done(2) | window start), cleared by a kernel at the start of every window and raised by kernels; the ordering hazards the flags cover are the plan's device copy (`plan`) and the staging slot's reader - the previous layer's grouped kernel (`pcie done`).

## Safety

- modes 0 and 1 keep their paths byte for byte; mode 2's decode-stream inline copy is replaced and the unreachable leftover `else if (pcie_mode == 2)` branch is gone.
- `wait_flag_ge_or` on a null `skip_` (device plan off) is guarded; the pre-existing decode-side waits (flag B, stage done) stay.
- The window still drains both streams at the end (`cudaStreamSynchronize(copy_)` covers the graph's copy-stream tail), so nothing leaks across windows.

## Measured (RTX 2070 + i7-8700K AVX-2, qwen3.8-flash-next 125B-MoE Q2_0, T up to 6)

- waitB on the decode stage line: **31-39 -> ~4 ms/window** - the mechanism works end to end and reverts cleanly.
- end-to-end: neutral (112.4/117.4/123.9 vs 114.2/116.5/117.3 ms/window across warmed paired passes, i.e. within +-2%). The CPU expert pool already covered the waitB stall, so this change alone does not move total time on this profile - it leaves the decode stream free for compute (this is also what a faster CPU pool or a GPU-side expert-cache policy would want; the decode-side `waitCPU` of 16-21 ms is the bottleneck here, not the PCIe staging).