Почему DeepSeek V4 Flash на RTX 4090 - это прорыв
Запустить DeepSeek V4 Flash на двух потребительских RTX 4090 с пропускной способностью 105 токенов в секунду - реально. Автор портировал эксклюзивные для Blackwell ядра (DeepGEMM, FlashInfer sparse-MLA, block-scaled FP8) на архитектуру Ada (sm89) с помощью Triton. Результат: в 2-3 раза быстрее, чем llama-server, особенно в параллельных агентных сценариях. Суммарный объём VRAM в 48 ГБ позволяет обслуживать контекст до 262K токенов без деградации скорости.
DeepSeek V4 Flash - это MoE-модель с разреженным вниманием и блочно-масштабируемой квантизацией FP8. Её архитектура заточена под тензорные ядра Blackwell (sm100). На Ada (sm89) стандартный vLLM не использует быстрые GEMM-ядра и sparse-MLA, из-за чего скорость падает до 30-40 t/s. Портирование через Triton снимает это ограничение. Сборка выполняется в Docker-образе vLLM-Moet, а для работы двух карт применяется P2P-патч ядра Linux.
Экономика тоже на стороне локального решения. Стресс-тестирование Artificial Analysis показало: DeepSeek V4 Flash стоит $14 за 657 сложных задач, тогда как Opus 4.8 от Anthropic - $1 тыс. Один миллион выходных токенов GLM-5.2 обходится в $1,92 против $50 у Claude Fable 5. Китайские модели в десятки раз дешевле американских аналогов. А локальный инференс на собственных 4090 окупается за 3-4 месяца при нагрузке от 1000 запросов в день.
Архитектурные вызовы: почему стандартный запуск неэффективен
DeepSeek V4 Flash опирается на три компонента, которые в стандартной сборке vLLM работают только на Blackwell. Без них инференс упирается в вычислительную эффективность тензорных ядер и пропускную способность памяти. На RTX 4090 стандартный llama-server выдаёт 35-40 t/s на одном потоке, а при параллельных запросах проваливается до 15-20 t/s. Причина - eager-режим для sparse MLA и отсутствие быстрых путей для block-scaled FP8 на sm89.
Портирование закрывает три узких места: GEMM-операции, разреженное внимание и формат квантизации. Triton позволяет перекомпилировать ядра под целевое оборудование без переписывания на CUDA. Это даёт прирост скорости, сопоставимый с нативной поддержкой Blackwell, но на потребительских картах.
Роль ядер DeepGEMM, FlashInfer sparse-MLA и block-scaled FP8
DeepGEMM - это специализированное ядро матричного умножения для MoE-слоёв. Оно группирует экспертов и выполняет GEMM с минимальными накладными расходами на переключение. На Ada без портирования vLLM использует стандартный torch.matmul, который не учитывает разреженность экспертов и даёт на 40-50% меньше утилизации тензорных ядер.
FlashInfer sparse-MLA реализует разреженное multi-head latent attention. DeepSeek V4 Flash использует MLA вместо стандартного MHA для сокращения KV-кэша. Без аппаратной поддержки sparse-операций vLLM переходит в eager-режим, отключая cuda graphs. Это убивает производительность при батчировании: каждый запрос проходит через полный цикл компиляции графа. Портированное ядро возвращает поддержку cuda graphs и снижает latency на 60%.
Block-scaled FP8 - формат квантизации, где масштабный коэффициент привязан к блоку тензора, а не ко всему слою. Это сохраняет точность, но требует специальной реализации деквантизации. На Ada без портирования vLLM либо падает в FP16 (рост VRAM на 30%), либо использует наивную эмуляцию FP8 с потерей скорости. Портированное ядро выполняет block-scaled FP8 напрямую на тензорных ядрах sm89.
Портирование Blackwell-ядер на Ada с помощью Triton
Triton выбран как инструмент кросс-архитектурной компиляции. Его промежуточное представление MLIR транслируется в PTX для конкретного SM-поколения. Для sm89 это означает генерацию инструкций, использующих тензорные ядра четвёртого поколения и асинхронное копирование памяти (TMA-подобные операции через cp.async).
Процесс портирования затронул три репозитория: DeepGEMM, FlashInfer и vLLM. В DeepGEMM переписан диспетчер экспертов - вместо статического распределения по SM100 используется динамическое с учётом 128 SM на RTX 4090. В FlashInfer sparse-MLA заменён бэкенд с cutlass на Triton с ручной развёрткой циклов для разреженных масок. Block-scaled FP8 реализован через Triton-язык с явным указанием layout для тензорных ядер sm89.
Ключевая особенность реализации - работа с разреженными структурами. DeepSeek V4 Flash использует блочную разреженность 1:4 для весов внимания. На Blackwell это поддерживается аппаратно через sparse-тензорные ядра. На Ada эмуляция через маски в Triton даёт 80% от нативной производительности при правильной настройке tile size.
Сборка Docker-образа vLLM-Moet
Базовый образ - nvidia/cuda:12.4.0-devel-ubuntu22.04. Поверх устанавливаются Triton 3.0, модифицированный FlashInfer и vLLM 0.6.3 с патчами. Сборка занимает около 40 минут на RTX 4090.
FROM nvidia/cuda:12.4.0-devel-ubuntu22.04
RUN pip install triton==3.0.0
RUN git clone https://github.com/vllm-project/vllm.git && cd vllm && git checkout v0.6.3
COPY patches/ /patches/
RUN cd /vllm && git apply /patches/ada-deepgemm.patch
RUN cd /vllm && git apply /patches/ada-flashinfer-mla.patch
RUN pip install -e /vllm
ENV VLLM_MOE_BACKEND=deep_gemm_mega_moe
ENV VLLM_SPARSE_MLA_BACKEND=flashinfer_tritonПатчи интегрируют портированные ядра в vLLM: ada-deepgemm.patch подменяет вызов стандартного GEMM на Triton-ядра, ada-flashinfer-mla.patch добавляет бэкенд для разреженного MLA. Переменные окружения VLLM_MOE_BACKEND и VLLM_SPARSE_MLA_BACKEND активируют портированные пути при запуске.
Для тех, кто уже работает с vLLM на других конфигурациях, будет полезен разбор бэкендов для инференса LLM - там детально сравниваются vLLM, LMDeploy и Triton с TensorRT-LLM на H100 и B200.
Настройка двух RTX 4090: P2P-патч и распределение памяти
RTX 4090 не имеют NVLink и не поддерживают нативный peer-to-peer доступ через PCIe на потребительских драйверах. Без P2P tensor parallelism работает через CPU-память: данные копируются с GPU0 в RAM, затем с RAM на GPU1. Это добавляет 2-3 мс задержки на каждый слой, что снижает общую пропускную способность до 25-30 t/s.
P2P-патч для ядра Linux включает прямой доступ между устройствами через PCIe BAR. Патч модифицирует nvidia-p2p.ko, разрешая DMA между потребительскими картами. После установки vLLM запускается с tensor_parallel_size=2, и обмен данными идёт напрямую через шину PCIe 4.0 x8 (в типичной конфигурации с двумя картами). Пропускная способность PCIe 4.0 x8 - около 16 ГБ/с в каждую сторону, что достаточно для MoE-модели с экспертами, распределёнными между GPU.
Распределение VRAM: 48 ГБ суммарно. Эксперты MoE распределяются поровну - по 12 экспертов на каждую карту. KV-кэш на 262K токенов занимает около 8 ГБ в FP8. Оставшиеся 16 ГБ на карту уходят на веса модели и временные буферы активаций. Без block-scaled FP8 кэш в FP16 занял бы 16 ГБ, что оставило бы только 8 ГБ на веса - недостаточно для полной загрузки модели.
Мониторинг и оптимизация использования VRAM
При контексте 4K токенов потребление VRAM составляет 18 ГБ на карту. С ростом до 262K оно увеличивается до 24 ГБ. Запас в 2-3 ГБ на карту необходим для пиковых аллокаций при батчировании. Рекомендуемая настройка gpu_memory_utilization=0.90 оставляет буфер безопасности.
Сравнение с llama-server: на том же оборудовании llama.cpp при контексте 128K занимает 26 ГБ на карту из-за FP16 KV-кэша и не умещается в 24 ГБ одной карты, требуя offloading на CPU. Это снижает скорость до 5-8 t/s. vLLM с портированными ядрами держит весь кэш в VRAM и не требует offloading до 262K.
Для тех, кто рассматривает альтернативные подходы к инференсу MoE-моделей с ограниченной памятью, рекомендуем материал по стриминг-инференсу через TensorRT-LLM и DeepSpeed - там разбирается запуск полновесных моделей на машинах с недостаточной памятью без квантизации.
Бенчмарки: 105 t/s и сравнение с llama-server
Методика тестирования: генерация текста с длиной ответа 512 токенов, batch size от 1 до 32, контекст 8K токенов. Замеры на прогретой модели после 50 итераций для стабилизации cuda graphs.
| Конфигурация | Batch 1 | Batch 4 | Batch 8 | Batch 16 |
|---|---|---|---|---|
| vLLM-Moet (порт) | 105 t/s | 340 t/s | 580 t/s | 720 t/s |
| vLLM стандартный | 38 t/s | 110 t/s | 180 t/s | 210 t/s |
| llama-server | 42 t/s | 95 t/s | 130 t/s | 140 t/s |
На одиночном запросе портированный vLLM-Moet выдаёт 105 t/s - в 2,5 раза быстрее стандартного vLLM и в 2,5 раза быстрее llama-server. С ростом batch size разрыв увеличивается: на batch 16 портированная версия достигает 720 t/s против 210 у стандартного vLLM и 140 у llama-server. Причина - эффективное батчирование через continuous batching и рабочие cuda graphs для sparse MLA.
Latency первого токена: 45 мс на batch 1, 120 мс на batch 8. Llama-server на batch 8 даёт 380 мс - втрое выше. Для агентных сценариев, где важна реакция на запрос, это критично.
Производительность в агентных сценариях
Агентные рабочие процессы характерны множеством одновременных запросов разной длины. Тест с 32 параллельными сессиями, каждая отправляет запрос каждые 2-5 секунд, контекст от 1K до 32K. vLLM-Moet держит среднюю пропускную способность 650 t/s при p99 latency 180 мс. Llama-server проваливается до 90 t/s с p99 latency 1,2 секунды.
Секрет vLLM - continuous batching и PagedAttention. Запросы не ждут завершения самого длинного в батче, а динамически добавляются и удаляются из обработки. KV-кэш выделяется страницами, что исключает фрагментацию памяти при переменной длине контекста. В агентном сценарии это даёт 3-кратное преимущество над llama-server, который использует статический батчинг.
Интересный контраст - опыт запуска DeepSeek V4 Flash на одном B300, где на пакетной обработке коротких текстов достигли 770 t/s, но столкнулись с отказом MoE-ядра без expert parallel и деградацией спекулятивного декодирования в насыщенном батче. На двух 4090 таких проблем нет: expert parallel работает из коробки при tensor_parallel_size=2.
Экономика локального инференса: окупаемость и сравнение с API
Две RTX 4090 стоят около $3200. При нагрузке 1000 запросов в день со средним ответом 500 токенов это 500 000 выходных токенов в день, 15 млн в месяц. Через API DeepSeek по $0,28 за миллион токенов это $4,2 в месяц. Через Claude API по $15 за миллион - $225 в месяц.
Локальный сервер потребляет около 700 Вт под нагрузкой. При цене $0,12 за кВт·ч это $2 в день, $60 в месяц. Амортизация оборудования за 3 года - $89 в месяц. Итого $149 в месяц против $225 у Claude. Окупаемость против Claude API - 21 месяц. Но если сравнивать с OpenAI o3 по $60 за миллион токенов, локальный сервер окупается за 4 месяца.
Главное преимущество - не цена, а контроль. Нет ограничений по rate limit, нет задержек облака, данные не покидают контур. Для агентных систем с сотнями параллельных сессий это решающий фактор.
Ограничения и подводные камни
Точность FP8 на Ada: тензорные ядра sm89 поддерживают FP8, но с ограничением - нет аппаратной поддержки block-scaled формата. Эмуляция через Triton даёт точность в пределах 0,5% отклонения от эталонного FP16. На задачах классификации и генерации текста разница незаметна. На математических бенчмарках возможна потеря 1-2% accuracy.
Ограничение контекста 262K: полный контекст DeepSeek V4 Flash - 1 млн токенов. На 48 ГБ VRAM помещается только 262K. Для задач с длинными документами (юридический анализ, научные обзоры) этого недостаточно. Портирование не снимает физический лимит памяти.
Зависимость от P2P-патча: патч модифицирует модуль ядра NVIDIA и может сломаться с обновлением драйверов. Рекомендуется зафиксировать версию драйвера 550.54.14 и ядра 6.5. Отсутствие официальной поддержки означает, что при проблемах рассчитывать придётся только на сообщество.
Для понимания альтернативных подходов к инференсу на нестандартном железе полезен разбор запуска GLM-5.2 на 16 AMD MI50 - там достигли 12,2 tok/s через llama.cpp RPC и 14 tok/s через vLLM-gfx906, но столкнулись с деградацией после 10K токенов.
Заключение: ваш собственный сервер DeepSeek V4 Flash
105 токенов в секунду на двух RTX 4090 - это реальность. Портирование Blackwell-ядер через Triton снимает архитектурные ограничения Ada. Docker-образ vLLM-Moet, P2P-патч и правильная конфигурация VRAM - три компонента, которые превращают потребительские карты в сервер инференса уровня A100.
Дальнейшие шаги: собирайте образ по инструкции выше, применяйте P2P-патч, запускайте бенчмарки на своих задачах. Сообщество активно портирует ядра и на другие архитектуры - RTX 3090 (sm86), A100 (sm80). С развитием Triton и открытых инференс-фреймворков разрыв между потребительским и серверным железом продолжит сокращаться.
Если вы только выбираете квантизацию для DeepSeek V4 Flash, обратите внимание на практический тест IQ3_XXS-AS против IQ2_S на RTX 3090 - там показано, что IQ3_XXS-AS даёт оптимальный компромисс между качеством и скоростью, а форк fairydreaming больше не нужен, так как mainline llama.cpp b10064 показывает идентичную производительность.