Telemetry 4: Filter the Events In Place

Compact the kept events into the front of the same array, preserve their order, and report the new count. The starter is correct but serial: one thread owns the read/write order, so it cannot overwrite an unread event. Replace it with a block-parallel scan while preserving that safety rule. For [800, 4200, 6100, 900, 7300] with threshold 4000, the valid prefix becomes [4200, 6100, 7300] and filteredCount is 3.

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

Lesson

The starter is safe because one thread owns the entire sequence: it reads the current event before writing to an earlier position. That implicit ordering is also why the starter cannot use the rest of the block.

Parallel in-place compaction is safe only if no write overwrites an element that another thread still needs to read. Each thread first loads its own event into a register. A block-wide barrier then guarantees that every original element has been read. After the scan computes destinations, kept threads write their register copies toward the front of the array.

The event data does not need a shared-memory staging array because threads never write values read by other threads. Shared memory holds only the scan. This proof is block-scoped; __syncthreads() cannot coordinate an in-place compaction across multiple blocks.

The read-before-write argument holds within one block, with a barrier before overwriting inputs. General multi-block compaction requires separate global coordination.

What you will learn

Replace a correct serial in-place loop with block-parallel work without changing its output.
State the parallel safety rule: every original input must be read before compacting writes begin.
Use registers and a block-wide barrier to separate the read and write phases.
Explain why the scan needs shared memory but the event data does not.
Preserve the same ordered output without allocating a second global array.

Starter code

    // Correct and in-place, but serial: one thread owns the read/write
    // order, so writing an earlier slot cannot race another reader.
    if (threadIdx.x != 0) {
        return;
    }

    unsigned int writePosition = 0;

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

        if (event.durationNs >= thresholdNs) {
            events[writePosition++] = event;
        }
    }

    *filteredCount = writePosition;