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.

DeroStorm 1.6.0 - zero developer fee

The 1.5 line ended up somewhere worth marking. 142,900 H/s on an RTX 5080 and 175,100 with the CPU beside it, a miner a rig manager can read, and a HiveOS package that has been tested on rigs it was originally broken on.

No kernel changed since 1.5.8. This is the line drawn under everything since.

Where the speed came from

	H/s
RTX 5080 at 1.4.0	91,870
at 1.5.2	100,640	+9.6%
at 1.5.5	119,463	+18.7%
at 1.5.6	127,930	+7.1%
at 1.5.8, unchanged since	142,900	+11.7%
CPU + GPU together	175,100

+55% on the card across the 1.5 line, every step measured interleaved against the build beside it on the same machine on the same day, and every one bit-identical to the CPU.

What 1.6.0 adds over 1.5.10

Nothing but this, and it is worth saying plainly:

No developer fee

Every hash pays the address you configure, and nobody else.

No fee, no fee period, no fee-off switch that turns itself back on, and no second address in the source. The wallet is read once from your config and handed to the engine, and it is the only address the miner has. grep -rn "dero1" cmd/ internal/ returns the placeholder in the setup prompt and nothing else — a short read, and worth doing yourself rather than believing a release note. MIT licensed, so it stays checkable in every version.

If DeroStorm earns for you and you would like to send some back, that is entirely optional and changes nothing about how it runs:

dero1qypj3sctlt7mefhvdhrvrygj55m40ugl7ml2dukzypxdtd2agpgsjqq2v3n6h

Copy it rather than typing it. DERO addresses carry a checksum so a typo is refused rather than lost, but the refusal is not obvious when it happens.

The rest of the 1.5 line, for anyone arriving here first

A miner you can watch from something else (1.5.9). --stats-file=<path> writes a JSON document every five seconds — hashrate, the split by device, temperatures, fans, power, miniblocks, rejects. Written whole and renamed into place so a reader never sees half a file, and removed on exit so a monitor can tell a stopped miner from one running at zero.

HiveOS (1.5.9, fixed in 1.5.10). derostorm-1.6.0.tar.gz is attached here; 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.

Rigs with more than eight cards (1.5.10). maxGPUs was 8, so a twelve-card rig started twelve workers, four of which quit on a bounds check while the console went on reporting twelve. The cap is 16 now — the nonce tagging always allowed it — --gpu=all trims to what it can address, and a device that will not be mined can no longer be counted as one that is.

CPU threads on a rig (1.5.10). The default was a desktop's. h-config.sh now sets CPUs − cards fed − 1, floored at one, so a four-thread rig with six cards runs one CPU miner instead of three fighting six GPU feeders.

The build refuses to ship stale kernels (1.5.9). 1.4.0's Linux archives carried the previous CUDA kernels because nothing checked. Both build scripts now stop if an embedded library is older than the source it is built from.

Verified

gpu/hash_parallel_test.exe matches all 512 reference vectors. go test ./cmd/derostorm/ is green. Three interleaved --bench --gpu=all runs read 171.8 / 175.1 / 174.6 KH/s. The HiveOS archive was unpacked and run end to end on the reference machine: flight sheet to config, per-card stats with temperatures and PCI bus mapping, and a clean stop.

Not measured here: a real multi-GPU rig. The cap is fixed and the rejection path is tested, but 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.

DeroStorm 1.5.10 - rigs with more than eight cards

Two bugs that only appear on a mining rig, and neither of them was visible on a desktop. Kernels, hashes and hashrate are unchanged from 1.5.8 — 142,600 H/s on an RTX 5080, 175,400 with the CPU beside it — and this release exists because 1.5.9 shipped a HiveOS package before anyone had thought hard about what a HiveOS box actually looks like.

A twelve-card rig was mining on eight of them and said it was mining on twelve

maxGPUs was 8. --gpu=all took every device the driver reported, so on a rig with more cards than that it started a worker per card, and the workers past the eighth hit a bounds check and returned immediately. The count the console reported came from the map those workers had already been added to.

The result: four cards idle, four error lines in a log nobody reads, and a headline figure claiming twelve GPUs. The hashrate looked low and nothing said why.

Three changes:

• maxGPUs is 16. The comment above it always said the nonce tagging allows it — GPUs take byte 47 in the range 0xf0..0xff — and 8 was simply conservative. Sixteen covers essentially every rig; past that the tagging scheme itself has to change, and that is not a quiet fix.
• --gpu=all trims to what the miner can address rather than handing out indices that are going to fail.
• SetGPUs refuses an out-of-range device before it is counted, so a card that is not being mined can never be reported as one that is.

