Atomic Histograms

The Challenge of Parallel Histograms

A histogram counts the number of data elements that fall into a set of discrete categories, or "bins." For example, we might count the frequency of each letter in a document or the number of pixels of each intensity value in an image.

Sequentially, this is trivial: you loop through the data and increment the counter for the corresponding bin.

// Sequential histogram for characters
void histogram_sequential(char* data, unsigned int* hist, unsigned int length) {
    for (unsigned int i = 0; i < length; i++) {
        // Assuming lowercase letters 'a' through 'z'
        hist[data[i] - 'a']++;
    }
}

The expression hist[...]++ looks like a single operation, but for the processor, it's a three-step read-modify-write sequence:

1. Read the current value of the bin's counter from memory.
2. Modify the value by incrementing it in a register.
3. Write the new value back to memory.

When we parallelize this naively—assigning one thread to each data element—this three-step sequence becomes a major problem. If two threads, Thread A and Thread B, are both processing the letter 'm', a race condition can occur:

1. Thread A reads the value of the 'm' bin (e.g., 100).
2. Simultaneously, Thread B reads the same value of the 'm' bin (still 100).
3. Thread A increments its value to 101.
4. Thread B increments its value to 101.
5. Thread A writes 101 back to the bin.
6. Thread B writes 101 back to the bin, overwriting Thread A's result.

The correct final count should be 102, but because of the interleaved execution, one of the updates is lost. This is a classic race condition.

To solve this, we need a way to make the read-modify-write sequence an indivisible, or atomic, operation.

Atomic Operations: Guaranteed Integrity

Atomic operations are special hardware instructions that perform a read-modify-write cycle on a memory location as a single, uninterruptible transaction. When one thread executes an atomic operation on a memory address, the hardware guarantees that no other thread can access that same address until the atomic operation is complete. This serializes access to the memory location, preventing race conditions and ensuring correctness.

Let's watch a short video that demonstrates both the problem of naive parallel updates and the solution using atomics.

As the video showed, CUDA provides a family of intrinsic functions for atomic operations, such as atomicAdd(), atomicMin(), and atomicCAS() (Compare-And-Swap). For our histogram, atomicAdd() is the perfect tool.

The following reading provides a concise summary of the problem and shows the code for a basic atomic histogram kernel.

The kernel you just saw uses atomicAdd on global memory for every single element.

__global__ void histogram_naive_atomic(unsigned char* data, unsigned int* histogram, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        unsigned char bin = data[idx];
        // PROBLEM: Global atomic on every element
        // This serializes threads hitting the same bin
        atomicAdd(&histogram[bin], 1);
    }
}

This approach is correct, but it can be extremely slow. Global memory atomics have high latency, and if the input data has "hot spots" (many elements mapping to the same bin), threads will form a long queue waiting to access that single memory location, completely negating the benefits of parallelism. As someone interested in high-performance systems, you can see this is an unacceptable bottleneck.

Optimization via Privatization

The bottleneck in our naive atomic histogram is contention on global memory. The solution is to reduce this contention. We can achieve this using a powerful pattern called privatization.

Instead of having thousands of threads across the entire grid competing for bins in a single global histogram, we can create a private histogram for each thread block.

1. Local Histograms: Each thread block computes a histogram for only the subset of data it is assigned to. Threads within a block will still contend for bins in their block's private histogram, but there is no contention between blocks.
2. Merge Results: After all blocks have finished, their private local histograms are merged (summed) into the final global histogram.

The best place to store these private histograms is in shared memory. Shared memory is on-chip, offering much higher bandwidth and lower latency than global memory. Atomic operations on shared memory are an order of magnitude faster than on global memory. This significantly reduces the cost of contention within a block. The final merge step still requires atomics on global memory, but we've reduced the number of expensive global atomics from one per element to one per bin, per block.

Here is what the optimized kernel looks like.

The shared memory privatization pattern is a cornerstone of high-performance CUDA development. It follows a general principle we've seen before: bring data into faster, on-chip memory (registers or shared memory), operate on it there, and only write the final results back to global memory when necessary. By moving the high-frequency atomic updates from global memory to shared memory, we drastically reduce the serialization bottleneck and allow the GPU's parallel architecture to shine.

While this technique is already very powerful, further optimizations like warp-aggregated atomics exist, which you can explore in the "Warp-Aggregated Atomics" section of the RightNow AI article. These techniques apply the warp-level communication principles from our previous lesson to further reduce shared memory contention.

Conclusion

In this lesson, we confronted the challenge of race conditions in parallel algorithms by using atomic operations. We saw how these hardware-guaranteed instructions ensure correctness but can introduce performance bottlenecks if used carelessly on global memory.

Key Takeaways:

• Output Interference: Naive parallel algorithms that perform read-modify-write operations on a shared output can suffer from race conditions, leading to incorrect results.
• Atomic Operations: Primitives like atomicAdd() guarantee that a read-modify-write sequence is indivisible, serializing access to a specific memory location and ensuring correctness in the face of contention.
• Performance Cost: Atomics, especially on global memory, can be slow. Heavy contention on a few memory locations can serialize execution and eliminate the benefits of parallelism.
• Privatization: A key optimization pattern to reduce contention is to create private copies of the output structure (e.g., a histogram) for each thread block. Using fast shared memory for these private copies dramatically improves performance by localizing most atomic operations to the on-chip memory.

We have now explored two fundamental parallel primitives: reduction and histogram. Each has revealed critical aspects of GPU programming—efficient communication with warp shuffles and safe contention management with atomics. In our next lesson, we will investigate a third essential primitive: the parallel scan. This algorithm, also known as prefix sum, forms the building block for many more complex parallel operations, including sorting, stream compaction, and even solvers for recurrence relations.