企业官网建设流程全解析

1. 这不是 benchmark 跑分，而是一次真实服务场景下的工程压测实录

最近在 Baseten 内部做模型服务架构迭代时，我们把一套刚上线的 LLM 推理引擎和社区广泛采用的 vLLM 做了一次横向对比——不是在 synthetic load 下跑吞吐、延迟曲线图，而是直接用生产环境的真实请求流：包含长上下文（平均 4200 token）、混合 batch（1–8 并发）、动态 prompt length（从 300 到 6800 token 不等）、带 streaming 输出的 API 调用。结果很意外：新推理引擎在 P95 延迟上比 vLLM 低 67%，在高并发（128 RPS）下吞吐提升 2.3 倍，最极端 case（长 context + max batch size）甚至快了 89.7%。这个“最多快 90%”不是修辞，是我们在三台 A100-80G 机器上连续 72 小时压测后取的实测极值。很多人第一反应是“是不是 trick？”——比如关了 KV cache、用了更激进的量化、或者只测了小模型？都不是。我们测的是同一套权重（Qwen2.5-7B-Instruct，AWQ 4-bit）、同一 CUDA 版本（12.4）、同一 Triton kernel 编译配置、同一 tokenizer（transformers 4.41），唯一变量就是后端推理 runtime。背后没有 magic，只有三个被 vLLM 默认忽略、但在真实业务中高频出现的瓶颈点：内存访问局部性断裂、prefill 阶段的 attention kernel 吞吐墙、以及 dynamic batching 下的 GPU 显存碎片化放大效应。这篇文章不讲理论推导，只讲我们怎么定位、怎么改、怎么验证——包括每一行关键 patch 的作用、为什么不能简单复用 HuggingFace Transformers 的原生 forward、以及为什么你在本地用 ollama 或 LM Studio 跑不出这个差距（它们根本没触发这些瓶颈）。如果你正在用 vLLM 部署 Qwen、DeepSeek、Llama3 等主流开源模型，且遇到“明明 GPU 利用率不到 60% 但延迟却飙升”的情况，这篇就是为你写的。它适合两类人：一是已经跑通 vLLM 但卡在性能天花板的工程师；二是正评估推理框架选型、想避开“文档写得漂亮、线上跑得心累”陷阱的技术负责人。

2. 为什么 vLLM 在真实负载下会“慢得合理”？——从设计哲学到工程现实的断层

2.1 vLLM 的核心优势与隐含假设

vLLM 成为事实标准，靠的是 PagedAttention 这一开创性设计。它把 KV cache 按 block 切片、用虚拟内存式管理，解决了传统自回归生成中显存浪费严重的问题。这在论文 benchmark（如 ShareGPT 数据集、固定 batch=1/2/4、prompt length ≤ 2048）下效果惊艳——吞吐翻倍、显存下降 40%。但它的成功建立在几个强假设上，而这些假设在真实 API 服务中往往不成立：

• 假设 1：prefill 和 decode 阶段计算负载均衡
  vLLM 默认将 prefill（一次性计算所有 prompt token 的 KV）和 decode（逐 token 生成）视为两个独立阶段，并复用同一 attention kernel。但在长 prompt 场景（如 4K+ token 的法律合同分析），prefill 占据整个请求 70% 以上耗时，而 decode 只占 30%。此时 GPU 计算单元大量空转等待 memory bandwidth，vLLM 却无法针对性优化 prefill 的 kernel launch granularity。

• 假设 2：batch 内 prompt length 高度同质
  PagedAttention 的 block 分配效率高度依赖 batch 内各 sequence 的长度接近。一旦混入一个 500-token query 和一个 5000-token query，短 query 的 block 会被长 query “拖垮”，导致大量 block 处于半填充状态，实际显存利用率反而比 naive KV cache 低 15–20%。我们线上流量中，length variance σ > 1200 是常态。

• 假设 3：GPU 显存带宽是瓶颈，而非 kernel launch overhead
  vLLM 重度依赖 CUDA Graph 捕获来降低 kernel launch 开销。但它默认只对 decode 阶段做 graph capture，prefill 阶段仍走动态 dispatch。而在混合 batch 场景下，每次 prefill 都要重新计算 block table、dispatch 不同 shape 的 matmul，launch overhead 占 prefill 总耗时 18–25%（A100 测得）。这个开销在 synthetic benchmark 中被平均掉，但在真实 burst 请求中直接暴露。

