Reductions and warp shuffles

Summing a block's values as a tree in shared memory, three ways, then register to register with warp shuffles.

Loading the animation…

Concept

Summing an array looks serial, but addition is associative, so a block can add its values as a tree: half the threads each add one pair, then a quarter, and so on, $\log_2 n$ steps for $n$ values, with a __syncthreads() between steps so every thread sees the previous step's sums. The animation sums 64 numbers with one block of 64 threads, which is two warps, four ways. They are the first three kernels of Mark Harris's classic "Optimizing Parallel Reduction in CUDA", then the warp-shuffle version that later CUDA made possible.

Divergent (if (tid % (2*s) == 0)): the threads that work are spread out, every other one, then every fourth... Both warps keep running at every step with most lanes masked off (chapter 3): the hatched cells.
Strided index (i = 2*s*tid): the working threads are now packed at the front, so whole warps drop out together and divergence is gone. But thread $t$ now touches word $2st$ , a stride of $2s$ words, and the reads collide in the banks: a 2-way conflict at every step (chapter 5).
Sequential addressing (if (tid < s) x[tid] += x[tid + s], with $s$ halving): threads read consecutive words (no conflicts) and the working threads are packed at the front (no divergence beyond the last warp). The first step adds the second warp's values onto the first's.
Warp shuffle: inside a warp, __shfl_down_sync reads another lane's register directly. Five shuffles (by 16, 8, 4, 2, 1) sum a warp's 32 values with no shared memory and no __syncthreads(); then lane 0 of each warp writes one value, and thread 0 adds the warps' sums.

Watch the counters: the three tree versions all make 189 shared-memory accesses and 6 barriers; the shuffle version makes 4 and 1.

Maths

A tree over $n$ values takes $\log_2 n$ steps and $n - 1$ additions in total, the same work as a serial loop, but its depth is $\log_2 n$ instead of $n - 1$ . At step $k$ (from 1) only $n / 2^k$ threads add, so the average share of threads busy is

\frac{1}{\log_2 n} \sum_{k=1}^{\log_2 n} 2^{-k} = \frac{1 - 1/n}{\log_2 n},

under 17% for $n = 64$ : reductions are bandwidth- and latency-bound, never compute-bound, which is why fewer barriers and no shared-memory round trips (the shuffle version) matter.

For the strided version, thread $t$ 's first read at stride $s$ is word $2st$ . Within a warp, lanes $t$ and $t'$ hit the same bank when $2s(t - t') \equiv 0 \pmod{32}$ , so as long as the warp's 32 (or fewer) active lanes span more than 32 words, two of them share each bank: a 2-way conflict, as the model counts.

Code

The shuffle step, cut from src/lib/gpu/model.ts: lane $\ell$ adds the value of lane $\ell + o$ , and a lane whose partner is off the end of the warp gets its own value back, as __shfl_down_sync does.

const src = lane + o < 32 ? lane + o : lane;
next[32 * w + lane] =
  (vals[32 * w + lane] as number) + (vals[32 * w + src] as number);

The same in CUDA, illustrative and not compiled by this site's CI:

__device__ float warp_sum(float v) {
  for (int o = 16; o > 0; o >>= 1)
    v += __shfl_down_sync(0xffffffff, v, o);  // lane l reads lane l + o
  return v;                                   // lane 0 holds the sum
}

The four versions, counted (model)

Version	Shared-memory accesses	`__syncthreads()`	Worst bank conflict	Divergent warps (first step)
Divergent (`tid % 2s`)	189	6	1-way (none)	2 of 2
Strided index	189	6	2-way	0 of 1
Sequential addressing	189	6	1-way (none)	0 of 1
Warp shuffle	4	1	–	none