海光 Z100 DCU (gfx906) 适配 vLLM 运行 Qwen3.5-27B 完整总结

环境信息

上述镜像可以去海光社区：https://developer.sourcefind.cn/modelzoo/list/qwen3.5_vllm/detail?post_id=46176b42-0c0b-11f1-a49f-0242ac150003 获得

Docker 镜像地址：harbor.sourcefind.cn:5443/dcu/admin/base/custom:vllm018-ubuntu22.04-dtk26.04-gemma4-0413

核心问题

该服务器上的容器镜像是为 BW1000 (gfx928) 构建的，而实际 GPU 是 Z100 (gfx906)。存在以下 4 类问题：

1. 库路径冲突：/opt/hyhal/lib 中的 librocm_smi64.so.2 缺少 rsmi_init 符号
2. Triton 未适配 gfx906：MLIR AMDGPU 后端不识别 gfx906
3. vLLM GCN 架构检测失败：amdsmi 返回垃圾值
4. 运行时缺失：_rocm_C 未安装；KV cache profiling 异常

修改一：LD_LIBRARY_PATH 修复（环境配置）

问题

LD_LIBRARY_PATH 中 /opt/hyhal/lib 排在前面，其 librocm_smi64.so.2 不导出 rsmi_init。

错误信息

ImportError: /opt/dtk/lib/librccl.so.1: undefined symbol: rsmi_init

ImportError: /opt/dtk/lib/librccl.so.1: undefined symbol: rsmi_init

解决方案

必须将 DTK 的 rocm_smi 路径放在 LD_LIBRARY_PATH 最前面：

export LD_LIBRARY_PATH=/opt/dtk/.hyhal/rocm_smi/lib:/opt/dtk/lib:/opt/dtk/lib64:/opt/hyhal/lib:/opt/hyhal/lib64:${LD_LIBRARY_PATH}

export LD_LIBRARY_PATH=/opt/dtk/.hyhal/rocm_smi/lib:/opt/dtk/lib:/opt/dtk/lib64:/opt/hyhal/lib:/opt/hyhal/lib64:${LD_LIBRARY_PATH}

路径优先级（从高到低）

1. /opt/dtk/.hyhal/rocm_smi/lib ← 包含 rsmi_init
2. /opt/dtk/lib
3. /opt/dtk/lib64
4. /opt/hyhal/lib ← 缺少 rsmi_init
5. /opt/hyhal/lib64

修改二：Triton 源码修改

前置步骤：备份旧 Triton，克隆源码

# 备份
cp -r /usr/local/lib/python3.10/dist-packages/triton /usr/local/lib/python3.10/dist-packages/triton.bak
cp -r /usr/local/lib/python3.10/dist-packages/triton-3.4.0+git1ef59765.dist-info /usr/local/lib/python3.10/dist-packages/triton-3.4.0+git1ef59765.dist-info.bak

# 克隆
cd /tmp
git clone https://github.com/triton-lang/triton.git --branch v3.4.0 --depth 1 triton-3.4.0-src
cd triton-3.4.0-src

# 备份
cp -r /usr/local/lib/python3.10/dist-packages/triton /usr/local/lib/python3.10/dist-packages/triton.bak
cp -r /usr/local/lib/python3.10/dist-packages/triton-3.4.0+git1ef59765.dist-info /usr/local/lib/python3.10/dist-packages/triton-3.4.0+git1ef59765.dist-info.bak

# 克隆
cd /tmp
git clone https://github.com/triton-lang/triton.git --branch v3.4.0 --depth 1 triton-3.4.0-src
cd triton-3.4.0-src

修改 2.1：TargetUtils.h — 添加 VEGA 到 ISAFamily 枚举

文件：third_party/amd/include/TritonAMDGPUToLLVM/TargetUtils.h

修改前：

enum class ISAFamily {
  Unknown,
  CDNA1,
  CDNA2,
  CDNA3,
  CDNA4,
  RDNA1,
  RDNA2,
  RDNA3,
};

enum class ISAFamily {
  Unknown,
  CDNA1,
  CDNA2,
  CDNA3,
  CDNA4,
  RDNA1,
  RDNA2,
  RDNA3,
};

修改后：

enum class ISAFamily {
  Unknown,
  CDNA1,
  CDNA2,
  CDNA3,
  CDNA4,
  VEGA,
  RDNA1,
  RDNA2,
  RDNA3,
};

