记录一次vLLM 0.27.1 2080ti x2 抄作业

annacomnena2026-09-05

16823

根据仓库 zyYuc/qwen3-8-27b-dual-2080ti-vllm 编译专属vllm版本，实际上0.24.0之前的几个版本可以直接使用（我记得我在0.19.0，0.20.0版本是可以的），但是之后开始出现sm75不允许使用fp8 kvcache了。这里记录一下我的抄作业方案（专用于 sm75设备启动 qwen3.8 27b），如果大家需要参考请使用自己的具体路径：

编译```
cd llms
git clone https://github.com/zyYuc/qwen3-8-27b-dual-2080ti-vllm.git
cd qwen3-8-27b-dual-2080ti-vllm
source vllm271-sm75/bin/activate
uv pip install -U setuptools wheel packaging ninja
uv pip install torch==2.13.0 \
--index-url https://download.pytorch.org/whl/cu130
``````
git clone https://github.com/vllm-project/vllm.git vllm-0.27.1-sm75
cd vllm-0.27.1-sm75
git checkout v0.27.1
git apply --check \
~/llms/qwen3-8-27b-dual-2080ti-vllm/patches/vllm-v0.27.1-sm75-qwen3.8.patch
``注意此时的位置``
(vllm271-sm75) annacomnena@annacomnena-Default-string:~/llms/qwen3-8-27b-dual-2080ti-vllm/vllm-0.27.1-sm75
``如果，没有报错旧apply patch``
git apply \
~/llms/qwen3-8-27b-dual-2080ti-vllm/patches/vllm-v0.27.1-sm75-qwen3.8.patch
```注意：CUDA-toolkit版本：

首先安装```
sudo apt update
sudo apt install cuda-toolkit-13-0
``然后``
export CUDA_HOME=/usr/local/cuda-13.0
export PATH="$CUDA_HOME/bin:$PATH"
export LD_LIBRARY_PATH="$CUDA_HOME/lib64:${LD_LIBRARY_PATH:-}"
export TORCH_CUDA_ARCH_LIST="7.5"
``此时验证cudatoolkit 与cuda版本：``
which nvcc
nvcc --version
python -c "import torch; print(torch.__version__, torch.version.cuda)"


patch完成之后可以安装依赖：```
uv pip install -e .
```他会自动构建，如果失败了需要清理缓存，基本必定失败因为uv会自己创造临时环境```
rm -rf build .deps CMakeCache.txt CMakeFiles
uv cache clean vllm
```因此uv需要锁定：```
uv pip install \
  "numpy==2.3.5" \
  "setuptools>=77.0.3,<81.0.0" \
  "setuptools-scm>=8.0" \
  "setuptools-rust>=1.9.0" \
  "packaging>=24.2" \
  cmake ninja wheel jinja2
```再构建：```
uv pip install --no-build-isolation -e .
```构建完成后拿FlashQLA的特定版本：（3ab27d77d8ca01d7a4718903b726add1a8886c0e）```
cd ~/llms/qwen3-8-27b-dual-2080ti-vllm

git clone https://github.com/weicj/FlashQLA-SM70-SM75.git
cd FlashQLA-SM70-SM75

git checkout 3ab27d77d8ca01d7a4718903b726add1a8886c0e
```验证patch：```
git apply --check ../patches/flashqla-sm70-sm75-local.patch
```如果没有输出，那么正常，继续```
git apply ../patches/flashqla-sm70-sm75-local.patch
```保持环境变量然后编译：```
export CUDA_HOME=/usr/local/cuda-13.0 
export PATH="$CUDA_HOME/bin:$PATH" 
export LD_LIBRARY_PATH="$CUDA_HOME/lib64:${LD_LIBRARY_PATH:-}" 
export TORCH_CUDA_ARCH_LIST="7.5"
uv pip install --no-build-isolation -e .
```下一步锁定flashinfer版本：```
cd ~/llms/qwen3-8-27b-dual-2080ti-vllm
uv pip install "flashinfer-python==0.6.16.post3"
```应当：

flashinfer-python 0.6.16.post3

torch 2.13.0+cu130

vllm 0.27.1

transformers 5.15.1

triton 3.7.1，检查：```
bash scripts/check_flashinfer_sm75.sh
```True 为通过；

总检察脚本：```
python - <<'PY'
import torch
import vllm
import transformers
import triton
import importlib.metadata as md

