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

Zaid Qureshi, Vikram Sharma Mailthody (equal contribution), Isaac Gelado, Seungwon Min, Amna Masood, Jeongmin Park, Jinjun Xiong, C. J. Newburn, Dmitri Vainbrand, I-Hsin Chung, Michael Garland, William Dally, Wen-mei Hwu | NVIDIA / UIUC / AMD / University at Buffalo / IBM Research / Stanford | ASPLOS '23, Vancouver, BC, pp. 325–339 | DOI: 10.1145/3575693.3575748

Per-section summary organized by the paper's headings. Each section includes paragraph-level bullet points and exact quantitative results where the paper provides them. "BaM" stands for Big accelerator Memory.

Abstract

• GPUs have traditionally relied on the host CPU to initiate access to data storage; this suits applications with known data-access patterns that can partition their dataset for pipelined GPU processing.
• Emerging applications — graph and data analytics, recommender systems, graph neural networks (GNNs) — require fine-grained, data-dependent access to storage, for which CPU-initiated access is unsuitable due to CPU–GPU synchronization overheads, I/O traffic amplification, and long CPU latencies.
• GPU-initiated storage access removes these overheads from the storage control path and can support such applications at much higher speed, but no system architecture or software stack existed to enable it efficiently — this work fills that gap.
• BaM features a fine-grained software cache that coalesces storage requests while minimizing I/O traffic amplification, communicating with the storage system through high-throughput queues that let the GPU's massive thread count issue I/O at a high rate to fully utilize the devices and the interconnect.
• Results: BaM delivers 1.0× and 1.49× end-to-end speedup for BFS and CC graph analytics while reducing hardware cost by up to 21.7× versus accessing graph data from host memory; and 5.3× speedup on data-analytics workloads over CPU-initiated storage access on the same hardware.

1. Introduction

Background and motivation:

• After a decade of growth in compute throughput and memory bandwidth, GPUs are popular for HPC/ML; emerging data-center workloads (graph/data analytics, GNNs, recommender systems) access massive array datasets ranging from tens of GBs to tens of TBs and growing.
• Storing datasets as in-memory objects lets applications process them naturally, but GPU memory capacity — despite a 53× increase from G80 to A100 — is far smaller than required, so state-of-the-art approaches use CPU user/OS code to partition datasets and orchestrate transfers into GPU memory.
• Another approach uses memory-mapped files plus GPU page faults to invoke the CPU page-fault handler on access to absent data. Both are termed the CPU-centric approach and are bottlenecked by CPU software and CPU/GPU synchronization overheads, yielding poor performance.

The DRAM-only alternative and its cost:

• To avoid CPU-centric inefficiency, some solutions place the entire dataset in host memory (typically 128GB–2TB) or pool multiple GPUs' memory — the DRAM-only solution.
• Extending host memory to tens of TBs is extremely expensive; pooling GPUs is also costly (an A100 has only 80GB, so a 10TB pool needs 125 A100 GPUs). Both require pre-loading the dataset, some of which may go unused under data-dependent access; graph-analytics results show the host-memory solution is comparable to or lower than BaM despite far higher cost.

Proposal and challenges:

• The paper proposes BaM (Big accelerator Memory), which capitalizes on recent improvements in storage latency, throughput, cost, density, and endurance to realize a new level of the accelerator memory hierarchy, providing efficient abstractions for GPU threads to make on-demand, fine-grained access to massive storage datasets.
• Prior GPU-initiated attempts (ActivePointers, GPUfs) achieved low throughput (~823K IOPs on an A100); BaM addresses two key technical challenges.
• First challenge: lack of a fast mechanism for GPU code to generate on-demand storage requests without CPU software bottlenecks (e.g., the OS page-fault handler). BaM provides a scalable, highly concurrent, high-throughput configurable software cache using the GPU's memory bandwidth and atomics to coalesce redundant accesses and exploit reuse.
• Second challenge: CPU-centric approaches suffer from low CPU thread-level parallelism for page-fault handlers/drivers. BaM provides a user-level GPU library of highly concurrent submission/completion queues in GPU memory, removing the page-fault bottleneck via fine-grained synchronization and minimal critical sections.

Positioning in the GPUDirect Async family:

• There is a trend toward GPU autonomy/asynchrony. The GPUDirect Async family accelerates the control path; each transaction involves initiation (work-queue structures created and written) and triggering (transmissions signaled). BaM is the first accelerator-centric approach where GPUs create on-demand data accesses without the CPU initiating or triggering them — a new variant called GPU kernel initiated (KI), GPUDirect Async KI Storage.
• User-level storage device queues raise security concerns in monolithic server architectures, but the shift toward zero-trust security models (guarantees via trusted hardware/software services) provides the framework for securing accelerator-centric user-level storage access like BaM.

