海光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 <cuda_runtime.h>，得用#include <hip/hip_runtime.h>
• 你不能cudaMalloc，得用hipMalloc
• 你不能<<<grid, block>>>启动CUDA kernel，但你可以用完全一样的<<<grid, block>>>语法启动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

环境模块加载很简单，就一行：

module load rocm/dtk-25.04.4

SLURM脚本里申请DCU资源：

#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可用，如果任务需要多卡并行，可以把数字调大。

编译命令长这样：

/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之后第一件事，肯定是先看看手里这块卡什么配置：

#include <hip/hip_runtime.h>

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一样，写个宏省事：

#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)

使用方法：

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。

// 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）：

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左右，几乎翻倍。

// 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在语法层面差异极小：

__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完全一致：

int batch_wg = (grid_cells + 15) / 16;  // workgroup数量
temporal_batch_kernel<<<batch_wg, 16>>>(
    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。