如何解决Compute Capabilities 7.x 和 8.x 如何协助协作组操作? 减速加速CC 8.0线程异步复制 (CC 8.0)
“合作组”机制出现在最新版本的 CUDA 中。其中一些涉及实际的硬件功能,这些功能不太明显(?)以其他方式使用;但很多基本上只是库代码;并且很难辨别硬件实际上在哪些方面有助于某些特殊功能。
这个问题是关于具有 Compute Capability 7.x 的 GPU。在 CUDA 编程指南中,我注意到以下功能依赖于计算能力:
合作群组功能 | 所需的最低 计算能力 |
---|---|
labeled_partition() |
7.0 |
binary_partition() |
7.0 |
async_memcpy() 实际上是异步的 |
8.0 |
wait() 的某种异步性 |
8.0 |
reduce() 的“加速” |
8.0 |
在约简中使用内在函数:plus、less、greater、bitwise and、bitwise or、bitwise_xor | 8.0 |
具体来说,CC 7.0 和 CC 8.0 引入了哪些硬件特性来启用此功能?它们的确切语义是什么?它们是否都通过 PTX 显式公开,或者其中一些仅在 SASS 中可见?
解决方法
您可以在最新的 PTX ISA reference 中找到作为新引入的 PTX 指令的部分或全部这些新硬件功能。
减速加速(CC 8.0)
现在有 PTX 级别的单操作数,它们可以减少扭曲中一些活动线程的值:
redux.sync.op.type dst,src,membermask;
.op = {.add,.min,.max}
.type = {.u32,.s32}
redux.sync.op.b32 dst,membermask;
.op = {.and,.or,.xor}
所有参与的线程都在 dst
寄存器中获得结果。超过 32 位的溢出被截断。
我想这些有点类似于 vote.sync
操作数。
线程异步复制 (CC 8.0)
这些实际上是几个相关的操作数:
cp.async.{ca,cg}.shared.global
cp.async.commit_group
cp.async.{wait_group,wait_all}
这允许为副本“注册”数据,将注册的数据一起打包到一个组中,然后等待该组中的数据准备就绪。但是 - 它似乎只适用于从全局内存复制到共享内存(甚至不适用于另一个方向)。
阅读有关此机制的更多信息 here。目前,您执行此操作的方式是以全局内存为源的 ld.whatever
指令,然后 st.whatever
进入共享内存。
(待补充:更多说明)
版权声明:本文内容由互联网用户自发贡献,该文观点与技术仅代表作者本人。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如发现本站有涉嫌侵权/违法违规的内容, 请发送邮件至 dio@foxmail.com 举报,一经查实,本站将立刻删除。