Prototype and contributions:

• The authors built a prototype from off-the-shelf hardware plus a novel custom software stack; evaluation shows BaM is on-par with a 21.7× more expensive host-memory DRAM-only solution and up to 5.3× faster than a state-of-the-art CPU-centric software solution.
• Four contributions: (1) propose BaM, an accelerator-centric architecture for GPU-thread on-demand access to array data where it is stored, without the CPU; (2) enable on-demand, high-throughput fine-grained storage access via a library of highly concurrent submission/completion protocol queues; (3) provide high-throughput scalable software-defined caching and a software API to exploit locality and control data placement; (4) construct and evaluate a cost-effective prototype.
• A full version with an appendix (limitations of current systems, extra evaluations) is published; BaM is fully open-source (hardware and software publicly accessible).

2. Background

• The section covers the limitations of keeping large datasets in active CPU memory or pooled multi-GPU memory (§2.1), then why GPUs can tolerate long storage access latency (§2.2).

2.1 Leveraging CPU or Pooled Multi-GPU Memory

• Applications can host large data structures in CPU memory or pooled GPU memory; prior work shows GPUs provide enough memory-level parallelism to tolerate this access latency, significantly out-performing UVM for graph-traversal apps.
• Two pitfalls regardless of CPU or pooled-GPU memory: (1) data must still be loaded from storage to memory before any GPU computation starts, and this initial load can be the main bottleneck (the Target "T" system of Figure 7); (2) hosting the dataset requires scaling memory (more CPU DRAM or more GPUs), prohibitively expensive for massive datasets.

2.2 Tolerating Storage Access Latency

• As device latency drops (Optane, Z-NAND), software overhead becomes a significant fraction of I/O latency. Even with the highly optimized io_uring CPU stack, OS kernel overhead severely limits throughput and reaches up to 36.4% of total storage access latency.
• Emerging storage systems allow direct user-level I/O: storage allocates user-level queue pairs (akin to NVMe submission SQ / completion CQ queues); threads enqueue requests and poll for completion, forgoing the userspace-to-kernel crossing. Isolation and file-system properties are provided by trusted user-level processes, kernel threads, or controller/firmware services.
• Parallelism to tolerate latency and reach full throughput follows Little's Law: T × L = Q_d, where T is target throughput, L is average latency, and Q_d the minimal queue depth needed.
  • For the critical resource — PCIe ×16 Gen4 at 26 GBps — T = 26GBps/512B = 51M/sec and 26GBps/4KB = 6.35M/sec.
  • Measured L: 11 μs (Intel Optane), 324 μs (Samsung 980pro).
  • Q_d to sustain 51M 512B accesses: Optane = 51M × 11μs = 561 requests (70 for 4KB); Samsung 980pro = 51M × 324μs = 16,524 (2057 for 4KB). Q_d can be spread across multiple physical device queues.
• Emerging queue-based storage poses challenges to sustaining T at massive parallelism: after enqueuing NVMe SQ requests, the queue's doorbell must be rung with the updated tail. Doorbell registers are write-only; only one thread may write at a time, and the value must be newer than any previous. Implementing queue insertion as a critical section is simple but serializes — fine for the CPU's limited parallelism but a substantial overhead for thousands of GPU threads. BaM's queue design addresses this.

3. BaM System and Architecture

• BaM's design goal: high-level abstractions for accelerators to make on-demand, fine-grained, high-throughput storage access while enhancing storage performance. It provisions storage I/O queues and buffers in GPU memory (Figure 1) and maps storage doorbell registers into the GPU address space. Three key challenges:
  1. Storage protocols/devices have significant latency, so BaM must use the GPU's massive parallelism (up to three orders of magnitude more than the CPU) to keep many requests in flight and overlap with compute (§3.3).
  2. Storage devices have low bandwidth and GPUs have limited memory capacity, so BaM must use these resources optimally (§3.4).
  3. GPU kernels don't expect to make storage accesses, so BaM must hide complexity behind high-level abstractions for easy integration (§3.5).

3.1 BaM System Overview

