海光 DCU 深度技术报告 · 上篇：硬件架构与软件生态

原创已于 2026-06-22 14:24:59 修改 · 659 阅读

5 ·

本内容遵循CC 4.0 BY-SA版权协议

GEO检测

标签

#c++

于 2026-06-22 13:43:20 首次发布

基于 MCC2026 OSTIA SST 气候态计算项目实战经验
硬件环境：Hygon C86 7185 32核 + 4×DCU Z100L × 2节点
软件栈：CentOS 7.6 / DTK 25.04 / hipcc (Clang) / HDF5 1.8 + NetCDF 4.3.0

文章目录

写在前面

某天深夜，盯着快要没电的电脑屏幕发呆，忽然想把 MCC2026 这些日子在海光 DCU 上摸爬滚打的经历整理出来。

我们在这次项目期间摔过的跟头、总结出的经验，如果只留在 CHANGELOG 里未免太可惜。

所以就有了三篇系列文章——上篇讲硬件架构和软件生态，中篇聊 HIP 编程和性能分析，下篇复盘系统级优化工程。内容都来自实战，不追求面面俱到，但保证说到的地方都是踩坑踩出来的。希望能给正在摸索 DCU 开发的同行们一点参考。

本期为上篇。

一、DCU 是什么？

DCU (Deep Computing Unit) 是海光推出的通用计算 GPU 加速卡，可以理解为国内对标 NVIDIA A100 / AMD MI50 的产品线。

DCU 的硬件架构源自 AMD CDNA (Compute DNA) 架构授权，软件栈基于 AMD 开源的 ROCm (Radeon Open Compute) 平台进行定制。关键关系：

NVIDIA:   CUDA 编程模型 → A100/H100 GPU → 闭源驱动
AMD:      HIP 编程模型  → MI50/MI250X   → ROCm 开源栈
海光 DCU: HIP 编程模型  → Z100/Z100L     → DTK (基于 ROCm 定制)

而在dcu上开发，就离不开HIP这个话题。

HIP (Heterogeneous-compute Interface for Portability) 是 AMD 设计的异构计算编程模型，语法 99% 对应 CUDA。基本对应关系：

CUDA	HIP	说明
`cudaMalloc`	`hipMalloc`	设备内存分配
`cudaMemcpy`	`hipMemcpy`	主机↔设备数据传输
`cudaStream`	`hipStream`	异步操作流
`__global__`	`__global__`	Kernel 函数声明（完全相同）
`threadIdx.x`	`threadIdx.x`	线程索引（完全相同）
`blockIdx.x`	`blockIdx.x`	块索引（完全相同）
`nvcc`	`hipcc`	编译器

一件重要的事：HIP 代码几乎就是 CUDA 代码，把 cuda 前缀换成 hip 即可编译运行在 DCU 上。这种接近 1:1 的对应关系极大降低了从 NVIDIA 生态迁移的成本。（CUDA 是 NVIDIA 的 GPU 编程模型，HIP 在语法上完全兼容它）

二、DCU Z100L 硬件微架构详解

2.1 “GPU” 和 “DCU” 的区别

传统 GPU (Graphics Processing Unit) 最初是为图形渲染设计的，包含大量固定功能的图形管线（纹理单元、光栅化器等）。DCU 则是一种通用计算加速器——它去掉了所有图形专用硬件，只保留计算单元和内存子系统，纯粹为 HPC 和 AI 计算设计。

NVIDIA 称这类产品为 “GPGPU”（如 A100 没有视频输出口），AMD 称 “CDNA 架构”，海光称 “DCU”。本质上都是同一类东西：大规模并行向量处理器。

2.2 核心概念：从计算单元到线程

理解 GPU 硬件需要先理解它的并行层次结构。这是一层一层嵌套的：

第一层：Compute Unit (CU，计算单元)

Z100L 有 60 个 CU。可以把每个 CU 想象成一个独立的"小 CPU 核心"——有自己的算术单元、寄存器和本地缓存。60 个 CU 同时工作，互不干扰。

第二层：SIMD 单元

每个 CU 内部有 4 组 SIMD (Single Instruction Multiple Data，单指令多数据流)。一组 SIMD 同时执行一条指令，但这条指令作用于 16 个不同的数据。

类比：老师对全班 64 个学生说"分小组朗诵课文段落"——所有学生做完全相同的动作，但每个人处理的是自己分配的不同的那个段落。这就是 SIMD：一个指令流，多个数据流。

参数	数值	说明
CU 数	60	独立计算核心
SIMD/CU	4	每 CU 内的并行流水线
SIMD 宽度	16	每条指令同时处理 16 个数据
总并行度	60×4×16=3840	理论上每周期可同时执行这么多操作

