CUDA Cooperative Groups
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
| Group Level | Size | Sync Support | Since |
|---|---|---|---|
| Thread (1) | 1 thread | N/A | — |
| Coalesced Group | 1-32 threads | Yes | CUDA 9 |
| Tile (sub-warp) | 1,2,4,8,16,32 | Yes | CUDA 9 |
| Thread Block | Up to 1024 | Yes (replaces __syncthreads) | CUDA 9 |
| Thread Block Cluster | Multiple blocks | Yes | CUDA 12 / H100 |
| Grid Group | Entire grid | Yes (cooperative launch) | CUDA 9 |
| Multi-Grid | Multiple GPUs | Yes | CUDA 9 |
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.