相关疑难解决方法(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 …

CUDA编程指南指出了这一点

条件代码中允许__syncthreads(),但仅当条件在整个线程块中进行相同的求值时,否则代码执行可能会挂起或产生意外的副作用.

因此,如果我需要在一个块上使用条件分支同步线程,其中一些线程可能会或可能不会采用包含该__syncthreads()调用的分支,这是否意味着它不起作用？

我想象可能存在各种各样的情况,你可能需要这样做; 例如,如果您有二进制掩码并需要有条件地对像素应用某个操作.比如说,if (mask(x, y) != 0)然后执行包含的代码__syncthreads(),否则什么都不做.怎么办？

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