Faster kernels, slower engine更快的核， 更慢的引擎

The result that did not fit对不上的那个结果

The kernel trace and the serving benchmark point in opposite directions. The useful result is not “the kernels explain the engine”; it is that they do not.

Kernel trace 与 serving benchmark 指向相反方向。 真正有用的结论不是「核函数解释了引擎」， 而是它们解释不了。

TokenSpeed is LightSeek's MIT-licensed inference engine. Its AMD path is concrete: an AMD image, a gfx950 Gluon MLA kernel, a Gluon SiTU MoE path, and per-commit MI35x CI. We already had SGLang traces for the same checkpoint on the same node, so both stacks could be measured on one day with one attribution method.

TokenSpeed 是 LightSeek 以 MIT 协议开源的推理引擎。 它的 AMD 路径不是口号： 有 AMD 镜像、 gfx950 Gluon MLA 核、 Gluon SiTU MoE 路径， 以及跑在 MI35x 上的 per-commit CI。 我们已有同一 checkpoint、 同一节点上的 SGLang trace， 因而能在同一天用同一套归因方法比较两套栈。

Four operation-level matched groups are faster on TokenSpeed; routed-expert GEMM is the one loss, and it is confounded by EP8/a16w4 versus TP8/a8w4. The aggregate non-collective device time is still lower and the total launch count is still smaller. Yet graph-on batch-1 serving is slower.

四组按功能匹配的算子在 TokenSpeed 上更快； routed-expert GEMM 是唯一输的一组， 但它混入了 EP8/a16w4 对 TP8/a8w4 的结构差异。 即便如此， TokenSpeed 的非集合通信 device 总时间仍更低、 总 launch 数仍更少， 而开 graph 的 batch-1 serving 却更慢。

What this rules out is a deficit in summed kernel work or an excess launch count. It does not rule out higher per-launch host cost, synchronization, different graph replay behavior, collective latency or poor overlap. Because the two time totals come from different execution modes, they cannot be subtracted into a measured “gap between kernels.” The host/runtime path is the leading hypothesis and the next profiling target.

这些数字排除的是「核函数工作量合计更大」和「launch 数量更多」， 排除不了单次 launch 的 host 成本、 同步、 graph replay 行为差异、 collective 延迟或 overlap 不足。 两组时间来自不同执行模式， 不能直接相减成一个已经测出的「核函数之间的空隙」。 Host/runtime 路径是首要假设， 也是下一步 profiling 目标。

Two stacks, one inverted flag两套栈， 一个反的 flag

One node, 8× AMD Instinct MI355X (gfx950, 288 GiB each), ROCm 7.2. The model is moonshotai/Kimi-K3: 2.78 T total parameters, 105.4 B active, 93 layers, 896 routed experts at top-16 plus 2 shared, situ activation, routed experts in compressed-tensors MXFP4 group-32 and everything else BF16. TokenSpeed's pinned config defines the layer and expert counts and the 69 KDA / 24 NoPE-MLA split.

一台机器， 8 张 AMD Instinct MI355X（gfx950， 每张 288 GiB）， ROCm 7.2。 模型是 moonshotai/Kimi-K3： 2.78 T 总参数、 105.4 B 激活、 93 层、 896 个 routed expert 走 top-16 外加 2 个 shared、 situ 激活； routed expert 以 compressed-tensors MXFP4 group-32 存储， 其余为 BF16。 TokenSpeed 固定版本的配置直接定义了 层数与 expert 数， 以及 69 层 KDA / 24 层 NoPE-MLA 的划分。

Each stack runs its own recommended configuration. That makes this best-config against best-config rather than a controlled A/B, which is the honest way to answer "which should I deploy" but means some differences are structural rather than incidental.

两套栈各跑各自推荐的配置。 所以这是 best-config 对 best-config， 而不是受控 A/B—— 这是回答「我该部署哪个」时最诚实的方式， 但也意味着有些差异是结构性的而非偶然的。

Three server boots, not one universal configuration三次 server boot， 不是一条万能配置

The serving table combines three deliberate profiles. Increasing --max-num-seqs reserves more graph state and reduced the reported token pool from 4.47 M to 3.56 M; the 8K-input workload also needs a 16K model limit because input plus output is 9,216 tokens. The exact invocations are parameterized in ts-bench.sh.

Serving 表格合并了三种有意区分的 profile。 提高 --max-num-seqs 会预留更多 graph 状态， 引擎报告的 token pool 因而从 447 万降到 356 万； 8K 输入负载还需要 16K model limit， 因为输入加输出一共 9,216 token。 精确参数由 ts-bench.sh 注入。

The fp8 KV cache deserves a note because it cuts the wrong way for the conclusion: fp8 halves the bytes moved per cached token relative to bf16, so it is a handicap given to TokenSpeed. Any deficit we find survives despite that advantage.