• BaM presents the bam::array high-level programming abstraction; applications call BaM APIs to map a bam::array to data on storage, akin to mmap'ing a file.
• Life of a GPU thread (Figure 2, markers ❶–❾): a thread accesses data with bam::array ❶, determines the offset / cache line ❷; threads in a warp coalesce their accesses ❸ when they touch the same cache line; for each unique cache line a single thread probes the line's metadata ❹ on behalf of the group.
• On a hit, the thread directly accesses GPU memory; on a miss, it fetches from backing memory. The software cache optimizes backing-memory bandwidth two ways: (1) eliminating redundant requests, (2) letting users configure the cache for application needs.
• If a storage device backs the data, the thread enters the BaM I/O stack to prepare a storage I/O request ❺, enqueues it to a submission queue ❻, and waits for the controller to post a completion entry ❼. The I/O stack amortizes software overhead via GPU thread-level parallelism, batching SQ/CQ entries to minimize expensive doorbell updates and shrink critical sections.
• On receiving the doorbell update (A), the storage controller fetches SQ entries (B), processes the command (C), transfers data between SSD and GPU memory (D), and posts a completion entry to the CQ head (E). The thread then updates the cache state ❽, updates SQ/CQ state, and accesses the fetched data in GPU memory ❾.

3.2 Comparison With the CPU-Centric Approaches

• First advantage over proactive tiling (Figure 3a vs 3b): with proactive tiling the CPU copies data and launches kernels repeatedly to cover a large dataset; each launch/termination incurs costly CPU–GPU synchronization. BaM lets GPU threads both compute and fetch, reducing synchronization and kernel launches, and overlapping some threads' storage latency with other threads' compute.
• Second advantage: in proactive tiling the CPU cannot accurately determine what data is needed and when, so it fetches many unneeded bytes; with BaM a GPU thread fetches bytes only when used, reducing I/O amplification.
• Third advantage: proactive tiling makes programmers partition data and hand-overlap compute with transfers; BaM lets the programmer naturally access data through the array abstraction and use GPU thread parallelism to hide latency.

3.3 High-Throughput I/O Queues

• BaM uses the GPU's massive parallelism and fast hardware scheduling to maintain the queue depths needed to hide latency. Because ringing doorbells and cleaning SQ entries require serialization, a single critical section over enqueue + doorbell would yield poor throughput; instead BaM uses fine-grained memory synchronization so many threads can enqueue, poll the CQ, or mark entries for cleanup in parallel without large critical sections.
• Per-SQ metadata kept in GPU memory: (1) local copies of head and tail, (2) an atomic ticket counter, (3) a turn_counter array (integer array as long as the queue), (4) a mark bit-vector (same length), (5) a lock.
• To enqueue, threads atomically increment the ticket counter by two; the returned ticket indexes a virtual queue of 2³² entries. Dividing by physical queue size yields an entry (remainder) and a turn with a valid bit (quotient). Each thread indexes the turn_counter by its entry and polls until the counter equals its turn — effectively creating sub-queues whose waiting threads are ordered by turn per physical entry.
• When it is the thread's turn (its virtual entry becomes active), it copies its I/O command into its assigned physical entry and sets that position's bit in the mark bit-vector; the turn_counter array lets as many threads as the queue size copy commands in parallel.
• Threads call move_tail to finish insertion: one thread acquires the lock, advances the tail past consecutive new entries, resets their mark bits with reset_marks, rings the doorbell once with the new tail (coalescing doorbell writes, which are expensive over PCIe), and releases the lock. Other calling threads return once their mark bits are reset; if not reset, they keep trying to take the lock. Once a thread's mark bit is reset, it atomically increments its position's turn_counter by one (to an odd value).
• After submission, a thread polls the CQ without any lock to find its completion entry, then marks the entry for dequeuing in the CQ's mark bit-vector. The CQ head and doorbell are managed like the SQ tail, except a thread stops resetting its CQ mark bit if it sees the head already moved past it; a winning thread rings the doorbell with the new head.
• The storage controller communicates progress by specifying a new SQ head in each CQ entry; the CQ-lock holder reads this from the last CQ entry it reset, iterates from the current SQ head to the new head incrementing each turn_counter by one (to an even value) — freeing waiting enqueuers — then updates the SQ head and releases the CQ lock.

3.4 BaM Software Cache

