Telemetry 1: Collect Slow Kernel Events

Keep every kernel event whose duration meets the threshold. The starter is correct, but every kept event reserves its output position through the same global counter — about 2.1 million contended atomics at the judged size. Change the reservation strategy so each block performs one global atomic without dropping or duplicating events. Output order does not matter in this lab.

Languages: CUDA C++, Rust

Lesson

The bottleneck is position reservation. Every kept event calls atomicAdd on the same four-byte global counter, so those reservations serialize even though the event copies are independent.

Privatize the work per block. Threads append survivors to a shared-memory list using a block-local counter. After the block finishes collecting, one thread reserves the entire output range with atomicAdd(slowEventCount, privateCount). The returned old value is the range start.

Every thread must reach every __syncthreads(), including threads past eventCount. Use guards instead of early returns. The result contains every kept event exactly once, but block reservations may arrive in any order; the next lab restores order with a scan.

Output order is intentionally unspecified. Check the selected event set and count, then compare reservation operations and elapsed time.

What you will learn

Recognize that reserving positions on one global counter creates the same contention as a global-atomic reduction.
Collect survivors with a shared-memory list and one block-local counter.
Reserve one contiguous global range per block using atomicAdd's return value.
Explain the trade-off: all events are preserved, but their order is not.

Starter code

    // Correct, and serialized: every qualifying event reserves its slot
    // on the same global counter.
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i >= eventCount) {
        return;
    }

    KernelEvent event = events[i];

    if (event.durationNs >= thresholdNs) {
        unsigned int position = atomicAdd(slowEventCount, 1);

        slowEvents[position] = event;
    }