1

我最近使用 CUDA 测试了减少算法（例如，您可以在http://www.cuvilib.com/Reduction.pdf找到的算法，第 16 页）。但最后，我遇到了不使用原子性的麻烦。所以基本上我做每个块的总和并将其存储到共享数组中。然后我将它返回到全局数组 x（tdx 是 threadIndex.x，i 是全局索引）。

if(i==0){
        *sum = 0.; // Initialize to 0
    }
__syncthreads();
if (tdx == 0){       
    x[blockIdx.x] = s_x[tdx]; //get the shared sums in global memory
}
__syncthreads();

然后我想对前 x 个元素求和（与我有块一样多）。当使用原子性时它工作正常（与 cpu 的结果相同），但是当我使用下面的注释行时它不起作用并且经常产生“nan”：

if(i == 0){    
    for(int k = 0; k < gridDim.x; k++){
        atomicAdd(sum, x[k]); //Works good
       //sum[0] += x[k]; //or *sum += x[k]; //Does not work, often results in nan
    }
}

现在实际上我直接使用 atomicadd 来对共享和求和，但我想了解为什么这不起作用。当将操作限制为单个线程时，原子添加是非常无意义的。简单的总和应该可以正常工作！

4

2 回答 2

4

__syncthreads()仅同步同一块中的线程，而不是跨块同步，并且 CUDA 没有跨块的安全同步机制。

不正确的结果是由于同步问题造成的。操作数x[k]是来自不同块的计算结果：x[0]是块的结果0，x[1]是块的结果1等。线程0可以在某些块真正完成计算之前开始将它们相加。

您应该将第二个代码片段放在不同的内核中，以便强制执行同步，并且该行sum[0] += x[k];现在可以工作。

于 2013-07-24T10:30:57.237 回答
0

正如已经指出的那样，您的问题是由于第一次通过后缺少同步，因为您无法在块之间同步。有一个很好的关于减少随工具包提供的示例代码的演练。

话虽如此，我强烈建议人们不要在库代码中存在此类原语的情况下编写归约内核（或其他原语，如扫描）。最好将您的精力投入到其他地方并在可用的地方重用现有的优化代码。如果您这样做当然是为了学习，这不适用！

我建议你看看Thrust和CUB。

于 2013-07-24T13:07:57.500 回答