Leveraging Chiplet-Locality for Efficient Memory Mapping in Multi-Chip Module GPUs

  • 作者/机构: JUNHYEOK PARK, Electronics and Telecommunications Research Institute; SUNGBIN JANG, OSANG KWON, YONGHO LEE, SEOKIN HONG, Sungkyunkwan University

A1 主要贡献

本文研究了多芯片模块（MCM）GPU中内存映射的页大小对性能的影响，并提出了一种名为CLAP（Chiplet-Locality Aware Page Placement）的高效内存映射机制。

核心问题与研究目标
- 问题: MCM设计虽然提高了GPU的可扩展性，但也引入了内存非均匀性问题，即访问远程chiplet上的数据会产生额外的延迟和能耗。页大小是影响内存映射的关键因素，但其在MCM GPU中的影响尚未被充分研究。
- 权衡: 大页面（如2MB）能通过单个TLB条目覆盖更大的内存区域，从而降低地址转换开销；但其粗粒度的数据放置策略可能导致数据被错误地分配到非本地的chiplet，增加远程访问。小页面（如4KB、64KB）允许更细粒度的数据放置，能更好地将数据映射到最常访问它的chiplet，但地址转换开销更高。
- 观察: 不同应用程序对页大小的敏感度不同。如图1所示，STE、3DC等应用在使用大页面时性能下降，因为远程访问比例增加；而SSSP、GPT3等应用则从大页面中受益。即使采用先进的远程缓存方案（如图2），选择合适的页大小仍然至关重要，因为不当的页大小会导致过多的远程访问，超出缓存系统的处理能力。
- 目标: 设计一种内存映射机制，能够为每个应用程序（甚至每个数据结构）确定合适的页面大小，从而在享受大页面带来的低地址转换开销的同时，不牺牲chiplet级别的内存局部性。


图 1: 不同页面大小下的性能（柱状图，以4KB页面为基准进行归一化）和内存指令的远程访问率（折线）。结果表明，每个应用程序对页面大小的敏感度各不相同。


图 2: 不同页面大小和远程缓存方案下的性能（柱状图）和内存指令的远程访问率（折线）。‘No_RC’表示“无远程缓存”。结果表明，即使有远程缓存，选择合适的页面大小仍然至关重要。

主要创新点
本文的主要贡献如下：
- 分析了页大小在MCM GPU中的影响：识别了地址转换效率和内存访问局部性之间的关键权衡。研究表明，GPU工作负载甚至在同一工作负载的不同数据结构之间也表现出多样的内存访问特性，导致不同的页大小偏好。因此，选择适合应用的页大小对实现MCM设计的高性能至关重要。
- 引入了“chiplet-locality”概念：这是GPU工作负载在MCM设计中表现出的一种独有特性，指GPU数据结构中存在一些页面组，这些组内的页面在虚拟地址上连续，且主要由同一个chiplet访问。这一特性为确定合适的数据放置粒度和策略提供了关键洞见。
- 提出了CLAP机制：这是一种在MCM GPU中实现高效页面映射的机制。CLAP能确保表现出chiplet-locality的页面组被连续地映射到合适的chiplet上，从而构建出类似大页面的区域。这种设计既能获得大页面的优势，又能保持chiplet内的内存局部性。

A3 背景知识

2.1 MCM GPU架构

基线MCM GPU架构。如图3所示，一个MCM GPU由四个GPU chiplet组成。每个chiplet包含多个流式多处理器（SMs）和内存分区。这些chiplet通过封装内互连技术相连，使得所有chiplet能够全局共享内存分区。在每个chiplet内部，交叉开关（Xbars）负责管理内存访问，将请求路由到本地或远程的内存分区。这种架构引入了NUMA效应，即远程内存访问会产生更高的延迟、消耗片外带宽并增加能耗。


图 3: 基准MCM GPU架构。

2.2 NUMA：仍然是个问题吗？

NUMA效应的持续影响。近年来，处理器供应商试图通过引入专用的I/O die进行内存集成和增加额外的缓存层来缓解MCM设计中的NUMA（非均匀内存访问）效应。同时，互连技术的进步也提升了片外内存带宽，同时降低了能耗和访问延迟。

物理限制导致NUMA无法消除。尽管有这些进步，由于基本的物理限制，NUMA效应无法被完全消除。例如，在由两个逻辑内存分区组组成的NVIDIA H100 GPU中，全局内存访问延迟大约在550到730个周期之间，具体取决于计算核心与内存分区之间的物理距离。随着MCM设计规模的扩大，这种延迟差异预计会加剧，因为chiplet间的通信成本增加会进一步恶化内存的非均匀性。这表明，解决NUMA效应对于维持新兴MCM架构的性能可扩展性至关重要。

2.3 与早期NUMA GPU的区别

MCM GPU与早期NUMA GPU的三个关键区别。虽然MCM GPU属于NUMA GPU的大类，但它们在三个关键方面有根本不同。

架构差异。首先，尽管早期的NUMA GPU（如多GPU系统）和MCM GPU都需要并受益于显式的工作负载（线程和数据）分配和一致性维护，但它们的架构方法不同。NUMA GPU由独立的设备组成，而MCM GPU则作为一个单一系统运行，对操作系统和程序员隐藏了硬件层面的复杂性。这种抽象减少了程序员驱动优化的机会，反而要求更强大的系统级支持。

地址空间差异。其次，与拥有独立物理地址空间和每个设备一个GPU页表的NUMA GPU不同，MCM GPU采用统一的地址空间和所有chiplet共享的单一GPU页表。因此，像页面复制【【65, Multi-GPU Memory Management with an Integrated Memory Model, 2021, MICRO】，【104, GRIT: Enhancing Multi-GPU Performance with Fine-Grained Dynamic Page Placement, 2024, HPCA】】这样的技术变得不适用，因为一个页表禁止将单个虚拟地址映射到多个物理位置。

性能权衡差异。最后，由于MCM GPU在封装层面而非板级或系统级集成chiplet，它们呈现出不同的性能权衡。NUMA GPU通常受益于页面迁移以减轻远程内存访问的高昂成本，而MCM设计可能因为封装内通信的开销相对较低而更倾向于远程访问【【13, MCM-GPU: Multi-Chip-Module GPUs for Continued Performance Scalability, 2017, ISCA】，【101, Design and Analysis of an APU for Exascale Computing, 2017, HPCA】】。这些根本差异凸显了为MCM设计特性量身定制架构策略的必要性。

2.4 虚拟内存系统设计

统一的虚拟内存系统。一个MCM GPU维护一个单一的GPU页表以支持统一的物理地址空间。不仅应用程序的数据结构，页表项（PTEs）本身也以页面粒度（通常是4KB）进行管理。因此，PTE页面可以根据PTE页面映射策略【【1, Mitosis: Transparently Self-Replicating Page-Tables for Large-Memory Machines, 2020, ASPLOS】，【81, Fast Local Page-Tables for Virtualized NUMA Servers with vMitosis, 2021, ASPLOS】，【87, Designing Virtual Memory System of MCM GPUs, 2022, MICRO】，【88, WASP: Workload-Aware Self-Replicating Page-Tables for NUMA Servers, 2024, ASPLOS】】分布在不同chiplet的内存分区中。

TLB层次结构与GMMU。在TLB层次结构中，L1 TLB位于每个SM内部，而L2 TLB要么在每个chiplet内的SMs之间共享（即chiplet-private），要么在GPU中所有SMs之间共享（即chiplet-shared）。我们的基线模型采用chiplet-private的L2 TLB配置，其中内存请求只访问其发起chiplet内的TLB。每个chiplet都配备了一个GPU内存管理单元（GMMU），它使用硬件页表遍历器（page table walker）来处理TLB未命中。页表遍历器会遍历多级内存页表【【93, Scheduling Page Table Walks for Irregular GPU Applications, 2018, ISCA】，【94, Neighborhood-Aware Address Translation for Irregular GPU Applications, 2018, MICRO】】；在每一级，页表访问可以是本地的或远程的，这取决于相应条目的物理位置【【1, Mitosis: Transparently Self-Replicating Page-Tables for Large-Memory Machines, 2020, ASPLOS】，【81, Fast Local Page-Tables for Virtualized NUMA Servers with vMitosis, 2021, ASPLOS】，【87, Designing Virtual Memory System of MCM GPUs, 2022, MICRO】，【88, WASP: Workload-Aware Self-Replicating Page-Tables for NUMA Servers, 2024, ASPLOS】】。

2.5 MCM GPU中的内存访问流程