fp8 KV cache 值得单独说一句， 因为它对结论是反向的： 相对 bf16， fp8 把每个缓存 token 搬运的字节数砍半， 所以这是*送给* TokenSpeed 的优势。 我们找到的任何劣势都是在这个优势之上仍然存在的。

Four pitfalls to a running server到跑起来之间的四个坑

None of these are exotic, all four are load-bearing, and the documented recipe walks into three of them. They are worth recording because each one is the kind of failure that costs an afternoon and leaves no trace in a benchmark table.

没有一个是稀奇古怪的问题， 四个都绕不过去， 而官方文档的 recipe 会踩中其中三个。 值得记下来， 是因为每一个都是那种「浪费一个下午、 却在 benchmark 表格里不留痕迹」的失败。

2.1 · A pinned dependency that was never published2.1 · 一个从未发布过的依赖被硬钉住

The pinned runtime dependency list requires tokenspeed-mooncake>=0.3.12.post20260725. That package had no resolvable release in this environment, so a normal install stopped at from versions: none. Mooncake serves P/D KV transfer and is not used by this aggregated single-node run with --disable-kvstore. We installed the remaining dependencies, omitted that one entry, then installed the runtime with --no-deps; the exact workaround is in ts-build-step6.sh.

固定版本的 runtime 依赖列表要求 tokenspeed-mooncake>=0.3.12.post20260725。 这个包在当时的环境里没有可解析的 release， 普通安装会停在 from versions: none。 Mooncake 服务于 P/D KV 传输， 而这次带 --disable-kvstore 的单机聚合式运行不会使用它。 我们安装其余依赖、 跳过这一项， 再以 --no-deps 安装 runtime； 完整 workaround 在 ts-build-step6.sh。

2.2 · xgrammar silently replaces the ROCm Triton2.2 · xgrammar 悄悄换掉了 ROCm 版 Triton

This is the subtle one, and it fails silently. torch 2.11.0+rocm7.2 requires triton-rocm==3.6.0. xgrammar declares a bare triton dependency. Both distributions install into the same triton/ directory, so pip installs stock triton 3.7.1 from PyPI on top of the ROCm build and whichever landed last wins:

这个坑很隐蔽， 而且是静默失败。 torch 2.11.0+rocm7.2 要求 triton-rocm==3.6.0， 而 xgrammar 声明了一个裸的 triton 依赖。 两个发行包装进的是*同一个* triton/ 目录， 于是 pip 把 PyPI 的 triton 3.7.1 装到 ROCm 版之上， 谁后落地谁赢：

$ python -c "import triton; print(triton.__version__)"
3.7.1          # stock PyPI build — not the ROCm one torch asked for

Nothing errors. You simply run a Triton that was not built for the platform. The fix is to uninstall triton and reinstall triton-rocm==3.6.0 after the dependency set settles. The Gluon gfx950 code is isolated from that collision because every compiler symbol enters through one namespaced indirection module.

全程不报错， 但实际运行的 Triton 并非为这个平台构建。 解法是等依赖安装完成后卸掉 triton， 再重装 triton-rocm==3.6.0。 Gluon gfx950 代码没有被这次冲突影响， 因为所有 compiler symbol 都经过 同一个独立命名空间的间接层。

# tokenspeed_kernel_amd/_triton.py
import tokenspeed_triton as triton
import tokenspeed_triton.experimental.gluon.language as gl
from tokenspeed_triton.experimental.gluon.language.amd.cdna4 import (
    async_copy as cdna4_async_copy,
)

Because tokenspeed_triton is a distinct top-level module name (version 3.8.10), it can never collide with whatever occupies triton/. Vendoring a compiler under its own namespace costs one file and buys immunity from every dependency resolver that will ever run against this environment.

因为 tokenspeed_triton 是一个独立的顶层模块名（版本 3.8.10）， 它永远不会跟占据 triton/ 的东西撞车。 把一个编译器 vendor 到自己的命名空间下， 代价是一个文件， 换来的是对将来任何一次依赖求解的免疫。

2.3 · The documented AMD recipe OOMs on boot2.3 · 文档里的 AMD recipe 一启动就 OOM

Following the pinned AMD K3 recipe exactly, our server loaded all 96 shards, captured decode graphs, then OOMed during prefill graph capture at 751 seconds. The timing and free-memory values below are observations from our server log, not source constants:

严格照固定版本的 AMD K3 recipe 启动时， 我们的 server 加载完 96 个 shard、 捕获完 decode graph， 随后在第 751 秒的 prefill graph capture 中 OOM。 下面的时间与空闲显存来自我们的 server log， 不是源码常量：

Capturing prefill buckets (bucket=2048 avail_mem=19.92 GB):   2%
Capturing prefill buckets (bucket=1024 avail_mem=14.64 GB):  22%
Capturing prefill buckets (bucket= 704 avail_mem= 4.65 GB):  35%
Capturing prefill buckets (bucket= 512 avail_mem= 0.03 GB):  42%
[FATAL ERROR]: HIP failure: 'out of memory'

