ROCm HIP GPU 编程实战

AMD ROCm 是 AMD 推出的开源 GPU 计算软件栈，旨在与 NVIDIA CUDA 生态直接竞争。HIP（Heterogeneous-computing Interface for Portability）作为其编程抽象层，允许开发者以接近 CUDA 的语法编写跨平台 GPU 代码。本文从架构到代码，拆解 ROCm/HIP 的完整实践路径。

ROCm 软件栈架构

+-----------------------------------------------------------+
|  User Application Layer (HIP / OpenMP / OpenCL / PyTorch)  |
+-----------------------------------------------------------+
|         HIP Runtime / ROCclr (compute language runtime)   |
+-----------------------------------------------------------+
|  ROCm GPU Driver (amdgpu)  |  ROCprofiler / ROCgdb        |
+-----------------------------------------------------------+
|     AMDGPU Kernel Driver  |  Thunk / ROCsMI (sysmgmt)       |
+-----------------------------------------------------------+
|                   AMD GPU Hardware (CDNA / RDNA)            |
+-----------------------------------------------------------+

核心组件说明：

HIP 编程模型

HIP 与 CUDA 在语法上高度同源：线程层级（grid / block / thread）、共享内存、同步栅栏、流队列等概念完全对齐。

向量加法示例

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

__global__ void vecAdd(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()
{
    int n = 1 << 20;
    size_t bytes = n * sizeof(float);
    float *h_A = new float[n], *h_B = new float[n], *h_C = new float[n];

    // 设备内存分配
    float *d_A, *d_B, *d_C;
    hipMalloc(&d_A, bytes);
    hipMalloc(&d_B, bytes);
    hipMalloc(&d_C, bytes);

    hipMemcpy(d_A, h_A, bytes, hipMemcpyHostToDevice);
    hipMemcpy(d_B, h_B, bytes, hipMemcpyHostToDevice);

    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, n);

    hipMemcpy(h_C, d_C, bytes, hipMemcpyDeviceToHost);

    hipFree(d_A); hipFree(d_B); hipFree(d_C);
    delete[] h_A; delete[] h_B; delete[] h_C;
    return 0;
}

编译命令：

hipcc vecAdd.hip -o vecAdd -O3
./vecAdd

与 nvcc 的区别仅在于前缀：hip 替代 cu，HIP 替代 CUDA。

共享内存与 Bank Conflict 优化

#define TILE_WIDTH 16

__global__ void matMulShared(const float *A, const float *B, float *C,
                             int M, int N, int K)
{
    __shared__ float As[TILE_WIDTH][TILE_WIDTH + 1]; // +1 消除 bank conflict
    __shared__ float Bs[TILE_WIDTH][TILE_WIDTH + 1];

    int row = blockIdx.y * TILE_WIDTH + threadIdx.y;
    int col = blockIdx.x * TILE_WIDTH + threadIdx.x;
    float val = 0.0f;

    for (int t = 0; t < (K + TILE_WIDTH - 1) / TILE_WIDTH; ++t) {
        As[threadIdx.y][threadIdx.x] = A[row * K + t * TILE_WIDTH + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(t * TILE_WIDTH + threadIdx.y) * N + col];
        __syncthreads();

        for (int i = 0; i < TILE_WIDTH; ++i)
            val += As[threadIdx.y][i] * Bs[i][threadIdx.x];
        __syncthreads();
    }

    if (row < M && col < N)
        C[row * N + col] = val;
}

在 MI210 / MI250X 这类 CDNA2 架构上，TILE_WIDTH 取 16 或 32 可充分利用 wavefront（64 线程）调度效率。

hipify：CUDA→HIP 迁移

ROCm 提供 hipify-perl 和 hipify-clang 两套工具，后者基于 Clang AST，转换更准确。

# 1. 单文件转换
hipify-perl cuda_kernel.cu > hip_kernel.hip

# 2. 目录批量转换（推荐 hipify-clang）
mkdir src_hip
hipify-clang cuda_project/ \
    --cuda-path=/usr/local/cuda \
    -o src_hip \
    -I./include \
    --skip-excluded-preprocessor-conditional-blocks

# 3. 生成转换报告（统计 API 映射与警告）
hipify-clang kernel.cu --print-stats

常见转换映射规律：

约 90% 以上的 CUDA 代码可通过 hipify 机械转换，剩余 10% 集中在纹理内存（Texture）、Tensor Core（WMMA）、以及 CUDA Graph API 等 NVIDIA 特化功能。