内存访问步骤。图3概述了MCM设计中的内存访问流程。当SIMD单元发起一个内存访问请求时，它首先查询多级TLB（❶, ❷）以进行虚拟到物理地址的转换。如果TLB未命中，请求被转发到本地的GMMU（❸）。然后，多线程的页表遍历器通过执行页表遍历（page walk）（❹）来处理该请求。地址转换完成后，请求从物理地址检索数据。根据数据的位置，请求通过交叉互连（❺）被路由到本地内存分区或远程chiplet。

按需分页（Demand Paging）流程。当系统采用按需分页方案时（类似于统一虚拟内存，UVM【【3, In-Depth Analyses of Unified Virtual Memory System for GPU Accelerated Computing, 2021, SC】，【11, Unified memory management, 2024, AMD】，【91, UNIFIED MEMORY ON PASCAL AND VOLTA, 2017, NVIDIA GTC】】），一个数据页可能不在GPU内存中，从而导致页错误（page fault）。这个错误在页表遍历过程中被检测到。如果页表遍历失败，它会触发一个页错误，该错误被记录在GMMU的故障缓冲区（Fault Buffer）（○A）。GMMU随后通过一个共享的封装上I/O模块通知主机GPU驱动程序该故障。GPU驱动程序通过将出错的页面从CPU内存迁移到GPU内存并更新GPU页表（○B）来解决该故障。一旦故障解决，GPU将恢复地址转换过程。

2.6 NUMA感知的内存交错

传统交错策略的局限性。传统GPU通过在亚页级粒度（即256B）上跨多个内存通道进行数据交错来增强通道级并行性。然而，在MCM设计中，这种方式限制了GPU驱动程序以NUMA感知的方式优化数据放置的能力。为缓解这些问题，交错设计必须避免干扰页面级别的局部性，从而允许GPU驱动程序将整个页面分配给特定的chiplet【111, NUBA: Non-Uniform Bandwidth GPUs, 2023, ASPLOS】。

NUMA感知交错策略。图4展示了一个用于4-chiplet MCM GPU的34位物理地址的内存交错策略，该GPU拥有64个内存通道（6个通道位）。在此策略中，通道位的两个最高有效位（MSBs）被放置在页偏移（本例中为2MB）之外，并用作chiplet标识符。这种方法使GPU驱动程序能够在保持每个chiplet内通道级并行性的同时，强制执行预期的数据放置。值得注意的是，类似策略已在近期的基于chiplet的设计中为NUMA优化的工作负载提供支持【【5, 4th Gen AMD EPYC Processor Architecture White Paper, 2023, AMD】，【7, AMD EPYC 9004 Series Processors BIOS and Workload Tuning Guide, 2023, AMD】，【8, 5th Gen AMD EPYC Processor Architecture White Paper, 2024, AMD】，【9, AMD EPYC 9005 Series Processors BIOS and Workload Tuning Guide, 2024, AMD】】。

性能影响。我们观察到，即使没有任何NUMA优化，这种交错策略也不会导致任何明显的性能下降，与单片GPU中使用的朴素交错策略相比仅有0.6%的差异。此外，通过一个简单的NUMA感知优化（本文用作基线）【13, MCM-GPU: Multi-Chip-Module GPUs for Continued Performance Scalability, 2017, ISCA】，其性能比朴素配置高出42.0%。

图 4: MCM设计中的内存交错策略。通道位的两个最高有效位在页偏移之外，并用作chiplet标识符。

2.7 线程块和数据布局

布局策略的重要性。线程块（TB）和数据布局在MCM设计中尤为重要，因为不准确的放置会加剧内存非均匀性【【13, MCM-GPU: Multi-Chip-Module GPUs for Continued Performance Scalability, 2017, ISCA】，【47, Locality-Centric Data and Threadblock Management for Massive GPUs, 2020, MICRO】】。我们从先前的工作中归类出两种代表性方法。

两种代表性方法。第一种方法，称为首次接触（First-Touch-based, FT），将连续的线程块调度在同一个chiplet内，并将数据放置在首次接触它的线程所在的chiplet中【13, MCM-GPU: Multi-Chip-Module GPUs for Continued Performance Scalability, 2017, ISCA】。这个策略确保了相邻线程块能利用空间局部性，同时数据被放置在访问它的线程附近。第二种方法，称为静态分析（Static-Analysis-based, SA），利用编译器辅助的静态分析来确定线程块与数据的局部性【47, Locality-Centric Data and Threadblock Management for Massive GPUs, 2020, MICRO】。该方法分析程序代码以识别局部性模式，并根据预定义的放置策略来安排线程块和数据。

A3 动机性研究

3.1 系统配置

线程块和数据布局。我们采用广泛使用的FT策略【【13, MCM-GPU: Multi-Chip-Module GPUs for Continued Performance Scalability, 2017, ISCA】，【27, CPElide: Efficient Multi-Chiplet GPU Coherence, 2024, MICRO】，【64, Beyond the Socket: NUMA-Aware GPUs, 2017, MICRO】，【90, HMG: Extending Cache Coherence Protocols Across Modern Hierarchical Multi-GPU Systems, 2020, HPCA】，【106, Combining HW/SW Mechanisms to Improve NUMA Performance of Multi-GPU Systems, 2018, MICRO】，【108, Characterizing Multi-Chip GPU Data Sharing, 2023, TACO】，【109, SAC: Sharing-Aware Caching in Multi-Chip GPUs, 2023, ISCA】】作为我们的基线，因为它能有效利用常规和不规则应用的数据局部性。尽管FT可能因按需分页而产生分页开销【113, Towards High Performance Paged Memory for GPUs, 2016, HPCA】，但这种成本本质上是一次性的数据传输，这在其他方案的初始数据放置中也会产生。

虚拟内存支持。我们的基线系统采用带物理帧预留的按需分页，这与当前的GPU统一虚拟内存系统【【3, In-Depth Analyses of Unified Virtual Memory System for GPU Accelerated Computing, 2021, SC】，【74, NVIDIA Linux Open GPU Kernel Moudle Source, 2023, NVIDIA】，【85, A Case for Speculative Address Translation with Rapid Validation for GPUs, 2024, MICRO】，【91, UNIFIED MEMORY ON PASCAL AND VOLTA, 2017, NVIDIA GTC】】和先前的分页策略【【68, Practical, transparent operating system support for superpages, 2003, SIGOPS Oper. Syst. Rev.】，【107, Transparent Superpages for FreeBSD on ARM, 2014, Bodek】】相似。图5展示了一个带物理帧预留的分页示例。对于一个大页面（例如2MB），GPU驱动程序首先预留一个相应大小的空闲物理帧。然后，子页面（例如64KB）按需映射到该预留帧中。一旦所有子页面都被填充，GPU驱动程序会将该区域提升为一个大页面。此策略为不同页面大小提供了一致的按需分页粒度，并保持了统一的硬件资源使用（例如CPU-GPU互连带宽），与页面大小无关。我们的基线支持三种页面大小——4KB、64KB和2MB——这些在现代GPU系统中很常用【【72, Pascal MMU Format Changes, 2020, NVIDIA】，【74, NVIDIA Linux Open GPU Kernel Moudle Source, 2023, NVIDIA】】。我们排除了更大的页面大小（例如512MB和1GB），因为它们被认为粒度过粗。


图 5: 带物理帧预留的分页。

3.2 方法论

模拟器与配置。我们扩展了GPGPU-Sim v4.2【48, Accel-Sim: An Extensible Simulation Framework for Validated GPU Modeling, 2020, ISCA】来模拟一个4-chiplet的MCM GPU。我们的模型基于先前的研究和商业GPU的规格【【13, MCM-GPU: Multi-Chip-Module GPUs for Continued Performance Scalability, 2017, ISCA】，【70, NVIDIA TESLA V100 GPU ARCHITECTURE, 2017, NVIDIA】】，详细参数见表1。为了支持多种页面大小，我们为每种页面大小分配一个独立的TLB。PTE页面被分布在各个chiplet上，如先前工作所建议【87, Designing Virtual Memory System of MCM GPUs, 2022, MICRO】，以减少页表遍历期间的远程访问。

