Block reduction

Warp tree

for (int off = 16; off > 0; off >>= 1)
  v += __shfl_down_sync(0xffffffffu, v, off);

Five steps, each halving the lanes that still hold a partial sum. Lane 0 ends with the warp total, and the dependency depth is five adds rather than 31. Each step is a register exchange, so nothing goes through memory.

Across warps

Each warp’s lane 0 writes to __shared__ float warp_sum[8], then __syncthreads(), then warp 0 loads warp_sum[threadIdx.x] for lanes that hold a warp sum and runs the same five shuffles. Thread 0 writes the block result. In a loop, the barrier before the next iteration’s write is not optional. __shfl_down_sync replaced the unsynchronised __shfl_down in CUDA 9, and its mask semantics assume Volta or later; check the signature in the guide for your toolkit.

Order changes the answer

A serial accumulation adds 32 values in 32 dependent steps; the tree does it in five. Both are correct and they round differently, because floating-point addition is not associative. The tree is faster, and its last bits are not the serial last bits. Record the order with the number and compare against a reference with a tolerance rather than for equality.

Things that break

The mask names the threads that participate. 0xffffffff in a block of 100 threads claims 32 lanes in a warp where four arrive, which is undefined and can hang. Build the mask with __ballot_sync from the same predicate the branch uses, and do not use __activemask as a membership mask. A block that is not a power of two needs if (tid < off) on every shared-memory step. And a reduction whose order changes between runs gives you a benchmark you cannot compare with itself, which is worse than a slow one.