--gpu=0,1,2,16 now says "16" is not a device index between 0 and 15 instead of starting something doomed.

The CPU thread default was a desktop default

The miner picks logical CPUs − 1 mining threads. That is right for a desktop and wrong for a rig, and DeroStorm's own measurements say so: every card needs a host thread to feed it, and a CPU miner competing with that feeder costs more than it earns. On a four-thread rig carrying six cards the old default started three CPU miners, starved six GPU feeders, and burned power to add a rounding error.

h-config.sh now sets --mining-threads to CPUs − cards fed − 1, floored at one:

rig	threads
4 CPUs, 6 GPUs	1
2 CPUs, 8 GPUs	1
12 CPUs, 6 GPUs	5
16 CPUs, 1 GPU	14
32 CPUs, 8 GPUs	23

Cards fed, not cards installed — it reads --gpu=off and --gpu=0,1 out of the extra arguments, because reserving a thread for a feeder that will never exist is the same mistake in the other direction. --mining-threads in the flight sheet still wins.

What was measured, and what could not be

The block count needs no tuning: the suffix kernel plateaus at four resident blocks per SM and stays flat above it — 134,209 H/s at 336 blocks against 134,214 at 672 and 133,597 at 1,344. The default is already the plateau.

Linux hashrate could not honestly be tuned here, and the package README says so. Every Linux figure available to this session was taken under WSL, which virtualises the GPU: four runs of the same binary read 136.4, 122.2, 117.1 and 137.6 KH/s, a 15% spread, against 5% for the same binary on Windows. That noise is the virtualisation, not Linux, and tuning against it would have been fitting to nothing. HiveOS runs on bare metal and should sit closer to the Windows figure — the last honest measurement of that gap, at 1.4.1, was about 2.5%. Numbers from a real rig are worth more than anything this session could produce.

Unchanged

Hashes, kernels and hashrate. gpu/hash_parallel_test.exe matches all 512 reference vectors, go test ./cmd/derostorm/ is green, three interleaved --bench --gpu=all runs read 175.4 / 175.0 / 176.1 KH/s, and the HiveOS package was tested end to end from the published archive: flight sheet to config, run, per-card stats with temperatures and PCI bus mapping, and a clean stop that removes the stats file so a stopped rig reads as stopped rather than frozen.

DeroStorm 1.5.9 - HiveOS, and a status file to run it from

Tooling, not speed. Hashes, kernels and hashrate are identical to 1.5.8 — 142,400 H/s on an RTX 5080, 175,400 with the CPU beside it — and the whole of this release is about running the miner on a rig you do not sit in front of.

A miner you can watch from something else

--stats-file=<path> writes a JSON document every five seconds and nothing else:

derostorm --no-tui --stats-file=/run/derostorm.json

{
  "version": "1.5.9", "state": "mining", "uptime": 3671,
  "hashrate": 175154, "hashrate_cpu": 33120, "hashrate_gpu": 142034,
  "miniblocks": 47, "blocks": 1, "rejected": 0,
  "devices": [
    { "label": "CPU",   "is_gpu": false, "hashrate": 33120,  "temp_c": 61.5 },
    { "label": "GPU 0", "is_gpu": true,  "hashrate": 142034, "temp_c": 55,
      "fan_pct": 42, "power_w": 225.1, "name": "NVIDIA GeForce RTX 5080" }
  ]
}

It exists because the alternative is parsing a console that was laid out for a person, and that breaks the first time a column is widened.

Two details that are the whole point of it. The file is written to a temporary name and renamed into place, so a monitor polling it gets the previous document or the next one and never half of one — a truncated read is how a shell script reports a hashrate of zero and wakes somebody at 3am. And the file is deleted when the miner exits, so a monitor can tell a stopped miner from one running at zero.

A missing sensor reads null, not 0. A card at 0% fan and a card with no fan sensor are different things, and only one of them is worth an alarm.

HiveOS

hiveos/derostorm-1.5.9.tar.gz is attached to this release: a HiveOS custom miner package built on the stats file. Point a flight sheet's Installation URL straight at it.

field	value
Miner	Custom
Miner name	derostorm
Installation URL	the tarball URL from this release
Wallet	your dero1… address
Pool URL	your derod node, e.g. 10.0.0.5:10100
Extra config arguments	--mining-threads=12 --gpu=all, or empty

