说到用DCU写程序，很多人的第一反应可能是：“这玩意儿跟CUDA差多少？代码要不要重写？”答案是：**比你想象的小得多**。 要理解为什么，关键就在于**HIP Runtime API**。它是DCU上最核心的编程接口，也是理解DCU软件生态的钥匙。这篇文章不谈虚的，直接以《DTK 26.04 HIP Runtime API开发手册》为蓝本，带你看看这套API到底长什么样、怎么用，以及它背后的设计逻辑。 --- ## 一、HIP到底是什么？一句话讲清楚 HIP的全称是**Heterogeneous-compute Interface for Portability**（异构计算可移植接口）。你可以把它理解为：**一套长得像CUDA、但能跑在AMD/DCU上的编程接口**。 HIP的Runtime API提供了一套C++风格的接口，涵盖了你写GPU程序需要的所有能力：设备管理、内存分配、Kernel启动、流和事件、图执行……基本你在CUDA里见过的，HIP里都有对应的版本。 区别在哪？**CUDA分了两套API**：Driver API（`cu`开头，细粒度控制）和Runtime API（`cuda`开头，自动管理）。**HIP把两者统一成了一套Runtime API**，同时保留了`hipModule`和`hipCtx`这类底层接口。换句话说，你写一套代码，就能覆盖CUDA Driver和Runtime两层的能力。 ## 二、从“Hello DCU”开始：最简代码长什么样？ 先看一个最简的DCU程序骨架： ```cpp #include // 注意：不是 cuda_runtime.h int main() { int deviceCount; hipGetDeviceCount(&deviceCount); // 获取DCU数量【手册1.22】 int deviceId; hipGetDevice(&deviceId); // 获取当前设备【手册1.21】 hipSetDevice(0); // 切换到0号设备【手册1.30】 float *d_data; hipMalloc(&d_data, 1024 * sizeof(float)); // 在DCU上分配内存【手册10.17】 // ... 启动Kernel、拷贝数据 ... hipFree(d_data); // 释放DCU内存【手册10.5】 return 0; } ``` **注意到什么了吗？** 除了头文件和函数名前缀从`cuda`换成`hip`，几乎一模一样。这就是HIP设计的核心哲学：**用最少的改动，让CUDA代码跑在DCU上**。 ## 三、API设计的三层逻辑 《DTK 26.04 HIP Runtime API开发手册》里列出了280多个API接口，看似庞大，但梳理清楚背后的设计逻辑，你会发现它们其实遵循着清晰的层次： ### 第一层：设备管理——告诉DCU“用哪块卡” 多卡场景下，首先要管理设备。HIP提供了一组简洁的设备管理API【手册第1章】： - `hipGetDeviceCount()`：系统里有多少块DCU【手册1.22】 - `hipGetDevice()` / `hipSetDevice()`：当前线程用哪块卡【手册1.21】【手册1.30】 - `hipGetDeviceProperties()`：查这块卡的显存大小、计算单元数量等详细信息【手册1.24】 - `hipDeviceSynchronize()`：让CPU等待当前设备上所有任务完成【手册1.20】 这些API的逻辑非常直观——**先查卡、再选卡、最后用卡**。多卡之间还能通过`hipDeviceCanAccessPeer()`检测是否支持直接访问对方显存【手册14.1】，以及`hipDeviceEnablePeerAccess()`开启P2P访问【手册14.3】——这些在分布式训练和多卡科学计算中非常关键。 ### 第二层：内存管理——数据怎么搬上DCU？ GPU编程最大的不同在于**显存是独立的**。数据得先从CPU内存搬到DCU显存，计算完再搬回来。HIP的内存管理API几乎1:1复刻了CUDA的命名【手册第10章】： **分配与释放：** - `hipMalloc()`：在DCU上分配显存【手册10.17】 - `hipFree()`：释放显存【手册10.5】 - `hipHostMalloc()`：分配**锁页内存**（pinned memory），这种内存可以被DCU直接访问，异步传输必备【手册10.21】 - `hipHostFree()`：释放锁页内存【手册10.7】 **数据搬运：** - `hipMemcpy()`：同步拷贝，CPU会等拷贝完成再继续【手册10.30】 - `hipMemcpyAsync()`：异步拷贝，需要配合流（Stream）使用【手册10.42】 - `hipMemcpyPeer()`：DCU之间的直接拷贝，不需要经过CPU内存【手册10.45】 **高级特性：** - `hipMallocManaged()`：统一内存，CPU和DCU共享同一块地址空间，系统自动管理数据迁移【手册10.22】 - `hipMemAdvise()`：给内存管理提建议（比如“这块数据主要被哪个设备访问”）【手册10.25】 - `hipMemPrefetchAsync()`：主动把数据预取到指定设备【手册10.27】 注意`hipFree()`有一个容易被忽略的行为：**它会隐式调用`hipDeviceSynchronize()`**。这意味着调用`hipFree()`时CPU会等待DCU上所有任务完成，如果不小心在性能敏感路径上频繁分配释放内存，可能会成为性能瓶颈。 ### 第三层：执行控制——Kernel怎么跑起来？ 在HIP里启动一个Kernel（在DCU上运行的函数），最常用的方式是`hipLaunchKernel`宏【手册7.10】： ```cpp hipLaunchKernel(myKernel, // 要运行的函数 dim3(blocks), // 网格大小（多少个Block） dim3(blockSize), // 每个Block里多少个线程 0, // 动态共享内存大小 0, // Stream（0表示默认流） arg1, arg2, ...); // 函数参数 ``` 如果你需要更精细的控制，还有`hipLaunchKernelExC`，支持传入额外的启动属性（比如`hipLaunchAttributeCooperative`标记协作式启动）【手册7.13】。 Kernel执行前后的同步靠两类API： - `hipDeviceSynchronize()`：等待当前设备上**所有**任务完成【手册1.20】 - `hipStreamSynchronize()`：等待**指定Stream**里的任务完成【手册4.19】 ## 四、进阶能力：流、事件与图 如果只掌握了上面三层，你已经能写DCU程序了。但要写出**高性能**的DCU程序，必须理解这三个概念。 ### 流（Stream）：让任务并行起来 Stream可以理解为**DCU上的任务队列**。不同Stream里的任务可以**并行执行**，同一Stream里的任务**按顺序执行**。 ```cpp hipStream_t stream1, stream2; hipStreamCreate(&stream1); // 创建两个流【手册4.6】 hipStreamCreate(&stream2); // 在两个流里同时执行不同的Kernel，它们可以在DCU上并行 hipLaunchKernel(kernelA, ..., stream1); hipLaunchKernel(kernelB, ..., stream2); hipStreamSynchronize(stream1); // 等待stream1完成【手册4.19】 hipStreamSynchronize(stream2); hipStreamDestroy(stream1); // 销毁流【手册4.9】 ``` Stream还可以设置优先级（`hipStreamCreateWithPriority`）【手册4.8】，甚至绑定到指定的计算单元（`hipExtStreamCreateWithCUMask`）【手册26.3】——这种细粒度控制在科学计算中非常有用。 ### 事件（Event）：测时间、做同步 Event主要有两个用途：**测时间**和**做同步**。 测时间： ```cpp hipEvent_t start, stop; hipEventCreate(&start); // 创建事件【手册5.1】 hipEventCreate(&stop); hipEventRecord(start, stream); // 在stream里记录开始【手册5.6】 // ... 执行Kernel ... hipEventRecord(stop, stream); hipEventSynchronize(stop); // 等待事件完成【手册5.8】 float ms; hipEventElapsedTime(&ms, start, stop); // 算时间差（毫秒）【手册5.4】 ``` 做同步：让一个Stream等待另一个Stream里的某个事件完成，实现Stream间的依赖关系。 ```cpp hipStreamWaitEvent(stream2, event1, 0); // stream2等待event1完成【手册4.21】 ``` ### 图（Graph）：把多个操作打包执行 Graph是HIP里相对较新的特性【手册第20章】，它允许你把多个Kernel、内存拷贝、内存设置操作**预先定义成一个有依赖关系的图**，然后一次性提交执行。好处是**减少CPU启动开销**，特别适合重复执行相同计算模式的场景。 基本用法： ```cpp hipGraph_t graph; hipGraphCreate(&graph, 0); // 创建图【手册20.22】 hipGraphNode_t kernelNode; hipKernelNodeParams params = {...}; hipGraphAddKernelNode(&kernelNode, graph, // 添加Kernel节点【手册20.12】 nullptr, 0, ¶ms); hipGraphExec_t exec; hipGraphInstantiate(&exec, graph, // 实例化图【手册20.53】 nullptr, nullptr, 0); hipGraphLaunch(exec, stream); // 启动图【手册20.60】 hipGraphExecDestroy(exec); // 销毁【手册20.31】 hipGraphDestroy(graph); // 销毁【手册20.24】 ``` DTK 25.04.1还新增了`hipGraphAddNode`通用接口【手册20.91】和`hipGraphInstantiateWithParams`【手册20.94】，让图构建更加灵活。 ## 五、DCU专属的“定制化API” 标准的HIP API已经很强大了，但海光DCU还提供了一套**定制化API**（`du`前缀），让你能触达更底层的硬件能力【手册第28章】： - `duDeviceGetHwQueueNum()`：查询当前DCU有多少个硬件队列【手册28.1】 - `duStreamGetHwQueue()`：查询某个Stream绑定到了哪个硬件队列【手册28.2】 - `duExtStreamCreateWithCUMask()`：创建Stream时指定使用哪些计算单元（CU）【手册26.3】 - `duDeviceFlushCache()`：主动刷新DCU缓存【手册28.4】 - `duKernelGetAql()` / `duKernelAqlLaunch()`：更底层的AQL包管理，可以实现批量Kernel启动【手册28.9】【手册28.10】 - `duGraphGetNodeByLabel()`：通过标签在图中查找节点【手册28.13】 这些API提供了**比标准HIP更细粒度的控制能力**，适合对性能有极致要求的场景——比如需要在多个DCU之间精细调配计算资源的科学计算任务。 ## 六、错误处理：别忽视返回值 HIP API的返回值都是`hipError_t`类型。**每个API调用都应该检查返回值**，这是写出健壮DCU程序的第一步。 ```cpp hipError_t err = hipMalloc(&d_data, size); if (err != hipSuccess) { printf("hipMalloc failed: %s\n", hipGetErrorString(err)); // 获取错误描述【手册3.2】 return -1; } ``` 常用的错误处理辅助函数： - `hipGetErrorName(err)`：返回错误码的名字，如`hipErrorOutOfMemory`【手册3.1】 - `hipGetErrorString(err)`：返回人类可读的错误描述【手册3.2】 - `hipGetLastError()`：返回上一个API调用的错误，并自动重置【手册3.3】 - `hipPeekAtLastError()`：同上，但不重置【手册3.4】 ## 七、从CUDA到HIP：迁移有多简单？ 说了这么多，回到最开始的问题：**从CUDA迁移到DCU到底有多麻烦？** 答案是：**比你想象的小得多**。 CUDA和HIP的API在函数名、参数、语义上都高度一致。大部分情况下，你只需要： 1. 把头文件从`cuda_runtime.h`换成`hip/hip_runtime.h` 2. 把所有`cuda`前缀换成`hip` 3. 编译器从`nvcc`换成`hipcc` AMD官方提供了**HIPIFY工具**，可以自动完成大部分CUDA代码到HIP代码的转换。海光的光合开发者社区也提供了“CUDA转码”的AI技能支持。 当然，深度优化时还是会遇到差异（比如warp size、内存层次结构等），但对于**绝大多数应用层代码，迁移成本极低**。 ## 结语 DTK 26.04的HIP Runtime API已经覆盖了从设备管理、内存管理、流和事件到图执行的完整编程模型。280多个API接口，每一行都在回答同一个问题：**怎么让开发者用最熟悉的方式，把计算任务搬到DCU上**。 对于写过CUDA的开发者来说，HIP的学习曲线几乎是一条平线——换几个前缀、改几个头文件，代码就能跑起来。对于没写过GPU程序的开发者来说，HIP的API设计直观、文档齐全（光合开发者社区提供了完整的API参考手册），入门门槛也远低于预期。 国产加速卡能不能用、好不好用，**API是第一步，也是最重要的一步**。从DTK 26.04的HIP Runtime API来看，这一步走得扎实。