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 进入共享内存。

（待补充：更多说明）