国产计算卡海光DCU保姆级教程
国产计算卡海光DCU保姆级教程
本博客基于曙光智算的超算互联网平台 https://www.scnet.cn/ 进行的实践与教学，其平台界面相较于北京超算云等平台更加现代，操作十分方便，几乎不需要使用vscode等ssh工具就可以获得相当便捷的运行环境。如果想要进行本地开发，也有一定的参考价值。目前我自己也在这一平台进行海洋相关科学计算和vllm推理加速的任务。
一、DCU简介
1.1 基本定义
DCU 全称Deep Computing Unit（深度计算单元），是海光信息自研的国产通用 GPGPU 并行加速芯片，对标英伟达 A/H 系列 GPU，主打 AI 大模型训练、推理、高性能科学计算，是国内成熟商用通用算力卡之一。
1.2 产品线迭代
-
深算一号（DCU 8000/Z100）
7nm 工艺，32/64GB HBM2 显存，支持 FP64 双精度，2021 年量产；兼顾超算仿真与中小模型训练，国产化算力起步主力。
-
深算二号（DCU 8200/K100）
迭代升级，96GB HBM3，FP32 峰值 256TFLOPS，大模型训练性能大幅提升，2025 年规模商用，适配千亿参数大模型训推一体。
-
深算三号（8300）
5nm 工艺，128GB 超大 HBM3E 显存，算力翻倍，对标 H100，研发推进中。
1.3 硬件核心特性
- 全精度算力：FP64/FP32/FP16/INT8/FP8 全覆盖，兼顾科学计算与 AI 混合精度训练；
- HBM 高带宽显存，单卡 TB/s 级带宽，适配大模型海量参数；
- PCIe4.0 + xGMI 多卡高速互联，支持单机多卡、分布式算力集群扩展；
- 服务器标准加速卡形态，风冷 / 液冷可选，适配国产 x86 服务器（海光 CPU 配套最优）。
1.4 软件生态 DTK（核心优势）
自研DTK 开发工具栈，兼容 ROCm 生态，提供类 CUDA 迁移接口：
- 少量修改即可迁移原有 CUDA 代码，迁移成本低；
- 原生适配 PyTorch、TensorFlow、飞桨等全部主流深度学习框架；
- 完整编译器、算子库、调试 / 性能分析工具；
- 全面适配 ChatGLM、Qwen、Llama、文心等国产 / 开源大模型。
1.5 应用场景
- AI 领域：大语言模型训练微调、AIGC、计算机视觉、语音识别；
- HPC 超算：气象、流体仿真、能源勘探、生物医药分子计算；
- 行业信创：金融、政务、运营商国产化智算中心算力底座。
二、基础入门：从0构建第一个DCU程序
本节在曙光超算互联网的 DCU 队列 下完成，所有命令均可在交互式
srun或sbatch提交方式下使用。如果你用的是别的平台（如北京超算云、阿里云弹性加速实例），逻辑一致，只是 DTK 路径和 SLURM 队列名不同。
2.1 第一步：写一个 SLURM 任务提交脚本
在超算上跑 DCU 程序，第一步永远不是写代码，而是先把资源申请好。海光集群几乎都使用 SLURM 作为调度器，我们用 sbatch 把任务扔到队列里执行。
最小可用模板 hello_dcu.sh
#!/bin/bash
#SBATCH -p dcu # 提交到 DCU 队列（不同集群名字可能不同：kshdmcc202x、dcuh100 等）
#SBATCH -N 1 # 申请 1 个计算节点
#SBATCH -n 8 # 8 个 CPU 核心（DCU 程序通常需要 CPU 线程做 I/O 与任务调度）
#SBATCH --gres=dcu:1 # 关键：每节点申请 1 张 DCU
#SBATCH -J hello_dcu # 作业名，方便 squeue 查询
#SBATCH --exclusive # 独占节点（避免别人和你抢 DCU 显存/带宽）
# 在这里放你的 DCU 程序
./hello_dcu
常用参数速查
| 参数 | 含义 | 备注 |
|---|---|---|
-p 队列名 |
提交到哪个分区 | DCU 队列一般叫 dcu / dcuh100 / kshdmcc2026 |
-N |
节点数 | 多机分布式才用 >1 |
-n |
CPU 任务总数 | 不等于核数，是 SLURM 任务槽数 |
--gres=dcu:N |
每节点 DCU 数量 | 整机申请要写成 --gres=dcu:4 |
--cpus-per-task |
每个任务占几个 CPU 核 | 配合 -n 一起算总核数 |
--mem=0 |
申请节点全部内存 | 大数据预加载时建议加上 |
提交与查看
# 提交
sbatch hello_dcu.sh
# 查看自己的作业
squeue -u $USER
# 取消（作业 ID 从 squeue 输出第一列获取）
scancel <作业ID>
提交后用 squeue 能看到状态是 PD（排队）或 R（运行）。等状态变成 R 后，程序就会在分配到的节点上开始执行，结果写入 hello_dcu.log。
2.2 第二步：寻找依赖（DTK、hipcc、运行时库）
DCU 的开发工具叫 DTK（DCU Toolkit），它本质上是 AMD ROCm 的国产化发行版。所以我们写代码用的 API 是 HIP（Heterogeneous-Compute Interface for Portability）——和 CUDA 几乎一一对应。
编译器在哪？
hipcc 一定藏在 DTK 安装目录里。不同集群的 DTK 版本不一定一样，先在交互节点上探测一下：
ls /public/software/compiler/rocm/
# 典型输出：dtk-25.04.4 dtk-24.04.1 dtk-22.10.1
找到版本后，编译器路径就是：
HIPCC=/public/software/compiler/rocm/dtk-25.04.4/hip/bin/hipcc
运行时库
DCU 程序运行时需要加载 ROCm 的动态库，否则会报 libamdhip64.so 找不到。统一加到 LD_LIBRARY_PATH 即可：
DTK=/public/software/compiler/rocm/dtk-25.04.4
export LD_LIBRARY_PATH=$DTK/lib:$DTK/lib64:$DTK/hip/lib:$LD_LIBRARY_PATH
一键探测脚本（推荐加到 ~/.bashrc）
把下面这一段加到作业脚本开头，自动找最新可用的 DTK，避免写死路径：
DTK_VER=""
for try_ver in "dtk-25.04.4" "dtk-24.04.1" "dtk-22.10.1"; do
candidate="/public/software/compiler/rocm/${try_ver}/hip/bin/hipcc"
if [ -f "$candidate" ]; then
DTK_VER="$try_ver"
break
fi
done
[ -z "$DTK_VER" ] && { echo "找不到 DTK，请联系超算管理员"; exit 1; }
DTK=/public/software/compiler/rocm/$DTK_VER
export HIPCC=$DTK/hip/bin/hipcc
export LD_LIBRARY_PATH=$DTK/lib:$DTK/lib64:$DTK/hip/lib:$LD_LIBRARY_PATH
echo "使用 DTK: $DTK_VER"
2.3 第三步：编译与链接
最小编译命令
假设你只有一个源文件 hello_dcu.cpp，里面包含 HIP kernel：
$HIPCC -std=c++14 -O2 \
--offload-arch=gfx906 \
-o hello_dcu hello_dcu.cpp
新手最常踩的坑：忘了加
--offload-arch=gfx906，编译能过，但运行时hipGetDeviceCount返回 0，程序直接退出而无明显报错。gfx906对应海光深算一号（Z100）和 AMD MI50/MI60。如果你是深算二号 K100，需要换成gfx940或gfx942（具体看集群文档）。
关键参数
| 参数 | 作用 |
|---|---|
-std=c++14 |
HIP 头文件要求至少 C++14 |
-O2 / -O3 |
优化等级，DCU 计算代码建议 -O3 |
--offload-arch=gfx906 |
目标 DCU 架构，必加 |
-I/path |
第三方头文件目录（如 netcdf、hdf5） |
-L/path -lxxx |
链接第三方库 |
-Wl,-rpath,/path |
把库路径烧进可执行文件，避免运行时找不到 |
完整实战例子（链接 NetCDF + HDF5）
这是真实项目里常见的"既要科学库又要 GPU"的编译命令：
NC_LIB=/public/home/achwjznh4b/install/bin/glnxa64
NC_INC=/public/sourcecode/amber18/AmberTools/src/netcdf-4.3.0/include
$HIPCC -std=c++14 -O3 \
--offload-arch=gfx906 \
-I"$NC_INC" \
-o myapp myapp.cpp \
-L"$NC_LIB" -l:libnetcdf.so.7 -l:libhdf5.so.8 \
-Wl,-rpath,"$NC_LIB"
注意：-l:libnetcdf.so.7 这种写法是直接链接具体的 .so 文件名，比 -lnetcdf 更稳（避免版本冲突），上面的"科学计算库"就是读取netCDF的库。
2.4 实例 1：查看当前节点的 DCU 设备信息
在写第一个 kernel 之前，先确认你申请到的节点上确实有 DCU，并且版本对。有两种方式：
方式 A：Shell 一行命令（最快）
rocm-smi
输出类似：
如果想看更详细的信息（频率、显存、温度）：
rocm-smi --showmeminfo vram --showuse --showtemp
方式 B：写一个 C++ 程序通过 HIP API 查询
文件 query_dcu.cpp：
#include <cstdio>
#include <hip/hip_runtime.h>
int main() {
int n = 0;
hipError_t err = hipGetDeviceCount(&n);
if (err != hipSuccess) {
printf("HIP 错误: %s\n", hipGetErrorString(err));
return 1;
}
printf("检测到 %d 张 DCU/GPU\n", n);
for (int i = 0; i < n; i++) {
hipDeviceProp_t prop;
hipGetDeviceProperties(&prop, i);
printf("\n=== DCU %d ===\n", i);
printf("名称 : %s\n", prop.name);
printf("架构 : %d.%d\n", prop.major, prop.minor);
printf("全局显存 : %.1f GB\n", prop.totalGlobalMem / 1e9);
printf("每个 block 的 shared mem: %zu KB\n", prop.sharedMemPerBlock / 1024);
printf("每个 block 的最大线程数: %d\n", prop.maxThreadsPerBlock);
printf("多处理器数量: %d\n", prop.multiProcessorCount);
}
return 0;
}
编译运行（记得按照第三步操作！）：
$HIPCC -std=c++14 -O2 --offload-arch=gfx906 -o query_dcu query_dcu.cpp
./query_dcu
预期输出（节选）：
看到这些数字基本就能判断这张卡的算力档位。
gfx906 + 60 CU就是海光深算一号/AMD MI50 这一档。
2.5 实例 2：第一个真正的 DCU 程序——向量加法
这是 GPU 编程的"Hello World"：计算 C = A + B，A、B、C 各 1024 万个 float。
文件 vec_add.cpp：
#include <cstdio>
#include <cstdlib>
#include <hip/hip_runtime.h>
// HIP_CHECK: 任何 HIP API 失败立刻报错退出
#define HIP_CHECK(cmd) do { \
hipError_t e = cmd; \
if (e != hipSuccess) { \
printf("HIP 错误 %s:%d: %s\n", \
__FILE__, __LINE__, hipGetErrorString(e)); \
exit(1); \
} \
} while(0)
// __global__ 标记的函数会被编译成 DCU 上的 kernel
__global__ void vec_add(const float* A, const float* B, float* C, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
C[i] = A[i] + B[i];
}
}
int main() {
const int N = 1 << 24; // 1677 万个元素
const size_t bytes = N * sizeof(float);
// 1) 分配主机端内存并初始化
float *hA = (float*)malloc(bytes);
float *hB = (float*)malloc(bytes);
float *hC = (float*)malloc(bytes);
for (int i = 0; i < N; i++) {
hA[i] = 1.0f;
hB[i] = 2.0f;
}
// 2) 分配 DCU 显存
float *dA, *dB, *dC;
HIP_CHECK(hipMalloc(&dA, bytes));
HIP_CHECK(hipMalloc(&dB, bytes));
HIP_CHECK(hipMalloc(&dC, bytes));
// 3) 把数据从内存拷到 DCU（H2D）
HIP_CHECK(hipMemcpy(dA, hA, bytes, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(dB, hB, bytes, hipMemcpyHostToDevice));
// 4) 启动 kernel
int block = 256;
int grid = (N + block - 1) / block;
vec_add<<<grid, block>>>(dA, dB, dC, N);
HIP_CHECK(hipGetLastError()); // 检查 kernel 启动是否成功
HIP_CHECK(hipDeviceSynchronize()); // 等待 DCU 算完
// 5) 把结果拷回内存（D2H）
HIP_CHECK(hipMemcpy(hC, dC, bytes, hipMemcpyDeviceToHost));
// 6) 验证结果
printf("C[0] = %f（应为 3）\n", hC[0]);
printf("C[N-1] = %f（应为 3）\n", hC[N - 1]);
// 7) 释放资源
HIP_CHECK(hipFree(dA));
HIP_CHECK(hipFree(dB));
HIP_CHECK(hipFree(dC));
free(hA);
free(hB);
free(hC);
return 0;
}
在 SLURM 作业里运行
把上面命令拼成 vec_add.sh 提交：
#!/bin/bash
#SBATCH -p kshdmcc2026
#SBATCH -N 2
#SBATCH -n 64
#SBATCH --exclusive
#SBATCH --gres=dcu:4
#SBATCH -J MCC_RUN
DTK_VER=""
for try_ver in "dtk-25.04.4" "dtk-24.04.1" "dtk-22.10.1"; do
candidate="/public/software/compiler/rocm/${try_ver}/hip/bin/hipcc"
if [ -f "$candidate" ]; then
DTK_VER="$try_ver"
break
fi
done
[ -z "$DTK_VER" ] && { echo "找不到 DTK，请联系超算管理员"; exit 1; }
DTK=/public/software/compiler/rocm/$DTK_VER
export HIPCC=$DTK/hip/bin/hipcc
export LD_LIBRARY_PATH=$DTK/lib:$DTK/lib64:$DTK/hip/lib:$LD_LIBRARY_PATH
echo "使用 DTK: $DTK_VER"
$HIPCC -std=c++14 -O2 --offload-arch=gfx906 -o vec_add vec_add.cpp
./vec_add
提交运行结果：
小结
到这里，你已经掌握了 DCU 编程的完整最小链路：
写 SLURM 脚本申请资源 → 探测/配置 DTK → hipcc 编译 → 提交作业运行
(2.1 节) (2.2 节) (2.3 节) (2.5 节)
下一节我们会从"向量加法"升级到真实的科学计算场景：利用 DCU 的多卡并行 + HIP 流并发，做一个滑动窗口气候态统计程序。
DAMO开发者矩阵，由阿里巴巴达摩院和中国互联网协会联合发起，致力于探讨最前沿的技术趋势与应用成果，搭建高质量的交流与分享平台，推动技术创新与产业应用链接，围绕“人工智能与新型计算”构建开放共享的开发者生态。
更多推荐
30
0- 0
已为社区贡献1条内容
所有评论(0)