CUDA Cooperative Groups: Thread Block Tiles, Grid Sync and Cooperative Launch

The group model

A group is a handle with three core operations: num_threads(), thread_rank() (this thread's index within the group) and sync(). Older code uses size(), which is equivalent to num_threads(). Because a group is an ordinary C++ object, you can pass it to a device function. The function then declares in its signature which threads have to call it together, where before that was a comment.

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

// Any group type works: the caller decides who participates.
template <typename Group>
__device__ float group_sum(Group g, float v) {
    return cg::reduce(g, v, cg::plus<float>());   // hardware-accelerated on cc 8.0+ where available
}

__global__ void demo(const float* x, float* out, int n) {
    cg::thread_block block = cg::this_thread_block();
    cg::thread_block_tile<32> warp = cg::tiled_partition<32>(block);
    int i = block.group_index().x * block.num_threads() + block.thread_rank();
    float v = (i < n) ? x[i] : 0.0f;
    float s = group_sum(warp, v);                 // every lane of the tile must call this
    if (warp.thread_rank() == 0) atomicAdd(out, s);
}

Thread block tiles and warp collectives

A thread_block_tile<N> with a compile-time size N (a power of two up to 32) is the type you will use most. Its size is known at compile time, so the compiler can generate shuffles and votes directly. It provides warp collectives as methods: shfl, shfl_down, shfl_up, shfl_xor, any, all, ballot and match_any. Each method is scoped to the tile, with no lane masks to compute by hand. meta_group_rank() tells a thread which tile of the parent it belongs to, which is useful for indexing per-tile output.

Here is the reduction written by hand, to show what cg::reduce replaces. In practice, call the library version.

template <unsigned N>
__device__ float tile_sum(cg::thread_block_tile<N> t, float v) {
    for (unsigned off = N / 2; off > 0; off /= 2)
        v += t.shfl_down(v, off);      // stays inside the tile, mask derived from N
    return v;                          // valid in t.thread_rank() == 0
}

// Block reduction: tiles of 32 reduce in registers, one slot per tile in shared memory.
__device__ float block_sum(cg::thread_block b, float v) {
    __shared__ float partial[32];                  // up to 1024 threads / 32
    auto w = cg::tiled_partition<32>(b);
    v = tile_sum(w, v);
    if (w.thread_rank() == 0) partial[w.meta_group_rank()] = v;
    b.sync();
    if (w.meta_group_rank() == 0) {
        v = (w.thread_rank() < w.meta_group_size()) ? partial[w.thread_rank()] : 0.0f;
        v = tile_sum(w, v);
    }
    return v;                                      // valid in thread 0 of the block
}

Smaller tiles help when the work per item is small. In a sparse matrix-vector product with short rows, giving each row a tile<8> instead of a whole warp keeps lanes busy without wasting 24 of 32 threads on a six-element row. Because the tile size is a template argument, you can choose it with a heuristic on the host and instantiate 4, 8, 16 and 32 variants.

Data-dependent groups: coalesced and labeled partitions

Sometimes the group you need is defined by the data, not the layout. cg::coalesced_threads() returns the threads that are active at that point, for example the ones inside an if. The guide states that it makes no guarantee about which threads are returned. Use it for work where any set of participants is acceptable. The classic example is warp-aggregated atomics. Instead of every thread that passes a filter doing its own atomicAdd on a global counter, the active threads elect a leader, which does one atomic for everyone and broadcasts the base offset.

__device__ int append_index(int* counter) {
    cg::coalesced_group g = cg::coalesced_threads();   // only threads that got here
    int base = 0;
    if (g.thread_rank() == 0) base = atomicAdd(counter, g.num_threads());
    base = g.shfl(base, 0);                             // broadcast leader's result
    return base + g.thread_rank();                      // unique slot per thread
}

__global__ void compact(const float* in, float* out, int* count, int n, float thr) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n && in[i] > thr) out[append_index(count)] = in[i];
}

In a filter that keeps 30% of elements, this reduces atomics on the counter by roughly the number of survivors per warp, about ten times, and it removes the contention that serialises them. cg::labeled_partition(g, key) generalises the idea. It splits a group into subgroups of threads that share a key, so threads updating the same histogram bin or hash bucket can combine their updates before touching memory. binary_partition is the cheaper special case for a boolean key.

Grid sync and cooperative launch

A grid_group from cg::this_grid() offers sync() across every thread of the kernel. This works only if every block is resident on the GPU at the same time. A block waiting at the barrier cannot finish, so a block that has not yet been scheduled would wait forever. CUDA therefore requires a cooperative launch, which checks co-residency up front and fails instead of deadlocking. Grid sync needs compute capability 6.0 or newer. The device attribute cudaDevAttrCooperativeLaunch tells you whether it is available. Older toolkit documentation required relocatable device code (-rdc=true) for grid sync. Check the guide for your toolkit version and add the flag if linking fails without it.

