GPU Challenges / CUDA

Parallel Histogram Computation: CUDA Exercise

medium Runs on a real NVIDIA T4 GPU in your browser

A histogram is the textbook example of a race condition on the GPU: many threads try to increment the same counter at once. This exercise teaches atomic operations, and shows why privatising bins in shared memory is the standard way to make them fast.

The problem

Compute a histogram of 256 bins for an array of 10,000 random byte values (0-255). Use atomic operations to avoid race conditions when multiple threads update the same bin.

Build a 256-bin histogram in parallel, using atomic updates so many threads can safely increment the same bin.

Your goal

  • Read one byte value per thread and map it to the matching histogram bin.
  • Use atomicAdd to avoid race conditions on shared output bins.
  • Keep the kernel bounds-safe for DATA_SIZE inputs.

How to approach it

  1. Assign each thread one or more input bytes using a grid-stride loop.
  2. Increment the matching bin with atomicAdd(&hist[value], 1) instead of hist[value]++.
  3. Optionally keep a per-block histogram in shared memory, then merge it into global memory once per block.
  4. Check that the bin counts add up to the number of input values.

Common mistakes

  • Using hist[value]++, which loses updates when two threads hit the same bin.
  • Not zeroing the histogram (and any shared-memory copy) before counting.
  • Merging a shared-memory histogram before every thread in the block has finished counting.

What success looks like

  • The GPU histogram should match the CPU histogram for the checked bins.
  • The histogram buffer must be zero-initialized before launching the kernel.
  • Each thread should process at most one valid input element in this starter version.

Starter code

Start from this scaffold. Fill in the TODOs, then run it against the hidden tests.

histogram.cu
1234567891011
#include <cuda_runtime.h>

// Kernel is launched as:  computeHistogram<<<ceil(size/256), 256>>>(data, histogram, size)
// gridDim.x  = ceil(size / 256)   blockDim.x = 256
// i = blockIdx.x * blockDim.x + threadIdx.x   (one thread per input byte)
__global__ void computeHistogram(const unsigned char* data, unsigned int* histogram, int size) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < size) {
        // TODO: Use atomicAdd to safely increment histogram[data[i]]
    }
}
Solve it in the editor →Free account, no GPU or CUDA install needed.

Concepts used

  • Atomics for correctness
  • High-contention write patterns
  • Comparing GPU results against a CPU reference

Learn them first:

More CUDA challenges