CUDA 数组到数组和答案

【问题标题】：CUDA array-to-array sumCUDA 数组到数组和
【发布时间】：2011-02-19 23:32:24
【问题描述】：

我有这样一小段代码：

typedef struct {
  double sX;
  double sY;
  double vX;
  double vY;
  int rX;
  int rY;
  int mass;
  int species;
  int boxnum;
} particle;

typedef struct {
  double mX;
  double mY;
  double count;
  int rotDir;
  double cX; 
  double cY; 
  int superDir;
} box;
//....
int i;
for(i=0;i<PART_COUNT;i++) {
    particles[i].boxnum = ((((int)(particles[i].sX+boxShiftX))/BOX_SIZE)%BWIDTH+BWIDTH*((((int)(particles[i].sY+boxShiftY))/BOX_SIZE)%BHEIGHT));
}
for(i=0;i<PART_COUNT;i++) {
    //sum the momenta
    boxnum = particles[i].boxnum;
    boxes[boxnum].mX += particles[i].vX*particles[i].mass;
    boxes[boxnum].mY += particles[i].vY*particles[i].mass;
    boxes[boxnum].count++;
}

现在，我想将它移植到 CUDA。第一步很简单；将计算分散到一堆线程中是没有问题的。问题是第二个。由于任何两个粒子在任何同一个盒子中的可能性都相同，因此我不确定如何对其进行分区以避免冲突。

粒子数在 10,000 到 10,000,000 之间，盒子的数量在 1024 到 1048576 之间。

想法？

【问题讨论】：

particle.boxnum 值是如何分布的？如果您从随后存储的粒子中访问particle.boxnum，那么boxnumbers 是否也是连续的？
随机。我在想我可能不得不重写这个东西的一些基本结构，以便我对它有更多的了解（例如，对那个粒子列表进行排序以确保这个条件）。

标签： cuda parallel-processing

【解决方案1】：

您可以尝试使用atomicAdd 操作来修改您的boxes 数组。全局内存上的原子操作非常慢，但同时由于两个原因，不可能进行任何涉及共享内存的优化：

假设粒子boxnum 的属性particles[0]..particles[n] 没有排序并且不位于任何小边界（在块大小范围内），您无法预测从全局加载哪些框内存变成共享内存。你必须先收集所有的 boxnumbers..
如果您尝试收集所有 boxnumber，则不能使用包含所有可能 boxnumber 的数组作为索引，因为有太多的 box 无法放入共享内存。因此，您必须使用队列收集索引（通过数组、指向下一个空闲槽的指针和原子操作来实现），但是您仍然会遇到冲突，因为相同的 boxnumber 可能会在您的队列中出现多次。

结论：atomicAdd 至少会给你正确的行为。尝试一下并测试性能。如果您对性能不满意，请考虑是否有另一种方法来执行相同的计算，从而从共享内存中获益。

【讨论】：

好的，谢谢。我担心我可能不得不做类似的事情，但不确定如何，确切地说。

【解决方案2】：

作为替代方案，您可以启动 2D 块网格。

blocks.x = numParticles / threadsPerBlock / repeatPerBlock.

blocks.y = numOfBoxes / 1024;

当且仅当 boxnum 位于 1024 * blockIdx.y 和 1024 * (blockIdx.y + 1) 之间时，每个块在共享内存中执行原子加法；

这之后是沿blocks.x的减少

这可能会也可能不会比 atomicAdd 在全局内存上更快，因为数据被读取 blocks.y 次。但是，如果“粒子”在排序过程中按 boxnum 排序，然后是分区过程，则可以解决此问题。

可能还有其他几种方法可以做到这一点，但由于问题大小差异很大，您最终可能不得不编写 2-3 种针对给定大小范围进行优化的不同方法。

【讨论】：