Cooperative Groups

A flexible, explicit model for defining and synchronizing groups of threads at any granularity.

Groups: thread_block, thread_block_tile<N> (N = 1..32 power of two), coalesced_group (cg::coalesced_threads()), grid_group (cg::this_grid()), cluster_group (Hopper), plus labeled_partition, binary_partition. Collectives: sync(), shfl, any, all, ballot, reduce, inclusive_scan, exclusive_scan, memcpy_async + wait.

#include <cooperative_groups.h>
#include <cooperative_groups/reduce.h>
namespace cg = cooperative_groups;

__global__ void cgReduce(const float* in, float* out, int n) {
    cg::thread_block block = cg::this_thread_block();
    cg::thread_block_tile<32> warp = cg::tiled_partition<32>(block);

    float v = 0;
    for (int i = block.group_index().x * block.size() + block.thread_rank(); i < n;
         i += block.size() * gridDim.x)
        v += in[i];

    v = cg::reduce(warp, v, cg::plus<float>());     // warp-level reduction
    if (warp.thread_rank() == 0) atomicAdd(out, v);
}

23.1 Grid-wide synchronization

All blocks must be co-resident (grid size ≤ max active blocks); use the occupancy API to size the grid.

__global__ void iterative(float* data, int iters) {
    cg::grid_group grid = cg::this_grid();
    for (int it = 0; it < iters; ++it) {
        step(data);           // phase 1
        grid.sync();          // all blocks in the grid synchronize
    }
}
void* args[] = { &d_data, &iters };
cudaLaunchCooperativeKernel((void*)iterative, gridDim, blockDim, args);