Task 2 of 5

Where Do I Land?

A survivor's output slot has a very short definition: how many survivors are in front of me. Element 9 with four survivors before it lands at index 4. That number, for every element at once, is the exclusive prefix sum — a scan — of the flags: cell i holds the total of flags 0 … i−1, not counting its own.

Scan is a whole subject of its own, and the Prefix Sums module derives the log-time ladder version properly. Sixty-four elements do not need it: every thread can simply walk the flags and count. That is n reads per thread — blunt, but perfectly parallel, and right now what matters is what the number means, not how fast you can get it.

Exclusive, not inclusive. Count yourself and every survivor lands one slot too far, with the first one shoved off the front of the array.

flags say who survives; the scan under them says where each one lands
Goal: for every index, return how many of the flags strictly before it are 1.

Requirements

Hint 1 — "strictly before"

The loop already visits every flag; it just needs to ignore the ones that are not in front of this thread. Guard the accumulation with if (i < this.thread.x) — note <, not <=.

Hint 2 — the loop body
if (i < this.thread.x) {
  seen += flags[i];
}
Hint 3 — check it by hand

For flags [1, 0, 1, 1] the destinations are [0, 1, 1, 2]: the survivor at index 0 goes to slot 0, the one at index 2 goes to slot 1, the one at index 3 goes to slot 2. Index 1 gets a number too — it just never uses it, because it does not survive.

Same idea elsewhere

Exclusive scan is one of the two or three primitives every GPU library is built on: thrust::exclusive_scan, CUB's DeviceScan::ExclusiveSum, rocPRIM's equivalent, and subgroupExclusiveAdd burned into WGSL and hardware. Blelloch's 1990 formulation gets it in O(n) work and O(log n) depth — the loop here is the honest, slow version of the same answer.

All tasks in Stream Compaction

  1. Filter Has No Kernel
  2. Where Do I Land?
  3. Turn the Scatter Around
  4. Find It in log n
  5. Payoff: How Many Survived

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