无法理解 __syncthreads()

Question

书名：

在 CUDA 中，如果存在 __syncthreads() 语句，则必须由块中的所有线程执行。当__syncthreads() 放在if 语句中时，块中的所有线程都执行包含__syncthreads() 的路径，或者它们都不执行。对于if-then-else 语句，如果每个路径都有一个__syncthreads()statement，则块中的所有线程都在then 路径上执行__syncthreads()，或者所有线程都执行elsepath。两个__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分支中，会出现死锁。

您应该避免在不同的 if 和 else 分支中使用__syncthreads，否则您必须确保整个线程块在同一个分支中运行！