8 The Histogram pattern

This chapter begins with a simple but important lesson: when multiple threads update the same value, a plain increment can fail because it is really a read-modify-write sequence that can be interrupted. Using a shared counter as the example, the text shows how race conditions silently lose updates and why the result can vary from run to run. Atomic operations solve this by making the update indivisible, so threads cannot overwrite one another’s work.

From there, the chapter generalizes the idea to histograms, where many threads must increment shared bins based on input values. A direct parallel version using ordinary increments is incorrect, so the first working solution replaces each bin update with an atomic add in global memory. This is safe, but it can become very slow when many threads target the same bin, because atomics to the same address serialize and create heavy contention.

To speed things up, the chapter introduces privatization: each block builds its own private histogram in shared memory, zeroes it cooperatively, tallies into it with shared-memory atomics, and then merges the local bins into the global histogram. This reduces contention dramatically because only block-level results reach global memory, and shared-memory atomics are much cheaper than global ones. The chapter also shows how grid-stride loops let each thread process multiple elements, and it closes by emphasizing the main design rule: protect shared writes with atomics, reduce contention with privatization, and use barriers carefully so every thread reaches them.

The race condition (Scenario A, top) and the correct sequential execution (Scenario B, bottom). When T0 and T1 run concurrently (A), both threads read count=5 and both write 6, so one increment is lost: the final count is 6. When T0 completes before T1 starts (B), the final count is correctly 7.

The histogram pattern on a CPU. A single thread T0 reads each element of the data array (length n) one at a time and writes to the corresponding bin in the 256-element hist array. With only one writer, there are no conflicts.

The histogram pattern in parallel, without atomics. n threads (T0 through Tn-1) read their data elements simultaneously and all write into hist at the same time. When multiple threads target the same bin, their writes race and updates are lost: the same race condition as in figure 8.1, scaled up across all bins.

The three-step shared memory histogram pipeline, with time flowing top to bottom. Step 1 (setup): each of the m blocks zeros its own 256-bin localHist in shared memory in parallel. Step 2 (compute): each block i reads its k-element slice of the input from global memory and atomically increments the matching bins of its own localHist. Step 3 (merge): each block adds the values in its localHist into the single global hist array. Only Step 2 reads global input and only Step 3 writes global output; the heavy traffic in between stays in fast shared memory.

Step 1 (setup) inside a single block. Block i has k threads, T0 through Tk−1. Thread Tj writes 0 into localHist entries at offset j and at every j + k, j + 2k, ... up to index 255, the stride-k pattern that guarantees all 256 bins are zeroed regardless of block size. Every block from 0 through m−1 does this independently in parallel.

Step 2 (compute) inside a single block. Block i is responsible for the slice data[i*k .. i*k+k−1] of the global input array. Thread T0 reads data[i*k], T1 reads data[i*k+1], and so on through Tk−1 reading data[i*k+k−1]. Each thread then atomically increments the localHist bin matching its element. The arrows converge because different threads in the block may target the same bin, but the atomicAdd lives in fast shared memory, not global.

Step 3 (merge) inside a single block. Block i’s k threads divide up its localHist using the same stride-k pattern as Step 1: thread Tj reads localHist[j], localHist[j+k], localHist[j+2k], ... and adds each value into the corresponding entry of the global hist with atomicAdd. All m blocks merge into the same hist concurrently, so the atomic is still required, but only m blocks contend per bin, instead of the n input elements that would have contended in the naive version.

• Atomic operations make read-modify-write indivisible, eliminating race conditions. atomicAdd, atomicSub, atomicMax, atomicMin, atomicExch, and atomicCAS make up your atomic toolkit. atomicCAS is the general one: it writes your value only if the current value is still what you expected, so it belongs inside a retry loop.
• atomicAdd takes an address, not a value: pass count when the counter is already a pointer, and &hist[bin] when the target is an array element. Zero the counter or the histogram on the device before you launch, or the count starts from whatever was already sitting in GPU memory.
• Atomics are safe, but they come with a cost. They serialize the threads that target the same address, so the cost tracks contention: values spread evenly across bins run at full speed, while one hot bin drops throughput to nearly sequential.
• The histogram pattern counts how often each value appears. On a GPU the fast version is privatized: each block accumulates into a private shared memory copy, then merges it into the global histogram with atomics.
• The privatized kernel runs in three steps, with a __syncthreads() after each of the first two: zero localHist in shared memory, tally into it with shared memory atomics, then merge it into the global histogram. Both barriers must sit where every thread reaches them.
• Instead of one global atomic per input element, each block issues at most one per bin, carrying that block’s accumulated count. The reduction factor is the block size, n / numBlocks, which equals 256 in our example only by coincidence with the 256 bins.
• “Reduce locally, combine globally” is privatization, the same shape as the per-block partial sums in chapter 5, and it answers contention outside CUDA too: OpenMP gives each CPU thread its own bins and sums them at the end.
• The histogram is the pattern where contention, not depth, sets performance. Its ideal depth is O(1), one independent element per thread, but atomics on a popular bin serialize threads and stretch it in practice.
• A grid-stride loop handles more elements than threads: each thread steps by blockDim.x * gridDim.x. Combined with privatization, fewer blocks each doing more work, it is a common design for production histogram kernels.
• In PyTorch, torch.bincount() and torch.histc() compute histograms; on a CUDA tensor they use optimized kernels built on the same ideas, privatization and atomics.

FAQ