Task 4 of 5

Find It in log n

The search you just wrote reads all 64 flags per thread. Across 64 threads that is 4,096 reads to move 30-odd values — worse than the CPU's single pass. It works, and it is the right shape, but it throws away the one thing that makes the array searchable: offsets never decreases.

Add the flag back to it and you get the running count — offsets[i] + flags[i], how many survived up to and including i. It is non-decreasing too, and it steps up by exactly one at each survivor. So the element for output cell j is the first index whose running count exceeds j, and a sorted array is something you can binary-search: seven halvings settle 64 elements instead of 64 reads.

This is a lower bound search — keep a window [lo, hi), look at its midpoint, and throw away the half that cannot contain the answer. When the window is empty, lo is the index you wanted.

the running count only ever goes up, so you can halve your way to it
Goal: replace the linear search with a binary search over the running count, and return the sample it lands on.

Requirements

Hint 1 — what you are searching for

For output cell j you want the smallest index whose running count is greater than j. Greater than, not equal to: the running count reaches j + 1 exactly at the survivor destined for slot j.

Hint 2 — one halving

If the midpoint's running count already exceeds this.thread.x, the answer is at mid or to its left, so hi = mid. Otherwise the answer is strictly to the right, so lo = mid + 1. Nothing else changes.

Hint 3 — the whole body
let lo = 0;
let hi = this.constants.n;
for (let s = 0; s < this.constants.steps; s++) {
  if (lo < hi) {
    const mid = Math.floor((lo + hi) / 2);
    if (offsets[mid] + flags[mid] > this.thread.x) {
      hi = mid;
    } else {
      lo = mid + 1;
    }
  }
}
return samples[Math.min(lo, this.constants.n - 1)];

The if (lo < hi) guard matters: the window can empty before the seventh halving, and a midpoint of an empty window is an out-of-bounds read.

Same idea elsewhere

Device-side binary search is a first-class primitive — thrust::lower_bound, CUB's DeviceSelect internals, and the merge-path partitioning that load-balances GPU merges and sparse-matrix products all lean on it. Production compactors often go one step further and build a scatter-address table instead: one pass writes each output slot's source index, a second gathers through it, trading a buffer for the search entirely. Same inversion, one more array.

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.