Fix MIGraphX EP decode nondeterminism from unsynchronized output reads

[aditya-dl]

Greedy decode through the MIGraphX EP returns different token ids on every iteration — same model, same params, same process. ep.9D is stable on the same model and script.

2f6c1d6 removed the per-Compute hipStreamSynchronize calls, on the basis that ORT flushes the compute stream at run end via DeviceStreamCollection::CleanUp -> SyncStream::Flush.

That flush does happen, but after the outputs have already been read — CopyOutputsAcrossDevices runs inside ExecuteGraphImpl (inference_session.cc:3231), while OnRunEnd and CleanUp are at :3244 and :3260. And because our compute stream is created hipStreamNonBlocking (shared/hip/stream_support.h), it is exempt from the null stream's implicit ordering — so nothing prevents the fetch from reading logits while kernels are still writing them.

Two commits:

• Reinstate hipStreamSynchronize(hip_stream) at the end of the three Compute paths. Only that hunk of 2f6c1d6 is undone — the staging and coalesce-IO work stays, and the OnRunEnd sync it dropped is correctly left out, since stream_ never carried compute work.
• Make the D2H branch of hip::DataTransfer::CopyTensors honour the stream ORT supplies (hipMemcpyWithStream), matching the D2D branch and the in-tree MIGraphX EP. Hardening, not a replacement: ORT passes a null stream on the single-pair path, and the staging path bypasses DataTransfer entirely.

Verified, 5 iterations each, warm cache, greedy:

Model	Tokens	Result
DeepSeek-1L	8	byte-identical to ep.9D, zero divergence
DeepSeek-R1-Distill-Qwen-1.5B_fp16	16	byte-identical to ep.9D, zero divergence

Also re-ran the two-session concurrent repro: both sessions run, no failure at data_transfer.cc:73.

Throughput is within run-to-run noise on both models, so the apparent gain from removing the syncs was skipped correctness. The env-gated deferred-sync experiment on main (e5ad9cb) is not needed here, and would not fix this in any case: it waits in OnRunEnd, which is on the wrong side of the output read.

Note src/migraphx/*.cc builds into migraphx-backend.dll, not amdgpu-ep.dll — both need deploying to test this.

[aditya-dl]

@tperry-amd could you take a look at this here? This is a gating issue for 10D since it affects all LLM workloads through OGA. I would also like your input in terms of is there a better solution that I am missing.