cuda block 同步

Question

我有 b 个块，每个块有 t 个线程。
我可以用

 __syncthreads()

同步特定块中的线程。例如

__global__ void aFunction()
{
    for(i=0;i<10;i++)
    {
       //execute something
        __syncthreads();
    }
}

但我的问题是同步所有块中的所有线程。我怎样才能做到这一点？

Answer 1

在 CUDA 9 中，NVIDIA 引入了协作组的概念，允许您同步属于该组的所有线程。这样的组可以跨越网格中的所有线程。通过这种方式，您将能够同步所有块中的所有线程:

#include <cuda_runtime_api.h> 
#include <cuda.h> 
#include <cooperative_groups.h>

cooperative_groups::grid_group g = cooperative_groups::this_grid(); 
g.sync();

您需要 Pascal(计算能力 60)或更新的架构来同步网格。此外，还有更具体的要求。见: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#grid-synchronization-cg

Basic functionality, such as synchronizing groups smaller than a thread block down to warp granularity, is supported on all architectures, while Pascal and Volta GPUs enable new grid-wide and multi-GPU synchronizing groups.

来源: https://devblogs.nvidia.com/parallelforall/cuda-9-features-revealed/

在 CUDA 9 之前，没有本地方法来同步所有块中的所有线程。事实上，CUDA 中块的概念是一些可能只有在其他一些块已经结束其工作后才会启动，例如，如果它运行的 GPU 太弱而无法并行处理它们。
如果您确保不会产生太多块，则可以尝试在它们之间同步所有块，例如通过使用原子操作主动等待。然而，这很慢，吃掉你的 GPU 内存 Controller ，被认为是“黑客”，应该避免。
因此，如果您不针对 Pascal(或更新的)架构，我建议的最佳方法是在同步点简单地终止您的内核，然后启动一个可以继续您的工作的新内核。在大多数情况下，它实际上比使用提到的 hack 执行得更快(或至少 - 以相似的速度)。