Between a ROCm and a Hard Place.

Scope

A Note On Scope

One piece of silicon: the Radeon 8060S iGPU in the AMD Ryzen AI MAX+ PRO 395, gfx1151, Strix Halo. Nothing below is a Strix Point or discrete RDNA4 number unless it says so.

The build that wins prefill on this box cannot find one sentence in a 16K-token prompt. That is most of the post. The rest is two questions about serving llama.cpp on Strix Halo: Vulkan or native ROCm as the backend, and Fedora's packages or AMD's TheRock SDK as the ROCm underneath it. Both have measured answers. Both answers carry a caveat that does not fit in a footnote.

Throughput

ROCm wins prefill, Vulkan wins decode, and decode wins by more

The ROCm side is llama.cpp at de699957b compiled against TheRock's ROCm 7.15.0a20260718 for gfx1151: no HSA_OVERRIDE_GFX_VERSION, no vendor image. The Vulkan side is the existing baseline on the same box, same model file, same two prompts, same flags. Greedy decoding, no prompt cache. When the ROCm build loaded, GTT usage jumped to about 15.5 GB against a 14.25 GB GGUF, so it was on the GPU and not quietly on the CPU.

Native ROCm takes prefill by 3.4% on word-salad and 16.0% on natural prose. Vulkan takes decode by 1.40× and 1.50×. That is the pattern people report for RDNA in general, ROCm ahead in the compute-bound phase and Vulkan ahead in the bandwidth-bound one, and on this box it holds for both prompt shapes.

Two things about the ROCm column. It lands within 0.1 to 9% of the vendor-bundled rocm-7.2 image I had used as a reference before (798 and 38.0 on word-salad, 279 and 46.9 on natural prose), so my own build is a faithful native engine and not a Containerfile artifact. And the word-salad decode number needed ignore_eos: greedy decode from that prompt emitted EOS as the very first token on this build, so without forcing n_predict there is no decode to measure. The Vulkan runs never needed it. The research note filed that as a measurement detail. It is the same class of effect, a backend-dependent argmax flip, as what follows.

This is two prompt shapes, one run each, one model, one box. The margins are large next to run-to-run noise for this kind of measurement, so the direction is solid. The exact percentages are what I saw on this model with these prompts, not a constant to plan capacity from.

Correctness

The same build gets the needle wrong, every time, differently

I ran a needle-in-a-haystack check against the native ROCm server: the 16,444-token word-salad prompt with one sentence inserted, "The secret access code for the vault is 77451-BASALT." With thinking on, the reply looped on "nabla, "nabla, "nabla until the token budget ran out. With thinking off it answered cleanly and wrong: "The code is 12345." A raw /completion with an "Answer: The secret access code is" primer gave " 12345." as well. The Strix Point box did the same.

On 1 August I re-ran it properly, on this box only, with the needle at 10, 30, 50, 70 and 90% depth, and with a Vulkan control on the same model file, same flags, same prompts: upstream's prebuilt b10219 Vulkan build, started after the ROCm container was stopped so the two never shared the memory pool.

ROCm was wrong at every depth. Vulkan was right at every depth, with byte-identical output on repeats. And the ROCm failures are not even stable. Repeating the identical request at temperature=0, no restart in between, gave " 12345." on one run and sixteen copies of <unused49> on the next at 10% depth, and flipped between the <unused49> loop and a partial " 77451." at 70%. Three failure shapes from one request that should be pure argmax: a fabricated number, a loop on a reserved vocabulary token, and the right prefix with the suffix dropped.

I have not found the kernel. Non-determinism at temperature zero points at a floating-point reduction whose order varies with dispatch, in attention softmax or a GEMM accumulation, on this specific TheRock snapshot; that narrows it, it does not pin it. What I can say is that on this build, native ROCm gets wrong a task that Vulkan gets right, on the same file, on the same silicon, and that is not a caveat you average into a throughput table.

Verdict

Serve on Vulkan

Vulkan is the default for llama.cpp on gfx1151 today. It is 1.4 to 1.5× faster where a chat workload spends its time, and it passed the check ROCm failed.

The prefill number is still real, and I am not going to hide it because it comes with an asterisk. If your workload is prefill-shaped, bulk ingestion for RAG or long-document summarisation with a short reply, native ROCm is 3 to 16% faster at the part you care about. It is also, on this build, wrong about what is in the document. Run the needle check on your own model and prompts before you route retrieval traffic through it. If it passes for you, take the speed. If you skip the check, you are trading a measured gain for a measured risk.

ROCm source

Fedora's ROCm or TheRock's

The second question has nothing to do with tokens per second. It is which ROCm you compile against, and the trade is version control against footprint and operational fuss.

Fedora ships ROCm 7.1.1 as ordinary packages, rocm-runtime, rocm-hip-libraries and friends. For gfx1151 that path builds clean, GGML_HIP_ROCWMMA_FATTN=ON included. The one gap I know of is on the smaller Strix Point part: Fedora's rocwmma-devel (rocWMMA 2.0.0, 7.1.0-5.fc44) has no architecture-table entry for gfx1150, so anything that includes the rocWMMA headers dies at compile time with static_assert(0, "Unsupported architecture"). AMD's own rocWMMA inside TheRock compiles the same code for gfx1150 without complaint, so that is Fedora's packaging, not upstream.

TheRock is AMD's relocatable ROCm SDK: a tarball you extract to any prefix, pinned to an exact build, independent of what your distro packages. The value of that is not abstract. The 7.14.0a20260612 snapshot crashed on --version, before any output, on both Strix boxes, stripped or unstripped, single-arch or fat binary. dmesg showed a jump to 0x100000001, a corrupted function pointer rather than a null dereference, and gdb put it inside TheRock's own libhsa-runtime64.so.1, in rocr::AMD::GpuAgent::InitDma(), during HSA init the first time a real GPU agent is enumerated. Before model load, before any of my code. HSA_ENABLE_SDMA=0 did not help. The same 7.14 line was already running fine on a discrete gfx1201 card, so this is a bug specific to Strix (RDNA 3.5) agents in one alpha snapshot. The fix was THEROCK_VERSION=7.15.0a20260718. No patch, no Containerfile edit, no waiting for Fedora to repackage anything. Every ROCm number above exists because of that one-line bump.

The cost is also concrete. Every version string above ends in a and a date; these are nightlies, not releases, and you manage them yourself. And the images are bigger:

A gfx1151-only TheRock image is within a few hundred megabytes of Fedora's. The cost shows up when you bundle arches for a mixed fleet: the LLVM and comgr closure (roughly 300 MB) is paid once, and the per-arch Tensile kernel data stacks on top. For a single Strix Halo host it is a modest tax, not a gigabyte one.

So: stay on Fedora's packages unless you need a ROCm version or a fix that has not reached them, then pin TheRock forward. The two axes are independent. A native ROCm llama.cpp builds from either source, and the Vulkan build touches neither.

Reference

Commands

rocm-smi --showmeminfo gtt
rocm-smi --showmeminfo vram

cmake -S . -B build \
  -DGGML_HIP=ON \
  -DGGML_HIP_ROCWMMA_FATTN=ON \
  -DAMDGPU_TARGETS=gfx1151 \
  ...

make therock-fetch FAMILY=gfx1151 THEROCK_VERSION=7.15.0a20260718
make therock-bases FAMILY=gfx1151 THEROCK_VERSION=7.15.0a20260718 GPU_TARGETS=gfx1151
make llamacpp-therock REF=master GPU_TARGETS=gfx1151

The thing I most want to close is the kernel behind the needle failure. Until it has a name, the ROCm prefill win on this box is a number I report, not one I serve from.