第三层：Wavefront（波前，简称 Wave）

在 AMD/DCU 术语中，一个 wavefront 是 64 个线程组成的最小调度单位。一个 SIMD 每周期只推进一个 wavefront 的其中 16 个线程，所以 64 个线程需要 4 个周期执行完一条指令。

64 个线程 ÷ 16 宽度 = 4 个周期/wavefront/指令。

这和 NVIDIA 的 “warp”（32 线程）概念类似但宽度不同：AMD/DCU 是 wavefront=64 线程，NVIDIA 是 warp=32 线程。

第四层：线程

线程是程序员看到的最小粒度。在我们的 kernel 中，每个线程处理一个海洋格点（pixel）的完整 330 样本统计。1,038,240 个格点 = 1,038,240 个线程。

2.3 寄存器文件与 Occupancy

每个 SIMD 有一组寄存器文件——一小块极快的片上存储器，用于存放线程的局部变量。

关键限制：寄存器文件大小是固定的。

Z100L (gfx906) 每 SIMD 有 256 个 VGPR（Vector General Purpose Register，向量通用寄存器），每个 32-bit = 4 bytes。总量 = 256 × 4 bytes = 1 KB / SIMD。

我们的 kernel 每个线程使用了 160 个 VGPR。每个 wavefront 有 64 个线程，所以一个 wavefront 需要的寄存器总数 = 64 × 160 = 10,240 个 VGPR。

但一个 SIMD 只有 256 个 VGPR！10,240 >> 256，怎么办？

答案是寄存器不是一次分配给整个 wavefront 的。GPU 采用一种叫 “barrel scheduling” 的技术——多个 wavefront 轮流使用 SIMD，每个 wavefront 活跃时占用全部 256 个 VGPR，被切换走时状态保存在其他地方。

每 SIMD 能同时驻留的 wavefront 数量，由寄存器文件大小决定：

每 SIMD 可驻留 wavefront = 256 VGPR ÷ (每线程 VGPR × 64 线程 / 4 周期)

160 VGPR/线程的情况下，每 SIMD 只能驻留约 1 个 wavefront。这就是 Occupancy = 25% 的来源——SIMD 理论上最多驻留 4 个 wavefront（256 / 64 = 4，当每线程只用一个 VGPR 时），而我们现在只能用 1 个。

Occupancy 低不一定是坏事。我们的 kernel 是内存瓶颈（每秒要传输的数据 > HBM 能提供的带宽），更多 wavefront 只会让 HBM 控制器排队更长，不会让计算变快。事实上，减少 wavefront 可以减少竞争、改善单个 wavefront 的 L2 缓存命中率。

CUDA 开发者经常会有一个概念误导——认为 occupancy 必须越高越好——但在纯内存瓶颈场景下，这条准则不成立。

2.4 缓存层级全解析

DCU 的缓存放缓有几个层次，和 CPU 的 L1/L2/L3 概念类似但细节不同：

寄存器 (0周期) → L1 (≈30周期) → L2 (≈200周期) → HBM (≈400周期)

VGPR (向量通用寄存器)

容量：256 × 32-bit / SIMD = 1 KB/SIMD
延迟：0 周期（编译器直接访问，无需 load/store 指令）
作用：存放线程局部变量、中间计算结果
注意：寄存器是编译器管理的，程序员无法直接控制。但数组大小、变量数量影响编译器分配
VGPR 不是缓存，不会出现 miss——编译器保证每个变量都能在寄存器文件中找到位置

L1 Vector Cache

容量：16 KB / CU
延迟：≈30 个周期
作用：缓存全局内存（HBM）的数据
关键特性：仅通过合并访问（coalesced access）才有效

什么是合并访问？当 64 个线程构成的 warp 访问的内存地址落在同一个或相邻的几个 cache line（64 bytes = 32 个 short 值）内时，L1 可以一次性载入。如果 64 个线程的访问散布在 64 个不同的 cache line 上，L1 完全失效——每次访问都要穿透到 L2 甚至 HBM。

LDS (Local Data Share，本地数据共享)

容量：64 KB / CU
延迟：≈20 个周期
作用：CU 内线程间的共享数据交换
特点：程序员显式管理（类似 CUDA 的 __shared__ 内存）
本项目未使用 LDS——每个线程独立处理自己的像素，不需要线程间通信

L2 Cache

容量：4 MB / 芯片（所有 CU 共享）
延迟：≈200 个周期
作用：HBM 之前的最后一级缓存
L2 命中率 67% 意味着 33% 的访问直接穿透到 HBM——这是数据集合太大（跨年 stride 2MB）的结果

