DeroStorm 1.6.2 - 202.87 KH/s mining, GPU +15.8% across the 1.6 line

202.87 KH/s mining, 170.39 of it on an RTX 5080. Two changes to how the GPU
suffix sort stores and moves bytes, +15.8% on the card against 1.6.0, and the
first real-path measurement since 1.5.6.

No developer fee. Every hash pays the address you configure — unchanged from
1.6.0 and still checkable in one grep.

Where it lands

Read off the miner's own dashboard, mining against a node, on a Ryzen 7 9800X3D
(8C/16T, DDR5-6000 CL30) and an RTX 5080:

 ┌─◈─MINING PERFORMANCE──────────────┐
 │ 202.87 KH/s                MINING │
 │ CPU  32.47 KH/s   ██░░░░░░░░  16% │
 │ GPU 170.39 KH/s   ████████░░  84% │
 └───────────────────────────────────┘

	H/s
RTX 5080 at 1.5.8 / 1.6.0	142,400
at 1.6.1	160,330	+12.6%
at 1.6.2, --bench	164,950	+2.9%
		+15.8% across the 1.6 line
at 1.6.2, real mining path	170,390
CPU + GPU, real mining path	202,870	+27.4% over 1.5.6

--bench and the mining path disagree by about 3% and the mining path is the
higher one, which is the opposite of the usual direction: the benchmark starts
cold and mining runs warm and steady. Both are given above rather than one being
presented as an estimate of the other.

1.6.1 — the arena stops moving bytes it does not need to, +10.2%

The descriptor sort's arena is one entry per suffix, read back by the scatter
that writes the suffix array. Two changes, neither touching the descriptor
word's width, the descriptor count, or the number of radix passes.

An entry is a block index, not a position. Every position one descriptor
emits sits at the same column within its 256-byte block — that is what a
descriptor is. So the position (block << 8) | column splits into a 16-bit
block index that differs per entry and an 8-bit column that does not. The column
moved into the descriptor's fourth key byte, free since 1.5.8 stopped sorting it,
and the entry narrowed from 32 bits to 16. +3.2%.

And a column whose order has not moved writes nothing at all. A constant
column prepends the same byte to every suffix in the run, so the order holds —
the premise the whole file rests on, and the walk already knew which columns
those are. It wrote the order out again for every one of them anyway. Since a
column writes exactly ord[0..len) into one contiguous slot in order, and a slot
is never written twice, a column can point its descriptors at an earlier slot and
skip the store entirely. +10.2%, three interleaved rounds with no overlap:
145.59 / 145.57 / 145.64 KH/s against 160.61 / 160.24 / 160.29.

1.6.2 — the radix sort, re-swept at the width the library ships, +3.3%

Every radix knob on record had been swept at the wrong block width. The test
binary builds at 1,024 threads; the shipped library builds at 256. The block's
shared memory is dominated by two arrays that scale with thread count, so a wider
digit costs four times as much at 1,024 as at 256 and the trade sits somewhere
else entirely.

Re-swept at 256, three interleaved rounds each:

digit bits	passes over the 24-bit key	GPU
5	5	160.95 / 160.97 / 160.94 KH/s
6	4	163.77 / 163.79 / 163.71 KH/s
7 — what shipped	4	160.44 / 160.00 / 160.56 KH/s
8	3	142.95 / 142.75 / 142.76 KH/s

+2.1% for one #define. And 8 is the interesting loss: it orders the key in
three passes instead of four, which the file's own note called the only lever
left, and it costs 12%. Occupancy beats pass count. BR_BLOCK was re-checked as
the knob that invalidated the others and 256 is a real peak.

Six barriers a tile became three. Between ranking a tile and staging it, the
sort ran two scans over the bins and spent four __syncthreads() on a prefix sum
over 64 numbers. One warp does both now, carrying the scan in shuffles, which
need no barrier. +1.1%, no overlap.

Shared memory per block falls from 11.0 KB to 6.8 KB along the way.

Verified

Every optimisation is a faster route to the same 32 bytes, and that is checked
rather than asserted:

• gpu/desc_test.exe — correct on all 512 reference suffix arrays
• gpu/hash_parallel_test.exe — all 512 hashes match the CPU exactly
• go test ./cmd/derostorm/ — green
• the miner's own start-up check — matches the CPU exactly, per device, before
  it will submit anything

All four embedded libraries were rebuilt for this release, so no archive carries
a stale kernel: the build stops if an embedded library is older than the source
it is compiled from.

HiveOS

derostorm-1.6.2.tar.gz is attached. Point a flight sheet's Installation URL
straight at it. DeroStorm is a solo miner — the "Pool URL" field takes a
derod node address, and accepted counts miniblocks and blocks, because there is
no pool.

Not measured here

A real multi-GPU rig. This machine has one card. If you run two or more, the
thing worth checking on the first run is that per-card hashrates line up with the
right cards.

Checksums

SHA256-1.6.2.txt is attached.