After weights and the KV pool there were roughly 20 GB left per GPU, not enough for all 40 prefill buckets. We added --disable-prefill-graph. This is not an invented tuning flag: TokenSpeed's own MI35x perf server command already passes it.

权重和 KV pool 占用后每张卡约剩 20 GB， 不足以捕获全部 40 个 prefill bucket。 我们加了 --disable-prefill-graph。 这不是临时发明的调优项： TokenSpeed 自己的 MI35x perf server 命令本来就带这个 flag。

2.4 · It cannot share a Python environment with SGLang2.4 · 它没法跟 SGLang 共用一个 Python 环境

The same dependency file pins torch==2.11.0 and transformers==5.12.0; the SGLang container uses a custom torch 2.9.1+rocm7.2.0. We therefore built TokenSpeed in its own venv, following the recorded installation order. This means the engine comparison is necessarily cross-environment, even though hardware, checkpoint and measurement day are held fixed.

同一份依赖文件还固定了 torch==2.11.0 与 transformers==5.12.0； SGLang 容器则使用定制的 torch 2.9.1+rocm7.2.0。 因此我们按记录下来的 安装顺序把 TokenSpeed 放进独立 venv。 硬件、 checkpoint 与测量日期相同， 但引擎比较在构造上仍是跨环境的。

After those fixes, each boot took about 11 minutes, dominated by reading 1.5 TB of weights. Profile A reported max_total_num_tokens=4466304; the two max-num-seqs=128 profiles reported 3556992, as listed in §01.

处理完这些问题后， 每次启动约 11 分钟， 主要时间用于读取 1.5 TB 权重。 Profile A 报告 max_total_num_tokens=4466304； 两个 max-num-seqs=128 profile 报告 3556992， 见 §01。

End to end端到端

Exact-length random workloads, temperature 0, ignore_eos, prefix caching off; output tokens/s is aggregated over all 8 GPUs. SGLang appears twice: nospec is the engine-level like-for-like row because TokenSpeed ran without speculation, while DSpark is the best tuned deployment result on this node. TokenSpeed's raw responses are archived as serving JSON.

精确长度 random 负载、 temperature 0、 ignore_eos、 关闭前缀缓存； 输出吞吐是 8 卡聚合 tokens/s。 SGLang 出现两次： TokenSpeed 没开投机， 所以 nospec 才是引擎层面的同类对比； DSpark 则是这台节点上调优后的最佳部署结果。 TokenSpeed 原始响应归档在 serving JSON。

Single stream is respectable. TokenSpeed's TPOT holds at 21.5–22.2 ms from 1 K to 8 K input, essentially flat, while DSpark's advantage decays as its accept length falls from 2.51 to 2.34. Concurrency is where it comes apart: 3.7x behind at both 8 and 32.

单流表现拿得出手。 TokenSpeed 的 TPOT 从 1 K 到 8 K 输入稳在 21.5–22.2 ms， 基本是一条水平线； 而 DSpark 的优势随输入变长衰减， 接受长度从 2.51 掉到 2.34。 崩掉的是并发： 在 8 和 32 上都落后 3.7 倍。

The more useful detail is what that CI file gates: --max-num-seqs 1, --cudagraph-capture-sizes 1, and one reference row, 1: [42, 5.2]. The AMD K3 gate therefore cannot detect a batched-throughput regression. That missing coverage does not explain the gap, but it explains why the gap can remain invisible to CI.

更有用的是那份 CI 实际 gate 了什么： --max-num-seqs 1、 --cudagraph-capture-sizes 1， 以及唯一一行参考值 1: [42, 5.2]。 因此 AMD K3 gate 无法发现批量吞吐回退。 这项覆盖缺口不能解释性能差距， 但能解释为什么差距在 CI 中不可见。

The wedge at 8K × 1288K × 128 上的卡死

At concurrency 128 with 8K prompts, SGLang delivered 890 output tok/s and 8,012 total tok/s. TokenSpeed's archived benchmark result records 1 completion and 127 failures after 3,600 seconds. This is not ordinary saturation: useful progress nearly stops.

在并发 128、 8K prompt 下， SGLang 达到 890 output tok/s、 8,012 total tok/s。 TokenSpeed 归档的 benchmark 结果记录了 3,600 秒内完成 1 个、 失败 127 个请求。 这不是普通饱和： 有效进度几乎停止。

The server log contains 9,256 flat retract ... to unwedge the pool events; the exact message comes from resolveFlatStarvation. The observed loop repeatedly retracts a holder, admits more work, and returns to starvation while instantaneous decode drops as low as 0.02 tok/s. The public JSON proves the end result; the retract count and instantaneous floor are server-log observations.

