#536 decode_cluster_parity: the graph replays wait for their uploads (the test failed now and then)

本文

`decode_cluster_parity` fails now and then on its own: a graph replay reports one differing token, from run to run, with no code change.

**Cause:** `run_graph_case` uploads each replay's scores, steps and logits with `Dev::put`, a pageable `cudaMemcpy` on the legacy stream. It then launches the graph on `cs`, a `cudaStreamNonBlocking` stream.
- A pageable host-to-device `cudaMemcpy` can return once the data is staged, before its DMA has landed.
- A non-blocking stream does not wait for the legacy stream.

So the captured cluster kernels sometimes read the previous replay's buffers. The one-CTA reference runs later on the same stream, after the copy has landed, and the two disagree. This is the same race 0.1.36 fixed in `prefill_fused_moe_test` ("the fused test waits for its uploads"). The test's other cases are not affected: the plain cases run on the legacy stream, ordered after the copy, and the timing cases compare nothing.

**Fix:** `cudaDeviceSynchronize()` after the uploads, before `cudaGraphLaunch`, as in the fused test. Test only; no engine code changes. The engine itself feeds these kernels from kernels on the same stream, not from host uploads, so the cluster kernels' bit-identity holds there.

**Measured** (RTX 5090, Windows 11, CUDA 13.3, driver 617.14; the test built from v0.1.37):

| | runs that failed | graph replays that differed |
| --- | ---: | ---: |
| v0.1.37 as is, first 20 runs | 4 of 20 | 4 of 120 |
| v0.1.37 as is, 20 runs alternating with the next row | 6 of 20 | 9 of 120 |
| with this change, alternating | 0 of 20 | 0 of 120 |
| this branch, 20 more runs | 0 of 20 | 0 of 120 |

Every failure was a graph replay with one differing token. The QSA top-k cases (434) and the argmax cases (216) passed in all runs.

🤖 Generated with [Claude Code](https://claude.com/claude-code)

https://claude.ai/code/session_01VZy1yKaDDiA8a7svdwaHio