# Payoff: Offsets Place the Data

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

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

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

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

## Starter code

```js
// 32 sessions, 128 seats, one flat list. Which session owns each seat?
const gpu = new GPU({ mode });
const ITEMS = 32;
const SLOTS = 128;

// Tasks 3 and 4, prewired: counts -> inclusive scan -> starting offsets.
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: [ITEMS] });

const toExclusive = gpu.createKernel(function (inclusive) {
  if (this.thread.x === 0) {
    return 0;
  }
  return inclusive[this.thread.x - 1];
}, { output: [ITEMS] });

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

// Your kernel: one thread per SLOT.
const ownerOf = gpu.createKernel(function (offsets, counts) {
  // TODO: search the sessions for the one whose block contains this slot.
  // Session i owns slot s while offsets[i] <= s < offsets[i] + counts[i].
  return 0;
}, { output: [SLOTS], constants: { items: ITEMS } });

const owners = await ownerOf(offsets, counts);
console.log('slots 0-9 belong to sessions:',
  owners[0], owners[1], owners[2], owners[3], owners[4],
  owners[5], owners[6], owners[7], owners[8], owners[9]);
console.log('the last slot belongs to session:', owners[SLOTS - 1]);
```

---

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

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