• The software cache enables optimal use of limited GPU memory and storage bandwidth. Traditional OS-kernel memory management uses large critical sections that limit multi-threaded effectiveness; BaM pre-allocates all virtual and physical memory for the cache at startup, reducing critical sections to a lock only on insert/evict and allowing many more concurrent accesses.
• To minimize redundant I/O: a thread probing the cache by offset checks the line's state. On a miss it locks the line, finds a victim, and requests the line from backing memory; on completion it validates the line and increments its reference count. Locking forces other threads on the same line to wait until insertion, preventing duplicate backing-memory requests (exploiting locality). On a valid hit, the thread atomically increments the reference count and decrements it when done.
• To avoid contention among concurrent evictions, the cache uses a clock replacement algorithm: a global counter is incremented to assign each thread a cache slot, letting threads evict in parallel. If the assigned slot is pinned (non-zero reference count), the thread increments the counter again until it finds an unpinned slot, marks that line invalid, and remaps the slot to the new line.
• Warp coalescing: threads in a warp may contend on the cache, especially when consecutive threads access contiguous bytes; this is costly even when the line is already resident. BaM coalesces in software with the __match_any_sync warp primitive to compute a mask grouping threads by offset; each group elects a leader who alone queries the cache and manipulates the line's state, broadcasting the offset address via __shfl_sync.
• Compared to prior work (ActivePointers) where unique probes within a warp are serialized over many instructions, BaM's coalescer groups all warp threads by accessed cache line with a single __match_any_sync instance; per-group leaders independently probe their line with no inter-group dependencies or synchronization.

3.5 BaM Abstraction and Software APIs

• BaM's stack provides an array-based high-level API, bam::array, consistent with array interfaces in C++/Python/Rust. Its overloaded subscript operator coalesces accesses, queries the cache, issues I/O on misses, and returns the element of type T (Listing 1). bam::array<T> can build higher-level abstractions (e.g., cache-line reference reuse so threads don't re-probe the same line). By contrast, the proactive-tiling CPU-centric model requires non-trivial full rewrites to decompose compute and transfers into GPU-memory-sized tiles.
• BaM initialization allocates a few internal data structures reused over the application's lifetime; it can happen implicitly via library construction if no customization is needed, otherwise the application specializes memory through C++ template parameters. In most cases specialization/fine-tuning is unnecessary — §5 uses only default parameters.

4. BaM Prototype

• Uses off-the-shelf hardware (NVIDIA GPUs, arrays of NVMe SSDs) to show that GPUs directly accessing storage can deliver enough random-access bandwidth to fully utilize a GPU's PCIe Gen4 ×16 link — at which point a storage-based solution is as good as host memory over PCIe but much cheaper. The prototype assumes bare-metal direct NVMe access.

4.1 Enable Direct NVMe Access From GPU Threads

• Using NVMe SSD controllers as the example, two things are needed: (1) move NVMe queues and I/O buffers from host CPU memory to GPU memory, and (2) let GPU threads write the queue doorbell registers in the SSD's BAR space. A custom Linux driver creates a character device per NVMe SSD (like SmartIO's); applications use BaM APIs to open the character device per SSD.
• In the driver, BaM leverages GPUDirect RDMA to pin and map NVMe queues and I/O buffers into GPU memory, enabling the SSD to do peer-to-peer reads/writes to GPU memory.
• BaM leverages GPUDirect Async to map NVMe doorbells into the CUDA address space so GPU threads can ring them on demand; this requires memory-mapping the SSD's BAR space into the application address space, then mapping it to CUDA via the cudaHostRegister API. Other storage systems can be enabled similarly.

4.2 Scalable Hardware

• Scaling BaM with the PCIe slots of a data-center 4U server is challenging: a Supermicro AS-4124 has five PCIe Gen4 ×16 slots per socket; a GPU in one slot can reach only four ×16 PCIe devices without crossing the inter-socket fabric (which degrades performance). No single NVMe SSD matches a PCIe ×16 Gen4 link, so BaM must scale the number of NVMe devices to match that bandwidth.
• The custom prototype (Table 1) uses a PCIe expansion chassis with a custom topology; PCIe switches give low-latency, high-throughput peer-to-peer access. The chassis has two identical drawers connectable independently to the host. One ×16 slot per drawer holds an NVIDIA A100; the rest hold SSDs. Each drawer has 8 ×16 PCIe slots but supports only 10 U.2 SSDs (size constraints); with PCIe bifurcation a multi-SSD riser can enable >16 M.2 SSDs per drawer.
• SSD technology trade-offs (Table 2): $/GB is based on current list price per device plus chassis and risers. Consumer-grade NAND Flash SSDs are inexpensive with harder characteristics; low-latency drives (Intel Optane, Samsung Z-NAND) are pricier with desirable characteristics. For write-intensive apps, Intel Optane gives the best throughput and endurance.
• Regardless of SSD technology, the prototype gives a 4.3–21.8× cost-per-GB advantage over a DRAM-only solution even with chassis and risers; this advantage grows with capacity added per device, so BaM scales well as SSD capacity and data size increase.

