AMD GPU开发入门指南：从环境搭建到首个CUDA程序

共计 2200 个字符，预计需要花费 6 分钟才能阅读完成。

AMD GPU 凭借开源的 ROCm 平台和 HSA 统一内存架构，正在成为异构计算领域的重要选择。其 SIMD 执行模型通过 wavefront 调度实现高效并行，同时保持与 CUDA 生态的兼容性。对于需要替代方案或跨平台部署的开发者，AMD 提供了从机器学习到科学计算的完整加速方案。

环境搭建：Ubuntu 系统下的 ROCm 安装

ROCm（Radeon Open Compute）软件栈包含以下核心组件：
– 内核驱动（amdgpu）
– 运行时库（HIP Runtime）
– 编译器（HIPCC）
– 调试工具（rocprofiler）

在 Ubuntu 20.04/22.04 上安装指定版本（以 5.7 为例）：

sudo apt update && sudo apt install -y rocm-opencl-runtime
sudo apt install -y rocm-hip-sdk=5.7.50700-1

验证安装成功的标准检查流程：

1. 确认设备识别

   rocminfo | grep -i 'gpu arch'

2. 测试编译器

   hipcc --version

HIP 编程实战：向量加法示例

以下是与 CUDA 语法对比的 HIP 实现（注意错误检查宏的使用）：

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

#define HIP_CHECK(cmd) \
{\n    hipError_t error = cmd;\
    if (error != hipSuccess) {\
        fprintf(stderr, "HIP error %d: %s\n", error, hipGetErrorString(error));\
        exit(EXIT_FAILURE);\
    }\
}

__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]; // SIMD 并行执行单元
}

int main() {
    const int N = 1<<20;
    std::vector<float> h_A(N), h_B(N), h_C(N);

    // 初始化主机数据
    for (int i=0; i<N; ++i) {h_A[i] = 1.0f; h_B[i] = 2.0f;
    }

    // 设备内存分配（注意 HIP 与 CUDA API 的对应关系）float *d_A, *d_B, *d_C;
    HIP_CHECK(hipMalloc(&d_A, N*sizeof(float)) ); // 设备内存分配
    HIP_CHECK(hipMalloc(&d_B, N*sizeof(float)) );
    HIP_CHECK(hipMalloc(&d_C, N*sizeof(float)) );

    // 数据拷贝（主机→设备）HIP_CHECK(hipMemcpy(d_A, h_A.data(), N*sizeof(float), hipMemcpyHostToDevice) );
    HIP_CHECK(hipMemcpy(d_B, h_B.data(), N*sizeof(float), hipMemcpyHostToDevice) );

    // 内核启动（注意 wavefront 与 warp 的尺寸差异）dim3 blocks(256);
    dim3 grids((N + blocks.x - 1) / blocks.x);
    vectorAdd<<<grids, blocks>>>(d_A, d_B, d_C, N);

    // 结果回传（设备→主机）HIP_CHECK(hipMemcpy(h_C.data(), d_C, N*sizeof(float), hipMemcpyDeviceToHost) );

    // 验证结果
    for (int i=0; i<N; ++i) assert(fabs(h_C[i] - 3.0f) < 1e-6);

    HIP_CHECK(hipFree(d_A) ); // 内存释放
    HIP_CHECK(hipFree(d_B) );
    HIP_CHECK(hipFree(d_C) );
    return 0;
}

开发避坑指南

内核参数配置

• AMD GPU 的 wavefront 大小通常为 64（NVIDIA warp 为 32），错误配置会导致计算单元利用率下降
• 建议使用 hipDeviceProp_t 动态获取 wavefrontSize：

  hipDeviceProp_t prop;
  HIP_CHECK(hipGetDeviceProperties(&prop, 0));
  int waveSize = prop.warpSize;

版本冲突解决

当出现 hipErrorNoBinaryForGpu 错误时：
1. 检查所有 ROCm 组件版本是否一致

dpkg -l | grep rocm

2. 若混合安装了不同版本，需完全卸载后重新安装

sudo apt autoremove rocm-*

性能优化思考方向

1. 如何根据 GPU 的 CU 数量（Compute Units）确定最优的 grid/blocksize 组合？
2. 在 HSA 架构下，统一内存模型对数据预取策略有何影响？
3. 当存在多个 kernel 连续执行时，怎样利用 HIP Stream 实现异步流水线？

通过这个基础示例，我们实现了从环境搭建到核心 API 调用的完整流程。AMD GPU 开发虽然生态仍在完善，但其开放的架构设计和逐渐丰富的工具链，为异构计算提供了更多可能性。建议下一步通过 ROCm 的官方示例库深入学习内存一致性模型等高级特性。