100 DAYS OF GPU Course preview
Day 29Module 3
in-technical-review

CUDA histograms, privatized and measured

Here is a complete CUDA histogram, minus the loop around it.

atomicAdd(&bins[in[i]], 1u);

It is correct. Its run time depends on the byte values, not only on the byte count. With 64 MiB of English text it is slow.

With 64 MiB of one repeated byte it is much slower, although it reads the same number of bytes and issues the same number of instructions.

The source line does not show that cost. This page runs the kernel over three equal-size buffers with different contents. It then gives each block its own copy of the counters and repeats the test.

What one counter costs when the whole grid wants it

An atomic add is a read, a modify, and a write that no other thread can interrupt.

NVIDIA's guide says atomic functions "perform read-modify-write operations on shared data, making them appear to execute in a single step" (https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html section 5.4.5, checked 2026-08-30). Two threads cannot update the same word at the same time, so one must wait. That wait is contention.

There are 256 bins, one per byte value, in a 1 KiB array. This program sends 67,109,475 atomic updates to that array from 262,144 threads.

The array size is not the problem: 1 KiB is eight 128-byte lines, which L2 can hold. The reads are coalesced as 32 consecutive bytes per warp.

The cost depends on how the updates spread across the bins. Uniform bytes give each bin about one update in 256, while English text uses a few dozen bins most often. One repeated byte sends all 67,109,475 updates to one address.

Privatization reduces the number of threads that share each counter. Give each thread block its own copy of the 256 counters in shared memory, count into it, and merge it into the global array at the end.

Each private counter then receives updates from one 256-thread block instead of the whole grid. The atomic updates to global memory fall from 67,109,475 to 1024 blocks times 256 bins, or 262,144. That is 256 times fewer.

Nothing was removed. The per-byte atomic still happens, somewhere smaller and more private, which is why the pattern is called privatization and not elimination.

Privatizing 256 bins across 1,024 blocks preserves 67,109,475 shared atomic calls but reduces global atomic calls from 67,109,475 to 262,144.

Contention is not something the launch configuration fixes

The instinct is to reach for the launch configuration, and it is a good CPU instinct. Fewer threads will collide less.

More blocks will spread the work out. Day 10 taught you to sweep the block size, so sweep it.

It does not help. The input fixes the number of atomic updates: 67,109,475 bytes require 67,109,475 updates for any thread count. Halving the thread count gives each thread more input elements, so it does not reduce the total updates to a busy counter.

The address is the contended resource. An occupancy value does not describe contention at one address.

Two things do move it: how many threads share one counter, which is privatization, and how many counters exist, which is the bin count and has a ceiling you will meet.

When 256 bins becomes 8,192

256 bins of 4 bytes is 1 KiB, so here nothing binds. Ask how many bins would fit and the numbers get interesting.

On a Tesla T4 a block gets 48 KiB of shared memory by default and 64 KiB if it asks through cudaFuncSetAttribute, and the SM has 64 KiB in total (FACT-SHEET.md section 3, from Table 31 of https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html , checked 2026-08-29). The program prints all three off your own device rather than trusting a table.

The per-block limit gives 12,288 4-byte bins in 48 KiB. A block that uses all 48 KiB may be the only resident block on that SM. A lower limit keeps more blocks resident.

A T4 SM holds 32 resident warps, or four 256-thread blocks. Splitting 64 KiB among four blocks gives 16 KiB per block, which holds 4,096 bins. More bins reduce occupancy when each block uses a private copy.

There are three options when the bins do not fit. Narrow the counters: an unsigned short bin uses half the space and can hold at most 65,535.

Every block here sees just over 65,536 bytes, so the worst input would overflow it by one. A static_assert should enforce that limit.

Partition the bins: privatize one slice at a time and read the input once per slice, ignoring bytes outside that slice. The exercise below uses this method. Privatize into global memory instead: keep a few full copies, one per group of blocks, to avoid the shared-memory limit.

The knob turns the other way too. One copy per warp cuts the threads sharing a counter from 256 to 32 at eight times the shared memory: 8 KiB a block on 256 bins, 128 KiB on 4,096.

Measuring it without measuring the memset

Full program in code/day29-histogram/histogram.cu. It follows three rules, and the third is where this problem fights the site's own code style.

The floor is measured here, not borrowed. sumBytes reads every byte and does nothing else with it, from the same grid over the same buffer in the same program, so no correct histogram can beat it. Borrowing a copy bandwidth from day 11 or day 12 would fold those days' block shapes, buffer sizes and clock states into this answer.

__global__ void sumBytes(const unsigned char* __restrict__ in,
                         unsigned int* __restrict__ partials, size_t n) {
    const size_t step = gridDim.x * static_cast<size_t>(blockDim.x);
    const size_t t = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    unsigned int sum = 0u;
    for (size_t i = t; i < n; i += step) {
        sum += in[i];
    }
    partials[t] = sum;
}

