返回
研究
2026年10月20日
加利福尼亚州旧金山
Zyphra 推理第二波：在 MI355X 上实现更快的 MoE 与长上下文服务Skyler Szot, Srivatsan Rajagopal, Abhinav Sarje, Varun Subramanian, Bhavana Mehta, Praneeth Medepalli, Quentin Anthony
Zyphra 推出面向 AMD Instinct MI355X GPU 的 Zyphra 优化引擎（ZOE）第二波推理优化。我们实测了 MI355X 的计算与内存带宽极限，在我们生产分支的 vLLM 和 SGLang 中为混合专家层和稀疏注意力编写了自定义 HIP kernel，并评估了 MXFP4 量化和投机解码。ZOE 将 GLM-5.3 在 512K token 上下文下的首 token 时间在无前缀缓存命中时加速 2.8 倍，在 90% 命中率时加速 3.6 倍，并将 DeepSeek-V4-Pro 在 1M token 上下文下的单用户解码吞吐提升 1.40 倍。
在前沿 MoE 模型中，大量推理时间花在几乎用不到 GPU 矩阵核心吞吐的操作上。这些操作包括行数不均的小型专家 GEMM、从高带宽内存（HBM）读取专家权重、稀疏注意力索引，以及长上下文下的 top-k 选择。对于 512K token 上下文下的 GLM-5.3，我们融合的 MoE 和 DSA 索引器 kernel 替代了占据大部分预填充时间的操作，从而将 TTFT 加速 2.8 倍至 3.6 倍（图 1）。
为了进一步提升我们用户的 Zyphra Cloud 性能，我们直接针对这些瓶颈。我们实测了 MI355X 的实际计算与带宽极限，为 MoE 和稀疏注意力编写了自定义 HIP kernel，并将它们集成到我们生产分支的 vLLM 和 SGLang 中。我们已将这些优化加入我们的核心 Zyphra 优化引擎（ZOE） ，它支撑着我们的推理云，也可本地部署。
这些结果建立在我们 第一轮 Zyphra 推理优化的基础之上，那一轮聚焦于调优、并行和通信。在第二阶段，我们深入模型前向传播。我们实测 MI355X 的经验 roofline，为最昂贵的操作（尤其是 MoE 前馈块和稀疏注意力（CSA/DSA））编写自定义 HIP kernel，并评估这些 kernel 如何与我们分支版本的 vLLM 和 SGLang. 本博客中的所有工作都针对 MI355X（gfx950 / CDNA4）。我们在本博客中的目标是：
a.) 展示我们使用 ZOE 在推理云中可交付的性能
b.) 分享我们的优化研究进展
c.) 预告我们即将开展的部分推理工作
3.6x3.6x3.6x更快的 TTFT
GLM-5.3 / 512K 上下文 / 90% 前缀缓存 / 对比 MI355X 上的 vLLM v0.29.0
关键结果
预填充： ZOE 将 GLM-5.3 在 512K token 下的首 token 时间（TTFT）从 178.8 秒降至 64.6 秒（无前缀缓存命中），即 2.8 倍 加速[^1]。在 90% 命中率下，它将 TTFT 从 42.3 秒降至 11.8 秒，即 3.6 倍 加速。这些结果使用 vLLM v0.29.0 及我们的 DSA lightning 索引器和融合 MoE kernel。
解码： 我们的 CSA top-K 选择 HIP kernel 将 DeepSeek-V4-Pro 在 1M token 上下文下的端到端单用户解码吞吐从 43.5 提升至 61.1 tokens/s（SGLang v0.5.19），即 1.40 倍 加速。
Kernel： 在 1M token 上下文下，我们的 CSA top-K 选择 kernel 对单请求比 SGLang kernel 快 11.1 倍，对 32 请求批次快 9.9 倍 。我们的 lightning 索引器比 AITER（各形状最佳）快 1.15 倍至 2.37 倍 ，吞吐为 AITER 的 Gluon kernel 的 0.73 倍至 1.37 倍 （图 11）。我们的融合 MoE kernel 在 Kimi-K2.6（a16w4）上比 vLLM Triton WNA16 kernel 快至多 2.0 倍 。在 DeepSeek-V3.2（FP8）上，对于最多 256 token 的前向传播，它比 AITER 快至多 1.17 倍 。
量化： GLM-5.3 的 AMD MXFP4 量化将模型大小从 756 GB 降至 410 GB，其在 AIME25、AIME26、Apex-Shortlist 和 HMMT 上的分数与 FP8 模型相差在 3 分以内。在低并发下，它的解码速度慢于 FP8 模型。请求很少的解码步骤只读取一小部分专家权重，因此将其大小减半节省的时间很少，而 vLLM 中当前的 MXFP4 kernel 增加的时间更多。在 32 个并发请求下，每一步读取更多专家权重，MXFP4 模型比 FP8 模型快 1.09 倍。
硬件： 我们对 MI355X 硬件进行了详细的 roofline 分析，以便 kernel 性能可以对照可实现的硬件极限来衡量，而不仅仅是数据手册标称的 FLOP/s 数字。在我们的微基准测试中，MI355X 达到其数据手册 FP8 MFMA 峰值的 98% 和数据手册 HBM 带宽的 78%。
图 1. ZOE 的最大收益出现在长上下文。 ZOE 相对于官方 vLLM 和 SGLang ROCm 镜像在各上下文长度下的单请求加速。1× 处的虚线为官方镜像。左：vLLM v0.29.0 中 GLM-5.3 在三种前缀缓存命中率下的 TTFT。右：SGLang v0.5.19 中 DeepSeek-V4-Pro 使用我们的 CSA top-K kernel 的单用户解码吞吐。完整结果见图 12 和图 14 及附录。
我们在本文末尾描述了我们的方法及其局限性。简而言之，所有结果均来自单个 MI355X 节点且无实时流量，投机解码结果取决于提示词，且我们的技术栈和基线都在快速变化，因此所有公布的結果均为特定时间点的结果。
本文涵盖的内容
硬件 Roofline: MI355X 在实践中交付的计算和 HBM 带宽，我们将其作为后续每个 kernel 的上限。
优化: 我们的四个新 HIP kernel（融合 MoE、DSA lightning 索引器、CSA top-K 选择和 CSA 双范围解码），以及我们评估的投机解码和量化方法。
Kernel 结果: 每个 kernel 单独测量，与其替代的 AITER、vLLM 和 SGLang kernel 对比。
服务结果: kernel 如何改变端到端预填充 TTFT 和解码吞吐。
量化: MXFP4 GLM-5.3 checkpoint 的准确性和解码吞吐，以及为什么目前它仅在高并发下快于 FP8。
方法与局限性：我们如何测量，以及如何解读这些数字。
引言
许多前沿开放权重模型已收敛到两种塑造推理成本的架构选择。 第一种是细粒度混合专家（MoE）模型 ，这类模型由诸如 DeepSeek-V3等模型推广开来。在 MoE 模型中，并非每个参数都在前向传播中被激活，而是由路由器从 N 个专家中选择 k 个激活。这使得参数数量可以增长，而不增加每 token 的 FLOPs。随着参数数量增长，总参数中更大比例被用于专家 MLP，而非注意力等序列混合器。细粒度 MoE 倾向于增加专家数量并减小专家宽度，从而产生许多更小的专家，而非更少更大的专家。虽然这在建模角度很有吸引力，但对推理而言却很别扭。具体来说，由于现在有许多小专家，细粒度 MoE 提供的可轻松并行化的工作更少[^2]。它们还有多得多的 kernel 需要启动，并且增加了专家间工作负载不均衡的可能性[^3]。因此，在解码阶段，MoE 块可能把大量时间花在从 HBM 搬运专家权重上；而在预填充阶段，它变成一堆独立的 GEMM，各专家的计算量差异很大。
第二种架构选择是 更廉价（长上下文下每 token）的注意力变体，如线性注意力和稀疏注意力。例如，DeepSeek-V3.2 和 GLM-5.1 使用 DeepSeek 稀疏注意力（DSA），其中索引器对过去位置打分，主注意力只读取被选中的 KV 条目。类似地，DeepSeek-V4 使用压缩稀疏注意力（CSA），它对得分最高的压缩位置加上约 128 个近期 token 的局部窗口进行注意力计算。在这两种设计中，注意力步骤对每个查询读取固定数量的 KV 条目（DSA 为 2,048 个，CSA 为 512 个压缩条目加上约 128 个 token 的窗口），因此其每查询成本随上下文增长保持恒定。选择步骤对每个查询为所有过去位置（CSA 中为所有压缩位置）打分。其每查询成本随上下文长度线性增长，而整个提示词上的成本随提示词长度的平方增长。随着上下文增长，选择步骤成为注意力成本中更大的部分。
DeepSeek 稀疏注意力（DSA）
图 2. DSA 将选择与注意力分离。两个面板都逐 token 解码同一个 16 token 序列。最新 token 产生一个查询（q）。注意力对其读取的值向量计算加权和，权重来自查询-键得分，该和即为该 token 的输出（o）。常规注意力（左）读取每个 KV 条目。在 DSA（右）中，索引器先用其自己的查询（qᵢ）和每个位置一个 128 维 FP8 索引键为所有过去位置打分。随后主注意力读取得分最高的 k 个条目。一个索引器得分的 FLOPs 更少，读取的字节也比对同一位置做注意力更少（主注意力使用 576 维键，并且还要对值向量求和）。索引器仍为所有过去位置打分，因此其每 token 成本随上下文长度线性增长。
压缩稀疏注意力（CSA）
图 3. CSA 同样让昂贵的注意力读取保持稀疏，但仍需要对压缩历史做一次 top-k 选择。CSA 通过在主注意力之前对压缩键执行选择，降低了索引器成本。
我们从硬件优先的原则出发处理这两种架构特性。在第一篇文章中，我们将受硬件限制的指标描述为性能目标，将受软件限制的指标描述为当前位置的度量。我们在此再次应用这一理念。我们首先测量 MI355X 硬件在实践中能提供什么，然后定位每个 kernel 相对于实际最优点的性能，最后把优化精力花在差距最大的地方。
硬件 Roofline
当厂商宣传加速器芯片时，他们发布的是基于经验理想化或理论工作负载的速率指标（FLOP/s、带宽）。随着后续硬件世代发布，宣传性能与实际工作负载性能之间的差距会扩大[^4]。Zyphra 使用经验 roofline 来指导所有与硬件相关的决策，但 roofline 依赖于工作负载且是特定时间点的。
为了更好地理解我们硬件的特性，我们在软件栈的三个不同层次上测量经验 roofline[^5]。第一层在循环中运行原始矩阵融合乘加（MFMA）指令，我们将结果视为硬件能提供的经验最大值。第二层在大量形状上扫描 PyTorch 和 Transformer Engine 的 GEMM，这更接近我们的应用。GEMM 必须 尺寸合理 才具有代表性，同样的道理也适用于内存带宽（ 我们早前报告的第 III.A 节）。第三层运行一个带宽受限的服务 kernel，即来自 AITER（AI Tensor Engine for ROCm，AMD 的优化推理 kernel 库）的分页注意力解码 kernel pa_fwd_asm，并记录达到的 KV 缓存读取带宽。
我们从硬件优先的原则出发处理这两种架构特性。在第一篇文章中，我们将受硬件限制的指标描述为性能目标，将受软件限制的指标描述为当前位置的度量。我们在此再次应用这一理念。我们首先测量 MI355X 硬件在实践中能提供什么，然后定位每个 kernel 相对于实际最优点的性能，最后把优化精力花在差距最大的地方。
硬件 Roofline
当厂商宣传加速器芯片时，他们发布的是基于经验理想化或理论工作负载的速率指标（FLOP/s、带宽）。随着后续硬件世代的发布，宣传性能与实际工作负载性能之间的差距会不断扩大[^4]。Zyphra 使用经验 roofline 来指导所有与硬件相关的决策，但 roofline 依赖于工作负载，且只反映特定时间点的情况。
为了更好地理解我们硬件的特性，我们在软件栈的三个不同层次上测量经验 roofline[^5]。第一层在循环中运行原始矩阵融合乘加（MFMA）指令，我们将结果视为硬件所能提供的经验最大值。第二层在大量形状上扫描 PyTorch 和 Transformer Engine 的 GEMM，这更接近我们的实际应用。GEMM 必须 尺寸合理 才具有代表性，同样的道理也适用于内存带宽（ 我们早前的报告第 III.A 节）。第三层运行一个带宽受限的服务 kernel，即来自 AITER（AI Tensor Engine for ROCm，AMD 的优化推理 kernel 库）的分页注意力解码 kernel pa_fwd_asm，并记录所达到的 KV 缓存读取带宽。
MI355X 经验 Roofline
图 4. MI355X 上的 roofline 模型[^6]。左：HBM 堆栈向计算 die 传输数据。三个仪表分别显示所使用的 HBM 带宽、MFMA 吞吐量以及 kernel 的算术强度。右：roofline 图。斜线是实测的 HBM 带宽，水平线是实测的 BF16、FP8 和 FP4 的 MFMA 峰值。动画首先将一个 kernel 的算术强度从内存受限区域扫到计算受限区域。然后降低一个内存受限 kernel 和一个计算受限 kernel 的精度。最后，将每种精度下实测最快的 GEMM 形状放到图上。
附录 B 列出了每项测量的指令、标志和扫描设置。
测试 | 实测 | 数据手册（稠密） | 占数据手册的百分比 |
|---|---|---|---|
FP4 MFMA 峰值（TFLOP/s） | 9830.2 | 10066.3 | 97.7% |
FP8 MFMA 峰值（TFLOP/s） | 4936.2 | 5033.2 | 98.1% |
BF16 MFMA 峰值（TFLOP/s） | 2363.1 | 2516.6 | 93.9% |
HBM 带宽，只读归约（GB/s） | 6271.2 | 8000.0 | 78.4% |
HBM 带宽，C++ 拷贝 kernel（GB/s） | 6068.8 | 8000.0 | 75.9% |
HBM 带宽，torch memcpy，多次运行的中位数（GB/s） | 5280.1 | 8000.0 | 66.0% |
表 1. 在单块 GPU 上实测的 MI355X 峰值与 AMD 数据手册中无稀疏性的峰值理论值对比。MFMA 峰值为 100 次实验的平均值。数据手册的计算峰值假设矩阵指令在 2.4 GHz 峰值时钟下每个周期都发射，数据手册的带宽假设 HBM 接口每个时钟周期都传输数据，而这在实践中往往无法实现。
图 5. MI355X 实测能力的可视化。左：MI355X 实测 roofline 图。横轴为算术强度，即 kernel 每向 HBM 移入或移出 1 字节所执行的 FLOPs 数。实线为 MFMA 上限，通过在循环中运行矩阵融合乘加（MFMA）指令测得（表 1）。虚线为每种数据类型下最快 GEMM（通用矩阵-矩阵乘法）形状的吞吐量，每条线上的圆圈标记该形状的算术强度。方块为表 2 中 AITER 分页注意力解码的结果。阴影带覆盖 26 至 137 FLOP/字节，涵盖表 3 中 MoE GEMM 路径算术强度的中位数范围。右：最快 GEMM 形状的吞吐量占 MFMA 峰值的百分比，分别展示 PyTorch（hipBLASLt）和 Transformer Engine 的结果。Transformer Engine 将 MXFP4 GEMM 分派给 AITER a4w4 kernel。
选择 | KV 数据类型 | KV 读取 GB/s | 配置 | 占 HBM 只读带宽的百分比 | 占数据手册 HBM 带宽的百分比 | 占 HBM 拷贝带宽的百分比 | TFLOP/s | FLOP/字节 |
|---|---|---|---|---|---|---|---|---|
最佳 KV 读取带宽 | BF16 | 6010.1 | B32, L16384, kvH16, gqa1 | 95.8% | 75.1% | 99.0% | 6.0 | 1.0 |
最佳 KV 读取带宽 | FP8 | 5574.1 | B32, L65536, kvH8, gqa4 | 88.9% | 69.7% | 91.8% | 43.2 | 7.8 |
最佳吞吐量 | BF16 | 5847.1 | B128, L16384, kvH4, gqa16 | 93.2% | 73.1% | 96.3% | 93.6 | 16.0 |
最佳吞吐量 | FP8 | 5556.0 | B32, L65536, kvH8, gqa16 | 88.6% | 69.5% | 91.6% | 172.4 | 31.0 |
表 2. MI355X GPU 0 上的 AITER pa_fwd_asm 分页注意力解码。头维度 128，块大小 16，连续块表。Config 列中，B 为批大小，L 为序列长度，kvH 为 KV 头数，gqa 为 GQA 比例。三个百分比列将 KV 读取带宽除以表 1 中的 HBM 带宽数值。只读归约 kernel 达到 6271.2 GB/s，C++ 拷贝 kernel 达到 6068.8 GB/s。数据手册值为 8000 GB/s。解码注意力的预期算术强度约为 2 × GQA 比例 / 每个 KV 元素的字节数。
表 1 将各工作负载的实测上限与 AMD 数据手册值进行对比。每个数据手册值描述的是一个理想化工作负载，我们试图展示从硬件最大值到软件现状的完整谱系。计算峰值等于 256 个计算单元 × 2.4 GHz 峰值时钟 × 每个计算单元的矩阵核心每时钟完成的 FLOPs，因此它假设矩阵指令在每个时钟周期都发射。带宽峰值假设 HBM 接口在每个周期都传输数据。
循环 MFMA 指令几乎就是数据手册中理想化的计算工作负载。它们在已在寄存器中的操作数上背靠背发射，没有内存流量或同步，因此只可能因运行在峰值时钟以下或留下空闲发射周期而达不到峰值。因此，MFMA 在 FP8 和 FP4 上达到数据手册峰值的 98%，在 BF16 上达到 94%。
任何内存工作负载都无法匹配标称带宽假设，因为 DRAM 无法在每个周期都传输数据。其存储体需要暂停以进行刷新，而打开新行或在读写之间切换总线都会消耗周期。在我们的三项带宽测试中，对大型连续缓冲区进行只读遍历最接近该假设，达到 6.27 TB/s（数据手册值的 78%）。拷贝内核还会执行写入，而读写之间的切换会带来一些开销，因此达到 6.07 TB/s（76%）。PyTorch 对 256 MiB 缓冲区的拷贝使用 PyTorch 的通用拷贝内核，达到 5.28 TB/s（66%）。因此，选择带宽基准测试，就是在硬件极限与当前软件之间选择所处的位置（见 我们早前的报告）。低级基准内核展示硬件性能，而应用基准既衡量该应用针对硬件的优化程度，也衡量该应用工作负载偏离硬件理想化条件的程度。
图 5 展示了 GEMM 和 AITER 分页注意力解码内核与实测上限的接近程度。最快的 GEMM 形状在 BF16 下达到实测 MFMA 峰值的 66%，在 FP8 下也达到 66%。在 FP4 下，使用 Transformer Engine 时达到 53%，使用 hipBLASLt 时达到 36%。GEMM 从 HBM 加载操作数，并通过 L2 和 LDS 逐级暂存，然后才执行 MFMA 指令；而循环 MFMA 基准测试则将操作数保留在寄存器中。AITER 分页注意力解码内核以 6,010 GB/s 的速度读取 BF16 KV 缓存，达到只读带宽的 96%（表 2）。其算术强度为 1 至 31 FLOP/字节。这低于所有脊点，因此该内核受内存带宽限制，HBM 带宽决定了内核可达到的最佳吞吐量。
我们使用 rocprofv3 硬件计数器，对每个 MoE GEMM 和辅助内核在实际服务形状下进行了性能剖析，并将每一个都放置在这张 roofline 图上（表 3）。在中位数上，这些 GEMM 使用了峰值算力的 4% 到 7%。它们的中位算术强度为 26 到 137 FLOP/字节（图 5 中的阴影带）。这低于所有脊点，因此它们能达到的最高吞吐量由 HBM 带宽决定。例如，稠密 BF16 路径的中位算术强度为 26.4 FLOP/字节。在该强度下，带宽上限为 6.07 TB/s × 26.4 FLOP/字节 = 160 TFLOP/s，即 BF16 MFMA 峰值的 7%。因此，表 3 还报告了每条路径达到的 HBM 带宽占峰值 HBM 带宽的百分比。在中位数上，稠密 BF16 路径为 69%，FP8 拆分路径为 32%。表 3 中的“受限于”列给出了每条路径的主要限制。稠密 BF16 路径受 HBM 带宽限制。FP8 块缩放路径在低 M 时受固定开销限制，在中 M 时受 HBM 带宽限制。a16w4 路径在低 M 时受 tile 填充限制。FP8 拆分路径受延迟和占用率限制。我们在下一节中围绕这些限制设计了融合 MoE 内核。
MoE GEMM 路径 | 算术强度（FLOP/字节），对已剖析 GEMM 形状取中位数 | 峰值算力占比（中位数 / 各形状最大值） | 峰值 HBM 带宽占比（中位数 / 各形状最大值） | Tile 填充（对各形状取中位数） | 受限于 |
|---|---|---|---|---|---|
稠密 BF16 | 26.4 | 5.1% / 6.6% | 69.0% / 88.8% | 1.5 倍 | HBM 带宽 |
FP8 块缩放（2 阶段） | 31.8 | 4.0% / 6.7% | 65.7% / 84.9% | 6.9 倍（M=1 时为 16 倍） | 中等 M 时受 HBM 带宽限制，低 M 时受开销限制 |
a16w4（BF16 x INT4） | 55.0 | 5.6% / 22.5% | 42.9% / 81.6% | 10.3 倍（M≤4 时为 16 倍） | 低 M 时受填充限制，其余情况为混合限制 |
FP8 拆分（3 内核） | 136.7 | 7.2% / 13.1% | 32.3% / 47.9% | 1.5 倍 | 延迟 / 占用率 |
表 3. 使用 rocprofv3 剖析的 MoE GEMM 代码路径。每一行对应一种 MoE GEMM 实现。对于每种实现，我们针对我们在 Zyphra Cloud 上服务的模型所使用的阶段（阶段 1 为 gate+up，阶段 2 为 down-projection）、M 和并行度设置的每一种组合，各剖析了一个 GEMM 形状。每一行中的中位数和最大值均取自这些形状。算术强度为每字节移入或移出 HBM 的 FLOPs。Tile 填充为矩阵核心计算的行数除以真实 token 行数。
在低 M 时，MFMA 浪费工作的最大来源是 tile 量化[^7]。我们使用的 MFMA 指令所计算的 tile 至少为 16 行，而分组 GEMM 无法将来自两个专家的行放入同一个 tile（tile 中的每一行都乘以同一专家的权重）。在低 M 时，每个活跃专家只持有一两行 token，而每个 16 行 tile 的其余部分都是填充。当 token 均匀分布时，填充系数为 min(16 × E / (M × k), 16)，因此对于 E = 256（DeepSeek-V3.2 和 GLM），tile 在 M 接近 512 时被填满；对于 E = 384（Kimi），则在 M 接近 768 时被填满。填充的时间成本取决于每条路径受什么限制（表 3）。在稠密 BF16 路径上，填充的 MFMA 工作在内核等待来自 HBM 的专家权重时运行，并且无论权重 tile 中有多少行持有真实 token，内核都只加载每个权重 tile 一次。在 a16w4 路径上，填充是低 M 时的主要限制。其 INT4 权重的大小是 BF16 权重的四分之一，因此每个权重 tile 的加载时间只有四分之一，而填充的 MFMA 工作和 INT4 到 BF16 的转换在每个 tile 的时间中占据更大份额。在 FP8 路径上，表 3 将固定开销（块缩放）以及延迟和占用率（拆分）列为低 M 时的限制。图 6 展示了 DeepSeek-V3.2 各形状的填充情况。
优化
roofline 测量给出了我们用来对比每个内核的硬件极限。我们的第二轮优化包括四个自定义内核，它们替换了 vLLM 和 SGLang 中的一个 MoE 内核和三个稀疏注意力内核（下文将逐一描述）。我们还在 MI355X 上评估了现有的量化和投机解码方法（MXFP4 方案来自 AMD Quark，MTP 和 DSpark 是已发表的方法）。
首先，我们描述这些新内核，以及它们提升利用率和性能的机制。
Zyphra 融合 MoE 内核的 Tile 调度
图 6. DeepSeek-V3.2 各形状下融合 MoE 内核中的专家 tile。该模型有 256 个专家，每个 token 被发送到其中 8 个。矩阵核心计算 16 行的 tile。当只有一个 token 时，排序内核会构建一个工作列表，为 8 个被选中的专家各分配一个 tile，而其余 248 个专家则没有 tile。每个 tile 中，该 token 占据一行，其余 15 行均为填充。在解码阶段，内核大部分时间都在等待每个专家的权重从 HBM 到达。无论权重 tile 中有多少行持有真实 token，内核都只加载每个权重 tile 一次，因此填充行只会增加 MFMA 工作，而不会增加 HBM 流量。随着批次增大，token 首先会分散到更多专家上，直到在 32 个 token 时全部 256 个专家都处于活跃状态。随后它们会填充已有的 tile，直到在 512 个 token 时每个 tile 都被填满。该动画假设 token 在专家之间均匀分布。
融合 MoE。 我们将专家 MLP 拆分为两个融合阶段。第一阶段是 gate+up GEMM，其 epilogue 在寄存器中计算 SwiGLU 激活，从而省去一次 HBM 往返。第二阶段是 down-projection GEMM，其 epilogue 应用 top-k 加权合并。一个排序内核按专家将 token 到专家的边分桶，并生成一个 tile 工作列表，因此没有收到任何 token 的专家不会分到 tile。该内核在当前 tile 计算的同时从 HBM 加载下一个 tile 的权重。同一结构适用于 BF16、FP8 块缩放、W4A16 INT4，以及 MXFP4 和 MXFP8 微缩放格式。微缩放格式将其 E8M0 块缩放直接送入硬件缩放 MFMA 指令。
Zyphra DSA 内核
图 7. DSA 闪电索引器内核。索引器将一个查询与每个过去位置的索引键进行打分（图中展示了 64 个位置）。两个内核产生相同的 64 个分数。朴素内核（上）在向量 ALU 上用标量乘加计算每个分数，并且为每个分数都再次从内存中读取每个 q 和 k 元素。我们的内核（下）将查询一次性加载到寄存器中。它从 HBM 将 32 个位置的 KV tile 交替复制到两个 LDS 缓冲区中，并在下一个 tile 加载的同时用 MFMA 指令计算每个 tile 的分数。
DSA 闪电索引器。 对于每个查询行 i 以及该行因果窗口内的每个 KV 位置 j，索引器计算 logits[i, j] = sum_h weights[i, h] * relu((Q[i, h, :] . KV[j, :]) * kv_scales[j])。DeepSeek-V3.2 和 GLM 通常以 TP 分片方式部署，索引器在各 rank 上复制，因此索引器在前向传播中所占的份额会随 TP 度增大而增长。我们的内核用 MFMA tile 替代标量点积，在整个行处理期间将查询和头权重保留在寄存器中，并在当前 tile 计算的同时把 KV tile 从 HBM 预取到 LDS。加载描述符携带该行的因果上限，因此超出窗口的位置在硬件层面读为零，内核无需软件边界检查。我们针对 D=128、H=32（GLM-5.1 和 GLM-5.3）以及 H=64（DeepSeek-V3.2）进行调优。
图 7. DSA 闪电索引器内核。索引器将一个查询与每个过去位置的索引键进行打分（图中展示了 64 个位置）。两个内核产生相同的 64 个分数。朴素内核（上）在向量 ALU 上用标量乘加计算每个分数，并且为每个分数都再次从内存中读取每个 q 和 k 元素。我们的内核（下）将查询一次性加载到寄存器中。它从 HBM 将 32 个位置的 KV tile 交替复制到两个 LDS 缓冲区中，并在下一个 tile 加载的同时用 MFMA 指令计算每个 tile 的分数。
DSA 闪电索引器。 对于每个查询行 i 以及该行因果窗口内的每个 KV 位置 j，索引器计算 logits[i, j] = sum_h weights[i, h] * relu((Q[i, h, :] . KV[j, :]) * kv_scales[j])。DeepSeek-V3.2 和 GLM 通常以 TP 分片方式部署，索引器在各 rank 上复制，因此索引器在前向传播中所占的份额会随 TP 度增大而增长。我们的内核用 MFMA tile 替代标量点积，在整个行处理期间将查询和头权重保留在寄存器中，并在当前 tile 计算的同时把 KV tile 从 HBM 预取到 LDS。加载描述符携带该行的因果上限，因此超出窗口的位置在硬件层面读为零，内核无需软件边界检查。我们针对 D=128、H=32（GLM-5.1 和 GLM-5.3）以及 H=64（DeepSeek-V3.2）进行调优。
Zyphra CSA Top-K 内核
图 8. 在 100 万 token 上下文下，单个请求的 CSA top-K 选择。SGLang 的 radix-select 内核（上）在 GPU 的 256 个计算单元之一上执行整个选择过程，每个线程一次只发出一个加载。我们的内核（下）每个线程一次发出八个加载，并将分数拆分到 32 个计算单元上的 32 个块中。每个单元在其块中找出最高的 512 个分数。随后由一个单元将这 32 个列表合并为整个上下文中最高的 512 个分数。计时器显示了实测的内核时间（748.0 微秒和 67.2 微秒）。
CSA top-K 选择。 SGLang 的 HIP top-K 内核采用基数选择（radix select），将单个请求的整个计算放在 256 个计算单元中的单个计算单元上运行，每次只加载一个分数。为改进这一点，我们做了两处改动。现在每个线程一次性发出八个加载，并在批次到达时进行处理。对于约 32K token 及以上的上下文，上下文会被划分到不同计算单元上的 8、16 或 32 个块中。每个单元找出其局部的 top 512，再由一个简短的合并阶段产生真正的 top 512。
CSA 双作用域解码。 我们的注意力内核将滑动窗口范围和压缩后的 top-K 范围融合为一次带有在线 softmax 的流式处理。SGLang 在 Triton 中实现了这一阶段。我们的实现是面向 gfx950 的自定义 HIP 内核，在数值容差范围内计算出完全相同的结果。对于单用户解码，每次调用的工作量很小，因此我们将选中的键分片到 GPU 上。每个分片产生一个部分结果和一个 log-sum-exp 统计量，再由一个合并步骤将各部分结果合并为与单次处理完全一致的答案。键从 FP8 反量化到 LDS 一次后即被复用。
投机解码
图 9. MI355X 上的投机解码。上：在相同时间跨度内，无投机、MTP、EAGLE 和 DSpark 的解码时间线，每种方法两轮。黄色块为草稿模型（drafter）前向，蓝色块为目标模型前向。每条时间线上方的标签是每轮产生的 token。中：token 经过目标模型的 transformer 层，每种方法的草稿模型从所标记的层读取隐藏状态。下：GPU 每次前向从 HBM 读取每个权重块一次，并将其应用于该次前向中的每个 token。因此，token 数更多的前向具有更高的算术强度，并在 roofline 上沿带宽线向上移动。最后一步将负载提升到 64 个并发请求，此时动画中的 DSpark 目标前向达到 FP8 MFMA 峰值。动画中的 MTP 时序与实测的 GLM-5.3 单用户 2.7 倍加速相符。动画中 EAGLE 和 DSpark 的接受率为示意性数据。
投机解码。 在低并发下，一个解码步骤的耗时随该步骤中 token 数量的增加而缓慢增长。对于禁用 MTP 的 vLLM 基线中的 GLM-5.3，1 个 token 的步骤耗时 17.7 ms，而 32 个 token 的步骤（32 个请求各一个）耗时 27.0 ms（表 A7）。该步骤从 HBM 中读取所需的每个权重一次，并将其应用于该步骤中的每个 token，且无论 token 数量多少，它都运行相同的内核序列和 all-reduce。投机解码（Leviathan 等; Chen 等）就使用了这一特性。一个小型草稿模型提出接下来的若干个 token。随后，完整模型（目标模型）在一次前向传播中处理当前 token 以及所有被提出的 token。目标模型按顺序用拒绝采样规则检查这些被提出的 token。它保留第一个被拒绝 token 之前的 token，然后自行采样一个 token。该规则使输出分布与仅用目标模型解码时完全一致。加速幅度取决于目标模型接受多少个被提出的 token，而接受率取决于提示词（见“解码”一节）。
这些方法在草稿生成方式上有所不同。多 token 预测（MTP）增加了一个随模型一起训练的 transformer 层。它接收目标模型的最后一个隐藏状态和最新 token 的嵌入。它使用目标模型的 LM head，并在每次草稿模型前向中提出一个 token。 EAGLE 在目标模型的隐藏状态上训练一个独立的小型头。它提出一棵候选续写树，目标模型在一次前向传播中检查所有分支。 在目标模型的隐藏状态上训练一个独立的小型头。它提出一棵候选续写树，目标模型在一次前向传播中检查所有分支。 DSpark 每次草稿模型前向传播中提出一组 token，并由一个小型顺序头让每个被提出的 token 以其之前的 token 为条件。置信度头估计每个被提出的 token 被接受的可能性。调度器利用这些估计值和服务器负载，决定每个请求向目标模型发送多少个被提出的 token。我们对 GLM-5.3 和 DeepSeek-V4 使用 MTP，对 Kimi-K3 使用 DSpark（对每个模型而言，这都是我们在 MI355X 上测试过的方法中解码吞吐量最高的方法）。
随着并发增加，加速比会下降（在 vLLM 基线中，GLM-5.3 在 1 个并发请求时为 2.7 倍，在 32 个并发请求时为 2.0 倍，见表 A7）。用一条直线拟合两个禁用 MTP 的步进时间，可得到每步约 17.4 ms，外加每 token 约 0.3 ms（(27.0 − 17.7) / 31）。在 1 个请求时，每个被提出的 token 会使步进时间增加约 2%。在 32 个请求时，每个被提出的位置会增加 32 个 token（每个请求一个），并使步进时间增加约 36%。
量化。 量化以更少的比特存储权重。更少的比特意味着权重占用的内存更少，每次前向传播从 HBM 读取的字节也更少，而在 MI355X 上，FP4 MFMA 的峰值是 FP8 峰值的两倍（表 1）。代价是精度更低，因此每个量化模型都必须检查是否损失了准确率。MXFP4 是微缩放（MX）格式之一，属于 开放计算项目 MX 规范 (Rouhani 等）。每 32 个值组成一个块，共享一个 8 位 2 的幂缩放因子，每个值是一个 4 位浮点数，取八种量级之一（0、0.5、1、1.5、2、3、4 或 6）。这样每个值占 4.25 比特，约为 FP8 大小的一半。MI355X 的缩放 MFMA 指令以 MXFP4 值及其块缩放因子作为输入。我们评估了 AMD 用其 Quark 工具包，用于保持推理能力和解码吞吐量。
内核结果
内核基准测试使用 Kimi-K2.6、DeepSeek-V3.2 和 GLM-5.1 的形状。服务结果使用我们目前所服务的模型。GLM-5.1 和 GLM-5.3 具有相同的 MoE 和索引器形状，因此 GLM 内核结果可直接沿用。
我们的融合 MoE 基线是用于 FP8 的 AMD AITER 内核，以及用于 a16w4 的 vLLM Triton WNA16 内核。一个稠密 BF16 PyTorch 实现作为我们的正确性参考。
图 10. 融合 MoE 内核相对基线内核的加速比，随 M（一次前向传播中的 token 数）变化。在解码阶段，M 是一步中解码的 token 数[^8]；在预填充阶段，M 是一次传播中的提示 token 数。大于 1 的值表示我们的内核更快[^9]。左：Kimi-K2.6 a16w4（H=7168，E=384，top-k=8）对比 vLLM Triton WNA16 内核。中：DeepSeek-V3.2 FP8（H=7168，E=256，top-k=8）对比 AITER。在这两个面板中，Full-N 使用完整的专家宽度，TP8 使用其八分之一，即每块 GPU 在 8 路张量并行下运行的形状。右：ZAYA1 FP8（E=24，top-k=1）对比 AITER。所有计时均使用 HIP graphs，与生产环境一致。
在 M 不超过 256 的大多数取值下，我们的内核都比基线更快。较小的 M 值对应解码步骤和较短的预填充块。在 Kimi-K2.6 a16w4 上，基线是 Triton 内核，而我们的内核在除全专家宽度下 M = 1 之外的每个 M 值上都更快，最高可达 2.0 倍。在 FP8 形状上，AITER 在较大的 M 下更快。在全专家宽度下，交叉点位于 M = 512 与 M = 1,024 之间；在 TP8 下，交叉点位于 M = 256 与 M = 512 之间。我们设计内核时针对的是解码场景，此时每个专家只接收少量 token 行（在 DeepSeek-V3.2 上，当 M 不超过 256 时，M × k / E 为 8 或更小）。ZOE 根据这些测量结果按 M 选择 MoE 内核。在交叉点以下它运行我们的内核，在交叉点以上则运行 AITER。
我们还测量了几种变体，它们的运行速度比“优化”一节中描述的内核更慢。我们将其列出，以便其他人可以跳过它们。一个将两个阶段融合在一起的单一内核产生了正确的输出，但在 M ≥ 256 时运行速度慢了 1.7 倍到 2 倍（融合内核需要同时占用两个阶段的寄存器，并将多余部分溢出到暂存内存）。重新排列网格，使最常用的专家先运行，在每个 M 下都更慢。这种重排旨在复用 L2 中的专家权重，却失去了默认列主序所提供的 L2 激活复用。对于 MX 格式，用原子加法对每个 token 的 k 个专家输出求和的收尾步骤比我们的默认收尾步骤更慢。另外三个变体也更慢。它们分别是更宽的 M 行权重复用、在 a16w4 的 M = 1 时为激活使用 LDS 缓存，以及 AITER 风格的 INT4 权重预混洗。
图 11. 在 MI355X 上，使用因果窗口时，我们的内核与两个 AITER 基线的 Lightning 索引器吞吐量（单位：TFLOP/s）。AITER Gluon 是用 Gluon 编写的 AITER 内核，Gluon 是 Triton 项目中的一种底层语言。AITER（各形状最佳）是每个形状上其他 AITER 内核中最快的。H 是索引器头数，D 是头维度。形状为查询行数 × KV 位置数。方形形状的查询行数与 KV 位置数相同（单个提示的预填充）。瘦长形状有 2,048 或 4,096 个查询行和 16K 到 64K 个 KV 位置（长缓存前缀之后短块的预填充）。我们每个柱上方的标签是我们的吞吐量除以 AITER Gluon 的吞吐量。
在方形形状上，我们的内核比 AITER（各形状最佳）快 1.9 倍到 2.4 倍；在瘦长形状上快 1.15 倍到 2.1 倍。在 H=32 时，AITER 的吞吐量在所有形状上保持在 445 到 492 TFLOP/s 之间；在 H=64 时保持在 670 到 782 TFLOP/s 之间。这些形状的查询行数从 2,048 到 65,536 不等。它们的 KV 数组从 0.5 MiB（可容纳于每个 XCD 的 4 MiB L2 中）到 8 MiB（无法容纳）不等。AITER 的吞吐量在这些形状之间变化很小，因此既不是内存也不是行数限制了它。它的速度取决于它为每个 KV 位置块所做的工作。
在方形形状上，当 H=32 时，我们的内核比 AITER Gluon 快 1.07 倍到 1.25 倍；当 H=64 时，其吞吐量是我们的 0.98 倍到 1.11 倍。Gluon 允许内核作者设置数据布局、分配 LDS 并发出异步拷贝。AITER 的 Gluon 内核利用这些能力，在用 MFMA 指令计算分数的同时将 KV 块拷贝到 LDS，这与我们的内核结构相同（见优化部分）。将因果边界检查移入缓冲区加载描述符后，我们在每个形状上的吞吐量都提升了约 10%。
在查询行数为 2,048、KV 位置数为 32K 或 64K 的瘦长形状上，AITER Gluon 比我们的内核快 1.12 倍到 1.37 倍。在 2,048 行、16K 位置时，两个内核的差距在 3% 以内。我们的内核按查询行在 GPU 上划分工作（每行一波），因此行数决定了有多少工作并行执行。在 64K 位置、H=32 时，从 2,048 行增加到 4,096 行，我们的吞吐量从 630 TFLOP/s 提升到 1,006 TFLOP/s，我们的内核比 Gluon 快 1.37 倍。两种形状下每行所做的工作大致相同，因此在 2,048 行时，我们的内核有一部分 GPU 未被利用。同样的变化下，Gluon 的吞吐量提升了 4%。在 H=64 时，该变化将我们的吞吐量从 1,045 TFLOP/s 提升到 1,176 TFLOP/s，Gluon 仍然快 1.10 倍。在大多数瘦长形状上，我们回退到 Gluon 内核。
有两个扩展针对这些形状。BLOCK_Q > 1 的变体会将多个查询行分组，并一次性读取它们 KV 窗口的并集，从而将 KV 读取量最多减少 BLOCK_Q 倍。该变体已实现并正在评估中。预取流水线目前保持两个 tile 在途，我们可以将其增加到三个或四个。
图 12. DeepSeek-V4 稀疏注意力内核在 MI355X 上与 SGLang 基线的对比。左：单次请求的 top-K 选择内核耗时。中：仅加入我们的 top-K 内核后，DeepSeek-V4-Pro 上的端到端单用户解码吞吐量。右：HIP 图下每次调用的双作用域解码内核耗时，与 SGLang 的 Triton 内核对比。
变更 | 测量 | 加速比 |
|---|---|---|
批处理下的 Top-K 内核 | 批大小为 32、上下文长度为 1M 时的内核耗时（1376.8 µs → 139.2 µs） | 9.89x |
双作用域解码内核 | 批大小为 1 / 4 / 8 时的端到端解码吞吐量 | 1.05x / 1.05x / 1.04x |
在两个内核之上实现 MTP | 256K / 512K / 1M 上下文下的单用户解码吞吐量 | 1.04x / 1.05x / 1.15x |
表 4. 在 8x MI355X（TP=8）上对 DeepSeek-V4-Pro 的额外测量。加速比是相对于未加入所列改动的同一配置而言的。
Top-K 选择会对每个压缩位置打分，因此其工作量随上下文线性增长。SGLang 内核在一个计算单元上为一个请求执行选择，而我们的内核将长上下文拆分到最多 32 个计算单元上，因此我们的加速比随上下文增长（图 12，左）。在 4K token 时，我们的内核略慢。在 1M token 时，单请求快 11.1 倍，批大小为 32 时快 9.9 倍（表 4）。
在端到端解码中，选择只是众多步骤之一，因此其效果取决于它在步骤时间中所占的份额。该份额随上下文增长。仅加入 top-K 内核，DeepSeek-V4-Pro 在 1M token 下的单用户吞吐量提升 1.40 倍，而在 32K 和 64K 下，变化处于运行间噪声范围内（图 12，中）。
双作用域解码内核将每次调用拆分到 GPU 上执行，因为单用户解码每次调用的工作量很小（见“优化”一节）。在批大小为 1 时，它每次调用比 SGLang 的 Triton 内核快 1.30 倍。随着批大小增大，每次调用的工作量增加，这一优势逐渐缩小，直到在批大小为 32 时我们的内核略慢（图 12，右）。该内核在批大小 1 到 8 时将端到端解码吞吐量提升约 5%，而在两个内核之上叠加 MTP 后，在 256K 到 1M token 时再提升 4% 到 15%（表 4）。
服务结果
微基准测试只能说明单个内核是否更快。我们的客户看到的是端到端的预填充 TTFT 和解码吞吐量（以每秒 token 数计），而这些还取决于调度、并行、缓存以及服务栈的其他部分。
为了衡量这些指标，我们使用固定的预填充和解码基准测试，其输入序列长度（ISL）、输出序列长度（OSL）和请求并发度均固定不变。这些基准测试并不代表生产流量。它们用于隔离我们的改动在不同模型和服务配方下的效果。我们的基准测试脚本以 vllm bench 和 sglang bench 工具为蓝本。我们自行编写脚本，是为了能够在预填充基准测试中控制前缀缓存命中率，并在两个推理引擎上使用完全相同的设置。
预填充
我们的预填充基准测试一次只运行一个请求，并生成单个输出 token。它遍历不同的输入序列长度，并记录在 1 次预热迭代后 10 次迭代的平均 TTFT。我们还遍历了对应不同类型工作负载的前缀缓存命中率。我们使用 0%（无缓存前缀）、60% 代表常见聊天，以及 90% 代表智能体工作负载。
图 13. 在我们找到的最佳 SGLang 配置下，Kimi-K3 在 MI355X 上的单请求 TTFT。随机输入，OSL=1，1 个并发请求，10 次运行的均值 ± 标准差。
图 13 展示了前缀缓存如何改变 Kimi-K3 的 TTFT。在没有缓存命中的情况下，长上下文时 TTFT 的增长快于输入长度的增长（在 128K token 以上，每翻一倍约增长 2.5 倍），因为在 Kimi-K3 的全注意力层中，每个 token 都要关注所有更早的 token。缓存前缀消除了缓存 token 的逐 token 计算，但每个新 token 仍然要关注完整上下文，包括缓存 token。当 90% 的提示被缓存时，逐 token 计算量下降 10 倍，而这些注意力层中的计算量仅下降约 5 倍[^10]，并且每个请求的固定开销完全没有下降。因此，TTFT 的降幅在中等长度时最大（128K token 时为 8 倍）。在短上下文时降幅较小，因为此时固定开销占主导（8K 时为 2.6 倍）；在长上下文时降幅也较小，因为此时注意力占主导（512K 时为 6 倍）。
对于 GLM-5.3，我们将开箱即用的 vllm/vllm-openai-rocm:v0.29.0 与在同一镜像上启用 ZOE 的版本进行对比，后者加入了我们的 DSA lightning indexer 和融合 MoE kernel。
图 14. GLM-5.3 在 MI355X 上的单请求 TTFT，vLLM v0.29.0 基线对比同一镜像启用 ZOE 的版本。随机输入，OSL=1，1 个并发请求，10 次运行的均值 ± 标准差。ZOE 始终达到或超过 vLLM 实现的性能，并在长上下文长度下明显拉开差距。
ZOE 为 GLM-5.3 增加了两个 kernel：lightning indexer 和融合 MoE kernel。在 DSA 模型中，indexer 是 prefill 中唯一每个 token 的工作量随上下文增长的部分。它将每个新 token 与所有更早的位置进行打分，而主注意力每个 token 只读取固定的 2,048 个条目，MoE 和投影层对每个 token 做相同的工作。在 16K token 及以下时，indexer 只占 prefill 的一小部分，加速比接近 1x，这也表明 MoE kernel 对 prefill 时间影响很小。随着上下文增长，indexer 占据了 prefill 时间的大部分，加速比也随之增长，在 512K token 时达到 2.8x。前缀缓存命中会进一步提高加速比。缓存 token 跳过了逐 token 的工作，但每个新 token 仍然要与完整上下文进行打分，因此 indexer 在剩余工作中所占比例更大（在 512K、90% 命中率时为 3.6x）。唯一低于 1x 的结果（在 32K、90% 命中率时为 0.80x）源于 ZOE 各次运行之间的高波动（标准差为 0.16 秒，均值为 0.30 秒）。
解码
我们的解码基准测试扫描了请求并发度。低并发度带来最佳的单用户吞吐量，高并发度则带来最佳的单节点吞吐量。每个解码步骤都会为其中的每个请求读取一次权重，因此更多请求可以共享每次读取，但每个步骤的耗时也会更长。我们排除了 TTFT，仅报告单用户和单节点吞吐量。解码性能在很大程度上取决于投机解码，因此对提示/输出分布较为敏感。我们的基准测试使用 temperature=1 和 top_p=0.95，由于投机器的接受率会发生变化，同一提示在多次运行之间可能产生较大方差。为展示这种敏感性，我们报告了 Kimi-K3 在两个提示上的结果，一个来自 Apex-Shortlist（重度推理），一个来自 GSM8K（简单数学）。
Apex。 “求在所有正整数三元组 \((a,b,c)\) 上，\(gcd(a2+b2+c2,abc)\) 可能取到的最小的 50 个值之和。让我们逐步思考，并将最终答案输出在 \boxed{} 中。”
GSM8K。 “Tobias 正在买一双价格为 95 美元的新鞋。过去三个月里，他每个月都在存钱。他每月得到 5 美元的零花钱。他还修剪草坪和铲除车道上的积雪。他修剪一块草坪收费 15 美元，铲雪收费 7 美元。买完鞋后，他还剩下 15 美元。如果他修剪了 4 块草坪，那么他铲了多少条车道？”
图 15. Kimi-K3 在 MI355X 上使用我们找到的最佳 SGLang 配置（包含 DSpark）时的解码吞吐量，OSL=1024，10 次运行的均值 ± 标准差。
模型和内核在两个提示上完全相同，因此两条曲线之间的差距来自 DSpark 接受率。目标模型对 GSM8K 答案中草稿模型提出的 token 接受得比 Apex 推理更多，因此每一步都会产生更多 token。因此，投机解码的结果必须结合提示分布来解读。
对于 GLM-5.3，我们报告两种设置。在禁用 MTP 时，基准测试的方差接近于零，从而将吞吐量变化与内核隔离开来。在启用 MTP 时，该配置与我们在生产环境中使用的配置一致，而投机解码带来的方差大于大多数内核增益。
图 16. GLM-5.3 在 MI355X 上的单节点解码吞吐量，vLLM v0.29.0 基线对比同一镜像启用 ZOE 的情况。Apex 提示，OSL=1024，100 次运行的均值。误差条表示一个标准差。
ZOE 对 GLM-5.3 解码的改动小于对预填充的改动。Apex 运行从短提示生成 1,024 个 token，因此上下文保持较短，索引器几乎不做什么工作。剩下的改动是融合 MoE 内核，它只加速每个解码步骤中的 MoE 部分。在禁用 MTP 时，ZOE 将吞吐量提升 1% 到 7%。在启用 MTP 时，改动范围从 −3% 到 +13%，且大多数都落在运行间波动的一个标准差之内。在短上下文下，投机解码对解码吞吐量的影响远大于本文中的内核（MTP 将 GLM-5.3 吞吐量提升 2 倍到 2.7 倍），而长上下文内核在上下文变长后才会起作用（图 12）。
量化
Kimi-K3 已经随附 MXFP4 专家权重，进一步压缩的空间很小，因此我们将重点放在随附 FP8 权重的 GLM-5.3 上。AMD 已发布 amd/GLM-5.3-Quark-MXFP4-AttnFP8该模型以 MXFP4 存储路由专家和共享专家，并将注意力保持在 FP8。这将模型大小从 756 GB 降至 410 GB。Hugging Face 模型卡报告的 GSM8K 得分为 91.28，而 FP8 模型为 91.89（恢复率 99.3%）。我们发现，仅凭 GSM8K 无法判断模型在训练后量化后是否仍保持推理能力，因此我们用自有的评测框架对两个模型在四个推理基准上进行了测试（表 5）。
数据来源 | 样本（问题 × 8 次重复） | zai-org/GLM-5.3 得分（%） | amd/GLM-5.3-Quark-MXFP4-AttnFP8（410 GB）得分（%） | 变化（百分点） |
|---|---|---|---|---|
AIME25 | 240 | 97.5 | 99.17 | +1.67 |
AIME25 | 240 | 96.67 | 96.67 | 0.00 |
Apex-Shortlist | 376 | 53.72 | 53.46 | −0.26 |
HMMT | 240 | 96.67 | 93.75 | −2.92 |
表 5. GLM-5.3 量化评测对比，使用 Zyphra 内部评测框架。每个提示 8 次重复，max-tokens 131072，temperature=1.0，top_p=0.95。变化以百分点计。镜像：vllm/vllm-openai-rocm:nightly-rocm100-9ea8f3ffc354901b740f0b31988900897b7221d7。
MXFP4 模型在全部四个基准上都与 FP8 模型相差不超过 3 分。每个基准有 30 道题（Apex-Shortlist 为 47 道），每题采样 8 次，因此一道从总是解出变为从未解出的题会使分数变动 3.3 分（Apex-Shortlist 为 2.1 分）。表 5 中的每个差异都小于该值。
对于服务吞吐量，我们禁用推测解码并使用固定的解码基准测试。在 410 GB 下，MXFP4 模型可容纳在四块 MI355X GPU（1,152 GB HBM）上，并为 KV 缓存留有余量，因此我们还尝试了在该节点上运行两个 TP=4 副本（DP=2）。这些运行使用表 5 中列出的 nightly 镜像，因此绝对数值无法与 Decode 部分中 v0.29.0 的数值直接比较。
图 17. GLM-5.3 在 MI355X 上使用 FP8 专家时的单节点解码吞吐量，分别展示使用和不使用我们的融合 MoE kernel，以及使用 MXFP4 专家的情况。所有运行均使用我们找到的最佳 vLLM 配置，以及表 5 中列出的 vLLM ROCm nightly 镜像。斜线阴影柱表示将节点拆分为两个 TP=4 副本。MTP 已禁用，Apex 提示，OSL=1024，取 10 次运行的平均值。
正如我们在 roofline 实验中所描述的，量化到 FP8 或 FP4 都能减少 HBM 传输时间，并提高可能的计算受限上限。这两种效果在实现上都有一些注意事项。对于 HBM 传输时间，可达到的 HBM 带宽取决于所传输缓冲区的大小，并且需要足够大的缓冲区才能饱和带宽[^11]。对于计算速度，尽管绝对吞吐量上限在提高，但随着逐级降到 FP8 和 FP4，相对可达到的吞吐量往往会下降。这源于若干与实现相关的具体原因：
FP8 和 FP4 量化并不像 BF16/FP16 与 FP32 之间那样只是简单的截断和舍入。它需要量化操作[^12]（通常融合进核心算子内核中），而这些操作会带来额外开销。由于计算和数据传输时间已经减少，这些开销在整体操作中所占的比例更大，而且我们大幅提高了所需的问题规模来隐藏这些操作。
重叠变得更加细粒度且更难实现，因为计算延迟已经降低，而且我们加入了更多无法在矩阵核心上运行、并且有严格顺序要求的量化算子。
形状与占用率往往会给 FP8 和 FP4 矩阵核心带来更多约束，即模型维度的除数更大，因此能被整除的模型更少，而且这些模型要大得多。以 MI355X 为例，与 BF16 相比，其 FP8 和 FP4 矩阵核心在 K GEMM 维度上的深度增加了 4 倍。以 16x16 MFMA 指令为例，FP8 和 FP4 的 K 深度为 128，而 BF16 仅为 32。这意味着，例如在 BF16 下可整除的 96 或 192 的 head 维度现在不再可整除，这需要在 FP4/FP8 内核中进行自定义处理（例如自定义填充或重叠多个工作流），才能达到极速性能。这种翻倍或减半的情况经常出现，例如低精度内核使用更大的 tile 来让更快的矩阵核心保持饱和，因此给定的 GEMM 会拆分成更少的 tile 和更少的 wave，而最后不满一个 wave 的部分在整体运行时间中所占比例更大。对于 Hopper 等较旧的 GPU 来说，这并不算大问题，因为增幅较小，仍然基本能整除任何数值，BF16 为 16，FP8 为 32；但对于 MI355X 上的 FP8/FP4、Blackwell 上的 FP4 等，情况已不再如此。
简而言之，MFMA 的速度显著提升，但大多数其他硬件操作并没有按比例变快，这就是为什么例如 Flash Attention 4 在设计时花了大量精力来消除 exp2f() 这一瓶颈，而内核开发者正转向融合尾声（例如 CuTe DSL 中的 QuACK）和巨型内核，以更好地隐藏非 GEMM 工作。硬件继续以远快于所有内存带宽的速度提升低精度 FLOPs。投机解码和扩散语言模型都能在每次权重读取中处理多个 token，这提高了解码的算术强度，并使其向计算受限区域移动。对于 BF16 来说这很直接，因为我们能得到人们习惯的实测与标称 TFLOP/s 之比，而且几乎所有模型形状都能受益。如果你把量化叠加在任何受内存带宽限制的东西之上，只要量化开销被隐藏（如果你的量化缓冲区仍超过带宽饱和点，则是线性加速），你就能获得加速。如果你把量化叠加在任何受计算限制的东西之上，你就要承担上述实现复杂性，而这通常是可以克服的。
至此，我们已经具备了理解量化结果的条件。量化主要通过减少从 HBM 读取的字节数来加速解码，因此其收益取决于每个解码步骤中有多大比例用于读取权重。在并发度为 1 的 MoE 模型中，每个 token 只使用其 Top-K 路由专家加上每层共享专家的数量，因此只有这些专家会从 HBM 传输。例如，在 1 个并发请求下，每个 GLM-5.3 MoE 层使用其 257 个专家中的 9 个（8 个路由专家和 1 个共享专家）。在 GLM-5.3 的 75 个 MoE 层中，这些专家以 FP8 格式占用 25.5 GB，节点的 8 块 GPU 在实测 HBM 带宽下约 0.5 ms 即可读取完毕（表 1）。解码步骤耗时约 16 ms（图 17），而 MXFP4 可将这 0.5 ms 减半。该步骤的大部分时间用于专家格式不会改变的工作，例如注意力、GPU 之间的通信以及每个内核的固定开销。将专家字节数减半，只能使这样的步骤缩短约 2%。随着并发度增长，一个步骤中的 token 会使用更多不同的专家，该步骤会读取更多专家权重，量化所减少的字节数在端到端时间中所占的比例也随之增大。
MXFP4 专家执行与 FP8 专家相同的操作（量化激活、运行两个专家 GEMM 并合并结果），但权重字节数减半，矩阵核心峰值翻倍。因此，MXFP4 计算时间较慢的原因归结为内核实现，以及围绕上述挑战进行优化的程度。我们的 Fused MoE 实现中的 FP4 支持（尤其是低批量场景）仍在进行中。在 vLLM 中，MXFP4 模型在 1 和 8 个并发请求时比 FP8 模型慢，在 32 个并发请求时更快（图 17）。具体而言，在低并发下，MXFP4 内核每层增加的时间超过了更小权重所节省的时间；而在 32 个请求时，节省的字节数超过了这一开销。将节点拆分为两个 TP=4 副本在所有并发度下都更慢（TP=4 副本中的每块 GPU 每步读取的权重字节数翻倍）。使用我们的 fused MoE 内核，FP8 模型在低并发下仍比 MXFP4 模型更快，在 32 个请求时差距约在 1% 以内，因此在低并发场景下我们继续在生产环境中提供 FP8 模型。
智能体性能基准
固定的预填充和解码基准无法可靠地反映生产环境中的性能。真实流量是多轮且长上下文的，其中包含子智能体突发、KV 缓存卸载和工具调用。在 Zyphra，我们使用 SemiAnalysis 的 AgentX 基准，这是一个开源的多轮智能体编码推理基准，上下文长度为 100 万 token。它可在 InferenceX 仓库，其中更详细地描述了它。我们将很快发布智能体基准测试结果。
未来工作
本文介绍我们在 Zyphra Inference 上进行的第二轮全栈推理优化。我们将继续改进推理平台，重点关注以下几个方面：
在 FP8 前沿 MoE 形状上弥合大 M 维度的差距，使预填充不再通过混合调度器回退到厂商内核。
完成 BLOCK_Q > 1 的闪电索引器变体，并为瘦长形状构建更深的预取流水线。
MXFP4 MoE 内核在低并发下更快，MXFP4 GEMM 达到 MI355X FP4 峰值的比例更高（图 5），因此 MXFP4 模型在每种并发下解码都比 FP8 模型更快。
量化方案源自我们的预训练经验，以及利用 MI355X 的 HBM 容量定制的推测器。
使用 AgentX 在 100 万 token 上下文下进行的智能体基准测试结果。
Zyphra Inference 现已上线，网址为 cloud.zyphra.com.
方法与局限性
与我们的第一篇博文一样，我们会介绍我们的方法及其局限性，因为推理领域存在许多不同的情形，我们不希望进行苹果与橘子的比较。所描述的局限性 在我们上一篇博客的开头 同样适用于此。我们在下面重述相关条目，以及针对我们内核基准测试的特定条目。
所有结果均在单节点上取得，且没有实时流量。请求以固定间隔而非随机到达，我们也没有在真实环境中测试调度，因为这超出了本工作的重点。
推测解码方法的性能对基准测试中使用的提示词很敏感，因为草稿接受率会随输入分布而变化。一个被提议的 token 被接受的概率等于草稿模型与目标模型下一 token 分布的重叠程度，而草稿模型在可预测文本（如算术或代码）上比在开放式文本上更接近目标模型。像 MTP 和 EAGLE 那样读取目标模型的隐藏状态可以改善这种匹配，但并不能消除这种依赖性。在温度为 1 时，EAGLE-2 在六个基准测试上平均每次目标模型前向通过接受 3.5 到 4.9 个 token（Li 等人），图 15 展示了 DSpark 的相同效果。
我们在各配置中固定输入提示词，以使 KV 缓存和前缀缓存的行为尽可能相似；对于敏感度较大的情况，我们报告了多个提示词下的解码结果。
内核基准测试在每次迭代之间使用 256 MB 缓冲区刷新 L2 缓存，并在 20 次预热迭代后报告 200 次迭代的中位数。
我们报告在 HIP 图捕获与重放下的计时结果，我们认为这反映了生产环境中的行为，因为逐次启动（eager）计时包含了图所消除的启动开销。加速比为 median(baseline) / median(ours)。
我们的推理软件栈和外部基线都在快速变化，因此我们在此展示的任何结果仅在某一时间点有效。这一时间点可能因模型而异，取决于实验运行的时间。这些博客是滑动窗口，总结自上一篇博客以来所完成的工作。服务结果使用 vLLM v0.29.0 和 SGLang ROCm 镜像 v0.5.19（lmsysorg/sglang-rocm:v0.5.19-rocm720-mi35x-20260910），量化结果使用 vLLM ROCm nightly 镜像。
附录 A：数据表
下表包含图中绘制的数值。
[wide]A1. 融合 MoE 加速（图 10）
[wide]Kimi-K2.6 a16w4（H=7168，E=384，top-k=8），在 HIP 图下相对 vLLM Triton WNA16 的加速比。
M | 1 | 8 | 16 | 64 | 256 | 512 | 1024 | 2048 | 4096 | 8192 |
|---|---|---|---|---|---|---|---|---|---|---|
Full-N | 0.71 | 1.07 | 1.04 | 1.20 | 2.03 | 1.81 | 1.51 | 1.24 | 1.24 | 1.16 |
TP81.091.031.001.171.060.840.750.560.550.55 | 1.81 | 1.48 | 1.50 | 1.45 | 1.96 | 1.68 | 1.39 | 1.78 | 1.51 | 1.65 |
[wide]DeepSeek-V3.2 FP8（H=7168，E=256，topk=8），在 HIP 图下相对 AITER 的加速比。
M1816642565121024204840968192Full-N0.711.071.041.202.031.811.511.241.241.16TP81.811.481.501.451.961.681.391.781.511.65[wide]DeepSeek-V3.2 FP8（H=7168，E=256，topk=8），在 HIP 图下相对 AITER 的加速比。 | 1 | 8 | 16 | 64 | 256 | 512 | 1024 | 2048 | 4096 | 8192 |
|---|---|---|---|---|---|---|---|---|---|---|
Full-N | 1.10 | 0.99 | 1.03 | 1.15 | 1.14 | 1.07 | 0.93 | 0.65 | 0.59 | 0.55 |
TP81.811.481.501.451.961.681.391.781.511.65 | 1.09 | 1.03 | 1.00 | 1.17 | 1.06 | 0.84 | 0.75 | 0.56 | 0.55 | 0.55 |
[wide]Zyphra ZAYA1 FP8（E=24，topk=1），相对 AITER 的加速比。
1816642565121024204840968192Graph1.071.040.981.421.461.231.000.650.580.54[wide]A2. 闪电索引器吞吐量（图 11）[wide]TFLOP/s，因果窗口，MI355X。高于 1.00x 的比值表示我们的 kernel 更快。H=32，D=128Shape（seq x kv）AITER（各形状最佳）AITER Gluon我们的Ours / AITER GluonOurs / AITER（各形状最佳）4096 x 40964457328521.16x1.91x8192 x 819248291411421.25x2.37x16384 x 1638449294910401.10x2.12x32768 x 3276846595310191.07x2.19x65536 x 6553646591610261.12x2.21x2048 x 1638447195092 | 1 | 8 | 16 | 64 | 256 | 512 | 1024 | 2048 | 4096 | 8192 |
|---|---|---|---|---|---|---|---|---|---|---|
图 | 1.07 | 1.04 | 0.98 | 1.42 | 1.46 | 1.23 | 1.00 | 0.65 | 0.58 | 0.54 |
[wide]A2. 闪电索引器吞吐量（图 11）
[wide]TFLOP/s，因果窗口，MI355X。高于 1.00x 的比值表示我们的 kernel 更快。H=32，D=128
形状（seq x kv） | AITER（各形状最佳） | AITER Gluon | 我们的 | 我们的 / AITER Gluon | 我们的 / AITER（各形状最佳） |
|---|---|---|---|---|---|
4096 x 4096 | 445 | 732 | 852 | 1.16x | 1.91x |
8192 x 8192 | 482 | 914 | 1142 | 1.25x | 2.37x |
16384 x 16384 | 492 | 949 | 1040 | 1.10x | 2.12x |
32768 x 32768 | 465 | 953 | 1019 | 1.07x | 2.19x |
65536 x 65536 | 465 | 916 | 1026 | 1.12x | 2.21x |
2048 x 16384 | 471 | 950 | 921 | 0.97x | 1.96x |
2048 x 32768 | 480 | 648 | 551 | 0.85x | 1.15x |
2048 x 65536 | 469 | 705 | 630 | 0.89x | 1.34x |
4096 x 65536 | 474 | 735 | 1006 | 1.37x | 2.12x |
[wide]H=64，D=128
形状（seq x kv） | AITER（各形状最佳） | AITER Gluon | 我们的 | 我们的 / AITER Gluon | 我们的 / AITER（各形状最佳） |
|---|---|---|---|---|---|
4096 x 4096 | 670 | 1273 | 1411 | 1.11x | 2.11x |
8192 x 8192 | 761 | 1550 | 1531 | 0.99x | 2.01x |
16384 x 16384 | 782 | 1618 | 1588 | 0.98x | 2.03x |
32768 x 32768 | 755 | 1623 | 1613 | 0.99x | 2.14x |
65536 x 65536 | 758 | 1600 | 1562 | 0.98x | 2.06x |
2048 x 16384 | 763 | 1448 | 1437 | 0.99x | 1.88x |
2048 x 32768 | 755 | 1233 | 898 | 0.73x | 1.19x |
2048 x 65536 | 751 | 1256 | 1045 | 0.83x | 1.39x |
4096 x 65536 | 747 | 1297 | 1176 | 0.91x | 1.57x |
[wide]A3. DeepSeek-V4 稀疏注意力（图 12）
[wide]Top-K kernel 时间，单请求，MI355X。
上下文 | 基线（µs） | 我们的（µs） | 加速比 |
|---|---|---|---|
4K | 9.7 | 11.2 | 0.87x |
64K | 58.2 | 30.5 | 1.91x |
256K | 178.4 | 43.1 | 4.14x |
1M | 748.0 | 67.2 | 11.1x |
[wide]端到端单用户解码，DeepSeek-V4-Pro，SGLang，tok/s。
上下文 | 32K | 64K | 256K | 512K | 1M |
|---|---|---|---|---|---|
基线 | 67.1 | 65.8 | 59.4 | 52.7 | 43.5 |
Zyphra Top-K | 66.0 | 65.6 | 64.3 | 63.0 | 61.1 |
[wide]HIP 图下双作用域逐调用解码延迟，对比 SGLang 生产环境 Triton 路径。
批处理 | 1 | 2 | 4 | 8 | 16 | 32 |
|---|---|---|---|---|---|---|
Triton（µs） | 22.9 | 26.9 | 28.4 | 34.1 | 43.9 | 51.7 |
我们的（µs） | 17.6 | 20.8 | 24.5 | 30.6 | 40.2 | 54.6 |
加速比 | 1.30x | 1.29x | 1.16x | 1.12x | 1.09x | 0.95x |
[wide]A4. Kimi-K3 预填充 TTFT（图 13）
[wide]MI355X、Kimi-K3、最佳 SGLang 配置、随机输入、OSL=1、固定前缀以模拟缓存命中率，含预热迭代，1 个并发请求。TTFT 以秒为单位。
输入长度（tokens） | 0% 缓存命中 | 60% 缓存命中 | 90% 缓存命中 |
|---|---|---|---|
8192 | 0.563±0.160 | 0.314±0.050 | 0.217±0.100 |
16384 | 1.017±0.155 | 0.522±0.155 | 0.234±0.020 |
32768 | 1.965±0.016 | 0.859±0.007 | 0.404±0.222 |
65536 | 4.284±0.275 | 1.964±0.412 | 0.620±0.183 |
131072 | 9.074±0.049 | 4.385±0.342 | 1.137±0.004 |
262144 | 22.861±0.148 | 10.730±0.504 | 2.905±0.002 |
524288 | 56.896±0.098 | 29.690±0.822 | 9.275±1.103 |
[wide]A5. GLM-5.3 预填充 TTFT（图 14）
[wide]MI355X、GLM-5.3、最佳 vLLM 配置（v0.29.0）、随机输入、OSL=1、固定前缀以模拟缓存命中率，含预热迭代，1 个并发请求。与 ZOE（DSA 索引器和融合 MoE）对比的基线。TTFT 以秒为单位。
输入长度（tokens） | 缓存命中率 | TTFT 基线 | 使用 ZOE 时的 TTFT | 变更 |
|---|---|---|---|---|
8192 | 0% | 0.436±0.135 | 0.407±0.017 | -6.65% |
16384 | 0% | 0.833±0.002 | 0.830±0.004 | -0.36% |
32768 | 0% | 1.752±0.002 | 1.838±0.405 | +4.91% |
65536 | 0% | 3.825±0.005 | 3.613±0.030 | -5.54% |
131072 | 0% | 10.122±0.451 | 8.113±0.006 | -19.85% |
262144 | 0% | 35.044±0.003 | 20.772±0.304 | -40.73% |
524288 | 0% | 178.770±0.790 | 64.569±0.317 | -63.88% |
8192 | 60% | 0.205±0.001 | 0.204±0.001 | -0.49% |
16384 | 60% | 0.367±0.001 | 0.360±0.001 | -1.91% |
32768 | 60% | 0.764±0.002 | 0.734±0.006 | -3.93% |
65536 | 60% | 1.715±0.001 | 1.564±0.003 | -8.80% |
131072 | 60% | 5.206±0.005 | 3.746±0.001 | -28.04% |
262144 | 60% | 21.763±0.008 | 10.443±0.010 | -52.01% |
524288 | 60% | 126.867±0.055 | 36.682±0.019 | -71.09% |
8192 | 90% | 0.129±0.003 | 0.129±0.001 | 0.00% |
16384 | 90% | 0.132±0.002 | 0.132±0.003 | 0.00% |
32768 | 90% | 0.240±0.001 | 0.301±0.161 | +25.42% |
65536 | 90% | 0.512±0.023 | 0.437±0.003 | -14.65% |
131072 | 90% | 1.552±0.004 | 1.038±0.004 | -33.12% |
262144 | 90% | 6.928±0.003 | 3.029±0.001 | -56.28% |
524288 | 90% | 42.301±0.190 | 11.821±0.735 | -72.06% |
[wide]A6. Kimi-K3 解码吞吐量（图 15）
[wide]MI355X、Kimi-K3、最佳 SGLang 配置（含 DSpark）、OSL=1024、10 次运行平均值，tok/s。
并发请求 | 每节点 Apex | 每用户 Apex | 每节点 GSM8K | 每用户 GSM8K |
|---|---|---|---|---|
1 | 91.6 ± 5.4 | 91.6 ± 5.4 | 247.9 ± 19.7 | 247.9 ± 19.7 |
2 | 161.8 ± 7.2 | 82.5 ± 3.5 | 426.1 ± 34.3 | 226.0 ± 15.6 |
4 | 285.8 ± 10.6 | 74.3 ± 2.5 | 761.4 ± 44.5 | 209.0 ± 9.3 |
8 | 449.7 ± 13.2 | 59.1 ± 1.7 | 1206.0 ± 80.9 | 170.2 ± 9.8 |
16 | 666.3 ± 16.4 | 43.9 ± 1.0 | 1790.2 ± 100.6 | 128.2 ± 6.1 |
32 | 1112.7 ± 24.2 | 36.9 ± 0.8 | 2964.1 ± 142.2 | 107.6 ± 4.5 |
[wide]A7. GLM-5.3 解码吞吐量（图 16）
[wide]MI355X、GLM-5.3、最佳 vLLM 配置（v0.29.0）、Apex 提示、OSL=1024、DSA + 融合 MoE、100 次运行平均值、每节点 tok/s。
[wide]MTP 已禁用
并发请求 | 基线 | ZOE | 变化 |
|---|---|---|---|
1 | 56.4 | 58.8 | +4.26% |
2 | 108.1 | 113.3 | +4.81% |
4 | 204.8 | 214.7 | +4.83% |
8 | 389.9 | 402.3 | +3.18% |
16 | 721.2 | 726.0 | +0.67% |
32 | 1186.5 | 1267.9 | +6.86% |
[wide]已启用 MTP（均值 ± 标准差）
并发请求 | 基线 | ZOE | 变化 |
|---|---|---|---|
1 | 152.7 ± 13.6 | 160.4 ± 14.1 | +5.04% |
2 | 282.7 ± 19.8 | 297.1 ± 19.9 | +5.09% |
4 | 461.8 ± 23.7 | 519.6 ± 28.3 | +12.52% |
8 | 810.1 ± 37.0 | 860.4 ± 41.4 | +6.21% |
16 | 1408.7 ± 59.4 | 1430.3 ± 57.6 | +1.53% |
32 | 2431.3 ± 86.5 | 2363.6 ± 81.8 | −2.78% |
[wide]A8. GLM-5.3 量化吞吐量（图 17）
[wide]MI355X、GLM-5.3、最佳 vLLM 配置（MTP 已禁用）、Apex 提示、OSL=1024、10 次运行平均值、每节点 tok/s。vLLM ROCm nightly 镜像（见表 5）。
量化方案 | 并行性 | 并发数=1 | 并发数=8 | 并发数=32 |
|---|---|---|---|---|
zai-org/GLM-5.3 | TP=8 | 62.4 | 412.3 | 1223.6 |
zai-org/GLM-5.3 + zyphra fused_moe | TP=8 | 64.6 | 430.8 | 1322.4 |
amd/GLM-5.3-Quark-MXFP4-AttnFP8 | TP=8 | 52.2 | 395.6 | 1338.7 |
amd/GLM-5.3-Quark-MXFP4-AttnFP8 | TP=4，DP=2 | 48.3 | 351.8 | 1185.3 |
附录 B：Roofline 测量设置
MI355X 只有通过新的 CDNA4 f8f6f4 指令，才能以全速率运行 FP8 和 FP4 矩阵运算。我们的 FP8 和 FP4 循环使用其中之一，即 v_mfma_scale_f32_32x32x64_f8f6f4，而 BF16 循环使用 v_mfma_f32_32x32x16_bf16。上游基准测试（rocm-amdgpu-bench）使用的是 v_mfma_f32_32x32x16_fp8_fp8，这是从 CDNA3 沿用下来的 FP8 指令，在 MI355X 上其 FP8 速率约为一半。MFMA 和 GEMM 基准测试将所有量化缩放系数设为 1。在每次计时的 GEMM 之前，我们会刷新 L2 并运行一段短暂的 GPU 忙等待，因此计时事件只测量 GEMM 内核。FP4 GEMM 运行中禁用了 TunableOp，因为其缩放 GEMM 调度器没有 FP4 路径。对于 FP8 和 BF16，启用与不启用 TunableOp 的结果差异小于 2%。对于注意力解码，pa_fwd_asm 要求头维度为 128，并支持最高 16 的 GQA 比例。我们扫描了批大小（1 到 256）、序列长度（1K 到 64K token）、KV 头数（4、8、16）以及 GQA 比例（1、4、8、16）。一个批次中的所有序列长度相同（生产批次包含不同长度的序列）。