> **海光 DCU 实战系列** > > 1. [从零跑通海光 DCU：完整配置与排错](https://blog.csdn.net/2403_86879992/article/details/162929518) > 2. [PyTorch + HIP 自定义融合算子](https://blog.csdn.net/2403_86879992/article/details/162934908) > 3. [用 hipprof 定位 PyTorch 性能瓶颈](https://blog.csdn.net/2403_86879992/article/details/162937545) > 本文面向第一次接触 DCU 的开发者。目标不是让终端里出现一次`torch.cuda.is_available() == True` 就结束，而是从设备、驱动、DTK、PyTorch一直验证到自定义 HIP 算子和性能分析。运行时、PyTorch、HIP 编译与 hipprof命令在一台 BW / `gfx936` DCU 容器中逐项实测；裸机安装和 Docker 宿主机部分以 DTK 25.04.2 对应官方文档为依据，没有在共享集群宿主机上执行安装操作。 第一次拿到 DCU 服务器时，我也犯过一个很典型的错误：把“能看到卡”“能导入PyTorch”“模型能跑”当成同一件事。实际上，它们属于不同层。某一层有问题时，盲目重装 Python 包通常只会把环境弄得更乱。 一套真正可用的 DCU 开发环境，至少应通过下面五层验收： | 层次 | 要回答的问题 | 最小验证 | | --- | --- | --- | | 硬件与驱动 | 操作系统是否识别 DCU | `hy-smi`、驱动模块 | | DTK Runtime | 运行时能否枚举计算设备 | `rocminfo` | | 编译工具链 | 能否把 HIP 源码编译并执行 | `hipcc` + vector add | | 深度学习框架 | PyTorch 是否连接到 HIP Runtime | BF16 GEMM 冒烟测试 | | 调试与分析 | 出错或变慢后能否定位 | `hipprof`、同步计时 | 后文会从零把这五层走完，并穿插我真正踩过的坑。  ## 一、先判断你是哪一种环境 开始安装前，先确认自己属于下面哪一类，同时弄清楚容器镜像从哪里来。宿主机、平台基础镜像和社区应用镜像是三层不同的东西，操作边界也完全不同。 ### 1. 共享集群或算力平台 平台通常已经在宿主机安装了驱动。以 SCNet 的创建页面为例，容器镜像又分成两条常见路径： | 镜像类型 | 通常已经包含什么 | 更适合谁 | | --- | --- | --- | | 平台基础/框架镜像 | Ubuntu、DTK、Python，以及选定版本的 PyTorch、TensorFlow、JAX 等 | 想获得干净、可控开发环境的人 | | “我的镜像”或社区复制镜像 | 在框架镜像上继续加入 OneScience、Qwen、vLLM、Transformers、DeepSpeed、启动脚本等 | 想快速跑通特定模型或应用的人 | 这里的“基础镜像”并不是只有一个空 Ubuntu。页面中类似`pytorch-2.5.1-ubuntu22.04-dtk25.04.2-py3.11` 的名字，已经明确包含操作系统、DTK、Python 和 PyTorch。社区里的 `onescience-0.3.0`、Qwen 方案镜像通常是在这类框架镜像之上继续叠加应用软件。  但是，“预装了很多组件”不等于“存在多个相互隔离的 Python 环境”。它可能只是在同一个系统 Python 里安装了很多包，也可能额外创建 Conda/venv。必须通过`which python`、`VIRTUAL_ENV`、`CONDA_PREFIX` 和包版本实际判断，不能只看镜像名称猜测。 普通用户不应在登录节点跑计算，也不应在已经分配好的平台容器里重新安装内核驱动或尝试再启动一层 Docker。 正确流程是： 1. 申请或调度到真正挂载 DCU 的计算节点； 2. 根据任务选择平台框架镜像或社区应用镜像； 3. 进入运行容器，检查 DTK 是否已经加载，必要时再 `source env.sh`； 4. 盘点 Python 和预装包，不要第一时间执行 `pip install -U`； 5. 从 `hy-smi`、`rocminfo`、PyTorch 和目标应用开始逐层验证。 “登录节点能 SSH”只证明你进入了集群，不证明当前 shell 能使用 DCU。很多 `torch.cuda.is_available() == False`，其实只是命令跑错了节点。 ### 2. 已装驱动的宿主机，准备使用 Docker 这种情况下，宿主机负责驱动，容器负责用户态 DTK 和框架。推荐使用 DTK Container Toolkit 把设备和驱动文件映射进容器，不要在容器里执行驱动安装器。 ### 3. 自己维护的裸机 只有这种情况才需要 root 权限安装 DCU 驱动、重启机器、安装 DTK。安装前必须核对操作系统、内核、驱动和 DTK 的兼容矩阵。生产服务器还应准备带外管理或救援入口，避免修改内核启动项后失联。 如果你不能确定自己是哪一种，先问平台管理员。不要从“安装驱动”开始试错。 ## 二、先采集环境指纹，不要急着改系统 建议把下面信息保存到一个文本文件。以后遇到问题，无论问管理员还是发社区帖，这些信息都比一句“DCU 跑不起来”有用得多。 ```bash hostname date uname -a cat /etc/os-release command -v hy-smi || true command -v rocminfo || true command -v hipcc || true command -v hipprof || true ``` 如果 DTK 已经在 `/opt`，先不要永久修改 `.bashrc`，临时加载： ```bash set +u if [[ -f /opt/dtk-25.04.2/env.sh ]]; then source /opt/dtk-25.04.2/env.sh else source /opt/dtk/env.sh fi set -u ``` 然后继续记录： ```bash hy-smi --showproductname rocminfo | grep -E 'Name:.*gfx|Marketing Name|Compute Unit' | head -n 30 hipcc --version ``` 容器还应增加一组“镜像继承与 Python 环境”指纹： ```bash PYTHON_BIN="$(command -v python || command -v python3)" echo "PYTHON_BIN=$PYTHON_BIN" command -v conda || true echo "VIRTUAL_ENV=${VIRTUAL_ENV-}" echo "CONDA_PREFIX=${CONDA_PREFIX-}" "$PYTHON_BIN" -m pip --version "$PYTHON_BIN" - <<'PY' import importlib.metadata as m import sys print("python executable:", sys.executable) for name in ( "torch", "onescience", "transformers", "vllm", "deepspeed", "triton", "onnxruntime" ): try: print(name, m.version(name)) except m.PackageNotFoundError: print(name, "not installed") PY "$PYTHON_BIN" -m pip check ``` 当前 `onescience-0.3.0` 社区镜像的实测结果很能说明问题：它使用`/usr/bin/python`，没有激活 `VIRTUAL_ENV` 或 `CONDA_PREFIX`，但同一个 Python里已经安装了： ```text torch 2.5.1+das.opt1.dtk25042 onescience 0.3.0 transformers 4.51.1 vllm 0.8.5.post1+das.opt1.dtk25041 deepspeed 0.14.2+das.opt1.dtk25041 triton 3.0.0+das.opt1.dtk25041 onnxruntime 1.19.2+das.opt1.dtk25042 ``` 这说明社区镜像确实能省掉大量安装工作，但 `pip check` 同时发现 vLLM 和若干融合算子的元数据仍声明依赖 Torch 2.4.1，而镜像实际使用 Torch 2.5.1。依赖声明不一致不等于应用一定不能运行——定制包可能经过额外适配——但它意味着必须用目标模型做最小冒烟测试（smoke test），不能因为软件“都装好了”就默认任意组合都兼容。 这里有三个容易误判的点。 第一，`rocminfo` 会同时列出 CPU agent 和 DCU agent。看到很多 CPU 核心不代表没有卡，应继续找 `gfx...` 架构项。 第二，安装目录名不一定等于编译器内部版本。我实测环境的路径是`/opt/dtk-25.04.2`，但 `hipcc --version` 输出的是 `dcc 25.08.0-0 / clang 17.0.0`。排查兼容性时应同时记录路径、编译器、PyTorch 和 HIP Runtime，不能只报一个“DTK 25.04.2”。 第三，容器内的 `uname -r` 是宿主机内核。它不必和容器的 Ubuntu 小版本看起来完全配套，也不能直接拿来套用裸机安装章节。 我把完整指纹封装成了一个不会修改系统的脚本：`dcu_env_check.sh`。它不会因为某条命令缺失就提前退出，适合新手直接运行。 ```bash chmod +x dcu_env_check.sh ./dcu_env_check.sh | tee dcu_env_report.txt ``` 完整脚本如下，复制保存为 `dcu_env_check.sh` 即可： ```bash #!/usr/bin/env bash # A non-destructive DCU environment fingerprint. It keeps checking after a # missing command so that one broken layer does not hide the remaining clues. set -o pipefail section() { printf '\n===== %s =====\n' "$1" } section "Operating system" hostname date uname -a if [[ -r /etc/os-release ]]; then sed -n '1,8p' /etc/os-release fi section "Load the DTK environment" # Some vendor env.sh versions reference variables that may not be defined. # Temporarily disable nounset while sourcing them. set +u if [[ -f /opt/dtk-25.04.2/env.sh ]]; then source /opt/dtk-25.04.2/env.sh echo "sourced /opt/dtk-25.04.2/env.sh" elif [[ -f /opt/dtk/env.sh ]]; then source /opt/dtk/env.sh echo "sourced /opt/dtk/env.sh" else echo "[missing] DTK env.sh under /opt" fi set -u section "Driver and device" lsmod | grep -E 'hycu|hydcu|hyttm|hykcl' || true if command -v hy-smi >/dev/null 2>&1; then hy-smi --showproductname 2>/dev/null || hy-smi | head -n 30 else echo "[missing] hy-smi" fi if command -v rocminfo >/dev/null 2>&1; then rocminfo | grep -E 'Name:.*(gfx|ZIFANG)|Marketing Name|Compute Unit' | head -n 24 else echo "[missing] rocminfo" fi section "Compiler and profiler" for tool in hipcc dcc hipprof; do if command -v "$tool" >/dev/null 2>&1; then printf '%-8s %s\n' "$tool" "$(command -v "$tool")" else printf '%-8s [missing]\n' "$tool" fi done if command -v hipcc >/dev/null 2>&1; then hipcc --version | head -n 8 fi PYTHON_BIN="" if command -v python >/dev/null 2>&1; then PYTHON_BIN="$(command -v python)" elif command -v python3 >/dev/null 2>&1; then PYTHON_BIN="$(command -v python3)" fi section "Python environment ownership" printf 'python %s\n' "${PYTHON_BIN:-[missing]}" if [[ -n "$PYTHON_BIN" ]]; then printf 'python -m pip ' "$PYTHON_BIN" -m pip --version 2>/dev/null || echo '[missing]' else echo 'python -m pip [missing]' fi printf 'conda %s\n' "$(command -v conda 2>/dev/null || echo '[not active]')" printf 'VIRTUAL_ENV %s\n' "${VIRTUAL_ENV-}" printf 'CONDA_PREFIX %s\n' "${CONDA_PREFIX-}" printf 'PYTHONPATH %s\n' "${PYTHONPATH-}" section "Python packages and PyTorch" if [[ -n "$PYTHON_BIN" ]]; then "$PYTHON_BIN" --version "$PYTHON_BIN" - <<'PY' import importlib.metadata as metadata import sys print("executable:", sys.executable) for package in ( "torch", "onescience", "transformers", "vllm", "accelerate", "deepspeed", "triton", "onnx", "onnxruntime", ): try: print(f"{package}:", metadata.version(package)) except metadata.PackageNotFoundError: print(f"{package}: not installed") try: import torch except Exception as exc: print("torch import failed:", repr(exc)) raise SystemExit(0) print("torch runtime:", torch.__version__) print("torch.version.hip:", torch.version.hip) print("torch.cuda.is_available():", torch.cuda.is_available()) print("device count:", torch.cuda.device_count()) if torch.cuda.is_available(): prop = torch.cuda.get_device_properties(0) print("device name:", prop.name) print("architecture:", getattr(prop, "gcnArchName", "unknown")) print("memory GiB:", round(prop.total_memory / 2**30, 3)) PY if [[ "${DCU_ENV_FULL_PIP_CHECK:-0}" == "1" ]]; then section "pip dependency check" "$PYTHON_BIN" -m pip check || true else echo "tip: run DCU_ENV_FULL_PIP_CHECK=1 ./dcu_env_check.sh to execute pip check" fi else echo "[missing] python" fi ```  如果还要检查社区镜像中的依赖约束，可以直接运行： ```bash DCU_ENV_FULL_PIP_CHECK=1 ./dcu_env_check.sh ``` ## 三、裸机安装：Ubuntu 22.04 的可靠顺序 如果平台已经提供容器，可以跳到下一节。下面只适用于自己维护的物理机。 ### 1. 先核对系统和启动项 DTK 的安装部署手册会为不同发行版列出明确的内核和依赖要求。以 25.04.2 文档中的 Ubuntu 22.04 章节为例，它给出了对应的内核要求，并特别提醒内核启动项中不要包含 `nomodeset`，否则驱动可能无法加载。 先看当前状态： ```bash cat /etc/os-release uname -r cat /proc/cmdline ``` 如果存在 `nomodeset`，不要在远程生产机上贸然改 GRUB。确认有控制台或带外管理后，再按当前发行版的方式修改，并保留可回退的旧启动项。 ### 2. 安装依赖 官方 Ubuntu 22.04 章节列出的核心依赖可以整理为： ```bash sudo apt-get update sudo apt-get install -y \ make gcc g++ gfortran cmake git wget curl autoconf \ python3 python3-pip python3-dev python3-wheel \ libdrm-dev libnuma-dev libelf-dev libpci-dev pciutils \ libpciaccess-dev pkg-config kmod rsync sqlite3 libsqlite3-dev ``` 不同系统和安装包可能还需要额外依赖，应以与你下载版本匹配的官方安装手册为准。不要为了“命令短”删除编译器、内核头文件或 `libdrm` 相关依赖。 ### 3. 安装驱动并重启 驱动包请登录光合开发者社区，从工具下载页选择与系统匹配的当前版本。页面路径和文件名会随版本更新，下面只展示安装顺序： ```bash DRIVER="rock-.run" # Replace with the exact downloaded file name. chmod 755 "$DRIVER" sudo "./$DRIVER" sudo reboot ``` 重启后先看模块，不要立刻装 Python： ```bash lsmod | grep -E 'hycu|hydcu|hyttm|hykcl' ``` 驱动模块名称可能随版本变化。我当前环境使用 `hycu`，一些文档和历史环境出现`hydcu`，所以诊断脚本同时匹配两者。 如果首次识别失败，检查与当前内核完全一致的 headers/modules 是否存在： ```bash sudo apt install "linux-headers-$(uname -r)" sudo apt install "linux-modules-extra-$(uname -r)" sudo apt install libdrm-dev ``` 然后按官方流程卸载并重装驱动。不要混装多个时期的驱动包。 ### 4. 安装 DTK 以 tar 包安装到 `/opt` 为例： ```bash sudo tar xvf DTK-25.04.2-ubuntu22.04-x86_64.tar.gz -C /opt ls -ld /opt/DTK-25.04.2 /opt/dtk 2>/dev/null || true sudo ln -s /opt/DTK-25.04.2 /opt/dtk ``` 如果 `/opt/dtk` 已经存在，不要直接强制覆盖。先用 `ls -ld /opt/dtk` 确认它是旧版本软链接还是实际目录，再由管理员决定是否切换链接。 第一次只在当前 shell 加载： ```bash set +u source /opt/dtk/env.sh set -u ``` 确认工具链正常后，再决定写入 `.bashrc` 或使用 environment modules。对多人共用服务器，我更推荐 module：项目明确选择版本，不让某个全局 `LD_LIBRARY_PATH`影响所有任务。 ### 5. 安装与 DTK 匹配的 PyTorch 安装 DTK 并不会自动得到可用的 PyTorch。裸机用户应从[光合开发者社区工具页](https://developer.sourcefind.cn/)，或对应版本的框架发布说明中，选择同时匹配下面三项的官方 DCU PyTorch 包： - DTK/HIP Runtime 版本； - Python 主版本和 ABI； - 操作系统与处理器架构。 不要直接运行普通的 `pip install torch`。PyPI 默认包可能是 CPU 或其他后端构建，即使能够 `import torch`，也不代表它能连接 DCU。更稳妥的方式是先创建独立环境，再安装已经下载并核对过的官方 wheel： ```bash python3 -m venv ~/venvs/dtk25042 source ~/venvs/dtk25042/bin/activate python -m pip install --upgrade pip # Replace this placeholder with the exact official wheel for the current # DTK, Python and operating-system combination. TORCH_WHEEL=/path/to/official-dcu-torch.whl python -m pip install "$TORCH_WHEEL" ``` 如果官方提供的是成套框架安装包或容器镜像，应优先按照该版本说明整体安装，不要自行混搭来自不同 DTK 版本的 `torch`、`torchvision`、Triton 和融合算子 wheel。安装完成后必须重新运行本文的环境指纹和 BF16 冒烟测试，并确认： ```bash python - <<'PY' import torch print(torch.__version__) print(torch.version.hip) print(torch.cuda.is_available()) PY ``` 只有 `torch.version.hip` 非空、设备可用且实际计算通过，才能说明框架安装完成。 ## 四、容器方式：先分清谁在管理容器 ### 1. 自己维护 Docker 宿主机 DTK Container Toolkit 的职责是把宿主机上的 DCU 设备和驱动文件正确映射给容器。它应安装在 Docker 宿主机： ```bash sudo dpkg -i dtk-container-toolkit_*_amd64.deb sudo systemctl restart docker ``` 给容器全部设备： ```bash IMAGE="your-dtk-image:tag" # Replace with the image selected for your DTK version. docker run --rm -it --gpus all "$IMAGE" bash ``` 选择部分设备时，注意 shell 引号： ```bash IMAGE="your-dtk-image:tag" docker run --rm -it --gpus '"device=0,1"' "$IMAGE" bash ``` 也可以使用： ```bash IMAGE="your-dtk-image:tag" docker run --rm -it \ -e DTK_VISIBLE_DEVICES=0,1 \ "$IMAGE" bash ``` 两个变量不要混淆： - `DTK_VISIBLE_DEVICES` 主要控制哪些设备被注入容器； - `HIP_VISIBLE_DEVICES` 控制当前 HIP 应用在已经可见的设备中使用哪些卡。 ### 2. SCNet 等平台已经替你创建好容器 这时上面的 `docker run` 命令是平台管理员或底层服务做的，普通用户不需要、通常也不应该在已分配容器里再执行。你需要选择的是镜像层次： - 写 HIP、做算子验证、希望环境尽量干净：优先平台 PyTorch 框架镜像； - 快速部署 Qwen/vLLM/OneScience：优先对应社区方案镜像； - 已经验证并想固化依赖：把自己的稳定容器保存为“我的镜像”。 社区方案镜像的优势是开箱速度快，代价是包更多、历史依赖更多、版本关系更难一眼看清。进入容器后的正确动作不是立刻升级 Torch、NumPy 或 Transformers，而是： 1. 保存镜像名称和页面显示的底层框架镜像； 2. 运行环境指纹脚本； 3. 执行 `python -m pip check`，记录而不是立刻修复所有告警； 4. 用目标模型跑最小加载、单步推理或官方冒烟测试； 5. 只有在确认具体冲突后，才复制一份镜像做变更。 为什么不建议直接“修到 pip check 全绿”？因为定制 DCU wheel 包可能带有与上游不同的版本后缀和补丁，盲目升级一个包，可能把已经适配好的 vLLM、Triton 或融合算子换成普通上游版本。先验证业务，再决定是否需要隔离新 venv 或重建镜像。 无论选哪一种镜像，进容器后仍然要重新跑 `hy-smi`、`rocminfo`、PyTorch 冒烟测试和目标应用冒烟测试。宿主机能看到卡、镜像名字里写着 Qwen，都不代表当前模型、量化格式和依赖组合已经验证成功。 ## 五、最隐蔽的坑：`env.sh` 遇到 `set -u` 很多训练脚本开头会写： ```bash set -euo pipefail source /opt/dtk-25.04.2/env.sh ``` 这看起来很严谨，却可能让脚本在真正业务代码之前直接退出。我在当前环境中实测： ```bash bash -uc 'source /opt/dtk-25.04.2/env.sh' ``` 得到： ```text /opt/dtk-25.04.2/env.sh: line 26: LD_LIBRARY_PATH: unbound variable exit code: 127 ``` 原因是环境脚本读取了尚未定义的 `LD_LIBRARY_PATH`，而 nounset 把它当成致命错误。安全写法是： ```bash set -eo pipefail set +u source /opt/dtk-25.04.2/env.sh set -u # 从这里开始再运行自己的程序 python your_program.py ``` 这类问题非常像“任务刚启动就没日志”。如果外层又用了 `nohup`，很容易误以为Python、数据集或调度器出错。排查长任务时至少确认三件事：进程还活着、日志在增长、输出目录在变化。 ## 六、验证 PyTorch：为什么接口仍然叫 `torch.cuda` 在 HIP 版 PyTorch 中，设备 API 仍复用 `torch.cuda` 命名空间。这是兼容层设计，并不表示程序偷偷使用了 NVIDIA CUDA。判断依据应是 `torch.version.hip`。 先运行下面的最小检查： ```python import torch print("torch:", torch.__version__) print("HIP runtime:", torch.version.hip) print("available:", torch.cuda.is_available()) print("device count:", torch.cuda.device_count()) if torch.cuda.is_available(): prop = torch.cuda.get_device_properties(0) print("name:", prop.name) print("arch:", getattr(prop, "gcnArchName", "unknown")) print("memory GiB:", prop.total_memory / 2**30) ``` 本次实测得到： ```text torch 2.5.1 HIP runtime 6.3.25405 device BW architecture gfx936:sramecc+:xnack- visible memory 63.984 GiB ``` 随后运行一个 BF16 GEMM，验证计算、数值和同步计时： ```bash python torch_smoke.py ``` `torch_smoke.py` 的完整内容如下： ```python """Small PyTorch/DCU correctness and synchronized timing smoke test.""" from __future__ import annotations import os import time import torch def main() -> None: print("torch:", torch.__version__) print("HIP runtime:", torch.version.hip) print("accelerator available:", torch.cuda.is_available()) if not torch.cuda.is_available() or torch.version.hip is None: raise SystemExit("PyTorch is not connected to a HIP/DCU runtime") prop = torch.cuda.get_device_properties(0) print("device:", prop.name) print("architecture:", getattr(prop, "gcnArchName", "unknown")) # A finite-value check alone can miss a numerically wrong GEMM. Compare a # small BF16 device result against a CPU FP32 matmul of the same quantized # inputs before running the larger timing smoke test. check_n = int(os.environ.get("DCU_SMOKE_CHECK_N", "64")) torch.manual_seed(0) check_a_cpu = torch.randn((check_n, check_n), dtype=torch.float32) check_b_cpu = torch.randn((check_n, check_n), dtype=torch.float32) check_a_bf16 = check_a_cpu.to(torch.bfloat16) check_b_bf16 = check_b_cpu.to(torch.bfloat16) expected = check_a_bf16.float() @ check_b_bf16.float() check_a = check_a_bf16.to("cuda") check_b = check_b_bf16.to("cuda") actual = (check_a @ check_b).float().cpu() torch.cuda.synchronize() max_abs = float((actual - expected).abs().max()) torch.testing.assert_close(actual, expected, rtol=3e-2, atol=2e-1) print(f"BF16 correctness check: n={check_n} max_abs={max_abs:.6f}") del check_a, check_b, actual torch.cuda.empty_cache() n = int(os.environ.get("DCU_SMOKE_N", "2048")) warmup = int(os.environ.get("DCU_SMOKE_WARMUP", "10")) repeats = int(os.environ.get("DCU_SMOKE_REPEATS", "50")) torch.manual_seed(0) a = torch.randn((n, n), device="cuda", dtype=torch.bfloat16) b = torch.randn((n, n), device="cuda", dtype=torch.bfloat16) for _ in range(warmup): c = a @ b torch.cuda.synchronize() start = time.perf_counter() for _ in range(repeats): c = a @ b torch.cuda.synchronize() mean_ms = (time.perf_counter() - start) * 1000.0 / repeats if not bool(torch.isfinite(c).all()): raise RuntimeError("The BF16 GEMM produced NaN or Inf") print("output:", tuple(c.shape), c.dtype) print(f"synchronized mean: {mean_ms:.4f} ms") print(f"allocated: {torch.cuda.memory_allocated() / 2**20:.3f} MiB") print(f"reserved: {torch.cuda.memory_reserved() / 2**20:.3f} MiB") if __name__ == "__main__": main() ``` 脚本先用 `64 x 64` 的 BF16 输入和 CPU FP32 参考结果做近似一致性检查，再运行大矩阵计时。在本次环境中，`2048 x 2048` BF16 GEMM 的同步均值为 `0.0618 ms`，输出全部有限，`allocated/reserved` 分别为 `24/62 MiB`。这个绝对速度只用于本机冒烟测试，不应拿来和不同型号、频率、功耗状态的机器横向排名。 计时时一定要同步： ```python torch.cuda.synchronize() start = time.perf_counter() for _ in range(repeats): output = model(input) torch.cuda.synchronize() elapsed = time.perf_counter() - start ``` HIP kernel 默认异步提交。如果没有末尾同步，测到的可能只是 Python 发射时间。 ## 七、绕开框架，编译第一个 HIP 程序 PyTorch 能运行还不够。自定义算子、框架迁移和底层排错都依赖编译工具链。最小的 vector add 可以同时验证 `hipcc`、Runtime、显存分配、主机到设备/设备到主机传输（H2D/D2H）、kernel launch 和 HIP event。 把下面内容保存为 `vector_add.hip`： ```cpp #include #include #include #include #include #include #define HIP_CHECK(call) \ do { \ const hipError_t status = (call); \ if (status != hipSuccess) { \ std::cerr << "HIP error at " << __FILE__ << ':' << __LINE__ << ": " \ << hipGetErrorString(status) << std::endl; \ std::exit(EXIT_FAILURE); \ } \ } while (0) __global__ void vector_add(const float* a, const float* b, float* c, int n) { const int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) { c[i] = a[i] + b[i]; } return; } int main() { int device = 0; HIP_CHECK(hipSetDevice(device)); hipDeviceProp_t prop{}; HIP_CHECK(hipGetDeviceProperties(&prop, device)); std::cout << "device=" << prop.name << " arch=" << prop.gcnArchName << '\n'; constexpr int n = 1 << 20; constexpr int threads = 256; constexpr int warmup = 20; constexpr int repeats = 100; const size_t bytes = static_cast(n) * sizeof(float); std::vector a(n), b(n), c(n); for (int i = 0; i < n; ++i) { a[i] = static_cast(i % 101) / 101.0f; b[i] = static_cast(i % 79) / 79.0f; } float *d_a = nullptr, *d_b = nullptr, *d_c = nullptr; HIP_CHECK(hipMalloc(&d_a, bytes)); HIP_CHECK(hipMalloc(&d_b, bytes)); HIP_CHECK(hipMalloc(&d_c, bytes)); HIP_CHECK(hipMemcpy(d_a, a.data(), bytes, hipMemcpyHostToDevice)); HIP_CHECK(hipMemcpy(d_b, b.data(), bytes, hipMemcpyHostToDevice)); const dim3 block(threads); const dim3 grid((n + threads - 1) / threads); for (int i = 0; i < warmup; ++i) { hipLaunchKernelGGL(vector_add, grid, block, 0, 0, d_a, d_b, d_c, n); } HIP_CHECK(hipGetLastError()); HIP_CHECK(hipDeviceSynchronize()); hipEvent_t start{}, stop{}; HIP_CHECK(hipEventCreate(&start)); HIP_CHECK(hipEventCreate(&stop)); HIP_CHECK(hipEventRecord(start, 0)); for (int i = 0; i < repeats; ++i) { hipLaunchKernelGGL(vector_add, grid, block, 0, 0, d_a, d_b, d_c, n); } HIP_CHECK(hipGetLastError()); HIP_CHECK(hipEventRecord(stop, 0)); HIP_CHECK(hipEventSynchronize(stop)); float elapsed_ms = 0.0f; HIP_CHECK(hipEventElapsedTime(&elapsed_ms, start, stop)); HIP_CHECK(hipMemcpy(c.data(), d_c, bytes, hipMemcpyDeviceToHost)); float max_abs_error = 0.0f; for (int i = 0; i < n; ++i) { max_abs_error = std::max(max_abs_error, std::abs(c[i] - (a[i] + b[i]))); } std::cout << "elements=" << n << " mean_kernel_ms=" << elapsed_ms / repeats << " max_abs_error=" << max_abs_error << '\n'; HIP_CHECK(hipEventDestroy(start)); HIP_CHECK(hipEventDestroy(stop)); HIP_CHECK(hipFree(d_a)); HIP_CHECK(hipFree(d_b)); HIP_CHECK(hipFree(d_c)); return max_abs_error == 0.0f ? EXIT_SUCCESS : EXIT_FAILURE; } ``` 编译前先从 `rocminfo` 确认目标架构： ```bash rocminfo | grep -E 'Name:.*gfx' ``` 本机是 `gfx936`，因此编译： ```bash hipcc -O3 -std=c++17 \ --offload-arch=gfx936 \ vector_add.hip -o vector_add ./vector_add ``` 实测输出： ```text device=BW arch=gfx936:sramecc+:xnack- elements=1048576 mean_kernel_ms=0.01875 max_abs_error=0 ```  这里又有一个真实的小坑。当前环境变量默认包含多个目标架构，不显式指定时，编译器为 `gfx906/gfx926/gfx928/gfx936` 都生成了代码，本次耗时 `4.132 s`；确认实机后指定 `gfx936`，编译耗时降到 `2.805 s`。绝对时间会随节点负载波动，这里更重要的是减少了不需要的目标代码生成。 不要照抄别人机器的 `--offload-arch`。指定错误的架构，轻则编译失败，重则在运行时出现 illegal instruction。 ## 八、从 HIP 程序走到 PyTorch 自定义扩展 真实项目往往不是独立 C++ 程序，而是希望在 PyTorch 图中替换某个热点算子。这里采用 source-only 方式：正文只提供源代码，由目标机器现场编译，不依赖提前生成的 `.so` 二进制。 先确认编译依赖： ```bash command -v c++ command -v hipcc command -v ninja || true python -c "import torch; print(torch.__version__, torch.version.hip)" ``` 当前 PyTorch 扩展流程会优先使用 Ninja。若镜像缺少它，应从平台软件源或当前Python 环境安装 Ninja，而不是为了补一个构建工具去整体升级 Torch。 示例 `torch_extension/` 包含三个文件： ```text add_ext.cpp pybind 接口 add_ext_kernel.cu HIP kernel 与张量检查 build_and_test.py 编译、缓存和数值验证 ``` `add_ext.cpp`： ```cpp #include torch::Tensor add_one_dcu(torch::Tensor input); PYBIND11_MODULE(TORCH_EXTENSION_NAME, module) { module.def("add_one", &add_one_dcu, "Add one with a HIP kernel"); } ``` `add_ext_kernel.cu`： ```cpp #include #include #include __global__ void add_one_kernel(const float* input, float* output, int64_t n) { const int64_t i = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; if (i < n) { output[i] = input[i] + 1.0f; } return; } torch::Tensor add_one_dcu(torch::Tensor input) { TORCH_CHECK(input.is_cuda(), "input must be on a HIP/DCU device"); TORCH_CHECK(input.is_contiguous(), "input must be contiguous"); TORCH_CHECK(input.scalar_type() == at::kFloat, "input must be float32"); auto output = torch::empty_like(input); const int64_t n = input.numel(); if (n == 0) { return output; } constexpr int threads = 256; const int blocks = static_cast((n + threads - 1) / threads); hipStream_t stream = at::cuda::getCurrentCUDAStream().stream(); hipLaunchKernelGGL(add_one_kernel, dim3(blocks), dim3(threads), 0, stream, input.data_ptr(), output.data_ptr(), n); const hipError_t status = hipGetLastError(); TORCH_CHECK(status == hipSuccess, "HIP launch failed: ", hipGetErrorString(status)); return output; } ``` `build_and_test.py`： ```python """Build a source-only PyTorch HIP extension and verify its output.""" from __future__ import annotations import hashlib import os from pathlib import Path import time import torch from torch.utils.cpp_extension import load def main() -> None: if not torch.cuda.is_available() or torch.version.hip is None: raise SystemExit("A HIP-enabled PyTorch runtime is required") here = Path(__file__).resolve().parent source_paths = [here / "add_ext.cpp", here / "add_ext_kernel.cu"] source_payloads = [(path.name, path.read_bytes()) for path in source_paths] prop = torch.cuda.get_device_properties(0) arch_name = getattr(prop, "gcnArchName", "") if not arch_name: raise SystemExit("The PyTorch runtime did not report a gfx architecture") arch = arch_name.split(":", 1)[0] digest = hashlib.sha256() for name, payload in source_payloads: digest.update(name.encode()) digest.update(payload) digest.update(torch.__version__.encode()) digest.update(str(torch.version.hip).encode()) digest.update(arch.encode()) digest.update(b"-O3") extension_name = f"dcu_add_one_{digest.hexdigest()[:12]}" os.environ.setdefault("PYTORCH_ROCM_ARCH", arch) os.environ.setdefault( "TORCH_EXTENSIONS_DIR", f"/tmp/torch_extensions_{os.getuid()}" ) # PyTorch HIPify can create a generated .hip file next to a .cu source. # Stage immutable source copies in /tmp so the project tree stays clean. stage_dir = Path( os.environ.get( "DCU_EXTENSION_STAGE_ROOT", f"/tmp/dcu_extension_sources_{os.getuid()}" ) ) / extension_name stage_dir.mkdir(parents=True, exist_ok=True) sources = [] for name, payload in source_payloads: staged_source = stage_dir / name if not staged_source.exists() or staged_source.read_bytes() != payload: staged_source.write_bytes(payload) sources.append(staged_source) build_start = time.perf_counter() extension = load( name=extension_name, sources=[str(path) for path in sources], extra_cflags=["-O3"], extra_cuda_cflags=["-O3"], verbose=True, ) build_seconds = time.perf_counter() - build_start x = torch.arange(1 << 20, device="cuda", dtype=torch.float32) expected = x + 1.0 for _ in range(20): actual = extension.add_one(x) torch.cuda.synchronize() start = time.perf_counter() for _ in range(100): actual = extension.add_one(x) torch.cuda.synchronize() mean_ms = (time.perf_counter() - start) * 1000.0 / 100 max_abs = float((actual - expected).abs().max()) if max_abs != 0.0: raise RuntimeError(f"extension validation failed: max_abs={max_abs}") print("architecture:", arch) print("extension:", extension_name) print(f"load/build time: {build_seconds:.3f} s") print(f"hot mean: {mean_ms:.4f} ms") print("max_abs_error:", max_abs) print("build root:", os.environ["TORCH_EXTENSIONS_DIR"]) print("staged sources:", stage_dir) if __name__ == "__main__": main() ``` 运行： ```bash cd torch_extension python build_and_test.py ``` 本次实测： ```text cold build/load 50.951 s second load 0.085 s hot kernel mean 0.0184 ms max_abs_error 0 staged sources /tmp/dcu_extension_sources_0/ ```  这组数字说明了两个经常被混为一谈的阶段：第一次需要 HIPify（把 `.cu` 转换成HIP 源码）、编译和链接；后续命中缓存后通常只需检查构建状态并加载共享库，部分PyTorch/HIPify 组合仍可能触发一次快速 relink。开发时必须分别记录冷启动（cold）和热态（hot），不能用第二次加载 掩盖首次部署成本，也不能把 50 秒编译误认为模型每次推理都这么慢。 例如，本次新容器复测 add-one 扩展时，第二次运行执行了一次快速 relink，耗时`0.301 s`；第二篇的 bias/GELU 扩展则显示 `ninja: no work to do`，热加载为`0.086 s`。因此 `0.085 s` 是一次实测结果，不是所有扩展都保证达到的固定常数。 我建议扩展缓存至少包含下面四个维度： ```text source hash + PyTorch version + HIP version + target architecture ``` 示例脚本把扩展放到： ```bash export TORCH_EXTENSIONS_DIR=/tmp/torch_extensions_$(id -u) ``` 扩展名称也包含哈希。原因是 `torch.utils.cpp_extension.load` 的构建目录并不适合让多个不同源码候选反复共用。做 A/B 测试时，如果 A、B 使用相同扩展名和缓存路径，绝对路径变化可能重写 `build.ninja`，造成重复编译甚至加载错版本。 另一个容易污染工程目录的细节是：ROCm PyTorch 会把 `.cu` HIPify 成 `.hip`。如果直接把正式源码目录传给 `load`，生成文件可能出现在源码旁。上面的脚本已经把不可变源码副本写入 `/tmp/dcu_extension_sources_//`，再从那里编译；实测原始源码目录没有生成 `.hip`、`.o` 或 `.so`，可以保持只读和干净。 ## 九、用 hipprof 找瓶颈，而不是凭感觉优化 环境跑通后，最常见的下一个误区是：“利用率不满，所以算子一定没吃满计算。”仅看 `hy-smi` 很难区分 kernel、数据搬运、内存分配、同步和 CPU 发射开销。 ### 1. 先看 HIP API 与 kernel 统计 ```bash hipprof --hip-trace --stats ./vector_add ``` 在这个最小程序的两次独立冷进程记录里，三次 `hipMalloc` 分别占 HIP API 时长的约 `81.12%` 和 `97.51%`，120 次 vector-add kernel 的平均执行时间约为`18.48–18.51 us`。百分比会受运行时初始化状态影响，不应被当成固定硬件常数；稳定结论是冷启动主要耗在分配和初始化，而不是这个简单 kernel 本身。 ### 2. kernel 名称太多时看调用栈 ```bash hipprof --kernel-stack ./your_program ``` 它能把底层 kernel 反查到上层调用位置。对于 PyTorch 模型，这比面对几百个自动生成的 kernel 名猜测来源有效得多。 ### 3. 对单个热点采性能监控计数器（PMC） ```bash hipprof --pmc --pmc-type 3 \ --kernel-name vector_add \ ./vector_add ``` 这次实测得到：workgroup 256、wave64、VGPR 8、SGPR 16、LDS 0、scratch 0、LDS bank conflict 0。其中 VGPR/SGPR 是向量/标量寄存器，LDS 类似片上共享内存，scratch 通常意味着寄存器溢出后的临时存储。PMC 的作用是解释资源和访存行为，不是提供发布成绩。 一个非常直观的证据是：普通运行时 kernel 均值约 `0.0184 ms`，打开 PMC 后变成约 `1.40 ms`，约慢 76 倍，接近两个数量级。不要把 profiler 下的墙钟时间（wall time）写进性能对比表。  ### 4. 不要一次叠满所有 profiler 选项 更稳的顺序是： 1. `--hip-trace --stats` 找 HIP API 和 kernel 热点； 2. `--kernel-stack` 找到源码调用者； 3. 只对少数 kernel 跑 PMC； 4. 退出 profiler，用包含同步的正常墙钟时间做最终 A/B。 显存诊断可以使用： ```bash hipprof --memory-usage ./your_program ``` 但进程峰值、缓存分配和业务稳定期平均占用是不同指标，报告时应说清口径。 ## 十、一个真实案例：气象 Transformer 为什么先优化搬运而不是 GEMM 在一个保持 69 通道完整输出的气象 Transformer 推理任务中，我们最初也以为GEMM 是唯一主角。`hipprof` 的稳态 kernel 分类却是： ```text direct copy 31.41% softmax 10.98% add 10.79% LayerNorm 9.39% roll 3.86% ``` 这改变了优化顺序。与其继续写一个只做搬运的新 kernel，不如把布局变化、bias、激活和归一化合到已有数据流中，减少中间张量与 kernel launch。 在一组严格 A/B 中，把 QKV 后处理与 MLP 的 bias/GELU 合入 source-only HIP路径后： ```text steady mean time 0.18896 s -> 0.17739 s (-6.12%) 2 ms sampled memory 2789.22 -> 2770.67 MiB (-18.55 MiB) 完整输出本地代理分 33.90079 -> 33.89462 (-0.00617) ```  这里的 2 ms 采样是同机 A/B 的诊断代理，只用于判断候选方向，不等同于任何平台 的正式显存统计口径。 这个案例最值得复用的不是某段模型专用代码，而是方法： 1. 先用真实 shape 做基准，不拿方阵 GEMM 代替业务形状； 2. 数值验证先于性能结论； 3. 单 kernel 快不等于全模型快； 4. 时间和显存一起测，避免用大 workspace 换来表面速度； 5. 设止损门槛，收益太小就停止，不让项目堆满不可维护的微优化。 ## 十一、常见故障对照表 | 现象 | 常见原因 | 建议检查 | | --- | --- | --- | | `hy-smi` 找不到 | 没加载环境、容器未注入设备、驱动未装 | 宿主机/容器边界、PATH、设备映射 | | `rocminfo` 只有 CPU | 当前节点无卡、设备权限缺失、Runtime/驱动不匹配 | 是否在计算节点、`/dev`、驱动日志 | | `torch.cuda.is_available()` 为 False | 装了普通 CPU/CUDA wheel 包或 HIP Runtime 不匹配 | `torch.version.hip`、wheel 来源、DTK 版本 | | `torch.cuda` 名字让人困惑 | HIP 版 PyTorch 复用 CUDA API 命名空间 | 看 `torch.version.hip`，不是看模块名 | | 脚本瞬间退出且没 Python 日志 | `set -u` 加载 `env.sh` 失败 | 先 `set +u` source，再恢复 `set -u` | | HIP kernel 报 VMFault | 越界、错误 stride、空指针、生命周期问题 | 缩小输入、检查索引、使用 VMFault/hipgdb 信息 | | 首次运行特别慢 | 即时编译（JIT）、扩展编译、库初始化、算法搜索 | 分开记录冷启动与热态，检查缓存目录 | | 修改源码却仍像旧结果 | 扩展名或缓存路径冲突 | 源码哈希命名，清晰隔离构建根目录 | | VRAM 很低仍报 HSA out of resources | HSA/设备上下文已损坏，不一定是显存 OOM | 新进程测试 `rocminfo` 和单张量；都失败则重置容器/节点 | | profiler 下速度突然变差 | trace/PMC 插桩扰动 | profiler 只做归因，退出后再做正常计时 A/B | | 编译出 illegal instruction | 抄错了目标架构或使用了不受支持的指令 | 从实机读取 `gfx...`，不要照抄其他卡参数 | 其中 HSA out-of-resources 很容易浪费时间。我们遇到过显存只占几 MiB、无人持有设备文件，但 `rocminfo` 和新启动的最小 PyTorch 进程都失败的情况。这时继续删Tensor、反复 `empty_cache()` 没有意义，设备上下文已经不是单个 Python 进程能修好的，应保存日志并重建容器或联系管理员重置节点。 ## 十二、最终验收：什么叫“环境已经配好” 我会用下面这张清单作为交付标准： - `hy-smi` 能识别预期设备； - `rocminfo` 能输出正确的 `gfx` 架构和计算单元（CU）数量； - `hipcc --version`、PyTorch、HIP Runtime 版本已记录； - BF16 PyTorch 冒烟测试数值正常，计时包含同步； - 独立 HIP vector add 编译运行、误差为 0； - PyTorch 自定义 HIP 扩展冷编译和热加载都成功； - `hipprof --stats` 能产出 API/kernel 统计； - 所有缓存、生成 `.hip`、`.so` 和 profiler 结果都写入明确的临时目录； - 长任务使用 `nohup` 或 `tmux`，并验证进程、日志、输出三者一致。 到这一步，环境才不仅是“能跑”，而是出了问题也知道从哪一层查。 ## 结语 DCU 环境配置最难的部分通常不是某一条安装命令，而是边界：驱动属于宿主机，DTK 属于用户态工具链，PyTorch wheel 包必须和 HIP Runtime 匹配，异步 kernel 的时间必须同步，profiler 的数字不能和正常运行混用。 把这几层拆开之后，排错会简单很多。先做环境指纹，再做最小 HIP/PyTorch 冒烟测试，最后才上真实模型；一旦真实模型变慢，也先用 trace 找热点，而不是凭感觉重装或重写算子。 本文正文已经完整给出环境诊断、PyTorch BF16、独立 HIP 程序和 PyTorch source-only 扩展的全部源码，可直接复制，作为新容器或新服务器的验收模板。 ## 参考资料 - 光合开发者社区工具下载（可能需要先登录）： - DTK 25.04.2《开发环境安装部署手册》 - DTK 25.04.2《开发环境使用手册》 - DTK 25.04.2《DTK Container Toolkit 使用手册》 - DTK 25.04.2《HIP C++ 编程指南》 - DTK 25.04.2《hipprof 使用手册》