【发布时间】：2012-07-20 19:39:58
【问题描述】：
我有一个内核，它进行一些比较并决定两个对象是否碰撞。我想将碰撞对象的 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。
【问题讨论】：
-
atomicAdd按定义进行序列化，因此只有在预测碰撞将是稀疏的时才应该依赖它。也许您可以重组您的计算以分层使用原子：首先，在每个线程块中累积到__shared__变量。在后处理中（例如，在上述内核的第三个__syncthreads之后），您可以将每个块的冲突累积到全局内存中的单个变量中。 -
事实上我确实有一个层次结构。但是同一块中的线程也在 atomicAdd 上为 shared 变量序列化，至少对于第一个极端情况，每个对象都相互碰撞。
-
www.cuvilib.com/Reduction.pdf 我找到了 M. Harris 的教程。我会尽量利用它。
-
其实原子性能取决于你的硬件。你在什么硬件上运行？
-
我有两个 Gtx460。我写了减少代码，但似乎我需要前缀和（几乎没有写，需要一些修复，银行冲突等）。完全碰撞示例花费了一半的时间。但是首先保存碰撞然后使用推力::remove_if 进行压缩也得到了很好的结果。不过我有点卡住了。
标签： cuda shared-memory atomic elapsedtime prefix-sum