DCU使用技术报告（上篇）：gfx936 DCU上的Qwen3.5-27B DTK/HIP环境、基线与性能画像实战

系列文章上篇。本文记录我们在单卡 gfx936 DCU 上优化 Qwen3.5-27B 推理服务的第一阶段：环境核对、评测基线、长上下文分块规律，以及如何拿到一份不掺启动噪声的 HIP Profile。

0. 先说结论

这次项目最重要的第一步，不是写 Triton，也不是翻 rocBLAS 的 solution library，而是把三个问题说清楚：

1. 服务实际加载的是哪份 vLLM 源码；
2. 8K、16K、32K 请求在 --max-num-batched-tokens 4096 下会拆成什么形状；
3. 正式推理期间，DCU 时间究竟花在 Attention、GEMM 还是别的算子上。

这三件事没搞清楚之前，任何“快了几个点”的结论都不可靠。我们前期就吃过这个亏：改了仓库里的文件，服务却从 site-packages 加载另一份；跑完一次吞吐就拿来比较，后来才发现第一次还混着 Triton 编译和形状预热；Profile 里看到一个 Fill 算子占了十几秒，继续追下去才发现它来自启动和预热，而不是正式五条请求。

最后我们形成了一套固定流程：先核对运行文件，再做同环境热身和前五条筛选，候选通过后才采集只覆盖正式推理窗口的 Profile。 后面的专用 Prefill Attention、Query 复用和 GEMV 优化，都是建立在这套基线上。

1. 项目背景：差距不是一个总分，而是三类负载

项目使用 Qwen3.5-27B，服务框架是 vLLM，硬件是一张 gfx936 DCU。排行榜上的领先队伍大致能做到：

桶	领先成绩区间（tok/s）	我们当时的成绩（tok/s）
4K-8K	约 20.4-22.0	16.19
8K-16K	约 18.8-20.0	14.69
16K-32K	约 15.1-16.9	11.56

官方计分权重不是平均分，而是：

4K-8K   20%
8K-16K  50%
16K-32K 30%

也就是说，中长上下文占了 80% 权重。只优化短请求，哪怕局部数字好看，对最终分数也未必有帮助。反过来，如果为了 TTFT 把 Decode 路径搞慢，TPOT 和输出吞吐又会把收益吃掉。

我们从一开始就把目标拆成三条：

• failed requests 必须为 0；
• 8K-16K 和 16K-32K 不能回退；
• 模型权重、结构、采样、最大输出长度、调度参数和请求覆盖保持不变。

这不是保守，而是为了防止“优化”变成改评测。

2. 机器和软件栈的基础事实

先把现场环境列出来。后面所有数据都只对这套组合负责。

项目	实际环境
DCU	BW gfx936:sramecc+:xnack-
显存	约 68.7 GB
PyTorch	2.10.0
HIP	6.3.26093
DTK	/opt/dtk-26.04-DCC2602-0317
vLLM	0.18.1
模型	Qwen3.5-27B，BF16
并行	单卡，tensor parallel size 1
调度关键参数	max_num_seqs=128、max_num_batched_tokens=4096
服务协议	OpenAI-compatible Chat Completions

PyTorch 在 HIP 平台上仍然大量复用 torch.cuda 命名空间，所以看到 torch.cuda 不等于底层在用 NVIDIA CUDA。判断环境应看 torch.version.hip、设备名和实际 kernel，而不是只看 Python API 的名字。

我们用下面这类最小命令确认 PyTorch 和 TunableOp 状态：

python3 -c "import torch; t=getattr(torch.cuda,'tunable',None); print('torch=',torch.__version__); print('hip=',torch.version.hip); print('device=',torch.cuda.get_device_name()); print('tunable=',t is not None)"

另外，DTK 的 benchmark 工具并不一定在 PATH 中。这台机器上实际位置是：

/opt/dtk-26.04-DCC2602-0317/lib/rocblas/benchmark_tool/rocblas-bench
/opt/dtk-26.04-DCC2602-0317/lib/rocblas/benchmark_tool/rocblas-gemm-tune
/opt/dtk-26.04-DCC2602-0317/lib/hipblaslt/benchmark_tool/hipblaslt-bench

刚找到它们时，我们还遇到过一次很直接的坑：三个 ELF 文件权限都是 644，执行就报 Permission denied。这类问题不应该上升成“库不兼容”，先看 stat 和 file，补上执行权限再说。

3. 第一个现场问题：你改的代码，真的是服务加载的代码吗

我们的仓库在共享目录中，同时镜像里已经安装了一份 vLLM。最开始直接运行：

python3 -c "import vllm, os; print(os.path.dirname(vllm.__file__))"

