100 DAYS OF GPU Course preview
Day 24Module 3
in-technical-review

Parallel reduction, versions 1 to 4

Day 15 measured a 32-way shared memory bank conflict on a Tesla T4 and it cost 31.10 times the conflict-free access. The common sequence of reduction kernels has a version 2 that removes a divergent branch and adds bank conflicts on purpose, and courses still teach it after version 1.

The result needs an explanation. Either the conflicts version 2 introduces are cheap here, which would sit badly next to day 15, or the divergence version 1 pays is expensive enough to cover them. Most accounts of this sequence do not say by how much.

This page builds all four kernels with one change between each neighbouring pair, counts what each change removes before anything runs, and then measures. By the end you can measure each change on your own card rather than relying on the 2007 hardware used for the original sequence.

What a block reduction actually costs

A block reduction sums one thread block's slice of an array into a single float. Every thread loads one element into shared memory, and then the block folds the tile in half over and over: 128 additions, then 64, then 32, down to one. For a 256-thread block that is 8 steps and 255 additions, which the program asserts at compile time rather than claiming in prose.

Two hundred and fifty-five additions take little GPU time. Each version below performs exactly the same 255 additions, so these versions do not reduce the addition count.

The versions change how many warps have to run the loop body at each step, and what those warps' addresses do to the banks. A warp is 32 threads issuing together. If one of its lanes is still active, the whole warp runs the body and masks the other 31 off, so the unit of cost is the warp and not the thread. Between every step sits a __syncthreads(), which every thread in the block must reach, so a block cannot leave a step behind until its slowest warp completes it.

That gives one number to count per version: warp-wide executions of the loop body, summed over the whole tree, for one block. Count it before you run anything, then see how much of it the clock agrees with.

The intuition that costs you version 2

The belief you arrive with is that removing a divergent branch makes a kernel faster, and that is usually true. Warp divergence is real, day 22 measures what it costs, and version 1's tid % (2 * s) == 0 is a textbook case: at the first step the active threads are the even ones, so every warp in the block runs the body with half its lanes switched off.

Version 2 fixes exactly that and nothing else. It moves the branch onto a packed index, so the active threads of a step are the first few and every warp above them retires. Count the body runs for a 256-thread block and version 1 comes to 47 while version 2 comes to 12, a factor of four.

Then look at where version 2's active lanes read. Lane L reads shared word 2s * L, which puts gcd(2s, 32) lanes in the same memory bank holding different words, and by step 16 every active lane is in bank 0. Weight each of those 12 body runs by how many ways it splits, and version 2 comes back to 47.

Version 2 adds the same weighted cost in bank conflicts that it removed in divergence, and it does so at every block size from 64 threads to 1024. That is arithmetic, not a measurement, and it is the first thing this page's run should test.

One change per rung, and a benchmark that cannot cheat

Full program in code/day24-reduction-1/reduction.cu.

Every row reduces the same input down to one float, not to a partial array. Version 4 launches half as many blocks, so it leaves half as many partials for the next pass, and timing only the first launch would credit it for work it moved rather than removed. That is day 11's benchmark bug in another form.

The counting model is a static_assert, not a paragraph. The body counts, the conflict weighting and the n - 1 additions all live in constexpr functions. If the model is wrong the file does not build.

Time with events and warm up each version. Day 9 covers why a host clock around an asynchronous launch measures the launch and not the kernel.

The input is 1.0f everywhere and the element count is below 2^24, so every partial sum in every version is a whole number float holds exactly. Reduction order cannot change the answer here, which is why the correctness check is an exact comparison with no tolerance. Day 68 is where that stops being true.

Version 1, the divergent branch:

    for (unsigned int s = 1; s < width; s *= 2) {
        if (tid % (2 * s) == 0) {
            tile[tid] += tile[tid + s];
        }
        __syncthreads();
    }

Version 2, the same tree with the branch on a packed index:

    for (unsigned int s = 1; s < width; s *= 2) {
        const unsigned int index = 2 * s * tid;
        if (index < width) {
            tile[index] += tile[index + s];
        }
        __syncthreads();
    }

Version 3, sequential addressing. The stride halves instead of doubling, so an active lane reads word tid and the warp covers 32 consecutive words, one per bank:

    for (unsigned int s = width / 2; s > 0; s >>= 1) {
        if (tid < s) {
            tile[tid] += tile[tid + s];
        }
        __syncthreads();
    }

Versions 2 and 3 activate the same threads at every step, in the same numbers. Only the address moves, which makes the step between them a clean experiment on bank conflicts and nothing else. The file asserts that equality.

Version 4 keeps version 3's tree and gives each thread two elements to add before the tree starts, which halves the grid:

    const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) * 2 + tid;

    float sum = (i < n) ? in[i] : 0.0f;
    if (i + blockDim.x < n) {
        sum += in[i + blockDim.x];
    }
    tile[tid] = sum;
    __syncthreads();

Note. The first version of that index was the tidier-looking one below, and it is not a partition. Simulating four blocks of 256 threads over 2,048 elements: every odd element is read zero times, 896 of the 1,024 even ones are read twice, and the block sums 1,920 reads where it should sum 2,048. Same family as the non-bijective index day 11 shipped and had to retract. The exact-answer check catches it, which is why the check is exact.

// Wrong. Gives one thread elements 2t and 2t + blockDim.x.
const size_t i = 2 * (blockIdx.x * static_cast<size_t>(blockDim.x) + tid);

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, warp size 32
Shared memory per block: 49152 bytes

