[sergqwer]

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.ai/code/session_01VZy1yKaDDiA8a7svdwaHio

[Niko1221]

Thanks for finding this and for the fix! The same change went into 0.1.38 as part of #548 (decode_cluster_parity waits for every upload before the graph replays), so this PR is superseded. Closing.