从 CUDA 到 HIP：思维切换与核心概念重构

很多刚从 NVIDIA 生态转向 AMD Instinct GPU 的同学，第一反应往往是："HIP 不就是改了个前缀的 CUDA 吗？”确实，HIP（Heterogeneous-Compute Interface for Portability）的设计初衷就是最大程度兼容 CUDA 语法，让 cudaMalloc 变成 hipMalloc，<<< >>> 启动符保持不变。但如果只停留在“查找替换”的层面，遇到复杂的内存模型或特定的硬件架构时，很容易踩坑。真正的跨平台开发，需要理解两者在底层运行时和编译链路上的微妙差异，建立起一套不依赖单一厂商的异构计算思维。

设备管理与显存操作：不仅仅是改名

在 CUDA 中，我们习惯用 cudaSetDevice 和 cudaGetDeviceProperties 来管理设备。HIP 保留了类似的 API 风格，但在头文件包含和命名空间上更加规范。编写 HIP 程序时，必须包含 <hip/hip_runtime.h>，这是所有功能的入口。

显存管理是迁移的第一道关卡。虽然 hipMalloc 和 hipMemcpy 的参数顺序与 CUDA 完全一致，但有一个细节值得注意：HIP 对指针类型的检查在某些编译器配置下更为严格。以下是一个标准的显存分配与数据拷贝示例，展示了如何安全地在主机与设备间传输数据：

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

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

int main() {
    const int size = 1024;
    float *h_data, *d_data;
    
    // 主机内存分配
    h_data = new float[size];
    for(int i = 0; i < size; i++) h_data[i] = 1.0f;

    // 设备内存分配 (对应 cudaMalloc)
    HIP_CHECK(hipMalloc(&d_data, size * sizeof(float)));

    // 数据拷贝 (对应 cudaMemcpy)
    HIP_CHECK(hipMemcpy(d_data, h_data, size * sizeof(float), hipMemcpyHostToDevice));

    // ... _kernel 启动 ...

    // 清理资源
    HIP_CHECK(hipFree(d_data));
    delete[] h_data;
    
    return 0;
}

这段代码看似平淡，但 HIP_CHECK 宏的使用是良好的工程习惯。AMD 的驱动在某些非致命错误上可能不会像 NVIDIA 那样立即终止程序，显式检查返回值能帮你快速定位是显存不足还是指针对齐问题。

线程网格与同步机制：深入执行模型

理解了内存管理，接下来是核心的并行执行模型。CUDA 开发者对 Grid、Block、Thread 的概念早已烂熟于心，HIP 完全复用了这一层级结构。但在实际编写 Kernel 时，共享内存（Shared Memory）的使用策略往往决定了性能上限。

在 HIP 中，声明共享内存的方式与 CUDA 几乎无异，使用 __shared__ 关键字。然而，在处理线程块内同步时，__syncthreads() 的行为需要格外小心。特别是在 Instinct MI300 系列这类拥有复杂缓存层级的架构上，错误的同步会导致数据竞争或死锁。

下面是一个向量相加的 Kernel 示例，展示了基本的线程索引计算和同步逻辑：

__global__ void vectorAddKernel(float *A, float *B, float *C, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 边界检查必不可少
    if (idx < n) {
        C[idx] = A[idx] + B[idx];
    }
}

// 启动配置
int threadsPerBlock = 256;
int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
vectorAddKernel<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, n);

对于更复杂的场景，比如矩阵分块计算，我们需要在共享内存中暂存数据。此时，必须在加载数据后、计算前调用 __syncthreads()，确保所有线程都完成了数据搬运。这一点在从 CUDA 迁移代码时容易被忽视，因为不同代际的 GPU 对未同步访问的容忍度不同，而 HIP 代码旨在运行于多种硬件，必须严格遵守同步规范。

实战演练：手写矩阵乘法与性能思考

为了真正掌握 HIP，我们来动手实现一个经典的矩阵乘法（GEMM）。这不仅是对语法的检验，更是理解内存访问模式的机会。我们将对比“朴素实现”与“共享内存优化版”的思路。

朴素版本直接全局内存读写，效率较低；优化版则将矩阵分块载入共享内存，减少全局内存访问次数。以下是优化版的核心逻辑片段：

#define TILE_SIZE 16

__global__ void matMulOptimized(float *A, float *B, float *C, int width) {
    __shared__ float sharedA[TILE_SIZE][TILE_SIZE];
    __shared__ float sharedB[TILE_SIZE][TILE_SIZE];

    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    float sum = 0.0f;

    for (int t = 0; t < (width + TILE_SIZE - 1) / TILE_SIZE; ++t) {
        // 协作加载数据到共享内存
        if (row < width && t * TILE_SIZE + threadIdx.x < width)
            sharedA[threadIdx.y][threadIdx.x] = A[row * width + t * TILE_SIZE + threadIdx.x];
        else
            sharedA[threadIdx.y][threadIdx.x] = 0.0f;

        if (col < width && t * TILE_SIZE + threadIdx.y < width)
            sharedB[threadIdx.y][threadIdx.x] = B[(t * TILE_SIZE + threadIdx.y) * width + col];
        else
            sharedB[threadIdx.y][threadIdx.x] = 0.0f;

        __syncthreads(); // 关键同步点

        // 在共享内存中进行计算
        for (int k = 0; k < TILE_SIZE; ++k)
            sum += sharedA[threadIdx.y][k] * sharedB[k][threadIdx.x];

        __syncthreads(); // 确保当前块计算完成再进入下一轮
    }

    if (row < width && col < width)
        C[row * width + col] = sum;
}

在实际测试中，这种手写优化在 Instinct GPU 上能带来显著的性能提升。有趣的是，如果你使用 HIPify 工具自动转换 CUDA 代码，它通常只能生成朴素版本。自动化工具擅长语法映射，却难以理解算法层面的内存局部性优化。这就是为什么初级工程师必须亲手写一遍 HIP 代码的原因：工具能帮你完成 80% 的迁移工作，但剩下 20% 的性能榨取和架构适配，全靠开发者对硬件的理解。

构建跨平台的开发思维

学习 HIP 的最终目的，不是为了写出只能在 AMD 卡上跑的代码，而是为了掌握一种“一次编写，多处编译”的能力。HIP 代码既可以在 AMD GPU 上通过 hipcc 编译原生运行，也可以在 NVIDIA GPU 上编译为 CUDA 后端（虽然生产环境较少这么做，但验证了其兼容性）。

当你开始关注 hipGetDeviceProperties 中的架构字段（如 gfx90a vs sm_80），当你学会根据缓存大小调整 TILE_SIZE，你就已经超越了简单的 API 调用者，成为了真正的异构计算工程师。这种思维模式让你在面对未来可能出现的任何新加速卡时，都能快速上手，不再被厂商锁定的生态所束缚。

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