The thread hierarchy of Section 7.2 has three levels: a grid of blocks of threads. The blocks of a grid are deliberately independent—the programming model makes no guarantee that any two of them run concurrently, and they can cooperate only through global memory (Section 7.2.2). Beginning with the Hopper architecture (compute capability 9.0), CUDA inserts an optional fourth level between the grid and the block: the thread block cluster.
A cluster is a small group of thread blocks that the hardware
guarantees to co-schedule on a single GPU Processing Cluster (GPC), so
that every block in the cluster is resident and running at the same
time. That concurrency guarantee—which ordinary blocks do not have—is
exactly what makes cooperation between blocks possible. Up to 8 blocks
may portably form a cluster; the H100 supports a nonportable size of 16
as an opt-in. The largest size a given kernel can use on a given device
is reported by cudaOccupancyMaxPotentialClusterSize().
A cluster’s shape is set either at compile time, with the
__cluster_dims__ attribute on the kernel (as in Listing
7-5), or—more flexibly—at launch time, by passing a
cudaLaunchAttributeClusterDimension attribute (with its
clusterDim.x/y/z fields) to
cudaLaunchKernelEx(). Within the kernel, the Cooperative
Groups cluster_group exposes the cluster:
cluster.num_blocks() and cluster.block_rank()
report the cluster’s size and this block’s position in it, and
cluster.sync() is a hardware-supported barrier across every
thread of every block in the cluster—a synchronization that, before
clusters, was available only within a single block (via __syncthreads())
or across an entire cooperative grid (Section 7.6).
The capability that makes clusters compelling, though, is distributed shared memory. The shared memory of all the blocks in a cluster is exposed as a single address space: a thread in one block can read, write, and perform atomics on the shared memory of any other block in the same cluster. A cluster therefore commands as much as eight times the fast on-chip memory of a single block—the per-block allocation times the number of blocks—as one shared cooperative working set.
A remote block’s shared memory is reached through the cluster group.
Given a pointer into this block’s own shared memory,
cluster.map_shared_rank() returns a pointer to the same
location in the shared memory of the block with a given cluster rank. On
Hopper, these accesses travel over a dedicated SM-to-SM network within
the GPC, so blocks exchange data on-chip, without staging it through
global memory.
Because a mapped pointer is only valid while the target block is
resident, remote accesses must be bracketed by cluster synchronization.
A cluster.sync() before the first remote access ensures
that every block has allocated and populated its shared memory; a second
cluster.sync() after the last remote access ensures that no
block exits—freeing its shared memory—while another block is still
reading it. Listing 7-5 shows the pattern for a two-block cluster: each
block loads a tile into its own shared memory, the cluster synchronizes,
and then each block reads its neighbor’s shared memory through a mapped
pointer.
#include <cooperative_groups.h>namespace cg = cooperative_groups; __global__ void __cluster_dims__(2, 1, 1)clusterExchange( const float *in, float *out, int n ){ __shared__ float smem[256]; cg::cluster_group cluster = cg::this_cluster(); // Each block loads its own tile into its own shared memory. int i = blockIdx.x*blockDim.x + threadIdx.x; smem[threadIdx.x] = (i < n) ? in[i] : 0.0f; cluster.sync(); // every block's shared memory is now allocated and loaded // Read the peer block's shared memory over the SM-to-SM network. unsigned peer = cluster.block_rank() ^ 1; // 0<->1 float *peerSmem = cluster.map_shared_rank( smem, peer ); float sum = smem[threadIdx.x] + peerSmem[threadIdx.x]; cluster.sync(); // no block may exit until all remote reads are done if ( i < n ) out[i] = sum;}Distributed shared memory relaxes one of the oldest rules in the CUDA programming model—that thread blocks cannot cooperate—for the special case of a handful of co-scheduled blocks. It is most useful when a working set is slightly too large for one block’s shared memory but fits comfortably in a cluster’s: the large tiles of GEMM and convolution, block-wide reductions and scans that would otherwise stage partial results through global memory, and stencil or histogram computations that share halo or bin data among neighboring blocks.