ROCm 与 CUDA 深度对比

华趣计算实践（国产加速器适配）

近年来，国产 AI/HPC 加速器（如燧原、天数、海光 DCU）多基于 AMD GCN/CDNA 架构派生，天然兼容 ROCm/HIP 工具链。华趣计算场景下的迁移实践：

1. 检测目标设备：通过 rocm-smi 或 rocminfo 确认 compute capability 与显存带宽。

rocminfo | grep "Name:"         # 查看设备名
rocm-smi --showmeminfo vram     # 查看显存占用

2. 调整 wavefront 与 warp 差异：CDNA 架构 wavefront 为 64 线程，CUDA warp 为 32 线程；在 warp-level primitive 代码（__shfl_sync 等）中需进行条件处理或改用 HIP warp primitives。

3. 内存分页锁定：

float *h_ptr;
hipHostMalloc(&h_ptr, bytes, hipHostMallocDefault); // 替代 cudaHostAlloc

4. MIOpen 深度学习推理：若从 cuDNN 迁移，注意 miopenConvolutionForward 的 workspace 分配策略略有不同，需通过 miopenConvolutionForwardGetWorkSpaceSize 动态查询。

ROCm/HIP 作为开源 GPU 计算的支柱，正逐步打破 CUDA 生态的单一垄断局面。对于已有 CUDA 代码库的 HPC 团队，借助 hipify 进行渐进式迁移，是降低平台锁定风险、拥抱多元算力基础设施的有效策略。如需进一步了解 Docker 在 HPC 场景的部署实践，可参考 https://plumephp.com/docker-compose-production/。

ROCm 与 CUDA 兼容性及迁移成本分析

HIP 设计的核心目标是让 CUDA 代码能够以极低的成本迁移到 AMD GPU 上运行。从语法层面看，HIP 几乎完整复刻了 CUDA 的线程层级模型，包括 grid、block、thread、共享内存、同步原语等概念都与 CUDA 一一对应。这种高度同源的抽象意味着，开发者不需要重新学习一套全新的编程范式，只需要熟悉 API 前缀的替换规则即可上手。

然而，兼容性与完全一致之间仍存在关键差异。首先是硬件层面的 warp 粒度差异：NVIDIA GPU 的 warp 为 32 线程，而 AMD CDNA 架构的 wavefront 为 64 线程。这意味着任何依赖 warp-level primitive（如 __shfl_sync、__ballot_sync）的代码都需要进行改写，因为 HIP 的 warp primitive 会自动适配 64 线程的 wavefront。其次，NVIDIA 的 Tensor Core（通过 WMMA API 暴露）在 AMD 端对应的是 rocWMMA 库，但两者的数据布局、寄存器分配策略和性能特征存在差异，需要额外调优。

以下代码展示了如何在 HIP 中处理 warp-level reduce 操作，同时兼容不同 warp 大小：

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

__device__ float warpReduceSum(float val) {
    // HIP 自动适配 wavefront 大小（64 线程）
    // 但代码本身对 32 或 64 线程都适用
    for (int offset = warpSize / 2; offset > 0; offset /= 2) {
        val += __shfl_xor(val, offset);
    }
    return val;
}

__global__ void sumKernel(const float *input, float *output, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    int laneId = threadIdx.x % warpSize;
    int warpId = threadIdx.x / warpSize;

    float val = (tid < n) ? input[tid] : 0.0f;
    val = warpReduceSum(val);

    // 每个 warp 的第一个线程将局部和写入共享内存
    __shared__ float warpSums[32]; // 最大支持 1024 线程 / block
    if (laneId == 0) warpSums[warpId] = val;
    __syncthreads();

    // block 级别的最终归约
    if (warpId == 0) {
        val = (laneId < blockDim.x / warpSize) ? warpSums[laneId] : 0.0f;
        val = warpReduceSum(val);
        if (laneId == 0) atomicAdd(output, val);
    }
}

性能差异方面，在矩阵乘法、FFT 等标准 HPC 负载上，MI250X 与 A100 的性能差距通常在 10% 到 30% 之间，具体取决于内存带宽利用率与算法对 wavefront 的适配程度。对于纯计算密集型任务（如蒙特卡洛模拟），差距往往更小。迁移成本的评估应综合考虑代码库规模、CUDA 特化 API 的使用密度以及性能敏感程度。