评估工作负载。我们评估了来自不同基准套件【【22, A Quantitative Study of Irregular Programs on GPUs, 2012, IISWC】，【25, Rodinia: A Benchmark Suite for Heterogeneous Computing, 2009, IISWC】，【35, Auto-tuning a High-Level Language Targeted to GPU Codes, 2012, InPar】，【67, GraphBIG: Understanding Graph Computing in the Context of Industrial Solutions, 2015, SC】，【75, CUDA C/C++ SDK Code Samples, 2024, NVIDIA】，【96, Parboil: A Revised Benchmark Suite for Scientific and Commercial Throughput Computing, 2012, Center for Reliable and High-Performance Computing】】的15个工作负载，选择了那些具有足够并行性以利用4-chiplet GPU系统的工作负载。我们还包括了从机器学习（ML）流水线中提取的几个矩阵运算，使用了来自各种ML模型【【21, Language Models are Few-Shot Learners, 2020, Advances in Neural Information Processing Systems】，【29, An Image is Worth 16x16 Words: Transformers for Image Recognition at Scale, 2020, arXiv preprint】，【38, Deep Residual Learning for Image Recognition, 2016, CVPR】】全连接层的矩阵维度。这些矩阵运算涵盖了一系列形状并表现出不同的内存访问特性。表2总结了工作负载的特性，包括总输入大小、主导内核中的线程块数量，以及在使用4KB、64KB和2MB页面进行内存映射时的L2缓存和L2 TLB的MPKI（每千个warp指令的未命中次数）。

表 1: 基线模拟配置。

表 2: 评估的工作负载。

3.3 MCM设计中页大小的影响

大页面对局部性的破坏。图7展示了一个场景，其中每个调度在不同chiplet上的线程块（TB0–TB3）访问一个专属的数据区域（P0–P3）。在小页面配置下（图7a），每个数据区域（P0–P3）被分割成独立的页面，允许GPU驱动程序将每个页面放置在请求数据的chiplet附近。然而，当使用大页面时，这种局部性会受到损害。如图7b所示，一个大页面P0*跨越了连续的区域P0–P3，假设大页面是小页面的四倍大。当chiplet-0中的TB0首次访问P0时，GPU驱动程序将整个大页面P0*放置在chiplet-0中。因此，TB1–TB3随后访问它们各自的区域（P1–P3）时，这些区域现在位于chiplet-0中，导致了远程访问。

页大小对性能的敏感度。我们分析了不同页大小对GPU工作负载的影响，包括标准大小（4KB、64KB、2MB）和介于64KB和2MB之间的假设大小。由于4KB到64KB之间的中间大小影响可以忽略不计，我们排除了它们。为了准确模拟假设页大小的性能影响，我们为每个大小添加了额外的TLB（L1为16个条目，L2为512个条目）。图6显示，工作负载对页大小表现出不同的敏感性。左侧的工作负载随着页大小接近2MB，远程访问率急剧增高，导致性能下降，尽管地址转换开销减少。相比之下，右侧的工作负载在不同页大小下远程访问率保持稳定，并且随着页大小的增加性能得到提升。这一发现强调了在MCM系统中自适应选择页大小的重要性。

确定合适页大小的挑战。然而，为不同的工作负载确定合适的页大小是具有挑战性的。首先，合适的页大小不一定与总输入大小或单个数据结构的大小相关，其中数据结构指的是GPU内存分配（例如，调用cudaMalloc()或cudaMallocManaged()）。例如，表现出不规则内存访问模式的PAF，即使其总输入接近2GB，最大数据结构为1.86GB，但在128KB页大小时性能最佳。相反，具有规则内存访问模式的BLK，尽管总输入仅为310MB，单个数据结构平均只有77MB，却偏好2MB的页大小。其次，像STE、LPS和SC这样的工作负载，在使用当前系统不支持的假设页大小时性能更高。最后，即使在单个工作负载中，由于特性不同，各个数据结构也可能偏好不同的页大小。图8显示了3DC和BFS中两个代表性数据结构在不同页大小下的远程访问率。在3DC中，数据结构对页大小变化的敏感度一致，趋势相似，而在BFS中，它们的敏感度则各不相同。


图 7: 不同页大小下的数据放置。使用大页面会增加内存非均匀性。


图 6: 不同页面大小下的性能（柱状图，以64KB页面为基准进行归一化）和内存指令的远程访问率（折线）。每个应用程序在使用其偏好的页面大小（通常是较大但远程访问率较低的尺寸）时性能得到提升，突显了对页面大小敏感度的可变性。


图 8: 不同页面大小下数据结构的内存指令远程访问率（折线）。即使在单个工作负载（如BFS）内，不同的数据结构也可能对页面大小表现出不同的敏感度。

3.4 GPU工作负载中的Chiplet-Locality

Chiplet-Locality的定义与示例。Chiplet-locality指的是特定的一组虚拟上连续的数据页主要由同一chiplet上的线程访问，并且这种分组粒度在同一个数据结构内保持一致的趋势。图9展示了GPU数据结构中chiplet-locality的一个例子。数据结构分布在不同的chiplet上，GPU驱动程序将每个数据页映射到请求访问它的chiplet。页表显示了虚拟页与物理帧之间的映射关系。值得注意的是，四个虚拟上连续的页（例如，VPN 0x100–0x103）被同一个chiplet访问并映射到该chiplet，并且这种模式在其他段（例如，VPN 0x104–0x107）中重复出现。

Chiplet-Locality的普遍性。我们观察到chiplet-locality是GPU工作负载的一个普遍特性。为了量化这一点，我们测量了不同GPU工作负载中每个数据结构的地址范围展现出chiplet-locality的比例。我们的分析涵盖了每个工作负载中几乎所有的数据结构，计算了每个结构的比例然后取平均值（分组粒度可能在不同结构间变化）。小于2MB的数据结构被排除，因为它们通常访问频率低，对整体性能影响可忽略。对于GEMM操作，三个矩阵之一（即C = A × B中的矩阵B）被内核中所有线程访问，呈现出全局访问模式。我们将这种情况归类为100%的chiplet-locality，因为从每个chiplet的角度看，该矩阵的整个地址空间被均匀访问。每个结构都使用小页面（即64KB）进行映射，以实现每个页面到其期望chiplet的细粒度放置。图10显示，GPU数据结构表现出高度的chiplet-locality，平均93.5%的内存范围展示了这种行为。这源于GPU工作负载的内在特性：一个内核内的线程执行相同的程序代码，导致相似的内存访问模式。特别是，相邻的线程（例如，位于一个chiplet内的线程块中的线程）倾向于访问空间上相邻的内存区域【【37, R2D2: Removing ReDunDancy Utilizing Linearity of Address Generation in GPUs, 2023, ISCA】，【51, Access Pattern-Aware Cache Management for Improving Data Utilization in GPU, 2017, ISCA】，【80, APRES: Improving Cache Efficiency by Exploiting Load Characteristics on GPUs, 2016, ISCA】，【103, Decoupled affine computation for SIMT GPUs, 2017, ISCA】】。


图 9: GPU数据结构中chiplet-locality的一个例子。四个连续的页面（VPNs 0x100–0x103）主要由chiplet-0访问，并且这种模式在其他区域（例如VPNs 0x104–0x107）中持续出现。


图 10: 表现出chiplet-locality的GPU数据结构地址范围的比例。数据结构通常高度表现出chiplet-locality。

3.5 回顾先前的NUMA解决方案

现有方案的局限性。在MCM-GPU中，许多研究已经解决了NUMA挑战【【13, MCM-GPU: Multi-Chip-Module GPUs for Continued Performance Scalability, 2017, ISCA】，【32, Barre Chord: Efficient Virtual Memory Translation for Multi-Chip-Module GPUs, 2024, ISCA】，【47, Locality-Centric Data and Threadblock Management for Massive GPUs, 2020, MICRO】，【87, Designing Virtual Memory System of MCM GPUs, 2022, MICRO】，【109, SAC: Sharing-Aware Caching in Multi-Chip GPUs, 2023, ISCA】，【111, NUBA: Non-Uniform Bandwidth GPUs, 2023, ASPLOS】】，但大多数主要关注内存的两个关键方面之一：物理数据访问（指地址转换后从缓存或GPU DRAM中获取数据）或地址转换本身。

物理数据访问优化。首先，各种线程块和数据布局策略【【13, MCM-GPU: Multi-Chip-Module GPUs for Continued Performance Scalability, 2017, ISCA】，【47, Locality-Centric Data and Threadblock Management for Massive GPUs, 2020, MICRO】，【49, CODA: Enabling Colocation of Computation and Data for Multiple GPU Systems, 2018, TACO】】通过增强chiplet内的数据局部性来改善物理数据访问。同样，远程数据缓存方案【【109, SAC: Sharing-Aware Caching in Multi-Chip GPUs, 2023, ISCA】，【111, NUBA: Non-Uniform Bandwidth GPUs, 2023, ASPLOS】】通过将远程映射的数据存储在本地资源中进一步增强了它。然而，这些方法没有充分考虑虚拟内存配置（如页大小和地址转换）如何影响其有效性。

