I’m seeing unexpectedly long cudaStreamSynchronize() times even though I only do GPU->CPU (DtoH) copies once per batch (not per inference), and I’m trying to understand what is actually blocking the stream. My assumption was that batching the memcpy_dtoh_async would reduce sync overhead, but the sync still takes a lot of time, which makes me think the delay is not the copy call itself but some earlier queued work in the same stream (kernels, implicit sync points, or previous memcpy operations) that only gets exposed when I call cuda Synchronize(). Since CUDA streams execute strictly in order, I suspect a backlog in the queue is accumulating and only becomes visible at sync time, but I’m not sure how to pinpoint which operation is responsible. Is there a reliable way (e.g., CUDA profiler / Nsight Systems / event timing) to break down a stream and identify exactly which kernel or memcpy in the queue is causing the stall before the sync?
cudaStreamSynchronize(stream) will wait until all work submitted to the stream before synchronization is completed.
Looking at the nsight systems timeline, I would guess that the kernel (blue bars) runs in the same stream. Looking even closer, you will notice that the long synchronization happens before the memcpy operation (slim red vertical bar)
Why do you believe there is an issue with that timeline? What behaviour would you expect or like to achieve?
so i hit a wall , and cant improve it more? the sync is obligatory and cant be more efficient?
You can determine whether you hit certain limits.
E.g. running the kernels separately, measuring their time.
Knowing the PCIe bandwidth and calculating how long memcpys take, etc.
Even within kernels, what operations are needed, what is the theoretical maximum=roofline.
The sync may be obligatory or not. But all the needed work has to be done at some point in time.
Some operations needing different resources and execution units may be run in parallel.