加入收藏 | 设为首页 | 会员中心 | 我要投稿 李大同 (https://www.lidatong.com.cn/)- 科技、建站、经验、云计算、5G、大数据,站长网!
当前位置: 首页 > 百科 > 正文

CUDA:__syncthreads()里面的if语句

发布时间:2020-12-16 03:05:39 所属栏目:百科 来源:网络整理
导读:我有一个关于CUDA同步的问题.特别是,我需要一些澄清if语句的同步.我的意思是,如果我把__syncthreads()放在一个if语句的范围之内,这个if语句是由块内的一小部分线程命中的,那会怎么样?我以为一些线程将保持“永远”等待不会击中同步点的其他线程.所以,我写了
我有一个关于CUDA同步的问题.特别是,我需要一些澄清if语句的同步.我的意思是,如果我把__syncthreads()放在一个if语句的范围之内,这个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.非常感谢.

解决方法

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

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

这当然会导致条件分支的问题,其中条件不会在整个warp中统一评估.通过执行这两个路径,一个接一个地解决这个问题,每个路径被禁用,不应该执行该路径. IIRC在现有硬件上首先采用分支,然后执行不采用分支的路径,但是这种行为是未定义的,因此不能保证.

路径的这种单独执行持续到某种程度,编译器可以确定其保证由两个单独的执行路径(“再融合点”或“同步点”)的所有线程达到.当第一代码路径的执行到达该点时,它被停止并且代替执行第二代码路径.当第二条路径到达同步点时,所有线程将再次启用,并且执行从那里均匀地继续.

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

插入同步点的位置是不确定的,甚至在体系结构之间略有不同,所以再也没有保证.从Nvidia获得的唯一(非官方)评论是,编译器非常适合找到最佳的同步点.然而,常常存在微妙的问题,可能会使您的最优点进一步下降,尤其是如果线程提前退出.

现在要了解__syncthreads()指令的行为(它转换成PTX中的一个bar.sync指令),重要的是要意识到这个指令不是每个线程都执行的,而是一次完整的转换(不管是否有任何的)线程被禁用或不被禁用),因为只有块的经线需要同步. warp的线程已经同步执行,并且进一步的同步将无效果(如果所有线程都被使能),或者当尝试从不同的条件代码路径同步线程时,会导致死锁.

您可以从此描述中了解您的特定代码行为.但请记住,所有这些都是未定义的,没有保证,依赖于具体行为可能会随时破坏您的代码.

您可能需要查看PTX manual的更多细节,特别是对于__syncthreads()编译的bar.sync指令.黄熙来的“Demystifying GPU Microarchitecture through Microbenchmarking” paper,以下由艾哈迈德参考,也值得一读.即使现在过时的架构和CUDA版本,有关条件分支和__syncthreads()的部分似乎仍然普遍有效.

(编辑:李大同)

【声明】本站内容均来自网络,其相关言论仅代表作者个人观点,不代表本站立场。若无意侵犯到您的权利,请及时与联系站长删除相关内容!

    推荐文章
      热点阅读