CUDA:atomicAdd 花费太多时间，序列化线程

Question

我有一个内核，可以进行一些比较并决定两个对象是否发生碰撞。我想将碰撞对象的 id 存储到输出缓冲区。我不想在输出缓冲区中有间隙。我想将每个碰撞记录到输出缓冲区中的唯一索引中。

所以我在共享内存(本地总和)和全局内存(全局总和)中创建了一个原子变量。下面的代码显示了在发现冲突时共享变量的递增。我现在在全局内存中增加原子变量没有问题。

__global__ void mykernel(..., unsigned int *gColCnt) {
    ...

    __shared__ unsigned int sColCnt;
    __shared__ unsigned int sIndex;

    if (threadIdx.x == 0) {
        sColCnt = 0;
    }

    __syncthreads();

    unsigned int index = 0;
    if (colliding)
        index = atomicAdd(&sColCnt, 1); //!!Time Consuming!!

    __syncthreads();

    if (threadIdx.x == 0)
        sIndex = atomicAdd(gColCnt, sColCnt);

    __syncthreads();

    if (sColCnt + sIndex > outputSize) { //output buffer is not enough
        //printf("Exceeds outputsize: %d + %d > %d\n", sColCnt, sIndex, outputSize);
        return;
    }

    if (colliding) {
        output[sIndex + index] = make_uint2(startId, toId);
    }
}

我的问题是，当许多线程尝试增加原子变量时，它们会被序列化。在写前缀总和之类的东西之前，我想问一下是否有一种方法可以有效地完成这项工作。

由于出现了这一行，我的内核的运行时间从 13 毫秒增加到 44 毫秒。

我找到了一个前缀和示例代码，但由于 NVIDIA 的讨论板已关闭，其引用的链接失败。
https://stackoverflow.com/a/3836944/596547

编辑:
我也在上面添加了我的代码的结尾。事实上，我确实有等级制度。为了查看每个代码行的影响，我设置了每个对象相互碰撞的场景，极端情况和另一个几乎没有对象碰撞的极端情况。

最后，我将共享原子变量添加到全局变量 (gColCnt) 中，以告知外部冲突的数量并找到正确的索引值。我想我必须以任何方式在这里使用 atomicAdd 。

Answer 1

考虑使用并行流压缩算法，例如 thrust::copy_if .