因为执行目录刚好位于源码树附近，Python 导入到了本地仓库。随后执行复制命令，得到：

cp: '.../vllm/envs.py' and '.../vllm/envs.py' are the same file

这条提示表面看没什么，实际上说明部署验证完全失效了：我们以为在更新系统安装包，实际源和目标是同一个目录。

后来我们不再动态猜测安装路径，而是显式使用：

/usr/local/lib/python3.10/dist-packages/vllm

部署后同时做三类检查：

python3 -m py_compile "$SITE/envs.py" "$SITE/model_executor/layers/utils.py"

cmp -s "$SRC/envs.py" "$SITE/envs.py" && echo OK || echo MISMATCH

cd /tmp && python3 -c "import vllm; print(vllm.__file__)"

为什么要切到 /tmp 再 import？因为这样可以避开当前工作目录对 Python 模块搜索顺序的干扰。这个动作很小，却替我们排除了多次“源码看起来没问题、服务就是不命中”的假故障。

还有一个更隐蔽的情况：比赛容器重建后，共享目录中的 Git 仓库还在，但 /usr/local/lib/python3.10/dist-packages 可能已经恢复成镜像初始状态。我们后来把“启动前检查 cmp”设为硬步骤，而不是出了问题才想起来复制。

4. 32K Prompt 并不会生成一个 q_len=32768 的 Prefill Kernel

这是理解本项目的关键。

服务固定使用：

--max-num-batched-tokens 4096

因此一条 32K 输入不会一次性进入 Attention。它更接近下面这个过程：

chunk 1: q_len ≈ 4096, kv_len ≈ 4096
chunk 2: q_len ≈ 4096, kv_len ≈ 8192
chunk 3: q_len ≈ 4096, kv_len ≈ 12288
...
chunk 8: q_len ≈ 4096, kv_len ≈ 32768
tail:    q_len < 4096, kv_len 已经很长

所以真正值得优化的不是“一个 32K kernel”，而是两类反复出现的形状：

• 完整 chunk：q_len 接近 4096，kv_len 从 4K 一路增长；
• 尾 chunk：q_len 较小，但已经携带很长的 KV。

这也解释了为什么很多全局参数调优不稳定。一个配置可能在 4096 x 32768 上很快，却在首个 8K chunk 或 512 x 32768 尾块上回退。若不按运行时形状分派，局部收益很容易被别的 chunk 抵消。

我们的分派只使用 Python 侧已经可见的 q.shape[0]、max_seqlen_k、dtype 和设备属性，不读取数据集名称，也不查看 prompt 内容，更不会为了选择 kernel 把 GPU 标量同步回 CPU。

5. 基准测试最难的不是跑，而是知道哪一轮能信

吞吐脚本固定使用：

temperature=0
max_concurrency=1
request_rate=1
custom_output_len=1024
disable_shuffle
num_warmups=2

快速筛选始终取每个数据集的前五条，并写入独立目录：

RESULT_ROOT=./candidate_xxx_5 ./run_throughput.sh all 5

我们不会用三条、随机抽样或更短输出去给候选“创造收益”。五条已经是小样本，再缩只会让 P99 和吞吐更容易被偶然波动支配。

5.1 为什么脚本自带 warmup，第一次 all-5 仍然可能偏慢

一次 Q10 Query 复用对照给了很典型的数据：

运行	8K-16K 输出吞吐
新服务第一次前五条	11.708551 tok/s
同一服务再次运行前五条	12.314362 tok/s

两次输出长度和文本完全一致，但吞吐差了约 5%。原因不是模型突然变快，而是两条 warmup 请求不一定覆盖正式五条中所有输入形状、Triton 编译键和缓存状态。

后来我们的做法是：

1. 服务启动并通过 API ready；
2. 跑一轮同分布热身，结果只留作观察；
3. 同一服务、不改环境，再跑固定前五条；
4. 小于 1% 的变化默认视为噪声，除非重复结果足够稳定。

这套规则让我们放弃了不少“首轮看起来不错”的候选，也避免错杀真正需要完整预热的 kernel。

6. 如何采一份不掺启动噪声的 HIP Profile

一开始我们直接用 hipprof --hip-trace --follow-fork 包住整个 start_vllm.sh。结果数据库里混进了模型加载、torch.compile、CUDA Graph 捕获、KV cache 初始化和正式推理，文件很大，热点却不代表比赛请求。

后来改成 retained session：服务照常启动，profile 默认关闭，完成预热后再从另一个终端开启采集，只包住正式的 8K-16K 前五条。

流程可以概括为：

启动 hipprof 服务（trace-off）
    ↓