地址转换优化。其他方案旨在通过在GPU页表和TLB条目放置以及页表遍历调度中应用NUMA感知优化来减少地址转换开销【【32, Barre Chord: Efficient Virtual Memory Translation for Multi-Chip-Module GPUs, 2024, ISCA】，【87, Designing Virtual Memory System of MCM GPUs, 2022, MICRO】】。虽然这些方法也考虑了物理数据访问，但它们主要扩展了现有的NUMA优化【47, Locality-Centric Data and Threadblock Management for Massive GPUs, 2020, MICRO】。因此，它们假设了有利的局部性，并且没有充分考虑在不同虚拟内存配置下转换和数据访问如何相互作用。

C-NUMA的局限性。在NUMA-CPU系统中，先前的工作试图通过页面迁移在运行时动态构建和拆分大页面来同时优化物理数据访问和地址转换——这种方法被称为C-NUMA【【28, Traffic Management: A Holistic Approach to Memory Placement on NUMA Systems, 2013, ASPLOS】，【34, Large Pages May Be Harmful on NUMA Systems, 2014, ATC】】。尽管有其优点，这种方法有四个主要限制：
1. 缺乏对数据结构的个性化考虑: C-NUMA没有考虑应用程序内单个数据结构的独特特性，因此可能表现出次优或错误的行为。
2. 对页面迁移的重度依赖: 它严重依赖页面迁移，这会产生巨大的开销。在chiplet内部或之间迁移页面需要昂贵的操作，如TLB shootdown和缓存刷新，这些操作会引入额外延迟，降低物理数据访问和地址转换的效率，从而降低整体GPU性能【【12, Don’t shoot down TLB shootdowns!, 2020, EuroSys】，【100, Observations and Opportunities in Architecting Shared Virtual Memory for Heterogeneous Systems, 2016, ISPASS】】。
3. 灵活性有限: 它只支持少数特定的页面大小（如4KB和2MB），因此无法适应那些受益于中间页面大小的工作负载。即使支持这种大小，它也缺乏确定给定应用程序合适页面大小的机制。
4. 被动策略不适用于GPU: 它采用被动策略，根据应用程序当前的执行状态调整其行为。这在CPU系统中是有效的——因为工作负载不可预测，进程可能在chiplet间重新调度，改变它们的内存访问模式——但这并不适合GPU。GPU调度器不会在核心间迁移线程，因为线程迁移需要重新启动线程，这会产生巨大开销。此外，如3.4节所讨论，GPU工作负载表现出由大规模并行性驱动的稳定且可预测的访问模式，使得C-NUMA的被动方法不适用于GPU。因此，必须重新评估更先进解决方案的设计方向。

我们的方法。在本文中，我们提出了一种新的解决方案，通过联合优化物理数据访问和虚拟内存管理来提高MCM GPU的性能。具体来说，我们的设计：(1) 通过低开销机制快速分析应用程序行为，(2) 适应每个数据结构的独特特性，(3) 最小化对昂贵的页面迁移的依赖，以及 (4) 提供超越系统支持的页面大小的内存映射灵活性。

表 3: 与先前的NUMA解决方案的比较。

A2 方法细节

4. Chiplet-Locality Aware Page Placement

CLAP的提出。我们提出了CLAP，它为每个数据结构提供合适的页大小——具体来说，是合适的页面连续性级别。

合适的页大小定义。合适的页大小的定义取决于系统。在这项工作中，我们为MCM GPU定义的合适页大小是能够最大化大页面优势（例如，减少地址转换开销）同时保持数据局部性的页大小。如第3.3节所示，每个工作负载性能最佳的配置都符合这一原则。为了实现这一点，CLAP利用了chiplet-locality，它反映了多少个虚拟上连续的页面可以被组合在一起并映射到同一个chiplet。这使得能够快速准确地选择映射大小。

CLAP的映射方式。图11展示了三种不同配置下的内存映射。左上角（①）显示了使用小页面的映射。细粒度的映射使得每个数据页可以被放置在其偏好的位置，表现出强烈的chiplet-locality（例如，VPN 0x100–0x103被映射到chiplet-0）。然而，相应的物理帧在chiplet内部是分散的。右上角（②）显示了使用大页面的映射。在这种情况下，为一个大页面预留了一个连续的物理帧，提高了地址转换效率。然而，这种粗粒度的分配迫使不相关的区域（例如，VPN 0x104–0x107）驻留在同一个chiplet（chiplet-0）中，增加了远程内存访问。相比之下，下半部分（③）展示了CLAP的方法。系统预留一个与chiplet-locality大小相匹配的连续物理帧。然后，CLAP将虚拟上连续的页面映射到这个预留的帧中，确保虚拟和物理连续性的一致。这种对齐有效地起到了大页面的作用，从而在不损害内存局部性的情况下减少了地址转换开销。


图 11: 无CLAP（上）和有CLAP（下）的内存映射示例。CLAP确保了chiplet-locality区域被映射到连续的物理帧。

CLAP概览。图12展示了CLAP的概览。GPU驱动程序中的内存管理器负责监督GPU应用程序的内存映射。在应用程序启动时，内存管理器开始映射数据结构（❶），但仅限于一部分页面。一旦映射的总大小达到预定义比例，该过程便暂停。这个初始阶段被称为部分内存映射（PMM）（❷）。在PMM期间，内存管理器使用小页面（例如64KB），并将每个页面映射到请求访问的chiplet。这个过程创建了一个样本映射，作为分析chiplet-locality的基础。同时，内存管理器将元数据（例如，分配ID）发送到GPU中的硬件组件——远程跟踪器（Remote Tracker, RT）（❸）。RT收集运行时指标，例如远程访问率，这些指标稍后用于评估chiplet-locality并确定合适的页面大小。一旦PMM完成，GPU驱动程序收集当前的映射信息和RT指标，以执行内存映射分析（MMA）（❹）。MMA评估数据结构的chiplet-locality并选择合适的页面大小。最后，内存管理器接收MMA的结果，并将其应用于数据结构的剩余（未映射）部分（❺）。


图 12: CLAP概览。

4.1 基于块的内存管理

VA块和PF块。内存管理器采用一种基于块的内存管理方法，将虚拟地址空间和物理地址空间划分为2MB的单元，分别称为VA（虚拟地址）块和PF（物理帧）块。

VA块作为页大小分配边界。每个VA块都作为页大小分配的边界。为了在虚拟地址空间内支持多种页大小，内存管理器为每个VA块分配一个特定的页大小，并强制该块内的所有映射都使用该大小。例如，为一个2MB的虚拟区域分配64KB的页大小，需要相应的VA块将其所有子区域都用64KB的页面进行映射。这种方法简化了对虚拟地址空间每个区域分配了哪种页大小的跟踪。

PF块作为物理帧大小分配边界。与VA块类似，每个PF块作为物理帧大小分配的边界。当内存管理器需要在特定chiplet中获取特定大小的物理帧时，该chiplet内的一个空闲PF块会被分割成多个该大小的帧。例如，对于64KB的页面，一个空闲的PF块会被分成32个64KB的帧，每个帧都被插入到相应的空闲列表中。空闲列表是按chiplet维护的，每个列表只包含其所在chiplet可用的帧。这个策略防止了在PF块内混合不同大小的帧，并将帧对齐到2MB的边界，从而简化了帧管理并减少了碎片。

4.2 部分内存映射 (PMM)

PMM的目标与过程。PMM的目标是为chiplet-locality分析创建一个样本映射。PMM将触发页错误的数据页放置到GPU chiplet中，直到映射的页面总数达到预定义的阈值。在此阶段，PMM使用小页面，允许每个数据页被放置在请求访问它的chiplet中，从而自然地形成chiplet-locality区域。

PMM的实现细节。在我们的实现中，PMM使用64KB的页面。这个选择是合理的：64KB的页面提供了与4KB页面相当的数据局部性，同时提供了更高的地址转换效率（如图6所示）。此外，64KB也与当前商用GPU支持的最小页面迁移粒度相匹配【74, NVIDIA Linux Open GPU Kernel Moudle Source, 2023, NVIDIA】。

