100 DAYS OF GPU Course preview
Day 77Module 8
draft

Thread block clusters and distributed shared memory

Day 24 reduced each block's input to one shared-memory value, then wrote the partials to global memory for a second kernel. Day 26 replaced that second pass with one global atomic per block.

Both methods follow the rule from day 13: shared memory belongs to one block. Compute capability 9.0 adds thread block clusters, which let blocks in a cluster access each other's shared memory. This lesson compares an eight-block cluster reduction with the atomic version.

The level between a block and the grid

A thread block cluster is a group of blocks with a residency contract. The programming guide draws the analogy itself: "Similar to how threads in a thread block are guaranteed to be co-scheduled on a streaming multiprocessor, thread blocks in a cluster are also guaranteed to be co-scheduled on a GPU Processing Cluster (GPC) in the GPU" (CUDA C++ Programming Guide 2.2.1, https://docs.nvidia.com/cuda/archive/12.6.2/cuda-c-programming-guide/index.html , checked 2026-09-01). A GPC is a hardware partition holding a handful of SMs, so the guarantee means every block of the cluster is running, on neighbouring SMs, at the same time.

That contract is what makes distributed shared memory legal. A block's shared memory exists only while the block is resident on its SM. The cluster keeps every member block resident while remote accesses occur.

Inside a cluster, the guide states that "threads that belong to a thread block cluster, can read, write or perform atomics in the distributed address space, regardless whether the address belongs to the local thread block or a remote thread block" (same guide, 3.2.5, checked 2026-09-01). The distributed space is nothing new being allocated: it is the cluster's ordinary per-block tiles, cross mapped, so a cluster of 8 blocks with 1 KiB each sees 8 KiB.

The contract has a cost and a limit. Co-scheduling blocks reduces the scheduler's freedom to hide latency, and the guide caps the portable size: "a maximum of 8 thread blocks in a cluster is supported as a portable cluster size in CUDA" (2.2.1). Bigger clusters are an architecture-specific opt-in, smaller GPUs and MIG slices may allow fewer, and cudaOccupancyMaxPotentialClusterSize answers for the card you are on; today's program prints it.

gridDim still counts plain blocks, and the grid dimension must be a multiple of the cluster size.

Diagram: where a cluster sits in the hierarchy. Three nested boxes left to right, matching scale to hardware: threads in a block above one SM, blocks in a cluster above one GPC, clusters in a grid above the whole GPU. Band 1, one block on one SM: 256 threads, one shared tile. Caption "the guarantee you know: one block, one SM, one tile." Band 2, one cluster on one GPC: 8 blocks, arrows crossing between their tiles. Caption "new at CC 9.0: 8 blocks co-resident, 8 tiles mutually readable." Band 3, the grid across the GPU: clusters with no arrows between them. Caption "between clusters, nothing changes: global memory or nothing." Alt text: "Three hierarchy levels. A block's 256 threads share one tile on one SM. A cluster's eight blocks, co-scheduled on one GPC, read each other's tiles. Separate clusters still meet only in global memory."

Hardware extends the shared-memory rule

Before CC 9.0, no instruction could form an address into another block's shared-memory tile. This is why day 24's tree ends at the block boundary. From 9.0, a mapping instruction and cluster barrier make cross-block access safe, and the cooperative groups API from day 28 grew a cluster_group to expose them.

The compute capability table marks thread block clusters and distributed shared memory present on 9.0, 10.x, 11.0 and 12.x (https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html , checked 2026-08-29), so an RTX 50 gaming card runs everything on this page.

What a GeForce card lacks is hardware TMA multicast, one copy feeding a tile to every block of a cluster at once, and that absence is why CUTLASS declines to use clusters there at all: "On Geforce series graphics card, there is no multicast feature therefore the cluster shape is fixed to 1x1x1" (https://github.com/NVIDIA/cutlass/blob/main/media/docs/cpp/blackwell_functionality.md#cluster-size , checked 2026-08-29). Cluster support does not imply support for Hopper-style cluster GEMMs. Day 78 maps the differences.

A cluster can be requested two ways. Baked into the kernel at compile time:

__global__ void __cluster_dims__(8, 1, 1) reduceClusterDsm(
    const float* in, float* sum, size_t n) { ... }

or chosen at launch with an attribute on cudaLaunchKernelEx. This program uses the launch attribute so it can test several cluster sizes.

One reduction tree, two completion paths

Full program: code/day77-clusters/cluster_reduction.cu.

Both kernels call the same block-level __device__ reduction function. The program keeps day 24's two-loads reduction kernel as its first stage, so only the completion path changes between rows.

Every comparison is exact by construction. The input is all ones and the count sits under 2^24, day 24's trick, so every partial and every running total is a whole number that float represents exactly and no atomic ordering can move a bit. A row that misses 16,776,605 fails the run, whatever its time was.

The baseline tail is day 26's move, one global atomic per block, 32,767 of them landing on one address:

    if (threadIdx.x == 0) {
        atomicAdd(sum, tile[0]);
    }

The cluster tail replaces it. Each block's partial stays in its own tile; after a cluster-wide barrier, rank 0 reads its neighbours' tiles through map_shared_rank and sends one sum onward:

    cg::cluster_group cluster = cg::this_cluster();

    // Every block's partial is in its own tile[0]. The barrier makes all
    // of them visible to every block in the cluster, and guarantees no
    // block has raced ahead into the combine.
    cluster.sync();

    if (cluster.block_rank() == 0 && threadIdx.x == 0) {
        float total = 0.0f;
        for (int r = 0; r < static_cast<int>(cluster.num_blocks()); ++r) {
            // The same shared array, seen through block r's mapping. Rank
            // 0's own tile comes back unchanged when r is its own rank.
            const float* remote = cluster.map_shared_rank(tile, r);
            total += remote[0];
        }
        atomicAdd(sum, total);
    }

    // Without this barrier a non-zero rank may exit while rank 0 is still
    // reading its shared memory, and the read lands in a dead address
    // space. The failure is a race, not an error message.
    cluster.sync();

Both cluster.sync() calls are required. The first is a cluster-wide version of __syncthreads(): partials written before partials read.

The second exists because of the residency contract. A block that exits takes its shared memory with it, so nobody may leave until the reads are done. The guide says it plainly: the user "needs to ensure that the shared memory read by remote thread block is completed before it can exit" (3.2.5).

The host side uses day 58's cudaLaunchKernelEx machinery with the cluster attribute in the slot PDL used:

    cudaLaunchAttribute attrs[1];
    attrs[0].id = cudaLaunchAttributeClusterDimension;
    attrs[0].val.clusterDim.x = static_cast<unsigned int>(clusterSize);
    attrs[0].val.clusterDim.y = 1;
    attrs[0].val.clusterDim.z = 1;

    cudaLaunchConfig_t cfg = {};
    cfg.gridDim = dim3(static_cast<unsigned int>(padded));
    cfg.blockDim = dim3(kThreadsPerBlock);
    cfg.dynamicSmemBytes = 0;
    cfg.stream = nullptr;  // the legacy default stream
    cfg.attrs = attrs;
    cfg.numAttrs = 1;

    CUDA_CHECK(cudaLaunchKernelEx(&cfg, reduceClusterDsm, d_in, d_sum, n));

The program attempts cluster sizes 1, 2, 4 and 8, skipping any size above cudaOccupancyMaxPotentialClusterSize for the current device. Size 1 is deliberate: every block is rank 0 of its own cluster, the atomic count does not move, and the row prices the cluster launch machinery itself.

Rank 0 combines eight values with serial reads by one thread. This costs a few hundred cycles beside 64 MiB of DRAM traffic. The exercise tests whether parallel reads change the result.

Results

Not yet run or compile-checked. The program targets sm_90 and needs a CC 9.0 or newer GPU to run. The README also requests a compile-only artifact, but none exists yet. This section contains expected results, not measurements, until the captures land in evidence/.

Mode Global atomics Time (ms) GB/s vs atomic
atomic baseline 32767 ? ? ?
dsm, cluster of 1 32767 ? ? ?
dsm, cluster of 2 16384 ? ? ?
dsm, cluster of 4 8192 ? ? ?
dsm, cluster of 8 4096 ? ? ?

Checks for an sm_90 run

  1. Every supported row sums to exactly 16,776,605 and the program exits 0. The input makes float addition exact here; any miss is a scheduling or mapping bug, not rounding.
  2. cudaOccupancyMaxPotentialClusterSize reports at least 8 for this kernel on a full H100. A smaller result is valid on a small GPU or MIG slice, and the program skips modes above whatever the API reports.
  3. All supported rows land within 10 percent of each other. Every mode reads the same 64 MiB of DRAM; the tails differ by at most 28,671 single-address atomics. If the cluster of 8 beats the baseline by more than 10 percent, atomic contention was a larger term than expected. If any cluster row loses by more than 5 percent, the residency contract cost real scheduling freedom, and that number is the interesting one.
  4. The cluster of 1 ties the baseline within 2 percent, pricing the extensible-launch path itself at roughly nothing.

The cluster moves the combine one level down the memory hierarchy. Its effect depends on whether the combine was the bottleneck. Day 34 showed the same constraint for tiling, as did the reduction tree on day 24.

Run it yourself

Run the program on a GPU with compute capability 9.0 or newer. An RTX 50 runs it as built because -arch=sm_90 embeds PTX that JIT-compiles forward to CC 12.x. The setup guide lists remote GPU options, and the planning snapshot is in FACT-SHEET.md.

Hardware. No card at all still gets you the compile: nvcc -std=c++17 -O3 -arch=sm_90 -c cluster_reduction.cu needs a toolkit, not a GPU, and the README has the exact command the T4 node uses for its compile-only check.

Exercise

Rank 0's combine is one thread walking eight tiles in sequence. Rewrite it so the first cluster.num_blocks() threads of the rank 0 block each map one remote tile and read one partial, then reduce those eight values inside the block. Before you run it, write down whether the five-row table will change, and why.

Time: 30 to 45 minutes. Submit: your modified combine, the five timed rows before and after, and one sentence saying what the change was worth and why.

Check: the program's own gate. Every mode must still sum to exactly 16,776,605; a wrong total prints the mode and the difference in elements and the run exits nonzero. A hang in your version means a thread that skipped a cluster.sync() every other thread reached.

Hint 1

Count what you are optimizing. Per timed run, how many distributed shared memory reads does the serial loop issue, and how many global memory reads did the tree below it just finish?

Hint 2

Eight parallel reads need a place to land before one thread can add them, and every thread of the block must still reach both cluster-wide barriers. Which existing shared array has eight free slots after the tree is done, and where must the block-level __syncthreads() go so thread 0 reads all eight?

Solution

Let threads 0 to 7 of the rank 0 block each call cluster.map_shared_rank(tile, r) with their own index as r, store the remote tile[0] into a second small shared array, __syncthreads(), then thread 0 adds the eight slots and issues the atomic. Both cluster.sync() calls stay exactly where they were, reached by every thread.

The table should not move. You parallelized eight loads in a kernel that reads 64 MiB from DRAM, so the before and after rows should differ only by measurement noise. Optimize the work that controls the run time.

Pitfalls

Your kernel is correct on every test run and still has a race. Dropping the second cluster.sync() lets a neighbour block exit while rank 0 is still reading its tile. On a lightly loaded card the blocks finish close together and the reads win the race for weeks. There is no error message; day 62's tools and a saturated card are how it surfaces.

The cluster launch fails before a single block runs. The grid dimension must be a multiple of the cluster size, and 9.0's launch checks it. Today's program pads the grid up to the next multiple and lets the padded blocks contribute an exact zero; the padding costs one line and removes the failure class.

You set __cluster_dims__ at compile time and the sweep stopped working. The guide is explicit: "If a kernel uses compile-time cluster size, the cluster size cannot be modified when launching the kernel" (2.2.1). Pick one mechanism per kernel; the runtime attribute is the one that can sweep 1, 2, 4, 8.

You bought an RTX 50 for the cluster speedups in a CUTLASS blog post. The card runs this page's program, but the GEMM wins in Hopper write-ups lean on TMA multicast, which GeForce Blackwell lacks; CUTLASS pins its cluster shape to 1x1x1 there. Day 78 is the map of which Blackwell has what.

You expected the 8x atomic cut to show up as an 8x speedup. The cut is real and structural, and prediction 3 still bets it is nearly invisible here, because the atomics were never the bottleneck at this size. Shrink the per-block work or fatten the contention and the same mechanism starts to matter; measure before believing either direction.

Go deeper

Next

Day 78 maps Blackwell capabilities, including wgmma and tcgen05, and the exact failure when an architecture-specific binary meets the wrong GPU. Hopper GEMM pipelines use the cluster mechanisms introduced here.