Compile vllm from scratch, AMD R9700, ROCm, gfx1201

README.md

vllm — ROCm build for AMD RDNA4 (gfx1201)

Build and runtime notes for the AMD Radeon AI PRO R9700 (and RX 9070 XT) on ROCm 7.2.

---

Why build from source?

AMD's RDNA4 architecture (gfx1201 / Navi 48) is new enough that pre-built vllm wheels do not target it. The official pip package is compiled for CUDA, and the AMD-published ROCm wheels are built for MI300-series datacenter GPUs (gfx942). Installing either will either fail at import or silently miscompile kernels for the wrong ISA.

Building from source with PYTORCH_ROCM_ARCH=gfx1201 produces HIP kernels compiled specifically for RDNA4's instruction set, which is meaningfully different from the MI300 and RDNA3 targets.

---

Prerequisites

Requirement	Version in use
ROCm	7.2.4 at /opt/rocm
Python	3.12 (.venv/)
uv	0.11+
Rust / cargo	1.95 (for vllm's Rust frontend)
cmake	3.26+ (system)

System packages ninja-build and protobuf-compiler are installed by the script via apt-get.

Directory layout expected

vllm/               ← this directory (working root)
├── .venv/          ← Python 3.12 virtual environment
├── build_rocm.sh   ← build script
└── vllm/           ← vllm source (git clone of vllm-project/vllm)

Clone the source if you haven't already:

git clone https://github.com/vllm-project/vllm.git

---

Build

./build_rocm.sh 2>&1 | tee build.log

The script runs five steps:

1. System tools — installs ninja-build and protobuf-compiler via apt-get
2. PyTorch — installs torch==2.11.0+rocm7.2 (cp312) from the official PyTorch wheel index, plus matching torchvision, torchaudio, and triton
3. Python build deps — setuptools, setuptools-scm, setuptools-rust, jinja2, tilelang, etc.
4. Rust frontend — builds vllm-rs and the _rust_tool_parser PyO3 extension via build_rust.sh
5. HIP compilation — uv pip install -e . with --no-build-isolation so the already-installed ROCm torch is used rather than a CPU-only fallback

--no-build-isolation is critical: without it pip creates a fresh isolated environment, finds no ROCm index, and pulls a CPU-only torch that breaks the entire HIP build.

amdsmi version

vllm's requirements/build/rocm.txt pins amdsmi==7.0.2, but ROCm 7.2.4 ships its own amdsmi with a different version scheme — AMD switched to calendar versioning, so the installed package is 26.2.2 (year 2026, month 2). The two version strings are not comparable; 26.2.2 is the newer, correct package for this ROCm installation.

The build script intentionally skips reinstalling amdsmi from pip. Letting pip resolve amdsmi==7.0.2 would either fail outright or downgrade to a package built against an older ROCm ABI. The version already in .venv (26.2.2+97f5574fe2, installed from the ROCm package tree) is what vllm uses at runtime to detect the GPU via amdsmi_init() / amdsmi_get_processor_handles() in platforms/__init__.py. If that call returns GPU handles, vllm knows it is on a ROCm system and loads the ROCm platform plugin. Do not reinstall or pin amdsmi from pip.

If amdsmi is not yet in your venv, the source tree is at /opt/rocm-7.2.4/share/amd_smi but it is owned by root, so uv pip install on that path directly will fail with a permission error when it tries to create a build cache. Copy it somewhere writable first:

cp -r /opt/rocm-7.2.4/share/amd_smi /tmp/amd_smi
uv pip install /tmp/amd_smi

Verify afterwards:

.venv/bin/python -c "import amdsmi; print(amdsmi.__version__)"

First build takes 20–40 minutes. The bottleneck is HIP kernel compilation. Subsequent builds reuse cmake's incremental build cache inside vllm/build/.

---

Runtime

source .venv/bin/activate
vllm serve <model>

Or without activating:

.venv/bin/vllm serve <model>

Environment variables worth knowing

Variable	Default	Effect
PYTORCH_ROCM_ARCH	gfx1201	Set by the build; leave it in your shell for any tools that re-compile kernels (e.g. triton JIT)
ROCBLAS_USE_HIPBLASLT	1	Prefers hipBLASLt over rocBLAS for GEMM on gfx12xx; measurably faster
VLLM_ROCM_USE_AITER	1 (but ineffective)	AITER is MI300-only; setting this has no effect on gfx1201
HIP_VISIBLE_DEVICES	unset	Same role as CUDA_VISIBLE_DEVICES
VLLM_WORKER_MULTIPROC_METHOD	fork	Default; fine for single-GPU

---

Attention backends on gfx1201

vllm selects an attention backend per attention layer based on what the layer type requires and what the hardware supports. On gfx1201 you get:

Standard decoder attention (Llama, Mistral, Qwen, Gemma, …)

ROCM_ATTN is selected automatically. This is vllm's native ROCm attention path: a C++ paged-attention kernel for decode, and a Triton chunked-prefill kernel for prefill. Block sizes must be multiples of 16; non-standard sizes (e.g. Qwen3's 544 or 784) are transparently routed to an optimised Triton kernel.

