📅 最后更新：2026-03-30
🎯 目标：帮助你从零开始掌握 AMD ROCm GPU 计算平台
📚 版本：基于 ROCm 7.2.1（生产版本）

📖 目录

1. ROCm 简介
2. 环境准备与安装
3. HIP 编程基础
4. GPU 架构与性能优化
5. ROCm 库与工具
6. AI/深度学习应用
7. 实战项目
8. 调试与故障排除​​​​​​​

---

1. ROCm 简介

1.1 什么是 ROCm？

ROCm（Radeon Open Compute Platform）是 AMD 开发的开源 GPU 计算软件栈，用于高性能计算（HPC）和人工智能（AI）工作负载。

核心特点：

• 🔓 开源：完全开放源代码
• 🚀 高性能：针对 AMD Instinct 和 Radeon GPU 优化
• 🔧 多语言支持：HIP、OpenCL、OpenMP 等
• 🌐 跨平台：支持 Linux 和 Windows

1.2 ROCm 生态系统

ROCm 软件栈
├── HIP (Heterogeneous-Compute Interface for Portability)
│   ├── C++ Runtime API
│   └── Kernel 语言
├── ROCm 库
│   ├── rocBLAS（线性代数）
│   ├── rocFFT（快速傅里叶变换）
│   ├── rocRAND（随机数生成）
│   └── MIOpen（机器学习）
├── 编译器
│   ├── HIPCC（HIP 编译器）
│   └── LLVM/Clang
└── 工具
    ├── rocprof（性能分析）
    ├── rocgdb（调试器）
    └── HIP Trace（跟踪工具）

1.3 支持的硬件

AMD Instinct GPU（数据中心）：

• MI300X、MI300A 系列
• MI250X、MI250 系列
• MI210、MI100 系列

AMD Radeon GPU（工作站）：

• Radeon RX 7000 系列
• Radeon RX 6000 系列
• Radeon PRO W 系列

AMD Ryzen APU：

• 支持 ROCm 的 Ryzen 处理器

---

2. 环境准备与安装

2.1 系统要求（Linux）

操作系统：

• Ubuntu 22.04.x / 24.04.x（推荐）
• RHEL 9.x
• SLES 15 SP6

硬件要求：

• x86_64 或 AMD64 处理器
• 至少 8GB RAM（推荐 16GB+）
• 支持 ROCm 的 AMD GPU
• PCIe 4.0 x16 插槽（推荐）

2.2 快速安装（Ubuntu）

# 1. 添加 ROCm 仓库密钥
wget https://repo.radeon.com/rocm/rocm.gpg.key
sudo apt-key add rocm.gpg.key

# 2. 添加 ROCm 仓库
echo 'deb [arch=amd64 signed-by=/usr/share/keyrings/rocm.gpg] \
  https://repo.radeon.com/rocm/apt/7.2.1 jammy main' | \
  sudo tee /etc/apt/sources.list.d/rocm.list

# 3. 更新包列表
sudo apt update

# 4. 安装 ROCm
sudo apt install rocm

# 5. 添加用户到 video 和 render 组
sudo usermod -a -G video $USER
sudo usermod -a -G render $USER

# 6. 重启或重新登录

2.3 验证安装

# 检查 ROCm 版本
rocminfo

# 检查 GPU 设备
rocm-smi

# 编译测试程序
hipcc --version

# 运行 ROCm 示例
/opt/rocm/share/hip-samples/cmake/samples/0_Intro/simpleHyperQ/simpleHyperQ

2.4 Docker 安装（推荐）

# 拉取 ROCm Docker 镜像
docker pull rocm/pytorch:latest

# 运行容器（需要 KFD 和 DRI 设备）
docker run -it --device /dev/kfd --device /dev/dri \
  --group-add video --cap-add=SYS_PTRACE \
  --security-opt seccomp=unconfined \
  rocm/pytorch:latest bash

# 在容器内验证
python -c "import torch; print(torch.cuda.is_available())"

2.5 Windows 安装（HIP SDK）

# 1. 下载 HIP SDK for Windows
# 访问：https://rocm.docs.amd.com/projects/install-on-windows/

# 2. 安装 Visual Studio 2019/2022
# 需要 "Desktop Development with C++" 工作负载

# 3. 安装 HIP SDK
# 运行下载的安装程序

# 4. 验证安装
hipcc --version

---

3. HIP 编程基础

3.1 什么是 HIP？

