Task 2 of 5
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.)
j holds the score of the element whose rank is j.output: [10] — one thread per result slotthis.constants.n ranks for the one equal to this.thread.xtop[0] is the largest score and top[9] the tenth largestThread 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.
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.
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.
This page is an interactive exercise — the editor, the GPU runner and your saved progress need JavaScript. The text above is the full brief.