▣ Source Reading · File № 002 An autopsy of the inference engine SGL-001 / 2026

SGLang

Two threads of control, six ZMQ sockets, twenty-seven attention backends, and a radix tree that turns a chat prefix into a cache hit. A reading of 729,000 lines.

LAB NOTEBOOK
SGLang
v0.5 · ROCm 7.0
Source
728,969loc · py
Scheduler
4,006lines · 1 file
Attn backends
27incl. AITER
HW targets
5NV · AMD · MUSA · Apple · NPU
Kernel pkg
sgl-kernelseparate wheel
Reading time
~6hthree sessions

§ 00 · PrologueWhat we're after.

SGLang 这次精读的目标,是把"高吞吐、低延迟、跨硬件"这三件事——分别由 RadixAttention、Continuous Batching with Overlap Scheduling、可插拔 Attention Backend 三大子系统承担——拆开来看个明白。它的源码量比 vLLM 还大;而且不像 PyTorch 那种"按 op/module 切片"的结构,SGLang 是按系统组件切片(tokenizer / scheduler / runner / cache / kernel),所以读起来更像读一个分布式系统。

精读分 10 个模块(M0 → M9)。本档案是读后的总览:Tokenizer/Scheduler/Workers 的三进程拓扑、ZMQ 消息总线、Overlap 调度器的时间线、RadixCache 的树形结构、27 个 attention backend 中 aiter_backend.py 给 MI300X 提供的入口,以及 sgl-kernel 这个独立 wheel 包的多平台构建系统。

4,006
scheduler.py
3,607
model_runner.py
3,284
aiter_backend.py
1,349
engine.py
828
radix_cache.py
27
attn backends
19
mem_cache files
20+
managers files
Plate I — Three processes, one engine tokenizer · scheduler · detokenizer, conjoined by ZMQ 1 : ∞
launch_server.py → 3 cooperating processes + N worker procs Process A · Tokenizer Manager TokenizerManager HTTP/FastAPI entry HF tokenizer · encode req → TokenizedGenerateReqInput async response router async event loop + uvicorn server 1 Process B · Scheduler (the heart) Scheduler scheduler.py · 4006 lines event_loop_normal() event_loop_overlap() ★ schedule_policy schedule_batch RadixCache prefill_delayer communicator · ZMQ I/O 2 Process C · Detokenizer DetokenizerManager incremental decode token_ids → text chunks stop-string detection stream chunk → TokMgr 3 Process(es) D · TP/EP workers ModelRunner model_runner.py · 3607 model.forward(batch) attn_backend (1 of 27) CUDAGraphRunner ★ memory_pool (KV cache) sampling (logits → tokens) N processes for TP/EP 4 ZMQ PUSH ZMQ PUSH ZMQ PUSH ZMQ PUSH tensor IPC + ZMQ (req batch) output tokens (per step)
四个进程(或进程组),由 ZMQ 消息总线串起来。Tokenizer 是 user-facing HTTP 入口;Scheduler 是大脑(4006 行);Detokenizer 增量解码;Worker 是 N 个真正跑模型的进程(TP/EP 拆分)。Tokenizer 与 Scheduler 之间走 ZMQ PUSH,Scheduler ↔ Worker 之间是更高带宽的 tensor IPC + ZMQ 控制信令。

§ M0 · 20 minThe lab, laid out.

SGLang 不是单包仓库 — 顶层就有 python/(Python 主包)+ sgl-kernel/(独立的 kernel wheel)+ rust/(一些 Rust 组件)+ sgl-model-gateway/(gateway 服务)+ proto/(gRPC 协议定义)。这种"主包 + 多 sub-package"结构反映了不同组件不同节奏发布的工程现实。

Python 主包 python/sglang/ 分两大块:

SubpackageWhat
sglang/lang/前端 DSL — Structured Generation Language(select/gen/fork/regex 等 primitives)。这是 SGLang 名字的来源。
sglang/srt/Serving Runtime — 推理引擎本体,下面所有模块都属于它

sglang/srt/ 进一步切分(这是真正的"系统"):

