...或者只是当前 warp 或 block 中的线程？
此外，当特定块中的线程遇到（在内核中）以下行时
__shared__ float srdMem[128];
他们会只声明这个空间一次（每个块）吗？
它们显然都是异步操作的，所以如果块 22 中的线程 23 是第一个到达该行的线程，然后块 22 中的线程 69 是最后一个到达该行的线程，线程 69 将知道它已经被声明?
【问题讨论】：
标签： cuda
__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 衡量的。因此，应该非常小心共享内存如何与启动配置一起使用。
【讨论】：
grid.sync() 一样简单，您还必须确保正确启动内核以避免@Bulat 提到的问题。工作正常，虽然它和你想象的一样慢！
我同意这里的所有答案，但我认为我们在第一个问题上遗漏了一个重要点。我没有回答第二个答案，因为它在上述答案中得到了完美的回答。
GPU 上的执行以扭曲为单位进行。 warp 是一组 32 个线程，并且在一次实例中，特定 warp 的每个线程都执行相同的指令。如果您在一个块中分配 128 个线程，则它的 (128/32 = ) 4 个扭曲用于 GPU。
现在问题变成了“如果所有线程都在执行相同的指令，那么为什么需要同步？”。答案是我们需要同步属于 SAME 块的经线。 __syncthreads 不同步经线中的线程，它们已经同步。它同步属于同一块的经线。
这就是为什么您的问题的答案是：__syncthreads 不会同步网格中的所有线程，而是同步属于一个块的线程，因为每个块都是独立执行的。
如果您想同步网格，则将您的内核 (K) 分成两个内核（K1 和 K2）并调用两者。它们将被同步（K2 将在 K1 完成后执行）。
【讨论】：
__syncthreads() 一直等待，直到同一块内的所有线程都到达命令和一个 warp 中的所有线程 - 这意味着属于一个线程块的所有 warp 必须到达该语句。
如果您在内核中声明共享内存，则该数组将仅对一个线程块可见。所以每个块都会有自己的共享内存块。
【讨论】：
shared 数组。
现有的答案已经很好地回答了__syncthreads() 的工作原理（它允许 块内 同步），我只是想添加一个更新，现在 inter 有更新的方法-块同步。自 CUDA 9.0 以来，引入了“合作组”，它允许同步整个块网格（如 Cuda Programming Guide 中所述）。这实现了与启动新内核（如上所述）相同的功能，但通常可以以较低的开销实现这一点，并使您的代码更具可读性。
【讨论】：
为了提供更多细节，除了答案，引用seibert：
更一般地说，__syncthreads() 是一种屏障原语，旨在保护您免受块内的 read-after-write 内存竞争条件的影响。
使用规则非常简单：
当线程可能读取另一个线程已写入的内存位置时，在写入之后和读取之前放置一个 __syncthreads()。
__syncthreads() 只是块内的屏障，因此它不能保护您免受全局内存中的 read-after-write 竞争条件，除非唯一可能的冲突是同一块中的线程之间。 __syncthreads() 几乎总是用于保护写后读的共享内存。
在您确定每个线程都会到达相同的 __syncthreads() 调用之前，不要在分支或循环中使用 __syncthreads() 调用。这有时可能需要您将 if 块分成几部分，以便将 __syncthread() 调用放在所有线程（包括那些未能通过 if 谓词的线程）将执行它们的顶层。
在寻找循环中的 read-after-write 情况时，在确定将 __syncthread() 调用放在何处时，它有助于在您的脑海中展开循环。例如，如果有来自不同线程的读取和写入到循环中相同的共享内存位置，您通常需要在循环结束时额外调用 __syncthreads()。
__syncthreads() 没有标记临界区，所以不要那样使用它。
不要将 __syncthreads() 放在内核调用的末尾。没必要。
许多内核根本不需要 __syncthreads()，因为两个不同的线程从不访问同一个内存位置。
【讨论】：