AMD Instinct MI100 微架构深度解析CDNA 架构的节点互联、计算单元与性能模型【免费下载链接】legacy-rocm-buildAMD ROCm™ Software - GitHub Home项目地址: https://gitcode.com/GitHub_Trending/ro/legacy-rocm-build本文基于 ROCm 官方文档 docs/reference/gpu-arch/mi100.md 编写结合 GPU 架构文档索引、AMD GPU 规格表、精度支持说明 与 GPU 原子操作支持说明 等仓库资料从节点级系统架构、芯片级微架构到单个计算单元CU的执行模型完整解读 AMD Instinct MI100 作为首代 CDNA 架构 GPU 的设计要点。读者阅读完本文后将能清晰掌握 MI100 的互联拓扑与带宽特性、各数据类型下的理论峰值性能模型以及寄存器分配与占用率occupancy之间的权衡关系为后续进行内核性能调优、profiling 数据解读和系统级性能分析奠定基础。概述MI100 与 CDNA 架构的定位AMD Instinct MI100 是 AMD 首款基于CDNACompute DNA架构的加速卡也是 ROCm 软件栈早期支持的核心硬件之一。CDNA 架构不再兼顾图形渲染任务而是专注面向计算应用——从单台服务器到全球最大规模的百亿亿次exascale超级计算机——进行极端可扩展性与计算性能设计主要目标场景为**高性能计算HPC与AI / 机器学习ML**负载。在 ROCm 文档体系中MI100 被归类为CDNA1架构其 LLVM 目标名为gfx908。这一点可以从以下仓库文件得到印证AMD GPU 规格表 中 MI100 一行明确记录Architecture 为CDNA、LLVM target name 为gfx908、VRAM 为 32 GiB、Compute Units 为 120、Wavefront Size 为 64硬件过渡指南The Rock 中架构与-gfx参数的对应表同样列出CDNA 对应-gfx908对应硬件为 AMD Instinct MI100GPU 选择器 中也包含AMD Instinct MI100 (gfx908)选项其 selector 值为mi100 gfxgfx908精度支持文档 在按硬件架构划分数据类型支持时将 MI100 列为 CDNA1 的代表产品。此外GPU 架构文档索引 中将 MI100 与 CDNAInstinct MI100一代对应并提供了 ISA 参考与白皮书的指引MI100 之后是 CDNA2MI200 系列gfx90a与更新的 CDNA3 / CDNA4 世代。适用前提说明本文所有规格与带宽数字均来自当前仓库文档所记载的 MI100CDNA代硬件设计参数适用于面向gfx908目标的内核开发与性能分析场景实际服务器系统的可用配置如 GPU 数量、是否安装 Infinity Fabric 桥接取决于认证服务器的具体部署。节点级系统架构双路 EPYC 与八卡 GPU 互联MI100 的典型节点配置围绕两颗 AMD EPYC 处理器与最多八颗 AMD Instinct GPU构建其互联拓扑由三个层次构成处理器间互联AMD Infinity Fabric两颗 EPYC 处理器通过AMD Infinity Fabric互连该互联提供最高 18 GT/sec的高带宽且缓存一致的链接使得每颗处理器都可以把节点内可用内存作为一个单一共享内存域以**非一致内存访问NUMA**方式访问。在双路2P配置下共有三条 AMD Infinity Fabric 链接用于连接两颗处理器每颗处理器额外提供一条PCIe Gen 4 x16链接用于挂接网络主机适配器HBA等附加 I/O 设备。处理器到 GPU 互联PCIe Gen 4在典型节点中每颗处理器最多可承载四颗 AMD Instinct GPU通过PCIe Gen 4链接以16 GT/sec速率挂接对应的峰值双向链路带宽为 32 GB/sec。GPU 到 GPU 互联AMD Infinity Fabric 与桥接每四颗 GPU 组成一个 hive蜂巢组可以通过完全互联且缓存一致的 AMD Instinct fabric 互连。组内 GPU 之间使用23 GT/sec的 AMD Infinity Fabric 链接其运行频率高于处理器间链接。需要特别注意的是这种 GPU 间链接只有在通过认证的服务器系统中才能建立要求 GPU 安装在相邻的 PCIe 插槽上并安装AMD Infinity Fabric 桥接器bridge用于 AMD Instinct GPU。综合来看MI100 节点的互联呈现 CPU 之间 18 GT/s、CPU-GPU 之间 PCIe Gen4 16 GT/s、GPU 之间 23 GT/s 的三级速率梯度GPU 之间通过 Infinity Fabric 获得最快的点对点通信能力这对多卡并行训练与集合通信如 RCCL的性能至关重要。芯片级微架构HBM2 内存与 Shader Engine 布局MI100 GPU 芯片的微架构延续了 CDNA 的整体设计。下图展示了 MI100 芯片的结构从图中可以清晰看出芯片的对外接口与内部布局底部为 PCIe Gen 4 x16 链接16 GT/sec用于连接宿主机处理器底部还分布着三个 AMD Infinity Fabric 端口23 GT/sec用于连接本地 hive 中的其他 GPU芯片左、右两侧为高带宽内存HBM通过 GPU 内存控制器memory controller挂接。内存子系统四堆栈 HBM2MI100 配备四堆栈 HBM 第二代HBM2总容量32 GB内存接口宽度达4096 bit。在1.2 GHz内存时钟频率下峰值内存带宽为1.228 TB/sec。注在 AMD GPU 规格表 的 MI100 行中L2 缓存记录为 8 MiB、L1 向量缓存为 16 KiB每 3 个 CU 共享 16 KiB、L1 标量缓存为 32 KiB每 3 个 CU 共享 32 KiB这些参数与本文微架构描述相互补充。计算单元组织120 CU / 8 SE / 4 SIMD芯片内部的执行单元为计算单元Compute Unit, CU共120 个计算单元物理上组织为8 个 Shader EngineSE每个 SE 包含15 个 CU每个 CU 进一步细分为4 个 SIMD 单元每个 SIMD 单元每条指令处理16 个数据元素因此单个 CU 每条指令可处理64 个数据元素——即一个wavefront波前峰值时钟频率为1.5 GHz。据此可以推导理论峰值 FP64 性能4 [SIMD units] × 16 [elements per instruction] × 120 [CU] × 1.5 [GHz] 11.5 TFLOPS不同数据类型的峰值性能MI100 的峰值性能因数据类型与执行单元向量 SIMD 单元 vs 矩阵单元/Matrix core而异下表完整记录了官方文档的性能模型Computation and Data TypeFLOPS/CLOCK/CUPeak TFLOPSVector FP646411.5Matrix FP3225646.1Vector FP3212823.1Matrix FP161024184.6Matrix BF1651292.3其中中间列表示单个 CU 在每个时钟周期内退休一条 SIMD或矩阵指令时所能处理的数据元素数量第三列为整颗 GPU 的理论峰值 TFLOPS。从表中可以看到Matrix矩阵类指令的吞吐显著高于 Vector向量类指令例如 Matrix FP32 为 256 FLOPS/CLOCK/CU而 Vector FP32 仅为 128Matrix FP16 高达 1024体现出 CDNA 架构为矩阵乘加等密集计算专门增强的设计取向FP16 的峰值184.6 TFLOPS约为 FP32 的 4 倍BF16 约为 FP32 的 2 倍符合混合精度训练对半精度计算能力的需求向量 FP64 与向量 FP32 的比值为 1:2保持了对双精度科学计算HPC 传统负载的完整支持。需要强调的是这些都是理论峰值实际可达到的性能还受内存带宽、内核访存模式、占用率与指令调度效率等因素制约不应将其等同于实际应用性能。单个计算单元CU内部结构指令流与执行资源上图展示了 MI100 单个 CU 的框图总结了指令在执行引擎中的流动路径指令获取CU 通过32 KB 指令缓存instruction cache获取指令指令分发指令经由**调度器dispatcher**向前传递到执行单元波前处理能力单个 CU 最多可同时处理10 个 wavefront并将它们的指令送入执行单元。寄存器堆VGPR 与 SGPR执行单元包含两类通用寄存器堆256 个向量通用寄存器VGPR800 个标量通用寄存器SGPR。VGPR 与 SGPR 会被动态分配给正在执行的 wavefront。需要特别注意寄存器使用的上限与代价一个 wavefront最多可访问 102 个标量寄存器。若标量寄存器使用量超过该上限将引发寄存器溢出register spilling从而可能影响执行性能。对照 AMD GPU 规格表MI100 的 VGPR 文件记录为 256 VGPR and 256 AccVGPR即在向量寄存器之外还包含 256 个累加寄存器供矩阵指令使用。占用率Occupancy与寄存器分配的权衡一个 wavefront 可以使用0 到 256 个 VGPR 中的任意数量这直接影响占用率——即 CU 内同时处于活动状态的 wavefront 数量。文档给出了一个具体示例例如当使用119 个 VGPR时同一时刻 CU 内只能有2 个 wavefront处于活动状态。由于每条 SIMD 指令存在4 个时钟周期的指令延迟占用率应当尽可能高这样计算单元才能通过从多个 wavefront 中调度指令来隐藏延迟、提升执行效率。这构成了内核性能调优中的经典权衡为单个 wavefront 分配更多寄存器提高单波前并行度会降低占用率减少可并发波前数进而可能影响延迟隐藏能力。对内核开发者而言可以通过__launch_bounds__或编译器的-maxrregcount类机制约束寄存器使用量来换取更高占用率但具体行为以实际编译器与驱动版本为准。MI100 的编程与系统优化要点数据类型支持精度支持文档 将 MI100 列为 CDNA1 的代表。从文档中的数据类型支持矩阵可以看到CDNA1 支持int8_t/uint8_t等整数类型以及对应的 HIP C 数据类型实现在开发以 MI100gfx908为目标的 HIP 内核时应以该文档的数据类型支持矩阵为准选择可用的计算类型。原子操作支持GPU 原子操作支持文档 中专门列出了 MI100 的原子操作行为这是编写并发内核时必须了解的关键约束L2 缓存与 Infinity Fabric 均支持常见的整数与位运算原子操作L2 缓存支持无返回NoReturn版本的 packed-FP16 与 FP32 原子加但这些操作无法返回数据Infinity Fabric 不支持 packed-FP16 或 FP32 原子加因此这些操作无法经由 Infinity Fabric 传输粗粒度coarse-grained内存被标记为可缓存原子操作将在L2 缓存中处理细粒度fine-grained内存通过页表被标记为不可缓存命中不可缓存内存的原子操作会被转发到 Infinity Fabric若不可缓存数据存放在宿主机内存中原子操作会从 Infinity Fabric 转发到 PCIePCIe 总线不支持的原子操作将变成 NOP 并产生错误结果若浮点原子加命中不可缓存内存由于无法转发到 Infinity Fabric同样会 NOP 并产生错误结果。这提醒开发者在 MI100 上使用浮点原子操作时应确保内存属性粗粒度/细粒度与操作类型匹配避免因转发路径不支持而静默得到错误结果。系统级优化指引ROCm 文档还通过 AMD Instinct 系统优化索引远端内容聚合自ROCm/amdgpu-docs仓库为 MI100 提供专门的系统优化指南对应mi100小节涵盖系统设置与调优建议可作为部署 MI100 节点的进一步参考资料。结语作为 AMD 首代专注计算的 CDNA 架构产品AMD Instinct MI100 确立了后续 CDNA 世代延续的设计主线以 Shader Engine 组织大规模计算单元、以多堆栈 HBM2 提供超高内存带宽、以 AMD Infinity Fabric 构建节点内 GPU 间的低延迟高带宽互联并在每个 CU 内通过 VGPR/SGPR 动态分配与多 wavefront 调度来支撑高占用率执行。理解这些微架构细节是解读 profiling 数据、制定内核优化策略以及评估系统级性能瓶颈的起点。若需进一步对比后续世代可参考仓库中的 MI250 微架构文档 与 MI300 微架构文档。【免费下载链接】legacy-rocm-buildAMD ROCm™ Software - GitHub Home项目地址: https://gitcode.com/GitHub_Trending/ro/legacy-rocm-build创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 AMD Instinct MI100 微架构深度解析：CDNA 架构的节点互联、计算单元与性能模型
AMD Instinct MI100 微架构深度解析CDNA 架构的节点互联、计算单元与性能模型【免费下载链接】legacy-rocm-buildAMD ROCm™ Software - GitHub Home项目地址: https://gitcode.com/GitHub_Trending/ro/legacy-rocm-build本文基于 ROCm 官方文档 docs/reference/gpu-arch/mi100.md 编写结合 GPU 架构文档索引、AMD GPU 规格表、精度支持说明 与 GPU 原子操作支持说明 等仓库资料从节点级系统架构、芯片级微架构到单个计算单元CU的执行模型完整解读 AMD Instinct MI100 作为首代 CDNA 架构 GPU 的设计要点。读者阅读完本文后将能清晰掌握 MI100 的互联拓扑与带宽特性、各数据类型下的理论峰值性能模型以及寄存器分配与占用率occupancy之间的权衡关系为后续进行内核性能调优、profiling 数据解读和系统级性能分析奠定基础。概述MI100 与 CDNA 架构的定位AMD Instinct MI100 是 AMD 首款基于CDNACompute DNA架构的加速卡也是 ROCm 软件栈早期支持的核心硬件之一。CDNA 架构不再兼顾图形渲染任务而是专注面向计算应用——从单台服务器到全球最大规模的百亿亿次exascale超级计算机——进行极端可扩展性与计算性能设计主要目标场景为**高性能计算HPC与AI / 机器学习ML**负载。在 ROCm 文档体系中MI100 被归类为CDNA1架构其 LLVM 目标名为gfx908。这一点可以从以下仓库文件得到印证AMD GPU 规格表 中 MI100 一行明确记录Architecture 为CDNA、LLVM target name 为gfx908、VRAM 为 32 GiB、Compute Units 为 120、Wavefront Size 为 64硬件过渡指南The Rock 中架构与-gfx参数的对应表同样列出CDNA 对应-gfx908对应硬件为 AMD Instinct MI100GPU 选择器 中也包含AMD Instinct MI100 (gfx908)选项其 selector 值为mi100 gfxgfx908精度支持文档 在按硬件架构划分数据类型支持时将 MI100 列为 CDNA1 的代表产品。此外GPU 架构文档索引 中将 MI100 与 CDNAInstinct MI100一代对应并提供了 ISA 参考与白皮书的指引MI100 之后是 CDNA2MI200 系列gfx90a与更新的 CDNA3 / CDNA4 世代。适用前提说明本文所有规格与带宽数字均来自当前仓库文档所记载的 MI100CDNA代硬件设计参数适用于面向gfx908目标的内核开发与性能分析场景实际服务器系统的可用配置如 GPU 数量、是否安装 Infinity Fabric 桥接取决于认证服务器的具体部署。节点级系统架构双路 EPYC 与八卡 GPU 互联MI100 的典型节点配置围绕两颗 AMD EPYC 处理器与最多八颗 AMD Instinct GPU构建其互联拓扑由三个层次构成处理器间互联AMD Infinity Fabric两颗 EPYC 处理器通过AMD Infinity Fabric互连该互联提供最高 18 GT/sec的高带宽且缓存一致的链接使得每颗处理器都可以把节点内可用内存作为一个单一共享内存域以**非一致内存访问NUMA**方式访问。在双路2P配置下共有三条 AMD Infinity Fabric 链接用于连接两颗处理器每颗处理器额外提供一条PCIe Gen 4 x16链接用于挂接网络主机适配器HBA等附加 I/O 设备。处理器到 GPU 互联PCIe Gen 4在典型节点中每颗处理器最多可承载四颗 AMD Instinct GPU通过PCIe Gen 4链接以16 GT/sec速率挂接对应的峰值双向链路带宽为 32 GB/sec。GPU 到 GPU 互联AMD Infinity Fabric 与桥接每四颗 GPU 组成一个 hive蜂巢组可以通过完全互联且缓存一致的 AMD Instinct fabric 互连。组内 GPU 之间使用23 GT/sec的 AMD Infinity Fabric 链接其运行频率高于处理器间链接。需要特别注意的是这种 GPU 间链接只有在通过认证的服务器系统中才能建立要求 GPU 安装在相邻的 PCIe 插槽上并安装AMD Infinity Fabric 桥接器bridge用于 AMD Instinct GPU。综合来看MI100 节点的互联呈现 CPU 之间 18 GT/s、CPU-GPU 之间 PCIe Gen4 16 GT/s、GPU 之间 23 GT/s 的三级速率梯度GPU 之间通过 Infinity Fabric 获得最快的点对点通信能力这对多卡并行训练与集合通信如 RCCL的性能至关重要。芯片级微架构HBM2 内存与 Shader Engine 布局MI100 GPU 芯片的微架构延续了 CDNA 的整体设计。下图展示了 MI100 芯片的结构从图中可以清晰看出芯片的对外接口与内部布局底部为 PCIe Gen 4 x16 链接16 GT/sec用于连接宿主机处理器底部还分布着三个 AMD Infinity Fabric 端口23 GT/sec用于连接本地 hive 中的其他 GPU芯片左、右两侧为高带宽内存HBM通过 GPU 内存控制器memory controller挂接。内存子系统四堆栈 HBM2MI100 配备四堆栈 HBM 第二代HBM2总容量32 GB内存接口宽度达4096 bit。在1.2 GHz内存时钟频率下峰值内存带宽为1.228 TB/sec。注在 AMD GPU 规格表 的 MI100 行中L2 缓存记录为 8 MiB、L1 向量缓存为 16 KiB每 3 个 CU 共享 16 KiB、L1 标量缓存为 32 KiB每 3 个 CU 共享 32 KiB这些参数与本文微架构描述相互补充。计算单元组织120 CU / 8 SE / 4 SIMD芯片内部的执行单元为计算单元Compute Unit, CU共120 个计算单元物理上组织为8 个 Shader EngineSE每个 SE 包含15 个 CU每个 CU 进一步细分为4 个 SIMD 单元每个 SIMD 单元每条指令处理16 个数据元素因此单个 CU 每条指令可处理64 个数据元素——即一个wavefront波前峰值时钟频率为1.5 GHz。据此可以推导理论峰值 FP64 性能4 [SIMD units] × 16 [elements per instruction] × 120 [CU] × 1.5 [GHz] 11.5 TFLOPS不同数据类型的峰值性能MI100 的峰值性能因数据类型与执行单元向量 SIMD 单元 vs 矩阵单元/Matrix core而异下表完整记录了官方文档的性能模型Computation and Data TypeFLOPS/CLOCK/CUPeak TFLOPSVector FP646411.5Matrix FP3225646.1Vector FP3212823.1Matrix FP161024184.6Matrix BF1651292.3其中中间列表示单个 CU 在每个时钟周期内退休一条 SIMD或矩阵指令时所能处理的数据元素数量第三列为整颗 GPU 的理论峰值 TFLOPS。从表中可以看到Matrix矩阵类指令的吞吐显著高于 Vector向量类指令例如 Matrix FP32 为 256 FLOPS/CLOCK/CU而 Vector FP32 仅为 128Matrix FP16 高达 1024体现出 CDNA 架构为矩阵乘加等密集计算专门增强的设计取向FP16 的峰值184.6 TFLOPS约为 FP32 的 4 倍BF16 约为 FP32 的 2 倍符合混合精度训练对半精度计算能力的需求向量 FP64 与向量 FP32 的比值为 1:2保持了对双精度科学计算HPC 传统负载的完整支持。需要强调的是这些都是理论峰值实际可达到的性能还受内存带宽、内核访存模式、占用率与指令调度效率等因素制约不应将其等同于实际应用性能。单个计算单元CU内部结构指令流与执行资源上图展示了 MI100 单个 CU 的框图总结了指令在执行引擎中的流动路径指令获取CU 通过32 KB 指令缓存instruction cache获取指令指令分发指令经由**调度器dispatcher**向前传递到执行单元波前处理能力单个 CU 最多可同时处理10 个 wavefront并将它们的指令送入执行单元。寄存器堆VGPR 与 SGPR执行单元包含两类通用寄存器堆256 个向量通用寄存器VGPR800 个标量通用寄存器SGPR。VGPR 与 SGPR 会被动态分配给正在执行的 wavefront。需要特别注意寄存器使用的上限与代价一个 wavefront最多可访问 102 个标量寄存器。若标量寄存器使用量超过该上限将引发寄存器溢出register spilling从而可能影响执行性能。对照 AMD GPU 规格表MI100 的 VGPR 文件记录为 256 VGPR and 256 AccVGPR即在向量寄存器之外还包含 256 个累加寄存器供矩阵指令使用。占用率Occupancy与寄存器分配的权衡一个 wavefront 可以使用0 到 256 个 VGPR 中的任意数量这直接影响占用率——即 CU 内同时处于活动状态的 wavefront 数量。文档给出了一个具体示例例如当使用119 个 VGPR时同一时刻 CU 内只能有2 个 wavefront处于活动状态。由于每条 SIMD 指令存在4 个时钟周期的指令延迟占用率应当尽可能高这样计算单元才能通过从多个 wavefront 中调度指令来隐藏延迟、提升执行效率。这构成了内核性能调优中的经典权衡为单个 wavefront 分配更多寄存器提高单波前并行度会降低占用率减少可并发波前数进而可能影响延迟隐藏能力。对内核开发者而言可以通过__launch_bounds__或编译器的-maxrregcount类机制约束寄存器使用量来换取更高占用率但具体行为以实际编译器与驱动版本为准。MI100 的编程与系统优化要点数据类型支持精度支持文档 将 MI100 列为 CDNA1 的代表。从文档中的数据类型支持矩阵可以看到CDNA1 支持int8_t/uint8_t等整数类型以及对应的 HIP C 数据类型实现在开发以 MI100gfx908为目标的 HIP 内核时应以该文档的数据类型支持矩阵为准选择可用的计算类型。原子操作支持GPU 原子操作支持文档 中专门列出了 MI100 的原子操作行为这是编写并发内核时必须了解的关键约束L2 缓存与 Infinity Fabric 均支持常见的整数与位运算原子操作L2 缓存支持无返回NoReturn版本的 packed-FP16 与 FP32 原子加但这些操作无法返回数据Infinity Fabric 不支持 packed-FP16 或 FP32 原子加因此这些操作无法经由 Infinity Fabric 传输粗粒度coarse-grained内存被标记为可缓存原子操作将在L2 缓存中处理细粒度fine-grained内存通过页表被标记为不可缓存命中不可缓存内存的原子操作会被转发到 Infinity Fabric若不可缓存数据存放在宿主机内存中原子操作会从 Infinity Fabric 转发到 PCIePCIe 总线不支持的原子操作将变成 NOP 并产生错误结果若浮点原子加命中不可缓存内存由于无法转发到 Infinity Fabric同样会 NOP 并产生错误结果。这提醒开发者在 MI100 上使用浮点原子操作时应确保内存属性粗粒度/细粒度与操作类型匹配避免因转发路径不支持而静默得到错误结果。系统级优化指引ROCm 文档还通过 AMD Instinct 系统优化索引远端内容聚合自ROCm/amdgpu-docs仓库为 MI100 提供专门的系统优化指南对应mi100小节涵盖系统设置与调优建议可作为部署 MI100 节点的进一步参考资料。结语作为 AMD 首代专注计算的 CDNA 架构产品AMD Instinct MI100 确立了后续 CDNA 世代延续的设计主线以 Shader Engine 组织大规模计算单元、以多堆栈 HBM2 提供超高内存带宽、以 AMD Infinity Fabric 构建节点内 GPU 间的低延迟高带宽互联并在每个 CU 内通过 VGPR/SGPR 动态分配与多 wavefront 调度来支撑高占用率执行。理解这些微架构细节是解读 profiling 数据、制定内核优化策略以及评估系统级性能瓶颈的起点。若需进一步对比后续世代可参考仓库中的 MI250 微架构文档 与 MI300 微架构文档。【免费下载链接】legacy-rocm-buildAMD ROCm™ Software - GitHub Home项目地址: https://gitcode.com/GitHub_Trending/ro/legacy-rocm-build创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考