1

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

看这段代码：

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

如果我为此更改第 4 行：

*dst -= 300000.0f;

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

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

我已经跑了好几次了，我总是因为这个变化而受到大约 20 秒的惩罚

更新 似乎没有使用原子的商店总是会在 L2 中产生缺失，而原子版本总是会产生命中......所以我想尝试写入一些被标记为“原子”的位置是不允许的L2 向 gmem 发出另一个请求

4

2 回答 2

2

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

如果另一个线程不打算使用该位置，那就太好了。并且它开启了让线程快速处理多个数据位置的可能性，因为如果您按顺序执行多个原子操作，线程可以将它们触发。将它们放在相邻的内存中，内存带宽可能会通过合并而减少。

于 2015-03-19T20:52:26.520 回答
1

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

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

正如开普勒白皮书所述，与费米相比，开普勒提高了全局原子的速度。

对公共全局内存地址的原子操作吞吐量提高了 9 倍，达到每个时钟一次操作。

于 2013-07-01T13:12:15.233 回答