DeroStorm is a solo miner. It speaks derod's getwork, so the "Pool URL" field holds a node address and there is no pool, no pool account and no pool share accounting. HiveOS's accepted column counts miniblocks and blocks instead. A URL with a scheme or a path is accepted and reduced to host:port.

The rig total includes the CPU; the per-card figures are GPUs only, so on a machine mining on both the cards will not add up to the total. The difference is the processor.

Three bugs found by testing the wrapper rather than reasoning about it, all of which would have shipped:

• h-stats.sh is sourced, not executed, so $0 is the caller and a cd in it moves HiveOS's own working directory. It now resolves paths from BASH_SOURCE and changes no directory.
• exec cmd | tee does not exec anything — a pipeline is forked, so $! was tee's pid and a stop signal went to the wrong process. The miner was then killed hard, without running its own shutdown, and left a stats file behind that looked like a running rig for a minute. The miner is now the direct child.
• bus_numbers could disagree with the card count, which happens whenever --gpu names a subset. HiveOS lines the per-card arrays up by that list, so the wrong length shows each card's hashrate against a different card. It is now one per card or empty.

The build refuses to ship stale kernels

1.4.0's Linux archives carried the previous CUDA kernels. nvcc cannot build a .so under Windows, the machine cutting that release had no Linux toolchain, and nothing objected because the stale file was sitting exactly where go:embed wanted it — Linux mined at 1.3.0 speed for a whole release while Windows had everything, and 1.4.1 exists only to undo it.

build.ps1 and build.sh checked that the four embedded libraries existed. They now check that each is newer than the sources it is compiled from, and stop the build if not:

cmd/derostorm/libderostorm_gpu.so is older than the source it is built from
  (gpu/desc.cuh) -- rebuild it with gpu/buildlib.sh, on Linux

It is a timestamp comparison and it will occasionally fire on a file that changed nothing. Rebuilding is cheap; shipping the wrong kernels is not.

Unchanged

The hash, the kernels and the hashrate. gpu/hash_parallel_test.exe matches all 512 reference vectors, gpu/desc_test.exe is correct on all 512 suffix arrays, go test ./cmd/derostorm/ is green, and the miner still re-verifies each device against the CPU before it will submit anything from it. The Linux archive was unpacked and benchmarked on the reference machine before release, which is the check 1.4.1 taught us to do.

DeroStorm 1.5.8 - GPU +11.2%, three key bytes not four

Two changes to the GPU descriptor sort. 142,400 H/s on an RTX 5080 against 128,000, measured interleaved against the shipped 1.5.6 binary on the same machine on the same day.

	1.5.8	1.5.6
RTX 5080, --bench --gpu=all	142,400 H/s	128,000	+11.2%
CPU + GPU, --bench sum	175,400 H/s	160,900	+9.0%

The real mining path was not re-measured; it needs a node and submits shares. Read the --bench sum as an upper bound on it. CPU mining is untouched.

Hashes are unchanged and still bit-identical to the CPU. gpu/hash_parallel_test.exe matches all 512 reference vectors, gpu/desc_test.exe is correct on all 512 suffix arrays, go test ./cmd/derostorm/ is green, and the miner still re-verifies the device against the CPU before it will submit anything.

---

One pass, asked twice (+3.7%)

A GPU hash sorts ~20,000 descriptors and then has to place their positions. Step 5a does the placing: it walks the sorted descriptors in windows of 256, stages each window in shared memory, and writes the output coalesced.

Step 5 then walked all 20,000 again, reading three global words for each one — the descriptor, its predecessor and its successor — to ask a single question: does anything else share this key? About 98.7% of the time the answer is no, and step 5a had already read all three words.

So the question is asked there instead, out of shared memory, and the ~250 groups that answer yes are pushed to a compact list. Step 5 iterates that list. The list grows down from the top of the merge's scratch array while the merge's own allocator grows up from the bottom, so it costs no memory and the overflow guard that was already there covers both.

phase	before	after
expand to sa	368.2M cycles	392.2M
find groups	326.8M cycles	137.8M
the pair	695.0M	530.0M

This is the same mistake 1.5.0 fixed once already and the 1.5.5 scatter rewrite re-introduced: two loops over all nd descriptors, both opening by asking the same question of the same words. It is worth checking for after any change that splits a pass in two.

The sort was ordering a byte it did not need (+6.4%)

The column walk builds a four-byte key and groups the suffixes that share all four, so the radix sort ordered 32 bits in five passes.

