AMD ROCm 从装好到跑通5 个动作装好环境、写出第一个 GPU 核函数【免费下载链接】legacy-rocm-buildAMD ROCm™ Software - GitHub Home项目地址: https://gitcode.com/GitHub_Trending/ro/legacy-rocm-build训练任务被 CPU 单核拖住、NVIDIA 的卡要排队、手上的 AMD 显卡却在吃灰——这种局面值得动一动了。AMD ROCm 就是 AMD 的开源 GPU 加速计算平台从内核驱动、数学库、HIP 运行时到 PyTorch 对接一层层配齐让深度学习与 HPC 任务直接落在 AMD GPU 上跑。这篇文章按装好 → 跑通 → 调快的顺序给你一份能直接复制执行的完整方案。ROCm 的覆盖范围Instinct、Radeon 到八卡互联先看它管得住多大的地盘。ROCm 支持的硬件分三档Instinct 数据中心卡MI300 系列是主力、Radeon 独显、Ryzen AI 的 APU支持的操作系统以主流 Linux 发行版为主Ubuntu 22.04/24.04、RHEL、SLES 等。能力上覆盖单卡推理、多卡并行训练和分布式集合通信最上层能直接对接 PyTorch 这类框架。AMD MI300X 节点架构8 张 Instinct 卡经 Infinity Fabric 全互联是 ROCm 上多卡并行训练的基本单元。知道它管到哪就开工装它。Ubuntu 上四条命令装好 ROCm附两个成功标志以 Ubuntu 24.04 为例apt 路线总共四步全部可复制粘贴。第一步先备好 ROCm 官方 apt 源和两个基础依赖库# 装依赖libatomic1/libquadmath0 是 ROCm 工具链运行时的硬需求缺了会报符号错误 sudo apt update sudo apt install -y libatomic1 libquadmath0 # 导入 ROCm 发布签名密钥并注册官方 apt 源 sudo mkdir -p --mode0755 /etc/apt/keyrings wget -O - https://repo.amd.com/rocm/packages-multi-arch/gpg/rocm.gpg | gpg --dearmor | sudo tee /etc/apt/keyrings/amdrocm.gpg /dev/null sudo tee /etc/apt/sources.list.d/rocm.list EOF deb [archamd64 signed-by/etc/apt/keyrings/amdrocm.gpg] https://repo.amd.com/rocm/packages-multi-arch/ubuntu2404 stable main EOF sudo apt update然后安装、放权、验证# 2. 安装 HIP SDKhipcc 编译器 HIP 运行时GPU 编程的最小可用集 sudo apt install -y rocm-hip-sdk # 3. 把当前用户加入 video/render 设备组否则普通用户程序打不开 GPU sudo usermod -aG video,render $USER # 4. 验证两条命令都要有输出才算装好 rocminfo | grep -E Marketing Name|gfx amd-smi version成功标志有两个rocminfo的 HSA Agents 段落能列出你显卡的 gfx 代号与商品名例如AMD Instinct MI350Xamd-smi version打印出 ROCm 版本与 amdgpu 驱动版本。第 3 步要注销重登或重启才生效。apt 不方便时离线机器、想连驱动一起装官方还提供 runfile 安装器ROCm 与 amdgpu 驱动打包在一起自带图形化引导菜单# 下载约 4GB 的一体化安装器双击或终端进入菜单逐项选择 wget https://repo.radeon.com/rocm/installer/rocm-runfile-installer/rocm-rel-7.14/rocm-installer-7.14.0-7.run chmod x rocm-installer-7.14.0-7.run ./rocm-installer-7.14.0-7.runROCm runfile 安装器主菜单ROCm 与 amdgpu 驱动打成一包支持离线安装与图形化选择。提醒一句系统里已有 ROCm 7.2.4 或更老版本的话先卸载干净再装新版混装会留下难查的库冲突。环境有了下一个动作就是写你自己的第一个核函数。第一个核函数SAXPYhipcc 一编就跑下面这段代码实现数值计算里的经典例程 SAXPYy a*x y。它不是玩具PyTorch 里成千上万个算子底层干的就是这类每个线程处理一个元素的活。把下面内容存成saxpy.cpp#include hip/hip_runtime.h #include iostream #include vector // 核函数一个线程只负责一个元素的 y[i] a*x[i] y[i] __global__ void saxpy(const float* x, const float* y, float* out, int n) { int i hipBlockIdx_x * hipBlockDim_x hipThreadIdx_x; if (i n) out[i] 4.0f * x[i] y[i]; } int main() { const int n 1 20; // 1048576 个元素规模足够让 GPU 忙起来 std::vectorfloat hX(n, 1.0f), hY(n, 2.0f), hO(n, 0.0f); float *dX, *dY, *dO; hipMalloc(dX, n * sizeof(float)); // 在设备显存上分配空间 hipMalloc(dY, n * sizeof(float)); hipMalloc(dO, n * sizeof(float)); hipMemcpy(dX, hX.data(), n * sizeof(float), hipMemcpyHostToDevice); hipMemcpy(dY, hY.data(), n * sizeof(float), hipMemcpyHostToDevice); // 2048 个线程块 x 256 线程共 524288 个线程同时开工 hipLaunchKernelGGL(saxpy, dim3(2048), dim3(256), 0, 0, dX, dY, dO, n); hipDeviceSynchronize(); // 核函数是异步启动的必须等它干完再取结果 hipMemcpy(hO.data(), dO, n * sizeof(float), hipMemcpyDeviceToHost); // 正确性自检4*126逐项验证 bool ok true; for (int i 0; i n; i) if (hO[i] ! 6.0f) { ok false; break; } std::cout (ok ? SAXPY 校验通过全部 1M 个元素等于 6.0 : 校验失败) std::endl; hipFree(dX); hipFree(dY); hipFree(dO); return ok ? 0 : 1; }编译运行只需两行hipcc -O2 -o saxpy saxpy.cpp ./saxpy预期输出SAXPY 校验通过全部 1M 个元素等于 6.0。想针对自己的卡优化先用rocminfo | grep gfx查出架构代号再补上参数例如hipcc -O2 -archgfx90a -o saxpy saxpy.cpp。跑通了只是及格线它到底凭什么快值得花三分钟搞清楚。为什么快把全局内存访问压进片上抽屉把 CUCompute Unit计算单元想成一家小型印刷厂一条 SIMD 流水线是一台机头同一张印版一条指令能同时印几十份页几十条数据这就是 GPU 吞吐的来源VGPR/SGPR 寄存器是机头旁的便签随手可取LDSLocal Data Share片上共享内存是厂里共享的印版台同一批工人围着它干不用再跑一趟仓库。单个 CU 的内部构成调度器在上SIMD 流水线居中LDS 与寄存器文件在底部——核函数优化基本都是围着这几块做文章。GPU 快就快在这条纪律让海量线程共用同一条指令流数据尽量从寄存器和 LDS 里取少碰全局显存。显存带宽有限每次从 DRAM 取数都要排队而同一批数据被多个线程反复取正是 GPU 程序里最常见的浪费。分块tiling优化干的就是这件事——数据取一次进 LDS块内大家共享一次全局读被摊薄几十次。代价呢寄存器是硬通货。每个线程分到的 VGPR 越多一张卡上同时驻留的线程就越少这叫占用率occupancy。占用率掉下去流水线就没有足够线程切换来填满等待时间机器白空转。寄存器占用与占用率的关系VGPR 用量决定每个计算单元能驻留多少线程是调核函数时的核心约束。单卡的故事讲完接下来把你的 Python 代码也搬上来。PyTorch 接入 ROCm两行 pip 加一次验证算子层的地基有了把活交给框架。ROCm 版 PyTorch 的安装是标准 pip 操作只是把索引源换成 ROCm 的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就成了。注意 ROCm 里照旧用torch.cuda这套 API名字是 CUDA 历史遗留代码一行不用改。纯 Python 工作流不想动系统环境还有另一条路用虚拟环境装打包好的 ROCm SDKpython -m venv .venv source .venv/bin/activate python -m pip install --index-url https://repo.amd.com/rocm/whl-multi-arch/ rocm[libraries,device-gfx950]7.14.0 pip freeze | grep rocm # 预期能列出 rocm-sdk 等已装包device-gfx950要换成rocminfo里你显卡的架构代号对应关系仓库的兼容矩阵里都有。框架能跑之后多卡就顺理成章了。多卡互联怎么验证一次 RCCL 集合通信测试八张卡挂在同一节点上数据要在卡与卡之间倒手带宽不够一切白搭。RCCLROCm 的集合通信库对标 NCCL是这条链路的底座而rccl-tests里的all_reduce就是它的体检表./rccl-tests/build/all_reduce_perf -b 8 -e 256M -f 2 -g 8八卡 all_reduce 集合通信测试输出数据规模、总线带宽busbw与延迟一目了然是分布式训练前验证卡间链路的常用手段。看两列就够busbw总线带宽应接近该卡规格里的互联峰值逐档数据量都稳跳数大、带宽掉得多的组合说明那两张卡没坐在高速链路的同一侧。能跑、能快接下来把最常翻车的几处排一排。故障速查5 类高频问题的 30 秒定位GPU 未被识别lspci | grep -i vga rocminfo | grep -i gfx预期lspci出现AMD/ATI的控制器rocminfo列出对应 gfx 代号。第一条有、第二条没有优先查内核驱动是否加载lsmod | grep amdgpu和内核版本是否匹配。hipcc 找不到which hipcc || ls /opt/rocm/hip/bin/hipcc预期which给出路径。找不到就确认是否装过rocm-hip-sdk或把/opt/rocm/bin加进PATH再试。kernel launch 报权限拒绝groups预期输出包含video和render两个组。缺哪个sudo usermod -aG 组名 $USER补上后注销重登。显存不足amd-smi showmem预期逐卡列出 VRAM 的已用与总量。Free还不少却报 OOM多半是碎片或别的服务占着卡先查进程。多卡带宽异常./rccl-tests/build/all_reduce_perf -g 8预期busbw贴近互联峰值。掉得厉害时回上一节检查拓扑跳数少的卡对才是好邻居。问题清完了剩下的是把它变成你自己的工具箱。接下来可以做的四件事把saxpy里的4.0f * x[i] y[i]换成你自己的算式改成分块版本进 LDS量一版前后的耗时差训练时开一个终端跑watch -n1 amd-smi盯 GPU 利用率与显存曲线判断瓶颈在计算还是在数据搬运把 hip/samples 目录下的官方示例挨个编一遍重点看多卡那几个用 RCCL 在自己机器上实测一次卡间带宽把busbw数字和卡规格对比存档仓库里还有两样东西值得翻GPU 架构参考手册在docs/reference/gpu-arch/系统级优化指南在docs/reference/system-optimization/。第一个核函数跑起来的那一刻后面就都是顺水推舟的事。【免费下载链接】legacy-rocm-buildAMD ROCm™ Software - GitHub Home项目地址: https://gitcode.com/GitHub_Trending/ro/legacy-rocm-build创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 AMD ROCm 从装好到跑通：5 个动作装好环境、写出第一个 GPU 核函数
AMD ROCm 从装好到跑通5 个动作装好环境、写出第一个 GPU 核函数【免费下载链接】legacy-rocm-buildAMD ROCm™ Software - GitHub Home项目地址: https://gitcode.com/GitHub_Trending/ro/legacy-rocm-build训练任务被 CPU 单核拖住、NVIDIA 的卡要排队、手上的 AMD 显卡却在吃灰——这种局面值得动一动了。AMD ROCm 就是 AMD 的开源 GPU 加速计算平台从内核驱动、数学库、HIP 运行时到 PyTorch 对接一层层配齐让深度学习与 HPC 任务直接落在 AMD GPU 上跑。这篇文章按装好 → 跑通 → 调快的顺序给你一份能直接复制执行的完整方案。ROCm 的覆盖范围Instinct、Radeon 到八卡互联先看它管得住多大的地盘。ROCm 支持的硬件分三档Instinct 数据中心卡MI300 系列是主力、Radeon 独显、Ryzen AI 的 APU支持的操作系统以主流 Linux 发行版为主Ubuntu 22.04/24.04、RHEL、SLES 等。能力上覆盖单卡推理、多卡并行训练和分布式集合通信最上层能直接对接 PyTorch 这类框架。AMD MI300X 节点架构8 张 Instinct 卡经 Infinity Fabric 全互联是 ROCm 上多卡并行训练的基本单元。知道它管到哪就开工装它。Ubuntu 上四条命令装好 ROCm附两个成功标志以 Ubuntu 24.04 为例apt 路线总共四步全部可复制粘贴。第一步先备好 ROCm 官方 apt 源和两个基础依赖库# 装依赖libatomic1/libquadmath0 是 ROCm 工具链运行时的硬需求缺了会报符号错误 sudo apt update sudo apt install -y libatomic1 libquadmath0 # 导入 ROCm 发布签名密钥并注册官方 apt 源 sudo mkdir -p --mode0755 /etc/apt/keyrings wget -O - https://repo.amd.com/rocm/packages-multi-arch/gpg/rocm.gpg | gpg --dearmor | sudo tee /etc/apt/keyrings/amdrocm.gpg /dev/null sudo tee /etc/apt/sources.list.d/rocm.list EOF deb [archamd64 signed-by/etc/apt/keyrings/amdrocm.gpg] https://repo.amd.com/rocm/packages-multi-arch/ubuntu2404 stable main EOF sudo apt update然后安装、放权、验证# 2. 安装 HIP SDKhipcc 编译器 HIP 运行时GPU 编程的最小可用集 sudo apt install -y rocm-hip-sdk # 3. 把当前用户加入 video/render 设备组否则普通用户程序打不开 GPU sudo usermod -aG video,render $USER # 4. 验证两条命令都要有输出才算装好 rocminfo | grep -E Marketing Name|gfx amd-smi version成功标志有两个rocminfo的 HSA Agents 段落能列出你显卡的 gfx 代号与商品名例如AMD Instinct MI350Xamd-smi version打印出 ROCm 版本与 amdgpu 驱动版本。第 3 步要注销重登或重启才生效。apt 不方便时离线机器、想连驱动一起装官方还提供 runfile 安装器ROCm 与 amdgpu 驱动打包在一起自带图形化引导菜单# 下载约 4GB 的一体化安装器双击或终端进入菜单逐项选择 wget https://repo.radeon.com/rocm/installer/rocm-runfile-installer/rocm-rel-7.14/rocm-installer-7.14.0-7.run chmod x rocm-installer-7.14.0-7.run ./rocm-installer-7.14.0-7.runROCm runfile 安装器主菜单ROCm 与 amdgpu 驱动打成一包支持离线安装与图形化选择。提醒一句系统里已有 ROCm 7.2.4 或更老版本的话先卸载干净再装新版混装会留下难查的库冲突。环境有了下一个动作就是写你自己的第一个核函数。第一个核函数SAXPYhipcc 一编就跑下面这段代码实现数值计算里的经典例程 SAXPYy a*x y。它不是玩具PyTorch 里成千上万个算子底层干的就是这类每个线程处理一个元素的活。把下面内容存成saxpy.cpp#include hip/hip_runtime.h #include iostream #include vector // 核函数一个线程只负责一个元素的 y[i] a*x[i] y[i] __global__ void saxpy(const float* x, const float* y, float* out, int n) { int i hipBlockIdx_x * hipBlockDim_x hipThreadIdx_x; if (i n) out[i] 4.0f * x[i] y[i]; } int main() { const int n 1 20; // 1048576 个元素规模足够让 GPU 忙起来 std::vectorfloat hX(n, 1.0f), hY(n, 2.0f), hO(n, 0.0f); float *dX, *dY, *dO; hipMalloc(dX, n * sizeof(float)); // 在设备显存上分配空间 hipMalloc(dY, n * sizeof(float)); hipMalloc(dO, n * sizeof(float)); hipMemcpy(dX, hX.data(), n * sizeof(float), hipMemcpyHostToDevice); hipMemcpy(dY, hY.data(), n * sizeof(float), hipMemcpyHostToDevice); // 2048 个线程块 x 256 线程共 524288 个线程同时开工 hipLaunchKernelGGL(saxpy, dim3(2048), dim3(256), 0, 0, dX, dY, dO, n); hipDeviceSynchronize(); // 核函数是异步启动的必须等它干完再取结果 hipMemcpy(hO.data(), dO, n * sizeof(float), hipMemcpyDeviceToHost); // 正确性自检4*126逐项验证 bool ok true; for (int i 0; i n; i) if (hO[i] ! 6.0f) { ok false; break; } std::cout (ok ? SAXPY 校验通过全部 1M 个元素等于 6.0 : 校验失败) std::endl; hipFree(dX); hipFree(dY); hipFree(dO); return ok ? 0 : 1; }编译运行只需两行hipcc -O2 -o saxpy saxpy.cpp ./saxpy预期输出SAXPY 校验通过全部 1M 个元素等于 6.0。想针对自己的卡优化先用rocminfo | grep gfx查出架构代号再补上参数例如hipcc -O2 -archgfx90a -o saxpy saxpy.cpp。跑通了只是及格线它到底凭什么快值得花三分钟搞清楚。为什么快把全局内存访问压进片上抽屉把 CUCompute Unit计算单元想成一家小型印刷厂一条 SIMD 流水线是一台机头同一张印版一条指令能同时印几十份页几十条数据这就是 GPU 吞吐的来源VGPR/SGPR 寄存器是机头旁的便签随手可取LDSLocal Data Share片上共享内存是厂里共享的印版台同一批工人围着它干不用再跑一趟仓库。单个 CU 的内部构成调度器在上SIMD 流水线居中LDS 与寄存器文件在底部——核函数优化基本都是围着这几块做文章。GPU 快就快在这条纪律让海量线程共用同一条指令流数据尽量从寄存器和 LDS 里取少碰全局显存。显存带宽有限每次从 DRAM 取数都要排队而同一批数据被多个线程反复取正是 GPU 程序里最常见的浪费。分块tiling优化干的就是这件事——数据取一次进 LDS块内大家共享一次全局读被摊薄几十次。代价呢寄存器是硬通货。每个线程分到的 VGPR 越多一张卡上同时驻留的线程就越少这叫占用率occupancy。占用率掉下去流水线就没有足够线程切换来填满等待时间机器白空转。寄存器占用与占用率的关系VGPR 用量决定每个计算单元能驻留多少线程是调核函数时的核心约束。单卡的故事讲完接下来把你的 Python 代码也搬上来。PyTorch 接入 ROCm两行 pip 加一次验证算子层的地基有了把活交给框架。ROCm 版 PyTorch 的安装是标准 pip 操作只是把索引源换成 ROCm 的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就成了。注意 ROCm 里照旧用torch.cuda这套 API名字是 CUDA 历史遗留代码一行不用改。纯 Python 工作流不想动系统环境还有另一条路用虚拟环境装打包好的 ROCm SDKpython -m venv .venv source .venv/bin/activate python -m pip install --index-url https://repo.amd.com/rocm/whl-multi-arch/ rocm[libraries,device-gfx950]7.14.0 pip freeze | grep rocm # 预期能列出 rocm-sdk 等已装包device-gfx950要换成rocminfo里你显卡的架构代号对应关系仓库的兼容矩阵里都有。框架能跑之后多卡就顺理成章了。多卡互联怎么验证一次 RCCL 集合通信测试八张卡挂在同一节点上数据要在卡与卡之间倒手带宽不够一切白搭。RCCLROCm 的集合通信库对标 NCCL是这条链路的底座而rccl-tests里的all_reduce就是它的体检表./rccl-tests/build/all_reduce_perf -b 8 -e 256M -f 2 -g 8八卡 all_reduce 集合通信测试输出数据规模、总线带宽busbw与延迟一目了然是分布式训练前验证卡间链路的常用手段。看两列就够busbw总线带宽应接近该卡规格里的互联峰值逐档数据量都稳跳数大、带宽掉得多的组合说明那两张卡没坐在高速链路的同一侧。能跑、能快接下来把最常翻车的几处排一排。故障速查5 类高频问题的 30 秒定位GPU 未被识别lspci | grep -i vga rocminfo | grep -i gfx预期lspci出现AMD/ATI的控制器rocminfo列出对应 gfx 代号。第一条有、第二条没有优先查内核驱动是否加载lsmod | grep amdgpu和内核版本是否匹配。hipcc 找不到which hipcc || ls /opt/rocm/hip/bin/hipcc预期which给出路径。找不到就确认是否装过rocm-hip-sdk或把/opt/rocm/bin加进PATH再试。kernel launch 报权限拒绝groups预期输出包含video和render两个组。缺哪个sudo usermod -aG 组名 $USER补上后注销重登。显存不足amd-smi showmem预期逐卡列出 VRAM 的已用与总量。Free还不少却报 OOM多半是碎片或别的服务占着卡先查进程。多卡带宽异常./rccl-tests/build/all_reduce_perf -g 8预期busbw贴近互联峰值。掉得厉害时回上一节检查拓扑跳数少的卡对才是好邻居。问题清完了剩下的是把它变成你自己的工具箱。接下来可以做的四件事把saxpy里的4.0f * x[i] y[i]换成你自己的算式改成分块版本进 LDS量一版前后的耗时差训练时开一个终端跑watch -n1 amd-smi盯 GPU 利用率与显存曲线判断瓶颈在计算还是在数据搬运把 hip/samples 目录下的官方示例挨个编一遍重点看多卡那几个用 RCCL 在自己机器上实测一次卡间带宽把busbw数字和卡规格对比存档仓库里还有两样东西值得翻GPU 架构参考手册在docs/reference/gpu-arch/系统级优化指南在docs/reference/system-optimization/。第一个核函数跑起来的那一刻后面就都是顺水推舟的事。【免费下载链接】legacy-rocm-buildAMD ROCm™ Software - GitHub Home项目地址: https://gitcode.com/GitHub_Trending/ro/legacy-rocm-build创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考