如果在条件分支内无法调用__syncthreads,如何减少CUDA？

Question

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循环体,它也是一个条件分支.

然而,一些在做题提高的问题,__syncthreads()内部条件的分支(例如我能已经下降线程后使用__syncthreads()？和条件syncthreads和死锁(或没有)),答说,__syncthreads()在条件分支可能导致陷入僵局.因此,NVIDIA建议的缩减方法可能会陷入僵局(如果相信答案所依据的文档).

此外,如果_syncthreads()不能在条件分支内部使用,那么我担心许多基本操作被阻止,减少只是一个例子.

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

Answer 1

限制不是

__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的条件已关闭.