无法理解__syncthreads()

Question

这本书引用:

在CUDA中,__syncthreads()语句(如果存在)必须由块中的所有线程执行.当a __syncthreads()放在一个if语句中时,块中的所有线程都会执行包含__syncthreads()或不包含它们的路径.对于if-then-else语句,如果每个路径都有一个__syncthreads()语句,则块中的所有线程都在路径__syncthreads()上执行,或者所有线程then都执行该else路径.这两个__syncthreads()是不同的屏障同步点.如果块中的线程执行then路径而另一个线程执行else路径,则它们将在不同的屏障同步点处等待.他们最终会永远等待对方.程序员有责任编写代码以满足这些要求.

没有给出例子if和if-else-then案例,所以我无法理解这个概念.请用简单的话语解释我的情况.

PS:我是并行编程和CUDA的初学者.

提前致谢 .

Answer 1

假设您有一个内核,该内核使用一个由32个线程组成的线程块启动.

kernel<<<1,32>>>()

内核的代码如下:

__global__ void kernel()
{
  if (threadIdx.x < 16)
  {
    // do something
    __syncthreads();
  }
  else
  {
    // do something
    __snycthreads();
  }
}

线程块的前16个线程将运行if语句.另外16个else语句.如果前16个线程中的每个线程都到达__syncthreads,则它们将阻塞,直到整个线程块到达该语句.但是这种情况永远不会出现,因为其他16个线程卡在了else分支中,并且会出现死锁.

你应该避免__syncthreads在不同的if和else分支中使用,或者你必须确保整个线程块在同一个分支中运行!