Cuda _sync 函数，如何处理未知的线程掩码？

Posted 2023-02-16

【中文标题】Cuda _sync 函数，如何处理未知的线程掩码？

【英文标题】：Cuda _sync functions, how to handle unknown thread mask?

【发布时间】：2020-11-10 19:35:16

【问题描述】：

这个问题是关于适应从锁定步骤到独立程序计数器的语义变化。从本质上讲，我可以将 int __all(int predicate); 之类的调用更改为 volta。

例如，int __all_sync(unsigned mask, int predicate);

有语义：

为掩码中所有未退出的线程计算谓词，当且仅当谓词对所有线程的计算结果都为非零时才返回非零。

文档假定调用者知道哪些线程处于活动状态，因此可以准确地填充掩码。

必须传递一个掩码，指定参与调用的线程

我不知道哪些线程处于活动状态。这是一个内联到用户代码不同位置的函数。这使得以下其中一项具有吸引力：

__all_sync(UINT32_MAX, predicate);
__all_sync(__activemask(), predicate);

第一个类似于https://forums.developer.nvidia.com/t/what-does-mask-mean-in-warp-shuffle-functions-shfl-sync/67697 宣布非法的案例，引用自那里：

例如，这是非法的（将导致 warp 0 的未定义行为）：

if (threadIdx.x > 3) __shfl_down_sync(0xFFFFFFFF, v, offset, 8);

第二个选择，这次引用__activemask() vs __ballot_sync()

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

操作语义似乎是：

有 N 个线程的经线 编译时控制流启用了 M (M D（M 的 D 个子集）线程作为运行时属性聚合 __activemask 返回恰好收敛的线程

这建议同步线程然后使用activemask，

__syncwarp();
__all_sync(__activemask(), predicate);

一篇 nvidia 博客文章说这也是未定义的，https://developer.nvidia.com/blog/using-cuda-warp-level-primitives/

在 __ballot() 之前的第 10 行调用新的 __syncwarp() 原语，如清单 11 所示，也不能解决问题。这又是隐式扭曲同步编程。它假设同一个warp中曾经同步过的线程将保持同步，直到下一个线程发散分支。虽然这通常是正确的，但在 CUDA 编程模型中并不能保证。

这标志着我的想法结束了。同一篇博客最后提供了一些关于选择掩码值的指导：

不要只使用 FULL_MASK（即 32 个线程的 0xffffffff）作为掩码值。如果按照程序逻辑，warp 中的所有线程都不能到达原语，那么使用 FULL_MASK 可能会导致程序挂起。

不要只使用 __activemask() 作为掩码值。 __activemask() 告诉您在调用函数时哪些线程碰巧会聚，这可能与您在集体操作中想要的不同。

请分析计划逻辑并了解会员要求。根据您的程序逻辑提前计算掩码。

但是，我无法计算掩码应该是什么。这取决于包含 __all_sync 的代码被内联到的调用站点的控制流，我不知道。我不想将每个函数都更改为采用无符号掩码参数。

如何在没有全局转换的情况下检索语义正确的行为？

【参考方案1】：

TL;DR：总之，正确的编程方法很可能是做你说你不想做的事情。

更长：

这个blog 特别建议了一种处理未知线程掩码的机会主义方法：在所需操作之前加上__activemask()，并将其用于所需操作。也就是说（从博客中逐字摘录）：

int mask = __match_any_sync(__activemask(), (unsigned long long)ptr);

这应该是完全合法的。

您可能会问“博客末尾提到的第 2 项呢？”我认为，如果您仔细阅读并考虑到我刚刚摘录的以前的用法，这表明“不要只使用__activemask()”如果您打算使用不同的东西。从那里的全文来看，这种解读似乎很明显。这并没有废除先前构造的合法性。

您可能会问“一路上偶然或强制的分歧怎么办？” （即在处理从 elsewhwere 调用的我的函数期间）

我认为你只有两个选择：

在函数入口处获取__activemask() 的值。稍后在您调用所需的同步操作时使用它。这是您对调用环境意图的最佳猜测。 CUDA 不保证这将是正确，但如果您在同步函数调用时没有强制发散，这肯定是合法。 p>

明确调用环境的意图 - 为您的函数添加一个掩码参数并在各处重写代码（您已声明不想这样做）。

如果您在进入函数之前允许扭曲发散的可能性，则无法从函数中推断出调用环境的意图，这会掩盖调用环境的意图。需要明确的是，带有 Volta 执行模型的 CUDA 允许随时出现扭曲发散的可能性。因此，正确的做法是重写代码以明确调用站点的意图，而不是试图从被调用函数中推断出来。

【讨论】：

谢谢。花了几天时间思考这个问题。我认为 volta 可能需要显式 CFG，其中掩码是在应用程序代码的分支上计算的。在 pre-volta 卡上，我可以通过 activemask() 设置相同的掩码，因此现有的入口点可以调用新的入口点。