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 ...
GPTQ
上文提到，OBQ 的致命痛点在于 “行依赖（Row-wise
dependency）”
由于采用了贪心策略，权重矩阵每一行选出的最优量化顺序都不一样，导致
Hessian 逆矩阵
的更新无法在不同输出通道（行）之间共享，复杂度高达 。
GPTQ (Generative Pre-trained Transformer Quantization)
正是为了解决这个工程灾难。它在 OBQ
的数学基础上做了极其巧妙的近似和系统级优化，将复杂度成功降到了 。
GPTQ 主要做了以下四大核心优化
优化一：统一量化顺序
OBQ使用贪心选择量化顺序，实验表明这对LLM带来的收益微乎其微。
GPTQ
团队通过实验和理论发现：如果不去贪心地寻找每次误差最小的权重，而是强行规定对所有行使用相同的、固定的顺序（例如从左到右，按列
依次量化），最终的精度损失几乎可以忽略不计。
这一步为什么是破局的关键？
回顾一下 Hessian 矩阵的定义：。
你会发现，
完全只依赖于输入激活值 ，和权重 没有任何关系！
因此，对于权重矩阵的每一行，初始的 是完全一模一样的。
在 OBQ 中，因 ...
