Telemetry 2: Preserve the Timeline Order

Keep the slow events in their original timeline order. The starter is correct but serial: one thread walks the full timeline while the rest of the block is idle. For [800, 4200, 6100, 900, 7300] with threshold 4000, the output must be [4200, 6100, 7300]. Replace the serial loop with a one-block exclusive prefix scan that assigns stable output positions without atomics.

Languages: CUDA C++, Rust, HIP C++

Lesson

The starter preserves order because one thread owns every read and write, but that leaves the other 255 threads idle. The optimization must keep the same ordered result while distributing the work across the block.

First convert each event into a keep flag: 1 means keep it, 0 means drop it. A kept event's stable output position is the number of 1s before it. That is an exclusive prefix sum.

Use a Hillis-Steele scan in shared memory. At each distance, read the neighbor into a register, wait, update scan[i], then wait again. The first barrier protects this step's reads; the second makes its writes visible to the next step. After the scan, scan[i - 1] is thread i's exclusive position. The last valid thread can derive the total count from its position plus its keep flag.

The stable-scan proof is limited to one block and N <= 256. Multi-block stable compaction needs block counts, an inter-block prefix stage and a separate scatter.

What you will learn

Replace a correct serial compaction loop with parallel work while preserving its stable output.
Turn 0/1 keep flags into stable output positions with an exclusive prefix sum.
Implement the read–barrier–write–barrier structure required by an in-place Hillis-Steele scan.
Convert the inclusive scan result into exclusive positions.
Derive the final output count without atomics or another kernel.

Starter code

    // Correct and stable, but serial: one thread performs every read,
    // comparison, and copy while the rest of the block stays idle.
    if (threadIdx.x != 0) {
        return;
    }

    unsigned int outputPosition = 0;

    for (unsigned int i = 0; i < eventCount; ++i) {
        KernelEvent event = events[i];

        if (event.durationNs >= thresholdNs) {
            slowEvents[outputPosition++] = event;
        }
    }

    *slowEventCount = outputPosition;