print("torch       :", torch.__version__)
print("torch CUDA  :", torch.version.cuda)
print("vllm        :", vllm.__version__)
print("transformers:", transformers.__version__)
print("triton      :", triton.__version__)
print("flashinfer  :", md.version("flashinfer-python"))

from flash_qla.ops.gated_delta_rule.legacy import chunk_gated_delta_rule_fwd_legacy
print("FlashQLA    : legacy import OK")
PY
```启动 3.8 27b```
cd ~/llms

conda deactivate
source qwen3-8-27b-dual-2080ti-vllm/vllm271-sm75/bin/activate

export CUDA_HOME=/usr/local/cuda-13.0
export PATH="$CUDA_HOME/bin:$PATH"
export LD_LIBRARY_PATH="$CUDA_HOME/lib64:${LD_LIBRARY_PATH:-}"
export PYTHONPATH=~/llms/qwen3-8-27b-dual-2080ti-vllm/FlashQLA-SM70-SM75

export VLLM_USE_DEEP_GEMM=0
export VLLM_USE_FLASHINFER_SAMPLER=0
export VLLM_QWOPUS_MTP_BF16_DRAFT=1
export VLLM_SM75_SPEC_SYNC_MODE=safe
export VLLM_USE_V2_MODEL_RUNNER=1

vllm serve ./Qwen3.8-27B-W4A16-AWQ \
  --port 3001 \
  --served-model-name medium \
  --dtype half \
  --tensor-parallel-size 2 \
  --device-ids 0,1 \
  --kv-cache-dtype fp8_e4m3 \
  --max-model-len 200000 \
  --max-num-seqs 3 \
  --max-num-batched-tokens 4096 \
  --enable-prefix-caching \
  --enable-chunked-prefill \
  --no-async-scheduling \
  --language-model-only \
  --skip-mm-profiling \
  --reasoning-parser qwen3 \
  --enable-auto-tool-choice \
  --tool-call-parser qwen3_coder \
  --gpu-memory-utilization 0.92 \
  --additional-config '{"gdn_prefill_backend":"flashqla_legacy"}' \
  --speculative-config '{"method":"mtp","num_speculative_tokens":3,"draft_sample_method":"probabilistic"}' \
  --compilation-config '{"cudagraph_capture_sizes":[4,8,12],"max_cudagraph_capture_size":12}'
```其中mtp默认 greedy :

–speculative-config ‘{“method”:“mtp”,“num_speculative_tokens”:2}’

可改成K=3+probabilistic:

–speculative-config ‘{“method”:“mtp”,“num_speculative_tokens”:3,“draft_sample_method”:“probabilistic”}’

另外的 ：

–kv-offloading-size 6

–kv-offloading-backend native

6是6GB，但是会造成额外的KV搬运性能损耗；有时候能开开，有时候开不开；

也可以替换为lmcache；

可以不开;常常开不开；

另外的， --compilation-config ‘{“cudagraph_capture_sizes”:[4,8,12],“max_cudagraph_capture_size”:12}’

需要开启，以降低3并发的时候可能造成的调度器错误导致推理尾部错误；（不确定是调度器还是哪里的问题，在推理快结束时候常常会只有2-3tps的解码速度）

`VLLM_QWOPUS_MTP_BF16_DRAFT` 的作用是mtp使用bf16，而不是一样的把他当成awq量化；0.28.0中官方已修复

但是最终测下来速度和lmdeploy直接启动，是一模一样的，单流下预填充913tps/解码36tps； 只是lmdepoly目前很尴尬，turbomind后端支持awq量化，不支持mtp；pytorch后端支持mtp，但是只支持bf16权重；

实际上做的比较好的1猫vllm，拿了turbomind嫁接到vllm上，效果是不错的，不过他那个仓库是给sm70设备专用的（看来v100要涨价了）

Ornith 1.5```
vllm serve ./Ornith-1.5-35B-A3B-AWQ-INT4 \
  --port 3001 \
  --served-model-name medium \
  --dtype half \
  --tensor-parallel-size 2 \
  --device-ids 0,1 \
  --kv-cache-dtype fp8_e4m3 \
  --max-model-len 200000 \
  --max-num-seqs 7 \
  --max-num-batched-tokens 4096 \
  --enable-prefix-caching \
  --enable-chunked-prefill \
  --no-async-scheduling \
  --language-model-only \
  --skip-mm-profiling \
  --reasoning-parser qwen3 \
  --enable-auto-tool-choice \
  --tool-call-parser qwen3_coder \
  --gpu-memory-utilization 0.92 \
  --additional-config '{"gdn_prefill_backend":"flashqla_legacy"}' \
