Skip to content

ROCm 10 evaluation for gfx1151, and a clang 23 prefill fix - #228

Closed
fedeizzo wants to merge 3 commits into
mainfrom
fedeizzo/rocm10-evaluation
Closed

fedeizzo wants to merge 3 commits into
mainfrom
fedeizzo/rocm10-evaluation

Conversation

@fedeizzo

@fedeizzo fedeizzo commented Sep 5, 2026

Copy link
Copy Markdown
Collaborator

What this is

An evaluation of ROCm 10.0.0 on Strix Halo (gfx1151), the packaging work needed
to run it at all, and two measured changes to the Qwen prefill path.

Net position: the 7.2.3 pin should stay. ROCm 10 wins shallow prefill by
~3% and loses deep prefill by ~3.5% and decode by ~5%. This branch records the
evidence and lands the parts that are worth keeping regardless of when the pin
moves.

Commits

  1. build(nix) — vendors nixpkgs' rocm-modules tree at therock-10.0.
    AMD's prebuilt tarball cannot create a HIP stream on this host (the runtime's
    internal blit program fails to link __amd_streamOpsIncrement, reported as
    hipErrorOutOfMemory); building the same sources through nixpkgs' packaging
    fixes it. ~50 changes across 14 modules.
  2. docs(rocm) — baseline.md (7.2.3 reference sweep) and
    rocm10-spike.md (the full investigation, including the dead ends).
  3. perf(qwen) — the two source/build changes below.

Measured (Qwen3.8-27B UD-Q4_K_XL, 3 repetitions)

clang 23's SLP vectorizer widens the blocked WMMA GEMM inner loops past the
256-VGPR budget: 67 of 214 instantiations spill (8570 VGPRs) against none under
clang 22. Disabling SLP for that one file:

Test 7.2.3 ROCm 10 before ROCm 10 after
pp2048 @ d4096 443.26 ± 0.16 373.50 461.16 ± 1.16
pp2048 @ d16384 389.33 ± 0.69 317.44 375.53 ± 3.06

The flag is scoped to the file and gated on CMAKE_HIP_COMPILER_VERSION >= 23,
so the pinned 7.2.3 build is untouched by construction.

roundf → __float2int_rn in activation quantization (18 → 10 instructions)
is worth +0.9% prefill on 7.2.3 (443.26 → 447.05) and +1.1% on ROCm 10.

Review focus

The rounding intrinsic changes numerics — ties round to even instead of away
from zero. Evidence gathered: prefill/sequential agreement improves
(max_abs_diff 0.5556 → 0.4168, cosine 0.99947 → 0.99963), the shift is
identical on both toolchains, and greedy output was unchanged on a sample
prompt. That is not a substitute for nix build .#checks.x86_64-linux.pr, which
has not been run here and should gate this before merge, since speculation
is documented as greedy-faithful.

Drafted for that reason. Splitting commit 3 out is easy if the numerics change
needs to wait.

Also recorded (negative results)

Min-register scheduling removes every spill and costs another 23%; -mcumode
looks good on paper and measures 10% worse; the existing prefill tile choice is
still best of six under clang 23. The remaining decode gap is localized with
rocprofv3 to the GEMV family and traced to a byte-load form change
(global_load_d16_u8 → global_load_i8) that no compiler flag or source
rewrite could reach — a candidate upstream clang report.

🤖 Generated with Claude Code

https://claude.ai/code/session_019QUqxmu9mfE9RHMmAjUdAm

fedeizzo and others added 3 commits September 5, 2026 09:11
AMD's prebuilt ROCm 10 distribution cannot create a HIP stream on this host:
the runtime JIT-compiles its internal blit/stream-ops program at the first
hipStreamCreate and the link fails on __amd_streamOpsIncrement, which HIP then
reports as hipErrorOutOfMemory. It reproduces with a twenty-line HIP program on
both the 10.0.0 release and the 10.1.0a nightly.

