无法理解 __syncthreads()

Posted 2023-04-15

【中文标题】无法理解 __syncthreads()

【英文标题】：not able to understand __syncthreads()

【发布时间】：2014-01-03 23:50:54

【问题描述】：

书名：

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

没有给出if和if-else-then案例的例子，所以我无法理解这个概念。请用简单的话解释一下这两种情况。

PS：我是并行编程和 CUDA 的初学者。

提前致谢。

【参考方案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，否则您必须确保整个线程块在同一个分支中运行！

【讨论】：

好吧，是不是就像这是两个等待点，前 16 个到达一个，后 16 个到达另一个，但这两个等待点彼此不知道。我这样想对吗

是的，你可以这样表达。前 16 位无法联系到其他的__syncthreads 来填写声明。