SubdirRole
entrypoints/Engine / HTTP server / gRPC server / OpenAI 兼容层
managers/★ scheduler.py / schedule_batch.py / schedule_policy.py — 调度核心,20+ 文件
model_executor/★ model_runner.py · 3607 行 — 跑模型 forward,管 CUDAGraph
mem_cache/★ radix_cache.py · KV cache 池 — RadixAttention 实现
layers/attention/★ 27 个 attention backend(aiter / flashinfer / triton / cutlass_mla / ...)
models/具体模型定义(Llama / Qwen / Kimi / DeepSeek / ...)
hardware_backend/gpu / mlx / musa / npu — 各硬件平台的 hook
speculative/投机解码(draft / target / EAGLE)
disaggregation/P/D 分离(prefill / decode 拆到不同节点)
distributed/ · elastic_ep/TP / EP / DP 通信和弹性扩展
checkpoint_engine/模型权重加载和热重启
compilation/编译器集成(torch.compile / TVM-like)
constrained/结构化生成(JSON schema / regex / EBNF)
connector/对接 vLLM、TensorRT-LLM 等外部 runtime
观察 · SUBPACKAGE STRATEGY

sgl-kernel 作为独立包是 SGLang 工程上最值得偷的设计。它有自己的 pyproject.toml / pyproject_cpu.toml / pyproject_musa.toml / pyproject_rocm.toml — 同一份 C++/Triton 源码,针对 5 个平台(CUDA / ROCm / MUSA / CPU / Metal)构建 5 个不同的 wheel。这样既允许 sgl-kernel 独立 release(不用等主包),又把"kernel 调优"和"runtime 调度"两件事的迭代节奏完全解耦。

这正是你 AMD 工作直接相关的模式 — 如果将来 multi-agent kernel optimization 系统要产出可被其他系统消费的 kernel,做成独立 wheel 包 + 多平台 pyproject 是范本。

§ M1 · 30 minThe Engine, started.

用户最常用的入口有两条路径:(1) HTTP 服务模式 — python -m sglang.launch_server --model X --tp 8;(2) Python in-process 模式 — from sglang import Engine; engine = Engine(model="X")。两条都最终落到 sky/entrypoints/engine.pyEngine 类(1349 行,line 174)。

class SchedulerInitResult:
    """Engine 启动期间收集每个 scheduler 子进程的初始化结果。"""

class Engine(EngineScoreMixin, EngineBase):
    """The main engine class. Holds the tokenizer manager,
    spawns scheduler procs, manages their lifecycle."""

Engine 启动序列(line 174-1158 大概结构):

  1. _set_envs_and_config(line 1159)— 设环境变量(CUDA_VISIBLE_DEVICES、ROCm 等),GC 调优
  2. init_tokenizer_manager(line 126)— 起 TokenizerManager 进程
  3. _compute_parallelism_ranks(line 1332)— 算 TP × DP × EP 拓扑下的 rank 分配
  4. spawn N 个 scheduler 子进程(multiprocessing),每个进程跑一个 Scheduler.event_loop()
  5. _wait_for_scheduler_ready(line 1265)— 等所有 scheduler 通过 ZMQ ready barrier
  6. HTTP / gRPC server 接管(如果是 server 模式)

关键细节:每个 TP rank 是一个独立 scheduler 进程(不是 thread),用 multiprocessing spawn。Tokenizer 是单独一个进程,Detokenizer 又是单独一个。所以一个 TP=8 的部署有 1 + 1 + 8 = 10 个 Python 进程 共同协作。

Easy mistake · DON'T

不要把 SGLang 当成"单进程 Python 服务" — 它是个 IPC-heavy 的多进程系统。任何对 scheduler 的 monkey-patch、log injection、observability 改动都要考虑"在哪个进程里生效"。kernel_api_logging.py 这种 hook 必须在所有 scheduler 进程里被 import 才能生效。

§ M2 · 25 minZMQ — the wires.

四个进程之间的消息总线是 ZeroMQ。SGLang 用 zmq 库的 PUSH/PULL 模式(不是 PUB/SUB)— 因为请求需要严格的 producer→consumer 单向流动,PUSH/PULL 比 PUB/SUB 多一层 fan-out / load-balancing 语义。

主要 socket 拓扑:

