77

何时cudaDeviceSynchronize真正需要调用该函数？

据我从 CUDA 文档中了解到，CUDA 内核是异步的，所以似乎我们应该cudaDeviceSynchronize在每次内核启动后调用。但是，我尝试了相同的代码（训练神经网络），有和没有任何cudaDeviceSynchronize，除了时间测量之前的一个。我发现我得到了相同的结果，但速度提高了 7-12 倍（取决于矩阵大小）。

所以，问题是是否有任何理由使用cudaDeviceSynchronize时间间隔测量。

例如：

  • 在将数据从 GPU 复制回主机之前是否需要cudaMemcpy？

  • 如果我做矩阵乘法，比如

    C = A * B
D = C * F

我应该放在cudaDeviceSynchronize两者之间吗？

从我的实验看来我没有。

为什么cudaDeviceSynchronize程序这么慢？

4

4 回答 4

69

尽管 CUDA 内核启动是异步的，但所有放在一个流中的与 GPU 相关的任务（这是默认行为）都是按顺序执行的。

所以，例如，

kernel1<<<X,Y>>>(...); // kernel start execution, CPU continues to next statement
kernel2<<<X,Y>>>(...); // kernel is placed in queue and will start after kernel1 finishes, CPU continues to next statement
cudaMemcpy(...); // CPU blocks until memory is copied, memory copy starts only after kernel2 finishes

因此，在您的示例中，不需要cudaDeviceSynchronize. 但是，它可能对调试有用，以检测哪个内核导致了错误（如果有的话）。

cudaDeviceSynchronize可能会导致一些减速，但 7-12x 似乎太多了。可能是时间测量存在问题，或者内核速度非常快，并且显式同步的开销相对于实际计算时间来说是巨大的。

于 2012-08-09T18:22:48.273 回答
18

适合使用的一种情况cudaDeviceSynchronize()是当您cudaStream运行多个 s 时，您希望它们交换一些信息。一个真实的例子是量子蒙特卡罗模拟中的并行回火。在这种情况下，我们希望确保每个流在开始相互传递消息之前已经完成了一组指令的运行并获得了一些结果，否则我们最终会传递垃圾信息。使用此命令会大大降低程序速度的原因是cudaDeviceSynchronize()强制程序在继续之前等待设备上所有流中所有先前发出的命令完成（来自 CUDA C 编程指南）。正如您所说，内核执行通常是异步的，因此当 GPU 设备正在执行您的内核时，CPU 可以继续处理一些其他命令，向设备发出更多指令等，而不是等待。但是，当您使用此同步命令时，CPU 会被迫闲置，直到所有 GPU 工作完成后才能执行任何其他操作。这种行为在调试时很有用，因为由于设备代码的异步执行（无论是在一个流中还是在多个流中），您可能会在看似“随机”的时间发生段错误。cudaDeviceSynchronize()将强制程序在继续之前确保流的内核/memcpys 是完整的，这样可以更容易地找出发生非法访问的位置（因为在同步期间会出现故障）。

于 2012-08-09T18:20:20.807 回答
12

当您希望您的 GPU 开始处理一些数据时，您通常会执行内核调用。当您这样做时，您的设备（GPU）将开始执行您告诉它执行的任何操作。但是，与主机（CPU）上的正常顺序程序不同，它将继续执行程序中的下一行代码。cudaDeviceSynchronize 使主机（CPU）等待，直到设备（GPU）完成执行您已启动的所有线程，因此您的程序将继续执行，就好像它是一个正常的顺序程序一样。

在小型简单程序中，当您使用 GPU 进行计算时，您通常会使用 cudaDeviceSynchronize，以避免请求结果的 CPU 和完成计算的 GPU 之间的时序不匹配。使用 cudaDeviceSynchronize 可以更轻松地编写程序，但有一个主要缺点：您的 CPU 一直处于空闲状态，而 GPU 进行计算。因此，在高性能计算中，您通常会努力让 CPU 在等待 GPU 完成时进行计算。

于 2014-01-20T16:45:14.303 回答
0

您可能还需要cudaDeviceSynchronize()在从内核启动内核后调用（动态并行）。

从这篇文章CUDA 动态并行 API 和原则：

如果父内核需要子内核计算的结果来完成自己的工作，则必须通过显式同步 using 来确保子网格已完成执行，然后再继续cudaDeviceSynchronize(void)。该函数等待之前由调用它的线程块启动的所有网格完成。由于嵌套，它还确保线程块启动的网格的任何后代都已完成。

...

请注意，执行内核启动构造时，全局内存的视图是不一致的。这意味着在下面的代码示例中，没有定义子内核是读取并打印值 1 还是 2。为了避免竞争条件，子内核可以读取的内存不应在内核启动后由父内核写入，而是在显式同步之前。

__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();
}
于 2021-12-04T22:16:25.317 回答