PMM阈值设定。默认情况下，PMM阈值设置为20%，意味着每个数据结构的20%在PMM期间被映射，而剩余的80%则推迟到使用合适的页大小进行映射。我们将PMM阈值设置为20%，这是一个根据经验得出的保守选择，以确保在所有评估的工作负载中进行可靠的chiplet-locality分析。虽然在我们的实验中15%的阈值已足够，但为了增强鲁棒性，我们增加了5%的余量。我们发现性能对PMM阈值基本不敏感。将阈值增加到30%仅导致平均1.3%的性能下降，表明其影响是非线性的且较弱。驱动程序中的内存管理器会监控每个数据结构的已映射比例，并在达到阈值时触发内存映射分析阶段。

机会性大页面（OLP）。由于PMM使用64KB页面，即使在2MB映射更优的情况下，此阶段映射的数据也无法享受到大页面的好处。为了解决这个问题，CLAP采用了机会性大页面（OLP），它在PMM期间动态地形成2MB页面，同时保留了细粒度的64KB映射。

OLP的内存映射流程。图13展示了PMM中带有OLP的内存映射流程。当一个VA块中的第一个64KB页面被映射到一个chiplet（例如，chiplet-0）时，内存管理器会为整个VA块预留一个2MB的物理帧（○a）。如果来自同一VA块的后续64KB页面被同一个chiplet请求，它们将被放置到预留的帧中（○b）。一旦该块中所有的64KB页面都映射到那个chiplet，预留的帧就会被提升为一个2MB的大页面。

处理不一致请求。然而，如果稍后的一个64KB页面被一个不同的chiplet（例如，chiplet-1）请求，它将被映射到那个chiplet（○c）。内存管理器随后会释放2MB帧的预留，并将剩余未使用的64KB帧放回空闲列表，从而防止碎片化并使其能被其他VA块重用。

OLP的禁用机制。如果超过5%的VA块释放了它们的2MB预留，PMM会禁用OLP并默认使用64KB映射。这个阈值可以防止在机会性大页面被证明无效时进行不必要的预留和释放操作。5%的限制是根据经验得出的，以确保从释放的预留中回收的64KB物理帧（即未使用的64KB帧）在后续的PMM操作中得到充分利用。

PMM阶段的持续时间。PMM阶段的持续时间因数据结构的使用模式而异——有些从执行开始就被使用，而有些则在稍后被访问。我们将PMM阶段的终点定义为所有数据结构完成分析的时刻。平均而言，PMM阶段在执行的早期部分结束，通常在执行的前6.92%的内核指令内完成。


图 13: PMM期间使用OLP的内存映射流程。

4.3 远程跟踪器 (RT)

RT的功能。虽然PMM为chiplet-locality分析建立了一个样本映射，但RT的作用是验证该映射是否真正对分配的chiplet是本地的。这是因为，即使使用小页面映射，一个页面内的不同缓存行也可能被不同chiplet上的线程访问，缓存行可能在chiplet间共享，或者整个数据结构可能是全局共享的（例如GEMM中的矩阵B）。

RT的操作流程。图14展示了RT的操作。当一个GPU数据结构被分配时，内存管理器会分配一个分配ID。这个ID被存储在末级页表项（PTE）的一个未使用位中【【10, AMD64 Architecture Programmer’s Manual Volumes 1-5, 2024, AMD】，【32, Barre Chord: Efficient Virtual Memory Translation for Multi-Chip-Module GPUs, 2024, ISCA】，【40, Intel 64 and IA-32 Architectures Software Developer’s Maunal Volume 3 (3A,3B,3C, and 3D): System Programming Guide, 2025, Intel】，【57, IDYLL: Enhancing Page Translation in Multi-GPUs via Light Weight PTE Invalidations, 2023, MICRO】，【72, Pascal MMU Format Changes, 2020, NVIDIA】，【104, GRIT: Enhancing Multi-GPU Performance with Fine-Grained Dynamic Page Placement, 2024, HPCA】】，使得RT能够识别哪些页面属于哪个分配。该ID同时被转发到每个chiplet的GMMU内嵌的RT，并随后记录在其内部表中。

远程访问检测与统计。一旦在地址转换期间页表遍历完成，就会使用从PTE中提取的分配ID和远程访问状态（指定页表遍历请求是否针对远程映射的页面）执行RT表查找（○A）。由于PTE中的PFN（物理帧号）编码了页面映射到的chiplet ID（即图4中通道位的两个最高有效位），通过将页面映射的chiplet ID与内存请求来源的请求者chiplet ID进行比较，可以将请求分类为本地或远程（○B）。如果在RT表中找到匹配项，RT会增加相应的访问计数器，并在访问被识别为远程时额外更新远程计数器（○C）。

与驱动程序的交互。当PMM阶段完成且驱动程序为给定的分配（即一个数据结构）启动chiplet-locality分析时，每个RT会将记录的统计数据转发给GPU驱动程序并清除相应的表项（○D）。RT通过主机-GMMU接口与驱动程序通信，这个接口在GPU中已经被使用【【2, Demystifying GPU UVM Cost with Deep Runtime and Workload Analysis, 2021, IPDPS】，【3, In-Depth Analyses of Unified Virtual Memory System for GPU Accelerated Computing, 2021, SC】】。

设计考量。RT根据页表遍历来收集统计数据。由于RT位于GMMU内部——页表遍历发生的地方，所有需要的元数据（如分配ID）都已可用，因此避免了将这些信息通过数据包传播到其他组件的需要。RT专注于估算远程访问的比例，而不是数据结构的绝对计数或热度，因此基于页表遍历的计数足以满足此目的。我们观察到，估算的比例与实际远程访问比例非常接近，相似度达到95.3%。在我们的基线中，每个RT包含一个32个条目的表。这个配置是根据我们的评估得出的；虽然16个条目的表对于评估的工作负载已足够，但32个条目的表提供了额外的灵活性。当表满且需要插入新的分配ID时，RT会根据远程计数器替换最近最少更新的条目。这种替换策略与RT跟踪具有较高远程访问强度的数据结构的目标非常吻合。被驱逐时，相应条目（即分配ID）的远程比例被视为零。作为一个可选功能，被驱逐的条目可以记录到GPU驱动程序以减轻数据丢失，尽管在我们的基线中此选项是禁用的。

优点和开销。这种基于硬件的方法比先前的基于软件的分析技术【【28, Traffic Management: A Holistic Approach to Memory Placement on NUMA Systems, 2013, ASPLOS】，【34, Large Pages May Be Harmful on NUMA Systems, 2014, ATC】】有两个优势。首先，RT以每个数据结构为基础收集统计数据。其次，它消除了频繁的软件干预，这种干预会降低GPU性能【56, SnakeByte: A TLB Design with Adaptive and Recursive Page Merging in GPUs, 2023, HPCA】。一个RT需要288字节，每个条目由一个8位的分配ID字段（基线）和两个32位的计数器组成。我们通过实现其Verilog模型并使用28nm UMC标准单元库【99, 28 Nanometer, 2022, UMC】进行综合来估算单个RT的面积开销。最终面积为0.0124mm²，约占现代GPU芯片面积（~800mm²）【【71, NVIDIA A100, 2020, NVIDIA】，【73, NVIDIA H100 Tensor Core GPU Architecture, 2023, NVIDIA】】的0.0015%。RT查找仅需两个周期：一个用于索引，一个用于更新。RT不在内存访问的关键路径上，也不需要RT间的同步。

PTE位的使用。现代PTE包含13个保留位（11个未使用 + 2个忽略）【【10, AMD64 Architecture Programmer’s Manual Volumes 1-5, 2024, AMD】，【40, Intel 64 and IA-32 Architectures Software Developer’s Maunal Volume 3 (3A,3B,3C, and 3D): System Programming Guide, 2025, Intel】，【57, IDYLL: Enhancing Page Translation in Multi-GPUs via Light Weight PTE Invalidations, 2023, MICRO】，【104, GRIT: Enhancing Multi-GPU Performance with Fine-Grained Dynamic Page Placement, 2024, HPCA】】。为了证明这些位足以编码分配ID，我们分析了各种GPU应用程序——特别是涉及多个GPU内存分配的LLM推理【54, Efficient Memory Management for Large Language Model Serving with PagedAttention, 2023, SOSP】。通过改变输入、模型类型【【43, Mistral 7B, 2023, arXiv preprint】，【98, LLaMa: Open and Efficient foundation language models, 2023, arXiv preprint】，【110, OPT: Open Pre-Trained Transformer Language Models, 2022, arXiv preprint】】和扩展参数至300亿，我们观察到最多约300个GPU分配，这完全在保留位的容量范围内。这个数字反映了整个执行过程中的总分配数，因此并发活动分配的数量可能要小得多。


图 14: RT监控每个数据结构的远程访问，收集到的统计数据用于chiplet-locality分析和页面大小选择。

