从 CUDA 到 HIP：一次真实的“一键迁移”实验

手里攥着一堆跑得很溜的 CUDA 代码，面对新入手的 AMD Instinct 显卡，心里多少有点打鼓：“难道要逐行重写？”这确实是很多开发者转向 ROCm 平台时的第一道心理门槛。但最近我亲自做了一次迁移实验，发现情况比预想中乐观得多。只要工具链用对，90% 的工作其实是可以自动完成的。今天就把这次从 hipify-clang 自动化转换到人工修补“硬骨头”的全过程记录下来，希望能给想尝试 AMD 显卡的朋友一点参考。

自动化转换：90% 的工作交给脚本

迁移的第一步，我并没有急着打开编辑器，而是直接祭出了 AMD 官方提供的 hipify-clang 工具。它的逻辑很直观：扫描源码，识别 CUDA 特有的 API 和语法糖，然后批量替换成 HIP 接口。

假设我的项目源码在 ./src 目录下，执行命令非常简单：

hipify-clang ./src --output-directory=./hip_src

运行结束后，./hip_src 里生成了一份全新的代码副本。我随手 diff 了一下修改前后的文件，变化确实令人舒适。原本熟悉的 cudaMalloc 变成了 hipMalloc，cudaMemcpy 变成了 hipMemcpy，就连内核启动语法 kernel<<<grid, block>>> 也被自动转换成了 HIP 兼容的形式。对于标准的算子实现，这种机械性的替换准确率极高，基本上扫一眼就能确认大部分文件已经“翻译”完毕。

这一刻你会觉得迁移不过如此，甚至想直接编译试试。但别急，真正的挑战往往藏在那剩下的 10% 里。

那 10% 的“硬骨头”：Thrust 与内联汇编

当我满怀信心地对着转换后的代码执行 make 时，编译器毫不留情地抛出了一堆错误。这就是自动化工具的边界所在：它能处理标准 API，但遇到复杂的模板特化、第三方库依赖或底层汇编时，就得靠人来填坑了。

第一个拦路虎是 Thrust 库。
原项目中大量使用了 thrust::device_vector 和一些特定的排序算法。虽然 HIP 也支持 Thrust，但在头文件包含路径和命名空间细节上存在差异。编译报错显示 undefined reference 或者模板参数不匹配。

解决思路并不复杂，但需要细心。我需要手动检查所有引用 Thrust 的地方，将 #include <thrust/...> 确认是否指向了 ROCm 版本的头文件，有时还需要显式指定 hip::thrust 命名空间。比如在某个自定义的排序算子中，原代码直接调用了 CUDA 特有的扩展接口，HIPify 没能识别出来。我不得不手动将其重构为标准的 HIP Thrust 调用，或者改用 rocPRIM 库中的对应实现。

更棘手的是内联汇编（Inline Assembly）。
项目中为了极致优化，有一段矩阵乘法的核心逻辑嵌入了 PTX 汇编。hipify-clang 对这部分完全无能为力，直接原样保留，结果当然是编译失败——AMD 显卡不认识 NVIDIA 的 PTX 指令。

这时候只能“硬改”。我得把那段 PTX 汇编翻译成 AMD 对应的 GCN 或 CDNA 汇编指令。这步工作量不小，需要对两种架构的寄存器模型和指令集有一定了解。如果不想深究汇编，另一个折中方案是直接用 HIP C++ intrinsic 函数重写这段逻辑，虽然可能损失一点点极限性能，但可维护性和兼容性会好很多。在这次实验中，我选择用 __builtin_amdgcn_... 系列内置函数替换了原有的 PTX 块，顺利通过了编译。

编译报错排查实录

在修补过程中，像下面这样的报错日志几乎天天见：

error: use of undeclared identifier 'cudaStream_t'; did you mean 'hipStream_t'?
error: no member named 'cudaDeviceSynchronize' in scope

看到这类错误，千万别慌着去搜通用答案。我的经验是：先看上下文。很多时候是因为某些头文件没有被正确包含，或者宏定义没有生效。例如，有些代码通过 #ifdef __CUDA_ARCH__ 来做设备端编译判断，在 HIP 环境下这个宏可能行为不同，需要改为 #ifdef __HIP_DEVICE_COMPILE__。

还有一个常见的坑是路径问题。编译系统可能还在链接旧版的 CUDA 库。务必检查 CMakeLists.txt 或 Makefile，确保 ROCM_PATH 环境变量已正确设置，并且链接器指向的是 rocblas 而不是 cublas。我在一次构建中就因为缓存问题，链接器找到了系统里的 CUDA 库，导致运行时直接 Segmentation Fault。清理构建目录（rm -rf build）并重新配置，往往能解决这类玄学问题。

迁移完成检查清单

当代码终于编译通过，并能成功运行 Hello World 级别的测试后，不要以为万事大吉。为了确保迁移彻底，避免在生产环境运行时崩溃，我整理了一份自查清单，建议大家在提交代码前逐项核对：

• API 替换完整性：全局搜索 cuda 关键字（排除注释），确认没有漏网的 API 调用，特别是那些不起眼的 cudaGetLastError 或事件管理函数。
• 内存管理逻辑：检查 malloc/free 与 hipMalloc/hipFree 是否混用，确保显存分配和释放成对出现，无内存泄漏。
• 流与同步：确认所有的 Stream 操作已转换为 hipStream_t，且同步逻辑（如 hipDeviceSynchronize）在关键路径上生效。
• 第三方库依赖：核实所有依赖库（如 FFT、BLAS）是否都已切换到 ROCm 版本（rocFFT, rocBLAS），并确认版本号兼容。
• 硬编码路径：检查代码中是否有写死的 /usr/local/cuda 路径，全部替换为环境变量或相对路径。
• 多卡通信：如果涉及多卡，确认 NCCL 已替换为 RCCL，并测试过 all_reduce 等通信原语在不同卡间的连通性。

这次迁移实验让我意识到，从 CUDA 到 HIP 并非不可逾越的鸿沟。工具链已经帮我们铺好了 90% 的路，剩下的 10% 虽然需要一些耐心和功底，但每解决一个报错，对异构计算的理解就加深一层。看着代码在 AMD 显卡上流畅跑起来，那种成就感，或许就是折腾技术的乐趣所在吧。

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