From → ToWhatMode
HTTP client → TokenizerManagerHTTP requestFastAPI / uvicorn
TokenizerManager → SchedulerTokenizedGenerateReqInputZMQ PUSH
Scheduler → Worker(s)ScheduleBatch (tensor IPC)ZMQ + shared memory
Worker → Scheduleroutput tokens / logitsZMQ PUSH
Scheduler → DetokenizerManagertoken_ids batchZMQ PUSH
DetokenizerManager → TokenizerManagerdecoded text chunksZMQ PUSH

注意是环形拓扑:Tokenizer → Scheduler → Worker → Scheduler → Detokenizer → Tokenizer。Tokenizer 既是入口也是出口,负责把 streamed 输出聚合后通过 HTTP 推给客户端。

为什么 ZMQ 而非 gRPC · DESIGN CHOICE

ZMQ 比 gRPC 轻 — 没有 protobuf 序列化开销(除非用户自己加),没有 HTTP/2 帧开销,socket 直接走 Unix domain 或 IPC。对于"同机多进程" + "高频小消息" + "大 tensor 走共享内存"的场景,ZMQ 比 gRPC 快很多。SGLang 选 ZMQ + 自己定义 dataclass payload,是经过权衡的工程选择。

代价是缺乏 schema 演进保护 — 改 payload 字段需要所有进程一起重启。但 SGLang 是单一二进制部署,这不是大问题。

§ M3 · 60 minThe Scheduler — 4006 lines in one class.

managers/scheduler.py 是 SGLang 整个 codebase 里最重的一个文件。class Scheduler 从 line 324 开始,__init__ 在 line 340,光构造函数就有几百行 — 它要初始化 ZMQ sockets、tokenizer、tp_group、model_runner、KV cache pool、radix cache、sampling、speculative decoding 等几乎所有子系统。

核心方法(按调用频率排):

MethodLineFrequencyWhat
event_loop_normal()1524per server普通调度循环 — sequential
event_loop_overlap()1551per server★ overlap 调度 — 计算/调度并发
process_input_requests1829per step从 ZMQ 收新请求
get_next_batch_to_run2461per step挑下一个要跑的 batch
update_running_batch2870per step合并新 prefill 进 running batch

Scheduler 类用了几个 Mixin 模式来切分代码(line 248 / 324 等):

  • SchedulerMlxOverlapMixin — Apple MLX 后端特化
  • scheduler_dp_attn_mixin — DP attention 特化
  • scheduler_output_processor_mixin — 输出后处理
  • scheduler_input_blocker — 输入背压

这种"主类 + Mixin"是 SGLang 切分 4006 行的方式。但这也是个学习成本 — 跳到一个方法时不一定知道它属于哪个 Mixin。读源码时记得 grep "def method_name" 而不是 grep "class .*:"

# 简化后的 event_loop_overlap 核心结构
def event_loop_overlap(self):
    """Overlap: while GPU runs forward(t), CPU plans forward(t+1)"""
    last_batch = None
    while True:
        # 1) 从 ZMQ 收新请求,加入 waiting queue
        recv_reqs = self.recv_requests()
        self.process_input_requests(recv_reqs)

        # 2) 决定下一个 batch (CPU work)
        batch = self.get_next_batch_to_run()

        # 3) 启动 forward (GPU work, async via cuda stream)
        if batch is not None:
            future = self.forward_async(batch)  # ← non-blocking

        # 4) 等上一个 batch 的 GPU 结果 (overlapped with #2/#3)
        if last_batch is not None:
            result = last_batch.future.wait()
            self.process_batch_result(last_batch, result)

        last_batch = batch
★ 关键创新 · OVERLAP SCHEDULING

"普通"循环是 plan(t+1) → run(t+1) → process(t+1) 严格串行的。Overlap 循环把"plan 下一步"和"GPU 执行当前步"并发起来 — CPU 在算下一个 batch 的 schedule 决策时,GPU 还在跑当前 forward。在 prefill-heavy workload 上这相当于把 CPU planning overhead 完全吸收掉,吞吐通常 +10-25%。

这跟 SkyPilot Optimizer 那种"乐观决策 + 锁内兜底"在精神上一致 — 都是用"提前做事"来藏掉延迟。但 SGLang 这里是 GPU/CPU 异步,靠的是 CUDA stream 和 future 抽象。

