如果不能在条件分支中调用 CUDA，如何在 CUDA 中减少__syncthreads？

How to reduce in CUDA if __syncthreads can't be called inside conditional branches?

本文关键字：CUDA 如果不 syncthreads 不能 分支 调用 如果 条件      更新时间：2023-10-16

NVIDIA建议的约简方法在条件分支内使用__syncthreads()，例如:

if (blockSize >= 512) { if (tid < 256) { sdata[tid] += sdata[tid + 256]; } __syncthreads(); }

或

for (unsigned int s=blockDim.x/2; s>32; s>>=1)
{
    if (tid < s)
        sdata[tid] += sdata[tid + s];
    __syncthreads();
}

在第二个例子中，__syncthreads()在for循环体内，这也是一个条件分支。

然而，关于SO的一些问题在条件分支内提出了__syncthreads()的问题(例如，我可以在丢弃线程后使用__syncthreads()吗?条件同步线程&死锁(或不死锁)，答案说条件分支中的__syncthreads()可能导致死锁。因此，NVIDIA建议的约简方法可能会死锁(如果相信答案所基于的文档)。

此外，如果_syncthreads()不能在条件分支中使用，那么恐怕很多基本的操作都会被阻塞，约简只是一个例子。

那么如何在没有使用__syncthreads()条件分支的情况下减少CUDA ?或者是文档中的错误?

限制不

__syncthreads不能用于条件分支

限制是

__syncthreads不能用于不会被所有线程同时遍历的分支

请注意，在和给出的示例中，__syncthreads没有被依赖于线程ID(或某些每线程数据)的条件所覆盖。在第一种情况下，blockSize是一个模板参数，它不依赖于线程ID。在第二种情况下，它同样位于if . 之后。

是的，for循环的s > 32是一个条件，但是它的真值不依赖于线程或它的数据。blockdim.x对于所有线程都是相同的。所有线程都将执行对s的完全相同的修改。这意味着所有线程将在其控制流的完全相同的点到达__syncthreads。这完全没问题。

另一种情况是，不能使用__syncthreads，这是一个对某些线程为真而对其他线程为假的条件。在这种情况下，您必须关闭所有条件才能使用__syncthreads。所以不用这个:

if (threadIdx.x < SOME_CONSTANT)
{
  operation1();
  __syncthreads();
  operation2();
}

你必须这样做:

if (threadIdx.x < SOME_CONSTANT)
{
  operation1();
}
__syncthreads();
if (threadIdx.x < SOME_CONSTANT)
{
  operation2();
}

您给出的两个示例也演示了这一点:在调用__syncthreads之前关闭依赖于线程id的条件。

相关文章：