TRITON_ATTN is also available and can be forced with --attention-backend TRITON_ATTN if you need sink tokens, KV connectors, or ENCODER_DECODER attention type in a custom model.

MTP and Eagle3 speculative decoding

Both use standard decoder attention on the draft model's layers. The draft runner (Step3p5MTPProposer / EagleProposer) inherits whatever backend was assigned to each attention layer — no special requirement. ROCM_ATTN handles it.

DFlash speculative decoding

DFlash is a parallel-drafting method where the draft model attends to the target model's hidden states via cross-attention (ENCODER_DECODER type). ROCM_ATTN explicitly does not support ENCODER_DECODER (it would produce incorrect results for cached cross-attention K/V). vllm detects this per-layer and falls back automatically:

• Decoder self-attention layers → ROCM_ATTN
• Cross-attention layers → TRITON_ATTN (fallback; supports all attention types)

You do not need to configure anything. The split is transparent at runtime; you will see log lines like:

Found incompatible backend(s) [ROCM_ATTN] with ENCODER_DECODER.
Overriding with TRITON_ATTN out of potential backends: ['TRITON_ATTN', 'TURBOQUANT'].

This is expected and correct behaviour.

MLA attention (DeepSeek-V2 / V3 / R1)

TRITON_MLA only. The AITER MLA backends (ROCM_AITER_MLA, ROCM_AITER_TRITON_MLA) require MI300-series hardware and are not available on gfx1201.

ViT attention (multimodal — LLaVA, Qwen-VL, …)

TORCH_SDPA by default.

To enable Flash Attention's Triton path for ViT layers, install a ROCm build of flash-attn and set:

FLASH_ATTENTION_TRITON_AMD_ENABLE=TRUE vllm serve <model>

---

What is NOT available on gfx1201

The AITER library (aiter package) contains precompiled ASM kernels tuned for MI300's matrix engines. Its entire backend family is hard-gated behind on_mi3xx() in vllm's platform code and will not activate on RDNA4, regardless of whether the package is installed:

• ROCM_AITER_FA — MI300 flash attention
• ROCM_AITER_MLA / ROCM_AITER_TRITON_MLA — MI300 MLA
• ROCM_AITER_UNIFIED_ATTN — unified attention with FP8 fused output quant

Installing aiter is harmless but pointless on this hardware.

---

FP8 inference

gfx1201 supports FP8 (float8_e4m3fn — standard IEEE format, not the MI300-specific e4m3fnuz). Models quantised to FP8 will run correctly. The supported quantisation schemes are listed in vllm/platforms/rocm.py:RocmPlatform.supported_quantization and include fp8, deepseek_v4_fp8, compressed-tensors, fbgemm_fp8, and others.

---

Rebuilding after a source update

cd vllm
git pull
cd ..
./build_rocm.sh 2>&1 | tee build.log

