如何在开普勒实施全球原子操作？我使用gmem而不是使用原子来降低性能

Question

我想知道开普勒全局原子的实现。

看到这段代码：

1. if (threadIdx.x < workers) {
2.    temp = atomicAdd(dst, temp + rangeOffset);
3.    if (isLastPartialCalc(temp)) {                            
4.        atomicAdd(dst,-300000.0f);
5.    }
6. }

如果我改变第4行：

* etc - = 300000.0f;

性能更低！更改是安全的，因为没有更多的线程会写入此值（输出相同）。

内核使用原子：~883us内核直接使用gmem：~903us

我跑过几次，总是得到这个〜20us的惩罚

更新似乎没有使用原子的商店总是产生L2的错过，而原子版总是产生一个命中...所以我想尝试写一些标记（或某事）与“原子”的位置是不允许的L2并且它向gmem发出另一个请求

Answer 1

与第二次全局原子访问相比，此缓存行更新显然更昂贵（在您的特定代码中）。

在Kepler GK110（例如K20）上从单个SM到全局存储器的单个全局原子访问实际上非常快。

如Kepler white paper所示，与费米相比，开普勒提高了全球原子的速度。

对于公共全局存储器地址的原子操作吞吐量每时钟提高9倍至一个操作。

Answer 2

原子学具有“即发即忘”的语义。这意味着内核调用原子操作并让实际的原子操作由缓存执行（而不是在SM上），内核将继续执行下一条指令而无需等待实际的原子操作完成。这仅在原子操作没有返回值时才有效，本例中就是这种情况。即发即弃的语义让SM继续进行计算，将原子的计算卸载到缓存中。

如果另一个线程不打算使用该位置，这很好。它开辟了一个让线程快速处理多个数据位置的可能性，因为如果顺序有几个原子操作，线程就可以解除它们。将它们放在相邻的内存中，通过合并可能会降低内存带宽。