Красные карты, наконец-то, не сидят на обочине
Долгое время владельцы AMD GPU чувствовали себя пасынками в мире LLM-инференса. Чтобы запустить тяжелую модель, приходилось изгаляться с llama.cpp, собирать его с ROCm или Vulkan, молиться на совместимость, а потом смотреть, как карта греется, а скорость ниже плинтуса. vLLM — мощный серверный фреймворк для пакетного инференса — тоже не баловал красную команду: поддержка квантованных моделей через HIP была, но криво, через общие fallback-пути, без настоящего профита.
Теперь всё меняется. В vLLM смерджили PR с нативным HIP-ядром для формата W4A16 (веса — 4 бита, активации — 16 бит). Это не просто очередной коммит — это заявка на то, что AMD перестает быть запасным вариантом.
Суть: раньше vLLM для W4A16 на AMD использовал универсальные CUDA-эмуляции через HIPIFY — медленно, неэффективно, с неоптимальным использованием регистров и shared memory. Теперь — чистый HIP-код, написанный ручками под архитектуру CDNA/RDNA.
Что дает W4A16 и почему это больно для AMD
Квантование до 4 бит — это стандарт для запуска больших моделей (вроде Mixtral 8x7B или Qwen2.5 72B) на потребительском железе. Модель съеживается в 4 раза без катастрофической потери качества. Но загвоздка в том, что все эти AWQ, GPTQ и прочие методы заточены под NVIDIA: их ядра написаны на CUDA с __half и uint4. AMD долгое время переваривала такие веса только через эмуляцию — медленно, с частыми ошибками сегментации, когда память заканчивалась.
Нативный HIP W4A16 kernel решает проблему в лоб: он умеет эффективно загружать 4-битные веса из VRAM, распаковывать их на лету в регистры и смешивать с FP16-активациями. Без копирования через shared memory, без лишних инструкций. Просто и быстро.
А как же llama.cpp? Он ведь давно умеет
Да, в llama.cpp B9387 добавили ускорение префилла для AMD ROCm, и Vulkan тоже неплохо разгоняется — вон, на Strix Halo Qwen 3.5 35B летает. Но у vLLM другое амплуа: он заточен на одновременную обработку кучи запросов (online serving), динамическое батчирование, управление ключами через PagedAttention. Для single-user десктопа llama.cpp — выбор очевидный. А для сервера, где нужно держать 100+ параллельных чатов с моделью 70B Q4, vLLM с новым ядром просто не оставляет шансов.
Плюс, настройка vLLM для связки двух AMD Radeon 7900 XTX теперь перестала быть танцами с бубном — переменные окружения вроде VLLM_AMD_USE_HIP_GEMM больше не требуют костылей.
Как это выглядит на практике
Всё, что нужно — собрать vLLM с поддержкой ROCm (последнюю версию из ветки main) и указать правильные параметры. Никаких специальных магических флагов.
# Запуск vLLM с W4A16 на AMD
docker run --rm --device=/dev/kfd --device=/dev/dri \
-v /models:/models \
vllm/vllm-openai:latest-rocm \
--model /models/Mixtral-8x7B-Instruct-v0.1-GPTQ \
--dtype half \
--quantization gptq \
--kv-cache-dtype auto \
--max-model-len 8192Важный нюанс: kernel пока работает только для симметричного W4A16 (групповой или per-channel). GPTQ, AWQ и даже простой RTN — все они подходят, если веса упакованы в 4 бита. Для fp8 или других форматов — пока нет, но команда vLLM уже анонсировала работу над W8A16 и гибридными ядрами под AMD.
Осторожно: на картах семейства RDNA 2 (RX 6000) поддержка ROCm ограничена, ядро может не запуститься или работать медленнее из-за отсутствия некоторых инструкций. Проверяйте на RDNA 3 (RX 7000) или AMD Instinct MI250/MI300.
В цифрах: что берем на карандаш
Приведу выжимку из внутренних тестов команды AMD AI (май 2026):
| Модель | GPU | Старый kernel (tok/s) | Новый HIP kernel (tok/s) | Ускорение |
|---|---|---|---|---|
| Qwen2.5 32B Q4 | RX 7900 XTX | 14.3 | 38.1 | 2.66x |
| Mixtral 8x7B Q4 | RX 7900 XTX | 9.7 | 26.2 | 2.70x |
| Mixtral 8x7B Q4 | Instinct MI300X | 68.2 | 205.0 | 3.01x |
Да, цифры не оторваны от реальности — профилировка показала, что старый kernel упирался в bandwidth на декодировании, а новый почти полностью его выбирает.
Кому это нужно прямо сейчас
- Владельцам двух и больше RX 7900 XTX — теперь нет смысла брать NVIDIA ради сервера. Две карты по 24 ГБ + vLLM с W4A16 = Mixtral 8x7B летает, Llama 3 70B влазит.
- Дата-центрам на AMD Instinct — если вы экономите на терах и не хотите переплачивать за NVIDIA, но раньше теряли в скорости — сейчас отрыв сокращается.
- Энтузиастам Synaptics/Minix/X1 Pro — настройка Minisforum AI X1 Pro под llama.cpp Vulkan уже рулит, а с vLLM ROCm можно получить серверную мощность на компактном устройстве.
Впрочем, есть и те, кому это не принесет пользы — например, владельцам eGPU через Thunderbolt. Там узкое место — шина, а не ядро. Как раз недавно eGPU убивает скорость, и новый HIP kernel не спасет.
А что дальше?
Этот PR — не финальная точка, а начало. AMD наконец заимела first-class citizen в vLLM, и теперь разработчики могут писать нативные ядра для квантования, а не эмулировать CUDA. В roadmap — поддержка W8A16 и смешанной точности. Плюс, ходят слухи, что AMD совместно с vLLM адаптирует PagedAttention под архитектуру CDNA3, что даст еще 20-30% к throughput.
Но уже сейчас можно смело утверждать: если у вас AMD и вы держались за NVIDIA только из-за vLLM — аргумент исчез. Скомпилируйте новую версию, подставьте модель и почувствуйте, как красная карта наконец дышит полной грудью.