Three bytes is enough, and it wins twice: the sort orders 24 bits in four passes, and it has fewer descriptors to order, because a coarser key merges neighbouring groups into one. What it costs is collisions — and the merge that resolves them was 1.6% of the kernel against the radix sort's 20.7%, so there was room for it to grow into.

DESC_KEY_BYTES	desc_test SA/s
4 (was shipped)	115,946 / 116,215
3	117,205 / 125,742
2	105,906 / 100,657

Two bytes is a clear loss, so three is a peak and not just a direction. DESC_CHUNKS, BR_BITS and DESC_MERGE_WIDE were all re-swept against the new shape over six interleaved rounds each and none of them moved.

This is the trade the CPU sort settled long ago — DSA_KEY_BYTES is 3 there, and a fourth byte measured a clear loss for exactly this reason — and the GPU had never been asked the same question.

What the hardware counters say, now that they can be read

ERR_NVGPUCTRPERM was cleared by hand this session (NVIDIA Control Panel → Desktop → Developer settings → Manage GPU Performance Counters → all users; no reboot needed, it applies to newly launched processes).

Nothing in suffix_kernel is saturated — DRAM 34%, SM 31% — so it is a latency kernel. Its scheduler finds an eligible warp on 23.8% of cycles, the L1 hit rate is 55.6%, only 17.25 of 32 threads in a warp are live, and 57% of every global sector fetched is excessive. The 27.8 cycles between two issues divide almost evenly between waiting at a barrier for sibling warps (35%) and waiting on an L1TEX load (33%).

By source line, the wasted sectors are 43% the text gather in descLoadBE64, 20% the arena write in the column walk, 10% the radix scatter. The largest single stall site is the first instruction after the barrier that ends the column walk — the walk's imbalance wearing a load's clothing.

Seven changes built against that data, and seven killed

hand the runs out longest-first, to even the walk	-9%
arena laid out column-major, so a warp's runs write adjacent spans	+0.5%, noise
warp-uniform chunks, so all 32 lanes walk one column	-8%
L2 persistence pinned on the texts	-8%
col_same as a 64-bit register mask instead of a 64-byte local array	-1.4%
five uint2 loads for a whole 32-byte suffix comparison step	-1.8%
a two-load eight-byte text load	-4.6%
seed the column walk by its key, full compare only on ties	-4.2%

Every one was exact and correct on all 512 vectors. They divide into two lessons.

This kernel will not trade text locality for anything — not for coalescing, not for balance, not for cache residency. The first four are that, and the losses are 8-9% where the gains are noise.

And it is register-saturated at exactly 64. 64 registers is what fits four blocks on an SM, which is the occupancy the whole kernel is tuned around, and ptxas will not go to 65 — it spills instead. The last four all added a live value, all spilled, and all paid more for the spill than they saved. An optimisation here has to be register-neutral or register-negative before anything else about it matters.

A note on sweeps

cmd.exe splits a batch argument on = as well as on spaces, so a build script taking its defines as %2 receives -DDESC_CHUNKS=8 as -DDESC_CHUNKS and defines it as 1. That produced a confident and wrong "DESC_CHUNKS 2 and 8 are 10% worse" — one build measured twice — and it nearly killed the three-byte key by making its own control five times slower than it is. Defines go through an environment variable now. Distrust any sweep where two settings land on the same number.

DeroStorm 1.5.6 - Stage 1 under the sort

Stage 1 now runs underneath the suffix sort instead of waiting for it. 127,930 H/s on an RTX 5080 against 119,463, and 159.0 – 159.5 KH/s with the CPU mining beside it.

	1.5.6	1.5.5
RTX 5080, --bench --gpu=all	127,930 H/s	119,463	+7.1%
CPU + GPU, real mining path	159.0 – 159.5 KH/s	153.9 – 154.7	+3.2%

Hashes are unchanged and still bit-identical to the CPU. gpu/hash_parallel_test.exe matches all 512 reference vectors, the pipelined path is checked against the one-batch path by go test -run GPU, and the miner still re-verifies the device against the CPU before it will submit anything.

---

The card was doing one thing at a time

A GPU hash is three kernels: stage 1 builds a ~71 KB text, the suffix sort sorts it, and a SHA-256 check turns the result into a difficulty test. They ran strictly one after another, because all three shared a single set of texts and suffix arrays and none of them could start until the last had finished with them.

Measured on a 5080 at the default 32,768-nonce batch:

kernel	ms a batch	share
suffix_kernel	207.5	77.8%
stage1_kernel	39.0	14.7%
sha_check_kernel	20.0	7.5%