Plate II — Overlap scheduling, in time CPU plans while GPU runs temporal
event_loop_overlap() vs. event_loop_normal() NORMAL CPU GPU plan t=1 forward t=1 (GPU) plan t=2 forward t=2 (GPU) plan t=3 forward t=3 (GPU) → GPU 利用率低 OVERLAP ★ CPU GPU plan t=1 forward t=1 (GPU) plan t=2 forward t=2 (GPU) plan t=3 forward t=3 (GPU) plan t=4 forward t=4 (GPU) → GPU 几乎不停 ★ overlap 区间 — CPU plan 完全藏在前一步 GPU forward 里 t₀ t →
Normal 循环 GPU 等 CPU,CPU 等 GPU,两边交替忙闲。Overlap 把"plan t+1"提前到"forward t"还在跑的时候 — CPU 时间被完全藏起来。代价是要管理一个 last_batch.future 状态,多了一层并发逻辑。SGLang 默认开 overlap,除非 backend 不支持 async forward。

§ M5 · 40 minRadixAttention, the prefix tree.

SGLang 在论文里的最大卖点是 RadixAttention — 把多个请求共享的 prompt prefix 自动复用 KV cache。实现在 mem_cache/radix_cache.py(828 行),三个核心类:

class RadixKey:        # 66  — token sequence → hashable key
class TreeNode:        # 206 — radix tree 的节点 (children, key, value, ref_count)
class RadixCache(KVCacheEventMixin, BasePrefixCache):  # 269

Tree 上每个节点:

  • key: 一段 token 序列(RadixKey,可哈希)
  • value: 这段 tokens 对应的 KV cache 的物理 slot indices
  • children: 后续 token 派生出的子节点
  • ref_count: 多少个活跃请求正在引用这段 prefix
  • last_access: LRU 时间戳

核心操作(来自 BasePrefixCache 抽象基类):

OperationWhat
match_prefix(key)沿树查最长匹配的 prefix,返回已 cache 的 KV slot 列表
insert(key, value)把新生成的 tokens 和它们的 KV slot 加进树(可能在中间节点分裂)
evict(num_tokens)当 KV pool 满时,按 LRU + ref_count=0 的策略驱逐
inc_lock_ref / dec_lock_ref请求开始/结束时增减 ref_count(防止驱逐还在用的 prefix)

mem_cache/ 还有 18 个其他文件 — 不同形态的 cache:hiradix_cache.py(多级 hierarchical)、chunk_cache.py(不用 radix 的简单 chunk)、mamba_radix_cache.py(Mamba/SSM 模型的特化)、hisparse_memory_pool.py(稀疏激活)等。还有 cpp_radix_tree/ — 一个 C++ 实现的 radix tree,给极端性能场景用。

Plate III — A RadixCache, populated three chat requests sharing two prefixes structural
"You are a helpful" tokens: [You, are, a, helpful] ref_count = 2 · slots: [0-3] "Translate this:" [Translate, this, :] ref_count = 1 · slots: [4-6] "def fibonacci(n):" [def, fibonacci, (, n, )] ref_count = 0 · evictable "assistant. User asks:" [assistant., User, asks, :] ref_count = 1 · slots: [7-10] "AI. Reply with care." [AI., Reply, with, care, .] ref_count = 1 · slots: [11-15] "How to print hi?" REQ #1 active leaf · slots [16-20] "What is 2+2?" REQ #2 active leaf · slots [21-23] Req #1 and #2 share the first two ancestor nodes → "You are a helpful" (4 tokens) "assistant. User asks:" (4 tokens) = 8 cached KV tokens reused free. LEGEND internal node branching point evictable (ref=0) leaf · active req
新请求来时调 match_prefix(tokens),沿树查最长公共前缀,已 cache 的 token 直接复用 KV slots — 跳过这部分的 prefill 计算。中间节点在新请求"分叉"时会动态分裂(split)。ref_count=0 的子树在 KV pool 不够时按 LRU 驱逐。这就是 RadixAttention 在多轮对话 / 系统 prompt 共享场景下的核心机制。

