chapter eight

8 The Histogram pattern

 

This chapter covers

  • Race conditions and how atomic operations fix them
  • The CUDA atomic toolkit
  • Contention: why atomics are safe but not free
  • Building a histogram with global atomics
  • A faster histogram with privatization: giving every block its own shared memory copy, then merging the copies
  • Processing more elements than you have threads with grid-stride loops

Imagine a busy classroom. Two students are grading exams and updating the class total on a shared whiteboard. The process is simple: check the current total, add the exams you just graded, write in the new total.

Student A checks the whiteboard: the total is 50. She graded 3 more, so she writes 53. But at the same moment, Student B also checks the whiteboard, sees the same 50, adds the 2 he graded, and writes 52. Together they graded 5 exams, so the total should be 55. Instead, it reads 52. The count is now off by 3.

The classroom is this chapter’s problem in miniature, the race condition we met in chapter 5: multiple threads write to the same memory and overwrite each other’s work. It can show up whenever parallel threads count, sum, or tally into a shared result, something GPU programs do all the time.

The conflict problem

One counter, a thousand threads

Making read-modify-write indivisible

The CUDA atomic toolkit

From one counter to 256 bins

Why parallel histograms are hard

A first fix: global atomics

A faster fix: privatization in shared memory

Common Mistakes

Mistake 1: Forgetting atomics

Mistake 2: Conditional __syncthreads() (the deadlock trap)

Summary