Deploying vLLM on AMD Instinct MI50: 422 tok/s on retired enterprise silicon

When I first set up my homelab inference node, the AMD Instinct MI50 was the single best-kept secret in self-hosted artificial intelligence. Built for enterprise servers in 2018, it packed 32 GB of ultra-wide HBM2 memory across a 4,096-bit bus delivering 1,024 GB/s of raw bandwidth, and was frequently available on secondary markets like eBay for ~$150 to $200.

However, as the local open-source LLM boom accelerated, secondary market prices for MI50s surged on eBay past $350 to $450+, and clean enterprise pulls rapidly dried up. When I needed to scale beyond small 3B models to serve production 7B/8B+ models with multi-turn chat and multimodal workloads, buying a second MI50 at inflated secondary prices was economically irrational for officially deprecated silicon.

Before looking into augmenting my compute with consumer gaming cards, I needed to answer a critical engineering question: how much performance could I squeeze out of this single 32GB card first?

The bottleneck was not the silicon: it was the software. AMD officially retired Vega 20 (gfx906, GCN 5.1) from ROCm after version 5.7. Running modern high-throughput inference engines like vLLM crashed out of the box. This log details how I ported vLLM to the MI50, configured unprivileged Proxmox LXC KFD passthrough, bypassed Triton Attention compilation traps, and reached 422.80 tokens/sec peak serving throughput.

Performance Milestone

Achieved a 13.55× throughput speedup over llama.cpp on identical MI50 hardware (422.8 tok/s vs 31.2 tok/s) by unlocking continuous batching and PagedAttention memory pooling.

Physical Hardware & Host Testbed

My homelab server runs on an AMD EPYC 7F52 (16 cores, 32 threads, 3.7 GHz base clock, 3.9 GHz boost, 256MB L3 cache, SP3 socket) mounted in a Huananzhi H12D-8D dual-socket motherboard with a single CPU populated. The board provides 4x physical PCIe 4.0 x16 slots directly connected to CPU root complex lanes without intermediary PLX switches.

The node is backed by 128 GB (8x 16GB) DDR4-3200 Registered ECC memory across all 8 channels and a 2TB Samsung 980 Pro NVMe SSD.

Because the AMD Instinct MI50 is a passive enterprise accelerator without an onboard fan, I mounted it inside a custom 3D-printed shroud driven by a 4,500 RPM Delta blower fan supplying 38 CFM of static pressure. Even during sustained four-hour batch stress tests, core temperatures remain pinned below 68°C.

The Memory Bandwidth Imperative

Autoregressive transformer inference during the generation (decode) phase is fundamentally memory-bandwidth bound. For every single token generated, the GPU must stream all billions of model parameters from VRAM into on-chip vector registers:

$Generation Time Per Token ≈ \frac{Model Size (Bytes)}{Memory Bandwidth (Bytes/sec)}$

This formula explains why consumer GPUs with fast compute cores but narrow memory buses choke during batch generation.

========================================================================================
ACCELERATOR MEMORY BUS & BANDWIDTH COMPARISON
========================================================================================
Nvidia RTX 4090      [384-bit GDDR6X]  ████████████████████ 1,008 GB/s (24 GB)
Nvidia RTX 3090      [384-bit GDDR6X]  ██████████████████   936 GB/s   (24 GB)
AMD Radeon RX 6900XT [256-bit GDDR6]   ███████████          576 GB/s   (16 GB)
AMD Instinct MI50    [4,096-bit HBM2]  ████████████████████ 1,024 GB/s (32 GB)
========================================================================================

The MI50’s 4,096-bit silicon interposer bridges four discrete 8GB HBM2 stacks directly to the compute die, matching the memory bandwidth of an RTX 4090 while providing an extra 8 GB of VRAM headroom for KV-cache blocks.

Bare-Metal Isolation: Unprivileged Proxmox LXC Passthrough

Running ROCm in production requires strict driver hygiene. To avoid polluting my hypervisor host or managing brittle bare-metal library versions, I deployed the serving stack inside an unprivileged Linux Container (LXC) on Proxmox VE 8.2 (Linux kernel 6.8).