§ M6 · 35 minThe ModelRunner — 3607 lines.

model_executor/model_runner.py 是另一个怪兽。class ModelRunner(ModelRunnerKVCacheMixin) 在 line 323,ModelRunnerOutput dataclass 在 315。它负责的事:

  1. 从 HF 或自定义 loader 加载模型权重
  2. 分配 KV cache memory pool(巨大 tensor,按 layer × block × head 切片)
  3. 选择并初始化 attention backend(从 27 个候选里挑)
  4. build / capture CUDA Graph(line 2339-2347 的关键分支)
  5. 每次 forward(batch) 时把 ScheduleBatch 翻译成 ForwardBatch(含位置编码、cache indices、mask 等),调 self.model.forward(...)
  6. 给 attention backend 调 init_forward_metadata(forward_batch) 准备元数据(page table / sliding window mask)
  7. 对接 sampling(top_k / top_p / 结构化生成)

CUDA Graph 这一段值得单独说 — SGLang 提供多种 capture 策略:

  • cuda_graph_runner.py — 普通整张图 capture
  • piecewise_cuda_graph_runner.py — 分段 capture,更灵活
  • breakable_cuda_graph_runner.py — 可中断的 capture(用于 disagg P/D)
  • cpu_graph_runner.py — CPU 版本(用于 host-side 模拟)
  • mindspore_runner.py — MindSpore 平台特化
为什么 ModelRunner 也是巨型类 · CONSTRAINT

3607 行单类的原因和 Scheduler 4006 行类似 — 它要协调"模型、KV cache、attention backend、CUDA Graph、sampling、speculative"六个子系统在每次 forward 时的精确顺序。任何把它拆成多个小类的尝试都要面对"这些子系统之间状态依赖严重"的问题。SGLang 选择了"一个大类 + 几个 Mixin",跟 Scheduler 是一脉相承的工程取舍。

对 reader 不友好。但对维护者来说,至少所有关键状态在一个 self.* 命名空间里。

§ M7-M8 · 50 minTwenty-seven backends, one interface.

layers/attention/ 是 SGLang 跨硬件能力的核心 — 27 个 attention 实现共享一个 AttentionBackend 抽象基类(base_attn_backend.py),attention_registry.py 维护"backend 名字 → class"映射。

FamilyBackendsHardware
FlashAttention 系flashattention_backend, flashinfer_backend, flashinfer_mla_backend, flashmla_backendNV (CUDA)
Tritontriton_backend, triton_ops/NV + AMD(Triton 跨平台)
CUTLASScutlass_mla_backendNV (Hopper / Blackwell)
TRT-LLMtrtllm_mha_backend, trtllm_mla_backendNV (TensorRT-LLM)
Torch nativetorch_native_backend, torch_flex_backendany (slow fallback)
AMD AITERaiter_backendAMD ROCm (MI300X+)
Intelintel_amx_backendIntel CPU + AMX
Mamba / State spacemamba/, linear/any
NSA (sparse)nsa_backend, nsa/NV
DSv4 (DeepSeek)deepseek_v4_backend, dsv4/NV + AMD
Hybridhybrid_attn_backend, hybrid_linear_attn_backend, dual_chunk_flashattention_backendany
TBOtbo_backend多 backend 二级路由

AITER 后端 — 给 MI300X 的入口

aiter_backend.py 一个文件就 3284 行class AiterAttnBackend(AttentionBackend) 在 line 117。它在 SGLang 里的位置:用户跑 sky launch --infra k8s ... sgl --attention-backend aiter,scheduler 用 AITER 的 forward / metadata 接口,背后调 aiter 这个独立 Python 包(github.com/ROCm/aiter)的 CK / hipBLASLt kernel。

class AiterAttnBackend(AttentionBackend):    # line 117
    def init_forward_metadata(self, forward_batch):
        # build page table, mask, sliding window indices
        ...
    def forward_extend(self, ...): ...    # prefill
    def forward_decode(self, ...): ...    # generation
    def forward_target_verify(self, ...): ...  # speculative decoding

class AiterIndicesUpdaterPrefill:            # line 2962
class AiterMlaIndicesUpdaterPrefill:         # line 3064  ← MLA variant
class AiterMultiStepDraftBackend:            # line 3146  ← speculative draft

