#548 decode_cluster_parity: graph replays race their input upload (pageable cudaMemcpy + non-blocking stream)

Description

At v0.1.37, `decode_cluster_parity --selftest` fails 5 of its 6 graph replays on an RTX 5060 Ti (sm_120). Each failing replay differs in one argmax token. The kernels look fine: the test's input upload races the graph launch, and with one synchronization after the upload the test passes every time.

## Environment

- RTX 5060 Ti 16 GB (cc 12.0, PCIe x4), running the test; an RTX 4070 SUPER (cc 8.9, x16) is also in the box, where the test skips
- Intel Core i5-13500, 48 GB RAM
- Windows 11 Pro 10.0.26200, NVIDIA driver 616.92
- Engine 0.1.37 (`db4f91a`), unmodified; built locally with CUDA 13.3, MSVC 19.44, `-DCMAKE_CUDA_ARCHITECTURES=89;120`; only the `decode_cluster_parity` target

## Measurements

```
CUDA_DEVICE_ORDER=PCI_BUS_ID CUDA_VISIBLE_DEVICES=1 decode_cluster_parity --selftest
```

Three consecutive runs all failed 5 of 6 replays and exited with 1 (an earlier run failed 4 of 6). This is one of them:

```
ok   graph replay 0 (ctx 100000): 0 ids, 0 tokens differ
FAIL graph replay 1 (ctx 104002): 0 ids, 1 tokens differ
FAIL graph replay 2 (ctx 108004): 0 ids, 1 tokens differ
FAIL graph replay 3 (ctx 112006): 0 ids, 1 tokens differ
FAIL graph replay 4 (ctx 116004): 0 ids, 1 tokens differ
FAIL graph replay 5 (ctx 120006): 0 ids, 1 tokens differ
QSA top-k: 434 cases, 0 failed
argmax: 216 cases, 0 failed
FAIL
```

With only the one-line fix below added to the same source and rebuilt, three runs out of three passed all 6 replays and exited with 0:

```
PASS: the cluster kernels are bitwise identical
```

## Likely cause

Not proven; it fits everything above:

- `run_graph_case` uploads each replay's inputs with `Dev::put`, a synchronous `cudaMemcpy` from a pageable `std::vector` (`src/kernels/decode_cluster_parity.cpp:60`). It then launches the graph on `cs` (`:279`), which is a `cudaStreamNonBlocking` stream (`:247`).
- For pageable host-to-device copies, `cudaMemcpy` may return once the data is staged, before the DMA reaches the device. A non-blocking stream is also not ordered behind the legacy stream. So the graph can start before the last upload has landed. The reference `sample_tokens` launch (`:283`) runs after the graph, so it reads the new logits.
- Only the argmax differs, and only for one of the 5 rows. The logits are the last and largest upload (5 x 248,320 floats, `:278`). The QSA inputs go up first and never differ.
- Replay 0 passes. Its first `cudaGraphLaunch` also uploads the executable graph, which may delay the kernels long enough for the copy to land.
- The single-launch cases never fail (216 argmax, 434 top-k).

## Suggested fix

Synchronize in `Dev::put`:

```cpp
void put(const std::vector<T>& h) {
    ck(cudaMemcpy(p, h.data(), h.size() * sizeof(T), cudaMemcpyHostToDevice), "up");
    ck(cudaDeviceSynchronize(), "up landed");
}
```

This also covers `bench()`, which has the same pattern (`d_l.put(l)` at `:359`, then work on its own non-blocking stream). Other options are `cudaMemcpyAsync` on `cs` from pinned buffers, or a sync before `cudaGraphLaunch`.

I have not checked whether any engine path uploads from pageable memory right before a launch on a non-blocking stream. As far as this test shows, the kernels are bitwise identical and the failure comes from the test itself.