# form G — cuda-radix-sort — build progress

## 2026-06-05 — landed live

Thin wrapper around `cub::DeviceRadixSort::SortKeys` for u64 ascending
sort. Day-1 wire freezes one op (0x01 sort-u64-asc); reserves op slots
0x02 (desc), 0x03 (u32), 0x04 (key-value) in our dispatcher switch
without wiring them.

### Wire

- Request: `BSRT` (4 B) + u32 op_id + u32 n + u64[n]
- Response: `BSRR` (4 B) + u32 status + u32 n + u64[n] sorted ascending

Magics distinct from BSHK/BSHR/BCGB/BCGR/BSCP/BSCR.

### Validation (3090-ai.foxhop.net, RTX 3090)

Byte-identity vs Python `sorted()` PASS at n ∈ {32, 1k, 100k, 1M, 10M}.
Bench table from `--bench` mode:

| n         | kernel ms | Gkeys/s |
|-----------|-----------|---------|
| 100,000   |     0.142 |   0.706 |
| 1,000,000 |     0.265 |   3.767 |
| 10,000,000|     1.817 |   5.504 |

Reference catalog number: 1.4 Gkeys/s on Titan. 3090 lands ~4x that at
saturation, which matches CUB's published Ampere scaling. Kernel time
at small N reflects fixed launch overhead, not algorithm cost — Gkeys/s
only stabilizes past n ≈ 1M.

### 4090 host (ai.foxhop.net) — capacity constraint, not a form-G defect

`./radix-sort --binary` OOMs on cudaMalloc for any N when all four
gpu-worker daemons are co-resident: the secp256k1 daemon allocates
its full scratch footprint at startup, leaving 47 MiB free of the
4090's 24 GiB. This is a pre-existing constraint of the ai box's
daemon mix; radix-sort itself built clean & registered with the worker
log line `gpu-worker: ready cuda-radix-sort <- ./radix-sort`.

Standalone validation on the 4090 will run once the secp daemon
shrinks its preallocation OR we add a `RADIX_SORT_ONLY=1` worker mode.
Out of scope for form-G day-1; tracked here as a follow-up.

### Files

- `radix-sort.cu` — daemon + binary mode, mirroring shake256-fanout.cu
- `Makefile` — `radix-sort`, `radix-sort-test`, `radix-sort-bench` targets
- `test_radix_sort_known_answers.py` — sweep + bench harness
- `gpu-worker.lsp` — `handle-binary-sort` + `BSRT` magic dispatch
  + `maybe-register-daemon!` of `cuda-radix-sort`
