4

我有一个内核，它进行一些比较并决定两个对象是否碰撞。我想将碰撞对象的 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 。

4

2 回答 2

1

nvidia 博客文章相关：http ://devblogs.nvidia.com/parallelforall/gpu-pro-tip-fast-histograms-using-shared-atomics-maxwell/

于 2015-03-21T07:11:18.583 回答
1

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

于 2012-07-22T14:24:24.670 回答