vllm-mi210

在 AMD Instinct MI210（gfx90a / CDNA2） 上跑 官方 vLLM 0.28，不换 davetha fork。

官方 ROCm 轮子里的 AITER gemm_a8w8 只编了 gfx942 / gfx950；fused GDN decode 核是 CUDA-only。本仓库补两件事：

1. 给官方 0.28 开 gfx90a 闸门，JIT 出 CK INT8 gemm_a8w8
2. 一份 HIP fused GDN decode 核（runtime GQA，Qwen3.8 是 3:1，不是 CUDA 核写死的 8:1）

English

实测（1× MI210 64 GB，2026-08-31）

模型：Freaksterz/Qwen3.8-27B-SmoothQuant-W8A8-INT8，MTP n=3，max-model-len 262144，thinking 关，temp 0，非流式 completion_tokens / wall（256 tok × 3）。

Linear kernel	GDN decode	tok/s	MTP 接受率
官方 Triton INT8	Triton	18.71	~43%
AITER CK INT8	Triton	54.32	~43%
AITER CK INT8	本仓库 HIP 核	53.54	~43%

对照：同卡 llama.cpp MTP 约 40 tok/s。DFlash2 n=8 在这份 W8A8 上接受率 0，不要用。MTP n=2 比 n=3 慢。

结论：真正的加速来自 CK INT8 GEMM（约 2.9× Triton）。HIP GDN 数值对得上，但 GEMM 吃上之后 48 层 GDN 已经不是瓶颈。

不要做的事

• 不要换 davetha / 其它 vLLM fork 当生产框架
• 不要 --kv-cache-dtype fp8（gfx90a 无 FP8 MFMA，会静默回退或崩）
• 不要 PYTORCH_CUDA_ALLOC_CONF=expandable_segments
• 不要用官方 FP8 / bitsandbytes / llama.cpp GGUF 给 vLLM
• 不要设 HF_ENDPOINT=hf-mirror.com 同时走 Hub
• JIT 完必须清 ~/.cache/vllm/torch_compile_cache，否则日志写 AITER、图还是 Triton
• 轮子自带 module_gemm_a8w8.so 是 gfx942/950，必须挪走再 JIT

环境

• GPU：MI210 / MI250（gfx90a），104 CU
• ROCm 7.2.x，Python 3.12
• vllm==0.28.0+rocm723（ROCm wheels）
• 新 venv，不要把旧 fork 的源码树塞进 PYTHONPATH

安装

python3.12 -m venv ~/vllm-028-env
source ~/vllm-028-env/bin/activate
pip install vllm==0.28.0+rocm723 --extra-index-url https://wheels.vllm.ai/rocm/0.28.0/rocm723

git clone https://github.com/wu1w/vllm-mi210.git
cd vllm-mi210
pip install -e .

权重（约 31 GB）：

hf download Freaksterz/Qwen3.8-27B-SmoothQuant-W8A8-INT8 \
  --local-dir /path/to/Qwen3.8-27B-SmoothQuant-W8A8-INT8

打补丁 + JIT

先停掉占用 GPU 的 vLLM。

# 1. 闸门：linear / GDN / CU 表 / 挡 FP8 实例
python scripts/apply_gfx90a.py --check   # 应全部 MISSING
python scripts/apply_gfx90a.py
python scripts/apply_gfx90a.py --check   # 应全部 ok / skip

# 2. 挪走官方 gfx942 so，否则 JIT 不会重编
SITE=$(python -c "import aiter, pathlib; print(pathlib.Path(aiter.__file__).parent)")
mv "$SITE/jit/module_gemm_a8w8.so" "$SITE/jit/module_gemm_a8w8.so.gfx942.bak"

# 3. JIT gfx90a CK INT8（约 4 分钟）。so_arch 必须是 ['gfx90a']，误差应为 0
python scripts/jit_gemm_a8w8.py

# 4. 可选：对拍 HIP GDN（Qwen3.8 形状 H=16, HV=48, L=3）
python fused_gdn_hip/build_and_test.py

# 5. 清 AOT，否则仍走 Triton 图
rm -rf ~/.cache/vllm/torch_compile_cache ~/.cache/vllm/torch_aot_compile

apply_gfx90a.py 只改当前 venv 的 site-packages，幂等。升级 vLLM 后要再跑一遍。

启动

export VENV=$HOME/vllm-028-env
export MODEL=/path/to/Qwen3.8-27B-SmoothQuant-W8A8-INT8
# 建议放到 systemd-run / 独立 session，不要塞进内存很小的 cgroup
chmod +x scripts/launch-w8a8.sh
scripts/launch-w8a8.sh

验收日志必须同时有：

• Selected AiterInt8ScaledMMLinearKernel for CompressedTensorsW8A8Int8
• [aiter] import [module_gemm_a8w8]
• GDN decode kernel: cuda（HIP 核挂上后；否则是 triton）

测速：

python scripts/bench.py freak-mtp-n3 Qwen3.8-27B-W8A8

gfx90a 要点

项	事实
INT8 MFMA	v_mfma_i32_16x16x16i8，181 TOPS，没有 CDNA3 那种 2×
FP8	无 FP8 MFMA；torch._scaled_mm 只认 MI300+
AITER 官方闸门	is_aiter_found_and_supported() 要求 CDNA3（get_cdna_version() > 2）
加载	必须 --safetensors-load-strategy eager（mmap 在部分盘上会卡死）
上下文	27B W8A8 + 262K BF16 KV 一张 64 GB 卡能放下；不要开 FP8 KV
GQA	Qwen3.8 GDN 是 16:48 = 3:1

许可证

Apache-2.0。HIP 核算法来自 vLLM CUDA fused_gdn_decode_kernel.cu。AITER 补丁针对 AMD aiter（MIT）在 site-packages 里的闸门，不是 fork。

---

English

Drop-in patches so stock vLLM 0.28.0+rocm723 runs real INT8 W8A8 on MI210 (gfx90a) without switching to a fork.

The ROCm wheel ships gemm_a8w8 for gfx942/950 only, and fused GDN decode is CUDA-only. This repo:

1. Opens the gfx90a AITER linear gate and JITs CK INT8 gemm_a8w8 (I8 only; FP8 instances hang LLVM on CDNA2).
2. Adds a HIP fused GDN decode kernel with runtime GQA (Qwen3.8 is 3:1, not the CUDA 8:1 hardcode).

On 1× MI210 64 GB, Freaksterz Qwen3.8-27B W8A8 + MTP n=3 + 262K: Triton INT8 18.71 tok/s → CK INT8 54.32 tok/s. The HIP GDN kernel matches numerics but does not beat Triton GDN once GEMM is on the CK path.

Do not use --kv-cache-dtype fp8. Clear the vLLM torch AOT cache after switching kernels. Move the wheel module_gemm_a8w8.so aside before JIT or you will keep the gfx942 image.