Task 4 of 6

Inclusive, Exclusive, and Why It Matters

Scans come in two flavours. The inclusive scan you just built answers "everything up to and including me". The exclusive scan answers "everything strictly before me": cell 0 is 0, and every other cell is the inclusive scan shifted one place right.

Exclusive is the one everything downstream actually wants, because it answers a different question — where does my run of output start? Here counts is a sign-up sheet: counts[i] people booked session i, and you are laying all of them out in one flat seating list. Session i's block begins at exclusive[i]. The inclusive scan would tell you where that block ends, which is exactly one seat too late.

Converting is a one-line gather: cell i reads inclusive[i − 1], and cell 0 returns 0 because it has nothing before it. One wrinkle worth knowing — an exclusive scan throws the grand total away. Its last cell holds everything except the last element, so keep the total separately: exclusive[n − 1] + counts[n − 1].

a zero goes in the front, the grand total drops off the back

Goal: turn the prewired inclusive scan into the exclusive scan — the starting offset of every session — and log the total number of seats.

Requirements

One kernel, taking the inclusive scan as its single argument
Cell 0 returns 0; cell i returns inclusive[i − 1]
Log the grand total, which the exclusive scan on its own no longer knows

Hint 1 — a shift is a gather

"Move everything one cell right" is a scatter, and kernels cannot scatter. Ask the inverted question instead — whose value lands in MY cell? — and it is a one-line read from this.thread.x - 1.

Hint 2 — the edge

Thread 0 must not read inclusive[-1]:

if (this.thread.x === 0) {
  return 0;
}
return inclusive[this.thread.x - 1];

Hint 3 — the total that got away

offsets[31] is where the LAST session starts, so the seat count is offsets[31] + counts[31]. (The inclusive scan's last cell had it all along — that is the reduction hiding inside every scan.)

Same idea elsewhere

Exclusive is the library default for exactly this reason: cub::DeviceScan::ExclusiveSum, thrust::exclusive_scan, WGSL's subgroupExclusiveAdd and Metal's simd_prefix_exclusive_sum all answer "where does my output begin?". And they all share the same wrinkle — CUB hands the aggregate back through a separate output, because the exclusive scan itself cannot carry it.

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.