So 22% of the card's time was spent on two kernels while the thing that actually does the work stood still.

They do not have to take turns

The three are short of completely different things. Stage 1 is capped at about 193 threads an SM — it needs 516 bytes of shared memory per thread for the state machine and the RC4 permutation, and no block size changes that, which is 9.4% occupancy and the whole reason it is slow. The suffix sort runs four 256-thread blocks an SM on 11 KB each and waits on memory latency. The SHA check waits on bandwidth.

gpu/overlap.cu is a new probe that puts two of them on two streams over separate storage and times the pair against each alone:

stage 1 alone      13.01 ms       suffix alone   29.88 ms
both, two streams  32.48 ms   ->  80% of stage 1 hidden

sha alone           9.95 ms       suffix + sha   31.91 ms
                              ->  80% of the SHA check hidden

Both look free. Only one of them is, and finding out which is most of this release.

Two banks, a stream each

The inter-kernel storage is now two banks. A chunk takes a bank and owns it from stage 1 through to its SHA check, and each bank has its own stream, so the ordering a chunk needs against the previous chunk on its own bank is the stream's own and costs no events at all.

One thing does need an event. The suffix sort's scratch pool is indexed by blockIdx and shared by every block in the grid, so two suffix kernels must never overlap however independent their data is; each bank waits for the other's sorted event before launching its own.

The banks together hold exactly what one bank held, because the chunk is divided by the same number they multiply it. The overlap is paid for in chunk size, not in VRAM.

The half that does not pay

Letting the SHA check run under the next chunk's sort as well makes the miner slower than never having started:

	H/s
one bank	123,265
two banks, SHA overlapping too	116,116
two banks, SHA kept out of the way	127,532

The SHA check reads back the suffix arrays the sort has just written, so it fights the sort for the one thing the sort is short of. Profiled, its own elapsed time nearly trebles — 20.0 ms a batch becomes 55.5 — and the sort gains nothing for it. Stage 1, which waits on shared memory rather than DRAM, costs the sort 6.5% of its own elapsed time and hands back four times that.

So the sort's turn is released after the SHA check rather than before it. DSG_OVERLAP_SHA is the knob, and 0 is the shipped answer.

Per batch, before and after:

one bank      suffix 207.5   stage1 39.0   sha 20.0   = 266.5 ms, all serial
two banks     suffix 221     stage1 55.1   sha 25.0   =   246 ms, stage 1 hidden

DSG_BANKS=1 rebuilds the old one-thing-at-a-time shape if you want to see it.

---

Measured and not kept

Three things were built or swept in this round and are not in the release. They are written up in gpu/derostorm_gpu.cu and the README so nobody repeats them.

Aligned loads in the stage-1 hashes. ld32le in crypto.cuh builds a 32-bit word out of four byte loads and ld64le out of eight — exactly the mistake the suffix kernel's descKey32 had, and every load xxhash64, siphash24 and fnv1a64 make walks the state buffer from a four-byte-aligned base. Fixing it is exactly neutral: 25.6 ms against 25.8, inside the noise. Stage 1 is not short of load slots, it is short of threads.

Ablating stage 1 to find its cost centres. This does not work at all, and it is worth knowing why. Every part of the loop feeds lhash, and lhash picks the next operation — so removing any part sends the state machine down a different path through a different number of iterations. Removing all three probabilistic hashes "saves" 69%; removing them one at a time saves 2% each. Both numbers are real measurements of something that is not the question.

S1_BLOCK on its own. 32, 64 and 128 are flat, and 96 and above will not launch — stage 1 never asked for the >48 KB dynamic shared-memory attribute. It would not help if it did: 516 bytes a thread pins an SM at ~193 threads whatever the block size. Under the overlap, 64 stays best (127,532) with 16 close behind (127,397) and 32 worst (125,623).

What is left

suffix_kernel is still 78% of the card, and its column walk still has the 1.98× load imbalance measured at 1.5.5 — the slowest thread takes twice the average, so a perfectly balanced walk would cost half of what this one does, which is ~16% of the kernel. Nothing here touched it, and the two ways of collecting it that were tried at 1.5.5 both failed. The SHA check is now the largest exposed cost at 25 ms a batch, running at about 465 GB/s of the card's ~960.

The CPU row of the README's headline table reads 31,150 H/s. Nothing on that path changed in this release; that is this machine on this day, and it is left as measured.

Checksums (SHA-256)

