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

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

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

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

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

这些线程与不同的存储体一起工作，但通常共享内存。它们是并行运行还是会排队？如果 Thread1 和 Thread2 来自单独的线程束或通用线程束，则此示例有什么区别吗？

我数了一下大概有10个问题。这使得回答变得相当困难。建议您每个问题问一个问题。

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

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

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

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

现在我们来看看您的问题：

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

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

有没有办法让其他线程执行某些工作，例如读取一些内存，以便原子操作将隐藏其延迟？

事实并非如此，原因就如我在开头所说的那样。

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

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

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

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

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