Server log 中出现了 9,256 次 flat retract ... to unwedge the pool； 这条消息来自 resolveFlatStarvation。 实际循环反复回收 holder、 接纳新工作、 再次进入 starvation， 瞬时 decode 一度降到 0.02 tok/s。 公开 JSON 证明最终结果； retract 次数与瞬时低点则是 server-log 观察。

The proposed fix is a hypothesis: make admission account for all four cache-group demands before accepting a batch, then queue requests that cannot fit. Success means bounded or zero retractions and monotonic progress under 8K × 128—not merely a higher average throughput number.

建议修法仍是假设： admission 在接纳 batch 前同时核算四个 cache group 的需求， 装不下的请求进入队列。 成功标准是 8K × 128 下 retract 次数有界或归零、 请求持续单调完成， 而不只是平均吞吐数字提高。

The documented --max-model-len 8192 is not the explanation: Profile C raised it to 16,384, booted with max_req_input_len=16383, and completed the same 8K/1K shape at concurrency 1. The failure appears only under concurrency in this experiment.

文档中的 --max-model-len 8192 不是解释： Profile C 将它提高到 16,384， 启动后报告 max_req_input_len=16383， 并在并发 1 时完成相同的 8K/1K 请求。 本实验中的失败只在高并发下出现。

How the traces were takentrace 是怎么抓的

Both sides use batch 1, ISL 4096, TP8, no speculative decoding, graphs off, device-kernel duration summed per decode step, and collective durations excluded. The complete capture sequence—not just a prose description—is in ts-profile.sh.

两边都是 batch 1、 ISL 4096、 TP8、 无投机解码、 关 graph， 按 decode step 加总 device-kernel 时长并排除 collective 时长。 完整 capture 顺序不只写在文字里， 还保存在 ts-profile.sh。

Why graph off. Inside a replayed HIP graph the individual kernels are not separately attributable — you get one opaque replay. The SGLang baseline was captured graph-off for the same reason, so the two remain comparable. The cost is that eager wall-clock time is meaningless here; only the summed device time is.

为什么关 graph。 在一次被 replay 的 HIP graph 内部， 单个核函数无法单独归因—— 你只会得到一次不透明的 replay。 SGLang 那份 baseline 也正是出于同样原因用 graph-off 抓的， 所以两者保持可比。 代价是 eager 的墙钟时间在这里没有意义， 只有 device 时间的加总有意义。

Why collective durations are excluded. In eager decode the small all-reduce kernels busy-wait for peers, so recorded duration is dominated by rank launch skew. TokenSpeed's collective rows sum to 67.6 ms/step against 17.1 ms of non-collective work—physically impossible as a wall-time decomposition. SGLang's graph-off baseline has the same issue. Dispatch identity and count remain evidence: TokenSpeed has 187 Iris layer all-reduces plus 2 RCCL calls; SGLang has 187 measured collectives.

为什么排除 collective 时长。 Eager decode 中的小 all-reduce 核会忙等其它 rank， 因此记录时长主要反映 rank launch 偏斜。 TokenSpeed 的 collective 行合计 67.6 ms/step， 对比 17.1 ms 非集合通信工作； 它不可能被当作墙钟分解。 SGLang 的 graph-off baseline 也有同样问题。 Dispatch 身份与次数仍可使用： TokenSpeed 是 187 次 Iris layer all-reduce 加 2 次 RCCL， SGLang 实测 187 次 collective。

How kernels were attributed. By the launching source module, not by substring. iris_stage_one_shot_allreduce_two_gluon_kernel contains “gluon” but is a collective; _stage1_a16w4_situ_warp_gemv contains “stage1” but is MoE. The explicit mapping is in ts_vs_sgl_kernels.py.

核函数如何归因。 按发起它的源码模块， 不按名字子串。 iris_stage_one_shot_allreduce_two_gluon_kernel 名字里有 “gluon”， 但它是 collective； _stage1_a16w4_situ_warp_gemv 名字里有 “stage1”， 但属于 MoE。 显式映射在 ts_vs_sgl_kernels.py。

One kernel genuinely spans blocks. _kimi3_projection_gemv_kernel backs four MoE call sites per layer plus kimi3_qkvfab_projection in each KDA layer. That predicts 4 × 92 + 69 = 437 dispatches, exactly matching the trace; its time is split 368/69. The 93-dispatch vendor GEMM is the once-per-layer attention output projection and is split 69/24.

确实有一个核函数跨 block。 _kimi3_projection_gemv_kernel 支撑每层四个 MoE 调用点， 再加每个 KDA 层中的 kimi3_qkvfab_projection。 这预测出 4 × 92 + 69 = 437 次 dispatch， 与 trace 完全吻合； 时间按 368/69 拆分。 另一个 93 次 dispatch 的 vendor GEMM 是每层一次的 attention output projection， 按 69/24 拆分。