提示：这不是 vLLM 的缺陷，而是其设计目标明确——最大化 throughput/GB 显存，而非最小化 P95 latency。当你用 vLLM 跑离线批处理（batch=64, fixed length），它依然是王者；但当你用它扛在线 API（batch=8, variable length, streaming），它就开始“合理地慢”。

2.2 Baseten 新推理引擎的破局思路：放弃通用性，专注服务链路

我们没重写 attention，也没发明新调度算法。而是把整条推理 pipeline 拆成四段，每一段都针对真实服务特征做定制：

1. Request Ingestion Layer：不等完整 HTTP body 到达就启动 tokenization，用 streaming tokenizer 预解析前 512 token，提前预分配 block；
2. Prefill Optimizer：为不同 length range（<1K / 1–3K / >3K）编译专用 Triton kernel，避免通用 kernel 的 branch divergence；
3. Dynamic Block Manager：放弃固定 block size（vLLM 默认 16），改为按 sequence length 动态选择 block size（8/16/32），并引入“block borrowing”机制——当长 sequence block 不足时，临时借用相邻短 sequence 的空闲 block；
4. Streaming Emitter：vLLM 的 output queue 是 per-request FIFO，我们改成 global priority queue，按 token arrival time 排序，确保高优先级请求（如付费用户）的首 token 延迟不受低优先级请求影响。

这四个模块加起来，代码量只有 vLLM 的 1/3，但每个都直击线上痛点。比如“block borrowing”机制，它让显存碎片率从 vLLM 的 34% 降到 9%，这意味着同样 80G 显存，vLLM 最多跑 42 个 4K-context 请求，而我们能跑 58 个——多出的 16 个并发，直接转化为吞吐提升。

2.3 关键决策背后的工程权衡：为什么不用 FlashAttention-3？

FlashAttention-3 确实更快，但它要求 CUDA 12.2+ 且仅支持 Hopper 架构（H100）。我们线上主力卡是 A100（Ampere），占比 78%。强行升级意味着：

• 所有模型需重训/重量化（FA3 的 kernel 对 weight layout 有强约束）；
• 现有 Triton custom op 全部失效（FA3 不开放 kernel source）；
• CI/CD 流程重构（需维护两套 CUDA 版本分支）。

我们测算过：在 A100 上，FA3 相比 FA2 的收益约 12%，但上述迁移成本会让团队至少损失 3 周交付周期。而通过定制 prefill kernel + dynamic block，我们在 A100 上拿到了 23% 的 prefill 加速——性价比更高。技术选型不是比谁用的库新，而是比谁更懂自己的硬件栈和交付节奏。

3. 实操拆解：如何复现这个 90% 的加速？——从环境准备到压测验证

3.1 环境搭建：最小可行对比组（非 docker，纯裸机）

我们拒绝用 docker 镜像做对比，因为容器网络栈、cgroup 限制、NVIDIA Container Toolkit 的 driver shim 都会引入不可控 variance。所有测试均在物理机上完成：

• 硬件：Dell R760，双路 AMD EPYC 7763，4×A100-80G（PCIe 4.0 x16），Ubuntu 22.04.4 LTS
• CUDA：12.4.0（必须！vLLM 0.4.2+ 要求 ≥12.1，但 12.4 在 A100 上比 12.1 稳定 17%）
• Driver：535.129.03（NVIDIA 官方推荐用于 A100 + CUDA 12.4）
• Python：3.10.12（vLLM 不支持 3.11+，Baseten 引擎暂未适配 3.12）

安装命令严格按顺序执行（顺序错会导致 CUDA context 冲突）：

# 1. 清理旧环境 sudo apt-get remove --purge nvidia-* && sudo apt autoremove sudo reboot # 2. 安装驱动（必须先装驱动，再装 CUDA） wget https://download.nvidia.com/XFree86/Linux-x86_64/535.129.03/NVIDIA-Linux-x86_64-535.129.03.run sudo ./NVIDIA-Linux-x86_64-535.129.03.run --no-opengl-files --no-nouveau-check # 3. 安装 CUDA 12.4（注意：不要选 driver！只选 runtime 和 toolkit） wget https://developer.download.nvidia.com/compute/cuda/12.4.0/local_installers/cuda_12.4.0_535.54.03_linux.run sudo sh cuda_12.4.0_535.54.03_linux.run --silent --override --toolkit --samples=false --no-opengl-libs # 4. 设置环境变量（写入 ~/.bashrc） export CUDA_HOME=/usr/local/cuda-12.4 export PATH=$CUDA_HOME/bin:$PATH export LD_LIBRARY_PATH=$CUDA_HOME/lib64:$LD_LIBRARY_PATH source ~/.bashrc # 5. 验证 nvidia-smi # 应显示 Driver Version: 535.129.03, CUDA Version: 12.4 nvcc --version # 应显示 release 12.4, V12.4.127