```ulkaa 这个量化 MTP head是原模型**逐字节搬过来的完整 BF16 head**

但是开启了`VLLM_QWOPUS_MTP_BF16_DRAFT` 也不能正确使用MTP；因为它只能处理dense模型的dense mtp；

对于qwen3_5 moe,不过ornith做的本来也不好用；

shisa-ai/Ornith-1.5-35B-A3B-MTP-ONLY 做了一个ornith专用的 MTP模型；

替换为更好的mtp模型

下载 shisa-ai/Ornith-1.5-35B-A3B-MTP-ONLY；

替换权重（注意备份原mtp权重model-mtp.safetensors和model.safetensors.index.json)之后,

修改model.safetensors.index.json：

（使用vscode）查找和替换，启用正则：```
^\s*"mtp\.[^"]+"\s*:\s*"[^"]+",?\s*\r?\n
```全部替换，删除原有的785行，加入：```
    "mtp.layers.0.self_attn.q_proj.weight": "model-mtp.safetensors",
    "mtp.layers.0.self_attn.k_proj.weight": "model-mtp.safetensors",
    "mtp.layers.0.self_attn.v_proj.weight": "model-mtp.safetensors",
    "mtp.layers.0.self_attn.o_proj.weight": "model-mtp.safetensors",
    "mtp.layers.0.self_attn.q_norm.weight": "model-mtp.safetensors",
    "mtp.layers.0.self_attn.k_norm.weight": "model-mtp.safetensors",
    "mtp.layers.0.input_layernorm.weight": "model-mtp.safetensors",
    "mtp.layers.0.mlp.gate.weight": "model-mtp.safetensors",
    "mtp.layers.0.mlp.shared_expert.gate_proj.weight": "model-mtp.safetensors",
    "mtp.layers.0.mlp.shared_expert.up_proj.weight": "model-mtp.safetensors",
    "mtp.layers.0.mlp.shared_expert.down_proj.weight": "model-mtp.safetensors",
    "mtp.layers.0.mlp.shared_expert_gate.weight": "model-mtp.safetensors",
    "mtp.layers.0.mlp.experts.gate_up_proj": "model-mtp.safetensors",
    "mtp.layers.0.mlp.experts.down_proj": "model-mtp.safetensors",
    "mtp.layers.0.post_attention_layernorm.weight": "model-mtp.safetensors",
    "mtp.pre_fc_norm_embedding.weight": "model-mtp.safetensors",
    "mtp.pre_fc_norm_hidden.weight": "model-mtp.safetensors",
    "mtp.fc.weight": "model-mtp.safetensors",
    "mtp.norm.weight": "model-mtp.safetensors",
```修复moe模型的moe mtp

上游新版 `qwen3_5_mtp.py` 的思路是：如果 MTP 被排除在 target 的量化之外，就在创建整个 MTP decoder layer 时临时把 `vllm_config.quant_config` 设为 `None`，创建完再恢复。

把现有的 `VLLM_QWOPUS_MTP_BF16_DRAFT=1` 扩展成：不仅 `mtp.fc` 用 BF16，整个 `mtp.layers.0` 都用 BF16。

打开：```
/home/annacomnena/llms/qwen3-8-27b-dual-2080ti-vllm/vllm-0.27.1-sm75/ vllm/model_executor/models/qwen3_5_mtp.py
```补丁：```
--- a/qwen3_5_mtp.py
+++ b/qwen3_5_mtp.py
@@ -119,25 +119,45 @@
             prefix=f"{prefix}.fc",
         )
 
-        # GPTQ: quantized checkpoints may exclude MTP from quantization via
-        # quantization_config.dynamic with "-:pattern" entries. When detected,
-        # disable quantization for MTP layers so they use unquantized params.
+        # Quantized checkpoints may exclude MTP from quantization. In addition
+        # to the upstream dynamic-exclude handling below, QWOPUS checkpoints can
+        # keep the *entire* MTP draft layer (including routed MoE experts) in
+        # BF16 while the target model uses compressed-tensors/AWQ. In that case
+        # the MTP decoder layer must be constructed with quantization disabled;
+        # otherwise RoutedExperts is built with packed/Marlin parameters and the
+        # BF16 fused MTP tensors (w13_weight / w2_weight) cannot be loaded.
         original_quant = vllm_config.quant_config