enum class ISAFamily {
  Unknown,
  CDNA1,
  CDNA2,
  CDNA3,
  CDNA4,
  VEGA,
  RDNA1,
  RDNA2,
  RDNA3,
};

操作：

sed -i '14a\  VEGA,' third_party/amd/include/TritonAMDGPUToLLVM/TargetUtils.h

sed -i '14a\  VEGA,' third_party/amd/include/TritonAMDGPUToLLVM/TargetUtils.h

修改 2.2：TargetUtils.cpp — 添加 gfx906 架构映射和 VDot 支持

文件：third_party/amd/lib/TritonAMDGPUToLLVM/TargetUtils.cpp

2.2a. deduceISAFamily() — 添加 gfx906 → VEGA 映射

修改前：

  switch (kind) {
  case GK_GFX950: return CDNA4;
  case GK_GFX942: return CDNA3;
  case GK_GFX90A: return CDNA2;
  case GK_GFX908: return CDNA1;
  default: break;
  }

  switch (kind) {
  case GK_GFX950: return CDNA4;
  case GK_GFX942: return CDNA3;
  case GK_GFX90A: return CDNA2;
  case GK_GFX908: return CDNA1;
  default: break;
  }

修改后：

  switch (kind) {
  case GK_GFX950: return CDNA4;
  case GK_GFX906: return VEGA;
  case GK_GFX942: return CDNA3;
  case GK_GFX90A: return CDNA2;
  case GK_GFX908: return CDNA1;
  default: break;
  }

  switch (kind) {
  case GK_GFX950: return CDNA4;
  case GK_GFX906: return VEGA;
  case GK_GFX942: return CDNA3;
  case GK_GFX90A: return CDNA2;
  case GK_GFX908: return CDNA1;
  default: break;
  }

2.2b. supportsVDot() — 添加 VEGA 支持

修改前：

bool supportsVDot() {
  switch (family) {
  case CDNA1: case CDNA2:
  case CDNA3: case CDNA4:
  case RDNA2: case RDNA3:
    return true;
  default: break;
  }
  return false;
}

bool supportsVDot() {
  switch (family) {
  case CDNA1: case CDNA2:
  case CDNA3: case CDNA4:
  case RDNA2: case RDNA3:
    return true;
  default: break;
  }
  return false;
}

修改后：

bool supportsVDot() {
  switch (family) {
  case CDNA1: case CDNA2: case VEGA:
  case CDNA3: case CDNA4:
  case RDNA2: case RDNA3:
    return true;
  default: break;
  }
  return false;
}

bool supportsVDot() {
  switch (family) {
  case CDNA1: case CDNA2: case VEGA:
  case CDNA3: case CDNA4:
  case RDNA2: case RDNA3:
    return true;
  default: break;
  }
  return false;
}

修改 2.3：TargetInfo.cpp — Warp Size 修正

文件：third_party/amd/lib/TritonAMDGPUToLLVM/TargetInfo.cpp

修改前：int TargetInfo::getWarpSize() const { return isCDNA() ? 64 : 32; }

修改后：int TargetInfo::getWarpSize() const { return (isCDNA() || getISAFamily() == ISAFamily::VEGA) ? 64 : 32; }

gfx906 和 CDNA 一样使用 64 线程/warp（wave64 模式）。

修改 2.4：TargetInfo.cpp — permlanex16 指令兼容性修复（关键！）

在 warp reduction 操作中，非 CDNA 路径使用了 permlanex16 指令（CDNA2+ 引入），gfx906 (VEGA) 不支持该指令。

错误信息：

LLVM ERROR: Cannot select: intrinsic %llvm.amdgcn.permlanex16

LLVM ERROR: Cannot select: intrinsic %llvm.amdgcn.permlanex16

修改：让 VEGA 走 DPP 广播路径而非 RDNA 的 permlanex16 路径

if (isCDNA() || getISAFamily() == ISAFamily::VEGA) {
  // DPP 广播路径
} else {
  // RDNA permlanex16 路径
}

if (isCDNA() || getISAFamily() == ISAFamily::VEGA) {
  // DPP 广播路径
} else {
  // RDNA permlanex16 路径
}

故意不修改的地方

