手里有张 AMD 显卡，想跑深度学习或者高性能计算，却面对一堆现成的 CUDA 代码发愁？这大概是很多从 N 卡转 A 卡的开发者最先遇到的痛点。其实，AMD 早就铺好了路，核心工具就是 HIPify。它不是让你重写代码，而是通过自动化脚本把 CUDA 语法“翻译”成 HIP（Heterogeneous-compute Interface for Portability），让代码能在 ROCm 平台上编译运行。

今天不聊虚的生态对比，直接上手实操。我们从一个最简单的向量加法 kernel 开始，完整走一遍 hipify-perl 转换、手动修正差异、再到 hipcc 编译运行的全流程。哪怕你之前只写过 CUDA，跟着步骤也能在 Linux 下让代码在 A 卡上跑起来。

准备环境与原始 CUDA 代码

开始前，确保你的环境是干净的 Linux 系统（ROCm 对 Windows 支持有限，生产环境建议 Ubuntu 20.04/22.04），并且已经安装了 ROCm 驱动和工具链。可以通过 rocm-smi 命令查看显卡状态，如果能看到 GPU 列表，说明底层驱动没问题。

我们需要一个典型的 CUDA 源文件作为靶子。创建一个名为 vector_add.cu 的文件，内容是最基础的并行向量加法：

#include <stdio.h>
#include <cuda_runtime.h>

__global__ void vectorAdd(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];
    }
}

int main() {
    int n = 1000000;
    size_t size = n * sizeof(float);

    float *h_a, *h_b, *h_c;
    h_a = (float*)malloc(size);
    h_b = (float*)malloc(size);
    h_c = (float*)malloc(size);

    for (int i = 0; i < n; i++) {
        h_a[i] = 1.0f;
        h_b[i] = 2.0f;
    }

    float *d_a, *d_b, *d_c;
    cudaMalloc((void**)&d_a, size);
    cudaMalloc((void**)&d_b, size);
    cudaMalloc((void**)&d_c, size);

    cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);

    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;

    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n);

    cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);

    printf("Result check: %f\n", h_c[0]);

    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);

    return 0;
}

这段代码在 NVIDIA 环境下用 nvcc 就能编过，但在 AMD 平台上，我们需要把它变成 HIP 代码。

使用 hipify-perl 自动转换

ROCm 工具链里自带了 hipify 系列工具，其中 hipify-perl 是基于正则表达式的文本替换工具，适合处理这种结构清晰的 C++ 代码。它会把 cuda 关键字批量替换为 hip。

在终端执行以下命令：

hipify-perl vector_add.cu > vector_add_hip.cpp

注意这里输出后缀我改成了 .cpp，因为转换后的代码本质上是带有 HIP API 调用的 C++ 文件，不再需要 NVCC 特有的 .cu 扩展名（虽然 .hip 也可以，但 .cpp 配合 hipcc 更通用）。

打开生成的 vector_add_hip.cpp，你会发现大部分内容已经变了：

• #include <cuda_runtime.h> 变成了 #include <hip/hip_runtime.h>
• cudaMalloc 变成了 hipMalloc
• cudaMemcpy 变成了 hipMemcpy
• cudaFree 变成了 hipFree
• 启动内核的 <<< >>> 语法也被保留了下来，这是 HIP 兼容 CUDA 语法的特性之一。

看起来好像大功告成了？别急，自动化工具只能解决 90% 的问题，剩下的 10% 往往藏在细节里。

手动修正与关键差异排查

拿着生成的代码直接去编译，大概率会报错或者行为异常。我们需要重点检查几个地方。

首先是头文件。虽然 hipify-perl 通常能替换对，但有时它会漏掉一些特定的数学函数头文件。确保文件顶部有：

#include <hip/hip_runtime.h>

如果有用到 printf，在 HIP 里通常不需要特殊处理，但如果涉及设备端打印，可能需要确认 hipPrintf 的支持情况（本例用的是主机端打印，无需改动）。

其次是内存拷贝的方向枚举值。在 CUDA 中我们使用 cudaMemcpyHostToDevice，转换后理论上会变成 hipMemcpyHostToDevice。HIP 确实定义了这些宏，但为了保险起见，建议打开文件搜索一下 Memcpy，确认所有枚举值都已正确前缀化。如果发现还残留着 cuda 开头的枚举，手动改成 hip 开头。

还有一个容易被忽视的点是错误检查。原代码为了简洁没加错误判断，但在迁移调试阶段，强烈建议在每个 HIP API 调用后加上 hipGetLastError() 检查，这样一旦编译通过但运行崩溃，能迅速定位是显存分配失败还是内核启动问题。

修改后的核心片段应该长这样：

hipMalloc((void**)&d_a, size);
hipMalloc((void**)&d_b, size);
hipMalloc((void**)&d_c, size);

hipMemcpy(d_a, h_a, size, hipMemcpyHostToDevice);
// ... 其他拷贝同理

使用 hipcc 编译与运行验证

一切就绪后，轮到编译器登场。AMD 对应的编译器是 hipcc，它基于 Clang/LLVM，用法和 nvcc 非常相似。

执行编译命令：

hipcc vector_add_hip.cpp -o vector_add_hip

如果过程中没有报错，恭喜你，二进制文件 vector_add_hip 已经生成。这时候可以直接运行：

./vector_add_hip

预期输出应该是：

Result check: 3.000000

因为我们是 1.0 + 2.0。如果看到了这个结果，说明你的第一段 CUDA 代码已经成功在 AMD GPU 上跑通了！

如果在运行时遇到 Segmentation fault 或者没有任何输出，先别慌。第一步先用 rocm-smi 确认显卡没掉驱动；第二步检查是否以普通用户权限运行（某些 ROCm 版本可能需要用户加入 render 或 video 用户组）；第三步就是在代码里加上前面提到的错误检查宏，看看到底是哪一步 HIP API 返回了错误码。

写在最后

从 cuda 到 hip，不仅仅是字符的替换，更是思维习惯的微调。通过这次向量加法的实战，你会发现 HIP 对 CUDA 的兼容性做得相当不错，大部分逻辑无需重构。对于更复杂的项目，比如涉及 cuBLAS、cuDNN 的深度学习模型，AMD 也提供了 rocBLAS、MIOpen 等对应库，迁移路径大同小异。

手里有 A 卡的朋友，不妨找个开源的小项目试试 hipify-perl，跨出这一步，你会发现 GPU 计算的天地其实比想象中更宽广。下次我们可以聊聊如何用 HIP 改写一些复杂的算子，或者如何在 PyTorch 中原生启用 ROCm 后端。

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