c++ - 在 CUDA 线程中填充计数 'buckets'

Question

在我的程序中，我通过体素网格跟踪大量粒子。粒子与体素的比例是任意的。在某个时刻，我需要知道哪些粒子位于哪些体素中，以及有多少粒子位于哪些体素中。具体来说，体素必须确切地知道其中包含哪些粒子。由于我不能在 CUDA 中使用任何类似 std::vector 的东西，我正在使用以下算法(在高级别):

• 分配一个整数数组，其大小为体素数
• 为所有粒子启动线程，确定每个粒子所在的体素，并在我的“桶”数组中增加适当的计数器
• 分配一个大小为粒子数的指针数组
• 计算每个体素在这个新数组中的偏移量(将其前面的体素中的粒子数相加)
• 将粒子以有序的方式放置在数组中(我使用此数据来加速稍后的操作。速度的提高非常值得增加内存使用量)。

虽然这在第二步就崩溃了。我在 CUDA 中编程的时间不长，只是发现线程之间同时写入全局内存中的同一位置会产生未定义的结果。这反射(reflect)在我在 buckets 中大多得到 1，偶尔得到 2。这是我在这一步中使用的代码草图:

__global__ void GPU_AssignParticles(Particle* particles, Voxel* voxels, int* buckets) {
    int tid = threadIdx.x + blockIdx.x*blockDim.x;

    if(tid < num_particles) { // <-- you can assume I actually passed this to the function :)
        // Some math to determine the index of the voxel which this particle
        // resides in.
        buckets[index] += 1;
    }
}

我的问题是，在 CUDA 中生成这些计数的正确方法是什么？

另外，有没有办法在体素内存储对粒子的引用？我看到的问题是体素内的粒子数量不断变化，因此几乎每一帧都必须重新分配和重新分配新阵列。

Answer 1

虽然计算桶计数可能有更有效的解决方案，但第一个可行的解决方案是使用您当前的方法，但使用原子增量。这样一次只有一个线程以原子方式递增桶计数(在整个网格上同步):

if(tid < num_particles) {
    // ...
    atomicAdd(&buckets[index], 1);
}