Task 5 of 6

The Whole Sort

Assemble it. 1,024 keys, all below 4,096 — three hex digits, so three passes. Each pass is the four kernels you have already written: histogram the digit, scan the counts into starting offsets, compute every element's destination, gather. The gathered array is the next pass's input.

Only one piece is left: the destination rule at radix 16. It is task 1's stable rank with a bucket offset in front of it — starts[digit] puts you at the head of your bucket, and counting the earlier elements that share your digit places you inside it. Stability is still the whole game, and now you can see why: the last pass sorts by the most significant digit, and everything the first two passes achieved survives only inside its ties.

Two things the driver must get right, both of them silent when they are wrong: the passes go least significant digit first, and the histogram and offsets are recomputed every pass — each one looks at a different digit of a differently ordered array.

low digit first, every pass stable — shown in base 10; the code counts in base 16

Goal: finish the destinations kernel and drive three passes over keys, then log the smallest, middle and largest of the result.

Requirements

destinations returns starts[digit] + how many earlier elements share that digit
Three passes with place = 1, then 16, then 256
Recompute the histogram and the offsets inside the loop — once per pass
console.log the sorted array's first, middle and last values

Hint 1 — the destination rule

Two halves. starts[mine] is where your bucket begins; the loop counts your rank inside it, exactly as in task 1 but restricted to your own digit:

if (d === mine && j < this.thread.x) rank++;

Hint 2 — the driver

for (let place = 1; place <= 256; place *= 16) {
  const counts = await histogram(values, place);
  const starts = await offsets(counts);
  const dest = await destinations(values, place, starts);
  values = await gather(values, dest);
}

Three iterations, and every line of it inside the loop.

Hint 3 — why low digit first

Each pass makes its own digit the primary sort key and demotes everything the previous passes did to a tie-break. So the digit you want to dominate — the most significant one — has to be sorted last. Run the passes the other way and the array comes out ordered by its ones digit.

Same idea elsewhere

This is the shape of the real thing. cub::DeviceRadixSort, thrust::sort on integers, rocPRIM, and the Vulkan/WebGPU sort libraries all run this loop: per-pass digit histogram, scan to global offsets, stable scatter, repeat for as many digits as the key has. They beat this version on the two lines you did not write — the rank inside a bucket comes from a parallel scan instead of an O(n) sweep, and the move is a scatter into shared memory rather than a search — but the algorithm on the page is the algorithm they run.

All tasks in Radix Sort

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