HIP（Heterogeneous-Compute Interface for Portability）是 AMD 开发的 GPU 编程模型，语法与 NVIDIA CUDA 高度相似，便于代码移植。

3.2 第一个 HIP 程序

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

// GPU Kernel（在设备上执行）
__global__ void helloFromGPU() {
    printf("Hello from GPU! Thread ID: %d\n", threadIdx.x);
}

int main() {
    printf("Hello from CPU!\n");
    
    // 启动 kernel：1 个 block，4 个线程
    helloFromGPU<<<1, 4>>>();
    
    // 同步等待 GPU 完成
    hipDeviceSynchronize();
    
    return 0;
}

编译和运行：

hipcc -o hello_world hello_world.hip
./hello_world

3.3 HIP 编程模型核心概念

3.3.1 线程层次结构

Grid（网格）
└── Block（块）
    └── Thread（线程）

关键变量：

• threadIdx.x/y/z：线程在块内的索引
• blockIdx.x/y/z：块在网格内的索引
• blockDim.x/y/z：每个块的线程数
• gridDim.x/y/z：网格的块数

全局线程 ID 计算：

int globalThreadId = blockIdx.x * blockDim.x + threadIdx.x;

3.3.2 内存类型

__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() {
    // 主机内存
    float *h_A, *h_B, *h_C;
    
    // 设备内存
    float *d_A, *d_B, *d_C;
    
    // 分配设备内存
    hipMalloc(&d_A, n * sizeof(float));
    hipMalloc(&d_B, n * sizeof(float));
    hipMalloc(&d_C, n * sizeof(float));
    
    // 主机到设备数据传输
    hipMemcpy(d_A, h_A, n * sizeof(float), hipMemcpyHostToDevice);
    hipMemcpy(d_B, h_B, n * sizeof(float), hipMemcpyHostToDevice);
    
    // 启动 kernel
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, n);
    
    // 设备到主机数据传输
    hipMemcpy(h_C, d_C, n * sizeof(float), hipMemcpyDeviceToHost);
    
    // 释放设备内存
    hipFree(d_A);
    hipFree(d_B);
    hipFree(d_C);
}

3.4 共享内存优化

__global__ void matrixTransposeShared(float *in, float *out, int width, int height) {
    // 共享内存声明
    __shared__ float tile[32][32];
    
    int x = blockIdx.x * 32 + threadIdx.x;
    int y = blockIdx.y * 32 + threadIdx.y;
    
    // 从全局内存加载到共享内存
    tile[threadIdx.y][threadIdx.x] = in[y * width + x];
    
    // 同步块内所有线程
    __syncthreads();
    
    // 转置写入
    int transX = blockIdx.y * 32 + threadIdx.x;
    int transY = blockIdx.x * 32 + threadIdx.y;
    out[transY * width + transX] = tile[threadIdx.x][threadIdx.y];
}

3.5 原子操作

__global__ void histogram(unsigned char *data, int *hist, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        unsigned char value = data[idx];
        // 原子加法，避免竞争条件
        atomicAdd(&hist[value], 1);
    }
}

3.6 流与异步执行

// 创建流
hipStream_t stream1, stream2;
hipStreamCreate(&stream1);
hipStreamCreate(&stream2);

// 在不同流中异步执行 kernel
kernel1<<<blocks, threads, 0, stream1>>>(...);
kernel2<<<blocks, threads, 0, stream2>>>(...);

// 流同步
hipStreamSynchronize(stream1);
hipStreamSynchronize(stream2);

// 销毁流
hipStreamDestroy(stream1);
hipStreamDestroy(stream2);

---

4. GPU 架构与性能优化

4.1 AMD GPU 架构概览

CDNA 架构（数据中心）：

• CDNA 3（MI300 系列）
• CDNA 2（MI250 系列）
• CDNA 1（MI100 系列）

RDNA 架构（游戏/工作站）：

• RDNA 3（RX 7000 系列）
• RDNA 2（RX 6000 系列）

4.2 性能优化原则

4.2.1 内存访问优化

// ❌ 低效：非合并访问
__global__ void inefficient(float *data) {
    int idx = threadIdx.x * blockDim.x + threadIdx.y;  // 跨步访问
    float val = data[idx];
}

// ✅ 高效：合并访问
__global__ void efficient(float *data) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;  // 连续访问
    float val = data[idx];
}

4.2.2 占用率优化

// 查询最大占用率
int minGridSize, blockSize;
hipOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize, 
                                   myKernel, 0, 0);