1233a0bb93e95354fdbbd5f9af5517bade21092e46dada3c24b2b6498ccc2881  derostorm-1.5.6-windows-amd64.zip
fee003daccd5cf3cc52b256cd05a477193482dccb8cc11fead7bc38f5dd8a4f3  derostorm-1.5.6-linux-amd64.zip
0b02e470f6ef9599e3664a728761965040042ca031caaad0823419d9c4bc4fa4  derostorm-1.5.6-linux-arm64.zip
11bfe70e4083ab96219a5534e5a8533342cadbd82b9c4ea9a8a42257fb779993  derostorm-1.5.6-macos-amd64.zip
7696351db085a61deb11c6d851cb285b98ef80179341889b951175f80de5d98d  derostorm-1.5.6-macos-arm64.zip

DeroStorm 1.5.5 - The GPU stops waiting

The GPU stops waiting. A batch pipeline so the card is never idle between launches, and a scatter that writes coalesced instead of a word at a time — 119,463 H/s on an RTX 5080 against 100,640, and 153.9 – 154.7 KH/s with the CPU mining beside it.

The card was waiting for a CPU

Taking two CPU mining threads away made the GPU faster, which should not happen: the threads and the card share nothing but a job. With every logical CPU busy the GPU ran at 106,429 H/s against 111,247 with the CPU idle, and the run-to-run spread went from 1% to 4.5%.

It is not the clocks — nvidia-smi reports 2,865 MHz either way at the same temperature. It is not the kernels — profiled with nsys, suffix_kernel takes 227 ms a batch whether the CPU is idle or saturated. The card was idle between batches, waiting for the one host thread that enqueues the next one to get a scheduler slot behind sixteen pinned mining threads that never yield.

The miner now keeps a second batch queued behind the running one, so the card starts it the instant the first ends and the host's wake-up happens with a whole batch of slack. Both batches ride one stream and share the same scratch; all the second slot costs is a few hundred bytes.

nonces per batch	before	after
32,768 (default)	106,429 H/s	112,548 H/s	+5.7%
4,096	55,504 H/s	72,728 H/s	+31%
1,024	17,740 H/s	31,604 H/s	+78%

The gap is a fixed cost per batch, so it is a rounding error against a batch that takes 290 ms and it is most of the time against one that takes 30 — which is why --gpu-batch, sold as a job-latency knob, was quietly also a throughput one. It is not any more: with the pipeline the GPU rate no longer depends on what the CPU is doing at all.

The scatter was writing a word at a time

Placing the sorted positions gave a whole descriptor — a run of ~5.7 positions — to one thread, which copied it one word at a time. Neighbouring threads wrote addresses ~23 bytes apart and read arena slots with no relation to each other, so a warp's 32 words cost ~28 memory transactions where 4 would do. It was 26.5% of the suffix kernel, the largest single phase in it.

Driving the loop by output position fixes the writes completely — output position q always lands at sa[q], so a warp writes 32 consecutive words. Finding which descriptor owns q is done by walking descriptors rather than output tiles: a window of 256 descriptors is read once, coalesced, into shared memory, and the ~1,460 output positions it covers are written from it.

	suffix ms
before	140.9
one thread per output position	136.7
two barriers a tile, not three	135.1
by descriptor window, not output tile	129.9

Both merges lose their gather loop as well — the positions are already where they want them. suffix_kernel is 226.4 → 207.2 ms a batch, and the run-to-run spread fell from ±5% to ±0.3%.

Headline

	1.5.5	1.5.0
RTX 5080, --bench	119,463 H/s	100,640	+18.7%
CPU + GPU, real mining path	153.9 – 154.7 KH/s	130.4 – 131.8	+17.6%

Hashes are bit-identical to the CPU. gpu/hash_parallel_test.exe matches all 512 reference vectors, and the miner still re-verifies the device against the CPU before it will submit anything.

What did not work

With the scatter fixed the column walk is the largest phase again, and most of it is threads doing nothing: the walk's slowest thread takes 1.98× the average, so a perfectly balanced walk would cost half of what this one does. Two ways to collect that were built, measured and thrown away — cutting the runs shorter is four times slower at every cap, and splitting the columns unevenly landed within noise of the flat split across five threshold settings. Both are written up in gpu/desc.cuh and the README so nobody tries them blind. BR_BITS, DESC_CHUNKS, DESC_MERGE_WIDE, DESC_CMP_WORDS and DESC_SPLIT were all re-swept and all stand.