注意：如果nvcc --version显示 12.1 或更低，说明你装了旧版 CUDA。必须彻底卸载/usr/local/cuda-*下所有目录，再重装。我们踩过坑：某次 CI 机器残留 CUDA 11.8，导致 vLLM 编译时链接错误，但 error message 只提示 "undefined symbol"，排查耗时 6 小时。

3.2 模型准备：统一权重，杜绝“模型差异”干扰

我们用 HuggingFace 官方 Qwen2.5-7B-Instruct（commit:a3f2b7d），量化方式为 AWQ（group_size=128, w_bit=4, version="gemm"）：

# 使用 awq-pytorch 量化（vLLM 和 Baseten 引擎均支持 AWQ） pip install awq-pytorch==0.1.6 python -c " from awq import AutoAWQForCausalLM from transformers import AutoTokenizer model = AutoAWQForCausalLM.from_pretrained( 'Qwen/Qwen2.5-7B-Instruct', safetensors=True, device_map='cpu' ) tokenizer = AutoTokenizer.from_pretrained('Qwen/Qwen2.5-7B-Instruct') model.quantize(tokenizer, quant_config={'zero_point': True, 'q_group_size': 128, 'w_bit': 4}) model.save_quantized('./qwen2.5-7b-awq') tokenizer.save_pretrained('./qwen2.5-7b-awq') "

关键点：

• 不使用 auto-gptq：GPTQ 量化在 vLLM 中需额外加载exllama2kernel，而 Baseten 引擎只支持 AWQ，为公平起见，全部用 AWQ；
• safetensors=True：避免 PyTorch 的 pickle 安全风险，且加载速度比 bin 快 22%；
• device_map='cpu'：防止量化时 GPU 显存溢出（7B 模型量化需 24G VRAM）。

量化后模型大小为 3.82 GB（原始 FP16 为 13.2 GB），这是后续所有测试的基础镜像。

3.3 vLLM 部署：标准配置，但启用所有优化项

vLLM 版本锁定为0.4.2（最新版0.4.3在 A100 上有 memory leak，已向官方提 issue）：

pip install vllm==0.4.2 --no-cache-dir

启动命令（关键参数已加注释）：

python -m vllm.entrypoints.api_server \ --model ./qwen2.5-7b-awq \ --tensor-parallel-size 4 \ # 四卡并行，每卡 load 1/4 权重 --dtype half \ # 必须用 half，AWQ kernel 依赖 FP16 input --gpu-memory-utilization 0.9 \ # 显存利用率设为 0.9，vLLM 默认 0.9，不调 --max-num-seqs 256 \ # 最大并发请求数，对应 max batch size --max-model-len 8192 \ # 支持最长 context，必须 ≥ 8K 才能测长 prompt --enable-prefix-caching \ # 启用 prefix cache，对重复 prompt 有效 --disable-async-output-proc \ # 关闭异步输出处理，避免 streaming 延迟抖动 --port 8000

实操心得：--max-num-seqs是 vLLM 的隐形瓶颈。它不是最大并发数，而是 scheduler 维护的 pending request queue 长度。如果设太小（如默认 256），burst 请求会排队，P95 延迟虚高。我们线上设为 512，但测试时保持 256 以对标 baseline。

3.4 Baseten 引擎部署：轻量级启动，无依赖污染

Baseten 引擎以 wheel 包形式提供（内部构建，不公开源码），安装即用：

pip install baseten-inference-engine-0.1.0-py3-none-any.whl --no-deps # 它不依赖 torch/vllm，只依赖 numpy、triton、nvidia-cublas-cu12

启动命令极简：

baseten-server \ --model-path ./qwen2.5-7b-awq \ --gpus 0,1,2,3 \ # 指定 GPU ID，非数量 --max-batch-size 64 \ # 实际最大 batch，比 vLLM 的 max-num-seqs 更直观 --context-length 8192 \ # 同 vLLM --quant-type awq \ # 明确指定量化类型 --port 8001

区别在于：

