CUDA:__syncthreads()里面的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范围. 解决方法
总之,行为是不明确的.所以它有时可以做你想要的,或者它可能不会,或者(很可能)会挂起或崩溃你的内核.
如果你真的很好奇内部的内容,你需要记住,线程不能独立执行,而是一个warp(一组32个线程). 这当然会导致条件分支的问题,其中条件不会在整个warp中统一评估.通过执行这两个路径,一个接一个地解决这个问题,每个路径被禁用,不应该执行该路径. IIRC在现有硬件上首先采用分支,然后执行不采用分支的路径,但是这种行为是未定义的,因此不能保证. 路径的这种单独执行持续到某种程度,编译器可以确定其保证由两个单独的执行路径(“再融合点”或“同步点”)的所有线程达到.当第一代码路径的执行到达该点时,它被停止并且代替执行第二代码路径.当第二条路径到达同步点时,所有线程将再次启用,并且执行从那里均匀地继续. 如果在同步之前遇到另一个条件分支,情况会变得更加复杂.这个问题是通过一堆仍然需要执行的路径解决的(幸运的是,堆栈的增长是有限的,因为我们可以为一个warp提供最多32个不同的代码路径). 插入同步点的位置是不确定的,甚至在体系结构之间略有不同,所以再也没有保证.从Nvidia获得的唯一(非官方)评论是,编译器非常适合找到最佳的同步点.然而,常常存在微妙的问题,可能会使您的最优点进一步下降,尤其是如果线程提前退出. 现在要了解__syncthreads()指令的行为(它转换成PTX中的一个bar.sync指令),重要的是要意识到这个指令不是每个线程都执行的,而是一次完整的转换(不管是否有任何的)线程被禁用或不被禁用),因为只有块的经线需要同步. warp的线程已经同步执行,并且进一步的同步将无效果(如果所有线程都被使能),或者当尝试从不同的条件代码路径同步线程时,会导致死锁. 您可以从此描述中了解您的特定代码行为.但请记住,所有这些都是未定义的,没有保证,依赖于具体行为可能会随时破坏您的代码. 您可能需要查看PTX manual的更多细节,特别是对于__syncthreads()编译的 (编辑:李大同) 【声明】本站内容均来自网络,其相关言论仅代表作者个人观点,不代表本站立场。若无意侵犯到您的权利,请及时与联系站长删除相关内容! |