打开浏览器，直奔 amd.com/developercloud，右上角「Get Started」。邮箱验证后，选「AI & HPC」→「GPU」→「MI210」即可。填项目描述时写一句「ROCm 入门实验」就能秒过。审核邮件通常在 10 min 内抵达，点邮件里的「Launch Portal」就能看见控制台。

第一次登录会提示绑定 GitHub，方便后续用公钥免密。控制台左侧「Instances」→「Create」：

• Image: ROCm 5.7 Ubuntu 22.04
• GPU: 1×MI210 64 GB
• CPU: 16 vCPU
• Disk: 100 GB
• SSH Key: 选刚才导入的 id_rsa.pub

点「Start」后 30 秒拿到公网 IP，本机执行

ssh ubuntu@<ip> -i ~/.ssh/id_rsa

就能进系统。rocm-smi 一下，看到 GFX90A 温度 38 ℃，心里踏实一半。

2. 启动 ROCm 容器：别在宿主机装驱动

官方镜像已经把驱动、编译器、调试器全打包好，宿主机只负责开机。执行：

docker run -it --device=/dev/kfd --device=/dev/dri \
           --group-add video --ipc=host \
           rocm/rocm-terminal:5.7 bash

容器里再 rocm-smi，输出和宿主机一致，说明直通成功。此时 apt update && apt install -y rocprofiler-dev 把火焰图工具装好，后面直接复用。

3. 向量加法 Hello World：25 行代码跑通

新建 vecadd.cpp：

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

__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 << 24;               // 16 M 元素
    size_t bytes = N * sizeof(float);
    float *dA, *dB, *dC;
    hipMalloc(&dA, bytes);         // 1. 与 cudaMalloc 签名完全一致
    hipMalloc(&dB, bytes);
    hipMalloc(&dC, bytes);

    // 初始化 Host 数据
    float *hA = new float[N], *hB = new float[N];
    for (int i = 0; i < N; ++i) { hA[i] = 1.0f; hB[i] = 2.0f; }

    hipMemcpy(dA, hA, bytes, hipMemcpyHostToDevice);
    hipMemcpy(dB, hB, bytes, hipMemcpyHostToDevice);

    int block = 256;
    int grid  = (N + block - 1) / block;
    hipLaunchKernelGGL(vecAdd, dim3(grid), dim3(block), 0, 0, dA, dB, dC, N);

    float *hC = new float[N];
    hipMemcpy(hC, dC, bytes, hipMemcpyDeviceToHost);

    printf("Last element = %f\n", hC[N-1]);   // 期望 3.0
    delete[] hA; delete[] hB; delete[] hC;
    hipFree(dA); hipFree(dB); hipFree(dC);
    return 0;
}

编译：

hipcc vecadd.cpp -o vecadd -O3

运行 ./vecadd，终端蹦出 Last element = 3.000000，世界安静。

4. hipMalloc 与 hipLaunchKernelGGL：CUDA 用户一眼懂

• hipMalloc 就是 cudaMalloc 的宏别名，参数顺序、字节单位都一样，直接替换。
• hipLaunchKernelGGL(kernel, g, b, s, stream, ...) 对应 CUDA 的 kernel<<<g,b,s,stream>>>(...)。宏展开后帮你把配置参数塞进 hipConfigureCall，所以写起来比 <<< >>> 啰嗦一点，却省掉了 __global__ 函数指针的强制转换。
• 线程索引 API 仍是 threadIdx.x / blockIdx.x，但头文件改成 <hip/hip_runtime.h>，老代码批量 sed 即可。

5. 用 rocprof 画火焰图：30 秒定位瓶颈

先让程序跑慢一点，把 N 调到 1 << 26，重新编译。接着：

rocprof --stats --timestamp on -o vectrace.csv ./vecadd

跑完会生成 vectrace.csv 和 results.stats.csv。我们用 rocprofv 把 CSV 转火焰图：

apt install -y python3-pip
pip install rocprofiler-python
rocprofv import vectrace.csv
rocprofv flamegraph -o vec.svg

把 vec.svg 拉到本地浏览器，一眼看到 vecAdd 占 98 % 时间，内存拷贝几乎不可见——说明瓶颈在计算而非 PCIe，后续可放心上优化 kernel。

6. VSCode 远程：边写边调

本地 VSCode 装 Remote-SSH 插件，点左下角「><」→「Connect to Host…」→ 填 ubuntu@<ip>。连上后打开 /workspace，装 ms-vscode.cpptools 与 amd-llvm.vscode-hip，F5 直接调用容器内 gdb-hip，断点可跳进 __global__ 函数。截图如下（图略，读者按步骤即可复现）。

7. 常见编译报错对照表

报错原文	原因	一键修复
error: unknown type name '__global__'	文件后缀 .c 或没 include <hip/hip_runtime.h>	改 .cpp 并确认头文件
use of undeclared identifier 'hipMalloc'	没链接 HIP runtime	用 hipcc 而不是 g++
invalid device function	架构不匹配	export HIP_PLATFORM=hcc 或加 --offload-arch=gfx90a
rocprof: No kernels found	程序运行太短	把数据量调大或加 --basetime on
Permission denied /dev/kfd	容器没加 --group-add video	重新 docker run 带参数

把上表贴到 README，基本能 cover 90 % 初阶踩坑。

8. 小结：30 分钟后的下一步

账号、容器、编译、采样、火焰图，一条线下来刚好半小时。接下来你可以：

• 把 vecAdd 换成 GEMM，试 rocblas；
• 用 hipify 把旧 CUDA 项目整体迁移；
• 在 ModelScope 创空间打包镜像，让伙伴直接点「一键复现」。

MI210 的 64 GB 显存足够 7 B 模型推理，火焰图告诉你时间花在哪，剩下的就是算法与内存排布的艺术。祝玩得开心，记得把踩到的新坑再写成博客，循环开源。