CUDA 护城河依然存在:SemiAnalysis 对 AMD MI300X 深度评测之训练篇
近日,SemiAnalysis 发布了对比 AMD MI300X 和 NVIDIA H100/H200 在训练场景的深度评测1,该报告内容非常丰富,涵盖性能、软件、网络、成本等多方面。报告指出,尽管 AMD MI300X 的纸面规格和 TCO 明显比 H100 和 H200 有更大优势,受限于软件制约,MI300X 的实际使用体验和实际训练性能并不及 H100/H200。
Dylan Patel 的报告2 引来了包括 AMD 苏妈3在内的众多业界大佬的关注。本文编译自该报告,并做了适当的精简和补充,欢迎指正。
TLDR
- 只通过算力 FLOP/s 和 HBM 的带宽/容量的纸面规格来比较不同芯片的性能并不可靠,就像只通过有几百万像素来比较相机的摄影性能一样。唯一可靠的方法是运行对应的基准测试。
- 相比于 NVIDIA 软硬件 Stack 的开箱即用,AMD 的开箱即用体验很差,即使是 AMD PyTorch 的公开稳定版本也依然存在问题。
- 与 H100/H200 相比,MI300X 的总拥有成本 TCO 更低,但是基于 AMD 公开的 PyTorch 版本上 MI300X 的 Training Performance per TCO 更差。如果是使用 AMD 定制的软件栈,性能会有所提升。
- 与 H100/H200 相比,AMD 的单机训练吞吐量更低,这与纸面规格数据不太一致。具体来说,MI300X 的 GEMM 性能较差,远不及规格的 TFLOP/s。尽管 NVIDIA 的真实性能也比纸面规格要低一些,但是差距不大。
- 与 NVIDIA 的集合通信库 NCCL 和 InfiniBand/Spectrum-X 紧密结合相比,AMD 的集合通信库 RCCL 与网络和交换硬件的集成度较低,从而限制了在 AMD Scale Out 的训练性能。
- 很多 AMD 的 AI 库都 fork 自 NVIDIA 的库,导致性能更低和一些兼容性问题
对 AMD 的可执行建议
- AMD 需要给自己的工程师更多自己的 GPU 芯片用于修复和改善 AMD 生态系统。相比于 NVIDIA 给工程师的资源,AMD 内部研发团队拥有的 GPU 太少。
- AMD 需要将上千的 MI300X、MI325X 提供给 PyTorch 的 CI/CD 用于自动化测试,以确保 AMD 的性能回归和功能错误检查。NVIDIA 能够做到足够好的开箱体验的重要一点是,他们已经 PyTorch CI/CD 提供了数千个 GPU。
- AMD 应该和 Meta 合作,来让有更多的 LLM 训练负载工作在 PyTorch ROCm 上,从而尽快发现更多可能的问题。
- 在 AMD 硬件上部署应用的时候,不要再过度依赖于各种各样的环境变量。更好的选择是让他们作为默认设置,这样可以极大改善 AMD 硬件的开箱即用。
- 不要再过度依赖为各种 VIP 用户设置 VIP 镜像,而是专注于构建更良好的开箱即用体验。
- 不要再期望用户去使用
PYTORCH_TUNABLE_OPS这个充满 bug 的功能,这个功能会让用户浪费众多的时间。 - AMD 应该提交 MLPerf 训练 GPT-3 175B 的结果。MLPerf 是一个端对端的基准测试方法,他通过训练的时间来作为性能比较的唯一参考。
AMD 令人振奋的性能预期和实际落差
MI300X 在 2023 年底推出,具有令人振奋的纸面规格:
- 1307 TFLOPs 的 fp16 算力,优于 H100/H200 的 989 TFLOPs
- 5.3 TB/s 的 HBM 内存带宽,优于 H100 的 3.35 TB/s 和 H200 的 4.8 TB/s
- 192 GB 的 HBM 容量,优于 H100 的 80GB 和 H200 的 141GB
除此之外,MI300X 的部署 TCO 与 NVIDIA 的 H100/H200 相比也低得多。这不仅仅是因为 MI300X 平均销售价格 ASP 比 NVIDIA 更低,还因为 MI300X 通常使用更便宜的以太网络进行部署。
如下图所示,将 16k H200 IB 集群和 16k AMD MI300X 以太网集群相比,可以节省将近 40% 的成本,每台服务器的成本 272k 美元 vs 177k 美元。与使用 NVIDIA 的 Quantum-2 交换机相比,使用白盒以太网交换机可以节省大量的成本,但真正的区别在于更便宜的收发器。因为 NVIDIA 品牌收发器的成本是典型收发器成本的 2-3 倍。
从表面上看,MI300X 似乎是两全其美的:更高的性能和更低的 TCO,这为市场对 AMD 相当期待,纷纷调高对其的预期。然而,2023 年底发布 MI300X 后,在 2024 年 AMD 数据中心的表现远低于预期,1Q24 到 3Q24 的收益指导仅从 40 亿美元提升到 50 亿美元,远低于投资者 60 亿美元到 80 亿美元的预期。
为什么呢?问题的关键在于软硬件栈的适配,更好的开箱体验,集群 Scale out 性能。相比于 NVIDIA 的 CUDA 和 NCCL 等软件栈,AMD 的 ROCm 和 RCCL 等都还有很长的路要走。与此同时,在 AMD 追赶的同时,NVIDIA 的众多工程师正在同步的加班加点,让 NVIDIA 的软硬件适配最新的模型、最新的框架,进一步扩大其 CUDA 护城河的优势。
本文深度测评了 AMD 当前的硬件,试图去分析理解纸面规格和实际性能的差距,并愿意给 AMD 团队更多的反馈。
GEMM 矩阵乘性能
基于 Transformer 架构的 ChatGPT、LLaMA、Claude 等众多模型训练中主要的 FLOPS 占用来自于矩阵乘法 GEMM。矩阵乘法 GEMM 典型构成如下:输入为 (M, K) 形状的矩阵 A 和 (K, N) 形状的矩阵 B,输出为形状为 (M, N) 的矩阵 C。
实际的矩阵计算效果如下所示,这是线性代数课程的第一课:
接下来,我们测试了真实世界中会用到的矩阵乘法的性能,其形状来自于 LLaMA 70B4:
|
|
具体的测试方法中,使用 OpenAI 的 do_bench 函数5,使用 warmup=30,rep=300,输入和输出矩阵都随机初始化均值为 0 和方差为 1 的正态分布,这是因为正态分布最接近于匹配现代神经网络中权重和激活的实际分布。
输入张量的分布将影响 TFLOP/s 性能基准测试的结果6。这是因为,GPU 中功率消耗的主要来源之一是晶体管从0->1和从1->0切换状态。保持相同值(即1->1或0->0)的晶体管比切换状态的晶体管使用更少的功率。正如 Horace 在他的博客文章中所说的那样,
|
|
因此,将两个零填充张量相乘可以获得比将两个正态分布张量相乘更高的性能。在实际工作负载中,零填充张量是否相乘?绝对不是。
尽管 Nvidia 硬件中没有针对非结构化稀疏度的定制电路,但它比正态分布张量更快。乘以均匀分布比乘以正态分布更快,因为晶体管翻转更少。我们的实验显示了 H100/H200/MI300X 的类似结果,因为这些 GPU 是功率受限的。我们还有2:4结构化稀疏度的实验。2:4 从未在当前的训练配方中使用,但在某种程度上用于推理。我们将在我们的 MI300X vs H100 vs H200 Inferencing 文章中讨论 2:4 结构化稀疏度。
测试结果显示:对于 BF16,H100 和 H200 实现了大约 720 TFLOP/s,对应的其规格为 989 TFLOP/s。而 MI300X 仅达到 620 TFLOP/s,但是其纸面规格为 1307 TFLOP/s。
也就是说,尽管 MI300X 纸面规格比 H100 要高得多,但是实际测试起来,MI300 X 比 H100 和 H200 还要慢 14%。不仅如此,要达到这个 620 TFLOP/s,还需要使用 AMD 的定制容器镜像并且要设置很多特殊的环境变量。
进一步的,对于 FP8 而言,AMD 的表现更糟糕。在 e4m3 格式输入下,H100/H200 达到了其宣称的 1989 TFLOP/s 的 1280 TFLOPs,也就是纸面规格的 65%。对于 AMD,MI300X 达到了其宣称的 2615 TFLOP/s 的 990 TFLOP/s,也就是纸面规格的 38%。
因此,对于 FP8,MI300X 比 H100 慢 22%。
除了上述提到的性能不及预期,在测试中还碰到了 AMD 的软件 bug。其中的一个主要 bug 是 torch.matmul 和 F.linear API 在 AMD 上提供了不同的性能,这和大家预期的两者应该相同的表现并不一致。F.linear 要比 torch.matmul 要慢得多,问题的原因是 AMD 使用了两个不同的底层 GEMM 库 rocBLAS 和 hipBLASLt,其中 hipBLASLt 对 MI300X 进行了进一步优化, torch.matmul 使用了 hipBLASLt 但是 F.linear 摸清情况下并没有使用,而使用了未经优化的 rocBLAS 库。
上述的结论和 GitHub 上的一些流行的测试结果7 并不一致,他们宣称 AMD MI300X 的性能接近于 H100,如下图所示:
然而,这个测试存在着两个主要问题:
- 没有正确执行 L2 Cache 清理
- 只是在迭代过程中获取最大性能,而不是特定形状的中位数/平均值 TFLOP/s
在迭代之间没有 L2 缓存清除,基准测试不能准确反映真实世界的 GEMM 性能。
由于 TFLOP/s 会根据迭代而变化,因此需要使用至少 100 次迭代的平均值/中位数作为准确 GEMM 基准测试的基础。OpenAI 的 do_bench 默认提供 L2 缓存和平均值/中位数,因此我们建议工程师将其用于微基准测试。下面,我们将基准测试简化为伪代码,并对上述问题进行了评论。
HBM 带宽性能
众所周知,AMD MI300X 比英伟达 H100 和 H200 具有更好的内存带宽,提供 5.3TB/s 的带宽,而 H200 为 4.8TB/s,H100 为 3.35TB/s。改进的 HBM 内存带宽在推理中非常有用,有时在训练中也很有用。在训练中,如果用户有更多的 HBM 内存容量和内存带宽,他们可以设置更大的 batch size。虽然如果使用更大的 global batch size,在一定大小之后,模型将需要更长的时间来收敛。很容易以大的全局批量大小快速运行,但在高水平上,它会损害收敛的时间。
从我们的 HBM 内存带宽基准测试中,我们发现 MI300X 的内存带宽确实比 H200 和 H100 好得多。我们用张量在 Pytorch 中测试了内存带宽。copy_ 并使用行业标准的 OpenAI do_bench 来确保准确性。
正如我们即将发布的 H100 vs H200 vs MI300X 推理文章中看到的那样,内存带宽对于推理非常重要。
对 HBM 带宽测试的代码如下所示:
开箱即用的镜像
需要注意的是,上述 MI300X 能够达到 H100 75% 左右的性能的基础是,我们获得了来自 AMD 众多团队的支持,专门定制了一个 Dockerfile,通过指定了特殊版本的 PyTorch、hipBLASLt、TransformerEngine、Triton、FlashAttention 等软件,才能够获得还不错的性能。目前公开的镜像还有各种各样的问题。
相比之下,NVIDIA 的开箱即用的镜像好用得多,包含了分析和调试所需的一整套开发人员工具,如 Nsight 计算和 Nsight 系统。相比之下,AMD 不包括开箱即用的 OmniTrace 开发人员工具。
下面是 AMD 12 月 21 日的开发构建 docker 镜像。如您所见,它使用了许多非稳定的开发分支作为依赖项,如 hipBLASlt,AOTriton,ROCm 注意,并从源代码安装了包括 PyTorch 在内的所有内容,构建时间长达 5 小时。这些版本的依赖项甚至还没有合并到 AMD 自己的主分支中。99.9%的用户不会从源代码安装 PyTorch 及其所有依赖项,而是使用公共稳定的 PyPi PyTorch。
模型训练性能评测方法
评测模型训练性能的方法有很多,最准确的方法是跑一个经典的模型,并在 512-1024 规模的集群运行它们。这样的一个端到端的任务压测将会考虑到软硬件系统中的 HBM 带宽、HBM 容量、算力、网络、H2D/D2H 等方方面面的性能。
MLPerf GPT3 175B 基准训练任务是衡量模型训练到特定收敛的一个很好的测试方法。MLPerf 基准会考虑 global batch size 以及混合精度训练是否会给模型收敛带来负优化。不幸的是,由于缺乏用户友好的文档和指令,MLPerf 很难运行,并且为了性能通常会专门为 MLPerf 调优而采用一些自定义配置,这些配置一般用户不会采用。
值得注意的是,NVIDIA 已经提交了 11k H100 的 MLPerf Training 结果。但是 AMD 并没有提交过任何 MLPerf Training,并不用说 MLPerf GPT3 175B Benchmark。AMD 在内部运行过 MLPerf Training,可能因为结果很弱,所以没有提交。
当设计 SemiAnalysis 的基准测试时,我们希望反应普通用户的模型实现,因此采用了 torch.scaled_dot_product_attention API 来使用 Flash Attention backend,使用 PyTorch DDP 或者带有 torch.compile 的 FSDP。我们任务这是典型用户工作负载最具代表性的。
我们使用通过 PyTorch native 实现这些模型,以使其更加接近典型 ML 科学家,并使其易于使用单行代码能够执行。与 MLPerf 相比,我们基准测试的目标时尽可能简单地运行,但是同时能够很好的反应硬件的性能。请注意,由于我们没有考虑收敛时间,因此此基准测试略微偏向 AMD,因为我们在 AMD 上设置的 micro batch size 高于 Nvidia。当考虑到收敛时间时,AMD 结果将比所述的更差。
顺便说一句,许多 AI 从业者表示,他们没有使用 Megatron、NeMo 或 3D Parallelism,因为这些库具有高度的复杂性和缺乏灵活性,其刚性和复杂性使其在 ML 研究中的使用实际上是不可能的。请注意,就 3D Parallelism 而言,假设 Nvidia 和 AMD 的软件堆栈正常工作,Nvidia 和 AMD 都将获得更高的性能,这是 AMD 的一个很大的假设。AMDMegatron 是 Nvidia Megatron 的一个分支,在 GitHub 上不到 10 个 star,这意味着它可能没有很好地进行内部测试。提交 bug 报告需要额外几个月的时间才能让 AMD Megatron 为简单的模型工作。
对于我们的 SemiAnalysis 模型训练基准,我们将测试四个模型:
- 简单的 GPT1.5B DDP,因为我们认为这代表了小规模实验/消融在扩展到更大的模型大小之前的样子
- 我们测试了标准的 Llama3-8B 和 Llama3-70B 4 层代理作为流行模型性能的 baseline
- 我们测试了 Mistral 7B v0.1,它评估了硬件在增加一点复杂性时是否表现良好,因为 Mistral 使用滑动窗口注意力而不是标准的因果注意力。现代模型,如 ChatGPT、Claude、Genimi、o1、o3 不使用标准因果注意力,而使用复杂的注意力机制
现代 GPT/Llama/Transformer 模型是通过一遍又一遍地堆叠相同的 Transformer 层来构建的。因此,仅测量 4 层的性能是模型整体性能的一个很好的代理。
每个训练的 token 的模型 FLOP 由以下公式定义:
|
|
密度是相对于完整掩码的注意力稀疏程度。例如,因果注意力具有 50%的稀疏度,而滑动窗口注意力的稀疏度甚至更低。
请注意,最初我们的测试线束使用 6*params 而不是 6*non_input_embedding_params 这是计算每个令牌的模型 FLOP 的错误方法。此外,我们使用 FSDP 的方式还有另一个 bug。此后,我们更新了我们的测试方法,并追溯性地重新测试以及更新了 H100、H200、MI300X AMD 开发版本等所有软件版本的所有基准测试结果。
单节点训练性能
我们看到,对于所有型号,H100/H200 相对于 MI300X 公开发布/公开夜间发布/11 月 25 日构建,从源 VIP 镜像中获胜。有趣的是,MI300X 在较小的型号上表现不佳,如 GPT 1.5B 或任何使用非因果注意力层的型号,如 Mistral 7B v0.1。这是由于 FlexAttention 在截止日期时没有完全运行,而在 Nvidia GPU 上,它自 2024 年 8 月以来一直在工作。因此,在 MI300X 公开发布/公开夜间发布/Nov25VIP 构建方面,H100/H200 的 TFLOP/s 超过 MI300X 2.5 倍。
三个月前,当我们试图在 AMD 上进行 FP8 训练时,我们碰到了分段故障和硬错误。如果它真的有效,事实上,它比使用 BF16 的相同运行速度慢。我们与 AMD 的 FP8 团队以及 AMD hipBLASlt 团队合作解决了这个问题,他们创建了修复 MI300X FP8 性能的调整8。FP8 训练很重要,因为与 BF16 相比,它加快了训练速度,目前大多数前沿实验室都使用 FP8 训练,比如 DeepSeek V3 的训练。
经过多次修复,我们可以看到 MI300X 11 月 25 日的 Llama3-8B 和 GPT 1.5B 吞吐量与 H100 有些竞争力。像往常一样,H200 在这一类别中获胜。然而,对于 Llama3-70B 4 层,11 月 25 日 AMD 结果再次败给 H200。
对于具有非因果注意力层的 Mistral 7B,AMD 11 月 25 日的表现接近 H100 的一半。这表明,对于任何不是简单模型的东西,即使经过数月的调整,由于模型结构的轻微调整,AMD 仍然没有竞争力。许多前沿模型和 AI 初创公司正在使用复杂的注意力层来实现长上下文和高效注意力,但是,AMD 在这些方面仍然远远落后。
不幸的是,AMD 上的 FP8 训练仅适用于自定义镜像,例如我们 11 月 25 日的 VIP 映像和 12 月 21 日进展中的开发分支映像。当我们第一次开始尝试 AMD FP8 训练时,它比公共版本上的 AMD BF16 训练慢。
在 AMD 的开发进展中,我们看到在 Llama3 8B 上,它战胜了 H100,但仍然比 H200 的公共稳定软件版本慢。即使在 12 月 21 日的开发分支上,H200 的性能也完全超过了 MI300X。
有趣的是,MI300X 在非因果注意力层上表现不佳,就像 Mistral 7B v0.1 一样,即使对于它们的内部构建也是如此。Mistral 使用一些前沿模型使用的滑动窗口注意力。似乎如果你想训练一个不使用因果注意力的模型,AMD MI300X 会显著如果 NVIDIA。
虽然很多人在硬件之间进行性能比较,但大多数人不会开源他们的测试代码,也不容易重现。我们采取了开源方法,我们已经开源了我们的单节点训练基准,并使其易于运行,只需几行代码即可。
多节点训练性能
对于多节点,我们对 H100 的两个节点和 MI300X 的两个节点进行了基准测试。不幸的是,我们没有及时获得多节点 H200 部署的访问权限。
与 MI300X 相比,H100 在这一基准测试中再次大获全胜,H100 的速度快了 10-25%。随着您将更多节点协同工作添加到单个训练工作负载中,这种差距会扩大。这是一个已知的问题,AMD 正试图通过部署新的内部 400G NIC 来解决这个问题。
集合通信性能
单节点集合通信 Scale Up
Scale Up 对于 GPU 集群非常重要,因为它为前沿模型训练中使用的张量和专家并行性提供了极快的路径。因此,我们进行了基准测试来衡量 scale up fabric 的性能。
H100 和 H200 上的 scale up fabric 称为 NVLink,每个 GPU 提供 450 GByte/s 的带宽,并将 8 个 GPU 连接在一起。在 MI300X 上,放大结构称为 xGMI,在纸面上,它连接 8 个 GPU,每个 GPU 提供 448 GByte/s 的带宽。表面上,MI300X 的放大网络与 H100/H200 非常相似,性能接近,仅有 0.5%的纸张带宽差距。不幸的是,实际情况截然不同。
相比之下,由于 Nvidia 的 NVLink 使用交换拓扑,一个 GPU 可以以完整的 450GByte/s 与另一个 GPU 通信。此外,H100/H200 中的四个 NVSwitch 支持 in-network reduction 这称为 NVLink SHARP(NVLS),默认启用,是一种通过在交换机内部执行 collectives/reduction 来减少数据移动的技术。
我们将展示英伟达 H100/H200 和 AMD MI300 的 scale up 和 scale out 网络基准。我们将测试的集合通信原语是前沿 LLM 训练中使用的主要集合:all_reduce、all_gather、reduce_scatter 和 all_to_all。
all_reduce用于数据并行和张量并行all_gather用于 ZeRO/FSDP 并行(以及张量并行)reduce_scatter用于 ZeRO/FSDP 并行
由于 compute-communication overlap 的工作方式,现实世界的消息大小范围从 16MiB 到 256MiB,默认的 PyTorch DDP 大小为 25MiB(NVIDIA 的 MLPerf 11k H100 GPT-3175B 运行使用最大 200MiB 的消息大小9)。
我们还测试了 8GiB 和 16GiB,只是为了看看峰值 bus bandwidth 是多少,尽管这些消息大小在现实世界中没有使用。上面讨论的所有这些集合通信都用于 3D 并行和 FSDP/ZeRO 并行,这是训练前沿模型的常见技术。
我们看到英伟达在每个集体通信的所有现实世界消息中做得比 AMD 好得多。这并不奇怪,因为与 MI300X 的 7x64GByte/s xGMI 点对点拓扑相比,H100/H200 的 450GByte/s NVLink 切换拓扑具有 NVLS。
注意:这里 16GB 下 NV 的 bus bandwidth 超越了 450GB/s 理论上线是因为默认使能了 NVLS
要重现此测试,您可以使用我们的开源 ClusterMax-NCCL/RCCL 基准测试,我们开发了该基准测试,以便使用一行 Bash 轻松运行。
多节点集合通信 Scale Out
在英伟达的 H100/H200 和 MI300X 上,每个 GPU 都使用 400G NIC 通过 Scale Out 网络连接到其他节点,直接连接每个 GPU。H100/H200 参考设计通常使用 ConnectX-7 网卡 for InfiniBand NDR 或 BlueField-3 for Spectrum-X 以太网。Spectrum-X 是英伟达专门为 AI 工作负载构建的定制以太网解决方案。在 MI300X 上,参考设计建议使用 RoCEv2 以太网和 Broadcom Thor-2 NIC。
典型的 GPU 集群几乎总是需要比单层网络更多的层,因为单层网络只能支持 128 个 GPU(在 Broadcom 以太网或 Nvidia Spectrum X 以太网的情况下)和 64 个 GPU(用于 H100/H200 InfiniBand)。在这样的多层网络中,部署通常使用 8-rail optimized fat tree 网络,其中 8 个 GPU 中的每一个都连接到单独的交换机(这样的连接称为 rail)
不幸的是,与默认启用的 NVLink SHARP 不同,UFM/IB 子网管理器中默认不启用 InfiniBand SHARP。我们已经与许多 Neocloud、H100 集群运营商和 AI 前沿实验室进行了交谈,大多数人表示,由于 NCCL_TIMEOUT 增加以及安装和配置网络的困难,他们没有启用 SHARP。我们询问 NVIDIA AI 客户使用 InfiniBand SHARP,但他们拒绝具体回答。人们可以推测,如果 InfiniBand SHARP 在 AI 生产工作负载中有用,NVIDIA 营销部门会竭尽全力促进其成功部署。鉴于目前 InfiniBand SHARP 的采用显然有限,我们在这里展示了英伟达在启用和未启用 SHARP 时的集体通信表现。
对于一些基准测试,我们还收集了英伟达内部集群 Israel-1 上的英伟达 Spectrum-X 以太网数据。英伟达 Spectrum-X 用于 xAI 的 200kH100/H200 集群,可以支持 Spectrum-X 参考架构 1.2 版中多达 100k GPU 的集群,但可能支持多达 512k 非参考定制设计的 GPU。
我们还在测试谷歌云(GCP)H100 的内部以太网,以及部署在 AWS 内部以太网(称为 EFAv2/EFAv3)上的 AWS H100 和 H200。我们将在即将发布的 Collective Deep Dive 文章中分享结果,该文章将提供不同类型集体的可视化,解释不同的 NCCL 协议(SIMPLE,LL,LL128),不同的 NCCL 算法(NVLS,NVLSTREE,RING,TREE,COLNETDIRECT,COLNETCHAIN,PAT),以及集体如何在 GCP H100 以太网,AWS H100/H200 EFA,InfiniBand H100,Spectrum-X 等上运行。
下面我们展示了一个 32 个 GPU 的 all reduce 集体测试。您可以看到,与普通的 InfiniBand H100 和启用 SHARP 的 InfiniBand H100 相比,MI300X RoCEv2 排名垫底。简单来说, all reduce 性能差会导致 Scale Out 训练差。
如果你 Scale out(即增加)参与集合通信的 GPU 数量,MI300X 的性能会降低。可以想象,现代前沿训练是在至少 100,000 个图形处理器的集群上进行的。与 InfiniBand Non-SHARP 的基线相比,MI300X RoCEv2 在 16MiB 至 256MiB 的所有真实世界消息大小中的运行速度是 NV 的一半。
根据下表,英伟达 Spectrum-X 以太网性能非常接近 InfiniBand Non-SHARP 的性能,这是由于 Spectrum-X 与 NCCL 集合通信库的垂直集成以及使用良好的拥塞控制和自适应路由。AMD 正试图明年与即将推出的 Pollara 400G NIC 垂直集成,该 NIC 支持 Ultra Ethernet 超以太网,希望 AMD 与英伟达竞争。一如既往,英伟达不会停滞不前,到明年年底,它将准备好投入生产其 800G ConnectX-8 网卡,其 Line Rate 是 AMD Pollara NIC 的两倍。
AMD RCCL 是英伟达 NCCL 的一个分支。AMD 的 RCCL 团队和 AMD 的许多其他团队拥有的资源有限,没有足够的计算或人员来改善 AMD 生态系统。AMD 的 RCCL 团队目前稳定地访问不到 32 个 MI300X 用于研发,这具有讽刺意味,因为改善集合通信就是要访问许多图形处理器。坦率地说,这很傻,AMD 应该花更多的钱在他们的软件团队上,让他们能够访问更多的图形处理器。
这与英伟达的 NCCL 团队形成对比,NCCL 团队可以获得英伟达 11000 个 H100 内部 EOS 集群的研发资源。此外,英伟达还有 Sylvain Jeaugey,他是集合通信的专家。英伟达还有许多其他世界级的集合通信专家,不幸的是,由于薪酬和资源不太有吸引力,AMD 在很大程度上未能吸引集合通信库人才,而英伟达的工程师则相反,由于 RSU 价值的升值,工程师每年的收入超过 100 万美元并不罕见。
为了帮助缓解这些问题,TensorWave 和 SemiAnalysis 目前正在与 AMD RCCL 团队合作,以提高集体性能。TensorWave 慷慨赞助 AMD 一个中型集群,以帮助 RCCL 团队有更多的资源来完成他们的工作。购买了许多 GPU 后,Tensorwave 反而不得不给 AMD 图形处理器让他们修复软件,很疯狂 🐶。
另一个值得注意的趋势是,对于非 SHARP 网络,随着 GPU 数量的翻倍,all reduce 的集合通信性能会呈对数降低。相比之下,对于 SHARP,速度/完成时间保持不变。我们有多达 1,024 个 H100 的结果表明,IB SHARP 的 all reduce 在任何数量的 GPU 上都是恒定时间。我们将在即将发表的 Collective Deep Dive 文章中发表这一点。
对于 all-gather、all-to-all 和 reduce-scatter,MI300X 比 InfiniBand 慢 2-4 倍。不幸的是,我们无法访问 Spectrum-X 或 InfiniBand SHARP 基准数据以获取所有 all-gather 或 reduce-scatter。
下面,我们提供了我们的 nccl/rccl 基准测试脚本。不幸的是,由于集群特定设置的性质,它不像单行代码那么简单。它确实需要你遵循 nccl/rccl 和 nccl-test/rccl-test 的 README.md 才能正常运行。在 AWS 和 Google Cloud 上,您可能还需要安装自定义 nccl plugin。
H100/H200/MI300X 网络 BoM 分析和 TCO 性能
使用更高基数的 51.2T 交换机可以简化网络部署,因为它可以支持给定网络层大小的更大 GPU 集群大小。而 25.6T 基数交换机,如 Nvidia Quantum-2 QM9700,可以连接到 2 层网络的 2,048 个 GPU(即连接到 GPU 的叶层和连接叶交换机的脊层),使用基于 Tomahawk 5 ASIC 或 Spectrum-X SN5600 的 51.2T 基数交换机,可以连接多达 8,192 个 GPU。
对于相同的8,192个 GPU 部署,使用25.6T 基数交换机需要3层网络和640个交换机,而使用51.2T 基数交换机只需要2层网络和192个交换机,以及更少的收发器。基于以太网的后端网络部署相对于基于 InfniBand 的部署的优势在于,当今市场上广泛部署的 InfiniBand 最高基数交换机仅为25.6T,而以太网可以选择使用51.2T 基数交换机。
我们首先分析 Neocloud Giant 上1,024个 GPU InfiniBand 部署的网络 BoM,因为这是许多 Neocloud 上发现的典型集群大小。假设后端结构为非阻塞网络,但同时处理存储网络(即融合结构)的2:1 收敛比前端结构,我们估计总集群网络 BoM 为520万美元,或每个8-GPU 服务器 40,705美元。
转向相同大小但使用以太网作为后端结构的部署,我们估计总集群网络成本为300万美元,即每台服务器 23,816 美元。这比基于 InfiniBand 的部署低得多,因为交换机和收发器的成本更低。1,024个 GPU 的集群大小需要使用25.6T 或51.2T 基数交换机的两层网络,因此使用基于 Arista Tomahawk 5 的 51.2T 基数交换机进行以太网部署并不能从较低的收发器数量或简化的网络拓扑中节省成本。
16,384个 GPU 的集群规模超过了使用51.2T 基数交换机时最大的两层节点数量8,192个 GPU,这需要转向使用三层网络作为后端结构。这大大增加了每台服务器的成本,达到34,214美元,因为交换机和收发器数量更多。我们注意到,由于需要额外的工作来调整 RoCEv2以在集合通信上良好运行,因此很少有 Neocloud 巨头会选择在以太网上部署。
超级规模化者可以利用更强的议价能力,以比 Neocloud Giant 更便宜的成本部署网络。使用基于 Whitebox Tomahawk 5的交换机,而不是来自 Arista 等提供商的品牌以太网交换机,可以节省40-50%的交换机 ASP。超级规模化者还可以以更低的价格购买收发器,在这里节省40-50%的单位定价。我们计算每台服务器的总网络资本支出为18,677美元。
价格最高的选择是部署基于 InfiniBand 的网络的 Neocloud Giants。在这里,他们将使用具有最高 ASP 的英伟达交换机,并购买英伟达品牌的收发器——在给定收发器速度下也是最高 ASP。
最后,我们分析了一个由 Hyperscaler 部署的16k 基于 InfiniBand 的集群。使用与 Neocloud 巨头相同的网络部署,仅更改网络设备的 ASP,我们看到每台服务器的总网络成本为47,470美元。
从集群的总前期资本支出来看,可以看出基于以太网的网络成本低得多,这与 AMD MI300X 服务器相对于 Nvidia H100或 H200服务器的低成本一样,是降低总拥有成本的一个因素。
综合考虑,Neocloud Giant 可以以每小时1.23美元的总成本部署 MI300X,而 H100和 H200的总成本分别为1.65美元和1.66美元,而 Hyperscaler 可以以每小时1.09美元的价格部署 MI300X,而 H100和 H200的总成本分别为1.5美元和1.51美元。
但是,如果我们将每小时的成本除以提供的训练有效 FLOP/s(由我们的 FP8单节点训练基准确定),以每有效 PFLOP/s 的单位$/hr 获得训练计算成本,我们发现 MI300X 的成本优势消失了。对于 Neocloud 巨头,MI300X 的计算成本为每有效 PFLOP/s 2.47美元/小时,高于 H200,而对于 Hyperscaler,MI300X 的计算成本为每有效 PFLOP 2.1美元/小时,高于 H200的1.85美元。
基准测试预热/重复效果
GPU 功率有限。这意味着它们永远无法在允许的热设计功率(TDP)内维持其最大时钟频率。因此,GPU 在早期迭代中总是比后期迭代更快,因为随着时间的推移,GPU 需要减速到稳定的频率。有趣的是,AMD 往往比 H100/H200更早达到稳定状态。
话虽如此,我们认为 warmup=30,rep=200 是基准测试的最佳设置,因为在实际工作负载中,用户不会被 compute bound 的 GEMM 执行 1000次。在实际使用中,执行 GEMM 后通常会继续执行 LayerNorm 或 softmax 或潜在的通信。所有这些 kernel 都允许 GPU 来 cooldown。
一个基准测试技巧是放弃预热并进行多次迭代以防止功率限制。另一种方法是设置基准测试,使 GPU 不需要三次迭代即可达到其峰值实现频率,而是锁定 GPU,使其立即以最高频率启动。然而,这并不会导致准确的锁定,反而会导致 GPU 达到稳定频率所需的时间更长。我们在基准测试时放弃这些技巧,因为在生产训练工作负载中,用户不会锁定 GPU 频率,这可能会导致时钟速度进一步降低。
sudo nvidia-smi -i 0 --lock-gpu-clocks=1830,1830
BF16 vs FP16
我们实验的另一个观察结果是,BF16的 TFLOP/s 往往比 FP16略高,尽管从业者已经期望非常相似的性能。在 NV IDIA 宣称的 H100的吞吐量中,BF16和 FP16都达到了989.5 TFLOP/s。我们的实验显示,对于 H100/H200/MI300X 中的每一个,BF16比 FP16略有优势,这是有道理的,因为 GPU 的功率有限。
为什么性能有差异?BF16是 E8M7(8个指数位和7个尾数位),而 FP16是 E4M10,这意味着 BF16的尾数位比 FP16少。对于浮点乘法,指数被相加,尾数被相乘。简单来说,加法电路通常使用 $O(n)$ 晶体管,而乘法电路使用 $O(nlogn)$ 到 $O (n^2)$,其中 n 是位数。
如前所述,GPU 功率受限。BF16的尾数位比 FP16少,因此使用的晶体管更少。这导致更高的时钟频率和更高的 TFLOP/s。逻辑延伸到比较 FP8 e4m3和 FP8 e5m2,但差异较小。
-
MI300X vs H100 vs H200 Benchmark Part 1: Training – CUDA Moat Still Alive, https://semianalysis.com/2024/12/22/mi300x-vs-h100-vs-h200-benchmark-part-1-training ↩︎
-
Dylan Patel talked with Lisa Su, https://x.com/dylan522p/status/1871287937268383867 ↩︎
-
Lisa Su’s response and George Hotz’s comment, https://x.com/realGeorgeHotz/status/1871556257410334723 ↩︎
-
https://github.com/pytorch-labs/float8_experimental/blob/fe6e08c867abf56b1acd0f34473c69cde624f0a3/benchmarks/bench_matmul.py#L57 ↩︎
-
https://triton-lang.org/main/python-api/generated/triton.testing.do_bench.html#triton.testing.do_bench ↩︎
-
https://www.thonking.ai/p/strangely-matrix-multiplications ↩︎
-
https://github.com/stas00/ml-engineering/blob/master/compute/accelerator/README.md#maximum-achievable-matmul-flops-comparison-table ↩︎
-
https://github.com/mlcommons/training_results_v4.1/blob/b87b9e396f771345d4ef122ba33456304f15228d/NVIDIA/benchmarks/gpt3/implementations/eos-dfw_n1452_ngc24.04_nemo/config_common.sh#L69 ↩︎
-
No backlinks found.