*之前我在论坛发过一篇情绪贴，吐槽海光以及 AMD 对老卡的适配太差，导致很难用于现代大模型推理。不过，光输出情绪是没有用的，所以不如现在就开始干。* 大家好，从这一篇开始，我准备把这段时间给海光 Z100 适配 vLLM 的过程陆续整理出来。vLLM 是面向大语言模型服务的高吞吐推理系统，其代表性工作 PagedAttention 通过分页式管理 KV Cache 来降低显存碎片和冗余占用[1]。这里说的“适配”并不是让 vLLM 识别出一张 Z100、把模型勉强启动起来就算完成，而是希望把现在大模型推理中常用的计算算子、显存访问、双卡通信、图执行和量化推理路径真正落到 gfx906 上，再在结果可靠的前提下继续往下压延迟、榨显存带宽。这一阶段我大量使用了现在常说的 Vibe Coding，也就是借助大模型辅助阅读代码、生成修改方案、设计实验和分析日志，但真正做到 GPU 底层以后，最后能相信的仍然只有实测：代码能编译不等于能算对，单独测试很快不等于完整模型会变快，完整模型能生成文字也不等于内部数值一定正确。 这件事刚开始的时候，我其实把难度想低了。Z100 在当前软件栈中被工具链识别成 `gfx906`，vLLM 本身又已经提供 ROCm 后端，所以第一反应很自然：把 gfx906 加进架构判断，修一修编译错误，剩下的现成 ROCm 算子应该就能接着用。真正跑起来以后才发现，问题远不只是“少了一个设备编号”。今天 vLLM 的 ROCm 高性能路径主要围绕较新的 AMD 数据中心 GPU 演进，很多代码里已经默认存在较新的指令、波前组织、矩阵计算能力和通信语义；gfx906 虽然名字也以 `gfx9` 开头，却属于更早一代 GCN 5.1 架构，而 MI300/MI350 已分别进入 CDNA3/CDNA4 世代，硬件能力和编程假设之间存在明显代际差异[3]。如果平台识别、Triton 代码生成、注意力、量化矩阵乘、集合通信或者多进程启动中的某一层错误地套用了新架构的前提，程序就可能出错，而且最麻烦的还不是直接崩溃，而是“能运行，但是偷偷算错”。 所以这篇文章不是一篇“如何安装 vLLM”的教程，也不准备在开头贴一条所谓万能启动命令。我更想把这一阶段真正遇到的问题串起来：为什么最开始只能使用即时执行模式，后来怎样把 HIP 图执行跑起来；为什么某些 Triton 内核明明编译成功，却不能直接相信；双卡自定义全归约为什么在微基准里快得很好看，放到完整模型以后却会读到上一轮留下来的旧数据；以及 AWQ INT4 的线性层为什么最后不得不针对 gfx906 的 64 线程波前重新写一条路径。后面如果继续写，W4A16、通信、GDN、MoE、FP8 和 FP4 都可以各自拆成单独一篇。 本文一直使用同一份官方 vLLM 源码作为基线： ```text vLLM: 0.27.2rc1.dev26+g399f97424.d20260819 源码位置: v0.27.2rc0 之后第 26 个提交 Commit: 399f97424e61ea918135bd80176491646cb69d17 ``` 该提交对应 vLLM 上游 `399f97424e61ea918135bd80176491646cb69d17`，便于读者复现实验时锁定完全相同的源码基线[2]。 主要测试环境如下： ```text GPU: Hygon Z100 GPU 架构: gfx906 DTK: 26.04 PyTorch: 2.9.0 HIP: 6.3 张量并行: TP2 ``` 除特别说明外，后文涉及的吞吐、延迟、逻辑带宽、错误复现和 token ID 一致性结果均来自上述环境中的实际测试，而不是引用文献中的性能数据。 这里先把几个后面反复出现的名字解释清楚。DTK，也就是 DCU Toolkit，是海光 DCU 使用的软件工具栈，里面包括编译器、HIP 运行时、数学库、通信库以及调试和性能分析工具。vLLM 上游把这一类设备放在 ROCm 平台分支里处理，所以后文提到“ROCm 路径”，指的是 vLLM 内部的 ROCm 后端，并不等于这台机器直接安装和使用 AMD 官方完整的 ROCm 发行版。HIP 是一套 GPU 并行编程接口和运行时，可以粗略理解成和 CUDA 编程模型比较接近的一套接口体系[4]；`gfx906` 则不是显卡商品型号，而是编译器识别 GPU 指令架构时使用的目标名称[3]。 这一阶段主要拿一个 27B 的 AWQ INT4 模型做回归。AWQ 的全称是 Activation-aware Weight Quantization，中文一般叫**激活感知权重量化**，核心思路是把模型中占显存大头的权重压到较低位宽，同时利用激活统计识别并保护更重要的权重通道，从而减小低比特量化带来的精度损失[5]。这里的 INT4 表示主要线性层权重以 4 比特整数保存，计算时再结合缩放因子和零点恢复其数值。到这篇文章写下来的时候，最直观的一组首尾数据是： ```text 最初未适配，只能即时执行： ~18.7 tok/s 当前完整 Decode 图执行： ~33.32 tok/s ``` 这里的 `tok/s` 是 tokens per second，也就是**每秒生成多少个 token**。token 可以粗略理解为模型内部处理文本时使用的基本符号单元，它不一定正好对应一个汉字或者一个英文单词。18.7 到 33.32 大约是 1.78 倍，但这个 78% 绝不能理解成“某一个神奇内核单独带来了 78%”。中间同时经历了正确性修复、图执行打通、双卡通信优化、量化线性层优化和若干小算子融合，后文凡是能够真正做到开关对照的项目，我会把单项数据分开写。 另外还要区分两个测试口径。一个是**微基准测试**，英文通常叫 Microbenchmark，也就是把某一个 GPU 内核或者一小段执行路径单独拿出来测；另一个是**端到端测试**，英文简称 E2E，也就是把完整模型、调度、通信和实际执行链全部跑起来测最终吞吐。微基准非常适合找瓶颈，但是它也特别容易制造错觉。这轮适配已经不止一次证明，某个内核快 30% 不代表完整模型也会快 30%，甚至可能出现“内核单独测试又快又正确，真正放进服务以后根本没有命中这条路径”的情况。 ## 一、第一阶段先不谈性能，先建立一条可信的正确性基线 大语言模型推理通常可以粗略分成两个阶段。第一个阶段叫**预填充**，英文是 Prefill。用户一次输入几百、几千甚至几万个 token，模型需要把整段提示词跑一遍，同时建立后续生成要使用的 KV Cache。KV Cache 中文可以叫**键值缓存**，它保存注意力机制里已经计算过的 Key 和 Value，这样生成下一个 token 时就不用把之前所有 token 的 K、V 全部重新算一遍。第二个阶段叫**逐 token 解码**，英文是 Decode。预填充完成以后，每一步通常只新增一个或者少量 token，然后模型不断生成下一个 token。如果用一句不严格但直观的话来形容，Prefill 像“一次把题目读完”，Decode 则像“读完题以后开始一个字一个字写答案”。 这两个阶段对 GPU 的要求不一样。Prefill 通常更接近大规模矩阵乘法，计算密度较高；Decode 在 batch 较小时，大量线性层会变成矩阵向量乘法，或者 M 很小的矩阵乘法。这里顺便解释一下后文经常出现的 M、N、K：假设有 `[M,K] × [K,N]` 这样一个矩阵乘法，M 可以粗略理解成当前同时处理多少行输入，K 是输入特征维度，N 是输出特征维度。单请求 Decode 时，M 经常只有 1，所以每生成一个 token 都要重新从显存中读取一大块权重，却只有很少的计算可以摊薄这些读取成本，这类低算术强度工作负载往往更容易受到显存带宽而不是峰值浮点算力的限制，这与经典 Roofline 性能模型对 memory-bound 工作负载的分析是一致的[6]；在实际服务中，GPU 内核启动延迟和双卡通信延迟还会进一步叠加。 HBM 是 High Bandwidth Memory，中文就是**高带宽显存**。可以把它理解成 GPU 旁边为高吞吐计算设计的高速显存。后文还会不断出现 GPU kernel，这里的 kernel 指**GPU 内核函数**，不是 Linux 的操作系统内核，而是一段被提交给 GPU 并行执行的程序。当前适配主要会遇到三类实现：PyTorch 原生算子组合、Triton 编写的 GPU 内核，以及手写 HIP C++ 内核。Triton 是一种面向分块张量程序的 GPU 编程语言与编译器，其设计目标之一就是以比直接手写底层 GPU 代码更高的开发效率生成高性能自定义算子[7]；但最后仍然需要编译器把它变成目标 GPU 真正执行的机器指令。 这次在 gfx906 上很早就碰到了一个必须时刻记住的问题：**Triton 编译成功，不等于结果一定正确。** 有些内核可以正常编译、正常启动，执行时间甚至也很合理，但数值就是错的。相比直接崩溃，这种错误更难排查，因为程序不会给你任何明显提示。尤其是大模型生成本身带有随机性，如果最后生成的文字还能看，很容易把底层数值错误误认为正常的采样波动。 最初遇到的问题很多。gfx906 会被部分架构判断错误地归到后续架构路径；一些预编译自定义算子能够运行，但是结果错误；部分 Triton 注意力内核可以编译，却产生明显的数值误差；GDN/Linear Attention 的一些内核出现 NaN 或错误数据；KV Cache 写入路径和 HIP 图捕获不兼容；多进程启动之前父进程已经提前碰过 HIP 运行时；vLLM 自定义全归约默认又不支持 gfx906；与此同时，AWQ/W4A16 的小 M 线性层也远远没有把 Z100 的 HBM 带宽利用起来。所以最早一段时间与其说是在“优化”，不如说是在建立一条可信的正确性基线。 这里最危险的一类问题叫**静默错误结果**，英文常写成 silent wrong result。意思是程序既不崩溃，也不报异常，输出张量却已经错了。后来我给自己定了一个很简单的原则：只要某条预编译快速路径在 gfx906 上出现过无法解释的数值错误，就先关掉，宁可退回速度稍慢但已经验证正确的实现，等完整模型可以稳定复现以后，再逐个把热点重新做快。这个原则看起来比较保守，但实际省了很多时间，因为底层性能优化最怕参考结果本身就是错的，那样你后面得到的所有“加速”都失去了意义。 最先需要处理的是架构判断。vLLM 官方 ROCm 路径主要面向更新的 AMD 数据中心 GPU，一些代码会根据 `gfx` 系列判断硬件能力。问题在于 gfx906 和后来的 CDNA GPU 虽然名字可能都落在 `gfx9` 这个大范围里，但不能认为它们具有相同的指令和内核前提。CDNA 是 AMD 面向高性能计算和人工智能负载的一条 GPU 架构路线，后来的 MI 系列就属于这条路线。如果某处简单写成“只要是 gfx9 就进入某一条 CDNA 快速路径”，后面就可能自动选择一个根本没有在 gfx906 上验证过的分页注意力或者其他内核。 分页注意力，英文是 Paged Attention，主要解决 KV Cache 如何高效管理的问题。它不要求一条序列的 KV Cache 在显存中占据一块完整连续区域，而是像操作系统管理内存页一样分块管理缓存，这样更适合服务大量不同长度的请求。问题是，如果平台能力判断错了，后面选择的分页注意力内核可能从一开始就不适合 gfx906。这类错误很烦，因为最后看到的可能只是某个 Attention 内核报错，真正的根因却发生在几十层调用以前的一次架构分类。 所以现在适配树里给 gfx906 单独增加了能力判断。后面选择 W4A16、注意力、MoE、自定义全归约等路径时，不再只是问“是不是 gfx9”，而是继续看当前架构、数据类型、M/N/K、量化格式、内存布局、零点形式以及对齐条件。我一直尽量避免另一种做法，就是根据模型名字写特判。如果今天 Qwen 能跑，明天换一个模型所有地方都重新加 `if model_name == ...`，那只能算把某一个模型临时扶起来，不能叫真正完成了 gfx906 后端适配。 编译侧同样有一层兼容问题。当前 DTK 26.04 使用 PyTorch 2.9.0，而本文这份较新的 vLLM 源码在稳定 libtorch 接口等地方已经面向更新版本，因此 `_C_stable_libtorch` 和部分 ROCm 扩展代码不能原样使用。这里会涉及 ABI，也就是 Application Binary Interface，中文叫**应用二进制接口**。简单说，Python 层看起来两个函数名字一样，不代表已经编译好的 C++ 模块就能直接互换，因为参数怎么传、符号叫什么、类型在内存里怎么表示，都属于 ABI 的一部分。这一层最后补了稳定张量接口、类型分派、ROCm BLAS 句柄获取、显存查询、CMake Torch 版本判断以及 MIOpen 链接差异等兼容代码。这些代码不会直接增加 tok/s，但如果它们没有解决，后面的自定义 HIP 内核根本接不到真正的运行路径里。 多进程启动则是另一类问题。为了做 TP2，也就是**两路张量并行**，模型的一部分矩阵会分别放在两张卡上计算，因此 vLLM 需要启动多个工作进程。把 Transformer 层内的大矩阵按设备切分、再通过集合通信合并结果，是大模型张量并行中的经典做法[8]。当时遇到过 `Cannot re-initialize CUDA in forked subprocess`。虽然报错里写 CUDA，但 PyTorch 的 ROCm 后端沿用了很多 `torch.cuda` 接口名称。开始我以为只是 worker 启动顺序有问题，后来逐层插入探针才发现，父进程在真正加载模型之前就已经提前触碰了 HIP 运行时。 原因之一是环境里的 `amdsmi` Python 包不可用以后，一些设备查询会退回 `torch.cuda.is_available()`、`torch.cuda.device_count()` 和 `torch.cuda.get_device_properties()`。这些函数看上去只是问“有没有 GPU”“有几张 GPU”，实际上可能已经初始化一部分加速器运行时。Linux 的 `fork()` 会复制父进程地址空间，但 GPU 运行时内部维护的状态并不是普通内存，不能指望复制以后子进程继续若无其事地初始化。可以通俗地理解成父进程先把 GPU 驱动“开机开到一半”，然后突然复制出一个子进程，子进程拿到的是一份看似存在、实际不完整的半初始化状态，后面再继续启动 GPU 就容易出问题。 后来启动阶段能够直接从 sysfs 获取的信息尽量全部改成 sysfs。sysfs 是 Linux 通过 `/sys` 暴露硬件设备信息的一套文件系统，读取设备架构、设备数量、显存容量等信息时不需要先建立 HIP 上下文。继续往下追，还找到过 `flash_attn` 在导入时查询设备属性、Triton/FLA 导入时获取当前目标、Transformers 图像处理链检查 FlashAttention 是否可用，以及 Inductor 性能测试代码间接触发 `torch.cuda.is_available()` 等入口。把这些点逐个处理以后，TP2 的工作进程才终于能够比较稳定地经过通信器初始化、模型加载、KV Cache 创建和图捕获。 ## 二、注意力、GDN 和 HIP 图执行：能跑起来以后，还要保证图里也能算对 当前 DTK/Z100 组合下，一部分 vLLM 预编译 C++ 自定义算子在 gfx906 上数值并不可靠，我实际碰到过 RMSNorm、融合加法 RMSNorm、SiLU-and-mul 和旋转位置编码。这里的“自定义算子”指 vLLM 自己编译的 C++/HIP 扩展，而不是直接由 PyTorch 内置算子完成的计算。最初先让这些算子退回已经验证正确的通用实现，等正确性稳定以后再重新融合。 RMSNorm 的全称是 Root Mean Square Normalization，中文一般叫**均方根归一化**，由 Zhang 和 Sennrich 提出，其核心是使用均方根进行重缩放而省去 LayerNorm 的重中心化步骤[9]。它可以粗略写成： ```text y = x / sqrt(mean(x²) + eps) * weight ``` 这个算子本身计算量不大，问题在于如果直接用 PyTorch 原生张量操作拼出来，可能被拆成类型转换、平方、归约、求均值、倒数平方根和多次乘法等一串很小的 GPU 内核。单个内核几微秒看上去无所谓，但几十层模型、每生成一个 token 都执行一遍，累计起来就是实打实的固定开销。因此正确性跑稳以后，又针对 gfx906 写了 Triton RMSNorm 和融合加法 RMSNorm，把原来的一串小操作重新合成一个内核。 这里又碰到过 gfx906 Triton 代码生成的问题，例如归约得到的标量向向量广播时出现错误，以及部分 `tl.atomic_add` 写法不能可靠工作。最后采用的处理方法并不“漂亮”：先把归约结果写到一个临时工作区，经过同步以后再重新读出来，避开已经确认有问题的广播形式。做底层适配以后我越来越不在意某段代码从形式上够不够优雅，目标硬件真正执行什么指令、结果是不是可靠，比代码看上去“高级”重要得多。 SiLU-and-mul 也是类似情况。SiLU 是一种激活函数，形式是 `x * sigmoid(x)`；很多大模型的前馈网络会先把一条分支做 SiLU，再和另一条分支相乘，所以通常会把这一串操作融合成一个内核。原来的预编译实现在 gfx906 上不可靠以后，也补了一份 Triton 版本，避免重新退化成多个逐元素小内核。 注意力部分的问题更典型。标准注意力在 Decode 时会不断读取前面提到的 KV Cache，而分页注意力负责高效管理这些缓存[1]。从更一般的 GPU 实现角度看，Attention 的性能与 HBM 和片上存储之间的数据搬运高度相关，FlashAttention 的 IO-aware 分析对此给出了系统性的论证[10]。我实际测到过某些 ROCm Triton 注意力路径虽然能够正常编译和启动，Prefill 的最大绝对误差却能达到大约 1.4，Decode 也出现过 0.1 量级误差，这已经远远超出了 FP16 正常的舍入误差。其中一个比较明确的问题和计算分块大小有关，某些 `BLOCK=128` 的配置在 gfx906 上可以成功生成内核，却不能稳定保证数值正确。后来 Prefill 的一条相对安全配置收敛到了 `BLOCK=32、num_warps=4`，并重新覆盖 FP16、BF16、因果注意力、非因果注意力以及不同序列长度做回归。 这里的 `num_warps` 虽然沿用了 Triton/CUDA 风格的名字，但在 AMD 后端最终会映射到目标 GPU 的线程组织方式。简单理解，它控制一个程序实例使用多少组协同执行的线程。分块太大、线程组织不合适，不仅可能影响寄存器和性能，在旧架构编译器上甚至可能暴露代码生成错误，所以不能因为某个参数在新卡上常见就直接照搬。 KV Cache 写入路径也曾经影响图执行。有一条原生 HIP C++ 的 `reshape_and_cache` 路径需要通过 PyTorch 稳定 C++ 接口传递 `kv_cache_dtype` 这样的字符串参数，在当前 DTK 组合下，这条调用与图捕获并不可靠，所以 gfx906 后来改用了 Triton cache writer，也就是由 Triton 内核直接完成 KV Cache 的写入。这件事情很容易查错方向，因为表面看最后出问题的是 Attention，实际上如果 K/V 一开始写进缓存时内存布局就已经错了，后面的注意力内核本身再正确也没有任何意义。 GDN/Linear Attention 比普通注意力更麻烦。这里的 GDN 指 Gated Delta Network / Gated Delta Rule 一类**带门控状态更新的线性注意力结构**。Gated DeltaNet 将门控机制与 delta rule 结合，用固定大小的递归状态承载历史信息，是这一类结构的代表性工作[11]。它和传统全注意力最大的系统差别之一，是除了当前输入输出之外，还需要维护一份递归状态。当前测试的 27B 模型结构并不是 64 层全部使用传统全注意力，而是 48 层 Linear Attention/GDN 加 16 层 Full Attention。可以把普通 Transformer 层想象成“这一页算完就结束”，而 GDN 更像一本连续记账的账本：当前 token 算完以后还必须把状态准确保存下来，下一个 token 要接着用。 当前随源码带来的 FLA 和 causal-conv1d 中，多条相关内核在 gfx906 上出现过 NaN 或错误结果，包括分块 gated delta、递归状态更新、卷积后处理和因果一维卷积。因此后来单独增加了一套 `gfx906_fallback.py`，先建立数学正确而且能够进入图执行的路径。这里所谓“能够进入图执行”要求比普通 Eager 路径严格得多：前向计算中不能随便调用 `.item()` 把 GPU 张量读回 CPU，不能根据 GPU 张量临时转成 Python 的布尔值或者整数，也尽量不能插入主机端同步，因为一旦 CPU 必须在中间等待 GPU 返回结果，整段图捕获就可能失效。后来又逐步补了 Decode 融合路径、Prefill 顺序扫描、分块路径、混合 Prefill/Decode 的图元数据以及静态状态缓冲区，做到这里以后，状态型 Attention 才算真正能够进入完整图执行。 这里再解释一下所谓 HIP Graph。普通**即时执行模式**，英文叫 Eager Execution，程序每运行到一个算子，就由 CPU 向 GPU 提交一次内核。Decode 一步里经常有大量很小的内核，CPU 调度和内核启动本身的固定成本很容易积累。HIP Graph 可以把一整段会重复执行的 GPU 工作预先“录制”下来，这一步叫图捕获；以后不再让 Python 和 CPU 一条一条重新发命令，而是直接把整张图重新执行，这一步通常叫图回放。HIP 官方文档将这套机制定义为通过 `hipStreamBeginCapture()` / `hipStreamEndCapture()` 把流上的操作捕获成图，再实例化和重复启动[4]。在 CUDA 平台上对应 CUDA Graph，所以 vLLM 源码和日志里即使在 ROCm 环境也经常还能看到 CUDA Graph 这个名字。 最初 Z100 上的状态其实很简单： ```text 图执行不可用 只能使用 --enforce-eager 27B TP2 Decode ~18.7 tok/s ``` 平台识别、注意力、GDN、KV Cache 和图执行相关问题陆续处理以后，完整 Decode 图终于能够稳定运行，中间基线大约到了 25.15 tok/s。从 18.7 到 25.15 大约提升了 34%，不过我并不愿意把它直接写成“HIP Graph 单项提升 34%”，因为两次测试之间同时修掉了不少错误路径。更准确的说法是：系统从“基本未适配、只能即时执行”走到了“经过第一轮底层修复，可以稳定执行完整 Decode Graph”。真正能够严格做开关对照的项目，才适合单独计算性能收益。 ## 三、M-RoPE 和双卡全归约：两次被漂亮微基准骗到的经历 M-RoPE 是这一阶段一个很典型的小内核链优化。RoPE 的全称是 Rotary Positional Embedding，中文一般叫**旋转位置编码**，它通过旋转矩阵把位置信息编码进 Query/Key，并在注意力中自然引入相对位置信息[12]。M-RoPE（Multimodal Rotary Position Embedding）则把 RoPE 扩展到多维位置索引；例如 Qwen2-VL 将位置分解到时间、高度和宽度等维度，以统一处理文本、图像和视频的位置关系[13]。当时一次 M-RoPE 在真实执行过程中会展开成大约 25 个即时执行的小内核，包括位置索引、切片和拷贝、cos/sin 选择、乘加以及拼接等操作。这些操作本身都不重，问题是内核启动太碎，所以后来把它们融合成了一个 Triton 内核。 第一版微基准非常漂亮，放进完整模型以后却几乎没有任何收益。继续追真实执行路径才发现，正式服务传进来的 `positions` 并不是底层内存连续的张量，而是一个带 stride 的 view。stride 中文通常叫**步长**，它描述张量在各个维度上移动一个元素时，底层内存地址需要跨过多少位置。PyTorch 的切片、转置经常只修改形状和步长，并不会真的把数据重新拷贝成连续排列。第一版包装层要求 `positions.is_contiguous() == True`，也就是只接受连续张量，因此真实服务每次都悄悄退回了原实现。内核写得再快也没有意义，因为生产路径根本没进去。 后来把内核改成能够直接读取真实 stride 的版本，图回放微基准从 84.3 微秒降到了 11.1 微秒，完整 Decode 则从大约 25.0 tok/s 提升到了大约 26.3 tok/s，整模提升约 5.5%。同时在固定 greedy decoding，也就是每一步都选择概率最高 token 的条件下，连续生成 60 个 token，与 Eager 基线的 token ID 逐个完全一致。这里所谓 exact token ID，就是不只比较“最后生成的文字大致相同”，而是要求每一步生成的 token 整数编号完全一致。做 GPU 内核正确性回归时，这比肉眼看两段文本是不是“差不多”严格得多。 双卡通信则给了我另外一次教训。TP 是 Tensor Parallel，中文叫**张量并行**。TP2 表示使用两张卡，把一层中的大矩阵拆到两张卡分别计算。有些层每张卡算完自己的局部结果以后，必须执行 All-Reduce，也就是**全归约**：各卡交换并合并结果，最后让两张卡都得到完整输出。可以通俗地理解成两个人各算了一半，最后还必须把答案交换、求和，再各自拿到最终结果。Decode 时这种通信会发生很多次，而且消息往往只有十几 KB，这时决定性能的未必是峰值带宽，几十微秒的固定通信延迟反而非常重要。 原 PyNCCL 路径的小消息延迟大致是： ```text 10 KB ~77.9 us 50 KB ~77.3 us 100 KB ~110.3 us ``` 这里的 PyNCCL 可以理解成 vLLM/PyTorch 侧调用 NCCL/RCCL 风格集合通信能力的一条 Python 封装路径。NCCL 是 NVIDIA 的 GPU 集合通信库，RCCL 则是 ROCm 生态中承担类似职责的通信库；RCCL 官方文档明确提供 All-Reduce、All-Gather、Reduce-Scatter、点对点收发等多 GPU/多节点通信原语[14]。vLLM 自己还有一套 Custom All-Reduce，也就是**自定义全归约**，主要针对小消息减少通用通信库的固定开销，不过上游默认主要向更新 AMD 架构开放，gfx906 并没有直接启用。 最开始把 gfx906 放进这条路径以后，10 KB 和 20 KB 的微基准很快就到了 30 微秒左右，看起来非常好。但真正放进完整模型以后，生成 token 很早就开始和基线不一致。最麻烦的是固定消息大小连续运行很多次又可以全部正确，于是后来专门写了一个不断改变消息大小的测试，让消息按照 8 KB、10 KB、16 KB 等尺寸连续变化，问题马上暴露出来。 有一次期望结果是 5，实际拿到的是 4。继续把各卡贡献拆开以后发现，这个 4 恰好等于“本卡这一轮的新数据，加上对端上一轮全归约留下来的旧数据”。也就是说，当前 GPU 读到了 stale data，中文可以叫**陈旧数据**。这里不是说显存里没有写入新值，而是由于缓存、一致性或者同步时序的问题，一张卡在某个时刻仍然看到了另一张卡上一次操作留下来的旧内容。ROCm 的内存模型明确区分 coarse-grained 与 fine-grained coherence：前者通常需要同步边界才能保证其他执行主体看到最新数据，而后者强调更强的系统级可见性[15]。普通 CPU 单线程程序很少需要直接面对这种情况，但多 GPU 的 P2P（Peer-to-Peer，设备间直接访问）和 IPC（Inter-Process Communication，进程间共享资源）场景必须认真考虑跨设备内存可见性。 最后这条路径使用了 fine-grained staging buffer，可以译成**细粒度中转缓冲区**。它使用更适合跨设备一致性访问的内存作为中转，但只准备一个地址仍然不够，因为下一轮全归约会立刻重新使用同一块内存。最终又做成了一个环形中转缓冲区： ```text slot0 -> slot1 -> slot2 -> ... -> slot15 -> slot0 ``` 也就是说，连续几轮通信不会马上重新落到同一个地址。图回放又带来了下一层问题：Python 在图捕获时执行过一次的槽位计数代码，真正回放时不会重新执行，因此槽位选择也不能只依赖 Python 端的计数，最后把这一部分代数放到了设备端。现在这条路径已经通过连续改变消息大小、不断改变输入数据、FP16/BF16/FP32、即时执行、图回放以及完整 27B token ID 一致性回归。正确版本的小消息延迟大约是 10 KB 29.2 微秒、20 KB 29.4 微秒、40 KB 29.3 微秒、100 KB 45.1 微秒；完整模型关闭自定义全归约时约 25.15 tok/s，开启后约 27.30 tok/s，提升大约 8.5%。 这次通信排障最后让我觉得，真正有价值的并不是“30 微秒”这个数字，而是测试方法本身。固定大小的微基准当时确实会骗人，只有“输入不断变化 + 消息大小不断变化”才把真正的跨卡旧数据问题暴露出来。如果只满足于第一组漂亮的延迟数字，这个错误很可能会一直带进实际服务。 ## 四、W4A16：固定开销降下来以后，真正开始碰 Z100 的显存带宽 前面的图执行和通信固定开销逐渐压下来以后，AWQ INT4 线性层在整体运行时间里的占比越来越明显，于是开始真正碰 W4A16。W4A16 的意思是 **Weight 4-bit / Activation 16-bit**，中文就是**4 比特权重、16 比特激活**。权重使用 INT4，激活仍然通常是 FP16 或 BF16。FP16 是 16 比特浮点数；BF16 也是 16 比特浮点格式，只是它给指数保留了更多位，因此数值范围更接近 FP32。 量化权重通常还需要 scale 和 zero point，中文分别叫**缩放因子**和**零点**。可以把 INT4 权重想象成只保存了一个很小的“刻度编号”，缩放因子告诉内核每个刻度有多大，零点则告诉它这个刻度表从什么位置开始。因此 W4A16 内核不能只是把 4 比特权重直接送进普通 FP16 矩阵乘法，而是要先读取打包后的 INT4、解包、读取缩放因子和零点、恢复权重对应的数值，再和激活做乘加。也正因为如此，“权重从 16 比特压到 4 比特”并不等于推理自动快四倍，如果解包、地址计算和量化参数处理太重，省下来的显存流量很快就会被额外算术指令吃掉。 为了知道还有多少空间，我先测了一下 Z100 自己的有效 HBM 上限。简单 copy/add 测到大约 890 GB/s。这里可以把这个数字看成一个**显存带宽参考上限**，也就是 memory-bound，也就是主要受显存读取限制的内核，理论上最多能够向这个数量级靠近，并不是说任何内核都应该跑到 890 GB/s。当时 M=1 的 Triton W4A16 代表数据如下： | 投影类型 | 矩阵形状 | 延迟 | 逻辑带宽 | | --- | --- | ---: | ---: | | QKV 投影 | N=8192, K=5120 | ~63.3 us | ~394 GB/s | | Gate/Up 投影 | N=17408, K=5120 | ~105 us | ~504 GB/s | | Down 投影 | N=5120, K=8704 | ~74.5 us | ~356 GB/s | 这里顺便解释一下几个名字。QKV 投影是注意力层把输入映射成 Query、Key 和 Value 的线性变换；Gate/Up 是大模型前馈网络中门控分支和扩维分支的线性投影；Down 则是把中间扩大的隐藏维度重新投影回模型隐藏维度。表里的“逻辑带宽”不是显存控制器实际统计到的真实流量，而是按照算法最低需要读写的数据量除以执行时间得到的估算值，主要用于比较不同内核。它不能直接说“显存控制器已经跑到了 504 GB/s”，但从 350～500 GB/s 与 890 GB/s 之间的差距仍然可以判断，这条路径还有明显优化空间。 一开始当然还是继续扫 Triton 参数，例如分块宽度、K 维分块、并行线程组数量和 `waves_per_eu`。但是扫到后面收益越来越小，性能分析和机器指令里开始出现一个明显信号：最快几组 Triton 内核使用了很多 VGPR。VGPR 的全称是 Vector General Purpose Register，中文叫**向量通用寄存器**，是 AMD GPU 每个波前用来暂存激活、权重、部分累加结果等数据的高速寄存器。寄存器很多当然方便，但一颗计算单元上的寄存器总量有限，一个波前吃掉太多寄存器，就会减少同一时间能够驻留的波前数量。 这里又会涉及 Occupancy，中文通常叫**驻留率**。它表示一个计算单元上实际能够同时驻留多少波前，相对于硬件能力处于什么水平。驻留率并不是越高越好，因为为了强行塞进更多波前，如果每个波前可用寄存器不够，编译器就可能发生 register spilling，也就是**寄存器溢出**：原本应该放在寄存器中的临时数据被迫放到 local memory。local memory 虽然名字有“local”，实际速度远远不等于寄存器，很多情况下仍然会落到显存体系。 当时在 gate 形状上强行提高 `waves_per_eu`，结果很夸张： ```text wpe=1 ~106 us wpe=2 ~798 us wpe=3 ~1256 us wpe=4 ~1773 us ``` 这个结果已经足够说明，问题根本不是“驻留率还不够高”，而是强行提高驻留波前以后发生了严重的寄存器溢出。继续在 Triton launch 参数上打转已经没有太大意义，内核结构本身需要换。 gfx906 的 wavefront，也就是**波前**，宽度是 64 个线程，所以通常叫 wave64；ROCm 的硬件规格表对 gfx906（GCN 5.1）给出的 Wavefront Size 即为 64，同时列出了 64 KiB LDS 和对应的 VGPR 文件规模[3]。AMD GPU 中一个波前是一组按照相同指令节奏协同执行的线程，NVIDIA 里与之类似的概念叫 warp，但常见 warp 是 32 线程。vLLM 原来的 skinny INT4 内核里其实已经有一些很不错的设计，例如打包 INT4、非临时加载、LDS 激活复用、half2 和 FP32 累加，但原来的快速路径更多面向 wave32 和更新架构，不能直接照搬到 gfx906。 LDS 的全称是 Local Data Share，中文可以叫**局部数据共享存储**，功能上可以类比 CUDA 的 Shared Memory。它是一个工作组内部共享的片上高速存储。如果把 HBM 比作远处的大仓库，那么 LDS 就像工位旁边的小料箱。一份 activation 后面会被多个输出行反复使用，与其每次都重新去 HBM 取，不如先搬进 LDS，再让同一工作组内的波前反复复用。 后来新增的 gfx906 wave64 W4A16 内核，大致就是先把 activation 搬到 LDS，打包 INT4 权重采用流式方式读取，每个波前负责一定数量的输出行，把缩放因子和零点提前整理成更适合输出维度访问的布局，用成对 FP16 数据做乘加，最后使用 FP32 累加保证数值稳定。M=1 时的代表结果如下： | 形状 | Triton | HIP wave64 | | --- | ---: | ---: | | QKV | ~63 us | ~50 us | | Gate/Up | ~105 us | ~96.6 us | | Down | ~74.5 us | ~51 us | 同一组内核又专门在图回放环境里重新测了一遍，因为即时执行微基准快，不代表捕获进图以后优势一定还在。图回放时 QKV 从 66.9 微秒降到 55.4 微秒，Gate 从 110.0 微秒降到 102.4 微秒，Down 从 91.7 微秒降到 55.3 微秒，这说明这条 wave64 路径在真实图执行方式下仍然保留了优势。 内核稳态速度提升以后，模型加载阶段又冒出了两个很重的权重重新打包开销。所谓 repack，就是**重新打包权重布局**。模型权重文件里的排布通常优先考虑存储、兼容和通用性，并不一定适合目标 GPU 最终读取。第一处是运行时 packed-weight 的构造，旧实现由 Python 沿 K 维分成很多小块，一层会启动大量内核，后来换成一个二维 Triton 重新打包内核后，QKV 从约 13.75 毫秒降到 0.47 毫秒，Gate 从约 8.02 毫秒降到 0.30 毫秒，Down 从约 12.95 毫秒降到 0.15 毫秒。 第二处是模型文件中 INT4 布局的转置和重新打包。旧逻辑近似先把 `[N,K/8] int32` 完整解包成 `[N,K] int32`，再转置、重新打包，对于大层来说，中间那个完整 int32 矩阵非常大。后来改成以 8×8 INT4 小块直接在 nibble 层面完成转置和重新打包。nibble 就是**半个字节，也就是 4 比特**，一个 INT4 数值恰好占一个 nibble。这样就不用为了转置先临时生成一整份巨大的 int32 权重。这类优化不会直接增加稳定 Decode 的 tok/s，但对模型加载时间和临时显存都很重要。 W4A16 这里还出现过一次很容易误判成显存泄漏的问题。完整图执行期间一度看到预计 Graph memory 只有约 0.09 GiB，实际 Graph pool 却占了约 7.35 GiB。第一反应当然是图内存分配器泄漏，继续沿分配时间追下去才发现，真正原因是 W4A16 的 packed weight、output-major scale 和 zero point 当时采用了 lazy build，也就是**延迟构造**：某一层第一次真正执行到时才创建。这些张量本来应该在整个模型运行期间长期保存，属于 persistent tensor，也就是**持久张量**，但它们第一次被创建时恰好处在图捕获阶段，于是图内存分配器把它们全部记到了 Graph pool 里。 换句话说，不是这些数据本身属于图，而是它们“出生的时间”刚好落在图捕获里，所以账被记错地方了。把这些表示提前移动到 `process_weights_after_loading()`，在模型权重加载结束以后就构造完成，预计 Graph memory 变成 0.09 GiB，实际 Graph pool 约 0.10 GiB，基本吻合。最终修的不是内存分配器，而是张量生命周期。 完成 wave64、图执行正确性和预打包生命周期处理以后，做了一组同一引擎下的多轮开关对照。开启 HIP wave64 的几次结果是 32.08、32.62、32.73、32.68、32.75 tok/s，中位数 32.68 tok/s；关闭以后中位数约 27.07 tok/s。也就是说，这条 W4A16 wave64 路径在完整 27B 测试里单项带来了大约 **20.7%** 的提升，而且所有测试都保持逐 token ID 完全一致。这是目前整个工程中最干净的一组端到端 A/B。 ## 五、单请求快了以后，还要处理 M=2、4、8 的并发情况 真实服务不可能永远只有一个请求。当 batch 从 1 增加到 2、4、8 时，前面矩阵乘法里的 M 也会逐渐增大。旧 Triton batch 路径中，QKV 的 M2、M4、M8 分别大约是 157、214、337 微秒，Gate 则大约是 273、374、617 微秒。如果最简单地把 M=1 的矩阵向量乘重复执行 M 次，同一份权重就会从 HBM 被反复读取很多遍，因此小 M 原生内核的核心思路变成：**权重只流一次，同时服务多行 activation。** 第一版 `BATCH_ROWS=4` 在 M=2 时表现很好，但到了 M=4 和 M=8，寄存器压力明显增大，后来收敛成 `BATCH_ROWS=2`，再按照具体矩阵形状做分派。继续分析以后发现 QKV 和 Gate 在 M≥2 时还有一些寄存器余量，于是又增加 `ROWS_PER_WAVE=2`，让一个波前同时计算两个输出行。目前一组代表结果是： | 形状 | M1 | M2 | M4 | M8 | | --- | ---: | ---: | ---: | ---: | | QKV | 55.06 us | 67.83 us | 129.82 us | 255.09 us | | Gate | 102.34 us | 123.62 us | 241.86 us | 478.17 us | | Down | 55.34 us | 131.89 us | 198.35 us | 319.24 us | Down projection 的情况不一样。它通常满足 K>N，也就是输入维度比输出维度大，M 增长以后，如果继续使用当前原生分块方案，很容易为了不同输入行重复读取权重，所以 M≥3 时反而不如 Triton batch。现在大致采用的策略是：QKV/Gate 在 M=1 时使用单行 wave64，在 M=2～8 时使用原生小 M 路径；Down 在 M=1 时使用 wave64，M=2 使用原生内核，M≥3 则重新交给 Triton batch。 这也是我现在越来越强调“按照形状选内核”的原因。英文源码里经常把这种行为叫 dispatch，正文里可以直接理解成**路径分派或者内核选择**。手写 HIP 并不天然比 Triton 高级，一个 HIP 内核在 M=1 上非常快，也不代表所有 M 都应该强制走 HIP。最终应该根据架构、数据类型、矩阵尺寸和内存布局决定哪条实现真正合适。 这一轮也做过不少最后没有留下来的实验。强制提高 `waves_per_eu` 会导致严重寄存器溢出；Split-K，也就是把 K 维拆给多个工作组分别计算再归约，理论上可以增加并行度，但实际工作区和最终归约成本把收益抵消了；把 zero point 从 int8 进一步压到 4 比特确实少读了一些显存，却增加了地址计算、位移和解包，Down 反而从约 55 微秒退化到约 70 微秒；把两个 32 位权重读取合成一个 64 位读取，看上去一次搬的数据更多，实际 QKV 和 Gate 都慢了几微秒；参数广播也尝试过，希望减少缩放因子和零点的重复读取，性能分析上看起来有道理，实测仍然是负收益。 这些负结果我觉得有必要保留，因为 GPU 优化里经常会看到一些听起来很正确的话，例如“读取越宽越快”“显存流量越少越快”“驻留率越高越快”。这些话单独拿出来都不成立。加载宽度、指令数量、地址计算、寄存器压力、驻留率、缓存行为和编译器调度必须放在一起看，最终还是只能通过真实目标硬件上的测量判断。 ## 六、FP8：新卡上的低精度优势，在 gfx906 上不能照搬 W4A16 第一轮基本闭环以后开始看 FP8。FP8 就是**8 比特浮点数**，常见格式包括 E4M3 和 E5M2。E4M3 表示 4 比特指数和 3 比特尾数，E5M2 则是 5 比特指数和 2 比特尾数；Micikevicius 等人的 FP8 工作系统讨论了这两类 8 比特浮点编码及其在深度学习中的使用[16]。另外在不同软硬件栈中还会遇到 FN、FNUZ 等具体编码规则。相比 FP16，FP8 可以进一步降低权重和激活的数据量，但前提是 GPU 本身具有高效的 FP8 转换和矩阵计算能力。 gfx906 在这里面对的最大问题不是“官方内核慢一点”，而是今天 ROCm 主流 FP8 快速路径基本围绕 MI300+ 这一代硬件建立。实际测试 `torch._scaled_mm` 时会直接遇到 ROCm MI300+ 的限制，vLLM 自己的 ROCm FP8 skinny kernel 也没有真正面向 gfx906 的实现。因此如果希望 Z100 跑 FP8，就必须重新评估：到底是让 FP8 权重一直保持压缩，每生成一个 token 时再软件转换，还是模型加载时一次性把它转换成 FP16/BF16。 为了确认第一条路有没有希望，我还是做了几种在线软件解码原型。在 N=8192、K=5120、M=1 的形状下，Triton 直接读取 FP8 字节，手工拆符号位、指数位和尾数位，最好大约 246.7 微秒，逻辑带宽只有 170.1 GB/s；改成 256 项查表，希望用查表代替大量位运算，结果约 277.6 微秒，反而更慢；HIP FP8 双元素转换再配合 `fdot2` 大约 403 微秒；手工把 FP8 转成 FP16 位模式以后再做 `fdot2`，甚至到了约 600 微秒。 `fdot2` 可以粗略理解成 AMD GPU 上一次对两个 FP16 通道做点积累加的指令能力。这里最有意思的是，写法越看起来“接近硬件”，速度反而越差。原因是当 GPU 没有真正高吞吐的原生 FP8 数据通路时，格式转换、打包、位操作和指令依赖已经把整个内核拖成了算术瓶颈。换句话说，这时真正卡住的不是 HBM 读取，而是为了把 FP8 软件“解释”成能参与计算的格式，ALU 已经忙不过来了。 所以现阶段比较现实的策略是：模型文件里仍然保存 FP8，加载模型时一次性反量化成 FP16/BF16，真正运行时直接使用已经优化过的 FP16/BF16 GEMV/GEMM。反量化就是把低精度表示恢复成更高精度表示。这样做的优点是每生成一个 token 不再重复支付 FP8 软件转换成本，缺点也很明显：FP8 原本节省显存的优势会大幅缩水。目前每张量缩放因子的合成数据正确性已经通过，但还没有拿真实 FP8 27B 权重做完整模型回归，所以现在还不能说“Z100 的 FP8 已经适配完成”。 有意思的是，FP8 加载时反量化又顺手暴露出了 FP16/BF16 的另一个性能缺口。N=8192、K=5120 时，原 gfx906 通用路径在 M2、M4、M8 上分别大约是 973.5、482.1、479.5 微秒，其中 M2 尤其差。后来补了一个**大 K、小 M 的分块多行 GEMV**：把 K 按大约 1024 一类大小切块，每个 GPU 工作组只保留一小段 activation，同时一次处理多个输出行，使用 FP32 完成乘法和累加。 第一轮测试里，M2 从 973.5 微秒降到约 139 微秒，M4 从 482.1 微秒降到约 138 微秒，M8 从 479.5 微秒降到约 234 微秒，大约分别是 7 倍、3.5 倍和 2 倍。后面又把 QKV、Gate、Down、FP16/BF16，以及每个程序实例同时处理多少输出行、K 维分块多大等组合补测了一轮，正式包装路径的 18 组 A/B 当前全部是正收益，代表提升约 1.25～4.7 倍，M2/M4 的逻辑带宽可以做到大约 0.6～0.75 TB/s。 M8 为什么没有 M2/M4 那么漂亮，用 rocprof 做性能分析以后就很清楚。rocprof 是 ROCm 生态里的性能分析工具，可以统计内核的寄存器使用、算术单元利用率和访存情况。当时 M8 的 VGPR 大约 112，VALU 利用率约 99.7%，MemUnitBusy 约 49%，MemUnitStalled 只有 2～3%。VALU 就是 Vector Arithmetic Logic Unit，中文可以叫**向量算术逻辑单元**，负责大量普通向量算术。接近 100% 的 VALU 利用率、很低的显存等待说明这里已经不是“HBM 没有吃满”，而是被算术指令和寄存器压力卡住了。后来又试过纯 FP16 `fdot2`、多波前工作组加 LDS 激活复用，都更慢，所以当前没有为了“手写 HIP 看起来更底层”就强行替换，仍然保留 Triton FP32-product 结构。 ## 七、Prefill 和 MoE 目前做到哪里，哪些地方还不能下结论 前面 Decode 写得比较多，是因为目前从 18.7 到 33.3 tok/s 这一条性能曲线最完整，但 Prefill 并没有被忽略。标准 Attention 的 Prefill 已经做过 gfx906 Triton 分块和线程组织修正、错误分页注意力路径的正确性回退、能够进入图执行的 KV Cache writer，以及混合 Prefill/Decode 场景下的元数据处理。GDN Prefill 当前已经有 gfx906 因果一维卷积路径、顺序扫描 Triton、分块 gated-delta、状态读写和混合批次完整图执行。 线性层方面，W4A16 现在最成熟的是 M=1 和小 M Decode，前面提到的 FP16/BF16 大 K、小 M 路径也已经开始进入正式执行路径。真正还没有完全收敛的是大 M Prefill。这里必须找到不同内核之间的 crossover，也就是**性能交叉点**。例如 M=2 时自定义 GEMV 可能明显更快，M=8 仍然领先，但 M 增长到 32 以后，通用的 rocBLAS 或 Triton GEMM 很可能已经反超。rocBLAS 是 ROCm 提供的 BLAS 数学库，其中包括高度优化的矩阵乘法实现；GEMM 是 General Matrix-Matrix Multiplication，也就是通用矩阵乘矩阵。 这个交叉点必须实测，不能因为小 M 内核在 Decode 很快，就一路强行使用到 Prefill。因此现在我不准备给一个“Z100 Prefill 已经多少 tok/s”的单数字。提示词只有 128 token 和上下文达到 32K 时，计算规模、Attention 开销和内存行为完全不是同一个问题。真正有意义的 Prefill 数据至少要同时列出提示词长度、分块大小、Attention/GDN 比例和主要矩阵形状。 MoE 也已经开始适配。MoE 的全称是 Mixture of Experts，中文叫**专家混合模型**。与普通 Dense 模型每个 token 都经过同一套前馈网络不同，MoE 会通过 router，也就是**路由器**，为 token 选择少数 expert，也就是**专家子网络**参与计算，从而形成稀疏激活的条件计算结构；Switch Transformer 是这一方向的代表性大规模实现之一[17]。这样可以让模型拥有非常大的总参数量，而每个 token 实际只激活其中一小部分参数。系统实现上却会复杂不少，因为除了专家矩阵乘法，还要处理 top-k 专家选择、token 按专家重新排列、对齐和填充，以及多卡之间的专家通信。 当前 gfx906 适配树已经处理了一些 MoE 问题，包括 `moe_align_block_size` 的回退、小批量对齐、小 top-k softmax/router、W4A16 专家小 M GEMV、打包 INT4 专家路径、`(w-zp)` 预计算、连续内存读取，以及错误的 `moe_sum` 和激活路径回退。有些局部微基准已经比早期内核快了不少，但真实 MoE 整模、Decode、Prefill、Graph、TP/EP 和多轮端到端还没有全部闭环。EP 是 Expert Parallel，中文叫**专家并行**，也就是把不同专家分布到不同 GPU 上。因此第一篇只能写“MoE 已经进入适配”，现在还不能把某几个局部内核数字包装成完整 MoE 性能。 ## 八、目前真正能够相信的性能数字，以及接下来还要做什么 为了避免不同开发阶段的数据混在一起，当前能够做比较干净开关 A/B 的项目主要有三项： | 优化项目 | 关闭 | 开启 | 变化 | | --- | ---: | ---: | ---: | | M-RoPE 融合 | ~25.0 tok/s | ~26.3 tok/s | ~+5.5% | | 自定义全归约 | 25.15 tok/s | 27.30 tok/s | +8.5% | | W4A16 HIP wave64 | 27.07 tok/s | 32.68 tok/s | +20.7% | 这三组数据来自不同开发阶段，因此不能简单把 5.5%、8.5% 和 20.7% 直接相乘，但每一项自己的开关对照是有参考意义的。最近一次完整 27B 回归使用 TP2、自定义全归约、完整 Decode HIP 图执行和 W4A16 正式路径，连续生成 60 个 token，三次运行分别得到 31.9529、33.3670 和 33.3235 tok/s，中位数为 33.3235 tok/s，同时三次生成的 token ID 完全一致。相比最初图执行不可用、只能 Eager 跑约 18.7 tok/s，累计大约是 1.78 倍，也就是提升约 78%。 需要特别说明的是，这组 33.32 tok/s 仍然是 batch=1，所以前面 M=2、4、8 的小 M 新路径基本不会反映在这个数字上，它们主要是在为真实并发服务做准备。也就是说，现在的单请求 Decode 曲线已经比较完整，并发服务还有单独的优化空间。 为了确认现在到底改到了什么程度，我重新把当前工作树和同一个官方 HEAD 做过一次差异统计，大约是 80 个文件发生变化、约 2.4 万行新增。不过这个数字里面包括完整补丁文件、安装脚本、构建脚本和兼容头文件，不能简单说成“写了 2.4 万行 GPU 内核”。真正和运行时有关的改动主要分布在平台和运行时、自定义算子、Attention/KV Cache、GDN/Linear Attention、RMSNorm/SiLU/M-RoPE、FP16/BF16 GEMV、W4A16、MoE、自定义全归约、HIP Graph 和多进程启动。做到现在，这已经不是简单给 vLLM 增加一个设备 ID，而是在补一个长期没有被完整覆盖的 gfx906 后端。 如果回头总结这一轮，我觉得真正值得记住的其实不是最后的 33.3 tok/s，而是一些排障经验。固定形状的微基准不能证明通信正确，自定义全归约只有不断改变输入和消息大小才暴露出旧数据；`torch.cuda.is_initialized() == False` 也不能证明父进程仍然适合 fork，真正需要追的是哪一行代码第一次触碰了 HIP 运行时；内核微基准正确不代表真实服务一定命中这条路径，M-RoPE 第一版就是典型例子；图捕获成功也不代表图回放一定正确，带状态模块、集合通信和延迟分配都必须专门测试不断变化的输入；HBM 没有跑满同样不代表一定是显存瓶颈，M8 大 K 场景已经出现向量算术单元接近满载、显存等待却只有几个百分点的情况。 同样，“读取更宽”和“读取更少”也都不是天然正确的优化方向。64 位读取实测反而比两个 32 位读取慢，4 比特零点虽然少搬了显存数据，却因为增加位移、解包和地址计算而变慢。GPU 性能优化到最后很少有一条可以脱离上下文成立的经验法则，真正能够相信的始终是目标硬件上带正确性验证的实测结果。 下一阶段会继续收 FP16/BF16 大 K 路径的图回放和大 M Prefill 交叉点，然后是真实 FP8 checkpoint、分块或逐通道缩放、加载时反量化、显存和 KV Cache 余量，再往后是 FP4/MXFP4、MoE 的完整 Decode/Prefill 和 TP/EP，以及 Attention/GDN 的分块 Prefill、FP8 KV Cache 和混合图执行。FP4 是 4 比特浮点数；MXFP4 可以粗略理解成带共享块级缩放（microscaling）的低比特浮点表示。Open Compute Project 的 MX 规范给出了这类 Microscaling Formats 的公开标准定义[18]。gfx906 没有新一代 GPU 上专门面向 FP4 的矩阵计算硬件，因此这里同样需要重新评估在线解包和加载时转换究竟哪条路线更现实。 自定义全归约也没有完全结束。目前主要验证的是 TP2，也就是两张卡，后面还需要继续看更多卡数以及不同硬件 topology。topology 中文就是**设备拓扑**，指 GPU 之间通过什么总线、交换芯片或者互连结构连接，不同拓扑会直接影响 P2P 和集合通信成本。与此同时还会继续拿 Qwen3.8-27B INT4 和其他模型交叉验证，确认现在的路径选择是不是真的由硬件架构、数据类型、量化格式、矩阵形状、内存布局和设备拓扑决定，而不是碰巧只对某一个模型有效。 ## 写在最后 如果只看最后结果，这一阶段当然可以压缩成一句话：Z100 上的 vLLM 从 HIP Graph 基本不可用、只能使用即时执行模式跑约 18.7 tok/s，做到 TP2 + 完整 Decode 图执行 + 自定义全归约 + gfx906 wave64 W4A16，当前 27B Decode 中位数约 33.32 tok/s。但真正占时间的地方，从来不是简单地把某个内核从 60 微秒改成 50 微秒，而是不断确认一些更基础的问题：这个内核在真实服务里到底有没有被调用？单独算是正确的，捕获进图以后还正确吗？固定输入正确，换输入以后会不会错？固定通信尺寸正常，连续切换尺寸以后还正常吗？父进程究竟在什么时候第一次碰到了 HIP？本来属于模型的长期权重为什么会被算进 Graph Pool？M=1 的快速内核到了 M=8、M=32 以后还应该继续使用吗？ 这些问题如果没有一层一层查清楚，很容易得到一套微基准数字非常漂亮、真正跑服务却不敢相信的代码。我觉得给老硬件做适配最有意思的地方也恰恰在这里。新 GPU 上很多问题已经被新指令、新编译器、新数学库和上游现成快速路径解决了，到了 gfx906 这种架构上，反而会逼着人重新回到最基本的问题：这个内核到底生成了什么机器指令，一块数据究竟从哪里读取，寄存器为什么会溢出，两个波前为什么会互相等待，一张卡写完的数据什么时候才能真正被另一张卡看到。很多平时被上层框架藏起来的东西，在这里又重新变成了必须亲手处理的问题。 下一篇我准备单独写 W4A16，从最开始的 Triton 基线和 HBM 带宽上限讲起，再展开 gfx906 wave64 内核、向量寄存器和驻留率、LDS 激活复用、INT4 解包、缩放因子和零点布局、小 M 权重复用，以及这一阶段那些“理论上应该更快、实际上却更慢”的实验。那一篇会比这一篇更偏 GPU 内核和机器指令，这一篇先把整个适配过程的第一阶段交代清楚。 ## 参考文献 > 参考文献采用顺序编码制，格式参照 GB/T 7714 的常用著录方式。论文优先给出正式会议/期刊或 DOI 链接；软件与硬件行为以官方源码、规范和文档为主。访问日期均为 2026-08-21。 [1] KWON W, LI Z, ZHUANG S, et al. Efficient memory management for large language model serving with PagedAttention[C]//Proceedings of the 29th Symposium on Operating Systems Principles (SOSP '23). New York: ACM, 2023: 611-626. DOI: [10.1145/3600006.3613165](https://doi.org/10.1145/3600006.3613165). [2] vLLM PROJECT. vLLM commit `399f97424e61ea918135bd80176491646cb69d17`[EB/OL]. GitHub, 2026-08-19. [https://github.com/vllm-project/vllm/commit/399f97424e61ea918135bd80176491646cb69d17](https://github.com/vllm-project/vllm/commit/399f97424e61ea918135bd80176491646cb69d17). [3] ADVANCED MICRO DEVICES, INC. GPU hardware specifications[EB/OL]. ROCm Documentation, 2026-02-20. [https://rocm.docs.amd.com/en/latest/reference/gpu-arch-specs.html](https://rocm.docs.amd.com/en/latest/reference/gpu-arch-specs.html). [4] ADVANCED MICRO DEVICES, INC. HIP Documentation: HIP runtime API, graph capture and memory coherence[EB/OL]. Release 6.2.41134. [https://rocm.docs.amd.com/_/downloads/HIP/en/docs-6.2.1/pdf/](https://rocm.docs.amd.com/_/downloads/HIP/en/docs-6.2.1/pdf/). [5] LIN J, TANG J, TANG H, et al. AWQ: Activation-aware weight quantization for on-device LLM compression and acceleration[C]//Proceedings of Machine Learning and Systems. 2024, 6. [https://proceedings.mlsys.org/paper_files/paper/2024/hash/42a452cbafa9dd64e9ba4aa95cc1ef21-Abstract-Conference.html](https://proceedings.mlsys.org/paper_files/paper/2024/hash/42a452cbafa9dd64e9ba4aa95cc1ef21-Abstract-Conference.html). [6] WILLIAMS S, WATERMAN A, PATTERSON D. Roofline: An insightful visual performance model for multicore architectures[J]. Communications of the ACM, 2009, 52(4): 65-76. DOI: [10.1145/1498765.1498785](https://doi.org/10.1145/1498765.1498785). [7] TILLET P, KUNG H T, COX D. Triton: An intermediate language and compiler for tiled neural network computations[C]//Proceedings of the 3rd ACM SIGPLAN International Workshop on Machine Learning and Programming Languages. New York: ACM, 2019: 10-19. DOI: [10.1145/3315508.3329973](https://doi.org/10.1145/3315508.3329973). [8] SHOEYBI M, PATWARY M, PURI R, et al. Megatron-LM: Training multi-billion parameter language models using model parallelism[EB/OL]. arXiv:1909.08053, 2019. [https://arxiv.org/abs/1909.08053](https://arxiv.org/abs/1909.08053). [9] ZHANG B, SENNRICH R. Root Mean Square Layer Normalization[C]//Advances in Neural Information Processing Systems 32. 2019. [https://proceedings.neurips.cc/paper/2019/hash/1e8a19426224ca89e83cef47f1e7f53b-Abstract.html](https://proceedings.neurips.cc/paper/2019/hash/1e8a19426224ca89e83cef47f1e7f53b-Abstract.html). [10] DAO T, FU D Y, ERMON S, et al. FlashAttention: Fast and memory-efficient exact attention with IO-awareness[C]//Advances in Neural Information Processing Systems 35. 2022. [https://arxiv.org/abs/2205.14135](https://arxiv.org/abs/2205.14135). [11] YANG S, KAUTZ J, HATAMIZADEH A. Gated Delta Networks: Improving Mamba2 with Delta Rule[C]//The Thirteenth International Conference on Learning Representations. 2025. [https://openreview.net/forum?id=r8H7xhYPwz](https://openreview.net/forum?id=r8H7xhYPwz). [12] SU J, LU Y, PAN S, et al. RoFormer: Enhanced Transformer with Rotary Position Embedding[EB/OL]. arXiv:2104.09864, 2021. [https://arxiv.org/abs/2104.09864](https://arxiv.org/abs/2104.09864). [13] WANG P, BAI S, TAN S, et al. Qwen2-VL: Enhancing vision-language model's perception of the world at any resolution[EB/OL]. arXiv:2409.12191, 2024. [https://arxiv.org/abs/2409.12191](https://arxiv.org/abs/2409.12191). [14] ADVANCED MICRO DEVICES, INC. RCCL Documentation[EB/OL]. Release 2.21.5. [https://rocm.docs.amd.com/_/downloads/rccl/en/docs-6.3.0/pdf/](https://rocm.docs.amd.com/_/downloads/rccl/en/docs-6.3.0/pdf/). [15] ADVANCED MICRO DEVICES, INC. GPU memory: Fine-grained and coarse-grained coherence[EB/OL]. ROCm Documentation 6.1.5. [https://rocm.docs.amd.com/en/docs-6.1.5/conceptual/gpu-memory.html](https://rocm.docs.amd.com/en/docs-6.1.5/conceptual/gpu-memory.html). [16] MICIKEVICIUS P, STOSIC D, BURGESS N, et al. FP8 formats for deep learning[EB/OL]. arXiv:2209.05433, 2022. [https://arxiv.org/abs/2209.05433](https://arxiv.org/abs/2209.05433). [17] FEDUS W, ZOPH B, SHAZEER N. Switch Transformers: Scaling to trillion parameter models with simple and efficient sparsity[J]. Journal of Machine Learning Research, 2022, 23(120): 1-39. [https://jmlr.org/papers/v23/21-0998.html](https://jmlr.org/papers/v23/21-0998.html). [18] OPEN COMPUTE PROJECT. OCP Microscaling Formats (MX) Specification, Version 1.0[S/OL]. 2023. [https://www.opencompute.org/documents/ocp-microscaling-formats-mx-v1-0-spec-final-pdf](https://www.opencompute.org/documents/ocp-microscaling-formats-mx-v1-0-spec-final-pdf).