• 无 tensor parallel 参数：引擎自动检测 GPU 数量并做最优切分（A100 四卡用 model parallel + data parallel 混合）；
• --max-batch-size 是硬上限：超过则直接 reject，不排队，保证 P95 可预测；
• 内置 health check endpoint：GET /health返回实时 GPU utilization、pending queue length、avg token/s。

3.5 压测工具：locust + 自定义 payload generator

我们不用 ab 或 wrk，因为它们无法模拟 streaming response 和 variable-length prompt。改用 locust，编写 custom task：

# locustfile.py from locust import HttpUser, task, between import random import json class LLMUser(HttpUser): wait_time = between(0.1, 0.5) # 模拟真实用户间隔 @task def generate(self): # 从真实日志抽样 1000 条 prompt，按 length 分组 prompts = [ ("简述量子纠缠原理", 32), ("请分析这份 20 页 PDF 的法律风险，重点标注第 7–12 页...", 4217), # ... 共 1000 条，length 分布符合线上 σ=1280 ] prompt, length = random.choice(prompts) payload = { "prompt": prompt, "max_tokens": 512, "stream": True, # 必须开启 streaming "temperature": 0.7 } with self.client.post("/generate", json=payload, catch_response=True) as resp: if resp.status_code != 200: resp.failure(f"HTTP {resp.status_code}") else: # 解析 streaming response，计算首 token delay 和 end-to-end delay tokens = [] for line in resp.iter_lines(): if line.startswith(b"data:"): try: data = json.loads(line[6:]) if "token" in data: tokens.append(data["token"]) except: pass if len(tokens) < 10: resp.failure("Too few tokens generated")

运行命令：

locust -f locustfile.py --headless -u 128 -r 10 -t 30m --host http://localhost:8000 # vLLM locust -f locustfile.py --headless -u 128 -r 10 -t 30m --host http://localhost:8001 # Baseten

• -u 128：目标并发用户数（即 RPS ≈ 128）；
• -r 10：每秒启动 10 个用户，模拟 ramp-up；
• -t 30m：持续 30 分钟，跳过 warm-up 阶段，直接测稳态。

3.6 关键指标采集：不止看 avg，更要盯 P95/P99

vLLM 和 Baseten 都暴露/metricsendpoint（Prometheus format），我们用 telegraf 抓取：

# telegraf.conf [[inputs.prometheus]] urls = ["http://localhost:8000/metrics", "http://localhost:8001/metrics"] metric_version = 2 [inputs.prometheus.tags] service = "vllm"

重点关注三个指标：

实测数据：在 128 RPS 下，vLLM 的 P95 end-to-end delay 为 3.21s，Baseten 为 1.03s，加速比 2.14×（即快 114%）。但标题说“最多快 90%”，是因为我们取的是相同 P95 延迟下的吞吐对比：当两者 P95 都控制在 2.0s 时，vLLM 最大 RPS 为 82，Baseten 为 154，提升 87.8% —— 四舍五入即“最多快 90%”。这是更合理的业务视角：你愿意为更低延迟多花多少硬件？还是为更高吞吐接受稍高延迟？

4. 核心环节深度解析：prefill 加速的 Triton kernel 如何写？

4.1 为什么通用 FlashAttention 在 prefill 下失效？

FlashAttention-2 的 kernel 是为 decode 阶段设计的：它假设 Q/K/V shape 为[B, H, T, D]，其中 T（sequence length）很小（通常 ≤ 128）。但在 prefill 阶段，T 可达 4096+，此时 kernel 的 shared memory usage 超出 SM limit（A100 为 164KB），导致大量 register spilling，性能暴跌。我们用 Nsight Compute 分析发现：prefill 时 kernel occupancy 从 decode 的 82% 降到 31%，SM 利用率不足 40%。

解决方案不是换 kernel，而是分治：把长 sequence 拆成多个子块，每个子块用 optimized small-T kernel 处理，再 merge result。

4.2 我们的 Triton kernel 设计（伪代码级说明）

核心思想：对 Q/K/V 做 block-wise attention，每个 block 大小为BLOCK_M=64, BLOCK_N=64（适配 A100 的 warp size），并利用tl.dot的 async load 机制隐藏 global memory latency。

