1

这是我的代码尝试减少以在块中找到最大 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);
//}
4

3 回答 3

3

这是使用 volatile 的更高效（至少在 Fermi 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 架构和更早的架构上，您的代码（具有 volatile 且没有 __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];
}
于 2011-06-30T04:06:58.917 回答
2

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

来自单个经纱（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()函数来确保全局内存可见性。

于 2011-06-28T20:10:41.460 回答
1

对于将来会偶然发现此线程的每个人，就像我所做的那样，除了 harrism 答案之外，这里还有一个建议 - 从性能的角度来看，考虑 shuffle 操作可能是值得的，所以更新的代码可以让 max out of 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;

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

于 2014-02-13T00:54:33.270 回答