100 DAYS OF GPU Course preview
← Glossary
CUDA glossaryPerformance
CC 7.5

What is privatization in CUDA?

Giving each block its own copy of a contended structure in shared memory, then merging the copies once at the end.

Nothing is deleted. Count the atomic adds in a privatized histogram and there are more of them than before, because every byte still costs one and the merge adds a few hundred per block on top. What changes is who is queueing with whom. An atomic update to an address the whole grid can reach queues behind every other thread that wants it, and day 29 launches 1024 blocks of 256 threads, so that is 262,143 rivals. The same update into a copy only one block can see has at most 255. The updates that reach global memory fall from 67,109,475, one per byte, to 262,144, one per bin per block. That second figure is not a measurement, it is 1024 times 256, which is why the ratio comes out at exactly 256.0 rather than near it.

The structure has to earn the copy. Two conditions: it fits in shared memory, and its merge operator is associative and commutative, so the order the blocks finish in cannot change the answer. Counters, sums, minima and maxima all qualify. A linked list or anything order-dependent does not, which is why the pattern turns up under histograms and reductions and almost nowhere else.

The merge is where the bugs live, and both of them are quiet. Writing bins[b] = privateBins[b] instead of an atomicAdd gives you one block's counts, whichever block finished last, and the total changes between runs. Guarding the zeroing and the merge with if (threadIdx.x < 256) is right at 256 threads a block and silently wrong at 128: the upper half of the bins is never zeroed, so it starts on whatever the previous block left in that shared memory, and never merged either. Stride both loops by blockDim.x and neither bug is expressible. The two __syncthreads() calls are load-bearing for the same reason: one orders the zeroing against the counting, the other the counting against the merge.

What you pay is shared memory, and the bill is per block, so it comes out of occupancy. On a T4 a block gets 49152 bytes by default and 65536 with the cudaFuncSetAttribute opt-in, against 65536 per SM. At 256 bins of 4 bytes the copy is 1 KiB and nothing binds. At a few thousand bins the copy is the reason only one block fits on an SM, and asking for more than the per-block cap fails the launch with too many resources requested for launch. Past that point the moves are narrower counters, a slice of the bins per pass, or a handful of full copies in global memory instead of one per block.

Measured

On a Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -O3 -arch=sm_75, day 29 runs two histogram kernels over three buffers of the same 67,109,475 bytes, from the same 1024 blocks of 256 threads. Only the contents of the buffer change between rows, so every difference down a column is contention.

Input Busiest bin Bins used Global atomics One copy per block Speedup
uniform bytes 0.39% 256 21.703 ms 0.759 ms 28.6
english text 18.76% 34 20.464 ms 0.783 ms 26.1
all one byte 100.00% 1 55.400 ms 2.386 ms 23.2

Reading every byte and doing nothing with it takes 0.631 ms on this card, 106.3 GB/s, measured by a third kernel in the same program. The privatized histogram on uniform bytes lands at 0.759 ms, which is most of the way to doing no work at all.

Two things in that table are worth more than the headline ratio. English prose puts nearly a fifth of its mass on a single bin and the global version barely notices, 20.464 ms against 21.703, so ordinary skew is not what hurts. The degenerate input is, and there the private copy degrades too, 0.759 ms to 2.386, because the queue moved into shared memory rather than disappearing. It is still 23.2 times faster and the curve has the same shape, which is the honest summary of the pattern: privatization changes the constant, not the behaviour.

Diagram

Original SVG: three bands, each a column of blocks on the left and a 256-cell bin array on the right, with arrows drawn heavy where they cross into global memory and light where they stay on the SM. Band one sends every per-byte arrow to the one global array. Band two gives each block its own 256-cell array, lands every per-byte arrow there, and sends one bundle of 256 heavy arrows out at the end. Band three shows the private array too wide for the block's shared memory box, split into slices, with the input redrawn under each slice.

Alt text: "One atomic per byte puts sixty-seven million updates on one global counter array. A private copy per block cuts the updates that reach global memory to two hundred and sixty-two thousand, one per bin per block, and the per-byte atomics stay on the SM."

Code

From code/day29-histogram/histogram.cu. The counting loop and the merge, which is the whole pattern:

    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]);
    }

One atomicAdd became two, and the expensive one is now the rare one. The merge loop strides by blockDim.x, which is what makes it correct at any block size, and it is an add rather than a store, which is what makes it correct with more than one block.

Related terms

Where you meet this

Sources

Byline

Written by: pending. Reviewed by: pending. Written on: pending. Last checked on: pending. This entry stays a draft until a named author and a different named reviewer sign it.