SGLang
KV Cache Retract 机制详解:普通模式 vs PD 分离模式
一条 input 长度逼近 KV cache 容量的请求,在 decode
阶段会触发一次看起来很费时间的 “retract” 交互。本文讲清楚 retract
到底是什么、为什么会被触发、怎么实现的,以及它在普通模式和
PD(prefill-decode)分离模式下走的是两条完全不同的路。
retract 解决什么问题
decode 阶段每生成一个 token,每个跨越 page 边界的请求都要从 KV cache
池子里申请新 page。当池子见底、下一步 decode
放不下时,系统有两种选择:
直接 OOM 崩掉 scheduler;
主动撤回若干个正在 decode 的请求,释放它们的 KV
cache,把腾出来的空间留给剩下的请求继续跑。
retract 就是第 2 种。它是 decode OOM
的兜底机制,用”牺牲少数请求”换”整个 batch 不崩”。
SGLang 作者自己对 retract 的评价很直白,在 PR #708 的描述里写着:
Retract de ...
HiSparse:把长上下文
decode 的容量瓶颈从 GPU 显存搬到 CPU 内存
长上下文推理里,稀疏注意力解决了”算不动”的问题,却没解决”装不下”的问题。一个
1M 上下文的请求,即便每步 decode 只看 top-k 个 token,它的完整 KV cache
仍然得常驻 GPU HBM,否则没法被快速访问。结果是 batch size
很快被显存撑爆,throughput 在并发还没拉起来时就见顶了。
HiSparse(Hierarchical Sparse Attention)就是 SGLang
社区针对这个容量瓶颈做的 decode 优化:把完整但冷的 KV 放进 CPU pinned
memory,GPU 上每个请求只留一个固定大小的”热”buffer,decode 时按需把 top-k
的 KV 从 host swap-in 回来。实测在 GLM-5.1-FP8 的长上下文场景下吞吐最高
5×,256 并发时约 3× 基线。
这篇文档面向想搞清楚 HiSparse
到底省了什么、代价是什么、能不能扩大系统 KV
容量的工程师。代码引用基于本仓库 r ...
coredump使用方式
首先新建test_coredump.py脚本，脚本内容如下：
123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657import os
import socket
import torch
from torch.utils.cpp_extension import load_inline
# 提前获取 hostname 和 pid，方便打印出触发 coredump 的准确命令
hostname = socket.gethostname()
pid = os.getpid()
pipe_path = f"/tmp/cuda_pipe_{hostname}_{pid}"
dump_path = f"/tmp/cuda_coredump_full_{hostname}_{pid}"
cuda_source = """
__global__ void hang_kernel(int* dummy) ...
如果把conda/miniconda换一个目录，会出现一些脚本bash
bang等硬编码写死路径的问题，需要修复
可以执行以下脚本修复
12345678910111213141516171819202122232425262728293031323334353637383940
#!/bin/bash
set -ex
OLD_PATH="/nfs/ofs-llm-ssd/user/gogongxt/miniconda3"
NEW_PATH="/nfs/gogongxt/miniconda3"
# 1. 修复 base 环境的 bin 目录
if [ -d "$NEW_PATH/bin" ]; then
cd "$NEW_PATH/bin"
# -r 参数表示如果 grep 没有输出，xargs 不会执行后面的 sed
grep -rIl "$OLD_PATH" . | xargs -r sed -i "s|$OLD_PATH|$NEW_PATH|g"
fi
# 2. 修复 Conda 自带的配置和环境激活脚本
if [ -d "$NEW_PAT ...
tbo-sbo
tbo 就是 two batch
overlap，最开始由deepseek提出，主要的目标是基于deepep的实现思路，优化计算和通信的overlap
在整体的计算流程中，有
attention -> alltoall -> moe -> alltoall -> ...，我们可以构造两个batch，实现一个做矩阵计算，另一个做alltoall通信，从而overlap
tbo是给decode准备的，主要原因是decode是带宽瓶颈，计算耗时短，两个batch可以相互干扰影响低
不过实际蚂蚁在h20上部署deepseek时发现了tbo在高并发下反而会更慢，因为h20算力太低，高并发会延迟爆炸不满足slo
因此也是联合deepseek，sglang提出了sbo：
两点优化：
计算shared experts时进行dispatch通信
计算down的时候就边计算边发送，按照block_m粒度进行处理，一个块好了就直接发送
sbo的pr链接：
SBO in SGLang: https://github.com/sgl-project/ ...
通信操作
主要介绍AllGather和AllReduce，别的通信操作可以参考下面的文档
NOTE
Ref:
NV
NCCL通信库
华为HCCL通信库
blog
blog
blog
nv blog
massively-scale-deep-learning-training-nccl
nv blog
fast-multi-gpu-collectives-nccl
AllGather
示意图：
每张卡上都有一个矩阵的一部分，AllGather让每张卡都有完整的矩阵
ring-AllGather 算法：
所谓ring，就是环形处理，每张卡把输出传到下一张卡：
假设一个
的矩阵，每个卡有P分之一参数，每次传输每张卡都是发送了
数据量，总共P-1次传输，总通信量就是：
注意我们讲的通信量一般都是单向的，计算效率是也是除上单向的带宽
AllReduce
示意图：
AllReduce操作是将通信域内所有节点的输入数据进行归约操作后（支持 ...
vllm-ascend 与 vllm 的关系
一句话总结：vllm-ascend 是 vllm 的硬件插件（out-of-tree
platform plugin），通过 Python entry_points 机制注入，覆盖 GPU
相关的算子/调度/通信层，复用 vllm 的 API
服务、模型定义、请求管理等上层逻辑。
各模块的职责分工
模块
归属
说明
API Server
vllm
OpenAI 兼容的 HTTP 服务，vllm serve 启动，vllm-ascend
不做任何修改
请求管理/Tokenization
vllm
请求接收、分词、结构化输出等
Scheduler 调度
vllm（默认）
vllm 的
v1/core/sched/scheduler.py，平台无关的请求调度、KV cache
管理
Scheduler 扩展
vllm-ascend（可选）
SchedulerDynamicBatch（动态batch调整SLO）、RecomputeScheduler（PD分离场景重计算），均继承自
vllm Sche ...
awq
AWQ 的研究团队发现了一个非常有意思的现象：对于 LLM 的权重量化，只有
0.1%-1% 的权重（Salient Weights）决定了最终的精度。
也就是在量化时，把1%的重要权重保留成bf16，其余是INT8，量化后的精度非常好
但是这里引入两个问题：
要怎么确定哪些是这1%重要的权重
硬件对这种W8/BF16混合精度计算很难高效实现
IMPORTANT
AWQ的解决思路：“不改变权重的存储格式，而是通过数学等价变换来保护显著权重。”
相比于直接保留
FP16，如果我们把显著权重在量化前按比例放大，它们在量化网格中的相对误差就会减小。同时，为了保证等价性，我们在计算时将对应的激活值按同等比例缩小。
这一块对激活值和权重的操作和SmoothQuant很像，只是SmoothQuant是想把激活值更平滑，AWQ是保护重要权重
核心算法与公式推导
权重的重要性在于：它不看权重本身的大小，而是看输入激活值的大小，因此每个权重通道的 ...
SmoothQuant
之前我们讲的算法是OBQ/GPTQ是量化权重的，把权重量化到W4/8，实际计算时还要转成bf16，并没有使用上INT8
tensor core加速计算
为了实现这个，我们就需要去量化激活值，这篇要介绍的SmoothQuant也是近几年来简单并且适用范围非常广的方法
对于LLM的激活值观察：
在LLM中，少数通道的 outlier
主导了量化范围，并且这部分都是分布在特定的通道上，普通值有一定比例关系，我们可以考虑把这部分异常比例关系移到权重上，让激活值变得平滑，Max更小更好量化
核心思想：等价缩放，迁移难度
回到线性层的基本运算 。我们引入一个对角缩放矩阵 ：
注意矩阵”左行右列”，就好理解了，数学上完全等价，输出 一模一样。
物理直觉：假设第 个通道的激活有严重的 outlier（max =
100），而对应通道的权重范围很小（max = 0.5）。如果我们选 ：
激活第 通道被缩小 10 倍，max
从 100 降到 10，量化精度大幅提升
权重第 通道被放大 10 倍，max
从 0.5 变成 5，权重本身范围小，放大后仍在 ...
NOTE
论文：
Medusa https://arxiv.org/pdf/2401.10774
EAGLE1 https://arxiv.org/pdf/2401.15077
EAGLE2 https://arxiv.org/pdf/2406.16858
EAGLE3 https://arxiv.org/pdf/2503.01840
在sglang
0.5.9版本中，实现了MTP/NEXTN，EAGLE，EAGLE3，单独draft小模型，NGRAM
的投机采样算法
其中MTP和NEXTN其实也是基于eagle的增加一些if,else修改实现的
没有实现Medusa
Algorithm
Enum Value
Description
EAGLE
EAGLE
Original EAGLE (Eagle 1) - draft model based speculative
decoding
EAGLE3
EA ...
