CUDA：if 语句中的 __syncthreads()答案

【问题标题】：CUDA: __syncthreads() inside if statementsCUDA：if 语句中的 __syncthreads()
【发布时间】：2012-09-13 05:30:30
【问题描述】：

我有一个关于 CUDA 同步的问题。特别是，我需要对 if 语句中的同步进行一些说明。我的意思是，如果我将 __syncthreads() 放在块内的一小部分线程命中的 if 语句的范围内，会发生什么？我认为一些线程将保持“永远”等待其他不会达到同步点的线程。所以，我编写并执行了一些示例代码来检查：

__global__ void kernel(float* vett, int n)
{
    int index = blockIdx.x*blockDim.x + threadIdx.x;
    int gridSize = blockDim.x*gridDim.x;

    while( index < n )
    {   
        vett[index] = 2;
        if(threadIdx.x < 10)
        {
            vett[index] = 100;
            __syncthreads();
        }
        __syncthreads();

        index += gridSize;
    }
}

令人惊讶的是，我观察到输出非常“正常”（64 个元素，块大小 32）：

100 100 100 100 100 100 100 100 100 100 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2
100 100 100 100 100 100 100 100 100 100 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2

所以我通过以下方式稍微修改了我的代码：

__global__ void kernel(float* vett, int n)
{
    int index = blockIdx.x*blockDim.x + threadIdx.x;
    int gridSize = blockDim.x*gridDim.x;

    while( index < n )
    {   
        vett[index] = 2;
        if(threadIdx.x < 10)
        {
            vett[index] = 100;
            __syncthreads();
        }
        __syncthreads();
            vett[index] = 3;
        __syncthreads();

        index += gridSize;
    }
}

输出是：

3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 
3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3

再次，我错了：我认为 if 语句中的线程，在修改了向量的元素后，会保持等待状态，永远不会超出 if 范围。所以……你能澄清一下发生了什么吗？在同步点之后获取的线程是否会解除阻塞在屏障处等待的线程？如果您需要重现我的情况，我使用了带有 SDK 4.2 的 CUDA Toolkit 5.0 RC。提前非常感谢。

【问题讨论】：

给回答你问题的人打勾（接受的答案）。

标签： c parallel-processing cuda synchronizing

【解决方案1】：

简而言之，行为是未定义。所以它有时可能会做你想做的事，也可能不会，或者（很可能）只会挂起或崩溃你的内核。

如果你真的很好奇事情是如何在内部工作的，你需要记住线程不是独立执行的，而是一次执行一个 warp（32 个线程的组）。

这当然会产生条件分支的问题，条件分支在整个经纱中不会统一评估。这个问题是通过执行两条路径来解决的，一个接一个，每个都禁用那些不应该执行该路径的线程。 IIRC 在现有硬件上首先采用分支，然后在未采用分支的地方执行路径，但这种行为是 undefined 的，因此不能保证。

这种单独的路径执行一直持续到编译器可以确定两个单独执行路径的所有线程都可以到达的某个点（“重新收敛点”或“同步点”）。当第一个代码路径的执行到达这一点时，它会停止并执行第二个代码路径。当第二条路径到达同步点时，所有线程再次启用，并从那里统一继续执行。

如果在同步之前遇到另一个条件分支，情况会变得更加复杂。这个问题通过一堆仍然需要执行的路径来解决（幸运的是，堆栈的增长是有限的，因为我们最多可以为一个 warp 拥有 32 个不同的代码路径）。

插入同步点的位置是未定义，甚至在架构之间略有不同，所以同样不能保证。您将从 Nvidia 获得的唯一（非官方）评论是编译器非常擅长寻找最佳同步点。然而，经常有一些微妙的问题可能会比您预期的更进一步地降低最佳点，尤其是在线程提前退出的情况下。

现在要了解 __syncthreads() 指令的行为（在 PTX 中转换为 bar.sync 指令），重要的是要意识到该指令不是每个线程执行的，而是一次对整个 warp 执行（不管是否禁用任何线程）因为只有块的扭曲需要同步。 warp 的线程已经在同步执行，进一步的同步要么无效（如果所有线程都启用），要么在尝试从不同的条件代码路径同步线程时导致死锁。

您可以按照自己的方式从这个描述到您的特定代码的行为方式。但请记住，这一切都是未定义，没有任何保证，依赖特定行为可能随时破坏您的代码。

您可能想查看PTX manual 以了解更多详细信息，尤其是__syncthreads() 编译成的bar.sync 指令。 Henry Wong 的"Demystifying GPU Microarchitecture through Microbenchmarking" paper，下面由ahmad 引用，也很值得一读。即使对于现在已经过时的架构和 CUDA 版本，关于条件分支和__syncthreads() 的部分似乎仍然普遍有效。

【讨论】：

【解决方案2】：

CUDA 模型是 MIMD，但当前的 NVIDIA GPU 以扭曲粒度而不是线程实现 __syncthreads()。这意味着，这些warps inside a thread-block 不一定是同步的threads inside a thread-block。 __syncthreds() 等待线程块的所有“扭曲”到达障碍或退出程序。有关详细信息，请参阅Henry Wong's Demistifying paper。

【讨论】：

那篇论文确实是一个很好的参考。我忘了它也包括条件分支。

【解决方案3】：

您不能使用__syncthreads()，除非该语句在一个线程块内的所有线程中都达到，始终如此。来自programming guide (B.6)：

__syncthreads() 允许在条件代码中使用，但前提是条件在整个线程块中的计算结果相同，否则代码执行可能会挂起或产生意外的副作用。

基本上，您的代码不是格式良好的 CUDA 程序。

【讨论】：

当然不是！但我写它只是为了检查它的行为。
@biagiop1986：嗯……你有一段库代码和硬件，附带一个文档，上面写着“你不能做 X”。现在你问我们，公众，如果你做X会发生什么——我们应该怎么知道？问卖家！知道程序将是非良构的还不够吗？
这取决于...说我应该在我的程序中避免这样的代码是正确的，因为它的格式不正确（而且，我发誓，我会的！），但我很好奇'如何'。而且，此外，我经常发现这里对问题的解释比供应商解释要好得多。所以，我会回到这里，而不是其他人，询问我将来遇到的每一个编码问题。堆栈溢出是最好的！顺便说一句，谢谢大家。

【解决方案4】：

__syncthreads() 用于同步块内的线程。这意味着块中的所有线程都将等待所有线程都完成后再继续。

考虑一个块中有一些线程进入 if 语句而一些没有的情况。那些等待的线程，将被阻塞；永远的等待。

一般来说，将同步放在 if 条件语句中并不是一种好的风格。最好避免它，如果你有它，重新设计你的代码。同步的目的是确保所有线程一起进行，为什么首先使用 if 语句将它们过滤掉？

添加，如果需要跨块同步。需要重新启动内核。

【讨论】：

【解决方案5】：

最好避免在 if 条件中使用 __syncthreads()。您可以使用 for 循环和__syncthreads() 在 for 循环后重写代码。

【讨论】：