CUDA Academy · 课时

投票与表决函数

在一个线程束中轮询谓词

第 3 / 4 课13 个步骤

投票与表决函数 是 CoddyKit 上的免费 CUDA Academy 课时。 这是第 3 节课,共 4 节。 你可以在下方免费阅读本课时的完整内容 — 然后在浏览器中使用内置代码编辑器和全天候 AI 导师进行实践。 这是 CUDA Academy 学习路径的一部分,你的进度在网页和 CoddyKit 应用中同步。 CUDA Academy 课程共包含 4 节课。

Lane 可以投票

除了移动数据之外,一个 warp 还可以共同回答是或否的问题。Vote 内置函数让每个 lane 用一个低开销步骤共享一个真或假的谓词结果。

认识 ballot_sync

最灵活的投票方式是 ballot_sync。它返回一个 32 位数字:当第 i 个 lane 的谓词为真时,第 i 位恰好为 1。

unsigned bits = __ballot_sync(mask, pred);

每个 lane 一个位

返回的整数就是这个 warp 的微型位图。读取第 i 位即可知道第 i 个 lane 是否投了赞成票,所有结果都打包在一个寄存器中。

统计真值

想知道有多少个 lane 回答“是”吗?请对 ballot 结果进行位计数。popc 会统计其中置位的位数,立即得到 warp 中真值投票的数量。

int yes = __popc(__ballot_sync(mask, pred));

全部或任意

两个快捷函数可以回答常见问题。只有当每个 lane 都同意时,all_sync 才返回 true;只要至少一个 lane 同意,any_sync 就会返回 true。

bool every = __all_sync(mask, pred);

任意表示至少一个

对于提前退出,请使用 any_sync:如果任意 lane 找到了匹配项,整个 warp 就可以共同响应,而不必逐个扫描 lane。

bool found = __any_sync(mask, pred);

流压缩

Ballot 可以实现压缩:每个 lane 决定是否保留自己的元素并进行投票,位图则告诉各个 lane 应将数据写入紧凑输出数组中的哪个位置。

找到我的位置

为了放置自己的值,一个 lane 会统计自身位置以下的真值位数量。这个由 ballot 得到的前缀计数就是它在压缩结果中的索引。

int slot = __popc(bits & ((1u << lane) - 1));

选择一个首领 lane

有时应由一个 lane 执行共享工作。使用 ffs 找到 ballot 中最低的置位位,它就是 warp 可以共同认可的自然首领。

始终传入掩码

与 shuffle 一样,vote 内置函数也需要活动 掩码。非活动 lane 会贡献一个 0 位,因此请传入正确的集合,避免统计不存在的 lane。

低开销协调

投票只需要一条指令,也不会访问内存。与原子操作或共享内存标志相比,它能以低得多的成本让 warp协调决策。

快速检查

回想一下,ballot_sync 会返回什么给每个 lane。

回顾

您使用 ballot 和 vote 对 warp 进行查询:统计真值、检查全部或任意条件,并以低成本压缩数据流。接下来学习用于灵活同步的协作组。✨

免费开始

用 AI 导师学习 C++ — 免费

在浏览器中编写并运行真实代码,获得全天候 AI 导师的即时帮助,并在网页或应用中继续学习。

课程
30
课程
120

常见问题解答

「投票与表决函数」课时是免费的吗?

是的 — 「投票与表决函数」的完整文本可在网页上免费阅读。要进行交互式练习(内置代码编辑器和全天候 AI 导师)并解锁 CUDA Academy 课程的其余内容,请升级到 CoddyKit PRO。 CUDA Academy 课程共包含 4 节课。

「投票与表决函数」这节课中我会学到什么?

在一个线程束中轮询谓词 你通过在浏览器中直接运行的动手代码来练习 CUDA Academy,全天候 AI 导师会在你学习这节课的过程中回答你的问题。

学习 CUDA Academy 需要有经验吗?

无需任何先前经验。CoddyKit 上的 CUDA Academy 课程适合初学者到高级学习者,你可以从这里开始或从头开始,按照自己的节奏学习。 这是第 3 节课,共 4 节。

「投票与表决函数」课时需要多长时间?

大多数 CoddyKit 课程大约需要 5–10 分钟。每节课都很精短且互动,所以你能稳步进步,并在网页和应用中从离开的地方继续。

我能在这节 CUDA Academy 课中编写并运行代码吗?

能。每节 CUDA Academy 课都包含内置代码编辑器,你可以在浏览器中直接编写并运行真实代码,并获得即时 AI 反馈 — 无需本地设置。

此课程中的所有课时

  1. 线程束、通道与掩码
  2. 使用 __shfl_down_sync 进行归约
  3. 投票与表决函数
  4. 协作组
← 返回 CUDA Academy