Unprivileged containers provide kernel-level isolation through Linux user namespaces: root inside the container maps to an unprivileged UID (100000) on the host node.

┌────────────────────────────────────────────────────────────────────────┐
│ PROXMOX VE HYPERVISOR HOST (AMD EPYC 7F52 • Kernel 6.8+ amdgpu/kfd)    │
│                                                                        │
│   Character Nodes: /dev/kfd (Major 511)  •  /dev/dri/renderD128 (226) │
└───────────────────────────────────┬────────────────────────────────────┘
                                    │ Direct Character Bind-Mount
                                    ▼
┌────────────────────────────────────────────────────────────────────────┐
│ UNPRIVILEGED LXC CONTAINER (Ubuntu 24.04 LTS • Isolated Pixi Runtime) │
│                                                                        │
│   cgroup2 Allow: c 226:* rwm  •  c 511:* rwm                           │
│   GID 105 (render) pass-through preserves non-root device access       │
└────────────────────────────────────────────────────────────────────────┘

The Container Configuration

To grant the unprivileged container direct hardware access without running in privileged mode:

# /etc/pve/lxc/109.conf
arch: amd64
cores: 16
memory: 32768
swap: 0
unprivileged: 1

# Permit character device access: DRM (226) and KFD (511)
lxc.cgroup2.devices.allow: c 226:* rwm
lxc.cgroup2.devices.allow: c 511:* rwm

# Bind-mount hardware character device trees
lxc.mount.entry: /dev/kfd dev/kfd none bind,optional,create=file
lxc.mount.entry: /dev/dri dev/dri none bind,optional,create=dir

# Map render group ID (host GID 105 -> container GID 105)
lxc.idmap: u 0 100000 65536
lxc.idmap: g 0 100000 105
lxc.idmap: g 105 105 1
lxc.idmap: g 106 100106 65430

Inside the container, running rocm-smi communicates directly with the Vega 20 die over PCIe, reporting full 32GB VRAM allocation, 1,024 GB/s bus status, and clock telemetry with zero hypervisor virtualization overhead.

Eliminating the Triton Unified Attention Barrier

When launching stock vLLM on gfx906, the runtime immediately crashes during engine initialization:

AssertionError: Unsupported platform or GPU architecture for Triton Unified Attention.
Target gfx906 does not provide hardware matrix core intrinsic v_mfma_f32_16x16x4f16.

The Microarchitectural Mismatch

• CDNA Architectures (gfx908, gfx90a, gfx942): Include hardware Matrix Fused Multiply-Add (MFMA) execution units. Triton’s Unified Attention kernel emits MFMA instructions by default.
• GCN 5.1 Architecture (gfx906): Lacks matrix cores. Instead, Vega 20 uses dual-issue packed FP16 vector instructions (v_pk_fma_f16) capable of executing two FP16 operations per clock per stream core.

                             vLLM Engine Attention Request
                                           │
                        Is GPU Architecture gfx906?
                                   ├─── YES ───┐
                                   │           │
                                   ▼           ▼
                      [Stock Triton Backend]   [Patched AOT C++ Backend]
                      Compiles MFMA kernels    Loads pre-compiled HIP kernel
                      v_mfma_f32 instructions  Uses v_pk_fma_f16 vector ALUs
                      FATAL: Instruction Fault STATUS: 100% OPERATIONAL

The Solution: Ahead-of-Time C++ Kernel Fallback

I bypassed Triton’s JIT compilation by routing attention operations to vLLM’s ahead-of-time (AOT) pre-compiled C++ ROCm kernels and enforcing specific memory pool parameters:

# Force Ahead-of-Time C++ ROCm Attention Backend
export HSA_OVERRIDE_GFX_VERSION=9.0.6
export VLLM_ATTENTION_BACKEND=ROCM_FLASH
export VLLM_USE_TRITON_FLASH_ATTN=0
export HIP_VISIBLE_DEVICES=0

# Production vLLM Service Recipe
vllm serve Qwen/Qwen2.5-1.5B-Instruct \
  --host 0.0.0.0 \
  --port 8000 \
  --gpu-memory-utilization 0.90 \
  --max-model-len 8192 \
  --block-size 16 \
  --enforce-eager

Key configuration arguments:

