cuda块同步

Question

我有b个块,每个块有t个线程.我可以用

 __syncthreads()

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

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

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

Answer 1

在CUDA 9中,NVIDIA引入了协作组的概念,允许您同步属于该组的所有线程.这样的组可以跨越网格中的所有线程.这样您就可以同步所有块中的所有线程:

#include <cooperative_groups.h>

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

您需要Pascal或更新的架构来同步网格.

所有体系结构都支持基本功能,例如将小于线程块的组同步到warp粒度,而Pascal和Volta GPU支持新的网格范围和多GPU同步组.

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

在CUDA 9之前,没有本地方法来同步所有块中的所有线程.实际上,CUDA中块的概念是某些块可能仅在其他块已经结束其工作之后启动,例如,如果它运行的GPU太弱而无法并行处理它们.

如果确保不会生成太多块,则可以尝试在它们之间同步所有块,例如通过主动等待使用原子操作.然而,这很慢,占用你的GPU内存控制器,被认为是"黑客",应该避免.

因此,如果您不针对Pascal(或更新版本)架构,我建议的最佳方法是在同步点简单地终止内核,然后启动一个新内核,这将继续您的工作.在大多数情况下,它实际上会比使用上面提到的黑客更快(或至少 - 以相同的速度)执行.