CUDA:__syncthreads() 在 if 语句中

Question

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

__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 个元素， block 大小 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 语句中的线程在修改 vector 的元素后，将保持等待状态，永远不会离开 if 范围。 所以……你能澄清一下发生了什么吗？在同步点之后获取的线程是否解除了在屏障处等待的线程的阻塞？ 如果您需要重现我的情况，我使用了带有 SDK 4.2 的 CUDA Toolkit 5.0 RC。非常感谢。

Answer 1

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

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

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

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

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

插入同步点的位置未定义，甚至在不同架构之间略有不同，因此同样无法保证。您将从 Nvidia 获得的唯一(非官方)评论是编译器非常擅长寻找最佳同步点。但是，通常存在一些细微的问题，可能会将最佳点向下移动得比您预期的更远，尤其是在线程提前退出的情况下。

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

您可以根据此描述了解您的特定代码段的行为方式。但请记住，所有这些都是未定义，没有任何保证，依赖特定行为可能随时破坏您的代码。

您可能想查看 PTX manual有关更多详细信息，特别是对于 bar.sync __syncthreads() 编译成的指令。亨利王的"Demystifying GPU Microarchitecture through Microbenchmarking" paper ，艾哈迈德在下面引用，也非常值得一读。即使对于现在过时的架构和 CUDA 版本，有关条件分支和 __syncthreads() 的部分似乎仍然普遍有效。