c++ - 扭曲如何与原子操作一起工作？

Question

warp 中的线程在物理上并行运行，所以如果其中一个(称为线程 X)启动原子操作，其他线程会做什么？等待？这是否意味着，当线程 X 被推送到原子队列时，所有线程都将等待，获取访问权限(互斥锁)并使用内存做一些事情，内存受到该互斥锁的保护，然后才是真正的互斥锁？

有没有办法让其他线程进行某些工作，比如读取一些内存，这样原子操作会隐藏它的延迟？我的意思是，有 15 个空闲线程......我猜不太好。 Atomic 真的很慢，是吗？我怎样才能加速它？有什么模式可以使用它吗？

共享内存的原子操作是否锁定银行或整个内存？
例如(没有互斥量)，有 __shared__ float smem[256];

线程 1 运行 atomicAdd(smem, 1);

线程 2 运行 atomicAdd(smem + 1, 1);

这些线程适用于不同的银行，但通常共享内存。他们是并行运行还是要排队？这个例子有什么区别，如果Thread1和Thread2来自分开的warp还是一般的warp？

Answer 1

我数了大约 10 个问题。这让人很难回答。建议您每个问题问一个问题。

一般来说，warp 中的所有线程都在执行相同的指令流。所以有两种情况我们可以考虑:

没有条件(例如 if...then...else)在这种情况下，所有线程都在执行相同的指令，这恰好是一条原子指令。然后所有 32 个线程都将执行一个原子，尽管不一定在同一位置。所有这些原子都将由 SM 处理，并且在某种程度上会序列化(如果它们更新相同的位置，它们将完全序列化)。

带条件 例如，假设我们有 if (!threadIdx.x) AtomicAdd(*data, 1);然后线程 0 将执行原子，并且
其他人不会。看起来我们可以让其他人做
别的东西，但锁步扭曲执行不允许这样做。
Warp 执行被序列化，以便所有线程都使用 if (true) path 将一起执行，并且所有执行该路径的线程if (false) path会一起执行，但是true和false
路径将被序列化。再说一次，我们真的不能有不同的
warp 中的线程同时执行不同的指令。

它的网络是，在一个经线内，我们不能让一个线程执行原子操作，而其他线程同时执行其他操作。

您的许多其他问题似乎期望内存事务在它们起源的指令周期结束时完成。事实并非如此。对于全局内存和共享内存，我们必须在代码中采取特殊步骤以确保之前的写入事务对其他线程可见(这可以作为事务完成的证据。)一种典型的方法是使用屏障说明，例如 __syncthreads()或 __threadfence()但是如果没有这些屏障指令，线程就不会“等待”写入完成。一次(依赖于 a 的操作)读取可以停止线程。写入通常不能停止线程。

现在让我们看看你的问题:

so if one of them start an atomic operation, what other will do? Wait?

不，他们不等。原子操作被分派(dispatch)到处理原子的 SM 上的一个功能单元，并且所有线程一起以锁步继续。由于原子通常意味着读取，是的，读取可以阻止扭曲。但是线程不会等到原子操作完成(即写入)。但是，对该位置的后续读取可能会再次阻止扭曲，等待原子(写入)完成。在保证更新全局内存的全局原子的情况下，它将使原始 SM 中的 L1(如果启用)和 L2 无效，如果它们包含该位置作为条目。

Is there any way to take other threads for some work, like reads some memory, so the atomic operation will hide it's latency?

不是真的，因为我在开头说的原因。

Atomic is really slow, does it? How can I accelerate it? Is there any pattern to work with it?

是的，如果原子支配事件(例如朴素的减少或朴素的直方图)，原子可以使程序运行得更慢。一般来说，加速原子操作的方法是不使用它们，或者谨慎地使用它们，以这样的方式不支配程序事件。例如，朴素的归约将使用原子将每个元素添加到全局总和中。对于线程块中完成的工作，智能并行缩减将根本不使用原子。在线程块减少结束时，可以使用单个原子将线程块部分总和更新为全局总和。这意味着我可以对任意数量的元素进行快速并行减少，大约 32 个原子添加或更少。这种对原子的 Thrift 使用在整个程序执行中基本上不会被注意到，除了它可以在单个内核调用而不是 2 个内核调用中完成并行减少。

Shared memory: Does they run parralel or they will be queued?

他们会排队。这样做的原因是，可以在共享内存上处理原子操作的功能单元数量有限，不足以在单个周期内为来自 warp 的所有请求提供服务。

我避免尝试回答与原子操作吞吐量相关的问题，因为文档 AFAIK 中没有很好地指定这些数据。可能是，如果您同时发出足够多的或几乎同时发生的原子操作，由于提供原子功能单元的队列已满，某些扭曲将在原子指令上停滞。我不知道这是真的，我无法回答有关它的问题。