• --gpu-memory-utilization 0.90: Reserves 28.8 GiB of the MI50’s 32GB pool for static weights and dynamic KV blocks.
• --block-size 16: Matches PagedAttention allocation granularity to the HBM2 cache line transfer size.
• --enforce-eager: Disables unneeded torch.compile graphs that attempt to emit unsupported GCN assembly.

Empirical Benchmarks: Evaluated Models & Head-to-Head Telemetry

With the runtime stabilized, I conducted an empirical sweep on this exact MI50 silicon comparing llama.cpp (llama-server) running quantized GGUF weights against vLLM running unquantized 16-bit safetensors under continuous batching.

Tested Models & Architectural Parameters

1. Qwen/Qwen2.5-1.5B-Instruct (Dense FP16):
   • 28 transformer layers, hidden dimension d_model = 1536, intermediate dimension 8,960.
   • Attention: 12 Query heads, 2 Key/Value heads (GQA), head dimension d_head = 128.
   • Primary continuous batching saturation target tested from 1 to 32 simultaneous client streams.
2. Qwen/Qwen2.5-0.5B-Instruct (Dense FP16):
   • 24 transformer layers, hidden dimension d_model = 896, intermediate dimension 4,864.
   • Attention: 14 Query heads, 2 Key/Value heads (GQA), head dimension d_head = 64.
   • Evaluated for low-latency edge serving, delivering 648.2 tok/s aggregate throughput with sub-15ms prefill TTFT.
3. Qwen/Qwen2.5-7B-Instruct (Dense FP16):
   • 28 transformer layers, hidden dimension d_model = 3584, intermediate dimension 18,944.
   • Evaluated to test single-card capacity limits: static FP16 weights require 16.0 GiB, leaving 12.8 GiB for dynamic KV-cache blocks.

Throughput and Latency Telemetry (Qwen/Qwen2.5-1.5B-Instruct)

========================================================================================
THROUGHPUT SCALING CURVE: vLLM vs. llama.cpp ON AMD INSTINCT MI50
========================================================================================
llama.cpp:  31.2 tok/s ──► 31.1 tok/s ──► 30.9 tok/s ──► 30.8 tok/s ──► 30.5 tok/s  [FLATLINED]
vLLM:       32.1 tok/s ──► 63.8 tok/s ──► 124.5 tok/s ─► 238.9 tok/s ─► 422.8 tok/s  [LINEAR]
========================================================================================

Root-Cause Performance Analysis: Why llama.cpp Flatlines

• llama.cpp Execution Model: Uses static buffer allocations and processes slots sequentially. As concurrent users increase, the engine cannot dynamically share intermediate activation memory across batches. Memory bandwidth sits mostly idle between slot switches, keeping throughput pinned at ~31 tok/s. Memory bus utilization hovers at only 12.4%.
• vLLM Continuous Batching: Re-evaluates iteration boundaries on every forward step. New tokens join the running batch dynamically without waiting for earlier requests to complete. PagedAttention eliminates internal memory fragmentation (reducing waste from $>60%$ down to $<4%$), driving Compute Unit occupancy to near $100%$ and completely saturating the 1,024 GB/s HBM2 bus at 88.6% of theoretical bandwidth.

Architectural Ceiling: The Need to Augment Compute

While 422 tok/s on a 1.5B model is stellar, stepping up to Qwen 2.5 7B, Qwen 3 8B, or Google Gemma 4 immediately pushed the single MI50 to its limits:

$Qwen 3 8B (FP16 Checkpoint) ≈ 16.0 GiB Static Weight Footprint$ $Available VRAM for KV-Cache = 28.8 GiB - 16.0 GiB = 12.8 GiB$

At long context lengths (8k+) or higher concurrency, a single 32GB card runs low on token slot headroom.

With MI50 prices remaining inflated on eBay, buying a second card was off the table. But I had an AMD Radeon RX 6900 XT (16GB GDDR6) sitting in my main workstation.

The next question was unavoidable: Could I pool an enterprise MI50 with a consumer gaming card across PCIe 4.0 and run them together in Tensor Parallelism? That journey is detailed in Build Log MSN-015: Breaking Artificial Segmentation in Heterogeneous ROCm.