Until CUDA 9, the set of threads that could cooperate was fixed by
the hardware and named only implicitly. Threads within a block
synchronized with __syncthreads() and exchanged data through shared
memory; threads within a warp were assumed to run in lockstep and
communicated through __shfl(). Neither the block nor the warp was a
value a program could name, pass to a function, or size at runtime.
Cooperative Groups, added in CUDA 9, makes the cooperating
group an explicit, typed object: a handle that knows how many threads it
contains and gives each a rank numbering it within the group,
and on which synchronization and data exchange are called as member
functions. The facility lives in the
<cooperative_groups.h> header, under the
cooperative_groups namespace (conventionally aliased
cg).
The group types form a ladder from the narrowest span to the widest:
coalesced_group – the threads of a
warp that are converged at the point
cg::coalesced_threads() is called. It is the disciplined
replacement for the pre-Volta habit of reasoning about a warp’s active
mask by hand (Section 7.5): the group is the set of lanes
executing together at that point, whatever branch divergence produced
that set.
thread_block_tile<N> – a
compile-time-sized tile of N threads, produced by
cg::tiled_partition<N>(). For N no larger
than a warp, the tile’s collectives lower to shuffle instructions, the
same code one would write by hand; the API also partitions a block into
larger tiles.
thread_block – the familiar CUDA
block, obtained with cg::this_thread_block(). Its
sync() is exactly __syncthreads().
cluster_group – the threads of a
Hopper thread block cluster, obtained with
cg::this_cluster() and covered in Section 7.8.
grid_group – every thread of the
launch, obtained with cg::this_grid().
Every group, whatever its width, exposes the same small interface:
size(), thread_rank(), and
sync(). A group can be subdivided –
tiled_partition<N>() cuts it into fixed-size tiles,
and binary_partition() splits it in two by a runtime
predicate, so the threads that took one side of a branch form one group
and the rest form another. Because a group carries its own membership
and size, the collective operations built on it –
cg::reduce(g, value, op), cg::inclusive_scan()
and cg::exclusive_scan(), and the communication primitives
g.shfl(), g.ballot(), g.any() –
work for any tile size and any block size, with no hand-written log-step
loop and no lane mask threaded through the code. The reduction of
Chapter 12 is built on exactly these collectives:
blockReduceCG() reduces a block with two calls to
cg::reduce(), one per warp and one across warps, and places
no constraint on the block size (Listing 12-1). The scans of Chapter 13
build their block scans on the same tile collectives.
A short example shows the handle in use. Warp-aggregated atomics funnel a whole warp’s increments of a shared counter into a single atomic: the lanes that arrive together elect one representative to add the group’s size, then share the base index back so each lane computes its own slot.
namespace cg = cooperative_groups;
__device__ uint32_t atomicAggInc( uint32_t *counter )
{
cg::coalesced_group active = cg::coalesced_threads();
uint32_t base = 0;
if ( active.thread_rank() == 0 )
base = atomicAdd( counter, active.size() );
base = active.shfl( base, 0 );
return base + active.thread_rank();
}The one atomic serves the entire group regardless of how many lanes are active, and the group adapts to whatever set of lanes actually reached the call – the property Independent Thread Scheduling (Section 7.5) makes it unsafe to assume.
The grid_group is the exception to the rule that any
group can synchronize. this_grid().sync() would be a
barrier across every thread of the launch, but it is sound only when
every block is resident at once – and a launch bigger than the hardware
can hold begins retiring early blocks before later ones start, so an
unconditional grid barrier deadlocks. (Section 5.2.11 noted the same
hazard behind using global-memory atomics as a grid-wide synchronization
primitive.) Cooperative kernel launches, added in CUDA 9,
address the problem at the root by guaranteeing co-residency: a kernel
launched with cudaLaunchCooperativeKernel() either fits on the machine
in its entirety or fails to launch with
cudaErrorCooperativeLaunchTooLarge.
In exchange, the kernel gains a legitimate grid-wide barrier, so
multiphase algorithms – a reduction whose result feeds a broadcast, an
iterative solver with a global convergence test – can run as a single
kernel launch instead of a sequence of launches with implicit barriers
between them. The price is that grid size is bounded by occupancy: the
application should size the grid using
cudaOccupancyMaxActiveBlocksPerMultiprocessor() (Section 7.3) multiplied
by the number of multiprocessors, and support should be confirmed via
the cooperativeLaunch device attribute. (A multidevice
variant, cudaLaunchCooperativeKernelMultiDevice(), was added alongside
it but has since been deprecated.)