Task 3 of 5

Turn the Scatter Around

You now have the two arrays a compaction needs: flags, and offsets — the exclusive scan that tells every survivor its slot. The obvious next line is the one you cannot write:

if (flags[i] === 1) out[offsets[i]] = samples[i];   // ✗ no scatter here

Thinking in Parallel spends a whole task on why: a thread writes one cell, its own, by returning a value. So turn the question inside out, exactly as it does. Not "where does my value go?" but "whose value lands in my cell?" — and output cell j can answer that itself. It goes looking for the index that is (a) a survivor and (b) carrying destination j. One thread, one pass over the flags, one value pulled home.

The output array is still 64 long, because that is what a kernel launch gives you. Only the first few cells will hold survivors; the rest hold whatever each thread's search failed to find. That is fine, and normal — you just have to know where the real data stops, which is the last task of this module.

same arrows, opposite owner — and only one of the two is legal

Goal: fill each output cell by searching for the element whose destination is this cell's index.

Requirements

Loop over all this.constants.n elements — the only write is the return
Take samples[i] when flags[i] is 1 and offsets[i] is this.thread.x
The survivors come out packed, in input order, starting at cell 0

Hint 1 — ask the other question

Thread j is not trying to place anything. It is trying to find something: the one index whose destination happens to be j. Keep a local value, overwrite it when the search hits, return it.

Hint 2 — both halves of the condition

Non-survivors have an offsets entry too — it just does not belong to them. So matching the offset alone is not enough; the flag has to be checked as well:

if (flags[i] === 1 && offsets[i] === this.thread.x) {
  value = samples[i];
}

Hint 3 — the whole body

let value = 0;
for (let i = 0; i < this.constants.n; i++) {
  if (flags[i] === 1 && offsets[i] === this.thread.x) {
    value = samples[i];
  }
}
return value;

Same idea elsewhere

"Turn the scatter into a gather" is the phrase GPU folklore compresses this into, and it is exactly how the libraries do it: thrust::copy_if and CUB's DeviceSelect both run a flag pass, a scan, and then a data movement driven by the scan. Compute APIs can scatter — CUDA and WebGPU threads may store anywhere — but two threads aiming at one address is a race, and the fix (atomics) serialises them. Scan-then-gather costs no atomics at all.

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.