Only the bytes change between rows. All three kernels read the same 67,109,475 bytes from the same grid with the same grid-stride loop, and neither histogram kernel learns which input it got.

__global__ void histogramGlobal(const unsigned char* __restrict__ in,
                                unsigned int* __restrict__ bins, size_t n) {
    const size_t step = gridDim.x * static_cast<size_t>(blockDim.x);
    for (size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
         i < n; i += step) {
        atomicAdd(&bins[in[i]], 1u);
    }
}

Correctness is checked before the timing, and the timing never zeroes. CUDA-CODE-STYLE.md says an in-place kernel needs a reset between timed runs, but that reset would become part of the measurement. A histogram must update its output buffer, so this program separates the check from the timing.

Before timing, it compares every bin against a CPU reference from a zeroed array. The timed loop does not zero or read its output. A static_assert proves that thirteen runs of the worst input fit in a 32-bit counter.

Here is the privatized kernel. The loops over the bins matter as much as the one over the bytes.

__global__ void histogramPrivate(const unsigned char* __restrict__ in,
                                 unsigned int* __restrict__ bins, size_t n) {
    __shared__ unsigned int privateBins[kBins];
    const int binStep = static_cast<int>(blockDim.x);

    for (int binIdx = static_cast<int>(threadIdx.x); binIdx < kBins;
         binIdx += binStep) {
        privateBins[binIdx] = 0u;
    }
    __syncthreads();

    const size_t step = gridDim.x * static_cast<size_t>(blockDim.x);
    for (size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
         i < n; i += step) {
        atomicAdd(&privateBins[in[i]], 1u);
    }
    __syncthreads();

    for (int binIdx = static_cast<int>(threadIdx.x); binIdx < kBins;
         binIdx += binStep) {
        atomicAdd(&bins[binIdx], privateBins[binIdx]);
    }
}

The merge is what people get wrong, and the version that reads more naturally has two bugs in it:

if (threadIdx.x < 256) {                        // wrong twice
    bins[threadIdx.x] = privateBins[threadIdx.x];
}

The assignment keeps only the values from the last block to update each bin. The literal 256 also works only when the block has at least 256 threads. With 128 threads, the upper half of the bins is not zeroed or merged.

Stride both loops by blockDim.x to handle any supported block size. The two __syncthreads() calls also matter: the first orders zeroing before counting, and the second orders counting before the merge.

Note. The GB/s column counts the input stream only. The bin traffic is 1 KiB against 64 MiB and the time goes on atomics rather than bytes, so the figure is there to make the histogram rows comparable to the floor row, not to describe the card's traffic. The private copy is also 256 words across 32 banks, so bins 32 apart share one; this program does not separate that from the atomic conflicts, and bank conflicts are day 15.

Results

Measured. Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -O3 -arch=sm_75. Captured 2026-08-30 on the project's verification node; full transcript in the page's evidence file.

GPU: Tesla T4 (compute capability 7.5)
Shared memory: 49152 bytes per block by default, 65536 with the opt-in,
               65536 per SM
Input: 67109475 bytes (64.00 MiB), 256 bins of 4 bytes
Grid: 1024 blocks x 256 threads. Mean of 10 runs after 3 warm-ups.

Atomic updates that reach global memory, per run
  histogramGlobal        67109475
  histogramPrivate         262144   (a factor of 256.0)

Reading every byte and doing nothing else with it
  sumBytes                0.631 ms     106.3 GB/s

input            max bin  bins used  global (ms)  private (ms)  speedup
--------------   -------  ---------  -----------  ------------  -------
uniform bytes      0.39%        256       21.703         0.759     28.6
english text      18.76%         34       20.464         0.783     26.1
all one byte     100.00%          1       55.400         2.386     23.2

Every row read the same 67109475 bytes from the same grid and the same
block, through the same two kernels. Only the contents of the buffer
changed between rows, so the spread down each column is contention.

CUDA 13.0 preserved the exact 256x reduction in global atomic updates and the 23x to 32x privatization speedups. The comparison reader changed: sumBytes moved from 0.631 ms and 106.3 GB/s to 0.943 ms and 71.2 GB/s. The privatized uniform row stayed at 0.760 ms.

In that run, the histogram was faster than this reader kernel. Treat sumBytes as a measured reference, not a fixed lower bound.

Privatization cut global atomic traffic by exactly 256x, from 67,109,475 updates to 262,144, which is one per bin per block. That number is not a measurement so much as arithmetic you can check: 1024 blocks times 256 bins. The speedup that follows is 23 to 29 times depending on the input.

Test an input with maximum contention. Uniform bytes spread updates across all 256 bins, and the global version takes 21.703 ms. English text puts 19 percent of its bytes in one bin and takes 20.464 ms.