// 使用计算出的 block 大小启动 kernel
myKernel<<<minGridSize, blockSize>>>(...);

4.2.3 指令级优化

// 使用向量化类型
float4 a, b, c;
c.x = a.x + b.x;
c.y = a.y + b.y;
c.z = a.z + b.z;
c.w = a.w + b.w;

// 或使用内置向量操作
float4 c = make_float4(a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w);

4.3 使用 rocprof 性能分析

# 基本性能分析
rocprof --hip-trace --sys-trace ./my_application

# 生成性能报告
rocprof --stats ./my_application

# 分析特定 kernel
rocprof --kernel-trace -i input.txt ./my_application

4.4 性能调优检查清单

•  确保内存访问是合并的
•  最大化共享内存使用
•  减少主机 - 设备数据传输
•  使用异步流重叠计算和传输
•  优化 block 和 grid 大小
•  避免分支发散
•  使用合适的精度（FP16/FP8 用于 AI）

---

5. ROCm 库与工具

5.1 rocBLAS（线性代数）

#include <rocblas/rocblas.h>

// 矩阵乘法 C = A * B
void matrixMultiply(float *A, float *B, float *C, int m, int n, int k) {
    rocblas_handle handle;
    rocblas_create_handle(&handle);
    
    float alpha = 1.0f, beta = 0.0f;
    
    rocblas_sgemm(handle,
                  rocblas_operation_none, rocblas_operation_none,
                  m, n, k,
                  &alpha,
                  A, m,
                  B, k,
                  &beta,
                  C, m);
    
    rocblas_destroy_handle(handle);
}

5.2 rocFFT（快速傅里叶变换）

#include <rocfft/rocfft.h>

// 创建 FFT 计划
rocfft_plan plan;
rocfft_plan_create(&plan,
                   rocfft_calculation_type_complex_to_complex,
                   rocfft_array_type_complex_interleaved,
                   rocfft_array_type_complex_interleaved,
                   1, &length, 1, nullptr);

// 执行 FFT
rocfft_execute(plan, in_buffer, out_buffer, nullptr);

// 销毁计划
rocfft_plan_destroy(plan);

5.3 rocRAND（随机数生成）

#include <rocrand/rocrand.h>

rocrand_generator generator;
rocrand_create_generator(&generator, ROCRAND_RNG_PSEUDO_DEFAULT);

// 生成均匀分布随机数
float *randomNumbers;
hipMalloc(&randomNumbers, n * sizeof(float));
rocrand_generate_uniform(generator, randomNumbers, n);

rocrand_destroy_generator(generator);

5.4 MIOpen（机器学习）

#include <miopen/miopen.h>

// 创建卷积描述符
miopenConvolutionDescriptor_t convDesc;
miopenCreateConvolutionDescriptor(&convDesc);

// 配置卷积参数
miopenInitConvolutionDescriptor(convDesc,
                                miopenConvolution,
                                pad_h, pad_w,
                                stride_h, stride_w,
                                dilation_h, dilation_w);

// ... 执行卷积操作

5.5 开发工具

roc-gdb（调试器）：

# 启动调试器
roc-gdb ./my_application

# 常用命令
(gdb) break kernel_name
(gdb) run
(gdb) info threads
(gdb) thread apply all bt

Omniperf（性能分析）：

# 收集性能数据
python3 omniperf.py analyze --name my_kernel --kernel my_kernel \
  --rocm-path /opt/rocm -- ./my_application

---

6. AI/深度学习应用

6.1 PyTorch on ROCm

安装：

# 使用 pip 安装
pip3 install torch torchvision torchaudio \
  --index-url https://download.pytorch.org/whl/rocm6.0

# 或使用 Docker
docker pull rocm/pytorch:latest

验证：

import torch

print(f"PyTorch version: {torch.__version__}")
print(f"CUDA available: {torch.cuda.is_available()}")
print(f"CUDA device count: {torch.cuda.device_count()}")
print(f"Current device: {torch.cuda.current_device()}")
print(f"Device name: {torch.cuda.get_device_name(0)}")

# 创建张量并移动到 GPU
x = torch.randn(3, 3).cuda()
y = torch.randn(3, 3).cuda()
z = x @ y  # 矩阵乘法
print(z)

6.2 TensorFlow on ROCm

安装：

pip install tensorflow-rocm

验证：

import tensorflow as tf

print(f"TensorFlow version: {tf.__version__}")
print(f"GPUs available: {tf.config.list_physical_devices('GPU')}")