@triton.jit def _fwd_kernel( Q, K, V, sm_scale, # 输入指针 L, M, # softmax 归一化中间变量 Out, # 输出 stride_qz, stride_qh, stride_qm, stride_qk, # Q 的 stride stride_kz, stride_kh, stride_kn, stride_kk, stride_vz, stride_vh, stride_vn, stride_vk, Z, H, N_CTX, # batch, head, seq_len BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, # 64x64 BLOCK_DMODEL: tl.constexpr, # head_dim=128 ): start_m = tl.program_id(0) # 计算当前 block 的 m 范围 [start_m*BLOCK_M, min((start_m+1)*BLOCK_M, N_CTX)] offs_m = start_m * BLOCK_M + tl.arange(0, BLOCK_M) offs_n = tl.arange(0, BLOCK_N) # 预加载 Q[offs_m, :] 到 shared memory（关键！减少 global load） q = tl.load(Q + offs_m[:, None] * stride_qm + offs_n[None, :] * stride_qk, mask=(offs_m[:, None] < N_CTX) & (offs_n[None, :] < BLOCK_DMODEL), other=0.0) # 对每个 m，计算 K[0:N_CTX, :] @ Q[m, :]^T → score[m, 0:N_CTX] # 但 N_CTX 太大，所以分段：每次 load K 的 BLOCK_N 行 lo = 0 hi = N_CTX for start_n in range(lo, hi, BLOCK_N): # load K[start_n:start_n+BLOCK_N, :] k = tl.load(K + (start_n + offs_n[:, None]) * stride_kn + offs_m[None, :] * stride_kk, mask=(start_n + offs_n[:, None] < N_CTX) & (offs_m[None, :] < BLOCK_DMODEL), other=0.0) # compute qk qk = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32) qk += tl.dot(q, k, trans_b=True) qk *= sm_scale # 归一化：减去 max（避免 exp overflow） m_i = tl.maximum(tl.max(qk, 1), lse_i) p = tl.exp(qk - m_i[:, None]) lse_i = tl.log(tl.sum(p, 1)) + m_i # load V[start_n:start_n+BLOCK_N, :] v = tl.load(V + (start_n + offs_n[:, None]) * stride_vn + offs_m[None, :] * stride_vk, mask=(start_n + offs_n[:, None] < N_CTX) & (offs_m[None, :] < BLOCK_DMODEL), other=0.0) # accumulate acc += tl.dot(p, v) # store output tl.store(Out + offs_m[:, None] * stride_qm + offs_n[None, :] * stride_qk, acc, mask=(offs_m[:, None] < N_CTX) & (offs_n[None, :] < BLOCK_DMODEL))

这个 kernel 的关键创新点：

• BLOCK_M=64：确保每个 warp 处理 64 行 Q，完美匹配 A100 的 32-warp SM；
• async load：tl.load的 mask 机制让无效地址不触发 fault，比ifbranch 更高效；
• shared memory reuse：Q 被预加载一次，反复用于与不同 K block 计算，减少 global memory traffic。

实测：在 4096-length prompt 下，该 kernel 比 FA2 快 3.2×，且 occupancy 保持在 79%。

4.3 如何集成到推理引擎？——不改模型，只换 backend

我们没动 transformers 的forward()，而是用torch.compile的 backend hook：

# 替换 Qwen2Model.forward 中的 attention call def patched_attn_forward(self, hidden_states, *args, **kwargs): if self.training: return original_attn(hidden_states, *args, **kwargs) else: # 提取 QKV q, k, v = self.q_proj(hidden_states), self.k_proj(hidden_states), self.v_proj(hidden_states) # 调用自研 Triton kernel out = triton_prefill_attention(q, k, v, self.scaling_factor) return self.o_proj(out) # 注入 hook for layer in model.layers: layer.self_attn.forward = MethodType(patched_attn_forward, layer.self_attn)

这样做的好处：

• 零模型修改：Qwen2、Llama3、DeepSeek 都适用，只需替换self_attn.forward；
• 无缝 fallback：当输入 length < 512 时，自动切回 FA2（短序列 FA2 更优）；
• 可热替换：无需重启服务，torch.compile会自动 recompile。

注意：torch.compile在 A100 上有时会生成 suboptimal kernel，我们强制指定 backend：torch._dynamo.config.cache_size_limit = 128，并禁用inductor的某些 fusion pass（具体 patch 见 internal doc #TRITON-2024）。

5. 常见问题与避坑指南：那些文档不会告诉你的细节

5.1 问题速查表：为什么我的 vLLM 比别人慢？

5.2 实操避坑：三个血泪教训

坑 1：CUDA Graph 在 vLLM 中的“伪优化”
vLLM 文档说“启用 CUDA Graph 可提升 20%”，但实测发现：它只对 decode 有效，且必须满足--enforce-eager（禁用 graph）才能稳定。我们曾在线上开启 graph，结果 burst 请求时 graph capture 失败，fallback 到 eager mode，延迟 spikes 达 300ms。结论：除非你 100% 确保 batch size 和 length 固定，否则关掉--enable-graph。

坑 2：tokenizer 的 decode 速度拖累整体
vLLM 默认用transformers的 tokenizer，但tokenizer.decode()在长 output 下极慢（O(n²)）。我们改用tokenizers库的 fast tokenizer，并缓存 decode 结果（LRU cache size=1024），首 token delay 降低 140ms。代码：

from tokenizers import Tokenizer fast_tokenizer = Tokenizer.from_file("./qwen2.5-7b-awq/tokenizer.json") # 替换 vLLM 的 tokenizer.decode

坑 3：Baseten 引擎的“冷启动”陷阱
Baseten 引擎首次请求会触发 Triton kernel 编译（JIT），耗时 8–12s。如果用 k8s rolling update，新 pod 的第一个请求必然超时。解决方案：启动后立即发一个 warmup request：

curl -X POST http://localhost:8001/generate \ -H "Content-Type: application/json" \ -d '{"prompt":"Hello","max_tokens":1,"stream":false}'

这个请求不返回 token，只触发 kernel compile，耗时 10.2s，之后所有请求稳定在 sub-100ms。

5.3 性能调优 checklist（A100 用户专属）

• [ ] 确认nvidia-smi显示P0power state（非P2），否则 GPU 降频；
• [ ]cat /sys/class/nvme/nvme0/nvme0n1/device/power_state应为D0（NVMe SSD 不休眠）；
• [ ] 关闭irqbalance服务，将 GPU IRQ 绑定到特定 CPU core（sudo irqtop查看，然后sudo echo 0-3 > /proc/irq/*/smp_affinity_list）；
• [ ] 在/etc/default/grub中添加intel_idle.max_cstate=1 rcu_nocbs=0-64，重启生效；
• [ ] 使用numactl --cpunodebind=0 --membind=0 python ...启动服务，避免跨 NUMA 访问。

这些调优项合计带来 11.3% 的 P95 降低。别小看这 11%，它让你在同等硬件下多承载 12% 的流量。

6. 这个加速能迁移到其他框架吗？——现实边界与扩展思考

6.1 对 ollama / LM Studio 用户的意义

ollama 和 LM Studio 本质是 vLLM 的封装层，它们的性能瓶颈完全继承自 vLLM。所以如果你用 ollama run qwen2.5:7b，底层仍是 vLLM 的 prefill kernel。本文的加速方案无法直接用于 ollama，但你可以：

• 在 ollama 的Modelfile中指定FROM ./qwen2.5-7b-awq，然后手动替换其内置的 vLLM 为 patch 版（需编译）；
• 或改用llama.cpp+gguf格式，它对 prefill 有类似优化（但仅限 llama 系列）。

LM Studio 同理，它用的是llama.cppbackend，所以长 prompt 加速已内置，但 multi-GPU 支持弱——这也是为什么我们坚持用 A100×4 而非单卡 H100。

6.2 对 DeepSeek-V2 / Qwen3 部署的启示

DeepSeek-V2 的 MoE 架构（16 experts，每次激活 2）让 prefill 更复杂：不仅要算 QKV，还要路由 expert。我们的 Triton kernel 已扩展支持 expert parallelism，但需要修改 routing logic。好消息是：MoE 的 prefill 计算密度更高，kernel 优化收益更大——在 DeepSeek-V2-236B 测试中，prefill 加速比达 4.1×（vLLM 原生为 1.8×）。

Qwen3 的 GroupRope 位置编码对 kernel 有微小影响，需调整tl.arange的 offset 计算，但整体结构不变。我们已验证 Qwen3-32B-AWQ 在 A100×4 上 P95 降至 1.42s（vLLM 为 3.89s）。

6.3 为什么不用 SGLang？——一个务实的选择

SGLang 是优秀的框架，它用 Python DSL 描述 pipeline，灵活性极高。但我们没选它，因为：

• SGLang 的 runtime 仍基于 vLLM 的 PagedAttention，prefill 瓶颈未解；
• 它的 Python DSL 在高并发下有 GIL contention，我们实测 128 RPS 时 CPU usage 达 92%；
• 它的 streaming support 依赖 asyncio，与我们现有的 sync HTTP server 不兼容。