Block by block逐块拆解

The two stacks bucket their per-block "misc" slightly differently, so read the block rows as indicative and the total as solid. §07 compares only the kernel groups whose mapping is unambiguous.

两套栈对每个 block 内部 misc 的归类方式略有差异， 所以 block 那几行当指示性数字看， 合计那行才是硬的。 §07 只比较映射毫无歧义的核函数组。

Glue is the one aggregate block TokenSpeed loses: 2.786 vs 1.290 ms. Its trace rows include RMSNorm, _add3, fp8 copies, Memcpy DtoD and a generic GEMV. SGLang's ROCm path calls AITER's fused residual-add + RMSNorm + fp8 quantization; add_rmsnorm_quant_kernel is the runtime-emitted AITER kernel name, not a symbol defined in the SGLang repository. At 93 layers, avoiding two extra tensor passes is material even when each individual launch is small.

Glue 是 TokenSpeed 唯一输掉的聚合 block： 2.786 对 1.290 ms。 Trace 行包括 RMSNorm、 _add3、 fp8 copy、 Memcpy DtoD 与通用 GEMV。 SGLang 的 ROCm 路径调用 AITER 的 残差加 + RMSNorm + fp8 量化融合算子； add_rmsnorm_quant_kernel 是运行时发出的 AITER kernel 名， 不是 SGLang 仓库中定义的 symbol。 在 93 层上， 即便单次 launch 很小， 少做两趟 tensor 遍历也会形成实质差距。

Gluon vs AITER, op by opGluon 对 AITER， 逐个算子

Four rows are matched by output role and layer count; they are not instruction-for-instruction identical. The fifth—routed-expert GEMM—is intentionally shown because it is important, but it also changes parallel layout and activation precision. Read the first four as operation-level implementation comparisons and the fifth as a strategy comparison.

前四行按输出功能与层数匹配， 并不代表逐条指令完全相同。 第五行 routed-expert GEMM 很重要， 所以仍然列出， 但它还改变了并行布局与激活精度。 前四行可读作算子级实现对比， 第五行只能读作策略对比。

7.1 · Gluon MLA decode — 2.77x, and the cleanest win7.1 · Gluon MLA decode—— 2.77 倍， 最干净的一场胜利

SGLang runs MLA decode as a two-stage split-KV Triton attention: _fwd_grouped_kernel_stage1 scans the KV cache in splits and _fwd_kernel_stage2 reduces them. The trace measures 50.5 µs per MLA layer. TokenSpeed's _mla_decode_gluon plus softmax/reduce-V measures 18.3 µs per layer.

SGLang 的 MLA decode 是两阶段 split-KV Triton attention： _fwd_grouped_kernel_stage1 分片扫描 KV cache， _fwd_kernel_stage2 做归约。 Trace 测得每个 MLA 层 50.5 µs。 TokenSpeed 的 _mla_decode_gluon 加 softmax/reduce-V 为每层 18.3 µs。

Why the gap is real and not a tuning artifact: at batch 1 with a 4 K context, decode attention has an arithmetic intensity near one — it is a pure streaming read of the KV cache, and the only thing that matters is how close to peak bandwidth the load instructions get. A split-K template must choose its split count without knowing the access pattern, materialise partial results, and reduce them in a second pass over memory. A kernel written against one architecture can size its tiles to the cache line and the LDS budget and skip the second pass. That is a structural advantage of hand-writing for a fixed target, and it is exactly the regime where it pays.

为什么这个差距是真实的而不是调参产物： 在 batch 1、 4 K 上下文下， decode attention 的算术强度接近 1—— 它就是对 KV cache 的一次纯流式读取， 唯一重要的是 load 指令能跑到峰值带宽的多少。 一个 split-K 模板必须在不知道访问模式的情况下选定分片数、 把部分结果物化出来、 再用第二趟遍历内存去归约。 而一个针对单一架构写的核函数可以把 tile 尺寸对齐到 cache line 和 LDS 预算， 并且省掉第二趟。 这是「为固定目标手写」的结构性优势， 而这里正是它最划算的场景。

7.2 · The router — 6.51x, and the largest prize on the board7.2 · Router—— 6.51 倍， 桌上最大的一笔

Selecting top-16 of 896 experts and preparing grouped-GEMM routing costs SGLang 2.96 ms per decode step across trace-emitted AITER kernels: grouped_topk, multi-phase opus_moe_sorting, fused_mx_quant_moe_sort and moe_reduction. TokenSpeed's measured routing row is one 0.45 ms _kimi3_sigmoid_bias_topk_kernel, which performs sigmoid, bias, top-k, normalization and optional logical-to-physical mapping.

