根据仓库 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)"
如果是13就对了
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能不能好一些
star 09-16 17:131楼