4.4 内存映射分析 (MMA)

MMA的目标。MMA根据PMM期间映射的页面推断数据结构的chiplet-locality，并确定合适的页面大小。

基于树的Chiplet-Locality分析。当一个数据结构的已映射数据页数量达到PMM阈值（20%）时，GPU驱动程序会调用MMA来评估chiplet-locality。MMA会检查该数据结构中所有64KB子页都已完全映射的VA块，并使用基于树的算法计算它们的chiplet-locality。

算法工作流程。图15用一个简化的512KB VA块示例说明了该算法的工作流程。树结构存储了页面到chiplet的映射信息，由主机显式维护，并在数据页映射到GPU时由内存管理器更新。树的叶节点代表连续的64KB VA区域，每个节点存储其对应区域被映射到的chiplet ID（❶）。在每个内部节点，其子节点的chiplet ID被聚合成计数（❷）。然后，每个内部节点被赋予一个局部性得分（$score(l)$），使用以下公式（1）计算，其中$C_n$表示系统中的chiplet总数，$l$表示内部节点的树层级（❸）。

局部性得分的计算与应用。每个节点的局部性得分量化了其后代叶节点中映射到同一chiplet的比例。得分为“1”表示所有对应的64KB页面都完全映射到了单个chiplet。MMA计算树的每个层级的平均局部性得分（$score_{avg}(l)$），这代表了在使用2的幂次页大小时，64KB页面被正确映射到其目标chiplet的比例（❹）。基于这些平均分，MMA通过选择树中平均分达到或超过预定义阈值（$thres$）的最高层级来确定数据结构的chiplet-locality，如公式（2）所示。

阈值设定与分析。默认情况下，该阈值设置为“1”，确保所选层级（即所选大小）下的所有64KB叶页面都完全映射到其相应的chiplet（❺）。这使得MMA能够识别出既能保持chiplet-locality又最大的页面大小。当分析多个VA块时，MMA会计算每个VA块的chiplet-locality，并选择最主导的级别作为该数据结构的chiplet-locality。

结合远程跟踪器（RT）。在某些情况下，由于线程间的共享特性（例如GEMM中的矩阵B），一个数据结构无论其映射策略如何，都可能固有地产生高远程访问。在这种情况下，增加页面大小成为一个更有效的解决方案，因为它即使在没有强局部性的情况下也能减少地址转换开销。为了支持这一点，MMA整合了来自RT的远程访问统计数据。

选择标准的调整。它聚合来自RTs的计数器统计数据，并计算数据结构的远程访问率（$ratio_{rt}$）。然后，MMA通过从阈值$thres$中减去远程访问率$ratio_{rt}$来完善公式（2）中的chiplet-locality选择标准。推导过程如下。

在公式（3）中，$k(thres - score_{avg}(l))$代表当系统使用层级$l$作为内存映射中的页面大小时的预期远程访问率，其中$k$是一个缩放参数。减去固有的远程访问率（$ratio_{rt}$）后，剩余的比率必须不大于CLAP的目标远程访问率（$ratio_{target}$）。重新整理公式（3）得到公式（4），它取代了公式（2）中的条件。在我们的实现中，我们设置$ratio_{target} = 0$和$k=1$。

$$score_{\text{avg}}(l) \geq thres - \frac{ratio_{\text{rt}} + ratio_{\text{target}}}{k}$$

调整效果。因为原始阈值“1”强制映射策略遵守chiplet-locality，修改选择条件放宽了这一约束，增加了选择更大页面大小的可能性。在图15中，如果数据结构的远程访问率为0.75，MMA会重新制定选择标准，使得有效阈值降低到0.25（即1.00 - 0.75），从而导致选择512KB级别作为合适的页面大小。

优点和开销。MMA在GPU驱动程序内部后台运行，允许GPU应用程序不中断或停顿地执行。具体来说，每当一个叶节点（即一个64KB页面）被映射到GPU内存时，每个2MB VA块的chiplet-locality树就会增量更新。我们使用CPU实现来估算这种更新开销，每次更新不到1µs（通常约为0.4µs）。这个延迟完全与页面映射过程（包括页表更新和GPU内存清零）重叠，后者根据在商用GPU上使用NVIDIA UVM的测量，大约需要3µs【【71, NVIDIA A100, 2020, NVIDIA】，【73, NVIDIA H100 Tensor Core GPU Architecture, 2023, NVIDIA】】。额外的操作，如在树结构上计算局部性得分和选择页面大小，开销不到1µs，也完全被隐藏。


图 15: 基于树的chiplet-locality分析算法在一个512KB VA区域上的示例。

4.5 应用所选页大小

应用流程。内存管理器接收由MMA选择的页大小，并用它来映射数据结构的剩余部分。所选页大小可以跨越很宽的范围，包括当前GPU本身不支持的中间大小（例如，128KB到1MB）。特别是，支持中间大小需要对页表和GMMU进行重大更改，降低了与现有GPU系统的兼容性。为了解决这个问题，CLAP通过组合常规的小页面（即64KB页面）来构建一个类似大页面的区域。

中间大小的实现。图16展示了一个使用256KB页大小的内存映射示例。一个VA块被分配了256KB的页大小，并逻辑上划分为多个256KB的VA区域（①）。当一个256KB VA区域中的第一个64KB页面需要映射到GPU内存时（例如，映射到chiplet-0），内存管理器从相应的空闲列表中取出一个空闲的256KB物理帧（②）。如果没有可用的空闲帧，管理器会通过将该chiplet中的一个空闲PF块分割成多个256KB的帧来填充列表（③）。取出的256KB帧被预留给该VA区域，该VA区域内的后续64KB子区域按需使用常规的64KB页面映射到预留的帧中（④）。这个策略构建了一个在虚拟和物理上都连续的256KB类页面块。

处理边缘情况。在极少数情况下，当PMM期间数据结构的任何VA块都未完全映射时，CLAP可能无法执行MMA。同样，小的分配（通常小于10MB）也可能缺乏足够的2MB VA块进行有意义的分析。在这种情况下，CLAP会回退到机会性大页面（OLP），如4.2节所述。OLP通过机会性地形成大页面同时保持细粒度映射，动态地利用chiplet-locality，从而有效地处理这些情况。在我们的评估中，在这些边缘情况下使用OLP与理想选择的页大小相比，性能差异仅为1.1%。


图 16: 使用256KB页大小的内存映射流程。

4.6 与TLB合并的协作

对架构支持的需求。由于CLAP支持多种页面大小（即连续页面的组），它需要相应的架构支持才能充分利用其优势。虽然一个2MB大小的连续页面组可以被提升为一个真正的2MB页面——这在当前GPU中已经支持【69, (Mis)Managed: A Novel TLB-Based Covert Channel on GPUs, 2021, ASIACCS】，但中间大小的组（例如256KB）则不能。即使支持这种中间大小，它们也需要额外的TLB结构，会产生存储开销并在不使用时造成资源浪费。为了解决这个问题，CLAP采用了一种CLAP特定的TLB合并机制。

TLB合并流程。图17展示了TLB合并如何与CLAP协作，以一个256KB大小的组为例。当为VPN 0x103发出页表遍历请求时，页表遍历器检索相应的PTE并将它们加载到L2缓存中。由于GPU使用128B的缓存行（由四个32B的扇区组成），十六个8字节的PTE被获取并打包到一个缓存行中（○A）。然后，TLB控制器中的合并逻辑检查获取的PTE以检测连续映射。一旦确认连续性，这些PTE就被合并并作为一个单一的合并条目插入到TLB中（○B）。该条目使用一组有效位记录了覆盖的虚拟范围，一个合并的TLB条目可以表示多达十六个连续的页面（最多1MB）。如果在合并条目中发生TLB命中，TLB控制器通过位拼接和加法重构相应的VPN和PPN（○C）。这种方法利用了有意映射的连续页面，有效地扩展了单个TLB条目的覆盖范围。对于2MB大小的组，内存管理器会将其提升为2MB页面【69, (Mis)Managed: A Novel TLB-Based Covert Channel on GPUs, 2021, ASIACCS】，并将其插入到用于2MB页面的TLB中。

