5

在阅读了 CUDA 开发人员博客上的这篇文章后，我很难理解什么时候可以安全\正确使用__activemask().__ballot_sync()

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

这是不正确的，因为它会导致部分和而不是总和。

之后，在Opportunistic Warp-level Programming部分中，他们使用该函数__activemask()是因为：

如果您想在库函数中使用 warp 级编程，但您无法更改函数接口，这可能会很困难。

4

1 回答 1

23

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

__activemask()只是一个查询。它询问“在这个循环中，warp 中的哪些线程当前正在执行这条指令？” 这相当于问“经线中的哪些线程目前正在收敛？”

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

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

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

我们这里所说的分歧是由于之前的条件执行造成的偶然分歧。由于显式编码导致的强制发散在 Volta 之前或之后是相同的。

让我们考虑一个例子：

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

假设线程块 X 维度大于 1，statement_A()则处于强制发散区域。执行时，warp 将处于发散状态statement_A()。

怎么样statement_B()？CUDA 执行模型没有特别说明warp 在执行时是否会处于发散状态statement_B()。在 Volta 之前的执行环境中，程序员通常会期望在前一个语句的右花括号处存在某种扭曲再收敛if（尽管 CUDA 不保证这一点）。因此，一般期望是statement_B()在非分歧状态下执行。

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

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

int mask = __activemask();

和

int mask = __ballot_sync(0xFFFFFFFF, 1);

因此，在伏打前的情况下，当我们通常不期望statement_B()在实践中出现分歧时，这两个表达式返回相同的值。

在 Volta 执行模型中，我们可以在 statement_B(). 因此，这两个表达式可能不会返回相同的结果。为什么？

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

但是，如果我们有偶然的分歧（仅对于这个例子），__ballot_sync()语义是首先重新收敛扭曲至少到掩码参数正在请求的程度，然后执行请求的投票操作。

该__activemask()操作没有这种重新收敛行为。它只是报告当前收敛的线程。如果某些线程发生了分歧，无论出于何种原因，它们都不会在返回值中报告。

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

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

__activemask关于“opportunistic warp-level programming”的博客文章中给出了一个正确使用的例子，这里：

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

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

res = __shfl_sync(mask, res, leader);

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

作为对mask参数用法的额外说明，请注意PTX 指南中的用法说明。特别是，该mask参数并非旨在作为排除方法。如果您希望将线程从 shuffle 操作中排除，则必须使用条件代码来执行此操作。鉴于 PTX 指南中的以下声明，这一点很重要：

如果执行线程不在成员掩码中，则 shfl.sync 的行为未定义。

此外，虽然与上述讨论没有直接关系，但对于__shfl_sync(). 编程指南指出，这在 volta 及更高版本上是可以接受的：

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

那里暗示了这样做的原因，我们可以从PTX 指南中进一步解释这种情况下的行为：

shfl.sync 将导致执行线程等待，直到与成员掩码对应的所有未退出线程都已使用相同的限定符和相同的成员掩码值执行 shfl.sync，然后再恢复执行。

这意味着-path 中的__shfl_sync()in和 -path 中的in在这种情况下有效地协同工作，为 warp 中的所有线程生成定义的结果。一些警告：if__shfl_sync()else

  • 本声明适用于cc7.0及更高版本

  • 其他构造不一定会起作用。例如这个：

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

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

这个问题/答案也可能很有趣。

于 2019-01-05T19:39:37.260 回答