CUDA，使用reduction查找最大值，错误

Question

这是我的代码，它试图在一个块中找到最大值为50的数组。我已将数组填充为64。

对于线程1-31，我有正确的maxVal打印输出，但对于线程32-49，这是一个完全随机的数字。我不知道我做错了什么。

顺便说一句。我认为我不需要在展开中_sync每一行，但显然我必须这样做。对此有什么建议吗？

提前感谢您的帮助。

//block size = 50


__syncthreads();

if (tid<32){

    cptmp[tid]=(cptmp[tid]< cptmp[tid+32]) ? cptmp[tid+32] : cptmp[tid];__syncthreads();    
    cptmp[tid]=(cptmp[tid]< cptmp[tid+16]) ? cptmp[tid+16] : cptmp[tid];__syncthreads();
    cptmp[tid]=(cptmp[tid]< cptmp[tid+8]) ? cptmp[tid+8] : cptmp[tid];  __syncthreads();    
    cptmp[tid]=(cptmp[tid]< cptmp[tid+4]) ? cptmp[tid+4] : cptmp[tid];  __syncthreads();
    cptmp[tid]=(cptmp[tid]< cptmp[tid+2]) ? cptmp[tid+2] : cptmp[tid];  __syncthreads();    
    cptmp[tid]=(cptmp[tid]< cptmp[tid+1]) ? cptmp[tid+1] : cptmp[tid];  __syncthreads();

}

__syncthreads();

//if (tid==0) {
    maxVal=cptmp[0];
    if(bix==0 && biy==0) cuPrintf(" max:%f x:%d y:%d\n", maxVal, blockIdx.x, blockIdx.y);
//}

Answer 1

这里是一个更有效的(至少在费米GPU上)和正确的代码使用易失性。将T替换为您的类型(或使用模板)：

if (tid<32) {
    volatile T *c = cptmp;
    T t = c[tid];
    c[tid] = t = (t < c[tid+32]) ? c[tid+32] : t;
    c[tid] = t = (t < c[tid+16]) ? c[tid+16] : t;
    c[tid] = t = (t < c[tid+ 8]) ? c[tid+ 8] : t;
    c[tid] = t = (t < c[tid+ 4]) ? c[tid+ 4] : t;
    c[tid] = t = (t < c[tid+ 2]) ? c[tid+ 2] : t;
    c[tid] = t = (t < c[tid+ 1]) ? c[tid+ 1] : t;
}

为什么这样更有效率呢？在没有__syncthreads()的情况下，为了保证正确性，我们必须使用指向共享内存的易失性指针。但这迫使编译器“遵守”所有对共享内存的读写操作--它不能优化寄存器中的任何内容并将其保存在寄存器中。因此，通过显式地始终将c[tid]保存在临时t中，我们每行代码就可以节省一个共享内存负载。由于Fermi是一种加载/存储体系结构，它只能使用寄存器作为指令操作数，这意味着我们每行可以节省一条指令，或者说总共节省6条指令(我估计总共大约25% )。

在旧的T10/GT200体系结构和更早的体系结构上，您的代码(具有易失性，没有__syncthreads())将同样有效，因为该体系结构可以直接从共享内存中获取每个指令的一个操作数。

如果您更喜欢if而不是?:，则此代码应该是等效的

if (tid<32) {
    volatile T *c = cptmp;
    T t = c[tid];
    if (t < c[tid+32]) c[tid] = t = c[tid+32];
    if (t < c[tid+16]) c[tid] = t = c[tid+16];
    if (t < c[tid+ 8]) c[tid] = t = c[tid+ 8];
    if (t < c[tid+ 4]) c[tid] = t = c[tid+ 4];
    if (t < c[tid+ 2]) c[tid] = t = c[tid+ 2];
    if (t < c[tid+ 1]) c[tid] = t = c[tid+ 1];
}

Answer 2

不要在不同的代码中使用__syncthreads()！来自给定块的所有线程或没有线程都应该到达同一位置上的每个__syncthreads()。

来自单个warp的所有线程(32个线程)都是隐式同步的，因此您不需要__syncthreads()来将它们放在一起。但是，如果您担心一个线程的共享内存写入可能不会被同一warp的另一个线程看到，请使用__threadfence_block()。

阐述__threadfence_block()的重要性。考虑以下两行：

cptmp[tid]=(cptmp[tid]< cptmp[tid+2]) ? cptmp[tid+2] : cptmp[tid];
cptmp[tid]=(cptmp[tid]< cptmp[tid+1]) ? cptmp[tid+1] : cptmp[tid];

它可能会编译成类似这样的内容：

int tmp; //assuming that cptmp is an array of int-s
tmp=cptmp[tid];
tmp=(tmp<cptmp[tid+2])?cptmp[tid+2]:tmp;
tmp=(tmp<cptmp[tid+1])?cptmp[tid+1]:tmp;
cptmp[tid]=tmp;

虽然它对于单线程代码是正确的，但对于CUDA显然是失败的。

为了防止这样的优化，您可以将cptmp数组声明为volatile，或者在代码行之间添加此__threadfence_block()。该函数确保在该函数存在之前，同一块中的所有线程都能看到当前线程写入的共享内存。

存在一个类似的__threadfence()函数来确保全局内存可见性。

Answer 3

对于像我一样在未来遇到这个线程的每个人，除了harrism的答案之外，还有一个建议-从性能的角度来看，考虑shuffle操作可能是值得的，所以使用单个warp从64个元素中获取最大值的更新代码将如下所示：

auto localMax = max(c[tid], c[tid + 32]);    
for (auto i = 16; i >= 1; i /= 2)
{
    localMax = max(localMax, __shfl_xor(localMax, i));
}
c[tid] = localMax;

只需要从全局内存中进行两次读取和一次写入，因此非常简洁。