性能 Profiling 与 Kernel 调优

写出可运行的 HIP 代码只是第一步，真正释放 AMD GPU 性能需要系统化的 profiling 和调优流程。ROCm 提供了 rocprof 作为与 NVIDIA Nsight Compute / Systems 对标的全栈性能分析工具，能够采集 kernel 执行时间、内存带宽利用率、L2 cache 命中率、 wavefront 占用率等关键指标。

rocprof 支持多种采集模式：

# 1. 基本 Kernel 计时
rocprof --stats ./vecAdd

# 2. 采集 Hardware Counters（需要 root 权限或特定系统配置）
rocprof -i counters.txt --stats ./matMul

# counters.txt 示例内容：
# pmc : Wavefronts VALUInsts SALUInsts
# pmc : FetchSize WriteSize L2CacheHit

# 3. 生成可视化时间线（JSON 格式，可用 chrome://tracing 查看）
rocprof --timestamp on --trace-period thread -o trace.json ./benchmark

性能调优的核心关注点包括：

1. 内存带宽瓶颈：大多数 GPU 计算任务受限于 HBM 带宽而非算力。通过合并全局内存访问（coalesced access）可以显著提升有效带宽。在 MI250X 上，HBM 带宽高达 1.6 TB/s，但非合并访问可能让有效带宽下降到理论值的 20% 以下。

2. Bank Conflict 优化：共享内存的 bank 数量在 CDNA 架构上为 32，当多个线程同时访问同一个 bank 的不同地址时会发生 bank conflict，导致串行化访问。通过在共享内存数组的列维度加 padding（如 TILE_WIDTH + 1），可以有效消除 conflict。

3. 寄存器压力与占用率：每个 wavefront 最多可使用 256 个 vector 寄存器（VGPR）。如果 kernel 使用了过多寄存器，会导致 GPU 无法同时调度足够的 wavefront 来隐藏内存延迟，从而降低占用率（occupancy）。可通过编译器选项 -ffp-contract=fast 或者手动减少临时变量来缓解。

# 查看编译器生成的寄存器使用报告
hipcc kernel.hip -o kernel -O3 --save-temps
# 在生成的 .s 文件中搜索 "vgpr" 和 "sgpr" 字段

4. 内核融合（Kernel Fusion）：减少多 kernel 之间的全局内存往返是提升整体吞吐量的有效手段。对于深度学习训练中的逐元素操作（elementwise op，如 ReLU、Scale、BiasAdd），可以将多个操作合并到同一个 kernel 中执行，减少 HBM 读写次数。

性能调优是一个迭代过程：先用 rocprof 定位热点 kernel，再针对性优化内存访问模式、寄存器使用和 kernel 调度策略。关于 HPC 中更通用的性能分析方法论，可以参考 https://plumephp.com/hpc-roofline-model/ 中的 Roofline 模型分析。

多 GPU 并行与分布式通信

现代 HPC 和 AI 训练任务几乎都离不开多 GPU 并行的支持。AMD ROCm 生态同样提供了与 NVIDIA NCCL 对应的通信库 rccl（ROCm Communication Collectives Library），它实现了 Ring AllReduce、AllGather、Broadcast 等标准集合通信原语，能够与 PyTorch Distributed、Horovod 等框架无缝集成。

多 GPU 并行主要有两种策略：数据并行（Data Parallelism）和模型并行（Model Parallelism）。数据并行将同一批数据的不同子集分发到多个 GPU 上，每个 GPU 独立计算梯度后通过 AllReduce 同步。这种方式实现简单，适合模型参数能放入单卡显存的场景。模型并行则是将模型的不同层或不同部分分布到多个 GPU 上，适用于超大模型（如 GPT-3 级别）训练。

以下是一个使用 rccl 进行 Ring AllReduce 的简化示例：

// rccl_allreduce.cpp
#include <rccl/rccl.h>
#include <hip/hip_runtime.h>
#include <vector>
#include <iostream>

