CUDA Cooperative Groups is the programming model extension that provides flexible, hierarchical thread synchronization and communication beyond the traditional warp and thread block boundaries — enabling grid-wide synchronization, dynamic sub-warp grouping, and multi-block cooperation that were previously impossible or required awkward workarounds in standard CUDA.

Traditional CUDA Hierarchy Limitations

• Warp (32 threads): Implicit synchronization (SIMT lockstep).
• Thread Block (up to 1024 threads): __syncthreads() — block-level barrier.
• Grid: NO synchronization primitives — blocks are independent.
• Problem: Algorithms needing grid-wide sync (global barrier) had to use kernel launch boundaries (expensive) or atomic-based hacks (error-prone).

Cooperative Groups Hierarchy

Key Features

Tiled Partition (Sub-Warp Groups)

auto tile = cg::tiled_partition<16>(this_thread_block());
tile.sync();  // Synchronize 16 threads
int val = tile.shfl(data, 0);  // Shuffle within 16-thread tile

• Enables warp-level primitives (shuffle, vote, reduce) on smaller groups.
• Useful for: Sub-warp reductions, cooperative matrix operations.

Grid-Wide Synchronization

auto grid = cg::this_grid();
// ... phase 1 computation ...
grid.sync();  // ALL blocks in grid synchronize
// ... phase 2 uses results from phase 1 ...

• Eliminates need for kernel launch boundary between phases.
• Requires cooperative launch: cudaLaunchCooperativeKernel().

Thread Block Clusters (CUDA 12 / H100)

• New hardware feature: Groups of thread blocks that can synchronize and access each other's shared memory.
• cluster.sync(): Barrier across all blocks in cluster.
• Distributed shared memory: Block A can read Block B's shared memory directly.
• Enables algorithms that need cross-block communication without going through global memory.

Use Cases

• Iterative algorithms: Grid sync eliminates kernel re-launch between iterations.
• Global reductions: Grid-wide reduce without atomic contention.
• Graph algorithms: BFS/SSSP with grid-level synchronization between frontier expansions.
• Physics simulations: Jacobi/Gauss-Seidel iterations with grid-wide convergence check.

CUDA Cooperative Groups is a fundamental evolution of the GPU programming model — it replaces the rigid warp/block/grid hierarchy with flexible, composable thread groups that enable algorithms requiring cross-block cooperation to run efficiently without kernel launch overhead.