Cooperative Groups
A modern API for flexible sync scopes.
Cooperative Groups is a free CUDA Academy lesson on CoddyKit — lesson 4 of 4. You can read the complete lesson below for free — then practise it hands-on in the browser with a built-in code editor and a 24/7 AI tutor. It is part of the CUDA Academy learning path, one of 4 lessons in the course, and your progress syncs across the web and the CoddyKit app.
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. ✨
Frequently asked questions
Is the “Cooperative Groups” lesson free?
Yes — the full text of “Cooperative Groups” is free to read here on the web, and the CUDA Academy course includes 4 lessons in total. To practise it interactively (a built-in code editor and a 24/7 AI tutor) and unlock the rest of the CUDA Academy course, upgrade to CoddyKit PRO.
What will I learn in “Cooperative Groups”?
A modern API for flexible sync scopes. You practise CUDA Academy with hands-on code you run directly in the browser, and a 24/7 AI tutor answers your questions as you work through the lesson.
Do I need any experience to start CUDA Academy?
No prior experience is required. CUDA Academy on CoddyKit is structured for beginners through advanced learners; this is — lesson 4 of 4, so you can start here or from the beginning and move at your own pace.
How long does the “Cooperative Groups” lesson take?
Most CoddyKit lessons take about 5–10 minutes. Each one is bite-sized and interactive, so you make steady progress and pick up exactly where you left off across the web and the app.
Can I write and run code in this CUDA Academy lesson?
Yes. Every CUDA Academy lesson includes a built-in code editor, so you write and run real code right in your browser and get instant AI feedback — no local setup required.