Building the same therock-10.0 sources through nixpkgs' packaging fixes it, so
this vendors nixpkgs' rocm-modules tree and bumps it to ROCm 10. Roughly fifty
changes were needed across fourteen modules, nearly all of them nixpkgs
carrying assumptions that no longer hold:

  * patches already merged upstream (verified per commit, not assumed)
  * patches that cannot be ported: rocm-libraries#2073 was never merged and
    ROCm 10 restructured TensileCreateLibrary around it, so hipblaslt now takes
    joblib from the python environment instead
  * new build requirements: python3 for the trap-handler and hipblas codegen,
    git for shared/origami, amd_smi for hipblaslt and rocblas clients, libnl
    and libmnl for amdsmi, rocprof-trace-decoder for rocprofiler-sdk
  * layout changes: MIOpen's ClangToolChain.cmake assumes TheRock's single
    prefix, so it is pointed at the split nix outputs
  * a regenerated device-libs cmake patch and an esmi bump to the commit ROCm
    10 pins, since the vendored 4.2 tag no longer matches the amdsmi API

MIOpen is built without Composable Kernel: nixpkgs marks CK broken without a
gfx9 target, and gufo only needs ck/ck.hpp headers, never a CK library.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_019QUqxmu9mfE9RHMmAjUdAm
baseline.md pins the ROCm 7.2.3 reference sweep for Qwen3.8-27B UD-Q4_K_XL so
any toolchain change can be compared against it.

rocm10-spike.md is the investigation: what broke in AMD's tarball and why, the
packaging work that fixed it, and the measurements that followed. The headline
result is that clang 23 SLP-vectorizes the blocked WMMA GEMM inner loops past
the 256-VGPR budget, so 67 of 214 instantiations spill where clang 22 spilled
none, costing 16-19% of prefill.

It also records what did not work, so it is not retried: min-register
scheduling removes every spill and costs another 23%, CU mode looks good on
paper and measures 10% worse, and the existing prefill tile choice is still the
best of the six variants under the new compiler. The remaining decode gap is
localized to the GEMV family with rocprofv3 and traced to a byte-load form
change that no flag and no source rewrite could reach; that one is a candidate
upstream report.

hip_probe.cpp and rocm10-env.sh are the reproducers those documents reference.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_019QUqxmu9mfE9RHMmAjUdAm
Two changes to the quantized prefill path, measured on Qwen3.8-27B UD-Q4_K_XL
at three repetitions.

ROCm 10 ships AMD clang 23, whose SLP vectorizer widens the blocked WMMA GEMM
inner loops past the 256-VGPR budget: 67 of 214 instantiations spill, 8570
VGPRs in total, against none under clang 22. Disabling SLP for that translation
unit cuts it to 25 and 2183, and takes pp2048@d4096 from 373.50 to 461.16 t/s,
which is 2.9% ahead of the 7.2.3 build rather than 16% behind it. The flag is
scoped to the one file and gated on CMAKE_HIP_COMPILER_VERSION >= 23, so the
pinned 7.2.3 toolchain is untouched by construction.

The activation quantization used roundf, which needs the half-away-from-zero
fixup sequence: 18 instructions against 10 for __float2int_rn. The intrinsic is
worth about 0.9% of prefill on 7.2.3 and 1.1% on ROCm 10.

That second change is not numerically free, since the intrinsic rounds ties to
even instead of away from zero. Prefill agreement with sequential execution
improves rather than degrades (max_abs_diff 0.5556 -> 0.4168, rmse 0.0951 ->
0.0794, cosine 0.99947 -> 0.99963), the shift is identical on both toolchains,
and greedy output was unchanged on a sample prompt. Speculation here is
documented as greedy-faithful, so this wants the full PR gate before it lands
on main.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_019QUqxmu9mfE9RHMmAjUdAm
@gabrielfreire

Copy link
Copy Markdown

Could you try again with ROCm 10.1 ?

https://rocm.blogs.amd.com/ecosystems-and-partners/rocm-10.1-blog/README.html

@jtsylve

jtsylve commented Oct 6, 2026

Copy link
Copy Markdown
Contributor

