从 CUDA 到 ROCm：我的第一次 HIPify 迁移实录

作为一个长期在 NVIDIA 生态里“打滚”的开发者，第一次把手伸向 AMD GPU 时，心里其实是有点打鼓的。习惯了 nvcc 的报错信息和 CUDA 的工具链，突然要面对 ROCm 这套新体系，最大的障碍往往不是算法本身，而是那成千上万行耦合了 CUDA API 的代码。最近为了验证一块新到的 AMD 显卡性能，我硬着头皮做了一次从 CUDA 到 HIP 的代码迁移。整个过程并非像文档里写的那样“一键完成”，中间踩了不少坑，但也总结出了一套相对稳妥的“手摸手”实操路径。今天就把这次真实的迁移经历记录下来，重点聊聊自动化工具的边界和那些必须人工介入的“硬骨头”。

自动化利器：hipify-perl 与 hipify-clang 的实战选择

迁移的第一步，肯定是能偷懒则偷懒。AMD 官方提供的 HIPify 工具链就是为此而生的，它能把大部分标准的 CUDA 调用自动映射为 HIP 接口。但在实际使用前，你得在 hipify-perl 和 hipify-clang 之间做个选择。

如果你手头的项目比较老旧，或者包含大量非标准的宏定义，hipify-perl 这种基于正则表达式替换的工具会更“宽容”一些。它的用法非常简单，直接指定源目录和目标目录即可：

hipify-perl -o ./hip_src ./cuda_src

这条命令会扫描 ./cuda_src 下的所有文件，将 cudaMalloc 变成 hipMalloc，cudaMemcpy 变成 hipMemcpy，甚至连 __global__ 这样的关键字也会保留（虽然在 HIP 中语义略有不同，但通常兼容）。对于简单的算子库，这一步能解决 90% 的工作。

但对于现代 C++ 项目，尤其是用到了模板元编程或复杂头文件依赖的工程，我更推荐 hipify-clang。它基于 Clang 编译器前端，能理解代码的语法树，转换准确率更高，不容易出现“误伤”。不过它的门槛也高，需要你能提供正确的编译参数（include 路径等）：

hipify-clang -p ./compile_commands.json ./cuda_src/main.cu --output-dir=./hip_src

在实际操作中，我发现 hipify-clang 对项目中引用的第三方 CUDA 库头文件非常敏感。如果编译数据库（compile_commands.json）没配置好，它会直接报错退出，而不是像 perl 版本那样强行替换。虽然前期配置麻烦点，但为了后续少改几处逻辑错误，这个时间花得值。

当自动化失效：一次典型的 cuBLAS 报错修复

自动化工具绝不是银弹。在我运行完转换脚本并尝试编译时， linker 阶段直接抛出了红色报错：

undefined reference to `cublasSgemm'
collect2: error: ld returned 1 exit status

这就是典型的“漏网之鱼”。HIPify 能够识别基础的内存操作和内核启动，但对于像 cuBLAS 这样的高级库，它往往只是机械地把 cublas 命名空间保留下来，却忘了你需要链接的是 AMD 对应的 rocBLAS 库，甚至某些特定的 API 在 rocBLAS 中签名都变了。

打开生成的 .hip 文件，我看到这样一段代码：

// 自动转换后的代码，依然调用了 cublas 句柄
cublasHandle_t handle;
cublasCreate(&handle);
cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, m, n, k, &alpha, d_A, lda, d_B, ldb, &beta, d_C, ldc);

这里有两个问题：一是头文件引用还是 <cublas_v2.h>，二是 API 调用前缀没变。手动修复的过程其实很标准化：

1. 替换头文件：将 #include <cublas_v2.h> 改为 #include <hipblas/hipblas.h>。
2. 重命名句柄与 API：cublasHandle_t 变为 hipblasHandle_t，cublasCreate 变为 hipblasCreate。
3. 调整枚举值：CUBLAS_OP_N 需要改为 HIPBLAS_OP_N。

修改后的代码长这样：

#include <hipblas/hipblas.h>

hipblasHandle_t handle;
hipblasCreate(&handle);
hipblasSgemm(handle, HIPBLAS_OP_N, HIPBLAS_OP_N, m, n, k, &alpha, d_A, lda, d_B, ldb, &beta, d_C, ldc);

除了 cuBLAS，像 cuFFT、cuRAND 等库也都有类似的对应关系（rocFFT、rocRAND）。我的经验是：不要指望一次转换成功。跑一遍 hipify，立刻尝试编译，根据报错信息定位残留的 CUDA 专有符号，然后批量替换。这种“转换 - 编译 - 修复”的小步快跑策略，比一次性改完再调试要高效得多。

最小化验证：Hello World 内核的对比启示

为了确认环境彻底打通，我写了一个最简单的向量加法内核作为"Hello World"。对比转换前后的代码，能很直观地看到 HIP 的兼容性设计哲学。

原始的 CUDA 版本：

__global__ void vecAdd(float *A, float *B, float *C, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) C[i] = A[i] + B[i];
}
// 主机端调用
vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, n);

经过 HIPify 处理后，代码几乎一模一样：

__global__ void vecAdd(float *A, float *B, float *C, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) C[i] = A[i] + B[i];
}
// 主机端调用 (注意：hipLaunchKernelGGL 宏在某些旧版本可能需要，但新版本大多兼容 <<< >>>)
vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, n);

你会发现，核心的内核逻辑（Kernel Logic）完全不需要动，连线程索引的计算方式都保持一致。这给了我很大的信心：大部分计算逻辑确实是可移植的。但要注意，如果在代码里硬编码了特定于 NVIDIA 架构的 Warp Size（比如写死 32），在 AMD 平台上虽然目前也是 64 线程为一组 Wavefront，但为了代码的纯粹性，建议改用 warpSize 变量或宏。

这次实践让我深刻体会到，迁移工作最忌讳“盲目全量替换”。很多人喜欢一口气跑完 hipify 然后指望直接运行，结果面对几百个编译错误无从下手。正确的姿势应该是：先拿一个最小的可编译单元（比如一个独立的算子文件）进行转换和编译验证，确保工具链参数配置无误，再逐步扩大到整个项目。

从 CUDA 到 ROCm，本质上是一次从“舒适区”走向“异构通用”的必经之路。工具已经足够好用，剩下的就是耐心和细致的排查。当你看到第一个内核在 AMD 显卡上正确跑出结果时，那种打破厂商锁定的成就感，绝对值得你花这几个小时去折腾。

200小时GPU算力已就位，快来领取：https://marketing.csdn.net/questions/Q2604140858304426315?utm_source=AIpaper