int main(int argc, char **argv) {
    int nGpus = 4;
    int localRank = 0; // 实际由 MPI / 环境变量注入

    // 初始化 RCCL
    ncclComm_t comm;
    ncclUniqueId id;
    hipStream_t stream;

    if (localRank == 0) ncclGetUniqueId(&id);
    // 假设已通过 MPI_Bcast 将 id 广播到所有 rank

    hipSetDevice(localRank);
    hipStreamCreate(&stream);
    ncclCommInitRank(&comm, nGpus, id, localRank);

    int size = 1024 * 1024;
    std::vector<float> h_data(size, localRank + 1.0f);
    float *d_data;
    hipMalloc(&d_data, size * sizeof(float));
    hipMemcpy(d_data, h_data.data(), size * sizeof(float), hipMemcpyHostToDevice);

    // 执行 AllReduce，将各 GPU 的 d_data 求和并同步到所有 GPU
    ncclAllReduce(d_data, d_data, size, ncclFloat, ncclSum, comm, stream);

    hipStreamSynchronize(stream);

    hipFree(d_data);
    ncclCommDestroy(comm);
    hipStreamDestroy(stream);
    return 0;
}

编译与链接时需要指定 rccl 路径：

hipcc rccl_allreduce.cpp -o rccl_demo -lrccl -I/opt/rocm/include -L/opt/rocm/lib

在 PyTorch 中，RCCL 作为后端被透明封装，用户只需设置环境变量即可切换：

export HSA_FORCE_FINE_GRAIN_PCIE=1
export NCCL_DEBUG=INFO
python -m torch.distributed.run --nproc_per_node=4 train.py

MI250X 的每个计算模块（Compute Die）内部通过 Infinity Fabric 连接，跨 die 通信延迟约为 3-5 微秒，与 NVIDIA NVLink 相当。若节点间通信通过 InfiniBand，ROCm 同样支持 GPUDirect RDMA，可实现 GPU 内存与网卡之间的零拷贝数据传输。关于 InfiniBand 网络配置的更多细节，请参阅 https://plumephp.com/hpc-infini-band/。

容器化部署：Docker + ROCm

在 HPC 集群中，容器化部署已经成为管理复杂依赖和确保环境一致性的标准做法。ROCm 官方提供了 rocm/pytorch、rocm/tensorflow 等预构建 Docker 镜像，极大简化了深度学习框架在 AMD GPU 上的部署流程。与 NVIDIA Docker（nvidia-docker）类似，AMD 提供了 rocker 或者直接在 Docker 中通过 --device 挂载 AMDGPU。

使用 ROCm Docker 镜像的基本流程如下：

# 1. 拉取官方 PyTorch ROCm 镜像
docker pull rocm/pytorch:rocm6.0_ubuntu22.04_py3.10_pytorch_2.1.2

# 2. 运行容器并挂载 AMD GPU
docker run -it --rm \
    --device=/dev/kfd \
    --device=/dev/dri \
    --group-add video \
    --cap-add=SYS_PTRACE \
    --security-opt seccomp=unconfined \
    -v $(pwd):/workspace \
    rocm/pytorch:rocm6.0_ubuntu22.04_py3.10_pytorch_2.1.2

# 3. 容器内验证 ROCm 可用性
rocminfo
python -c "import torch; print(torch.cuda.is_available()); print(torch.version.hip)"

对于需要使用 Singularity / Apptainer 的 HPC 环境（常见于 Slurm 调度集群），可以直接从 Docker 镜像转换：

# 构建 Singularity 镜像
singularity pull rocm-pytorch.sif docker://rocm/pytorch:rocm6.0_ubuntu22.04_py3.10_pytorch_2.1.2

# 提交到 Slurm 集群运行
singularity exec --rocm rocm-pytorch.sif python train.py

容器化部署的关键注意事项包括：

• 内核版本兼容性：ROCm 驱动 amdgpu 对内核版本有一定要求，通常建议使用 Ubuntu 20.04/22.04 的 LTS 内核或对应发行版的 HWE 内核。容器内的 ROCm 版本应与宿主机驱动版本匹配，否则可能出现运行时错误。
• 显存隔离：与 NVIDIA MIG（Multi-Instance GPU）类似，AMD GPU 目前主要通过 rocm-smi 的显存限制功能进行粗粒度隔离，细粒度的多租户隔离仍在发展中。
• 存储挂载策略：训练数据通常存储在 Lustre GPFS 或 NFS 上，通过 -v 挂载进容器时要确保 ACL 和权限正确。关于高性能文件系统的配置，可参考 https://plumephp.com/hpc-lustre-filesystem/。

更多 Docker 生产环境最佳实践，包括多阶段构建和镜像安全加固，请参见 https://plumephp.com/docker-compose-production/。

ML 框架支持矩阵与选型建议

