Task 5 of 5

Payoff: How Many Survived

Three kernels, wired together: flags, then the scan, then the gather. All of it is below, finished — the pipeline is yours already. What is missing is the number that makes the result usable.

The output is 64 cells long because the launch was 64 threads wide. Only the first count of them hold survivors; the rest hold whatever the search failed to find, and reading them as data is the classic way to ship a bug. So where does count come from? The end of the scan — but not from offsets[63] alone. An exclusive scan at the last index counts everyone before the last element, so the last element's own flag is still outstanding:

const count = offsets[n - 1] + flags[n - 1];

Drop the + flags[n - 1] and the pipeline quietly loses its final element, but only when that element happens to survive — which is exactly the kind of bug that passes every test you wrote by hand. And note what this number costs: it has to come back to JavaScript before anything can use it. That single readback is why compaction is the awkward step in an otherwise fully on-device pipeline.

Goal: compute the survivor count from the end of the scan, trim the packed output to it, and log both.

Requirements

count = offsets[63] + flags[63] — the last offset plus the last flag
Trim the 64-cell output down to those count values
console.log the count on its own line, and the kept values as a list

Hint 1 — the count lives at the end of the scan

offsets[63] is "how many survived among elements 0…62". Element 63 is not in that total — its flag is. Add them.

Hint 2 — trimming

The kernel returns a Float32Array; turn it into a plain array and cut it at the count:

const kept = Array.from(packed).slice(0, count);

Hint 3 — check yourself

kept.length should equal count, and every value in it should be at least 50. If the last one is missing, you dropped the + flags[63].

Same idea elsewhere

Every real compaction API hands the length back separately, and for the same reason: thrust::copy_if returns an end iterator, CUB's DeviceSelect::Flagged writes d_num_selected_out to device memory, and Vulkan/WebGPU pipelines that want to avoid the readback entirely feed that counter straight into an indirect dispatch or draw — the GPU deciding its own launch size from a number the CPU never sees.

All tasks in Stream Compaction

This page is an interactive exercise — the editor, the GPU runner and your saved progress need JavaScript. The text above is the full brief.