When to call cudaDeviceSynchronize?

user_1588226 0
Viewed 88909

when is calling to the cudaDeviceSynchronize function really needed?.

As far as I understand from the CUDA documentation, CUDA kernels are asynchronous, so it seems that we should call cudaDeviceSynchronize after each kernel launch. However, I have tried the same code (training neural networks) with and without any cudaDeviceSynchronize, except one before the time measurement. I have found that I get the same result but with a speed up between 7-12x (depending on the matrix sizes).

So, the question is if there are any reasons to use cudaDeviceSynchronize apart of time measurement.

For example:

  • Is it needed before copying data from the GPU back to the host with cudaMemcpy?

  • If I do matrix multiplications like

    C = A * B
D = C * F

should I put cudaDeviceSynchronize between both?

From my experiment It seems that I don't.

Why does cudaDeviceSynchronize slow the program so much?

gpgpu cuda gpu
4 Answers

You might also need to call cudaDeviceSynchronize() after launching kernels from kernels (Dynamic Parallelism).

From this post CUDA Dynamic Parallelism API and Principles:

If the parent kernel needs results computed by the child kernel to do its own work, it must ensure that the child grid has finished execution before continuing by explicitly synchronizing using cudaDeviceSynchronize(void). This function waits for completion of all grids previously launched by the thread block from which it has been called. Because of nesting, it also ensures that any descendants of grids launched by the thread block have completed.

...

Note that the view of global memory is not consistent when the kernel launch construct is executed. That means that in the following code example, it is not defined whether the child kernel reads and prints the value 1 or 2. To avoid race conditions, memory which can be read by the child should not be written by the parent after kernel launch but before explicit synchronization.

__device__ int v = 0;

__global__ void child_k(void) {
  printf("v = %d\n", v);
}

__global__ void parent_k(void) {
  v = 1;
  child_k <<< 1, 1 >>>> ();
  v = 2; // RACE CONDITION
  cudaDeviceSynchronize();
}
user_30410680
Related