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

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

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

它的网络是,在经线内,我们不能让@L_489_0@线程做原子,而其他线程同时做其他事情.

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

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

不,他们不等.原子操作被分派到处理原子的SM上的@L_489_0@功能单元,并且所有线程一起以锁步方式继续.由于原子通常意味着读取,是的,读取可以使经线停滞.但是线程不会等到原子操作完成(即写入).但是,随后读取此位置可能会使warp停顿,等待原子(写入)完成.对于保证更新全局内存的全局原子,它将使原始SM(如果启用)中的L1和L2(如果它们包含该位置作为条目)无效.

不是真的,因为我在开始时说的原因.

是的,原子可以使程序运行得更慢,如果它们支配活动(例如天真减少或天真的直方图).一般来说,加速原子操作的方法是不使用它们,或者谨慎地使用它们,方式是不支配计划活动.例如,天真减少将使用原子将每个元素添加到全局和.对于在threadblock中完成的工作,智能并行缩减将根本不使用原子.在线程块减少结束时,可以使用单个原子将线程块部分和更新为全局和.这意味着我可以快速并行减少任意数量的元素,可能大约32个原子级或更少.除了能够在单个内核调用而不是2中完成并行缩减之外,原始程序执行中基本上不会注意到原子的节约使用.

他们将排队等候.其原因在于,可以在共享内存上处理原子操作的功能单元数量有限,不足以在单个周期内为来自warp的所有请求提供服务.

我避免尝试回答与原子操作吞吐量相关的问题,因为文档AFAIK中没有详细说明这些数据.可能是因为如果你发出足够的同时或几乎同时的原子操作,一些warp会因原子功能单元满的队列而停止在原子指令上.我不知道这是真的,我不能回答有关它的问题.