Task 1 of 5

Rank by Counting

"Give me the ten largest of these four thousand scores." On a CPU you keep a heap of ten and walk the data once — and that plan does not port, because the heap's contents after element i depend on every element before it. Serial by construction.

So ask a question every element can answer alone: how many scores beat me? That count is the element's rank, rank 0 means nothing beats it, and anything with a rank below k is in the top k. No sorting, no shared state, one thread per element — each of them reading the whole array, which makes this O(n²) work and gloriously parallel.

Ties are where it bites. Two equal scores each counting the other come back with the same rank: two elements claim one slot, and the slot after it is claimed by nobody. The fix is a tie-break on the index — an element earlier in the array outranks you when the scores are equal, a later one does not. That turns the ranks into a permutation of 0…4095, exactly one element per slot. These scores repeat constantly, so you will feel it immediately.

a rank is a count, and a count is something every element can do alone — Eight scores in a row. One of them, highlighted, counts the scores that outrank it: two strictly larger scores count, and an equal score at a lower index counts, while an equal score at a higher index does not. The total, three, is its rank and its output slot.

Goal: return, for each element, the number of scores that outrank it — strictly larger anywhere, or equal at a lower index.

Requirements

One thread per score: output: [4096], loop bound this.constants.n
A strictly larger score always counts
An equal score counts only when its index is below this.thread.x
The largest score must come back with rank 0

Hint 1 — one loop, two comparisons

Split on the index, not on the value. For j < this.thread.x an equal score wins, so that side tests >=; for every other j an equal score loses, so that side tests >.

Hint 2 — the loop body

const other = scores[j];
if (j < this.thread.x) {
  if (other >= mine) ahead++;
} else if (other > mine) {
  ahead++;
}

Hint 3 — checking yourself

Every rank from 0 to 4095 should appear exactly once. If two elements share a rank, then somewhere a > is doing a >='s job (or the other way round).

Same idea elsewhere

Counting ranks is how a GPU sorts small things — it is the first sort in every CUDA and WebGPU tutorial, and the reason CUB's DeviceRadixSort and bitonic networks exist is that O(n²) stops being free somewhere above a few thousand elements. The index tie-break is what makes such a sort stable, the same guarantee thrust::stable_sort and std::stable_sort sell.

All tasks in Top-K Selection

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