The CPU row of the headline table reads 32,640 H/s against 34,293 at 1.5.0. Nothing on that path got slower — it is this machine on this day, and it is left as measured rather than carried forward from a better one.

Checksums (SHA-256)

c786dca2467323b71e211a218488296576d5e4cc5dcbe912ad2ebb32f9656a41  derostorm-1.5.5-windows-amd64.zip
a4ee77fef21ec8d6e1e7910a23a45ab6481900a979cd651d6f53c167667b9086  derostorm-1.5.5-linux-amd64.zip
999b37b9b09551f8ae3a1c0be2d8279c3f61bbe98a412717bc506a6c65339d13  derostorm-1.5.5-linux-arm64.zip
75bd9859b080e5d7bdd6a44935dff5600e9e14acea3738224eed804527d10fd3  derostorm-1.5.5-macos-amd64.zip
b22893f3752ef469b431ea1011ef1b69489a3ae58ac6c966661274f4e7d92030  derostorm-1.5.5-macos-arm64.zip

DeroStorm 1.5.2 - Fill the GPU

Fill the GPU. The batch now uses the VRAM it already had, and mining sits at four suffix blocks per SM instead of sweeping up to 1,252.

GPU

--gpu-batch=0 already claimed to size the launch from free VRAM; the library still launched 8,192 hashes, which left stage 1 and SHA on half-empty SMs. It now fills the card, capped at 32,768 so job latency stays under about half a second. On a 5080 that is 30,016 hashes a batch, ~350 ms, and 85.87 KH/s at 336 suffix blocks against 74.9k at 8,192.

Stage 1 writes each 256-byte append as sixteen-byte stores.

Mining defaults to four suffix blocks per SM (the occupancy this kernel actually reaches) instead of sweeping from a few KH/s up through 1,252, which measured slower under a display. --bench still sweeps. Pin --gpu-blocks=336 on a 5080 if you want the same number in the logs.

Windows and Linux GPU archives both ship the new kernels (fat cubins, Turing through Blackwell).

Checksums (SHA-256)

686269c4cd6219d1e4438ecae1ac775eebe45e789f35b24028cb952d1e95ef63  derostorm-1.5.2-windows-amd64.zip
f013e4cf93ab0d2efeb9858cbd5f1fdc7b3f207b922de27305d282062b9c3587  derostorm-1.5.2-linux-amd64.zip
04fabd1297b351b5fc6519d028435c74f641f7d86ad52ac0ab0228724d1494dd  derostorm-1.5.2-linux-arm64.zip
bef36f1c719415c7e4769f3625d7272630acc6f87f353ddb4b7c2fcd37d3e60d  derostorm-1.5.2-macos-amd64.zip
40162b89e35b908cc167650bdd76f88bd1b7f54d9d7c930cc0188cabf47ae736  derostorm-1.5.2-macos-arm64.zip

DeroStorm 1.5.1 - The descriptor sort on Mac

The fast suffix sort on Mac and arm64. Windows and Linux already had it; those builds were still on the Go SA-IS, about a quarter of the same silicon.

Mac and arm64

The descriptor sort is the algorithm that is the 3×, not the language it is written in. macOS and linux/arm64 now use it.

• GitHub archives are cross-compiled with cgo off, so they ship the same algorithm in Go plus hardware SHA pairing.
• ./build.sh on the Mac compiles the C sort and ARMv8 SHA-2 pairing into the binary.
• Darwin mining threads request user-interactive QoS so they prefer P-cores.

Startup should say suffix sort: descriptor (and paired SHA-256 if pairing is live). If it still says SA-IS, you are on the old path.

Also in this release

• DERO difficulty is already hashes per second. The console was dividing it by the eighteen-second block target and showing the network eighteen times slower than it was.
• If the native library is missing on Windows or Linux, the fallback is the portable descriptor, not the Go SA-IS.
• Dashboard: CPU/GPU split bar, share flash, storm pulse.

GPU kernels are unchanged from 1.5.0. No configuration change. Upgrading is replacing the binary.

Checksums (SHA-256)

ed3481c677fe350823908fafb34bfaaf7bfca87a96e644de8dc6978c2efffd47  derostorm-1.5.1-windows-amd64.zip
85f5eeddc47065850cf69f277fd73a95e401bc3a2ab0e1d91a1c2e95eb2b46b7  derostorm-1.5.1-linux-amd64.zip
a30cfcb385bf836cb4c084e9fc0c20311f68321045d2e95aa7238ca22cebd125  derostorm-1.5.1-linux-arm64.zip
31e6a8a2ba63ad5f6affd6d64a007807f56ac2b2d0c9e51c83914a10d59a55bf  derostorm-1.5.1-macos-amd64.zip
1370b316001d1aaa7b58618a759fbadf2e6ce2f1ea8bf650ac9c0bf403d9f1cd  derostorm-1.5.1-macos-arm64.zip