等待 vLLM API ready
    ↓
运行同分布热身
    ↓
session-client --start
    ↓
运行正式前五条
    ↓
session-client --stop
    ↓
优雅结束服务，等待 CSV 导出

这里有三个现场经验。

第一，输出名必须唯一。重复使用同一个名字会直接报：

File xxx.db already exists, please delete it first.

第二，结束 profile 时不能一上来 kill -KILL。数据库导出需要时间，我们有一次停服务后等了几分钟，看起来像“卡住”，最后正常生成了约 144 MB 的 .db、hipkernel.csv 和 hiptrace.csv。如果强杀，最有价值的结果可能只剩一个损坏数据库。

第三，不要只看 hiptrace.csv 的总时长。真正做算子归类时要以 hipkernel.csv 为主，按 kernel 名称、调用次数和累计 GPU 时间分组。

7. 干净 Profile 之后，热点和我们想的不一样

在保留专用 Prefill Attention、Chunk-aware Pipeline、Q10 Query 复用、长 KV Q20 复用和 Gate/Up Triton GEMV 后，我们采集了一次干净的 8K-16K 前五条 Profile。

GPU kernel 总时间约为 87.579 s，主要构成为：

类别	GPU 时间占比
通用 GEMM	56.57%
Gate/Up GEMV	25.78%
专用 Prefill Attention	7.13%
Gated DeltaNet	3.89%
Decode Attention	2.76%

这个结果改变了后续路线。

项目早期的 Profile 中，Attention 一度接近一半，看起来所有精力都应该继续投到 Attention 上。但 Query 复用落地后，同为 432 次调用的专用 Prefill Attention 已从 14.015 s 降到 7.631 s，后续长 KV Q20 又把它从 7.631 s 降到 6.248 s。热点发生了迁移：线性层合计已经超过 80%。

这时如果还围着 Attention tile size 做微调，就算局部再快 10%，整机收益也只有不到 1%。Profile 的价值不只是告诉我们“哪里慢”，还会告诉我们“这个方向已经优化到该停手了”。

8. 那些看起来像算子问题、实际是运维问题的故障

这次项目的时间并不全花在 kernel 上。下面几类问题重复出现过。

8.1 E-Shell 断开，不等于计算实例真的死了

SCNet Web 终端偶尔会弹出“当前实例已停止服务”，SSH 随后 logout。重新进入后我们却看到：

vllm serve      仍在运行
VLLM::EngineCore 仍在运行
GPU memory use  96%
API             仍可访问

这说明有些情况只是 WebSocket 或 PTY 会话中断。后来服务和长测试统一用 nohup 输出到日志，浏览器断线后可以重新连接查看，不再把模型进程绑在脆弱的前台终端上。

8.2 GPU 显存 96% 在当前配置下是正常状态

模型加载约占 51.2 GiB，KV cache 和图捕获继续占用显存，gpu-memory-utilization=0.95 下看到 94%-96% 并不自动等于泄漏。真正需要警惕的是重复启动第二个 vLLM、残留 hipprof 或 standalone benchmark 与服务同时占卡。

8.3 dmesg 里的 OOM 不一定属于当前容器

宿主机的 dmesg 能看到其他 Docker cgroup 的历史 OOM。必须核对时间和 cgroup ID，不能看到 oom-kill 就认定是当前测试。我们遇到的那次断线发生在 13:30，而日志最后一条 OOM 还停留在凌晨，系统内存也有数百 GiB 可用，显然不是同一件事。

9. 上篇小结：先把问题缩到一个可测的形状

到这个阶段，我们还没有讨论具体 kernel 写法，但已经得到几条决定后续成败的事实：

• 长上下文主要由多个 q_len≈4096 的 Prefill chunk 构成；
• 第一次 all-5 可能仍含形状编译和缓存预热，不能直接当正式对照；
• 正式 Profile 必须排除加载、编译和图捕获；
• Query 复用后，Attention 已不再是唯一主热点，线性层成为下一阶段重点；
• 源码仓库、安装包和运行时环境必须逐文件核对。

中篇会进入真正的算子优化：为什么我们没有继续全局扫 tile，而是做了 Q10/Q20 Query 复用；Gate/Up GEMV 为什么值得手写 Triton；以及 rocBLAS 调优跑了八分多钟，最后为什么一个 solution 都没留下。

---

系列导航

• 上篇：环境、基线、分块负载与干净 Profile（本文）
• 中篇：专用 Prefill Attention、Query 复用与 GEMM 筛选
• 下篇：默认关闭、严格回退、精度门禁与可复现部署

建议标签： DCU、ROCm、HIP、vLLM、Qwen3.5、Triton、大模型推理优化