从 CUDA 到 ROCm：用 HIPify 跑通第一个向量加法

很多习惯了 NVIDIA CUDA 生态的开发者，在面对 AMD GPU 时往往望而却步。其实，迁移的第一步并没有想象中那么复杂。AMD 官方提供的 hipify 工具链就是专门为了解决这个问题而生的，它能帮你完成大部分机械性的代码替换工作。今天我们就抛开那些宏大的架构叙事，只聚焦最基础的实操：如何安装工具、扫描代码、修复报错，并最终在 AMD 显卡上跑通一个简单的向量加法程序。

安装与扫描：让工具替你干脏活

如果你已经配置好了 ROCm 开发环境，那么获取 hipify-perl 通常非常简单。在大多数基于 Debian/Ubuntu 的系统上，它包含在 rocmi-hipify-tools 包中。安装完成后，我们不需要手动去逐个文件查找 cudaMalloc 或 __global__ 关键字，只需一条命令即可启动批量转换。

假设你有一个名为 vector_add 的旧项目目录，里面全是标准的 CUDA C++ 代码。进入该目录后，执行以下命令：

hipify-perl vector_add.cu

这条命令会原地修改 vector_add.cu 文件（建议先 git commit 备份），将所有的 CUDA API 调用映射为 HIP 接口。例如，cudaMalloc 会变成 hipMalloc，cudaMemcpy 变为 hipMemcpy，而内核启动语法 <<< >>> 也会被保留，因为 HIP 兼容这种写法。对于大多数标准算子，这种自动化转换的准确率极高，基本能搞定 90% 以上的内容。

转换完成后，你会看到源文件中的头文件引用发生了变化。原本熟悉的 #include <cuda_runtime.h> 被替换成了 #include <hip/hip_runtime.h>。这时候不要急着高兴，直接编译通常会失败，因为真正的挑战才刚刚开始。

典型报错与手动修正：头文件与路径陷阱

当你尝试使用 hipcc 编译器构建项目时，很可能会遇到类似以下的报错：

fatal error: 'cuda_runtime.h' file not found
#include <cuda_runtime.h>
         ^~~~~~~~~~~~~~~~

或者更隐蔽的链接错误，提示找不到 cudart 库。这通常是因为 hipify-perl 在某些复杂的 include 逻辑或注释中漏掉了替换，或者是你的项目中硬编码了 CUDA 的路径。

案例一：遗漏的头文件
检查你的代码，有时候工具会忽略宏定义中的 CUDA 引用。你需要手动全局搜索 .cu 和 .h 文件，确保没有任何残留的 cuda_ 前缀头文件。正确的做法是统一改为：

#include <hip/hip_runtime.h>
// 如果用了 math 函数，确保引入 hip 版本
#include <hip/hip_math.h> 

案例二：API 参数差异
虽然大部分 API 是一对一的，但少数函数在 ROCm 下的行为略有不同。比如在处理流（Stream）相关操作时，某些旧版 CUDA 代码可能隐式依赖了默认流的特定行为。如果在编译后运行出现 Segmentation Fault，请检查 hipStreamCreate 和 hipStreamSynchronize 的返回值。

一个常见的修正片段如下，用于显式检查错误码，这在调试初期非常有用：

#define HIP_CHECK(error) \
    if (error != hipSuccess) { \
        std::cerr << "HIP error at " << __FILE__ << ":" << __LINE__ << " - " << hipGetErrorString(error) << std::endl; \
        exit(EXIT_FAILURE); \
    }

// 使用时
hipError_t err = hipMalloc(&d_data, size);
HIP_CHECK(err);

加上这段宏定义，能让你在遇到“硬骨头”时快速定位是哪一行 HIP 调用出了问题，而不是对着黑屏猜谜。

实战验证：跑通向量加法 Hello World

理论再多不如跑个 Demo。下面是一个经过 HIPify 处理后的简化版向量加法代码，你可以直接保存为 vector_add_hip.cpp。它完成了从内存分配到内核启动的全过程。

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

const int SIZE = 1024;

__global__ void vectorAdd(const float *A, const float *B, float *C, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        C[i] = A[i] + B[i];
    }
}

int main() {
    std::vector<float> h_A(SIZE), h_B(SIZE), h_C(SIZE);
    for (int i = 0; i < SIZE; i++) {
        h_A[i] = 1.0f;
        h_B[i] = 2.0f;
    }

    float *d_A, *d_B, *d_C;
    // 分配显存
    hipMalloc(&d_A, SIZE * sizeof(float));
    hipMalloc(&d_B, SIZE * sizeof(float));
    hipMalloc(&d_C, SIZE * sizeof(float));

    // 拷贝数据到设备
    hipMemcpy(d_A, h_A.data(), SIZE * sizeof(float), hipMemcpyHostToDevice);
    hipMemcpy(d_B, h_B.data(), SIZE * sizeof(float), hipMemcpyHostToDevice);

    // 启动内核
    int threadsPerBlock = 256;
    int blocksPerGrid = (SIZE + threadsPerBlock - 1) / threadsPerBlock;
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, SIZE);

    // 拷贝结果回主机
    hipMemcpy(h_C.data(), d_C, SIZE * sizeof(float), hipMemcpyDeviceToHost);

    // 验证结果
    bool success = true;
    for (int i = 0; i < SIZE; i++) {
        if (h_C[i] != 3.0f) {
            success = false;
            break;
        }
    }

    std::cout << (success ? "Success: Vector addition works on ROCm!" : "Failed") << std::endl;

    // 释放资源
    hipFree(d_A);
    hipFree(d_B);
    hipFree(d_C);

    return 0;
}

编译这段代码非常简单，只需调用 ROCm 自带的编译器包装器：

hipcc vector_add_hip.cpp -o vector_add_hip

如果没有报错，直接运行 ./vector_add_hip。如果终端输出了 Success: Vector addition works on ROCm!，恭喜你，你已经成功迈出了从 CUDA 到 ROCm 的第一步。

这个过程看似简单，却是后续迁移大型项目的基础。一旦你熟悉了 hipify 的工作流和常见的修补套路，面对更复杂的深度学习框架迁移时，心里也就有了底。接下来，你可以尝试用同样的方法扫描更大的项目目录，逐步解决那些特定的算子适配问题。

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