chapter six

6 The Scan pattern: Hillis-Steele

 

This chapter covers

  • Inclusive versus exclusive scan
  • Comparing parallel algorithms
  • The Hillis-Steele algorithm
  • Using double buffering to avoid collisions

Imagine you’re checking your bank statement. You start the month with $500. Then you spend $20 on coffee, receive a $1,000 paycheck, spend $60 on groceries, and so on. The bank shows you a running balance on every single line.

Now imagine you’re the bank, and you have a million transactions to process. You could compute each running balance one at a time: transaction 1, then 2, then 3, all the way to a million. That works, but it’s painfully slow. What you need is to split the work across thousands of helpers and still get every running balance correct.

That’s a scan, also called a prefix sum.

In chapter 5, you learned the Reduction pattern: combining a million values into one result. Instead of throwing away the intermediate results, scan keeps every single one. Because of this property, scan lies behind sorting, filtering, memory allocation, and many other GPU algorithms; chapter 7 shows how it unlocks the first three. Learn this one pattern well and a large part of GPU programming falls into place.

Understanding the Scan pattern

The Hillis-Steele algorithm

Round 0: Starting state

Round 1: stride = 1

Round 2: stride = 2

Round 3: stride = 4

Why does this work?

Programming the Hillis-Steele algorithm with CUDA

Part 1: Setup with shared memory and the double buffer

Part 2: Load input into shared memory

Part 3: The main loop, doubling stride

Part 4: The parallel add with boundary check

Part 5: Write the result to global memory

Complete Hillis-Steele inclusive scan

Common mistakes

Mistake 1: Reading and writing the same array

Mistake 2: Missing boundary check

Mistake 3: Forgetting __syncthreads()

Summary