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.