相关疑难解决方法(0)

我想知道如何退出一个线程,其线程索引很大.我看到两种可能性:

int i = threadIdx.x;
if(i >= count)
    return;
// do logic

要么

int i = threadIdx.x;
if(i < count) {
    // do logic
}

我知道,两者都是正确的,但哪一个会影响性能呢？

跟进Q:EarlyExit和DroppedThreads

根据以上链接,下面的代码应该是死锁.
请解释为什么这不会死锁.(费米的Cuda 5)

__device__ int add[144];
__device__ int result;

add<<<1,96>>>();  // the calling 

__global__ void add() {
 for(idx=72>>1; idx>0; idx>>=1) {
  if(thrdIdx < idx) 
   add[thrdIdx]+= add[thrdIdx+idx];
  else
   return;
  __syncthreads();
 }

 if(thrdIdx == 0)
  result= add[0];
}

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 呢？或者它是文档中的错误？