从 896 个 expert 中选 top-16 并准备 grouped-GEMM routing， SGLang 每个 decode step 花 2.96 ms， 分散在 trace 发出的 AITER 核上： grouped_topk、 多阶段 opus_moe_sorting、 fused_mx_quant_moe_sort 与 moe_reduction。 TokenSpeed 的 routing 行只有一个 0.45 ms 的 _kimi3_sigmoid_bias_topk_kernel， 在一处完成 sigmoid、 bias、 top-k、 归一化与可选 logical-to-physical 映射。

That single kernel is worth 2.5 ms per decode step, which is more than the entire end-to-end deficit TokenSpeed carries. Nothing about it is AMD-specific or Gluon-specific: it is the observation that when intermediate state is proportional to a large constant — here 896 — the cost of materialising it dominates the cost of producing it, so the boundary between kernels should not fall there. Delete the proper nouns and the idea still stands.

单这一个核函数就值 2.5 ms/step， 比 TokenSpeed 背着的整个端到端赤字还大。 它跟 AMD 无关， 跟 Gluon 也无关： 它讲的是， 当中间状态正比于一个很大的常数时—— 这里是 896—— 把它物化出来的成本会压过产生它的成本， 所以核函数的边界不该切在那里。 把专有名词删掉， 这个道理照样成立。

7.3 · Where AITER wins, and why the row is confounded7.3 · AITER 赢在哪里， 以及这一行为什么被污染

SGLang's trace-emitted mfma_moe1/mfma_moe2 MXFP4 kernels run the routed experts in 1.44 ms, against 2.39 ms for TokenSpeed's Gluon stage 1 and stage 2. Two caveats make this the weakest row in the table:

SGLang trace 发出的 mfma_moe1 / mfma_moe2 MXFP4 核跑 routed expert 用 1.44 ms； TokenSpeed 的 Gluon stage 1 与 stage 2 用 2.39 ms。 两个前提让这成为表里最弱的一行：

• Different parallelism, different shapes. TokenSpeed runs EP8, so each rank owns 112 whole experts and at batch 1 touches only the few its tokens route to. SGLang runs pure TP8, where every rank runs all 16 experts at one-eighth width. Those are not the same GEMM, so this compares two strategies as much as two kernels.
• 并行方式不同， shape 就不同。 TokenSpeed 跑 EP8， 每个 rank 拥有 112 个完整 expert， batch 1 时只碰它的 token 路由到的那几个。 SGLang 跑纯 TP8， 每个 rank 都跑全部 16 个 expert， 但只做 1/8 的宽度。 这不是同一个 GEMM， 所以这既是在比两个核函数， 也是在比两种策略。
• Different activation precision. AITER's path is a8w4 — 8-bit activations — against Gluon's a16w4. AITER is moving half the activation bytes, which at batch 1 is most of what the kernel does.
• 激活精度不同。 AITER 那条路是 a8w4—— 8-bit 激活—— 而 Gluon 是 a16w4。 AITER 搬的激活字节数只有一半， 而在 batch 1 下这几乎就是这个核函数做的全部事情。

The honest reading is not "AITER's GEMM is better" but "8-bit activations plus TP sharding beat 16-bit activations plus EP sharding at batch 1 on this shape". That is a useful finding, and it points at quantising the activations rather than rewriting the kernel.

诚实的读法不是「AITER 的 GEMM 更好」， 而是「在 batch 1、 这个 shape 下， 8-bit 激活加 TP 切分胜过 16-bit 激活加 EP 切分」。 这仍然是个有用的结论， 而且它指向的是去量化激活， 而不是重写核函数。

7.4 · KDA is a wash, and both are Triton7.4 · KDA 打平， 而且两边都是 Triton

TokenSpeed's batch-1 KDA recurrence dispatches to a Triton implementation, like SGLang's path. The measured recurrence + short-conv rows land within 14%. Our occupancy estimate is 48 active workgroups on 256 CUs for this fixed-size batch-1 update; that is an analysis of the launch shape, not a source constant. Batching is the direct experiment for testing whether exposed parallelism—not kernel language—is the limit.

TokenSpeed 的 batch-1 KDA recurrence dispatch 到 Triton 实现， 与 SGLang 路径相同。 实测 recurrence + short-conv 行相差在 14% 内。 我们根据 launch shape 估算， 这次固定大小的 batch-1 update 在 256 个 CU 上只有 48 个活跃 workgroup； 这是分析结果， 不是源码常量。 要验证限制来自暴露的并行度而不是 kernel language， 直接实验是扩大 batch。

What the evidence does—and does not—locate证据定位了什么、 没定位什么

Concurrency strengthens the prioritization but not the causal claim. At concurrency 8 and 32, TokenSpeed is about 3.7× behind SGLang nospec. That shape is consistent with scheduler, batching, graph-shape or overlap inefficiency; it is not evidence that one fixed per-step tax simply grew by 3.7×. The FlatKV failure at 8K × 128 provides a concrete scheduler bug, while the smaller-workload gap still needs localization.

