0

如果没有 CUDA，我的代码只有两个 for 循环，用于计算系统中所有坐标对之间的距离，并将这些距离分类到 bin 中。

我的 CUDA 版本的问题是显然线程不能同时写入相同的全局内存位置（竞争条件？）。我最终为每个 bin 获得的值是不正确的，因为只有一个线程最终写入每个 bin。

__global__ void computePcf(
        double const * const atoms,
        double * bins,
        int numParticles,
        double dr) {

    int i = blockDim.x * blockIdx.x + threadIdx.x;

    if (i < numParticles - 1) {
        for (int j = i + 1; j < numParticles; j++) {
            double r = distance(&atoms[3*i + 0], &atoms[3*j + 0]);

            int binNumber = floor(r/dr);

            // Problem line right here.
            // This memory address is modified by multiple threads
            bins[binNumber] += 2.0;
        }
    }
}

所以......我不知道该怎么做。我一直在谷歌搜索和阅读有关共享内存的信息，但问题是我不知道我将要访问哪个内存区域，直到我进行距离计算！

我知道这是可能的，因为一个叫做 VMD 的程序使用 GPU 来加速这个计算。任何帮助（甚至想法）将不胜感激。我不需要这个优化，只是功能。

4

1 回答 1

1

有多少bins[]？有什么理由bins[]需要 typedouble吗？从您的代码中并不明显。你所拥有的本质上是一个直方图操作，你可能想看看快速并行直方图技术。 推力可能是有意义的。

您的代码有几种可能的途径：

  1. 看看是否有一种方法可以重组你的算法来安排计算，使得给定的一组线程（或 bin 计算）不会相互影响。这也许可以基于排序距离来完成。

  2. 使用原子 这应该可以解决您的问题，但在执行时间方面可能会很昂贵（但由于它非常简单，您可能想尝试一下。）代替这个：

    bins[binNumber] += 2.0;

    像这样的东西：

    int * bins,
...
atomicAdd(bins+binNumber, 2);

    如果bins是 type double，你仍然可以这样做，只是有点复杂。有关如何在. _atomicAdddouble

  3. 如果数量bins很少（可能是几千或更少），那么您可以创建几组由多个线程块更新的 bin，然后在处理序列的结束。在这种情况下，您可能需要考虑使用较少数量的线程或线程块，每个线程或线程块处理多个元素，方法是在内核代码中放置一个额外的循环，以便在每个粒子处理完成后，循环跳转到下一个粒子通过添加gridDim.x*blockDim.x到i变量，并重复该过程。由于每个线程或线程块都有自己的 bin 本地副本，因此它可以在不影响其他线程访问的情况下执行此操作。

    例如，假设我只需要 1000 个 int 类型的 bin。我可以创建 1000 组垃圾箱，它们只占用大约 4 兆字节。然后，我可以为 1000 个线程中的每一个线程设置自己的 bin 集，然后 1000 个线程中的每一个线程都将拥有自己的 bin 集来更新，并且不需要原子，因为它不会干扰任何其他线程。通过让每个线程循环通过多个粒子，我仍然可以通过这种方式有效地保持机器忙碌。当所有的粒子分箱完成后，我必须将我的 1000 个分箱集加在一起，也许需要一个单独的内核调用。

于 2013-07-10T01:25:39.620 回答