GPU Cooperative Groups

GPU Cooperative Groups is the CUDA programming model extension that provides flexible, hierarchical thread grouping and synchronization primitives beyond the fixed thread-block model — enabling grid-level synchronization, dynamic sub-warp partitioning, and multi-GPU cooperative launches that allow algorithm designers to express synchronization patterns matching their computation's natural structure rather than being forced into the rigid block/grid hierarchy.

Why Cooperative Groups Exist

Classic CUDA provides two synchronization scopes: __syncthreads() within a thread block, and kernel launch boundaries for grid-level synchronization. This forces algorithms requiring global synchronization to split into multiple kernel launches (expensive: 5-20 μs overhead each) or use unreliable atomic-based ad-hoc synchronization. Cooperative Groups fills the gap.

Group Hierarchy

• Thread (1 thread): The fundamental unit. Useful as a parameter to templated algorithms that accept any group type.
• Coalesced Group: Dynamically-formed group of converged threads within a warp. Created by tiled_partition or coalesced_threads() — only threads that are actually active participate. Enables efficient sub-warp algorithms.
• Thread Block: Equivalent to the traditional block — all threads launched in the same block. sync() is equivalent to __syncthreads().
• Thread Block Cluster (Hopper+): A group of up to 16 thread blocks guaranteed to execute concurrently on the same GPC (Graphics Processing Cluster). Enables direct shared-memory access across blocks via distributed shared memory.
• Grid Group: ALL thread blocks in the grid. grid.sync() provides a true global barrier — all blocks synchronize before proceeding. Requires cooperative launch (cudaLaunchCooperativeKernel) which guarantees all blocks execute concurrently.
• Multi-Grid Group: Synchronization across multiple GPUs in a multi-GPU cooperative launch. Enables single-kernel multi-GPU algorithms without CPU-side synchronization.

Tiled Partition

Split a group into fixed-size tiles for warp-level algorithms:

auto warp = cooperative_groups::tiled_partition<32>(this_thread_block());
auto half_warp = cooperative_groups::tiled_partition<16>(warp);
int sum = half_warp.shfl_down(val, 8) + val;  // 16-thread reduction

This enables portable warp-level algorithms that work with any tile size (1, 2, 4, 8, 16, 32).

Use Cases

• Persistent Kernels: A single kernel that runs for the lifetime of the application, processing work items from a global queue. Grid-level sync separates phases. Avoids repeated kernel launch overhead.
• Graph Algorithms: BFS/SSSP iterations require global synchronization between levels. Cooperative grid sync enables single-kernel BFS — 5-10× faster than multi-kernel approaches on small graphs.
• Iterative Solvers: Conjugate gradient and Jacobi iterations require a global reduction (dot product) between iterations. Grid sync enables single-kernel iterative solvers.

Cooperative Groups is the synchronization abstraction that unlocks algorithm patterns impossible in the classic CUDA model — providing the flexibility to synchronize at any granularity from sub-warp to multi-GPU, enabling persistent kernels and global-barrier algorithms that were previously impractical.