协作组
用于灵活同步范围的现代 API
协作组 是 CoddyKit 上的免费 CUDA Academy 课时。 这是第 4 节课，共 4 节。 你可以在下方免费阅读本课时的完整内容 — 然后在浏览器中使用内置代码编辑器和全天候 AI 导师进行实践。 这是 CUDA Academy 学习路径的一部分，你的进度在网页和 CoddyKit 应用中同步。 CUDA Academy 课程共包含 4 节课。
本课时的部分内容尚未翻译，以英文显示。
A Cleaner Sync API
Raw masks and intrinsics work, but they are fiddly. Cooperative groups wrap threads into objects you can name, size, and synchronize explicitly.
#include <cooperative_groups.h>
namespace cg = cooperative_groups;Grab the Whole Block
Start by getting a handle to your block of threads. this_thread_block returns a group you can sync just like syncthreads, but as an object.
cg::thread_block block = cg::this_thread_block();Sync Through the Group
Calling sync on the block is the modern barrier. It does exactly what syncthreads does, but reads clearly as a method on the group you mean.
block.sync();Carve Out a Warp Tile
You can split a block into fixed-size tiles. A 32-lane tile gives you a warp-sized group with clean methods instead of raw shuffle masks.
auto warp = cg::tiled_partition<32>(block);Methods Replace Masks
A tiled group offers shfl_down and friends without a mask argument. The group already knows its members, so the API stays short and safe.
val += warp.shfl_down(val, offset);Know Your Position
Every group exposes its members and your spot. thread_rank gives your index inside the group, and size returns how many threads it holds.
int rank = warp.thread_rank();Smaller Tiles Too
Tiles need not be 32 wide. A tiled_partition of 8 or 16 makes sub-warp groups, handy when your data naturally clusters in small sets.
Group-Level Reductions
The library ships ready-made collectives. A group reduce sums a tile in one call, hiding the offset loop you wrote by hand earlier.
int total = cg::reduce(warp, val, cg::plus<int>());Grids That Sync
The boldest group is the grid group. With a cooperative launch, every block can sync at one barrier, something a normal kernel cannot do.
Cooperative Launch Required
Grid-wide sync only works if you start the kernel with cudaLaunchCooperativeKernel and the GPU supports it. A normal launch will not allow it.
Why Bother
Cooperative groups make warp code readable and portable: no hand-managed masks, clear scopes, and reusable collectives that match what you mean.
Quick Check
Recall how you create a warp-sized cooperative group from a block.
Recap
Cooperative groups turn masks into named objects: tile a block, call reduce, even sync a whole grid. You now own warp-level CUDA. ✨
常见问题解答
「协作组」课时是免费的吗？
是的 — 「协作组」的完整文本可在网页上免费阅读。要进行交互式练习（内置代码编辑器和全天候 AI 导师）并解锁 CUDA Academy 课程的其余内容，请升级到 CoddyKit PRO。 CUDA Academy 课程共包含 4 节课。
「协作组」这节课中我会学到什么？
用于灵活同步范围的现代 API 你通过在浏览器中直接运行的动手代码来练习 CUDA Academy，全天候 AI 导师会在你学习这节课的过程中回答你的问题。
学习 CUDA Academy 需要有经验吗？
无需任何先前经验。CoddyKit 上的 CUDA Academy 课程适合初学者到高级学习者，你可以从这里开始或从头开始，按照自己的节奏学习。 这是第 4 节课，共 4 节。
「协作组」课时需要多长时间？
大多数 CoddyKit 课程大约需要 5–10 分钟。每节课都很精短且互动，所以你能稳步进步，并在网页和应用中从离开的地方继续。
我能在这节 CUDA Academy 课中编写并运行代码吗？
能。每节 CUDA Academy 课都包含内置代码编辑器，你可以在浏览器中直接编写并运行真实代码，并获得即时 AI 反馈 — 无需本地设置。