并发结果加强了优先级判断， 但没有加强到因果结论。 在并发 8 与 32 上， TokenSpeed 约落后 SGLang nospec 3.7 倍。 这个形状与 scheduler、 batching、 graph shape 或 overlap 效率问题都一致； 它不能证明某一笔固定 per-step 税简单放大了 3.7 倍。 8K × 128 的 FlatKV 失败提供了一个具体 scheduler bug， 较小负载上的差距仍需定位。

Three optimization tracks三个优化方向

These are ordered by evidence and blast radius. Each is a testable engineering loop: observed signal → concrete code boundary → success criterion.

下面按证据强度与影响范围排序。 每一项都是可验证的工程闭环： 观察信号 → 具体代码边界 → 成功标准。

Shared guardrail. Extend AMD K3 CI beyond concurrency 1: gate at least conc 1, 8 and 32 for 4K/1K, plus a bounded-progress 8K/1K admission test. CI coverage is not an optimization; it is what keeps each optimization from silently regressing.

共同护栏。 AMD K3 CI 不能只测并发 1： 至少 gate 4K/1K 的并发 1、 8、 32， 再加一项 8K/1K 的有界进度 admission 测试。 CI 覆盖本身不是优化， 但它保证每项优化不会静默回退。

Deployment decision for this measured setup对这组实测配置的部署判断

Choose SGLang for concurrent serving on this exact 8× MI355X setup and these workloads. TokenSpeed is close enough at single stream to be worth fixing, but its measured batched throughput and 8K admission behavior are not production-ready. This conclusion is scoped to the pinned commits, flags and workloads above.

在这套 8× MI355X 配置和这些负载上， 并发 serving 选 SGLang。 TokenSpeed 的单流表现足够接近， 值得继续修， 但实测批量吞吐与 8K admission 行为尚不适合生产。 这个判断只适用于上面固定的 commit、 flag 与 workload。

Exactly how we served Kimi-K3 with TokenSpeed我们实际如何用 TokenSpeed 起 Kimi-K3

This is the operational recipe used for every TokenSpeed number above: pinned source, isolated ROCm environment, one base serve command, and three explicit capacity profiles.

上面每一个 TokenSpeed 数字都来自这套操作路径： 固定源码、 独立 ROCm 环境、 一条基础 serve 命令， 再明确区分三种容量 profile。

10.1 · Build the pinned source in an isolated venv10.1 · 在独立 venv 中构建固定版本

We used TokenSpeed d50bb481. Package order matters: the in-tree AMD package precedes the generic kernel build; the scheduler is compiled with FlatKV; runtime dependencies omit unavailable Mooncake; ROCm Triton is restored last. The full executable records are ts-build.sh and ts-build-step6.sh.

我们使用 TokenSpeed d50bb481。 安装顺序不能乱： 先装仓库内 AMD 包， 再编译通用 kernel； scheduler 打开 FlatKV； runtime 依赖跳过不可用的 Mooncake； 最后恢复 ROCm Triton。 完整可执行记录是 ts-build.sh 与 ts-build-step6.sh。

TS_ROOT=/sgl-workspace/tokenspeed
python3 -m venv "$TS_ROOT/.venv"
source "$TS_ROOT/.venv/bin/activate"
pip install --upgrade pip "setuptools<82" wheel cmake ninja
pip install --index-url https://download.pytorch.org/whl/rocm7.2 \
  torch==2.11.0+rocm7.2

pip install --force-reinstall --no-deps "$TS_ROOT/tokenspeed-kernel-amd" \
  --no-build-isolation
TOKENSPEED_KERNEL_BACKEND=rocm pip install \
  "$TS_ROOT/tokenspeed-kernel/python/" --no-build-isolation
pip install "$TS_ROOT/tokenspeed-scheduler/" \
  --config-settings=cmake.define.TOKENSPEED_FLAT_KVCACHE=ON

# Install pyproject dependencies except torch and tokenspeed-mooncake, then:
pip install -e "$TS_ROOT/python" --no-build-isolation --no-deps
pip uninstall -y triton
pip install --force-reinstall --no-deps triton-rocm==3.6.0 \
  --index-url https://download.pytorch.org/whl/rocm7.2

python -c "import torch,tokenspeed_scheduler; \
assert torch.cuda.device_count()==8 and tokenspeed_scheduler.FLAT_KVCACHE"

10.2 · Set the actual model environment10.2 · 设置实际模型环境

The node already held the Hugging Face checkpoint and a flattened tokenizer path. Runs were offline and isolated from SGLang's Python environment.

节点上已经有 Hugging Face checkpoint 与 flatten 后的 tokenizer 路径。 所有运行均为 offline， 并与 SGLang Python 环境隔离。

