Telemetry 3: Build a Kernel-Latency Distribution

Count kernel durations into five latency bins. The starter is correct, but all 4,194,304 events update only five global counters, and a short-kernel-heavy workload funnels most updates into bin 0. Give each block a private shared-memory histogram, then merge each private bin into the global result once.

Languages: CUDA C++, Rust

Lesson

A histogram is five reductions running at once. The small number of destination addresses creates contention, and skew makes it worse by sending most events to the same bin.

Create five shared counters per block. Zero them cooperatively, synchronize, and let the grid-stride loop update those private bins. Synchronize again before the first five threads merge one bin each into global memory.

The harness tests evenly spread, short-kernel-heavy, and long-kernel-heavy inputs. Compare the profiles: the starter changes dramatically with skew, while the privatized version keeps almost all contention on-chip.

Shared histograms reduce global atomics but skewed values can still contend on shared bins. Compare uniform and skewed distributions with correctness and runtime.

What you will learn

Recognize a histogram as a set of contended reductions whose cost depends on input skew.
Accumulate into five block-private shared counters and merge once per bin.
Separate shared-state initialization, counting, and merging with barriers.
Use a grid-stride loop and compare behavior under different latency distributions.

Starter code

    // Correct, and contended: every event updates one of five global
    // counters directly. latencyBin() maps a duration to its bin.
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int stride = blockDim.x * gridDim.x;

    while (i < eventCount) {
        unsigned int bin = latencyBin(durationsNs[i]);

        atomicAdd(&latencyBins[bin], 1);

        i += stride;
    }