cmake will perform an incremental rebuild; only changed files are recompiled. Pure Python changes (no .cu / .hip / .cpp modifications) are picked up immediately without rebuilding because the install is editable (-e).

build_rocm_vllm.sh

#!/usr/bin/env bash
# Build vllm from source for AMD Radeon AI PRO R9700 (gfx1201 / RDNA4) with ROCm 7.2
#
# Prerequisites (satisfied on this machine):
#   - ROCm 7.2.4 at /opt/rocm
#   - uv, cargo/rustc 1.95, cmake 3.31 already installed
#   - .venv with Python 3.12 in the same directory as this script
#   - vllm source in ./vllm/
set -euo pipefail

SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
VLLM_SRC="$SCRIPT_DIR/vllm"
VENV="$SCRIPT_DIR/.venv"
ROCM_WHEEL_INDEX="https://download.pytorch.org/whl/rocm7.2"

# ── Sanity checks ─────────────────────────────────────────────
[[ -d "$VLLM_SRC" ]]           || { echo "ERROR: $VLLM_SRC not found"; exit 1; }
[[ -f "$VENV/bin/activate" ]]  || { echo "ERROR: $VENV venv not found"; exit 1; }
[[ -d /opt/rocm ]]             || { echo "ERROR: /opt/rocm not found"; exit 1; }

echo "========================================================"
echo "  vllm ROCm build — gfx1201 (RDNA4 R9700) + ROCm 7.2"
echo "  $(date)"
echo "========================================================"

# ── ROCm environment ──────────────────────────────────────────
export ROCM_HOME="/opt/rocm"
export PATH="$ROCM_HOME/bin:$PATH"
export LD_LIBRARY_PATH="$ROCM_HOME/lib:${LD_LIBRARY_PATH:-}"
# Compile vllm HIP extensions only for the GPU in this machine.
# gfx1201 = RDNA4 Navi 48 (R9700 / RX 9070 XT)
export PYTORCH_ROCM_ARCH="gfx1201"
export VLLM_TARGET_DEVICE="rocm"
# hipBLASLt is preferred over rocBLAS for gfx12xx performance
export ROCBLAS_USE_HIPBLASLT=1
# Keep UV network timeout generous — ROCm wheels are large
export UV_HTTP_TIMEOUT=500
export UV_INDEX_STRATEGY="unsafe-best-match"

# ── Activate venv ─────────────────────────────────────────────
# shellcheck source=/dev/null
source "$VENV/bin/activate"
PYTHON="$VENV/bin/python"
echo "Python: $($PYTHON --version)"

# ── Step 1: System build tools ────────────────────────────────
echo ""
echo "[1/5] System build tools (ninja, protoc)"
# ninja: parallel HIP compilation
# protoc: required by prost-build / tonic-build in the Rust frontend
sudo apt-get install -y -q ninja-build protobuf-compiler
echo "  ninja $(ninja --version)  protoc $(protoc --version)"

# ── Step 2: PyTorch for ROCm 7.2 ──────────────────────────────
echo ""
echo "[2/5] PyTorch 2.11.0 + ROCm 7.2 wheels (cp312, ~2 GB)"
# These wheels are built against ROCm 7.2 and support gfx1201 natively.
# rocm7.1 wheels (the project default) would mismatch our installed 7.2.4.
uv pip install \
  "torch==2.11.0" \
  "torchvision==0.26.0" \
  "torchaudio==2.11.0" \
  --extra-index-url "$ROCM_WHEEL_INDEX"

# triton is platform-agnostic but lives on the same index
uv pip install "triton==3.6.0" \
  --extra-index-url "$ROCM_WHEEL_INDEX"

echo ""
echo "── PyTorch GPU smoke test ──"
"$PYTHON" - <<'PYEOF'
import os, torch
print(f"  torch {torch.__version__}  HIP {torch.version.hip}")
n = torch.cuda.device_count()
print(f"  {n} GPU(s) visible")
for i in range(n):
    print(f"    [{i}] {torch.cuda.get_device_name(i)}")
if n == 0:
    print(f"  WARNING: no GPUs detected — ROCM_HOME={os.environ.get('ROCM_HOME','?')}")
