CUDA cooperative groups and grid-wide sync
Someone learning CUDA from the cooperative groups documentation, on a compute capability 7.5 card:
"when I currently call this function from my kernel, it runs only on the first block and ignores the other. I.e., when I try to run this kernel on n=8 array {1,2,3,4,5,6,7,8} with 2 blocks and 4 threads per block, I get results only from the first half of the array"
https://forums.developer.nvidia.com/t/how-to-launch-cuda-cooperative-groups-standard-deviation-example-kernel/241166 (checked 2026-08-30)
The compiler does not reject that program. The header compiles, the kernel
links, and an ordinary launch still returns cudaSuccess. The kernel produces
a per-block result because the launch did not enable a grid-wide barrier.
Cooperative groups can provide a barrier across every block in the grid. You must enable this barrier through the launch API before the kernel starts. This page shows how to do that, size the grid with an occupancy query, and handle an oversized launch.
Four groups, and only one of them costs anything
Cooperative groups is a header, <cooperative_groups.h>, and a set of types
that name the threads you want to talk to. They all carry sync(),
thread_rank() and num_threads(). What differs is who is in the group and
what agreement costs.
cg::this_thread_block() returns the current
block, and its sync() has the same scope as
__syncthreads(). A thread_block argument also
shows at the call site that a device function may use a block-wide barrier.
Day 14 covers what happens when a caller cannot see
that barrier.
cg::tiled_partition<32>(block) cuts the block into fixed-size tiles, eight
of them in a 256-thread block. A tile is a warp-shaped
group carrying sync(), shfl_down(), ballot() and cg::reduce(), which
is day 23's warp shuffle code with the mask
arithmetic done for you. The guide limits the sizes to "native hardware
sizes, 1/2/4/8/16/32"; above 32 you also need a shared-memory allocation on
compute capability 7.5 and below, so this course stays at 32.
cg::coalesced_threads() names the lanes that actually took this branch, and
it is opportunistic: the call "returns the set of active threads at that point
in time, and makes no guarantee about which threads are returned (as long as
they are active)".
cg::this_grid() names every thread in the launch. Its sync() is the one
that changes the launch contract: "synchronizing the entire grid requires using the
cudaLaunchCooperativeKernel runtime launch API".
Why a grid barrier changes the launch contract
The intuition you arrive with is that a barrier is a barrier, so grid.sync()
is __syncthreads() over a bigger set. On a CPU that is true. Here it is not,
and day 2 already gave you the reason.
A block runs as a unit on one SM. By the time its kernel code runs, every thread in that block is resident, so its barrier does not depend on an unscheduled block. The barrier can still make resident threads wait.
A grid-wide barrier has a different limit.
The measured T4 has 40 SMs. Each SM supports at most 16 blocks and 32 resident warps, and the block size decides which limit applies.
At 256 threads, each block uses 8 warps, so an SM can hold 4 blocks. On this GPU, 4 times 40 is 160 resident blocks.
The 16-block limit applies only at 64 threads per block or fewer on this GPU. Day 2's sweep prints both values: 16 blocks per SM at 48 and 64 threads, and 4 blocks per SM at 256.
Launch 16,000 blocks of 256 and 15,840 of them are queued until a resident block finishes. A barrier where most participants have not started cannot complete, and nothing can start them: the blocks holding the SMs are the ones waiting at it.
So residency has to be arranged before the kernel starts. A cooperative launch is that arrangement, and it has a size.
// Device: compiles, links, and runs.
cg::grid_group grid = cg::this_grid();
firstStage(grid);
grid.sync(); // only a barrier if the launch promised residency
// Host: the launch that made no such promise.
myKernel<<<blocks, 256>>>(args); // returns cudaSuccess
The guide requires the cooperative launch API and does not describe what
happens if you skip it. cg::grid_group carries is_valid(), which "Returns
whether the grid_group can synchronize", and that is the question to ask from
inside the kernel rather than guess at from outside it.
Sizing a grid you are allowed to launch
Full program in
code/day28-cooperative-groups/cooperative_reduce.cu.
It sums 4,194,915 floats twice: once as day 25's two launches, once as one
launch with a grid barrier in the middle. Three rules hold the comparison
together.
Both paths run the same grid. They use the same block size, grid-stride
loop, block reduction, and arithmetic. Only the step between the two stages
differs: a kernel boundary or grid.sync().
Different grids would measure a grid-size change too.
The grid comes from the occupancy API, not from the problem size. The runtime documents the ceiling in exactly those terms: "The total number of blocks launched cannot exceed the maximum number of blocks per multiprocessor as returned by cudaOccupancyMaxActiveBlocksPerMultiprocessor ... times the number of multiprocessors as specified by the device attribute cudaDevAttrMultiProcessorCount". It is a per-kernel number, so the program asks it separately for each kernel it launches.
Nothing here runs a barrier that might not complete. The oversized launch
uses a probe kernel with no barrier in it, so the experiment is safe whichever
way the runtime answers. Day 14 made the same call about a non-uniform
__syncthreads().
cudaLaunchKernelEx takes a config struct instead of the
triple chevron. Its grid, block,
shared-memory and stream fields are what <<<>>> already gives you. The
attribute array is what <<<>>> cannot express:
cudaLaunchAttribute coopAttr[1] = {};
coopAttr[0].id = cudaLaunchAttributeCooperative;
coopAttr[0].val.cooperative = 1;
cudaLaunchConfig_t config = {};
config.gridDim = dim3(static_cast<unsigned int>(blocks));
config.blockDim = dim3(static_cast<unsigned int>(kThreadsPerBlock));
config.dynamicSmemBytes = 0;
config.stream = nullptr;
config.attrs = coopAttr;
config.numAttrs = 1;
A non-zero value there gives a launch "with exactly the same usage and semantics of cuLaunchCooperativeKernel". Learn the shape once: days 58 and 77 pass different attributes through the same struct.
The kernel it launches is one function where day 25 needed two:
__global__ void reduceCooperativeGrid(const float* __restrict__ in,
float* __restrict__ partials,
float* __restrict__ out, size_t n) {
cg::grid_group grid = cg::this_grid();
cg::thread_block block = cg::this_thread_block();
cg::thread_block_tile<32> tile = cg::tiled_partition<32>(block);
const size_t step = grid.num_threads();
float v = 0.0f;
for (size_t i = grid.thread_rank(); i < n; i += step) {
v += in[i];
}
float total = blockSum(block, tile, v);
if (block.thread_rank() == 0) {
partials[blockIdx.x] = total;
}
// Every write above is visible to every thread below. That is the second
// half of the sync() contract and it is what replaces the kernel
// boundary the two-launch version pays for.
grid.sync();
if (blockIdx.x == 0) {
const size_t m = grid.num_blocks();
float w = 0.0f;
for (size_t i = block.thread_rank(); i < m; i += block.num_threads()) {
w += partials[i];
}
total = blockSum(block, tile, w);
if (block.thread_rank() == 0) {
out[0] = total;
}
}
}
grid.thread_rank() and grid.num_threads() are the cooperative-groups
spelling of the two index lines you have written since day 4. The names
change, the addresses do not.
Note. Holding the two-launch version to the cooperative grid is a choice. At 256 threads a block that grid already fills every warp slot on every SM, so the baseline is not handicapped, but it is not the grid you would have written by hand for that path.
Results
Measured. Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -O3 -arch=sm_75. Captured 2026-08-30 on the project's verification node; full transcript in the page's evidence file.
GPU: Tesla T4 (compute capability 7.5), 40 SMs
cudaDevAttrCooperativeLaunch: 1
n = 4194915 floats, 16.0 MiB, 256 threads/block
How big a cooperative grid fits
reduceCooperativeGrid 4 blocks/SM x 40 SMs = 160
probeCooperative 4 blocks/SM x 40 SMs = 160
a <<<>>> launch sized to n would use 16387 blocks
both sums below run 160 blocks, the same grid
What cg::this_grid() reports
<<<160, 256>>> is_valid=0 blocks=160
cooperative, 160 blocks is_valid=1 blocks=160
cooperative, 161 blocks cudaErrorCooperativeLaunchTooLarge: too many blocks in cooperative launch
Sum of 4194915 floats
CPU reference, double 3670548.750000
two launches 3670548.750000
one cooperative launch 3670548.750000
What cg::coalesced_threads() returns
warp lane group size rank in group
0 2 3 0
0 4 3 1
0 8 3 2
1 0 16 0
1 1 16 1
1 2 16 2
1 3 16 3
1 4 16 4
1 5 16 5
1 6 16 6
1 7 16 7
1 8 16 8
1 9 16 9
1 10 16 10
1 11 16 11
1 12 16 12
1 13 16 13
1 14 16 14
1 15 16 15
Time, mean of 10 runs after 3 warm-ups, copies excluded
two launches 0.0819 ms
one cooperative launch 0.0830 ms
all checks passed
A grid-wide barrier limits the grid size. A cooperative launch on this GPU fits 160 blocks, 4 per SM across 40 SMs. An ordinary launch sized to the same input would use 16,387 blocks.
Grid sync requires every block to be resident at the same time because a block that has not started cannot reach the barrier. The device's resident block count, not the input size, sets the limit.
Ask for one more block and the launch fails. A launch of 161 blocks returns
cudaErrorCooperativeLaunchTooLarge. The runtime rejects it instead of
starting a grid whose barrier cannot complete.
The program uses a probe kernel without a barrier to test this error without risking a hang.
cg::this_grid().is_valid() is how you tell. It returns 0 for an ordinary
<<<>>> launch and 1 for a cooperative one, on the same kernel with the same
grid. Calling sync() on an invalid group is undefined, so this is the guard
to write rather than assume.
Both paths produce the same sum. The cooperative version is not faster here; it is one launch instead of two, and it costs a grid two orders of magnitude smaller. That trade is the lesson.
Run it yourself
A free Colab T4 or any card you own, on Linux. There is no Compiler Explorer embed: the program is past this site's 60-line embed budget, and the exercise wants a shell. The build line is in the repo's README:
nvcc -std=c++17 -O3 -arch=sm_75 -o cooperative_reduce cooperative_reduce.cu
Hardware. Cooperative launch needs compute capability 6.0 or higher and one of three platforms the guide names: Linux without MPS, Linux with MPS on a device of compute capability 7.0 or higher, or the latest Windows. Ask the device rather than assume, with
cudaDeviceGetAttributeandcudaDevAttrCooperativeLaunch. If the answer is 0, or you have no card at all, start at /setup/learn-cuda-without-a-gpu.
Exercise
Write solve() for a single-launch sum: one cooperative kernel, a grid you
size yourself, the total in out[0]. The harness gives you the stream and
never tells you n in advance.
Time: 30 to 40 minutes. Submit: your solve(), plus the two block
counts your own card reports for your kernel.
Check: the harness runs the nine case-ladder sizes, 0 through 67108864,
against a double-precision reference and reports the smallest failing case
first. Two failures belong to this day: a grid sized from n fails the two
largest cases with a launch error rather than a wrong number, and a zero-block
grid fails the empty case before the kernel runs. A wrong sum at 611 and a
right one at 1024 is the bounds bug from
day 8.
Hint 1
Your grid can no longer be a function of n alone. Something else has to be
true of it, and that something is not in your source file and not on your
card's spec sheet.
Hint 2
Use two values. Query the active blocks per SM for your kernel and block size,
then read the SM count from cudaDeviceProp. Multiply them and use the
smaller of that value and the block count needed for n.
Then decide what the code should do when n = 0.
Solution
The grid is min(ceil(n / threads), blocksPerSM * multiProcessorCount), with
blocksPerSM from cudaOccupancyMaxActiveBlocksPerMultiprocessor called on
your kernel, at your block size, with your dynamic shared memory size. Ask
about the kernel you are launching, not a similar one: registers and shared
memory move the answer.
At n = 0 the grid is 0 blocks. Return without launching; the harness has
already zeroed out.
The reference is reduceCooperativeGrid in
cooperative_reduce.cu,
and the counts it reaches on a T4 are in the Results table above: 4 blocks
per SM across 40 SMs, so 160. Your kernel will report its own number, and if
it differs the difference is your registers and your shared memory.
What outlives this: a cooperative launch turns the grid from a way of describing your problem into a resource you have to fit inside.
Pitfalls
grid.sync() under an ordinary launch. It compiles, links, and returns
cudaSuccess because the launch API, not the header, enables cooperative
residency. The guide does not define grid.sync() under an ordinary launch.
Call grid.is_valid() inside the kernel before using the grid group.
Sizing the cooperative grid from n. The ceiling is blocks per SM times
SM count, and past it the launch is refused with too many blocks in cooperative launch. This is the good failure: it happens at the launch, with
a name, before anything runs.
Reading that ceiling as a property of the card. It belongs to the kernel: add a register array and it drops, which is day 17's subject, and day 45 turns the same query into occupancy.
Building a group inside a branch that not every thread reaches.
tiled_partition, sync, and cg::reduce are collectives.
The guide says: "Partitioning a group is a collective operation and all threads in the group must participate. If the group was created in a conditional branch that not all threads reach, this can lead to deadlocks or data corruption." Day 14 is the block-scoped version of the same rule.
Expecting cg::reduce to be one instruction. The hardware path covers
4-byte integer add, min, max and the bitwise operators on compute capability
8.0 and higher. The guide's own example notes that cg::plus<float> "fails to
match with an accelerator and instead performs a standard shuffle based
reduction", so a float reduction is shuffles on every card.
Assuming one launch must beat two. Launch overhead is small, but a grid barrier also takes time and limits the grid size. Measure both versions as shown on day 9.
Go deeper
- CUDA Programming Guide 4.4, "Cooperative Groups", for the group types and the cooperative launch rules: https://docs.nvidia.com/cuda/cuda-programming-guide/04-special-topics/cooperative-groups.html (checked 2026-08-30)
- CUDA Programming Guide 5.6.3, "Cooperative Groups API", for
grid_group::is_valid(), the tile-size limits andcg::reduce: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/device-callable-apis.html (checked 2026-08-30) - CUDA Runtime API,
cudaLaunchCooperativeKernelandcudaLaunchKernelExC, for the block-count ceiling and the config struct: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__EXECUTION.html (checked 2026-08-30) cuda-samples,cpp/2_Concepts_and_Techniques/reductionMultiBlockCG, NVIDIA's own cooperative reduction: https://github.com/NVIDIA/cuda-samples/tree/master/cpp/2_Concepts_and_Techniques/reductionMultiBlockCG (checked 2026-08-30)- Programming Massively Parallel Processors, 4th edition, chapter 10, "Reduction: And minimizing divergence": https://shop.elsevier.com/books/programming-massively-parallel-processors/hwu/978-0-323-91231-0 (checked 2026-08-30)
Next
Day 29 privatizes a histogram, where a per-block copy of the counters and a
merge beat one global counter array and none of this is needed. Day 56 is the other answer to the launch
cost you just measured: capture the chain as a graph and pay once.
cudaLaunchKernelEx returns on days 58 and 77 with different attributes in
the same struct, which is why it was worth learning here.