// Host side: size the grid to what can be resident, then launch cooperatively.
int dev = 0, coop = 0, sms = 0, perSM = 0;
cudaDeviceGetAttribute(&coop, cudaDevAttrCooperativeLaunch, dev);
cudaDeviceGetAttribute(&sms, cudaDevAttrMultiProcessorCount, dev);
const int threads = 256;
const size_t smem = 0;                          // this kernel uses no shared memory
cudaOccupancyMaxActiveBlocksPerMultiprocessor(&perSM, jacobi_kernel, threads, smem);
if (!coop || perSM == 0) { /* fall back to one launch per iteration */ }

dim3 grid(sms * perSM), block(threads);         // never more than can co-reside
void* args[] = { &d_a, &d_b, &d_rhs, &d_err, &n, &iters };
cudaError_t e = cudaLaunchCooperativeKernel((void*)jacobi_kernel, grid, block, args, smem, stream);
// cudaErrorCooperativeLaunchTooLarge means the grid could not be made resident.

The kernel then becomes a persistent loop. Each thread covers several elements with a grid-stride loop, because the grid is now sized by occupancy rather than by the problem size.

__global__ void jacobi_kernel(float* a, float* b, const float* rhs, float* err, int n, int iters) {
    cg::grid_group grid = cg::this_grid();
    for (int it = 0; it < iters; ++it) {
        for (int i = grid.thread_rank(); i < n; i += grid.num_threads())
            if (i > 0 && i < n - 1) b[i] = 0.5f * (a[i - 1] + a[i + 1] - rhs[i]);
        grid.sync();                                   // all of b written before anyone reads it
        float* t = a; a = b; b = t;                    // every thread swaps identically
    }
}

Worked example: a Jacobi solver in one kernel

Take a 1D Jacobi solve with n = 2²⁰ points and 2,000 iterations. The naive version launches one kernel per iteration. Each launch does about 1 million cheap updates, a few microseconds of real work, plus launch overhead and the drain and refill of the GPU between kernels. At that size, overhead can be a large share of the time per iteration.

The cooperative version launches once. The grid is SM count times resident blocks per SM, for example 132 × 8 = 1,056 blocks of 256 threads. Each thread handles about four points per iteration, and grid.sync() replaces the kernel boundary. You save the per-launch overhead and the inter-kernel bubble. You keep the cost of a grid-wide barrier, which is far from free: it is implemented through global memory and must wait for the slowest block.

Measure before you decide. Time both versions with CUDA events over the same iteration count, and profile with Nsight Systems to see the gaps between launches. Then compare with a CUDA Graph that captures the 2,000 launches. Graphs cut launch overhead without requiring co-residency. As a rule of thumb, grid sync wins when each phase is short and phases are many, and when data can stay in registers or shared memory across a barrier, which is impossible across a kernel boundary. It loses when the problem would want many more blocks than can be resident, because occupancy-limited grids underuse a large GPU on a large problem.

Clusters, and what happened to multi-GPU groups

Hopper (compute capability 9.0) adds a level between the block and the grid: the thread block cluster. cg::this_cluster() returns it. Blocks in a cluster are guaranteed to be co-scheduled on a GPU processing cluster, can synchronise with cluster.sync() or with the split barrier_arrive()/barrier_wait() pair, and can read each other's shared memory (distributed shared memory). When an algorithm needs neighbouring blocks to exchange data, such as a halo or a split-K reduction, a cluster gives you this without a cooperative launch or global-memory round trips. When launched without clusters, this_cluster() behaves as a 1x1x1 cluster. Multi-device cooperative groups, which earlier toolkits offered for synchronising across GPUs, were removed as of CUDA 13. Use NCCL or explicit inter-GPU signalling instead.

Failure modes

Most of these come down to one rule. Every thread in a group must reach each collective call on that group. CG makes the group explicit in the code, but the compiler does not check that every thread actually reaches the call.

What to do next

1. Replace hand-written __shfl_down_sync(0xffffffff, ...) loops with thread_block_tile<32> and cg::reduce, and delete the masks.
2. Change device helpers that must be called by a whole warp or block to take the group as a parameter.
3. Use warp-aggregated atomics in any filter, compaction or histogram that hammers a shared counter.
4. Before using grid sync, read GPU occupancy and size the grid from cudaOccupancyMaxActiveBlocksPerMultiprocessor.
5. Benchmark grid sync against relaunching and against a CUDA Graph on your real problem size, then keep the fastest.
6. Review warp scheduling, shared memory and async memcpy for the other half of the picture: what threads wait on.