source /sgl-workspace/tokenspeed/.venv/bin/activate
export HF_HOME=/sgl-workspace/models
export HF_HUB_OFFLINE=1
export HF_MODULES_CACHE=/sgl-workspace/workspace/.ts_hf_modules
MODEL=moonshotai/Kimi-K3
TOKENIZER=/sgl-workspace/models/Kimi-K3-flat
PORT=8100

10.3 · Start the base server10.3 · 启动基础 server

MAXLEN and MAXSEQS are the only profile variables. Everything else stays fixed. The model's K3 notes explain FlatKV, fp8 KV and MoonViT data mode; the pinned AMD recipe supplies the TP8/EP8 backend choices. Our two deviations are highlighted.

MAXLEN 与 MAXSEQS 是仅有的 profile 变量， 其余保持不变。 模型的 K3 说明解释 FlatKV、 fp8 KV 与 MoonViT data mode； 固定版本的 AMD recipe提供 TP8/EP8 backend 选择。 两处偏离已高亮。

MAXLEN=8192    # profile-specific; see 10.4
MAXSEQS=32     # profile-specific; see 10.4

tokenspeed serve "$MODEL" \
  --served-model-name kimi-k3 --trust-remote-code \
  --max-model-len "$MAXLEN" --max-num-seqs "$MAXSEQS" \
  --kv-cache-dtype fp8 --gpu-memory-utilization 0.92 \
  --tensor-parallel-size 8 --mm-encoder-tp-mode data \
  --enable-expert-parallel --attention-backend mla --moe-backend auto \
  --disable-kvstore \
  --disable-prefill-graph \
  --no-enable-prefix-caching \
  --host 127.0.0.1 --port "$PORT"

10.4 · Run the three profiles actually measured10.4 · 运行实际测量的三个 profile

The harness takes the shared-node lock, starts the server above, waits on /readiness, sends an OpenAI-compatible chat smoke test, runs the listed workloads, archives JSON, then stops the server. These are the exact profile inputs:

Harness 会取得共享节点锁、 启动上述 server、 等待 /readiness、 发送一次 OpenAI-compatible chat smoke test、 运行指定负载、 归档 JSON， 最后停止 server。 三组实际输入如下：

HARNESS=/sgl-workspace/jinnpan.com/data/tokenspeed-vs-sglang-kimi-k3/scripts/ts-bench.sh

# A · 1K/4K serving + the base configuration used for traces
MAXLEN=8192 MAXSEQS=32 \
WORKLOADS="4096:1024:1:3:1 1024:1024:1:4:1 1024:1024:8:16:2 1024:1024:32:64:4" \
  bash "$HARNESS"

# B · 1K high concurrency
MAXLEN=8192 MAXSEQS=128 \
WORKLOADS="1024:1024:32:64:4 1024:1024:64:64:4 1024:1024:128:128:8" \
  bash "$HARNESS"

# C · 8K input + 1K output requires a limit above 9,216 tokens
MAXLEN=16384 MAXSEQS=128 \
WORKLOADS="8192:1024:1:2:1 8192:1024:128:128:8" \
  bash "$HARNESS"

10.5 · Benchmark and profile semantics10.5 · Benchmark 与 profile 语义

Serving results use normal graph replay. Exact lengths require TokenSpeed ratio 0. Kernel attribution is a separate restart with --enforce-eager; it is not the serving configuration.

Serving 结果使用正常 graph replay。 精确长度必须传 TokenSpeed ratio 0。 Kernel 归因会另行带 --enforce-eager 重启； 它不是 serving 配置。

# Serving benchmark: graph on, exact 4K/1K, batch 1
tokenspeed bench serve --base-url http://127.0.0.1:8100 --model kimi-k3 \
  --tokenizer /sgl-workspace/models/Kimi-K3-flat --dataset-name random \
  --input-len 4096 --output-len 1024 --random-range-ratio 0 \
  --num-prompts 3 --max-concurrency 1 --num-warmups 1 \
  --ignore-eos --extra-body '{"temperature": 0}'

# Attribution trace: restart Profile A with --enforce-eager, warm up, then:
curl -sS -X POST http://127.0.0.1:8100/start_profile \
  -H 'Content-Type: application/json' \
  -d '{"num_steps":16,"activities":["GPU"],
       "profile_by_stage":true,"profile_id":"kimi-k3-tp8"}'

The exact profiling restart and warmup are in ts-profile.sh. Trace conversion is in ts_trace_report.py; attribution and comparison are in ts_vs_sgl_kernels.py. Derived CSVs and serving JSON are committed in the experiment archive; raw per-rank traces and server logs are not.

精确的 profiling 重启与 warmup 在 ts-profile.sh； trace 转换在 ts_trace_report.py； 归因与对比在 ts_vs_sgl_kernels.py。 派生 CSV 与 serving JSON 已提交到 实验归档； 原始 per-rank trace 与 server log 未提交。