[jtsylve]

Summary

AMD clang 23 is the HIP compiler of ROCm 7.14 and ROCm 10.0. With clang 23, the
blocked WMMA prefill GEMMs of Qwen3.8-27B spill registers. One Flash-Next GEMM keeps a private
array in scratch memory. On current main (05688f9), a clang 23 build has these results,
compared with the published 0.8.1 image (clang 22):

gufo bench, tok/s	0.8.1 image (clang 22)	clang 23, before	clang 23, after
Qwen3.8-27B Q8 pp2048	454	291	464
Qwen3.8-27B Q8 pp64	322	112	308
Qwen3.8-27B Q8 pp16	93	30	92
Qwen3.8-27B Q4 pp64	327	230	335
Qwen3.8 Flash-Next Q4 pp2048	1598	1592	1603

Without this PR, a clang 23 build makes Q8 prefill of 2048 tokens 36% slower, and prompts of
9 to 95 tokens about 3 times slower. Decode does not change.

This PR removes the spills with source constraints. It does not use hidden compiler options or
a compiler version test. It also adds a kernel resource check: a CTest that fails if a kernel
uses more scratch memory than a baseline allows. Thus a later compiler cannot add a spill and
go unseen.

The greedy output does not change. On clang 22, the pinned toolchain, pp2048 and decode
change by 1.1% or less, and prefill of 16 to 64 tokens is 1.0% to 1.9% slower (see "Cost on
clang 22").

Changes

File	Change
src/models/qwen/hip/kernels/prefill_quant_gemm.hip	W8A8BlockedWmmaGEMMKernel: one __builtin_amdgcn_sched_barrier(0) at the end of each K block. The K-block loop uses #pragma unroll(BK <= 2 ? BK : 1).
src/models/qwen/hip/kernels/prefill_quant_gemm.hpp	WKQuantA8BlockedWmmaGEMMKernel: the same #pragma unroll(BK <= 2 ? BK : 1). No fence: the fence makes clang 22 spill in the Q3_K tile.
src/models/qwen/CMakeLists.txt	-fno-slp-vectorize for prefill_quant_gemm.hip and prefill_quant_wave64.hip, on all compilers.
src/models/qwen38_flash_next/kernels/rocm/kernels.hip.cpp	RoutedF16GEMMKernel: blend the four cached code chunks with masks. Before, a conditional over uint4 lvalues selected one chunk.
tools/ci/check-kernel-resources.py, tools/ci/kernel-resources.json	The kernel resource check and its baseline.
CMakeLists.txt	The CTest kernel_resources_test (labels cpu;kernels, HIP builds only).
docs/PERFORMANCE.md, docs/models/qwen3.8-27b/EXPERIMENTS.md, docs/models/qwen3.8-flash-next/EXPERIMENTS.md	How to use the check, and the retained results.

The arithmetic order does not change.

Causes

Kernels	Cause with clang 23	Compiler stage
W8A8 128x128 tiles (BK=2)	The scheduler moves the LDS reads of the next K block above the WMMAs of the current K block. These values stay live across the WMMA block.	Backend scheduler
W8A8 and WKQuantA8 BK=4 tiles (batches of 9-95 and 1-8 tokens)	The same, over four unrolled K blocks.	Backend scheduler
WKQuantA8 Q3_K and Q6_K 128x128 tiles	SLP vectorization puts the per-block scales and offsets into wide register tuples.	Optimizer (SLP)
Flash-Next RoutedF16GEMMKernel pair tiles	The conditional over uint4 lvalues becomes an address select, and code_cache stays in scratch. An if/else chain does not help: clang 23 merges it into one load through a pointer phi.	Optimizer

To find the stage of the W8A8 spill, we gave the optimized IR of clang 22 to the backend of
clang 23 (-Xclang -disable-llvm-optzns). The spill occurred again with the same numbers.

Why source constraints

A fix must continue to work with the next compiler. These constraints do not depend on one
compiler:

• __builtin_amdgcn_sched_barrier is a documented builtin.
• The scheduler does not move instructions across a loop iteration that is not unrolled.
• -fno-slp-vectorize is a documented driver flag.
• A mask blend has no address to select.

We rejected these alternatives. We measured them at 8a46cd9 with clang 23 (Q8_K_XL). The
Qwen3.8-27B kernel files are the same at 05688f9.

Alternative	Result	Problem
-mllvm -amdgpu-load-store-vectorizer=false	Q8 pp2048 465	DFlash2 decode 33% slower (24.3 to 16.4 tok/s). Hidden option.
-fno-slp-vectorize -mllvm -amdgpu-use-amdgpu-trackers=1	Q8 pp2048 434	Partial. Hidden option.
-fno-slp-vectorize alone (from #228)	Q8 pp2048 297	No effect on the W8A8 spill.
The #255 fence alone	Q8 pp2048 462, pp64 125	The BK=4 tiles still spill.
#pragma unroll 1 on all K-block loops	Q8 pp2048 437	The 128x128 tiles need 244 VGPRs (5 waves, not 8).
#if __clang_major__ >= 23 (as #255)	—	A version test guesses which compilers need the fix.

If a later compiler adds a spill, kernel_resources_test fails and names the kernel.

Results

All numbers are gufo bench -r 3 in tok/s, at 05688f9. Two rounds, builds alternated. The
machine and the builds are in "Test setup". The DFlash2 columns use the Q4_K_M draft. "Before"
is 05688f9; "after" is 05688f9 with this PR.

Qwen3.8-27B UD-Q8_K_XL, DFlash2:

Build	pp2048	pp64	pp16	pp8	tg64	DFlash2 pp2048	DFlash2 tg128
clang 22, published 0.8.1 image	455.4 / 453.1	322.8 / 321.5	93.3 / 93.1	50.9 / 50.8	7.17 / 7.16	431.7 / 433.5	24.47 / 24.49
clang 23, before	291.5 / 290.8	112.1 / 112.0	29.6 / 29.6	42.9 / 42.8	7.26 / 7.26	281.6 / 280.9	24.26 / 24.26
clang 23, after	465.4 / 462.8	308.4 / 307.8	92.0 / 91.9	51.1 / 51.1	7.26 / 7.26	439.5 / 438.6	24.26 / 24.23
clang 22 over ROCm 10, before	475.8 / 476.1	322.9 / 323.1	93.6 / 93.3	51.7 / 51.7	7.23 / 7.23	451.7 / 450.7	24.27 / 24.31
clang 22 over ROCm 10, after	480.3 / 480.4	319.1 / 319.1	92.3 / 92.1	51.4 / 51.4	7.23 / 7.24	457.0 / 455.3	24.31 / 24.29
clang 23, after vs before	+59.4%	+175.0%	+210.4%	+19.3%	+0.0%	+56.1%	-0.1%
clang 23 after vs 0.8.1 image	+2.2%	-4.4%	-1.4%	+0.5%	+1.3%	+1.5%	-1.0%
clang 22, after vs before	+0.9%	-1.2%	-1.4%	-0.6%	+0.1%	+1.1%	+0.0%

Qwen3.8-27B UD-Q4_K_XL, DFlash2:

Build	pp2048	pp64	pp16	pp8	tg64	DFlash2 pp2048	DFlash2 tg128
clang 22, published 0.8.1 image	649.5 / 648.0	325.6 / 328.0	88.1 / 87.2	65.2 / 65.1	12.29 / 12.28	605.1 / 603.2	26.77 / 26.76
clang 23, before	669.3 / 666.6	230.3 / 229.7	59.9 / 59.9	63.5 / 63.2	12.17 / 12.17	617.9 / 617.2	26.74 / 26.70
clang 23, after	669.3 / 667.5	334.4 / 335.7	89.0 / 89.3	68.3 / 68.1	12.17 / 12.17	617.5 / 616.7	26.73 / 26.72
clang 22 over ROCm 10, before	669.3 / 668.6	344.6 / 342.3	91.2 / 91.2	67.7 / 67.9	12.10 / 12.09	622.1 / 622.2	26.68 / 26.67
clang 22 over ROCm 10, after	669.9 / 668.6	338.0 / 337.6	89.6 / 89.8	68.5 / 68.5	12.09 / 12.09	621.9 / 622.0	26.69 / 26.66
clang 23, after vs before	+0.1%	+45.7%	+48.9%	+7.6%	+0.0%	-0.1%	+0.0%
clang 23 after vs 0.8.1 image	+3.0%	+2.5%	+1.7%	+4.7%	-0.9%	+2.1%	-0.1%
clang 22, after vs before	+0.0%	-1.7%	-1.6%	+1.0%	-0.0%	-0.0%	+0.0%

Qwen3.8 Flash-Next UD-Q4_K_XL, MTP:

Build	pp2048	pp64	pp16	pp8	tg64	MTP pp2048	MTP tg128
clang 22, published 0.8.1 image	1599.5 / 1596.3	363.6 / 363.4	139.4 / 138.1	88.4 / 88.3	27.05 / 27.03	1570.2 / 1567.8	44.94 / 44.96
clang 23, before	1595.5 / 1588.2	376.8 / 373.2	146.9 / 142.6	91.3 / 89.9	26.97 / 26.97	1554.1 / 1561.0	45.93 / 46.13
clang 23, after	1604.0 / 1601.7	380.9 / 383.1	146.9 / 147.2	93.1 / 93.0	27.00 / 26.98	1568.8 / 1572.8	45.99 / 45.96
clang 22 over ROCm 10, before	1616.4 / 1615.7	365.9 / 364.7	140.3 / 139.7	88.5 / 86.8	27.02 / 27.02	1583.7 / 1582.2	44.85 / 44.71
clang 22 over ROCm 10, after	1605.7 / 1602.6	365.2 / 351.7	139.9 / 137.2	88.4 / 87.4	27.03 / 27.00	1574.2 / 1566.6	44.90 / 44.73
clang 23, after vs before	+0.7%	+1.9%	+1.6%	+2.7%	+0.1%	+0.9%	-0.1%
clang 23 after vs 0.8.1 image	+0.3%	+5.1%	+6.0%	+5.3%	-0.2%	+0.1%	+2.3%
clang 22, after vs before	-0.7%	-1.9%	-1.0%	+0.3%	-0.0%	-0.8%	+0.1%

Each cell shows the two rounds. The rows in italics give the change of the mean.

With this PR, clang 23 prefill is faster than the 0.8.1 image in all columns except Q8 pp64
(-4.4%) and Q8 pp16 (-1.4%). Decode is within 1.3% of the image, except Flash-Next MTP tg128
(+2.3%). That difference comes from the text, not from the kernels: the drafts follow the text,
and clang 23 gives a different Flash-Next text (see "Numerics").

Q4 pp2048 was not slow before this PR. At pp2048, Q4_K_XL uses HalfPrefillGemmKernel for most
of its GPU time (rocprofv3 at 8a46cd9), and clang 23 compiles that kernel as well as clang 22.
The Q4 short prompts use the BK=4 tiles that spilled.

Cost on clang 22

Upstream builds with clang 22. Thus the "clang 22, after vs before" rows above are the cost of
this PR for the current toolchain:

• pp2048 (also with DFlash2 or MTP) and decode change by 1.1% or less on the three models.
• Prefill of 16 to 64 tokens is 1.0% to 1.9% slower.
• Flash-Next pp2048 is 0.7% slower, and MTP pp2048 is 0.8% slower.

The cause is in the register use (see "Register use"). With clang 22, four WKQuantA8 128x64
BK=4 tiles use more VGPRs without the K-block unroll: Q3_K (179 to 230 VGPRs, 8 to 6 waves),
and IQ4_NL, IQ3_S and IQ4_XS (9 to 8 waves). No kernel spills. Eleven kernels are better.

We chose constraints that do not depend on one compiler, and they have this small cost. If you
do not want this cost on clang 22, you can put the BK=4 unroll and the Flash-Next blend behind
a gate. The BK=2 fence does not change the occupancy of the W8A8 128x128 tiles on clang 22
(180 to 183 VGPRs, 8 waves), and Q8 pp2048 does not become slower. kernel_resources_test
protects the kernels in all cases.

The clang 22 build over the ROCm 10 userspace is faster than the published 0.8.1 image (Q8
pp2048 476 against 454, without this PR), and its output is the same. Thus the ROCm 10
libraries are not the cause of the clang 23 slowdown.

Numerics

Greedy output: the first 12 hex digits of the SHA-256 of 160 generated tokens
(gufo prompt -t 0, a 553-word prompt).

Model	0.8.1 image	clang 23, before	clang 23, after	clang 22, before	clang 22, after
Qwen3.8-27B Q8	1e504ca34bf6	same	same	same	same
Qwen3.8-27B Q4	821f519124de	same	same	same	same
Qwen3.8 Flash-Next Q4	ea43a269d4aa	4b063b5eeb2c	4b063b5eeb2c	ea43a269d4aa	ea43a269d4aa

• This PR does not change the output with either compiler.
• Qwen3.8-27B Q8 --validate-prefill 1024 gives the same result on all five builds
  (max_abs_diff 0.78307986, rmse 0.09793035, cosine 0.99944276, top-1 match).
• Clang 23 gives a different Flash-Next text than clang 22, before and after this PR. The
  Flash-Next unit compiles with -O3 -ffast-math, and clang 23 makes different arithmetic. A
  clang 22 build over the ROCm 10 libraries gives the text of the image. Thus the compiler causes
  the difference, not the libraries. A move to clang 23 needs the Flash-Next and DeepSeek V4
  Flash quality gates. This PR does not change that.

Kernel resource check

tools/ci/check-kernel-resources.py reads the .hip_fatbin section of a binary. It decodes
the AMDGPU metadata note of each code object. It needs no GPU. It needs no LLVM tool, except
clang-offload-bundler when the bundles are compressed and Python has no zstd. It runs in
about 1.5 seconds on the 228 MB release binary.

Command	Result
check-kernel-resources.py BINARY --baseline tools/ci/kernel-resources.json	Exit status 1 if a kernel uses more scratch than its allowance.
check-kernel-resources.py NEW --compare OLD	Lists each kernel with more scratch, more spilled VGPRs or less VGPR occupancy.
check-kernel-resources.py BINARY --report [FILTER]	Lists the kernels that use scratch.
check-kernel-resources.py BINARY --write-baseline FILE --note TEXT	Writes a new baseline.

The baseline comes from the release binary of the published clang 22 image. It allows the
scratch of the 247 kernels that use scratch in that build. Mangled kernel names are the keys,
and the tool removes the version from the rocPRIM inline namespace. Thus a ROCm update does not
change the key of most kernels. An allowance for a kernel that no longer exists gives a note,
not a failure. Results at 05688f9:

Build	Kernels	Kernels with scratch	Over the baseline
clang 22, published 0.8.1 image	2586	247	0
clang 22 over ROCm 10, before	2519	243	0
clang 22 over ROCm 10, after	2519	243	0
clang 23, before	2519	255	31
clang 23, after	2519	237	13

Thus kernel_resources_test passes on the current toolchain. On a clang 23 build it fails,
also with this PR: 13 kernels remain (see "Remaining clang 23 items"). That is intended. The
test shows the work that a move to clang 23 needs. A clang 23 build can keep its own baseline
(--write-baseline).

--compare of the "after" build with the "before" build:

• clang 23: 28 kernels are better. 2 are worse: the Flash-Next Q5_K pair tiles now use no
  scratch, but they lose one occupancy step (9 to 8 and 12 to 10 waves). The Flash-Next results
  include this effect.
• clang 22: 11 kernels are better. 4 are worse (see "Cost on clang 22"). No kernel spills.

Not verified

• nix build .#checks.x86_64-linux.pr and the model quality gates. The machine has no Nix.
  We ran tools/ci/check-docs.py, tools/ci/check-dependencies.py and
  tools/ci/check-pr-title.py (all pass), and clang-format --dry-run --Werror 21.1.8 on the
  three changed C++ files (clean).
• The speed on the full Nix ROCm 7.2.3 toolchain. We used its clang 22 over the ROCm 10
  userspace, and the published image.
• kernel_resources_test through CTest. We ran the same command by hand on all five builds.
• Speed of other models. Only the Qwen3.8-27B and Flash-Next code calls the changed kernels
  (checked in the source), so DeepSeek V4 Flash, ASR, TTS and image models do not change.

Related work

• perf(qwen-flash): fence each K block in the W8A8 WMMA GEMM #255 (merged) added the same K-block fence to the Flash-Next copy of
  W8A8BlockedWmmaGEMMKernel, with #if __clang_major__ >= 23. The Qwen3.8-27B copy did not
  get the fence. This PR adds it without the version test.
• ROCm 10 evaluation for gfx1151, and a clang 23 prefill fix #228 (open) evaluates ROCm 10. Its -fno-slp-vectorize change is in this PR, for two units
  and on all compilers. Its roundf to __float2int_rn change is not in this PR, because it
  changes the numerics. A comment on ROCm 10 evaluation for gfx1151, and a clang 23 prefill fix #228 asks for a ROCm 10.1 test. We did not test ROCm 10.1.
  kernel_resources_test gives the result for any compiler without a GPU.

Test setup

• Machine: Strix Halo, gfx1151, 128 GiB. Fedora, Linux 7.2.7. No desktop session. The tuned
  profile was accelerator-performance. No other GPU process ran; each run recorded this.
• Models: Qwen3.8-27B UD-Q8_K_XL and UD-Q4_K_XL with the DFlash2 Q4_K_M draft; Qwen3.8
  Flash-Next UD-Q4_K_XL with the shared MTP Q8_0 head. All at the Hub revisions that the model
  cards pin.
• Containers: rootless podman, --device /dev/kfd --device /dev/dri/renderD128 --group-add keep-groups, models mounted read-only.

Build	Source	Compiler and libraries
clang 22, published 0.8.1 image	ghcr.io/gufo-org/toolboxes/gufo-runtime:0.8.1 (543300e). 05688f9 adds one commit to it, #449, which changes server code only.	AMD clang 22.0.0, ROCm 7.2.3, GCC 15.3 (Nix)
clang 23, before / after	05688f9 without / with this PR. CMake release preset in AMD's rocm/dev-ubuntu-26.04:10.0.0-full image.	AMD clang 23.0.0, ROCm 10.0.0, GCC 15.2
clang 22 over ROCm 10, before / after	The same source. The Nix clang 22 of the published image is CMAKE_HIP_COMPILER, over the ROCm 10.0.0 userspace.	AMD clang 22.0.0, ROCm 10.0.0 libraries

Commands (M, D, F and T are the model, draft, Flash-Next and MTP files):

gufo bench -m "$M" -p 2048,64,16,8 -n 64 -r 3
gufo bench -m "$M" -p 2048 -n 128 -r 3 --speculative dflash2 --dflash-model "$D"
gufo bench -m "$F" -p 2048 -n 128 -r 3 --speculative mtp --mtp-model "$T"
gufo prompt -m "$M" -f prompt.txt -n 160 -t 0 --no-display-prompt
gufo bench -m "$M" --validate-prefill 1024 -p 1024 -n 0 -r 1
python3 tools/ci/check-kernel-resources.py BINARY --baseline tools/ci/kernel-resources.json

Method

1. We ran gufo bench with a clang 22 and a clang 23 build. We used the same flags, the same
   model files and alternate runs. We measured pp2048, pp64, pp16 and pp8, because batches of
   9 to 95 tokens and of 8 tokens or fewer use different kernel tiles.
2. We recorded kernel times with rocprofv3 --kernel-trace. One kernel,
   W8A8BlockedWmmaGEMMKernel<128,128,2,4,2,true,true>, used 47% of the GPU time of Q8 pp2048.
3. We read the AMDGPU metadata of all code objects in the two binaries. With clang 23 this
   kernel used 256 VGPRs, 149 spilled VGPRs and 480 bytes of scratch. With clang 22 it used
   181 VGPRs and no scratch. Its clang 23 code has 150 scratch_load and 120 scratch_store
   instructions. The clang 22 code has none.
4. We compiled all 74 HIP units with the two compilers from the same source and the same flags
   (--offload-device-only). With clang 23, 33 kernels have new or larger spills, 95 kernels
   have less VGPR occupancy, and 355 kernels are better.
5. We gave the optimized IR of clang 22 to the backend of clang 23. The spill occurred again.
   Thus the backend scheduler causes the W8A8 spill, not the optimizer.
6. We read the disassembly. Clang 23 changes 64 ds_load_b32 instructions into 32
   ds_load_2addr_stride64_b32 instructions, and it moves them above the WMMAs of the current
   K block.
7. We tested each candidate fix with the two compilers: register use, speed at four prompt
   sizes, speculative decode and the greedy output.
8. We made a clang 22 build over the ROCm 10 userspace. It isolates the compiler from the ROCm
   libraries. It is faster than the published image, and its output is the same. Thus the
   slowdown comes only from clang 23.

Steps 1 to 8 were done at 8a46cd9. The Qwen3.8-27B kernel files did not change from
8a46cd9 to 05688f9. The measurements in "Results" are new, at 05688f9.

Candidates (measured at 8a46cd9, clang 23, Q8_K_XL)

W8A8 kernel only. The K-quant changes are not in these builds, except in the last row.

Candidate	pp2048	pp64	pp16	pp8	tg64	DFlash2 pp2048	DFlash2 tg128
none	291.4	112.0	29.6	42.7	7.19	279.7	24.02
fence at the end of each K block (as #255)	461.6	125.0	33.4	50.2	7.20	435.2	24.02
fence at the end of each K block and each token tile	457.1	294.3	86.2	50.2	7.19	433.0	24.04
#pragma unroll 1 on all K-block loops	436.9	302.6	89.6	50.6	7.18	413.5	24.04
K-block fence and token-tile fence, unroll(BK <= 2 ? BK : 1)	459.8	301.7	88.0	50.4	7.15	433.9	23.98
K-block fence, unroll(BK <= 2 ? BK : 1) (chosen)	458.3	305.3	91.0	50.7	7.16	433.2	23.97
K-block fence, 128x64 tile launched with BK=2	458.6	289.5	85.0	49.7	7.19	432.5	24.02
chosen + K-quant changes + -fno-slp-vectorize (this PR)	465.5	308.6	90.8	50.7	7.14	436.3	23.98

K-quant kernels in prefill_quant_gemm.hip, total spilled VGPRs of the 33 WKQuantA8
instantiations:

Candidate	clang 22	clang 23
none	0	1022
unroll(BK <= 2 ? BK : 1)	0	265
unroll(BK <= 2 ? BK : 1) + -fno-slp-vectorize (chosen)	0	0
K-block fence + unroll(BK <= 2 ? BK : 1)	134	265
K-block fence + unroll(...) + -fno-slp-vectorize	134	134

An unconditional fence in SmallBatchKQuantExactFp32GEMMKernel also made one clang 22
instantiation spill (30 VGPRs). Thus this PR does not change that kernel.

Flash-Next RoutedF16GEMMKernel pair tiles, clang 23 scratch:

Candidate	Scratch
none (conditional over uint4 lvalues)	80 B
four named variables instead of an array	80 B
if/else chain	80 B
native ext_vector_type storage	does not compile
mask blend (chosen)	0 B

Register use

VGPRs / spilled VGPRs at 05688f9, from the release binaries. A value in brackets is
scratch without spills.

Kernel	clang 22 before	clang 22 after	clang 23 before	clang 23 after
W8A8 128x128	180/0	183/0	256/156	183/0
W8A8 128x128, fused SwiGLU	181/0	184/0	256/149	184/0
W8A8 128x64	166/0	154/0	256/148	153/0
W8A8 128x32	148/0	154/0	256/64	146/0
WKQuantA8 128x64 Q3_K	179/0	230/0	256/200	230/0
WKQuantA8 128x64 Q6_K	185/0	162/0	256/201	161/0
WKQuantA8 128x64 IQ4_XS	159/0	180/0	256/31	180/0
WKQuantA8 128x128 Q3_K	233/0	233/0	256/133	196/0
WKQuantA8 128x128 Q6_K	227/0	235/0	256/132	197/0
WKQuantA8 128x128 Q6_K, wave64 unit	175/0	175/0	256/90	167/0
RoutedF16GEMM Q4_K 128x128 pair	170/0	166/0	155/0 (80 B)	167/0
RoutedF16GEMM Q5_K 128x128 pair	179/0	175/0	164/0 (80 B)	176/0
RoutedF16GEMM Q5_K 128x64 pair	133/0	129/0	117/0 (80 B)	129/0

Remaining clang 23 items

With this PR, a clang 23 build has 13 kernels over the clang 22 baseline. In the measurements
above, we saw no cost from them. They are a list for a move to ROCm 10 (#228):

• qfn_mmq::mul_mat_q<type 7, 96|112|128> (6 tiles, up to 58 spilled VGPRs). The Flash-Next
  Q4 workloads do not use them.
• DS4 mul_mat_q<Q2_K> and mul_mat_q<IQ2_XXS> (2 tiles). They spill with clang 22 too.
  Clang 23 spills more.
• HalfPrefillGemmKernel<256,256,8,4> Q4_K, Q6_K and IQ4_NL tiles (2 to 6 spilled VGPRs).
• SmallBatchKQuantExactFp32GEMMKernel<16,7,3,Q4_K> (6 spilled VGPRs).
• rms_norm_plain_regs_kernel<64> (1 spilled VGPR).

Also, many BatchedExactBf16GEMMFp32VecKernel instances use more VGPRs with clang 23. For
example, <16,12,2,true,32,1,false> uses 61 VGPRs with clang 22 and 100 with clang 23 (16
waves against 12). They do not spill. Decode and verification use them too, so this PR does
not change them.

[Copilot]

Copilot review overview

🟡 Changes recommended

The resource test can silently succeed when it decodes zero kernels.

Review effort: Balanced
Findings: 1

Open (1)

• Reject empty kernel extraction to prevent false-passing baseline checks · New

What changed in this PR

Eliminates clang 23 register spills in Qwen prefill kernels and adds regression detection.

Changes:

• Constrains kernel scheduling, unrolling, and SLP vectorization.
• Reworks Flash-Next cache selection to avoid scratch memory.
• Adds kernel-resource validation and performance documentation.

File	Description
CMakeLists.txt	Registers the kernel-resource CTest.
src/​models/​qwen/​CMakeLists.txt	Disables SLP for affected HIP units.
src/​models/​qwen/​hip/​kernels/​prefill_quant_gemm.hip	Bounds unrolling and scheduling.
src/​models/​qwen/​hip/​kernels/​prefill_quant_gemm.hpp	Bounds K-block unrolling.
src/​models/​qwen38_flash_next/​kernels/​rocm/​kernels.hip.cpp	Replaces indexed cache selection with mask blending.
tools/​ci/​check-kernel-resources.py	Implements resource extraction and comparison.
tools/​ci/​kernel-resources.json	Defines accepted scratch-memory limits.
docs/​PERFORMANCE.md	Documents resource-check usage.
docs/​models/​qwen3.8-27b/​EXPERIMENTS.md	Records Qwen performance results.
docs/​models/​qwen3.8-flash-next/​EXPERIMENTS.md	Records Flash-Next results.

---

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

[francescobozzo]

Interesting, similar to #255 . It looks like ROCm 10 has also some small benefit of around ~20 tok/s on Qwen 27B

[jtsylve]

@francescobozzo is there anything else you'd like me to address to get this merged?

[francescobozzo]

Overall looks good, but my Strix Halo is currently doing some long running benchmarks and I'd like to verify locally a few of these numbers on it.

About DeepSeek V4 Flash, ASR, TTS, Qwen Image 2.1: are we spill-free there or you didn't check these models yet?

[jtsylve]

VGPR spills by model:

Model	Kernels	clang 22 (0.8.1 image)	clang 23 (this PR)
Qwen-Image-2.1	63	none	none
MiniMax-H3	85	none	none
ASR	28	none	none
TTS	68	none	none
DeepSeek V4 Flash	1010	28 kernels, 3388 VGPRs	23 kernels, 3206 VGPRs

Qwen-Image-2.1, MiniMax-H3, ASR and TTS are spill-free with the two compilers.

• Qwen-Image-2.1 and MiniMax-H3 use no scratch at all.
• ASR SegmentedSelfAttentionKernel<true> uses 272 B of scratch with the two compilers. This is
  not a spill: float4 query_row[16] is a private array with a runtime index. The 8 B of scratch
  in ASR TextQkNormRoPEKernel<128> and TTS Bfloat16QkNormRoPEKernel<128> on clang 22 is 0 B
  on clang 23.
• On clang 23, four kernels lose one occupancy step (12 to 10 waves) but do not spill: ASR
  TextDecodeGemvKernel and three Qwen-Image-2.1 ImplicitConvolution3<576, 32, 32, 6, 2, ...>
  tiles. We did not measure their speed.

DeepSeek V4 Flash is not spill-free with either compiler. Its mul_mat_q tiles spill on the
published 0.8.1 image too: Q2_K, IQ2_XXS, Q8_0 and MXFP4, up to 699 spilled VGPRs (2792 B of
scratch) in mul_mat_q<Q2_K, 80>. Of the 29 DS4 kernels that spill with one of the compilers:

With clang 23	Kernels
The spill goes away	6 (four Q2_K tiles of 96-128, two Q8_0 128 tiles)
Fewer spilled VGPRs	18
The same	2 (Q8_0 16 tiles, 1 VGPR)
More spilled VGPRs	3

The three worse kernels are the DS4 items in "Remaining clang 23 items" of the PR description:

Kernel	clang 22	clang 23
mul_mat_q<Q2_K, 16, true>	77 VGPRs, 312 B	86 VGPRs, 348 B
mul_mat_q<IQ2_XXS, 32, true>	12 VGPRs, 52 B	16 VGPRs, 68 B
rms_norm_plain_regs_kernel<64> (DS4 backend.hip.cpp)	0	1 VGPR, 8 B

Since this is not a regression, it's probably better to fix in a separate PR.

[francescobozzo]

LGTM, thanks! This will make the migration to ROCm 10 easier. Also, nice findings on ds4 flash