Data from a ROCm 10.0.0 build of main at 05688f9, compared with the published 0.8.1 image on
the same Strix Halo machine. #459 has the full data.

  1. With clang 23, Q8 pp2048 is 291 tok/s against 454, and Q8 pp64 is 112 against 322. Q4
    pp2048 is not slower (668 against 649). -fno-slp-vectorize alone does not change the Q8
    numbers. The backend scheduler causes the Q8 spill: the clang 23 backend spills with the
    optimized IR of clang 22 too.
  2. perf(hip): keep prefill GEMMs spill-free on clang 23 #459 removes the spills with source constraints. With it, clang 23 gives Q8 pp2048 464 and
    pp64 308.
  3. A clang 22 build over the ROCm 10 userspace is faster than the 0.8.1 image (Q8 pp2048 476
    against 454), and its output is the same. Thus the ROCm 10 libraries do not cause the
    slowdown. Only clang 23 does.
  4. Clang 23 changes the greedy text of Flash-Next (the unit uses -ffast-math). The Flash-Next
    and DeepSeek V4 Flash quality gates need a new run before a move to clang 23.
  5. About ROCm 10.1 (the comment above): perf(hip): keep prefill GEMMs spill-free on clang 23 #459 adds tools/ci/check-kernel-resources.py. It reads
    the scratch use of each kernel from a binary, with no GPU. Run it on a ROCm 10.1 build to see
    if its compiler adds or removes spills.

@sdwilsh

sdwilsh commented Oct 22, 2026

Copy link
Copy Markdown
Contributor

I tasked an agent to re-measure ROCm 10.1.0 against the 7.2.3 pin on current main, find what slows it down, and wrote three small commits that remove most of the slowness (for one model). With them, Qwen3.8-27B Q8 prefill is 1% to 2% faster than 7.2.3, Q4 prefill is 1% to 2% slower, and Q4 decode is still 3% slower.

Setup

ROCm 10.1.0 is not in nixpkgs, so I used the open nixpkgs PR 573885 at 925eadb600eb, which builds ROCm's LLVM 24 (this is unlikely to merge as is, IMO, and it might behoove us to help get it merged in a more sensible way if we want to pursue this). The device compiler is now clang 24, and builds use nix build. I had to hack together a commit that exists only for these runs that adds a package that uses nixpkgs' gfx1151-only ROCm set, because the 10.1 libraries have no binary cache. This table tries to better explain what I tested:

Label ROCm Device compiler Source
A 7.2.3 clang 22 main f020bd3a, unchanged
B 7.2.4 clang 22 same source, newer nixpkgs, as a control
D 10.1.0 clang 24 same source, no CWSR patch
FIN 10.1.0 clang 24 main 53ebb51e plus my local changes

Each pod ran gufo bench --n-prompt 64,2048 --n-gen 128 --n-depth 0,4096,16384 --repetitions 3 on an otherwise idle GPU. A and B agree within 0.6% on every pp2048 and tg128 cell, so the nixpkgs bump alone does not appear to impact things.

The CWSR patch is gufo's rocr-fixed-cwsr.patch, which keeps the GPU queue save areas out of movable pages so that memory compaction cannot suspend inference. It does not apply to ROCr 1.21 as written, so FIN carries a port of it and D has no patch at all.

Q4 prefill on A drifted down 1.2% between day one and day two. Q8 and decode did not drift. D ran on day one, so I compare it to A from day one. FIN ran on day two, so I compare it to A from day two.

Results

Qwen3.8-27B, tokens per second.

Test A, day 1 D D vs A A, day 2 FIN FIN vs A
Q4 pp2048 624.8 590.7 -5.5% 617.3 611.0 -1.0%
Q4 pp2048 @ d4096 593.7 556.8 -6.2% 586.2 576.9 -1.6%
Q4 pp2048 @ d16384 533.1 502.0 -5.8% 529.6 518.8 -2.0%
Q4 pp64 @ d16384 267.9 248.8 -7.1% 265.9 272.4 +2.4%
Q4 tg128 12.30 11.96 -2.8% 12.27 11.91 -2.9%
Q4 tg128 @ d16384 11.65 11.36 -2.5% 11.62 11.31 -2.7%
Q8 pp2048 445.7 438.8 -1.5% 443.2 452.9 +2.2%
Q8 pp2048 @ d16384 398.7 391.7 -1.7% 398.9 403.0 +1.0%
Q8 pp64 @ d4096 296.2 279.4 -5.7% 297.4 297.3 0.0%
Q8 pp64 @ d16384 269.0 233.9 -13.0% 268.9 266.4 -0.9%
Q8 tg128 7.21 7.29 +1.1% 7.21 7.30 +1.2%

This PR's hunks

