hrx: long-context decode 9.70 tok/s at 18.8k ctx (prefill 74-93) vs 31.15 / 459 on a gfx1151 fork — same GGUF, same hour, short-prompt parity 28.6 vs 28.9

Summary

On Qwen3.8-Flash-Next UD-IQ3_XXS (76.3 GiB), our HRX route loses three quarters of its decode rate on a
long context: 9.70 tok/s at an 18,782-token prompt against 31.15 tok/s for a llama.cpp fork built for
gfx1151, on the same GGUF, same requests, same hour. Prefill falls from 74–93 tok/s to 459. Short-prompt
decode is parity (28.6 vs 28.9), which is what makes this a long-context finding rather than a backend
finding.

Numbers

Protocol for every row: 512 greedy tokens (ignore_eos), median of 3, short 216-token / long 18,782-token
prompt, -c 32768 --parallel 1, one gpu-hold claim per arm, strixhalo gfx1151, 22:31–23:20 ADT.
Bold = the two rows being compared.

arm	config	short tg	short pp	long tg (18.8k)	long pp	GTT peak	peak RSS
c1 — our HRX route	house route's own server, `--device HRX0 -ngl 99 -np 1 --cache-reuse 256`, house defaults, no `-fa`, no MTP	28.61 (28.53/28.61/28.67)	465.7 (warm)	9.70 (9.23/9.85/9.70)	74 – 93	106.5 GiB	45.3 GiB
b1 — fork, no MTP	`-fa on -ctk q8_0 -ctv q4_0`, shipped default (its expert cache resolves off at this much device memory)	28.95	336.6	31.15	459.0	50.4 GiB	28.1 GiB
b3 — fork, `--moe-cache off`	same config as b1	31.72	350.9	31.14	465.5	50.4 GiB	28.1 GiB
c3 — fork + its cut-vocab MTP head, n=2	as b1 + drafting	44.86	329.8	47.87	458.4	52.2 GiB	28.4 GiB

Short prompt: 28.6 vs 28.9 — parity. With neither side drafting, the fork is not faster on this box.
Long prompt: 9.70 vs 31.15 — 3.2×, and prefill 74–93 vs 459 — 5–6×.
b1 vs b3 differ only in a flag whose value resolved off in both, so their 29–32 t/s spread is run-to-run
noise on this box, not an effect. The c1 numbers are medians of three, spread 9.23–9.85.

The configuration difference, stated so it is auditable

Our arm ran the house route's own server with the argv/env that 1bit serve --device hrx --print-backend
prints (engine hrxls): --device HRX0 -ngl 99 --jinja -c 32768 -np 1 --cache-reuse 256,
GGML_HRX_REPORT=1, IREE_HAL_AMDGPU_LIBHSA_PATH=…, plus --no-strict-device because 1bit refuses
ffn_moe_gate-47 (MUL_MAT_ID, f32, ne=[640,10,0,1]) under strict placement — the house route leaves 2 of
6060 nodes on the CPU for this model. The fork's arms ran -fa on -ctk q8_0 -ctv q4_0.

The -fa on control on our route was never run. Arm d1 (-fa on -ctk q8_0 -ctv q4_0 on c1's exact
server) was killed during its load when the sequence was stopped, so there is no number for "our route with
flash attention at 18.8k". That is the one missing row, and it decides how the 3.2× is attributed:

if -fa on closes most of the gap → the finding is "the house defaults for this model must enable
flash attention at long context", and it is a one-line default, not a backend defect;
if it does not → the long-context attention path (not the backend as a whole, given short-prompt parity)
is where the work is.

What is already excluded by the data above: "the fork is faster" (it is not, at short context, without
drafting), and "the file or the quant is the difference" (same GGUF, byte-identical).

Why this matters

This is the first file where a published long-context number exists for the house route, and it is the
kind of row a user will hit first: a 76 GiB model at a 32k context. It also sits next to #572 (Flash-Next
prefill 285 tok/s on UD-IQ4_XS vs 1,679 for a tuned Vulkan engine), which is the same family at a short
context — different file, different mechanism, and the two should not be conflated: #572's blocker is a CPU
GET_ROWS node and f32 dispatch overhead; this one is decode attention over 18.8k of history.

Reproduce

# house route (one gpu-hold claim)
llama-server -m Qwen3.8-Flash-Next-UD-IQ3_XXS-00001-of-00003.gguf \
    --device HRX0 -ngl 99 --jinja -c 32768 -np 1 --cache-reuse 256 --no-strict-device
    # env: GGML_HRX_REPORT=1 IREE_HAL_AMDGPU_LIBHSA_PATH=<TheRock libhsa-runtime64.so.1>
# fork, same box, same file, same hour
llama-server ... -fa on -ctk q8_0 -ctv q4_0      # + the fork's own cached build for gfx1151
# client: 512 greedy tokens, ignore_eos, short 216 / long 18,782 tokens, medians of 3

