# 海光DCU深度使用手记（上）：从零开始——初识DCU与HIP移植实战 ## 前言 今年参加MCC2026超算竞赛时，我们的赛题之一是OSTIA SST海表温度的气候态计算——用30年（1991-2020）逐日卫星观测数据，对6到8月每一天计算出气候态平均值和P90极端阈值。原始MATLAB代码在超算单节点上要跑**17个半小时**（62768秒），非常夸张。我们的最终优化版本跑到了9秒出头，加速比接近7000倍。 在这个过程中，除了CPU端的极限优化（滑动窗口、多进程并行I/O、异步预取流水线等），我们还专门做了一轮GPU加速——用的不是NVIDIA的卡，而是国产的海光DCU。这篇文章就基于这次实战经历，聊聊海光DCU到底怎么用、坑在哪里、跟NVIDIA GPU有什么异同。 三篇文章的结构是这样： - **上篇（本篇）**：海光DCU是什么、HIP/ROCm开发环境怎么搭、从CPU代码到DCU的首次移植过程 - **中篇**：DCU深度优化——显存管理、内核设计、I/O流水线、时序分块策略，以及各种踩坑记录 - **下篇**：DCU研发的经验总结、与CPU方案的对比分析、给DCU初学者的实战建议 ## 一、海光DCU到底是什么 先说个很多人容易搞混的事。海光DCU不是GPU——至少官方不这么叫。DCU的全称是Deep Computing Unit（深度计算器），本质上是一块通用计算加速卡。但从程序员的角度看，你可以把它理解成一块兼容AMD ROCm生态的GPGPU。 关键点来了：**海光DCU的软件栈走的是AMD ROCm/HIP路线**，不是NVIDIA的CUDA。这意味着什么？ - 你不能用 `nvcc`编译，得用 `hipcc` - 你不能 `#include `，得用 `#include ` - 你不能 `cudaMalloc`，得用 `hipMalloc` - 你不能 `<<>>`启动CUDA kernel，但你**可以**用完全一样的 `<<>>`语法启动HIP kernel 也就是说，HIP在语法层面做了大量兼容——如果你之前写过CUDA，迁移到DCU的代码改动比你想象的小得多。很多情况下甚至可以直接写一个脚本把 `cuda`全局替换成 `hip`就完事。当然，实际优化的时候远没有这么简单，后面会细说。 ## 二、我们的超算环境 先交代一下我们这次用的具体环境，方便读者有个参照： - **DCU型号**：海光Z100系列（`gfx906`架构，兼容ROCm） - **DTK版本**：DTK 25.04.4（海光的ROCm定制工具链） - **编译器**：`hipcc`（基于clang，路径在 `/public/software/compiler/rocm/dtk-25.04.4/bin/hipcc`） - **数学库**：NetCDF 4.4.1 + GCC - **调度系统**：SLURM，需要 `--gres=dcu:1`申请DCU资源 - **节点配置**：单节点4块DCU，搭配32核CPU 环境模块加载很简单，就一行： ```bash module load rocm/dtk-25.04.4 ``` SLURM脚本里申请DCU资源： ```bash #SBATCH -p kshdmcc2026 #SBATCH -N 1 #SBATCH -n 4 #SBATCH --gres=dcu:1 #SBATCH --exclusive ``` 注意这里用的是 `--gres=dcu:1`而非NVIDIA GPU的 `--gres=gpu:1`——SLURM通过不同的资源类型名称区分DCU和GPU，需要根据集群的具体配置来确定。一个节点上有4块DCU可用，如果任务需要多卡并行，可以把数字调大。 编译命令长这样： ```bash /public/software/compiler/rocm/dtk-25.04.4/bin/hipcc \ -O3 -std=c++11 -fopenmp \ --offload-arch=gfx906 \ -I/opt/hpc/software/mathlib/netcdf/4.4.1/gcc/include \ -L/opt/hpc/software/mathlib/netcdf/4.4.1/gcc/lib \ -o get_climatology_gpu get_climatology_gpu.cpp \ -lnetcdf -lm ``` 注意 `--offload-arch=gfx906`——这个是AMD GPU的架构代号，海光DCU因为兼容ROCm所以也用这套命名。如果你用的是其他型号的DCU，这个参数可能会不一样，要看具体硬件对应的架构代号。 ## 三、从CPU到DCU：第一次移植 我们不是一上来就直接写GPU代码的。在此之前，我们已经在CPU上做了十几轮的迭代优化（具体过程可以看我之前的优化报告），CPU版本最终稳定在9-13秒。做DCU加速的动机很简单：CPU版本虽然已经很快了，但I/O等待时间始终占大头——在共享的Lustre文件系统上，3060个NetCDF文件的读取时间波动很大，有时候I/O能占到总时间的60%以上。我们想看看GPU/DCU能不能在计算侧进一步压榨，同时通过一次性的全量预读策略把I/O波动隔离掉。 ### 3.1 代码结构设计 第一次移植的时候，我们没有简单粗暴地把整个程序扔到DCU上。而是做了一个分阶段的流水线设计： ``` Phase 1 (CPU): 一次性预读全部3060个文件到CPU内存（~16GB） Phase 2 (H2D): 单次大块上传到DCU显存 Phase 3 (DCU): 在DCU上完成全部92天的计算 Phase 4 (D2H): 单次大块下载结果回CPU Phase 5 (CPU): 批量写出NetCDF结果文件 ``` 这个设计的核心思路是**I/O和计算彻底分离**——所有文件读取在计算开始前一次性完成，DCU计算期间零I/O。这样做的好处很明显：DCU的计算时间是完全确定的，不受文件系统状态影响。 对于海光DCU来说，这个策略还有一个额外的好处：虽然DCU卡上的HBM2显存带宽很高（~1 TB/s级别），但CPU到DCU的PCIe带宽是有限的。如果把数据切成小块反复搬运，PCIe开销会吃掉很多加速收益。**一次大块传输**才能充分利用PCIe带宽。 ### 3.2 关键代码：HIP初始化与设备信息 拿到DCU之后第一件事，肯定是先看看手里这块卡什么配置： ```cpp #include int gpu_count; hipGetDeviceCount(&gpu_count); printf("检测到 %d 块DCU设备\n", gpu_count); hipDeviceProp_t dev_prop; hipGetDeviceProperties(&dev_prop, 0); printf("DCU: %s\n", dev_prop.name); printf("显存: %.1f GB\n", dev_prop.totalGlobalMem / 1e9); printf("计算单元数: %d\n", dev_prop.multiProcessorCount); hipSetDevice(0); // 选择第0块DCU ``` 这段代码跟CUDA的 `cudaGetDeviceProperties`几乎一模一样。实际跑出来的信息大概是：显存16GB或32GB HBM2，几十个计算单元。具体数字取决于你用的DCU型号。 ### 3.3 HIP错误处理宏 HIP的API调用基本都返回 `hipError_t`，每个调用都要检查返回值。跟CUDA一样，写个宏省事： ```cpp #define HIP_CHECK(call) do { \ hipError_t err = (call); \ if (err != hipSuccess) { \ fprintf(stderr, "HIP error %s:%d — %s\n", \ __FILE__, __LINE__, hipGetErrorString(err)); \ exit(1); \ } \ } while(0) ``` 使用方法： ```cpp HIP_CHECK(hipMalloc(&gpu_pool, total_bytes)); HIP_CHECK(hipMemcpy(gpu_pool, cpu_data, total_bytes, hipMemcpyHostToDevice)); HIP_CHECK(hipDeviceSynchronize()); HIP_CHECK(hipFree(gpu_pool)); ``` 跟CUDA的区别几乎只有前缀——`cuda`换成 `hip`。 ### 3.4 显存分配与数据传输 我们这次的数据规模：30年 × 102天（DOY 147-248，窗口覆盖范围）× 1440×721格点 × 4字节(float) = 约12.7GB。加上输出缓冲区（92天×2变量×1440×721×8字节(double)）约1.5GB，DCU显存总需求约14-15GB。 ```cpp // DCU显存分配：输入数据池 float *gpu_pool; HIP_CHECK(hipMalloc(&gpu_pool, N_TOTAL_GRIDS * grid_cells * sizeof(float))); // DCU显存分配：输出缓冲区 double *gpu_clim_out, *gpu_p90_out; HIP_CHECK(hipMalloc(&gpu_clim_out, TOTAL_DOYS * grid_cells * sizeof(double))); HIP_CHECK(hipMalloc(&gpu_p90_out, TOTAL_DOYS * grid_cells * sizeof(double))); // DCU显存分配：陆地掩码 unsigned char *gpu_land_mask; HIP_CHECK(hipMalloc(&gpu_land_mask, grid_cells * sizeof(unsigned char))); ``` 数据上传——用 `hipMemcpy`从主机到设备（H2D）： ```cpp HIP_CHECK(hipMemcpy(gpu_pool, full_data, full_buf_cells * sizeof(float), hipMemcpyHostToDevice)); HIP_CHECK(hipMemcpy(gpu_land_mask, land_mask, grid_cells * sizeof(unsigned char), hipMemcpyHostToDevice)); ``` 这里有一个优化细节不能不说：**`hipHostRegister`**。如果CPU端的缓冲区是普通 `malloc`出来的可分页内存，H2D传输时DCU驱动要先把数据拷到内部的一个pin住的内存区域再做DMA，多了一次拷贝。我们用的是 `mmap`分配的共享内存，配上 `hipHostRegister`直接pin住，让DCU的DMA引擎直接读——在我们的环境里，这个操作让H2D带宽从大约6-7 GB/s提升到了13 GB/s左右，几乎翻倍。 ```cpp // Pin mmap缓冲区，让DCU做DMA直接读取 HIP_CHECK(hipHostRegister(full_data, full_buf_cells * sizeof(float), hipHostRegisterDefault)); // ... H2D传输 ... // 传输完成后解除pin HIP_CHECK(hipHostUnregister(full_data)); ``` 注意一点：`hipHostRegister`之后这块内存就不能page out了，所以传输完成之后记得 `hipHostUnregister`，不然会一直占着物理内存。 ### 3.5 第一个DCU Kernel 我们第一个DCU kernel是计算气候态统计的——每个格点从330个样本（30年×11天滑动窗口）中算平均值和P90分位数。写HIP kernel和写CUDA kernel在语法层面差异极小： ```cpp __global__ void temporal_batch_kernel( const float *__restrict__ pool, // 输入数据池 const unsigned char *__restrict__ land_mask, // 陆地掩码 double *__restrict__ clim_out, // 输出：平均值 double *__restrict__ p90_out, // 输出：P90分位数 size_t grid_cells, // 总格点数 int B, // 批次包含的DOY数 int pool_offset, // 池偏移 int batch_total_days, // 批次覆盖的天数 int n_unique, // 唯一网格数 float fill_val, // 填充值 int has_fill // 是否有填充值 ) { const int T = 16; // 每个workgroup处理16个cell size_t c = (size_t)blockIdx.x * T + threadIdx.x; if (c >= grid_cells) return; for (int j = 0; j < B; j++) { size_t out_idx = (size_t)j * grid_cells + c; if (land_mask[c]) { // 陆地格点直接填NaN clim_out[out_idx] = -9999.0; p90_out[out_idx] = -9999.0; } else { clim_out[out_idx] = 0.0; p90_out[out_idx] = 0.0; } } } ``` 上面这个是简化版。完整的kernel还包含了共享内存协同加载和quickselect排序。启动方式跟CUDA完全一致： ```cpp int batch_wg = (grid_cells + 15) / 16; // workgroup数量 temporal_batch_kernel<<>>( gpu_pool, gpu_land_mask, gpu_clim_out + offset, gpu_p90_out + offset, grid_cells, B, batch_start, batch_total_days, n_unique, fill_val, has_fill ); HIP_CHECK(hipGetLastError()); HIP_CHECK(hipDeviceSynchronize()); ``` 这里 `<<< >>>`的语法、`hipDeviceSynchronize()`的语义、`hipGetLastError()`的用法，跟CUDA完全一样。第一次在DCU上看到kernel跑通、出正确结果的时候，我的感受是：**国产加速卡的程序员体验已经能做到跟NVIDIA CUDA近乎无缝切换的程度了**，这比我想象的要成熟得多。 ## 四、第一次移植的性能数据 第一版DCU程序的性能表现： | 阶段 | 耗时 | 说明 | | -------------------- | ---------------- | ------------------------- | | Upfront I/O | ~12s | CPU端一次性读取3060个文件 | | H2D上传 | ~1.1s | 12.7GB数据传到DCU显存 | | DCU计算 | **4.4s** | 92个DOY在DCU上全部算完 | | D2H下载 | ~1.7s | 结果从DCU传回CPU | | NetCDF写出 | ~2s | 92个结果文件写出 | | **Wall Clock** | **~21.5s** | 端到端 | DCU的计算时间只有4.4秒——平均每个DOY才48毫秒。对比我们CPU版本计算阶段约12-14秒，DCU纯计算加速了约3倍。不过因为CPU版本已经把I/O和计算做了深度流水线重叠，所以端到端对比的时候，第一版GPU方案（21.5s）反而不如CPU最终版（9-13s）。这不是DCU的问题，而是**I/O策略还没跟上**——第一版的所有文件读取都在" upfront I/O"阶段串行等完了才给DCU喂数据。 这个发现直接引出了我们后续的优化方向（中篇会详细讲），也让我意识到一个很重要的原则：**加速卡编程，瓶颈往往不在卡上，而在怎么把数据高效地喂给卡**。 ## 五、第一次移植的几点感受 1. **HIP生态的成熟度超出预期**。从CUDA迁移到HIP的代码改动量很小，很多地方直接替换前缀就行。`hipcc`编译器用起来跟 `nvcc`体验接近，编译速度也不慢。 2. **海光DCU跟AMD ROCm的兼容性做得不错**。我们用AMD的ROCm文档来查API基本没问题，DTK工具链的命令行参数也跟ROCm保持一致。这对于习惯了查英文文档的开发者来说很友好。 3. **DTK版本很重要**。我们用的是DTK 25.04.4，不同版本的DTK可能在 `--offload-arch`支持上有差异。如果你用的是更早的版本（比如23.x），可能需要不同的编译参数。 4. **显存是实实在在的限制**。我们的数据规模需要约15GB显存，刚好卡在16GB DCU的极限。如果你的数据更大，可能得做显存优化——分批次处理或者用统一内存（managed memory）。这也是我们中篇要重点讨论的话题。 5. **`hipHostRegister`的效果很明显**。如果你的数据是 `mmap`或者 `malloc`出来的大块连续内存，配上host register让DCU DMA直接访问，H2D带宽可以提升近一倍。 --- 下一篇（中篇）我会展开讲DCU的深度优化——包括时序分块策略怎么把计算效率再翻倍、kernel内部用共享内存做协同加载、以及我们在DCU上踩过的各种坑：quickselect在GPU上的实现细节、显存不足怎么处理、多进程I/O和DCU的配合等等。都是真实踩过的坑，qwq。