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

Question

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

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

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

Answer 1

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

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

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