Compute Capabilities 7.x 和 8.x 如何协助协作组操作？答案

【问题标题】：How do Compute Capabilities 7.x & 8.x assist cooperative group operations?Compute Capabilities 7.x 和 8.x 如何协助协作组操作？
【发布时间】：2021-05-14 15:49:46
【问题描述】：

“合作组”机制出现在最新版本的 CUDA 中。其中一些涉及实际的硬件功能，这些功能不太明显（？）以其他方式使用；但其中很多基本上只是库代码；并且很难辨别硬件实际上在哪些地方提供了一些特殊功能。

这个问题是关于具有 Compute Capability 7.x 的 GPU。在 CUDA 编程指南中，我注意到以下特性依赖于计算能力：

Cooperative group feature	Required minimum Compute Capability
`labeled_partition()`	7.0
`binary_partition()`	7.0
`async_memcpy()` actually being asynchronous	8.0
Some kind of asynchronicity of `wait()`	8.0
"acceleration" of `reduce()`	8.0
use of intrinsics in reductions for: plus, less, greater, bitwise and, bitwise or, bitwise_xor	8.0

CC 7.0 和 CC 8.0 引入了哪些具体的硬件功能来启用此功能？它们的确切语义是什么？它们都是通过 PTX 显式公开的，还是其中一些仅在 SASS 中可见？

【问题讨论】：

标签： cuda gpu nvidia intrinsics

【解决方案1】：

您可以在最新的PTX ISA reference 中找到部分或全部这些新硬件功能作为新引入的 PTX 指令。

减速加速（CC 8.0）

现在有 PTX 级别的单个操作数，可以减少 warp 中某些活动线程的值：

redux.sync.op.type dst, src, membermask;
.op   = {.add, .min, .max}
.type = {.u32, .s32}

redux.sync.op.b32 dst, src, 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 放入共享内存中。

（待补充：更多说明）

【讨论】：