activemask() 对比 ballot_sync()

Question

阅读 CUDA 开发者博客上的 this post 后，我很难理解什么时候 safe\correct 使用 __activemask() 代替 __ballot_sync()。

在 Active Mask Query 部分，作者写道：

This is incorrect, as it would result in partial sums instead of a total sum.

及之后，在 机会性 Warp 级编程 部分中，他们使用函数 __activemask() 因为：

This may be difficult if you want to use warp-level programming inside a library function but you cannot change the function interface.

Answer 1

CUDA中没有__active_mask()。这是一个错字（在博客文章中）。应该是__activemask().

__activemask() 是 only a query。它询问“warp 中的哪些线程当前正在执行此指令，在这个周期中？”这个问题。这相当于询问“此时 warp 中的哪些线程当前会聚？”

对收敛没有影响。它不会导致线程收敛。它没有扭曲同步行为。

另一方面，

__ballot_sync() 具有收敛行为（根据提供的 mask）。

这里的主要区别应该根据 Volta warp 执行模型来考虑。 Volta 及更高版本，由于 warp 执行引擎的硬件变化，可以支持 warp 中的线程在更多场景中发散，并且比以前的架构支持更长的时间。

我们这里所说的背离是由于之前的条件执行导致的附带背离。由于显式编码而导致的强制分歧在 Volta 之前或之后是相同的。

让我们考虑一个例子：

if (threadIdx.x < 1){
   statement_A();}
statement_B();

假设线程块X维度大于1，statement_A()处于enforced发散区域。当执行statement_A()时，扭曲将处于发散状态。

statement_B() 呢？ CUDA 执行模型没有具体说明执行 statement_B() 时 warp 是否处于发散状态。在 Volta 之前的执行环境中，程序员通常会期望在前一个 if 语句的结束花括号处有某种扭曲重新收敛（尽管 CUDA 对此不作任何保证）。因此，一般预期 statement_B() 将在非分歧状态下执行。

然而在 Volta 执行模型中，不仅没有 CUDA 提供的保证，而且在实践中我们可能会观察到 warp 在 statement_B() 处处于发散状态。 Divergence at statement_B() 不是代码正确性所必需的（而在 statement_A() 是必需的），convergence 也不是在 CUDA 执行模型所需的 statement_B() 处。如果 Volta 执行模型中可能出现 statement_B() 处的分歧，我将其称为偶然分歧。分歧不是出于代码的某些要求，而是由于某种先前的条件执行行为。

如果我们在 statement_B() 处没有分歧，那么这两个表达式（如果它们在 statement_B() 处）应该 return 相同的结果：

int mask = __activemask();

和

int mask = __ballot_sync(0xFFFFFFFF, 1);

所以在 pre-volta 的情况下，当我们通常不期望在 statement_B() 处出现分歧时，实际上这两个表达式 return 具有相同的值。

在 Volta 执行模型中，我们可以在 statement_B() 处有偶然分歧。因此这两个表达式可能不会 return 相同的结果。为什么？

__ballot_sync() 指令与所有其他具有掩码参数的 CUDA 9+ 扭曲级内在函数一样，具有同步效果。如果我们有代码强制分歧，如果不能满足掩码参数指示的同步“请求”（就像上面我们请求完全收敛的情况），那将代表非法代码。

然而，如果我们有 偶然的 分歧（仅对于此示例），__ballot_sync() 语义首先 重新收敛扭曲至少到掩码参数请求的范围，然后执行请求的投票操作。

__activemask() 操作没有这种重新收敛行为。它只是报告当前收敛的线程。如果某些线程发散，无论出于 什么原因，它们都不会在 return 值中报告。

如果您随后创建了执行某些 warp 级操作的代码（例如博客文章中建议的 warp 级总和减少）并根据 __activemask() 与 __ballot_sync(0xFFFFFFFF, 1)，在偶然分歧的情况下，你可以想象得到不同的结果。在存在偶然分歧的情况下，__activemask() 实现将计算不包括所有线程的结果（即，它将计算“部分”和）。另一方面，__ballot_sync(0xFFFFFFFF, 1) 实现，因为它会首先消除偶然的分歧，所以会强制所有线程参与（计算“总和”）。

博客文章中的清单 10 附近给出了与我在此处给出的内容类似的示例和描述。

有关“机会性 warp 级编程”的博客文章中给出了可以正确使用 __activemask 的示例，此处：

int mask = __match_all_sync(__activemask(), ptr, &pred);

这个语句是说“告诉我哪些线程被收敛”（即 __activemask() 请求），然后“使用（至少）那些线程来执行 __match_all 操作。这是完全合法，并且将使用恰好在该点收敛的任何线程。随着清单 9 示例的继续，在上述步骤中计算的 mask 用于唯一的其他 warp-cooperative 原语：

res = __shfl_sync(mask, res, leader);

（恰好在一段条件代码之后）。这会确定哪些线程可用，然后强制使用这些线程，而不管可能存在什么偶然的分歧，以产生可预测的结果。

作为对 mask 参数用法的补充说明，请注意 usage statements in the PTX guide。特别是，mask 参数并不是一种排除方法。如果您希望线程被排除在洗牌操作之外，您必须使用条件代码来做到这一点。根据 PTX 指南中的以下声明，这一点很重要：

The behavior of shfl.sync is undefined if the executing thread is not in the membermask.

此外，虽然与上述讨论没有直接关系，但__shfl_sync()的强制发散思想有一个“例外”。 programming guide 声明这在 volta 及更高版本上是可以接受的：

if (tid % warpSize < 16) {
    ...
    float swapped = __shfl_xor_sync(0xffffffff, val, 16);
    ...
} else {
    ...
    float swapped = __shfl_xor_sync(0xffffffff, val, 16);
    ...
}

那里暗示了这样做的原因，我们可以从 PTX guide:

中获得对这种情况下行为的进一步解释

shfl.sync will cause executing thread to wait until all non-exited threads corresponding to membermask have executed shfl.sync with the same qualifiers and same membermask value before resuming execution.

这意味着 if 路径中的 __shfl_sync() 和 else 路径中的 __shfl_sync() 在这种情况下有效地协同工作以产生定义的warp 中所有线程的结果。一些注意事项：

此声明适用于cc7.0及higer

其他构造不一定有效。例如：

  if (tid % warpSize < 16) {
      ...
      float swapped = __shfl_xor_sync(0xffffffff, val, 16);
      ...
  } else {
  }

不会为 warp 中的任何线程提供有趣的结果。

question/answer 可能也有兴趣。

activemask() 对比 ballot_sync()

activemask() vs ballot_sync()

cuda

gpu-warp

__activemask() 对比 __ballot_sync()

__activemask() vs __ballot_sync()

cuda

gpu-warp

activemask() 对比 ballot_sync()

activemask() vs ballot_sync()