HBM2 (High Bandwidth Memory)

容量：32 GB / 卡
带宽：约 1 TB/s (理论上，ECC 开启后实际约 900 GB/s)
延迟：≈400 个周期
特点：通过硅中介层（silicon interposer）与 GPU die 垂直堆叠，1024-bit 超宽总线
MemUnitBusy 99.7% 意味着 HBM 控制器在几乎所有时钟周期内都在传输数据

HBM 和普通 DDR 的区别：DDR 内存条是水平插在主板上的 DIMM，64-bit 总线，带宽约 50 GB/s。HBM 是垂直堆叠在 GPU 芯片旁的 DRAM 堆栈，1024-bit 总线，带宽约 1 TB/s——宽 16 倍，快 20 倍。代价是不能扩展容量（32 GB 定死）且成本极高。

2.5 CFX906 ISA

Z100L 的指令集架构（ISA）代号为 gfx906，与 AMD MI50/MI60 完全相同。这意味着：

为 MI50/MI60 编译的二进制可以直接在 Z100L 上运行（前提是 ROCm 内核驱动兼容）
所有 AMD ROCm 的 gfx906 kernel 优化技巧完全适用于 Z100L
编译器目标：--offload-arch=gfx906

ISA 的一些关键特性：

sramecc+：片上 SRAM 和 HBM 均启用 ECC（错误校验），数据可靠性提升但有效带宽损失约 3-5%
xnack-：不支持自动页面迁移。host 端内存由程序员通过 hipHostRegister 显式锁定
flat addressing：host 和 device 共享 64-bit 地址空间（但我们未使用此特性）

三、DTK (DCU Toolkit) 软件栈详解

3.1 软件栈层次

一个 HIP 程序从源码到在 DCU 上运行，经过以下层次：

用户代码 (get_climatology_hip.cpp)
    │
    ▼ hipcc (Clang/LLVM 编译器)
AMD GPU 设备码 (gfx906)  +  x86-64 主机码
    │
    ▼ libgalaxyhip.so (海光 HIP 运行时)
HSA (Heterogeneous System Architecture) 运行时
    │
    ▼ 用户态：libhsa-runtime64.so
内核态：hydcu.ko + hykcl.ko (海光定制 ROCm 驱动)
    │
    ▼ PCIe 总线
DCU Z100L 硬件

几个关键点：

hipcc 只是 Clang/LLVM 的包装器。它自动调用 Clang 编译 CPU 端代码，调用 AMDGCN 后端编译 GPU 端代码，然后链接在一起。
ROCm 是 AMD 开源的全栈 GPU 计算平台（类比 NVIDIA 的 CUDA Toolkit）。海光将其定制为 DTK，主要改动在内核驱动层（hydcu/hykcl 替代 amdgpu）。
HSA 是底层硬件抽象层。HIP runtime 的所有 API 调用最终都转化为 HSA 调用。

我们项目实际的编译命令：

$HIPCC -std=c++14 -O3 -march=native \
    --offload-arch=gfx906 \
    -I"${NC_INC}" \
    ${H5_INC:+-I"${H5_INC}"} \
    -o get_climatology_hip get_climatology_hip.hip \
    -L"${NC_LIB}" -l:libnetcdf.so.7 -l:libhdf5.so.8 -Wl,-rpath,"${NC_LIB}"

--offload-arch=gfx906 告诉 Clang/LLVM 的 AMDGCN 后端为 Z100L 的 ISA 生成设备码。编译产出是一个fat binary——同时包含 x86-64 主机码和 gfx906 GPU 设备码，运行时自动加载对应部分。

SLURM 提交脚本头部展示了完整的资源申请：

#!/bin/bash
#SBATCH -p kshdmcc2026
#SBATCH -N 2
#SBATCH -n 64
#SBATCH --gres=dcu:4
#SBATCH --exclusive
#SBATCH --mem=0
#SBATCH -J MCC_HIP

--gres=dcu:4 申请每节点 4 张 DCU 加速卡，--mem=0 申请节点全部内存（约 128 GB DDR4），--exclusive 独占节点避免其他作业干扰 PCIe 和 Lustre I/O。

3.2 DTK 版本兼容性——实测踩坑

DTK 的版本兼容性问题是我们遇到的最隐蔽的 bug 之一。问题的根源在于海光的内核驱动 hydcu 和用户态 DTK 版本是独立升级的：

比赛时集群管理员可能升级了内核驱动（hydcu.ko），但没有同时升级 /public/software/compiler/rocm/ 下的 DTK 用户态库
旧版 DTK 22.10.1（基于 ROCm 5.2）的用户态库与新版内核驱动（基于 ROCm 5.7）不兼容
导致了以下症状：所有 kernel launch 都报 “invalid device function (98)”——连一个最简单的 out[i] = i 内核都不行