4.3 BaM Raw Throughput

• To establish that BaM can saturate the storage system, raw throughput is measured (Intel Optane SSDs + A100). All available SQ/CQ queues are allocated in GPU memory at queue depth 1024; a CUDA kernel has each thread request a random 512-byte block via a designated queue, distributed round-robin across queues (and across SSDs for the multi-SSD case). IOPS = requests submitted / kernel execution time.
• Results (Figure 4, 512B random read/write): BaM reaches peak IOPs per SSD and scales linearly with added SSDs for reads and writes. A single Optane needs only ~16K–64K GPU threads to near-peak IOPs. With ten Optane SSDs, BaM achieves 45.8M random read IOPs and 10.6M random write IOPs, i.e., 22.9 GBps read (90% of measured peak Gen4 ×16 bandwidth) and 5.3 GBps write. Write bandwidth can grow further with more SSDs (not yet PCIe-limited). Similar scaling holds for Samsung SSDs and at 4KB. These confirm BaM's software can match the underlying storage's peak.

4.4 Discussion

• GPUDirect RDMA I/O consistency: when a third-party device writes GPU memory via GPUDirect RDMA over PCIe, without a following PCIe read the write order may not be preserved from concurrent GPU threads' view. In BaM, submitting a second I/O after polling the first's completion guarantees a PCIe read completed after all the first command's writes (since the storage must read the submitted request before writing the completion) — but this incurs 100% performance overhead.
• To reduce extra I/O, BaM uses a shared global virtual queue with a lock: threads that find their CQ entries race for the lock; the winner submits one additional I/O on behalf of all coalescable competing threads, notifies them on completion, and releases the lock. Others poll to see if their second request was covered. This incurs less than 8% performance overhead.
• Programming model: BaM gives GPU threads a memory-like abstraction with the same memory consistency properties as GPU memory; if threads read/write the same words in BaM-backed memory, the application must implement synchronization to avoid races, as with GPU memory.
• BaM supports writes at all stack levels: block write I/O requests, dirty-line tracking, and a high-level write method that sets the dirty bit. The BaM cache is write-back with APIs to flush a specific or all dirty lines.
• If the system crashes mid-kernel during writes, there are no guarantees unless the application does its own check-pointing (common in accelerated applications).
• CPU–GPU data sharing: if the application instantiates one BaM cache in GPU memory and another in host memory, it must implement synchronization to keep a consistent shared view, just as when using GPU memory without BaM.

5. Evaluation

• BaM excels on data-dependent access patterns (graph analytics, RAPIDS frameworks). The evaluation shows: (a) BaM is on-par with or outperforms state-of-the-art solutions (§5.2, §5.3); (b) BaM's design is agnostic to SSD medium, enabling cost-effective application-specific solutions; (c) BaM reduces I/O amplification and CPU orchestration overhead for data-analytics workloads (§5.3).

5.1 Comparison With GDS and ActivePointers

• BaM is compared against NVIDIA GDS (GPUDirect Storage) for I/O block sizes 4KB–1MB. For GDS, fio benchmarks sequential reads of 128GB from four SSDs to GPU memory with 16 CPU threads. For BaM, each warp is assigned a cache-line of the same size as the GDS I/O block, with consecutive warps accessing consecutive cache-lines; the cache-line size defines BaM's I/O granularity.
• Figure 5: GDS saturates the GPU PCIe link only at a large 32KB granularity and reaches only 23.6% of PCIe bandwidth at 4KB, limited by the Linux software stack regardless of CPU thread count. BaM easily reaches 25 GBps with four SSDs — the measured peak of the GPU's PCIe link.
• BaM is compared against ActivePointers. A warp reads 1024 contiguous 8-byte elements in a file with coalesced accesses. For ActivePointers the file is pinned in the Linux page cache (favoring it, since misses only need CPU→GPU transfers, avoiding storage latency); for BaM, data lives on four SSDs and misses incur storage latency. Results are shown for 64K and 1 Million threads (Figure 6); up to 32 Million threads show similar trends.
• Figure 6: with a cold cache, ActivePointers' I/O is hindered by the GPUfs mechanism, achieving only 4.4 Gbps effective bandwidth for 8KB transfers and a peak miss-handling throughput of 823 KIOPs (512-byte lines). BaM nearly saturates the PCIe link at ~24 Gbps with 4KB and 8KB lines, and even with 512-byte lines reaches 85% of the 4-SSD peak at 17 MIOPS. With both caches hitting, BaM delivers up to 430 Gbps effective bandwidth — 11.2× ActivePointers' peak. BaM beats ActivePointers by more than an order of magnitude in both miss handling and hit delivery.