The -fno-slp-vectorize hunk for prefill_quant_gemm.hip does nothing on current main. D with the hunk gives the same numbers as D, and the same register and scratch report for every kernel. D has no spills in that file for the hunk to remove. I did not measure the roundf hunk or run the numerical checks it needs.

What slows 10.1 down

I used rocprofv3 --kernel-trace, the disassembly of the hot kernels, and tools/ci/check-kernel-resources.py --compare.

  1. Real true16 code generation costs Q4 prefill about 4%. HalfPrefillGemmKernel has 79 instantiations in prefill_fp16.hip, and it is the largest stage. clang 24 stops emitting v_pack_b32_f16 and emits v_lshrrev_b32 plus 16-bit v_mov_b16 instead. Each instantiation then uses 25 to 30 more VGPRs and loses one wave per SIMD. -Xclang -target-feature -Xclang -real-true16 removes the change.
  2. The SLP vectorizer costs Q8 prefill about 2.5%. clang 24 widens the inner loops of BatchedExactBf16GEMMFp32VecKernel, and some instantiations drop from 16 to 12 waves. -fno-slp-vectorize on prefill_gemm.hip and small_batch_wave64.hip restores them. It makes Q8 prefill faster than 7.2.3 and costs Q4 nothing.
  3. A register spill costs short prompts at depth up to 13%. WmmaCausalAttention<32,16,false> uses 256 VGPRs and spills 14, which is 60 bytes of scratch. clang 22 used 253 VGPRs and no spill. A call now takes 3.8 ms instead of 2.1 ms. The 208 calls in the timed repetitions add about 27 of the 36 extra milliseconds per repetition. No compiler flag removes the spill that the investigating agent could find, unfortunately. Two live ranges appear to cause it: LICM hoists 13 per-lane limits out of the key loop, and clang 24 keeps four probability fragments live across both dimension-tile passes.
  4. Q4 decode is 3% slower. The gemm: gemv stage is 1.9% slower. In Q8KBlockGEMVKernel_2Rows, clang 24 turns 16 global_load_d16_u8 into 8 of those plus 8 global_load_i8 (the same change the 10.0.0 write-up found on clang 23). Wave32FusedSSMInputProjectionsKernel_1Row is 7% slower with almost the same instruction mix, and I do not know why. It is not the CWSR patch, because 7.2.3 without it matches 7.2.3 with it.

Checks

  • On 7.2.3, my change to the attention kernel gave byte-identical --validate-prefill 1024 numbers and byte-identical greedy output on three prompts for Q4 and Q8. I tested it before limiting it to compilers newer than clang 22. After that the disassembly on clang 22 is identical to the original.
  • On 10.1, all six greedy outputs match 7.2.3. --validate-prefill is identical to 7.2.3 for Q8. For Q4 it shifts slightly, with rmse from 0.00430 to 0.00569 and cosine from 0.99999893 to 0.99999809, and the top-1 token stays the same. Stock 10.1 gives the same shifted numbers as FIN, so the 10.1 toolchain causes the shift and the commits do not.
  • With HSAKMT_DEBUG_LEVEL=6, stock 10.1 prints "Allocating GTT for CWSR" and 7.2.3 with the patch does not. FIN does not print it either, so the port moves the queue save area off the movable path.

Not checked

  • Compaction itself. I cannot trigger it in the pods (my two strix halo machines run kubernetes), so I know the CWSR path is the fixed one but not that inference survives compaction.
  • Whether the port's amd_memory_region.cpp hunk still keeps HIP's runtime arenas resident. I had to make it clear the NonPaged flag for host registration, because ROCr 1.21 rejects it there. The model loads and runs.
  • Whether -fno-slp-vectorize is needed on both files.
  • The roundf hunk in this PR, on either toolchain.
  • The remaining gaps in Q4 decode (-3%), Q4 prefill (-1% to -2%), and Q8 pp64 at depth (about -1%).

I can share the work if needed, but it's local-only for now. Hope this helps!

@fedeizzo

Copy link
Copy Markdown
Collaborator Author

Closing in favor of #741

@fedeizzo fedeizzo closed this Oct 23, 2026
@francescobozzo
francescobozzo deleted the fedeizzo/rocm10-evaluation branch October 31, 2026 18:51
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants