Introduction

Work in progress: This wiki is being actively expanded from Discord research and fork release notes. Many tuning recipes and throughput numbers are community-reported and have not been independently reproduced here. See Verification status for what is solid vs what still needs checking on your hardware.

Welcome to the GFX1030 Wiki — a focused, hands-on knowledge base for running LLM inference on AMD gfx1030 (RDNA2 / Navi 21) GPUs, with a strong bias toward the Radeon PRO V620 and the purpose-built tooling collected here.

gfx1030 is the LLVM/ROCm target for the Navi 21 “Sienna Cichlid” die. It powers the consumer Radeon RX 6800 / 6800 XT / 6900 XT / 6950 XT and the workstation/data-center PRO W6800 and PRO V620. These cards are officially supported by ROCm on Linux, which makes them a cost-effective platform for modern LLMs — but getting the most out of them (power tuning, multi-GPU P2P, and RDNA-tuned kernels) takes a bit of extra work. That’s what this wiki documents.

What this wiki focuses on

• Tuning — the v620_toolbox recipes: lowering the V620’s VBIOS-locked 250 W floor to 120 W, enabling GPU↔GPU PCIe Peer-to-Peer between multiple V620s (Fedora + AMD hosts; power tuning also on Ubuntu 26.04; Intel Ice Lake often sees no P2P win — A/B it), and optionally disabling Pro-card ECC for ~2 GB extra VRAM.
• llama.cpp — stock builds plus the RDNA2 fork (tensor parallel, MTP / DFlash). Usually the easiest first stack on V620.
• vLLM on RDNA — ready-to-run Docker images; recipes for Hub vs recipe container vs Flash-Next (blivioniag/vllm-rdna on a blivioniag/rocm-rdna PyTorch base), how they are built with vllm-rdna-docker, and the official vLLM extras fork opengfx1030/vllm-rdna (rdna_extras) plus the separate Flash-Next fork (better for multi-stream / agentic loads once llama.cpp is working).

New here? Start with Supported Hardware and Getting Started.

Community Discord

Much of this wiki is distilled from the gfx1030 club Discord — live tuning tips, fork releases, and multi-GPU benches land there first.

• Server: gfx1030 club
• Join: https://discord.gg/mESex2aBp
• Useful channels: #vllm-rdna, #llamacpp, #general, #benchmarks, #harnesses, #lmcache

When turning Discord threads into docs, follow the privacy rules in Contributing.

Disclaimer: This is a community wiki, not affiliated with or endorsed by AMD. Kernel patches and power-cap changes are done at your own risk. Always cross-check against the official ROCm documentation. Contributions welcome — see Contributing.

Supported Hardware

gfx1030 is the LLVM target name for AMD’s Navi 21 GPU (RDNA2 generation, codename Sienna Cichlid). Several retail and workstation cards use this die and therefore report themselves as gfx1030.

Cards that natively report as gfx1030

Pro cards hide ~2 GB behind ECC. See Disabling ECC if you want the full 32 GB.

Do not flash W6800 VBIOS onto a V620. #general (Sep 2026): pre-modded eBay cards and DIY flashes that load a W6800 BIOS on a V620 drop active CUs from 72 → 54 (W6800’s different CU layout). That is a large compute hit — stay on stock V620 BIOS. Cooling shroud mods are fine; BIOS swaps are not a free “Pro unlock.”

Need OverDrive / unlocked clocks in a passthrough VM without flashing? Prefer the reversible 4-byte PowerPlay soft unlock (Tamalero/amd-v620-soft-unlock) over a W6800 flash — see Power tuning → Soft unlock.

All of these are on the officially supported list for recent ROCm releases on Linux.

AMD BC-250 (gfx1013)

Headless mining APU board (Cyan Skillfish). RDNA2, reports as gfx1013 — not Navi 21 / gfx1030. Community ROCm and Vulkan work exists; most of this wiki still targets Navi 21.

Community references: AMD BC-250 docs, akandr/bc250-rocm.

RDNA2 relatives that can run gfx1030 code

The rest of the RDNA2 line uses a different LLVM target but shares the same ISA family. They are not on the official support matrix, but in practice they run gfx1030 kernels once you set HSA_OVERRIDE_GFX_VERSION=10.3.0 (see HSA_OVERRIDE for RDNA2 Cousins).

Architecture highlights (Navi 21)

• RDNA2 compute units with a native wavefront size of 32 (wave32), unlike GCN’s wave64.
• Infinity Cache (128 MB on Navi 21) that dramatically raises effective memory bandwidth.
• No dedicated matrix/tensor cores — RDNA2 predates the WMMA/matrix instructions added in RDNA3 (gfx11xx). Matrix math runs on the regular vector ALUs, so expect lower peak throughput than RDNA3 or CDNA cards, but very good price/performance for inference.
• FP16 is well supported; BF16 has limited/emulated support and is best avoided for hot paths (prefer float16).

How to check your GPU target

rocminfo | grep -i 'gfx\|Name'
# or, more directly:
rocminfo | grep -m1 -o 'gfx[0-9]*'

If the output shows gfx1030, everything in this wiki applies directly. If it shows gfx1031, gfx1032, etc., head to HSA_OVERRIDE for RDNA2 Cousins first.

Getting Started

This page gets you from a bare Linux install to a working ROCm + PyTorch stack on a gfx1030 card, and explains how to preview this wiki locally.

1. Confirm your hardware

Make sure your card is a Navi 21 gfx1030 (or a related RDNA2 card):

lspci | grep -i vga
# After ROCm is installed:
rocminfo | grep -m1 -o 'gfx[0-9]*'

See Supported Hardware for the full card list.

2. Install ROCm

Follow Installing ROCm. At a high level:

# Add the amdgpu repo, then:
sudo apt install rocm
sudo usermod -aG render,video "$LOGNAME"
# Reboot, then verify:
rocminfo
clinfo | grep -i 'gfx\|Board'

3. (Optional) Tune the card

If you run one or more Radeon PRO V620 (or PRO W6800) cards, these tweaks are worth doing before you load models (Fedora or Ubuntu 26.04 for power cap; P2P is Fedora + AMD-validated — A/B on Intel, see When P2P helps):

• Power Tuning — drop the VBIOS-locked 250 W floor to 120 W and boot-cap at 180 W. Passthrough VM and need OD clocks instead? See Soft unlock (amd-v620-soft-unlock).
• Disabling ECC — Pro cards hide ~2 GB behind ECC; optional extra VRAM.
• Multi-GPU PCIe P2P — enable GPU↔GPU peer-to-peer for multi-card setups.

Baremetal power floor + P2P come from the v620_toolbox repo.

4. Run an inference stack (Docker)

The fastest path is the prebuilt images — no local ROCm/PyTorch/vLLM build required:

• Running vLLM (Docker) — blivioniag/vllm-rdna on a blivioniag/rocm-rdna PyTorch base.
• vLLM recipes — Hub -extras vs recipe container vs Flash-Next by card count.

Prefer GGUF and building from source? See Building & Running llama.cpp (ROCm or Vulkan).

Want to build the images yourself, or use the RDNA-tuned kernels? See Building the Images and The rdna_extras fork.

5. Smoke test

# Inside a rocm-rdna / vllm-rdna container, or a local ROCm PyTorch env:
python -c "import torch; print(torch.cuda.is_available(), torch.cuda.get_device_name(0))"

You should see True and your Radeon card’s name. If not, see Troubleshooting.

6. Join the community Discord

Questions, live benches, and fork updates land first on gfx1030 club: https://discord.gg/mESex2aBp (#vllm-rdna, #llamacpp, #general, #benchmarks). See Useful resources.

Previewing this wiki locally

This site is built with mdBook. You do not need a GPU or ROCm to work on the docs.

Prerequisites

• Git — to clone the repository.
• mdBook — a single static binary (no runtime dependencies).

Install mdBook

The quickest way is to grab a prebuilt binary from the mdBook releases page:

mkdir -p "$HOME/.local/bin"
MDBOOK_VERSION=v0.5.4
curl -sL "https://github.com/rust-lang/mdBook/releases/download/${MDBOOK_VERSION}/mdbook-${MDBOOK_VERSION}-x86_64-unknown-linux-gnu.tar.gz" \
  | tar -xz -C "$HOME/.local/bin"
export PATH="$HOME/.local/bin:$PATH"
mdbook --version

If you have a Rust toolchain, cargo install mdbook also works.

Build and preview

git clone https://github.com/blivioniag/gfx1030-wiki.git
cd gfx1030-wiki

mdbook build      # outputs static HTML to ./book
mdbook serve      # live-reloading preview at http://localhost:3000

Add content

1. Create a new markdown file in src/, e.g. src/my_page.md.
2. Add an entry for it in src/SUMMARY.md.
3. Re-run mdbook serve to preview.
4. Open a pull request against master.

See Contributing for the full guidelines.

Installing ROCm

gfx1030 (Navi 21) is on AMD’s officially supported list for recent ROCm releases on Linux, so in most cases a stock install “just works” — no override hacks required.

Always follow the official ROCm install guide for your exact distro and ROCm version. The commands below are a convenience summary and may drift over time.

Supported operating systems

ROCm’s RDNA2 support targets specific LTS releases. As of recent ROCm versions that typically means:

• Ubuntu 22.04 / 24.04 LTS
• RHEL / Rocky 9.x
• Debian 12 (community, less tested)
• Fedora 43 (wiki-validated for power / P2P); Fedora Server 44 + kernel 6.19 reported working for the V620 powerfix (community)

Check the system requirements page for the version you plan to install.

Ubuntu quick install (amdgpu-install)

# 1. Add the amdgpu package repository (replace VERSION with the ROCm release you want, e.g. 6.4.60400-1)
sudo apt update
wget https://repo.radeon.com/amdgpu-install/latest/ubuntu/jammy/amdgpu-install_VERSION_all.deb
sudo apt install ./amdgpu-install_VERSION_all.deb
sudo apt update

# 2. Install ROCm (compute use case)
sudo amdgpu-install --usecase=rocm

# 3. Add yourself to the render/video groups so you can access the GPU without root
sudo usermod -aG render,video "$LOGNAME"

# 4. Reboot
sudo reboot

Verify the install

rocminfo                       # should list your Navi 21 card and "Name: gfx1030"
clinfo | grep -i 'Board\|gfx'  # OpenCL view
rocm-smi                       # live clocks, temps, VRAM, power

If rocminfo shows Name: gfx1030, you’re done. If it shows gfx1031/gfx1032/etc., your card is a smaller RDNA2 die — continue to HSA_OVERRIDE for RDNA2 Cousins.

Multi-GPU: pin ROCm 7.2.0 or 7.14.0

For more than one card, stay on ROCm 7.2.0 (not 7.2.1+) or jump to 7.14.0. #vllm-rdna / #llamacpp report an RCCL bug from 7.2.1 upward (including mid-line builds such as 7.2.4) that shows up as soon as you leave a single GPU — tensor-parallel hangs, comm failures, AllReduce init failed, or cards dropping offline. The published vllm-rdna images already sit on those two bases for that reason.

Do not “upgrade within 7.2.x” on a multi-GPU box. If you are already on a broken 7.2.1–7.13 userspace, rebuild or pull a 7.2.0 or 7.14.0 image rather than debugging RCCL on the in-between releases. ROCm 10 / TheRock packaging is interesting for newer cards; community llama.cpp benches on Ice Lake have run ROCm 10.0 with numbers close to 7.1 on that host, but day-to-day advice for gfx1030 multi-GPU is still: prefer 7.14.0 (or 7.2.0) unless you are deliberately lab-testing 10.x.

Notes & gotchas

• Kernel driver: ROCm relies on the amdgpu kernel module. Very new kernels sometimes ship a driver newer than your ROCm userspace expects; the amdgpu-dkms package from the amdgpu repo keeps them in sync.
• Secure Boot: if Secure Boot is enabled, the DKMS module must be signed or it will fail to load. The easiest path for a dev box is to disable Secure Boot.
• Multiple GPUs / iGPU present: if your CPU also has an RDNA2 iGPU, ROCm may enumerate it. Pin the discrete card with HIP_VISIBLE_DEVICES (see Reference).
• hipBLASLt: some libraries assume hipBLASLt, which historically had gaps on Navi 21. If a workload complains, try TORCH_BLAS_PREFER_HIPBLASLT=0 (PyTorch) or the workload’s equivalent flag.

HSA_OVERRIDE for RDNA2 Cousins

Only Navi 21 cards report as gfx1030. The smaller RDNA2 dies use different LLVM targets:

Many ROCm libraries only ship precompiled kernels for a subset of targets. Because every RDNA2 card shares the same instruction set family, you can tell ROCm to treat your card as gfx1030 and reuse the gfx1030 kernels.

The override

export HSA_OVERRIDE_GFX_VERSION=10.3.0

10.3.0 maps to gfx1030. Set it in the shell (or systemd unit / container env) before launching any HIP/ROCm program:

HSA_OVERRIDE_GFX_VERSION=10.3.0 python my_inference_script.py

To make it permanent for your user:

echo 'export HSA_OVERRIDE_GFX_VERSION=10.3.0' >> ~/.bashrc

Why this works (and its limits)

• RDNA2 GPUs (gfx1030–gfx1036) are binary-compatible enough that gfx1030 kernels execute correctly on the smaller dies for the vast majority of ML ops.
• It is still a workaround. AMD does not officially validate it, and you may hit edge cases — particularly in hand-tuned assembly kernels or libraries that query the exact arch at runtime.
• Performance-tuned kernels (e.g. in rocBLAS/Tensile) were tuned for Navi 21’s CU count and cache; on a smaller die they run correctly but may be sub-optimal.

Building instead of overriding

If you compile a library yourself, prefer building for your real target so you get correctly-tuned kernels, e.g.:

# Build for multiple RDNA2 targets at once
export PYTORCH_ROCM_ARCH="gfx1030;gfx1031;gfx1032"
# or for a single card
export AMDGPU_TARGETS=gfx1032

For llama.cpp:

cmake -B build -DGGML_HIP=ON -DAMDGPU_TARGETS=gfx1030 -DCMAKE_BUILD_TYPE=Release

Verifying it took effect

HSA_OVERRIDE_GFX_VERSION=10.3.0 python -c \
  "import torch; print(torch.cuda.is_available(), torch.cuda.get_device_name(0))"

If PyTorch previously errored with "no kernel image is available for execution on the device" and now prints True, the override is working.

Power Tuning (120 W floor)

The Radeon PRO V620 (gfx1030) ships with its power limit floor locked to 250 W in the VBIOS. That’s wasteful for inference, where the card spends most of its time memory-bound. This page summarizes how to unlock a 120 W floor and apply a 180 W boot cap, following the powertuning/ feature of the v620_toolbox repo.

Validated on Fedora 43 (kernels 6.17.6 pre-7 / 7.1.7 post-7) and Ubuntu 26.04 LTS (kernel 7.0.0-30-generic). Community confirmation: Fedora Server 44 + kernel 6.19, 4× V620, powerfix + 180 W cap (cap_min=120 W). This involves patching and rebuilding a kernel module — do it at your own risk.

Why it’s needed

The V620’s VBIOS declares 250 W as its minimum power limit, and the amdgpu driver trusts that number, so any lower cap is rejected:

echo 180000000 | sudo tee /sys/class/hwmon/hwmonX/power1_cap
# -> Invalid argument

The SMU firmware actually accepts values down to 120 W — only the kernel is in the way.

The fix

A tiny patch to sienna_cichlid_get_power_limit() in drivers/gpu/drm/amd/pm/swsmu/smu11/sienna_cichlid_ppt.c clamps the reported minimum to 120 W — but only when the GPU matches the V620 reference board by PCI identity:

• PCI device 1002:73a1, subsystem 1002:0e34

It does not touch the VBIOS, pp_table, or pp_features, and works on any kernel ≥ 5.15 with any number of V620s in the host. The canonical patch is patches/v620-powercap-min-120W.patch (portable across pre-7 and post-7 kernels via noinline).