5.2 Performance Benefit for Graph Analytics

• Graphs are listed in Table 3; K, U, F, M are the four largest graphs from the SuiteSparse Matrix collection, and Uk is from LAW.
• The goal is competitive performance against host-memory DRAM-only graph analytics. The optimistic target baseline T lets GPU threads directly do coalesced fine-grain access to graph data in host memory, with enough CPU memory for the inputs, enabling a direct comparison.
• Two algorithms — Breadth-First Search (BFS) and Connected Components (CC) — run on T and on BaM with different SSDs. Porting state-of-the-art GPU implementations to BaM needs minimal code changes (§3.5). BFS reports average runtime over at least 32 source nodes with >2 neighbors; CC is not run on Uk (CC needs undirected graphs). Default config: 8GB cache, 4KB cache-line, four Intel Optane SSDs, 128 queue-pairs at depth 1024.
• Overall performance with Intel SSD (Figure 7): with a single Optane SSD (B_1I), BaM is on average slower by 1.43× (BFS) and 1.27× (CC), due to the limited ×4 PCIe Gen4 throughput of one SSD.
• Scaling to four SSDs (B_4I) with replicated data raises aggregate bandwidth to ×16 Gen4-like levels; versus T (including its file-loading time), BaM gives on average 1.00× (BFS) and 1.49× (CC) speedup.
• In both workloads BaM overlaps some threads' SSD transfers with others' compute, fully using PCIe Gen4 ×16 while incurring much less I/O amplification; T must wait for the whole file to load before offloading compute, so its superior host-memory bandwidth cannot overcome the initial loading latency — BaM is on-par or higher end-to-end.
• Versus single-SSD BaM, the four-SSD config scales 3.48× (BFS) and 4× (CC). The main BFS scaling detractor is Uk: many nodes have tiny neighborhoods, giving deep traversal (>100 iterations) with small frontiers, so the total overlapping I/O per iteration is insufficient to tolerate latency.
• BaM performance breakdown (Figure 7): loading the dataset into GPU HBM and measuring gives Compute (green); routing every access through the BaM cache API while all data is in HBM (no I/O) gives the best-case cache cost, and subtracting Compute yields cache API overhead (orange); constraining the cache to 8GB and subtracting Compute and cache time gives storage I/O time (grey).
• Observations: cache overhead is about 2–15% with one SSD and 4–45% with four SSDs (single-SSD perf is storage-bound; more SSDs alleviate this). Remaining cache overhead comes from metadata contention, long-latency atomics, and warp scheduling among polling threads. With four SSDs BaM is still storage-I/O bound at 5–6.2M IOPs (>80% of peak) — an upper limit beyond which adding SSDs won't help without application changes to trigger cache misses earlier.
• Sources of improvement (Figure 8): a naive cache (no warp coalescing or reference reuse) gives 11.9× (BFS) and 12.65× (CC) over a no-cache implementation (the cache minimizes I/O amplification); adding warp coalescing and reference reuse gives a further 6.07× (BFS) and 11.24× (CC).
• Impact of SSD type (Figure 9): Samsung DC PM1735 (S1735) ≈ Intel Optane for almost all workloads (similar peak 4KB read IOPs). With Samsung 980pro, BaM is on average 3.21× (BFS) and 2.68× (CC) slower than Optane — encouraging, since consumer-grade SSDs are by far the cheapest.
• Sensitivity — cache size (Figure 10, K dataset): even a 1GB cache shows no degradation versus 8GB (same spatial/temporal locality captured); 32GB and 64GB let BaM keep the entire working set, incurring only cold misses.
• Number of queue pairs (Figure 11): performance holds as queue pairs decrease, degrading only at 40 or fewer pairs, where queue contention and per-queue NVMe serialization begin to matter.

5.3 I/O Amplification Benefit for Data Analytics

