Task 6 of 6

Payoff: Offsets Place the Data

Now the reason scan is the primitive everything else is built on. Each of the 32 sessions produces a variable number of output rows — counts[i] of them — and they all have to land in one flat 128-slot list, in order, with no gaps. The exclusive scan of counts is exactly the array of starting offsets, and it is prewired for you here out of tasks 3 and 4.

On a CPU you would loop the sessions and write each block — a scatter, which kernels cannot do. So invert it, the way a gather always inverts a scatter: one thread per output slot, each asking "which session owns me?". Slot s belongs to session i when offsets[i] <= s < offsets[i] + counts[i] — the session whose block has already started and has not yet run out. Six sessions here booked nobody; their blocks are empty, contain no slot at all, and drop out of the search on their own. An offset exists whether or not anything lands on it, which is exactly why the scan has to produce one for every session.

Count, scan, place. That is stream compaction, run-length decoding, sparse-matrix assembly, and every "each thread emits a different number of results" problem on a GPU — all of them a scan wearing a hat.

Goal: fill 128 slots, each one returning the index of the session that owns it.

Requirements

One thread per output slot — output: [128]
Search all 32 sessions with a compile-time loop bound (this.constants.items)
Slot s belongs to session i when offsets[i] <= s and s < offsets[i] + counts[i]

Hint 1 — invert the question

You cannot push a session's rows into the list. Ask the other question — whose row lands in MY slot? — and every slot searches the 32 sessions for the one whose block contains it.

Hint 2 — mind the first seat

Session i owns slot offsets[i] itself, so the lower test needs the equals sign: offsets[i] <= slot, not <. Get that wrong and every block's opening seat comes back ownerless.

Hint 3 — the loop

const slot = this.thread.x;
let found = 0;
for (let i = 0; i < this.constants.items; i++) {
  if (offsets[i] <= slot && slot < offsets[i] + counts[i]) {
    found = i;
  }
}
return found;

Same idea elsewhere

Count, scan, place is the standard three-kernel recipe for variable-sized output on every platform. thrust::copy_if and cub::DeviceSelect::Flagged are a scan of a 0/1 flag array with a gather bolted on; a WebGPU or Metal particle system whose sources each emit a different number of fragments uses the same scan to decide where each one writes; GPU sparse-matrix builders scan row lengths to get row pointers. Without a scan, none of it is parallel.

All tasks in Prefix Sums (Scan)

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