# 创建简单模型
model = tf.keras.Sequential([
    tf.keras.layers.Dense(64, activation='relu'),
    tf.keras.layers.Dense(10, activation='softmax')
])

model.compile(optimizer='adam', loss='sparse_categorical_crossentropy')

6.3 vLLM 推理部署

# 安装 vLLM for ROCm
pip install vllm

# 运行模型推理
python -m vllm.entrypoints.api_server \
  --model meta-llama/Llama-2-7b-chat-hf \
  --port 8000 \
  --tensor-parallel-size 1

6.4 AI 开发者教程资源

官方教程库： https://github.com/ROCm/gpuaidev

教程分类：

• 推理：vLLM、SGLang、Ollama 部署
• 微调：LoRA、GRPO、Unsloth
• 预训练：Megatron-LM、TorchTitan
• 优化：量化、内核开发、性能分析

---

7. 实战项目

7.1 项目 1：向量加法（入门）

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

#define N 1000000

__global__ void vectorAdd(float *a, float *b, float *c) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        c[i] = a[i] + b[i];
    }
}

int main() {
    float *h_a, *h_b, *h_c;
    float *d_a, *d_b, *d_c;
    
    // 分配主机内存
    h_a = (float*)malloc(N * sizeof(float));
    h_b = (float*)malloc(N * sizeof(float));
    h_c = (float*)malloc(N * sizeof(float));
    
    // 初始化数据
    for (int i = 0; i < N; i++) {
        h_a[i] = i * 1.0f;
        h_b[i] = i * 2.0f;
    }
    
    // 分配设备内存
    hipMalloc(&d_a, N * sizeof(float));
    hipMalloc(&d_b, N * sizeof(float));
    hipMalloc(&d_c, N * sizeof(float));
    
    // 复制数据到设备
    hipMemcpy(d_a, h_a, N * sizeof(float), hipMemcpyHostToDevice);
    hipMemcpy(d_b, h_b, N * sizeof(float), hipMemcpyHostToDevice);
    
    // 启动 kernel
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c);
    
    // 复制结果回主机
    hipMemcpy(h_c, d_c, N * sizeof(float), hipMemcpyDeviceToHost);
    
    // 验证结果
    printf("Result[0] = %f (expected: 0.0)\n", h_c[0]);
    printf("Result[999] = %f (expected: 2997.0)\n", h_c[999]);
    
    // 清理
    hipFree(d_a);
    hipFree(d_b);
    hipFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);
    
    return 0;
}

7.2 项目 2：矩阵乘法（中级）

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

#define TILE_SIZE 32

__global__ void matrixMulShared(float *A, float *B, float *C, int width) {
    __shared__ float As[TILE_SIZE][TILE_SIZE];
    __shared__ float Bs[TILE_SIZE][TILE_SIZE];
    
    int bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;
    
    int row = by * TILE_SIZE + ty;
    int col = bx * TILE_SIZE + tx;
    
    float sum = 0.0f;
    
    for (int t = 0; t < (width + TILE_SIZE - 1) / TILE_SIZE; t++) {
        // 加载数据到共享内存
        if (row < width && t * TILE_SIZE + tx < width)
            As[ty][tx] = A[row * width + t * TILE_SIZE + tx];
        else
            As[ty][tx] = 0.0f;
            
        if (t * TILE_SIZE + ty < width && col < width)
            Bs[ty][tx] = B[(t * TILE_SIZE + ty) * width + col];
        else
            Bs[ty][tx] = 0.0f;
        
        __syncthreads();
        
        // 计算部分和
        for (int k = 0; k < TILE_SIZE; k++)
            sum += As[ty][k] * Bs[k][tx];
        
        __syncthreads();
    }
    
    if (row < width && col < width)
        C[row * width + col] = sum;
}

int main() {
    int width = 1024;
    int size = width * width * sizeof(float);
    
    // ... 内存分配和初始化（类似向量加法）
    
    dim3 threadsPerBlock(TILE_SIZE, TILE_SIZE);
    dim3 blocksPerGrid((width + TILE_SIZE - 1) / TILE_SIZE,
                       (width + TILE_SIZE - 1) / TILE_SIZE);
    
    matrixMulShared<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, width);
    
    // ... 结果验证和清理
}

7.3 项目 3：神经网络推理（高级）

# simple_nn_inference.py
import torch
import torch.nn as nn
import time

