ROCm 快速上手从跑通第一个 HIP 核函数到把算力用满的完整路径【免费下载链接】legacy-rocm-buildAMD ROCm™ Software - GitHub Home项目地址: https://gitcode.com/GitHub_Trending/ro/legacy-rocm-build训练任务被 CPU 卡住、GPU 机时费越涨越狠AMD 的开源 GPU 加速计算平台 ROCm 是一个现实的备选它覆盖驱动到 AI 框架的全家桶让深度学习与 HPC 任务在 AMD GPU 上高效落地。这篇文章带你走完安装 → 跑通 → 调优全流程并讲清每一步为什么这样设计。三步装好在 Ubuntu 上落地 ROCm 这一节回答一个问题怎么最快把 ROCm 装到机器上并且当场验证可用。先对一遍前置条件免得返工系统Ubuntu 22.04/24.04、RHEL 8/9、SLES 15 SP5 等主流发行版GPUInstinct MI 系列或支持列表内的 Radeon 卡资源16GB 内存磁盘预留 20GB 以上# 1. 导入官方 GPG 密钥并配置 apt 源否则 apt 认不到 ROCm 包 wget -qO - https://repo.radeon.com/rocm/rocm.gpg.key | gpg --dearmor | sudo tee /etc/apt/trusted.gpg.d/rocm.gpg echo deb [archamd64] https://repo.radeon.com/rocm/apt/apt/ stable main | sudo tee /etc/apt/sources.list.d/rocm.list # 2. 只装 HIP SDK含 hipcc 编译器与 HIP 运行时够用且最轻 sudo apt update sudo apt install -y rocm-hip-sdk # 3. 把当前用户加入 video/render 组否则程序打不开 GPU sudo usermod -aG video,render $USER执行完注销再登录让用户组生效。ROCm Runfile Installer 的主菜单无网络、无包管理器的内网环境可以走这条路离线安装仓库docs/install/里有完整的 runfile 安装截图流程。装完立刻验证三条命令确认 ROCm 安装就绪# ① 你的 GPU 对应 gfx target 应出现在输出里 rocminfo | grep -i gfx # ② 确认 hipcc 编译器可用 hipcc --version # ③ 编译运行官方 sample 做冒烟测试 cd /opt/rocm/share/hip/samples/0_Intro/bit_extract make ./bit_extractbit_extract能正常打印结果说明 ROCm 安装已经能接住后面的事。环境就绪先别急着写代码花两分钟看懂它为什么快。计算单元怎么把数据嚼碎这一节回答一个问题同样的算法搬上 AMD GPU 凭什么能快一个量级。把一个计算单元CU想象成一个小车间顶上的调度器是工头不停派发任务4 个 SIMD 单元是四条流水线一条指令同时处理几十份数据这是 GPU 吞吐的主来源标量单元干标量与控制的细活每条流水线配大容量寄存器文件整个 CU 还共享一块叫 LDSLocal Data Share的高速本地内存。单个 CU 的内部构成调度器居上4 个 SIMD 居中LDS 与 VGPR/SGPR 寄存器文件在底部。本质是海量线程共用同一条指令流数据尽量从寄存器和 LDS 里抓几乎不等慢速的全局内存。你写的核函数只要让每个线程都有活干、少等待就天然贴合这个结构。看懂结构就该动手写代码了。两个 HIP 核函数向量加法到分块矩阵乘本节是核心实战写两个 HIP 核函数从向量加法一路到矩阵乘法全部能编译、能运行。新建vector_add.cpp你只定义每个线程干什么其余交给平台#include hip/hip_runtime.h #include iostream #include vector #include algorithm // 核函数每个线程负责一个元素的加法 __global__ void vecAdd(const float* a, const float* b, float* c, int n) { int idx hipBlockIdx_x * hipBlockDim_x hipThreadIdx_x; if (idx n) c[idx] a[idx] b[idx]; } int main() { const int n 1024; std::vectorfloat hA(n, 1.0f), hB(n, 2.0f), hC(n, 0.0f); float *dA, *dB, *dC; hipMalloc(dA, n * sizeof(float)); // 申请设备侧显存 hipMalloc(dB, n * sizeof(float)); hipMalloc(dC, n * sizeof(float)); hipMemcpy(dA, hA.data(), n * sizeof(float), hipMemcpyHostToDevice); hipMemcpy(dB, hB.data(), n * sizeof(float), hipMemcpyHostToDevice); // 4 块 x 256 线程 1024 线程恰好覆盖 n 个元素 hipLaunchKernelGGL(vecAdd, dim3(4), dim3(256), 0, 0, dA, dB, dC, n); hipDeviceSynchronize(); // 阻塞等 GPU 算完否则读到的还是旧值 hipMemcpy(hC.data(), dC, n * sizeof(float), hipMemcpyDeviceToHost); bool ok std::all_of(hC.begin(), hC.end(), [](float v){ return v 3.0f; }); std::cout (ok ? SUCCESS: all 3.0 : FAIL) std::endl; hipFree(dA); hipFree(dB); hipFree(dC); return ok ? 0 : 1; }编译运行预期输出SUCCESS: all 3.0hipcc -O2 -o vec_add vector_add.cpp ./vec_addAPI 只需记住四件事hipMalloc申请显存hipMemcpy负责主机与设备间数据搬运hipLaunchKernelGGL指定块数、线程数和参数来启动核函数hipDeviceSynchronize阻塞到 GPU 忙完。向量加法只是热身矩阵乘法才是暴露内存瓶颈的典型。朴素写法让每个线程直接从全局内存读 A 和 B同一份数据会被反复从 DRAM 拉取几千次算力被带宽拖死。标准解法是共享内存分块#define TILE 16 __global__ void matmulTiled(const float* A, const float* B, float* C, int M, int N, int K) { __shared__ float tileA[TILE][TILE]; // LDS块内共享的工作台 __shared__ float tileB[TILE][TILE]; // 1) 线程协作把一片 tile 从全局内存搬进 LDS用前 __syncthreads() 同步 // 2) 沿 K 循环在寄存器里累加把全局读取摊薄到 1/TILE // 3) 最后才写回 C减少显存写次数 }before/after 差在哪朴素版每线程各自去 DRAM 捞数据分块版捞一次、块内共享一次全局读取被摊薄 16 次以上。还有一点常被忽略——HIP 核函数写法与 CUDA 同构同一份源码可在 AMD 与 NVIDIA 双平台编译仓库工具链里的 HIPIFY 还能把既有 CUDA 工程自动改写迁移将来换平台不必重写计算逻辑。算子层面的地基打好了该把活儿交给框架。两条 pip 命令让框架认出你的显卡这一节解决上层框架怎么无感切换。ROCm PyTorch 的安装就两条命令pip3 install torch torchvision --index-url https://download.pytorch.org/whl/rocm6.2 python3 -c import torch; print(torch.__version__, torch.cuda.is_available())第二行打印出版本号和True说明 PyTorch 已自动认出 AMD GPU——框架里照旧用torch.cuda这套 API代码不用改。TensorFlow 同理pip3 install tensorflow-rocm后tf.config.list_physical_devices(GPU)返回非空即可。RCCL 八卡集合通信测试多卡通信由 RCCL 集合通信库兜底分布式训练前用它验证卡间带宽是标准动作。多机多卡场景下集合通信走 RCCL单机内拓扑关系则靠 SMI 工具查清。代码能跑不等于跑得快下面三个杠杆决定吞吐上限。三个杠杆把瓶颈换成速度 这一节给出可量化的性能调优方向逐个拉效果立现。杠杆一异步化数据搬运。主机与设备之间的拷贝往往是最大的隐形开销。把hipMemcpy换成hipMemcpyAsync交给 stream让拷贝和核函数计算并行推进hipStream_t s; hipStreamCreate(s); hipMemcpyAsync(dA, hA, bytes, hipMemcpyHostToDevice, s); // 核函数挂在同一 stream 上自然排队拷贝延迟被计算掩盖 hipStreamSynchronize(s);配合上一节的 LDS 分块减少全局内存访问两者叠加拷贝开销基本被藏掉。杠杆二核函数调度与占用率。块大小、tile 尺寸、寄存器压力都影响每个 CU 能同时容纳多少工作组。原则和 CU 结构一致让线程尽量从寄存器与 LDS 抓数据别给 SIMD 流水线留空闲。占用率与寄存器VGPR用量的关系寄存器用太多会压低占用率是调 tile 尺寸时的核心约束。杠杆三用剖析工具量别猜。拿不准就先量化rocprof --stats ./vec_add # 核函数级耗时与统计 rocm-smi --showtopo # 卡间链路类型、跳数与 NUMA 亲和rocm-smi --showtopo实际输出GPU 之间的 XGMI/PCIe 链路类型、跳数与 NUMA 亲和是多卡调度绑核的直接依据。瓶颈定位清楚了生产部署也就顺理成章——但生产部署难免出状况。30 秒排查三类高频 ROCm 故障这一节把定位问题变成三行命令的事。排查 ① GPU 未识别lspci | grep -iE amd|ati # 预期出现 AMD/ATI 的 3D controller rocminfo | grep -i gfx # 预期看到对应 gfx target为空优先查驱动与用户组排查 ② HIP 编译报错which hipcc hipcc --version # 预期路径指向 /opt/rocm/hip/bin/hipcc 并打印版本 source /opt/rocm/setenv.sh # 找不到 hipcc 时先补环境变量再试排查 ③ 显存不足out of memoryrocm-smi --showmeminfo vram # 预期逐卡列出 VRAM Free/Total一眼分辨真占满还是碎片再备两件深挖工具rocgdb ./vec_add断点单步看核函数启动路径rocprof -i input.txt -o trace.csv ./vec_add生成 trace 数据导入剖析工具。环境、代码、性能三条线都理顺后值得知道仓库里还能往哪挖。往哪深挖仓库关键目录与下一步清单回头看这条路径装环境、懂并行模型、写核函数、接框架、把算力抠满这就是 ROCm 上 GPU 加速计算从 0 到跑通的完整路线。仓库里这几个目录是继续深入的入口安装指南docs/install/含 runfile 离线安装全流程截图在docs/data/how-to/GPU 架构参考docs/reference/gpu-arch/CDNA/RDNA 各代结构图系统优化指南docs/reference/system-optimization/NUMA 亲和、GPU 隔离等机器已经热好了下一步挑一个就能开始git clone https://gitcode.com/GitHub_Trending/ro/legacy-rocm-build拉下最新文档与示例跑通share/hip/samples/下的全部示例重点看多卡那几个用不同规模矩阵乘法产出一张 CPU vs GPU 耗时对照表验证你机器上的加速比用 RCCL 实测一下卡间带宽为分布式训练做预检【免费下载链接】legacy-rocm-buildAMD ROCm™ Software - GitHub Home项目地址: https://gitcode.com/GitHub_Trending/ro/legacy-rocm-build创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 
长河君 · 2026/8/24 23:03:00 · 阅读 1.2k ROCm 快速上手：从跑通第一个 HIP 核函数到把算力用满的完整路径
ROCm 快速上手从跑通第一个 HIP 核函数到把算力用满的完整路径【免费下载链接】legacy-rocm-buildAMD ROCm™ Software - GitHub Home项目地址: https://gitcode.com/GitHub_Trending/ro/legacy-rocm-build训练任务被 CPU 卡住、GPU 机时费越涨越狠AMD 的开源 GPU 加速计算平台 ROCm 是一个现实的备选它覆盖驱动到 AI 框架的全家桶让深度学习与 HPC 任务在 AMD GPU 上高效落地。这篇文章带你走完安装 → 跑通 → 调优全流程并讲清每一步为什么这样设计。三步装好在 Ubuntu 上落地 ROCm 这一节回答一个问题怎么最快把 ROCm 装到机器上并且当场验证可用。先对一遍前置条件免得返工系统Ubuntu 22.04/24.04、RHEL 8/9、SLES 15 SP5 等主流发行版GPUInstinct MI 系列或支持列表内的 Radeon 卡资源16GB 内存磁盘预留 20GB 以上# 1. 导入官方 GPG 密钥并配置 apt 源否则 apt 认不到 ROCm 包 wget -qO - https://repo.radeon.com/rocm/rocm.gpg.key | gpg --dearmor | sudo tee /etc/apt/trusted.gpg.d/rocm.gpg echo deb [archamd64] https://repo.radeon.com/rocm/apt/apt/ stable main | sudo tee /etc/apt/sources.list.d/rocm.list # 2. 只装 HIP SDK含 hipcc 编译器与 HIP 运行时够用且最轻 sudo apt update sudo apt install -y rocm-hip-sdk # 3. 把当前用户加入 video/render 组否则程序打不开 GPU sudo usermod -aG video,render $USER执行完注销再登录让用户组生效。ROCm Runfile Installer 的主菜单无网络、无包管理器的内网环境可以走这条路离线安装仓库docs/install/里有完整的 runfile 安装截图流程。装完立刻验证三条命令确认 ROCm 安装就绪# ① 你的 GPU 对应 gfx target 应出现在输出里 rocminfo | grep -i gfx # ② 确认 hipcc 编译器可用 hipcc --version # ③ 编译运行官方 sample 做冒烟测试 cd /opt/rocm/share/hip/samples/0_Intro/bit_extract make ./bit_extractbit_extract能正常打印结果说明 ROCm 安装已经能接住后面的事。环境就绪先别急着写代码花两分钟看懂它为什么快。计算单元怎么把数据嚼碎这一节回答一个问题同样的算法搬上 AMD GPU 凭什么能快一个量级。把一个计算单元CU想象成一个小车间顶上的调度器是工头不停派发任务4 个 SIMD 单元是四条流水线一条指令同时处理几十份数据这是 GPU 吞吐的主来源标量单元干标量与控制的细活每条流水线配大容量寄存器文件整个 CU 还共享一块叫 LDSLocal Data Share的高速本地内存。单个 CU 的内部构成调度器居上4 个 SIMD 居中LDS 与 VGPR/SGPR 寄存器文件在底部。本质是海量线程共用同一条指令流数据尽量从寄存器和 LDS 里抓几乎不等慢速的全局内存。你写的核函数只要让每个线程都有活干、少等待就天然贴合这个结构。看懂结构就该动手写代码了。两个 HIP 核函数向量加法到分块矩阵乘本节是核心实战写两个 HIP 核函数从向量加法一路到矩阵乘法全部能编译、能运行。新建vector_add.cpp你只定义每个线程干什么其余交给平台#include hip/hip_runtime.h #include iostream #include vector #include algorithm // 核函数每个线程负责一个元素的加法 __global__ void vecAdd(const float* a, const float* b, float* c, int n) { int idx hipBlockIdx_x * hipBlockDim_x hipThreadIdx_x; if (idx n) c[idx] a[idx] b[idx]; } int main() { const int n 1024; std::vectorfloat hA(n, 1.0f), hB(n, 2.0f), hC(n, 0.0f); float *dA, *dB, *dC; hipMalloc(dA, n * sizeof(float)); // 申请设备侧显存 hipMalloc(dB, n * sizeof(float)); hipMalloc(dC, n * sizeof(float)); hipMemcpy(dA, hA.data(), n * sizeof(float), hipMemcpyHostToDevice); hipMemcpy(dB, hB.data(), n * sizeof(float), hipMemcpyHostToDevice); // 4 块 x 256 线程 1024 线程恰好覆盖 n 个元素 hipLaunchKernelGGL(vecAdd, dim3(4), dim3(256), 0, 0, dA, dB, dC, n); hipDeviceSynchronize(); // 阻塞等 GPU 算完否则读到的还是旧值 hipMemcpy(hC.data(), dC, n * sizeof(float), hipMemcpyDeviceToHost); bool ok std::all_of(hC.begin(), hC.end(), [](float v){ return v 3.0f; }); std::cout (ok ? SUCCESS: all 3.0 : FAIL) std::endl; hipFree(dA); hipFree(dB); hipFree(dC); return ok ? 0 : 1; }编译运行预期输出SUCCESS: all 3.0hipcc -O2 -o vec_add vector_add.cpp ./vec_addAPI 只需记住四件事hipMalloc申请显存hipMemcpy负责主机与设备间数据搬运hipLaunchKernelGGL指定块数、线程数和参数来启动核函数hipDeviceSynchronize阻塞到 GPU 忙完。向量加法只是热身矩阵乘法才是暴露内存瓶颈的典型。朴素写法让每个线程直接从全局内存读 A 和 B同一份数据会被反复从 DRAM 拉取几千次算力被带宽拖死。标准解法是共享内存分块#define TILE 16 __global__ void matmulTiled(const float* A, const float* B, float* C, int M, int N, int K) { __shared__ float tileA[TILE][TILE]; // LDS块内共享的工作台 __shared__ float tileB[TILE][TILE]; // 1) 线程协作把一片 tile 从全局内存搬进 LDS用前 __syncthreads() 同步 // 2) 沿 K 循环在寄存器里累加把全局读取摊薄到 1/TILE // 3) 最后才写回 C减少显存写次数 }before/after 差在哪朴素版每线程各自去 DRAM 捞数据分块版捞一次、块内共享一次全局读取被摊薄 16 次以上。还有一点常被忽略——HIP 核函数写法与 CUDA 同构同一份源码可在 AMD 与 NVIDIA 双平台编译仓库工具链里的 HIPIFY 还能把既有 CUDA 工程自动改写迁移将来换平台不必重写计算逻辑。算子层面的地基打好了该把活儿交给框架。两条 pip 命令让框架认出你的显卡这一节解决上层框架怎么无感切换。ROCm PyTorch 的安装就两条命令pip3 install torch torchvision --index-url https://download.pytorch.org/whl/rocm6.2 python3 -c import torch; print(torch.__version__, torch.cuda.is_available())第二行打印出版本号和True说明 PyTorch 已自动认出 AMD GPU——框架里照旧用torch.cuda这套 API代码不用改。TensorFlow 同理pip3 install tensorflow-rocm后tf.config.list_physical_devices(GPU)返回非空即可。RCCL 八卡集合通信测试多卡通信由 RCCL 集合通信库兜底分布式训练前用它验证卡间带宽是标准动作。多机多卡场景下集合通信走 RCCL单机内拓扑关系则靠 SMI 工具查清。代码能跑不等于跑得快下面三个杠杆决定吞吐上限。三个杠杆把瓶颈换成速度 这一节给出可量化的性能调优方向逐个拉效果立现。杠杆一异步化数据搬运。主机与设备之间的拷贝往往是最大的隐形开销。把hipMemcpy换成hipMemcpyAsync交给 stream让拷贝和核函数计算并行推进hipStream_t s; hipStreamCreate(s); hipMemcpyAsync(dA, hA, bytes, hipMemcpyHostToDevice, s); // 核函数挂在同一 stream 上自然排队拷贝延迟被计算掩盖 hipStreamSynchronize(s);配合上一节的 LDS 分块减少全局内存访问两者叠加拷贝开销基本被藏掉。杠杆二核函数调度与占用率。块大小、tile 尺寸、寄存器压力都影响每个 CU 能同时容纳多少工作组。原则和 CU 结构一致让线程尽量从寄存器与 LDS 抓数据别给 SIMD 流水线留空闲。占用率与寄存器VGPR用量的关系寄存器用太多会压低占用率是调 tile 尺寸时的核心约束。杠杆三用剖析工具量别猜。拿不准就先量化rocprof --stats ./vec_add # 核函数级耗时与统计 rocm-smi --showtopo # 卡间链路类型、跳数与 NUMA 亲和rocm-smi --showtopo实际输出GPU 之间的 XGMI/PCIe 链路类型、跳数与 NUMA 亲和是多卡调度绑核的直接依据。瓶颈定位清楚了生产部署也就顺理成章——但生产部署难免出状况。30 秒排查三类高频 ROCm 故障这一节把定位问题变成三行命令的事。排查 ① GPU 未识别lspci | grep -iE amd|ati # 预期出现 AMD/ATI 的 3D controller rocminfo | grep -i gfx # 预期看到对应 gfx target为空优先查驱动与用户组排查 ② HIP 编译报错which hipcc hipcc --version # 预期路径指向 /opt/rocm/hip/bin/hipcc 并打印版本 source /opt/rocm/setenv.sh # 找不到 hipcc 时先补环境变量再试排查 ③ 显存不足out of memoryrocm-smi --showmeminfo vram # 预期逐卡列出 VRAM Free/Total一眼分辨真占满还是碎片再备两件深挖工具rocgdb ./vec_add断点单步看核函数启动路径rocprof -i input.txt -o trace.csv ./vec_add生成 trace 数据导入剖析工具。环境、代码、性能三条线都理顺后值得知道仓库里还能往哪挖。往哪深挖仓库关键目录与下一步清单回头看这条路径装环境、懂并行模型、写核函数、接框架、把算力抠满这就是 ROCm 上 GPU 加速计算从 0 到跑通的完整路线。仓库里这几个目录是继续深入的入口安装指南docs/install/含 runfile 离线安装全流程截图在docs/data/how-to/GPU 架构参考docs/reference/gpu-arch/CDNA/RDNA 各代结构图系统优化指南docs/reference/system-optimization/NUMA 亲和、GPU 隔离等机器已经热好了下一步挑一个就能开始git clone https://gitcode.com/GitHub_Trending/ro/legacy-rocm-build拉下最新文档与示例跑通share/hip/samples/下的全部示例重点看多卡那几个用不同规模矩阵乘法产出一张 CPU vs GPU 耗时对照表验证你机器上的加速比用 RCCL 实测一下卡间带宽为分布式训练做预检【免费下载链接】legacy-rocm-buildAMD ROCm™ Software - GitHub Home项目地址: https://gitcode.com/GitHub_Trending/ro/legacy-rocm-build创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考