What does cudaDeviceSynchronize() do?

Blocks the calling host thread until every previously launched piece of GPU work has finished.

Three different barriers get called "sync" and they sit at three different levels. __syncthreads() is a device-side barrier inside one thread block: it runs in the kernel, it costs nanoseconds, and it knows nothing about the host. cudaStreamSynchronize is a host call that waits for one stream and lets the others keep going. cudaDeviceSynchronize is the blunt one at the top: the host thread stops until the whole device is idle, every stream, every launch, every async copy. Reach for the widest of the three by default and you have thrown away the concurrency you asked for by launching asynchronously.

You need it because a launch does not wait. NVIDIA's own wording: "Kernel launches are asynchronous with respect to the host thread. That is, the kernel will be setup for execution on the GPU, but the host code will not wait for the kernel to complete (or even start) executing on the GPU before proceeding." So the <<<>>> statement is an enqueue. Control reaches your next line while the GPU may not have started, and any host code that reads the kernel's output, or stops a stopwatch, or exits main, has to be told to wait first. A hello world whose device printf never appears is this, and so is a std::vector that still holds its old contents.

The place it earns its keep is error reporting. A launch hands back no cudaError_t, so a fault inside the kernel surfaces at whichever runtime call looks next, which is usually a copy that did nothing wrong. Putting cudaGetLastError() and then cudaDeviceSynchronize() under a launch pins the report to the launch you are looking at: the first catches a configuration the driver refused, the second waits for the kernel and reports what it did. That pairing is what error checking looks like on day 6, and it is why no macro on this site rolls the two into one.

What it is not is a timer. Wrapping a launch in a host clock and adding a synchronize gives you an honest number for that one launch and destroys the overlap in the process, which is why the best-practices guide points host timers at CUDA events instead. It is also not something to sprinkle. cudaMemcpy already blocks, so the copy back at the end of a program is a synchronization point you already own, and a cudaDeviceSynchronize above it does nothing but serialize whatever else was in flight. The same goes for CUDA_LAUNCH_BLOCKING=1: right for finding which launch faulted, wrong to leave on, since a benchmark run under it measures a program you did not write.

Measured

Day 9 times one vectorAdd over 16,777,827 floats five ways. Three of those rows are about this call. Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), nvcc -O3 -arch=sm_75, captured 2026-08-30.

What the clock covered	ms
host clock, launch only	0.016
host clock, launch plus cudaDeviceSynchronize()	0.806
events, same kernel, after warm-up	0.786

The first row is the cost of queueing work, not the cost of doing it, and it comes in well under the kernel it claims to have timed. Add the synchronize and the same clock lands within 3 percent of the event measurement, which is the argument for the call in one line: the host stopwatch was never wrong, it was answering a different question.

There is a much larger synchronization in the same run that nobody writes. cudaSetDevice(0), the first CUDA call in the process, took 255.966 ms building the primary context before any arithmetic ran. A whole-process clock contains that, which is most of why the same program's wall time was 1426.993 ms against a 0.786 ms kernel.

Diagram

timeline-host-device, preset whole-process-vs-events. Two host lanes over one device lane. In the first, the host line runs straight past <<<>>> and the clock closes while the device bar has not started. In the second, a synchronize stretches the host bar out to the end of the device bar and the two right edges line up.

Alt text: "A host clock around a launch closes before the kernel starts. The same clock with a synchronize closes when the kernel does, and matches the event measurement."

Code

From code/day09-timing/timing.cu. One clock, read twice, so both numbers come from the same launch.

    const double tLaunch0 = hostMs();
    vectorAdd<<<blocks, kThreadsPerBlock>>>(d_a, d_b, d_out, kElems);
    const double tLaunch1 = hostMs();
    CUDA_CHECK(cudaGetLastError());
    CUDA_CHECK(cudaDeviceSynchronize());
    const double launchOnlyMs = tLaunch1 - tLaunch0;
    const double launchPlusSyncMs = hostMs() - tLaunch0;

Related terms

• __syncthreads()
• CUDA stream
• CUDA event
• kernel
• launch overhead
• CUDA error checking
• cudaMemcpy

Where you meet this

• Day 1, your first CUDA kernel, where a hello world prints nothing until this call flushes the device buffer.
• Day 6, how to check for errors in CUDA, for the two-line discipline under every launch.
• Day 9, why your GPU code looks slower than your CPU, the lesson that owns this term.
• Day 14, __syncthreads(), the barrier one level down, inside a block.
• an illegal memory access was encountered, the error this call usually surfaces and a copy usually gets blamed for.

Sources

• CUDA Programming Guide 2.1, on launches being asynchronous and on cudaSetDevice creating the primary context: https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/intro-to-cuda-cpp.html (checked 2026-08-30)
• CUDA Runtime API, Device Management, for what cudaDeviceSynchronize waits on: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__DEVICE.html (checked 2026-08-30)
• CUDA C++ Best Practices Guide, "Using CPU Timers" and "Using CUDA GPU Timers": https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html (checked 2026-08-30)
• "When to call cudaDeviceSynchronize?", 122,577 views: https://stackoverflow.com/questions/11888772/when-to-call-cudadevicesynchronize (checked 2026-08-29)

Byline

Written by: pending. Reviewed by: pending. Written on: pending. Last checked on: pending. This entry stays a draft until a named author and a different named reviewer sign it.