100 DAYS OF GPU Course preview
← Glossary
CUDA glossaryExecution model
CC 7.5

What are cooperative groups in CUDA?

A C++ API for naming the group you want to synchronize, from a 32-lane tile up to the whole grid.

Put grid.sync() in a kernel, launch it with <<<blocks, 256>>>, and everything works except the barrier. The header compiles, the link succeeds, the launch returns cudaSuccess, and each block walks past the call on its own. Somebody on the NVIDIA forums hit exactly that and reported it as a correctness bug: "it runs only on the first block and ignores the other". The API is the part that makes such a call mean something, and the meaning gets arranged at the launch rather than inside the kernel.

Four types cover almost all use of it. cg::this_thread_block() hands back the block you already have, cg::tiled_partition<32>(block) cuts that block into warp-shaped tiles, cg::coalesced_threads() names whichever lanes reached this branch, and cg::this_grid() names every thread in the launch. All of them carry sync(), thread_rank() and num_threads(), which is most of the appeal: __syncthreads() and the shuffle intrinsics stop being loose functions and become methods on a group whose size a function signature can state.

Three of those four barriers are free, and the reason is residency. A block goes to one SM whole or it does not start, so by the time your kernel body runs, waiting is all a block barrier has to do. A grid is not resident that way. Ask for 16,387 blocks on a card that holds a few hundred at once and most of them sit queued behind the blocks standing at your barrier, which will never finish, so nothing can release them. A cooperative launch is the promise that every block is resident before the kernel starts, and the promise costs you the grid: its size comes from cudaOccupancyMaxActiveBlocksPerMultiprocessor for that kernel times the SM count, not from your data.

So the grid stops describing the problem and becomes a resource you fit inside, with a grid-stride loop covering the rest of the elements. Two habits follow. Size the grid from the occupancy query rather than from n, and read that ceiling as a property of the kernel rather than the card: add a register array and it drops, which is day 17's subject. Then guard with grid.is_valid() instead of assuming, because it returns 0 under an ordinary launch and calling sync() on an invalid group is undefined.

Measured

On a Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -O3 -arch=sm_75, day 28 sums 4,194,915 floats twice: once as two ordinary launches, once as one launch with a grid barrier in the middle. cudaDevAttrCooperativeLaunch reads 1 on this card.

Question asked of the runtime Answer
blocks a cooperative launch fits, reduceCooperativeGrid 4 per SM x 40 SMs = 160
blocks an ordinary launch sized to n would use 16387
cg::this_grid().is_valid() under <<<160, 256>>> 0
the same kernel and grid, launched cooperatively 1
a cooperative launch of 161 blocks cudaErrorCooperativeLaunchTooLarge

Two orders of magnitude between the grid the data wants and the grid the barrier allows, on one kernel and one card. The refusal is the good part. 161 blocks fails at the launch with a name on it, before anything runs, where the alternative would be a barrier waiting forever on blocks that were never scheduled.

Both routes returned 3670548.750000, matching a CPU double reference, and the timing says the barrier is not a speedup: two launches took 0.0819 ms and the single cooperative launch 0.0830 ms, mean of 10 runs after 3 warm-ups with copies excluded. One launch instead of two, at a hundredth of the grid. That trade is the thing to weigh, and day 28 measures it rather than assuming it.

Diagram

Original SVG: three barrier scopes stacked as bands over the same grid. Band one brackets 32 lane boxes inside one block, band two brackets all eight rows of that block, band three has to reach every block outline on the row, with the blocks past the resident set drawn greyed and outside the bracket.

Alt text: "Three barrier scopes over one launch. A tile barrier covers 32 lanes and a block barrier covers 256 threads, both free because those threads are already running. A grid barrier has to cover every block, so the launch itself has to promise they are all resident."

Code

From code/day28-cooperative-groups/cooperative_reduce.cu. cudaLaunchKernelEx takes a config struct where the triple chevron takes four values, and the attribute array is the part <<<>>> cannot express:

    cudaLaunchAttribute coopAttr[1] = {};
    coopAttr[0].id = cudaLaunchAttributeCooperative;
    coopAttr[0].val.cooperative = 1;

    cudaLaunchConfig_t config = {};
    config.gridDim = dim3(static_cast<unsigned int>(blocks));
    config.blockDim = dim3(static_cast<unsigned int>(kThreadsPerBlock));
    config.dynamicSmemBytes = 0;
    config.stream = nullptr;
    config.attrs = coopAttr;
    config.numAttrs = 1;

Grid, block, shared bytes and stream are what you already write. The one non-zero attribute is what turns a launch into a promise about residency.

Related terms

Where you meet this

Sources

Byline

Written by: pending. Reviewed by: pending. Written on: pending. Last checked on: pending. This entry stays a draft until a named author and a different named reviewer sign it.