warp 如何与原子操作一起工作？

Question

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

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

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

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

这些线程适用于不同的银行，但通常是共享内存。他们是并行运行还是排队？如果 Thread1 和 Thread2 来自单独的经纱或一般经纱，这个例子有什么不同吗？

Answer 1

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

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

1. 没有条件（例如 if...then...else） 在这种情况下，所有线程都在执行相同的指令，这恰好是一条原子指令。然后所有 32 个线程将执行一个原子，尽管不一定在同一位置。所有这些原子都将由 SM 处理，并且在某种程度上会序列化（如果它们更新相同的位置，它们将完全序列化）。
2. with conditionals 例如，假设我们有if (!threadIdx.x) AtomicAdd(*data, 1); 然后线程 0 将执行原子，并且 其他人不会。看起来我们可以让其他人去做 别的东西，但锁步扭曲执行不允许这样做。 Warp 执行是序列化的，这样所有采用if (true) 路径的线程将一起执行，并且所有执行 if (false)路径会一起执行，但是真假 路径将被序列化。所以再一次，我们真的不能有不同的 warp 中的线程同时执行不同的指令。

它的本质是，在一个 warp 中，我们不能让一个线程执行原子操作，而其他线程同时执行其他操作。

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

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

所以如果其中一个启动原子操作，其他的会做什么？等待？

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

有没有办法让其他线程做一些工作，比如读取一些内存，所以原子操作会隐藏它的延迟？

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

Atomic 真的很慢，是吗？我怎样才能加速它？有什么模式可以使用它吗？

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

共享内存：它们是并行运行还是排队？

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

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