三个 helper 类分别处理:

  • Indices updater — 把 SGLang 的 RadixCache slot indices 翻译成 AITER 期望的 page table 布局(每个 kernel 都要求一种特定的内存布局)
  • MLA indices updater — DeepSeek / Kimi-K2 的 Multi-head Latent Attention 特化(压缩 KV)
  • MultiStep draft — 投机解码 draft model 的多步 forward 优化
★ 直接相关你的工作 · DEBUG ENTRY POINT

你在 CLAUDE.md 提到的 AITER 调试 agent team — "Kernel Debugger" 负责 aiter/csrc/py_itfs_ck/,"SGLang Integrator" 负责 python/sglang/srt/layers/attention/。后者具体的入口就是这里:aiter_backend.py + AiterIndicesUpdaterPrefill + AiterMlaIndicesUpdaterPrefill。任何"参数组合不被 kernel 接受"的报错都先从这里追,看 indices updater 构造的张量是否符合 AITER kernel 的预期。

Plate IV — Attention backend constellation one interface, twenty-seven implementations cross-hardware
AttentionBackend base_attn_backend.py .init_forward_metadata(.) .forward_extend / _decode one interface — 27 implementations NVIDIA · CUDA flashattention flashinfer flashinfer_mla flashmla cutlass_mla trtllm_mha trtllm_mla nsa + dual_chunk_flashattention AMD · ROCm ★ aiter_backend 3284 lines · MI300X / MI355X AiterIndicesUpdaterPrefill (line 2962) AiterMlaIndicesUpdaterPrefill (3064) AiterMultiStepDraftBackend (3146) Cross-platform · Triton triton_backend triton_ops/ torch_flex_backend Model-specific paths deepseek_v4 mamba1/2 hybrid_attn / hybrid_linear_attn / fla Fallback · slow torch_native intel_amx tbo / linear / nsa_backend → AITER 是你 AMD 工作的 主入口点。MI300X 上跑 Kimi-K2 / Qwen3-Next 时它是 forward 路径 里最热的代码。
27 个 backend 不是平均分布 — NV 系最丰富(FlashAttention 各种变体 + CUTLASS + TRT-LLM),AMD 主力是 AITER 一根扁担挑全场(外加 Triton 跨平台兜底)。对你的工作含义:AITER 上每多一种 kernel 优化(FA variant、MLA、sparse、speculative),都要在 aiter_backend.py 里加对应的 IndicesUpdater,因为 SGLang scheduler 还是按 NV 那套 page table 抽象在喂参数。

§ M9 · 25 minsgl-kernel — five wheels, one tree.

sgl-kernel/ 是 SGLang 工程组织最值得偷的部分 — 把 kernel 代码做成独立 wheel 包,5 个平台各一份 pyproject_*.toml,但共享同一份 C++ / Triton 源码。

Build targetpyprojectTooling
CUDA wheelpyproject.toml(默认)CMake + nvcc + ninja
ROCm wheelpyproject_rocm.tomlCMake + hipcc + composable_kernel
CPU wheelpyproject_cpu.tomlCMake + cc
MUSA wheelpyproject_musa.tomlCMake + Moore Threads tooling
Metal binarysetup_metal.pyXcode / Metal compiler

包结构:

sgl-kernel/
├── csrc/             # C++ / CUDA / HIP 源码
├── include/          # 头文件
├── python/           # Python binding(pybind11 / cppextension)
├── cmake/            # 各平台的 CMake module
├── CMakeLists.txt    # 主 CMake
├── build.sh          # 构建脚本
├── Makefile
├── pyproject.toml    ← CUDA
├── pyproject_rocm.toml
├── pyproject_cpu.toml
├── pyproject_musa.toml
├── setup_metal.py
└── setup_musa.py
值得偷的设计 · INDEPENDENT KERNEL WHEEL

把 kernel 做成独立 wheel 的工程价值:(1) 独立 release 节奏 — kernel 团队可以发新 wheel 不用等主包;(2) 多平台 wheel 同名,靠 pip 安装时的 platform tag 自动挑对;(3) 用户可以在没装 SGLang 主包的情况下单独用 kernels(比如别的 inference engine 想借 SGLang 的高性能 attention kernel)。