• Enterprise data-analytics workloads illustrate reduced I/O amplification and software overhead on large structured datasets; though a small subset of GPU applications today, the market was \(205 billion in 2020** and **\)230 billion in 2021.
• Setup: the NYC taxi-ride dataset and six queries compare BaM against RAPIDS v21.12. The queries answer "Q5: average dollar/mile for trips ≥ 30 miles?", starting with a distance-column scan (Q0) and adding total cost (Q1), surcharges (Q2), hail fee (Q3), tolls (Q4), and taxes (Q5) for trips ≥ 30 miles. For RAPIDS the whole file is pinned in the Linux page cache (so RAPIDS reads from CPU DRAM with no storage I/O — best-case for modern systems). BaM replicates data across up to four Intel Optane P5800X SSDs with 4KB cache-lines and 8GB cache.
• Results (Figure 12): even with a single SSD, BaM outperforms RAPIDS on all queries. For Q0 BaM is up to 1.22× faster despite no I/O-amplification benefit, because RAPIDS suffers CPU software overhead to find/move data and manage GPU memory even with the data preloaded.
• With each added data-dependent metric (Q1→Q5) BaM's advantage grows, from its reduced I/O amplification under on-demand fetching while RAPIDS must transfer entire columns; the baseline's I/O amplification increases with more metrics, whereas BaM handles multiple data-dependent columns nearly as efficiently as one.
• BaM's end-to-end time scales 1.46× with two SSDs and 1.62× with four SSDs over single-SSD BaM; the sub-linear scaling is from setup overhead (pinning/mapping queues and I/O buffers for DMA) growing with SSD count. With four SSDs BaM reaches up to 5.3× speedup over the baseline.

5.4 VectorAdd Workload

• vectorAdd is a write-intensive workload: two input arrays of four billion 8-byte elements each produce one four-billion-element output, with inputs in storage and output written to storage. The baseline uses proactive tiling, splitting into five tiles to fully overlap the current tile's output write with loading the next tile's inputs; BaM works on the whole dataset (no tiling), each warp assigned to an output cache-line.
• BaM is 1.51× slower than the tiling baseline because it does not yet overlap read-miss handling with write-back, exposing the full write latency; enabling asynchronous write-back is left as future work.

5.5 SM Resource Utilization

• Figure 13 shows per-thread register usage with and without BaM. Register spilling is observed in RAPIDS. Neither BaM nor the apps use shared memory and no occupancy constraints are enforced. Register counts (w/o BaM → BaM): BFS 28→68; CC 36→66; RAPIDS Q0 12→90; RAPIDS Q5 21→255; VecAdd 32→112. Despite higher register usage, all studied applications are storage-I/O bound, so register spilling/reduced occupancy is not a bottleneck.

6. Related Work

6.1 Optimized CPU-Centric Model

• Most GPU programming models assume the working dataset fits in GPU memory; otherwise application-specific techniques process large data on GPUs.
• SPIN and NVMMU enable peer-to-peer DMA from SSD to GPU via GPUDirect RDMA, excluding the CPU from the data path; SPIN integrates P2P into the OS file stack with page cache and read-ahead for sequential reads. GAIA extends SPIN's page cache to GPU memory. Gullfoss gives a high-level interface to use GPUDirect APIs efficiently. Hippogriffdb adds P2P transfer to an OLAP database. GPUDirect Storage removes the CPU from the SSD↔︎GPU data path via GPUDirect RDMA; AMD RADEON-SSG has similar efforts. All still require the CPU to orchestrate data movement, whereas BaM lets any GPU thread initiate SSD access.

6.2 Prior Accelerator-Centric Systems

• ActivePointers, GPUfs, GPUNet, and Syscalls for GPU previously attempted accelerator-centric data orchestration. GPUfs and Syscalls for GPU first let GPUs request file data from the host CPU; ActivePointers added a memory-like abstraction over GPUfs. Dragon folded storage access into the UVM page-faulting mechanism. All use less-parallel CPUs to serve a massively parallel GPU, giving poor performance (§5 and the full version).
• The authors acknowledge prior work moving network control to the GPU, but treat storage as a new set of challenges; with BaM, GPUs efficiently and directly access storage with demonstrated performance and cost benefits.

6.3 Hardware Extensions

• Prior work proposed replacing or integrating GPU global memory with non-volatile memories. DCS enabled direct access among storage, network, and accelerators with an FPGA providing translation for coarse-grain transfers; enabling persistence within the GPU has also been proposed. These validate the need for large memory capacity, but BaM uses emerging disaggregated storage components to provide end-to-end performance on very large real-world datasets.

7. Conclusion

• The paper argues for enabling GPUs to orchestrate high-throughput, fine-grained storage access without GPU software overhead, in the new BaM architecture. BaM mitigates I/O amplification by letting GPU compute code read/write at finer granularities on demand, and by supporting caching, translation, and protocol queues on the GPU it avoids costly CPU–GPU synchronization, OS kernel crossings, and software bottlenecks. Using off-the-shelf hardware, the prototype shows on multiple applications/datasets that BaM is a viable, superior alternative to DRAM-only and other state-of-the-art solutions.

8. System Architecture (synthesized)

        GPU                                         Storage (per SSD)