+		if (smu->adev->pdev->vendor == 0x1002 &&
+		    smu->adev->pdev->device == 0x73a1 &&
+		    smu->adev->pdev->subsystem_vendor == 0x1002 &&
+		    smu->adev->pdev->subsystem_device == 0x0e34)
+			sienna_cichlid_v620_min_powercap_fix(smu, min_power_limit);

Other gfx1030 boards (RX 6900 XT / 6800 use device 0x73bf) have different PCI IDs. To power-tune those, change the identity match in the patch accordingly.

Platform paths

Both platforms share the same 120 W floor patch logic, the same v620-cap-apply.sh runtime script, and the same v620-powercap.service boot cap. Follow the README in the folder for your distro.

Ubuntu quick path

git clone https://github.com/blivioniag/v620_toolbox.git
cd v620_toolbox/ubuntu_powertuning

# Install deps (see ubuntu_powertuning/README.md), then:
sudo cp v620-rebuild-amdgpu /usr/local/sbin/
sudo cp ../powertuning/scripts/v620-cap-apply.sh /usr/local/sbin/
sudo chmod 755 /usr/local/sbin/v620-rebuild-amdgpu /usr/local/sbin/v620-cap-apply.sh

sudo /usr/local/sbin/v620-rebuild-amdgpu "$(uname -r)"
sudo reboot

After reboot: sudo dmesg | grep -i 'V620 powerfix' and ../powertuning/scripts/v620-verify.sh. Install the systemd unit and optional kernel/postinst.d/v620-amdgpu hook so future kernel updates rebuild the module — see ubuntu_powertuning/README.md.

Secure Boot: disable it or sign the rebuilt module — unsigned overrides won’t load with SB on.

Two ways to apply it (Fedora)

The toolbox provides scripts under powertuning/scripts/:

1. Bake a kernel RPM — v620-kernel-bake.sh builds a Fedora dist-git kernel RPM with the patch (and the P2P kernel-config delta) baked in. Survives cleanly across reboots.
2. Out-of-tree module override — v620-module-install.sh builds a patched amdgpu.ko against your running kernel and installs it as an override. Re-run after every kernel update.

Then apply the runtime cap at boot:

• scripts/v620-cap-apply.sh — writes the 180 W cap (matches the V620 by PCI ID).
• systemd/v620-powercap.service — a oneshot unit that runs v620-cap-apply.sh 180 at boot.

Verify

# The kernel logs the fix on match:
sudo dmesg | grep -i 'V620 powerfix'

# The reported minimum is now 120 W (120000000 µW):
cat /sys/class/hwmon/hwmon*/power1_cap_min

# The toolbox's own check:
sudo ./powertuning/scripts/v620-verify.sh

v620-verify.sh confirms both the V620 powerfix dmesg marker and power1_cap_min=120000000. After the override, idle around ~7 W per card has been reported (Fedora 44, 4× V620). If idle is still ~50 W, the floor/cap path is not active — re-run verify.

Token cost vs stock 250 W

#llamacpp: at the stock 250 W floor, V620 inference can look ~15% more expensive per token than a high-end Nvidia reference; with a 180 W cap after the floor unlock, the same report called V620 ~20% cheaper per token, with falloff “undetectable” down to ~200 W and only small percentages below that. Cap at 160 / 140 W mainly for PSU headroom, not because 180 W is slow. #general (Sep 2026) similarly calls ~120–170 W the practical sweet spot — little gained above that for the extra watts.

Undervolt experiments exist; community: about −25 mV was stable for one host, while more aggressive settings started emitting random tokens. Treat UV as optional and A/B carefully.

Slot power and PSU transients

Unlike many gaming cards, a V620 reaches TDP from the PCIe slot, not from extra 8-pin cables. That makes the +12 V rail and slot power delivery more sensitive than a 3080-class swap-in:

• Prefill / tensor-split can trip old or miner PSUs even when average watts look fine — see llama.cpp PSU troubleshooting.

• If a new card hard-reboots the host on load, try a gentler SMU ramp before blaming the kernel:

  sudo rocm-smi --setperflevel standard   # community: ~80 W min / ~150 W max, softer ramp
  

  (STANDARD / standard — check rocm-smi --help on your ROCm; the enum name varies.)

• Cap at 160 W or 140 W if 180 W still trips protection.

Multi-card PSU sizing (community)

#forum / #general (Sep 2026): an 8× V620 @ 180 W build targets ~1440 W for GPUs alone, leaving headroom for CPU/board. One reported pick that avoids sketchy splitters: SilverStone HELA 2050 Platinum (SST-AX2050MCPT-A). US hosts often need a 240 V circuit for that class of load — stock 250 W × 8 on a 120 V / 20 A breaker is a non-starter. Prefer the 120–180 W caps for both thermals and wall power.

See the full recipe, prerequisites, and the deep-dive docs (docs/POWERCAP.md) in the repo.

Soft unlock (passthrough VMs)

Different problem, different tool: Tamalero/amd-v620-soft-unlock unlocks OverDrive clock + power controls on a V620 that is PCIe-passthrough into a Linux VM (Proxmox / QEMU / libvirt). Stock passthrough often leaves OverDrive empty (LACT / CoreCtrl / sysfs), a fixed ~1825 MHz core pstate, and a hard 250 W power cap.

How soft unlock works (summary — follow the upstream README):