对你 multi-agent kernel optimization 系统:把你做出来的高性能 AMD kernel 也做成独立 wheel,让 vLLM / SGLang / 任何 PyTorch 用户都能直接 pip install amd-kernels 用。这是把工作 leverage 出去的最直接路径。

§ Traps · five of themWhat I wish I'd known.

读 SGLang 时最容易栽的几个坑。每个坑都对应一类具体的失误模式。

Trap N° 1 · Scheduler is multi-process

不是 multi-thread。你 monkey-patch Scheduler.event_loop_normal 在主进程是无效的 — 每个 TP rank 是一个独立 multiprocessing.Process,它们 spawn 时从头 import 整个 sglang 包,之前在主进程注入的 hook 完全丢失。要影响所有 scheduler,hook 必须放在 sglang.srt.__init__ 或 module-level,import 时就生效。

Trap N° 2 · overlap loop 改坏会 deadlock

event_loop_overlaplast_batch.future 异步等 GPU 结果。如果你改 schedule 逻辑导致某次 future 永不 resolve(比如 batch 被错误地 drop),整个进程会卡死。建议任何改 overlap 路径的人先打开 SGLANG_DEBUG=1 看 future state。

Trap N° 3 · RadixCache 的 ref_count 极易泄漏

RadixCache 每个匹配到 prefix 的请求要 inc_lock_ref,请求结束要 dec_lock_ref。如果一个 request 在 setup 时崩了,没人帮你 dec,那段 KV slots 永远不会被驱逐,最终 OOM。所有 catch 异常分支都要小心 ref 是否清理。cache_controller.py 有 helper 但不强制。

Trap N° 4 · 27 个 backend 各有"必填"字段

不同 backend 对 forward_batch 的元数据要求不一样 — 有的要 seq_lens_cpu,有的要 extend_seq_lens_int32,有的要 page-aligned tensor。init_forward_metadata 是每个 backend 自己实现的,遗漏一个字段就会在 kernel 里炸(且报错信息常常很神秘)。切 backend 前先看目标 backend 的 init_forward_metadata 签名和断言。

Trap N° 5 · CUDA Graph capture 时数据必须 reproducible

CUDA Graph capture 时所有输入张量的形状必须固定。如果你的 batch 形状是动态的(变 padding),cuda_graph_runner.py 会按 padding bucket 准备多张 graph。改 batch 构造逻辑时,确认 padding 策略不要打破 bucket 假设,否则会触发 piecewise_cuda_graph_runner.py 的 fallback,性能掉一半。

§ Red lines · threeThe big questions.