Model per block, computed at compile time
version                           warp runs   serialised
---------------------------- -------------- ------------
1  interleaved, divergent                47           47
2  interleaved, strided                  12           47
3  sequential addressing                 12           12
4  sequential, two loads                 12           12

16776605 elements, 64 MiB in, every element 1.0f
version                       passes    time (ms)         GB/s     vs v1
----------------------------  ------   ----------   ----------  --------
1  interleaved, divergent          3        0.574        117.3      1.00
2  interleaved, strided            3        0.718         93.9      0.80
3  sequential addressing           3        0.567        118.9      1.01
4  sequential, two loads           3        0.312        215.5      1.84

step speedups
  v1 -> v2  0.80x
  v2 -> v3  1.27x
  v3 -> v4  1.82x
largest single step: v3 -> v4

CUDA 13.0 made every row roughly twice as slow in absolute time, but preserved the useful ratios and every exact-sum gate. Version 2 remained a loss at 0.79x version 1, version 3 recovered to 1.01x, and version 4 remained the largest step at 1.90x version 3. The release comparison therefore changes the throughput numbers, not the conclusion.

The model predicted the version order, but version 2 was slower than version 1. Version 2 removes the divergence that version 1 has, and the compile-time model says it trades that for bank conflicts one for one: 47 serialised warp runs either way. The measurement is worse than that. Version 2 runs at 93.9 GB/s against version 1's 117.3, so on this card removing divergence and adding conflicts is a net loss, not a wash.

This result shows why you should measure each version. Version 2 is an intermediate form used to explain version 3, not a kernel to ship.

Version 3 matches the model. Sequential addressing removes both problems and the serialised count drops from 47 to 12, but the time barely moves, 118.9 against 117.3. The kernel is bandwidth-bound and the instruction savings do not reduce its time.

Version 4 reaches 215.5 GB/s, a 1.84x gain. Adding a load during the first reduction halves the number of blocks and therefore halves the reads. Fewer bytes is the only thing on this page that moved the clock, which is the lesson the rest of the module keeps repeating.

Run it yourself

Colab's free T4 or any card you own. The build line is in the repo's README:

nvcc -std=c++17 -O3 -arch=sm_75 -o reduction reduction.cu

The program allocates a 64 MiB device buffer, two small scratch buffers and a 64 MiB host vector, then runs about 170 launches. That is comfortably inside Compiler Explorer's 20 second run cap on paper, but the input has to be built and copied first and the runner is a shared spot instance, so a card you control gives a table you can trust. Target sm_75 or lower either way.

Exercise

Rebuild the program at 64, 256 and 1024 threads per block, and report the ratio of version 1's time to version 3's at each. Then say in two sentences why the ratio moves the way it does.

Time: 25 to 40 minutes. Submit: the three ratios, and the two sentences.

Check: the program compares each version's sum against the element count exactly, with no tolerance, at whatever block size you build it with. On failure it names the version and what it summed to, and returns EXIT_FAILURE.

The model table it prints recomputes itself from the new block size, so you can read the counted ratio off the first two columns before you look at the clock. Your three ratios should not be equal.

Hint 1

The tree's shape does not change when the block grows: it is still one addition per element, still arranged in halvings. Something else is changing, and it is countable before you time anything.

Hint 2

At step s, version 1 activates the threads whose id is a multiple of 2s. How many of the block's warps hold at least one of those, and how does that number behave differently when the block has 2 warps instead of 32?

Solution

Divergence costs whole warps, so its cost grows with the number of warps in the block. While 2s is at or below 32, version 1's active threads sit in every warp and every warp runs the body, while version 3's are packed into the low warps and the rest retire. A bigger block leaves more warps partly idle, so the gap widens: the counted ratio of body runs is 11 to 6 at 64 threads, 47 to 12 at 256 and 191 to 36 at 1024.

Your measured ratios will be far smaller, because the kernel is bounded by how fast the card reads 64 MiB.

The exercise shows that removing divergence saves more work in a large block than in a small one.

Pitfalls

You time the first launch and call it the reduction. Version 4 halves the blocks, so it halves the partials the next pass must sum. Compare complete reductions, n floats down to one, or you credit a version for work it only moved.

You put the __syncthreads() inside the if. The barrier belongs after the branch, because not every thread takes the branch and a barrier some threads never reach makes the program undefined. It will often still print the right answer, which is what makes it dangerous rather than obvious. Day 14 reproduces it and proves it with a tool that does not care which answer you got.

You assume the reduction tree takes most of the time. This kernel reads 64 MiB and performs one addition per element, so it is bound by memory bandwidth long before it is bound by adds. Day 42 is where you find out which bound you are on.

You decide tid % (2 * s) is the expensive part. With a compile-time loop bound the loop unrolls, 2 * s becomes a constant power of two and the modulo becomes a mask. Whether that happened on your toolkit is a question for day 46, not for a guess. Write the bound as a runtime blockDim.x and the loop cannot unroll, and then the modulo really is expensive, which is a confound and not a finding.

You expect the numbers from the slides. The four version names come from Mark Harris's deck, measured on a G80 in 2007. The scheduler, the caches and the ratio of arithmetic to bandwidth have all moved since. Reproduce the version sequence, not the table.

Go deeper

Next

Day 25 finishes the comparison: unrolling the last warp, unrolling all of it, giving each thread many elements instead of two, and replacing the last five steps with a warp shuffle. The first of those removes a barrier, which is where copying a 2007 slide stops being safe.