PYEOF

# ── Step 3: Python build dependencies ─────────────────────────
echo ""
echo "[3/5] Python build dependencies"
uv pip install \
  "cmake>=3.26.1,<4" \
  ninja \
  packaging \
  "setuptools>=77.0.3,<80.0.0" \
  "setuptools-scm>=8" \
  "setuptools-rust>=1.9.0" \
  wheel \
  "jinja2>=3.1.6" \
  "timm>=1.0.17" \
  "tilelang==0.1.10"
# NOTE: amdsmi is intentionally NOT reinstalled from pip.
# The venv already has amdsmi 26.x from the ROCm 7.2 package tree,
# which is newer than the pinned 7.0.2 in requirements/build/rocm.txt.

# ── Step 4: Rust frontend ──────────────────────────────────────
echo ""
echo "[4/5] Building Rust frontend (vllm-rs + tool parser)"
cd "$VLLM_SRC"
# The venv activation ensures python3 → .venv/bin/python3 (3.12)
bash build_rust.sh

# ── Step 5: vllm C++ / HIP extensions ─────────────────────────
echo ""
echo "[5/5] Building vllm (HIP kernel compilation)"
echo "      PYTORCH_ROCM_ARCH = $PYTORCH_ROCM_ARCH"
echo "      MAX_JOBS          = $(nproc)"
echo "      ETA               : 20-40+ minutes (first build)"
echo ""
cd "$VLLM_SRC"
# --no-build-isolation: use the venv's torch/setuptools instead of letting
#   pip create an isolated env and pull in a CPU-only torch from PyPI.
MAX_JOBS=$(nproc) \
uv pip install -e . \
  --no-build-isolation \
  --extra-index-url "$ROCM_WHEEL_INDEX"

# ── Verification ──────────────────────────────────────────────
echo ""
echo "========================================================"
echo "  Build complete!  $(date)"
echo "========================================================"
"$PYTHON" -c "
import vllm, torch
print(f'vllm  {vllm.__version__}')
print(f'torch {torch.__version__}  HIP {torch.version.hip}')
n = torch.cuda.device_count()
print(f'{n} GPU(s) visible')
"
echo ""
echo "Usage:"
echo "  source .venv/bin/activate"
echo "  vllm serve <model>"
echo ""
echo "If you hit AITER kernel errors at runtime on gfx1201, disable it:"
echo "  VLLM_ROCM_USE_AITER=0 vllm serve <model>"

install_flash_attn_rocm.sh

#!/usr/bin/env bash
# Install flash-attn from ROCm/flash-attention (Triton backend) for gfx1201 / ROCm 7.2
#
# Why the Triton backend?
#   The default CK (Composable Kernel) backend has hand-tuned ASM for MI200/MI300.
#   For RDNA4 (gfx1201) use the Triton backend instead: it JIT-compiles at first
#   use via Triton, so it works on any architecture in the allowed list including
#   gfx1201 — with no C++ compilation needed during install.
#
# What this unlocks in vllm:
#   - flash_attn.flash_attn_triton_amd is present → ViT flash attention on multimodal
#     models (LLaVA, Qwen-VL, etc.) with FLASH_ATTENTION_TRITON_AMD_ENABLE=TRUE
#   - FLASH_ATTN backend becomes available (not auto-selected on ROCm, but forceable)
#   - Does NOT unlock FLASH_ATTN_MLA (requires compute capability major == 9, i.e. MI300)
set -euo pipefail

SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
VENV="$SCRIPT_DIR/.venv"
SRC_DIR="$SCRIPT_DIR/flash-attention"

# ── Sanity checks ─────────────────────────────────────────────
[[ -f "$VENV/bin/activate" ]] || { echo "ERROR: $VENV not found — run build_rocm.sh first"; exit 1; }
[[ -d /opt/rocm ]]            || { echo "ERROR: /opt/rocm not found"; exit 1; }

echo "======================================================"
echo "  flash-attn ROCm (Triton) — gfx1201 / ROCm 7.2"
echo "  $(date)"
echo "======================================================"