如果要给 SGLang 加一个新的 attention backend(比如自家 AMD 的某个新 kernel),要改哪些文件?
  1. sky/srt/layers/attention/<new>_backend.py — 实现 AttentionBackend 抽象(init_forward_metadata / forward_extend / forward_decode / forward_target_verify
  2. sky/srt/layers/attention/attention_registry.py — 注册名字 → class 映射
  3. sky/srt/server_args.py — 加入 --attention-backend <new> 的选项校验
  4. sky/srt/model_executor/model_runner.py_get_attention_backend() 方法 — 必要的话加 fallback 逻辑
  5. (推荐)test/srt/test_attention_backend.py — 加 smoke test 跑短 prefill + decode
  6. (如果是 ROCm-specific)sgl-kernel/ 里加配套的低层 kernel

关键参考:aiter_backend.py 是最近也是最大型的"加新 backend"实例,可以照抄结构。

一个 chat completion 请求,从 HTTP POST 到 token 返回,跨越多少个进程边界、多少种通信形式?

进程:HTTP client → TokenizerManager → Scheduler → ModelRunner(N 个 TP workers)→ Scheduler → DetokenizerManager → TokenizerManager → HTTP client。共 ~7 个 Python 进程(如果 TP=8,N=8)。

通信形式:HTTP/JSON、ZMQ PUSH/PULL(结构化 message)、tensor IPC(shared memory)、CUDA stream events(在 worker 内部 GPU 异步)、NCCL collective(TP 内 all-reduce)。5 类通信

这跟 SkyPilot 的"4 进程 + HTTP + spawn"复杂度相当 — 但 SGLang 多了 tensor IPC 和 NCCL 这两层 GPU-aware 通信,因为它要传巨大的 KV cache 张量。

SGLang 比 vLLM 快的关键在哪里?

四个具体来源:

  1. RadixAttention 自动 prefix 共享 — 多轮对话 / shared system prompt 场景下 prefill 计算被 cache 大量复用。vLLM 也有 prefix caching 但默认开关和粒度不如 SGLang aggressive。
  2. Overlap scheduling — CPU plan 完全 overlap 在 GPU forward 里。vLLM 的 v1 scheduler 在做类似优化但尚未完全 catch up。
  3. Structured generation 的 outlines 编译 — JSON / regex 约束生成时把 FSM 编译成 GPU mask 而不是 CPU rejection sampling。
  4. 27 个 attention backend 自动选最优 — SGLang 几乎是即时跟进每个新 attention kernel 的 SOTA 实现,用户不用自己选。

注意:vLLM 在 v0.6+ 的 v1 架构重写后差距迅速缩小。SGLang 优势主要在 multi-turn dialog 和结构化生成场景。

§ AMD · for JhinWhere to start touching.

基于这次精读,给你 AMD 工作的具体 actionable 起手点。

Workstream 1 — AITER kernel 调试

你 CLAUDE.md 里描述的 agent team:"Kernel Debugger" / "SGLang Integrator" / "Arch Tracer"。具体到代码位置:

TeammateOwnsModify · Don't
Kernel Debuggeraiter/csrc/py_itfs_ck/(aiter 仓库)Modify CK kernel; don't touch SGLang
SGLang Integratorpython/sglang/srt/layers/attention/aiter_backend.pyModify backend dispatch + indices updater; don't touch kernel
Arch Tracerpython/sglang/srt/models/(target 模型文件)Trace shape at crash; don't modify kernel or backend

Workstream 2 — Grid search for MI300X / MI355X

当你跑 Kimi-K2.5 / Qwen3-Coder-Next 的 throughput 实验时:

  • Tier 1 attention backend sweep: --attention-backend{aiter, triton, fa3}。固定 TP=8,看哪个吞吐高。
  • Tier 2 parallelism: 在最优 backend 下扫 TP / EP / DP(EP=2/4/8)。
  • Tier 3 CUDA Graph + mem: --enable-torch-compile / --enable-cuda-graph / --mem-fraction-static

每次实验都用 SGLang 自带的 bench_serving.py(你已经在用)。结果存进性能数据库 — 这正是你 CLAUDE.md 提到的"performance database is the moat"。

Workstream 3 — 长期:把你的 kernel 做成独立 wheel

如果 multi-agent kernel optimization 系统产出了若干"AMD 上更快的 kernel",仿照 sgl-kernel/ 把它们做成 amd-kernels wheel:

  • 独立 pyproject.toml,only ROCm platform tag
  • CMake + hipcc 构建管线
  • 暴露 Python API 给 SGLang / vLLM / 任何 PyTorch 程序直接调用
  • 独立 release 节奏 — kernel 改进不用等 SGLang 主包发版

这是把你的工作 leverage 出 AMD 团队、影响整个开源 ecosystem 的最直接路径。SGLang 的 sgl-kernel、vLLM 的 vllm-flash-attn 都是这种独立 kernel 包,已经形成工业惯例。

§

§ EpilogueWhat we found, and what's next.

SGLang 不是一个 fast LLM server — 它是一个用四进程 + ZMQ 总线 + 27 个 attention backend + RadixAttention 共同实现的 inference orchestration system。它的工程价值在三个层面:(1) 多进程 IPC 的精巧解耦,(2) 跨硬件 attention 多态的优雅抽象,(3) 把"用户 prompt prefix"这种应用层语义带进系统层做共享的洞察。

下一份 File № 003 是 vLLM 的对应分析 — 它走的是不同路径(PagedAttention 而非 RadixCache、v1 重构后的 EngineCore vs Worker 拆分)。两份合起来看才能理解"现代 LLM serving 的设计空间"。

— Fin.