诊断方法：

# 查看内核驱动版本
lsmod | grep -E "hydcu|hykcl"

# 查看可用的 DTK 版本
ls /public/software/compiler/rocm/

# 确认 hipcc 实际使用的库
ldd ./get_climatology_hip | grep -E "hip|hsa|galaxy"

解决方案：后续我们在提交脚本中按优先级自动探测可用 DTK 版本，优先使用最新版：

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

经验法则：如果 invalid device function 出现了，大概率是 DTK 版本和内核驱动不匹配，不是代码问题。

3.3 与 NVIDIA CUDA 生态的对比

方面	NVIDIA	AMD/DCU
编程模型	CUDA	HIP (≈ CUDA)
编译器	nvcc	hipcc
运行时	libcudart.so	libamdhip64.so
Profiler	nvprof / Nsight	rocprof
管理工具	nvidia-smi	rocm-smi
内核驱动	nvidia.ko	hydcu.ko + hykcl.ko
线程束	warp = 32 线程	wavefront = 64 线程
硬件架构	SM (Streaming Multiprocessor)	CU (Compute Unit)
开源程度	闭源驱动 + 部分开源库	ROCm 全栈开源（海光版驱动闭源）

迁移 CUDA 代码到 DCU 的基本操作：

cudaMalloc → hipMalloc
cudaMemcpy → hipMemcpy
cudaStream → hipStream
kernel 源码完全不动（__global__、threadIdx.x 等完全兼容）
更换编译器：nvcc → hipcc

四、Hygon C86 7185 主机处理器

虽然 GPU 是主角，但 host CPU 也值得了解：

参数	数值
架构	Zen1 (AMD 授权)
核心数	32 核
线程	1 线程/核（无 SMT/超线程）
频率	2.0 GHz
NUMA	4 个节点，每节点 8 核 + 32 GB DDR4
L1d/L1i	32 KB / 64 KB (每核)
L2	512 KB (每核独占)
L3	8 MB (每 4 核共享)
SIMD	MMX, SSE, SSE2, SSE4.1/4.2, AVX, AVX2
缺失	无 AVX-512 (Zen1 通病)

关键限制：无 AVX-512。AVX-512 是 Intel 在 Skylake-X 之后引入的 512-bit SIMD 指令集，单指令可同时处理 16 个 float 或 8 个 double。没有 AVX-512 意味着 CPU 端的向量化只能到 256-bit（一次 8 个 float / 4 个 double），浮点吞吐理论仅为 AVX-512 的一半。

不过对我们的项目影响不大——计算主力在 GPU 上，CPU 只做 I/O 和数据格式转换。

五、ParaStor (Lustre) 并行文件系统

我们的数据存储在 /public 挂载的 ParaStor 文件系统上，底层是 Lustre，ParaStor 是中科曙光基于开源 Lustre 的闭源商业软件。

Lustre 是 HPC 领域最主流的并行文件系统，设计目标是对成千上万个节点同时提供高聚合 I/O 带宽。它的关键概念：

OSS (Object Storage Server)：存储数据的服务器，管理物理磁盘
OST (Object Storage Target)：OSS 上的存储卷
MDS (Metadata Server)：管理文件元数据（目录结构、文件属性）
Client：计算节点上的 Lustre 客户端内核模块

Lustre 的文件数据被条带化（stripe）分散在多个 OST 上，大文件可以同时从多个 OST 并行读写。但小文件（我们的 NC 文件每个 ~4 MB）可能只落在一个 OST 上，并发读大量小文件会导致对特定 OST 的竞争。

从 f16r4n13 的 df -h 输出可以看到集群的 Lustre 规模：

Filesystem       Size  Used Avail Use% Mounted on
ks_p300s_public   61P   42P   19P  70% /public

61 PB 的总容量。42 PB 已使用。这是典型的大规模 HPC 存储配置。

六、附：探查硬件信息的命令

lscpu -C                          # CPU 架构、缓存
numactl --hardware                 # NUMA 拓扑和内存分布
rocminfo                          # DCU 详细信息
rocm-smi --showhw                 # DCU 硬件概览
rocm-smi --showtopo               # DCU 间的拓扑连接
lspci | grep -iE "vga|display|co-process"  # PCIe 总线上的 GPU
cat /sys/devices/system/node/node*/distance  # NUMA 距离矩阵
cat /sys/devices/system/cpu/cpu0/cache/index*/size  # 缓存大小
lstopo-no-graphics --no-io        # 拓扑图 (hwloc 工具)