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.