# Inclusive, Exclusive, and Why It Matters

*Task 4 of 6 · [Prefix Sums (Scan)](https://gpu.rocks/learn/prefix-sum-351cfa41.md) · GPU.js Learn*

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

## Figures

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

## Goal

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

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

## Starter code

```js
// counts[i] people booked session i. Where does each session's block start?
const gpu = new GPU({ mode });
const N = 32;

// Last task's ladder, prewired: inclusive[i] = counts[0] + ... + counts[i].
const scanStep = gpu.createKernel(function (data, stride) {
  if (this.thread.x >= stride) {
    return data[this.thread.x] + data[this.thread.x - stride];
  }
  return data[this.thread.x];
}, { output: [N] });

let v = Float32Array.from(counts);
for (let stride = 1; stride < N; stride *= 2) {
  v = await scanStep(v, stride);
}
const inclusive = v;

const toExclusive = gpu.createKernel(function (inclusive) {
  // TODO: cell i should hold the total of everything BEFORE session i.
  // Cell 0 has nothing before it.
  return inclusive[this.thread.x];
}, { output: [N] });

const offsets = await toExclusive(inclusive);

console.log('counts [0..3]:', counts[0], counts[1], counts[2], counts[3]);
console.log('offsets[0..3]:', offsets[0], offsets[1], offsets[2], offsets[3]);
// TODO: the exclusive scan dropped the grand total. Log it.
console.log('total seats:', 0);
```

---

Interactive version: https://gpu.rocks/learn/prefix-sum-351cfa41/4

[Previous task](https://gpu.rocks/learn/prefix-sum-351cfa41/3.md) · [Next task](https://gpu.rocks/learn/prefix-sum-351cfa41/5.md)