Cuda _sync 函数,如何处理未知的线程掩码?

Posted

技术标签:

【中文标题】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() 设置相同的掩码,因此现有的入口点可以调用新的入口点。

以上是关于Cuda _sync 函数,如何处理未知的线程掩码?的主要内容,如果未能解决你的问题,请参考以下文章

CMake + CUDA + 可分离编译->“nvcc 不知道如何处理''”

如何处理未知实体引用?

PyTorch 在加载图像/掩码文件以进行图像分割时如何处理标签?

在 cuda 线程之间共享大量常量数据

Apache,如何处理未知的 php 文件,如标准 URL?

如何处理语义分割中未知类的平均交集(mIOU)?