在CUDA Developer Blog上阅读了这篇文章后，我努力了解何时可以安全\正确使用__activemask()代替__ballot_sync()。
在活动掩码查询部分中，作者写道：
这是不正确的，因为它将导致部分和而不是总计。
之后，在“ 机会扭曲级编程”部分中，他们使用此函数__activemask()是因为：
如果要在库函数中使用扭曲级编程，但是不能更改函数接口，则可能会很困难。
Rob*_*lla 11
__active_mask()CUDA中没有。这是一个错字（在博客文章中）。应该是__activemask()。
__activemask()是唯一一个查询。它询问问题“在这个周期中，warp中的哪些线程当前正在执行该指令？” 这等效于询问“此时扭曲中的哪些线程当前已收敛？”
它对收敛没有影响。这不会导致线程收敛。它没有扭曲同步行为。
__ballot_sync()另一方面具有收敛行为（根据提供的mask）。
应该根据Volta变形执行模型来考虑此处的主要区别。由于warp执行引擎中的硬件发生了变化，因此Volta及其以后版本可以支持warp中的线程在比以前的体系结构更多的场景中使用更长的时间，并且可以在更长的时间内使用。
我们这里所说的背离是由于先前的条件执行而引起的偶然背离。在Volta之前或之后，由于显式编码而引起的强制分歧是相同的。
让我们考虑一个例子：
if (threadIdx.x < 1){
statement_A();}
statement_B();
假设线程块X尺寸大于1，statement_A()则处于强制发散区域。执行时，经纱将处于发散状态statement_A()。
那statement_B()呢 CUDA执行模型在执行时不对经纱是否处于发散状态做出任何特定声明statement_B()。在Volta之前的执行环境中，程序员通常会期望在上一条if语句的结尾大括号处存在某种形式的翘曲重新收敛（尽管CUDA对此不作任何保证）。因此，一般的期望是statement_B()将在非分歧状态下执行。
但是，在Volta执行模型中，CUDA不仅不提供任何保证，而且在实践中，我们可能会观察到翘曲处于发散状态statement_B()。 发散在statement_B()不需要代码的正确性（而它需要在statement_A()），也不是收敛在statement_B()所要求的CUDA执行模型。如果statement_B()在Volta执行模型中可能会发生分歧，我将其称为偶然分歧。产生这种差异并不是由于代码的某些要求，而是由于某种先前的条件执行行为而引起的。
如果我们在处没有分歧statement_B()，那么这两个表达式（如果它们在处statement_B()）应该返回相同的结果：
int mask = __activemask();
和
int mask = __ballot_sync(0xFFFFFFFF, 1);
因此，在伏特伏特情况下，当我们通常不希望statement_B()在实际中出现分歧时，这两个表达式将返回相同的值。
在Volta执行模型中，我们可以在处产生偶发散度 statement_B()。因此，这两个表达式可能不会返回相同的结果。为什么？
该__ballot_sync()指令与所有其他具有mask参数的CUDA 9+扭曲级别内部函数一样，具有同步效果。如果我们存在代码强制的分歧，则如果mask参数指示的同步“请求”无法实现（就像上面我们要求完全收敛的情况），那将代表非法代码。
但是，如果我们有偶然的分歧（仅对于本例而言），则__ballot_sync()语义至少应首先重新收敛扭曲，直到mask参数要求的程度为止，然后执行请求的投票操作。
该__activemask()操作没有这种收敛行为。它仅报告当前聚合的线程。如果某些线程因任何原因而发散，它们将不会在返回值中报告。
如果然后创建执行某些扭曲级别操作（例如博客文章中建议的扭曲级别总和减少）的代码，并基于__activemask()vs 选择了要参与的线程__ballot_sync(0xFFFFFFFF, 1)，那么可以想象得到一个不同的结果的偶然分歧。__activemask()在存在偶然差异的情况下，该实现将计算不包括所有线程的结果（即它将计算“部分”和）。另一方面，__ballot_sync(0xFFFFFFFF, 1)由于此实现将首先消除偶然的差异，因此将强制所有线程参与（计算“总”和）。
在博客文章的清单10中给出了与我在此处给出的类似示例和说明。
__activemask在“机会主义扭曲级编程”的博客文章中提供了一个可能正确使用的示例，这里：
int mask = __match_all_sync(__activemask(), ptr, &pred);
该语句的意思是“告诉我哪些线程已收敛”（即__activemask()请求），然后“至少使用这些线程来执行该__match_all操作。这是完全合法的，并且将使用此时恰好会聚的所有线程。继续清单9的示例，mask在上面的步骤中计算出的值仅用于其他扭曲合作原语：
res = __shfl_sync(mask, res, leader);
（恰好在一段条件代码之后）。这将确定哪些线程可用，然后强制使用那些线程，而不管可能存在什么偶然差异，以产生可预测的结果。