__syncthreads()是否同步网格中的所有线程？

Question

...或只是当前warp或块中的线程？

此外,当特定块中的线程遇到(在内核中)以下行时

__shared__  float srdMem[128];

他们只会宣布一次这个空间(每个街区)吗？

它们都显然是异步操作所以如果块22中的线程23是到达该线的第一个线程,然后块22中的线程69是到达该线的最后一个线程,则线程69将知道它已经被声明了？

Answer 1

该__syncthreads()命令是块级同步屏障.这意味着当块中的所有线程到达屏障时使用它是安全的.它也可以__syncthreads()在条件代码中使用,但只有当所有线程都相同地评估这样的代码时,否则执行可能会挂起或产生意外的副作用[4].

使用示例__syncthreads():( 来源)

__global__ void globFunction(int *arr, int N) 
{
    __shared__ int local_array[THREADS_PER_BLOCK];  //local block memory cache           
    int idx = blockIdx.x* blockDim.x+ threadIdx.x;

    //...calculate results
    local_array[threadIdx.x] = results;

    //synchronize the local threads writing to the local memory cache
    __syncthreads();

    // read the results of another thread in the current thread
    int val = local_array[(threadIdx.x + 1) % THREADS_PER_BLOCK];

    //write back the value to global memory
    arr[idx] = val;        
}

要同步网格中的所有线程,目前没有本机API调用.在网格级别上同步线程的一种方法是使用连续的内核调用,因为此时所有线程都从同一点结束并重新开始.它通常也称为CPU同步或隐式同步.因此它们都是同步的.

使用此技术的示例(来源):

关于第二个问题.是的,它确实声明了每个块指定的共享内存量.考虑每个SM测量可用共享内存的数量.所以要非常小心的是如何共享内存与一起使用的启动配置.

Answer 2

__syncthreads() 等待直到同一个块中的所有线程都到达命令并且warp中的所有线程 - 这意味着属于threadblock的所有warp必须到达该语句.

如果在内核中声明共享内存,则该数组仅对一个线程块可见.所以每个块都有自己的共享内存块.

Answer 3

我同意这里的所有答案,但我认为我们在第一个问题中遗漏了一个重要的观点.我没有回答第二个答案,因为它在上面的答案中得到了完美的回答.

GPU上的执行以warp为单位进行.warp是一组32个线程,并且在一次实例中,特定warp的每个线程执行相同的指令.如果在一个块中分配128个线程,则为GPU(128/32 =)4个warp.

现在问题变成"如果所有线程都执行相同的指令,那么为什么需要同步？".答案是我们需要同步属于SAME块的warp .__syncthreads不会同步warp中的线程,它们已经同步.它同步属于同一块的warp.

这就是为什么你的问题的答案是:__syncthreads不同步网格中的所有线程,但属于一个块的线程作为每个块独立执行.

如果要同步网格,则将内核(K)分成两个内核(K1和K2)并同时调用它们.它们将被同步(K2将在K1完成后执行).

Answer 4

现有答案在回答__syncthreads()工作原理方面做得很好（它允许块内同步），我只想添加一个更新，现在有更新的块间同步方法。从 CUDA 9.0 开始，引入了“合作组”，它允许同步整个块网格（如Cuda 编程指南中所述）。这实现了与启动新内核相同的功能（如上所述），但通常可以以较低的开销这样做并使您的代码更具可读性。