CUDA 中的线程同步答案

【问题标题】：Threads syncronization in CUDACUDA 中的线程同步
【发布时间】：2016-02-28 14:50:33
【问题描述】：

我有一个 3D 块的 3D 网格，在每个块中，我需要在块的“z”层上按顺序计算。换句话说，我想首先执行所有 (x,y,0) 线程，然后执行所有 (x,y,1) 等。我需要逐层执行我的线程（按 z 轴计算层数）。我知道函数 __syncthreads()，但我不知道如何按照我想要的方式同步线程。

更新：

 __global__ void Kernel(//some params)
 {


      //some code 
      __syncthreads();
 }

它同步块中的所有线程。但是我需要执行z = 0的所有线程，然后是z = 1的所有线程，等等。

【问题讨论】：

这个问题需要更多细节，但我怀疑答案是：将你的线程块定义为需要同步的线程组的大小。换句话说，将您的图层拆分为单独的块。
@harrism，我知道我可以这样做。但在这种情况下，我将不得不启动内核几次，这会对性能产生负面影响。
为什么需要多次启动内核？您可以有一个 2D 网格，其中 Y 维度表示块中的（前）Z 层。并在网格内使用 2D x/y 块。那么你只需要块内的__syncthreads()，blockIdx.y对应你之前在threadIdx.z中的...
@harrism，因为我的算法中确实有一个 3D 网格。
好的，感谢您的澄清。你应该把它放在问题中。

标签： c++ cuda

【解决方案1】：

您可以使用一个简单的循环，并指定您希望在每次迭代中完成工作的线程。比如：

for (int z = 0; z < zmax; z++) {
    if (threadIdx.z == z) { 
        //do whatever with x and y
    }
    __syncthreads();
}

在每次迭代中，具有特定 z-index 的线程执行指令，而其他线程处于空闲状态；在每次迭代结束时，所有线程都会同步。

【讨论】：

你确定它会起作用吗？这似乎可以保证所有具有相同z的线程将被同步，但不能保证它们将按z的升序执行。首先，z = 0，然后 z = 1 等的线程。
@harrism，据我所知，这与安全或不安全使用无关，这可能会导致错误的结果。很有可能所有 z = 1 的线程在 z = 0 的线程之前到达其同步点。因此第 1 层将比第 0 层更早执行。如果我错了，请纠正我。
@Max __syncthreads() 在继续执行之前等待所有线程到达该行。所以所有线程都进入循环，在第一次迭代中，只有threadIdx.z==0 的线程完成工作，而其他线程处于空闲状态。在循环结束时，所有线程同步。至于它的安全性，我不明白为什么它不安全（假设你所有的线程都进入循环），但我又不是专家。
按照你的代码编写方式，你的 __syncthreads() 是 if() 中唯一的东西，这就是为什么我说它不安全（只有 threadId.z == z 线程会同步）。我误解了您的意思是您的 C++ 注释是可执行代码。我不擅长像编译器一样解析！我已经编辑了您的代码以使其更正确，并将我的反对票改为赞成票。但是，我要求您在答案中添加一个散文解释，以阐明代码如何解决问题。
@Noel，非常感谢。它真的很有帮助。我自己总是使用括号，这就是我误解你的原因。

【解决方案2】：

__syncthreads() 阻塞块内的线程。如果您坚持使用__syncthreads()，则必须确保块大小与a（x * y）相同，这不一定是可能的，即使这样，也不能保证块的顺序。实现目标的另一种方法是每层启动一个内核，并在内核之间同步（即阻止内核并发）。这当然真的取决于你的确切内核做了什么，以及这样分解它是否可行。

【讨论】：

我不确定你说得对。例如，我有一个 2 x 2 x 2 块。首先我需要执行以下线程：（0,0,0），（1,0,0），（0,1,0），（1,1,0）。这是第一层。第二个是（0,0,1），（1,0,1），（0,1,1），（1,1,1）。所以对于第一层 z = 0，对于第二层 z = 1。
@Max 是的，我明白了。第一个内核将处理第一层（我假设每一层都略大于 2x2）；第二层，第二层等等。第 n 层处理 (x,y,n)。
我认为这可能有点矫枉过正，请参阅我上面关于使用 2D 块的 2D 网格而不是多次启动的评论。
@harrism 很难与你争论，但我会试一试。据我了解，OP 要求处理层n 的所有线程在层n+1 的第一个线程之前完成。在这种情况下，即使是幼稚的表示 (fig. 5) 也会在每个 SM 上运行多个块。因此，不同的层将并行运行。我认为您对 OP 的第一条评论错过了每一层单独运行的要求（？）。我同意我的建议可能有点矫枉过正并且会产生大量开销。
是的，我错过了 OP 中的“then”这个词！