CUDA:if语句中的__syncthreads()

Question

我有一个关于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范围.那么......你能澄清一下发生了什么吗？在同步点之后获取的线程是否会阻塞在屏障处等待的线程？如果您需要重现我的情况,我使用CUDA Toolkit 5.0 RC和SDK 4.2.非常感谢提前.

Answer 1

简而言之,行为是不确定的.所以它有时可能会做你想要的,或者它可能没有,或者(很可能)只会挂起或崩溃你的内核.

如果你真的很好奇内部的工作方式,你需要记住线程不是独立执行,而是一次一个warp(32个线程组).

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

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

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

插入同步点的位置未定义,甚至在架构之间略有不同,因此无法保证.您将从Nvidia获得的唯一(非官方)评论是编译器非常擅长找到最佳同步点.然而,通常存在微妙的问题,可能会使最佳点进一步下降,尤其是如果线程提前退出.

现在要了解__syncthreads()指令的行为(转换为bar.syncPTX中的指令),重要的是要意识到每个线程不执行该指令,而是立即执行整个warp(无论是否禁用任何线程)或不)因为只需要同步一个块的warp.warp的线程已经同步执行,并且进一步同步将无效(如果所有线程都已启用)或在尝试同步来自不同条件代码路径的线程时导致死锁.

您可以按照此描述的方式使用您的特定代码行为.但请记住,所有这些都是未定义的,没有任何保证,并且依赖于特定行为可能会随时破坏您的代码.

您可能需要查看PTX手册以获取更多详细信息,尤其bar.sync是__syncthreads()编译指令.Henry Wong的"通过Microbenchmarking揭开GPU微架构"的论文,下面由ahmad引用,也值得一读.即使对于现在过时的体系结构和CUDA版本,关于条件分支的部分__syncthreads()看起来仍然普遍有效.

Answer 2

CUDA模型是MIMD,但是当前的NVIDIA GPU __syncthreads()以warp粒度而不是线程实现.这意味着,这些warps inside a thread-block人不一定同步threads inside a thread-block.__syncthreds()等待所有'warp'的线程块击中障碍或退出程序.有关详细信息,请参阅Henry Wong的Demistifying文章.

Answer 3

__syncthreads()除非在一个线程块内的所有线程中都到达该语句，否则不得使用该语句。来自编程指南(B.6)：

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

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