工具选型：hipify-perl 还是 hipify-clang？

对于习惯了 NVIDIA CUDA 生态的开发者来说，初次面对 AMD ROCm 平台，最直观的门槛就是代码迁移。别被“迁移”这个词吓到，AMD 官方提供的 HIPify 工具链已经能帮我们完成绝大部分机械性的工作。但在动手之前，得先搞清楚两个核心工具的区别：hipify-perl 和 hipify-clang。

简单来说，hipify-perl 是基于正则表达式的文本替换工具。它的优势在于快、依赖少，不需要完整的编译环境就能跑，适合快速扫描项目看看有多少 CUDA API 需要改。但它的缺点也很明显：不懂语法结构，容易误伤注释里的代码，或者在复杂的宏定义面前束手无策。

相比之下，hipify-clang 则是基于 Clang 编译器的 AST（抽象语法树）进行解析。它能真正“理解”C++ 代码结构，转换准确率极高，尤其是处理模板、宏和复杂类型时更靠谱。如果你的项目代码规范、能正常通过 Clang 编译，强烈建议直接使用 hipify-clang。虽然它配置稍麻烦（需要指定 CUDA include 路径），但能省去后续大量手动修 bug 的时间。

实际操作中，我通常先用 hipify-perl 做个快速预估：

hipify-perl -stat -print-stats src/

这会输出一个统计报告，告诉你哪些文件改了、哪些 API 被替换了。确认无误后，再用 hipify-clang 进行正式转换：

hipify-clang -o=hipified_src/ -cuda-path=/usr/local/cuda src/*.cu

注意，这里 -cuda-path 指向的是你本地安装的 CUDA toolkit 路径，因为 hipify-clang 需要读取 CUDA 的头文件来解析语法。

避坑实录：从自动转换到手动修补

运行完转换工具，千万别以为万事大吉。自动化只能解决 90% 的问题，剩下的 10% 往往是硬骨头。最常见的情况是编译报错，提示找不到符号。这通常是因为某些 CUDA 特有的库函数没有对应的直接映射。

比如，很多深度学习项目依赖 cuBLAS 做矩阵运算。HIPify 会把 cublasSgemm 转换成 hipblasSgemm，但如果你没安装 rocBLAS 或者链接参数没改，编译器直接会报 undefined reference。这时候需要手动检查 CMakeLists.txt 或 Makefile，把链接库从 -lcublas 改为 -lrocblas，并确保 include 路径指向了 ROCm 的头文件目录。

还有一个高频坑点是未转换标记。HIPify 遇到它不认识的非标准 CUDA 写法时，会在代码里留下 // HIPIFY: ... 这样的注释标记，或者直接跳过不转。我有一次迁移一个自定义算子，发现内核启动配置 <<<grid, block>>> 没变，原因是代码里用了非标准的宏包裹。这种时候必须人工介入，逐行审查生成的 .hip 文件。

下面是一个简单的矩阵乘法内核迁移对比，大家感受一下细节差异：

迁移前 (CUDA):

__global__ void matMulKernel(float *A, float *B, float *C, int N) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < N && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < N; ++k) {
            sum += A[row * N + k] * B[k * N + col];
        }
        C[row * N + col] = sum;
    }
}
// 调用
matMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C, N);

迁移后 (HIP):

__global__ void matMulKernel(float *A, float *B, float *C, int N) {
    // 语法基本一致，但需注意内存模型差异
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < N && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < N; ++k) {
            sum += A[row * N + k] * B[k * N + col];
        }
        C[row * N + col] = sum;
    }
}
// 调用 hipLaunchKernelGGL 或直接使用 <<< >>> (HIP 也支持后者，但推荐前者以兼容旧版)
hipLaunchKernelGGL(matMulKernel, dimGrid, dimBlock, 0, 0, d_A, d_B, d_C, N);

看起来变化不大？重点在于背后的运行时库变了。如果编译时报错说 hipLaunchKernelGGL 未定义，记得检查是否包含了 <hip/hip_runtime.h>。另外，有些老版本的 HIPify 可能不会自动添加这个头文件，需要手动补上。

环境隔离：用 Conda 跑通第一个 HIP 程序

很多开发者在迁移初期崩溃，不是因为代码难改，而是环境乱套了。系统里同时装着 CUDA 和 ROCm 的驱动、库文件，Python 包管理工具又默认去抓 CUDA 版本的 PyTorch，结果就是各种 ImportError 和 Segmentation Fault。

我的血泪经验是：务必使用 Conda 进行环境隔离。不要试图在系统全局环境下混用两套生态。创建一个干净的虚拟环境，专门用于 ROCm 开发：

conda create -n rocm-dev python=3.9
conda activate rocm-dev

接下来安装 PyTorch 时，一定要去 PyTorch 官网查找明确标注 “ROCm” 的安装命令，千万别手滑装了 CUDA 版本。例如：

pip3 install torch torchvision torchaudio --index-url https://download.pytorch.org/whl/rocm6.0

（具体版本号请根据你安装的 ROCm 版本调整）

安装完成后，写个最简单的 Hello World 验证一下。不需要复杂的矩阵运算，只要能调用 GPU 就算成功：

#include <hip/hip_runtime.h>
#include <iostream>

__global__ void helloFromGPU() {
    printf("Hello from GPU thread %d!\n", threadIdx.x);
}

int main() {
    helloFromGPU<<<1, 5>>>();
    hipDeviceSynchronize();
    std::cout << "Check dmesg or stdout for GPU output." << std::endl;
    return 0;
}

编译命令也要换成 hipcc：

hipcc -o hello_gpu hello_gpu.cpp
./hello_gpu

如果能看到线程打印信息，恭喜你，第一个 HIP 程序跑通了！这一步看似简单，实则验证了驱动、编译器、运行时库全链路是否正常。只有在这个基础上，后续复杂的算子迁移和框架适配才有意义。

迁移之路确实有些琐碎，但只要把工具选对、环境隔离好、耐心处理那 10% 的手动修改，你会发现 ROCm 并没有想象中那么遥远。当你看到代码在 AMD 显卡上流畅运行时，那种成就感绝对值得这番折腾。

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