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

Question

__syncthreads()函数是否同步网格中的所有线程，还是只同步当前warp或块中的线程？

此外，当特定块中的线程在内核中遇到以下行时：

__shared__  float srdMem[128];

他们只会在每个块中声明这个空间一次吗？

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

- Wuschelbeutel Kartoffelhuhn

---

共享内存为每个块单独分配，但不是同时分配。当SM实际开始执行块时，共享内存会在那时分配。 - sgarizvi

Answer 1

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

使用 __syncthreads() 的示例：(source)

__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同步或隐式同步。因此，它们都被同步了。

使用此技术的示例（source）：

关于第二个问题。是的，它确实声明了每个块指定的共享内存量。请注意，可用的共享内存数量是按SM计算的。因此，在使用共享内存时，应非常小心地与启动配置一起使用。

- KiaMorot

---

警告，这是危险的代码。来自@harrism在您所引用的同一来源中。 - Nikos Yotis

在网格中同步所有线程是有问题的，因为不能保证它们会同时执行。GPU 只能运行有限数量的线程，如果内核执行需要太多的线程块，则必须在启动新块之前完成其中一些块。这个限制取决于 GPU 型号以及软件环境（用户可能同时执行多个 GPU 程序），因此试图同步所有线程块的内核非常危险。正确的方法是完成一个内核并启动另一个。 - Bulat

@Bulat，我还没有机会使用比Fermi更新的硬件。您知道自Kepler以来引入的动态并行性和同时执行多个内核是否可以在某种程度上解决这个问题吗？ - KiaMorot

D.P. 允许在内核内运行内核，并等待其执行。虽然它可以用于实现更复杂的同步场景，但它无法避免基本问题 - GPU 实现基于任务的并行性，并且您永远不知道两个任务（内核实例）是并行还是顺序执行。如果您确实需要块间同步，请参见 http://eprints.cs.vt.edu/archive/00001087/01/TR_GPU_synchronization.pdf - Bulat

2016年，@Bulat在写道，同步网格中的所有线程存在问题，他是正确的。现在我们有了合作式网格，可以让您安全地进行同步。只需使用“grid.sync()”就可以轻松完成，还必须确保正确启动内核以避免@Bulat提到的问题。虽然速度较慢，但运行良好，正如您所预期的那样。 - Eyal

Answer 2

我同意这里所有的答案，但是我认为我们在第一个问题上缺少了一个重要的点。我不回答第二个问题，因为它已经在以上答案中得到了完美的回答。

GPU上的执行以warp为单位进行。一个warp是32个线程的一组，在某个时间点上，每个warp的每个线程执行相同的指令。如果您在一个块中分配了128个线程，则对于GPU而言，它是（128/32 =）4个warps。

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

这就是为什么你的问题的答案是：__syncthreads不会同步网格中的所有线程，而是同步属于一个块的线程，因为每个块都是独立执行的。

如果您想同步一个网格，请将您的内核（K）分成两个内核（K1和K2）并调用两个内核。它们将被同步执行（K2将在K1完成后执行）。

- Adarsh

---

__syncthreads不会同步warp中的线程，因为它们已经被隐式地同步了。这个说法在Volta架构中已经过时了。现在不能再假定warp会被隐式地同步。 - paleonix

Answer 3

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

如果在内核中声明共享内存，则该数组仅对一个线程块可见。因此，每个块都有自己的共享内存块。

- hubs

---

这实际上是不正确的。shared 数组为设备中的每个块分配了空间。 - KiaMorot

@KiaMorot：我觉得你误解了一些东西。这个答案没有任何问题。共享内存是块作用域，这就是答案所说的，也是你的常规说法。哪里有矛盾之处？ - talonmies

Answer 4

现有的答案已经很好地解释了__syncthreads()的工作方式（它允许块内同步），我只想补充一下最新的方法用于块间同步。自CUDA 9.0以来，引入了“协作组”，可以同步整个块网格（如Cuda编程指南中所述）。这实现了与启动新内核相同的功能（如上所述），但通常可以在更低的开销下完成，并使您的代码更易读。

- TimD1

Answer 5

为了提供更多细节，除了答案之外，引用seibert的话：

更一般地说，__syncthreads()是一个屏障原语，旨在保护您免受块内读写竞争条件的影响。

使用规则非常简单：

1. 当有可能一个线程读取了另一个线程写入的内存位置时，在写操作后和读操作前加上__syncthreads()。

2. __syncthreads()只是块内的屏障，因此除非唯一可能的冲突是在同一块中的线程之间，否则它无法保护您免受全局内存读写竞争条件的影响。__syncthreads()几乎总是用于保护共享内存的读写操作。

3. 不要在分支或循环中使用__syncthreads()，直到您确定每个线程都将到达相同的__syncthreads()调用为止。这有时需要将if块分成多个部分，以便在所有线程（包括未通过if谓词的线程）都执行它们的顶级处放置__syncthread()调用。

4. 在查找循环中的读写后情况时，有助于在脑海中展开循环，以找出在哪里放置__syncthread()调用。例如，如果循环中来自不同线程的读写访问了相同的共享内存位置，则通常需要在循环末尾额外添加一个__syncthreads()调用。

5. __syncthreads()不会标记关键部分，因此不要像那样使用它。

6. 不要在内核调用的末尾放置__syncthreads()。没有必要这样做。

7. 许多内核根本不需要__syncthreads()，因为两个不同的线程从不访问相同的内存位置。

- Soleil