文章标题: NetZIP: 算法/硬件协同设计的网络内无损压缩,用于分布式大模型训练
作者/机构:
- JINGHAN HUANG, University of Illinois Urbana-Champaign
- HYUNGYO KIM, University of Illinois Urbana-Champaign
- NACHUAN WANG, University of Illinois Urbana-Champaign
- JAEYOUNG KANG, University of Illinois Urbana-Champaign
- HRISHI SHAH, University of Illinois Urbana-Champaign
- EUN-KYUNG LEE, IBM Research
- MINJIA ZHANG, University of Illinois Urbana-Champaign
- FAN LAI, University of Illinois Urbana-Champaign
- NAM SUNG KIM, University of Illinois Urbana-Champaign
核心问题: 在分布式大模型训练中,GPU之间交换大量梯度和激活值所需的通信时间主导了总训练时间。尤其是在使用网络带宽有限的经济型计算集群时,这一问题更为严重。
现有方案的局限性:
1. 有损压缩: 对梯度进行有损压缩通常需要更多训练迭代才能达到同等模型精度;而对激活值进行有損压缩,可能导致收敛效果不佳甚至失败。
2. 无损压缩: 现有的无损压缩算法(如LZ4、Snappy、Zstd、Deflate)在当前计算平台(CPU、GPU、SmartNIC)上,无法充分减小梯度和激活值(bfloat16格式)的数据量,其压缩和解压带来的巨大延迟开销远超过了因数据量减少而节省的传输时间,反而导致总通信时间增加。
研究目标与创新点 (NetZIP):
本文提出了 NetZIP,一种算法/硬件协同设计的方案,旨在通过网络内无损压缩梯度和激活值来解决上述挑战。NetZIP 包含两个核心组件:
NetZIP-算法:
创新方法: 提出了两种转换算法,使轻量级无损压缩算法能更有效地压缩梯度和激活值。
效果: NetZIP-算法使LZ4等轻量级算法对Llama-3 70B、GPT-3 175B和Llama-3 405B模型的梯度和激活值的压缩率比重量级标准无损压缩算法(如Zstd)高出40-63和43-75个百分点。
NetZIP-加速器:
总体贡献: NetZIP通过更高的压缩率和更低的压缩/解压延迟,最终将大模型训练时间降低了35%。
大模型训练的特点。训练像Llama-3 405B这样的大模型通常需要数千万次迭代,每次迭代都包括一次前向传播(计算预测损失)和一次反向传播(根据梯度调整模型权重)。为了在可接受的时间内完成训练,现有的分布式训练框架,如Megatron-LM【索引54,Megatron-LM: Training Multi-Billion Parameter Language Models Using Model Parallelism,2019,arXiv】和DeepSpeed【索引52,DeepSpeed: System Optimizations Enable Training Deep Learning Models with Over 100 Billion Parameters,2020,SIGKDD】,主要利用三种并行方式来扩展训练规模。
三种并行方式。
网络基础设施的多样性。用于分布式大模型训练的网络基础设施差异很大。为了在成百上千个GPU之间快速交换大量计算出的梯度和激活值(统称为中间值),顶级的超大规模集群为每个计算节点配置了多个400 Gbps的网络链路【索引13,RDMA over Ethernet for Distributed AI Training at Meta Scale,2024,SIGCOMM】、【索引49,Alibaba HPN: A Data Center Network for Large Language Model Training,2024,SIGCOMM】。然而,公有云中经济实惠的集群仍然依赖于中低带宽的网络链路。例如,CloudLab集群【索引9,The Design and Operation of CloudLab,2019,USENIX ATC】提供了一个低成本甚至免费的模型训练环境,但其网络链路带宽通常在10 Gbps到常见的50 Gbps之间【索引26,Sol: Fast Distributed Computation Over Slow Networks,2020,NSDI】、【索引51,CASSINI: Network-Aware Job Scheduling in Machine Learning Clusters,2024,NSDI】。
典型的经济型集群架构。图1展示了这样一个集群,它采用了常见的数据中心网络拓扑(DCN+【索引49,Alibaba HPN: A Data Center Network for Large Language Model Training,2024,SIGCOMM】)和一个基于高端CloudLab实例(nvidiagh【索引8,CloudLab Portal Hardware,accessed in 2025】)的计算节点架构。网络拓扑是一个分层的三层Clos网络,由架顶(ToR)和核心(spine)交换机组成一个无阻塞的网络结构。每个节点包含一个服务器级CPU、一个上一代或当前代的GPU(例如,Intel Xeon CPU搭配NVIDIA A100或H100 GPU)以及一个Mellanox ConnectX-5 50 Gbps NIC。
图 1: 常见的数据中心网络拓扑,并放大了基于CloudLab实例的计算节点视图。
网络流量特征与瓶颈。分布式训练框架(§2.1)在交换梯度和激活值时会产生高度集中的大网络包突发流量【索引13,RDMA over Ethernet for Distributed AI Training at Meta Scale,2024,SIGCOMM】。这种网络流量模式给网络基础设施带来了巨大压力,导致严重的拥塞,从而大幅增加了GPU之间交换梯度和激活值所需的时间。当网络带宽受限时,这个问题尤其严重,这在公有云的经济型集群中很常见。因此,分布式训练时间显著增加,因为GPU在收到这些梯度和激活值之前无法继续计算后续阶段或层。
主流无损压缩算法原理。无损压缩算法通常基于字典匹配(如LZ77【索引68,LZ77 and LZ78,accessed in 2025,Wikipedia】)和/或熵编码(如霍夫曼编码【索引18,A Method for the Construction of Minimum-Redundancy Codes,1952,Proceedings of the IRE】),如图2所示。为了在压缩吞吐量、延迟和压缩率之间取得平衡,已经开发了多种无损压缩算法。
不同算法的权衡。例如,LZ4【索引33,lz4,accessed in 2025,GitHub】和Snappy【索引14,snappy,accessed in 2025,GitHub】采用字典匹配。它们使用一个滑动窗口来检测和编码重复的字符串,用对该字符串先前位置的引用来替换后续出现。这种方法实现了高压缩/解压吞吐量,但压缩率相对较低。相比之下,Deflate【索引34,zlib,accessed in 2025,GitHub】和Zstd【索引11,zstd,accessed in 2025,GitHub】结合了字典匹配和熵编码。在使用LZ77定位重复字符串后,霍夫曼编码为更频繁出现的字符分配更短的比特序列。
图 2: 代表性的无损压缩算法。
实验设置。为了评估有限网络带宽在公有云经济型集群中对分布式大模型训练的影响并保证可复现性(§2.2),我们使用了SimAI【索引63,SimAI: Unifying Architecture Design and Performance Tunning for Large-Scale Large Language Model Training with Scalability and Precision,2025,USENIX Symposium Yiding Wang, Decang Sun, Kai Chen, Fan Lai, and Mosharaf Chowdhury. 2023. Egeria: Efficient DNN Training with Knowledge-Guided Layer Freezing. In Proceedings of the European conference on computer systems (EuroSys). 851–866. Zeqin Wang, Ming Wen, Yuedong Xu, Yipeng Zhou, Jessie Hui Wang, and Liang Zhang. 2023. Communication compression techniques in distributed deep learning: A survey. Journal of Systems Architecture 142 (2023), 102927.】,一个由工业界开发的全栈、高精度模拟器。SimAI能够全面模拟大规模分布式模型训练的整个生命周期,不仅能模拟AllReduce和ReduceScatter等集体通信原语,还能模拟网络交换机和NIC的拥塞等底层网络行为。我们将其配置为模拟一个包含512个节点的集群,每个节点包含一个NVIDIA H100 GPU,网络拓扑如图1所示。表1描述了由SimAI评估的三个模型——Llama-3 70B、GPT-3 175B和Llama-3 405B——以及它们的分布式训练策略、参数数量、层数和注意力头数。为了在多个GPU上高效训练,这些模型利用了所有三种并行方式:TP、DP和PP(§2.1)。DP由WS/(PP×TP)计算得出,其中WS表示世界大小,即使用的GPU总数。这些模型使用Megatron-LM框架进行训练。
表 1: SimAI的模型配置。WS表示世界大小,即使用的GPU数量。
通信时间与网络带宽。我们首先研究网络带宽对每次训练迭代的节点间通信时间的影响,通过改变网络链路带宽并按比例缩放网络交换机带宽(统称为网络带宽)。图3(a)显示,对于Llama-3 70B、GPT-3 175B和Llama-3 405B,当网络带宽低于200 Gbps时,每次训练迭代的通信时间开始超线性增长。同时,图3(b)显示,在公有云的2561个实例中,有2350个实例的网络带宽为50 Gbps或更低。这些实例通过CloudLab中相对中等带宽的网络链路连接【索引8,CloudLab Portal Hardware,accessed in 2025】。此外,在Amazon EC2的850种实例类型中,约有800种的网络带宽低于50 Gbps【索引3,Amazon EC2 Instance Types,accessed in 2025】。这表明,网络带宽有限的实例在公有云中很常见,并非每个人都能轻易获得网络带宽超过100 Gbps的高端实例。基于这一观察,当网络带宽降至50 Gbps时,Llama-3 70B、GPT-3 175B和Llama-3 405B的每次训练迭代通信时间分别比400 Gbps网络带宽时增加了2.2倍、2.0倍和1.7倍。值得注意的是,每次训练迭代的总时间主要由通信时间决定,对于Llama-3 70B和GPT-3 175B,当带宽超过200 Gbps时,通信时间不再减少;对于Llama-3 405B,则是在400 Gbps以上。这是由于通过网络链路和交换机交换梯度和激活值的固定延迟所致,该延迟由给定节点数的网络拓扑决定。
图 3: (a) 不同网络带宽下每次训练迭代的通信时间,以及 (b) 公有云中实例的NIC带宽分布【索引8,CloudLab Portal Hardware,accessed in 2025】。
总训练时间与有损压缩。为了在有限的网络带宽下减少通信时间,我们可以压缩节点间交换的梯度和激活值。图4(a)显示,假设网络带宽为50 Gbps,通信时间随着压缩率的降低(即梯度和激活值的数据量减少)而减少。例如,当压缩率降至75%时,Llama-3 8B每次训练迭代的通信时间减少了23%。然而,当我们对梯度应用最新的有损压缩技术——top-k梯度稀疏化【索引55,Optimus-CC: Efficient Large NLP Model Training with 3D Parallelism Aware Communication Compression,2023,ASPLOS】——时,我们观察到当压缩率低于50%时,Llama-3 8B的总训练时间反而增加了。这是因为使用这种有损压缩进行训练,需要更多的迭代次数才能达到与无压缩训练相同的模型精度,因为每次迭代都会产生精度损失,这一观点也得到了近期工作的支持【索引10,Accuracy is Not All You Need,2024,arXiv】、【索引20,Compressing LLMs: The Truth is Rarely Pure and Never Simple,2023,arXiv】、【索引57,Communication Compression for Decentralized Training,2018,Advances in Neural Information Processing Systems】、【索引71,Beyond Perplexity: Multi-dimensional Safety Evaluation of LLM Compression,2024,arXiv】。
要点-1: 先前的有损压缩工作可以显著减少每次训练迭代的通信时间,但可能会增加总训练时间。
激活值压缩的挑战。在拥有数百个节点的分布式大模型训练中,通常会利用所有三种并行方式——DP、TP和PP——这不仅需要交换梯度,还需要交换激活值。我们观察到,对于Llama-3 70B、GPT-3 175B和Llama-3 405B,由交换激活值贡献的通信时间百分比分别为27%、21%和49%。尽管交换激活值贡献了很大一部分通信时间,但无论是过去的工作【索引28,A Network-Centric Hardware/Algorithm Co-Design to Accelerate Distributed Training of Deep Neural Networks,2018,MICRO】还是最新的工作【索引55,Optimus-CC: Efficient Large NLP Model Training with 3D Parallelism Aware Communication Compression,2023,ASPLOS】,都未能成功地将有损压缩应用于激活值。这是因为与梯度不同,激活值中即使是微小的精度损失也会对模型收敛产生有害影响。
要点-2: 激活值的交换也占了通信时间的很大一部分,但有损压缩不能应用于激活值。
图 4: (a) 每次训练迭代的通信时间 和 (b) 总训练时间,均相对于无压缩情况进行归一化,随压缩率变化。虚线柱表示收敛失败。
本节首先评估了四种流行的无损压缩算法的压缩率,考虑到有损压缩算法可能会增加模型的总训练时间或导致其无法收敛。然后,我们测量了在三种代表性计算平台(CPU、GPU和SNIC)上使用这些算法压缩和解压中间值的吞吐量和延迟。
实验设置。为了评估压缩和解压的吞吐量和延迟,我们配置了一台服务器,配备了第四代Intel Xeon可扩展处理器、NVIDIA H100 GPU和NVIDIA BlueField-2 SNIC(详见6.1节)。我们使用NVIDIA BlueField-2 SNIC,因为最新的NVIDIA BlueField-3只有一个解压加速器。请注意,我们使用BlueField-2专门用于分析压缩加速器和节点间通信。然而,我们的主要关注点是低带宽NIC(例如50 Gbps),如第3节所讨论和分析的。我们首先收集每个节点在每次训练迭代期间根据表1中所示的DP、TP和PP发送给其他节点的梯度和/或激活值。CPU使用标准的C++库运行LZ4【索引33,lz4,accessed in 2025,GitHub】、Snappy【索引14,snappy,accessed in 2025,GitHub】、Zstd【索引11,zstd,accessed in 2025,GitHub】和Deflate【索引34,zlib,accessed in 2025,GitHub】,而GPU使用nvCOMP【索引45,NVIDIA nvCOMP,accessed in 2025】库运行它们。SNIC只支持Deflate硬件加速器,使用DOCA Compress库【索引41,DOCA SDK Document,accessed in 2025】运行Deflate。对于基于CPU的压缩和解压,我们将微调迭代中的梯度和激活值分成128 kB的块,并使用所有40个CPU核心并行压缩和解压每个块,以最大化吞吐量。我们还忠实地捕捉了端到端的节点间通信时间,包括梯度和激活值的压缩、通过网络链路的传输以及解压。例如,基于CPU的压缩路径如下:发送方GPU内存→CPU内存→CPU压缩→CPU内存→NIC→接收方NIC→CPU内存→CPU解压→CPU内存→GPU内存。
压缩率。表2展示了LZ4、Snappy、Zstd和Deflate对Llama-3 70B、GPT-3 175B和Llama-3 405B的梯度和激活值的平均压缩率。结果显示,依赖字典匹配的LZ4和Snappy完全无法减少梯度或激活值的数据量(即压缩率为100%),甚至由于存储字典的开销而增加了数据量。这是因为LZ4和Snappy使用带滑动窗口的字典匹配来查找重复的8位字符串,而bfloat16格式的梯度和激活值的低8位(被LZ4和Snappy识别为字符)根据我们的分析(§5.1)大多是随机的,这使得很难找到重复的长8位字符串。与此同时,同时使用字典匹配和霍夫曼编码的Zstd和Deflate将梯度和激活值的数据量减少了20%–28%,因为霍夫曼编码可以识别出它们之间一些常见的上8位。平均而言,Zstd对梯度和激活值的压缩效果比Deflate好2-7%,并且对激活值的压缩效果比梯度好7%。Deflate和Zstd比LZ4和Snappy提供更高压缩率,以及激活值比梯度更易压缩的原因将在后面讨论(第5.1节)。
表 2: 梯度(G)和激活值(A)的平均压缩率,以及它们在节点间交换的梯度和激活值总量中所占的百分比(括号内)。
吞吐量和延迟。表3提供了使用CPU、GPU和SNIC对Llama-3 70B、GPT-3 175B和Llama-3 405B的中间值(即梯度和激活值)进行压缩和解压的吞吐量和延迟。由于LZ4和Snappy完全无法压缩梯度和激活值,无论它们在CPU还是GPU上运行,都只会因为增加的压缩/解压延迟而增加节点间的通信时间。然而,我们评估了它们的吞吐量和延迟,因为在NetZIP-算法的帮助下,它们可以高度压缩梯度和激活值(第5.1节)。
各平台性能分析。基于CPU的LZ4与基于CPU的Snappy提供了相似的吞吐量,这里的“吞吐量”指的是压缩/解压吞吐量的平均值,除非特别说明。凭借更简单的算法,基于CPU的LZ4比Zstd的吞吐量高1.6-1.8倍,比基于CPU的Deflate高6.6-11.0倍。基于GPU的Snappy比LZ4的吞吐量高1.1-6.1倍,因为它被设计为能更好地利用并行性,将给定的梯度或激活值分成更小的块(例如32 kB)并并行压缩它们。这种更高的吞吐量通常以牺牲压缩率为代价(即在更小的块中找到重复字符串的概率更低),尽管在这项工作中LZ4和Snappy都完全没有压缩梯度和激活值。基于GPU的Snappy比基于GPU的Zstd的吞吞吐量高9.6-13.6倍,比基于GPU的Deflate高7.8-9.6倍。最后,基于SNIC的Deflate比基于GPU的Deflate提供了高7-12%的压缩吞吐量,但解压吞吐量低25-88%。
GPU压缩的优缺点。在所有无损压缩算法中,GPU利用其大规模并行性,可以提供比CPU高一个数量级的压缩和解压吞吐量。基于GPU的LZ4、Snappy、Zstd和Deflate的吞吐量分别是基于CPU的LZ4、Snappy、Zstd和Deflate的2.9倍、8.3倍、1.2倍和7.8倍。然而,使用nvCOMP进行基于GPU的压缩和解压存在一个缺点。例如,在GPU将梯度和/或激活值发送到NIC之前,它必须从GPU内存中读取未压缩的梯度和/或激活值,然后将压缩后的梯度和/或激活值写回GPU内存,之后GPU才能将它们发送到NIC。在这种情况下,如果没有复杂的优化,GPU可能需要同时在GPU内存中存储未压缩/压缩的梯度和/或激活值,从而减少了可用于存储张量和/或层的可用内存容量。在单个GPU内存容量的限制下,基于GPU的压缩可能需要更多的GPU来进行分布式大模型训练。
总通信时间分析。表3还展示了两个节点之间的总通信时间,及其在中间值压缩、压缩后的中间值通过50 Gbps网络链路传输以及这些值的解压方面的分解。压缩/解压延迟与压缩/解压吞吐量成反比,而网络链路延迟与压缩率成正比。Zstd和Deflate减少了网络链路延迟,这与其对梯度和激活值的压缩率直接相关(表2),尽管这个延迟测量没有捕捉到因网络拥塞减少而带来的通信时间减少。尽管如此,即使我们为Zstd和Deflate选择了能提供最低总通信时间的平台(即GPU),我们观察到,与不进行压缩相比,无损压缩分别使Llama-3 70B、GPT-3 175B和Llama-3 405B的总通信时间增加了95%、66%和74%,这是因为在压缩率不高的情况下,压缩和解压带来了很长的延迟。
要点-3: 使用标准的无损压缩算法和平台来压缩和解压中间值,会因为压缩/解压延迟过长而增加通信延迟,而不是减少它,而压缩率最多也只是中等水平。
表 3: 在CPU、GPU和SNIC上运行的无损压缩算法的吞吐量和延迟。
本节我们将介绍NetZIP-算法和NetZIP-加速器。具体来说,NetZIP-算法在比特和值级别上对梯度和激活值进行转换,以帮助标准无损压缩算法更好地压缩它们,利用了它们独特的数值特性。其次,NetZIP-加速器是一个“线上传输中处理”的加速器,经过优化以在硬件资源受限的情况下实现高压缩/解压吞吐量。
字节和比特分组以实现更多字典匹配。图5绘制了(a)bfloat16格式的梯度和激活值中每个比特索引处1的百分比,以及(b)它们在Llama-3 70B微调迭代期间的数值大小。在bfloat16格式中,第15位、第14-7位和第6-0位分别表示符号(S)、指数(E)和尾数(M)。当指数满足以下条件:$1 \le E \le 254$时,bfloat16数表示的数值计算为$(-1)^S \times 2^{E-127} \times (1 + M/2^7)$。梯度和激活值的尾数位中1的概率约为50%,这使得它们几乎不可压缩。相比之下,梯度和激活值的指数位中1的概率很高,因为它们中很大一部分通常接近于零(例如,$1.07169 \times 10^{-8} = 0_01100100_0111000$),它们的数值大小相似且范围有限。具体来说,梯度的指数位通常是011????????,其中?是1或0,概率随机,范围从96到127。这导致梯度的范围从$2^{96-127}$(即$2^{-31}$)到$2^{127-127}$(即$2^0$)。激活值的指数位通常是0111???????,范围在112到127之间,对应于$2^{-15}$到$2^0$范围内的激活值。换句话说,指数位比尾数位具有更低的熵,使其更易于压缩。
图 5: (a) Llama-3 70B所有梯度和激活值中每个比特位为1的概率,以及 (b) 梯度和激活值的数值分布。
分组算法的提出。基于上述观察,我们提出了NetZIP-算法。该算法首先将bfloat16格式的梯度或激活值的高位字节作为一个字节组(图6左侧),并使用标准无损压缩算法对其进行压缩。剩余的低位字节形成另一个字节组,我们可以选择不压缩它,因为它们压缩效果不佳,只会增加压缩延迟。以这种方式对高位字节进行分组,即使是之前未能压缩它们的LZ4和Snappy,也能对梯度和激活值实现更多的压缩。这是因为它有助于向LZ4和Snappy暴露更多重复的字符串以进行字典匹配,而不会被每隔一个字节就出现的随机值(即低位字节)所阻碍(§2.3)。此外,这种方法对于重量级的无损压缩算法(如Zstd和Deflate)也有效,因为它们同时采用字典匹配和霍夫曼编码。更进一步,我们可以将分组粒度细化到每个梯度或激活值中具有相同比特索引的1个比特,即总共16个比特组(图6右侧)。然而,CPU和GPU都难以处理这种比特级操作,而ASIC或FPGA则擅长此道。
图 6: 16位bfloat数值的字节分组(左)和比特分组(右)。不同数值的相同字节和比特位置被组合在一起。
增量值压缩。直观上,尤其是在微调期间,训练模型会随着训练迭代逐步调整其激活值。也就是说,一个给定的梯度或激活值在迭代$i$和$i+1$之间的差异(或增量)应该很小——通常比原始值本身更接近于零。这反过来又使得无损压缩算法能够找到更多重复的字符串,从而实现更高的压缩率。图7绘制了(a)梯度和激活值增量中每个比特位为1的百分比,以及(b)Llama-3 70B微调迭代期间增量值的分布。为了获得每个梯度或激活值的增量值,我们用它在第$i+1$次训练迭代的值减去它在第$i$次训练迭代的值。结果显示,梯度和激活值的增量在其指数和尾数中分别有更多的1和0。
图 7: (a) 梯度和激活值的增量值中每个比特位为1的概率,以及 (b) 微调迭代期间增量值的分布。
增量值压缩的实现。基于上述观察,我们提出了一种增量值压缩方法。理想情况下,我们存储前一次迭代的所有梯度和激活值,用当前迭代的相应值减去它们,然后压缩这些差值。我们称之为“真实增量压缩”。然而,这样做是不切实际的,因为它将使存储前一次和当前迭代的梯度和激活值所需的内存容量加倍。作为替代方案,考虑到梯度和激活值的分布很窄,例如,我们选择一个单一的梯度作为过去迭代的基准值,并用当前迭代的每个梯度值减去该基准值。例如,通过使用最小激活值作为基准值,我们可以如下获得激活值的增量值:$delta_value[i+1][j] = activation[i+1][j] - min(activation[i])$,其中$i$和$j$分别代表迭代次数和激活向量索引。最后,由于我们观察到不同层的梯度和激活值范围不同(图7(b)),我们为每一层使用一个基准值来获取该层梯度和激活值的增量。这消除了存储前一次迭代所有中间值的内存容量限制。
SNIC压缩加速器的局限性。使用NVIDIA BlueField-2 SNIC,我们已经证明,与不进行压缩的训练相比,采用压缩的训练最终会增加节点间的通信时间,原因是压缩/解压延迟过长。除了加速器本身的固有压缩/解压延迟外,其他因素也导致了长延迟。具体来说,Deflate加速器连接到SNIC的嵌入式PCIe交换机。因此,为了进行压缩,SNIC的Arm CPU必须通过其内部PCIe互连将未压缩的中间值从其内存传输到加速器的内存,这与GPU等其他PCIe连接的加速器情况相同。由于压缩加速器的内存容量有限,它一次只能压缩128MB的梯度或激活值。压缩后,SNIC的Arm CPU需要将压缩后的梯度和激活值复制到其内存,然后才能将它们发送到另一个节点。我们的评估显示,由于SNIC的Arm CPU内存带宽有限,这些传输占了压缩延迟的56%。此外,如果我们使用NetZIP-算法来进一步降低压缩率,SNIC的Arm CPU必须从其内存中读取梯度和激活值,计算它们的增量值,对它们进行字节或比特分组,然后通过其内部PCIe互连将它们发送到加速器。由于SNIC的Arm CPU核心比服务器级CPU功能更弱、数量更少【索引17,HAL: Hardwareassisted Load Balancing for Energy-efficient SNIC-Host Cooperative Computing,2024,ISCA】,这只会进一步增加压缩延迟。我们在NVIDIA BlueField-2 SNIC上实现了NetZIP,包括通过SNIC的Arm CPU进行数据传输和压缩/解压过程。这个实现与无压缩相比,将节点间通信时间增加了11倍。
NVIDIA BlueField-3 SNIC仅支持解压加速器的原因。NVIDIA BlueField-2 SNIC实现了一个Deflate加速器。由于Deflate是一种重量级、资源密集型的无损压缩算法,在硬件资源和功耗限制下,NVIDIA BlueField-2 SNIC中的Deflate加速器只能提供50 Gbps的压缩吞吐量,这仅是SNIC 100 Gbps线路速率的一半。这种相对于线路速率较低的压缩吞吐量使得Deflate加速器对于需要高网络带宽的应用不太有用。此外,即使在50 Gbps下,Deflate加速器也消耗大量的硬件资源和功耗(§4)。考虑到NVIDIA BlueField-3 SNIC的网络链路带宽翻了一番,它可能没有能力在硬件资源和功耗限制下集成一个吞吐量更高的Deflate压缩器。
设计空间探索。为了解决这些限制,我们提出了NetZIP-加速器,一个集成到NIC ASIC中的经济高效的“线上传输中处理”压缩加速器。首先,作为一个“线上传输中处理”加速器,它消除了因在SNIC的CPU、内存和加速器之间传输梯度和激活值而产生的长延迟。其次,它实现了NetZIP-算法和一个标准压缩算法,并将它们集成在一个NIC ASIC内。在为NetZIP-加速器选择标准压缩算法时,我们考虑了压缩率、吞吐量、延迟、硬件资源和功耗之间的权衡。
硬件实现对比。尽管像Zstd和Deflate这样的重量级算法通常提供更高的压缩率,但当在CPU或GPU上运行时,它们的吞吐量更低,延迟更高,相比之下,像LZ4和Snappy这样的轻量级算法则表现更好。也就是说,在硬件中实现的Zstd和Deflate也将产生更长的延迟、更低的吞吐量和/或更高的功耗,同时需要比LZ4和Snappy更多的硬件资源。因此,如果NetZIP-算法即使对于LZ4和Snappy也证明有效,NetZIP-加速器就可以实现一个轻量级压缩算法,并比重量级、资源密集型的压缩算法实例化更多的压缩器/解压器,从而在硬件资源和功耗限制下提供更高的吞吐量和更低的延迟。
硬件实现性能比较。表4比较了基于FPGA的LZ4、Snappy、Zstd和Deflate的吞吐量、延迟和功耗,其中延迟定义为对于给定的输入流输出第一个压缩或解压字节的时间。为了综合这些压缩加速器,我们使用了Vitis数据压缩库【索引5,Vitis Data Compression Library,accessed in 2025】并将它们映射到AMD Alveo U280数据中心加速卡【索引4,Alveo U280 Data Center Accelerator Card,accessed in 2025】上。由于LZ4和Snappy提供了几乎相同的吞吐量、延迟和FPGA资源消耗,我们将在本节中以LZ4与Zstd和Deflate进行比较。
表 4: 在与LZ4提供50 Gbps吞吐量相同的FPGA资源限制下,LZ4、Snappy、Zstd和Deflate加速器的吞吐量、延迟和功耗。
硬件资源与性能分析。为了达到50 Gbps的吞吐量,LZ4消耗了17%的FPGA资源。在相同或相似的资源限制下,LZ4的吞吐量比Zstd和Deflate分别高1.7倍和2.1倍,后两者仅能达到30 Gbps和24 Gbps的吞吐量。这是因为LZ4用于压缩器和解压器实例的资源仅分别占Zstd所用资源的13%和20%。LZ4的压缩和解压延迟也分别比Zstd低51%和86%。
图 8: 集成在标准NIC中的NetZIP-加速器。
设计。NetZIP-加速器作为一个“线上传输中处理”的加速器,位于NIC ASIC中的DMA引擎+缓冲区和协议引擎之间,如图8所示,我们对每个4 kB的数据包进行梯度和激活值的压缩。由于分布式训练框架使用RDMA协议并以突发方式发送大量梯度和激活值,它们对所有数据包都使用巨型数据包大小(即4 kB),以最小化数据包头开销。为了选择性地仅压缩和解压包含梯度或激活值的数据包有效载荷,我们调整了RDMA verbs,允许分布式训练框架在每个数据包中附加一个自定义头部——一个1位的压缩标志和一个15位的层ID。这可以支持32,768(=$2^{15}$)个层ID,而Llama-3 405B有128个层。
NetZIP-加速器的组成与工作流程。NetZIP-加速器包括压缩和解压路径,如图8所示。压缩路径由1个增量编码器(Delta Encoder)、2个分组暂存缓冲区(Grouping Staging Buffer)、3个分组器(Grouper)、4个压缩器(Compressor)和5个发送(解)复用器(Tx (DE)MUX)组成。由于并非每个数据包都是存储中间值的RDMA数据包,为了防止不必要地压缩这类数据包,Tx DEMUX会根据自定义头部中的压缩标志,将数据包引导至协议引擎或增量编码器。增量编码器由三个单元组成:(1)一个基准值计算单元,(2)一个增量值计算单元,以及(3)基准值寄存器。为了使用给定层中梯度或激活值的最小值为基准值,基准值计算单元首先在增量值寄存器中存储bfloat16格式的最大可能值,将其视为当前最小值。当它通过Tx DEMUX从DMA引擎+缓冲区接收数据包时,它会将当前最小值与这些与同一层ID关联的数据包有效载荷中的每个bfloat16值进行比较。如果比较的值小于寄存器中的当前最小值,它将被存储在寄存器中作为新的当前最小值。增量值计算单元获取数据包有效载荷中每个中间值(即梯度或激活值)的增量值。在计算完所有中间值的增量值后,它将这些增量值连同基准值一起传递给分组暂存缓冲区。
压缩路径细节。为了提供高效的压缩和最小的延迟,我们为分组暂存缓冲区采用了双缓冲架构。当一个缓冲区正在接收一个4 kB的数据包时,另一个存储前一个数据包的缓冲区被分组器用来对增量值进行字节或比特粒度的分组,然后将它们发送给压缩器。压缩器压缩分组后的值,并将基准值附加到有效载荷的末尾。如果压缩后的有效载荷大小超过4 kB——这在我们的评估中从未发生过——它会简单地丢弃压缩后的数据包,并让Tx MUX在移除自定义头部后将未压缩的数据包发送到协议引擎。
解压路径细节。解压路径是压缩路径的逆过程,由接收解复用器(Rx DEMUX)、解压器(Decompressor)、解组暂存缓冲区(Ungrouping Staging Buffer)、解组器(Ungrouper)、增量解码器(Delta Decoder)和接收复用器(Rx MUX)组成。当从协议引擎接收到一个数据包时,Rx DEMUX检查其自定义头部以确定该数据包是否被压缩。如果发现数据包未被压缩,它会通过Rx MUX将数据包直接发送到DMA引擎+缓冲区。否则,它会将数据包发送到解压器。解压有效载荷后,解压器将其写入解组暂存缓冲区中的一个缓冲区。解组器将比特或字节的顺序重新排列回原始顺序,并将有效载荷连同基准值一起发送到增量解码器。随后,增量解码器将基准值加到每个增量值上,完成有效载荷中所有中间值的恢复。最后,它将数据包发送到DMA引擎+缓冲区。
ASIC实现分析。我们的FPGA实现基于Xilinx U280(16nm),目前以300 MHz的时钟频率运行。通过采用现代NIC ASIC中常见的7nm工艺进行ASIC实现【索引21,A summary of High Speed Ethernet ASICs,accessed in 2025】,我们估计设计可以优化到运行频率高达1 GHz。使用Kuon和Rose的方法来预测FPGA特性到ASIC【索引6,GlobalFoundries Details 7 nm Plans: Three Generations, 700 mm2, HVM in 2018,accessed in 2025】、【索引24,Quantifying and Exploring the Gap Between FPGAs and ASICs,2010,Springer Science & Business Media】,估计的面积约为2mm²(不到CX-7 NIC物理面积的0.1%【索引43,NVIDIA ConnectX-7 Adapter Cards User Manual,accessed in 2025】),功耗约为0.5 W(约为CX-7 NIC热限制25 W的2%【索引43,NVIDIA ConnectX-7 Adapter Cards User Manual,accessed in 2025】)。估计的流水线深度为170。NetZIP与不断发展的NIC架构高度兼容。
与其他分布式训练框架的兼容性。NetZIP可以透明地应用于当今领先的分布式训练框架,如PyTorch【索引48,PyTorch,accessed in 2025】、DeepSpeed【索引52,DeepSpeed: System Optimizations Enable Training Deep Learning Models with Over 100 Billion Parameters,2020,SIGKDD】和TensorFlow【索引1,TensorFlow: Large-Scale Machine Learning on Heterogeneous Distributed Systems,2016,arXiv】,因为NetZIP只改变网络通信层;核心训练代码无需任何更改。具体来说,所有这些框架都通过像NCCL【索引42,NVIDIA Collective Communications Library (NCCL),accessed in 2025】这样的后端来处理它们的通信调用,而NCCL又会发出RDMA verbs。这些框架可以通过扩展NCCL调用(如ncclAllReduce())增加一个额外字段来暴露当前层ID。在NCCL内部,我们只需要在调用ibv_post_send()(提交工作请求到NIC的RDMA verb)之前,为每个数据包附加一个自定义头部。这个插入操作发生在NCCL的ncclIbIsend()函数中,因此该点以上的所有内容,包括核心训练逻辑,都保持不变。
模型与数据集:
* 模型: Llama-3 70B、GPT-3 175B (使用BLOOM 176B【索引69,BLOOM: A 176B-Parameter Open-Access Multilingual Language Model,2023,arXiv】收集数据进行近似) 和 Llama-3 405B。
* 数据集: 使用Alpaca数据集【索引58,Alpaca: A Strong, Replicable Instruction-Following Model,2023,Stanford Center for Research on Foundation Models】对模型进行微调,以收集梯度和激活值。
硬件配置:
* 评估服务器 (表5):
* CPU: 2 x 第四代Intel Xeon可扩展处理器 (Sapphire Rapids),共80核。
* 内存: 1 TB DDR5-4800。
* GPU: 1 x NVIDIA H100 (80GB HBM3)。
* 网络: 1 x NVIDIA ConnectX-6 NIC, 1 x NVIDIA BlueField-2 SmartNIC。
NetZIP-加速器原型 (图9):
大规模集群仿真:
表 5: 评估系统配置。
图 9: 将基于FPGA的NetZIP-加速器以“线上传输中处理”方式连接到标准NIC的设置。
软件配置:
压缩率评估 (图10, 图11):
NetZIP-算法效果显著:
通用性: 该方法在Mixtral 8×7B模型上,对bfloat16和bfloat8精度的数据同样有效,证明了其通用性。
图 10: (a) 梯度和 (b) 激活值的压缩率比较,包括:原始(Orig)、字节分组(Byte)、比特分组(Bit)、真实增量压缩(D)、最小值增量压缩(Dmin)及其组合。
图 11: Mixtral 8×7B模型使用bfloat16和bfloat8时的梯度和激活值压缩率;字节分组不适用于bfloat8。
通信时间评估 (图12, 图13):
* 节点间通信 (图12):
* 延迟: NetZIP-加速器的压缩/解压延迟比最快的商用平台方案(GPU-based Snappy)低了几个数量级。NetZIP-LZ4的延迟比NetZIP-Zstd低约2-7倍。
* 总时间: NetZIP-LZ4将节点间通信时间平均减少了43%-44%,而GPU-based Snappy反而使通信时间增加了10%-18%。
集群级通信 (图13):
总训练时间:
图 12: NetZIP-加速器的压缩/解压延迟(相对于GPU-based Snappy归一化)和通信时间(相对于无压缩归一化)。
图 13: (a) 在512节点集群中,NetZIP相对于无压缩的平均每次迭代节点间通信时间(柱顶浅色部分表示网络拥塞的贡献);(b) 通信时间减少在比特分组和增量压缩之间的分解。
与有损压缩的比较 (图14):
图 14: Llama-3 70B在三种设置下训练损失随迭代次数的变化比较。
本文提出了NetZIP,一个算法/硬件协同设计的方案,以解决分布式大模型训练中的通信瓶颈问题。首先,我们提出了NetZIP-算法,它通过比特/字节分组和增量值压缩等技术,显著增强了标准无损压缩算法对梯度和激活值的压缩能力。其次,我们设计了NetZIP-加速器,它将NetZIP-算法与一个轻量级压缩算法(如LZ4或Snappy)的硬件加速器以“线上传输中处理”的方式集成到NIC ASIC中。这种设计最大化了压缩/解压的吞吐量和延迟,同时满足了NIC ASIC的硬件资源限制。我们的评估表明,在众多标准无损压缩算法中,LZ4或Snappy是实现NetZIP-加速器的最佳选择。最终,NetZIP将总训练时间减少了35%,且完全不影响模型精度。
本节说明了如何运行实验以复现论文的关键结果。步骤分为三部分:(1)在微调期间收集中间值(梯度/激活值)以支持其他实验;(2)评估NetZIP-算法(字节/比特分组和增量值压缩)与LZ4/Snappy/Zstd/Deflate结合的压缩率,对应图10;(3)运行SimAI大规模实验以量化端到端影响,对应图3、4和13。
A.3.1 如何访问。克隆仓库并进入项目目录:
git clone https://github.com/ece-fast-lab/MICRO-2025-NetZIP.git
cd MICRO-2025-NetZIP
A.3.2 硬件依赖。评估平台包括一台配备双Intel Xeon Sapphire Rapids处理器、1 TB DDR5内存、NVIDIA H100 GPU和Mellanox ConnectX-6 (CX-6) NIC的CPU服务器。可使用NVIDIA DGX平台(8×H100 GPU)来加速实验。
A.3.3 软件依赖。使用基于NVIDIA PyTorch Docker镜像的容器化环境。软件栈包括用于模型和数据管理的Hugging Face Transformers和Datasets库,以及用于数值运算的NumPy。为评估压缩方案,使用了LZ4、Snappy、Zstd和Deflate的库。SimAI用于大规模模拟多GPU环境。
要运行实验,请按照以下步骤使用Docker和Python包设置所需环境。
Docker。我们使用NVIDIA官方的PyTorch容器。挂载当前目录并启动容器。
环境。进入容器后,安装必要的Python包,包括用于模型训练、数据集处理、数值运算和压缩的库:
pip install torch transformers datasets huggingface_hub numpy lz4 snappy zstandard tqdm
huggingface-cli login
本部分我们使用较小的模型进行功能性测试。
1) 数据收集。我们从预训练的语言模型中收集中间训练数据(例如,梯度或激活值):
cd data_collection
./run1.sh
model_name指定要加载的Hugging Face模型标识符。num_steps设置收集数据的迭代次数。
2) 使用NetZIP-算法进行压缩分析。用保存的中间数据评估压缩率:
dir_step0和dir_step1标志指定了来自两个连续训练迭代的数据的输入目录。output_dir设置结果的目标目录,num_workers定义了用于加速计算的并行线程数。
3) 大规模模拟。运行SimAI以模拟不同带宽和模型大小的大规模训练。
第一个参数设置模型大小。第二个参数指定网络拓扑。第三个参数是全局批量大小。其余值为Gbps单位的带宽设置。
例如,由NetZIP-算法增强的LZ4,对梯度和激活值的压缩效果比标准LZ4分别高出60和70个百分点。使用NetZIP的端到端训练比无压缩训练的训练时间低35%。