GPU-Initiated On-Demand High-Throughput Storage Access in the BaM System Architecture

发表时间: 2022-03 · arXiv:2203.04910 (UIUC & NVIDIA, ASPLOS '23)

原文: https://arxiv.org/abs/2203.04910

作者/机构: Zaid Qureshi (NVIDIA/UIUC), Vikram Sharma Mailthody (NVIDIA/UIUC), Isaac Gelado (NVIDIA), Seungwon Min (NVIDIA/UIUC), Amna Masood (AMDU/IUC), Jeongmin Park (UIUC), Jinjun Xiong (University at Buffalo), C. J. Newburn (NVIDIA), Dmitri Vainbrand (NVIDIA), I-Hsin Chung (IBM Research), Michael Garland (NVIDIA), William Dally (NVIDIA/Stanford), Wen-mei Hwu (NVIDIA/UIUC)

速读

一句话结论
本文提出了一种名为 BaM 的系统架构，允许 GPU 线程绕过 CPU 直接向存储设备发起按需、高吞吐的数据访问，在大幅降低硬件成本的同时，实现了比传统 CPU 编排方案高达 5.3 倍的端到端性能提升。

要解决什么问题
图分析、推荐系统等新兴应用的数据集规模（数十 GB 到数十 TB）远超现代 GPU 的显存容量上限（如 A100 仅有 80GB），且这类应用对数据有着细粒度、数据依赖的随机访问模式。原有的应对做法主要卡在两个方向：一是“CPU 中心”路线（如主动分块传输或统一虚拟内存 UVM 缺页中断），这类做法由 CPU 负责编排数据移动，其机制卡点在于 CPU 与 GPU 之间高昂的同步开销、处理缺页中断时 CPU 线程级并行度不足，以及严重的 I/O 流量放大（CPU 无法精准预判 GPU 线程究竟需要哪些细粒度数据，只能整块搬运，导致大量无效传输）；二是“纯 DRAM”路线（如扩容主机内存或多 GPU 显存池化），这种做法虽然规避了软件开销，但将内存扩展至 TB 级的硬件成本极其高昂。因此，如何让 GPU 在不依赖 CPU 的情况下高效、按需地访问海量廉价存储，成为了亟待解决的卡点。

怎么做的
BaM 的核心思路是实现 GPU 发起的按需存储访问，让 GPU 线程直接与 NVMe SSD 交互。它之所以能绕开 CPU 编排的卡点，是因为它利用了 GPU 海量的线程级并行能力来维持极高的 I/O 队列深度，从而掩盖存储设备的访问延迟并打满 PCIe 带宽。根据利特尔法则（Little's Law），系统维持目标吞吐量所需的最小队列深度满足：

$$T \times L = Q_d$$


其中 $T$ 为目标吞吐量，$L$ 为平均访问延迟，$Q_d$ 为所需的并发请求数。为了在 GPU 上高效支撑庞大的 $Q_d$，BaM 的关键设计由以下几个部件构成：
首先是位于 GPU 显存中的高吞吐 I/O 提交与完成队列（SQ/CQ）。传统存储协议在更新队列门铃（Doorbell）寄存器时需要极大的串行临界区，BaM 采用细粒度内存同步机制，通过原子计数器为每个线程分配虚拟队列的票据，让成千上万个线程并行写入请求，最后由单个获胜线程批量敲响门铃，大幅摊销了 PCIe 上的昂贵写入开销。
其次是高度并发的软件缓存。它在应用启动时预分配所有内存以消除系统调用瓶颈，并使用时钟替换算法管理驱逐。为了解决同一个线程束（Warp）内多个线程访问相邻数据的竞争冲突，BaM 利用 __match_any_sync 原语在软件层实现了线程束合并，选出一个领导线程代表整个组去探测缓存和发起 I/O，极大地消除了对底层存储的冗余请求。
最后是提供给开发者的 bam::array<T> 抽象接口。它重载了数组下标操作符，将缓存探测、未命中时的 I/O 请求和数据返回全部封装在底层，使得 GPU 线程可以像访问普通内存一样按需获取存储中的数据。

效果如何
实验在搭载 AMD EPYC CPU、1TB 内存、单张 NVIDIA A100 GPU 以及通过自定义 PCIe 扩展底板连接的多块 NVMe SSD（包括企业级 Intel Optane 和消费级 Samsung 980pro）的硬件上进行。对比基线包括：代表“纯 DRAM”路线的 Target 系统（将数据全部放在主机内存中通过 PCIe 访问），代表“CPU 中心”软件路线的 RAPIDS 框架，以及代表现有存储直通路线的 NVIDIA GDS 和 ActivePointers。在图分析任务（BFS 和 CC，处理高达 1.34 亿节点和 66 亿条边的数据集）中，配备 4 块 Optane SSD 的 BaM 取得了与极其昂贵的“纯 DRAM”基线相当甚至更好的结果（BFS 提速 1.0 倍，CC 提速 1.49 倍），同时将硬件成本降低了 21.8 倍。在纽约出租车数据集（17 亿行）的数据分析查询任务中，BaM 相比 RAPIDS 框架实现了高达 5.3 倍的提速，这主要归功于按需获取列数据带来的 I/O 放大率骤降。作者也承认了该方法的局限性：在诸如 vectorAdd 这种写入密集型任务中，BaM 的表现比主动分块的基线慢了 1.51 倍，因为当前的 BaM 尚未支持异步写回机制，导致应用必须完全承受存储的写入延迟。

1. 主要贡献

图形处理单元（GPU）传统上依赖主机CPU来启动对数据存储的访问。这种方法非常适合具有已知数据访问模式的GPU应用程序，但对于图分析、数据分析、推荐系统或图神经网络（GNN）等需要细粒度、数据依赖型存储访问的新兴应用程序而言，由CPU启动存储访问是不合适的。这是因为CPU-GPU同步开销高、I/O流量放大以及CPU处理延迟长。虽然将整个数据集放入扩展的主机内存或多GPU池化内存（DRAM-only解决方案）可以避免这些问题，但成本极其高昂。

为了解决这一核心问题，本文提出了BaM（Big accelerator Memory）系统架构，其研究目标是为GPU线程提供高效的抽象，使其能够轻松地对存储中的海量数据集进行按需、细粒度的访问，从而实现远超现有解决方案的应用程序性能。本文的创新点包括：
1. 提出了一种以加速器为中心的系统架构BaM，GPU线程无需依赖CPU启动，即可直接对存储或内存中的数组数据进行按需访问。
2. 通过在GPU内存中设计一种新颖的高度并发的提交/完成协议队列库，实现了对存储的按需、高吞吐量、细粒度访问。
3. 提供了一个高吞吐量、可扩展的软件定义缓存和软件API，供程序员利用局部性并控制其应用程序的数据放置。
4. 使用现成的硬件组件构建并评估了BaM原型设计，证明了GPU以经济高效的方式访问海量存储数据的可行性。实验表明，BaM的性能与成本高出$21.7\times$的纯DRAM解决方案相当，并且比最先进的以CPU为中心的软件解决方案快$5.3\times$。

2. 背景知识/关键Observation/设计原则

依赖CPU或多GPU内存的局限性。应用程序可以利用CPU内存或池化多个GPU的内存来托管大型数据结构。GPU提供了足够的内存级并行性来容忍这些内存的访问延迟，其性能显著优于统一虚拟内存（UVM）[48, Unified Memory for CUDA Beginners, 2022]。然而，无论使用何种内存，数据都必须在任何GPU计算开始之前从存储加载到内存中，这种初始数据加载通常是主要的性能瓶颈。此外，随着数据集规模的增加，扩展CPU DRAM或系统中的GPU数量极其昂贵。

容忍存储访问延迟。随着Optane和Z-NAND等存储设备延迟的降低，OS内核软件开销严重限制了存储访问吞吐量，占总延迟的比例高达$36.4\%$。新兴存储系统允许应用程序通过用户级队列对（类似于NVMe的SQ和CQ）直接进行用户级I/O访问 [17, The ZoFS User-Space NVM File System, 2019], [25, SplitFS, 2019], [26, DevFS, 2018], [31, Strata, 2017], [34, Filesystem Semi-Microkernel, 2021], [35, DAOS, 2016], [53, CrossFS, 2020], [67, WekaFS, 2021], [69, Spdk, 2017]。根据利特尔法则（Little's Law）$T \times L = Q_d$（$T$为目标吞吐量，$L$为平均延迟，$Q_d$为队列深度），要充分利用PCIe $\times 16$ Gen4（$\sim 26\mathrm{GBps}$），对于512B和4KB的访问粒度，$T$分别为$51\mathrm{M/sec}$和$6.35\mathrm{M/sec}$。对于Optane SSD（$L=11\mu s$），维持512B访问需要$Q_d=561$（4KB需70）；对于Samsung 980pro（$L=324\mu s$），维持512B访问需要$Q_d=16,524$（4KB需2057）。在基于队列的系统中，入队后更新门铃寄存器需要序列化的临界区，这对于成千上万的GPU线程并发访问会产生巨大的序列化延迟，BaM队列的设计正是为了应对这些挑战。

3. 方法细节

BaM设计目标与挑战。BaM的设计目标是为加速器提供高级抽象，以进行按需、细粒度、高吞吐量的存储访问。为此，BaM在GPU内存中配置了存储I/O队列和缓冲区，并利用GPU的内存映射功能将存储门铃寄存器映射到GPU地址空间。BaM解决了三个挑战：1）利用GPU的大规模并行性来保持大量请求在途以容忍存储延迟；2）在存储带宽低和GPU内存有限的情况下实现资源的最优利用；3）提供高级抽象，隐藏复杂性，使程序员易于集成。
BaM设计的逻辑视图

BaM系统概述。BaM提供了bam::array高级编程抽象，应用程序可调用API将其映射到存储数据。当GPU线程访问数据时，首先使用抽象计算所访问数据的缓存行偏移量（❶）。warp中的线程如果访问相同的缓存行，可以合并其访问（❷）。对于每个唯一的缓存行，由单个线程代表其他线程探测元数据（❸）。如果命中，直接在GPU内存中访问数据。如果未命中，线程进入BaM I/O栈：准备存储I/O请求（❺），将其排入提交队列SQ（❻），然后等待存储控制器在完成队列CQ中发布完成条目（❼）。存储控制器接收门铃更新（A），获取SQ条目（B），处理命令（C），传输数据（D），并在CQ中发布条目（E）。完成后，线程更新缓存状态（❽），更新SQ/CQ状态（❾），最后访问获取的数据。
GPU线程在BaM中的生命周期

与以CPU为中心的方法比较。与传统的CPU主动分块模型相比，BaM具有三个优势。首先，BaM允许GPU线程同时进行计算和从存储获取数据，减少了CPU-GPU同步和GPU内核启动的频率，实现了计算与I/O的细粒度重叠。其次，CPU主动分块由于无法准确预测所需数据，会导致取回大量未使用的字节；而BaM仅在数据被使用时才获取，显著降低了I/O放大开销。最后，BaM通过数组抽象让程序员自然地访问数据，无需手动划分数据集。
传统的以CPU为中心的计算模型与BaM计算模型的比较

高吞吐量I/O队列设计原则。现有的存储I/O协议在入队或清理SQ条目后敲响门铃需要序列化，简单的临界区会导致数千个GPU线程并发时的吞吐量极低。BaM使用细粒度的内存同步，允许许多线程并行执行SQ入队、CQ轮询或标记清理，避免了大型临界区。

SQ元数据结构。BaM在GPU内存中为每个SQ维护以下元数据：1）队列头和尾的本地副本；2）原子票据计数器（ticket counter）；3）turn_counter数组（长度与队列相同）；4）mark位向量（长度与队列相同）；5）一个锁。

入队操作流程。为了将请求入队，线程首先原子地将ticket counter增加2。返回的ticket值是具有$2^{32}$个条目的虚拟队列的索引。将该值除以物理队列大小，余数分配为线程在物理队列中的条目位置，商作为该条目的turn值（轮次）。线程使用其分配的条目索引到turn_counter数组中，并在该位置轮询，直到计数器等于线程的turn值。这有效地为每个物理队列条目创建了按turn值排序的等待线程子队列。

命令拷贝与标记。当轮到该线程时（即分配给该线程的虚拟队列条目变为活动状态），线程将其I/O访问命令复制到物理队列中分配的位置。随后，线程在mark位向量中设置该位置对应的位。turn_counter数组允许多达物理队列大小的线程并行复制其命令。

尾部移动与门铃更新。完成插入后，线程调用move_tail例程。其中一个线程将成功获取锁，将队列尾部移动过所有由调用线程插入的连续新条目，并通过reset_marks例程重置这些条目关联的mark位。然后，该线程使用新的尾部值在存储控制器处敲响门铃，并释放锁。这合并了多个线程在PCIe互连上昂贵的门铃写入操作。如果其他调用线程发现其mark位已被重置，则直接返回；如果未被重置，则继续尝试获取锁。一旦线程知道其位置的mark位已被重置，它会原子地将其位置的turn_counter值加1（变为奇数）。

完成队列(CQ)轮询与出队。命令提交后，线程可以在没有任何锁的情况下轮询CQ以寻找完成条目。找到后，它在CQ的mark位向量中标记该条目，以便CQ头部下一次移动时将其出队。CQ的头部和门铃管理与SQ尾部类似，区别在于：如果线程发现CQ头部已经越过了它标记的条目，它将停止尝试重置其CQ条目的mark位。调用线程中的获胜者用新的头部位置敲响门铃，向存储系统传达进度。

存储控制器进度通信。存储控制器在每个CQ条目中指定一个新的SQ头部来传达进度。持有CQ锁的线程从它重置的最后一个CQ条目中读取此字段，然后从当前的SQ头部迭代到指定的新头部，将每个位置的turn_counter值加1（变为偶数），从而允许等待这些位置的线程将其命令入队。最后，线程更新SQ头部并释放CQ锁。

BaM软件缓存的内存分配。传统的OS内核模式内存管理包含大型临界区，限制了多线程效率。BaM通过在每个应用程序启动时分配软件缓存所需的所有虚拟和物理内存来解决此瓶颈，将临界区缩小到仅在插入或逐出缓存行时需要锁，从而支持更多并发访问。

缓存探测与未命中处理。当线程探测缓存时，它检查缓存行的状态。如果不在缓存中，线程锁定该缓存行，寻找牺牲者进行逐出，并向后备内存请求该缓存行。请求完成后，线程通过将状态设为有效并增加引用计数来解锁该缓存行。这种锁定迫使访问同一缓存行的其他线程等待，从而消除了对后备内存的冗余请求。如果探测时缓存行有效，线程原子地增加其引用计数，使用完毕后再减少引用计数。

时钟替换算法。为避免并发逐出时的争用，BaM缓存使用时钟替换算法 [14, A Paging Experiment With The Multics System, 1968]。缓存有一个全局计数器，线程通过递增它来获取分配的缓存槽。如果分配的槽映射到引用计数非零（即被固定）的缓存行，线程将再次递增计数器，直到找到未固定的缓存槽。然后，线程将该缓存行标记为无效，并将缓存槽的映射更新为新插入的缓存行。

Warp合并机制。Warp中的线程在访问连续字节时可能会产生缓存争用。BaM使用__match_any_sync warp原语在软件中实现合并。线程进行同步，并计算掩码以确定哪些线程正在访问相同的偏移量。线程组决定一个领导者，只有领导者查询缓存并操作状态。组内线程使用__shfl_sync同步，领导者将GPU内存中的地址广播给整个组。与先前工作 [57, ActivePointers, 2016] 中序列化探测不同，BaM的合并器允许warp中所有线程并行分组，每个领导者独立探测缓存，组间无需依赖或同步。

BaM抽象与软件API。BaM提供基于数组的高级API（bam::array<T>）。重载的下标运算符使线程能够合并访问、查询缓存、在未命中时发出I/O请求，并返回适当的元素。程序员只需极少的工作即可调整内核（如Listing 1所示）。此外，BaM初始化需要分配内部数据结构，通常通过默认参数隐式完成，但也可以通过C++模板参数进行专业化配置。

__global__
void kernel(bam::array<float> data, size_t n, 
            bam::array<float> out, bam::array<int> randidx) {
    size_t tid = ...;
    ...
    for(; tid < n; tid += (blockIdx.x * blockDim.x))
        out[tid] = data[randidx[tid]];
};

原型构建。使用现成的NVIDIA GPU和NVMe SSD阵列构建BaM原型。为了使GPU线程能够直接访问NVMe SSD上的数据，必须：1）将NVMe队列和I/O缓冲区移动到GPU内存中；2）使GPU线程能够写入SSD BAR空间中的队列门铃寄存器。

直接NVMe访问机制。BaM创建了一个自定义Linux驱动程序，为系统中的每个NVMe SSD创建一个字符设备。该驱动程序利用GPUDirect RDMA API在GPU内存中固定并映射NVMe队列和I/O缓冲区，使SSD能够执行点对点的数据读写。此外，利用GPUDirect Async [42] 技术将NVMe SSD门铃映射到CUDA地址空间（通过cudaHostRegister API），以便GPU线程按需敲响门铃。

硬件扩展性。标准的4U服务器PCIe插槽有限。BaM原型使用具有自定义PCIe拓扑的PCIe扩展机箱来扩展SSD。扩展机箱有两个抽屉，每个抽屉支持8个$\times 16$ PCIe插槽（1个分配给NVIDIA A100 GPU，其余用于SSD），提供低延迟的点对点访问。

SSD技术权衡。消费级NAND Flash SSD便宜但特性较差；低延迟驱动器（如Intel Optane和Samsung Z-NAND）昂贵但吞吐量和耐用性极佳。无论底层SSD技术如何，BaM原型在每GB成本上比仅使用DRAM的解决方案具有$4.3 - 21.8\times$的优势，且随着设备容量的增加具有高度可扩展性。

BaM原始吞吐量测试。使用Intel Optane SSD和A100 GPU进行微基准测试。队列深度为1024，线程通过指定队列向SSD请求随机的512字节块。结果显示，BaM能够达到每个SSD的峰值IOPs，并随SSD数量线性扩展。使用10个Optane SSD，BaM实现了$45.8\mathrm{M}$的随机读取IOPs和$10.6\mathrm{M}$的随机写入IOPs（这是512B访问的物理极限），相当于$22.9\mathrm{GBps}$（Gen4 $\times 16$ PCIe链路实测峰值带宽的90%）。这验证了BaM的基础设施软件能够匹配底层存储系统的峰值性能。
在英特尔Optane P5800X SSD上的512B随机读（上）和写（下）基准测试扩展

GPUDirect RDMA I/O一致性。先前的工作 [43, 61] 指出，当第三方设备通过PCIe向GPU内存写入数据时，如果没有随后的PCIe读取，并发GPU线程可能无法观察到写入的顺序。在BaM中，允许GPU线程在成功轮询到第一个命令的完成条目后，提交第二个I/O请求（这迫使存储设备执行PCIe读取），但这会产生100%的性能开销。为了优化，BaM实现了一个带锁的共享全局虚拟队列。发现CQ条目的线程竞争该锁，获胜者代表所有可以合并的线程提交一个额外的I/O请求。完成后通知其他线程。该方案将性能开销降低到不到8%。

编程模型与数据共享。BaM为GPU线程提供了与GPU内存相同的内存一致性属性。如果线程读写同一地址，应用程序必须自行实现同步以避免竞争。BaM缓存是写回缓存，提供刷新API。系统崩溃时的状态检查点由应用程序负责。如果应用程序在CPU和GPU内存中同时实例化BaM缓存，必须自行实现同步以保持一致性。

4. 实验环境

5. 实验结果

1. 与NVIDIA GDS的比较：
* 实验内容：使用fio基准测试，将128GB数据从4个SSD传输到GPU内存，测试4KB到1MB的不同I/O块大小。BaM分配与I/O块大小相同的缓存行。
* 结果与分析：GDS仅在32KB的大I/O粒度下才能饱和PCIe链路，在4KB时仅达到PCIe带宽的23.6%，因为受限于Linux软件栈的高开销。相反，BaM在4KB粒度下轻松达到$25\mathrm{GBps}$，饱和了GPU的PCIe链路。
* 引用图表：Fig 5。
BaM与NVIDIA GDS的性能比较

2. 与ActivePointers的比较：
* 实验内容：测试64K和100万个GPU线程。ActivePointers将文件固定在CPU内存的页面缓存中（避免了存储延迟），而BaM将数据保留在4个SSD上。
* 结果与分析：在冷缓存状态下，ActivePointers的未命中处理峰值吞吐量仅为823 KIOPs（4.4 GBps）。而BaM在4KB和8KB缓存行下接近饱和PCIe链路（24 GBps），在512字节下达到17 MIOPs（$20.7\times$的提升）。在热缓存状态下，BaM提供高达430 GBps的带宽，比ActivePointers高出$11.2\times$。
* 引用图表：Fig 6。
BaM与ActivePointers的性能比较

3. 图分析性能优势：
* 实验内容：在不同数据集上运行广度优先搜索（BFS）和连通分量（CC）。基线Target (T)系统将数据存储在主机内存中。BaM配置为8GB缓存，4KB缓存行，使用Intel Optane SSD。
* 结果与分析：使用单个SSD时，BaM由于PCIe $\times 4$吞吐量限制，BFS和CC分别慢$1.43\times$和$1.27\times$。扩展到4个SSD时，BaM的端到端时间分别比Target系统快$1.00\times$（BFS）和$1.49\times$（CC）。这是因为BaM将I/O与计算重叠，而Target系统必须等待文件完全加载到内存后才能计算。
* 性能提升来源分析：使用朴素缓存（无warp合并或引用重用）可使BFS和CC分别获得$11.9\times$和$12.65\times$的加速；启用warp合并和引用重用后，进一步获得$6.07\times$和$11.24\times$的加速。
* SSD类型影响：Samsung DC 1735与Optane性能相近。消费级的Samsung 980pro虽然分别慢$3.21\times$和$2.68\times$，但成本极低。
* 参数敏感性：将缓存缩小到1GB不会降低性能（Fig 10）；队列对数量减少到40以下时性能才开始下降（Fig 11）。
* 引用图表：Fig 7, Fig 8, Fig 9, Fig 10, Fig 11。
BaM的图分析性能
BaM性能提升的来源
不同SSD对BaM性能的影响
BaM缓存容量对性能的影响
队列对数量对性能的影响

4. 数据分析中的I/O放大优势：
* 实验内容：在纽约出租车数据集上执行逐步增加数据依赖列的查询（Q0至Q5）。与RAPIDSv21.12框架（数据预加载在CPU内存中）进行对比。
* 结果与分析：即使使用单个SSD，BaM在Q0上也比基线快$1.22\times$，因为基线受限于CPU查找、移动数据和管理GPU内存的软件开销。随着数据依赖列的增加，RAPIDS必须将整个列传输到GPU，导致I/O放大超过$6\times$。BaM通过按需获取数据，最大限度地减少了I/O放大，使用4个SSD时，BaM比基线快$5.3\times$。
* 引用图表：Fig 12。
BaM和RAPIDS在数据分析查询上的性能

5. VectorAdd工作负载：
* 实验内容：测试写入密集型工作负载。基线使用主动分块，重叠输出写入与下一块的加载。
* 结果与分析：BaM比基线慢$1.51\times$。原因是BaM目前不支持读未命中处理与写回活动的异步重叠，导致整个写入延迟暴露给应用程序。这可以通过在未来启用异步写回来解决。

6. SM资源利用率：
* 实验内容：评估BaM对每个线程寄存器使用的影响。
* 结果与分析：虽然BaM需要更多的寄存器，导致RAPIDS工作负载中出现寄存器溢出，但由于所有被研究的应用程序都受限于存储I/O，寄存器溢出或占用率降低并没有成为性能瓶颈。
* 引用图表：Fig 13。
应用程序中每个线程的寄存器使用情况

6. 补充细节

优化的以CPU为中心的模型。现有工作如SPIN [8]、NVMMU [70]、GAIA [12]和GPUDirect Storage [47] 提出了使用GPUDirect RDMA在SSD和GPU之间启用点对点(P2P)直接内存访问，从而将CPU从数据路径中排除。然而，这些工作仍然需要CPU来编排数据移动。BaM的不同之处在于允许任何GPU线程直接启动对SSD的数据访问。

先前的以加速器为中心的系统。ActivePointers [57]、GPUfs [58]、GPUNet [29] 和 Syscalls for GPU [64] 先前尝试启用以加速器为中心的数据编排模型。然而，这些方法使用并行度较低的CPU来处理来自大规模并行GPU的数据需求，导致整体性能低下。Dragon [37] 将存储访问纳入了UVM缺页异常机制，但同样受限于CPU处理瓶颈。

硬件扩展。先前的工作提出用非易失性内存替换或集成GPU全局内存，或使用FPGA实现存储和加速器之间的直接访问。BaM原型旨在利用新兴的解耦存储硬件组件（现成的NVMe SSD），为具有超大真实数据集的端到端应用程序提供显著的性能和成本优势。

7. 附录细节

主动分块(Proactive Tiling)的局限性。主动分块是一种以CPU为中心的解决方案，要求程序员显式地将数据分解并划分为适合GPU内存的块。CPU代码编排数据移动以主动将块预加载到GPU内存中，并启动计算内核。这种主动访问对于具有动态、数据依赖访问模式的新兴应用程序（如数据分析）存在问题。以纽约出租车数据集的查询为例，CPU初始化行组（查找、分配内存并加载到GPU）和清理行组的代码分别占端到端应用时间的$73\%$和$23\%$以上（Fig 14）。此外，由于CPU无法确定需要哪些数据，它必须获取整个列。例如，当只有$0.03\%$的距离满足条件时，CPU仍获取相关的整个成本列，导致$2.02\times$的I/O放大。随着数据依赖指标数量的增加，CPU中心模型的I/O放大呈线性扩展，超过$6\times$。
使用最先进的RAPIDS系统的GPU加速数据分析应用程序中的执行时间明细和I/O放大

响应式缺页错误(Reactive Page Faults)的局限性。对于图遍历等无法干净划分数据集的应用程序，通常倾向于将整个数据结构保留在GPU的地址空间中。统一虚拟内存（UVM）允许GPU访问驻留在主机内存中的对象，并在数据不在GPU内存时生成缺页异常，由CPU上的UVM驱动程序响应式地进行存储I/O请求。然而，这种方法引入了显著的软件开销。在没有存储访问（数据在主机内存中）的乐观测试中（Fig 15），UVM缺页处理机制实现的平均PCIe数据传输带宽仅为$\sim 14.52\mathrm{GBps}$（实测PCIe Gen4峰值带宽的55.2%）。分析表明，CPU上的UVM故障处理程序利用率达到$100\%$，最大处理速率饱和在$\sim 500\mathrm{K}$ IOPs。这种速率甚至无法充分利用单个消费级SSD的带宽。
运行BFS图遍历时UVM和Zerocopy机制的平均带宽和执行时间

8. 结论

本文提出了BaM系统架构，旨在使GPU能够在没有CPU软件开销的情况下编排对存储的高吞吐量、细粒度访问。BaM允许GPU应用程序计算代码按需以更细的粒度读取或写入数据，从而缓解了I/O放大问题。由于BaM在GPU上支持存储访问控制平面功能（包括缓存、地址转换和协议队列），它避免了昂贵的CPU-GPU同步、OS内核跨越以及限制存储访问吞吐量的软件瓶颈。使用现成硬件组件构建的BaM原型在多个应用程序和数据集上的评估表明，BaM是纯DRAM系统和其他最先进解决方案的强大且具有成本效益的替代方案。未来的工作将包括在BaM系统中启用异步写回功能，以进一步优化写入密集型工作负载的性能。