从一条命令开始的迁移：HIPify 实战记录

对于习惯了 NVIDIA CUDA 生态的开发者来说，初次面对 AMD ROCm 平台时，最大的心理障碍往往不是硬件本身的差异，而是那种“代码要推倒重来”的焦虑。看着项目里成千上万行写着 cudaMalloc、kernel<<<>>> 的代码，下意识会觉得换个平台就得手动翻译一遍，这工程量想想都头大。

但实际动手后发现，这种顾虑完全是多余的。ROCm 生态中的 HIPify 工具链已经非常成熟，它能把这场看似浩大的“翻译工程”简化成终端里的一行命令。今天就来记录一下我最近把一个小型 CUDA 项目迁移到 ROCm 的真实过程，重点聊聊自动转换后遇到的第一个“拦路虎”，以及如何通过调整 Thrust 库的调用让代码真正跑起来。

一键转换：hipify-clang 的实际表现

迁移的第一步非常简单。假设你有一个名为 my_cuda_proj 的目录，里面包含了所有的 .cu 源文件和头文件。你不需要逐行打开文件去替换 API，只需要在终端执行：

hipify-clang ./my_cuda_proj/src --output-directory=./my_hip_proj

这条命令会启动 hipify-clang，它基于 Clang 编译器前端，能够深入解析 C++ 语法树。它会扫描整个目录，识别标准的 CUDA API（如 cudaMemcpy、cudaFree）以及内核启动语法，并将它们自动映射为对应的 HIP 接口（如 hipMemcpy、hipFree）。

在我的实践中，对于绝大多数标准算子和基础的内存操作，这个工具的准确率极高，基本上完成了 90% 以上的机械性工作。转换完成后，你会在 ./my_hip_proj 目录下看到一份全新的代码副本，文件后缀通常会被标记为 .hip 或者保持 .cpp 但内容已变更。

这时候很多人会觉得“大功告成”了，直接跑去编译。但现实往往是：编译报错。别慌，这正是人工介入的最佳时机。

第一个编译报错：Thrust 库的“水土不服”

在我尝试编译生成的 HIP 代码时，立刻遇到了一个典型的错误。报错信息指向一段使用了 Thrust 库的代码。Thrust 是 CUDA 中常用的并行算法库，类似于 C++ 的 STL，但在 HIP 环境下，它的命名空间和部分行为略有不同。

报错现场：

error: 'thrust::device_vector' is not a member of 'thrust'; did you mean 'hip thrust'?
...
undefined reference to `thrust::sort'

检查源码发现，自动转换后的代码保留了原有的 include 路径和命名空间写法，但在 ROCm 环境下，Thrust 的实现已经被整合进 HIP 运行时，或者需要显式地包含特定的头文件并调整命名空间。

修改前的代码（CUDA 风格）：

#include <thrust/device_vector.h>
#include <thrust/sort.h>

void sort_data(float* d_data, int size) {
    thrust::device_vector<float> vec(d_data, d_data + size);
    thrust::sort(vec.begin(), vec.end());
    // ...
}

修正后的代码（HIP 风格）：

在 ROCm 中，我们需要确保引入了正确的 HIP Thrust 头文件，并且有时需要显式指定执行策略（execution policy），以明确告诉编译器这是在 GPU 上执行。

#include <hip/hip_runtime.h>
// 关键修改：引入 hip/thrust 相关头文件，确保路径正确
#include <thrust/device_vector.h>
#include <thrust/sort.h>
#include <thrust/execution_policy.h>

void sort_data(float* d_data, int size) {
    // 构造 device_vector
    thrust::device_vector<float> vec(d_data, d_data + size);
    
    // 关键修改：显式添加 execution_policy，指定为 device 端
    thrust::sort(thrust::device, vec.begin(), vec.end());
    
    // 如果需要拷回数据，继续使用 hipMemcpy 或 thrust::copy
}

除了添加 thrust::device 执行策略外，还需要检查 CMakeLists.txt 或编译脚本。在 CUDA 项目中，我们习惯链接 cudart，而在 HIP 项目中，必须确保链接的是 hiprtc 和 rocthrust（如果单独使用）。通常 hipcc 编译器会自动处理大部分链接选项，但显式指定 -D__HIP_PLATFORM_AMD__ 宏定义有助于消除一些歧义。

这个插曲让我意识到：HIPify 负责“翻译”语法，但“逻辑校验”必须由人来完成。 特别是涉及第三方库或特定硬件特性的地方，自动化工具很难做到 100% 完美。

验证时刻：跑通 Hello World

修复了 Thrust 库的问题后，再次运行编译命令：

hipcc -o hello_rocm main.hip

这次终于没有报错了！生成的可执行文件 hello_rocm 静静躺在目录下。接下来就是激动人心的验证环节。

在运行之前，先确认一下环境是否识别到了 AMD 显卡：

rocminfo

如果输出了详细的 GPU 架构信息（比如 gfx942 对应 MI300 系列），说明驱动层没问题。然后直接运行程序：

./hello_rocm

终端输出了预期的结果：

Hello from HIP!
Device ID: 0
Device Name: AMD Instinct MI250X
Memory Usage: OK
Computation Result: 42.000000

看到 AMD Instinct 的字样出现在输出里，那一刻的成就感是不言而喻的。这不仅仅是一个简单的打印，它意味着你的代码已经成功脱离了 NVIDIA 的绑定，能够在 AMD 的硬件上独立运行。

写在最后

从运行 hipify-clang 到手动修正 Thrust 调用，再到最终看到输出结果，整个迁移过程比我预想中要顺畅得多。原本以为会是一场漫长的“逐行重写”，实际上核心工作只是集中在那些自动化工具无法覆盖的 10% 的边缘情况上。

对于有 CUDA 基础的开发者来说，跨入 ROCm 生态的门槛其实并没有想象中那么高。HIPify 帮你解决了最繁琐的语法转换，而你只需要把精力放在理解 HIP 的执行模型、调整库的调用差异上。一旦跑通了第一个 Hello World，后续的算子优化、多卡通信等工作就有了坚实的基础。

别再被“迁移成本”吓退了，打开终端，试着对你的项目运行那条命令吧。也许你会发现，新的硬件世界比你想象的更开放、更有趣。

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