hip-rocm

HIP / ROCm

用途

指导 Agent 使用 HIP 进行 AMD GPU 编程：HIP 运行时 API、hipcc 编译、使用 HIPIFY（hipify-perl、hipify-clang）移植 CUDA 代码、ROCm 工具链设置、使用 rocprof 进行性能分析、使用 rocgdb 调试、HIP 与 CUDA API 映射，以及针对 MI300X 的优化。

适用场景

• 将现有 CUDA 代码库移植到 AMD GPU
• 在 Linux 上为 MI200/MI300 硬件设置 ROCm
• 为 AMD 数据中心 GPU 编写原生 HIP 内核
• 使用 rocprof 或 rocprofiler-sdk 对 HIP 应用进行性能分析
• 使用 rocgdb 或计算消毒器调试设备故障
• 使用 HIP 可移植性宏构建多厂商 GPU 代码

工作流

1. 安装并验证 ROCm

# Ubuntu/Debian (check ROCm docs for your distro version)
sudo apt install rocm-dev rocm-libs hip-dev

# Verify
rocminfo | head -30
hipconfig --version
hipcc --version

# List devices
rocm-smi

设置编译用的 GPU 目标：

export AMDGPU_TARGETS=gfx942   # MI300X
export HIP_PLATFORM=amd

2. 最简 HIP 内核

// vector_add.hip
#include <hip/hip_runtime.h>
#include <stdio.h>

__global__ void vector_add(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(void) {
    const int n = 1 << 20;
    size_t bytes = n * sizeof(float);
    float *d_a, *d_b, *d_c;

    hipMalloc(&d_a, bytes);
    hipMalloc(&d_b, bytes);
    hipMalloc(&d_c, bytes);

    int threads = 256;
    int blocks = (n + threads - 1) / threads;
    hipLaunchKernelGGL(vector_add, dim3(blocks), dim3(threads), 0, 0,
                       d_a, d_b, d_c, n);
    hipDeviceSynchronize();

    hipFree(d_a); hipFree(d_b); hipFree(d_c);
    return 0;
}

hipcc -O3 --offload-arch=gfx942 -o vector_add vector_add.hip
./vector_add

3. 使用 HIPIFY 移植 CUDA → HIP

# Perl-based batch converter (quick port)
hipify-perl cuda_kernel.cu > cuda_kernel.hip

# Clang-based (more accurate, preserves structure)
hipify-clang cuda_project/ -o hip_project/ --cuda-path=/usr/local/cuda

# Convert single file in place
hipify-clang -inplace --cuda-path=/usr/local/cuda main.cu

常见 API 映射：

CUDA	HIP
cudaMalloc	hipMalloc
cudaMemcpy	hipMemcpy
cudaMemcpyAsync	hipMemcpyAsync
cudaStream_t	hipStream_t
<<<grid, block>>>	hipLaunchKernelGGL 或 <<<>>>（HIP 支持 CUDA 语法）
__syncthreads()	__syncthreads()（相同）
threadIdx / blockIdx	相同的内置变量

用于双重编译的可移植性头文件：

#ifdef __HIP_PLATFORM_AMD__
#include <hip/hip_runtime.h>
#else
#include <cuda_runtime.h>
#define hipMalloc cudaMalloc
#define hipMemcpy cudaMemcpy
// ... more macros
#endif

4. hipcc 标志

# Target specific GPU architecture
hipcc --offload-arch=gfx942 -O3 -o app main.hip

# Multiple architectures
hipcc --offload-arch=gfx90a --offload-arch=gfx942 -o app main.hip

# Debug
hipcc -g -O0 --offload-arch=gfx942 -o app_debug main.hip

# Link with rocBLAS
hipcc -lrocblas -o app main.hip

5. 使用 rocprof 进行性能分析

# Basic kernel trace
rocprof --stats ./app

# CSV metrics output
rocprof -i input.csv -o output.csv ./app

# input.csv example:
# pmc: SQ_INSTS_VALU_ADD_F32,SQ_INSTS_VALU_MUL_F32,GRBM_COUNT

# ROCm 6.x rocprofiler-sdk (preferred for new projects)
rocprofv3 --kernel-trace -- ./app

关键指标（AMD 术语）：

• VALU 利用率 — 计算单元活动度
• LDS bank 冲突 — 共享内存（LDS）停顿
• 内存吞吐量 — HBM 带宽利用率

6. 使用 rocgdb 调试

# Build with debug symbols
hipcc -g -O0 --offload-arch=gfx942 -o app_debug main.hip

rocgdb ./app_debug

(rocgdb) break vector_add
(rocgdb) run
(rocgdb) info rocm kernels
(rocgdb) rocm thread 0 0 0
(rocgdb) print i

AMD 还通过 ROCm 的 roc-obj-extract 和可用的内存检查工具提供与 compute-sanitizer 等效的工具。

7. MI300X 优化

# Enable MFMA (matrix fused multiply-add) instructions
hipcc --offload-arch=gfx942 -munsafe-fp-atomics -O3 -o app main.hip

优化	MI300X 注意事项
矩阵运算	GEMM 使用 rocBLAS/hipBLASLt；自定义运算使用 MFMA 内部函数
HBM 带宽	峰值约为 5.3 TB/s（MI300X）——尽可能提高内存合并访问率，以接近峰值
Wavefront 大小	64 个线程（CUDA warp 为 32）——调整归约模式
LDS（共享内存）	每个 CU 为 64 KB；注意 bank 冲突

考虑 Wavefront 的归约：

__device__ float warp_reduce_sum(float val) {
    // AMD wavefront = 64 lanes
    for (int offset = 32; offset > 0; offset >>= 1)
        val += __shfl_down(val, offset);
    return val;
}

8. 库生态

NVIDIA	AMD ROCm
cuBLAS	rocBLAS / hipBLAS
cuDNN	MIOpen
NCCL	rccl
Thrust	hipCUB（可移植）
cuFFT	rocFFT

hipcc -lrocblas -o gemm_test gemm.hip

常见问题

症状	原因	解决方法
hipErrorNoDevice	ROCm 驱动未加载	检查 rocm-smi；将用户添加到 render 组
架构二进制文件错误	gfx* 目标不匹配	使用 rocminfo → 设置 --offload-arch
hipify 移植不完整	CUDA 专属 API	手动修复：协作组、纹理引用
比 CUDA 参考实现慢	Wavefront 为 64，warp 为 32	将块大小调整为 64 的倍数
HSA_STATUS_ERROR	GPU 忙或内存不足	使用 rocm-smi --showmeminfo；减少分配
rocprof 输出为空	未启动内核	启动后验证 hipGetLastError()

相关技能

• skills/gpu/cuda — 被移植的 CUDA 源代码模式
• skills/gpu/cuda-profiling — Nsight 概念与 rocprof 的对应关系
• skills/gpu/gpu-memory-model — Wavefront 与 warp、合并访问规则
• skills/gpu/triton-lang — Triton 通过 ROCm 后端支持 AMD
• skills/compilers/llvm — HIP 使用 Clang/LLVM 工具链