§ 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 包的多平台构建系统。
§ 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/ 分两大块:
| Subpackage | What |
|---|---|
sglang/lang/ | 前端 DSL — Structured Generation Language(select/gen/fork/regex 等 primitives)。这是 SGLang 名字的来源。 |
sglang/srt/ | Serving Runtime — 推理引擎本体,下面所有模块都属于它 |
sglang/srt/ 进一步切分(这是真正的"系统"):
| Subdir | Role |
|---|---|
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 |
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.py 的 Engine 类(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 大概结构):
_set_envs_and_config(line 1159)— 设环境变量(CUDA_VISIBLE_DEVICES、ROCm 等),GC 调优init_tokenizer_manager(line 126)— 起 TokenizerManager 进程_compute_parallelism_ranks(line 1332)— 算 TP × DP × EP 拓扑下的 rank 分配- spawn N 个 scheduler 子进程(multiprocessing),每个进程跑一个
Scheduler.event_loop() _wait_for_scheduler_ready(line 1265)— 等所有 scheduler 通过 ZMQ ready barrier- HTTP / gRPC server 接管(如果是 server 模式)
关键细节:每个 TP rank 是一个独立 scheduler 进程(不是 thread),用 multiprocessing spawn。Tokenizer 是单独一个进程,Detokenizer 又是单独一个。所以一个 TP=8 的部署有 1 + 1 + 8 = 10 个 Python 进程 共同协作。
不要把 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 → To | What | Mode |
|---|---|---|
| HTTP client → TokenizerManager | HTTP request | FastAPI / uvicorn |
| TokenizerManager → Scheduler | TokenizedGenerateReqInput | ZMQ PUSH |
| Scheduler → Worker(s) | ScheduleBatch (tensor IPC) | ZMQ + shared memory |
| Worker → Scheduler | output tokens / logits | ZMQ PUSH |
| Scheduler → DetokenizerManager | token_ids batch | ZMQ PUSH |
| DetokenizerManager → TokenizerManager | decoded text chunks | ZMQ PUSH |
注意是环形拓扑:Tokenizer → Scheduler → Worker → Scheduler → Detokenizer → Tokenizer。Tokenizer 既是入口也是出口,负责把 streamed 输出聚合后通过 HTTP 推给客户端。
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 等几乎所有子系统。
核心方法(按调用频率排):
| Method | Line | Frequency | What |
|---|---|---|---|
event_loop_normal() | 1524 | per server | 普通调度循环 — sequential |
event_loop_overlap() | 1551 | per server | ★ overlap 调度 — 计算/调度并发 |
process_input_requests | 1829 | per step | 从 ZMQ 收新请求 |
get_next_batch_to_run | 2461 | per step | 挑下一个要跑的 batch |
update_running_batch | 2870 | per 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
"普通"循环是 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 抽象。
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 抽象基类):
| Operation | What |
|---|---|
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,给极端性能场景用。
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。它负责的事:
- 从 HF 或自定义 loader 加载模型权重
- 分配 KV cache memory pool(巨大 tensor,按 layer × block × head 切片)
- 选择并初始化 attention backend(从 27 个候选里挑)
- build / capture CUDA Graph(line 2339-2347 的关键分支)
- 每次
forward(batch)时把 ScheduleBatch 翻译成 ForwardBatch(含位置编码、cache indices、mask 等),调self.model.forward(...) - 给 attention backend 调
init_forward_metadata(forward_batch)准备元数据(page table / sliding window mask) - 对接 sampling(top_k / top_p / 结构化生成)
CUDA Graph 这一段值得单独说 — SGLang 提供多种 capture 策略:
cuda_graph_runner.py— 普通整张图 capturepiecewise_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 平台特化
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"映射。
| Family | Backends | Hardware |
|---|---|---|
| FlashAttention 系 | flashattention_backend, flashinfer_backend, flashinfer_mla_backend, flashmla_backend | NV (CUDA) |
| Triton | triton_backend, triton_ops/ | NV + AMD(Triton 跨平台) |
| CUTLASS | cutlass_mla_backend | NV (Hopper / Blackwell) |
| TRT-LLM | trtllm_mha_backend, trtllm_mla_backend | NV (TensorRT-LLM) |
| Torch native | torch_native_backend, torch_flex_backend | any (slow fallback) |
| AMD AITER | aiter_backend | AMD ROCm (MI300X+) |
| Intel | intel_amx_backend | Intel CPU + AMX |
| Mamba / State space | mamba/, linear/ | any |
| NSA (sparse) | nsa_backend, nsa/ | NV |
| DSv4 (DeepSeek) | deepseek_v4_backend, dsv4/ | NV + AMD |
| Hybrid | hybrid_attn_backend, hybrid_linear_attn_backend, dual_chunk_flashattention_backend | any |
| TBO | tbo_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 优化
你在 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 的预期。
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 target | pyproject | Tooling |
|---|---|---|
| CUDA wheel | pyproject.toml(默认) | CMake + nvcc + ninja |
| ROCm wheel | pyproject_rocm.toml | CMake + hipcc + composable_kernel |
| CPU wheel | pyproject_cpu.toml | CMake + cc |
| MUSA wheel | pyproject_musa.toml | CMake + Moore Threads tooling |
| Metal binary | setup_metal.py | Xcode / 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
把 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 时最容易栽的几个坑。每个坑都对应一类具体的失误模式。
不是 multi-thread。你 monkey-patch Scheduler.event_loop_normal 在主进程是无效的 — 每个 TP rank 是一个独立 multiprocessing.Process,它们 spawn 时从头 import 整个 sglang 包,之前在主进程注入的 hook 完全丢失。要影响所有 scheduler,hook 必须放在 sglang.srt.__init__ 或 module-level,import 时就生效。
event_loop_overlap 用 last_batch.future 异步等 GPU 结果。如果你改 schedule 逻辑导致某次 future 永不 resolve(比如 batch 被错误地 drop),整个进程会卡死。建议任何改 overlap 路径的人先打开 SGLANG_DEBUG=1 看 future state。
RadixCache 每个匹配到 prefix 的请求要 inc_lock_ref,请求结束要 dec_lock_ref。如果一个 request 在 setup 时崩了,没人帮你 dec,那段 KV slots 永远不会被驱逐,最终 OOM。所有 catch 异常分支都要小心 ref 是否清理。cache_controller.py 有 helper 但不强制。
不同 backend 对 forward_batch 的元数据要求不一样 — 有的要 seq_lens_cpu,有的要 extend_seq_lens_int32,有的要 page-aligned tensor。init_forward_metadata 是每个 backend 自己实现的,遗漏一个字段就会在 kernel 里炸(且报错信息常常很神秘)。切 backend 前先看目标 backend 的 init_forward_metadata 签名和断言。
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.
sky/srt/layers/attention/<new>_backend.py— 实现AttentionBackend抽象(init_forward_metadata/forward_extend/forward_decode/forward_target_verify)sky/srt/layers/attention/attention_registry.py— 注册名字 → class 映射sky/srt/server_args.py— 加入--attention-backend <new>的选项校验sky/srt/model_executor/model_runner.py的_get_attention_backend()方法 — 必要的话加 fallback 逻辑- (推荐)
test/srt/test_attention_backend.py— 加 smoke test 跑短 prefill + decode - (如果是 ROCm-specific)
sgl-kernel/里加配套的低层 kernel
关键参考:aiter_backend.py 是最近也是最大型的"加新 backend"实例,可以照抄结构。
进程: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 张量。
四个具体来源:
- RadixAttention 自动 prefix 共享 — 多轮对话 / shared system prompt 场景下 prefill 计算被 cache 大量复用。vLLM 也有 prefix caching 但默认开关和粒度不如 SGLang aggressive。
- Overlap scheduling — CPU plan 完全 overlap 在 GPU forward 里。vLLM 的 v1 scheduler 在做类似优化但尚未完全 catch up。
- Structured generation 的 outlines 编译 — JSON / regex 约束生成时把 FSM 编译成 GPU mask 而不是 CPU rejection sampling。
- 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"。具体到代码位置:
| Teammate | Owns | Modify · Don't |
|---|---|---|
| Kernel Debugger | aiter/csrc/py_itfs_ck/(aiter 仓库) | Modify CK kernel; don't touch SGLang |
| SGLang Integrator | python/sglang/srt/layers/attention/aiter_backend.py | Modify backend dispatch + indices updater; don't touch kernel |
| Arch Tracer | python/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.