• isCDNA()：VEGA 不属于 CDNA，返回 false 正确，避免添加 MFMA 优化
• isRDNA()：VEGA 也不属于 RDNA
• supportsDirectToLdsLoadBitWidth()：走 default 返回 false 是正确保守策略

Triton 编译安装

source /opt/dtk/env.sh
export LD_LIBRARY_PATH=/opt/dtk/.hyhal/rocm_smi/lib:/opt/dtk/lib:/opt/dtk/lib64:/opt/hyhal/lib:/opt/hyhal/lib64:${LD_LIBRARY_PATH}
cd /tmp/triton-3.4.0-src
rm -rf build
pip3 install . --no-build-isolation

source /opt/dtk/env.sh
export LD_LIBRARY_PATH=/opt/dtk/.hyhal/rocm_smi/lib:/opt/dtk/lib:/opt/dtk/lib64:/opt/hyhal/lib:/opt/hyhal/lib64:${LD_LIBRARY_PATH}
cd /tmp/triton-3.4.0-src
rm -rf build
pip3 install . --no-build-isolation

编译时间较长（libtriton.so 约 600MB），约需 5-10 分钟。

修改三：vLLM 源码修改

cp /usr/local/lib/python3.10/dist-packages/vllm/platforms/rocm.py /usr/local/lib/python3.10/dist-packages/vllm/platforms/rocm.py.bak

cp /usr/local/lib/python3.10/dist-packages/vllm/platforms/rocm.py /usr/local/lib/python3.10/dist-packages/vllm/platforms/rocm.py.bak

修改 3.1：rocm.py — 修复 GCN 架构查询

问题：DTK 的 amdsmi 返回垃圾值 gfx18446744073709551615

修复 1：添加正则验证拒绝无效架构名

修复 2：fallback 从 get_device_properties("cuda") 改为 get_device_properties(0)，避免死锁

修改 3.2：rocm.py — 添加 gfx906 到 GFX9 架构识别

修改前：_ON_GFX9 = any(arch in _GCN_ARCH for arch in ["gfx90a", "gfx942", "gfx950"])

修改后：_ON_GFX9 = any(arch in _GCN_ARCH for arch in ["gfx906", "gfx90a", "gfx942", "gfx950"])

gfx906 属于 GFX9 架构族。不加会导致 ViT attention backend 回退等问题。

修改四：禁用 SKINNY_GEMM（环境变量）

问题：预编译 vLLM wheel 中 _rocm_C 自定义扩展未安装

AttributeError: '_OpNamespace' '_rocm_C' object has no attribute 'wvSplitK'

AttributeError: '_OpNamespace' '_rocm_C' object has no attribute 'wvSplitK'

解决：

export VLLM_ROCM_USE_SKINNY_GEMM=0

export VLLM_ROCM_USE_SKINNY_GEMM=0

修改五：手动指定 KV Cache 内存（启动参数）

问题：DTK ROCm 环境下 torch.cuda.memory_stats() 峰值统计异常，计算出负的可用 KV cache

ValueError: No available memory for the cache blocks.

ValueError: No available memory for the cache blocks.

解决：使用 --kv-cache-memory-bytes 手动指定 KV cache 为 1 GiB

--kv-cache-memory-bytes 1073741824

--kv-cache-memory-bytes 1073741824

最终启动命令

source /opt/dtk/env.sh
export LD_LIBRARY_PATH=/opt/dtk/.hyhal/rocm_smi/lib:/opt/dtk/lib:/opt/dtk/lib64:/opt/hyhal/lib:/opt/hyhal/lib64:${LD_LIBRARY_PATH}
export TRITON_HIP_LLD_PATH=/opt/dtk/llvm/bin/ld.lld
export TRITON_CACHE_DIR=/tmp/triton_cache
export VLLM_ROCM_USE_SKINNY_GEMM=0

python3 -m vllm.entrypoints.openai.api_server \
  --model /public/home/acchmzvrrk/models/Qwen3.5-27B \
  --tensor-parallel-size 4 \
  --dtype bfloat16 \
  --max-model-len 1024 \
  --gpu-memory-utilization 0.96 \
  --enforce-eager \
  --max-num-seqs 1 \
  --port 8000 \
  --host 0.0.0.0 \
  --skip-mm-profiling \
  --kv-cache-memory-bytes 1073741824