# ── ROCm environment ──────────────────────────────────────────
export ROCM_HOME="/opt/rocm"
export PATH="$ROCM_HOME/bin:$PATH"
export LD_LIBRARY_PATH="$ROCM_HOME/lib:${LD_LIBRARY_PATH:-}"

# ── flash-attn build flags ────────────────────────────────────
# Triton backend: no CK composable_kernel C++ compilation, JIT at runtime instead.
export FLASH_ATTENTION_TRITON_AMD_ENABLE="TRUE"
# Scope the GPU architecture so no native-detect fallback is attempted.
export GPU_ARCHS="gfx1201"
# Force ROCm build path in case hipcc is present alongside nvcc.
export BUILD_TARGET="rocm"
# Suppress version mismatch warnings between the bundled aiter and torch.
export UV_HTTP_TIMEOUT=500

# ── Activate venv ─────────────────────────────────────────────
source "$VENV/bin/activate"
echo "Python: $(python --version)"

# ── Step 1: Source ────────────────────────────────────────────
echo ""
echo "[1/4] Fetching ROCm/flash-attention source (main branch)"
if [[ -d "$SRC_DIR/.git" ]]; then
    echo "      Source already present — updating"
    git -C "$SRC_DIR" fetch --depth 1 origin main
    git -C "$SRC_DIR" reset --hard origin/main
else
    git clone https://github.com/ROCm/flash-attention.git "$SRC_DIR" \
        --branch main --depth 1
fi

cd "$SRC_DIR"

# ── Step 2: aiter submodule ───────────────────────────────────
# setup.py also does this, but doing it here gives a clear progress line
# and avoids a confusing mid-install pause.
echo ""
echo "[2/4] Initialising third_party/aiter submodule (Triton backend)"
git submodule update --init --depth 1 third_party/aiter

# ── Step 3: Install flash-attn ────────────────────────────────
echo ""
echo "[3/4] Installing flash-attn (Triton backend — no C++ compile, should be fast)"
echo "      setup.py will also pip-install third_party/aiter into the venv"
# --no-build-isolation: lets setup.py see the venv's torch and the env vars above.
# The internal 'pip install third_party/aiter' inside setup.py uses sys.executable
# which resolves to .venv/bin/python because the venv is activated.
uv pip install -e . --no-build-isolation

# ── Step 4: Verify ────────────────────────────────────────────
echo ""
echo "[4/4] Verifying"
python - <<'PYEOF'
import importlib, sys

# flash_attn itself
import flash_attn
print(f"  flash-attn {flash_attn.__version__}")

# The Triton AMD submodule — this is what vllm checks before enabling ViT flash attn
if importlib.util.find_spec("flash_attn.flash_attn_triton_amd") is not None:
    print("  flash_attn.flash_attn_triton_amd: FOUND  ✓")
else:
    print("  flash_attn.flash_attn_triton_amd: NOT FOUND  ✗")
    print("  (built without Triton backend? check FLASH_ATTENTION_TRITON_AMD_ENABLE=TRUE)")

# fa_utils._ROCM_FLASH_ATTN_AVAILABLE confirms vllm can use it
try:
    from vllm.v1.attention.backends.fa_utils import _ROCM_FLASH_ATTN_AVAILABLE
    print(f"  vllm _ROCM_FLASH_ATTN_AVAILABLE: {_ROCM_FLASH_ATTN_AVAILABLE}  {'✓' if _ROCM_FLASH_ATTN_AVAILABLE else '✗'}")
except ImportError:
    print("  (vllm not importable from here — expected if run standalone)")
PYEOF

echo ""
echo "======================================================"
echo "  Done!"
echo "======================================================"
echo ""
echo "To activate ViT flash attention for multimodal models:"
echo "  FLASH_ATTENTION_TRITON_AMD_ENABLE=TRUE vllm serve <model>"
echo ""
echo "Note: Triton kernels JIT-compile on first use — expect a short"
echo "      pause the first time a model is loaded."