When every byte is equal, every thread updates one counter. That input takes 55.400 ms, more than twice as long.

The privatized version degrades too on that input, 0.759 ms to 2.386 ms, because the contention moves into shared memory rather than disappearing. It is still 23 times faster than the global version, and the shape of the degradation is the same, which tells you privatization changes the constant and not the behaviour.

Compare all of it against sumBytes, which reads exactly the same bytes and does nothing else with them. On CUDA 12.6 it measured 0.631 ms, putting the 0.759 ms privatized histogram within 21 percent.

On CUDA 13.0 the reader itself measured 0.943 ms while the histogram remained at 0.760 ms. The reference is useful within one run, but the reversal shows why it is not a portable floor.

Run it yourself

A free Colab T4 or any card you own. The build line is in the repo's README:

nvcc -std=c++17 -O3 -arch=sm_75 -o histogram histogram.cu

There is no Compiler Explorer embed on this page. The program refills one 64 MiB host buffer three times, checks 256 bins six times against a CPU reference, and runs about a hundred launches. This exceeds Compiler Explorer's 20 second run limit.

Shrinking the buffer enough to meet the limit would also reduce the contention being measured. If you have no GPU, read /setup/learn-cuda-without-a-gpu.

Exercise

Add a fourth kernel, histogramPartitioned, that privatizes 64 bins at a time and makes four passes over the input, skipping any byte outside the slice a pass owns. Then say what those four passes cost you.

Time: 30 to 45 minutes. Submit: histogram.cu with the new kernel, your table with its three new rows, and one sentence on when you would ship it.

Check: the harness compares all 256 bins against the CPU reference on all three inputs. It prints the first mismatched bin and both counts.

It then checks that the bins sum to 67,109,475, which catches a merge written as an assignment.

The harness reports your kernel's time against histogramPrivate on every input whether the checks pass or not. A partitioned kernel that is faster than histogramPrivate has probably skipped input bytes.

Hint 1

You cannot shrink a counter below four bytes without risking an overflow, so do not shrink the counters. Shrink how many of them are present at once, then ask what that costs on the other side of the ledger.

Hint 2

Count the reads. A pass has to look at every byte to find out whether that byte belongs to it. How many bytes does one pass increment, and how many does it read?

Compare the second number against the floor row.

Solution

With P passes over N bytes you read P times N bytes and perform N shared atomics in total, spread across the passes. The atomic work does not grow. The reads do, by a factor of P.

So partitioning is a loss whenever the whole histogram already fits: four passes over 256 bins cost close to four times the input bandwidth and save nothing. Expect your rows near four times the histogramPrivate rows on the uniform and prose inputs, and rather better on the one-byte input, where three passes in four find nothing to increment and still have to read.

The diff against the shipped histogram.cu is one kernel, one extra argument giving the pass its bin range, and one loop in main. Nothing else moves: the private copy shrinks and everything around it stays put.

Privatization pays for global atomics with shared memory, and partitioning pays for shared memory with input bandwidth. Every version of this pattern is that trade priced differently. Spend the currency you have spare.

Pitfalls

Your counts come out ten times too big. You timed the kernel in a loop and the kernel accumulates into the buffer being timed. Check correctness once from a zeroed array before timing.

Do not read the bins produced by the timed loop. Zeroing inside the loop is the other option, but it puts a memset in your measurement.

You zeroed the private copy with if (threadIdx.x < 256). Correct at 256 threads a block, silently wrong at 128: bins 128 to 255 start on whatever the previous block left there, so the counts are wrong by an amount that moves with the schedule. Stride the loop by blockDim.x.

You merged with bins[b] = privateBins[b]. Every block overwrites the one before it, the answer is about one block's worth of counts, and it changes between runs. One line catches it and it needs no reference: the bins have to sum to the number of bytes.

You dropped the second __syncthreads(). The merge then reads a private copy that other warps in the same block are still adding to. It does not hang or crash.

The counts come out low, and the value changes between runs. Day 14 explains what a barrier does and does not promise.

You tuned on the wrong file. Uniform random bytes are the easiest case anyone will hand you, and they are what a synthetic benchmark generates by default. Measure the degenerate case too: a log file that is one repeated token is not a hypothetical.

You concluded that atomics are always the problem. With enough distinct bins and a flat distribution a global-atomic histogram can already sit near the read speed, and privatization then spends shared memory for nothing. The program prints the busiest bin's share beside every row so you can see which case you are in before optimising.

Go deeper

Next

Day 30 closes module 3 by putting five kernels from the first twenty-nine days on a roofline and asking which optimisation would move each, which is the question this page answered by measurement for exactly one of them. Day 42 comes back with the profiler counter that separates "the atomics got cheaper" from "there are fewer of them", the one thing this program cannot tell you.