1. Guest kernel: amdgpu.ppfeaturemask=0xffffffff
2. Dump the card’s own VBIOS from the guest (/sys/kernel/debug/dri/*/amdgpu_vbios)
3. Patch with make_odcaps_rom.py (sets OD caps 0–3 + ATOM checksum; no ROMs shipped)
4. Attach as QEMU romfile= / Proxmox hostpci…,romfile= / libvirt <rom file=…/>
5. Verify pp_od_clk_voltage shows an OD_RANGE, then tune with sysfs / tools/v620 / LACT

Warnings from upstream (do not skip):

• Never write guest pp_features or sysfs pp_table on this card — can wedge the SMU; V620 has no FLR, so recovery is a host reboot.
• Soft unlock is per-VM (romfile=). Nothing is written to the physical flash; remove romfile= to revert.
• Passive baremetal ACPI VFCT delivery path is documented upstream as experimental / untested.
• Passive server cards need real airflow before raising clocks / 275 W.

#general: community reports of reliable use under Proxmox passthrough (weeks-long). This wiki has not independently validated TFLOPS or power-range claims — treat benches in the upstream README as community-reported.

Disabling ECC (Pro VRAM)

WIP: Procedure is documented for Radeon PRO W6800 in lunnova’s guide. #vllm-rdna and #general report the same ECC-on-by-default behavior on V620. Treat the two-reboot kernel-param path as community until you confirm rocm-smi on your own cards. See Verification status.

Workstation/data-center Navi 21 cards (PRO V620, PRO W6800) ship with on-board ECC enabled. That costs about 7% of VRAM (~2 GB on a 32 GB card). Consumer RX 6800 / 6900 cards have no ECC and already show the full 16 GB.

#vllm-rdna (pinned): extra ~2 GB per GPU if you disable it. Some members keep ECC on for bit-flip protection; others disable it for long-context KV cache.

Check whether ECC is on

sudo dmesg | grep -i ecc
# "MEM ECC is active" / "GECC is enabled"  → ECC is on
# "MEM ECC is not presented"               → consumer card / no ECC

rocm-smi --showmeminfo vram
# ~30.7 GB total on a 32 GB Pro card → ECC is eating the rest

Linux: amdgpu.ras_enable=0 (two reboots)

There is no Radeon control-panel toggle on Linux. The community path (lunnova, 6.x kernels) is:

1. Add the kernel parameter amdgpu.ras_enable=0.
2. Reboot twice, leaving the parameter in place both times.

GRUB (Ubuntu / Debian / Fedora):

# Ubuntu/Debian: edit GRUB_CMDLINE_LINUX_DEFAULT in /etc/default/grub
# Fedora: sudo grubby --update-kernel=ALL --args="amdgpu.ras_enable=0"
sudo nano /etc/default/grub
# …add amdgpu.ras_enable=0 to GRUB_CMDLINE_LINUX_DEFAULT, then:
sudo update-grub   # Ubuntu/Debian
# Fedora uses grubby above instead of update-grub
sudo reboot

After the first reboot, dmesg should mention that GECC will be disabled on the next boot:

GECC will be disabled in next boot cycle if set amdgpu_ras_enable and/or amdgpu_ras_mask to 0x0

Reboot again. On the second boot you want:

amdgpu: GECC is disabled

and rocm-smi --showmeminfo vram should report the full ~32 768 MiB.

To turn ECC back on: remove the kernel parameter and reboot twice.

An older amdgpu-no-ecc.patch is not needed on 6.x kernels. Keep it only if you are still on 5.x and the two-reboot path does nothing — details in lunnova’s article.

Notes

• Consumer gfx1030 (RX 6800 / 6800 XT / 6900 XT / 6950 XT): skip this page.
• Secure Boot / signed kernels: same caveats as power tuning — unsigned module overrides will not load.
• This does not replace the 120 W power-cap patch. ECC and the VBIOS power floor are independent.

Related

• Supported Hardware
• Power Tuning
• Getting Started

Multi-GPU PCIe P2P

For multi-card V620 rigs, enabling GPU↔GPU PCIe Peer-to-Peer (P2P) lets the GPUs DMA directly to each other’s VRAM instead of bouncing through system memory — important for tensor/pipeline parallel serving. This page summarizes the pcie_p2p/ feature of the v620_toolbox repo.

Validated end-to-end on Fedora 43 + AMD CPU (EPYC 7452), 4× Radeon PRO V620. Result: rocminfo enumerates 5 HSA agents and amd-smi topology shows 12/12 GPU↔GPU P2P ENABLED.

Power-cap tuning on Ubuntu 26.04 is supported separately via ubuntu_powertuning/ — P2P recipes here remain Fedora-validated for now. Community Intel notes (Ice Lake) are in When P2P helps — and when it does not.

Two validated kernel paths

Both kernels can be installed on the same host simultaneously; toggle the default with grubby --set-default /boot/vmlinuz-<evr>. Both boot with the 180 W power cap from Power Tuning.

Key prerequisites

• AMD CPU — the KFD P2P path is wiki-validated on AMD (EPYC 7452). Intel Ice Lake can enumerate P2P (community) but often does not speed up inference — see below.
• ≥ 2 Radeon PRO V620 for a meaningful P2P topology (bench validated on 4).
• No pre-gfx1030 AMD GPU installed alongside the V620. ROCm 7.x dropped support for older ASIC families (gfx8xx Polaris, gfx900 Vega, …). If any pre-gfx1030 AMD GPU is present at boot, KFD registration fails and rocminfo bails with Failed to map remapped mmio page on gpu_mem 0. This only matters if your host actually has older AMD cards — remove them first.

Identify your V620 reference boards:

lspci -nn | grep '1002:73a1'    # one line per V620, subsystem 1002:0e34

The full 4-tuple 1002:73a1:1002:0e34 uniquely identifies the V620 reference board and is what the toolbox scripts match on. Other gfx1030 boards (RX 6900 XT / 6800) use device 0x73bf and different subsystem IDs — they’d need the identity match adjusted.

Readiness & verification

The toolbox ships diagnostics that gate on kernel config, hardware, ACS/IOMMU, and runtime state:

# Four-gate readiness check (kernel + hardware + ACS/IOMMU + runtime)
sudo ./powertuning/scripts/v620-p2p-readiness.sh

# End-to-end P2P verification
sudo ./pcie_p2p/scripts/verify-p2p.sh

A healthy system shows:

rocminfo | grep -c '^  Name:.*gfx1030'     # one per V620
amd-smi topology                            # all GPU<->GPU pairs: P2P ENABLED

For the full recipe (kernel config deltas, BIOS/IOMMU settings, what can go wrong), read pcie_p2p/README.md and the knowledge base powertuning/docs/AMD_P2P.md.

When P2P helps — and when it does not

amd-smi topology saying P2P ENABLED is not the same as faster tokens. Always A/B with NCCL_P2P_DISABLE=1. Community reports (#general, Ice Lake 4× V620 host):

If bandwidth tests pass but inference regresses:

export NCCL_P2P_DISABLE=1          # llama.cpp / RCCL tensor parallel

On the RDNA2 fork, the README also documents a flag to disable P2P all-reduce fusion (GGML_HIP_GFX1030_P2P_ALLREDUCE=off). Verified ~25 GB/s P2P between V620 pairs does not guarantee faster token generation if links are narrow or the host is a single monolithic PCIe root.

Intel: do not disable IOMMU to “help” P2P

On Ice Lake, disabling VT-d / IOMMU broke P2P even though some generic docs say IOMMU-off is more permissive. Leave IOMMU on, then run the toolbox readiness script and ACS checks. Fedora Server 44 + kernel 6.19 + the powerfix has been reported working on 4× V620 (community; not the Fedora 43 AMD validation).

PLX / PCIe switches

• Intra-switch P2P can stay full-width (e.g. 4× Gen4 x16 behind one PLX 88096). The host↔switch uplink (typically one x16) is the bottleneck.
• Without tensor parallel, a PLX box is usually slower than native CPU lanes: higher latency and less aggregate host bandwidth.
• Do not tensor-split across a daisy-chained pair of switches — that single inter-switch link is a TP bottleneck. Prefer TP inside one switch and pipeline-parallel between switches (vLLM can do TP+PP that way; llama.cpp generally cannot).
• Community: ACS often needs extra kernel cmdline fiddling; pcie surprise link down crashes were fixed by putting a small fan on the PLX heatsink (these boards often ship with no airflow notes).
• External PCIe expansion (#general, Sep 2026 — e.g. Cubix Xpander–class enclosures): each group of cards may get full-width slots behind the switch, while only one x16 (often PCIe 3.0) returns to the host. Community: 8× GPU on that pattern was fine at low context but fell to ~3 t/s around 50–60k context. Prefer 4 GPUs on one switch + P2P / TP inside the switch over stretching TP across a narrow host uplink.

Host topology

Community reports, not wiki-benched:

Layer split vs tensor split (quick)

Community explainer (#benchmarks thread, Sep 2026):

• Layer / pipeline split — VRAM grows with each card, but only one GPU works the model at a time (sequential stages). Less GPU↔GPU traffic.
• Tensor (row) split — all GPUs work together each step; needs much more interconnect bandwidth and often becomes PCIe-limited on narrow links.

On fat CPU roots (gen4 x16 × N), try tensor first. Behind a narrow switch uplink, prefer smaller TP domains or layer/PP splits.

Cabling / risers (community)

Not wiki-benched — common #general notes:

• SlimSAS / SFF-8654 cables: PCIe gen3 is usually fine at ~70 cm; gen4 needs testing per cable/insulation. Prefer known-good gen4 kits over the cheapest Amazon/eBay passive ribbon risers.
• Passive PCIe risers: several community reports of timeouts / inability to hold gen4 x16 with cheap passive risers under multi-GPU load, while the same slots work with cards seated directly. Brand-name gen4 risers (e.g. ADT-class) are often the next step — A/B one card first.
• V620 + blower shroud length: community measure ~37 cm with a common EFH-08E12W-style fan shroud installed — plan chassis / PLX slot spacing accordingly.

vLLM on RDNA — Overview

WIP: This section is actively expanded from Discord and fork release notes. See Verification status before treating benchmarks as gospel.

The quickest way to serve LLMs on gfx1030 / RDNA is the prebuilt blivioniag/vllm-rdna images on a blivioniag/rocm-rdna ROCm + PyTorch base. A single image targets seven RDNA architectures (gfx1030 through RDNA4).

Where to start

Image variants

Image tags are refreshed in place when fixes land — always docker pull before debugging. Confirm your -extras image includes the latest extras commits (AWQ dispatch, GDN HIP, TP graph fix). #vllm-rdna (Aug 31 2026): blivioniag/vllm-rdna:v0.27.1-extras was refreshed again — re-pull even if you already had that tag.

Official source moved. Kernel work lives in opengfx1030/vllm-rdna (rdna_extras). Day-to-day serving is still Hub blivioniag/vllm-rdna:*-extras (v0.27.1; bake still clones the historical personal fork). Flash-Next = leapdragon/vllm-rdna2-qwen until that line merges into the org. Details: Fork landscape.

Upstream vLLM 0.28.x: Official GPU docs still omit Navi 21 / gfx1030. Keep community forks until upstream documents it.

Qwen3.8 Flash-Next on vLLM

llama.cpp still struggles with Flash-Next on gfx1030 (upstream gaps). Community production path is the leapdragon/vllm-rdna2-qwen fork — not the Hub -extras image. Prefer that stack over llama.cpp for Flash-Next until the RDNA2 llama.cpp fork catches up. Flash-Next work is expected to land in opengfx1030/vllm-rdna after the 0.28 rebase / merge.

#vllm-rdna (Sep 2026) ballpark on 4× V620 (host-dependent; fork author + community):

Docs live under docs/rdna2/ (README.md, RESULTS.md, TROUBLESHOOTING.md, ROCR-CPU-FIX.md). Pull latest before re-benching — tags and container latest move with the prefill campaign.

Long-prompt stalls / timeouts: if large agentic prompts (tens of k tokens) hang or take many minutes while short prompts are fine, try VLLM_USE_V2_MODEL_RUNNER=0 — community report of stable ~68 t/s with dense INT8 + custom all-reduce after that switch; fork docs now call it out. See vLLM troubleshooting.

Intel AutoRound Flash-Next (draft, #vllm-rdna Sep 10–11 2026)

A second Flash-Next track is the Intel W4A16 AutoRound checkpoint (Intel/Qwen3.8-Flash-Next-W4A16-AutoRound; community also cites the RTN sibling Intel/Qwen3.8-Flash-Next-W4A16-RTN-AutoRound). Weights are ~75 GB; plan on 4× V620. Community quality notes: Intel AutoRound beats some other INT4 Flash-Next packs on published tool-calling scores (example cited: 80.5 vs 76.0 for a different INT4). Treat those as publisher / community numbers, not wiki benches.

Draft integration: opengfx1030/vllm-rdna#5 (open, not in Hub -extras). PR validation used FP16 activations, CPU PLE / n-gram offload, and --max-num-batched-tokens 4096. Published short-run figures on 4× V620 (no P2P in the Discord report):

Gotchas (do not treat as a drop-in Hub image):

• Known-good PLE is the embedded BF16 n-gram table (~95 GiB tensor data) with --engram-config '{"cpu_offload":true}'. Quantized CPU PLE (group-16 INT4 sidecar) is not end-to-end validated; #vllm-rdna saw incoherent generation on that path.
• --max-num-batched-tokens 2048 + chunked prefill was the suspected cause of the 128k PP cliff. Raise to 4096 (or see troubleshooting).
• Concurrent varlen GDN prefill still has a fork kernel addressing bug (gdn_prefill_o_rdna2.cu reuses a global chunk index). Single-stream benches can look fine.
• The PR is huge and dirty against rdna_extras — wait for a cleaned cherry-pick / image before calling it production. Day-to-day Flash-Next remains leapdragon/vllm-rdna2-qwen.

Also watch TheRock ROCR idle-CPU spin on ROCm 7.14 hosts.

llama.cpp vs vLLM

For a short comparison table (battle-tested GGUF vs agentic / Flash-Next), see llama.cpp overview.

What fits well on V620

#general (Sep 2026) consensus — RDNA2 has no matrix / tensor cores, so prefill on dense models is the weak spot (agentic “read a pile of files” workloads frustrate people even when decode looks fine):

Related

• Multi-GPU PCIe P2P — important for tensor parallel.
• Environment variables — cheat-sheet.
• Recipes (wiki) — pick Hub -extras vs recipe container vs Flash-Next by card count.
• Community recipe book (mirror opengfx1030/vllm-rdna2-recipe) — presets, patches, concurrent MTP PRs.

Running vLLM (Docker)

WIP: See Verification status.

Image matrix

Base images — blivioniag/rocm-rdna

These are a general-purpose ROCm PyTorch base for RDNA — useful on their own if you just want a working torch on a Radeon card.

vLLM images — blivioniag/vllm-rdna

The -extras tags use the rdna_extras fork lineage (historical bake still clones blivioniag/vllm rdna2_extras), which adds hand-written RDNA2 HIP kernels (FlashAttention, quantized GEMM, MoE, GDN, …). Check Docker Hub for the current tag list. Tags are refreshed in place — docker pull before debugging.

Multi-GPU: pick a 7.2.0 or 7.14.0 tag, not a host ROCm in the 7.2.1–7.13 gap. RCCL on those in-between releases is reported broken with more than one card — see Installing ROCm.

Every image bakes these PYTORCH_ROCM_ARCH targets: gfx1030;gfx1100;gfx1101;gfx1150;gfx1151;gfx1200;gfx1201.

Run it

docker run -it --rm \
  --device /dev/kfd --device /dev/dri \
  --group-add video --group-add render \
  --security-opt seccomp=unconfined \
  --ipc host \
  -p 8000:8000 \
  -v "$HOME/.cache/huggingface:/root/.cache/huggingface" \
  docker.io/blivioniag/vllm-rdna:v0.27.1-extras \
  vllm serve cyankiwi/Qwen3.8-27B-AWQ-INT4 \
    --dtype float16 \
    --max-model-len 8192 \
    --language-model-only --skip-mm-profiling --trust-remote-code

• Give the container the GPU with --device /dev/kfd --device /dev/dri and the video and render groups. Missing render is a common #vllm-rdna cause of Failed to infer device type.
• Prefer --dtype float16. RDNA2 has weak/emulated BF16; letting vLLM pick bf16 from a model’s config.json can trigger slow float32 fallbacks.
• The 27B AWQ example is the current #vllm-rdna day-to-day dense pick (Sep 2026). For a smaller smoke test, swap in any instruct model that fits VRAM. Model / card-count matrix: Recipes.
• For a non-Navi-21 RDNA2 card (gfx1031/1032/…), add -e HSA_OVERRIDE_GFX_VERSION=10.3.0. See HSA_OVERRIDE.
• Multi-GPU: add --tensor-parallel-size N; enabling PCIe P2P helps a lot here.

Query the OpenAI-compatible endpoint:

curl http://localhost:8000/v1/chat/completions \
  -H "Content-Type: application/json" \
  -d '{"model":"cyankiwi/Qwen3.8-27B-AWQ-INT4","messages":[{"role":"user","content":"Hi from gfx1030!"}]}'

Next steps

• Choose Hub vs recipe container vs Flash-Next: Recipes
• Tune env vars and CUDA graphs: Configuration
• Pick a quant format: Quantization
• Kernel details: rdna_extras fork

vLLM recipes (pick a path)

WIP: Stack picks and tok/s figures are community-reported (#vllm-rdna, Sep 2026). See Verification status. Prefer this page when you want “which image / which model / how many cards” — not a kernel deep-dive.

Discord keeps asking the same questions: which Docker stack, which Hugging Face weights, and whether 1× V620 is enough. This page consolidates those answers. Details for env vars and CUDA graphs live on Configuration; fork history on vLLM forks.

Which stack?

#vllm-rdna (Sep 2026): the recipe repo is treated as a parts pile (compose/env/patches against pristine vLLM 0.27.1). New Flash-Next work lives in vllm-rdna2-qwen. Official kernel development is opengfx1030/vllm-rdna rdna_extras — Hub -extras tags may lag HEAD until bake retargets.

Do not mix host ROCm userspace into the recipe container — the image carries its own stack. Mounting host ROCm into it is a common break (recipe TROUBLESHOOTING.md).

Pick by card count

Also see What fits well on V620.

Gemma 4 note

Community reports Gemma 4 ~26B still fails or is unfinished on current gfx1030 vLLM paths (#vllm-rdna, Sep 13 2026). Prefer Qwen / Ornith until someone posts a working recipe.

Model cheat sheet

Small-VRAM experiment: Ornith 9B EXL3 (~6.8 GB) vs AWQ (~9 GB) — experimental, not in published -extras tags yet. See Quantization.

Path A — Hub -extras (recommended default)

docker pull docker.io/blivioniag/vllm-rdna:v0.27.1-extras

docker run -it --rm \
  --device /dev/kfd --device /dev/dri \
  --group-add video --group-add render \
  --security-opt seccomp=unconfined \
  --ipc host \
  -p 8000:8000 \
  -v "$HOME/.cache/huggingface:/root/.cache/huggingface" \
  docker.io/blivioniag/vllm-rdna:v0.27.1-extras \
  vllm serve cyankiwi/Qwen3.8-27B-AWQ-INT4 \
    --dtype float16 \
    --max-model-len 8192 \
    --language-model-only --skip-mm-profiling --trust-remote-code

• Add --tensor-parallel-size N for multi-GPU. Enable PCIe P2P when possible.
• Prefer --dtype float16. Mount Triton / torch-compile caches for faster restarts — Configuration.
• Force the native W4A16 path when needed: VLLM_DISABLED_KERNELS=ExllamaLinearKernel,TritonW4A16LinearKernel — confirm Using RDNA2W4A16LinearKernel in logs.

Full env block and Compose (GPTQ + MTP): Configuration.

Hub -extras TP4 Qwen3.8-27B AWQ

Sanitized from a #vllm-rdna (Aug 31 2026) bench recipe on 4× V620 with working P2P / custom all-reduce. Drop the custom-AR block if P2P is broken on your board — use VLLM_DISABLE_CUSTOM_ALL_REDUCE=1 instead (Configuration).

export VLLM_USE_V2_MODEL_RUNNER=1
export VLLM_ROCM_USE_AITER=0
export VLLM_ROCM_USE_AITER_MOE=0
export FLASH_ATTENTION_TRITON_AMD_ENABLE=TRUE
export VLLM_RDNA_FORCE_FP16=1
export VLLM_USE_RDNA2_FA=1
export TORCH_BLAS_PREFER_HIPBLASLT=0
export PYTORCH_TUNABLEOP_ENABLED=1
export PYTORCH_TUNABLEOP_HIPBLASLT_ENABLED=0
export VLLM_WORKER_MULTIPROC_METHOD=spawn
export VLLM_BATCH_INVARIANT=0
export GPU_MAX_HW_QUEUES=2

# Only when GPU↔GPU P2P works:
export VLLM_FORCE_CUSTOM_ALL_REDUCE=1
export NCCL_P2P_LEVEL=pix
export RCCL_P2P_NET_DISABLE=1
export RCCL_P2P_BATCH_ENABLE=1
export NCCL_PROTO=Simple
export RCCL_MSCCL_ENABLE=0

cd /tmp   # avoid sys.path collisions with a local vllm checkout
vllm serve cyankiwi/Qwen3.8-27B-AWQ-INT4 \
  --port 8000 \
  --tensor-parallel-size 4 \
  --max-model-len 20480 \
  --max-num-seqs 8 \
  --gpu-memory-utilization 0.88 \
  --dtype float16 \
  --language-model-only --skip-mm-profiling --trust-remote-code \
  --enable-prefix-caching --enable-chunked-prefill \
  --compilation-config '{"cudagraph_mode": "FULL_AND_PIECEWISE", "compile_ranges_endpoints": []}'

Path B — Recipe container (27B / 122B presets)

Fastest path when you want the recipe book’s measured knobs without installing ROCm on the host:

IMG=ghcr.io/leapdragon/vllm-rdna2-recipe:0.27.1-rocm7.2.3-gfx1030
docker pull "$IMG"
docker run --rm "$IMG" list-presets

docker run -d --name vllm-rdna2 --network=host \
  --device /dev/kfd --device /dev/dri \
  --group-add "$(getent group render | cut -d: -f3)" \
  --group-add "$(getent group video | cut -d: -f3)" \
  --ipc=host --ulimit memlock=-1 --security-opt seccomp=unconfined \
  -e ROCR_VISIBLE_DEVICES=0,1 \
  -v "$HOME/.cache/huggingface:/root/.cache/huggingface" \
  "$IMG" preset:qwen38-27b-gptq

First boot cold-compiles for ~10–15 minutes; mount /compile-cache and Triton caches for ~3 min warm boots — see the recipe containers/README.md. Community ballpark on 2× V620: ~40–49 tok/s decode when TunableOp rows are seeded; ~27 t/s flat if they are missing (recipe troubleshooting).

122B builds need the repo wrapper / one-time weight prep — not a bare preset: one-liner.

Path C — Flash-Next (vllm-rdna2-qwen)

For agentic / long-context Flash-Next on 4× V620, use the dedicated fork and its docs — not Hub -extras and not the 27B recipe presets:

• Repo: leapdragon/vllm-rdna2-qwen
• Docs: docs/rdna2/
• Weights: wtdcode/Qwen3.8-Flash-Next-AWQ-W4A16
• PLE sidecar: primitive-ai/Qwen3.8-Flash-Next-PLE-quant

Expect large host DRAM for the n-gram / PLE store (~64 GB class called out in-channel). Long-prompt stalls: try VLLM_USE_V2_MODEL_RUNNER=0 — troubleshooting. Throughput snapshot: overview.

Related

• Running (Docker) — image matrix and minimal docker run
• Configuration — env vars, CUDA graphs, Compose
• Quantization — GPTQ/AWQ, KV, MTP
• vLLM forks — rdna_extras vs Flash-Next vs recipe ports
• Useful resources — external links

vLLM Configuration

WIP: Env recipes and throughput numbers are community-reported. See Verification status.

Recommended environment (#vllm-rdna)

These settings are commonly used in the gfx1030 Discord for -extras images on ROCm 7.14:

export VLLM_TARGET_DEVICE=rocm
export VLLM_ROCM_USE_AITER=0
export VLLM_ROCM_USE_AITER_MOE=0
export VLLM_RDNA_FORCE_FP16=1
export TORCH_BLAS_PREFER_HIPBLASLT=0
export PYTORCH_TUNABLEOP_ENABLED=0          # or 1 for autotuning (see compose below)
export PYTORCH_TUNABLEOP_HIPBLASLT_ENABLED=0
export GPU_MAX_HW_QUEUES=2
export VLLM_WORKER_MULTIPROC_METHOD=spawn
export VLLM_BATCH_INVARIANT=0
export HIP_FORCE_DEV_KERNARG=1
export RCCL_MSCCL_ENABLE=0
export VLLM_USE_RDNA2_FA=1                  # extras images: native RDNA2 FlashAttention
export VLLM_USE_V2_MODEL_RUNNER=1           # +17% vs V1 reported on gfx1030 `-extras`
export FLASH_ATTENTION_TRITON_AMD_ENABLE=TRUE
export SAFETENSORS_FAST_GPU=1               # faster safetensors → GPU load (`#vllm-rdna` Sep 2026)

VLLM_USE_V2_MODEL_RUNNER=1 remains the usual default on Hub -extras. On the Flash-Next fork, long agentic prompts that stall or timeout with V2 enabled have been fixed in-community by setting VLLM_USE_V2_MODEL_RUNNER=0 instead — see Flash-Next long-prompt stalls.

Custom all-reduce / P2P (two community stacks)

Pick one path — do not mix both:

If custom all-reduce misbehaves (stalls, bad latency on V620), fall back to the disable path. Some community recipes now run a short all-reduce self-test at boot and auto-fallback — pull latest recipe / fork notes.

Prefer --attention-backend RDNA_ATTN (or VLLM_USE_RDNA2_FA=1) over ROCM_ATTN on -extras. #vllm-rdna reports ROCM_ATTN sitting in AMD Triton flash-attention compile for hours (RCCL and Triton also fight each other). FA_RDNA2 may not show up on older -extras images or GPTQ models that still auto-select ROCM_ATTN — that is expected on hybrid GDN; see fork.

For multi-GPU TP on current -extras images, if AOT compile cache replay causes device-bound errors, add:

export VLLM_USE_AOT_COMPILE=0
export VLLM_DISABLE_COMPILE_CACHE=1

RDNA_ATTN / VLLM_USE_RDNA2_FA steer vLLM away from the generic AMD Triton flash-attention path, which can be slower or crash on some Qwen head sizes. See rdna_extras fork for kernel details. Full env cheat-sheet: Reference.

Don’t force backends or quantization unless you’re A/B testing — or avoiding a ROCM_ATTN Triton hang. Let vLLM read the model’s config.json unless that auto-selects the slow AMD Triton FA path.

CUDA graphs (preferred over --enforce-eager)

On current -extras images, CUDA graphs are the fast path — you generally should not use --enforce-eager. Graph capture can take a while on first boot (Triton JIT + torch-compile cache), but steady-state throughput is much higher once warmed up.

Recommended graph config (fork-author reported on Qwen3.8-27B-AWQ-INT4, TP4, 4× V620 — not wiki-reproduced):

vllm serve /path/to/model \
  --dtype float16 \
  --tensor-parallel-size 4 \
  --enable-chunked-prefill \
  --enable-prefix-caching \
  --language-model-only \
  --skip-mm-profiling \
  --trust-remote-code \
  --compilation-config '{"cudagraph_mode": "FULL_AND_PIECEWISE", "compile_ranges_endpoints": []}'

Reported steady-state throughput on current -extras (1024/512 bench, TP4):

Alternative that also avoids --enforce-eager:

--compilation-config '{"mode": "NONE", "cudagraph_mode": "FULL", "compile_ranges_endpoints": []}'

To disable graphs entirely (debugging only):

--compilation-config '{"cudagraph_mode": "NONE"}'

If graph capture still crashes, fall back to --enforce-eager — but try the updated image first (docker pull blivioniag/vllm-rdna:v0.27.1-extras-rocm7.14.0). See vLLM troubleshooting.

Cache volumes (first boot is slow)

Mount these so Triton and torch-compile artifacts persist across container restarts:

volumes:
  - ~/.cache/huggingface:/root/.cache/huggingface
  - ~/.triton/cache:/root/.triton/cache          # ~3 GB compiled Triton kernels
  - ~/.triton/dump:/root/.triton/dump
  - ~/.triton/llvm:/root/.triton/llvm
  - ~/.cache/vllm/torch_compile_cache:/root/.cache/vllm/torch_compile_cache  # ~700 MB

First startup compiles kernels and can take many minutes. Subsequent boots reuse the cache.

For recipe / Flash-Next stacks that want the compile cache (not the multi-GPU AOT-disable workaround above), #vllm-rdna (Sep 2026) recommends keeping the cache on and pointing it at a persistent directory:

export VLLM_DISABLE_COMPILE_CACHE=0
export VLLM_CACHE_ROOT=/path/to/persistent/vllm-cache   # mount this in Docker

Community reports ~5 min vs ~10 min subsequent startups once the cache is warm — still slow cold, but better than rebuilding every boot. Do not mix this with the TP VLLM_DISABLE_COMPILE_CACHE=1 workaround unless you have verified your image needs that disable path.

On Flash-Next long-context serves, if prefill falls off a cliff around 128k while decode stays flat, raise --max-num-batched-tokens to 4096 before blaming kernels — see 128k prefill cliff.

SAFETENSORS_FAST_GPU=1 is also commonly set (and already present in some vllm-rdna Dockerfiles) to speed weight load into GPU memory — see AMD vLLM optimization notes.

Docker Compose example (GPTQ + MTP + CUDA graphs)

This #vllm-rdna setup reached ~24 output t/s on TP2 with a GPTQ model. Key points: GPTQ hits the native RDNA2W4A16LinearKernel, CUDA graphs via --compilation-config, and VLLM_DISABLED_KERNELS forces the RDNA2 quant path:

services:
  server:
    image: blivioniag/vllm-rdna:v0.27.1-extras-rocm7.14.0
    network_mode: host
    ipc: host
    devices: [/dev/kfd, /dev/dri]
    group_add: [video, render]
    security_opt: [label=disable]
    volumes:
      - ~/.cache/huggingface:/root/.cache/huggingface
      - ~/.triton/cache:/root/.triton/cache
      - ~/.triton/dump:/root/.triton/dump
      - ~/.triton/llvm:/root/.triton/llvm
      - ~/.cache/vllm/torch_compile_cache:/root/.cache/vllm/torch_compile_cache
    environment:
      VLLM_TARGET_DEVICE: rocm
      VLLM_ROCM_USE_AITER: "0"
      VLLM_USE_RDNA2_FA: "1"
      PYTORCH_TUNABLEOP_ENABLED: "1"
      PYTORCH_TUNABLEOP_HIPBLASLT_ENABLED: "0"
      VLLM_WORKER_MULTIPROC_METHOD: spawn
      GPU_MAX_HW_QUEUES: "2"
      VLLM_USE_V2_MODEL_RUNNER: "1"
      VLLM_DISABLED_KERNELS: ExllamaLinearKernel,TritonW4A16LinearKernel
    command: >
      serve btbtyler09/Qwen3.8-27B-GPTQ-4bit
      --served-model-name "Qwen 27B"
      --host 0.0.0.0 --port 8091
      --max-model-len 262144
      --gpu-memory-utilization 0.85
      --kv-cache-dtype float16
      --speculative-config '{"method":"mtp","num_speculative_tokens":4}'
      --tensor-parallel-size 2
      --dtype float16
      --max-num-seqs 4
      --language-model-only --skip-mm-profiling --trust-remote-code
      --enable-auto-tool-choice --tool-call-parser qwen3_coder
      --enable-prefix-caching --enable-chunked-prefill
      --compilation-config '{"cudagraph_mode": "FULL_AND_PIECEWISE", "compile_ranges_endpoints": []}'

Confirm the kernel is active in logs: Using RDNA2W4A16LinearKernel for AutoGPTQLinearMethod.

vLLM Quantization on gfx1030

WIP: Throughput numbers are community-reported. See Verification status.

GPTQ vs AWQ

On older images (before the AWQ→RDNA2 dispatch fix), AWQ fell through to Triton/Exllama and could stall at ~4–5 t/s on a 27B. Pull the latest -extras image and confirm Using RDNA2W4A16LinearKernel in startup logs. Qwen3.8-27B AWQ also needs the fork’s head_size=256 FlashAttention path — without it, FA falls back or never lists RDNA_ATTN.

Kernel dispatch details: rdna_extras fork.

KV-cache dtype

Prefer float16 unless you are A/B testing a quantized KV for a non-agentic workload.

MTP speculative decoding

MTP (--speculative-config '{"method":"mtp","num_speculative_tokens":N}') can boost throughput on GPTQ models with CUDA graphs enabled. Acceptance rates dropped after a v0.27.1 speculator update (~0.25), but base decode speed remains good — worth testing on your model. Example in Configuration.

MTP is not free at high concurrency. A #vllm-rdna TP4 matrix on Qwen3.6-35B-A3B-FP16 (4× V620, --enforce-eager, 16k/1k-style bench) reported MTP-2 +17% output tok/s at c=1, but −53% at c=8. Use MTP for latency-critical single-stream; leave it off for batched throughput.

INT4 on gfx1030 (no native int4 ALUs)

RDNA2 has no hardware int4 matrix units. The -extras W4A16 kernels use vdot2 on fp16 with on-the-fly dequant — int4 weights packed and processed via dp4a-style instructions. Both GPTQ and AWQ dense now hit the same native HIP kernel on current -extras images.

Recent fork work on hybrid GDN models (Qwen3.8-27B-AWQ-INT4, TP4) reported ~93 output tok/s with CUDA graphs (1024/512), ~331 total tok/s at 8 concurrent requests (16k/512), and prefill peaks of 1450–1573 tok/s — with the full HIP GDN prefill + decode chain replacing Triton JIT.

Tips

• Lower --gpu-memory-utilization (e.g. 0.9 → 0.8) if KV-cache allocation OOMs on 16 GB cards.
• For GPTQ or AWQ on -extras: set VLLM_DISABLED_KERNELS=ExllamaLinearKernel,TritonW4A16LinearKernel and watch logs for RDNA2W4A16LinearKernel.
• On older images, AWQ could fall through to Triton (~4–5 t/s on a 27B). Pull latest -extras and confirm the native kernel is active before blaming the quant format.
• Don’t force --attention-backend or --quantization — let vLLM auto-select unless A/B testing.
• Mount Triton and torch-compile caches (see Configuration).

Experimental: EXL3 and Quark (#vllm-rdna, Sep 2026)

These paths are not in the published v0.27.1 -extras image matrix yet. EXL3 HIP kernels are in opengfx1030/vllm-rdna rdna_extras HEAD — track Discord and rebuild from that branch, or wait for a tagged image.

Prefer GPTQ/AWQ on published images until EXL3/Quark land in a tagged Docker build.

Intel AutoRound W4A16 (Flash-Next, #vllm-rdna Sep 2026)

Not in published -extras tags. Community + draft opengfx1030/vllm-rdna#5 are exercising Intel W4A16 AutoRound Flash-Next:

Community take: Intel AutoRound sits between Unsloth Q5_XL and Q6_XL on quality while the weight pack is tens of GB smaller (~75 GB). That still wants four 32 GB cards plus a large CPU PLE / n-gram table (embedded BF16 table is ~95 GiB — offload to RAM).

Do not assume group-16 INT4 PLE sidecars match the published benches. See overview.

vLLM forks (RDNA2 / gfx1030)

WIP: Kernel behavior below is traced to fork source commits; published -extras Docker images may lag HEAD until rebuilt. See Verification status.

Fork landscape

#vllm-rdna (Sep 2026): the official gfx1030 vLLM source is the shared-org repo. File new PRs and issues there. Flash-Next is still a separate production fork until that work lands in the org. Published Docker Hub tags have not moved yet.

Day-to-day serving: keep pulling blivioniag/vllm-rdna -extras tags (v0.27.1). Those images still come from the personal Docker Hub namespace and the historical rdna2_extras clone URL. Source of truth for new kernel work is opengfx1030/vllm-rdna. Do not assume org GHCR/Docker tags exist yet.

Companion (not an engine fork): leapdragon/vllm-rdna2-recipe (mirror opengfx1030/vllm-rdna2-recipe) collects compose/env recipes and open PRs (concurrent MTP, etc.). Wiki entry point: Recipes. #vllm-rdna (Sep 2026): the Flash-Next author treats the recipes repo as a parts pile (less active new work); day-to-day Flash-Next optimization lives in vllm-rdna2-qwen. Concurrent-MTP PRs on the recipes repo are still worth cherry-picking into either stack.

Consolidation status (Sep 3–11 2026)

#vllm-rdna snapshot — expect this to move quickly:

Official-fork feature set called out in-channel (HIP): MXFP4, AWQ INT4, GPTQ INT4, GDN prefill/decode linear attention, and FlashAttention-equivalent kernels. Treat as fork-source / WIP on 0.28 HEAD until Hub -extras retargets.

The rdna_extras fork

The official extras source is opengfx1030/vllm-rdna on the default rdna_extras branch (formerly blivioniag/vllm rdna2_extras). It adds hand-written RDNA2 HIP kernels and the vLLM plumbing to dispatch to them.

Published Hub -extras images are still built from the historical clone (VLLM_REPOSITORY=https://github.com/BlivionIaG/vllm.git, VLLM_REF=rdna2_extras) until vllm-rdna-docker is retargeted. Pull those images for serving; clone the org repo to contribute.

Why it exists: RDNA2 (gfx1030) has no matrix/WMMA cores — those arrived with RDNA3 (gfx11xx). So quantized GEMM, attention, and MoE have to be implemented efficiently on RDNA2’s regular vector ALUs. This fork does exactly that (with RDNA3 WMMA variants where it makes sense).

INT4/W4A16 on gfx1030 does not use native int4 compute. The RDNA2 kernels dequant on the fly using vdot2 fp16 instructions with int4 packed via dp4a. Community testing found no faster INT8-based HIP kernel vs fp16 for these paths — fp16 dequant is the winning approach on V620.

Branches

Historical names on the old personal fork (feat/enable-gfx1030, perf/rdna2_w4a16, rdna2_extras) are superseded. Hub -extras bake targets still pin VLLM_REF=rdna2_extras on that old repo.

What’s in it

The custom device code lives under csrc/rocm/*_rdna2.cu (with _rdna3 variants where WMMA applies), wired into vLLM through Python kernel/layer modules and covered by targeted tests under tests/kernels/.

Attention

• fa_rdna2.cu — a FlashAttention kernel tuned for RDNA2, exposed via vllm/v1/attention/ops/fa_rdna2_backend.py and the vllm/v1/attention/backends/rdna_attn.py backend. Now supports head_size=256 on gfx10x (needed for Qwen3.8-27B-AWQ-INT4; previously capped at 128 and caused Triton GDN compile hangs).
• sparse_mla_rdna2.cu + rocm_rdna2_mla_sparse.py — sparse MLA (multi-head latent attention, DeepSeek-style).
• indexer_paged_mqa_rdna2.cu — paged MQA indexer.

Quantized GEMM (mixed-precision linear)

• q_gemm_rdna2.cu / q_gemm_rdna2_prefill.cu (+ q_gemm_rdna2_common.cuh) — W4A16 quantized GEMM for decode and prefill.
• q_gemm_w8a16_fp8_rdna2.cu, gemm_w8a8_fp8_dense_rdna2.cu — FP8 W8A16 / W8A8 paths.
• qdq_4_rdna2.cuh, qdq_8_rdna2.cuh, qdq_fp8_rdna2.cuh — quant/dequant helpers.
• Python: model_executor/kernels/linear/mixed_precision/rdna2_w4a16.py, .../scaled_mm/rdna2_w8a16_fp8*.py, rdna2_w8a8_fp8.py, and the rdna_hybrid_w4a16.py selector.
• EXL3 (in-tree on rdna_extras, not in published v0.27.1 -extras images yet): exl3_hadamard.cu, exl3_dot2_*.cu, plus vllm/.../quantization/exl3.py. Experimental; see Quantization.

MoE (mixture of experts)

• moe_q_gemm_rdna2.cu, moe_w8a16_rdna2.cu, moe_w8a16_fp8_rdna2.cu — quantized expert GEMMs.
• Python experts fused_moe/experts/rdna2_mxfp4_moe.py, rdna2_w8a16_fp8_moe.py, and compressed_tensors MoE glue (..._fp8_rdna2, ..._w4a4_mxfp4_rdna2, ..._wna16_rdna2).

GDN (gated delta-net / linear attention)

Kernels for gated-delta-net models (e.g. Qwen3.8-27B hybrid linear attention). As of Aug 2026 the full GDN chain is hand-ported to HIP — no Triton JIT on the hot path:

This replaces the Triton FLA GDN path that previously caused 16k/1k prefill hangs and slow decode on hybrid models. Recent tuning commits improved delta_h register pressure, o-kernel block-vector width, and prep-kernel vectorization.

Using it

The easiest path is a prebuilt -extras image — no compilation required:

docker run -it --rm \
  --device /dev/kfd --device /dev/dri --group-add video --group-add render \
  --security-opt seccomp=unconfined --ipc host -p 8000:8000 \
  docker.io/blivioniag/vllm-rdna:v0.27.1-extras \
  vllm serve <model> --dtype float16

See Running (Docker) for the full run recipe and Building images for how the -extras variant is produced (VLLM_VARIANT=extras-fork, currently VLLM_REF=rdna2_extras from the historical personal clone until bake is retargeted at opengfx1030/vllm-rdna).

Kernel dispatch on gfx1030

On -extras images, vLLM picks kernels based on quantization format:

Check startup logs for lines like Using RDNA2W4A16LinearKernel for AutoGPTQLinearMethod. If you see TritonW4A16LinearKernel or ExllamaLinearKernel instead, the RDNA2 quant path isn’t active.

Attention backends

• VLLM_USE_RDNA2_FA=1 — enables the custom fa_rdna2.cu FlashAttention backend.
• --attention-backend RDNA_ATTN — alternative RDNA-tuned attention path (useful for Qwen models with head size 256 where generic AMD Triton FA is slow or broken).
• FLASH_ATTENTION_TRITON_AMD_ENABLE=TRUE — enables AMD Triton FA as a fallback; often slower on gfx1030.

On v0.27.1, hybrid GDN models may still auto-select ROCM_ATTN even with VLLM_USE_RDNA2_FA=1 for the attention layers — that’s expected. The GDN linear-attention layers now use the native HIP kernels above (not Triton FLA). For graph capture issues, see CUDA graphs and vLLM troubleshooting.

CUDA graph capture (TP comm fix)

A common _SimpleCData.__new__ crash during V2 cudagraph capture on multi-GPU TP setups was fixed by marking TP communication wrappers with allow_in_graph. With current -extras images, CUDA graphs are the preferred fast path — you should not need --enforce-eager for this class of failure anymore. On multi-GPU, if AOT compile cache replay misbehaves, try VLLM_USE_AOT_COMPILE=0 VLLM_DISABLE_COMPILE_CACHE=1.

Disabling fallback kernels

export VLLM_DISABLED_KERNELS=ExllamaLinearKernel,TritonW4A16LinearKernel

This is the main lever for forcing GPTQ onto the native RDNA2 W4A16 path. The variable accepts a comma-separated list of kernel class names registered in vLLM’s linear-kernel registry.

To test Exllama instead (still relevant on some quants per #general):

--linear-backend exllama

Omit VLLM_DISABLED_KERNELS when testing Exllama — it competes with the RDNA2 kernel for dispatch.

Building from source (advanced)

git clone -b rdna_extras https://github.com/opengfx1030/vllm-rdna.git
cd vllm-rdna
export PYTORCH_ROCM_ARCH=gfx1030
pip install -r requirements/rocm.txt
pip install --no-build-isolation -e .

The kernels have their own tests, e.g.:

pytest tests/kernels/quantization/test_rdna2_w4a16.py
pytest tests/kernels/attention/test_fa_rdna2_shape_sweep.py

These kernels are actively evolving. Treat the fork as experimental, pin to a known-good image tag for reproducible serving, and file issues on opengfx1030/vllm-rdna if you hit correctness or performance problems.

Building the Images

The blivioniag/rocm-rdna and blivioniag/vllm-rdna images are produced by the vllm-rdna-docker repo. It’s deliberately small: two Dockerfiles, one bake graph, one CI workflow — no Python layer, no custom linter.

Layout

The seven RDNA archs gfx1030;gfx1100;gfx1101;gfx1150;gfx1151;gfx1200;gfx1201 are baked into every image via PYTORCH_ROCM_ARCH.

Build locally

Everything runs through stock docker buildx bake:

# Build everything (bases + vLLM images)
docker buildx bake --file docker-bake.hcl all

# Just the base images, or just the vLLM images
docker buildx bake --file docker-bake.hcl all-bases
docker buildx bake --file docker-bake.hcl all-vllm

# A single target
docker buildx bake --file docker-bake.hcl vllm-0271-rocm720

# Print the plan without building
docker buildx bake --file docker-bake.hcl --print all

Targets are named vllm-<source>-<base>, e.g. vllm-0271-rocm720, vllm-0271-rocm720-extras. Groups: all, all-bases, all-vllm.

Key build arguments (Dockerfile.vllm)

Set per-target in docker-bake.hcl:

The patches/ directory

vLLM occasionally needs small RDNA fixes (e.g. platform detection so consumer Radeon cards are picked up). The repo keeps these as versioned diffs, for example:

• patches/v0.26.0-rocm-platforms.patch
• patches/v0.27.1-rocm-platform-detect.patch
• patches/v0.27.1-amdsmi-wrapper-guard.patch
• patches/v0.27.1-extras-*.patch

To add one for a new vLLM release: reproduce the fix against a fresh clone, git diff the changed file(s) into patches/, verify it applies against a clean checkout of that tag, then point the target’s VLLM_PATCH_FILE at it.

Adding a new base or vLLM source

• New base: copy a target "base-<id>" block, set the ROCm/PyTorch/Triton versions and index URL, and add the id to all-bases + all-vllm.
• New vLLM source: add a target "vllm-<source>-<base>" block per base with VLLM_REPOSITORY, VLLM_REF, VLLM_COMMIT, VLLM_VARIANT, IMAGE_TAG, and add the ids to all-vllm.

See the vllm-rdna-docker README for the full, authoritative instructions.

CI

.github/workflows/build.yml builds and pushes to docker.io/blivioniag/ on tag push (v*), and supports manual dispatch of a single target (optionally pushing). It uses stock docker/setup-buildx, docker/login, docker/metadata, and docker/bake actions with GitHub Actions cache.

llama.cpp on gfx1030 — Overview

llama.cpp is a fast, low-dependency way to run GGUF LLMs on gfx1030. This section covers stock builds and the community RDNA2-optimized fork.

Where to start

Which path?

• Stock llama.cpp — good starting point, single-GPU. Vulkan without ROCm is possible but #llamacpp prefers HIP for --split-mode tensor (RADV has crashed V620 hosts).
• edwinbrowwn/llama.cpp-rdna2 — multi-GPU V620 rigs, tensor parallel, RCCL all-reduce, DFlash2. Most #llamacpp performance work happens here. A matched A/B on Qwen3.8-27B Q6_K + MTP reported +57% vs stock with identical outputs — see Benchmarks.

llama.cpp vs vLLM on V620 (#llamacpp / #vllm-rdna)

Community rule of thumb (Aug 2026):

Neither stack is “finished” for every model. New to the cards? Start with RDNA2 serving, then try vLLM when you need concurrency.

Multi-GPU tensor parallel benefits greatly from PCIe P2P.

llama.cpp on gfx1030

WIP: Tips marked community-validated in this page come from #llamacpp — see Verification status.

llama.cpp is a fast, low-dependency way to run GGUF LLMs on gfx1030. This page is a build-from-source recipe validated by the gfx1030 community (the #llamacpp Discord channel), on Fedora with ROCm 7.2.0 targeting gfx1030. A Vulkan path is included as an alternative that doesn’t require ROCm.

Commands are shown as used on Fedora. Adjust package names for your distro and tweak versions/paths as needed.

Looking for heavy multi-GPU tuning? See the RDNA2-optimized fork (edwinbrowwn/llama.cpp-rdna2) with RDNA2/V620 tensor-parallel and MMQ optimizations.

Dependencies (Fedora)

sudo dnf install @development-tools glm-devel cmake libpng-devel wayland-devel libpciaccess-devel \
  libX11-devel libXpresent libxcb xcb-util libxcb-devel libXrandr-devel xcb-util-keysyms-devel \
  xcb-util-wm-devel python3 git lz4-devel libzstd-devel python3-distutils-extra qt gcc-g++ \
  wayland-protocols-devel ninja-build python3-jsonschema qt5-qtbase-devel qt6-qtbase-devel \
  libcurl-devel xinput libXinerama xcb-util-cursor

ROCm (recommended)

Install ROCm

Example: Fedora with ROCm 7.2.0. Create /etc/yum.repos.d/rocm.repo:

[rocm720]
name=ROCm 7.2.0 repository
baseurl=https://repo.radeon.com/rocm/el10/7.2/main
enabled=1
gpgcheck=1
priority=50
gpgkey=https://repo.radeon.com/rocm/rocm.gpg.key

Then install ROCm and add yourself to the GPU access groups:

sudo dnf clean all
sudo dnf makecache
sudo rpm --import https://repo.radeon.com/rocm/rocm.gpg.key
sudo dnf install rocm rocm-hip-runtime-devel
sudo usermod -a -G render,video $LOGNAME
# log out / back in (or reboot) so the group change takes effect

See Installing ROCm for more detail and for non-Fedora distros.

Build llama.cpp with ROCm (HIP)

git clone https://github.com/ggml-org/llama.cpp.git
export MAX_JOBS=8            # adjust to your CPU cores / available RAM
export ROCM_HOME=/opt/rocm
export PATH=${ROCM_HOME}/bin:${PATH}

HIPCXX="$(hipconfig -l)/clang" HIP_PATH="$(hipconfig -R)" \
  cmake -S llama.cpp -B build \
    -DLLAMA_CURL=ON -DGGML_HIP=ON -DCMAKE_BUILD_TYPE=Release -DGPU_TARGETS=gfx1030 && \
  cmake --build build --config Release -- -j ${MAX_JOBS}

-DGPU_TARGETS=gfx1030 targets Navi 21. For a non-Navi-21 RDNA2 card, build for its real target (e.g. gfx1031/gfx1032) or add it to the list; see HSA_OVERRIDE for RDNA2 Cousins.

Vulkan (alternative)

The Vulkan backend works without ROCm and runs across many GPUs/drivers. For multi-GPU tensor split, prefer the ROCm build (or the RDNA2 fork) — #llamacpp finds HIP/-sm tensor faster, and RCCL TP is a ROCm path.

Mesa RADV has been reported to hard-crash V620 llama.cpp; AMDVLK can stay up but is severely slower. If you must use Vulkan, pin the ICD explicitly:

export VK_ICD_FILENAMES=/etc/vulkan/icd.d/amd_icd64.json   # AMDVLK, not RADV
export GGML_VULKAN_DEVICE=0

Vulkan SDK from your distro

sudo dnf install mesa-vulkan-drivers vulkan-devel glslc spirv-headers-devel

Vulkan SDK from source

Example with version 1.4.350.1, assuming you keep things in ~/Apps/llama.cpp:

export VULKAN_VERSION=1.4.350.1
wget https://sdk.lunarg.com/sdk/download/${VULKAN_VERSION}/linux/vulkansdk-linux-x86_64-${VULKAN_VERSION}.tar.xz
mkdir vulkan
cd vulkan
tar xf ../vulkansdk-linux-x86_64-${VULKAN_VERSION}.tar.xz
export VULKAN_SDK=~/Apps/llama.cpp/vulkan/${VULKAN_VERSION}/x86_64
export PATH=${VULKAN_SDK}/bin:${PATH}
export LD_LIBRARY_PATH=$VULKAN_SDK/lib${LD_LIBRARY_PATH:+:$LD_LIBRARY_PATH}
export VK_LAYER_PATH=${VULKAN_SDK}/share/vulkan/explicit_layer.d
export VK_ADD_LAYER_PATH=${VULKAN_SDK}/share/vulkan/explicit_layer.d
export PKG_CONFIG_PATH=$VULKAN_SDK/share/pkgconfig:$VULKAN_SDK/lib/pkgconfig${PKG_CONFIG_PATH:+:$PKG_CONFIG_PATH}
export CMAKE_PREFIX_PATH=${VULKAN_SDK}:${VULKAN_SDK}/lib/VulkanLoader

Build with Vulkan

git clone https://github.com/ggml-org/llama.cpp.git
export MAX_JOBS=8            # adjust to your CPU cores / available RAM

cmake -S llama.cpp -B build \
  -DLLAMA_CURL=ON -DGGML_VULKAN=ON -DCMAKE_BUILD_TYPE=Release && \
  cmake --build build --config Release -- -j ${MAX_JOBS}

Usage

Multi-GPU llama-server examples with speculative decoding (MTP) and tensor-split across four cards.

ROCm:

llama-server -hf unsloth/Qwen3.5-122B-A10B-MTP-GGUF:UD-Q4_K_XL \
  --no-mmap -dio -fa on -ngl 999 -np 1 \
  --spec-type draft-mtp --spec-draft-n-max 2 \
  --device ROCm0,ROCm1,ROCm2,ROCm3 --split-mode tensor --host 0.0.0.0

Vulkan:

llama-server -hf unsloth/Qwen3.5-122B-A10B-MTP-GGUF:UD-Q4_K_XL \
  --no-mmap -dio -fa on -ngl 999 -np 1 \
  --spec-type draft-mtp --spec-draft-n-max 6 \
  --device Vulkan0,Vulkan1,Vulkan2,Vulkan3 --split-mode tensor --host 0.0.0.0

Flag notes (tune to your setup):

• -ngl 999 — offload all layers to the GPU(s).
• -fa on — flash attention.
• --device ROCm0,ROCm1,… / Vulkan0,Vulkan1,… — select the backend devices to use.
• --split-mode tensor — split each tensor across the selected GPUs (needs good inter-GPU bandwidth; see Multi-GPU PCIe P2P).
• --spec-type draft-mtp --spec-draft-n-max N — Multi-Token-Prediction speculative decoding; the Vulkan example above uses a larger N (6) than the ROCm one (2).
• --spec-draft-device ROCm0 — run the MTP draft model on a single GPU while the main model is tensor-split across multiple cards (community-validated on TP2).
• --no-mmap, -dio — memory/IO tuning; -np 1 sets the number of parallel sequences.

For RDNA2-tuned multi-GPU serving (DFlash2, RCCL all-reduce, higher throughput), see the RDNA2-optimized fork — that’s where most #llamacpp performance work happens.

Adjust the model, quant, device list, and speculative-decoding settings for your hardware.

RDNA2 fork — Overview

WIP: See Verification status.

edwinbrowwn/llama.cpp-rdna2 is a community fork of llama.cpp with RDNA2-specific optimization work, developed on ROCm 7.14 / Ubuntu Server 26 and validated primarily on four Radeon PRO V620 (gfx1030) GPUs. The focus is tensor parallel (TP) and MMQ/MMVQ (quantized matmul) kernels.

Actively evolving and experimental. Treat throughput claims as needing your own matched before/after runs. See Benchmarks. #llamacpp (Sep 2026): RDNA3 (gfx11xx) support is landing on the fork soon — this wiki still focuses on gfx1030 / V620 until that path is documented.

What it does

A native RDNA2 profile activates automatically at runtime — unsupported models, shapes, quants, and topologies fall back to stock llama.cpp. Headline areas:

• MMQ / MMVQ: RDNA2 expert-width MMQ, Q4_0 DOT8 MMVQ, MXFP4/NVFP4 native arithmetic, MTP/DFlash paths.
• FlashAttention: native tiled RDNA2 arithmetic/reductions.
• Tensor parallel: RCCL tuner + P2P all-reduce schedules, embedding-sharded LM head, TP4 P2P fusion.
• Graph fusion: ADD/RMSNorm, Q8_1 reuse, SwiGLU→Q8_1, GDN sibling projection (Qwen3.5/3.6 MoE).

Optimization highlights (author-reported)

• RCCL tensor-parallel all-reduce — GGML_CUDA_ALLREDUCE=nccl; reported +10% tgen / +20% prefill on Qwen 122B.
• DFlash2 speculative decoding with ngram helpers — see Speculative decoding.
• RCCL autotuner, parallel multi-GPU weight uploads, AMD checkpoint backports for Qwen.

See the fork’s README.md and docs/gfx1030-* / docs/rdna2-* for the authoritative list.

Requirements

• Linux, CMake, ROCm with HIP clang and RCCL.
• Validated path: four V620 / gfx1030 with tensor splitting.
• Compatible main GGUF; optional DFlash/MTP draft GGUF.
• Fastest fork quants in #llamacpp / #harnesses are Q4_0 and Q8_0 (native MMVQ paths). Q6 / UD-Q4_K can be slower even when they “fit”. MXFP4 has been reported a bit faster than Q4_0 on some 27B runs — A/B on your cards.
• For GPU-side sampling paths, install hipcub-devel (package name varies by distro) before building — otherwise TOP_K / draft sampling falls back to CPU.

Build

git clone https://github.com/edwinbrowwn/llama.cpp-rdna2.git
cd llama.cpp-rdna2
./scripts/build-rdna2-portable.sh

Override discovery when needed:

ROCM_PATH=/path/to/rocm TARGET_ARCH=gfx1030 BUILD_DIR=build ./scripts/build-rdna2-portable.sh

Maintainer helper for gfx1030 / ROCm 7.14: scripts/build-rdna2-rocm.sh (defaults: ROCm /opt/rocm/core-7.14, target gfx1030).

Next steps

• Benchmarks — author-reported numbers and PR #10 status.
• Serving — launch commands and Docker.
• Speculative decoding — DFlash2, MTP, ngram configs.

RDNA2 fork — Benchmarks

WIP: All numbers below are author- or Discord-reported. Run matched before/after on your hardware. See Verification status.

Point-in-time numbers using edwinbrowwn/llama.cpp-rdna2, ROCm/RCCL, Flash Attention on, F16 KV, batch/ubatch 2048/256, with the env prefix GGML_CUDA_DISABLE_GRAPHS=1 GGML_CUDA_ALLREDUCE=nccl HSA_OVERRIDE_GFX_VERSION=10.3.0 HSA_NO_SCRATCH_RECLAIM=1 (ROCm 7.14 libs on LD_LIBRARY_PATH).

4× V620 (24 CPU threads)

2-GPU tensor split (TP2)

Latest fork update (Aug 2026, #llamacpp)

MTP-4 TP4 on a pelican prompt:

DFlash2 on TP2: ~52 t/s coding at 0–31k context, ~40 t/s at 64k. Checkpoint fix in PR #12 (in review).

Stock llama.cpp vs this fork (community A/B)

#llamacpp matched run on Qwen3.8-27B Q6_K + MTP (same prompts, temp 0, seed 42, 400 tok, 3 runs). Outputs were byte-identical; the fork was faster:

+56.8% on that workload. Stock + llama-swap in the mid-20s t/s on 27B is a common “I have not switched to the fork” report; 2× V620 + fork + Q4_0 has been quoted at 500–600 PP and 40–50 t/s decode.

pp512 = prefill (512-token prompt); tg128 = generation (128 tokens). For A3B MoE, layer split beats tensor split; dense models use tensor split.

Long-context community sweeps (Aug 29–30 2026)

Community #benchmarks posts (4× V620, RDNA2 fork, short env stack, --ubatch-size 1024, MTP). Treat as single-host snapshots — not wiki-reproduced.

Quant comparison (one host, tensor-split + MTP)

Prompt-eval and generation across context tiers; Q8_0 often wins on this fork (native MMVQ path):

Prefill ranking on that host: Q8_0 > Q4 ≈ Q6 > bf16. A separate real ~107k-token prose prompt on Q8_0 reported ~667 PP / ~23 tg with ~45% MTP accept.

ROCm 7.1 vs 10.0 (Ice Lake 4× V620, Q6_K_XL + MTP)

Same fork commit / command shape; community numbers are close across ROCm 7.1 and 10.0 on that host (regular 16k ~963–965 PP / ~46 tg). Prefer ROCm 7.2.0 or 7.14.0 for multi-GPU RCCL unless you are deliberately lab-testing TheRock/10.x.

Reproduce long-context benches with fixed non-repeating prompts (16/32/64/128k) and full command + commit — community tool: GeorgeMA-Strong/llm-context-bench.

Upcoming optimizations (PR #10)

Status: in review — PR #10. Env vars below will be consolidated post-merge.

Reported improvements: 15.9–61.2% lower MMVQ latency, up to 17.7% faster GDN prefill, +4.39% tgen from DeltaNet sibling fusion. MTP works up to 6 with minimal hit on low acceptance.

Peak throughput (4× V620 @ 140 W)

See Power Tuning for the 140 W cap. Long-context PP HIP crash fix tracks ROCm/rocm-systems#4817.

Environment used for PR #10 results

GGML_HIP_SAFE_STATE_IO=1 GGML_HIP_GFX1030_Q8_CACHE=1 GGML_HIP_GFX1030_GDN_SIBLING_FUSION=1 \
GGML_HIP_GFX1030_Q8_1_FUSION=1 GGML_HIP_GFX1030_NATIVE=1 NCCL_P2P_DISABLE=0 NCCL_P2P_LEVEL=PXB \
GGML_TP_SHARDED_OUTPUT=1 GGML_CUDA_ALLREDUCE=nccl HSA_OVERRIDE_GFX_VERSION=10.3.0 \
HSA_NO_SCRATCH_RECLAIM=1 GGML_CUDA_P2P=1 GGML_HIP_GRAPHS=1

RDNA2 fork — Speculative decoding

WIP: Community-tested configs. See Verification status.

DFlash / MTP basics

  --spec-type draft-dflash \
  --spec-draft-model /path/to/dflash.gguf \
  --spec-draft-n-max 6

--spec-draft-n-max is a workload knob; start at the draft model’s block size and tune.

DFlash2 + ngram (community-tested)

DFlash2 shines on real workloads (coding, agents, long context) more than short synthetic benches — community members saw DFlash2 alone score lower than MTP on a 20k bench (~30 vs ~40 t/s) but win on actual agent/coding sessions (~47 t/s with 700+ t/s prefill).

Draft model quant: use Q4_K_M for the DFlash2 draft GGUF — do not use Q8_0 for the drafter (same acceptance, slower).

Recommended draft GGUFs (Aug 2026):

• incoai/Qwen3.8-27B-DFlash2-GGUF
• webhie/Qwen3.8-27B-Q4-AutoRound-Code-GGUF

Coding / ngram-map-k4v:

--spec-type draft-dflash,ngram-map-k4v \
  --spec-draft-n-max 5 \
  --spec-ngram-map-k4v-size-n 12 \
  --spec-ngram-map-k4v-size-m 48

General / ngram-mod:

--spec-type draft-dflash,ngram-mod \
  --spec-draft-n-max 5 \
  --spec-ngram-mod-n-match 24 \
  --spec-ngram-mod-n-min 48 \
  --spec-ngram-mod-n-max 64

Use tg1024+ (not tg32) for realistic decode benchmarks. DFlash2 is more consistent at long context; MTP acceptance tends to fall off. #llamacpp (Aug 26 2026): prefer the DFlash2 Q4_K_M drafter over MTP when context grows — same advice as the draft-quant row above.

Prefill cost: DFlash tends to hurt prompt processing more than MTP. If PP is the bottleneck (short chats, many reconnects), A/B MTP first; if long-context decode matters more, keep DFlash2.

Qwen3.8 Flash-Next (experimental)

Community benches (4× V620, layer vs tensor split) show Flash-Next layer-split can run, while tensor split on experimental qwen4exp / Flash-Next branches is still rough. Fork maintainers have deferred dedicated Flash-Next work until upstream settles — expect experimental support via upstream merges only, not a polished gfx1030 profile.

#llamacpp / forum (Sep 2026): Flash-Next on llama.cpp often lands around ~15–30 t/s on multi-V620 hosts and is widely called out as weaker than the vLLM Flash-Next recipe (~60–100+ t/s decode class after the Sep prefill campaign). Prefer stable Qwen3.8-27B / MoE recipes for production TP on llama.cpp; use vLLM for Flash-Next until upstream/fork gaps close.

#llamacpp (Sep 10–11 2026) community snapshots — still not a polished gfx1030 profile:

Host RAM for Flash-Next n-gram: keeping the n-gram table off storage is on the order of ~50 GB of system RAM (community, 4× V620). Budget that on top of OS + any CPU offload — or keep the table on fast NVMe and accept the slower path (APEX report: NVMe-resident n-gram was fine).

Full DFlash2 serve example (TP4, Qwen3.8-27B)

HSA_NO_SCRATCH_RECLAIM=1 GGML_HIP_RDNA2_AUTO=1 GGML_HIP_SAFE_STATE_IO=1 \
GGML_TP_SHARDED_OUTPUT=1 HSA_OVERRIDE_GFX_VERSION=10.3.0 \
./build/bin/llama-server \
  -m ./models/qwen38-27b-q4s8/autoround/Qwen3.8-27B-Q4_0.gguf \
  -ngl all --split-mode tensor --tensor-split 1,1,1,1 \
  --device ROCm0,ROCm1,ROCm2,ROCm3 --flash-attn on \
  --ctx-size 262144 --batch-size 8192 --ubatch-size 4096 \
  --host 0.0.0.0 --port 8080 --metrics \
  --reasoning-effort xhigh --reasoning-preserve \
  --spec-type draft-dflash,ngram-map-k4v \
  --spec-ngram-map-k4v-size-n 12 --spec-ngram-map-k4v-size-m 48 \
  --spec-draft-n-max 5 \
  -md ./models/qwen38-27b-q4s8/dflash2/Qwen3.8-27B-DFlash2-Q4_K_M.gguf \
  --device-draft ROCm0 --parallel 1 --spec-draft-ubatch-size 4096 \
  --cache-ram 65535

MTP on tensor-split setups

Pin the draft to one GPU while the main model stays tensor-split:

--spec-type draft-mtp --spec-draft-n-max 4 --spec-draft-ngl 999 --spec-draft-device ROCm0

Community tuning tip (Qwen3.8-27B Q8, TP4, real prompts): --spec-draft-n-max 3 often beat n=4 (higher acceptance / mean accepted length). A/B on your workload — do not trust 100% draft acceptance in synthetic benches (usually a bad prompt).

MXFP4 quants (e.g. quark75/Qwen3.8-27B-MXFP4-GGUF) pair well with MTP on TP2 — see Serving.

HIP sidecar speculative decoding is now an opt-in fork path (#llamacpp, Sep 2026):

export SPEC_SIDECAR=1   # moves MTP/DFlash logic outside the main llama.cpp process

Pair with --spec-draft-p-min 0.0 (fork maintainer tip). Sidecar work is still evolving — pull latest edwinbrowwn/llama.cpp-rdna2 and prefer MTP over DFlash on the sidecar path while DFlash remains slow. Track the fork README for flag churn; do not treat SPEC_SIDECAR as required for the built-in MTP flags above.

Sidecar / DFlash gotchas (Sep 2026)

#llamacpp reports to know about:

RDNA2 fork — Serving

Recommended env stack

HSA_OVERRIDE_GFX_VERSION=10.3.0 is required for the tested V620/gfx1030 native profile.

Prefer this short stack (#llamacpp, Aug 2026). Older long lists of GGML_HIP_GFX1030_* knobs are mostly redundant with the override and can clash with the RCCL autotune path:

HSA_OVERRIDE_GFX_VERSION=10.3.0 \
HSA_NO_SCRATCH_RECLAIM=1 \
GGML_HIP_RDNA2_AUTO=1 \
GGML_HIP_SAFE_STATE_IO=1 \
GGML_TP_SHARDED_OUTPUT=1   # TP2+ only

• GGML_HIP_SAFE_STATE_IO=1 — recommended default; mitigates a known ROCm FA crash class.
• GGML_CUDA_ALLREDUCE=nccl — community reports +10% tgen / +20% prefill on Qwen 122B when RCCL is healthy.
• Optional TP4 mode: GGML_HIP_GFX1030_P2P_ALLREDUCE=auto-expanded (topology-gated).
• If you see internal AllReduce init failed (n_devices != 2) or wild PP variance, strip custom GGML_HIP_GFX1030_* / P2P knobs back to the short stack and re-test on ROCm 7.2.0 or 7.14.0 (not mid-7.2.x such as 7.2.4). See Installing ROCm.

Launch

TP2+ (four-GPU example):

HSA_OVERRIDE_GFX_VERSION=10.3.0 \
HSA_NO_SCRATCH_RECLAIM=1 \
GGML_HIP_RDNA2_AUTO=1 \
GGML_HIP_SAFE_STATE_IO=1 \
GGML_TP_SHARDED_OUTPUT=1 \
./build/bin/llama-server \
  -m /path/to/main.gguf \
  -ngl all \
  --split-mode tensor \
  --tensor-split 1,1,1,1 \
  --flash-attn on \
  --host 0.0.0.0 --port 8080

TP1 — omit tensor split and GGML_TP_SHARDED_OUTPUT:

HSA_OVERRIDE_GFX_VERSION=10.3.0 \
HSA_NO_SCRATCH_RECLAIM=1 \
GGML_HIP_SAFE_STATE_IO=1 \
./build/bin/llama-server -m /path/to/main.gguf -ngl all --flash-attn on --host 0.0.0.0 --port 8080

Batch / ubatch tips

For tensor-split prefill, #llamacpp often does better with larger ubatch than the old 2048/256 bench defaults — roughly ~1024 ubatch per GPU (e.g. TP4 → --ubatch-size 4096) while keeping --batch-size ≥ ubatch. Small prompts may regress slightly; long prompts usually win.

Recent long-context community recipes commonly use --batch-size 16384 --ubatch-size 1024 on TP4 (see Benchmarks).

Host tips that affect llama-server

• CPU governor: if any hot path stays on the host (Flash-Next n-gram tables, --override-tensor …=CPU, MoE/KV offload), powersave can lag bursty PP. Community: switching Intel intel_pstate to performance improved Flash-Next PP ~33% while VRAM-resident 27B was unchanged. Check /sys/devices/system/cpu/cpu0/cpufreq/scaling_governor. See General troubleshooting.
• DAX / Optane model store: fast loads are nice for swap-testing; never mmap GGUFs from dax=always mounts into ROCm — use --no-mmap / --load-mode none. See troubleshooting.

Full validated example (4× V620, Qwen3.5-122B-A10B-MTP)

GGML_CUDA_DISABLE_GRAPHS=1 GGML_CUDA_ALLREDUCE=nccl HSA_OVERRIDE_GFX_VERSION=10.3.0 \
HSA_NO_SCRATCH_RECLAIM=1 LD_LIBRARY_PATH=/opt/rocm/core-7.14/lib${LD_LIBRARY_PATH:+:$LD_LIBRARY_PATH} \
./build/bin/llama-server \
  -m /path/to/Qwen3.5-122B-A10B-MTP-UD-Q4_K_M.gguf \
  -ngl all --split-mode tensor --tensor-split 1,1,1,1 --main-gpu 0 \
  --ctx-size 262144 --cache-type-k f16 --cache-type-v f16 --kv-unified \
  --batch-size 2048 --ubatch-size 256 --parallel 4 \
  --host 0.0.0.0 --port 8080 --temp 0.6 --metrics --jinja --numa distribute \
  --reasoning on --reasoning-format deepseek --chat-template-kwargs '{"preserve_thinking": true}' \
  --spec-type draft-mtp --spec-draft-n-max 3 --spec-draft-type-k f16 --spec-draft-type-v f16 \
  --spec-draft-p-min 0.0 --spec-draft-p-split 0.10 --no-spec-draft-backend-sampling --cache-ram 0

Docker example

TP2 + MTP4 from #llamacpp (ROCm 7.14, MXFP4 Qwen3.8-27B):

environment:
  HSA_NO_SCRATCH_RECLAIM: "1"
  GGML_HIP_RDNA2_AUTO: "1"
  GGML_HIP_SAFE_STATE_IO: "1"
  GGML_TP_SHARDED_OUTPUT: "1"
  HSA_OVERRIDE_GFX_VERSION: "10.3.0"
command: >
  llama-server
  -hf quark75/Qwen3.8-27B-MXFP4-GGUF:MXFP4
  -ngl 999 -fa on -sm tensor -ts 1,1 -fit off
  --flash-attn on --ctx-size 32768
  --cache-type-k f16 --cache-type-v f16
  --batch-size 4096 --ubatch-size 4096 --parallel 1 --cont-batching
  --host 0.0.0.0 --port 8091 --temp 0.7 --jinja --metrics -kvu
  --reasoning-preserve
  --chat-template-kwargs '{"preserve_thinking": true}'
  --spec-type draft-mtp --spec-draft-n-max 4 --spec-draft-ngl 999
  --ctx-checkpoints 0

Build with ./scripts/build-rdna2-portable.sh. RCCL needs working PCIe P2P. If P2P is enabled but slower, set NCCL_P2P_DISABLE=1. If all-reduce fails, try GGML_HIP_GFX1030_P2P_ALLREDUCE=off or GGML_CUDA_ALLREDUCE=none.

The RDNA2 fork’s recent HIP/RCCL work is validated primarily against ROCm 7.14 (#llamacpp, Sep 2026). Mid-7.2.x (e.g. 7.2.4) is not a confident target — upgrade or pin per Installing ROCm.

Single-GPU community ballpark (Qwen3.8-27B)

#llamacpp (Sep 2026), 1× V620, RDNA2 fork, short env stack — treat as single-host snapshots:

On 4× V620, long-context MTP benches (community suite) saw decode hold ~60–65 t/s at 32k on coding prompts but drop harder by 64k — ngram / sidecar MTP work aims to lift acceptance; A/B with llm-context-bench.

Idle power with embedded servers

Running an embedding server (nomic, etc.) alongside chat on the same cards can add ~40 W per GPU even when the embedder looks idle (#llamacpp). Park embedders on a spare card or stop them when not needed if you care about power.

Concurrent slots / multi-agent (V620)

#llamacpp (Sep 2026): on V620 (no matrix cores), --parallel > 1 can crush per-stream decode — community report of ~35–50 t/s → ~5–9 t/s when two generations overlap, and prefill of a second request can stall the first. Prefer one stream per process and scale out with separate instances (different GPUs / aliases) if you need multi-agent throughput. KV-aware load balancers (e.g. paddler) and shared KV layers (LMCache on the vLLM side) are discussed in #llamacpp / #lmcache but are not wiki-validated recipes yet.

Faster GGUF loads (upstream WIP)

Community (#vllm-rdna / #llamacpp, Sep 2026): upstream llama.cpp PR #22466 (async pinned upload for -sm tensor loads) dropped a 122B-Q4 load from >2 min to <35 s on one host. Watch that PR / fork merges if cold-start time matters.

Notable limits

• Validated primarily on 4× V620 gfx1030, ROCm 7.14; other systems use conservative fallbacks. TP4 does work on known-good server boards (e.g. Gigabyte MC62-G40) and on some dual-socket Broadwell-EP hosts. Most successful TP4 + P2P reports are on AMD CPUs.
• Multi-socket hosts can hurt TP — single-socket EPYC preferred. Dual-Xeon UPI traffic has halved prefill in #general even when theoretical cross-socket bandwidth looked fine.
• KV checkpoints crash on tensor split — use --ctx-checkpoints 0.
• GGML_TP_SHARDED_OUTPUT and GGML_TP_VOCAB_SHARDED_OUTPUT are incompatible modes.
• Tensor-split prefill spikes current — layer split can look fine while TP shuts the PSU off at prefill start (transient on the +12 V rail, classic with old miner PSUs). Try 160 W (~2–4% slower) or 140 W (~8–10% slower vs unlocked) before blaming the fork. See Power Tuning.
• TP3 (three cards) has caused driver crashes; stick to 2 or 4 when possible. On boards with three full CPU x16 slots, community still prefers trying tensor split first if the links are PCIe 4.0 x16 — see Host topology.
• Most GGML_HIP_GFX1030_* flags are redundant with HSA_OVERRIDE_GFX_VERSION=10.3.0 unless A/B testing — prefer the short env stack.
• FA / q8 KV: quantized V cache needs FA on; FA occupancy asserts on some head-256 models — see troubleshooting.
• GPU sampling (--spec-draft-backend-sampling and related): install hipcub-devel (or distro equivalent) at build time. Without it, expect device 'Meta()' does not have support for op TOP_K and fall back to slower CPU sampling.
• Long-context quant quality: community (#llamacpp): below Q6 often looks fine on short benches then loops / stalls as context grows; Q8 is noticeably more stable. Pair with sidecar gotchas.

For stock builds see Building & running. Speculative decoding configs: RDNA2 speculative decoding.

Environment Variables & Quick Reference

WIP: vLLM and llama.cpp tables include Discord-sourced defaults — see Verification status.

A cheat-sheet of the settings that matter most when running ML workloads on gfx1030 / RDNA2.

Key environment variables

vLLM on gfx1030 (#vllm-rdna)

llama.cpp RDNA2 fork (#llamacpp)

Prefer the short stack below. Long lists of GGML_HIP_GFX1030_* knobs are usually redundant with HSA_OVERRIDE_GFX_VERSION=10.3.0 and can clash with RCCL autotune — see Serving.

vLLM compilation / CUDA graphs

10.3.0 is the magic value for gfx1030 because the target decodes as gfx + 10 (major) 3 (minor) 0 (stepping) → gfx1030.

V620 / gfx1030 PCI identity

Used by the tuning scripts to match the right board:

lspci -nn | grep '1002:73a1'   # find V620 reference boards

Handy commands

rocminfo                              # full agent/GPU info; look for "Name: gfx1030"
rocminfo | grep -m1 -o 'gfx[0-9]*'    # just the target name
rocm-smi                              # live clocks, temps, VRAM, power, utilization
rocm-smi --showmeminfo vram           # VRAM usage
clinfo | grep -i board                # OpenCL board name
/opt/rocm/bin/rocminfo | grep -i wavefront   # confirm wave32 on RDNA2

VRAM rules of thumb (LLMs)

Pro cards show ~30 GB until you disable ECC.

FP16 vs BF16

RDNA2 has no fast BF16. Always prefer FP16 for hot paths:

• PyTorch: pass dtype=torch.float16.
• vLLM: --dtype float16 (see Running vLLM).
• The rdna_extras fork adds quantized (W4A16 / FP8) RDNA2 kernels to cut VRAM and sidestep BF16 entirely.

Verification status (WIP)

This wiki is work in progress. Pages expanded from Discord (#vllm-rdna, #llamacpp, #general) and fork release notes mix solid facts, fork-source claims, and community benchmarks that have not been independently reproduced in this repo.

Legend

Global gaps (read first)

vLLM (vllm/)

running.md

recipes.md

configuration.md

quantization.md

fork.md

llama.cpp (llama-cpp/)

building.md

rdna2-benchmarks.md

rdna2-speculative.md

rdna2-serving.md

Reference & troubleshooting

Checklist before production

1. docker pull latest -extras image; confirm fork commit in build metadata.
2. Grep logs for Using RDNA2W4A16LinearKernel.
3. Try CUDA graphs before --enforce-eager.
4. One matched A/B on your hardware.
5. Update this page when you confirm or refute a claim.

See Contributing and Wiki structure.

Troubleshooting

Common errors on gfx1030 / RDNA2 and how to fix them. Pick the section that matches your stack:

WIP: Fixes involving latest -extras images assume a current image pull — see Verification status.

Still stuck?

• ROCm troubleshooting docs
• Community resources
• Contributing — add your fix so others don’t rediscover it

General troubleshooting

torch.cuda.is_available() returns False

1. Confirm ROCm sees the card: rocminfo | grep -m1 -o 'gfx[0-9]*'.
2. Confirm you installed the ROCm build of PyTorch (torch.__version__ should end in +rocmX.Y).
3. Make sure your user is in the render and video groups: groups | grep -E 'render|video'. If not: sudo usermod -aG render,video "$LOGNAME" and re-login.

"no kernel image is available for execution on the device"

Set the override:

export HSA_OVERRIDE_GFX_VERSION=10.3.0

See HSA_OVERRIDE. If building from source, use PYTORCH_ROCM_ARCH / AMDGPU_TARGETS for your real target.

hipErrorNoBinaryForGpu / Memory access fault

• Usually wrong/missing arch — apply the override or rebuild.
• Can also mean OOM; check rocm-smi and reduce batch size / context.

hipBLASLt errors

export TORCH_BLAS_PREFER_HIPBLASLT=0

BF16 is extremely slow

Force FP16 everywhere (--dtype float16 in vLLM, dtype=torch.float16 in PyTorch). See Environment variables.

The iGPU is being selected instead of my discrete card

export HIP_VISIBLE_DEVICES=0     # index of your discrete GPU in `rocminfo`

First diffusion / MIOpen run is very slow

MIOpen compiles kernels on first use in ~/.cache/miopen. Later runs are fast.

Secure Boot blocks the amdgpu-dkms module

Either sign the module or disable Secure Boot.

CPU governor hurts host-resident models

Symptom: Flash-Next (or any model with large CPU-side tables / offload) has weak prefill; VRAM-only 27B is fine.

On Intel intel_pstate, default powersave still boosts but ramps lazily. Bursty host work (n-gram hash + gather from a multi-GB host table) finishes before the governor reacts. Community: performance improved Flash-Next PP ~33% on a ~6k prompt; GPU-bound Qwen3.8-27B unchanged; idle clocks still drop.

cat /sys/devices/system/cpu/cpu0/cpufreq/scaling_governor
echo performance | sudo tee /sys/devices/system/cpu/cpu*/cpufreq/scaling_governor

Persist with a systemd oneshot if needed (resets on reboot). Skip this if every tensor stays in VRAM.

Unsupported AMD GPU in the box “steals” ROCm

Symptom: env overrides look correct, but ROCm / vLLM acts as if only an old unsupported AMDGPU exists (e.g. Polaris / WX4100) and ignores the V620s — or rocminfo crashes with topology errors about an unsupported KFD node.

#general / #llamacpp: some ROCm code paths punt the entire stack when they see an unsupported AMD device, instead of skipping that one card. ROCR_VISIBLE_DEVICES alone is unreliable for this on mixed hosts.

Workarounds (pick one):

1. Unbind the display/legacy card from amdgpu after boot (keep it for POST / recovery, but out of the compute topology).
2. Use a non-amdgpu display GPU (old NVIDIA / Intel / Terascale HD that never binds amdgpu) or run the host headless.
3. Hide devices carefully with HIP_VISIBLE_DEVICES / ROCR_VISIBLE_DEVICES and re-check rocminfo still enumerates every V620.

Docker device passthrough of only the V620 does not always isolate an unsupported host GPU from HSA topology init — VM passthrough or unbind is the reliable fix when mixed AMD cards fight ROCm.

Overheating / thermal paste on V620

#general: V620s are passive server cards — airflow matters more than most desktop GPUs.

PCB photos for cooler compatibility: community link Linus Tech Tips V620 thread. W6800 waterblocks are not a guaranteed fit.

vLLM troubleshooting

CUDA graph capture crashes

Symptom: crash at GDN _output_projection all-reduce / _SimpleCData.__new__, or OOM during graph capture.

Root cause (fixed in current extras): TP comm wrappers were not allow_in_graph. Fork fix: fix(distributed): allow TP comm ops in torch.compile graph capture.

First, pull the latest image:

docker pull blivioniag/vllm-rdna:v0.27.1-extras-rocm7.14.0

Then try CUDA graphs (fast path):

--compilation-config '{"cudagraph_mode": "FULL_AND_PIECEWISE", "compile_ranges_endpoints": []}'

Mount cache volumes — see Configuration.

Fallback: --enforce-eager

On multi-GPU AOT cache issues: VLLM_USE_AOT_COMPILE=0 VLLM_DISABLE_COMPILE_CACHE=1

Low throughput (~4–5 t/s on 27B)? Check for older image where AWQ still used Triton — see Quantization.

First boot is extremely slow

Triton and torch.compile JIT on first run. Typical cache: ~/.triton/cache (~3 GB), ~/.cache/vllm/torch_compile_cache (~700 MB).

GPTQ/AWQ not using RDNA2 kernels

Check logs for Using RDNA2W4A16LinearKernel. If you see Triton/Exllama instead:

export VLLM_DISABLED_KERNELS=ExllamaLinearKernel,TritonW4A16LinearKernel

Confirm -extras image from current extras. See Fork kernel dispatch.

vLLM picks the wrong platform / doesn’t see my Radeon

Use published blivioniag/vllm-rdna images with patches/*rocm-platform* fixes rather than stock upstream builds.

Failed to infer device type / AMDSMI_STATUS_NOT_INIT

Symptom: vLLM logs ROCm platform is not available because no GPU is found and AMDSMI_STATUS_NOT_INIT - Device not initialized (often right after CUDA/NVML is also missing — that part is expected on AMD).

Check the Docker device/group block exactly — #vllm-rdna hits this when render is missing:

devices: [/dev/kfd, /dev/dri]
group_add: [video, render]    # both — video alone is not enough on many hosts
ipc: host
security_opt: [label=disable]  # or seccomp=unconfined on docker run

Host user still needs render/video as in General troubleshooting. Debug with VLLM_LOGGING_LEVEL=DEBUG. If the same compose worked on an older tag, docker pull a known-good image — a bad rebuild can also fail AMDSMI init.

ROCM_ATTN hangs for hours (Triton compile)

AMD Triton flash-attention compile on gfx1030 can sit there for hours (RCCL vs Triton). On -extras, use --attention-backend RDNA_ATTN and/or VLLM_USE_RDNA2_FA=1 instead of ROCM_ATTN. See Configuration.

FA_RDNA2 / RDNA_ATTN may not appear in the backend list on older -extras images or some GPTQ models (logs only show Triton / ROCM / TurboQuant). Pull the latest -extras tag and confirm Using RDNA2W4A16LinearKernel / native FA in startup logs. Qwen3.8-27B AWQ needs head size 256 on the fork — see Quantization.

shm_broadcast / one GPU + one CPU pegged

Symptom (#vllm-rdna Sep 2026, often Flash-Next / multi-GPU): serve looks wedged; logs repeat something like:

No available shared memory broadcast block found in 60 seconds.
This typically happens when some processes are hanging or doing some time-consuming work
(e.g. compilation, weight/kv cache quantization).

Community diagnosis: classic Triton JIT compile fighting RCCL — not necessarily a dead process. Guidance:

1. Wait — first boots can sit like this a long time; keep the Triton / compile cache volumes mounted (Configuration).
2. Stay on ROCm 7.2.0 or 7.14.x — avoid mid-7.2.x (same pin as multi-GPU RCCL).
3. Prefer HIP / RDNA_ATTN paths over AMD Triton FA where the fork offers them — less Triton means fewer of these stalls.
4. Next cold start with a warm cache should be much shorter; if it never recovers after hours, A/B VLLM_USE_V2_MODEL_RUNNER=0 and the long-prompt notes.

Multi-GPU RCCL hangs or cards drop offline

If TP works on one image and dies after a host ROCm bump, check the ROCm version before the model. 7.2.1 through ~7.13 are reported to have a multi-card RCCL bug. Stay on 7.2.0 or 7.14.0 — see Installing ROCm.

Flash-Next long-prompt stalls

Symptom (Flash-Next fork / recipe containers, #vllm-rdna Sep 2026): short prompts decode fine, but large prompts (tens of k tokens — agentic coding, session resume) take many minutes, timeout, or appear wedged. Temps and power caps look healthy.

Community fix that unblocked one 4× V620 host:

export VLLM_USE_V2_MODEL_RUNNER=0

Reporter then saw stable ~68 tok/s with dense INT8 + custom all-reduce, including large prompts. The Flash-Next fork author added this to their docs / troubleshooting. Official-extras authors note separate Dense-on-V2 fixes in progress on the org rebase — A/B both values on your image.

Also rule out thermal / power first (Power tuning), and measure expected prefill time (~1k tok/s class ⇒ ~40 s for 40k tokens, not minutes). Prefill campaign numbers: vLLM overview.

Flash-Next 128k prefill cliff

Symptom (#vllm-rdna Sep 2026, Intel AutoRound / draft opengfx1030/vllm-rdna#5): prefill holds ~950 tok/s through 64k, then falls to ~375 tok/s at 128k while decode stays flat (~48–55 tok/s). Looks like a scheduler / chunking misconfig more than a kernel cliff.

Community fix on that host: raise the batch cap:

--max-num-batched-tokens 4096

The suspected combination was enable_chunked_prefill=True with 2048 scheduled / batched tokens. After 4096, 128k prose/code PP returned to the ~950 tok/s class (same 4× V620, FP16 KV, CPU PLE offload). See overview.

If 128k is still slow after 4096, A/B V2 runner and confirm you are not on quantized PLE (known-good is the embedded BF16 n-gram table).

Prefill blocks decode / MTP stalls under concurrency

Symptom: with speculative decode (MTP) and multiple in-flight requests, generation stalls while prefill runs; or graph + MTP3 reaches “Application startup complete” then hangs on PLE lookup / sample_tokens timeout.

Community notes (#vllm-rdna):

• Prefer GPTQ + RDNA2 W4A16 (or AWQ HIP) paths over GGUF-in-vLLM for these cards.
• Concurrent MTP / prefill-vs-decode fixes land in community recipes first — see open PRs on leapdragon/vllm-rdna2-recipe.
• MTP=0 vs MTP=3 are different bug surfaces; a commit that “works” at MTP=3 can still emit spurious tokens at MTP=0. A/B and pin a known-good recipe commit.
• Slow or broken P2P + custom all-reduce can look like MTP latency bugs — A/B the disable vs PIX custom AR stacks.

ROCR idle CPU spin (TheRock 7.14)

Symptom: after starting a multi-GPU vLLM serve on TheRock / ROCm 7.14 (ROCR 1.21), the host holds several CPU cores at high utilization even when the GPUs are idle (~one core per HIP process, plus more once RCCL initializes).

Root cause: AsyncEventsLoop / signal-wait paths busy-spin without backoff (ROCm/TheRock#7051, ROCm/ROCm#6522). Stock env knobs (HSA_ENABLE_INTERRUPT, etc.) do not fix multi-GPU cases.

Fix: rebuild only libhsa-runtime64.so with the poll-backoff patch and LD_PRELOAD it. Step-by-step for host builds (and note that recipe containers already bake the patch): docs/rdna2/ROCR-CPU-FIX.md in leapdragon/vllm-rdna2-qwen. Pull latest recipe / container before re-debugging idle CPU.

llama.cpp troubleshooting

KV checkpoint crash on tensor split

Symptom: fatal error in ggml-backend-meta.cpp during warmup with tensor split.

Fix: --ctx-checkpoints 0. Known on stock llama.cpp and the RDNA2 fork. See RDNA2 serving limits.

FlashAttention abort: max_blocks_per_sm > 0

Symptom: server aborts with something like:

fattn-common.cuh:…: GGML_ASSERT(max_blocks_per_sm > 0) failed
  launch_fattn<256, …>

#llamacpp / community reports this on gfx1030 when the HIP occupancy query returns 0 for the head-size 256 FA tile kernel (common on Qwen3.8 dense and some MoE paths).

Constraints that make this painful:

Mitigations to try (in order):

1. Prefer the RDNA2 fork build script (./scripts/build-rdna2-portable.sh) so FA paths match the fork’s gfx1030 profile.
2. Keep the simplified env stack — especially GGML_HIP_SAFE_STATE_IO=1 (known ROCm FA crash workaround). Do not pile on every GGML_HIP_GFX1030_* flag; that set can clash with the RCCL autotune path. See Serving.
3. If FA still aborts on head-256 models under TP: fall back to f16 KV + FA on only after a fork update / local occupancy patch, or temporarily use layer split for that model until FA occupancy is fixed upstream/fork-side.

DAX-backed mmap oopses amdgpu SVM

Symptom: loading a GGUF from an Optane / pmem dax=always mount with default mmap → kernel fault in svm_range_dma_map_dev, process becomes unkillable, VRAM leaked on all cards, reboot required.

Fix: --no-mmap and/or --load-mode none whenever the model file lives on a DAX mount. Storage only affects load time (community: Optane DAX ~4 GB/s vs NVMe hundreds of MB/s for a ~30 GB GGUF); once weights are in VRAM, inference is unchanged. Great for swap-testing models; useless for a single long-lived production load.

RCCL all-reduce fails (HIP “operation cannot be performed”)

Symptom: ggml_backend_cuda_comm_allreduce_nccl crash, NCCL WARN HIP failure.

Try in order:

1. Confirm PCIe P2P is working.
2. Set NCCL_P2P_LEVEL=PHB or NCCL_P2P_DISABLE=1.
3. On the RDNA2 fork: GGML_HIP_GFX1030_P2P_ALLREDUCE=off or GGML_CUDA_ALLREDUCE=none.
4. Check ACS — CPU root-port ACS can block GPU-direct P2P.

P2P enabled but slower inference

Bandwidth tests can pass while inference regresses on gen3 x4 or ACS-blocked topologies. A/B with NCCL_P2P_DISABLE=1. See When P2P hurts.

PSU dies the moment tensor-split prefill starts

Symptom: layer split is stable; --split-mode tensor kills power (no HIP error in logs). #llamacpp traced this to PSU transients, not the kernels — especially old miner PSUs and Lenovo P620 proprietary GPU cables (that chassis PSU often only feeds two cards).

1. Cap at 160 W or 140 W (Power Tuning).
2. A/B card pairs — one slot pair can trip protection while others do not.
3. Prefer 2 or 4 GPUs; TP3 has caused driver crashes after a “successful” run.
4. Split GPU power off the motherboard PSU if the board only has two GPU power ports.

Dual-socket / NUMA is slow

Pin llama.cpp to one socket (numactl --cpunodebind=0 --membind=0) and keep all TP GPUs on that socket. Crossing NUMA for tensor split is a known prefill killer — see Host topology.

Vulkan RADV hard-crashes; AMDVLK is slow

#llamacpp: Mesa RADV can hard-reboot or crash the host on V620 llama.cpp; switching the ICD to AMDVLK (VK_ICD_FILENAMES=/etc/vulkan/icd.d/amd_icd64.json) can get inference running but is much slower. For multi-GPU --split-mode tensor, the community path is ROCm / HIP, not Vulkan — tensor parallel needs RCCL. See Building llama.cpp.

If a new V620 hard-reboots a box that was stable with a 3080, read Slot power and PSU transients before chasing Vulkan ICDs.

DFlash2 / sidecar crashes

Symptoms (#llamacpp Sep 2026):

• HSA_STATUS_ERROR_MEMORY_APERTURE_VIOLATION in a DFlash gemv_mmvq2_* kernel, then [dflash-sidecar] an illegal memory access was encountered and ROCm abort.
• Sidecar probe: qwen35-mtp target mismatch: target GGUF model identity differs → falls back off sidecar.

See Sidecar / DFlash gotchas: match publisher families for target/draft GGUFs, keep --spec-draft-p-min 0, pull latest fork, A/B MTP.

MTP slot wedge: LCP / prompt-cache position desync

Symptom (#llamacpp, Sep 2026): llama-server stays up and the HTTP API returns 200, but no tokens are generated (clients such as Hermes show 0 tok/s). Common on long agentic sessions (>~40k tokens; community failures clustered around 32–89k). Restarting the process is not required — a later new-context request often recovers the wedged slot.

Typical stderr fingerprint (M-RoPE / draft context):

W find_slot: non-consecutive token position … after … for sequence …
E init: the tokens of sequence … have inconsistent sequence positions:
  - the last position stored in the memory module … is X = …
  - the tokens … have a starting position of Y = …
  for M-RoPE, it is required that the position satisfies: X < Y
E spec process: llama_decode(ctx_dft) head=0 failed rc=-1
E srv decode: failed to process speculative batch

X = Y and X > Y both show up. The main decode can accept the batch; the MTP draft context (ctx_dft) is what rejects it.

This is not the same as the KV checkpoint crash. Reporters already had --ctx-checkpoints 0. The trigger is the LCP / prompt-cache slot-reuse path (selected slot by LCP similarity) plus --spec-type draft-mtp: a reused prefix is re-spanned at a stale position while the draft context holds another. It also is not specific to quantized KV — the same wedge was seen after switching KV from q8_0 back to f16.

Community setup that reproduced it daily: RDNA2 fork, 3× V620 tensor-split, Qwen3.8-27B Quark-AWQ-MXFP4, --parallel 2, Flash Attention on, --spec-type draft-mtp --spec-draft-n-max 4, --ctx-checkpoints 0.

Mitigations to try (in order):

1. Turn MTP off for that slot / session (or gate it when logs show non-consecutive token position).
2. Avoid prefix-cache reuse: new session / --parallel 1, or force a fresh slot after a 200-with-no-tokens event.
3. Prefer Q6+ / Q8 long-session quants over Quark-AWQ-MXFP4 — see Sidecar / DFlash gotchas.
4. Pull latest edwinbrowwn/llama.cpp-rdna2 and report the three-step fingerprint on #llamacpp if it still wedges.

If the hang is a failed tool call or a premature stop token (raw payload looks complete), that is a different class — inspect the harness payload before blaming MTP.

Concurrent decode collapses

Symptom: single-stream is ~35–50 t/s; with two overlapping generations (or a second prefill) decode falls to single-digit t/s.

Expected on V620 for many llama.cpp configs — see Concurrent slots. Prefer separate instances / GPUs over high --parallel for multi-agent.

Wiki structure

How this mdBook is organized — for editors adding or moving pages.

Directory layout

src/
  intro.md                 # Landing page (WIP banner)
  SUMMARY.md               # Sidebar navigation — edit this when adding pages

  setup/                   # Hardware and ROCm install
    hardware.md
    getting-started.md
    installing-rocm.md
    hsa-override.md

  tuning/                  # V620-specific host tuning
    power.md
    ecc.md                 # Optional Pro-card ECC disable (~2 GB)
    p2p.md

  vllm/                    # vLLM serving on RDNA
    overview.md            # Section hub — start here
    running.md             # Images, docker run, quick start
    recipes.md             # Which stack / model / card count
    configuration.md       # Env vars, CUDA graphs, compose
    quantization.md        # GPTQ/AWQ, KV, MTP, INT4
    fork.md                # rdna_extras kernels
    images.md              # Building Docker images

  llama-cpp/               # llama.cpp and RDNA2 fork
    overview.md            # Section hub
    building.md            # Stock llama.cpp build
    rdna2-overview.md      # Fork intro, build, requirements
    rdna2-benchmarks.md    # Author-reported numbers
    rdna2-speculative.md   # DFlash2, MTP, ngram
    rdna2-serving.md         # Launch, Docker, limits

  reference/               # Lookup tables
    env-vars.md            # Environment variable cheat-sheet
    verification.md        # What's verified vs community-reported

  troubleshooting/         # Problem → fix, by stack
    index.md               # Hub
    general.md
    vllm.md
    llama-cpp.md

  meta/                    # Wiki maintenance
    resources.md           # External links
    contributing.md
    structure.md           # This page

Conventions

Adding a new page

1. Pick the folder (or create one if it’s a new major topic).
2. Create src/<folder>/<topic>.md with a # title and optional WIP banner.
3. Add an entry to src/SUMMARY.md under the right section.
4. Link from the section overview.md if one exists.
5. Add verification rows if the page contains benchmarks or unconfirmed claims.
6. Run mdbook build — fix broken links before merging.

Splitting an existing page

1. Move shared intro to overview.md or trim the original.
2. Create focused sub-pages; link between them at the bottom (“Next steps”).
3. Update SUMMARY.md to nest sub-pages under the section.
4. Grep for old filename links: rg 'old-name\.md' src/
5. Update reference/verification.md section headers to match new paths.

What goes where?

Useful Resources

Curated links for gfx1030 / RDNA2 work. Grouped by type — not everything here is “official” or wiki-maintained; see Verification status for claim audits on guides that reference these projects.

Community Discord

Primary hangout for V620 / gfx1030 inference work. Wiki pages often cite channel names (#vllm-rdna, #llamacpp, #general, #benchmarks, #lmcache) — that is where the claim came from.

Official documentation

• ROCm documentation
• ROCm install on Linux
• System requirements & supported GPUs
• GPU architecture specs (Navi 21 = gfx1030)
• ROCm Device Support Wishlist — community-tracked support matrix on the official ROCm repo

Docker images

Prebuilt images for RDNA (gfx1030 through RDNA4). See Running vLLM for run commands and Building images for how they are produced.

Tags are listed on Docker Hub and in Running (Docker). Re-pull before debugging — tags are refreshed in place when fixes land.

Upstream projects

Stock projects this wiki builds on or tracks. Use these when you want upstream behavior or to compare against forks.

• vLLM — upstream serving engine
• llama.cpp — GGUF inference (HIP/ROCm and Vulkan)
• vLLM PR #52391 — RDNA gfx1030 platform detection (upstream CI enablement for consumer Radeon)
• TheRock: RDNA2 (gfx103X) support — ROCm/TheRock gfx1030 enablement work

Community projects

Forks, recipes, and tooling maintained outside (or alongside) upstream. Often where gfx1030-specific performance work happens first.

Tools & write-ups

• Disabling ECC on Radeon Pro GPUs (lunnova.dev) — amdgpu.ras_enable=0 + two reboots; pinned in #vllm-rdna. See Disabling ECC.
• rocminfo, rocm-smi, amd-smi, clinfo — GPU visibility, topology, and P2P checks. See Environment variables.

Know a gfx1030 resource that belongs here? Contribute a link and say which section it fits (official / Docker / upstream / community).

Contributing

This is a community wiki for the gfx1030 (RDNA2 / Navi 21) GPU ecosystem. Contributions of all sizes are welcome — fixing a typo, correcting an out-of-date command, or adding a whole new guide.

Live discussion happens on the gfx1030 club Discord: https://discord.gg/mESex2aBp. Propose wiki fixes there (#wiki) or open a GitHub PR — both are fine. See Useful resources.

A daily Cursor Automation (when enabled) reads Discord and opens draft PRs titled wiki: discord digest YYYY-MM-DD. Treat those like any other contribution: review privacy and accuracy before merging. Setup lives in the repo at .cursor/automations/daily-discord-wiki.md (not loaded by Cursor automatically — activate once at cursor.com/automations/new).

Two ways to contribute

1. Directly on GitHub (no command line)

1. Open the page you want to edit on GitHub and click the edit (pencil) button, or use the “Suggest an edit” link on the mdBook page. You’ll be prompted to fork the repo — do so.
2. Make your changes, keeping the markdown clean (see style below).
3. Commit with a descriptive message.
4. Open a pull request targeting the master branch.

2. Locally with mdBook (lets you preview)

git clone https://github.com/blivioniag/gfx1030-wiki.git
cd gfx1030-wiki

# Install mdBook (prebuilt binary, no Rust required)
mkdir -p "$HOME/.local/bin"
MDBOOK_VERSION=v0.5.4
curl -sL "https://github.com/rust-lang/mdBook/releases/download/${MDBOOK_VERSION}/mdbook-${MDBOOK_VERSION}-x86_64-unknown-linux-gnu.tar.gz" \
  | tar -xz -C "$HOME/.local/bin"
export PATH="$HOME/.local/bin:$PATH"

mdbook serve   # live preview at http://localhost:3000

Adding a new page

1. Read Wiki structure — pick the right folder and naming convention.
2. Create a markdown file under src/<section>/, e.g. src/vllm/my-topic.md.
3. Add a link in src/SUMMARY.md and the section overview.md if one exists.
4. Run mdbook build to confirm it renders and there are no broken links.
5. Update Verification status for benchmarks or community claims.

Markdown style

• Use ATX headings (#, ##, …) and start each page with a single # title.
• Wrap shell commands in fenced code blocks with a language hint (```sh).
• Prefer relative links between wiki pages (e.g. ../vllm/overview.md) so they work locally and deployed.
• Keep lines readable; hard-wrapping around ~100 columns is fine but not required.

Accuracy

RDNA2 tooling moves quickly. When you add a command, note the ROCm / library version you tested it with if it might matter, and prefer linking to official docs over pasting version-specific numbers that will age.

For pages sourced from Discord or fork release notes, add or update entries in Verification status so readers know what is solid, fork-source, community-reported, or needs verify. Mark unconfirmed throughput claims as community-reported rather than presenting them as wiki-tested facts.

Privacy (Discord → wiki)

This wiki is public on GitHub Pages. When turning Discord messages into docs:

Public GitHub / Hugging Face org or repo names are fine to link (they are already public attribution). Prefer “community report” / “fork README” over naming Discord people even when the same handle appears on GitHub.

Never commit: DISCORD_BOT_TOKEN, DISCORD_GUILD_ID, .env files, or screenshots of private channels. Cloud Agents should set token + guild ID in Secrets only (see repo README.md).

Automated Discord digest PRs

When the daily digest automation is active, expect occasional draft PRs that only touch src/ (and SUMMARY.md if a page was added). Review them like a human patch:

• Reject transcripts, @mentions, message links, and personal host/path leaks.
• Keep unconfirmed tok/s and one-off anecdotes in Verification status as Community / Needs verify.
• Empty calendar days should produce no PR; do not “fill” the wiki from stale chat.

The agent instructions are .cursor/skills/discord-wiki-digest/SKILL.md. You can still ask a Cloud Agent for an on-demand pass; it should follow the same skill.

Deployment

Merges to master are built and published to GitHub Pages automatically by the .github/workflows/mdbook.yml workflow — no manual deploy step is needed.