source /opt/dtk/env.sh
export LD_LIBRARY_PATH=/opt/dtk/.hyhal/rocm_smi/lib:/opt/dtk/lib:/opt/dtk/lib64:/opt/hyhal/lib:/opt/hyhal/lib64:${LD_LIBRARY_PATH}
export TRITON_HIP_LLD_PATH=/opt/dtk/llvm/bin/ld.lld
export TRITON_CACHE_DIR=/tmp/triton_cache
export VLLM_ROCM_USE_SKINNY_GEMM=0

python3 -m vllm.entrypoints.openai.api_server \
  --model /public/home/acchmzvrrk/models/Qwen3.5-27B \
  --tensor-parallel-size 4 \
  --dtype bfloat16 \
  --max-model-len 1024 \
  --gpu-memory-utilization 0.96 \
  --enforce-eager \
  --max-num-seqs 1 \
  --port 8000 \
  --host 0.0.0.0 \
  --skip-mm-profiling \
  --kv-cache-memory-bytes 1073741824

启动参数说明

内存使用分析

关键设计决策详解

为什么新建 VEGA 族而非映射到 CDNA1？

1. 不支持 MFMA 指令：gfx906 使用 FMA/DPP，映射到 CDNA1 会导致 LLVM 生成 MFMA 内部函数
2. Warp Size：gfx906 使用 64 线程/warp，非 CDNA 分支默认返回 32
3. AccelerateAMDMatmul：VEGA 不在 CDNA 中，不会添加 MFMA 优化
4. DPP 支持：gfx906 支持 DPP 广播，要走 CDNA 的 DPP 路径而非 RDNA 的 permlanex16

为什么 gfx906→gfx928 映射失败？

DTK 26.04 的 LLVM 是为 gfx928 构建的。生成的 HSACO 包含 gfx928 指令，在 gfx906 上无法运行。

vLLM 架构检测的坑

amdsmi 返回 gfx18446744073709551615，而 torch.cuda.get_device_properties("cuda") 在 DTK 上会死锁，必须用 get_device_properties(0)。

Qwen3.5-27B 多模态模型处理

使用 --skip-mm-profiling 跳过 vision encoder。纯文本推理完全正常。

_rocm_C 模块缺失

设置 VLLM_ROCM_USE_SKINNY_GEMM=0 回退到标准 PyTorch linear。

首次推理注意事项

• 首次编译需 10-20 分钟，生成 1000+ 个缓存文件
• 后续推理直接使用 /tmp/triton_cache 缓存

完整文件修改清单

测试验证

# 文本补全测试
curl http://localhost:8000/v1/completions \
  -H "Content-Type: application/json" \
  -d '{"model":"...Qwen3.5-27B","prompt":"1+1=","max_tokens":5}'

# 对话测试
curl http://localhost:8000/v1/chat/completions \
  -H "Content-Type: application/json" \
  -d '{"model":"...Qwen3.5-27B","messages":[{"role":"user","content":"你好"}],"max_tokens":50}'

# 文本补全测试
curl http://localhost:8000/v1/completions \
  -H "Content-Type: application/json" \
  -d '{"model":"...Qwen3.5-27B","prompt":"1+1=","max_tokens":5}'

# 对话测试
curl http://localhost:8000/v1/chat/completions \
  -H "Content-Type: application/json" \
  -d '{"model":"...Qwen3.5-27B","messages":[{"role":"user","content":"你好"}],"max_tokens":50}'

故障排查

问题 1：import torch 报 undefined symbol: rsmi_init

原因：LD_LIBRARY_PATH 顺序错误。解决：按修改一调整。

问题 2：import triton 死锁

原因：Triton 为 BW1000 编译。解决：按修改二从源码重编。

问题 3：GCN arch gfx18446744073709551615

原因：amdsmi 返回垃圾值。解决：按修改 3.1 加验证。

问题 4：LLVM ERROR: permlanex16

原因：Triton 生成不支持指令。解决：按修改 2.4 走 DPP 路径。

问题 5：HIP error: invalid device function

解决：完成修改 2.4 + 使用 --skip-mm-profiling。

问题 6：AttributeError: _rocm_C has no wvSplitK

解决：设置 VLLM_ROCM_USE_SKINNY_GEMM=0。

问题 7：ValueError: No available memory for cache

解决：使用 --kv-cache-memory-bytes 手动指定。

问题 8：Free memory less than desired

解决：pkill -9 -f vllm 清理残留进程。

备份文件清单