-        if quant_config and quant_config.get_name() not in ("modelopt_fp4",):
+        force_unquantized_mtp_layer = qwopus_bf16_mtp and quant_name != "fp8"
+
+        if force_unquantized_mtp_layer:
+            logger.info(
+                "VLLM_QWOPUS_MTP_BF16_DRAFT=1: building the entire "
+                "Qwen3.5 MTP decoder layer unquantized; target quantization=%s",
+                quant_name or "disabled",
+            )
+            vllm_config.quant_config = None
+        elif quant_config and quant_config.get_name() not in ("modelopt_fp4",):
+            # GPTQ / other quantized checkpoints may exclude MTP via
+            # quantization_config.dynamic with "-:pattern" entries.
             hf_qc = getattr(model_config.hf_config, "quantization_config", None)
             if isinstance(hf_qc, dict):
                 dynamic = hf_qc.get("dynamic", {})
                 if any(k.startswith("-:") and "mtp" in k for k in dynamic):
                     vllm_config.quant_config = None
-        self.layers = torch.nn.ModuleList(
-            Qwen3_5DecoderLayer(
-                vllm_config,
-                layer_type="full_attention",
-                prefix=f"{prefix}.layers.{idx}",
+
+        try:
+            self.layers = torch.nn.ModuleList(
+                Qwen3_5DecoderLayer(
+                    vllm_config,
+                    layer_type="full_attention",
+                    prefix=f"{prefix}.layers.{idx}",
+                )
+                for idx in range(self.num_mtp_layers)
             )
-            for idx in range(self.num_mtp_layers)
-        )
-        vllm_config.quant_config = original_quant
+        finally:
+            # Do not leak the draft-layer override into the target model or
+            # later modules such as lm_head.
+            vllm_config.quant_config = original_quant
         self.make_empty_intermediate_tensors = make_empty_intermediate_tensors_factory(
             ["hidden_states", "residual"], config.hidden_size
         )
```合适的话会看到：```
(Worker_TP1 pid=2168845) INFO 09-04 15:46:45 [qwen3_5_mtp.py:106] VLLM_QWOPUS_MTP_BF16_DRAFT=1: loading Qwen3.5 MTP draft fc without the target quantization config when needed; draft layer quantization=compressed-tensors
(Worker_TP1 pid=2168845) INFO 09-04 15:46:45 [qwen3_5_mtp.py:133] VLLM_QWOPUS_MTP_BF16_DRAFT=1: building the entire Qwen3.5 MTP decoder layer unquantized; target quantization=compressed-tensors
```如果MTP开2，并发开6，最终CUDA图捕捉size要重新计算，1生成+2草稿=3，6并发应当是：```
--compilation-config '{"cudagraph_capture_sizes":[3,6,9,12,15,18],"max_cudagraph_capture_size":18}'
```并且max-num-seqs开7给调度器留一个余量；

如果MTP=2，那么：```
--max-num-seqs 7 \
--compilation-config '{"cudagraph_capture_sizes":[4,8,12,16,20,24,28],"max_cudagraph_capture_size":28}'
```,“draft_tensor_parallel_size”:1

可以不TP2，减小卡间通信成本,但是可能起不来

完整启动命令：```
vllm serve ./Ornith-1.5-35B-A3B-AWQ-INT4 \
  --port 3001 \
  --served-model-name medium \
  --dtype half \
  --tensor-parallel-size 2 \
  --device-ids 0,1 \
  --kv-cache-dtype fp8_e4m3 \
  --max-model-len 200000 \
  --max-num-seqs 7 \
  --max-num-batched-tokens 4096 \
  --enable-prefix-caching \
  --enable-chunked-prefill \
  --no-async-scheduling \
  --language-model-only \
  --skip-mm-profiling \
  --reasoning-parser qwen3 \
  --enable-auto-tool-choice \
  --tool-call-parser qwen3_coder \
  --gpu-memory-utilization 0.92 \
  --additional-config '{"gdn_prefill_backend":"flashqla_legacy"}' \
  --speculative-config '{"method":"mtp","num_speculative_tokens":1,"draft_sample_method":"probabilistic"}' \
  --compilation-config '{"cudagraph_capture_sizes":[2,4,6,8,10,12],"max_cudagraph_capture_size":12}'
```其实我最终测下来可能mtp开1就够了，多了也没啥用，第二节接受度很低，所以此时把cuda图编译的设置做相应的更改；

由于没有nvlink，最终decode 单流也只有50-60tps，不知道上了nvlink能不能好一些

主楼

发表回复