了解 CUDA 序列化和再收敛点

Question

编辑：我意识到，不幸的是，我忽略了第一个示例代码中 while 语句末尾的分号，并自己误解了它。所以实际上对于threadIdx.x != s 的线程有一个空循环，在该循环之后有一个收敛点，并且在这一点上等待所有其他线程而不增加s 变量。我将下面的原始（未更正）问题留给任何对此感兴趣的人。请注意，在第一个示例中，第二行末尾缺少一个分号，因此 s++ 与循环体没有任何共同之处。

--

我们在 CUDA 课上学习序列化，我们的老师告诉我们这样的代码：

__shared__ int s = 0;
while (s != threadIdx.x)
    s++; // serialized code

最终会导致硬件死锁，因为 nvcc 编译器在 while (s != threadIdx.x) 和 s++ 语句之间放置了一个重新收敛点。如果我理解正确，这意味着一旦一个线程达到再收敛点，该线程就会停止执行并等待其他线程，直到它们也到达该点。然而，在这个例子中，这永远不会发生，因为线程 #0 进入了 while 循环的主体，在没有增加 s 变量的情况下到达了重新收敛点，并且其他线程陷入了无限循环。

一个可行的解决方案应该如下：

__shared__ int s = 0;
while (s < blockDim.x)
    if (threadIdx.x == s)
        s++; // serialized code

这里，一个块中的所有线程都进入循环体，都评估条件，只有线程#0在第一次迭代中增加s变量（循环继续）。

我的问题是，如果第一个示例挂起，为什么第二个示例有效？更具体地说，if 语句只是另一个分歧点，就汇编语言而言，应该编译成与循环中的条件相同的条件跳转指令。那么为什么在第二个示例中s++ 之前没有任何重新收敛点，而实际上它在语句之后立即消失了呢？

在其他来源中，我只发现每个分支独立计算不同的代码 - 例如在if/else 语句中，首先计算if 分支，将所有else 分支线程屏蔽在同一个warp 中，然后其他线程在第一次等待时计算else 分支。在 if/else 语句之后有一个再收敛点。那么为什么第一个示例冻结，没有将循环分成两个分支（一个线程的true 分支和一个等待中的所有其他线程的false 分支）？

谢谢。

Answer 1

将再收敛点放在对while (s != threadIdx.x) 和s++; 的调用之间是没有意义的。它扰乱了程序流程，因为一段代码的重新收敛点应该可以在编译时被所有线程访问。下图显示了您的第一段代码的流程图以及可能和不可能的重新收敛点。

关于this answer 关于通过SSY 指令记录收敛点，我在下面创建了类似于您的第一段代码的简单内核

__global__ void kernel_1() {
    __shared__ int s;
    if(threadIdx.x==0)
        s = 0;
    __syncthreads();
    while (s == threadIdx.x)
        s++; // serialized code
}

并使用 -O3 将其编译为 CC=3.5。下面是使用cuobjdumbinary 工具输出观察CUDA 程序集的结果。结果是：

我不是阅读 CUDA 程序集的专家，但我可以在 0038 和 00a0 行中看到 while 循环条件检查。在00a8 行，如果满足while 循环条件，则分支到0x80 并再次执行代码块。再收敛点的引入在0058行引入0xb8行作为再收敛点，在出口附近的循环条件检查之后。

总体而言，尚不清楚您要通过这段代码实现什么目标。同样在第二段代码中，再收敛点应该在while循环代码块之后（我不是指while和if之间）。

Answer 2

它“挂起”的原因既不是硬件死锁也不是分支，至少不是直接的。您为一个或多个线程生成了一个无限循环（正如已经怀疑的那样）。

在您的示例中，实际上并没有收敛点。由于您不使用任何同步，因此实际上没有任何线程在等待。 while循环在这里发生的事情几乎是一个忙碌的等待。 只有当所有线程都返回时，内核才会完成。由于您有一个（或多个）无限循环（偶然甚至可能没有 - 然而这不太可能）内核永远不会完成。

您声明了一个共享变量 s。该变量对块内的所有线程都是已知的。 使用您的 while 语句，您基本上会说（对每个线程）：递增 s 直到它达到您的（本地）线程 ID 的值。由于所有线程都在并行递增 s，因此您引入了竞争条件。 示例：

1. 列表项
2. 线程 5 正在循环并检查 s 是否变为 5
3. s 是 4
4. 两个线程递增s，变为6
5. 同时线程 5 仅到达其循环的末尾。
6. 现在它到达下一个循环迭代并检查 s，它不是 5。
7. 线程 5 将永远无法完成，因为您通过 == 检查并且 s 的值已经超过了线程 id 的值。

您的解决方案也很混乱，因为每个线程都连续执行序列化代码（这可能毕竟是意图 - 尽管这实际上很奇怪）：

1. 线程 0 将执行序列化代码
2. 之后，线程1将执行序列化代码
3. 等等

大多数示例显示了一个程序，其中每个线程处理一些代码，然后所有线程同步，只有单个线程执行更多代码（也许它需要所有线程的结果）。 因此，您的第二个示例“有效”，因为没有线程陷入无限循环，但是我想不出有人会使用这样的代码的原因， 因为它令人困惑，而且根本不平行。