__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()的部分似乎仍然普遍有效.