ROCm 生态的一个核心吸引力在于对主流深度学习框架的原生支持。截至 ROCm 6.x 版本，PyTorch、TensorFlow 和 JAX 均提供了官方或社区的 ROCm 后端支持，但各框架的成熟度、功能覆盖度和更新节奏存在差异。

PyTorch + ROCm 是目前最成熟的组合。AMD 官方维护 rocm/pytorch Docker 镜像，并与 PyTorch 上游社区保持密切同步。绝大多数 PyTorch 模型可以在 ROCm 后端上直接运行而无需修改代码，包括 torchvision、torchaudio 和 torchtext。分布式训练通过 torch.distributed + RCCL backend 支持，DDP 和 FSDP 均已可用。尚未完全支持的特性主要是 CUDA Graph 和某些自定义 CUDA 扩展，这些需要等待 hipify 兼容。

TensorFlow + ROCm 方面，AMD 提供了 rocm/tensorflow 镜像和 pip 包 tensorflow-rocm。功能上覆盖了 Keras API、SavedModel 导出和 TF Serving，但 XLA 编译器（用于图优化和 JIT）的 ROCm 后端成熟度略低于 PyTorch 的 torch.compile。对于推理场景，TensorFlow Lite 的 ROCm delegate 也在积极开发中。

JAX + ROCm 的支持相对最新。Google 与 AMD 合作开发了 jax-rocm 插件，利用 XLA 的 AMDGPU 编译器后端。JAX 的组合变换（vmap、grad、pmap）在 ROCm 上可用，但部分高级特性（如 Pallas 内核）的支持度还在追赶中。

以下是在各框架中检测 ROCm 可用性的示例代码：

# pytorch_rocm_check.py
import torch

print(f"PyTorch version: {torch.__version__}")
print(f"HIP version: {torch.version.hip}")
print(f"GPU available: {torch.cuda.is_available()}")
if torch.cuda.is_available():
    print(f"GPU count: {torch.cuda.device_count()}")
    print(f"GPU name: {torch.cuda.get_device_name(0)}")

# 简单 GEMM 基准测试
a = torch.randn(4096, 4096, device='cuda')
b = torch.randn(4096, 4096, device='cuda')
import time
torch.cuda.synchronize()
start = time.time()
c = torch.matmul(a, b)
torch.cuda.synchronize()
print(f"GEMM 耗时: {(time.time() - start) * 1000:.2f} ms")

# tensorflow_rocm_check.py
import tensorflow as tf

print(f"TensorFlow version: {tf.__version__}")
print(f"Built with ROCm: {tf.test.is_built_with_rocm()}")
print(f"GPU available: {tf.config.list_physical_devices('GPU')}")

# 简单卷积基准
x = tf.random.normal([64, 224, 224, 3])
w = tf.random.normal([7, 7, 3, 64])
_ = tf.nn.conv2d(x, w, strides=[1, 2, 2, 1], padding='SAME')

选型建议方面：

• 从头训练深度学习模型：首选 PyTorch + ROCm，社区生态最活跃，迁移成本最低。
• 传统 HPC + 深度学习混合负载：如果同时涉及 PETSc 求解器等传统 HPC 库与深度学习推理，需要考虑 ROCm 与 MPI 的协同。可以参考 https://plumephp.com/hpc-mpi-basics/ 了解 MPI 的基础知识。
• 生产推理服务：TensorFlow Serving + ROCm 的部署流程已相对稳定，但监控和自动扩缩容工具链不如 NVIDIA Triton 成熟。
• 科研快速原型：JAX 的功能式编程模型在 ROCm 上运行良好，适合需要大量自动微分和并行化的研究工作。

框架支持的持续演进意味着 ROCm 生态的可用性正在快速提升。对于已有 Python 机器学习代码的团队，通过将 torch.cuda 替换为 torch.hip（实际上 PyTorch 已经透明处理），结合 rocm/pytorch Docker 镜像，通常可以在一天内完成基础的验证和迁移工作。关于 Python 现代工具链的更多内容，可以参考 https://plumephp.com/python-modern-toolchain/。

ROCm/HIP 作为开源 GPU 计算的支柱，正逐步打破 CUDA 生态的单一垄断局面。对于已有 CUDA 代码库的 HPC 团队，借助 hipify 进行渐进式迁移，是降低平台锁定风险、拥抱多元算力基础设施的有效策略。