class SimpleNN(nn.Module):
    def __init__(self, input_size, hidden_size, output_size):
        super(SimpleNN, self).__init__()
        self.fc1 = nn.Linear(input_size, hidden_size)
        self.relu = nn.ReLU()
        self.fc2 = nn.Linear(hidden_size, output_size)
        self.softmax = nn.Softmax(dim=1)
    
    def forward(self, x):
        x = self.fc1(x)
        x = self.relu(x)
        x = self.fc2(x)
        return self.softmax(x)

# 创建模型
model = SimpleNN(784, 256, 10)

# 移动到 GPU
device = torch.device("cuda:0" if torch.cuda.is_available() else "cpu")
model.to(device)

# 创建随机输入
batch_size = 32
input_data = torch.randn(batch_size, 784).to(device)

# 预热
for _ in range(10):
    _ = model(input_data)

# 性能测试
torch.cuda.synchronize()
start = time.time()

for _ in range(100):
    output = model(input_data)

torch.cuda.synchronize()
end = time.time()

print(f"Average inference time: {(end - start) / 100 * 1000:.2f} ms")
print(f"Throughput: {batch_size / ((end - start) / 100):.2f} samples/sec")

---

8. 调试与故障排除

8.1 常见错误及解决方案

错误 1：hipErrorNoDevice

原因： 未检测到 GPU 设备

解决方案：

# 检查用户组
groups $USER

# 添加用户到必要组
sudo usermod -a -G video,render $USER

# 重启或重新登录
# 验证设备
rocminfo

错误 2：hipErrorOutOfMemory

原因： GPU 显存不足

解决方案：

// 检查可用显存
size_t free, total;
hipMemGetInfo(&free, &total);
printf("Free: %zu MB, Total: %zu MB\n", free / 1024 / 1024, total / 1024 / 1024);

// 减少 batch size 或模型大小
// 使用梯度检查点
// 启用混合精度训练

错误 3：内核执行失败

原因： Kernel 配置错误或资源不足

调试方法：

// 检查 kernel 启动错误
vectorAdd<<<blocks, threads>>>(...);
hipError_t err = hipGetLastError();
if (err != hipSuccess) {
    printf("Kernel launch error: %s\n", hipGetErrorString(err));
}

// 同步并检查执行错误
err = hipDeviceSynchronize();
if (err != hipSuccess) {
    printf("Kernel execution error: %s\n", hipGetErrorString(err));
}

8.2 调试工具

使用 roc-gdb：

# 编译时添加调试信息
hipcc -g -G -o my_app my_app.hip

# 启动调试器
roc-gdb ./my_app

# 常用命令
(gdb) break myKernel
(gdb) run
(gdb) info threads
(gdb) thread 1
(gdb) frame
(gdb) print variable

使用 HIP Trace：

# 启用 HIP 跟踪
export HIP_TRACE=1
export HIP_TRACE_FILE=trace.log

# 运行程序
./my_app

# 分析跟踪文件
cat trace.log

8.3 性能问题排查

# 1. 使用 rocprof 分析
rocprof --stats ./my_app

# 2. 检查 GPU 利用率
rocm-smi --showuse

# 3. 监控温度
rocm-smi --showtemp

# 4. 查看显存使用
rocm-smi --showmeminfo vram

---

附录 A：快速参考卡

HIP vs CUDA 语法对比

功能	CUDA	HIP
头文件	#include <cuda_runtime.h>	#include <hip/hip_runtime.h>
内存分配	cudaMalloc	hipMalloc
内存复制	cudaMemcpy	hipMemcpy
内核启动	kernel<<<blocks, threads>>>()	kernel<<<blocks, threads>>>()
同步	cudaDeviceSynchronize()	hipDeviceSynchronize()
错误检查	cudaGetLastError()	hipGetLastError()

常用编译选项

# 基本编译
hipcc -o app app.hip

# 优化编译
hipcc -O3 -o app app.hip

# 调试编译
hipcc -g -G -o app app.hip

# 指定 GPU 架构
hipcc --amdgpu-target=gfx908,gfx90a -o app app.hip

# 显示编译命令
hipcc -v -o app app.hip

性能优化速查

✅ 合并内存访问
✅ 最大化共享内存使用
✅ 减少主机 - 设备传输
✅ 使用异步流
✅ 优化 block/grid 大小
✅ 避免分支发散
✅ 使用合适精度

❌ 随机内存访问
❌ 过度使用全局内存
❌ 频繁数据传输
❌ 同步阻塞
❌ 小 block 大小
❌ 复杂条件分支
❌ 不必要的双精度

---

祝你 GPU 编程之旅顺利！🚀