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

Question

warp中的线程在物理上是并行的,所以如果其中一个(称为线程X)启动一个原子操作,那么其他什么呢？等待？这是否意味着,当线程X被推送到原子队列时,所有线程都将等待,获取访问权限(互斥量)并使用该互斥锁保护内存,然后重新安排互斥锁？

有没有办法把其他线程用于某些工作,比如读取一些内存,所以原子操作会隐藏它的延迟？我的意思是,15个闲置的线程......不太好,我想.Atomic真的很慢,不是吗？我怎样才能加速它呢？是否有任何模式可以使用它？

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

• Thread1运行 atomicAdd(smem, 1);
• Thread2运行 atomicAdd(smem + 1, 1);

这些线程适用于不同的银行,但通常是共享内存.他们是运行parralel还是排队？如果Thread1和Thread2来自分离的warp或一般的warp,那么这个例子有什么不同吗？

Answer 1

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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