相关疑难解决方法(0)

__syncthreads()在我有意使用线程丢弃的块中使用是否安全return？

文档说明__syncthreads() 必须由块中的每个线程调用,否则它将导致死锁,但实际上我从未经历过这样的行为.

示例代码:

__global__ void kernel(float* data, size_t size) {
    // Drop excess threads if user put too many in kernel call.
    // After the return, there are `size` active threads.
    if (threadIdx.x >= size) {
        return;
    }

    // ... do some work ...

    __syncthreads(); // Is this safe?

    // For the rest of the kernel, we need to drop one excess thread
    // After the return, there are `size - 1` active threads …

跟进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];
}