如何有效地从 CUDA 中的线程收集数据？答案

【问题标题】：How to efficiently gather data from threads in CUDA?如何有效地从 CUDA 中的线程收集数据？
【发布时间】：2012-06-24 07:39:54
【问题描述】：

我有一个在 CUDA 中求解方程组的应用程序，我确信每个线程最多可以找到 4 个解，但是我如何将其复制回主机？

我将一个具有足够空间的巨大数组传递给所有线程存储 4 个解决方案（每个解决方案 4 个双倍），另一个包含每个线程的解决方案数量，但这是一个幼稚的解决方案，并且是当前的瓶颈我的内核。

我真的很喜欢优化这个。主要问题是将每个线程的可变数量的解决方案连接到一个数组中。

【问题讨论】：

如果我对你的程序有更多的了解，我会更容易提供帮助。据我所知（自从我弄乱 cuda 以来已经有一年了，所以我可能错了），memcopies 是检索信息的唯一方法，而且速度很慢。什么卡上的 cuda 版本是什么？
我有一个 CUDA 4.0 和 4.2。
代码太大，放在这里。我同意 cudaMemCpy 这是获得结果的唯一方法，但我可以避免垃圾复制。

标签： optimization cuda gpu gpgpu

【解决方案1】：

您正在寻找的功能称为流压缩。

您可能确实需要提供一个数组，该数组包含每个线程 4 个解决方案的空间，因为尝试以紧凑的形式直接存储结果可能会在线程之间创建如此多的依赖关系，以至于能够通过更少的复制获得性能较长的内核执行时间会丢失返回主机的数据。例外情况是，如果几乎所有线程都找不到解决方案。在这种情况下，您也许可以使用原子操作来维护数组的索引。因此，对于找到的每个解决方案，您将其存储在索引处的数组中，然后使用原子操作来增加索引。我认为为此使用 atomicAdd() 是安全的。在存储结果之前，线程将使用 atomicAdd() 将索引增加一。 atomicAdd() 返回旧值，线程可以将旧值作为索引存储结果。

但是，考虑到更常见的情况，即有相当数量的结果，最好的解决方案是将压缩操作作为单独的步骤执行。一种方法是使用thrust::copy_if。有关更多背景信息，请参阅this question。

【讨论】：