> 基于 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 硬件 ``` **几个关键点**： 1. hipcc 只是 **Clang/LLVM 的包装器**。它自动调用 Clang 编译 CPU 端代码，调用 AMDGCN 后端编译 GPU 端代码，然后链接在一起。 2. ROCm 是 AMD 开源的全栈 GPU 计算平台（类比 NVIDIA 的 CUDA Toolkit）。海光将其定制为 DTK，主要改动在内核驱动层（`hydcu`/`hykcl` 替代 `amdgpu`）。 3. HSA 是底层硬件抽象层。HIP runtime 的所有 API 调用最终都转化为 HSA 调用。 我们项目实际的编译命令： ```bash $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 提交脚本头部展示了完整的资源申请： ```bash #!/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` 内核都不行 **诊断方法**： ```bash # 查看内核驱动版本 lsmod | grep -E "hydcu|hykcl" # 查看可用的 DTK 版本 ls /public/software/compiler/rocm/ # 确认 hipcc 实际使用的库 ldd ./get_climatology_hip | grep -E "hip|hsa|galaxy" ``` **解决方案**：后续我们在提交脚本中按优先级自动探测可用 DTK 版本，优先使用最新版： ```bash 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 的基本操作： 1. `cudaMalloc` → `hipMalloc` 2. `cudaMemcpy` → `hipMemcpy` 3. `cudaStream` → `hipStream` 4. kernel 源码完全不动（`__global__`、`threadIdx.x` 等完全兼容） 5. 更换编译器：`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 存储配置。 --- ## 六、附：探查硬件信息的命令 ```bash 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 工具) ``` --- *下一篇：中篇《HIP 编程实战与性能分析》—— kernel 设计、rocprof 调优、H2D/D2H 管线、常见 Bug 排查。*