CUDA Reduction - 原子与单线程求和

Posted 2023-02-16

【中文标题】CUDA Reduction - 原子与单线程求和

【英文标题】：CUDA Reduction - atomic vs single thread summation

【发布时间】：2013-07-24 09:46:06

【问题描述】：

我最近使用 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 来对共享和求和，但我想了解为什么这不起作用。当将操作限制为单个线程时，原子添加是非常无意义的。简单的总和应该可以正常工作！

【问题讨论】：

__syncthreads() 只同步同一块中的线程，而不是跨不同块。我认为不正确的结果是由于同步问题。通过atomicAdd，您正在执行__syncthreads() 缺少的不同块之间的同步。

确实，当我在 for 循环中添加一个 __syncthreads() 时，简单的求和就起作用了！但我不明白。我在全局数组上只使用一个线程进行求和，那么我为什么要关心在 for 循环中同步呢？

好的，我想我明白了！进入循环时不一定会写入全局数组，因为所有块都不会同步。那么“全局”同步线程的命令是什么？

操作数x[k]是不同块的计算结果：x[0]是块0的结果，x[1]是块1的结果，等等。怀疑线程0 可能在某些块真正完成计算之前开始将它们相加。试试下面的。将第二个代码 sn -p 放在不同的内核中，以便强制同步，然后尝试 sum[0] += x[k]; 行是否有效。

关于你的新问题，CUDA 没有跨区块的安全同步机制。

【参考方案1】：

__syncthreads()只同步同一个block中的线程，不跨不同block，CUDA没有跨block的安全同步机制。

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

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

【讨论】：

另请注意，原子并不是确保块之间同步的安全方法，尤其是当您的块数量增加时。当我尝试使用 10^5 个元素，每个块 512 个线程因此 196 个块时，我得到“nan”结果。它只是有助于减慢求和过程，让其他线程有时间从其他块写入它们的结果，但这绝对不是应对它的好方法。另一个内核更好

你绝对是对的。我已经修改了我的答案，删除了关于 atomicAdd 的句子。

还有一点补充：事实上，我在没有将内核一分为二的情况下搜索了一个解决方案，因为在设备函数中将在其他操作中使用缩减。但是在设备功能结束时都没有同步。不过我找到了一个很好的方法，它使用 10^5 个元素：在将共享内存检索到全局内存时在 if 块中放置一个 __threadfence()，这样每个块的每个线程 0 将确保所有线程都能够看到它的写作。你可以添加更健壮的东西。事实上，在 cuda prog 中有一个例子。指南，B5 节。

【参考方案2】：

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

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

我建议你看看Thrust 和CUB。