在内核中使用 cudaDeviceSynchronize() 进行全局同步

Question

我刚刚阅读了有关动态并行性的文档。我想知道的是：我可以在内核中使用cudaDeviceSynchronize() 来同步当前在该设备上运行的所有块吗？

文档说：

来自任何线程的 CUDA 运行时操作，包括内核启动，在线程块中都是可见的。这意味着父网格中的调用线程可以在由该线程启动的网格、线程块中的其他线程或在同一线程块中创建的流上执行同步。

此外：

在网格内创建的流和事件存在于线程块范围内，但在创建它们的线程块之外使用时具有未定义的行为。

我的问题基本上是NO。但是由于cudaDeviceSynchronize() 为整个设备使用全局流，我不确定该流是否对设备上的所有线程可见并且相同，无论它们属于哪个块或启动。这样我就可以在内核中使用cudaDeviceSynchronize() 进行全局同步。

Answer 1

没有。没有办法安全地进行设备范围的同步。

编程指南 (Link) 的 C.3.1.4 节：

cudaDeviceSynchronize() 函数将同步线程块中任何线程启动的所有工作，直到调用 cudaDeviceSynchronize() 为止。

它没有说明与其他线程块交互。

一般来说，CUDA 中的全局同步会导致问题，因为超额订阅方法最常用于填充 GPU 的工作。要同步的块的数量通常比设备上容纳的要多得多，因此每个块的上下文都必须交换进出全局内存，从而破坏性能。

如果您知道自己有特殊情况，可以使用一些技巧来解决此问题，但通常，同步块的最简单和最有效的方法是退出内核并启动一个新内核。