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

概述warp中的线程在物理上是并行的,所以如果其中一个(称为线程X)启动一个原子 *** 作,那么其他什么呢？等待？这是否意味着,当线程X被推送到原子队列时,所有线程都将等待,获取访问权限(互斥量)并使用该互斥锁保护内存,然后重新安排互斥锁？ 有没有办法把其他线程用于某些工作,比如读取一些内存,所以原子 *** 作会隐藏它的延迟？我的意思是,15个闲置的线程……不太好,我想. Atomic真的很慢,不是吗？我怎样才能 warp中的线程在物理上是并行的,所以如果其中一个(称为线程X)启动一个原子 *** 作,那么其他什么呢？等待？这是否意味着,当线程X被推送到原子队列时,所有线程都将等待,获取访问权限(互斥量)并使用该互斥锁保护内存,然后重新安排互斥锁？

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

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

> Thread1运行atomicAdd(smem,1);
> Thread2运行atomicAdd(smem 1,1);

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

解决方法 我算上10个问题.这让人很难回答.建议你在每个问题上提出一个问题.

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

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

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

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

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

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

不,他们不等.原子 *** 作被分派到处理原子的SM上的一个功能单元,并且所有线程一起以锁步方式继续.由于原子通常意味着读取,是的,读取可以使经线停滞.但是线程不会等到原子 *** 作完成(即写入).但是,随后读取此位置可能会使warp停顿,等待原子(写入)完成.对于保证更新全局内存的全局原子,它将使原始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?

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

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

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

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

总结

以上是内存溢出为你收集整理的c – warp如何与原子 *** 作一起工作？全部内容，希望文章能够帮你解决c – warp如何与原子 *** 作一起工作？所遇到的程序开发问题。

如果觉得内存溢出网站内容还不错，欢迎将内存溢出网站推荐给程序员好友。