优点和开销。由于CLAP采用基于预留的映射，预留区域内的页面保持相同的虚拟到物理偏移。这使得硬件即使是部分连续的PTE也能合并成一个条目。为了估算支持CLAP合并机制的硬件成本，我们评估了两个组件。首先，我们实现了额外TLB控制器逻辑（即合并、拼接和加法单元）的Verilog模型，并使用28nm UMC标准单元库【99, 28 Nanometer, 2022, UMC】进行综合。综合后的设计占用了0.0024mm²，占GPU芯片面积的0.0003%【【71, NVIDIA A100, 2020, NVIDIA】，【73, NVIDIA H100 Tensor Core GPU Architecture, 2023, NVIDIA】】。其次，我们使用CACTI【19, CACTI 7: New Tools for Interconnect Exploration in Innovative Off-Chip Memories, 2017, TACO】估算了扩展的L2 TLB内存单元面积，发现修改使面积增加了8.98%。相比之下，将L2 TLB大小增加同等的面积开销仅能带来2.31%的性能提升，这表明CLAP实现了更高的效率。


图 17: TLB合并流程。CLAP利用合并来利用有意映射的连续页面，从而实现更大有效页面的优势。

4.7 讨论

内存碎片。CLAP通过以下设计选择来减轻外部和内部分区。由于CLAP为每个数据结构确定了合适的页大小，内存管理器仅使用所选大小的物理帧来映射该结构。因此，管理器可以为每个数据结构维护一个专用的空闲列表。当一个空闲的PF块被分割成多个空闲帧时，所有帧都被插入到同一个列表中，确保该PF块被单个数据结构使用。因此，当该结构被释放时，管理器可以回收整个PF块并将其重新分配给其他数据结构，而不会产生碎片。当预留的物理帧未被完全利用时，可能会发生内部分区。然而，CLAP仅为至少部分映射到GPU内存的虚拟地址区域预留物理帧。这样的区域在执行期间被完全利用的可能性要大得多。我们评估了不同分页方案下消耗的PF块总数，发现CLAP相对于64KB和2MB静态分页方案，平均内存使用量仅增加了0.57%和1.27%。

可扩展性。尽管我们的基线实现采用64KB作为基础页大小，但其核心机制——包括基于块的映射、物理帧预留和基于树的chiplet-locality分析——对于更小或更大的基础页大小都具有可扩展性。例如，CLAP可以适应支持4KB基础页面，从而实现4KB到64KB之间更细粒度的页大小。反之，通过将基础页大小增加到2MB，并将VA和PF块大小调整到1GB，CLAP可以支持高达1GB的大页面，从而最大化大页面映射的好处。

数据初始化。对于一个数据结构，其内存访问模式在初始化和计算阶段可能不同。然而，GPU工作负载通常在CPU上初始化输入数据。输出数据结构本身是计算阶段的一部分，因此自然地遵循内核的内存访问模式。这些特性使得CLAP能够在分析阶段捕捉到代表性的内存访问模式。

Chiplet内存耗尽。一个潜在的极端情况是，当一组页面需要映射到一个物理帧已完全占用的chiplet时。在这种情况下，如果另一个chiplet上有空闲帧，CLAP会将页面映射到那些帧上。虽然这降低了内存局部性，但迁移已经映射的页面既会降低局部性，又会产生额外的开销，如TLB shootdown。这种情况在实践中很少见，因为GPU数据结构通常被各个chiplet上的线程以平衡的方式访问。因此，当映射遵循局部性时——如CLAP中那样——它们自然会导致页面在chiplet间分布良好，这与之前的观察结果一致【108, Characterizing Multi-Chip GPU Data Sharing, 2023, TACO】。

内存超售或分析阶段内存耗尽。另一种情况可能发生在 (1) 所有chiplet的内存空间都被完全利用时，例如在UVM【【3, In-Depth Analyses of Unified Virtual Memory System for GPU Accelerated Computing, 2021, SC】，【11, Unified memory management, 2024, AMD】，【17, SUV: Static Analysis Guided Unified Virtual Memory, 2024, MICRO】，【33, Interplay between Hardware Prefetcher and Page Eviction Policy in CPU-GPU Unified Virtual Memory, 2019, ISCA】，【50, MOST: Memory Oversubscription-Aware Scheduling for Tensor Migration on GPU Unified Storage, 2025, CAL】，【62, Forest: Accessaware GPU UVM Management, 2025, ISCA】，【91, UNIFIED MEMORY ON PASCAL AND VOLTA, 2017, NVIDIA GTC】】中的内存超售情况下，或 (2) 一个数据结构的页面必须映射到特定的chiplet，因为该结构正处于分析阶段。在第一种情况下，CLAP会将大小与当前正在映射的组相匹配的页组迁移到主机内存。内存管理器会从已经完成分析的页面中选择迁移目标，并优先选择最近最少映射到GPU的页面【3, In-Depth Analyses of Unified Virtual Memory System for GPU Accelerated Computing, 2021, SC】。在第二种情况下，CLAP采用相同的迁移机制，但如果其他chiplet上有可用的物理帧，页面可以被迁移到那里。目标chiplet被选为映射页面最少的那个，从而改善chiplet间的内存使用平衡。

与多页TLB的合作。虽然我们的基线假设每种页大小都有独立的TLB（例如，一个64KB TLB和一个2MB TLB），但CLAP也可以与多页TLB设计【【26, Efficient Address Translation for Architectures with Multiple Page Sizes, 2017, ASPLOS】，【84, Prediction-based superpage-friendly TLB designs, 2015, HPCA】，【92, Concurrent Support of Multiple page Sizes on a Skewed Associative TLB, 2004, TC】】协同工作，这些设计在单一结构中存储不同页大小的TLB条目。由于先前的工作已经提出了如何将TLB合并机制【86, CoLT: Coalesced Large-Reach TLBs, 2012, MICRO】应用于多页TLB【26, Efficient Address Translation for Architectures with Multiple Page Sizes, 2017, ASPLOS】，CLAP可以自然地在这些技术的基础上扩展自己的合并机制到此类设计中。

A4 实验环境

  • 硬件配置:

    • GPU: 模拟了一个4-chiplet的MCM GPU。每个chiplet包含64个SM，总共256个SM，频率为1132MHz。每个SM最多支持64个warp。
    • 缓存/TLB: 每个SM配备128KB的L1缓存（20周期延迟）和独立的L1 TLB（32项4KB, 16项64KB, 8项2MB，10周期延迟）。每个chiplet配备4MB的L2缓存（160周期延迟）和chiplet私有的L2 TLB（1024项4KB, 512项64KB, 256项2MB，80周期延迟）。
    • 内存: HBM2内存，每个chiplet有8个内存栈，每个栈8个通道，总共16个通道/chiplet。总带宽为1.8TB/s。
    • 互连: 环形拓扑，GPU内部总带宽为768GB/s，延迟为32ns。

  • 软件配置:

    • 模拟器: GPGPU-Sim v4.2。
    • 基线策略: 采用基于首次接触（First-Touch, FT）的线程块和数据布局策略，并支持按需分页。

  • 数据集与模型:

    • 工作负载: 评估了15个来自多个基准测试套件（如Rodinia, Pannotia, Parboil）的GPU应用，涵盖了规则和不规则的访存模式。
    • ML模型: 包含了从ResNet-50, ViT, GPT-3等模型中提取的GEMM操作，输入数据大小从60MB到4.00GB不等。

  • 对比方案:

    1. 静态分页 (S-64KB & S-2MB): 固定使用64KB或2MB页面进行内存映射。
    2. 理想C-NUMA (Ideal_C-NUMA & Ideal_C-NUMA+inter): 理想化的C-NUMA实现，假设页面迁移相关操作（如TLB shootdown）零延迟。+inter版本支持中间大小页面。
    3. GRIT: 一种为多GPU系统设计的页面迁移策略，已适配MCM GPU架构（去除页面复制），并假设理想的零延迟迁移。
    4. MGvm: 专注于通过优化PTE页面和TLB条目放置来提升MCM GPU地址转换效率的方案。
    5. Barre-Chord (F-Barre): 通过合并跨多个chiplet均匀交错的页面的地址转换来降低开销。
    6. CLAP (本文方案): 利用chiplet-locality为每个数据结构确定合适的页大小，并使用TLB合并技术支持中间大小页面。
    7. Ideal: 一个理想化的配置，使用小页面（64KB）进行细粒度数据放置，但地址转换效率与大页面（2MB）相同。

A4 实验结果

5.1 性能

总体性能。如图18所示，CLAP在性能上优于静态分页方案，与64KB页面相比平均性能提升17.5%，与2MB页面相比平均提升19.2%。CLAP的性能接近理想情况，差距仅为5.78%。

