Task 2 of 5

Gather the Winners

Ranks are the answer in an unusable shape: the ten you want are scattered somewhere among 4,096 slots. What you want is packed — top[0] the biggest score, top[9] the tenth biggest.

The obvious move is a scatter: element i writes itself into top[ranks[i]]. Kernels cannot do that — a thread writes one cell, its own. So turn it inside out, the way every scatter gets turned inside out. Instead of "where does my value go?", output slot j asks "who has rank j?" and goes looking. Ten threads, each scanning 4,096 ranks: a gather.

Exactly one element answers each slot — which is what last task's tie-break bought you. (Turning a rank array into a packed result is a pattern in its own right, and the Stream Compaction module develops it properly, with the prefix sum that makes it O(n) instead of O(k·n). You do not need that here: k is ten.)

Goal: fill ten output slots with the ten largest scores, largest first — slot j holds the score of the element whose rank is j.

Requirements

Hint 1 — which element is mine?

Thread j owns output slot j, and the element it wants is the one whose rank happens to be j. There is no way to know where that element sits, so look at all of them — a loop over the whole ranks array.

Hint 2 — the scan
let best = 0;
for (let i = 0; i < this.constants.n; i++) {
  if (ranks[i] === this.thread.x) best = scores[i];
}
return best;

No break needed — exactly one i matches.

Same idea elsewhere

Gather-by-rank is the back half of every GPU sort: compute a destination for each element, then have each destination fetch its element — thrust::gather, cub::DeviceRadixSort's final scatter pass, a WebGPU compute shader indexing a storage buffer. Production k-selection libraries (FAISS, RAFT's select_k) do exactly this once the candidates are down to a manageable few.

All tasks in Top-K Selection

  1. Rank by Counting
  2. Gather the Winners
  3. The Brightest Pixels
  4. Find the Cutoff Instead
  5. Which One Wins?

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