+---------------------------+    PCIe Gen4    +-----------------------+
|  GPU Thread               |  <----------->  |  Doorbell Reg (DB)    |
|   ❶ val = bam::array[tid] |   control path  |  DMA engine           |
|   ❷ offset / cache line   |   (doorbell)    |  Controller (CTRL)    |
|   ❸ warp coalescer        |                 |  Media (NAND/Optane)  |
|      __match_any_sync     |   data path     +-----------------------+
|      __shfl_sync (leader) |  <===========>  (peer-to-peer DMA)
|   ❹ probe cache metadata  |
+------------+--------------+
   HIT |     | MISS
       v     v
+---------------------------+   In GPU memory (GPUDirect RDMA pinned):
|  BaM Software Cache       |     - Queue metadata (head/tail, ticket,
|   write-back, clock evict |       turn_counter[], mark bit-vec, lock)
|   per-line: state, lock,  |     - SQ/CQ queue pairs
|   ref count, dirty bit    |     - Cache metadata
+------------+--------------+     - I/O buffers (DMA target)
       | MISS -> I/O Stack
       v
+---------------------------+   Enqueue protocol (lock-minimized):
|  Submission Queue (SQ)    |     ticket += 2 -> (entry, turn)
|   ❺ build NVMe cmd        |     poll turn_counter[entry]==turn
|   ❻ enqueue + set mark    |     write cmd, set mark bit
|   (A) move_tail: 1 thread |     move_tail -> reset_marks -> ring DB
|       rings doorbell once |
+---------------------------+
|  Completion Queue (CQ)    |   (B) ctrl fetches SQ  (C) processes
|   ❼ lock-free poll        |   (D) DMA to I/O buffer (E) post CQ entry
|   ❽ update cache + SQ/CQ  |   ctrl returns new SQ head via CQ entry
|   ❾ read data in GPU mem  |
+---------------------------+

9. Cross-Cutting Empirical Take-Aways

10. Limitations of BaM (stated by authors)

• No overlap of read-miss handling with write-back; full write latency is exposed, making write-intensive vectorAdd 1.51× slower than the tiling baseline (§5.4).
• GPUDirect RDMA write-ordering needs either a dependent I/O request (100% overhead) or a coalesced shared-queue scheme (<8% overhead) (§4.4).
• Cache overhead from metadata contention, long-latency atomics, and warp scheduling among polling threads (4–45% with four SSDs) (§5.2).
• Four-SSD BaM is bounded by I/O request generation rate (5–6.2M IOPs); beyond ~80% of peak, adding SSDs doesn't help without application changes to trigger misses earlier (§5.2).
• BFS scaling is limited on Uk due to deep traversal, small frontiers, and insufficient concurrent I/O (§5.2).
• No crash-consistency guarantees unless the application self-checkpoints (§4.4).
• CPU–GPU data sharing across two BaM caches needs application-managed synchronization (§4.4).
• BaM raises per-thread register usage (up to 255 for RAPIDS Q5), causing register spilling in RAPIDS, though not performance-limiting since workloads are I/O bound (§5.5).
• The prototype is constrained by server PCIe slot count and needs a custom expansion chassis to match PCIe ×16 bandwidth; the provided artifact system is limited to ~2 SSDs (§4.2).

11. Future Work / Open Problems

• Enable asynchronous write-back to overlap read-miss handling with write-back and hide write latency (§5.4).
• Application-level changes to work assignment/scheduling to trigger BaM cache misses earlier and improve I/O-request generation rate and SSD scaling (§5.2).
• Use of emerging disaggregated storage hardware components for very large real-world datasets (§6.3).
• Broader adoption of zero-trust / trusted-service security models to secure user-level storage queues at data-center scale (§1).

12. Relation to NCCL / GPU Collective Communication / HPC

• NCCL and GPU collective communication: not discussed. The paper does not mention collective libraries.
• Multi-GPU appears only as a baseline argued against — pooling many GPUs' memory (DRAM-only) is dismissed as prohibitively expensive (125 A100s for 10TB). BaM itself is single-GPU per drawer and performs no GPU-to-GPU communication.
• HPC is mentioned only as a general GPU application domain in the Introduction; there is no HPC collective/communication evaluation.
• The mechanisms most conceptually adjacent to GPU-initiated communication are BaM's lock-minimized GPU-resident SQ/CQ protocol (ticket / turn / mark / move_tail) and GPUDirect Async KI doorbell ringing from GPU threads — but these are storage control paths; prior work on moving network control to the GPU is acknowledged (§6.2) yet storage is treated as a distinct problem.