Sampler: peak GTT from /sys/class/drm/card0/device/mem_info_gtt_used at 1 Hz, peak RSS from /proc.
Per-arm server logs, bench.json and the client sample.txt are under ~/lb/aval/*/ on strixhalo.

HRX: Qwen3.8-Flash-Next prefill is 287 tok/s on UD-IQ4_XS vs 1,679 for gufo on the same box (and 500-2,400 across published Strix Halo engines) #572 — Flash-Next prefill on UD-IQ4_XS (short context, CPU GET_ROWS node); different mechanism.
The same arm produced two [gfxhub] page fault + MODE2 GPU reset bursts at GTT 106.5 of 120 GiB; filed as hrx: gfxhub page fault + MODE2 GPU reset when GTT nears the 120 GiB ceiling (Flash-Next UD-IQ3_XXS, 18.8k ctx, 106.5 GiB GTT) #680,
separately to keep this issue about throughput.

Attribution for #679 — long-context decode on HRX (Qwen3.8-Flash-Next UD-IQ3_XXS)

Everything below is measured on strixhalo (gfx1151), medians of 3 unless the row says otherwise, one gpu-hold
claim per arm, GTT sampled at 1 Hz with a hard 95 GiB abort guard.

1. The "one missing row" cannot be measured as written, and it is not the mechanism

Arm d1 (-fa on -ctk q8_0 -ctv q4_0 on c1's exact server argv) does not run: the server loads (49.7 GiB peak
GTT) and the first decode dies with

E graph_compute: unsupported HRX node: SET_ROWS output=621:q4_0[512,8192,1,1]<-SET_ROWS
E llama_decode: failed to decode, ret = -3

HRX's only quantized-KV writer (match_set_rows_addon, ggml/src/ggml-hrx/dispatch_registration/common/dispatch-rope-set-rows.cpp)
accepts GGML_TYPE_Q8_0 only, so a q4_0 V cache has no writer at all; and the HRX flash-attention matchers require
KV element size sizeof(ggml_fp16_t), while llama.cpp force-enables FA for a quantized V cache. So the fork's KV
configuration is doubly unreachable on our route, and -ctk q8_0 -ctv q4_0 would not use the HRX FA kernel even if the
write existed.

The -fa half of the difference is also a non-issue, measured on the same arm at -c 8192, 5.4k prompt, 64 tokens:

arm	flags	tg	pp
d2	`-fa on`	14.65	205
c1 argv	no `-fa` (llama.cpp default is `auto`)	15.91	443
c1 argv	`-fa off`	20.65 (short ctx only)	segfault during the 5.4k prefill

i.e. c1 already ran the HRX FA kernel, and turning it off is both slower and unstable.

2. What actually scales with context: the QSA selection path

Qwen3.8-Flash-Next's full-attention layers are indexer-sparse (QSA, src/models/qwen4exp.cpp). While
n_kv <= indexer_top_k + r - 1 = 2051 the top-k graph is skipped entirely (dense); above it, every such layer, every
step, runs: indexer gather over the whole indexer KV -> block pooling -> ggml_top_k(2051 of n_kv) -> a full-n_kv
mask rebuild -> flash_attn_ext.

One server, one process, decode 64 greedy tokens, median of 3 (-c 8192, f16 KV, no -fa override):

n_kv	pp t/s	tg t/s	ms/step
1,001	694.9	28.06	35.6
1,498	676.1	25.71	38.9
1,986	683.8	25.49	39.2
2,096	595.7	20.42	49.0
2,398	620.6	19.40	51.5
3,005	293.1	17.86	56.0
3,995	491.5	17.33	57.7
5,391	414.6	16.43	60.9
6,978	351.7	14.84	67.4

stepping from 1,986 to 2,096 tokens (the only change is crossing 2051) costs +9.7 ms per decode step;
above it the slope is ~3.8 us per KV token per step, which extrapolates to ~112 ms/step at 18,782 tokens against
the 103 ms (9.70 tok/s) that c1 measured there -- the whole gap to the fork's flat 31.15;
forcing the same path below the threshold (LLAMA_QSA_DENSE=0) at fixed n_kv ~1,480 costs 13x prefill
(676 -> 52.1 tok/s).

GGML_OP_TOP_K is the only op in that path with no HRX dispatch: it is absent from eager_capability_declared()
(ggml-hrx.cpp), no matcher registers for it, so device_supports_op is false, the scheduler puts it on the CPU
backend and splits every full-attention layer's graph around it. The CPU implementation is a single-row
std::partial_sort with an indirect comparator (ggml-cpu/ops.cpp), and QSA has one row per layer per step.

3. The fix and what it buys

PR 1bit-MONSTER/llama.cpp#161: GGML_OP_TOP_K on HRX as a radix select on a monotone uint32 key map (three levels
of 11/11/10 bits to the exact rank-k key, then one emit pass), one workgroup per row, 8 KiB LDS histogram, no global
scratch, no host round trip; plus a bounded claim in device_supports_op so shapes the kernel cannot take keep their
CPU fallback.

test-backend-ops -o TOP_K -b HRX0: 484/484 with ggml_hrx dispatch: hip.top_k_f32 (single), and the
GGML_HRX_DISABLE_DISPATCH=hip.top_k_f32 control degrades to the CPU backend instead of failing the graph.

Long-context rows with the kernel, same binary, same pin, paired against GGML_HRX_DISABLE_DISPATCH=hip.top_k_f32
(identical to the pre-PR behaviour):

arm	n_kv	tg t/s	pp t/s
kernel ON (`hip.top_k_f32`)	pending
kernel OFF (`GGML_HRX_DISABLE_DISPATCH=hip.top_k_f32`, = pre-PR behaviour)	pending

Those rows are queued behind a release-gate window on the box (a 76.3 GiB Flash-Next load needs ~88+ GiB of
MemAvailable and the box is currently hosting other work), together with the 18,782-token protocol row (512 greedy
tokens, median of 3, -c 19456, 95 GiB GTT cap). I will post them here as a follow-up rather than guess.

What is measured for the kernel so far, at the op level (test-backend-ops):

-o TOP_K -b HRX0: 484/484, ggml_hrx dispatch: hip.top_k_f32 (single) in the dispatch log, and the
GGML_HRX_DISABLE_DISPATCH control falls back to the CPU backend instead of failing the graph;
perf -o TOP_K -b HRX0: 106 us (n=1000), 230 us (n=65,000), 623 us (n=200,000) per single row -- so the QSA decode
shape (12 layers, one row each, n=n_kv) should cost ~1.5-2 ms per step against the ~49 ms the CPU path costs at
18.8k. The fixed part is the serial 2048-bucket scan in the radix select, which the GEAK loop will attack next.

4. Not the cause (so no one re-does these)

flash-attention kernel coverage on the full-attention layers: matched and HRX-placed at f16 KV; no rejected nodes
in any server log of these runs.
KV layout / KV bytes: 462 MB of KV traffic per decode step at 18.8k is ~5 ms at 100 GB/s -- not the ~49 ms that
was missing.
The GDN / linear-attention state path: fixed-size state, 36 of 48 layers, no context term in the curve.
The 2 CPU MUL_MAT_ID expert nodes (--no-strict-device): they are per step, not per layer, and have no
context term.

5. Residual, after the kernel

The O(n_kv) mask rebuild that follows the selection (a [1,n_kv] f32 zeros + set_rows + add, then a cast for FA)
is still per-layer-per-step work that grows with context, on HRX but serial. If the numbers above leave a slope, that
is where it is, and the next step is either a fused qsa_mask_build dispatch or an FA variant that reads the selected
indices directly instead of rebuilding a dense mask.

Follow-up on #679 — the long-context rows with the fix in (and what it does not fix)

Same box (strixhalo, gfx1151), same model (Qwen3.8-Flash-Next UD-IQ3_XXS), one gpu-hold claim per arm, GTT sampled at
1 Hz under a 95 GiB abort guard, 18,782-token prompt, 512 greedy tokens, ignore_eos, median of 3.

The fix: fork PR 1bit-MONSTER/llama.cpp#161 — GGML_OP_TOP_K on HRX (radix select on a monotone uint32 key map of
the f32 bit pattern, one workgroup per row, 8 KiB LDS histogram, no global scratch, no host round trip) plus a bounded
claim in device_supports_op. test-backend-ops -o TOP_K -b HRX0: 484/484 with
ggml_hrx dispatch: hip.top_k_f32 (single), and GGML_HRX_DISABLE_DISPATCH=hip.top_k_f32 degrades to the CPU backend
instead of failing the graph — which is what makes a clean pre/post comparison possible on one binary.

The rows

Both rows: same binary, same pin (4a58da9b0 = hrx-base + #159 + #160 + #161), same flags
(-c 19456 --no-repack, c1's argv otherwise). The only difference is the environment variable.

arm	short tg (216)	long tg (18,782)	long pp	peak GTT
`GGML_HRX_DISABLE_DISPATCH=hip.top_k_f32` (pre-PR behaviour)	28.61	10.90 (11.02, 10.78)	131.0 (134.1, 127.9)	84.17 GiB
kernel in (`hip.top_k_f32`)	28.43	12.67 (12.56/12.94/12.67)	388.7 (388.7/411.3/251.4)	83.82 GiB
c1, published above (old pin, `-c 32768`, no `--no-repack`)	28.61	9.70	74–93 (84.5)	106.5 GiB

(the disabled row has two long reps rather than three: I cut it short to hand the GPU to two lanes already queued behind
my claim; the two agree to 2%, and the enabled row's three agree to 3%.)

What the dispatch buys, measured on one binary

	disabled	kernel	factor
long-context prefill	131.0 t/s	388.7 t/s	2.97x
long-context decode	10.90 t/s	12.67 t/s	1.16x
short-context decode	28.61	28.43	flat

Prefill is the win, and it is the expected one: the QSA layers call TOP_K once per layer per ubatch with one row per
token, so prefill was 512 serial std::partial_sorts on 4 CPU threads (with a graph split around every layer); on device
it is 512 workgroups. Long-context prefill now sits at 389 t/s, i.e. within reach of the gfx1151 fork's 459 rather than
5–6x under it.

Decode at 18.8k only moves 1.16x, so the selection was not where decode time went. What remains is the rest of the
QSA path, all O(n_kv) per full-attention layer per step: the indexer get_rows over the whole indexer cache, the
block→cell expand, and the [1, n_kv] f32 mask rebuild (zeros + set_rows + add, then a cast for FA) ahead of a
dense FA over all n_kv. Consistently, the 2051-token cliff survives with the kernel in (26.86 → 21.55 tok/s across it in
the same run): crossing 2051 switches the whole path on, not just the sort.

Correction on the memory side (#680)

I expected the ~24 GiB long-prompt GTT transient (c1: 106.5 GiB) to be the per-layer graph splits, and therefore to
disappear with the dispatch. It does not: the disabled row on the new pin/flags also peaks at 84.17 GiB. So the kernel
has no GTT effect (83.82 vs 84.17, noise), and the 106.5 → 83.8 GiB difference comes from the pin move and/or
-c 19456 + --no-repack, not from this PR. The operationally useful fact for #680 stands either way: at
-c 19456 --no-repack on the current pin, this model's 18.8k row never exceeds 84.2 GiB, ~11 GiB clear of the 95 GiB
rule, with or without the kernel.

Next, in order

The residual decode term — the O(n_kv) mask rebuild and indexer gathers, not the selection: either a fused
qsa_mask_build dispatch or an FA variant that consumes the selected indices directly instead of rebuilding a dense
mask.
GEAK on the new kernel — its fixed cost is the serial 2,048-bucket scan (~100 us/row); the plan is cooperative
scan / warp-aggregated histogram / 8-bit levels.

hyperloom custom: size the long-context prompt by what the tokeniser …

hyperloom custom: a long-context decode leg, so the tuner can score e…

TOP_K on HRX is merged on the fork (#161 → 199165047), but this issue stays open for two reasons, both stated so it is not closed on the merge:

it is in no build: engine main pins third_party/llama.cpp = 4a58da9b0, 8 commits behind the fork tip 199165047, and there is no open bump PR (#759 → 37ca4d442 was closed as superseded at 15:54:53Z).
the row still needs a fault-free arm: with the TheRock nightly libhsa, any run without HSA_USE_SVM=0 HSA_USERPTR_FOR_PAGED_MEM=0 HRX_ALLOW_NO_SVM=1 faults on its first graph replay (hrx: gfxhub permission page faults (UTCL2 client CPF) on the rebuilt box with the TheRock nightly libhsa — the no-SVM env trio is the only fault-free arm #715, gfxhub … UTCL2 client CPF), and test-backend-ops hangs on the fault. No fault-free measurement of the 18.8k row exists yet.

Closes on: the pin bump, then arm d1 (-fa on -ctk q8_0 -ctv q4_0) at 18.8k, memory-capped and alone, fault delta 0, beating the 28.43 tok/s short-prompt baseline.

Follow-up: the TOP_K kernel from #161 now has a 2.38x geomean improvement (op-level), branch 1bit/hrx-top-k-coop-scan ready.

The dispatch fixed the cliff and the prefill gap, but it left one measured defect: inside the radix select the descending
bucket scan ran serial in lane 0 — 3 levels x 2,048 buckets = 6,144 dependent LDS reads plus integer work,
independent of n. That is the ~100 us fixed cost test-backend-ops perf -o TOP_K -b HRX0 shows at n=1000, and it is
paid by all 12 full-attention layers on every decode step.

Each lane now owns 8 consecutive buckets; the 256 per-lane chunk sums get a reverse (suffix) inclusive scan as an LDS
Hillis-Steele tree (8 steps, 2 syncs each); the crossing lane is the one with suffix_incl >= remaining > suffix_excl,
and only that lane walks its own 8 buckets. Same bucket, same residue, same tie behaviour as the serial scan.

Same-hold A/B against the merged kernel, arms interleaved A/B/A/B, best of 3 rounds per shape, only that one .hip
differs:

shape	before	after	ratio
n=1,000, rows=1, k=400	111 us	13 us	8.74x
n=1,000, rows=16, k=400	108 us	14 us	7.48x
n=65,000, rows=1, k=400	246 us	155 us	1.50x
n=65,000, rows=16, k=400	341 us	258 us	1.29x
n=200,000, rows=1, k=400	643 us	511 us	1.25x
n=200,000, rows=16, k=400	908 us	766 us	1.16x

geomean 2.38x, worst 1.29x. The 111 -> 13 us at small n is the fixed cost disappearing; the n-proportional slope is
untouched, which is why the large-n rows only gain 1.16-1.25x. Correctness: -o TOP_K -b HRX0 484/484 with
hip.top_k_f32 matched, the GGML_HRX_DISABLE_DISPATCH control still declining every case as unsupported, CPU
reference 529/529. Registers 17 VGPRs (was 8), 0 spills, occupancy 16 waves/SIMD, LDS 9,228 B (was 8,204). Only the
kernel changed — the matcher, its shape bounds and the device_supports_op claim are byte-identical, and the add-on ABI
check still reports 336 kernels matching their code objects.

Two things worth recording for whoever touches this next:

Do not use __shfl_down_sync here. HIP rejects a 32-bit mask on this toolchain, and with a 64-bit mask an
unguarded add inflates every suffix sum at the top of the wave, because shfl_down returns the caller's own value
for an out-of-range source lane rather than zero: ERR = 4.0 on n=2/rows=2, ERR = 0.667 after adding the range
guard. The LDS tree is wave-agnostic and cheap; the shuffle is not worth the tuning.
The harness's "484/484" is 484 of 529 cases. The other 45 are shapes the bounded claim declines by design
(k > n, non-f32, n above the template range) and they correctly fall to the CPU backend — not failures.

Not measured, and it is the one thing left: the model row (18.8k decode, kernel ON vs
GGML_HRX_DISABLE_DISPATCH=hip.top_k_f32). It needs the 76 GiB Flash-Next file at -c 19456 with ~92 GiB of
MemAvailable and the owner's serve holds ~35 GiB of this 122 GiB box, so the 18.8k rows for #679 are still
memory-blocked. Every number above is op-level.

Branch 1bit/hrx-top-k-coop-scan @ 8b7966f90 (base 199165047, one commit, one file) is pushed and ready for a PR;
the round log with the measurement protocol (including why a single 81-shape sweep is not usable: a byte-identical
candidate "measured" 0.59x-2.05x on it) is at ~/geak-tasks/tattn-top-k/ROUNDS.md on strixhalo.

Addendum to the comment above — one more in-contract change is in the same branch: geomean 2.38x -> 2.77x.

The emit phase ran two full passes over the row (keys strictly above the rank-k key, then keys equal to it). The
split is already known when the level loop ends — exactly top_k - remaining keys are above, remaining are equal — so
both groups can be written in one pass into disjoint slot ranges, [0, above_slots) and [above_slots, top_k),
with the per-group slot guard keeping the result duplicate-free (the ties variant's requirement). Five passes over the
row become four.

shape	before (merged #161)	+ cooperative scan	+ single emit pass	vs before
n=1,000, rows=1	106 us	13 us	11 us	9.36x
n=1,000, rows=16	106 us	14 us	12 us	9.05x
n=65,000, rows=1	225 us	155 us	133 us	1.70x
n=65,000, rows=16	328 us	258 us	211 us	1.55x
n=200,000, rows=1	617 us	511 us	423 us	1.46x
n=200,000, rows=16	865 us	766 us	620 us	1.40x

geomean 2.77x, worst 1.40x, baseline's own per-round spread 1.01x (quiet box). Correctness re-verified on this exact
candidate: 484/484, control declining all 529, CPU 529/529, 17 VGPRs, 0 spills. Branch head is now 1b292a11b (still one
commit, one file).

One thing the measurement says that changes the next step, worth flagging rather than leaving for a later PR: the
remaining cost is parallelism, not atomics or bandwidth. The kernel reads the row 5 times (now 4), and at
rows=1 it moves ~8 GB/s on a ~100+ GB/s box while rows=16 reaches ~83 GB/s — the 16-row case does 16x the work in
1.54x the time, so throughput scales with workgroup count and rows=1 (one workgroup, ~1/40th of the GPU) is starved.
Splitting a long row across P workgroups is therefore the biggest remaining lever for the QSA shape — but it is not
doable inside the kernel: it needs a global histogram and a device-scope counter, i.e. either an extra binding for
scratch (the ABI is fixed) or a cooperative launch (the matcher launches one workgroup per row). That is a
matcher-level task, so it should be scheduled as one instead of being smuggled into the kernel work.

TOP_K on HRX, round 3 (vec4 row reads): 1.909x over round 1, 5.436x cumulative against the #161 pin. Branch
1bit/hrx-top-k-coop-scan @ b25b0485d (2 commits, 1 file, +101/-28, base 199165047). Op-level only — the 18.8k
model row is still unmeasured and is not being claimed.

Why vec4, and the number that pointed there

Round 1 (1b292a11b: cooperative bucket scan + single emit pass) removed the fixed ~100 µs and one pass, and measured
2.77x against the pin. I then tried candidate compaction (after level 0, write the higher-bucket indices straight to
dst and collect the selected bucket into an LDS candidate list, so levels 1/2 and the emit run over ≤2048 entries — 4
row passes to 2). It was rejected on measurement: 0.973x geomean, worst 0.95x against round 1, on a quiet box. The
passes it removes are cheap because the row is cache-warm by then; its restructured base was ~5-10% slower to start with.

That failure is the useful part: at rows=1 one pass is 200 000/256 = 781 iterations per lane and ~143 ns each
(~300 cycles) = one un-hidden DRAM latency per iteration, because a lane has one outstanding load and the workgroup is
only 4 waves. So the limit is memory-level parallelism, not work. Every pass (3 histograms + emit) now reads float4
when the row base is 16-byte aligned (src + row*column_count, so the guard covers pointer and pitch), with a scalar
tail for column_count % 4 and a fully scalar fallback. Index order is unchanged, so the ties rule and slot numbering
are untouched; the ABI, the matcher bounds and the device_supports_op claim are byte-identical to the pin.

shape	round 1	round 3	ratio
n=1 000, rows=1, k=400	12 µs	11 µs	1.03x
n=1 000, rows=16, k=400	12 µs	11 µs	1.07x
n=65 000, rows=1, k=400	127 µs	55 µs	2.30x
n=65 000, rows=16, k=400	211 µs	81 µs	2.61x
n=200 000, rows=1, k=400	414 µs	165 µs	2.51x
n=200 000, rows=16, k=400	622 µs	212 µs	2.93x

geomean 1.909x, worst 1.03x; cumulative against the pin on the same gate 5.436x, worst 4.23x. Protocol:
same-hold, arms interleaved A/B/A/B, best of 3 rounds per shape, raw per-round values printed (this box's per-sweep
noise is one-sided and reaches 1.6x, so a phase-ordered or single-sample protocol measures the box, not the kernel).

Gates (all re-run on the pushed head)

test-backend-ops -o TOP_K -b HRX0 484/484 with hip.top_k_f32 (single) taken; GGML_HRX_DISABLE_DISPATCH control
declines all 529 cases as not supported [HRX0] with nothing failing; CPU reference 529/529; whole-backend guard
identical 4289 per-case outcomes (the pin's own 320 [VIEW]/TOPMOE failures are present on both arms and their
error text is normalised). Compiler: 18 VGPRs, 34 SGPRs, 0 spills, occupancy 16, LDS 8 228 B. The algorithm and the
rejected design were both modelled in scripts/ref_topk.py against a sort reference first (616 cases, 0 failures).

What this does and does not mean for the #679 gap

There is no expected decode jump: top_k at 18.8k is ~12 x 50 µs ≈ 0.6 ms of a ~79 ms step (~0.8%), so the decode
residue stays where the attribution put it — the O(n_kv) mask rebuild and the dense FA over n_kv (P2 in the kernel
plan). Where it should show is prefill: the same dispatch was worth 2.97x there, and it is now ~5x faster at the
n=8192..262144 shapes prefill produces. The row that would confirm both (18 782 tokens, 512 greedy, median of 3,
-c 19456 --no-repack, kernel on vs GGML_HRX_DISABLE_DISPATCH) needs the 76.33 GiB qwen4exp file and ~92 GiB of
available memory; it is queued behind a headroom gate and no model-level number is claimed here.

Measured tonight: the attention path is ~14 % of this gap, not the gap

The 9.70 tok/s row is from c1 (pre-#161). Since then the QSA path has been fixed in three measured steps
(TOP_K on device #161, the mask rebuild's SET_ROWS, and now sparse index-list attention), so it is worth saying
plainly where the long-context decode time actually goes, because the answer is not the attention.

The measurement that settles it (new tonight, no model needed): the model's entire full-attention set for one
decode step — 12 QSA layers, each with indexer gather + pool + top_k(2051) + the mask rebuild + attention — built as
one graph at the GGUF's real shapes, and timed as one step:

n_kv	12 QSA layers, today's path (Loom decode split)	with sparse index-list attention (#49)
18,816	13.00 ms	9.95 ms (829 us/layer)
2,048	8.02 ms	9.74 ms → guarded back to 7.75 ms

So at 18.8k the whole attention path — the thing this issue was opened about — costs ~10 ms of a decode step. The
model decodes at 12.67 tok/s on today's main, i.e. ~79 ms/step, of which the short-prompt floor is ~35 ms; the
context-dependent slope is therefore ~44 ms, and the 12 full-attention layers account for 10 ms of it (~23 % of the
slope, ~13 % of the step). The other ~34 ms/step of context-dependent cost is not in these layers.

That matches the per-stage probes rather than contradicting them: indexer 0.60 ms, top_k 0.07, mask rebuild 0.07,
attention 0.81 per layer at 18.8k — and 12 layers in one graph cost exactly 12x that, so there is no hidden per-node or
per-split overhead inside the attention graph (the earlier "1.5 ms measured against 5.9 ms/layer implied" gap was me
dividing the model's whole slope by 12 QSA layers, which assumed the slope lives there; it does not).

What is left, with what each candidate needs to be settled

The remaining context-dependent cost has to be one of these, and every one of them needs the model to attribute:

The 36 GDN (linear-attention) layers. Their state is fixed-size, so this is expected to be context-flat — but it
is the largest unmeasured block in the step, and "expected" is not measured.
Per-step host/JIT overhead. This model needs 66 graph programs and the cache holds 64
(GGML_HRX_GRAPH_PROGRAM_CACHE, earlier sessions counted 933 trims in 40 min), so at least one program is evicted and
re-JITed per step. GGML_HRX_GRAPH_PROGRAM_CACHE=128 was measured to give 0 trims on another model and was never
measured here — the row never ran.
Memory pressure. The 18.8k arm peaked at 83.8 GiB GTT against ~50 GiB at 8k, on a box where GTT is DRAM-backed
and the weights are 76.33 GiB. A context-dependent fault rate is exactly the shape of a slope that no kernel-level
profile shows.

Why the model row still has not run

It needs the 76.33 GiB file with ~92 GiB MemAvailable (the arm's own GTT peak is 83.8 GiB). Tonight the box has
49-51 GiB available with the owner's forward-serve holding 31.7 GiB, and the waiter is armed for the moment that
clears. I am not lowering the gate to manufacture the row — the last time this arm ran near the ceiling it took the box
down (engine #680).

Ours, fixed tonight

gpu-kernels #49 @ 22dc9f6: sparse index-list attention for these layers (3.0 ms/step off the attention at 18.8k,
verified against the dense path at NMSE 1.7e-07), including the guard the 12-layer row exposed — at n_kv ~2048 the
selection is the whole cache and the same dispatch was 1.7 ms/step slower, so it now declines when the selection covers
half the cache or more.

What I would do next, in this order, once the box has the memory: (a) the same 18.8k row with
GGML_HRX_GRAPH_PROGRAM_CACHE=128 — one env var, and it either moves the step or kills candidate 2 for good; (b) the
same row with the 36 GDN layers' graph timed separately (they can be built model-free the way the QSA set was); (c) only
then a kernel, if what is left is arithmetic.

Long-context follow-up: multi-token prediction is not a long-context lever (paired A/B, one claim, both orders)

Different file from this issue's headline row, said up front so the numbers are not moved onto Flash-Next: this is
Qwen3.8-27B UD-Q4_K_XL (17.56 GB), 20,480 ctx, 18,762-token prompt, via the Hyperloom custom workload on strixhalo
(tools/hyperloom/custom/benchmark.py, engine 7d7c46f >= d3d2667a6, llama.cpp 1d32c226, hrx-system 82804197bcbe,
kernel 7.3.0-rc4-1bit-perfopt+). Metric: decode tok/s at depth (the serve long leg), plus the short serve legs.

Why paired, and why two orders

The arm-to-arm spread of this metric on this box is 19.3 % (12.938 / 15.435 / 13.342 on three identical no-lever
arms) -- larger than any lever effect measured here. So the lever was measured as A and B inside one gpu-hold claim,
then repeated with the order reversed, because "B second, onto a hotter die" is a real confound (the 2026-11-01 11:13
drafter fault reset the box).

pair (date)	order	A: no levers	B: `HL_MTP_MAX=6`	B vs A
1 (11-02 19:21-19:27)	AB (B second, die 43-64 C)	14.72 decode	13.29 decode	-9.7 % decode; prose -23.5 %, short -26.8 %
2 (11-03 15:19-15:27)	BA (B first, die 53-66 C)	14.5893 decode, geomean 48.8571	11.9825 decode, geomean 47.9897	-17.9 % decode; prose -23.6 %, short -27.3 %, code +3.1 %, geomean -1.8 %

Both halves of both pairs gate PASS at KLD 0 (greedy texts identical), gfxhub fault lines 1 -> 1 (MES banner only).

The sign did not flip, and the prose/short penalty replicates to within 0.7 pp across the two orders (-23.5/-26.8 vs
-23.6/-27.3). Running B first, onto a cooler box, makes the decode penalty larger, so the effect is the lever, not a
thermal ramp; the pair-1 ordering was masking part of it. The decode magnitude itself (-9.7 vs -17.9 %) is not stable
across pairs because the two halves' prefill differed by 27 % in pair 2 (250.85 -> 318.77, second half faster) -- and that
drift runs against the conclusion (the faster half is the one that won), so -17.9 % is conservative. Prefill is not
comparable across days at all (139.57 / 279.93 / 318.77 on the same binary), so decode is the metric of record here.

Mechanism, captured in the same runs: draft acceptance falls 0.6746 -> 0.4381 with depth (drafts 85/126 -> 92/210,
mean draft length 3.02 -> 3.63). Deeper drafts are converted less often, so at depth the lever spends draft work it cannot
turn into tokens.

What this means for this issue

HL_MTP_MAX=6 must not be carried onto long context, even though it is a +13.8 % win on the short-prompt geomean
(50.87 vs 44.69 clean). The two metrics disagree; only the short-prompt number stands.
The long-context gain this issue is after is not going to come from the config levers the tuning loop can sweep:
drafting is measured harmful (above); HL_DFLASH at long context is refused in committed code (hyperloom custom: refuse the drafter at long context (#764's GTT ceiling does not cover it) #773) because it is the
arm that faulted the graph on 2026-11-01; HL_UB/HL_B move the llama-bench legs but not the serve decode leg
(1bit serve rejects -b/-ub/-ctk/-ctv and sizes those per route); ONEBIT_HRX_DECODE_SPLIT=0 is documented as a loss
at ctx >= 512. That points at the kernels -- the attention work in this lane's own thread -- not at the tuner.
Number of record for the shape on the current binary: 14.59 tok/s decode at 18,762 tokens (no levers, pair-2 half A,
prefill 318.77). Use that, not the 14.72 from the older binary, and not the 12.94 or 15.44 single arms.

#679 follow-up comment (draft, 2026-11-03 night)

Posted by pi-attn (strixhalo). Numbers are from this box, this pin, one gpu-hold claim per row; nothing below is a
projection unless it says so.

Two things changed since the last comment, and one candidate is now dead. All rows are on the fork build
b2b545d06 (= base d8bd0ff60 + the top_k GEAK branch).

1. The "our route with `-ctk q8_0 -ctv q4_0`" row is runnable now — and it is 3x to 17x SLOWER

The 10-31 comment said that row could not run: a quantized V had no SET_ROWS writer at all (match_set_rows_addon
accepts GGML_Q8_0 only), so the graph was refused outright. hip.set_rows (PR-READY branch 1bit/hrx-set-rows)
supplies that writer, so quantized KV now runs end to end on HRX with 0 rejected nodes. Measured on Qwen3-4B-Q4_K_M,
-c 32768, 29,150-token prompt, 64 greedy tokens:

KV	fa dispatch	prefill t/s	decode t/s
f16	`hip.fa_prompt_f16`	979.5	25.26
`-ctk q8_0`	`hip.fa_prompt_q8`	57.1	14.77
`-ctk q8_0 -ctv q8_0`	`hip.fa_prompt_q8`	55.7	18.92

Same three at -c 8192 / 3,257 tokens: f16 1153.9 / 59.90, q8k 378.1 / 53.62, q8k8v 350.8 / 46.70.

So the answer to "is quantized KV the missing row" is no, by a wide margin: it costs 3.0x prefill at 3.3k and 17.2x
at 29k on this route, in the same binary, with the quantization itself doing the expected thing (KV bytes halved - the
f16 row also shows prefill falling from 1154 to 979 t/s over the same range, i.e. KV bytes are not the limit). This is a
kernel gap, not a policy one, and not explained by bytes: the q8 cache is half the traffic of f16, so the 17x is in
hip.fa_prompt_q8's own schedule (one workgroup per (query head, 32-query tile), each walking the whole cache and
dequantizing its chunk into LDS - so the cache is dequantized once per query tile per head), not in the quantization.
The KLD/needle A/B for quantized KV is moot at these numbers.

2. Long context is not generically slow on HRX — control row on a non-QSA model

Qwen3.6-35B-A3B-UD-IQ4_XS (17.7 GB, 3B active), same binary, 64 greedy tokens:

ctx	prompt_n	prefill t/s	decode t/s	ms/step
8,192	4,909	607.4	31.53	31.7
32,768	21,227	755.2	33.99	29.4
131,072	82,017	387.0	28.15	35.5

A 35B-A3B MoE decodes at 26-34 tok/s at 82k tokens (30-38 ms/step, 1.4 GiB of KV) where Flash-Next at 18,782 tokens
measures 79 ms/step. "HRX decode degrades with context" is not a general property of the route; the residual is
Flash-Next's own layer mix, which is why the model-free layer probes below are the right instrument.

3. Candidate killed: the HRX program-cache trims do not cost decode time

Flash-Next at 18.8k needed 66 programs against a 64-entry cache and showed 933 trims in a 40-minute run, so
GGML_HRX_GRAPH_PROGRAM_CACHE=128 was one of the three ranked candidates. Measured on Qwen3.6 at 82k with decode-only
reps (the first request pays the prefill, the next two reuse the KV), arms alternating so the box's 8-15%/hour drift
lands on both:

arm	trims	decode t/s (2 reps)
cache default (64)	99	31.47, 33.20
cache default (64), second server	99	30.97, 32.22
`GGML_HRX_GRAPH_PROGRAM_CACHE=128`	89	33.47, 33.35
`GGML_HRX_GRAPH_PROGRAM_CACHE=128`, second server	89	30.30, 31.58

Indistinguishable. Caveat stated so this is not over-read: that model does not overflow the cache the way Flash-Next was
counted doing, so this weakens the candidate rather than closing it on the real model. It also means the 128 setting is
not worth a model row before something else does.

4. The 36 GDN layers: 15.0 ms of the step, and an instrument that is not yet numerics-clean

~/wt/hrx-probe/gdn_layer.cpp rebuilds one GDN layer's decode step from src/models/qwen4exp.cpp
(build_layer_attn_linear + build_conv_state_at) at the GGUF's real shapes (state 128, K heads 16, V heads 48, inner
6144, conv 4, channels 10240, the file's own Q6_K projections) and sums N of them into one graph:

layers	HRX0 p50	per layer	CPU p50
1	582 us	582	913 us
12	5,597 us	466	8,043 us
36	14,973 us	416	21,866 us

36 layers = 15.0 ms of the decode step, context-flat by construction (fixed recurrent state).
Drift over a fixed 36-layer graph, 600 reps: first quarter 14,993 us vs last quarter 14,976 us, ratio 0.999.
With the 12 QSA layers at 10-13 ms and the ~13 ms short-context floor, ~40 ms of the 79 ms 18.8k step remains
unattributed, which is now the whole question.

The caveat, measured rather than assumed. That graph's output is wrong on HRX: 1,400 of 2,560 values are nonfinite
and the finite ones are 29 distinct garbage values repeated 40 times each (~2e8), byte-identically across processes,
and unchanged when I shrink the probe's leaf ranges to mild values. I stage-bisected it by copying each layer-0
intermediate into a leaf (the only reliable read-back on HRX):

stage	HRX nonfinite	max\|err\|/max\|ref\| vs CPU
`llm.ssm_conv` out	0	3.3e-3
v (convolved)	0	3.3e-3
recurrent-state gather (786k)	0	6.1e-3
`common.rmsnorm_gate` out (normed)	0	4.7e-3
gated (ngated)	0	5.8e-3
the activation into the closing `ssm_out` matvec	0	5.8e-3
layer output (`ggml_mul_mat(ssm_out, final)`, Q6_K x 1-token f32)	1400	-

So the recurrent math, the conv state, the state gather and the gated norm are all correct, and the corruption is in the
last node: a quantized-weight matvec with a one-token f32 activation. The shapes are not the issue either - I parsed both
GGUFs, and Bonsai-2-27B (qwen35) has the identical GDN config (group_count 16, time_step_rank 48, inner 6144, state
128) and runs correctly on HRX with the same fused kernel. Switching that closing projection from Q6_K to f32 (same
graph, same values, one env var) makes the output finite and equal to the CPU reference (rms 4.20265 vs 4.20273,
NMSE 1.8e-4) and removes kquant.mul_mat.decode_f32 from the dispatch list, while the same kernel on a leaf activation
is correct in the same graph. So the recurrent math is right and the single broken node is the quantized-weight matvec
whose activation comes from another fused dispatch rather than from a leaf - the shape of a missing activation staging,
not a coverage or shape problem.

What that costs the numbers above: the 15.0 ms is a dispatch/shape timing, and with the file's own Q6_K weights it is
timed on a graph whose last node mis-stages its activation. I am quoting it as provisional, not as a validated row.
The ~40 ms unattributed arithmetic is unchanged either way.

And one thing this lane should say out loud about its own evidence: every number in this issue and in my earlier
comments is a speed row - the 76 GiB model has been memory-blocked for four days, so no KLD, needle or even
coherence check has ever been run on Flash-Next on HRX. If the model's GDN layers take the path found broken above (in
the model the projection may instead be fused into the GDN dispatch - ..._projection_epilogue exists in the kernel
catalog, and GGML_HRX_LOG_DISPATCH=1 on a real load would answer it in one line), the model would still generate tokens
at 12.67 tok/s and nobody would notice. Before the next parity number is quoted, the speed row and a quality check should
be taken in the same window.

Not run, and why: the 18.8k model row. Its own waiter logged 18:27 -> 19:50 on pool >= 88 GiB and
MemAvailable >= 92 GiB; the pool never exceeded 72 GiB and MemAvailable peaked at 66 GiB, because the owner's
1bit forward-serve holds ~30 GiB of GTT. The gate is not being lowered: the one attempt near the ceiling was the
#680 reset.

The long-context config surface is empty, and here is the log evidence — so nobody spends a GPU window on these two names.

This lane measured the 27B UD-Q4_K_XL at 20,480 ctx (HL_LONGCTX=18800, 18,762 tokens as the server counted them). Its numbers of record: 14.5893 tok/s decode (pair-2 half A, one claim) — not Flash-Next, not the 12.67 of the issue. Two sweepable names were left untested at long context (GGML_HRX_GRAPH_PROGRAM_CACHE, GGML_HRX_Q8_PREFILL_RELAX); both are now ruled out from the artifacts rather than by another arm:

The graph-program cache never trims in the long arm. grep -c 'graph program cache trimmed' over all 11 leg logs (serve.log, llama-bench.log, kld.log) of the four pair-2 arm dirs → 0 in every one, including the arm that did the 18,762-token prefill and the 128-token decode. A lever that sets the count cap to 0 (graph_cache_unbounded) cannot move a cache that never evicts, and it does not touch the separate 8 GiB byte budget (plan_evictions(budget, max_programs, …) applies both; GGML_HRX_GRAPH_PROGRAM_CACHE_MB defaults to 8192 and is not in the grid).
Trims happen in the KLD/perplexity leg, not the scored legs: 163 / 164 / 92 in the HL_UB=1024/2048 explore runs — trimmed (1 of 39 programs, 8583 -> 8323 MiB of arenas, budget 8192 MiB, 64 max). A quality-leg cost at widened chunk shapes; no score depends on it.
q8_prefill_relax_off is a prefill-routing knob (docs/hrx.md: prompt matmuls on the q8_1 x4 kernel, 94 % of HRX prompt time, prompt chunks 256–2048), and =0 restores the older, slower routing. The long-context score here is the decode rate, so it can only cost the 320 tok/s prefill and cannot raise decode.

Combined with what this issue's thread already has — HL_MTP_MAX=6 measured harmful at depth in both orders (−10…−18 % decode, −24…−27 % prose/short, draft acceptance 0.6746 → 0.4381), HL_DFLASH refused by committed code (#773) after faulting the HRX graph at 18.8k on 2026-11-01, and HL_UB/HL_B/-ctk/-ctv not reaching 1bit serve at all — there is no configuration lever left for the 18.8k decode rate. The remaining gain is attention-kernel work, which is not this lane's.

Separately, so the artifacts above are attributable from now on: #947 → PR #948 adds bench_summary.device_stack = {kernel, hsa_runtime, hsa_runtime_sha256} to every measurement, because the report named the binary (#933) but not the kernel, and this box rebooted five times in 18 h.

hyperloom custom: refuse a long-context leg that does not fit the pin…

#679 follow-up (pi-attn, strixhalo): the re-JIT price is measured, and it is the first thing in this issue with a number in the right order of magnitude

Model-free work only; the 18.8k row is still blocked (see the end). One number matters here:

A graph-program re-JIT costs ~41 ms, and it is flat in graph size

Measured by timing the first graph_compute of a fresh graph against the immediately following one - the second is
served from the HRX program cache - same graph, same leaves, one gpu-hold claim, 0 faults:

graph	nodes	first	second	price of a re-JIT	steady p50
1 GDN layer	44	41.9 ms	0.94 ms	40.9 ms	0.65 ms
12 GDN layers	528	52.0	6.04	46.0 ms	5.39 ms
36 GDN layers	1584	58.4	15.68	42.7 ms	15.29 ms
36 layers, cache limit 1, twice	1584	56.8 / 59.6	15.8 / 18.5	41.0 / 41.1 ms	15.5 / 15.3 ms

It is compile+load overhead, not per-command emission, so any single trim costs the same whether it is a small decode
program or a large prefill one; and it is 3-4x the whole steady decode of a 36-layer GDN graph. (First compute in a
cold process is 110-136 ms; 41 ms is the warm figure, cross-checked by the steady p50 in those same runs scaling with
the clock ramp and the JIT delta scaling with it.)

Why this matters for this issue, and what it does not say

This model's own 18.8k server log carries 933 trim_graph_programs ... (1 of 66 programs, ... 64 max) lines in one
40-minute run, and 779 in the pre-rebuild run of the same row; each trim evicts one program and frees ~6 MiB, so the
8192 MiB byte budget is never the constraint. Candidate (a) was "raise GGML_HRX_GRAPH_PROGRAM_CACHE; one env var, and
it either moves the step or kills the candidate" - ranked on intuition because nobody had priced a re-JIT.

It is priced now, and it sits inside the unexplained ~44 ms/step: at the observed rate of ~1 trim per 10 decode steps it
is ~4 ms/step; at 1 trim per step it would be ~41 ms/step, i.e. the whole residual. Which of those it is, is
still unmeasured, because it needs the 18.8k row itself.

So I shipped the cheap half as data rather than guess: engine #1006 adds one built-in recipe,
hrx-graph-program-cache-128 (hrx + qwen4exp -> GGML_HRX_GRAPH_PROGRAM_CACHE=128), with the trim counts and this
price table in its measured field and an explicit "the tok/s row is not measured" sentence. It is a recipe and not a
default change for the same reason. Raising the count does not preallocate: the byte budget still caps the cache, and
128 slots of ~6 MiB are ~768 MiB against a 120 GiB GTT budget.

One retraction that removes a red herring from this issue's history

The "GDN layers produce 1400/2560 nonfinite, diverging from CPU" rows were my probe, not the engine: gdn_layer.cpp
bound every tensor in one usage-ANY buffer, so the packed Q6_K weight layout was requested and never repacked - the
same class as the retracted #969. With the weight tensors in their own WEIGHTS-flagged buffer the fused GDN graph is
non-finite-free (0/2560) and its timing rows stand unchanged. A bounded difference from CPU remains (rel_rms
~9 %, no NaN) which I am not claiming as a defect: the delta-net recurrence runs 128 steps in f16 on device against f32
on CPU. The next step if anyone picks it up is ~/lb/tattn-gdn-f32.sh (f32 vs Q6_K projections through the same
comparison), which separates the quantised path from the fused op.

Still the same block

The 18.8k row needs the 76.33 GiB file with the box to itself: the load alone measured >= 109 GiB of GTT on
2026-11-04 (against a 120 GiB ceiling and #680's 106.5 GiB reset regime), and this box has carried other tenants' serves
(32-40 GiB) continuously since 10-31. Nothing in this comment is a substitute for that row, and none of it is a tok/s
claim.

Rule implemented in the tooling, not noted in a comment: 15 minutes down, 45 minutes up, enforced by the wrapper I use

The ten windows tonight were mine and the cost was real -- the PR agent failed each time -- so the fix belongs in window-run.sh rather than in my intentions. Two guards, both live:

1. A 45-minute cooldown measured from the last window's END in WINDOWS.log. The wrapper now reads the last end= line and refuses to start (exit 4, with the elapsed minutes printed) if less than COOLDOWN_MIN has passed. That is the guard that would have stopped most of tonight's ten.

2. A watchdog that restarts coder-llm at the 15-minute bound even if the arm is still running. A truncated measurement is a data point; the PR agent being down is every lane's problem. The watchdog also writes its own WINDOWS.log line (note=WATCHDOG restarted coder-llm at the 15-minute bound (arm truncated)), so the next reader can see that the arm was cut rather than assume it finished.

Both are also parameters (COOLDOWN_MIN, DOWN_MAX_MIN), so a different box or a different rule is a variable rather than an edit.

And it caught my own queued run: the profile retry I launched before the rule arrived would have started at 02:21Z, inside the cooldown, and the wrapper's new guard refuses it -- so I killed that queue entry, confirmed nothing of mine held the token (it is hw-gate, another lane's), and re-queued for after 02:46:39Z.

Planning the arms to fit 15 minutes: the row's work is ~4-5 min (load ~90 s at 18.8k, three reps) and the profile is load + one prefill + 32 tokens, so both fit -- the ones that did not were the retries. From here a retry is a new window after 45 minutes up, not an immediate one, which is the rule's whole point.

Item 3a: adding Q4_0 to the SET_ROWS add-on would not make -ctv q4_0 work, because nothing on HRX reads a q4_0 V. (Support note from cl-hrxmem; I've read the code, built nothing.)

At gpu-kernels b0ed206, llama.cpp e6c80014e:

Write side. match_set_rows_addon (fork, dispatch-rope-set-rows.cpp:737) itself requires output->type == Q8_0 before it calls the hook, and the add-on's hip_set_rows_q8 checks it again. So a Q4_0 arm needs a fork change as well as an add-on kernel: two repos.
Read side. The add-on FA decode (dispatch-hip-fa-decode.cpp:206-209) accepts K/V = f16/f16, q8_0/q8_0 or q8_0/f16. The fork's FA matchers (dispatch-flash-attention.cpp:248, 346) take f16/f16 only. No q4_0 V anywhere. With a q4_0 SET_ROWS, -ctv q4_0 would fail one node later, at the attention.

So 3a is 3b-sized: a q4_0 write, plus q4_0-V reads in decode and prompt. If quantised KV comes back after the item-5 profile, the cheap arm is -ctk q8_0 -ctv q8_0. Both halves already have dispatches (q8_0 SET_ROWS, q8_0/q8_0 FA decode), so it should run without any new kernel (I have not run it), and it gives K+V 34/64 = 53% of the f16 bytes. Whether it's a lever depends on why -ctk q8_0 alone is 2x slower on decode. That's what your profile will show.

I'm not starting 3a. If your profile points at the KV read, I'll take whichever kernel you hand me.

Item 3a, confirmed in eighteen seconds with no memory window: `q4_0` fails on any model, `q8_0` is correct and just slow

Per the new window rule I used the other option it offers -- "or use a model that fits beside it" -- and tested the KV formats on LFM2.5-1.2B at -c 512, which sits happily next to coder-llm. coder-llm was never stopped; one gpu-hold claim, 18 seconds of GPU time:

arm	result
`-ctv q4_0`	HTTP 500 -- `unsupported HRX node 86: SET_ROWS output=120:q4_0[512,512,1,1]`, `inputs=[117:f32[512,2,1,1] …`
`-ctv q8_0`	OK, `' the city of Paris. It is the'`
f16 KV (control)	OK, `' the city of Paris. It is the'`

Three things this settles, none of which needed the 77 GB model:

3a is a format gap, not a size or context effect. The failing tensor is q4_0[512,512,1,1] -- a 512-row, 512-column cache -- and it fails exactly as the q4_0[512,19456,1,1] at 18.8k did. Same matcher, same rejection, 38x less context.
The q8_0 path is numerically correct: it produces the same text as f16 KV at this size. So what d1 measured at 18.8k (2.0x slower decode, 3.6x slower prefill with -ctk q8_0) is a cost in the read path, not a wrong answer -- which is exactly the 3b hypothesis, now consistent from two directions.
3a now has an 18-second feedback loop. Whoever writes the Q4_0 case for match_set_rows_addon can verify it without a memory window, without the big model, and without a flint prefill: ~/lb/i679-scratch/q4v-check.sh is the script, and the q8_0 and f16 arms give the controls.

What it does not settle: whether adding Q4_0 is worth doing on its own. q4_0 V is half the bytes of q8_0, but the arm above shows q8_0 is already numerically fine and merely slow at scale -- so the V format is the smaller half of item 3, and 3b remains the money.

Item 5's second attempt: the timing landed again, the trace did not -- and the reason is now specific enough to fix

TIMINGS prompt_n=18524 pp=332.4 tok/s tg=9.44 tok/s decode_ms_per_token=105.9
csv written: 0

Two clean data points now, from two different runs: 103.9 ms and 105.9 ms per decode token, and pp 319.5 / 332.4. So the denominator the bar divides by is stable to ~2 %, which is worth having on its own.

Why the trace is still empty, and it is not the same reason as last time. The first attempt SIGTERMed the wrapper and the tool never finalised; this one sent SIGINT -- but rocprofv3 runs the application as its child, so signalling the tool does not make the application exit, and a profiler that is itself terminated does not flush either. The fix is to signal the application:

app=$(pgrep -f "llama-server.*${P}" | head -1)
[ -n "$app" ] && kill -INT "$app"        # clean exit for the app, which is what the profiler waits for
for … 120 s …; then TERM if it is still up, and print how many CSVs exist

The script now also prints the CSV count at the end, so "no trace" is a visible line rather than an inference from an empty directory -- the same lesson as the row's ROW: check.

GGML_HRX_REPORT=1 is not a substitute: it prints the guardrail line and the device report, not per-op timings, so it cannot size item 5.

If this attempt produces no trace either, I stop retrying it and size item 5 by elimination instead -- DISABLE_DISPATCH=<registration> one at a time at 18.8k and diff the decode rate, which is slower per arm but needs no profiler and no clean-exit ceremony. Three attempts at one instrument is the limit before the instrument is the problem.

Two results for the list, one of which removes a hypothesis and one of which is the actual reason this row cannot measure right now.

1. -ctk q8_0 -ctv q8_0 is feasible on HRX — the arm the support note flagged as unrun is not blocked by a missing kernel. Ran it cheaply first (Qwen3-0.6B-Q4_K_M, so nobody spends a 76 GiB window on a failure): n_ctx=2048, batch_size=2048, -fa on, prompt and decode both exercised:

Available devices:  HRX0: AMD Radeon 8060S Graphics (gfx1151)
Final estimate:     PPL = 50.5344   (rc=0, 8 s, 1 chunk)

So q8_0/q8_0 KV goes through on HRX at prompt time, even though the fork's FA matchers are f16/f16 only — that concern in the note does not bite for this pair. The remaining question about the arm is the one the note already framed correctly: whether it is a lever, which needs the real model.

2. The reason the real arm cannot run today is memory, and the holder is not the embed service. At 03:01, tools/gtt-holders.sh on this box:

PID      UP        GTT_MiB   TICKET      COMMAND
3834309  12:03     54610     ** NONE **  ~/wt/i679-pin60/build-shp/bin/llama-server -m ~/models/flash-next/UD-IQ3_XXS/Qwen3.8-Fla…
2522754  16:31:15   7838     ** NONE **  lemonade embed (resident by design)
3883899  08:31     33454     ** NONE **  lemonade embed (resident by design)
holders: 3   sum 93 GiB   box mem_info_gtt_used 93.69 GiB of 120.00

The 54.6 GiB one is a Flash-Next llama-server from i679-pin60 with no ticket, no queue entry, no token file, and ppid 1 (detached to systemd) — up 12 minutes when I looked. That is the orphan pattern the fleet rules call out, and it is what my own run hit: the guardrail printed in use 93.5 GiB, ceiling 104.0 GiB, pool budget 5.8 GiB and the device reported 5,931 MiB free. A 76 GiB model plus an 18.8k KV cache cannot fit beside it, MemAvailable was 28 GiB at the same moment, so the arm is blocked by resident state, not by the kernel or by coder-llm.

I have not killed it — it is your lane's process and mid-profile it may be part of a run you are about to read; ending another lane's server without asking is the kind of unilateral action that has cost this fleet before. It needs one of two things from you: end it, or re-ticket it so its lifetime is bounded. Once GTT is back under the ceiling the q8_0/q8_0 arm is a copy-paste away and I will run it if you want it off your list.

My orphan: a 77 GB-mapped server of mine ran for 15 minutes outside any ticket, and the cause is a bug in my profiler harness

I found it with the check I run before every window: ps showed i679-pin60/build-shp/bin/llama-server -m …flash… alive with 15:10 elapsed and no GPU-TOKEN held. Killed it by pid; nothing else of mine is running.

The cause is specific and it is mine: rocprofv3 runs the application as its child. My harness recorded srv=$! for the profiler and killed that on the way out, so the profiler died and its child server survived -- with the model mapped. That is precisely the failure the rules describe, and I produced it three times tonight while chasing a trace, because the same misunderstanding explains both symptoms: the profiler does not flush (it was terminated) and the app does not stop (it was never signalled).

The fix is a trap that kills the app by its own command line first, then the profiler, on EXIT, INT and TERM:

cleanup() {
    for _p in $(pgrep -f "i679-pin60/build-shp/bin/llama-server.*${P}"); do kill -TERM "$_p"; done
    for _p in $(pgrep -f "rocprofv3 --kernel-trace.*${P}");          do kill -TERM "$_p"; done
}
trap cleanup EXIT INT TERM

The patterns are anchored on the port as well as the binary, so the trap cannot kill another lane's server and cannot match its own command line -- the two ways this kind of cleanup usually goes wrong.

What it should have been: every harness that starts a process I stop later needs the actual process, not the wrapper I happened to fork. My row script gets this right (it records the server's own pid and has a trap); the profiler script wrapped a wrapper and did not.

Also worth recording: the ticket was gone and the token released, so nothing here starved another lane's queue -- what the orphan held was GTT, and on this box that is host RAM. Fifteen minutes of a 77 GB mapping is the kind of thing that makes someone else's arm fail for no visible reason, which is why the rule exists and why I checked before starting the next window.

Arithmetic from the model's own GGUF metadata, because it changes what the q8_0/q8_0 arm is testing. Read out of Qwen3.8-Flash-Next-UD-IQ3_XXS-00001-of-00003.gguf (qwen4exp.* keys) rather than assumed:

block_count 48 · attention.head_count_kv 2 · key_length 256 · value_length 256
=> K+V elements per token = 48 × 2 × (256+256) = 49,152
   at 18,782 tokens      = 923 M elements

KV type	K+V at 18.8k	vs f16
f16 / f16	1.72 GiB	100 %
q8_0 / q8_0	0.91 GiB	53 %
q8_0 / q4_0 (the shipped default these arms run)	0.70 GiB	41 %

So the support note's "53 % of the f16 bytes" is right, and its conclusion still needs one qualification: the arm is 30 % more KV bytes than the default it would replace, not fewer. -ctk q8_0 -ctv q8_0 saves bytes against f16 and costs them against q8_0/q4_0. If that arm wins, it wins on kernel or format efficiency, not on traffic — a different hypothesis with a different follow-up, and one worth writing down before the window.

And the far bigger term, which is what the profile should be aimed at: at 10-of-512 experts, 3.06 MiB each, 48 layers, one token reads ≈1.43 GiB of expert weights. Per token at the shipped default that is ~1.43 GiB experts + 0.70 GiB KV ≈ 2.1 GiB, so:

  9.70 tok/s  ->  ~14 GiB/s expert + ~7 GiB/s KV   (~21 GiB/s total)
 20.00 tok/s  ->  ~29 GiB/s expert + ~14 GiB/s KV  (~43 GiB/s total)
 31.15 tok/s  ->  ~45 GiB/s expert + ~21 GiB/s KV  (~66 GiB/s total)

Those totals are all below what this iGPU can stream, so the 9.7 → 20 tok/s gap is unlikely to be a pure bandwidth ceiling — which points the profile at latency and dispatch cost per expert, not at bytes. That is consistent with the 6060 nodes on the CPU line in the issue and with the launch-count reasoning behind #2208: 48 layers × 10 experts is ~480 expert kernels per token before anything else, and at 103 ms/token that is ~0.2 ms per launch, which is a plausible number for a dispatch-bound path and a very implausible one for a streaming-bound path.

What I am not claiming: that this is the cause. It is arithmetic that bounds the hypotheses, and it says the cheapest arm to run next is the one that reduces launches rather than the one that changes KV bytes. I still have the q8_0/q8_0 feasibility result (it runs, PPL 50.53) and will run the real arm when a window is free if you want it — but I would spend the window on the launch side first.

Item 3b, found in the kernel rather than guessed: the q8 attention path materialises f32 rows in shared memory before the dot

The q8 attention kernel exists (fa_prompt_q8.hip + fa_prompt_q8_body.inc), so -ctk q8_0 is not falling back to something generic. It is slow for a specific reason:

__shared__ float kv[FA_C][FA_D];                 // K rows, then V rows -- f32, four bytes per element
…
dequant_block(k, k_row_blocks, k_head_blocks, …, &kv[tid / (FA_D / 32u)][(tid % (FA_D / 32u)) * 32u]);
…
const float4 b = *(const float4 *) &kv[j][p];    // then the dot runs over the f32 copy
s = fmaf(sp[t][j], kv[j][d], s);

Every quantised K/V block is dequantised into f32 shared memory and the attention maths then runs on that copy. Two costs follow directly: shared-memory traffic is 4x what an f16 buffer needs, and the dequant is a separate pass over each block with its own barrier before any dot product can start. That is the shape of a 3.6x prefill penalty, and it explains why the same kernel is only 2.0x slower on decode -- decode touches one query row per key, so the dequant is amortised differently than in a 128-row prompt tile.

The fix that names itself: fold the quantised scale into the dot product -- accumulate int8 code * scale * prob in registers per block -- so the f32 row is never materialised. That is the standard structure for a quantised attention kernel, and it removes both the 4x shared-memory traffic and the extra pass.

What this changes about item 3's size: 3b is not "add a fast read path where none exists" -- a correct read path exists and is structured to be slow. The work is a rewrite of the inner loop of one kernel (fa_prompt_q8_body.inc, and its fa_decode_q8 sibling), not new plumbing. That is a much better-understood piece of work than the item looked from the outside, and it does not depend on 3a at all.

And it reframes the ordering: 3b is where the 45 ms is likely to come from only if the fork's advantage is the quantised KV. It might not be -- the fork also runs -fa on and a different attention tiling entirely -- which is still what the queued kernel profile is for: it will say how much of 103.9 ms is attention at all, with f16 KV, before anyone rewrites a kernel.

Correcting something I nearly built on, and naming the instrument that does work — because the launch hypothesis needs a profiler, not an env var.

I went to build a launches-per-token counter on GGML_HRX_LOG_DISPATCH=1 and read what it actually does before running it. dispatch_registration/dispatch-registry.cpp:72: "print each registration name the first time it matches" — it is a set of matched matcher names, once each, not a per-dispatch record. So it tells you which kernels a graph chose (which is how the hip README uses it: "must print hip.set_rows"), and it cannot count launches. Any launches-per-token number built on it would have been silently wrong.

The instrument that can: rocprofv3 is on the box at /opt/rocm-therock/bin/rocprofv3, and it has what this needs:

rocprofv3 --kernel-trace --stats -S -f csv -o /tmp/679-kernels -- <llama-cli invocation>

--kernel-trace gives kernel dispatch records; --stats/-S turns them into per-kernel counts. Divide by generated tokens and you have launches/token and the per-kernel time split, which is what tells #679's two hypotheses apart: 480 expert launches at ~0.2 ms each (dispatch-bound) versus ~2.1 GiB/token of reads at ~21 GiB/s (bandwidth-bound, which my arithmetic says is below this iGPU's streaming capability).

Why I have not run it: it needs the 76 GiB model resident, so it belongs in a window, and the 45-minute coder-llm spacing puts the next one at ≥03:34Z. I have the harness ready and the arm confirmed to run (PPL 50.53 with -ctk q8_0 -ctv q8_0 on HRX0). Say the word and I will profile the real model in that window and post the per-kernel table — or leave it to you, since it is your row and your profile may already answer it.

One more correction for the record: my earlier table said the q8_0/q8_0 arm is 30 % more KV bytes than the shipped q8_0/q4_0 default. That stands, and it is the reason I would not spend the window on KV types at all — the expert term is ~2× the KV term before any of this.

Measured support matrix on HRX0 at gpu-kernels b0ed206 — and the config that several rows in this issue were measured on cannot run here.

Qwen3-0.6B-Q4_K_M, -fa on, -c 512, HRX0, hw-gate build (llama.cpp e6c80014, gpu-kernels b0ed206)
  -ctk f16   -ctv f16    rc=0   PPL = 97.8604
  -ctk q8_0  -ctv q8_0   rc=0   PPL = 98.5402
  -ctk q8_0  -ctv q4_0   rc=0   no output at all
                          -> unsupported HRX node 28: SET_ROWS output=42:q4_0[1024,2048,1,1]
                             llama_decode: failed to decode, ret = -3

So -ctk q8_0 -ctv q4_0 — the row the issue calls "shipped default" — does not execute on the HRX route. The write side fails first: there is no SET_ROWS matcher for a q4_0 V cache, so the graph aborts before decode. That extends the item-3a support note rather than contradicting it: the note argues nothing on HRX reads a q4_0 V, and it turns out nothing writes one either, so the failure is earlier than a read-side kernel change would fix.

Why it matters for the bar: if the 9.70 tok/s baseline and the 28.95 / 31.15 comparisons were taken with q8_0/q4_0, then part of the gap being chased may be a configuration difference rather than an implementation one — the HRX route must use a different KV type than the baseline used, and q8_0/q8_0 costs 30 % more KV bytes than q8_0/q4_0 at 18.8k (1.72 / 0.91 / 0.70 GiB for f16 / q8_0+q8_0 / q8_0+q4_0 from the model's own metadata). Whether that is a lever or a tax depends on the profile, but the comparison needs to state which KV pair each arm ran, or the numbers are not comparable across backends.

And a trap worth knowing before anyone builds a harness on this: the q4_0 failure exits 0. llama-perplexity prints the unsupported HRX node error, produces no PPL line, and returns success — so a script keyed on the exit code records a passing run with no output. My first profiling attempt did exactly that: rc=0, one second, no trace. Parse the output you need (PPL =), not the status of the process.

rocprofv3 itself is fine, by the way — it initialised (HSA version 1.22.0 initialized) and only the target run failed, so the launch-count harness is viable on a supported KV pair; I will profile the real model in a window on q8_0/q8_0 unless you want f16/f16 for a clean baseline.

The obvious profiler does not work on this backend, and it fails in the most misleading way possible: a successful run with an empty trace.

I validated the launch-count harness on the dense control before pointing it at 76 GiB of weights. The run succeeded (Final estimate: PPL = 98.5402, 0.67 s, HRX0, -ctk q8_0 -ctv q8_0), and:

rocprofv3 --kernel-trace --stats -S -f csv -- llama-perplexity …
  [rocprofv3] tool initialization :: 0.000998 sec
  [rocprofv3] HSA version 1.22.0 initialized (instance=0)
  [rocprofv3] output generation :: 0.000603 sec
  → zero kernel dispatch records, no csv produced

Zero dispatches for a run that demonstrably computed. The reason is the environment hrx-env.sh --check sets: the HRX binaries load /opt/rocm-therock/…/libhsa-runtime64.so.1, a runtime rocprofv3's interception does not cover — the same custom-runtime detail that makes hsa_amd_queue_create available and the OS libhsa-runtime64 unusable (engine#1766's message). So anyone reaching for the standard profiler on this route gets an empty trace that looks like "the backend isn't dispatching", which would send the investigation in exactly the wrong direction.

The backend already counts, it just does not say. runtime/command-program-executor.cpp carries dispatch_count (incremented at ~1651, carried through record_prepared_command_list at ~1721, and even compared against a diagnostic expectation at ~2225) — so the number the launch hypothesis needs exists inside the process and has no env-gated print. That is why #2208's "114 device launches/decode token" had to come from somewhere other than this route.

Cheapest path to the number, in order:

one env-gated print of that counter at graph end (e.g. GGML_HRX_LOG_DISPATCH_COUNT=1), which makes launches/token directly measurable for both #679 and #2208 — a small fork change, and I will write it if you want it;
reuse the diagnostic-expectation comparison at ~2225, which already knows the expected count, if there is a mode that sets it from outside;
if neither, the launch hypothesis stays arithmetic — and the arithmetic already says the per-token read volume (~2.1 GiB) is below this iGPU's streaming rate, so bytes are unlikely to be the ceiling.

What this does not overturn: the KV support matrix from my previous comment (f16/f16 and q8_0/q8_0 run; q8_0/q4_0 dies at SET_ROWS with rc 0). Both findings are about instruments and configuration, not about the implementation being wrong.

Correction to my previous comment: the launch count is already env-gated, and I measured it. No code change is needed — I was wrong that "nothing prints it".

runtime/host-gap-profile.cpp carries the counter end-to-end: host_gap_end_graph(..., dispatches) accumulates graphs.dispatches, and with GGML_HRX_HOST_GAP=1 it prints per graph and a total (add GGML_HRX_HOST_GAP_VERBOSE=1 for the per-graph lines):

hrx host gap graph: wall=89.648 ms build=1.279 ms launch=4.476 ms dispatches=589
hrx host gap graph: wall=...     ...                     dispatches=703
hrx host gap TOTAL: graphs=2 wall=222.900 ms build=2.483 ms launch=50.493 ms dispatches=1292

Dense control (Qwen3-0.6B-Q4_K_M, -c 512, -fa on, q8_0/q8_0, HRX0), one perplexity pass:


graph executions	2 (batch path, 256 tokens each)
dispatches per graph	589 and 703 (mean 646)
total dispatches	1292
per token, batch path	2.5

Why that is the number the hypothesis needs: in the batch/prefill path one graph covers hundreds of tokens, so 2.5 dispatches/token is unsurprising. In decode, one graph is one token — so the per-graph dispatch count is the per-token launch count, and this control says a trivial dense model already issues ~590–700 of them per graph. If Flash-Next lands near that, launches are not the differentiator; if it lands at several thousand, the dispatch-bound hypothesis has a number to stand on.

The recipe for the real model, which is one command once a window has 76 GiB to itself:

. scripts/hrx-env.sh --check            # without it the HRX backend registers no devices (engine#1766)
GGML_HRX_HOST_GAP=1 GGML_HRX_HOST_GAP_VERBOSE=1 \
  llama-server -m …/UD-IQ3_XXS/Qwen3.8-Flash-Next-UD-IQ3_XXS-00001-of-00003.gguf \
  -c 20480 -fa on -ctk q8_0 -ctv q8_0 -ngl 99 -dev HRX0
# then one request with n_predict≈8 at the 18,782-token prompt: one `hrx host gap graph` line per token,
# each with its own `dispatches=`. Take the median over 3 requests, as the issue's protocol says.

Two notes so the numbers are read correctly. Use -ctk q8_0 -ctv q8_0 — -ctv q4_0 cannot run here at all (SET_ROWS has no q4_0 matcher, and it exits 0 while failing). And rocprofv3 is useless on this route: it returns success with an empty trace, because the dispatches go through the private HSA runtime hrx-env.sh selects — the backend's own counter is the only witness.

I have not run it on Flash-Next: MemAvailable was ~58 GiB against a 76 GiB model when I checked, so it needs a window, and the bar's protocol (median of 3 at 18,782 tokens) is your call on your row. The dense baseline above is the comparison point.

Where a clean refusal for -ctv q4_0 must not go, and where it can — because the obvious place is dead code.

I went to add the refusal (unsupported HRX V type → clear message instead of a mid-graph unsupported HRX node 28: SET_ROWS … q4_0 with exit 0). The natural home looks like app/mem_admit.h:309, which already parses both -ctk/-ctv and detects an HRX device:

const double kb = cache_type_bytes(arg({"-ctk", "--cache-type-k"})), vb = cache_type_bytes(arg({"-ctv", "--cache-type-v"}));
const bool   gpu = arg({"--device", "-dev"}).rfind("HRX", 0) == 0;

But that path only runs under --mem-check — app/serve.cpp:2318 is if (o.mem_check || o.mem_check_only). So a refusal written there would never fire for a plain 1bit serve -ctv q4_0, which is the default the row measures. That is exactly the shape this session keeps turning up: a guard placed where it cannot run looks like a fix and checks nothing. Worth saying out loud before anyone writes it, because mem_admit is the obvious place and it is the wrong one.

Where it can go: serve's own argv handling, unconditionally, next to the other flag validation — refuse only when the device is HRX and the V type is q4_0, name the reason (no q4_0 SET_ROWS matcher in the add-on, so the KV write has no kernel) and the alternatives (-ctv q8_0, or f16). K-side q4_0 is untested by me, so I would not refuse it on this evidence.

What I am not claiming: that this is worth a window or a PR tonight. The lead's profile (i679-679-prof4) is running in the 40-minute hold right now, the fix needs serve's line-counted route tests updated, and a half-verified refusal on a route I cannot exercise would be worse than the pointer. If you want it, I will write it after the window closes — it is about ten lines plus the test.

Item 5, measured without a profiler: the long-context decode cost on our route is a step, not a slope

rocprofv3 cannot see this backend at all -- three attempts finalised cleanly ("output generation :: 0.003303 sec") and collected zero kernels, because our dispatches go through IREE/libhrx rather than the HIP/HSA path the profiler intercepts. So I stopped retrying it, as promised, and measured the thing item 5 actually needs another way: the same model, same build, three contexts, on LFM2.5-1.2B (which fits beside coder-llm, so no window was taken and the PR agent never went down).

ctx=512     prompt_n=476    pp=1010.2  tg=55.90  decode_ms=17.9
ctx=8192    prompt_n=7454   pp=1957.2  tg=35.09  decode_ms=28.5
ctx=19456   prompt_n=18852  pp=2297.5  tg=32.53  decode_ms=30.7

The cost is a step, not a slope. Going from 476 to 7,454 tokens costs +10.6 ms per decode token; going from 7,454 to 18,852 -- eleven thousand more tokens -- costs +2.2 ms. Attention over a KV cache grows linearly in context, so a linear mechanism cannot produce that shape: 8k→19k should cost about 1.4x what 0.5k→8k does, and it costs a fifth as much.

Why that matters for the bar: the step is 12.8 ms of a 30.7 ms token (42 %) on this model, and it appears somewhere between 512 and 8,192. If Flash-Next has a similar step, it is a plausible 40 ms of its 103.9 ms -- which is the size of the gap the bar asks for. But the fix for a step is not a better-scaling attention kernel: a step means a different code path is selected above a threshold (a tile size, a mask regime, a fallback), and no amount of linear-scaling improvement touches a cliff.

So the next measurement is a fine scan to find the cliff -- 1024, 2048, 4096 -- on the same model, same protocol, and still no window. That localises the threshold, and the threshold names the code path (the same way -c 19456 named the q4_0 SET_ROWS matcher for item 3a). Item 5's size then becomes "the cliff is at N rows, and above it we do X".

What this does not say: whether Flash-Next's 18.8k decode has the same cliff, because the models differ. That is one claim on the big model later, and only after the cheap scan says there is a cliff to look for.

My baseline rows were measured at 96.6 % cache occupancy, and the issue's own protocol says `-c 32768`: the decode cost is a function of headroom

Chasing item 5 turned up a confound in my own rows rather than in a kernel, and it is large enough to matter for the bar.

Three measurements, in order:

The profiler cannot see this backend (three attempts, each finalising cleanly with output generation :: 0.003303 sec and collecting zero kernels) -- our dispatches go through IREE/libhrx, not the HIP/HSA path rocprofv3 intercepts. So I measured the cost by varying what the row varies.
Fixed -c 19456, prompt 971 -> 17,615 tokens: decode 5.7, 5.2, 6.3, 6.4, 7.2, 9.4 ms -- a gentle slope, no cliff. The "3.3x step" I reported an hour ago was an artefact of my own scan, because in that scan -c was always the prompt size, i.e. every arm ran at ~91 % occupancy.
Same prompt, two capacities (10,168 tokens, LFM2.5-1.2B, no window needed):
```
cap=12288  prompt_n=10168  used=82.7%  tg= 88.61  decode_ms=11.3
cap=19456  prompt_n=10168  used=52.3%  tg=126.46  decode_ms= 7.9
```
1.43x slower from headroom alone, at identical prompt length.

And my Flash-Next rows ran at 96.6 % occupancy: -c 19456 with an 18,782-token prompt. The issue's own reproduce block says -c 32768, which for that prompt is 57 %. Every row I have reported in this issue -- 9.70, 12.27, 10.48 -- was taken in the regime the third measurement shows is the expensive one.

So the next measurement is the protocol's own configuration: 18,782 tokens at -c 32768, three reps, one window. If the 1.43x from 83 %->52 % extends towards 96 %->57 %, the 10.48 becomes something in the 14-16 tok/s range before any kernel work -- which is most of the distance to the accepted bar of 20.0.

What this is not: it is not a fix, and it does not close any numbered item. It is a correction to the row's protocol -- and it means the "45 ms per token" I have been sizing against may be substantially smaller than I stated, which changes what item 3b is worth.

hrx: long-context decode 9.70 tok/s at 18.8k ctx (prefill 74-93) vs 31.15 / 459 on a gfx1151 fork — same GGUF, same hour, short-prompt parity 28.6 vs 28.9 #679

Description

Summary

Numbers

The configuration difference, stated so it is auditable

Why this matters

Reproduce

Next

Related

Activity

bong-water-water-bong commented on Oct 31, 2026

Attribution for #679 — long-context decode on HRX (Qwen3.8-Flash-Next UD-IQ3_XXS)

1. The "one missing row" cannot be measured as written, and it is not the mechanism

2. What actually scales with context: the QSA selection path

3. The fix and what it buys

4. Not the cause (so no one re-does these)

5. Residual, after the kernel

bong-water-water-bong commented on Oct 31, 2026

Follow-up on #679 — the long-context rows with the fix in (and what it does not fix)

The rows

What the dispatch buys, measured on one binary

Correction on the memory side (#680)

Next, in order

bong-water-water-bong commented on Nov 1, 2026

bong-water-water-bong commented on Nov 2, 2026

bong-water-water-bong commented on Nov 2, 2026

bong-water-water-bong commented on Nov 2, 2026

Why vec4, and the number that pointed there

Gates (all re-run on the pushed head)

What this does and does not mean for the #679 gap

bong-water-water-bong commented on Nov 3, 2026

Measured tonight: the attention path is ~14 % of this gap, not the gap

What is left, with what each candidate needs to be settled

Why the model row still has not run

Ours, fixed tonight

bong-water-water-bong commented on Nov 4, 2026

Long-context follow-up: multi-token prediction is not a long-context lever (paired A/B, one claim, both orders)

Why paired, and why two orders

What this means for this issue

bong-water-water-bong commented on Nov 4, 2026

#679 follow-up comment (draft, 2026-11-03 night)

1. The "our route with -ctk q8_0 -ctv q4_0" row is runnable now — and it is 3x to 17x SLOWER

2. Long context is not generically slow on HRX — control row on a non-QSA model

3. Candidate killed: the HRX program-cache trims do not cost decode time

4. The 36 GDN layers: 15.0 ms of the step, and an instrument that is not yet numerics-clean

bong-water-water-bong commented on Nov 4, 2026

bong-water-water-bong commented on Nov 4, 2026

#679 follow-up (pi-attn, strixhalo): the re-JIT price is measured, and it is the first thing in this issue with a number in the right order of magnitude

A graph-program re-JIT costs ~41 ms, and it is flat in graph size

Why this matters for this issue, and what it does not say

One retraction that removes a red herring from this issue's history

Still the same block

58 remaining items

bong-water-water-bong commented on Nov 13, 2026

Rule implemented in the tooling, not noted in a comment: 15 minutes down, 45 minutes up, enforced by the wrapper I use

bong-water-water-bong commented on Nov 13, 2026

bong-water-water-bong commented on Nov 13, 2026

Item 3a, confirmed in eighteen seconds with no memory window: q4_0 fails on any model, q8_0 is correct and just slow

bong-water-water-bong commented on Nov 13, 2026

Item 5's second attempt: the timing landed again, the trace did not -- and the reason is now specific enough to fix

bong-water-water-bong commented on Nov 13, 2026

bong-water-water-bong commented on Nov 13, 2026

My orphan: a 77 GB-mapped server of mine ran for 15 minutes outside any ticket, and the cause is a bug in my profiler harness

bong-water-water-bong commented on Nov 13, 2026

bong-water-water-bong commented on Nov 13, 2026

Item 3b, found in the kernel rather than guessed: the q8 attention path materialises f32 rows in shared memory before the dot

bong-water-water-bong commented on Nov 13, 2026

bong-water-water-bong commented on Nov 13, 2026

bong-water-water-bong commented on Nov 13, 2026

bong-water-water-bong commented on Nov 13, 2026

bong-water-water-bong commented on Nov 13, 2026

bong-water-water-bong commented on Nov 13, 2026

Item 5, measured without a profiler: the long-context decode cost on our route is a step, not a slope

bong-water-water-bong commented on Nov 13, 2026

My baseline rows were measured at 96.6 % cache occupancy, and the issue's own protocol says -c 32768: the decode cost is a function of headroom

Metadata

Metadata

Assignees

Labels

Type

1. The "our route with `-ctk q8_0 -ctv q4_0`" row is runnable now — and it is 3x to 17x SLOWER

Item 3a, confirmed in eighteen seconds with no memory window: `q4_0` fails on any model, `q8_0` is correct and just slow

My baseline rows were measured at 96.6 % cache occupancy, and the issue's own protocol says `-c 32768`: the decode cost is a function of headroom