填充计数'buckets' CUDA 线程

Filling counting 'buckets' in CUDA threads

本文关键字：CUDA 线程 buckets 填充      更新时间：2023-10-16

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

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

不过，这在第二步中被分解了。 我在 CUDA 中编程的时间不长，只是发现线程之间同时写入全局内存中的同一位置会产生未定义的结果。 这反映在我大多在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 中生成这些计数的正确方法是什么？

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

尽管可能有更有效的解决方案来计算存储桶计数，但第一个有效的解决方案是使用当前的方法，但使用原子增量。这样，一次只有一个线程以原子方式递增存储桶计数（在整个网格上同步）：

if(tid < num_particles) {
    // ...
    atomicAdd(&buckets[index], 1);
}
相关文章：