55

...或者只是当前扭曲或块中的线程？

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

__shared__  float srdMem[128];

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

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

4

5 回答 5

74

该__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 同步或隐式同步。因此它们都是同步的。

使用此技术的示例（来源）：

CPU同步

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

于 2013-03-06T08:27:36.697 回答
19

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

GPU 上的执行以经纱为单位进行。warp 是一组 32 个线程，并且在一次实例中，特定 warp 的每个线程都执行相同的指令。如果您在一个块中分配 128 个线程，则其 (128/32 = ) 4 个扭曲用于 GPU。

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

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

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

于 2016-12-31T15:20:45.647 回答
17

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

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

于 2013-03-06T06:28:50.787 回答
7

现有的答案已经很好地回答了如何__syncthreads()工作（它允许块内同步），我只是想添加一个更新，现在有更新的块间同步方法。自 CUDA 9.0 以来，引入了“合作组”，它允许同步整个块网格（如Cuda Programming Guide中所述）。这实现了与启动新内核相同的功能（如上所述），但通常可以以较低的开销实现这一点，并使您的代码更具可读性。

于 2020-04-05T17:35:44.247 回答
2

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

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

使用规则非常简单：

  1. 当线程可能读取另一个线程已写入的内存位置时，在写入之后和读取之前放置一个 __syncthreads()。

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

  3. 不要在分支或循环中使用 __syncthreads() 调用，直到您确定每个线程都会到达相同的 __syncthreads() 调用。这有时可能需要您将 if 块分成几个部分，以将 __syncthread() 调用放在所有线程（包括那些 if 谓词失败的线程）将执行它们的顶层。

  4. 在寻找循环中的 read-after-write 情况时，在确定将 __syncthread() 调用放在哪里时，它有助于在您的脑海中展开循环。例如，如果有来自不同线程的读取和写入到循环中相同的共享内存位置，您通常需要在循环结束时额外调用 __syncthreads()。

  5. __syncthreads() 没有标记关键部分，所以不要那样使用它。

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

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

于 2021-01-08T21:41:42.647 回答