Qwen3.8-27B on RX 7900 XTX: 128K-context vision chat on ROCm with llama.cpp

What You'll Build

A local OpenAI-compatible server running Qwen3.8-27B — Qwen's 27B dense vision-language model — on a single RX 7900 XTX, answering both text and image prompts at a 131,072-token context window. The install is llama.cpp built against ROCm/HIP, with a 4-bit GGUF plus its separate vision projector.

Hardware data: RX 7900 XTX (24GB VRAM) · ~21.2 GiB derived working set at 128K context · See benchmark data

⚠️ Known issue: on ROCm, overshooting VRAM does not fail loudly. Buffers spill into GTT — system memory addressed through the GPU — and throughput collapses mid-request with no error at load. A reporter running the previous generation of this same qwen35 architecture on an RX 7900 XT watched GTT go from 0.59 GB to 14.65 GB during a single prefill (llama.cpp #26432). Budget headroom accordingly; Troubleshooting has the monitoring command.

Requirements

What ROCm changes, and what it doesn't

This page exists because a Radeon is not an NVIDIA card with a different logo. Three things about this model behave differently here, and one thing that looks like it should does not.

One backend, compiled twice. llama.cpp has no separate AMD kernel tree. ggml/src/ggml-hip/CMakeLists.txt globs ../ggml-cuda/*.cu and compiles it through HIP with GGML_USE_HIP defined. So the model-level reasoning on this site's NVIDIA pages for Qwen3.8-27B — the layer split, the KV arithmetic, the flash-attention type gates — is reading the same source file your build compiles. What differs is not the code but the branch it takes at runtime, and that branch is keyed on the device's compute capability.

Matrix-multiply dispatch is type-dependent here, and it is not on NVIDIA. In ggml/src/ggml-cuda/mmq.cu, ggml_cuda_should_use_mmq first checks whether the quantization type has an MMQ kernel at all, then returns true outright if turing_mma_available(cc). That predicate is defined in common.cuh as GGML_CUDA_CC_IS_NVIDIA(cc) && …, so on a Radeon it is false and control falls through to the AMD branches. RDNA 3's branch — reached via amd_wmma_available(cc), which covers RDNA 3 and RDNA 4 — carries a per-type switch: Q2_K uses MMQ only up to a batch of 128, Q6_K likewise on this generation, IQ2_XS/IQ2_S likewise, and everything else falls to default: return true. Both quant types this recipe considers, Q4_K and IQ4_NL, land in that default, so the kernel choice does not separate them on this card — but the shape of the rule is inverted from the NVIDIA one, and a quant tier that is named in that switch would behave differently here than on a GeForce.

The WMMA flash-attention kernel does not apply to this model. ggml/src/ggml-cuda/fattn.cu selects the AMD matrix-core path only when amd_wmma_available(cc) && gqa_opt_applies && Q->ne[0] <= 128. Qwen3.8-27B's head dimension is 256 (head_dim: 256 in the config, qwen35.attention.key_length = 256 in the GGUF header), so that condition is false and attention runs on llama.cpp's generic vector kernel during decode and its tile kernel during prefill. That is a statement about which kernel is selected, not a measured slowdown — but it is why you should not assume this model inherits the throughput of a 128-head-dim model on the same card.

What does not transfer at all: the FP8 and FP4 escape hatches. RDNA 3's WMMA units take FP16, BF16, INT8 and INT4 — there is no FP8 and no FP4 in hardware, which arrived with RDNA 4. So Qwen/Qwen3.8-27B-FP8 and the NVFP4 conversions that exist for this model are not memory-saving paths on this card; an FP8 tensor upcasts on load and buys nothing. Nor is there a flash-attn wheel to install, an ExLlamaV2/EXL2 path, or a Marlin kernel — subtract all of it. The one place RDNA 3 is favoured: llama.cpp keeps BF16 compute output on RDNA 3 and CDNA instead of falling back to F32, which is why this recipe uses the BF16 vision projector rather than the F16 one. Full-BF16 weights are not an option regardless — unsloth's split BF16 GGUF is 54.66 GB across two files.

Why a 27B model leaves room for 128K of context

Qwen3.8-27B is a hybrid: the model card gives its stack as "16 × (3 × (Gated DeltaNet → FFN) → 1 × (Gated Attention → FFN))". Its config.json says the same thing in a different form: a layer_types array of 64 entries, 48 linear_attention and 16 full_attention, at indices 3, 7, 11 … 63. Only those 16 layers hold a KV cache; the other 48 carry a fixed-size recurrent state that does not grow with context. llama.cpp derives the identical split in src/models/qwen35.cpp, marking a layer recurrent when (i + 1) % full_attention_interval != 0.

With head_count_kv = 4 and key_length = value_length = 256, one token costs 2 × 4 × 256 × 2 bytes = 4 KiB per attention layer — 64 KiB per token at f16 across 16 layers, a quarter of what a conventional 64-layer 27B would charge. At q8_0 (34 bytes per 32-element block) it is 34,816 bytes per token.

The 48 recurrent layers cost a constant instead. llama_hparams sizes them at (d_conv − 1) × (d_inner + 2 × n_group × d_state) for the convolution state and d_state × d_inner for the recurrent state — 30,720 + 786,432 elements per layer, held in F32 regardless of --cache-type-k/v, and multiplied by the number of parallel sequences. At --parallel 1 that is 156,893,184 B, about 0.146 GiB.

Derived budget at -c 131072 with q8_0 K and V:

(The GiB column is rounded for display; the total is computed from the unrounded byte column and rounded once.)

That leaves 2.805 GiB of the card's 24 GiB for graph and compute buffers, the HIP context and the vision graph, none of which is derived here. This is a derived envelope from cited file sizes and the runtime's own allocation formulas, not a measured peak — see /check/qwen3-8-27b/rx-7900-xtx for measured data as it lands.

Two deductions come out of that headroom rather than out of the components, and the first is bigger on AMD than the equivalent line on an NVIDIA page. The card does not present 24 GiB to the runtime: the vulkaninfo dump posted with llama.cpp #27097, from an RX 7900 XTX running this exact model under Linux, reports a device-local heap of 25,501,368,320 B — 23.75 GiB, with a 22.44 GiB budget once the desktop is resident. Treat ~23.7 GiB as the ceiling, not 24. Second, the display driver's own reservation comes out of the same pool if a monitor is attached. Check yours before committing to a context size:

amd-smi monitor --vram-usage

Two limits are worth knowing before you pick a window. The same weights with an unquantized f16 KV at 131072 need 24.945 GiB and do not fit; drop to -c 65536 and f16 KV fits at 20.945 GiB. And the model's full native window — the card advertises "Context Length: 262,144 natively and extensible up to 1,000,000 tokens." — needs 25.445 GiB even at q8_0, so 262K does not fit on a 24 GB card by this path.

Picking a quant on a 24 GB Radeon

There is no first-party GGUF. Qwen's Hugging Face organisation carries exactly four Qwen3.8 repositories — Qwen3.8-27B, Qwen3.8-27B-FP8, Qwen3.8-2.4T-A95B and Qwen3.8-2.4T-A95B-FP8 — and none is a GGUF conversion, even though the same org publishes 54 -GGUF repositories for other model families. Every GGUF below is a community conversion, re-checked by enumeration on 2026-08-16.

This recipe leads unsloth's Q4_K_M at 15.932 GiB, the same file the RTX 4090 page here leads, for the same reason: with 2.8 GiB of headroom the less lossy of the two 4-bit builds that fit 128K is the better default, and on this card the kernel path does not distinguish them (see the MMQ note above). Two alternatives are worth naming:

• UD-Q4_K_XL, 16.692 GiB. Totals 21.956 GiB in the same configuration, leaving 2.044 GiB. It has one thing no other file has: it is the build in the only public benchmark of this model on this card (#27097) — though that run is on the Vulkan backend and without a projector, so it is not evidence about this recipe's configuration. If you take it, note that llama-bench prints it as Q4_K - Small: unsloth's Dynamic builds carry a wrong general.file_type, which unsloth acknowledged — a display defect, not the tensors.
• IQ4_NL, 15.216 GiB. Totals 20.479 GiB, buying back 0.716 GiB against roughly half a bit per weight. The documented step down if your headroom proves tighter than the arithmetic suggests.

Do not climb: Q5_K_M (19,834,055,648 B) totals 23.735 GiB in this configuration. That is inside 24 GiB on paper, and inside the 23.75 GiB heap as well — by 15 MiB, which is not a margin anyone can plan on. What actually rules it out is the next line of the same dump: a 22.44 GiB budget once a desktop is resident, which 23.735 GiB exceeds by 1.295 GiB.

Two AMD-shaped detours that look like the right answer

AMD publishes its own quantizations of this model, and they are not for this card. The amd org ships Qwen3.8-27B-Quark-AWQ-INT4-W4A16, a Qronos INT4 sibling and an MXFP4 build. They are safetensors for vLLM-class runtimes, not GGUF: the INT4 card states it requires a runtime with W4A16Int4 scheme support and links vLLM PR #48606, which was still open and unmerged when this page was written. The weights are 19,535,853,721 B (18.194 GiB) before any cache, and AMD names no consumer Radeon anywhere on those cards. Nothing here establishes a released runtime that loads them on a 7900 XTX.

Nine repositories whose names begin "ROCm" will surface if you search, and stock llama.cpp cannot load them. The ROCmFP4 / ROCmFPX family — Qwen3.8-27B-ROCmFP4-STRIX-MTP-GGUF, -ROCmFPX-GGUF and their siblings — uses tensor formats from a third-party fork. That fork's own documentation is honest about it: "The ROCmFP4 / ROCmFPX tensor formats (ggml types 100–106) exist only in this fork." The -STRIX variants additionally target gfx1151, the Ryzen AI Max+ APU, not a discrete RDNA 3 card. Treat the prefix as a fork marker, not as "the AMD build".

Installation

1. Confirm your gfx target

Every build flag below keys off it, and it is the single easiest thing to get wrong on AMD:

rocminfo | grep gfx | head -1 | awk '{print $2}'

This card is gfx1100, per AMD's own ROCm 7.14.0 compatibility matrix, which lists AMD Radeon RX 7900 XTX (gfx1100) alongside the 7900 XT and 7900 GRE. Use that matrix rather than the family-level "7000 Series" compatibility page — it is the document that distinguishes gfx1100 from gfx1101 (RX 7800 XT / 7700 XT) and gfx1102 (RX 7600), and those are different code objects.

2. Install ROCm

Follow AMD's ROCm quick start for Linux. Two constraints worth knowing before you start: llama.cpp's HIP build asserts a floor in CMake — "At least ROCM/HIP V6.1 is required" — and AMD's system-requirements table restricts consumer Radeon cards to Ubuntu 24.04.4, Ubuntu 22.04.5, RHEL 10.1 and RHEL 9.7 (footnote 7 there; the shorter compatibility matrix carries the gfx targets but not this restriction). ROCm 7.14.0 is the current production release.

You should not need HSA_OVERRIDE_GFX_VERSION. llama.cpp's build documentation frames it as a fallback — "If your GPU is not officially supported you can use the environment variable" — and this card is officially supported. Setting it makes your card masquerade as another ISA, which is a good way to turn a working configuration into an unexplained one.

3. Build llama.cpp with HIP

There is no prebuilt Linux ROCm binary: the project's release assets carry llama-<build>-bin-win-rocm-7.14-x64.zip for Windows and no ROCm tarball for Linux, so on Linux you build (or use the container below).

git clone https://github.com/ggml-org/llama.cpp
cd llama.cpp
HIPCXX="$(hipconfig -l)/clang" HIP_PATH="$(hipconfig -R)" \
    cmake -S . -B build -DGGML_HIP=ON -DGPU_TARGETS=gfx1100 -DCMAKE_BUILD_TYPE=Release \
    && cmake --build build --config Release -- -j 16

That is llama.cpp's own documented invocation with the target switched to this card. The docs note that GPU_TARGETS "is optional, omitting it will build the code for all GPUs in the current system" — naming it is faster to compile and removes any doubt about which ISA you got.

Container alternative. ghcr.io/ggml-org/llama.cpp:server-rocm is a linux/amd64 image built from .devops/rocm.Dockerfile, whose ROCM_DOCKER_ARCH is a fat list including gfx1100, on a rocm/dev-ubuntu-24.04:7.2.1-complete base. llama.cpp's docker documentation is candid that the GPU images are "not currently tested by CI beyond being built".

On build numbers. The two things this recipe structurally needs — the qwen35 architecture in src/llama-arch.cpp and the qwen3vl_merger projector type in tools/mtmd/clip-impl.h — are both present as far back as b8001, thousands of builds before this model shipped, so the support floor is old. One flag in the Running section is newer: --reasoning-effort is absent at b10433 and present at b10434. The model is days old and several bugs against it are open, so prefer a recent build and pin the one you tested.

4. Download the weights and the vision projector

Vision is a separate file. The main GGUF alone gives you a text-only model.

pip install -U "huggingface_hub[cli]"
hf download unsloth/Qwen3.8-27B-GGUF \
  Qwen3.8-27B-Q4_K_M.gguf mmproj-BF16.gguf \
  --local-dir ./qwen3.8-27b

mmproj-BF16.gguf declares clip.projector_type = qwen3vl_merger, which llama.cpp maps to its PROJECTOR_TYPE_QWEN3VL handler — the projector is structurally loadable by this runtime, not merely published alongside it.

Running

./build/bin/llama-server \
  -m ./qwen3.8-27b/Qwen3.8-27B-Q4_K_M.gguf \
  --mmproj ./qwen3.8-27b/mmproj-BF16.gguf \
  -ngl 99 -c 131072 --parallel 1 \
  --cache-type-k q8_0 --cache-type-v q8_0 \
  --flash-attn on -b 2048 -ub 512 \
  --image-min-tokens 1024 \
  --jinja --reasoning-effort medium \
  --temp 1.0 --top-p 0.95 --top-k 20 --min-p 0.0 \
  --presence-penalty 0.0 --repeat-penalty 1.0 \
  --host 127.0.0.1 --port 8080

The server listens on http://127.0.0.1:8080 with an OpenAI-compatible API and a web UI that accepts image uploads. On a HIP build the device enumerates as ROCm0; on load the server logs loaded multimodal model once the projector is accepted.

Four of those flags are doing load-bearing work and are not defaults:

• --parallel 1 is worth 0.438 GiB and is easy to omit. common/common.h declares n_parallel = 1, but that is not what llama-server runs: common/arg.cpp rewrites it to -1 for the server example before your arguments are parsed, and tools/server/server.cpp resolves a negative value to n_parallel = 4 with kv_unified = true. Because the recurrent state is allocated per sequence, the default costs 627,572,736 bytes instead of 156,893,184. This is backend-independent, and it is observable on AMD hardware: the startup log in llama.cpp #27124, from an AMD box whose command line carries no -np flag at all, prints n_slots = 4, n_ctx_slot = 65536, kv_unified = 'true'. The KV cache itself is not multiplied — under kv_unified the four slots share one full-size window — so what you buy is one copy of the recurrent state and a window that cannot be evicted by a concurrent request.
• --cache-type-k q8_0 --cache-type-v q8_0 halves KV against f16. Do not go to 4-bit KV, and do not mix the two types — see Troubleshooting.
• --image-min-tokens 1024 is what llama.cpp itself asks for: it warns at load that Qwen-VL models need at least 1024 image tokens to work correctly on grounding tasks. At q8_0 KV, an image at that floor costs roughly 35 MB of cache.
• --reasoning-effort medium overrides the model card's default of xhigh (the card documents that "supported levels are xhigh, medium, and low"). On a 128K window a runaway trace is affordable; on any smaller one it is the thing that exhausts your context. On a build older than b10434 the flag does not parse — use --chat-template-kwargs '{"reasoning_effort":"medium"}' instead.

One default worth knowing about even though it does not change this command: llama.cpp now fits parameters to free device memory unless you pass --fit off. common/fit.h documents the scope — "only parameters that have the same value as in llama_default_model_params are modified", with context size changed only if it is left at 0 — so an explicit -c and -ngl are untouched. Its sibling flag is useful here precisely because our budget is derived rather than measured: --fit-print on asks llama.cpp for its own estimate of the memory the configuration needs, which you can compare against the table above.

Images

curl -s http://127.0.0.1:8080/v1/chat/completions -H "Content-Type: application/json" -d '{
  "messages": [{"role":"user","content":[
    {"type":"text","text":"What is in this image?"},
    {"type":"image_url","image_url":{"url":"data:image/jpeg;base64,'"$(base64 -w0 photo.jpg)"'"}}
  ]}]
}'

Does the vision path actually work on ROCm? Yes, and the evidence is specific rather than structural-only. On 2026-08-16, llama.cpp #27124 — an image-input crash report against this exact model — was closed after a second reporter ran the same weights and the same mmproj-BF16.gguf on a ROCm build and answered "I am able to input images and the model correctly describes their content", with a log showing the device as ROCm0; the original reporter switched from Vulkan to the ROCm build and confirmed the fix. Read the scope honestly: that machine is a Ryzen AI Max+ 395, a gfx1151 APU, not this discrete gfx1100 card, and it is a functional confirmation rather than a measurement. What underwrites the transfer is that llama.cpp's projector is not backend-specific code — tools/mtmd/clip.cpp contains no CUDA/HIP branches and initialises through the generic ggml_backend_init_by_type GPU path — so the projector's status here is one backend compiled twice, exactly like the rest of the tree. If it fails on your build, that is new information worth filing upstream and sending to /contribute.

Results

• Speed: no ROCm measurement of this model on this card exists, and this recipe will not manufacture one. What does exist is a Vulkan measurement, on this exact card and model: the reporter of llama.cpp #27097, running Qwen3.8-27B-UD-Q4_K_XL.gguf under llama-bench on an RX 7900 XTX (RADV NAVI31) on Linux at build 10438, recorded pp512 877.84 ± 1.13 t/s and tg128 37.09 ± 0.05 t/s, holding at 812.34 and 36.33 at a depth of 4096 — with the workaround in Troubleshooting applied. Three reasons not to read those as this page's numbers: a different backend from the one installed here, no vision projector loaded, and llama-bench rather than a served session. They are the right order of magnitude to sanity-check against, and nothing more. If you measure the ROCm path, please send it via /contribute so it lands on /check/qwen3-8-27b/rx-7900-xtx.
• AMD's own day-0 numbers are also Vulkan, and on other hardware. AMD's launch post for this model quotes up to 24.5 tok/s on a Ryzen AI Max+ 395 and up to 51.8 tok/s on a Radeon AI PRO R9700, under a footnote that reads "Testing as of August 2026 by AMD using preliminary performance results for Qwen 3.8 27B in llama.cpp on Windows with the Vulkan backend and MTP=4." Neither machine is this card. The same post's guidance on capacity is worth having, though, because it is the vendor's own: "This model can also run on supported AMD hardware with more than 24 GB of Variable Graphics Memory or VRAM", and for LM Studio it says the model "requires roughly 24GB of variable graphics memory". This card has exactly 24 GB, which is why the headroom line above is the one to watch.
• VRAM usage: ~21.2 GiB derived working set at 128K context (table above), against a device-local heap the card reports as 23.75 GiB. Derived from cited file sizes and llama.cpp's allocation formulas, not measured.
• Quality notes: this is a thinking model with xhigh reasoning effort by default; expect long traces on short questions unless you lower the effort. The card's non-thinking preset is temperature=0.7, top_p=0.80, top_k=20, presence_penalty=1.5.

For the full benchmark data, see /check/qwen3-8-27b/rx-7900-xtx.

Troubleshooting

Throughput collapses mid-request and nothing reports an error

This is the AMD failure mode that replaces "it OOMs at startup", and it is the reason to keep the headroom above. llama.cpp #26432 documents it on an RX 7900 XT running the previous generation of this same qwen35 architecture under HIP: the model loads cleanly, then the first large request pushes buffers into GTT and decode falls from 42–46 tok/s to 18. The reporter watched /sys/class/drm/card1/device/mem_info_gtt_used climb from 0.59 GB to 14.65 GB during a single prefill and fall back afterwards, and — having tested hipMalloc directly and found it fails cleanly rather than spilling — concluded that llama.cpp's load-time accounting misses runtime-peak buffers, so "on this platform the overage shows up as a silent GTT cliff instead of a load-time error". Watch for it:

watch -n1 'cat /sys/class/drm/card0/device/mem_info_gtt_used; amd-smi monitor --vram-usage'

If GTT moves during a request, you are over the line. Take back the context size or the quant tier; the fact that the model loaded proves nothing. Note the reporter's configuration also had multi-token prediction on, which is extra footprint this recipe does not spend.

Prompt processing collapses after switching to a 4-bit KV cache

Not a CUDA-only trap, despite where it was reported. The gates live in fattn.cu, which your HIP build compiles from the same source: ggml_cuda_fattn_kv_type_supported returns false for Q4_1/Q5_0/Q5_1 under #ifndef GGML_CUDA_FA_ALL_QUANTS, and a few lines above, the same guard returns BEST_FATTN_KERNEL_NONE whenever K->type != V->type. When no kernel is available the scheduler quietly relocates the attention op to the CPU backend, with no error printed. The default HIP build compiles exactly the same four FlashAttention vector instances as the default CUDA one — f16-f16, q4_0-q4_0, q8_0-q8_0 and bf16-bf16 — as the else branch of ggml-hip/CMakeLists.txt shows. It was measured on an RTX 3090 in llama.cpp #27109, where prefill fell from 991–1276 t/s at q8_0/q8_0 to 34–106 t/s with K=q4_1, V=q8_0; read those figures as the size of the cliff on another vendor's card, not as throughput to expect here. The rule that survives is match the two types. This recipe's q8_0/q8_0 is in the default set, so no flag and no newer build is needed; if you want the exotic combinations, the fix is a compile flag rather than a version bump:

cmake -S . -B build -DGGML_HIP=ON -DGPU_TARGETS=gfx1100 -DGGML_CUDA_FA_ALL_QUANTS=ON

Token generation is 2.7× slower than it should be on the Vulkan backend

If you run the Vulkan build instead of the ROCm one, check Resizable BAR in your BIOS. The reporter of llama.cpp #27097 measured this model at tg128 13.89 t/s on this card with ReBAR disabled and 37.09 t/s with GGML_VK_DISABLE_HOST_VISIBLE_VIDMEM=1 set, on the same build and file — the small host-visible VRAM window becomes the bottleneck. A candidate patch to detect the case automatically is linked in that thread and was not merged when this page was written. The HIP path in this recipe does not use that heap, so the workaround is Vulkan-specific.

Images crash the server

Which backend you are on matters more than which card you have. Every image-input crash report found for this model is on a Vulkan build (#27124, on an AMD Ryzen AI MAX+ 395 under Windows), and that same thread was closed by moving to a ROCm build. If you hit an image crash, try the ROCm binary before you try a different quant.

llama-server --list-devices prints nothing on Windows

A packaging defect in the Windows ROCm release zip, not your driver: llama.cpp #26996 shows ggml-hip.dll importing hipblas.dll, which the win-rocm-7.14-x64 archive does not ship, so device enumeration silently returns none and exits 0. A project contributor's answer in that thread is to install ROCm 7.14 from AMD's Windows tarball and complete its post-installation environment-variable steps; a reporter who copied the whole ROCm bin\ directory next to the binaries got his GPU recognised. Linux builds from source are unaffected.

Ollama, if you would rather not build anything

qwen3.8:27b is an 18 GB tag and fits this card. Two AMD-specific conditions: Ollama's documentation lists the RX 7900 XTX among supported cards but requires the AMD ROCm v7 driver on Linux, installed with AMD's amdgpu-install utility — it does not bundle ROCm. And the tag defaults are not what you may expect: qwen3.8:27b and qwen3.8:27b-mtp-q4_K_M resolve to byte-identical manifests on the Ollama registry, whose params blob carries "draft_num_predict":4, so you get multi-token prediction whether or not you asked. The only no-speculation tag is qwen3.8:27b-q4_K_M. Ollama will also not give you 131,072 tokens of context by default — raise num_ctx explicitly.

Multi-token prediction: leave it off here

The model ships an MTP head inline in unsloth's file (block_count = 65, nextn_predict_layers = 1), so --spec-type draft-mtp needs no extra download, and on AMD hardware it is genuinely fast in the right conditions — AMD's own launch numbers use it. But every acceptance-rate figure found on a ggml-cuda-family backend is poor: the one ROCm log of this model, in #27124, records draft acceptance = 0.35929, and the CUDA reports in #26750 sit in the same band against roughly 92% on Vulkan. Enabling it also costs VRAM — llama_memory_recurrent allocates mem_size * (1 + n_rs_seq) rows, so a draft length of 4 takes the recurrent line from 156,893,184 to 784,465,920 bytes — and adds a 17th attention layer's worth of KV, since llama.cpp gives the MTP block a plain attention cache rather than the hybrid one. On a budget with 2.805 GiB spare that is worth measuring before you commit to it.