The CUDA scan algorithm: the tree and the three kernels
Here is the fix for the problem this page is about, as GPU Gems 3 chapter 39 prints it today (https://developer.nvidia.com/gpugems/gpugems3/part-vi-gpu-computing/chapter-39-parallel-prefix-sum-scan-cuda , checked 2026-08-30):
#define NUM_BANKS 16
#define LOG_NUM_BANKS 4
#define CONFLICT_FREE_OFFSET(n)((n) >>
NUM_BANKS + (n) >>
(2 * LOG_NUM_BANKS))
Two things in three lines. NUM_BANKS is 16, and every GPU this course
targets has 32. And the first shift names NUM_BANKS where the second names
LOG_NUM_BANKS, inside an expression where + binds tighter than >>.
The chapter is from 2007, it is still the clearest written explanation of the scan, and the idea under that macro is exactly right. The arithmetic is for a card with 16 banks that served memory a half warp at a time, and you have neither. This page rebuilds it for 32, puts a number on what it buys at each of the twelve levels of the tree, and then takes the scan past one block to ten million elements.
Two sweeps over one tree
An exclusive scan turns [2 1 1 1] into [0 2 3 4]: every output is the sum
of everything to its left. Day 31 built one the
Hillis-Steele way, where at step s every element adds the element s places
behind it, and after log2(n) steps everyone is done.
That works and it is wasteful. GPU Gems 3 names the standard it fails:
"A parallel computation is work-efficient if it does asymptotically no more work (add operations, in this case) than the sequential version."
Blelloch's scan meets it with a balanced binary tree and two passes over it. The up-sweep is exactly day 24's tree reduction, and it leaves behind more than the total: after it, every internal node holds the sum of the leaves below it.
The root is then cleared to zero, and the down-sweep walks the same levels in reverse. At each node the left child receives the node's own value and the right child receives that value plus the left child's old one. When the sweep reaches the leaves, every leaf holds the sum of everything to its left, which is the answer.
Counted per tile of 4,096 elements, the chapter's claim holds: the tree makes 8,190 adds against Hillis-Steele's 45,057, five and a half times fewer.
Note. That widget is this page's only figure. A day with an interactive writes no separate diagram, so the widget carries its own static frame for a reader with no JavaScript, and
blelloch-upsweepis that frame.
The stride doubles, and so does the conflict
The chapter says what happens next, and says it plainly:
"Binary tree algorithms such as our work-efficient scan double the stride between memory accesses at each level of the tree, simultaneously doubling the number of threads that access the same bank."
At level stride, active lane L touches shared
words (2L + 1) * stride - 1 and (2L + 2) * stride - 1, so consecutive
lanes sit 2 * stride words apart. That is the case
day 15 measured: a warp whose
lanes are s words apart splits into gcd(s, 32) requests, because there are
32 banks and successive words go to successive ones.
The doubling stops, though, and that is the part the 2007 sentence cannot
tell you. gcd(2 * stride, 32) doubles up to stride 16 and then pins at 32,
because a warp only has 32 lanes to put in one bank. From there the
conflict degree falls again, not because the
addresses improve but because the tree runs out of active lanes.
So count warps, not levels. Sum the warp-wide requests each level costs after its conflicts split it, over both sweeps, and a 4,096-element tile comes to 1,534 for the plain tree. The Hillis-Steele scan, whose lanes stay adjacent and never conflict, comes to 1,536.
The tree saves five and a half times the adds but loses that gain to bank requests. That is the finding this page exists for, and it is arithmetic, not a measurement: the program computes both numbers at compile time and prints them before it times anything.
Padding is what removes the conflicts. Whether that buys any time is a question for the clock, and the Results section below compares the result with this page's expectation.
Give the tile one spare word every 32 and the spacing between
lanes goes from 2 * stride to 2 * stride + 2 * stride / 32. At stride 16
that is 33 rather than 32, and gcd(33, 32) is 1, the same result as day 15's
stride-33 case. It stops working once the added term is itself
even: the spacing at stride 32 is 66 and at stride 64 it is 132.
One spare word every 32 therefore cuts requests from 1,534 to 294 but leaves 2-way conflicts at strides 32 and 1,024 and 4-way ones from 64 to 512. A second spare word every 1,024 clears those as well, at 264. That second term is the one the chapter's macro includes, but most modern copies omit it.
One block scans 4,096. Ten million needs three kernels.
The tree lives in shared memory, so a tile is as big as one block can hold. At 4,096 floats that is 16 KiB, and ten million elements is 2,442 tiles. No tile knows what came before it.
The fix is three launches. Day 24's reduction also took more than one, but every launch there did the same job on a smaller array. This is the first algorithm in the course whose passes do different things, and whose last pass writes back over the first one's output:
- Scan every tile, and leave that tile's total in
blockSums[blockIdx.x]. - Scan the 2,442 totals. One block, same kernel, smaller array.
- Add tile
b's scanned total to every element of tileb. That last pass reads and writes the whole array in coalesced 128-byte runs and touches no shared memory at all.
Step 2 sets the method's size limit. It is one block, so it can only
scan kTileElems totals, and two levels therefore reach kTileElems squared
elements: 16,777,216 here. Ten million fits.
Set the tile to 1,024 and it does not. The three-kernel plan then needs a recursive scan with a stated base case. The program refuses to run rather than print a scan missing 8,742 tiles of offsets.
The obvious alternative is one kernel with a grid-wide barrier between the phases. Cooperative groups give you that, and day 28 shows the residency limit: every block must stay resident until it reaches the barrier.
The measured cooperative launch fit 160 blocks of 256 threads, while this grid wants 2,442. Three launches and two extra passes over global memory are cheaper than a grid that will not fit.
Measuring three scans that give the same answer
Full program in
code/day32-scan-2/scan_2.cu. Four
decisions in the harness carry the result.
Every row runs the whole pipeline and is checked before it is timed. All three scans go through the same three launches and produce the same 10,000,000 outputs, so the table compares algorithms and not amounts of work finished. Day 11 is the lesson about the other kind of benchmark.
The check is exact, with no tolerance. Every input is 1.0f except every
fourth, which is 2.0f, so the total is 12,500,000, under 2^24, and every
prefix lands on a float exactly.
No ordering of the adds can move a bit, so a kernel that drops one element is off by exactly one instead of hiding inside a relative tolerance. The period of 4 matters too: an all-ones input cannot catch a kernel that reads the right count of elements from the wrong addresses, and this one can.
The model is checked by the compiler. The conflict degrees above are
constexpr, and static_assert fails the build if a 4,096-element tree stops
putting all 32 lanes in one bank at stride 16, or if padding stops fixing it.
Events, and a warm-up per kernel. Day 9 covers why a host clock around a launch measures the launch.
Here are both sweeps. first is the thread's own index, span is the block
size, and the inner loop exists because a level can have more adds than the
block has threads:
// Up-sweep. Level `stride` has kTileElems / (2 * stride) adds to make,
// and a block of kThreadsPerBlock threads takes them in as many passes
// as that needs.
for (int stride = 1; stride < kTileElems; stride <<= 1) {
const int adds = kTileElems / (2 * stride);
for (int k = first; k < adds; k += span) {
const int a = (2 * k + 1) * stride - 1;
const int b = (2 * k + 2) * stride - 1;
tile[b] += tile[a];
}
__syncthreads();
}
// The tile total leaves before the root is cleared. Clearing it is what
// makes the down-sweep produce an exclusive scan, and the total is what
// the second kernel scans.
if (tid == 0) {
blockSums[blockIdx.x] = tile[kTileElems - 1];
tile[kTileElems - 1] = 0.0f;
}
__syncthreads();
// Down-sweep. The same levels in reverse: every node hands its own value
// to its left child and the sum of both children to its right child.
for (int stride = kTileElems / 2; stride > 0; stride >>= 1) {
const int adds = kTileElems / (2 * stride);
for (int k = first; k < adds; k += span) {
const int a = (2 * k + 1) * stride - 1;
const int b = (2 * k + 2) * stride - 1;
const float left = tile[a];
tile[a] = tile[b];
tile[b] += left;
}
__syncthreads();
}
The barrier sits outside the inner loop, so every
thread reaches every one of them whether or not it had an add to make. The
padded kernel is that code with each index sent through paddedWord and one
declaration changed:
__shared__ float tile[kTileWords];
And the three kernels are three lines:
static void scanAll(int variant, const float* d_in, float* d_out, float* d_sums,
float* d_offsets, float* d_total, size_t n, int blocks) {
launchTileScan(variant, blocks, d_in, d_out, d_sums, n);
launchTileScan(variant, 1, d_sums, d_offsets, d_total,
static_cast<size_t>(blocks));
addBlockOffsets<<<blocks, kThreadsPerBlock>>>(d_offsets, d_out, n);
}
One honest note about the block size. It is 1024 here, not the course default of 256, because the tile has to be big enough that one launch's block sums fit back inside a single tile.
That large block can limit how many blocks stay resident on each SM, which reduces latency hiding. The measured table below includes that cost, while the copy row does not.
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
Max threads per block: 1024
Shared memory per block: 49152 bytes
Tile model, computed at compile time
4096 elements per tile, 1024 threads per block, 12 levels
kernel adds/tile requests/tile
-------------------- ----------- ------------------
hillis-steele 45057 1536
blelloch, plain 8190 1534
blelloch, padded 8190 264
Up-sweep levels. The down-sweep repeats them in reverse.
stride lanes warps plain i/32 two-term
-------- -------- ------- -------- -------- --------
1 2048 64 2 1 1
2 1024 32 4 1 1
4 512 16 8 1 1
8 256 8 16 1 1
16 128 4 32 1 1
32 64 2 32 2 1
64 32 1 32 4 1
128 16 1 16 4 1
256 8 1 8 4 1
512 4 1 4 4 1
1024 2 1 2 2 1
2048 1 1 1 1 1
10000000 elements, 2442 tiles of 4096, 38 MiB in
every fourth element is 2.0f, the rest 1.0f, total 12500000
copy, no scan: 0.325 ms, 245.9 GB/s over 76 MiB
kernel tile (ms) pipeline (ms) GB/s % of copy
-------------------- ------------ -------------- ---------- -----------
hillis-steele 1.436 1.798 89.0 36.2
blelloch, plain 1.359 1.739 92.0 37.4
blelloch, padded 1.810 2.180 73.4 29.9
padded is 1.25x the plain tree's pipeline time
the copy row moves 76 MiB and a scan pipeline moves 153, so the last column compares rates and not times
Blelloch does five and a half times fewer adds. 8,190 per tile against Hillis-Steele's 45,057, for the same 4096-element scan. That is the difference between O(n) and O(n log n) written as a number you can check.
The padding cut the requests 5.8x and lost on time. Plain Blelloch
issues 1,534 shared memory requests per tile and the padded version issues
264, for the reason day 15 measured: the
up-sweep's spacing doubles at every level, so by stride 16 every active lane
is hitting the same bank, and gcd(2 * stride, 32) is 32.
Yet the padded tile runs 1.810 ms against the plain tree's 1.359, and the pipeline 2.180 against 1.739, 1.25x slower.
The request count is right, but the timing rejects the first guess. In this run, the tile scan is not bound by shared-memory bank conflicts, and the other two launches never touch shared memory. The padded kernel's extra index work costs more than the conflicts it removes.
Keep the request count as the model and let the clock test it. Day 34 does the same for tiling.
Notice that Hillis-Steele's request count, 1,536, is almost identical to plain Blelloch's. Doing less arithmetic did not by itself make the memory behaviour better, and a work-efficient algorithm with a conflicting access pattern can lose to a wasteful one with a clean pattern.
CUDA 13 narrowed the padding penalty without changing the result. On the same T4, the padded pipeline took 1.792 ms against 1.573 ms for the plain tree, 1.14x slower rather than CUDA 12.6's 1.25x. All three scans still matched the reference exactly.
The request model is unchanged. Padding still loses on elapsed time, but the size of that loss depends on the runtime.
Run it yourself
Use a CUDA GPU that supports the lesson's minimum compute capability. The README has the exact build line. The Colab setup is one way to get such a GPU.
nvcc -std=c++17 -O3 -arch=sm_75 -o scan_2 scan_2.cu
Compiler Explorer would take the kernels but not the harness. The widest tile here uses 32 KiB of shared memory. What does not fit is two 40 MB device buffers, a CPU scan of 10,000,000 elements for the reference, and seventy timed runs across seven measurements, far more than an embed's 20 second budget will take.
Exercise
The padded kernel offsets word i by i / 32 + i / 1024. Cut it to i / 32,
which is what most modern copies of the macro use. Before you build, read the
per-level table the program prints and write down which of the twelve levels
change and by how much.
Then run it and report both pipeline times.
Time: 25 to 40 minutes. Submit: the levels you predicted, the two pipeline times, and one sentence saying why the change is small.
Check: the harness re-runs the exact scan comparison on your build, so a padding function that broke an index fails with the element, the block and the position in the tile named, rather than winning by scanning less. It then reports both pipeline times whether they pass or not, because the ratio is the answer and the gate is only the gate.
Hint 1
The one-term offset and the two-term offset agree everywhere i is below
1,024. Which tree levels touch words above that, and how many lanes are still
active by the time you get there?
Hint 2
The program prints a warps column next to the degree columns. Multiply, do
not compare. A 4-way conflict on a level running one warp costs three extra
requests; a 2-way conflict on a level running 64 warps costs 64.
Solution
Six levels change. Strides 64, 128, 256 and 512 go from degree 1 to degree 4, and strides 32 and 1,024 go from 1 to 2. Every one of them runs a single warp except stride 32, which runs two.
That is 30 extra requests out of 264 for the pair of sweeps, and the tile scan is one part of a three-launch pipeline whose other two launches never touch shared memory. The measured difference should be small enough that run-to-run variation is a fair share of it.
Pitfalls
Your scan is right inside each tile and wrong across them. Every tile starts at zero, so the output is 2,442 correct scans laid end to end. The offset pass is not an optimisation, it is part of the algorithm.
Check element
kTileElems before you check element 0.
You clear the root before you save the tile total. The down-sweep needs a
zero at the root and the second kernel needs the sum that was there. Read it
out and write the zero in the same if (tid == 0) block, in that order, or
the block sums are all zero and every tile after the first is short.
You grow the array and forget one index. __shared__ float tile[kTileWords] with tile[b] in one place and tile[paddedWord(b)] in
another compiles, runs and reads a word nobody wrote. Pad every index or none;
the padded width belongs in one function every access goes through.
Day 15's version of this is the flat array indexed at the unpadded row stride.
You return early instead of guarding the loop. The deep levels leave most
threads with no add to make, and if (k >= adds) return; above the barrier
looks like the tidy way to say so.
It does not hang. A thread that returned has exited.
__syncthreads() waits only for the threads that have not, so
the barrier releases and the next level reads words nobody finished writing.
Guard the loop and leave the barrier alone;
day 14 has the Programming Guide's exact wording.
You believe a degree without counting the warps. A 32-way conflict on the level where one warp is still active is 31 wasted requests. The same degree at stride 16, where 4 warps are active, is 124.
Fix the shallow end of the tree first.
You profile as a normal user and get nothing.
Nsight Compute reports ERR_NVGPUCTRPERM, in
full: The user running <tool_name/application_name> does not have permission to access NVIDIA GPU Performance Counters or the Hardware Event System on the target device
(https://developer.nvidia.com/nvidia-development-tools-solutions-err_nvgpuctrperm-permission-issue-performance-counters
, checked 2026-08-30). On a machine where you have root, sudo ncu works,
because the stock driver default RmProfilingAdminOnly is 1.
Go deeper
- GPU Gems 3, chapter 39, sections 39.2.2 "A Work-Efficient Parallel Scan", 39.2.3 "Avoiding Bank Conflicts" and 39.2.4 "Arrays of Arbitrary Size": https://developer.nvidia.com/gpugems/gpugems3/part-vi-gpu-computing/chapter-39-parallel-prefix-sum-scan-cuda (checked 2026-08-30)
- CUDA C++ Best Practices Guide 10.2.3.1, "Shared Memory and Memory Banks", for the bank map and the serialization rule: https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html (checked 2026-08-30)
- Guy Blelloch, "Prefix Sums and Their Applications", CMU-CS-90-190, the paper the two sweeps come from: https://www.cs.cmu.edu/~guyb/papers/Ble93.pdf (checked 2026-08-30)
- Programming Massively Parallel Processors, 4th edition, chapter 11, on scan and its hierarchical form: https://shop.elsevier.com/books/programming-massively-parallel-processors/hwu/978-0-323-91231-0
Next
Day 33 uses a scan for the thing it was invented for: a scan over a vector of ones and zeros gives every surviving element its output address, which is stream compaction in two steps. Day 35 builds a radix sort the same way, one scan per pass of the digit.
Both keep day 31's simpler Hillis-Steele tile inside their pipelines rather than this page's tree. The timing table shows why: the pipeline's cost sits in the offset passes, not in the tile scan, so the work-efficient tile buys nothing there.
Both inherit the multi-level structure in step 2.