详细分析。

  • 页大小选择: 表4展示了CLAP为每个工作负载中最大的三个数据结构选择的页大小。这些选择与图6中观察到的性能趋势基本一致。例如，对于BFS，静态2MB分页因对第三大数据结构的不当映射而导致远程访问率增加，而CLAP通过为其分配64KB页面缓解了此问题。
  • 边缘情况处理: 对于因数据结构过小（如PAF、SC）或PMM阶段映射不充分（如LUD）而无法进行可靠分析的情况，CLAP回退到机会性大页面（OLP）机制。OLP能动态地形成大页面或维持细粒度映射，从而有效处理这些边缘情况。
  • 与CPU NUMA方案对比: CLAP的性能优于理想化的C-NUMA（提升11.9%）及其变体（提升8.5%）。这是因为C-NUMA采用全局页大小调整策略，无法适应同一应用内不同数据结构的差异化需求，且其对页面迁移的依赖会引入额外开销。
  • 与多GPU方案对比: CLAP比GRIT性能高17.1%。GRIT虽然能通过页面放置保持高数据局部性，但因其页大小固定，无法利用大页面的优势，性能与静态64KB分页方案相近。
  • 与MCM GPU方案对比: CLAP比MGvm性能高24.8%，比F-Barre性能高13.8%。MGvm和F-Barre虽然优化了地址转换，但缺乏对页大小的自适应调整能力，特别是对中间页大小的支持，而CLAP通过构建连续页面组并结合TLB合并，实现了更高的效率和灵活性。
  • 与理想情况对比: CLAP与理想情况的性能差距主要源于那些偏好小页面以保持局部性的工作负载（如STE、LPS）。即使CLAP通过中间页大小缩小了差距，小页面在地址转换效率上仍无法完全媲美2MB大页面。


图 18: 不同配置下的性能（柱状图，以S-64KB为基准归一化）和内存指令的远程访问率（折线）。


表 4: CLAP为每个工作负载中最大的三个数据结构确定的页大小。斜体加粗文本表示使用OLP的分页结果。

5.2 敏感性分析

  • 应用于基于静态分析的方法: CLAP的核心思想与具体的页面映射策略无关。实验将CLAP与基于静态分析（SA）的布局策略结合（CLAP-SA），结果如图19所示。CLAP-SA相比于固定使用64KB和2MB页面的SA策略，性能分别平均提升了18.8%和16.1%，证明了CLAP在不同布局策略下的有效性。进一步地，通过结合运行时分析处理不规则访存模式（CLAP-SA++），性能得到进一步提升，展示了静态分析与运行时分析混合的潜力。


图 19: 基于静态方法[17, 47]的配置性能（柱状图，以SA-64KB为基准归一化）和远程访问率（折线）。CLAP-SA通过合适的页大小选择最大化了静态方法的优势，而CLAP-SA++通过运行时分析扩展了对不规则工作负载的支持。

  • 与页面迁移的协作: 在数据结构使用模式发生变化的动态场景下，页面迁移是有效的。如图20所示，在一个GEMM操作复用部分输出矩阵的场景中，基础CLAP因无法重映射数据而导致高远程访问。而结合了C-NUMA页面迁移的CLAP（CLAP+migration）能有效重映射共享数据，性能提升了34.7%，展示了CLAP与选择性迁移结合的优势。


图 20: 不同配置下的性能（柱状图）和远程访问率（折线）。工作负载（GEMM, 8192 × 768 × 1024）复用了前一个内核的部分数据结构，导致了不同的内存访问模式。

  • 与现有缓存方案的协作: CLAP是正交于现有远程缓存技术的。如图21所示，将CLAP与两种先进的远程缓存方案（NUBA和SAC）结合，可以显著提升它们的效率。CLAP通过在数据到达缓存前就减少了远程访问，使得缓存方案能更好地发挥其潜力（NUBA性能提升从4.8%增至23.9%，SAC从4.3%增至24.1%）。


图 21: 在静态2MB分页和CLAP下各种远程缓存方案的性能。结果以无远程缓存的静态2MB分页为基准进行归一化。

  • Chiplet数量扩展性: 在8-chiplet MCM GPU上的评估显示（图22），CLAP相比于静态64KB和2MB分页方案分别取得了13.3%和21.5%的性能提升。值得注意的是，随着系统规模扩大，CLAP相对于2MB分页的性能优势进一步扩大，这表明不加选择地使用大页面在更大规模系统中的危害更大，从而凸显了CLAP自适应页大小管理的重要性。


图 22: 在8-chip MCM设计中的性能（柱状图，以S-64KB为基准归一化）和远程访问率（折线）比较。

A7 补充细节

6. 相关工作

多GPU系统优化。大量研究提出了通过优化线程和数据放置【【49, CODA: Enabling Colocation of Computation and Data for Multiple GPU Systems, 2018, TACO】，【102, The Locality Descriptor: A Holistic Cross-Layer Abstraction to Express Data Locality In GPUs, 2018, ISCA】】、改进GPU间通信【【30, NetCrafter: Tailoring Network Traffic for Non-Uniform Bandwidth Multi-GPU Systems, 2025, ISCA】，【65, Multi-GPU Memory Management with an Integrated Memory Model, 2021, MICRO】，【66, FinePack: Transparently Improving the Efficiency of Fine-Grained Transfers in Multi-GPU Systems, 2023, HPCA】】、远程缓存机制【【64, Beyond the Socket: NUMA-Aware GPUs, 2017, MICRO】，【106, Combining HW/SW Mechanisms to Improve NUMA Performance of Multi-GPU Systems, 2018, MICRO】】以及页面迁移策略【【20, Griffin: Hardware-Software Support for Efficient Page Migration in Multi-GPU Systems, 2020, HPCA】，【104, GRIT: Enhancing Multi-GPU Performance with Fine-Grained Dynamic Page Placement, 2024, HPCA】，【105, OASIS: Object-Aware Page Management for Multi-GPU Systems, 2025, HPCA】】来缓解多GPU系统中的内存非均匀性问题。本文的研究成果可以补充并增强这些方法在不同类型NUMA GPU系统中的效果。

GPU地址转换优化。先前的工作主要集中于通过扩展TLB覆盖范围【【15, Mosaic: A GPU Memory Manager with Application-Transparent Support for Multiple Page Sizes, 2017, MICRO】，【16, MASK: Redesigning the GPU Memory Hierarchy to Support Multi-Application Concurrency, 2018, ASPLOS】，【41, DUCATI: Highperformance Address Translation by Extending TLB Reach of GPU-accelerated Systems, 2019, TACO】，【52, Increasing GPU Translation Reach by Leveraging Under-Utilized On-Chip Resources, 2021, MICRO】，【56, SnakeByte: A TLB Design with Adaptive and Recursive Page Merging in GPUs, 2023, HPCA】，【58, Orchestrated Scheduling and Partitioning for Improved Address Translation in GPUs, 2023, DAC】，【59, STAR: Sub-Entry Sharing-Aware TLB for Multi-Instance GPU, 2024, MICRO】，【61, Improving Address Translation in Multi-GPUs via Sharing and Spilling Aware TLB Design, 2021, MICRO】】、优化页表遍历【【31, Heliostat: Harnessing Ray Tracing Accelerators for Page Table Walks, 2025, ISCA】，【42, Rethinking Page Table Structure for Fast Address Translation in GPUs: A Fixed-Size Hashed Page Table, 2024, PACT】，【55, Marching Page Walks: Batching and Concurrent Page Table Walks for Enhancing GPU Throughput, 2025, HPCA】，【57, IDYLL: Enhancing Page Translation in Multi-GPUs via Light Weight PTE Invalidations, 2023, MICRO】，【60, Trans-FW: Short Circuiting Page Table Walk in Multi-GPU Systems via Remote Forwarding, 2023, HPCA】，【93, Scheduling Page Table Walks for Irregular GPU Applications, 2018, ISCA】，【94, Neighborhood-Aware Address Translation for Irregular GPU Applications, 2018, MICRO】】和应用推测性转换【85, A Case for Speculative Address Translation with Rapid Validation for GPUs, 2024, MICRO】】来降低GPU中的地址转换开销。这些技术与CLAP是正交的，可以集成以进一步降低地址转换成本。

A5 结论

本文揭示了在MCM GPU中，内存映射的页大小选择扮演着至关重要的角色，因为不合适的页大小会加剧数据在chiplet间的错位，从而放大内存非均匀性问题。为应对这一挑战，我们提出了CLAP，一种能够优化页大小以在不牺牲局部性的前提下利用大页面优势的机制。CLAP利用了GPU数据结构的一种独有属性——chiplet-locality，来实现准确高效的页大小选择。我们的评估结果显示，与先前的分页方案相比，CLAP能够将性能提升高达19.2%。