DeroStorm 1.5.0 - A new console, GPU +9.5%, and the native sort on Linux

A new full-screen console, a faster GPU merge, and the native suffix sort on Linux at last.

The console

Eight screens, one key each: M S N T C L P H for mining,
statistics, network, threads, config, logs, pools and help, Tab to cycle,
Esc back to the dashboard, : for the command line, Q to quit.

It reflows rather than assuming a size. Panels drop out of each row as the
window narrows, columns collapse, and labels shorten before they truncate. It is
readable from about 100x32 and opens up fully around 150 columns.

• --preview draws one real frame with sample data and exits — no node, no
  wallet. --screen=<name> picks which of the eight and --size=<WxH> a window
  you do not have.
• --classic keeps the compact in-place panel from 1.4.x, which scrolls with
  the shell so the scrollback above it survives.
• --no-tui is plain scrolling lines for tmux logs, systemd units and CI.
  Colour and cursor control are off automatically when output is not a terminal
  or NO_COLOR is set.
• --termdiag reports what every source says the terminal size is and rules a
  line that wide, for the one failure that leaves the screen unreadable.
• A sixth theme, cyber, and it is now the default.

The GPU: +9.5%

The merge that resolves colliding descriptor groups gave each group to one
thread. The average group holds seven positions — which is why it was written
that way, and why the average was the wrong thing to measure. Every thread waits
at the barrier afterwards, so the phase costs whatever the busiest thread
costs, and over the 512 reference vectors that thread does 43.5% of the whole
text's merge comparisons. Of ~247 colliding groups in a text, the five or so
holding more than 32 positions carry 62% of them.

Those now go through a block-wide merge, one group at a time with all 256
threads on it; the per-thread loop keeps the ~242 small ones. The scatter was
folded into the same pass while the code was open — it and the merge were asking
the same question of the same words, about 50,000 loads a text spent re-deriving
what the previous loop had already worked out.

	1.5.0	1.4.x
RTX 5080, --bench --gpu=all	100,640 H/s	91,870 H/s
CPU, 15 threads	34,293 H/s	33,700 – 34,500 H/s
together, real mining path	130,400 – 131,800 H/s	124,700 – 125,400 H/s

Ryzen 7 9800X3D + RTX 5080. The combined figure is
--run-for=55 --gpu=all --gpu-blocks=1252 over three runs, not the sum of the
two above.

Output is bit-identical. gpu/hash_parallel_test.exe still matches all 512 CPU
vectors, and the miner still re-verifies each device against the CPU at start-up
before it will submit anything from it.

Linux gets the native suffix sort

It was packaged as a DLL and nothing had been built for anything else, so Linux
rigs were running the Go sort and reading about a quarter of the hashrate the
same hardware managed on Windows. It was never a tuning problem.

native/buildlib.sh now builds libderostorm_sa.so, and --bench on
linux/amd64 reads +248.9% at 8 threads against the Windows +260.0% on the same
row. macOS and arm64 still use the Go sort: the descriptor merge is AVX2 and the
paired hash is the x86 SHA extensions, which is a NEON port and not a recompile.

Notes

• All four embedded native libraries were rebuilt from current source before
  this release was cut — the check 1.4.1 existed to make up for.
• No configuration change. Upgrading is replacing the binary.
• Linux arm64 still ships no CUDA .so, as in every prior release.

Checksums (SHA-256)

35f037f6771b078405b59e1b12b8958b0783ba3e7cee9f7f301b7be2d38a3102  derostorm-1.5.0-windows-amd64.zip
f0fe65bac5f661e29a5f80bba95bec1d1144cba6afc1c95b88a0633e550b1097  derostorm-1.5.0-linux-amd64.zip
6fd24cf9aeb73a648145bfa3c5041168b36ddfec7fc62a00271548a6d4613da8  derostorm-1.5.0-linux-arm64.zip
26dd302930d01a712f7d79191777a129cda9c4f661ba03e4301a0123bff85717  derostorm-1.5.0-macos-amd64.zip
7cf32a997313c06cf1556175e88d3350a11f1ad6f44587e7f08e57c81d888b16  derostorm-1.5.0-macos-arm64.zip