# Payoff: How Many Survived

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

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:

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

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

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

## Starter code

```js
// The finished pipeline: flags → scan → gather. Only the count is missing.
const gpu = new GPU({ mode });

const flag = gpu.createKernel(function (samples) {
  return samples[this.thread.x] >= this.constants.threshold ? 1 : 0;
}, { output: [64], constants: { threshold: 50 } });

const destination = gpu.createKernel(function (flags) {
  let seen = 0;
  for (let i = 0; i < this.constants.n; i++) {
    if (i < this.thread.x) {
      seen += flags[i];
    }
  }
  return seen;
}, { output: [64], constants: { n: 64 } });

// The linear search from task 3 — task 4's binary search drops straight in.
const compact = gpu.createKernel(function (samples, flags, offsets) {
  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;
}, { output: [64], constants: { n: 64 } });

const flags = await flag(samples);
const offsets = await destination(flags);
const packed = await compact(samples, flags, offsets);

// TODO: the scan stops one short — offsets[63] counts everyone BEFORE
// element 63, so element 63's own flag is still missing from the total.
const count = offsets[63];

// TODO: keep only the cells that actually hold survivors.
const kept = Array.from(packed);

console.log('survivors:', count);
console.log('kept:', kept.join(', '));
```

---

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

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