# Turn the Scatter Around

*Task 3 of 5 · [Stream Compaction](https://gpu.rocks/learn/stream-compaction-0aed2e43.md) · GPU.js Learn*

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:

```js
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.

## Figures

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

## Goal

**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:

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

## Hint 3 — the whole body

```js
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.

## Starter code

```js
// No scatter. Cell j goes looking for the element destined for j.
const gpu = new GPU({ mode });

const compact = gpu.createKernel(function (samples, flags, offsets) {
  // TODO: search the whole array for the element whose destination is
  // this thread's index — a survivor (flags[i] === 1) whose offsets[i]
  // equals this.thread.x — and return its sample.
  return samples[this.thread.x];
}, {
  output: [64],
  constants: { n: 64 },
});

const packed = await compact(samples, flags, offsets);
console.log('samples:', samples.slice(0, 12).join(', '));
console.log('packed: ', Array.from(packed).slice(0, 12).join(', '));
```

---

Interactive version: https://gpu.rocks/learn/stream-compaction-0aed2e43/3

[Previous task](https://gpu.rocks/learn/stream-compaction-0aed2e43/2.md) · [Next task](https://gpu.rocks/learn/stream-compaction-0aed2e43/4.md)