# The Whole Sort

*Task 5 of 6 · [Radix Sort](https://gpu.rocks/learn/radix-sort-fd3ff796.md) · GPU.js Learn*

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.

## Figures

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

## Goal

**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:

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

## Hint 2 — the driver

```js
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.

## Starter code

```js
// Four kernels, three passes. Only the destination rule is missing.
const gpu = new GPU({ mode });

const histogram = gpu.createKernel(function (keys, place) {
  let count = 0;
  for (let i = 0; i < this.constants.n; i++) {
    const d = Math.floor(keys[i] / place) % this.constants.radix;
    if (d === this.thread.x) {
      count++;
    }
  }
  return count;
}, { output: [16], constants: { n: 1024, radix: 16 } });

const offsets = gpu.createKernel(function (counts) {
  let start = 0;
  for (let b = 0; b < this.constants.radix; b++) {
    if (b < this.thread.x) {
      start += counts[b];
    }
  }
  return start;
}, { output: [16], constants: { radix: 16 } });

const destinations = gpu.createKernel(function (keys, place, starts) {
  // TODO: your digit is Math.floor(keys[this.thread.x] / place) % this.constants.radix.
  // Return starts[digit], plus how many EARLIER elements carry the same digit.
  return 0;
}, { output: [1024], constants: { n: 1024, radix: 16 } });

const gather = gpu.createKernel(function (keys, dest) {
  let value = 0;
  for (let i = 0; i < this.constants.n; i++) {
    if (dest[i] === this.thread.x) {
      value = keys[i];
    }
  }
  return value;
}, { output: [1024], constants: { n: 1024 } });

// gpu.js locks an argument's type on a kernel's first call, and every kernel
// here is fed another kernel's output — so the chain starts as a Float32Array.
let values = Float32Array.from(keys);

// TODO: three passes, least significant digit first — place = 1, then 16,
// then 256. Each pass: counts → starts → destinations → gather, and the
// gathered array becomes the next pass's input.

console.log('smallest:', values[0], '| middle:', values[512], '| largest:', values[1023]);
```

---

Interactive version: https://gpu.rocks/learn/radix-sort-fd3ff796/5

[Previous task](https://gpu.rocks/learn/radix-sort-fd3ff796/4.md) · [Next task](https://gpu.rocks/learn/radix-sort-fd3ff796/6.md)