CUDA：__syncthreads()里面的if语句

概述我有一个关于CUDA同步的问题.特别是,我需要一些澄清if语句的同步.我的意思是,如果我把__syncthreads()放在一个if语句的范围之内,这个if语句是由块内的一小部分线程命中的,那会怎么样？我以为一些线程将保持“永远”等待不会击中同步点的其他线程.所以,我写了并执行了一些示例代码来检查： __global__ void kernel(float* vett, int n){ 我有一个关于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 2100 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()的部分似乎仍然普遍有效.

总结

以上是内存溢出为你收集整理的CUDA：__syncthreads()里面的if语句全部内容，希望文章能够帮你解决CUDA：__syncthreads()里面的if语句所遇到的程序开发问题。

如果觉得内存溢出网站内容还不错，欢迎将内存溢出网站推荐给程序员好友。